device_reduce_threadwise_multi_d.hpp Source File#
device_reduce_threadwise_multi_d.hpp
Go to the documentation of this file.
float launch_and_time_kernel(const StreamConfig &stream_config, F kernel, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, Args... args)
Definition host_utility/kernel_launch.hpp:14
__host__ __device__ constexpr auto integer_least_multiple(X x, Y y)
Definition utility/math.hpp:78
Definition convolution_backward_data_specialization.hpp:8
std::pair< long_index_t, long_index_t > get_2d_lengths(const std::vector< index_t > &inLengths)
Definition device_reduce_common.hpp:20
std::vector< index_t > shuffle_tensor_dimensions(const std::vector< index_t > &origLengthsStrides, const std::vector< int > &reduceDims)
Definition device_reduce_common.hpp:75
Definition convolution_backward_data_specialization.hpp:7
Definition ck.hpp:268
__global__ void kernel_reduce_threadwise_multi_d(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_value_global, OutDataType *const __restrict__ p_out_value_global)
Definition gridwise_2d_reduction_threadwise_multi_d.hpp:28
__host__ __device__ constexpr auto make_naive_tensor_descriptor(const Tuple< Lengths... > &lengths, const Tuple< Strides... > &strides)
Definition tensor_descriptor_helper.hpp:49
__host__ __device__ constexpr auto make_right_pad_transform(const LowLength &low_length, const RightPadLength &right_pad, integral_constant< bool, SkipIsValidCheck >=integral_constant< bool, false >{})
Definition multi_index_transform_helper.hpp:37
typename tuple_element< I, TTuple >::type tuple_element_t
Definition utility/tuple.hpp:208
__host__ __device__ constexpr auto make_merge_transform(const LowLengths &low_lengths)
Definition multi_index_transform_helper.hpp:55
__host__ __device__ constexpr auto generate_tuple(F &&f, Number< N >)
Definition tuple_helper.hpp:21
__host__ __device__ constexpr auto make_tuple(Xs &&... xs)
Definition utility/tuple.hpp:211
__host__ __device__ constexpr auto transform_tensor_descriptor(const OldTensorDescriptor &old_tensor_desc, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)
Definition tensor_description/tensor_descriptor.hpp:319
__host__ __device__ constexpr auto make_unmerge_transform(const UpLengths &up_lengths, integral_constant< bool, Use24BitIntegerCalculation >=integral_constant< bool, false >{})
Definition multi_index_transform_helper.hpp:90
Definition ck/stream_config.hpp:10
Definition gridwise_2d_reduction_threadwise_multi_d.hpp:66
decltype(MakeDsGridPointer()) DsGridPointer
Definition gridwise_2d_reduction_threadwise_multi_d.hpp:98
Definition utility/sequence.hpp:43
typename conditional< kHasContent, type0, type1 >::type type
Definition utility/sequence.hpp:271
Definition functional2.hpp:33
Definition device_base.hpp:197
BaseArgument()=default
BaseInvoker()=default
Definition device_reduce_multi_d.hpp:26
Definition device_reduce_threadwise_multi_d.hpp:198
long_index_t invariant_total_length
Definition device_reduce_threadwise_multi_d.hpp:267
OutElementwiseOperation out_elementwise_op_
Definition device_reduce_threadwise_multi_d.hpp:261
size_t gridSize
Definition device_reduce_threadwise_multi_d.hpp:271
std::array< index_t, NumDstDim > outLengths_
Definition device_reduce_threadwise_multi_d.hpp:252
const InDataType * in_dev_
Definition device_reduce_threadwise_multi_d.hpp:255
std::array< std::array< index_t, NumDstDim >, NumDTensor > DsStrides_
Definition device_reduce_threadwise_multi_d.hpp:250
DsGridPointer p_ds_grid_
Definition device_reduce_threadwise_multi_d.hpp:258
index_t invariant_lowest_length
Definition device_reduce_threadwise_multi_d.hpp:265
std::array< index_t, Rank > inStrides_
Definition device_reduce_threadwise_multi_d.hpp:247
index_t reduce_lowest_length
Definition device_reduce_threadwise_multi_d.hpp:266
int numBlockTileIteration
Definition device_reduce_threadwise_multi_d.hpp:270
std::array< index_t, Rank > inLengths_
Definition device_reduce_threadwise_multi_d.hpp:246
InElementwiseOperation in_elementwise_op_
Definition device_reduce_threadwise_multi_d.hpp:260
OutDataType * out_dev_
Definition device_reduce_threadwise_multi_d.hpp:256
std::array< std::array< index_t, NumDstDim >, NumDTensor > DsLengths_
Definition device_reduce_threadwise_multi_d.hpp:249
long_index_t reduce_total_length
Definition device_reduce_threadwise_multi_d.hpp:268
std::array< index_t, NumDstDim > outStrides_
Definition device_reduce_threadwise_multi_d.hpp:253
DsGridDesc_M ds_grid_desc_m_
Definition device_reduce_threadwise_multi_d.hpp:263
Argument(const std::array< index_t, Rank > inLengths, const std::array< index_t, Rank > inStrides, const std::array< std::array< index_t, NumDstDim >, NumDTensor > DsLengths, const std::array< std::array< index_t, NumDstDim >, NumDTensor > DsStrides, const std::array< index_t, NumDstDim > outLengths, const std::array< index_t, NumDstDim > outStrides, const std::array< int, NumReduceDim > reduceDims, const InDataType *in_dev, const std::array< const void *, NumDTensor > ds_dev, OutDataType *out_dev, const InElementwiseOperation in_elementwise_op, const OutElementwiseOperation out_elementwise_op)
Definition device_reduce_threadwise_multi_d.hpp:199
Definition device_reduce_threadwise_multi_d.hpp:275
float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition device_reduce_threadwise_multi_d.hpp:276
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition device_reduce_threadwise_multi_d.hpp:313
Definition device_reduce_threadwise_multi_d.hpp:47
static auto MakeSrc2dDescriptor(const std::array< index_t, Rank > &inLengths, const std::array< index_t, Rank > &inStrides)
Definition device_reduce_threadwise_multi_d.hpp:68
std::string GetTypeString() const override
Definition device_reduce_threadwise_multi_d.hpp:395
int32_t IndexDataType
Definition device_reduce_threadwise_multi_d.hpp:55
static constexpr index_t K_BlockTileSize
Definition device_reduce_threadwise_multi_d.hpp:66
decltype(MakeDsDescriptor({}, {})) DsGridDesc_M
Definition device_reduce_threadwise_multi_d.hpp:173
static constexpr index_t NumInvariantDim
Definition device_reduce_threadwise_multi_d.hpp:57
std::unique_ptr< BaseArgument > MakeArgumentPointer(const std::array< index_t, Rank > inLengths, const std::array< index_t, Rank > inStrides, const std::array< std::array< index_t, NumDstDim >, NumDTensor > DsLengths, const std::array< std::array< index_t, NumDstDim >, NumDTensor > DsStrides, const std::array< index_t, NumDstDim > outLengths, const std::array< index_t, NumDstDim > outStrides, const std::array< int, NumReduceDim > reduceDims, const void *in_dev, const std::array< const void *, NumDTensor > ds_dev, void *out_dev, const InElementwiseOperation in_elementwise_op, const OutElementwiseOperation out_elementwise_op) override
Definition device_reduce_threadwise_multi_d.hpp:363
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition device_reduce_threadwise_multi_d.hpp:320
static constexpr index_t NumSrcDim
Definition device_reduce_threadwise_multi_d.hpp:61
static constexpr index_t NumDstDim
Definition device_reduce_threadwise_multi_d.hpp:62
static auto MakeDst1dDescriptor(const std::array< index_t, NumDstDim > &outLengths, const std::array< index_t, NumDstDim > &outStrides)
Definition device_reduce_threadwise_multi_d.hpp:130
GridwiseReduction_mk_to_m_threadwise_multi_d< ReduceDataType, DsDataType, CDataType, GemmAccDataType, InGridDesc_M_K, DsGridDesc_M, OutGridDesc_M, ReduceAdd, PassThrough, OutElementwiseOperation, InMemoryDataOperationEnum::Set, BlockSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, decltype(DsVectorLengthSequence) > GridwiseReduce
Definition device_reduce_threadwise_multi_d.hpp:175
decltype(MakeDst1dDescriptor({}, {})) OutGridDesc_M
Definition device_reduce_threadwise_multi_d.hpp:172
decltype(MakeSrc2dDescriptor({}, {})) InGridDesc_M_K
Definition device_reduce_threadwise_multi_d.hpp:171
static auto MakeDsDescriptor(const std::array< std::array< index_t, NumDstDim >, NumDTensor > DsLengths, std::array< std::array< index_t, NumDstDim >, NumDTensor > DsStrides)
Definition device_reduce_threadwise_multi_d.hpp:160
static constexpr bool reduceAllDim
Definition device_reduce_threadwise_multi_d.hpp:63
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition device_reduce_threadwise_multi_d.hpp:390
static constexpr index_t NumDTensor
Definition device_reduce_threadwise_multi_d.hpp:59
static constexpr index_t M_BlockTileSize
Definition device_reduce_threadwise_multi_d.hpp:65
typename GridwiseReduce::DsGridPointer DsGridPointer
Definition device_reduce_threadwise_multi_d.hpp:195