WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ > Struct Template Reference

WarpGemmAttributeMfma&lt; WarpGemmAttributeMfmaImpl_, AttrNumAccess_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ > Struct Template Reference

#include <warp_gemm_attribute_mfma.hpp>

Public Types

using Impl = remove_cvref_t<WarpGemmAttributeMfmaImpl_>
using ADataType = typename Impl::ADataType
using BDataType = typename Impl::BDataType
using CDataType = typename Impl::CDataType
using AVecType = typename Impl::AVecType
using BVecType = typename Impl::BVecType
using CVecType = typename Impl::CVecType
using AWarpDstrEncoding = decltype(get_warp_dstr_encoding<Impl::kAMLane>())
using BWarpDstrEncoding = decltype(get_warp_dstr_encoding<Impl::kBNLane>())
using CWarpDstrEncoding

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
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
CK_TILE_DEVICE CVecType operator() (const AVecType &a_vec, const BVecType &b_vec) 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 Member Functions

static CK_TILE_HOST_DEVICE constexpr auto get_num_of_access ()
template<index_t kMNLane>
static constexpr auto get_warp_dstr_encoding ()

Static Public Attributes

static constexpr auto AttrNumAccess = AttrNumAccess_
static constexpr auto AttrNumAccessV = static_cast<index_t>(AttrNumAccess)
static constexpr index_t kM = Impl::kM
static constexpr index_t kN = Impl::kN
static constexpr index_t kK = Impl::kK
static constexpr index_t kKPerThread = Impl::kABKPerLane
static constexpr index_t kCMLane = Impl::kCMLane

Member Typedef Documentation

◆ ADataType

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::ADataType = typename Impl::ADataType

◆ AVecType

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::AVecType = typename Impl::AVecType

◆ AWarpDstrEncoding

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::AWarpDstrEncoding = decltype(get_warp_dstr_encoding<Impl::kAMLane>())

◆ BDataType

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::BDataType = typename Impl::BDataType

◆ BVecType

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::BVecType = typename Impl::BVecType

◆ BWarpDstrEncoding

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::BWarpDstrEncoding = decltype(get_warp_dstr_encoding<Impl::kBNLane>())

◆ CDataType

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::CDataType = typename Impl::CDataType

◆ CVecType

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::CVecType = typename Impl::CVecType

◆ CWarpDstrEncoding

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::CWarpDstrEncoding
Initial value:
tile_distribution_encoding<
sequence<>,
tuple<sequence<Impl::kCM0PerLane, Impl::kCMLane, Impl::kCM1PerLane>,
sequence<Impl::kCNLane>>,
tuple<sequence<1, 2>>,
tuple<sequence<1, 0>>,
sequence<1, 1>,
sequence<0, 2>>
ck_tile::sequence<>
ck_tile::tile_distribution_encoding
Definition tile_distribution_encoding.hpp:26
ck_tile::tuple
Definition tile/core/container/tuple.hpp:192

◆ Impl

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::Impl = remove_cvref_t<WarpGemmAttributeMfmaImpl_>

Member Function Documentation

◆ get_num_of_access()

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::get_num_of_access ( )
inlinestaticconstexpr

◆ get_warp_dstr_encoding()

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
template<index_t kMNLane>
constexpr auto ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::get_warp_dstr_encoding ( )
inlinestaticconstexpr

◆ operator()() [1/4]

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
CK_TILE_DEVICE CVecType ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::operator() ( const AVecType & a_vec,
const BVecType & b_vec ) const
inline

◆ operator()() [2/4]

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
template<index_t opselA, index_t opselB>
CK_TILE_DEVICE CVecType ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::operator() ( const AVecType & a_vec,
const int32_t & a_scale,
const BVecType & b_vec,
const int32_t & b_scale ) const
inline

◆ operator()() [3/4]

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
template<bool post_nop_ = false>
CK_TILE_DEVICE void ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::operator() ( CVecType & c_vec,
const AVecType & a_vec,
const BVecType & b_vec,
bool_constant< post_nop_ > = {} ) const
inline

◆ operator()() [4/4]

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
template<index_t opselA, index_t opselB, bool post_nop_ = false>
CK_TILE_DEVICE void ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::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

◆ AttrNumAccess

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
auto ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::AttrNumAccess = AttrNumAccess_
staticconstexpr

◆ AttrNumAccessV

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
auto ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::AttrNumAccessV = static_cast<index_t>(AttrNumAccess)
staticconstexpr

◆ kCMLane

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
index_t ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::kCMLane = Impl::kCMLane
staticconstexpr

◆ kK

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
index_t ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::kK = Impl::kK
staticconstexpr

◆ kKPerThread

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
index_t ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::kKPerThread = Impl::kABKPerLane
staticconstexpr

◆ kM

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
index_t ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::kM = Impl::kM
staticconstexpr

◆ kN

template<typename WarpGemmAttributeMfmaImpl_, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
index_t ck_tile::WarpGemmAttributeMfma< WarpGemmAttributeMfmaImpl_, AttrNumAccess_ >::kN = Impl::kN
staticconstexpr
The documentation for this struct was generated from the following file: