From fc1ffe626278e75ba11c803280a6ef46a5bd1ad6 Mon Sep 17 00:00:00 2001 From: Pablo Tello Date: Tue, 23 Jan 2018 08:31:41 +0000 Subject: 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 Reviewed-by: Georgios Pinitas Reviewed-by: Anthony Barbier --- src/core/NEON/kernels/NERemapKernel.cpp | 50 +++++++++++++++----------------- tests/validation/NEON/Remap.cpp | 6 ++-- tests/validation/fixtures/RemapFixture.h | 21 +++++++------- tests/validation/reference/Remap.cpp | 19 +++++++----- 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 tolerance_value(1); -constexpr float tolerance_number = 0.2f; +constexpr AbsoluteTolerance 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 distribution(0, 255); - T constant_border_value = static_cast(distribution(gen)); + const T constant_border_value = static_cast(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 - 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 { 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(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 remap(const SimpleTensor &in, SimpleTensor &map_x, Sim T constant_border_value) { ARM_COMPUTE_ERROR_ON_MSG(border_mode == BorderMode::REPLICATE, "BorderMode not supported"); - SimpleTensor 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(std::floor(map_x[idx]))); src_idx.set(1, static_cast(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 remap(const SimpleTensor &in, SimpleTensor &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"); -- cgit v1.2.1