GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector > Struct Template Reference

GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector > Struct Template Reference#

Composable Kernel: ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector > Struct Template Reference

#include <gridwise_permute.hpp>

Classes
struct	Block2TileMap

Public Types
using	ThisThreadBlock = ThisThreadBlock<BlockSize>
using	DefaultBlock2TileMap = Block2TileMap

Static Public Member Functions
__host__ static __device__ constexpr auto	GetInBlockDesc_NPerBlock_HPerBlock_WPerBlock ()
template<typename GridDesc>
__host__ static __device__ constexpr auto	GetMergedDesc (const GridDesc &desc)
__host__ static __device__ constexpr index_t	GetSharedMemoryNumberOfByte ()
__host__ static __device__ constexpr auto	MakeDefaultBlock2TileMap (const InGridDesc &desc)
__host__ static __device__ constexpr bool	CheckValidity (const InGridDesc &in_grid_desc, const OutGridDesc &out_grid_desc)
template<typename Block2TileMap>
static __device__ void	Run (const InGridDesc in_grid_desc, const OutGridDesc out_grid_desc, const InDataType p_in_global, OutDataType p_out_global, void *__restrict__ p_shared, const ElementwiseOperation elementwise_op, const Block2TileMap &block_2_tile_map)

Static Public Attributes
static constexpr auto	I0 = Number<0>{}
static constexpr auto	I1 = Number<1>{}
static constexpr auto	I2 = Number<2>{}

Member Typedef Documentation

◆ DefaultBlock2TileMap

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

using ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::DefaultBlock2TileMap = Block2TileMap

◆ ThisThreadBlock

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

using ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::ThisThreadBlock = ThisThreadBlock<BlockSize>

Member Function Documentation

◆ CheckValidity()

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

__host__ static __device__ constexpr bool ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::CheckValidity	(	const InGridDesc &	in_grid_desc,
		const OutGridDesc &	out_grid_desc )

inlinestaticconstexpr

◆ GetInBlockDesc_NPerBlock_HPerBlock_WPerBlock()

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

__host__ static __device__ constexpr auto ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::GetInBlockDesc_NPerBlock_HPerBlock_WPerBlock ( )

inlinestaticconstexpr

◆ GetMergedDesc()

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

template<typename GridDesc>

__host__ static __device__ constexpr auto ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::GetMergedDesc ( const GridDesc & desc )

inlinestaticconstexpr

◆ GetSharedMemoryNumberOfByte()

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

__host__ static __device__ constexpr index_t ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::GetSharedMemoryNumberOfByte ( )

inlinestaticconstexpr

◆ MakeDefaultBlock2TileMap()

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

__host__ static __device__ constexpr auto ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::MakeDefaultBlock2TileMap ( const InGridDesc & desc )

inlinestaticconstexpr

◆ Run()

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

template<typename Block2TileMap>

__device__ void ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::Run	(	const InGridDesc	in_grid_desc,
		const OutGridDesc	out_grid_desc,
		const InDataType *	p_in_global,
		OutDataType *	p_out_global,
		void *__restrict__	p_shared,
		const ElementwiseOperation	elementwise_op,
		const Block2TileMap &	block_2_tile_map )

inlinestatic

Member Data Documentation

◆ I0

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

auto ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::I0 = Number<0>{}

staticconstexpr

◆ I1

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

auto ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::I1 = Number<1>{}

staticconstexpr

◆ I2

template<typename InGridDesc, typename OutGridDesc, typename InDataType, typename OutDataType, typename ElementwiseOperation, index_t BlockSize, index_t NPerBlock, index_t HPerBlock, index_t WPerBlock, index_t InBlockLdsExtraW, typename InBlockTransferThreadClusterLengths, typename InBlockTransferThreadClusterArrangeOrder, index_t SrcVectorDim, index_t DstVectorDim, index_t SrcScalarPerVector, index_t DstScalarPerVector>

auto ck::GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim, DstVectorDim, SrcScalarPerVector, DstScalarPerVector >::I2 = Number<2>{}

staticconstexpr

The documentation for this struct was generated from the following file:

gridwise_permute.hpp

Classes

Public Types

Static Public Member Functions

Static Public Attributes

Member Typedef Documentation

◆ DefaultBlock2TileMap

◆ ThisThreadBlock

Member Function Documentation

◆ CheckValidity()

◆ GetInBlockDesc_NPerBlock_HPerBlock_WPerBlock()

◆ GetMergedDesc()

◆ GetSharedMemoryNumberOfByte()

◆ MakeDefaultBlock2TileMap()

◆ Run()

Member Data Documentation

◆ I0

◆ I1

◆ I2