COMPMID-837: Fixed remap tests failures in Valgrind.

Some minor improvements in the test fixture, for example making sure
the values in the mapx and mapy tensors are in the range of [-5, in_width+5]
and [-5,in_height].

Tolerance was changed to 0, no mismatches expected.

Change-Id: I2fad06defb293bf9fdd1988799b19547c102dee5
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/118044
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>

4 files changed, 49 insertions, 47 deletions

diff --git a/src/core/NEON/kernels/NERemapKernel.cpp b/src/core/NEON/kernels/NERemapKernel.cpp
index 9b8d931b39..66115bb8fc 100644
--- a/src/core/NEON/kernels/NERemapKernel.cpp
+++ b/src/core/NEON/kernels/NERemapKernel.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -46,7 +46,7 @@ namespace
 {
 inline int32x4_t offset_nearest_interpolation(const float *mapx_ptr, const float *mapy_ptr, const float32x4_t &width, const float32x4_t &height, const int32x4_t &stride)
 {
-    static const float32x4_t lowerxy = vdupq_n_f32(-1.0f);
+    const float32x4_t lowerxy = vdupq_n_f32(-1.f);
 
     float32x4_t x = vld1q_f32(mapx_ptr);
     float32x4_t y = vld1q_f32(mapy_ptr);
@@ -113,11 +113,10 @@ void NERemapKernel::configure(const ITensor *input, const ITensor *map_x, const
     AccessWindowStatic input_access(input->info(), -border_size().left, -border_size().top, access_right, input->info()->dimension(1) + border_size().bottom);
 
     AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration);
+    AccessWindowHorizontal mapx_access(map_x->info(), 0, 0, num_elems_processed_per_iteration);
+    AccessWindowHorizontal mapy_access(map_y->info(), 0, 0, num_elems_processed_per_iteration);
 
-    update_window_and_padding(win, input_access,
-                              AccessWindowRectangle(map_x->info(), 0, 0, num_elems_processed_per_iteration, 1),
-                              AccessWindowRectangle(map_y->info(), 0, 0, num_elems_processed_per_iteration, 1),
-                              output_access);
+    update_window_and_padding(win, input_access, mapx_access, mapy_access, output_access);
 
     output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape()));
 
@@ -152,27 +151,24 @@ void NERemapKernel::remap_nearest(const Window &window)
         const int32x4_t offset2 = offset_nearest_interpolation(mapx_ptr + 8, mapy_ptr + 8, width, height, in_stride);
         const int32x4_t offset3 = offset_nearest_interpolation(mapx_ptr + 12, mapy_ptr + 12, width, height, in_stride);
 
-        uint8x8_t tmp0 = vdup_n_u8(0);
-        tmp0           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset0, 0)], tmp0, 0);
-        tmp0           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset0, 1)], tmp0, 1);
-        tmp0           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset0, 2)], tmp0, 2);
-        tmp0           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset0, 3)], tmp0, 3);
-        tmp0           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset1, 0)], tmp0, 4);
-        tmp0           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset1, 1)], tmp0, 5);
-        tmp0           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset1, 2)], tmp0, 6);
-        tmp0           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset1, 3)], tmp0, 7);
-
-        uint8x8_t tmp1 = vdup_n_u8(0);
-        tmp1           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset2, 0)], tmp1, 0);
-        tmp1           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset2, 1)], tmp1, 1);
-        tmp1           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset2, 2)], tmp1, 2);
-        tmp1           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset2, 3)], tmp1, 3);
-        tmp1           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset3, 0)], tmp1, 4);
-        tmp1           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset3, 1)], tmp1, 5);
-        tmp1           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset3, 2)], tmp1, 6);
-        tmp1           = vset_lane_u8(in_ptr[vgetq_lane_s32(offset3, 3)], tmp1, 7);
-
-        vst1q_u8(out.ptr(), vcombine_u8(tmp0, tmp1));
+        uint8x16_t tmp = vdupq_n_u8(0);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset0, 0)], tmp, 0);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset0, 1)], tmp, 1);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset0, 2)], tmp, 2);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset0, 3)], tmp, 3);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset1, 0)], tmp, 4);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset1, 1)], tmp, 5);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset1, 2)], tmp, 6);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset1, 3)], tmp, 7);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset2, 0)], tmp, 8);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset2, 1)], tmp, 9);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset2, 2)], tmp, 10);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset2, 3)], tmp, 11);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset3, 0)], tmp, 12);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset3, 1)], tmp, 13);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset3, 2)], tmp, 14);
+        tmp            = vsetq_lane_u8(in_ptr[vgetq_lane_s32(offset3, 3)], tmp, 15);
+        vst1q_u8(out.ptr(), tmp);
     },
     in, out, mapx, mapy);
 }
diff --git a/tests/validation/NEON/Remap.cpp b/tests/validation/NEON/Remap.cpp
index 6e58000d54..2e54b1152e 100644
--- a/tests/validation/NEON/Remap.cpp
+++ b/tests/validation/NEON/Remap.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -43,8 +43,8 @@ namespace validation
 {
 namespace
 {
-constexpr AbsoluteTolerance<uint8_t> tolerance_value(1);
-constexpr float                      tolerance_number = 0.2f;
+constexpr AbsoluteTolerance<uint8_t> tolerance_value(0);
+constexpr float                      tolerance_number = 0.f;
 } // namespace
 
 TEST_SUITE(NEON)
diff --git a/tests/validation/fixtures/RemapFixture.h b/tests/validation/fixtures/RemapFixture.h
index 846ebf44a2..78b30151ac 100644
--- a/tests/validation/fixtures/RemapFixture.h
+++ b/tests/validation/fixtures/RemapFixture.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -50,7 +50,7 @@ public:
     {
         std::mt19937                           gen(library->seed());
         std::uniform_int_distribution<uint8_t> distribution(0, 255);
-        T                                      constant_border_value = static_cast<T>(distribution(gen));
+        const T                                constant_border_value = static_cast<T>(distribution(gen));
 
         _target    = compute_target(shape, policy, data_type, border_mode, constant_border_value);
         _reference = compute_reference(shape, policy, data_type, border_mode, constant_border_value);
@@ -58,9 +58,10 @@ public:
 
 protected:
     template <typename U>
-    void fill(U &&tensor, int i)
+    void fill(U &&tensor, int i, float min, float max)
     {
-        library->fill_tensor_uniform(tensor, i);
+        std::uniform_int_distribution<> distribution((int)min, (int)max);
+        library->fill(tensor, distribution, i);
     }
 
     TensorType compute_target(const TensorShape &shape, InterpolationPolicy policy, DataType data_type, BorderMode border_mode, T constant_border_value)
@@ -92,9 +93,9 @@ protected:
         ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS);
 
         // Fill tensors
-        fill(AccessorType(src), 0);
-        fill(AccessorType(map_x), 1);
-        fill(AccessorType(map_y), 2);
+        fill(AccessorType(src), 0, 0, 255);
+        fill(AccessorType(map_x), 1, -5, shape.x() + 5);
+        fill(AccessorType(map_y), 2, -5, shape.y() + 5);
 
         // Compute function
         remap.run();
@@ -115,9 +116,9 @@ protected:
         _valid_mask = SimpleTensor<T> { shape, data_type };
 
         // Fill reference
-        fill(src, 0);
-        fill(map_x, 1);
-        fill(map_y, 2);
+        fill(src, 0, 0, 255);
+        fill(map_x, 1, -5, shape.x() + 5);
+        fill(map_y, 2, -5, shape.y() + 5);
 
         // Compute reference
         return reference::remap<T>(src, map_x, map_y, _valid_mask, policy, border_mode, constant_border_value);
diff --git a/tests/validation/reference/Remap.cpp b/tests/validation/reference/Remap.cpp
index bef5962fbf..f862c13700 100644
--- a/tests/validation/reference/Remap.cpp
+++ b/tests/validation/reference/Remap.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2018 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -42,16 +42,15 @@ SimpleTensor<T> remap(const SimpleTensor<T> &in, SimpleTensor<float> &map_x, Sim
                       T constant_border_value)
 {
     ARM_COMPUTE_ERROR_ON_MSG(border_mode == BorderMode::REPLICATE, "BorderMode not supported");
-
     SimpleTensor<T> out(in.shape(), in.data_type());
-
+    ARM_COMPUTE_ERROR_ON(out.num_elements() != map_x.num_elements());
     const int width  = in.shape().x();
     const int height = in.shape().y();
-
     for(int idx = 0; idx < out.num_elements(); idx++)
     {
-        valid_mask[idx] = 1;
-        Coordinates src_idx;
+        const Coordinates id_out = index2coord(out.shape(), idx);
+        valid_mask[idx]          = 1;
+        Coordinates src_idx      = id_out; // need to setup all coordinates and not just xy
         src_idx.set(0, static_cast<int>(std::floor(map_x[idx])));
         src_idx.set(1, static_cast<int>(std::floor(map_y[idx])));
         if((0 <= map_y[idx]) && (map_y[idx] < height) && (0 <= map_x[idx]) && (map_x[idx] < width))
@@ -59,11 +58,17 @@ SimpleTensor<T> remap(const SimpleTensor<T> &in, SimpleTensor<float> &map_x, Sim
             switch(policy)
             {
                 case InterpolationPolicy::NEAREST_NEIGHBOR:
+                {
                     out[idx] = tensor_elem_at(in, src_idx, border_mode, constant_border_value);
                     break;
+                }
                 case InterpolationPolicy::BILINEAR:
-                    (valid_bilinear_policy(map_x[idx], map_y[idx], width, height, border_mode)) ? out[idx] = bilinear_policy(in, src_idx, map_x[idx], map_y[idx], border_mode, constant_border_value) : valid_mask[idx] = 0;
+                {
+                    (valid_bilinear_policy(map_x[idx], map_y[idx], width, height, border_mode)) ?
+                    out[idx]        = bilinear_policy(in, src_idx, map_x[idx], map_y[idx], border_mode, constant_border_value) :
+                                      valid_mask[idx] = 0;
                     break;
+                }
                 case InterpolationPolicy::AREA:
                 default:
                     ARM_COMPUTE_ERROR("Interpolation not supported");