/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/config.hpp Source File
/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/config.hpp Source File
Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/config.hpp Source File
Go to the documentation of this file.
6 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__) || \
7 defined(__gfx9_4_generic__)
10 #if defined(__gfx942__) || defined(__gfx950__) || defined(__gfx9_4_generic__)
13 #if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \
14 defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \
15 defined(__gfx10_3_generic__)
18 #if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \
19 defined(__gfx1103__) || defined(__gfx1150__) || defined(__gfx1151__) || \
20 defined(__gfx1152__) || defined(__gfx11_generic__)
23 #if defined(__gfx1200__) || defined(__gfx1201__) || defined(__gfx12_generic__)
27 #include "hip/hip_version.h"
28 #ifndef CK_TILE_DONT_USE_HIP_RUNTIME_HEADERS
29 #include "hip/hip_runtime.h"
30 #include "hip/hip_fp16.h"
34 #define CK_TILE_HOST inline __host__
35 #define CK_TILE_DEVICE inline __device__
36 #define CK_TILE_HOST_DEVICE inline __host__ __device__
37 #define CK_TILE_DEVICE_EXTERN __device__
38 #define CK_TILE_HOST_DEVICE_EXTERN __host__ __device__
40 #define CK_TILE_HOST inline
41 #define CK_TILE_DEVICE inline
42 #define CK_TILE_HOST_DEVICE inline
43 #define CK_TILE_DEVICE_EXTERN
44 #define CK_TILE_HOST_DEVICE_EXTERN
51 #define CK_TILE_GENERIC_ADDR __attribute__((address_space(0)))
52 #define CK_TILE_GLOBAL_ADDR __attribute__((address_space(1)))
53 #define CK_TILE_LDS_ADDR __attribute__((address_space(3)))
54 #define CK_TILE_BUF_RES_ADDR __attribute__((address_space(8)))
56 #define CK_TILE_GENERIC_ADDR
57 #define CK_TILE_GLOBAL_ADDR
58 #define CK_TILE_LDS_ADDR
59 #define CK_TILE_BUF_RES_ADDR
61 #ifndef CK_TILE_USE_CUSTOM_DATA_TYPE
62 #define CK_TILE_USE_CUSTOM_DATA_TYPE 0
65 #define CK_TILE_FLOAT_TO_BFLOAT16_STANDARD 0
66 #define CK_TILE_FLOAT_TO_BFLOAT16_TRUNCATE_WITH_NAN 1
67 #define CK_TILE_FLOAT_TO_BFLOAT16_TRUNCATE 2
68 #define CK_TILE_FLOAT_TO_BFLOAT16_STANDARD_ASM 3
69 #define CK_TILE_FLOAT_TO_BFLOAT16_RTA_ASM 4
71 #ifndef CK_TILE_FLOAT_TO_BFLOAT16_DEFAULT
72 #define CK_TILE_FLOAT_TO_BFLOAT16_DEFAULT CK_TILE_FLOAT_TO_BFLOAT16_TRUNCATE
75 #define CK_TILE_FLOAT_TO_FP8_STANDARD 0
76 #define CK_TILE_FLOAT_TO_FP8_STOCHASTIC 1
78 #ifndef CK_TILE_FLOAT_TO_FP8_DEFAULT
79 #define CK_TILE_FLOAT_TO_FP8_DEFAULT CK_TILE_FLOAT_TO_FP8_STANDARD
84 #define CK_TILE_STATICALLY_INDEXED_ARRAY_USE_ARRAY 0
85 #define CK_TILE_STATICALLY_INDEXED_ARRAY_USE_TUPLE 1
86 #ifndef CK_TILE_STATICALLY_INDEXED_ARRAY_DEFAULT
87 #define CK_TILE_STATICALLY_INDEXED_ARRAY_DEFAULT CK_TILE_STATICALLY_INDEXED_ARRAY_USE_TUPLE
90 #define CK_TILE_THREAD_BUFFER_USE_ARRAY 0
91 #define CK_TILE_THREAD_BUFFER_USE_TUPLE 1
92 #ifndef CK_TILE_THREAD_BUFFER_DEFAULT
93 #define CK_TILE_THREAD_BUFFER_DEFAULT CK_TILE_THREAD_BUFFER_USE_ARRAY
96 #ifndef CK_TILE_TUPLE_CTOR_WITH_INITIALIZER_LIST
97 #if CK_TILE_THREAD_BUFFER_DEFAULT == CK_TILE_THREAD_BUFFER_USE_TUPLE
100 #define CK_TILE_TUPLE_CTOR_WITH_INITIALIZER_LIST 1
102 #define CK_TILE_TUPLE_CTOR_WITH_INITIALIZER_LIST 0
106 #ifndef CK_TILE_USE_LAUNCH_BOUNDS
107 #define CK_TILE_USE_LAUNCH_BOUNDS 1
110 #ifndef CK_TILE_TIME_KERNEL
111 #define CK_TILE_TIME_KERNEL 1
114 #define CK_TILE_MAX_THREAD_PER_BLOCK 256
115 #define CK_TILE_MIN_BLOCK_PER_CU 2
117 #ifndef CK_TILE_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK
118 #define CK_TILE_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK 0
121 #ifndef CK_TILE_EXPERIMENTAL_USE_BUFFER_STORE_OOB_CHECK_OFFSET_TRICK
122 #define CK_TILE_EXPERIMENTAL_USE_BUFFER_STORE_OOB_CHECK_OFFSET_TRICK 1
125 #ifndef CK_TILE_EXPERIMENTAL_USE_BUFFER_ATOMIC_ADD_OOB_CHECK_OFFSET_TRICK
126 #define CK_TILE_EXPERIMENTAL_USE_BUFFER_ATOMIC_ADD_OOB_CHECK_OFFSET_TRICK 1
129 #ifndef CK_TILE_EXPERIMENTAL_USE_BUFFER_ATOMIC_MAX_OOB_CHECK_OFFSET_TRICK
130 #define CK_TILE_EXPERIMENTAL_USE_BUFFER_ATOMIC_MAX_OOB_CHECK_OFFSET_TRICK 1
133 #ifndef CK_TILE_USE_AMD_LDS_DIRECT_LOAD_INLINE_ASM
134 #define CK_TILE_USE_AMD_LDS_DIRECT_LOAD_INLINE_ASM 1
137 #ifndef CK_TILE_USE_AMD_BUFFER_LOAD
138 #define CK_TILE_USE_AMD_BUFFER_LOAD 1
141 #ifndef CK_TILE_USE_AMD_BUFFER_STORE
142 #define CK_TILE_USE_AMD_BUFFER_STORE 1
145 #ifndef CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_INTEGER
146 #define CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_INTEGER 1
149 #ifndef CK_TILE_USE_PK4_LAYOUT_SHUFFLE
150 #define CK_TILE_USE_PK4_LAYOUT_SHUFFLE 1
154 #ifndef __HIP_DEVICE_COMPILE__
155 #define CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1
156 #elif defined(__gfx9__) || defined(__gfx12__)
157 #define CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1
159 #define CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 0
162 #if(defined(__gfx90a__) || defined(__gfx94__))
163 #define CK_TILE_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 1
165 #define CK_TILE_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 0
168 #ifndef CK_TILE_EXPERIMENTAL_USE_MEMCPY_FOR_VECTOR_ACCESS
169 #define CK_TILE_EXPERIMENTAL_USE_MEMCPY_FOR_VECTOR_ACCESS 0
172 #ifndef CK_TILE_WORKAROUND_SWDEV_XXXXXX_INT8_DS_WRITE_ISSUE
173 #define CK_TILE_WORKAROUND_SWDEV_XXXXXX_INT8_DS_WRITE_ISSUE 1
176 #ifndef CK_TILE_WORKAROUND_ROCM_6_1_SCRATCH_MEMORY_ISSUE
177 #if HIP_VERSION_MAJOR == 6 && HIP_VERSION_MINOR == 1 && HIP_VERSION_PATCH >= 40091
178 #define CK_TILE_WORKAROUND_ROCM_6_1_SCRATCH_MEMORY_ISSUE 1
180 #define CK_TILE_WORKAROUND_ROCM_6_1_SCRATCH_MEMORY_ISSUE 0
185 #ifndef CK_TILE_WORKAROUND_ROCM_6_2_SCRATCH_MEMORY_ISSUE
186 #if(HIP_VERSION_MAJOR == 6 && HIP_VERSION_MINOR == 2 && HIP_VERSION_PATCH >= 41133) || \
187 (HIP_VERSION_MAJOR == 6 && HIP_VERSION_MINOR == 3 && HIP_VERSION_PATCH >= 42131) || \
188 (HIP_VERSION_MAJOR == 6 && HIP_VERSION_MINOR > 3)
189 #define CK_TILE_WORKAROUND_ROCM_6_2_SCRATCH_MEMORY_ISSUE 1
191 #define CK_TILE_WORKAROUND_ROCM_6_2_SCRATCH_MEMORY_ISSUE 0
196 #ifndef CK_TILE_USE_LLVM_BUILTIN_BF16
197 #if(HIP_VERSION_MAJOR == 6 && HIP_VERSION_MINOR == 5 && HIP_VERSION_PATCH >= 50421) || \
198 (HIP_VERSION_MAJOR >= 7)
199 #define CK_TILE_USE_LLVM_BUILTIN_BF16 1
201 #define CK_TILE_USE_LLVM_BUILTIN_BF16 0
205 #ifndef CK_TILE_DEBUG_LOG
206 #define CK_TILE_DEBUG_LOG 0
209 #ifndef __HIP_DEVICE_COMPILE__
210 #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0xffffffff
211 #elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || \
213 #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x00020000
214 #elif defined(__gfx103__)
215 #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x31014000
216 #elif defined(__gfx11__) || defined(__gfx12__)
217 #define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x31004000
220 #ifndef CK_TILE_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM
221 #define CK_TILE_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM 1
224 #ifndef CK_TILE_USE_SUBDWORD_TILE_CAST
225 #define CK_TILE_USE_SUBDWORD_TILE_CAST 0
228 #ifndef CK_TILE_USE_PK_FP16_TILE_CAST
229 #define CK_TILE_USE_PK_FP16_TILE_CAST 0
233 #ifndef CK_TILE_FMHA_FWD_FAST_EXP2
234 #define CK_TILE_FMHA_FWD_FAST_EXP2 0
237 #ifndef CK_TILE_FMHA_FLOAT_TO_FLOAT16_RTN
238 #define CK_TILE_FMHA_FLOAT_TO_FLOAT16_RTN 0
241 #ifndef CK_TILE_BUFFER_LOAD_RAW_BF16_WA
242 #define CK_TILE_BUFFER_LOAD_RAW_BF16_WA 1
246 #ifndef CK_TILE_WORKAROUND_SWDEV_383542
247 #define CK_TILE_WORKAROUND_SWDEV_383542 1
250 #ifndef CK_TILE_REFERENCE_MOE_SORTING_MOCK_ID
251 #define CK_TILE_REFERENCE_MOE_SORTING_MOCK_ID 1
254 #ifndef CK_TILE_USE_OCP_FP8
255 #if defined(__HIP_DEVICE_COMPILE__)
256 #if defined(__gfx950__) || defined(__gfx12__)
257 #define CK_TILE_USE_OCP_FP8 1
259 #define CK_TILE_USE_OCP_FP8 0
262 #define CK_TILE_USE_OCP_FP8 0
266 #ifndef CK_TILE_USE_BUFFER_ADDRESSING_BUILTIN
267 #if __clang_major__ >= 20
268 #define CK_TILE_USE_BUFFER_ADDRESSING_BUILTIN 1
270 #define CK_TILE_USE_BUFFER_ADDRESSING_BUILTIN 0
274 #ifndef CK_TILE_WA_ISSUE_2028
275 #define CK_TILE_WA_ISSUE_2028 0
280 #ifndef CK_TILE_ENC_SUPPORT_Y_TO_R
281 #define CK_TILE_ENC_SUPPORT_Y_TO_R 0