Improve start-up time for winograd_input_transform_*_nhwc
- pass tensor's dimensions at runtime rather than compile time
- Add guard macro to compile only kernel(s) of internest
Resolves: COMPMID-5119
Signed-off-by: Ramy Elgammal <ramy.elgammal@arm.com>
Change-Id: Ib01098e397011a1201c2800c62a8954ec70e63e8
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/7083
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Giorgio Arena <giorgio.arena@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
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 4865982..ba7b13b 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-2021 Arm Limited.
+ * Copyright (c) 2018-2022 Arm Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -62,16 +62,16 @@
out.s7 = -36.0f * tmp.s1 + 0.0f * tmp.s2 + 49.0f * tmp.s3 - 14.0f * tmp.s5 + tmp.s7; \
})
-#if defined(NUM_TILES_X) && defined(PAD_LEFT) && defined(PAD_TOP) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H)
+#if defined(PAD_LEFT) && defined(PAD_TOP) && defined(OUTPUT_TILE_W) && defined(OUTPUT_TILE_H)
-#if defined(NHWC) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(NUM_TILES_X) && defined(NUM_TILES_Y)
+#if defined(NHWC)
+#if defined(WINOGRAD_INPUT_TRANSFORM_4X4_3X3_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_4X1_3X1_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_1X4_1X3_STEPZ1_NHWC)
//! @cond Doxygen_Suppress
/** This OpenCL kernel computes the input transform when the output tile is 4x4, 4x1 or 1x4, the filter size 3x3, 3x1 or 1x3 and the data layout is NHWC
*
* @note Data layout supported: NHWC
* @note Data type supported: F32/F16
* @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=half)
- * @note The number of tiles in the X and Y axes must be passed at compile time using -DNUM_TILES_X and -DNUM_TILES_Y (i.e.-DNUM_TILES_X=5, -DNUM_TILES_Y=3).
* @note The convolution padding (left and top) must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2)
* @note The spatial dimensions of the source tensor must be passed at compile time using -DSRC_WIDTH and -DSRC_HEIGHT (e.g. -DSRC_WIDTH=96, -DSRC_HEIGHT=64)
* @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
@@ -99,23 +99,24 @@
* @param[in] dst_stride_w Stride of the destination 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_WIDTH The src tensor's width
+ * @param[in] _ISRC_HEIGHT The src tensor's height
+ * @param[in] _INUM_TILES_X The number of tiles in the X dimension
+ * @param[in] _INUM_TILES_Y The number of tiles in the Y dimension
*/
//! @endcond
__kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc(
TENSOR4D(src, BUFFER),
- TENSOR4D(dst, BUFFER))
+ TENSOR4D(dst, BUFFER),
+ const int _ISRC_WIDTH,
+ const int _ISRC_HEIGHT,
+ 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 bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX
- // All the tensor dimensions are passed at compile time.
- // In case of dynamic tensor support, the following dimensions should be passed as function argument.
-#define _ISRC_WIDTH SRC_WIDTH
-#define _ISRC_HEIGHT SRC_HEIGHT
-#define _INUM_TILES_X NUM_TILES_X
-#define _INUM_TILES_Y NUM_TILES_Y
-
int x = (mout % _INUM_TILES_X) * OUTPUT_TILE_W;
int y = (mout / _INUM_TILES_X) * OUTPUT_TILE_H;
x -= PAD_LEFT;
@@ -234,14 +235,15 @@
T_STORE_INDIRECT_WIDTH_SELECT(DATA_TYPE, 36, 1, 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)
+#if defined(WINOGRAD_INPUT_TRANSFORM_4X4_5X5_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_4X1_5X1_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_1X4_1X5_STEPZ1_NHWC)
//! @cond Doxygen_Suppress
/** This OpenCL kernel computes the input transform when the kernel size is 5x5/5x1 or 1x5 and the output tile is 4x4/4x1 or 1x4 when the data layout is NHWC
*
* @note Data layout supported: NHWC
* @note Data type supported: F32/F16
* @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=half)
- * @note The number of tiles in the X and Y axes must be passed at compile time using -DNUM_TILES_X and -DNUM_TILES_Y (i.e.-DNUM_TILES_X=5, -DNUM_TILES_Y=3).
* @note The convolution padding (left and top) must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2)
* @note The spatial dimensions of the source tensor must be passed at compile time using -DSRC_WIDTH and -DSRC_HEIGHT (e.g. -DSRC_WIDTH=96, -DSRC_HEIGHT=64)
* @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
@@ -269,23 +271,24 @@
* @param[in] dst_stride_w Stride of the destination 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_WIDTH The src tensor's width
+ * @param[in] _ISRC_HEIGHT The src tensor's height
+ * @param[in] _INUM_TILES_X The number of tiles in the X dimension
+ * @param[in] _INUM_TILES_Y The number of tiles in the Y dimension
*/
//! @endcond
__kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc(
TENSOR4D(src, BUFFER),
- TENSOR4D(dst, BUFFER))
+ TENSOR4D(dst, BUFFER),
+ const int _ISRC_WIDTH,
+ const int _ISRC_HEIGHT,
+ 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 bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX
- // All the tensor dimensions are passed at compile time.
- // In case of dynamic tensor support, the following dimensions should be passed as function argument.
-#define _ISRC_WIDTH SRC_WIDTH
-#define _ISRC_HEIGHT SRC_HEIGHT
-#define _INUM_TILES_X NUM_TILES_X
-#define _INUM_TILES_Y NUM_TILES_Y
-
int x = (mout % _INUM_TILES_X) * OUTPUT_TILE_W;
int y = (mout / _INUM_TILES_X) * OUTPUT_TILE_H;
x -= PAD_LEFT;
@@ -403,14 +406,15 @@
#endif // !defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) && !defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL)
}
+#endif // defined(WINOGRAD_INPUT_TRANSFORM_4X4_5X5_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_4X1_5X1_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_1X4_1X5_STEPZ1_NHWC)
+#if defined(WINOGRAD_INPUT_TRANSFORM_2X2_7X7_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_2X1_7X1_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_1X2_1X7_STEPZ1_NHWC)
//! @cond Doxygen_Suppress
/** This OpenCL kernel computes the input transform when the kernel size is 7x7/7x1/1x7 and the output tile is 2x2/7x1/1x7 when the data layout is NHWC
*
* @note Data layout supported: NHWC
* @note Data type supported: F32/F16
* @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=half)
- * @note The number of tiles in the X and Y axes must be passed at compile time using -DNUM_TILES_X and -DNUM_TILES_Y (i.e.-DNUM_TILES_X=5, -DNUM_TILES_Y=3).
* @note The convolution padding (left and top) must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2)
* @note The spatial dimensions of the source tensor must be passed at compile time using -DSRC_WIDTH and -DSRC_HEIGHT (e.g. -DSRC_WIDTH=96, -DSRC_HEIGHT=64)
* @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
@@ -438,23 +442,24 @@
* @param[in] dst_stride_w Stride of the destination 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_WIDTH The src tensor's width
+ * @param[in] _ISRC_HEIGHT The src tensor's height
+ * @param[in] _INUM_TILES_X The number of tiles in the X dimension
+ * @param[in] _INUM_TILES_Y The number of tiles in the Y dimension
*/
//! @endcond
__kernel void winograd_input_transform_2x2_7x7_stepz1_nhwc(
TENSOR4D(src, BUFFER),
- TENSOR4D(dst, BUFFER))
+ TENSOR4D(dst, BUFFER),
+ const int _ISRC_WIDTH,
+ const int _ISRC_HEIGHT,
+ 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 bout = GET_SPATIAL_IDX(2, 1, 0); // BATCH SIZE IDX
- // All the tensor dimensions are passed at compile time.
- // In case of dynamic tensor support, the following dimensions should be passed as function argument.
-#define _ISRC_WIDTH SRC_WIDTH
-#define _ISRC_HEIGHT SRC_HEIGHT
-#define _INUM_TILES_X NUM_TILES_X
-#define _INUM_TILES_Y NUM_TILES_Y
-
int x = (mout % _INUM_TILES_X) * OUTPUT_TILE_W;
int y = (mout / _INUM_TILES_X) * OUTPUT_TILE_H;
x -= PAD_LEFT;
@@ -577,14 +582,15 @@
#endif // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL)
}
+#endif // defined(WINOGRAD_INPUT_TRANSFORM_2X2_7X7_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_2X1_7X1_STEPZ1_NHWC) || defined(WINOGRAD_INPUT_TRANSFORM_1X2_1X7_STEPZ1_NHWC)
+#if defined(WINOGRAD_INPUT_TRANSFORM_4X1_3X1_STEPZ1_NHWC)
//! @cond Doxygen_Suppress
/** This OpenCL kernel computes the input transform when the kernel size is 3x1 and the output tile is 4x1 for data layout NHWC
*
* @note Data layout supported: NHWC
* @note Data type supported: F32/F16
* @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=half)
- * @note The number of tiles in the X and Y axes must be passed at compile time using -DNUM_TILES_X and -DNUM_TILES_Y (i.e.-DNUM_TILES_X=5, -DNUM_TILES_Y=3).
* @note The convolution padding (left and top) must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2)
* @note The spatial dimensions of the source tensor must be passed at compile time using -DSRC_WIDTH and -DSRC_HEIGHT (e.g. -DSRC_WIDTH=96, -DSRC_HEIGHT=64)
* @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
@@ -612,11 +618,19 @@
* @param[in] dst_stride_w Stride of the destination 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_WIDTH The src tensor's width
+ * @param[in] _ISRC_HEIGHT The src tensor's height
+ * @param[in] _INUM_TILES_X The number of tiles in the X dimension
+ * @param[in] _INUM_TILES_Y The number of tiles in the Y dimension
*/
//! @endcond
__kernel void winograd_input_transform_4x1_3x1_stepz1_nhwc(
TENSOR4D(src, BUFFER),
- TENSOR4D(dst, BUFFER))
+ TENSOR4D(dst, BUFFER),
+ const int _ISRC_WIDTH,
+ const int _ISRC_HEIGHT,
+ const int _INUM_TILES_X,
+ const int _INUM_TILES_Y)
{
winograd_input_transform_4x4_3x3_stepz1_nhwc(src_ptr,
src_stride_x,
@@ -637,16 +651,21 @@
dst_step_z,
dst_stride_w,
dst_step_w,
- dst_offset_first_element_in_bytes);
+ dst_offset_first_element_in_bytes,
+ _ISRC_WIDTH,
+ _ISRC_HEIGHT,
+ _INUM_TILES_X,
+ _INUM_TILES_Y);
}
+#endif // defined(WINOGRAD_INPUT_TRANSFORM_4X1_3X1_STEPZ1_NHWC)
+#if defined(WINOGRAD_INPUT_TRANSFORM_4X1_5X1_STEPZ1_NHWC)
//! @cond Doxygen_Suppress
/** This OpenCL kernel computes the input transform when the kernel size is 5x1 and the output tile is 4x1 for data layout NHWC
*
* @note Data layout supported: NHWC
* @note Data type supported: F32/F16
* @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=half)
- * @note The number of tiles in the X and Y axes must be passed at compile time using -DNUM_TILES_X and -DNUM_TILES_Y (i.e.-DNUM_TILES_X=5, -DNUM_TILES_Y=3).
* @note The convolution padding (left and top) must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2)
* @note The spatial dimensions of the source tensor must be passed at compile time using -DSRC_WIDTH and -DSRC_HEIGHT (e.g. -DSRC_WIDTH=96, -DSRC_HEIGHT=64)
* @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
@@ -674,11 +693,19 @@
* @param[in] dst_stride_w Stride of the destination 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_WIDTH The src tensor's width
+ * @param[in] _ISRC_HEIGHT The src tensor's height
+ * @param[in] _INUM_TILES_X The number of tiles in the X dimension
+ * @param[in] _INUM_TILES_Y The number of tiles in the Y dimension
*/
//! @endcond
__kernel void winograd_input_transform_4x1_5x1_stepz1_nhwc(
TENSOR4D(src, BUFFER),
- TENSOR4D(dst, BUFFER))
+ TENSOR4D(dst, BUFFER),
+ const int _ISRC_WIDTH,
+ const int _ISRC_HEIGHT,
+ const int _INUM_TILES_X,
+ const int _INUM_TILES_Y)
{
winograd_input_transform_4x4_5x5_stepz1_nhwc(src_ptr,
src_stride_x,
@@ -699,16 +726,21 @@
dst_step_z,
dst_stride_w,
dst_step_w,
- dst_offset_first_element_in_bytes);
+ dst_offset_first_element_in_bytes,
+ _ISRC_WIDTH,
+ _ISRC_HEIGHT,
+ _INUM_TILES_X,
+ _INUM_TILES_Y);
}
+#endif // defined(WINOGRAD_INPUT_TRANSFORM_4X1_5X1_STEPZ1_NHWC)
+#if defined(WINOGRAD_INPUT_TRANSFORM_2X1_7X1_STEPZ1_NHWC)
//! @cond Doxygen_Suppress
/** This OpenCL kernel computes the input transform when the kernel size is 7x1 and the output tile is 2x1 for data layout NHWC
*
* @note Data layout supported: NHWC
* @note Data type supported: F32/F16
* @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=half)
- * @note The number of tiles in the X and Y axes must be passed at compile time using -DNUM_TILES_X and -DNUM_TILES_Y (i.e.-DNUM_TILES_X=5, -DNUM_TILES_Y=3).
* @note The convolution padding (left and top) must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2)
* @note The spatial dimensions of the source tensor must be passed at compile time using -DSRC_WIDTH and -DSRC_HEIGHT (e.g. -DSRC_WIDTH=96, -DSRC_HEIGHT=64)
* @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
@@ -736,11 +768,19 @@
* @param[in] dst_stride_w Stride of the destination 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_WIDTH The src tensor's width
+ * @param[in] _ISRC_HEIGHT The src tensor's height
+ * @param[in] _INUM_TILES_X The number of tiles in the X dimension
+ * @param[in] _INUM_TILES_Y The number of tiles in the Y dimension
*/
//! @endcond
__kernel void winograd_input_transform_2x1_7x1_stepz1_nhwc(
TENSOR4D(src, BUFFER),
- TENSOR4D(dst, BUFFER))
+ TENSOR4D(dst, BUFFER),
+ const int _ISRC_WIDTH,
+ const int _ISRC_HEIGHT,
+ const int _INUM_TILES_X,
+ const int _INUM_TILES_Y)
{
winograd_input_transform_2x2_7x7_stepz1_nhwc(src_ptr,
src_stride_x,
@@ -761,16 +801,22 @@
dst_step_z,
dst_stride_w,
dst_step_w,
- dst_offset_first_element_in_bytes);
+ dst_offset_first_element_in_bytes,
+ _ISRC_WIDTH,
+ _ISRC_HEIGHT,
+ _INUM_TILES_X,
+ _INUM_TILES_Y);
}
+#endif // defined(WINOGRAD_INPUT_TRANSFORM_2X1_7X1_STEPZ1_NHWC)
+#if defined(WINOGRAD_INPUT_TRANSFORM_1X4_1X3_STEPZ1_NHWC)
//! @cond Doxygen_Suppress
/** This OpenCL kernel computes the input transform when the kernel size is 1x3 and the output tile is 1x4 for data layout NHWC
*
* @note Data layout supported: NHWC
* @note Data type supported: F32/F16
* @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=half)
- * @note The number of tiles in the X and Y axes must be passed at compile time using -DNUM_TILES_X and -DNUM_TILES_Y (i.e.-DNUM_TILES_X=5, -DNUM_TILES_Y=3).
+ *
* @note The convolution padding (left and top) must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2)
* @note The spatial dimensions of the source tensor must be passed at compile time using -DSRC_WIDTH and -DSRC_HEIGHT (e.g. -DSRC_WIDTH=96, -DSRC_HEIGHT=64)
* @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
@@ -798,11 +844,19 @@
* @param[in] dst_stride_w Stride of the destination 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_WIDTH The src tensor's width
+ * @param[in] _ISRC_HEIGHT The src tensor's height
+ * @param[in] _INUM_TILES_X The number of tiles in the X dimension
+ * @param[in] _INUM_TILES_Y The number of tiles in the Y dimension
*/
//! @endcond
__kernel void winograd_input_transform_1x4_1x3_stepz1_nhwc(
TENSOR4D(src, BUFFER),
- TENSOR4D(dst, BUFFER))
+ TENSOR4D(dst, BUFFER),
+ const int _ISRC_WIDTH,
+ const int _ISRC_HEIGHT,
+ const int _INUM_TILES_X,
+ const int _INUM_TILES_Y)
{
winograd_input_transform_4x4_3x3_stepz1_nhwc(src_ptr,
src_stride_x,
@@ -823,16 +877,21 @@
dst_step_z,
dst_stride_w,
dst_step_w,
- dst_offset_first_element_in_bytes);
+ dst_offset_first_element_in_bytes,
+ _ISRC_WIDTH,
+ _ISRC_HEIGHT,
+ _INUM_TILES_X,
+ _INUM_TILES_Y);
}
+#endif // defined(WINOGRAD_INPUT_TRANSFORM_1X4_1X3_STEPZ1_NHWC)
+#if defined(WINOGRAD_INPUT_TRANSFORM_1X4_1X5_STEPZ1_NHWC)
//! @cond Doxygen_Suppress
/** This OpenCL kernel computes the input transform when the kernel size is 1x5 and the output tile is 1x4 for data layout NHWC
*
* @note Data layout supported: NHWC
* @note Data type supported: F32/F16
* @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=half)
- * @note The number of tiles in the X and Y axes must be passed at compile time using -DNUM_TILES_X and -DNUM_TILES_Y (i.e.-DNUM_TILES_X=5, -DNUM_TILES_Y=3).
* @note The convolution padding (left and top) must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2)
* @note The spatial dimensions of the source tensor must be passed at compile time using -DSRC_WIDTH and -DSRC_HEIGHT (e.g. -DSRC_WIDTH=96, -DSRC_HEIGHT=64)
* @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
@@ -860,11 +919,19 @@
* @param[in] dst_stride_w Stride of the destination 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_WIDTH The src tensor's width
+ * @param[in] _ISRC_HEIGHT The src tensor's height
+ * @param[in] _INUM_TILES_X The number of tiles in the X dimension
+ * @param[in] _INUM_TILES_Y The number of tiles in the Y dimension
*/
//! @endcond
__kernel void winograd_input_transform_1x4_1x5_stepz1_nhwc(
TENSOR4D(src, BUFFER),
- TENSOR4D(dst, BUFFER))
+ TENSOR4D(dst, BUFFER),
+ const int _ISRC_WIDTH,
+ const int _ISRC_HEIGHT,
+ const int _INUM_TILES_X,
+ const int _INUM_TILES_Y)
{
winograd_input_transform_4x4_5x5_stepz1_nhwc(src_ptr,
src_stride_x,
@@ -885,16 +952,21 @@
dst_step_z,
dst_stride_w,
dst_step_w,
- dst_offset_first_element_in_bytes);
+ dst_offset_first_element_in_bytes,
+ _ISRC_WIDTH,
+ _ISRC_HEIGHT,
+ _INUM_TILES_X,
+ _INUM_TILES_Y);
}
+#endif // defined(WINOGRAD_INPUT_TRANSFORM_1X4_1X5_STEPZ1_NHWC)
+#if defined(WINOGRAD_INPUT_TRANSFORM_1X2_1X7_STEPZ1_NHWC)
//! @cond Doxygen_Suppress
/** This OpenCL kernel computes the input transform when the kernel size is 1x7 and the output tile is 1x2 for data layout NHWC
*
* @note Data layout supported: NHWC
* @note Data type supported: F32/F16
* @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=half)
- * @note The number of tiles in the X and Y axes must be passed at compile time using -DNUM_TILES_X and -DNUM_TILES_Y (i.e.-DNUM_TILES_X=5, -DNUM_TILES_Y=3).
* @note The convolution padding (left and top) must be passed at compile time using -DPAD_LEFT and -DPAD_TOP (e.g. -DPAD_LEFT=2, -DPAD_TOP=2)
* @note The spatial dimensions of the source tensor must be passed at compile time using -DSRC_WIDTH and -DSRC_HEIGHT (e.g. -DSRC_WIDTH=96, -DSRC_HEIGHT=64)
* @note The width of the output tile must be passed at compile time using -DOUTPUT_TILE_W: e.g. -DOUTPUT_TILE_W=4
@@ -922,11 +994,19 @@
* @param[in] dst_stride_w Stride of the destination 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_WIDTH The src tensor's width
+ * @param[in] _ISRC_HEIGHT The src tensor's height
+ * @param[in] _INUM_TILES_X The number of tiles in the X dimension
+ * @param[in] _INUM_TILES_Y The number of tiles in the Y dimension
*/
//! @endcond
__kernel void winograd_input_transform_1x2_1x7_stepz1_nhwc(
TENSOR4D(src, BUFFER),
- TENSOR4D(dst, BUFFER))
+ TENSOR4D(dst, BUFFER),
+ const int _ISRC_WIDTH,
+ const int _ISRC_HEIGHT,
+ const int _INUM_TILES_X,
+ const int _INUM_TILES_Y)
{
winograd_input_transform_2x2_7x7_stepz1_nhwc(src_ptr,
src_stride_x,
@@ -947,7 +1027,12 @@
dst_step_z,
dst_stride_w,
dst_step_w,
- dst_offset_first_element_in_bytes);
+ dst_offset_first_element_in_bytes,
+ _ISRC_WIDTH,
+ _ISRC_HEIGHT,
+ _INUM_TILES_X,
+ _INUM_TILES_Y);
}
-#endif // defined(NHWC) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(NUM_TILES_X) && defined(NUM_TILES_Y)
+#endif // defined(WINOGRAD_INPUT_TRANSFORM_1X2_1X7_STEPZ1_NHWC)
+#endif // defined(NHWC)
#endif // defined(NUM_TILES_X) && defined(PAD_LEFT) && defined(PAD_TOP) && 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 5887421..d6b038f 100644
--- a/src/gpu/cl/kernels/ClWinogradInputTransformKernel.cpp
+++ b/src/gpu/cl/kernels/ClWinogradInputTransformKernel.cpp
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2018-2021 Arm Limited.
+ * Copyright (c) 2018-2022 Arm Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -147,10 +147,8 @@
if(_data_layout == DataLayout::NHWC)
{
build_opts.add_option("-DNHWC");
- build_opts.add_option("-DSRC_WIDTH=" + support::cpp11::to_string(src->dimension(idx_w)));
- build_opts.add_option("-DSRC_HEIGHT=" + support::cpp11::to_string(src->dimension(idx_h)));
- build_opts.add_option("-DNUM_TILES_X=" + support::cpp11::to_string(_num_tiles_x));
- build_opts.add_option("-DNUM_TILES_Y=" + support::cpp11::to_string(_num_tiles_y));
+ _src_width = src->dimension(idx_w);
+ _src_height = src->dimension(idx_h);
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));
@@ -189,6 +187,8 @@
kernel_name += support::cpp11::to_string(_step_z);
kernel_name += "_" + lower_string(string_from_data_layout(_data_layout));
+ // 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());
// Create window and update padding
@@ -247,6 +247,10 @@
unsigned int idx = 0;
add_4D_tensor_argument(idx, src, slice);
add_4D_tensor_argument(idx, dst, slice);
+ _kernel.setArg<cl_uint>(idx++, _src_width);
+ _kernel.setArg<cl_uint>(idx++, _src_height);
+ _kernel.setArg<cl_uint>(idx++, _num_tiles_x);
+ _kernel.setArg<cl_uint>(idx++, _num_tiles_y);
enqueue(queue, *this, slice, lws_hint());
}
else
@@ -275,4 +279,4 @@
}
} // namespace kernels
} // namespace opencl
-} // namespace arm_compute
\ No newline at end of file
+} // namespace arm_compute
diff --git a/src/gpu/cl/kernels/ClWinogradInputTransformKernel.h b/src/gpu/cl/kernels/ClWinogradInputTransformKernel.h
index 631f427..c10c528 100644
--- a/src/gpu/cl/kernels/ClWinogradInputTransformKernel.h
+++ b/src/gpu/cl/kernels/ClWinogradInputTransformKernel.h
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2018-2021 Arm Limited.
+ * Copyright (c) 2018-2022 Arm Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -80,6 +80,8 @@
int _num_tiles_x{ 0 };
int _num_tiles_y{ 0 };
unsigned int _step_z{ 1 };
+ int32_t _src_width{ 0 };
+ int32_t _src_height{ 0 };
};
} // namespace kernels
} // namespace opencl