diff --git a/src/core/CL/cl_kernels/reduction_operation.cl b/src/core/CL/cl_kernels/reduction_operation.cl
index d46a226..aa7403b 100644
--- a/src/core/CL/cl_kernels/reduction_operation.cl
+++ b/src/core/CL/cl_kernels/reduction_operation.cl
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016, 2017 ARM Limited.
+ * Copyright (c) 2016-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -70,39 +70,46 @@
  * @param[in] src_ptr                                   Pointer to the source tensor. Supported data types: 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_offset_first_element_in_bytes         The offset of the first element in the source tensor
  * @param[in] partial_sum_ptr                           The local buffer to hold sumed values. Supported data types: same as @p src_ptt
- * @param[in] partial_sum_stride_x                      Stride of the source tensor in X dimension (in bytes)
+ * @param[in] partial_sum_stride_x                      Stride of the output tensor in X dimension (in bytes)
  * @param[in] partial_sum_step_x                        partial_sum_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] partial_sum_stride_y                      Stride of the output tensor in Y dimension (in bytes)
+ * @param[in] partial_sum_step_y                        partial_sum_stride_y * number of elements along Y processed per workitem(in bytes)
  * @param[in] partial_sum_offset_first_element_in_bytes The offset of the first element in the source tensor
- * @param[in] local_sums                                Local buffer for storing the partioal sum
+ * @param[in] local_sums                                Local buffer for storing the partial sum
  */
 __kernel void reduction_operation(
-    VECTOR_DECLARATION(src),
-    VECTOR_DECLARATION(partial_sum),
+    IMAGE_DECLARATION(src),
+    IMAGE_DECLARATION(partial_sum),
     __local DATA_TYPE *local_sums)
 {
-    Vector src         = CONVERT_TO_VECTOR_STRUCT(src);
-    Vector partial_sum = CONVERT_TO_VECTOR_STRUCT(partial_sum);
+    Image src         = CONVERT_TO_IMAGE_STRUCT(src);
+    Image partial_sum = CONVERT_TO_IMAGE_STRUCT(partial_sum);
 
     unsigned int lsize = get_local_size(0);
     unsigned int lid   = get_local_id(0);
 
-    local_sums[lid] = OPERATION((__global DATA_TYPE *)src.ptr);
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    // Perform parallel reduction
-    for(unsigned int i = lsize >> 1; i > 0; i >>= 1)
+    for(unsigned int y = 0; y < get_local_size(1); ++y)
     {
-        if(lid < i)
-        {
-            local_sums[lid] += local_sums[lid + i];
-        }
+        local_sums[lid] = OPERATION((__global DATA_TYPE *)offset(&src, 0, y));
         barrier(CLK_LOCAL_MEM_FENCE);
-    }
 
-    if(lid == 0)
-    {
-        ((__global DATA_TYPE *)partial_sum.ptr + get_group_id(0))[0] = local_sums[0];
+        // Perform parallel reduction
+        for(unsigned int i = lsize >> 1; i > 0; i >>= 1)
+        {
+            if(lid < i)
+            {
+                local_sums[lid] += local_sums[lid + i];
+            }
+            barrier(CLK_LOCAL_MEM_FENCE);
+        }
+
+        if(lid == 0)
+        {
+            ((__global DATA_TYPE *)offset(&partial_sum, get_group_id(0), y))[0] = local_sums[0];
+        }
     }
 }
\ No newline at end of file
diff --git a/src/core/CL/kernels/CLReductionOperationKernel.cpp b/src/core/CL/kernels/CLReductionOperationKernel.cpp
index 18a8e35..1dd5eb9 100644
--- a/src/core/CL/kernels/CLReductionOperationKernel.cpp
+++ b/src/core/CL/kernels/CLReductionOperationKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -118,8 +118,8 @@
     out_window.set(Window::DimX, Window::Dimension(0, 0, 0));
 
     // Get first input and output slices
-    Window in_slice  = window.first_slice_window_1D();
-    Window out_slice = out_window.first_slice_window_1D();
+    Window in_slice  = window.first_slice_window_2D();
+    Window out_slice = out_window.first_slice_window_2D();
 
     // Reshape window
     const unsigned int border_width = ((in_slice.x().end() % 128) != 0) ? 128 - in_slice.x().end() % 128 : 0;
@@ -127,14 +127,14 @@
 
     // Set local sums buffer
     unsigned int local_sum_size = _lws_hint[0] * _input->info()->element_size();
-    _kernel.setArg(num_arguments_per_1D_tensor() * 2, local_sum_size, nullptr);
+    _kernel.setArg(num_arguments_per_2D_tensor() * 2, local_sum_size, nullptr);
 
     do
     {
         unsigned int idx = 0;
-        add_1D_tensor_argument(idx, _input, in_slice);
-        add_1D_tensor_argument(idx, _output, out_slice);
+        add_2D_tensor_argument(idx, _input, in_slice);
+        add_2D_tensor_argument(idx, _output, out_slice);
         enqueue(queue, *this, in_slice, _lws_hint);
     }
-    while(window.slide_window_slice_1D(in_slice) && window.slide_window_slice_1D(out_slice));
+    while(window.slide_window_slice_2D(in_slice) && window.slide_window_slice_2D(out_slice));
 }