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 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
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_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__ __device__ constexpr auto make_merge_transform(const LowLengths &low_lengths)
Definition multi_index_transform_helper.hpp:55
__host__ __device__ constexpr auto generate_sequence_v2(F &&f, Number< N >)
Definition sequence_helper.hpp:25
__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
Definition ck/stream_config.hpp:10
Definition gridwise_elementwise_1d_scale.hpp:49
Definition utility/sequence.hpp:43
Definition functional2.hpp:33
Definition device_base.hpp:197
Definition device_base.hpp:208
Definition device_elementwise.hpp:21
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
Definition device_elementwise_dynamic_vector_dims_impl.hpp:37
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
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
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
decltype(GenerateInDataTypePointerTuple()) InDataTypePointerTuple
Definition device_elementwise_dynamic_vector_dims_impl.hpp:70
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::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition device_elementwise_scale_impl.hpp:325
GridwiseElementwise_1D< InGrid1dDescTuple, OutGrid1dDescTuple, InDataTypePointerTuple, OutDataTypePointerTuple, ElementwiseOperation, UnaryOperation, Scale, MPerThread, InScalarPerVectorSeq, OutScalarPerVectorSeq > GridwiseElementwise
Definition device_elementwise_scale_impl.hpp:136
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
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition device_elementwise_scale_impl.hpp:279