diff options
author | Manuel Bottini <manuel.bottini@arm.com> | 2019-10-17 18:37:26 +0100 |
---|---|---|
committer | Pablo Marquez <pablo.tello@arm.com> | 2019-10-24 09:11:40 +0000 |
commit | 07263980e66059a91ce57612e4ca8f4b2a2a206a (patch) | |
tree | 138dc3ecf835df9f38a60959379a52eca08f8b0f /src/core | |
parent | 05069f07bcf95676597698a79926327555276362 (diff) | |
download | ComputeLibrary-07263980e66059a91ce57612e4ca8f4b2a2a206a.tar.gz |
COMPMID-2501: Support multiplier > 1 during QASYMM8 requantization for Quantized LSTM
Change-Id: I7eddbdf77881f313b707b9e59428245f1330a2cf
Signed-off-by: Manuel Bottini <manuel.bottini@arm.com>
Reviewed-on: https://review.mlplatform.org/c/2119
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Pablo Marquez <pablo.tello@arm.com>
Diffstat (limited to 'src/core')
-rw-r--r-- | src/core/CL/cl_kernels/gemmlowp.cl | 4 | ||||
-rw-r--r-- | src/core/utils/quantization/AsymmHelpers.cpp | 14 |
2 files changed, 18 insertions, 0 deletions
diff --git a/src/core/CL/cl_kernels/gemmlowp.cl b/src/core/CL/cl_kernels/gemmlowp.cl index fc90dbd16c..214c7a4825 100644 --- a/src/core/CL/cl_kernels/gemmlowp.cl +++ b/src/core/CL/cl_kernels/gemmlowp.cl @@ -1888,7 +1888,11 @@ __kernel void gemmlowp_output_stage_quantize_down_fixedpoint_qsymm16(TENSOR3D_DE #endif // defined(ADD_BIAS) // Multiply by result_mult_int and shift +#if RESULT_SHIFT < 0 + input_values = ASYMM_MULT(input_values * (1 << (-RESULT_SHIFT)), RESULT_FIXEDPOINT_MULTIPLIER, 4); +#else // RESULT_SHIFT >= 0 input_values = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(input_values, RESULT_FIXEDPOINT_MULTIPLIER, RESULT_SHIFT, 4); +#endif // RESULT_SHIFT < 0 short4 res = convert_short4_sat(input_values); diff --git a/src/core/utils/quantization/AsymmHelpers.cpp b/src/core/utils/quantization/AsymmHelpers.cpp index 59052449af..42bd84db47 100644 --- a/src/core/utils/quantization/AsymmHelpers.cpp +++ b/src/core/utils/quantization/AsymmHelpers.cpp @@ -34,6 +34,20 @@ namespace quantization constexpr int64_t fixed_point_one_Q0 = (1LL << 31); constexpr float epsilon = 0.00001f; +Status calculate_quantized_multiplier(float multiplier, int *quant_multiplier, int *shift) +{ + if(multiplier > 1.f) + { + Status status = calculate_quantized_multiplier_greater_than_one(multiplier, quant_multiplier, shift); + *shift *= -1; + return status; + } + else + { + return calculate_quantized_multiplier_less_than_one(multiplier, quant_multiplier, shift); + } +} + Status calculate_quantized_multiplier_less_than_one(float multiplier, int *quant_multiplier, int *right_shift) |