include/ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp Source File

include/ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp Source File#

Composable Kernel: include/ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp Source File
warp_gemm_dispatcher.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.hpp"
8 
9 namespace ck_tile {
10 
11 namespace impl {
12 template <typename AType,
13  typename BType,
14  typename AccType,
15  index_t MPerWave,
16  index_t NPerWave,
17  index_t KPerWave,
18  bool TransposeC,
19  bool SwizzleA = false,
20  bool UseStructuredSparsity = false,
21  WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
22 struct WarpGemmMfmaDispatcher;
23 
24 // clang-format off
25 // fp16
26 // ADataType, BDataType, AccDataType, MPerWave, NPerWave, KPerWave, TransposeC, SwizzleA, UseStructuredSparsity
27 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 32, 32, 8, false> { using Type = WarpGemmMfmaF16F16F32M32N32K8; };
28 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 32, 32, 8, true> { using Type = WarpGemmMfmaF16F16F32M32N32K8TransposedCDistribution; };
29 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 32, 32, 16, false> { using Type = WarpGemmMfmaF16F16F32M32N32K16<>; };
30 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 32, 32, 16, true> { using Type = WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution<>; };
31 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 32, 32, 16, false, false, false, WGAttrNumAccessEnum::Double> {
32  using Type = WarpGemmMfmaF16F16F32M32N32K16<WGAttrNumAccessEnum::Double>; };
33 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 32, 32, 16, true, false, false, WGAttrNumAccessEnum::Double> {
34  using Type = WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution<WGAttrNumAccessEnum::Double>; };
35 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 16, 16, 16, false> { using Type = WarpGemmMfmaF16F16F32M16N16K16; };
36 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 16, 16, 16, true> { using Type = WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution; };
37 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 16, 16, 32, false> { using Type = WarpGemmMfmaF16F16F32M16N16K32<>; };
38 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 16, 16, 32, true> { using Type = WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution<>; };
39 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 16, 16, 32, false, false, false, WGAttrNumAccessEnum::Double> {
40  using Type = WarpGemmMfmaF16F16F32M16N16K32<WGAttrNumAccessEnum::Double>; };
41 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 16, 16, 32, true, false, false, WGAttrNumAccessEnum::Double> {
42  using Type = WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution<WGAttrNumAccessEnum::Double>; };
43 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 4, 64, 16, false> { using Type = WarpGemmMfmaF16F16F32M4N64K16; };
44 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 64, 4, 16, false> { using Type = WarpGemmMfmaF16F16F32M64N4K16; };
45 
46 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 32, 32, 8, false, true> { using Type = WarpGemmMfmaF16F16F32M32N32K8SwizzleA; };
47 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 32, 32, 16, false, true> { using Type = WarpGemmMfmaF16F16F32M32N32K16SwizzleA; };
48 
49 // fp16 2:4 structural sparsity
50 // ADataType, BDataType, AccDataType, MPerWave, NPerWave, KPerWave, TransposeC, SwizzleA, UseStructuredSparsity
51 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 32, 32, 16, false, false, true> { using Type = WarpGemmSmfmacF16F16F32M32N32K16; };
52 template<> struct WarpGemmMfmaDispatcher<ck_tile::half_t, ck_tile::half_t, float, 16, 16, 32, false, false, true> { using Type = WarpGemmSmfmacF16F16F32M16N16K32; };
53 
54 // bf16
55 // ADataType, BDataType, AccDataType, MPerWave, NPerWave, KPerWave, TransposeC, SwizzleA, UseStructuredSparsity
56 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 32, 32, 8, false> { using Type = WarpGemmMfmaBf16Bf16F32M32N32K8; };
57 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 32, 32, 8, true> { using Type = WarpGemmMfmaBf16Bf16F32M32N32K8TransposedCDistribution; };
58 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 32, 32, 16, false> { using Type = WarpGemmMfmaBf16Bf16F32M32N32K16<>; };
59 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 32, 32, 16, true> { using Type = WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution<>; };
60 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 32, 32, 16, false, false, false, WGAttrNumAccessEnum::Double> {
61  using Type = WarpGemmMfmaBf16Bf16F32M32N32K16<WGAttrNumAccessEnum::Double>; };
62 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 32, 32, 16, true, false, false, WGAttrNumAccessEnum::Double> {
63  using Type = WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution<WGAttrNumAccessEnum::Double>; };
64 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 16, 16, 16, false> { using Type = WarpGemmMfmaBf16Bf16F32M16N16K16; };
65 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 16, 16, 16, true> { using Type = WarpGemmMfmaBf16Bf16F32M16N16K16TransposedCDistribution; };
66 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 16, 16, 32, false> { using Type = WarpGemmMfmaBf16Bf16F32M16N16K32<>; };
67 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 16, 16, 32, true> { using Type = WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution<>; };
68 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 16, 16, 32, false, false, false, WGAttrNumAccessEnum::Double> {
69  using Type = WarpGemmMfmaBf16Bf16F32M16N16K32<WGAttrNumAccessEnum::Double>; };
70 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 16, 16, 32, true, false, false, WGAttrNumAccessEnum::Double> {
71  using Type = WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution<WGAttrNumAccessEnum::Double>; };
72 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 4, 64, 16, false> { using Type = WarpGemmMfmaBf16Bf16F32M4N64K16; };
73 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 64, 4, 16, false> { using Type = WarpGemmMfmaBf16Bf16F32M64N4K16; };
74 
75 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 32, 32, 8, false, true> { using Type = WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleA; };
76 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf16_t, ck_tile::bf16_t, float, 32, 32, 16, false, true> { using Type = WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleA; };
77 
78 // fp8
79 // ADataType, BDataType, AccDataType, MPerWave, NPerWave, KPerWave, TransposeC, SwizzleA, UseStructuredSparsity
80 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::fp8_t, float, 32, 32, 16, false> { using Type = WarpGemmMfma_f32_32x32x16_fp8_fp8; };
81 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::fp8_t, float, 32, 32, 32, false> { using Type = WarpGemmMfma_f32_32x32x32_fp8_fp8; };
82 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::fp8_t, float, 16, 16, 32, false> { using Type = WarpGemmMfma_f32_16x16x32_fp8_fp8; };
83 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::fp8_t, float, 16, 16, 64, false> { using Type = WarpGemmMfma_f32_16x16x64_fp8_fp8; };
84 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::fp8_t, float, 32, 32, 16, true> { using Type = WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed; };
85 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::bf8_t, float, 32, 32, 16, false> { using Type = WarpGemmMfma_f32_32x32x16_fp8_bf8; };
86 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::bf8_t, float, 32, 32, 16, true> { using Type = WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed; };
87 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::fp8_t, float, 32, 32, 16, false> { using Type = WarpGemmMfma_f32_32x32x16_bf8_fp8; };
88 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::fp8_t, float, 32, 32, 16, true> { using Type = WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed; };
89 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::bf8_t, float, 32, 32, 16, false> { using Type = WarpGemmMfma_f32_32x32x16_bf8_bf8; };
90 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::bf8_t, float, 32, 32, 32, false> { using Type = WarpGemmMfma_f32_32x32x32_bf8_bf8; };
91 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::bf8_t, float, 16, 16, 32, false> { using Type = WarpGemmMfma_f32_16x16x32_bf8_bf8; };
92 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::bf8_t, float, 16, 16, 64, false> { using Type = WarpGemmMfma_f32_16x16x64_bf8_bf8; };
93 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::bf8_t, float, 32, 32, 16, true> { using Type = WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed; };
94 
95 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::fp8_t, float, 16, 16, 128, false> { using Type = WarpGemmMfma_f32_16x16x128_fp8_fp8<>; };
96 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::bf8_t, float, 16, 16, 128, false> { using Type = WarpGemmMfma_f32_16x16x128_fp8_bf8<>; };
97 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::fp8_t, float, 16, 16, 128, false> { using Type = WarpGemmMfma_f32_16x16x128_bf8_fp8<>; };
98 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::bf8_t, float, 16, 16, 128, false> { using Type = WarpGemmMfma_f32_16x16x128_bf8_bf8<>; };
99 
100 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::fp8_t, float, 32, 32, 64, false> { using Type = WarpGemmMfma_f32_32x32x64_fp8_fp8<>; };
101 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::bf8_t, float, 32, 32, 64, false> { using Type = WarpGemmMfma_f32_32x32x64_fp8_bf8<>; };
102 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::fp8_t, float, 32, 32, 64, false> { using Type = WarpGemmMfma_f32_32x32x64_bf8_fp8<>; };
103 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::bf8_t, float, 32, 32, 64, false> { using Type = WarpGemmMfma_f32_32x32x64_bf8_bf8<>; };
104 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::fp8_t, float, 32, 32, 64, false, false, false, WGAttrNumAccessEnum::Quad> {
105  using Type = WarpGemmMfma_f32_32x32x64_fp8_fp8<WGAttrNumAccessEnum::Quad>; };
106 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::bf8_t, float, 32, 32, 64, false, false, false, WGAttrNumAccessEnum::Quad> {
107  using Type = WarpGemmMfma_f32_32x32x64_fp8_bf8<WGAttrNumAccessEnum::Quad>; };
108 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::fp8_t, float, 32, 32, 64, false, false, false, WGAttrNumAccessEnum::Quad> {
109  using Type = WarpGemmMfma_f32_32x32x64_bf8_fp8<WGAttrNumAccessEnum::Quad>; };
110 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::bf8_t, float, 32, 32, 64, false, false, false, WGAttrNumAccessEnum::Quad> {
111  using Type = WarpGemmMfma_f32_32x32x64_bf8_bf8<WGAttrNumAccessEnum::Quad>; };
112 
113 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::fp8_t, float, 16, 16, 128, false, false, false, WGAttrNumAccessEnum::Quad> {
114  using Type = WarpGemmMfma_f32_16x16x128_fp8_fp8<WGAttrNumAccessEnum::Quad>; };
115 template<> struct WarpGemmMfmaDispatcher<ck_tile::fp8_t, ck_tile::bf8_t, float, 16, 16, 128, false, false, false, WGAttrNumAccessEnum::Quad> {
116  using Type = WarpGemmMfma_f32_16x16x128_fp8_bf8<WGAttrNumAccessEnum::Quad>; };
117 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::fp8_t, float, 16, 16, 128, false, false, false, WGAttrNumAccessEnum::Quad> {
118  using Type = WarpGemmMfma_f32_16x16x128_bf8_fp8<WGAttrNumAccessEnum::Quad>; };
119 template<> struct WarpGemmMfmaDispatcher<ck_tile::bf8_t, ck_tile::bf8_t, float, 16, 16, 128, false, false, false, WGAttrNumAccessEnum::Quad> {
120  using Type = WarpGemmMfma_f32_16x16x128_bf8_bf8<WGAttrNumAccessEnum::Quad>; };
121 // int8
122 // ADataType, BDataType, AccDataType, MPerWave, NPerWave, KPerWave, TransposeC, SwizzleA, UseStructuredSparsity
123 template<> struct WarpGemmMfmaDispatcher<ck_tile::int8_t, ck_tile::int8_t, ck_tile::int32_t, 32, 32, 16, false> { using Type = WarpGemmMfma_i32_32x32x16_i8_i8; };
124 template<> struct WarpGemmMfmaDispatcher<ck_tile::int8_t, ck_tile::int8_t, ck_tile::int32_t, 32, 32, 16, true> { using Type = WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed; };
125 template<> struct WarpGemmMfmaDispatcher<ck_tile::int8_t, ck_tile::int8_t, ck_tile::int32_t, 16, 16, 32, false> { using Type = WarpGemmMfma_i32_16x16x32_i8_i8; };
126 template<> struct WarpGemmMfmaDispatcher<ck_tile::int8_t, ck_tile::int8_t, ck_tile::int32_t, 16, 16, 32, true> { using Type = WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed; };
127 
128 // clang-format on
129 } // namespace impl
130 
131 template <typename AType,
132  typename BType,
133  typename AccType,
134  index_t MPerWave,
135  index_t NPerWave,
136  index_t KPerWave,
137  bool TransposeC,
138  bool SwizzleA = false,
139  bool UseStructuredSparsity = false,
140  WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single>
141 using WarpGemmMfmaDispatcher = typename impl::WarpGemmMfmaDispatcher<AType,
142  BType,
143  AccType,
144  MPerWave,
145  NPerWave,
146  KPerWave,
147  TransposeC,
148  SwizzleA,
149  UseStructuredSparsity,
150  AttrNumAccess>::Type;
151 
152 } // namespace ck_tile
ck_tile
Definition: cluster_descriptor.hpp:13
ck_tile::WarpGemmMfma_f32_32x32x16_bf8_fp8
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_f32_32x32x16_bf8_fp8
Definition: warp_gemm.hpp:241
ck_tile::WGAttrNumAccessEnum
WGAttrNumAccessEnum
Definition: warp_gemm_attribute_mfma.hpp:13
ck_tile::WGAttrNumAccessEnum::Single
@ Single
ck_tile::WGAttrNumAccessEnum::Double
@ Double
ck_tile::WGAttrNumAccessEnum::Quad
@ Quad
ck_tile::WarpGemmMfmaF16F16F32M4N64K16
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK< WarpGemmAttributeMfmaImplF16F16F32M4N64K4< WGAttrCtlEnum::Default_ >, 4 > > WarpGemmMfmaF16F16F32M4N64K16
Definition: warp_gemm.hpp:120
ck_tile::WarpGemmMfmaF16F16F32M32N32K8
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ > >> WarpGemmMfmaF16F16F32M32N32K8
Definition: warp_gemm.hpp:18
ck_tile::WarpGemmMfmaDispatcher
typename impl::WarpGemmMfmaDispatcher< AType, BType, AccType, MPerWave, NPerWave, KPerWave, TransposeC, SwizzleA, UseStructuredSparsity, AttrNumAccess >::Type WarpGemmMfmaDispatcher
Definition: warp_gemm_dispatcher.hpp:150
ck_tile::WarpGemmMfmaF16F16F32M64N4K16
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK< WarpGemmAttributeMfmaImplF16F16F32M64N4K4< WGAttrCtlEnum::Default_ >, 4 > > WarpGemmMfmaF16F16F32M64N4K16
Definition: warp_gemm.hpp:124
ck_tile::fp8_t
_BitInt(8) fp8_t
Definition: float8.hpp:204
ck_tile::WarpGemmMfmaF16F16F32M32N32K16SwizzleA
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK_SwizzleA< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ >, 2 > > WarpGemmMfmaF16F16F32M32N32K16SwizzleA
Definition: warp_gemm.hpp:55
ck_tile::WarpGemmMfmaBf16Bf16F32M4N64K16
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK< WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4< WGAttrCtlEnum::Default_ >, 4 > > WarpGemmMfmaBf16Bf16F32M4N64K16
Definition: warp_gemm.hpp:226
ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K8
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ > >> WarpGemmMfmaBf16Bf16F32M32N32K8
Definition: warp_gemm.hpp:135
ck_tile::WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution
WarpGemmImpl< WarpGemmAtrributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImplF16F16F32M16N16K16< WGAttrCtlEnum::Default_ > >> WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution
Definition: warp_gemm.hpp:63
ck_tile::WarpGemmMfma_f32_32x32x32_fp8_fp8
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK< WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8< WGAttrCtlEnum::Default_ >, 2 > > WarpGemmMfma_f32_32x32x32_fp8_fp8
Definition: warp_gemm.hpp:248
ck_tile::WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed
WarpGemmImpl< WarpGemmAtrributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed
Definition: warp_gemm.hpp:310
ck_tile::WarpGemmMfma_i32_16x16x32_i8_i8
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImpl_i32_16x16x32_i8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_i32_16x16x32_i8_i8
Definition: warp_gemm.hpp:340
ck_tile::int8_t
int8_t int8_t
Definition: int8.hpp:20
ck_tile::WarpGemmMfma_f32_16x16x32_bf8_bf8
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_f32_16x16x32_bf8_bf8
Definition: warp_gemm.hpp:258
ck_tile::WarpGemmMfma_f32_16x16x64_bf8_bf8
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK< WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8< WGAttrCtlEnum::Default_ >, 2 > > WarpGemmMfma_f32_16x16x64_bf8_bf8
Definition: warp_gemm.hpp:266
ck_tile::bf16_t
bfloat16_t bf16_t
Definition: bfloat16.hpp:106
ck_tile::WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed
WarpGemmImpl< WarpGemmAtrributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_i32_32x32x16_i8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed
Definition: warp_gemm.hpp:337
ck_tile::WarpGemmMfmaF16F16F32M32N32K8SwizzleA
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK_SwizzleA< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ >, 1 > > WarpGemmMfmaF16F16F32M32N32K8SwizzleA
Definition: warp_gemm.hpp:51
ck_tile::index_t
int32_t index_t
Definition: integer.hpp:9
ck_tile::WarpGemmSmfmacF16F16F32M32N32K16
WarpGemmSmfmacImpl< WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImplF16F16F32M32N32K16< WGAttrCtlEnum::Default_ > >> WarpGemmSmfmacF16F16F32M32N32K16
Definition: warp_gemm.hpp:128
ck_tile::WarpGemmMfma_f32_32x32x16_fp8_bf8
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_f32_32x32x16_fp8_bf8
Definition: warp_gemm.hpp:238
ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K8TransposedCDistribution
WarpGemmImpl< WarpGemmAtrributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ > >> WarpGemmMfmaBf16Bf16F32M32N32K8TransposedCDistribution
Definition: warp_gemm.hpp:177
ck_tile::WarpGemmMfma_i32_32x32x16_i8_i8
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImpl_i32_32x32x16_i8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_i32_32x32x16_i8_i8
Definition: warp_gemm.hpp:333
ck_tile::WarpGemmMfma_f32_16x16x32_fp8_fp8
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_f32_16x16x32_fp8_fp8
Definition: warp_gemm.hpp:255
ck_tile::WarpGemmMfmaBf16Bf16F32M16N16K16
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16< WGAttrCtlEnum::Default_ > >> WarpGemmMfmaBf16Bf16F32M16N16K16
Definition: warp_gemm.hpp:138
ck_tile::WarpGemmMfma_f32_16x16x64_fp8_fp8
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK< WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8< WGAttrCtlEnum::Default_ >, 2 > > WarpGemmMfma_f32_16x16x64_fp8_fp8
Definition: warp_gemm.hpp:262
ck_tile::WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed
WarpGemmImpl< WarpGemmAtrributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed
Definition: warp_gemm.hpp:314
ck_tile::int32_t
int32_t int32_t
Definition: integer.hpp:10
ck_tile::WarpGemmMfmaBf16Bf16F32M64N4K16
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK< WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< WGAttrCtlEnum::Default_ >, 4 > > WarpGemmMfmaBf16Bf16F32M64N4K16
Definition: warp_gemm.hpp:230
ck_tile::WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed
WarpGemmImpl< WarpGemmAtrributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed
Definition: warp_gemm.hpp:318
ck_tile::WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed
WarpGemmImpl< WarpGemmAtrributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_i32_16x16x32_i8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed
Definition: warp_gemm.hpp:344
ck_tile::bf8_t
unsigned _BitInt(8) bf8_t
Definition: float8.hpp:206
ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleA
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK_SwizzleA< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ >, 2 > > WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleA
Definition: warp_gemm.hpp:173
ck_tile::WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed
WarpGemmImpl< WarpGemmAtrributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed
Definition: warp_gemm.hpp:322
ck_tile::WarpGemmMfma_f32_32x32x32_bf8_bf8
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK< WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8< WGAttrCtlEnum::Default_ >, 2 > > WarpGemmMfma_f32_32x32x32_bf8_bf8
Definition: warp_gemm.hpp:252
ck_tile::WarpGemmSmfmacF16F16F32M16N16K32
WarpGemmSmfmacImpl< WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImplF16F16F32M16N16K32< WGAttrCtlEnum::Default_ > >> WarpGemmSmfmacF16F16F32M16N16K32
Definition: warp_gemm.hpp:131
ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleA
WarpGemmImpl< WarpGemmAtrributeMfmaIterateK_SwizzleA< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ >, 1 > > WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleA
Definition: warp_gemm.hpp:168
ck_tile::WarpGemmMfma_f32_32x32x16_fp8_fp8
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_f32_32x32x16_fp8_fp8
Definition: warp_gemm.hpp:235
ck_tile::half_t
_Float16 half_t
Definition: half.hpp:111
ck_tile::WarpGemmMfmaBf16Bf16F32M16N16K16TransposedCDistribution
WarpGemmImpl< WarpGemmAtrributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16< WGAttrCtlEnum::Default_ > >> WarpGemmMfmaBf16Bf16F32M16N16K16TransposedCDistribution
Definition: warp_gemm.hpp:181
ck_tile::WarpGemmMfmaF16F16F32M32N32K8TransposedCDistribution
WarpGemmImpl< WarpGemmAtrributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ > >> WarpGemmMfmaF16F16F32M32N32K8TransposedCDistribution
Definition: warp_gemm.hpp:59
ck_tile::WarpGemmMfmaF16F16F32M16N16K16
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImplF16F16F32M16N16K16< WGAttrCtlEnum::Default_ > >> WarpGemmMfmaF16F16F32M16N16K16
Definition: warp_gemm.hpp:21
ck_tile::WarpGemmMfma_f32_32x32x16_bf8_bf8
WarpGemmImpl< WarpGemmAtrributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8< WGAttrCtlEnum::Default_ > >> WarpGemmMfma_f32_32x32x16_bf8_bf8
Definition: warp_gemm.hpp:244
ck_tile::WarpGemmImpl
Definition: warp_gemm_impl.hpp:11
ck_tile::WarpGemmSmfmacImpl
Definition: warp_gemm_smfmac_impl.hpp:11
ck_tile::impl::WarpGemmMfmaDispatcher
Definition: warp_gemm_dispatcher.hpp:22
warp_gemm.hpp