From cbe39055f61bc7a3f6aef4588de47d48b77cf354 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Tue, 2 Jul 2019 11:39:10 +0100 Subject: COMPMID-2433: (Nightly) Bug in CLDequantizationLayer The problem was a typo in the OpenCL kernel. Also, removed VEC_SIZE from internal ifdef because it is already checked outside. Change-Id: I24721996f7b5ccbc3b98a80045c45aba765f7522 Signed-off-by: Michele Di Giorgio Reviewed-on: https://review.mlplatform.org/c/1454 Reviewed-by: Michalis Spyrou Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/dequantization_layer.cl | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/src/core/CL/cl_kernels/dequantization_layer.cl b/src/core/CL/cl_kernels/dequantization_layer.cl index ad3ed35480..7d87dc6a2d 100644 --- a/src/core/CL/cl_kernels/dequantization_layer.cl +++ b/src/core/CL/cl_kernels/dequantization_layer.cl @@ -58,7 +58,7 @@ __kernel void dequantization_layer( Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input); Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output); -#if defined(VEC_SIZE) && defined(LAST_ACCESSED_X) +#if defined(LAST_ACCESSED_X) // Check if access on width gets out of bounds // If it does shift access vector to access elements within bounds const int xi = (int)(get_global_id(0) * VEC_SIZE); @@ -83,9 +83,9 @@ __kernel void dequantization_layer( // Store result VSTORE(VEC_SIZE) (CONVERT(res, VEC_DATA_TYPE(DATA_TYPE_DST, VEC_SIZE)), 0, (__global DATA_TYPE_DST *)output.ptr); -#else // !defined(VEC_SIZE) || !defined(LAST_ACCESSED_X) - *((__global DATA_TYPE_DST *)(output.ptr)) = (DATA_TYPE)((float)((int)(*((__global DATA_TYPE_SRC *)(input.ptr))) - (int)(OFFSET)) * (float)(SCALE)); -#endif // defined(VEC_SIZE) && defined(LAST_ACCESSED_X) +#else // !defined(LAST_ACCESSED_X) + *((__global DATA_TYPE_DST *)(output.ptr)) = (DATA_TYPE_DST)((float)((int)(*((__global DATA_TYPE_SRC *)(input.ptr))) - (int)(OFFSET)) * (float)(SCALE)); +#endif // defined(LAST_ACCESSED_X) } #endif // defined(VEC_SIZE) && defined(DATA_TYPE_SRC) && defined(DATA_TYPE_DST) && defined(SCALE) && defined(OFFSET) -- cgit v1.2.1