/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_pool3d_fwd_ndhwc_ndhwc.hpp Source File#
device_pool3d_fwd_ndhwc_ndhwc.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: 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 make_merge_transform(const LowLengths &low_lengths)
Definition: multi_index_transform_helper.hpp:55
__global__ void kernel_reduce_threadwise(const InGridDesc_M_K in_grid_desc_m_k, const OutGridDesc_M out_grid_desc_m, const InElementwiseOperation in_elementwise_op, const AccElementwiseOperation acc_elementwise_op, AccDataType alpha, const InDataType *const __restrict__ p_in_value_global, const IndexDataType *const __restrict__ p_in_index_global, AccDataType beta, OutDataType *const __restrict__ p_out_value_global, IndexDataType *const __restrict__ p_out_index_global)
Definition: gridwise_2d_reduction_threadwise.hpp:28
__host__ constexpr __device__ auto make_embed_transform(const UpLengths &up_lengths, const Coefficients &coefficients)
Definition: multi_index_transform_helper.hpp:48
__host__ constexpr __device__ auto make_pass_through_transform(const LowLength &low_length)
Definition: multi_index_transform_helper.hpp:12
__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_pad_transform(const LowLength &low_length, const LeftPad &left_pad, const RightPad &right_pad, integral_constant< bool, SkipIsValidCheck >=integral_constant< bool, false >{})
Definition: multi_index_transform_helper.hpp:19
__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_2d_reduction_threadwise.hpp:84
Definition: sequence.hpp:43
Definition: integral_constant.hpp:10
Definition: reduction_operator_mapping.hpp:20
Definition: reduction_operator_mapping.hpp:91
static std::tuple< InElementwiseOperation, AccElementwiseOperation > GetElementwiseOperator(int32_t reduceLength)
Definition: reduction_operator_mapping.hpp:96
Definition: device_base.hpp:50
Definition: device_base.hpp:61
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:194
BGridDesc_M b_grid_desc_m_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:242
Argument(const InDataType *p_in_dev, OutDataType *p_out_dev, IndexDataType *p_out_indices_dev, std::vector< ck::index_t > &input_ncdhw_lengths, std::vector< ck::index_t > &output_ncdhw_lengths, std::vector< ck::index_t > &input_ncdhw_stride, std::vector< ck::index_t > &output_ncdhw_stride, std::vector< ck::index_t > &, std::vector< ck::index_t > &window_spatial_zyx_lengths, std::vector< ck::index_t > &window_zyx_strides, std::vector< ck::index_t > &window_zyx_dilations, std::vector< ck::index_t > &input_left_dhw_pads, std::vector< ck::index_t > &input_right_dhw_pads)
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:195
IndexDataType * p_out_indices_dev_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:240
std::vector< ck::index_t > input_ncdhw_lengths_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:248
std::vector< ck::index_t > output_ncdhw_lengths_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:249
const InDataType * p_in_dev_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:238
std::vector< ck::index_t > output_ncdhw_stride_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:251
InElementwiseOperation in_element_op_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:244
AGridDesc_M_K a_grid_desc_m_k_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:241
std::vector< ck::index_t > input_ncdhw_stride_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:250
AccElementwiseOperation acc_element_op_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:245
OutDataType * p_out_dev_
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:239
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:255
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:316
float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:256
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:44
static constexpr index_t WindowRank
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:53
static constexpr auto I2
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:47
static constexpr auto I1
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:46
static constexpr auto I3
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:48
static constexpr ck::index_t M_BlockTileSize
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:63
remove_cvref_t< decltype(ABGridDescs{}[I0])> AGridDesc_M_K
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:190
static auto MakeABGridDescriptor_A_M_K_B_M(std::vector< ck::index_t > input_ncdhw_lengths, std::vector< ck::index_t > output_ncdhw_lengths, std::vector< ck::index_t > input_ncdhw_stride, std::vector< ck::index_t > output_ncdhw_stride, std::vector< ck::index_t > window_spatial_zyx_lengths, std::vector< ck::index_t > window_zyx_strides, std::vector< ck::index_t > window_zyx_dilations, std::vector< ck::index_t > input_left_dhw_pads, std::vector< ck::index_t > input_right_dhw_pads)
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:66
static constexpr ck::index_t K_BlockTileSize
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:64
decltype(MakeABGridDescriptor_A_M_K_B_M({}, {}, {}, {}, {}, {}, {}, {}, {})) ABGridDescs
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:188
static constexpr index_t InOutRank
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:52
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:323
static constexpr auto I4
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:49
typename reduce_unary_operator< ReduceOpId, true, true >::AccElementwiseOperation AccElementwiseOperation
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:61
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:389
virtual std::unique_ptr< BaseArgument > MakeArgumentPointer(const void *p_in_dev, void *p_out_dev, void *p_out_indices_dev, std::vector< ck::index_t > input_ncdhw_lengths, std::vector< ck::index_t > window_zyx_lengths, std::vector< ck::index_t > output_ncdhw_lengths, std::vector< ck::index_t > input_ncdhw_stride, std::vector< ck::index_t > output_ncdhw_stride, std::vector< ck::index_t > indices_ncdhw_stride, std::vector< ck::index_t > window_zyx_strides, std::vector< ck::index_t > window_zyx_dilations, std::vector< ck::index_t > input_left_dhw_pads, std::vector< ck::index_t > input_right_dhw_pads, std::vector< ck::index_t > pooling_dims) override
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:346
typename reduce_unary_operator< ReduceOpId, true, true >::InElementwiseOperation InElementwiseOperation
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:58
remove_cvref_t< decltype(ABGridDescs{}[I1])> BGridDesc_M
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:191
std::string GetTypeString() const override
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:394
typename reduce_binary_operator< ReduceOpId >::opType ReduceOperation
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:55
static constexpr auto I5
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:50
static constexpr auto I0
Definition: device_pool3d_fwd_ndhwc_ndhwc.hpp:45
Definition: device_pool_fwd.hpp:25