sync 2e7068faf55e

Browse files

Files changed (13) hide show

README.md +90 -0
build/webgpu/bench.json +823 -0
build/webgpu/manifest.json +1480 -0
build/webgpu/matmul-nbits-dp4a-quantize.wgsl.jinja +46 -0
build/webgpu/matmul-nbits-gemv-q4.wgsl.jinja +144 -0
build/webgpu/matmul-nbits-q4-dp4a-prefill.wgsl.jinja +104 -0
build/webgpu/matmul-nbits-q4-prefill-tile4x4.wgsl.jinja +171 -0
build/webgpu/matmul-nbits-q4-prefill-tiled-reg.wgsl.jinja +200 -0
build/webgpu/matmul-nbits-q4-prefill-tiled.wgsl.jinja +130 -0
build/webgpu/matmul-nbits-q4-sgmat.wgsl.jinja +193 -0
build/webgpu/matmul-nbits.wgsl.jinja +69 -0
build/webgpu/metadata.json +25 -0
build/webgpu/test.json +1967 -0

README.md CHANGED Viewed

@@ -1,3 +1,93 @@
 ---
 license: apache-2.0
 ---

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.MatMulNBits
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Matrix multiplication with `B` block-quantized along K and dequantized as `(code - zero_point) * scale`. Each power-of-two `block_size` group has a scale and optional zero point; optional bias is added afterward. Two-, four-, and eight-bit codes are packed low-first, and `A` may have rank 2 or 3. This package supports standard unpacked zero points with the same dtype as `A`. Deprecated `g_idx`, prepacked weights, and bfloat16 tensors are not implemented.
+See the [ONNX Runtime `MatMulNBits` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.MatMulNBits) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `A` | `aT` | `T1` | — | — | Float input matrix, not quantized. Rank 2 has shape `(M, K)` and rank 3 has shape `(batch, sequence, K)`; only the last axis is the reduction axis and the leading axes fold into the row count, so the ordinary activation needs no surrounding Reshape. | required |
+| `B` | `bT` | `uint8` | `3` | — | Bit-packed uint8 weight matrix of shape `(N, k_blocks, blob_size)`, where `k_blocks = ceil(K / block_size)` and `blob_size = block_size * bits / 8`. Codes are packed low-first along K. | required |
+| `scales` | `scalesT` | `T1` | `2` | — | Per-block dequantization scale factors of shape `(N, k_blocks)`, with the same dtype as `A`. | required |
+| `zero_points` | `zeroPointsT` | `T3` | `2` | — | Standard unpacked per-block zero points with shape `(N, k_blocks)` and the same dtype as `A`. Omission uses `2^(bits - 1)`. | optional |
+| `bias` | `biasT` | `T1` | `1` | — | Optional bias vector of shape `[N]` added to the output. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `yT` | `T1` | same as `A` | derived; see description | Result of A multiplied by the dequantized weight matrix, with optional bias, same dtype and rank as A: the leading axes of A with a trailing N. | required |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `accuracy_level` | `0` | Minimum internal accuracy level: 0 (unset), 1 (float32), 2 (float16), 3 (bfloat16), or 4 (int8). |
+| `bits` | `4` | Bit width used to quantize B; this package supports 2, 4, and 8. |
+| `K` | — | Input feature dimension of the weight matrix. |
+| `N` | — | Output feature dimension of the weight matrix. |
+| `block_size` | — | Power-of-two quantization block size along K; it must be at least 16. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T1` | `float32`, `float16` |
+| `T3` | `float32`, `float16` |
+## Device requirements
+Some implementation variants require `subgroup-matrix` and `subgroups`. These are route-specific capabilities, not package-wide requirements; availability also depends on the request shape and dtype.
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`matmul-nbits-dp4a-quantize.wgsl.jinja`](build/webgpu/matmul-nbits-dp4a-quantize.wgsl.jinja)
+- [`matmul-nbits-gemv-q4.wgsl.jinja`](build/webgpu/matmul-nbits-gemv-q4.wgsl.jinja)
+- [`matmul-nbits-q4-dp4a-prefill.wgsl.jinja`](build/webgpu/matmul-nbits-q4-dp4a-prefill.wgsl.jinja)
+- [`matmul-nbits-q4-prefill-tile4x4.wgsl.jinja`](build/webgpu/matmul-nbits-q4-prefill-tile4x4.wgsl.jinja)
+- [`matmul-nbits-q4-prefill-tiled-reg.wgsl.jinja`](build/webgpu/matmul-nbits-q4-prefill-tiled-reg.wgsl.jinja)
+- [`matmul-nbits-q4-prefill-tiled.wgsl.jinja`](build/webgpu/matmul-nbits-q4-prefill-tiled.wgsl.jinja)
+- [`matmul-nbits-q4-sgmat.wgsl.jinja`](build/webgpu/matmul-nbits-q4-sgmat.wgsl.jinja)
+- [`matmul-nbits.wgsl.jinja`](build/webgpu/matmul-nbits.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/com.microsoft.MatMulNBits", { version: 1 });
+const { yT } = await kernel({
+  aT: { data: aTData, shape: [2, 17] },
+  bT: { data: bTData, shape: [2, 2, 8] },
+  scalesT: { data: scalesTData, shape: [2, 2] },
+}, {
+  attrs: { K: 17, N: 2, block_size: 16 },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,823 @@

+{
+  "op": "com.microsoft.MatMulNBits",
+  "tunableSpace": { "workgroupSize": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "qwen3-qproj-decode-m1-q8g32-zero-bias-gemv",
+      "preset": "smoke",
+      "vars": { "M": 1, "K": 2560, "N": 4096, "bits": 8, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [1, 2560], "dtype": "float32", "dist": "normal", "seed": 360, "scale": 0.2 },
+        "bT": { "shape": [4096, 80, 32], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [4096, 80],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 362,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        },
+        "zeroPointsT": {
+          "shape": [4096, 80],
+          "dtype": "float32",
+          "dist": "randint",
+          "seed": 363,
+          "min": 120,
+          "max": 140
+        },
+        "biasT": { "shape": [4096], "dtype": "float32", "dist": "normal", "seed": 364, "scale": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [1, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + numel(shapes.zeroPointsT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.biasT) + numel(shapes.yT)) * 4"
+          }
+        ]
+      },
+      "attrs": { "K": 2560, "N": 4096, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "qwen3-qproj-decode-m1-q8g32-zero-bias-gemv-ncols8",
+      "tunables": { "GEMV_N_COLS": 8 },
+      "preset": "smoke",
+      "vars": { "M": 1, "K": 2560, "N": 4096, "bits": 8, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [1, 2560], "dtype": "float32", "dist": "normal", "seed": 360, "scale": 0.2 },
+        "bT": { "shape": [4096, 80, 32], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [4096, 80],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 362,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        },
+        "zeroPointsT": {
+          "shape": [4096, 80],
+          "dtype": "float32",
+          "dist": "randint",
+          "seed": 363,
+          "min": 120,
+          "max": 140
+        },
+        "biasT": { "shape": [4096], "dtype": "float32", "dist": "normal", "seed": 364, "scale": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [1, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + numel(shapes.zeroPointsT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.biasT) + numel(shapes.yT)) * 4"
+          }
+        ]
+      },
+      "attrs": { "K": 2560, "N": 4096, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "matmul-nbits-q4-64x128x128",
+      "preset": "smoke",
+      "vars": { "M": 64, "K": 128, "N": 128, "bits": 4, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [64, 128], "dtype": "float32", "dist": "normal", "seed": 210, "scale": 0.2 },
+        "bT": { "shape": [128, 4, 16], "dtype": "uint8", "dist": "q4pair", "seed": 211 },
+        "scalesT": {
+          "shape": [128, 4],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 212,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [64, 128], "dtype": "float32" } },
+      "bench": { "primary": true, "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 128, "N": 128, "block_size": 32 }
+    },
+    {
+      "name": "matmul-nbits-q4-zero-bias-64x128x128",
+      "preset": "smoke",
+      "vars": { "M": 64, "K": 128, "N": 128, "bits": 4, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [64, 128], "dtype": "float32", "dist": "normal", "seed": 211, "scale": 0.2 },
+        "bT": { "shape": [128, 4, 16], "dtype": "uint8", "dist": "q4pair", "seed": 212 },
+        "scalesT": {
+          "shape": [128, 4],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 213,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        },
+        "zeroPointsT": { "shape": [128, 4], "dtype": "float32", "dist": "randint", "seed": 214, "min": 6, "max": 10 },
+        "biasT": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 215, "scale": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [64, 128], "dtype": "float32" } },
+      "bench": { "primary": true, "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 128, "N": 128, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "qwen3-mlp-gate-decode-m1-q4g64",
+      "preset": "model",
+      "vars": { "M": 1, "K": 2560, "N": 9728, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [1, 2560], "dtype": "float32", "dist": "normal", "seed": 350, "scale": 0.2 },
+        "bT": { "shape": [9728, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 351 },
+        "scalesT": {
+          "shape": [9728, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 352,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [1, 9728], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.yT)) * 4"
+          }
+        ]
+      },
+      "attrs": { "K": 2560, "N": 9728, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m512-q4g64",
+      "preset": "model",
+      "vars": { "M": 512, "K": 2560, "N": 4096, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [512, 2560], "dtype": "float32", "dist": "normal", "seed": 353, "scale": 0.2 },
+        "bT": { "shape": [4096, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 354 },
+        "scalesT": {
+          "shape": [4096, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 355,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [512, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m512-q4g64-zero-bias",
+      "preset": "model",
+      "vars": { "M": 512, "K": 2560, "N": 4096, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [512, 2560], "dtype": "float32", "dist": "normal", "seed": 359, "scale": 0.2 },
+        "bT": { "shape": [4096, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 360 },
+        "scalesT": {
+          "shape": [4096, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 361,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        },
+        "zeroPointsT": { "shape": [4096, 40], "dtype": "float32", "dist": "randint", "seed": 362, "min": 6, "max": 10 },
+        "biasT": { "shape": [4096], "dtype": "float32", "dist": "normal", "seed": 363, "scale": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [512, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "llama-qproj-decode-m1-q4g32",
+      "preset": "model",
+      "vars": { "M": 1, "K": 4096, "N": 4096, "bits": 4, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [1, 4096], "dtype": "float32", "dist": "normal", "seed": 356, "scale": 0.2 },
+        "bT": { "shape": [4096, 128, 16], "dtype": "uint8", "dist": "q4pair", "seed": 357 },
+        "scalesT": {
+          "shape": [4096, 128],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 358,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [1, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.yT)) * 4"
+          }
+        ]
+      },
+      "attrs": { "K": 4096, "N": 4096, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "qwen3-qproj-decode-m1-q8g32",
+      "preset": "model",
+      "vars": { "M": 1, "K": 2560, "N": 4096, "bits": 8, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [1, 2560], "dtype": "float32", "dist": "normal", "seed": 350, "scale": 0.2 },
+        "bT": { "shape": [4096, 80, 32], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [4096, 80],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 352,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [1, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.yT)) * 4"
+          }
+        ]
+      },
+      "attrs": { "K": 2560, "N": 4096, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m512-q8g32",
+      "preset": "model",
+      "vars": { "M": 512, "K": 2560, "N": 4096, "bits": 8, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [512, 2560], "dtype": "float32", "dist": "normal", "seed": 350, "scale": 0.2 },
+        "bT": { "shape": [4096, 80, 32], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [4096, 80],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 352,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [512, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "qwen3-prefill-m512-q8g32-unaligned-n2052",
+      "preset": "edge",
+      "vars": { "M": 512, "K": 2560, "N": 2052, "bits": 8, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [512, 2560], "dtype": "float32", "dist": "normal", "seed": 350, "scale": 0.2 },
+        "bT": { "shape": [2052, 80, 32], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [2052, 80],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 352,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [512, 2052], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 2052, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m512-q4g32-f16",
+      "preset": "model",
+      "vars": { "M": 512, "K": 2560, "N": 4096, "bits": 4, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [512, 2560], "dtype": "float16", "dist": "normal", "seed": 350, "scale": 0.2 },
+        "bT": { "shape": [4096, 80, 16], "dtype": "uint8", "dist": "q4pair", "seed": 351 },
+        "scalesT": {
+          "shape": [4096, 80],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 352,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [512, 4096], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "qwen3-mlp-gate-decode-m1-q4g64-zero-bias-generic",
+      "preset": "smoke",
+      "vars": { "M": 1, "K": 2560, "N": 9728, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [1, 2560], "dtype": "float32", "dist": "normal", "seed": 410, "scale": 0.2 },
+        "bT": { "shape": [9728, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 411 },
+        "scalesT": {
+          "shape": [9728, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 412,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        },
+        "zeroPointsT": { "shape": [9728, 40], "dtype": "float32", "dist": "randint", "seed": 413, "min": 6, "max": 10 },
+        "biasT": { "shape": [9728], "dtype": "float32", "dist": "normal", "seed": 414, "scale": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [1, 9728], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + numel(shapes.zeroPointsT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.biasT) + numel(shapes.yT)) * 4"
+          }
+        ]
+      },
+      "attrs": { "K": 2560, "N": 9728, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m8-q4g64-zero-bias-floor-generic",
+      "preset": "smoke",
+      "vars": { "M": 8, "K": 2560, "N": 4096, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [8, 2560], "dtype": "float32", "dist": "normal", "seed": 420, "scale": 0.2 },
+        "bT": { "shape": [4096, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 421 },
+        "scalesT": {
+          "shape": [4096, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 422,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        },
+        "zeroPointsT": { "shape": [4096, 40], "dtype": "float32", "dist": "randint", "seed": 423, "min": 6, "max": 10 },
+        "biasT": { "shape": [4096], "dtype": "float32", "dist": "normal", "seed": 424, "scale": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [8, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m8-q4g64-default-zero-tile4x4",
+      "preset": "smoke",
+      "vars": { "M": 8, "K": 2560, "N": 4096, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [8, 2560], "dtype": "float32", "dist": "normal", "seed": 420, "scale": 0.2 },
+        "bT": { "shape": [4096, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 421 },
+        "scalesT": {
+          "shape": [4096, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 422,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [8, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m32-q4g64-default-zero",
+      "preset": "model",
+      "vars": { "M": 32, "K": 2560, "N": 4096, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [32, 2560], "dtype": "float32", "dist": "normal", "seed": 430, "scale": 0.2 },
+        "bT": { "shape": [4096, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 431 },
+        "scalesT": {
+          "shape": [4096, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 432,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [32, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m64-q4g64-default-zero",
+      "preset": "model",
+      "vars": { "M": 64, "K": 2560, "N": 4096, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [64, 2560], "dtype": "float32", "dist": "normal", "seed": 440, "scale": 0.2 },
+        "bT": { "shape": [4096, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 441 },
+        "scalesT": {
+          "shape": [4096, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 442,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [64, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m128-q4g64-default-zero",
+      "preset": "model",
+      "vars": { "M": 128, "K": 2560, "N": 4096, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [128, 2560], "dtype": "float32", "dist": "normal", "seed": 450, "scale": 0.2 },
+        "bT": { "shape": [4096, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 451 },
+        "scalesT": {
+          "shape": [4096, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 452,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [128, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "qwen3-qproj-decode-m1-q4g32-f16-gemv",
+      "preset": "smoke",
+      "vars": { "M": 1, "K": 4096, "N": 4096, "bits": 4, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [1, 4096], "dtype": "float16", "dist": "normal", "seed": 430, "scale": 0.2 },
+        "bT": { "shape": [4096, 128, 16], "dtype": "uint8", "dist": "q4pair", "seed": 431 },
+        "scalesT": {
+          "shape": [4096, 128],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 432,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [1, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.yT)) * 2"
+          }
+        ]
+      },
+      "attrs": { "K": 4096, "N": 4096, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "qwen3-qproj-decode-m1-q4g32-f16-gemv-ncols8",
+      "tunables": { "GEMV_N_COLS": 8 },
+      "preset": "smoke",
+      "vars": { "M": 1, "K": 4096, "N": 4096, "bits": 4, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [1, 4096], "dtype": "float16", "dist": "normal", "seed": 430, "scale": 0.2 },
+        "bT": { "shape": [4096, 128, 16], "dtype": "uint8", "dist": "q4pair", "seed": 431 },
+        "scalesT": {
+          "shape": [4096, 128],
+          "dtype": "float16",
+          "dist": "uniform",
+          "seed": 432,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [1, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.yT)) * 2"
+          }
+        ]
+      },
+      "attrs": { "K": 4096, "N": 4096, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "matmul-nbits-q4-naive-fallback-tailK-m32-k2561-n4096",
+      "preset": "stress",
+      "vars": { "M": 32, "K": 2561, "N": 4096, "bits": 4, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [32, 2561], "dtype": "float32", "dist": "normal", "seed": 610, "scale": 0.2 },
+        "bT": { "shape": [4096, 81, 16], "dtype": "uint8", "dist": "q4pair", "seed": 611 },
+        "scalesT": {
+          "shape": [4096, 81],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 612,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [32, 4096], "dtype": "float32" } },
+      "bench": { "primary": true, "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2561, "N": 4096, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "matmul-nbits-q8-naive-fallback-tailK-m48-k2561-n1021",
+      "preset": "stress",
+      "vars": { "M": 48, "K": 2561, "N": 1021, "bits": 8, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [48, 2561], "dtype": "float32", "dist": "normal", "seed": 620, "scale": 0.2 },
+        "bT": { "shape": [1021, 81, 32], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [1021, 81],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 622,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [48, 1021], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2561, "N": 1021, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "matmul-nbits-q8-oddN-compact-fallback-m8-k64-n17",
+      "preset": "smoke",
+      "vars": { "M": 8, "K": 64, "N": 17, "bits": 8, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [8, 64], "dtype": "float32", "dist": "normal", "seed": 630, "scale": 0.2 },
+        "bT": { "shape": [17, 2, 32], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [17, 2],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 632,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [8, 17], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 64, "N": 17, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "rfdetr-qkvo-m1300-q4g64",
+      "preset": "smoke",
+      "vars": { "M": 1300, "K": 384, "N": 384, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [1300, 384], "dtype": "float32", "dist": "normal", "seed": 353, "scale": 0.2 },
+        "bT": { "shape": [384, 6, 32], "dtype": "uint8", "dist": "q4pair", "seed": 354 },
+        "scalesT": {
+          "shape": [384, 6],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 355,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [1300, 384], "dtype": "float32" } },
+      "tunableSpace": { "REGISTER_TILE_TALL_MIN_N": [256, 1024] },
+      "bench": { "primary": true, "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 384, "N": 384, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "rfdetr-fc1-m1300-q4g64",
+      "preset": "smoke",
+      "vars": { "M": 1300, "K": 384, "N": 1536, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [1300, 384], "dtype": "float32", "dist": "normal", "seed": 353, "scale": 0.2 },
+        "bT": { "shape": [1536, 6, 32], "dtype": "uint8", "dist": "q4pair", "seed": 354 },
+        "scalesT": {
+          "shape": [1536, 6],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 355,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [1300, 1536], "dtype": "float32" } },
+      "bench": { "primary": true, "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 384, "N": 1536, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "rfdetr-fc2-m1300-q4g64",
+      "preset": "smoke",
+      "vars": { "M": 1300, "K": 1536, "N": 384, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [1300, 1536], "dtype": "float32", "dist": "normal", "seed": 353, "scale": 0.2 },
+        "bT": { "shape": [384, 24, 32], "dtype": "uint8", "dist": "q4pair", "seed": 354 },
+        "scalesT": {
+          "shape": [384, 24],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 355,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [1300, 384], "dtype": "float32" } },
+      "tunableSpace": { "REGISTER_TILE_TALL_MIN_N": [256, 1024] },
+      "bench": { "primary": true, "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 1536, "N": 384, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "qwen3-qproj-decode-m1-q8g32-zero-only-gemv",
+      "preset": "smoke",
+      "vars": { "M": 1, "K": 2560, "N": 4096, "bits": 8, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [1, 2560], "dtype": "float32", "dist": "normal", "seed": 360, "scale": 0.2 },
+        "bT": { "shape": [4096, 80, 32], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [4096, 80],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 362,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        },
+        "zeroPointsT": {
+          "shape": [4096, 80],
+          "dtype": "float32",
+          "dist": "randint",
+          "seed": 363,
+          "min": 120,
+          "max": 140
+        }
+      },
+      "outputs": { "yT": { "shape": [1, 4096], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + numel(shapes.zeroPointsT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.yT)) * 4"
+          }
+        ]
+      },
+      "attrs": { "K": 2560, "N": 4096, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m512-q4g64-zero-only",
+      "preset": "model",
+      "vars": { "M": 512, "K": 2560, "N": 4096, "bits": 4, "blockSize": 64 },
+      "inputs": {
+        "aT": { "shape": [512, 2560], "dtype": "float32", "dist": "normal", "seed": 359, "scale": 0.2 },
+        "bT": { "shape": [4096, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 360 },
+        "scalesT": {
+          "shape": [4096, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 361,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        },
+        "zeroPointsT": { "shape": [4096, 40], "dtype": "float32", "dist": "randint", "seed": 362, "min": 6, "max": 10 }
+      },
+      "outputs": { "yT": { "shape": [512, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "qwen3-qproj-prefill-m512-q4g64-al4-dp4a",
+      "preset": "model",
+      "vars": { "M": 512, "K": 2560, "N": 4096, "bits": 4, "blockSize": 64, "accuracyLevel": 4 },
+      "inputs": {
+        "aT": { "shape": [512, 2560], "dtype": "float32", "dist": "normal", "seed": 353, "scale": 0.2 },
+        "bT": { "shape": [4096, 40, 32], "dtype": "uint8", "dist": "q4pair", "seed": 354 },
+        "scalesT": {
+          "shape": [4096, 40],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 355,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [512, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 2560, "N": 4096, "bits": 4, "block_size": 64, "accuracy_level": 4 }
+    },
+    {
+      "name": "mmnb-q4-dp4a-prefill-smoke-m128-k1024-n1024-al4",
+      "preset": "smoke",
+      "vars": { "M": 128, "K": 1024, "N": 1024, "bits": 4, "blockSize": 32, "accuracyLevel": 4 },
+      "inputs": {
+        "aT": { "shape": [128, 1024], "dtype": "float32", "dist": "normal", "seed": 353, "scale": 0.2 },
+        "bT": { "shape": [1024, 32, 16], "dtype": "uint8", "dist": "q4pair", "seed": 354 },
+        "scalesT": {
+          "shape": [1024, 32],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 355,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [128, 1024], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] },
+      "attrs": { "K": 1024, "N": 1024, "bits": 4, "block_size": 32, "accuracy_level": 4 }
+    },
+    {
+      "name": "gemma-lmhead-decode-m1-q4g32-n262144-gemv-fold",
+      "preset": "smoke",
+      "vars": { "M": 1, "K": 128, "N": 262144, "bits": 4, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [1, 128], "dtype": "float32", "dist": "normal", "seed": 470, "scale": 0.2 },
+        "bT": { "shape": [262144, 4, 16], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [262144, 4],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 472,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [1, 262144], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.yT)) * 4"
+          }
+        ]
+      },
+      "attrs": { "K": 128, "N": 262144, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "gemma-lmhead-decode-m1-q4g32-n262144-gemv-fold-ncols8",
+      "tunables": { "GEMV_N_COLS": 8 },
+      "preset": "smoke",
+      "vars": { "M": 1, "K": 128, "N": 262144, "bits": 4, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [1, 128], "dtype": "float32", "dist": "normal", "seed": 470, "scale": 0.2 },
+        "bT": { "shape": [262144, 4, 16], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [262144, 4],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 472,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [1, 262144], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.yT)) * 4"
+          }
+        ]
+      },
+      "attrs": { "K": 128, "N": 262144, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "mmnb-prefill-m4-q4g32-n262144-tile4x4-fold",
+      "preset": "smoke",
+      "vars": { "M": 4, "K": 128, "N": 262144, "bits": 4, "blockSize": 32 },
+      "inputs": {
+        "aT": { "shape": [4, 128], "dtype": "float32", "dist": "normal", "seed": 470, "scale": 0.2 },
+        "bT": { "shape": [262144, 4, 16], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "scalesT": {
+          "shape": [262144, 4],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 472,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [4, 262144], "dtype": "float32" } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "bandwidth",
+            "value": "numel(shapes.bT) * 4 + (numel(shapes.aT) + numel(shapes.scalesT) + numel(shapes.yT)) * 4"
+          }
+        ]
+      },
+      "attrs": { "K": 128, "N": 262144, "bits": 4, "block_size": 32 }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,1480 @@

+{
+  "domain": "com.microsoft",
+  "name": "MatMulNBits",
+  "sinceVersion": 1,
+  "description": "Matrix multiplication with `B` block-quantized along K and dequantized as `(code - zero_point) * scale`. Each power-of-two `block_size` group has a scale and optional zero point; optional bias is added afterward. Two-, four-, and eight-bit codes are packed low-first, and `A` may have rank 2 or 3. This package supports standard unpacked zero points with the same dtype as `A`. Deprecated `g_idx`, prepacked weights, and bfloat16 tensors are not implemented.",
+  "inputs": [
+    {
+      "role": "A",
+      "dtype": "T1",
+      "description": "Float input matrix, not quantized. Rank 2 has shape `(M, K)` and rank 3 has shape `(batch, sequence, K)`; only the last axis is the reduction axis and the leading axes fold into the row count, so the ordinary activation needs no surrounding Reshape."
+    },
+    {
+      "role": "B",
+      "dtype": "uint8",
+      "rank": 3,
+      "description": "Bit-packed uint8 weight matrix of shape `(N, k_blocks, blob_size)`, where `k_blocks = ceil(K / block_size)` and `blob_size = block_size * bits / 8`. Codes are packed low-first along K."
+    },
+    {
+      "role": "scales",
+      "dtype": "T1",
+      "rank": 2,
+      "description": "Per-block dequantization scale factors of shape `(N, k_blocks)`, with the same dtype as `A`."
+    },
+    {
+      "role": "zero_points",
+      "dtype": "T3",
+      "rank": 2,
+      "optional": true,
+      "description": "Standard unpacked per-block zero points with shape `(N, k_blocks)` and the same dtype as `A`. Omission uses `2^(bits - 1)`."
+    },
+    {
+      "role": "bias",
+      "dtype": "T1",
+      "rank": 1,
+      "optional": true,
+      "description": "Optional bias vector of shape `[N]` added to the output."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T1",
+      "rank": "ranks.A",
+      "shape": "shapes.A[:-1] + [attrs.N]",
+      "description": "Result of A multiplied by the dequantized weight matrix, with optional bias, same dtype and rank as A: the leading axes of A with a trailing N."
+    }
+  ],
+  "attributes": { "accuracy_level": 0, "bits": 4 },
+  "attributeDescriptions": {
+    "K": "Input feature dimension of the weight matrix.",
+    "N": "Output feature dimension of the weight matrix.",
+    "accuracy_level": "Minimum internal accuracy level: 0 (unset), 1 (float32), 2 (float16), 3 (bfloat16), or 4 (int8).",
+    "bits": "Bit width used to quantize B; this package supports 2, 4, and 8.",
+    "block_size": "Power-of-two quantization block size along K; it must be at least 16."
+  },
+  "attributeConstraints": {
+    "K": { "required": true },
+    "N": { "required": true },
+    "accuracy_level": { "values": [0, 1, 2, 3, 4] },
+    "bits": { "values": [2, 4, 8] },
+    "block_size": { "required": true }
+  },
+  "typeConstraints": { "T1": ["float32", "float16"], "T3": ["float32", "float16"] },
+  "args": {
+    "aT": { "kind": "tensor", "semantic": "A", "role": "input" },
+    "bT": { "kind": "tensor", "semantic": "B", "role": "input" },
+    "scalesT": { "kind": "tensor", "semantic": "scales", "role": "input" },
+    "zeroPointsT": { "kind": "tensor", "semantic": "zero_points", "role": "input", "required": false },
+    "biasT": { "kind": "tensor", "semantic": "bias", "role": "input", "required": false },
+    "yT": { "kind": "tensor", "semantic": "Y", "role": "output" }
+  },
+  "tunables": {
+    "REGISTER_TILE_TAILK_MIN_N": 256,
+    "WORKGROUP_SIZE": 64,
+    "GEMV_N_COLS": 4,
+    "REGISTER_TILE_MIN_N": 1024,
+    "REGISTER_TILE_TALL_MIN_M": 512,
+    "REGISTER_TILE_TALL_MIN_N": 256,
+    "REGISTER_TILE_MIN_WORKGROUPS": 64,
+    "REGISTER_TILE_LARGE_M": 64,
+    "REGISTER_TILE_BK32_MIN_M": 128,
+    "REGISTER_TILE_MIN_M": 16
+  },
+  "constants": {
+    "B_LEN": "attrs.N * kBlocksExpected * blobSizeExpected",
+    "SCALES_LEN": "attrs.N * kBlocksExpected",
+    "BIAS_LEN": "attrs.N"
+  },
+  "bindingSets": {
+    "dp4aQuantize": [
+      { "name": "a", "arg": "aT", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      { "name": "a_quant", "semantic": "aQuant", "buffer": { "type": "storage" }, "elementType": "u32" },
+      { "name": "a_scales", "semantic": "aScales", "buffer": { "type": "storage" }, "elementType": "f32" }
+    ],
+    "dp4aGemm": [
+      { "name": "a_quant", "semantic": "aQuant", "buffer": { "type": "read-only-storage" }, "elementType": "u32" },
+      { "name": "a_scales", "semantic": "aScales", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      { "name": "b", "arg": "bT", "buffer": { "type": "read-only-storage" }, "elementType": "u32", "length": "$B_LEN" },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32",
+        "length": "$SCALES_LEN"
+      },
+      { "name": "y", "arg": "yT", "buffer": { "type": "storage" }, "elementType": "f32" }
+    ],
+    "genericZeroBias": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "zero_points",
+        "arg": "zeroPointsT",
+        "semantic": "zero_points",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$BIAS_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "aRows" },
+            { "name": "K", "type": "u32", "value": "attrs.K" },
+            { "name": "N", "type": "u32", "value": "attrs.N" },
+            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" },
+            { "name": "kBlocks", "type": "u32", "value": "dim(shapes.B, 1)" },
+            { "name": "blobSize", "type": "u32", "value": "dim(shapes.B, 2)" }
+          ]
+        }
+      }
+    ],
+    "genericZeroOnly": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "zero_points",
+        "arg": "zeroPointsT",
+        "semantic": "zero_points",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$SCALES_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "aRows" },
+            { "name": "K", "type": "u32", "value": "attrs.K" },
+            { "name": "N", "type": "u32", "value": "attrs.N" },
+            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" },
+            { "name": "kBlocks", "type": "u32", "value": "dim(shapes.B, 1)" },
+            { "name": "blobSize", "type": "u32", "value": "dim(shapes.B, 2)" }
+          ]
+        }
+      }
+    ],
+    "genericBiasOnly": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$BIAS_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "aRows" },
+            { "name": "K", "type": "u32", "value": "attrs.K" },
+            { "name": "N", "type": "u32", "value": "attrs.N" },
+            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" },
+            { "name": "kBlocks", "type": "u32", "value": "dim(shapes.B, 1)" },
+            { "name": "blobSize", "type": "u32", "value": "dim(shapes.B, 2)" }
+          ]
+        }
+      }
+    ],
+    "genericDefaultZero": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "rows", "type": "u32", "value": "aRows" },
+            { "name": "K", "type": "u32", "value": "attrs.K" },
+            { "name": "N", "type": "u32", "value": "attrs.N" },
+            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" },
+            { "name": "kBlocks", "type": "u32", "value": "dim(shapes.B, 1)" },
+            { "name": "blobSize", "type": "u32", "value": "dim(shapes.B, 2)" }
+          ]
+        }
+      }
+    ],
+    "gemvZeroBias": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "zero_points",
+        "arg": "zeroPointsT",
+        "semantic": "zero_points",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$BIAS_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "K", "type": "u32", "value": "attrs.K" },
+            { "name": "N", "type": "u32", "value": "attrs.N" },
+            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" },
+            { "name": "kBlocks", "type": "u32", "value": "dim(shapes.B, 1)" },
+            { "name": "blobSize", "type": "u32", "value": "dim(shapes.B, 2)" }
+          ]
+        }
+      }
+    ],
+    "gemvZeroOnly": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "zero_points",
+        "arg": "zeroPointsT",
+        "semantic": "zero_points",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$SCALES_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "K", "type": "u32", "value": "attrs.K" },
+            { "name": "N", "type": "u32", "value": "attrs.N" },
+            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" },
+            { "name": "kBlocks", "type": "u32", "value": "dim(shapes.B, 1)" },
+            { "name": "blobSize", "type": "u32", "value": "dim(shapes.B, 2)" }
+          ]
+        }
+      }
+    ],
+    "gemvBiasOnly": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$BIAS_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "K", "type": "u32", "value": "attrs.K" },
+            { "name": "N", "type": "u32", "value": "attrs.N" },
+            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" },
+            { "name": "kBlocks", "type": "u32", "value": "dim(shapes.B, 1)" },
+            { "name": "blobSize", "type": "u32", "value": "dim(shapes.B, 2)" }
+          ]
+        }
+      }
+    ],
+    "gemvDefaultZero": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "K", "type": "u32", "value": "attrs.K" },
+            { "name": "N", "type": "u32", "value": "attrs.N" },
+            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" },
+            { "name": "kBlocks", "type": "u32", "value": "dim(shapes.B, 1)" },
+            { "name": "blobSize", "type": "u32", "value": "dim(shapes.B, 2)" }
+          ]
+        }
+      }
+    ],
+    "prefillZeroBias": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "zero_points",
+        "arg": "zeroPointsT",
+        "semantic": "zero_points",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$BIAS_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" }
+    ],
+    "prefillZeroOnly": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "zero_points",
+        "arg": "zeroPointsT",
+        "semantic": "zero_points",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$SCALES_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" }
+    ],
+    "prefillBiasOnly": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar",
+        "length": "$BIAS_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" }
+    ],
+    "prefillDefaultZero": [
+      {
+        "name": "a",
+        "arg": "aT",
+        "semantic": "A",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$aScalar"
+      },
+      {
+        "name": "b",
+        "arg": "bT",
+        "semantic": "B",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$bScalar",
+        "length": "$B_LEN"
+      },
+      {
+        "name": "scales",
+        "arg": "scalesT",
+        "semantic": "scales",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$scaleScalar",
+        "length": "$SCALES_LEN"
+      },
+      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outputScalar" }
+    ]
+  },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "narrowSubgroupRange": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize < device.adapterInfo.subgroupMaxSize and device.adapterInfo.subgroupMaxSize <= 16",
+    "packedFeature": "device.wgslLanguageFeatures.has(\"packed_4x8_integer_dot_product\")",
+    "kBlocksExpected": "ceilDiv(attrs.K, attrs.block_size)",
+    "blobSizeExpected": "ceilDiv(attrs.block_size * attrs.bits, 8)",
+    "aRows": "numel(shapes.A) / max(1, attrs.K)",
+    "aRankOk": "(ranks.A == 2 or ranks.A == 3) and ranks.Y == ranks.A and dim(shapes.A, ranks.A - 1) == attrs.K and dim(shapes.Y, ranks.Y - 1) == attrs.N and dim(shapes.Y, 0) == dim(shapes.A, 0) and (ranks.A == 2 or dim(shapes.Y, 1) == dim(shapes.A, 1))",
+    "dispatchN4": "ceilDiv(attrs.N, 4)",
+    "gemvDispatchN": "ceilDiv(attrs.N, tunables.GEMV_N_COLS)",
+    "dispatchN32": "ceilDiv(attrs.N, 32)",
+    "dispatchN64": "ceilDiv(attrs.N, 64)",
+    "dispatchM32": "ceilDiv(aRows, 32)",
+    "dispatchM64": "ceilDiv(aRows, 64)",
+    "sgmatTileRows": "64 if aRows >= 64 else 32",
+    "sgmatWorkgroupSize": "256 if aRows >= 64 else 128",
+    "sgmatRowSubtiles": "4 if aRows >= 64 else 2",
+    "sgmatBLoadsPerRow": "sgmatRowSubtiles",
+    "sgmatBLoadWidth": "8 if aRows >= 64 else 16",
+    "sgmatNumSubgroups": "8 if aRows >= 64 else 4",
+    "sgmatWorkgroupStorageBytes": "24576 if aRows >= 64 else 16384",
+    "sgmatDispatchM": "dispatchM64 if aRows >= 64 else dispatchM32",
+    "tiledRegBK": "32 if aRows >= tunables.REGISTER_TILE_BK32_MIN_M else 16",
+    "aFloatOk": "(tensorDtypes.A == \"float32\" or tensorDtypes.A == \"float16\") and f16Ok(tensorDtypes.A)",
+    "portableWorkgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "bitsSupported": "attrs.bits == 2 or attrs.bits == 4 or attrs.bits == 8",
+    "blockSizeSupported": "attrs.block_size >= 16 and attrs.block_size == pow2ceil(attrs.block_size)",
+    "commonShapeValid": "aRankOk and ranks.B == 3 and ranks.scales == 2 and aFloatOk and blockSizeSupported and tensorDtypes.B == \"uint8\" and tensorDtypes.scales == tensorDtypes.A and tensorDtypes.Y == tensorDtypes.A and attrs.K > 0 and attrs.N > 0 and dim(shapes.B, 0) == attrs.N and dim(shapes.B, 1) == kBlocksExpected and dim(shapes.B, 2) == blobSizeExpected and dim(shapes.scales, 0) == attrs.N and dim(shapes.scales, 1) == dim(shapes.B, 1)",
+    "gemvShapeValid": "commonShapeValid and aRows == 1",
+    "zeroPointsValid": "present.zeroPointsT and ranks.zero_points == 2 and tensorDtypes.zero_points == tensorDtypes.A and dim(shapes.zero_points, 0) == attrs.N and dim(shapes.zero_points, 1) == dim(shapes.B, 1)",
+    "biasValid": "present.biasT and ranks.bias == 1 and tensorDtypes.bias == tensorDtypes.A and dim(shapes.bias, 0) == attrs.N",
+    "defaultEpilogue": "not present.zeroPointsT and not present.biasT",
+    "zeroBiasEpilogue": "zeroPointsValid and biasValid",
+    "zeroOnlyEpilogue": "zeroPointsValid and not present.biasT",
+    "biasOnlyEpilogue": "not present.zeroPointsT and biasValid",
+    "portableWorkgroupFits": "portableWorkgroupSize > 0 and portableWorkgroupSize * 64 <= device.limits.maxComputeWorkgroupStorageSize",
+    "tiledWorkgroupFits": "16 <= device.limits.maxComputeWorkgroupSizeX and 16 <= device.limits.maxComputeWorkgroupSizeY and 256 <= device.limits.maxComputeInvocationsPerWorkgroup and 4096 <= device.limits.maxComputeWorkgroupStorageSize",
+    "tiledRegWorkgroupFits": "tiledWorkgroupFits and 16384 <= device.limits.maxComputeWorkgroupStorageSize",
+    "mediumTiledRegWorkgroupFits": "tiledWorkgroupFits and 6144 <= device.limits.maxComputeWorkgroupStorageSize",
+    "sgmatWorkgroupFits": "sgmatWorkgroupSize <= deviceWorkgroupCap and sgmatWorkgroupStorageBytes <= device.limits.maxComputeWorkgroupStorageSize",
+    "registerTileShape": "aRows >= tunables.REGISTER_TILE_MIN_M and ((attrs.N >= tunables.REGISTER_TILE_MIN_N or (aRows >= tunables.REGISTER_TILE_TALL_MIN_M and attrs.N >= tunables.REGISTER_TILE_TALL_MIN_N) or (attrs.K % attrs.block_size != 0 and attrs.N >= tunables.REGISTER_TILE_TAILK_MIN_N)) and (attrs.K % attrs.block_size != 0 or aRows >= tunables.REGISTER_TILE_LARGE_M or dispatchM64 * dispatchN64 >= tunables.REGISTER_TILE_MIN_WORKGROUPS))",
+    "portableTile4Preferred": "registerTileShape and attrs.K % attrs.block_size == 0 and (not device.features.has(\"subgroups\") or narrowSubgroupRange) and has(device.adapterInfo, \"subgroupMinSize\") and device.adapterInfo.subgroupMinSize * 2 < portableWorkgroupSize",
+    "portableMediumRegisterPreferred": "portableTile4Preferred and aRows >= 128 and attrs.K >= 128",
+    "registerTilePreferred": "registerTileShape and not portableTile4Preferred",
+    "mediumRegisterEligible": "registerTilePreferred or portableMediumRegisterPreferred",
+    "largeTiledRegEligible": "registerTilePreferred and tiledRegWorkgroupFits and dispatchM64 <= device.limits.maxComputeWorkgroupsPerDimension",
+    "mediumTiledRegEligible": "mediumRegisterEligible and mediumTiledRegWorkgroupFits and dispatchM32 <= device.limits.maxComputeWorkgroupsPerDimension",
+    "tiledRegVariantEligible": "largeTiledRegEligible or mediumTiledRegEligible",
+    "tiledRegSelectedBK": "tiledRegBK if largeTiledRegEligible else 16",
+    "tiledRegSelectedTileRows": "64 if largeTiledRegEligible else 32",
+    "tiledRegSelectedThreadRows": "4 if largeTiledRegEligible else 2",
+    "tiledRegSelectedDispatchM": "dispatchM64 if largeTiledRegEligible else dispatchM32"
+  },
+  "variants": [
+    {
+      "id": "q4_dp4a_prefill",
+      "priority": 19,
+      "when": ["packedFeature", "commonShapeValid", "defaultEpilogue", "attrs.bits == 4", "attrs.accuracy_level == 4", "tensorDtypes.A == \"float32\"", "attrs.block_size % 32 == 0", "attrs.K % 128 == 0", "attrs.N % 16 == 0", "aRows >= 32", "ceilDiv(attrs.N, 64) <= device.limits.maxComputeWorkgroupsPerDimension", "ceilDiv(aRows, 64) <= device.limits.maxComputeWorkgroupsPerDimension", "16 <= device.limits.maxComputeWorkgroupSizeX", "16 <= device.limits.maxComputeWorkgroupSizeY", "64 <= device.limits.maxComputeWorkgroupSizeX", "256 <= device.limits.maxComputeInvocationsPerWorkgroup", "4608 <= device.limits.maxComputeWorkgroupStorageSize"],
+      "demoteWhen": ["has(device.adapterInfo, \"architecture\") and device.adapterInfo.architecture == \"maxwell\"", "device.features.has(\"chromium-experimental-subgroup-matrix\")", "device.adapterInfo.vendor == \"apple\""],
+      "constants": {
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "vec4Count": "aRows * attrs.K / 4",
+        "blockCount": "aRows * attrs.K / 128"
+      },
+      "intermediates": [
+        { "id": "aQuant", "dtype": "uint32", "shape": "[aRows * attrs.K / 4]" },
+        { "id": "aScales", "dtype": "float32", "shape": "[aRows * attrs.K / 128]" }
+      ],
+      "passes": [
+        {
+          "id": "quantize",
+          "name": "MatMulNBits.Dp4aQuantizeA",
+          "shader": "matmul-nbits-dp4a-quantize.wgsl.jinja",
+          "bindings": "dp4aQuantize",
+          "dispatch": { "threads": "aRows * attrs.K / 4", "workgroupSize": 64 }
+        },
+        {
+          "id": "main",
+          "name": "MatMulNBits.Dp4aPrefill",
+          "shader": "matmul-nbits-q4-dp4a-prefill.wgsl.jinja",
+          "bindings": "dp4aGemm",
+          "dispatch": { "x": "ceilDiv(attrs.N, 64)", "y": "ceilDiv(aRows, 64)" }
+        }
+      ]
+    },
+    {
+      "id": "gemv_default_zero",
+      "priority": 20,
+      "when": ["gemvShapeValid", "defaultEpilogue", "bitsSupported", "portableWorkgroupFits"],
+      "constants": {
+        "gemvNCols": "tunables.GEMV_N_COLS",
+        "hasZero": false,
+        "hasBias": false,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-gemv-q4.wgsl.jinja",
+          "bindings": "gemvDefaultZero",
+          "dispatch": { "workgroups": "gemvDispatchN" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_sgmat_default_zero",
+      "priority": 18,
+      "requires": {
+        "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
+        "limits": { "maxComputeWorkgroupStorageSize": 16384 },
+        "subgroupMinSize": 32,
+        "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
+      },
+      "when": ["commonShapeValid", "defaultEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "attrs.block_size % 16 == 0", "attrs.K % 32 == 0", "aRows >= 32", "has(device.adapterInfo, \"subgroupMinSize\")", "has(device.adapterInfo, \"subgroupMaxSize\")", "device.adapterInfo.subgroupMinSize == 32", "device.adapterInfo.subgroupMaxSize == 32", "dispatchN64 <= device.limits.maxComputeWorkgroupsPerDimension", "sgmatDispatchM <= device.limits.maxComputeWorkgroupsPerDimension", "portableWorkgroupFits", "sgmatWorkgroupFits"],
+      "constants": {
+        "hasZero": false,
+        "hasBias": false,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\"",
+        "tileRows": "sgmatTileRows",
+        "workgroupSize": "sgmatWorkgroupSize",
+        "rowSubtiles": "sgmatRowSubtiles",
+        "bLoadsPerRow": "sgmatBLoadsPerRow",
+        "bLoadWidth": "sgmatBLoadWidth",
+        "numSubgroups": "sgmatNumSubgroups"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-sgmat.wgsl.jinja",
+          "bindings": "prefillDefaultZero",
+          "dispatch": { "x": "dispatchN64", "y": "sgmatDispatchM" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tiled_reg_default_zero",
+      "priority": 17,
+      "when": ["commonShapeValid", "defaultEpilogue", "bitsSupported", "dispatchN64 <= device.limits.maxComputeWorkgroupsPerDimension", "tiledRegVariantEligible"],
+      "constants": {
+        "hasZero": false,
+        "hasBias": false,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\"",
+        "bk": "tiledRegSelectedBK",
+        "tileRows": "tiledRegSelectedTileRows",
+        "tileCols": 64,
+        "threadRows": "tiledRegSelectedThreadRows",
+        "threadCols": 4,
+        "alignedBlockLoads": "aRows >= 128 and attrs.block_size % tiledRegSelectedBK == 0"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tiled-reg.wgsl.jinja",
+          "bindings": "prefillDefaultZero",
+          "dispatch": { "x": "dispatchN64", "y": "tiledRegSelectedDispatchM" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tiled_default_zero",
+      "priority": 16,
+      "when": ["commonShapeValid", "defaultEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "aRows >= 64", "dispatchN32 <= device.limits.maxComputeWorkgroupsPerDimension", "dispatchM32 <= device.limits.maxComputeWorkgroupsPerDimension", "tiledWorkgroupFits"],
+      "constants": {
+        "hasZero": false,
+        "hasBias": false,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tiled.wgsl.jinja",
+          "bindings": "prefillDefaultZero",
+          "dispatch": { "x": "dispatchN32", "y": "dispatchM32" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tile4x4_default_zero",
+      "priority": 15,
+      "when": ["commonShapeValid", "defaultEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "aRows >= 2", "portableWorkgroupFits"],
+      "constants": {
+        "hasZero": false,
+        "hasBias": false,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tile4x4.wgsl.jinja",
+          "bindings": "genericDefaultZero",
+          "dispatch": {
+            "workgroups": "dispatchN4",
+            "y": "min(ceilDiv(aRows, 4), device.limits.maxComputeWorkgroupsPerDimension)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "default_zero",
+      "priority": 0,
+      "when": ["commonShapeValid", "defaultEpilogue", "bitsSupported", "portableWorkgroupFits"],
+      "constants": {
+        "hasZero": false,
+        "hasBias": false,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits.wgsl.jinja",
+          "bindings": "genericDefaultZero",
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "constants.workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "gemv_zero_bias",
+      "priority": 20,
+      "when": ["gemvShapeValid", "zeroBiasEpilogue", "bitsSupported", "portableWorkgroupFits"],
+      "constants": {
+        "gemvNCols": "tunables.GEMV_N_COLS",
+        "hasZero": true,
+        "hasBias": true,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-gemv-q4.wgsl.jinja",
+          "bindings": "gemvZeroBias",
+          "dispatch": { "workgroups": "gemvDispatchN" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_sgmat_zero_bias",
+      "priority": 18,
+      "requires": {
+        "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
+        "limits": { "maxComputeWorkgroupStorageSize": 16384 },
+        "subgroupMinSize": 32,
+        "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
+      },
+      "when": ["commonShapeValid", "zeroBiasEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "attrs.block_size % 16 == 0", "attrs.K % 32 == 0", "aRows >= 32", "has(device.adapterInfo, \"subgroupMinSize\")", "has(device.adapterInfo, \"subgroupMaxSize\")", "device.adapterInfo.subgroupMinSize == 32", "device.adapterInfo.subgroupMaxSize == 32", "dispatchN64 <= device.limits.maxComputeWorkgroupsPerDimension", "sgmatDispatchM <= device.limits.maxComputeWorkgroupsPerDimension", "portableWorkgroupFits", "sgmatWorkgroupFits"],
+      "constants": {
+        "hasZero": true,
+        "hasBias": true,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\"",
+        "tileRows": "sgmatTileRows",
+        "workgroupSize": "sgmatWorkgroupSize",
+        "rowSubtiles": "sgmatRowSubtiles",
+        "bLoadsPerRow": "sgmatBLoadsPerRow",
+        "bLoadWidth": "sgmatBLoadWidth",
+        "numSubgroups": "sgmatNumSubgroups"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-sgmat.wgsl.jinja",
+          "bindings": "prefillZeroBias",
+          "dispatch": { "x": "dispatchN64", "y": "sgmatDispatchM" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tiled_reg_zero_bias",
+      "priority": 17,
+      "when": ["commonShapeValid", "zeroBiasEpilogue", "bitsSupported", "dispatchN64 <= device.limits.maxComputeWorkgroupsPerDimension", "tiledRegVariantEligible"],
+      "constants": {
+        "hasZero": true,
+        "hasBias": true,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\"",
+        "bk": "tiledRegSelectedBK",
+        "tileRows": "tiledRegSelectedTileRows",
+        "tileCols": 64,
+        "threadRows": "tiledRegSelectedThreadRows",
+        "threadCols": 4,
+        "alignedBlockLoads": "aRows >= 128 and attrs.block_size % tiledRegSelectedBK == 0"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tiled-reg.wgsl.jinja",
+          "bindings": "prefillZeroBias",
+          "dispatch": { "x": "dispatchN64", "y": "tiledRegSelectedDispatchM" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tiled_zero_bias",
+      "priority": 16,
+      "when": ["commonShapeValid", "zeroBiasEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "aRows >= 64", "dispatchN32 <= device.limits.maxComputeWorkgroupsPerDimension", "dispatchM32 <= device.limits.maxComputeWorkgroupsPerDimension", "tiledWorkgroupFits"],
+      "constants": {
+        "hasZero": true,
+        "hasBias": true,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tiled.wgsl.jinja",
+          "bindings": "prefillZeroBias",
+          "dispatch": { "x": "dispatchN32", "y": "dispatchM32" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tile4x4_zero_bias",
+      "priority": 15,
+      "when": ["commonShapeValid", "zeroBiasEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "aRows >= 2", "portableWorkgroupFits"],
+      "constants": {
+        "hasZero": true,
+        "hasBias": true,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tile4x4.wgsl.jinja",
+          "bindings": "genericZeroBias",
+          "dispatch": {
+            "workgroups": "dispatchN4",
+            "y": "min(ceilDiv(aRows, 4), device.limits.maxComputeWorkgroupsPerDimension)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "zero_bias",
+      "priority": 0,
+      "when": ["commonShapeValid", "zeroBiasEpilogue", "bitsSupported", "portableWorkgroupFits"],
+      "constants": {
+        "hasZero": true,
+        "hasBias": true,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits.wgsl.jinja",
+          "bindings": "genericZeroBias",
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "constants.workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "gemv_zero_only",
+      "priority": 20,
+      "when": ["gemvShapeValid", "zeroOnlyEpilogue", "bitsSupported", "portableWorkgroupFits"],
+      "constants": {
+        "gemvNCols": "tunables.GEMV_N_COLS",
+        "hasZero": true,
+        "hasBias": false,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-gemv-q4.wgsl.jinja",
+          "bindings": "gemvZeroOnly",
+          "dispatch": { "workgroups": "gemvDispatchN" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_sgmat_zero_only",
+      "priority": 18,
+      "requires": {
+        "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
+        "limits": { "maxComputeWorkgroupStorageSize": 16384 },
+        "subgroupMinSize": 32,
+        "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
+      },
+      "when": ["commonShapeValid", "zeroOnlyEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "attrs.block_size % 16 == 0", "attrs.K % 32 == 0", "aRows >= 32", "has(device.adapterInfo, \"subgroupMinSize\")", "has(device.adapterInfo, \"subgroupMaxSize\")", "device.adapterInfo.subgroupMinSize == 32", "device.adapterInfo.subgroupMaxSize == 32", "dispatchN64 <= device.limits.maxComputeWorkgroupsPerDimension", "sgmatDispatchM <= device.limits.maxComputeWorkgroupsPerDimension", "portableWorkgroupFits", "sgmatWorkgroupFits"],
+      "constants": {
+        "hasZero": true,
+        "hasBias": false,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\"",
+        "tileRows": "sgmatTileRows",
+        "workgroupSize": "sgmatWorkgroupSize",
+        "rowSubtiles": "sgmatRowSubtiles",
+        "bLoadsPerRow": "sgmatBLoadsPerRow",
+        "bLoadWidth": "sgmatBLoadWidth",
+        "numSubgroups": "sgmatNumSubgroups"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-sgmat.wgsl.jinja",
+          "bindings": "prefillZeroOnly",
+          "dispatch": { "x": "dispatchN64", "y": "sgmatDispatchM" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tiled_reg_zero_only",
+      "priority": 17,
+      "when": ["commonShapeValid", "zeroOnlyEpilogue", "bitsSupported", "dispatchN64 <= device.limits.maxComputeWorkgroupsPerDimension", "tiledRegVariantEligible"],
+      "constants": {
+        "hasZero": true,
+        "hasBias": false,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\"",
+        "bk": "tiledRegSelectedBK",
+        "tileRows": "tiledRegSelectedTileRows",
+        "tileCols": 64,
+        "threadRows": "tiledRegSelectedThreadRows",
+        "threadCols": 4,
+        "alignedBlockLoads": "aRows >= 128 and attrs.block_size % tiledRegSelectedBK == 0"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tiled-reg.wgsl.jinja",
+          "bindings": "prefillZeroOnly",
+          "dispatch": { "x": "dispatchN64", "y": "tiledRegSelectedDispatchM" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tiled_zero_only",
+      "priority": 16,
+      "when": ["commonShapeValid", "zeroOnlyEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "aRows >= 64", "dispatchN32 <= device.limits.maxComputeWorkgroupsPerDimension", "dispatchM32 <= device.limits.maxComputeWorkgroupsPerDimension", "tiledWorkgroupFits"],
+      "constants": {
+        "hasZero": true,
+        "hasBias": false,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tiled.wgsl.jinja",
+          "bindings": "prefillZeroOnly",
+          "dispatch": { "x": "dispatchN32", "y": "dispatchM32" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tile4x4_zero_only",
+      "priority": 15,
+      "when": ["commonShapeValid", "zeroOnlyEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "aRows >= 2", "portableWorkgroupFits"],
+      "constants": {
+        "hasZero": true,
+        "hasBias": false,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tile4x4.wgsl.jinja",
+          "bindings": "genericZeroOnly",
+          "dispatch": {
+            "workgroups": "dispatchN4",
+            "y": "min(ceilDiv(aRows, 4), device.limits.maxComputeWorkgroupsPerDimension)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "zero_only",
+      "priority": 0,
+      "when": ["commonShapeValid", "zeroOnlyEpilogue", "bitsSupported", "portableWorkgroupFits"],
+      "constants": {
+        "hasZero": true,
+        "hasBias": false,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits.wgsl.jinja",
+          "bindings": "genericZeroOnly",
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "constants.workgroupSize" }
+        }
+      ]
+    },
+    {
+      "id": "gemv_bias_only",
+      "priority": 20,
+      "when": ["gemvShapeValid", "biasOnlyEpilogue", "bitsSupported", "portableWorkgroupFits"],
+      "constants": {
+        "gemvNCols": "tunables.GEMV_N_COLS",
+        "hasZero": false,
+        "hasBias": true,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-gemv-q4.wgsl.jinja",
+          "bindings": "gemvBiasOnly",
+          "dispatch": { "workgroups": "gemvDispatchN" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_sgmat_bias_only",
+      "priority": 18,
+      "requires": {
+        "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
+        "limits": { "maxComputeWorkgroupStorageSize": 16384 },
+        "subgroupMinSize": 32,
+        "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
+      },
+      "when": ["commonShapeValid", "biasOnlyEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "attrs.block_size % 16 == 0", "attrs.K % 32 == 0", "aRows >= 32", "has(device.adapterInfo, \"subgroupMinSize\")", "has(device.adapterInfo, \"subgroupMaxSize\")", "device.adapterInfo.subgroupMinSize == 32", "device.adapterInfo.subgroupMaxSize == 32", "dispatchN64 <= device.limits.maxComputeWorkgroupsPerDimension", "sgmatDispatchM <= device.limits.maxComputeWorkgroupsPerDimension", "portableWorkgroupFits", "sgmatWorkgroupFits"],
+      "constants": {
+        "hasZero": false,
+        "hasBias": true,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\"",
+        "tileRows": "sgmatTileRows",
+        "workgroupSize": "sgmatWorkgroupSize",
+        "rowSubtiles": "sgmatRowSubtiles",
+        "bLoadsPerRow": "sgmatBLoadsPerRow",
+        "bLoadWidth": "sgmatBLoadWidth",
+        "numSubgroups": "sgmatNumSubgroups"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-sgmat.wgsl.jinja",
+          "bindings": "prefillBiasOnly",
+          "dispatch": { "x": "dispatchN64", "y": "sgmatDispatchM" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tiled_reg_bias_only",
+      "priority": 17,
+      "when": ["commonShapeValid", "biasOnlyEpilogue", "bitsSupported", "dispatchN64 <= device.limits.maxComputeWorkgroupsPerDimension", "tiledRegVariantEligible"],
+      "constants": {
+        "hasZero": false,
+        "hasBias": true,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\"",
+        "bk": "tiledRegSelectedBK",
+        "tileRows": "tiledRegSelectedTileRows",
+        "tileCols": 64,
+        "threadRows": "tiledRegSelectedThreadRows",
+        "threadCols": 4,
+        "alignedBlockLoads": "aRows >= 128 and attrs.block_size % tiledRegSelectedBK == 0"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tiled-reg.wgsl.jinja",
+          "bindings": "prefillBiasOnly",
+          "dispatch": { "x": "dispatchN64", "y": "tiledRegSelectedDispatchM" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tiled_bias_only",
+      "priority": 16,
+      "when": ["commonShapeValid", "biasOnlyEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "aRows >= 64", "dispatchN32 <= device.limits.maxComputeWorkgroupsPerDimension", "dispatchM32 <= device.limits.maxComputeWorkgroupsPerDimension", "tiledWorkgroupFits"],
+      "constants": {
+        "hasZero": false,
+        "hasBias": true,
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "M": "aRows",
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "kBlocks": "dim(shapes.B, 1)",
+        "blockSize": "attrs.block_size",
+        "blobSize": "dim(shapes.B, 2)",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tiled.wgsl.jinja",
+          "bindings": "prefillBiasOnly",
+          "dispatch": { "x": "dispatchN32", "y": "dispatchM32" }
+        }
+      ]
+    },
+    {
+      "id": "prefill_tile4x4_bias_only",
+      "priority": 15,
+      "when": ["commonShapeValid", "biasOnlyEpilogue", "bitsSupported", "attrs.K % attrs.block_size == 0", "aRows >= 2", "portableWorkgroupFits"],
+      "constants": {
+        "hasZero": false,
+        "hasBias": true,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits-q4-prefill-tile4x4.wgsl.jinja",
+          "bindings": "genericBiasOnly",
+          "dispatch": {
+            "workgroups": "dispatchN4",
+            "y": "min(ceilDiv(aRows, 4), device.limits.maxComputeWorkgroupsPerDimension)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "bias_only",
+      "priority": 0,
+      "when": ["commonShapeValid", "biasOnlyEpilogue", "bitsSupported", "portableWorkgroupFits"],
+      "constants": {
+        "hasZero": false,
+        "hasBias": true,
+        "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bScalar": "\"u32\"",
+        "scaleScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "outputScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
+        "bits": "attrs.bits",
+        "defaultZero": "\"2.0\" if attrs.bits == 2 else (\"8.0\" if attrs.bits == 4 else \"128.0\")",
+        "usesF16": "tensorDtypes.A == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-nbits.wgsl.jinja",
+          "bindings": "genericBiasOnly",
+          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "constants.workgroupSize" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/matmul-nbits-dp4a-quantize.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,46 @@

+{{ env.wgsl.resourceDeclarations }}
+// Quantize A rows to signed int8 for the packed-integer-dot kernel: each
+// 128-element block shares one scale (max_abs / 127) and stores 4 values per
+// u32 via pack4x8snorm. All-zero blocks store zeros with a zero scale. K is a
+// multiple of 128, so blocks never straddle rows and the flat layout is exact.
+const VEC4_COUNT: u32 = {{ vec4Count }}u;
+const BLOCK_COUNT: u32 = {{ blockCount }}u;
+const WG: u32 = 64u;
+var<workgroup> maxAbs: array<f32, WG>;
+@compute @workgroup_size(WG)
+fn main(
+  @builtin(workgroup_id) wid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  // 2D-folded dispatch: wid.y carries the high bits past the per-dimension limit.
+  let wgFlat = wid.x + wid.y * nwg.x;
+  let i = wgFlat * WG + lid.x;
+  var v = vec4<f32>(0.0);
+  if (i < VEC4_COUNT) {
+    v = vec4<f32>(f32(a[i * 4u]), f32(a[i * 4u + 1u]), f32(a[i * 4u + 2u]), f32(a[i * 4u + 3u]));
+  }
+  let av = abs(v);
+  maxAbs[lid.x] = max(max(av.x, av.y), max(av.z, av.w));
+  workgroupBarrier();
+  // Each 32-lane half of the workgroup covers one 128-element quant block.
+  let halfBase = (lid.x / 32u) * 32u;
+  var scale = 0.0;
+  for (var l = 0u; l < 32u; l = l + 1u) { scale = max(scale, maxAbs[halfBase + l]); }
+  if (i < VEC4_COUNT) {
+    if (scale > 0.0) {
+      a_quant[i] = pack4x8snorm(v / scale);
+    } else {
+      a_quant[i] = 0u;
+    }
+  }
+  let blockIdx = wgFlat * 2u + lid.x / 32u;
+  if (lid.x % 32u == 0u && blockIdx < BLOCK_COUNT) {
+    a_scales[blockIdx] = scale / 127.0;
+  }
+}

build/webgpu/matmul-nbits-gemv-q4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,144 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Each workgroup handles gemvNCols columns in groups of four, with one vec4
+// accumulator and partial array per group. Every column after the first is
+// bounds-guarded, so arbitrary N and group counts compose.
+{% set nCols = gemvNCols %}
+{% set groups = 2 if nCols == 8 else 1 %}
+{% set comps = ["x", "y", "z", "w"] %}
+const WG: u32 = {{ workgroupSize }}u;
+const N_COLS: u32 = {{ nCols }}u;
+{% for g in range(groups) %}
+{% set sfx = "" if groups == 1 else g %}
+var<workgroup> partials{{ sfx }}: array<vec4<f32>, {{ workgroupSize }}>;
+{% endfor %}
+// {{ bits }}-bit unpack: q2 packs four codes per byte, q4 two nibbles, q8 one byte/weight.
+{% macro matmul_nbits_packed_code(fn="packed_weight", buffer="b", kBlocks="params.kBlocks", blobSize="params.blobSize", bits=4) %}
+fn {{ fn }}(n: u32, block: u32, offset: u32) -> u32 {
+{% if bits == 2 %}
+  let byte_index = offset / 4u;
+  let shift = (offset % 4u) * 2u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 3u;
+{% elif bits == 4 %}
+  let byte_index = offset / 2u;
+  let shift = (offset % 2u) * 4u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 15u;
+{% else %}
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + offset;
+  return {{ buffer }}[packed_index] & 255u;
+{% endif %}
+}
+{%- endmacro %}
+{{- matmul_nbits_packed_code(bits=bits) }}
+// Per-(column, block) zero point. Symmetric weights use the baked default
+// ({{ defaultZero }}); asymmetric weights carry one widened u32 per (col, block)
+// in row-major [N, kBlocks] layout.
+fn zero_point({% if hasZero %}n: u32, block: u32{% endif %}) -> f32 {
+{% if hasZero %}
+  return f32(zero_points[n * params.kBlocks + block]);
+{% else %}
+  return {{ defaultZero }};
+{% endif %}
+}
+@compute @workgroup_size({{ workgroupSize }}, 1, 1)
+fn main(
+  @builtin(workgroup_id) wid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let tid = lid.x;
+  // 2D-folded dispatch: wid.y carries the high bits past the per-dimension
+  // limit (vocab-sized N). The fold over-dispatches, so the tail workgroups
+  // return before any barrier.
+  let col_base = (wid.x + wid.y * nwg.x) * N_COLS;
+  if (col_base >= params.N) {
+    return;
+  }
+{% for g in range(groups) %}
+{% set sfx = "" if groups == 1 else g %}
+  var acc{{ sfx }} = vec4<f32>(0.0);
+{% endfor %}
+  for (var k = tid; k < params.K; k = k + WG) {
+    let block = k / params.blockSize;
+    let offset = k % params.blockSize;
+    let av = f32(a[k]);
+{% for g in range(groups) %}
+{% set sfx = "" if groups == 1 else g %}
+{% for c in range(4) %}
+{% set i = g * 4 + c %}
+{% if i == 0 %}
+    let scale0 = f32(scales[(col_base + 0u) * params.kBlocks + block]);
+    let q0 = f32(packed_weight(col_base + 0u, block, offset)) - zero_point({% if hasZero %}col_base + 0u, block{% endif %});
+    acc{{ sfx }}.x = acc{{ sfx }}.x + av * q0 * scale0;
+{% else %}
+    if (col_base + {{ i }}u < params.N) {
+      let scale{{ i }} = f32(scales[(col_base + {{ i }}u) * params.kBlocks + block]);
+      let q{{ i }} = f32(packed_weight(col_base + {{ i }}u, block, offset)) - zero_point({% if hasZero %}col_base + {{ i }}u, block{% endif %});
+      acc{{ sfx }}.{{ comps[c] }} = acc{{ sfx }}.{{ comps[c] }} + av * q{{ i }} * scale{{ i }};
+    }
+{% endif %}
+{% endfor %}
+{% endfor %}
+  }
+{% for g in range(groups) %}
+{% set sfx = "" if groups == 1 else g %}
+  partials{{ sfx }}[tid] = acc{{ sfx }};
+{% endfor %}
+  workgroupBarrier();
+  var stride = WG / 2u;
+  loop {
+    if (stride == 0u) {
+      break;
+    }
+    if (tid < stride) {
+{% for g in range(groups) %}
+{% set sfx = "" if groups == 1 else g %}
+      partials{{ sfx }}[tid] = partials{{ sfx }}[tid] + partials{{ sfx }}[tid + stride];
+{% endfor %}
+    }
+    stride = stride / 2u;
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+{% for g in range(groups) %}
+{% set sfx = "" if groups == 1 else g %}
+    let total{{ sfx }} = partials{{ sfx }}[0];
+{% endfor %}
+{% for g in range(groups) %}
+{% set sfx = "" if groups == 1 else g %}
+{% for c in range(4) %}
+{% set i = g * 4 + c %}
+{% if hasBias %}
+{% if i == 0 %}
+    y[col_base + 0u] = {{ outputScalar }}(total{{ sfx }}.x + f32(bias[col_base + 0u]));
+{% else %}
+    if (col_base + {{ i }}u < params.N) {
+      y[col_base + {{ i }}u] = {{ outputScalar }}(total{{ sfx }}.{{ comps[c] }} + f32(bias[col_base + {{ i }}u]));
+    }
+{% endif %}
+{% else %}
+{% if i == 0 %}
+    y[col_base + 0u] = {{ outputScalar }}(total{{ sfx }}.x);
+{% else %}
+    if (col_base + {{ i }}u < params.N) {
+      y[col_base + {{ i }}u] = {{ outputScalar }}(total{{ sfx }}.{{ comps[c] }});
+    }
+{% endif %}
+{% endif %}
+{% endfor %}
+{% endfor %}
+  }
+}

build/webgpu/matmul-nbits-q4-dp4a-prefill.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,104 @@

+{{ env.wgsl.resourceDeclarations }}
+// Centralize the signed packed-int8 dot expression so unrolled callers cannot
+// collide in lowering-generated temporaries.
+fn dot4_packed(a_word: u32, b_word: u32) -> i32 {
+  return dot4I8Packed(a_word, b_word);
+}
+// com.microsoft.MatMulNBits q4 prefill with int8-quantized activations
+// (accuracy_level 4). A arrives pre-quantized as packed int8 words with one
+// scale per 128-element block; each weight nibble is rebiased by -8 and packed
+// so the whole inner product runs on packed-integer dots. One 64x64 output
+// tile per workgroup, 4x4 per thread, K stepped 32 elements (8 words) at a
+// time — a step never straddles a weight block (blockSize % 32 == 0) or an
+// activation block (128 % 32 == 0), so both scales are per-step constants.
+const M: u32 = {{ M }}u;
+const K: u32 = {{ K }}u;
+const N: u32 = {{ N }}u;
+const KBLOCKS: u32 = {{ kBlocks }}u;
+const BLOCK_SIZE: u32 = {{ blockSize }}u;
+const BLOB_SIZE: u32 = {{ blobSize }}u;
+const K4: u32 = K / 4u;
+const KSCALE: u32 = K / 128u;
+var<workgroup> tA: array<array<u32, 8u>, 64u>;
+var<workgroup> tB: array<array<u32, 8u>, 64u>;
+var<workgroup> tAscale: array<f32, 64u>;
+var<workgroup> tBscale: array<f32, 64u>;
+@compute @workgroup_size(16, 16, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let mBase = wg.y * 64u;
+  let nBase = wg.x * 64u;
+  let li = lid.y * 16u + lid.x;
+  var acc: array<f32, 16u>;
+  for (var t = 0u; t < 16u; t = t + 1u) { acc[t] = 0.0; }
+  let numTiles = K / 32u;
+  for (var kt = 0u; kt < numTiles; kt = kt + 1u) {
+    let kBase = kt * 32u;
+    let block = kBase / BLOCK_SIZE;
+    // Stage 64 rows x 8 packed A words and 64 cols x 8 packed B words; each of
+    // the 256 threads loads two of each.
+    for (var e = 0u; e < 2u; e = e + 1u) {
+      let idx = li + e * 256u;
+      let r = idx / 8u;
+      let w = idx % 8u;
+      let am = mBase + r;
+      tA[r][w] = select(0u, a_quant[am * K4 + kt * 8u + w], am < M);
+      let bn = nBase + r;
+      if (bn < N) {
+        // Two blob bytes hold this word's four sequential weights
+        // (low nibble = even k). Rebias by the implicit zero point 8.
+        let byteBase = (bn * KBLOCKS + block) * BLOB_SIZE + ((kBase % BLOCK_SIZE) + w * 4u) / 2u;
+        let b0 = b[byteBase];
+        let b1 = b[byteBase + 1u];
+        let quads = vec4<i32>(i32(b0 & 15u), i32((b0 >> 4u) & 15u), i32(b1 & 15u), i32((b1 >> 4u) & 15u));
+        tB[r][w] = pack4xI8(quads - vec4<i32>(8));
+      } else {
+        tB[r][w] = 0u;
+      }
+    }
+    if (li < 64u) {
+      let am = mBase + li;
+      tAscale[li] = select(0.0, a_scales[am * KSCALE + kBase / 128u], am < M);
+      let bn = nBase + li;
+      tBscale[li] = select(0.0, scales[bn * KBLOCKS + block], bn < N);
+    }
+    workgroupBarrier();
+    let aRow = lid.y * 4u;
+    let bCol = lid.x * 4u;
+    for (var i = 0u; i < 4u; i = i + 1u) {
+      let ar = aRow + i;
+      var aw: array<u32, 8u>;
+      for (var w = 0u; w < 8u; w = w + 1u) { aw[w] = tA[ar][w]; }
+      let sa = tAscale[ar];
+      for (var j = 0u; j < 4u; j = j + 1u) {
+        let bc = bCol + j;
+        var d = 0i;
+        for (var w = 0u; w < 8u; w = w + 1u) { d = d + dot4_packed(aw[w], tB[bc][w]); }
+        acc[i * 4u + j] = acc[i * 4u + j] + f32(d) * sa * tBscale[bc];
+      }
+    }
+    workgroupBarrier();
+  }
+  let m0 = mBase + lid.y * 4u;
+  let n0 = nBase + lid.x * 4u;
+  for (var i = 0u; i < 4u; i = i + 1u) {
+    let m = m0 + i;
+    if (m >= M) { continue; }
+    for (var j = 0u; j < 4u; j = j + 1u) {
+      let n = n0 + j;
+      if (n < N) {
+        y[m * N + n] = acc[i * 4u + j];
+      }
+    }
+  }
+}

build/webgpu/matmul-nbits-q4-prefill-tile4x4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,171 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+const TILE_M: u32 = 4u;
+const TILE_N: u32 = 4u;
+var<workgroup> partial0: array<vec4<f32>, {{ workgroupSize }}>;
+var<workgroup> partial1: array<vec4<f32>, {{ workgroupSize }}>;
+var<workgroup> partial2: array<vec4<f32>, {{ workgroupSize }}>;
+var<workgroup> partial3: array<vec4<f32>, {{ workgroupSize }}>;
+// {{ bits }}-bit unpack: q2 stores four codes per byte, q4 two, and q8 one.
+{% macro matmul_nbits_packed_code(fn="packed_weight", buffer="b", kBlocks="params.kBlocks", blobSize="params.blobSize", bits=4) %}
+fn {{ fn }}(n: u32, block: u32, offset: u32) -> u32 {
+{% if bits == 2 %}
+  let byte_index = offset / 4u;
+  let shift = (offset % 4u) * 2u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 3u;
+{% elif bits == 4 %}
+  let byte_index = offset / 2u;
+  let shift = (offset % 2u) * 4u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 15u;
+{% else %}
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + offset;
+  return {{ buffer }}[packed_index] & 255u;
+{% endif %}
+}
+{%- endmacro %}
+{{- matmul_nbits_packed_code(bits=bits) }}
+fn store_row(row: u32, col_base: u32, value: vec4<f32>) {
+  if (row >= params.rows) {
+    return;
+  }
+  let out_base = row * params.N + col_base;
+  // N need not be a multiple of TILE_N: the final column tile is partial, so
+  // guard lanes 1..3 against params.N (col_base itself is always < N). Without
+  // this an odd-N tail would spill into the next row. Lane 0 is unconditional.
+  y[out_base + 0u] = {{ outputScalar }}(value.x);
+  if (col_base + 1u < params.N) { y[out_base + 1u] = {{ outputScalar }}(value.y); }
+  if (col_base + 2u < params.N) { y[out_base + 2u] = {{ outputScalar }}(value.z); }
+  if (col_base + 3u < params.N) { y[out_base + 3u] = {{ outputScalar }}(value.w); }
+}
+@compute @workgroup_size({{ workgroupSize }}, 1, 1)
+fn main(
+  @builtin(workgroup_id) wid: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  let tid = lid.x;
+  // Column tiles fold into z past the per-dimension limit (vocab-sized N,
+  // where y already carries the row-tile grid-stride). Over-dispatched tail
+  // workgroups return before any barrier.
+  let col_base = (wid.x + wid.z * nwg.x) * TILE_N;
+  if (col_base >= params.N) {
+    return;
+  }
+  // Grid-stride over row tiles because dispatch y is device-clamped. This loop
+  // reduces to one iteration when the dispatch covers every tile directly.
+  for (var row_tile = wid.y; row_tile * TILE_M < params.rows; row_tile = row_tile + nwg.y) {
+  let row_base = row_tile * TILE_M;
+  var acc0 = vec4<f32>(0.0);
+  var acc1 = vec4<f32>(0.0);
+  var acc2 = vec4<f32>(0.0);
+  var acc3 = vec4<f32>(0.0);
+  for (var block: u32 = 0u; block < params.kBlocks; block = block + 1u) {
+    let scale = vec4<f32>(
+      f32(scales[(col_base + 0u) * params.kBlocks + block]),
+      f32(scales[(col_base + 1u) * params.kBlocks + block]),
+      f32(scales[(col_base + 2u) * params.kBlocks + block]),
+      f32(scales[(col_base + 3u) * params.kBlocks + block])
+    );
+    let block_k = block * params.blockSize;
+    for (var offset: u32 = tid; offset < params.blockSize; offset = offset + WG) {
+      let k = block_k + offset;
+{% if hasZero %}
+      let zero = vec4<f32>(
+        f32(zero_points[(col_base + 0u) * params.kBlocks + block]),
+        f32(zero_points[(col_base + 1u) * params.kBlocks + block]),
+        f32(zero_points[(col_base + 2u) * params.kBlocks + block]),
+        f32(zero_points[(col_base + 3u) * params.kBlocks + block])
+      );
+{% else %}
+      let zero = vec4<f32>({{ defaultZero }});
+{% endif %}
+      let w = (vec4<f32>(
+        f32(packed_weight(col_base + 0u, block, offset)),
+        f32(packed_weight(col_base + 1u, block, offset)),
+        f32(packed_weight(col_base + 2u, block, offset)),
+        f32(packed_weight(col_base + 3u, block, offset))
+      ) - zero) * scale;
+      var a0 = 0.0;
+      var a1 = 0.0;
+      var a2 = 0.0;
+      var a3 = 0.0;
+      if (row_base + 0u < params.rows) {
+        a0 = f32(a[(row_base + 0u) * params.K + k]);
+      }
+      if (row_base + 1u < params.rows) {
+        a1 = f32(a[(row_base + 1u) * params.K + k]);
+      }
+      if (row_base + 2u < params.rows) {
+        a2 = f32(a[(row_base + 2u) * params.K + k]);
+      }
+      if (row_base + 3u < params.rows) {
+        a3 = f32(a[(row_base + 3u) * params.K + k]);
+      }
+      acc0 = acc0 + a0 * w;
+      acc1 = acc1 + a1 * w;
+      acc2 = acc2 + a2 * w;
+      acc3 = acc3 + a3 * w;
+    }
+  }
+  // Sync before reusing the shared partial buffers: on a subsequent grid-stride
+  // iteration the previous iteration's tid==0 store may still be reading them.
+  workgroupBarrier();
+  partial0[tid] = acc0;
+  partial1[tid] = acc1;
+  partial2[tid] = acc2;
+  partial3[tid] = acc3;
+  workgroupBarrier();
+  var stride = WG / 2u;
+  loop {
+    if (stride == 0u) {
+      break;
+    }
+    if (tid < stride) {
+      partial0[tid] = partial0[tid] + partial0[tid + stride];
+      partial1[tid] = partial1[tid] + partial1[tid + stride];
+      partial2[tid] = partial2[tid] + partial2[tid + stride];
+      partial3[tid] = partial3[tid] + partial3[tid + stride];
+    }
+    stride = stride / 2u;
+    workgroupBarrier();
+  }
+  if (tid == 0u) {
+{% if hasBias %}
+    let biasv = vec4<f32>(
+      f32(bias[col_base + 0u]),
+      f32(bias[col_base + 1u]),
+      f32(bias[col_base + 2u]),
+      f32(bias[col_base + 3u])
+    );
+    store_row(row_base + 0u, col_base, partial0[0] + biasv);
+    store_row(row_base + 1u, col_base, partial1[0] + biasv);
+    store_row(row_base + 2u, col_base, partial2[0] + biasv);
+    store_row(row_base + 3u, col_base, partial3[0] + biasv);
+{% else %}
+    store_row(row_base + 0u, col_base, partial0[0]);
+    store_row(row_base + 1u, col_base, partial1[0]);
+    store_row(row_base + 2u, col_base, partial2[0]);
+    store_row(row_base + 3u, col_base, partial3[0]);
+{% endif %}
+  }
+  }
+}

build/webgpu/matmul-nbits-q4-prefill-tiled-reg.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,200 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// com.microsoft.MatMulNBits q4/q8 prefill, no-subgroup-matrix tier — register-blocked.
+//   Y[M,N] = A[M,K] @ W[K,N] (+ bias),  W[k,n] = (q(n,block,offset) - zero) * scale
+// Dequantizes into a shared [BN][BK/4] tile with device-selected register
+// blocking. Both tiles are indexed by their own output axis and group four K
+// values per vector word, so the micro-tile accumulates through dot() and the
+// column-owning loader writes whole words instead of a BN-strided column.
+// The full geometry computes a 4x4 micro-tile over a 64x64 output tile. The
+// portable geometry computes 2x4 over 32x64 to bound the per-lane accumulator
+// set. K_TILE specializes the K tile. For standard 32/64-element quant blocks,
+// one lane owns one output column and the full BK slice: scale and zero are
+// loaded once, and each stored byte is read once for the K-adjacent codes it
+// carries.
+// Fully bounds-checked (any M/N; K % blockSize == 0).
+const M: u32 = {{ M }}u;
+const K: u32 = {{ K }}u;
+const N: u32 = {{ N }}u;
+const KBLOCKS: u32 = {{ kBlocks }}u;
+const BLOCK_SIZE: u32 = {{ blockSize }}u;
+const BLOB_SIZE: u32 = {{ blobSize }}u;
+const BK: u32 = {{ bk }}u;
+const BM: u32 = {{ tileRows }}u;
+const BN: u32 = {{ tileCols }}u;
+const TM: u32 = {{ threadRows }}u;
+const TN: u32 = {{ threadCols }}u;
+const WG_X: u32 = 16u;
+const WG_Y: u32 = 16u;
+const WG_THREADS: u32 = WG_X * WG_Y;
+// When A uses f16 storage, staging tileA at f16 is bit-identical to widening it
+// here. tileB stays f32 because the dequantized (code - zero)*scale product
+// needs more mantissa than f16 provides.
+{% set aTile = "f16" if usesF16 else "f32" %}
+{% set aRead = "vec4<f32>(" if usesF16 else "" %}
+{% set aReadEnd = ")" if usesF16 else "" %}
+const K_VECS: u32 = BK / 4u;
+var<workgroup> tileA: array<array<vec4<{{ aTile }}>, K_VECS>, BM>;   // [BM][BK/4]
+var<workgroup> tileB: array<array<vec4<f32>, K_VECS>, BN>;   // [BN][BK/4]
+{% if not alignedBlockLoads %}
+// Generic unpack used only when a quant block does not contain a whole BK tile.
+{% macro matmul_nbits_packed_code(fn="packed_weight", buffer="b", kBlocks="params.kBlocks", blobSize="params.blobSize", bits=4) %}
+fn {{ fn }}(n: u32, block: u32, offset: u32) -> u32 {
+{% if bits == 2 %}
+  let byte_index = offset / 4u;
+  let shift = (offset % 4u) * 2u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 3u;
+{% elif bits == 4 %}
+  let byte_index = offset / 2u;
+  let shift = (offset % 2u) * 4u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 15u;
+{% else %}
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + offset;
+  return {{ buffer }}[packed_index] & 255u;
+{% endif %}
+}
+{%- endmacro %}
+{{- matmul_nbits_packed_code(bits=bits, kBlocks="KBLOCKS", blobSize="BLOB_SIZE") }}
+{% endif %}
+{% macro zero_of(blockExpr) %}{% if hasZero %}f32(zero_points[bn * KBLOCKS + {{ blockExpr }}]){% else %}{{ defaultZero }}{% endif %}{% endmacro %}
+@compute @workgroup_size(16, 16, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let mBase = wg.y * BM;
+  let nBase = wg.x * BN;
+  let li = lid.y * WG_X + lid.x;
+{% for i in range(threadRows) %}
+{% for j in range(threadCols) %}
+  var acc{{ i }}{{ j }}: f32 = 0.0;
+{% endfor %}
+{% endfor %}
+  let numTiles = (K + BK - 1u) / BK;
+  for (var kt: u32 = 0u; kt < numTiles; kt = kt + 1u) {
+    let kBase = kt * BK;
+    // A and B may have different tile areas (32x64 portable geometry), so stage
+    // them independently. Every configured area is an exact WG_THREADS multiple.
+    for (var idx: u32 = li; idx < BM * K_VECS; idx = idx + WG_THREADS) {
+      let ar = idx / K_VECS;
+      let ac4 = idx % K_VECS;
+      let am = mBase + ar;
+      let ak = kBase + ac4 * 4u;
+      var aWord = vec4<{{ aTile }}>({{ aTile }}(0.0));
+      if (am < M) {
+{% for component in range(4) %}
+        if (ak + {{ component }}u < K) { aWord[{{ component }}u] = {{ aTile }}(a[am * K + ak + {{ component }}u]); }
+{% endfor %}
+      }
+      tileA[ar][ac4] = aWord;
+    }
+{% if alignedBlockLoads %}
+    // Every lane materializes one vector word of one output column. BLOCK_SIZE is
+    // a multiple of BK, so the whole tile slice of a column shares one scale and
+    // zero point and each stored byte is read once for the codes it carries.
+    for (var idx: u32 = li; idx < BN * K_VECS; idx = idx + WG_THREADS) {
+      let bc = idx / K_VECS;
+      let kv = idx % K_VECS;
+      let bn = nBase + bc;
+      var word = vec4<f32>(0.0);
+      if (bn < N && kBase < K) {
+        let block = kBase / BLOCK_SIZE;
+        let offset0 = kBase % BLOCK_SIZE;
+        let scale = f32(scales[bn * KBLOCKS + block]);
+        let zero = {{ zero_of("block") }};
+        let blob_base = (bn * KBLOCKS + block) * BLOB_SIZE;
+        let kk = kv * 4u;
+{% if bits == 2 %}
+        // Four 2-bit codes share a stored byte, and an aligned tile starts a
+        // word on a byte boundary, so one byte fills one vector word.
+        let packed = b[blob_base + (offset0 + kk) / 4u];
+{% for component in range(4) %}
+        if (kBase + kk + {{ component }}u < K) {
+          word[{{ component }}u] = (f32((packed >> {{ component * 2 }}u) & 3u) - zero) * scale;
+        }
+{% endfor %}
+{% elif bits == 4 %}
+        // Two bytes carry the four K-adjacent codes of one vector word.
+        let byte_index = blob_base + (offset0 + kk) / 2u;
+        let packed0 = b[byte_index];
+        let packed1 = b[byte_index + 1u];
+{% for component in range(4) %}
+        if (kBase + kk + {{ component }}u < K) {
+          word[{{ component }}u] = (f32(({{ "packed0" if component < 2 else "packed1" }} >> {{ (component % 2) * 4 }}u) & 15u) - zero) * scale;
+        }
+{% endfor %}
+{% else %}
+        let byte_index = blob_base + offset0 + kk;
+{% for component in range(4) %}
+        if (kBase + kk + {{ component }}u < K) {
+          word[{{ component }}u] = (f32(b[byte_index + {{ component }}u] & 255u) - zero) * scale;
+        }
+{% endfor %}
+{% endif %}
+      }
+      tileB[bc][kv] = word;
+    }
+{% else %}
+    // Quant blocks that do not contain a whole BK tile use the element-wise loader.
+    for (var idx: u32 = li; idx < BN * K_VECS; idx = idx + WG_THREADS) {
+      let bc = idx / K_VECS;
+      let br4 = idx % K_VECS;
+      let bk = kBase + br4 * 4u;
+      let bn = nBase + bc;
+      var word = vec4<f32>(0.0);
+      if (bn < N) {
+{% for component in range(4) %}
+        if (bk + {{ component }}u < K) {
+          let block{{ component }} = (bk + {{ component }}u) / BLOCK_SIZE;
+          let offset{{ component }} = (bk + {{ component }}u) % BLOCK_SIZE;
+          let scale{{ component }} = f32(scales[bn * KBLOCKS + block{{ component }}]);
+          let zero{{ component }} = {{ zero_of("block" ~ component) }};
+          word[{{ component }}u] = (f32(packed_weight(bn, block{{ component }}, offset{{ component }})) - zero{{ component }}) * scale{{ component }};
+        }
+{% endfor %}
+      }
+      tileB[bc][br4] = word;
+    }
+{% endif %}
+    workgroupBarrier();
+    let aRow = lid.y * TM;
+    let bCol = lid.x * TN;
+    for (var kv: u32 = 0u; kv < K_VECS; kv = kv + 1u) {
+{% for i in range(threadRows) %}
+      let av{{ i }} = {{ aRead }}tileA[aRow + {{ i }}u][kv]{{ aReadEnd }};
+{% endfor %}
+{% for j in range(threadCols) %}
+      let bv{{ j }} = tileB[bCol + {{ j }}u][kv];
+{% endfor %}
+{% for i in range(threadRows) %}
+{% for j in range(threadCols) %}
+      acc{{ i }}{{ j }} = acc{{ i }}{{ j }} + dot(av{{ i }}, bv{{ j }});
+{% endfor %}
+{% endfor %}
+    }
+    workgroupBarrier();
+  }
+  let m0 = mBase + lid.y * TM;
+  let n0 = nBase + lid.x * TN;
+{% for i in range(threadRows) %}
+  if (m0 + {{ i }}u < M) {
+{% for j in range(threadCols) %}
+    if (n0 + {{ j }}u < N) {
+      y[(m0 + {{ i }}u) * N + n0 + {{ j }}u] = {{ outputScalar }}(acc{{ i }}{{ j }}{% if hasBias %} + f32(bias[n0 + {{ j }}u]){% endif %});
+    }
+{% endfor %}
+  }
+{% endfor %}
+}

build/webgpu/matmul-nbits-q4-prefill-tiled.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,130 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// com.microsoft.MatMulNBits q4/q8 prefill, no-subgroup-matrix tier:
+//   Y[M,N] = A[M,K] @ W[K,N] (+ bias),  W[k,n] = (q(n,block,offset) - zero) * scale
+// Register-blocked 32x32 shared-memory tiled GEMM (16x16 threads, each computes a
+// 2x2 micro-tile, K stepped in BK=16 chunks). The B operand is dequantized from
+// the packed blob straight into the shared B tile (no f32 weight matrix stored).
+// Each staged operand is reused across the micro-tile, and f32 accumulation is
+// exact for the 4/8-bit dequantized weight. Fully bounds-checked, so any M/N is
+// supported when K % blockSize == 0.
+const M: u32 = {{ M }}u;
+const K: u32 = {{ K }}u;
+const N: u32 = {{ N }}u;
+const KBLOCKS: u32 = {{ kBlocks }}u;
+const BLOCK_SIZE: u32 = {{ blockSize }}u;
+const BLOB_SIZE: u32 = {{ blobSize }}u;
+const BK: u32 = 16u;
+const BM: u32 = 32u;
+const BN: u32 = 32u;
+// When A uses f16 storage, staging tileA at f16 is bit-identical to widening it
+// here. tileB stays f32 because the dequantized (code - zero)*scale product
+// needs more mantissa than f16 provides.
+{% set aTile = "f16" if usesF16 else "f32" %}
+{% set aRead = "f32(" if usesF16 else "" %}
+{% set aReadEnd = ")" if usesF16 else "" %}
+var<workgroup> tileA: array<array<{{ aTile }}, 16>, 32>;
+var<workgroup> tileB: array<array<f32, 32>, 16>;
+// {{ bits }}-bit unpack: q2 stores four codes per byte, q4 two, and q8 one.
+{% macro matmul_nbits_packed_code(fn="packed_weight", buffer="b", kBlocks="params.kBlocks", blobSize="params.blobSize", bits=4) %}
+fn {{ fn }}(n: u32, block: u32, offset: u32) -> u32 {
+{% if bits == 2 %}
+  let byte_index = offset / 4u;
+  let shift = (offset % 4u) * 2u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 3u;
+{% elif bits == 4 %}
+  let byte_index = offset / 2u;
+  let shift = (offset % 2u) * 4u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 15u;
+{% else %}
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + offset;
+  return {{ buffer }}[packed_index] & 255u;
+{% endif %}
+}
+{%- endmacro %}
+{{- matmul_nbits_packed_code(bits=bits, kBlocks="KBLOCKS", blobSize="BLOB_SIZE") }}
+@compute @workgroup_size(16, 16, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let mBase = wg.y * BM;
+  let nBase = wg.x * BN;
+  let li = lid.y * 16u + lid.x;
+  var acc00: f32 = 0.0;
+  var acc01: f32 = 0.0;
+  var acc10: f32 = 0.0;
+  var acc11: f32 = 0.0;
+  let numTiles = (K + BK - 1u) / BK;
+  for (var kt: u32 = 0u; kt < numTiles; kt = kt + 1u) {
+    let kBase = kt * BK;
+    // Cooperative load: 32x16 A tile + 16x32 B tile, 256 threads x 2 each. The B
+    // tile is dequantized from the packed q blob during the load.
+    for (var e: u32 = 0u; e < 2u; e = e + 1u) {
+      let idx = li + e * 256u;
+      let ar = idx / BK;
+      let ac = idx % BK;
+      let am = mBase + ar;
+      let ak = kBase + ac;
+      if (am < M && ak < K) {
+        tileA[ar][ac] = {{ aTile }}(a[am * K + ak]);
+      } else {
+        tileA[ar][ac] = 0.0;
+      }
+      let br = idx / BN;
+      let bc = idx % BN;
+      let bk = kBase + br;
+      let bn = nBase + bc;
+      if (bk < K && bn < N) {
+        let block = bk / BLOCK_SIZE;
+        let offset = bk % BLOCK_SIZE;
+        let scale = f32(scales[bn * KBLOCKS + block]);
+{% if hasZero %}
+        let zero = f32(zero_points[bn * KBLOCKS + block]);
+{% else %}
+        let zero = {{ defaultZero }};
+{% endif %}
+        tileB[br][bc] = (f32(packed_weight(bn, block, offset)) - zero) * scale;
+      } else {
+        tileB[br][bc] = 0.0;
+      }
+    }
+    workgroupBarrier();
+    for (var kk: u32 = 0u; kk < BK; kk = kk + 1u) {
+      let a0 = {{ aRead }}tileA[lid.y * 2u][kk]{{ aReadEnd }};
+      let a1 = {{ aRead }}tileA[lid.y * 2u + 1u][kk]{{ aReadEnd }};
+      let b0 = tileB[kk][lid.x * 2u];
+      let b1 = tileB[kk][lid.x * 2u + 1u];
+      acc00 = acc00 + a0 * b0;
+      acc01 = acc01 + a0 * b1;
+      acc10 = acc10 + a1 * b0;
+      acc11 = acc11 + a1 * b1;
+    }
+    workgroupBarrier();
+  }
+  let m0 = mBase + lid.y * 2u;
+  let m1 = m0 + 1u;
+  let n0 = nBase + lid.x * 2u;
+  let n1 = n0 + 1u;
+{% if hasBias %}
+  if (m0 < M && n0 < N) { y[m0 * N + n0] = {{ outputScalar }}(acc00 + f32(bias[n0])); }
+  if (m0 < M && n1 < N) { y[m0 * N + n1] = {{ outputScalar }}(acc01 + f32(bias[n1])); }
+  if (m1 < M && n0 < N) { y[m1 * N + n0] = {{ outputScalar }}(acc10 + f32(bias[n0])); }
+  if (m1 < M && n1 < N) { y[m1 * N + n1] = {{ outputScalar }}(acc11 + f32(bias[n1])); }
+{% else %}
+  if (m0 < M && n0 < N) { y[m0 * N + n0] = {{ outputScalar }}(acc00); }
+  if (m0 < M && n1 < N) { y[m0 * N + n1] = {{ outputScalar }}(acc01); }
+  if (m1 < M && n0 < N) { y[m1 * N + n0] = {{ outputScalar }}(acc10); }
+  if (m1 < M && n1 < N) { y[m1 * N + n1] = {{ outputScalar }}(acc11); }
+{% endif %}
+}

build/webgpu/matmul-nbits-q4-sgmat.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,193 @@

+// com.microsoft.MatMulNBits q4 prefill subgroup-matrix tier:
+//   Y[M,N] = A[M,K] @ W[K,N] + bias?,  W[k,n] = (q4(n,block,offset) - zero) * scale
+// Uses a 32x64 or 64x64 subgroup-matrix output tile with an M-tail row_limit guard; the
+// B operand is materialized on the fly by dequantizing the packed q4 blob into
+// the shared B tile, so no f32 weight matrix is stored. f32 operands and f32
+// accumulation preserve precision; the dequantized weight is a 4-bit value, so
+// f32 is exact for it.
+// Alignment gates keep K % 32 == 0 and N % 64 == 0 so subgroupMatrixLoad never
+// sees partial 8x8 tiles; M tails handled by row_limit + the loadSHMA zero-fill.
+enable subgroups;
+enable chromium_experimental_subgroup_matrix;
+{% if usesF16 %}
+enable f16;
+{% endif %}
+diagnostic(off, chromium.subgroup_matrix_uniformity);
+{{ env.wgsl.resourceDeclarations }}
+const M: u32 = {{ M }}u;
+const K: u32 = {{ K }}u;
+const N: u32 = {{ N }}u;
+const KBLOCKS: u32 = {{ kBlocks }}u;
+const BLOCK_SIZE: u32 = {{ blockSize }}u;
+const BLOB_SIZE: u32 = {{ blobSize }}u;
+const TILE_COLS: u32 = 64u;
+const TILE_ROWS: u32 = {{ tileRows }}u;
+const TILE_K: u32 = 32u;
+const SUB_COLS: u32 = 32u;
+const SUB_ROWS: u32 = 16u;
+var<workgroup> tile_A: array<f32, {{ tileRows }} * 32>;
+var<workgroup> tile_B: array<f32, 64 * 32>;
+var<workgroup> scratch: array<array<array<f32, 64>, 4>, {{ numSubgroups }}>;
+fn loadSHMA(tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
+  let a_global = tile_base + row;
+  let col = c_idx * 8u;
+  for (var col_offset = 0u; col_offset < 8u; col_offset = col_offset + 1u) {
+    let k = k_idx + col + col_offset;
+    if (a_global < M) {
+      tile_A[row * TILE_K + col + col_offset] = f32(a[a_global * K + k]);
+    } else {
+      tile_A[row * TILE_K + col + col_offset] = 0.0;
+    }
+  }
+}
+fn loadSHMB(tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
+  let n = tile_base + row;            // output column
+  let col = c_idx * {{ bLoadWidth }}u;
+  // When N % 64 != 0 the last column-tile is padded past N. Those columns
+  // dequantize to 0 (so they contribute 0 to the matrix-multiply and are never
+  // stored), and their scale/zero/packed reads would be out of bounds — skip them.
+  let valid_n = n < N;
+  if (valid_n) {
+    // The aligned loader requires BLOCK_SIZE % 16 == 0. k_idx is 32-aligned and
+    // each lane owns an aligned 8/16-element slice, so the whole slice shares
+    // one scale/zero-point. Hoist those reads and unpack each q4 byte once.
+    let k0 = k_idx + col;
+    let block = k0 / BLOCK_SIZE;
+    let offset0 = k0 % BLOCK_SIZE;
+    let scale = f32(scales[n * KBLOCKS + block]);
+{% if hasZero %}
+    let zero = f32(zero_points[n * KBLOCKS + block]);
+{% else %}
+    let zero = {{ defaultZero }};
+{% endif %}
+{% if bits == 2 %}
+    // Four 2-bit codes share a stored byte; one code per iteration needs no assumption
+    // about where offset0 lands inside a byte.
+    for (var i = 0u; i < {{ bLoadWidth }}u; i = i + 1u) {
+      let code_offset = offset0 + i;
+      let packed_index = (n * KBLOCKS + block) * BLOB_SIZE + code_offset / 4u;
+      tile_B[row * TILE_K + col + i] = (f32((b[packed_index] >> ((code_offset % 4u) * 2u)) & 3u) - zero) * scale;
+    }
+{% elif bits == 4 %}
+    for (var i = 0u; i < {{ bLoadWidth }}u; i = i + 2u) {
+      let packed_index = (n * KBLOCKS + block) * BLOB_SIZE + (offset0 + i) / 2u;
+      let packed = b[packed_index];
+      tile_B[row * TILE_K + col + i] = (f32(packed & 15u) - zero) * scale;
+      tile_B[row * TILE_K + col + i + 1u] = (f32((packed >> 4u) & 15u) - zero) * scale;
+    }
+{% else %}
+    for (var i = 0u; i < {{ bLoadWidth }}u; i = i + 1u) {
+      let packed_index = (n * KBLOCKS + block) * BLOB_SIZE + offset0 + i;
+      tile_B[row * TILE_K + col + i] = (f32(b[packed_index] & 255u) - zero) * scale;
+    }
+{% endif %}
+  } else {
+    for (var i = 0u; i < {{ bLoadWidth }}u; i = i + 1u) {
+      tile_B[row * TILE_K + col + i] = 0.0;
+    }
+  }
+}
+// Store one output column: slot ∈ {0,1,2,3} selects the 8-wide result block, cc
+// ∈ {col, col+1} the lane pair. The absolute column is col_base + cc + slot*8;
+// when N % 64 != 0 the padded tile columns (>= N) are skipped (the guard is
+// always true for aligned N, so this is bit-identical there).
+fn storeCol(offset: u32, col_base: u32, row: u32, cc: u32, slot: u32, src_slot: u32) {
+  let c = cc + slot * 8u;
+  if (col_base + c < N) {
+{% if hasBias %}
+    y[offset + row * N + c] = {{ outputScalar }}(scratch[src_slot][slot][row * 8u + cc] + f32(bias[col_base + c]));
+{% else %}
+    y[offset + row * N + c] = {{ outputScalar }}(scratch[src_slot][slot][row * 8u + cc]);
+{% endif %}
+  }
+}
+fn storeOutput(offset: u32, col_base: u32, row: u32, col: u32, src_slot: u32, row_limit: i32) {
+  if (row_limit > 0 && row < u32(row_limit)) {
+    let col2 = col + 1u;
+{% for blockCol in [0, 1, 2, 3] %}
+    storeCol(offset, col_base, row, col,  {{ blockCol }}u, src_slot);
+    storeCol(offset, col_base, row, col2, {{ blockCol }}u, src_slot);
+{% endfor %}
+  }
+}
+@compute @workgroup_size({{ workgroupSize }}, 1, 1)
+fn main(
+  @builtin(workgroup_id) workgroup_id: vec3<u32>,
+  @builtin(local_invocation_index) local_idx: u32,
+  @builtin(subgroup_invocation_id) sg_id: u32,
+  @builtin(subgroup_size) sg_size: u32
+) {
+  let a_global_base = workgroup_id.y * TILE_ROWS;
+  let b_global_base = workgroup_id.x * TILE_COLS;
+  let subtile_id = local_idx / sg_size;
+  let subtile_idx = subtile_id / {{ rowSubtiles }}u;
+  let subtile_idy = subtile_id % {{ rowSubtiles }}u;
+  let base_A = subtile_idy * SUB_ROWS;
+  let base_B = subtile_idx * SUB_COLS;
+{% for row in [0, 1] %}{% for col in [0, 1, 2, 3] %}
+  var matC{{ row }}{{ col }}: subgroup_matrix_result<f32, 8, 8>;
+{% endfor %}{% endfor %}
+  for (var kidx = 0u; kidx < K; kidx = kidx + TILE_K) {
+    loadSHMA(a_global_base, kidx, local_idx / 4u, local_idx % 4u);
+    loadSHMB(b_global_base, kidx, local_idx / {{ bLoadsPerRow }}u, local_idx % {{ bLoadsPerRow }}u);
+    workgroupBarrier();
+    for (var step = 0u; step < TILE_K; step = step + 8u) {
+{% set operandScalar = "f32" %}      let matrix_a_offset = subtile_idy * SUB_ROWS * TILE_K + step;
+      var matA0: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>>(&tile_A, matrix_a_offset, false, TILE_K);
+      var matA1: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>>(&tile_A, matrix_a_offset + 8u * TILE_K, false, TILE_K);
+      let matrix_b_offset = subtile_idx * SUB_COLS * TILE_K + step;
+      var matB0: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>>(&tile_B, matrix_b_offset, true, TILE_K);
+      var matB1: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>>(&tile_B, matrix_b_offset + 8u * TILE_K, true, TILE_K);
+      var matB2: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>>(&tile_B, matrix_b_offset + 16u * TILE_K, true, TILE_K);
+      var matB3: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>>(&tile_B, matrix_b_offset + 24u * TILE_K, true, TILE_K);
+      matC00 = subgroupMatrixMultiplyAccumulate(matA0, matB0, matC00);
+      matC01 = subgroupMatrixMultiplyAccumulate(matA0, matB1, matC01);
+      matC02 = subgroupMatrixMultiplyAccumulate(matA0, matB2, matC02);
+      matC03 = subgroupMatrixMultiplyAccumulate(matA0, matB3, matC03);
+      matC10 = subgroupMatrixMultiplyAccumulate(matA1, matB0, matC10);
+      matC11 = subgroupMatrixMultiplyAccumulate(matA1, matB1, matC11);
+      matC12 = subgroupMatrixMultiplyAccumulate(matA1, matB2, matC12);
+      matC13 = subgroupMatrixMultiplyAccumulate(matA1, matB3, matC13);
+    }
+    workgroupBarrier();
+  }
+  // The four scratch banks are reused across the two row-groups, and each is written
+  // by a collective subgroupMatrixStore then read across lanes by storeOutput. Barriers
+  // give the reads visibility of the store and stop the second row-group's store from
+  // clobbering the first's still-in-flight readback when a partial final M-tile
+  // diverges storeOutput's guard. Without both barriers the last valid row can be corrupted.
+{% for col in [0, 1, 2, 3] %}
+  subgroupMatrixStore(&scratch[subtile_id][{{ col }}], 0u, matC0{{ col }}, false, 8u);
+{% endfor %}
+  workgroupBarrier();
+  let row = sg_id / 4u;
+  let col = (sg_id % 4u) * 2u;
+  let col_base = b_global_base + base_B;
+  var matrix_c_offset = (a_global_base + base_A) * N + col_base;
+  var row_limit = i32(M) - i32(a_global_base + base_A);
+  storeOutput(matrix_c_offset, col_base, row, col, subtile_id, row_limit);
+  workgroupBarrier();
+{% for col in [0, 1, 2, 3] %}
+  subgroupMatrixStore(&scratch[subtile_id][{{ col }}], 0u, matC1{{ col }}, false, 8u);
+{% endfor %}
+  workgroupBarrier();
+  matrix_c_offset = matrix_c_offset + 8u * N;
+  row_limit = i32(M) - i32(a_global_base + base_A + 8u);
+  storeOutput(matrix_c_offset, col_base, row, col, subtile_id, row_limit);
+}

build/webgpu/matmul-nbits.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,69 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+// Unpack one {{ bits }}-bit quantized weight, low codes first along K. q2 packs four
+// codes per stored byte, q4 two nibbles, and q8 is one byte per weight
+// (blobSize == blockSize). `b` is bound u32-per-byte in every case, so all three
+// index the stored byte directly.
+{% macro matmul_nbits_packed_code(fn="packed_weight", buffer="b", kBlocks="params.kBlocks", blobSize="params.blobSize", bits=4) %}
+fn {{ fn }}(n: u32, block: u32, offset: u32) -> u32 {
+{% if bits == 2 %}
+  let byte_index = offset / 4u;
+  let shift = (offset % 4u) * 2u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 3u;
+{% elif bits == 4 %}
+  let byte_index = offset / 2u;
+  let shift = (offset % 2u) * 4u;
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + byte_index;
+  return ({{ buffer }}[packed_index] >> shift) & 15u;
+{% else %}
+  let packed_index = (n * {{ kBlocks }} + block) * {{ blobSize }} + offset;
+  return {{ buffer }}[packed_index] & 255u;
+{% endif %}
+}
+{%- endmacro %}
+{{- matmul_nbits_packed_code(bits=bits) }}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded flat output-element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit. When nwg.y == 1 this
+  // reduces to gid.x; the index >= total guard drops the tail.
+  let index = gid.x + gid.y * nwg.x * WG;
+  let total = params.rows * params.N;
+  if (index >= total) {
+    return;
+  }
+  let row = index / params.N;
+  let col = index % params.N;
+{% if hasBias %}
+  var acc = f32(bias[col]);
+{% else %}
+  var acc = 0.0;
+{% endif %}
+  for (var k: u32 = 0u; k < params.K; k = k + 1u) {
+    let block = k / params.blockSize;
+    let offset = k % params.blockSize;
+    let q = packed_weight(col, block, offset);
+{% if hasZero %}
+    let zero = f32(zero_points[col * params.kBlocks + block]);
+{% else %}
+    let zero = {{ defaultZero }};
+{% endif %}
+    let w = (f32(q) - zero) * f32(scales[col * params.kBlocks + block]);
+    acc = acc + f32(a[row * params.K + k]) * w;
+  }
+  y[index] = {{ outputScalar }}(acc);
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,25 @@

+{
+  "name": "com.microsoft.MatMulNBits",
+  "id": "_com_microsoft_matmulnbits_webgpu_81b817c",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "BbpRtHBH/vkJDaybyRBNKhuUCpe4XMJtKmW3bhRcEZE=",
+      "manifest.json": "Gz92F8En4SeyVGrSZBtQghet4kVvjHxGIplFWSIq7AM=",
+      "matmul-nbits-dp4a-quantize.wgsl.jinja": "WEUGSDH5qWYg10f/mkAHxq8AkIkMW9R/Wn/Y23Q0Jw0=",
+      "matmul-nbits-gemv-q4.wgsl.jinja": "r9eGhKkZPaNaUHFcDZh1JpS+w6z4Wow/7wGrLMVVgsA=",
+      "matmul-nbits-q4-dp4a-prefill.wgsl.jinja": "qUrURZDpPzMpABnTgESkHgLCQiAWxEZ8wvzkAPTW2kI=",
+      "matmul-nbits-q4-prefill-tile4x4.wgsl.jinja": "HjZraSjOpaSl5Cjns4scselOIn+ibxFEu+9LTeUI0QI=",
+      "matmul-nbits-q4-prefill-tiled-reg.wgsl.jinja": "oU7BcxJnKIEfgnkWHPq77RrCmrE0iIm6ED7cgVfZUYs=",
+      "matmul-nbits-q4-prefill-tiled.wgsl.jinja": "TdetCiv5o88hvWCGNDewpFX+M+NZrz9rekU/d5kMZTM=",
+      "matmul-nbits-q4-sgmat.wgsl.jinja": "8E4yGRvezICz5ctdW4g2w/HBfLvqJF2hA+jOobcuJio=",
+      "matmul-nbits.wgsl.jinja": "DRkwODEER11pYQxt2ssS9YvFggO3CbUqQ9yag53J5lY=",
+      "test.json": "2mpJ5fx2GnlZycPygUyHbOM37n7X6VZmUMGHNkJa7qw="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.MatMulNBits" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1967 @@

+{
+  "op": "com.microsoft.MatMulNBits",
+  "fixtureArrays": {
+    "q4_prefill_tiled_reg_tailk_m32_k33_n4096_input_bT": [16, 50, 84, 118, 152, 186, 220, 254, 135, 119, 102, 85, 68, 51, 34, 17],
+    "q4_weight_cycle_b_t": [16, 50, 84, 118, 152, 186, 220, 254, 33, 67, 101, 135, 169, 203, 237, 15],
+    "quant_scale_cycle_t": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07, 0.085],
+    "mixed_weight_cycle_b_t": [17, 200, 91, 45, 233, 128, 7, 176, 250, 33, 142, 99, 210, 64, 188, 121],
+    "q8_zero_bias_gemv_m1_tail_n5_input_bT": [19, 56, 93, 130, 167, 204, 241, 22, 59, 96, 133, 170, 207, 244, 25, 62]
+  },
+  "cases": [
+    {
+      "name": "q4_zero_bias_prefill_tile4x4_large",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [6.0, 7.0, 8.0, 9.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.1 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.0001 } },
+      "attrs": { "K": 128, "N": 64, "block_size": 32 }
+    },
+    {
+      "name": "q4_zero_bias_prefill_tile4x4_small_m8",
+      "provenance": {
+        "notes": "Small-M (M=8) q4 prefill with bias and zero points. M<64 excludes prefill_tiled_zero_bias, while the row-guarded prefill_tile4x4_zero_bias route admits M>=2 when N is divisible by 4. This pins the tile4x4 lower-bound contract and its bias/zero-point arithmetic."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [6.0, 7.0, 8.0, 9.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.1 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 64], "tolerance": 0.0001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_bias_prefill_tile4x4_small_m8",
+      "provenance": {
+        "notes": "q8 companion to q4_zero_bias_prefill_tile4x4_small_m8: the relaxed tile4x4 bias floor (M>=2) also makes prefill_tile4x4_zero_bias selectable for q8 prefill+bias at M in [2,64). Validates the q8 unpack branch + bias epilogue on the tile4x4 path at small M."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0, 144.0, 126.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "ort_projection_q4_block16_tail_zero_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/matmul_4bits_test.cc",
+        "test": "MatMulNBits.Float32_4b_Accuracy0",
+        "notes": "Rank-2 projection of ORT q4 block-size coverage with per-block zero points and bias."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [2, 17],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.37 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2, 2, 8],
+          "data": {
+            "kind": "values",
+            "values": [16, 50, 84, 118, 152, 186, 220, 254, 33, 0, 0, 0, 0, 0, 0, 0, 255, 238, 221, 204, 187, 170, 153, 136, 240, 0, 0, 0, 0, 0, 0, 0]
+          }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.04, 0.08, 0.06, 0.12] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [7.0, 1.0, 8.0, 0.0] }
+        },
+        "biasT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.5, -1.0] } }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.00001 } },
+      "attrs": { "K": 17, "N": 2, "bits": 4, "block_size": 16 }
+    },
+    {
+      "name": "ort_projection_q4_block32_tail_no_zero",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/matmul_4bits_test.cc",
+        "test": "MatMulNBits.Float32_4b_Accuracy0",
+        "notes": "Rank-2 deterministic projection of ORT q4 coverage with blockSize=32 and tail K."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [2, 33],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.75 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2, 2, 16],
+          "data": {
+            "kind": "values",
+            "values": [16, 50, 84, 118, 152, 186, 220, 254, 16, 50, 84, 118, 152, 186, 220, 254, 135, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 255, 238, 221, 204, 187, 170, 153, 136, 119, 102, 85, 68, 51, 34, 17, 0, 248, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]
+          }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [0.03, 0.06, 0.04, 0.07] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.00001 } },
+      "attrs": { "K": 33, "N": 2, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_block512_schema_legal_cpu_provider_gap_exact",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/matmul_4bits_test.cc",
+        "test": "MatMulNBits.UnsupportedBlockSize_512",
+        "notes": "The public schema permits every power-of-two block size of at least 16, and this package implements block_size=512. ORT's installed, pinned, and inspected current CPU providers instead restrict execution to 16, 32, 64, 128, or 256. With A=1, packed q4 codes=9, default zero point=8, and scale=0.25, the exact result is 512 * (9 - 8) * 0.25 = 128."
+      },
+      "inputs": {
+        "aT": { "dtype": "float32", "shape": [1, 512], "data": { "kind": "constant", "value": 1.0 } },
+        "bT": { "dtype": "uint8", "shape": [1, 1, 256], "data": { "kind": "constant", "value": 153 } },
+        "scalesT": { "dtype": "float32", "shape": [1, 1], "data": { "kind": "constant", "value": 0.25 } }
+      },
+      "outputs": {
+        "yT": { "dtype": "float32", "shape": [1, 1], "tolerance": 0, "data": { "kind": "values", "values": [128.0] } }
+      },
+      "attrs": { "K": 512, "N": 1, "bits": 4, "block_size": 512 }
+    },
+    {
+      "name": "q4_no_zero_prefill_sgmat_large",
+      "inputs": {
+        "aT": {
+          "shape": [128, 128],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "shape": [64, 4, 16],
+          "dtype": "uint8",
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "shape": [64, 4],
+          "dtype": "float32",
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.0001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_no_zero_prefill_sgmat_m32_storage_floor",
+      "provenance": {
+        "notes": "The 32-row subgroup-matrix tile fits the WebGPU minimum 16 KiB workgroup-storage limit."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.0001 } },
+      "attrs": { "K": 32, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_zero_bias_prefill_sgmat_m32_storage_floor",
+      "provenance": { "notes": "Zero-point/bias sibling of the 32-row subgroup-matrix storage-floor regression." },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [7.0, 8.0, 6.0, 9.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.0001 } },
+      "attrs": { "K": 32, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_no_zero_prefill_g64_sgmat",
+      "provenance": {
+        "notes": "Compact blockSize=64 q4 default-zero prefill sibling for model g64 benchmarks. Full tier routes through prefill_sgmat_default_zero; no-MMA/no-subgroups/baseline route through prefill_tiled_default_zero."
+      },
+      "inputs": {
+        "aT": {
+          "shape": [64, 128],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "shape": [64, 2, 32],
+          "dtype": "uint8",
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "shape": [64, 2],
+          "dtype": "float32",
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [64, 64], "tolerance": 0.0001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 4, "block_size": 64 }
+    },
+    {
+      "name": "q8_no_zero_prefill_sgmat_large",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_no_zero_prefill_sgmat_m32_storage_floor",
+      "provenance": {
+        "notes": "q8 sibling proving the 32-row subgroup-matrix tile remains available at the 16 KiB storage floor."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 1, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 32, "N": 64, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_bias_prefill_sgmat_large",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0, 144.0, 126.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_bias_prefill_sgmat_m32_storage_floor",
+      "provenance": { "notes": "q8 zero-point/bias sibling of the 32-row subgroup-matrix storage-floor regression." },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 1, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 32, "N": 64, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_no_zero_prefill_tile4x4_unaligned",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [8, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [12, 2, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [12, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 12], "tolerance": 0.001 } },
+      "attrs": { "K": 64, "N": 12, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q4_f16_zero_points_prefill_sgmat",
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float16",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float16",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [7.0, 8.0, 9.0, 6.0] }
+        },
+        "biasT": { "dtype": "float16", "shape": [64], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [128, 64], "tolerance": 0.05 } },
+      "attrs": { "K": 128, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_no_zero_prefill_sgmat_unaligned_n72",
+      "inputs": {
+        "aT": {
+          "shape": [128, 128],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "shape": [72, 4, 16],
+          "dtype": "uint8",
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "shape": [72, 4],
+          "dtype": "float32",
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 72], "tolerance": 0.0001 } },
+      "attrs": { "K": 128, "N": 72, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q8_no_zero_prefill_sgmat_unaligned_n72",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [72, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [72, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 72], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 72, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_no_zero_prefill_sgmat_unaligned_n72_m32",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [72, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [72, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 72], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 72, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_bias_prefill_sgmat_unaligned_n72",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [72, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [72, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [72, 4],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0, 144.0, 126.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [72],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 72], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 72, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_bias_prefill_sgmat_unaligned_n65_m33",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [33, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [65, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [65, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [65, 4],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0, 144.0, 126.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [65],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [33, 65], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 65, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_bias_prefill_sgmat_unaligned_n72_m32_bs16",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [72, 8, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [72, 8],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [72, 8],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0, 144.0, 126.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [72],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 72], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 72, "bits": 8, "block_size": 16 }
+    },
+    {
+      "name": "empty_zero_dim",
+      "provenance": {
+        "notes": "Zero-length M (rows) axis of A (empty input). With M=0 the N/K-driven packed uint8 B[4,1,16] and scales[4,1] stay normal and Y=[0,N] is empty. K=32,N=4,bits=4,block_size=32."
+      },
+      "inputs": {
+        "aT": { "dtype": "float32", "shape": [0, 32], "data": { "kind": "values", "values": [] } },
+        "bT": { "dtype": "uint8", "shape": [4, 1, 16], "data": { "kind": "constant", "value": 17 } },
+        "scalesT": { "dtype": "float32", "shape": [4, 1], "data": { "kind": "constant", "value": 0.05 } }
+      },
+      "outputs": {
+        "yT": { "dtype": "float32", "shape": [0, 4], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      },
+      "attrs": { "K": 32, "N": 4, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "empty_zero_dim_f16",
+      "provenance": { "notes": "f16 variant of the empty zero-M-axis case." },
+      "inputs": {
+        "aT": { "dtype": "float16", "shape": [0, 32], "data": { "kind": "values", "values": [] } },
+        "bT": { "dtype": "uint8", "shape": [4, 1, 16], "data": { "kind": "constant", "value": 17 } },
+        "scalesT": { "dtype": "float16", "shape": [4, 1], "data": { "kind": "constant", "value": 0.05 } }
+      },
+      "outputs": {
+        "yT": { "dtype": "float16", "shape": [0, 4], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      },
+      "attrs": { "K": 32, "N": 4, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_no_zero_prefill_reg_m512_n1024",
+      "inputs": {
+        "aT": {
+          "shape": [512, 128],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "shape": [1024, 4, 16],
+          "dtype": "uint8",
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "shape": [1024, 4],
+          "dtype": "float32",
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 1024], "tolerance": 0.0001 } },
+      "attrs": { "K": 128, "N": 1024, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_f16_prefill_reg_m512_n1024",
+      "inputs": {
+        "aT": {
+          "shape": [512, 128],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "shape": [1024, 4, 16],
+          "dtype": "uint8",
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "shape": [1024, 4],
+          "dtype": "float16",
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [512, 1024], "tolerance": 0.0001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "f16 activations on the register-tiled prefill path: this is the only shape class where the shared A/B tiles are staged in f16, so it is the fixture that covers that staging."
+      },
+      "attrs": { "K": 128, "N": 1024, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_zero_bias_prefill_reg_m512_n1024",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [512, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024, 4, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": [6.0, 7.0, 8.0, 9.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.1 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 1024], "tolerance": 0.0001 } },
+      "attrs": { "K": 128, "N": 1024, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q8_no_zero_prefill_reg_m512_n1024",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [512, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 1024], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 1024, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_f16_prefill_reg_m512_n1024",
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [512, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float16",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [512, 1024], "tolerance": 0.0001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "q8 sibling of q4_f16_prefill_reg_m512_n1024. 8-bit codes give the widest dequantized range, so the f16 tile store rounds hardest here."
+      },
+      "attrs": { "K": 128, "N": 1024, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_bias_prefill_reg_m512_n1024",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [512, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0, 144.0, 126.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 1024], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 1024, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q4_no_zero_prefill_tile4x4_partial_row_tile_m6",
+      "provenance": {
+        "notes": "M=6 is not a multiple of TILE_M=4, so the tile4x4 kernel's second row-tile (row_base=4) has valid rows 4,5 and guarded rows 6,7. Verifies the store_row partial-row-tile guard writes rows 4,5 correctly and does not corrupt/OOB rows 6,7. N=8 (%4==0), K=64, blockSize=32 routes to prefill_tile4x4_default_zero."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [6, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [8, 2, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [6, 8], "tolerance": 0.0001 } },
+      "attrs": { "K": 64, "N": 8, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_gemv_default_zero_m1_tail_n7",
+      "provenance": {
+        "notes": "M=1 GEMV with N=7 (not a multiple of 4). Exercises the vec4 tail-column store guards (col_base+1/2/3 < N) in matmul-nbits-gemv-q4.wgsl.jinja; existing M==1 tests all use N=4. Routes to gemv_default_zero (prio 20)."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [7, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [7, 1],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 7], "tolerance": 0.00001 } },
+      "attrs": { "K": 32, "N": 7, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_gemv_default_zero_m1_n13_ncols8",
+      "tunables": { "GEMV_N_COLS": 8 },
+      "provenance": {
+        "notes": "GEMV_N_COLS=8 with N=13: two workgroups, first fully live, second with a partially live first group and one live column in the second."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [13, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [13, 1],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 13], "tolerance": 0.00001 } },
+      "attrs": { "K": 32, "N": 13, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_gemv_default_zero_m1_tail_n7_ncols8",
+      "tunables": { "GEMV_N_COLS": 8 },
+      "provenance": {
+        "notes": "GEMV_N_COLS=8 twin of tail_n7: one workgroup, second column group entirely tail-guarded (cols 4-6 active, 7 out of range)."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [7, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [7, 1],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 7], "tolerance": 0.00001 } },
+      "attrs": { "K": 32, "N": 7, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_bias_gemv_m1_tail_n5",
+      "provenance": {
+        "notes": "M=1 q8 GEMV with N=5 (not a multiple of 4), WITH zero_points and bias. Exercises the vec4 tail-column store guards together with the q8 unpack and bias epilogue in matmul-nbits-gemv-q4.wgsl.jinja. Routes to gemv_zero_bias (prio 20)."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [5, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q8_zero_bias_gemv_m1_tail_n5_input_bT" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [5, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [5, 1],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 5], "tolerance": 0.00001 } },
+      "attrs": { "K": 16, "N": 5, "bits": 8, "block_size": 16 }
+    },
+    {
+      "name": "q8_zero_only_gemv_m1_tail_n5",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "M=1 q8 GEMV with explicit zero points and no bias. This pins the zero-points-only binding layout and epilogue instead of relying on the combined zero-point-plus-bias case."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [5, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q8_zero_bias_gemv_m1_tail_n5_input_bT" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [5, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [5, 1],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 5], "tolerance": 0.00001 } },
+      "attrs": { "K": 16, "N": 5, "bits": 8, "block_size": 16 }
+    },
+    {
+      "name": "q8_bias_only_gemv_m1_tail_n5",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "M=1 q8 GEMV with bias and the schema-default zero point. This pins the bias-only binding layout and epilogue instead of relying on the combined zero-point-plus-bias case."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [5, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q8_zero_bias_gemv_m1_tail_n5_input_bT" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [5, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 5], "tolerance": 0.00001 } },
+      "attrs": { "K": 16, "N": 5, "bits": 8, "block_size": 16 }
+    },
+    {
+      "name": "q8_zero_bias_gemv_m1_tail_n5_ncols8",
+      "tunables": { "GEMV_N_COLS": 8 },
+      "provenance": {
+        "notes": "GEMV_N_COLS=8 twin of tail_n5: q8 unpack + zero_points + bias epilogue with only one live column in the second group."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [1, 16],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [5, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q8_zero_bias_gemv_m1_tail_n5_input_bT" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [5, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [5, 1],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [1, 5], "tolerance": 0.00001 } },
+      "attrs": { "K": 16, "N": 5, "bits": 8, "block_size": 16 }
+    },
+    {
+      "name": "q8_zero_bias_naive_fallback_tailK_m3_n6",
+      "provenance": {
+        "notes": "K=17,blockSize=16 gives a tail block; N=6 is below the register-tiled floor and the aligned tile paths are ineligible, while M=3 bypasses GEMV. With zero_points+bias present this routes to the priority-0 zero_bias scalar fallback and covers q8 unpack, per-block zero points, bias, and a partial K block."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [3, 17],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.75 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [6, 2, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q8_zero_bias_gemv_m1_tail_n5_input_bT" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [6, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [6, 2],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 6], "tolerance": 0.001 } },
+      "attrs": { "K": 17, "N": 6, "bits": 8, "block_size": 16 }
+    },
+    {
+      "name": "q4_prefill_tiled_reg_tailk_m32_k33_n4096",
+      "provenance": {
+        "notes": "Compact tail-block lock for the register-tiled prefill path used by the realistic K=2561 benchmark; the final 31 padded weights must not read past A."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 33],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096, 2, 16],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/q4_prefill_tiled_reg_tailk_m32_k33_n4096_input_bT" }
+          }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [4096, 2],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 4096], "tolerance": 0.001 } },
+      "attrs": { "K": 33, "N": 4096, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_prefill_tiled_reg_aligned_m32_k32_n4096_at_workgroup_floor",
+      "provenance": {
+        "source": "authored for route coverage",
+        "notes": "Aligned-K twin of the tail-K case at the same M and N. A tail-K shape skips the register tile's workgroup floor entirely, so no fixture exercised that floor; this one sits exactly on it at 64 workgroups."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096, 1, 16],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/q4_prefill_tiled_reg_tailk_m32_k33_n4096_input_bT" }
+          }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [4096, 1],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 4096], "tolerance": 0.001 } },
+      "attrs": { "K": 32, "N": 4096, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_prefill_tiled_reg_tailk_m32_k33_n512_below_wide_n_floor",
+      "provenance": {
+        "notes": "Tail-block coverage below the ordinary wide register-tile N floor. K=33 exercises the padded final block, while N=512 keeps the case inside the relaxed band that the larger N=4096 case cannot distinguish."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 33],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [512, 2, 16],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/q4_prefill_tiled_reg_tailk_m32_k33_n4096_input_bT" }
+          }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [512, 2],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 512], "tolerance": 0.001 } },
+      "attrs": { "K": 33, "N": 512, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_f16_prefill_reg_tailk_m32_k33_n4096",
+      "provenance": {
+        "notes": "f16 activations through the register-tiled general (non-block-aligned) loader and the M=32 portable geometry, plus the padded tail K."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [32, 33],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096, 2, 16],
+          "data": {
+            "kind": "cycle",
+            "values": { "$ref": "#/fixtureArrays/q4_prefill_tiled_reg_tailk_m32_k33_n4096_input_bT" }
+          }
+        },
+        "scalesT": {
+          "dtype": "float16",
+          "shape": [4096, 2],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [32, 4096], "tolerance": 0.0001, "relTolerance": 0.001 } },
+      "attrs": { "K": 33, "N": 4096, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q8_no_zero_prefill_odd_n_fallback",
+      "provenance": {
+        "notes": "M>1 q8 prefill with N=17 and no zero_points/bias. Odd N excludes subgroup-matrix execution; the portable tile4x4 tail guards handle the final output column used by the odd-column benchmark guardrail."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [8, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [17, 2, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [17, 2],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 17], "tolerance": 0.001 } },
+      "attrs": { "K": 64, "N": 17, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q4_zero_only_prefill_sgmat_m32_storage_floor",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [7.0, 8.0, 6.0, 9.0] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.0001 } },
+      "attrs": { "K": 32, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_bias_only_prefill_sgmat_m32_storage_floor",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 1, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.0001 } },
+      "attrs": { "K": 32, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_only_prefill_sgmat_large",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0, 144.0, 126.0] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.001 } },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "attrs": { "K": 128, "N": 64, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_bias_only_prefill_sgmat_large",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.001 } },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "attrs": { "K": 128, "N": 64, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q4_zero_only_prefill_reg_m512_n1024",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [512, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024, 4, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": [6.0, 7.0, 8.0, 9.0] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 1024], "tolerance": 0.0001 } },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "attrs": { "K": 128, "N": 1024, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_bias_only_prefill_reg_m512_n1024",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [512, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024, 4, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.1 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 1024], "tolerance": 0.0001 } },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "attrs": { "K": 128, "N": 1024, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_only_prefill_reg_m512_n1024",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [512, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0, 144.0, 126.0] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 1024], "tolerance": 0.001 } },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "attrs": { "K": 128, "N": 1024, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_bias_only_prefill_reg_m512_n1024",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [512, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 1024], "tolerance": 0.001 } },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "attrs": { "K": 128, "N": 1024, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q4_zero_only_prefill_tile4x4_small_m8",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [6.0, 7.0, 8.0, 9.0] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 64], "tolerance": 0.0001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_bias_only_prefill_tile4x4_small_m8",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.1 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 64], "tolerance": 0.0001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_only_prefill_tile4x4_small_m8",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0, 144.0, 126.0] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_bias_only_prefill_tile4x4_small_m8",
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 32],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/mixed_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q4_zero_only_prefill_tile4x4_large",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [6.0, 7.0, 8.0, 9.0] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.0001 } },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "attrs": { "K": 128, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_bias_only_prefill_tile4x4_large",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q4_weight_cycle_b_t" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/quant_scale_cycle_t" } }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.1 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.0001 } },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/matmul_nbits.cc",
+        "test": "independent has_zero_points/has_bias epilogues",
+        "notes": "zero_points-only and bias-only inputs must stay on the fast kernels instead of falling to the naive per-element path."
+      },
+      "attrs": { "K": 128, "N": 64, "bits": 4, "block_size": 32 }
+    },
+    {
+      "name": "q4_dp4a_prefill_al4_m64_k256_n64_b32",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [64, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 8, 16],
+          "data": {
+            "kind": "cycle",
+            "values": [37, 201, 116, 88, 250, 13, 164, 95, 142, 7, 233, 60, 178, 121, 44, 209, 86, 155]
+          }
+        },
+        "scalesT": { "dtype": "float32", "shape": [64, 8], "data": { "kind": "linspace", "start": 0.01, "end": 0.05 } }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [64, 64], "tolerance": 0.03 } },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul_nbits.cc",
+        "test": "accuracy_level=4 int8 activation quantization (DP4AMatMulQuantize + DP4AMatMulNBits)",
+        "notes": "accuracy_level 4 permits int8 activation math: A is quantized per 128-element block (scale = max_abs / 127, values pack4x8snorm-rounded) and multiplied against the 4-bit weights with packed integer dots. The reference stays float, so tolerances absorb the permitted quantization error."
+      },
+      "attrs": { "K": 256, "N": 64, "bits": 4, "block_size": 32, "accuracy_level": 4 }
+    },
+    {
+      "name": "q4_dp4a_prefill_al4_m48_k128_n16_b64",
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [48, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.027 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [16, 2, 32],
+          "data": { "kind": "cycle", "values": [37, 201, 116, 88, 250, 13, 164, 95, 142, 7, 233] }
+        },
+        "scalesT": { "dtype": "float32", "shape": [16, 2], "data": { "kind": "linspace", "start": 0.02, "end": 0.06 } }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [48, 16], "tolerance": 0.03 } },
+      "provenance": {
+        "source": "onnxruntime/contrib_ops/webgpu/quantization/dp4a_matmul_nbits.cc",
+        "test": "accuracy_level=4 int8 activation quantization (DP4AMatMulQuantize + DP4AMatMulNBits)",
+        "notes": "accuracy_level 4 permits int8 activation math: A is quantized per 128-element block (scale = max_abs / 127, values pack4x8snorm-rounded) and multiplied against the 4-bit weights with packed integer dots. The reference stays float, so tolerances absorb the permitted quantization error."
+      },
+      "attrs": { "K": 128, "N": 16, "bits": 4, "block_size": 64, "accuracy_level": 4 }
+    },
+    {
+      "name": "q2_zero_bias_prefill_tile4x4_small_m8",
+      "provenance": {
+        "notes": "2-bit weights pack four codes per stored byte, low codes first along K, with a default zero point of 2. Small-M prefill on the 4x4 output tile."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [8, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195, 84, 237] }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [1.0, 2.0, 3.0, 0.0, 2.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.1 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 2, "block_size": 32 }
+    },
+    {
+      "name": "q2_no_zero_prefill_sgmat_large",
+      "provenance": {
+        "notes": "2-bit weights pack four codes per stored byte, low codes first along K, with a default zero point of 2. Prefill wide enough for the subgroup-matrix tile."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195, 84, 237] }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 2, "block_size": 32 }
+    },
+    {
+      "name": "q2_zero_bias_prefill_sgmat_m32_storage_floor",
+      "provenance": {
+        "notes": "2-bit weights pack four codes per stored byte, low codes first along K, with a default zero point of 2. Subgroup-matrix tile at the M=32 shared-storage floor."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [32, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 1, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195, 84, 237] }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "cycle", "values": [1.0, 2.0, 3.0, 0.0, 2.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.1 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 32, "N": 64, "bits": 2, "block_size": 32 }
+    },
+    {
+      "name": "q2_no_zero_prefill_reg_m512_n1024",
+      "provenance": {
+        "notes": "2-bit weights pack four codes per stored byte, low codes first along K, with a default zero point of 2. Register-tile prefill, whose loader unrolls by K and so needs its own unpack."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [512, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024, 4, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195, 84, 237] }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 1024], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 1024, "bits": 2, "block_size": 32 }
+    },
+    {
+      "name": "q2_zero_bias_prefill_reg_m512_n1024",
+      "provenance": {
+        "notes": "2-bit weights pack four codes per stored byte, low codes first along K, with a default zero point of 2. Register-tile prefill with zero points and bias."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [512, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024, 4, 8],
+          "data": { "kind": "cycle", "values": [27, 180, 75, 226, 33, 150, 201, 108, 57, 246, 129, 66, 195, 84, 237] }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [1024, 4],
+          "data": { "kind": "cycle", "values": [1.0, 2.0, 3.0, 0.0, 2.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.1 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [512, 1024], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 1024, "bits": 2, "block_size": 32 }
+    },
+    {
+      "name": "q8_rank3_prefill_tile4x4",
+      "provenance": { "notes": "Rank-3 A folding to eight rows, wide enough for the prefill output tile." },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [2, 4, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.35 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [64, 4, 32],
+          "data": { "kind": "cycle", "values": [17, 200, 91, 45, 233, 128, 7, 176, 250, 33, 142, 99, 210, 64, 188] }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075, 0.07] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [64, 4],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0, 144.0] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.1 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [2, 4, 64], "tolerance": 0.001 } },
+      "attrs": { "K": 128, "N": 64, "bits": 8, "block_size": 32 }
+    },
+    {
+      "name": "q8_zero_only_naive_fallback_tailk_m3_n6",
+      "provenance": {
+        "notes": "K=17 leaves a partial final block and M=3 bypasses GEMV. With explicit zero points but no bias, the aligned prefill paths are ineligible and the zero-only scalar fallback handles the tail block."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [3, 17],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.75 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [6, 2, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q8_zero_bias_gemv_m1_tail_n5_input_bT" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [6, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075] }
+        },
+        "zeroPointsT": {
+          "dtype": "float32",
+          "shape": [6, 2],
+          "data": { "kind": "cycle", "values": [120.0, 132.0, 128.0, 140.0, 135.0, 118.0] }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 6], "tolerance": 0.001 } },
+      "attrs": { "K": 17, "N": 6, "bits": 8, "block_size": 16 }
+    },
+    {
+      "name": "q8_bias_only_naive_fallback_tailk_m3_n6",
+      "provenance": {
+        "notes": "K=17 leaves a partial final block and M=3 bypasses GEMV. With bias and the schema-default zero point, the aligned prefill paths are ineligible and the bias-only scalar fallback handles the tail block."
+      },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [3, 17],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.75 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [6, 2, 16],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/q8_zero_bias_gemv_m1_tail_n5_input_bT" } }
+        },
+        "scalesT": {
+          "dtype": "float32",
+          "shape": [6, 2],
+          "data": { "kind": "cycle", "values": [0.04, 0.055, 0.05, 0.065, 0.06, 0.075] }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.2 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 6], "tolerance": 0.001 } },
+      "attrs": { "K": 17, "N": 6, "bits": 8, "block_size": 16 }
+    }
+  ]
+}