From b28905010a95044c7a1c0a5665fc886521a56541 Mon Sep 17 00:00:00 2001 From: Usama Arif Date: Tue, 21 May 2019 11:48:37 +0100 Subject: COMPMID-2281: Implement REDUCE_MIN operator for CL Change-Id: I60fd3affad0ab5a2f1cef18aaa46dc2dc448caeb Signed-off-by: Usama Arif Reviewed-on: https://review.mlplatform.org/c/1194 Comments-Addressed: Arm Jenkins Reviewed-by: Pablo Marquez Reviewed-by: Michalis Spyrou Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/reduction_operation.cl | 12 ++++++++++++ 1 file changed, 12 insertions(+) (limited to 'src/core/CL/cl_kernels/reduction_operation.cl') diff --git a/src/core/CL/cl_kernels/reduction_operation.cl b/src/core/CL/cl_kernels/reduction_operation.cl index 749e3cdaa3..86cf37e491 100644 --- a/src/core/CL/cl_kernels/reduction_operation.cl +++ b/src/core/CL/cl_kernels/reduction_operation.cl @@ -197,6 +197,8 @@ __kernel void reduction_operation_non_parallel_x( #elif defined(ARG_MIN) indx = select(indx, x, ISLESS(in, res)); res = select(res, in, CONVERT(ISLESS(in, res), COND_DATA_TYPE)); +#elif defined(MIN) + res = select(res, in, CONVERT(ISLESS(in, res), COND_DATA_TYPE)); #else // !(defined(ARG_MAX) || defined(ARG_MIN)) res += in; #endif // defined(ARG_MAX) || defined(ARG_MIN) @@ -209,7 +211,11 @@ __kernel void reduction_operation_non_parallel_x( #if defined(MEAN) res /= WIDTH; #endif // defined(MEAN) +#if defined(MIN) + *((__global DATA_TYPE_PROMOTED *)output.ptr) = res; +#else // defined(MIN) *((__global uchar *)output.ptr) = convert_uchar(res); +#endif // defined(MIN) #endif // defined(ARG_MAX) || defined(ARG_MIN) } #endif // defined(WIDTH) @@ -263,6 +269,8 @@ __kernel void reduction_operation_y( uint16 cond_conv = CONVERT(ISLESS(in, res), uint16); indx = select(indx, y, cond_conv); res = select(res, in, ISLESS(in, res)); +#elif defined(MIN) + res = select(res, in, ISLESS(in, res)); #else // !(defined(ARG_MAX) || defined(ARG_MIN)) #if defined(SUM_SQUARE) in *= in; @@ -350,6 +358,8 @@ __kernel void reduction_operation_z( uint16 cond_conv = CONVERT(ISLESS(in, res), uint16); indx = select(indx, z, cond_conv); res = select(res, in, ISLESS(in, res)); +#elif defined(MIN) + res = select(res, in, ISLESS(in, res)); #else // !(defined(ARG_MAX) || defined(ARG_MIN)) #if defined(SUM_SQUARE) in *= in; @@ -439,6 +449,8 @@ __kernel void reduction_operation_w( uint16 cond_conv = CONVERT(ISLESS(in, res), uint16); indx = select(indx, w, cond_conv); res = select(res, in, ISLESS(in, res)); +#elif defined(MIN) + res = select(res, in, ISLESS(in, res)); #else // !(defined(ARG_MAX) || defined(ARG_MIN)) #if defined(SUM_SQUARE) in *= in; -- cgit v1.2.1