/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_elementwise_scale_impl.hpp Source File#
device_elementwise_scale_impl.hpp
Go to the documentation of this file.
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: 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
__global__ void kernel_elementwise_1d(const InGrid1dDescTuple in_grid_1d_desc_tuple, const OutGrid1dDescTuple out_grid_1d_desc_tuple, const InDataTypePointerTuple p_in_global_tuple, const OutDataTypePointerTuple p_out_global_tuple, const ElementwiseOperation elementwise_op, const UnaryOperation unary_op, const Scale scale_op)
Definition: gridwise_elementwise_1d_scale.hpp:21
__host__ constexpr __device__ auto generate_sequence_v2(F &&f, Number< N >)
Definition: sequence_helper.hpp:25
__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_elementwise_1d_scale.hpp:49
Definition: sequence.hpp:43
Definition: integral_constant.hpp:10
Definition: functional2.hpp:31
Definition: device_base.hpp:50
Definition: device_base.hpp:61
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:214
Scale scale_op_
Definition: device_elementwise_scale_impl.hpp:190
InDataTypePointerTuple in_dev_buffers_
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:242
UnaryOperation unary_op_
Definition: device_elementwise_scale_impl.hpp:189
std::array< index_t, NumDim > lengths_
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:245
Argument(const std::array< index_t, NumDim > lengths, const std::array< std::array< index_t, NumDim >, NumInput > inStridesArray, const std::array< std::array< index_t, NumDim >, NumOutput > outStridesArray, const std::array< const void *, NumInput > in_dev_buffers, const std::array< void *, NumOutput > out_dev_buffers, ElementwiseOperation elementwise_op, UnaryOperation unary_op, Scale scale_op)
Definition: device_elementwise_scale_impl.hpp:149
OutDataTypePointerTuple out_dev_buffers_
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:243
index_t blockSize_
Definition: device_elementwise_scale_impl.hpp:191
ElementwiseOperation elementwise_op_
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:249
std::array< std::array< index_t, NumDim >, NumInput > inStridesArray_
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:246
std::array< std::array< index_t, NumDim >, NumOutput > outStridesArray_
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:247
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:253
float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition: device_elementwise_scale_impl.hpp:196
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition: device_elementwise_scale_impl.hpp:239
static auto MakeInvoker()
Definition: device_elementwise_scale_impl.hpp:324
decltype(GenerateInOutGrid1dDescTuple(Number< NumInput >{})) InGrid1dDescTuple
Definition: device_elementwise_scale_impl.hpp:133
static auto MakeArgument(const std::array< index_t, NumDim > lengths, const std::array< std::array< index_t, NumDim >, NumInput > inStridesArray, const std::array< std::array< index_t, NumDim >, NumOutput > outStridesArray, const std::array< const void *, NumInput > in_dev_buffers, const std::array< void *, NumOutput > out_dev_buffers, ElementwiseOperation elementwise_op, UnaryOperation unary_op, Scale scale_op)
Definition: device_elementwise_scale_impl.hpp:285
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition: device_elementwise_scale_impl.hpp:325
static auto PadDescriptor_M_1d(Desc_M desc_m, index_t gridSize, index_t blockSize)
Definition: device_elementwise_scale_impl.hpp:75
static constexpr auto I0
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:41
decltype(GenerateInDataTypePointerTuple()) InDataTypePointerTuple
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:70
GridwiseElementwise_1D< InGrid1dDescTuple, OutGrid1dDescTuple, InDataTypePointerTuple, OutDataTypePointerTuple, ElementwiseOperation, UnaryOperation, Scale, MPerThread, InScalarPerVectorSeq, OutScalarPerVectorSeq > GridwiseElementwise
Definition: device_elementwise_scale_impl.hpp:145
static auto MakeDescriptor_M(const std::array< index_t, NumDim > &lengths, const std::array< index_t, NumDim > &stride, index_t gridSize, index_t blockSize)
Definition: device_elementwise_scale_impl.hpp:90
decltype(GenerateInOutGrid1dDescTuple(Number< NumOutput >{})) OutGrid1dDescTuple
Definition: device_elementwise_scale_impl.hpp:134
decltype(GenerateOutDataTypePointerTuple()) OutDataTypePointerTuple
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:71
static auto GenerateInOutGrid1dDescTuple(Number< TupleSize >)
Definition: device_elementwise_scale_impl.hpp:117
static constexpr int NumInput
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:38
static bool IsSupportedArgument(const Argument &arg)
Definition: device_elementwise_scale_impl.hpp:246
std::string GetTypeString() const override
Definition: device_elementwise_scale_impl.hpp:330
static constexpr int NumOutput
Definition: device_elementwise_dynamic_vector_dims_impl.hpp:39
static auto GenerateInDataTypePointerTuple()
Definition: device_elementwise_scale_impl.hpp:49
static auto GenerateOutDataTypePointerTuple()
Definition: device_elementwise_scale_impl.hpp:60
std::unique_ptr< BaseArgument > MakeArgumentPointer(const std::array< index_t, NumDim > lengths, const std::array< std::array< index_t, NumDim >, NumInput > inStridesArray, const std::array< std::array< index_t, NumDim >, NumOutput > outStridesArray, const std::array< const void *, NumInput > in_dev_buffers, const std::array< void *, NumOutput > out_dev_buffers, ElementwiseOperation elementwise_op, UnaryOperation unary_op, Scale scale_op) override
Definition: device_elementwise_scale_impl.hpp:305
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition: device_elementwise_scale_impl.hpp:279