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>

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);