/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_max_pool_bwd_impl.hpp Source File#
device_max_pool_bwd_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
@ AtomicAdd
__host__ constexpr __device__ auto make_naive_tensor_descriptor_packed(const Tuple< Lengths... > &lengths)
Definition: tensor_descriptor_helper.hpp:101
typename conditional< predicate, X, Y >::type conditional_t
Definition: functional.hpp:115
__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 make_unmerge_transform(const UpLengths &up_lengths, integral_constant< bool, Use24BitIntegerCalculation >=integral_constant< bool, false >{})
Definition: multi_index_transform_helper.hpp:90
__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
__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: stream_config.hpp:10
Definition: gridwise_elementwise_2d.hpp:162
Definition: gridwise_put_element_1d.hpp:36
Definition: multi_index_transform.hpp:13
Definition: sequence.hpp:43
Definition: tuple.hpp:117
Definition: integral_constant.hpp:10
Definition: device_base.hpp:50
void * p_workspace_
Definition: device_base.hpp:57
Definition: device_base.hpp:61
Definition: device_max_pool_bwd.hpp:17
Definition: device_max_pool_bwd_impl.hpp:118
index_t dout_length_raw_
Definition: device_max_pool_bwd_impl.hpp:145
index_t din_length_raw_
Definition: device_max_pool_bwd_impl.hpp:146
index_t blockSize_
Definition: device_max_pool_bwd_impl.hpp:147
const IndexDataType * p_indices_
Definition: device_max_pool_bwd_impl.hpp:143
DInDataType * p_din_
Definition: device_max_pool_bwd_impl.hpp:144
bool windowOverlap_
Definition: device_max_pool_bwd_impl.hpp:148
const DOutDataType * p_dout_
Definition: device_max_pool_bwd_impl.hpp:142
Argument(const DOutDataType *p_dout, const IndexDataType *p_indices, DInDataType *p_din, index_t dout_length, index_t din_length, const std::vector< ck::index_t > &window_lengths, const std::vector< ck::index_t > &window_strides, const std::vector< ck::index_t > &window_dilations)
Definition: device_max_pool_bwd_impl.hpp:119
Definition: device_max_pool_bwd_impl.hpp:152
float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override
Definition: device_max_pool_bwd_impl.hpp:305
float Run(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})
Definition: device_max_pool_bwd_impl.hpp:153
Definition: device_max_pool_bwd_impl.hpp:32
decltype(MakeDescriptor_M(1, 1)) InOutGrid1dDesc
Definition: device_max_pool_bwd_impl.hpp:73
static auto ExpendDescFirstDim(Desc_M desc_m)
Definition: device_max_pool_bwd_impl.hpp:64
bool IsSupportedArgument(const BaseArgument *p_arg) override
Definition: device_max_pool_bwd_impl.hpp:325
std::unique_ptr< BaseArgument > MakeArgumentPointer(const void *p_dout, const void *p_indices, void *p_din, index_t dout_length, index_t din_length, std::vector< ck::index_t > window_lengths, std::vector< ck::index_t > window_strides, std::vector< ck::index_t > window_dilations) override
Definition: device_max_pool_bwd_impl.hpp:337
static constexpr auto I1
Definition: device_max_pool_bwd_impl.hpp:42
ck::tensor_operation::element_wise::UnaryConvert UnaryConvert
Definition: device_max_pool_bwd_impl.hpp:39
ck::tensor_operation::element_wise::PassThrough PassThrough
Definition: device_max_pool_bwd_impl.hpp:38
decltype(ExpendDescFirstDim(InOutGrid1dDesc{})) InOutGrid2dDesc
Definition: device_max_pool_bwd_impl.hpp:74
static constexpr index_t NPerThread
Definition: device_max_pool_bwd_impl.hpp:94
static constexpr auto I0
Definition: device_max_pool_bwd_impl.hpp:41
GridwisePutElement_1D< InOutGrid1dDesc, DOutDataType, IndexDataType, DInDataType, PassThrough, InMemoryDataOperationEnum::Set, InOutVectorSize > GridwisePutElementSet
Definition: device_max_pool_bwd_impl.hpp:82
static constexpr index_t MPerThread
Definition: device_max_pool_bwd_impl.hpp:93
conditional_t< is_same_v< DInDataType, float >||is_same_v< DInDataType, double >, DInDataType, float > DInDataType_AutomicAddPreCast
Definition: device_max_pool_bwd_impl.hpp:36
std::unique_ptr< BaseInvoker > MakeInvokerPointer() override
Definition: device_max_pool_bwd_impl.hpp:358
static constexpr index_t BlockSize
Definition: device_max_pool_bwd_impl.hpp:92
GridwisePutElement_1D< InOutGrid1dDesc, DOutDataType, IndexDataType, DInDataType_AutomicAddPreCast, PassThrough, InMemoryDataOperationEnum::AtomicAdd, InOutVectorSize > GridwisePutElementAtomicAdd
Definition: device_max_pool_bwd_impl.hpp:90
static auto MakeDescriptor_M(index_t length, index_t loop_step)
Definition: device_max_pool_bwd_impl.hpp:57
BlockToCTileMap_M00_N0_M01Adapt< MPerBlock, NPerBlock > Block2TileMap
Definition: device_max_pool_bwd_impl.hpp:98
GridwiseElementwise< Tuple< InOutGrid2dDesc >, Tuple< InOutGrid2dDesc >, Tuple< const DInDataType_AutomicAddPreCast * >, Tuple< DInDataType * >, Block2TileMap, UnaryConvert, BlockSize, MPerBlock, NPerBlock, MPerThread, NPerThread, Sequence< 0, 1 >, Sequence< InOutVectorSize >, Sequence< InOutVectorSize >, I1, I1 > GridwiseCasting
Definition: device_max_pool_bwd_impl.hpp:115
static constexpr index_t NPerBlock
Definition: device_max_pool_bwd_impl.hpp:96
static auto PadDescriptor_M_1d(Desc_M &desc_m, index_t loop_step)
Definition: device_max_pool_bwd_impl.hpp:45
size_t GetWorkSpaceSize(const BaseArgument *pArg) const override
Definition: device_max_pool_bwd_impl.hpp:312
static constexpr index_t MPerBlock
Definition: device_max_pool_bwd_impl.hpp:95
Definition: unary_element_wise_operation.hpp:241
Definition: unary_element_wise_operation.hpp:446