fix(webgpu): use f32 accumulators in fp16 MatMul/MatMulNBits to prevent overflow

[RobertoReale]

Summary

Promotes the MatMul / MatMulNBits dot-product accumulators from f16 to f32 in the JSEP shader generators and the native WebGPU EP WGSL templates. Weights and activations stay f16/q4 in global memory — only register/workgroup accumulators change.

Correction (2026-07-14). This PR originally claimed to fix #26732 and listed Gemma 3 270M / SmolLM2 / kokoro-js as affected models. That attribution was wrong and I have retracted it — see the writeup on #29611. Root cause of #26732 is a model-level fp16 residual-stream overflow in Gemma 3, not the matmul accumulator: on the pre-mitigation revision (2950c41f) the residual crosses 65504 between layers 7 and 8, while the f16 running partial sum inside every one of the model's 126 MatMulNBits nodes peaks at ~324. That is why onnx-community's 2025-12-02 re-export fixed it with Clip(±32752) nodes on the residual adds. An f32 matmul accumulator does not address it. Details below.

The overflow this PR does fix

The WGSL shaders accumulate the dot-product inner loop in the output element type, which is f16 for fp16/q4f16 models. When the partial sums along K exceed 65504 they saturate to +Inf and propagate NaN through subsequent LayerNorm / Softmax. CPU/WASM is unaffected because MLAS accumulates internally in f32 even with f16 inputs.

Whether a given model actually hits this depends on its activation and weight magnitudes — it is not universal, and I no longer claim it is. The demonstrated case is below.

Evidence

whisper-small q4, end-to-end (full methodology): unusable on WebGPU on the released build (my Chrome extension force-disables WebGPU for q4 models because of it), and word-for-word identical to the WASM EP with this branch, ~2x faster. The accumulator promotion is the only change. This exercises the JSEP MatMulNBits path.

Portability of f16 accumulation (raw-WGSL probes): WGSL permits extra intermediate precision. Intel's D3D12 compiler evaluates straight-line/unrolled acc += chains at higher precision and rounds only across loop iterations; NVIDIA/Vulkan rounds strictly. So "f16 accumulation is safe here" is a property of vendor × backend × code shape, not of a vendor — which is why a vendor-string gate cannot express it, and why kernels that carry the accumulator across a K loop have no headroom guarantee on strict-rounding configurations.

Not claimed: I have no real model that overflows the native wide-tile accumulator specifically. @daijh is building A/B onnxruntime-web packages (f16 vs f32 accumulator) to settle that empirically across vendors; I'll run them on Intel Iris Xe (D3D12 + Vulkan) and NVIDIA RTX 3050 Ti and report the results, whichever way they go.

Fix details

JSEP (js/web/lib/wasm/jsep/webgpu/ops/):

• matmulnbits.ts: both kernels (default and BlockwiseMatMulNBits32) accumulate in f32 (workgroup_shared / inter_results), downcast to the output type only at the final write.
• matmul-shaders.ts (naive MatMul): values accumulator promoted to f32.
• 3rd-party/matmul_packed_webgpu.ts (tiled MatMul, vec4 + scalar variants): acc promoted to f32. Shared-memory tiles (mm_Asub/mm_Bsub) stay in the input dtype, so shared-memory usage and bandwidth are unchanged.

Explicit vec4() casts on each multiply operand are required — Dawn/D3D12 re-demotes temporaries to f16 when enable f16; is active, even if the accumulator variable is declared f32.

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

• matmul_nbits.wgsl.template: inter_results and the final reduction accumulate in f32 along K; downcast at the output write. This aligns the kernel with matmul_nbits_wide_tile.wgsl.template, which already accumulates in f32 on main.
• matmul_nbits_mlp.wgsl.template / matmul_nbits_qkv.wgsl.template (fused MLP / QKV, added per review feedback): gate/up and q/k/v projection sums plus the workgroup inter_results accumulate in f32; bias and SiLU applied in f32; downcast only at the final store.
• dp4a_matmul_small_m.wgsl.template / dp4a_matmul.wgsl.template: the cross-tile accumulators (inter_results, lane_outputs / lane_output1..4) accumulate in f32. SDP8AI in dp4a_matmul_common.wgsl.template is unchanged — it is an exact integer dot4I8Packed dot product; the accumulation across K tiles is the path that can saturate.
• Intentionally excluded: subgroup_matrix_matmul_nbits_* (accumulator precision is tied to the hardware cooperative-matrix configuration negotiated on the C++ side — see review thread).

Global memory is unchanged: weights and activations remain in their original dtype (f16 / q4) in global GPU memory, so bandwidth is unaffected. Shared memory is not entirely free, and the original wording here overstated that: the accumulator workgroup arrays (inter_results, workgroup_shared) do double in size for fp16/q4f16 outputs, which can lower occupancy on SLM-bound dispatch configurations. The A/B tiles stay in the input dtype.

Testing

• WGSL validation: every modified shader pattern compile-checked with naga (wgpu-native) in both output_element_t = f16 (with enable f16;) and f32 variants, including the vector conversion constructors at the final stores.
• cd js/web && npx tsc --noEmit: clean; prettier: clean.
• End-to-end: whisper-small q4 through onnxruntime-web built from this branch (above).
• MatMulNBits.Float16_LargeK_AccumulatorOverflow in matmul_4bits_test.cc: A=8 with dequantized B at +112 over the first half of K and -112 over the second, K=8192, M=1 so the dispatch lands on matmul_nbits.wgsl.template. A lane peaks at 114688 halfway along K and cancels back to an exact 0, so the f32 accumulator returns exactly 0 while an f16 accumulator 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. Note it is a regression guard, not a universal detector: Intel's D3D12 compiler promotes unrolled f16 chains on its own, so it can pass there with an f16 accumulator.

Performance

f32 accumulators have negligible ALU overhead on most GPUs; memory bandwidth — the actual bottleneck in LLM inference — is unaffected. @daijh reports measurable register-pressure cost on Intel, which is the open question the A/B packages are meant to answer. If the cost is real and the correctness risk is not, accuracy_level = 2 as an explicit opt-down (default staying f32, mirroring the existing dp4a gate on level 4) is the compromise I'd propose, and I'm happy to implement it here.

Related: #26732, #26367, #29611

[RobertoReale]

@microsoft-github-policy-service agree

[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

[tianleiwu]

@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

[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.

[tianleiwu]

This correctly fixes the generic kernel, but the same f16-accumulation pattern is left unfixed in the fused native kernels that get dispatched for the very models this PR targets:

• 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
• dp4a_matmul_small_m.wgsl.template: inter_results : array<array<output_element_t, ...>>

The MLP/QKV kernels are selected for fused MLP / QKV subgraphs (common in Gemma 3 / SmolLM), so those paths can still overflow. Please extend the f32 accumulation to them or note explicitly why they are excluded.

[tianleiwu]

Please drop this file from the PR. It introduces a new top-level scripts/ directory for a static regex check over shader source (not a runtime test), it already fails lintrunner (RUFF-FORMAT), and it will silently bit-rot the moment the shaders are refactored because it asserts on exact source substrings. If you want regression coverage, prefer a numerical unit test (large-K fp16 MatMul/MatMulNBits vs. an f32 reference) in the existing WebGPU op test suite, which exercises real runtime behavior.

[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.

[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?

[RobertoReale]

@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.

[RobertoReale]

@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.

[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.

[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?

[RobertoReale]

@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.

[daijh]

@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.

@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?

[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
}

[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.

[RobertoReale]

Thanks @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.

[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.

[RobertoReale]

Thanks @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.

[RobertoReale]

@tianleiwu @qjia7 @hariharans29 @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 @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 @daijh's A/B packages report), say so and I'll do it.

[RobertoReale]

Two things I owed this PR from @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.