From 0333704d15f68fe18551b00d00839a9efbbe858f Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Wed, 26 Feb 2020 17:47:55 +0000 Subject: COMPMID-3223: Fix -14 error code on Firefly for CLReduction Wrong data type was used for MIN/MAX reduction causing segfaults. This patch also simplifies window calculation for non-parallel reduction. Change-Id: I0abd9e84540801f92b306accd1439ff5df126a9e Signed-off-by: Michele Di Giorgio Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2792 Reviewed-by: Georgios Pinitas Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/reduction_operation.cl | 12 ++++-------- 1 file changed, 4 insertions(+), 8 deletions(-) (limited to 'src/core/CL/cl_kernels') diff --git a/src/core/CL/cl_kernels/reduction_operation.cl b/src/core/CL/cl_kernels/reduction_operation.cl index a5fd0b3622..b9d33bdd98 100644 --- a/src/core/CL/cl_kernels/reduction_operation.cl +++ b/src/core/CL/cl_kernels/reduction_operation.cl @@ -196,7 +196,7 @@ __kernel void reduction_operation_non_parallel_x( #if defined(MIN) res = select(res, in, CONVERT(ISLESS(in, res), COND_DATA_TYPE)); #elif defined(MAX) - res = select(res, in, CONVERT(ISGREATER(in, res), COND_DATA_TYPE)); + res = select(res, in, CONVERT(ISGREATER(in, res), COND_DATA_TYPE)); #elif defined(PROD) #if defined(OFFSET) && defined(SCALE) res_f *= DEQUANTIZE(in, OFFSET, SCALE, DATA_TYPE_PROMOTED, 1); @@ -223,11 +223,7 @@ __kernel void reduction_operation_non_parallel_x( res = QUANTIZE(res_f, OFFSET, SCALE, DATA_TYPE_PROMOTED, 1); #endif // defined(PROD) && defined(OFFSET) && defined(SCALE) -#if defined(MIN) || defined(MAX) - *((__global DATA_TYPE_PROMOTED *)output.ptr) = res; -#else // !(defined(MIN) || defined(MAX)) *((__global DATA_TYPE *)output.ptr) = CONVERT_SAT(res, DATA_TYPE); -#endif // defined(MIN) || defined(MAX) } #endif // defined(WIDTH) @@ -276,7 +272,7 @@ __kernel void reduction_operation_y( #if defined(MIN) res = select(res, in, ISLESS(in, res)); #elif defined(MAX) - res = select(res, in, ISGREATER(in, res)); + res = select(res, in, ISGREATER(in, res)); #else // !(defined(MAX) || defined(MIN)) #if defined(SUM_SQUARE) in *= in; @@ -373,7 +369,7 @@ __kernel void reduction_operation_z( #if defined(MIN) res = select(res, in, ISLESS(in, res)); #elif defined(MAX) - res = select(res, in, ISGREATER(in, res)); + res = select(res, in, ISGREATER(in, res)); #else // !(defined(MAX) || defined(MIN)) #if defined(SUM_SQUARE) in *= in; @@ -472,7 +468,7 @@ __kernel void reduction_operation_w( #if defined(MIN) res = select(res, in, ISLESS(in, res)); #elif defined(MAX) - res = select(res, in, ISGREATER(in, res)); + res = select(res, in, ISGREATER(in, res)); #else // !(defined(MAX) || defined(MIN)) #if defined(SUM_SQUARE) in *= in; -- cgit v1.2.1