device_avgpool3d_bwd_ndhwc_ndhwc.hpp Source File#
device_avgpool3d_bwd_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 host_utility/kernel_launch.hpp:14
__host__ __device__ constexpr index_t gcd(index_t x, index_t y)
Definition utility/math.hpp:154
__host__ __device__ constexpr auto integer_least_multiple(X x, Y y)
Definition utility/math.hpp:78
__host__ __device__ constexpr auto integer_divide_floor(X x, Y y)
Definition utility/math.hpp:66
__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 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_slice_transform(const LowLength &low_length, const SliceBegin &slice_begin, const SliceEnd &slice_end)
Definition multi_index_transform_helper.hpp:163
__host__ __device__ constexpr auto make_freeze_transform(const LowerIndex &low_idx)
Definition multi_index_transform_helper.hpp:151
__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_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__ __device__ constexpr auto make_embed_transform(const UpLengths &up_lengths, const Coefficients &coefficients)
Definition multi_index_transform_helper.hpp:48
__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__ __device__ constexpr auto make_merge_transform(const LowLengths &low_lengths)
Definition multi_index_transform_helper.hpp:55
__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_2d_reduction_threadwise.hpp:84
Definition utility/sequence.hpp:43
Definition reduction_operator.hpp:37
Definition device_base.hpp:197
BaseArgument()=default
BaseInvoker()=default
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:349
int num_reduce_
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:423
Div div_element_op_
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:427
std::vector< ck::index_t > dout_n_c_wos_strides_
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:420
const DOutDataType * p_dout_grid_
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:416
std::vector< ck::index_t > din_n_c_wos_length_
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:419
DInDataType * p_din_grid_
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:417
Argument(const DOutDataType *p_dout, DInDataType *p_din, std::vector< ck::index_t > dout_n_c_wos_lengths, std::vector< ck::index_t > din_n_c_wos_length, std::vector< ck::index_t > dout_n_c_wos_strides, std::vector< ck::index_t > din_n_c_wos_strides, std::vector< ck::index_t > window_lengths, std::vector< ck::index_t > window_strides, std::vector< ck::index_t > window_dilations, std::vector< ck::index_t > input_left_pads, std::vector< ck::index_t > input_right_pads)
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:350
std::vector< ck::index_t > dout_n_c_wos_lengths_
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:418
std::vector< DinGridDesc_M > din_grid_desc_m_container_
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:425
std::vector< ck::index_t > din_n_c_wos_strides_
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:421
std::vector< DoutGridDesc_M_K > dout_grid_desc_m_k_container_
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:424
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:431
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:474
float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:432
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:41
static constexpr auto I0
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:44
std::unique_ptr< BaseArgument > MakeArgumentPointer(const void *p_dout, void *p_din, std::vector< ck::index_t > dout_n_c_wos_lengths, std::vector< ck::index_t > din_n_c_wos_length, std::vector< ck::index_t > dout_n_c_wos_strides, std::vector< ck::index_t > din_n_c_wos_strides, std::vector< ck::index_t > window_lengths, std::vector< ck::index_t > window_strides, std::vector< ck::index_t > window_dilations, std::vector< ck::index_t > input_left_pads, std::vector< ck::index_t > input_right_pads) override
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:517
tensor_operation::element_wise::UnaryDivide Div
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:328
GridwiseReduction_mk_to_m_threadwise< DOutDataType, DInDataType, ComputeDataType, int, DoutGridDesc_M_K, DinGridDesc_M, reduce::Add, PassThrough, Div, InMemoryDataOperationEnum::Set, false, BlockSize, MThreadSliceSize, KThreadSliceSize, OutSrcInDstVectorDim, InSrcOutDstVectorSize, InSrcOutDstVectorSize > gridwise_reduce
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:330
static auto Make3DGridDescriptor_Out_M_K_In_M(const std::vector< ck::index_t > &dout_n_c_wos_lengths, const std::vector< ck::index_t > &din_n_c_wos_length, const std::vector< ck::index_t > &dout_n_c_wos_strides, const std::vector< ck::index_t > &din_n_c_wos_strides, const std::vector< ck::index_t > &window_lengths, const std::vector< ck::index_t > &window_strides, const std::vector< ck::index_t > &window_dilations, const std::vector< ck::index_t > &input_left_pads, const std::vector< ck::index_t > &input_right_pads, const std::vector< ck::index_t > &tildes)
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:51
decltype(Make3DGridDescriptor_Out_M_K_In_M({0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0, 0, 0}, {0, 0, 0}, {0, 0, 0}, {0, 0, 0}, {0, 0, 0}, {0, 0, 0}, {0, 0, 0})) DoutDinGridDesc
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:308
remove_cvref_t< tuple_element_t< 0, DoutDinGridDesc > > DoutGridDesc_M_K
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:319
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:511
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:553
static constexpr ck::index_t M_BlockTileSize
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:47
tensor_operation::element_wise::PassThrough PassThrough
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:327
std::string GetTypeString() const override
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:558
static constexpr ck::index_t K_BlockTileSize
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:48
static bool IsSupportedArgument(const Argument &arg)
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:481
static constexpr index_t OutSrcInDstVectorDim
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:325
remove_cvref_t< tuple_element_t< 1, DoutDinGridDesc > > DinGridDesc_M
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:320
static constexpr auto I1
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:45
static constexpr ck::index_t NDimSpatial
Definition device_avgpool3d_bwd_ndhwc_ndhwc.hpp:42
Definition device_avgpool_bwd.hpp:20
Definition tensor_operation/gpu/element/unary_element_wise_operation.hpp:340
Definition tensor_operation/gpu/element/unary_element_wise_operation.hpp:701