/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.
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") \
}