diff options
| author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2020-02-26 17:47:55 +0000 |
|---|---|---|
| committer | Michele Di Giorgio <michele.digiorgio@arm.com> | 2020-02-27 07:53:37 +0000 |
| commit | 0333704d15f68fe18551b00d00839a9efbbe858f (patch) | |
| tree | 2bf7c0e684d089ff86af675057701fc1faf3aaba /src/core/CL/cl_kernels | |
| parent | 0b192e84510c006d87cee3c29f95511ad088b704 (diff) | |
| download | ComputeLibrary-0333704d15f68fe18551b00d00839a9efbbe858f.tar.gz | |
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 <michele.digiorgio@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2792
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels')
| -rw-r--r-- | src/core/CL/cl_kernels/reduction_operation.cl | 12 |
1 files changed, 4 insertions, 8 deletions
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; |