diff options
| author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2019-07-02 11:39:10 +0100 |
|---|---|---|
| committer | Michele Di Giorgio <michele.digiorgio@arm.com> | 2019-07-02 13:00:26 +0000 |
| commit | cbe39055f61bc7a3f6aef4588de47d48b77cf354 (patch) | |
| tree | d194a35eda0a8d43024dc4ce93e29a742d66640b | |
| parent | a7484365725133b22408aeccf16824025fbba16a (diff) | |
| download | ComputeLibrary-cbe39055f61bc7a3f6aef4588de47d48b77cf354.tar.gz | |
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 <michele.digiorgio@arm.com>
Reviewed-on: https://review.mlplatform.org/c/1454
Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/dequantization_layer.cl | 8 |
1 files 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) |