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 = 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 > > |
Static Public Member Functions | |
| 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) |
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 = 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> > |
◆ 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 = SpaceFillingCurve< Sequence<1, MPerBlock, 1, NPerBlock>, Sequence<0, 2, 1, 3>, 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:
- /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/tensor_operation/gpu/grid/epilogue_cshuffle_v3_wmma_base.hpp