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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/ck.hpp Source File
ck.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include "ck/config.h"
7 #include "ck/utility/env.hpp"
8 #ifndef CK_CODE_GEN_RTC
9 #ifndef CK_DONT_USE_HIP_RUNTIME_HEADERS
10 #include "hip/hip_runtime.h"
11 #include "hip/hip_fp16.h"
12 #endif
13 
14 // environment variable to enable logging:
15 // export CK_LOGGING=ON or CK_LOGGING=1 or CK_LOGGING=ENABLED
16 CK_DECLARE_ENV_VAR_BOOL(CK_LOGGING)
17 #endif
18 // to do: add various levels of logging with CK_LOG_LEVEL
19 
20 #ifndef CK_TIME_KERNEL
21 #define CK_TIME_KERNEL 1
22 #endif
23 
24 // constant address space for kernel parameter
25 // https://llvm.org/docs/AMDGPUUsage.html#address-spaces
26 #define CK_CONSTANT_ADDRESS_SPACE __attribute__((address_space(4)))
27 
28 // launch bounds
29 #define CK_USE_LAUNCH_BOUNDS 1
30 
31 #ifdef CK_USE_LAUNCH_BOUNDS
32 // for most kernels
33 #define CK_MAX_THREAD_PER_BLOCK 256
34 #define CK_MIN_BLOCK_PER_CU 2
35 
36 // for wavelet GEMM kernel
37 #define CK_WAVELET_MAX_THREAD_PER_BLOCK 512
38 #define CK_WAVELET_MIN_BLOCK_PER_CU 2
39 #endif
40 
41 // kernel attribute: amdgpu_waves_per_eu()
42 #ifdef CK_USE_WAVES_PER_EU
43 // for 1-wave kernels, control arguments of amdgpu_waves_per_eu() attribute
44 #ifndef CK_MIN_WAVES_PER_EU
45 #define CK_MIN_WAVES_PER_EU 0
46 #endif
47 
48 #ifndef CK_MAX_WAVES_PER_EU
49 #define CK_MAX_WAVES_PER_EU 0
50 #endif
51 
52 #else
53 #define CK_USE_WAVES_PER_EU 0
54 #endif
55 
56 // define general macros for various architectures
57 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx940__) || defined(__gfx941__) || \
58  defined(__gfx942__) || defined(__gfx950__)
59 #define __gfx9__
60 #endif
61 #if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
62 #define __gfx94__
63 #endif
64 #if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__)
65 #define __gfx101__
66 #endif
67 #if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \
68  defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \
69  defined(__gfx10_3_generic__)
70 #define __gfx103__
71 #endif
72 #if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \
73  defined(__gfx1103__) || defined(__gfx11_generic__)
74 #define __gfx11__
75 #endif
76 #if defined(__gfx1200__) || defined(__gfx1201__) || defined(__gfx12_generic__)
77 #define __gfx12__
78 #endif
79 
80 // buffer resource
81 #ifndef __HIP_DEVICE_COMPILE__ // for host code
82 #define CK_BUFFER_RESOURCE_3RD_DWORD -1
83 #elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx9__)
84 #define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000
85 #elif defined(__gfx103__)
86 #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000
87 #elif defined(__gfx11__) || defined(__gfx12__)
88 #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31004000
89 #endif
90 
91 // FMA instruction
92 #ifndef __HIP_DEVICE_COMPILE__ // for host code, define nothing
93 #elif defined(__gfx803__) || defined(__gfx900__) // for GPU code
94 #define CK_USE_AMD_V_MAC_F32
95 #elif defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__) // for GPU code
96 #define CK_USE_AMD_V_FMAC_F32
97 #define CK_USE_AMD_V_DOT2_F32_F16
98 #define CK_USE_AMD_V_DOT4_I32_I8
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__) // 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 // denorm test fix, necessary for gfx90a
249 #ifndef CK_GFX90A_DENORM_WORKAROUND
250 #define CK_GFX90A_DENORM_WORKAROUND 0
251 #endif // CK_GFX90A_DENORM_WORKAROUND
252 // Enable only for gfx90a
253 #if defined(__gfx90a__)
254 #if CK_GFX90A_DENORM_WORKAROUND
255 #define CK_GFX90A_DENORM_WORKAROUND 1
256 #endif // CK_GFX90A_DENORM_WORKAROUND is set to 1
257 #else
258 #define CK_GFX90A_DENORM_WORKAROUND 0
259 #endif // gfx90a
260 
261 // set flag to 1 to build deprecated instances
262 #define CK_BUILD_DEPRECATED 1
263 
264 namespace ck {
265 
267 {
268  Set,
269  AtomicAdd,
270  AtomicMax,
271  Add
272 };
273 
274 // FIXME: use regular Sequence and remove this
275 template <InMemoryDataOperationEnum... Is>
277 {
278  static constexpr int mSize = sizeof...(Is);
279 
280  __host__ __device__ static constexpr InMemoryDataOperationEnum At(int I)
281  {
282  // the last dummy element is to prevent compiler complain about empty array, when mSize = 0
284  return mData[I];
285  }
286 };
287 
288 // index type
289 using index_t = int32_t;
291 
292 } // namespace ck
#define CK_DECLARE_ENV_VAR_BOOL(name)
Definition: env.hpp:122
Definition: ck.hpp:264
InMemoryDataOperationEnum
Definition: ck.hpp:267
int64_t long_index_t
Definition: ck.hpp:290
long int64_t
Definition: data_type.hpp:2474
int32_t index_t
Definition: ck.hpp:289
__host__ static constexpr __device__ InMemoryDataOperationEnum At(int I)
Definition: ck.hpp:280
static constexpr int mSize
Definition: ck.hpp:278