WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ > Struct Template Reference

WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4&lt; Ctrl_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Attributes | List of all members
ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ > Struct Template Reference

#include <warp_gemm_attribute_mfma_impl.hpp>

Public Types

using ADataType = bf16_t
using BDataType = bf16_t
using CDataType = float
using AVecType = ext_vector_t<bf16_t, 4>
using BVecType = ext_vector_t<bf16_t, 4>
using CVecType = ext_vector_t<float, 4>

Public Member Functions

template<bool post_nop_ = false>
CK_TILE_DEVICE void operator() (CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const
CK_TILE_DEVICE CVecType operator() (const AVecType &a_vec, const BVecType &b_vec) const

Static Public Attributes

static constexpr WGAttrCtlEnum Ctrl = Ctrl_
static constexpr index_t kM = 64
static constexpr index_t kN = 4
static constexpr index_t kK = 4
static constexpr index_t kAMBlock = 16
static constexpr index_t kBNBlock = 1
static constexpr index_t kAMLane = 4
static constexpr index_t kBNLane = 4
static constexpr index_t kABKLane = 1
static constexpr index_t kABKPerLane = 4
static constexpr index_t kCMLane = 1
static constexpr index_t kCNLane = 4
static constexpr index_t kCM0PerLane = 1
static constexpr index_t kCM1PerLane = 4

Member Typedef Documentation

◆ ADataType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::ADataType = bf16_t

◆ AVecType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::AVecType = ext_vector_t<bf16_t, 4>

◆ BDataType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::BDataType = bf16_t

◆ BVecType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::BVecType = ext_vector_t<bf16_t, 4>

◆ CDataType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::CDataType = float

◆ CVecType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::CVecType = ext_vector_t<float, 4>

Member Function Documentation

◆ operator()() [1/2]

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
CK_TILE_DEVICE CVecType ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::operator() ( const AVecType & a_vec,
const BVecType & b_vec ) const
inline

◆ operator()() [2/2]

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
template<bool post_nop_ = false>
CK_TILE_DEVICE void ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::operator() ( CVecType & c_vec,
const AVecType & a_vec,
const BVecType & b_vec,
bool_constant< post_nop_ > = {} ) const
inline

Member Data Documentation

◆ Ctrl

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
WGAttrCtlEnum ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::Ctrl = Ctrl_
staticconstexpr

◆ kABKLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kABKLane = 1
staticconstexpr

◆ kABKPerLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kABKPerLane = 4
staticconstexpr

◆ kAMBlock

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kAMBlock = 16
staticconstexpr

◆ kAMLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kAMLane = 4
staticconstexpr

◆ kBNBlock

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kBNBlock = 1
staticconstexpr

◆ kBNLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kBNLane = 4
staticconstexpr

◆ kCM0PerLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kCM0PerLane = 1
staticconstexpr

◆ kCM1PerLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kCM1PerLane = 4
staticconstexpr

◆ kCMLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kCMLane = 1
staticconstexpr

◆ kCNLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kCNLane = 4
staticconstexpr

◆ kK

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kK = 4
staticconstexpr

◆ kM

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kM = 64
staticconstexpr

◆ kN

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4< Ctrl_ >::kN = 4
staticconstexpr
The documentation for this struct was generated from the following file: