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

• 3643
• include
• ck_tile
• ops
• gemm
• warp

warp_gemm.hpp

Go to the documentation of this file.

1 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.

2 // SPDX-License-Identifier: MIT

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 WarpGemmMfmaF32F32F32M16N16K8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

28  WarpGemmAttributeMfmaImplF32F32F32M16N16K4<WGAttrCtlEnum::Default_>,

29  2,

30  AttrNumAccess>>;

31 

32 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

33 using WarpGemmMfmaF32F32F32M32N32K8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

34  WarpGemmAttributeMfmaImplF32F32F32M32N32K2<WGAttrCtlEnum::Default_>,

35  4,

36  AttrNumAccess>>;

37 

38 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

39 using WarpGemmMfmaF32F32F32M16N16K16TransposedCDistribution =

40  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

41  WarpGemmAttributeMfmaImplF32F32F32M16N16K4<WGAttrCtlEnum::Default_>,

42  4,

43  AttrNumAccess>>;

44 

45 // fp16

46 

47 using WarpGemmMfmaF16F16F32M32N32K8 = WarpGemmImpl<

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

49 

50 using WarpGemmMfmaF16F16F32M16N16K16 = WarpGemmImpl<

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

52 

53 #if defined(__gfx950__)

54 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

55 using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl<

56  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>,

57  AttrNumAccess>>;

58 #else

59 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

60 using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

61  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

62  2,

63  AttrNumAccess>>;

64 #endif

65 

66 #if defined(__gfx950__)

67 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

68 using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl<

69  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,

70  AttrNumAccess>>;

71 #else

72 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

73 using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

74  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>,

75  2,

76  AttrNumAccess>>;

77 #endif

78 

79 using WarpGemmMfmaF16F16F32M32N32K8SwizzleA = WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

80  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

81  1>>;

82 

83 using WarpGemmMfmaF16F16F32M32N32K16SwizzleA = WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

84  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

85  2>>;

86 

87 using WarpGemmMfmaF16F16F32M32N32K8TransposedCDistribution =

88  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

89  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

90 

91 using WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution =

92  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

93  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

94 

95 #if defined(__gfx950__)

96 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

97 using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution =

98  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

99  WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>,

100  AttrNumAccess>>;

101 #else

102 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

103 using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution =

104  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

105  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

106  2,

107  AttrNumAccess>>;

108 #endif

109 

110 #if defined(__gfx950__)

111 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

112 using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution =

113  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

114  WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,

115  AttrNumAccess>>;

116 #else

117 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

118 using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution =

119  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

120  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>,

121  2,

122  AttrNumAccess>>;

123 #endif

124 

125 #if defined(__gfx950__)

126 using WarpGemmMfmaF16F16F32M16N16K32SwizzleBTransposedCDistribution =

127  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

128  WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,

129  1>>;

130 

131 using WarpGemmMfmaBf16Bf16F32M16N16K32SwizzleBTransposedCDistribution =

132  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

133  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,

134  1>>;

135 #endif

136 

137 using WarpGemmMfmaF16F16F32M32N32K8SwizzleBTransposedCDistribution =

138  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

139  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

140 

141 #if defined(__gfx950__)

142 using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution =

143  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

144  WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

145 #else

146 using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution =

147  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

148  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

149  2>>;

150 #endif

151 

152 using WarpGemmMfmaF16F16F32M4N64K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

153  WarpGemmAttributeMfmaImplF16F16F32M4N64K4<WGAttrCtlEnum::Default_>,

154  4>>;

155 

156 using WarpGemmMfmaF16F16F32M64N4K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

157  WarpGemmAttributeMfmaImplF16F16F32M64N4K4<WGAttrCtlEnum::Default_>,

158  4>>;

159 

160 // fp16 2:4 structured sparsity

161 using WarpGemmSmfmacF16F16F32M32N32K16 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<

162  WarpGemmAttributeSmfmacImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

163 

164 using WarpGemmSmfmacF16F16F32M16N16K32 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<

165  WarpGemmAttributeSmfmacImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>>>;

166 

167 // bf16

168 using WarpGemmMfmaBf16Bf16F32M32N32K8 = WarpGemmImpl<

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

170 

171 using WarpGemmMfmaBf16Bf16F32M16N16K16 = WarpGemmImpl<

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

173 

174 #if defined(__gfx950__)

175 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

176 using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl<

177  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>,

178  AttrNumAccess>>;

179 #else

180 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

181 using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

182  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

183  2,

184  AttrNumAccess>>;

185 #endif

186 

187 #if defined(__gfx950__)

188 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

189 using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl<

190  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,

191  AttrNumAccess>>;

192 #else

193 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

194 using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

195  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,

196  2,

197  AttrNumAccess>>;

198 #endif

199 

200 using WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleA = WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

201  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

202  1>>;

203 

204 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleA =

205  WarpGemmImpl<WarpGemmAttributeMfmaIterateK_SwizzleA<

206  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

207  2>>;

208 

209 using WarpGemmMfmaBf16Bf16F32M32N32K8TransposedCDistribution =

210  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

211  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

212 

213 using WarpGemmMfmaBf16Bf16F32M16N16K16TransposedCDistribution =

214  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

215  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

216 

217 #if defined(__gfx950__)

218 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

219 using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution =

220  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

221  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>,

222  AttrNumAccess>>;

223 #else

224 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

225 using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution =

226  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

227  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

228  2,

229  AttrNumAccess>>;

230 #endif

231 

232 #if defined(__gfx950__)

233 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

234 using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution =

235  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

236  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,

237  AttrNumAccess>>;

238 #else

239 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

240 using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution =

241  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

242  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,

243  2,

244  AttrNumAccess>>;

245 #endif

246 

247 using WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleBTransposedCDistribution =

248  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

249  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

250 

251 #if defined(__gfx950__)

252 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution =

253  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB<

254  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

255 #else

256 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution =

257  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

258  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

259  2>>;

260 #endif

261 

262 using WarpGemmMfmaBf16Bf16F32M4N64K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

263  WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4<WGAttrCtlEnum::Default_>,

264  4>>;

265 

266 using WarpGemmMfmaBf16Bf16F32M64N4K16 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

267  WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4<WGAttrCtlEnum::Default_>,

268  4>>;

269 

270 // fp8

271 

272 using WarpGemmMfma_f32_32x32x16_fp8_fp8 = WarpGemmImpl<

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

274 

275 using WarpGemmMfma_f32_32x32x16_fp8_bf8 = WarpGemmImpl<

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

277 

278 using WarpGemmMfma_f32_16x16x32_fp8_bf8 = WarpGemmImpl<

279  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_bf8<WGAttrCtlEnum::Default_>>>;

280 

281 using WarpGemmMfma_f32_32x32x16_bf8_fp8 = WarpGemmImpl<

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

283 

284 using WarpGemmMfma_f32_32x32x16_bf8_bf8 = WarpGemmImpl<

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

286 

287 using WarpGemmMfma_f32_32x32x32_fp8_fp8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

288  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>,

289  2>>;

290 

291 using WarpGemmMfma_f32_32x32x32_bf8_bf8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

292  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>,

293  2>>;

294 

295 using WarpGemmMfma_f32_32x32x32_fp8_bf8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

296  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>,

297  2>>;

298 

299 using WarpGemmMfma_f32_16x16x32_fp8_fp8 = WarpGemmImpl<

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

301 

302 using WarpGemmMfma_f32_16x16x32_fp8_fp8_CTransposed =

303  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

304  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>>>;

305 

306 using WarpGemmMfma_f32_16x16x32_bf8_bf8 = WarpGemmImpl<

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

308 

309 using WarpGemmMfma_f32_16x16x32_bf8_bf8_CTransposed =

310  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

311  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>>>;

312 

313 using WarpGemmMfma_f32_16x16x64_fp8_fp8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

314  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>,

315  2>>;

316 

317 using WarpGemmMfma_f32_16x16x64_bf8_bf8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

318  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>,

319  2>>;

320 

321 using WarpGemmMfma_f32_16x16x64_fp8_fp8_CTransposed =

322  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

323  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>,

324  2>>;

325 

326 using WarpGemmMfma_f32_16x16x64_bf8_bf8_CTransposed =

327  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

328  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>,

329  2>>;

330 

331 template <typename A, typename B, WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

332 using WarpGemmMfma_f32_16x16x128_f8f6f4 = WarpGemmImpl<

333  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_f8f6f4<A, B>, AttrNumAccess>>;

334 

335 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

336 using WarpGemmMfma_f32_16x16x128_fp8_fp8_CTransposed =

337  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

338  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8f6f4<fp8_t, fp8_t>,

339  AttrNumAccess>>;

340 

341 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

342 using WarpGemmMfma_f32_16x16x128_fp8_bf8_CTransposed =

343  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

344  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8f6f4<fp8_t, bf8_t>,

345  AttrNumAccess>>;

346 

347 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

348 using WarpGemmMfma_f32_16x16x128_bf8_fp8_CTransposed =

349  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

350  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8f6f4<bf8_t, fp8_t>,

351  AttrNumAccess>>;

352 

353 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

354 using WarpGemmMfma_f32_16x16x128_bf8_bf8_CTransposed =

355  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

356  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8f6f4<bf8_t, bf8_t>,

357  AttrNumAccess>>;

358 

359 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

360 using WarpGemmMfma_f32_32x32x64_fp8_fp8 = WarpGemmImpl<

361  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_fp8<WGAttrCtlEnum::Default_>,

362  AttrNumAccess>>;

363 

364 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

365 using WarpGemmMfma_f32_32x32x64_fp8_bf8 = WarpGemmImpl<

366  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_bf8<WGAttrCtlEnum::Default_>,

367  AttrNumAccess>>;

368 

369 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

370 using WarpGemmMfma_f32_32x32x64_bf8_fp8 = WarpGemmImpl<

371  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_fp8<WGAttrCtlEnum::Default_>,

372  AttrNumAccess>>;

373 

374 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

375 using WarpGemmMfma_f32_32x32x64_bf8_bf8 = WarpGemmImpl<

376  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_bf8<WGAttrCtlEnum::Default_>,

377  AttrNumAccess>>;

378 

379 using WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed =

380  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

381  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>>>;

382 

383 using WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed =

384  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

385  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>>>;

386 

387 using WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed =

388  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

389  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8<WGAttrCtlEnum::Default_>>>;

390 

391 using WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed =

392  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

393  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>>>;

394 

395 template <index_t swizzle_factor = 2>

396 using WarpGemmMfmaFp8Fp8F32M32N32K32SwizzleBTransposedCDistribution =

397  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

398  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>,

399  2,

400  swizzle_factor>>;

401 

402 // int8

403 using WarpGemmMfma_i32_32x32x16_i8_i8 = WarpGemmImpl<

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

405 

406 using WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed =

407  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

408  WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;

409 

410 using WarpGemmMfma_i32_16x16x32_i8_i8 = WarpGemmImpl<

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

412 

413 using WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed =

414  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

415  WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;

416 

417 } // 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_f8f6f4

Definition: warp_gemm_attribute_mfma_impl.hpp:1531

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:1808

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:1738

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::WarpGemmAttributeMfmaImplF32F32F32M32N32K2

Definition: warp_gemm_attribute_mfma_impl.hpp:131

ck_tile::WarpGemmAttributeMfmaIterateK_SwizzleA

Definition: warp_gemm_attribute_mfma.hpp:708

ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB

Definition: warp_gemm_attribute_mfma.hpp:581

ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution

Definition: warp_gemm_attribute_mfma.hpp:461

ck_tile::WarpGemmAttributeMfmaIterateK

Definition: warp_gemm_attribute_mfma.hpp:126

ck_tile::WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB

Definition: warp_gemm_attribute_mfma.hpp:365

ck_tile::WarpGemmAttributeMfmaTransposedCDistribution

Definition: warp_gemm_attribute_mfma.hpp:306

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