/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm/warp/warp_gemm.hpp File Reference#
warp_gemm.hpp File Reference
#include "ck_tile/core.hpp"#include "ck_tile/ops/gemm/warp/warp_gemm_impl.hpp"#include "ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma.hpp"#include "ck_tile/ops/gemm/warp/warp_gemm_smfmac_impl.hpp"#include "ck_tile/ops/gemm/warp/warp_gemm_attribute_smfmac.hpp"Go to the source code of this file.
Namespaces | |
| ck_tile | |
Typedefs | |
| using | ck_tile::WarpGemmMfmaF32F32F32M16N16K4 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImplF32F32F32M16N16K4< WGAttrCtlEnum::Default_ > >> |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfmaF32F32F32M16N16K16 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImplF32F32F32M16N16K4< WGAttrCtlEnum::Default_ >, 4, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfmaF32F32F32M16N16K16TransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImplF32F32F32M16N16K4< WGAttrCtlEnum::Default_ >, 4, AttrNumAccess > > |
| using | ck_tile::WarpGemmMfmaF16F16F32M32N32K8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfmaF16F16F32M16N16K16 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImplF16F16F32M16N16K16< WGAttrCtlEnum::Default_ > >> |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfmaF16F16F32M32N32K16 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ >, 2, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfmaF16F16F32M16N16K32 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImplF16F16F32M16N16K16< WGAttrCtlEnum::Default_ >, 2, AttrNumAccess > > |
| using | ck_tile::WarpGemmMfmaF16F16F32M32N32K8SwizzleA = WarpGemmImpl< WarpGemmAttributeMfmaIterateK_SwizzleA< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ >, 1 > > |
| using | ck_tile::WarpGemmMfmaF16F16F32M32N32K16SwizzleA = WarpGemmImpl< WarpGemmAttributeMfmaIterateK_SwizzleA< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ >, 2 > > |
| using | ck_tile::WarpGemmMfmaF16F16F32M32N32K8TransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfmaF16F16F32M16N16K16TransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImplF16F16F32M16N16K16< WGAttrCtlEnum::Default_ > >> |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfmaF16F16F32M32N32K16TransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ >, 2, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfmaF16F16F32M16N16K32TransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImplF16F16F32M16N16K16< WGAttrCtlEnum::Default_ >, 2, AttrNumAccess > > |
| using | ck_tile::WarpGemmMfmaF16F16F32M32N32K8SwizzleBTransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfmaF16F16F32M32N32K16SwizzleBTransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImplF16F16F32M32N32K8< WGAttrCtlEnum::Default_ >, 2 > > |
| using | ck_tile::WarpGemmMfmaF16F16F32M4N64K16 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImplF16F16F32M4N64K4< WGAttrCtlEnum::Default_ >, 4 > > |
| using | ck_tile::WarpGemmMfmaF16F16F32M64N4K16 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImplF16F16F32M64N4K4< WGAttrCtlEnum::Default_ >, 4 > > |
| using | ck_tile::WarpGemmSmfmacF16F16F32M32N32K16 = WarpGemmSmfmacImpl< WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImplF16F16F32M32N32K16< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmSmfmacF16F16F32M16N16K32 = WarpGemmSmfmacImpl< WarpGemmAttributeSmfmac< WarpGemmAttributeSmfmacImplF16F16F32M16N16K32< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M16N16K16 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16< WGAttrCtlEnum::Default_ > >> |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K16 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ >, 2, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M16N16K32 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16< WGAttrCtlEnum::Default_ >, 2, AttrNumAccess > > |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleA = WarpGemmImpl< WarpGemmAttributeMfmaIterateK_SwizzleA< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ >, 1 > > |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleA = WarpGemmImpl< WarpGemmAttributeMfmaIterateK_SwizzleA< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ >, 2 > > |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K8TransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M16N16K16TransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16< WGAttrCtlEnum::Default_ > >> |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K16TransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ >, 2, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M16N16K32TransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16< WGAttrCtlEnum::Default_ >, 2, AttrNumAccess > > |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K8SwizzleBTransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M32N32K16SwizzleBTransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< WGAttrCtlEnum::Default_ >, 2 > > |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M4N64K16 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4< WGAttrCtlEnum::Default_ >, 4 > > |
| using | ck_tile::WarpGemmMfmaBf16Bf16F32M64N4K16 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< WGAttrCtlEnum::Default_ >, 4 > > |
| using | ck_tile::WarpGemmMfma_f32_32x32x16_fp8_fp8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_32x32x16_fp8_bf8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_32x32x16_bf8_fp8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_32x32x16_bf8_bf8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_32x32x32_fp8_fp8 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8< WGAttrCtlEnum::Default_ >, 2 > > |
| using | ck_tile::WarpGemmMfma_f32_32x32x32_bf8_bf8 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8< WGAttrCtlEnum::Default_ >, 2 > > |
| using | ck_tile::WarpGemmMfma_f32_16x16x32_fp8_fp8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_16x16x32_fp8_fp8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_16x16x32_bf8_bf8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_16x16x32_bf8_bf8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_16x16x64_fp8_fp8 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8< WGAttrCtlEnum::Default_ >, 2 > > |
| using | ck_tile::WarpGemmMfma_f32_16x16x64_bf8_bf8 = WarpGemmImpl< WarpGemmAttributeMfmaIterateK< WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8< WGAttrCtlEnum::Default_ >, 2 > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_16x16x128_fp8_fp8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_16x16x128_fp8_bf8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_16x16x128_bf8_fp8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_16x16x128_bf8_bf8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_16x16x128_fp8_fp8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_16x16x128_fp8_bf8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_16x16x128_bf8_fp8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_16x16x128_bf8_bf8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_32x32x64_fp8_fp8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_fp8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_32x32x64_fp8_bf8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_bf8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_32x32x64_bf8_fp8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_fp8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| template<WGAttrNumAccessEnum AttrNumAccess = WGAttrNumAccessEnum::Single> | |
| using | ck_tile::WarpGemmMfma_f32_32x32x64_bf8_bf8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_bf8< WGAttrCtlEnum::Default_ >, AttrNumAccess > > |
| using | ck_tile::WarpGemmMfma_f32_32x32x16_fp8_fp8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_32x32x16_fp8_bf8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_32x32x16_bf8_fp8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_f32_32x32x16_bf8_bf8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8< WGAttrCtlEnum::Default_ > >> |
| template<index_t swizzle_factor = 2> | |
| using | ck_tile::WarpGemmMfmaFp8Fp8F32M32N32K16SwizzleBTransposedCDistribution = WarpGemmImpl< WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base< fp8_t, fp8_t, WGAttrCtlEnum::Default_ >, 2, swizzle_factor > > |
| using | ck_tile::WarpGemmMfma_i32_32x32x16_i8_i8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_i32_32x32x16_i8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_i32_32x32x16_i8_i8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_i32_32x32x16_i8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_i32_16x16x32_i8_i8 = WarpGemmImpl< WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_i32_16x16x32_i8< WGAttrCtlEnum::Default_ > >> |
| using | ck_tile::WarpGemmMfma_i32_16x16x32_i8_i8_CTransposed = WarpGemmImpl< WarpGemmAttributeMfmaTransposedCDistribution< WarpGemmAttributeMfmaImpl_i32_16x16x32_i8< WGAttrCtlEnum::Default_ > >> |