[aciddelgado]

Description

Add a vendor-neutral, single-dispatch MatMulGemv specialization for single-row FP16 projections with a long reduction dimension and a small output width.

The specialization is narrowly gated to:

• Effective M = 1 and batch = 1.
• FP16 inputs and output.
• 2048 ≤ K ≤ 8192.
• 16 ≤ N ≤ 64, with N divisible by four.
• No bias or fused activation.

Each workgroup computes four output columns using 128 invocations. Threads partition K, accumulate in FP32, and perform a workgroup-local tree reduction before converting the result to FP16. The kernel uses 2 KiB of shared memory and requires no atomics, subgroup operations, external scratch buffers, or additional dispatches. Existing optimized paths retain their precedence, and cases outside the gate retain their existing dispatch behavior.

Add correctness tests covering:

• Exact Qwen decode, speculative-style, and prefill shapes: M = 1, 8, and 66 with K = 5120 and N = 48.
• Explicit GEMV/non-GEMV route assertions using existing WebGPU program-name logs, without adding runtime instrumentation.
• Immediate M/N/K/type selector boundaries, vector promotion, singleton dimensions, and batch/broadcast semantics.
• Bias and fused-activation exclusions through the existing pointwise Conv caller.
• Odd-K nonzero tails, distinct final-vector values, mixed magnitudes, cancellation, repeated-run determinism, changed inputs, and graph capture/replay.
• Decode and tail cases with robustness disabled in an isolated process while retaining full validation.

References are computed on CPU in float64 from the actual FP16-rounded inputs. Tests do not depend on CUDA/Triton or require bitwise equality with the previous FP16-accumulating implementation.

Validated on NVIDIA A100 / Dawn Vulkan: 11 tests and 43 subtests passed.

Motivation and Context

Profiling Qwen3.8-27B batch-size-1 decode identified 96 recurrent [1,5120] × [5120,48] projections per token as a major bottleneck, accounting for approximately 15.6 ms/token and 31% of measured shader GPU time.

For this shape, the generic MatMul path launches only two workgroups using 8×32 output tiles. Most of the row tile is padding, and each group traverses 160 K tiles. The GEMV specialization eliminates padded-row work, parallelizes K within each workgroup, and launches 12 workgroups in one dispatch.

Development microbenchmarks on A100 / Dawn Vulkan measured:

Weight reuse	Generic MatMul	GEMV	Speedup
Repeated weights	114.50 µs	21.05 µs	5.44×
96 rotating weights	159.34 µs	25.10 µs	6.35×

These are mean GPU shader times over 9,600 measured dispatches per case after warmup, with resident device buffers and isolated benchmark processes. GPU clocks were not locked. FP32 accumulation also reduced the target's maximum absolute error against the float64 reference from 0.0391 to 0.000972 in the rotating-weight case.

The measurement-only profiling bridge and benchmark harness are not included in this PR. This change contains only the production GEMV implementation and correctness tests; it does not introduce split-K, GenAI changes, or model-export changes.

Full-model end-to-end ITL and Apple Metal performance have not been measured. The shader is vendor-neutral, but Metal correctness has not yet been validated on Apple hardware.

[Copilot]

Warning

Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.

Copilot review overview

Review effort: Lite
Findings: 4

Open (4)

• The WGSL reduction loop halves workgroup_size_y each iteration, which requires kWorkgroupSizeY… · New
• The dispatch size is computed from dim_b_outer (N) and kColumnsPerWorkgroup, but the shader’s… · New
• The native logger writes raw bytes to fd=2, but the capture file is opened in text mode with UTF-8… · New
• Output OrtValues are allocated using a.dtype rather than the output’s declared element type. In… · New

What changed in this PR

Adds a WebGPU FP16 GEMV shader specialization for single-row MatMul decode-style projections and introduces focused correctness tests to validate routing, gating boundaries, and numerical behavior.

Changes:

• Add MatMulGemv WebGPU program and gate it for M=1, batch=1, FP16 I/O, 2048≤K≤8192, and 16≤N≤64 (N%4==0) with no bias/activation.
• Add WGSL GEMV template that accumulates in FP32 and performs a workgroup-local reduction.
• Add Python correctness tests asserting GEMV vs non-GEMV routing plus robustness/capture/update edge cases.

File	Description
onnxruntime/​test/​python/​transformers/​test_webgpu_matmul.py	New targeted WebGPU MatMul/GEMV tests that validate routing and numerics across gate boundaries and execution modes.
onnxruntime/​core/​providers/​webgpu/​math/​matmul_gemv.wgsl.template	New WGSL kernel template implementing vec4 FP16 GEMV with FP32 accumulation and workgroup reduction.
onnxruntime/​core/​providers/​webgpu/​math/​matmul.cc	Adds MatMulGemvProgram and gating logic to dispatch the GEMV specialization for qualifying shapes/types.

---

💡 Configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[Baiju Meswani (baijumeswani)]

The GEMV indexing and reduction look sound. I have two follow-up requests about test coverage and performance on other GPUs.