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

• docs-7.0.0
• 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  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

19 

20 using WarpGemmMfmaF16F16F32M16N16K16 = WarpGemmImpl<

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

22 

23 #if defined(__gfx950__)

24 using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl<

25  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

26 

27 #else

28 using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

29  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

30  2>>;

31 #endif

32 

33 #if defined(__gfx950__)

34 using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl<

35  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>>>;

36 #else

37 using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

38  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>,

39  2>>;

40 #endif

41 

42 using WarpGemmMfmaF16F16F32M32N32K8SwizzleA = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK_SwizzleA<

43  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

44  1>>;

45 

46 using WarpGemmMfmaF16F16F32M32N32K16SwizzleA = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK_SwizzleA<

47  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

48  2>>;

49 

50 using WarpGemmMfmaF16F16F32M32N32K8TransposedCDistribution =

51  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

52  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

53 

54 using WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution =

55  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

56  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

57 

58 #if defined(__gfx950__)

59 using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution =

60  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

61  WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

62 #else

63 using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution =

64  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution<

65  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

66  2>>;

67 #endif

68 

69 #if defined(__gfx950__)

70 using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution =

71  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

72  WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>>>;

73 #else

74 using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution =

75  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution<

76  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>,

77  2>>;

78 #endif

79 

80 #if defined(__gfx950__)

81 using WarpGemmMfmaF16F16F32M16N16K32SwizzleBTransposedCDistribution =

82  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB<

83  WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,

84  1>>;

85 

86 using WarpGemmMfmaBf16Bf16F32M16N16K32SwizzleBTransposedCDistribution =

87  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB<

88  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,

89  1>>;

90 #endif

91 

92 #if defined(__gfx950__)

93 using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution =

94  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB<

95  WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

96 #else

97 using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution =

98  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

99  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,

100  2>>;

101 #endif

102 

103 using WarpGemmMfmaF16F16F32M4N64K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

104  WarpGemmAttributeMfmaImplF16F16F32M4N64K4<WGAttrCtlEnum::Default_>,

105  4>>;

106 

107 using WarpGemmMfmaF16F16F32M64N4K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

108  WarpGemmAttributeMfmaImplF16F16F32M64N4K4<WGAttrCtlEnum::Default_>,

109  4>>;

110 

111 // fp16 2:4 structured sparsity

112 using WarpGemmSmfmacF16F16F32M32N32K16 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<

113  WarpGemmAttributeSmfmacImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

114 

115 using WarpGemmSmfmacF16F16F32M16N16K32 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<

116  WarpGemmAttributeSmfmacImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>>>;

117 

118 // bf16

119 using WarpGemmMfmaBf16Bf16F32M32N32K8 = WarpGemmImpl<

120  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

121 

122 using WarpGemmMfmaBf16Bf16F32M16N16K16 = WarpGemmImpl<

123  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

124 

125 #if defined(__gfx950__)

126 using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl<

127  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

128 

129 #else

130 using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

131  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

132  2>>;

133 #endif

134 

135 #if defined(__gfx950__)

136 using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl<

137  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>>>;

138 #else

139 using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

140  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,

141  2>>;

142 #endif

143 

144 using WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleA = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK_SwizzleA<

145  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

146  1>>;

147 

148 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleA =

149  WarpGemmImpl<WarpGemmAtrributeMfmaIterateK_SwizzleA<

150  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

151  2>>;

152 

153 using WarpGemmMfmaBf16Bf16F32M32N32K8TransposedCDistribution =

154  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

155  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

156 

157 using WarpGemmMfmaBf16Bf16F32M16N16K16TransposedCDistribution =

158  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

159  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>>>;

160 

161 #if defined(__gfx950__)

162 using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution =

163  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

164  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>>>;

165 #else

166 using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution =

167  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution<

168  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

169  2>>;

170 #endif

171 

172 #if defined(__gfx950__)

173 using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution =

174  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

175  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>>>;

176 #else

177 using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution =

178  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution<

179  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,

180  2>>;

181 #endif

182 

183 #if defined(__gfx950__)

184 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution =

185  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB<

186  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;

187 #else

188 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution =

189  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

190  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,

191  2>>;

192 #endif

193 

194 using WarpGemmMfmaBf16Bf16F32M4N64K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

195  WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4<WGAttrCtlEnum::Default_>,

196  4>>;

197 

198 using WarpGemmMfmaBf16Bf16F32M64N4K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

199  WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4<WGAttrCtlEnum::Default_>,

200  4>>;

201 

202 // fp8

203 

204 using WarpGemmMfma_f32_32x32x16_fp8_fp8 = WarpGemmImpl<

205  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>>>;

206 

207 using WarpGemmMfma_f32_32x32x16_fp8_bf8 = WarpGemmImpl<

208  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>>>;

209 

210 using WarpGemmMfma_f32_32x32x16_bf8_fp8 = WarpGemmImpl<

211  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8<WGAttrCtlEnum::Default_>>>;

212 

213 using WarpGemmMfma_f32_32x32x16_bf8_bf8 = WarpGemmImpl<

214  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>>>;

215 

216 using WarpGemmMfma_f32_32x32x32_fp8_fp8 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

217  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>,

218  2>>;

219 

220 using WarpGemmMfma_f32_32x32x32_bf8_bf8 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

221  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>,

222  2>>;

223 

224 using WarpGemmMfma_f32_16x16x32_fp8_fp8 = WarpGemmImpl<

225  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>>>;

226 

227 using WarpGemmMfma_f32_16x16x32_bf8_bf8 = WarpGemmImpl<

228  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>>>;

229 

230 using WarpGemmMfma_f32_16x16x64_fp8_fp8 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

231  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>,

232  2>>;

233 

234 using WarpGemmMfma_f32_16x16x64_bf8_bf8 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<

235  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>,

236  2>>;

237 

238 using WarpGemmMfma_f32_16x16x128_fp8_fp8 = WarpGemmImpl<WarpGemmAtrributeMfma<

239  WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8<WGAttrCtlEnum::Default_>>>;

240 

241 using WarpGemmMfma_f32_16x16x128_fp8_bf8 = WarpGemmImpl<WarpGemmAtrributeMfma<

242  WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8<WGAttrCtlEnum::Default_>>>;

243 

244 using WarpGemmMfma_f32_16x16x128_bf8_fp8 = WarpGemmImpl<WarpGemmAtrributeMfma<

245  WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8<WGAttrCtlEnum::Default_>>>;

246 

247 using WarpGemmMfma_f32_16x16x128_bf8_bf8 = WarpGemmImpl<WarpGemmAtrributeMfma<

248  WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8<WGAttrCtlEnum::Default_>>>;

249 

250 using WarpGemmMfma_f32_32x32x64_fp8_fp8 = WarpGemmImpl<

251  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_fp8<WGAttrCtlEnum::Default_>>>;

252 

253 using WarpGemmMfma_f32_32x32x64_fp8_bf8 = WarpGemmImpl<

254  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_bf8<WGAttrCtlEnum::Default_>>>;

255 

256 using WarpGemmMfma_f32_32x32x64_bf8_fp8 = WarpGemmImpl<

257  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_fp8<WGAttrCtlEnum::Default_>>>;

258 

259 using WarpGemmMfma_f32_32x32x64_bf8_bf8 = WarpGemmImpl<

260  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_bf8<WGAttrCtlEnum::Default_>>>;

261 

262 using WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed =

263  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

264  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>>>;

265 

266 using WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed =

267  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

268  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>>>;

269 

270 using WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed =

271  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

272  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8<WGAttrCtlEnum::Default_>>>;

273 

274 using WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed =

275  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

276  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>>>;

277 

278 template <index_t swizzle_factor = 2>

279 using WarpGemmMfmaFp8Fp8F32M32N32K16SwizzleBTransposedCDistribution =

280  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution_SwizzleB<

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

282  2,

283  swizzle_factor>>;

284 

285 // int8

286 using WarpGemmMfma_i32_32x32x16_i8_i8 = WarpGemmImpl<

287  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;

288 

289 using WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed =

290  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

291  WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;

292 

293 using WarpGemmMfma_i32_16x16x32_i8_i8 = WarpGemmImpl<

294  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;

295 

296 using WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed =

297  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<

298  WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;

299 

300 } // namespace ck_tile

core.hpp

ck_tile

Definition: cluster_descriptor.hpp:13

ck_tile::WarpGemmAtrributeMfma

Definition: warp_gemm_attribute_mfma.hpp:13

ck_tile::WarpGemmAtrributeMfmaIterateK_SwizzleA

Definition: warp_gemm_attribute_mfma.hpp:814

ck_tile::WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution_SwizzleB

Definition: warp_gemm_attribute_mfma.hpp:670

ck_tile::WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution

Definition: warp_gemm_attribute_mfma.hpp:455

ck_tile::WarpGemmAtrributeMfmaIterateK

Definition: warp_gemm_attribute_mfma.hpp:78

ck_tile::WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB

Definition: warp_gemm_attribute_mfma.hpp:361

ck_tile::WarpGemmAtrributeMfmaTransposedCDistribution

Definition: warp_gemm_attribute_mfma.hpp:294

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1347

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:983

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1142

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:1614

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:1544

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

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

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

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