aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--src/core/CL/cl_kernels/reduction_operation.cl28
-rw-r--r--src/core/CL/kernels/CLReductionOperationKernel.cpp11
-rw-r--r--src/runtime/CL/functions/CLReductionOperation.cpp32
-rw-r--r--tests/validation/CL/ReductionOperation.cpp1
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<Status, Window> 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<float>::max());
+ break;
+ }
+ case DataType::F16:
+ {
+ pixelValue = PixelValue(static_cast<half>(-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,
});
generated by cgit v1.2.1 (git 2.23.0) at 2024-05-29 02:16:07 +0000