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

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-7.1.0/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/docs-7.1.0/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 // 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 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
25 using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl<
26  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>,
27  AttrNumAccess>>;
28 #else
29 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
30 using WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
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  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,
40  AttrNumAccess>>;
41 #else
42 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
43 using WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
44  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>,
45  2,
46  AttrNumAccess>>;
47 #endif
48 
49 using WarpGemmMfmaF16F16F32M32N32K8SwizzleA = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK_SwizzleA<
50  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,
51  1>>;
52 
53 using WarpGemmMfmaF16F16F32M32N32K16SwizzleA = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK_SwizzleA<
54  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,
55  2>>;
56 
57 using WarpGemmMfmaF16F16F32M32N32K8TransposedCDistribution =
58  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
59  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>>>;
60 
61 using WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution =
62  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
63  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>>>;
64 
65 #if defined(__gfx950__)
66 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
67 using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution =
68  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
69  WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>,
70  AttrNumAccess>>;
71 #else
72 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
73 using WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution =
74  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution<
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<WarpGemmAtrributeMfmaTransposedCDistribution<
84  WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,
85  AttrNumAccess>>;
86 #else
87 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
88 using WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution =
89  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution<
90  WarpGemmAttributeMfmaImplF16F16F32M16N16K16<WGAttrCtlEnum::Default_>,
91  2,
92  AttrNumAccess>>;
93 #endif
94 
95 #if defined(__gfx950__)
96 using WarpGemmMfmaF16F16F32M16N16K32SwizzleBTransposedCDistribution =
97  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB<
98  WarpGemmAttributeMfmaImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>,
99  1>>;
100 
101 using WarpGemmMfmaBf16Bf16F32M16N16K32SwizzleBTransposedCDistribution =
102  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB<
103  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,
104  1>>;
105 #endif
106 
107 #if defined(__gfx950__)
108 using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution =
109  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB<
110  WarpGemmAttributeMfmaImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;
111 #else
112 using WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution =
113  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution_SwizzleB<
114  WarpGemmAttributeMfmaImplF16F16F32M32N32K8<WGAttrCtlEnum::Default_>,
115  2>>;
116 #endif
117 
118 using WarpGemmMfmaF16F16F32M4N64K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
119  WarpGemmAttributeMfmaImplF16F16F32M4N64K4<WGAttrCtlEnum::Default_>,
120  4>>;
121 
122 using WarpGemmMfmaF16F16F32M64N4K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
123  WarpGemmAttributeMfmaImplF16F16F32M64N4K4<WGAttrCtlEnum::Default_>,
124  4>>;
125 
126 // fp16 2:4 structured sparsity
127 using WarpGemmSmfmacF16F16F32M32N32K16 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<
128  WarpGemmAttributeSmfmacImplF16F16F32M32N32K16<WGAttrCtlEnum::Default_>>>;
129 
130 using WarpGemmSmfmacF16F16F32M16N16K32 = WarpGemmSmfmacImpl<WarpGemmAttributeSmfmac<
131  WarpGemmAttributeSmfmacImplF16F16F32M16N16K32<WGAttrCtlEnum::Default_>>>;
132 
133 // bf16
134 using WarpGemmMfmaBf16Bf16F32M32N32K8 = WarpGemmImpl<
135  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;
136 
137 using WarpGemmMfmaBf16Bf16F32M16N16K16 = WarpGemmImpl<
138  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>>>;
139 
140 #if defined(__gfx950__)
141 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
142 using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl<
143  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>,
144  AttrNumAccess>>;
145 #else
146 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
147 using WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
148  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,
149  2,
150  AttrNumAccess>>;
151 #endif
152 
153 #if defined(__gfx950__)
154 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
155 using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl<
156  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,
157  AttrNumAccess>>;
158 #else
159 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
160 using WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
161  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,
162  2,
163  AttrNumAccess>>;
164 #endif
165 
166 using WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleA = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK_SwizzleA<
167  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,
168  1>>;
169 
170 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleA =
171  WarpGemmImpl<WarpGemmAtrributeMfmaIterateK_SwizzleA<
172  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,
173  2>>;
174 
175 using WarpGemmMfmaBf16Bf16F32M32N32K8TransposedCDistribution =
176  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
177  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>>>;
178 
179 using WarpGemmMfmaBf16Bf16F32M16N16K16TransposedCDistribution =
180  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
181  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>>>;
182 
183 #if defined(__gfx950__)
184 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
185 using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution =
186  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
187  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>,
188  AttrNumAccess>>;
189 #else
190 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
191 using WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution =
192  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution<
193  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,
194  2,
195  AttrNumAccess>>;
196 #endif
197 
198 #if defined(__gfx950__)
199 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
200 using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution =
201  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
202  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32<WGAttrCtlEnum::Default_>,
203  AttrNumAccess>>;
204 #else
205 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
206 using WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution =
207  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution<
208  WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16<WGAttrCtlEnum::Default_>,
209  2,
210  AttrNumAccess>>;
211 #endif
212 
213 #if defined(__gfx950__)
214 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution =
215  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB<
216  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16<WGAttrCtlEnum::Default_>>>;
217 #else
218 using WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution =
219  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution_SwizzleB<
220  WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8<WGAttrCtlEnum::Default_>,
221  2>>;
222 #endif
223 
224 using WarpGemmMfmaBf16Bf16F32M4N64K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
225  WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4<WGAttrCtlEnum::Default_>,
226  4>>;
227 
228 using WarpGemmMfmaBf16Bf16F32M64N4K16 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
229  WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4<WGAttrCtlEnum::Default_>,
230  4>>;
231 
232 // fp8
233 
234 using WarpGemmMfma_f32_32x32x16_fp8_fp8 = WarpGemmImpl<
235  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>>>;
236 
237 using WarpGemmMfma_f32_32x32x16_fp8_bf8 = WarpGemmImpl<
238  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>>>;
239 
240 using WarpGemmMfma_f32_32x32x16_bf8_fp8 = WarpGemmImpl<
241  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8<WGAttrCtlEnum::Default_>>>;
242 
243 using WarpGemmMfma_f32_32x32x16_bf8_bf8 = WarpGemmImpl<
244  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>>>;
245 
246 using WarpGemmMfma_f32_32x32x32_fp8_fp8 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
247  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>,
248  2>>;
249 
250 using WarpGemmMfma_f32_32x32x32_bf8_bf8 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
251  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>,
252  2>>;
253 
254 using WarpGemmMfma_f32_16x16x32_fp8_fp8 = WarpGemmImpl<
255  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>>>;
256 
257 using WarpGemmMfma_f32_16x16x32_bf8_bf8 = WarpGemmImpl<
258  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>>>;
259 
260 using WarpGemmMfma_f32_16x16x64_fp8_fp8 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
261  WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8<WGAttrCtlEnum::Default_>,
262  2>>;
263 
264 using WarpGemmMfma_f32_16x16x64_bf8_bf8 = WarpGemmImpl<WarpGemmAtrributeMfmaIterateK<
265  WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8<WGAttrCtlEnum::Default_>,
266  2>>;
267 
268 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
269 using WarpGemmMfma_f32_16x16x128_fp8_fp8 = WarpGemmImpl<
270  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8<WGAttrCtlEnum::Default_>,
271  AttrNumAccess>>;
272 
273 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
274 using WarpGemmMfma_f32_16x16x128_fp8_bf8 = WarpGemmImpl<
275  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8<WGAttrCtlEnum::Default_>,
276  AttrNumAccess>>;
277 
278 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
279 using WarpGemmMfma_f32_16x16x128_bf8_fp8 = WarpGemmImpl<
280  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8<WGAttrCtlEnum::Default_>,
281  AttrNumAccess>>;
282 
283 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
284 using WarpGemmMfma_f32_16x16x128_bf8_bf8 = WarpGemmImpl<
285  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8<WGAttrCtlEnum::Default_>,
286  AttrNumAccess>>;
287 
288 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
289 using WarpGemmMfma_f32_32x32x64_fp8_fp8 = WarpGemmImpl<
290  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_fp8<WGAttrCtlEnum::Default_>,
291  AttrNumAccess>>;
292 
293 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
294 using WarpGemmMfma_f32_32x32x64_fp8_bf8 = WarpGemmImpl<
295  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_bf8<WGAttrCtlEnum::Default_>,
296  AttrNumAccess>>;
297 
298 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
299 using WarpGemmMfma_f32_32x32x64_bf8_fp8 = WarpGemmImpl<
300  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_fp8<WGAttrCtlEnum::Default_>,
301  AttrNumAccess>>;
302 
303 template <WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
304 using WarpGemmMfma_f32_32x32x64_bf8_bf8 = WarpGemmImpl<
305  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_bf8<WGAttrCtlEnum::Default_>,
306  AttrNumAccess>>;
307 
308 using WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed =
309  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
310  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8<WGAttrCtlEnum::Default_>>>;
311 
312 using WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed =
313  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
314  WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8<WGAttrCtlEnum::Default_>>>;
315 
316 using WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed =
317  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
318  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8<WGAttrCtlEnum::Default_>>>;
319 
320 using WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed =
321  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
322  WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8<WGAttrCtlEnum::Default_>>>;
323 
324 template <index_t swizzle_factor = 2>
325 using WarpGemmMfmaFp8Fp8F32M32N32K16SwizzleBTransposedCDistribution =
326  WarpGemmImpl<WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution_SwizzleB<
327  WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base<fp8_t, fp8_t, WGAttrCtlEnum::Default_>,
328  2,
329  swizzle_factor>>;
330 
331 // int8
332 using WarpGemmMfma_i32_32x32x16_i8_i8 = WarpGemmImpl<
333  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;
334 
335 using WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed =
336  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
337  WarpGemmAttributeMfmaImpl_i32_32x32x16_i8<WGAttrCtlEnum::Default_>>>;
338 
339 using WarpGemmMfma_i32_16x16x32_i8_i8 = WarpGemmImpl<
340  WarpGemmAtrributeMfma<WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;
341 
342 using WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed =
343  WarpGemmImpl<WarpGemmAtrributeMfmaTransposedCDistribution<
344  WarpGemmAttributeMfmaImpl_i32_16x16x32_i8<WGAttrCtlEnum::Default_>>>;
345 
346 } // namespace ck_tile
ck_tile
Definition: cluster_descriptor.hpp:13
ck_tile::WarpGemmAtrributeMfma
Definition: warp_gemm_attribute_mfma.hpp:23
ck_tile::WarpGemmAtrributeMfmaIterateK_SwizzleA
Definition: warp_gemm_attribute_mfma.hpp:844
ck_tile::WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution_SwizzleB
Definition: warp_gemm_attribute_mfma.hpp:700
ck_tile::WarpGemmAtrributeMfmaIterateKAndTransposedCDistribution
Definition: warp_gemm_attribute_mfma.hpp:550
ck_tile::WarpGemmAtrributeMfmaIterateK
Definition: warp_gemm_attribute_mfma.hpp:107
ck_tile::WarpGemmAtrributeMfmaTransposedCDistribution_SwizzleB
Definition: warp_gemm_attribute_mfma.hpp:454
ck_tile::WarpGemmAtrributeMfmaTransposedCDistribution
Definition: warp_gemm_attribute_mfma.hpp:371
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