Improve Winograd performance on OpenCL

- Performs more output elements per work-item in the case of Fp16
computation in Winograd Input/Output transform

Resolves COMPMID-6018

Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Change-Id: If5e6f5182eff8c1f05a3505c437d0a997490f0bd
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/9447
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Jakub Sujak <jakub.sujak@arm.com>
Reviewed-by: Viet-Hoa Do <viet-hoa.do@arm.com>
Benchmark: Arm Jenkins <bsgcomp@arm.com>

4 files changed, 186 insertions, 121 deletions

diff --git a/src/core/CL/cl_kernels/nhwc/winograd_input_transform.cl b/src/core/CL/cl_kernels/nhwc/winograd_input_transform.cl
index ba7b13b774..7341336b92 100644
--- a/src/core/CL/cl_kernels/nhwc/winograd_input_transform.cl
+++ b/src/core/CL/cl_kernels/nhwc/winograd_input_transform.cl
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2022 Arm Limited.
+ * Copyright (c) 2018-2023 Arm Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -24,42 +24,42 @@
 #include "helpers.h"
 #include "tile_helpers.h"
 
-#define OUTPUT_ROW_4x4_5x5(out, tmp, comm_fact)                     \
-    ({                                                              \
-        comm_fact.s0 = tmp.s2 - 4.25f * tmp.s4 + tmp.s6;            \
-        comm_fact.s1 = tmp.s1 - 4.25f * tmp.s3 + tmp.s5;            \
-        comm_fact.s2 = 2.5f * tmp.s3;                               \
-        comm_fact.s3 = 0.5f * tmp.s1 + 2.f * tmp.s5 - comm_fact.s2; \
-        comm_fact.s4 = 0.25f * tmp.s2 - 1.25f * tmp.s4 + tmp.s6;    \
-        comm_fact.s5 = 4.f * tmp.s2 + tmp.s6 - 5.f * tmp.s4;        \
-        comm_fact.s6 = 2.f * tmp.s1 + 0.5f * tmp.s5 - comm_fact.s2; \
+#define OUTPUT_ROW_4x4_5x5(out, tmp, comm_fact)                                           \
+    ({                                                                                    \
+        comm_fact.s0 = tmp.s2 - (DATA_TYPE)4.25f * tmp.s4 + tmp.s6;                       \
+        comm_fact.s1 = tmp.s1 - (DATA_TYPE)4.25f * tmp.s3 + tmp.s5;                       \
+        comm_fact.s2 = (DATA_TYPE)2.5f * tmp.s3;                                          \
+        comm_fact.s3 = (DATA_TYPE)0.5f * tmp.s1 + (DATA_TYPE)2.f * tmp.s5 - comm_fact.s2; \
+        comm_fact.s4 = (DATA_TYPE)0.25f * tmp.s2 - (DATA_TYPE)1.25f * tmp.s4 + tmp.s6;    \
+        comm_fact.s5 = (DATA_TYPE)4.f * tmp.s2 + tmp.s6 - (DATA_TYPE)5.f * tmp.s4;        \
+        comm_fact.s6 = (DATA_TYPE)2.f * tmp.s1 + (DATA_TYPE)0.5f * tmp.s5 - comm_fact.s2; \
         \
-        out.s0 = tmp.s0 - tmp.s6 + 5.25f * tmp.s4 - 5.25f * tmp.s2; \
-        out.s1 = comm_fact.s0 + comm_fact.s1;                       \
-        out.s2 = comm_fact.s0 - comm_fact.s1;                       \
-        out.s3 = comm_fact.s3 + comm_fact.s4;                       \
-        out.s4 = comm_fact.s4 - comm_fact.s3;                       \
-        out.s5 = comm_fact.s5 + comm_fact.s6;                       \
-        out.s6 = comm_fact.s5 - comm_fact.s6;                       \
-        out.s7 = tmp.s7 - tmp.s1 + 5.25f * tmp.s3 - 5.25f * tmp.s5; \
+        out.s0 = tmp.s0 - tmp.s6 + (DATA_TYPE)5.25f * tmp.s4 - (DATA_TYPE)5.25f * tmp.s2; \
+        out.s1 = comm_fact.s0 + comm_fact.s1;                                             \
+        out.s2 = comm_fact.s0 - comm_fact.s1;                                             \
+        out.s3 = comm_fact.s3 + comm_fact.s4;                                             \
+        out.s4 = comm_fact.s4 - comm_fact.s3;                                             \
+        out.s5 = comm_fact.s5 + comm_fact.s6;                                             \
+        out.s6 = comm_fact.s5 - comm_fact.s6;                                             \
+        out.s7 = tmp.s7 - tmp.s1 + (DATA_TYPE)5.25f * tmp.s3 - (DATA_TYPE)5.25f * tmp.s5; \
     })
 
-#define OUTPUT_ROW_2x2_7x7(out, tmp, comm_fact)                                                    \
-    ({                                                                                             \
-        comm_fact.s0 = 36.0f * tmp.s2 - 13.0f * tmp.s4 + tmp.s6;                                   \
-        comm_fact.s1 = 36.0f * tmp.s1 - 13.0f * tmp.s3 + 1.0f * tmp.s5;                            \
-        comm_fact.s2 = 9.0f * tmp.s2 - 10.0f * tmp.s4 + tmp.s6;                                    \
-        comm_fact.s3 = 18.0f * tmp.s1 - 20.0f * tmp.s3 + 2.0f * tmp.s5;                            \
-        comm_fact.s4 = 4.0f * tmp.s2 - 5.0f * tmp.s4 + tmp.s6;                                     \
-        comm_fact.s5 = 12.0f * tmp.s1 - 15.0f * tmp.s3 + 3.0f * tmp.s5;                            \
-        out.s0       = -36.0f * tmp.s0 + 49.0f * tmp.s2 + -14.0f * tmp.s4 + tmp.s6;                \
-        out.s1       = comm_fact.s0 - comm_fact.s1;                                                \
-        out.s2       = comm_fact.s0 + comm_fact.s1;                                                \
-        out.s3       = comm_fact.s2 - comm_fact.s3;                                                \
-        out.s4       = comm_fact.s2 + comm_fact.s3;                                                \
-        out.s5       = comm_fact.s4 - comm_fact.s5;                                                \
-        out.s6       = comm_fact.s4 + comm_fact.s5;                                                \
-        out.s7       = -36.0f * tmp.s1 + 0.0f * tmp.s2 + 49.0f * tmp.s3 - 14.0f * tmp.s5 + tmp.s7; \
+#define OUTPUT_ROW_2x2_7x7(out, tmp, comm_fact)                                                                                                \
+    ({                                                                                                                                         \
+        comm_fact.s0 = (DATA_TYPE)36.0f * tmp.s2 - (DATA_TYPE)13.0f * tmp.s4 + tmp.s6;                                                         \
+        comm_fact.s1 = (DATA_TYPE)36.0f * tmp.s1 - (DATA_TYPE)13.0f * tmp.s3 + (DATA_TYPE)1.0f * tmp.s5;                                       \
+        comm_fact.s2 = (DATA_TYPE)9.0f * tmp.s2 - (DATA_TYPE)10.0f * tmp.s4 + tmp.s6;                                                          \
+        comm_fact.s3 = (DATA_TYPE)18.0f * tmp.s1 - (DATA_TYPE)20.0f * tmp.s3 + (DATA_TYPE)2.0f * tmp.s5;                                       \
+        comm_fact.s4 = (DATA_TYPE)4.0f * tmp.s2 - (DATA_TYPE)5.0f * tmp.s4 + tmp.s6;                                                           \
+        comm_fact.s5 = (DATA_TYPE)12.0f * tmp.s1 - (DATA_TYPE)15.0f * tmp.s3 + (DATA_TYPE)3.0f * tmp.s5;                                       \
+        out.s0       = -(DATA_TYPE)36.0f * tmp.s0 + (DATA_TYPE)49.0f * tmp.s2 + -(DATA_TYPE)14.0f * tmp.s4 + tmp.s6;                           \
+        out.s1       = comm_fact.s0 - comm_fact.s1;                                                                                            \
+        out.s2       = comm_fact.s0 + comm_fact.s1;                                                                                            \
+        out.s3       = comm_fact.s2 - comm_fact.s3;                                                                                            \
+        out.s4       = comm_fact.s2 + comm_fact.s3;                                                                                            \
+        out.s5       = comm_fact.s4 - comm_fact.s5;                                                                                            \
+        out.s6       = comm_fact.s4 + comm_fact.s5;                                                                                            \
+        out.s7       = -(DATA_TYPE)36.0f * tmp.s1 + (DATA_TYPE)0.0f * tmp.s2 + (DATA_TYPE)49.0f * tmp.s3 - (DATA_TYPE)14.0f * tmp.s5 + tmp.s7; \
     })
 
 #if defined(PAD_LEFT) && defined(PAD_TOP) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H)
@@ -113,9 +113,13 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc(
     const int _INUM_TILES_X,
     const int _INUM_TILES_Y)
 {
-    const int cout = GET_SPATIAL_IDX(0, 1, 0); // OFM
-    const int mout = GET_SPATIAL_IDX(1, 1, 0); // NUM_TILES_X x NUM_TILES_Y
+    const int cout = GET_SPATIAL_IDX(0, N0, 0); // OFM
+    const int mout = GET_SPATIAL_IDX(1, 1, 0);  // NUM_TILES_X x NUM_TILES_Y
+#if defined(IS_BATCHED)
     const int bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX
+#else                                          // defined(IS_BATCHED)
+    const int bout = 0; // BATCH SIZE IDX
+#endif                                         // defined(IS_BATCHED)
 
     int x = (mout % _INUM_TILES_X) * OUTPUT_TILE_W;
     int y = (mout / _INUM_TILES_X) * OUTPUT_TILE_H;
@@ -124,8 +128,8 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc(
 
 #if defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL)
 
-    TILE(DATA_TYPE, 6, 1, in);
-    TILE(DATA_TYPE, 6, 1, out);
+    TILE(DATA_TYPE, 6, N0, in);
+    TILE(DATA_TYPE, 6, N0, out);
 
     // Initialize the input tile
     LOOP_UNROLLING(int, i, 0, 1, 6,
@@ -134,22 +138,22 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc(
     })
 
 #if defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL)
-    T_LOAD_NHWC(DATA_TYPE, 1, 6, 1, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in);
+    T_LOAD_NHWC(DATA_TYPE, 1, 6, N0, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in);
 #else  // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL)
-    T_LOAD_NHWC(DATA_TYPE, 6, 1, 1, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in);
+    T_LOAD_NHWC(DATA_TYPE, 6, 1, N0, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in);
 #endif // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL)
 
-    TILE(DATA_TYPE, 6, 1, com);
+    TILE(DATA_TYPE, 6, N0, com);
 
     LOOP_UNROLLING(int, i, 0, 1, 6,
     {
-        in[i].v *= 4.0f;
+        in[i].v *= (DATA_TYPE)4.0f;
     })
 
-    com[0].v = in[2].v - 4.f * in[0].v;
-    com[1].v = in[3].v - 4.f * in[1].v;
-    com[2].v = in[4].v - 4.f * in[2].v;
-    com[3].v = in[5].v - 4.f * in[3].v;
+    com[0].v = in[2].v - (DATA_TYPE)4.f * in[0].v;
+    com[1].v = in[3].v - (DATA_TYPE)4.f * in[1].v;
+    com[2].v = in[4].v - (DATA_TYPE)4.f * in[2].v;
+    com[3].v = in[5].v - (DATA_TYPE)4.f * in[3].v;
     com[4].v = in[3].v - in[1].v;
     com[4].v = com[4].v + com[4].v;
     com[5].v = in[4].v - in[2].v;
@@ -169,11 +173,11 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc(
         dst_indirect_y[i].v += bout *_INUM_TILES_X *_INUM_TILES_Y * 6;
     })
 
-    T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 6, 1, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y);
+    T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 6, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y);
 
 #else  // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL)
 
-    TILE(DATA_TYPE, 36, 1, in);
+    TILE(DATA_TYPE, 36, N0, in);
 
     // Initialize the input tile
     LOOP_UNROLLING(int, i, 0, 1, 36,
@@ -182,10 +186,10 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc(
     })
 
     // Load the tile from a NHWC tensor
-    T_LOAD_NHWC(DATA_TYPE, 6, 6, 1, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in);
+    T_LOAD_NHWC(DATA_TYPE, 6, 6, N0, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in);
 
-    TILE(DATA_TYPE, 6, 1, com);
-    TILE(DATA_TYPE, 36, 1, tmp);
+    TILE(DATA_TYPE, 6, N0, com);
+    TILE(DATA_TYPE, 36, N0, tmp);
 
     LOOP_UNROLLING(int, i, 0, 1, 6,
     {
@@ -204,14 +208,14 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc(
         tmp[i + 5 * 6].v = com[3].v - com[1].v;
     })
 
-    TILE(DATA_TYPE, 36, 1, out);
+    TILE(DATA_TYPE, 36, N0, out);
 
     LOOP_UNROLLING(int, i, 0, 1, 6,
     {
-        com[0].v         = tmp[i * 6 + 2].v - 4.f *tmp[i * 6 + 0].v;
-        com[1].v         = tmp[i * 6 + 3].v - 4.f *tmp[i * 6 + 1].v;
-        com[2].v         = tmp[i * 6 + 4].v - 4.f *tmp[i * 6 + 2].v;
-        com[3].v         = tmp[i * 6 + 5].v - 4.f *tmp[i * 6 + 3].v;
+        com[0].v         = tmp[i * 6 + 2].v - (DATA_TYPE)4.f *tmp[i * 6 + 0].v;
+        com[1].v         = tmp[i * 6 + 3].v - (DATA_TYPE)4.f *tmp[i * 6 + 1].v;
+        com[2].v         = tmp[i * 6 + 4].v - (DATA_TYPE)4.f *tmp[i * 6 + 2].v;
+        com[3].v         = tmp[i * 6 + 5].v - (DATA_TYPE)4.f *tmp[i * 6 + 3].v;
         com[4].v         = tmp[i * 6 + 3].v - tmp[i * 6 + 1].v;
         com[4].v         = com[4].v + com[4].v;
         com[5].v         = tmp[i * 6 + 4].v - tmp[i * 6 + 2].v;
@@ -232,7 +236,7 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc(
         dst_indirect_y[i].v += bout *_INUM_TILES_X *_INUM_TILES_Y * 36;
     })
 
-    T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 36, 1, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y);
+    T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 36, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y);
 #endif // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL)
 }
 #endif // defined(WINOGRAD_INPUT_TRANSFORM_4X4_3X3_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_4X1_3X1_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_1X4_1X3_STEPZ1_NHWC)
@@ -287,7 +291,11 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc(
 {
     const int cout = GET_SPATIAL_IDX(0, 1, 0); // OFM
     const int mout = GET_SPATIAL_IDX(1, 1, 0); // NUM_TILES_X x NUM_TILES_Y
+#if defined(IS_BATCHED)
     const int bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX
+#else                                          // defined(IS_BATCHED)
+    const int bout = 0; // BATCH SIZE IDX
+#endif                                         // defined(IS_BATCHED)
 
     int x = (mout % _INUM_TILES_X) * OUTPUT_TILE_W;
     int y = (mout / _INUM_TILES_X) * OUTPUT_TILE_H;
@@ -306,27 +314,27 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc(
     })
 
 #if defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL)
-    T_LOAD_NHWC(DATA_TYPE, 1, 8, 1, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in);
+    T_LOAD_NHWC(DATA_TYPE, 1, 8, N0, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in);
 #else  // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL)
-    T_LOAD_NHWC(DATA_TYPE, 8, 1, 1, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in);
+    T_LOAD_NHWC(DATA_TYPE, 8, 1, N0, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in);
 #endif // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL)
 
     TILE(DATA_TYPE, 1, 8, com);
 
-    com[0].s[0] = in[2].v - 4.25f * in[4].v + in[6].v;
-    com[0].s[1] = in[1].v - 4.25f * in[3].v + in[5].v;
-    com[0].s[2] = 0.5f * in[1].v - 2.5f * in[3].v + 2.0f * in[5].v;
-    com[0].s[3] = 0.25f * in[2].v - 1.25f * in[4].v + in[6].v;
-    com[0].s[4] = 4.0f * in[2].v - 5.0f * in[4].v + in[6].v;
-    com[0].s[5] = 2.0f * in[1].v - 2.5f * in[3].v + 0.5f * in[5].v;
-    out[0].s[0] = in[0].v - 5.25f * in[2].v + 5.25f * in[4].v - in[6].v;
+    com[0].s[0] = in[2].v - (DATA_TYPE)4.25f * in[4].v + in[6].v;
+    com[0].s[1] = in[1].v - (DATA_TYPE)4.25f * in[3].v + in[5].v;
+    com[0].s[2] = (DATA_TYPE)0.5f * in[1].v - (DATA_TYPE)2.5f * in[3].v + (DATA_TYPE)2.0f * in[5].v;
+    com[0].s[3] = (DATA_TYPE)0.25f * in[2].v - (DATA_TYPE)1.25f * in[4].v + in[6].v;
+    com[0].s[4] = (DATA_TYPE)4.0f * in[2].v - (DATA_TYPE)5.0f * in[4].v + in[6].v;
+    com[0].s[5] = (DATA_TYPE)2.0f * in[1].v - (DATA_TYPE)2.5f * in[3].v + (DATA_TYPE)0.5f * in[5].v;
+    out[0].s[0] = in[0].v - 5.25f * in[2].v + (DATA_TYPE)5.25f * in[4].v - in[6].v;
     out[1].s[0] = com[0].s[0] + com[0].s[1];
     out[2].s[0] = com[0].s[0] - com[0].s[1];
     out[3].s[0] = com[0].s[3] + com[0].s[2];
     out[4].s[0] = com[0].s[3] - com[0].s[2];
     out[5].s[0] = com[0].s[4] + com[0].s[5];
     out[6].s[0] = com[0].s[4] - com[0].s[5];
-    out[7].s[0] = -in[1].v + 5.25f * in[3].v - 5.25f * in[5].v + in[7].v;
+    out[7].s[0] = -in[1].v + (DATA_TYPE)5.25f * in[3].v - (DATA_TYPE)5.25f * in[5].v + in[7].v;
 
     TILE(uint, 8, 1, dst_indirect_y);
 
@@ -378,20 +386,20 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc(
 
     LOOP_UNROLLING(int, i, 0, 1, 8,
     {
-        com[0].s[0]         = tmp[i].s[2] - 4.25f * tmp[i].s[4] + tmp[i].s[6];
-        com[0].s[1]         = tmp[i].s[1] - 4.25f * tmp[i].s[3] + tmp[i].s[5];
-        com[0].s[2]         = 0.5f * tmp[i].s[1] - 2.5f * tmp[i].s[3] + 2.0f * tmp[i].s[5];
-        com[0].s[3]         = 0.25f * tmp[i].s[2] - 1.25f * tmp[i].s[4] + tmp[i].s[6];
-        com[0].s[4]         = 4.0f * tmp[i].s[2] - 5.0f * tmp[i].s[4] + tmp[i].s[6];
-        com[0].s[5]         = 2.0f * tmp[i].s[1] - 2.5f * tmp[i].s[3] + 0.5f * tmp[i].s[5];
-        out[i * 8 + 0].s[0] = tmp[i].s[0] - 5.25f * tmp[i].s[2] + 5.25f * tmp[i].s[4] - tmp[i].s[6];
+        com[0].s[0]         = tmp[i].s[2] - (DATA_TYPE)4.25f * tmp[i].s[4] + tmp[i].s[6];
+        com[0].s[1]         = tmp[i].s[1] - (DATA_TYPE)4.25f * tmp[i].s[3] + tmp[i].s[5];
+        com[0].s[2]         = (DATA_TYPE)0.5f * tmp[i].s[1] - (DATA_TYPE)2.5f * tmp[i].s[3] + (DATA_TYPE)2.0f * tmp[i].s[5];
+        com[0].s[3]         = (DATA_TYPE)0.25f * tmp[i].s[2] - (DATA_TYPE)1.25f * tmp[i].s[4] + tmp[i].s[6];
+        com[0].s[4]         = (DATA_TYPE)4.0f * tmp[i].s[2] - (DATA_TYPE)5.0f * tmp[i].s[4] + tmp[i].s[6];
+        com[0].s[5]         = (DATA_TYPE)2.0f * tmp[i].s[1] - (DATA_TYPE)2.5f * tmp[i].s[3] + (DATA_TYPE)0.5f * tmp[i].s[5];
+        out[i * 8 + 0].s[0] = tmp[i].s[0] - (DATA_TYPE)5.25f * tmp[i].s[2] + (DATA_TYPE)5.25f * tmp[i].s[4] - tmp[i].s[6];
         out[i * 8 + 1].s[0] = com[0].s[0] + com[0].s[1];
         out[i * 8 + 2].s[0] = com[0].s[0] - com[0].s[1];
         out[i * 8 + 3].s[0] = com[0].s[3] + com[0].s[2];
         out[i * 8 + 4].s[0] = com[0].s[3] - com[0].s[2];
         out[i * 8 + 5].s[0] = com[0].s[4] + com[0].s[5];
         out[i * 8 + 6].s[0] = com[0].s[4] - com[0].s[5];
-        out[i * 8 + 7].s[0] = -tmp[i].s[1] + 5.25f * tmp[i].s[3] - 5.25f * tmp[i].s[5] + tmp[i].s[7];
+        out[i * 8 + 7].s[0] = -tmp[i].s[1] + (DATA_TYPE)5.25f * tmp[i].s[3] - (DATA_TYPE)5.25f * tmp[i].s[5] + tmp[i].s[7];
     })
 
     TILE(uint, 64, 1, dst_indirect_y);
@@ -458,7 +466,11 @@ __kernel void winograd_input_transform_2x2_7x7_stepz1_nhwc(
 {
     const int cout = GET_SPATIAL_IDX(0, 1, 0); // OFM
     const int mout = GET_SPATIAL_IDX(1, 1, 0); // NUM_TILES_X x NUM_TILES_Y
+#if defined(IS_BATCHED)
     const int bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX
+#else                                          // defined(IS_BATCHED)
+    const int bout = 0; // BATCH SIZE IDX
+#endif                                         // defined(IS_BATCHED)
 
     int x = (mout % _INUM_TILES_X) * OUTPUT_TILE_W;
     int y = (mout / _INUM_TILES_X) * OUTPUT_TILE_H;
@@ -489,20 +501,20 @@ __kernel void winograd_input_transform_2x2_7x7_stepz1_nhwc(
 
     TILE(DATA_TYPE, 1, 8, com) = { { { 0 } } };
 
-    com[0].s[0] = 36.0f * in[2].v - 13.0f * in[4].v + in[6].v;
-    com[0].s[1] = 36.0f * in[1].v - 13.0f * in[3].v + 1.0f * in[5].v;
-    com[0].s[2] = 9.0f * in[2].v - 10.0f * in[4].v + in[6].v;
-    com[0].s[3] = 18.0f * in[1].v - 20.0f * in[3].v + 2.0f * in[5].v;
-    com[0].s[4] = 4.0f * in[2].v - 5.0f * in[4].v + in[6].v;
-    com[0].s[5] = 12.0f * in[1].v - 15.0f * in[3].v + 3.0f * in[5].v;
-    out[0].s[0] = -36.0f * in[0].v + 49.0f * in[2].v + -14.0f * in[4].v + in[6].v;
+    com[0].s[0] = (DATA_TYPE)36.0f * in[2].v - (DATA_TYPE)13.0f * in[4].v + in[6].v;
+    com[0].s[1] = (DATA_TYPE)36.0f * in[1].v - (DATA_TYPE)13.0f * in[3].v + (DATA_TYPE)1.0f * in[5].v;
+    com[0].s[2] = (DATA_TYPE)9.0f * in[2].v - (DATA_TYPE)10.0f * in[4].v + in[6].v;
+    com[0].s[3] = (DATA_TYPE)18.0f * in[1].v - (DATA_TYPE)20.0f * in[3].v + (DATA_TYPE)2.0f * in[5].v;
+    com[0].s[4] = (DATA_TYPE)4.0f * in[2].v - (DATA_TYPE)5.0f * in[4].v + in[6].v;
+    com[0].s[5] = (DATA_TYPE)12.0f * in[1].v - (DATA_TYPE)15.0f * in[3].v + (DATA_TYPE)3.0f * in[5].v;
+    out[0].s[0] = (DATA_TYPE) - 36.0f * in[0].v + (DATA_TYPE)49.0f * in[2].v + -(DATA_TYPE)14.0f * in[4].v + in[6].v;
     out[1].s[0] = com[0].s[0] - com[0].s[1];
     out[2].s[0] = com[0].s[0] + com[0].s[1];
     out[3].s[0] = com[0].s[2] - com[0].s[3];
     out[4].s[0] = com[0].s[2] + com[0].s[3];
     out[5].s[0] = com[0].s[4] - com[0].s[5];
     out[6].s[0] = com[0].s[4] + com[0].s[5];
-    out[7].s[0] = -36.0f * in[1].v + 0.0f * in[2].v + 49.0f * in[3].v - 14.0f * in[5].v + in[7].v;
+    out[7].s[0] = -(DATA_TYPE)36.0f * in[1].v + (DATA_TYPE)0.0f * in[2].v + (DATA_TYPE)49.0f * in[3].v - (DATA_TYPE)14.0f * in[5].v + in[7].v;
 
     TILE(uint, 8, 1, dst_indirect_y);
 
@@ -554,20 +566,20 @@ __kernel void winograd_input_transform_2x2_7x7_stepz1_nhwc(
 
     LOOP_UNROLLING(int, i, 0, 1, 8,
     {
-        com[0].s[0]         = 36.0f * tmp[i].s[2] - 13.0f * tmp[i].s[4] + tmp[i].s[6];
-        com[0].s[1]         = 36.0f * tmp[i].s[1] - 13.0f * tmp[i].s[3] + 1.0f * tmp[i].s[5];
-        com[0].s[2]         = 9.0f * tmp[i].s[2] - 10.0f * tmp[i].s[4] + tmp[i].s[6];
-        com[0].s[3]         = 18.0f * tmp[i].s[1] - 20.0f * tmp[i].s[3] + 2.0f * tmp[i].s[5];
-        com[0].s[4]         = 4.0f * tmp[i].s[2] - 5.0f * tmp[i].s[4] + tmp[i].s[6];
-        com[0].s[5]         = 12.0f * tmp[i].s[1] - 15.0f * tmp[i].s[3] + 3.0f * tmp[i].s[5];
-        out[i * 8 + 0].s[0] = -36.0f * tmp[i].s[0] + 49.0f * tmp[i].s[2] + -14.0f * tmp[i].s[4] + tmp[i].s[6];
+        com[0].s[0]         = (DATA_TYPE)36.0f * tmp[i].s[2] - (DATA_TYPE)13.0f * tmp[i].s[4] + tmp[i].s[6];
+        com[0].s[1]         = (DATA_TYPE)36.0f * tmp[i].s[1] - (DATA_TYPE)13.0f * tmp[i].s[3] + (DATA_TYPE)1.0f * tmp[i].s[5];
+        com[0].s[2]         = (DATA_TYPE)9.0f * tmp[i].s[2] - (DATA_TYPE)10.0f * tmp[i].s[4] + tmp[i].s[6];
+        com[0].s[3]         = (DATA_TYPE)18.0f * tmp[i].s[1] - (DATA_TYPE)20.0f * tmp[i].s[3] + (DATA_TYPE)2.0f * tmp[i].s[5];
+        com[0].s[4]         = (DATA_TYPE)4.0f * tmp[i].s[2] - (DATA_TYPE)5.0f * tmp[i].s[4] + tmp[i].s[6];
+        com[0].s[5]         = (DATA_TYPE)12.0f * tmp[i].s[1] - (DATA_TYPE)15.0f * tmp[i].s[3] + (DATA_TYPE)3.0f * tmp[i].s[5];
+        out[i * 8 + 0].s[0] = (DATA_TYPE) - 36.0f * tmp[i].s[0] + (DATA_TYPE)49.0f * tmp[i].s[2] + -(DATA_TYPE)14.0f * tmp[i].s[4] + tmp[i].s[6];
         out[i * 8 + 1].s[0] = com[0].s[0] - com[0].s[1];
         out[i * 8 + 2].s[0] = com[0].s[0] + com[0].s[1];
         out[i * 8 + 3].s[0] = com[0].s[2] - com[0].s[3];
         out[i * 8 + 4].s[0] = com[0].s[2] + com[0].s[3];
         out[i * 8 + 5].s[0] = com[0].s[4] - com[0].s[5];
         out[i * 8 + 6].s[0] = com[0].s[4] + com[0].s[5];
-        out[i * 8 + 7].s[0] = -36.0f * tmp[i].s[1] + 0.0f * tmp[i].s[2] + 49.0f * tmp[i].s[3] - 14.0f * tmp[i].s[5] + tmp[i].s[7];
+        out[i * 8 + 7].s[0] = -(DATA_TYPE)36.0f * tmp[i].s[1] + (DATA_TYPE)0.0f * tmp[i].s[2] + (DATA_TYPE)49.0f * tmp[i].s[3] - (DATA_TYPE)14.0f * tmp[i].s[5] + tmp[i].s[7];
     })
 
     TILE(uint, 64, 1, dst_indirect_y);
diff --git a/src/core/CL/cl_kernels/nhwc/winograd_output_transform.cl b/src/core/CL/cl_kernels/nhwc/winograd_output_transform.cl
index 0883cd99c8..9eb995fbb2 100644
--- a/src/core/CL/cl_kernels/nhwc/winograd_output_transform.cl
+++ b/src/core/CL/cl_kernels/nhwc/winograd_output_transform.cl
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2022 Arm Limited.
+ * Copyright (c) 2018-2023 Arm Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -75,7 +75,11 @@ __kernel void winograd_output_transform_2x2_7x7_nhwc(
 {
     const int cout = GET_SPATIAL_IDX(0, N0, 0); // OFM
     const int mout = GET_SPATIAL_IDX(1, 1, 0);  // WINOGRAD OUTPUT TILES
-    const int bout = GET_SPATIAL_IDX(2, 1, 0);  // BATCH SIZE IDX
+#if defined(IS_BATCHED)
+    const int bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX
+#else                                          // defined(IS_BATCHED)
+    const int bout = 0; // BATCH SIZE IDX
+#endif                                         // defined(IS_BATCHED)
 
     int x_out = (mout % NUM_TILES_X) * OUTPUT_TILE_W;
     int y_out = (mout / NUM_TILES_X) * OUTPUT_TILE_H;
@@ -103,7 +107,7 @@ __kernel void winograd_output_transform_2x2_7x7_nhwc(
 
     // Compute out0 and out01
     out[0].v = in[0].v + in[1].v + in[2].v + in[3].v + in[4].v + in[5].v + in[6].v;
-    out[1].v = -in[1].v + in[2].v - 2.f * in[3].v + 2.0f * in[4].v - 3.0f * in[5].v + 3.0f * in[6].v + in[7].v;
+    out[1].v = -in[1].v + in[2].v - (DATA_TYPE)2.f * in[3].v + (DATA_TYPE)2.0f * in[4].v - (DATA_TYPE)3.0f * in[5].v + (DATA_TYPE)3.0f * in[6].v + in[7].v;
 
 #if defined(HAS_BIAS)
     // Add bias
@@ -161,14 +165,14 @@ __kernel void winograd_output_transform_2x2_7x7_nhwc(
     LOOP_UNROLLING(int, i, 0, 1, 8,
     {
         tmp[i * 2].v     = in[0 + i].v + in[8 + i].v + in[16 + i].v + in[24 + i].v + in[32 + i].v + in[40 + i].v + in[48 + i].v;
-        tmp[i * 2 + 1].v = -in[8 + i].v + in[16 + i].v - 2 * in[24 + i].v + 2 * in[32 + i].v + -3 * in[40 + i].v + 3 * in[48 + i].v + in[56 + i].v;
+        tmp[i * 2 + 1].v = -in[8 + i].v + in[16 + i].v - (DATA_TYPE)2 * in[24 + i].v + (DATA_TYPE)2 * in[32 + i].v + (DATA_TYPE) - 3 * in[40 + i].v + (DATA_TYPE)3 * in[48 + i].v + in[56 + i].v;
     })
 
     // Compute the 2x2 output tile
     LOOP_UNROLLING(int, i, 0, 1, 2,
     {
         out[i * 2].v     = tmp[0 + i].v + tmp[2 + i].v + tmp[4 + i].v + tmp[6 + i].v + tmp[8 + i].v + tmp[10 + i].v + tmp[12 + i].v;
-        out[i * 2 + 1].v = -tmp[2 + i].v + tmp[4 + i].v - 2 * tmp[6 + i].v + 2 * tmp[8 + i].v - 3 * tmp[10 + i].v + 3 * tmp[12 + i].v + tmp[14 + i].v;
+        out[i * 2 + 1].v = -tmp[2 + i].v + tmp[4 + i].v - (DATA_TYPE)2 * tmp[6 + i].v + (DATA_TYPE)2 * tmp[8 + i].v - (DATA_TYPE)3 * tmp[10 + i].v + (DATA_TYPE)3 * tmp[12 + i].v + tmp[14 + i].v;
     })
 
 #if defined(HAS_BIAS)
@@ -252,7 +256,11 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc(
 {
     const int cout = GET_SPATIAL_IDX(0, N0, 0); // OFM
     const int mout = GET_SPATIAL_IDX(1, 1, 0);  // WINOGRAD OUTPUT TILES
-    const int bout = GET_SPATIAL_IDX(2, 1, 0);  // BATCH SIZE IDX
+#if defined(IS_BATCHED)
+    const int bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX
+#else                                          // defined(IS_BATCHED)
+    const int bout = 0; // BATCH SIZE IDX
+#endif                                         // defined(IS_BATCHED)
 
 #if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL)
 
@@ -277,9 +285,9 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc(
 
     // Compute out00, out01, out02 and out03
     out[0].v = in[0].v + in[1].v + in[2].v + in[3].v + in[4].v;
-    out[1].v = in[1].v - in[2].v + 2.0f * in[3].v - 2.0f * in[4].v;
-    out[2].v = in[1].v + in[2].v + 4.0f * in[3].v + 4.0f * in[4].v;
-    out[3].v = in[1].v - in[2].v + 8.0f * in[3].v - 8.0f * in[4].v + in[5].v;
+    out[1].v = in[1].v - in[2].v + (DATA_TYPE)2.0f * in[3].v - (DATA_TYPE)2.0f * in[4].v;
+    out[2].v = in[1].v + in[2].v + (DATA_TYPE)4.0f * in[3].v + (DATA_TYPE)4.0f * in[4].v;
+    out[3].v = in[1].v - in[2].v + (DATA_TYPE)8.0f * in[3].v - (DATA_TYPE)8.0f * in[4].v + in[5].v;
 
 #if defined(HAS_BIAS)
     TILE(DATA_TYPE, 1, N0, b);
@@ -449,7 +457,11 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc(
 {
     const int cout = GET_SPATIAL_IDX(0, N0, 0); // OFM
     const int mout = GET_SPATIAL_IDX(1, 1, 0);  // WINOGRAD OUTPUT TILES
-    const int bout = GET_SPATIAL_IDX(2, 1, 0);  // BATCH SIZE IDX
+#if defined(IS_BATCHED)
+    const int bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX
+#else                                          // defined(IS_BATCHED)
+    const int bout = 0; // BATCH SIZE IDX
+#endif                                         // defined(IS_BATCHED)
 
 #if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL)
     TILE(DATA_TYPE, 8, N0, in);
@@ -474,13 +486,13 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc(
 
     // A^T * in, and in this degenerate case out consists of 1 column/row
     tmp[0].v = in[1].v - in[2].v;
-    tmp[1].v = 2.0f * (in[3].v - in[4].v);
-    tmp[2].v = 2.0f * (in[5].v + in[6].v);
+    tmp[1].v = (DATA_TYPE)2.0f * (in[3].v - in[4].v);
+    tmp[2].v = (DATA_TYPE)2.0f * (in[5].v + in[6].v);
     tmp[3].v = in[3].v + in[4].v;
-    out[0].v = in[0].v + in[1].v + in[2].v + tmp[3].v + 4.0f * tmp[2].v;
-    out[1].v = tmp[0].v + tmp[1].v + 4.0f * (in[5].v - in[6].v);
-    out[2].v = in[1].v + in[2].v + 4.0f * tmp[3].v + tmp[2].v;
-    out[3].v = tmp[0].v + 4.0f * tmp[1].v + in[5].v - in[6].v + in[7].v;
+    out[0].v = in[0].v + in[1].v + in[2].v + tmp[3].v + (DATA_TYPE)4.0f * tmp[2].v;
+    out[1].v = tmp[0].v + tmp[1].v + (DATA_TYPE)4.0f * (in[5].v - in[6].v);
+    out[2].v = in[1].v + in[2].v + (DATA_TYPE)4.0f * tmp[3].v + tmp[2].v;
+    out[3].v = tmp[0].v + (DATA_TYPE)4.0f * tmp[1].v + in[5].v - in[6].v + in[7].v;
 
 #if defined(HAS_BIAS)
     TILE(DATA_TYPE, 1, N0, b);
@@ -1094,4 +1106,4 @@ __kernel void winograd_output_transform_1x4_1x5_nhwc(
 #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X5_NHWC)
 #endif // defined(VEC_SIZE) && VEC_SIZE == 4
 #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL)
-#endif // defined(NUM_TILES_X) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H)
\ No newline at end of file
+#endif // defined(NUM_TILES_X) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H)
diff --git a/src/gpu/cl/kernels/ClWinogradInputTransformKernel.cpp b/src/gpu/cl/kernels/ClWinogradInputTransformKernel.cpp
index d6b038f0f8..48d806dc7c 100644
--- a/src/gpu/cl/kernels/ClWinogradInputTransformKernel.cpp
+++ b/src/gpu/cl/kernels/ClWinogradInputTransformKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2022 Arm Limited.
+ * Copyright (c) 2018-2023 Arm Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -79,8 +79,30 @@ std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITen
     ARM_COMPUTE_UNUSED(output);
     ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
 
-    bool   window_changed = false;
-    Window win            = calculate_max_window(*input, Steps(1, 1));
+    bool window_changed                    = false;
+    int  num_elems_processed_per_iteration = 1;
+
+    if(input->data_layout() == DataLayout::NHWC)
+    {
+        // In the case of FP16 computation, we can perform more
+        // output feature maps in a single work-item.
+        // From experiments, num_elems_processed_per_iteration = 2 looks good for fp16 to
+        // improve the performance. However, in order to make the implementation simpler,
+        // we set num_elems_processed_per_iteration = 2 only when the OFMs are multiple of 2.
+        // Note: At the moment, only Winograd Input Transform 3x3 can support N0 != 1
+        const DataType dt   = input->data_type();
+        const size_t   dim0 = input->dimension(0);
+        const size_t   k_sz = winograd_info.kernel_size.area();
+        const bool     cond = dt == DataType::F16 && ((dim0 % 2) == 0);
+        if(cond)
+        {
+            if(k_sz == 3 || k_sz == 9)
+            {
+                num_elems_processed_per_iteration = 2;
+            }
+        }
+    }
+    Window win = calculate_max_window(*input, Steps(num_elems_processed_per_iteration));
 
     if(input->data_layout() == DataLayout::NCHW)
     {
@@ -143,12 +165,19 @@ void ClWinogradInputTransformKernel::configure(const ClCompileContext &compile_c
     ARM_COMPUTE_ERROR_ON(_num_tiles_x * _num_tiles_y != static_cast<int>(dst->dimension(1)));
     const size_t total_batches = src->tensor_shape().total_size_upper(3);
 
+    // Create window and update padding
+    auto win_config = validate_and_configure_window(src, dst, winograd_info);
+    ARM_COMPUTE_ERROR_THROW_ON(win_config.first);
+    IClKernel::configure_internal(win_config.second, cl::NDRange(1, 1, 8));
+
+    _src_width  = src->dimension(idx_w);
+    _src_height = src->dimension(idx_h);
+
     CLBuildOptions build_opts;
     if(_data_layout == DataLayout::NHWC)
     {
         build_opts.add_option("-DNHWC");
-        _src_width  = src->dimension(idx_w);
-        _src_height = src->dimension(idx_h);
+        build_opts.add_option("-DN0=" + support::cpp11::to_string(win_config.second.x().step()));
         build_opts.add_option("-DPAD_LEFT=" + support::cpp11::to_string(conv_info.pad_left()));
         build_opts.add_option("-DPAD_TOP=" + support::cpp11::to_string(conv_info.pad_top()));
         build_opts.add_option("-DOUTPUT_TILE_W=" + support::cpp11::to_string(output_tile_size.width));
@@ -156,6 +185,7 @@ void ClWinogradInputTransformKernel::configure(const ClCompileContext &compile_c
         build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(src->data_type()));
         build_opts.add_option_if(winograd_info.kernel_size.height == 1, "-DWINOGRAD_INPUT_TRANSFORM_HORIZONTAL");
         build_opts.add_option_if(winograd_info.kernel_size.width == 1, "-DWINOGRAD_INPUT_TRANSFORM_VERTICAL");
+        build_opts.add_option_if(total_batches > 1, "-DIS_BATCHED");
     }
     else
     {
@@ -191,11 +221,6 @@ void ClWinogradInputTransformKernel::configure(const ClCompileContext &compile_c
     build_opts.add_option("-D" + upper_string(kernel_name));
     _kernel = create_kernel(compile_context, kernel_name, build_opts.options());
 
-    // Create window and update padding
-    auto win_config = validate_and_configure_window(src, dst, winograd_info);
-    ARM_COMPUTE_ERROR_THROW_ON(win_config.first);
-    IClKernel::configure_internal(win_config.second, cl::NDRange(1, 1, 8));
-
     _border_size = BorderSize(src->padding());
 
     ARM_COMPUTE_ERROR_ON((src->data_layout() == DataLayout::NHWC) && has_padding_changed(padding_info));
diff --git a/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.cpp b/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.cpp
index 9eb249a66a..c5c24886bd 100644
--- a/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.cpp
+++ b/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2022 Arm Limited.
+ * Copyright (c) 2018-2023 Arm Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -102,7 +102,23 @@ std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITen
     ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
     ARM_COMPUTE_UNUSED(bias);
 
-    constexpr unsigned int num_elems_processed_per_iteration = 1;
+    unsigned int num_elems_processed_per_iteration = 1;
+
+    if(input->data_layout() == DataLayout::NHWC)
+    {
+        // In the case of FP16 computation, we can perform more
+        // output feature maps in a single work-item.
+        // From experiments, num_elems_processed_per_iteration = 2 looks good for fp16 to
+        // improve the performance. However, in order to make the implementation simpler,
+        // we set num_elems_processed_per_iteration = 2 only when the OFMs are multiple of 2.
+        const DataType dt   = input->data_type();
+        const size_t   dim0 = input->dimension(0);
+        const bool     cond = dt == DataType::F16 && ((dim0 % 2) == 0);
+        if(cond)
+        {
+            num_elems_processed_per_iteration = 2;
+        }
+    }
 
     Window win            = calculate_max_window(*input, Steps(num_elems_processed_per_iteration));
     bool   window_changed = false;
@@ -203,7 +219,7 @@ void ClWinogradOutputTransformKernel::configure(const ClCompileContext &compile_
         build_opts.add_option("-DOUTPUT_TILE_W=" + support::cpp11::to_string(output_tile_size.width));
         build_opts.add_option("-DOUTPUT_TILE_H=" + support::cpp11::to_string(output_tile_size.height));
         build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(src_data_type));
-        build_opts.add_option_if(total_batches > 1, "-DSRC_DEPTH=" + support::cpp11::to_string(src->dimension(2)));
+        build_opts.add_option_if(total_batches > 1, "-DIS_BATCHED");
         build_opts.add_option_if(winograd_info.kernel_size.height == 1, "-DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL");
         build_opts.add_option_if(winograd_info.kernel_size.width == 1, "-DWINOGRAD_OUTPUT_TRANSFORM_VERTICAL");
         build_opts.add_option("-DNUM_TILES_X=" + support::cpp11::to_string(_num_tiles_x));