COMPMID-358 Implement OpenCL ROI Pooling

* Implement OpenCL ROI Pooling
* Add CLROIPoolingLayer benchmarks

Change-Id: I8786d01d551850a1b4d599a48fabe3925e0a27d0
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/79833
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>

28 files changed, 1296 insertions, 112 deletions

diff --git a/arm_compute/core/CL/ICLArray.h b/arm_compute/core/CL/ICLArray.h
index 1b676ed5a3..e12695f206 100644
--- a/arm_compute/core/CL/ICLArray.h
+++ b/arm_compute/core/CL/ICLArray.h
@@ -107,6 +107,7 @@ private:
 using ICLKeyPointArray        = ICLArray<KeyPoint>;
 using ICLCoordinates2DArray   = ICLArray<Coordinates2D>;
 using ICLDetectionWindowArray = ICLArray<DetectionWindow>;
+using ICLROIArray             = ICLArray<ROI>;
 using ICLSize2DArray          = ICLArray<Size2D>;
 using ICLUInt8Array           = ICLArray<cl_uchar>;
 using ICLUInt16Array          = ICLArray<cl_ushort>;
diff --git a/arm_compute/core/CL/ICLKernel.h b/arm_compute/core/CL/ICLKernel.h
index cfbf760f1e..1334c54a6c 100644
--- a/arm_compute/core/CL/ICLKernel.h
+++ b/arm_compute/core/CL/ICLKernel.h
@@ -31,6 +31,8 @@
 
 namespace arm_compute
 {
+template <typename T>
+class ICLArray;
 class ICLTensor;
 class Window;
 
@@ -45,6 +47,16 @@ public:
      * @return A reference to the OpenCL kernel of this object.
      */
     cl::Kernel &kernel();
+    /** Add the passed 1D array's parameters to the object's kernel's arguments starting from the index idx.
+     *
+     * @param[in,out] idx            Index at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
+     * @param[in]     array          Array to set as an argument of the object's kernel.
+     * @param[in]     strides        @ref Strides object containing stride of each dimension in bytes.
+     * @param[in]     num_dimensions Number of dimensions of the @p array.
+     * @param[in]     window         Window the kernel will be executed on.
+     */
+    template <typename T>
+    void add_1D_array_argument(unsigned int &idx, const ICLArray<T> *array, const Strides &strides, unsigned int num_dimensions, const Window &window);
     /** Add the passed 1D tensor's parameters to the object's kernel's arguments starting from the index idx.
      *
      * @param[in,out] idx    Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
@@ -73,6 +85,11 @@ public:
      * @param[in]     window Window the kernel will be executed on.
      */
     void add_4D_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window);
+    /** Returns the number of arguments enqueued per 1D array object.
+     *
+     * @return The number of arguments enqueues per 1D array object.
+     */
+    unsigned int num_arguments_per_1D_array() const;
     /** Returns the number of arguments enqueued per 1D tensor object.
      *
      * @return The number of arguments enqueues per 1D tensor object.
@@ -142,6 +159,16 @@ public:
     GPUTarget get_target() const;
 
 private:
+    /** Add the passed array's parameters to the object's kernel's arguments starting from the index idx.
+     *
+     * @param[in,out] idx            Index at which to start adding the array's arguments. Will be incremented by the number of kernel arguments set.
+     * @param[in]     array          Array to set as an argument of the object's kernel.
+     * @param[in]     strides        @ref Strides object containing stride of each dimension in bytes.
+     * @param[in]     num_dimensions Number of dimensions of the @p array.
+     * @param[in]     window         Window the kernel will be executed on.
+     */
+    template <typename T, unsigned int dimension_size>
+    void add_array_argument(unsigned int &idx, const ICLArray<T> *array, const Strides &strides, unsigned int num_dimensions, const Window &window);
     /** Add the passed tensor's parameters to the object's kernel's arguments starting from the index idx.
      *
      * @param[in,out] idx    Index at which to start adding the tensor's arguments. Will be incremented by the number of kernel arguments set.
@@ -150,6 +177,12 @@ private:
      */
     template <unsigned int dimension_size>
     void add_tensor_argument(unsigned int &idx, const ICLTensor *tensor, const Window &window);
+    /** Returns the number of arguments enqueued per array object.
+     *
+     * @return The number of arguments enqueued per array object.
+     */
+    template <unsigned int dimension_size>
+    unsigned int           num_arguments_per_array() const;
     /** Returns the number of arguments enqueued per tensor object.
      *
      * @return The number of arguments enqueued per tensor object.
@@ -177,5 +210,50 @@ protected:
  * @note If any dimension of the lws is greater than the global workgroup size then no lws will be passed.
  */
 void enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Window &window, const cl::NDRange &lws_hint = CLKernelLibrary::get().default_ndrange());
+
+template <typename T, unsigned int dimension_size>
+void ICLKernel::add_array_argument(unsigned &idx, const ICLArray<T> *array, const Strides &strides, unsigned int num_dimensions, const Window &window)
+{
+    // Calculate offset to the start of the window
+    unsigned int offset_first_element = 0;
+
+    for(unsigned int n = 0; n < num_dimensions; ++n)
+    {
+        offset_first_element += window[n].start() * strides[n];
+    }
+
+    unsigned int idx_start = idx;
+    _kernel.setArg(idx++, array->cl_buffer());
+
+    for(unsigned int dimension = 0; dimension < dimension_size; dimension++)
+    {
+        _kernel.setArg<cl_uint>(idx++, strides[dimension]);
+        _kernel.setArg<cl_uint>(idx++, strides[dimension] * window[dimension].step());
+    }
+
+    _kernel.setArg<cl_uint>(idx++, offset_first_element);
+
+    ARM_COMPUTE_ERROR_ON_MSG(idx_start + num_arguments_per_array<dimension_size>() != idx,
+                             "add_%dD_array_argument() is supposed to add exactly %d arguments to the kernel", dimension_size, num_arguments_per_array<dimension_size>());
+    ARM_COMPUTE_UNUSED(idx_start);
+}
+
+template <typename T>
+void ICLKernel::add_1D_array_argument(unsigned int &idx, const ICLArray<T> *array, const Strides &strides, unsigned int num_dimensions, const Window &window)
+{
+    add_array_argument<T, 1>(idx, array, strides, num_dimensions, window);
+}
+
+template <unsigned int dimension_size>
+unsigned int           ICLKernel::num_arguments_per_array() const
+{
+    return num_arguments_per_tensor<dimension_size>();
+}
+
+template <unsigned int dimension_size>
+unsigned int           ICLKernel::num_arguments_per_tensor() const
+{
+    return 2 + 2 * dimension_size;
+}
 }
 #endif /*__ARM_COMPUTE_ICLKERNEL_H__ */
diff --git a/arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h b/arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h
new file mode 100644
index 0000000000..51aae30561
--- /dev/null
+++ b/arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h
@@ -0,0 +1,76 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_CLROIPOOLINGLAYERKERNEL_H__
+#define __ARM_COMPUTE_CLROIPOOLINGLAYERKERNEL_H__
+
+#include "arm_compute/core/CL/ICLKernel.h"
+
+#include "arm_compute/core/CL/ICLArray.h"
+
+namespace arm_compute
+{
+class ICLTensor;
+
+/** Interface for the ROI pooling layer kernel */
+class CLROIPoolingLayerKernel : public ICLKernel
+{
+public:
+    /** Default constructor */
+    CLROIPoolingLayerKernel();
+    /** Prevent instances of this class from being copied (As this class contains pointers) */
+    CLROIPoolingLayerKernel(const CLROIPoolingLayerKernel &) = delete;
+    /** Prevent instances of this class from being copied (As this class contains pointers) */
+    CLROIPoolingLayerKernel &operator=(const CLROIPoolingLayerKernel &) = delete;
+    /** Allow instances of this class to be moved */
+    CLROIPoolingLayerKernel(CLROIPoolingLayerKernel &&) = default;
+    /** Allow instances of this class to be moved */
+    CLROIPoolingLayerKernel &operator=(CLROIPoolingLayerKernel &&) = default;
+    /** Default destructor */
+    ~CLROIPoolingLayerKernel() = default;
+
+    /** Set the input and output tensors.
+     *
+     * @param[in]  input     Source tensor. Data types supported: F16/F32.
+     * @param[in]  rois      Array containing @ref ROI.
+     * @param[out] output    Destination tensor. Data types supported: Same as @p input.
+     * @param[in]  pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo.
+     *
+     * @note The x and y dimensions of @p output tensor must be the same as @p pool_info 's pooled
+     * width and pooled height.
+     * @note The z dimensions of @p output tensor and @p input tensor must be the same.
+     * @note The fourth dimension of @p output tensor must be the same as the number of elements in @p rois array.
+     */
+    void configure(const ICLTensor *input, const ICLROIArray *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info);
+
+    // Inherited methods overridden:
+    void run(const Window &window, cl::CommandQueue &queue) override;
+
+private:
+    const ICLTensor    *_input;
+    const ICLROIArray *_rois;
+    ICLTensor          *_output;
+    ROIPoolingLayerInfo _pool_info;
+};
+}
+#endif /*__ARM_COMPUTE_CLROIPOOLINGLAYERKERNEL_H__ */
diff --git a/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h b/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h
index 3a2f761370..40f79acc79 100644
--- a/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h
+++ b/arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h
@@ -52,9 +52,14 @@ public:
     /** Set the input and output tensors.
      *
      * @param[in]  input     Source tensor. Data types supported: F32.
-     * @param[in]  rois      Array containing the regions of interest.
+     * @param[in]  rois      Array containing @ref ROI.
      * @param[out] output    Destination tensor. Data types supported: Same as @p input.
      * @param[in]  pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo.
+     *
+     * @note The x and y dimensions of @p output tensor must be the same as that specified by @p pool_info 's pooled
+     * width and pooled height.
+     * @note The z dimensions of @p output tensor and @p input tensor must be the same.
+     * @note The fourth dimension of @p output tensor must be the same as the number of elements in @p rois array.
      */
     void configure(const ITensor *input, const IROIArray *rois, ITensor *output, const ROIPoolingLayerInfo &pool_info);
 
diff --git a/arm_compute/runtime/CL/CLArray.h b/arm_compute/runtime/CL/CLArray.h
index f4c2ef06d9..3dc7f19bc7 100644
--- a/arm_compute/runtime/CL/CLArray.h
+++ b/arm_compute/runtime/CL/CLArray.h
@@ -97,6 +97,7 @@ private:
 using CLKeyPointArray        = CLArray<KeyPoint>;
 using CLCoordinates2DArray   = CLArray<Coordinates2D>;
 using CLDetectionWindowArray = CLArray<DetectionWindow>;
+using CLROIArray             = CLArray<ROI>;
 using CLSize2DArray          = CLArray<Size2D>;
 using CLUInt8Array           = CLArray<cl_uchar>;
 using CLUInt16Array          = CLArray<cl_ushort>;
diff --git a/arm_compute/runtime/CL/functions/CLROIPoolingLayer.h b/arm_compute/runtime/CL/functions/CLROIPoolingLayer.h
new file mode 100644
index 0000000000..f089375e51
--- /dev/null
+++ b/arm_compute/runtime/CL/functions/CLROIPoolingLayer.h
@@ -0,0 +1,60 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_CLROIPOOLINGLAYER_H__
+#define __ARM_COMPUTE_CLROIPOOLINGLAYER_H__
+
+#include "arm_compute/runtime/CL/ICLSimpleFunction.h"
+
+#include "arm_compute/core/CL/ICLArray.h"
+#include "arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h"
+
+namespace arm_compute
+{
+class ICLTensor;
+
+/** Basic function to run @ref CLROIPoolingLayerKernel.
+ *
+ * This function calls the following OpenCL kernels:
+ * -# @ref CLROIPoolingLayerKernel
+ *
+ */
+class CLROIPoolingLayer : public ICLSimpleFunction
+{
+public:
+    /** Set the input and output tensors.
+     *
+     * @param[in]  input     Source tensor. Data types supported: F16/F32.
+     * @param[in]  rois      Array containing @ref ROI.
+     * @param[out] output    Destination tensor. Data types supported: Same as @p input.
+     * @param[in]  pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo.
+     *
+     * @note The x and y dimensions of @p output tensor must be the same as @p pool_info 's pooled
+     * width and pooled height.
+     * @note The z dimensions of @p output tensor and @p input tensor must be the same.
+     * @note The fourth dimension of @p output tensor must be the same as the number of elements in @p rois array.
+     */
+    void configure(const ICLTensor *input, const ICLROIArray *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info);
+};
+}
+#endif /* __ARM_COMPUTE_CLROIPOOLINGLAYER_H__ */
diff --git a/arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h b/arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h
index 04b5c35150..5adc1110d5 100644
--- a/arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h
+++ b/arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h
@@ -47,9 +47,14 @@ public:
     /** Set the input and output tensors.
      *
      * @param[in]  input     Source tensor. Data types supported: F32.
-     * @param[in]  rois      Array containing the regions of interest.
+     * @param[in]  rois      Array containing @ref ROI.
      * @param[out] output    Destination tensor. Data types supported: Same as @p input.
      * @param[in]  pool_info Contains pooling operation information described in @ref ROIPoolingLayerInfo.
+     *
+     * @note The x and y dimensions of @p output tensor must be the same as that specified by @p pool_info 's pooled
+     * width and pooled height.
+     * @note The z dimensions of @p output tensor and @p input tensor must be the same.
+     * @note The fourth dimension of @p output tensor must be the same as the number of elements in @p rois array.
      */
     void configure(const ITensor *input, const IROIArray *rois, ITensor *output, const ROIPoolingLayerInfo &pool_info);
 
diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp
index 2589bd12b5..ce2cfef67a 100644
--- a/src/core/CL/CLKernelLibrary.cpp
+++ b/src/core/CL/CLKernelLibrary.cpp
@@ -246,6 +246,7 @@ const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map =
     { "RGBA8888_to_NV12_bt709", "color_convert.cl" },
     { "RGBA8888_to_RGB888_bt709", "color_convert.cl" },
     { "RGBA8888_to_YUV444_bt709", "color_convert.cl" },
+    { "roi_pooling_layer", "roi_pooling_layer.cl" },
     { "scale_nearest_neighbour", "scale.cl" },
     { "scale_bilinear", "scale.cl" },
     { "scharr3x3", "scharr_filter.cl" },
@@ -477,6 +478,10 @@ const std::map<std::string, std::string> CLKernelLibrary::_program_source_map =
 #include "./cl_kernels/remap.clembed"
     },
     {
+        "roi_pooling_layer.cl",
+#include "./cl_kernels/roi_pooling_layer.clembed"
+    },
+    {
         "scale.cl",
 #include "./cl_kernels/scale.clembed"
     },
diff --git a/src/core/CL/ICLKernel.cpp b/src/core/CL/ICLKernel.cpp
index 12af8c68c1..7a95374bbf 100644
--- a/src/core/CL/ICLKernel.cpp
+++ b/src/core/CL/ICLKernel.cpp
@@ -70,12 +70,6 @@ cl::Kernel &ICLKernel::kernel()
 }
 
 template <unsigned int dimension_size>
-unsigned int           ICLKernel::num_arguments_per_tensor() const
-{
-    return 2 + 2 * dimension_size;
-}
-
-template <unsigned int dimension_size>
 void ICLKernel::add_tensor_argument(unsigned &idx, const ICLTensor *tensor, const Window &window)
 {
     ARM_COMPUTE_ERROR_ON(tensor == nullptr);
@@ -127,6 +121,11 @@ void ICLKernel::add_4D_tensor_argument(unsigned int &idx, const ICLTensor *tenso
     add_tensor_argument<4>(idx, tensor, window);
 }
 
+unsigned int ICLKernel::num_arguments_per_1D_array() const
+{
+    return num_arguments_per_array<1>();
+}
+
 unsigned int ICLKernel::num_arguments_per_1D_tensor() const
 {
     return num_arguments_per_tensor<1>();
diff --git a/src/core/CL/cl_kernels/roi_pooling_layer.cl b/src/core/CL/cl_kernels/roi_pooling_layer.cl
new file mode 100644
index 0000000000..35a9c0a21f
--- /dev/null
+++ b/src/core/CL/cl_kernels/roi_pooling_layer.cl
@@ -0,0 +1,169 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "helpers.h"
+
+#if DATA_SIZE == 32
+#define VEC_SIZE 4
+#define VEC_MAX vec4_max
+#elif DATA_SIZE == 16
+#define VEC_SIZE 8
+#define VEC_MAX vec8_max
+#else /* DATA_SIZE not equals 32 or 16 */
+#error "Unsupported data size"
+#endif /* DATA_SIZE == 32 */
+
+inline DATA_TYPE vec4_max(VEC_DATA_TYPE(DATA_TYPE, 4) vec)
+{
+    VEC_DATA_TYPE(DATA_TYPE, 2)
+    temp = fmax(vec.lo, vec.hi);
+    return fmax(temp.x, temp.y);
+}
+
+inline DATA_TYPE vec8_max(VEC_DATA_TYPE(DATA_TYPE, 8) vec)
+{
+    VEC_DATA_TYPE(DATA_TYPE, 4)
+    temp = fmax(vec.lo, vec.hi);
+    return vec4_max(temp);
+}
+
+/** Performs a roi pooling on a single output pixel.
+ *
+ * @param[in] input          Pointer to input Tensor3D struct.
+ * @param[in] region_start_x Start x index projected onto the input tensor.
+ * @param[in] region_end_x   End x index projected onto the input tensor.
+ * @param[in] region_start_y Start y index projected onto the input tensor.
+ * @param[in] region_end_y   End y index projected onto the input tensor.
+ * @param[in] pz             z index of the input tensor.
+ *
+ * @return A max pooled value from the region specified in the input tensor.
+ */
+inline DATA_TYPE roi_pool_1x1(const Tensor3D *input, int region_start_x, int region_end_x, int region_start_y, int region_end_y, int pz)
+{
+    // Iterate through the pooling region
+    if((region_end_x <= region_start_x) || (region_end_y <= region_start_y))
+    {
+        return (DATA_TYPE)0;
+    }
+    else
+    {
+        int num_iter = (int)((region_end_x - region_start_x) / VEC_SIZE);
+        VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
+        curr_max = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))(-FLT_MAX);
+        for(int j = region_start_y; j < region_end_y; ++j)
+        {
+            int i = region_start_x;
+            for(; i < region_start_x + num_iter * VEC_SIZE; i += VEC_SIZE)
+            {
+                VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
+                val      = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)tensor3D_offset(input, i, j, pz));
+                curr_max = fmax(val, curr_max);
+            }
+            for(; i < region_end_x; ++i)
+            {
+                DATA_TYPE val = *(__global DATA_TYPE *)tensor3D_offset(input, i, j, pz);
+                curr_max      = fmax(curr_max, val);
+            }
+        }
+        return (DATA_TYPE)VEC_MAX(curr_max);
+    }
+}
+
+/** Performs a roi pooling function.
+ *
+ * @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16, F32;
+ * @note Datasize must be passed using -DDATA_SIZE e.g. -DDATA_SIZE=32;
+ * @note Input dimensions must be passed using -DMAX_DIM_X, -DMAX_DIM_Y and -DMAX_DIM_Z;
+ * @note Pooled region dimensions must be passed using -DPOOLED_DIM_X and -DPOOLED_DIM_Y;
+ * @note Spatial scale must be passed using -DSPATIAL_SCALE;
+ *
+ * @param[in]  input_ptr                            Pointer to the source image. Supported data types: F16, F32
+ * @param[in]  input_stride_x                       Stride of the source image in X dimension (in bytes)
+ * @param[in]  input_step_x                         input_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  input_stride_y                       Stride of the source image in Y dimension (in bytes)
+ * @param[in]  input_step_y                         input_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  input_stride_z                       Stride of the source tensor in Z dimension (in bytes)
+ * @param[in]  input_step_z                         input_stride_z * number of elements along Z processed per workitem(in bytes)
+ * @param[in]  input_offset_first_element_in_bytes  The offset of the first element in the pooled region of the source image as specifed by ROI
+ * @param[in]  rois_ptr                             Pointer to the rois array. Layout: {x, y, width, height, batch_indx}
+ * @param[in]  rois_stride_x                        Stride of the rois array in X dimension (in bytes)
+ * @param[in]  rois_step_x                          rois_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  rois_offset_first_element_in_bytes   The offset of the first element in the rois array
+ * @param[out] output_ptr                           Pointer to the destination image. Supported data types: F16, F32
+ * @param[in]  output_stride_x                      Stride of the destination image in X dimension (in bytes)
+ * @param[in]  output_step_x                        output_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  output_stride_y                      Stride of the destination image in Y dimension (in bytes)
+ * @param[in]  output_step_y                        output_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  output_stride_z                      Stride of the destination tensor in Z dimension (in bytes)
+ * @param[in]  output_step_z                        output_stride_z * number of elements along Z processed per workitem(in bytes)
+ * @param[in]  output_offset_first_element_in_bytes The offset of the first element in the destination image
+ * @param[in]  input_stride_w                       Stride of the source image in W dimension (in bytes)
+ * @param[in]  output_stride_w                      Stride of the destination image in W dimension (in bytes)
+ */
+__kernel void roi_pooling_layer(
+    TENSOR3D_DECLARATION(input),
+    VECTOR_DECLARATION(rois),
+    TENSOR3D_DECLARATION(output),
+    unsigned int input_stride_w, unsigned int output_stride_w)
+{
+    // Get pixels pointer
+    Tensor3D input  = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);
+    Vector   rois   = CONVERT_TO_VECTOR_STRUCT_NO_STEP(rois);
+    Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);
+
+    const int px = get_global_id(0);
+    const int py = get_global_id(1);
+    const int pw = get_global_id(2);
+
+    // Load roi parameters
+    // roi is laid out as follows:
+    // { x, y, width, height, batch_index }
+    const ushort8 roi     = vload8(0, (__global ushort *)vector_offset(&rois, pw));
+    const int2 roi_anchor = convert_int2_sat(round(convert_float2(roi.s01) * (float)SPATIAL_SCALE));
+    const int2 roi_dims   = convert_int2_sat(fmax(round(convert_float2(roi.s23) * (float)SPATIAL_SCALE), 1.f));
+
+    // Determine pooled region in input image to pooled region in output image ratio
+    const float2 pool_region_ratio = convert_float2(roi_dims) / (float2)(POOLED_DIM_X, POOLED_DIM_Y);
+
+    // Calculate pooled region start and end
+    const float2 spatial_indx     = (float2)(px, py);
+    const int2   max_spatial_dims = (int2)(MAX_DIM_X, MAX_DIM_Y);
+    int2         region_start     = convert_int2_sat(floor(spatial_indx * pool_region_ratio)) + roi_anchor;
+    int2         region_end       = convert_int2_sat(ceil((spatial_indx + 1) * pool_region_ratio)) + roi_anchor;
+
+    region_start = clamp(region_start, 0, max_spatial_dims);
+    region_end   = clamp(region_end, 0, max_spatial_dims);
+
+    // Move input and output pointer across the fourth dimension
+    input.ptr += roi.s4 * input_stride_w;
+    output.ptr += pw * output_stride_w;
+
+    for(int pz = 0; pz < MAX_DIM_Z; ++pz)
+    {
+        *(__global DATA_TYPE *)tensor3D_offset(&output, px, py, pz) = (__global DATA_TYPE)roi_pool_1x1(&input,
+                                                                                                       region_start.x,
+                                                                                                       region_end.x,
+                                                                                                       region_start.y,
+                                                                                                       region_end.y, pz);
+    }
+}
diff --git a/src/core/CL/kernels/CLROIPoolingLayerKernel.cpp b/src/core/CL/kernels/CLROIPoolingLayerKernel.cpp
new file mode 100644
index 0000000000..4e000c61b1
--- /dev/null
+++ b/src/core/CL/kernels/CLROIPoolingLayerKernel.cpp
@@ -0,0 +1,121 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h"
+
+#include "arm_compute/core/AccessWindowStatic.h"
+#include "arm_compute/core/CL/CLHelpers.h"
+#include "arm_compute/core/CL/CLKernelLibrary.h"
+#include "arm_compute/core/CL/ICLArray.h"
+#include "arm_compute/core/CL/ICLTensor.h"
+#include "arm_compute/core/CL/OpenCL.h"
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/TensorInfo.h"
+#include "arm_compute/core/Utils.h"
+#include "arm_compute/core/Validate.h"
+#include "arm_compute/core/Window.h"
+
+#include <cmath>
+#include <set>
+#include <string>
+
+using namespace arm_compute;
+
+CLROIPoolingLayerKernel::CLROIPoolingLayerKernel()
+    : _input(nullptr), _rois(nullptr), _output(nullptr), _pool_info(0, 0, 0.f)
+{
+}
+
+void CLROIPoolingLayerKernel::configure(const ICLTensor *input, const ICLROIArray *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info)
+{
+    ARM_COMPUTE_ERROR_ON_NULLPTR(input, rois, output);
+    ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32);
+    ARM_COMPUTE_ERROR_ON((pool_info.pooled_width() == 0) || (pool_info.pooled_height() == 0));
+    ARM_COMPUTE_ERROR_ON(rois->num_values() == 0);
+
+    // Output auto inizialitation if not yet initialized
+    TensorShape output_shape(pool_info.pooled_width(), pool_info.pooled_height(), input->info()->dimension(2), rois->num_values());
+    auto_init_if_empty(*output->info(), output_shape, 1, input->info()->data_type(), input->info()->fixed_point_position());
+
+    ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output);
+    ARM_COMPUTE_ERROR_ON((output->info()->dimension(0) != pool_info.pooled_width()) || (output->info()->dimension(1) != pool_info.pooled_height()));
+    ARM_COMPUTE_ERROR_ON(input->info()->dimension(2) != output->info()->dimension(2));
+    ARM_COMPUTE_ERROR_ON(rois->num_values() != output->info()->dimension(3));
+
+    // Set instance variables
+    _input     = input;
+    _rois      = rois;
+    _output    = output;
+    _pool_info = pool_info;
+
+    // Set build options
+    std::set<std::string> build_opts;
+    build_opts.emplace(("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())));
+    build_opts.emplace(("-DDATA_SIZE=" + get_data_size_from_data_type(input->info()->data_type())));
+    build_opts.emplace(("-DMAX_DIM_X=" + support::cpp11::to_string(_input->info()->dimension(Window::DimX))));
+    build_opts.emplace(("-DMAX_DIM_Y=" + support::cpp11::to_string(_input->info()->dimension(Window::DimY))));
+    build_opts.emplace(("-DMAX_DIM_Z=" + support::cpp11::to_string(_input->info()->dimension(Window::DimZ))));
+    build_opts.emplace(("-DPOOLED_DIM_X=" + support::cpp11::to_string(pool_info.pooled_width())));
+    build_opts.emplace(("-DPOOLED_DIM_Y=" + support::cpp11::to_string(pool_info.pooled_height())));
+    build_opts.emplace(("-DSPATIAL_SCALE=" + support::cpp11::to_string(pool_info.spatial_scale())));
+
+    // Create kernel
+    std::string kernel_name = "roi_pooling_layer";
+    _kernel                 = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name, build_opts));
+
+    // Set static kernel arguments
+    unsigned int idx = 2 * num_arguments_per_3D_tensor() + num_arguments_per_1D_array();
+    add_argument<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(window);
+}
+
+void CLROIPoolingLayerKernel::run(const Window &window, cl::CommandQueue &queue)
+{
+    ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
+    ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window);
+
+    Window slice(window);
+    // Parallelize spatially and across the fourth dimension of the output tensor (also across ROIArray)
+    slice.set(Window::DimZ, window[3]);
+
+    // Set arguments
+    unsigned int idx = 0;
+    add_3D_tensor_argument(idx, _input, slice);
+    add_1D_array_argument<ROI>(idx, _rois, Strides(sizeof(ROI)), 1U, slice);
+    add_3D_tensor_argument(idx, _output, slice);
+    enqueue(queue, *this, slice);
+}
diff --git a/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp b/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp
index ec2b6a5079..fd6210021c 100644
--- a/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp
+++ b/src/core/NEON/kernels/NEROIPoolingLayerKernel.cpp
@@ -50,7 +50,7 @@ void NEROIPoolingLayerKernel::configure(const ITensor *input, const IROIArray *r
     ARM_COMPUTE_ERROR_ON(rois->num_values() == 0);
 
     // Output auto inizialitation if not yet initialized
-    TensorShape output_shape(pool_info.pooled_width(), pool_info.pooled_height(), rois->num_values());
+    TensorShape output_shape(pool_info.pooled_width(), pool_info.pooled_height(), input->info()->dimension(2), rois->num_values());
     auto_init_if_empty(*output->info(), output_shape, 1, input->info()->data_type(), input->info()->fixed_point_position());
 
     ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output);
diff --git a/src/runtime/CL/functions/CLROIPoolingLayer.cpp b/src/runtime/CL/functions/CLROIPoolingLayer.cpp
new file mode 100644
index 0000000000..0f480eeac9
--- /dev/null
+++ b/src/runtime/CL/functions/CLROIPoolingLayer.cpp
@@ -0,0 +1,39 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/runtime/CL/functions/CLROIPoolingLayer.h"
+
+#include "arm_compute/core/CL/ICLArray.h"
+
+#include "arm_compute/core/CL/kernels/CLROIPoolingLayerKernel.h"
+#include "support/ToolchainSupport.h"
+
+using namespace arm_compute;
+
+void CLROIPoolingLayer::configure(const ICLTensor *input, const ICLROIArray *rois, ICLTensor *output, const ROIPoolingLayerInfo &pool_info)
+{
+    // Configure ROI pooling kernel
+    auto k = arm_compute::support::cpp14::make_unique<CLROIPoolingLayerKernel>();
+    k->configure(input, rois, output, pool_info);
+    _kernel = std::move(k);
+}
diff --git a/tests/CL/CLArrayAccessor.h b/tests/CL/CLArrayAccessor.h
new file mode 100644
index 0000000000..a38907dd78
--- /dev/null
+++ b/tests/CL/CLArrayAccessor.h
@@ -0,0 +1,83 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_TEST_CLARRAYACCESSOR_H__
+#define __ARM_COMPUTE_TEST_CLARRAYACCESSOR_H__
+
+#include "arm_compute/runtime/CL/CLArray.h"
+#include "tests/IArrayAccessor.h"
+
+namespace arm_compute
+{
+namespace test
+{
+/** Accessor implementation for @ref CLArray objects. */
+template <typename T>
+class CLArrayAccessor : public IArrayAccessor<T>
+{
+public:
+    /** Create an accessor for the given @p array.
+     *
+     * @param[in, out] array To be accessed array.
+     *
+     * @note The CL memory is mapped by the constructor.
+     *
+     */
+    CLArrayAccessor(CLArray<T> &array)
+        : _array{ array }
+    {
+        _array.map();
+    }
+
+    CLArrayAccessor(const CLArrayAccessor &) = delete;
+    CLArrayAccessor &operator=(const CLArrayAccessor &) = delete;
+    CLArrayAccessor(CLArrayAccessor &&)                 = default;
+    CLArrayAccessor &operator=(CLArrayAccessor &&) = default;
+
+    /** Destructor that unmaps the CL memory. */
+    ~CLArrayAccessor()
+    {
+        _array.unmap();
+    }
+
+    int num_values() const override
+    {
+        return _array.num_values();
+    }
+
+    T *buffer() override
+    {
+        return _array.buffer();
+    }
+
+    void resize(size_t num) override
+    {
+        _array.resize(num);
+    }
+
+private:
+    CLArray<T> &_array;
+};
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_CLARRAYACCESSOR_H__ */
diff --git a/tests/IArrayAccessor.h b/tests/IArrayAccessor.h
new file mode 100644
index 0000000000..73f8829ed2
--- /dev/null
+++ b/tests/IArrayAccessor.h
@@ -0,0 +1,60 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_TEST_IARRAYACCESSOR_H__
+#define __ARM_COMPUTE_TEST_IARRAYACCESSOR_H__
+
+namespace arm_compute
+{
+namespace test
+{
+/** Common interface to provide information and access to array like
+ * structures.
+ */
+template <typename T>
+class IArrayAccessor
+{
+public:
+    using value_type = T;
+
+    /** Virtual destructor. */
+    virtual ~IArrayAccessor() = default;
+
+    /** Number of elements of the tensor. */
+    virtual int num_values() const = 0;
+
+    /** Access to the buffer.
+     *
+     * @return A pointer to the first element in the buffer.
+     */
+    virtual T *buffer() = 0;
+
+    /** Resize array.
+     *
+     *  @param[in] num The new array size in number of elements
+     */
+    virtual void resize(size_t num) = 0;
+};
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_IARRAYACCESSOR_H__ */
diff --git a/tests/NEON/ArrayAccessor.h b/tests/NEON/ArrayAccessor.h
new file mode 100644
index 0000000000..e56e6538ba
--- /dev/null
+++ b/tests/NEON/ArrayAccessor.h
@@ -0,0 +1,73 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_TEST_ARRAYACCESSOR_H__
+#define __ARM_COMPUTE_TEST_ARRAYACCESSOR_H__
+
+#include "arm_compute/runtime/Array.h"
+#include "tests/IArrayAccessor.h"
+
+namespace arm_compute
+{
+namespace test
+{
+/** ArrayAccessor implementation for @ref Array objects. */
+template <typename T>
+class ArrayAccessor : public IArrayAccessor<T>
+{
+public:
+    /** Create an accessor for the given @p array.
+     *
+     * @param[in, out] array To be accessed array.
+     */
+    ArrayAccessor(Array<T> &array)
+        : _array{ array }
+    {
+    }
+
+    ArrayAccessor(const ArrayAccessor &) = delete;
+    ArrayAccessor &operator=(const ArrayAccessor &) = delete;
+    ArrayAccessor(ArrayAccessor &&)                 = default;
+    ArrayAccessor &operator=(ArrayAccessor &&) = default;
+
+    int num_values() const override
+    {
+        return _array.num_values();
+    }
+
+    T *buffer() override
+    {
+        return _array.buffer();
+    }
+
+    void resize(size_t num) override
+    {
+        _array.resize(num);
+    }
+
+private:
+    Array<T> &_array;
+};
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_ARRAYACCESSOR_H__ */
diff --git a/tests/NEON/Helper.h b/tests/NEON/Helper.h
index 9e60a48ecd..4efab17fca 100644
--- a/tests/NEON/Helper.h
+++ b/tests/NEON/Helper.h
@@ -35,18 +35,6 @@ namespace arm_compute
 {
 namespace test
 {
-template <typename T>
-Array<T> create_array(const std::vector<T> &v)
-{
-    const size_t v_size = v.size();
-    Array<T>     array(v_size);
-
-    array.resize(v_size);
-    std::copy(v.begin(), v.end(), array.buffer());
-
-    return array;
-}
-
 template <typename D, typename T, typename... Ts>
 void fill_tensors(D &&dist, std::initializer_list<int> seeds, T &&tensor, Ts &&... other_tensors)
 {
diff --git a/tests/Utils.h b/tests/Utils.h
index 93831aa30b..f1725d7d74 100644
--- a/tests/Utils.h
+++ b/tests/Utils.h
@@ -37,9 +37,11 @@
 #include <cstddef>
 #include <limits>
 #include <memory>
+#include <random>
 #include <sstream>
 #include <string>
 #include <type_traits>
+#include <vector>
 
 namespace arm_compute
 {
@@ -421,6 +423,59 @@ inline T create_tensor(const TensorShape &shape, DataType data_type, int num_cha
 
     return tensor;
 }
+
+/** Create a vector of random ROIs.
+ *
+ * @param[in] shape     The shape of the input tensor.
+ * @param[in] pool_info The ROI pooling information.
+ * @param[in] num_rois  The number of ROIs to be created.
+ * @param[in] seed      The random seed to be used.
+ *
+ * @return A vector that contains the requested number of random ROIs
+ */
+inline std::vector<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;
+}
+
+template <typename T, typename ArrayAccessor_T>
+inline void fill_array(ArrayAccessor_T &&array, const std::vector<T> &v)
+{
+    array.resize(v.size());
+    std::memcpy(array.buffer(), v.data(), v.size() * sizeof(T));
+}
 } // namespace test
 } // namespace arm_compute
 #endif /* __ARM_COMPUTE_TEST_UTILS_H__ */
diff --git a/tests/benchmark_new/CL/ROIPoolingLayer.cpp b/tests/benchmark_new/CL/ROIPoolingLayer.cpp
new file mode 100644
index 0000000000..260f4d7b0d
--- /dev/null
+++ b/tests/benchmark_new/CL/ROIPoolingLayer.cpp
@@ -0,0 +1,53 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/CLArray.h"
+#include "arm_compute/runtime/CL/CLTensor.h"
+#include "arm_compute/runtime/CL/CLTensorAllocator.h"
+#include "arm_compute/runtime/CL/functions/CLROIPoolingLayer.h"
+#include "framework/Macros.h"
+#include "framework/datasets/Datasets.h"
+#include "tests/CL/CLAccessor.h"
+#include "tests/CL/CLArrayAccessor.h"
+#include "tests/TypePrinter.h"
+#include "tests/datasets_new/ROIPoolingLayerDataset.h"
+#include "tests/fixtures_new/ROIPoolingLayerFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+using CLROIPoolingLayerFixture = ROIPoolingLayerFixture<CLTensor, CLROIPoolingLayer, CLAccessor, CLArray<ROI>, CLArrayAccessor<ROI>>;
+
+TEST_SUITE(CL)
+
+REGISTER_FIXTURE_DATA_TEST_CASE(SmallROIPoolingLayer, CLROIPoolingLayerFixture, framework::DatasetMode::ALL,
+                                framework::dataset::combine(framework::dataset::combine(datasets::SmallROIPoolingLayerDataset(),
+                                                                                        framework::dataset::make("DataType", { DataType::F16, DataType::F32 })),
+                                                            framework::dataset::make("Batches", { 1, 4, 8 })));
+
+TEST_SUITE_END()
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/benchmark_new/NEON/ROIPoolingLayer.cpp b/tests/benchmark_new/NEON/ROIPoolingLayer.cpp
new file mode 100644
index 0000000000..e8c5ba3109
--- /dev/null
+++ b/tests/benchmark_new/NEON/ROIPoolingLayer.cpp
@@ -0,0 +1,53 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Array.h"
+#include "arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "framework/Macros.h"
+#include "framework/datasets/Datasets.h"
+#include "tests/NEON/Accessor.h"
+#include "tests/NEON/ArrayAccessor.h"
+#include "tests/TypePrinter.h"
+#include "tests/datasets_new/ROIPoolingLayerDataset.h"
+#include "tests/fixtures_new/ROIPoolingLayerFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+using NEROIPoolingLayerFixture = ROIPoolingLayerFixture<Tensor, NEROIPoolingLayer, Accessor, Array<ROI>, ArrayAccessor<ROI>>;
+
+TEST_SUITE(NEON)
+
+REGISTER_FIXTURE_DATA_TEST_CASE(SmallROIPoolingLayer, NEROIPoolingLayerFixture, framework::DatasetMode::ALL,
+                                framework::dataset::combine(framework::dataset::combine(datasets::SmallROIPoolingLayerDataset(),
+                                                                                        framework::dataset::make("DataType", { DataType::F32 })),
+                                                            framework::dataset::make("Batches", { 1, 4, 8 })));
+
+TEST_SUITE_END()
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/datasets_new/ROIPoolingLayerDataset.h b/tests/datasets_new/ROIPoolingLayerDataset.h
new file mode 100644
index 0000000000..c68a33fe8f
--- /dev/null
+++ b/tests/datasets_new/ROIPoolingLayerDataset.h
@@ -0,0 +1,127 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef ARM_COMPUTE_TEST_ROI_POOLING_LAYER_DATASET
+#define ARM_COMPUTE_TEST_ROI_POOLING_LAYER_DATASET
+
+#include "tests/TypePrinter.h"
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace datasets
+{
+class ROIPoolingLayerDataset
+{
+public:
+    using type = std::tuple<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, 3U), ROIPoolingLayerInfo(7U, 7U, 1.f / 8.f), 40U);
+        add_config(TensorShape(50U, 47U, 10U), ROIPoolingLayerInfo(7U, 7U, 1.f / 8.f), 80U);
+        add_config(TensorShape(50U, 47U, 80U), ROIPoolingLayerInfo(7U, 7U, 1.f / 8.f), 80U);
+        add_config(TensorShape(50U, 47U, 3U), ROIPoolingLayerInfo(9U, 9U, 1.f / 8.f), 40U);
+        add_config(TensorShape(50U, 47U, 10U), ROIPoolingLayerInfo(9U, 9U, 1.f / 8.f), 80U);
+        add_config(TensorShape(50U, 47U, 80U), ROIPoolingLayerInfo(9U, 9U, 1.f / 8.f), 80U);
+    }
+};
+
+} // namespace datasets
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_ROI_POOLING_LAYER_DATASET */
diff --git a/tests/fixtures_new/ROIPoolingLayerFixture.h b/tests/fixtures_new/ROIPoolingLayerFixture.h
new file mode 100644
index 0000000000..54dac51ee4
--- /dev/null
+++ b/tests/fixtures_new/ROIPoolingLayerFixture.h
@@ -0,0 +1,95 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef ARM_COMPUTE_TEST_ROIPOOLINGLAYERFIXTURE
+#define ARM_COMPUTE_TEST_ROIPOOLINGLAYERFIXTURE
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "framework/Fixture.h"
+#include "tests/Globals.h"
+#include "tests/Utils.h"
+
+#include <vector>
+
+namespace arm_compute
+{
+namespace test
+{
+/** Fixture that can be used for NEON and CL */
+template <typename TensorType, typename Function, typename Accessor, typename Array_T, typename ArrayAccessor>
+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)
+    {
+        // Set batched in source and destination shapes
+        const unsigned int fixed_point_position = 4;
+        TensorShape        shape_dst;
+        shape.set(shape.num_dimensions(), batches);
+        shape_dst.set(0, pool_info.pooled_width());
+        shape_dst.set(1, pool_info.pooled_height());
+        shape_dst.set(2, shape.z());
+        shape_dst.set(3, num_rois);
+
+        // Create tensors
+        src = create_tensor<TensorType>(shape, data_type, 1, fixed_point_position);
+        dst = create_tensor<TensorType>(shape_dst, data_type, 1, fixed_point_position);
+
+        // 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);
+
+        // Allocate tensors
+        src.allocator()->allocate();
+        dst.allocator()->allocate();
+
+        // Fill tensors
+        library->fill_tensor_uniform(Accessor(src), 0);
+    }
+
+    void run()
+    {
+        roi_pool.run();
+    }
+
+    void teardown()
+    {
+        src.allocator()->free();
+        dst.allocator()->free();
+    }
+
+private:
+    TensorType               src{};
+    TensorType               dst{};
+    std::unique_ptr<Array_T> rois_array{};
+    Function                 roi_pool{};
+};
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_ROIPOOLINGLAYERFIXTURE */
diff --git a/tests/validation/CL/CMakeLists.txt b/tests/validation/CL/CMakeLists.txt
index 67900b3edc..f4477f621f 100644
--- a/tests/validation/CL/CMakeLists.txt
+++ b/tests/validation/CL/CMakeLists.txt
@@ -27,15 +27,24 @@ set(arm_compute_test_validation_OPENCL_SOURCE_FILES
     ${CMAKE_SOURCE_DIR}/CL/CLAccessor.h
     ${CMAKE_CURRENT_SOURCE_DIR}/CLFixture.h
     ${CMAKE_CURRENT_SOURCE_DIR}/CLFixture.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/ActivationLayer.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/BitwiseAnd.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/Box3x3.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/ConvolutionLayer.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/DepthConvert.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/FillBorder.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/FixedPoint/FixedPoint_QS8.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/FullyConnectedLayer.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/Gaussian3x3.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/Gaussian5x5.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/GEMM.cpp
 `   ${CMAKE_CURRENT_SOURCE_DIR}/IntegralImage.cpp
+`   ${CMAKE_CURRENT_SOURCE_DIR}/NonLinearFilter.cpp
+`   ${CMAKE_CURRENT_SOURCE_DIR}/PoolingLayer.cpp
+`   ${CMAKE_CURRENT_SOURCE_DIR}/ROIPoolingLayer.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/Sobel3x3.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/Sobel5x5.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/SoftmaxLayer.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/Threshold.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/DirectConvolutionLayer.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/MeanStdDev.cpp
diff --git a/tests/validation/CL/ROIPoolingLayer.cpp b/tests/validation/CL/ROIPoolingLayer.cpp
new file mode 100644
index 0000000000..19d7903128
--- /dev/null
+++ b/tests/validation/CL/ROIPoolingLayer.cpp
@@ -0,0 +1,112 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "CL/CLAccessor.h"
+#include "CL/CLArrayAccessor.h"
+#include "TypePrinter.h"
+#include "arm_compute/runtime/CL/CLArray.h"
+#include "arm_compute/runtime/CL/functions/CLROIPoolingLayer.h"
+#include "tests/Globals.h"
+#include "tests/Utils.h"
+#include "validation/Datasets.h"
+#include "validation/Reference.h"
+#include "validation/Validation.h"
+#include "validation/ValidationUserConfiguration.h"
+
+#include <random>
+#include <vector>
+
+using namespace arm_compute;
+using namespace arm_compute::test;
+using namespace arm_compute::test::validation;
+
+namespace
+{
+CLTensor compute_roi_pooling_layer(const TensorShape &shape, DataType dt, const std::vector<ROI> &rois, ROIPoolingLayerInfo pool_info)
+{
+    TensorShape shape_dst;
+    shape_dst.set(0, pool_info.pooled_width());
+    shape_dst.set(1, pool_info.pooled_height());
+    shape_dst.set(2, shape.z());
+    shape_dst.set(3, rois.size());
+
+    // Create tensors
+    CLTensor src = create_tensor<CLTensor>(shape, dt);
+    CLTensor dst = create_tensor<CLTensor>(shape_dst, dt);
+
+    // Create ROI array
+    CLArray<ROI> rois_array(rois.size());
+    fill_array(CLArrayAccessor<ROI>(rois_array), rois);
+
+    // Create and configure function
+    CLROIPoolingLayer roi_pool;
+    roi_pool.configure(&src, &rois_array, &dst, pool_info);
+
+    // Allocate tensors
+    src.allocator()->allocate();
+    dst.allocator()->allocate();
+
+    BOOST_TEST(!src.info()->is_resizable());
+    BOOST_TEST(!dst.info()->is_resizable());
+
+    // Fill tensors
+    std::uniform_real_distribution<> distribution(-1, 1);
+    library->fill(CLAccessor(src), distribution, 0);
+
+    // Compute function
+    roi_pool.run();
+
+    return dst;
+}
+} // namespace
+
+#ifndef DOXYGEN_SKIP_THIS
+BOOST_AUTO_TEST_SUITE(CL)
+BOOST_AUTO_TEST_SUITE(ROIPoolingLayer)
+
+BOOST_AUTO_TEST_SUITE(Float)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, boost::unit_test::data::make({ DataType::F16, DataType::F32 }) * boost::unit_test::data::make({ 10, 20, 40 }) * boost::unit_test::data::make({ 7, 9 }) *
+                     boost::unit_test::data::make({ 1.f / 8.f, 1.f / 16.f }),
+                     dt, num_rois, roi_pool_size, roi_scale)
+{
+    TensorShape         shape(50U, 47U, 2U, 3U);
+    ROIPoolingLayerInfo pool_info(roi_pool_size, roi_pool_size, roi_scale);
+
+    // Construct ROI vector
+    std::vector<ROI> rois = generate_random_rois(shape, pool_info, num_rois, user_config.seed);
+
+    // Compute function
+    CLTensor dst = compute_roi_pooling_layer(shape, dt, rois, pool_info);
+
+    // Compute reference
+    RawTensor ref_dst = Reference::compute_reference_roi_pooling_layer(shape, dt, rois, pool_info);
+
+    // Validate output
+    validate(CLAccessor(dst), ref_dst);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
+BOOST_AUTO_TEST_SUITE_END()
+BOOST_AUTO_TEST_SUITE_END()
+#endif /* DOXYGEN_SKIP_THIS */
diff --git a/tests/validation/Helpers.cpp b/tests/validation/Helpers.cpp
deleted file mode 100644
index 747260ea99..0000000000
--- a/tests/validation/Helpers.cpp
+++ /dev/null
@@ -1,72 +0,0 @@
-/*
- * Copyright (c) 2017 ARM Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-#include "validation/Helpers.h"
-
-using namespace arm_compute::test;
-
-namespace arm_compute
-{
-namespace test
-{
-namespace validation
-{
-std::vector<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;
-}
-} // namespace validation
-} // namespace test
-} // namespace arm_compute
diff --git a/tests/validation/Helpers.h b/tests/validation/Helpers.h
index 09ffda8957..19a0c4105c 100644
--- a/tests/validation/Helpers.h
+++ b/tests/validation/Helpers.h
@@ -32,6 +32,7 @@
 #include "tests/validation/half.h"
 
 #include <array>
+#include <cstring>
 #include <random>
 #include <type_traits>
 #include <utility>
@@ -250,17 +251,6 @@ inline void fill_warp_matrix(std::array<float, SIZE> &matrix, int cols, int rows
     }
 }
 
-/** Create a vector of random ROIs.
- *
- * @param[in] shape     The shape of the input tensor.
- * @param[in] pool_info The ROI pooling information.
- * @param[in] num_rois  The number of ROIs to be created.
- * @param[in] seed      The random seed to be used.
- *
- * @return A vector that contains the requested number of random ROIs
- */
-std::vector<ROI> generate_random_rois(const TensorShape &shape, const ROIPoolingLayerInfo &pool_info, unsigned int num_rois, std::random_device::result_type seed);
-
 /** Helper function to fill the Lut random by a ILutAccessor.
  *
  * @param[in,out] table Accessor at the Lut.
@@ -277,7 +267,6 @@ void fill_lookuptable(T &&table)
         table[i] = distribution(generator);
     }
 }
-
 } // namespace validation
 } // namespace test
 } // namespace arm_compute
diff --git a/tests/validation/NEON/Gaussian5x5.cpp b/tests/validation/NEON/Gaussian5x5.cpp
index 522ececee9..7727340f66 100644
--- a/tests/validation/NEON/Gaussian5x5.cpp
+++ b/tests/validation/NEON/Gaussian5x5.cpp
@@ -41,8 +41,6 @@
 
 #include "boost_wrapper.h"
 
-#include <iostream>
-
 #include <random>
 #include <string>
 
diff --git a/tests/validation/NEON/ROIPoolingLayer.cpp b/tests/validation/NEON/ROIPoolingLayer.cpp
index 5775db7e03..523885d908 100644
--- a/tests/validation/NEON/ROIPoolingLayer.cpp
+++ b/tests/validation/NEON/ROIPoolingLayer.cpp
@@ -22,13 +22,12 @@
  * SOFTWARE.
  */
 #include "NEON/Accessor.h"
-#include "NEON/Helper.h"
+#include "NEON/ArrayAccessor.h"
 #include "TypePrinter.h"
 #include "arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h"
 #include "tests/Globals.h"
 #include "tests/Utils.h"
 #include "validation/Datasets.h"
-#include "validation/Helpers.h"
 #include "validation/Reference.h"
 #include "validation/Validation.h"
 #include "validation/ValidationUserConfiguration.h"
@@ -51,9 +50,12 @@ Tensor compute_roi_pooling_layer(const TensorShape &shape, DataType dt, const st
     shape_dst.set(3, rois.size());
 
     // Create tensors
-    Tensor     src        = create_tensor<Tensor>(shape, dt);
-    Tensor     dst        = create_tensor<Tensor>(shape_dst, dt);
-    Array<ROI> rois_array = create_array(rois);
+    Tensor src = create_tensor<Tensor>(shape, dt);
+    Tensor dst = create_tensor<Tensor>(shape_dst, dt);
+
+    // Create ROI array
+    Array<ROI> rois_array(rois.size());
+    fill_array(ArrayAccessor<ROI>(rois_array), rois);
 
     // Create and configure function
     NEROIPoolingLayer roi_pool;