device_multiple_reduce_multiblock.hpp Source File#
device_multiple_reduce_multiblock.hpp
Go to the documentation of this file.
225 static auto MakeDst1dDescriptorForBufferSet(const std::array<index_t, NumOutputDim>& outLengths,
580 str << (OutMemoryDataOperation == InMemoryDataOperationEnum::Set? "DeviceMultipleReduceBlockWise<" : "DeviceMultipleReduceMultiBlock<") << BlockSize << ",";
585 static_for<0, OutDstVectorSizeSeq::Size(), 1>{}([&](auto I) {str << "_" << OutDstVectorSizeSeq::At(I); });
auto pad(ck::index_t mpb, ck::index_t npb, ck::index_t kpb, ck::tensor_operation::device::GemmSpecialization gemm, CDesc_MRaw_NRaw conv)
Definition helper.hpp:70
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
constexpr T GetIdentityValueForInMemoryDataOperation(InMemoryDataOperationEnum operation)
Definition reduction_operator.hpp:473
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
__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
__global__ void kernel_multiple_buffer_set_value(const Grid1dBufferDescTuple grid_1d_buffer_desc_tuple, DataTypePointerTuple p_global_tuple, DataTypeTuple value_tuple)
Definition gridwise_set_multiple_buffer_value.hpp:17
__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
__global__ void kernel_multiple_reduce_multiblock(const InGridDesc_M_K in_grid_desc_m_k, const OutGridDesc_M_Tuple out_grid_desc_m_tuple, const InElementwiseOperationTuple in_elementwise_op_tuple, const AccElementwiseOperationTuple acc_elementwise_op_tuple, index_t block_group_size, index_t num_k_block_tile_iteration, Array< AccDataType, NumReduction > alpha_values, const InDataType *const __restrict__ p_in_value_global, Array< AccDataType, NumReduction > beta_values, OutDataTypePointerTuple p_out_value_global_tuple)
Definition gridwise_2d_multiple_reduction_multiblock.hpp:26
__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
__host__ __device__ constexpr bool sequence_all_of(Seq, F f)
Definition utility/sequence.hpp:912
Definition ck/stream_config.hpp:10
Definition utility/array.hpp:14
Definition gridwise_2d_multiple_reduction_multiblock.hpp:69
Definition utility/sequence.hpp:43
typename conditional< kHasContent, type0, type1 >::type type
Definition utility/sequence.hpp:271
Definition reduction_operator.hpp:485
Definition functional2.hpp:33
Definition device_base.hpp:197
BaseArgument()=default
BaseInvoker()=default
Definition device_multiple_reduce.hpp:25
Definition device_multiple_reduce_multiblock.hpp:267
OutGridDesc_M_Tuple_2 out_grid_desc_m_tuple_2
Definition device_multiple_reduce_multiblock.hpp:370
std::array< index_t, NumInputDim > inLengths_
Definition device_multiple_reduce_multiblock.hpp:356
InGridDesc_M_K in_grid_desc_m_k
Definition device_multiple_reduce_multiblock.hpp:368
long_index_t invariant_total_length
Definition device_multiple_reduce_multiblock.hpp:375
Array< AccDataType, NumReduction > beta_values_
Definition device_multiple_reduce_multiblock.hpp:363
long_index_t reduce_total_length
Definition device_multiple_reduce_multiblock.hpp:376
size_t gridSize_pre
Definition device_multiple_reduce_multiblock.hpp:382
int blkGroupSize
Definition device_multiple_reduce_multiblock.hpp:378
int numBlockTileIteration
Definition device_multiple_reduce_multiblock.hpp:379
const InDataType * in_dev_
Definition device_multiple_reduce_multiblock.hpp:365
size_t gridSize
Definition device_multiple_reduce_multiblock.hpp:380
OutGridDesc_M_Tuple out_grid_desc_m_tuple
Definition device_multiple_reduce_multiblock.hpp:369
OutDataTypePointerTuple out_dev_buffers_
Definition device_multiple_reduce_multiblock.hpp:366
Argument(const std::array< index_t, NumInputDim > &inLengths, const std::array< index_t, NumInputDim > &inStrides, const std::array< index_t, NumOutputDim > &outLengths, const std::array< std::array< index_t, NumOutputDim >, NumReduction > &outStridesArray, const std::array< int, NumReduceDim > &reduceDims, const std::array< double, NumReduction > &alphas, const std::array< double, NumReduction > &betas, const void *in_dev, const std::array< void *, NumReduction > &out_dev_buffers, const InElementwiseOperationTuple in_elementwise_op_tuple, const AccElementwiseOperationTuple acc_elementwise_op_tuple)
Definition device_multiple_reduce_multiblock.hpp:268
std::array< std::array< index_t, NumOutputDim >, NumReduction > outStridesArray_
Definition device_multiple_reduce_multiblock.hpp:360
std::array< index_t, NumOutputDim > outLengths_
Definition device_multiple_reduce_multiblock.hpp:359
std::array< index_t, NumInputDim > inStrides_
Definition device_multiple_reduce_multiblock.hpp:357
InElementwiseOperationTuple in_elementwise_op_tuple_
Definition device_multiple_reduce_multiblock.hpp:372
Array< AccDataType, NumReduction > alpha_values_
Definition device_multiple_reduce_multiblock.hpp:362
AccElementwiseOperationTuple acc_elementwise_op_tuple_
Definition device_multiple_reduce_multiblock.hpp:373
Definition device_multiple_reduce_multiblock.hpp:386
float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition device_multiple_reduce_multiblock.hpp:387
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition device_multiple_reduce_multiblock.hpp:468
Definition device_multiple_reduce_multiblock.hpp:48
static constexpr index_t NumInvariantDim
Definition device_multiple_reduce_multiblock.hpp:86
decltype(GenerateOutGrid1dDescTuple_2()) OutGridDesc_M_Tuple_2
Definition device_multiple_reduce_multiblock.hpp:264
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition device_multiple_reduce_multiblock.hpp:475
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition device_multiple_reduce_multiblock.hpp:570
static constexpr bool use_multiblock
Definition device_multiple_reduce_multiblock.hpp:94
static auto GenerateOutGrid1dDescTuple()
Definition device_multiple_reduce_multiblock.hpp:210
static auto GenerateOutGrid1dDescTuple_2()
Definition device_multiple_reduce_multiblock.hpp:253
static constexpr bool CheckDataTypeTuple()
Definition device_multiple_reduce_multiblock.hpp:69
static auto MakeDst1dDescriptor(const std::array< index_t, NumOutputDim > &outLengths, const std::array< index_t, NumOutputDim > &outStrides)
Definition device_multiple_reduce_multiblock.hpp:181
std::string GetTypeString() const override
Definition device_multiple_reduce_multiblock.hpp:575
std::unique_ptr< BaseArgument > MakeArgumentPointer(const std::array< index_t, NumInputDim > inLengths, const std::array< index_t, NumInputDim > inStrides, const std::array< index_t, NumOutputDim > outLengths, const std::array< std::array< index_t, NumOutputDim >, NumReduction > outStridesArray, const std::array< int, NumReduceDim > reduceDims, const std::array< double, NumReduction > alphas, const std::array< double, NumReduction > betas, const void *in_dev, const std::array< void *, NumReduction > out_dev_buffers, const InElementwiseOperationTuple in_elementwise_op_tuple, const AccElementwiseOperationTuple acc_elementwise_op_tuple) override
Definition device_multiple_reduce_multiblock.hpp:544
decltype(GenerateOutGrid1dDescTuple()) OutGridDesc_M_Tuple
Definition device_multiple_reduce_multiblock.hpp:223
static constexpr index_t K_BlockTileSize
Definition device_multiple_reduce_multiblock.hpp:102
static constexpr bool reduceAllDim
Definition device_multiple_reduce_multiblock.hpp:90
static auto MakeSrc2dDescriptor(const std::array< index_t, NumInputDim > &inLengths, const std::array< index_t, NumInputDim > &inStrides, int blkGroupSize, int numBlockTileIteration)
Definition device_multiple_reduce_multiblock.hpp:117
static auto GenerateOutDataTypePointerTuple()
Definition device_multiple_reduce_multiblock.hpp:104
static constexpr index_t NumInputDim
Definition device_multiple_reduce_multiblock.hpp:88
static constexpr index_t M_BlockTileSize
Definition device_multiple_reduce_multiblock.hpp:101
decltype(GenerateOutDataTypePointerTuple()) OutDataTypePointerTuple
Definition device_multiple_reduce_multiblock.hpp:115
static constexpr index_t NumOutputDim
Definition device_multiple_reduce_multiblock.hpp:89
static auto MakeDst1dDescriptorForBufferSet(const std::array< index_t, NumOutputDim > &outLengths, const std::array< index_t, NumOutputDim > &outStrides)
Definition device_multiple_reduce_multiblock.hpp:225
decltype(MakeSrc2dDescriptor( std::array< index_t, NumInputDim >{}, std::array< index_t, NumInputDim >{}, 1, 1)) InGridDesc_M_K
Definition device_multiple_reduce_multiblock.hpp:221