WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ > Struct Template Reference

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

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

#include <warp_gemm_attribute_mfma_impl.hpp>

Public Types

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

Public Member Functions

template<index_t opselA, index_t opselB, bool post_nop_ = false>
CK_TILE_DEVICE void operator() (CVecType &c_vec, const AVecType &a_vec, const int32_t &a_scale, const BVecType &b_vec, const int32_t &b_scale, bool_constant< post_nop_ >={}) const
template<index_t opselA, index_t opselB>
CK_TILE_DEVICE CVecType operator() (const AVecType &a_vec, const int32_t &a_scale, const BVecType &b_vec, const int32_t &b_scale) 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 = 128
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 = 32
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<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::ADataType = pk_fp4_t

◆ AVecType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::AVecType = ext_vector_t<ADataType, 16>

◆ BDataType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::BDataType = pk_fp4_t

◆ BVecType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::BVecType = ext_vector_t<BDataType, 16>

◆ CDataType

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

◆ CVecType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::CVecType = ext_vector_t<CDataType, 4>

Member Function Documentation

◆ operator()() [1/2]

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
template<index_t opselA, index_t opselB>
CK_TILE_DEVICE CVecType ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::operator() ( const AVecType & a_vec,
const int32_t & a_scale,
const BVecType & b_vec,
const int32_t & b_scale ) const
inline

◆ operator()() [2/2]

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
template<index_t opselA, index_t opselB, bool post_nop_ = false>
CK_TILE_DEVICE void ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::operator() ( CVecType & c_vec,
const AVecType & a_vec,
const int32_t & a_scale,
const BVecType & b_vec,
const int32_t & b_scale,
bool_constant< post_nop_ > = {} ) const
inline

Member Data Documentation

◆ Ctrl

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

◆ kABKLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::kABKLane = 4
staticconstexpr

◆ kABKPerLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::kABKPerLane = 32
staticconstexpr

◆ kAMBlock

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::kAMBlock = 1
staticconstexpr

◆ kAMLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::kAMLane = 16
staticconstexpr

◆ kBNBlock

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

◆ kBNLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::kBNLane = 16
staticconstexpr

◆ kCM0PerLane

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

◆ kCM1PerLane

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

◆ kCMLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::kCMLane = 4
staticconstexpr

◆ kCNLane

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::kCNLane = 16
staticconstexpr

◆ kK

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::kK = 128
staticconstexpr

◆ kM

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
index_t ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4< Ctrl_ >::kM = 16
staticconstexpr

◆ kN

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