diff options
author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2020-02-26 17:47:55 +0000 |
---|---|---|
committer | Michele Di Giorgio <michele.digiorgio@arm.com> | 2020-02-27 07:53:37 +0000 |
commit | 0333704d15f68fe18551b00d00839a9efbbe858f (patch) | |
tree | 2bf7c0e684d089ff86af675057701fc1faf3aaba | |
parent | 0b192e84510c006d87cee3c29f95511ad088b704 (diff) | |
download | ComputeLibrary-0333704d15f68fe18551b00d00839a9efbbe858f.tar.gz |
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 <michele.digiorgio@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2792
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
-rw-r--r-- | src/core/CL/cl_kernels/reduction_operation.cl | 12 | ||||
-rw-r--r-- | src/core/CL/kernels/CLReductionOperationKernel.cpp | 18 |
2 files changed, 11 insertions, 19 deletions
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<Status, Window> 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 { |