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"
20 #ifndef CK_TIME_KERNEL
21 #define CK_TIME_KERNEL 1
26 #define CK_CONSTANT_ADDRESS_SPACE __attribute__((address_space(4)))
29 #define CK_USE_LAUNCH_BOUNDS 1
31 #ifdef CK_USE_LAUNCH_BOUNDS
33 #define CK_MAX_THREAD_PER_BLOCK 256
34 #define CK_MIN_BLOCK_PER_CU 2
37 #define CK_WAVELET_MAX_THREAD_PER_BLOCK 512
38 #define CK_WAVELET_MIN_BLOCK_PER_CU 2
42 #ifdef CK_USE_WAVES_PER_EU
44 #ifndef CK_MIN_WAVES_PER_EU
45 #define CK_MIN_WAVES_PER_EU 0
48 #ifndef CK_MAX_WAVES_PER_EU
49 #define CK_MAX_WAVES_PER_EU 0
53 #define CK_USE_WAVES_PER_EU 0
57 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx940__) || defined(__gfx941__) || \
58 defined(__gfx942__) || defined(__gfx950__)
61 #if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
64 #if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__)
67 #if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \
68 defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \
69 defined(__gfx10_3_generic__)
72 #if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \
73 defined(__gfx1103__) || defined(__gfx11_generic__)
76 #if defined(__gfx1200__) || defined(__gfx1201__) || defined(__gfx12_generic__)
81 #ifndef __HIP_DEVICE_COMPILE__
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
92 #ifndef __HIP_DEVICE_COMPILE__
93 #elif defined(__gfx803__) || defined(__gfx900__)
94 #define CK_USE_AMD_V_MAC_F32
95 #elif defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__)
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
106 #ifndef __HIP_DEVICE_COMPILE__
107 #define CK_USE_AMD_MFMA
108 #elif defined(__gfx9__)
109 #define CK_USE_AMD_MFMA
112 #if(defined(__gfx90a__) || defined(__gfx94__))
113 #define CK_USE_AMD_MFMA_BF16_1K_OP
116 #if defined(__gfx94__)
117 #define CK_USE_AMD_MFMA_GFX940
121 #define CK_USE_AMD_BUFFER_LOAD 1
124 #define CK_USE_AMD_BUFFER_STORE 1
127 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_INTEGER 1
130 #ifndef __HIP_DEVICE_COMPILE__
131 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1
132 #elif defined(__gfx9__)
133 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1
135 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 0
138 #if(defined(__gfx90a__) || defined(__gfx94__))
139 #define CK_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 1
141 #define CK_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 0
145 #define CK_USE_AMD_INLINE_ASM 1
148 #define CK_USE_AMD_V_MAC_INLINE_ASM 1
152 #define CK_USE_AMD_V_DOT_INLINE_ASM 0
155 #define CK_USE_AMD_V_DOT_DPP8_INLINE_ASM 1
158 #define CK_USE_AMD_LDS_DIRECT_LOAD_INLINE_ASM 0
161 #define CK_USE_RNE_BF16_CONVERSION 1
164 #define CK_USE_SR_F8_CONVERSION 0
167 #define CK_USE_SR_F6_CONVERSION 0
170 #define CK_USE_SR_F4_CONVERSION 0
174 #define CK_USE_PK4_LAYOUT_SHUFFLE 1
177 #define CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM 1
180 #define CK_EXPERIMENTAL_USE_DYNAMICALLY_INDEXED_MULTI_INDEX 0
183 #define CK_EXPERIMENTAL_STATIC_TENSOR_DESCRIPTOR 0
189 #ifndef CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK
190 #define CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK 0
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
197 #define CK_EXPERIMENTAL_USE_IN_REGISTER_SUB_DWORD_TRANSPOSE 1
200 #define CK_EXPERIMENTAL_MERGE_USE_MAGIC_DIVISION 1
204 #define CK_EXPERIMENTAL_USE_MEMCPY_FOR_VECTOR_ACCESS 0
207 #define CK_EXPERIMENTAL_USE_MEMCPY_FOR_BIT_CAST 1
210 #define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING 1
211 #define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING_MAC_CLUSTERS 1
213 #define CK_EXPERIMENTAL_DEFAULT_TO_INTER_WAVE_SCHEDULING 0
215 #define CK_EXPERIMENTAL_INTER_WAVE_INSTANCES 1
217 #define CK_EXPERIMENTAL_PIPELINE_V2_INSTANCES 1
219 #ifndef CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT
220 #define CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT 0
227 #define CK_HACK_MERGE_CALCULATE_IDX_DIFF_LOW_CONST_USE_AMD_GCN_READ_FIRST_LANE 0
230 #define CK_WORKAROUND_SWDEV_275126 1
233 #define CK_WORKAROUND_SWDEV_XXXXXX_INT8_BUFFER_LOAD_STORE_ISSUE 1
236 #define CK_WORKAROUND_SWDEV_XXXXXX_INT8_DS_WRITE_ISSUE 1
240 #define CK_WORKAROUND_SWDEV_325164 0
243 #define CK_WORKAROUND_SWDEV_383542 1
246 #define CK_WORKAROUND_SWDEV_388832 1
249 #ifndef CK_GFX90A_DENORM_WORKAROUND
250 #define CK_GFX90A_DENORM_WORKAROUND 0
253 #if defined(__gfx90a__)
254 #if CK_GFX90A_DENORM_WORKAROUND
255 #define CK_GFX90A_DENORM_WORKAROUND 1
258 #define CK_GFX90A_DENORM_WORKAROUND 0
262 #define CK_BUILD_DEPRECATED 1
278 static constexpr
int mSize =
sizeof...(Is);
#define CK_DECLARE_ENV_VAR_BOOL(name)
Definition: env.hpp:122
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