WarpGemmImpl< WarpGemmAttribute_ > Struct Template Reference

WarpGemmImpl&lt; WarpGemmAttribute_ &gt; Struct Template Reference#

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

#include <warp_gemm_impl.hpp>

Public Types

using WarpGemmAttribute = remove_cvref_t<WarpGemmAttribute_>
using ADataType = typename WarpGemmAttribute::ADataType
using BDataType = typename WarpGemmAttribute::BDataType
using CDataType = typename WarpGemmAttribute::CDataType
using AWarpDstrEncoding = typename WarpGemmAttribute::AWarpDstrEncoding
using BWarpDstrEncoding = typename WarpGemmAttribute::BWarpDstrEncoding
using CWarpDstrEncoding = typename WarpGemmAttribute::CWarpDstrEncoding
using AWarpDstr = remove_cvref_t<decltype(make_static_tile_distribution(AWarpDstrEncoding{}))>
using BWarpDstr = remove_cvref_t<decltype(make_static_tile_distribution(BWarpDstrEncoding{}))>
using CWarpDstr = remove_cvref_t<decltype(make_static_tile_distribution(CWarpDstrEncoding{}))>
using AWarpTensor = static_distributed_tensor<ADataType, AWarpDstr>
using BWarpTensor = static_distributed_tensor<BDataType, BWarpDstr>
using CWarpTensor = static_distributed_tensor<CDataType, CWarpDstr>

Public Member Functions

template<typename CTensor, typename ATensor, typename BTensor, bool post_nop_ = false>
CK_TILE_DEVICE void operator() (CTensor &c, const ATensor &a, const BTensor &b, bool_constant< post_nop_ >={}) const
template<typename CTensor, typename ATensor, typename BTensor, index_t i_subk, bool post_nop_ = false>
CK_TILE_DEVICE void operator() (CTensor &c, const ATensor &a, const BTensor &b, number< i_subk >, bool_constant< post_nop_ >={}) const
template<index_t opselA, index_t opselB, typename CTensor, typename ATensor, typename BTensor, bool post_nop_ = false>
CK_TILE_DEVICE void operator() (CTensor &c, const ATensor &a, const BTensor &b, const int32_t &a_scale, const int32_t &b_scale, bool_constant< post_nop_ >={}) const
template<typename ATensor, typename BTensor>
CK_TILE_DEVICE auto operator() (const ATensor &a, const BTensor &b) const
template<index_t opselA, index_t opselB, typename ATensor, typename BTensor>
CK_TILE_DEVICE auto operator() (const ATensor &a, const BTensor &b, const int32_t &a_scale, const int32_t &b_scale) const

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr auto get_num_of_access ()

Static Public Attributes

static constexpr index_t kM = WarpGemmAttribute::kM
static constexpr index_t kN = WarpGemmAttribute::kN
static constexpr index_t kK = WarpGemmAttribute::kK
static constexpr index_t kCMLane = WarpGemmAttribute::kCMLane
static constexpr index_t kKPerThread = WarpGemmAttribute::kKPerThread
 The number of elements in K dimension processed by single thread in wavefront.

Member Typedef Documentation

◆ ADataType

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::ADataType = typename WarpGemmAttribute::ADataType

◆ AWarpDstr

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::AWarpDstr = remove_cvref_t<decltype(make_static_tile_distribution(AWarpDstrEncoding{}))>

◆ AWarpDstrEncoding

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::AWarpDstrEncoding = typename WarpGemmAttribute::AWarpDstrEncoding

◆ AWarpTensor

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::AWarpTensor = static_distributed_tensor<ADataType, AWarpDstr>

◆ BDataType

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::BDataType = typename WarpGemmAttribute::BDataType

◆ BWarpDstr

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::BWarpDstr = remove_cvref_t<decltype(make_static_tile_distribution(BWarpDstrEncoding{}))>

◆ BWarpDstrEncoding

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::BWarpDstrEncoding = typename WarpGemmAttribute::BWarpDstrEncoding

◆ BWarpTensor

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::BWarpTensor = static_distributed_tensor<BDataType, BWarpDstr>

◆ CDataType

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::CDataType = typename WarpGemmAttribute::CDataType

◆ CWarpDstr

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::CWarpDstr = remove_cvref_t<decltype(make_static_tile_distribution(CWarpDstrEncoding{}))>

◆ CWarpDstrEncoding

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::CWarpDstrEncoding = typename WarpGemmAttribute::CWarpDstrEncoding

◆ CWarpTensor

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::CWarpTensor = static_distributed_tensor<CDataType, CWarpDstr>

◆ WarpGemmAttribute

template<typename WarpGemmAttribute_>
using ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::WarpGemmAttribute = remove_cvref_t<WarpGemmAttribute_>

Member Function Documentation

◆ get_num_of_access()

template<typename WarpGemmAttribute_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::get_num_of_access ( )
inlinestaticconstexpr

◆ operator()() [1/5]

template<typename WarpGemmAttribute_>
template<typename ATensor, typename BTensor>
CK_TILE_DEVICE auto ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::operator() ( const ATensor & a,
const BTensor & b ) const
inline

◆ operator()() [2/5]

template<typename WarpGemmAttribute_>
template<index_t opselA, index_t opselB, typename ATensor, typename BTensor>
CK_TILE_DEVICE auto ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::operator() ( const ATensor & a,
const BTensor & b,
const int32_t & a_scale,
const int32_t & b_scale ) const
inline

◆ operator()() [3/5]

template<typename WarpGemmAttribute_>
template<typename CTensor, typename ATensor, typename BTensor, bool post_nop_ = false>
CK_TILE_DEVICE void ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::operator() ( CTensor & c,
const ATensor & a,
const BTensor & b,
bool_constant< post_nop_ > = {} ) const
inline

◆ operator()() [4/5]

template<typename WarpGemmAttribute_>
template<index_t opselA, index_t opselB, typename CTensor, typename ATensor, typename BTensor, bool post_nop_ = false>
CK_TILE_DEVICE void ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::operator() ( CTensor & c,
const ATensor & a,
const BTensor & b,
const int32_t & a_scale,
const int32_t & b_scale,
bool_constant< post_nop_ > = {} ) const
inline

◆ operator()() [5/5]

template<typename WarpGemmAttribute_>
template<typename CTensor, typename ATensor, typename BTensor, index_t i_subk, bool post_nop_ = false>
CK_TILE_DEVICE void ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::operator() ( CTensor & c,
const ATensor & a,
const BTensor & b,
number< i_subk > ,
bool_constant< post_nop_ > = {} ) const
inline

Member Data Documentation

◆ kCMLane

template<typename WarpGemmAttribute_>
index_t ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::kCMLane = WarpGemmAttribute::kCMLane
staticconstexpr

◆ kK

template<typename WarpGemmAttribute_>
index_t ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::kK = WarpGemmAttribute::kK
staticconstexpr

◆ kKPerThread

template<typename WarpGemmAttribute_>
index_t ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::kKPerThread = WarpGemmAttribute::kKPerThread
staticconstexpr

The number of elements in K dimension processed by single thread in wavefront.

Note
Note that WarpGemm may run MFMA instruction multiple times (on different K). In such situation this value reflects this fact.

◆ kM

template<typename WarpGemmAttribute_>
index_t ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::kM = WarpGemmAttribute::kM
staticconstexpr

◆ kN

template<typename WarpGemmAttribute_>
index_t ck_tile::WarpGemmImpl< WarpGemmAttribute_ >::kN = WarpGemmAttribute::kN
staticconstexpr
The documentation for this struct was generated from the following file: