WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ > Struct Template Reference

WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB&lt; WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ > 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
using BWarpDstrEncoding
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 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 Public Attributes

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
static constexpr index_t SFactor = SFactor_

Member Typedef Documentation

◆ ADataType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::ADataType = typename Impl::BDataType

◆ AVecType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::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, index_t SFactor_ = 2>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::AWarpDstrEncoding
Initial value:
tile_distribution_encoding<
sequence<>,
tuple<sequence<Impl::kBNLane>, sequence<Impl::kABKLane, Impl::kABKPerLane * kKIter>>,
tuple<sequence<2, 1>>,
tuple<sequence<0, 0>>,
sequence<2>,
sequence<1>>
ck_tile::sequence<>
ck_tile::tile_distribution_encoding
Definition tile_distribution_encoding.hpp:26
ck_tile::tuple
Definition tile/core/container/tuple.hpp:192

◆ BDataType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::BDataType = typename Impl::ADataType

◆ BVecType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::BVecType
Initial value:
ext_vector_t<BDataType, vector_traits<typename Impl::BVecType>::vector_size * kKIter>

◆ BWarpDstrEncoding

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::BWarpDstrEncoding
Initial value:
tile_distribution_encoding<
sequence<>,
tuple<sequence<Impl::kAMLane / (Impl::kCMLane * SFactor * Impl::kCM1PerLane),
Impl::kCMLane,
SFactor,
Impl::kCM1PerLane>,
sequence<Impl::kABKLane, Impl::kABKPerLane * kKIter>>,
tuple<sequence<2, 1, 1, 1, 1>>,
tuple<sequence<0, 0, 2, 1, 3>>,
sequence<2>,
sequence<1>>
ck_tile::WarpGemmAttributeMfmaTransposedCDistribution_SwizzleB::SFactor
static constexpr index_t SFactor
Definition warp_gemm_attribute_mfma.hpp:494

◆ CDataType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::CDataType = typename Impl::CDataType

◆ CVecType

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::CVecType = typename Impl::CVecType

◆ CWarpDstrEncoding

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::CWarpDstrEncoding
Initial value:
tile_distribution_encoding<
sequence<>,
tuple<sequence<Impl::kCNLane>,
sequence<Impl::kCM0PerLane / SFactor, Impl::kCMLane, Impl::kCM1PerLane * SFactor>>,
tuple<sequence<2, 1>>,
tuple<sequence<1, 0>>,
sequence<2, 2>,
sequence<0, 2>>

◆ Impl

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
using ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::Impl = remove_cvref_t<WarpGemmAttributeMfmaImpl_>

Member Function Documentation

◆ get_num_of_access()

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::get_num_of_access ( )
inlinestaticconstexpr

◆ operator()() [1/3]

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
CK_TILE_DEVICE CVecType ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::operator() ( const AVecType & a_vec,
const BVecType & b_vec ) const
inline

◆ operator()() [2/3]

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
template<bool post_nop_ = false>
CK_TILE_DEVICE void ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::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, index_t SFactor_ = 2>
template<index_t iKIter, bool post_nop_ = false>
CK_TILE_DEVICE void ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::operator() ( CVecType & c_vec,
const AVecType & a_vec,
const BVecType & b_vec,
number< iKIter > ,
bool_constant< post_nop_ > = {} ) const
inline

Member Data Documentation

◆ kK

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
index_t ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::kK = Impl::kK * kKIter
staticconstexpr

◆ kKPerThread

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
index_t ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::kKPerThread = Impl::kABKPerLane * kKIter
staticconstexpr

◆ kM

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
index_t ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::kM = Impl::kN
staticconstexpr

◆ kN

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
index_t ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::kN = Impl::kM
staticconstexpr

◆ SFactor

template<typename WarpGemmAttributeMfmaImpl_, index_t kKIter, index_t SFactor_ = 2>
index_t ck_tile::WarpGemmAttributeMfmaIterateKAndTransposedCDistribution_SwizzleB< WarpGemmAttributeMfmaImpl_, kKIter, SFactor_ >::SFactor = SFactor_
staticconstexpr
The documentation for this struct was generated from the following file: