/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm/warp/warp_gemm.hpp Source File

• develop
• include
• ck_tile
• ops
• gemm
• warp

warp_gemm.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_tile/core.hpp"

7 #include "ck_tile/ops/gemm/warp/warp_gemm_impl.hpp"

8 #include "ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp"

9 

10 #include "ck_tile/ops/gemm/warp/warp_gemm_smfmac_impl.hpp"

11 #include "ck_tile/ops/gemm/warp/warp_gemm_attribute_smfmac.hpp"

12 

13 namespace ck_tile {

14 

15 // fp32

16 

17 using WarpGemmMfmaF32F32F32M16N16K4 = WarpGemmImpl<

18  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplF32F32F32M16N16K4<WGAttrCtlEnum::Default_>>>;

19 

20 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

21 using WarpGemmMfmaF32F32F32M16N16K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

22  WarpGemmAttributeMfmaImplF32F32F32M16N16K4<WGAttrCtlEnum::Default_>,

23  4,

24  AttrNumAccess>>;

25 

26 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

27 using WarpGemmMfmaF32F32F32M16N16K16TransposedCDistribution =

28  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

29  WarpGemmAttributeMfmaImplF32F32F32M16N16K4<WGAttrCtlEnum::Default_>,

30  4,

31  AttrNumAccess>>;

32 

33 // fp16

34 

35 using WarpGemmMfmaF16F16F32M32N32K8 = WarpGemmImpl<

36  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

37 

38 using WarpGemmMfmaF16F16F32M16N16K16 = WarpGemmImpl<

39  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

40 

41 #if defined(__gfx950__)

42 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

43 using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl<

44  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>,

45  AttrNumAccess>>;

46 #else

47 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

48 using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

49  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

50  2,

51  AttrNumAccess>>;

52 #endif

53 

54 #if defined(__gfx950__)

55 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

56 using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl<

57  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,

58  AttrNumAccess>>;

59 #else

60 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

61 using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

62  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>,

63  2,

64  AttrNumAccess>>;

65 #endif

66 

67 using WarpGemmMfmaF16F16F32M32N32K8SwizzleA = WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

68  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

69  1>>;

70 

71 using WarpGemmMfmaF16F16F32M32N32K16SwizzleA = WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

72  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

73  2>>;

74 

75 using WarpGemmMfmaF16F16F32M32N32K8TransposedCDistribution =

76  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

77  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

78 

79 using WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution =

80  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

81  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

82 

83 #if defined(__gfx950__)

84 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

85 using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution =

86  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

87  WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>,

88  AttrNumAccess>>;

89 #else

90 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

91 using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution =

92  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

93  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

94  2,

95  AttrNumAccess>>;

96 #endif

97 

98 #if defined(__gfx950__)

99 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

100 using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution =

101  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

102  WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,

103  AttrNumAccess>>;

104 #else

105 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

106 using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution =

107  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

108  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>,

109  2,

110  AttrNumAccess>>;

111 #endif

112 

113 #if defined(__gfx950__)

114 using WarpGemmMfmaF16F16F32M16N16K32SwizzleBTransposedCDistribution =

115  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

116  WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,

117  1>>;

118 

119 using WarpGemmMfmaBf16Bf16F32M16N16K32SwizzleBTransposedCDistribution =

120  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

121  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,

122  1>>;

123 #endif

124 

125 using WarpGemmMfmaF16F16F32M32N32K8SwizzleBTransposedCDistribution =

126  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

127  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

128 

129 #if defined(__gfx950__)

130 using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution =

131  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

132  WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

133 #else

134 using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution =

135  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

136  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

137  2>>;

138 #endif

139 

140 using WarpGemmMfmaF16F16F32M4N64K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

141  WarpGemmAttributeMfmaImplF16F16F32M4N64K4<WGAttrCtlEnum::Default_>,

142  4>>;

143 

144 using WarpGemmMfmaF16F16F32M64N4K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

145  WarpGemmAttributeMfmaImplF16F16F32M64N4K4<WGAttrCtlEnum::Default_>,

146  4>>;

147 

148 // fp16 2:4 structured sparsity

149 using WarpGemmSmfmacF16F16F32M32N32K16 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<

150  WarpGemmAttributeSmfmacImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

151 

152 using WarpGemmSmfmacF16F16F32M16N16K32 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<

153  WarpGemmAttributeSmfmacImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>>>;

154 

155 // bf16

156 using WarpGemmMfmaBf16Bf16F32M32N32K8 = WarpGemmImpl<

157  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

158 

159 using WarpGemmMfmaBf16Bf16F32M16N16K16 = WarpGemmImpl<

160  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

161 

162 #if defined(__gfx950__)

163 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

164 using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl<

165  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>,

166  AttrNumAccess>>;

167 #else

168 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

169 using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

170  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

171  2,

172  AttrNumAccess>>;

173 #endif

174 

175 #if defined(__gfx950__)

176 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

177 using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl<

178  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,

179  AttrNumAccess>>;

180 #else

181 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

182 using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

183  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,

184  2,

185  AttrNumAccess>>;

186 #endif

187 

188 using WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleA = WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

189  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

190  1>>;

191 

192 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleA =

193  WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

194  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

195  2>>;

196 

197 using WarpGemmMfmaBf16Bf16F32M32N32K8TransposedCDistribution =

198  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

199  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

200 

201 using WarpGemmMfmaBf16Bf16F32M16N16K16TransposedCDistribution =

202  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

203  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

204 

205 #if defined(__gfx950__)

206 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

207 using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution =

208  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

209  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>,

210  AttrNumAccess>>;

211 #else

212 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

213 using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution =

214  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

215  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

216  2,

217  AttrNumAccess>>;

218 #endif

219 

220 #if defined(__gfx950__)

221 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

222 using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution =

223  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

224  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,

225  AttrNumAccess>>;

226 #else

227 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

228 using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution =

229  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

230  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,

231  2,

232  AttrNumAccess>>;

233 #endif

234 

235 using WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleBTransposedCDistribution =

236  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

237  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

238 

239 #if defined(__gfx950__)

240 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution =

241  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

242  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

243 #else

244 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution =

245  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

246  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

247  2>>;

248 #endif

249 

250 using WarpGemmMfmaBf16Bf16F32M4N64K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

251  WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4<WGAttrCtlEnum::Default_>,

252  4>>;

253 

254 using WarpGemmMfmaBf16Bf16F32M64N4K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

255  WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4<WGAttrCtlEnum::Default_>,

256  4>>;

257 

258 // fp8

259 

260 using WarpGemmMfma_f32_32x32x16_fp8_fp8 = WarpGemmImpl<

261  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>>>;

262 

263 using WarpGemmMfma_f32_32x32x16_fp8_bf8 = WarpGemmImpl<

264  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>>>;

265 

266 using WarpGemmMfma_f32_32x32x16_bf8_fp8 = WarpGemmImpl<

267  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8<WGAttrCtlEnum::Default_>>>;

268 

269 using WarpGemmMfma_f32_32x32x16_bf8_bf8 = WarpGemmImpl<

270  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>>>;

271 

272 using WarpGemmMfma_f32_32x32x32_fp8_fp8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

273  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>,

274  2>>;

275 

276 using WarpGemmMfma_f32_32x32x32_bf8_bf8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

277  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>,

278  2>>;

279 

280 using WarpGemmMfma_f32_16x16x32_fp8_fp8 = WarpGemmImpl<

281  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>>>;

282 

283 using WarpGemmMfma_f32_16x16x32_fp8_fp8_CTransposed =

284  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

285  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>>>;

286 

287 using WarpGemmMfma_f32_16x16x32_bf8_bf8 = WarpGemmImpl<

288  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>>>;

289 

290 using WarpGemmMfma_f32_16x16x32_bf8_bf8_CTransposed =

291  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

292  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>>>;

293 

294 using WarpGemmMfma_f32_16x16x64_fp8_fp8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

295  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>,

296  2>>;

297 

298 using WarpGemmMfma_f32_16x16x64_bf8_bf8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

299  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>,

300  2>>;

301 

302 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

303 using WarpGemmMfma_f32_16x16x128_fp8_fp8 = WarpGemmImpl<

304  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8<WGAttrCtlEnum::Default_>,

305  AttrNumAccess>>;

306 

307 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

308 using WarpGemmMfma_f32_16x16x128_fp8_bf8 = WarpGemmImpl<

309  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8<WGAttrCtlEnum::Default_>,

310  AttrNumAccess>>;

311 

312 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

313 using WarpGemmMfma_f32_16x16x128_bf8_fp8 = WarpGemmImpl<

314  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8<WGAttrCtlEnum::Default_>,

315  AttrNumAccess>>;

316 

317 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

318 using WarpGemmMfma_f32_16x16x128_bf8_bf8 = WarpGemmImpl<

319  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8<WGAttrCtlEnum::Default_>,

320  AttrNumAccess>>;

321 

322 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

323 using WarpGemmMfma_f32_16x16x128_fp8_fp8_CTransposed =

324  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

325  WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8<WGAttrCtlEnum::Default_>,

326  AttrNumAccess>>;

327 

328 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

329 using WarpGemmMfma_f32_16x16x128_fp8_bf8_CTransposed =

330  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

331  WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8<WGAttrCtlEnum::Default_>,

332  AttrNumAccess>>;

333 

334 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

335 using WarpGemmMfma_f32_16x16x128_bf8_fp8_CTransposed =

336  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

337  WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8<WGAttrCtlEnum::Default_>,

338  AttrNumAccess>>;

339 

340 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

341 using WarpGemmMfma_f32_16x16x128_bf8_bf8_CTransposed =

342  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

343  WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8<WGAttrCtlEnum::Default_>,

344  AttrNumAccess>>;

345 

346 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

347 using WarpGemmMfma_f32_32x32x64_fp8_fp8 = WarpGemmImpl<

348  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_fp8<WGAttrCtlEnum::Default_>,

349  AttrNumAccess>>;

350 

351 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

352 using WarpGemmMfma_f32_32x32x64_fp8_bf8 = WarpGemmImpl<

353  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_bf8<WGAttrCtlEnum::Default_>,

354  AttrNumAccess>>;

355 

356 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

357 using WarpGemmMfma_f32_32x32x64_bf8_fp8 = WarpGemmImpl<

358  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_fp8<WGAttrCtlEnum::Default_>,

359  AttrNumAccess>>;

360 

361 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

362 using WarpGemmMfma_f32_32x32x64_bf8_bf8 = WarpGemmImpl<

363  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_bf8<WGAttrCtlEnum::Default_>,

364  AttrNumAccess>>;

365 

366 using WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed =

367  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

368  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>>>;

369 

370 using WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed =

371  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

372  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>>>;

373 

374 using WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed =

375  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

376  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8<WGAttrCtlEnum::Default_>>>;

377 

378 using WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed =

379  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

380  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>>>;

381 

382 template <index_t swizzle_factor = 2>

383 using WarpGemmMfmaFp8Fp8F32M32N32K16SwizzleBTransposedCDistribution =

384  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

385  WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base<fp8_t, fp8_t, WGAttrCtlEnum::Default_>,

386  2,

387  swizzle_factor>>;

388 

389 // int8

390 using WarpGemmMfma_i32_32x32x16_i8_i8 = WarpGemmImpl<

391  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;

392 

393 using WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed =

394  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

395  WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;

396 

397 using WarpGemmMfma_i32_16x16x32_i8_i8 = WarpGemmImpl<

398  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;

399 

400 using WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed =

401  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

402  WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;

403 

404 } // namespace ck_tile

core.hpp

ck_tile

Definition: cluster_descriptor.hpp:13

ck_tile::WarpGemmAttributeMfma

Definition: warp_gemm_attribute_mfma.hpp:23

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1528

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1164

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1323

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:1795

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:1725

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16

Definition: warp_gemm_attribute_mfma_impl.hpp:666

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32

Definition: warp_gemm_attribute_mfma_impl.hpp:196

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16

Definition: warp_gemm_attribute_mfma_impl.hpp:1049

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8

Definition: warp_gemm_attribute_mfma_impl.hpp:577

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4

Definition: warp_gemm_attribute_mfma_impl.hpp:754

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4

Definition: warp_gemm_attribute_mfma_impl.hpp:844

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16

Definition: warp_gemm_attribute_mfma_impl.hpp:322

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32

Definition: warp_gemm_attribute_mfma_impl.hpp:385

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16

Definition: warp_gemm_attribute_mfma_impl.hpp:935

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8

Definition: warp_gemm_attribute_mfma_impl.hpp:259

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4

Definition: warp_gemm_attribute_mfma_impl.hpp:448

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4

Definition: warp_gemm_attribute_mfma_impl.hpp:512

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4

Definition: warp_gemm_attribute_mfma_impl.hpp:67

ck_tile::WarpGemmAttributeMfmaIterateK_SwizzleA

Definition: warp_gemm_attribute_mfma.hpp:846

ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB

Definition: warp_gemm_attribute_mfma.hpp:702

ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution

Definition: warp_gemm_attribute_mfma.hpp:552

ck_tile::WarpGemmAttributeMfmaIterateK

Definition: warp_gemm_attribute_mfma.hpp:107

ck_tile::WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB

Definition: warp_gemm_attribute_mfma.hpp:456

ck_tile::WarpGemmAttributeMfmaTransposedCDistribution

Definition: warp_gemm_attribute_mfma.hpp:372

ck_tile::WarpGemmAttributeSmfmac

Class describing structured sparsity mfma instructions.

Definition: warp_gemm_attribute_smfmac.hpp:26

ck_tile::WarpGemmAttributeSmfmacImplF16F16F32M16N16K32

Definition: warp_gemm_attribute_smfmac_impl.hpp:65

ck_tile::WarpGemmAttributeSmfmacImplF16F16F32M32N32K16

Definition: warp_gemm_attribute_smfmac_impl.hpp:14

ck_tile::WarpGemmImpl

Definition: warp_gemm_impl.hpp:11

ck_tile::WarpGemmSmfmacImpl

Definition: warp_gemm_smfmac_impl.hpp:11

warp_gemm_attribute_mfma.hpp

warp_gemm_attribute_smfmac.hpp

warp_gemm_impl.hpp

warp_gemm_smfmac_impl.hpp