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