/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 // 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 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_16x16x32_fp8_bf8 = WarpGemmImpl<

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

268 

269 using WarpGemmMfma_f32_32x32x16_bf8_fp8 = WarpGemmImpl<

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

271 

272 using WarpGemmMfma_f32_32x32x16_bf8_bf8 = WarpGemmImpl<

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

274 

275 using WarpGemmMfma_f32_32x32x32_fp8_fp8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

276  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>,

277  2>>;

278 

279 using WarpGemmMfma_f32_32x32x32_bf8_bf8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

280  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>,

281  2>>;

282 

283 using WarpGemmMfma_f32_32x32x32_fp8_bf8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

284  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>,

285  2>>;

286 

287 using WarpGemmMfma_f32_16x16x32_fp8_fp8 = WarpGemmImpl<

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

289 

290 using WarpGemmMfma_f32_16x16x32_fp8_fp8_CTransposed =

291  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

292  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>>>;

293 

294 using WarpGemmMfma_f32_16x16x32_bf8_bf8 = WarpGemmImpl<

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

296 

297 using WarpGemmMfma_f32_16x16x32_bf8_bf8_CTransposed =

298  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

299  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>>>;

300 

301 using WarpGemmMfma_f32_16x16x64_fp8_fp8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

302  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>,

303  2>>;

304 

305 using WarpGemmMfma_f32_16x16x64_bf8_bf8 = WarpGemmImpl<WarpGemmAttributeMfmaIterateK<

306  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>,

307  2>>;

308 

309 using WarpGemmMfma_f32_16x16x64_fp8_fp8_CTransposed =

310  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

311  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>,

312  2>>;

313 

314 using WarpGemmMfma_f32_16x16x64_bf8_bf8_CTransposed =

315  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution<

316  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>,

317  2>>;

318 

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

320 using WarpGemmMfma_f32_16x16x128_f8f6f4 = WarpGemmImpl<

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

322 

323 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

324 using WarpGemmMfma_f32_16x16x128_fp8_fp8_CTransposed =

325  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

326  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8f6f4<fp8_t, fp8_t>,

327  AttrNumAccess>>;

328 

329 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

330 using WarpGemmMfma_f32_16x16x128_fp8_bf8_CTransposed =

331  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

332  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8f6f4<fp8_t, bf8_t>,

333  AttrNumAccess>>;

334 

335 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

336 using WarpGemmMfma_f32_16x16x128_bf8_fp8_CTransposed =

337  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

338  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8f6f4<bf8_t, fp8_t>,

339  AttrNumAccess>>;

340 

341 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

342 using WarpGemmMfma_f32_16x16x128_bf8_bf8_CTransposed =

343  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

344  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8f6f4<bf8_t, bf8_t>,

345  AttrNumAccess>>;

346 

347 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

348 using WarpGemmMfma_f32_32x32x64_fp8_fp8 = WarpGemmImpl<

349  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_fp8<WGAttrCtlEnum::Default_>,

350  AttrNumAccess>>;

351 

352 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

353 using WarpGemmMfma_f32_32x32x64_fp8_bf8 = WarpGemmImpl<

354  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_bf8<WGAttrCtlEnum::Default_>,

355  AttrNumAccess>>;

356 

357 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

358 using WarpGemmMfma_f32_32x32x64_bf8_fp8 = WarpGemmImpl<

359  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_fp8<WGAttrCtlEnum::Default_>,

360  AttrNumAccess>>;

361 

362 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>

363 using WarpGemmMfma_f32_32x32x64_bf8_bf8 = WarpGemmImpl<

364  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_bf8<WGAttrCtlEnum::Default_>,

365  AttrNumAccess>>;

366 

367 using WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed =

368  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

369  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>>>;

370 

371 using WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed =

372  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

373  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>>>;

374 

375 using WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed =

376  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

377  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8<WGAttrCtlEnum::Default_>>>;

378 

379 using WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed =

380  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

381  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>>>;

382 

383 template <index_t swizzle_factor = 2>

384 using WarpGemmMfmaFp8Fp8F32M32N32K32SwizzleBTransposedCDistribution =

385  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

386  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>,

387  2,

388  swizzle_factor>>;

389 

390 // int8

391 using WarpGemmMfma_i32_32x32x16_i8_i8 = WarpGemmImpl<

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

393 

394 using WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed =

395  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

396  WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;

397 

398 using WarpGemmMfma_i32_16x16x32_i8_i8 = WarpGemmImpl<

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

400 

401 using WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed =

402  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<

403  WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;

404 

405 } // 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::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