[Roberto (RobertoReale)]

Summary

Adds the WebGPU EP provider option enableMatmulFp32Accumulation (boolean, default false), and routes every MatMulNBits kernel this PR touches through it instead of hardcoding the accumulator type. Weights and activations are untouched: only the register/workgroup accumulators change, so global memory traffic is identical either way.

Default behaviour on main is unchanged by this PR.

This description was rewritten on 2026-08-06 (scope refreshed 2026-08-12, corrected 2026-08-28, renamed 2026-09-02, test coverage updated 2026-09-04). The PR started life as an unconditional f16→f32 accumulator promotion across both the JSEP generators and the native EP, justified by Gemma 3 270M and later by whisper-small q4. Both of those attributions were wrong and I retracted them (Gemma, whisper). What is left is below, and it is deliberately narrower.

Scope of this PR, and what follows it

This is the first of the three changes agreed in #29599 (comment):

1. This PR: enableMatmulFp32Accumulation covering MatMulNBits and its fused variants. The six kernels below, plus the EP plumbing and the JS bindings. The unquantized MatMul, Gemm and Conv kernels under core/providers/webgpu/math/ and core/providers/webgpu/nn/ do not read this option and are unchanged by it.
2. Follow-up: mixed-precision subgroup-matrix support. Use the subgroup-matrix path with componentType = F16 / resultComponentType = F32 when the adapter reports exactly that configuration, and keep the fallback below for adapters that do not.
3. Follow-up: extend the same option to the unquantized MatMul family — MatMul, Gemm, the shared packed implementations and the MatMul-backed Conv paths — with the covered set defined and documented there.

The option is named for the final scope rather than for what it covers today, so that (3) does not need a second rename. Until (3) lands, "enable f32 accumulation" means MatMulNBits and its fused variants only.

Session-scoped, not a model attribute. It is an EP provider option, so nobody has to re-export a model to use it and it does not depend on the model carrying an accuracy_level.

What I am not claiming

I have no real model that overflows one of these accumulators. Not Gemma 3 270M (the f16 partial sum inside all 126 MatMulNBits nodes peaks at ~324; the overflow there is the model-level fp16 residual stream). Not whisper-small q4 (that model has no fp16 tensor at all, so these hunks are inert on it; its real saturation is a Pow inside an unfused decomposed LayerNorm, filed separately as #31626). A/B runs on Qwen2.5-0.5B q4f16, gemma-3-270m and Phi-4-mini across two GPUs showed zero Inf and zero NaN either way, only rounding drift.

What holds up is portability, and only that. WGSL permits extra intermediate precision: Intel's D3D12 compiler (dxc) promotes unrolled acc += chains to f32 on its own but rounds to f16 between loop iterations, while NVIDIA/Vulkan rounds strictly (probes). So "f16 accumulation is safe" is a property of vendor × backend × code shape, not of a vendor — which is why a vendor-string gate cannot express it, and why a kernel carrying an accumulator across a K loop has no headroom guarantee on strict-rounding configurations. That argues for a way back, which is what this option is. It does not argue for changing the default, and I am not asking for that.

What changed

Native WebGPU EP (onnxruntime/contrib_ops/webgpu/quantization/):

• matmul_nbits, matmul_nbits_wide_tile, matmul_nbits_mlp, matmul_nbits_qkv, dp4a_matmul, dp4a_matmul_small_m. Each template gains a compile-time acc_f32 parameter selecting alias acc_element_t = f32 or the output element type; the accumulators, their reduction and the bias adds use that alias, and the downcast happens only at the store. acc_f32 is added to every affected CacheHint so the two variants cannot collide in the program cache.
• SDP8AI in dp4a_matmul_common.wgsl.template returns acc_element_t — it is an exact integer dot4I8Packed dot product, but it must not downcast before the value reaches the accumulator. The mul_precision alias in the same file, which scales that dot product, follows acc_element_t too for 4- and 2-bit weights; 8-bit keeps its unconditional f32 because dot4I8Packed can return an int32 that f16 cannot represent whatever the accumulator is. Without that, the scaled tile result rounded (and could saturate) in f16 on its way into an f32 accumulator, so the main DP4A path was not honouring the option.
• subgroup_matrix_matmul_nbits_* cannot honour an f32 request today: every fp16 entry in supported_subgroup_matrix_configs has resultComponentType == F16 and the kernels declare subgroup_matrix_result<f16, ...>. Rather than ignore the option there, CanApplySubgroupMatrixMatMulNBits declines when f32 accumulation is requested for an fp16 output, so the dispatch falls through to a kernel that does honour it. fp32 outputs already accumulate in f32 and are unaffected. Follow-up (2) above replaces the decline with a mixed-precision config where the adapter offers one.
• Plumbing follows the existing kvCacheQuantizationBits option: constant and "0"/"1" value strings in webgpu_provider_options.h, parsing in webgpu_provider_factory.cc, a config field plus accessor on the EP, and an inline forwarder on ComputeContext.

JS bindings: the option is typed as a boolean in js/common/lib/inference-session.ts, forwarded as "0"/"1" from js/web/lib/wasm/session-options.ts following the enableGraphCapture precedent, and accepted as a boolean by the node bridge in js/node/src/session_options_helper.cc following enableRobustness, so onnxruntime-web and onnxruntime-node users can actually set it.

The JSEP hunks have been removed (matmulnbits.ts, matmul-shaders.ts, 3rd-party/matmul_packed_webgpu.ts are back to main). JSEP shaders read env.webgpu and never see EP provider options, so the option cannot reach them; leaving them in would have meant shipping a PR where the native side is opt-in and the JSEP side is switched unconditionally for everyone. Adding a parallel env.webgpu flag to restore the symmetry means putting an option into a path #29716 is removing, so if JSEP needs this before it goes away it should be its own change.

Testing

• MatMulNBits.Float16_LargeK_AccumulatorOverflow in matmul_4bits_test.cc, built with the option on. A=8, dequantized B at +112 over the first half of K and -112 over the second, K=8192. It runs at M=1 and M=8, so it covers both matmul_nbits.wgsl.template (K is split over tile_size_k_vec lanes, a lane peaks at 114688) and matmul_nbits_wide_tile.wgsl.template (no cross-lane split, results[m] carries the whole prefix and peaks at 3670016, still exact in f32). Either way the sum cancels back to an exact 0, so f32 returns exactly 0 while f16 saturates to +Inf. Every other MatMulNBits test in that file draws its inputs from Gaussian(0, 0.25) and cannot reach 65504 at any K, which is why this class of bug had no coverage. It is a regression guard, not a universal detector: on Intel D3D12 it can pass with an f16 accumulator, and the comment above it says so.
• MatMulNBits.Float16_AccumulatorPrecisionOption_AllPaths runs the generic, wide-tile and dp4a dispatches with the option in both states on ordinary Gaussian inputs, so a variant that fails to compile, an unpropagated flag or a cache hint that cannot separate the two variants shows up as a test failure. On Gaussian inputs both accumulator widths agree inside the tolerance, so this one is a smoke test over the dispatches; the numerical discrimination is in the case above, and the dp4a paths have none because A is int8-quantized before the kernel sees it.
• The fused MLP and QKV decode kernels are reached through graph fusion rather than a MatMulNBits node, so they are covered from the fusion tests instead: MatMulNBitsMlpFusionMatchesUnfusedSimplifiedWebGpuResultsWithFp32Accumulation and MatMulNBitsQkvFusionMatchesUnfusedWebGpuResultsWithFp32Accumulation run the existing fused-vs-unfused comparison with the option on for both the baseline and the fused session.
• WGSL: tools/python/wgsl_gen.py over contrib_ops/webgpu parses all six templates and emits both alias branches for each. The changed shader shapes are compile-checked with naga in both accumulator states, with negative controls confirming the harness rejects an f32 accumulator stored to an f16 output without conversion.
• JS: tsc --noEmit is clean on js/common and js/web, and both changed TypeScript files pass prettier --check.
• CI built this for the first time on 2026-09-02, and it caught a real bug. acc_f32 was passed first in three WGSL_TEMPLATE_APPLY lists, but wgsl_gen.py sorts parameters and a_length_per_tile sorts ahead of acc_f32, so the designators were out of declaration order. GCC, Clang and MSVC all rejected it; fixed in 9a2154314.
• Both new op tests have now run and passed, on webgpu / build-and-test (macOS arm64), which builds without -Werror and so got past the ordering bug: Float16_LargeK_AccumulatorOverflow (216 ms) and Float16_AccumulatorPrecisionOption_AllPaths (3004 ms) on Metal. I still have no local build tree, so the rest of the matrix is CI's word rather than mine.

Performance

f32 accumulators cost registers and shared memory, not bandwidth — the accumulator workgroup arrays double for fp16 outputs, which can lower occupancy on SLM-bound dispatch configurations. Jianhui Dai (@daijh) measured that cost on Intel, and #29611 is the conclusion drawn from it. Since the option defaults to off, nobody pays for it unless they ask.

Related: #29611, #29716, #31626, #26732, #26367, #31703 (both touch matmul_nbits_wide_tile.wgsl.template)

[Roberto (RobertoReale)]

@microsoft-github-policy-service agree

[Roberto (RobertoReale)]

Additional real-world validation

I validated this fix end-to-end against a production consumer of onnxruntime-web (my Chrome extension, Voice Message Transcriber, which runs Whisper via @huggingface/transformers in the browser and currently force-disables WebGPU for q4 models specifically because of #26732).

Method: since this fix only touches the JSEP TypeScript path (js/web/lib/wasm/jsep/...), no wasm binary rebuild was needed. I sparse-cloned js/web + js/common, checked out this PR's branch, built onnxruntime-web from source (pulling the prebuilt wasm artifacts unchanged), and swapped the resulting dist/ into the extension's node_modules/onnxruntime-web in place of the published 1.22.0-dev build. I then temporarily removed the extension's q4/q8 WebGPU exclusion and ran the same voice message through both WASM and WebGPU on Windows, clearing the transcription cache between runs so each device actually re-ran inference.

Result (whisper-small, q4, ~36s of Italian audio):

Device	Time	Output
WASM (baseline)	26.7s	correct, full transcript
WebGPU (this PR's onnxruntime-web build)	13.2s	word-for-word identical text

No [Music]/hallucination, no NaN garbage, no crash — console confirmed device: webgpu used throughout with no fallback to WASM. ~2x speedup with zero accuracy regression, consistent with the "zero bandwidth regression" claim in the PR description (only register/workgroup accumulators promoted to f32).

Only whisper-small was tested this way so far; the fix is generic to the shared MatMul/MatMulNBits kernels rather than model-specific, so I'd expect the same result on tiny/large-v3-turbo (also q4) — happy to confirm those too if useful.

Full write-up (methodology + how I'll re-enable WebGPU once this merges): https://github.com/RobertoReale/Voice-Message-Transcriber/blob/main/docs/webgpu-onnxruntime-fix.md

[Tianlei Wu (tianleiwu)]

Roberto (@RobertoReale), could you run lintrunner -a.
See

onnxruntime/docs/Coding_Conventions_and_Standards.md

Line 156 in a06675e

This project uses [lintrunner](https://github.com/suo/lintrunner) for linting. It provides a consistent linting experience locally and in CI. You can install the dependencies and initialize with

[Tianlei Wu (tianleiwu)]

Thanks for the fix — the root-cause analysis (f16 dot-product accumulation saturating past 65504 for D >= 2048) is correct and the JSEP-side changes (promoting values/acc/workgroup_shared/inter_results to f32 and downcasting only at the final write) look right. A few points before this is complete:

1. The native WebGPU EP fix is incomplete (high priority). Only matmul_nbits.wgsl.template is patched, but the same overflow-prone f16 accumulation pattern exists in the sibling fused kernels that ORT dispatches for the same models:

• matmul_nbits_mlp.wgsl.template — gate_sum/up_sum = output_element_t(0) and gate_inter_results/up_inter_results : array<array<output_element_t, ...>>
• matmul_nbits_qkv.wgsl.template — sum = q_output_element_t(0) and q/k/v_inter_results : array<array<q_output_element_t, ...>>
• dp4a_matmul_small_m.wgsl.template — inter_results : array<array<output_element_t, ...>>

The MLP and QKV fused kernels are selected precisely for the fused MLP / attention-QKV subgraphs that Gemma 3 / SmolLM produce, so those paths can still emit garbage for the models this PR targets. Either extend the same f32-accumulation change to these templates or explicitly document why they are out of scope.

2. Shared-memory / occupancy claim is imprecise. "Zero bandwidth regression / no shared-memory regression" is accurate for the A/B tiles (they stay in the input dtype), but the accumulator workgroup arrays (inter_results, workgroup_shared) do double in size (f16 -> f32) for fp16/q4f16 outputs. On SLM-bound dispatch configs this can lower occupancy. Correctness rightly wins here, but the description/comments should acknowledge the accumulator SLM growth rather than claim zero regression.

3. Missing automated regression coverage. Validation currently relies on manual browser runs plus a static grep script. A numerical unit test (large-K fp16 MatMul / MatMulNBits compared against an f32 reference) in the existing WebGPU op test suite would guard against this regressing again and would exercise the actual runtime rather than source strings.

Inline comments below.

[Jianhui Dai (daijh)]

Hi everyone,
Forcing f32 accumulators for f16 data types can cause register pressure on Intel GPUs, leading to significant performance degradation.
Could we reconsider this approach, or perhaps gate it by vendor device?
In the meantime, I will try to reproduce the reported overflow on my end.

[Jianhui Dai (daijh)]

Fixes the numerical overflow that causes fp16 and q4f16 models (Gemma 3 270M, Whisper Q4, SmolLM2, kokoro-js) to produce garbage output on WebGPU, by promoting the MatMul / MatMulNBits accumulators from f16 to f32 in both the JSEP shader generators and the native WebGPU EP WGSL template.

Could you share the links to these models?

[Roberto (RobertoReale)]

Tianlei Wu (@tianleiwu) Thanks for the review — addressed in 2f2a4b3 and 53c9320:

Fused/native kernels extended to f32 accumulation (same pattern as the generic kernel: buffers and dequantized weights stay fp16/q4, only accumulators are promoted, downcast at the final store):

• matmul_nbits_mlp.wgsl.template — gate_sum/up_sum, gate/up_inter_results, bias and SiLU activation now computed in f32.
• matmul_nbits_qkv.wgsl.template — compute_projection_sum and q/k/v_inter_results in f32.
• dp4a_matmul_small_m.wgsl.template — inter_results and the final reduction in f32. SDP8AI itself is exact (integer dot4I8Packed), so the overflow path was the cross-tile accumulation over K; dp4a_matmul_common.wgsl.template is intentionally unchanged since it is shared.
• While auditing for the same pattern I found that the generic dp4a_matmul.wgsl.template also accumulates lane_outputs/lane_output1..4 in output_element_t across the whole K loop — fixed the same way in 53c9320.

Explicitly excluded: the subgroup_matrix_matmul_nbits_* kernels declare subgroup_matrix_result<f16, ...>; accumulator precision there is tied to the cooperative-matrix configuration negotiated on the C++ side, so moving to an f16×f16→f32 config is a larger change (and MMA units typically accumulate internally at higher precision already). I'd prefer to handle that in a follow-up if overflow is ever reported on those paths.

scripts/verify_f16_fix.py dropped in 2f2a4b3, which also resolves the lintrunner RUFF-FORMAT finding (the remaining diff is TS — prettier-clean — plus WGSL templates; if CI lint still flags anything I'll fix it). On regression coverage: agreed a numerical test is the right long-term answer — e.g. a large-K fp16 MatMulNBits case with inputs arranged so partial sums exceed 65504 while the true result stays representable, compared against an f32 reference. Happy to add that to onnxruntime/test/contrib_ops/matmul_4bits_test.cc in this PR or as a follow-up — whichever you prefer.

[Roberto (RobertoReale)]

Jianhui Dai (@daijh) Thanks for taking a look. Model links (these are the reports collected in #26732 / #26367):

• Gemma 3 270M fp16 — https://huggingface.co/onnx-community/gemma-3-270m-it-ONNX, transformers.js with dtype: "fp16", device: "webgpu" → <unused56> garbage tokens.
• SmolLM2 360M q4f16 — https://huggingface.co/HuggingFaceTB/SmolLM2-360M-Instruct (ONNX weights in the onnx-community mirror), dtype: "q4f16".
• Whisper q4 — https://huggingface.co/onnx-community/whisper-small and https://huggingface.co/onnx-community/whisper-large-v3-turbo, dtype: "q4" → hallucination loops / [Music]. This is the one I validated end-to-end above (word-for-word parity with the WASM EP after the fix, ~2x faster).
• kokoro-js (Kokoro-82M TTS) — https://huggingface.co/onnx-community/Kokoro-82M-v1.0-ONNX, fp16 → corrupted audio.

On register pressure — a few data points for the discussion:

1. The promotion is accumulator-only: buffer traffic, shared-memory tile loads and dequantized weights all stay fp16/q4. Per thread it is a handful of registers (e.g. 16×f32 instead of 16×f16 lane outputs in the generic dp4a kernel); the workgroup inter_results arrays grow by 2 bytes/element.
2. There is precedent in-tree: matmul_nbits_wide_tile.wgsl.template already accumulates in f32, and the dp4a 8-bit path already promotes to f32 before scaling for exactly this overflow reason (mul_precision = f32 in dp4a_matmul_common.wgsl.template).
3. fp32 models are unaffected — with output_element_t == f32 the added casts are identity.
4. The failure mode this fixes is not a minor accuracy loss but Inf/NaN propagation that makes fp16/q4f16 output unusable on WebGPU regardless of speed.

That said, if your measurements on Intel show a significant regression, I'm open to gating — either per-vendor, or (probably better, since overflow risk scales with reduction length) promoting only when K exceeds a threshold. Looking forward to your repro results and numbers; happy to iterate on whatever the data shows.

[Roberto (RobertoReale)]

Follow-up on 53c9320: while compile-validating the changed shader patterns with naga (wgpu-native), I caught a type error I had introduced in the is_qualcomm store path of dp4a_matmul.wgsl.template - WGSL has no implicit conversions, so building vec4<output_element_t> directly from the now-f32 lane scalars is rejected when output_element_t is f16. Fixed in 2cf1e37 by constructing a vec4<f32> first and using the whole-vector conversion constructor. All modified patterns (MLP, QKV, dp4a small-M, dp4a generic, plus the final-store conversions) now compile clean under naga in both f16 and f32 variants; PR description updated to reflect the current scope.

[Jianhui Dai (daijh)]

Gemma 3 270M fp16 — https://huggingface.co/onnx-community/gemma-3-270m-it-ONNX, transformers.js with dtype: "fp16", device: "webgpu" → garbage tokens.
SmolLM2 360M q4f16 — https://huggingface.co/HuggingFaceTB/SmolLM2-360M-Instruct (ONNX weights in the onnx-community mirror), dtype: "q4f16".

I tested these two models using the WebGPU build with the fp16 accumulator from #29611, but I am unable to reproduce the reported issues.

• dtype: "fp16" (MatMul Op)

• dtype: "q4fp16" (MatMulNBitsWideTile Shader)

Could you verify if this still happens on a recent ONNX Runtime build and share a test page to help us reproduce the issue?

[Roberto (RobertoReale)]

Jianhui Dai (@daijh) Thanks for testing — I believe I can explain the non-repro: the fp16 Gemma weights on the Hub today are not the ones the bug was reported against.

On 2025-12-02 the fp16 exports of onnx-community/gemma-3-270m-it-ONNX were re-uploaded with Clip nodes inserted into the graph (update fp16 models with clips) — a model-side mitigation for exactly this overflow (see xenova's comment in #26732 asking reporters to clear the model cache and retry). So a fresh download today reproduces nothing: the graph clamps activations so that f16 partial sums stay in range. Any fp16 export that predates that upload — or that users produce themselves with standard tooling — still hits the overflow.

Reproduction with the same model, pinned to the pre-mitigation revision

transformers.js accepts a revision option, so the original export is still directly reachable:

const gen = await pipeline('text-generation', 'onnx-community/gemma-3-270m-it-ONNX', {
  dtype: 'fp16',
  device: 'webgpu',
  revision: '2950c41f44e1799fa9a42c2dc9fd2648fee8fb4f', // last commit before the clip mitigation
});

I just ran this on stock onnxruntime-web 1.22.0-dev.20250409 (as bundled by transformers.js 3.8.1), Chrome 150 / Windows 11. Notably, the WebGPU adapter Chrome selected is an Intel Iris Xe (vendor=intel, architecture=gen-12lp, shader-f16 supported) — so this overflow reproduces on Intel hardware:

Revision	Device	Output for "What is the capital of France? Reply in one short sentence."
2950c41f (pre-clip)	webgpu	"" — empty/garbage
2950c41f (pre-clip)	wasm	"Paris"
main (with clip mitigation)	webgpu	"Paris"

Same weights: correct on WASM, garbage on WebGPU. The current main weights only pass because of the inserted Clips.

Self-contained test page (serve locally, open in Chrome; ?rev=…&device=webgpu|wasm)

<!doctype html>
<html lang="en">
<head><meta charset="utf-8"><title>Repro onnxruntime#26732</title></head>
<body>
<p id="status">Starting…</p>
<pre id="log"></pre>
<script type="module">
  const params = new URLSearchParams(location.search);
  const rev = params.get('rev') || 'main';
  const device = params.get('device') || 'webgpu';
  const log = s => { document.getElementById('log').textContent += s + '\n'; };
  try {
    log(`revision=${rev} device=${device}`);
    if (navigator.gpu) {
      const adapter = await navigator.gpu.requestAdapter();
      const i = adapter?.info || {};
      log(`GPU adapter: vendor=${i.vendor} architecture=${i.architecture}`);
      log(`shader-f16 supported: ${adapter?.features.has('shader-f16')}`);
    }
    const { pipeline } = await import('https://cdn.jsdelivr.net/npm/@huggingface/transformers@3.8.1/dist/transformers.min.js');
    const gen = await pipeline('text-generation', 'onnx-community/gemma-3-270m-it-ONNX',
                               { dtype: 'fp16', device, revision: rev });
    const out = await gen([{ role: 'user', content: 'What is the capital of France? Reply in one short sentence.' }],
                          { max_new_tokens: 24, do_sample: false });
    const text = out[0].generated_text.at(-1).content.trim();
    log('OUTPUT: ' + JSON.stringify(text));
    const bug = text === '' || /<unused\d+>/.test(text) || text.includes('NaN');
    document.getElementById('status').textContent = bug
      ? 'BUG: garbage output (fp16 overflow)' : 'OK: coherent output';
  } catch (e) { log('ERROR: ' + e.message); }
</script>
</body>
</html>

Reproductions that don't depend on re-exported weights

• Kokoro-82M fp16 (kokoro-js) — never re-exported (last weight upload 2025-02), and reported still broken on WebGPU on 2026-01-20: [Web] fp16 and q4f16 Gemma 3 models produce invalid outputs on WebGPU due to overflow in ONNX runtime #26732 (comment)
• Whisper q4 — the end-to-end validation earlier in this thread: hallucination loops on stock WebGPU builds (the reason my extension force-disables WebGPU for q4 models today), word-for-word parity with WASM after this PR.

On SmolLM2-360M specifically: its reduction lengths are comparatively short (d_model 960, FFN 2560), so failures there are prompt/data-dependent rather than deterministic — pre-clip Gemma and Kokoro are the reliable reproducers.

Re accuracy_level: replied in #29611 to keep that discussion in one place — short version: it works as an explicit opt-down for models that knowingly tolerate f16 accumulation, but the unset default needs to be f32, since none of the affected models in the wild set the attribute, and the plain fp16 MatMul path (this Gemma case) has no such attribute at all.

[Jianhui Dai (daijh)]

Jianhui Dai (@daijh) Thanks for testing — I believe I can explain the non-repro: the fp16 Gemma weights on the Hub today are not the ones the bug was reported against.

On 2025-12-02 the fp16 exports of onnx-community/gemma-3-270m-it-ONNX were re-uploaded with Clip nodes inserted into the graph (update fp16 models with clips) — a model-side mitigation for exactly this overflow (see xenova's comment in #26732 asking reporters to clear the model cache and retry). So a fresh download today reproduces nothing: the graph clamps activations so that f16 partial sums stay in range. Any fp16 export that predates that upload — or that users produce themselves with standard tooling — still hits the overflow.

Reproduction with the same model, pinned to the pre-mitigation revision

transformers.js accepts a revision option, so the original export is still directly reachable:

const gen = await pipeline('text-generation', 'onnx-community/gemma-3-270m-it-ONNX', {
  dtype: 'fp16',
  device: 'webgpu',
  revision: '2950c41f44e1799fa9a42c2dc9fd2648fee8fb4f', // last commit before the clip mitigation
});

I just ran this on stock onnxruntime-web 1.22.0-dev.20250409 (as bundled by transformers.js 3.8.1), Chrome 150 / Windows 11. Notably, the WebGPU adapter Chrome selected is an Intel Iris Xe (vendor=intel, architecture=gen-12lp, shader-f16 supported) — so this overflow reproduces on Intel hardware:

Revision Device Output for "What is the capital of France? Reply in one short sentence."
2950c41f (pre-clip) webgpu "" — empty/garbage
2950c41f (pre-clip) wasm "Paris"
main (with clip mitigation) webgpu "Paris"
Same weights: correct on WASM, garbage on WebGPU. The current main weights only pass because of the inserted Clips.

Self-contained test page (serve locally, open in Chrome; ?rev=…&device=webgpu|wasm)

<!doctype html>
<html lang="en">
<head><meta charset="utf-8"><title>Repro onnxruntime#26732</title></head>
<body>
<p id="status">Starting…</p>
<pre id="log"></pre>
<script type="module">
  const params = new URLSearchParams(location.search);
  const rev = params.get('rev') || 'main';
  const device = params.get('device') || 'webgpu';
  const log = s => { document.getElementById('log').textContent += s + '\n'; };
  try {
    log(`revision=${rev} device=${device}`);
    if (navigator.gpu) {
      const adapter = await navigator.gpu.requestAdapter();
      const i = adapter?.info || {};
      log(`GPU adapter: vendor=${i.vendor} architecture=${i.architecture}`);
      log(`shader-f16 supported: ${adapter?.features.has('shader-f16')}`);
    }
    const { pipeline } = await import('https://cdn.jsdelivr.net/npm/@huggingface/transformers@3.8.1/dist/transformers.min.js');
    const gen = await pipeline('text-generation', 'onnx-community/gemma-3-270m-it-ONNX',
                               { dtype: 'fp16', device, revision: rev });
    const out = await gen([{ role: 'user', content: 'What is the capital of France? Reply in one short sentence.' }],
                          { max_new_tokens: 24, do_sample: false });
    const text = out[0].generated_text.at(-1).content.trim();
    log('OUTPUT: ' + JSON.stringify(text));
    const bug = text === '' || /<unused\d+>/.test(text) || text.includes('NaN');
    document.getElementById('status').textContent = bug
      ? 'BUG: garbage output (fp16 overflow)' : 'OK: coherent output';
  } catch (e) { log('ERROR: ' + e.message); }
</script>
</body>
</html>

Reproductions that don't depend on re-exported weights

• Kokoro-82M fp16 (kokoro-js) — never re-exported (last weight upload 2025-02), and reported still broken on WebGPU on 2026-01-20: [Web] fp16 and q4f16 Gemma 3 models produce invalid outputs on WebGPU due to overflow in ONNX runtime #26732 (comment)
• Whisper q4 — the end-to-end validation earlier in this thread: hallucination loops on stock WebGPU builds (the reason my extension force-disables WebGPU for q4 models today), word-for-word parity with WASM after this PR.

On SmolLM2-360M specifically: its reduction lengths are comparatively short (d_model 960, FFN 2560), so failures there are prompt/data-dependent rather than deterministic — pre-clip Gemma and Kokoro are the reliable reproducers.

Re accuracy_level: replied in #29611 to keep that discussion in one place — short version: it works as an explicit opt-down for models that knowingly tolerate f16 accumulation, but the unset default needs to be f32, since none of the affected models in the wild set the attribute, and the plain fp16 MatMul path (this Gemma case) has no such attribute at all.

Roberto (@RobertoReale)
It would be helpful to have a simple test page with easy steps to reproduce this using a recent ONNX Runtime build.

By the way, regarding the previous issue, were you using the legacy JSEP or the WebGPU EP?

[Roberto (RobertoReale)]

By the way, regarding the previous issue, were you using the legacy JSEP or the WebGPU EP?

Legacy JSEP — transformers.js 3.8.1 bundles onnxruntime-web@1.22.0-dev.20250409, which only ships the JSEP artifact (ort-wasm-simd-threaded.jsep.wasm; verified from the network log of the run).

So I re-ran everything on the native WebGPU EP with a recent build — transformers.js 4.2.0 → onnxruntime-web@1.26.0-dev.20260416 (loads ort-wasm-simd-threaded.asyncify.wasm), plus the standalone tests below on onnxruntime-web@1.27.0 stable. Same machine, same Intel Iris Xe (gen-12lp) adapter, Chrome 150 / Windows 11. The results explain rather precisely why you can't reproduce on Intel, and I think they're important for the f16-accumulator discussion in #29611:

1. On the native EP, the pre-clip Gemma model fails before it can overflow

• main revision (with clip mitigation): works, outputs "Paris".

• pre-clip revision 2950c41f: shader compilation error —

  LayerNorm: 46:16 cannot assign 'vec4<f16>' to 'vec4<f32>'
  → OrtRun failed: SimplifiedLayerNormalization node '/model/layers.18/final_norm_layernorm/LayerNorm'
  

  The pre-clip export's final SimplifiedLayerNormalization has X = f16 but scale/Y = f32 (the exporter fused a Cast into the LN output), and the shader generator assumes X and Y have the same type (layer_norm.cc#L116: y[offset + i] = input_value;). JSEP runs this graph fine. Happy to file this as a separate issue — but it means the original fp16 exports can't currently be used to test the overflow on the native EP at all.

2. A minimal test that needs no model download — and why it passes on Intel

Self-contained page (below): a 119-byte embedded ONNX model, MatMul(A[M,4096] f16 × B[4096,N] f16) → Y f16, with inputs chosen so every product is ±36000 in a + + − − pattern along K: any two consecutive same-sign terms exceed f16 max (72000 > 65504), while the exact result is 0. Run on onnxruntime-web@1.27.0, WebGPU vs WASM.

On my Intel Iris Xe: WebGPU returns exact 0 — no overflow, even though the generated MatMulSubgroup shader accumulates in vec4<f16> (var acc_0 = vec4<output_element_t>(0) with alias output_element_t = f16).

That looks like evidence that f16 accumulators are fine — until you probe the same device with raw WGSL:

WGSL code shape (same ±36000 data, exact sum = 0)	Intel Iris Xe, D3D12 (Chrome)
acc += f16(16)*inp[k] in a loop over 4096 terms	Inf
f16(16)*inp[0] + f16(16)*inp[1] (one expression)	72000 — i.e. > f16 max, not rounded
32 unrolled acc += statements per tile in a tile loop	0

The Intel shader compiler evaluates straight-line f16 chains at higher precision and only rounds to f16 across loop iterations. ORT's native MatMul kernels are heavily unrolled, so on Intel they get f32 accumulation silently, from the driver — which is why the f16-accumulator builds look correct there. The JSEP kernels accumulate in for-loops, get real per-iteration f16 rounding, and overflow — on the same physical GPU. That is exactly the JSEP-garbage / native-EP-OK split I'm seeing, with one machine.

And it does not transfer across vendors. The same probes via wgpu-native:

	loop	2-term expr	32 unrolled
Intel Iris Xe / Vulkan	0 (promoted)	72000 (promoted)	Inf
NVIDIA RTX 3050 Ti / Vulkan	Inf	Inf (strict rounding)	0 (reassociated)

Six configurations, six behaviors. This is all WGSL-conformant — implementations are allowed to evaluate floating-point expressions with extra intermediate precision — and that's precisely the problem: whether an f16 accumulator overflows is an accident of vendor × driver × code shape. A kernel validated with f16 accumulators on Intel says nothing about NVIDIA/Qualcomm/Apple, and even on Intel it flips between D3D12 and Vulkan and between looped and unrolled codegen.

The explicit f32 accumulator is the only way to make the guarantee portable — and on hardware whose compiler already promotes intermediates (Intel), it changes essentially nothing, which is consistent with the neutral perf numbers discussed earlier.

(Side note for scope: the native EP's plain fp16 MatMul/Gemm shaders also declare f16 accumulators — gemm_utils.cc#L284 var acc: array<vec4<a_element_t>, …> — and are currently protected only by this driver behavior (plus the split-K path for small shapes, which keeps per-slice partial sums small). This PR covers JSEP + the quantized native templates; I'm happy to extend the same accumulator-only change to gemm_utils.cc here or as a follow-up.)

Test page — save as .html, serve locally, open in Chrome (?ver=1.27.0&m=256&n=512)

<!doctype html>
<html lang="en">
<head><meta charset="utf-8"><title>fp16 MatMul accumulator overflow — onnxruntime#26732</title></head>
<body>
<p id="status">Running…</p>
<pre id="log"></pre>
<script type="module">
  const params = new URLSearchParams(location.search);
  const ver = params.get('ver') || '1.27.0';
  const M = Number(params.get('m')) || 256, N = Number(params.get('n')) || 512, K = 4096;
  const log = s => { document.getElementById('log').textContent += s + '\n'; };

  // 119-byte ONNX model: MatMul(A[M,4096] f16, B[4096,N] f16) -> Y[M,N] f16, opset 17, M/N dynamic
  const MODEL_B64 = 'CAg6bQoRCgFBCgFCEgFZIgZNYXRNdWwSE21hdG11bF9mMTZfb3ZlcmZsb3daFQoBQRIQCg4IChIKCgMSAU0KAwiAIFoVCgFCEhAKDggKEgoKAwiAIAoDEgFOYhUKAVkSEAoOCAoSCgoDEgFNCgMSAU5CBAoAEBE=';
  const modelBytes = Uint8Array.from(atob(MODEL_B64), c => c.charCodeAt(0));

  // A = 16.0 (0x4C00); B alternates +2250,+2250,-2250,-2250 (0x6865/0xE865) along K.
  // Products are ±36000: two consecutive same-sign terms exceed f16 max; exact row sum = 0.
  const A = new Uint16Array(M * K).fill(0x4C00);
  const B = new Uint16Array(K * N);
  for (let k = 0; k < K; k++) for (let n = 0; n < N; n++) B[k * N + n] = (k % 4 < 2) ? 0x6865 : 0xE865;

  const f16ToF32 = h => {
    const s = (h & 0x8000) ? -1 : 1, e = (h >> 10) & 0x1f, m = h & 0x3ff;
    if (e === 0x1f) return m ? NaN : s * Infinity;
    return e === 0 ? s * m * 2 ** -24 : s * (1 + m / 1024) * 2 ** (e - 15);
  };

  async function runOn(ort, ep) {
    const session = await ort.InferenceSession.create(modelBytes, { executionProviders: [ep] });
    const { Y } = await session.run({ A: new ort.Tensor('float16', A, [M, K]),
                                      B: new ort.Tensor('float16', B, [K, N]) });
    let nonFinite = 0, maxAbs = 0;
    for (const h of Y.data) { const v = f16ToF32(h);
      if (!Number.isFinite(v)) nonFinite++; else maxAbs = Math.max(maxAbs, Math.abs(v)); }
    log(`${ep}: nonFinite=${nonFinite}/${Y.data.length} maxAbsFinite=${maxAbs} (exact result: all zeros)`);
    await session.release();
    return nonFinite;
  }

  const a = await navigator.gpu?.requestAdapter();
  log(`GPU: ${a?.info?.vendor} ${a?.info?.architecture} shader-f16=${a?.features.has('shader-f16')}`);
  const ort = await import(`https://cdn.jsdelivr.net/npm/onnxruntime-web@${ver}/dist/ort.webgpu.bundle.min.mjs`);
  await runOn(ort, 'wasm');
  const bad = await runOn(ort, 'webgpu');
  document.getElementById('status').textContent =
    bad ? 'BUG: f16 accumulator overflow on WebGPU' : 'OK on this adapter (see notes on driver-dependent f16 precision)';
</script>
</body>
</html>

Raw WGSL probe (same data), for checking any adapter's f16 rounding behavior

enable f16;
@group(0) @binding(0) var<storage, read> inp: array<f16>;        // +2250,+2250,-2250,-2250 repeated
@group(0) @binding(1) var<storage, read_write> outp: array<f32>;
@compute @workgroup_size(1)
fn main() {
  var acc: f16 = f16(0);                                          // 1) looped accumulation
  for (var k = 0u; k < 4096u; k++) { acc += f16(16.0) * inp[k]; }
  outp[0] = f32(acc);                                             // Inf ⇒ strict f16; 0 ⇒ promoted
  outp[1] = f32(f16(16.0) * inp[0] + f16(16.0) * inp[1]);         // 2) Inf ⇒ strict; 72000 ⇒ promoted
}

[Jianhui Dai (daijh)]

This is a great deep dive!
I don't intend to block the current PR.
Since this is an NVIDIA-specific issue, could you gate the fix behind a vendor string check?
Please let me know if any issues arise on Intel hardware, and I'd be happy to follow up.

[Roberto (RobertoReale)]

Thanks Jianhui Dai (@daijh) — really appreciate you not blocking, and the offer to follow up.

One thing I want to correct before we settle on the gating, because I think it changes the conclusion: this isn't NVIDIA-specific. The failure that #26732 was actually reported against runs on Intel too — my original garbage repro (onnx-community/gemma-3-270m-it-ONNX @ 2950c41f, pre-clip) was on an Intel Iris Xe (gen-12lp) adapter via legacy JSEP, output "". So an NVIDIA-only vendor gate would leave the reporting users unfixed on the exact hardware I reproduced on.

The reason your test passes on Intel while mine fails on the same vendor is the piece I probed in the deep-dive, and it's the key point for gating:

• Intel's D3D12 shader compiler evaluates straight-line / unrolled f16 acc += chains at higher internal precision and only rounds across loop iterations. ORT's native kernels are heavily unrolled → they get f32-like accumulation for free from the driver. The JSEP kernels are looped → they round every iteration → overflow. Same Intel GPU, opposite result, purely a function of code shape.
• And it's not even stable within Intel: the same Iris Xe on the Vulkan backend rounds the two-term f16 expression strictly → Inf, where D3D12 promotes it. Both are WGSL-conformant (extra precision is permitted, not guaranteed).

So the property "f16 accumulation is safe here" isn't vendor == Intel — it's vendor × backend × whether-this-particular-kernel-is-unrolled, and it flips on a driver update or a codegen change. A static vendor string can't express that, and an NVIDIA allowlist additionally excludes AMD / Qualcomm / Apple / Intel-on-Vulkan, all of which can round strictly.

On cost: I don't think there's a measured regression to gate around yet. The change is accumulator-only (weights/activations stay f16/q4 in VRAM, downcast at store), and the in-tree wide-tile kernel already accumulated in f32 with the comment "minimal performance impact compared to an f16 accumulator" — which #29611 is now removing. If you do have Panther Lake numbers showing f32 accumulation costs measurably on a specific fast path, I'd genuinely like to see them — that's exactly the input that should drive any narrowing.

And that's the constructive version I'd propose instead of a vendor gate: default to f32 (correct everywhere), and opt down to f16 only where it's both proven safe and proven to cost. Your accuracy_level idea fits this perfectly for MatMulNBits — a model that tolerates f16 accumulation sets accuracy_level = 2 and keeps the fast path, mirroring how the dp4a kernels already gate on accuracy_level = 4. The one invariant is that unset (0) must mean f32, since none of the affected exports in the wild set the attribute. A K-threshold promotion (overflow risk scales with reduction length) is another data-driven knob if you'd prefer. Happy to implement either on top of this PR.

Last thing, tying the two PRs together: if #29611 lands the wide-tile accumulator as output_element_t and #29599 is gated to NVIDIA, then q4f16 models on any non-NVIDIA strict-rounding config hit the wide-tile path with no f32 fallback — which is precisely the SmolLM2 / Gemma-q4f16 failure mode. That regression only exists as the combination of the two changes, so I wanted to flag it here rather than in isolation.

[Jianhui Dai (daijh)]

IMO, JSEP actually gets less maintenance.

I'm really looking forward to a test page where a language model inference session causes an overflow on an Intel GPU via either D3D or Vulkan, because of a f16 accumulators in MatMul/MatMulNBits.

[Roberto (RobertoReale)]

Thanks Jianhui Dai (@daijh) — completely agree on JSEP getting less maintenance than the C++ Native WebGPU EP, and that is why this PR updates both paths (the TS generators in js/web/... for legacy/high-level wrappers like transformers.js v3 where users hit this overflow on Intel Iris Xe today, and the C++ contrib_ops/webgpu/quantization/*.wgsl.template shaders for modern ort-web / .asyncify.wasm).

On why f16 accumulators in MatMul/MatMulNBits on the Native WebGPU EP overflow on Intel via Vulkan (and why a vendor check vendor != 'nvidia' isn't safe for Intel):

As explored in the raw-WGSL probes (#issuecomment-4923584353), the reason f16 accumulators in unrolled MatMul templates happen to pass on Intel D3D12 (dxc) without f32 is that dxc promotes straight-line / unrolled acc += chains across kTileK to higher internal precision before emitting DXIL. Because ORT's Native EP MatMul/MatMulNBits templates (matmul_nbits.wgsl.template, dp4a_matmul*.wgsl.template) are heavily unrolled along K, they get f32-like accumulation silently from the Intel D3D12 driver.

However, on Intel via Vulkan (SPIR-V), that unrolled promotion does NOT occur.
As shown in our empirical Vulkan probe on Iris Xe (#issuecomment-4923584353), SPIR-V (OpFAdd %half) evaluates unrolled 16-bit float additions strictly. When K >= 2048, any MatMul / MatMulNBits inner loop with f16 accumulators whose partial sums exceed 65504 deterministically saturates to +Inf (unrolled -> Inf) on Intel GPU via Vulkan and propagates NaN through subsequent LayerNorm / Softmax.
(On D3D12, the Native EP Gemma 3 270M pre-clip currently hits a separate shader compile crash cannot assign vec4<f16> to vec4<f32> inside LayerNorm due to layer_norm.cc:116 assuming equal types when X=f16, scale=f32, so MatMul isn't reached there yet).

Because f16 accumulation safety on Intel is purely a consequence of dxc unroll promotion (vendor × backend × codeshape), an NVIDIA-only vendor gate would leave Intel users on Vulkan (as well as AMD RDNA, Qualcomm Adreno, and Apple Silicon users) exposed to MatMul/MatMulNBits overflow when K is large.

This is why @qjia7 and @hariharans29 noted on #29611 that universally accumulating in f16 across all workloads is risky and that f32 accumulators is the right direction (and notably, the existing in-tree wide_tile template matmul_nbits_wide_tile.wgsl.template already uses f32 accumulators today (var results : array<f32, kTileM>)).

I am 100% supportive of your idea to use accuracy_level (MatMulNBits attribute) as an opt-down (accuracy_level = 2) to enable f16 accumulators when a quantized model explicitly opts in to maximum ALU throughput over dynamic range (just as dp4a gates on level = 4). But to ensure out-of-the-box correctness across all vendors and backends (including Intel Vulkan), the default when accuracy_level is unset (0) should remain f32 accumulation as implemented here.

[Roberto (RobertoReale)]

Tianlei Wu (@tianleiwu) Jiajia Qin (@qjia7) Hariharan Seshadri (@hariharans29) Jianhui Dai (@daijh) — I've corrected this PR's description, because its central claim was wrong.

I had been citing Gemma 3 270M as the real-model proof of an f16 accumulator overflow, and asserting that the Clip nodes added to the Hub model on 2025-12-02 were a workaround for it. I finally checked that instead of arguing it, and it is false:

• On the pre-mitigation revision (2950c41f), with real activations and real dequantized q4 weights, the f16 running partial sum inside all 126 MatMulNBits nodes peaks at ~324 — three orders of magnitude below 65504. down_proj (K=2048) peaks at 37.
• The Clip nodes are not on any matmul output. They are 18 nodes on the residual stream (/model/layers.N/Add_1), bounded at ±32752.
• The fp16 residual on that revision crosses 65504 between layers 7 and 8 (L7/Add_1 = 52,352 → Inf). That is the known Gemma-3 fp16 activation-range problem — which the original reporter of [Web] fp16 and q4f16 Gemma 3 models produce invalid outputs on WebGPU due to overflow in ONNX runtime #26732 had actually flagged himself.

So #26732's root cause is a model-level fp16 residual overflow, and this PR does not fix it. I've removed Fixes #26732 and the Gemma / SmolLM2 / kokoro claims. Full data on #29611.

What this PR still rests on: whisper-small q4, which is a genuine end-to-end accumulator overflow on the JSEP path (garbage on WebGPU → word-for-word WASM parity with only the accumulator promoted), plus the portability argument for strict-rounding backends. That is narrower than what I originally wrote, and Jianhui Dai (@daijh) was right to keep pushing for a real test case rather than accept the synthetic ones.

I'd rather you review this on what it actually demonstrates. If you'd prefer I re-scope it further (e.g. JSEP-only, or hold the native-EP hunks until Jianhui Dai (@daijh)'s A/B packages report), say so and I'll do it.

[Roberto (RobertoReale)]

Two things I owed this PR from Tianlei Wu (@tianleiwu)'s review.

The unit test is in, Float16_LargeK_AccumulatorOverflow. Writing it turned up why this class of bug never showed here: every MatMulNBits test in that file draws its inputs from Gaussian(0, 0.25), so even Float16_Large at K=11008 keeps the running partial sum under 5. The tests cannot reach 65504 by accident. The new case sets A=8 and dequantized B to +112 over the first half of K and -112 over the second, K=8192, M=1 so it dispatches to matmul_nbits.wgsl.template. A lane walks 256 elements, peaks at 114688 halfway, then cancels; the exact output is 0. With the f32 accumulator it returns exactly 0, with f16 it returns Inf. The block-local sum peaks at 28672 and stays in range, so only the cross-K accumulator is under test.

It's a guard, not a detector: on Intel D3D12 dxc promotes unrolled f16 chains by itself, so it can pass there even with an f16 accumulator. That's stated in the comment above the test.

Second, the description claimed zero regression, which was only true of global memory. The accumulator workgroup arrays do double for fp16 outputs. Corrected.

[Roberto (RobertoReale)]

Done, pushed as 8a8d9305d. The option is the boolean enableMatmulFp32Accumulation, default false, and the description now states the MatMulNBits-and-fused scope with your steps 2 and 3 written out as the follow-ups. I kept your name rather than putting "NBits" into it so step 3 doesn't need a second rename, and said explicitly that until step 3 lands "enable" means MatMulNBits and its fused variants only.

Native side is the constant with "0"/"1" values, the config field and the ComputeContext forwarder. JS is a boolean in js/common, forwarded from js/web the way enableGraphCapture is, and a boolean branch in the node bridge next to enableRobustness. tsc --noEmit is clean on both JS packages, and wgsl_gen.py still emits both alias branches for all six templates. The C++ is still uncompiled here.

One thing I ran into reading ahead to step 2. All four entries in supported_subgroup_matrix_configs have componentType == resultComponentType, and IsSubgroupMatrixConfigSupported filters on componentType against is_fp16 alone (subgroup_matrix_config.cc:23-24). An F16/F32 entry would match an fp16 output whether the flag is set or not, so that selector needs the request threaded into it, not just a new row. That's for the follow-up PR.

Eight weeks in, license/cla is still the only check that has ever reported here. A /azp run from someone authorised is the one thing left that I can't do myself.

[Copilot]

🟡 Changes recommended

DP4A still performs some scaling in f16 before widening, and several option-specific test paths are ineffective or uncovered.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Review details

Suppressed comments (3)

onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul_common.wgsl.template:638

• This overload has the same premature f16 multiplication for 4-bit and 2-bit paths: casting the already-scaled result to acc_element_t cannot recover from overflow. Make the shared multiplication precision f32 whenever acc_f32 is enabled, without regressing the existing 8-bit f32 multiplication.

        return acc_element_t(mul_precision(local_sum) * mul_precision(scale));

onnxruntime/test/contrib_ops/matmul_4bits_test.cc:1009

• In plugin/adapters builds this helper can return null when WebGPU is unavailable, so RunTest receives an invalid EP and fails instead of skipping. Check the provider before adding it to the vector, as other WebGPU-specific tests do.

      execution_providers.push_back(WebGpuExecutionProviderWithOptions(config_options));

onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.wgsl.template:240

• Issue #26732 does not establish this Dawn/D3D12 temporary-demotion behavior, and the PR description explicitly retracts that issue's MatMulNBits attribution. Remove the stale issue reference so the implementation comment does not misstate the report's cause.

        // Explicit cast: Dawn/D3D12 re-demotes temporaries to f16 when
        // 'enable f16;' is active (issue #26732).
        inter_results[local_row_offset + idy][idx] += acc_element_t(sum);

• Files reviewed: 23/23 changed files
• Comments generated: 6
• Review effort level: Balanced

[Roberto (RobertoReale)]

The pipelines ran, and they caught a real bug on the first attempt. Fix pushed as 9a2154314.

wgsl_gen.py sorts template parameters, so in the generated struct a_length_per_tile comes before acc_f32 (_ sorts under c). I had put acc_f32 at the head of the WGSL_TEMPLATE_APPLY lists, which left the designators out of declaration order in matmul_nbits.cc, matmul_nbits_mlp.cc and matmul_nbits_qkv.cc. C++20 wants declaration order and all three compilers said so independently: GCC -Werror=missing-field-initializers, Clang -Wreorder-init-list, MSVC C7560. The wide-tile and the two dp4a call sites were correct by accident, since with no a_length_per_tile in the list acc_f32 genuinely does sort first. That is exactly why those three compiled and the other three did not. Also applied clang-format to the two files lintrunner flagged.

The part worth recording: on webgpu / build-and-test (macOS arm64), which builds without -Werror and so got past the ordering bug, both new tests ran and passed on Metal. Float16_LargeK_AccumulatorOverflow in 216 ms, Float16_AccumulatorPrecisionOption_AllPaths in 3004 ms. First time either has executed anywhere, and it is the answer to the thing I have been asking for since July.

I also merged main (ce96b9a58). 1. Build Full ORT and Generate ORT Files was red on uuid/uuid.h, which is the coremltools problem Lapis0x0 (@Lapis0x0) chased down on #32294; 2a5b485 adds the libuuid vcpkg port to the coreml-ep feature, so that job should go green on the new head rather than needing a re-run.

[Jiajia Qin (qjia7)]

Roberto (@RobertoReale) Please resolve Copilot's comments, thanks.

[Roberto (RobertoReale)]

Pushed 4f40a04e2.

The dp4a one was the consequential one. mul_precision stayed output_element_t for 4 and 2 bits, so mul_precision(local_sum) * mul_precision(scale) rounded in f16 before the result ever reached the accumulator: with the option on, the main DP4A path was widening the accumulator and then feeding it values that had already lost the range. It follows acc_element_t there now, and 8 bits keeps its unconditional f32 for the int32-to-f16 reason that was already documented. naga compiles all six (n_bits, acc_f32, output type) combinations, and rejects the two negative controls where an f32 value reaches an f16 output unconverted.

WebGpuExecutionProviderWithOptions is now built and null-checked before the OpTester exists, in both tests.

The fused MLP and QKV comments were fair, and I was wrong in August when I said I had no entry point for them. RunWebGpuFusionTransformerTest takes an EP factory, so passing one with the option on covers the baseline and the fused session both. There is a case in each fusion test file now.

I also dropped the #26732 references from the shader comments. They contradicted the description, which retracted that issue as motivation.

The one I have not fully done is the discrimination in Float16_AccumulatorPrecisionOption_AllPaths. I extended Float16_LargeK_AccumulatorOverflow to M = 8 so the wide-tile dispatch gets the same numerical case as the generic one (results[m] carries the whole prefix there, peaking at 3670016), but dp4a has none: A is int8-quantized before the kernel sees it, and I have not built a construction that saturates an f16 accumulator on that path. I left that thread open and wrote the limitation into the comment above the test rather than papering over it.

[Roberto (RobertoReale)]

CI finally ran on 4f40a04e2 yesterday and two tests failed, both mine. Fix pushed as 19f3697cf.

Float16_AccumulatorPrecisionOption_AllPaths fails only on {M=8, N=128, K=4096, accuracy_level=4}, and it fails identically with the option off and on, so the accumulator is not what is moving. That is the one case that actually reaches CanApplyDP4AMatrixMatMulNBits — the M=2 case does not, since an fp16 output needs M >= kMinMForTileOptimization — and dp4a int8-quantizes A while the CPU reference does not. At K=4096 that sets the floor around 0.08 against my 0.055 bound. The accuracy_level 4 cases now get 0.15/0.03 and the rest keep 0.055/0.02; Float16_Large already allows 0.1 at this K for accuracy_level 0. I fixed the dispatch comment too, it claimed M=2 covered dp4a small-M.

The MLP fusion failure is my mistake in the other direction: I copied 1e-3 from the option-off test. With acc_element_t = f32 the fused kernel rounds once, at the store, while the baseline rounds at gate_out, up_out, the Sigmoid output and the SiLU Mul. Four roundings against one with a nonlinearity between them is a 1.2% gap on the largest outputs (4.10156 against 4.15234). Now 2e-2. The QKV test passed unchanged at 2e-3/5e-3, which fits, there is no activation in that fusion.

These lanes are non-Apple, so dp4a executed for the first time. The other red WebGPU jobs are the gradle wrapper timing out on services.gradle.org, not us. I have no WebGPU build tree here, so the numbers above are read off the CI logs and I could not run the new bounds before pushing.

[Copilot]

🟡 Changes recommended

The generic kernel can still overflow locally before widening, and DP4A behavior lacks a discriminating regression test.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Review details

Suppressed comments (1)

onnxruntime/test/contrib_ops/matmul_4bits_test.cc:1022

• This test explicitly admits that an ignored flag still passes and that no DP4A case numerically distinguishes the two accumulator modes. Consequently, a missing DP4A propagation or cache-key discriminator can pass despite the earlier claim that this test catches it. Add a DP4A-specific overflow/precision case, or directly assert the generated variant/cache key, so the changed DP4A behavior is actually regression-tested.

// answer to within the tolerance, so a run that quietly ignored the flag would still pass here. The
// numerical discrimination lives in the test above, and only for the two dispatches it reaches. The
// dp4a paths have no equivalent case: A is int8-quantized before the kernel sees it, so a construction
// whose exact result stays in range while the f16 prefix saturates needs a different setup, and I have
// not written one.

• Files reviewed: 25/25 changed files
• Comments generated: 3
• Review effort level: Balanced

[Roberto (RobertoReale)]

Pushed 861788261. Copilot's first point is right and I took it.

matmul_nbits.wgsl.template built its block-local sum in output_element_t and widened only when adding into inter_results, so with the option on a prefix of component_b dot4 pairs still ran in f16. The prefix is short and I have no case that saturates it, but the wide-tile, QKV and MLP variants all widen at the dot operands already, so the generic path was the only one where turning the option on still left a saturation point the caller cannot do anything about. It matches them now. The operands go through an explicit vec4<output_element_t> constructor first, since WGSL has no implicit f16 to f32 conversion and the shorter vec4<f32>(vec2<f16>, vec2<f16>) form does not compile. naga takes both branches of the alias across the three component_a shapes and rejects that negative control. I have not measured the extra registers; it only costs anything with the option on.

The other two were fair. "Only the accumulators change" is wrong for the fused MLP, where the bias, the SiLU and the gate/up product all run on the f32 accumulators and round once at the store, which is exactly what the 2e-2 tolerance is paying for. Corrected in the provider option header and in the JS binding, and in the comment above Float16_LargeK_AccumulatorOverflow, which claimed the block-local sum stayed in output_element_t by design. The arithmetic in that comment is unchanged, so the case still isolates the cross-K accumulator.

Unrelated: the 33 workflow runs on 19f3697cf are still action_required and never started. The 4f40a04e2 batch sat the same way from the 4th to the 7th.

[Roberto (RobertoReale)]

The full matrix ran on 861788261, approved about three hours after the push, and the two tests that failed on 4f40a04e2 are green. On webgpu_build_x64_RelWithDebInfo (novcpkg, static): Float16_AccumulatorPrecisionOption_AllPaths 2845 ms, Float16_LargeK_AccumulatorOverflow 550 ms, MatMulNBitsMlpFusionMatchesUnfusedSimplifiedWebGpuResultsWithFp32Accumulation 686 ms, MatMulNBitsQkvFusionMatchesUnfusedWebGpuResultsWithFp32Accumulation 718 ms, and no failures in the lane. The (vcpkg, static) lane gives 2922/556 ms for the two MatMulNBits cases. All five webgpu_*_x64_RelWithDebInfo jobs, Linux WebGPU x64 Release and both WGSL template jobs are green, so the 0.15/0.03 and 2e-2 bounds from 19f3697cf have been through a run now instead of being read off a log, which was the part I could not do from here.

The one red check is Lint, and it is not mine. Reviewdog's misspell action builds a debian:bullseye-slim image, and its apt-get update fails with Release file for http://deb.debian.org/debian-security/dists/bullseye-security/InRelease is expired (invalid since 1d 4h 37min). It retries three times and dies before any file of the PR is looked at. Optional Lint C++ and Python format in the same workflow are green.

That also retires what I wrote yesterday about the runs never starting: it was true of 19f3697cf, not of this head.

[Copilot]

🟡 Changes recommended

DP4A variants lack deterministic coverage, and the public documentation overstates the option’s scope.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Review details

• Files reviewed: 25/25 changed files
• Comments generated: 5
• Review effort level: Balanced

[Jiajia Qin (qjia7)]

Overall, this looks good to me. I only have a few cleanup comments.

Copilot’s documentation comments are valid: please scope the fused-MLP epilogue description to the decode fast path and clarify that this option currently applies only to the native WebGPU EP. JSEP is being deprecated, so I do not expect this option to be implemented or followed up there; the public documentation only needs to make the current limitation clear.

Regarding Copilot’s deterministic-coverage comments, I consider this a broader pre-existing testability limitation in WebGPU MatMulNBits. There are many implementations, and dispatch depends on the shape, output type, adapter capabilities, and vendor. We therefore cannot reliably select a particular algorithm across CI using shapes alone.

A better long-term solution is to provide a test mechanism that explicitly selects or forces the intended implementation and verifies which program was dispatched. That mechanism can then cover the small-M DP4A path and the 2/4/8-bit accumulator variants deterministically. I think this should be tracked as follow-up work and should not block this PR.

[Roberto (RobertoReale)]

Pushed b55316e96 with the three cleanups.

acc_f32 is now read once before the wide-tile branch and reused by both programs and both cache hints.

You are right about SDP8AI: it is declared -> acc_element_t, so the outer conversion was converting the result to the type it already had. That was 68 call sites in dp4a_matmul.wgsl.template and 4 in dp4a_matmul_small_m.wgsl.template. I put the unwrapped form through naga for the twelve combinations of output type, acc_f32 and n_bits the generator can emit, with a control that declares the accumulator f16 against an f32 return. That one is rejected, so the check discriminates rather than just passing everything.

On the docs, I scoped the epilogue paragraph to the decode fast path in both webgpu_provider_options.h and the JS option: ApplyUnfusedMlp creates gate_output and up_output with a->DataType() and runs SiLU on those, so the epilogue rounds at output precision there whatever the option says. I also checked the JSEP branch of session-options.ts before writing anything — it forwards preferredLayout and nothing else — and documented the option as native-EP-only.

Agreed on the coverage comments as follow-up. The dp4a conversation from the earlier round is still open for the same reason as before: A is int8-quantized ahead of the kernel, so I have no numeric case that separates the two accumulators on that path.

[Jiajia Qin (qjia7)]

LGTM if all CI pipelines are happy.
Hariharan Seshadri (@hariharans29) Please take a final look.

[Roberto (RobertoReale)]

The one red check on b55316e96 is web_Release / build_onnxruntime_web, and it is a runner dropout, not a test failure. The job's annotation is "The self-hosted runner lost communication with the server", and it hit at step 25, E2E package consuming test. Every test step before it passed in that same job: ort-web ALL backends, Suite1 webgpu with both IO-binding modes, WebGL packed mode, the WebAssembly proxy, and Run ort-web tests - WebGPU EP.

Two other runs of that job died the same way earlier today, on branches that are not mine: 34309494961 at 05:36 UTC and 34314249011 at 06:45 UTC. Same annotation, same step 25, and in both cases the job had run about 31 minutes.

For what it is worth on the js side, the only difference between 861788261, where the Web CI pipeline was green, and b55316e96 is the doc comment on enableMatmulFp32Accumulation in js/common/lib/inference-session.ts. No code changed there.

A re-run of that single job should clear it. I cannot trigger one from a fork.