diff options
author | Isabella Gottardi <isabella.gottardi@arm.com> | 2017-07-28 17:24:08 +0100 |
---|---|---|
committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-09-17 14:16:42 +0100 |
commit | f9bae2e9c4926e91b1bd89af899fec4301d05eb3 (patch) | |
tree | 446f0545552a0473a1716b33202f73e642aa23d4 | |
parent | 02dfb2cd0bfab1bcff141db4598c23055e67311b (diff) | |
download | ComputeLibrary-f9bae2e9c4926e91b1bd89af899fec4301d05eb3.tar.gz |
COMPMID-417 - Bug Fix WarpPerspective kernel
Change-Id: Ic26fb3b1b60c1a1f4848d683862a25bd1ebc2cc8
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/82053
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Steven Niu <steven.niu@arm.com>
-rw-r--r-- | arm_compute/core/NEON/kernels/NEWarpKernel.h | 3 | ||||
-rw-r--r-- | src/core/CL/cl_kernels/warp_perspective.cl | 4 | ||||
-rw-r--r-- | src/core/CL/kernels/CLWarpPerspectiveKernel.cpp | 4 | ||||
-rw-r--r-- | 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); |