diff options
| author | Sheri Zhang <sheri.zhang@arm.com> | 2021-04-08 18:37:24 +0100 |
|---|---|---|
| committer | Sheri Zhang <sheri.zhang@arm.com> | 2021-04-09 10:06:14 +0000 |
| commit | ec82b95f26198dc538f94aa90f0febfaf0eb2751 (patch) | |
| tree | bacfe3eeb3bab067f61bc314b29000b00ceeded5 | |
| parent | d340eb2015fbfca00769076a7169d5398f78bac9 (diff) | |
| download | ComputeLibrary-ec82b95f26198dc538f94aa90f0febfaf0eb2751.tar.gz | |
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 <sheri.zhang@arm.com>
Change-Id: I2d8fc6400f32af5406fbf2d2556127a53b2ce918
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/5392
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/direct_convolution.cl | 10 | ||||
| -rw-r--r-- | 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) |