device_elementwise_dynamic_vector_dims_impl.hpp Source File#
device_elementwise_dynamic_vector_dims_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 host_utility/kernel_launch.hpp:14
__host__ __device__ constexpr auto integer_divide_ceil(X x, Y y)
Definition utility/math.hpp:72
Definition convolution_backward_data_specialization.hpp:8
Definition convolution_backward_data_specialization.hpp:7
Definition ck.hpp:268
__host__ __device__ constexpr auto concat_tuple(const Tuple< X... > &tx, const Tuple< Y... > &ty)
Definition tuple_helper.hpp:52
__host__ __device__ constexpr auto make_pass_through_transform(const LowLength &low_length)
Definition multi_index_transform_helper.hpp:12
__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
__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
__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
__global__ void kernel_elementwise(const InGridDescTuple in_grid_desc_tuple, const OutGridDescTuple out_grid_desc_tuple, const InDataTypePointerTuple p_in_global_tuple, const OutDataTypePointerTuple p_out_global_tuple, const Block2TileMap block_2_tile_map, const ElementwiseOperation elementwise_op)
Definition gridwise_elementwise_2d.hpp:29
Definition ck/stream_config.hpp:10
Definition block_to_ctile_map.hpp:261
Definition utility/sequence.hpp:43
Definition utility/sequence.hpp:256
Definition functional2.hpp:33
Definition device_base.hpp:197
BaseArgument()=default
BaseInvoker()=default
Definition device_elementwise.hpp:21
Definition device_elementwise_dynamic_vector_dims_impl.hpp:214
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)
Definition device_elementwise_dynamic_vector_dims_impl.hpp:215
InDataTypePointerTuple in_dev_buffers_
Definition device_elementwise_dynamic_vector_dims_impl.hpp:242
std::array< index_t, NumDim > lengths_
Definition device_elementwise_dynamic_vector_dims_impl.hpp:245
OutDataTypePointerTuple out_dev_buffers_
Definition device_elementwise_dynamic_vector_dims_impl.hpp:243
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_dynamic_vector_dims_impl.hpp:254
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition device_elementwise_dynamic_vector_dims_impl.hpp:316
Definition device_elementwise_dynamic_vector_dims_impl.hpp:37
static auto MakeInvoker()
Definition device_elementwise_dynamic_vector_dims_impl.hpp:398
static auto PadInputOutputDescriptor(const InOutDescriptor &desc)
Definition device_elementwise_dynamic_vector_dims_impl.hpp:89
decltype(GenerateInOutGridDescTuple< NumOutput >()) OutGridDescTuple
Definition device_elementwise_dynamic_vector_dims_impl.hpp:175
GridwiseElementwise< InGridDescTuple, OutGridDescTuple, InDataTypePointerTuple, OutDataTypePointerTuple, Block2TileMap, ElementwiseOperation, BlockSize, M0PerBlock, M1PerBlock, M0PerThread, M1PerThread, ThreadClusterArrangeOrder, InScalarPerVectorSeq, OutScalarPerVectorSeq, I1, I0 > GridwiseElementwiseOp
Definition device_elementwise_dynamic_vector_dims_impl.hpp:179
static auto GenerateBatchDimsLenghtsTuple(const std::array< index_t, NumDim > &lengths, const index_t M0_dim, const index_t M1_dim)
Definition device_elementwise_dynamic_vector_dims_impl.hpp:105
static constexpr auto I1
Definition device_elementwise_dynamic_vector_dims_impl.hpp:42
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
decltype(GenerateOutDataTypePointerTuple()) OutDataTypePointerTuple
Definition device_elementwise_dynamic_vector_dims_impl.hpp:71
BlockToCTileMap_M00_N0_M01Adapt< M0PerBlock, M1PerBlock > Block2TileMap
Definition device_elementwise_dynamic_vector_dims_impl.hpp:177
static constexpr int NumInput
Definition device_elementwise_dynamic_vector_dims_impl.hpp:38
static bool IsSupportedArgument(const Argument &arg)
Definition device_elementwise_dynamic_vector_dims_impl.hpp:323
static auto GenerateInOutGridDescTuple()
Definition device_elementwise_dynamic_vector_dims_impl.hpp:162
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition device_elementwise_dynamic_vector_dims_impl.hpp:399
GridwiseElementwise< InGridDescTuple, OutGridDescTuple, InDataTypePointerTuple, OutDataTypePointerTuple, Block2TileMap, ElementwiseOperation, BlockSize, M0PerBlock, M1PerBlock, M0PerThread, M1PerThread, ThreadClusterArrangeOrder, InScalarPerVectorSeq, OutScalarPerVectorSeq, I1, I1 > GridwiseElementwiseOpSameInOutVectorDim
Definition device_elementwise_dynamic_vector_dims_impl.hpp:196
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_dynamic_vector_dims_impl.hpp:404
static auto MakeDescriptor(const std::array< index_t, NumDim > &lengths, const std::array< index_t, NumDim > &in_strides, const std::array< index_t, NumDim > &out_strides, const std::array< index_t, NumDim > &desc_strides)
Definition device_elementwise_dynamic_vector_dims_impl.hpp:128
static index_t GetLowestStrideDim(const std::array< index_t, NumDim > &strides)
Definition device_elementwise_dynamic_vector_dims_impl.hpp:73
static constexpr int NumOutput
Definition device_elementwise_dynamic_vector_dims_impl.hpp:39
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) override
Definition device_elementwise_dynamic_vector_dims_impl.hpp:383
static auto GenerateInDataTypePointerTuple()
Definition device_elementwise_dynamic_vector_dims_impl.hpp:48
static auto GenerateOutDataTypePointerTuple()
Definition device_elementwise_dynamic_vector_dims_impl.hpp:59
decltype(GenerateInOutGridDescTuple< NumInput >()) InGridDescTuple
Definition device_elementwise_dynamic_vector_dims_impl.hpp:174
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)
Definition device_elementwise_dynamic_vector_dims_impl.hpp:367
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition device_elementwise_dynamic_vector_dims_impl.hpp:361