GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize > Struct Template Reference

GridwiseReduction_mk_to_m_threadwise_multi_d&lt; InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize &gt; Struct Template Reference#

Composable Kernel: ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize > Struct Template Reference

#include <gridwise_2d_reduction_threadwise_multi_d.hpp>

Public Types

using ThreadBufferDimAccessOrder
using ThreadReduceSrcDesc_M_K
using ThreadReduceDstDesc_M
using PassThrough = tensor_operation::element_wise::PassThrough
using DsGridPointer = decltype(MakeDsGridPointer())

Static Public Member Functions

static constexpr auto MakeDsGridPointer ()
static __device__ void Run (const InGridDesc_M_K &in_grid_desc_m_k, const DsGridDesc_M &ds_grid_desc_m, const OutGridDesc_M &out_grid_desc_m, const InElementwiseOperation &in_elementwise_op, const OutElementwiseOperation &out_elementwise_op, const InDataType *const __restrict__ p_in_value_global, const DsGridPointer p_ds_grid, OutDataType *const __restrict__ p_out_value_global)

Static Public Attributes

static constexpr auto I0 = Number<0>{}
static constexpr index_t NumDTensor = DsDataType::Size()

Member Typedef Documentation

◆ DsGridPointer

template<typename InDataType, typename DsDataType, typename OutDataType, typename AccDataType, typename InGridDesc_M_K, typename DsGridDesc_M, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename OutElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, index_t BlockSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize, typename DsVectorSize>
using ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize >::DsGridPointer = decltype(MakeDsGridPointer())

◆ PassThrough

template<typename InDataType, typename DsDataType, typename OutDataType, typename AccDataType, typename InGridDesc_M_K, typename DsGridDesc_M, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename OutElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, index_t BlockSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize, typename DsVectorSize>
using ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize >::PassThrough = tensor_operation::element_wise::PassThrough

◆ ThreadBufferDimAccessOrder

template<typename InDataType, typename DsDataType, typename OutDataType, typename AccDataType, typename InGridDesc_M_K, typename DsGridDesc_M, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename OutElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, index_t BlockSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize, typename DsVectorSize>
using ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize >::ThreadBufferDimAccessOrder
Initial value:
typename conditional<InSrcVectorDim == 0, Sequence<1, 0>, Sequence<0, 1>>::type
ck::Sequence
Definition utility/sequence.hpp:43
ck::conditional
Definition utility/functional.hpp:100

◆ ThreadReduceDstDesc_M

template<typename InDataType, typename DsDataType, typename OutDataType, typename AccDataType, typename InGridDesc_M_K, typename DsGridDesc_M, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename OutElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, index_t BlockSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize, typename DsVectorSize>
using ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize >::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 DsDataType, typename OutDataType, typename AccDataType, typename InGridDesc_M_K, typename DsGridDesc_M, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename OutElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, index_t BlockSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize, typename DsVectorSize>
using ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize >::ThreadReduceSrcDesc_M_K
Initial value:
decltype(make_naive_tensor_descriptor_packed(
make_tuple(Number<MThreadSliceSize>{}, Number<KThreadSliceSize>{})))

Member Function Documentation

◆ MakeDsGridPointer()

template<typename InDataType, typename DsDataType, typename OutDataType, typename AccDataType, typename InGridDesc_M_K, typename DsGridDesc_M, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename OutElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, index_t BlockSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize, typename DsVectorSize>
constexpr auto ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize >::MakeDsGridPointer ( )
inlinestaticconstexpr

◆ Run()

template<typename InDataType, typename DsDataType, typename OutDataType, typename AccDataType, typename InGridDesc_M_K, typename DsGridDesc_M, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename OutElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, index_t BlockSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize, typename DsVectorSize>
__device__ void ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize >::Run ( const InGridDesc_M_K & in_grid_desc_m_k,
const DsGridDesc_M & ds_grid_desc_m,
const OutGridDesc_M & out_grid_desc_m,
const InElementwiseOperation & in_elementwise_op,
const OutElementwiseOperation & out_elementwise_op,
const InDataType *const __restrict__ p_in_value_global,
const DsGridPointer p_ds_grid,
OutDataType *const __restrict__ p_out_value_global )
inlinestatic

Member Data Documentation

◆ I0

template<typename InDataType, typename DsDataType, typename OutDataType, typename AccDataType, typename InGridDesc_M_K, typename DsGridDesc_M, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename OutElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, index_t BlockSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize, typename DsVectorSize>
auto ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize >::I0 = Number<0>{}
staticconstexpr

◆ NumDTensor

template<typename InDataType, typename DsDataType, typename OutDataType, typename AccDataType, typename InGridDesc_M_K, typename DsGridDesc_M, typename OutGridDesc_M, typename ReduceOperation, typename InElementwiseOperation, typename OutElementwiseOperation, InMemoryDataOperationEnum OutMemoryDataOperation, index_t BlockSize, index_t MThreadSliceSize, index_t KThreadSliceSize, index_t InSrcVectorDim, index_t InSrcVectorSize, index_t OutDstVectorSize, typename DsVectorSize>
index_t ck::GridwiseReduction_mk_to_m_threadwise_multi_d< InDataType, DsDataType, OutDataType, AccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceOperation, InElementwiseOperation, OutElementwiseOperation, OutMemoryDataOperation, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, DsVectorSize >::NumDTensor = DsDataType::Size()
staticconstexpr
The documentation for this struct was generated from the following file: