sync 91d990483a17

Browse files

Files changed (10) hide show

README.md +57 -11
build/webgpu/bench.json +112 -2
build/webgpu/manifest.json +0 -0
build/webgpu/metadata.json +79 -8
build/webgpu/quant-dp4a-matmul.wgsl.jinja +19 -17
build/webgpu/quant-exact-matrix.wgsl.jinja +126 -0
build/webgpu/quant-exact-portable.wgsl.jinja +143 -0
build/webgpu/quant-exact-prepare.wgsl.jinja +33 -0
build/webgpu/quant-matmul-accumulate-rank4.wgsl.jinja +1 -1
build/webgpu/test.json +731 -24

README.md CHANGED Viewed

@@ -12,24 +12,24 @@ tags:
 ## Description
-Computes an integer matrix product with 8-bit inputs, `int32` accumulation, and independently optional zero points that default to 0. The package implements rank-1 dot products, rank-2 products, rank-2/rank-3 broadcasting, rank-3 products, and rank-4-by-rank-4 products. Scalar zero points are supported throughout; `b_zero_point` additionally supports `[N]` for rank-2 B and `[batch, 1, N]` for non-broadcast rank-3 B. Other standard ONNX matmul rank combinations and N-D per-row/per-column zero-point layouts are not yet implemented.
 See the [ONNX `MatMulInteger` spec](https://onnx.ai/onnx/operators/onnx__MatMulInteger.html) for the reference semantics.
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `A` | `a` | `TA` | — | — | N-dimensional integer matrix A (int8 or uint8). | required |
-| `B` | `b` | `TB` | — | — | N-dimensional integer matrix B (int8 or uint8). | required |
-| `a_zero_point` | `a_zero_point` | `TA` | — | — | Optional scalar zero point for A; defaults to 0. Standard N-D per-row layouts are not yet implemented. | optional |
-| `b_zero_point` | `b_zero_point` | `TB` | — | — | Optional zero point for B; defaults to 0. Supports a scalar, `[N]` for rank-2 B, or `[batch, 1, N]` for non-broadcast rank-3 B; other standard N-D per-column layouts are not yet implemented. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `y` | `TY` | derived | ONNX MatMul result of `A` and `B` | int32 matrix product result of A * B. | required |
 ## Type constraints
@@ -39,23 +39,69 @@ See the [ONNX `MatMulInteger` spec](https://onnx.ai/onnx/operators/onnx__MatMulI
 | `TB` | `uint8`, `int8` |
 | `TY` | `int32` |
 ## Files
-- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 - [`matmul-integer-batched.wgsl.jinja`](build/webgpu/matmul-integer-batched.wgsl.jinja)
 - [`quant-dp4a-matmul.wgsl.jinja`](build/webgpu/quant-dp4a-matmul.wgsl.jinja)
 - [`quant-matmul-accumulate-rank2.wgsl.jinja`](build/webgpu/quant-matmul-accumulate-rank2.wgsl.jinja)
 - [`quant-matmul-accumulate-rank4.wgsl.jinja`](build/webgpu/quant-matmul-accumulate-rank4.wgsl.jinja)
 ## Use with `@huggingface/kernels`
-The loader derives every required output's shape and logical dtype from the manifest contract and this call.
-It then allocates the result tensors automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

 ## Description
+Computes an integer matrix product with 8-bit inputs, `int32` accumulation, and independently optional zero points that default to 0. The package implements rank-1 dot products, rank-2 products, rank-2/rank-3 broadcasting, rank-3 products, and rank-4-by-rank-4 products. Scalar zero points are supported throughout; `b_zero_point` additionally supports `[N]` for rank-2 B and `[batch, 1, N]` for non-broadcast rank-3 B. Other standard ONNX matmul rank combinations and N-D per-row/per-column zero-point layouts are unsupported.
 See the [ONNX `MatMulInteger` spec](https://onnx.ai/onnx/operators/onnx__MatMulInteger.html) for the reference semantics.
 ## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `a` | `A` | `TA` | — | — | N-dimensional integer matrix A (int8 or uint8). | required |
+| `b` | `B` | `TB` | — | — | N-dimensional integer matrix B (int8 or uint8). | required |
+| `a_zero_point` | — | `TA` | — | — | Optional scalar zero point for A; defaults to 0. Standard N-D per-row layouts are unsupported. | optional |
+| `b_zero_point` | — | `TB` | — | — | Optional zero point for B; defaults to 0. Supports a scalar, `[N]` for rank-2 B, or `[batch, 1, N]` for non-broadcast rank-3 B; other standard N-D per-column layouts are unsupported. | optional |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `Y` | `TY` | derived | ONNX MatMul result of `a` and `b` | int32 matrix product result of A * B. | required |
 ## Type constraints
 | `TB` | `uint8`, `int8` |
 | `TY` | `int32` |
+## 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_precast_a0_bnone_f16` — Prepare exact f16 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a0_bnone_f16` — Prepare lossless f16 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a0_bnone` — Prepare exact f32 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a0_bnone` — Prepare lossless f32 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a0_bscalar_f16` — Prepare exact f16 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a0_bscalar_f16` — Prepare lossless f16 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a0_bscalar` — Prepare exact f32 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a0_bscalar` — Prepare lossless f32 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a0_bcolumn_f16` — Prepare exact f16 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a0_bcolumn_f16` — Prepare lossless f16 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a0_bcolumn` — Prepare exact f32 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a0_bcolumn` — Prepare lossless f32 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a1_bnone_f16` — Prepare exact f16 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a1_bnone_f16` — Prepare lossless f16 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a1_bnone` — Prepare exact f32 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a1_bnone` — Prepare lossless f32 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a1_bscalar_f16` — Prepare exact f16 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a1_bscalar_f16` — Prepare lossless f16 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a1_bscalar` — Prepare exact f32 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a1_bscalar` — Prepare lossless f32 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a1_bcolumn_f16` — Prepare exact f16 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a1_bcolumn_f16` — Prepare lossless f16 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_precast_a1_bcolumn` — Prepare exact f32 representations of adjusted byte integers once, then load them directly into f32-accumulating subgroup matrices. Row groups follow matrix geometry, output rows, invocation limits and workgroup storage; bounded partial sums retain exact modular int32 arithmetic.
+- `portable_precast_a1_bcolumn` — Prepare lossless f32 representations of adjusted byte integers, then use vector workgroup GEMM with exact bounded f32 partials and modular int32 totals. Geometry follows output rows, invocation and storage limits; f32 storage requires enough row reuse to amortize preparation.
+- `sgmat_exact_a0_bnone` — Exact integer products through bounded f32 subgroup-matrix partials accumulated with modular int32 additions. Tile height follows workgroup memory and invocation limits.
+- `sgmat_exact_a0_bscalar` — Exact integer products through bounded f32 subgroup-matrix partials accumulated with modular int32 additions. Tile height follows workgroup memory and invocation limits.
+- `sgmat_exact_a0_bcolumn` — Exact integer products through bounded f32 subgroup-matrix partials accumulated with modular int32 additions. Tile height follows workgroup memory and invocation limits.
+- `sgmat_exact_a1_bnone` — Exact integer products through bounded f32 subgroup-matrix partials accumulated with modular int32 additions. Tile height follows workgroup memory and invocation limits.
+- `sgmat_exact_a1_bscalar` — Exact integer products through bounded f32 subgroup-matrix partials accumulated with modular int32 additions. Tile height follows workgroup memory and invocation limits.
+- `sgmat_exact_a1_bcolumn` — Exact integer products through bounded f32 subgroup-matrix partials accumulated with modular int32 additions. Tile height follows workgroup memory and invocation limits.
+## Device requirements
+Some implementation variants require `subgroup-matrix`, `shader-f16`, 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
 - [`matmul-integer-batched.wgsl.jinja`](build/webgpu/matmul-integer-batched.wgsl.jinja)
 - [`quant-dp4a-matmul.wgsl.jinja`](build/webgpu/quant-dp4a-matmul.wgsl.jinja)
+- [`quant-exact-matrix.wgsl.jinja`](build/webgpu/quant-exact-matrix.wgsl.jinja)
+- [`quant-exact-portable.wgsl.jinja`](build/webgpu/quant-exact-portable.wgsl.jinja)
+- [`quant-exact-prepare.wgsl.jinja`](build/webgpu/quant-exact-prepare.wgsl.jinja)
 - [`quant-matmul-accumulate-rank2.wgsl.jinja`](build/webgpu/quant-matmul-accumulate-rank2.wgsl.jinja)
 - [`quant-matmul-accumulate-rank4.wgsl.jinja`](build/webgpu/quant-matmul-accumulate-rank4.wgsl.jinja)
 ## 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": "ai.onnx.MatMulInteger",
   "cases": [
     {
       "name": "u8_128x128",
@@ -302,7 +301,7 @@
       "name": "manifest-gap-rank2-u8s8-k63-dp4a-floor-m512n4096",
       "preset": "stress",
       "provenance": {
-        "source": "authored for boundary coverage",
         "notes": "Realistic narrow projection immediately below the rank-2 DP4A K >= 64 guard; the paired K64 case isolates the selector boundary."
       },
       "vars": { "M": 512, "K": 63, "N": 4096 },
@@ -314,6 +313,117 @@
       },
       "outputs": { "y": { "dtype": "int32", "shape": [512, 4096], "dist": "empty" } },
       "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
     }
   ]
 }

 {
   "cases": [
     {
       "name": "u8_128x128",
       "name": "manifest-gap-rank2-u8s8-k63-dp4a-floor-m512n4096",
       "preset": "stress",
       "provenance": {
+        "source": "synthetic benchmark",
         "notes": "Realistic narrow projection immediately below the rank-2 DP4A K >= 64 guard; the paired K64 case isolates the selector boundary."
       },
       "vars": { "M": 512, "K": 63, "N": 4096 },
       },
       "outputs": { "y": { "dtype": "int32", "shape": [512, 4096], "dist": "empty" } },
       "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
+    },
+    {
+      "name": "matrix_storage_m32_k128_n8192",
+      "preset": "stress",
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [32, 128], "dist": "randint", "seed": 1303, "min": 0, "max": 255 },
+        "b": { "dtype": "int8", "shape": [128, 8192], "dist": "randint", "seed": 1304, "min": -127, "max": 127 },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [32, 8192] } },
+      "tunableSpace": {
+        "DP4A_LARGE_TILE_M": [16, 32],
+        "DP4A_LARGE_TILE_N": [64, 128],
+        "DP4A_LARGE_K_TILE_WORDS": [16, 32, 64]
+      },
+      "bench": { "metrics": [{ "type": "gflops", "value": 67108864 }] }
+    },
+    {
+      "name": "matrix_storage_m32_k1024_n8192",
+      "preset": "stress",
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [32, 1024], "dist": "randint", "seed": 1303, "min": 0, "max": 255 },
+        "b": { "dtype": "int8", "shape": [1024, 8192], "dist": "randint", "seed": 1304, "min": -127, "max": 127 },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [32, 8192] } },
+      "tunableSpace": {
+        "DP4A_LARGE_TILE_M": [16, 32],
+        "DP4A_LARGE_TILE_N": [64, 128],
+        "DP4A_LARGE_K_TILE_WORDS": [16, 32, 64]
+      },
+      "bench": { "metrics": [{ "type": "gflops", "value": 536870912 }] }
+    },
+    {
+      "name": "matrix_storage_m48_k1024_n8192",
+      "preset": "stress",
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [48, 1024], "dist": "randint", "seed": 1303, "min": 0, "max": 255 },
+        "b": { "dtype": "int8", "shape": [1024, 8192], "dist": "randint", "seed": 1304, "min": -127, "max": 127 },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
+        "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [48, 8192] } },
+      "tunableSpace": {
+        "DP4A_LARGE_TILE_M": [16, 32],
+        "DP4A_LARGE_TILE_N": [64, 128],
+        "DP4A_LARGE_K_TILE_WORDS": [16, 32, 64]
+      },
+      "bench": { "metrics": [{ "type": "gflops", "value": 805306368 }] }
+    },
+    {
+      "name": "exact-precast-padded-m127-k257-n4097-column-zp",
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [127, 257], "dist": "randint", "min": 0, "max": 255, "seed": 1307 },
+        "b": { "dtype": "int8", "shape": [257, 4097], "dist": "randint", "min": -128, "max": 127, "seed": 1308 },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "dist": "constant", "value": 127 },
+        "b_zero_point": { "dtype": "int8", "shape": [4097], "dist": "randint", "min": -8, "max": 8, "seed": 1309 }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [127, 4097] } },
+      "vars": { "M": 127, "K": 257, "N": 4097 },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
+    },
+    {
+      "name": "exact-precast-padded-m129-k1025-n2049-column-zp",
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [129, 1025], "dist": "randint", "min": 0, "max": 255, "seed": 1307 },
+        "b": { "dtype": "int8", "shape": [1025, 2049], "dist": "randint", "min": -128, "max": 127, "seed": 1308 },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "dist": "constant", "value": 127 },
+        "b_zero_point": { "dtype": "int8", "shape": [2049], "dist": "randint", "min": -8, "max": 8, "seed": 1309 }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [129, 2049] } },
+      "vars": { "M": 129, "K": 1025, "N": 2049 },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
+    },
+    {
+      "name": "exact-portable-row-reuse-m63-k1025-n8193",
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [63, 1025], "dist": "randint", "min": 0, "max": 255, "seed": 1307 },
+        "b": { "dtype": "int8", "shape": [1025, 8193], "dist": "randint", "min": -128, "max": 127, "seed": 1308 },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "dist": "constant", "value": 127 },
+        "b_zero_point": { "dtype": "int8", "shape": [8193], "dist": "randint", "min": -8, "max": 8, "seed": 1309 }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [63, 8193] } },
+      "vars": { "M": 63, "K": 1025, "N": 8193 },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
+    },
+    {
+      "name": "exact-portable-row-reuse-m64-k1025-n8193",
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [64, 1025], "dist": "randint", "min": 0, "max": 255, "seed": 1307 },
+        "b": { "dtype": "int8", "shape": [1025, 8193], "dist": "randint", "min": -128, "max": 127, "seed": 1308 },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "dist": "constant", "value": 127 },
+        "b_zero_point": { "dtype": "int8", "shape": [8193], "dist": "randint", "min": -8, "max": 8, "seed": 1309 }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 8193] } },
+      "vars": { "M": 64, "K": 1025, "N": 8193 },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
+    },
+    {
+      "name": "exact-portable-row-reuse-m65-k1025-n8193",
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [65, 1025], "dist": "randint", "min": 0, "max": 255, "seed": 1307 },
+        "b": { "dtype": "int8", "shape": [1025, 8193], "dist": "randint", "min": -128, "max": 127, "seed": 1308 },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "dist": "constant", "value": 127 },
+        "b_zero_point": { "dtype": "int8", "shape": [8193], "dist": "randint", "min": -8, "max": 8, "seed": 1309 }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [65, 8193] } },
+      "vars": { "M": 65, "K": 1025, "N": 8193 },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
     }
   ]
 }

build/webgpu/manifest.json CHANGED Viewed

The diff for this file is too large to render. See raw diff

build/webgpu/metadata.json CHANGED Viewed

@@ -1,21 +1,92 @@
 {
   "name": "ai.onnx.MatMulInteger",
-  "id": "_ai_onnx_matmulinteger_webgpu_49a7f67",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "FnGzHstTdNzZUAKkS4uTFK6rlFK9ZlpPlI3CowZMTYo=",
-      "manifest.json": "4Z43Up2znCaOeLQGxVbSRZ8lMck5cRA9Dh9wYJoEyhY=",
       "matmul-integer-batched.wgsl.jinja": "pKWtzKeL62KPieYc+I4FzUkq7JlCuZg3n7/JNjoWTHo=",
-      "quant-dp4a-matmul.wgsl.jinja": "Cn4P+iAsCchDRKrfqRJCz7fwjxB9dUBMeA5O+SaZcS0=",
       "quant-matmul-accumulate-rank2.wgsl.jinja": "S7SS4pG+oB2O1EmKrPiNg9VHkqQwJ8mrA5M68wciLE0=",
-      "quant-matmul-accumulate-rank4.wgsl.jinja": "8nSBI898d9CXoR7teRTn37WPRfYQNpHKBiB/1XD2LBo=",
-      "test.json": "ne4X/jrV8sQ610RNxaGWFTWppPB64cOp7y4KHoMljKk="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.MatMulInteger" }
 }

 {
   "name": "ai.onnx.MatMulInteger",
+  "id": "_ai_onnx_matmulinteger_webgpu_6f2ef1c",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "oZ95YN2Dm5Bwzz2Zt4UkOCHxJyia3aXTiAzZ+QnFRoo=",
+      "manifest.json": "B1e4/D6YPkD3jYTM1UmR3Ikz8VpTGxK95JNzlqEGhZI=",
       "matmul-integer-batched.wgsl.jinja": "pKWtzKeL62KPieYc+I4FzUkq7JlCuZg3n7/JNjoWTHo=",
+      "quant-dp4a-matmul.wgsl.jinja": "uOf2BCHfSx0u60XKY89imRx6cjw+Ey0V4o/sbrrSeu4=",
+      "quant-exact-matrix.wgsl.jinja": "efr11+suguVwWDMoNI3JqosWvEXa3vbMbeA0b7at14A=",
+      "quant-exact-portable.wgsl.jinja": "DaqBS74omKoDp+d77Uxv95rECoElAArVImPpDA4v5jY=",
+      "quant-exact-prepare.wgsl.jinja": "JeFBHxhxPiMqRbSJ/tfqFsgb4T5o3q3/RZGPcHDn6rk=",
       "quant-matmul-accumulate-rank2.wgsl.jinja": "S7SS4pG+oB2O1EmKrPiNg9VHkqQwJ8mrA5M68wciLE0=",
+      "quant-matmul-accumulate-rank4.wgsl.jinja": "JBwbT3KMpqRNRXwjkCBPkpRCjLd0BfOn5hAUBm3b33o=",
+      "test.json": "RFp6aUPoekPtgawzlEd6e+iSWtmo0UsVU+PLYdU92z4="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "sgmat_precast_a0_bnone_f16": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a0_bnone_f16": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a0_bnone": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a0_bnone": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a0_bscalar_f16": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a0_bscalar_f16": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a0_bscalar": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a0_bscalar": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a0_bcolumn_f16": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a0_bcolumn_f16": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a0_bcolumn": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a0_bcolumn": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a1_bnone_f16": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a1_bnone_f16": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a1_bnone": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a1_bnone": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a1_bscalar_f16": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a1_bscalar_f16": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a1_bscalar": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a1_bscalar": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a1_bcolumn_f16": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a1_bcolumn_f16": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_precast_a1_bcolumn": ["quant-exact-matrix.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "portable_precast_a1_bcolumn": ["quant-exact-portable.wgsl.jinja", "quant-exact-prepare.wgsl.jinja"],
+      "sgmat_exact_a0_bnone": ["quant-exact-matrix.wgsl.jinja"],
+      "sgmat_exact_a0_bscalar": ["quant-exact-matrix.wgsl.jinja"],
+      "sgmat_exact_a0_bcolumn": ["quant-exact-matrix.wgsl.jinja"],
+      "sgmat_exact_a1_bnone": ["quant-exact-matrix.wgsl.jinja"],
+      "sgmat_exact_a1_bscalar": ["quant-exact-matrix.wgsl.jinja"],
+      "sgmat_exact_a1_bcolumn": ["quant-exact-matrix.wgsl.jinja"],
+      "dp4a_rank2_scalar_zero_points_tail": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_rank2_no_zero_points_tail": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_rank3_by_rank2_scalar_zero_points_tail": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_rank2_scalar_zero_points": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_rank2_a_zero_point_only": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_rank2_b_zero_point_only": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_rank2_no_zero_points": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_rank2_b_zero_point_per_column": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_rank2_a_zero_point_b_per_column": ["quant-dp4a-matmul.wgsl.jinja"],
+      "rank2_scalar_zero_points": ["quant-matmul-accumulate-rank2.wgsl.jinja"],
+      "rank2_a_zero_point_only": ["quant-matmul-accumulate-rank2.wgsl.jinja"],
+      "rank2_b_zero_point_only": ["quant-matmul-accumulate-rank2.wgsl.jinja"],
+      "rank1_dot_b_zero_point_only": ["quant-matmul-accumulate-rank2.wgsl.jinja"],
+      "rank1_dot_no_zero_points": ["quant-matmul-accumulate-rank2.wgsl.jinja"],
+      "rank2_no_zero_points": ["quant-matmul-accumulate-rank2.wgsl.jinja"],
+      "rank2_b_zero_point_per_column": ["quant-matmul-accumulate-rank2.wgsl.jinja"],
+      "rank2_b_zero_point_per_column_no_azp": ["quant-matmul-accumulate-rank2.wgsl.jinja"],
+      "rank3_by_rank3_scalar_zero_points": ["matmul-integer-batched.wgsl.jinja"],
+      "rank3_by_rank3_no_zero_points": ["matmul-integer-batched.wgsl.jinja"],
+      "rank3_by_rank3_a_zero_point_only": ["matmul-integer-batched.wgsl.jinja"],
+      "rank3_by_rank3_b_zero_point_only": ["matmul-integer-batched.wgsl.jinja"],
+      "rank3_by_rank3_b_zero_point_per_column_batched": ["matmul-integer-batched.wgsl.jinja"],
+      "rank3_by_rank3_b_zero_point_per_column_batched_no_azp": ["matmul-integer-batched.wgsl.jinja"],
+      "rank2_by_rank3_scalar_zero_points": ["matmul-integer-batched.wgsl.jinja"],
+      "rank2_by_rank3_a_zero_point_only": ["matmul-integer-batched.wgsl.jinja"],
+      "rank2_by_rank3_b_zero_point_only": ["matmul-integer-batched.wgsl.jinja"],
+      "rank2_by_rank3_no_zero_points": ["matmul-integer-batched.wgsl.jinja"],
+      "rank4_scalar_zero_points": ["quant-matmul-accumulate-rank4.wgsl.jinja"],
+      "rank4_a_zero_point_only": ["quant-matmul-accumulate-rank4.wgsl.jinja"],
+      "rank4_b_zero_point_only": ["quant-matmul-accumulate-rank4.wgsl.jinja"],
+      "rank4_no_zero_points": ["quant-matmul-accumulate-rank4.wgsl.jinja"],
+      "dp4a_rank3_by_rank3_scalar_zero_points": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_rank4_scalar_zero_points": ["quant-dp4a-matmul.wgsl.jinja"],
+      "dp4a_rank4_no_zero_points": ["quant-dp4a-matmul.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/quant-dp4a-matmul.wgsl.jinja CHANGED Viewed

@@ -1,8 +1,25 @@
-{% if requant != "none" and scaleScalar == "f16" %}
-enable f16;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const TILE_M: u32 = {{ tileM }}u;
 const TILE_N: u32 = {{ tileN }}u;
 const KT_WORDS: u32 = {{ ktWords }}u;
@@ -92,12 +109,6 @@ fn read_b_zero({% if bZeroPerColumn %}col: u32{% endif %}) -> i32 {
 {% endif %}
 }
-{% endif %}
-{% if requant != "none" %}
-fn read_y_zero() -> i32 {
-  return y_zero_point[0];
-}
 {% endif %}
 @compute @workgroup_size(TILE_N, 1, 1)
 fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
@@ -303,16 +314,7 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid:
       acc = acc + ca * {% if tileM == 1 %}sum_b_total{% else %}sum_b{% endif %};
 {% endif %}
       acc = acc + kcc;
-{% if requant == "none" %}
       y[y_base + mr * params.N + col] = acc;
-{% else %}
-      let scaled = f32(acc) * f32(a_scale[0]) * f32(b_scale[{% if requant == "per_column" %}col{% elif requant == "per_row" %}mr{% else %}0u{% endif %}]) / f32(y_scale[0]);
-      // Clamp the rounded float into an i32-representable range before the cast
-      // so a huge-finite / Inf requantized value saturates (the final clamp pins
-      // it to [qMin,qMax]) instead of invoking undefined i32(huge_float).
-      let q = clamp(i32(clamp(round(scaled), -2.0e9, 2.0e9)) + read_y_zero(), {{ qMin }}, {{ qMax }});
-      y[y_base + mr * params.N + col] = q;
-{% endif %}
     }
   }
 {% endfor %}

+{% if not env.device.wgslLanguageFeatures.has("packed_4x8_integer_dot_product") %}
+// Runtime source-compatible fallbacks for the packed 4x8 integer built-ins used
+// by quantized kernels. Native implementations remain unshadowed when the
+// language feature is present. The user functions retain the built-ins'
+// @must_use contract, but cannot be @const; every call is runtime.
+// Inputs are function parameters, so replacing a built-in call with one of
+// these same-name functions preserves WGSL's single evaluation of each
+// argument. Byte extraction is unsigned. Signed lanes use the exact
+// i32(byte ^ 0x80u) - 128i sign extension, avoiding signed-left-shift overflow
+// while preserving every 8-bit two's-complement value.
+@must_use fn dot4I8Packed(e1: u32, e2: u32) -> i32 {
+  let shifts = vec4<u32>(0u, 8u, 16u, 24u);
+  let a_bytes = (vec4<u32>(e1) >> shifts) & vec4<u32>(0xffu);
+  let b_bytes = (vec4<u32>(e2) >> shifts) & vec4<u32>(0xffu);
+  let a = vec4<i32>(a_bytes ^ vec4<u32>(0x80u)) - vec4<i32>(128i);
+  let b = vec4<i32>(b_bytes ^ vec4<u32>(0x80u)) - vec4<i32>(128i);
+  return dot(a, b);
+}
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const TILE_M: u32 = {{ tileM }}u;
 const TILE_N: u32 = {{ tileN }}u;
 const KT_WORDS: u32 = {{ ktWords }}u;
 {% endif %}
 }
 {% endif %}
 @compute @workgroup_size(TILE_N, 1, 1)
 fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
       acc = acc + ca * {% if tileM == 1 %}sum_b_total{% else %}sum_b{% endif %};
 {% endif %}
       acc = acc + kcc;
       y[y_base + mr * params.N + col] = acc;
     }
   }
 {% endfor %}

build/webgpu/quant-exact-matrix.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,126 @@

+{% set precastOperands = precastOperands | default(false) %}
+{% set precastScalar = precastScalar | default("f32") %}
+{% if (precastOperands and precastScalar == "f16") or (requantize and scaleScalar == "f16") %}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 }}
+// Adjusted byte integers fit exactly in f32. Each bounded f32 dot product
+// has absolute magnitude at most 2^24; int32 adds then preserve wraparound.
+const TILE_M: u32 = {{ sgmatTileM }}u;
+const TILE_N: u32 = {{ sgmatTileN }}u;
+const TILE_K: u32 = {{ sgmatTileK }}u;
+{% if precastOperands %}
+const PADDED_K: u32 = {{ preparedK }}u;
+const PADDED_N: u32 = {{ preparedN }}u;
+var<workgroup> matrix_out: array<f32, {{ sgmatSubgroups * 8 * 64 }}>;
+{% else %}
+var<workgroup> tile_A: array<f32, {{ sgmatTileM * sgmatTileK }}>;
+var<workgroup> tile_B: array<f32, {{ sgmatTileK * sgmatTileN }}>;
+{% if not sgmatReuseStaging %}
+var<workgroup> scratch: array<array<array<f32, 64>, 4>, {{ sgmatSubgroups }}>;
+{% endif %}
+{% endif %}
+@compute @workgroup_size({{ sgmatWorkgroup }}){{ " @subgroup_size(32)" if pinSubgroupSize32 else "" }}
+fn main(@builtin(workgroup_id) wid: vec3<u32>,
+        @builtin(local_invocation_index) tid: u32,
+        @builtin(subgroup_invocation_id) lane: u32,
+        @builtin(subgroup_size) subgroup_size: u32) {
+  let subgroup = tid / subgroup_size;
+  let base_A = (subgroup % {{ sgmatRowGroups }}u) * 16u;
+  let base_B = (subgroup / {{ sgmatRowGroups }}u) * 32u;
+  let row = lane / 4u;
+  let col = (lane % 4u) * 2u;
+  // Bound adjusted byte magnitudes from scalar zero points. A column-wise
+  // B zero point keeps the full byte-difference bound for the whole tile.
+  let max_a = {% if hasAZero %}u32(max(abs({{ 0 if aUnsigned else (0 - 128) }} - i32(a_zero_point[0])), abs({{ 255 if aUnsigned else 127 }} - i32(a_zero_point[0])))){% else %}{{ 255 if aUnsigned else 128 }}u{% endif %};
+  let max_b = {% if hasBZero and not bZeroPerColumn %}u32(max(abs({{ 0 if bUnsigned else (0 - 128) }} - i32(b_zero_point[0])), abs({{ 255 if bUnsigned else 127 }} - i32(b_zero_point[0])))){% else %}{{ 255 if hasBZero or bUnsigned else 128 }}u{% endif %};
+  let partial_k = (16777216u / (max_a * max_b) / TILE_K) * TILE_K;
+  var totals: array<i32, 16>;
+  for (var partial = 0u; partial < params.K; partial += partial_k) {
+{% for r in range(2) %}{% for c in range(4) %}
+    var matC{{ r }}{{ c }}: subgroup_matrix_result<f32, 8, 8>;
+{% endfor %}{% endfor %}
+    let partial_end = min(params.K, partial + partial_k);
+    for (var kb = partial; kb < partial_end; kb += TILE_K) {
+{% if not precastOperands %}
+      for (var index = tid; index < TILE_M * TILE_K; index += {{ sgmatWorkgroup }}u) {
+        let m = wid.y * TILE_M + index / TILE_K;
+        let k = kb + index % TILE_K;
+        var value = f32(0.0);
+        if (m < params.M && k < params.K) {
+          value = f32(i32(a[m * params.K + k]){% if hasAZero %} - i32(a_zero_point[0]){% endif %});
+        }
+        tile_A[index] = value;
+      }
+      for (var index = tid; index < TILE_K * TILE_N; index += {{ sgmatWorkgroup }}u) {
+        let k = kb + index / TILE_N;
+        let n = wid.x * TILE_N + index % TILE_N;
+        var value = f32(0.0);
+        if (k < params.K && n < params.N) {
+          value = f32(i32(b[k * params.N + n]){% if hasBZero %} - i32(b_zero_point[{{ "n" if bZeroPerColumn else "0" }}]){% endif %});
+        }
+        tile_B[index] = value;
+      }
+      workgroupBarrier();
+{% endif %}
+      for (var step = 0u; step < TILE_K; step += 8u) {
+{% for r in range(2) %}
+        let matA{{ r }} = subgroupMatrixLoad<subgroup_matrix_left<{{ precastScalar if precastOperands else "f32" }}, 8, 8>, row_major>(&{{ "a" if precastOperands else "tile_A" }}, {% if precastOperands %}(wid.y * TILE_M + base_A + {{ r * 8 }}u) * PADDED_K + kb + step{% else %}(base_A + {{ r * 8 }}u) * TILE_K + step{% endif %}, {{ "PADDED_K" if precastOperands else "TILE_K" }});
+{% endfor %}
+{% for c in range(4) %}
+        let matB{{ c }} = subgroupMatrixLoad<subgroup_matrix_right<{{ precastScalar if precastOperands else "f32" }}, 8, 8>, row_major>(&{{ "b" if precastOperands else "tile_B" }}, {% if precastOperands %}(kb + step) * PADDED_N + wid.x * TILE_N + base_B + {{ c * 8 }}u{% else %}step * TILE_N + base_B + {{ c * 8 }}u{% endif %}, {{ "PADDED_N" if precastOperands else "TILE_N" }});
+{% endfor %}
+{% for r in range(2) %}{% for c in range(4) %}
+        matC{{ r }}{{ c }} = subgroupMatrixMultiplyAccumulate(matA{{ r }}, matB{{ c }}, matC{{ r }}{{ c }});
+{% endfor %}{% endfor %}
+      }
+{% if not precastOperands %}
+      workgroupBarrier();
+{% endif %}
+    }
+{% if precastOperands %}
+    // The prepared operands need no staging barriers. Publish all eight
+    // matrices together before converting the bounded partial to int32.
+{% for r in range(2) %}{% for c in range(4) %}
+    subgroupMatrixStore<row_major>(&matrix_out, subgroup * 512u + {{ r * 256 + c * 64 }}u, matC{{ r }}{{ c }}, 8u);
+{% endfor %}{% endfor %}
+    workgroupBarrier();
+{% for r in range(2) %}{% for c in range(4) %}{% for h in range(2) %}
+    totals[{{ r * 8 + c * 2 + h }}] += i32(matrix_out[subgroup * 512u + {{ r * 256 + c * 64 }}u + row * 8u + col + {{ h }}u]);
+{% endfor %}{% endfor %}{% endfor %}
+    workgroupBarrier();
+{% else %}
+{% if sgmatReuseStaging %}
+    // All staged A reads have completed. Reuse its storage for one result band
+    // at a time: each subgroup needs 4 * 64 floats, exactly its A-tile share.
+{% endif %}
+{% for r in range(2) %}
+{% for c in range(4) %}
+    subgroupMatrixStore<row_major>({% if sgmatReuseStaging %}&tile_A{% else %}&scratch[subgroup][{{ c }}]{% endif %}, {% if sgmatReuseStaging %}subgroup * 256u + {{ c * 64 }}u{% else %}0u{% endif %}, matC{{ r }}{{ c }}, 8u);
+{% endfor %}
+    workgroupBarrier();
+{% for c in range(4) %}{% for h in range(2) %}
+    totals[{{ r * 8 + c * 2 + h }}] += i32({% if sgmatReuseStaging %}tile_A[subgroup * 256u + {{ c * 64 }}u + row * 8u + col + {{ h }}u]{% else %}scratch[subgroup][{{ c }}][row * 8u + col + {{ h }}u]{% endif %});
+{% endfor %}{% endfor %}
+    workgroupBarrier();
+{% endfor %}
+{% endif %}
+  }
+{% for r in range(2) %}{% for c in range(4) %}{% for h in range(2) %}
+  {
+    let m = wid.y * TILE_M + base_A + {{ r * 8 }}u + row;
+    let n = wid.x * TILE_N + base_B + {{ c * 8 + h }}u + col;
+    if (m < params.M && n < params.N) {
+      y[m * params.N + n] = totals[{{ r * 8 + c * 2 + h }}];
+    }
+  }
+{% endfor %}{% endfor %}{% endfor %}
+}

build/webgpu/quant-exact-portable.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,143 @@

+{% if precastScalar == "f16" or (requantize and scaleScalar == "f16") %}enable f16;{% endif %}
+{% set exactIntegerPartials = true %}
+// Both integer-output and requantizing consumers share the vector GEMM body.
+// Its callback stores scalar outputs so padded operands also cover column tails.
+fn store_integer_result(value: i32, output_index: u32) {
+  y[output_index] = value;
+}
+{% set exactIntegerPartialsValue = exactIntegerPartials is defined and exactIntegerPartials %}
+{% set preparedNValue = preparedN | default(0) %}
+{% set integerPartialKValue = integerPartialK | default(0) %}
+{{ env.wgsl.resourceDeclarations }}
+{% set kTileValue = kTile %}
+{% set blockRowsValue = blockRows if blockRows is defined else 128 %}
+{% set blockColsValue = blockCols if blockCols is defined else 64 %}
+{% set accumulatorScalarValue = accumulatorScalar if accumulatorScalar is defined else "f32" %}
+{% set workgroupColsValue = workgroupCols if workgroupCols is defined else 16 %}
+{% set workgroupRowsValue = workgroupRows if workgroupRows is defined else 16 %}
+{% set laneCount = workgroupColsValue * workgroupRowsValue %}
+{% set threadRows = (blockRowsValue / workgroupRowsValue)|int %}
+{% set threadColVectors = (blockColsValue / (workgroupColsValue * 4))|int %}
+{% set threadCols = threadColVectors * 4 %}
+{% set aVectorCount = (blockRowsValue * kTileValue / (4 * laneCount))|int %}
+{% set bVectorCount = (kTileValue * blockColsValue / (4 * laneCount))|int %}
+{% set batchStridesValue = batchStrides if batchStrides is defined else false %}
+// High-intensity matrix multiplication for the large, untransposed layout:
+//   Y[M,N] = alpha * A[M,K] @ B[K,N].
+// B may alternatively be stored physically as [N,K]; that mode retains vec4
+// K loads and transposes only while staging the tile in workgroup memory.
+//
+// Both inputs and the output are four-wide along their contiguous dimension.
+// The output-tile geometry and BK are selected from the shape and device's
+// workgroup-storage limit. f16 operands stay packed in global and workgroup
+// memory, then widen at the FMA so accumulation remains f32.
+const M: u32 = {{ M }}u;
+const K: u32 = {{ K }}u;
+const N: u32 = {{ N }}u;
+const K_V4: u32 = K / 4u;
+const B_N_V4: u32 = {{ preparedNValue }}u / 4u;
+const BK: u32 = {{ kTileValue }}u;
+const BM: u32 = {{ blockRowsValue }}u;
+const BN: u32 = {{ blockColsValue }}u;
+const TM: u32 = {{ threadRows }}u;
+const TN: u32 = {{ threadCols }}u;
+// A is loaded and consumed four K values at a time. Keeping that packing in
+// workgroup memory reduces A-side shared-memory instructions and inner-loop
+// control without changing the four scalar FMA operations or their order.
+var<workgroup> tileA: array<array<vec4<{{ scalar }}>, BK / 4u>, BM>;
+var<workgroup> tileB: array<vec4<{{ scalar }}>, BK * (BN / 4u)>;
+@compute @workgroup_size({{ workgroupColsValue }}, {{ workgroupRowsValue }}, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let li = lid.y * {{ workgroupColsValue }}u + lid.x;
+  let mBase = wg.y * BM;
+  let nBase = wg.x * BN;
+{% set aOff = "aBatch + " if batchStridesValue else "" %}
+{% set bOff = "bBatch + " if batchStridesValue else "" %}
+{% for i in range(threadRows) %}
+{% for j in range(threadColVectors) %}
+  var acc{{ i }}_{{ j }} = vec4<{{ accumulatorScalarValue }}>(0.0);
+  var total{{ i }}_{{ j }} = vec4<i32>(0);
+{% endfor %}
+{% endfor %}
+  for (var kBase: u32 = 0u; kBase < K; kBase = kBase + BK) {
+{% for av in range(aVectorCount) %}
+    let aVector{{ av }} = li + {{ av * laneCount }}u;
+    let aRow{{ av }} = aVector{{ av }} / (BK / 4u);
+    let aColV4_{{ av }} = aVector{{ av }} % (BK / 4u);
+    let am{{ av }} = mBase + aRow{{ av }};
+    var av4_{{ av }} = vec4<{{ scalar }}>({{ scalar }}(0));
+    if (am{{ av }} < M) {
+      av4_{{ av }} = a[{{ aOff }}am{{ av }} * K_V4 + kBase / 4u + aColV4_{{ av }}];
+    }
+    tileA[aRow{{ av }}][aColV4_{{ av }}] = av4_{{ av }};
+{% endfor %}
+{% for bv in range(bVectorCount) %}
+    let bVector{{ bv }} = li + {{ bv * laneCount }}u;
+    let bRow{{ bv }} = bVector{{ bv }} / (BN / 4u);
+    let bColV4_{{ bv }} = bVector{{ bv }} % (BN / 4u);
+    let bn{{ bv }} = nBase + bColV4_{{ bv }} * 4u;
+    var bv4_{{ bv }} = vec4<{{ scalar }}>({{ scalar }}(0));
+    if (bn{{ bv }} < N) {
+      bv4_{{ bv }} = b[{{ bOff }}(kBase + bRow{{ bv }}) * {{ "B_N_V4" if exactIntegerPartialsValue else "N_V4" }} + bn{{ bv }} / 4u];
+    }
+    tileB[bRow{{ bv }} * (BN / 4u) + bColV4_{{ bv }}] = bv4_{{ bv }};
+{% endfor %}
+    workgroupBarrier();
+    let aRow = lid.y * TM;
+    let bCol4 = lid.x * {{ threadColVectors }}u;
+    for (var kv: u32 = 0u; kv < BK / 4u; kv = kv + 1u) {
+      let kk = kv * 4u;
+{% for j in range(threadColVectors) %}
+      let bv0_{{ j }} = vec4<{{ accumulatorScalarValue }}>(tileB[(kk + 0u) * (BN / 4u) + bCol4 + {{ j }}u]);
+      let bv1_{{ j }} = vec4<{{ accumulatorScalarValue }}>(tileB[(kk + 1u) * (BN / 4u) + bCol4 + {{ j }}u]);
+      let bv2_{{ j }} = vec4<{{ accumulatorScalarValue }}>(tileB[(kk + 2u) * (BN / 4u) + bCol4 + {{ j }}u]);
+      let bv3_{{ j }} = vec4<{{ accumulatorScalarValue }}>(tileB[(kk + 3u) * (BN / 4u) + bCol4 + {{ j }}u]);
+{% endfor %}
+{% for i in range(threadRows) %}
+      let av{{ i }} = tileA[aRow + {{ i }}u][kv];
+{% for j in range(threadColVectors) %}
+      acc{{ i }}_{{ j }} = acc{{ i }}_{{ j }} + vec4<{{ accumulatorScalarValue }}>({{ accumulatorScalarValue }}(av{{ i }}.x)) * bv0_{{ j }};
+      acc{{ i }}_{{ j }} = acc{{ i }}_{{ j }} + vec4<{{ accumulatorScalarValue }}>({{ accumulatorScalarValue }}(av{{ i }}.y)) * bv1_{{ j }};
+      acc{{ i }}_{{ j }} = acc{{ i }}_{{ j }} + vec4<{{ accumulatorScalarValue }}>({{ accumulatorScalarValue }}(av{{ i }}.z)) * bv2_{{ j }};
+      acc{{ i }}_{{ j }} = acc{{ i }}_{{ j }} + vec4<{{ accumulatorScalarValue }}>({{ accumulatorScalarValue }}(av{{ i }}.w)) * bv3_{{ j }};
+{% endfor %}
+{% endfor %}
+    }
+    workgroupBarrier();
+    // Adjusted byte integers and each bounded partial are exact in f32.
+    // Convert before exceeding that bound; int32 additions retain wraparound.
+    if ((kBase + BK) % {{ integerPartialKValue }}u == 0u || kBase + BK == K) {
+{% for i in range(threadRows) %}{% for j in range(threadColVectors) %}
+      total{{ i }}_{{ j }} += vec4<i32>(acc{{ i }}_{{ j }});
+      acc{{ i }}_{{ j }} = vec4<f32>(0.0);
+{% endfor %}{% endfor %}
+    }
+  }
+  let m0 = mBase + lid.y * TM;
+{% for j in range(threadColVectors) %}
+  let n{{ j }} = nBase + lid.x * TN + {{ j * 4 }}u;
+{% endfor %}
+{% for i in range(threadRows) %}
+{% for j in range(threadColVectors) %}
+  if (m0 + {{ i }}u < M && n{{ j }} < N) {
+{% for h in range(4) %}
+    if (n{{ j }} + {{ h }}u < N) {
+      store_integer_result(total{{ i }}_{{ j }}[{{ h }}], (m0 + {{ i }}u) * N + n{{ j }} + {{ h }}u);
+    }
+{% endfor %}
+  }
+{% endfor %}
+{% endfor %}
+}

build/webgpu/quant-exact-prepare.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,33 @@

+{% if precastScalar == "f16" %}enable f16;{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+@compute @workgroup_size({{ prepareWorkgroup }})
+fn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) tid: u32) {
+  let offset = ((wid.y * params.groupsX + wid.x) * {{ prepareWorkgroup }}u + tid) * 4u;
+  if (offset >= params.count) { return; }
+  if (offset < params.aCount) {
+    let m = offset / params.paddedK;
+    let k = offset % params.paddedK;
+{% for h in range(4) %}
+    {
+      var value = 0.0;
+      if (m < params.M && k + {{ h }}u < params.K) {
+        value = f32(i32(a[m * params.K + k + {{ h }}u]){% if hasAZero %} - i32(a_zero_point[0]){% endif %});
+      }
+      aFloat[offset + {{ h }}u] = {% if precastScalar == "f16" %}f16(value){% else %}value{% endif %};
+    }
+{% endfor %}
+  } else {
+    let bOffset = offset - params.aCount;
+    let k = bOffset / params.paddedN;
+    let n = bOffset % params.paddedN;
+{% for h in range(4) %}
+    {
+      var value = 0.0;
+      if (k < params.K && n + {{ h }}u < params.N) {
+        value = f32(i32(b[k * params.N + n + {{ h }}u]){% if hasBZero %} - i32(b_zero_point[{% if bZeroPerColumn %}n + {{ h }}u{% else %}0{% endif %}]){% endif %});
+      }
+      bFloat[bOffset + {{ h }}u] = {% if precastScalar == "f16" %}f16(value){% else %}value{% endif %};
+    }
+{% endfor %}
+  }
+}

build/webgpu/quant-matmul-accumulate-rank4.wgsl.jinja CHANGED Viewed

@@ -30,7 +30,7 @@ fn {{ fn_name }}() -> i32 {
 {%- endmacro %}
-{{ quant_read_scalar_zero("read_a_zero", "a_zero_point", aUnsigned, true if source.outputToScratch else hasAZero) }}
 {{ quant_read_scalar_zero("read_b_zero", "b_zero_point", bUnsigned, hasBZero) }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE_X }}, {{ tunables.WORKGROUP_SIZE_Y }})

 {%- endmacro %}
+{{ quant_read_scalar_zero("read_a_zero", "a_zero_point", aUnsigned, true if outputToScratch else hasAZero) }}
 {{ quant_read_scalar_zero("read_b_zero", "b_zero_point", bUnsigned, hasBZero) }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE_X }}, {{ tunables.WORKGROUP_SIZE_Y }})

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.MatMulInteger",
   "fixtureArrays": {
     "dp4a_s8_pattern": [-128, 127, 0, -1, 56, -77, 3, 100, -100, 42, -5, 19, -64, 88, -33, 7, 125, -90],
     "dp4a_u8_pattern": [0, 255, 128, 7, 200, 31, 129, 64, 99, 3, 254, 127, 1, 88, 173, 240, 16, 5],
@@ -642,7 +641,7 @@
         }
       },
       "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
-      "tunables": { "DP4A_SHORT_K_MIN_MN": 64 }
     },
     {
       "name": "dp4a_s8s8_medium_k_m64_k96_n128",
@@ -662,7 +661,12 @@
         }
       },
       "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
-      "tunables": { "DP4A_SHORT_K_MAX": 64, "DP4A_MEDIUM_K_MAX": 128, "DP4A_MEDIUM_K_MIN_MN": 64 }
     },
     {
       "name": "dp4a_u8s8_k256_a_zero_point_only_large",
@@ -679,7 +683,8 @@
         },
         "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [131] } }
       },
-      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } }
     },
     {
       "name": "dp4a_s8u8_k256_b_zero_point_only_large",
@@ -696,7 +701,8 @@
         },
         "b_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [200] } }
       },
-      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } }
     },
     {
       "name": "dp4a_s8s8_k260_both_zero_points_large",
@@ -718,7 +724,12 @@
         "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-7] } }
       },
       "outputs": { "y": { "dtype": "int32", "shape": [65, 129], "tolerance": 0 } },
-      "tunables": { "DP4A_LARGE_MIN_M": 64, "DP4A_LARGE_MIN_N": 128, "DP4A_LARGE_MIN_K": 256 }
     },
     {
       "name": "dp4a_rank3_by_rank3_u8s8_k256_large",
@@ -932,7 +943,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/matmul_integer_test.cc",
         "test": "MatmulIntegerOpTest.MatMulInteger_int8_uint8_A_ND_B_ND",
-        "notes": "Compact 4D MatMul broadcast extension of ORT's ND MatMulInteger coverage."
       },
       "inputs": {
         "a": {
@@ -988,7 +999,7 @@
     {
       "name": "dispatch_cliff_mtile_rows_524296",
       "provenance": {
-        "notes": "M=524296 rows -> M-tiles=ceil(524296/8)=65537 > 65535. dp4a_rank2 is capped (ceil(M/8)<=65535) so this falls through to the folded matmul-integer rank2 (gid.y + gid.z*nwg.y*WORKGROUP_SIZE_Y at nwg.z=2). Validates both the fold and the dp4a cap-fallthrough; before the fold the plan guard throws."
       },
       "inputs": {
         "a": { "dtype": "uint8", "shape": [524296, 4], "data": { "kind": "cycle", "values": [1, 2, 3, 4, 5, 6, 7] } },
@@ -1022,7 +1033,8 @@
           "data": { "kind": "cycle", "values": [1, -2, 2, -1, 3, -3, 0, -4, 4, -5, 5] }
         }
       },
-      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } }
     },
     {
       "name": "percol_b_zero_point_no_azp_below_dp4a_floor",
@@ -1068,7 +1080,8 @@
           "data": { "kind": "cycle", "values": [1, -2, 2, -1, 3, -3, 0, -4, 4, -5, 5] }
         }
       },
-      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } }
     },
     {
       "name": "k258_unaligned_k_scalar_zp_u8u8",
@@ -1089,12 +1102,13 @@
         "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
         "b_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } }
       },
-      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } }
     },
     {
       "name": "large_accumulator_dp4a_k1024_u8u8_maxmag",
       "provenance": {
-        "notes": "Confirms the dp4a path and the naive path agree at large int32 accumulation. M=64,K=1024,N=128 -> numel(Y)=8192 and K%4==0 and K>=256, so this selects the dp4a (dot4I8Packed) kernel. Inputs include the max-magnitude bytes (255 for u8) so per-row accumulation reaches ~1024*255*255~=66.6M, well within int32, exercising the bias/zero-point fold (ca,cb,kcc) in quant-dp4a-matmul at large magnitude. ORT/CPU validates."
       },
       "inputs": {
         "a": {
@@ -1110,7 +1124,8 @@
         "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [3] } },
         "b_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [7] } }
       },
-      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } }
     },
     {
       "name": "dp4a_rank4_scalar_s8s8_2x2x64x128",
@@ -1149,7 +1164,7 @@
     {
       "name": "dp4a_gemv_m1_u8s8_k257_n512_scalar_zp_tail",
       "provenance": {
-        "notes": "M=1 int8 decode GEMV with a one-lane K tail routes to the tileM=1 ILP DP4A config; validates its unrolled full-word loop and scalar tail against the reference."
       },
       "inputs": {
         "a": {
@@ -1224,7 +1239,8 @@
           "data": { "kind": "cycle", "values": [1, -2, 2, -1, 3, -3, 0, -4, 4, -5, 5] }
         }
       },
-      "outputs": { "y": { "dtype": "int32", "shape": [64, 130], "tolerance": 0 } }
     },
     {
       "name": "dp4a_partial_m63_n130_tile_edge_s8s8_k256",
@@ -1282,7 +1298,8 @@
           "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/dp4a_s8_pattern" } }
         }
       },
-      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } }
     },
     {
       "name": "dp4a_tail_k130_rank2_u8u8_scalar_zp_exact",
@@ -1309,7 +1326,8 @@
         "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [121] } },
         "b_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [117] } }
       },
-      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } }
     },
     {
       "name": "dp4a_tail_k131_rank2_s8s8_scalar_zp_exact",
@@ -1333,7 +1351,8 @@
         "a_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [5] } },
         "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-7] } }
       },
-      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } }
     },
     {
       "name": "dp4a_tail_k129_batched3_by_rank2_exact",
@@ -1361,9 +1380,7 @@
     },
     {
       "name": "dp4a_tail_k63_rank2_u8s8_scalar_zp_exact",
-      "provenance": {
-        "notes": "Exactness and direct-route lock for the 32<=K<128 packed-tail range opened by the K63 projection pathology fix."
-      },
       "inputs": {
         "a": { "dtype": "uint8", "shape": [16, 63], "data": { "kind": "cycle", "values": [0, 255, 121, 7, 200, 31] } },
         "b": { "dtype": "int8", "shape": [63, 512], "data": { "kind": "cycle", "values": [-128, 127, 0, -1, 56, -77] } },
@@ -1495,7 +1512,7 @@
     {
       "name": "dp4a_gemv_no_zero_s8s8_m1_k64_n512",
       "provenance": {
-        "notes": "M=1 selects the dp4a GEMV tile (tileM 1) and signed inputs with both zero points omitted make the B-sum correction unnecessary, which is the only combination that renders the GEMV loop without its running row sum. N=512 and K=64 clear the route's minimum-work guards."
       },
       "inputs": {
         "a": {
@@ -1517,7 +1534,7 @@
     {
       "name": "dp4a_gemv_b_per_column_zero_u8u8_m1_k64_n512",
       "provenance": {
-        "notes": "A per-column B zero point on an unsigned B, which is the only combination that renders the widening read in the per-column accessor; every other per-column case had a signed B. The zero point varies per column on a period coprime with N so a column-index error changes the result."
       },
       "inputs": {
         "a": {
@@ -1541,7 +1558,7 @@
     {
       "name": "dp4a_tail_no_zero_s8s8_m1_k66_n512",
       "provenance": {
-        "notes": "K=66 leaves a two-lane K tail, and signed inputs with both zero points omitted drop both running sums, which is the only combination that renders the tail block without them. The tail route previously existed only with both zero points present."
       },
       "inputs": {
         "a": {
@@ -1559,6 +1576,696 @@
         }
       },
       "outputs": { "y": { "dtype": "int32", "shape": [1, 512], "tolerance": 0 } }
     }
   ]
 }

 {
   "fixtureArrays": {
     "dp4a_s8_pattern": [-128, 127, 0, -1, 56, -77, 3, 100, -100, 42, -5, 19, -64, 88, -33, 7, 125, -90],
     "dp4a_u8_pattern": [0, 255, 128, 7, 200, 31, 129, 64, 99, 3, 254, 127, 1, 88, 173, 240, 16, 5],
         }
       },
       "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": { "DP4A_SHORT_K_MIN_MN": 64, "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "dp4a_s8s8_medium_k_m64_k96_n128",
         }
       },
       "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": {
+        "DP4A_SHORT_K_MAX": 64,
+        "DP4A_MEDIUM_K_MAX": 128,
+        "DP4A_MEDIUM_K_MIN_MN": 64,
+        "EXACT_MATRIX_MIN_INVOCATIONS": 0
+      }
     },
     {
       "name": "dp4a_u8s8_k256_a_zero_point_only_large",
         },
         "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [131] } }
       },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "dp4a_s8u8_k256_b_zero_point_only_large",
         },
         "b_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [200] } }
       },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "dp4a_s8s8_k260_both_zero_points_large",
         "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-7] } }
       },
       "outputs": { "y": { "dtype": "int32", "shape": [65, 129], "tolerance": 0 } },
+      "tunables": {
+        "DP4A_LARGE_MIN_M": 64,
+        "DP4A_LARGE_MIN_N": 128,
+        "DP4A_LARGE_MIN_K": 256,
+        "EXACT_MATRIX_MIN_INVOCATIONS": 0
+      }
     },
     {
       "name": "dp4a_rank3_by_rank3_u8s8_k256_large",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/matmul_integer_test.cc",
         "test": "MatmulIntegerOpTest.MatMulInteger_int8_uint8_A_ND_B_ND",
+        "notes": "A compact four-dimensional uint8 MatMulInteger exercises multidimensional batch broadcasting."
       },
       "inputs": {
         "a": {
     {
       "name": "dispatch_cliff_mtile_rows_524296",
       "provenance": {
+        "notes": "M=524,296 produces 65,537 eight-row tiles, exceeding the dp4a rank-2 dispatch cap and selecting a folded y=65,535, z=2 integer MatMul. Distinct rows make incorrect y/z reconstruction observable."
       },
       "inputs": {
         "a": { "dtype": "uint8", "shape": [524296, 4], "data": { "kind": "cycle", "values": [1, 2, 3, 4, 5, 6, 7] } },
           "data": { "kind": "cycle", "values": [1, -2, 2, -1, 3, -3, 0, -4, 4, -5, 5] }
         }
       },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "percol_b_zero_point_no_azp_below_dp4a_floor",
           "data": { "kind": "cycle", "values": [1, -2, 2, -1, 3, -3, 0, -4, 4, -5, 5] }
         }
       },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "k258_unaligned_k_scalar_zp_u8u8",
         "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [128] } },
         "b_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } }
       },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "large_accumulator_dp4a_k1024_u8u8_maxmag",
       "provenance": {
+        "notes": "Confirms the dp4a path and the naive path agree at large int32 accumulation. M=64,K=1024,N=128 -> numel(Y)=8192 and K%4==0 and K>=256, so this selects the dp4a (dot4I8Packed) kernel. Inputs include the max-magnitude bytes (255 for u8) so per-row accumulation reaches ~1024*255*255~=66.6M, well within int32, exercising the bias/zero-point fold (ca,cb,kcc) in quant-dp4a-matmul at large magnitude."
       },
       "inputs": {
         "a": {
         "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [3] } },
         "b_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [7] } }
       },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "dp4a_rank4_scalar_s8s8_2x2x64x128",
     {
       "name": "dp4a_gemv_m1_u8s8_k257_n512_scalar_zp_tail",
       "provenance": {
+        "notes": "An M=1 int8 GEMV with K=257 exercises the tileM=1 DP4A full-word loop and its one-lane scalar tail."
       },
       "inputs": {
         "a": {
           "data": { "kind": "cycle", "values": [1, -2, 2, -1, 3, -3, 0, -4, 4, -5, 5] }
         }
       },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 130], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "dp4a_partial_m63_n130_tile_edge_s8s8_k256",
           "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/dp4a_s8_pattern" } }
         }
       },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "dp4a_tail_k130_rank2_u8u8_scalar_zp_exact",
         "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [121] } },
         "b_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [117] } }
       },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "dp4a_tail_k131_rank2_s8s8_scalar_zp_exact",
         "a_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [5] } },
         "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [-7] } }
       },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "dp4a_tail_k129_batched3_by_rank2_exact",
     },
     {
       "name": "dp4a_tail_k63_rank2_u8s8_scalar_zp_exact",
+      "provenance": { "notes": "K=63 exercises exact packed-tail handling in the 32<=K<128 DP4A range." },
       "inputs": {
         "a": { "dtype": "uint8", "shape": [16, 63], "data": { "kind": "cycle", "values": [0, 255, 121, 7, 200, 31] } },
         "b": { "dtype": "int8", "shape": [63, 512], "data": { "kind": "cycle", "values": [-128, 127, 0, -1, 56, -77] } },
     {
       "name": "dp4a_gemv_no_zero_s8s8_m1_k64_n512",
       "provenance": {
+        "notes": "M=1, N=512, and K=64 select the packed signed GEMV path. Omitting both zero points makes the B-sum correction and its running row sum unnecessary."
       },
       "inputs": {
         "a": {
     {
       "name": "dp4a_gemv_b_per_column_zero_u8u8_m1_k64_n512",
       "provenance": {
+        "notes": "An unsigned B tensor uses per-column zero points that vary on a period coprime with N=512. Each zero point must be widened and indexed by its output column."
       },
       "inputs": {
         "a": {
     {
       "name": "dp4a_tail_no_zero_s8s8_m1_k66_n512",
       "provenance": {
+        "notes": "K=66 leaves a two-lane reduction tail. Signed inputs with both zero points omitted exercise that tail without either running-sum correction."
       },
       "inputs": {
         "a": {
         }
       },
       "outputs": { "y": { "dtype": "int32", "shape": [1, 512], "tolerance": 0 } }
+    },
+    {
+      "name": "exact_integer_partial_k255_uint8_maxmag",
+      "provenance": {
+        "notes": "Maximum adjusted-byte products across exact-f32 partial-sum boundaries. K=33027 also crosses signed-int32 overflow; the integer result must wrap exactly."
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [64, 255], "data": { "kind": "cycle", "values": [255] } },
+        "b": { "dtype": "uint8", "shape": [255, 32], "data": { "kind": "cycle", "values": [255] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 32], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_integer_partial_k256_uint8_maxmag",
+      "provenance": {
+        "notes": "Maximum adjusted-byte products across exact-f32 partial-sum boundaries. K=33027 also crosses signed-int32 overflow; the integer result must wrap exactly."
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [64, 256], "data": { "kind": "cycle", "values": [255] } },
+        "b": { "dtype": "uint8", "shape": [256, 32], "data": { "kind": "cycle", "values": [255] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 32], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_integer_partial_k259_uint8_maxmag",
+      "provenance": {
+        "notes": "Maximum adjusted-byte products across exact-f32 partial-sum boundaries. K=33027 also crosses signed-int32 overflow; the integer result must wrap exactly."
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [64, 259], "data": { "kind": "cycle", "values": [255] } },
+        "b": { "dtype": "uint8", "shape": [259, 32], "data": { "kind": "cycle", "values": [255] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 32], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_integer_partial_k1025_int8_maxmag",
+      "provenance": {
+        "notes": "Maximum adjusted-byte products across exact-f32 partial-sum boundaries. K=33027 also crosses signed-int32 overflow; the integer result must wrap exactly."
+      },
+      "inputs": {
+        "a": { "dtype": "int8", "shape": [64, 1025], "data": { "kind": "cycle", "values": [127] } },
+        "b": { "dtype": "int8", "shape": [1025, 32], "data": { "kind": "cycle", "values": [127] } },
+        "a_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-128] } },
+        "b_zero_point": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-128] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 32], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_integer_partial_k33027_uint8_maxmag",
+      "provenance": {
+        "notes": "Maximum adjusted-byte products across exact-f32 partial-sum boundaries. K=33027 also crosses signed-int32 overflow; the integer result must wrap exactly."
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [64, 33027], "data": { "kind": "cycle", "values": [255] } },
+        "b": { "dtype": "uint8", "shape": [33027, 32], "data": { "kind": "cycle", "values": [255] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 32], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_integer_partial_k513_m65_n65_column_zero_points",
+      "provenance": {
+        "notes": "Mixed signedness, per-column zero points and partial M/N/K tiles must retain exact integer arithmetic."
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [65, 513], "data": { "kind": "cycle", "values": [0, 255, 128, 1, 254] } },
+        "b": { "dtype": "int8", "shape": [513, 65], "data": { "kind": "cycle", "values": [-128, 127, 0, -1, 1] } },
+        "a_zero_point": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [17] } },
+        "b_zero_point": { "dtype": "int8", "shape": [65], "data": { "kind": "cycle", "values": [-128, 127, 3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [65, 65], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_integer_partial_k1041_int8_no_zp_signed_outputs",
+      "provenance": {
+        "notes": "Without zero points, signed-byte bounds permit 1024-term exact f32 partials. The odd final sums +/-16790289 are not representable in f32; converting each partial before the int32 addition preserves both signs exactly."
+      },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 },
+      "inputs": {
+        "a": { "dtype": "int8", "shape": [64, 1041], "data": { "kind": "constant", "value": 127 } },
+        "b": { "dtype": "int8", "shape": [1041, 64], "data": { "kind": "cycle", "values": [127, -127] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 64], "tolerance": 0 } }
+    },
+    {
+      "name": "exact_integer_storage_growth_m32_k128_n65",
+      "provenance": {
+        "notes": "Maximum adjusted-byte products across exact-f32 partial-sum boundaries. K=33027 also crosses signed-int32 overflow; the integer result must wrap exactly."
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [32, 128], "data": { "kind": "cycle", "values": [255] } },
+        "b": { "dtype": "uint8", "shape": [128, 65], "data": { "kind": "cycle", "values": [255] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [32, 65], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_integer_storage_growth_m48_k259_n65",
+      "provenance": {
+        "notes": "Maximum adjusted-byte products across exact-f32 partial-sum boundaries. K=33027 also crosses signed-int32 overflow; the integer result must wrap exactly."
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [48, 259], "data": { "kind": "cycle", "values": [255] } },
+        "b": { "dtype": "uint8", "shape": [259, 65], "data": { "kind": "cycle", "values": [255] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [48, 65], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m65_k33_n33_a0_bnone",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [65, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [65, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m65_k33_n33_a0_bscalar",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [65, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "cycle", "values": [3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [65, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m65_k33_n33_a0_bcolumn",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [65, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [65, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m65_k33_n33_a1_bnone",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [65, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [65, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m65_k33_n33_a1_bscalar",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [65, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "cycle", "values": [3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [65, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m65_k33_n33_a1_bcolumn",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [65, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [65, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m129_k65_n63_a0_bnone",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [129, 65],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [65, 63],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [129, 63], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m129_k65_n63_a0_bscalar",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [129, 65],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [65, 63],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "cycle", "values": [3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [129, 63], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m129_k65_n63_a0_bcolumn",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [129, 65],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [65, 63],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "b_zero_point": { "dtype": "int8", "shape": [63], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [129, 63], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m129_k65_n63_a1_bnone",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [129, 65],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [65, 63],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [129, 63], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m129_k65_n63_a1_bscalar",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [129, 65],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [65, 63],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "cycle", "values": [3] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [129, 63], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_padding_m129_k65_n63_a1_bcolumn",
+      "provenance": {
+        "notes": "Exact operand conversion with simultaneous M/K/N padding; per-column zero points and scales vary independently. The original integer partial and requantization contract is unchanged."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [129, 65],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [65, 63],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [63], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [129, 63], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_row_band_m15_k33_n33",
+      "provenance": {
+        "notes": "Matrix row-band boundaries with all three padded axes and varying column zero points. M15 retains the integer fallback; subsequent bands use one to four effective 32-wide subgroups."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [15, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [15, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_row_band_m16_k33_n33",
+      "provenance": {
+        "notes": "Matrix row-band boundaries with all three padded axes and varying column zero points. M15 retains the integer fallback; subsequent bands use one to four effective 32-wide subgroups."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [16, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [16, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_row_band_m17_k33_n33",
+      "provenance": {
+        "notes": "Matrix row-band boundaries with all three padded axes and varying column zero points. M15 retains the integer fallback; subsequent bands use one to four effective 32-wide subgroups."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [17, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [17, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_row_band_m31_k33_n33",
+      "provenance": {
+        "notes": "Matrix row-band boundaries with all three padded axes and varying column zero points. M15 retains the integer fallback; subsequent bands use one to four effective 32-wide subgroups."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [31, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [31, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_row_band_m32_k33_n33",
+      "provenance": {
+        "notes": "Matrix row-band boundaries with all three padded axes and varying column zero points. M15 retains the integer fallback; subsequent bands use one to four effective 32-wide subgroups."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [32, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [32, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_row_band_m33_k33_n33",
+      "provenance": {
+        "notes": "Matrix row-band boundaries with all three padded axes and varying column zero points. M15 retains the integer fallback; subsequent bands use one to four effective 32-wide subgroups."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [33, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_row_band_m47_k33_n33",
+      "provenance": {
+        "notes": "Matrix row-band boundaries with all three padded axes and varying column zero points. M15 retains the integer fallback; subsequent bands use one to four effective 32-wide subgroups."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [47, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [47, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_row_band_m48_k33_n33",
+      "provenance": {
+        "notes": "Matrix row-band boundaries with all three padded axes and varying column zero points. M15 retains the integer fallback; subsequent bands use one to four effective 32-wide subgroups."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [48, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [48, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_row_band_m49_k33_n33",
+      "provenance": {
+        "notes": "Matrix row-band boundaries with all three padded axes and varying column zero points. M15 retains the integer fallback; subsequent bands use one to four effective 32-wide subgroups."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [49, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [49, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_precast_row_band_m63_k33_n33",
+      "provenance": {
+        "notes": "Matrix row-band boundaries with all three padded axes and varying column zero points. M15 retains the integer fallback; subsequent bands use one to four effective 32-wide subgroups."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [63, 33],
+          "data": { "kind": "cycle", "values": [109, 151, 128, 137, 102, 165, 124] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [33, 33],
+          "data": { "kind": "cycle", "values": [-9, 7, -3, 12, 5, -11, 2, 1, -6, 4, 8] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": { "dtype": "int8", "shape": [33], "data": { "kind": "cycle", "values": [-3, 2, 7, -5, 1] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [63, 33], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_portable_m63_k255_n35",
+      "provenance": {
+        "notes": "Exact partial-sum and row-reuse boundaries, signed and unsigned extremes, column zero points, and simultaneous padded axes."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [63, 255],
+          "data": { "kind": "cycle", "values": [0, 255, 17, 239, 128, 1, 254] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [255, 35],
+          "data": { "kind": "cycle", "values": [-128, 127, -127, 63, -64, 0, 126] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [255] } },
+        "b_zero_point": {
+          "dtype": "int8",
+          "shape": [35],
+          "data": { "kind": "cycle", "values": [-128, 127, 0, -111, 111] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [63, 35], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_portable_m64_k256_n35",
+      "provenance": {
+        "notes": "Exact partial-sum and row-reuse boundaries, signed and unsigned extremes, column zero points, and simultaneous padded axes."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [64, 256],
+          "data": { "kind": "cycle", "values": [0, 255, 17, 239, 128, 1, 254] }
+        },
+        "b": {
+          "dtype": "uint8",
+          "shape": [256, 35],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 191, 64, 128, 254] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "cycle", "values": [255] } },
+        "b_zero_point": {
+          "dtype": "uint8",
+          "shape": [35],
+          "data": { "kind": "cycle", "values": [0, 255, 128, 17, 239] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 35], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_portable_m65_k257_n67",
+      "provenance": {
+        "notes": "Exact partial-sum and row-reuse boundaries, signed and unsigned extremes, column zero points, and simultaneous padded axes."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "int8",
+          "shape": [65, 257],
+          "data": { "kind": "cycle", "values": [-128, 127, -111, 111, 0, -127, 126] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [257, 67],
+          "data": { "kind": "cycle", "values": [-128, 127, -127, 63, -64, 0, 126] }
+        },
+        "a_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": {
+          "dtype": "int8",
+          "shape": [67],
+          "data": { "kind": "cycle", "values": [-128, 127, 0, -111, 111] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [65, 67], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "exact_portable_m65_k1025_n35",
+      "provenance": {
+        "notes": "Exact partial-sum and row-reuse boundaries, signed and unsigned extremes, column zero points, and simultaneous padded axes."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "int8",
+          "shape": [65, 1025],
+          "data": { "kind": "cycle", "values": [-128, 127, -111, 111, 0, -127, 126] }
+        },
+        "b": {
+          "dtype": "uint8",
+          "shape": [1025, 35],
+          "data": { "kind": "cycle", "values": [0, 255, 1, 191, 64, 128, 254] }
+        },
+        "a_zero_point": { "dtype": "int8", "shape": [1], "data": { "kind": "cycle", "values": [127] } },
+        "b_zero_point": {
+          "dtype": "uint8",
+          "shape": [35],
+          "data": { "kind": "cycle", "values": [0, 255, 128, 17, 239] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [65, 35], "tolerance": 0 } },
+      "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     }
   ]
 }