/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 // fp16

16 

17 using WarpGemmMfmaF16F16F32M32N32K8 = WarpGemmImpl<

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

19 

20 using WarpGemmMfmaF16F16F32M16N16K16 = WarpGemmImpl<

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

22 

23 #if defined(__gfx950__)

24 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

25 using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl<

26  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>,

27  AttrNumAccess>>;

28 #else

29 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

30 using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

31  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

32  2,

33  AttrNumAccess>>;

34 #endif

35 

36 #if defined(__gfx950__)

37 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

38 using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl<

39  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,

40  AttrNumAccess>>;

41 #else

42 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

43 using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

44  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>,

45  2,

46  AttrNumAccess>>;

47 #endif

48 

49 using WarpGemmMfmaF16F16F32M32N32K8SwizzleA = WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

50  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

51  1>>;

52 

53 using WarpGemmMfmaF16F16F32M32N32K16SwizzleA = WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

54  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

55  2>>;

56 

57 using WarpGemmMfmaF16F16F32M32N32K8TransposedCDistribution =

58  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

59  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

60 

61 using WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution =

62  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

63  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

64 

65 #if defined(__gfx950__)

66 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

67 using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution =

68  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

69  WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>,

70  AttrNumAccess>>;

71 #else

72 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

73 using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution =

74  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

75  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

76  2,

77  AttrNumAccess>>;

78 #endif

79 

80 #if defined(__gfx950__)

81 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

82 using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution =

83  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

84  WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,

85  AttrNumAccess>>;

86 #else

87 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

88 using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution =

89  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

90  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>,

91  2,

92  AttrNumAccess>>;

93 #endif

94 

95 #if defined(__gfx950__)

96 using WarpGemmMfmaF16F16F32M16N16K32SwizzleBTransposedCDistribution =

97  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

98  WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,

99  1>>;

100 

101 using WarpGemmMfmaBf16Bf16F32M16N16K32SwizzleBTransposedCDistribution =

102  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

103  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,

104  1>>;

105 #endif

106 

107 using WarpGemmMfmaF16F16F32M32N32K8SwizzleBTransposedCDistribution =

108  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

109  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

110 

111 #if defined(__gfx950__)

112 using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution =

113  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

114  WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

115 #else

116 using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution =

117  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

118  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

119  2>>;

120 #endif

121 

122 using WarpGemmMfmaF16F16F32M4N64K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

123  WarpGemmAttributeMfmaImplF16F16F32M4N64K4<WGAttrCtlEnum::Default_>,

124  4>>;

125 

126 using WarpGemmMfmaF16F16F32M64N4K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

127  WarpGemmAttributeMfmaImplF16F16F32M64N4K4<WGAttrCtlEnum::Default_>,

128  4>>;

129 

130 // fp16 2:4 structured sparsity

131 using WarpGemmSmfmacF16F16F32M32N32K16 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<

132  WarpGemmAttributeSmfmacImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

133 

134 using WarpGemmSmfmacF16F16F32M16N16K32 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<

135  WarpGemmAttributeSmfmacImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>>>;

136 

137 // bf16

138 using WarpGemmMfmaBf16Bf16F32M32N32K8 = WarpGemmImpl<

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

140 

141 using WarpGemmMfmaBf16Bf16F32M16N16K16 = WarpGemmImpl<

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

143 

144 #if defined(__gfx950__)

145 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

146 using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl<

147  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>,

148  AttrNumAccess>>;

149 #else

150 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

151 using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

152  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

153  2,

154  AttrNumAccess>>;

155 #endif

156 

157 #if defined(__gfx950__)

158 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

159 using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl<

160  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,

161  AttrNumAccess>>;

162 #else

163 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

164 using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

165  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,

166  2,

167  AttrNumAccess>>;

168 #endif

169 

170 using WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleA = WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

171  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

172  1>>;

173 

174 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleA =

175  WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

176  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

177  2>>;

178 

179 using WarpGemmMfmaBf16Bf16F32M32N32K8TransposedCDistribution =

180  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

181  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

182 

183 using WarpGemmMfmaBf16Bf16F32M16N16K16TransposedCDistribution =

184  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

185  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

186 

187 #if defined(__gfx950__)

188 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

189 using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution =

190  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

191  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>,

192  AttrNumAccess>>;

193 #else

194 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

195 using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution =

196  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

197  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

198  2,

199  AttrNumAccess>>;

200 #endif

201 

202 #if defined(__gfx950__)

203 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

204 using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution =

205  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

206  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,

207  AttrNumAccess>>;

208 #else

209 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

210 using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution =

211  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

212  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,

213  2,

214  AttrNumAccess>>;

215 #endif

216 

217 using WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleBTransposedCDistribution =

218  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

219  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

220 

221 #if defined(__gfx950__)

222 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution =

223  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

224  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

225 #else

226 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution =

227  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

228  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

229  2>>;

230 #endif

231 

232 using WarpGemmMfmaBf16Bf16F32M4N64K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

233  WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4<WGAttrCtlEnum::Default_>,

234  4>>;

235 

236 using WarpGemmMfmaBf16Bf16F32M64N4K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

237  WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4<WGAttrCtlEnum::Default_>,

238  4>>;

239 

240 // fp8

241 

242 using WarpGemmMfma_f32_32x32x16_fp8_fp8 = WarpGemmImpl<

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

244 

245 using WarpGemmMfma_f32_32x32x16_fp8_bf8 = WarpGemmImpl<

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

247 

248 using WarpGemmMfma_f32_32x32x16_bf8_fp8 = WarpGemmImpl<

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

250 

251 using WarpGemmMfma_f32_32x32x16_bf8_bf8 = WarpGemmImpl<

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

253 

254 using WarpGemmMfma_f32_32x32x32_fp8_fp8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

255  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>,

256  2>>;

257 

258 using WarpGemmMfma_f32_32x32x32_bf8_bf8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

259  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>,

260  2>>;

261 

262 using WarpGemmMfma_f32_16x16x32_fp8_fp8 = WarpGemmImpl<

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

264 

265 using WarpGemmMfma_f32_16x16x32_fp8_fp8_CTransposed =

266  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

267  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>>>;

268 

269 using WarpGemmMfma_f32_16x16x32_bf8_bf8 = WarpGemmImpl<

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

271 

272 using WarpGemmMfma_f32_16x16x32_bf8_bf8_CTransposed =

273  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

274  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>>>;

275 

276 using WarpGemmMfma_f32_16x16x64_fp8_fp8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

277  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>,

278  2>>;

279 

280 using WarpGemmMfma_f32_16x16x64_bf8_bf8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

281  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>,

282  2>>;

283 

284 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

285 using WarpGemmMfma_f32_16x16x128_fp8_fp8 = WarpGemmImpl<

286  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8<WGAttrCtlEnum::Default_>,

287  AttrNumAccess>>;

288 

289 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

290 using WarpGemmMfma_f32_16x16x128_fp8_bf8 = WarpGemmImpl<

291  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8<WGAttrCtlEnum::Default_>,

292  AttrNumAccess>>;

293 

294 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

295 using WarpGemmMfma_f32_16x16x128_bf8_fp8 = WarpGemmImpl<

296  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8<WGAttrCtlEnum::Default_>,

297  AttrNumAccess>>;

298 

299 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

300 using WarpGemmMfma_f32_16x16x128_bf8_bf8 = WarpGemmImpl<

301  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8<WGAttrCtlEnum::Default_>,

302  AttrNumAccess>>;

303 

304 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

305 using WarpGemmMfma_f32_16x16x128_fp8_fp8_CTransposed =

306  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

307  WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8<WGAttrCtlEnum::Default_>,

308  AttrNumAccess>>;

309 

310 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

311 using WarpGemmMfma_f32_16x16x128_fp8_bf8_CTransposed =

312  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

313  WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8<WGAttrCtlEnum::Default_>,

314  AttrNumAccess>>;

315 

316 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

317 using WarpGemmMfma_f32_16x16x128_bf8_fp8_CTransposed =

318  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

319  WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8<WGAttrCtlEnum::Default_>,

320  AttrNumAccess>>;

321 

322 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

323 using WarpGemmMfma_f32_16x16x128_bf8_bf8_CTransposed =

324  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

325  WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8<WGAttrCtlEnum::Default_>,

326  AttrNumAccess>>;

327 

328 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

329 using WarpGemmMfma_f32_32x32x64_fp8_fp8 = WarpGemmImpl<

330  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_fp8<WGAttrCtlEnum::Default_>,

331  AttrNumAccess>>;

332 

333 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

334 using WarpGemmMfma_f32_32x32x64_fp8_bf8 = WarpGemmImpl<

335  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_bf8<WGAttrCtlEnum::Default_>,

336  AttrNumAccess>>;

337 

338 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

339 using WarpGemmMfma_f32_32x32x64_bf8_fp8 = WarpGemmImpl<

340  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_fp8<WGAttrCtlEnum::Default_>,

341  AttrNumAccess>>;

342 

343 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

344 using WarpGemmMfma_f32_32x32x64_bf8_bf8 = WarpGemmImpl<

345  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_bf8<WGAttrCtlEnum::Default_>,

346  AttrNumAccess>>;

347 

348 using WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed =

349  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

350  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>>>;

351 

352 using WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed =

353  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

354  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>>>;

355 

356 using WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed =

357  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

358  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8<WGAttrCtlEnum::Default_>>>;

359 

360 using WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed =

361  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

362  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>>>;

363 

364 template <index_t swizzle_factor = 2>

365 using WarpGemmMfmaFp8Fp8F32M32N32K16SwizzleBTransposedCDistribution =

366  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

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

368  2,

369  swizzle_factor>>;

370 

371 // int8

372 using WarpGemmMfma_i32_32x32x16_i8_i8 = WarpGemmImpl<

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

374 

375 using WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed =

376  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

377  WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;

378 

379 using WarpGemmMfma_i32_16x16x32_i8_i8 = WarpGemmImpl<

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

381 

382 using WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed =

383  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

384  WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;

385 

386 } // 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:1399

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1035

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1194

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:1666

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:1596

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16

Definition: warp_gemm_attribute_mfma_impl.hpp:537

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32

Definition: warp_gemm_attribute_mfma_impl.hpp:67

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16

Definition: warp_gemm_attribute_mfma_impl.hpp:920

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8

Definition: warp_gemm_attribute_mfma_impl.hpp:448

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4

Definition: warp_gemm_attribute_mfma_impl.hpp:625

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4

Definition: warp_gemm_attribute_mfma_impl.hpp:715

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16

Definition: warp_gemm_attribute_mfma_impl.hpp:193

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32

Definition: warp_gemm_attribute_mfma_impl.hpp:256

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16

Definition: warp_gemm_attribute_mfma_impl.hpp:806

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8

Definition: warp_gemm_attribute_mfma_impl.hpp:130

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4

Definition: warp_gemm_attribute_mfma_impl.hpp:319

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4

Definition: warp_gemm_attribute_mfma_impl.hpp:383

ck_tile::WarpGemmAttributeMfmaIterateK_SwizzleA

Definition: warp_gemm_attribute_mfma.hpp:845

ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB

Definition: warp_gemm_attribute_mfma.hpp:701

ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution

Definition: warp_gemm_attribute_mfma.hpp:551

ck_tile::WarpGemmAttributeMfmaIterateK

Definition: warp_gemm_attribute_mfma.hpp:107

ck_tile::WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB

Definition: warp_gemm_attribute_mfma.hpp:455

ck_tile::WarpGemmAttributeMfmaTransposedCDistribution

Definition: warp_gemm_attribute_mfma.hpp:371

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