/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

• develop
• include
• ck_tile
• ops
• gemm
• warp

Classes | Namespaces | Macros | Typedefs | Enumerations

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::WarpGemmAttributeMfmaImplF32F32F32M16N16K4< Ctrl_ >
struct	ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2< Ctrl_ >
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::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< 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_fp8_bf8 = WarpGemmAttributeMfmaImpl_f32_16x16x32_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") \

}