ci: restore a red main — refresh the conformance negative control and formatting

.github/workflows/ci.yml

@@ -160,6 +160,15 @@ jobs:
         run: python scripts/benchmark_muse_native_local.py --self-test --require-numpy
       - name: Check formatting
         run: cargo fmt --all -- --check
+      # This lane builds a subset of the workspace, so only that subset's
+      # packages land in the registry cache. The MLAS-free merge gate in
+      # onnx-runtime-ep-cpu-plugin asks `cargo metadata --offline` for the
+      # *whole* workspace resolve, and fails loud rather than skipping when it
+      # cannot read one — so a package no built crate happens to need (via
+      # tokenizers, say) turns that gate red for a reason unrelated to MLAS.
+      # Fetching the locked graph first keeps the gate hermetic and honest.
+      - name: Fetch the locked dependency graph
+        run: cargo fetch --locked
       - name: Build Linux offline crates
         run: >-
           cargo build --locked
@@ -486,6 +495,11 @@ jobs:
           else
             cargo install cargo-llvm-cov --version "$CARGO_LLVM_COV_VERSION" --locked --force
           fi
+      # Same reason as the fast-linux lane: the MLAS-free gate reads the whole
+      # workspace resolve with `cargo metadata --offline`, so every package in
+      # the lockfile has to be present before the subset build starts.
+      - name: Fetch the locked dependency graph
+        run: cargo fetch --locked
       # Never replace these explicit packages with bare `cargo test` or
       # `--workspace`: that selects onnx-genai-ort-sys, whose build script
       # downloads a native ONNX Runtime distribution. The package list is

conformance/run_onnx_tests.py

@@ -44,7 +44,20 @@
     "Gemm",
 ]
 ADDITIONAL_PASS_OPS = ["Abs", "Conv", "Sigmoid"]
-SYNTHETIC_UNSUPPORTED_OP = "SyntheticUnsupportedCelu"
+# A real ONNX op that this EP declines, used as a live negative control: it
+# proves the harness can still *detect* an unsupported operator, which every
+# other case would pass silently if the decline path broke.
+#
+# It names a real op, so it goes stale the moment that op is implemented -- and
+# it did. This was `Celu` until the CPU EP grew a `Celu` kernel, at which point
+# the control started reporting "expected UNSUPPORTED, got PASS" and turned
+# every PR's conformance job red. That is the control working, not failing.
+# When it happens again, do not relax the expectation: move this to another op
+# the EP genuinely declines (the runner must exit 2 with `UNSUPPORTED_OP`, and
+# `onnx.reference` must still be able to evaluate it). `Det` is chosen because
+# it is a linear-algebra op well outside this EP's dispatch and activation
+# roadmap, so it should stay declined far longer than an elementwise op will.
+SYNTHETIC_UNSUPPORTED_OP = "SyntheticUnsupportedDet"
 EXPECTED_STATUS = {
     **{op: "PASS" for op in CPU_OPS + ADDITIONAL_PASS_OPS},
     SYNTHETIC_UNSUPPORTED_OP: "UNSUPPORTED",
@@ -71,8 +84,20 @@ def conformance_failures(
         if result is None:
             failures.append(f"{op}: missing result; expected {expected}")
         elif result.status != expected:
+            hint = ""
+            if op == SYNTHETIC_UNSUPPORTED_OP and result.status == "PASS":
+                # The control is doing its job: an op it assumed we decline is
+                # now implemented. Say so, because "expected UNSUPPORTED, got
+                # PASS" reads like a regression when it is the opposite.
+                hint = (
+                    " -- this EP now implements the op behind the negative"
+                    " control, so the control is stale. Point"
+                    " SYNTHETIC_UNSUPPORTED_OP at an op the EP still declines"
+                    " rather than relaxing this expectation."
+                )
             failures.append(
-                f"{op}: expected {expected}, got {result.status} ({result.detail})"
+                f"{op}: expected {expected}, got {result.status}"
+                f" ({result.detail}){hint}"
             )
     for result in results:
         if result.op not in expected_status:
@@ -219,7 +244,7 @@ def manual_model(op: str) -> onnx.ModelProto:
         "ReduceMean": np.array([[1.0, 3.0], [5.0, 7.0]], dtype=np.float32),
         "Unsqueeze": np.array([[1.0, 2.0]], dtype=np.float32),
         "Conv": np.arange(9, dtype=np.float32).reshape(1, 1, 3, 3),
-        SYNTHETIC_UNSUPPORTED_OP: np.array([0.25, -0.5], dtype=np.float32),
+        SYNTHETIC_UNSUPPORTED_OP: np.array([[1.0, 2.0], [3.0, 4.0]], dtype=np.float32),
     }
     if op == "Constant":
         node = helper.make_node(
@@ -307,12 +332,12 @@ def manual_model(op: str) -> onnx.ModelProto:
         )
     elif op == SYNTHETIC_UNSUPPORTED_OP:
         x = numpy_helper.from_array(values[op], "X")
-        node = helper.make_node("Celu", ["X"], ["Y"])
+        node = helper.make_node("Det", ["X"], ["Y"])
         graph = helper.make_graph(
             [node],
             op,
             [],
-            [helper.make_tensor_value_info("Y", TensorProto.FLOAT, [2])],
+            [helper.make_tensor_value_info("Y", TensorProto.FLOAT, [])],
             [x],
         )
     else:

crates/onnx-runtime-ep-cpu-plugin/tests/default_artifacts_are_mlas_free.rs

@@ -738,10 +738,22 @@ fn gate_must_run() -> bool {
 /// failure mode as measuring a configuration nobody ships. Where the gate is
 /// required, an unusable probe is therefore a failure and not a skip.
 fn gate_probe_failed(why: &str) -> Option<String> {
+    // A lane that builds only a subset of the workspace caches only that
+    // subset's packages, but the resolve probe reads every manifest in the
+    // lockfile. That shows up here as an `--offline` download failure for some
+    // crate with no connection to MLAS, which reads like a mystery unless the
+    // remedy is named at the point of failure.
+    let hint = if why.contains("--offline was specified") {
+        "\nhint: this is a missing package in the local registry, not an MLAS \
+         finding. A lane that builds part of the workspace must run \
+         `cargo fetch --locked` before the probe reads the whole resolve."
+    } else {
+        ""
+    };
     assert!(
         !gate_must_run(),
         "the MLAS-free gate could not run its probe on the configuration it \
-         exists to check ({why}); passing here would prove nothing"
+         exists to check ({why}); passing here would prove nothing{hint}"
     );
     eprintln!("⏭ skipped — {why}");
     None

crates/onnx-runtime-ep-cuda/src/kernels/matmul_nbits.rs

@@ -14159,11 +14159,7 @@ extern "C" __global__ void matmul_nbits_gemv_f16_scales_f16_down_staged_referenc
         );
         assert!(
             use_scales_f16_zp_splitk_gate(ZP_SPLITK_BANDWIDTH_MIN_K, saturated_n, a100)
-                && !use_scales_f16_zp_splitk_gate(
-                    ZP_SPLITK_BANDWIDTH_MIN_K - 1,
-                    saturated_n,
-                    a100
-                ),
+                && !use_scales_f16_zp_splitk_gate(ZP_SPLITK_BANDWIDTH_MIN_K - 1, saturated_n, a100),
             "the K threshold is inclusive at the measured depth and does not \
              extrapolate below it"
         );

crates/onnx-runtime-ep-cuda/src/kernels/normalization.rs

@@ -4706,8 +4706,6 @@ mod tests {
             "tree mean-square error {tree_err:e} exceeds serial error {serial_err:e}"
         );
     }
-
-
 }
 
 #[cfg(test)]

crates/onnx-runtime-ep-cuda/src/optimizer.rs

@@ -749,10 +749,12 @@ impl CudaLinearAttentionGatingFusion {
             (a0, a1, false)
         } else if graph.initializers.contains_key(&a1) {
             (a1, a0, false)
-        } else if let Some((a_log, chain)) = Self::match_neg_exp_chain(graph, a0, mul_id, io_dtype) {
+        } else if let Some((a_log, chain)) = Self::match_neg_exp_chain(graph, a0, mul_id, io_dtype)
+        {
             dead.extend(chain);
             (a_log, a1, true)
-        } else if let Some((a_log, chain)) = Self::match_neg_exp_chain(graph, a1, mul_id, io_dtype) {
+        } else if let Some((a_log, chain)) = Self::match_neg_exp_chain(graph, a1, mul_id, io_dtype)
+        {
             dead.extend(chain);
             (a_log, a0, true)
         } else {
@@ -7306,7 +7308,11 @@ mod tests {
     /// `neg_exp_from_a_log`: when true the decay coefficient is the exported
     /// `Neg(Exp(A_log))` constant chain (A_log the initializer) instead of a
     /// ready `neg_exp_A` initializer, exercising the inline constant-chain fold.
-    fn gated_delta_la_graph_ext(heads: usize, trailing_cast: bool, neg_exp_from_a_log: bool) -> Graph {
+    fn gated_delta_la_graph_ext(
+        heads: usize,
+        trailing_cast: bool,
+        neg_exp_from_a_log: bool,
+    ) -> Graph {
         let mut graph = Graph::new();
         graph.opset_imports.insert(String::new(), 17);
         graph.opset_imports.insert(MICROSOFT_DOMAIN.into(), 1);
@@ -7325,7 +7331,11 @@ mod tests {
         // a ready `neg_exp_A` initializer, or `A_log` (initializer) feeding an
         // `Exp → Neg` chain that produces the `neg_exp_A` value.
         let dt_bias = vec1d(&mut graph, "dt_bias", DataType::Float32, heads);
-        let coeff_init_name = if neg_exp_from_a_log { "A_log" } else { "neg_exp_A" };
+        let coeff_init_name = if neg_exp_from_a_log {
+            "A_log"
+        } else {
+            "neg_exp_A"
+        };
         let coeff_init = vec1d(&mut graph, coeff_init_name, DataType::Float32, heads);
         for init in [dt_bias, coeff_init] {
             graph.set_initializer(
@@ -7339,9 +7349,19 @@ mod tests {
         }
         let neg_exp_a = if neg_exp_from_a_log {
             let exp_out = value(&mut graph, "exp_out", DataType::Float32, heads);
-            graph.insert_node(Node::new(NodeId(0), "Exp", vec![Some(coeff_init)], vec![exp_out]));
+            graph.insert_node(Node::new(
+                NodeId(0),
+                "Exp",
+                vec![Some(coeff_init)],
+                vec![exp_out],
+            ));
             let neg_out = value(&mut graph, "neg_exp_A", DataType::Float32, heads);
-            graph.insert_node(Node::new(NodeId(0), "Neg", vec![Some(exp_out)], vec![neg_out]));
+            graph.insert_node(Node::new(
+                NodeId(0),
+                "Neg",
+                vec![Some(exp_out)],
+                vec![neg_out],
+            ));
             neg_out
         } else {
             coeff_init

docs/execution/EP_CONFORMANCE.md

@@ -80,7 +80,7 @@ Current CI baseline against `cbourjau/onnx-tests` commit `856e89b`.
 | Status | Operators |
 |---|---|
 | PASS | MatMul, Add, Relu, Reshape, Transpose, Gather, LayerNormalization, Sub, Mul, Div, Pow, Min, Max, Sqrt, Erf, Tanh, Cast, ReduceMean, Softmax, Shape, Unsqueeze, Expand, Slice, Constant, Gemm, Abs, Conv, Sigmoid |
-| UNSUPPORTED | SyntheticUnsupportedCelu (`Celu`) |
+| UNSUPPORTED | SyntheticUnsupportedDet (`Det`) |
 | MISMATCH | None |
 | ERROR | None |