From 3e36369a5511c3028c30fc820752dc1248bddf5c Mon Sep 17 00:00:00 2001 From: SiCong Li Date: Tue, 4 Jul 2017 15:02:10 +0100 Subject: COMPMID-358 Implement OpenCL ROI Pooling * Implement OpenCL ROI Pooling * Add CLROIPoolingLayer benchmarks Change-Id: I8786d01d551850a1b4d599a48fabe3925e0a27d0 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/79833 Reviewed-by: Anthony Barbier Tested-by: Kaizen --- arm_compute/core/CL/ICLArray.h | 1 + arm_compute/core/CL/ICLKernel.h | 78 ++++++++++ .../core/CL/kernels/CLROIPoolingLayerKernel.h | 76 +++++++++ .../core/NEON/kernels/NEROIPoolingLayerKernel.h | 7 +- arm_compute/runtime/CL/CLArray.h | 1 + .../runtime/CL/functions/CLROIPoolingLayer.h | 60 ++++++++ .../runtime/NEON/functions/NEROIPoolingLayer.h | 7 +- src/core/CL/CLKernelLibrary.cpp | 5 + src/core/CL/ICLKernel.cpp | 11 +- src/core/CL/cl_kernels/roi_pooling_layer.cl | 169 +++++++++++++++++++++ src/core/CL/kernels/CLROIPoolingLayerKernel.cpp | 121 +++++++++++++++ src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp | 2 +- src/runtime/CL/functions/CLROIPoolingLayer.cpp | 39 +++++ tests/CL/CLArrayAccessor.h | 83 ++++++++++ tests/IArrayAccessor.h | 60 ++++++++ tests/NEON/ArrayAccessor.h | 73 +++++++++ tests/NEON/Helper.h | 12 -- tests/Utils.h | 55 +++++++ tests/benchmark_new/CL/ROIPoolingLayer.cpp | 53 +++++++ tests/benchmark_new/NEON/ROIPoolingLayer.cpp | 53 +++++++ tests/datasets_new/ROIPoolingLayerDataset.h | 127 ++++++++++++++++ tests/fixtures_new/ROIPoolingLayerFixture.h | 95 ++++++++++++ tests/validation/CL/CMakeLists.txt | 9 ++ tests/validation/CL/ROIPoolingLayer.cpp | 112 ++++++++++++++ tests/validation/Helpers.cpp | 72 --------- tests/validation/Helpers.h | 13 +- tests/validation/NEON/Gaussian5x5.cpp | 2 - tests/validation/NEON/ROIPoolingLayer.cpp | 12 +- 28 files changed, 1296 insertions(+), 112 deletions(-) create mode 100644 arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h create mode 100644 arm_compute/runtime/CL/functions/CLROIPoolingLayer.h create mode 100644 src/core/CL/cl_kernels/roi_pooling_layer.cl create mode 100644 src/core/CL/kernels/CLROIPoolingLayerKernel.cpp create mode 100644 src/runtime/CL/functions/CLROIPoolingLayer.cpp create mode 100644 tests/CL/CLArrayAccessor.h create mode 100644 tests/IArrayAccessor.h create mode 100644 tests/NEON/ArrayAccessor.h create mode 100644 tests/benchmark_new/CL/ROIPoolingLayer.cpp create mode 100644 tests/benchmark_new/NEON/ROIPoolingLayer.cpp create mode 100644 tests/datasets_new/ROIPoolingLayerDataset.h create mode 100644 tests/fixtures_new/ROIPoolingLayerFixture.h create mode 100644 tests/validation/CL/ROIPoolingLayer.cpp delete mode 100644 tests/validation/Helpers.cpp diff --git a/arm_compute/core/CL/ICLArray.h b/arm_compute/core/CL/ICLArray.h index 1b676ed5a3..e12695f206 100644 --- a/arm_compute/core/CL/ICLArray.h +++ b/arm_compute/core/CL/ICLArray.h @@ -107,6 +107,7 @@ private: using ICLKeyPointArray = ICLArray; using ICLCoordinates2DArray = ICLArray; using ICLDetectionWindowArray = ICLArray; +using ICLROIArray = ICLArray; using ICLSize2DArray = ICLArray; using ICLUInt8Array = ICLArray; using ICLUInt16Array = ICLArray; diff --git a/arm_compute/core/CL/ICLKernel.h b/arm_compute/core/CL/ICLKernel.h index cfbf760f1e..1334c54a6c 100644 --- a/arm_compute/core/CL/ICLKernel.h +++ b/arm_compute/core/CL/ICLKernel.h @@ -31,6 +31,8 @@ namespace arm_compute { +template +class ICLArray; class ICLTensor; class Window; @@ -45,6 +47,16 @@ public: * @return A reference to the OpenCL kernel of this object. */ cl::Kernel &kernel(); + /** Add the passed 1D array'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 array's arguments. Will be incremented by the number of kernel arguments set. + * @param[in] array Array to set as an argument of the object's kernel. + * @param[in] strides @ref Strides object containing stride of each dimension in bytes. + * @param[in] num_dimensions Number of dimensions of the @p array. + * @param[in] window Window the kernel will be executed on. + */ + template + void add_1D_array_argument(unsigned int &idx, const ICLArray *array, const Strides &strides, unsigned int num_dimensions, const Window &window); /** Add the passed 1D 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. @@ -73,6 +85,11 @@ public: * @param[in] window Window the kernel will be executed on. */ void add_4D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window); + /** Returns the number of arguments enqueued per 1D array object. + * + * @return The number of arguments enqueues per 1D array object. + */ + unsigned int num_arguments_per_1D_array() const; /** Returns the number of arguments enqueued per 1D tensor object. * * @return The number of arguments enqueues per 1D tensor object. @@ -142,6 +159,16 @@ public: GPUTarget get_target() const; private: + /** Add the passed array'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 array's arguments. Will be incremented by the number of kernel arguments set. + * @param[in] array Array to set as an argument of the object's kernel. + * @param[in] strides @ref Strides object containing stride of each dimension in bytes. + * @param[in] num_dimensions Number of dimensions of the @p array. + * @param[in] window Window the kernel will be executed on. + */ + template + void add_array_argument(unsigned int &idx, const ICLArray *array, const Strides &strides, unsigned int num_dimensions, const Window &window); /** Add the passed 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. @@ -150,6 +177,12 @@ private: */ template void add_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window); + /** Returns the number of arguments enqueued per array object. + * + * @return The number of arguments enqueued per array object. + */ + template + unsigned int num_arguments_per_array() const; /** Returns the number of arguments enqueued per tensor object. * * @return The number of arguments enqueued per tensor object. @@ -177,5 +210,50 @@ protected: * @note If any dimension of the lws is greater than the global workgroup size then no lws will be passed. */ void enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Window &window, const cl::NDRange &lws_hint = CLKernelLibrary::get().default_ndrange()); + +template +void ICLKernel::add_array_argument(unsigned &idx, const ICLArray *array, const Strides &strides, unsigned int num_dimensions, const Window &window) +{ + // Calculate offset to the start of the window + unsigned int offset_first_element = 0; + + for(unsigned int n = 0; n < num_dimensions; ++n) + { + offset_first_element += window[n].start() * strides[n]; + } + + unsigned int idx_start = idx; + _kernel.setArg(idx++, array->cl_buffer()); + + for(unsigned int dimension = 0; dimension < dimension_size; dimension++) + { + _kernel.setArg(idx++, strides[dimension]); + _kernel.setArg(idx++, strides[dimension] * window[dimension].step()); + } + + _kernel.setArg(idx++, offset_first_element); + + ARM_COMPUTE_ERROR_ON_MSG(idx_start + num_arguments_per_array() != idx, + "add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_array()); + ARM_COMPUTE_UNUSED(idx_start); +} + +template +void ICLKernel::add_1D_array_argument(unsigned int &idx, const ICLArray *array, const Strides &strides, unsigned int num_dimensions, const Window &window) +{ + add_array_argument(idx, array, strides, num_dimensions, window); +} + +template +unsigned int ICLKernel::num_arguments_per_array() const +{ + return num_arguments_per_tensor(); +} + +template +unsigned int ICLKernel::num_arguments_per_tensor() const +{ + return 2 + 2 * dimension_size; +} } #endif /*__ARM_COMPUTE_ICLKERNEL_H__ */ diff --git a/arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h b/arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h new file mode 100644 index 0000000000..51aae30561 --- /dev/null +++ b/arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h @@ -0,0 +1,76 @@ +/* + * 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_CLROIPOOLINGLAYERKERNEL_H__ +#define __ARM_COMPUTE_CLROIPOOLINGLAYERKERNEL_H__ + +#include "arm_compute/core/CL/ICLKernel.h" + +#include "arm_compute/core/CL/ICLArray.h" + +namespace arm_compute +{ +class ICLTensor; + +/** Interface for the ROI pooling layer kernel */ +class CLROIPoolingLayerKernel : public ICLKernel +{ +public: + /** Default constructor */ + CLROIPoolingLayerKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLROIPoolingLayerKernel(const CLROIPoolingLayerKernel &) = delete; + /** Prevent instances of this class from being copied (As this class contains pointers) */ + CLROIPoolingLayerKernel &operator=(const CLROIPoolingLayerKernel &) = delete; + /** Allow instances of this class to be moved */ + CLROIPoolingLayerKernel(CLROIPoolingLayerKernel &&) = default; + /** Allow instances of this class to be moved */ + CLROIPoolingLayerKernel &operator=(CLROIPoolingLayerKernel &&) = default; + /** Default destructor */ + ~CLROIPoolingLayerKernel() = default; + + /** Set the input and output tensors. + * + * @param[in] input Source tensor. Data types supported: F16/F32. + * @param[in] rois Array containing @ref ROI. + * @param[out] output Destination tensor. Data types supported: Same as @p input. + * @param[in] pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo. + * + * @note The x and y dimensions of @p output tensor must be the same as @p pool_info 's pooled + * width and pooled height. + * @note The z dimensions of @p output tensor and @p input tensor must be the same. + * @note The fourth dimension of @p output tensor must be the same as the number of elements in @p rois array. + */ + void configure(const ICLTensor *input, const ICLROIArray *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info); + + // Inherited methods overridden: + void run(const Window &window, cl::CommandQueue &queue) override; + +private: + const ICLTensor *_input; + const ICLROIArray *_rois; + ICLTensor *_output; + ROIPoolingLayerInfo _pool_info; +}; +} +#endif /*__ARM_COMPUTE_CLROIPOOLINGLAYERKERNEL_H__ */ diff --git a/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h b/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h index 3a2f761370..40f79acc79 100644 --- a/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h +++ b/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h @@ -52,9 +52,14 @@ public: /** Set the input and output tensors. * * @param[in] input Source tensor. Data types supported: F32. - * @param[in] rois Array containing the regions of interest. + * @param[in] rois Array containing @ref ROI. * @param[out] output Destination tensor. Data types supported: Same as @p input. * @param[in] pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo. + * + * @note The x and y dimensions of @p output tensor must be the same as that specified by @p pool_info 's pooled + * width and pooled height. + * @note The z dimensions of @p output tensor and @p input tensor must be the same. + * @note The fourth dimension of @p output tensor must be the same as the number of elements in @p rois array. */ void configure(const ITensor *input, const IROIArray *rois, ITensor *output, const ROIPoolingLayerInfo &pool_info); diff --git a/arm_compute/runtime/CL/CLArray.h b/arm_compute/runtime/CL/CLArray.h index f4c2ef06d9..3dc7f19bc7 100644 --- a/arm_compute/runtime/CL/CLArray.h +++ b/arm_compute/runtime/CL/CLArray.h @@ -97,6 +97,7 @@ private: using CLKeyPointArray = CLArray; using CLCoordinates2DArray = CLArray; using CLDetectionWindowArray = CLArray; +using CLROIArray = CLArray; using CLSize2DArray = CLArray; using CLUInt8Array = CLArray; using CLUInt16Array = CLArray; diff --git a/arm_compute/runtime/CL/functions/CLROIPoolingLayer.h b/arm_compute/runtime/CL/functions/CLROIPoolingLayer.h new file mode 100644 index 0000000000..f089375e51 --- /dev/null +++ b/arm_compute/runtime/CL/functions/CLROIPoolingLayer.h @@ -0,0 +1,60 @@ +/* + * 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_CLROIPOOLINGLAYER_H__ +#define __ARM_COMPUTE_CLROIPOOLINGLAYER_H__ + +#include "arm_compute/runtime/CL/ICLSimpleFunction.h" + +#include "arm_compute/core/CL/ICLArray.h" +#include "arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h" + +namespace arm_compute +{ +class ICLTensor; + +/** Basic function to run @ref CLROIPoolingLayerKernel. + * + * This function calls the following OpenCL kernels: + * -# @ref CLROIPoolingLayerKernel + * + */ +class CLROIPoolingLayer : public ICLSimpleFunction +{ +public: + /** Set the input and output tensors. + * + * @param[in] input Source tensor. Data types supported: F16/F32. + * @param[in] rois Array containing @ref ROI. + * @param[out] output Destination tensor. Data types supported: Same as @p input. + * @param[in] pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo. + * + * @note The x and y dimensions of @p output tensor must be the same as @p pool_info 's pooled + * width and pooled height. + * @note The z dimensions of @p output tensor and @p input tensor must be the same. + * @note The fourth dimension of @p output tensor must be the same as the number of elements in @p rois array. + */ + void configure(const ICLTensor *input, const ICLROIArray *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info); +}; +} +#endif /* __ARM_COMPUTE_CLROIPOOLINGLAYER_H__ */ diff --git a/arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h b/arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h index 04b5c35150..5adc1110d5 100644 --- a/arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h +++ b/arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h @@ -47,9 +47,14 @@ public: /** Set the input and output tensors. * * @param[in] input Source tensor. Data types supported: F32. - * @param[in] rois Array containing the regions of interest. + * @param[in] rois Array containing @ref ROI. * @param[out] output Destination tensor. Data types supported: Same as @p input. * @param[in] pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo. + * + * @note The x and y dimensions of @p output tensor must be the same as that specified by @p pool_info 's pooled + * width and pooled height. + * @note The z dimensions of @p output tensor and @p input tensor must be the same. + * @note The fourth dimension of @p output tensor must be the same as the number of elements in @p rois array. */ void configure(const ITensor *input, const IROIArray *rois, ITensor *output, const ROIPoolingLayerInfo &pool_info); diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 2589bd12b5..ce2cfef67a 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -246,6 +246,7 @@ const std::map CLKernelLibrary::_kernel_program_map = { "RGBA8888_to_NV12_bt709", "color_convert.cl" }, { "RGBA8888_to_RGB888_bt709", "color_convert.cl" }, { "RGBA8888_to_YUV444_bt709", "color_convert.cl" }, + { "roi_pooling_layer", "roi_pooling_layer.cl" }, { "scale_nearest_neighbour", "scale.cl" }, { "scale_bilinear", "scale.cl" }, { "scharr3x3", "scharr_filter.cl" }, @@ -475,6 +476,10 @@ const std::map CLKernelLibrary::_program_source_map = { "remap.cl", #include "./cl_kernels/remap.clembed" + }, + { + "roi_pooling_layer.cl", +#include "./cl_kernels/roi_pooling_layer.clembed" }, { "scale.cl", diff --git a/src/core/CL/ICLKernel.cpp b/src/core/CL/ICLKernel.cpp index 12af8c68c1..7a95374bbf 100644 --- a/src/core/CL/ICLKernel.cpp +++ b/src/core/CL/ICLKernel.cpp @@ -69,12 +69,6 @@ cl::Kernel &ICLKernel::kernel() return _kernel; } -template -unsigned int ICLKernel::num_arguments_per_tensor() const -{ - return 2 + 2 * dimension_size; -} - template void ICLKernel::add_tensor_argument(unsigned &idx, const ICLTensor *tensor, const Window &window) { @@ -127,6 +121,11 @@ void ICLKernel::add_4D_tensor_argument(unsigned int &idx, const ICLTensor *tenso add_tensor_argument<4>(idx, tensor, window); } +unsigned int ICLKernel::num_arguments_per_1D_array() const +{ + return num_arguments_per_array<1>(); +} + unsigned int ICLKernel::num_arguments_per_1D_tensor() const { return num_arguments_per_tensor<1>(); diff --git a/src/core/CL/cl_kernels/roi_pooling_layer.cl b/src/core/CL/cl_kernels/roi_pooling_layer.cl new file mode 100644 index 0000000000..35a9c0a21f --- /dev/null +++ b/src/core/CL/cl_kernels/roi_pooling_layer.cl @@ -0,0 +1,169 @@ +/* + * 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 "helpers.h" + +#if DATA_SIZE == 32 +#define VEC_SIZE 4 +#define VEC_MAX vec4_max +#elif DATA_SIZE == 16 +#define VEC_SIZE 8 +#define VEC_MAX vec8_max +#else /* DATA_SIZE not equals 32 or 16 */ +#error "Unsupported data size" +#endif /* DATA_SIZE == 32 */ + +inline DATA_TYPE vec4_max(VEC_DATA_TYPE(DATA_TYPE, 4) vec) +{ + VEC_DATA_TYPE(DATA_TYPE, 2) + temp = fmax(vec.lo, vec.hi); + return fmax(temp.x, temp.y); +} + +inline DATA_TYPE vec8_max(VEC_DATA_TYPE(DATA_TYPE, 8) vec) +{ + VEC_DATA_TYPE(DATA_TYPE, 4) + temp = fmax(vec.lo, vec.hi); + return vec4_max(temp); +} + +/** Performs a roi pooling on a single output pixel. + * + * @param[in] input Pointer to input Tensor3D struct. + * @param[in] region_start_x Start x index projected onto the input tensor. + * @param[in] region_end_x End x index projected onto the input tensor. + * @param[in] region_start_y Start y index projected onto the input tensor. + * @param[in] region_end_y End y index projected onto the input tensor. + * @param[in] pz z index of the input tensor. + * + * @return A max pooled value from the region specified in the input tensor. + */ +inline DATA_TYPE roi_pool_1x1(const Tensor3D *input, int region_start_x, int region_end_x, int region_start_y, int region_end_y, int pz) +{ + // Iterate through the pooling region + if((region_end_x <= region_start_x) || (region_end_y <= region_start_y)) + { + return (DATA_TYPE)0; + } + else + { + int num_iter = (int)((region_end_x - region_start_x) / VEC_SIZE); + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + curr_max = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))(-FLT_MAX); + for(int j = region_start_y; j < region_end_y; ++j) + { + int i = region_start_x; + for(; i < region_start_x + num_iter * VEC_SIZE; i += VEC_SIZE) + { + VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) + val = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)tensor3D_offset(input, i, j, pz)); + curr_max = fmax(val, curr_max); + } + for(; i < region_end_x; ++i) + { + DATA_TYPE val = *(__global DATA_TYPE *)tensor3D_offset(input, i, j, pz); + curr_max = fmax(curr_max, val); + } + } + return (DATA_TYPE)VEC_MAX(curr_max); + } +} + +/** Performs a roi pooling function. + * + * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16, F32; + * @note Datasize must be passed using -DDATA_SIZE e.g. -DDATA_SIZE=32; + * @note Input dimensions must be passed using -DMAX_DIM_X, -DMAX_DIM_Y and -DMAX_DIM_Z; + * @note Pooled region dimensions must be passed using -DPOOLED_DIM_X and -DPOOLED_DIM_Y; + * @note Spatial scale must be passed using -DSPATIAL_SCALE; + * + * @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 pooled region of the source image as specifed by ROI + * @param[in] rois_ptr Pointer to the rois array. Layout: {x, y, width, height, batch_indx} + * @param[in] rois_stride_x Stride of the rois array in X dimension (in bytes) + * @param[in] rois_step_x rois_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] rois_offset_first_element_in_bytes The offset of the first element in the rois array + * @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 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_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[in] input_stride_w Stride of the source image in W dimension (in bytes) + * @param[in] output_stride_w Stride of the destination image in W dimension (in bytes) + */ +__kernel void roi_pooling_layer( + TENSOR3D_DECLARATION(input), + VECTOR_DECLARATION(rois), + TENSOR3D_DECLARATION(output), + unsigned int input_stride_w, unsigned int output_stride_w) +{ + // Get pixels pointer + Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input); + Vector rois = CONVERT_TO_VECTOR_STRUCT_NO_STEP(rois); + Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output); + + const int px = get_global_id(0); + const int py = get_global_id(1); + const int pw = get_global_id(2); + + // Load roi parameters + // roi is laid out as follows: + // { x, y, width, height, batch_index } + const ushort8 roi = vload8(0, (__global ushort *)vector_offset(&rois, pw)); + const int2 roi_anchor = convert_int2_sat(round(convert_float2(roi.s01) * (float)SPATIAL_SCALE)); + const int2 roi_dims = convert_int2_sat(fmax(round(convert_float2(roi.s23) * (float)SPATIAL_SCALE), 1.f)); + + // Determine pooled region in input image to pooled region in output image ratio + const float2 pool_region_ratio = convert_float2(roi_dims) / (float2)(POOLED_DIM_X, POOLED_DIM_Y); + + // Calculate pooled region start and end + const float2 spatial_indx = (float2)(px, py); + const int2 max_spatial_dims = (int2)(MAX_DIM_X, MAX_DIM_Y); + int2 region_start = convert_int2_sat(floor(spatial_indx * pool_region_ratio)) + roi_anchor; + int2 region_end = convert_int2_sat(ceil((spatial_indx + 1) * pool_region_ratio)) + roi_anchor; + + region_start = clamp(region_start, 0, max_spatial_dims); + region_end = clamp(region_end, 0, max_spatial_dims); + + // Move input and output pointer across the fourth dimension + input.ptr += roi.s4 * input_stride_w; + output.ptr += pw * output_stride_w; + + for(int pz = 0; pz < MAX_DIM_Z; ++pz) + { + *(__global DATA_TYPE *)tensor3D_offset(&output, px, py, pz) = (__global DATA_TYPE)roi_pool_1x1(&input, + region_start.x, + region_end.x, + region_start.y, + region_end.y, pz); + } +} diff --git a/src/core/CL/kernels/CLROIPoolingLayerKernel.cpp b/src/core/CL/kernels/CLROIPoolingLayerKernel.cpp new file mode 100644 index 0000000000..4e000c61b1 --- /dev/null +++ b/src/core/CL/kernels/CLROIPoolingLayerKernel.cpp @@ -0,0 +1,121 @@ +/* + * 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 "arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h" + +#include "arm_compute/core/AccessWindowStatic.h" +#include "arm_compute/core/CL/CLHelpers.h" +#include "arm_compute/core/CL/CLKernelLibrary.h" +#include "arm_compute/core/CL/ICLArray.h" +#include "arm_compute/core/CL/ICLTensor.h" +#include "arm_compute/core/CL/OpenCL.h" +#include "arm_compute/core/Helpers.h" +#include "arm_compute/core/TensorInfo.h" +#include "arm_compute/core/Utils.h" +#include "arm_compute/core/Validate.h" +#include "arm_compute/core/Window.h" + +#include +#include +#include + +using namespace arm_compute; + +CLROIPoolingLayerKernel::CLROIPoolingLayerKernel() + : _input(nullptr), _rois(nullptr), _output(nullptr), _pool_info(0, 0, 0.f) +{ +} + +void CLROIPoolingLayerKernel::configure(const ICLTensor *input, const ICLROIArray *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info) +{ + ARM_COMPUTE_ERROR_ON_NULLPTR(input, rois, output); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32); + ARM_COMPUTE_ERROR_ON((pool_info.pooled_width() == 0) || (pool_info.pooled_height() == 0)); + ARM_COMPUTE_ERROR_ON(rois->num_values() == 0); + + // Output auto inizialitation if not yet initialized + TensorShape output_shape(pool_info.pooled_width(), pool_info.pooled_height(), input->info()->dimension(2), rois->num_values()); + auto_init_if_empty(*output->info(), output_shape, 1, input->info()->data_type(), input->info()->fixed_point_position()); + + ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); + ARM_COMPUTE_ERROR_ON((output->info()->dimension(0) != pool_info.pooled_width()) || (output->info()->dimension(1) != pool_info.pooled_height())); + ARM_COMPUTE_ERROR_ON(input->info()->dimension(2) != output->info()->dimension(2)); + ARM_COMPUTE_ERROR_ON(rois->num_values() != output->info()->dimension(3)); + + // Set instance variables + _input = input; + _rois = rois; + _output = output; + _pool_info = pool_info; + + // Set build options + std::set build_opts; + build_opts.emplace(("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type()))); + build_opts.emplace(("-DDATA_SIZE=" + get_data_size_from_data_type(input->info()->data_type()))); + build_opts.emplace(("-DMAX_DIM_X=" + support::cpp11::to_string(_input->info()->dimension(Window::DimX)))); + build_opts.emplace(("-DMAX_DIM_Y=" + support::cpp11::to_string(_input->info()->dimension(Window::DimY)))); + build_opts.emplace(("-DMAX_DIM_Z=" + support::cpp11::to_string(_input->info()->dimension(Window::DimZ)))); + build_opts.emplace(("-DPOOLED_DIM_X=" + support::cpp11::to_string(pool_info.pooled_width()))); + build_opts.emplace(("-DPOOLED_DIM_Y=" + support::cpp11::to_string(pool_info.pooled_height()))); + build_opts.emplace(("-DSPATIAL_SCALE=" + support::cpp11::to_string(pool_info.spatial_scale()))); + + // Create kernel + std::string kernel_name = "roi_pooling_layer"; + _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts)); + + // Set static kernel arguments + unsigned int idx = 2 * num_arguments_per_3D_tensor() + num_arguments_per_1D_array(); + add_argument(idx, _input->info()->strides_in_bytes()[3]); + add_argument(idx, _output->info()->strides_in_bytes()[3]); + + // Configure kernel window + const unsigned int num_elems_processed_per_iteration = 1; + Window window = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration)); + AccessWindowStatic input_access(input->info(), + input->info()->valid_region().start(0), + input->info()->valid_region().start(1), + input->info()->valid_region().end(0), + input->info()->valid_region().end(1)); + AccessWindowStatic output_access(output->info(), 0, 0, pool_info.pooled_width(), pool_info.pooled_height()); + + update_window_and_padding(window, input_access, output_access); + output_access.set_valid_region(window, ValidRegion(Coordinates(), output->info()->tensor_shape())); + ICLKernel::configure(window); +} + +void CLROIPoolingLayerKernel::run(const Window &window, cl::CommandQueue &queue) +{ + ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window); + + Window slice(window); + // Parallelize spatially and across the fourth dimension of the output tensor (also across ROIArray) + slice.set(Window::DimZ, window[3]); + + // Set arguments + unsigned int idx = 0; + add_3D_tensor_argument(idx, _input, slice); + add_1D_array_argument(idx, _rois, Strides(sizeof(ROI)), 1U, slice); + add_3D_tensor_argument(idx, _output, slice); + enqueue(queue, *this, slice); +} diff --git a/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp b/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp index ec2b6a5079..fd6210021c 100644 --- a/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp +++ b/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp @@ -50,7 +50,7 @@ void NEROIPoolingLayerKernel::configure(const ITensor *input, const IROIArray *r ARM_COMPUTE_ERROR_ON(rois->num_values() == 0); // Output auto inizialitation if not yet initialized - TensorShape output_shape(pool_info.pooled_width(), pool_info.pooled_height(), rois->num_values()); + TensorShape output_shape(pool_info.pooled_width(), pool_info.pooled_height(), input->info()->dimension(2), rois->num_values()); auto_init_if_empty(*output->info(), output_shape, 1, input->info()->data_type(), input->info()->fixed_point_position()); ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); diff --git a/src/runtime/CL/functions/CLROIPoolingLayer.cpp b/src/runtime/CL/functions/CLROIPoolingLayer.cpp new file mode 100644 index 0000000000..0f480eeac9 --- /dev/null +++ b/src/runtime/CL/functions/CLROIPoolingLayer.cpp @@ -0,0 +1,39 @@ +/* + * 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 "arm_compute/runtime/CL/functions/CLROIPoolingLayer.h" + +#include "arm_compute/core/CL/ICLArray.h" + +#include "arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h" +#include "support/ToolchainSupport.h" + +using namespace arm_compute; + +void CLROIPoolingLayer::configure(const ICLTensor *input, const ICLROIArray *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info) +{ + // Configure ROI pooling kernel + auto k = arm_compute::support::cpp14::make_unique(); + k->configure(input, rois, output, pool_info); + _kernel = std::move(k); +} diff --git a/tests/CL/CLArrayAccessor.h b/tests/CL/CLArrayAccessor.h new file mode 100644 index 0000000000..a38907dd78 --- /dev/null +++ b/tests/CL/CLArrayAccessor.h @@ -0,0 +1,83 @@ +/* + * 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_CLARRAYACCESSOR_H__ +#define __ARM_COMPUTE_TEST_CLARRAYACCESSOR_H__ + +#include "arm_compute/runtime/CL/CLArray.h" +#include "tests/IArrayAccessor.h" + +namespace arm_compute +{ +namespace test +{ +/** Accessor implementation for @ref CLArray objects. */ +template +class CLArrayAccessor : public IArrayAccessor +{ +public: + /** Create an accessor for the given @p array. + * + * @param[in, out] array To be accessed array. + * + * @note The CL memory is mapped by the constructor. + * + */ + CLArrayAccessor(CLArray &array) + : _array{ array } + { + _array.map(); + } + + CLArrayAccessor(const CLArrayAccessor &) = delete; + CLArrayAccessor &operator=(const CLArrayAccessor &) = delete; + CLArrayAccessor(CLArrayAccessor &&) = default; + CLArrayAccessor &operator=(CLArrayAccessor &&) = default; + + /** Destructor that unmaps the CL memory. */ + ~CLArrayAccessor() + { + _array.unmap(); + } + + int num_values() const override + { + return _array.num_values(); + } + + T *buffer() override + { + return _array.buffer(); + } + + void resize(size_t num) override + { + _array.resize(num); + } + +private: + CLArray &_array; +}; +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_CLARRAYACCESSOR_H__ */ diff --git a/tests/IArrayAccessor.h b/tests/IArrayAccessor.h new file mode 100644 index 0000000000..73f8829ed2 --- /dev/null +++ b/tests/IArrayAccessor.h @@ -0,0 +1,60 @@ +/* + * 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_IARRAYACCESSOR_H__ +#define __ARM_COMPUTE_TEST_IARRAYACCESSOR_H__ + +namespace arm_compute +{ +namespace test +{ +/** Common interface to provide information and access to array like + * structures. + */ +template +class IArrayAccessor +{ +public: + using value_type = T; + + /** Virtual destructor. */ + virtual ~IArrayAccessor() = default; + + /** Number of elements of the tensor. */ + virtual int num_values() const = 0; + + /** Access to the buffer. + * + * @return A pointer to the first element in the buffer. + */ + virtual T *buffer() = 0; + + /** Resize array. + * + * @param[in] num The new array size in number of elements + */ + virtual void resize(size_t num) = 0; +}; +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_IARRAYACCESSOR_H__ */ diff --git a/tests/NEON/ArrayAccessor.h b/tests/NEON/ArrayAccessor.h new file mode 100644 index 0000000000..e56e6538ba --- /dev/null +++ b/tests/NEON/ArrayAccessor.h @@ -0,0 +1,73 @@ +/* + * 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_ARRAYACCESSOR_H__ +#define __ARM_COMPUTE_TEST_ARRAYACCESSOR_H__ + +#include "arm_compute/runtime/Array.h" +#include "tests/IArrayAccessor.h" + +namespace arm_compute +{ +namespace test +{ +/** ArrayAccessor implementation for @ref Array objects. */ +template +class ArrayAccessor : public IArrayAccessor +{ +public: + /** Create an accessor for the given @p array. + * + * @param[in, out] array To be accessed array. + */ + ArrayAccessor(Array &array) + : _array{ array } + { + } + + ArrayAccessor(const ArrayAccessor &) = delete; + ArrayAccessor &operator=(const ArrayAccessor &) = delete; + ArrayAccessor(ArrayAccessor &&) = default; + ArrayAccessor &operator=(ArrayAccessor &&) = default; + + int num_values() const override + { + return _array.num_values(); + } + + T *buffer() override + { + return _array.buffer(); + } + + void resize(size_t num) override + { + _array.resize(num); + } + +private: + Array &_array; +}; +} // namespace test +} // namespace arm_compute +#endif /* __ARM_COMPUTE_TEST_ARRAYACCESSOR_H__ */ diff --git a/tests/NEON/Helper.h b/tests/NEON/Helper.h index 9e60a48ecd..4efab17fca 100644 --- a/tests/NEON/Helper.h +++ b/tests/NEON/Helper.h @@ -35,18 +35,6 @@ namespace arm_compute { namespace test { -template -Array create_array(const std::vector &v) -{ - const size_t v_size = v.size(); - Array array(v_size); - - array.resize(v_size); - std::copy(v.begin(), v.end(), array.buffer()); - - return array; -} - template void fill_tensors(D &&dist, std::initializer_list seeds, T &&tensor, Ts &&... other_tensors) { diff --git a/tests/Utils.h b/tests/Utils.h index 93831aa30b..f1725d7d74 100644 --- a/tests/Utils.h +++ b/tests/Utils.h @@ -37,9 +37,11 @@ #include #include #include +#include #include #include #include +#include namespace arm_compute { @@ -421,6 +423,59 @@ inline T create_tensor(const TensorShape &shape, DataType data_type, int num_cha return tensor; } + +/** Create a vector of random ROIs. + * + * @param[in] shape The shape of the input tensor. + * @param[in] pool_info The ROI pooling information. + * @param[in] num_rois The number of ROIs to be created. + * @param[in] seed The random seed to be used. + * + * @return A vector that contains the requested number of random ROIs + */ +inline std::vector generate_random_rois(const TensorShape &shape, const ROIPoolingLayerInfo &pool_info, unsigned int num_rois, std::random_device::result_type seed) +{ + ARM_COMPUTE_ERROR_ON((pool_info.pooled_width() < 4) || (pool_info.pooled_height() < 4)); + + std::vector rois; + std::mt19937 gen(seed); + const int pool_width = pool_info.pooled_width(); + const int pool_height = pool_info.pooled_height(); + const float roi_scale = pool_info.spatial_scale(); + + // Calculate distribution bounds + const auto scaled_width = static_cast((shape.x() / roi_scale) / pool_width); + const auto scaled_height = static_cast((shape.y() / roi_scale) / pool_height); + const auto min_width = static_cast(pool_width / roi_scale); + const auto min_height = static_cast(pool_height / roi_scale); + + // Create distributions + std::uniform_int_distribution dist_batch(0, shape[3] - 1); + std::uniform_int_distribution dist_x(0, scaled_width); + std::uniform_int_distribution dist_y(0, scaled_height); + std::uniform_int_distribution dist_w(min_width, std::max(min_width, (pool_width - 2) * scaled_width)); + std::uniform_int_distribution dist_h(min_height, std::max(min_height, (pool_height - 2) * scaled_height)); + + for(unsigned int r = 0; r < num_rois; ++r) + { + ROI roi; + roi.batch_idx = dist_batch(gen); + roi.rect.x = dist_x(gen); + roi.rect.y = dist_y(gen); + roi.rect.width = dist_w(gen); + roi.rect.height = dist_h(gen); + rois.push_back(roi); + } + + return rois; +} + +template +inline void fill_array(ArrayAccessor_T &&array, const std::vector &v) +{ + array.resize(v.size()); + std::memcpy(array.buffer(), v.data(), v.size() * sizeof(T)); +} } // namespace test } // namespace arm_compute #endif /* __ARM_COMPUTE_TEST_UTILS_H__ */ diff --git a/tests/benchmark_new/CL/ROIPoolingLayer.cpp b/tests/benchmark_new/CL/ROIPoolingLayer.cpp new file mode 100644 index 0000000000..260f4d7b0d --- /dev/null +++ b/tests/benchmark_new/CL/ROIPoolingLayer.cpp @@ -0,0 +1,53 @@ +/* + * 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 "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/CL/CLArray.h" +#include "arm_compute/runtime/CL/CLTensor.h" +#include "arm_compute/runtime/CL/CLTensorAllocator.h" +#include "arm_compute/runtime/CL/functions/CLROIPoolingLayer.h" +#include "framework/Macros.h" +#include "framework/datasets/Datasets.h" +#include "tests/CL/CLAccessor.h" +#include "tests/CL/CLArrayAccessor.h" +#include "tests/TypePrinter.h" +#include "tests/datasets_new/ROIPoolingLayerDataset.h" +#include "tests/fixtures_new/ROIPoolingLayerFixture.h" + +namespace arm_compute +{ +namespace test +{ +using CLROIPoolingLayerFixture = ROIPoolingLayerFixture, CLArrayAccessor>; + +TEST_SUITE(CL) + +REGISTER_FIXTURE_DATA_TEST_CASE(SmallROIPoolingLayer, CLROIPoolingLayerFixture, framework::DatasetMode::ALL, + framework::dataset::combine(framework::dataset::combine(datasets::SmallROIPoolingLayerDataset(), + framework::dataset::make("DataType", { DataType::F16, DataType::F32 })), + framework::dataset::make("Batches", { 1, 4, 8 }))); + +TEST_SUITE_END() +} // namespace test +} // namespace arm_compute diff --git a/tests/benchmark_new/NEON/ROIPoolingLayer.cpp b/tests/benchmark_new/NEON/ROIPoolingLayer.cpp new file mode 100644 index 0000000000..e8c5ba3109 --- /dev/null +++ b/tests/benchmark_new/NEON/ROIPoolingLayer.cpp @@ -0,0 +1,53 @@ +/* + * 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 "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "arm_compute/runtime/Array.h" +#include "arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h" +#include "arm_compute/runtime/Tensor.h" +#include "arm_compute/runtime/TensorAllocator.h" +#include "framework/Macros.h" +#include "framework/datasets/Datasets.h" +#include "tests/NEON/Accessor.h" +#include "tests/NEON/ArrayAccessor.h" +#include "tests/TypePrinter.h" +#include "tests/datasets_new/ROIPoolingLayerDataset.h" +#include "tests/fixtures_new/ROIPoolingLayerFixture.h" + +namespace arm_compute +{ +namespace test +{ +using NEROIPoolingLayerFixture = ROIPoolingLayerFixture, ArrayAccessor>; + +TEST_SUITE(NEON) + +REGISTER_FIXTURE_DATA_TEST_CASE(SmallROIPoolingLayer, NEROIPoolingLayerFixture, framework::DatasetMode::ALL, + framework::dataset::combine(framework::dataset::combine(datasets::SmallROIPoolingLayerDataset(), + framework::dataset::make("DataType", { DataType::F32 })), + framework::dataset::make("Batches", { 1, 4, 8 }))); + +TEST_SUITE_END() +} // namespace test +} // namespace arm_compute diff --git a/tests/datasets_new/ROIPoolingLayerDataset.h b/tests/datasets_new/ROIPoolingLayerDataset.h new file mode 100644 index 0000000000..c68a33fe8f --- /dev/null +++ b/tests/datasets_new/ROIPoolingLayerDataset.h @@ -0,0 +1,127 @@ +/* + * 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_ROI_POOLING_LAYER_DATASET +#define ARM_COMPUTE_TEST_ROI_POOLING_LAYER_DATASET + +#include "tests/TypePrinter.h" + +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" + +namespace arm_compute +{ +namespace test +{ +namespace datasets +{ +class ROIPoolingLayerDataset +{ +public: + using type = std::tuple; + + struct iterator + { + iterator(std::vector::const_iterator tensor_shape_it, + std::vector::const_iterator infos_it, + std::vector::const_iterator num_rois_it) + : _tensor_shape_it{ std::move(tensor_shape_it) }, + _infos_it{ std::move(infos_it) }, + _num_rois_it{ std::move(num_rois_it) } + { + } + + std::string description() const + { + std::stringstream description; + description << "In=" << *_tensor_shape_it << ":"; + description << "Info=" << *_infos_it; + description << "NumROIS=" << *_num_rois_it; + return description.str(); + } + + ROIPoolingLayerDataset::type operator*() const + { + return std::make_tuple(*_tensor_shape_it, *_infos_it, *_num_rois_it); + } + + iterator &operator++() + { + ++_tensor_shape_it; + ++_infos_it; + ++_num_rois_it; + + return *this; + } + + private: + std::vector::const_iterator _tensor_shape_it; + std::vector::const_iterator _infos_it; + std::vector::const_iterator _num_rois_it; + }; + + iterator begin() const + { + return iterator(_tensor_shapes.begin(), _infos.begin(), _num_rois.begin()); + } + + int size() const + { + return std::min(std::min(_tensor_shapes.size(), _infos.size()), _num_rois.size()); + } + + void add_config(TensorShape tensor_shape, ROIPoolingLayerInfo info, unsigned int num_rois) + { + _tensor_shapes.emplace_back(std::move(tensor_shape)); + _infos.emplace_back(std::move(info)); + _num_rois.emplace_back(std::move(num_rois)); + } + +protected: + ROIPoolingLayerDataset() = default; + ROIPoolingLayerDataset(ROIPoolingLayerDataset &&) = default; + +private: + std::vector _tensor_shapes{}; + std::vector _infos{}; + std::vector _num_rois{}; +}; + +class SmallROIPoolingLayerDataset final : public ROIPoolingLayerDataset +{ +public: + SmallROIPoolingLayerDataset() + { + add_config(TensorShape(50U, 47U, 3U), ROIPoolingLayerInfo(7U, 7U, 1.f / 8.f), 40U); + add_config(TensorShape(50U, 47U, 10U), ROIPoolingLayerInfo(7U, 7U, 1.f / 8.f), 80U); + add_config(TensorShape(50U, 47U, 80U), ROIPoolingLayerInfo(7U, 7U, 1.f / 8.f), 80U); + add_config(TensorShape(50U, 47U, 3U), ROIPoolingLayerInfo(9U, 9U, 1.f / 8.f), 40U); + add_config(TensorShape(50U, 47U, 10U), ROIPoolingLayerInfo(9U, 9U, 1.f / 8.f), 80U); + add_config(TensorShape(50U, 47U, 80U), ROIPoolingLayerInfo(9U, 9U, 1.f / 8.f), 80U); + } +}; + +} // namespace datasets +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_ROI_POOLING_LAYER_DATASET */ diff --git a/tests/fixtures_new/ROIPoolingLayerFixture.h b/tests/fixtures_new/ROIPoolingLayerFixture.h new file mode 100644 index 0000000000..54dac51ee4 --- /dev/null +++ b/tests/fixtures_new/ROIPoolingLayerFixture.h @@ -0,0 +1,95 @@ +/* + * 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_ROIPOOLINGLAYERFIXTURE +#define ARM_COMPUTE_TEST_ROIPOOLINGLAYERFIXTURE + +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "framework/Fixture.h" +#include "tests/Globals.h" +#include "tests/Utils.h" + +#include + +namespace arm_compute +{ +namespace test +{ +/** Fixture that can be used for NEON and CL */ +template +class ROIPoolingLayerFixture : public framework::Fixture +{ +public: + template + void setup(TensorShape shape, const ROIPoolingLayerInfo pool_info, unsigned int num_rois, DataType data_type, int batches) + { + // Set batched in source and destination shapes + const unsigned int fixed_point_position = 4; + TensorShape shape_dst; + shape.set(shape.num_dimensions(), batches); + shape_dst.set(0, pool_info.pooled_width()); + shape_dst.set(1, pool_info.pooled_height()); + shape_dst.set(2, shape.z()); + shape_dst.set(3, num_rois); + + // Create tensors + src = create_tensor(shape, data_type, 1, fixed_point_position); + dst = create_tensor(shape_dst, data_type, 1, fixed_point_position); + + // Create random ROIs + std::vector rois = generate_random_rois(shape, pool_info, num_rois, 0U); + rois_array = arm_compute::support::cpp14::make_unique(num_rois); + fill_array(ArrayAccessor(*rois_array), rois); + + // Create and configure function + roi_pool.configure(&src, rois_array.get(), &dst, pool_info); + + // Allocate tensors + src.allocator()->allocate(); + dst.allocator()->allocate(); + + // Fill tensors + library->fill_tensor_uniform(Accessor(src), 0); + } + + void run() + { + roi_pool.run(); + } + + void teardown() + { + src.allocator()->free(); + dst.allocator()->free(); + } + +private: + TensorType src{}; + TensorType dst{}; + std::unique_ptr rois_array{}; + Function roi_pool{}; +}; +} // namespace test +} // namespace arm_compute +#endif /* ARM_COMPUTE_TEST_ROIPOOLINGLAYERFIXTURE */ diff --git a/tests/validation/CL/CMakeLists.txt b/tests/validation/CL/CMakeLists.txt index 67900b3edc..f4477f621f 100644 --- a/tests/validation/CL/CMakeLists.txt +++ b/tests/validation/CL/CMakeLists.txt @@ -27,15 +27,24 @@ set(arm_compute_test_validation_OPENCL_SOURCE_FILES ${CMAKE_SOURCE_DIR}/CL/CLAccessor.h ${CMAKE_CURRENT_SOURCE_DIR}/CLFixture.h ${CMAKE_CURRENT_SOURCE_DIR}/CLFixture.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/ActivationLayer.cpp ${CMAKE_CURRENT_SOURCE_DIR}/BitwiseAnd.cpp ${CMAKE_CURRENT_SOURCE_DIR}/Box3x3.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/ConvolutionLayer.cpp ${CMAKE_CURRENT_SOURCE_DIR}/DepthConvert.cpp ${CMAKE_CURRENT_SOURCE_DIR}/FillBorder.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/FixedPoint/FixedPoint_QS8.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/FullyConnectedLayer.cpp ${CMAKE_CURRENT_SOURCE_DIR}/Gaussian3x3.cpp ${CMAKE_CURRENT_SOURCE_DIR}/Gaussian5x5.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/GEMM.cpp ` ${CMAKE_CURRENT_SOURCE_DIR}/IntegralImage.cpp +` ${CMAKE_CURRENT_SOURCE_DIR}/NonLinearFilter.cpp +` ${CMAKE_CURRENT_SOURCE_DIR}/PoolingLayer.cpp +` ${CMAKE_CURRENT_SOURCE_DIR}/ROIPoolingLayer.cpp ${CMAKE_CURRENT_SOURCE_DIR}/Sobel3x3.cpp ${CMAKE_CURRENT_SOURCE_DIR}/Sobel5x5.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/SoftmaxLayer.cpp ${CMAKE_CURRENT_SOURCE_DIR}/Threshold.cpp ${CMAKE_CURRENT_SOURCE_DIR}/DirectConvolutionLayer.cpp ${CMAKE_CURRENT_SOURCE_DIR}/MeanStdDev.cpp diff --git a/tests/validation/CL/ROIPoolingLayer.cpp b/tests/validation/CL/ROIPoolingLayer.cpp new file mode 100644 index 0000000000..19d7903128 --- /dev/null +++ b/tests/validation/CL/ROIPoolingLayer.cpp @@ -0,0 +1,112 @@ +/* + * 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/CLArrayAccessor.h" +#include "TypePrinter.h" +#include "arm_compute/runtime/CL/CLArray.h" +#include "arm_compute/runtime/CL/functions/CLROIPoolingLayer.h" +#include "tests/Globals.h" +#include "tests/Utils.h" +#include "validation/Datasets.h" +#include "validation/Reference.h" +#include "validation/Validation.h" +#include "validation/ValidationUserConfiguration.h" + +#include +#include + +using namespace arm_compute; +using namespace arm_compute::test; +using namespace arm_compute::test::validation; + +namespace +{ +CLTensor compute_roi_pooling_layer(const TensorShape &shape, DataType dt, const std::vector &rois, ROIPoolingLayerInfo pool_info) +{ + TensorShape shape_dst; + shape_dst.set(0, pool_info.pooled_width()); + shape_dst.set(1, pool_info.pooled_height()); + shape_dst.set(2, shape.z()); + shape_dst.set(3, rois.size()); + + // Create tensors + CLTensor src = create_tensor(shape, dt); + CLTensor dst = create_tensor(shape_dst, dt); + + // Create ROI array + CLArray rois_array(rois.size()); + fill_array(CLArrayAccessor(rois_array), rois); + + // Create and configure function + CLROIPoolingLayer roi_pool; + roi_pool.configure(&src, &rois_array, &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 + roi_pool.run(); + + return dst; +} +} // namespace + +#ifndef DOXYGEN_SKIP_THIS +BOOST_AUTO_TEST_SUITE(CL) +BOOST_AUTO_TEST_SUITE(ROIPoolingLayer) + +BOOST_AUTO_TEST_SUITE(Float) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, boost::unit_test::data::make({ DataType::F16, DataType::F32 }) * boost::unit_test::data::make({ 10, 20, 40 }) * boost::unit_test::data::make({ 7, 9 }) * + boost::unit_test::data::make({ 1.f / 8.f, 1.f / 16.f }), + dt, num_rois, roi_pool_size, roi_scale) +{ + TensorShape shape(50U, 47U, 2U, 3U); + ROIPoolingLayerInfo pool_info(roi_pool_size, roi_pool_size, roi_scale); + + // Construct ROI vector + std::vector rois = generate_random_rois(shape, pool_info, num_rois, user_config.seed); + + // Compute function + CLTensor dst = compute_roi_pooling_layer(shape, dt, rois, pool_info); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_roi_pooling_layer(shape, dt, rois, pool_info); + + // Validate output + validate(CLAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE_END() +BOOST_AUTO_TEST_SUITE_END() +#endif /* DOXYGEN_SKIP_THIS */ diff --git a/tests/validation/Helpers.cpp b/tests/validation/Helpers.cpp deleted file mode 100644 index 747260ea99..0000000000 --- a/tests/validation/Helpers.cpp +++ /dev/null @@ -1,72 +0,0 @@ -/* - * 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 "validation/Helpers.h" - -using namespace arm_compute::test; - -namespace arm_compute -{ -namespace test -{ -namespace validation -{ -std::vector generate_random_rois(const TensorShape &shape, const ROIPoolingLayerInfo &pool_info, unsigned int num_rois, std::random_device::result_type seed) -{ - ARM_COMPUTE_ERROR_ON((pool_info.pooled_width() < 4) || (pool_info.pooled_height() < 4)); - - std::vector rois; - std::mt19937 gen(seed); - const int pool_width = pool_info.pooled_width(); - const int pool_height = pool_info.pooled_height(); - const float roi_scale = pool_info.spatial_scale(); - - // Calculate distribution bounds - const auto scaled_width = static_cast((shape.x() / roi_scale) / pool_width); - const auto scaled_height = static_cast((shape.y() / roi_scale) / pool_height); - const auto min_width = static_cast(pool_width / roi_scale); - const auto min_height = static_cast(pool_height / roi_scale); - - // Create distributions - std::uniform_int_distribution dist_batch(0, shape[3] - 1); - std::uniform_int_distribution dist_x(0, scaled_width); - std::uniform_int_distribution dist_y(0, scaled_height); - std::uniform_int_distribution dist_w(min_width, std::max(min_width, (pool_width - 2) * scaled_width)); - std::uniform_int_distribution dist_h(min_height, std::max(min_height, (pool_height - 2) * scaled_height)); - - for(unsigned int r = 0; r < num_rois; ++r) - { - ROI roi{}; - roi.batch_idx = dist_batch(gen); - roi.rect.x = dist_x(gen); - roi.rect.y = dist_y(gen); - roi.rect.width = dist_w(gen); - roi.rect.height = dist_h(gen); - rois.push_back(roi); - } - - return rois; -} -} // namespace validation -} // namespace test -} // namespace arm_compute diff --git a/tests/validation/Helpers.h b/tests/validation/Helpers.h index 09ffda8957..19a0c4105c 100644 --- a/tests/validation/Helpers.h +++ b/tests/validation/Helpers.h @@ -32,6 +32,7 @@ #include "tests/validation/half.h" #include +#include #include #include #include @@ -250,17 +251,6 @@ inline void fill_warp_matrix(std::array &matrix, int cols, int rows } } -/** Create a vector of random ROIs. - * - * @param[in] shape The shape of the input tensor. - * @param[in] pool_info The ROI pooling information. - * @param[in] num_rois The number of ROIs to be created. - * @param[in] seed The random seed to be used. - * - * @return A vector that contains the requested number of random ROIs - */ -std::vector generate_random_rois(const TensorShape &shape, const ROIPoolingLayerInfo &pool_info, unsigned int num_rois, std::random_device::result_type seed); - /** Helper function to fill the Lut random by a ILutAccessor. * * @param[in,out] table Accessor at the Lut. @@ -277,7 +267,6 @@ void fill_lookuptable(T &&table) table[i] = distribution(generator); } } - } // namespace validation } // namespace test } // namespace arm_compute diff --git a/tests/validation/NEON/Gaussian5x5.cpp b/tests/validation/NEON/Gaussian5x5.cpp index 522ececee9..7727340f66 100644 --- a/tests/validation/NEON/Gaussian5x5.cpp +++ b/tests/validation/NEON/Gaussian5x5.cpp @@ -41,8 +41,6 @@ #include "boost_wrapper.h" -#include - #include #include diff --git a/tests/validation/NEON/ROIPoolingLayer.cpp b/tests/validation/NEON/ROIPoolingLayer.cpp index 5775db7e03..523885d908 100644 --- a/tests/validation/NEON/ROIPoolingLayer.cpp +++ b/tests/validation/NEON/ROIPoolingLayer.cpp @@ -22,13 +22,12 @@ * SOFTWARE. */ #include "NEON/Accessor.h" -#include "NEON/Helper.h" +#include "NEON/ArrayAccessor.h" #include "TypePrinter.h" #include "arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h" #include "tests/Globals.h" #include "tests/Utils.h" #include "validation/Datasets.h" -#include "validation/Helpers.h" #include "validation/Reference.h" #include "validation/Validation.h" #include "validation/ValidationUserConfiguration.h" @@ -51,9 +50,12 @@ Tensor compute_roi_pooling_layer(const TensorShape &shape, DataType dt, const st shape_dst.set(3, rois.size()); // Create tensors - Tensor src = create_tensor(shape, dt); - Tensor dst = create_tensor(shape_dst, dt); - Array rois_array = create_array(rois); + Tensor src = create_tensor(shape, dt); + Tensor dst = create_tensor(shape_dst, dt); + + // Create ROI array + Array rois_array(rois.size()); + fill_array(ArrayAccessor(rois_array), rois); // Create and configure function NEROIPoolingLayer roi_pool; -- cgit v1.2.1