EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe > Struct Template Reference#
ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe > Struct Template Reference
#include <epilogue_cshuffle_v3_wmma_base.hpp>
Inheritance diagram for 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 | SpaceFillingCurveVgpr |
| using | SpaceFillingCurveVmem |
Static Public Member Functions | |
| static __device__ constexpr auto | GetCShuffleBlockDescriptor_MShRepeat_MPerShRepeat_NShRepeat_NPerShRepeat () |
| static __device__ constexpr 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) |
Static Public Attributes | |
| 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
◆ SpaceFillingCurveVgpr
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>
| using ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >::SpaceFillingCurveVgpr |
Initial value:
Sequence<CShuffleMRepeatPerShuffle,
1,
1,
CShuffleNRepeatPerShuffle,
1,
1,
BlockwiseGemmPipe::MAccVgprs>>
Definition utility/sequence.hpp:43
Definition tensor_space_filling_curve.hpp:20
◆ SpaceFillingCurveVmem
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>
| using ck::EpilogueCShuffleBase< DsDataType, EDataType, AccDataType, CShuffleDataType, MPerBlock, NPerBlock, MPerWmma, NPerWmma, MRepeat, NRepeat, CShuffleMRepeatPerShuffle, CShuffleNRepeatPerShuffle, CDEShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CDEShuffleBlockTransferScalarPerVectors, CDEElementwiseOperation, ThisThreadBlock, BlockwiseGemmPipe >::SpaceFillingCurveVmem |
Initial value:
Sequence<1,
CShuffleMRepeatPerShuffle * BlockwiseGemmPipe::MWaves * MPerWmma,
1,
CShuffleNRepeatPerShuffle * BlockwiseGemmPipe::NWaves * NPerWmma>>
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>
|
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>
|
inlinestaticconstexpr |
◆ GetLDSToVmemEpilogueDescriptor()
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>
template<InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename InterDataType, typename CDsDescRefs, typename EGridDesc>
|
inlinestatic |
◆ 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>
|
inlinestatic |
Member Data Documentation
◆ EShuffleBlockTransferScalarPerVector
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>
|
staticconstexpr |
Initial value:
=
CDEShuffleBlockTransferScalarPerVectors{}[I0]
static constexpr auto I0
Definition epilogue_cshuffle_v3_wmma_base.hpp:30
◆ 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>
|
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>
|
staticconstexpr |
◆ I2
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>
|
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>
|
staticconstexpr |
◆ I4
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>
|
staticconstexpr |
◆ I5
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>
|
staticconstexpr |
◆ I6
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>
|
staticconstexpr |
◆ 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>
|
staticconstexpr |
The documentation for this struct was generated from the following file: