From ef6ec50eff04adb3e5d60ba96a1a7d9118b10c5a Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Fri, 31 Jul 2020 11:38:36 +0100 Subject: COMPMID-3324: Fix oclgrind warnings Signed-off-by: Michalis Spyrou Change-Id: Ib14d158b9c5568981835312dcd9d5b9ca116649a Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/3637 Tested-by: Arm Jenkins Reviewed-by: Michele Di Giorgio Comments-Addressed: Arm Jenkins --- src/core/CL/cl_kernels/batchnormalization_layer.cl | 8 ++++---- src/core/CL/cl_kernels/winograd_output_transform.cl | 20 +++++++++++--------- 2 files changed, 15 insertions(+), 13 deletions(-) diff --git a/src/core/CL/cl_kernels/batchnormalization_layer.cl b/src/core/CL/cl_kernels/batchnormalization_layer.cl index 6f51db85e7..ad27aa386c 100644 --- a/src/core/CL/cl_kernels/batchnormalization_layer.cl +++ b/src/core/CL/cl_kernels/batchnormalization_layer.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 Arm Limited. + * Copyright (c) 2017-2020 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -346,9 +346,9 @@ __kernel void fuse_batchnormalization_layer(TENSOR3D_DECLARATION(w), int c0 = z % DIM2; int c1 = z / DIM2; #else // ! defined(DIM2) - int c0 = 0; + int c0 = 0; #if defined(NHWC) - int c1 = x; + int c1 = x; #else // defined(NHWC) int c1 = z; #endif // defined(NHWC) @@ -386,7 +386,7 @@ __kernel void fuse_batchnormalization_layer(TENSOR3D_DECLARATION(w), // Compute bias #if !defined(DIM2) && defined(NHWC) if(z == 0 && y == 0) -#else !defined(DIM2) && defined(NHWC) +#else // !defined(DIM2) && defined(NHWC) if(x == 0 && y == 0 && c0 == 0) #endif // !defined(DIM2) && defined(NHWC) { diff --git a/src/core/CL/cl_kernels/winograd_output_transform.cl b/src/core/CL/cl_kernels/winograd_output_transform.cl index 99888edd20..efd8502657 100644 --- a/src/core/CL/cl_kernels/winograd_output_transform.cl +++ b/src/core/CL/cl_kernels/winograd_output_transform.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 Arm Limited. + * Copyright (c) 2018-2020 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -157,8 +157,8 @@ __kernel void winograd_output_transform_2x2_3x3_nchw( // Store the output tile #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) - const const VEC_DATA_TYPE(DATA_TYPE, 2) - out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 2))(out00, out01), VEC_DATA_TYPE(DATA_TYPE, 2)), A_VAL, B_VAL); + const VEC_DATA_TYPE(DATA_TYPE, 2) + out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 2))(out00, out01), VEC_DATA_TYPE(DATA_TYPE, 2)), A_VAL, B_VAL); *((__global DATA_TYPE *)(dst_addr + 0 * dst_stride_y)) = out0_dt.s0; *((__global DATA_TYPE *)(dst_addr + 1 * dst_stride_y)) = out0_dt.s1; #else // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) @@ -288,7 +288,7 @@ __kernel void winograd_output_transform_2x2_7x7_nhwc( // Get output address int offset = dst_offset_first_element_in_bytes + x_out * sizeof(DATA_TYPE) + y_out * dst_stride_y + z_out * dst_stride_z; VEC_DATA_TYPE(DATA_TYPE, 2) - out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 2))(out00, out01), VEC_DATA_TYPE(DATA_TYPE, 2)), A_VAL, B_VAL); + out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 2))(out00, out01), VEC_DATA_TYPE(DATA_TYPE, 2)), A_VAL, B_VAL); *(__global DATA_TYPE *)(dst_ptr + 0 * dst_stride_y + offset) = out0_dt.s0; *(__global DATA_TYPE *)(dst_ptr + 1 * dst_stride_y + offset) = out0_dt.s1; #endif // defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) @@ -599,7 +599,8 @@ __kernel void winograd_output_transform_4x4_3x3_nchw( // Store the output tile #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) VEC_DATA_TYPE(DATA_TYPE, 4) - out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), VEC_DATA_TYPE(DATA_TYPE, 4)), A_VAL, B_VAL); + out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), VEC_DATA_TYPE(DATA_TYPE, 4)), A_VAL, + B_VAL); *((__global DATA_TYPE *)(dst_addr + 0 * dst_stride_y)) = out0_dt.s0; *((__global DATA_TYPE *)(dst_addr + 1 * dst_stride_y)) = out0_dt.s1; *((__global DATA_TYPE *)(dst_addr + 2 * dst_stride_y)) = out0_dt.s2; @@ -839,7 +840,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( // Store the 1x4 output tile VEC_DATA_TYPE(DATA_TYPE, 4) - out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), VEC_DATA_TYPE(DATA_TYPE, 4)), A_VAL, B_VAL); + out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), VEC_DATA_TYPE(DATA_TYPE, 4)), A_VAL, B_VAL); *((__global DATA_TYPE *)(dst_ptr + offset.s0)) = out0_dt.s0; *((__global DATA_TYPE *)(dst_ptr + offset.s1)) = out0_dt.s1; *((__global DATA_TYPE *)(dst_ptr + offset.s2)) = out0_dt.s2; @@ -875,7 +876,7 @@ __kernel void winograd_output_transform_4x4_3x3_nhwc( out2_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 4))(out20, out21, out22, out23), VEC_DATA_TYPE(DATA_TYPE, 4)), A_VAL, B_VAL); VEC_DATA_TYPE(DATA_TYPE, 4) out3_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 4))(out30, out31, out32, out33), - VEC_DATA_TYPE(DATA_TYPE, 4)), + VEC_DATA_TYPE(DATA_TYPE, 4)), A_VAL, B_VAL); *((__global DATA_TYPE *)(dst_ptr + mult_y.s0 * 0 * dst_stride_y + offset.s0)) = out0_dt.s0; *((__global DATA_TYPE *)(dst_ptr + mult_y.s0 * 1 * dst_stride_y + offset.s0)) = out0_dt.s1; @@ -1011,7 +1012,8 @@ __kernel void winograd_output_transform_4x4_5x5_nchw( // Store the output tile #if defined(WINOGRAD_OUTPUT_TRANSFORM_VERTICAL) VEC_DATA_TYPE(DATA_TYPE, 4) - out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), VEC_DATA_TYPE(DATA_TYPE, 4)), A_VAL, B_VAL); + out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), VEC_DATA_TYPE(DATA_TYPE, 4)), A_VAL, + B_VAL); *((__global DATA_TYPE *)(dst_addr + 0 * dst_stride_y)) = out0_dt.s0; *((__global DATA_TYPE *)(dst_addr + 1 * dst_stride_y)) = out0_dt.s1; *((__global DATA_TYPE *)(dst_addr + 2 * dst_stride_y)) = out0_dt.s2; @@ -1239,7 +1241,7 @@ __kernel void winograd_output_transform_4x4_5x5_nhwc( offset = min(offset + (int4)(0, 1, 2, 3) * (int4)dst_stride_z, (int4)dst_size); // If address is beyond the last plane, clamp it to dst_size (which points to the last padding). VEC_DATA_TYPE(DATA_TYPE, 4) - out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), VEC_DATA_TYPE(DATA_TYPE, 4)), A_VAL, B_VAL); + out0_dt = ACTIVATION(ACTIVATION_TYPE, DATA_TYPE, CONVERT((VEC_DATA_TYPE(float, 4))(out00, out01, out02, out03), VEC_DATA_TYPE(DATA_TYPE, 4)), A_VAL, B_VAL); *(__global DATA_TYPE *)(dst_ptr + offset.s0) = out0_dt.s0; *(__global DATA_TYPE *)(dst_ptr + offset.s1) = out0_dt.s1; *(__global DATA_TYPE *)(dst_ptr + offset.s2) = out0_dt.s2; -- cgit v1.2.1