From 9428a182911802cf6e6df6eb751a7c7eb43602f9 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Mon, 30 Mar 2020 14:10:20 +0100 Subject: COMPMID-3237: Add support for QSYMM16 into S32 NEPixelwiseMultiplicationKernel Change-Id: I8dc3348db37b041f442639ac0d072740ca639878 Signed-off-by: Michele Di Giorgio Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2960 Tested-by: Arm Jenkins Reviewed-by: Georgios Pinitas Reviewed-by: Sang-Hoon Park Comments-Addressed: Arm Jenkins --- .../NEON/kernels/NEPixelWiseMultiplicationKernel.h | 22 ++- .../NEON/functions/NEPixelWiseMultiplication.h | 22 ++- .../kernels/NEPixelWiseMultiplicationKernel.cpp | 175 ++++++++++++++++----- tests/validation/CL/PixelWiseMultiplication.cpp | 76 +++++---- tests/validation/NEON/PixelWiseMultiplication.cpp | 147 ++++++++--------- tests/validation/fixtures/LSTMLayerFixture.h | 22 +-- .../fixtures/PixelWiseMultiplicationFixture.h | 46 ++---- .../reference/PixelWiseMultiplication.cpp | 57 +++---- .../validation/reference/PixelWiseMultiplication.h | 9 +- .../reference/QLSTMLayerNormalization.cpp | 2 +- 10 files changed, 352 insertions(+), 226 deletions(-) diff --git a/arm_compute/core/NEON/kernels/NEPixelWiseMultiplicationKernel.h b/arm_compute/core/NEON/kernels/NEPixelWiseMultiplicationKernel.h index 9b71ac81cf..1a9dd6be2e 100644 --- a/arm_compute/core/NEON/kernels/NEPixelWiseMultiplicationKernel.h +++ b/arm_compute/core/NEON/kernels/NEPixelWiseMultiplicationKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -58,7 +58,15 @@ public: * * @param[in] input1 An input tensor. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 * @param[in] input2 An input tensor. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), QASYMM8_SIGNED (only if @p input1 is QASYMM8_SIGNED), S16, QSYMM16 (only if @p input1 is QSYMM16), F16 (only if @p input1 is F16), F32 (only if @p input1 is F32). - * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), QASYMM8_SIGNED (only if @p input1 is QASYMM8_SIGNED), S16, QSYMM16 (only if both inputs are QSYMM16), F16 (only if @p input1 is F16), F32 (only if both inputs are F32). + * @param[out] output Output tensor. Data types supported: + * - U8, only if both inputs are U8. + * - QASYMM8, only if both inputs are QASYMM8. + * - QASYMM8_SIGNED, only if @p input1 is QASYMM8_SIGNED. + * - S16. + * - QSYMM16, only if both inputs are QSYMM16. + * - S32, only if both inputs are QSYMM16. + * - F16, only if @p input1 is F16. + * - F32, only if both inputs are F32. * @param[in] scale Scale to apply after multiplication. * Scale must be positive and its value must be either 1/255 or 1/2^n where n is between 0 and 15. * @param[in] overflow_policy Overflow policy. ConvertPolicy cannot be WRAP if datatype is QASYMM8, QASYMM8_SIGNED or QSYMM16. @@ -72,7 +80,15 @@ public: * * @param[in] input1 An input tensor info. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/QSYMM16/S16/F16/F32 * @param[in] input2 An input tensor info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), QASYMM8_SIGNED (only if @p input1 is QASYMM8_SIGNED), S16, QSYMM16 (only if @p input1 is QSYMM16), F16 (only if @p input1 is F16), F32 (only if @p input1 is F32). - * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), QASYMM8_SIGNED (only if @p input1 is QASYMM8_SIGNED) , S16, QSYMM16 (only if both inputs are QSYMM16), F16 (only if @p input1 is F16), F32 (only if both inputs are F32). + * @param[in] output Output tensor info. Data types supported: + * - U8, only if both inputs are U8. + * - QASYMM8, only if both inputs are QASYMM8. + * - QASYMM8_SIGNED, only if @p input1 is QASYMM8_SIGNED. + * - S16. + * - QSYMM16, only if both inputs are QSYMM16. + * - S32, only if both inputs are QSYMM16. + * - F16, only if @p input1 is F16. + * - F32, only if both inputs are F32. * @param[in] scale Scale to apply after multiplication. * Scale must be positive and its value must be either 1/255 or 1/2^n where n is between 0 and 15. * @param[in] overflow_policy Overflow policy. ConvertPolicy cannot be WRAP if datatype is QASYMM8, QASYMM8_SIGNED or QSYMM16. diff --git a/arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h b/arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h index 25f409871b..ede4327bfb 100644 --- a/arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h +++ b/arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -44,7 +44,15 @@ public: * This input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. * @param[in, out] input2 An input tensor. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), QASYMM8_SIGNED (only if @p input1 is QASYMM8_SIGNED), S16, QSYMM16 (only if @p input1 is QSYMM16), F16 (only if @p input1 is F16), F32 (only if @p input1 is F32). * This input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), QASYMM8_SIGNED (only if @p input1 is QASYMM8_SIGNED), S16, QSYMM16 (only if both inputs are QSYMM16), F16 (only if @p input1 is F16), F32 (only if both inputs are F32). + * @param[out] output Output tensor. Data types supported: + * - U8, only if both inputs are U8. + * - QASYMM8, only if both inputs are QASYMM8. + * - QASYMM8_SIGNED, only if @p input1 is QASYMM8_SIGNED. + * - S16. + * - QSYMM16, only if both inputs are QSYMM16. + * - S32, only if both inputs are QSYMM16. + * - F16, only if @p input1 is F16. + * - F32, only if both inputs are F32. * @param[in] scale Scale to apply after multiplication. * Scale must be positive and its value must be either 1/255 or 1/2^n where n is between 0 and 15. * @param[in] overflow_policy Overflow policy. ConvertPolicy cannot be WRAP if datatype is QASYMM8, QASYMM8_SIGNED or QSYMM16. @@ -58,7 +66,15 @@ public: * * @param[in] input1 An input tensor info. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 * @param[in] input2 An input tensor info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), QASYMM8_SIGNED (only if @p input1 is QASYMM8_SIGNED), S16, QSYMM16 (only if both inputs are QSYMM16), F16 (only if @p input1 is F16), F32 (only if @p input1 is F32). - * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), QASYMM8_SIGNED (only if @p input1 is QASYMM8_SIGNED), S16, QSYMM16 (only if both inputs are QSYMM16), F16 (only if @p input1 is F16), F32 (only if both inputs are F32). + * @param[in] output Output tensor info. Data types supported: + * - U8, only if both inputs are U8. + * - QASYMM8, only if both inputs are QASYMM8. + * - QASYMM8_SIGNED, only if @p input1 is QASYMM8_SIGNED. + * - S16. + * - QSYMM16, only if both inputs are QSYMM16. + * - S32, only if both inputs are QSYMM16. + * - F16, only if @p input1 is F16. + * - F32, only if both inputs are F32. * @param[in] scale Scale to apply after multiplication. * Scale must be positive and its value must be either 1/255 or 1/2^n where n is between 0 and 15. * @param[in] overflow_policy Overflow policy. ConvertPolicy cannot be WRAP if datatype is QASYMM8, QASYMM8_SIGNED or QSYMM16. diff --git a/src/core/NEON/kernels/NEPixelWiseMultiplicationKernel.cpp b/src/core/NEON/kernels/NEPixelWiseMultiplicationKernel.cpp index a87588dbb3..ca59e66293 100644 --- a/src/core/NEON/kernels/NEPixelWiseMultiplicationKernel.cpp +++ b/src/core/NEON/kernels/NEPixelWiseMultiplicationKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -24,23 +24,12 @@ #include "arm_compute/core/NEON/kernels/NEPixelWiseMultiplicationKernel.h" #include "arm_compute/core/CPP/Validate.h" -#include "arm_compute/core/Error.h" -#include "arm_compute/core/Helpers.h" -#include "arm_compute/core/IAccessWindow.h" -#include "arm_compute/core/ITensor.h" #include "arm_compute/core/NEON/NEAsymm.h" -#include "arm_compute/core/NEON/NEFixedPoint.h" #include "arm_compute/core/NEON/NESymm.h" #include "arm_compute/core/NEON/wrapper/wrapper.h" #include "arm_compute/core/TensorInfo.h" -#include "arm_compute/core/Types.h" -#include "arm_compute/core/Validate.h" #include -#include -#include -#include -#include #if __ARM_FEATURE_FP16_VECTOR_ARITHMETIC #include // needed for float16_t @@ -48,8 +37,6 @@ namespace arm_compute { -class Coordinates; - namespace { const float scale255_constant = 1.f / 255.f; @@ -66,10 +53,9 @@ inline Status validate_arguments(const ITensorInfo *input1, const ITensorInfo *i ARM_COMPUTE_RETURN_ERROR_ON_CPU_F16_UNSUPPORTED(input1); ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input1, 1, DataType::U8, DataType::QASYMM8, DataType::QASYMM8_SIGNED, DataType::S16, DataType::QSYMM16, DataType::F16, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input2, 1, DataType::U8, DataType::QASYMM8, DataType::QASYMM8_SIGNED, DataType::S16, DataType::QSYMM16, DataType::F16, DataType::F32); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8, DataType::QASYMM8, DataType::QASYMM8_SIGNED, DataType::S16, DataType::QSYMM16, DataType::F16, DataType::F32); - ARM_COMPUTE_RETURN_ERROR_ON_MSG(output->data_type() == DataType::U8 && (input1->data_type() != DataType::U8 || input2->data_type() != DataType::U8), - "Output can only be U8 if both inputs are U8"); - + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8, DataType::QASYMM8, DataType::QASYMM8_SIGNED, + DataType::S16, DataType::QSYMM16, + DataType::S32, DataType::F16, DataType::F32); if(is_data_type_quantized(input1->data_type()) || is_data_type_quantized(input2->data_type())) { ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input1, input2); @@ -86,6 +72,12 @@ inline Status validate_arguments(const ITensorInfo *input1, const ITensorInfo *i const TensorShape &out_shape = TensorShape::broadcast_shape(input1->tensor_shape(), input2->tensor_shape()); ARM_COMPUTE_RETURN_ERROR_ON_MSG(detail::have_different_dimensions(out_shape, output->tensor_shape(), 0), "Wrong shape for output"); ARM_COMPUTE_RETURN_ERROR_ON_MSG(out_shape.total_size() == 0, "Inputs are not broadcast compatible"); + + ARM_COMPUTE_RETURN_ERROR_ON_MSG(output->data_type() == DataType::U8 && (input1->data_type() != DataType::U8 || input2->data_type() != DataType::U8), + "Output can only be U8 if both inputs are U8"); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(output->data_type() == DataType::S32 && (input1->data_type() != DataType::QSYMM16 || input2->data_type() != DataType::QSYMM16), + "Output can only be S32 if both inputs are QSYMM16"); + ARM_COMPUTE_RETURN_ERROR_ON_MSG(output->data_type() == DataType::S32 && scale != 1.f, "Unsupported scale for QSYMM16 inputs and S32 output"); } if(std::abs(scale - scale255_constant) < 0.00001f) @@ -266,8 +258,20 @@ void mul_saturate_QSYMM16_QSYMM16_QSYMM16_n(const void *__restrict input1_ptr, c const auto input2 = static_cast(input2_ptr); const auto output = static_cast(output_ptr); - const qsymm16x8x2_t input1_q = vld2q_s16(input1); - const qsymm16x8x2_t input2_q = vld2q_s16(input2); + const qsymm16x8x2_t input1_q = + { + { + vld1q_s16(input1), + vld1q_s16(input1 + 8), + } + }; + const qsymm16x8x2_t input2_q = + { + { + vld1q_s16(input2), + vld1q_s16(input2 + 8), + } + }; // Dequantitize inputs const float32x4x4_t in1_f32x4x4 = vdequantize(input1_q, input1_qua_info); @@ -284,7 +288,65 @@ void mul_saturate_QSYMM16_QSYMM16_QSYMM16_n(const void *__restrict input1_ptr, c }; const qsymm16x8x2_t result = vquantize_qsymm16(out_f32x4x4, tmp_qua_info); - vst2q_s16(output, result); + vst1q_s16(output, result.val[0]); + vst1q_s16(output + 8, result.val[1]); +} + +void mul_QSYMM16_QSYMM16_S32_n(const void *__restrict input1_ptr, const void *__restrict input2_ptr, void *__restrict output_ptr, int scale) +{ + ARM_COMPUTE_UNUSED(scale); + const auto input1 = static_cast(input1_ptr); + const auto input2 = static_cast(input2_ptr); + const auto output = static_cast(output_ptr); + + const qsymm16x8x2_t input1_q = + { + { + vld1q_s16(input1), + vld1q_s16(input1 + 8), + } + }; + const qsymm16x8x2_t input2_q = + { + { + vld1q_s16(input2), + vld1q_s16(input2 + 8), + } + }; + + const int32x4x4_t in1_s32 = + { + { + vmovl_s16(vget_low_s16(input1_q.val[0])), + vmovl_s16(vget_high_s16(input1_q.val[0])), + vmovl_s16(vget_low_s16(input1_q.val[1])), + vmovl_s16(vget_high_s16(input1_q.val[1])), + } + }; + const int32x4x4_t in2_s32 = + { + { + vmovl_s16(vget_low_s16(input2_q.val[0])), + vmovl_s16(vget_high_s16(input2_q.val[0])), + vmovl_s16(vget_low_s16(input2_q.val[1])), + vmovl_s16(vget_high_s16(input2_q.val[1])), + } + }; + + const int32x4x4_t result = + { + { + vmulq_s32(in1_s32.val[0], in2_s32.val[0]), + vmulq_s32(in1_s32.val[1], in2_s32.val[1]), + vmulq_s32(in1_s32.val[2], in2_s32.val[2]), + vmulq_s32(in1_s32.val[3], in2_s32.val[3]), + } + }; + + vst1q_s32(output, result.val[0]); + vst1q_s32(output + 4, result.val[1]); + vst1q_s32(output + 8, result.val[2]); + vst1q_s32(output + 12, result.val[3]); } template @@ -412,11 +474,24 @@ void mul_S16_S16_S16_n(const void *__restrict input1_ptr, const void *__restrict const auto input2 = static_cast(input2_ptr); const auto output = static_cast(output_ptr); - const int16x8x2_t ta1 = vld2q_s16(input1); - const int16x8x2_t ta2 = vld2q_s16(input2); + const int16x8x2_t ta1 = + { + { + vld1q_s16(input1), + vld1q_s16(input1 + 8), + } + }; + const int16x8x2_t ta2 = + { + { + vld1q_s16(input2), + vld1q_s16(input2 + 8), + } + }; const int16x8x2_t result = mul_S16_S16_S16_n_k(ta1, ta2, n); - vst2q_s16(output, result); + vst1q_s16(output, result.val[0]); + vst1q_s16(output + 8, result.val[1]); } void mul_F32_F32_F32_n(const void *__restrict input1_ptr, const void *__restrict input2_ptr, void *__restrict output_ptr, float scale) @@ -472,11 +547,23 @@ void c_mul_F32_F32_F32_n(const void *__restrict input1_ptr, const void *__restri void mul_F16_F16_F16_n(const void *__restrict input1_ptr, const void *__restrict input2_ptr, void *__restrict output_ptr, float scale) { #ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC - const auto input1 = static_cast(input1_ptr); - const auto input2 = static_cast(input2_ptr); - const auto output = static_cast(output_ptr); - const float16x8x2_t ta1 = vld2q_f16(input1); - const float16x8x2_t ta2 = vld2q_f16(input2); + const auto input1 = static_cast(input1_ptr); + const auto input2 = static_cast(input2_ptr); + const auto output = static_cast(output_ptr); + const float16x8x2_t ta1 = + { + { + vld1q_f16(input1), + vld1q_f16(input1 + 8), + } + }; + const float16x8x2_t ta2 = + { + { + vld1q_f16(input2), + vld1q_f16(input2 + 8), + } + }; const float16x8_t scale_vec = vdupq_n_f16(scale); const float16x8x2_t result = { @@ -485,7 +572,8 @@ void mul_F16_F16_F16_n(const void *__restrict input1_ptr, const void *__restrict vmulq_f16(vmulq_f16(ta1.val[1], ta2.val[1]), scale_vec), } }; - vst2q_f16(output, result); + vst1q_f16(output, result.val[0]); + vst1q_f16(output + 8, result.val[1]); #else /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */ ARM_COMPUTE_UNUSED(input1_ptr); ARM_COMPUTE_UNUSED(input2_ptr); @@ -550,8 +638,20 @@ void mul_S16_U8_S16_n(const void *__restrict input1_ptr, const void *__restrict const auto input2 = static_cast(input2_ptr); const auto output = static_cast(output_ptr); - const int16x8x2_t ta1 = vld2q_s16(input1); - const uint8x8x2_t ta2u = vld2_u8(input2); + const int16x8x2_t ta1 = + { + { + vld1q_s16(input1), + vld1q_s16(input1 + 8), + } + }; + const uint8x8x2_t ta2u = + { + { + vld1_u8(input2), + vld1_u8(input2 + 8), + } + }; const int16x8x2_t ta2 = { { @@ -562,7 +662,8 @@ void mul_S16_U8_S16_n(const void *__restrict input1_ptr, const void *__restrict const int16x8x2_t result = mul_S16_S16_S16_n_k(ta1, ta2, n); - vst2q_s16(output, result); + vst1q_s16(output, result.val[0]); + vst1q_s16(output + 8, result.val[1]); } template @@ -629,10 +730,14 @@ void NEPixelWiseMultiplicationKernel::configure(const ITensor *input1, const ITe { _func_quantized = &mul_saturate_QASYMM8_SIGNED_QASYMM8_SIGNED_QASYMM8_SIGNED_n; } - else if(dt_input1 == DataType::QSYMM16 && dt_input2 == DataType::QSYMM16) + else if(dt_input1 == DataType::QSYMM16 && dt_input2 == DataType::QSYMM16 && dt_output == DataType::QSYMM16) { _func_quantized = &mul_saturate_QSYMM16_QSYMM16_QSYMM16_n; } + else if(dt_input1 == DataType::QSYMM16 && dt_input2 == DataType::QSYMM16 && dt_output == DataType::S32) + { + _func_int = &mul_QSYMM16_QSYMM16_S32_n; + } else if(DataType::U8 == dt_input1 && DataType::U8 == dt_input2 && DataType::U8 == dt_output) { if(is_scale_255) @@ -750,7 +855,7 @@ void NEPixelWiseMultiplicationKernel::run(const Window &window, const ThreadInfo Iterator input2(_input2, slice_input2); Iterator output(_output, slice); - if(is_data_type_quantized(_input1->info()->data_type())) + if((_run_optimized_qasymm8) || (_func_quantized != nullptr)) { if(_run_optimized_qasymm8) { diff --git a/tests/validation/CL/PixelWiseMultiplication.cpp b/tests/validation/CL/PixelWiseMultiplication.cpp index 3b55e25f37..ff9101a997 100644 --- a/tests/validation/CL/PixelWiseMultiplication.cpp +++ b/tests/validation/CL/PixelWiseMultiplication.cpp @@ -127,14 +127,17 @@ using CLPixelWiseMultiplicationQuantizedFixture = PixelWiseMultiplicationValidat TEST_SUITE(Quantized) TEST_SUITE(QASYMM8) -FIXTURE_DATA_TEST_CASE(RunSmall, CLPixelWiseMultiplicationQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::QASYMM8)), - framework::dataset::make("Scale", { 1.f, 2.f })), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), - framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_EVEN)), - framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), - framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OUtQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLPixelWiseMultiplicationQuantizedFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8)), + framework::dataset::make("Scale", { 1.f, 2.f })), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_EVEN)), + framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("OUtQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_qasymm8); @@ -142,14 +145,17 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLPixelWiseMultiplicationQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), - framework::dataset::make("Scale", { 1.f, 2.f })), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), - framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_EVEN)), - framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), - framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OUtQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLPixelWiseMultiplicationQuantizedFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8_SIGNED)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8_SIGNED)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8_SIGNED)), + framework::dataset::make("Scale", { 1.f, 2.f })), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_EVEN)), + framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("OUtQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_qasymm8); @@ -157,26 +163,32 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLPixelWiseMultiplicationQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::QSYMM16)), - framework::dataset::make("Scale", { 1.f, 2.f })), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), - framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_EVEN)), - framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 32768.f, 0) })), - framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLPixelWiseMultiplicationQuantizedFixture, framework::DatasetMode::PRECOMMIT, + combine(combine(combine(combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QSYMM16)), + framework::dataset::make("DataTypeIn2", DataType::QSYMM16)), + framework::dataset::make("DataTypeOut", DataType::QSYMM16)), + framework::dataset::make("Scale", { 1.f, 2.f })), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_EVEN)), + framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 32768.f, 0) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_qsymm16); } -FIXTURE_DATA_TEST_CASE(RunLarge, CLPixelWiseMultiplicationQuantizedFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(combine(combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::QSYMM16)), - framework::dataset::make("Scale", { 1.f, 2.f })), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), - framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_EVEN)), - framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 32768.f, 0) })), - framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) }))) +FIXTURE_DATA_TEST_CASE(RunLarge, CLPixelWiseMultiplicationQuantizedFixture, framework::DatasetMode::NIGHTLY, + combine(combine(combine(combine(combine(combine(combine(combine(combine(datasets::LargeShapes(), + framework::dataset::make("DataTypeIn1", DataType::QSYMM16)), + framework::dataset::make("DataTypeIn2", DataType::QSYMM16)), + framework::dataset::make("DataTypeOut", DataType::QSYMM16)), + framework::dataset::make("Scale", { 1.f, 2.f })), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE })), + framework::dataset::make("RoundingPolicy", RoundingPolicy::TO_NEAREST_EVEN)), + framework::dataset::make("Src0QInfo", { QuantizationInfo(1.f / 32768.f, 0) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 32768.f, 0) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(5.f / 32768.f, 0) }))) { // Validate output validate(CLAccessor(_target), _reference, tolerance_qsymm16); diff --git a/tests/validation/NEON/PixelWiseMultiplication.cpp b/tests/validation/NEON/PixelWiseMultiplication.cpp index fd54e42083..6a75b00b9b 100644 --- a/tests/validation/NEON/PixelWiseMultiplication.cpp +++ b/tests/validation/NEON/PixelWiseMultiplication.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -70,20 +70,6 @@ const auto PixelWiseMultiplicationPolicySTZDataset = combine( // *INDENT-OFF* // clang-format off -#define PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(DT1, DT2, SCALE, RP) \ - DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, \ - combine(combine(combine(combine(combine( \ - concat(datasets::SmallShapes(), datasets::LargeShapes()), \ - framework::dataset::make("DataType1", DataType::DT1)), \ - framework::dataset::make("DataType2", DataType::DT2)), \ - framework::dataset::make("Scale", std::move(SCALE))), \ - datasets::ConvertPolicies()), \ - framework::dataset::make("RoundingPolicy", RoundingPolicy::RP)), \ - shape, dt1, dt2, scale, convert_policy, rounding_policy) \ - { \ - validate_configuration(shape, dt1, dt2, scale, convert_policy, rounding_policy); \ - } - #define PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(TEST_NAME, FIXTURE, MODE, SHAPES, DT1, DT2, SCALE, RP, VALIDATE) \ FIXTURE_DATA_TEST_CASE(TEST_NAME, NEPixelWiseMultiplication##FIXTURE, framework::DatasetMode::MODE, \ combine(combine(combine(combine(combine( \ @@ -99,38 +85,12 @@ const auto PixelWiseMultiplicationPolicySTZDataset = combine( // *INDENT-ON* // clang-format on - -void validate_configuration(TensorShape shape, DataType dt1, DataType dt2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy) -{ - Tensor src1 = create_tensor(shape, dt1); - Tensor src2 = create_tensor(shape, dt2); - Tensor dst = create_tensor(shape, dt2); - - ARM_COMPUTE_EXPECT(src1.info()->is_resizable(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(src2.info()->is_resizable(), framework::LogLevel::ERRORS); - ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); - - // Create and configure function - NEPixelWiseMultiplication multiply; - multiply.configure(&src1, &src2, &dst, scale, convert_policy, rounding_policy); - - // Validate valid region - const ValidRegion valid_region = shape_to_valid_region(shape); - validate(src1.info()->valid_region(), valid_region); - validate(src2.info()->valid_region(), valid_region); - validate(dst.info()->valid_region(), valid_region); - - // Validate padding - const PaddingSize padding = PaddingCalculator(shape.x(), 16).required_padding(); - validate(src1.info()->padding(), padding); - validate(src2.info()->padding(), padding); - validate(dst.info()->padding(), padding); -} } // namespace using NEPixelWiseMultiplicationQASYMM8Fixture = PixelWiseMultiplicationValidationQuantizedFixture; using NEPixelWiseMultiplicationQASYMM8SignedFixture = PixelWiseMultiplicationValidationQuantizedFixture; using NEPixelWiseMultiplicationQSYMM16Fixture = PixelWiseMultiplicationValidationQuantizedFixture; +using NEPixelWiseMultiplicationQSYMM16ToS32Fixture = PixelWiseMultiplicationValidationQuantizedFixture; template using NEPixelWiseMultiplicationToU8Fixture = PixelWiseMultiplicationValidationFixture; template @@ -231,8 +191,10 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip(zip( TEST_SUITE(Quantized) TEST_SUITE(QASYMM8_SIGNED) TEST_SUITE(Scale255) -FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8SignedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), +FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8SignedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8_SIGNED)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8_SIGNED)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8_SIGNED)), framework::dataset::make("Scale", { scale_unity })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQASYMM8QuantDataset)) @@ -245,8 +207,10 @@ TEST_SUITE_END() // QASYMM8 TEST_SUITE(QASYMM8) TEST_SUITE(Scale255) -FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::QASYMM8)), +FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8)), framework::dataset::make("Scale", { scale_255 })), PixelWiseMultiplicationPolicySTNUDataset), PixelWiseMultiplicationQASYMM8QuantDataset)) @@ -254,8 +218,10 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framew // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::QASYMM8)), +FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(combine(combine(datasets::LargeShapes(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8)), framework::dataset::make("Scale", { scale_255 })), PixelWiseMultiplicationPolicySTNUDataset), PixelWiseMultiplicationQASYMM8QuantDataset)) @@ -265,8 +231,10 @@ FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQASYMM8Fixture, framew } TEST_SUITE_END() // Scale255 TEST_SUITE(ScaleUnity) -FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::QASYMM8)), +FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8)), framework::dataset::make("Scale", { scale_unity })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQASYMM8QuantDataset)) @@ -274,8 +242,10 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framew // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::QASYMM8)), +FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(combine(combine(datasets::LargeShapes(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8)), framework::dataset::make("Scale", { scale_unity })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQASYMM8QuantDataset)) @@ -285,8 +255,10 @@ FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQASYMM8Fixture, framew } TEST_SUITE_END() // ScaleUnity TEST_SUITE(ScaleOther) -FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::QASYMM8)), +FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8)), framework::dataset::make("Scale", { scale_other })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQASYMM8QuantDataset)) @@ -294,8 +266,10 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQASYMM8Fixture, framew // Validate output validate(Accessor(_target), _reference, tolerance_qasymm8); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::QASYMM8)), +FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQASYMM8Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(combine(combine(datasets::LargeShapes(), + framework::dataset::make("DataTypeIn1", DataType::QASYMM8)), + framework::dataset::make("DataTypeIn2", DataType::QASYMM8)), + framework::dataset::make("DataTypeOut", DataType::QASYMM8)), framework::dataset::make("Scale", { scale_other })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQASYMM8QuantDataset)) @@ -307,8 +281,10 @@ TEST_SUITE_END() // ScaleOther TEST_SUITE_END() // QASYMM8 TEST_SUITE(QSYMM16) TEST_SUITE(Scale255) -FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::QSYMM16)), +FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QSYMM16)), + framework::dataset::make("DataTypeIn2", DataType::QSYMM16)), + framework::dataset::make("DataTypeOut", DataType::QSYMM16)), framework::dataset::make("Scale", { scale_255 })), PixelWiseMultiplicationPolicySTNUDataset), PixelWiseMultiplicationQSYMM16QuantDataset)) @@ -316,8 +292,10 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framew // Validate output validate(Accessor(_target), _reference, tolerance_qsymm16); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::QSYMM16)), +FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(combine(combine(datasets::LargeShapes(), + framework::dataset::make("DataTypeIn1", DataType::QSYMM16)), + framework::dataset::make("DataTypeIn2", DataType::QSYMM16)), + framework::dataset::make("DataTypeOut", DataType::QSYMM16)), framework::dataset::make("Scale", { scale_255 })), PixelWiseMultiplicationPolicySTNUDataset), PixelWiseMultiplicationQSYMM16QuantDataset)) @@ -327,8 +305,10 @@ FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQSYMM16Fixture, framew } TEST_SUITE_END() // Scale255 TEST_SUITE(ScaleUnity) -FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::QSYMM16)), +FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QSYMM16)), + framework::dataset::make("DataTypeIn2", DataType::QSYMM16)), + framework::dataset::make("DataTypeOut", DataType::QSYMM16)), framework::dataset::make("Scale", { scale_unity })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQSYMM16QuantDataset)) @@ -336,8 +316,10 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framew // Validate output validate(Accessor(_target), _reference, tolerance_qsymm16); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::QSYMM16)), +FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(combine(combine(datasets::LargeShapes(), + framework::dataset::make("DataTypeIn1", DataType::QSYMM16)), + framework::dataset::make("DataTypeIn2", DataType::QSYMM16)), + framework::dataset::make("DataTypeOut", DataType::QSYMM16)), framework::dataset::make("Scale", { scale_unity })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQSYMM16QuantDataset)) @@ -347,8 +329,10 @@ FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQSYMM16Fixture, framew } TEST_SUITE_END() // ScaleUnity TEST_SUITE(ScaleOther) -FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), - framework::dataset::make("DataType", DataType::QSYMM16)), +FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QSYMM16)), + framework::dataset::make("DataTypeIn2", DataType::QSYMM16)), + framework::dataset::make("DataTypeOut", DataType::QSYMM16)), framework::dataset::make("Scale", { scale_other })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQSYMM16QuantDataset)) @@ -356,8 +340,10 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16Fixture, framew // Validate output validate(Accessor(_target), _reference, tolerance_qsymm16); } -FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(datasets::LargeShapes(), - framework::dataset::make("DataType", DataType::QSYMM16)), +FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQSYMM16Fixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(combine(combine(combine(datasets::LargeShapes(), + framework::dataset::make("DataTypeIn1", DataType::QSYMM16)), + framework::dataset::make("DataTypeIn2", DataType::QSYMM16)), + framework::dataset::make("DataTypeOut", DataType::QSYMM16)), framework::dataset::make("Scale", { scale_other })), PixelWiseMultiplicationPolicySTZDataset), PixelWiseMultiplicationQSYMM16QuantDataset)) @@ -367,24 +353,34 @@ FIXTURE_DATA_TEST_CASE(RunLarge, NEPixelWiseMultiplicationQSYMM16Fixture, framew } TEST_SUITE_END() // ScaleOther TEST_SUITE_END() // QSYMM16 +TEST_SUITE(QSYMM16toS32) +FIXTURE_DATA_TEST_CASE(RunSmall, NEPixelWiseMultiplicationQSYMM16ToS32Fixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(combine(combine(datasets::SmallShapes(), + framework::dataset::make("DataTypeIn1", DataType::QSYMM16)), + framework::dataset::make("DataTypeIn2", DataType::QSYMM16)), + framework::dataset::make("DataTypeOut", DataType::S32)), + framework::dataset::make("Scale", { scale_unity })), + PixelWiseMultiplicationPolicySTZDataset), + PixelWiseMultiplicationQSYMM16QuantDataset)) +{ + // Validate output + validate(Accessor(_target), _reference); +} +TEST_SUITE_END() // QSYMM16toS32 TEST_SUITE_END() // Quantized TEST_SUITE(U8toU8) TEST_SUITE(Scale255) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(U8, U8, scale_255, TO_NEAREST_UP) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToU8Fixture, PRECOMMIT, SmallShapes(), U8, U8, scale_255, TO_NEAREST_UP, WRAP_VALIDATE(uint8_t, 1)) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToU8Fixture, NIGHTLY, LargeShapes(), U8, U8, scale_255, TO_NEAREST_UP, WRAP_VALIDATE(uint8_t, 1)) TEST_SUITE_END() // Scale255 TEST_SUITE(ScaleUnity) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(U8, U8, scale_unity, TO_ZERO) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToU8Fixture, PRECOMMIT, SmallShapes(), U8, U8, scale_unity, TO_ZERO, DEFAULT_VALIDATE) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToU8Fixture, NIGHTLY, LargeShapes(), U8, U8, scale_unity, TO_ZERO, DEFAULT_VALIDATE) TEST_SUITE_END() // ScaleUnity TEST_SUITE(ScaleOther) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(U8, U8, scale_other, TO_ZERO) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToU8Fixture, PRECOMMIT, SmallShapes(), U8, U8, scale_other, TO_ZERO, DEFAULT_VALIDATE) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToU8Fixture, NIGHTLY, LargeShapes(), U8, U8, scale_other, TO_ZERO, DEFAULT_VALIDATE) TEST_SUITE_END() // ScaleOther @@ -394,19 +390,16 @@ TEST_SUITE_END() // U8toU8 TEST_SUITE(U8toS16) TEST_SUITE(Scale255) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(U8, S16, scale_255, TO_NEAREST_UP) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, PRECOMMIT, SmallShapes(), U8, S16, scale_255, TO_NEAREST_UP, WRAP_VALIDATE(int16_t, 2)) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToS16Fixture, NIGHTLY, LargeShapes(), U8, S16, scale_255, TO_NEAREST_UP, WRAP_VALIDATE(int16_t, 2)) TEST_SUITE_END() // Scale255 TEST_SUITE(ScaleUnity) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(U8, S16, scale_unity, TO_ZERO) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, PRECOMMIT, SmallShapes(), U8, S16, scale_unity, TO_ZERO, DEFAULT_VALIDATE) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToS16Fixture, NIGHTLY, LargeShapes(), U8, S16, scale_unity, TO_ZERO, DEFAULT_VALIDATE) TEST_SUITE_END() // ScaleUnity TEST_SUITE(ScaleOther) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(U8, S16, scale_other, TO_ZERO) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, PRECOMMIT, SmallShapes(), U8, S16, scale_other, TO_ZERO, DEFAULT_VALIDATE) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToS16Fixture, NIGHTLY, LargeShapes(), U8, S16, scale_other, TO_ZERO, DEFAULT_VALIDATE) TEST_SUITE_END() // ScaleOther @@ -416,19 +409,16 @@ TEST_SUITE_END() // U8toS16 TEST_SUITE(S16toS16) TEST_SUITE(Scale255) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(S16, S16, scale_255, TO_NEAREST_UP) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, PRECOMMIT, SmallShapes(), S16, S16, scale_255, TO_NEAREST_UP, WRAP_VALIDATE(int16_t, 2)) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToS16Fixture, NIGHTLY, LargeShapes(), S16, S16, scale_255, TO_NEAREST_UP, WRAP_VALIDATE(int16_t, 2)) TEST_SUITE_END() // Scale255 TEST_SUITE(ScaleUnity) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(S16, S16, scale_unity, TO_ZERO) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, PRECOMMIT, SmallShapes(), S16, S16, scale_unity, TO_ZERO, DEFAULT_VALIDATE) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToS16Fixture, NIGHTLY, LargeShapes(), S16, S16, scale_unity, TO_ZERO, DEFAULT_VALIDATE) TEST_SUITE_END() // ScaleUnity TEST_SUITE(ScaleOther) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(S16, S16, scale_other, TO_ZERO) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToS16Fixture, PRECOMMIT, SmallShapes(), S16, S16, scale_other, TO_ZERO, DEFAULT_VALIDATE) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToS16Fixture, NIGHTLY, LargeShapes(), S16, S16, scale_other, TO_ZERO, DEFAULT_VALIDATE) TEST_SUITE_END() // ScaleOther @@ -448,19 +438,16 @@ TEST_SUITE_END() // F16toF16 TEST_SUITE(F32toF32) TEST_SUITE(Scale255) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(F32, F32, scale_255, TO_NEAREST_UP) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToF32Fixture, PRECOMMIT, SmallShapes(), F32, F32, scale_255, TO_NEAREST_UP, VALIDATE(float, 1.f)) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToF32Fixture, NIGHTLY, LargeShapes(), F32, F32, scale_255, TO_NEAREST_UP, VALIDATE(float, 1.f)) TEST_SUITE_END() // Scale255 TEST_SUITE(ScaleUnity) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(F32, F32, scale_unity, TO_ZERO) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToF32Fixture, PRECOMMIT, SmallShapes(), F32, F32, scale_unity, TO_ZERO, DEFAULT_VALIDATE) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToF32Fixture, NIGHTLY, LargeShapes(), F32, F32, scale_unity, TO_ZERO, DEFAULT_VALIDATE) TEST_SUITE_END() // ScaleUnity TEST_SUITE(ScaleOther) -PIXEL_WISE_MULTIPLICATION_DATA_TEST_CASE(F32, F32, scale_other, TO_ZERO) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToF32Fixture, PRECOMMIT, SmallShapes(), F32, F32, scale_other, TO_ZERO, DEFAULT_VALIDATE) PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunLarge, ToF32Fixture, NIGHTLY, LargeShapes(), F32, F32, scale_other, TO_ZERO, DEFAULT_VALIDATE) TEST_SUITE_END() // ScaleOther diff --git a/tests/validation/fixtures/LSTMLayerFixture.h b/tests/validation/fixtures/LSTMLayerFixture.h index 9260686d56..858ee07d3e 100644 --- a/tests/validation/fixtures/LSTMLayerFixture.h +++ b/tests/validation/fixtures/LSTMLayerFixture.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -407,7 +407,7 @@ protected: if(peephole_opt) { - SimpleTensor pixelwise_mul_forget_gate = reference::pixel_wise_multiplication(cell_state_in, cell_to_forget_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_ZERO); + SimpleTensor pixelwise_mul_forget_gate = reference::pixel_wise_multiplication(cell_state_in, cell_to_forget_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_ZERO, data_type); forget_gate = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, forget_gate, pixelwise_mul_forget_gate, data_type, ConvertPolicy::SATURATE); } @@ -416,7 +416,7 @@ protected: SimpleTensor forget_layer_norm_w{ cell_bias_shape, data_type }; fill(forget_layer_norm_w, 23); forget_gate = reference::mean_std_normalization_layer(forget_gate); - forget_gate = reference::pixel_wise_multiplication(forget_gate, forget_layer_norm_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN); + forget_gate = reference::pixel_wise_multiplication(forget_gate, forget_layer_norm_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN, data_type); fill(forget_gate_bias, 7); forget_gate = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, forget_gate, forget_gate_bias, data_type, ConvertPolicy::SATURATE); } @@ -438,7 +438,7 @@ protected: input_gate = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, fully_connected_input, gemm, data_type, ConvertPolicy::SATURATE); if(peephole_opt) { - SimpleTensor pixelwise_mul_input_gate = reference::pixel_wise_multiplication(cell_state_in, cell_to_input_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN); + SimpleTensor pixelwise_mul_input_gate = reference::pixel_wise_multiplication(cell_state_in, cell_to_input_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN, data_type); input_gate = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, input_gate, pixelwise_mul_input_gate, data_type, ConvertPolicy::SATURATE); } if(use_layer_norm) @@ -446,7 +446,7 @@ protected: SimpleTensor input_layer_norm_w{ cell_bias_shape, data_type }; fill(input_layer_norm_w, 22); input_gate = reference::mean_std_normalization_layer(input_gate); - input_gate = reference::pixel_wise_multiplication(input_gate, input_layer_norm_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN); + input_gate = reference::pixel_wise_multiplication(input_gate, input_layer_norm_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN, data_type); fill(input_gate_bias, 17); input_gate = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, input_gate, input_gate_bias, data_type, ConvertPolicy::SATURATE); } @@ -457,19 +457,19 @@ protected: SimpleTensor fully_connected_cell_state = reference::fully_connected_layer(input, input_to_cell_w, cell_bias, output_cell_shape); transposed_weights = reference::transpose(recurrent_to_cell_w); gemm = reference::gemm(output_state_in, transposed_weights, cell_state_out, 1.f, 0.f); - SimpleTensor pixelwise_mul = reference::pixel_wise_multiplication(cell_state_in, forget_gate, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN); + SimpleTensor pixelwise_mul = reference::pixel_wise_multiplication(cell_state_in, forget_gate, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN, data_type); cell_state_out = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, fully_connected_cell_state, gemm, data_type, ConvertPolicy::SATURATE); if(use_layer_norm) { SimpleTensor cell_layer_norm_w{ cell_bias_shape, data_type }; fill(cell_layer_norm_w, 24); cell_state_out = reference::mean_std_normalization_layer(cell_state_out); - cell_state_out = reference::pixel_wise_multiplication(cell_state_out, cell_layer_norm_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN); + cell_state_out = reference::pixel_wise_multiplication(cell_state_out, cell_layer_norm_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN, data_type); fill(cell_bias, 8); cell_state_out = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, cell_state_out, cell_bias, data_type, ConvertPolicy::SATURATE); } cell_state_out = reference::activation_layer(cell_state_out, ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC)); - cell_state_out = reference::pixel_wise_multiplication(cell_state_out, input_gate, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN); + cell_state_out = reference::pixel_wise_multiplication(cell_state_out, input_gate, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN, data_type); cell_state_out = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, cell_state_out, pixelwise_mul, data_type, ConvertPolicy::SATURATE); if(cell_threshold != 0.f) { @@ -483,7 +483,7 @@ protected: output = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, fully_connected_output, gemm, data_type, ConvertPolicy::SATURATE); if(peephole_opt) { - pixelwise_mul = reference::pixel_wise_multiplication(cell_state_out, cell_to_output_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN); + pixelwise_mul = reference::pixel_wise_multiplication(cell_state_out, cell_to_output_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN, data_type); output = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, output, pixelwise_mul, data_type, ConvertPolicy::SATURATE); } if(use_layer_norm) @@ -491,7 +491,7 @@ protected: SimpleTensor output_layer_norm_w{ cell_bias_shape, data_type }; fill(output_layer_norm_w, 25); output = reference::mean_std_normalization_layer(output); - output = reference::pixel_wise_multiplication(output, output_layer_norm_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN); + output = reference::pixel_wise_multiplication(output, output_layer_norm_w, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN, data_type); fill(output_gate_bias, 9); output = reference::arithmetic_operation(reference::ArithmeticOperation::ADD, output, output_gate_bias, data_type, ConvertPolicy::SATURATE); } @@ -499,7 +499,7 @@ protected: // Compute output state SimpleTensor cell_state_activation = reference::activation_layer(cell_state_out, info); - output_state_out = reference::pixel_wise_multiplication(output, cell_state_activation, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN); + output_state_out = reference::pixel_wise_multiplication(output, cell_state_activation, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_NEAREST_EVEN, data_type); if(projection_opt) { diff --git a/tests/validation/fixtures/PixelWiseMultiplicationFixture.h b/tests/validation/fixtures/PixelWiseMultiplicationFixture.h index efdf5d078e..37359f421b 100644 --- a/tests/validation/fixtures/PixelWiseMultiplicationFixture.h +++ b/tests/validation/fixtures/PixelWiseMultiplicationFixture.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -39,7 +39,7 @@ namespace test { namespace validation { -template +template class PixelWiseMultiplicationGenericValidationFixture : public framework::Fixture { public: @@ -48,6 +48,7 @@ public: const TensorShape &shape1, DataType dt_in1, DataType dt_in2, + DataType dt_out, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, @@ -55,8 +56,8 @@ public: QuantizationInfo qinfo1, QuantizationInfo qinfo_out) { - _target = compute_target(shape0, shape1, dt_in1, dt_in2, scale, convert_policy, rounding_policy, qinfo0, qinfo1, qinfo_out); - _reference = compute_reference(shape0, shape1, dt_in1, dt_in2, scale, convert_policy, rounding_policy, qinfo0, qinfo1, qinfo_out); + _target = compute_target(shape0, shape1, dt_in1, dt_in2, dt_out, scale, convert_policy, rounding_policy, qinfo0, qinfo1, qinfo_out); + _reference = compute_reference(shape0, shape1, dt_in1, dt_in2, dt_out, scale, convert_policy, rounding_policy, qinfo0, qinfo1, qinfo_out); } protected: @@ -66,14 +67,14 @@ protected: library->fill_tensor_uniform(tensor, seed_offset); } - TensorType compute_target(const TensorShape &shape0, const TensorShape &shape1, DataType dt_in1, DataType dt_in2, + TensorType compute_target(const TensorShape &shape0, const TensorShape &shape1, DataType dt_in1, DataType dt_in2, DataType dt_out, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) { // Create tensors TensorType src1 = create_tensor(shape0, dt_in1, 1, qinfo0); TensorType src2 = create_tensor(shape1, dt_in2, 1, qinfo1); - TensorType dst = create_tensor(TensorShape::broadcast_shape(shape0, shape1), dt_in2, 1, qinfo_out); + TensorType dst = create_tensor(TensorShape::broadcast_shape(shape0, shape1), dt_out, 1, qinfo_out); // Create and configure function FunctionType multiply; @@ -102,7 +103,7 @@ protected: return dst; } - SimpleTensor compute_reference(const TensorShape &shape0, const TensorShape &shape1, DataType dt_in1, DataType dt_in2, + SimpleTensor compute_reference(const TensorShape &shape0, const TensorShape &shape1, DataType dt_in1, DataType dt_in2, DataType dt_out, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) { @@ -114,24 +115,11 @@ protected: fill(src1, 0); fill(src2, 1); - return reference::pixel_wise_multiplication(src1, src2, scale, convert_policy, rounding_policy, qinfo_out); + return reference::pixel_wise_multiplication(src1, src2, scale, convert_policy, rounding_policy, dt_out, qinfo_out); } TensorType _target{}; - SimpleTensor _reference{}; -}; - -template -class PixelWiseMultiplicationQuatizedValidationFixture : public PixelWiseMultiplicationGenericValidationFixture -{ -public: - template - void setup(const TensorShape &shape, DataType dt_in1, DataType dt_in2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, - QuantizationInfo in1_qua_info, QuantizationInfo in2_qua_info, QuantizationInfo out_qua_info) - { - PixelWiseMultiplicationGenericValidationFixture::setup(shape, shape, dt_in1, dt_in2, scale, convert_policy, rounding_policy, - in1_qua_info, in2_qua_info, out_qua_info); - } + SimpleTensor _reference{}; }; template @@ -141,7 +129,7 @@ public: template void setup(const TensorShape &shape, DataType dt_in1, DataType dt_in2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy) { - PixelWiseMultiplicationGenericValidationFixture::setup(shape, shape, dt_in1, dt_in2, scale, convert_policy, rounding_policy, + PixelWiseMultiplicationGenericValidationFixture::setup(shape, shape, dt_in1, dt_in2, dt_in2, scale, convert_policy, rounding_policy, QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); } }; @@ -153,21 +141,21 @@ public: template void setup(const TensorShape &shape0, const TensorShape &shape1, DataType dt_in1, DataType dt_in2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy) { - PixelWiseMultiplicationGenericValidationFixture::setup(shape0, shape1, dt_in1, dt_in2, scale, convert_policy, rounding_policy, + PixelWiseMultiplicationGenericValidationFixture::setup(shape0, shape1, dt_in1, dt_in2, dt_in2, scale, convert_policy, rounding_policy, QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); } }; -template -class PixelWiseMultiplicationValidationQuantizedFixture : public PixelWiseMultiplicationGenericValidationFixture +template +class PixelWiseMultiplicationValidationQuantizedFixture : public PixelWiseMultiplicationGenericValidationFixture { public: template - void setup(const TensorShape &shape, DataType dt, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, + void setup(const TensorShape &shape, DataType dt_in1, DataType dt_in2, DataType dt_out, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) { - PixelWiseMultiplicationGenericValidationFixture::setup(shape, shape, dt, dt, scale, convert_policy, rounding_policy, - qinfo0, qinfo1, qinfo_out); + PixelWiseMultiplicationGenericValidationFixture::setup(shape, shape, dt_in1, dt_in2, dt_out, scale, convert_policy, + rounding_policy, qinfo0, qinfo1, qinfo_out); } }; } // namespace validation diff --git a/tests/validation/reference/PixelWiseMultiplication.cpp b/tests/validation/reference/PixelWiseMultiplication.cpp index 2b4c849c39..3e21fca72a 100644 --- a/tests/validation/reference/PixelWiseMultiplication.cpp +++ b/tests/validation/reference/PixelWiseMultiplication.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -52,16 +52,16 @@ namespace * @param[in] convert_policy Overflow policy. Supported overflow policies: Wrap, Saturate * @param[in] rounding_policy Rounding policy. Supported rounding modes: to zero, to nearest even. */ -template -T2 mul(const T1 src1, const T2 src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy) +template +T3 mul(const T1 src1, const T2 src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy) { - using intermediate_type = typename common_promoted_signed_type::intermediate_type; + using intermediate_type = typename common_promoted_signed_type::intermediate_type; const double val = static_cast(src1) * static_cast(src2) * static_cast(scale); - if(is_floating_point::value) + if(is_floating_point::value) { - const auto result = static_cast(val); + const auto result = static_cast(val); return result; } @@ -83,7 +83,7 @@ T2 mul(const T1 src1, const T2 src2, float scale, ConvertPolicy convert_policy, ARM_COMPUTE_ERROR("Unsupported rounding policy"); } - const auto result = static_cast((convert_policy == ConvertPolicy::SATURATE) ? saturate_cast(rounded_val) : rounded_val); + const auto result = static_cast((convert_policy == ConvertPolicy::SATURATE) ? saturate_cast(rounded_val) : rounded_val); return result; } @@ -92,8 +92,8 @@ T2 mul(const T1 src1, const T2 src2, float scale, ConvertPolicy convert_policy, template struct BroadcastUnroll { - template - static void unroll(const SimpleTensor &src1, const SimpleTensor &src2, SimpleTensor &dst, + template + static void unroll(const SimpleTensor &src1, const SimpleTensor &src2, SimpleTensor &dst, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, Coordinates &id_src1, Coordinates &id_src2, Coordinates &id_dst) { @@ -117,23 +117,23 @@ struct BroadcastUnroll template <> struct BroadcastUnroll<0> { - template - static void unroll(const SimpleTensor &src1, const SimpleTensor &src2, SimpleTensor &dst, + template + static void unroll(const SimpleTensor &src1, const SimpleTensor &src2, SimpleTensor &dst, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, Coordinates &id_src1, Coordinates &id_src2, Coordinates &id_dst) { - dst[coord2index(dst.shape(), id_dst)] = mul(src1[coord2index(src1.shape(), id_src1)], src2[coord2index(src2.shape(), id_src2)], scale, convert_policy, rounding_policy); + dst[coord2index(dst.shape(), id_dst)] = mul(src1[coord2index(src1.shape(), id_src1)], src2[coord2index(src2.shape(), id_src2)], scale, convert_policy, rounding_policy); } }; } // namespace -template -SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, - const QuantizationInfo &qout) +template +SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, + DataType dt_out, const QuantizationInfo &qout) { ARM_COMPUTE_UNUSED(qout); - SimpleTensor dst(TensorShape::broadcast_shape(src1.shape(), src2.shape()), src2.data_type()); + SimpleTensor dst(TensorShape::broadcast_shape(src1.shape(), src2.shape()), dt_out); if(scale < 0) { @@ -151,15 +151,15 @@ SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const S template <> SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, - const QuantizationInfo &qout) + DataType dt_out, const QuantizationInfo &qout) { - SimpleTensor dst(TensorShape::broadcast_shape(src1.shape(), src2.shape()), src2.data_type(), 1, qout); + SimpleTensor dst(TensorShape::broadcast_shape(src1.shape(), src2.shape()), dt_out, 1, qout); if(src1.data_type() == DataType::QASYMM8 && src2.data_type() == DataType::QASYMM8) { SimpleTensor src1_tmp = convert_from_asymmetric(src1); SimpleTensor src2_tmp = convert_from_asymmetric(src2); - SimpleTensor dst_tmp = pixel_wise_multiplication(src1_tmp, src2_tmp, scale, convert_policy, rounding_policy, qout); + SimpleTensor dst_tmp = pixel_wise_multiplication(src1_tmp, src2_tmp, scale, convert_policy, rounding_policy, DataType::F32, qout); dst = convert_to_asymmetric(dst_tmp, qout); } else @@ -179,15 +179,15 @@ SimpleTensor pixel_wise_multiplication(const SimpleTensor &src template <> SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, - const QuantizationInfo &qout) + DataType dt_out, const QuantizationInfo &qout) { - SimpleTensor dst(TensorShape::broadcast_shape(src1.shape(), src2.shape()), src2.data_type(), 1, qout); + SimpleTensor dst(TensorShape::broadcast_shape(src1.shape(), src2.shape()), dt_out, 1, qout); if(src1.data_type() == DataType::QASYMM8_SIGNED && src2.data_type() == DataType::QASYMM8_SIGNED) { SimpleTensor src1_tmp = convert_from_asymmetric(src1); SimpleTensor src2_tmp = convert_from_asymmetric(src2); - SimpleTensor dst_tmp = pixel_wise_multiplication(src1_tmp, src2_tmp, scale, convert_policy, rounding_policy, qout); + SimpleTensor dst_tmp = pixel_wise_multiplication(src1_tmp, src2_tmp, scale, convert_policy, rounding_policy, DataType::F32, qout); dst = convert_to_asymmetric(dst_tmp, qout); } else @@ -207,15 +207,15 @@ SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, template <> SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, - const QuantizationInfo &qout) + DataType dt_out, const QuantizationInfo &qout) { - SimpleTensor dst(TensorShape::broadcast_shape(src1.shape(), src2.shape()), src2.data_type(), 1, qout); + SimpleTensor dst(TensorShape::broadcast_shape(src1.shape(), src2.shape()), dt_out, 1, qout); if(src1.data_type() == DataType::QSYMM16 && src2.data_type() == DataType::QSYMM16) { SimpleTensor src1_tmp = convert_from_symmetric(src1); SimpleTensor src2_tmp = convert_from_symmetric(src2); - SimpleTensor dst_tmp = pixel_wise_multiplication(src1_tmp, src2_tmp, scale, convert_policy, rounding_policy, qout); + SimpleTensor dst_tmp = pixel_wise_multiplication(src1_tmp, src2_tmp, scale, convert_policy, rounding_policy, DataType::F32, qout); dst = convert_to_symmetric(dst_tmp, qout); } else @@ -234,9 +234,10 @@ SimpleTensor pixel_wise_multiplication(const SimpleTensor &src } // *INDENT-OFF* // clang-format off -template SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, const QuantizationInfo &qout); -template SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, const QuantizationInfo &qout); -template SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, const QuantizationInfo &qout); +template SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, DataType dt_out, const QuantizationInfo &qout); +template SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, DataType dt_out, const QuantizationInfo &qout); +template SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, DataType dt_out, const QuantizationInfo &qout); +template SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, DataType dt_out, const QuantizationInfo &qout); // clang-format on // *INDENT-ON* } // namespace reference diff --git a/tests/validation/reference/PixelWiseMultiplication.h b/tests/validation/reference/PixelWiseMultiplication.h index f5b8e777fb..f8afa0384b 100644 --- a/tests/validation/reference/PixelWiseMultiplication.h +++ b/tests/validation/reference/PixelWiseMultiplication.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -34,9 +34,10 @@ namespace validation { namespace reference { -template -SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, - ConvertPolicy convert_policy, RoundingPolicy rounding_policy, const QuantizationInfo &qout = QuantizationInfo()); +template +SimpleTensor pixel_wise_multiplication(const SimpleTensor &src1, const SimpleTensor &src2, float scale, + ConvertPolicy convert_policy, RoundingPolicy rounding_policy, DataType dt_out, + const QuantizationInfo &qout = QuantizationInfo()); } // namespace reference } // namespace validation } // namespace test diff --git a/tests/validation/reference/QLSTMLayerNormalization.cpp b/tests/validation/reference/QLSTMLayerNormalization.cpp index 90d59b93ad..0e24de6584 100644 --- a/tests/validation/reference/QLSTMLayerNormalization.cpp +++ b/tests/validation/reference/QLSTMLayerNormalization.cpp @@ -41,7 +41,7 @@ namespace reference SimpleTensor qlstm_layer_normalization_float_compute(SimpleTensor src, SimpleTensor weight, SimpleTensor bias) { SimpleTensor output = mean_std_normalization_layer(src); - output = pixel_wise_multiplication(output, weight, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_ZERO); + output = pixel_wise_multiplication(output, weight, 1, ConvertPolicy::SATURATE, RoundingPolicy::TO_ZERO, DataType::F32); return arithmetic_operation(ArithmeticOperation::ADD, output, bias, DataType::F32, ConvertPolicy::SATURATE); } -- cgit v1.2.1