From 6dbcc0e4d2fd0c61602a1a0c4a0ac548da713087 Mon Sep 17 00:00:00 2001 From: Sheri Zhang Date: Mon, 12 Apr 2021 10:53:57 +0100 Subject: Fix OpenCL kernel compiling failure with array initializer The issue is related with clang version, clang 3.9 has the problem, clange 4.0 works. The workaround is to add an extra {} to make this work. Partial resolves: COMPMID-4348 Signed-off-by: Sheri Zhang Change-Id: Ia079cbb3c44d617b1b42cb2af758b5a8ba1a032e Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5399 Tested-by: Arm Jenkins Reviewed-by: Manuel Bottini Comments-Addressed: Arm Jenkins --- src/core/CL/cl_kernels/winograd_input_transform.cl | 66 +++++++++++----------- 1 file changed, 33 insertions(+), 33 deletions(-) (limited to 'src/core/CL/cl_kernels/winograd_input_transform.cl') diff --git a/src/core/CL/cl_kernels/winograd_input_transform.cl b/src/core/CL/cl_kernels/winograd_input_transform.cl index 93ce878def..60750f054d 100644 --- a/src/core/CL/cl_kernels/winograd_input_transform.cl +++ b/src/core/CL/cl_kernels/winograd_input_transform.cl @@ -967,10 +967,10 @@ __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) = { { 0 } }; - TILE(DATA_TYPE, 6, 1, out) = { { 0 } }; - TILE(int, 6, 1, src_indirect_y) = { { 0 } }; - TILE(int, 6, 1, src_indirect_mask) = { { 0 } }; + TILE(DATA_TYPE, 6, 1, in) = { { { 0 } } }; + TILE(DATA_TYPE, 6, 1, out) = { { { 0 } } }; + TILE(int, 6, 1, src_indirect_y) = { { { 0 } } }; + TILE(int, 6, 1, src_indirect_mask) = { { { 0 } } }; #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); @@ -978,7 +978,7 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc( T_LOAD_NHWC(DATA_TYPE, 6, 1, 1, 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) = { { 0 } }; + TILE(DATA_TYPE, 6, 1, com) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 6, 1) { @@ -1000,7 +1000,7 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc( out[4].v = com[5].v - com[4].v; out[5].v = com[3].v - com[1].v; - TILE(uint, 6, 1, dst_indirect_y) = { { 0 } }; + TILE(uint, 6, 1, dst_indirect_y) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 6, 1) { @@ -1012,13 +1012,13 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc( #else // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - TILE(DATA_TYPE, 36, 1, in) = { { 0 } }; + TILE(DATA_TYPE, 36, 1, in) = { { { 0 } } }; // 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); - TILE(DATA_TYPE, 6, 1, com) = { { 0 } }; - TILE(DATA_TYPE, 36, 1, tmp) = { { 0 } }; + TILE(DATA_TYPE, 6, 1, com) = { { { 0 } } }; + TILE(DATA_TYPE, 36, 1, tmp) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 6, 1) { @@ -1037,7 +1037,7 @@ __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) = { { 0 } }; + TILE(DATA_TYPE, 36, 1, out) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 6, 1) { @@ -1057,7 +1057,7 @@ __kernel void winograd_input_transform_4x4_3x3_stepz1_nhwc( } // Compute destination address - TILE(uint, 36, 1, dst_indirect_y) = { { 0 } }; + TILE(uint, 36, 1, dst_indirect_y) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 36, 1) { @@ -1127,10 +1127,10 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc( #if defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - TILE(DATA_TYPE, 8, 1, in) = { { 0 } }; - TILE(DATA_TYPE, 8, 1, out) = { { 0 } }; - TILE(int, 8, 1, src_indirect_y) = { { 0 } }; - TILE(int, 8, 1, src_indirect_mask) = { { 0 } }; + TILE(DATA_TYPE, 8, 1, in) = { { { 0 } } }; + TILE(DATA_TYPE, 8, 1, out) = { { { 0 } } }; + TILE(int, 8, 1, src_indirect_y) = { { { 0 } } }; + TILE(int, 8, 1, src_indirect_mask) = { { { 0 } } }; #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); @@ -1138,7 +1138,7 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc( T_LOAD_NHWC(DATA_TYPE, 8, 1, 1, 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) = { { 0 } }; + TILE(DATA_TYPE, 1, 8, com) = { { { 0 } } }; 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; @@ -1155,7 +1155,7 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc( 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; - TILE(uint, 8, 1, dst_indirect_y) = { { 0 } }; + TILE(uint, 8, 1, dst_indirect_y) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 8, 1) { @@ -1167,13 +1167,13 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc( #else // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - TILE(DATA_TYPE, 64, 1, in) = { { 0 } }; - TILE(DATA_TYPE, 64, 1, out) = { { 0 } }; + TILE(DATA_TYPE, 64, 1, in) = { { { 0 } } }; + TILE(DATA_TYPE, 64, 1, out) = { { { 0 } } }; // Load the tile from a NHWC tensor T_LOAD_NHWC(DATA_TYPE, 8, 8, 1, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in); - TILE(DATA_TYPE, 8, 8, com) = { { 0 } }; + TILE(DATA_TYPE, 8, 8, com) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 8, 1) { @@ -1187,7 +1187,7 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc( com[7].s[i] = -in[1 * 8 + i].s[0] + (DATA_TYPE)5.25f * in[3 * 8 + i].s[0] - (DATA_TYPE)5.25f * in[5 * 8 + i].s[0] + in[7 * 8 + i].s[0]; } - TILE(DATA_TYPE, 8, 8, tmp) = { { 0 } }; + TILE(DATA_TYPE, 8, 8, tmp) = { { { 0 } } }; tmp[0].v = com[6].v; tmp[1].v = com[0].v + com[1].v; tmp[2].v = com[0].v - com[1].v; @@ -1215,7 +1215,7 @@ __kernel void winograd_input_transform_4x4_5x5_stepz1_nhwc( 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]; } - TILE(uint, 64, 1, dst_indirect_y) = { { 0 } }; + TILE(uint, 64, 1, dst_indirect_y) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 64, 1) { @@ -1286,10 +1286,10 @@ __kernel void winograd_input_transform_2x2_7x7_stepz1_nhwc( #if defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - TILE(DATA_TYPE, 8, 1, in) = { { 0 } }; - TILE(DATA_TYPE, 8, 1, out) = { { 0 } }; - TILE(int, 8, 1, src_indirect_y) = { { 0 } }; - TILE(int, 8, 1, src_indirect_mask) = { { 0 } }; + TILE(DATA_TYPE, 8, 1, in) = { { { 0 } } }; + TILE(DATA_TYPE, 8, 1, out) = { { { 0 } } }; + TILE(int, 8, 1, src_indirect_y) = { { { 0 } } }; + TILE(int, 8, 1, src_indirect_mask) = { { { 0 } } }; #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); @@ -1302,7 +1302,7 @@ __kernel void winograd_input_transform_2x2_7x7_stepz1_nhwc( in[i].v *= (DATA_TYPE) - 36.0f; } - TILE(DATA_TYPE, 1, 8, com) = { { 0 } }; + 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; @@ -1319,7 +1319,7 @@ __kernel void winograd_input_transform_2x2_7x7_stepz1_nhwc( 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; - TILE(uint, 8, 1, dst_indirect_y) = { { 0 } }; + TILE(uint, 8, 1, dst_indirect_y) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 8, 1) { @@ -1331,13 +1331,13 @@ __kernel void winograd_input_transform_2x2_7x7_stepz1_nhwc( #else // defined(WINOGRAD_INPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_INPUT_TRANSFORM_VERTICAL) - TILE(DATA_TYPE, 64, 1, in) = { { 0 } }; - TILE(DATA_TYPE, 64, 1, out) = { { 0 } }; + TILE(DATA_TYPE, 64, 1, in) = { { { 0 } } }; + TILE(DATA_TYPE, 64, 1, out) = { { { 0 } } }; // Load the tile from a NHWC tensor T_LOAD_NHWC(DATA_TYPE, 8, 8, 1, BUFFER, src, bout, y, x, cout, _ISRC_WIDTH, _ISRC_HEIGHT, src_stride_y, in); - TILE(DATA_TYPE, 8, 8, com) = { { 0 } }; + TILE(DATA_TYPE, 8, 8, com) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 8, 1) { @@ -1351,7 +1351,7 @@ __kernel void winograd_input_transform_2x2_7x7_stepz1_nhwc( com[7].s[i] = (DATA_TYPE)49.0f * in[3 * 8 + i].s[0] - (DATA_TYPE)36.0f * in[1 * 8 + i].s[0] + in[7 * 8 + i].s[0] - (DATA_TYPE)14.0f * in[5 * 8 + i].s[0]; } - TILE(DATA_TYPE, 8, 8, tmp) = { { 0 } }; + TILE(DATA_TYPE, 8, 8, tmp) = { { { 0 } } }; tmp[0].v = com[6].v; tmp[1].v = com[0].v - com[1].v; tmp[2].v = com[0].v + com[1].v; @@ -1379,7 +1379,7 @@ __kernel void winograd_input_transform_2x2_7x7_stepz1_nhwc( 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]; } - TILE(uint, 64, 1, dst_indirect_y) = { { 0 } }; + TILE(uint, 64, 1, dst_indirect_y) = { { { 0 } } }; LOOP_UNROLLING(int, i, 0, 64, 1) { -- cgit v1.2.1