From f9bae2e9c4926e91b1bd89af899fec4301d05eb3 Mon Sep 17 00:00:00 2001 From: Isabella Gottardi Date: Fri, 28 Jul 2017 17:24:08 +0100 Subject: COMPMID-417 - Bug Fix WarpPerspective kernel Change-Id: Ic26fb3b1b60c1a1f4848d683862a25bd1ebc2cc8 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/82053 Tested-by: Kaizen Reviewed-by: Steven Niu --- arm_compute/core/NEON/kernels/NEWarpKernel.h | 3 +++ src/core/CL/cl_kernels/warp_perspective.cl | 4 ++-- src/core/CL/kernels/CLWarpPerspectiveKernel.cpp | 4 ++-- src/core/NEON/kernels/NEWarpKernel.cpp | 11 ++++++++--- 4 files changed, 15 insertions(+), 7 deletions(-) diff --git a/arm_compute/core/NEON/kernels/NEWarpKernel.h b/arm_compute/core/NEON/kernels/NEWarpKernel.h index 10fed1d450..39f6496edb 100644 --- a/arm_compute/core/NEON/kernels/NEWarpKernel.h +++ b/arm_compute/core/NEON/kernels/NEWarpKernel.h @@ -60,6 +60,9 @@ public: // Inherited methods overridden: void run(const Window &window) override; + // Inherited methods overridden: + BorderSize border_size() const override; + protected: /** function to perform warp affine or warp perspective on the given window when border mode == UNDEFINED * diff --git a/src/core/CL/cl_kernels/warp_perspective.cl b/src/core/CL/cl_kernels/warp_perspective.cl index 863b6c9e96..d955e427c4 100644 --- a/src/core/CL/cl_kernels/warp_perspective.cl +++ b/src/core/CL/cl_kernels/warp_perspective.cl @@ -92,7 +92,7 @@ __kernel void warp_perspective_nearest_neighbour( { Image in = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in); Image out = CONVERT_TO_IMAGE_STRUCT(out); - vstore4(read_texels4(&in, convert_int8(clamp_to_border(apply_perspective_transform(get_current_coords(), build_perspective_mtx()), width, height))), 0, out.ptr); + vstore4(read_texels4(&in, convert_int8_rtn(clamp_to_border(apply_perspective_transform(get_current_coords(), build_perspective_mtx()), width, height))), 0, out.ptr); } /** Performs a perspective transform on an image interpolating with the BILINEAR method. Input and output are single channel U8. @@ -124,5 +124,5 @@ __kernel void warp_perspective_bilinear( { Image in = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in); Image out = CONVERT_TO_IMAGE_STRUCT(out); - vstore4(bilinear_interpolate(&in, clamp_to_border(apply_perspective_transform(get_current_coords(), build_perspective_mtx()), width, height), width, height), 0, out.ptr); + vstore4(bilinear_interpolate(&in, apply_perspective_transform(get_current_coords(), build_perspective_mtx()), width, height), 0, out.ptr); } diff --git a/src/core/CL/kernels/CLWarpPerspectiveKernel.cpp b/src/core/CL/kernels/CLWarpPerspectiveKernel.cpp index fddb580750..a47952fc6b 100644 --- a/src/core/CL/kernels/CLWarpPerspectiveKernel.cpp +++ b/src/core/CL/kernels/CLWarpPerspectiveKernel.cpp @@ -88,8 +88,8 @@ void CLWarpPerspectiveKernel::configure(const ICLTensor *input, ICLTensor *outpu Window win = calculate_max_window(*output->info(), Steps(num_elems_processed_per_iteration)); - AccessWindowHorizontal input_access(input->info(), 0, num_elems_processed_per_iteration); - AccessWindowStatic output_access(output->info(), 0, 0, output->info()->dimension(0), output->info()->dimension(1)); + AccessWindowStatic input_access(input->info(), -border_size().left, -border_size().top, input->info()->dimension(0) + border_size().right, input->info()->dimension(1) + border_size().bottom); + AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration); update_window_and_padding(win, input_access, output_access); diff --git a/src/core/NEON/kernels/NEWarpKernel.cpp b/src/core/NEON/kernels/NEWarpKernel.cpp index 6c90a334af..8f8c852672 100644 --- a/src/core/NEON/kernels/NEWarpKernel.cpp +++ b/src/core/NEON/kernels/NEWarpKernel.cpp @@ -49,6 +49,11 @@ INEWarpKernel::INEWarpKernel() { } +BorderSize INEWarpKernel::border_size() const +{ + return BorderSize(1); +} + void INEWarpKernel::run(const Window &window) { ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); @@ -93,9 +98,9 @@ void INEWarpKernel::configure(const ITensor *input, ITensor *output, const float // Reads can occur within the valid region of the input AccessWindowStatic input_access(input->info(), - input_valid_region.anchor[0], input_valid_region.anchor[1], - input_valid_region.anchor[0] + input_valid_region.shape[0], - input_valid_region.anchor[1] + input_valid_region.shape[1]); + input_valid_region.anchor[0] - border_size().left, input_valid_region.anchor[1] - border_size().top, + input_valid_region.anchor[0] + input_valid_region.shape[0] + border_size().right, + input_valid_region.anchor[1] + input_valid_region.shape[1] + border_size().bottom); AccessWindowHorizontal output_access(output->info(), 0, 1); update_window_and_padding(win, input_access, output_access); -- cgit v1.2.1