From 048b0f300ee729cac1b3019311589c654771fb8f Mon Sep 17 00:00:00 2001 From: Usama Arif Date: Wed, 22 May 2019 16:32:27 +0100 Subject: COMPMID-2278: Implement REDUCE_MAX operator for CL. Change-Id: Ie23e3ddc45d6f5506a63f935758a215ba7412bf5 Signed-off-by: Usama Arif Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/186031 Tested-by: bsgcomp Reviewed-by: Georgios Pinitas Comments-Addressed: bsgcomp Reviewed-on: https://review.mlplatform.org/c/1214 Tested-by: Arm Jenkins Comments-Addressed: Arm Jenkins --- src/core/CL/cl_kernels/reduction_operation.cl | 28 ++++++++++++------- src/core/CL/kernels/CLReductionOperationKernel.cpp | 11 ++++++-- src/runtime/CL/functions/CLReductionOperation.cpp | 32 ++++++++++++++++++++++ tests/validation/CL/ReductionOperation.cpp | 1 + 4 files changed, 59 insertions(+), 13 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; diff --git a/src/core/CL/kernels/CLReductionOperationKernel.cpp b/src/core/CL/kernels/CLReductionOperationKernel.cpp index eb76349a02..9db8ae6cde 100644 --- a/src/core/CL/kernels/CLReductionOperationKernel.cpp +++ b/src/core/CL/kernels/CLReductionOperationKernel.cpp @@ -89,7 +89,8 @@ std::tuple validate_and_configure_window(ITensorInfo *input, ITe const unsigned int num_elems_processed_per_iteration = (is_data_type_quantized(input->data_type()) && (axis == 0)) ? 1 : 16; Window win = calculate_max_window(*input, Steps(num_elems_processed_per_iteration)); bool window_changed = false; - const bool is_serial_op = (op == ReductionOperation::ARG_IDX_MAX || op == ReductionOperation::ARG_IDX_MIN || op == ReductionOperation::MIN || is_data_type_quantized(input->data_type())); + const bool is_serial_op = (op == ReductionOperation::ARG_IDX_MAX || op == ReductionOperation::ARG_IDX_MIN || op == ReductionOperation::MIN + || op == ReductionOperation::MAX || is_data_type_quantized(input->data_type())); switch(axis) { @@ -170,6 +171,7 @@ void CLReductionOperationKernel::configure(const ICLTensor *input, ICLTensor *ou build_opts.add_option_if(op == ReductionOperation::ARG_IDX_MIN, "-DARG_MIN"); build_opts.add_option_if(op == ReductionOperation::PROD, "-DPROD"); build_opts.add_option_if(op == ReductionOperation::MIN, "-DMIN"); + build_opts.add_option_if(op == ReductionOperation::MAX, "-DMAX"); build_opts.add_option_if(input->info()->num_channels() == 2, "-DCOMPLEX"); switch(op) @@ -184,6 +186,7 @@ void CLReductionOperationKernel::configure(const ICLTensor *input, ICLTensor *ou case ReductionOperation::ARG_IDX_MAX: case ReductionOperation::ARG_IDX_MIN: case ReductionOperation::MIN: + case ReductionOperation::MAX: break; case ReductionOperation::PROD: build_opts.add_option(("-DOPERATION=product")); @@ -195,7 +198,8 @@ void CLReductionOperationKernel::configure(const ICLTensor *input, ICLTensor *ou // Create kernel cl::NDRange lws_hint = CLKernelLibrary::get().default_ndrange(); std::string kernel_axis_name; - const bool is_serial_op = (op == ReductionOperation::ARG_IDX_MAX || op == ReductionOperation::ARG_IDX_MIN || op == ReductionOperation::MIN || is_data_type_quantized(input->info()->data_type())); + const bool is_serial_op = (op == ReductionOperation::ARG_IDX_MAX || op == ReductionOperation::ARG_IDX_MIN || op == ReductionOperation::MIN || op == ReductionOperation::MAX + || is_data_type_quantized(input->info()->data_type())); switch(axis) { case 0: @@ -260,7 +264,8 @@ void CLReductionOperationKernel::run(const Window &window, cl::CommandQueue &que ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window); - const bool is_serial_op = (_op == ReductionOperation::ARG_IDX_MAX || _op == ReductionOperation::ARG_IDX_MIN || _op == ReductionOperation::MIN || is_data_type_quantized(_input->info()->data_type())); + const bool is_serial_op = (_op == ReductionOperation::ARG_IDX_MAX || _op == ReductionOperation::ARG_IDX_MIN || _op == ReductionOperation::MIN || _op == ReductionOperation::MAX + || is_data_type_quantized(_input->info()->data_type())); switch(_reduction_axis) { case 0: diff --git a/src/runtime/CL/functions/CLReductionOperation.cpp b/src/runtime/CL/functions/CLReductionOperation.cpp index 2e48897323..38f0a7523c 100644 --- a/src/runtime/CL/functions/CLReductionOperation.cpp +++ b/src/runtime/CL/functions/CLReductionOperation.cpp @@ -110,6 +110,11 @@ Status CLReductionOperation::validate(const ITensorInfo *input, const ITensorInf intermediate_kernel_op = ReductionOperation::MIN; last_kernel_op = ReductionOperation::MIN; break; + case ReductionOperation::MAX: + first_kernel_op = ReductionOperation::MAX; + intermediate_kernel_op = ReductionOperation::MAX; + last_kernel_op = ReductionOperation::MAX; + break; default: ARM_COMPUTE_ERROR("Not supported"); } @@ -211,6 +216,33 @@ void CLReductionOperation::configure(ICLTensor *input, ICLTensor *output, unsign } } break; + case ReductionOperation::MAX: + first_kernel_op = ReductionOperation::MAX; + intermediate_kernel_op = ReductionOperation::MAX; + last_kernel_op = ReductionOperation::MAX; + switch(input->info()->data_type()) + { + case DataType::F32: + { + pixelValue = PixelValue(-std::numeric_limits::max()); + break; + } + case DataType::F16: + { + pixelValue = PixelValue(static_cast(-65504.0f)); + break; + } + case DataType::QASYMM8: + { + pixelValue = PixelValue(0, input->info()->data_type(), input->info()->quantization_info()); + break; + } + default: + { + ARM_COMPUTE_ERROR("Unsupported DataType"); + } + } + break; default: ARM_COMPUTE_ERROR("Not supported"); } diff --git a/tests/validation/CL/ReductionOperation.cpp b/tests/validation/CL/ReductionOperation.cpp index 8fc2911a2b..9a3cd996fa 100644 --- a/tests/validation/CL/ReductionOperation.cpp +++ b/tests/validation/CL/ReductionOperation.cpp @@ -53,6 +53,7 @@ const auto ReductionOperations = framework::dataset::make("ReductionOperation", ReductionOperation::SUM, ReductionOperation::PROD, ReductionOperation::MIN, + ReductionOperation::MAX, }); -- cgit v1.2.1