diff options
| -rw-r--r-- | arm_compute/core/NEON/kernels/NEDepthConvertKernel.h | 33 | ||||
| -rw-r--r-- | arm_compute/runtime/NEON/functions/NEDepthConvert.h | 21 | ||||
| -rw-r--r-- | src/core/NEON/kernels/NEDepthConvertKernel.cpp | 192 | ||||
| -rw-r--r-- | src/runtime/NEON/functions/NEDepthConvert.cpp | 2 | ||||
| -rw-r--r-- | tests/validation/CL/DepthConvert.cpp | 4 | ||||
| -rw-r--r-- | tests/validation/FixedPoint.h | 9 | ||||
| -rw-r--r-- | tests/validation/NEON/DepthConvert.cpp | 201 | ||||
| -rw-r--r-- | tests/validation/Reference.cpp | 7 | ||||
| -rw-r--r-- | tests/validation/Reference.h | 16 | ||||
| -rw-r--r-- | tests/validation/TensorOperations.h | 26 |
10 files changed, 342 insertions, 169 deletions
diff --git a/arm_compute/core/NEON/kernels/NEDepthConvertKernel.h b/arm_compute/core/NEON/kernels/NEDepthConvertKernel.h index 0c5c29e4db..ad8d152bbf 100644 --- a/arm_compute/core/NEON/kernels/NEDepthConvertKernel.h +++ b/arm_compute/core/NEON/kernels/NEDepthConvertKernel.h @@ -24,7 +24,7 @@ #ifndef __ARM_COMPUTE_DEPTHCONVERTKERNEL_H__ #define __ARM_COMPUTE_DEPTHCONVERTKERNEL_H__ -#include "arm_compute/core/NEON/INESimpleKernel.h" +#include "arm_compute/core/NEON/INEKernel.h" #include "arm_compute/core/Types.h" #include <cstdint> @@ -34,35 +34,52 @@ namespace arm_compute class ITensor; /** Depth conversion kernel */ -class NEDepthConvertKernel : public INESimpleKernel +class NEDepthConvertKernel : public INEKernel { public: /** Default constructor*/ NEDepthConvertKernel(); + /** Prevent instances of this class from being copied (As this class contains pointers) */ + NEDepthConvertKernel(const NEDepthConvertKernel &) = delete; + /** Default move constructor */ + NEDepthConvertKernel(NEDepthConvertKernel &&) = default; + /** Prevent instances of this class from being copied (As this class contains pointers) */ + NEDepthConvertKernel &operator=(const NEDepthConvertKernel &) = delete; + /** Default move assignment operator */ + NEDepthConvertKernel &operator=(NEDepthConvertKernel &&) = default; /** Set the input and output of the kernel * * Valid conversions Input -> Output : * - * - QS8 -> F32 + * - QS8 -> QS8, F32 * - U8 -> U16, S16, S32 * - U16 -> U8, U32 * - S16 -> U8, S32 + * - QS16 -> QS16, F32 * - F32 -> QS8 * + * @warning In case of in-place fixed point position conversion make sure that configure has been called + * before the updated tensor is used in other functions, as the TensorInfo of the tensor will be + * altered. In-place is only supported for QS8 -> QS8, QS16 -> QS16. * - * @param[in] input The input tensor to convert. Data types supported: U8/QS8/U16/S16/F32. - * @param[out] output The output tensor. Data types supported: U8/QS8/U16/S16/U32/S32/F32. - * @param[in] policy Conversion policy. - * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8. + * @param[in, out] input The input tensor to convert (Written in case of in-place computation). Data types supported: U8/QS8/U16/S16/F32. + * @param[out] output The output tensor. Can be null in case of in-place computation. Data types supported: U8/QS8/U16/S16/U32/S32/F32. + * @param[in] policy Conversion policy. + * @param[in] shift (Optional) Value for down/up conversions. Must be 0 <= shift < 8. + * In case of fixed point position conversion, it specifies the new fixed point position, if operation is in-place. */ - void configure(const ITensor *input, ITensor *output, ConvertPolicy policy, uint32_t shift); + void configure(ITensor *input, ITensor *output, ConvertPolicy policy, uint32_t shift = 0); // Inherited methods overridden: void run(const Window &window) override; private: + ITensor *_input; + ITensor *_output; ConvertPolicy _policy; uint32_t _shift; + int _fixed_point_position_input; + int _fixed_point_position_output; }; } #endif /*__ARM_COMPUTE_NEDEPTHCONVERTKERNEL_H__ */ diff --git a/arm_compute/runtime/NEON/functions/NEDepthConvert.h b/arm_compute/runtime/NEON/functions/NEDepthConvert.h index 47b3a7e6f6..37f7293fb3 100644 --- a/arm_compute/runtime/NEON/functions/NEDepthConvert.h +++ b/arm_compute/runtime/NEON/functions/NEDepthConvert.h @@ -45,24 +45,25 @@ public: const NEDepthConvert &operator=(const NEDepthConvert &) = delete; /** Initialize the function's source, destination * - * Input format must be different than output format. - * * Valid conversions Input -> Output : - * QS8 -> F32 + * QS8 -> QS8, F32 * U8 -> U16, S16, S32 * U16 -> U8, U32 * S16 -> U8, S32 - * QS16 -> F32 + * QS16 -> QS16, F32 * F32 -> QS8, QS16 * + * @warning In case of in-place fixed point position conversion make sure that configure has been called + * before the updated tensor is used in other functions, as the TensorInfo of the tensor will be + * altered. In-place is only supported for QS8 -> QS8, QS16 -> QS16. * - * @param[in] input The input tensor to convert. Data type supported: QS8/U8/U16/S16/QS16/F32. - * @param[out] output The output tensor. Data type supported: QS8/U8/U16/S16/QS16/U32/S32/F32. - * @param[in] policy Conversion policy. - * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8. - * It is not used on fixed point conversion. + * @param[in, out] input The input tensor to convert (Written in case of in-place computation). Data types supported: U8/QS8/U16/S16/F32. + * @param[out] output The output tensor. Can be null in case of in-place computation. Data types supported: U8/QS8/U16/S16/U32/S32/F32. + * @param[in] policy Conversion policy. + * @param[in] shift (Optional) Value for down/up conversions. Must be 0 <= shift < 8. + * In case of fixed point position conversion, it specifies the new fixed point position, if operation is in-place. */ - void configure(const ITensor *input, ITensor *output, ConvertPolicy policy, uint32_t shift); + void configure(ITensor *input, ITensor *output, ConvertPolicy policy, uint32_t shift = 0); }; } #endif /*__ARM_COMPUTE_NEDEPTHCONVERT_H__*/ diff --git a/src/core/NEON/kernels/NEDepthConvertKernel.cpp b/src/core/NEON/kernels/NEDepthConvertKernel.cpp index 3c1a94df74..f7203701e7 100644 --- a/src/core/NEON/kernels/NEDepthConvertKernel.cpp +++ b/src/core/NEON/kernels/NEDepthConvertKernel.cpp @@ -40,24 +40,53 @@ class Coordinates; } // namespace arm_compute NEDepthConvertKernel::NEDepthConvertKernel() - : _policy(), _shift(0) + : _input(nullptr), _output(nullptr), _policy(), _shift(0), _fixed_point_position_input(0), _fixed_point_position_output(0) { } -void NEDepthConvertKernel::configure(const ITensor *input, ITensor *output, ConvertPolicy policy, uint32_t shift) +void NEDepthConvertKernel::configure(ITensor *input, ITensor *output, ConvertPolicy policy, uint32_t shift) { ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::QS8, DataType::S16, DataType::U16, DataType::QS16, DataType::F32); - ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8, DataType::QS8, DataType::S16, DataType::U16, DataType::QS16, DataType::U32, DataType::S32, DataType::F32); - ARM_COMPUTE_ERROR_ON(shift >= 8); - ARM_COMPUTE_ERROR_ON(input == output); - ARM_COMPUTE_ERROR_ON_MSG(input->info()->data_type() == output->info()->data_type(), "Input and output data_types must be different"); + + _input = input; + _output = input; + _policy = policy; + _shift = shift; + + if(output != nullptr) + { + // Auto initialize output shape if not initialized (We can only auto-configure the shape, datatype must be given) + set_shape_if_empty(*output->info(), input->info()->tensor_shape()); + + ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8, DataType::QS8, DataType::S16, DataType::U16, DataType::QS16, DataType::U32, DataType::S32, DataType::F32); + ARM_COMPUTE_ERROR_ON_MISMATCHING_SHAPES(input, output); + + // Set output + _output = output; + } + + // Set initial fixed point position of input and output + _fixed_point_position_input = input->info()->fixed_point_position(); + _fixed_point_position_output = _output->info()->fixed_point_position(); + + // Set the fixed point position to the output tensor if needed + if(is_data_type_fixed_point(input->info()->data_type()) && is_data_type_fixed_point(_output->info()->data_type())) + { + // If in-place set the fixed point position of the output tensor to be equal to shift + _fixed_point_position_output = (_input == _output) ? static_cast<int>(_shift) : _fixed_point_position_output; + // Set fixed point position to output tensor + _output->info()->set_fixed_point_position(_fixed_point_position_output); + } + + ARM_COMPUTE_ERROR_ON(shift >= 8 && (!is_data_type_fixed_point(input->info()->data_type()) && !is_data_type_fixed_point(output->info()->data_type()))); + ARM_COMPUTE_ERROR_ON(input == output && (data_size_from_type(input->info()->data_type()) != data_size_from_type(output->info()->data_type()))); ARM_COMPUTE_ERROR_ON_MSG(input->info()->data_type() == DataType::U8 && (output->info()->data_type() != DataType::S16 && output->info()->data_type() != DataType::U16 && output->info()->data_type() != DataType::S32), "Only data_types supported [in] U8 -> [out] U16, S16, S32"); - ARM_COMPUTE_ERROR_ON_MSG(input->info()->data_type() == DataType::QS8 && output->info()->data_type() != DataType::F32, - "Only data_types supported [in] QS8 -> [out] F32"); + ARM_COMPUTE_ERROR_ON_MSG(input->info()->data_type() == DataType::QS8 && (output->info()->data_type() != DataType::QS8 && output->info()->data_type() != DataType::F32), + "Only data_types supported [in] QS8 -> [out] QS8, F32"); ARM_COMPUTE_ERROR_ON_MSG(input->info()->data_type() == DataType::U16 && (output->info()->data_type() != DataType::U8 && output->info()->data_type() != DataType::U32), "Only data_types supported [in] U16 -> [out] U8, U32"); @@ -65,28 +94,36 @@ void NEDepthConvertKernel::configure(const ITensor *input, ITensor *output, Conv ARM_COMPUTE_ERROR_ON_MSG(input->info()->data_type() == DataType::S16 && (output->info()->data_type() != DataType::U8 && output->info()->data_type() != DataType::S32), "Only data_types supported [in] S16 -> [out] U8, S32"); - ARM_COMPUTE_ERROR_ON_MSG(input->info()->data_type() == DataType::QS16 && output->info()->data_type() != DataType::F32, - "Only data_types supported [in] QS16 -> [out] F32"); + ARM_COMPUTE_ERROR_ON_MSG(input->info()->data_type() == DataType::QS16 && (output->info()->data_type() != DataType::QS16 && output->info()->data_type() != DataType::F32), + "Only data_types supported [in] QS16 -> [out] QS16, F32"); ARM_COMPUTE_ERROR_ON_MSG(input->info()->data_type() == DataType::F32 && (output->info()->data_type() != DataType::QS8 && output->info()->data_type() != DataType::QS16), "Only data_types supported [in] F32 -> [out] QS8, QS16"); - // Auto initialize output shape if not initialized (We can only auto-configure the shape, datatype must be given) - set_shape_if_empty(*output->info(), input->info()->tensor_shape()); - - ARM_COMPUTE_ERROR_ON_MISMATCHING_SHAPES(input, output); + constexpr unsigned int num_elems_processed_per_iteration = 16; - _policy = policy; - _shift = shift; + // Configure kernel window + Window win = calculate_max_window(*input->info(), Steps(num_elems_processed_per_iteration)); - constexpr unsigned int num_elems_processed_per_iteration = 16; - INESimpleKernel::configure(input, output, num_elems_processed_per_iteration); + AccessWindowHorizontal input_access(input->info(), 0, num_elems_processed_per_iteration); + if(output != nullptr) + { + 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, input->info()->valid_region()); + } + else + { + // In-place computation + update_window_and_padding(win, input_access); + } + ICPPKernel::configure(win); } void NEDepthConvertKernel::run(const Window &window) { ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this); - ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INESimpleKernel::window(), window); + ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window); ARM_COMPUTE_ERROR_ON(nullptr == _input); ARM_COMPUTE_ERROR_ON(nullptr == _output); ARM_COMPUTE_ERROR_ON(_input == _output); @@ -94,37 +131,10 @@ void NEDepthConvertKernel::run(const Window &window) Iterator input(_input, window); Iterator output(_output, window); + bool in_place = (_input == _output); + switch(_input->info()->data_type()) { - case DataType::QS8: - { - const int fixed_point_position = _input->info()->fixed_point_position(); - - switch(_output->info()->data_type()) - { - case DataType::F32: - { - /* Up-conversion QS8 -> F32 */ - execute_window_loop(window, [&](const Coordinates & id) - { - const int8x16_t texels_s8 = vld1q_s8(reinterpret_cast<const int8_t *>(input.ptr())); - - float32x4x2_t texels_low = vcvt_f32_qs8(vget_low_s8(texels_s8), fixed_point_position); - float32x4x2_t texels_high = vcvt_f32_qs8(vget_high_s8(texels_s8), fixed_point_position); - - vst1q_f32(reinterpret_cast<float *>(output.ptr()), texels_low.val[0]); - vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 4, texels_low.val[1]); - vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 8, texels_high.val[0]); - vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 12, texels_high.val[1]); - }, - input, output); - break; - } - default: - ARM_COMPUTE_ERROR("Output data type not supported"); - } - break; - } case DataType::U8: { const int16x8_t b = vdupq_n_s16(_shift); @@ -201,6 +211,49 @@ void NEDepthConvertKernel::run(const Window &window) } break; } + case DataType::QS8: + { + switch(_output->info()->data_type()) + { + case DataType::QS8: + { + const int relative_shift = _fixed_point_position_output - _fixed_point_position_input; + /* Fixed point position conversion QS8 -> QS8 */ + if(relative_shift != 0 || !in_place) + { + const auto relative_shift_vec = vdupq_n_qs8(relative_shift); + execute_window_loop(window, [&](const Coordinates & id) + { + const qint8x16_t texels_qs8 = vld1q_qs8(reinterpret_cast<const qint8_t *>(input.ptr())); + vst1q_qs8(reinterpret_cast<qint8_t *>(output.ptr()), vqrshlq_s8(texels_qs8, relative_shift_vec)); + }, + input, output); + } + break; + } + case DataType::F32: + { + /* Up-conversion QS8 -> F32 */ + execute_window_loop(window, [&](const Coordinates & id) + { + const qint8x16_t texels_qs8 = vld1q_qs8(reinterpret_cast<const qint8_t *>(input.ptr())); + + float32x4x2_t texels_low = vcvt_f32_qs8(vget_low_s8(texels_qs8), _fixed_point_position_input); + float32x4x2_t texels_high = vcvt_f32_qs8(vget_high_s8(texels_qs8), _fixed_point_position_input); + + vst1q_f32(reinterpret_cast<float *>(output.ptr()), texels_low.val[0]); + vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 4, texels_low.val[1]); + vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 8, texels_high.val[0]); + vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 12, texels_high.val[1]); + }, + input, output); + break; + } + default: + ARM_COMPUTE_ERROR("Output data type not supported"); + } + break; + } case DataType::S16: { switch(_output->info()->data_type()) @@ -356,16 +409,37 @@ void NEDepthConvertKernel::run(const Window &window) } case DataType::QS16: { - const int fixed_point_position = _input->info()->fixed_point_position(); - switch(_output->info()->data_type()) { + case DataType::QS16: + { + const int relative_shift = _fixed_point_position_output - _fixed_point_position_input; + /* Fixed point position conversion QS16 -> QS16 */ + if(relative_shift != 0 || !in_place) + { + const auto relative_shift_vec = vdupq_n_qs16(relative_shift); + execute_window_loop(window, [&](const Coordinates & id) + { + const qint16x8x2_t texels_qs16 = + { + { + vld1q_qs16(reinterpret_cast<qint16_t *>(input.ptr())), + vld1q_qs16(reinterpret_cast<qint16_t *>(input.ptr()) + 8) + } + }; + vst1q_qs16(reinterpret_cast<qint16_t *>(output.ptr()), vqrshlq_s16(texels_qs16.val[0], relative_shift_vec)); + vst1q_qs16(reinterpret_cast<qint16_t *>(output.ptr()) + 8, vqrshlq_s16(texels_qs16.val[1], relative_shift_vec)); + }, + input, output); + } + break; + } case DataType::F32: { /* Up-conversion QS16 -> F32 */ execute_window_loop(window, [&](const Coordinates & id) { - const int16x8x2_t texels = + const int16x8x2_t texels_qs16 = { { vld1q_s16(reinterpret_cast<qint16_t *>(input.ptr())), @@ -373,10 +447,10 @@ void NEDepthConvertKernel::run(const Window &window) } }; - vst1q_f32(reinterpret_cast<float *>(output.ptr()), vcvt_f32_qs16(vget_low_s16(texels.val[0]), fixed_point_position)); - vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 4, vcvt_f32_qs16(vget_high_s16(texels.val[0]), fixed_point_position)); - vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 8, vcvt_f32_qs16(vget_low_s16(texels.val[1]), fixed_point_position)); - vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 12, vcvt_f32_qs16(vget_high_s16(texels.val[1]), fixed_point_position)); + vst1q_f32(reinterpret_cast<float *>(output.ptr()), vcvt_f32_qs16(vget_low_s16(texels_qs16.val[0]), _fixed_point_position_input)); + vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 4, vcvt_f32_qs16(vget_high_s16(texels_qs16.val[0]), _fixed_point_position_input)); + vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 8, vcvt_f32_qs16(vget_low_s16(texels_qs16.val[1]), _fixed_point_position_input)); + vst1q_f32(reinterpret_cast<float *>(output.ptr()) + 12, vcvt_f32_qs16(vget_high_s16(texels_qs16.val[1]), _fixed_point_position_input)); }, input, output); break; @@ -392,7 +466,6 @@ void NEDepthConvertKernel::run(const Window &window) { case DataType::QS8: { - const int fixed_point_position = _output->info()->fixed_point_position(); /* Down-conversion F32 -> QS8 */ execute_window_loop(window, [&](const Coordinates & id) { @@ -406,7 +479,7 @@ void NEDepthConvertKernel::run(const Window &window) } }; - const qint8x16_t texels_s8 = vqcvtq_qs8_f32(texels_f32, fixed_point_position); + const qint8x16_t texels_s8 = vqcvtq_qs8_f32(texels_f32, _fixed_point_position_output); vst1q_s8(reinterpret_cast<int8_t *>(output.ptr()), texels_s8); }, @@ -415,7 +488,6 @@ void NEDepthConvertKernel::run(const Window &window) } case DataType::QS16: { - const int fixed_point_position = _output->info()->fixed_point_position(); /* Down-conversion F32 -> QS16 */ execute_window_loop(window, [&](const Coordinates & id) { @@ -434,8 +506,8 @@ void NEDepthConvertKernel::run(const Window &window) } }; - vst1q_s16(reinterpret_cast<qint16_t *>(output.ptr()), vqcvtq_qs16_f32(texels_f32_1, fixed_point_position)); - vst1q_s16(reinterpret_cast<qint16_t *>(output.ptr()) + 8, vqcvtq_qs16_f32(texels_f32_2, fixed_point_position)); + vst1q_s16(reinterpret_cast<qint16_t *>(output.ptr()), vqcvtq_qs16_f32(texels_f32_1, _fixed_point_position_output)); + vst1q_s16(reinterpret_cast<qint16_t *>(output.ptr()) + 8, vqcvtq_qs16_f32(texels_f32_2, _fixed_point_position_output)); }, input, output); break; diff --git a/src/runtime/NEON/functions/NEDepthConvert.cpp b/src/runtime/NEON/functions/NEDepthConvert.cpp index 24b51493c6..37857b6534 100644 --- a/src/runtime/NEON/functions/NEDepthConvert.cpp +++ b/src/runtime/NEON/functions/NEDepthConvert.cpp @@ -30,7 +30,7 @@ using namespace arm_compute; -void NEDepthConvert::configure(const ITensor *input, ITensor *output, ConvertPolicy policy, uint32_t shift) +void NEDepthConvert::configure(ITensor *input, ITensor *output, ConvertPolicy policy, uint32_t shift) { auto k = arm_compute::support::cpp14::make_unique<NEDepthConvertKernel>(); k->configure(input, output, policy, shift); diff --git a/tests/validation/CL/DepthConvert.cpp b/tests/validation/CL/DepthConvert.cpp index 2655f0024a..547820e5d7 100644 --- a/tests/validation/CL/DepthConvert.cpp +++ b/tests/validation/CL/DepthConvert.cpp @@ -469,7 +469,7 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Da CLTensor dst = compute_depth_convert(shape, DataType::F32, dt, policy, 0, fixed_point_position); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, dt, policy, 0, fixed_point_position); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, dt, policy, 0, fixed_point_position, fixed_point_position); // Validate output validate(CLAccessor(dst), ref_dst); @@ -484,7 +484,7 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Da CLTensor dst = compute_depth_convert(shape, DataType::F32, dt, policy, 0, fixed_point_position); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, dt, policy, 0, fixed_point_position); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, dt, policy, 0, fixed_point_position, fixed_point_position); // Validate output validate(CLAccessor(dst), ref_dst); diff --git a/tests/validation/FixedPoint.h b/tests/validation/FixedPoint.h index ab6f14a49b..12ffcdfc3d 100644 --- a/tests/validation/FixedPoint.h +++ b/tests/validation/FixedPoint.h @@ -244,15 +244,20 @@ public: { assert(p > 0 && p < std::numeric_limits<T>::digits); + using promoted_T = typename traits::promote<T>::type; + promoted_T val = _value; if(p > _fixed_point_position) { - _value <<= (p - _fixed_point_position); + val <<= (p - _fixed_point_position); } else if(p < _fixed_point_position) { - _value >>= (_fixed_point_position - p); + uint8_t pbar = _fixed_point_position - p; + val += (pbar != 0) ? (1 << (pbar - 1)) : 0; + val >>= pbar; } + _value = detail::constant_expr<T>::saturate_cast(val); _fixed_point_position = p; } diff --git a/tests/validation/NEON/DepthConvert.cpp b/tests/validation/NEON/DepthConvert.cpp index 65d3ab1be7..8a30c746e9 100644 --- a/tests/validation/NEON/DepthConvert.cpp +++ b/tests/validation/NEON/DepthConvert.cpp @@ -51,20 +51,22 @@ namespace { /** Compute Neon depth convert function. * - * @param[in] shape Shape of the input and output tensors. - * @param[in] dt_in Data type of input tensor. - * @param[in] dt_out Data type of the output tensor. - * @param[in] policy Conversion policy. - * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8. - * @param[in] fixed_point_position Fixed point position. + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in Data type of input tensor. + * @param[in] dt_out Data type of the output tensor. + * @param[in] policy Conversion policy. + * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8. + * @param[in] fixed_point_position_in (Optional) Fixed point position for the input tensor. + * @param[in] fixed_point_position_out (Optional) Fixed point position for the output tensor. * * @return Computed output tensor. */ -Tensor compute_depth_convert(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, uint32_t shift, uint32_t fixed_point_position) +Tensor compute_depth_convert(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, + uint32_t shift, uint32_t fixed_point_position_in = 0, uint32_t fixed_point_position_out = 0) { // Create tensors - Tensor src = create_tensor<Tensor>(shape, dt_in, 1, fixed_point_position); - Tensor dst = create_tensor<Tensor>(shape, dt_out, 1, fixed_point_position); + Tensor src = create_tensor<Tensor>(shape, dt_in, 1, fixed_point_position_in); + Tensor dst = create_tensor<Tensor>(shape, dt_out, 1, fixed_point_position_out); // Create and configure function NEDepthConvert depth_convert; @@ -87,20 +89,22 @@ Tensor compute_depth_convert(const TensorShape &shape, DataType dt_in, DataType } /** Configure and validate region/padding function. * - * @param[in] shape Shape of the input and output tensors. - * @param[in] dt_in Data type of input tensor. - * @param[in] dt_out Data type of the output tensor. - * @param[in] policy Conversion policy. - * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8. - * @param[in] fixed_point_position Fixed point position. + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in Data type of input tensor. + * @param[in] dt_out Data type of the output tensor. + * @param[in] policy Conversion policy. + * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8. + * @param[in] fixed_point_position_in (Optional) Fixed point position for the input tensor. + * @param[in] fixed_point_position_out (Optional) Fixed point position for the output tensor. * */ -void compute_configure_validate(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, uint32_t shift, uint32_t fixed_point_position) +void compute_configure_validate(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, + uint32_t shift, uint32_t fixed_point_position_in = 0, uint32_t fixed_point_position_out = 0) { // Create tensors - Tensor src = create_tensor<Tensor>(shape, dt_in, 1, fixed_point_position); - Tensor dst = create_tensor<Tensor>(shape, dt_out, 1, fixed_point_position); + Tensor src = create_tensor<Tensor>(shape, dt_in, 1, fixed_point_position_in); + Tensor dst = create_tensor<Tensor>(shape, dt_out, 1, fixed_point_position_out); BOOST_TEST(src.info()->is_resizable()); BOOST_TEST(dst.info()->is_resizable()); @@ -125,6 +129,32 @@ void compute_configure_validate(const TensorShape &shape, DataType dt_in, DataTy BOOST_AUTO_TEST_SUITE(NEON) BOOST_AUTO_TEST_SUITE(DepthConvert) +BOOST_AUTO_TEST_SUITE(QS8_to_QS8) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * (boost::unit_test::data::make({ 1, 3, 5, 6 }) ^ boost::unit_test::data::make({ 6, 5, 1, 3 })), + shape, policy, fixed_point_position_in, fixed_point_position_out) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::QS8, DataType::QS8, policy, 0, fixed_point_position_in, fixed_point_position_out); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * (boost::unit_test::data::make({ 1, 3, 5, 6 }) ^ boost::unit_test::data::make({ 6, 5, 1, 3 })), + shape, policy, fixed_point_position_in, fixed_point_position_out) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::QS8, policy, 0, fixed_point_position_in, fixed_point_position_out); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::QS8, policy, 0, fixed_point_position_in, fixed_point_position_out); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + BOOST_AUTO_TEST_SUITE(QS8_to_F32) BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) @@ -132,7 +162,7 @@ BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::uni shape, policy, fixed_point_position) { // Compute configure and validate region/padding - compute_configure_validate(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position); + compute_configure_validate(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); } BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) @@ -141,10 +171,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, fixed_point_position) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position); + Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); // Validate output validate(NEAccessor(dst), ref_dst); @@ -156,15 +186,14 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, fixed_point_position) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position); + Tensor dst = compute_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS8, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); // Validate output validate(NEAccessor(dst), ref_dst); } - BOOST_AUTO_TEST_SUITE_END() BOOST_AUTO_TEST_SUITE(F32_to_QS8) @@ -174,7 +203,7 @@ BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::uni shape, policy, fixed_point_position) { // Compute configure and validate region/padding - compute_configure_validate(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position); + compute_configure_validate(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position); } BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) @@ -183,10 +212,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, fixed_point_position) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position); + Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position); // Validate output validate(NEAccessor(dst), ref_dst); @@ -198,10 +227,36 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, fixed_point_position) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position); + Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position); + + // Compute reference + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position, fixed_point_position); + + // Validate output + validate(NEAccessor(dst), ref_dst); +} +BOOST_AUTO_TEST_SUITE_END() + +BOOST_AUTO_TEST_SUITE(QS16_to_QS16) +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit") * boost::unit_test::label("nightly")) +BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * (boost::unit_test::data::make({ 3, 6, 7, 13, 14 }) ^ boost::unit_test::data::make({ 5, 10, 14, 4, 7 })), + shape, policy, fixed_point_position_in, fixed_point_position_out) +{ + // Compute configure and validate region/padding + compute_configure_validate(shape, DataType::QS16, DataType::QS16, policy, 0, fixed_point_position_in, fixed_point_position_out); +} + +BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) +BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ ConvertPolicy::SATURATE }) + * (boost::unit_test::data::make({ 3, 6, 7, 13, 14 }) ^ boost::unit_test::data::make({ 5, 10, 14, 4, 7 })), + shape, policy, fixed_point_position_in, fixed_point_position_out) +{ + // Compute function + Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::QS16, policy, 0, fixed_point_position_in, fixed_point_position_out); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS8, policy, 0, fixed_point_position); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::QS16, policy, 0, fixed_point_position_in, fixed_point_position_out); // Validate output validate(NEAccessor(dst), ref_dst); @@ -215,7 +270,7 @@ BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::uni shape, policy, fixed_point_position) { // Compute configure and validate region/padding - compute_configure_validate(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position); + compute_configure_validate(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); } BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) @@ -224,10 +279,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, fixed_point_position) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position); + Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); // Validate output validate(NEAccessor(dst), ref_dst); @@ -239,10 +294,10 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, fixed_point_position) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position); + Tensor dst = compute_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::QS16, DataType::F32, policy, 0, fixed_point_position, fixed_point_position); // Validate output validate(NEAccessor(dst), ref_dst); @@ -257,7 +312,7 @@ BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::uni shape, policy, fixed_point_position) { // Compute configure and validate region/padding - compute_configure_validate(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position); + compute_configure_validate(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position); } BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) @@ -266,10 +321,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, fixed_point_position) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position); + Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position); // Validate output validate(NEAccessor(dst), ref_dst); @@ -281,10 +336,10 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, fixed_point_position) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position); + Tensor dst = compute_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::F32, DataType::QS16, policy, 0, fixed_point_position, fixed_point_position); // Validate output validate(NEAccessor(dst), ref_dst); @@ -308,10 +363,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::U16, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::U16, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -322,10 +377,10 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::U16, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::U16, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::U16, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -339,7 +394,7 @@ BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::uni shape, policy, shift) { // Compute configure and validate region/padding - compute_configure_validate(shape, DataType::U8, DataType::S16, policy, shift, 0); + compute_configure_validate(shape, DataType::U8, DataType::S16, policy, shift); } BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) @@ -348,10 +403,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S16, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S16, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -363,10 +418,10 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S16, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S16, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S16, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -380,7 +435,7 @@ BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::uni shape, policy, shift) { // Compute configure and validate region/padding - compute_configure_validate(shape, DataType::U8, DataType::S32, policy, shift, 0); + compute_configure_validate(shape, DataType::U8, DataType::S32, policy, shift); } BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) @@ -389,10 +444,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S32, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S32, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -404,10 +459,10 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S32, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::U8, DataType::S32, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U8, DataType::S32, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -421,7 +476,7 @@ BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::uni shape, policy, shift) { // Compute configure and validate region/padding - compute_configure_validate(shape, DataType::U16, DataType::U8, policy, shift, 0); + compute_configure_validate(shape, DataType::U16, DataType::U8, policy, shift); } BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) @@ -430,10 +485,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U8, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U8, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -445,10 +500,10 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U8, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U8, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U8, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -462,7 +517,7 @@ BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::uni shape, policy, shift) { // Compute configure and validate region/padding - compute_configure_validate(shape, DataType::U16, DataType::U32, policy, shift, 0); + compute_configure_validate(shape, DataType::U16, DataType::U32, policy, shift); } BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) @@ -471,10 +526,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U32, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U32, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -486,10 +541,10 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U32, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::U16, DataType::U32, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::U16, DataType::U32, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -503,7 +558,7 @@ BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::uni shape, policy, shift) { // Compute configure and validate region/padding - compute_configure_validate(shape, DataType::S16, DataType::U8, policy, shift, 0); + compute_configure_validate(shape, DataType::S16, DataType::U8, policy, shift); } BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) @@ -512,10 +567,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::U8, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::U8, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -527,10 +582,10 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::U8, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::U8, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::U8, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -544,7 +599,7 @@ BOOST_DATA_TEST_CASE(Configuration, (SmallShapes() + LargeShapes()) * boost::uni shape, policy, shift) { // Compute configure and validate region/padding - compute_configure_validate(shape, DataType::S16, DataType::S32, policy, shift, 0); + compute_configure_validate(shape, DataType::S16, DataType::S32, policy, shift); } BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit")) @@ -553,10 +608,10 @@ BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::S32, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::S32, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); @@ -568,10 +623,10 @@ BOOST_DATA_TEST_CASE(RunLarge, LargeShapes() * boost::unit_test::data::make({ Co shape, policy, shift) { // Compute function - Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::S32, policy, shift, 0); + Tensor dst = compute_depth_convert(shape, DataType::S16, DataType::S32, policy, shift); // Compute reference - RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift, 0); + RawTensor ref_dst = Reference::compute_reference_depth_convert(shape, DataType::S16, DataType::S32, policy, shift); // Validate output validate(NEAccessor(dst), ref_dst); diff --git a/tests/validation/Reference.cpp b/tests/validation/Reference.cpp index acf010e7b7..4b4919fd28 100644 --- a/tests/validation/Reference.cpp +++ b/tests/validation/Reference.cpp @@ -283,10 +283,11 @@ RawTensor Reference::compute_reference_box3x3(const TensorShape &shape, BorderMo return ref_dst; } -RawTensor Reference::compute_reference_depth_convert(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, uint32_t shift, uint32_t fixed_point_position) +RawTensor Reference::compute_reference_depth_convert(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, + uint32_t shift, uint32_t fixed_point_position_in, uint32_t fixed_point_position_out) { - RawTensor ref_src = library->get(shape, dt_in, 1, fixed_point_position); - RawTensor ref_dst = library->get(shape, dt_out, 1, fixed_point_position); + RawTensor ref_src = library->get(shape, dt_in, 1, fixed_point_position_in); + RawTensor ref_dst = library->get(shape, dt_out, 1, fixed_point_position_out); // Fill reference library->fill_tensor_uniform(ref_src, 0); diff --git a/tests/validation/Reference.h b/tests/validation/Reference.h index 3ad9814439..a3ae3b6d02 100644 --- a/tests/validation/Reference.h +++ b/tests/validation/Reference.h @@ -168,16 +168,18 @@ public: static RawTensor compute_reference_box3x3(const TensorShape &shape, BorderMode border_mode, uint8_t constant_border_value); /** Compute reference depth convert. * - * @param[in] shape Shape of the input and output tensors. - * @param[in] dt_in Data type of input tensor. - * @param[in] dt_out Data type of the output tensor. - * @param[in] policy Overflow policy of the operation. - * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8. - * @param[in] fixed_point_position Fixed point position. + * @param[in] shape Shape of the input and output tensors. + * @param[in] dt_in Data type of input tensor. + * @param[in] dt_out Data type of the output tensor. + * @param[in] policy Overflow policy of the operation. + * @param[in] shift Value for down/up conversions. Must be 0 <= shift < 8. + * @param[in] fixed_point_position_in (Optional) Fixed point position for the input tensor. + * @param[in] fixed_point_position_out (Optional) Fixed point position for the output tensor. * * @return Computed raw tensor. */ - static RawTensor compute_reference_depth_convert(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, uint32_t shift, uint32_t fixed_point_position); + static RawTensor compute_reference_depth_convert(const TensorShape &shape, DataType dt_in, DataType dt_out, ConvertPolicy policy, + uint32_t shift, uint32_t fixed_point_position_in = 0, uint32_t fixed_point_position_out = 0); /** Compute reference gaussian3x3 filter. * * @param[in] shape Shape of the input and output tensors. diff --git a/tests/validation/TensorOperations.h b/tests/validation/TensorOperations.h index 9e6f5cf5d1..882c9e07e1 100644 --- a/tests/validation/TensorOperations.h +++ b/tests/validation/TensorOperations.h @@ -519,7 +519,7 @@ void box3x3(const Tensor<T> &in, Tensor<T> &out, BorderMode border_mode, T const } // Depth conversion -template < typename T1, typename T2, typename std::enable_if < std::is_integral<T1>::value &&is_floating_point<T2>::value, int >::type = 0 > +template < typename T1, typename T2, typename std::enable_if < std::is_integral<T1>::value &&std::is_floating_point<T2>::value, int >::type = 0 > void depth_convert(const Tensor<T1> &in, Tensor<T2> &out, ConvertPolicy policy, uint32_t shift) { using namespace fixed_point_arithmetic; @@ -531,7 +531,7 @@ void depth_convert(const Tensor<T1> &in, Tensor<T2> &out, ConvertPolicy policy, } } -template < typename T1, typename T2, typename std::enable_if < is_floating_point<T1>::value &&std::is_integral<T2>::value, int >::type = 0 > +template < typename T1, typename T2, typename std::enable_if < std::is_floating_point<T1>::value &&std::is_integral<T2>::value, int >::type = 0 > void depth_convert(const Tensor<T1> &in, Tensor<T2> &out, ConvertPolicy policy, uint32_t shift) { using namespace fixed_point_arithmetic; @@ -543,7 +543,7 @@ void depth_convert(const Tensor<T1> &in, Tensor<T2> &out, ConvertPolicy policy, } } -template < typename T1, typename T2, typename std::enable_if < std::is_integral<T1>::value &&std::is_integral<T2>::value, int >::type = 0 > +template < typename T1, typename T2, typename std::enable_if < std::is_integral<T1>::value &&std::is_integral<T2>::value &&!std::is_same<T1, T2>::value, int >::type = 0 > void depth_convert(const Tensor<T1> &in, Tensor<T2> &out, ConvertPolicy policy, uint32_t shift) { // Up-casting @@ -565,6 +565,26 @@ void depth_convert(const Tensor<T1> &in, Tensor<T2> &out, ConvertPolicy policy, } } +template < typename T1, typename T2, typename std::enable_if < std::is_integral<T1>::value &&std::is_integral<T2>::value &&std::is_same<T1, T2>::value, int >::type = 0 > +void depth_convert(const Tensor<T1> &in, Tensor<T2> &out, ConvertPolicy policy, uint32_t shift) +{ + using namespace fixed_point_arithmetic; + bool is_in_place = (&in == &out); + + const int fixed_point_position_in = in.fixed_point_position(); + const int fixed_point_position_out = (is_in_place) ? static_cast<int>(shift) : out.fixed_point_position(); + + if(!is_in_place || (fixed_point_position_in != fixed_point_position_out)) + { + for(int i = 0; i < in.num_elements(); ++i) + { + auto x = fixed_point<T2>(in[i], fixed_point_position_in, true); + x.rescale(fixed_point_position_out); + out[i] = x.raw(); + } + } +} + template < typename T1, typename T2, typename std::enable_if < is_floating_point<T1>::value &&is_floating_point<T2>::value, int >::type = 0 > void depth_convert(const Tensor<T1> &in, Tensor<T2> &out, ConvertPolicy policy, uint32_t shift) { |