[CPU] Fuse 16-bit QuantizeLinear -> DequantizeLinear into QuantizeDeq…

docs/ContribOperators.md

@@ -111,6 +111,7 @@ Do not modify directly.*
   * <a href="#com.microsoft.QOrderedLongformerAttention">com.microsoft.QOrderedLongformerAttention</a>
   * <a href="#com.microsoft.QOrderedMatMul">com.microsoft.QOrderedMatMul</a>
   * <a href="#com.microsoft.QuantizeBFP">com.microsoft.QuantizeBFP</a>
+  * <a href="#com.microsoft.QuantizeDequantizeLinear">com.microsoft.QuantizeDequantizeLinear</a>
   * <a href="#com.microsoft.QuantizeLinear">com.microsoft.QuantizeLinear</a>
   * <a href="#com.microsoft.QuantizeWithOrder">com.microsoft.QuantizeWithOrder</a>
   * <a href="#com.microsoft.QuickGelu">com.microsoft.QuickGelu</a>
@@ -7123,6 +7124,46 @@ This version of the operator has been available since version 1 of the 'com.micr
 </dl>
 
 
+### <a name="com.microsoft.QuantizeDequantizeLinear"></a><a name="com.microsoft.quantizedequantizelinear">**com.microsoft.QuantizeDequantizeLinear**</a>
+
+  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
+
+<dl>
+<dt><tt>x</tt> : T1</dt>
+<dd>N-D full precision input tensor.</dd>
+<dt><tt>y_scale</tt> : T1</dt>
+<dd>Scalar quantization scale.</dd>
+<dt><tt>y_zero_point</tt> : T2</dt>
+<dd>Scalar quantization zero point. Its type is the quantized type.</dd>
+</dl>
+
+#### Outputs
+
+<dl>
+<dt><tt>y</tt> : T1</dt>
+<dd>N-D full precision output tensor with the same shape as 'x'.</dd>
+</dl>
+
+#### Type Constraints
+
+<dl>
+<dt><tt>T1</tt> : tensor(float)</dt>
+<dd>Constrain 'x', 'y_scale' and 'y' to float tensors.</dd>
+<dt><tt>T2</tt> : tensor(int16), tensor(uint16)</dt>
+<dd>Constrain 'y_zero_point' to 16-bit integer tensors.</dd>
+</dl>
+
+
 ### <a name="com.microsoft.QuantizeLinear"></a><a name="com.microsoft.quantizelinear">**com.microsoft.QuantizeLinear**</a>
 
   The linear quantization operator. It consumes a full precision data, a scale, a zero point to compute the low precision / quantized tensor.

docs/OperatorKernels.md

@@ -639,6 +639,7 @@ The **OpSet Version** column uses the following notation:
 |QLinearSoftmax|*in* X:**T**<br> *in* X_scale:**tensor(float)**<br> *in* x_zero_point:**T**<br> *in* y_scale:**tensor(float)**<br> *in* y_zero_point:**T**<br> *out* Y:**T**|1+|**T** = tensor(int8), tensor(uint8)|
 |QLinearWhere|*in* condition:**B**<br> *in* X:**T**<br> *in* x_scale:**TF**<br> *in* x_zero_point:**T**<br> *in* Y:**T**<br> *in* y_scale:**TF**<br> *in* y_zero_point:**T**<br> *in* z_scale:**TF**<br> *in* z_zero_point:**T**<br> *out* Z:**T**|1+|**T** = tensor(int8), tensor(uint8)|
 |QMoE|*in* input:**T**<br> *in* router_probs:**T**<br> *in* fc1_experts_weights:**T1**<br> *in* fc1_scales:**T2**<br> *in* fc1_experts_bias:**T**<br> *in* fc2_experts_weights:**T1**<br> *in* fc2_scales:**T2**<br> *in* fc2_experts_bias:**T**<br> *in* fc3_experts_weights:**T1**<br> *in* fc3_scales:**T2**<br> *in* fc3_experts_bias:**T**<br> *in* fc1_zero_points:**T1**<br> *in* fc2_zero_points:**T1**<br> *in* fc3_zero_points:**T1**<br> *in* router_weights:**T**<br> *in* fc1_global_scale:**T4**<br> *in* fc2_global_scale:**T4**<br> *in* fc1_act_scale:**T4**<br> *in* fc2_act_scale:**T4**<br> *in* fc1_act_block_scale:**T2**<br> *in* fc2_act_block_scale:**T2**<br> *out* output:**T**|1+|**T** = tensor(bfloat16), tensor(float), tensor(float16)<br/> **T1** = tensor(uint8)<br/> **T2** = tensor(bfloat16), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e8m0)<br/> **T4** = tensor(float)|
+|QuantizeDequantizeLinear|*in* x:**T1**<br> *in* y_scale:**T1**<br> *in* y_zero_point:**T2**<br> *out* y:**T1**|1+|**T1** = tensor(float)<br/> **T2** = tensor(int16), tensor(uint16)|
 |QuantizeLinear|*in* x:**T1**<br> *in* y_scale:**T1**<br> *in* y_zero_point:**T2**<br> *out* y:**T2**|1+|**T1** = tensor(float)<br/> **T2** = tensor(int16), tensor(int4), tensor(int8), tensor(uint16), tensor(uint4), tensor(uint8)|
 |QuickGelu|*in* X:**T**<br> *out* Y:**T**|1+|**T** = tensor(float)|
 |Range|*in* start:**T**<br> *in* limit:**T**<br> *in* delta:**T**<br> *out* Y:**T**|1+|**T** = tensor(double), tensor(float), tensor(int16), tensor(int32), tensor(int64)|

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);
@@ -310,6 +312,10 @@ Status RegisterQuantizationKernels(KernelRegistry& kernel_registry) {
                                                                   QuantizeLinear)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, Int4x2,
                                                                   QuantizeLinear)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, uint16_t,
+                                                                  QuantizeDequantizeLinear)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, int16_t,
+                                                                  QuantizeDequantizeLinear)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, uint8_t, QLinearLeakyRelu)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, int8_t, QLinearLeakyRelu)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, uint8_t, QLinearSigmoid)>,

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 <algorithm>
+
+#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 <typename T>
+class QuantizeDequantizeLinear final : public OpKernel {
+ public:
+  explicit QuantizeDequantizeLinear(const OpKernelInfo& info) : OpKernel(info) {}
+
+  Status Compute(OpKernelContext* context) const override;
+};
+
+template <typename T>
+Status QuantizeDequantizeLinear<T>::Compute(OpKernelContext* context) const {
+  const Tensor& x = *context->Input<Tensor>(0);
+  const Tensor& y_scale = *context->Input<Tensor>(1);
+  const Tensor& y_zero_point = *context->Input<Tensor>(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<std::ptrdiff_t>(x.Shape().Size());
+  if (count == 0) {
+    return Status::OK();
+  }
+
+  const float* input = x.Data<float>();
+  float* output = y.MutableData<float>();
+  const float scale = *y_scale.Data<float>();
+  const T zero_point = *y_zero_point.Data<T>();
+
+  // 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<double>(kBlockSize * sizeof(float)),
+                               static_cast<double>(kBlockSize * sizeof(float)),
+                               static_cast<double>(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<size_t>(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<float>()) \
+          .TypeConstraint("T2", DataTypeImpl::GetTensorType<T>()),    \
+      QuantizeDequantizeLinear<T>);
+
+REGISTER_QUANTIZE_DEQUANTIZE_LINEAR(uint16_t)
+REGISTER_QUANTIZE_DEQUANTIZE_LINEAR(int16_t)
+
+}  // namespace contrib
+}  // namespace onnxruntime

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);
@@ -175,6 +176,7 @@ class OpSet_Microsoft_ver1 {
     fn(GetOpSchema<ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QLinearSigmoid)>());
     fn(GetOpSchema<ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QLinearSoftmax)>());
     fn(GetOpSchema<ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QuantizeLinear)>());
+    fn(GetOpSchema<ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QuantizeDequantizeLinear)>());
     fn(GetOpSchema<ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QuantizeBFP)>());
     fn(GetOpSchema<ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, ReduceSumInteger)>());
     fn(GetOpSchema<ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, QOrderedMatMul)>());

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";

onnxruntime/core/mlas/inc/mlas.h

@@ -1458,6 +1458,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<typename QuantType>
+void
+MLASCALL
+MlasQuantizeDequantizeLinear(
+    const float* Input,
+    float* Output,
+    size_t N,
+    float Scale,
+    QuantType ZeroPoint
+    );
+
 void
 MLASCALL
 MlasQuantizeLinearU4(

onnxruntime/core/mlas/lib/quantize.cpp

@@ -1003,6 +1003,147 @@ MlasQuantizeLinearU4(
 
 #endif
 
+template<typename QuantType>
+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<float>(static_cast<int32_t>(Quantized[n]) - ZeroPointValue) * Scale;
+        }
+        Input += Count;
+        Output += Count;
+        N -= Count;
+    }
+}
+
+#if defined(MLAS_NEON64_INTRINSICS) || defined(MLAS_SSE2_INTRINSICS) || \
+    defined(MLAS_LSX_INTRINSICS)
+
+template<typename QuantType>
+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<QuantType>::lowest();
+    constexpr int32_t MaximumValue = std::numeric_limits<QuantType>::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<typename QuantType>
+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<QuantType, uint16_t>) {
+        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<uint16_t>(
+    const float* Input,
+    float* Output,
+    size_t N,
+    float Scale,
+    uint16_t ZeroPoint
+    );
+
+template
+void
+MLASCALL
+MlasQuantizeDequantizeLinear<int16_t>(
+    const float* Input,
+    float* Output,
+    size_t N,
+    float Scale,
+    int16_t ZeroPoint
+    );
+
 #if defined(MLAS_SSE2_INTRINSICS)
 
 template <typename OutputType>

onnxruntime/core/optimizer/graph_transformer_utils.cc

@@ -80,6 +80,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"
@@ -512,6 +513,11 @@ InlinedVector<std::unique_ptr<GraphTransformer>> 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<ConvAddActivationFusion>(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<QuantizeDequantizeLinearFusion>(cpu_ep));
+      }
 #else
       ORT_UNUSED_PARAMETER(logger);
 #endif