diff options
| author | Usama Arif <usama.arif@arm.com> | 2019-05-22 16:32:27 +0100 |
|---|---|---|
| committer | Usama Arif <usama.arif@arm.com> | 2019-05-24 09:20:23 +0000 |
| commit | 048b0f300ee729cac1b3019311589c654771fb8f (patch) | |
| tree | cbf5f8cb2baee37292bc4dfe5ab218e40a464b54 /src/core/CL/cl_kernels | |
| parent | 9e631c204444e7b095510c54819e944f9be8d342 (diff) | |
| download | ComputeLibrary-048b0f300ee729cac1b3019311589c654771fb8f.tar.gz | |
COMPMID-2278: Implement REDUCE_MAX operator for CL.
Change-Id: Ie23e3ddc45d6f5506a63f935758a215ba7412bf5
Signed-off-by: Usama Arif <usama.arif@arm.com>
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/186031
Tested-by: bsgcomp <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Comments-Addressed: bsgcomp <bsgcomp@arm.com>
Reviewed-on: https://review.mlplatform.org/c/1214
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels')
| -rw-r--r-- | src/core/CL/cl_kernels/reduction_operation.cl | 28 |
1 files changed, 18 insertions, 10 deletions
diff --git a/src/core/CL/cl_kernels/reduction_operation.cl b/src/core/CL/cl_kernels/reduction_operation.cl index 86cf37e491..5a4bb9ff4c 100644 --- a/src/core/CL/cl_kernels/reduction_operation.cl +++ b/src/core/CL/cl_kernels/reduction_operation.cl @@ -198,7 +198,9 @@ __kernel void reduction_operation_non_parallel_x( 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)); + 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)); #else // !(defined(ARG_MAX) || defined(ARG_MIN)) res += in; #endif // defined(ARG_MAX) || defined(ARG_MIN) @@ -211,11 +213,11 @@ __kernel void reduction_operation_non_parallel_x( #if defined(MEAN) res /= WIDTH; #endif // defined(MEAN) -#if defined(MIN) +#if defined(MIN) || defined(MAX) *((__global DATA_TYPE_PROMOTED *)output.ptr) = res; -#else // defined(MIN) +#else // defined(MIN) || defined(MAX) *((__global uchar *)output.ptr) = convert_uchar(res); -#endif // defined(MIN) +#endif // defined(MIN) || defined(MAX) #endif // defined(ARG_MAX) || defined(ARG_MIN) } #endif // defined(WIDTH) @@ -266,11 +268,13 @@ __kernel void reduction_operation_y( indx = select(indx, y, cond_conv); res = select(res, in, ISGREATER(in, res)); #elif defined(ARG_MIN) - uint16 cond_conv = CONVERT(ISLESS(in, res), uint16); - indx = select(indx, y, cond_conv); - res = select(res, in, ISLESS(in, res)); + 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)); + res = select(res, in, ISLESS(in, res)); +#elif defined(MAX) + res = select(res, in, ISGREATER(in, res)); #else // !(defined(ARG_MAX) || defined(ARG_MIN)) #if defined(SUM_SQUARE) in *= in; @@ -359,7 +363,9 @@ __kernel void reduction_operation_z( indx = select(indx, z, cond_conv); res = select(res, in, ISLESS(in, res)); #elif defined(MIN) - res = select(res, in, ISLESS(in, res)); + res = select(res, in, ISLESS(in, res)); +#elif defined(MAX) + res = select(res, in, ISGREATER(in, res)); #else // !(defined(ARG_MAX) || defined(ARG_MIN)) #if defined(SUM_SQUARE) in *= in; @@ -450,7 +456,9 @@ __kernel void reduction_operation_w( indx = select(indx, w, cond_conv); res = select(res, in, ISLESS(in, res)); #elif defined(MIN) - res = select(res, in, ISLESS(in, res)); + res = select(res, in, ISLESS(in, res)); +#elif defined(MAX) + res = select(res, in, ISGREATER(in, res)); #else // !(defined(ARG_MAX) || defined(ARG_MIN)) #if defined(SUM_SQUARE) in *= in; |