diff options
author | Manuel Bottini <manuel.bottini@arm.com> | 2019-01-09 17:04:39 +0000 |
---|---|---|
committer | Georgios Pinitas <georgios.pinitas@arm.com> | 2019-01-11 15:17:31 +0000 |
commit | cc5171b85654b9f19a5f52bbe8abea0572ee0163 (patch) | |
tree | 032999b06aeced4f5e2963ab60c833acb951c3c8 | |
parent | 587708b05ca63fa88118daec82e2c39d63e60086 (diff) | |
download | ComputeLibrary-cc5171b85654b9f19a5f52bbe8abea0572ee0163.tar.gz |
COMPMID-1677: Change ROIPooling layer interface to accept ROIs as tensors
Change-Id: If16b572a4d906187b77f32133a72a44316fa74e4
Reviewed-on: https://review.mlplatform.org/490
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
22 files changed, 231 insertions, 319 deletions
diff --git a/arm_compute/core/CL/ICLArray.h b/arm_compute/core/CL/ICLArray.h index 22fc7cf32e..eb57ea4ce3 100644 --- a/arm_compute/core/CL/ICLArray.h +++ b/arm_compute/core/CL/ICLArray.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -121,8 +121,6 @@ using ICLKeyPointArray = ICLArray<KeyPoint>; using ICLCoordinates2DArray = ICLArray<Coordinates2D>; /** Interface for OpenCL Array of Detection Windows. */ using ICLDetectionWindowArray = ICLArray<DetectionWindow>; -/** Interface for OpenCL Array of ROIs. */ -using ICLROIArray = ICLArray<ROI>; /** Interface for OpenCL Array of 2D Sizes. */ using ICLSize2DArray = ICLArray<Size2D>; /** Interface for OpenCL Array of uint8s. */ diff --git a/arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h b/arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h index 93bfb3004c..106a4b9b6d 100644 --- a/arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h +++ b/arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -52,7 +52,8 @@ 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[in] rois ROIs tensor, it is a 2D tensor of size [5, N] (where N is the number of ROIs) containing top left and bottom right corner + * as coordinate of an image and batch_id of ROI [ batch_id, x1, y1, x2, y2 ]. Data types supported: U16 * @param[out] output Destination tensor. Data types supported: Same as @p input. * @param[in] pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo. * @@ -61,14 +62,14 @@ public: * @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); + void configure(const ICLTensor *input, const ICLTensor *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; + const ICLTensor *_rois; ICLTensor *_output; ROIPoolingLayerInfo _pool_info; }; diff --git a/arm_compute/core/IArray.h b/arm_compute/core/IArray.h index f9e09a308b..35ab16c22a 100644 --- a/arm_compute/core/IArray.h +++ b/arm_compute/core/IArray.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -34,7 +34,6 @@ struct KeyPoint; struct Coordinates2D; struct DetectionWindow; class Size2D; -struct ROI; /** Array of type T */ template <class T> @@ -142,8 +141,6 @@ using IKeyPointArray = IArray<KeyPoint>; using ICoordinates2DArray = IArray<Coordinates2D>; /** Interface for Array of Detection Windows. */ using IDetectionWindowArray = IArray<DetectionWindow>; -/** Interface for Array of ROIs. */ -using IROIArray = IArray<ROI>; /** Interface for Array of 2D Sizes. */ using ISize2DArray = IArray<Size2D>; /** Interface for Array of uint8s. */ diff --git a/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h b/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h index 5d9a7cfbf6..cae305ba43 100644 --- a/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h +++ b/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2018 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -56,23 +56,24 @@ public: /** Set the input and output tensors. * * @param[in] input Source tensor. Data types supported: F32. - * @param[in] rois Array containing @ref ROI. + * @param[in] rois ROIs tensor, it is a 2D tensor of size [5, N] (where N is the number of ROIs) containing top left and bottom right corner + * as coordinate of an image and batch_id of ROI [ batch_id, x1, y1, x2, y2 ]. Data types supported: U16 * @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. + * @note The fourth dimension of @p output tensor must be the same as the number of elements in @p rois tensor. */ - void configure(const ITensor *input, const IROIArray *rois, ITensor *output, const ROIPoolingLayerInfo &pool_info); + void configure(const ITensor *input, const ITensor *rois, ITensor *output, const ROIPoolingLayerInfo &pool_info); // Inherited methods overridden: void run(const Window &window, const ThreadInfo &info) override; private: const ITensor *_input; - const IROIArray *_rois; + const ITensor *_rois; ITensor *_output; ROIPoolingLayerInfo _pool_info; }; diff --git a/arm_compute/core/Types.h b/arm_compute/core/Types.h index 02001a2438..dc87617f55 100644 --- a/arm_compute/core/Types.h +++ b/arm_compute/core/Types.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -501,13 +501,6 @@ using PaddingList = std::vector<PaddingInfo>; /** Information to produce a tiled version of a Tensor */ using Multiples = std::vector<uint32_t>; -/** Region of interest */ -struct ROI -{ - Rectangle rect; /**< Rectangle specifying the region of interest */ - uint16_t batch_idx; /**< The batch index of the region of interest */ -}; - /** Available channels */ enum class Channel { diff --git a/arm_compute/graph/GraphBuilder.h b/arm_compute/graph/GraphBuilder.h index cb905e700e..b73f4f23ca 100644 --- a/arm_compute/graph/GraphBuilder.h +++ b/arm_compute/graph/GraphBuilder.h @@ -369,7 +369,7 @@ public: * @param[in] g Graph to add the node to * @param[in] params Common node parameters * @param[in] input Input to the reshape layer node as a NodeID-Index pair - * @param[in] rois Input containing @ref ROI. + * @param[in] rois Input containing the ROIs. * @param[in] pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo. * * @return Node ID of the created node, EmptyNodeID in case of error diff --git a/arm_compute/runtime/Array.h b/arm_compute/runtime/Array.h index 4fc79026e0..0fe6dda047 100644 --- a/arm_compute/runtime/Array.h +++ b/arm_compute/runtime/Array.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -67,8 +67,6 @@ using KeyPointArray = Array<KeyPoint>; using Coordinates2DArray = Array<Coordinates2D>; /** Array of Detection Windows. */ using DetectionWindowArray = Array<DetectionWindow>; -/** Array of ROIs. */ -using ROIArray = Array<ROI>; /** Array of 2D Sizes. */ using Size2DArray = Array<Size2D>; /** Array of uint8s. */ diff --git a/arm_compute/runtime/CL/CLArray.h b/arm_compute/runtime/CL/CLArray.h index 01c6d8df3d..3ff9eb801c 100644 --- a/arm_compute/runtime/CL/CLArray.h +++ b/arm_compute/runtime/CL/CLArray.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -109,8 +109,6 @@ using CLKeyPointArray = CLArray<KeyPoint>; using CLCoordinates2DArray = CLArray<Coordinates2D>; /** OpenCL Array of Detection Windows. */ using CLDetectionWindowArray = CLArray<DetectionWindow>; -/** OpenCL Array of ROIs. */ -using CLROIArray = CLArray<ROI>; /** OpenCL Array of 2D Sizes. */ using CLSize2DArray = CLArray<Size2D>; /** OpenCL Array of uint8s. */ diff --git a/arm_compute/runtime/CL/functions/CLROIPoolingLayer.h b/arm_compute/runtime/CL/functions/CLROIPoolingLayer.h index f089375e51..70a3ba9c95 100644 --- a/arm_compute/runtime/CL/functions/CLROIPoolingLayer.h +++ b/arm_compute/runtime/CL/functions/CLROIPoolingLayer.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -45,7 +45,8 @@ 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[in] rois ROIs tensor, it is a 2D tensor of size [5, N] (where N is the number of ROIs) containing top left and bottom right corner + * as coordinate of an image and batch_id of ROI [ batch_id, x1, y1, x2, y2 ]. Data types supported: U16 * @param[out] output Destination tensor. Data types supported: Same as @p input. * @param[in] pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo. * @@ -54,7 +55,7 @@ public: * @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); + void configure(const ICLTensor *input, const ICLTensor *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 69a90dd89a..cf41552694 100644 --- a/arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h +++ b/arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -47,7 +47,8 @@ public: /** Set the input and output tensors. * * @param[in] input Source tensor. Data types supported: F32. - * @param[in] rois Array containing @ref ROI. + * @param[in] rois ROIs tensor, it is a 2D tensor of size [5, N] (where N is the number of ROIs) containing top left and bottom right corner + * as coordinate of an image and batch_id of ROI [ batch_id, x1, y1, x2, y2 ]. Data types supported: U16 * @param[out] output Destination tensor. Data types supported: Same as @p input. * @param[in] pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo. * @@ -56,7 +57,7 @@ public: * @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); + void configure(const ITensor *input, const ITensor *rois, ITensor *output, const ROIPoolingLayerInfo &pool_info); // Inherited methods overridden: void run() override; @@ -64,5 +65,5 @@ public: private: NEROIPoolingLayerKernel _roi_kernel; }; -} +} // namespace arm_compute #endif /* __ARM_COMPUTE_NEROIPOOLINGLAYER_H__ */ diff --git a/src/core/CL/cl_kernels/roi_pooling_layer.cl b/src/core/CL/cl_kernels/roi_pooling_layer.cl index 042b102a15..0cf296c011 100644 --- a/src/core/CL/cl_kernels/roi_pooling_layer.cl +++ b/src/core/CL/cl_kernels/roi_pooling_layer.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -105,10 +105,12 @@ inline DATA_TYPE roi_pool_1x1(const Tensor3D *input, int region_start_x, int reg * @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[in] rois_ptr Pointer to the ROIs tensor. Layout: { batch_index, x1, y1, x2, y2 }. Supported data types: same as @p input_ptr + * @param[in] rois_stride_x Stride of the ROIs tensor in X dimension (in bytes) + * @param[in] rois_step_x Step of the ROIs tensor in X dimension (in bytes) + * @param[in] rois_stride_y Stride of the ROIs tensor in Y dimension (in bytes) + * @param[in] rois_step_y Step of the ROIs tensor in Y dimension (in bytes) + * @param[in] rois_offset_first_element_in_bytes The offset of the first element in the ROIs tensor * @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) @@ -122,13 +124,13 @@ inline DATA_TYPE roi_pool_1x1(const Tensor3D *input, int region_start_x, int reg */ __kernel void roi_pooling_layer( TENSOR3D_DECLARATION(input), - VECTOR_DECLARATION(rois), + IMAGE_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); + Image rois = CONVERT_TO_IMAGE_STRUCT_NO_STEP(rois); Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output); const int px = get_global_id(0); @@ -136,12 +138,12 @@ __kernel void roi_pooling_layer( const int pw = get_global_id(2); // Load roi parameters - // roi is laid out as follows: - // { x, y, width, height, batch_index } - const ushort4 roi = vload4(0, (__global ushort *)vector_offset(&rois, pw)); - const ushort roi_batch = *((__global ushort *)vector_offset(&rois, pw) + 4); - 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)); + // roi is laid out as follows { batch_index, x1, y1, x2, y2 } + const ushort roi_batch = (ushort) * ((__global DATA_TYPE *)offset(&rois, 0, pw)); + const VEC_DATA_TYPE(DATA_TYPE, 4) + roi = vload4(0, (__global DATA_TYPE *)offset(&rois, 1, 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 - roi.s01) * (float)SPATIAL_SCALE), 1.f)); // Calculate pooled region start and end const float2 spatial_indx = (float2)(px, py); diff --git a/src/core/CL/kernels/CLROIPoolingLayerKernel.cpp b/src/core/CL/kernels/CLROIPoolingLayerKernel.cpp index 23676942a6..df7687edea 100644 --- a/src/core/CL/kernels/CLROIPoolingLayerKernel.cpp +++ b/src/core/CL/kernels/CLROIPoolingLayerKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2018 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -39,29 +39,61 @@ #include <set> #include <string> -using namespace arm_compute; +namespace arm_compute +{ +namespace +{ +std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITensorInfo *rois, ITensorInfo *output, const ROIPoolingLayerInfo &pool_info) +{ + ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); + + // Output auto initialization if not yet initialized + TensorShape output_shape(pool_info.pooled_width(), pool_info.pooled_height(), input->dimension(2), rois->dimension(1)); + auto_init_if_empty((*output), output_shape, 1, input->data_type()); + + // Configure kernel window + const unsigned int num_elems_processed_per_iteration = 1; + Window win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration)); + + AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration); + AccessWindowHorizontal input_access(input, input->valid_region().start(0), num_elems_processed_per_iteration); + + bool window_changed = update_window_and_padding(win, input_access, output_access); + output_access.set_valid_region(win, ValidRegion(Coordinates(), output->tensor_shape())); + Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; + return std::make_pair(err, win); +} +} // namespace 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) +void CLROIPoolingLayerKernel::configure(const ICLTensor *input, const ICLTensor *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info) { ARM_COMPUTE_ERROR_ON_NULLPTR(input, rois, output); + + //Validate arguments + ARM_COMPUTE_ERROR_ON_NULLPTR(input->info(), rois->info(), output->info()); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(rois, 1, DataType::U16); + ARM_COMPUTE_ERROR_ON(rois->info()->dimension(0) != 5); + ARM_COMPUTE_ERROR_ON(rois->info()->num_dimensions() > 2); ARM_COMPUTE_ERROR_ON_F16_UNSUPPORTED(input); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32, DataType::F16); 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()); + if(output->info()->total_size() != 0) + { + 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->info()->dimension(1) != output->info()->dimension(3)); + } - 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)); + // Configure kernel window + auto win_config = validate_and_configure_window(input->info(), rois->info(), output->info(), pool_info); + ARM_COMPUTE_ERROR_THROW_ON(win_config.first); // Set instance variables _input = input; @@ -89,19 +121,7 @@ void CLROIPoolingLayerKernel::configure(const ICLTensor *input, const ICLROIArra add_argument<cl_uint>(idx, _input->info()->strides_in_bytes()[3]); add_argument<cl_uint>(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_internal(window); + ICLKernel::configure_internal(win_config.second); } void CLROIPoolingLayerKernel::run(const Window &window, cl::CommandQueue &queue) @@ -109,14 +129,20 @@ 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.first_slice_window_3D(); - // Parallelize spatially and across the fourth dimension of the output tensor (also across ROIArray) + Window slice = window.first_slice_window_3D(); + Window slice_rois = slice; + // Parallelize spatially and across the fourth dimension of the output tensor (also across ROITensor) + slice_rois.set_dimension_step(Window::DimX, _rois->info()->dimension(0)); slice.set(Window::DimZ, window[3]); // Set arguments unsigned int idx = 0; add_3D_tensor_argument(idx, _input, slice); - add_1D_array_argument<ROI>(idx, _rois, Strides(sizeof(ROI)), 1U, slice); + add_2D_tensor_argument(idx, _rois, slice_rois); add_3D_tensor_argument(idx, _output, slice); + add_argument<cl_uint>(idx, _input->info()->strides_in_bytes()[3]); + add_argument<cl_uint>(idx, _output->info()->strides_in_bytes()[3]); + enqueue(queue, *this, slice); } +} // namespace arm_compute
\ No newline at end of file diff --git a/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp b/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp index 4d908db77b..6fd6792ff8 100644 --- a/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp +++ b/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2018 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -35,22 +35,35 @@ #include <cfloat> #include <cmath> -using namespace arm_compute; - +namespace arm_compute +{ NEROIPoolingLayerKernel::NEROIPoolingLayerKernel() : _input(nullptr), _rois(nullptr), _output(nullptr), _pool_info(0, 0, 0.f) { } -void NEROIPoolingLayerKernel::configure(const ITensor *input, const IROIArray *rois, ITensor *output, const ROIPoolingLayerInfo &pool_info) +void NEROIPoolingLayerKernel::configure(const ITensor *input, const ITensor *rois, ITensor *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::F32); + ARM_COMPUTE_ERROR_ON_NULLPTR(input, output, rois); + + //Validate arguments + ARM_COMPUTE_ERROR_ON_NULLPTR(input->info(), rois->info(), output->info()); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(rois, 1, DataType::U16); + ARM_COMPUTE_ERROR_ON(rois->info()->dimension(0) != 5); + ARM_COMPUTE_ERROR_ON(rois->info()->num_dimensions() > 2); + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32, DataType::F16); 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()); + if(output->info()->total_size() != 0) + { + 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->info()->dimension(1) != output->info()->dimension(3)); + } + + // Output auto initialization if not yet initialized + TensorShape output_shape(pool_info.pooled_width(), pool_info.pooled_height(), input->info()->dimension(2), rois->info()->dimension(1)); auto_init_if_empty(*output->info(), output_shape, 1, input->info()->data_type()); ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); @@ -64,7 +77,7 @@ void NEROIPoolingLayerKernel::configure(const ITensor *input, const IROIArray *r // Configure kernel window Window window; - window.set(Window::DimX, Window::Dimension(0, rois->num_values())); + window.set(Window::DimX, Window::Dimension(0, rois->info()->dimension(1))); window.set(Window::DimY, Window::Dimension(0, 1)); AccessWindowStatic input_access(input->info(), @@ -85,6 +98,8 @@ void NEROIPoolingLayerKernel::run(const Window &window, const ThreadInfo &info) ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window); + const size_t values_per_roi = _rois->info()->dimension(0); + const int roi_list_start = window.x().start(); const int roi_list_end = window.x().end(); const int width = _input->info()->dimension(Window::DimX); @@ -94,16 +109,21 @@ void NEROIPoolingLayerKernel::run(const Window &window, const ThreadInfo &info) const int pooled_h = _pool_info.pooled_height(); const float spatial_scale = _pool_info.spatial_scale(); + const auto *rois_ptr = reinterpret_cast<const uint16_t *>(_rois->buffer()); + for(int roi_indx = roi_list_start; roi_indx < roi_list_end; ++roi_indx) { - const ROI &curr_roi = _rois->at(roi_indx); + const unsigned int roi_batch = rois_ptr[values_per_roi * roi_indx]; + const auto x1 = rois_ptr[values_per_roi * roi_indx + 1]; + const auto y1 = rois_ptr[values_per_roi * roi_indx + 2]; + const auto x2 = rois_ptr[values_per_roi * roi_indx + 3]; + const auto y2 = rois_ptr[values_per_roi * roi_indx + 4]; // Scale ROI - const int roi_batch = curr_roi.batch_idx; - const int roi_anchor_x = support::cpp11::round(curr_roi.rect.x * spatial_scale); - const int roi_anchor_y = support::cpp11::round(curr_roi.rect.y * spatial_scale); - const int roi_width = std::max(support::cpp11::round(curr_roi.rect.width * spatial_scale), 1.f); - const int roi_height = std::max(support::cpp11::round(curr_roi.rect.height * spatial_scale), 1.f); + const int roi_anchor_x = support::cpp11::round(x1 * spatial_scale); + const int roi_anchor_y = support::cpp11::round(y1 * spatial_scale); + const int roi_width = std::max(support::cpp11::round((x2 - x1) * spatial_scale), 1.f); + const int roi_height = std::max(support::cpp11::round((y2 - y1) * spatial_scale), 1.f); // Iterate through all feature maps for(int fm = 0; fm < fms; ++fm) @@ -146,3 +166,4 @@ void NEROIPoolingLayerKernel::run(const Window &window, const ThreadInfo &info) } } } +} // namespace arm_compute
\ No newline at end of file diff --git a/src/runtime/CL/functions/CLROIPoolingLayer.cpp b/src/runtime/CL/functions/CLROIPoolingLayer.cpp index 0f480eeac9..7bb41784ac 100644 --- a/src/runtime/CL/functions/CLROIPoolingLayer.cpp +++ b/src/runtime/CL/functions/CLROIPoolingLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -30,7 +30,7 @@ using namespace arm_compute; -void CLROIPoolingLayer::configure(const ICLTensor *input, const ICLROIArray *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info) +void CLROIPoolingLayer::configure(const ICLTensor *input, const ICLTensor *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info) { // Configure ROI pooling kernel auto k = arm_compute::support::cpp14::make_unique<CLROIPoolingLayerKernel>(); diff --git a/src/runtime/NEON/functions/NEROIPoolingLayer.cpp b/src/runtime/NEON/functions/NEROIPoolingLayer.cpp index 1f1400cf42..3aca4b7b60 100644 --- a/src/runtime/NEON/functions/NEROIPoolingLayer.cpp +++ b/src/runtime/NEON/functions/NEROIPoolingLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -27,14 +27,14 @@ #include "arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h" #include "arm_compute/runtime/NEON/NEScheduler.h" -using namespace arm_compute; - +namespace arm_compute +{ NEROIPoolingLayer::NEROIPoolingLayer() : _roi_kernel() { } -void NEROIPoolingLayer::configure(const ITensor *input, const IROIArray *rois, ITensor *output, const ROIPoolingLayerInfo &pool_info) +void NEROIPoolingLayer::configure(const ITensor *input, const ITensor *rois, ITensor *output, const ROIPoolingLayerInfo &pool_info) { _roi_kernel.configure(input, rois, output, pool_info); } @@ -43,3 +43,4 @@ void NEROIPoolingLayer::run() { NEScheduler::get().schedule(&_roi_kernel, Window::DimX); } +} // namespace arm_compute
\ No newline at end of file diff --git a/tests/Utils.h b/tests/Utils.h index 111d80fdfe..7c55a3ef50 100644 --- a/tests/Utils.h +++ b/tests/Utils.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2018 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -647,52 +647,6 @@ inline void init_separable_conv(int16_t *conv, unsigned int width, unsigned int } } -/** 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<ROI> 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<ROI> 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<int>((shape.x() / roi_scale) / pool_width); - const auto scaled_height = static_cast<int>((shape.y() / roi_scale) / pool_height); - const auto min_width = static_cast<int>(pool_width / roi_scale); - const auto min_height = static_cast<int>(pool_height / roi_scale); - - // Create distributions - std::uniform_int_distribution<int> dist_batch(0, shape[3] - 1); - std::uniform_int_distribution<int> dist_x(0, scaled_width); - std::uniform_int_distribution<int> dist_y(0, scaled_height); - std::uniform_int_distribution<int> dist_w(min_width, std::max(min_width, (pool_width - 2) * scaled_width)); - std::uniform_int_distribution<int> 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; -} - /** Create a vector with a uniform distribution of floating point values across the specified range. * * @param[in] num_values The number of values to be created. diff --git a/tests/benchmark/CL/ROIPoolingLayer.cpp b/tests/benchmark/CL/ROIPoolingLayer.cpp index 2ef91d45fd..eadb027b75 100644 --- a/tests/benchmark/CL/ROIPoolingLayer.cpp +++ b/tests/benchmark/CL/ROIPoolingLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2018 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -30,7 +30,7 @@ #include "tests/CL/CLAccessor.h" #include "tests/CL/CLArrayAccessor.h" #include "tests/benchmark/fixtures/ROIPoolingLayerFixture.h" -#include "tests/datasets/ROIPoolingLayerDataset.h" +#include "tests/datasets/ROIDataset.h" #include "tests/framework/Macros.h" #include "tests/framework/datasets/Datasets.h" #include "utils/TypePrinter.h" @@ -41,16 +41,20 @@ namespace test { namespace benchmark { -using CLROIPoolingLayerFixture = ROIPoolingLayerFixture<CLTensor, CLROIPoolingLayer, CLAccessor, CLArray<ROI>, CLArrayAccessor<ROI>>; +template <typename T> +using CLROIPoolingLayerFixture = ROIPoolingLayerFixture<CLTensor, CLROIPoolingLayer, CLAccessor, T>; 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 })), +REGISTER_FIXTURE_DATA_TEST_CASE(SmallROIPoolingLayerHalf, CLROIPoolingLayerFixture<half>, framework::DatasetMode::ALL, + framework::dataset::combine(framework::dataset::combine(datasets::SmallROIDataset(), + framework::dataset::make("DataType", { DataType::F16 })), framework::dataset::make("Batches", { 1, 4, 8 }))); -TEST_SUITE_END() +REGISTER_FIXTURE_DATA_TEST_CASE(SmallROIPoolingLayerFloat, CLROIPoolingLayerFixture<float>, framework::DatasetMode::ALL, + framework::dataset::combine(framework::dataset::combine(datasets::SmallROIDataset(), + framework::dataset::make("DataType", { DataType::F32 })), + framework::dataset::make("Batches", { 1, 4, 8 }))); +TEST_SUITE_END() // CL } // namespace benchmark } // namespace test } // namespace arm_compute diff --git a/tests/benchmark/NEON/ROIPoolingLayer.cpp b/tests/benchmark/NEON/ROIPoolingLayer.cpp index 02cf47b189..10652a5606 100644 --- a/tests/benchmark/NEON/ROIPoolingLayer.cpp +++ b/tests/benchmark/NEON/ROIPoolingLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2018 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -30,7 +30,7 @@ #include "tests/NEON/Accessor.h" #include "tests/NEON/ArrayAccessor.h" #include "tests/benchmark/fixtures/ROIPoolingLayerFixture.h" -#include "tests/datasets/ROIPoolingLayerDataset.h" +#include "tests/datasets/ROIDataset.h" #include "tests/framework/Macros.h" #include "tests/framework/datasets/Datasets.h" #include "utils/TypePrinter.h" @@ -41,16 +41,17 @@ namespace test { namespace benchmark { -using NEROIPoolingLayerFixture = ROIPoolingLayerFixture<Tensor, NEROIPoolingLayer, Accessor, Array<ROI>, ArrayAccessor<ROI>>; +template <typename T> +using NEROIPoolingLayerFixture = ROIPoolingLayerFixture<Tensor, NEROIPoolingLayer, Accessor, T>; TEST_SUITE(NEON) -REGISTER_FIXTURE_DATA_TEST_CASE(SmallROIPoolingLayer, NEROIPoolingLayerFixture, framework::DatasetMode::ALL, - framework::dataset::combine(framework::dataset::combine(datasets::SmallROIPoolingLayerDataset(), +REGISTER_FIXTURE_DATA_TEST_CASE(SmallROIPoolingLayerFloat, NEROIPoolingLayerFixture<float>, framework::DatasetMode::ALL, + framework::dataset::combine(framework::dataset::combine(datasets::SmallROIDataset(), framework::dataset::make("DataType", { DataType::F32 })), framework::dataset::make("Batches", { 1, 4, 8 }))); -TEST_SUITE_END() +TEST_SUITE_END() // NEON } // namespace benchmark } // namespace test } // namespace arm_compute diff --git a/tests/benchmark/fixtures/ROIPoolingLayerFixture.h b/tests/benchmark/fixtures/ROIPoolingLayerFixture.h index fa4a5b7044..2c828272de 100644 --- a/tests/benchmark/fixtures/ROIPoolingLayerFixture.h +++ b/tests/benchmark/fixtures/ROIPoolingLayerFixture.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2018 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -39,37 +39,38 @@ namespace test namespace benchmark { /** Fixture that can be used for NEON and CL */ -template <typename TensorType, typename Function, typename Accessor, typename Array_T, typename ArrayAccessor> +template <typename TensorType, typename Function, typename AccessorType, typename T> class ROIPoolingLayerFixture : public framework::Fixture { public: template <typename...> - void setup(TensorShape shape, const ROIPoolingLayerInfo pool_info, unsigned int num_rois, DataType data_type, int batches) + void setup(TensorShape input_shape, const ROIPoolingLayerInfo pool_info, TensorShape rois_shape, DataType data_type, int batches) { // Set batched in source and destination shapes TensorShape shape_dst; - shape.set(shape.num_dimensions(), batches); + rois_tensor = create_tensor<TensorType>(rois_shape, DataType::U16); + + input_shape.set(input_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); + shape_dst.set(2, input_shape.z()); + shape_dst.set(3, rois_shape[1]); // Create tensors - src = create_tensor<TensorType>(shape, data_type, 1); + src = create_tensor<TensorType>(input_shape, data_type, 1); dst = create_tensor<TensorType>(shape_dst, data_type, 1); - // Create random ROIs - std::vector<ROI> rois = generate_random_rois(shape, pool_info, num_rois, 0U); - rois_array = arm_compute::support::cpp14::make_unique<Array_T>(num_rois); - fill_array(ArrayAccessor(*rois_array), rois); - // Create and configure function - roi_pool.configure(&src, rois_array.get(), &dst, pool_info); + roi_pool.configure(&src, &rois_tensor, &dst, pool_info); // Allocate tensors + rois_tensor.allocator()->allocate(); src.allocator()->allocate(); dst.allocator()->allocate(); + + // Create random ROIs + generate_rois(AccessorType(rois_tensor), input_shape, pool_info, rois_shape); } void run() @@ -89,11 +90,54 @@ public: dst.allocator()->free(); } +protected: + template <typename U> + void generate_rois(U &&rois, const TensorShape &shape, const ROIPoolingLayerInfo &pool_info, TensorShape rois_shape) + { + const size_t values_per_roi = rois_shape.x(); + const size_t num_rois = rois_shape.y(); + + std::mt19937 gen(library->seed()); + uint16_t *rois_ptr = static_cast<uint16_t *>(rois.data()); + + const float pool_width = pool_info.pooled_width(); + const float pool_height = pool_info.pooled_height(); + const float roi_scale = pool_info.spatial_scale(); + + // Calculate distribution bounds + const auto scaled_width = static_cast<uint16_t>((shape.x() / roi_scale) / pool_width); + const auto scaled_height = static_cast<uint16_t>((shape.y() / roi_scale) / pool_height); + const auto min_width = static_cast<uint16_t>(pool_width / roi_scale); + const auto min_height = static_cast<uint16_t>(pool_height / roi_scale); + + // Create distributions + std::uniform_int_distribution<int> dist_batch(0, shape[3] - 1); + std::uniform_int_distribution<uint16_t> dist_x1(0, scaled_width); + std::uniform_int_distribution<uint16_t> dist_y1(0, scaled_height); + std::uniform_int_distribution<uint16_t> dist_w(min_width, std::max(float(min_width), (pool_width - 2) * scaled_width)); + std::uniform_int_distribution<uint16_t> dist_h(min_height, std::max(float(min_height), (pool_height - 2) * scaled_height)); + + for(unsigned int pw = 0; pw < num_rois; ++pw) + { + const auto batch_idx = dist_batch(gen); + const auto x1 = dist_x1(gen); + const auto y1 = dist_y1(gen); + const auto x2 = x1 + dist_w(gen); + const auto y2 = y1 + dist_h(gen); + + rois_ptr[values_per_roi * pw] = batch_idx; + rois_ptr[values_per_roi * pw + 1] = x1; + rois_ptr[values_per_roi * pw + 2] = y1; + rois_ptr[values_per_roi * pw + 3] = x2; + rois_ptr[values_per_roi * pw + 4] = y2; + } + } + private: - TensorType src{}; - TensorType dst{}; - std::unique_ptr<Array_T> rois_array{}; - Function roi_pool{}; + TensorType src{}; + TensorType dst{}; + TensorType rois_tensor{}; + Function roi_pool{}; }; } // namespace benchmark } // namespace test diff --git a/tests/datasets/ROIAlignLayerDataset.h b/tests/datasets/ROIDataset.h index 27c6ee4d91..9e21ab119c 100644 --- a/tests/datasets/ROIAlignLayerDataset.h +++ b/tests/datasets/ROIDataset.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018 ARM Limited. + * Copyright (c) 2018-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -35,7 +35,7 @@ namespace test { namespace datasets { -class ROIAlignLayerDataset +class ROIDataset { public: using type = std::tuple<TensorShape, ROIPoolingLayerInfo, TensorShape>; @@ -60,7 +60,7 @@ public: return description.str(); } - ROIAlignLayerDataset::type operator*() const + ROIDataset::type operator*() const { return std::make_tuple(*_tensor_shape_it, *_infos_it, *_rois_shape_it); } @@ -98,8 +98,8 @@ public: } protected: - ROIAlignLayerDataset() = default; - ROIAlignLayerDataset(ROIAlignLayerDataset &&) = default; + ROIDataset() = default; + ROIDataset(ROIDataset &&) = default; private: std::vector<TensorShape> _tensor_shapes{}; @@ -107,10 +107,10 @@ private: std::vector<TensorShape> _rois_shape{}; }; -class SmallROIAlignLayerDataset final : public ROIAlignLayerDataset +class SmallROIDataset final : public ROIDataset { public: - SmallROIAlignLayerDataset() + SmallROIDataset() { add_config(TensorShape(50U, 47U, 1U, 1U), ROIPoolingLayerInfo(7U, 7U, 1.f / 4.f), TensorShape(5U, 1U)); add_config(TensorShape(50U, 47U, 3U, 4U), ROIPoolingLayerInfo(7U, 7U, 1.f / 4.f), TensorShape(5U, 1U)); diff --git a/tests/datasets/ROIPoolingLayerDataset.h b/tests/datasets/ROIPoolingLayerDataset.h deleted file mode 100644 index eb1d165202..0000000000 --- a/tests/datasets/ROIPoolingLayerDataset.h +++ /dev/null @@ -1,129 +0,0 @@ -/* - * Copyright (c) 2017-2018 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 "utils/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<TensorShape, ROIPoolingLayerInfo, int>; - - struct iterator - { - iterator(std::vector<TensorShape>::const_iterator tensor_shape_it, - std::vector<ROIPoolingLayerInfo>::const_iterator infos_it, - std::vector<unsigned int>::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<TensorShape>::const_iterator _tensor_shape_it; - std::vector<ROIPoolingLayerInfo>::const_iterator _infos_it; - std::vector<unsigned int>::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<TensorShape> _tensor_shapes{}; - std::vector<ROIPoolingLayerInfo> _infos{}; - std::vector<unsigned int> _num_rois{}; -}; - -class SmallROIPoolingLayerDataset final : public ROIPoolingLayerDataset -{ -public: - SmallROIPoolingLayerDataset() - { - add_config(TensorShape(50U, 47U, 1U), ROIPoolingLayerInfo(7U, 7U, 1.f / 8.f), 1U); - add_config(TensorShape(50U, 47U, 3U), ROIPoolingLayerInfo(7U, 7U, 1.f / 8.f), 1U); - 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/validation/CL/ROIAlignLayer.cpp b/tests/validation/CL/ROIAlignLayer.cpp index f3fc3818f2..926a3de68d 100644 --- a/tests/validation/CL/ROIAlignLayer.cpp +++ b/tests/validation/CL/ROIAlignLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018 ARM Limited. + * Copyright (c) 2018-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -25,7 +25,7 @@ #include "arm_compute/runtime/CL/functions/CLROIAlignLayer.h" #include "tests/CL/CLAccessor.h" #include "tests/Globals.h" -#include "tests/datasets/ROIAlignLayerDataset.h" +#include "tests/datasets/ROIDataset.h" #include "tests/datasets/ShapeDatasets.h" #include "tests/framework/Macros.h" #include "tests/framework/datasets/Datasets.h" @@ -100,14 +100,14 @@ using CLROIAlignLayerFixture = ROIAlignLayerFixture<CLTensor, CLAccessor, CLROIA TEST_SUITE(Float) FIXTURE_DATA_TEST_CASE(SmallROIAlignLayerFloat, CLROIAlignLayerFixture<float>, framework::DatasetMode::ALL, - framework::dataset::combine(datasets::SmallROIAlignLayerDataset(), + framework::dataset::combine(datasets::SmallROIDataset(), framework::dataset::make("DataType", { DataType::F32 }))) { // Validate output validate(CLAccessor(_target), _reference, relative_tolerance_f32, .02f, absolute_tolerance_f32); } FIXTURE_DATA_TEST_CASE(SmallROIAlignLayerHalf, CLROIAlignLayerFixture<half>, framework::DatasetMode::ALL, - framework::dataset::combine(datasets::SmallROIAlignLayerDataset(), + framework::dataset::combine(datasets::SmallROIDataset(), framework::dataset::make("DataType", { DataType::F16 }))) { // Validate output |