[WebGPU] Accumulate MatMulNBits wide-tile in output_element_t - #29611
[WebGPU] Accumulate MatMulNBits wide-tile in output_element_t#29611daijh wants to merge 2 commits into
Conversation
|
@qjia7 @hariharans29 PTAL |
Please check #29599, I think using f32 accumulators is the right direction. Could you split the refactoring changes into a separate PR? That would make it much easier to review and focus on the optimization work. |
daijh
changed the title
|
Let's leave this refactoring for a follow-up PR. |
I agree - isn't it safer to accumulate in fp32 ? Accumulating in fp16 universally across all workloads (if the op is fp16) seems risky ? |
|
Since #29599 was referenced here — a data point on the accumulator part of this change (the tile-layout optimization itself looks like a clear win and is orthogonal): The f32 accumulator in the wide-tile kernel isn't just precision hygiene — it's load-bearing for correctness. With f16 accumulation, partial dot-product sums overflow 65504 and saturate to +Inf once K approaches ~2048 (deterministically, not as a gradual accuracy loss), which then propagates NaN through LayerNorm/Softmax. That's the failure class of #26732. Notably, both benchmark models here are in that range (gpt-oss-20b K≈2880, Phi-4-mini K=3072) — the table reports prefill TPS, but was output correctness checked against the f32-accumulator baseline on those runs? Also, since the PR changes tile layout and accumulator precision together, it would be useful to benchmark them separately — the memory-access refactoring may account for most of the gain, in which case you could keep the speedup without reintroducing the overflow. If f32 accumulation does turn out to be a measurable cost on Intel in isolation, gating the promotion on K above a threshold (overflow risk scales with reduction length) was floated in #29599 as a compromise that keeps small-K workloads untouched — happy to coordinate there. |
I think it would be helpful to reproduce the overflow issue in the LLM first, before we decide on our next steps. Alternatively, we could use the |
Done — details in #29599 (comment) (the non-repro had a concrete cause: the Hub's fp16 Gemma weights were re-exported on 2025-12-02 with Clip nodes inserted as a model-side workaround for this exact overflow). Summary: pinning On
|
|
Data point relevant to the f16-accumulator question here: following up on the repro request, I ran raw-WGSL f16 accumulation probes across adapters and backends — full write-up in #29599 (comment). Short version: on Intel (D3D12) the shader compiler evaluates straight-line/unrolled f16 The implication for this PR: correctness results obtained with f16 accumulators on Intel don't transfer to other vendors (or even to Intel on a different backend, or to looped codegen). Demoting the wide-tile accumulator to |
Since the NVIDIA-specific logic requires fp32 accumulators and will be gated by a vendor string check. |
|
Flagging one interaction for whoever reviews this PR, since the accumulator part is being deferred to #29599 (full reasoning in #29599 (comment)): The proposal on #29599 is to gate the f32 accumulator behind an NVIDIA vendor-string check. If that lands and this PR demotes the wide-tile accumulator to The tile-layout half of this PR looks like a clean, orthogonal win. The only ask is to keep the accumulator in f32 (or make the demotion opt-in via |
Could you provide a link to the test page where you observed the failure with SmolLM2 / Gemma-q4f16? |
|
@daijh For the overall 121% improvement for Phi-4 1024, how much did each of these two optimizations contribute individually? I'd like to understand the impact of using FP32 on Intel devices. |
|
Let me further divide the current PR into multiple smaller ones, with each focusing on a specific optimization. |
|
Save original description Description
Intel Panther Lake
[1] https://huggingface.co/onnx-community/gpt-oss-20b-ONNX Motivation and ContextSee above. |
|
@daijh @qjia7 — you asked for an Intel repro of an f16 accumulator overflow in MatMulNBits. Here it is, at the level that matters for this PR: the wide-tile kernel's own accumulation shape, on Intel Iris Xe (gen-12lp) through Chrome/Dawn/D3D12 — the same vendor and backend you benchmark on. Why an LLM-level repro of this cannot exist todayWorth stating plainly, because I think it explains the loop we've been in: The probeI took the inner loop of Chrome 1xx, Dawn → D3D12, adapter
Same shader shape on wgpu-native/Vulkan: Intel Iris Xe → Inf, NVIDIA RTX 3050 Ti → Inf. So this is not NVIDIA-specific, and it is not Vulkan-specific. Note how ordinary the inputs are. At K = 3072 you only need a mean |a·w| ≥ 21.3 to cross 65504 — an activation of 6 against a dequantized weight of 4 does it. Case B is the more insidious one: the mathematically correct output is 0, well inside f16 range, and the f16 accumulator still returns Inf, which then propagates NaN through LayerNorm/Softmax. That is exactly the #26732 failure signature, and no output-value inspection of the result would predict it. Why this doesn't contradict your non-repro on IntelIt's consistent with the WGSL probes in #29599 (comment): Intel's D3D12 compiler evaluates straight-line / unrolled f16 On the benchmark numbersBranch That leaves the accumulator as a standalone question with, as of now, no isolated number attached. The comment this PR deletes says "Utilizing an f32 accumulator mitigated precision loss with minimal performance impact compared to an f16 accumulator" — if Panther Lake now shows otherwise, those numbers would be very interesting and would genuinely change the discussion. Absent them, the trade is a measurable-but-unquantified speedup against a deterministic Inf at K ≥ ~2048 on the exact hardware in the benchmark table. If f16 accumulation is worth real throughput on Intel, Self-contained repro page (save as .html, open in Chrome; no model download, runs in ~1s)<!doctype html>
<html><head><meta charset="utf-8"><title>MatMulNBits wide-tile accumulator probe</title></head>
<body><pre id="log">running</pre>
<script type="module">
const log = s => { document.getElementById('log').textContent += '\n' + s; };
const K = 3072, NB = K / 32, TM = 4; // K = Phi-4-mini reduction length; block32
// Inner loop of matmul_nbits_wide_tile.wgsl.template, accumulator parameterized.
const shader = (acc) => `
enable f16;
@group(0) @binding(0) var<storage, read> params : array<f16>;
@group(0) @binding(1) var<storage, read_write> outp : array<f32>;
const KAVecSizeForBlock32 = 8u;
const kTileM : u32 = ${TM}u;
var<workgroup> a_data_tile : array<array<vec4<f16>, KAVecSizeForBlock32>, kTileM>;
@compute @workgroup_size(64)
fn main(@builtin(local_invocation_index) local_idx : u32) {
let a_val = params[0]; let w_val = params[1]; let flip = params[2];
var results : array<${acc}, kTileM>;
for (var block_idx = 0u; block_idx < ${NB}u; block_idx++) {
let a_row_idx = local_idx / KAVecSizeForBlock32;
let a_col_idx = local_idx % KAVecSizeForBlock32;
if (a_row_idx < kTileM) { a_data_tile[a_row_idx][a_col_idx] = vec4<f16>(a_val); }
workgroupBarrier();
var s = f16(1.0);
if (flip != f16(0.0) && block_idx >= ${NB}u / 2u) { s = f16(-1.0); }
let w = w_val * s;
let b_dequantized = mat2x4<f16>(w, w, w, w, w, w, w, w);
for (var b_idx = 0u; b_idx < 4u; b_idx++) {
for (var m_idx = 0u; m_idx < kTileM; m_idx++) {
let a_data0 = a_data_tile[m_idx][b_idx * 2u];
let a_data1 = a_data_tile[m_idx][b_idx * 2u + 1u];
results[m_idx] += ${acc}(dot(a_data0, b_dequantized[0])) +
${acc}(dot(a_data1, b_dequantized[1]));
}
}
workgroupBarrier();
}
if (local_idx == 0u) {
for (var m_idx = 0u; m_idx < kTileM; m_idx++) { outp[m_idx] = f32(results[m_idx]); }
}
}`;
async function run(device, acc, params) {
const module = device.createShaderModule({ code: shader(acc) });
const f32tof16 = v => { const f = new Float32Array([v]); const u = new Uint32Array(f.buffer)[0];
const s = (u >> 16) & 0x8000, e = ((u >> 23) & 0xff) - 112, m = (u >> 13) & 0x3ff;
return v === 0 ? 0 : s | (e << 10) | m; };
const p = device.createBuffer({ size: 8, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
device.queue.writeBuffer(p, 0, new Uint16Array([...params.map(f32tof16), 0]));
const o = device.createBuffer({ size: TM * 4, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC });
const rd = device.createBuffer({ size: TM * 4, usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST });
const pipe = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
const bg = device.createBindGroup({ layout: pipe.getBindGroupLayout(0),
entries: [{ binding: 0, resource: { buffer: p } }, { binding: 1, resource: { buffer: o } }] });
const enc = device.createCommandEncoder();
const cp = enc.beginComputePass(); cp.setPipeline(pipe); cp.setBindGroup(0, bg);
cp.dispatchWorkgroups(1); cp.end();
enc.copyBufferToBuffer(o, 0, rd, 0, TM * 4);
device.queue.submit([enc.finish()]);
await rd.mapAsync(GPUMapMode.READ);
const v = new Float32Array(rd.getMappedRange().slice(0))[0]; rd.unmap();
return v;
}
try {
const adapter = await navigator.gpu.requestAdapter();
const i = adapter.info || {};
log(`adapter: vendor=${i.vendor} architecture=${i.architecture} shader-f16=${adapter.features.has('shader-f16')}`);
const device = await adapter.requestDevice({ requiredFeatures: ['shader-f16'] });
for (const [name, params, exact] of [
['A all-positive a=6 w=4', [6, 4, 0], 32 * 6 * 4 * NB],
['B half+ half- a=6 w=8', [6, 8, 1], 0],
]) {
const f32 = await run(device, 'f32', params); // main
const f16 = await run(device, 'f16', params); // this PR (output_element_t = f16)
log(`${name}: exact=${exact} acc_f32=${f32} acc_f16=${f16}`);
}
log('DONE');
} catch (e) { log('ERROR: ' + e.message); }
</script></body></html> |
daijh
changed the title
I can build two onnxruntime-web packages—one with f16-accumulator and another with f32-accumulator—to confirm the issue on GPUs from Intel and other vendors. It is worth noting that MatMulNbits-DP4A in the released onnxruntime-web package also uses an f16 accumulator, which might exhibit the same overflow issue. |
That's exactly the right experiment — thank you. Two A/B packages differing only in the accumulator, run on real models across vendors, settles both the correctness question and the cost question at once. If it helps, my falsifiable prediction for it: the f16 package produces Inf/NaN (garbage tokens, or a hard
Confirmed — and it does. Same probe as above, this time replicating the shipped Chrome / Dawn / D3D12, Intel Iris Xe (gen-12lp), K = 3072, int8 A, dequantized q4 B, scale = 1:
Note these are unremarkable quantized values — an int8 activation of 6 or 12 against a dequantized 4-bit weight of 4. This is already fixed in #29599: DP4A repro page (save as .html, open in Chrome; no model, no download)<!doctype html>
<html><head><meta charset="utf-8"><title>dp4a f16 accumulator probe</title></head>
<body><pre id="log">running</pre>
<script type="module">
const log = s => { document.getElementById('log').textContent += '\n' + s; };
const K = 3072, ITERS = K / 32; // each SDP8AI covers 32 k-values (a1+a2 = 8 u32 = 32 int8)
// Verbatim shape of the shipped dp4a_matmul.wgsl.template accumulation.
const shader = (acc) => `
enable f16;
@group(0) @binding(0) var<storage, read> ab : array<u32>;
@group(0) @binding(1) var<storage, read> sc : array<f16>;
@group(0) @binding(2) var<storage, read_write> outp : array<f32>;
alias output_element_t = f16; // f16 model
alias mul_precision = output_element_t; // n_bits == 4, per dp4a_matmul_common.wgsl.template
fn SDP8AI(a1:vec4<u32>, b1:vec4<u32>, a2:vec4<u32>, b2:vec4<u32>, scale:output_element_t) -> output_element_t {
var local_sum = dot4I8Packed(a1[0], b1[0]);
local_sum += dot4I8Packed(a1[1], b1[1]);
local_sum += dot4I8Packed(a1[2], b1[2]);
local_sum += dot4I8Packed(a1[3], b1[3]);
local_sum += dot4I8Packed(a2[0], b2[0]);
local_sum += dot4I8Packed(a2[1], b2[1]);
local_sum += dot4I8Packed(a2[2], b2[2]);
local_sum += dot4I8Packed(a2[3], b2[3]);
return output_element_t(mul_precision(local_sum) * mul_precision(scale));
}
@compute @workgroup_size(1)
fn main() {
let a = vec4<u32>(ab[0], ab[1], ab[2], ab[3]);
let b = vec4<u32>(ab[4], ab[5], ab[6], ab[7]);
var lane_output1 : vec4<${acc}>;
for (var kidx = 0u; kidx < ${ITERS}u; kidx++) {
var s = sc[0];
if (sc[1] != f16(0.0) && kidx >= ${ITERS}u / 2u) { s = -s; }
lane_output1[0] += ${acc}(SDP8AI(a, b, a, b, s));
}
outp[0] = f32(lane_output1[0]);
}`;
async function run(device, acc, aVal, bVal, scale, flip) {
const module = device.createShaderModule({ code: shader(acc) });
const pack = v => { const b = (v & 0xff) >>> 0; return (b | (b << 8) | (b << 16) | (b << 24)) >>> 0; };
const f32tof16 = v => { const f = new Float32Array([v]); const u = new Uint32Array(f.buffer)[0];
const s = (u >> 16) & 0x8000, e = ((u >> 23) & 0xff) - 112, m = (u >> 13) & 0x3ff;
return v === 0 ? 0 : s | (e << 10) | m; };
const abBuf = device.createBuffer({ size: 32, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
device.queue.writeBuffer(abBuf, 0, new Uint32Array([pack(aVal), pack(aVal), pack(aVal), pack(aVal),
pack(bVal), pack(bVal), pack(bVal), pack(bVal)]));
const scBuf = device.createBuffer({ size: 8, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
device.queue.writeBuffer(scBuf, 0, new Uint16Array([f32tof16(scale), f32tof16(flip ? 1 : 0), 0, 0]));
const o = device.createBuffer({ size: 16, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC });
const rd = device.createBuffer({ size: 16, usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST });
const pipe = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
const bg = device.createBindGroup({ layout: pipe.getBindGroupLayout(0), entries: [
{ binding: 0, resource: { buffer: abBuf } }, { binding: 1, resource: { buffer: scBuf } },
{ binding: 2, resource: { buffer: o } }] });
const enc = device.createCommandEncoder();
const cp = enc.beginComputePass(); cp.setPipeline(pipe); cp.setBindGroup(0, bg);
cp.dispatchWorkgroups(1); cp.end();
enc.copyBufferToBuffer(o, 0, rd, 0, 16);
device.queue.submit([enc.finish()]);
await rd.mapAsync(GPUMapMode.READ);
const v = new Float32Array(rd.getMappedRange().slice(0))[0]; rd.unmap();
return v;
}
try {
const adapter = await navigator.gpu.requestAdapter();
const i = adapter.info || {};
log(`adapter: vendor=${i.vendor} architecture=${i.architecture} shader-f16=${adapter.features.has('shader-f16')}`);
const device = await adapter.requestDevice({ requiredFeatures: ['shader-f16'] });
const partial = 32 * 6 * 4;
for (const [name, a8, flip, exact] of [
['A all-positive (a_i8=6, b_i8=4)', 6, false, partial * ITERS],
['B half+ half- (a_i8=12, b_i8=4)', 12, true, 0],
]) {
const f32 = await run(device, 'f32', a8, 4, 1, flip); // #29599
const f16 = await run(device, 'f16', a8, 4, 1, flip); // shipped
log(`${name}: exact=${exact} lane_output_f32=${f32} lane_output_f16=${f16}`);
}
log('DONE');
} catch (e) { log('ERROR: ' + e.message); }
</script></body></html> |
I insist that we reproduce the reported issue using real LLM inference with the MatMulNBits operation and an FP16 accumulator. |
There is one — a real model, real MatMulNBits, no synthetic values. It's already in this thread (comment); here it is as a runnable recipe:
The reason the current Hub model doesn't repro isn't that the overflow is absent — it's that those fp16 Gemma weights were re-exported on 2025-12-02 with Clip nodes inserted specifically to clamp this overflow model-side. Pinning to On normalization: LayerNorm/RMSNorm bounds the activations going into MatMulNBits, but the overflow is neither in the activations nor in the final output — it's in the partial sum accumulated across K inside the kernel. At K ≈ 2048–3072 a mean |a·w| ≈ 21 is enough for the running sum to cross 65504 even when every input and the true output sit comfortably in f16 range. Case B in the probe is exactly that: normalized-scale inputs, true output 0, f16 accumulator still returns Inf. Q/K/V normalization can't prevent it because it doesn't touch the reduction — and MatMulNBits here is the quantized weight matmul (MLP / projections), not the attention-score matmul, so the normalized Q/K/V tensors aren't even the operands. That said, your A/B onnxruntime-web packages remain the cleanest way to settle it end-to-end, and I'm glad you're building them. The pinned- |
Correction — I was wrong about Gemma, and @daijh was right to insistI went back and checked my own claim against the model instead of arguing for it. It does not hold, and I'm retracting it. What I actually measuredI loaded Peak running partial sum across all 126 nodes: ~324. Not one comes within two orders of magnitude of 65504. Your normalization argument is correct for this model. My "mean |a·w| ≈ 21 is enough" was arithmetically true and practically irrelevant: the real mean |a·w| here is ~0.01–0.1. What the 2025-12-02 Clip nodes actually clampNot MatMulNBits. They are 18 nodes, one per layer, on And when I trace that residual in fp16 on the pre-Clip revision, it grows monotonically and crosses the fp16 ceiling between layers 7 and 8:
That is a model-level fp16 activation-range overflow in the residual stream — the known Gemma-3 fp16 behaviour that the original reporter of #26732 flagged himself in the issue (he linked the unsloth writeup: "activations become infinity for float16"). The Clips fix that. An f32 matmul accumulator does not, and I should not have cited them as evidence for it. Consequences I accept
What still stands — stated narrowly this timeThe one real-model, end-to-end case I have is whisper-small q4 (details): unusable on WebGPU, then word-for-word parity with the WASM EP after promoting only the accumulator, ~2x faster, nothing else changed. That is a genuine f16-accumulator overflow on a real model — and it is the "genuine LLM test case" you asked for. But it exercises the JSEP MatMulNBits path, not the native wide-tile kernel this PR touches, and I'm not going to stretch it into something it doesn't cover. So for the wide-tile accumulator specifically, what I have is the portability argument (WGSL permits extra intermediate precision; Intel/D3D12 promotes unrolled Your A/B packages remain the right way to settle it, and my prediction is now explicitly falsifiable: if the f16 package matches the f32 package on real models across vendors, the wide-tile demotion is safe and I have no objection to it. I'll run both on Intel Iris Xe (D3D12 + Vulkan) and NVIDIA RTX 3050 Ti and report whatever they say, including if it contradicts me. Apologies for the noise on the Gemma line — it cost you several rounds of review time. |
|
@RobertoReale |
|
@daijh Thanks! Just to confirm on next steps - are you currently preparing the two A/B onnxruntime-web packages (f16 vs f32 accumulators) you mentioned in #29611 (comment)? As soon as you share the links or instructions, I will run the multi-vendor suite across both Intel Iris Xe and NVIDIA RTX 3050 Ti and report the exact numbers here. |
@qjia7 @hariharans29 |
|
Please get the package from the CI with this PR, and compare with default package. |
|
Before running the A/B I set your package up locally, and the build linked from run 29236183541 doesn't actually exercise the f16 accumulator yet — it still compiles the f32 wide-tile shader. I dropped its Could you re-trigger CI and confirm the shipped shader has the f16 accumulator? If you can also drop the matching f32 build from the same commit, I'll run both on Iris Xe (D3D12 + Vulkan) and RTX 3050 Ti and post numbers right away. |
|
Please follow the steps below: Build onnxruntime-webonnxruntime-web llm demogit clone https://github.com/daijh/onnxruntime_llm
# Download models into `onnxruntime_llm\models`：
# `onnxruntime_llm\models\HuggingFaceTB\SmolLM2-360M-Instruct`
# `onnxruntime_llm\models\onnx-community\gemma-3-270m-it-ONNX`
# `onnxruntime_llm\models\onnx-community\Phi-4-mini-instruct-ONNX`
# `onnxruntime_llm\models\onnx-community\gpt-oss-20b-ONNX`
cd onnxruntime_llm/web
# Install JavaScript dependencies (onnxruntime-web, @huggingface/transformers)
npm install
# Serve the repository root over HTTP (the script serves the parent dir)
npm run serve
# Open http://localhost:8080/web/llm_runner/, it will scan models under `onnxruntime_llm\models` A/B TestingCopy your onnxruntime-web Restart npm/browser, disable cache to ensure the new Enable the |
|
I ran the A/B, and the f16 accumulator does not saturate on the two models I tested. I skipped the two source builds. The WGSL sits in Two controls, because a patch you haven't verified is worthless. The verbose dump prints 256-token prefill of real text, single forward pass, Chrome/Dawn/D3D12.
On Qwen, So I have no wide-tile counterexample, and on this evidence the demotion is safe. What is left is rounding, not saturation: 12 of 256 argmax positions flip on Qwen. Worth knowing, but it is not the failure I argued for. I'm not quoting timings, one run each proves nothing about throughput. If you want a longer reduction than 4864 I can run Phi-4-mini the same way, and the binary patch above lets you reproduce the A/B without waiting on CI. |
|
@qjia7 @hariharans29 |
@qjia7 is currently OOF. I prefer is she takes a look at this. Please allow some time. Meanwhile, I have kicked off CI. |
|
@daijh @RobertoReale @tianleiwu @hariharans29 — combining a few threads here. I checked PyTorch and the CUDA EP: both default to f32 accumulation for fp16 GEMM and expose an opt-in for f16. I don't think that translates directly to WebGPU though. On modern CUDA hardware (Ampere and later), f16-accumulate and f32-accumulate tensor-core MMAs run at the same rate, so "safe by default" costs little to nothing there. WebGPU has no such hardware guarantee — f32 accumulation is a real ALU/register cost, especially on mobile/integrated GPUs, and @daijh's numbers show a meaningful f16 win on Intel. Also worth noting: every other MatMulNBits variant in the WebGPU EP (default, dp4a, subgroup) already accumulates in On correctness: the repros in the thread are synthetic inputs, not real-model divergence (Gemma / SmolLM2 were walked back). Without a real model showing bad output, I'd rather keep f16 as the default and match the rest of the family. So I'm good with approving this PR (#29611) as-is. For @RobertoReale's #29599, I'd suggest we land it behind a new WebGPU EP provider option — e.g. If we later encounter a real model with a correctness issue, we can simply change the default value of the provider option from |
|
Thanks. |
|
The provider option is a better shape than the On the default, one correction. Gemma and SmolLM2 were walked back, by me, and I'm not reopening either. whisper-small q4 was not: unusable on WebGPU on the released build, then word-for-word identical to the WASM EP once only the accumulator is promoted (methodology). That is a real model producing bad output today, not a synthetic input. It runs through the JSEP MatMulNBits path rather than the wide-tile kernel, so it is no argument against this PR, and I've already said I have no wide-tile counterexample. With So: could the default be |
|
Correcting something in my last comment before anyone acts on it. I went to check how the option would actually reach the kernels: the JSEP shaders read That matters for what I argued. whisper-small q4 runs through JSEP. If the option is EP-only and the JSEP accumulator promotion in #29599 lands unconditionally, then whisper is fixed whatever the default is, and my "an f16 default leaves that model broken" only holds if the option is also meant to gate JSEP. I don't know which you had in mind, and I stated it as though I did. So the question I need answered first is what the option governs. If it should gate the whole PR, JSEP needs a parallel I have the EP-side plumbing written locally, modelled on |
|
@RobertoReale, can you build the latest ort-web based on wasm+ webgpu instead of wasm+jsep to see whether the issue still exists for whisper-small q4? The jsep will be deprecated soon and not maintained activity. See #29716. |
|
The answer is no, and for a worse reason than I expected: whisper-small q4 has no f16 accumulator to overflow. I read the file this time instead of assuming. So the whisper argument for an f32 default is withdrawn, and I ran it on wasm+webgpu anyway, 1.29.0-dev jspi package, on 11 s and 30 s of audio: token-for-token identical to the WASM EP in the same package and to the CPU EP, no NaN, no Inf, on Iris Xe (gen-12lp) and RTX 3050 Ti. One loose end, which I am not calling a bug. The q4f16 export of the same model does diverge on webgpu: a comma becomes a period, then |
|
Isolated, and it corrects what I said about it. There is saturation in that model — just not in an accumulator. whisper-small q4f16 keeps LayerNorm decomposed: Two controls, both through the EP's own So it is not the matmul accumulator, #29599 would not fix it, and it is not an argument for an f32 default. It is an fp16 activation overflow, the #26732 family, and WASM and CPU only look correct because ORT inserts fp32 casts around those ops. Happy to file it separately if that is useful. |
|
@RobertoReale @qjia7 @hariharans29 |
|
Please merge latest main. Then rerun CI again. |
- Store the A workgroup tile as [kTileM][KAVecSizeForBlock32/2][2] pairs of vecs so the two vecs consumed per dequantized weight column are read in one access. - Accumulate directly in output_element_t (e.g., f16 for f16 models) instead of hardcoding to f32.
Accumulate directly in output_element_t (e.g., f16 for f16 models) instead of hardcoding an f32 accumulator in the MatMulNBits wide-tile shader.
|
Copilot checked the CI failure:
Could you please re-run the CI to see if it passes this time? |
|
Re-running won't clear it, and updating dependencies is what caused it. The E2E step fails inside It isn't your change. Web CI Pipeline on #31652 already adds |
4 participants
Description
Accumulate directly in output_element_t (e.g., f16 for f16 models)
instead of hardcoding an f32 accumulator in the MatMulNBits wide-tile
shader.
Intel Panther Lake
[1] https://huggingface.co/onnx-community/gpt-oss-20b-ONNX
[2] https://huggingface.co/onnx-community/Phi-4-mini-instruct-ONNX
Motivation and Context
See above.