/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/ck.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/ck.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/ck.hpp Source File
ck.hpp
Go to the documentation of this file.
1 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
2 // SPDX-License-Identifier: MIT
3 
4 #pragma once
5 
6 #include "ck/config.h"
7 #include <stdint.h>
8 
9 #if !defined(__HIPCC_RTC__) || !defined(CK_CODE_GEN_RTC)
10 #ifndef CK_DONT_USE_HIP_RUNTIME_HEADERS
11 #include "hip/hip_runtime.h"
12 #include "hip/hip_fp16.h"
13 #endif
14 #endif
15 // to do: add various levels of logging with CK_LOG_LEVEL
16 
17 #ifndef CK_TIME_KERNEL
18 #define CK_TIME_KERNEL 1
19 #endif
20 
21 // constant address space for kernel parameter
22 // https://llvm.org/docs/AMDGPUUsage.html#address-spaces
23 #define CK_CONSTANT_ADDRESS_SPACE __attribute__((address_space(4)))
24 
25 // launch bounds
26 #define CK_USE_LAUNCH_BOUNDS 1
27 
28 #ifdef CK_USE_LAUNCH_BOUNDS
29 // for most kernels
30 #define CK_MAX_THREAD_PER_BLOCK 256
31 #define CK_MIN_BLOCK_PER_CU 2
32 
33 // for wavelet GEMM kernel
34 #define CK_WAVELET_MAX_THREAD_PER_BLOCK 512
35 #define CK_WAVELET_MIN_BLOCK_PER_CU 2
36 #endif
37 
38 // kernel attribute: amdgpu_waves_per_eu()
39 #ifdef CK_USE_WAVES_PER_EU
40 // for 1-wave kernels, control arguments of amdgpu_waves_per_eu() attribute
41 #ifndef CK_MIN_WAVES_PER_EU
42 #define CK_MIN_WAVES_PER_EU 0
43 #endif
44 
45 #ifndef CK_MAX_WAVES_PER_EU
46 #define CK_MAX_WAVES_PER_EU 0
47 #endif
48 
49 #else
50 #define CK_USE_WAVES_PER_EU 0
51 #endif
52 
53 // define general macros for various architectures
54 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__) || \
55  defined(__gfx9_4_generic__)
56 #define __gfx9__
57 #endif
58 #if defined(__gfx942__) || defined(__gfx950__) || defined(__gfx9_4_generic__)
59 #define __gfx94__
60 #endif
61 #if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) || \
62  defined(__gfx1013__) || defined(__gfx10_1_generic__)
63 #define __gfx101__
64 #endif
65 #if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \
66  defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \
67  defined(__gfx10_3_generic__)
68 #define __gfx103__
69 #endif
70 #if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \
71  defined(__gfx1103__) || defined(__gfx1150__) || defined(__gfx1151__) || \
72  defined(__gfx1152__) || defined(__gfx1153__) || defined(__gfx11_generic__)
73 #define __gfx11__
74 #endif
75 #if defined(__gfx1200__) || defined(__gfx1201__) || defined(__gfx12_generic__)
76 #define __gfx12__
77 #endif
78 
79 // buffer resource
80 #ifndef __HIP_DEVICE_COMPILE__ // for host code
81 #define CK_BUFFER_RESOURCE_3RD_DWORD -1
82 #elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx9__)
83 #define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000
84 #elif defined(__gfx101__) || defined(__gfx103__)
85 #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000
86 #elif defined(__gfx11__) || defined(__gfx12__)
87 #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31004000
88 #endif
89 
90 // FMA instruction
91 #ifndef __HIP_DEVICE_COMPILE__ // for host code, define nothing
92 #elif defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__) || defined(__gfx1011__) || \
93  defined(__gfx1012__) // for GPU code
94 #define CK_USE_AMD_V_FMAC_F32
95 #define CK_USE_AMD_V_DOT2_F32_F16
96 #define CK_USE_AMD_V_DOT4_I32_I8
97 #elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx101__)
98 #define CK_USE_AMD_V_MAC_F32
99 #elif defined(__gfx11__) || defined(__gfx12__)
100 #define CK_USE_AMD_V_FMAC_F32
101 #define CK_USE_AMD_V_DOT2_F32_F16
102 #define CK_USE_AMD_V_DOT4_I32_I8_GFX11
103 #endif
104 
105 // MFMA instruction
106 #ifndef __HIP_DEVICE_COMPILE__ // for host code
107 #define CK_USE_AMD_MFMA
108 #elif defined(__gfx9__) // for GPU code
109 #define CK_USE_AMD_MFMA
110 #endif
111 
112 #if(defined(__gfx90a__) || defined(__gfx94__))
113 #define CK_USE_AMD_MFMA_BF16_1K_OP
114 #endif
115 
116 #if defined(__gfx94__)
117 #define CK_USE_AMD_MFMA_GFX940
118 #endif
119 
120 // buffer load
121 #define CK_USE_AMD_BUFFER_LOAD 1
122 
123 // buffer store
124 #define CK_USE_AMD_BUFFER_STORE 1
125 
126 // buffer atomic add: integer
127 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_INTEGER 1
128 
129 // buffer atomic add: floating point
130 #ifndef __HIP_DEVICE_COMPILE__ // for host code
131 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1
132 #elif defined(__gfx9__) || defined(__gfx12__) // for GPU code
133 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1
134 #else // for GPU code
135 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 0
136 #endif
137 
138 #if(defined(__gfx90a__) || defined(__gfx94__)) // for GPU code
139 #define CK_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 1
140 #else
141 #define CK_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 0
142 #endif
143 
144 // inline asm
145 #define CK_USE_AMD_INLINE_ASM 1
146 
147 // inner product (V_MAC/V_FMAC)
148 #define CK_USE_AMD_V_MAC_INLINE_ASM 1
149 
150 // V_DOT inline instructions, less efficient since they require adding
151 // `s_nop`s to avoid hazard
152 #define CK_USE_AMD_V_DOT_INLINE_ASM 0
153 
154 // inner product using V_DOT with DPP8 modifiers
155 #define CK_USE_AMD_V_DOT_DPP8_INLINE_ASM 1
156 
157 // LDS direct loads using inline assembly
158 #define CK_USE_AMD_LDS_DIRECT_LOAD_INLINE_ASM 0
159 
160 // set rounding to nearest even as default for bf16 conversions
161 #define CK_USE_RNE_BF16_CONVERSION 1
162 
163 // set rounding to nearest even as default for f8 conversions
164 #define CK_USE_SR_F8_CONVERSION 0
165 
166 // set rounding to nearest even as default for f6 conversions
167 #define CK_USE_SR_F6_CONVERSION 0
168 
169 // set rounding to nearest even as default for f4 conversions
170 #define CK_USE_SR_F4_CONVERSION 0
171 
172 // shuffle pk_i4 values during conversion to optimize number of binary
173 // operations
174 #define CK_USE_PK4_LAYOUT_SHUFFLE 1
175 
176 // block synchronization only s_wait lgkmcnt(0), not vmcnt(0)
177 #define CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM 1
178 
179 // experimental feature: multi index implemented as array
180 #define CK_EXPERIMENTAL_USE_DYNAMICALLY_INDEXED_MULTI_INDEX 0
181 
182 // experimental feature: static tensor descriptor
183 #define CK_EXPERIMENTAL_STATIC_TENSOR_DESCRIPTOR 0
184 
185 // experimental feature: buffer load/store/atomic-add/ OOB trick
186 // This (ifndef) is a hack to use customized behavior for buffer load rather than using default
187 // setting. Don't use this hack unless absolutely necessary!
188 // FIXME: make the behavior of buffer load a configurable (template) parameter for each usage
189 #ifndef CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK
190 #define CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK 0
191 #endif
192 #define CK_EXPERIMENTAL_USE_BUFFER_STORE_OOB_CHECK_OFFSET_TRICK 1
193 #define CK_EXPERIMENTAL_USE_BUFFER_ATOMIC_ADD_OOB_CHECK_OFFSET_TRICK 1
194 #define CK_EXPERIMENTAL_USE_BUFFER_ATOMIC_MAX_OOB_CHECK_OFFSET_TRICK 1
195 
196 // experimental feature: in-regsiter sub-dword transpose
197 #define CK_EXPERIMENTAL_USE_IN_REGISTER_SUB_DWORD_TRANSPOSE 1
198 
199 // experimental feature: merge transformation use magic number division
200 #define CK_EXPERIMENTAL_MERGE_USE_MAGIC_DIVISION 1
201 
202 // experimental feature: use __builtin_memcpy instead of pointer cast to access a vector from
203 // pointer of scalar
204 #define CK_EXPERIMENTAL_USE_MEMCPY_FOR_VECTOR_ACCESS 0
205 
206 // experimental feature: use __builtin_memcpy instead of union to do bit_cast
207 #define CK_EXPERIMENTAL_USE_MEMCPY_FOR_BIT_CAST 1
208 
209 // experimental feature: optimize for inter-wave scheduling policy
210 #define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING 1
211 #define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING_MAC_CLUSTERS 1
212 // this will let make_default_loop_scheduler() return interwave scheduling flag by default
213 #define CK_EXPERIMENTAL_DEFAULT_TO_INTER_WAVE_SCHEDULING 0
214 // experimental feature: add instances using interwave scheduling
215 #define CK_EXPERIMENTAL_INTER_WAVE_INSTANCES 1
216 // experimental feature: add instances using pipeline v2
217 #define CK_EXPERIMENTAL_PIPELINE_V2_INSTANCES 1
218 // experimental feature: optimize pipeline v2 by IGLP strategy (value=ID of strategy)
219 #ifndef CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT
220 #define CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT 0
221 #endif
222 
223 // hack: have underlying assumption that need to be satsified, otherwise it's a bug
224 // hack for forcing register to keep idx_diff_low_const in SGPR. idx_diff_low_const must be
225 // thread-invariant, otherwise it's a bug
226 // TODO: separate index calculation into "compile-time", "global", "block", "wave", "thread"
227 #define CK_HACK_MERGE_CALCULATE_IDX_DIFF_LOW_CONST_USE_AMD_GCN_READ_FIRST_LANE 0
228 
229 // workaround: compiler crash when compiling recursive lambda
230 #define CK_WORKAROUND_SWDEV_275126 1
231 
232 // workaround: compiler crash when using buffer load/store for i8
233 #define CK_WORKAROUND_SWDEV_XXXXXX_INT8_BUFFER_LOAD_STORE_ISSUE 1
234 
235 // workaround: compiler gnerating inefficient ds_write instructions
236 #define CK_WORKAROUND_SWDEV_XXXXXX_INT8_DS_WRITE_ISSUE 1
237 
238 // workaround: verifaction failure, due to compiler regression, for conv bwd-data fp16 using some
239 // tuning parameter
240 #define CK_WORKAROUND_SWDEV_325164 0
241 
242 // workaround: compiler not emiting reciprocal instruction frm __frcp_rn()
243 #define CK_WORKAROUND_SWDEV_383542 1
244 
245 // workaround: compiler issue on gfx908
246 #define CK_WORKAROUND_SWDEV_388832 1
247 
248 // workaround: compiler issue on gfx950
249 #define CK_WORKAROUND_FP16_TO_FP8_CONVERSION 1
250 
251 // workaround: compiler issue on gfx950
252 #define CK_WORKAROUND_BF16_TO_FP8_CONVERSION 1
253 
254 // denorm test fix, necessary for gfx90a
255 #ifndef CK_GFX90A_DENORM_WORKAROUND
256 #define CK_GFX90A_DENORM_WORKAROUND 0
257 #endif // CK_GFX90A_DENORM_WORKAROUND
258 // Enable only for gfx90a
259 #if defined(__gfx90a__)
260 #if CK_GFX90A_DENORM_WORKAROUND
261 #define CK_GFX90A_DENORM_WORKAROUND 1
262 #endif // CK_GFX90A_DENORM_WORKAROUND is set to 1
263 #else
264 #define CK_GFX90A_DENORM_WORKAROUND 0
265 #endif // gfx90a
266 
267 // set flag to 1 to build deprecated instances
268 #define CK_BUILD_DEPRECATED 1
269 
270 namespace ck {
271 
272 #if defined(__GFX9__) || !defined(__HIP_DEVICE_COMPILE__)
273 __device__ static constexpr int WarpSize = 64;
274 #else
275 __device__ static constexpr int WarpSize = 32;
276 #endif
277 
278 enum struct InMemoryDataOperationEnum
279 {
280  Set,
281  AtomicAdd,
282  AtomicMax,
283  Add
284 };
285 
286 // FIXME: use regular Sequence and remove this
287 template <InMemoryDataOperationEnum... Is>
288 struct InMemoryDataOperationEnumSequence
289 {
290  static constexpr int mSize = sizeof...(Is);
291 
292  __host__ __device__ static constexpr InMemoryDataOperationEnum At(int I)
293  {
294  // the last dummy element is to prevent compiler complain about empty array, when mSize = 0
295  const InMemoryDataOperationEnum mData[mSize + 1] = {Is..., InMemoryDataOperationEnum::Set};
296  return mData[I];
297  }
298 };
299 
300 // index type
301 using index_t = int32_t;
302 using long_index_t = int64_t;
303 
304 } // namespace ck
ck
Definition: ck.hpp:270
ck::InMemoryDataOperationEnum
InMemoryDataOperationEnum
Definition: ck.hpp:279
ck::long_index_t
int64_t long_index_t
Definition: ck.hpp:302
ck::int64_t
long int64_t
Definition: data_type.hpp:464
ck::index_t
int32_t index_t
Definition: ck.hpp:301
int32_t
signed int int32_t
Definition: stdint.h:123
ck::InMemoryDataOperationEnumSequence::At
__host__ static constexpr __device__ InMemoryDataOperationEnum At(int I)
Definition: ck.hpp:292
ck::InMemoryDataOperationEnumSequence::mSize
static constexpr int mSize
Definition: ck.hpp:290