WarpGemmAttributeMfmaImplF16F16F32M64N4K4< Ctrl_ > Struct Template Reference

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

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

#include <warp_gemm_attribute_mfma_impl.hpp>

Public Types

using ADataType = fp16_t
 
using BDataType = fp16_t
 
using CDataType = float
 
using AVecType = ext_vector_t< fp16_t, 4 >
 
using BVecType = ext_vector_t< fp16_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::WarpGemmAttributeMfmaImplF16F16F32M64N4K4< Ctrl_ >::ADataType = fp16_t

◆ AVecType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4< Ctrl_ >::AVecType = ext_vector_t<fp16_t, 4>

◆ BDataType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4< Ctrl_ >::BDataType = fp16_t

◆ BVecType

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
using ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4< Ctrl_ >::BVecType = ext_vector_t<fp16_t, 4>

◆ CDataType

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

◆ CVecType

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

Member Function Documentation

◆ operator()() [1/2]

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
CK_TILE_DEVICE CVecType ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4< 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::WarpGemmAttributeMfmaImplF16F16F32M64N4K4< 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_>
constexpr WGAttrCtlEnum ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4< Ctrl_ >::Ctrl = Ctrl_
staticconstexpr

◆ kABKLane

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

◆ kABKPerLane

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

◆ kAMBlock

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

◆ kAMLane

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

◆ kBNBlock

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

◆ kBNLane

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

◆ kCM0PerLane

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

◆ kCM1PerLane

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

◆ kCMLane

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

◆ kCNLane

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

◆ kK

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

◆ kM

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

◆ kN

template<WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>
constexpr index_t ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4< Ctrl_ >::kN = 4
staticconstexpr
The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-7.1.1/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp