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

• develop
• include
• ck

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 

8 #if !defined(__HIPCC_RTC__) || !defined(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 #endif

14 // to do: add various levels of logging with CK_LOG_LEVEL

15 

16 #ifndef CK_TIME_KERNEL

17 #define CK_TIME_KERNEL 1

18 #endif

19 

20 // constant address space for kernel parameter

21 // https://llvm.org/docs/AMDGPUUsage.html#address-spaces

22 #define CK_CONSTANT_ADDRESS_SPACE __attribute__((address_space(4)))

23 

24 // launch bounds

25 #define CK_USE_LAUNCH_BOUNDS 1

26 

27 #ifdef CK_USE_LAUNCH_BOUNDS

28 // for most kernels

29 #define CK_MAX_THREAD_PER_BLOCK 256

30 #define CK_MIN_BLOCK_PER_CU 2

31 

32 // for wavelet GEMM kernel

33 #define CK_WAVELET_MAX_THREAD_PER_BLOCK 512

34 #define CK_WAVELET_MIN_BLOCK_PER_CU 2

35 #endif

36 

37 // kernel attribute: amdgpu_waves_per_eu()

38 #ifdef CK_USE_WAVES_PER_EU

39 // for 1-wave kernels, control arguments of amdgpu_waves_per_eu() attribute

40 #ifndef CK_MIN_WAVES_PER_EU

41 #define CK_MIN_WAVES_PER_EU 0

42 #endif

43 

44 #ifndef CK_MAX_WAVES_PER_EU

45 #define CK_MAX_WAVES_PER_EU 0

46 #endif

47 

48 #else

49 #define CK_USE_WAVES_PER_EU 0

50 #endif

51 

52 // define general macros for various architectures

53 #if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__) || \

54  defined(__gfx9_4_generic__)

55 #define __gfx9__

56 #endif

57 #if defined(__gfx942__) || defined(__gfx950__) || defined(__gfx9_4_generic__)

58 #define __gfx94__

59 #endif

60 #if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__)

61 #define __gfx101__

62 #endif

63 #if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \

64  defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \

65  defined(__gfx10_3_generic__)

66 #define __gfx103__

67 #endif

68 #if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \

69  defined(__gfx1103__) || defined(__gfx1150__) || defined(__gfx1151__) || \

70  defined(__gfx1152__) || defined(__gfx11_generic__)

71 #define __gfx11__

72 #endif

73 #if defined(__gfx1200__) || defined(__gfx1201__) || defined(__gfx12_generic__)

74 #define __gfx12__

75 #endif

76 

77 // buffer resource

78 #ifndef __HIP_DEVICE_COMPILE__ // for host code

79 #define CK_BUFFER_RESOURCE_3RD_DWORD -1

80 #elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx9__)

81 #define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000

82 #elif defined(__gfx103__)

83 #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000

84 #elif defined(__gfx11__) || defined(__gfx12__)

85 #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31004000

86 #endif

87 

88 // FMA instruction

89 #ifndef __HIP_DEVICE_COMPILE__ // for host code, define nothing

90 #elif defined(__gfx803__) || defined(__gfx900__) // for GPU code

91 #define CK_USE_AMD_V_MAC_F32

92 #elif defined(__gfx906__) || defined(__gfx9__) || defined(__gfx103__) // for GPU code

93 #define CK_USE_AMD_V_FMAC_F32

94 #define CK_USE_AMD_V_DOT2_F32_F16

95 #define CK_USE_AMD_V_DOT4_I32_I8

96 #elif defined(__gfx11__) || defined(__gfx12__)

97 #define CK_USE_AMD_V_FMAC_F32

98 #define CK_USE_AMD_V_DOT2_F32_F16

99 #define CK_USE_AMD_V_DOT4_I32_I8_GFX11

100 #endif

101 

102 // MFMA instruction

103 #ifndef __HIP_DEVICE_COMPILE__ // for host code

104 #define CK_USE_AMD_MFMA

105 #elif defined(__gfx9__) // for GPU code

106 #define CK_USE_AMD_MFMA

107 #endif

108 

109 #if(defined(__gfx90a__) || defined(__gfx94__))

110 #define CK_USE_AMD_MFMA_BF16_1K_OP

111 #endif

112 

113 #if defined(__gfx94__)

114 #define CK_USE_AMD_MFMA_GFX940

115 #endif

116 

117 // buffer load

118 #define CK_USE_AMD_BUFFER_LOAD 1

119 

120 // buffer store

121 #define CK_USE_AMD_BUFFER_STORE 1

122 

123 // buffer atomic add: integer

124 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_INTEGER 1

125 

126 // buffer atomic add: floating point

127 #ifndef __HIP_DEVICE_COMPILE__ // for host code

128 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1

129 #elif defined(__gfx9__) || defined(__gfx12__) // for GPU code

130 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 1

131 #else // for GPU code

132 #define CK_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT 0

133 #endif

134 

135 #if(defined(__gfx90a__) || defined(__gfx94__)) // for GPU code

136 #define CK_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 1

137 #else

138 #define CK_USE_AMD_BUFFER_ATOMIC_MAX_FLOAT64 0

139 #endif

140 

141 // inline asm

142 #define CK_USE_AMD_INLINE_ASM 1

143 

144 // inner product (V_MAC/V_FMAC)

145 #define CK_USE_AMD_V_MAC_INLINE_ASM 1

146 

147 // V_DOT inline instructions, less efficient since they require adding

148 // `s_nop`s to avoid hazard

149 #define CK_USE_AMD_V_DOT_INLINE_ASM 0

150 

151 // inner product using V_DOT with DPP8 modifiers

152 #define CK_USE_AMD_V_DOT_DPP8_INLINE_ASM 1

153 

154 // LDS direct loads using inline assembly

155 #define CK_USE_AMD_LDS_DIRECT_LOAD_INLINE_ASM 0

156 

157 // set rounding to nearest even as default for bf16 conversions

158 #define CK_USE_RNE_BF16_CONVERSION 1

159 

160 // set rounding to nearest even as default for f8 conversions

161 #define CK_USE_SR_F8_CONVERSION 0

162 

163 // set rounding to nearest even as default for f6 conversions

164 #define CK_USE_SR_F6_CONVERSION 0

165 

166 // set rounding to nearest even as default for f4 conversions

167 #define CK_USE_SR_F4_CONVERSION 0

168 

169 // shuffle pk_i4 values during conversion to optimize number of binary

170 // operations

171 #define CK_USE_PK4_LAYOUT_SHUFFLE 1

172 

173 // block synchronization only s_wait lgkmcnt(0), not vmcnt(0)

174 #define CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM 1

175 

176 // experimental feature: multi index implemented as array

177 #define CK_EXPERIMENTAL_USE_DYNAMICALLY_INDEXED_MULTI_INDEX 0

178 

179 // experimental feature: static tensor descriptor

180 #define CK_EXPERIMENTAL_STATIC_TENSOR_DESCRIPTOR 0

181 

182 // experimental feature: buffer load/store/atomic-add/ OOB trick

183 // This (ifndef) is a hack to use customized behavior for buffer load rather than using default

184 // setting. Don't use this hack unless absolutely necessary!

185 // FIXME: make the behavior of buffer load a configurable (template) parameter for each usage

186 #ifndef CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK

187 #define CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK 0

188 #endif

189 #define CK_EXPERIMENTAL_USE_BUFFER_STORE_OOB_CHECK_OFFSET_TRICK 1

190 #define CK_EXPERIMENTAL_USE_BUFFER_ATOMIC_ADD_OOB_CHECK_OFFSET_TRICK 1

191 #define CK_EXPERIMENTAL_USE_BUFFER_ATOMIC_MAX_OOB_CHECK_OFFSET_TRICK 1

192 

193 // experimental feature: in-regsiter sub-dword transpose

194 #define CK_EXPERIMENTAL_USE_IN_REGISTER_SUB_DWORD_TRANSPOSE 1

195 

196 // experimental feature: merge transformation use magic number division

197 #define CK_EXPERIMENTAL_MERGE_USE_MAGIC_DIVISION 1

198 

199 // experimental feature: use __builtin_memcpy instead of pointer cast to access a vector from

200 // pointer of scalar

201 #define CK_EXPERIMENTAL_USE_MEMCPY_FOR_VECTOR_ACCESS 0

202 

203 // experimental feature: use __builtin_memcpy instead of union to do bit_cast

204 #define CK_EXPERIMENTAL_USE_MEMCPY_FOR_BIT_CAST 1

205 

206 // experimental feature: optimize for inter-wave scheduling policy

207 #define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING 1

208 #define CK_EXPERIMENTAL_INTER_WAVE_SCHEDULING_MAC_CLUSTERS 1

209 // this will let make_default_loop_scheduler() return interwave scheduling flag by default

210 #define CK_EXPERIMENTAL_DEFAULT_TO_INTER_WAVE_SCHEDULING 0

211 // experimental feature: add instances using interwave scheduling

212 #define CK_EXPERIMENTAL_INTER_WAVE_INSTANCES 1

213 // experimental feature: add instances using pipeline v2

214 #define CK_EXPERIMENTAL_PIPELINE_V2_INSTANCES 1

215 // experimental feature: optimize pipeline v2 by IGLP strategy (value=ID of strategy)

216 #ifndef CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT

217 #define CK_EXPERIMENTAL_PIPELINE_V2_IGLP_OPT 0

218 #endif

219 

220 // hack: have underlying assumption that need to be satsified, otherwise it's a bug

221 // hack for forcing register to keep idx_diff_low_const in SGPR. idx_diff_low_const must be

222 // thread-invariant, otherwise it's a bug

223 // TODO: separate index calculation into "compile-time", "global", "block", "wave", "thread"

224 #define CK_HACK_MERGE_CALCULATE_IDX_DIFF_LOW_CONST_USE_AMD_GCN_READ_FIRST_LANE 0

225 

226 // workaround: compiler crash when compiling recursive lambda

227 #define CK_WORKAROUND_SWDEV_275126 1

228 

229 // workaround: compiler crash when using buffer load/store for i8

230 #define CK_WORKAROUND_SWDEV_XXXXXX_INT8_BUFFER_LOAD_STORE_ISSUE 1

231 

232 // workaround: compiler gnerating inefficient ds_write instructions

233 #define CK_WORKAROUND_SWDEV_XXXXXX_INT8_DS_WRITE_ISSUE 1

234 

235 // workaround: verifaction failure, due to compiler regression, for conv bwd-data fp16 using some

236 // tuning parameter

237 #define CK_WORKAROUND_SWDEV_325164 0

238 

239 // workaround: compiler not emiting reciprocal instruction frm __frcp_rn()

240 #define CK_WORKAROUND_SWDEV_383542 1

241 

242 // workaround: compiler issue on gfx908

243 #define CK_WORKAROUND_SWDEV_388832 1

244 

245 // workaround: compiler issue on gfx950

246 #define CK_WORKAROUND_FP16_TO_FP8_CONVERSION 1

247 

248 // workaround: compiler issue on gfx950

249 #define CK_WORKAROUND_BF16_TO_FP8_CONVERSION 1

250 

251 // denorm test fix, necessary for gfx90a

252 #ifndef CK_GFX90A_DENORM_WORKAROUND

253 #define CK_GFX90A_DENORM_WORKAROUND 0

254 #endif // CK_GFX90A_DENORM_WORKAROUND

255 // Enable only for gfx90a

256 #if defined(__gfx90a__)

257 #if CK_GFX90A_DENORM_WORKAROUND

258 #define CK_GFX90A_DENORM_WORKAROUND 1

259 #endif // CK_GFX90A_DENORM_WORKAROUND is set to 1

260 #else

261 #define CK_GFX90A_DENORM_WORKAROUND 0

262 #endif // gfx90a

263 

264 // set flag to 1 to build deprecated instances

265 #define CK_BUILD_DEPRECATED 1

266 

267 namespace ck {

268 

269 #if defined(__GFX9__) || !defined(__HIP_DEVICE_COMPILE__)

270 __device__ static constexpr int WarpSize = 64;

271 #else

272 __device__ static constexpr int WarpSize = 32;

273 #endif

274 

275 enum struct InMemoryDataOperationEnum

276 {

277  Set,

278  AtomicAdd,

279  AtomicMax,

280  Add

281 };

282 

283 // FIXME: use regular Sequence and remove this

284 template <InMemoryDataOperationEnum... Is>

285 struct InMemoryDataOperationEnumSequence

286 {

287  static constexpr int mSize = sizeof...(Is);

288 

289  __host__ __device__ static constexpr InMemoryDataOperationEnum At(int I)

290  {

291  // the last dummy element is to prevent compiler complain about empty array, when mSize = 0

292  const InMemoryDataOperationEnum mData[mSize + 1] = {Is..., InMemoryDataOperationEnum::Set};

293  return mData[I];

294  }

295 };

296 

297 // index type

298 using index_t = int32_t;

299 using long_index_t = int64_t;

300 

301 } // namespace ck

ck

Definition: ck.hpp:267

ck::InMemoryDataOperationEnum

InMemoryDataOperationEnum

Definition: ck.hpp:276

ck::InMemoryDataOperationEnum::Set

@ Set

ck::InMemoryDataOperationEnum::AtomicMax

@ AtomicMax

ck::InMemoryDataOperationEnum::AtomicAdd

@ AtomicAdd

ck::InMemoryDataOperationEnum::Add

@ Add

ck::long_index_t

int64_t long_index_t

Definition: ck.hpp:299

ck::int64_t

long int64_t

Definition: data_type.hpp:461

ck::index_t

int32_t index_t

Definition: ck.hpp:298

int32_t

signed int int32_t

Definition: stdint.h:123

ck::InMemoryDataOperationEnumSequence

Definition: ck.hpp:286

ck::InMemoryDataOperationEnumSequence::At

__host__ static constexpr __device__ InMemoryDataOperationEnum At(int I)

Definition: ck.hpp:289

ck::InMemoryDataOperationEnumSequence::mSize

static constexpr int mSize

Definition: ck.hpp:287