/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_put_element_impl.hpp Source File#
device_put_element_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_packed(const Tuple< Lengths... > &lengths)
Definition: tensor_descriptor_helper.hpp:101
__global__ void kernel_put_element_1d(const InGrid1dDesc in_grid_1d_desc, const InDataType *__restrict__ p_in_global, const IndexDataType *__restrict__ p_indices_global, OutDataType *__restrict__ p_out_global, const ElementwiseOperation elementwise_op)
Definition: gridwise_put_element_1d.hpp:17
__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_put_element_1d.hpp:36
Definition: sequence.hpp:43
Definition: integral_constant.hpp:10
Definition: device_base.hpp:50
Definition: device_base.hpp:61
Definition: device_put_element.hpp:22
Definition: device_put_element_impl.hpp:64
const IndexDataType * p_indices_
Definition: device_put_element_impl.hpp:80
Argument(const InDataType *p_input, const IndexDataType *p_indices, OutDataType *p_output, index_t input_length, ElementwiseOperation elementwise_op)
Definition: device_put_element_impl.hpp:65
ElementwiseOperation elementwise_op_
Definition: device_put_element_impl.hpp:83
index_t blockSize_
Definition: device_put_element_impl.hpp:84
index_t input_length_raw_
Definition: device_put_element_impl.hpp:82
OutDataType * p_output_
Definition: device_put_element_impl.hpp:81
const InDataType * p_input_
Definition: device_put_element_impl.hpp:79
Definition: device_put_element_impl.hpp:88
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition: device_put_element_impl.hpp:115
float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition: device_put_element_impl.hpp:89
Definition: device_put_element_impl.hpp:30
decltype(MakeDescriptor_M(1, 1, 1)) InGrid1dDesc
Definition: device_put_element_impl.hpp:53
std::unique_ptr< BaseArgument > MakeArgumentPointer(const void *p_input, const void *p_indices, void *p_output, index_t input_length, index_t, ElementwiseOperation elementwise_op) override
Definition: device_put_element_impl.hpp:133
static auto MakeDescriptor_M(index_t length, index_t gridSize, index_t blockSize)
Definition: device_put_element_impl.hpp:47
GridwisePutElement_1D< InGrid1dDesc, InDataType, IndexDataType, OutDataType, ElementwiseOperation, MemOp, InVectorSize > GridwisePutElement
Definition: device_put_element_impl.hpp:61
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition: device_put_element_impl.hpp:147
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition: device_put_element_impl.hpp:122
static auto PadDescriptor_M_1d(Desc_M desc_m, index_t gridSize, index_t blockSize)
Definition: device_put_element_impl.hpp:32