From 62fc3a305952811fdf703cf3e2ddeccb06a57032 Mon Sep 17 00:00:00 2001 From: mirounga Date: Fri, 30 Oct 2026 13:13:33 -0700 Subject: [PATCH 1/2] [CPU] Fuse 16-bit QuantizeLinear -> DequantizeLinear into QuantizeDequantizeLinear A16 models keep 16-bit activation Q/DQ boundaries around every quantized op, including the MatMulNBits nodes fused in #32885. On CPU each pair runs as two kernels and materializes an int16/uint16 tensor that is immediately dequantized again. Add the com.microsoft.QuantizeDequantizeLinear op (float input, int16 or uint16 zero point) and a CPU kernel for it: - MlasQuantizeDequantizeLinear runs the per-element steps of the MLAS QuantizeLinear kernel and then (q - zp) * scale in one vector loop (SSE2, NEON, LSX). On other platforms, or when a non-default x64 quantize kernel is selected, it falls back to two passes through MlasQuantizeLinear. - Results are bitwise identical to the unfused QuantizeLinear -> DequantizeLinear pair. Add QuantizeDequantizeLinearFusion, a Level3 CPU transformer that replaces a 16-bit QuantizeLinear whose consumers are all DequantizeLinear nodes with the same constant scalar scale and zero point by one fused node. It also follows the quantized tensor through data-movement ops (Transpose, Reshape, Slice, Gather, Split, Unsqueeze, ...) to matching DQs and runs those ops on the float values. It runs after the Level2 Q/DQ node-unit fusions, so it never removes a Q/DQ they could still use, and session.disable_quant_qdq disables it. Tests: - QuantizeDequantizeLinearOpTest: int16 and uint16 round-half-to-even and saturation, a large tensor against a reference, empty input, and rejection of per-axis parameters. - MLAS QuantizeDequantizeLinearS16/U16: bitwise comparison with MlasQuantizeLinear followed by dequantization. - QuantizeDequantizeLinearFusionTests: interior pairs are bit-exact, fusion through data-movement ops, fan-out and graph-output DQs, A16 MatMulNBits boundaries, session.disable_quant_qdq, and the cases that must not fuse (8-bit, mismatched or non-constant or per-axis parameters, other uses of the quantized value, non-CPU EPs). Co-authored-by: Jovan Bubonja --- docs/ContribOperators.md | 41 ++ docs/OperatorKernels.md | 1 + .../contrib_ops/cpu/cpu_contrib_kernels.cc | 6 + .../quantize_dequantize_linear.cc | 75 ++++ onnxruntime/core/graph/contrib_ops/ms_opset.h | 2 + .../graph/contrib_ops/quantization_defs.cc | 23 ++ onnxruntime/core/mlas/inc/mlas.h | 16 + onnxruntime/core/mlas/lib/quantize.cpp | 141 +++++++ .../core/optimizer/graph_transformer_utils.cc | 6 + .../quantize_dequantize_linear_fusion.cc | 215 ++++++++++ .../quantize_dequantize_linear_fusion.h | 30 ++ .../quantize_dequantize_linear_op_test.cc | 87 ++++ .../mlas/unittest/test_quantizelinear.cpp | 74 ++++ .../quantize_dequantize_linear_fusion_test.cc | 388 ++++++++++++++++++ 14 files changed, 1105 insertions(+) create mode 100644 onnxruntime/contrib_ops/cpu/quantization/quantize_dequantize_linear.cc create mode 100644 onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.cc create mode 100644 onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.h create mode 100644 onnxruntime/test/contrib_ops/quantize_dequantize_linear_op_test.cc create mode 100644 onnxruntime/test/optimizer/quantize_dequantize_linear_fusion_test.cc diff --git a/docs/ContribOperators.md b/docs/ContribOperators.md index 2589e5878a808..05a0ca1908f90 100644 --- a/docs/ContribOperators.md +++ b/docs/ContribOperators.md @@ -110,6 +110,7 @@ Do not modify directly.* * com.microsoft.QOrderedLongformerAttention * com.microsoft.QOrderedMatMul * com.microsoft.QuantizeBFP + * com.microsoft.QuantizeDequantizeLinear * com.microsoft.QuantizeLinear * com.microsoft.QuantizeWithOrder * com.microsoft.QuickGelu @@ -7039,6 +7040,46 @@ This version of the operator has been available since version 1 of the 'com.micr +### **com.microsoft.QuantizeDequantizeLinear** + + Computes DequantizeLinear(QuantizeLinear(x, y_scale, y_zero_point), y_scale, y_zero_point) with per-tensor + quantization parameters: y = (saturate(round(x / y_scale) + y_zero_point) - y_zero_point) * y_scale, where round is + round half to even and saturate clamps to the range of the zero point type. The quantized tensor is not materialized. + The CPU EP graph optimizer creates this operator from a QuantizeLinear whose output reaches only DequantizeLinear + nodes with the same scale and zero point, directly or through data-movement ops such as Transpose or Reshape. + +#### Version + +This version of the operator has been available since version 1 of the 'com.microsoft' operator set. + +#### Inputs + +
+
x : T1
+
N-D full precision input tensor.
+
y_scale : T1
+
Scalar quantization scale.
+
y_zero_point : T2
+
Scalar quantization zero point. Its type is the quantized type.
+
+ +#### Outputs + +
+
y : T1
+
N-D full precision output tensor with the same shape as 'x'.
+
+ +#### Type Constraints + +
+
T1 : tensor(float)
+
Constrain 'x', 'y_scale' and 'y' to float tensors.
+
T2 : tensor(int16), tensor(uint16)
+
Constrain 'y_zero_point' to 16-bit integer tensors.
+
+ + ### **com.microsoft.QuantizeLinear** The linear quantization operator. It consumes a full precision data, a scale, a zero point to compute the low precision / quantized tensor. diff --git a/docs/OperatorKernels.md b/docs/OperatorKernels.md index 0c07221078719..8258811b2f6c8 100644 --- a/docs/OperatorKernels.md +++ b/docs/OperatorKernels.md @@ -638,6 +638,7 @@ The **OpSet Version** column uses the following notation: |QLinearSoftmax|*in* X:**T**
*in* X_scale:**tensor(float)**
*in* x_zero_point:**T**
*in* y_scale:**tensor(float)**
*in* y_zero_point:**T**
*out* Y:**T**|1+|**T** = tensor(int8), tensor(uint8)| |QLinearWhere|*in* condition:**B**
*in* X:**T**
*in* x_scale:**TF**
*in* x_zero_point:**T**
*in* Y:**T**
*in* y_scale:**TF**
*in* y_zero_point:**T**
*in* z_scale:**TF**
*in* z_zero_point:**T**
*out* Z:**T**|1+|**T** = tensor(int8), tensor(uint8)| |QMoE|*in* input:**T**
*in* router_probs:**T**
*in* fc1_experts_weights:**T1**
*in* fc1_scales:**T2**
*in* fc1_experts_bias:**T**
*in* fc2_experts_weights:**T1**
*in* fc2_scales:**T2**
*in* fc2_experts_bias:**T**
*in* fc3_experts_weights:**T1**
*in* fc3_scales:**T2**
*in* fc3_experts_bias:**T**
*in* fc1_zero_points:**T1**
*in* fc2_zero_points:**T1**
*in* fc3_zero_points:**T1**
*in* router_weights:**T**
*in* fc1_global_scale:**T4**
*in* fc2_global_scale:**T4**
*in* fc1_act_scale:**T4**
*in* fc2_act_scale:**T4**
*in* fc1_act_block_scale:**T2**
*in* fc2_act_block_scale:**T2**
*out* output:**T**|1+|**T** = tensor(float), tensor(float16)
**T1** = tensor(uint8)
**T2** = tensor(float), tensor(float16)| +|QuantizeDequantizeLinear|*in* x:**T1**
*in* y_scale:**T1**
*in* y_zero_point:**T2**
*out* y:**T1**|1+|**T1** = tensor(float)
**T2** = tensor(int16), tensor(uint16)| |QuantizeLinear|*in* x:**T1**
*in* y_scale:**T1**
*in* y_zero_point:**T2**
*out* y:**T2**|1+|**T1** = tensor(float)
**T2** = tensor(int16), tensor(int4), tensor(int8), tensor(uint16), tensor(uint4), tensor(uint8)| |QuickGelu|*in* X:**T**
*out* Y:**T**|1+|**T** = tensor(float)| |Range|*in* start:**T**
*in* limit:**T**
*in* delta:**T**
*out* Y:**T**|1+|**T** = tensor(double), tensor(float), tensor(int16), tensor(int32), tensor(int64)| diff --git a/onnxruntime/contrib_ops/cpu/cpu_contrib_kernels.cc b/onnxruntime/contrib_ops/cpu/cpu_contrib_kernels.cc index 194c4dafbe518..7e13ea9fc0792 100644 --- a/onnxruntime/contrib_ops/cpu/cpu_contrib_kernels.cc +++ b/onnxruntime/contrib_ops/cpu/cpu_contrib_kernels.cc @@ -137,6 +137,8 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, int16_t, QuantizeLinear); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, UInt4x2, QuantizeLinear); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, Int4x2, QuantizeLinear); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, uint16_t, QuantizeDequantizeLinear); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, int16_t, QuantizeDequantizeLinear); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, uint8_t, QLinearLeakyRelu); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, int8_t, QLinearLeakyRelu); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, uint8_t, QLinearSigmoid); @@ -309,6 +311,10 @@ Status RegisterQuantizationKernels(KernelRegistry& kernel_registry) { QuantizeLinear)>, BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, diff --git a/onnxruntime/contrib_ops/cpu/quantization/quantize_dequantize_linear.cc b/onnxruntime/contrib_ops/cpu/quantization/quantize_dequantize_linear.cc new file mode 100644 index 0000000000000..da2b9dbb0e2fa --- /dev/null +++ b/onnxruntime/contrib_ops/cpu/quantization/quantize_dequantize_linear.cc @@ -0,0 +1,75 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include + +#include "core/common/narrow.h" +#include "core/framework/op_kernel.h" +#include "core/mlas/inc/mlas.h" +#include "core/platform/threadpool.h" +#include "core/providers/common.h" + +namespace onnxruntime { +namespace contrib { + +// Bitwise identical to the CPU QuantizeLinear kernel followed by the CPU DequantizeLinear kernel. +template +class QuantizeDequantizeLinear final : public OpKernel { + public: + explicit QuantizeDequantizeLinear(const OpKernelInfo& info) : OpKernel(info) {} + + Status Compute(OpKernelContext* context) const override; +}; + +template +Status QuantizeDequantizeLinear::Compute(OpKernelContext* context) const { + const Tensor& x = *context->Input(0); + const Tensor& y_scale = *context->Input(1); + const Tensor& y_zero_point = *context->Input(2); + ORT_RETURN_IF_NOT(IsScalarOr1ElementVector(&y_scale) && IsScalarOr1ElementVector(&y_zero_point), + "QuantizeDequantizeLinear requires a scalar y_scale and y_zero_point."); + + Tensor& y = *context->Output(0, x.Shape()); + const std::ptrdiff_t count = narrow(x.Shape().Size()); + if (count == 0) { + return Status::OK(); + } + + const float* input = x.Data(); + float* output = y.MutableData(); + const float scale = *y_scale.Data(); + const T zero_point = *y_zero_point.Data(); + + // Same partitioning as ParQuantizeLinearStd, so every MLAS call starts on the same 128-element boundary as in + // QuantizeLinear; MlasQuantizeDequantizeLinear matches MlasQuantizeLinear over the same range. + constexpr std::ptrdiff_t kBlockSize = 128; + const std::ptrdiff_t num_blocks = (count + kBlockSize - 1) / kBlockSize; + const TensorOpCost unit_cost{static_cast(kBlockSize * sizeof(float)), + static_cast(kBlockSize * sizeof(float)), + static_cast(kBlockSize) * 3.0}; + concurrency::ThreadPool::TryParallelFor( + context->GetOperatorThreadPool(), num_blocks, unit_cost, [&](std::ptrdiff_t begin, std::ptrdiff_t end) { + const std::ptrdiff_t start = begin * kBlockSize; + const std::ptrdiff_t stop = std::min(count, end * kBlockSize); + MlasQuantizeDequantizeLinear(input + start, output + start, static_cast(stop - start), scale, + zero_point); + }); + + return Status::OK(); +} + +#define REGISTER_QUANTIZE_DEQUANTIZE_LINEAR(T) \ + ONNX_CPU_OPERATOR_TYPED_MS_KERNEL( \ + QuantizeDequantizeLinear, \ + 1, \ + T, \ + KernelDefBuilder() \ + .TypeConstraint("T1", DataTypeImpl::GetTensorType()) \ + .TypeConstraint("T2", DataTypeImpl::GetTensorType()), \ + QuantizeDequantizeLinear); + +REGISTER_QUANTIZE_DEQUANTIZE_LINEAR(uint16_t) +REGISTER_QUANTIZE_DEQUANTIZE_LINEAR(int16_t) + +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/core/graph/contrib_ops/ms_opset.h b/onnxruntime/core/graph/contrib_ops/ms_opset.h index 7bfbc49c762c0..7a96d4df2c733 100644 --- a/onnxruntime/core/graph/contrib_ops/ms_opset.h +++ b/onnxruntime/core/graph/contrib_ops/ms_opset.h @@ -34,6 +34,7 @@ class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QLinearReduceMean); class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QLinearSigmoid); class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QLinearSoftmax); class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QuantizeLinear); +class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QuantizeDequantizeLinear); class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QuantizeBFP); class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, ReduceSumInteger); class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QOrderedMatMul); @@ -174,6 +175,7 @@ class OpSet_Microsoft_ver1 { fn(GetOpSchema()); fn(GetOpSchema()); fn(GetOpSchema()); + fn(GetOpSchema()); fn(GetOpSchema()); fn(GetOpSchema()); fn(GetOpSchema()); diff --git a/onnxruntime/core/graph/contrib_ops/quantization_defs.cc b/onnxruntime/core/graph/contrib_ops/quantization_defs.cc index 70cd688832b02..02492c0a78bb8 100644 --- a/onnxruntime/core/graph/contrib_ops/quantization_defs.cc +++ b/onnxruntime/core/graph/contrib_ops/quantization_defs.cc @@ -225,6 +225,29 @@ ONNX_MS_OPERATOR_SET_SCHEMA(DequantizeLinear, 1, updateOutputShape(ctx, 0, input_shape); })); +static const char* QuantizeDequantizeLinear_ver1_doc = R"DOC( +Computes DequantizeLinear(QuantizeLinear(x, y_scale, y_zero_point), y_scale, y_zero_point) with per-tensor +quantization parameters: y = (saturate(round(x / y_scale) + y_zero_point) - y_zero_point) * y_scale, where round is +round half to even and saturate clamps to the range of the zero point type. The quantized tensor is not materialized. +The CPU EP graph optimizer creates this operator from a QuantizeLinear whose output reaches only DequantizeLinear +nodes with the same scale and zero point, directly or through data-movement ops such as Transpose or Reshape.)DOC"; + +ONNX_MS_OPERATOR_SET_SCHEMA( + QuantizeDequantizeLinear, 1, + OpSchema() + .Input(0, "x", "N-D full precision input tensor.", "T1") + .Input(1, "y_scale", "Scalar quantization scale.", "T1") + .Input(2, "y_zero_point", "Scalar quantization zero point. Its type is the quantized type.", "T2") + .Output(0, "y", "N-D full precision output tensor with the same shape as 'x'.", "T1") + .TypeConstraint("T1", {"tensor(float)"}, "Constrain 'x', 'y_scale' and 'y' to float tensors.") + .TypeConstraint("T2", {"tensor(int16)", "tensor(uint16)"}, + "Constrain 'y_zero_point' to 16-bit integer tensors.") + .SetDoc(QuantizeDequantizeLinear_ver1_doc) + .TypeAndShapeInferenceFunction([](ONNX_NAMESPACE::InferenceContext& ctx) { + propagateElemTypeFromInputToOutput(ctx, 0, 0); + propagateShapeFromInputToOutput(ctx, 0, 0); + })); + static const char* QuantizeBFP_ver1_doc = R"DOC( The BFP quantization operator. It consumes a full precision tensor and computes an BFP tensor. More documentation on the BFP format can be found in this paper: https://www.microsoft.com/en-us/research/publication/pushing-the-limits-of-narrow-precision-inferencing-at-cloud-scale-with-microsoft-floating-point/)DOC"; diff --git a/onnxruntime/core/mlas/inc/mlas.h b/onnxruntime/core/mlas/inc/mlas.h index 66d02d0d19c8c..ca034342d1cf0 100644 --- a/onnxruntime/core/mlas/inc/mlas.h +++ b/onnxruntime/core/mlas/inc/mlas.h @@ -1448,6 +1448,22 @@ MlasQuantizeLinear( OutputType ZeroPoint ); +// +// QuantizeLinear followed by DequantizeLinear with the same per-tensor parameters, for 16-bit QuantType. +// Bitwise equal to MlasQuantizeLinear over the same range followed by (Quantized - ZeroPoint) * Scale. +// + +template +void +MLASCALL +MlasQuantizeDequantizeLinear( + const float* Input, + float* Output, + size_t N, + float Scale, + QuantType ZeroPoint + ); + void MLASCALL MlasQuantizeLinearU4( diff --git a/onnxruntime/core/mlas/lib/quantize.cpp b/onnxruntime/core/mlas/lib/quantize.cpp index 0c4ecfa8bd0ff..4b46847d5bb25 100644 --- a/onnxruntime/core/mlas/lib/quantize.cpp +++ b/onnxruntime/core/mlas/lib/quantize.cpp @@ -1003,6 +1003,147 @@ MlasQuantizeLinearU4( #endif +template +void +MlasQuantizeDequantizeLinearTwoPass( + const float* Input, + float* Output, + size_t N, + float Scale, + QuantType ZeroPoint + ) +{ + // A multiple of 128 so each MlasQuantizeLinear call starts where QuantizeLinear's 128-element split would. + constexpr size_t ChunkSize = 512; + QuantType Quantized[ChunkSize]; + const int32_t ZeroPointValue = ZeroPoint; + + while (N > 0) { + const size_t Count = std::min(N, ChunkSize); + MlasQuantizeLinear(Input, Quantized, Count, Scale, ZeroPoint); + for (size_t n = 0; n < Count; n++) { + Output[n] = static_cast(static_cast(Quantized[n]) - ZeroPointValue) * Scale; + } + Input += Count; + Output += Count; + N -= Count; + } +} + +#if defined(MLAS_NEON64_INTRINSICS) || defined(MLAS_SSE2_INTRINSICS) || \ + defined(MLAS_LSX_INTRINSICS) + +template +void +MlasQuantizeDequantizeLinearKernel( + const float* Input, + float* Output, + size_t N, + float Scale, + QuantType ZeroPoint + ) +/*++ + +Routine Description: + + This routine runs the per-element steps of MlasQuantizeLinearKernel and then + computes (Quantized - ZeroPoint) * Scale, without storing the quantized values. + +--*/ +{ + constexpr int32_t MinimumValue = std::numeric_limits::lowest(); + constexpr int32_t MaximumValue = std::numeric_limits::max(); + + auto ScaleVector = MlasBroadcastFloat32x4(Scale); + auto MinimumValueVector = MlasBroadcastFloat32x4(float(MinimumValue - ZeroPoint)); + auto MaximumValueVector = MlasBroadcastFloat32x4(float(MaximumValue - ZeroPoint)); + auto ZeroPointVector = MlasBroadcastInt32x4(ZeroPoint); + + while (N >= 4) { + + auto IntegerVector = MlasQuantizeLinearVector(MlasLoadFloat32x4(Input), ScaleVector, + MinimumValueVector, MaximumValueVector, ZeroPointVector); + IntegerVector = MlasSubtractInt32x4(IntegerVector, ZeroPointVector); + MlasStoreFloat32x4(Output, MlasMultiplyFloat32x4(MlasCastToFloat32x4(IntegerVector), ScaleVector)); + + Input += 4; + Output += 4; + N -= 4; + } + + for (size_t n = 0; n < N; n++) { + +#if defined(MLAS_NEON64_INTRINSICS) + auto FloatVector = vld1q_dup_f32(Input + n); +#elif defined(MLAS_LSX_INTRINSICS) + MLAS_FLOAT32X4 FloatVector = (MLAS_FLOAT32X4)__lsx_vldrepl_w(Input+n, 0); +#else + auto FloatVector = _mm_load_ss(Input + n); +#endif + auto IntegerVector = MlasQuantizeLinearVector(FloatVector, ScaleVector, + MinimumValueVector, MaximumValueVector, ZeroPointVector); + IntegerVector = MlasSubtractInt32x4(IntegerVector, ZeroPointVector); + MlasStoreLaneFloat32x4<0>(&Output[n], + MlasMultiplyFloat32x4(MlasCastToFloat32x4(IntegerVector), ScaleVector)); + } +} + +#endif + +template +void +MLASCALL +MlasQuantizeDequantizeLinear( + const float* Input, + float* Output, + size_t N, + float Scale, + QuantType ZeroPoint + ) +{ +#if defined(MLAS_NEON64_INTRINSICS) || defined(MLAS_SSE2_INTRINSICS) || \ + defined(MLAS_LSX_INTRINSICS) +#if defined(MLAS_TARGET_AMD64) + // The single-pass kernel mirrors the default quantize kernel, so use two passes if another one is selected. + bool IsDefaultQuantizeKernel; + if constexpr (std::is_same_v) { + IsDefaultQuantizeKernel = GetMlasPlatform().QuantizeLinearU16Kernel == MlasQuantizeLinearU16Kernel; + } else { + IsDefaultQuantizeKernel = GetMlasPlatform().QuantizeLinearS16Kernel == MlasQuantizeLinearS16Kernel; + } + if (!IsDefaultQuantizeKernel) { + MlasQuantizeDequantizeLinearTwoPass(Input, Output, N, Scale, ZeroPoint); + return; + } +#endif + MlasQuantizeDequantizeLinearKernel(Input, Output, N, Scale, ZeroPoint); +#else + MlasQuantizeDequantizeLinearTwoPass(Input, Output, N, Scale, ZeroPoint); +#endif +} + +template +void +MLASCALL +MlasQuantizeDequantizeLinear( + const float* Input, + float* Output, + size_t N, + float Scale, + uint16_t ZeroPoint + ); + +template +void +MLASCALL +MlasQuantizeDequantizeLinear( + const float* Input, + float* Output, + size_t N, + float Scale, + int16_t ZeroPoint + ); + #if defined(MLAS_SSE2_INTRINSICS) template diff --git a/onnxruntime/core/optimizer/graph_transformer_utils.cc b/onnxruntime/core/optimizer/graph_transformer_utils.cc index 64372c52e1a20..d197eb199d1cb 100644 --- a/onnxruntime/core/optimizer/graph_transformer_utils.cc +++ b/onnxruntime/core/optimizer/graph_transformer_utils.cc @@ -79,6 +79,7 @@ #include "core/optimizer/qdq_transformer/ensure_unique_dq_for_node_unit.h" #include "core/optimizer/qdq_transformer/qdq_propagation.h" #include "core/optimizer/qdq_transformer/qdq_s8_to_u8.h" +#include "core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.h" #include "core/optimizer/qdq_transformer/relu_quantizelinear.h" #include "core/optimizer/quick_gelu_fusion.h" #include "core/optimizer/relu_clip_fusion.h" @@ -525,6 +526,11 @@ InlinedVector> GenerateTransformers( // PR #6351 implemented similar fusion-pattern for CUDA only, and can only fuse conv-add-relu, // while we can fuse more activation. transformers.emplace_back(std::make_unique(cpu_ep)); + + // Runs after the Level2 Q/DQ node-unit fusions have converged, so it cannot hide a Q/DQ they could still use. + if (!disable_quant_qdq) { + transformers.emplace_back(std::make_unique(cpu_ep)); + } #else ORT_UNUSED_PARAMETER(logger); #endif diff --git a/onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.cc b/onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.cc new file mode 100644 index 0000000000000..878cd1b7b6afc --- /dev/null +++ b/onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.cc @@ -0,0 +1,215 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.h" + +#include +#include +#include +#include + +#include "core/graph/constants.h" +#include "core/graph/graph_utils.h" +#include "core/optimizer/qdq_transformer/qdq_util.h" + +namespace onnxruntime { +namespace { + +// Ops that only move or select elements of their first input, so they give the same values on dequantized data. +constexpr std::array kDataMovementOps = { + "DepthToSpace", "Expand", "Flatten", "Gather", "Identity", "Reshape", "Slice", + "SpaceToDepth", "Split", "Squeeze", "Tile", "Transpose", "Unsqueeze"}; + +int32_t ElemType(const NodeArg& arg) { + const auto* type = arg.TypeAsProto(); + return type != nullptr && type->has_tensor_type() ? type->tensor_type().elem_type() + : ONNX_NAMESPACE::TensorProto_DataType_UNDEFINED; +} + +bool HasNoBlockSize(const Node& node) { + const auto* block_size = graph_utils::GetNodeAttribute(node, "block_size"); + return block_size == nullptr || block_size->i() == 0; +} + +bool IsDataMovementOp(const Node& node) { + return node.Domain() == kOnnxDomain && + std::find(kDataMovementOps.begin(), kDataMovementOps.end(), node.OpType()) != kDataMovementOps.end(); +} + +bool IsFusableQ(const Graph& graph, const Node& q_node, const QDQ::GetConstantInitializerFn& get_constant_initializer, + const InlinedHashSet& compatible_eps) { + if (!QDQ::MatchQNode(q_node) || !graph_utils::IsSupportedProvider(q_node, compatible_eps) || + !HasNoBlockSize(q_node) || q_node.GetOutputEdgesCount() == 0 || graph.NodeProducesGraphOutput(q_node)) { + return false; + } + + const auto& inputs = q_node.InputDefs(); + if (inputs.size() != QDQ::InputIndex::TOTAL_COUNT || + ElemType(*inputs[QDQ::InputIndex::INPUT_ID]) != ONNX_NAMESPACE::TensorProto_DataType_FLOAT || + ElemType(*inputs[QDQ::InputIndex::SCALE_ID]) != ONNX_NAMESPACE::TensorProto_DataType_FLOAT) { + return false; + } + + const int32_t quantized_type = ElemType(*inputs[QDQ::InputIndex::ZERO_POINT_ID]); + if (quantized_type != ONNX_NAMESPACE::TensorProto_DataType_UINT16 && + quantized_type != ONNX_NAMESPACE::TensorProto_DataType_INT16) { + return false; + } + + // A positive scale that compares equal to the DQ scale is also bitwise equal to it. + return QDQ::IsQOrDQScalePositiveConstantScalar(graph, q_node, get_constant_initializer, graph.ModelPath()); +} + +struct DQLeaf { + Node* dq_node; + NodeIndex producer; + int producer_output; +}; + +struct QuantizedTree { + // The DQ that produces a graph output, if any, is first. + InlinedVector dq_leaves; + // Each node's quantized input comes from the Q node or an earlier entry. + InlinedVector data_movement_nodes; +}; + +// Succeeds if every use of the Q output reaches a DQ with the same scale and zero point, directly or through +// data-movement ops. Only a direct DQ may produce a graph output, since the fused node takes over its output. +std::optional CollectQuantizedTree(Graph& graph, const Node& q_node, + const QDQ::GetConstantInitializerFn& get_constant_initializer) { + const int32_t quantized_type = ElemType(*q_node.InputDefs()[QDQ::InputIndex::ZERO_POINT_ID]); + QuantizedTree tree; + bool has_graph_output_dq = false; + InlinedVector producers{&q_node}; + for (size_t i = 0; i < producers.size(); ++i) { + const Node& producer = *producers[i]; + for (auto it = producer.OutputEdgesBegin(), end = producer.OutputEdgesEnd(); it != end; ++it) { + Node* consumer = graph.GetNode(it->GetNode().Index()); + if (it->GetDstArgIndex() != 0 || consumer->GetExecutionProviderType() != q_node.GetExecutionProviderType()) { + return std::nullopt; + } + + if (QDQ::MatchDQNode(*consumer)) { + if (!HasNoBlockSize(*consumer) || + ElemType(*consumer->OutputDefs()[0]) != ONNX_NAMESPACE::TensorProto_DataType_FLOAT || + !QDQ::IsQDQPairSupported(graph, q_node, *consumer, get_constant_initializer, graph.ModelPath(), false)) { + return std::nullopt; + } + + const DQLeaf leaf{consumer, producer.Index(), it->GetSrcArgIndex()}; + if (graph.NodeProducesGraphOutput(*consumer)) { + if (&producer != &q_node || has_graph_output_dq) { + return std::nullopt; + } + has_graph_output_dq = true; + tree.dq_leaves.insert(tree.dq_leaves.begin(), leaf); + } else { + tree.dq_leaves.push_back(leaf); + } + continue; + } + + if (!IsDataMovementOp(*consumer) || graph.NodeProducesGraphOutput(*consumer)) { + return std::nullopt; + } + for (const NodeArg* output : consumer->OutputDefs()) { + if (output->Exists() && ElemType(*output) != quantized_type) { + return std::nullopt; + } + } + tree.data_movement_nodes.push_back(consumer); + producers.push_back(consumer); + } + } + + return tree; +} + +NodeArg& MakeFloatNodeArg(Graph& graph, const NodeArg& like) { + ONNX_NAMESPACE::TypeProto type = *like.TypeAsProto(); + type.mutable_tensor_type()->set_elem_type(ONNX_NAMESPACE::TensorProto_DataType_FLOAT); + return graph.GetOrCreateNodeArg(graph.GenerateNodeArgName(like.Name()), &type); +} + +} // namespace + +Status QuantizeDequantizeLinearFusion::ApplyImpl(Graph& graph, bool& modified, int graph_level, + const logging::Logger& logger) const { + const auto get_constant_initializer = [&graph](const std::string& name) { + return graph.GetConstantInitializer(name, true); + }; + + GraphViewer graph_viewer(graph); + for (NodeIndex node_index : graph_viewer.GetNodesInTopologicalOrder()) { + Node* q_node = graph.GetNode(node_index); + if (q_node == nullptr) { + continue; + } + + ORT_RETURN_IF_ERROR(Recurse(*q_node, modified, graph_level, logger)); + + if (!IsFusableQ(graph, *q_node, get_constant_initializer, GetCompatibleExecutionProviders())) { + continue; + } + + std::optional tree = CollectQuantizedTree(graph, *q_node, get_constant_initializer); + if (!tree.has_value() || tree->dq_leaves.empty()) { + continue; + } + + // Reuse the output of a direct DQ (the graph output, if any) so its consumers keep their input. + const DQLeaf& first = tree->dq_leaves.front(); + auto& q_inputs = q_node->MutableInputDefs(); + NodeArg& fused_output = first.producer == q_node->Index() ? *first.dq_node->MutableOutputDefs()[0] + : MakeFloatNodeArg(graph, *q_inputs[0]); + Node& fused = graph.AddNode(graph.GenerateNodeName(q_node->Name() + "_QuantizeDequantizeLinear"), + "QuantizeDequantizeLinear", "Fused QuantizeLinear and DequantizeLinear", + {q_inputs[0], q_inputs[1], q_inputs[2]}, {&fused_output}, nullptr, kMSDomain); + fused.SetExecutionProviderType(q_node->GetExecutionProviderType()); + + // Disconnect the Q and DQ nodes while the edges still match the quantized node args. + InlinedVector source_nodes{q_node->Index()}; + graph_utils::MoveAllNodeInputEdges(graph, *q_node, fused); + graph_utils::RemoveNodeOutputEdges(graph, *q_node); + for (const DQLeaf& leaf : tree->dq_leaves) { + if (leaf.producer != q_node->Index()) { + graph.RemoveEdge(leaf.producer, leaf.dq_node->Index(), leaf.producer_output, 0); + } + source_nodes.push_back(leaf.dq_node->Index()); + } + + // The data-movement nodes now run on the dequantized float values. + InlinedHashMap float_args{{q_node->OutputDefs()[0], &fused_output}}; + for (Node* node : tree->data_movement_nodes) { + NodeArg*& input = node->MutableInputDefs()[0]; + const bool reads_q_output = input == q_node->OutputDefs()[0]; + input = float_args.at(input); + for (NodeArg*& output : node->MutableOutputDefs()) { + if (output->Exists()) { + NodeArg& float_output = MakeFloatNodeArg(graph, *output); + float_args.emplace(output, &float_output); + output = &float_output; + } + } + if (reads_q_output) { + graph.AddEdge(fused.Index(), node->Index(), 0, 0); + } + } + + for (const DQLeaf& leaf : tree->dq_leaves) { + Node& producer = leaf.producer == q_node->Index() ? fused : *graph.GetNode(leaf.producer); + graph_utils::ReplaceDownstreamNodeInput(graph, *leaf.dq_node, 0, producer, leaf.producer_output); + } + + for (NodeIndex index : source_nodes) { + graph.RemoveNode(index); + } + graph.NotifyNodeReplacement(source_nodes, fused.Index()); + + modified = true; + } + + return Status::OK(); +} + +} // namespace onnxruntime diff --git a/onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.h b/onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.h new file mode 100644 index 0000000000000..b8e8f31d396b5 --- /dev/null +++ b/onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.h @@ -0,0 +1,30 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include "core/optimizer/graph_transformer.h" + +namespace onnxruntime { + +/** + @Class QuantizeDequantizeLinearFusion + + Replaces a 16-bit QuantizeLinear whose output reaches only DequantizeLinear nodes with the same constant scalar + scale and zero point, directly or through data-movement ops such as Transpose or Reshape, by a single + com.microsoft.QuantizeDequantizeLinear node. The data-movement ops then run on the float values. The result is + bitwise identical, but the quantized tensors and the extra nodes and kernel launches are removed. + + It must run after the transformers that fuse Q/DQ node units, since it removes the DQ/Q nodes they match. +*/ +class QuantizeDequantizeLinearFusion : public GraphTransformer { + public: + explicit QuantizeDequantizeLinearFusion( + const InlinedHashSet& compatible_execution_providers = {}) noexcept + : GraphTransformer("QuantizeDequantizeLinearFusion", compatible_execution_providers) {} + + private: + Status ApplyImpl(Graph& graph, bool& modified, int graph_level, const logging::Logger& logger) const override; +}; + +} // namespace onnxruntime diff --git a/onnxruntime/test/contrib_ops/quantize_dequantize_linear_op_test.cc b/onnxruntime/test/contrib_ops/quantize_dequantize_linear_op_test.cc new file mode 100644 index 0000000000000..46817770d2c21 --- /dev/null +++ b/onnxruntime/test/contrib_ops/quantize_dequantize_linear_op_test.cc @@ -0,0 +1,87 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include +#include +#include +#include +#include + +#include "gtest/gtest.h" +#include "test/common/random_generator.h" +#include "test/providers/provider_test_utils.h" + +namespace onnxruntime { +namespace test { + +namespace { + +// Q then DQ with round half to even, written independently of the kernels. +template +std::vector ReferenceQuantizeDequantize(const std::vector& x, float scale, T zero_point) { + std::vector y; + y.reserve(x.size()); + const float lo = static_cast(std::numeric_limits::lowest()); + const float hi = static_cast(std::numeric_limits::max()); + for (float v : x) { + const float q = std::min(hi, std::max(lo, std::nearbyintf(v / scale) + static_cast(zero_point))); + y.push_back(static_cast(static_cast(q) - static_cast(zero_point)) * scale); + } + return y; +} + +template +void RunQuantizeDequantizeLinear(const std::vector& dims, const std::vector& x, float scale, + T zero_point, const std::vector& expected) { + OpTester test("QuantizeDequantizeLinear", 1, kMSDomain); + test.AddInput("x", dims, x); + test.AddInput("y_scale", {}, {scale}, true); + test.AddInput("y_zero_point", {}, {zero_point}, true); + test.AddOutput("y", dims, expected); + test.SetOutputTolerance(0.0f, 0.0f); + test.Run(); +} + +} // namespace + +TEST(QuantizeDequantizeLinearOpTest, UInt16RoundsHalfToEvenAndSaturates) { + // x / scale = 0.5, 1.5, 2.5, -0.5, -1.5, and two values far outside [0, 65535] - 32768. + const std::vector x{1.0f, 3.0f, 5.0f, -1.0f, -3.0f, 1.0e6f, -1.0e6f}; + const std::vector expected{0.0f, 4.0f, 4.0f, 0.0f, -4.0f, 65534.0f, -65536.0f}; + EXPECT_EQ(ReferenceQuantizeDequantize(x, 2.0f, 32768), expected); + RunQuantizeDequantizeLinear({7}, x, 2.0f, 32768, expected); +} + +TEST(QuantizeDequantizeLinearOpTest, Int16RoundsHalfToEvenAndSaturates) { + // x / scale = 0.5, 1.5, -1.5, then values beyond [-32768, 32767] + 3. + const std::vector x{0.25f, 0.75f, -0.75f, 1.0e6f, -1.0e6f}; + const std::vector expected{0.0f, 1.0f, -1.0f, 16385.0f, -16382.5f}; + EXPECT_EQ(ReferenceQuantizeDequantize(x, 0.5f, -3), expected); + RunQuantizeDequantizeLinear({5}, x, 0.5f, -3, expected); +} + +TEST(QuantizeDequantizeLinearOpTest, LargeTensorMatchesReference) { + // Not a multiple of the 128-element block, so the parallel split and the tail are both exercised. + const std::vector dims{3, 1037}; + RandomValueGenerator random{optional{1234}}; + const std::vector x = random.Uniform(dims, -16.0f, 16.0f); + RunQuantizeDequantizeLinear(dims, x, 0.0004f, 32768, + ReferenceQuantizeDequantize(x, 0.0004f, 32768)); + RunQuantizeDequantizeLinear(dims, x, 0.0003f, 12, ReferenceQuantizeDequantize(x, 0.0003f, 12)); +} + +TEST(QuantizeDequantizeLinearOpTest, EmptyInput) { + RunQuantizeDequantizeLinear({0, 4}, {}, 0.5f, 7, {}); +} + +TEST(QuantizeDequantizeLinearOpTest, RejectsPerAxisParameters) { + OpTester test("QuantizeDequantizeLinear", 1, kMSDomain); + test.AddInput("x", {2, 2}, {1.0f, 2.0f, 3.0f, 4.0f}); + test.AddInput("y_scale", {2}, {0.5f, 0.25f}); + test.AddInput("y_zero_point", {2}, {0, 0}); + test.AddOutput("y", {2, 2}, {1.0f, 2.0f, 3.0f, 4.0f}); + test.Run(OpTester::ExpectResult::kExpectFailure, "requires a scalar y_scale and y_zero_point"); +} + +} // namespace test +} // namespace onnxruntime diff --git a/onnxruntime/test/mlas/unittest/test_quantizelinear.cpp b/onnxruntime/test/mlas/unittest/test_quantizelinear.cpp index 7c160b6696265..e87ce7cbcb196 100644 --- a/onnxruntime/test/mlas/unittest/test_quantizelinear.cpp +++ b/onnxruntime/test/mlas/unittest/test_quantizelinear.cpp @@ -1,6 +1,8 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include + #include "test_util.h" template @@ -181,6 +183,76 @@ class MlasQuantizeLinear4BitTest : public MlasTestBase { } }; +template +class MlasQuantizeDequantizeLinearTest : public MlasTestBase { + private: + MatrixGuardBuffer BufferInput; + MatrixGuardBuffer BufferOutput; + MatrixGuardBuffer BufferOutputReference; + MatrixGuardBuffer BufferQuantized; + + void Test(size_t N) { + float* Input = BufferInput.GetBuffer(N); + float* Output = BufferOutput.GetBuffer(N); + float* OutputReference = BufferOutputReference.GetBuffer(N); + QuantInt* Quantized = BufferQuantized.GetBuffer(N); + + std::default_random_engine generator(static_cast(N)); + const float Scale = std::uniform_real_distribution(1e-4f, 1e-2f)(generator); + std::uniform_int_distribution zp_distribution(std::numeric_limits::min(), + std::numeric_limits::max()); + const QuantInt ZeroPoint = static_cast(zp_distribution(generator)); + + // Values well past both saturation limits, rounding ties, and non-finite / signed-zero inputs. + const float Range = 70000.0f * Scale; + std::uniform_real_distribution distribution(-Range, Range); + std::uniform_int_distribution tie_distribution(-40000, 40000); + const float Specials[] = {std::numeric_limits::quiet_NaN(), std::numeric_limits::infinity(), + -std::numeric_limits::infinity(), -0.0f, 0.0f, -0.4f * Scale}; + for (size_t n = 0; n < N; n++) { + switch (n % 8) { + case 3: + Input[n] = (static_cast(tie_distribution(generator)) + 0.5f) * Scale; + break; + case 5: + Input[n] = Specials[(n / 8) % (sizeof(Specials) / sizeof(Specials[0]))]; + break; + default: + Input[n] = distribution(generator); + } + } + + MlasQuantizeLinear(Input, Quantized, N, Scale, ZeroPoint); + for (size_t n = 0; n < N; n++) { + OutputReference[n] = static_cast(static_cast(Quantized[n]) - static_cast(ZeroPoint)) * + Scale; + } + MlasQuantizeDequantizeLinear(Input, Output, N, Scale, ZeroPoint); + + for (size_t n = 0; n < N; n++) { + ASSERT_EQ(std::memcmp(&Output[n], &OutputReference[n], sizeof(float)), 0) + << "size=" << N << ", index=" << n << ", input=" << Input[n] << ", output=" << Output[n] + << ", expected=" << OutputReference[n]; + } + } + + public: + static const char* GetTestSuiteName() { + if constexpr (std::is_same_v) { + return "QuantizeDequantizeLinearS16"; + } else { + return "QuantizeDequantizeLinearU16"; + } + } + + void ExecuteShort(void) override { + for (size_t n = 1; n <= 512; n++) { + Test(n); + } + Test(4099); + } +}; + static UNUSED_VARIABLE bool added_to_main = AddTestRegister([](bool is_short_execute) { size_t count = 0; if (is_short_execute) { @@ -190,6 +262,8 @@ static UNUSED_VARIABLE bool added_to_main = AddTestRegister([](bool is_short_exe count += MlasDirectShortExecuteTests>::RegisterShortExecute(); count += MlasDirectShortExecuteTests>::RegisterShortExecute(); count += MlasDirectShortExecuteTests>::RegisterShortExecute(); + count += MlasDirectShortExecuteTests>::RegisterShortExecute(); + count += MlasDirectShortExecuteTests>::RegisterShortExecute(); } return count; }); diff --git a/onnxruntime/test/optimizer/quantize_dequantize_linear_fusion_test.cc b/onnxruntime/test/optimizer/quantize_dequantize_linear_fusion_test.cc new file mode 100644 index 0000000000000..883328597431e --- /dev/null +++ b/onnxruntime/test/optimizer/quantize_dequantize_linear_fusion_test.cc @@ -0,0 +1,388 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include +#include +#include +#include + +#include "core/framework/int4.h" +#include "core/graph/constants.h" +#include "core/graph/node_attr_utils.h" +#include "core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.h" +#include "core/session/onnxruntime_session_options_config_keys.h" + +#include "test/test_environment.h" +#include "test/unittest_util/framework_test_utils.h" +#include "test/unittest_util/graph_transform_test_builder.h" +#include "test/util/include/asserts.h" +#include "test/util/include/inference_session_wrapper.h" + +#include "gtest/gtest.h" + +namespace onnxruntime { +namespace test { + +#if !defined(DISABLE_CONTRIB_OPS) + +namespace { + +constexpr const char* kFusedOp = "com.microsoft.QuantizeDequantizeLinear"; + +template +T TestZeroPoint() { + if constexpr (std::is_signed_v) { + return static_cast(-5); + } else { + return static_cast(32768); + } +} + +// X -> Mul -> Q -> DQ -> Add -> Y. The input range saturates part of the values. +template +std::function BuildInteriorPair(const std::vector& shape) { + return [shape](ModelTestBuilder& builder) { + NodeArg* x = builder.MakeInput(shape, -20.0f, 20.0f); + NodeArg* scaled = builder.MakeIntermediate(); + builder.AddNode("Mul", {x, builder.MakeScalarInitializer(0.75f)}, {scaled}); + NodeArg* q = builder.MakeIntermediate(); + NodeArg* dq = builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(scaled, 0.0004f, TestZeroPoint(), q); + builder.AddDequantizeLinearNode(q, 0.0004f, TestZeroPoint(), dq); + builder.AddNode("Add", {dq, builder.MakeScalarInitializer(0.25f)}, {builder.MakeOutput()}); + }; +} + +void ExpectFused(const Graph& graph, int expected_fused) { + const auto op_count = CountOpsInGraph(graph); + EXPECT_EQ(OpCount(op_count, kFusedOp), expected_fused); + EXPECT_EQ(OpCount(op_count, "QuantizeLinear"), 0); + EXPECT_EQ(OpCount(op_count, "DequantizeLinear"), 0); + for (const auto& node : graph.Nodes()) { + if (node.OpType() == "QuantizeDequantizeLinear") { + EXPECT_EQ(node.GetExecutionProviderType(), kCpuExecutionProvider); + } + } +} + +// Level2 keeps the Q/DQ nodes; Level3 adds the fusion. Outputs must be bitwise equal. +void RunPipeline(const std::function& build, + const std::function& check, + const std::function& add_session_options = {}) { + TransformerTester(build, check, TransformerLevel::Level2, TransformerLevel::Level3, 21, 0.0, 0.0, nullptr, + add_session_options); +} + +Status AssignAllNodesToCpu(Graph& graph) { + for (auto& node : graph.Nodes()) { + node.SetExecutionProviderType(kCpuExecutionProvider); + } + return Status::OK(); +} + +// Applies only the fusion and checks that the graph is unchanged. +void ExpectNotFused(const std::function& build, + const std::function& pre = AssignAllNodesToCpu) { + OpCountMap before; + auto pre_checker = [&](Graph& graph) { + ORT_RETURN_IF_ERROR(pre(graph)); + before = CountOpsInGraph(graph); + return Status::OK(); + }; + auto post_checker = [&](Graph& graph) { + const auto after = CountOpsInGraph(graph); + EXPECT_EQ(OpCount(after, kFusedOp), 0); + EXPECT_EQ(after, before); + return Status::OK(); + }; + ASSERT_STATUS_OK(TestGraphTransformer( + build, 21, DefaultLoggingManager().DefaultLogger(), + std::make_unique(InlinedHashSet{kCpuExecutionProvider}), + TransformerLevel::Level3, 1, pre_checker, post_checker)); +} + +// X -> Mul -> Q16 -> {DQ, Transpose -> DQ, Slice -> Unsqueeze -> DQ, Split -> (DQ, DQ), Gather -> DQ, +// Reshape -> DQ}, each DQ followed by Relu -> output. +std::function BuildDataMovementTree() { + return [](ModelTestBuilder& builder) { + NodeArg* x = builder.MakeInput({2, 6, 40}, -15.0f, 15.0f); + NodeArg* scaled = builder.MakeIntermediate(); + builder.AddNode("Mul", {x, builder.MakeScalarInitializer(0.75f)}, {scaled}); + NodeArg* q = builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(scaled, 0.0004f, static_cast(30000), q); + auto add_dq_relu = [&](NodeArg* quantized) { + NodeArg* dq = builder.MakeIntermediate(); + builder.AddDequantizeLinearNode(quantized, 0.0004f, static_cast(30000), dq); + builder.AddNode("Relu", {dq}, {builder.MakeOutput()}); + }; + auto add_op = [&](const std::string& op, std::vector inputs, size_t num_outputs = 1, + const NodeAttributes* attrs = nullptr) { + std::vector outputs; + for (size_t i = 0; i < num_outputs; ++i) { + outputs.push_back(builder.MakeIntermediate()); + } + builder.AddNode(op, inputs, outputs, "", attrs); + return outputs; + }; + + add_dq_relu(q); + + NodeAttributes perm; + utils::SetNodeAttribute(utils::MakeAttribute("perm", std::vector{0, 2, 1}), perm); + add_dq_relu(add_op("Transpose", {q}, 1, &perm)[0]); + + NodeArg* sliced = add_op("Slice", {q, builder.Make1DInitializer({1}), builder.Make1DInitializer({5}), + builder.Make1DInitializer({1})})[0]; + add_dq_relu(add_op("Unsqueeze", {sliced, builder.Make1DInitializer({0})})[0]); + + NodeAttributes split; + utils::SetNodeAttribute(utils::MakeAttribute("axis", static_cast(2)), split); + utils::SetNodeAttribute(utils::MakeAttribute("num_outputs", static_cast(2)), split); + for (NodeArg* part : add_op("Split", {q}, 2, &split)) { + add_dq_relu(part); + } + + NodeAttributes gather; + utils::SetNodeAttribute(utils::MakeAttribute("axis", static_cast(1)), gather); + add_dq_relu(add_op("Gather", {q, builder.Make1DInitializer({0, 2})}, 1, &gather)[0]); + + add_dq_relu(add_op("Reshape", {q, builder.Make1DInitializer({2, 240})})[0]); + }; +} + +} // namespace + +TEST(QuantizeDequantizeLinearFusionTests, ThroughDataMovementOps) { + auto post = [](Graph& graph) { + ExpectFused(graph, 1); + const auto op_count = CountOpsInGraph(graph); + for (const char* op : {"Transpose", "Slice", "Unsqueeze", "Split", "Gather", "Reshape"}) { + EXPECT_EQ(OpCount(op_count, op), 1) << op; + } + for (const auto& node : graph.Nodes()) { + if (node.OpType() == "Relu" || node.OpType() == "Mul") { + continue; + } + for (const NodeArg* output : node.OutputDefs()) { + EXPECT_EQ(output->TypeAsProto()->tensor_type().elem_type(), ONNX_NAMESPACE::TensorProto_DataType_FLOAT) + << node.OpType() << " " << output->Name(); + } + } + return Status::OK(); + }; + ASSERT_STATUS_OK(TestGraphTransformer( + BuildDataMovementTree(), 21, DefaultLoggingManager().DefaultLogger(), + std::make_unique(InlinedHashSet{kCpuExecutionProvider}), + TransformerLevel::Level3, 1, AssignAllNodesToCpu, post)); + + RunPipeline(BuildDataMovementTree(), [](InferenceSessionWrapper& session) { ExpectFused(session.GetGraph(), 1); }); +} + +TEST(QuantizeDequantizeLinearFusionTests, NotFused_DataMovementOutputsNotOnlyDQ) { + // Q -> Transpose -> X, where X is a uint16 graph output, a DQ graph output, a DQ with another scale, or a Cast. + for (const char* variant : {"quantized output", "dq output", "dq scale", "cast"}) { + SCOPED_TRACE(variant); + ExpectNotFused([variant](ModelTestBuilder& builder) { + const std::string v = variant; + NodeArg* q = builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(builder.MakeInput({2, 8}, -1.0f, 1.0f), 0.001f, + static_cast(32768), q); + NodeArg* dq = builder.MakeIntermediate(); + builder.AddDequantizeLinearNode(q, 0.001f, static_cast(32768), dq); + builder.AddNode("Relu", {dq}, {builder.MakeOutput()}); + + NodeArg* transposed = v == "quantized output" ? builder.MakeOutput() : builder.MakeIntermediate(); + builder.AddNode("Transpose", {q}, {transposed}); + if (v == "dq output" || v == "dq scale") { + NodeArg* dq_out = v == "dq output" ? builder.MakeOutput() : builder.MakeIntermediate(); + builder.AddDequantizeLinearNode(transposed, v == "dq scale" ? 0.002f : 0.001f, + static_cast(32768), dq_out); + if (v == "dq scale") { + builder.AddNode("Relu", {dq_out}, {builder.MakeOutput()}); + } + } else if (v == "cast") { + NodeAttributes to; + utils::SetNodeAttribute(utils::MakeAttribute("to", static_cast(ONNX_NAMESPACE::TensorProto_DataType_FLOAT)), + to); + builder.AddNode("Cast", {transposed}, {builder.MakeOutput()}, "", &to); + } + }); + } +} + +TEST(QuantizeDequantizeLinearFusionTests, InteriorPairIsBitExact) { + // The larger shape is not a multiple of 128 elements and runs on several threads. + for (const std::vector& shape : {std::vector{1, 7}, std::vector{3, 1037}}) { + SCOPED_TRACE(::testing::Message() << "elements=" << shape[0] * shape[1]); + auto check = [](InferenceSessionWrapper& session) { ExpectFused(session.GetGraph(), 1); }; + RunPipeline(BuildInteriorPair(shape), check); + RunPipeline(BuildInteriorPair(shape), check); + } +} + +TEST(QuantizeDequantizeLinearFusionTests, FanOutAndGraphOutputDQ) { + // Q -> {DQ -> Relu -> Y0, DQ -> Sigmoid -> Y1, DQ -> Y2}: one fused node feeds both ops and produces Y2. + auto build = [](ModelTestBuilder& builder) { + NodeArg* x = builder.MakeInput({4, 300}, -15.0f, 15.0f); + NodeArg* q = builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(x, 0.0004f, static_cast(30000), q); + for (const char* op : {"Relu", "Sigmoid"}) { + NodeArg* dq = builder.MakeIntermediate(); + builder.AddDequantizeLinearNode(q, 0.0004f, static_cast(30000), dq); + builder.AddNode(op, {dq}, {builder.MakeOutput()}); + } + builder.AddDequantizeLinearNode(q, 0.0004f, static_cast(30000), builder.MakeOutput()); + }; + auto check = [](InferenceSessionWrapper& session) { + const Graph& graph = session.GetGraph(); + ExpectFused(graph, 1); + const Node* fused = nullptr; + for (const auto& node : graph.Nodes()) { + if (node.OpType() == "QuantizeDequantizeLinear") { + fused = &node; + } + } + ASSERT_NE(fused, nullptr); + EXPECT_TRUE(graph.IsOutput(fused->OutputDefs()[0])); + EXPECT_EQ(fused->GetOutputEdgesCount(), 2u); + EXPECT_TRUE(graph.IsInputsIncludingInitializers(fused->InputDefs()[0])); + }; + RunPipeline(build, check); +} + +TEST(QuantizeDequantizeLinearFusionTests, A16MatMulNBitsBoundaries) { + // X -> Q16 -> DQ16 -> MatMul(DQ(W4)) -> Q16 -> DQ16 -> Y: both A16 boundaries of the MatMulNBits are fused. + auto build = [](ModelTestBuilder& builder) { + NodeArg* x = builder.MakeInput({5, 64}, -1.0f, 1.0f); + NodeArg* xq = builder.MakeIntermediate(); + NodeArg* xdq = builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(x, 0.0004f, static_cast(32768), xq); + builder.AddDequantizeLinearNode(xq, 0.0004f, static_cast(32768), xdq); + + NodeArg* w = builder.MakeInitializer({64, 16}, Int4x2(Int4x2::min_val, 0), Int4x2(Int4x2::max_val, 0)); + NodeArg* w_dq = builder.MakeIntermediate(); + NodeAttributes attrs; + utils::SetNodeAttribute(utils::MakeAttribute("axis", static_cast(1)), attrs); + builder.AddNode("DequantizeLinear", {w, builder.MakeInitializer({16}, 0.01f, 0.02f)}, {w_dq}, "", &attrs); + + NodeArg* mm = builder.MakeIntermediate(); + builder.AddNode("MatMul", {xdq, w_dq}, {mm}); + NodeArg* yq = builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(mm, 0.0005f, static_cast(32768), yq); + builder.AddDequantizeLinearNode(yq, 0.0005f, static_cast(32768), builder.MakeOutput()); + }; + auto check = [](InferenceSessionWrapper& session) { + const auto op_count = CountOpsInGraph(session.GetGraph()); + EXPECT_EQ(OpCount(op_count, "com.microsoft.MatMulNBits"), 1); + EXPECT_EQ(OpCount(op_count, "MatMul"), 0); + ExpectFused(session.GetGraph(), 2); + }; + RunPipeline(build, check); +} + +TEST(QuantizeDequantizeLinearFusionTests, DisabledByDisableQuantQDQ) { + auto check = [](InferenceSessionWrapper& session) { + const auto op_count = CountOpsInGraph(session.GetGraph()); + EXPECT_EQ(OpCount(op_count, kFusedOp), 0); + EXPECT_EQ(OpCount(op_count, "QuantizeLinear"), 1); + EXPECT_EQ(OpCount(op_count, "DequantizeLinear"), 1); + }; + auto options = [](SessionOptions& so) { + ASSERT_STATUS_OK(so.config_options.AddConfigEntry(kOrtSessionOptionsDisableQuantQDQ, "1")); + }; + RunPipeline(BuildInteriorPair({2, 8}), check, options); +} + +TEST(QuantizeDequantizeLinearFusionTests, NotFused_8Bit) { + ExpectNotFused([](ModelTestBuilder& builder) { + NodeArg* q = builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(builder.MakeInput({2, 8}, -1.0f, 1.0f), 0.01f, + static_cast(128), q); + NodeArg* dq = builder.MakeIntermediate(); + builder.AddDequantizeLinearNode(q, 0.01f, static_cast(128), dq); + builder.AddNode("Relu", {dq}, {builder.MakeOutput()}); + }); +} + +TEST(QuantizeDequantizeLinearFusionTests, NotFused_DifferentDQParameters) { + for (bool different_scale : {true, false}) { + SCOPED_TRACE(different_scale ? "scale" : "zero point"); + ExpectNotFused([different_scale](ModelTestBuilder& builder) { + NodeArg* q = builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(builder.MakeInput({2, 8}, -1.0f, 1.0f), 0.001f, + static_cast(32768), q); + NodeArg* dq = builder.MakeIntermediate(); + builder.AddDequantizeLinearNode(q, different_scale ? 0.002f : 0.001f, + static_cast(different_scale ? 32768 : 32767), dq); + builder.AddNode("Relu", {dq}, {builder.MakeOutput()}); + }); + } +} + +TEST(QuantizeDequantizeLinearFusionTests, NotFused_QuantizedValueUsedElsewhere) { + // The Q output is also a graph output, or feeds a non-DQ consumer. + for (bool q_is_graph_output : {true, false}) { + SCOPED_TRACE(q_is_graph_output ? "graph output" : "non-DQ consumer"); + ExpectNotFused([q_is_graph_output](ModelTestBuilder& builder) { + NodeArg* q = q_is_graph_output ? builder.MakeOutput() : builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(builder.MakeInput({2, 8}, -1.0f, 1.0f), 0.001f, + static_cast(32768), q); + if (!q_is_graph_output) { + builder.AddNode("Identity", {q}, {builder.MakeOutput()}); + } + NodeArg* dq = builder.MakeIntermediate(); + builder.AddDequantizeLinearNode(q, 0.001f, static_cast(32768), dq); + builder.AddNode("Relu", {dq}, {builder.MakeOutput()}); + }); + } +} + +TEST(QuantizeDequantizeLinearFusionTests, NotFused_TwoDQGraphOutputs) { + ExpectNotFused([](ModelTestBuilder& builder) { + NodeArg* q = builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(builder.MakeInput({2, 8}, -1.0f, 1.0f), 0.001f, + static_cast(0), q); + builder.AddDequantizeLinearNode(q, 0.001f, static_cast(0), builder.MakeOutput()); + builder.AddDequantizeLinearNode(q, 0.001f, static_cast(0), builder.MakeOutput()); + }); +} + +TEST(QuantizeDequantizeLinearFusionTests, NotFused_NonConstantOrPerAxisParameters) { + // Scale is a graph input. + ExpectNotFused([](ModelTestBuilder& builder) { + NodeArg* scale = builder.MakeInput(std::vector{}, std::vector{0.001f}); + NodeArg* zp = builder.MakeScalarInitializer(32768); + NodeArg* q = builder.MakeIntermediate(); + builder.AddNode("QuantizeLinear", {builder.MakeInput({2, 8}, -1.0f, 1.0f), scale, zp}, {q}); + NodeArg* dq = builder.MakeIntermediate(); + builder.AddNode("DequantizeLinear", {q, scale, zp}, {dq}); + builder.AddNode("Relu", {dq}, {builder.MakeOutput()}); + }); + + // Per-axis scale and zero point. + ExpectNotFused([](ModelTestBuilder& builder) { + NodeAttributes attrs; + utils::SetNodeAttribute(utils::MakeAttribute("axis", static_cast(1)), attrs); + const std::vector scales{0.001f, 0.002f}; + const std::vector zps{32768, 32768}; + NodeArg* q = builder.MakeIntermediate(); + builder.AddQuantizeLinearNode(builder.MakeInput({4, 2}, -1.0f, 1.0f), scales, zps, q, &attrs); + NodeArg* dq = builder.MakeIntermediate(); + builder.AddDequantizeLinearNode(q, scales, zps, dq, &attrs); + builder.AddNode("Relu", {dq}, {builder.MakeOutput()}); + }); +} + +TEST(QuantizeDequantizeLinearFusionTests, NotFused_NonCpuEp) { + ExpectNotFused(BuildInteriorPair({2, 8}), [](Graph& graph) { + for (auto& node : graph.Nodes()) { + node.SetExecutionProviderType(kCudaExecutionProvider); + } + return Status::OK(); + }); +} + +#endif // !defined(DISABLE_CONTRIB_OPS) + +} // namespace test +} // namespace onnxruntime From c6537ea182f3f2f66b61fd7b72085d356c827d6b Mon Sep 17 00:00:00 2001 From: mirounga Date: Thu, 5 Nov 2026 06:29:23 -0800 Subject: [PATCH 2/2] Skip QuantizeLinear with non-float precision in QuantizeDequantizeLinearFusion From opset 23 the CPU QuantizeLinear kernel honors precision=FLOAT16, but the fused kernel always computes in float, so fusing such a node changes the outputs. Only fuse when precision is 0 or FLOAT, and add opset-23 tests for both the rejected and the still-fused cases. --- .../quantize_dequantize_linear_fusion.cc | 14 +++++- .../quantize_dequantize_linear_fusion_test.cc | 45 +++++++++++++++++-- 2 files changed, 54 insertions(+), 5 deletions(-) diff --git a/onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.cc b/onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.cc index 878cd1b7b6afc..6bc4445e73575 100644 --- a/onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.cc +++ b/onnxruntime/core/optimizer/qdq_transformer/quantize_dequantize_linear_fusion.cc @@ -31,6 +31,17 @@ bool HasNoBlockSize(const Node& node) { return block_size == nullptr || block_size->i() == 0; } +// From opset 23 QuantizeLinear can quantize in another precision, but the fused kernel always computes in float. +// Precision 0 means the scale type, which IsFusableQ requires to be float. +bool HasFloatPrecision(const Node& node) { + if (node.SinceVersion() < 23) { + return true; + } + const auto* precision = graph_utils::GetNodeAttribute(node, "precision"); + return precision == nullptr || precision->i() == 0 || + precision->i() == ONNX_NAMESPACE::TensorProto_DataType_FLOAT; +} + bool IsDataMovementOp(const Node& node) { return node.Domain() == kOnnxDomain && std::find(kDataMovementOps.begin(), kDataMovementOps.end(), node.OpType()) != kDataMovementOps.end(); @@ -39,7 +50,8 @@ bool IsDataMovementOp(const Node& node) { bool IsFusableQ(const Graph& graph, const Node& q_node, const QDQ::GetConstantInitializerFn& get_constant_initializer, const InlinedHashSet& compatible_eps) { if (!QDQ::MatchQNode(q_node) || !graph_utils::IsSupportedProvider(q_node, compatible_eps) || - !HasNoBlockSize(q_node) || q_node.GetOutputEdgesCount() == 0 || graph.NodeProducesGraphOutput(q_node)) { + !HasNoBlockSize(q_node) || !HasFloatPrecision(q_node) || q_node.GetOutputEdgesCount() == 0 || + graph.NodeProducesGraphOutput(q_node)) { return false; } diff --git a/onnxruntime/test/optimizer/quantize_dequantize_linear_fusion_test.cc b/onnxruntime/test/optimizer/quantize_dequantize_linear_fusion_test.cc index 883328597431e..15efa63ea7c59 100644 --- a/onnxruntime/test/optimizer/quantize_dequantize_linear_fusion_test.cc +++ b/onnxruntime/test/optimizer/quantize_dequantize_linear_fusion_test.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include +#include #include #include #include @@ -68,8 +69,8 @@ void ExpectFused(const Graph& graph, int expected_fused) { // Level2 keeps the Q/DQ nodes; Level3 adds the fusion. Outputs must be bitwise equal. void RunPipeline(const std::function& build, const std::function& check, - const std::function& add_session_options = {}) { - TransformerTester(build, check, TransformerLevel::Level2, TransformerLevel::Level3, 21, 0.0, 0.0, nullptr, + const std::function& add_session_options = {}, int opset = 21) { + TransformerTester(build, check, TransformerLevel::Level2, TransformerLevel::Level3, opset, 0.0, 0.0, nullptr, add_session_options); } @@ -82,7 +83,7 @@ Status AssignAllNodesToCpu(Graph& graph) { // Applies only the fusion and checks that the graph is unchanged. void ExpectNotFused(const std::function& build, - const std::function& pre = AssignAllNodesToCpu) { + const std::function& pre = AssignAllNodesToCpu, int opset = 21) { OpCountMap before; auto pre_checker = [&](Graph& graph) { ORT_RETURN_IF_ERROR(pre(graph)); @@ -96,11 +97,30 @@ void ExpectNotFused(const std::function& build, return Status::OK(); }; ASSERT_STATUS_OK(TestGraphTransformer( - build, 21, DefaultLoggingManager().DefaultLogger(), + build, opset, DefaultLoggingManager().DefaultLogger(), std::make_unique(InlinedHashSet{kCpuExecutionProvider}), TransformerLevel::Level3, 1, pre_checker, post_checker)); } +// X -> Q (with the given precision attribute, if any) -> DQ -> Relu -> Y, for opset 23+ models. +template +std::function BuildPairWithPrecision(std::optional precision) { + return [precision](ModelTestBuilder& builder) { + NodeArg* x = builder.MakeInput({2, 8}, -30.0f, 30.0f); + NodeArg* q = builder.MakeIntermediate(); + NodeAttributes attrs; + if (precision.has_value()) { + utils::SetNodeAttribute(utils::MakeAttribute("precision", *precision), attrs); + } + NodeArg* scale = builder.MakeScalarInitializer(0.01f); + NodeArg* zero_point = builder.MakeScalarInitializer(TestZeroPoint()); + builder.AddNode("QuantizeLinear", {x, scale, zero_point}, {q}, "", &attrs); + NodeArg* dq = builder.MakeIntermediate(); + builder.AddDequantizeLinearNode(q, 0.01f, TestZeroPoint(), dq); + builder.AddNode("Relu", {dq}, {builder.MakeOutput()}); + }; +} + // X -> Mul -> Q16 -> {DQ, Transpose -> DQ, Slice -> Unsqueeze -> DQ, Split -> (DQ, DQ), Gather -> DQ, // Reshape -> DQ}, each DQ followed by Relu -> output. std::function BuildDataMovementTree() { @@ -373,6 +393,23 @@ TEST(QuantizeDequantizeLinearFusionTests, NotFused_NonConstantOrPerAxisParameter }); } +TEST(QuantizeDequantizeLinearFusionTests, Opset23FloatPrecisionIsFused) { + auto check = [](InferenceSessionWrapper& session) { ExpectFused(session.GetGraph(), 1); }; + for (std::optional precision : {std::optional{}, std::optional{0}, + std::optional{ONNX_NAMESPACE::TensorProto_DataType_FLOAT}}) { + SCOPED_TRACE(::testing::Message() << "precision=" << precision.value_or(-1)); + RunPipeline(BuildPairWithPrecision(precision), check, {}, 23); + RunPipeline(BuildPairWithPrecision(precision), check, {}, 23); + } +} + +TEST(QuantizeDequantizeLinearFusionTests, NotFused_Float16Precision) { + // The fused kernel computes in float and would not match a float16-precision QuantizeLinear. + constexpr int64_t kFloat16 = ONNX_NAMESPACE::TensorProto_DataType_FLOAT16; + ExpectNotFused(BuildPairWithPrecision(kFloat16), AssignAllNodesToCpu, 23); + ExpectNotFused(BuildPairWithPrecision(kFloat16), AssignAllNodesToCpu, 23); +} + TEST(QuantizeDequantizeLinearFusionTests, NotFused_NonCpuEp) { ExpectNotFused(BuildInteriorPair({2, 8}), [](Graph& graph) { for (auto& node : graph.Nodes()) {