WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ > Struct Template Reference

WarpGemmAttributeMfmaIterateKAndTransposedCDistribution&lt; WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ &gt; Struct Template Reference#

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

#include <warp_gemm_attribute_mfma.hpp>

Public Types

using Impl = remove_cvref_t<WarpGemmAttributeMfmaImpl_>
using ADataType = typename Impl::BDataType
using BDataType = typename Impl::ADataType
using CDataType = typename Impl::CDataType
using AVecType
using BVecType
using CVecType = typename Impl::CVecType
using AWarpDstrEncoding = decltype(get_awarp_dstr_encoding())
using BWarpDstrEncoding = decltype(get_bwarp_dstr_encoding())
using CWarpDstrEncoding = decltype(get_cwarp_dstr_encoding())

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 iKIter, bool post_nop_ = false>
CK_TILE_DEVICE void operator() (CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, number< iKIter >, bool_constant< post_nop_ >={}) const
CK_TILE_DEVICE CVecType operator() (const AVecType &a_vec, const BVecType &b_vec) const

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr auto get_num_of_access ()
static CK_TILE_DEVICE constexpr auto get_awarp_dstr_encoding ()
static CK_TILE_DEVICE constexpr auto get_bwarp_dstr_encoding ()
static CK_TILE_DEVICE constexpr auto get_cwarp_dstr_encoding ()

Static Public Attributes

static constexpr auto AttrNumAccess = AttrNumAccess_
static constexpr index_t kM = Impl::kN
static constexpr index_t kN = Impl::kM
static constexpr index_t kK = Impl::kK * kKIter
static constexpr index_t kKPerThread = Impl::kABKPerLane * kKIter

Member Typedef Documentation

◆ ADataType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::ADataType = typename Impl::BDataType

◆ AVecType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::AVecType
Initial value:
ext_vector_t<ADataType, vector_traits<typename Impl::AVecType>::vector_size * kKIter>
ck_tile::ext_vector_t
typename impl::ext_vector< T, N >::type ext_vector_t
Definition vector_type.hpp:84

◆ AWarpDstrEncoding

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::AWarpDstrEncoding = decltype(get_awarp_dstr_encoding())

◆ BDataType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::BDataType = typename Impl::ADataType

◆ BVecType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::BVecType
Initial value:
ext_vector_t<BDataType, vector_traits<typename Impl::BVecType>::vector_size * kKIter>

◆ BWarpDstrEncoding

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::BWarpDstrEncoding = decltype(get_bwarp_dstr_encoding())

◆ CDataType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::CDataType = typename Impl::CDataType

◆ CVecType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::CVecType = typename Impl::CVecType

◆ CWarpDstrEncoding

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::CWarpDstrEncoding = decltype(get_cwarp_dstr_encoding())

◆ Impl

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::Impl = remove_cvref_t<WarpGemmAttributeMfmaImpl_>

Member Function Documentation

◆ get_awarp_dstr_encoding()

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
CK_TILE_DEVICE constexpr auto ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::get_awarp_dstr_encoding ( )
inlinestaticconstexpr

◆ get_bwarp_dstr_encoding()

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
CK_TILE_DEVICE constexpr auto ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::get_bwarp_dstr_encoding ( )
inlinestaticconstexpr

◆ get_cwarp_dstr_encoding()

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
CK_TILE_DEVICE constexpr auto ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::get_cwarp_dstr_encoding ( )
inlinestaticconstexpr

◆ get_num_of_access()

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

◆ operator()() [1/3]

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

◆ operator()() [2/3]

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

◆ operator()() [3/3]

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

Member Data Documentation

◆ AttrNumAccess

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, WGAttrNumAccessEnum AttrNumAccess_ = WGAttrNumAccessEnum::Single>
auto ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution< WarpGemmAttributeMfmaImpl_, kKIter, AttrNumAccess_ >::AttrNumAccess = AttrNumAccess_
staticconstexpr

◆ kK

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

◆ kKPerThread

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

◆ kM

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

◆ kN

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