COMPMID-1343: Add grouping support to CLCol2ImKernel

Change-Id: I5188a2163e7341f1915d98c21464fea13a9a7faf
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/143330
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Reviewed-by: Giorgio Arena <giorgio.arena@arm.com>

9 files changed, 577 insertions, 27 deletions

diff --git a/arm_compute/core/CL/kernels/CLCol2ImKernel.h b/arm_compute/core/CL/kernels/CLCol2ImKernel.h
index 94f21b1ebc..5c047ca091 100644
--- a/arm_compute/core/CL/kernels/CLCol2ImKernel.h
+++ b/arm_compute/core/CL/kernels/CLCol2ImKernel.h
@@ -70,18 +70,20 @@ public:
      * @param[out] output         The output tensor. 3 lower dimensions represent a single output [width, height, OFM],
      *                            while the rest represent batch of outputs. Data types supported: Same as @p input
      * @param[in]  convolved_dims Output convolved dimensions.
+     * @param[in]  num_groups     (Optional) Number of groups when performing a grouped convolution
      */
-    void configure(const ICLTensor *input, ICLTensor *output, std::pair<unsigned int, unsigned int> convolved_dims);
+    void configure(const ICLTensor *input, ICLTensor *output, std::pair<unsigned int, unsigned int> convolved_dims, unsigned int num_groups = 1);
     /** Static function to check if given info will lead to a valid configuration of @ref CLCol2ImKernel
      *
      * @param[in] input          The input tensor to convert. Data types supported: QASYMM8/F16/F32
      * @param[in] output         The output tensor. 3 lower dimensions represent a single output [width, height, OFM],
      *                           while the rest represent batch of outputs. Data types supported: Same as @p input
      * @param[in] convolved_dims Output convolved dimensions.
+     * @param[in] num_groups     (Optional) Number of groups when performing a grouped convolution
      *
      * @return a status
      */
-    static Status validate(const ITensorInfo *input, const ITensorInfo *output, std::pair<unsigned int, unsigned int> convolved_dims);
+    static Status validate(const ITensorInfo *input, const ITensorInfo *output, std::pair<unsigned int, unsigned int> convolved_dims, unsigned int num_groups = 1);
 
     // Inherited methods overridden:
     void run(const Window &window, cl::CommandQueue &queue) override;
diff --git a/arm_compute/core/utils/misc/ShapeCalculator.h b/arm_compute/core/utils/misc/ShapeCalculator.h
index fc6abf95f3..8a00c22306 100644
--- a/arm_compute/core/utils/misc/ShapeCalculator.h
+++ b/arm_compute/core/utils/misc/ShapeCalculator.h
@@ -140,13 +140,19 @@ inline TensorShape compute_reductionB_shape(const ITensorInfo &a)
 
     return shape_vector_sum_row;
 }
-inline TensorShape compute_col2im_shape(const ITensorInfo &input, std::pair<unsigned int, unsigned int> convolved_dims)
+inline TensorShape compute_col2im_shape(const ITensorInfo &input, std::pair<unsigned int, unsigned int> convolved_dims, unsigned int num_groups = 1)
 {
+    ARM_COMPUTE_ERROR_ON(num_groups == 0);
+    ARM_COMPUTE_ERROR_ON(input.tensor_shape()[1] != (convolved_dims.first * convolved_dims.second));
+    ARM_COMPUTE_ERROR_ON((num_groups > 1) && input.tensor_shape()[2] != num_groups);
+
     TensorShape col2im_shape{ input.tensor_shape() };
-    col2im_shape.shift_right(1);
     col2im_shape.set(0, convolved_dims.first);
     col2im_shape.set(1, convolved_dims.second);
-    col2im_shape.set(2, input.tensor_shape()[0]);
+    col2im_shape.set(2, input.tensor_shape()[0] * num_groups);
+
+    const unsigned int batch_idx = (num_groups == 1) ? 2 : 3;
+    col2im_shape.set(3, input.tensor_shape()[batch_idx]);
 
     return col2im_shape;
 }
diff --git a/src/core/CL/cl_kernels/col2im.cl b/src/core/CL/cl_kernels/col2im.cl
index 98bf8d1ed4..5e52127f27 100644
--- a/src/core/CL/cl_kernels/col2im.cl
+++ b/src/core/CL/cl_kernels/col2im.cl
@@ -41,12 +41,15 @@
  * @note The width of the input tensor must be passed at compile time using -DWIDTH_INPUT: e.g. -DWIDTH_INPUT=320
  * @note The width of the output tensor must be passed at compile time using -DWIDTH_OUTPUT: e.g. -DWIDTH_OUTPUT=600
  * @note The element size must be passed at compile time using -DELEMENT_SIZE: e.g. -DELEMENT_SIZE=4
+ * @note In case of grouping the GROUPING flag must be passed at compile time using -DGROUPING
  *
  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: QASYMM8/F16/F32
  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
  * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  src_stride_z                      Stride of the source tensor in Z dimension (in bytes)
+ * @param[in]  src_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
  * @param[out] dst_ptr                           Pointer to the destination tensor. Supported data types: same as @p src_ptr
  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
@@ -59,11 +62,14 @@
  * @param[in]  dst_stride_w                      Stride of the destination tensor in W dimension (in bytes)
  */
 __kernel void col2im(
-    IMAGE_DECLARATION(src),
+    TENSOR3D_DECLARATION(src),
     TENSOR3D_DECLARATION(dst),
     uint dst_stride_w)
 {
-    Image src = CONVERT_TO_IMAGE_STRUCT(src);
+    Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
+
+    const uint xd = get_global_id(1) % WIDTH_OUTPUT; // x coordinate of the destination tensor
+    const uint yd = get_global_id(1) / WIDTH_OUTPUT; // y coordinate of the destination tensor
 
     VEC_DATA_TYPE(DATA_TYPE, 8)
     data = vload8(0, (__global DATA_TYPE *)src.ptr);
@@ -82,8 +88,16 @@ __kernel void col2im(
 
     __global uchar *output_ptr = dst_ptr + dst_offset_first_element_in_bytes;
 
-    // Compute output offset
-    int idx = (get_global_id(1) / WIDTH_OUTPUT) * dst_stride_y + (get_global_id(1) % WIDTH_OUTPUT) * dst_stride_x + get_global_id(2) * dst_stride_w;
+#if defined(GROUPING)
+    // Compute output offset (batches on 4th dimension, no need to compute manually)
+    int idx = yd * dst_stride_y + xd * dst_stride_x;
+
+    const uint group = get_global_id(2); // group ID
+    x_clamped += group * WIDTH_INPUT;
+#else  /* defined(GROUPING) */
+    // Compute output offset (batches on 3rd dimension)
+    int idx = yd * dst_stride_y + xd * dst_stride_x + get_global_id(2) * dst_stride_w;
+#endif /* GROUPING */
 
     // Store value
     *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s0 * dst_stride_z)) = data.s0;
@@ -95,4 +109,4 @@ __kernel void col2im(
     *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s6 * dst_stride_z)) = data.s6;
     *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s7 * dst_stride_z)) = data.s7;
 }
-#endif // defined(DATA_TYPE) && defined(WIDTH_OUTPUT) && defined(ELEMENT_SIZE) && defined(WIDTH_INPUT)
\ No newline at end of file
+#endif // defined(DATA_TYPE) && defined(WIDTH_OUTPUT) && defined(ELEMENT_SIZE) && defined(WIDTH_INPUT)
diff --git a/src/core/CL/kernels/CLCol2ImKernel.cpp b/src/core/CL/kernels/CLCol2ImKernel.cpp
index 6fd3be7f6a..d7582dc943 100644
--- a/src/core/CL/kernels/CLCol2ImKernel.cpp
+++ b/src/core/CL/kernels/CLCol2ImKernel.cpp
@@ -40,7 +40,7 @@ using namespace arm_compute::misc::shape_calculator;
 
 namespace
 {
-Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, std::pair<unsigned int, unsigned int> convolved_dims)
+Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, std::pair<unsigned int, unsigned int> convolved_dims, unsigned int num_groups)
 {
     ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, output);
     ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(input);
@@ -49,19 +49,20 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, s
     // Checks performed when output is configured
     if(output->total_size() != 0)
     {
-        ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(output->tensor_shape(), compute_col2im_shape(*input, convolved_dims));
+        ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(output->tensor_shape(), compute_col2im_shape(*input, convolved_dims, num_groups));
         ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output);
         ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_QUANTIZATION_INFO(input, output);
+        ARM_COMPUTE_RETURN_ERROR_ON_MSG(output->data_layout() != DataLayout::NCHW, "Col2Im output's data layout must always be NCHW");
     }
 
     return Status{};
 }
 
-std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITensorInfo *output, std::pair<unsigned int, unsigned int> convolved_dims)
+std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITensorInfo *output, std::pair<unsigned int, unsigned int> convolved_dims, unsigned int num_groups)
 {
     ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
     // Output auto inizialitation if not yet initialized
-    auto_init_if_empty(*output, input->clone()->set_tensor_shape(compute_col2im_shape(*input, convolved_dims)));
+    auto_init_if_empty(*output, input->clone()->set_tensor_shape(compute_col2im_shape(*input, convolved_dims, num_groups)).set_data_layout(DataLayout::NCHW));
 
     const unsigned int num_elems_read_per_iteration = 8;
 
@@ -86,12 +87,12 @@ CLCol2ImKernel::CLCol2ImKernel()
 {
 }
 
-void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::pair<unsigned int, unsigned int> convolved_dims)
+void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::pair<unsigned int, unsigned int> convolved_dims, unsigned int num_groups)
 {
     ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
 
     // Perform validation step
-    ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), convolved_dims));
+    ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), output->info(), convolved_dims, num_groups));
 
     _input          = input;
     _output         = output;
@@ -105,11 +106,12 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p
     build_opts.add_option("-DELEMENT_SIZE=" + support::cpp11::to_string(input->info()->element_size()));
     build_opts.add_option("-DWIDTH_INPUT=" + support::cpp11::to_string(input->info()->dimension(0)));
     build_opts.add_option("-DWIDTH_OUTPUT=" + support::cpp11::to_string(_convolved_dims.first));
+    build_opts.add_option_if(num_groups > 1, "-DGROUPING");
 
     _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("col2im", build_opts.options()));
 
     // Configure kernel window
-    auto win_config = validate_and_configure_window(input->info(), output->info(), _convolved_dims);
+    auto win_config = validate_and_configure_window(input->info(), output->info(), _convolved_dims, num_groups);
     ARM_COMPUTE_ERROR_THROW_ON(win_config.first);
     ICLKernel::configure_internal(win_config.second);
 
@@ -117,6 +119,7 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p
     _config_id = "col2im_";
     _config_id += lower_string(string_from_data_type(input->info()->data_type()));
     _config_id += "_";
+    _config_id += (num_groups > 1) ? "grouping_" : "";
     _config_id += support::cpp11::to_string(input->info()->dimension(0));
     _config_id += "_";
     _config_id += support::cpp11::to_string(input->info()->dimension(1));
@@ -126,11 +129,11 @@ void CLCol2ImKernel::configure(const ICLTensor *input, ICLTensor *output, std::p
     _config_id += support::cpp11::to_string(output->info()->dimension(1));
 }
 
-Status CLCol2ImKernel::validate(const ITensorInfo *input, const ITensorInfo *output, std::pair<unsigned int, unsigned int> convolved_dims)
+Status CLCol2ImKernel::validate(const ITensorInfo *input, const ITensorInfo *output, std::pair<unsigned int, unsigned int> convolved_dims, unsigned int num_groups)
 {
     ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
-    ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, convolved_dims));
-    ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), output->clone().get(), convolved_dims).first);
+    ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, convolved_dims, num_groups));
+    ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), output->clone().get(), convolved_dims, num_groups).first);
     return Status{};
 }
 
@@ -142,21 +145,19 @@ void CLCol2ImKernel::run(const Window &window, cl::CommandQueue &queue)
     Window out_window;
     out_window.use_tensor_dimensions(_output->info()->tensor_shape());
 
-    Window collapsed_window = window.collapse_if_possible(ICLKernel::window(), Window::DimZ);
-    Window slice            = collapsed_window.first_slice_window_2D();
-    Window slice_out        = out_window.first_slice_window_3D();
+    Window slice     = window.first_slice_window_3D();
+    Window slice_out = out_window.first_slice_window_3D();
 
-    // Set static kernel arguments
-    unsigned int idx = num_arguments_per_2D_tensor() + num_arguments_per_3D_tensor();
+    unsigned int idx = 2 * num_arguments_per_3D_tensor();
     _kernel.setArg<cl_uint>(idx++, _output->info()->strides_in_bytes()[3]);
 
     do
     {
         // Set inputs
         unsigned int idx = 0;
-        add_2D_tensor_argument(idx, _input, slice);
+        add_3D_tensor_argument(idx, _input, slice);
         add_3D_tensor_argument(idx, _output, slice_out);
         enqueue(queue, *this, slice, lws_hint());
     }
-    while(collapsed_window.slide_window_slice_2D(slice) && out_window.slide_window_slice_3D(slice_out));
+    while(window.slide_window_slice_3D(slice) && out_window.slide_window_slice_3D(slice_out));
 }
diff --git a/tests/datasets/Col2ImLayerDataset.h b/tests/datasets/Col2ImLayerDataset.h
new file mode 100644
index 0000000000..96a3cab134
--- /dev/null
+++ b/tests/datasets/Col2ImLayerDataset.h
@@ -0,0 +1,155 @@
+/*
+ * Copyright (c) 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_COL2IM_DATASET
+#define ARM_COMPUTE_TEST_COL2IM_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 Col2ImLayerDataset
+{
+public:
+    using type = std::tuple<TensorShape, unsigned int, unsigned int, unsigned int>;
+
+    struct iterator
+    {
+        iterator(std::vector<TensorShape>::const_iterator  src_it,
+                 std::vector<unsigned int>::const_iterator convolved_width_it,
+                 std::vector<unsigned int>::const_iterator convolved_height_it,
+                 std::vector<unsigned int>::const_iterator num_groups_it)
+            : _src_it{ std::move(src_it) },
+              _convolved_width_it{ std::move(convolved_width_it) },
+              _convolved_height_it{ std::move(convolved_height_it) },
+              _num_groups_it{ std::move(num_groups_it) }
+        {
+        }
+
+        std::string description() const
+        {
+            std::stringstream description;
+            description << "In=" << *_src_it << ":";
+            description << "ConvolvedWidth=" << *_convolved_width_it << ":";
+            description << "ConvolvedHeight=" << *_convolved_height_it << ":";
+            description << "NumGroups=" << *_num_groups_it;
+            return description.str();
+        }
+
+        Col2ImLayerDataset::type operator*() const
+        {
+            return std::make_tuple(*_src_it, *_convolved_width_it, *_convolved_height_it, *_num_groups_it);
+        }
+
+        iterator &operator++()
+        {
+            ++_src_it;
+            ++_convolved_width_it;
+            ++_convolved_height_it;
+            ++_num_groups_it;
+
+            return *this;
+        }
+
+    private:
+        std::vector<TensorShape>::const_iterator  _src_it;
+        std::vector<unsigned int>::const_iterator _convolved_width_it;
+        std::vector<unsigned int>::const_iterator _convolved_height_it;
+        std::vector<unsigned int>::const_iterator _num_groups_it;
+    };
+
+    iterator begin() const
+    {
+        return iterator(_src_shapes.begin(), _convolved_widths.begin(), _convolved_heights.begin(), _num_groups.begin());
+    }
+
+    int size() const
+    {
+        return std::min(_src_shapes.size(), std::min(_convolved_widths.size(), std::min(_convolved_heights.size(), _num_groups.size())));
+    }
+
+    void add_config(TensorShape src, unsigned int convolved_width, unsigned int convolved_height, unsigned int info)
+    {
+        _src_shapes.emplace_back(std::move(src));
+        _convolved_widths.emplace_back(std::move(convolved_width));
+        _convolved_heights.emplace_back(std::move(convolved_height));
+        _num_groups.emplace_back(std::move(info));
+    }
+
+protected:
+    Col2ImLayerDataset()                      = default;
+    Col2ImLayerDataset(Col2ImLayerDataset &&) = default;
+
+private:
+    std::vector<TensorShape>  _src_shapes{};
+    std::vector<unsigned int> _convolved_widths{};
+    std::vector<unsigned int> _convolved_heights{};
+    std::vector<unsigned int> _num_groups{};
+};
+
+/** Dataset containing small grouped col2im shapes. */
+class SmallGroupedCol2ImLayerDataset final : public Col2ImLayerDataset
+{
+public:
+    SmallGroupedCol2ImLayerDataset()
+    {
+        add_config(TensorShape(10U, 12U, 1U, 1U), 3U, 4U, 1U);
+        add_config(TensorShape(12U, 30U, 1U, 2U), 5U, 6U, 1U);
+        add_config(TensorShape(12U, 30U, 4U, 1U), 5U, 6U, 1U);
+        add_config(TensorShape(10U, 12U, 2U, 4U), 3U, 4U, 2U);
+        add_config(TensorShape(10U, 12U, 2U, 4U), 3U, 4U, 2U);
+        add_config(TensorShape(8U, 16U, 3U, 1U), 4U, 4U, 3U);
+        add_config(TensorShape(8U, 16U, 3U, 3U), 4U, 4U, 3U);
+        add_config(TensorShape(12U, 20U, 4U, 1U), 5U, 4U, 4U);
+        add_config(TensorShape(12U, 20U, 4U, 3U), 5U, 4U, 4U);
+    }
+};
+
+/** Dataset containing large grouped col2im shapes. */
+class LargeGroupedCol2ImLayerDataset final : public Col2ImLayerDataset
+{
+public:
+    LargeGroupedCol2ImLayerDataset()
+    {
+        add_config(TensorShape(233U, 280U, 1U, 55U), 14U, 20U, 1U);
+        add_config(TensorShape(333U, 280U, 1U, 77U), 14U, 20U, 1U);
+        add_config(TensorShape(333U, 280U, 77U, 1U), 14U, 20U, 1U);
+        add_config(TensorShape(120U, 300U, 8U, 3U), 20U, 15U, 8U);
+        add_config(TensorShape(233U, 300U, 8U, 3U), 20U, 15U, 8U);
+        add_config(TensorShape(333U, 280U, 12U, 5U), 20U, 14U, 12U);
+        add_config(TensorShape(177U, 300U, 12U, 5U), 15U, 20U, 12U);
+        add_config(TensorShape(450U, 400U, 16U, 5U), 20U, 20U, 16U);
+        add_config(TensorShape(220U, 400U, 16U, 5U), 20U, 20U, 16U);
+    }
+};
+} // namespace datasets
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_COL2IM_DATASET */
diff --git a/tests/validation/CL/Col2Im.cpp b/tests/validation/CL/Col2Im.cpp
new file mode 100644
index 0000000000..6f1163c278
--- /dev/null
+++ b/tests/validation/CL/Col2Im.cpp
@@ -0,0 +1,126 @@
+/*
+ * Copyright (c) 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.
+ */
+#include "arm_compute/core/CL/kernels/CLCol2ImKernel.h"
+#include "arm_compute/core/Types.h"
+#include "tests/CL/Helper.h"
+
+#include "tests/CL/CLAccessor.h"
+#include "tests/datasets/Col2ImLayerDataset.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/Col2ImFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+TEST_SUITE(CL)
+TEST_SUITE(Col2Im)
+
+using CLCol2Im = CLSynthetizeFunction<CLCol2ImKernel>;
+
+// *INDENT-OFF*
+// clang-format off
+DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip(zip(
+               framework::dataset::make("InputInfo", { TensorInfo(TensorShape(10U, 12U, 1U, 2U), 1, DataType::S64),    // Unsupported data type
+                                                       TensorInfo(TensorShape(10U, 12U, 1U, 2U), 1, DataType::F32),    // Mismatching data type
+                                                       TensorInfo(TensorShape(10U, 12U, 1U, 2U), 1, DataType::F32),    // Invalid output shape
+                                                       TensorInfo(TensorShape(3U, 12U, 4U, 2U), 1, DataType::F32),
+                                                     }),
+               framework::dataset::make("OutputInfo",{ TensorInfo(TensorShape(3U, 4U, 10U, 2U), 1, DataType::F16),
+                                                       TensorInfo(TensorShape(3U, 4U, 10U, 2U), 1, DataType::F16),
+                                                       TensorInfo(TensorShape(3U, 3U, 10U, 2U), 1, DataType::F32),
+                                                       TensorInfo(TensorShape(3U, 4U, 12U, 2U), 1, DataType::F32),
+                                                     })),
+               framework::dataset::make("ConvolvedWidth", { 3, 3, 3, 3 })),
+               framework::dataset::make("ConvolvedHeight", { 4, 4, 4, 4 })),
+               framework::dataset::make("NumGroups", { 1, 1, 1, 4 })),
+               framework::dataset::make("Expected", { false, false, false, true })),
+               input_info, output_info, convolved_width, convolved_height, num_groups, expected)
+{
+    bool status = bool(CLCol2Im::validate(&input_info, &output_info, std::make_pair(convolved_width, convolved_height), num_groups));
+    ARM_COMPUTE_EXPECT(status == expected, framework::LogLevel::ERRORS);
+}
+// clang-format on
+// *INDENT-ON*
+
+template <typename T>
+using CLCol2ImFixture = Col2ImValidationFixture<CLTensor, CLAccessor, CLCol2Im, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLCol2ImFixture<float>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallGroupedCol2ImLayerDataset(), framework::dataset::make("DataType", DataType::F32)))
+{
+    // Validate output
+    validate(CLAccessor(_target), _reference);
+}
+
+FIXTURE_DATA_TEST_CASE(RunLarge, CLCol2ImFixture<float>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGroupedCol2ImLayerDataset(), framework::dataset::make("DataType", DataType::F32)))
+{
+    // Validate output
+    validate(CLAccessor(_target), _reference);
+}
+TEST_SUITE_END()
+
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLCol2ImFixture<half>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallGroupedCol2ImLayerDataset(), framework::dataset::make("DataType", DataType::F16)))
+{
+    // Validate output
+    validate(CLAccessor(_target), _reference);
+}
+
+FIXTURE_DATA_TEST_CASE(RunLarge, CLCol2ImFixture<half>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGroupedCol2ImLayerDataset(), framework::dataset::make("DataType", DataType::F16)))
+{
+    // Validate output
+    validate(CLAccessor(_target), _reference);
+}
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+
+TEST_SUITE(QASYMM8)
+FIXTURE_DATA_TEST_CASE(RunSmall, CLCol2ImFixture<uint8_t>, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallGroupedCol2ImLayerDataset(), framework::dataset::make("DataType",
+                                                                                                      DataType::QASYMM8)))
+{
+    // Validate output
+    validate(CLAccessor(_target), _reference);
+}
+
+FIXTURE_DATA_TEST_CASE(RunLarge, CLCol2ImFixture<uint8_t>, framework::DatasetMode::NIGHTLY, combine(datasets::LargeGroupedCol2ImLayerDataset(), framework::dataset::make("DataType",
+                                                                                                    DataType::QASYMM8)))
+{
+    // Validate output
+    validate(CLAccessor(_target), _reference);
+}
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/fixtures/Col2ImFixture.h b/tests/validation/fixtures/Col2ImFixture.h
new file mode 100644
index 0000000000..ddc78a5032
--- /dev/null
+++ b/tests/validation/fixtures/Col2ImFixture.h
@@ -0,0 +1,114 @@
+/*
+ * Copyright (c) 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_COL2IM_FIXTURE
+#define ARM_COMPUTE_TEST_COL2IM_FIXTURE
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/core/utils/misc/ShapeCalculator.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/framework/Asserts.h"
+#include "tests/framework/Fixture.h"
+#include "tests/validation/reference/Col2Im.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+using namespace arm_compute::misc::shape_calculator;
+
+template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
+class Col2ImValidationFixture : public framework::Fixture
+{
+public:
+    template <typename...>
+    void setup(TensorShape input_shape, const unsigned int convolved_width, unsigned int convolved_height, unsigned int num_groups, DataType data_type)
+    {
+        const std::pair<unsigned int, unsigned int> convolved_dims(convolved_width, convolved_height);
+
+        const TensorShape output_shape = compute_col2im_shape(TensorInfo(input_shape, 1, data_type), convolved_dims, num_groups);
+
+        _target    = compute_target(input_shape, output_shape, convolved_dims, num_groups, data_type);
+        _reference = compute_reference(input_shape, output_shape, num_groups, data_type);
+    }
+
+protected:
+    template <typename U>
+    void fill(U &&tensor, const int seed)
+    {
+        library->fill_tensor_uniform(tensor, seed);
+    }
+
+    TensorType compute_target(const TensorShape &input_shape, const TensorShape &output_shape, std::pair<unsigned int, unsigned int> convolved_dims, unsigned int num_groups, DataType data_type)
+    {
+        // Create tensors
+        TensorType src = create_tensor<TensorType>(input_shape, data_type);
+        TensorType dst = create_tensor<TensorType>(output_shape, data_type);
+
+        // Create and configure function
+        FunctionType col2im_func;
+        col2im_func.configure(&src, &dst, convolved_dims, num_groups);
+
+        ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+        ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+        // Allocate tensors
+        src.allocator()->allocate();
+        dst.allocator()->allocate();
+
+        ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS);
+        ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+        // Fill tensors
+        fill(AccessorType(src), 0);
+
+        // Compute function
+        col2im_func.run();
+
+        return dst;
+    }
+
+    SimpleTensor<T> compute_reference(const TensorShape &input_shape, const TensorShape &output_shape, unsigned int num_groups, DataType data_type)
+    {
+        // Create reference
+        SimpleTensor<T> src{ input_shape, data_type };
+
+        // Fill reference
+        fill(src, 0);
+
+        return reference::col2im<T>(src, output_shape, num_groups);
+    }
+    TensorType      _target{};
+    SimpleTensor<T> _reference{};
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_COL2IM_FIXTURE */
diff --git a/tests/validation/reference/Col2Im.cpp b/tests/validation/reference/Col2Im.cpp
new file mode 100644
index 0000000000..90e488f928
--- /dev/null
+++ b/tests/validation/reference/Col2Im.cpp
@@ -0,0 +1,88 @@
+/*
+ * Copyright (c) 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.
+ */
+#include "Col2Im.h"
+
+#include "tests/validation/Helpers.h"
+#include "tests/validation/reference/Utils.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> col2im(const SimpleTensor<T> &src, const TensorShape &dst_shape, unsigned int num_groups)
+{
+    SimpleTensor<T> dst{ dst_shape, src.data_type(), 1 };
+
+    // Compute reference
+    const size_t batches    = dst_shape[3];
+    const size_t src_width  = src.shape().x();
+    const size_t src_height = src.shape().y();
+
+    if(num_groups == 1)
+    {
+        // Batches are on the 3rd dimension of the input tensor
+        int dst_idx = 0;
+        for(size_t b = 0; b < batches; ++b)
+        {
+            for(size_t x = 0; x < src_width; ++x)
+            {
+                for(size_t y = 0; y < src_height; ++y)
+                {
+                    dst[dst_idx++] = src[coord2index(src.shape(), Coordinates(x, y, b))];
+                }
+            }
+        }
+    }
+    else
+    {
+        int dst_idx = 0;
+        for(size_t b = 0; b < batches; ++b)
+        {
+            for(size_t g = 0; g < num_groups; ++g)
+            {
+                for(size_t x = 0; x < src_width; ++x)
+                {
+                    for(size_t y = 0; y < src_height; ++y)
+                    {
+                        dst[dst_idx++] = src[coord2index(src.shape(), Coordinates(x, y, g, b))];
+                    }
+                }
+            }
+        }
+    }
+    return dst;
+}
+
+template SimpleTensor<float> col2im(const SimpleTensor<float> &src, const TensorShape &dst_shape, unsigned int num_groups);
+template SimpleTensor<half> col2im(const SimpleTensor<half> &src, const TensorShape &dst_shape, unsigned int num_groups);
+template SimpleTensor<uint8_t> col2im(const SimpleTensor<uint8_t> &src, const TensorShape &dst_shape, unsigned int num_groups);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation/reference/Col2Im.h b/tests/validation/reference/Col2Im.h
new file mode 100644
index 0000000000..608261035d
--- /dev/null
+++ b/tests/validation/reference/Col2Im.h
@@ -0,0 +1,44 @@
+/*
+ * Copyright (c) 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_COL2IM_H__
+#define __ARM_COMPUTE_TEST_COL2IM_H__
+
+#include "tests/SimpleTensor.h"
+#include "tests/validation/Helpers.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> col2im(const SimpleTensor<T> &src, const TensorShape &dst_shape, unsigned int num_groups);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_COL2IM_H__ */