sync 91d990483a17

Browse files

Files changed (10) hide show

README.md +24 -11
build/webgpu/bench.json +282 -1
build/webgpu/cast-scalar-x4.wgsl.jinja +3 -16
build/webgpu/manifest.json +321 -227
build/webgpu/matmul-bnb4-gemv.wgsl.jinja +1 -5
build/webgpu/matmul-bnb4-sgmat.wgsl.jinja +84 -19
build/webgpu/matmul-bnb4-tiled.wgsl.jinja +7 -10
build/webgpu/matmul-bnb4.wgsl.jinja +3 -6
build/webgpu/metadata.json +23 -11
build/webgpu/test.json +1444 -2

README.md CHANGED Viewed

@@ -18,17 +18,17 @@ See the [ONNX Runtime `MatMulBnb4` contrib-operator spec](https://github.com/mic
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `A` | `aT` | `T1` | `2` | — | Float input matrix of shape `(M, K)`, not quantized. | required |
-| `B` | `bT` | `T2` | `1` | — | The `[N, K]` weight, flattened and quantized to 4 bits, stored as `(N * K + 1) / 2` bytes; the ONNX type is uint8 (this WebGPU implementation reads one widened u32 per stored byte). | required |
-| `absmax` | `absmaxT` | `T1` | `1` | — | Per-block absolute-maximum dequantization scales of shape `((N * K + block_size - 1) / block_size)`, same dtype as A. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `yT` | `T1` | `2` | `[A[0], N]` | Result of `A` multiplied by the dequantized, transposed weight matrix, with shape `(M, N)` and the same dtype as `A`. | required |
 ## Attributes
@@ -36,12 +36,12 @@ Attributes and default values (overridable per request):
 | Attribute | Default | Description |
 | --- | --- | --- |
-| `training_mode` | `0` | Whether training outputs are requested. This inference-only implementation supports the standard default value 0. |
-| `transB` | `1` | Whether the quantized weight is stored transposed. This implementation supports the standard default value 1. |
 | `K` | — | Input feature count (the shared dimension). |
 | `N` | — | Output feature count. |
 | `block_size` | — | Number of weights sharing one absmax scale; a power of two, at least 16. |
 | `quant_type` | — | Codebook selector: 0 = FP4, 1 = NF4. |
 ## Type constraints
@@ -50,13 +50,22 @@ Attributes and default values (overridable per request):
 | `T1` | `float32`, `float16` |
 | `T2` | `uint8` |
 ## 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
@@ -68,10 +77,14 @@ Some implementation variants require `subgroup-matrix` and `subgroups`. These ar
 ## 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.

 ## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `aT` | `A` | `T1` | `2` | — | Float input matrix of shape `(M, K)`, not quantized. | required |
+| `bT` | `B` | `T2` | `1` | — | The `[N, K]` weight, flattened and quantized to 4 bits, stored as `(N * K + 1) / 2` bytes; the ONNX type is uint8 (this WebGPU implementation reads one widened u32 per stored byte). | required |
+| `absmaxT` | `absmax` | `T1` | `1` | — | Per-block absolute-maximum dequantization scales of shape `((N * K + block_size - 1) / block_size)`, same dtype as A. | required |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `yT` | `Y` | `T1` | `2` | `[aT[0], N]` | Result of `A` multiplied by the dequantized, transposed weight matrix, with shape `(M, N)` and the same dtype as `A`. | required |
 ## Attributes
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `K` | — | Input feature count (the shared dimension). |
 | `N` | — | Output feature count. |
 | `block_size` | — | Number of weights sharing one absmax scale; a power of two, at least 16. |
 | `quant_type` | — | Codebook selector: 0 = FP4, 1 = NF4. |
+| `training_mode` | `0` | Whether training outputs are requested. This inference-only implementation supports the standard default value 0. |
+| `transB` | `1` | Whether the quantized weight is stored transposed. This implementation supports the standard default value 1. |
 ## Type constraints
 | `T1` | `float32`, `float16` |
 | `T2` | `uint8` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `sgmat_hybrid_rows` — Combines full and partial row tiles in one matrix dispatch. Complete tiles read activations directly; partial tiles stage bounded rows in f32 and reuse weight staging for output. Float16 requests widen before the multiply and narrow once afterward.
+- `sgmat_hybrid_rows_widened` — Combines full and partial row tiles in one matrix dispatch. Complete tiles read activations directly; partial tiles stage bounded rows in f32 and reuse weight staging for output. Float16 requests widen before the multiply and narrow once afterward.
+- `sgmat_staged_rows` — Stages partial activation rows in f32, dequantizes the shared weight tile, and reuses dead weight staging for masked output stores.
+- `sgmat_widened` — Runs the subgroup-matrix tier for a float16 request by widening the activations to float32 either side of the multiply. The matrix units this operator uses accumulate in float32 from float32 operands; the device's float16 configuration returns a float16 result, so feeding them float16 directly would drop the accumulator's precision.
 ## 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, per-variant templates, 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
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.MatMulBnb4",
   "cases": [
     {
       "name": "bnb4-nf4-decode-m1-k2048-n2048-b64",
@@ -119,6 +118,288 @@
       "outputs": { "yT": { "shape": [512, 4096], "dtype": "float16" } },
       "bench": { "metrics": [{ "type": "gflops", "value": "2 * dim(shapes.aT, 0) * args.K * args.N" }] },
       "tunableSpace": { "PORTABLE_TILE_K": [16, 32, 64] }
     }
   ]
 }

 {
   "cases": [
     {
       "name": "bnb4-nf4-decode-m1-k2048-n2048-b64",
       "outputs": { "yT": { "shape": [512, 4096], "dtype": "float16" } },
       "bench": { "metrics": [{ "type": "gflops", "value": "2 * dim(shapes.aT, 0) * args.K * args.N" }] },
       "tunableSpace": { "PORTABLE_TILE_K": [16, 32, 64] }
+    },
+    {
+      "name": "bnb4-nf4-speculative-m15-k4096-n4096-b64-band",
+      "provenance": {
+        "notes": "Fifteen rows represent a speculative decoder's verification width and exercise the tiled route immediately below M=16. `nf4-tiled-speculative-band-m16-control` brackets that boundary."
+      },
+      "preset": "model",
+      "vars": { "M": 15, "K": 4096, "N": 4096, "blockSize": 64 },
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [15, 4096], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": {
+          "shape": [262144],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 427,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [15, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * dim(shapes.aT, 0) * args.K * args.N" }] }
+    },
+    {
+      "name": "bnb4-nf4-speculative-m16-k4096-n4096-b64-tile-floor-control",
+      "preset": "model",
+      "vars": { "M": 16, "K": 4096, "N": 4096, "blockSize": 64 },
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [16, 4096], "dtype": "float32", "dist": "normal", "seed": 428, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": {
+          "shape": [262144],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 429,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [16, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * dim(shapes.aT, 0) * args.K * args.N" }] }
+    },
+    {
+      "name": "bnb4-nf4-prefill-m64-k4096-n4096-b64-control",
+      "preset": "model",
+      "vars": { "M": 64, "K": 4096, "N": 4096, "blockSize": 64 },
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [64, 4096], "dtype": "float32", "dist": "normal", "seed": 430, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": {
+          "shape": [262144],
+          "dtype": "float32",
+          "dist": "uniform",
+          "seed": 431,
+          "offset": 0.04,
+          "scale": 0.01,
+          "signed": false
+        }
+      },
+      "outputs": { "yT": { "shape": [64, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * dim(shapes.aT, 0) * args.K * args.N" }] }
+    },
+    {
+      "name": "bnb4-probe-float32-q1-m8-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [8, 4096], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float32", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [8, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 268435456 }] }
+    },
+    {
+      "name": "bnb4-probe-float32-q1-m32-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [32, 4096], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float32", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [32, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 1073741824 }] }
+    },
+    {
+      "name": "bnb4-probe-float32-q1-m33-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [33, 4096], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float32", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [33, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 1107296256 }] }
+    },
+    {
+      "name": "bnb4-probe-float32-q1-m65-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [65, 4096], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float32", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [65, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 2181038080 }] }
+    },
+    {
+      "name": "bnb4-probe-float16-q1-m15-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [15, 4096], "dtype": "float16", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float16", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [15, 4096], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 503316480 }] }
+    },
+    {
+      "name": "bnb4-probe-float16-q1-m16-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [16, 4096], "dtype": "float16", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float16", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [16, 4096], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 536870912 }] }
+    },
+    {
+      "name": "bnb4-probe-float16-q1-m33-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [33, 4096], "dtype": "float16", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float16", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [33, 4096], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 1107296256 }] }
+    },
+    {
+      "name": "bnb4-probe-float32-q0-m16-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 0 },
+      "inputs": {
+        "aT": { "shape": [16, 4096], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float32", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [16, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 536870912 }] }
+    },
+    {
+      "name": "bnb4-probe-float32-q1-m15-k1024-n128-b64",
+      "preset": "stress",
+      "attrs": { "K": 1024, "N": 128, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [15, 1024], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [65536], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [2048], "dtype": "float32", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [15, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 3932160 }] }
+    },
+    {
+      "name": "bnb4-probe-float16-q1-m16-k1024-n128-b64",
+      "preset": "stress",
+      "attrs": { "K": 1024, "N": 128, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [16, 1024], "dtype": "float16", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [65536], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [2048], "dtype": "float16", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [16, 128], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 4194304 }] }
+    },
+    {
+      "name": "bnb4-probe-float32-q1-m15-k1023-n257-b64",
+      "preset": "stress",
+      "attrs": { "K": 1023, "N": 257, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [15, 1023], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [131456], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [4108], "dtype": "float32", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [15, 257], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 7887330 }] }
+    },
+    {
+      "name": "bnb4-probe-float16-q1-m16-k1023-n257-b64",
+      "preset": "stress",
+      "attrs": { "K": 1023, "N": 257, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [16, 1023], "dtype": "float16", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [131456], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [4108], "dtype": "float16", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [16, 257], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 8413152 }] }
+    },
+    {
+      "name": "bnb4-probe-float32-q1-m3-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [3, 4096], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float32", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [3, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 100663296 }] }
+    },
+    {
+      "name": "bnb4-probe-float32-q1-m129-k1024-n1024-b64",
+      "preset": "stress",
+      "attrs": { "K": 1024, "N": 1024, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [129, 1024], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [524288], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [16384], "dtype": "float32", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [129, 1024], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 270532608 }] }
+    },
+    {
+      "name": "bnb4-probe-float16-q1-m191-k1024-n1024-b64",
+      "preset": "stress",
+      "attrs": { "K": 1024, "N": 1024, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [191, 1024], "dtype": "float16", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [524288], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [16384], "dtype": "float16", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [191, 1024], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 400556032 }] }
+    },
+    {
+      "name": "bnb4-probe-float16-q1-m255-k1024-n1024-b64",
+      "preset": "stress",
+      "attrs": { "K": 1024, "N": 1024, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [255, 1024], "dtype": "float16", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [524288], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [16384], "dtype": "float16", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [255, 1024], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 534773760 }] }
+    },
+    {
+      "name": "bnb4-probe-float32-q1-m255-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [255, 4096], "dtype": "float32", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float32", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [255, 4096], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 8556380160 }] }
+    },
+    {
+      "name": "bnb4-probe-float16-q1-m511-k4096-n4096-b64",
+      "preset": "stress",
+      "attrs": { "K": 4096, "N": 4096, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": { "shape": [511, 4096], "dtype": "float16", "dist": "normal", "seed": 426, "scale": 0.2 },
+        "bT": { "shape": [8388608], "dtype": "uint8", "dist": "linearMod", "mod": 256, "step": 7 },
+        "absmaxT": { "shape": [262144], "dtype": "float16", "dist": "uniform", "seed": 427, "min": 0.04, "max": 0.05 }
+      },
+      "outputs": { "yT": { "shape": [511, 4096], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": 17146314752 }] }
     }
   ]
 }

build/webgpu/cast-scalar-x4.wgsl.jinja CHANGED Viewed

@@ -1,27 +1,14 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-{% if wrapNarrowInt %}
-// ONNX float-to-int8/uint8 narrowing wraps modulo 256 rather than saturating.
-// Preserve the low byte after truncation; the signed path then sign-extends it.
-fn cast_wrap(v: f32) -> {{ outScalar }} {
-  let low = i32(v) & 0xFF;
-  return {{ outScalar }}({% if wrapSigned %}select(low, low - 256, low > 127){% else %}low{% endif %});
-}
-{% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
   // Four scalar conversions per invocation retain vector-like dispatch density even when the
   // logical element count is not vec4 aligned. For a vec4 bulk + scalar tail plan, only lane zero
   // of the tail dispatch runs and starts at the first element not covered by the packed pass.
-  let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   let base = invocation * 4u;
 {% for lane in range(4) %}
   if (base + {{ lane }}u < params.count) {
-    y[base + {{ lane }}u] = {% if wrapNarrowInt %}cast_wrap(f32(x[base + {{ lane }}u])){% else %}{{ outScalar }}(x[base + {{ lane }}u]){% endif %};
   }
 {% endfor %}
 }

 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // Four scalar conversions per invocation retain vector-like dispatch density even when the
   // logical element count is not vec4 aligned. For a vec4 bulk + scalar tail plan, only lane zero
   // of the tail dispatch runs and starts at the first element not covered by the packed pass.
+  let invocation = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   let base = invocation * 4u;
 {% for lane in range(4) %}
   if (base + {{ lane }}u < params.count) {
+    y[base + {{ lane }}u] = {{ outScalar }}(x[base + {{ lane }}u]);
   }
 {% endfor %}
 }

build/webgpu/manifest.json CHANGED Viewed

@@ -2,32 +2,20 @@
   "domain": "com.microsoft",
   "name": "MatMulBnb4",
   "sinceVersion": 1,
-  "description": "Computes `A @ dequant(B)^T` where `B` uses bitsandbytes 4-bit quantization: `quant_type = 0` selects FP4 and `quant_type = 1` selects NF4. Supports rank-2 float16/float32 `A`, `transB = 1`, and `training_mode = 0`; rank-1 and rank-3-or-higher `A`, bfloat16, `transB = 0`, and training are not implemented. `B` is the flattened `[N, K]` weight, two codes per byte with the even flat index in the high nibble. Each code indexes a fixed 16-entry codebook, and the value is `codebook[code] * absmax[flat_index / block_size]`.",
-  "inputs": [
-    { "role": "A", "dtype": "T1", "rank": 2, "description": "Float input matrix of shape `(M, K)`, not quantized." },
-    {
-      "role": "B",
-      "dtype": "T2",
-      "rank": 1,
-      "description": "The `[N, K]` weight, flattened and quantized to 4 bits, stored as `(N * K + 1) / 2` bytes; the ONNX type is uint8 (this WebGPU implementation reads one widened u32 per stored byte)."
-    },
-    {
-      "role": "absmax",
-      "dtype": "T1",
-      "rank": 1,
-      "description": "Per-block absolute-maximum dequantization scales of shape `((N * K + block_size - 1) / block_size)`, same dtype as A."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "Y",
-      "dtype": "T1",
-      "rank": 2,
-      "shape": "[dim(shapes.A, 0), attrs.N]",
-      "description": "Result of `A` multiplied by the dequantized, transposed weight matrix, with shape `(M, N)` and the same dtype as `A`."
-    }
-  ],
-  "attributes": { "training_mode": 0, "transB": 1 },
   "attributeConstraints": {
     "K": { "required": true },
     "N": { "required": true },
@@ -36,158 +24,16 @@
     "training_mode": { "values": [0] },
     "transB": { "values": [1] }
   },
-  "attributeDescriptions": {
-    "K": "Input feature count (the shared dimension).",
-    "N": "Output feature count.",
-    "block_size": "Number of weights sharing one absmax scale; a power of two, at least 16.",
-    "quant_type": "Codebook selector: 0 = FP4, 1 = NF4.",
-    "training_mode": "Whether training outputs are requested. This inference-only implementation supports the standard default value 0.",
-    "transB": "Whether the quantized weight is stored transposed. This implementation supports the standard default value 1."
-  },
   "typeConstraints": { "T1": ["float32", "float16"], "T2": ["uint8"] },
-  "args": {
-    "aT": { "kind": "tensor", "semantic": "A", "role": "input" },
-    "bT": { "kind": "tensor", "semantic": "B", "role": "input" },
-    "absmaxT": { "kind": "tensor", "semantic": "absmax", "role": "input" },
-    "yT": { "kind": "tensor", "semantic": "Y", "role": "output" }
-  },
   "tunables": {
-    "WORKGROUP_SIZE": 64,
-    "TILE_MIN_M": 16,
-    "PORTABLE_TILE_K": 16,
-    "SGMAT_TILE_ROWS": 64,
-    "SGMAT_TALL_TILE_ROWS": 128,
-    "SGMAT_TALL_MIN_M": 128,
-    "SGMAT_TILE_COLS": 64,
-    "SGMAT_TILE_K": 32
-  },
-  "bindingSets": {
-    "main": [
-      {
-        "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": "u32" },
-      {
-        "name": "absmax",
-        "arg": "absmaxT",
-        "semantic": "absmax",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$absmaxScalar",
-        "length": "$ABSMAX_LEN"
-      },
-      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "rows", "type": "u32", "value": "dim(shapes.A, 0)" },
-            { "name": "K", "type": "u32", "value": "attrs.K" },
-            { "name": "N", "type": "u32", "value": "attrs.N" },
-            { "name": "blockSize", "type": "u32", "value": "attrs.block_size" }
-          ]
-        }
-      }
-    ],
-    "gemv": [
-      {
-        "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": "u32" },
-      {
-        "name": "absmax",
-        "arg": "absmaxT",
-        "semantic": "absmax",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$absmaxScalar",
-        "length": "$ABSMAX_LEN"
-      },
-      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" },
-      {
-        "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" }
-          ]
-        }
-      }
-    ],
-    "sgmat": [
-      {
-        "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": "u32" },
-      {
-        "name": "absmax",
-        "arg": "absmaxT",
-        "semantic": "absmax",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$absmaxScalar",
-        "length": "$ABSMAX_LEN"
-      },
-      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$aScalar" }
-    ],
-    "castA": [
-      {
-        "name": "x",
-        "arg": "aT",
-        "semantic": "A",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$srcScalar"
-      },
-      { "name": "y", "semantic": "aF32", "buffer": { "type": "storage" }, "elementType": "f32" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.A)" }] }
-      }
-    ],
-    "sgmatWiden": [
-      { "name": "a", "semantic": "aF32", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      { "name": "b", "arg": "bT", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "u32" },
-      {
-        "name": "absmax",
-        "arg": "absmaxT",
-        "semantic": "absmax",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$absmaxScalar",
-        "length": "$ABSMAX_LEN"
-      },
-      { "name": "y", "semantic": "yF32", "buffer": { "type": "storage" }, "elementType": "f32" }
-    ],
-    "castY": [
-      { "name": "x", "semantic": "yF32", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
-      { "name": "y", "arg": "yT", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$srcScalar" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [{ "name": "count", "type": "u32", "value": "dim(shapes.A, 0) * attrs.N" }]
-        }
-      }
-    ]
   },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
@@ -195,67 +41,293 @@
     "canPinSubgroupSize32": "device.features.has(\"subgroups\") and device.features.has(\"subgroup-size-control\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize <= 32 and device.adapterInfo.subgroupMaxSize >= 32",
     "pinSubgroupSize32": "canPinSubgroupSize32 and not wave32Adapter",
     "wave32Effective": "wave32Adapter or pinSubgroupSize32",
     "packedBytesExpected": "ceilDiv(attrs.N * attrs.K, 2)",
     "absmaxCountExpected": "ceilDiv(attrs.N * attrs.K, attrs.block_size)",
-    "aFloatOk": "(tensorDtypes.A == \"float32\" or tensorDtypes.A == \"float16\") and f16Ok(tensorDtypes.A)",
     "portableWorkgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
-    "commonShapeValid": "ranks.A == 2 and ranks.B == 1 and ranks.absmax == 1 and ranks.Y == 2 and aFloatOk and tensorDtypes.B == \"uint8\" and tensorDtypes.absmax == tensorDtypes.A and tensorDtypes.Y == tensorDtypes.A and attrs.K > 0 and attrs.N > 0 and attrs.block_size >= 16 and pow2ceil(attrs.block_size) == attrs.block_size and dim(shapes.A, 1) == attrs.K and dim(shapes.B, 0) == packedBytesExpected and dim(shapes.absmax, 0) == absmaxCountExpected and dim(shapes.Y, 0) == dim(shapes.A, 0) and dim(shapes.Y, 1) == attrs.N",
-    "gemvShapeValid": "commonShapeValid and dim(shapes.A, 0) == 1",
     "portableWorkgroupFits": "portableWorkgroupSize > 0 and portableWorkgroupSize * 16 <= device.limits.maxComputeWorkgroupStorageSize",
     "portableTileKValid": "tunables.PORTABLE_TILE_K >= 8 and tunables.PORTABLE_TILE_K % 8 == 0",
-    "tileEligible": "commonShapeValid and dim(shapes.A, 0) >= tunables.TILE_MIN_M and portableTileKValid",
-    "tileWorkgroupStorageBytes": "8 * 64 * tunables.PORTABLE_TILE_K",
-    "tileWorkgroupFits": "16 <= device.limits.maxComputeWorkgroupSizeX and 8 <= device.limits.maxComputeWorkgroupSizeY and 128 <= device.limits.maxComputeInvocationsPerWorkgroup and tileWorkgroupStorageBytes <= device.limits.maxComputeWorkgroupStorageSize and ceilDiv(attrs.N, 64) <= device.limits.maxComputeWorkgroupsPerDimension and ceilDiv(dim(shapes.A, 0), 64) <= device.limits.maxComputeWorkgroupsPerDimension",
     "sgmatMatrixSize": "8",
-    "sgmatTileRows": "tunables.SGMAT_TALL_TILE_ROWS if dim(shapes.A, 0) >= tunables.SGMAT_TALL_MIN_M else tunables.SGMAT_TILE_ROWS",
     "sgmatRowSubtiles": "4",
     "sgmatSubRows": "sgmatTileRows / sgmatRowSubtiles",
     "sgmatSubCols": "4 * sgmatMatrixSize",
     "sgmatLoadWidth": "sgmatMatrixSize",
     "sgmatColSubtiles": "tunables.SGMAT_TILE_COLS / sgmatSubCols",
     "sgmatNumSubgroups": "sgmatRowSubtiles * sgmatColSubtiles",
-    "sgmatSubgroupSize": "device.adapterInfo.subgroupMinSize if has(device.adapterInfo, \"subgroupMinSize\") else 1",
     "sgmatWorkgroupSize": "sgmatNumSubgroups * sgmatSubgroupSize",
     "sgmatBLoadsPerRow": "tunables.SGMAT_TILE_K / sgmatLoadWidth",
     "sgmatWorkgroupStorageBytes": "4 * tunables.SGMAT_TILE_COLS * tunables.SGMAT_TILE_K",
     "sgmatDispatchN": "ceilDiv(attrs.N, tunables.SGMAT_TILE_COLS)",
-    "sgmatDispatchM": "ceilDiv(dim(shapes.A, 0), sgmatTileRows)",
-    "sgmatWorkgroupFits": "sgmatWorkgroupSize <= deviceWorkgroupCap and sgmatWorkgroupStorageBytes <= device.limits.maxComputeWorkgroupStorageSize and sgmatDispatchN <= device.limits.maxComputeWorkgroupsPerDimension and sgmatDispatchM <= device.limits.maxComputeWorkgroupsPerDimension",
-    "sgmatWidenBytes": "numel(shapes.A) * 4",
-    "sgmatWidenOutBytes": "dim(shapes.A, 0) * attrs.N * 4",
-    "sgmatWidenFits": "sgmatWidenBytes <= device.limits.maxStorageBufferBindingSize and sgmatWidenBytes <= device.limits.maxBufferSize and sgmatWidenOutBytes <= device.limits.maxStorageBufferBindingSize and sgmatWidenOutBytes <= device.limits.maxBufferSize"
   },
-  "constants": { "quantType": "attrs.quant_type", "ABSMAX_LEN": "ceilDiv(attrs.N * attrs.K, attrs.block_size)" },
   "variants": [
     {
       "id": "gemv",
       "priority": 20,
       "when": ["gemvShapeValid", "portableWorkgroupFits"],
-      "constants": {
         "workgroupSize": "portableWorkgroupSize",
-        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "usesF16": "tensorDtypes.A == \"float16\"",
-        "absmaxScalar": "\"f16\" if tensorDtypes.absmax == \"float16\" else \"f32\""
       },
       "passes": [
         {
           "id": "main",
           "shader": "matmul-bnb4-gemv.wgsl.jinja",
-          "bindings": "gemv",
-          "dispatch": { "workgroups": "attrs.N" }
         }
       ]
     },
     {
       "id": "sgmat",
       "priority": 15,
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "limits": { "maxComputeWorkgroupStorageSize": 8192 },
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
-      "when": ["commonShapeValid", "tensorDtypes.A == \"float32\"", "dim(shapes.A, 0) >= sgmatTileRows", "dim(shapes.A, 0) % sgmatTileRows == 0", "attrs.N % tunables.SGMAT_TILE_COLS == 0", "attrs.K % tunables.SGMAT_TILE_K == 0", "attrs.K % attrs.block_size == 0", "attrs.block_size % sgmatLoadWidth == 0", "wave32Effective", "sgmatWorkgroupFits"],
-      "constants": {
         "K": "attrs.K",
         "N": "attrs.N",
         "blockSize": "attrs.block_size",
@@ -273,13 +345,13 @@
         "bLoadsPerRow": "sgmatBLoadsPerRow",
         "aScalar": "\"f32\"",
         "usesF16": false,
-        "absmaxScalar": "\"f16\" if tensorDtypes.absmax == \"float16\" else \"f32\""
       },
       "passes": [
         {
           "id": "main",
           "shader": "matmul-bnb4-sgmat.wgsl.jinja",
-          "bindings": "sgmat",
           "dispatch": { "x": "sgmatDispatchN", "y": "sgmatDispatchM" }
         }
       ]
@@ -287,13 +359,13 @@
     {
       "id": "sgmat_widened",
       "priority": 16,
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "limits": { "maxComputeWorkgroupStorageSize": 8192 },
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
-      "when": ["commonShapeValid", "tensorDtypes.A == \"float16\"", "dim(shapes.A, 0) >= sgmatTileRows", "dim(shapes.A, 0) % sgmatTileRows == 0", "attrs.N % tunables.SGMAT_TILE_COLS == 0", "attrs.K % tunables.SGMAT_TILE_K == 0", "attrs.K % attrs.block_size == 0", "attrs.block_size % sgmatLoadWidth == 0", "wave32Effective", "sgmatWorkgroupFits", "sgmatWidenFits"],
-      "constants": {
         "K": "attrs.K",
         "N": "attrs.N",
         "blockSize": "attrs.block_size",
@@ -311,61 +383,79 @@
         "bLoadsPerRow": "sgmatBLoadsPerRow",
         "aScalar": "\"f32\"",
         "usesF16": true,
-        "absmaxScalar": "\"f16\" if tensorDtypes.absmax == \"float16\" else \"f32\"",
         "srcScalar": "\"f16\"",
         "outScalar": "\"f32\"",
         "wrapNarrowInt": false,
         "wrapSigned": false
       },
       "passes": [
         {
           "id": "widen_a",
           "name": "MatMulBnb4.WidenActivations",
           "shader": "cast-scalar-x4.wgsl.jinja",
-          "bindings": "castA",
-          "dispatch": { "threads": "ceilDiv(numel(shapes.A), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         },
         {
           "id": "main",
           "name": "MatMulBnb4.SubgroupMatrixWidened",
           "shader": "matmul-bnb4-sgmat.wgsl.jinja",
-          "bindings": "sgmatWiden",
           "dispatch": { "x": "sgmatDispatchN", "y": "sgmatDispatchM" }
         },
         {
           "id": "narrow_y",
           "name": "MatMulBnb4.NarrowOutput",
           "shader": "cast-scalar-x4.wgsl.jinja",
-          "bindings": "castY",
-          "constants": { "outScalar": "\"f16\"" },
           "dispatch": {
-            "threads": "ceilDiv(dim(shapes.A, 0) * attrs.N, 4)",
-            "workgroupSize": "tunables.WORKGROUP_SIZE"
           }
         }
-      ],
-      "description": "Runs the subgroup-matrix tier for a float16 request by widening the activations to float32 either side of the multiply. The matrix units this operator uses accumulate in float32 from float32 operands; the device's float16 configuration returns a float16 result, so feeding them float16 directly would drop the accumulator's precision.",
-      "intermediates": [
-        { "id": "aF32", "dtype": "float32", "shape": "[numel(shapes.A)]" },
-        { "id": "yF32", "dtype": "float32", "shape": "[dim(shapes.A, 0) * attrs.N]" }
       ]
     },
     {
       "id": "tiled",
       "priority": 10,
       "when": ["tileEligible", "tileWorkgroupFits"],
-      "constants": {
         "tileK": "tunables.PORTABLE_TILE_K",
-        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "usesF16": "tensorDtypes.A == \"float16\"",
-        "absmaxScalar": "\"f16\" if tensorDtypes.absmax == \"float16\" else \"f32\""
       },
       "passes": [
         {
           "id": "main",
           "shader": "matmul-bnb4-tiled.wgsl.jinja",
-          "bindings": "main",
-          "dispatch": { "x": "ceilDiv(attrs.N, 64)", "y": "ceilDiv(dim(shapes.A, 0), 64)" }
         }
       ]
     },
@@ -373,18 +463,22 @@
       "id": "scalar",
       "priority": 0,
       "when": ["commonShapeValid", "portableWorkgroupSize > 0"],
-      "constants": {
         "workgroupSize": "portableWorkgroupSize",
-        "aScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "usesF16": "tensorDtypes.A == \"float16\"",
-        "absmaxScalar": "\"f16\" if tensorDtypes.absmax == \"float16\" else \"f32\""
       },
       "passes": [
         {
           "id": "main",
           "shader": "matmul-bnb4.wgsl.jinja",
-          "bindings": "main",
-          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "constants.workgroupSize" }
         }
       ]
     }

   "domain": "com.microsoft",
   "name": "MatMulBnb4",
   "sinceVersion": 1,
+  "inputs": {
+    "aT": { "onnx": "A", "dtype": "T1", "rank": 2 },
+    "bT": { "onnx": "B", "dtype": "T2", "rank": 1 },
+    "absmaxT": { "onnx": "absmax", "dtype": "T1", "rank": 1 }
+  },
+  "outputs": { "yT": { "onnx": "Y", "dtype": "T1", "rank": 2, "shape": "[dim(shapes.aT, 0), attrs.N]" } },
+  "attributes": {
+    "training_mode": { "default": 0 },
+    "transB": { "default": 1 },
+    "K": {},
+    "N": {},
+    "block_size": {},
+    "quant_type": {}
+  },
   "attributeConstraints": {
     "K": { "required": true },
     "N": { "required": true },
     "training_mode": { "values": [0] },
     "transB": { "values": [1] }
   },
   "typeConstraints": { "T1": ["float32", "float16"], "T2": ["uint8"] },
   "tunables": {
+    "WORKGROUP_SIZE": { "default": 64 },
+    "TILE_MIN_M": { "default": 4 },
+    "PORTABLE_TILE_K": { "default": 16 },
+    "SGMAT_TILE_ROWS": { "default": 64 },
+    "SGMAT_TALL_TILE_ROWS": { "default": 128 },
+    "SGMAT_TALL_MIN_M": { "default": 128 },
+    "SGMAT_TILE_COLS": { "default": 64 },
+    "SGMAT_TILE_K": { "default": 32 }
   },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "canPinSubgroupSize32": "device.features.has(\"subgroups\") and device.features.has(\"subgroup-size-control\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize <= 32 and device.adapterInfo.subgroupMaxSize >= 32",
     "pinSubgroupSize32": "canPinSubgroupSize32 and not wave32Adapter",
     "wave32Effective": "wave32Adapter or pinSubgroupSize32",
+    "portableTileRows": "8 * min(8, max(1, ceilDiv(dim(shapes.aT, 0), 8)))",
     "packedBytesExpected": "ceilDiv(attrs.N * attrs.K, 2)",
     "absmaxCountExpected": "ceilDiv(attrs.N * attrs.K, attrs.block_size)",
+    "aFloatOk": "(tensorDtypes.aT == \"float32\" or tensorDtypes.aT == \"float16\") and f16Ok(tensorDtypes.aT)",
     "portableWorkgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
+    "commonShapeValid": "ranks.aT == 2 and ranks.bT == 1 and ranks.absmaxT == 1 and ranks.yT == 2 and aFloatOk and tensorDtypes.bT == \"uint8\" and tensorDtypes.absmaxT == tensorDtypes.aT and tensorDtypes.yT == tensorDtypes.aT and attrs.K > 0 and attrs.N > 0 and attrs.block_size >= 16 and pow2ceil(attrs.block_size) == attrs.block_size and dim(shapes.aT, 1) == attrs.K and dim(shapes.bT, 0) == packedBytesExpected and dim(shapes.absmaxT, 0) == absmaxCountExpected and dim(shapes.yT, 0) == dim(shapes.aT, 0) and dim(shapes.yT, 1) == attrs.N",
+    "gemvShapeValid": "commonShapeValid and dim(shapes.aT, 0) == 1",
     "portableWorkgroupFits": "portableWorkgroupSize > 0 and portableWorkgroupSize * 16 <= device.limits.maxComputeWorkgroupStorageSize",
     "portableTileKValid": "tunables.PORTABLE_TILE_K >= 8 and tunables.PORTABLE_TILE_K % 8 == 0",
+    "tileEligible": "commonShapeValid and dim(shapes.aT, 0) >= tunables.TILE_MIN_M and portableTileKValid",
+    "tileWorkgroupStorageBytes": "(portableTileRows * dtypeBytes(tensorDtypes.aT) + 64 * 4) * tunables.PORTABLE_TILE_K",
+    "tileWorkgroupFits": "16 <= device.limits.maxComputeWorkgroupSizeX and 8 <= device.limits.maxComputeWorkgroupSizeY and 128 <= device.limits.maxComputeInvocationsPerWorkgroup and tileWorkgroupStorageBytes <= device.limits.maxComputeWorkgroupStorageSize and ceilDiv(attrs.N, 64) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and ceilDiv(dim(shapes.aT, 0), portableTileRows) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "sgmatMatrixSize": "8",
+    "sgmatTileRows": "tunables.SGMAT_TALL_TILE_ROWS if dim(shapes.aT, 0) >= tunables.SGMAT_TALL_MIN_M else tunables.SGMAT_TILE_ROWS",
     "sgmatRowSubtiles": "4",
     "sgmatSubRows": "sgmatTileRows / sgmatRowSubtiles",
     "sgmatSubCols": "4 * sgmatMatrixSize",
     "sgmatLoadWidth": "sgmatMatrixSize",
     "sgmatColSubtiles": "tunables.SGMAT_TILE_COLS / sgmatSubCols",
     "sgmatNumSubgroups": "sgmatRowSubtiles * sgmatColSubtiles",
+    "sgmatSubgroupSize": "32",
     "sgmatWorkgroupSize": "sgmatNumSubgroups * sgmatSubgroupSize",
     "sgmatBLoadsPerRow": "tunables.SGMAT_TILE_K / sgmatLoadWidth",
     "sgmatWorkgroupStorageBytes": "4 * tunables.SGMAT_TILE_COLS * tunables.SGMAT_TILE_K",
     "sgmatDispatchN": "ceilDiv(attrs.N, tunables.SGMAT_TILE_COLS)",
+    "sgmatDispatchM": "ceilDiv(dim(shapes.aT, 0), sgmatTileRows)",
+    "sgmatWorkgroupFits": "sgmatWorkgroupSize <= deviceWorkgroupCap and sgmatWorkgroupStorageBytes <= device.limits.maxComputeWorkgroupStorageSize and sgmatDispatchN <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and sgmatDispatchM <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "sgmatWidenBytes": "numel(shapes.aT) * 4",
+    "sgmatWidenOutBytes": "dim(shapes.aT, 0) * attrs.N * 4",
+    "sgmatWidenFits": "sgmatWidenBytes <= device.limits.maxStorageBufferBindingSize and sgmatWidenBytes <= device.limits.maxBufferSize and sgmatWidenOutBytes <= device.limits.maxStorageBufferBindingSize and sgmatWidenOutBytes <= device.limits.maxBufferSize",
+    "quantType": "attrs.quant_type",
+    "ABSMAX_LEN": "ceilDiv(attrs.N * attrs.K, attrs.block_size)",
+    "stagedGeometryValid": "tunables.SGMAT_TILE_COLS >= sgmatSubCols and tunables.SGMAT_TILE_COLS % sgmatSubCols == 0 and tunables.SGMAT_TILE_K >= sgmatMatrixSize and tunables.SGMAT_TILE_K % sgmatMatrixSize == 0",
+    "stagedRowSubtiles": "max(1, min(sgmatRowSubtiles, pow2ceil(ceilDiv(dim(shapes.aT, 0), sgmatMatrixSize)), floor(deviceWorkgroupCap / (32 * max(1, sgmatColSubtiles))), floor(tunables.SGMAT_TILE_K / sgmatMatrixSize), floor((device.limits.maxComputeWorkgroupStorageSize - sgmatWorkgroupStorageBytes) / (4 * sgmatMatrixSize * max(1, tunables.SGMAT_TILE_K)))))",
+    "stagedTileRows": "stagedRowSubtiles * sgmatMatrixSize",
+    "stagedLoadWorkgroupTarget": "max(1, min(deviceWorkgroupCap, tunables.SGMAT_TILE_COLS * tunables.SGMAT_TILE_K / sgmatLoadWidth))",
+    "stagedColumnMatrices": "min(4, pow2ceil(max(1, ceilDiv(stagedRowSubtiles * tunables.SGMAT_TILE_COLS * 32, stagedLoadWorkgroupTarget * sgmatMatrixSize))))",
+    "stagedSubCols": "stagedColumnMatrices * sgmatMatrixSize",
+    "stagedWorkgroupSize": "stagedRowSubtiles * (tunables.SGMAT_TILE_COLS / stagedSubCols) * 32",
+    "stagedStorageBytes": "4 * (stagedTileRows + tunables.SGMAT_TILE_COLS) * tunables.SGMAT_TILE_K",
+    "stagedResourcesFit": "stagedWorkgroupSize <= deviceWorkgroupCap and stagedStorageBytes <= device.limits.maxComputeWorkgroupStorageSize and sgmatDispatchN <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and ceilDiv(dim(shapes.aT, 0), stagedTileRows) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "sgmatGeometryValid": "tunables.SGMAT_TILE_COLS >= sgmatSubCols and tunables.SGMAT_TILE_COLS % sgmatSubCols == 0 and tunables.SGMAT_TILE_K >= sgmatMatrixSize and tunables.SGMAT_TILE_K % sgmatMatrixSize == 0 and sgmatTileRows >= sgmatRowSubtiles * sgmatMatrixSize and sgmatTileRows % (sgmatRowSubtiles * sgmatMatrixSize) == 0",
+    "prefixTileRows": "tunables.SGMAT_TILE_ROWS",
+    "sgmatFullRows": "floor(dim(shapes.aT, 0) / max(1, prefixTileRows)) * prefixTileRows",
+    "sgmatTailRows": "dim(shapes.aT, 0) - sgmatFullRows",
+    "tailGeometryValid": "tunables.SGMAT_TILE_COLS >= sgmatSubCols and tunables.SGMAT_TILE_COLS % sgmatSubCols == 0 and tunables.SGMAT_TILE_K >= sgmatMatrixSize and tunables.SGMAT_TILE_K % sgmatMatrixSize == 0",
+    "tailRowSubtiles": "max(1, min(sgmatRowSubtiles, pow2ceil(ceilDiv(sgmatTailRows, sgmatMatrixSize)), floor(deviceWorkgroupCap / (32 * max(1, sgmatColSubtiles))), floor(tunables.SGMAT_TILE_K / sgmatMatrixSize), floor((device.limits.maxComputeWorkgroupStorageSize - sgmatWorkgroupStorageBytes) / (4 * sgmatMatrixSize * max(1, tunables.SGMAT_TILE_K)))))",
+    "prefixGeometryValid": "tunables.SGMAT_TILE_COLS >= sgmatSubCols and tunables.SGMAT_TILE_COLS % sgmatSubCols == 0 and tunables.SGMAT_TILE_K >= sgmatMatrixSize and tunables.SGMAT_TILE_K % sgmatMatrixSize == 0 and prefixTileRows >= sgmatRowSubtiles * sgmatMatrixSize and prefixTileRows % (sgmatRowSubtiles * sgmatMatrixSize) == 0",
+    "prefixResourcesFit": "sgmatWorkgroupFits and ceilDiv(sgmatFullRows, max(1, prefixTileRows)) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
+    "hybridTailRowSubtiles": "pow2ceil(tailRowSubtiles) / (2 if pow2ceil(tailRowSubtiles) > tailRowSubtiles else 1)",
+    "hybridTailTileRows": "hybridTailRowSubtiles * sgmatMatrixSize",
+    "hybridTailSubCols": "sgmatSubCols * hybridTailRowSubtiles / sgmatRowSubtiles",
+    "hybridStorageBytes": "4 * (tunables.SGMAT_TILE_COLS + hybridTailTileRows) * tunables.SGMAT_TILE_K",
+    "hybridDispatchRows": "sgmatFullRows / max(1, prefixTileRows) + ceilDiv(sgmatTailRows, hybridTailTileRows)",
+    "hybridResourcesFit": "prefixResourcesFit and hybridStorageBytes <= device.limits.maxComputeWorkgroupStorageSize and hybridDispatchRows <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)"
+  },
+  "bindings": {
+    "a": { "arg": "aT", "buffer": "read-only-storage", "elementType": "$aScalar" },
+    "b": { "arg": "bT", "buffer": "read-only-storage", "elementType": "u32" },
+    "absmax": {
+      "arg": "absmaxT",
+      "buffer": "read-only-storage",
+      "elementType": "$absmaxScalar",
+      "length": "$ABSMAX_LEN"
+    },
+    "y": { "arg": "yT", "buffer": "storage", "elementType": "$aScalar" },
+    "params": {
+      "buffer": "uniform",
+      "struct": [
+        { "name": "rows", "type": "u32", "value": "dim(shapes.aT, 0)" },
+        { "name": "K", "type": "u32", "value": "attrs.K" },
+        { "name": "N", "type": "u32", "value": "attrs.N" },
+        { "name": "blockSize", "type": "u32", "value": "attrs.block_size" }
+      ]
+    }
   },
   "variants": [
+    {
+      "id": "sgmat_hybrid_rows",
+      "priority": 21,
+      "when": ["prefixGeometryValid", "commonShapeValid", "tensorDtypes.aT == \"float32\"", "attrs.N % tunables.SGMAT_TILE_COLS == 0", "attrs.K % tunables.SGMAT_TILE_K == 0", "attrs.K % attrs.block_size == 0", "attrs.block_size % sgmatLoadWidth == 0", "wave32Effective", "prefixResourcesFit", "sgmatFullRows > 0", "sgmatTailRows > 0", "tailGeometryValid", "hybridResourcesFit"],
+      "requires": {
+        "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
+        "limits": { "maxComputeWorkgroupStorageSize": 8192 },
+        "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
+      },
+      "derive": {
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "blockSize": "attrs.block_size",
+        "tileRows": "prefixTileRows",
+        "tileCols": "tunables.SGMAT_TILE_COLS",
+        "tileK": "tunables.SGMAT_TILE_K",
+        "subRows": "prefixTileRows / sgmatRowSubtiles",
+        "subCols": "sgmatSubCols",
+        "matrixSize": "sgmatMatrixSize",
+        "rowMatrices": "prefixTileRows / (sgmatRowSubtiles * sgmatMatrixSize)",
+        "colMatrices": "sgmatSubCols / sgmatMatrixSize",
+        "loadWidth": "sgmatLoadWidth",
+        "rowSubtiles": "sgmatRowSubtiles",
+        "workgroupSize": "sgmatWorkgroupSize",
+        "bLoadsPerRow": "sgmatBLoadsPerRow",
+        "aScalar": "\"f32\"",
+        "usesF16": false,
+        "absmaxScalar": "\"f16\" if tensorDtypes.absmaxT == \"float16\" else \"f32\"",
+        "hybridRows": true,
+        "rowCount": "dim(shapes.aT, 0)",
+        "fullDispatchRows": "sgmatFullRows / prefixTileRows",
+        "rowOffset": "sgmatFullRows"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-bnb4-sgmat.wgsl.jinja",
+          "bindings": ["a", "b", "absmax", "y"],
+          "dispatch": { "x": "sgmatDispatchN", "y": "hybridDispatchRows" }
+        }
+      ]
+    },
+    {
+      "id": "sgmat_hybrid_rows_widened",
+      "priority": 21,
+      "when": ["prefixGeometryValid", "commonShapeValid", "tensorDtypes.aT == \"float16\"", "attrs.N % tunables.SGMAT_TILE_COLS == 0", "attrs.K % tunables.SGMAT_TILE_K == 0", "attrs.K % attrs.block_size == 0", "attrs.block_size % sgmatLoadWidth == 0", "wave32Effective", "prefixResourcesFit", "sgmatWidenFits", "sgmatFullRows > 0", "sgmatTailRows > 0", "tailGeometryValid", "hybridResourcesFit"],
+      "requires": {
+        "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
+        "limits": { "maxComputeWorkgroupStorageSize": 8192 },
+        "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
+      },
+      "derive": {
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "blockSize": "attrs.block_size",
+        "tileRows": "prefixTileRows",
+        "tileCols": "tunables.SGMAT_TILE_COLS",
+        "tileK": "tunables.SGMAT_TILE_K",
+        "subRows": "prefixTileRows / sgmatRowSubtiles",
+        "subCols": "sgmatSubCols",
+        "matrixSize": "sgmatMatrixSize",
+        "rowMatrices": "prefixTileRows / (sgmatRowSubtiles * sgmatMatrixSize)",
+        "colMatrices": "sgmatSubCols / sgmatMatrixSize",
+        "loadWidth": "sgmatLoadWidth",
+        "rowSubtiles": "sgmatRowSubtiles",
+        "workgroupSize": "sgmatWorkgroupSize",
+        "bLoadsPerRow": "sgmatBLoadsPerRow",
+        "aScalar": "\"f32\"",
+        "usesF16": true,
+        "absmaxScalar": "\"f16\" if tensorDtypes.absmaxT == \"float16\" else \"f32\"",
+        "srcScalar": "\"f16\"",
+        "outScalar": "\"f32\"",
+        "wrapNarrowInt": false,
+        "wrapSigned": false,
+        "hybridRows": true,
+        "rowCount": "dim(shapes.aT, 0)",
+        "fullDispatchRows": "sgmatFullRows / prefixTileRows",
+        "rowOffset": "sgmatFullRows"
+      },
+      "intermediates": [
+        { "id": "aF32", "dtype": "float32", "shape": "[numel(shapes.aT)]" },
+        { "id": "yF32", "dtype": "float32", "shape": "[numel(shapes.yT)]" }
+      ],
+      "passes": [
+        {
+          "id": "widen_a",
+          "name": "MatMulBnb4.WidenActivations",
+          "shader": "cast-scalar-x4.wgsl.jinja",
+          "bindings": [
+            { "arg": "aT", "name": "x", "elementType": "$srcScalar" },
+            { "scratch": "aF32", "name": "y", "elementType": "f32" },
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.aT)" }] }
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.aT), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.aT), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "main",
+          "name": "MatMulBnb4.SubgroupMatrixWidened",
+          "shader": "matmul-bnb4-sgmat.wgsl.jinja",
+          "bindings": [
+            { "scratch": "aF32", "name": "a", "buffer": "read-only-storage", "elementType": "f32" },
+            "b",
+            "absmax",
+            { "scratch": "yF32", "name": "y", "elementType": "f32" }
+          ],
+          "dispatch": { "x": "sgmatDispatchN", "y": "hybridDispatchRows" }
+        },
+        {
+          "id": "narrow_y",
+          "name": "MatMulBnb4.NarrowOutput",
+          "shader": "cast-scalar-x4.wgsl.jinja",
+          "derive": { "outScalar": "\"f16\"" },
+          "bindings": [
+            { "scratch": "yF32", "name": "x", "buffer": "read-only-storage", "elementType": "f32" },
+            { "arg": "yT", "name": "y", "elementType": "$srcScalar" },
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.yT)" }] }
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.yT), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.yT), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "sgmat_staged_rows",
+      "priority": 19,
+      "when": ["commonShapeValid", "dim(shapes.aT, 0) > 1", "dim(shapes.aT, 0) < sgmatTileRows or dim(shapes.aT, 0) % sgmatTileRows != 0", "stagedGeometryValid", "attrs.N % tunables.SGMAT_TILE_COLS == 0", "attrs.K % tunables.SGMAT_TILE_K == 0", "attrs.K % attrs.block_size == 0", "attrs.block_size % sgmatLoadWidth == 0", "wave32Effective", "stagedResourcesFit"],
+      "requires": {
+        "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
+        "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
+      },
+      "derive": {
+        "K": "attrs.K",
+        "N": "attrs.N",
+        "blockSize": "attrs.block_size",
+        "tileRows": "stagedTileRows",
+        "tileCols": "tunables.SGMAT_TILE_COLS",
+        "tileK": "tunables.SGMAT_TILE_K",
+        "subRows": "sgmatMatrixSize",
+        "subCols": "stagedSubCols",
+        "matrixSize": "sgmatMatrixSize",
+        "rowMatrices": 1,
+        "colMatrices": "stagedColumnMatrices",
+        "loadWidth": "sgmatLoadWidth",
+        "rowSubtiles": "stagedRowSubtiles",
+        "workgroupSize": "stagedWorkgroupSize",
+        "bLoadsPerRow": "sgmatBLoadsPerRow",
+        "aScalar": "\"f16\" if tensorDtypes.aT == \"float16\" else \"f32\"",
+        "usesF16": "tensorDtypes.aT == \"float16\"",
+        "absmaxScalar": "\"f16\" if tensorDtypes.absmaxT == \"float16\" else \"f32\"",
+        "stageActivations": true,
+        "rowCount": "dim(shapes.aT, 0)"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "shader": "matmul-bnb4-sgmat.wgsl.jinja",
+          "bindings": ["a", "b", "absmax", "y"],
+          "dispatch": { "x": "sgmatDispatchN", "y": "ceilDiv(dim(shapes.aT, 0), stagedTileRows)" }
+        }
+      ]
+    },
     {
       "id": "gemv",
       "priority": 20,
       "when": ["gemvShapeValid", "portableWorkgroupFits"],
+      "derive": {
         "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.aT == \"float16\" else \"f32\"",
+        "usesF16": "tensorDtypes.aT == \"float16\"",
+        "absmaxScalar": "\"f16\" if tensorDtypes.absmaxT == \"float16\" else \"f32\""
       },
       "passes": [
         {
           "id": "main",
           "shader": "matmul-bnb4-gemv.wgsl.jinja",
+          "bindings": [
+            "a",
+            "b",
+            "absmax",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "name": "K", "type": "u32", "value": "attrs.K" },
+                { "name": "N", "type": "u32", "value": "attrs.N" },
+                { "name": "blockSize", "type": "u32", "value": "attrs.block_size" }
+              ]
+            }
+          ],
+          "dispatch": { "x": "min(attrs.N, 65535)", "y": "ceilDiv(attrs.N, 65535)", "z": 1 }
         }
       ]
     },
     {
       "id": "sgmat",
       "priority": 15,
+      "when": ["sgmatGeometryValid", "commonShapeValid", "tensorDtypes.aT == \"float32\"", "dim(shapes.aT, 0) >= sgmatTileRows", "dim(shapes.aT, 0) % sgmatTileRows == 0", "attrs.N % tunables.SGMAT_TILE_COLS == 0", "attrs.K % tunables.SGMAT_TILE_K == 0", "attrs.K % attrs.block_size == 0", "attrs.block_size % sgmatLoadWidth == 0", "wave32Effective", "sgmatWorkgroupFits"],
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "limits": { "maxComputeWorkgroupStorageSize": 8192 },
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
+      "derive": {
         "K": "attrs.K",
         "N": "attrs.N",
         "blockSize": "attrs.block_size",
         "bLoadsPerRow": "sgmatBLoadsPerRow",
         "aScalar": "\"f32\"",
         "usesF16": false,
+        "absmaxScalar": "\"f16\" if tensorDtypes.absmaxT == \"float16\" else \"f32\""
       },
       "passes": [
         {
           "id": "main",
           "shader": "matmul-bnb4-sgmat.wgsl.jinja",
+          "bindings": ["a", "b", "absmax", "y"],
           "dispatch": { "x": "sgmatDispatchN", "y": "sgmatDispatchM" }
         }
       ]
     {
       "id": "sgmat_widened",
       "priority": 16,
+      "when": ["sgmatGeometryValid", "commonShapeValid", "tensorDtypes.aT == \"float16\"", "dim(shapes.aT, 0) >= sgmatTileRows", "dim(shapes.aT, 0) % sgmatTileRows == 0", "attrs.N % tunables.SGMAT_TILE_COLS == 0", "attrs.K % tunables.SGMAT_TILE_K == 0", "attrs.K % attrs.block_size == 0", "attrs.block_size % sgmatLoadWidth == 0", "wave32Effective", "sgmatWorkgroupFits", "sgmatWidenFits"],
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "limits": { "maxComputeWorkgroupStorageSize": 8192 },
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
+      "derive": {
         "K": "attrs.K",
         "N": "attrs.N",
         "blockSize": "attrs.block_size",
         "bLoadsPerRow": "sgmatBLoadsPerRow",
         "aScalar": "\"f32\"",
         "usesF16": true,
+        "absmaxScalar": "\"f16\" if tensorDtypes.absmaxT == \"float16\" else \"f32\"",
         "srcScalar": "\"f16\"",
         "outScalar": "\"f32\"",
         "wrapNarrowInt": false,
         "wrapSigned": false
       },
+      "intermediates": [
+        { "id": "aF32", "dtype": "float32", "shape": "[numel(shapes.aT)]" },
+        { "id": "yF32", "dtype": "float32", "shape": "[dim(shapes.aT, 0) * attrs.N]" }
+      ],
       "passes": [
         {
           "id": "widen_a",
           "name": "MatMulBnb4.WidenActivations",
           "shader": "cast-scalar-x4.wgsl.jinja",
+          "bindings": [
+            { "arg": "aT", "name": "x", "elementType": "$srcScalar" },
+            { "scratch": "aF32", "name": "y", "elementType": "f32" },
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.aT)" }] }
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.aT), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.aT), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         },
         {
           "id": "main",
           "name": "MatMulBnb4.SubgroupMatrixWidened",
           "shader": "matmul-bnb4-sgmat.wgsl.jinja",
+          "bindings": [
+            { "scratch": "aF32", "name": "a", "buffer": "read-only-storage", "elementType": "f32" },
+            "b",
+            "absmax",
+            { "scratch": "yF32", "name": "y", "elementType": "f32" }
+          ],
           "dispatch": { "x": "sgmatDispatchN", "y": "sgmatDispatchM" }
         },
         {
           "id": "narrow_y",
           "name": "MatMulBnb4.NarrowOutput",
           "shader": "cast-scalar-x4.wgsl.jinja",
+          "derive": { "outScalar": "\"f16\"" },
+          "bindings": [
+            { "scratch": "yF32", "name": "x", "buffer": "read-only-storage", "elementType": "f32" },
+            { "arg": "yT", "name": "y", "elementType": "$srcScalar" },
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "dim(shapes.aT, 0) * attrs.N" }] }
+          ],
           "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(dim(shapes.aT, 0) * attrs.N, 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(dim(shapes.aT, 0) * attrs.N, 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
           }
         }
       ]
     },
     {
       "id": "tiled",
       "priority": 10,
       "when": ["tileEligible", "tileWorkgroupFits"],
+      "derive": {
         "tileK": "tunables.PORTABLE_TILE_K",
+        "aScalar": "\"f16\" if tensorDtypes.aT == \"float16\" else \"f32\"",
+        "usesF16": "tensorDtypes.aT == \"float16\"",
+        "absmaxScalar": "\"f16\" if tensorDtypes.absmaxT == \"float16\" else \"f32\"",
+        "tileRows": "portableTileRows"
       },
       "passes": [
         {
           "id": "main",
           "shader": "matmul-bnb4-tiled.wgsl.jinja",
+          "bindings": ["a", "b", "absmax", "y", "params"],
+          "dispatch": { "x": "ceilDiv(attrs.N, 64)", "y": "ceilDiv(dim(shapes.aT, 0), portableTileRows)" }
         }
       ]
     },
       "id": "scalar",
       "priority": 0,
       "when": ["commonShapeValid", "portableWorkgroupSize > 0"],
+      "derive": {
         "workgroupSize": "portableWorkgroupSize",
+        "aScalar": "\"f16\" if tensorDtypes.aT == \"float16\" else \"f32\"",
+        "usesF16": "tensorDtypes.aT == \"float16\"",
+        "absmaxScalar": "\"f16\" if tensorDtypes.absmaxT == \"float16\" else \"f32\""
       },
       "passes": [
         {
           "id": "main",
           "shader": "matmul-bnb4.wgsl.jinja",
+          "bindings": ["a", "b", "absmax", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.yT)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.yT)), (workgroupSize)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/matmul-bnb4-gemv.wgsl.jinja CHANGED Viewed

@@ -1,6 +1,3 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ workgroupSize }}u;
@@ -47,14 +44,13 @@ fn dequant_weight(flat: u32) -> f32 {
 @compute @workgroup_size(WG, 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. The fold over-dispatches, so tail workgroups must leave before the
   // first barrier — hence the early return here rather than a guarded write.
-  let col = wid.x + wid.y * nwg.x;
   if (col >= params.N) {
     return;
   }

 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
 fn main(
   @builtin(workgroup_id) wid: 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. The fold over-dispatches, so tail workgroups must leave before the
   // first barrier — hence the early return here rather than a guarded write.
+  let col = wid.x + wid.y * {{ DISPATCH_FOLD_WIDTH }}u;
   if (col >= params.N) {
     return;
   }

build/webgpu/matmul-bnb4-sgmat.wgsl.jinja CHANGED Viewed

@@ -1,32 +1,46 @@
 // Large aligned prefill tier. The packed B tile is dequantized once per
 // workgroup, then reused by subgroup-matrix f32 multiply-accumulate operations.
 // This specialization operates on full A, B, and output tiles, so both
 // subgroupMatrixLoad and the direct output stores are in bounds.
-enable subgroups;
-{% if pinSubgroupSize32 %}
-enable subgroup_size_control;
-{% endif %}
-enable chromium_experimental_subgroup_matrix;
 {% if usesF16 %}
 // The widened path keeps the activations and the output in f32 for the matrix
 // units, but the quantization scales stay in the request's own dtype.
 enable f16;
 {% endif %}
 diagnostic(off, chromium.subgroup_matrix_uniformity);
 {{ env.wgsl.resourceDeclarations }}
 const K: u32 = {{ K }}u;
 const N: u32 = {{ N }}u;
 const BLOCK_SIZE: u32 = {{ blockSize }}u;
 const TILE_ROWS: u32 = {{ tileRows }}u;
 const TILE_COLS: u32 = {{ tileCols }}u;
 const TILE_K: u32 = {{ tileK }}u;
 const SUB_ROWS: u32 = {{ subRows }}u;
 const SUB_COLS: u32 = {{ subCols }}u;
 const LOAD_WIDTH: u32 = {{ loadWidth }}u;
 var<workgroup> tile_B: array<f32, {{ tileCols * tileK }}>;
 /* bitsandbytes 4-bit weight decoding.
@@ -64,7 +78,7 @@ fn load_b(tile_base: u32, k_base: u32, row: u32, col_group: u32) {
   let tile_offset = row * TILE_K + col;
   // Eligibility guarantees that each aligned LOAD_WIDTH slice stays in one
-  // quantization block and begins on an even weight. The loader widens each
   // packed byte once and reuses both nibbles.
   for (var i = 0u; i < LOAD_WIDTH; i = i + 2u) {
     let packed = b[(flat + i) >> 1u];
@@ -75,13 +89,7 @@ fn load_b(tile_base: u32, k_base: u32, row: u32, col_group: u32) {
   }
 }
-@compute @workgroup_size({{ workgroupSize }}, 1, 1){{ " @subgroup_size(32)" if pinSubgroupSize32 else "" }}
-fn main(
-  @builtin(workgroup_id) workgroup_id: vec3<u32>,
-  @builtin(local_invocation_index) local_idx: u32,
-  @builtin(subgroup_size) subgroup_size: u32
-) {
-  let a_base = workgroup_id.y * TILE_ROWS;
   let b_base = workgroup_id.x * TILE_COLS;
   let subgroup_id = local_idx / subgroup_size;
@@ -95,14 +103,30 @@ fn main(
 {% endfor %}{% endfor %}
   for (var k_base = 0u; k_base < K; k_base += TILE_K) {
     load_b(b_base, k_base, local_idx / {{ bLoadsPerRow }}u, local_idx % {{ bLoadsPerRow }}u);
     workgroupBarrier();
     for (var step = 0u; step < TILE_K; step += {{ matrixSize }}u) {
-      let matrix_a_offset = (a_base + base_A) * K + k_base + step;
 {% for row in range(rowMatrices) %}
       let matA{{ row }} = subgroupMatrixLoad<subgroup_matrix_left<f32, {{ matrixSize }}, {{ matrixSize }}>, row_major>(
-        &a, matrix_a_offset + {{ row * matrixSize }}u * K, K
       );
 {% endfor %}
       let matrix_b_offset = subtile_idx * SUB_COLS * TILE_K + step;
@@ -118,13 +142,54 @@ fn main(
     workgroupBarrier();
   }
-  let output_offset = (a_base + base_A) * N + b_base + base_B;
 {% for row in range(rowMatrices) %}{% for col in range(colMatrices) %}
   subgroupMatrixStore<row_major>(
-    &y,
-    output_offset + {{ row * matrixSize }}u * N + {{ col * matrixSize }}u,
     matC{{ row }}{{ col }},
-    N
   );
 {% endfor %}{% endfor %}
 }

+{% set hybridRows = hybridRows | default(false) %}
+{% set stageActivations = stageActivations | default(false) %}
+{% set aStride = "TILE_K" if stageActivations else "K" %}
+{% set yStride = "TILE_COLS" if stageActivations else "N" %}
 // Large aligned prefill tier. The packed B tile is dequantized once per
 // workgroup, then reused by subgroup-matrix f32 multiply-accumulate operations.
 // This specialization operates on full A, B, and output tiles, so both
 // subgroupMatrixLoad and the direct output stores are in bounds.
 {% if usesF16 %}
 // The widened path keeps the activations and the output in f32 for the matrix
 // units, but the quantization scales stay in the request's own dtype.
 enable f16;
 {% endif %}
+enable subgroups;
+{% if pinSubgroupSize32 %}
+enable subgroup_size_control;
+{% endif %}
+enable chromium_experimental_subgroup_matrix;
 diagnostic(off, chromium.subgroup_matrix_uniformity);
 {{ env.wgsl.resourceDeclarations }}
 const K: u32 = {{ K }}u;
 const N: u32 = {{ N }}u;
 const BLOCK_SIZE: u32 = {{ blockSize }}u;
+{% if not hybridRows %}
 const TILE_ROWS: u32 = {{ tileRows }}u;
+{% endif %}
 const TILE_COLS: u32 = {{ tileCols }}u;
 const TILE_K: u32 = {{ tileK }}u;
+{% if not hybridRows %}
 const SUB_ROWS: u32 = {{ subRows }}u;
+{% endif %}
+{% if not hybridRows %}
 const SUB_COLS: u32 = {{ subCols }}u;
+{% endif %}
 const LOAD_WIDTH: u32 = {{ loadWidth }}u;
 var<workgroup> tile_B: array<f32, {{ tileCols * tileK }}>;
+{% if stageActivations or hybridRows %}
+var<workgroup> tile_A: array<f32, {{ ((hybridTailTileRows | default(0)) if hybridRows else tileRows) * tileK }}>;
+{% endif %}
 /* bitsandbytes 4-bit weight decoding.
   let tile_offset = row * TILE_K + col;
   // Eligibility guarantees that each aligned LOAD_WIDTH slice stays in one
+  // quantization block and begins on an even weight. This shader widens each
   // packed byte once and reuses both nibbles.
   for (var i = 0u; i < LOAD_WIDTH; i = i + 2u) {
     let packed = b[(flat + i) >> 1u];
   }
 }
+{% macro tile_body(stageActivations, tileRows, subRows, subCols, rowMatrices, colMatrices, rowSubtiles, workgroupSize, rowOffset, aStride, yStride) %}  let a_base = workgroup_id.y * TILE_ROWS{% if (rowOffset | default(0)) > 0 %} + {{ rowOffset | default(0) }}u{% endif %};
   let b_base = workgroup_id.x * TILE_COLS;
   let subgroup_id = local_idx / subgroup_size;
 {% endfor %}{% endfor %}
   for (var k_base = 0u; k_base < K; k_base += TILE_K) {
+{% if workgroupSize == ((tileCols * tileK / loadWidth) | int) %}
     load_b(b_base, k_base, local_idx / {{ bLoadsPerRow }}u, local_idx % {{ bLoadsPerRow }}u);
+{% else %}
+    for (var group = local_idx; group < {{ ((tileCols * tileK / loadWidth) | int) }}u; group += {{ workgroupSize }}u) {
+      load_b(b_base, k_base, group / {{ bLoadsPerRow }}u, group % {{ bLoadsPerRow }}u);
+    }
+{% endif %}
+{% if stageActivations %}
+    for (var linear = local_idx; linear < {{ tileRows * tileK }}u; linear += {{ workgroupSize }}u) {
+      let row = a_base + linear / TILE_K;
+      var value = 0.0;
+      if (row < {{ rowCount | default(0) }}u) {
+        value = f32(a[row * K + k_base + linear % TILE_K]);
+      }
+      tile_A[linear] = value;
+    }
+{% endif %}
     workgroupBarrier();
     for (var step = 0u; step < TILE_K; step += {{ matrixSize }}u) {
+      let matrix_a_offset = {% if stageActivations %}base_A * TILE_K + step{% else %}(a_base + base_A) * K + k_base + step{% endif %};
 {% for row in range(rowMatrices) %}
       let matA{{ row }} = subgroupMatrixLoad<subgroup_matrix_left<f32, {{ matrixSize }}, {{ matrixSize }}>, row_major>(
+        &{{ "tile_A" if stageActivations else "a" }}, matrix_a_offset + {{ row * matrixSize }}u * {{ aStride }}, {{ aStride }}
       );
 {% endfor %}
       let matrix_b_offset = subtile_idx * SUB_COLS * TILE_K + step;
     workgroupBarrier();
   }
+  let output_offset = {% if stageActivations %}base_A * TILE_COLS + base_B{% else %}(a_base + base_A) * N + b_base + base_B{% endif %};
 {% for row in range(rowMatrices) %}{% for col in range(colMatrices) %}
   subgroupMatrixStore<row_major>(
+    &{{ "tile_B" if stageActivations else "y" }},
+    output_offset + {{ row * matrixSize }}u * {{ yStride }} + {{ col * matrixSize }}u,
     matC{{ row }}{{ col }},
+    {{ yStride }}
   );
 {% endfor %}{% endfor %}
+{% if stageActivations %}
+  // The closing K-loop barrier ends all weight reads. Output banks reuse B
+  // staging, then a barrier publishes them for disjoint masked global stores.
+  workgroupBarrier();
+  for (var linear = local_idx; linear < {{ tileRows * tileCols }}u; linear += {{ workgroupSize }}u) {
+    let row = a_base + linear / TILE_COLS;
+    if (row < {{ rowCount | default(0) }}u) {
+      y[row * N + b_base + linear % TILE_COLS] = {{ aScalar }}(tile_B[linear]);
+    }
+  }
+{% endif %}{% endmacro %}
+{% if hybridRows %}
+fn full_tile(workgroup_id: vec3<u32>, local_idx: u32, subgroup_size: u32) {
+  const TILE_ROWS: u32 = {{ tileRows }}u;
+  const SUB_ROWS: u32 = {{ subRows }}u;
+  const SUB_COLS: u32 = {{ subCols }}u;
+{{ tile_body(false, tileRows, subRows, subCols, rowMatrices, colMatrices, rowSubtiles, workgroupSize, 0, "K", "N") }}}
+fn partial_tile(workgroup_id: vec3<u32>, local_idx: u32, subgroup_size: u32) {
+  const TILE_ROWS: u32 = {{ hybridTailTileRows }}u;
+  const SUB_ROWS: u32 = {{ matrixSize }}u;
+  const SUB_COLS: u32 = {{ hybridTailSubCols }}u;
+{{ tile_body(true, hybridTailTileRows, matrixSize, hybridTailSubCols, 1, ((hybridTailSubCols / matrixSize) | int), hybridTailRowSubtiles, workgroupSize, rowOffset, "TILE_K", "TILE_COLS") }}}
+@compute @workgroup_size({{ workgroupSize }}, 1, 1){{ " @subgroup_size(32)" if pinSubgroupSize32 else "" }}
+fn main(
+  @builtin(workgroup_id) workgroup_id: vec3<u32>,
+  @builtin(local_invocation_index) local_idx: u32,
+  @builtin(subgroup_size) subgroup_size: u32
+) {
+  if (workgroup_id.y < {{ fullDispatchRows }}u) {
+    full_tile(workgroup_id, local_idx, subgroup_size);
+  } else {
+    partial_tile(vec3<u32>(workgroup_id.x, workgroup_id.y - {{ fullDispatchRows }}u, workgroup_id.z), local_idx, subgroup_size);
+  }
 }
+{% else %}
+@compute @workgroup_size({{ workgroupSize }}, 1, 1){{ " @subgroup_size(32)" if pinSubgroupSize32 else "" }}
+fn main(
+  @builtin(workgroup_id) workgroup_id: vec3<u32>,
+  @builtin(local_invocation_index) local_idx: u32,
+  @builtin(subgroup_size) subgroup_size: u32
+) {
+{{ tile_body(stageActivations, tileRows, subRows, subCols, rowMatrices, colMatrices, rowSubtiles, workgroupSize, (rowOffset | default(0)), aStride, yStride) }}}{% endif %}

build/webgpu/matmul-bnb4-tiled.wgsl.jinja CHANGED Viewed

@@ -1,10 +1,7 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-// Prefill: register-blocked 64x64 shared-memory tiled GEMM (16x8 threads, each
-// holding an 8x4 micro-tile, with K stepped in BK chunks).
 //
 // Each weight is dequantized once into the shared B tile and reused across the
 // row block. Both operands are stored K-major — A is [M, K] and the packed
@@ -12,9 +9,9 @@ enable f16;
 // staging read is contiguous, and one micro-tile step reads TM + TN vector
 // words to issue TM * TN * 4 fused multiply-adds through dot().
 const BK: u32 = {{ tileK }}u;
-const BM: u32 = 64u;
 const BN: u32 = 64u;
-const TM: u32 = 8u;
 const TN: u32 = 4u;
 const WG_X: u32 = 16u;
 const WG_SIZE: u32 = 128u;
@@ -97,7 +94,7 @@ fn main(
   let aRow = lid.y * TM;
   let bCol = lid.x * TN;
-{% for row in range(8) %}
   var acc{{ row }} = vec4<f32>(0.0);
 {% endfor %}
   let numTiles = (params.K + BK - 1u) / BK;
@@ -148,7 +145,7 @@ fn main(
 {% for column in range(4) %}
       let bv{{ column }} = tileB[bCol + {{ column }}u][kv];
 {% endfor %}
-{% for row in range(8) %}
       let av{{ row }} = vec4<f32>(tileA[aRow + {{ row }}u][kv]);
       acc{{ row }} += vec4<f32>({% for column in range(4) %}dot(av{{ row }}, bv{{ column }}){% if not loop.last %}, {% endif %}{% endfor %});
 {% endfor %}
@@ -156,7 +153,7 @@ fn main(
     workgroupBarrier();
   }
-{% for row in range(8) %}
 {% for column in range(4) %}
   {
     let m = mBase + aRow + {{ row }}u;

 {{ env.wgsl.resourceDeclarations }}
+// Prefill: register-blocked {{ tileRows }}x64 shared-memory tiled GEMM (16x8 threads, each
+// holding an {{ ((tileRows / 8) | int) }}x4 micro-tile, with K stepped in BK chunks).
 //
 // Each weight is dequantized once into the shared B tile and reused across the
 // row block. Both operands are stored K-major — A is [M, K] and the packed
 // staging read is contiguous, and one micro-tile step reads TM + TN vector
 // words to issue TM * TN * 4 fused multiply-adds through dot().
 const BK: u32 = {{ tileK }}u;
+const BM: u32 = {{ tileRows }}u;
 const BN: u32 = 64u;
+const TM: u32 = {{ ((tileRows / 8) | int) }}u;
 const TN: u32 = 4u;
 const WG_X: u32 = 16u;
 const WG_SIZE: u32 = 128u;
   let aRow = lid.y * TM;
   let bCol = lid.x * TN;
+{% for row in range(((tileRows / 8) | int)) %}
   var acc{{ row }} = vec4<f32>(0.0);
 {% endfor %}
   let numTiles = (params.K + BK - 1u) / BK;
 {% for column in range(4) %}
       let bv{{ column }} = tileB[bCol + {{ column }}u][kv];
 {% endfor %}
+{% for row in range(((tileRows / 8) | int)) %}
       let av{{ row }} = vec4<f32>(tileA[aRow + {{ row }}u][kv]);
       acc{{ row }} += vec4<f32>({% for column in range(4) %}dot(av{{ row }}, bv{{ column }}){% if not loop.last %}, {% endif %}{% endfor %});
 {% endfor %}
     workgroupBarrier();
   }
+{% for row in range(((tileRows / 8) | int)) %}
 {% for column in range(4) %}
   {
     let m = mBase + aRow + {{ row }}u;

build/webgpu/matmul-bnb4.wgsl.jinja CHANGED Viewed

@@ -1,6 +1,3 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ workgroupSize }}u;
@@ -41,11 +38,11 @@ fn dequant_weight(flat: u32) -> f32 {
 @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) {

 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat output-element index: gid.y carries the high bits past the
+  // per-axis dispatch fold width. With no fold this
   // reduces to gid.x; the index >= total guard drops the tail.
+  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
   let total = params.rows * params.N;
   if (index >= total) {

build/webgpu/metadata.json CHANGED Viewed

@@ -1,22 +1,34 @@
 {
   "name": "com.microsoft.MatMulBnb4",
-  "id": "_com_microsoft_matmulbnb4_webgpu_aab3ed9",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "65V0CAoAl7vsenxQG3P5y9j0T4tJEprrb5BwhczZ/So=",
-      "cast-scalar-x4.wgsl.jinja": "uplwFdE8dJSffJop4txQHZgeGDDJ5K9+F3V4n1CotVo=",
-      "manifest.json": "TFYtgeuCmTA2Ei4kkTZQ2PUcZP7JT9DEynhMUofG+BU=",
-      "matmul-bnb4-gemv.wgsl.jinja": "7e4ZY+raLAMG/wmXm72XpSzPsR6adHvvrcKxv2ZKj/4=",
-      "matmul-bnb4-sgmat.wgsl.jinja": "x3RRM9yb0++oJ2o6qqmhx+qwCqb2MAZui1Pdlof6nsU=",
-      "matmul-bnb4-tiled.wgsl.jinja": "+pb1pUgFU+RLITtJ+zrnTDl3LUmU7sZDGqIdyKtAmn4=",
-      "matmul-bnb4.wgsl.jinja": "xzeJm4jFqbXih7plcQup4jUQsRw4h8HpLzcAxWR2aic=",
-      "test.json": "1oYDh6veLTUP+o7nE+G7nU1FMK88nLKArXE/FcbOfTo="
     }
   },
-  "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.MatMulBnb4" }
 }

 {
   "name": "com.microsoft.MatMulBnb4",
+  "id": "_com_microsoft_matmulbnb4_webgpu_237048a",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "/Z1zFaT4D+0TEFMxG/zqun5CjI9irI93GNGwLZdLXHQ=",
+      "cast-scalar-x4.wgsl.jinja": "s8Y9E2+D7cFgYTUEJhiuPsUzvHjv2mXVoiRmv0A0vTU=",
+      "manifest.json": "TEJfUwu6Isit+t0fBCRp0Ny+uC8hUjSe2jgG5kCFcEg=",
+      "matmul-bnb4-gemv.wgsl.jinja": "7W3L5RH+5dk5/dKNnjwVFtBt+xgCxorJJ04/8zHrySw=",
+      "matmul-bnb4-sgmat.wgsl.jinja": "9WS5F6zd/KFztTpbZjBgJe7WIl0l7uiRelcDv/9YbBc=",
+      "matmul-bnb4-tiled.wgsl.jinja": "bADcd5OFuhrZ33Qntqh1PXREZolJZ+RKS+VlweAjPo0=",
+      "matmul-bnb4.wgsl.jinja": "6rtf2K0op1IJv2+l66HJjq8XnBATGyefxit0Hj5QREI=",
+      "test.json": "KPAf0hjucty64auyD+xOS5RAGDEPYeAlUSJPhbig7CA="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "sgmat_hybrid_rows": ["matmul-bnb4-sgmat.wgsl.jinja"],
+      "sgmat_hybrid_rows_widened": ["cast-scalar-x4.wgsl.jinja", "matmul-bnb4-sgmat.wgsl.jinja"],
+      "sgmat_staged_rows": ["matmul-bnb4-sgmat.wgsl.jinja"],
+      "gemv": ["matmul-bnb4-gemv.wgsl.jinja"],
+      "sgmat": ["matmul-bnb4-sgmat.wgsl.jinja"],
+      "sgmat_widened": ["cast-scalar-x4.wgsl.jinja", "matmul-bnb4-sgmat.wgsl.jinja"],
+      "tiled": ["matmul-bnb4-tiled.wgsl.jinja"],
+      "scalar": ["matmul-bnb4.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.MatMulBnb4",
   "fixtureArrays": {
     "bnb4_weight_cycle_b_t": [1, 35, 69, 103, 137, 171, 205, 239, 254, 220, 186, 152, 118, 84, 50, 16, 90]
   },
@@ -408,7 +407,7 @@
     {
       "name": "nf4_sgmat_widened_f16",
       "provenance": {
-        "notes": "A float16 request on the aligned subgroup-matrix shape. The matrix units take float32 operands and accumulate in float32, so this path widens the activations either side of the multiply rather than using the device's float16 configuration, whose result component is also float16. It is the only case that reaches the widening passes."
       },
       "attrs": { "K": 64, "N": 64, "block_size": 32, "quant_type": 1 },
       "inputs": {
@@ -429,6 +428,1449 @@
         }
       },
       "outputs": { "yT": { "dtype": "float16", "shape": [128, 64], "tolerance": 0.002, "relTolerance": 0.002 } }
     }
   ]
 }

 {
   "fixtureArrays": {
     "bnb4_weight_cycle_b_t": [1, 35, 69, 103, 137, 171, 205, 239, 254, 220, 186, 152, 118, 84, 50, 16, 90]
   },
     {
       "name": "nf4_sgmat_widened_f16",
       "provenance": {
+        "notes": "A float16 request on an aligned subgroup-matrix shape widens activations to float32 around the multiply, accumulates in float32, and narrows the result afterward."
       },
       "attrs": { "K": 64, "N": 64, "block_size": 32, "quant_type": 1 },
       "inputs": {
         }
       },
       "outputs": { "yT": { "dtype": "float16", "shape": [128, 64], "tolerance": 0.002, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "nf4_tiled_speculative_band_m5",
+      "provenance": {
+        "notes": "A five-row speculative-decoder multiply exercises the tiled route within the 2-to-15-row band, using a ragged K dimension and NF4 block-quantized weights."
+      },
+      "attrs": { "K": 33, "N": 19, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [5, 33],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [314],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [40],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [5, 19], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "rows_float32_q0_m2_k64_n64",
+      "attrs": { "K": 64, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [2, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [2, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q1_m3_k128_n128",
+      "attrs": { "K": 128, "N": 128, "block_size": 32, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [3, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [8192],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [3, 128], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q0_m7_k512_n64",
+      "attrs": { "K": 512, "N": 64, "block_size": 64, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [7, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [16384],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [7, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q1_m8_k64_n128",
+      "attrs": { "K": 64, "N": 128, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [8, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [8, 128], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q0_m9_k128_n64",
+      "attrs": { "K": 128, "N": 64, "block_size": 32, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [9, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [9, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q1_m15_k512_n128",
+      "attrs": { "K": 512, "N": 128, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [15, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [32768],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [15, 128], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q0_m16_k64_n64",
+      "attrs": { "K": 64, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [16, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [16, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q1_m17_k128_n128",
+      "attrs": { "K": 128, "N": 128, "block_size": 32, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [17, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [8192],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [17, 128], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q0_m24_k512_n64",
+      "attrs": { "K": 512, "N": 64, "block_size": 64, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [24, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [16384],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [24, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q1_m31_k64_n128",
+      "attrs": { "K": 64, "N": 128, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [31, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [31, 128], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q0_m33_k128_n64",
+      "attrs": { "K": 128, "N": 64, "block_size": 32, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [33, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [33, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q1_m63_k512_n128",
+      "attrs": { "K": 512, "N": 128, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [63, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [32768],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [63, 128], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q0_m65_k64_n64",
+      "attrs": { "K": 64, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [65, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [65, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float32_q1_m129_k128_n128",
+      "attrs": { "K": 128, "N": 128, "block_size": 32, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [129, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [8192],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [129, 128], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q0_m2_k64_n64",
+      "attrs": { "K": 64, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [2, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [2, 64], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q1_m3_k128_n128",
+      "attrs": { "K": 128, "N": 128, "block_size": 32, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [3, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [8192],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [3, 128], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q0_m7_k512_n64",
+      "attrs": { "K": 512, "N": 64, "block_size": 64, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [7, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [16384],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [7, 64], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q1_m8_k64_n128",
+      "attrs": { "K": 64, "N": 128, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [8, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [8, 128], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q0_m9_k128_n64",
+      "attrs": { "K": 128, "N": 64, "block_size": 32, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [9, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [9, 64], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q1_m15_k512_n128",
+      "attrs": { "K": 512, "N": 128, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [15, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [32768],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [15, 128], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q0_m16_k64_n64",
+      "attrs": { "K": 64, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [16, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [16, 64], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q1_m17_k128_n128",
+      "attrs": { "K": 128, "N": 128, "block_size": 32, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [17, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [8192],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [17, 128], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q0_m24_k512_n64",
+      "attrs": { "K": 512, "N": 64, "block_size": 64, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [24, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [16384],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [24, 64], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q1_m31_k64_n128",
+      "attrs": { "K": 64, "N": 128, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [31, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [31, 128], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q0_m33_k128_n64",
+      "attrs": { "K": 128, "N": 64, "block_size": 32, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [33, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4096],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [33, 64], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q1_m63_k512_n128",
+      "attrs": { "K": 512, "N": 128, "block_size": 64, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [63, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [32768],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [63, 128], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q0_m65_k64_n64",
+      "attrs": { "K": 64, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [65, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [65, 64], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "rows_float16_q1_m129_k128_n128",
+      "attrs": { "K": 128, "N": 128, "block_size": 32, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [129, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [8192],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [129, 128], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "matrix_tunable_float32_q0_m15_k192_n64_32_8",
+      "attrs": { "K": 192, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [15, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [6144],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [15, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 32, "SGMAT_TILE_K": 8 }
+    },
+    {
+      "name": "matrix_tunable_float32_q0_m128_k192_n64_32_8",
+      "attrs": { "K": 192, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [6144],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 32, "SGMAT_TILE_K": 8 }
+    },
+    {
+      "name": "matrix_tunable_float16_q1_m15_k192_n128_64_16",
+      "attrs": { "K": 192, "N": 128, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [15, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [12288],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [1536],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [15, 128], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 64, "SGMAT_TILE_K": 16 }
+    },
+    {
+      "name": "matrix_tunable_float16_q1_m128_k192_n128_64_16",
+      "attrs": { "K": 192, "N": 128, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [128, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [12288],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [1536],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [128, 128], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 64, "SGMAT_TILE_K": 16 }
+    },
+    {
+      "name": "matrix_tunable_float32_q0_m15_k192_n192_96_24",
+      "attrs": { "K": 192, "N": 192, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [15, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [18432],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [2304],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [15, 192], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 96, "SGMAT_TILE_K": 24 }
+    },
+    {
+      "name": "matrix_tunable_float32_q0_m128_k192_n192_96_24",
+      "attrs": { "K": 192, "N": 192, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [18432],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [2304],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 192], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 96, "SGMAT_TILE_K": 24 }
+    },
+    {
+      "name": "matrix_tunable_float16_q1_m15_k192_n256_128_64",
+      "attrs": { "K": 192, "N": 256, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [15, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [24576],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [3072],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [15, 256], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 128, "SGMAT_TILE_K": 64 }
+    },
+    {
+      "name": "matrix_tunable_float16_q1_m128_k192_n256_128_64",
+      "attrs": { "K": 192, "N": 256, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [128, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [24576],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [3072],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [128, 256], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 128, "SGMAT_TILE_K": 64 }
+    },
+    {
+      "name": "matrix_tunable_float32_q0_m15_k192_n64_32_64",
+      "attrs": { "K": 192, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [15, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [6144],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [15, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 32, "SGMAT_TILE_K": 64 }
+    },
+    {
+      "name": "matrix_tunable_float32_q0_m128_k192_n64_32_64",
+      "attrs": { "K": 192, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [6144],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [768],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 32, "SGMAT_TILE_K": 64 }
+    },
+    {
+      "name": "matrix_tunable_float16_q1_m15_k192_n256_128_8",
+      "attrs": { "K": 192, "N": 256, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [15, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [24576],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [3072],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [15, 256], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 128, "SGMAT_TILE_K": 8 }
+    },
+    {
+      "name": "matrix_tunable_float16_q1_m128_k192_n256_128_8",
+      "attrs": { "K": 192, "N": 256, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [128, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [24576],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [3072],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [128, 256], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 128, "SGMAT_TILE_K": 8 }
+    },
+    {
+      "name": "matrix_tunable_float32_q0_m15_k192_n128_64_24",
+      "attrs": { "K": 192, "N": 128, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [15, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [12288],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [1536],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [15, 128], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 64, "SGMAT_TILE_K": 24 }
+    },
+    {
+      "name": "matrix_tunable_float32_q0_m128_k192_n128_64_24",
+      "attrs": { "K": 192, "N": 128, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [128, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [12288],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [1536],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [128, 128], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 64, "SGMAT_TILE_K": 24 }
+    },
+    {
+      "name": "matrix_tunable_float16_q1_m15_k192_n192_96_16",
+      "attrs": { "K": 192, "N": 192, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [15, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [18432],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [2304],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [15, 192], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 96, "SGMAT_TILE_K": 16 }
+    },
+    {
+      "name": "matrix_tunable_float16_q1_m128_k192_n192_96_16",
+      "attrs": { "K": 192, "N": 192, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [128, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [18432],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [2304],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [128, 192], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "SGMAT_TILE_COLS": 96, "SGMAT_TILE_K": 16 }
+    },
+    {
+      "name": "row_and_weight_tails_float32_q1_m9_k31_n65",
+      "attrs": { "K": 31, "N": 65, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [9, 31],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1008],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [126],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [9, 65], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "row_and_weight_tails_float32_q1_m33_k129_n127",
+      "attrs": { "K": 129, "N": 127, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [33, 129],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [8192],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [33, 127], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "row_and_weight_tails_float16_q1_m9_k31_n65",
+      "attrs": { "K": 31, "N": 65, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [9, 31],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1008],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [126],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [9, 65], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "row_and_weight_tails_float16_q1_m33_k129_n127",
+      "attrs": { "K": 129, "N": 127, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [33, 129],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [8192],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [33, 127], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "portable_f16_tunable_float16_q0_m17_k129_n65_32",
+      "attrs": { "K": 129, "N": 65, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [17, 129],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4193],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [525],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [17, 65], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "PORTABLE_TILE_K": 32 }
+    },
+    {
+      "name": "portable_f16_tunable_float16_q0_m17_k129_n65_64",
+      "attrs": { "K": 129, "N": 65, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [17, 129],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [4193],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [525],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [17, 65], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      },
+      "tunables": { "PORTABLE_TILE_K": 64 }
+    },
+    {
+      "name": "empty_rows_float32_q1_m0_k32_n64",
+      "attrs": { "K": 32, "N": 64, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [0, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [0, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "empty_rows_float16_q1_m0_k32_n64",
+      "attrs": { "K": 32, "N": 64, "block_size": 16, "quant_type": 1 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [0, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [1024],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [0, 64], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "Covers row-tile boundaries, f32 accumulation and position-varying NF4/FP4 weights on matrix and portable capability tiers."
+      }
+    },
+    {
+      "name": "prefix_tail_float32_m95_k64_n64",
+      "attrs": { "K": 64, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [95, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [95, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "The full-row prefix and a multi-fragment tail write disjoint rows, including the last output, with f32 accumulation."
+      }
+    },
+    {
+      "name": "prefix_tail_float32_m191_k64_n64",
+      "attrs": { "K": 64, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float32",
+          "shape": [191, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float32",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float32", "shape": [191, 64], "tolerance": 0.00002, "relTolerance": 0.00001 } },
+      "provenance": {
+        "notes": "The full-row prefix and a multi-fragment tail write disjoint rows, including the last output, with f32 accumulation."
+      }
+    },
+    {
+      "name": "prefix_tail_float16_m95_k64_n64",
+      "attrs": { "K": 64, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [95, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [95, 64], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "The full-row prefix and a multi-fragment tail write disjoint rows, including the last output, with f32 accumulation."
+      }
+    },
+    {
+      "name": "prefix_tail_float16_m191_k64_n64",
+      "attrs": { "K": 64, "N": 64, "block_size": 16, "quant_type": 0 },
+      "inputs": {
+        "aT": {
+          "dtype": "float16",
+          "shape": [191, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "bT": {
+          "dtype": "uint8",
+          "shape": [2048],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/bnb4_weight_cycle_b_t" } }
+        },
+        "absmaxT": {
+          "dtype": "float16",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.07, "scale": 0.02, "offset": 0.06 }
+        }
+      },
+      "outputs": { "yT": { "dtype": "float16", "shape": [191, 64], "tolerance": 0.001, "relTolerance": 0.001 } },
+      "provenance": {
+        "notes": "The full-row prefix and a multi-fragment tail write disjoint rows, including the last output, with f32 accumulation."
+      }
     }
   ]
 }