diff options
| author | ramelg01 <ramy.elgammal@arm.com> | 2022-02-26 22:06:20 +0000 |
|---|---|---|
| committer | Ramy Elgammal <ramy.elgammal@arm.com> | 2022-03-15 16:18:13 +0000 |
| commit | 375156937a0783432c5d18e199b5d8d2b3ec33f7 (patch) | |
| tree | 4c755770ac569c286179e5ce40df1c820baa73d8 | |
| parent | 4e66d707a292b90a344e32c59eb1dacb67a0e4c1 (diff) | |
| download | ComputeLibrary-375156937a0783432c5d18e199b5d8d2b3ec33f7.tar.gz | |
Implementation of ClPooling3d
- For NDHWC layout
- For F16 and F32 data types
- Mixed Precision stil not supported
Resolves: COMPMID-4670
Signed-off-by: ramy.elgammal@arm.com
Change-Id: I0e14a13e4625569e8e5ee67e6033bd1efe0da469
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/7262
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: SiCong Li <sicong.li@arm.com>
Reviewed-by: Gunes Bayir <gunes.bayir@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
27 files changed, 1762 insertions, 178 deletions
diff --git a/Android.bp b/Android.bp index 491dbca165..c80bba0998 100644 --- a/Android.bp +++ b/Android.bp @@ -109,6 +109,7 @@ opencl_srcs = [ "src/core/CL/cl_kernels/nhwc/normalization_layer.cl", "src/core/CL/cl_kernels/nhwc/normalize_planar_yuv_layer.cl", "src/core/CL/cl_kernels/nhwc/normalize_planar_yuv_layer_quantized.cl", + "src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl", "src/core/CL/cl_kernels/nhwc/pooling_layer.cl", "src/core/CL/cl_kernels/nhwc/pooling_layer_quantized.cl", "src/core/CL/cl_kernels/nhwc/reorg_layer.cl", @@ -604,6 +605,7 @@ cc_library_static { "src/gpu/cl/kernels/ClMulKernel.cpp", "src/gpu/cl/kernels/ClPermuteKernel.cpp", "src/gpu/cl/kernels/ClPool2dKernel.cpp", + "src/gpu/cl/kernels/ClPool3dKernel.cpp", "src/gpu/cl/kernels/ClQuantizeKernel.cpp", "src/gpu/cl/kernels/ClReshapeKernel.cpp", "src/gpu/cl/kernels/ClScaleKernel.cpp", @@ -651,6 +653,7 @@ cc_library_static { "src/gpu/cl/operators/ClPRelu.cpp", "src/gpu/cl/operators/ClPermute.cpp", "src/gpu/cl/operators/ClPool2d.cpp", + "src/gpu/cl/operators/ClPool3d.cpp", "src/gpu/cl/operators/ClQuantize.cpp", "src/gpu/cl/operators/ClReshape.cpp", "src/gpu/cl/operators/ClScale.cpp", @@ -734,6 +737,7 @@ cc_library_static { "src/runtime/CL/functions/CLPadLayer.cpp", "src/runtime/CL/functions/CLPermute.cpp", "src/runtime/CL/functions/CLPixelWiseMultiplication.cpp", + "src/runtime/CL/functions/CLPooling3dLayer.cpp", "src/runtime/CL/functions/CLPoolingLayer.cpp", "src/runtime/CL/functions/CLPriorBoxLayer.cpp", "src/runtime/CL/functions/CLQLSTMLayer.cpp", diff --git a/SConscript b/SConscript index 13ef37a7e2..342c79c7a4 100644 --- a/SConscript +++ b/SConscript @@ -443,6 +443,7 @@ if env['opencl'] and env['embed_kernels']: 'src/core/CL/cl_kernels/nhwc/normalize_planar_yuv_layer.cl', 'src/core/CL/cl_kernels/nhwc/normalize_planar_yuv_layer_quantized.cl', 'src/core/CL/cl_kernels/nhwc/pooling_layer.cl', + 'src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl', 'src/core/CL/cl_kernels/nhwc/pooling_layer_quantized.cl', 'src/core/CL/cl_kernels/nhwc/reorg_layer.cl', 'src/core/CL/cl_kernels/nhwc/scale.cl', diff --git a/arm_compute/core/Types.h b/arm_compute/core/Types.h index 6c49cc35c0..47b9485d80 100644 --- a/arm_compute/core/Types.h +++ b/arm_compute/core/Types.h @@ -1264,6 +1264,109 @@ struct PoolingLayerInfo bool fp_mixed_precision; }; +/** Pooling Layer Information struct*/ +struct Pooling3dLayerInfo +{ + /** Default Constructor */ + Pooling3dLayerInfo() + : pool_type(PoolingType::MAX), + pool_size(Size3D()), + stride(Size3D()), + padding(Padding3D()), + exclude_padding(false), + is_global_pooling(false), + fp_mixed_precision(false), + round_type(DimensionRoundingType::FLOOR) + { + } + /** Constructor + * + * @param[in] pool_type Pooling type @ref PoolingType. + * @param[in] pool_size Pooling size, in elements, across x, y and z. + * @param[in] stride (Optional) stride information @ref Size3D + * @param[in] padding (Optional) padding information @ref Padding3D + * @param[in] exclude_padding (Optional) Strategy when accounting padding in calculations. + * True will exclude padding while false will not (Used in AVG/L2 pooling to determine the pooling area). + * Defaults to false; + * @param[in] fp_mixed_precision (Optional) Use wider accumulators (32 bit instead of 16 for FP16) to improve accuracy. + * @param[in] round_type (Optional) Dimensions rounding. Defaults to @ref FLOOR + */ + explicit Pooling3dLayerInfo(PoolingType pool_type, + unsigned int pool_size, + Size3D stride = Size3D(1U, 1U, 1U), + Padding3D padding = Padding3D(), + bool exclude_padding = false, + bool fp_mixed_precision = false, + DimensionRoundingType round_type = DimensionRoundingType::FLOOR) + : pool_type(pool_type), + pool_size(Size3D(pool_size, pool_size, pool_size)), + stride(stride), + padding(padding), + exclude_padding(exclude_padding), + is_global_pooling(false), + fp_mixed_precision(fp_mixed_precision), + round_type(round_type) + { + } + + /** Constructor + * + * @param[in] pool_type Pooling type @ref PoolingType. + * @param[in] pool_size Pooling size, in elements, across x, y and z. + * @param[in] stride (Optional) stride information @ref Size3D + * @param[in] padding (Optional) padding information @ref Padding3D + * @param[in] exclude_padding (Optional) Strategy when accounting padding in calculations. + * True will exclude padding while false will not (Used in AVG/L2 pooling to determine the pooling area). + * Defaults to false; + * @param[in] fp_mixed_precision (Optional) Use wider accumulators (32 bit instead of 16 for FP16) to improve accuracy. + * @param[in] round_type (Optional) Dimensions rounding. Defaults to @ref FLOOR + */ + explicit Pooling3dLayerInfo(PoolingType pool_type, + Size3D pool_size, + Size3D stride = Size3D(1U, 1U, 1U), + Padding3D padding = Padding3D(), + bool exclude_padding = false, + bool fp_mixed_precision = false, + DimensionRoundingType round_type = DimensionRoundingType::FLOOR) + : pool_type(pool_type), + pool_size(pool_size), + stride(stride), + padding(padding), + exclude_padding(exclude_padding), + is_global_pooling(false), + fp_mixed_precision(fp_mixed_precision), + round_type(round_type) + { + } + + /** Constructor + * + * @note This constructor is used for global pooling + * + * @param[in] pool_type Pooling type @ref PoolingType. + */ + explicit Pooling3dLayerInfo(PoolingType pool_type) + : pool_type(pool_type), + pool_size(Size3D()), + stride(Size3D(1U, 1U, 1U)), + padding(Padding3D(0, 0, 0)), + exclude_padding(false), + is_global_pooling(true), + fp_mixed_precision(false), + round_type(DimensionRoundingType::FLOOR) + { + } + + PoolingType pool_type; + Size3D pool_size; + Size3D stride; + Padding3D padding; + bool exclude_padding; + bool is_global_pooling; + bool fp_mixed_precision; + DimensionRoundingType round_type; +}; + /** ROI Pooling Layer Information class */ class ROIPoolingLayerInfo final { @@ -1307,42 +1410,6 @@ private: unsigned int _sampling_ratio; }; -struct Pool3DInfo -{ - Pool3DInfo() = default; - - /** Constructor - * - * @param[in] pool_type Pooling type @ref PoolingType. - * @param[in] pool_size Pooling size, in elements, across x, y and z @ref Size3D - * @param[in] padding Paddings in x, y and z dimensions - * @param[in] strides Strides in x, y and z dimensions @ref Size3D - * @param[in] round_type Dimension rounding type (ceil or floor) - * @param[in] exclude_padding Strategy when accounting padding in calculations. - * True will exclude padding while false will not (Used in AVG/L2 pooling to determine the pooling area). - * @param[in] is_global_pooling Sets the pool size to the input size if True - */ - Pool3DInfo(const PoolingType pool_type, - const Size3D pool_size, - const Padding3D padding, - const Size3D strides, - const DimensionRoundingType round_type, - bool exclude_padding, - bool is_global_pooling) - : pool_type(pool_type), pool_size(pool_size), padding(padding), strides(strides), round_type(round_type), exclude_padding(exclude_padding), is_global_pooling(is_global_pooling) - - { - } - - PoolingType pool_type{ PoolingType::MAX }; - Size3D pool_size{ 1U, 1U, 1U }; - Padding3D padding{}; - Size3D strides{ 1U, 1U, 1U }; - DimensionRoundingType round_type{ DimensionRoundingType::FLOOR }; - bool exclude_padding{ false }; - bool is_global_pooling{ false }; -}; - /** Generate Proposals Information class */ class GenerateProposalsInfo { diff --git a/arm_compute/core/Utils.h b/arm_compute/core/Utils.h index 2d774770ae..a47cfbdec6 100644 --- a/arm_compute/core/Utils.h +++ b/arm_compute/core/Utils.h @@ -793,6 +793,23 @@ std::pair<int, int> scaled_dimensions_signed(int width, int height, int kernel_width, int kernel_height, const PadStrideInfo &pad_stride_info); +/** Returns calculated width, height and depth of output scaled tensor depending on dimensions rounding mode. + * + * @param[in] width Width of input tensor + * @param[in] height Height of input tensor + * @param[in] depth Depth of input tensor + * @param[in] kernel_width Kernel width. + * @param[in] kernel_height Kernel height. + * @param[in] kernel_depth Kernel depth. + * @param[in] pool3d_info Pad and stride and round information for 3d pooling + * + * @return A tuple with the new width in the first position, the new height in the second, and the new depth in the third. + * Returned values can be < 1 + */ +std::tuple<int, int, int> scaled_3d_dimensions_signed(int width, int height, int depth, + int kernel_width, int kernel_height, int kernel_depth, + const Pooling3dLayerInfo &pool3d_info); + /** Check if the given reduction operation should be handled in a serial way. * * @param[in] op Reduction operation to perform @@ -893,6 +910,23 @@ const std::string &string_from_pooling_type(PoolingType type); * @return True if the pool region is entirely outside the input tensor, False otherwise. */ bool is_pool_region_entirely_outside_input(const PoolingLayerInfo &info); +/** Check if the 3d pool region is entirely outside the input tensor + * + * @param[in] info @ref Pooling3dLayerInfo to be checked. + * + * @return True if the pool region is entirely outside the input tensor, False otherwise. + */ +bool is_pool_3d_region_entirely_outside_input(const Pooling3dLayerInfo &info); +/** Check if the 3D padding is symmetric i.e. padding in each opposite sides are euqal (left=right, top=bottom and front=back) + * + * @param[in] info @ref Padding3D input 3D padding object to check if it is symmetric + * + * @return True if padding is symmetric + */ +inline bool is_symmetric(const Padding3D& info) +{ + return ((info.left == info.right) && (info.top == info.bottom) && (info.front == info.back)); +} /** Translates a given GEMMLowp output stage to a string. * * @param[in] output_stage @ref GEMMLowpOutputStageInfo to be translated to string. diff --git a/arm_compute/core/Window.h b/arm_compute/core/Window.h index 150320a90e..f603e6c148 100644 --- a/arm_compute/core/Window.h +++ b/arm_compute/core/Window.h @@ -47,6 +47,8 @@ public: static constexpr size_t DimZ = 2; /** Alias for dimension 3 also known as W dimension */ static constexpr size_t DimW = 3; + /** Alias for dimension 4 also known as V dimension */ + static constexpr size_t DimV = 4; /** Default constructor: create a window containing a single element. */ constexpr Window() @@ -348,7 +350,6 @@ public: { return slide_window_slice<4>(slice); } - /** Collapse the dimensions between @p first and @p last if possible. * * A dimension is collapsable if it starts from 0 and matches the corresponding dimension in the full_window diff --git a/arm_compute/core/utils/misc/ShapeCalculator.h b/arm_compute/core/utils/misc/ShapeCalculator.h index ee4fe0c02f..df907c106e 100644 --- a/arm_compute/core/utils/misc/ShapeCalculator.h +++ b/arm_compute/core/utils/misc/ShapeCalculator.h @@ -1467,52 +1467,29 @@ inline TensorShape compute_conv3d_shape(const TensorShape &src, const TensorShap * * @return the calculated shape */ -inline TensorShape compute_pool3d_shape(const TensorShape &src, Pool3DInfo pool3d_info) +inline TensorShape compute_pool3d_shape(const TensorShape &src, Pooling3dLayerInfo pool3d_info) { TensorShape output_shape{ src }; - const int idx_width = 1; - const int idx_height = 2; - const int idx_depth = 3; - const int pool_size_width = pool3d_info.is_global_pooling ? src[idx_width] : pool3d_info.pool_size.width; - const int pool_size_height = pool3d_info.is_global_pooling ? src[idx_height] : pool3d_info.pool_size.height; - const int pool_size_depth = pool3d_info.is_global_pooling ? src[idx_depth] : pool3d_info.pool_size.depth; - const int pool_stride_width = pool3d_info.strides.width; - const int pool_stride_height = pool3d_info.strides.height; - const int pool_stride_depth = pool3d_info.strides.depth; - - int output_width_size = 0; - int output_height_size = 0; - int output_depth_size = 0; - - const size_t pad_left = pool3d_info.padding.left; - const size_t pad_right = pool3d_info.padding.right; - const size_t pad_top = pool3d_info.padding.top; - const size_t pad_bottom = pool3d_info.padding.bottom; - const size_t pad_front = pool3d_info.padding.front; - const size_t pad_back = pool3d_info.padding.back; - - switch(pool3d_info.round_type) - { - case DimensionRoundingType::FLOOR: - output_width_size = static_cast<int>(std::floor((static_cast<float>(src[idx_width] + pad_left + pad_right - pool_size_width)) / pool_stride_width) + 1); - output_height_size = static_cast<int>(std::floor((static_cast<float>(src[idx_height] + pad_top + pad_bottom - pool_size_height)) / pool_stride_height) + 1); - output_depth_size = static_cast<int>(std::floor((static_cast<float>(src[idx_depth] + pad_front + pad_back - pool_size_depth)) / pool_stride_depth) + 1); - break; - case DimensionRoundingType::CEIL: - output_width_size = static_cast<int>(std::ceil((static_cast<float>(src[idx_width] + pad_left + pad_right - pool_size_width)) / pool_stride_width) + 1); - output_height_size = static_cast<int>(std::ceil((static_cast<float>(src[idx_height] + pad_top + pad_bottom - pool_size_height)) / pool_stride_height) + 1); - output_depth_size = static_cast<int>(std::ceil((static_cast<float>(src[idx_depth] + pad_front + pad_back - pool_size_depth)) / pool_stride_depth) + 1); - break; - default: - ARM_COMPUTE_ERROR("Unsupported rounding type"); - } - - ARM_COMPUTE_ERROR_ON_MSG((output_width_size < 1 || output_height_size < 1 || output_depth_size < 1), "Calculated output dimension size is invalid"); - - output_shape.set(idx_width, static_cast<size_t>(output_width_size)); - output_shape.set(idx_height, static_cast<size_t>(output_height_size)); - output_shape.set(idx_depth, static_cast<size_t>(output_depth_size)); + const auto data_layout = DataLayout::NDHWC; + const int idx_width = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH); + const int idx_height = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT); + const int idx_depth = get_data_layout_dimension_index(data_layout, DataLayoutDimension::DEPTH); + const int pool_size_width = pool3d_info.is_global_pooling ? src[idx_width] : pool3d_info.pool_size.width; + const int pool_size_height = pool3d_info.is_global_pooling ? src[idx_height] : pool3d_info.pool_size.height; + const int pool_size_depth = pool3d_info.is_global_pooling ? src[idx_depth] : pool3d_info.pool_size.depth; + int output_width = 0; + int output_height = 0; + int output_depth = 0; + + std::tie(output_width, output_height, output_depth) = scaled_3d_dimensions_signed(src[idx_width], src[idx_height], src[idx_depth], pool_size_width, pool_size_height, + pool_size_depth, pool3d_info); + + ARM_COMPUTE_ERROR_ON_MSG((output_width < 1 || output_height < 1 || output_depth < 1), "Calculated output dimension size is invalid"); + + output_shape.set(idx_width, static_cast<size_t>(output_width)); + output_shape.set(idx_height, static_cast<size_t>(output_height)); + output_shape.set(idx_depth, static_cast<size_t>(output_depth)); return output_shape; } diff --git a/arm_compute/runtime/CL/CLFunctions.h b/arm_compute/runtime/CL/CLFunctions.h index 9f67f2e05b..f42da5801c 100644 --- a/arm_compute/runtime/CL/CLFunctions.h +++ b/arm_compute/runtime/CL/CLFunctions.h @@ -84,6 +84,7 @@ #include "arm_compute/runtime/CL/functions/CLPadLayer.h" #include "arm_compute/runtime/CL/functions/CLPermute.h" #include "arm_compute/runtime/CL/functions/CLPixelWiseMultiplication.h" +#include "arm_compute/runtime/CL/functions/CLPooling3dLayer.h" #include "arm_compute/runtime/CL/functions/CLPoolingLayer.h" #include "arm_compute/runtime/CL/functions/CLPriorBoxLayer.h" #include "arm_compute/runtime/CL/functions/CLQLSTMLayer.h" diff --git a/arm_compute/runtime/CL/functions/CLPooling3dLayer.h b/arm_compute/runtime/CL/functions/CLPooling3dLayer.h new file mode 100644 index 0000000000..2e4823756d --- /dev/null +++ b/arm_compute/runtime/CL/functions/CLPooling3dLayer.h @@ -0,0 +1,102 @@ +/* + * Copyright (c) 2022 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef ARM_COMPUTE_CLPOOLING3DLAYER_H +#define ARM_COMPUTE_CLPOOLING3DLAYER_H + +#include "arm_compute/runtime/IFunction.h" + +#include "arm_compute/core/Types.h" + +#include <memory> + +namespace arm_compute +{ +class CLCompileContext; +class ICLTensor; +class ITensorInfo; + +/** Basic function to run @ref opencl::ClPool3d */ +class CLPooling3dLayer : public IFunction +{ +public: + /** Default Constructor */ + CLPooling3dLayer(); + /** Default Destructor */ + ~CLPooling3dLayer(); + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLPooling3dLayer(const CLPooling3dLayer &) = delete; + /** Default move constructor */ + CLPooling3dLayer(CLPooling3dLayer &&) = default; + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLPooling3dLayer &operator=(const CLPooling3dLayer &) = delete; + /** Default move assignment operator */ + CLPooling3dLayer &operator=(CLPooling3dLayer &&) = default; + /** Set the input and output tensors. + * + * Valid data layout: + * - NDHWC + * + * Valid data type configurations: + * |src |dst | + * |:--------------|:--------------| + * |F16 |F16 | + * |F32 |F32 | + * + * @note Source tensor is padded with -inf for MAX pooling and 0 otherwise + * Cases where pooling region is completely outside input tensor are not supported + * + * @note Asymmetric padding is not supported when dimension rounding type == CEIL. + * + * @param[in,out] input Source tensor. Data types supported: F16/F32. + * @param[out] output Destination tensor. Data types supported: Same as @p input. + * @param[in] pool_info Contains 3d pooling operation information described in @ref Pooling3dLayerInfo. + */ + void configure(const ICLTensor *input, ICLTensor *output, const Pooling3dLayerInfo &pool_info); + /** Set the input and output tensors. + * + * @param[in] compile_context The compile context to be used. + * @param[in,out] input Source tensor. Data types supported: F16/F32. + * @param[out] output Destination tensor. Data types supported: Same as @p input. + * @param[in] pool_info Contains 3d pooling operation information described in @ref Pooling3dLayerInfo. + */ + void configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, const Pooling3dLayerInfo &pool_info); + /** Static function to check if given info will lead to a valid configuration of @ref CLPooling3dLayer + * + * @param[in] input Source tensor info. Data types supported: F16/F32. + * @param[in] output Destination tensor info. Data types supported: Same as @p input. + * @param[in] pool_info Contains 3d pooling operation information described in @ref Pooling3dLayerInfo. + * + * @return a status + */ + static Status validate(const ITensorInfo *input, const ITensorInfo *output, const Pooling3dLayerInfo &pool_info); + + // Inherited methods overridden: + void run() override; + +private: + struct Impl; + std::unique_ptr<Impl> _impl; +}; +} // namespace arm_compute +#endif /* ARM_COMPUTE_CLPOOLING3DLAYER_H */ diff --git a/filelist.json b/filelist.json index b72b7b4591..fa43d86d78 100644 --- a/filelist.json +++ b/filelist.json @@ -619,6 +619,15 @@ ] } }, + "Pool3d": { + "files": { + "common": [ + "src/gpu/cl/kernels/ClPool3dKernel.cpp", + "src/gpu/cl/operators/ClPool3d.cpp", + "src/runtime/CL/functions/CLPooling3dLayer.cpp" + ] + } + }, "PRelu": { "deps": [ "ElementwiseBinary" ], "files": { diff --git a/src/core/CL/ICLKernel.cpp b/src/core/CL/ICLKernel.cpp index 9bbc710c88..109a076e9a 100644 --- a/src/core/CL/ICLKernel.cpp +++ b/src/core/CL/ICLKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2021 Arm Limited. + * Copyright (c) 2016-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -173,6 +173,7 @@ template void ICLKernel::add_tensor_argument<1>(unsigned &idx, const ICLTensor * template void ICLKernel::add_tensor_argument<2>(unsigned &idx, const ICLTensor *tensor, const Window &window); template void ICLKernel::add_tensor_argument<3>(unsigned &idx, const ICLTensor *tensor, const Window &window); template void ICLKernel::add_tensor_argument<4>(unsigned &idx, const ICLTensor *tensor, const Window &window); +template void ICLKernel::add_tensor_argument<5>(unsigned &idx, const ICLTensor *tensor, const Window &window); #endif /* DOXYGEN_SKIP_THIS */ void ICLKernel::set_target(cl::Device &device) diff --git a/src/core/CL/ICLKernel.h b/src/core/CL/ICLKernel.h index 4c8028e42a..046679e34e 100644 --- a/src/core/CL/ICLKernel.h +++ b/src/core/CL/ICLKernel.h @@ -238,6 +238,16 @@ public: { add_tensor_argument<4>(idx, tensor, window); } + /** Add the passed 5D tensor's parameters to the object's kernel's arguments starting from the index idx. + * + * @param[in,out] idx Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set. + * @param[in] tensor Tensor to set as an argument of the object's kernel. + * @param[in] window Window the kernel will be executed on. + */ + void add_5D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window) + { + add_tensor_argument<5>(idx, tensor, window); + } /** Add the passed NHW 3D tensor's parameters to the object's kernel's arguments by passing strides, dimensions and the offset to the first valid element in bytes. * diff --git a/src/core/CL/cl_kernels/helpers.h b/src/core/CL/cl_kernels/helpers.h index bfb693e376..4018c40b16 100644 --- a/src/core/CL/cl_kernels/helpers.h +++ b/src/core/CL/cl_kernels/helpers.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2021 Arm Limited. + * Copyright (c) 2016-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -831,6 +831,20 @@ uint name##_step_w, \ uint name##_offset_first_element_in_bytes +#define TENSOR5D_DECLARATION(name) \ + __global uchar *name##_ptr, \ + uint name##_stride_x, \ + uint name##_step_x, \ + uint name##_stride_y, \ + uint name##_step_y, \ + uint name##_stride_z, \ + uint name##_step_z, \ + uint name##_stride_w, \ + uint name##_step_w, \ + uint name##_stride_v, \ + uint name##_step_v, \ + uint name##_offset_first_element_in_bytes + #define CONVERT_TO_VECTOR_STRUCT(name) \ update_vector_workitem_ptr(name##_ptr, name##_offset_first_element_in_bytes, name##_stride_x, name##_step_x) diff --git a/src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl b/src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl new file mode 100644 index 0000000000..7c6414312f --- /dev/null +++ b/src/core/CL/cl_kernels/nhwc/pooling_3d_layer.cl @@ -0,0 +1,190 @@ +/* + * Copyright (c) 2022 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "helpers.h" +#include "tile_helpers.h" // Needed for GET_SPATIAL_IDX() + +#if defined(POOL_AVG) || defined(POOL_L2) +#define POOL_OP(x, y) ((x) + (y)) +#else /* defined(POOL_AVG) || defined(POOL_L2) */ +#define POOL_OP(x, y) (fmax((x), (y))) +#endif /* defined(POOL_AVG) || defined(POOL_L2) */ + +#define SQRT_OP(x) sqrt((x)) + +#if defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(SRC_DEPTH) && defined(DST_CHANNELS) && defined(DST_HEIGHT) && defined(DST_DEPTH) && defined(DST_BATCH_SIZE) && defined(ACC_DATA_TYPE) + +#if defined(POOL_SIZE_X) && defined(POOL_SIZE_Y) && defined(POOL_SIZE_Z) + +/** Performs 3d pooling layer of size equal to MxNXD. This OpenCL kernel can perform the following pooling types: + * -# max, -DPOOL_MAX must be passed at compile time + * -# average, -DPOOL_AVG must be passed at compile time. If padding has to be excluded, -DEXCLUDE_PADDING should be passed at compile time + * -# l2 normalisation, -DPOOL_L2 must be passed at compile time + * + * @note Datatype must be passed at compile type using -DDATA_TYPE e.g. -DDATA_TYPE=half. Supported data types are F32/F16 + * @note Accumulation data type must be passed at compile time using -DACC_DATA_TYPE e.g. -DACC_DATA_TYPE=float + * @note If -DFP_MIXED_PRECISION is passed at compile time, the kernel will use F32 for the partial result + * @note Pool size must be passed at compile time using -DPOOL_SIZE_X, -DPOOL_SIZE_Y, and -DPOOL_SIZE_Z. e.g. -DPOOL_SIZE_X=4, -DPOOL_SIZE_Y=4, -DPOOL_SIZE_Z=2 + * @note Input tensor width, height and depth must be passed at compile time using -DSRC_WIDTH, -DSRC_HEIGHT, and -DSRC_DEPTH + * @note Output tensor height, channels, depth, and batch size must be passed at compile time using -DDST_HEIGHT, -DDST_CHANNELS, -DDST_DEPTH, and -DDST_BATCH_SIZE + * @note Pool strides must be passed at compile time using -DSTRIDE_X, -DSTRIDE_Y and -DSTRIDE_Z which are the steps of the window along the x, y and z directions + * @note Pool pads must be passed at compile time using -DPAD_X, -DPAD_Y, -DPAD_Z + * @note Vector size must be passed at compile time using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16 + * @note Leftover vector size must be passed at compile time using -DVEC_SIZE_LEFTOVER. e.g. -DVEC_SIZE_LEFTOVER=3. It is defined as the remainder between the input's first dimension and VEC_SIZE + * @note The initial value for the pooling operation must be passed at compile time using -DINITIAL_VALUE e.g. -DINITIAL_VALUE=0 + * + * @param[in] input_ptr Pointer to the source tensor. Supported data types: F32/F16 + * @param[in] input_stride_x Stride of the source tensor in X dimension (in bytes) + * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] input_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes) + * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] input_stride_w Stride of the source tensor in W dimension (in bytes) + * @param[in] input_step_w input_stride_w * number of elements along W processed per workitem(in bytes) + * @param[in] input_stride_v Stride of the source tensor in V dimension (in bytes) + * @param[in] input_step_v input_stride_v * number of elements along V processed per workitem(in bytes) + * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source tensor + * @param[out] output_ptr Pointer to the destination tensor. Supported data types: same as @p input_ptr + * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes) + * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] output_stride_z Stride of the destination tensor in Z dimension (in bytes) + * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] output_stride_w Stride of the destination tensor in W dimension (in bytes) + * @param[in] output_step_w output_stride_w * number of elements along W processed per workitem(in bytes) + * @param[in] output_stride_v Stride of the destination tensor in V dimension (in bytes) + * @param[in] output_step_v output_stride_v * number of elements along V processed per workitem(in bytes) + * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor + */ +__kernel void pooling_3d_layer_MxN_ndhwc( + TENSOR5D_DECLARATION(input), + TENSOR5D_DECLARATION(output)) +{ + // Note: If C is not multiple of VEC_SIZE, we shift back of VEC_SIZE_LEFTOVER elements to compute the leftover elements for get_global_id(0) == 0 + // Note: If C is less than VEC_SIZE, VEC_SIZE should be SHRINKED to the closest smaller VEC_SIZE. This operation is performed on the host side + int idx_out_c = GET_SPATIAL_IDX(0, VEC_SIZE, VEC_SIZE_LEFTOVER); + int idx_out_w = GET_SPATIAL_IDX(1, 1, 0); + + // The depth size dimension and the batch size dimension are collapsed over the height dimension + int idx_out_h = GET_SPATIAL_IDX(2, 1, 0) % DST_HEIGHT; + int idx_out_d = (GET_SPATIAL_IDX(2, 1, 0) / DST_HEIGHT) % DST_DEPTH; + int idx_out_n = (GET_SPATIAL_IDX(2, 1, 0) / DST_HEIGHT) / DST_DEPTH; + + __global unsigned char *in_base_ptr = input_ptr + input_offset_first_element_in_bytes + idx_out_c * sizeof(DATA_TYPE) + idx_out_n * input_stride_v; + + __global unsigned char *out_base_ptr = output_ptr + output_offset_first_element_in_bytes + idx_out_c * sizeof(DATA_TYPE) + idx_out_w * output_stride_y + idx_out_h * output_stride_z + idx_out_d * + output_stride_w + idx_out_n * output_stride_v; + + VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE) + res0 = INITIAL_VALUE; + + int idx_in_w = idx_out_w * STRIDE_X - (int)PAD_X; + int idx_in_h = idx_out_h * STRIDE_Y - (int)PAD_Y; + int idx_in_d = idx_out_d * STRIDE_Z - (int)PAD_Z; + + // The start of width to consider in calculation should exclude padding + int pool_x_s = max((int)0, -idx_in_w); + // Assumed Symmetric Padding (left padding = right padding = PAD_X), the filter end should be either the pool width or what is remaining from current pos to the (src width + pad right) + int pool_x_e = min((int)POOL_SIZE_X, (int)SRC_WIDTH + PAD_X - idx_in_w); + int pool_y_s = max((int)0, -idx_in_h); + int pool_y_e = min((int)POOL_SIZE_Y, (int)SRC_HEIGHT + PAD_Y - idx_in_h); + int pool_z_s = max((int)0, -idx_in_d); + int pool_z_e = min((int)POOL_SIZE_Z, (int)SRC_DEPTH + PAD_Z - idx_in_d); + + // The filter size with all padding in all directions considered. + int filter_size = pool_z_e * pool_y_e * pool_x_e; + + // The end of width to consider in calculation should exclude PAD_X + pool_x_e = min(pool_x_e, SRC_WIDTH - idx_in_w); + pool_y_e = min(pool_y_e, SRC_HEIGHT - idx_in_h); + pool_z_e = min(pool_z_e, SRC_DEPTH - idx_in_d); + +#if defined(EXCLUDE_PADDING) + filter_size = (pool_z_e - pool_z_s) * (pool_y_e - pool_y_s) * (pool_x_e - pool_x_s); +#endif // defined(EXCLUDE_PADDING) + +#if POOL_SIZE_X == SRC_WIDTH && POOL_SIZE_Y == SRC_HEIGHT && POOL_SIZE_Z == SRC_DEPTH && PAD_X == 0 && PAD_Y == 0 && PAD_Z == 0 + // Global pooling path + for(int z = 0; z < POOL_SIZE_Z; ++z) + { + int depth_offset_src = (z + idx_in_d) * input_stride_w; + for(int y = 0; y < POOL_SIZE_Y; ++y) + { + int height_offset_src = (y + idx_in_h) * input_stride_z; +#pragma unroll 8 + for(int x = 0; x < POOL_SIZE_X; ++x) + { + int width_offset_src = (x + idx_in_w) * input_stride_y; +#else // POOL_SIZE_X == SRC_WIDTH && POOL_SIZE_Y == SRC_HEIGHT && POOL_SIZE_Z == SRC_DEPTH && PAD_X == 0 && PAD_Y == 0 && PAD_Z == 0 + for(int z = pool_z_s; z < pool_z_e; ++z) + { + int depth_offset_src = (z + idx_in_d) * input_stride_w; + for(int y = pool_y_s; y < pool_y_e; ++y) + { + int height_offset_src = (y + idx_in_h) * input_stride_z; +#pragma unroll 8 + for(int x = pool_x_s; x < pool_x_e; ++x) + { + int width_offset_src = (x + idx_in_w) * input_stride_y; +#endif // POOL_SIZE_X == SRC_WIDTH && POOL_SIZE_Y == SRC_HEIGHT && POOL_SIZE_Z == SRC_DEPTH && PAD_X == 0 && PAD_Y == 0 && PAD_Z == 0 + VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE) + data0; +#if defined(FP_MIXED_PRECISION) + // In case of FP_MIXED_PRECISION, ACC_DATA_TYPE is != DATA_TYPE + data0 = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(in_base_ptr + width_offset_src + height_offset_src + depth_offset_src)), + VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE)); +#else // defined(FP_MIXED_PRECISION) + data0 = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(in_base_ptr + width_offset_src + height_offset_src + depth_offset_src)); +#endif // defined(FP_MIXED_PRECISION) + +#if defined(POOL_L2) + // Raise to power of 2 for L2 Pooling + data0 *= data0; +#endif // defined(POOL_L2) + res0 = POOL_OP(res0, data0); + } + } + } + +#if defined(POOL_AVG) || defined(POOL_L2) + res0 /= (VEC_DATA_TYPE(ACC_DATA_TYPE, VEC_SIZE))filter_size; +#endif // defined(POOL_AVG) || defined(POOL_L2) + +#if defined(POOL_L2) + // Take square root of the result in L2 pooling + res0 = SQRT_OP(res0); +#endif // defined(POOL_L2) + + // Store result +#if defined(FP_MIXED_PRECISION) + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + res_converted0 = CONVERT(res0, VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)); + STORE_VECTOR_SELECT(res_converted, DATA_TYPE, out_base_ptr, VEC_SIZE, VEC_SIZE_LEFTOVER, (VEC_SIZE_LEFTOVER != 0) && get_global_id(0) == 0); +#else // defined(FP_MIXED_PRECISION) + STORE_VECTOR_SELECT(res, DATA_TYPE, out_base_ptr, VEC_SIZE, VEC_SIZE_LEFTOVER, (VEC_SIZE_LEFTOVER != 0) && get_global_id(0) == 0); +#endif // defined(FP_MIXED_PRECISION) +} +#endif // defined(POOL_SIZE_X) && defined(POOL_SIZE_Y) && defined(POOL_SIZE_Z) +#endif // defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(SRC_DEPTH) && defined(DST_CHANNELS) && defined(DST_HEIGHT) && defined(DST_DEPTH) && defined(DST_BATCH_SIZE) && defined(ACC_DATA_TYPE) diff --git a/src/core/Utils.cpp b/src/core/Utils.cpp index ef16cb5e75..904362e1b4 100644 --- a/src/core/Utils.cpp +++ b/src/core/Utils.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2021 Arm Limited. + * Copyright (c) 2016-2022 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -244,6 +244,19 @@ bool is_pool_region_entirely_outside_input(const PoolingLayerInfo &info) return pool_le_padding_x || pool_le_padding_y; } +bool is_pool_3d_region_entirely_outside_input(const Pooling3dLayerInfo &info) +{ + if(info.is_global_pooling || info.pool_size.x() == 0 || info.pool_size.y() == 0 || info.pool_size.z() == 0) + { + return false; + } + const auto ps = info.padding; + const auto pool_le_padding_x = info.pool_size.x() <= std::max({ ps.left, ps.right }); + const auto pool_le_padding_y = info.pool_size.y() <= std::max({ ps.top, ps.bottom }); + const auto pool_le_padding_z = info.pool_size.z() <= std::max({ ps.front, ps.back }); + return pool_le_padding_x || pool_le_padding_y || pool_le_padding_z; +} + const std::string &string_from_gemmlowp_output_stage(GEMMLowpOutputStageType output_stage) { static std::map<GEMMLowpOutputStageType, const std::string> output_stage_map = @@ -474,6 +487,42 @@ std::pair<int, int> scaled_dimensions_signed(int width, int height, return std::make_pair(static_cast<int>(w), static_cast<int>(h)); } +std::tuple<int, int, int> scaled_3d_dimensions_signed(int width, int height, int depth, + int kernel_width, int kernel_height, int kernel_depth, + const Pooling3dLayerInfo &pool3d_info) +{ + const int pad_left = pool3d_info.padding.left; + const int pad_top = pool3d_info.padding.top; + const int pad_right = pool3d_info.padding.right; + const int pad_bottom = pool3d_info.padding.bottom; + const int pad_front = pool3d_info.padding.front; + const int pad_back = pool3d_info.padding.back; + const int stride_x = pool3d_info.stride.x(); + const int stride_y = pool3d_info.stride.y(); + const int stride_z = pool3d_info.stride.z(); + int w = 0; + int h = 0; + int d = 0; + + switch(pool3d_info.round_type) + { + case DimensionRoundingType::FLOOR: + w = static_cast<int>(std::floor((static_cast<float>(width + pad_left + pad_right - kernel_width) / stride_x) + 1)); + h = static_cast<int>(std::floor((static_cast<float>(height + pad_top + pad_bottom - kernel_height) / stride_y) + 1)); + d = static_cast<int>(std::floor((static_cast<float>(depth + pad_front + pad_back - kernel_depth) / stride_z) + 1)); + break; + case DimensionRoundingType::CEIL: + w = static_cast<int>(std::ceil((static_cast<float>(width + pad_left + pad_right - kernel_width) / stride_x) + 1)); + h = static_cast<int>(std::ceil((static_cast<float>(height + pad_top + pad_bottom - kernel_height) / stride_y) + 1)); + d = static_cast<int>(std::ceil((static_cast<float>(depth + pad_front + pad_back - kernel_depth) / stride_z) + 1)); + break; + default: + ARM_COMPUTE_ERROR("Unsupported rounding type"); + } + + return std::make_tuple(static_cast<int>(w), static_cast<int>(h), static_cast<int>(d)); +} + bool needs_serialized_reduction(ReductionOperation op, DataType dt, unsigned int axis) { const bool is_min_max = (op == ReductionOperation::MAX || op == ReductionOperation::MIN); diff --git a/src/gpu/cl/ClKernelLibrary.cpp b/src/gpu/cl/ClKernelLibrary.cpp index bab5342168..a5d37f49c4 100644 --- a/src/gpu/cl/ClKernelLibrary.cpp +++ b/src/gpu/cl/ClKernelLibrary.cpp @@ -437,6 +437,7 @@ const std::map<std::string, std::string> ClKernelLibrary::_kernel_program_map = { "pooling_layer_MxN_nhwc", "nhwc/pooling_layer.cl" }, { "pooling_layer_2x2_nhwc", "nhwc/pooling_layer.cl" }, { "pooling_layer_MxN_quantized_nhwc", "nhwc/pooling_layer_quantized.cl" }, + { "pooling_3d_layer_MxN_ndhwc", "nhwc/pooling_3d_layer.cl" }, { "reorg_layer_nhwc", "nhwc/reorg_layer.cl" }, { "scale_nearest_neighbour_nhwc", "nhwc/scale.cl" }, { "scale_bilinear_nhwc", "nhwc/scale.cl" }, @@ -882,6 +883,10 @@ const std::map<std::string, std::string> ClKernelLibrary::_program_source_map = #include "./cl_kernels/nhwc/pooling_layer.clembed" }, { + "nhwc/pooling_3d_layer.cl", +#include "./cl_kernels/nhwc/pooling_3d_layer.clembed" + }, + { "nhwc/pooling_layer_quantized.cl", #include "./cl_kernels/nhwc/pooling_layer_quantized.clembed" }, diff --git a/src/gpu/cl/kernels/ClPool3dKernel.cpp b/src/gpu/cl/kernels/ClPool3dKernel.cpp new file mode 100644 index 0000000000..929ccf7cb6 --- /dev/null +++ b/src/gpu/cl/kernels/ClPool3dKernel.cpp @@ -0,0 +1,225 @@ +/* + * Copyright (c) 2022 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "src/gpu/cl/kernels/ClPool3dKernel.h" + +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/utils/misc/ShapeCalculator.h" +#include "src/core/CL/CLValidate.h" +#include "src/core/helpers/AutoConfiguration.h" +#include "src/core/helpers/WindowHelpers.h" +#include "support/Cast.h" +#include "utils/TypePrinter.h" + +namespace arm_compute +{ +namespace opencl +{ +namespace kernels +{ +using namespace arm_compute::misc::shape_calculator; + +namespace +{ +Status validate_arguments(const ITensorInfo *src, const ITensorInfo *dst, const Pooling3dLayerInfo &pool_info) +{ + ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(src, dst); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(src->data_layout() != DataLayout::NDHWC, "Only NDHWC layout supported"); + + ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(src); + ARM_COMPUTE_RETURN_ERROR_ON_MSG((pool_info.stride.x() == 0 || pool_info.stride.y() == 0 || pool_info.stride.z() == 0), "Strides cannot be zero."); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(src, 1, DataType::F16, DataType::F32); + + const auto data_layout = src->data_layout(); + const int idx_width = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH); + const int idx_height = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT); + const int idx_depth = get_data_layout_dimension_index(data_layout, DataLayoutDimension::DEPTH); + const bool is_global_pooling = pool_info.is_global_pooling; + const unsigned int pool_size_x = is_global_pooling ? src->dimension(idx_width) : pool_info.pool_size.width; + const unsigned int pool_size_y = is_global_pooling ? src->dimension(idx_height) : pool_info.pool_size.height; + const unsigned int pool_size_z = is_global_pooling ? src->dimension(idx_depth) : pool_info.pool_size.depth; + int output_width = 0; + int output_height = 0; + int output_depth = 0; + + bool round_type_ceil_with_asymm_padding = (pool_info.round_type == DimensionRoundingType::CEIL) && (!is_symmetric(pool_info.padding)); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(round_type_ceil_with_asymm_padding, "Cannot use dimension round type CEIL when padding is asymmetric."); + + ARM_COMPUTE_RETURN_ERROR_ON_MSG(is_pool_3d_region_entirely_outside_input(pool_info), "Pooling region that is entirely outside input tensor is unsupported"); + std::tie(output_width, output_height, output_depth) = scaled_3d_dimensions_signed(src->tensor_shape()[idx_width], src->tensor_shape()[idx_height], + src->tensor_shape()[idx_depth], pool_size_x, pool_size_y, + pool_size_z, pool_info); + + ARM_COMPUTE_RETURN_ERROR_ON_MSG((output_width < 1 || output_height < 1 || output_depth < 1), "Calculated output dimension size is invalid"); + // Checks performed when dst is configured + if(dst->total_size() != 0) + { + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(src, dst); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_LAYOUT(src, dst); + TensorInfo out_info(TensorInfo(compute_pool3d_shape(src->tensor_shape(), pool_info), 1, dst->data_type())); + ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(dst, &out_info); + } + + return Status{}; +} +} // namespace + +ClPool3dKernel::ClPool3dKernel() +{ + _type = CLKernelType::POOL; +} + +void ClPool3dKernel::configure(const ClCompileContext &compile_context, const ITensorInfo *src, ITensorInfo *dst, const Pooling3dLayerInfo &pool_info) +{ + ARM_COMPUTE_ERROR_ON_NULLPTR(src, dst); + ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(src, dst, pool_info)); + auto padding_info = get_padding_info({ src, dst }); + + // Auto init if empty + TensorShape out_shape = compute_pool3d_shape(src->tensor_shape(), pool_info); + auto_init_if_empty(*dst, src->clone()->set_tensor_shape(out_shape)); + + // Set instance variables + _pool_info = pool_info; + _data_layout = src->data_layout(); + + _num_elems_processed_per_iteration = (dst->data_type() == DataType::F32) ? 2 : 4; + _num_elems_processed_per_iteration = adjust_vec_size(_num_elems_processed_per_iteration, dst->dimension(0)); + + const PoolingType pool_type = pool_info.pool_type; + const int idx_width = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::WIDTH); + const int idx_height = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::HEIGHT); + const int idx_depth = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::DEPTH); + const int idx_channel = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::CHANNEL); + const int idx_batch_size = get_data_layout_dimension_index(_data_layout, DataLayoutDimension::BATCHES); + const int pool_size_x = pool_info.is_global_pooling ? src->dimension(idx_width) : pool_info.pool_size.width; + const int pool_size_y = pool_info.is_global_pooling ? src->dimension(idx_height) : pool_info.pool_size.height; + const int pool_size_z = pool_info.is_global_pooling ? src->dimension(idx_depth) : pool_info.pool_size.depth; + const bool exclude_padding = pool_info.exclude_padding; + const int pool_stride_x = pool_info.stride.x(); + const int pool_stride_y = pool_info.stride.y(); + const int pool_stride_z = pool_info.stride.z(); + const int pool_pad_top = pool_info.padding.top; + const int pool_pad_left = pool_info.padding.left; + const int pool_pad_front = pool_info.padding.front; + const DataType data_type = src->data_type(); + + // Set build options + CLBuildOptions build_opts; + build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(_num_elems_processed_per_iteration)); + build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(data_type)); + build_opts.add_option("-DPOOL_" + string_from_pooling_type(pool_type)); + build_opts.add_option("-DSTRIDE_X=" + support::cpp11::to_string(pool_stride_x)); + build_opts.add_option("-DSTRIDE_Y=" + support::cpp11::to_string(pool_stride_y)); + build_opts.add_option("-DSTRIDE_Z=" + support::cpp11::to_string(pool_stride_z)); + build_opts.add_option("-DPAD_X=" + support::cpp11::to_string(pool_pad_left)); + build_opts.add_option("-DPAD_Y=" + support::cpp11::to_string(pool_pad_top)); + build_opts.add_option("-DPAD_Z=" + support::cpp11::to_string(pool_pad_front)); + build_opts.add_option("-DPOOL_SIZE_X=" + support::cpp11::to_string(pool_size_x)); + build_opts.add_option("-DPOOL_SIZE_Y=" + support::cpp11::to_string(pool_size_y)); + build_opts.add_option("-DPOOL_SIZE_Z=" + support::cpp11::to_string(pool_size_z)); + build_opts.add_option("-DSRC_WIDTH=" + support::cpp11::to_string(src->dimension(idx_width))); + build_opts.add_option("-DSRC_HEIGHT=" + support::cpp11::to_string(src->dimension(idx_height))); + build_opts.add_option("-DSRC_DEPTH=" + support::cpp11::to_string(src->dimension(idx_depth))); + + // Set the initial value for the pooling operation accordingly with the data type + if(pool_type == PoolingType::MAX) + { + build_opts.add_option("-DINITIAL_VALUE=" + float_to_string_with_full_precision(std::numeric_limits<float>::lowest())); + } + else + { + // Pool AVG and Pool L2 initial value + build_opts.add_option("-DINITIAL_VALUE=0"); + } + // Create kernel + // Floating point mixed precision is support on F16 only + const auto use_fp_mixed_precision = (data_type == DataType::F16) && pool_info.fp_mixed_precision && pool_type != PoolingType::MAX; + + // Wider accumulation is required to avoid accuracy loss + // Case 1: Floating point mixed precision (fp16 src data and fp32 accumulation) + DataType acc_data_type = data_type; + if(use_fp_mixed_precision) + { + acc_data_type = DataType::F32; + } + build_opts.add_option("-DACC_DATA_TYPE=" + get_cl_type_from_data_type(acc_data_type)); + build_opts.add_option_if(use_fp_mixed_precision, "-DFP_MIXED_PRECISION"); + build_opts.add_option_if(exclude_padding, "-DEXCLUDE_PADDING"); + build_opts.add_option("-DDST_HEIGHT=" + support::cpp11::to_string(dst->dimension(idx_height))); + build_opts.add_option("-DDST_DEPTH=" + support::cpp11::to_string(dst->dimension(idx_depth))); + build_opts.add_option("-DDST_CHANNELS=" + support::cpp11::to_string(dst->dimension(idx_channel))); + build_opts.add_option("-DDST_BATCH_SIZE=" + support::cpp11::to_string(dst->dimension(idx_batch_size))); + build_opts.add_option("-DVEC_SIZE_LEFTOVER=" + support::cpp11::to_string(src->dimension(0) % _num_elems_processed_per_iteration)); + std::string kernel_name = "pooling_3d_layer_MxN_ndhwc"; + _kernel = create_kernel(compile_context, kernel_name, build_opts.options()); + + // Configure kernel window + Window win = calculate_max_window(*dst, Steps(_num_elems_processed_per_iteration)); + ICLKernel::configure_internal(win); + + // Set config_id for enabling LWS tuning + _config_id = "pooling_layer_3d"; + _config_id += lower_string(string_from_data_type(data_type)); + _config_id += "_"; + _config_id += lower_string(string_from_data_layout(_data_layout)); + _config_id += "_"; + _config_id += support::cpp11::to_string(dst->dimension(idx_width)); + _config_id += "_"; + _config_id += support::cpp11::to_string(dst->dimension(idx_height)); + _config_id += "_"; + _config_id += support::cpp11::to_string(dst->dimension(idx_channel)); + _config_id += "_"; + _config_id += lower_string(string_from_data_layout(src->data_layout())); + + ARM_COMPUTE_ERROR_ON(has_padding_changed(padding_info)); +} + +Status ClPool3dKernel::validate(const ITensorInfo *src, const ITensorInfo *dst, const Pooling3dLayerInfo &pool_info) +{ + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(src, dst, pool_info)); + return Status{}; +} + +void ClPool3dKernel::run_op(ITensorPack &tensors, const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window); + + const auto src = utils::cast::polymorphic_downcast<const ICLTensor *>(tensors.get_const_tensor(TensorType::ACL_SRC)); + auto dst = utils::cast::polymorphic_downcast<ICLTensor *>(tensors.get_tensor(TensorType::ACL_DST_0)); + + // Collapse 3D window + Window window_collapsed = window.collapse_if_possible(ICLKernel::window(), Window::DimZ); + + // Set CL kernel arguments + unsigned int idx = 0; + // Passing of the window not needed, as the steps are not used for the pool3d kernel + add_5D_tensor_argument(idx, src, window); + add_5D_tensor_argument(idx, dst, window); + enqueue(queue, *this, window_collapsed, lws_hint()); +} +} // namespace kernels +} // namespace opencl +} // namespace arm_compute diff --git a/src/gpu/cl/kernels/ClPool3dKernel.h b/src/gpu/cl/kernels/ClPool3dKernel.h new file mode 100644 index 0000000000..30c76ed632 --- /dev/null +++ b/src/gpu/cl/kernels/ClPool3dKernel.h @@ -0,0 +1,73 @@ +/* + * Copyright (c) 2022 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef ARM_COMPUTE_CL_POOL3D_KERNEL_H +#define ARM_COMPUTE_CL_POOL3D_KERNEL_H + +#include "src/core/common/Macros.h" +#include "src/gpu/cl/ClCompileContext.h" +#include "src/gpu/cl/IClKernel.h" + +namespace arm_compute +{ +namespace opencl +{ +namespace kernels +{ +/** Interface for the pooling layer kernel */ +class ClPool3dKernel : public IClKernel +{ +public: + ClPool3dKernel(); + ARM_COMPUTE_DISALLOW_COPY_ALLOW_MOVE(ClPool3dKernel); + + /** Configure kernel for a given list of arguments + * + * @note Asymmetric padding is not supported when dimension rounding type == CEIL. + * + * @param[in] compile_context The compile context to be used. + * @param[in] src Source tensor info. Data types supported: F16/F32. + * @param[out] dst Destination tensor info. Data types supported: same as @p src. + * @param[in] pool_info Contains pooling operation information described in @ref Pooling3dLayerInfo. + */ + void configure(const ClCompileContext &compile_context, const ITensorInfo *src, ITensorInfo *dst, const Pooling3dLayerInfo &pool_info); + /** Static function to check if given info will lead to a valid configuration + * + * Similar to ClPool3dKernel::configure() + * + * @return a status + */ + static Status validate(const ITensorInfo *src, const ITensorInfo *dst, const Pooling3dLayerInfo &pool_info); + + // Inherited methods overridden: + void run_op(ITensorPack &tensors, const Window &window, cl::CommandQueue &queue) override; + +private: + Pooling3dLayerInfo _pool_info{}; + DataLayout _data_layout{ DataLayout::UNKNOWN }; + unsigned int _num_elems_processed_per_iteration{ 1 }; +}; +} // namespace kernels +} // namespace opencl +} // namespace arm_compute +#endif /* ARM_COMPUTE_CL_POOL3D_KERNEL_H */ diff --git a/src/gpu/cl/operators/ClPool3d.cpp b/src/gpu/cl/operators/ClPool3d.cpp new file mode 100644 index 0000000000..7dec6c5958 --- /dev/null +++ b/src/gpu/cl/operators/ClPool3d.cpp @@ -0,0 +1,57 @@ +/* + * Copyright (c) 2022 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "src/gpu/cl/operators/ClPool3d.h" + +#include "arm_compute/runtime/CL/CLScheduler.h" + +#include "src/gpu/cl/ClCompileContext.h" +#include "src/gpu/cl/kernels/ClPool3dKernel.h" + +#include "src/common/utils/Log.h" + +namespace arm_compute +{ +namespace opencl +{ +void ClPool3d::configure(const ClCompileContext &compile_context, const ITensorInfo *src, ITensorInfo *dst, const Pooling3dLayerInfo &info) +{ + ARM_COMPUTE_ERROR_ON_NULLPTR(src); + ARM_COMPUTE_LOG_PARAMS(src, dst, info); + + // Configure pooling kernel + auto k = std::make_unique<kernels::ClPool3dKernel>(); + k->set_target(CLScheduler::get().target()); + k->configure(compile_context, src, dst, info); + _kernel = std::move(k); + + // Tune kernels + CLScheduler::get().tune_kernel_static(*_kernel); +} + +Status ClPool3d::validate(const ITensorInfo *src, const ITensorInfo *dst, const Pooling3dLayerInfo &info) +{ + return kernels::ClPool3dKernel::validate(src, dst, info); +} +} // namespace opencl +} // namespace arm_compute diff --git a/src/gpu/cl/operators/ClPool3d.h b/src/gpu/cl/operators/ClPool3d.h new file mode 100644 index 0000000000..7d994fd194 --- /dev/null +++ b/src/gpu/cl/operators/ClPool3d.h @@ -0,0 +1,65 @@ +/* + * Copyright (c) 2022 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef ARM_COMPUTE_CL_POOL3D_H +#define ARM_COMPUTE_CL_POOL3D_H + +#include "src/gpu/cl/ClCompileContext.h" +#include "src/gpu/cl/IClOperator.h" + +#include <memory> + +namespace arm_compute +{ +namespace opencl +{ +/** Basic function to simulate a pooling layer with the specified pooling operation. This function calls the following OpenCL kernels: + * + * -# @ref opencl::ClPool3d + */ +class ClPool3d : public IClOperator +{ +public: + /** Constructor */ + ClPool3d() = default; + /** Configure operator for a given list of arguments + * + * @note Asymmetric padding is not supported when dimension rounding type == CEIL. + * + * @param[in] compile_context The compile context to be used. + * @param[in] src Source tensor info. + * @param[out] dst Destination tensor info. + * @param[in] info 3d Pooling layer parameters. + */ + void configure(const ClCompileContext &compile_context, const ITensorInfo *src, ITensorInfo *dst, const Pooling3dLayerInfo &info); + /** Static function to check if given info will lead to a valid configuration + * + * Similar to ClPool3d::configure() + * + * @return a status + */ + static Status validate(const ITensorInfo *src, const ITensorInfo *dst, const Pooling3dLayerInfo &info); +}; +} // namespace opencl +} // namespace arm_compute +#endif /* ARM_COMPUTE_CL_POOL3D_H */ diff --git a/src/runtime/CL/functions/CLPooling3dLayer.cpp b/src/runtime/CL/functions/CLPooling3dLayer.cpp new file mode 100644 index 0000000000..11ae1d0fe6 --- /dev/null +++ b/src/runtime/CL/functions/CLPooling3dLayer.cpp @@ -0,0 +1,73 @@ +/* + * Copyright (c) 2022 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "arm_compute/runtime/CL/functions/CLPooling3dLayer.h" + +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "src/core/CL/ICLKernel.h" +#include "src/gpu/cl/operators/ClPool3d.h" + +namespace arm_compute +{ +struct CLPooling3dLayer::Impl +{ + const ICLTensor *src{ nullptr }; + ICLTensor *dst{ nullptr }; + ICLTensor *indices{ nullptr }; + std::unique_ptr<opencl::ClPool3d> op{ nullptr }; +}; + +CLPooling3dLayer::CLPooling3dLayer() + : _impl(std::make_unique<Impl>()) +{ +} +CLPooling3dLayer::~CLPooling3dLayer() = default; + +void CLPooling3dLayer::configure(const ICLTensor *input, ICLTensor *output, const Pooling3dLayerInfo &pool_info) +{ + configure(CLKernelLibrary::get().get_compile_context(), input, output, pool_info); +} + +void CLPooling3dLayer::configure(const CLCompileContext &compile_context, const ICLTensor *input, ICLTensor *output, const Pooling3dLayerInfo &pool_info) +{ + _impl->src = input; + _impl->dst = output; + + _impl->op = std::make_unique<opencl::ClPool3d>(); + _impl->op->configure(compile_context, input->info(), output->info(), pool_info); +} + +Status CLPooling3dLayer::validate(const ITensorInfo *input, const ITensorInfo *output, const Pooling3dLayerInfo &pool_info) +{ + return opencl::ClPool3d::validate(input, output, pool_info); +} + +void CLPooling3dLayer::run() +{ + ITensorPack pack; + pack.add_tensor(TensorType::ACL_SRC, _impl->src); + pack.add_tensor(TensorType::ACL_DST_0, _impl->dst); + _impl->op->run(pack); +} +} // namespace arm_compute diff --git a/tests/datasets/Pooling3dLayerDataset.h b/tests/datasets/Pooling3dLayerDataset.h new file mode 100644 index 0000000000..cfe970e8be --- /dev/null +++ b/tests/datasets/Pooling3dLayerDataset.h @@ -0,0 +1,121 @@ +/* + * Copyright (c) 2022 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef ARM_COMPUTE_TEST_POOLING_3D_LAYER_DATASET +#define ARM_COMPUTE_TEST_POOLING_3D_LAYER_DATASET + +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "utils/TypePrinter.h" + +namespace arm_compute +{ +namespace test +{ +namespace datasets +{ +class Pooling3dLayerDataset +{ +public: + using type = std::tuple<TensorShape, Pooling3dLayerInfo>; + + struct iterator + { + iterator(std::vector<TensorShape>::const_iterator src_it, + std::vector<Pooling3dLayerInfo>::const_iterator infos_it) + : _src_it{ std::move(src_it) }, + _infos_it{ std::move(infos_it) } + { + } + + std::string description() const + { + std::stringstream description; + description << "In=" << *_src_it << ":"; + description << "Info=" << *_infos_it << ":"; + return description.str(); + } + + Pooling3dLayerDataset::type operator*() const + { + return std::make_tuple(*_src_it, *_infos_it); + } + + iterator &operator++() + { + ++_src_it; + ++_infos_it; + + return *this; + } + + private: + std::vector<TensorShape>::const_iterator _src_it; + std::vector<Pooling3dLayerInfo>::const_iterator _infos_it; + }; + + iterator begin() const + { + return iterator(_src_shapes.begin(), _infos.begin()); + } + + int size() const + { + return std::min(_src_shapes.size(), _infos.size()); + } + + void add_config(TensorShape src, Pooling3dLayerInfo info) + { + _src_shapes.emplace_back(std::move(src)); + _infos.emplace_back(std::move(info)); + } + +protected: + Pooling3dLayerDataset() = default; + Pooling3dLayerDataset(Pooling3dLayerDataset &&) = default; + +private: + std::vector<TensorShape> _src_shapes{}; + std::vector<Pooling3dLayerInfo> _infos{}; +}; + +// Special pooling dataset +class Pooling3dLayerDatasetSpecial final : public Pooling3dLayerDataset +{ +public: + Pooling3dLayerDatasetSpecial() + { + // Special cases + add_config(TensorShape(2U, 3U, 4U, 2U, 4U), Pooling3dLayerInfo(PoolingType::AVG, /*pool size*/ Size3D(2, 2, 1), /*pool strides*/ Size3D(3, 3, 1), /*pool padding*/ Padding3D(0, 0, 0), true)); + add_config(TensorShape(20U, 22U, 10U, 2U), Pooling3dLayerInfo(PoolingType::AVG, Size3D(100, 100, 100), Size3D(5, 5, 5), Padding3D(50, 50, 50), true)); + add_config(TensorShape(10U, 20U, 32U, 3U, 2U), Pooling3dLayerInfo(PoolingType::MAX, /*pool size*/ 3, /*pool strides*/ Size3D(2, 2, 2), Padding3D(1, 1, 1, 1, 1, 1), false, false, + DimensionRoundingType::FLOOR)); + add_config(TensorShape(14U, 10U, 10U, 3U, 5U), Pooling3dLayerInfo(PoolingType::AVG, Size3D(3, 3, 3), /*pool strides*/ Size3D(3, 3, 3), Padding3D(2, 1, 2), true, false, DimensionRoundingType::CEIL)); + add_config(TensorShape(14U, 10U, 10U, 2U, 4U), Pooling3dLayerInfo(PoolingType::AVG, Size3D(3, 3, 3), /*pool strides*/ Size3D(3, 3, 3), Padding3D(2, 1, 2), false, false, DimensionRoundingType::CEIL)); + add_config(TensorShape(15U, 13U, 13U, 3U, 5U), Pooling3dLayerInfo(PoolingType::AVG, Size3D(4, 4, 4), /*pool strides*/ Size3D(2, 2, 2), Padding3D(2, 2, 2), true, false, DimensionRoundingType::CEIL)); + } +}; +} // namespace datasets +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_POOLING_3D_LAYER_DATASET */ diff --git a/tests/datasets/ShapeDatasets.h b/tests/datasets/ShapeDatasets.h index e21589946b..31c9b9e051 100644 --- a/tests/datasets/ShapeDatasets.h +++ b/tests/datasets/ShapeDatasets.h @@ -540,6 +540,21 @@ public: } }; +/** Data set containing small 5D tensor shapes. */ +class Small5dShapes final : public ShapeDataset +{ +public: + Small5dShapes() + : ShapeDataset("Shape", + { + TensorShape{ 5U, 5U, 7U, 4U, 3U }, + TensorShape{ 5U, 5U, 4U, 13U, 2U }, + TensorShape{ 5U, 5U, 3U, 5U , 2U}, + }) + { + } +}; + /** Data set containing large 5x5 tensor shapes. */ class Large5x5Shapes final : public ShapeDataset { @@ -553,6 +568,19 @@ public: } }; +/** Data set containing large 5D tensor shapes. */ +class Large5dShapes final : public ShapeDataset +{ +public: + Large5dShapes() + : ShapeDataset("Shape", + { + TensorShape{ 30U, 40U, 30U, 32U, 3U } + }) + { + } +}; + /** Data set containing small 5x1 tensor shapes. */ class Small5x1Shapes final : public ShapeDataset { diff --git a/tests/validation/CL/Pooling3dLayer.cpp b/tests/validation/CL/Pooling3dLayer.cpp new file mode 100644 index 0000000000..5c80351da2 --- /dev/null +++ b/tests/validation/CL/Pooling3dLayer.cpp @@ -0,0 +1,283 @@ +/* + * Copyright (c) 2022 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ + +#include "arm_compute/core/TensorShape.h" +#include "tests/framework/datasets/Datasets.h" + +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/CLTensor.h" +#include "arm_compute/runtime/CL/CLTensorAllocator.h" +#include "arm_compute/runtime/CL/functions/CLPooling3dLayer.h" +#include "tests/CL/CLAccessor.h" +#include "tests/PaddingCalculator.h" +#include "tests/datasets/Pooling3dLayerDataset.h" +#include "tests/datasets/PoolingTypesDataset.h" +#include "tests/datasets/ShapeDatasets.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Macros.h" +#include "tests/framework/datasets/Datasets.h" +#include "tests/validation/Validation.h" +#include "tests/validation/fixtures/Pooling3dLayerFixture.h" + +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +namespace +{ +/** Input data sets for floating-point data types */ +const auto Pooling3dLayerDatasetFP = combine(combine(combine(combine(datasets::PoolingTypes(), framework::dataset::make("PoolingSize", { Size3D(2, 3, 2) })), + framework::dataset::make("Stride", { Size3D(1, 1, 1), Size3D(2, 1, 1), Size3D(1, 2, 1), Size3D(2, 2, 1) })), + framework::dataset::make("Padding", { Padding3D(0, 1, 0), Padding3D(1, 1, 1) })), + framework::dataset::make("ExcludePadding", { true, false })); + +const auto Pooling3dLayerDatasetFPSmall = combine(combine(combine(combine(datasets::PoolingTypes(), framework::dataset::make("PoolingSize", { Size3D(2, 2, 2), Size3D(3, 3, 3) })), + framework::dataset::make("Stride", { Size3D(2, 2, 2), Size3D(2, 1, 1) })), + framework::dataset::make("Padding", { Padding3D(0, 0, 0), Padding3D(1, 1, 1), Padding3D(1, 0, 0) })), + framework::dataset::make("ExcludePadding", { true, false })); + +using ShapeDataset = framework::dataset::ContainerDataset<std::vector<TensorShape>>; + +constexpr AbsoluteTolerance<float> tolerance_f32(0.001f); /**< Tolerance value for comparing reference's output against implementation's output for 32-bit floating-point type */ +constexpr AbsoluteTolerance<float> tolerance_f16(0.1f); /**< Tolerance value for comparing reference's output against implementation's output for 16-bit floating-point type */ + +} // namespace + +TEST_SUITE(CL) +TEST_SUITE(Pooling3dLayer) + +// *INDENT-OFF* +// clang-format off +DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( + framework::dataset::make("InputInfo", { TensorInfo(TensorShape(2U, 27U, 13U, 4U, 3U), 1, DataType::F32, DataLayout::NDHWC), // Mismatching data type + TensorInfo(TensorShape(2U, 27U, 13U, 4U, 2U), 1, DataType::F32, DataLayout::NDHWC), // Invalid pad/size combination + TensorInfo(TensorShape(2U, 27U, 13U, 4U, 2U), 1, DataType::F32, DataLayout::NDHWC), // Invalid pad/size combination + TensorInfo(TensorShape(2U, 27U, 13U, 4U, 3U), 1, DataType::F32, DataLayout::NDHWC), // Invalid output shape + TensorInfo(TensorShape(5U, 13U, 15U, 2U, 3U), 1, DataType::F32, DataLayout::NDHWC), // Global Pooling + TensorInfo(TensorShape(13U,13U, 5U, 1U, 2U), 1, DataType::F32, DataLayout::NDHWC), // Invalid output Global Pooling + TensorInfo(TensorShape(5U, 13U, 13U, 4U, 4U), 1, DataType::F32, DataLayout::NDHWC), // Invalid data type + TensorInfo(TensorShape(5U, 13U, 13U, 4U, 4U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(5U, 13U, 13U, 5U, 4U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(1U, 16U, 1U, 3U, 4U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(5U, 13U, 13U, 4U, 3U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(5U, 13U, 13U, 4U, 2U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(5U, 13U, 13U, 4U, 3U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(5U, 13U, 13U, 4U, 3U), 1, DataType::F32, DataLayout::NDHWC), + }), + framework::dataset::make("OutputInfo",{ TensorInfo(TensorShape(2U, 25U, 11U, 3U, 3U), 1, DataType::F16, DataLayout::NDHWC), + TensorInfo(TensorShape(2U, 30U, 11U, 3U, 2U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(2U, 25U, 16U, 3U, 2U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(2U, 27U, 13U, 3U, 3U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(5U, 1U, 1U, 1U, 3U), 1, DataType::F32, DataLayout::NDHWC), // Global pooling applied + TensorInfo(TensorShape(5U, 2U, 2U, 2U, 2U), 1, DataType::F32, DataLayout::NDHWC), // Invalid output Global Pooling + TensorInfo(TensorShape(5U, 12U, 12U, 3U, 4U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(5U, 12U, 12U, 3U, 4U), 1, DataType::QASYMM8, DataLayout::NDHWC), // Invalid data type + TensorInfo(TensorShape(5U, 1U, 1U, 1U, 4U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(1U, 15U, 1U, 2U, 4U), 1, DataType::F32, DataLayout::NDHWC), // Output width larger than input + TensorInfo(TensorShape(5U, 6U, 6U, 2U, 3U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(5U, 6U, 6U, 2U, 2U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(5U, 6U, 6U, 2U, 3U), 1, DataType::F32, DataLayout::NDHWC), + TensorInfo(TensorShape(5U, 6U, 6U, 2U, 3U), 1, DataType::F32, DataLayout::NDHWC), + })), + framework::dataset::make("PoolInfo", { Pooling3dLayerInfo(PoolingType::AVG, 3, Size3D(1, 1, 1), Padding3D(0, 0, 0)), + Pooling3dLayerInfo(PoolingType::AVG, 2, Size3D(1, 1, 1), Padding3D(2, 0, 0)), + Pooling3dLayerInfo(PoolingType::AVG, 2, Size3D(1, 1, 1), Padding3D(0, 0, 0)), + Pooling3dLayerInfo(PoolingType::L2, 3, Size3D(1, 1, 1), Padding3D(0, 0, 0)), + Pooling3dLayerInfo(PoolingType::AVG), + Pooling3dLayerInfo(PoolingType::MAX), + Pooling3dLayerInfo(PoolingType::AVG, 2, Size3D(), Padding3D(), false), + Pooling3dLayerInfo(PoolingType::AVG, 2, Size3D(1U, 1U, 1U), Padding3D(), false), + Pooling3dLayerInfo(PoolingType::AVG), + Pooling3dLayerInfo(PoolingType::MAX, 2, Size3D(1, 1, 2), Padding3D(0, 0, 0), false), + Pooling3dLayerInfo(PoolingType::AVG, 2, Size3D(2U, 2U, 2U), Padding3D(), false), + Pooling3dLayerInfo(PoolingType::AVG, 1, Size3D(2U, 2U, 2U), Padding3D(2, 2, 2), true), // Pool size is smaller than the padding size with padding excluded + Pooling3dLayerInfo(PoolingType::AVG, 1, Size3D(2U, 2U, 2U), Padding3D(2, 2, 2), false), // Pool size is smaller than the padding size with padding included + Pooling3dLayerInfo(PoolingType::AVG, 3, Size3D(2U, 2U, 2U), Padding3D(2,1,2,2,1,2), false, false, DimensionRoundingType::CEIL), // CEIL with asymmetric Padding + })), + framework::dataset::make("Expected", { false, false, false, false, true, false, false, false, true , false, true, false, false, false})), + input_info, output_info, pool_info, expected) +{ + ARM_COMPUTE_EXPECT(bool(CLPooling3dLayer::validate(&input_info.clone()->set_is_resizable(false), &output_info.clone()->set_is_resizable(false), pool_info)) == expected, framework::LogLevel::ERRORS); +} + + +template <typename T> +using CLPooling3dLayerFixture = Pooling3dLayerValidationFixture<CLTensor, CLAccessor, CLPooling3dLayer, T>; + +template <typename T> +using CLSpecialPooling3dLayerFixture = SpecialPooling3dLayerValidationFixture<CLTensor, CLAccessor, CLPooling3dLayer, T>; + +template <typename T> +using CLPooling3dLayerGlobalFixture = Pooling3dLayerGlobalValidationFixture<CLTensor, CLAccessor, CLPooling3dLayer, T>; + +// clang-format on +// *INDENT-ON* +TEST_SUITE(Float) +TEST_SUITE(FP32) + +FIXTURE_DATA_TEST_CASE(RunSpecial, CLSpecialPooling3dLayerFixture<float>, framework::DatasetMode::ALL, datasets::Pooling3dLayerDatasetSpecial() * framework::dataset::make("DataType", DataType::F32)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f32); +} + +FIXTURE_DATA_TEST_CASE(RunSmall, CLPooling3dLayerFixture<float>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small5dShapes(), combine(Pooling3dLayerDatasetFPSmall, + framework::dataset::make("DataType", DataType::F32)))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f32); +} + +FIXTURE_DATA_TEST_CASE(RunLarge, CLPooling3dLayerFixture<float>, framework::DatasetMode::NIGHTLY, combine(datasets::Large5dShapes(), combine(Pooling3dLayerDatasetFP, + framework::dataset::make("DataType", + DataType::F32)))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f32); +} + +TEST_SUITE(GlobalPooling) +// *INDENT-OFF* +// clang-format off +FIXTURE_DATA_TEST_CASE(RunSmall, CLPooling3dLayerFixture<float>, framework::DatasetMode::ALL, + combine(combine(combine(combine(combine(combine( + framework::dataset::make("InputShape", { TensorShape(3U, 27U, 13U, 4U), + TensorShape(4U, 27U, 13U, 4U, 2U) + }), + framework::dataset::make("PoolingType", { PoolingType::AVG, PoolingType::L2, PoolingType::MAX })), + framework::dataset::make("PoolingSize", { Size3D(27, 13, 4) })), + framework::dataset::make("Strides", Size3D(1, 1, 1))), + framework::dataset::make("Paddings", Padding3D(0, 0, 0))), + framework::dataset::make("ExcludePadding", false)), + framework::dataset::make("DataType", DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f32); +} + +FIXTURE_DATA_TEST_CASE(RunSmallGlobal, CLPooling3dLayerGlobalFixture<float>, framework::DatasetMode::ALL, + combine(combine( + framework::dataset::make("InputShape", { TensorShape(27U, 13U, 4U, 3U), + TensorShape(27U, 13U, 4U, 4U, 2U) + }), + framework::dataset::make("PoolingType", { PoolingType::AVG, PoolingType::L2, PoolingType::MAX })), + framework::dataset::make("DataType", DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f32); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLPooling3dLayerFixture<float>, framework::DatasetMode::NIGHTLY, + combine(combine(combine(combine(combine(combine( + framework::dataset::make("InputShape", { TensorShape(4U, 79U, 37U, 11U), + TensorShape(4U, 79U, 37U, 11U, 2U) + }), + framework::dataset::make("PoolingType", { PoolingType::AVG, PoolingType::L2, PoolingType::MAX })), + framework::dataset::make("PoolingSize", { Size3D(79, 37, 11) })), + framework::dataset::make("Strides", Size3D(1, 1, 1))), + framework::dataset::make("Paddings", Padding3D(0, 0, 0))), + framework::dataset::make("ExcludePadding", false)), + framework::dataset::make("DataType", DataType::F32))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f32); +} +// clang-format on +// *INDENT-ON* +TEST_SUITE_END() // GlobalPooling +TEST_SUITE_END() // FP32 + +TEST_SUITE(FP16) + +FIXTURE_DATA_TEST_CASE(RunSmall, CLPooling3dLayerFixture<half>, framework::DatasetMode::PRECOMMIT, combine(datasets::Small5x5Shapes(), combine(Pooling3dLayerDatasetFPSmall, + framework::dataset::make("DataType", DataType::F16)))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f16); +} + +FIXTURE_DATA_TEST_CASE(RunLarge, CLPooling3dLayerFixture<half>, framework::DatasetMode::NIGHTLY, combine(datasets::Large5dShapes(), combine(Pooling3dLayerDatasetFP, + framework::dataset::make("DataType", + DataType::F16)))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f16); +} + +TEST_SUITE(GlobalPooling) +// *INDENT-OFF* +// clang-format off +FIXTURE_DATA_TEST_CASE(RunSmall, CLPooling3dLayerFixture<half>, framework::DatasetMode::ALL, + combine(combine(combine(combine(combine(combine( + framework::dataset::make("InputShape", { TensorShape(3U, 27U, 13U, 4U), + TensorShape(4U, 27U, 13U, 4U, 2U) + }), + framework::dataset::make("PoolingType", { PoolingType::AVG, PoolingType::L2, PoolingType::MAX })), + framework::dataset::make("PoolingSize", { Size3D(27, 13, 4) })), + framework::dataset::make("Strides", Size3D(1, 1, 1))), + framework::dataset::make("Paddings", Padding3D(0, 0, 0))), + framework::dataset::make("ExcludePadding", false)), + framework::dataset::make("DataType", DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f16); +} + +FIXTURE_DATA_TEST_CASE(RunSmallGlobal, CLPooling3dLayerGlobalFixture<half>, framework::DatasetMode::ALL, + combine(combine( + framework::dataset::make("InputShape", { TensorShape(27U, 13U, 4U, 3U), + TensorShape(27U, 13U, 4U, 4U, 2U) + }), + framework::dataset::make("PoolingType", { PoolingType::AVG, PoolingType::L2, PoolingType::MAX })), + framework::dataset::make("DataType", DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f16); +} +FIXTURE_DATA_TEST_CASE(RunLarge, CLPooling3dLayerFixture<half>, framework::DatasetMode::NIGHTLY, + combine(combine(combine(combine(combine(combine( + framework::dataset::make("InputShape", { TensorShape(4U, 79U, 37U, 11U), + TensorShape(4U, 79U, 37U, 11U, 2U) + }), + framework::dataset::make("PoolingType", { PoolingType::AVG, PoolingType::L2, PoolingType::MAX })), + framework::dataset::make("PoolingSize", { Size3D(79, 37, 11) })), + framework::dataset::make("Strides", Size3D(1, 1, 1))), + framework::dataset::make("Paddings", Padding3D(0, 0, 0))), + framework::dataset::make("ExcludePadding", false)), + framework::dataset::make("DataType", DataType::F16))) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_f16); +} +// clang-format on +// *INDENT-ON* +TEST_SUITE_END() // GlobalPooling +TEST_SUITE_END() // FP16 +TEST_SUITE_END() // Float +TEST_SUITE_END() // Pooling3dLayer +TEST_SUITE_END() // CL +} // namespace validation +} // namespace test +} // namespace arm_compute diff --git a/tests/validation/fixtures/Pooling3dLayerFixture.h b/tests/validation/fixtures/Pooling3dLayerFixture.h new file mode 100644 index 0000000000..c1b3519e80 --- /dev/null +++ b/tests/validation/fixtures/Pooling3dLayerFixture.h @@ -0,0 +1,156 @@ +/* + * Copyright (c) 2022 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#ifndef ARM_COMPUTE_TEST_POOLING_3D_LAYER_FIXTURE +#define ARM_COMPUTE_TEST_POOLING_3D_LAYER_FIXTURE + +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/core/utils/misc/ShapeCalculator.h" +#include "arm_compute/runtime/Tensor.h" +#include "tests/AssetsLibrary.h" +#include "tests/Globals.h" +#include "tests/IAccessor.h" +#include "tests/framework/Asserts.h" +#include "tests/framework/Fixture.h" +#include "tests/validation/reference/Pooling3dLayer.h" +#include <random> +namespace arm_compute +{ +namespace test +{ +namespace validation +{ +template <typename TensorType, typename AccessorType, typename FunctionType, typename T> +class Pooling3dLayerValidationGenericFixture : public framework::Fixture +{ +public: + template <typename...> + void setup(TensorShape shape, Pooling3dLayerInfo pool_info, DataType data_type) + { + _target = compute_target(shape, pool_info, data_type); + _reference = compute_reference(shape, pool_info, data_type); + } + +protected: + template <typename U> + void fill(U &&tensor) + { + if(tensor.data_type() == DataType::F32) + { + std::uniform_real_distribution<float> distribution(-1.0f, 1.0f); + library->fill(tensor, distribution, 0); + } + else if(tensor.data_type() == DataType::F16) + { + arm_compute::utils::uniform_real_distribution_16bit<half> distribution{ -1.0f, 1.0f }; + library->fill(tensor, distribution, 0); + } + else // data type is quantized_asymmetric + { + ARM_COMPUTE_ERROR("Passed Type Not Supported"); + } + } + + TensorType compute_target(TensorShape shape, Pooling3dLayerInfo info, + DataType data_type) + { + // Create tensors + TensorType src = create_tensor<TensorType>(shape, data_type, 1, QuantizationInfo(), DataLayout::NDHWC); + const TensorShape dst_shape = misc::shape_calculator::compute_pool3d_shape((src.info()->tensor_shape()), info); + TensorType dst = create_tensor<TensorType>(dst_shape, data_type, 1, QuantizationInfo(), DataLayout::NDHWC); + + // Create and configure function + FunctionType pool_layer; + pool_layer.validate(src.info(), dst.info(), info); + pool_layer.configure(&src, &dst, info); + + ARM_COMPUTE_ASSERT(src.info()->is_resizable()); + ARM_COMPUTE_ASSERT(dst.info()->is_resizable()); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + ARM_COMPUTE_ASSERT(!src.info()->is_resizable()); + ARM_COMPUTE_ASSERT(!dst.info()->is_resizable()); + + // Fill tensors + fill(AccessorType(src)); + + // Compute function + pool_layer.run(); + return dst; + } + + SimpleTensor<T> compute_reference(TensorShape shape, Pooling3dLayerInfo info, DataType data_type) + { + // Create reference + SimpleTensor<T> src(shape, data_type, 1, QuantizationInfo(), DataLayout::NDHWC); + // Fill reference + fill(src); + return reference::pooling_3d_layer<T>(src, info); + } + + TensorType _target{}; + SimpleTensor<T> _reference{}; +}; + +template <typename TensorType, typename AccessorType, typename FunctionType, typename T> +class Pooling3dLayerValidationFixture : public Pooling3dLayerValidationGenericFixture<TensorType, AccessorType, FunctionType, T> +{ +public: + template <typename...> + void setup(TensorShape shape, PoolingType pool_type, Size3D pool_size, Size3D stride, Padding3D padding, bool exclude_padding, DataType data_type) + { + Pooling3dLayerValidationGenericFixture<TensorType, AccessorType, FunctionType, T>::setup(shape, Pooling3dLayerInfo(pool_type, pool_size, stride, padding, exclude_padding), + data_type); + } +}; + +template <typename TensorType, typename AccessorType, typename FunctionType, typename T> +class Pooling3dLayerGlobalValidationFixture : public Pooling3dLayerValidationGenericFixture<TensorType, AccessorType, FunctionType, T> +{ +public: + template <typename...> + void setup(TensorShape shape, PoolingType pool_type, DataType data_type) + { + Pooling3dLayerValidationGenericFixture<TensorType, AccessorType, FunctionType, T>::setup(shape, Pooling3dLayerInfo(pool_type), data_type); + } +}; + +template <typename TensorType, typename AccessorType, typename FunctionType, typename T> +class SpecialPooling3dLayerValidationFixture : public Pooling3dLayerValidationGenericFixture<TensorType, AccessorType, FunctionType, T> +{ +public: + template <typename...> + void setup(TensorShape src_shape, Pooling3dLayerInfo pool_info, DataType data_type) + { + Pooling3dLayerValidationGenericFixture<TensorType, AccessorType, FunctionType, T>::setup(src_shape, pool_info, data_type); + } +}; + +} // namespace validation +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_POOLING_3D_LAYER_FIXTURE */ diff --git a/tests/validation/reference/Pool3D.cpp b/tests/validation/reference/Pooling3dLayer.cpp index 85a594e262..2e8f3a0b92 100644 --- a/tests/validation/reference/Pool3D.cpp +++ b/tests/validation/reference/Pooling3dLayer.cpp @@ -22,7 +22,7 @@ * SOFTWARE. */ -#include "Pool3D.h" +#include "Pooling3dLayer.h" #include "arm_compute/core/utils/misc/ShapeCalculator.h" #include "tests/validation/Helpers.h" @@ -37,7 +37,7 @@ namespace reference using namespace arm_compute::misc::shape_calculator; template <typename T> -SimpleTensor<T> pool3d_internal(const SimpleTensor<T> &src, const Pool3DInfo &pool3d_info, SimpleTensor<uint32_t> *indices) +SimpleTensor<T> pooling_3d_layer_internal(const SimpleTensor<T> &src, const Pooling3dLayerInfo &pool3d_info, SimpleTensor<uint32_t> *indices) { TensorShape pooled_shape = compute_pool3d_shape(src.shape(), pool3d_info); SimpleTensor<T> dst{ pooled_shape, src.data_type(), 1 }; @@ -57,9 +57,9 @@ SimpleTensor<T> pool3d_internal(const SimpleTensor<T> &src, const Pool3DInfo &po const int pool_size_height = pool3d_info.is_global_pooling ? src.shape()[idx_height] : pool3d_info.pool_size.height; const int pool_size_depth = pool3d_info.is_global_pooling ? src.shape()[idx_depth] : pool3d_info.pool_size.depth; - const int pool_stride_width = static_cast<int>(pool3d_info.strides.width); - const int pool_stride_height = static_cast<int>(pool3d_info.strides.height); - const int pool_stride_depth = static_cast<int>(pool3d_info.strides.depth); + const int pool_stride_width = static_cast<int>(pool3d_info.stride.width); + const int pool_stride_height = static_cast<int>(pool3d_info.stride.height); + const int pool_stride_depth = static_cast<int>(pool3d_info.stride.depth); const int pad_left = static_cast<int>(pool3d_info.padding.left); const int pad_top = static_cast<int>(pool3d_info.padding.top); @@ -144,7 +144,6 @@ SimpleTensor<T> pool3d_internal(const SimpleTensor<T> &src, const Pool3DInfo &po { const auto val = static_cast<T>( src[batch_offset_src + depth_offset_src + height_offset_src + x * num_channels + c]); - if(val > max_val) { max_val = val; @@ -189,33 +188,33 @@ SimpleTensor<T> pool3d_internal(const SimpleTensor<T> &src, const Pool3DInfo &po return dst; } -template SimpleTensor<float> pool3d(const SimpleTensor<float> &src, const Pool3DInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices); -template SimpleTensor<half> pool3d(const SimpleTensor<half> &src, const Pool3DInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices); +template SimpleTensor<float> pooling_3d_layer(const SimpleTensor<float> &src, const Pooling3dLayerInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices); +template SimpleTensor<half> pooling_3d_layer(const SimpleTensor<half> &src, const Pooling3dLayerInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices); template <typename T> -SimpleTensor<T> pool3d(const SimpleTensor<T> &src, const Pool3DInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices) +SimpleTensor<T> pooling_3d_layer(const SimpleTensor<T> &src, const Pooling3dLayerInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices) { ARM_COMPUTE_UNUSED(output_qinfo); - return pool3d_internal<T>(src, pool3d_info, indices); + return pooling_3d_layer_internal<T>(src, pool3d_info, indices); } template <> -SimpleTensor<int8_t> pool3d<int8_t>(const SimpleTensor<int8_t> &src, const Pool3DInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices) +SimpleTensor<int8_t> pooling_3d_layer<int8_t>(const SimpleTensor<int8_t> &src, const Pooling3dLayerInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices) { SimpleTensor<float> src_tmp = convert_from_asymmetric(src); - SimpleTensor<float> dst_tmp = pool3d_internal<float>(src_tmp, pool3d_info, indices); + SimpleTensor<float> dst_tmp = pooling_3d_layer_internal<float>(src_tmp, pool3d_info, indices); return convert_to_asymmetric<int8_t>(dst_tmp, output_qinfo); } template <> -SimpleTensor<uint8_t> pool3d<uint8_t>(const SimpleTensor<uint8_t> &src, const Pool3DInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices) +SimpleTensor<uint8_t> pooling_3d_layer<uint8_t>(const SimpleTensor<uint8_t> &src, const Pooling3dLayerInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices) { SimpleTensor<float> src_tmp = convert_from_asymmetric(src); - SimpleTensor<float> dst_tmp = pool3d_internal<float>(src_tmp, pool3d_info, indices); + SimpleTensor<float> dst_tmp = pooling_3d_layer_internal<float>(src_tmp, pool3d_info, indices); return convert_to_asymmetric<uint8_t>(dst_tmp, output_qinfo); } } // namespace reference } // namespace validation } // namespace test -} // namespace arm_compute
\ No newline at end of file +} // namespace arm_compute diff --git a/tests/validation/reference/Pool3D.h b/tests/validation/reference/Pooling3dLayer.h index bdb5744ecc..481a0d3024 100644 --- a/tests/validation/reference/Pool3D.h +++ b/tests/validation/reference/Pooling3dLayer.h @@ -38,10 +38,11 @@ namespace validation namespace reference { template <typename T> -SimpleTensor<T> pool3d_internal(const SimpleTensor<T> &src, const Pool3DInfo &pool3d_info, SimpleTensor<uint32_t> *indices); +SimpleTensor<T> pooling_3d_layer_internal(const SimpleTensor<T> &src, const Pooling3dLayerInfo &pool3d_info, SimpleTensor<uint32_t> *indices = nullptr); template <typename T> -SimpleTensor<T> pool3d(const SimpleTensor<T> &src, const Pool3DInfo &pool3d_info, const QuantizationInfo &output_qinfo, SimpleTensor<uint32_t> *indices); +SimpleTensor<T> pooling_3d_layer(const SimpleTensor<T> &src, const Pooling3dLayerInfo &pool3d_info, const QuantizationInfo &output_qinfo = QuantizationInfo(), + SimpleTensor<uint32_t> *indices = nullptr); } // namespace reference } // namespace validation } // namespace test diff --git a/utils/TypePrinter.h b/utils/TypePrinter.h index 664a39d150..dae81e4a5a 100644 --- a/utils/TypePrinter.h +++ b/utils/TypePrinter.h @@ -2083,6 +2083,122 @@ inline std::string to_string(const PoolingLayerInfo &info) return str.str(); } +/** Formatted output of the Size3D type. + * + * @param[out] os Output stream + * @param[in] size Type to output + * + * @return Modified output stream. + */ +inline ::std::ostream &operator<<(::std::ostream &os, const Size3D &size) +{ + os << size.width << "x" << size.height << "x" << size.depth; + + return os; +} + +/** Formatted output of the Size3D type. + * + * @param[in] type Type to output + * + * @return Formatted string. + */ +inline std::string to_string(const Size3D &type) +{ + std::stringstream str; + str << type; + return str.str(); +} + +/** Formatted output of the Padding3D type. + * + * @param[out] os Output stream. + * @param[in] padding3d Padding info for 3D spatial dimension shape. + * + * @return Modified output stream. + */ +inline ::std::ostream &operator<<(::std::ostream &os, const Padding3D &padding3d) +{ + os << padding3d.left << "," << padding3d.right << "," + << padding3d.top << "," << padding3d.bottom << "," + << padding3d.front << "," << padding3d.back; + return os; +} + +/** Converts a @ref Padding3D to string + * + * @param[in] padding3d Padding3D value to be converted + * + * @return String representing the corresponding Padding3D + */ +inline std::string to_string(const Padding3D &padding3d) +{ + std::stringstream str; + str << padding3d; + return str.str(); +} + +/** Formatted output of the DimensionRoundingType type. + * + * @param[out] os Output stream. + * @param[in] rounding_type DimensionRoundingType Dimension rounding type when down-scaling, or compute output shape of pooling(2D or 3D). + * + * @return Modified output stream. + */ +inline ::std::ostream &operator<<(::std::ostream &os, const DimensionRoundingType &rounding_type) +{ + switch(rounding_type) + { + case DimensionRoundingType::CEIL: + os << "CEIL"; + break; + case DimensionRoundingType::FLOOR: + os << "FLOOR"; + break; + default: + ARM_COMPUTE_ERROR("NOT_SUPPORTED!"); + } + return os; +} + +/** Formatted output of the Pooling 3d Layer Info. + * + * @param[out] os Output stream. + * @param[in] info Pooling 3D layer info to print to output stream. + * + * @return Modified output stream. + */ +inline ::std::ostream &operator<<(::std::ostream &os, const Pooling3dLayerInfo &info) +{ + os << "{Type=" << info.pool_type << "," + << "IsGlobalPooling=" << info.is_global_pooling; + if(!info.is_global_pooling) + { + os << "," + << "PoolSize=" << info.pool_size << ", " + << "Stride=" << info.stride << ", " + << "Padding=" << info.padding << ", " + << "Exclude Padding=" << info.exclude_padding << ", " + << "fp_mixed_precision=" << info.fp_mixed_precision << ", " + << "DimensionRoundingType=" << info.round_type; + } + os << "}"; + return os; +} + +/** Formatted output of the Pooling 3d Layer Info. + * + * @param[in] info Type to output. + * + * @return Formatted string. + */ +inline std::string to_string(const Pooling3dLayerInfo &info) +{ + std::stringstream str; + str << info; + return str.str(); +} + /** Formatted output of the PriorBoxLayerInfo. * * @param[in] info Type to output. @@ -2134,33 +2250,6 @@ inline std::string to_string(const Size2D &type) return str.str(); } -/** Formatted output of the Size3D type. - * - * @param[out] os Output stream - * @param[in] size Type to output - * - * @return Modified output stream. - */ -inline ::std::ostream &operator<<(::std::ostream &os, const Size3D &size) -{ - os << size.width << "x" << size.height << "x" << size.depth; - - return os; -} - -/** Formatted output of the Size2D type. - * - * @param[in] type Type to output - * - * @return Formatted string. - */ -inline std::string to_string(const Size3D &type) -{ - std::stringstream str; - str << type; - return str.str(); -} - /** Formatted output of the ConvolutionMethod type. * * @param[out] os Output stream @@ -3097,29 +3186,6 @@ inline std::string to_string(const BoxNMSLimitInfo &info) return str.str(); } -/** Formatted output of the DimensionRoundingType type. - * - * @param[out] os Output stream. - * @param[in] rounding_type DimensionRoundingType to output. - * - * @return Modified output stream. - */ -inline ::std::ostream &operator<<(::std::ostream &os, const DimensionRoundingType &rounding_type) -{ - switch(rounding_type) - { - case DimensionRoundingType::CEIL: - os << "CEIL"; - break; - case DimensionRoundingType::FLOOR: - os << "FLOOR"; - break; - default: - ARM_COMPUTE_ERROR("NOT_SUPPORTED!"); - } - return os; -} - /** Converts a @ref DimensionRoundingType to string * * @param[in] rounding_type DimensionRoundingType value to be converted @@ -3133,34 +3199,6 @@ inline std::string to_string(const DimensionRoundingType &rounding_type) return str.str(); } -/** Formatted output of the Padding3D type. - * - * @param[out] os Output stream. - * @param[in] padding3d Padding3D to output. - * - * @return Modified output stream. - */ -inline ::std::ostream &operator<<(::std::ostream &os, const Padding3D &padding3d) -{ - os << padding3d.left << "," << padding3d.right << "," - << padding3d.top << "," << padding3d.bottom << "," - << padding3d.front << "," << padding3d.back; - return os; -} - -/** Converts a @ref Padding3D to string - * - * @param[in] padding3d Padding3D value to be converted - * - * @return String representing the corresponding Padding3D - */ -inline std::string to_string(const Padding3D &padding3d) -{ - std::stringstream str; - str << padding3d; - return str.str(); -} - /** Formatted output of the Conv3dInfo type. * * @param[out] os Output stream. |