/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp File Reference

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp File Reference#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp File Reference
warp_gemm_attribute_mfma_impl.hpp File Reference
#include "ck_tile/core.hpp"

Go to the source code of this file.

Classes

struct  ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base< AType_, BType_, Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base< AType_, BType_, Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base< AType_, BType_, Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8< Ctrl_ >
 
struct  ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8< Ctrl_ >
 

Namespaces

 ck_tile
 

Macros

#define DISPATCH_MFMA_(mfma_, dmod_, amod_, bmod_, cmod_)
 
#define DISPATCH_MFMA_CTRL_(mfma_, ctrl_)
 

Typedefs

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8 = WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base< fp8_t, fp8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8 = WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< fp8_t, fp8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8 = WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base< fp8_t, bf8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8 = WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< bf8_t, bf8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8 = WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base< bf8_t, fp8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8 = WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base< bf8_t, bf8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8 = WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base< fp8_t, fp8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8 = WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base< fp8_t, bf8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8 = WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base< bf8_t, fp8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8 = WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base< bf8_t, bf8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_fp8 = WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base< fp8_t, fp8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_bf8 = WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base< fp8_t, bf8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_fp8 = WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base< bf8_t, fp8_t, Ctrl_ >
 
template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_bf8 = WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base< bf8_t, bf8_t, Ctrl_ >
 

Enumerations

enum class  ck_tile::WGAttrCtlEnum {
  ck_tile::Default_ = 0 ,
  ck_tile::Raw_vvv = 1 ,
  ck_tile::Raw_vaa = 2 ,
  ck_tile::Raw_vav = 3 ,
  ck_tile::Raw_vva = 4 ,
  ck_tile::Raw_avv = 5
}
 

Macro Definition Documentation

◆ DISPATCH_MFMA_

#define DISPATCH_MFMA_ (   mfma_,
  dmod_,
  amod_,
  bmod_,
  cmod_ 
)
Value:
if constexpr(post_nop_) \
{ \
asm volatile(mfma_ " %0, %1, %2, %3 ; yyy\n" \
"s_nop 3" \
: dmod_(c_vec) \
: amod_(a_vec), bmod_(b_vec), cmod_(c_vec) \
:); \
} \
else \
{ \
asm volatile(mfma_ " %0, %1, %2, %3\n" \
: dmod_(c_vec) \
: amod_(a_vec), bmod_(b_vec), cmod_(c_vec) \
:); \
}

◆ DISPATCH_MFMA_CTRL_

#define DISPATCH_MFMA_CTRL_ (   mfma_,
  ctrl_ 
)
Value:
if constexpr(ctrl_ == WGAttrCtlEnum::Raw_vvv) \
{ \
DISPATCH_MFMA_(mfma_, "+v", "v", "v", "v") \
} \
else if constexpr(ctrl_ == WGAttrCtlEnum::Raw_vaa) \
{ \
DISPATCH_MFMA_(mfma_, "+v", "a", "a", "v") \
} \
else if constexpr(ctrl_ == WGAttrCtlEnum::Raw_vav) \
{ \
DISPATCH_MFMA_(mfma_, "+v", "a", "v", "v") \
} \
else if constexpr(ctrl_ == WGAttrCtlEnum::Raw_vva) \
{ \
DISPATCH_MFMA_(mfma_, "+v", "v", "a", "v") \
} \
else if constexpr(ctrl_ == WGAttrCtlEnum::Raw_avv) \
{ \
DISPATCH_MFMA_(mfma_, "+a", "v", "v", "a") \
}