WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ > Struct Template Reference

WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base&lt; AType_, BType_, Ctrl_ &gt; Struct Template Reference#

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

#include <warp_gemm_attribute_mfma_impl.hpp>

Public Types

using ADataType = AType_
using BDataType = BType_
using CDataType = float
using AVecType = ext_vector_t<ADataType, 8>
using BVecType = ext_vector_t<BDataType, 8>
using CVecType = ext_vector_t<CDataType, 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 = 16
static constexpr index_t kN = 16
static constexpr index_t kK = 32
static constexpr index_t kAMBlock = 1
static constexpr index_t kBNBlock = 1
static constexpr index_t kAMLane = 16
static constexpr index_t kBNLane = 16
static constexpr index_t kABKLane = 4
static constexpr index_t kABKPerLane = 8
static constexpr index_t kCMLane = 4
static constexpr index_t kCNLane = 16
static constexpr index_t kCM0PerLane = 1
static constexpr index_t kCM1PerLane = 4

Member Typedef Documentation

◆ ADataType

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::ADataType = AType_

◆ AVecType

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::AVecType = ext_vector_t<ADataType, 8>

◆ BDataType

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::BDataType = BType_

◆ BVecType

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::BVecType = ext_vector_t<BDataType, 8>

◆ CDataType

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::CDataType = float

◆ CVecType

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::CVecType = ext_vector_t<CDataType, 4>

Member Function Documentation

◆ operator()() [1/2]

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

◆ operator()() [2/2]

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

Member Data Documentation

◆ Ctrl

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
WGAttrCtlEnum ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::Ctrl = Ctrl_
staticconstexpr

◆ kABKLane

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kABKLane = 4
staticconstexpr

◆ kABKPerLane

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kABKPerLane = 8
staticconstexpr

◆ kAMBlock

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kAMBlock = 1
staticconstexpr

◆ kAMLane

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kAMLane = 16
staticconstexpr

◆ kBNBlock

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kBNBlock = 1
staticconstexpr

◆ kBNLane

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kBNLane = 16
staticconstexpr

◆ kCM0PerLane

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kCM0PerLane = 1
staticconstexpr

◆ kCM1PerLane

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kCM1PerLane = 4
staticconstexpr

◆ kCMLane

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kCMLane = 4
staticconstexpr

◆ kCNLane

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kCNLane = 16
staticconstexpr

◆ kK

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kK = 32
staticconstexpr

◆ kM

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kM = 16
staticconstexpr

◆ kN

template<typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base< AType_, BType_, Ctrl_ >::kN = 16
staticconstexpr
The documentation for this struct was generated from the following file: