/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

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

Composable Kernel: /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
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_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 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
310 using WarpGemmMfma_f32_16x16x128_fp4 = WarpGemmImpl<
311  WarpGemmAttributeMfma<WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4<WGAttrCtlEnum::Default_>,
312  AttrNumAccess>>;
313 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
314 using WarpGemmMfma_f32_16x16x128_fp8_fp8 = WarpGemmImpl<
315  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8<WGAttrCtlEnum::Default_>,
316  AttrNumAccess>>;
317 
318 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
319 using WarpGemmMfma_f32_16x16x128_fp8_bf8 = WarpGemmImpl<
320  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8<WGAttrCtlEnum::Default_>,
321  AttrNumAccess>>;
322 
323 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
324 using WarpGemmMfma_f32_16x16x128_bf8_fp8 = WarpGemmImpl<
325  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8<WGAttrCtlEnum::Default_>,
326  AttrNumAccess>>;
327 
328 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
329 using WarpGemmMfma_f32_16x16x128_bf8_bf8 = WarpGemmImpl<
330  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8<WGAttrCtlEnum::Default_>,
331  AttrNumAccess>>;
332 
333 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
334 using WarpGemmMfma_f32_16x16x128_fp8_fp8_CTransposed =
335  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<
336  WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8<WGAttrCtlEnum::Default_>,
337  AttrNumAccess>>;
338 
339 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
340 using WarpGemmMfma_f32_16x16x128_fp8_bf8_CTransposed =
341  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<
342  WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8<WGAttrCtlEnum::Default_>,
343  AttrNumAccess>>;
344 
345 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
346 using WarpGemmMfma_f32_16x16x128_bf8_fp8_CTransposed =
347  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<
348  WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8<WGAttrCtlEnum::Default_>,
349  AttrNumAccess>>;
350 
351 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
352 using WarpGemmMfma_f32_16x16x128_bf8_bf8_CTransposed =
353  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<
354  WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8<WGAttrCtlEnum::Default_>,
355  AttrNumAccess>>;
356 
357 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
358 using WarpGemmMfma_f32_32x32x64_fp8_fp8 = WarpGemmImpl<
359  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_fp8<WGAttrCtlEnum::Default_>,
360  AttrNumAccess>>;
361 
362 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
363 using WarpGemmMfma_f32_32x32x64_fp8_bf8 = WarpGemmImpl<
364  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_bf8<WGAttrCtlEnum::Default_>,
365  AttrNumAccess>>;
366 
367 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
368 using WarpGemmMfma_f32_32x32x64_bf8_fp8 = WarpGemmImpl<
369  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_fp8<WGAttrCtlEnum::Default_>,
370  AttrNumAccess>>;
371 
372 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
373 using WarpGemmMfma_f32_32x32x64_bf8_bf8 = WarpGemmImpl<
374  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_bf8<WGAttrCtlEnum::Default_>,
375  AttrNumAccess>>;
376 
377 using WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed =
378  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<
379  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>>>;
380 
381 using WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed =
382  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<
383  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>>>;
384 
385 using WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed =
386  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<
387  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8<WGAttrCtlEnum::Default_>>>;
388 
389 using WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed =
390  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<
391  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>>>;
392 
393 template <index_t swizzle_factor = 2>
394 using WarpGemmMfmaFp8Fp8F32M32N32K32SwizzleBTransposedCDistribution =
395  WarpGemmImpl<WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB<
396  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>,
397  2,
398  swizzle_factor>>;
399 
400 // int8
401 using WarpGemmMfma_i32_32x32x16_i8_i8 = WarpGemmImpl<
402  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;
403 
404 using WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed =
405  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<
406  WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;
407 
408 using WarpGemmMfma_i32_16x16x32_i8_i8 = WarpGemmImpl<
409  WarpGemmAttributeMfma<WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;
410 
411 using WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed =
412  WarpGemmImpl<WarpGemmAttributeMfmaTransposedCDistribution<
413  WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;
414 
415 } // namespace ck_tile
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: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:1890
ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8
Definition: warp_gemm_attribute_mfma_impl.hpp:1820
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:869
ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB
Definition: warp_gemm_attribute_mfma.hpp:725
ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution
Definition: warp_gemm_attribute_mfma.hpp:575
ck_tile::WarpGemmAttributeMfmaIterateK
Definition: warp_gemm_attribute_mfma.hpp:130
ck_tile::WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB
Definition: warp_gemm_attribute_mfma.hpp:479
ck_tile::WarpGemmAttributeMfmaTransposedCDistribution
Definition: warp_gemm_attribute_mfma.hpp:395
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