ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun > Struct Template Reference

ABTransferThreadTiles&lt; ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun &gt; Struct Template Reference#

Composable Kernel: ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun > Struct Template Reference

#include <gridwise_ab_transfer_thread_tiles.hpp>

Public Types

using ThisThreadBlock = ThisThreadBlock<BlockSize>

Static Public Member Functions

template<bool PadMN, bool PadK, typename GridDescriptorBase>
__host__ static __device__ auto MakeGridDescriptor (const GridDescriptorBase &ab_grid_desc, index_t MN, index_t MNPad, index_t K, index_t KPad, index_t StrideAB, index_t ABK0)
static __device__ constexpr auto GetBlockDescriptor ()
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 ABK0Number = Number<KPerBlock / ABK1Value>{}
static constexpr auto ABK1Number = Number<ABK1Value>{}
static constexpr auto I0 = Number<0>{}
static constexpr auto I1 = Number<1>{}
static constexpr auto I2 = Number<2>{}
static constexpr index_t ABPackedSize

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 ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
using ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::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 ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
__device__ constexpr auto ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::GetBlockDescriptor ( )
inlinestaticconstexpr

◆ GetBlockStep()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
__device__ constexpr auto ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::GetBlockStep ( )
inlinestaticconstexpr

◆ GetBlockTransfer()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
template<typename GridDescriptor, typename BlockDescriptor, typename ABsDataType, typename ABElementwiseOperation, index_t GlobalBufferNum>
__device__ auto ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::GetBlockTransfer ( GridDescriptor & grid_descriptor,
BlockDescriptor & block_descriptor,
ABElementwiseOperation & ab_element_op,
const index_t block_mn_id )
inlinestatic

◆ GetKDimension()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
template<typename GridDescriptor>
__device__ constexpr index_t ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::GetKDimension ( const GridDescriptor & grid_desc)
inlinestaticconstexpr

◆ MakeGridDescriptor()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
template<bool PadMN, bool PadK, typename GridDescriptorBase>
__host__ static __device__ auto ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::MakeGridDescriptor ( const GridDescriptorBase & ab_grid_desc,
index_t MN,
index_t MNPad,
index_t K,
index_t KPad,
index_t StrideAB,
index_t ABK0 )
inlinestatic

◆ MakeWmmaTileDescriptor()

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
template<index_t MNRepeat, index_t MNWaves>
__host__ static __device__ constexpr auto ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::MakeWmmaTileDescriptor ( )
inlinestaticconstexpr

Member Data Documentation

◆ ABK0Number

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
auto ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::ABK0Number = Number<KPerBlock / ABK1Value>{}
staticconstexpr

◆ ABK1Number

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
auto ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::ABK1Number = Number<ABK1Value>{}
staticconstexpr

◆ ABPackedSize

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
index_t ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::ABPackedSize
staticconstexpr
Initial value:
= []() {
if constexpr(is_same_v<remove_cvref_t<LDSTypeAB>, pk_i4_t>)
return 2;
else
return 1;
}()
ck::is_same_v
constexpr bool is_same_v
Definition type.hpp:283
ck::pk_i4_t
Definition data_type.hpp:187

◆ I0

template<typename ABLayout, typename ABMajorLayout, typename LDSTypeAB, index_t BlockSize, index_t MNPerBlock, index_t KPerBlock, index_t MNPerWmma, index_t ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
auto ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::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 ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
auto ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::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 ABK1Value, bool UseBlockPaddingAB, bool PermuteAB, typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, typename ABBlockTransferThreadClusterArrangeOrder, typename ABBlockTransferSrcAccessOrder, index_t ABBlockTransferSrcVectorDim, index_t ABBlockTransferSrcScalarPerVector, index_t ABBlockTransferDstScalarPerVector_ABK1, bool ABThreadTransferSrcResetCoordinateAfterRun>
auto ck::ABTransferThreadTiles< ABLayout, ABMajorLayout, LDSTypeAB, BlockSize, MNPerBlock, KPerBlock, MNPerWmma, ABK1Value, UseBlockPaddingAB, PermuteAB, ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1, ABBlockTransferThreadClusterArrangeOrder, ABBlockTransferSrcAccessOrder, ABBlockTransferSrcVectorDim, ABBlockTransferSrcScalarPerVector, ABBlockTransferDstScalarPerVector_ABK1, ABThreadTransferSrcResetCoordinateAfterRun >::I2 = Number<2>{}
staticconstexpr
The documentation for this struct was generated from the following file: