/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_softmax_impl.hpp Source File#
device_softmax_impl.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: 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
__host__ constexpr __device__ auto make_merge_transform(const LowLengths &low_lengths)
Definition: multi_index_transform_helper.hpp:55
__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
__global__ void kernel_softmax(const GridDesc_M_K in_grid_desc_m_k, const GridDesc_M_K out_grid_desc_m_k, 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)
Definition: gridwise_softmax.hpp:22
__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_softmax.hpp:55
Definition: sequence.hpp:43
typename conditional< kHasContent, type0, type1 >::type type
Definition: sequence.hpp:256
Definition: integral_constant.hpp:10
Definition: device_base.hpp:50
Definition: device_base.hpp:61
Definition: device_softmax.hpp:24
Definition: device_softmax_impl.hpp:148
std::vector< index_t > inLengths_
Definition: device_softmax_impl.hpp:211
AccDataType alpha_
Definition: device_softmax_impl.hpp:214
index_t invariant_lowest_length_
Definition: device_softmax_impl.hpp:223
AccElementwiseOp acc_elementwise_op_
Definition: device_softmax_impl.hpp:221
const InDataType * in_dev_
Definition: device_softmax_impl.hpp:217
AccDataType beta_
Definition: device_softmax_impl.hpp:215
size_t gridSize
Definition: device_softmax_impl.hpp:227
Argument(const std::vector< index_t > inLengths, const std::vector< index_t > inStrides, const std::vector< index_t > reduceDims, double alpha, double beta, const InDataType *in_dev, OutDataType *out_dev, InElementwiseOp in_elementwise_op, AccElementwiseOp acc_elementwise_op)
Definition: device_softmax_impl.hpp:149
int blkGroupSize
Definition: device_softmax_impl.hpp:225
InElementwiseOp in_elementwise_op_
Definition: device_softmax_impl.hpp:220
OutDataType * out_dev_
Definition: device_softmax_impl.hpp:218
int numBlockTileIteration
Definition: device_softmax_impl.hpp:226
std::vector< index_t > inStrides_
Definition: device_softmax_impl.hpp:212
Definition: device_softmax_impl.hpp:231
float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition: device_softmax_impl.hpp:232
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition: device_softmax_impl.hpp:272
Definition: device_softmax_impl.hpp:43
static auto MakeSrc2dDescriptor(const std::vector< index_t > &inLengths, const std::vector< index_t > &inStrides, int blkGroupSize, int numBlockTileIteration)
Definition: device_softmax_impl.hpp:53
static constexpr index_t NumInvariantDim
Definition: device_softmax_impl.hpp:44
static auto MakeInvoker()
Definition: device_softmax_impl.hpp:390
static bool IsSupportedArgument(const Argument &arg)
Definition: device_softmax_impl.hpp:279
static constexpr index_t NumSrcDim
Definition: device_softmax_impl.hpp:46
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition: device_softmax_impl.hpp:392
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition: device_softmax_impl.hpp:325
std::unique_ptr< BaseArgument > MakeArgumentPointer(const std::vector< index_t > inLengths, const std::vector< index_t > inStrides, const std::vector< int > reduceDims, double alpha, double beta, const void *in_dev, void *out_dev, InElementwiseOp in_elementwise_op, AccElementwiseOp acc_elementwise_op) override
Definition: device_softmax_impl.hpp:369
static constexpr index_t M_BlockTileSize
Definition: device_softmax_impl.hpp:50
decltype(MakeSrc2dDescriptor({1}, {1}, 1, 1)) GridDesc_M_K
Definition: device_softmax_impl.hpp:117
GridwiseSoftmax_mk_to_mk< InDataType, OutDataType, AccDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, false > GridwiseSoftmaxGeneric
Definition: device_softmax_impl.hpp:131
static auto MakeArgument(const std::vector< index_t > inLengths, const std::vector< index_t > inStrides, const std::vector< int > reduceDims, double alpha, double beta, const InDataType *in_dev, OutDataType *out_dev, InElementwiseOp in_elementwise_op, AccElementwiseOp acc_elementwise_op)
Definition: device_softmax_impl.hpp:330
static constexpr index_t NumDstDim
Definition: device_softmax_impl.hpp:47
static constexpr index_t K_BlockTileSize
Definition: device_softmax_impl.hpp:51
std::string GetTypeString() const override
Definition: device_softmax_impl.hpp:397
GridwiseSoftmax_mk_to_mk< InDataType, OutDataType, AccDataType, GridDesc_M_K, BlockSize, MThreadClusterSize, KThreadClusterSize, MThreadSliceSize, KThreadSliceSize, InSrcVectorDim, InSrcVectorSize, OutDstVectorSize, true > GridwiseSoftmaxSweepOnce
Definition: device_softmax_impl.hpp:145
static constexpr bool reduceAllDim
Definition: device_softmax_impl.hpp:48