From 5c09ae8d3f82ddf8bf9438fc67e1991b62ae88e3 Mon Sep 17 00:00:00 2001 From: Luca Foschiani Date: Tue, 21 Jan 2020 16:24:16 +0000 Subject: COMPMID-2809 Added support for QASYMM8_SIGNED in NEUpsampleLayer Signed-off-by: Luca Foschiani Change-Id: I61638cb1e4f65130b150a456dc1cdb90099b0efa Reviewed-on: https://review.mlplatform.org/c/2615 Comments-Addressed: Arm Jenkins Reviewed-by: Michele Di Giorgio Tested-by: Arm Jenkins --- .../core/NEON/kernels/NEUpsampleLayerKernel.h | 37 +-- .../runtime/NEON/functions/NEUpsampleLayer.h | 6 +- src/core/NEON/kernels/NEUpsampleLayerKernel.cpp | 248 ++++++--------------- tests/validation/NEON/Upsample.cpp | 16 +- tests/validation/fixtures/UpsampleLayerFixture.h | 13 +- 5 files changed, 100 insertions(+), 220 deletions(-) diff --git a/arm_compute/core/NEON/kernels/NEUpsampleLayerKernel.h b/arm_compute/core/NEON/kernels/NEUpsampleLayerKernel.h index 50c3e342e3..1ea3f974e7 100644 --- a/arm_compute/core/NEON/kernels/NEUpsampleLayerKernel.h +++ b/arm_compute/core/NEON/kernels/NEUpsampleLayerKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -74,40 +74,19 @@ public: void run(const Window &window, const ThreadInfo &info) override; private: - /** Function to run upsample layer for FP32 (NCHW) + /** Function to run upsample layer (NCHW) * * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()). */ - void upsample_f32_nchw(const Window &window); - /** Function to run upsample layer for FP32 (NHWC) + template + void upsample_nchw(const Window &window); + /** Function to run upsample layer (NHWC) * * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()). */ - void upsample_f32_nhwc(const Window &window); - /** Function to run upsample layer for FP16 (NCHW) - * - * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()). - */ - void upsample_f16_nchw(const Window &window); - /** Function to run upsample layer for FP16 (NHWC) - * - * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()). - */ - void upsample_f16_nhwc(const Window &window); - /** Function to run upsample layer for QASYMM8 (NCHW) - * - * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()). - */ - void upsample_qasymm8_nchw(const Window &window); - /** Function to run upsample layer for QASYMM8 (NHWC) - * - * @param[in] window Region on which to execute the kernel. (Must be a valid region of the window returned by window()). - */ - void upsample_qasymm8_nhwc(const Window &window); - /** Common signature for all the upsample layer functions - * - * @param[in] window Region on which to execute the kernel. - */ + template + void upsample_nhwc(const Window &window); + using UpsampleFunctionPtr = void (NEUpsampleLayerKernel::*)(const Window &window); private: diff --git a/arm_compute/runtime/NEON/functions/NEUpsampleLayer.h b/arm_compute/runtime/NEON/functions/NEUpsampleLayer.h index 6cd6ba3e9a..ff465e54a0 100644 --- a/arm_compute/runtime/NEON/functions/NEUpsampleLayer.h +++ b/arm_compute/runtime/NEON/functions/NEUpsampleLayer.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -42,7 +42,7 @@ public: NEUpsampleLayer(); /** Set the input output tensors. * - * @param[in] input Source tensor. Data types supported: QASYMM8/F16/F32. + * @param[in] input Source tensor. Data types supported: QASYMM8_SIGNED/QASYMM8/F16/F32. * @param[out] output Destination tensor. Data types supported: same as @p input. * @param[in] info Contains stride information described in @ref Size2D. * @param[in] policy Defines the policy to fill the intermediate pixels. @@ -52,7 +52,7 @@ public: const InterpolationPolicy &policy); /** Static function to check if given info will lead to a valid configuration of @ref NEUpsampleLayer * - * @param[in] input Source tensor info. Data types supported: QASYMM8/F16/F32. + * @param[in] input Source tensor info. Data types supported: QASYMM8_SIGNED/QASYMM8/F16/F32. * @param[out] output Destination tensor info. Data types supported: same as @p input. * @param[in] info Contains stride information described in @ref Size2D. * @param[in] policy Defines the policy to fill the intermediate pixels. diff --git a/src/core/NEON/kernels/NEUpsampleLayerKernel.cpp b/src/core/NEON/kernels/NEUpsampleLayerKernel.cpp index d3d88b3bf0..3b6faea2c1 100644 --- a/src/core/NEON/kernels/NEUpsampleLayerKernel.cpp +++ b/src/core/NEON/kernels/NEUpsampleLayerKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -31,6 +31,7 @@ #include "arm_compute/core/Validate.h" #include "arm_compute/core/Window.h" #include "arm_compute/core/utils/misc/ShapeCalculator.h" +#include "arm_compute/core/NEON/wrapper/wrapper.h" #include @@ -38,30 +39,16 @@ namespace arm_compute { namespace { -std::pair validate_and_configure_window_nchw(ITensorInfo *input, ITensorInfo *output, int num_elems_processed_per_iteration_x, const Size2D &info) +template +inline T get_data_out(T data, int offset) { - const int num_elems_processed_per_iteration_x_out = num_elems_processed_per_iteration_x * info.x(); - Window win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration_x_out)); - AccessWindowRectangle input_access(input, 0, 0, num_elems_processed_per_iteration_x, 1, 0.5f, 0.5f); - AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration_x_out); - bool window_changed = update_window_and_padding(win, input_access, output_access); - output_access.set_valid_region(win, output->valid_region()); - - Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; - return std::make_pair(err, win); -} - -std::pair validate_and_configure_window_nhwc(ITensorInfo *input, ITensorInfo *output, int num_elems_processed_per_iteration_x, const Size2D &info) -{ - ARM_COMPUTE_UNUSED(info); - Window win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration_x)); - AccessWindowHorizontal input_access(input, 0, num_elems_processed_per_iteration_x); - AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration_x); - bool window_changed = update_window_and_padding(win, input_access, output_access); - output_access.set_valid_region(win, output->valid_region()); - - Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; - return std::make_pair(err, win); + T out{ 0 }; + for(int i = 0; i < S / 2; ++i) + { + out[2 * i] = wrapper::vgetlane(data, i + offset); + out[2 * i + 1] = wrapper::vgetlane(data, i + offset); + } + return out; } std::pair validate_and_configure_window(ITensorInfo *input, ITensorInfo *output, int num_elems_processed_per_iteration_x, const Size2D &info) @@ -70,13 +57,34 @@ std::pair validate_and_configure_window(ITensorInfo *input, ITen switch(input->data_layout()) { case DataLayout::NCHW: - win_config = validate_and_configure_window_nchw(input, output, num_elems_processed_per_iteration_x, info); + { + const int num_elems_processed_per_iteration_x_out = num_elems_processed_per_iteration_x * info.x(); + Window win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration_x_out)); + AccessWindowRectangle input_access(input, 0, 0, num_elems_processed_per_iteration_x, 1, 0.5f, 0.5f); + AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration_x_out); + bool window_changed = update_window_and_padding(win, input_access, output_access); + output_access.set_valid_region(win, output->valid_region()); + + Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; + win_config = std::make_pair(err, win); break; + } case DataLayout::NHWC: - win_config = validate_and_configure_window_nhwc(input, output, num_elems_processed_per_iteration_x, info); + { + Window win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration_x)); + AccessWindowHorizontal input_access(input, 0, num_elems_processed_per_iteration_x); + AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration_x); + bool window_changed = update_window_and_padding(win, input_access, output_access); + output_access.set_valid_region(win, output->valid_region()); + + Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{}; + win_config = std::make_pair(err, win); break; + } default: + { win_config = std::make_pair(ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Unsupported data layout!"), Window{}); + } } return win_config; @@ -97,7 +105,7 @@ Status NEUpsampleLayerKernel::validate(const ITensorInfo *input, const ITensorIn const int idx_height = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT); ARM_COMPUTE_RETURN_ERROR_ON_CPU_F16_UNSUPPORTED(input); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8, DataType::F16, DataType::F32); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8_SIGNED, DataType::QASYMM8, DataType::F16, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON_MSG(info.x() != 2 || info.y() != 2, "Only stride 2 is supported"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(policy != InterpolationPolicy::NEAREST_NEIGHBOR, "Only nearest neighbor policy supported"); @@ -118,60 +126,11 @@ Status NEUpsampleLayerKernel::validate(const ITensorInfo *input, const ITensorIn return Status{}; } -void NEUpsampleLayerKernel::upsample_f32_nchw(const arm_compute::Window &window) -{ - Window window_in(window); - window_in.set(Window::DimX, Window::Dimension(0, _input->info()->dimension(0), _num_elems_processed_per_iteration_x)); - - Window window_out(window); - window_out.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), _info.y())); - - Iterator input(_input, window_in); - Iterator output(_output, window_out); - const int offset_y_out = _output->info()->strides_in_bytes().y() / sizeof(float); - - execute_window_loop(window_out, [&](const Coordinates &) - { - const float32x4_t data = vld1q_f32(reinterpret_cast(input.ptr())); - const float32x4_t data_out1 = { vgetq_lane_f32(data, 0), vgetq_lane_f32(data, 0), vgetq_lane_f32(data, 1), vgetq_lane_f32(data, 1) }; - const float32x4_t data_out2 = { vgetq_lane_f32(data, 2), vgetq_lane_f32(data, 2), vgetq_lane_f32(data, 3), vgetq_lane_f32(data, 3) }; - auto out = reinterpret_cast(output.ptr()); - - vst1q_f32(out, data_out1); - vst1q_f32(out + 4, data_out2); - vst1q_f32(out + offset_y_out, data_out1); - vst1q_f32(out + offset_y_out + 4, data_out2); - }, - input, output); -} - -void NEUpsampleLayerKernel::upsample_f32_nhwc(const arm_compute::Window &window) +template +void NEUpsampleLayerKernel::upsample_nchw(const arm_compute::Window &window) { - Window window_out(window); - window_out.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), _info.x())); - window_out.set(Window::DimZ, Window::Dimension(0, _output->info()->dimension(2), _info.y())); - - Iterator input(_input, window); - Iterator output(_output, window_out); - - const int offset_y_out = _output->info()->strides_in_bytes().y() / sizeof(float); - const int offset_z_out = _output->info()->strides_in_bytes().z() / sizeof(float); + using VectorType = typename wrapper::traits::neon_vector::type; - execute_window_loop(window_out, [&](const Coordinates &) - { - const float32x4_t data = vld1q_f32(reinterpret_cast(input.ptr())); - auto out = reinterpret_cast(output.ptr()); - - vst1q_f32(out, data); - vst1q_f32(out + offset_y_out, data); - vst1q_f32(out + offset_z_out, data); - vst1q_f32(out + offset_y_out + offset_z_out, data); - }, - input, output); -} - -void NEUpsampleLayerKernel::upsample_qasymm8_nchw(const arm_compute::Window &window) -{ Window window_in(window); window_in.set(Window::DimX, Window::Dimension(0, _input->info()->dimension(0), _num_elems_processed_per_iteration_x)); @@ -180,35 +139,28 @@ void NEUpsampleLayerKernel::upsample_qasymm8_nchw(const arm_compute::Window &win Iterator input(_input, window_in); Iterator output(_output, window_out); - const int offset_y_out = _output->info()->strides_in_bytes().y() / sizeof(uint8_t); + const int offset_y_out = _output->info()->strides_in_bytes().y() / sizeof(T); execute_window_loop(window_out, [&](const Coordinates &) { - const uint8x16_t data = vld1q_u8(reinterpret_cast(input.ptr())); - const uint8x16_t data_out1 = { vgetq_lane_u8(data, 0), vgetq_lane_u8(data, 0), vgetq_lane_u8(data, 1), vgetq_lane_u8(data, 1), - vgetq_lane_u8(data, 2), vgetq_lane_u8(data, 2), vgetq_lane_u8(data, 3), vgetq_lane_u8(data, 3), - vgetq_lane_u8(data, 4), vgetq_lane_u8(data, 4), vgetq_lane_u8(data, 5), vgetq_lane_u8(data, 5), - vgetq_lane_u8(data, 6), vgetq_lane_u8(data, 6), vgetq_lane_u8(data, 7), vgetq_lane_u8(data, 7) - }; - const uint8x16_t data_out2 = - { - vgetq_lane_u8(data, 8), vgetq_lane_u8(data, 8), vgetq_lane_u8(data, 9), vgetq_lane_u8(data, 9), - vgetq_lane_u8(data, 10), vgetq_lane_u8(data, 10), vgetq_lane_u8(data, 11), vgetq_lane_u8(data, 11), - vgetq_lane_u8(data, 12), vgetq_lane_u8(data, 12), vgetq_lane_u8(data, 13), vgetq_lane_u8(data, 13), - vgetq_lane_u8(data, 14), vgetq_lane_u8(data, 14), vgetq_lane_u8(data, 15), vgetq_lane_u8(data, 15) - }; - auto out = reinterpret_cast(output.ptr()); - - vst1q_u8(out, data_out1); - vst1q_u8(out + 16, data_out2); - vst1q_u8(out + offset_y_out, data_out1); - vst1q_u8(out + offset_y_out + 16, data_out2); + const VectorType data = wrapper::vloadq(reinterpret_cast(input.ptr())); + const VectorType data_out1 = get_data_out(data, 0); + const VectorType data_out2 = get_data_out(data, S / 2); + auto out = reinterpret_cast(output.ptr()); + + wrapper::vstore(out, data_out1); + wrapper::vstore(out + S, data_out2); + wrapper::vstore(out + offset_y_out, data_out1); + wrapper::vstore(out + offset_y_out + S, data_out2); }, input, output); } -void NEUpsampleLayerKernel::upsample_qasymm8_nhwc(const arm_compute::Window &window) +template +void NEUpsampleLayerKernel::upsample_nhwc(const arm_compute::Window &window) { + using VectorType = typename wrapper::traits::neon_vector::type; + Window window_out(window); window_out.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), _info.x())); window_out.set(Window::DimZ, Window::Dimension(0, _output->info()->dimension(2), _info.y())); @@ -216,82 +168,21 @@ void NEUpsampleLayerKernel::upsample_qasymm8_nhwc(const arm_compute::Window &win Iterator input(_input, window); Iterator output(_output, window_out); - const int offset_y_out = _output->info()->strides_in_bytes().y() / sizeof(uint8_t); - const int offset_z_out = _output->info()->strides_in_bytes().z() / sizeof(uint8_t); + const int offset_y_out = _output->info()->strides_in_bytes().y() / sizeof(T); + const int offset_z_out = _output->info()->strides_in_bytes().z() / sizeof(T); execute_window_loop(window_out, [&](const Coordinates &) { - const uint8x16_t data = vld1q_u8(reinterpret_cast(input.ptr())); - auto out = reinterpret_cast(output.ptr()); + const VectorType data = wrapper::vloadq(reinterpret_cast(input.ptr())); + auto out = reinterpret_cast(output.ptr()); - vst1q_u8(out, data); - vst1q_u8(out + offset_y_out, data); - vst1q_u8(out + offset_z_out, data); - vst1q_u8(out + offset_y_out + offset_z_out, data); + wrapper::vstore(out, data); + wrapper::vstore(out + offset_y_out, data); + wrapper::vstore(out + offset_z_out, data); + wrapper::vstore(out + offset_y_out + offset_z_out, data); }, input, output); } -void NEUpsampleLayerKernel::upsample_f16_nchw(const arm_compute::Window &window) -{ - ARM_COMPUTE_UNUSED(window); -#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC - Window window_in(window); - window_in.set(Window::DimX, Window::Dimension(0, _input->info()->dimension(0), _num_elems_processed_per_iteration_x)); - - Window window_out(window); - window_out.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), _info.y())); - - Iterator input(_input, window_in); - Iterator output(_output, window_out); - const int offset_y_out = _output->info()->strides_in_bytes().y() / sizeof(float16_t); - - execute_window_loop(window_out, [&](const Coordinates &) - { - const float16x8_t data = vld1q_f16(reinterpret_cast(input.ptr())); - const float16x8_t data_out1 = { vgetq_lane_f16(data, 0), vgetq_lane_f16(data, 0), vgetq_lane_f16(data, 1), vgetq_lane_f16(data, 1), - vgetq_lane_f16(data, 2), vgetq_lane_f16(data, 2), vgetq_lane_f16(data, 3), vgetq_lane_f16(data, 3) - }; - const float16x8_t data_out2 = { vgetq_lane_f16(data, 4), vgetq_lane_f16(data, 4), vgetq_lane_f16(data, 5), vgetq_lane_f16(data, 5), - vgetq_lane_f16(data, 6), vgetq_lane_f16(data, 6), vgetq_lane_f16(data, 7), vgetq_lane_f16(data, 7) - }; - auto out = reinterpret_cast(output.ptr()); - - vst1q_f16(out, data_out1); - vst1q_f16(out + 8, data_out2); - vst1q_f16(out + offset_y_out, data_out1); - vst1q_f16(out + offset_y_out + 8, data_out2); - }, - input, output); -#endif // __ARM_FEATURE_FP16_VECTOR_ARITHMETIC -} - -void NEUpsampleLayerKernel::upsample_f16_nhwc(const arm_compute::Window &window) -{ - ARM_COMPUTE_UNUSED(window); -#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC - Window window_out(window); - window_out.set(Window::DimY, Window::Dimension(0, _output->info()->dimension(1), _info.x())); - window_out.set(Window::DimZ, Window::Dimension(0, _output->info()->dimension(2), _info.y())); - - Iterator input(_input, window); - Iterator output(_output, window_out); - const int offset_y_out = _output->info()->strides_in_bytes().y() / sizeof(float16_t); - const int offset_z_out = _output->info()->strides_in_bytes().z() / sizeof(float16_t); - - execute_window_loop(window_out, [&](const Coordinates &) - { - const float16x8_t data = vld1q_f16(reinterpret_cast(input.ptr())); - auto out = reinterpret_cast(output.ptr()); - - vst1q_f16(out, data); - vst1q_f16(out + offset_y_out, data); - vst1q_f16(out + offset_z_out, data); - vst1q_f16(out + offset_y_out + offset_z_out, data); - }, - input, output); -#endif // __ARM_FEATURE_FP16_VECTOR_ARITHMETIC -} - void NEUpsampleLayerKernel::configure(const ITensor *input, ITensor *output, const Size2D &info, const InterpolationPolicy policy) { ARM_COMPUTE_ERROR_ON_NULLPTR(input, output); @@ -318,15 +209,19 @@ void NEUpsampleLayerKernel::configure(const ITensor *input, ITensor *output, con { switch(input->info()->data_type()) { + case DataType::QASYMM8_SIGNED: + _func = &NEUpsampleLayerKernel::upsample_nchw; + break; case DataType::QASYMM8: - _func = &NEUpsampleLayerKernel::upsample_qasymm8_nchw; + _func = &NEUpsampleLayerKernel::upsample_nchw; break; case DataType::F32: - _func = &NEUpsampleLayerKernel::upsample_f32_nchw; + _func = &NEUpsampleLayerKernel::upsample_nchw; break; #ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC case DataType::F16: - _func = &NEUpsampleLayerKernel::upsample_f16_nchw; + _func = &NEUpsampleLayerKernel::upsample_nchw; + ; break; #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ default: @@ -338,15 +233,18 @@ void NEUpsampleLayerKernel::configure(const ITensor *input, ITensor *output, con { switch(input->info()->data_type()) { + case DataType::QASYMM8_SIGNED: + _func = &NEUpsampleLayerKernel::upsample_nhwc; + break; case DataType::QASYMM8: - _func = &NEUpsampleLayerKernel::upsample_qasymm8_nhwc; + _func = &NEUpsampleLayerKernel::upsample_nhwc; break; case DataType::F32: - _func = &NEUpsampleLayerKernel::upsample_f32_nhwc; + _func = &NEUpsampleLayerKernel::upsample_nhwc; break; #ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC case DataType::F16: - _func = &NEUpsampleLayerKernel::upsample_f16_nhwc; + _func = &NEUpsampleLayerKernel::upsample_nhwc; break; #endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ default: diff --git a/tests/validation/NEON/Upsample.cpp b/tests/validation/NEON/Upsample.cpp index 9ddfbe0f8e..50aadeff05 100644 --- a/tests/validation/NEON/Upsample.cpp +++ b/tests/validation/NEON/Upsample.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -149,8 +149,20 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEUpsampleLayerQuantizedFixture, frame // Validate output validate(Accessor(_target), _reference); } -TEST_SUITE_END() // Quantized TEST_SUITE_END() // QASYMM8 +TEST_SUITE(QASYMM8_SIGNED) +FIXTURE_DATA_TEST_CASE(RunSmall, NEUpsampleLayerQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), + framework::dataset::make("DataLayout", { DataLayout::NCHW, DataLayout::NHWC })), + framework::dataset::make("PadInfo", { Size2D(2, 2) })), + framework::dataset::make("UpsamplingPolicy", { InterpolationPolicy::NEAREST_NEIGHBOR })), + framework::dataset::make("QuantizationInfo", QuantizationInfo(2.f / 255.f, 10)))) +{ + // Validate output + validate(Accessor(_target), _reference); +} +TEST_SUITE_END() // QASYMM8_SIGNED +TEST_SUITE_END() // Quantized TEST_SUITE_END() // UpsampleLayer TEST_SUITE_END() // NEON diff --git a/tests/validation/fixtures/UpsampleLayerFixture.h b/tests/validation/fixtures/UpsampleLayerFixture.h index 0a72e44dad..e46e19a635 100644 --- a/tests/validation/fixtures/UpsampleLayerFixture.h +++ b/tests/validation/fixtures/UpsampleLayerFixture.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -57,16 +57,7 @@ protected: template void fill(U &&tensor, int i) { - if(_data_type == DataType::QASYMM8) - { - const auto bounds = get_quantized_bounds(tensor.quantization_info(), -1.0f, 1.0f); - std::uniform_int_distribution distribution(bounds.first, bounds.second); - library->fill(tensor, distribution, i); - } - else - { - library->fill_tensor_uniform(tensor, i); - } + library->fill_tensor_uniform(tensor, i); } TensorType compute_target(TensorShape input_shape, const Size2D &info, const InterpolationPolicy &policy, -- cgit v1.2.1