Skip to content

feat(webgpu): add preferredMatmulAccumulatorPrecision provider option for MatMulNBits - #29599

Open
Roberto (RobertoReale) wants to merge 13 commits into
microsoft:mainfrom
RobertoReale:fix/webgpu-f16-overflow
Open

feat(webgpu): add preferredMatmulAccumulatorPrecision provider option for MatMulNBits#29599
Roberto (RobertoReale) wants to merge 13 commits into
microsoft:mainfrom
RobertoReale:fix/webgpu-f16-overflow

Conversation

@RobertoReale

@RobertoReale Roberto (RobertoReale) commented Jul 7, 2026

Copy link
Copy Markdown

Summary

Adds the WebGPU EP provider option preferredMatmulAccumulatorPrecision ("f16" default, "f32" opt-in) proposed by Jiajia Qin (@qjia7) on #29611, 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). 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.

What this is now

An opt-in, not a default change. The shipped default is "f16", which is what #29611 landed. Setting "f32" promotes the accumulators carried along K in the six MatMulNBits kernels below. The block-local dot products keep the precision they have today either way, so the f16 setting rounds at the block boundary — the same shape that was measured in the Gemma writeup.

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 also covers the fp16 MatMul path, which has no accuracy_level to hang off.

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.

Scope

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.
  • subgroup_matrix_matmul_nbits_* cannot honour an f32 request: 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 &#34;f32&#34; 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.
  • Plumbing follows the existing kvCacheQuantizationBits option: constant and 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 in js/common/lib/inference-session.ts, validated and forwarded to the EP in js/web/lib/wasm/session-options.ts, and accepted by the node bridge in js/node/src/session_options_helper.cc, 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 set to &#34;f32&#34;. A=8, 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 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. The fused MLP and QKV decode kernels are reached through graph fusion rather than a MatMulNBits node and are not covered by it.
  • 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.
  • Not compiled. I have no ORT build tree available, so the C++ has not been through a compiler and the tests have never run. That is the main thing I need CI, or a reviewer with a build, to tell me.

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 default here stays &#34;f16&#34;, nobody pays for it unless they ask.

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

…nt overflow

Problem:
- fp16 and q4f16 models (Gemma 3 270M, Whisper Q4, SmolLM2, kokoro-js) produce
  garbage output on WebGPU. f16 max is ~65504; summing 2048+ dot-product terms
  overflows to +Inf, which propagates as NaN through LayerNorm/Softmax.
- CPU/WASM path is unaffected (MLAS accumulates in f32 internally).

Fix (JSEP path):
- matmulnbits.ts: both kernels (default and BlockwiseMatMulNBits32) now
  accumulate in f32 (workgroup_shared / inter_results). Explicit vec<f32>()
  casts per operand are required: Dawn/D3D12 re-demotes temporaries to f16
  when 'enable f16;' is active. Output downcast to f16 only at the 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; tiles (mm_Asub/mm_Bsub) stay in f16, so there is no
  shared-memory or bandwidth regression.

Fix (native WebGPU EP):
- contrib_ops/webgpu/quantization/matmul_nbits.wgsl.template: inter_results
  and the final reduction accumulate in f32 along K; downcast at output write.
  This aligns the kernel with matmul_nbits_wide_tile.wgsl.template, which
  already accumulates in f32.

Testing:
- scripts/verify_f16_fix.py: static source check - all checks PASS
- cd js/web && npx tsc --noEmit: clean; prettier: clean
- Manual browser test on Chrome 121+ with Gemma 3 270M FP16 and
  SmolLM2 360M Q4F16 via transformers.js

Performance note: f32 accumulators only affect register-level computation;
weights/activations remain f16/q4 in GPU memory, so memory bandwidth (the
actual bottleneck) is unaffected.

Fixes microsoft#26732
Related: microsoft#26367
@RobertoReale

Copy link
Copy Markdown
Author

@microsoft-github-policy-service agree

@RobertoReale

Copy link
Copy Markdown
Author

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

Comment thread scripts/verify_f16_fix.py Fixed
@tianleiwu

Copy link
Copy Markdown
Contributor

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

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 Tianlei Wu (tianleiwu) left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.templategate_sum/up_sum = output_element_t(0) and gate_inter_results/up_inter_results : array<array<output_element_t, ...>>
  • matmul_nbits_qkv.wgsl.templatesum = q_output_element_t(0) and q/k/v_inter_results : array<array<q_output_element_t, ...>>
  • dp4a_matmul_small_m.wgsl.templateinter_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.

// Accumulate partial sums along K in f32: with fp16 outputs, summing 2048+ f16
// products overflows the f16 max (65504) and poisons the output with +Inf/NaN
// (issue #26732). Only the block-local `sum` stays in output_element_t.
var<workgroup> inter_results: array<array<f32, tile_size_k_vec>, tile_size>;

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Comment thread scripts/verify_f16_fix.py Outdated
@@ -0,0 +1,116 @@
#!/usr/bin/env python3

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Copy Markdown
Contributor

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

Copy link
Copy Markdown
Contributor

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?

…els, drop verify script

- matmul_nbits_mlp.wgsl.template: accumulate gate/up projection sums and
  workgroup inter_results in f32; apply bias and SiLU activation in f32;
  downcast to output_element_t only at the final store.
- matmul_nbits_qkv.wgsl.template: accumulate q/k/v projection sums and
  workgroup inter_results in f32; downcast at the final store.
- dp4a_matmul_small_m.wgsl.template: accumulate inter_results and the
  final reduction in f32 (SDP8AI partial products are exact integer dots;
  the cross-tile accumulation over K was the overflow path); downcast at
  the final store. dp4a_matmul_common.wgsl.template is shared with other
  kernels and is intentionally left unchanged.
- Remove scripts/verify_f16_fix.py per review feedback (static source
  check, not a runtime test; also failed lintrunner RUFF-FORMAT).

Weights, activations and all buffer traffic remain fp16/q4; only
register/workgroup accumulators are promoted, so memory bandwidth is
unchanged.
Same overflow pattern as dp4a_matmul_small_m: per-lane register
accumulators (lane_outputs / lane_output1..4) summed output_element_t
partial dot products across the whole K loop. Promote them to f32 and
downcast to output_element_t only at the final store; SDP8AI itself and
all buffer types are unchanged.
@RobertoReale

Copy link
Copy Markdown
Author

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.templategate_sum/up_sum, gate/up_inter_results, bias and SiLU activation now computed in f32.
  • matmul_nbits_qkv.wgsl.templatecompute_projection_sum and q/k/v_inter_results in f32.
  • dp4a_matmul_small_m.wgsl.templateinter_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

Copy link
Copy Markdown
Author

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

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.

WGSL has no implicit conversions: constructing vec4<output_element_t>
directly from f32 scalar components is a compile error when
output_element_t is f16. Build a vec4<f32> first, then use the
whole-vector conversion constructor. Verified with naga (wgpu-native):
the direct mixed-scalar constructor is rejected, the vector-to-vector
conversion compiles for both f16 and f32 output types.
@RobertoReale

Copy link
Copy Markdown
Author

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

Copy link
Copy Markdown
Contributor

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

Copy link
Copy Markdown
Author

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

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

Copy link
Copy Markdown
Contributor

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

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?

@RobertoReale

Copy link
Copy Markdown
Author

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

Copy link
Copy Markdown
Contributor

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

Copy link
Copy Markdown
Author

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.

@daijh

Copy link
Copy Markdown
Contributor

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

Copy link
Copy Markdown
Author

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.

@RobertoReale Roberto (RobertoReale) changed the title fix(webgpu): use f32 accumulators in fp16 MatMul/MatMulNBits to prevent overflow (#26732) fix(webgpu): use f32 accumulators in fp16 MatMul/MatMulNBits to prevent overflow Jul 14, 2026
@RobertoReale

Copy link
Copy Markdown
Author

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.

Every MatMulNBits test in matmul_4bits_test.cc generates its inputs from
Gaussian(0, 0.25), which keeps the running partial sum along K in the single
digits even at K=11008, so none of them can reach the f16 ceiling of 65504.

Add a case where the partial sum crosses it while the inputs and the exact
result stay inside f16 range: A=8, dequantized B=+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 through K and
cancels back to an exact 0. With an f32 accumulator the output is exactly 0;
with an f16 accumulator it saturates to +Inf.

The block-local sum stays in output_element_t and peaks at 28672, so the test
isolates the accumulator carried across K.
@RobertoReale

Copy link
Copy Markdown
Author

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.

…tion

Adds the WebGPU EP provider option "preferredMatmulAccumulatorPrecision"
("f16" / "f32") proposed in microsoft#29611, and makes every MatMulNBits kernel this
branch touches read it instead of hardcoding f32.

Plumbing follows the existing kvCacheQuantizationBits option: constant and
value strings in webgpu_provider_options.h, parsing in
webgpu_provider_factory.cc, config field plus accessor on the EP, and an
inline forwarder on ComputeContext so kernels can read it.

The five WGSL templates gain a compile-time `acc_f32` parameter selecting
`alias acc_element_t = f32` or `= output_element_t`, and the accumulators,
their reduction and the bias adds use that alias. Only the accumulator moves:
the block-local dot products keep the precision they have today, so the f16
setting rounds at the block boundary, which is the same shape measured in the
Gemma writeup. acc_f32 is added to every affected CacheHint so the two
variants cannot collide in the program cache.

The shipped default lives on a single line
(WebGpuExecutionProviderConfig::matmul_accumulator_precision_f32) and is f32
here, which keeps current behaviour; flipping it to f16 is a one-line change
if that is where the review lands.

Not compiled: no build tree available. Verified by running the in-tree
tools/python/wgsl_template generator over contrib_ops/webgpu, which parses all
five templates and emits both alias branches for each. The JSEP shaders are
untouched: they read env.webgpu and never see EP provider options, so this
option does not reach them.
Matches the default agreed on microsoft#29611: the option ships opt-in to f32, the
native kernels keep accumulating in output_element_t unless asked otherwise.
Roberto (RobertoReale) added a commit to RobertoReale/Voice-Message-Transcriber that referenced this pull request Aug 4, 2026
The exclusion has been there since the first commit, on the grounds that
ONNX Runtime's WebGPU backend returned "[Music]" instead of speech on q4.
Re-measured against the pinned stack, which has not moved since day one
(transformers 3.8.1 -> onnxruntime-web 1.22.0-dev.20250409): on
whisper-small q4, three real voice messages of 26 / 33 / 39 s, WebGPU
output is character-identical to WASM and 1.5x faster on an Intel Iris Xe,
more on a discrete GPU. The failure does not reproduce.

q8 stays excluded: its crash was never retested and no model here uses it.
The crash flag and the WASM fallback are untouched.

docs/webgpu-onnxruntime-fix.md rewritten. It claimed the cause was an f16
accumulator overflow that microsoft/onnxruntime#29599 would fix; both
halves were wrong. whisper-small q4 contains zero fp16 tensors, so that PR
is inert on this model, and the July A/B compared two builds a release
cycle apart rather than the patch. What actually fixed it is not known -
neither the packages nor the model changed - so the re-enable rests on the
measurement, not on a theory.
Float16_LargeK_AccumulatorOverflow was written when this branch hardcoded the
f32 accumulator, so it built the EP with no provider options. Now that the
option ships defaulting to "f16", a plain DefaultWebGpuExecutionProvider()
accumulates in output_element_t and the lane saturates at 114688 exactly as the
comment describes: the test would fail on the shipped default, and fail for the
right reason, which makes it useless as a guard.

Build the EP with preferredMatmulAccumulatorPrecision = "f32" instead, using
WebGpuExecutionProviderWithOptions the same way Float32_Large already does a few
lines above. The test now guards the opt-in path, which is the one that is
supposed to be exact.
The three JSEP shader generators (matmulnbits.ts, matmul-shaders.ts,
3rd-party/matmul_packed_webgpu.ts) go back to what main has. Two reasons, and
the first one is the one that forces it.

The provider option is a WebGPU EP option. JSEP shaders read env.webgpu and
never see EP provider options, so the option cannot reach them. Leaving those
hunks in would ship a PR where the native kernels are opt-in and the JSEP path
is switched to f32 unconditionally for everyone, which is not a coherent thing
to ask a reviewer to approve, and is not what was agreed on microsoft#29611.

The second reason is that plumbing a parallel env.webgpu flag through JSEP to
restore the symmetry means adding an option to a path microsoft#29716 is removing. If the
JSEP path needs this before it goes away, it is a separate change with its own
justification.

This is the retraction of the whisper-small q4 evidence carried through to the
diff: that measurement was the reason the JSEP hunks were here, and it did not
hold up (whisper-small q4 has no fp16 tensor at all, so the hunks were inert on
the model they were credited to).
@RobertoReale Roberto (RobertoReale) changed the title fix(webgpu): use f32 accumulators in fp16 MatMul/MatMulNBits to prevent overflow feat(webgpu): add preferredMatmulAccumulatorPrecision provider option for MatMulNBits Aug 6, 2026
@RobertoReale

Copy link
Copy Markdown
Author

Jiajia Qin (@qjia7) Jianhui Dai (@daijh) Tianlei Wu (@tianleiwu) Hariharan Seshadri (@hariharans29) — I've pushed the rework we settled on over on #29611, and rewritten the description to match. Now that #29611 has merged, that thread isn't the right place for this any more, so I'm moving it back here.

This is now the preferredMatmulAccumulatorPrecision provider option, "f16" default and "f32" opt-in, wired into matmul_nbits, _mlp, _qkv, dp4a_matmul and dp4a_matmul_small_m through a compile-time acc_f32 template parameter, with acc_f32 in every affected CacheHint. Plumbing follows kvCacheQuantizationBits. main's behaviour is unchanged.

Two things I decided rather than kept asking about. The JSEP hunks are gone — the files are back to main. I'd asked whether to drop them and never got an answer before that thread closed, but the option is an EP provider option and JSEP shaders only ever see env.webgpu, so keeping them would have shipped a PR that's opt-in on one side and unconditional on the other. Adding a parallel flag to a path #29716 is deleting didn't seem worth it. And Float16_LargeK_AccumulatorOverflow now builds its EP with the option set to "f32" — as written it would have failed on the shipped default, correctly.

The honest gap: I still have no ORT build tree here, so none of this has seen a compiler and the test has never run. The templates go through wgsl_gen.py cleanly and emit both alias branches, which is as far as I can check on my own.

If the shape is wrong, say so and I'll redo it.

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Adds a new WebGPU Execution Provider option, preferredMatmulAccumulatorPrecision ("f16" default, "f32" opt-in), and plumbs it through the native WebGPU MatMulNBits family so kernels can select f16 vs f32 accumulation across the K loop. This is a targeted numerical-stability/portability knob for strict-rounding backends, while keeping the shipped default behavior unchanged.

Changes:

  • Introduces and parses the new provider option, storing it on the WebGPU EP config and exposing it via ComputeContext.
  • Threads an acc_f32 compile-time parameter through multiple MatMulNBits/DP4A WGSL templates and incorporates it into cache hints to avoid shader-cache collisions.
  • Adds a WebGPU-focused unit test that exercises accumulator overflow behavior at large K when f16 accumulation is used.

Reviewed changes

Copilot reviewed 17 out of 17 changed files in this pull request and generated 6 comments.

Show a summary per file
File Description
onnxruntime/test/contrib_ops/matmul_4bits_test.cc Adds a regression test that requires the f32 accumulator option to avoid Inf/NaN on strict-rounding backends.
onnxruntime/core/providers/webgpu/webgpu_provider_options.h Defines the new provider option key and its "f16"/"f32" values.
onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc Parses preferredMatmulAccumulatorPrecision into the EP config.
onnxruntime/core/providers/webgpu/webgpu_execution_provider.h Adds config storage + accessor for the chosen accumulator precision.
onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc Wires the parsed config into the EP instance.
onnxruntime/core/providers/webgpu/compute_context.h Exposes the EP setting to kernels via ComputeContextBase.
onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.wgsl.template Adds acc_f32 template param and switches the cross-K accumulator storage/reduction accordingly.
onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.h Extends MatMulNBitsProgram to carry acc_f32 into shader generation.
onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc Plumbs ComputeContext option into program construction and cache hint.
onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_qkv.wgsl.template Adds acc_f32 param and updates accumulator storage for QKV fused decode path.
onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_qkv.cc Passes option into template expansion and cache hints for the QKV decode program.
onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_mlp.wgsl.template Adds acc_f32 param and updates accumulator storage for MLP fused decode path.
onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_mlp.cc Passes option into template expansion and cache hints for the MLP decode program.
onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul.wgsl.template Adds acc_f32 param and changes accumulator types for DP4A matmul variants.
onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul_small_m.wgsl.template Adds acc_f32 param and adjusts cross-K accumulator and final store casting.
onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul_nbits.h Extends DP4A program classes to carry acc_f32 for shader generation.
onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul_nbits.cc Plumbs ComputeContext option into DP4A program construction and cache hints.
Suppressed comments (3)

onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_qkv.cc:44

  • acc_f32_ is used as a WGSL template parameter, but it is never initialized from the acc_f32 constructor argument. That leaves the shader variant selection non-deterministic and can also desync the program-cache key from the generated shader code.
      : Program{"MatMulNBitsQkvDecode"},
        tile_size_(tile_size),
        single_scale_weights_(single_scale_weights),
        tile_size_k_vec_(tile_size_k_vec),
        k_unroll_tiles_(k_unroll_tiles),
        has_norm_(has_norm),
        has_skip_input_(has_skip_input),
        has_skip_output_(has_skip_output) {

onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits_mlp.cc:75

  • acc_f32_ is referenced as a WGSL template parameter but is never initialized from the acc_f32 constructor argument. This makes shader variant selection and caching non-deterministic.
                              bool acc_f32)
      : Program{"MatMulNBitsMlpDecode"},
        tile_size_(tile_size),
        has_gate_bias_(has_gate_bias),
        has_up_bias_(has_up_bias),
        has_norm_input_(has_norm_input),
        has_skip_input_(has_skip_input),
        has_skip_output_(has_skip_output),
        single_scale_weights_(single_scale_weights),
        tile_size_k_vec_(tile_size_k_vec),
        k_unroll_tiles_(k_unroll_tiles),
        activation_kind_(activation_kind) {}

onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul.wgsl.template:508

  • Bias addition is currently done after casting lane_output* down to output_element_t, so enabling acc_f32 doesn't actually keep the bias add in f32. Add bias in acc_element_t and downcast only once at the final store.
        output.setByOffset(output_idx, vec4<output_element_t>(lane_output1) + bias_vec1);
        output.setByOffset(output_idx+1, vec4<output_element_t>(lane_output2) + bias_vec2);
        output.setByOffset(output_idx+2, vec4<output_element_t>(lane_output3) + bias_vec3);
        output.setByOffset(output_idx+3, vec4<output_element_t>(lane_output4) + bias_vec4);

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

Comment on lines +25 to +33
// Precision of the partial sums accumulated along K, selected by the
// "preferredMatmulAccumulatorPrecision" provider option. The block-local dot products stay in
// f32 either way; this only changes the accumulator carried across K, which is what can
// overflow the f16 max (~65504).
#if acc_f32
alias acc_element_t = f32;
#else
alias acc_element_t = output_element_t;
#endif
Comment on lines +468 to +471
output.setByOffset(output_idx, vec4<output_element_t>(vec4<acc_element_t>(lane_outputs[0], lane_outputs[1], lane_outputs[2], lane_outputs[3])) + bias_vec1);
output.setByOffset(output_idx+1, vec4<output_element_t>(vec4<acc_element_t>(lane_outputs[4], lane_outputs[5], lane_outputs[6], lane_outputs[7])) + bias_vec2);
output.setByOffset(output_idx+2, vec4<output_element_t>(vec4<acc_element_t>(lane_outputs[8], lane_outputs[9], lane_outputs[10], lane_outputs[11])) + bias_vec3);
output.setByOffset(output_idx+3, vec4<output_element_t>(vec4<acc_element_t>(lane_outputs[12], lane_outputs[13], lane_outputs[14], lane_outputs[15])) + bias_vec4);
Comment on lines 64 to +78
@@ -70,7 +74,8 @@ class DP4AMatMulNBitsSmallMProgram final : public Program<DP4AMatMulNBitsSmallMP
has_weight_idx_(has_weight_idx),
has_weight_idx_indirect_(has_weight_idx_indirect),
single_scale_weights_(single_scale_weights),
broadcast_a_row_(broadcast_a_row) {}
broadcast_a_row_(broadcast_a_row),
acc_f32_(acc_f32) {}
Comment on lines +44 to +45
MatMulNBitsProgram(uint32_t tile_size, uint32_t nbits, bool has_zero_points, bool has_bias, bool has_weight_idx, bool has_weight_idx_indirect, bool single_scale_weights, uint32_t tile_size_k_vec = 16, bool broadcast_a_row = false, bool acc_f32 = true)
: Program{"MatMulNBits"}, tile_size_(tile_size), nbits_(nbits), has_zero_points_(has_zero_points), has_bias_(has_bias), has_weight_idx_{has_weight_idx}, has_weight_idx_indirect_{has_weight_idx_indirect}, single_scale_weights_(single_scale_weights), tile_size_k_vec_(tile_size_k_vec), broadcast_a_row_(broadcast_a_row), acc_f32_(acc_f32) {}
Comment on lines 106 to 120
@@ -106,15 +116,15 @@ fn compute_projection_sum(weight: q_b_value_t,
#if component_a == 1
let a0 = load_a_vec4(a_offset);
let a1 = load_a_vec4(a_offset + 4);
sum += dot(a0, w0) + dot(a1, w1);
sum += dot(vec4<f32>(a0), vec4<f32>(w0)) + dot(vec4<f32>(a1), vec4<f32>(w1));
#elif component_a == 2
Comment on lines +77 to +115
@@ -76,8 +86,8 @@ fn compute_gate_up_sums(b_global: u32, kidx: u32, idx: u32, k_offset: u32) -> ve
let gate_b_value = gate_b.getByOffset(b_global * uniforms.K_of_b + k_offset);
let up_b_value = up_b.getByOffset(b_global * uniforms.K_of_b + k_offset);

var gate_sum = output_element_t(0);
var up_sum = output_element_t(0);
var gate_sum = f32(0);
var up_sum = f32(0);
var a_offset = idx * (8 / component_a) * component_b;
#if component_b == 1
let gate_b_value_lower = vec4<output_element_t>(unpack4xU8(gate_b_value & 0x0F0F0F0Fu)) - vec4<output_element_t>(default_zero_point);
@@ -91,18 +101,18 @@ fn compute_gate_up_sums(b_global: u32, kidx: u32, idx: u32, k_offset: u32) -> ve
#if component_a == 1
let a0 = vec4<output_element_t>(tile_A[a_offset], tile_A[a_offset + 1], tile_A[a_offset + 2], tile_A[a_offset + 3]);
let a1 = vec4<output_element_t>(tile_A[a_offset + 4], tile_A[a_offset + 5], tile_A[a_offset + 6], tile_A[a_offset + 7]);
gate_sum += dot(a0, gate_b0) + dot(a1, gate_b1);
up_sum += dot(a0, up_b0) + dot(a1, up_b1);
gate_sum += dot(vec4<f32>(a0), vec4<f32>(gate_b0)) + dot(vec4<f32>(a1), vec4<f32>(gate_b1));
up_sum += dot(vec4<f32>(a0), vec4<f32>(up_b0)) + dot(vec4<f32>(a1), vec4<f32>(up_b1));
#elif component_a == 2
let a0 = vec4<output_element_t>(tile_A[a_offset], tile_A[a_offset + 1]);
let a1 = vec4<output_element_t>(tile_A[a_offset + 2], tile_A[a_offset + 3]);
gate_sum += dot(a0, gate_b0) + dot(a1, gate_b1);
up_sum += dot(a0, up_b0) + dot(a1, up_b1);
gate_sum += dot(vec4<f32>(a0), vec4<f32>(gate_b0)) + dot(vec4<f32>(a1), vec4<f32>(gate_b1));
up_sum += dot(vec4<f32>(a0), vec4<f32>(up_b0)) + dot(vec4<f32>(a1), vec4<f32>(up_b1));
#elif component_a == 4
let a0 = tile_A[a_offset];
let a1 = tile_A[a_offset + 1];
gate_sum += dot(a0, gate_b0) + dot(a1, gate_b1);
up_sum += dot(a0, up_b0) + dot(a1, up_b1);
gate_sum += dot(vec4<f32>(a0), vec4<f32>(gate_b0)) + dot(vec4<f32>(a1), vec4<f32>(gate_b1));
up_sum += dot(vec4<f32>(a0), vec4<f32>(up_b0)) + dot(vec4<f32>(a1), vec4<f32>(up_b1));

@qjia7 Jiajia Qin (qjia7) left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for narrowing this change to an opt-in native WebGPU EP option. Exposing the accumulator choice through the provider configuration is preferable to changing precision globally. I found several correctness issues and some missing shader-path coverage.

1. Initialize acc_f32_ in the QKV and MLP decode programs

MatMulNBitsQkvDecodeProgram and MatMulNBitsMlpDecodeProgram accept bool acc_f32, but their constructors do not assign it to acc_f32_.

GenerateShaderCode() later uses that member:

WGSL_TEMPLATE_PARAMETER(acc_f32, acc_f32_)

This means shader generation reads an uninitialized value. It can also cause the generated WGSL variant to disagree with the acc_f32 value included in CacheHint.

Please add acc_f32_(acc_f32) to both constructor initializer lists.

This is a pre-merge blocker.


2. Fix the undefined output_element_t alias in the QKV template

The false branch in the QKV template currently uses:

#if acc_f32
alias acc_element_t = f32;
#else
alias acc_element_t = output_element_t;
#endif

However, the program declares outputs named q_output, k_output, and v_output. Their generated aliases should be q_output_element_t, k_output_element_t, and v_output_element_t. There is no output named output that would generate output_element_t.

Consequently, the default accumulator branch can produce invalid WGSL.

Please use an available output alias—likely q_output_element_t if all three outputs are guaranteed to have the same element type—or introduce an explicit common alias.

This is a pre-merge blocker.


3. Preserve f32 precision through the DP4A bias addition

The DP4A bias paths convert accumulated values to output_element_t before adding the bias:

vec4<output_element_t>(vec4<acc_element_t>(...)) + bias_vec

When acc_f32=true and the output type is fp16, this rounds or saturates the accumulated result to fp16 before adding the bias. The f32 accumulator option is therefore only partially effective.

Please convert the bias to acc_element_t, perform the addition in accumulator precision, and convert to the output type only at the final store:

f32 accumulator + f32 bias -> final output conversion

The same issue appears in both DP4A output-store blocks.


4. Consider exposing this as a boolean provider option

preferredMatmulAccumulatorPrecision = "f16" | "f32" is more general than the behavior implemented here. The actual choice is binary:

  • enable f32 accumulation; or
  • use the normal accumulator type, which follows the output element type.

Consider naming the option:

enableMatmulNBitsF32Accumulator

with the provider key:

ep.webgpuexecutionprovider.enableMatmulNBitsF32Accumulator

Including MatmulNBits clarifies that this affects the MatMulNBits family rather than every WebGPU MatMul kernel.

Describing the false case as "f16" is also misleading when the output is fp32. The false behavior means “follow the output type,” not “force f16.”


5. Make constructor defaults consistent with the provider default

Some program constructors currently use:

bool acc_f32 = true

while the provider configuration defaults the option to false.

Existing call sites may pass the value explicitly, but the mismatch creates a future footgun. A new call site could silently enable f32 accumulation without consulting the provider configuration.

Please either:

  • remove the default argument and require every caller to provide the option; or
  • change the default to false.

Removing the default is safer because it makes incomplete option plumbing a compile-time error.


6. Clarify the precision of block-local dot products

The QKV and MLP templates make their block-local helpers return f32 and explicitly convert operands to vec4<f32>, including when acc_f32=false.

That appears broader than the PR description, which states that block-local dot products retain their previous precision and only the accumulator carried across K changes.

If always-f32 block-local dot products are intentional, please update the PR description and shader comments and provide performance and numerical validation for the default path. Otherwise, these helpers should preserve output-type computation when acc_f32=false.


7. Add focused tests for every accumulator-dependent shader path

Please add small WebGPU tests that instantiate and execute both acc_f32=false and acc_f32=true for every affected implementation.

The coverage should include:

  • generic MatMulNBits
  • fused QKV
  • fused MLP
  • DP4A regular path
  • DP4A small-M path
  • wide-tile path
  • subgroup-matrix path
  • bias and no-bias variants where applicable
  • default/output-type accumulation and f32 accumulation

Small deterministic inputs should be sufficient. Each case should verify:

  1. session and pipeline creation succeeds;
  2. the selected WGSL variant compiles;
  3. dispatch completes;
  4. output is finite and numerically reasonable;
  5. changing the provider option selects the correct cached shader variant.

This coverage would detect:

  • undefined aliases in an individual WGSL branch;
  • missing acc_f32_ initialization;
  • incomplete provider-option propagation;
  • an affected implementation that ignores the option;
  • cache hints that do not distinguish accumulator variants;
  • shader compilation failures limited to wide-tile or subgroup-matrix paths.

The existing large-K cancellation test can remain as the focused numerical regression demonstrating why f32 accumulation is needed. The smaller tests would complement it by ensuring every generated shader variant compiles and executes.

Verdict

I would request changes before approval.

The primary blockers are:

  1. uninitialized acc_f32_ in the QKV and MLP programs;
  2. undefined output_element_t in the QKV default branch;
  3. premature downcasting in the DP4A bias path;
  4. missing compilation and runtime coverage across all affected implementations, including wide-tile and subgroup-matrix paths.

The provider-option design, constructor defaults, and block-local dot-product precision should also be resolved or documented before merge.

for (var i = 0u; i < 16u; i++)
{
lane_outputs[i] += SDP8AI(own_a0, subgroupShuffle(own_b0, i), own_a1, subgroupShuffle(own_b1, i), subgroupShuffle(own_scale_b, i) * own_scale_a, subgroupShuffle(zero, i));
lane_outputs[i] += acc_element_t(SDP8AI(own_a0, subgroupShuffle(own_b0, i), own_a1, subgroupShuffle(own_b1, i), subgroupShuffle(own_scale_b, i) * own_scale_a, subgroupShuffle(zero, i)));

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is it better if we change SDP8AI definition to return acc_element_t type instead of output_element_t type so that the whole code changes are small?

}

fn compute_gate_up_sums(b_global: u32, kidx: u32, idx: u32, k_offset: u32) -> vec2<output_element_t> {
fn compute_gate_up_sums(b_global: u32, kidx: u32, idx: u32, k_offset: u32) -> vec2<f32> {

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why f32 not acc_element_t?

idx: u32) -> q_output_element_t {
var sum = q_output_element_t(0);
idx: u32) -> f32 {
var sum = f32(0);

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Similar here.

…in QKV template, bias downcast ordering in dp4a, SDP8AI return type, constructor defaults

matmul_nbits_qkv.wgsl.template has no single "output" binding (q_output/k_output/v_output
instead), so acc_element_t's f16 branch aliased to an undefined output_element_t and would
fail shader compilation on the shipped default ("f16"). Aliases to q_output_element_t instead,
since all three outputs share one dtype (single T1 constraint).

dp4a_matmul.wgsl.template added bias after downcasting the accumulator to output_element_t,
which could saturate to f16 range before the bias add even when acc_f32 selected an f32
accumulator, defeating the point of the option. Bias is now added in acc_element_t and the
result downcast to output_element_t only at the store, in both the is_qualcomm and
non-qualcomm paths.

SDP8AI in dp4a_matmul_common.wgsl.template downcast its result to output_element_t before
returning, so for n_bits==8 (where mul_precision is f32) an individual dot-product term could
already saturate before ever reaching the accumulator, regardless of acc_f32. It now returns
acc_element_t directly; every call site already wraps the call in acc_element_t(...), so this
needed no caller changes.

MatMulNBitsProgram and DP4AMatMulNBitsSmallMProgram defaulted their acc_f32 constructor
parameter to true, inconsistent with the shipped provider-option default of "f16". Current call
sites pass the value explicitly, but the default is corrected to match.

Validated with tools/python/wgsl_gen.py (exit 0, both alias branches emitted for all five
templates) and with naga/wgpu-native on the four changed patterns in both acc_f32 states, plus
a negative control confirming the harness rejects an undefined identifier.
@RobertoReale

Copy link
Copy Markdown
Author

Pushed (b5e48ef77) fixing three real problems Copilot and Jiajia Qin (@qjia7) found, plus one latent one I found while looking at the first.

matmul_nbits_qkv.wgsl.template has no single output binding — it's q_output/k_output/v_output — so alias acc_element_t = output_element_t; referenced an identifier that doesn't exist in that shader. That would have failed to compile on the shipped default ("f16"). Now aliases to q_output_element_t; all three outputs share one T1 type constraint, so any of the three works.

On the bias ordering in dp4a_matmul.wgsl.template: correct, the store was downcasting lane_outputs/lane_output1..4 to output_element_t before adding bias, in both the qualcomm and non-qualcomm paths. That meant an accumulator sum safely inside f32 range could saturate to f16 the moment bias was added, defeating the reason to set acc_f32 in the first place. Bias now adds in acc_element_t, downcast only at the store.

Jiajia Qin (@qjia7) on SDP8AI: yes, and it turned up something in the same family. It downcast to output_element_t before returning, so for n_bits==8 (mul_precision = f32 there) a single dot-product term could already saturate before reaching the accumulator, independent of acc_f32. It now returns acc_element_t directly; every call site already wraps it in acc_element_t(...), so no caller changes needed.

On compute_gate_up_sums/compute_projection_sum staying f32: intentional, not an oversight — that's the block-local dot product over a handful of dequantized values, bounded by block_size, not the cross-K accumulator acc_f32 is meant to gate. Left as is.

Still not compiled — no build tree here. Checked with wgsl_gen.py (both alias branches, all five templates) and naga on the four changed patterns in both accumulator states, plus a negative control confirming the harness catches the original QKV bug.

@qjia7 Jiajia Qin (qjia7) left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Review frame

  • Exact head reviewed: b5e48ef776313b6a1e96c62a42c0aee460b93d7d.
  • Feature validity: valid. A session-scoped opt-in to f32 accumulation is a reasonable portability fallback while retaining the current f16 default.
  • Risk/scope: deep. This adds a provider option and changes precision, shader generation, dispatch variants, program-cache identity, and public WebGPU configuration.
  • Direction gate: the provider-option direction is appropriate, but the current implementation does not yet provide a consistent or publicly reachable contract.

qjia7 comment status

Item Status at the latest head
Initialize acc_f32_ in QKV and MLP Not resolved. Both members remain uninitialized.
QKV undefined output_element_t Resolved. It now uses q_output_element_t, and Q/K/V share T1.
DP4A bias addition precision Resolved. Bias is converted to acc_element_t and the result is downcast only at the store.
Boolean, MatMulNBits-scoped option Not resolved. The general string option remains.
Constructor defaults Resolved as requested: the remaining defaults now match the provider default (false); the regular DP4A constructor requires the argument.
Block-local MLP/QKV dot-product precision Not resolved. The default path still computes these helpers explicitly in f32.
Coverage of accumulator-dependent paths Not resolved. The only new runtime case covers the generic f32 path.

GitHub still reports all three qjia7 inline threads as unresolved. The SDP8AI request is fixed in code; the MLP and QKV f32 threads remain substantively open.

Confirmed findings

C1: Initialize the QKV and MLP acc_f32_ members

MatMulNBitsQkvDecodeProgram accepts acc_f32, but its initializer list ends at has_skip_output_(has_skip_output) (matmul_nbits_qkv.cc:44). MatMulNBitsMlpDecodeProgram similarly ends at activation_kind_(activation_kind) (matmul_nbits_mlp.cc:75). Both later pass the uninitialized member to WGSL_TEMPLATE_PARAMETER, while their CacheHint uses the correctly initialized local value from the caller.

This is undefined behavior and can compile an f16 shader under an f32 cache key, or the reverse. Please add acc_f32_(acc_f32) to both initializer lists. This remains a merge blocker.

C2: Make the f32 option effective for optimized MatMulNBits dispatches

ApplyMatMulNBits selects subgroup-matrix and wide-tile programs before it reads MatmulAccumulatorPrecisionF32() for the generic program (matmul_nbits.cc:228-326). The wide-tile shader still declares results : array<output_element_t, ...>, and subgroup-matrix configurations can use f16 results. Consequently, common M >= 4 shapes can ignore an explicit "f32" request.

Please either make those paths honor the option, or skip/fall back from a path that cannot honor it when f32 is requested. If the intended contract is deliberately narrower, rename the option to describe that scope and document exactly which paths may ignore it. Simply excluding these paths leaves the current provider-option contract misleading.

S1: Expose and forward the option through the supported WebGPU APIs

The native parser accepts the new config key, but WebGpuExecutionProviderOption in js/common/lib/inference-session.ts has no corresponding property. The browser bridge in js/web/lib/wasm/session-options.ts only forwards explicitly handled fields, so an extra property is ignored; the Node bridge in js/node/src/session_options_helper.cc rejects it as unrecognized. The new C++ test constructs the EP internally and therefore does not cover either public path.

Please add the typed option plus validation/forwarding in the browser and Node bridges, with a public-path test. Otherwise the primary onnxruntime-web users cannot enable the feature.

S2: Preserve the f16 default in the fused block-local helpers

Relative to main, compute_gate_up_sums and compute_projection_sum were changed from output-element computation to explicit f32 returns and vec4<f32> operands. This happens even when acc_f32=false, contradicting the stated unchanged default and potentially changing default numerics and performance.

Roberto's explanation that these sums are block-local and bounded explains why overflow is less concerning there, but it does not justify changing their default precision. Please use acc_element_t consistently for the helper return/operands, so the option selects f32 and the default retains output-type computation.

Roberto's latest reply

The reply is partially valid:

  • QKV alias fix: valid.
  • DP4A bias-order fix: valid in both Qualcomm and non-Qualcomm store paths.
  • SDP8AI return-type fix: valid; it now preserves f32 through the scaled term when requested.
  • Constructor-default fix: valid.
  • Always-f32 MLP/QKV helper explanation: incomplete for compatibility; it does not address the default behavior change.

The reply does not mention or fix the uninitialized QKV/MLP members, public option reachability, optimized-path bypass, or remaining coverage gap.

Test coverage

Float16_LargeK_AccumulatorOverflow deterministically selects the generic MatMulNBits path and requests f32, which is useful numerical coverage. It does not exercise the default/f16 variant, fused MLP, fused QKV, either DP4A path, public option forwarding, or cache separation. Existing MLP/QKV fusion tests and accuracy-level-4 MatMulNBits tests should be parameterized/reused with both option values instead of adding duplicate WebGPU-only test bodies.

For wide-tile/subgroup-matrix, first settle C2: either test that f32 is honored/falls back, or narrow the option's contract. wgsl_gen.py succeeds and git diff --check is clean, but the PR has no compile/test CI beyond CLA, and the author states that the C++ and runtime test have not been built or run.

Verdict

Request changes. The feature and provider-option direction are valid, but C1, C2, S1, and S2 block merge. qjia7's alias, DP4A bias, SDP8AI, and constructor-default points are fixed; initialization, helper precision, option scope, and coverage remain open. The option naming/boolean shape should be resolved while the API is still new. Formatting is cleanup only.

qjia7's review of b5e48ef:

- MatMulNBitsQkvDecodeProgram and MatMulNBitsMlpDecodeProgram took acc_f32 but
  never assigned it, so shader generation read an uninitialized member while the
  cache hint used the caller's value. Both now initialize it.
- The wide-tile program is selected before the generic one and ignored the option
  entirely, so M >= 4 prefill shapes silently kept an f16 accumulator. It now
  takes acc_f32, aliases acc_element_t, adds the bias in accumulator precision and
  downcasts only at the store, with acc_f32 in its cache hint.
- The subgroup-matrix path cannot honour an f32 request: every fp16 entry in
  supported_subgroup_matrix_configs has resultComponentType F16 and the kernels
  declare subgroup_matrix_result<f16, ...>. It now declines itself when f32 is
  requested for an fp16 output instead of ignoring the option.
- compute_gate_up_sums and compute_projection_sum were changed to unconditional
  f32 in an earlier commit, which altered the default path relative to main. They
  follow acc_element_t again, so acc_f32=false is identical to main.
- Expose preferredMatmulAccumulatorPrecision through the JS APIs: typed option in
  js/common, validation and forwarding in js/web, accepted key in js/node.
- Add Float16_AccumulatorPrecisionOption_AllPaths, which runs the generic,
  wide-tile and dp4a dispatches with the option in both states.
@RobertoReale

Copy link
Copy Markdown
Author

Pushed 352c87856.

C1 was undefined behaviour and you're right that the cache hint made it worse: both decode programs took acc_f32 and never assigned it, while CacheHint used the caller's copy, so the compiled variant and the key could disagree. Initialized in both.

C2: the wide-tile program is selected before the generic one, so the shapes it takes were ignoring the option outright. It now carries acc_f32, adds the bias in acc_element_t and downcasts at write_output only. One coordination note: this touches matmul_nbits_wide_tile.wgsl.template, which #31703 is rewriting at the same time. Whichever lands second gets a small conflict, and I'd rather rebase onto Jianhui Dai (@daijh)'s work than make him rebase onto mine.

For subgroup matrix I took the other branch of your suggestion and declined the path. Every fp16 entry in supported_subgroup_matrix_configs has resultComponentType == F16 and the kernels declare subgroup_matrix_result<f16, ...>, so there is no variant to switch to; CanApplySubgroupMatrixMatMulNBits now returns false when f32 is requested for an fp16 output, and the dispatch falls through to a kernel that honours it. fp32 outputs are unaffected.

S2: correct, and I should have fixed it when you first raised it instead of explaining why it was safe. The helpers follow acc_element_t again, so with the default the casts are identities and the shader computes what main computes.

S1: typed option in js/common, validation and forwarding in js/web, key accepted in the node bridge.

On the name, I'd keep preferredMatmulAccumulatorPrecision since it was your proposal on #29611, but if you want the boolean form say so and it's a one-commit rename.

New test Float16_AccumulatorPrecisionOption_AllPaths runs the generic, wide-tile and dp4a dispatches with the option in both states. The fused MLP and QKV kernels are reached through fusion rather than a MatMulNBits node, so they are still uncovered and I haven't found a clean entry point for them.

Still nothing compiled here: wgsl_gen.py across all six templates emitting both alias branches, and naga on the new wide-tile, MLP and QKV shapes in both accumulator states, plus two negative controls where the f32 accumulator reaches the f16 output without the conversion (both correctly rejected). The C++ has never seen a compiler, which is why enabling the pipelines would be worth more than anything else I can check locally.

microsoft#31703 rewrote the MatMulNBits wide-tile compute stage: the A tile is stored
as pairs of vecs and the reduction now has three variants picked from
subgroup_min_size (full-tile subgroupShuffle band, chunked bands, or a direct
read). Re-apply the acc_element_t accumulator on top of that shape.

- results is array<acc_element_t, kTileM>, and all three reduction branches
  widen the dot products; the subgroupShuffle itself stays in the
  workgroup-memory element type, so the broadcast payload does not grow.
- MatMulNBitsWideTileProgram carries both subgroup_min_size and acc_f32, and
  the cache hint carries both, so the variants cannot collide.
- With acc_f32 = false the casts are identities and the generated shader is
  main's, checked by diffing wgsl_gen output against 1bc68c1.
@RobertoReale

Copy link
Copy Markdown
Author

Now that #31703 has landed I've merged main in and taken the conflict on my side as offered: 62836761d, and the PR is mergeable again.

The wide-tile hunk needed redoing rather than re-applying. Jianhui Dai (@daijh)'s reduction has three variants selected from subgroup_min_size, so acc_element_t had to go on all three: the full-tile shuffle band, the chunked bands and the direct read. The subgroupShuffle stays in the workgroup-memory element type and only the dot products widen, so the broadcast payload is the same in either state. MatMulNBitsWideTileProgram carries both subgroup_min_size and acc_f32 now, and so does the cache hint.

With the default the casts are identities, and I checked that instead of assuming it: generating both templates with wgsl_gen.py and diffing against 1bc68c1d2, the only differences are the alias block, the type of results, the three dot sites and the two stores. naga compiles all three branches in both accumulator states and rejects the two controls where the f32 value reaches the f16 output unconverted.

One thing I had wrong: I expected shuffling the already-widened vec4<f32> to be a type error, and it compiles. So keeping the casts outside the shuffle is a bandwidth argument, not a correctness one.

Worth flagging for the cost side: on fp16 tile_m is 32 now rather than 16, so "f32" asks for 32 f32 accumulator registers per lane where the default keeps 32 f16. Still nothing compiled here.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants