/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/tensor_operation/gpu/device/impl/device_multiple_reduce_threadwise.hpp Source File#
device_multiple_reduce_threadwise.hpp
Go to the documentation of this file.
412 static_for<0, OutDstVectorSizeSeq::Size(), 1>{}([&](auto I) {str << "_" << OutDstVectorSizeSeq::At(I); });
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: kernel_launch.hpp:13
__host__ constexpr __device__ auto integer_least_multiple(X x, Y y)
Definition: math.hpp:78
Definition: ck.hpp:264
__host__ constexpr __device__ auto make_naive_tensor_descriptor(const Tuple< Lengths... > &lengths, const Tuple< Strides... > &strides)
Definition: tensor_descriptor_helper.hpp:49
__host__ constexpr __device__ auto generate_tuple(F &&f, Number< N >)
Definition: tuple_helper.hpp:15
__global__ void kernel_multiple_reduce_threadwise(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, 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_threadwise.hpp:26
__host__ constexpr __device__ auto make_merge_transform(const LowLengths &low_lengths)
Definition: multi_index_transform_helper.hpp:55
__host__ constexpr __device__ bool sequence_all_of(Seq, F f)
Definition: sequence.hpp:885
__host__ constexpr __device__ auto make_unmerge_transform(const UpLengths &up_lengths, integral_constant< bool, Use24BitIntegerCalculation >=integral_constant< bool, false >{})
Definition: multi_index_transform_helper.hpp:90
__host__ constexpr __device__ auto transform_tensor_descriptor(const OldTensorDescriptor &old_tensor_desc, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)
Definition: tensor_descriptor.hpp:319
__host__ constexpr __device__ 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
Definition: stream_config.hpp:10
Definition: gridwise_2d_multiple_reduction_threadwise.hpp:63
Definition: sequence.hpp:43
typename conditional< kHasContent, type0, type1 >::type type
Definition: sequence.hpp:256
Definition: integral_constant.hpp:10
Definition: functional2.hpp:31
Definition: device_base.hpp:50
Definition: device_base.hpp:61
Definition: device_multiple_reduce.hpp:25
Definition: device_multiple_reduce_threadwise.hpp:192
InGridDesc_M_K in_grid_desc_m_k
Definition: device_multiple_reduce_threadwise.hpp:254
long_index_t invariant_total_length
Definition: device_multiple_reduce_threadwise.hpp:260
long_index_t reduce_total_length
Definition: device_multiple_reduce_threadwise.hpp:261
Array< AccDataType, NumReduction > beta_values_
Definition: device_multiple_reduce_threadwise.hpp:249
Array< AccDataType, NumReduction > alpha_values_
Definition: device_multiple_reduce_threadwise.hpp:248
const InDataType * in_dev_
Definition: device_multiple_reduce_threadwise.hpp:251
std::array< index_t, NumInputDim > inLengths_
Definition: device_multiple_reduce_threadwise.hpp:242
std::array< index_t, NumInputDim > inStrides_
Definition: device_multiple_reduce_threadwise.hpp:243
size_t gridSize
Definition: device_multiple_reduce_threadwise.hpp:263
InElementwiseOperationTuple in_elementwise_op_tuple_
Definition: device_multiple_reduce_threadwise.hpp:257
std::array< index_t, NumOutputDim > outLengths_
Definition: device_multiple_reduce_threadwise.hpp:245
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_threadwise.hpp:193
AccElementwiseOperationTuple acc_elementwise_op_tuple_
Definition: device_multiple_reduce_threadwise.hpp:258
OutGridDesc_M_Tuple out_grid_desc_m_tuple
Definition: device_multiple_reduce_threadwise.hpp:255
std::array< std::array< index_t, NumOutputDim >, NumReduction > outStridesArray_
Definition: device_multiple_reduce_threadwise.hpp:246
OutDataTypePointerTuple out_dev_buffers_
Definition: device_multiple_reduce_threadwise.hpp:252
Definition: device_multiple_reduce_threadwise.hpp:267
float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition: device_multiple_reduce_threadwise.hpp:268
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition: device_multiple_reduce_threadwise.hpp:319
Definition: device_multiple_reduce_threadwise.hpp:44
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition: device_multiple_reduce_threadwise.hpp:326
static constexpr index_t NumInputDim
Definition: device_multiple_reduce_threadwise.hpp:65
static constexpr index_t NumInvariantDim
Definition: device_multiple_reduce_threadwise.hpp:63
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_threadwise.hpp:371
static auto MakeSrc2dDescriptor(const std::array< index_t, NumInputDim > &inLengths, const std::array< index_t, NumInputDim > &inStrides)
Definition: device_multiple_reduce_threadwise.hpp:85
static auto GenerateOutGrid1dDescTuple()
Definition: device_multiple_reduce_threadwise.hpp:176
static auto MakeDst1dDescriptor(const std::array< index_t, NumOutputDim > &outLengths, const std::array< index_t, NumOutputDim > &outStrides)
Definition: device_multiple_reduce_threadwise.hpp:147
std::string GetTypeString() const override
Definition: device_multiple_reduce_threadwise.hpp:402
decltype(GenerateOutDataTypePointerTuple()) OutDataTypePointerTuple
Definition: device_multiple_reduce_threadwise.hpp:83
static constexpr index_t K_BlockTileSize
Definition: device_multiple_reduce_threadwise.hpp:70
static auto GenerateOutDataTypePointerTuple()
Definition: device_multiple_reduce_threadwise.hpp:72
decltype(GenerateOutGrid1dDescTuple()) OutGridDesc_M_Tuple
Definition: device_multiple_reduce_threadwise.hpp:189
static constexpr bool reduceAllDim
Definition: device_multiple_reduce_threadwise.hpp:67
decltype(MakeSrc2dDescriptor(std::array< index_t, NumInputDim >{}, std::array< index_t, NumInputDim >{})) InGridDesc_M_K
Definition: device_multiple_reduce_threadwise.hpp:188
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition: device_multiple_reduce_threadwise.hpp:397
static constexpr index_t NumOutputDim
Definition: device_multiple_reduce_threadwise.hpp:66
static constexpr index_t M_BlockTileSize
Definition: device_multiple_reduce_threadwise.hpp:69