From de60a14c16bf384d4a30c7e16a62ed2c89da333f Mon Sep 17 00:00:00 2001 From: Tracy Sharpe <42477615+tracysh@users.noreply.github.com> Date: Tue, 7 Apr 2020 15:01:20 -0700 Subject: [PATCH] Fix output range for int8_t QuantizeLinear op (#3445) --- onnxruntime/core/mlas/inc/mlas.h | 15 +- onnxruntime/core/mlas/lib/quantize.cpp | 148 +++++++++--------- .../providers/cuda/tensor/quantize_linear.cu | 2 +- .../cpu/tensor/quantize_linear_test.cc | 12 +- 4 files changed, 83 insertions(+), 94 deletions(-) diff --git a/onnxruntime/core/mlas/inc/mlas.h b/onnxruntime/core/mlas/inc/mlas.h index 98722863b9..9a6f068b1f 100644 --- a/onnxruntime/core/mlas/inc/mlas.h +++ b/onnxruntime/core/mlas/inc/mlas.h @@ -420,24 +420,15 @@ MlasNchwcUpsample( // Linear quantization routines. // +template void MLASCALL MlasQuantizeLinear( const float* Input, - uint8_t* Output, + OutputType* Output, size_t N, float Scale, - uint8_t ZeroPoint - ); - -void -MLASCALL -MlasQuantizeLinear( - const float* Input, - int8_t* Output, - size_t N, - float Scale, - int8_t ZeroPoint + OutputType ZeroPoint ); void diff --git a/onnxruntime/core/mlas/lib/quantize.cpp b/onnxruntime/core/mlas/lib/quantize.cpp index fea003031a..2f8658e563 100644 --- a/onnxruntime/core/mlas/lib/quantize.cpp +++ b/onnxruntime/core/mlas/lib/quantize.cpp @@ -130,17 +130,44 @@ MlasQuantizeLinearPackBytes( #endif -template -MLAS_FORCEINLINE +template void -MlasQuantizeLinearKernel( +MLASCALL +MlasQuantizeLinear( const float* Input, OutputType* Output, size_t N, float Scale, - int32_t ZeroPoint + OutputType ZeroPoint ) +/*++ + +Routine Description: + + This routine quantizes the input buffer using the supplied quantization + parameters. + +Arguments: + + Input - Supplies the input buffer. + + Output - Supplies the output buffer. + + N - Supplies the number of elements to process. + + Scale - Supplies the quantization scale. + + ZeroPoint - Supplies the quantization zero point value. + +Return Value: + + None. + +--*/ { + constexpr int32_t MinimumValue = std::numeric_limits::min(); + constexpr int32_t MaximumValue = std::numeric_limits::max(); + auto ScaleVector = MlasBroadcastFloat32x4(Scale); auto MinimumValueVector = MlasBroadcastFloat32x4(float(MinimumValue - ZeroPoint)); auto MaximumValueVector = MlasBroadcastFloat32x4(float(MaximumValue - ZeroPoint)); @@ -189,17 +216,44 @@ MlasQuantizeLinearKernel( // QuantizeLinear implementation using the C++ runtime. // -template -MLAS_FORCEINLINE +template void -MlasQuantizeLinearKernel( +MLASCALL +MlasQuantizeLinear( const float* Input, OutputType* Output, size_t N, float Scale, - int32_t ZeroPoint + OutputType ZeroPoint ) +/*++ + +Routine Description: + + This routine quantizes the input buffer using the supplied quantization + parameters. + +Arguments: + + Input - Supplies the input buffer. + + Output - Supplies the output buffer. + + N - Supplies the number of elements to process. + + Scale - Supplies the quantization scale. + + ZeroPoint - Supplies the quantization zero point value. + +Return Value: + + None. + +--*/ { + constexpr int32_t MinimumValue = std::numeric_limits::min(); + constexpr int32_t MaximumValue = std::numeric_limits::max(); + for (size_t n = 0; n < N; n++) { float FloatValue = std::nearbyintf(Input[n] / Scale) + float(ZeroPoint); @@ -211,79 +265,27 @@ MlasQuantizeLinearKernel( #endif +template void MLASCALL -MlasQuantizeLinear( - const float* Input, - uint8_t* Output, - size_t N, - float Scale, - uint8_t ZeroPoint - ) -/*++ - -Routine Description: - - This routine quantizes the input buffer using the supplied quantization - parameters. - -Arguments: - - Input - Supplies the input buffer. - - Output - Supplies the output buffer. - - N - Supplies the number of elements to process. - - Scale - Supplies the quantization scale. - - ZeroPoint - Supplies the quantization zero point value. - -Return Value: - - None. - ---*/ -{ - return MlasQuantizeLinearKernel(Input, Output, N, Scale, ZeroPoint); -} - -void -MLASCALL -MlasQuantizeLinear( +MlasQuantizeLinear( const float* Input, int8_t* Output, size_t N, float Scale, int8_t ZeroPoint - ) -/*++ + ); -Routine Description: - - This routine quantizes the input buffer using the supplied quantization - parameters. - -Arguments: - - Input - Supplies the input buffer. - - Output - Supplies the output buffer. - - N - Supplies the number of elements to process. - - Scale - Supplies the quantization scale. - - ZeroPoint - Supplies the quantization zero point value. - -Return Value: - - None. - ---*/ -{ - return MlasQuantizeLinearKernel(Input, Output, N, Scale, ZeroPoint); -} +template +void +MLASCALL +MlasQuantizeLinear( + const float* Input, + uint8_t* Output, + size_t N, + float Scale, + uint8_t ZeroPoint + ); #if defined(MLAS_SSE2_INTRINSICS) diff --git a/onnxruntime/core/providers/cuda/tensor/quantize_linear.cu b/onnxruntime/core/providers/cuda/tensor/quantize_linear.cu index 89f0e40dcc..ec49023ced 100644 --- a/onnxruntime/core/providers/cuda/tensor/quantize_linear.cu +++ b/onnxruntime/core/providers/cuda/tensor/quantize_linear.cu @@ -16,7 +16,7 @@ __global__ void QuantizeLinearKernel(const float* input, int8_t* output, const f for (int i = 0; i < NumElementsPerThread; i++) { if (id < N) { int value = __float2int_rn(input[id] / (*scale)) + *zero_point; - output[id] = static_cast(max(-127, min(127, value))); + output[id] = static_cast(max(-128, min(127, value))); id += NumThreadsPerBlock; } } diff --git a/onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc b/onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc index f9b15105cf..fdea852b93 100644 --- a/onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc +++ b/onnxruntime/test/providers/cpu/tensor/quantize_linear_test.cc @@ -75,10 +75,8 @@ TEST(QuantizeLinearOpTest, QuantizeLinear_int8_NegativeZeroPoint) { test.AddInput("x", dims, {0, 2, 3, 5, 6, -2, -5, -6}); test.AddInput("y_scale", {}, {.039215686f}); test.AddInput("y_zero_point", {}, {-23}); - test.AddOutput("y", dims, {-23, 28, 53, 104, 127, -74, -127, -127}); - // TODO: nGraph returns a range of [-128,127] but the default CPU provider - // returns a range of [-127,127]. Resolve this difference in behavior later. - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider}); + test.AddOutput("y", dims, {-23, 28, 53, 104, 127, -74, -128, -128}); + test.Run(); } // quantize with scalar zero point and scale @@ -88,10 +86,8 @@ TEST(QuantizeLinearOpTest, QuantizeLinear_int8_PositiveZeroPoint) { test.AddInput("x", dims, {0, 2, 3, 5, 6, -2, -5, -6}); test.AddInput("y_scale", {}, {.039215686f}); test.AddInput("y_zero_point", {}, {23}); - test.AddOutput("y", dims, {23, 74, 99, 127, 127, -28, -104, -127}); - // TODO: nGraph returns a range of [-128,127] but the default CPU provider - // returns a range of [-127,127]. Resolve this difference in behavior later. - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider}); + test.AddOutput("y", dims, {23, 74, 99, 127, 127, -28, -104, -128}); + test.Run(); } // quantize with 2D data