CShuffleEpilogue< Problem_, Policy_ > Struct Template Reference

CShuffleEpilogue&lt; Problem_, Policy_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::CShuffleEpilogue< Problem_, Policy_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::CShuffleEpilogue< Problem_, Policy_ > Struct Template Reference

#include <cshuffle_epilogue.hpp>

Public Types

using Problem = remove_cvref_t< Problem_ >
 
using AccDataType = remove_cvref_t< typename Problem::AccDataType >
 
using ODataType = remove_cvref_t< typename Problem::ODataType >
 
using CLayout = remove_cvref_t< typename Problem::CLayout >
 
using WG = WarpGemmMfmaDispatcher< ODataType, ODataType, AccDataType, kMPerXdl, kNPerXdl, kKPerXdl, isCTransposed >
 
using CWarpDstr = typename WG::CWarpDstr
 
using CWarpTensor = typename WG::CWarpTensor
 

Public Member Functions

template<typename ODramWindow , typename OAccTile , memory_operation_enum out_memory_data_op = memory_operation_enum::set>
CK_TILE_DEVICE auto operator() (ODramWindow &out_dram_window, const OAccTile &o_acc_tile, void *p_smem)
 

Static Public Member Functions

template<typename ODataType >
static constexpr CK_TILE_HOST_DEVICE auto GetVectorSizeC ()
 Get the vector store size for C tensor. More...
 
template<typename Problem >
static constexpr CK_TILE_HOST_DEVICE auto MakeLdsBlockDescriptor ()
 
static constexpr CK_TILE_HOST_DEVICE index_t GetSmemSize ()
 

Static Public Attributes

static constexpr index_t kBlockSize = Problem::kBlockSize
 
static constexpr index_t kMPerBlock = Problem::kMPerBlock
 
static constexpr index_t kNPerBlock = Problem::kNPerBlock
 
static constexpr index_t kMWave = Problem::kMWave
 
static constexpr index_t kNWave = Problem::kNWave
 
static constexpr index_t kMPerXdl = Problem::kMPerXdl
 
static constexpr index_t kNPerXdl = Problem::kNPerXdl
 
static constexpr index_t kKPerXdl = Problem::kKPerXdl
 
static constexpr index_t isCTransposed = Problem::isCTransposed
 
static constexpr index_t kMPerIteration = kMPerXdl * kMWave
 
static constexpr index_t kNPerIteration = kNPerXdl * kNWave
 

Member Typedef Documentation

◆ AccDataType

template<typename Problem_ , typename Policy_ = void>
using ck_tile::CShuffleEpilogue< Problem_, Policy_ >::AccDataType = remove_cvref_t<typename Problem::AccDataType>

◆ CLayout

template<typename Problem_ , typename Policy_ = void>
using ck_tile::CShuffleEpilogue< Problem_, Policy_ >::CLayout = remove_cvref_t<typename Problem::CLayout>

◆ CWarpDstr

template<typename Problem_ , typename Policy_ = void>
using ck_tile::CShuffleEpilogue< Problem_, Policy_ >::CWarpDstr = typename WG::CWarpDstr

◆ CWarpTensor

template<typename Problem_ , typename Policy_ = void>
using ck_tile::CShuffleEpilogue< Problem_, Policy_ >::CWarpTensor = typename WG::CWarpTensor

◆ ODataType

template<typename Problem_ , typename Policy_ = void>
using ck_tile::CShuffleEpilogue< Problem_, Policy_ >::ODataType = remove_cvref_t<typename Problem::ODataType>

◆ Problem

template<typename Problem_ , typename Policy_ = void>
using ck_tile::CShuffleEpilogue< Problem_, Policy_ >::Problem = remove_cvref_t<Problem_>

◆ WG

template<typename Problem_ , typename Policy_ = void>
using ck_tile::CShuffleEpilogue< Problem_, Policy_ >::WG = WarpGemmMfmaDispatcher<ODataType, ODataType, AccDataType, kMPerXdl, kNPerXdl, kKPerXdl, isCTransposed>

Member Function Documentation

◆ GetSmemSize()

template<typename Problem_ , typename Policy_ = void>
static constexpr CK_TILE_HOST_DEVICE index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::GetSmemSize ( )
inlinestaticconstexpr

◆ GetVectorSizeC()

template<typename Problem_ , typename Policy_ = void>
template<typename ODataType >
static constexpr CK_TILE_HOST_DEVICE auto ck_tile::CShuffleEpilogue< Problem_, Policy_ >::GetVectorSizeC ( )
inlinestaticconstexpr

Get the vector store size for C tensor.

Note
The vector store size for output C tensor would depend on multiple factors like its data layout and warp gemm C transposition. In general it would be the number of consecutive elements in contiguous C dimension hold by single thread.
Returns
The vector store size for C tensor.

◆ MakeLdsBlockDescriptor()

template<typename Problem_ , typename Policy_ = void>
template<typename Problem >
static constexpr CK_TILE_HOST_DEVICE auto ck_tile::CShuffleEpilogue< Problem_, Policy_ >::MakeLdsBlockDescriptor ( )
inlinestaticconstexpr

◆ operator()()

template<typename Problem_ , typename Policy_ = void>
template<typename ODramWindow , typename OAccTile , memory_operation_enum out_memory_data_op = memory_operation_enum::set>
CK_TILE_DEVICE auto ck_tile::CShuffleEpilogue< Problem_, Policy_ >::operator() ( ODramWindow &  out_dram_window,
const OAccTile &  o_acc_tile,
void *  p_smem 
)
inline

Member Data Documentation

◆ isCTransposed

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::isCTransposed = Problem::isCTransposed
staticconstexpr

◆ kBlockSize

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::kBlockSize = Problem::kBlockSize
staticconstexpr

◆ kKPerXdl

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::kKPerXdl = Problem::kKPerXdl
staticconstexpr

◆ kMPerBlock

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::kMPerBlock = Problem::kMPerBlock
staticconstexpr

◆ kMPerIteration

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::kMPerIteration = kMPerXdl * kMWave
staticconstexpr

◆ kMPerXdl

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::kMPerXdl = Problem::kMPerXdl
staticconstexpr

◆ kMWave

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::kMWave = Problem::kMWave
staticconstexpr

◆ kNPerBlock

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::kNPerBlock = Problem::kNPerBlock
staticconstexpr

◆ kNPerIteration

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::kNPerIteration = kNPerXdl * kNWave
staticconstexpr

◆ kNPerXdl

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::kNPerXdl = Problem::kNPerXdl
staticconstexpr

◆ kNWave

template<typename Problem_ , typename Policy_ = void>
constexpr index_t ck_tile::CShuffleEpilogue< Problem_, Policy_ >::kNWave = Problem::kNWave
staticconstexpr
The documentation for this struct was generated from the following file: