Improve start-up time for winograd_output_transform_*_nhwc

 - pass tensor's dimensions at runtime rather than compile time
 - Add guard macro to compile only kernel of internest

Resolves: COMPMID-5120
Signed-off-by: Ramy Elgammal <ramy.elgammal@arm.com>
Change-Id: I87c3b56ce0cd3c97ffdeabdd9c5d433f361bb005
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/7101
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>

3 files changed, 182 insertions, 59 deletions

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 0fcd04e713..ed6da9fd12 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-2021 Arm Limited.
+ * Copyright (c) 2018-2022 Arm Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -25,16 +25,14 @@
 #include "helpers.h"
 #include "tile_helpers.h"
 
-#if defined(NUM_TILES_X) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H)
+#if defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H)
 #if defined(VEC_SIZE) && VEC_SIZE == 2
+#if defined(WINOGRAD_OUTPUT_TRANSFORM_2X2_7X7_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_2X1_7X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X2_1X7_NHWC)
 /** This OpenCL kernel performs Winograd output transform when the output tile is 2x2/2x1 or 1x2, the filter size 7x7/7x1 or 1x7 and the data layout is NHWC
  *
- * @note The number of tiles along the X direction must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16
+ * @note  must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16
  * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=2
  * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=2
- * @note The height of the input tensor must be passed at compile time using -DSRC_HEIGHT: e.g. -DSRC_HEIGHT=32
- * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24
- * @note The height of the output tensor must be passed at compile time using -DDST_HEIGHT: e.g. -DDST_HEIGHT=32
  * @note If this kernel is used to perform Winograd output transform 7x1, -DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL has to be passed at compile time
  * @note If this kernel is used to perform Winograd output transform 1x7, -DWINOGRAD_OUTPUT_TRANSFORM_VERTICAL has to be passed at compile time
  * @note The data type must be passed at compile time using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types: float/half.
@@ -60,6 +58,10 @@
  * @param[in]  dst_stride_w                      Stride of the source tensor in W dimension (in bytes)
  * @param[in]  dst_step_w                        dst_stride_w * number of elements along W processed per workitem(in bytes)
  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in]  _ISRC_HEIGHT                      The source tensor's height
+ * @param[in]  _IDST_WIDTH                       The destination tensor's width
+ * @param[in]  _IDST_HEIGHT                      The destination tensor's height
+ * @param[in]  _INUM_TILES_X                     The number of tiles along the X direction
  */
 __kernel void winograd_output_transform_2x2_7x7_nhwc(
     TENSOR4D(src, BUFFER),
@@ -67,13 +69,12 @@ __kernel void winograd_output_transform_2x2_7x7_nhwc(
 #if defined(HAS_BIAS)
     VECTOR_DECLARATION(bias),
 #endif // defined(HAS_BIAS)
-    int dst_size)
+    int       dst_size,
+    const int _ISRC_HEIGHT,
+    const int _IDST_WIDTH,
+    const int _IDST_HEIGHT,
+    const int _INUM_TILES_X)
 {
-#define _ISRC_HEIGHT SRC_HEIGHT
-#define _IDST_WIDTH DST_WIDTH
-#define _IDST_HEIGHT DST_HEIGHT
-#define _INUM_TILES_X NUM_TILES_X
-
     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
@@ -201,17 +202,15 @@ __kernel void winograd_output_transform_2x2_7x7_nhwc(
     T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 4, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y);
 #endif // !defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL)
 }
+#endif // defined(WINOGRAD_OUTPUT_TRANSFORM_2X2_7X7_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_2X1_7X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X2_1X7_NHWC)
 #endif // defined(VEC_SIZE) && VEC_SIZE == 2
 
 #if defined(VEC_SIZE) && VEC_SIZE == 4
+#if defined(WINOGRAD_OUTPUT_TRANSFORM_4X4_3X3_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_3X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X3_NHWC)
 /** This OpenCL kernel performs Winograd output transform when the output tile is 4x4, 4x1 or 1x4, the filter size 3x3, 3x1 or 1x3 and the data layout is NHWC
  *
- * @note The number of tiles along the X direction must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16
  * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
  * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=4
- * @note The height of the input tensor must be passed at compile time using -DSRC_HEIGHT: e.g. -DSRC_HEIGHT=32
- * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24
- * @note The height of the output tensor must be passed at compile time using -DDST_HEIGHT: e.g. -DDST_HEIGHT=32
  * @note If this kernel is used to perform Winograd output transform 3x1, -DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL has to be passed at compile time
  * @note If this kernel is used to perform Winograd output transform 1x3, -DWINOGRAD_OUTPUT_TRANSFORM_VERTICAL has to be passed at compile time
  * @note The data type must be passed at compile time using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types: float/half.
@@ -238,6 +237,10 @@ __kernel void winograd_output_transform_2x2_7x7_nhwc(
  * @param[in]  dst_step_w                        dst_stride_w * number of elements along W processed per workitem(in bytes)
  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
  * @param[in]  dst_size                          Size of the destination tensor, minus the last padding
+ * @param[in]  SRC_HEIGHT                        The source tensor's height
+ * @param[in]  DST_WIDTH                         The destination tensor's width
+ * @param[in]  DST_HEIGHT                        The destination tensor's height
+ * @param[in]  NUM_TILES_X                       The number of tiles along the X direction
  */
 __kernel void winograd_output_transform_4x4_3x3_nhwc(
     TENSOR4D(src, BUFFER),
@@ -245,7 +248,11 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc(
 #if defined(HAS_BIAS)
     VECTOR_DECLARATION(bias),
 #endif // defined(HAS_BIAS)
-    int dst_size)
+    int       dst_size,
+    const int SRC_HEIGHT,
+    const int DST_WIDTH,
+    const int DST_HEIGHT,
+    const int NUM_TILES_X)
 {
     const int cout = GET_SPATIAL_IDX(0, N0, 0); // OFM
     const int mout = GET_SPATIAL_IDX(1, 1, 0);  // WINOGRAD OUTPUT TILES
@@ -397,15 +404,13 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc(
     T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 16, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y);
 #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL)
 }
+#endif // defined(WINOGRAD_OUTPUT_TRANSFORM_4X4_3X3_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_3X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X3_NHWC)
 
+#if defined(WINOGRAD_OUTPUT_TRANSFORM_4X4_5X5_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_5X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X5_NHWC)
 /** This OpenCL kernel performs Winograd output transform when the output tile is 4x4/4x1 or 1x4, the filter size 5x5/5x1 or 1x5 and the data layout is NHWC
  *
- * @note The number of tiles along the X direction must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16
  * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
  * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=4
- * @note The height of the input tensor must be passed at compile time using -DSRC_HEIGHT: e.g. -DSRC_HEIGHT=32
- * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24
- * @note The height of the output tensor must be passed at compile time using -DDST_HEIGHT: e.g. -DDST_HEIGHT=32
  * @note If this kernel is used to perform Winograd output transform 5x1, -DWINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL has to be passed at compile time
  * @note If this kernel is used to perform Winograd output transform 1x5, -DWINOGRAD_OUTPUT_TRANSFORM_VERTICAL has to be passed at compile time
  * @note The data type must be passed at compile time using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types: float/half.
@@ -431,6 +436,10 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc(
  * @param[in]  dst_stride_w                      Stride of the source tensor in W dimension (in bytes)
  * @param[in]  dst_step_w                        dst_stride_w * number of elements along W processed per workitem(in bytes)
  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in]  SRC_HEIGHT                        The source tensor's height
+ * @param[in]  DST_WIDTH                         The destination tensor's width
+ * @param[in]  DST_HEIGHT                        The destination tensor's height
+ * @param[in]  NUM_TILES_X                       The number of tiles along the X direction
  */
 __kernel void winograd_output_transform_4x4_5x5_nhwc(
     TENSOR4D(src, BUFFER),
@@ -438,7 +447,11 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc(
 #if defined(HAS_BIAS)
     VECTOR_DECLARATION(bias),
 #endif // defined(HAS_BIAS)
-    int dst_size)
+    int       dst_size,
+    const int SRC_HEIGHT,
+    const int DST_WIDTH,
+    const int DST_HEIGHT,
+    const int NUM_TILES_X)
 {
     const int cout = GET_SPATIAL_IDX(0, N0, 0); // OFM
     const int mout = GET_SPATIAL_IDX(1, 1, 0);  // WINOGRAD OUTPUT TILES
@@ -605,13 +618,14 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc(
     T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 16, N0, 0, BUFFER, dst, cout, dst_stride_y, false, out, dst_indirect_y);
 #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL)
 }
+#endif // defined(WINOGRAD_OUTPUT_TRANSFORM_4X4_5X5_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_5X1_NHWC) || defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X5_NHWC)
 #endif // defined(VEC_SIZE) && VEC_SIZE == 4
 
 #if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL)
 #if defined(VEC_SIZE) && VEC_SIZE == 2
+#if defined(WINOGRAD_OUTPUT_TRANSFORM_2X1_7X1_NHWC)
 /** This OpenCL kernel performs Winograd output transform when the output tile is 2x1, the filter size 7x1 and the data layout is NHWC
  *
- * @note The number of tiles along the X direction must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16
  * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=2
  * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=1
  * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24
@@ -639,6 +653,10 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc(
  * @param[in]  dst_stride_w                      Stride of the source tensor in W dimension (in bytes)
  * @param[in]  dst_step_w                        dst_stride_w * number of elements along W processed per workitem(in bytes)
  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in]  SRC_HEIGHT                        The source tensor's height
+ * @param[in]  DST_WIDTH                         The destination tensor's width
+ * @param[in]  DST_HEIGHT                        The destination tensor's height
+ * @param[in]  NUM_TILES_X                       The number of tiles along the X direction
  */
 __kernel void winograd_output_transform_2x1_7x1_nhwc(
     TENSOR4D_DECLARATION(src),
@@ -646,7 +664,11 @@ __kernel void winograd_output_transform_2x1_7x1_nhwc(
 #if defined(HAS_BIAS)
     VECTOR_DECLARATION(bias),
 #endif // defined(HAS_BIAS)
-    int dst_size)
+    int       dst_size,
+    const int SRC_HEIGHT,
+    const int DST_WIDTH,
+    const int DST_HEIGHT,
+    const int NUM_TILES_X)
 {
     winograd_output_transform_2x2_7x7_nhwc(src_ptr,
                                            src_stride_x,
@@ -674,15 +696,19 @@ __kernel void winograd_output_transform_2x1_7x1_nhwc(
                                            bias_step_x,
                                            bias_offset_first_element_in_bytes,
 #endif // defined(HAS_BIAS)
-                                           dst_size);
+                                           dst_size,
+                                           SRC_HEIGHT,
+                                           DST_WIDTH,
+                                           DST_HEIGHT,
+                                           NUM_TILES_X);
 }
+#endif // defined(WINOGRAD_OUTPUT_TRANSFORM_2X1_7X1_NHWC)
 #endif // defined(VEC_SIZE) && VEC_SIZE == 2
 
 #if defined(VEC_SIZE) && VEC_SIZE == 4
-
+#if defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_3X1_NHWC)
 /** This OpenCL kernel performs Winograd output transform when the output tile is 4x1, the filter size 3x1 and the data layout is NHWC
  *
- * @note The number of tiles along the X direction must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16
  * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
  * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=1
  * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24
@@ -710,6 +736,10 @@ __kernel void winograd_output_transform_2x1_7x1_nhwc(
  * @param[in]  dst_stride_w                      Stride of the source tensor in W dimension (in bytes)
  * @param[in]  dst_step_w                        dst_stride_w * number of elements along W processed per workitem(in bytes)
  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in]  SRC_HEIGHT                        The source tensor's height
+ * @param[in]  DST_WIDTH                         The destination tensor's width
+ * @param[in]  DST_HEIGHT                        The destination tensor's height
+ * @param[in]  NUM_TILES_X                       The number of tiles along the X direction
  */
 __kernel void winograd_output_transform_4x1_3x1_nhwc(
     TENSOR4D_DECLARATION(src),
@@ -717,7 +747,11 @@ __kernel void winograd_output_transform_4x1_3x1_nhwc(
 #if defined(HAS_BIAS)
     VECTOR_DECLARATION(bias),
 #endif // defined(HAS_BIAS)
-    int dst_size)
+    int       dst_size,
+    const int SRC_HEIGHT,
+    const int DST_WIDTH,
+    const int DST_HEIGHT,
+    const int NUM_TILES_X)
 {
     winograd_output_transform_4x4_3x3_nhwc(src_ptr,
                                            src_stride_x,
@@ -745,12 +779,17 @@ __kernel void winograd_output_transform_4x1_3x1_nhwc(
                                            bias_step_x,
                                            bias_offset_first_element_in_bytes,
 #endif // defined(HAS_BIAS)
-                                           dst_size);
+                                           dst_size,
+                                           SRC_HEIGHT,
+                                           DST_WIDTH,
+                                           DST_HEIGHT,
+                                           NUM_TILES_X);
 }
+#endif // defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_3X1_NHWC)
 
+#if defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_5X1_NHWC)
 /** This OpenCL kernel performs Winograd output transform when the output tile is 4x1, the filter size 5x1 and the data layout is NHWC
  *
- * @note The number of tiles along the X direction must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16
  * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
  * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=1
  * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24
@@ -778,6 +817,10 @@ __kernel void winograd_output_transform_4x1_3x1_nhwc(
  * @param[in]  dst_stride_w                      Stride of the source tensor in W dimension (in bytes)
  * @param[in]  dst_step_w                        dst_stride_w * number of elements along W processed per workitem(in bytes)
  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in]  SRC_HEIGHT                        The source tensor's height
+ * @param[in]  DST_WIDTH                         The destination tensor's width
+ * @param[in]  DST_HEIGHT                        The destination tensor's height
+ * @param[in]  NUM_TILES_X                       The number of tiles along the X direction
  */
 __kernel void winograd_output_transform_4x1_5x1_nhwc(
     TENSOR4D_DECLARATION(src),
@@ -785,7 +828,11 @@ __kernel void winograd_output_transform_4x1_5x1_nhwc(
 #if defined(HAS_BIAS)
     VECTOR_DECLARATION(bias),
 #endif // defined(HAS_BIAS)
-    int dst_size)
+    int       dst_size,
+    const int SRC_HEIGHT,
+    const int DST_WIDTH,
+    const int DST_HEIGHT,
+    const int NUM_TILES_X)
 {
     winograd_output_transform_4x4_5x5_nhwc(src_ptr,
                                            src_stride_x,
@@ -813,16 +860,21 @@ __kernel void winograd_output_transform_4x1_5x1_nhwc(
                                            bias_step_x,
                                            bias_offset_first_element_in_bytes,
 #endif // defined(HAS_BIAS)
-                                           dst_size);
+                                           dst_size,
+                                           SRC_HEIGHT,
+                                           DST_WIDTH,
+                                           DST_HEIGHT,
+                                           NUM_TILES_X);
 }
+#endif // defined(WINOGRAD_OUTPUT_TRANSFORM_4X1_5X1_NHWC)
 #endif // defined(VEC_SIZE) && VEC_SIZE == 4
 #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL)
 
 #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL)
 #if defined(VEC_SIZE) && VEC_SIZE == 2
+#if defined(WINOGRAD_OUTPUT_TRANSFORM_1X2_1X7_NHWC)
 /** This OpenCL kernel performs Winograd output transform when the output tile is 1x2, the filter size 1x7 and the data layout is NHWC
  *
- * @note The number of tiles along the X direction must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16
  * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=1
  * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=2
  * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24
@@ -850,6 +902,10 @@ __kernel void winograd_output_transform_4x1_5x1_nhwc(
  * @param[in]  dst_stride_w                      Stride of the source tensor in W dimension (in bytes)
  * @param[in]  dst_step_w                        dst_stride_w * number of elements along W processed per workitem(in bytes)
  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in]  SRC_HEIGHT                        The source tensor's height
+ * @param[in]  DST_WIDTH                         The destination tensor's width
+ * @param[in]  DST_HEIGHT                        The destination tensor's height
+ * @param[in]  NUM_TILES_X                       The number of tiles along the X direction
  */
 __kernel void winograd_output_transform_1x2_1x7_nhwc(
     TENSOR4D_DECLARATION(src),
@@ -857,7 +913,11 @@ __kernel void winograd_output_transform_1x2_1x7_nhwc(
 #if defined(HAS_BIAS)
     VECTOR_DECLARATION(bias),
 #endif // defined(HAS_BIAS)
-    int dst_size)
+    int       dst_size,
+    const int SRC_HEIGHT,
+    const int DST_WIDTH,
+    const int DST_HEIGHT,
+    const int NUM_TILES_X)
 {
     winograd_output_transform_2x2_7x7_nhwc(src_ptr,
                                            src_stride_x,
@@ -885,14 +945,19 @@ __kernel void winograd_output_transform_1x2_1x7_nhwc(
                                            bias_step_x,
                                            bias_offset_first_element_in_bytes,
 #endif // defined(HAS_BIAS)
-                                           dst_size);
+                                           dst_size,
+                                           SRC_HEIGHT,
+                                           DST_WIDTH,
+                                           DST_HEIGHT,
+                                           NUM_TILES_X);
 }
+#endif // defined(WINOGRAD_OUTPUT_TRANSFORM_1X2_1X7_NHWC)
 #endif // defined(VEC_SIZE) && VEC_SIZE == 2
 
 #if defined(VEC_SIZE) && VEC_SIZE == 4
+#if defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X3_NHWC)
 /** This OpenCL kernel performs Winograd output transform when the output tile is 1x4, the filter size 1x3 and the data layout is NHWC
  *
- * @note The number of tiles along the X direction must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16
  * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=1
  * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=4
  * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24
@@ -920,6 +985,10 @@ __kernel void winograd_output_transform_1x2_1x7_nhwc(
  * @param[in]  dst_stride_w                      Stride of the source tensor in W dimension (in bytes)
  * @param[in]  dst_step_w                        dst_stride_w * number of elements along W processed per workitem(in bytes)
  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in]  SRC_HEIGHT                        The source tensor's height
+ * @param[in]  DST_WIDTH                         The destination tensor's width
+ * @param[in]  DST_HEIGHT                        The destination tensor's height
+ * @param[in]  NUM_TILES_X                       The number of tiles along the X direction
  */
 __kernel void winograd_output_transform_1x4_1x3_nhwc(
     TENSOR4D_DECLARATION(src),
@@ -927,7 +996,11 @@ __kernel void winograd_output_transform_1x4_1x3_nhwc(
 #if defined(HAS_BIAS)
     VECTOR_DECLARATION(bias),
 #endif // defined(HAS_BIAS)
-    int dst_size)
+    int       dst_size,
+    const int SRC_HEIGHT,
+    const int DST_WIDTH,
+    const int DST_HEIGHT,
+    const int NUM_TILES_X)
 {
     winograd_output_transform_4x4_3x3_nhwc(src_ptr,
                                            src_stride_x,
@@ -955,12 +1028,17 @@ __kernel void winograd_output_transform_1x4_1x3_nhwc(
                                            bias_step_x,
                                            bias_offset_first_element_in_bytes,
 #endif // defined(HAS_BIAS)
-                                           dst_size);
+                                           dst_size,
+                                           SRC_HEIGHT,
+                                           DST_WIDTH,
+                                           DST_HEIGHT,
+                                           NUM_TILES_X);
 }
+#endif // defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X3_NHWC)
 
+#if defined(WINOGRAD_OUTPUT_TRANSFORM_1X4_1X5_NHWC)
 /** This OpenCL kernel performs Winograd output transform when the output tile is 1x4, the filter size 1x5 and the data layout is NHWC
  *
- * @note The number of tiles along the X direction must be passed at compile time using -DNUM_TILES_X: e.g. -DNUM_TILES_X=16
  * @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=1
  * @note The height of the output tile must be passed at compile time using -DOUTPUT_TILE_H: e.g. -DOUTPUT_TILE_H=4
  * @note The width of the output tensor must be passed at compile time using -DDST_WIDTH: e.g. -DDST_WIDTH=24
@@ -988,6 +1066,10 @@ __kernel void winograd_output_transform_1x4_1x3_nhwc(
  * @param[in]  dst_stride_w                      Stride of the source tensor in W dimension (in bytes)
  * @param[in]  dst_step_w                        dst_stride_w * number of elements along W processed per workitem(in bytes)
  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
+ * @param[in]  SRC_HEIGHT                        The source tensor's height
+ * @param[in]  DST_WIDTH                         The destination tensor's width
+ * @param[in]  DST_HEIGHT                        The destination tensor's height
+ * @param[in]  NUM_TILES_X                       The number of tiles along the X direction
  */
 __kernel void winograd_output_transform_1x4_1x5_nhwc(
     TENSOR4D_DECLARATION(src),
@@ -995,7 +1077,11 @@ __kernel void winograd_output_transform_1x4_1x5_nhwc(
 #if defined(HAS_BIAS)
     VECTOR_DECLARATION(bias),
 #endif // defined(HAS_BIAS)
-    int dst_size)
+    int       dst_size,
+    const int SRC_HEIGHT,
+    const int DST_WIDTH,
+    const int DST_HEIGHT,
+    const int NUM_TILES_X)
 {
     winograd_output_transform_4x4_5x5_nhwc(src_ptr,
                                            src_stride_x,
@@ -1023,8 +1109,13 @@ __kernel void winograd_output_transform_1x4_1x5_nhwc(
                                            bias_step_x,
                                            bias_offset_first_element_in_bytes,
 #endif // defined(HAS_BIAS)
-                                           dst_size);
+                                           dst_size,
+                                           SRC_HEIGHT,
+                                           DST_WIDTH,
+                                           DST_HEIGHT,
+                                           NUM_TILES_X);
 }
+#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
diff --git a/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.cpp b/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.cpp
index a8cf8234ad..ff57c83959 100644
--- a/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.cpp
+++ b/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2021 Arm Limited.
+ * Copyright (c) 2018-2022 Arm Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -176,23 +176,47 @@ void ClWinogradOutputTransformKernel::configure(const ClCompileContext &compile_
         build_opts.add_option("-DVEC_SIZE=4");
     }
 
-    build_opts.add_option_if(bias != nullptr, std::string("-DHAS_BIAS"));
-    build_opts.add_option("-cl-fast-relaxed-math");
-    build_opts.add_option("-DN0=" + support::cpp11::to_string(win_config.second.x().step()));
-    build_opts.add_option("-DNUM_TILES_X=" + support::cpp11::to_string(num_tiles.width));
-    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("-DSRC_HEIGHT=" + support::cpp11::to_string(src->dimension(1)));
-    build_opts.add_option("-DDST_WIDTH=" + support::cpp11::to_string(dst->dimension(idx_width)));
-    build_opts.add_option("-DDST_HEIGHT=" + support::cpp11::to_string(dst->dimension(idx_height)));
-    build_opts.add_option_if(total_batches > 1, "-DSRC_DEPTH=" + support::cpp11::to_string(src->dimension(2)));
-    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");
+    if(_is_nhwc)
+    {
+        build_opts.add_option_if(bias != nullptr, std::string("-DHAS_BIAS"));
+        build_opts.add_option("-cl-fast-relaxed-math");
+        build_opts.add_option("-DN0=" + support::cpp11::to_string(win_config.second.x().step()));
+        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(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");
+    }
+    else
+    {
+        build_opts.add_option_if(bias != nullptr, std::string("-DHAS_BIAS"));
+        build_opts.add_option("-cl-fast-relaxed-math");
+        build_opts.add_option("-DN0=" + support::cpp11::to_string(win_config.second.x().step()));
+        build_opts.add_option("-DNUM_TILES_X=" + support::cpp11::to_string(num_tiles.width));
+        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("-DSRC_HEIGHT=" + support::cpp11::to_string(src->dimension(1)));
+        build_opts.add_option("-DDST_WIDTH=" + support::cpp11::to_string(dst->dimension(idx_width)));
+        build_opts.add_option("-DDST_HEIGHT=" + support::cpp11::to_string(dst->dimension(idx_height)));
+        build_opts.add_option_if(total_batches > 1, "-DSRC_DEPTH=" + support::cpp11::to_string(src->dimension(2)));
+        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");
+    }
+
+    // Storing tensor dimensions to be sent later as kernel arguments
+    _src_height  = src->dimension(1);
+    _dst_width   = dst->dimension(idx_width);
+    _dst_height  = dst->dimension(idx_height);
+    _num_tiles_x = num_tiles.width;
 
     // Create kernel
     std::string kernel_name = "winograd_output_transform_" + output_tile_size.to_string() + "_" + kernel_size.to_string() + "_" + lower_string(string_from_data_layout(winograd_info.output_data_layout));
-    _kernel                 = create_kernel(compile_context, kernel_name, build_opts.options());
+
+    // A macro guard to compile ONLY the kernel of interest
+    build_opts.add_option("-D" + upper_string(kernel_name));
+    _kernel = create_kernel(compile_context, kernel_name, build_opts.options());
 
     // Set config_id for enabling LWS tuning
     _config_id = kernel_name;
@@ -251,7 +275,11 @@ void ClWinogradOutputTransformKernel::run_op(ITensorPack &tensors, const Window
     if(_is_nhwc)
     {
         unsigned int idx2 = 2 * num_arguments_per_4D_tensor() + ((bias != nullptr) ? num_arguments_per_1D_tensor() : 0);
-        _kernel.setArg(idx2, static_cast<int>(dst->info()->total_size() - dst->info()->strides_in_bytes().y()));
+        _kernel.setArg(idx2++, static_cast<int>(dst->info()->total_size() - dst->info()->strides_in_bytes().y()));
+        _kernel.setArg<cl_int>(idx2++, _src_height);
+        _kernel.setArg<cl_int>(idx2++, _dst_width);
+        _kernel.setArg<cl_int>(idx2++, _dst_height);
+        _kernel.setArg<cl_int>(idx2++, _num_tiles_x);
     }
 
     do
diff --git a/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.h b/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.h
index 674d52c904..6f018967d0 100644
--- a/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.h
+++ b/src/gpu/cl/kernels/ClWinogradOutputTransformKernel.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2021 Arm Limited.
+ * Copyright (c) 2018-2022 Arm Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -77,7 +77,11 @@ public:
 private:
     using WinogradKey = std::pair<std::pair<int, int>, std::pair<int, int>>;
 
-    bool _is_nhwc{ false };
+    bool    _is_nhwc{ false };
+    int32_t _src_height{ 0 };
+    int32_t _dst_width{ 0 };
+    int32_t _dst_height{ 0 };
+    int32_t _num_tiles_x{ 0 };
 };
 } // namespace kernels
 } // namespace opencl