COMPMID-661: issue# 23 Scale border fix (#26)

Changes in CL and reference in terms of border handling.

Change-Id: I5bed95b1f4c308629d7113455dc8a55d74500bcd
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/95742
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>

6 files changed, 92 insertions, 47 deletions

diff --git a/src/core/CL/cl_kernels/scale.cl b/src/core/CL/cl_kernels/scale.cl
index 0106ce095c..d533d970c7 100644
--- a/src/core/CL/cl_kernels/scale.cl
+++ b/src/core/CL/cl_kernels/scale.cl
@@ -84,7 +84,7 @@ __kernel void scale_nearest_neighbour(
     Image        in  = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in);
     Image        out = CONVERT_TO_IMAGE_STRUCT(out);
     const float2 r   = (float2)(scale_x, scale_y);
-    const float8 tc  = clamp_to_border(transform_nearest(get_current_coords(), r), input_width, input_height);
+    const float8 tc  = clamp_to_border_with_size(transform_nearest(get_current_coords(), r), input_width, input_height, BORDER_SIZE);
     vstore4(read_texels4(&in, convert_int8(tc)), 0, (__global DATA_TYPE *)out.ptr);
 }
 
@@ -119,5 +119,5 @@ __kernel void scale_bilinear(
     Image        out = CONVERT_TO_IMAGE_STRUCT(out);
     const float2 r   = (float2)(scale_x, scale_y);
     const float8 tc  = transform_bilinear(get_current_coords(), r);
-    vstore4(bilinear_interpolate(&in, tc, input_width, input_height), 0, (__global DATA_TYPE *)out.ptr);
+    vstore4(bilinear_interpolate_with_border(&in, tc, input_width, input_height, BORDER_SIZE), 0, (__global DATA_TYPE *)out.ptr);
 }
diff --git a/src/core/CL/cl_kernels/warp_helpers.h b/src/core/CL/cl_kernels/warp_helpers.h
index 26a8b859a4..76f0a4a104 100644
--- a/src/core/CL/cl_kernels/warp_helpers.h
+++ b/src/core/CL/cl_kernels/warp_helpers.h
@@ -23,6 +23,22 @@
  */
 #include "helpers.h"
 
+/** Clamps the given coordinates to the borders according to the border size.
+ *
+ * @param[in] coords      Vector of 2D coordinates to clamp. Even positions are X coords, odd positions are Y coords.
+ * @param[in] width       Width of the image
+ * @param[in] height      Height of the image
+ * @param[in] border_size Border size of the image
+ *
+ */
+inline const float8 clamp_to_border_with_size(float8 coords, const float width, const float height, const float border_size)
+{
+    const float4 clamped_x = clamp(coords.even, 0.0f - border_size, width - 1 + border_size);
+    const float4 clamped_y = clamp(coords.odd, 0.0f - border_size, height - 1 + border_size);
+    return (float8)(clamped_x.s0, clamped_y.s0, clamped_x.s1, clamped_y.s1, clamped_x.s2, clamped_y.s2, clamped_x.s3, clamped_y.s3);
+}
+
+/* FIXME(COMPMID-682): Clamp border properly in UNDEFINED border mode in Warp, Scale, Remap */
 /** Clamps the given coordinates to the borders.
  *
  * @param[in] coords Vector of 2D coordinates to clamp. Even positions are X coords, odd positions are Y coords.
@@ -32,9 +48,7 @@
  */
 inline const float8 clamp_to_border(float8 coords, const float width, const float height)
 {
-    const float4 clamped_x = clamp(coords.even, -1.0f, width);
-    const float4 clamped_y = clamp(coords.odd, -1.0f, height);
-    return (float8)(clamped_x.s0, clamped_y.s0, clamped_x.s1, clamped_y.s1, clamped_x.s2, clamped_y.s2, clamped_x.s3, clamped_y.s3);
+    return clamp_to_border_with_size(coords, width, height, 1);
 }
 
 /** Reads four texels from the input image. The coords vector is used to determine which texels to be read.
@@ -72,23 +86,25 @@ inline const float8 get_neighbour_coords(const float2 coord)
 
 /** Computes the bilinear interpolation for each set of coordinates in the vector coords and returns the values
  *
- * @param[in] in     Pointer to the source image.
- * @param[in] coords Vector of four 2D coordinates. Even pos is x and odd y.
- * @param[in] width  Width of the image
- * @param[in] height Height of the image
+ * @param[in] in          Pointer to the source image.
+ * @param[in] coords      Vector of four 2D coordinates. Even pos is x and odd y.
+ * @param[in] width       Width of the image
+ * @param[in] height      Height of the image
+ * @param[in] border_size Border size
 */
-inline const VEC_DATA_TYPE(DATA_TYPE, 4) bilinear_interpolate(const Image *in, const float8 coords, const float width, const float height)
+inline const VEC_DATA_TYPE(DATA_TYPE, 4) bilinear_interpolate_with_border(const Image *in, const float8 coords, const float width, const float height, const float border_size)
 {
     // If any of the 4 texels is out of the image's boundaries we use the border value (REPLICATE or CONSTANT) for any texel out of the image.
 
     // Sets the 4x4 coordinates for each of the four input texels
     const float8  fc = floor(coords);
     const float16 c1 = (float16)(
-                           clamp_to_border(get_neighbour_coords((float2)(fc.s0, fc.s1)), width, height),
-                           clamp_to_border(get_neighbour_coords((float2)(fc.s2, fc.s3)), width, height));
+                           clamp_to_border_with_size(get_neighbour_coords((float2)(fc.s0, fc.s1)), width, height, border_size),
+                           clamp_to_border_with_size(get_neighbour_coords((float2)(fc.s2, fc.s3)), width, height, border_size));
     const float16 c2 = (float16)(
-                           clamp_to_border(get_neighbour_coords((float2)(fc.s4, fc.s5)), width, height),
-                           clamp_to_border(get_neighbour_coords((float2)(fc.s6, fc.s7)), width, height));
+                           clamp_to_border_with_size(get_neighbour_coords((float2)(fc.s4, fc.s5)), width, height, border_size),
+                           clamp_to_border_with_size(get_neighbour_coords((float2)(fc.s6, fc.s7)), width, height, border_size));
+
     // Loads the values from the input image
     const float16 t = (float16)(
                           /* tl, tr, bl, br */
@@ -109,3 +125,16 @@ inline const VEC_DATA_TYPE(DATA_TYPE, 4) bilinear_interpolate(const Image *in, c
                           ((t.sc * b.s6 * b.s7) + (t.sd * a.s6 * b.s7) + (t.se * b.s6 * a.s7) + (t.sf * a.s6 * a.s7)));
     return CONVERT(fr, VEC_DATA_TYPE(DATA_TYPE, 4));
 }
+
+/* FIXME(COMPMID-682): Clamp border properly in UNDEFINED border mode in Warp, Scale, Remap */
+/** Computes the bilinear interpolation for each set of coordinates in the vector coords and returns the values
+ *
+ * @param[in] in     Pointer to the source image.
+ * @param[in] coords Vector of four 2D coordinates. Even pos is x and odd y.
+ * @param[in] width  Width of the image
+ * @param[in] height Height of the image
+*/
+inline const VEC_DATA_TYPE(DATA_TYPE, 4) bilinear_interpolate(const Image *in, const float8 coords, const float width, const float height)
+{
+    return bilinear_interpolate_with_border(in, coords, width, height, 1);
+}
diff --git a/src/core/CL/kernels/CLScaleKernel.cpp b/src/core/CL/kernels/CLScaleKernel.cpp
index 82ebe644ea..6a5d24c943 100644
--- a/src/core/CL/kernels/CLScaleKernel.cpp
+++ b/src/core/CL/kernels/CLScaleKernel.cpp
@@ -54,11 +54,14 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo
     _input  = input;
     _output = output;
 
-    /* Compute the ratio between source width/height and destination width/height */
+    // Compute the ratio between source width/height and destination width/height
     const auto wr = static_cast<float>(input->info()->dimension(0)) / static_cast<float>(output->info()->dimension(0));
     const auto hr = static_cast<float>(input->info()->dimension(1)) / static_cast<float>(output->info()->dimension(1));
 
-    /* Area interpolation behaves as Nearest Neighbour in case of up-sampling */
+    // Compute actual border size
+    BorderSize border = border_undefined ? BorderSize(0) : border_size();
+
+    // Area interpolation behaves as Nearest Neighbour in case of up-sampling
     if(policy == InterpolationPolicy::AREA && wr <= 1.f && hr <= 1.f)
     {
         policy = InterpolationPolicy::NEAREST_NEIGHBOR;
@@ -69,11 +72,14 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo
     }
 
     // Create kernel
-    std::set<std::string> build_opts         = { ("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())) };
-    std::string           interpolation_name = string_from_interpolation_policy(policy);
+    CLBuildOptions build_opts;
+    build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type()));
+    build_opts.add_option("-DBORDER_SIZE=" + support::cpp11::to_string(border.right));
+
+    std::string interpolation_name = string_from_interpolation_policy(policy);
     std::transform(interpolation_name.begin(), interpolation_name.end(), interpolation_name.begin(), ::tolower);
     std::string kernel_name = "scale_" + interpolation_name;
-    _kernel                 = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name, build_opts));
+    _kernel                 = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options()));
 
     // Configure kernel window
     constexpr unsigned int num_elems_processed_per_iteration = 4;
@@ -84,15 +90,18 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo
 
     // Reads can occur within the valid region of the input
     AccessWindowStatic input_access(input->info(),
-                                    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);
+                                    input_valid_region.anchor[0] - border.left, input_valid_region.anchor[1] - border.top,
+                                    input_valid_region.anchor[0] + input_valid_region.shape[0] + border.right,
+                                    input_valid_region.anchor[1] + input_valid_region.shape[1] + border.bottom);
 
     AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration);
 
     update_window_and_padding(win, input_access, output_access);
 
-    output_access.set_valid_region(win, calculate_valid_region_scale(*(input->info()), output->info()->tensor_shape(), policy, border_size(),
+    output_access.set_valid_region(win, calculate_valid_region_scale(*(input->info()),
+                                                                     output->info()->tensor_shape(),
+                                                                     policy,
+                                                                     border,
                                                                      border_undefined));
 
     ICLKernel::configure(win);
diff --git a/tests/validation/CL/Scale.cpp b/tests/validation/CL/Scale.cpp
index f43f2aeeb9..1ddf03a74c 100644
--- a/tests/validation/CL/Scale.cpp
+++ b/tests/validation/CL/Scale.cpp
@@ -92,16 +92,18 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(combine(combi
     CLScale clscale;
     clscale.configure(&src, &dst, policy, border_mode, constant_border_value);
 
-    // Validate valid region
-    const ValidRegion dst_valid_region = calculate_valid_region_scale(*(src.info()), shape_scaled, policy, BorderSize(1), (border_mode == BorderMode::UNDEFINED));
+    // Get border size depending on border mode
+    const BorderSize border_size(border_mode == BorderMode::UNDEFINED ? 0 : 1);
 
+    // Validate valid region
+    const ValidRegion dst_valid_region = calculate_valid_region_scale(*(src.info()), shape_scaled, policy, border_size, (border_mode == BorderMode::UNDEFINED));
     validate(dst.info()->valid_region(), dst_valid_region);
 
     // Validate padding
     PaddingCalculator calculator(shape_scaled.x(), 4);
     calculator.set_border_mode(border_mode);
 
-    const PaddingSize read_padding(1);
+    const PaddingSize read_padding(border_size);
     const PaddingSize write_padding = calculator.required_padding(PaddingCalculator::Option::EXCLUDE_BORDER);
     validate(src.info()->padding(), read_padding);
     validate(dst.info()->padding(), write_padding);
diff --git a/tests/validation/CPP/Scale.cpp b/tests/validation/CPP/Scale.cpp
index 74489aaa96..0da7497277 100644
--- a/tests/validation/CPP/Scale.cpp
+++ b/tests/validation/CPP/Scale.cpp
@@ -50,6 +50,9 @@ SimpleTensor<T> scale(const SimpleTensor<T> &in, float scale_x, float scale_y, I
     const auto width  = static_cast<int>(in.shape().x());
     const auto height = static_cast<int>(in.shape().y());
 
+    // Determine border size
+    const int border_size = (border_mode == BorderMode::UNDEFINED) ? 0 : 1;
+
     // Area interpolation behaves as Nearest Neighbour in case of up-sampling
     if(policy == InterpolationPolicy::AREA && wr <= 1.f && hr <= 1.f)
     {
@@ -75,30 +78,17 @@ SimpleTensor<T> scale(const SimpleTensor<T> &in, float scale_x, float scale_y, I
                 id.set(1, y_src);
 
                 // If coordinates in range of tensor's width or height
-                if(x_src >= -1 || y_src >= -1 || x_src <= width || y_src <= height)
+                if(is_valid_pixel_index(x_src, y_src, width, height, border_size))
                 {
                     out[element_idx] = tensor_elem_at(in, id, border_mode, constant_border_value);
                 }
-                else
-                {
-                    if(border_mode == BorderMode::CONSTANT)
-                    {
-                        out[element_idx] = constant_border_value;
-                    }
-                    else if(border_mode == BorderMode::REPLICATE)
-                    {
-                        id.set(0, clamp(static_cast<int>(x_src), 0, width - 1));
-                        id.set(1, clamp(static_cast<int>(y_src), 0, height - 1));
-                        out[element_idx] = in[coord2index(in.shape(), id)];
-                    }
-                }
                 break;
             }
             case InterpolationPolicy::BILINEAR:
             {
                 id.set(0, std::floor(x_src));
                 id.set(1, std::floor(y_src));
-                if(x_src >= -1 || y_src >= -1 || x_src <= width || y_src <= height)
+                if(is_valid_pixel_index(x_src, y_src, width, height, border_size))
                 {
                     out[element_idx] = bilinear_policy(in, id, x_src, y_src, border_mode, constant_border_value);
                 }
@@ -127,14 +117,14 @@ SimpleTensor<T> scale(const SimpleTensor<T> &in, float scale_x, float scale_y, I
                 const int yi     = std::floor(y_src);
 
                 // Clamp position to borders
-                x_src = std::max(-1.f, std::min(x_src, static_cast<float>(width)));
-                y_src = std::max(-1.f, std::min(y_src, static_cast<float>(height)));
+                x_src = std::max(-static_cast<float>(border_size), std::min(x_src, static_cast<float>(width - 1 + border_size)));
+                y_src = std::max(-static_cast<float>(border_size), std::min(y_src, static_cast<float>(height - 1 + border_size)));
 
                 // Clamp bounding box offsets to borders
-                x_from = ((x_src + x_from) < -1) ? -1 : x_from;
-                y_from = ((y_src + y_from) < -1) ? -1 : y_from;
-                x_to   = ((x_src + x_to) > width) ? (width - x_src) : x_to;
-                y_to   = ((y_src + y_to) > height) ? (height - y_src) : y_to;
+                x_from = ((x_src + x_from) < -border_size) ? -border_size : x_from;
+                y_from = ((y_src + y_from) < -border_size) ? -border_size : y_from;
+                x_to   = ((x_src + x_to) >= (width + border_size)) ? (width - 1 + border_size) : x_to;
+                y_to   = ((y_src + y_to) >= (height + border_size)) ? (height - 1 + border_size) : y_to;
                 ARM_COMPUTE_ERROR_ON((x_to - x_from + 1) == 0 || (y_to - y_from + 1) == 0);
 
                 float sum = 0;
diff --git a/tests/validation/CPP/Utils.h b/tests/validation/CPP/Utils.h
index 0733411462..2aa77c6ff7 100644
--- a/tests/validation/CPP/Utils.h
+++ b/tests/validation/CPP/Utils.h
@@ -41,6 +41,21 @@ namespace test
 {
 namespace validation
 {
+/** Checks if a pixel has valid coordinates
+ *
+ * @param x            X coordinate
+ * @param y            Y coordinate
+ * @param width        Width of the image
+ * @param height       Height of the image
+ * @param border_size  Border size
+ *
+ * @return True if pixel is valid else false
+ */
+inline bool is_valid_pixel_index(int x, int y, int width, int height, int border_size)
+{
+    return ((x >= -border_size) && (y >= -border_size) && (x < (width + border_size)) && (y < height + border_size));
+}
+
 // Return a tensor element at a specified coordinate with different border modes
 template <typename T>
 T tensor_elem_at(const SimpleTensor<T> &src, Coordinates coord, BorderMode border_mode, T constant_border_value)