device_permute_impl.hpp Source File#
device_permute_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
__host__ __device__ constexpr auto PadTensorDescriptor(const TensorDesc &desc, const TileLengths &tile_lengths, DoPads)
Definition matrix_padder.hpp:19
Definition convolution_backward_data_specialization.hpp:7
Definition ck.hpp:268
__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
__global__ void kernel_nd_permute(const InGridDesc in_grid_desc, const OutGridDesc out_grid_desc, const InDataType *p_in_global, OutDataType *p_out_global, const ElementwiseOperation elementwise_op, const Block2TileMap block_2_tile_map)
Definition gridwise_permute.hpp:25
__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
STL namespace.
Definition ck/stream_config.hpp:10
__host__ static __device__ constexpr bool CheckValidity(const InGridDesc &in_grid_desc, const OutGridDesc &out_grid_desc)
Definition gridwise_permute.hpp:182
Block2TileMap DefaultBlock2TileMap
Definition gridwise_permute.hpp:134
Definition utility/sequence.hpp:43
Definition device_base.hpp:197
BaseArgument()=default
BaseInvoker()=default
Definition device_permute.hpp:18
std::array< index_t, NumDim > Lengths
Definition device_permute.hpp:19
Definition device_permute_impl.hpp:113
Strides in_strides_
Definition device_permute_impl.hpp:140
ElementwiseOperation elementwise_op_
Definition device_permute_impl.hpp:144
Argument(const Lengths &in_lengths, const Strides &in_strides, const Lengths &out_lengths, const Strides &out_strides, const void *in_dev_buffer, void *out_dev_buffer, ElementwiseOperation elementwise_op)
Definition device_permute_impl.hpp:114
OutGridDesc out_grid_desc_
Definition device_permute_impl.hpp:137
Block2TileMap block_2_tile_map_
Definition device_permute_impl.hpp:146
InGridDesc in_grid_desc_
Definition device_permute_impl.hpp:136
const InDataType * in_dev_buffer_
Definition device_permute_impl.hpp:134
Lengths in_lengths_
Definition device_permute_impl.hpp:139
Strides out_strides_
Definition device_permute_impl.hpp:142
OutDataType * out_dev_buffer_
Definition device_permute_impl.hpp:135
Lengths out_lengths_
Definition device_permute_impl.hpp:141
Definition device_permute_impl.hpp:150
static float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition device_permute_impl.hpp:151
float Run(const BaseArgument *arg, const StreamConfig &stream_config=StreamConfig{}) override final
Definition device_permute_impl.hpp:177
Definition device_permute_impl.hpp:45
InGridDesc OutGridDesc
Definition device_permute_impl.hpp:90
decltype(MakeDescriptor_N_H_W({1, 1}, {1, 1})) InGridDesc
Definition device_permute_impl.hpp:89
static std::enable_if_t< std::is_constructible_v< Argument, Args... >, Argument > MakeArgument(Args &&... args) noexcept(std::is_nothrow_constructible_v< Argument, Args... >)
Definition device_permute_impl.hpp:268
static auto MakeDescriptor_N_H_W(const Lengths &lengths, const Strides &stride)
Definition device_permute_impl.hpp:63
GridwisePermute< InGridDesc, OutGridDesc, InDataType, OutDataType, ElementwiseOperation, BlockSize, NPerBlock, HPerBlock, WPerBlock, InBlockLdsExtraW, InBlockTransferThreadClusterLengths, InBlockTransferThreadClusterArrangeOrder, SrcVectorDim -(NumDim - 3), DstVectorDim -(NumDim - 3), SrcScalarPerVector, DstScalarPerVector > GridwisePermute
Definition device_permute_impl.hpp:92
Lengths Strides
Definition device_permute.hpp:20
bool IsSupportedArgument(const BaseArgument *arg) override final
Definition device_permute_impl.hpp:230
std::unique_ptr< BaseArgument > MakeArgumentPointer(const Lengths &in_lengths, const Strides &in_strides, const Lengths &out_lengths, const Strides &out_strides, const void *in_dev_buffer, void *out_dev_buffer, ElementwiseOperation elementwise_op) override final
Definition device_permute_impl.hpp:243
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override final
Definition device_permute_impl.hpp:260
static auto ConvertArrayToTuple(const std::array< index_t, NumDim > &array)
Definition device_permute_impl.hpp:56
typename GridwisePermute::DefaultBlock2TileMap Block2TileMap
Definition device_permute_impl.hpp:110
static bool IsSupportedArgument(const Argument &arg)
Definition device_permute_impl.hpp:190
DevicePermute< NumDim, InDataType, OutDataType, ElementwiseOperation > BaseType
Definition device_permute_impl.hpp:46
std::array< index_t, NumDim > Lengths
Definition device_permute.hpp:19
static std::enable_if_t< std::is_default_constructible_v< Invoker >, Invoker > MakeInvoker() noexcept(std::is_nothrow_default_constructible_v< Invoker >)
Definition device_permute_impl.hpp:274