From ec82b95f26198dc538f94aa90f0febfaf0eb2751 Mon Sep 17 00:00:00 2001 From: Sheri Zhang Date: Thu, 8 Apr 2021 18:37:24 +0100 Subject: Fix OpenCL kernel compiling failure with array initilizer 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. Resolves: COMPMID-4348 Signed-off-by: Sheri Zhang Change-Id: I2d8fc6400f32af5406fbf2d2556127a53b2ce918 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5392 Reviewed-by: Georgios Pinitas Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/direct_convolution.cl | 10 +++++----- src/core/CL/cl_kernels/winograd_output_transform.cl | 18 +++++++++--------- 2 files changed, 14 insertions(+), 14 deletions(-) diff --git a/src/core/CL/cl_kernels/direct_convolution.cl b/src/core/CL/cl_kernels/direct_convolution.cl index dde024faa4..96196bda8d 100644 --- a/src/core/CL/cl_kernels/direct_convolution.cl +++ b/src/core/CL/cl_kernels/direct_convolution.cl @@ -136,8 +136,8 @@ __kernel void direct_convolution_nhwc( // .v = access the whole vector (OpenCL vector) // .s[x] = access the vector element at position x (scalar access) - TILE(int, M0, 1, xi) = { { 0 } }; - TILE(int, M0, 1, yi) = { { 0 } }; + TILE(int, M0, 1, xi) = {{ { 0 } }}; + TILE(int, M0, 1, yi) = {{ { 0 } }}; // Convert the linear index to coordinate LOOP_UNROLLING(int, i, 0, M0, 1) @@ -151,7 +151,7 @@ __kernel void direct_convolution_nhwc( uint wei_x = 0; // Initialize the accumulators - TILE(ACC_DATA_TYPE, M0, N0, c) = { { 0 } }; + TILE(ACC_DATA_TYPE, M0, N0, c) = {{ { 0 } }}; for(int i = 0; i < (_IWEI_WIDTH * _IWEI_HEIGHT); ++i) { @@ -159,8 +159,8 @@ __kernel void direct_convolution_nhwc( int xk = i % _IWEI_WIDTH; int yk = i / _IWEI_WIDTH; - TILE(int, M0, 1, src_indirect_y) = { { 0 } }; - TILE(int, M0, 1, src_indirect_mask) = { { 0 } }; + TILE(int, M0, 1, src_indirect_y) = {{ { 0 } }}; + TILE(int, M0, 1, src_indirect_mask) = {{ { 0 } }}; // Calculate the source indirect Y and the source indirect mask // Since the indirect Y is clamped when out-of-bound, the mask is used to diff --git a/src/core/CL/cl_kernels/winograd_output_transform.cl b/src/core/CL/cl_kernels/winograd_output_transform.cl index a56ce3de6f..59402bf9b3 100644 --- a/src/core/CL/cl_kernels/winograd_output_transform.cl +++ b/src/core/CL/cl_kernels/winograd_output_transform.cl @@ -690,9 +690,9 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( #if defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) - TILE(DATA_TYPE, 6, N0, in) = { { 0 } }; - TILE(DATA_TYPE, 4, N0, out) = { { 0 } }; - TILE(uint, 6, 1, src_indirect_y) = { { 0 } }; + TILE(DATA_TYPE, 6, N0, in) = {{ { 0 } }}; + TILE(DATA_TYPE, 4, N0, out) = {{ { 0 } }}; + TILE(uint, 6, 1, src_indirect_y) = {{ { 0 } }}; LOOP_UNROLLING(int, i, 0, 6, 1) { @@ -723,7 +723,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( T_ACTIVATION(DATA_TYPE, 4, N0, ACTIVATION_TYPE, A_VAL, B_VAL, out, out); - TILE(uint, 4, 1, dst_indirect_y) = { { 0 } }; + TILE(uint, 4, 1, dst_indirect_y) = {{ { 0 } }}; // Calculate the destination indirect Y #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) @@ -748,9 +748,9 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( #else // defined(WINOGRAD_OUTPUT_TRANSFORM_HORIZONTAL) || defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) // Calculate the indirect Y for the source tensor - TILE(DATA_TYPE, 36, N0, in) = { { 0 } }; - TILE(DATA_TYPE, 4, N0, tmp) = { { 0 } }; - TILE(uint, 36, 1, src_indirect_y) = { { 0 } }; + TILE(DATA_TYPE, 36, N0, in) = {{ { 0 } }}; + TILE(DATA_TYPE, 4, N0, tmp) = {{ { 0 } }}; + TILE(uint, 36, 1, src_indirect_y) = {{ { 0 } }}; LOOP_UNROLLING(int, i, 0, 36, 1) { @@ -775,7 +775,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( } // Compute the output tile - TILE(DATA_TYPE, 16, N0, out) = { { 0 } }; + TILE(DATA_TYPE, 16, N0, out) = {{ { 0 } }}; LOOP_UNROLLING(int, i, 0, 4, 1) { @@ -804,7 +804,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( T_ACTIVATION(DATA_TYPE, 16, N0, ACTIVATION_TYPE, A_VAL, B_VAL, out, out); - TILE(uint, 16, 1, dst_indirect_y) = { { 0 } }; + TILE(uint, 16, 1, dst_indirect_y) = {{ { 0 } }}; // Calculate the destination indirect Y LOOP_UNROLLING(int, yk, 0, 4, 1) -- cgit v1.2.1