From 0333704d15f68fe18551b00d00839a9efbbe858f Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Wed, 26 Feb 2020 17:47:55 +0000 Subject: 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 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2792 Reviewed-by: Georgios Pinitas Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins --- src/core/CL/cl_kernels/reduction_operation.cl | 12 ++++-------- src/core/CL/kernels/CLReductionOperationKernel.cpp | 18 +++++++----------- 2 files changed, 11 insertions(+), 19 deletions(-) (limited to 'src') 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; diff --git a/src/core/CL/kernels/CLReductionOperationKernel.cpp b/src/core/CL/kernels/CLReductionOperationKernel.cpp index ff824fce33..a0f9a49fff 100644 --- a/src/core/CL/kernels/CLReductionOperationKernel.cpp +++ b/src/core/CL/kernels/CLReductionOperationKernel.cpp @@ -87,14 +87,7 @@ std::tuple validate_and_configure_window(ITensorInfo *input, ITe { case 0: { - if(is_serial_op) - { - AccessWindowHorizontal input_access(input, 0, input->dimension(0)); - AccessWindowHorizontal output_access(output, 0, 1); - window_changed = update_window_and_padding(win, input_access, output_access); - output_access.set_valid_region(win, ValidRegion(Coordinates(), output->tensor_shape())); - } - else + if(!is_serial_op) { const unsigned int border_width = ((input->dimension(0) % border_val) != 0) ? border_val - input->dimension(0) % border_val : 0; AccessWindowStatic input_access(input, 0, 0, input->dimension(0) + border_width, 1); @@ -269,8 +262,11 @@ void CLReductionOperationKernel::run(const Window &window, cl::CommandQueue &que Window window_in{ window }; window_in.set(Window::DimX, Window::Dimension(0, _input->info()->dimension(0), _input->info()->dimension(0))); - Window in_slice = window.first_slice_window_1D(); - Window out_slice = window.first_slice_window_1D(); + Window out_window{ window }; + out_window.set(Window::DimX, Window::Dimension(0, 0, 0)); + + Window in_slice = window_in.first_slice_window_1D(); + Window out_slice = out_window.first_slice_window_1D(); do { @@ -279,7 +275,7 @@ void CLReductionOperationKernel::run(const Window &window, cl::CommandQueue &que add_1D_tensor_argument(idx, _output, out_slice); enqueue(queue, *this, in_slice, lws_hint()); } - while(window_in.slide_window_slice_1D(in_slice) && window.slide_window_slice_1D(out_slice)); + while(window_in.slide_window_slice_1D(in_slice) && out_window.slide_window_slice_1D(out_slice)); } else { -- cgit v1.2.1