diff options
-rw-r--r-- | arm_compute/core/CL/kernels/CLPoolingLayerKernel.h | 1 | ||||
-rw-r--r-- | src/core/CL/CLKernelLibrary.cpp | 1 | ||||
-rw-r--r-- | src/core/CL/cl_kernels/pooling_layer.cl | 91 | ||||
-rw-r--r-- | src/core/CL/kernels/CLPoolingLayerKernel.cpp | 23 | ||||
-rw-r--r-- | tests/Utils.h | 40 | ||||
-rw-r--r-- | tests/dataset/PoolingTypesDataset.h | 55 | ||||
-rw-r--r-- | tests/dataset/ShapeDatasets.h | 2 | ||||
-rw-r--r-- | tests/dataset/ThresholdDataset.h | 4 | ||||
-rw-r--r-- | tests/validation/CL/PoolingLayer.cpp | 119 | ||||
-rw-r--r-- | tests/validation/Datasets.h | 7 | ||||
-rw-r--r-- | tests/validation/NEON/PoolingLayer.cpp (renamed from tests/validation/NEON/Pooling/PoolingLayer.cpp) | 2 | ||||
-rw-r--r-- | tests/validation/TensorOperations.h | 84 |
12 files changed, 367 insertions, 62 deletions
diff --git a/arm_compute/core/CL/kernels/CLPoolingLayerKernel.h b/arm_compute/core/CL/kernels/CLPoolingLayerKernel.h index 546a40b15e..6c5091ff9e 100644 --- a/arm_compute/core/CL/kernels/CLPoolingLayerKernel.h +++ b/arm_compute/core/CL/kernels/CLPoolingLayerKernel.h @@ -52,6 +52,7 @@ public: * @param[in] 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 pooling operation information described in @ref PoolingLayerInfo. + * Supported pooling sizes : 2, 3 and 7 */ void configure(const ICLTensor *input, ICLTensor *output, const PoolingLayerInfo &pool_info); diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 15a5d90835..3070d4817e 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -219,6 +219,7 @@ const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map = { "pixelwise_mul_int", "pixelwise_mul_int.cl" }, { "pooling_layer_2", "pooling_layer.cl" }, { "pooling_layer_3", "pooling_layer.cl" }, + { "pooling_layer_7", "pooling_layer.cl" }, { "remap_nearest_neighbour", "remap.cl" }, { "remap_bilinear", "remap.cl" }, { "reshape_to_columns", "convolution_layer.cl" }, diff --git a/src/core/CL/cl_kernels/pooling_layer.cl b/src/core/CL/cl_kernels/pooling_layer.cl index 1902df9b7d..6bdb174235 100644 --- a/src/core/CL/cl_kernels/pooling_layer.cl +++ b/src/core/CL/cl_kernels/pooling_layer.cl @@ -41,7 +41,6 @@ float calculate_avg_scale(const int pool_size, const int upper_bound_w, const in /** Performs a pooling function of pool size equal to 2. * - * @note Pooling stride must be passed using -DPOOL_STRIDE e.g -DPOOL_STRIDE=2. Supported strides are 1,2,3 * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16, F32; * @note In case of average pooling -DPOOL_AVG must be provided otherwise max pooling will be performed. * @@ -88,7 +87,7 @@ __kernel void pooling_layer_2( data0 = POOL_OP(data0, data1); DATA_TYPE res = POOL_OP(data0.s0, data0.s1); - // Divide by 4 in case of average pooling + // Divide by pool region in case of average pooling #ifdef POOL_AVG res *= calculate_avg_scale(2, max_dims.x, max_dims.y, paddings.x, paddings.y, strides.x, strides.y); #endif @@ -99,7 +98,6 @@ __kernel void pooling_layer_2( /** Performs a pooling function of pool size equal to 3. * - * @note Pooling stride must be passed using -DPOOL_STRIDE e.g -DPOOL_STRIDE=2. Supported strides are 1,2,3 * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16, F32; * @note In case of average pooling -DPOOL_AVG must be provided otherwise max pooling will be performed. * @@ -149,7 +147,7 @@ __kernel void pooling_layer_3( data0 = POOL_OP(data0, data2); DATA_TYPE res = POOL_OP(POOL_OP(data0.s0, data0.s1), data0.s2); - // Divide by 4 in case of average pooling + // Divide by pool region in case of average pooling #ifdef POOL_AVG res *= calculate_avg_scale(3, max_dims.x, max_dims.y, paddings.x, paddings.y, strides.x, strides.y); #endif @@ -157,3 +155,88 @@ __kernel void pooling_layer_3( // Store result *(__global DATA_TYPE *)output.ptr = res; } + +/** Performs a pooling function of pool size equal to 7. + * + * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16, F32; + * @note In case of average pooling -DPOOL_AVG must be provided otherwise max pooling will be performed. + * + * @param[in] input_ptr Pointer to the source image. Supported data types: F16, F32 + * @param[in] input_stride_x Stride of the source image 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 image 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_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] output_ptr Pointer to the destination image. Supported data types: F16, F32 + * @param[in] output_stride_x Stride of the destination image 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 image 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 source 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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[in] max_dims The maximum index that can be accessed in x and y dimension (width + pad) + * @param[in] strides The pooling operation strides in each dimension + * @param[in] paddings The pooling operation paddings in each dimension + */ +__kernel void pooling_layer_7( + TENSOR3D_DECLARATION(input), + TENSOR3D_DECLARATION(output) +#ifdef POOL_AVG + , + int2 max_dims, int2 strides, int2 paddings +#endif +) +{ + // Get pixels pointer + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input); + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); + + // Load data + VEC_DATA_TYPE(DATA_TYPE, 8) + data0 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 0, 0)); + VEC_DATA_TYPE(DATA_TYPE, 8) + data1 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 1, 0)); + VEC_DATA_TYPE(DATA_TYPE, 8) + data2 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 2, 0)); + VEC_DATA_TYPE(DATA_TYPE, 8) + data3 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 3, 0)); + VEC_DATA_TYPE(DATA_TYPE, 8) + data4 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 4, 0)); + VEC_DATA_TYPE(DATA_TYPE, 8) + data5 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 5, 0)); + VEC_DATA_TYPE(DATA_TYPE, 8) + data6 = vload8(0, (__global DATA_TYPE *)tensor3D_offset(&input, 0, 6, 0)); + + // Pool operation of all rows + data0 = POOL_OP(data0, data1); + data2 = POOL_OP(data2, data3); + data4 = POOL_OP(data4, data5); + data0 = POOL_OP(data0, data2); + data4 = POOL_OP(data4, data6); + data0 = POOL_OP(data0, data4); + + // Set last element +#ifdef POOL_AVG + data0.s7 = 0; +#else + data0.s7 = data0.s6; +#endif + + // Reduce result + VEC_DATA_TYPE(DATA_TYPE, 4) + reduce4 = POOL_OP(data0.s0123, data0.s4567); + VEC_DATA_TYPE(DATA_TYPE, 2) + reduce2 = POOL_OP(reduce4.s01, reduce4.s23); + DATA_TYPE res = POOL_OP(reduce2.s0, reduce2.s1); + + // Divide by pool region in case of average pooling +#ifdef POOL_AVG + res *= calculate_avg_scale(7, max_dims.x, max_dims.y, paddings.x, paddings.y, strides.x, strides.y); +#endif + + // Store result + *(__global DATA_TYPE *)output.ptr = res; +} diff --git a/src/core/CL/kernels/CLPoolingLayerKernel.cpp b/src/core/CL/kernels/CLPoolingLayerKernel.cpp index dc5ae4ec7a..7648025caa 100644 --- a/src/core/CL/kernels/CLPoolingLayerKernel.cpp +++ b/src/core/CL/kernels/CLPoolingLayerKernel.cpp @@ -65,10 +65,13 @@ void CLPoolingLayerKernel::configure(const ICLTensor *input, ICLTensor *output, std::tie(pool_pad_x, pool_pad_y) = pad_stride_info.pad(); std::tie(pool_stride_x, pool_stride_y) = pad_stride_info.stride(); + static const std::set<int> supported_pool_sizes = { 2, 3, 7 }; + ARM_COMPUTE_UNUSED(supported_pool_sizes); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32); ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::F16, DataType::F32); ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); - ARM_COMPUTE_ERROR_ON(2 != pool_size && 3 != pool_size); + ARM_COMPUTE_ERROR_ON(supported_pool_sizes.find(pool_size) == supported_pool_sizes.end()); ARM_COMPUTE_ERROR_ON(pool_pad_x >= pool_size || pool_pad_y >= pool_size); // Check output dimensions @@ -82,10 +85,11 @@ void CLPoolingLayerKernel::configure(const ICLTensor *input, ICLTensor *output, ARM_COMPUTE_UNUSED(pooled_h); ARM_COMPUTE_ERROR_ON((output->info()->dimension(0) != pooled_w) || (output->info()->dimension(1) != pooled_h)); - const int input_width = input->info()->dimension(0); - const int input_height = input->info()->dimension(1); - const int upper_bound_w = ((pooled_w - 1) * pool_stride_x - pool_pad_x + pool_size) - input_width; - const int upper_bound_h = ((pooled_h - 1) * pool_stride_y - pool_pad_y + pool_size) - input_height; + const int num_elements_read_per_iteration = (pool_size == 7) ? 8 : pool_size; + const int input_width = input->info()->dimension(0); + const int input_height = input->info()->dimension(1); + const int upper_bound_w = ((pooled_w - 1) * pool_stride_x - pool_pad_x + num_elements_read_per_iteration) - input_width; + const int upper_bound_h = ((pooled_h - 1) * pool_stride_y - pool_pad_y + pool_size) - input_height; // Set instance variables _input = input; @@ -138,17 +142,12 @@ void CLPoolingLayerKernel::configure(const ICLTensor *input, ICLTensor *output, } // Configure kernel window - const unsigned int num_elems_processed_per_iteration = 1; - - Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration)); - + const unsigned int num_elems_processed_per_iteration = 1; + Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration)); AccessWindowStatic input_access(input->info(), -pool_pad_x, -pool_pad_y, input_width + _border_size.right, input_height + _border_size.bottom); AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); - update_window_and_padding(win, input_access, output_access); - output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape())); - ICLKernel::configure(win); } diff --git a/tests/Utils.h b/tests/Utils.h index 53f749df48..f3622cafaa 100644 --- a/tests/Utils.h +++ b/tests/Utils.h @@ -615,13 +615,47 @@ inline int coord2index(const TensorShape &shape, const Coordinates &coord) return index; } +/** Check if Coordinates dimensionality can match the respective shape one. + * + * @param coords Coordinates + * @param shape Shape to match dimensionality + * + * @return True if Coordinates can match the dimensionality of the shape else false. + */ +inline bool match_shape(Coordinates &coords, const TensorShape &shape) +{ + auto check_nz = [](unsigned int i) + { + return i != 0; + }; + + unsigned int coords_dims = coords.num_dimensions(); + unsigned int shape_dims = shape.num_dimensions(); + + // Increase coordinates scenario + if(coords_dims < shape_dims) + { + coords.set_num_dimensions(shape_dims); + return true; + } + // Decrease coordinates scenario + if(coords_dims > shape_dims && !std::any_of(coords.begin() + shape_dims, coords.end(), check_nz)) + { + coords.set_num_dimensions(shape_dims); + return true; + } + + return (coords_dims == shape_dims); +} + /** Check if a coordinate is within a valid region */ inline bool is_in_valid_region(const ValidRegion &valid_region, const Coordinates &coord) { - ARM_COMPUTE_ERROR_ON_MSG(valid_region.shape.num_dimensions() != coord.num_dimensions(), "Shapes of valid region and coordinates do not agree"); - for(int d = 0; static_cast<size_t>(d) < coord.num_dimensions(); ++d) + Coordinates coords(coord); + ARM_COMPUTE_ERROR_ON_MSG(!match_shape(coords, valid_region.shape), "Shapes of valid region and coordinates do not agree"); + for(int d = 0; static_cast<size_t>(d) < coords.num_dimensions(); ++d) { - if(coord[d] < valid_region.start(d) || coord[d] >= valid_region.end(d)) + if(coords[d] < valid_region.start(d) || coords[d] >= valid_region.end(d)) { return false; } diff --git a/tests/dataset/PoolingTypesDataset.h b/tests/dataset/PoolingTypesDataset.h new file mode 100644 index 0000000000..c78a20b809 --- /dev/null +++ b/tests/dataset/PoolingTypesDataset.h @@ -0,0 +1,55 @@ +/* + * Copyright (c) 2017 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_DATASET_POOLING_TYPE_DATASET_H__ +#define __ARM_COMPUTE_TEST_DATASET_POOLING_TYPE_DATASET_H__ + +#include "arm_compute/core/Types.h" +#include "dataset/GenericDataset.h" + +#ifdef BOOST +#include "boost_wrapper.h" +#endif + +namespace arm_compute +{ +namespace test +{ +/** Data set containing all possible pooling types. + * + * Can be used as input for Boost data test cases to automatically run a test + * case on all pooling types. + */ +class PoolingTypes final : public GenericDataset<PoolingType, 2> +{ +public: + PoolingTypes() + : GenericDataset{ PoolingType::MAX, PoolingType::AVG } + { + } + + ~PoolingTypes() = default; +}; +} // namespace test +} // namespace arm_compute +#endif //__ARM_COMPUTE_TEST_DATASET_POOLING_TYPE_DATASET_H__ diff --git a/tests/dataset/ShapeDatasets.h b/tests/dataset/ShapeDatasets.h index d2b82cae40..ecb478dbf0 100644 --- a/tests/dataset/ShapeDatasets.h +++ b/tests/dataset/ShapeDatasets.h @@ -118,7 +118,7 @@ class SmallShapes final : public ShapeDataset<3> { public: SmallShapes() - : ShapeDataset(TensorShape(5U, 5U), + : ShapeDataset(TensorShape(7U, 7U), TensorShape(27U, 13U, 2U), TensorShape(128U, 64U, 1U, 3U)) { diff --git a/tests/dataset/ThresholdDataset.h b/tests/dataset/ThresholdDataset.h index 956cf3d54d..a2d76e3c48 100644 --- a/tests/dataset/ThresholdDataset.h +++ b/tests/dataset/ThresholdDataset.h @@ -58,8 +58,8 @@ public: std::stringstream ss; ss << "Threshold"; ss << "_threshold_value" << threshold; - ss << "_false_value" << false_value; - ss << "_true_value" << true_value; + ss << "_false_value" << std::boolalpha << false_value; + ss << "_true_value" << std::boolalpha << true_value; ss << "_type"; ss << ((type == ThresholdType::BINARY) ? "binary" : "range"); ss << "_upper" << upper; diff --git a/tests/validation/CL/PoolingLayer.cpp b/tests/validation/CL/PoolingLayer.cpp new file mode 100644 index 0000000000..1d0e745088 --- /dev/null +++ b/tests/validation/CL/PoolingLayer.cpp @@ -0,0 +1,119 @@ +/* + * Copyright (c) 2017 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 "CL/CLAccessor.h" +#include "CL/Helper.h" +#include "TypePrinter.h" +#include "arm_compute/runtime/CL/functions/CLPoolingLayer.h" +#include "tests/dataset/PoolingLayerDataset.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" + +#include <random> + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::cl; +using namespace arm_compute::test::validation; + +namespace +{ +const float tolerance_f = 1e-05; /**< Tolerance value for comparing reference's output against implementation's output for float input */ + +/** Compute CL pooling layer function. + * + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt Data type of input and output tensors. + * @param[in] pool_info Pooling Layer information. + * + * @return Computed output tensor. + */ +CLTensor compute_pooling_layer(const TensorShape &shape_in, const TensorShape &shape_out, DataType dt, PoolingLayerInfo pool_info) +{ + // Create tensors + CLTensor src = create_tensor(shape_in, dt); + CLTensor dst = create_tensor(shape_out, dt); + + // Create and configure function + CLPoolingLayer pool; + pool.configure(&src, &dst, pool_info); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + BOOST_TEST(!src.info()->is_resizable()); + BOOST_TEST(!dst.info()->is_resizable()); + + // Fill tensors + std::uniform_real_distribution<> distribution(-1, 1); + library->fill(CLAccessor(src), distribution, 0); + + // Compute function + pool.run(); + + return dst; +} + +TensorShape get_output_shape(TensorShape in_shape, const PoolingLayerInfo &pool_info) +{ + TensorShape out_shape(in_shape); + const std::pair<unsigned int, unsigned int> scaled_dims = arm_compute::scaled_dimensions(in_shape.x(), + in_shape.y(), + pool_info.pool_size(), + pool_info.pad_stride_info().stride().first, pool_info.pad_stride_info().stride().second, + pool_info.pad_stride_info().pad().first, pool_info.pad_stride_info().pad().second, + pool_info.pad_stride_info().round()); + out_shape.set(0, scaled_dims.first); + out_shape.set(1, scaled_dims.second); + return out_shape; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(CL) +BOOST_AUTO_TEST_SUITE(PoolingLayer) + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * CNNFloatDataTypes() * PoolingTypes() * boost::unit_test::data::make({ 2, 3, 7 }) * boost::unit_test::data::make({ 1, 2 }) * boost::unit_test::data::make({ 0, 1 }), + src_shape, dt, pool_type, pool_size, pool_stride, pool_pad) +{ + PoolingLayerInfo pool_info(pool_type, pool_size, PadStrideInfo(pool_stride, pool_stride, pool_pad, pool_pad, DimensionRoundingType::CEIL)); + TensorShape dst_shape = get_output_shape(src_shape, pool_info); + + // Compute function + CLTensor dst = compute_pooling_layer(src_shape, dst_shape, dt, pool_info); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_pooling_layer(src_shape, dst_shape, dt, pool_info); + + // Validate output + validate(CLAccessor(dst), ref_dst, tolerance_f); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif diff --git a/tests/validation/Datasets.h b/tests/validation/Datasets.h index ab21787f45..33776d2e44 100644 --- a/tests/validation/Datasets.h +++ b/tests/validation/Datasets.h @@ -36,6 +36,7 @@ #include "dataset/InterpolationPolicyDataset.h" #include "dataset/NormalizationTypeDataset.h" #include "dataset/PoolingLayerDataset.h" +#include "dataset/PoolingTypesDataset.h" #include "dataset/RoundingPolicyDataset.h" #include "dataset/ShapeDatasets.h" #include "dataset/ThresholdDataset.h" @@ -186,6 +187,12 @@ struct is_dataset<arm_compute::test::RoundingPolicies> : boost::mpl::true_ /// Register the data set with Boost template <> +struct is_dataset<arm_compute::test::PoolingTypes> : boost::mpl::true_ +{ +}; + +/// Register the data set with Boost +template <> struct is_dataset<arm_compute::test::AlexNetConvolutionLayerDataset> : boost::mpl::true_ { }; diff --git a/tests/validation/NEON/Pooling/PoolingLayer.cpp b/tests/validation/NEON/PoolingLayer.cpp index b15ad1c5e6..10b9a5250e 100644 --- a/tests/validation/NEON/Pooling/PoolingLayer.cpp +++ b/tests/validation/NEON/PoolingLayer.cpp @@ -96,7 +96,6 @@ Tensor compute_pooling_layer(const TensorShape &shape_in, const TensorShape &sha #ifndef DOXYGEN_SKIP_THIS BOOST_AUTO_TEST_SUITE(NEON) -BOOST_AUTO_TEST_SUITE(Pooling) BOOST_AUTO_TEST_SUITE(PoolingLayer) BOOST_AUTO_TEST_SUITE(Float) @@ -135,5 +134,4 @@ BOOST_AUTO_TEST_SUITE_END() BOOST_AUTO_TEST_SUITE_END() BOOST_AUTO_TEST_SUITE_END() -BOOST_AUTO_TEST_SUITE_END() #endif diff --git a/tests/validation/TensorOperations.h b/tests/validation/TensorOperations.h index 119be02423..7337924b47 100644 --- a/tests/validation/TensorOperations.h +++ b/tests/validation/TensorOperations.h @@ -1154,47 +1154,57 @@ void pooling_layer(const Tensor<T> &in, Tensor<T> &out, PoolingLayerInfo pool_in std::tie(pool_stride_x, pool_stride_y) = pool_info.pad_stride_info().stride(); std::tie(pad_x, pad_y) = pool_info.pad_stride_info().pad(); - const int cols_in = static_cast<int>(in.shape()[0]); - const int rows_in = static_cast<int>(in.shape()[1]); + const int w_in = static_cast<int>(in.shape()[0]); + const int h_in = static_cast<int>(in.shape()[1]); - const int cols_out = static_cast<int>(out.shape()[0]); - const int rows_out = static_cast<int>(out.shape()[1]); + const int w_out = static_cast<int>(out.shape()[0]); + const int h_out = static_cast<int>(out.shape()[1]); - int upper_dims = in.shape().total_size() / (cols_in * rows_in); + int upper_dims = in.shape().total_size() / (w_in * h_in); - int pooled_height = static_cast<int>(ceil(static_cast<float>(rows_in + 2 * pad_x - pool_size) / pool_stride_x)) + 1; - int pooled_width = static_cast<int>(ceil(static_cast<float>(cols_in + 2 * pad_y - pool_size) / pool_stride_y)) + 1; + int pooled_w = 0; + int pooled_h = 0; + if(pool_info.pad_stride_info().round() == DimensionRoundingType::CEIL) + { + pooled_w = static_cast<int>(ceil(static_cast<float>(w_in + 2 * pad_x - pool_size) / pool_stride_x)) + 1; + pooled_h = static_cast<int>(ceil(static_cast<float>(h_in + 2 * pad_y - pool_size) / pool_stride_y)) + 1; + } + else + { + pooled_w = static_cast<int>(floor(static_cast<float>(w_in + 2 * pad_x - pool_size) / pool_stride_x)) + 1; + pooled_h = static_cast<int>(floor(static_cast<float>(h_in + 2 * pad_y - pool_size) / pool_stride_y)) + 1; + } - if((pooled_height - 1) * pool_stride_x >= rows_in + pad_x) + if((pooled_w - 1) * pool_stride_x >= w_in + pad_x) { - --pooled_height; + --pooled_w; } - if((pooled_width - 1) * pool_stride_y >= cols_in + pad_y) + if((pooled_h - 1) * pool_stride_y >= h_in + pad_y) { - --pooled_width; + --pooled_h; } if(type == PoolingType::MAX) { for(int r = 0; r < upper_dims; ++r) { - for(int i = 0; i < pooled_height; ++i) + for(int h = 0; h < pooled_h; ++h) { - for(int k = 0; k < pooled_width; ++k) + for(int w = 0; w < pooled_w; ++w) { - int hstart = i * pool_stride_x - pad_x; - int wstart = k * pool_stride_y - pad_y; - int hend = std::min(hstart + pool_size, rows_in); - int wend = std::min(wstart + pool_size, cols_in); - hstart = std::max(hstart, 0); + int wstart = w * pool_stride_x - pad_x; + int hstart = h * pool_stride_y - pad_y; + int wend = std::min(wstart + pool_size, w_in); + int hend = std::min(hstart + pool_size, h_in); wstart = std::max(wstart, 0); + hstart = std::max(hstart, 0); T max_val = std::numeric_limits<T>::lowest(); for(int y = hstart; y < hend; ++y) { for(int x = wstart; x < wend; ++x) { - T val = in[r * cols_in * rows_in + y * cols_in + x]; + T val = in[r * h_in * w_in + y * w_in + x]; if(val > max_val) { max_val = val; @@ -1202,7 +1212,7 @@ void pooling_layer(const Tensor<T> &in, Tensor<T> &out, PoolingLayerInfo pool_in } } - out[r * rows_out * cols_out + i * pooled_width + k] = max_val; + out[r * h_out * w_out + h * pooled_w + w] = max_val; } } } @@ -1211,32 +1221,30 @@ void pooling_layer(const Tensor<T> &in, Tensor<T> &out, PoolingLayerInfo pool_in { for(int r = 0; r < upper_dims; ++r) { - for(int i = 0; i < pooled_height; ++i) + for(int h = 0; h < pooled_h; ++h) { - for(int k = 0; k < pooled_width; ++k) + for(int w = 0; w < pooled_w; ++w) { - T avg_val = 0; - - int hstart = i * pool_stride_x - pad_x; - int wstart = k * pool_stride_y - pad_y; - int hend = std::min(hstart + pool_size, cols_in + pad_x); - int wend = std::min(wstart + pool_size, rows_in + pad_y); - int pool = (hend - hstart) * (wend - wstart); - hstart = std::max(hstart, 0); - wstart = std::max(wstart, 0); - hend = std::min(hend, rows_in); - wend = std::min(wend, cols_in); - + T avg_val = 0; + int wstart = w * pool_stride_x - pad_x; + int hstart = h * pool_stride_y - pad_y; + int wend = std::min(wstart + pool_size, w_in + pad_x); + int hend = std::min(hstart + pool_size, h_in + pad_y); + int pool = (hend - hstart) * (wend - wstart); + wstart = std::max(wstart, 0); + hstart = std::max(hstart, 0); + wend = std::min(wend, w_in); + hend = std::min(hend, h_in); if(std::is_floating_point<T>::value) { for(int y = hstart; y < hend; ++y) { for(int x = wstart; x < wend; ++x) { - avg_val += in[r * cols_in * rows_in + y * cols_in + x]; + avg_val += in[r * h_in * w_in + y * w_in + x]; } } - out[r * rows_out * cols_out + i * pooled_width + k] = avg_val / pool; + out[r * h_out * w_out + h * pooled_w + w] = avg_val / pool; } else { @@ -1247,10 +1255,10 @@ void pooling_layer(const Tensor<T> &in, Tensor<T> &out, PoolingLayerInfo pool_in { for(int x = wstart; x < wend; ++x) { - avg_val = sqadd_qs8(avg_val, in[r * cols_in * rows_in + y * cols_in + x]); + avg_val = sqadd_qs8(avg_val, in[r * h_in * w_in + y * w_in + x]); } } - out[r * rows_out * cols_out + i * pooled_width + k] = sqmul_qs8(avg_val, (scale_values_q8[pool] >> (7 - fixed_point_position)), fixed_point_position); + out[r * h_out * w_out + h * pooled_w + w] = sqmul_qs8(avg_val, (scale_values_q8[pool] >> (7 - fixed_point_position)), fixed_point_position); } } } |