ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize > Struct Template Reference

ABTransferWaveTiles&lt; ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize &gt; Struct Template Reference#

Composable Kernel: ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize > Struct Template Reference

#include <gridwise_ab_transfer_wave_tiles.hpp>

Public Types

using ThisThreadBlock = ThisThreadBlock<BlockSize>

Static Public Member Functions

template<bool PadMN, bool PadK, typename GridDescriptorBase>
__host__ static __device__ auto MakeGridDescriptor (GridDescriptorBase &base_desc, index_t sizeMN, index_t, index_t sizeK, index_t, index_t, index_t)
static __device__ constexpr auto GetBlockDescriptor ()
static __device__ auto GetWaveIdx ()
static __device__ auto GetBlockLaneIdx ()
template<typename ABDataType>
static __device__ auto GetGridLaneIdx ()
template<typename GridDescriptor, typename BlockDescriptor, typename ABsDataType, typename ABElementwiseOperation, index_t GlobalBufferNum>
static __device__ auto GetBlockTransfer (GridDescriptor &grid_descriptor, BlockDescriptor &block_descriptor, ABElementwiseOperation &ab_element_op, const index_t block_mn_id)
template<index_t MNRepeat, index_t MNWaves>
__host__ static __device__ constexpr auto MakeWmmaTileDescriptor ()
static __device__ constexpr auto GetBlockStep ()
template<typename GridDescriptor>
static __device__ constexpr index_t GetKDimension (const GridDescriptor &grid_desc)

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 index_t MNKRow = 2
static constexpr index_t NumberOfWaves = BlockSize / WaveSize
static constexpr index_t MNMajorWaves_
static constexpr index_t KMajorWaves_
static constexpr bool ABDoTranspose = !is_same_v<ABLayout, ABMajorLayout>
static constexpr index_t MNWaves_
static constexpr index_t KWaves_ = ABDoTranspose ? KMajorWaves_ : NumberOfWaves / MNMajorWaves_
static constexpr index_t KRepeat_ = KPerBlock / (KWaves_ * KPack)
static constexpr index_t MNRepeat_ = MNPerBlock / (MNWaves_ * MNPerWmma)

Member Typedef Documentation

◆ ThisThreadBlock

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
using ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::ThisThreadBlock = ThisThreadBlock<BlockSize>

Member Function Documentation

◆ GetBlockDescriptor()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
__device__ constexpr auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::GetBlockDescriptor ( )
inlinestaticconstexpr

◆ GetBlockLaneIdx()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
__device__ auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::GetBlockLaneIdx ( )
inlinestatic

◆ GetBlockStep()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
__device__ constexpr auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::GetBlockStep ( )
inlinestaticconstexpr

◆ GetBlockTransfer()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
template<typename GridDescriptor, typename BlockDescriptor, typename ABsDataType, typename ABElementwiseOperation, index_t GlobalBufferNum>
__device__ auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::GetBlockTransfer ( GridDescriptor & grid_descriptor,
BlockDescriptor & block_descriptor,
ABElementwiseOperation & ab_element_op,
const index_t block_mn_id )
inlinestatic

◆ GetGridLaneIdx()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
template<typename ABDataType>
__device__ auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::GetGridLaneIdx ( )
inlinestatic

◆ GetKDimension()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
template<typename GridDescriptor>
__device__ constexpr index_t ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::GetKDimension ( const GridDescriptor & grid_desc)
inlinestaticconstexpr

◆ GetWaveIdx()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
__device__ auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::GetWaveIdx ( )
inlinestatic

◆ MakeGridDescriptor()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
template<bool PadMN, bool PadK, typename GridDescriptorBase>
__host__ static __device__ auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::MakeGridDescriptor ( GridDescriptorBase & base_desc,
index_t sizeMN,
index_t ,
index_t sizeK,
index_t ,
index_t ,
index_t  )
inlinestatic

◆ MakeWmmaTileDescriptor()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
template<index_t MNRepeat, index_t MNWaves>
__host__ static __device__ constexpr auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::MakeWmmaTileDescriptor ( )
inlinestaticconstexpr

Member Data Documentation

◆ ABDoTranspose

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
bool ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::ABDoTranspose = !is_same_v<ABLayout, ABMajorLayout>
staticconstexpr

◆ I0

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::I0 = Number<0>{}
staticconstexpr

◆ I1

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::I1 = Number<1>{}
staticconstexpr

◆ I2

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::I2 = Number<2>{}
staticconstexpr

◆ I3

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
auto ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::I3 = Number<3>{}
staticconstexpr

◆ KMajorWaves_

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
index_t ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::KMajorWaves_
staticconstexpr
Initial value:
=
KPerBlock / KPack % std::min(KPerBlock / KPack, NumberOfWaves) == 0
? std::min(KPerBlock / KPack, NumberOfWaves)
: (KPerBlock / KPack % 2 == 0 ? 2 : 1)
ck::ABTransferWaveTiles::NumberOfWaves
static constexpr index_t NumberOfWaves
Definition gridwise_ab_transfer_wave_tiles.hpp:60

◆ KRepeat_

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
index_t ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::KRepeat_ = KPerBlock / (KWaves_ * KPack)
staticconstexpr

◆ KWaves_

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
index_t ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::KWaves_ = ABDoTranspose ? KMajorWaves_ : NumberOfWaves / MNMajorWaves_
staticconstexpr

◆ MNKRow

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
index_t ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::MNKRow = 2
staticconstexpr

◆ MNMajorWaves_

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
index_t ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::MNMajorWaves_
staticconstexpr
Initial value:
=
MNPerBlock / MNPerWmma % std::min(MNPerBlock / MNPerWmma, NumberOfWaves) == 0
? std::min(MNPerBlock / MNPerWmma, NumberOfWaves)
: (MNPerBlock / MNPerWmma % 2 == 0 ? 2 : 1)

◆ MNRepeat_

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
index_t ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::MNRepeat_ = MNPerBlock / (MNWaves_ * MNPerWmma)
staticconstexpr

◆ MNWaves_

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
index_t ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::MNWaves_
staticconstexpr
Initial value:
=
ABDoTranspose ? NumberOfWaves / KMajorWaves_ : MNMajorWaves_
ck::ABTransferWaveTiles::KMajorWaves_
static constexpr index_t KMajorWaves_
Definition gridwise_ab_transfer_wave_tiles.hpp:65
ck::ABTransferWaveTiles::MNMajorWaves_
static constexpr index_t MNMajorWaves_
Definition gridwise_ab_transfer_wave_tiles.hpp:61
ck::ABTransferWaveTiles::ABDoTranspose
static constexpr bool ABDoTranspose
Definition gridwise_ab_transfer_wave_tiles.hpp:70

◆ NumberOfWaves

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t KPack, index_t ABK1Value, index_t WaveSize>
index_t ck::ABTransferWaveTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, KPack, ABK1Value, WaveSize >::NumberOfWaves = BlockSize / WaveSize
staticconstexpr
The documentation for this struct was generated from the following file: