GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize > Struct Template Reference

GridwiseReduction_mk_to_m_multiblock&lt; InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize &gt; Struct Template Reference#

Composable Kernel: ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize > Struct Template Reference

#include <gridwise_2d_reduction_multiblock.hpp>

Public Types

using ThreadClusterLengths_M_K = Sequence<MThreadClusterSize, KThreadClusterSize>
using ThreadBufferDimAccessOrder
using ThreadClusterArrangeOrder
using ThreadReduceSrcDesc_M_K
using ThreadReduceDstDesc_M
using BlockwiseReduce
using ThreadwiseReduce
using PassThroughOp = tensor_operation::element_wise::PassThrough
using Accumulation = detail::AccumulateWithNanCheck<PropagateNan, ReduceOperation, AccDataType>

Static Public Member Functions

static __device__ void Run (const InGridDesc_M_K &in_grid_desc_m_k, const OutGridDesc_M &out_grid_desc_m, const InElementwiseOperation &in_elementwise_op, const AccElementwiseOperation &acc_elementwise_op, index_t block_group_size, index_t num_k_block_tile_iteration, AccDataType alpha, const InDataType *const __restrict__ p_in_value_global, AccDataType beta, OutDataType *const __restrict__ p_out_value_global)
template<bool HaveIndexInput>
static __device__ void RunWithIndex (const InGridDesc_M_K &in_grid_desc_m_k, const OutGridDesc_M &out_grid_desc_m, const InElementwiseOperation in_elementwise_op, const AccElementwiseOperation acc_elementwise_op, index_t num_k_block_tile_iteration, AccDataType alpha, const InDataType *const __restrict__ p_in_value_global, const IndexDataType *const __restrict__ p_in_index_global, AccDataType beta, OutDataType *const __restrict__ p_out_value_global, IndexDataType *const __restrict__ p_out_index_global)

Static Public Attributes

static constexpr bool reorder_thread_cluster = (InSrcVectorDim == 0)
static constexpr auto thread_cluster_desc
static constexpr auto I0 = Number<0>{}
static constexpr auto I1 = Number<1>{}
static constexpr index_t M_BlockTileSize = MThreadClusterSize * MThreadSliceSize
static constexpr index_t K_BlockTileSize = KThreadClusterSize * KThreadSliceSize

Member Typedef Documentation

◆ Accumulation

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
using ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::Accumulation = detail::AccumulateWithNanCheck<PropagateNan, ReduceOperation, AccDataType>

◆ BlockwiseReduce

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
using ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::BlockwiseReduce
Initial value:
PartitionedBlockwiseReduction<AccDataType,
BlockSize,
ThreadClusterLengths_M_K,
ThreadClusterArrangeOrder,
ReduceOperation,
PropagateNan>
ck::GridwiseReduceSecondHalfBatchNormBackwardFinal::ThreadClusterLengths_M_K
Sequence< MThreadClusterSize, KThreadClusterSize > ThreadClusterLengths_M_K
Definition gridwise_multiblock_reduce_second_half_batchnorm_backward_final.hpp:110
ck::GridwiseReduceSecondHalfBatchNormBackwardFinal::ThreadClusterArrangeOrder
typename conditional< reorder_thread_cluster, Sequence< 1, 0 >, Sequence< 0, 1 > >::type ThreadClusterArrangeOrder
Definition gridwise_multiblock_reduce_second_half_batchnorm_backward_final.hpp:115
ck::PartitionedBlockwiseReduction
Definition reduction_functions_blockwise.hpp:28

◆ PassThroughOp

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
using ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::PassThroughOp = tensor_operation::element_wise::PassThrough

◆ ThreadBufferDimAccessOrder

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
using ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::ThreadBufferDimAccessOrder
Initial value:
typename conditional<reorder_thread_cluster, Sequence<1, 0>, Sequence<0, 1>>::type
ck::Sequence
Definition utility/sequence.hpp:43
ck::conditional
Definition utility/functional.hpp:100

◆ ThreadClusterArrangeOrder

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
using ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::ThreadClusterArrangeOrder
Initial value:
typename conditional<reorder_thread_cluster, Sequence<1, 0>, Sequence<0, 1>>::type

◆ ThreadClusterLengths_M_K

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
using ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::ThreadClusterLengths_M_K = Sequence<MThreadClusterSize, KThreadClusterSize>

◆ ThreadReduceDstDesc_M

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
using ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::ThreadReduceDstDesc_M
Initial value:
decltype(make_naive_tensor_descriptor_packed(make_tuple(Number<MThreadSliceSize>{})))
ck::Number
integral_constant< index_t, N > Number
Definition number.hpp:12
ck::make_naive_tensor_descriptor_packed
__host__ __device__ constexpr auto make_naive_tensor_descriptor_packed(const Tuple< Lengths... > &lengths)
Definition tensor_descriptor_helper.hpp:101
ck::make_tuple
__host__ __device__ constexpr auto make_tuple(Xs &&... xs)
Definition utility/tuple.hpp:211

◆ ThreadReduceSrcDesc_M_K

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
using ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::ThreadReduceSrcDesc_M_K
Initial value:
decltype(make_naive_tensor_descriptor_packed(
make_tuple(Number<MThreadSliceSize>{}, Number<KThreadSliceSize>{})))

◆ ThreadwiseReduce

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
using ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::ThreadwiseReduce
Initial value:
ThreadwiseReduction<AccDataType,
ThreadReduceSrcDesc_M_K,
ThreadReduceDstDesc_M,
ReduceOperation,
PropagateNan>
ck::GridwiseReduceSecondHalfBatchNormBackwardFinal::ThreadReduceDstDesc_M
decltype(make_naive_tensor_descriptor_packed(make_tuple(Number< MThreadSliceSize >{}))) ThreadReduceDstDesc_M
Definition gridwise_multiblock_reduce_second_half_batchnorm_backward_final.hpp:123
ck::GridwiseWelfordSecondHalfReduceFirstHalf::ThreadReduceSrcDesc_M_K
decltype(make_naive_tensor_descriptor_packed( make_tuple(Number< MThreadSliceSize >{}, Number< KThreadSliceSize >{}))) ThreadReduceSrcDesc_M_K
Definition gridwise_multiblock_welford_second_half_multiblock_reduce_first_half.hpp:116
ck::ThreadwiseReduction
Definition reduction_functions_threadwise.hpp:23

Member Function Documentation

◆ Run()

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
__device__ void ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::Run ( const InGridDesc_M_K & in_grid_desc_m_k,
const OutGridDesc_M & out_grid_desc_m,
const InElementwiseOperation & in_elementwise_op,
const AccElementwiseOperation & acc_elementwise_op,
index_t block_group_size,
index_t num_k_block_tile_iteration,
AccDataType alpha,
const InDataType *const __restrict__ p_in_value_global,
AccDataType beta,
OutDataType *const __restrict__ p_out_value_global )
inlinestatic

◆ RunWithIndex()

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
template<bool HaveIndexInput>
__device__ void ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::RunWithIndex ( const InGridDesc_M_K & in_grid_desc_m_k,
const OutGridDesc_M & out_grid_desc_m,
const InElementwiseOperation in_elementwise_op,
const AccElementwiseOperation acc_elementwise_op,
index_t num_k_block_tile_iteration,
AccDataType alpha,
const InDataType *const __restrict__ p_in_value_global,
const IndexDataType *const __restrict__ p_in_index_global,
AccDataType beta,
OutDataType *const __restrict__ p_out_value_global,
IndexDataType *const __restrict__ p_out_index_global )
inlinestatic

Member Data Documentation

◆ I0

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
auto ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::I0 = Number<0>{}
staticconstexpr

◆ I1

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
auto ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::I1 = Number<1>{}
staticconstexpr

◆ K_BlockTileSize

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
index_t ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::K_BlockTileSize = KThreadClusterSize * KThreadSliceSize
staticconstexpr

◆ M_BlockTileSize

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
index_t ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::M_BlockTileSize = MThreadClusterSize * MThreadSliceSize
staticconstexpr

◆ reorder_thread_cluster

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
bool ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::reorder_thread_cluster = (InSrcVectorDim == 0)
staticconstexpr

◆ thread_cluster_desc

template<typename InDataType, typename OutDataType, typename AccDataType, typename IndexDataType, typename InGridDesc_M_K, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename AccElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, bool PropagateNan, index_t BlockSize, index_t MThreadClusterSize, index_t KThreadClusterSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize>
auto ck::GridwiseReduction_mk_to_m_multiblock< InDataType, OutDataType, AccDataType, IndexDataType, InGridDesc_M_K, OutGridDesc_M, ReduceOperation, InElementwiseOperation, AccElementwiseOperation, OutMemoryDataOperation, PropagateNan, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize >::thread_cluster_desc
staticconstexpr
Initial value:
=
make_cluster_descriptor(ThreadClusterLengths_M_K{}, ThreadClusterArrangeOrder{})
ck::make_cluster_descriptor
__host__ __device__ constexpr auto make_cluster_descriptor(const Lengths &lengths, ArrangeOrder order=typename arithmetic_sequence_gen< 0, Lengths::Size(), 1 >::type{})
Definition tensor_description/cluster_descriptor.hpp:13
ck::GridwiseReduction_mk_to_m_multiblock::ThreadClusterArrangeOrder
typename conditional< reorder_thread_cluster, Sequence< 1, 0 >, Sequence< 0, 1 > >::type ThreadClusterArrangeOrder
Definition gridwise_2d_reduction_multiblock.hpp:105
ck::GridwiseReduction_mk_to_m_multiblock::ThreadClusterLengths_M_K
Sequence< MThreadClusterSize, KThreadClusterSize > ThreadClusterLengths_M_K
Definition gridwise_2d_reduction_multiblock.hpp:100
The documentation for this struct was generated from the following file: