EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait > Struct Template Reference

EpilogueReduceCShuffle&lt; DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait &gt; Struct Template Reference#

Composable Kernel: ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Public Attributes | Static Public Attributes | List of all members
ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait > Struct Template Reference

#include <epilogue_cshuffle_v3_reduce_wmma.hpp>

Inheritance diagram for ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >:
ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >

Public Types

using Base = EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >
 
using ReduceGridDesc_M = decltype(MakeReduceGridDescriptor_M(1))
 
- Public Types inherited from ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >
using SpaceFillingCurveVgpr = SpaceFillingCurve< Sequence< MRepeat, 1, 1, NRepeat, 1, 1, BlockwiseGemmPipe::MAccVgprs >, Sequence< 0, 1, 2, 3, 4, 5, 6 >, Sequence< CShuffleMRepeatPerShuffle, 1, 1, CShuffleNRepeatPerShuffle, 1, 1, BlockwiseGemmPipe::MAccVgprs > >
 
using SpaceFillingCurveVmem = SpaceFillingCurve< Sequence< 1, MPerBlock, 1, NPerBlock >, Sequence< 0, 2, 1, 3 >, Sequence< 1, CShuffleMRepeatPerShuffle *BlockwiseGemmPipe::MWaves *MPerWmma, 1, CShuffleNRepeatPerShuffle *BlockwiseGemmPipe::NWaves *NPerWmma > >
 

Public Member Functions

__device__ EpilogueReduceCShuffle (typename ReduceTrait::ReducePtrsGlobal_ p_reduces_grid_, const typename ReduceTrait::ReduceInElementwiseOperations_ reduce_in_element_ops_, const typename ReduceTrait::ReduceAccElementwiseOperations_ reduce_out_element_ops_, const index_t MRaw_)
 
template<InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename CThreadBuf , typename DsGridPointer , typename DsGridDesc_MBlock_MPerBlock_NBlock_NPerBlock , typename EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock >
__device__ void Run (CThreadBuf &c_thread_buf, DsGridPointer p_ds_grid, EDataType *p_e_grid, void *p_shared, const DsGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &ds_grid_desc_mblock_mperblock_nblock_nperblock, const EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &e_grid_desc_mblock_mperblock_nblock_nperblock, CDEElementwiseOperation &cde_element_op, const index_t &block_m_id, const index_t &block_n_id)
 

Static Public Member Functions

static __device__ auto MakeReduceGridDescriptor_M (index_t MRaw)
 
static constexpr __device__ auto MakeReduceGridDescriptor_MBlock_MPerBlock (const ReduceGridDesc_M &d_grid_desc_m)
 
static constexpr __device__ auto GetCShuffleBlockDescriptor_MShRepeat_MPerShRepeat_NShRepeat_NPerShRepeat ()
 
static constexpr __device__ auto GetCShuffleLDSDescriptor ()
 
static __device__ auto GetVgprToLDSEpilogueDescriptor ()
 
- Static Public Member Functions inherited from ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >
static constexpr __device__ auto GetCShuffleBlockDescriptor_MShRepeat_MPerShRepeat_NShRepeat_NPerShRepeat ()
 
static constexpr __device__ auto GetCShuffleLDSDescriptor ()
 
static __device__ auto GetVgprToLDSEpilogueDescriptor ()
 
template<InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename InterDataType , typename CDsDescRefs , typename EGridDesc >
static __device__ auto GetLDSToVmemEpilogueDescriptor (CDsDescRefs &c_ds_desc_refs, EGridDesc &e_grid_desc_mblock_mperblock_nblock_nperblock, CDEElementwiseOperation &cde_element_op, const index_t &block_m_id, const index_t &block_n_id)
 

Public Attributes

ReduceTrait::ReducePtrsGlobal_ p_reduces_grid
 
ReduceTrait::ReduceInElementwiseOperations_ reduce_in_element_ops
 
ReduceTrait::ReduceAccElementwiseOperations_ reduce_out_element_ops
 
index_t MRaw
 
ReduceGridDesc_M reduce_grid_desc_m
 

Static Public Attributes

static constexpr auto I0
 
static constexpr auto I1
 
static constexpr auto I3
 
static constexpr index_t NumDTensor
 
- Static Public Attributes inherited from ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >
static constexpr auto I0 = Number<0>{}
 
static constexpr auto I1 = Number<1>{}
 
static constexpr auto I2 = Number<2>{}
 
static constexpr auto I3 = Number<3>{}
 
static constexpr auto I4 = Number<4>{}
 
static constexpr auto I5 = Number<5>{}
 
static constexpr auto I6 = Number<6>{}
 
static constexpr index_t NumDTensor = DsDataType::Size()
 
static constexpr auto EShuffleBlockTransferScalarPerVector
 

Member Typedef Documentation

◆ Base

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
using ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::Base = EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe>

◆ ReduceGridDesc_M

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
using ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::ReduceGridDesc_M = decltype(MakeReduceGridDescriptor_M(1))

Constructor & Destructor Documentation

◆ EpilogueReduceCShuffle()

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
__device__ ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::EpilogueReduceCShuffle ( typename ReduceTrait::ReducePtrsGlobal_  p_reduces_grid_,
const typename ReduceTrait::ReduceInElementwiseOperations_  reduce_in_element_ops_,
const typename ReduceTrait::ReduceAccElementwiseOperations_  reduce_out_element_ops_,
const index_t  MRaw_ 
)
inline

Member Function Documentation

◆ GetCShuffleBlockDescriptor_MShRepeat_MPerShRepeat_NShRepeat_NPerShRepeat()

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
static constexpr __device__ auto ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >::GetCShuffleBlockDescriptor_MShRepeat_MPerShRepeat_NShRepeat_NPerShRepeat
inlinestaticconstexpr

◆ GetCShuffleLDSDescriptor()

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
static constexpr __device__ auto ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >::GetCShuffleLDSDescriptor
inlinestaticconstexpr

◆ GetVgprToLDSEpilogueDescriptor()

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
static __device__ auto ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >::GetVgprToLDSEpilogueDescriptor
inlinestatic

◆ MakeReduceGridDescriptor_M()

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
static __device__ auto ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::MakeReduceGridDescriptor_M ( index_t  MRaw)
inlinestatic

◆ MakeReduceGridDescriptor_MBlock_MPerBlock()

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
static constexpr __device__ auto ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::MakeReduceGridDescriptor_MBlock_MPerBlock ( const ReduceGridDesc_M &  d_grid_desc_m)
inlinestaticconstexpr

◆ Run()

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
template<InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename CThreadBuf , typename DsGridPointer , typename DsGridDesc_MBlock_MPerBlock_NBlock_NPerBlock , typename EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock >
__device__ void ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::Run ( CThreadBuf &  c_thread_buf,
DsGridPointer  p_ds_grid,
EDataType *  p_e_grid,
void *  p_shared,
const DsGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &  ds_grid_desc_mblock_mperblock_nblock_nperblock,
const EGridDesc_MBlock_MPerBlock_NBlock_NPerBlock &  e_grid_desc_mblock_mperblock_nblock_nperblock,
CDEElementwiseOperation &  cde_element_op,
const index_t &  block_m_id,
const index_t &  block_n_id 
)
inline

Member Data Documentation

◆ I0

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
constexpr auto ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >::I0
staticconstexpr

◆ I1

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
constexpr auto ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >::I1
staticconstexpr

◆ I3

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
constexpr auto ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >::I3
staticconstexpr

◆ MRaw

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
index_t ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::MRaw

◆ NumDTensor

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
constexpr index_t ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >::NumDTensor
staticconstexpr

◆ p_reduces_grid

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
ReduceTrait::ReducePtrsGlobal_ ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::p_reduces_grid

◆ reduce_grid_desc_m

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
ReduceGridDesc_M ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::reduce_grid_desc_m

◆ reduce_in_element_ops

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
ReduceTrait::ReduceInElementwiseOperations_ ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::reduce_in_element_ops

◆ reduce_out_element_ops

template<typename DsDataType , typename EDataType , typename AccDataType , typename CShuffleDataType , index_t MPerBlock, index_t NPerBlock, index_t MPerWmma, index_t NPerWmma, index_t MRepeat, index_t NRepeat, index_t CShuffleMRepeatPerShuffle, index_t CShuffleNRepeatPerShuffle, typename CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock , typename CDEShuffleBlockTransferScalarPerVectors , typename CDEElementwiseOperation , typename ThisThreadBlock , typename BlockwiseGemmPipe , tensor_operation::device::GemmSpecialization GemmSpec, index_t BlockSize, typename ReduceTrait >
ReduceTrait::ReduceAccElementwiseOperations_ ck::EpilogueReduceCShuffle< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe, GemmSpec, BlockSize, ReduceTrait >::reduce_out_element_ops
The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/tensor_operation/gpu/grid/epilogue_cshuffle_v3_reduce_wmma.hpp