sync 6fdf6301e2bb

Browse files

Files changed (12) hide show

README.md +4 -2
build/webgpu/bench.json +28 -2
build/webgpu/manifest.json +225 -355
build/webgpu/matmul-integer-batched.wgsl.jinja +3 -6
build/webgpu/metadata.json +13 -13
build/webgpu/quant-dp4a-matmul.wgsl.jinja +18 -24
build/webgpu/quant-exact-matrix.wgsl.jinja +3 -4
build/webgpu/quant-exact-portable.wgsl.jinja +4 -7
build/webgpu/quant-exact-prepare.wgsl.jinja +5 -4
build/webgpu/quant-matmul-accumulate-rank2.wgsl.jinja +7 -8
build/webgpu/quant-matmul-accumulate-rank4.wgsl.jinja +1 -3
build/webgpu/test.json +96 -21

README.md CHANGED Viewed

@@ -73,6 +73,8 @@ One implementation is selected per call from the device capabilities, the reques
 - `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
@@ -83,7 +85,7 @@ Some implementation variants require `subgroup-matrix`, `shader-f16`, and `subgr
 - [`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)
@@ -95,7 +97,7 @@ Some implementation variants require `subgroup-matrix`, `shader-f16`, and `subgr
 ## 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.

 - `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.
+- `dp4a_rank2_b_zero_point_per_column` — Packed-dot rank-2 matmul with per-column B zero points. Large prefill tiles use twice as many output-column lanes on variable 16–32-lane devices with sufficient workgroup capacity.
+- `dp4a_rank2_a_zero_point_b_per_column` — Packed-dot rank-2 matmul with per-column B zero points. Large prefill tiles use twice as many output-column lanes on variable 16–32-lane devices with sufficient workgroup capacity.
 ## Device requirements
 - [`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 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)
 ## Use with `@huggingface/kernels`
 ```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.

build/webgpu/bench.json CHANGED Viewed

@@ -136,6 +136,32 @@
       "tunableSpace": { "TILE_M": [8, 16, 32], "TILE_N": [32, 64, 128], "KT_WORDS": [32, 64] },
       "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
     },
     {
       "name": "kfloor_k128_2048x128x2048_u8s8_naive",
       "preset": "smoke",
@@ -298,11 +324,11 @@
       "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.batch * args.M * args.K * args.N" }] }
     },
     {
-      "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 },
       "inputs": {

       "tunableSpace": { "TILE_M": [8, 16, 32], "TILE_N": [32, 64, 128], "KT_WORDS": [32, 64] },
       "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
     },
+    {
+      "name": "percol_b_zp_no_azp_1024x4096x4096_u8s8_dp4a",
+      "preset": "model",
+      "vars": { "M": 1024, "K": 4096, "N": 4096 },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [1024, 4096], "dist": "randint", "seed": 2339, "min": 0, "max": 255 },
+        "b": { "dtype": "int8", "shape": [4096, 4096], "dist": "randint", "seed": 1316, "min": -127, "max": 127 },
+        "b_zero_point": { "dtype": "int8", "shape": [4096], "dist": "randint", "seed": 1317, "min": -5, "max": 5 }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1024, 4096] } },
+      "tunableSpace": { "TILE_M": [8, 16, 32], "TILE_N": [32, 64, 128], "KT_WORDS": [32, 64] },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
+    },
+    {
+      "name": "percol_b_zp_no_azp_256x4096x4096_u8s8_dp4a",
+      "preset": "model",
+      "vars": { "M": 256, "K": 4096, "N": 4096 },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [256, 4096], "dist": "randint", "seed": 1571, "min": 0, "max": 255 },
+        "b": { "dtype": "int8", "shape": [4096, 4096], "dist": "randint", "seed": 1316, "min": -127, "max": 127 },
+        "b_zero_point": { "dtype": "int8", "shape": [4096], "dist": "randint", "seed": 1317, "min": -5, "max": 5 }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [256, 4096] } },
+      "tunableSpace": { "TILE_M": [8, 16, 32], "TILE_N": [32, 64, 128], "KT_WORDS": [32, 64] },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.M * args.K * args.N" }] }
+    },
     {
       "name": "kfloor_k128_2048x128x2048_u8s8_naive",
       "preset": "smoke",
       "bench": { "metrics": [{ "type": "gflops", "value": "2 * args.batch * args.M * args.K * args.N" }] }
     },
     {
+      "name": "rank2-u8s8-k63-dp4a-floor-m512n4096",
       "preset": "stress",
       "provenance": {
         "source": "synthetic benchmark",
+        "notes": "Measures MatMulInteger over a 512x4096 uint8 by int8 output with an unaligned K=63 reduction length."
       },
       "vars": { "M": 512, "K": 63, "N": 4096 },
       "inputs": {

build/webgpu/manifest.json CHANGED Viewed

@@ -44,16 +44,25 @@
     "EXACT_MATRIX_MIN_INVOCATIONS": { "default": 16384 }
   },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "wave32Adapter": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 32 and device.adapterInfo.subgroupMaxSize == 32",
     "canPinSubgroupSize32": "device.features.has(\"subgroups\") and device.features.has(\"subgroup-size-control\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize <= 32 and device.adapterInfo.subgroupMaxSize >= 32",
     "pinSubgroupSize32": "canPinSubgroupSize32 and not wave32Adapter",
     "wave32Effective": "wave32Adapter or pinSubgroupSize32",
     "sgmatTileN": "64",
     "sgmatTileK": "32",
     "sgmatColumnGroups": "2",
     "sgmatStorageRowGroups": "floor((device.limits.maxComputeWorkgroupStorageSize - sgmatTileN * sgmatTileK * 4) / (16 * sgmatTileK * 4))",
-    "sgmatRowGroups": "max(1, min(4, floor(deviceWorkgroupCap / (32 * sgmatColumnGroups)), sgmatStorageRowGroups, max(floor(sgmatStorageRowGroups / 2), floor(dim(shapes.a, 0) / 16))))",
     "sgmatTileM": "16 * sgmatRowGroups",
     "sgmatSubgroups": "sgmatRowGroups * sgmatColumnGroups",
     "sgmatWorkgroup": "32 * sgmatSubgroups",
@@ -62,31 +71,31 @@
     "sgmatReuseStaging": "sgmatStagingBytes + sgmatScratchBytes > device.limits.maxComputeWorkgroupStorageSize",
     "sgmatSharedBytes": "sgmatStagingBytes + (0 if sgmatReuseStaging else sgmatScratchBytes)",
     "sgmatDeviceOk": "sgmatWorkgroup <= deviceWorkgroupCap and sgmatSharedBytes <= device.limits.maxComputeWorkgroupStorageSize",
-    "precastSubgroups": "max(1, min(4, ceilDiv(dim(shapes.a, 0), 16), floor(deviceWorkgroupCap / 32), floor(device.limits.maxComputeWorkgroupStorageSize / (8 * 64 * 4))))",
     "precastTileM": "16 * precastSubgroups",
     "precastTileN": "4 * 8",
     "precastTileK": "4 * 8",
     "precastWorkgroup": "32 * precastSubgroups",
-    "preparedK": "ceilDiv(dim(shapes.a, 1), precastTileK) * precastTileK",
-    "preparedN": "ceilDiv(dim(shapes.b, 1), precastTileN) * precastTileN",
-    "preparedElements": "ceilDiv(dim(shapes.a, 0), precastTileM) * precastTileM * preparedK",
-    "preparedBElements": "preparedK * preparedN",
     "prepareWorkgroup": "min(256, deviceWorkgroupCap)",
     "prepareTotalElements": "preparedElements + preparedBElements",
     "prepareGroups": "ceilDiv(prepareTotalElements, 4 * prepareWorkgroup)",
     "prepareDispatchElements": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535)) * min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535)) * prepareWorkgroup * 4",
     "precastResourcesOk": "deviceWorkgroupCap >= precastWorkgroup and precastSubgroups * 8 * 64 * 4 <= device.limits.maxComputeWorkgroupStorageSize",
-    "precastDispatchOk": "prepareTotalElements <= 4294967295 and prepareDispatchElements <= 4294967296 and ceilDiv(dim(shapes.b, 1), precastTileN) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and ceilDiv(dim(shapes.a, 0), precastTileM) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and prepareGroups <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "precastF32BuffersOk": "preparedElements * 4 <= device.limits.maxStorageBufferBindingSize and preparedElements * 4 <= device.limits.maxBufferSize and preparedBElements * 4 <= device.limits.maxStorageBufferBindingSize and preparedBElements * 4 <= device.limits.maxBufferSize",
     "precastF16BuffersOk": "preparedElements * 2 <= device.limits.maxStorageBufferBindingSize and preparedElements * 2 <= device.limits.maxBufferSize and preparedBElements * 2 <= device.limits.maxStorageBufferBindingSize and preparedBElements * 2 <= device.limits.maxBufferSize",
     "precastWorthwhile": "dim(shapes.a, 0) >= 16 and dim(shapes.b, 1) >= 32 and dim(shapes.a, 1) >= 32 and floor(dim(shapes.a, 0) / 16) * floor(dim(shapes.b, 1) / 32) * 32 >= tunables.EXACT_MATRIX_MIN_INVOCATIONS",
     "portableWorkgroupCols": 16,
-    "portableWorkgroupRows": "pow2ceil(min(16, max(4, pow2ceil(ceilDiv(dim(shapes.a, 0), 4))), floor(deviceWorkgroupCap / portableWorkgroupCols), device.limits.maxComputeWorkgroupSizeY) + 1) / 2",
     "portableTileM": "4 * portableWorkgroupRows",
     "portableTileN": "4 * portableWorkgroupCols",
     "portableTileK": 32,
     "portablePartialK": "floor(16777216 / (255 * 255) / portableTileK) * portableTileK",
-    "portableGeometryOk": "portableWorkgroupRows >= 4 and portableWorkgroupCols <= device.limits.maxComputeWorkgroupSizeX and portablePartialK >= portableTileK and ceilDiv(dim(shapes.a, 0), portableTileM) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "rank2MatMulShapeOk": "ranks.a == 2 and ranks.b == 2 and ranks.y == 2 and dim(shapes.a, 1) == dim(shapes.b, 0) and dim(shapes.y, 0) == dim(shapes.a, 0) and dim(shapes.y, 1) == dim(shapes.b, 1)",
     "rank3ByRank3ShapeOk": "ranks.a == 3 and ranks.b == 3 and ranks.y == 3 and dim(shapes.a, 2) == dim(shapes.b, 1) and dim(shapes.y, 0) == max(dim(shapes.a, 0), dim(shapes.b, 0)) and (dim(shapes.a, 0) == dim(shapes.y, 0) or dim(shapes.a, 0) == 1) and (dim(shapes.b, 0) == dim(shapes.y, 0) or dim(shapes.b, 0) == 1) and dim(shapes.y, 1) == dim(shapes.a, 1) and dim(shapes.y, 2) == dim(shapes.b, 2)",
     "rank3ByRank2ShapeOk": "ranks.a == 3 and ranks.b == 2 and ranks.y == 3 and dim(shapes.a, 2) == dim(shapes.b, 0) and dim(shapes.y, 0) == dim(shapes.a, 0) and dim(shapes.y, 1) == dim(shapes.a, 1) and dim(shapes.y, 2) == dim(shapes.b, 1)",
@@ -99,6 +108,7 @@
     "requiredScalarBZeroOk": "present.b_zero_point and scalarBZeroOk",
     "rank2PerColumnBZeroOk": "ranks.b_zero_point == 1 and dim(shapes.b_zero_point, 0) == dim(shapes.y, 1)",
     "rank2Dp4aLargePrefill": "ranks.y == 2 and ranks.a == 2 and dim(shapes.y, 0) >= tunables.DP4A_LARGE_MIN_M and dim(shapes.y, 1) >= tunables.DP4A_LARGE_MIN_N and dim(shapes.a, 1) >= tunables.DP4A_LARGE_MIN_K",
     "rank2Dp4aWideShortK": "ranks.y == 2 and ranks.a == 2 and dim(shapes.y, 0) >= tunables.DP4A_SHORT_K_MIN_MN and dim(shapes.y, 1) >= tunables.DP4A_SHORT_K_MIN_MN and dim(shapes.a, 1) <= tunables.DP4A_SHORT_K_MAX",
     "rank2Dp4aMediumK": "ranks.y == 2 and ranks.a == 2 and dim(shapes.y, 0) >= tunables.DP4A_MEDIUM_K_MIN_MN and dim(shapes.y, 1) >= tunables.DP4A_MEDIUM_K_MIN_MN and dim(shapes.a, 1) > tunables.DP4A_SHORT_K_MAX and dim(shapes.a, 1) <= tunables.DP4A_MEDIUM_K_MAX",
     "rank2Dp4aTileM": "1 if ranks.y < 2 or dim(shapes.y, 0) == 1 else (tunables.DP4A_SHORT_K_TILE_M if rank2Dp4aWideShortK else (tunables.DP4A_MEDIUM_K_TILE_M if rank2Dp4aMediumK else (tunables.DP4A_LARGE_TILE_M if rank2Dp4aLargePrefill else tunables.DP4A_TILE_M)))",
@@ -127,25 +137,24 @@
     "rank4YBatchStride": "dim(shapes.y, 2) * dim(shapes.y, 3) if ranks.y >= 4 else 1",
     "aScalar": "\"u32\" if tensorDtypes.a == \"uint8\" else \"i32\"",
     "bScalar": "\"u32\" if tensorDtypes.b == \"uint8\" else \"i32\"",
-    "aUnsigned": "tensorDtypes.a == \"uint8\"",
-    "bUnsigned": "tensorDtypes.b == \"uint8\"",
-    "aZeroPerRow": false,
-    "hasBias": false,
     "sgmatM": "matrixM",
     "sgmatN": "matrixN",
     "sgmatK": "matrixK",
     "sgmatDispatchOk": "ceilDiv(sgmatN, sgmatTileN) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and ceilDiv(sgmatM, sgmatTileM) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "sgmatParallelInvocations": "floor(sgmatM / sgmatTileM) * floor(sgmatN / sgmatTileN) * sgmatWorkgroup",
-    "sgmatWorthwhile": "sgmatM >= sgmatTileM and sgmatN >= sgmatTileN / 2 and sgmatK >= sgmatTileK and sgmatParallelInvocations >= tunables.EXACT_MATRIX_MIN_INVOCATIONS"
   },
   "bindings": {
-    "a": { "buffer": "read-only-storage", "elementType": "$aStorage" },
-    "b": { "buffer": "read-only-storage", "elementType": "$bScalar" },
-    "a_zero_point": { "buffer": "read-only-storage", "elementType": "$aScalar", "length": 1 },
-    "b_zero_point": { "buffer": "read-only-storage", "elementType": "$bScalar" },
-    "y": { "buffer": "storage", "elementType": "i32" },
     "params": {
-      "buffer": "uniform",
       "struct": [
         { "name": "M", "type": "u32", "value": "matrixM" },
         { "name": "N", "type": "u32", "value": "matrixN" },
@@ -155,29 +164,26 @@
         { "name": "yBatchStride", "type": "u32", "value": 0 }
       ]
     },
-    "a_2": { "name": "a", "buffer": "read-only-storage", "elementType": "$aScalar" },
-    "a_3": { "name": "a", "buffer": "read-only-storage", "elementType": "$aVec4" },
-    "params_3": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "M", "type": "u32", "value": "matrixM" },
         { "name": "N", "type": "u32", "value": "matrixN" },
         { "name": "K", "type": "u32", "value": "matrixK" }
       ]
     },
-    "params_4": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "M", "type": "u32", "value": 1 },
         { "name": "N", "type": "u32", "value": 1 },
         { "name": "K", "type": "u32", "value": "dim(shapes.a, 0)" }
       ]
     },
-    "params_5": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "M", "type": "u32", "value": "batchM" },
         { "name": "N", "type": "u32", "value": "batchN" },
@@ -188,9 +194,8 @@
         { "name": "batchCount", "type": "u32", "value": "batchCount" }
       ]
     },
-    "params_6": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "M", "type": "u32", "value": "rank4M" },
         { "name": "N", "type": "u32", "value": "rank4N" },
@@ -219,9 +224,8 @@
         }
       ]
     },
-    "params_7": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "M", "type": "u32", "value": "batchM" },
         { "name": "N", "type": "u32", "value": "batchN" },
@@ -231,9 +235,8 @@
         { "name": "yBatchStride", "type": "u32", "value": "batchYStride" }
       ]
     },
-    "params_8": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "M", "type": "u32", "value": "rank4M" },
         { "name": "N", "type": "u32", "value": "rank4N" },
@@ -249,7 +252,6 @@
     "precast_f32_bFloat": { "name": "bFloat", "buffer": "storage", "elementType": "f32" },
     "precast_prepare_params": {
       "name": "params",
-      "buffer": "uniform",
       "struct": [
         { "name": "M", "type": "u32", "value": "sgmatM" },
         { "name": "K", "type": "u32", "value": "sgmatK" },
@@ -291,20 +293,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f16\"",
         "hasAZero": false,
         "hasBZero": false,
-        "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f16_aFloat", "b", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -312,9 +311,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f16_a", "precast_f16_b", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -327,6 +326,7 @@
       "id": "portable_precast_a0_bnone_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "not present.b_zero_point", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
@@ -334,8 +334,6 @@
         "hasAZero": false,
         "hasBZero": false,
         "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -344,17 +342,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f16_aFloat", "b", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -362,7 +357,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -389,20 +384,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f32\"",
         "hasAZero": false,
         "hasBZero": false,
-        "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f32_aFloat", "b", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -410,9 +402,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f32_a", "precast_f32_b", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -425,6 +417,7 @@
       "id": "portable_precast_a0_bnone",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "not present.b_zero_point", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
@@ -432,8 +425,6 @@
         "hasAZero": false,
         "hasBZero": false,
         "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -442,17 +433,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f32_aFloat", "b", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -460,7 +448,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -487,20 +475,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f16\"",
         "hasAZero": false,
         "hasBZero": true,
-        "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -508,9 +493,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f16_a", "precast_f16_b", "b_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -523,6 +508,7 @@
       "id": "portable_precast_a0_bscalar_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "requiredScalarBZeroOk", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
@@ -530,8 +516,6 @@
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -540,17 +524,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -558,7 +539,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -585,20 +566,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f32\"",
         "hasAZero": false,
         "hasBZero": true,
-        "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -606,9 +584,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f32_a", "precast_f32_b", "b_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -621,6 +599,7 @@
       "id": "portable_precast_a0_bscalar",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "requiredScalarBZeroOk", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
@@ -628,8 +607,6 @@
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -638,17 +615,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -656,7 +630,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -683,20 +657,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f16\"",
         "hasAZero": false,
         "hasBZero": true,
-        "bZeroPerColumn": true,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -704,9 +675,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f16_a", "precast_f16_b", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -719,6 +690,7 @@
       "id": "portable_precast_a0_bcolumn_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "present.b_zero_point and rank2PerColumnBZeroOk", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
@@ -726,8 +698,6 @@
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": true,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -736,17 +706,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -754,7 +721,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -781,20 +748,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f32\"",
         "hasAZero": false,
         "hasBZero": true,
-        "bZeroPerColumn": true,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -802,9 +766,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f32_a", "precast_f32_b", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -817,6 +781,7 @@
       "id": "portable_precast_a0_bcolumn",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "present.b_zero_point and rank2PerColumnBZeroOk", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
@@ -824,8 +789,6 @@
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": true,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -834,17 +797,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -852,7 +812,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -879,20 +839,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f16\"",
         "hasAZero": true,
         "hasBZero": false,
-        "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f16_aFloat", "b", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -900,9 +857,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f16_a", "precast_f16_b", "a_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -915,6 +872,7 @@
       "id": "portable_precast_a1_bnone_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "not present.b_zero_point", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
@@ -922,8 +880,6 @@
         "hasAZero": true,
         "hasBZero": false,
         "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -932,17 +888,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f16_aFloat", "b", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -950,7 +903,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -977,20 +930,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f32\"",
         "hasAZero": true,
         "hasBZero": false,
-        "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f32_aFloat", "b", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -998,9 +948,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f32_a", "precast_f32_b", "a_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -1013,6 +963,7 @@
       "id": "portable_precast_a1_bnone",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "not present.b_zero_point", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
@@ -1020,8 +971,6 @@
         "hasAZero": true,
         "hasBZero": false,
         "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -1030,17 +979,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f32_aFloat", "b", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -1048,7 +994,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -1075,20 +1021,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f16\"",
         "hasAZero": true,
         "hasBZero": true,
-        "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -1096,9 +1039,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f16_a", "precast_f16_b", "a_zero_point", "b_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -1111,6 +1054,7 @@
       "id": "portable_precast_a1_bscalar_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "requiredScalarBZeroOk", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
@@ -1118,8 +1062,6 @@
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -1128,17 +1070,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -1146,7 +1085,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -1173,20 +1112,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f32\"",
         "hasAZero": true,
         "hasBZero": true,
-        "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -1194,9 +1130,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f32_a", "precast_f32_b", "a_zero_point", "b_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -1209,6 +1145,7 @@
       "id": "portable_precast_a1_bscalar",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "requiredScalarBZeroOk", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
@@ -1216,8 +1153,6 @@
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -1226,17 +1161,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -1244,7 +1176,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -1271,20 +1203,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f16\"",
         "hasAZero": true,
         "hasBZero": true,
-        "bZeroPerColumn": true,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -1292,9 +1221,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f16_a", "precast_f16_b", "a_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -1307,6 +1236,7 @@
       "id": "portable_precast_a1_bcolumn_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "present.b_zero_point and rank2PerColumnBZeroOk", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
@@ -1314,8 +1244,6 @@
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": true,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -1324,17 +1252,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -1342,7 +1267,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -1369,20 +1294,17 @@
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
-        "sgmatSharedBytes": "sgmatSubgroups * 8 * 64 * 4",
         "precastScalar": "\"f32\"",
         "hasAZero": true,
         "hasBZero": true,
-        "bZeroPerColumn": true,
-        "bScalePerColumn": false,
-        "requantize": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -1390,9 +1312,9 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.Multiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["precast_f32_a", "precast_f32_b", "a_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
@@ -1405,6 +1327,7 @@
       "id": "portable_precast_a1_bcolumn",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "present.b_zero_point and rank2PerColumnBZeroOk", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
@@ -1412,8 +1335,6 @@
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": true,
-        "bScalePerColumn": false,
-        "requantize": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
@@ -1422,17 +1343,14 @@
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
-        "integerPartialK": "portablePartialK",
-        "packedATile": true,
-        "vec4BTile": true,
-        "hasBias": false
       },
       "passes": [
         {
           "id": "prepare",
-          "name": "ExactMatrix.PrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
-          "bindings": ["a_2", "a_zero_point", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
@@ -1440,7 +1358,7 @@
         },
         {
           "id": "main",
-          "name": "ExactMatrix.PortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
@@ -1459,13 +1377,13 @@
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
-      "derive": { "hasAZero": false, "hasBZero": false, "bZeroPerColumn": false, "requantize": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["a_2", "b", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
@@ -1478,13 +1396,13 @@
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
-      "derive": { "hasAZero": false, "hasBZero": true, "bZeroPerColumn": false, "requantize": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["a_2", "b", "b_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
@@ -1497,13 +1415,13 @@
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
-      "derive": { "hasAZero": false, "hasBZero": true, "bZeroPerColumn": true, "requantize": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["a_2", "b", "b_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
@@ -1516,13 +1434,13 @@
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
-      "derive": { "hasAZero": true, "hasBZero": false, "bZeroPerColumn": false, "requantize": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
@@ -1535,13 +1453,13 @@
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
-      "derive": { "hasAZero": true, "hasBZero": true, "bZeroPerColumn": false, "requantize": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "b_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
@@ -1554,13 +1472,13 @@
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
-      "derive": { "hasAZero": true, "hasBZero": true, "bZeroPerColumn": true, "requantize": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "b_zero_point", "y", "params_3"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
@@ -1575,12 +1493,7 @@
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
-        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL",
-        "tileM": "rank2Dp4aTileM",
-        "tileN": "rank2Dp4aTileN",
-        "ktWords": "rank2Dp4aKtWords"
       },
       "passes": [
         {
@@ -1605,12 +1518,7 @@
         "hasAZero": false,
         "hasBZero": false,
         "bZeroPerColumn": false,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
-        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL",
-        "tileM": "rank2Dp4aTileM",
-        "tileN": "rank2Dp4aTileN",
-        "ktWords": "rank2Dp4aKtWords"
       },
       "passes": [
         {
@@ -1634,8 +1542,6 @@
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
         "tileM": 8,
         "tileN": 64,
         "ktWords": 32
@@ -1646,7 +1552,7 @@
           "name": "MatMulInteger.Dp4aBatched3x2Tail",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
           "bindings": [
-            "a_2",
             "b",
             "a_zero_point",
             "b_zero_point",
@@ -1676,12 +1582,7 @@
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
-        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL",
-        "tileM": "rank2Dp4aTileM",
-        "tileN": "rank2Dp4aTileN",
-        "ktWords": "rank2Dp4aKtWords"
       },
       "passes": [
         {
@@ -1705,19 +1606,14 @@
         "hasAZero": true,
         "hasBZero": false,
         "bZeroPerColumn": false,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
-        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL",
-        "tileM": "rank2Dp4aTileM",
-        "tileN": "rank2Dp4aTileN",
-        "ktWords": "rank2Dp4aKtWords"
       },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Dp4a",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
-          "bindings": ["a_3", "b", "a_zero_point", "y", "params"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.y, 1), rank2Dp4aTileN)",
             "y": "ceilDiv(dim(shapes.y, 0), rank2Dp4aTileM)"
@@ -1734,12 +1630,7 @@
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": false,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
-        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL",
-        "tileM": "rank2Dp4aTileM",
-        "tileN": "rank2Dp4aTileN",
-        "ktWords": "rank2Dp4aKtWords"
       },
       "passes": [
         {
@@ -1763,19 +1654,14 @@
         "hasAZero": false,
         "hasBZero": false,
         "bZeroPerColumn": false,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
-        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL",
-        "tileM": "rank2Dp4aTileM",
-        "tileN": "rank2Dp4aTileN",
-        "ktWords": "rank2Dp4aKtWords"
       },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Dp4a",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
-          "bindings": ["a_3", "b", "y", "params"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.y, 1), rank2Dp4aTileN)",
             "y": "ceilDiv(dim(shapes.y, 0), rank2Dp4aTileM)"
@@ -1800,11 +1686,9 @@
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": true,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
         "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL",
-        "tileM": "rank2Dp4aTileM if dim(shapes.y, 0) == 1 else tunables.TILE_M",
-        "tileN": "rank2Dp4aTileN if dim(shapes.y, 0) == 1 else tunables.TILE_N",
         "ktWords": "rank2Dp4aKtWords if dim(shapes.y, 0) == 1 else tunables.KT_WORDS"
       },
       "passes": [
@@ -1813,10 +1697,7 @@
           "name": "MatMulInteger.Dp4aBPerColumn",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
           "bindings": ["a", "b", "b_zero_point", "y", "params"],
-          "dispatch": {
-            "x": "ceilDiv(dim(shapes.y, 1), rank2Dp4aTileN if dim(shapes.y, 0) == 1 else tunables.TILE_N)",
-            "y": "ceilDiv(dim(shapes.y, 0), rank2Dp4aTileM if dim(shapes.y, 0) == 1 else tunables.TILE_M)"
-          }
         }
       ]
     },
@@ -1837,11 +1718,9 @@
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": true,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
         "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL",
-        "tileM": "rank2Dp4aTileM if dim(shapes.y, 0) == 1 else tunables.TILE_M",
-        "tileN": "rank2Dp4aTileN if dim(shapes.y, 0) == 1 else tunables.TILE_N",
         "ktWords": "rank2Dp4aKtWords if dim(shapes.y, 0) == 1 else tunables.KT_WORDS"
       },
       "passes": [
@@ -1850,10 +1729,7 @@
           "name": "MatMulInteger.Dp4aAScalarBPerColumn",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
           "bindings": ["a", "b", "a_zero_point", "b_zero_point", "y", "params"],
-          "dispatch": {
-            "x": "ceilDiv(dim(shapes.y, 1), rank2Dp4aTileN if dim(shapes.y, 0) == 1 else tunables.TILE_N)",
-            "y": "ceilDiv(dim(shapes.y, 0), rank2Dp4aTileM if dim(shapes.y, 0) == 1 else tunables.TILE_M)"
-          }
         }
       ]
     },
@@ -1866,11 +1742,11 @@
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "b_zero_point", "y", "params_3"],
           "dispatch": {
             "x": "ceil(matrixN / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((matrixM), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
-            "z": "ceilDiv(ceilDiv((matrixM), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)"
           }
         }
       ]
@@ -1884,11 +1760,11 @@
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "y", "params_3"],
           "dispatch": {
             "x": "ceil(matrixN / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((matrixM), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
-            "z": "ceilDiv(ceilDiv((matrixM), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)"
           }
         }
       ]
@@ -1902,11 +1778,11 @@
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
-          "bindings": ["a_2", "b", "b_zero_point", "y", "params_3"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
-            "z": "ceilDiv(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)"
           }
         }
       ]
@@ -1920,7 +1796,7 @@
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
-          "bindings": ["a_2", "b", "b_zero_point", "y", "params_4"],
           "dispatch": { "x": 1 }
         }
       ]
@@ -1934,7 +1810,7 @@
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
-          "bindings": ["a_2", "b", "y", "params_4"],
           "dispatch": { "x": 1 }
         }
       ]
@@ -1948,11 +1824,11 @@
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
-          "bindings": ["a_2", "b", "y", "params_3"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
-            "z": "ceilDiv(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)"
           }
         }
       ]
@@ -1966,11 +1842,11 @@
           "id": "main",
           "name": "MatMulInteger.PerColumn",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "b_zero_point", "y", "params_3"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
-            "z": "ceilDiv(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)"
           }
         }
       ]
@@ -1984,11 +1860,11 @@
           "id": "main",
           "name": "MatMulInteger.PerColumnNoAZero",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
-          "bindings": ["a_2", "b", "b_zero_point", "y", "params_3"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
-            "z": "ceilDiv(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)"
           }
         }
       ]
@@ -1996,13 +1872,13 @@
     {
       "id": "rank3_by_rank3_scalar_zero_points",
       "when": ["present.a_zero_point", "present.b_zero_point", "(rank3ByRank3ShapeOk or rank3ByRank2ShapeOk)", "scalarAZeroOk and scalarBZeroOk"],
-      "derive": { "hasAZero": true, "hasBZero": true, "bZeroPerColumn": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "b_zero_point", "y", "params_5"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2014,13 +1890,13 @@
     {
       "id": "rank3_by_rank3_no_zero_points",
       "when": ["not present.a_zero_point", "not present.b_zero_point", "(rank3ByRank3ShapeOk or rank3ByRank2ShapeOk)", "true"],
-      "derive": { "hasAZero": false, "hasBZero": false, "bZeroPerColumn": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
-          "bindings": ["a_2", "b", "y", "params_5"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2032,13 +1908,13 @@
     {
       "id": "rank3_by_rank3_a_zero_point_only",
       "when": ["present.a_zero_point", "not present.b_zero_point", "(rank3ByRank3ShapeOk or rank3ByRank2ShapeOk)", "scalarAZeroOk"],
-      "derive": { "hasAZero": true, "hasBZero": false, "bZeroPerColumn": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "y", "params_5"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2050,13 +1926,13 @@
     {
       "id": "rank3_by_rank3_b_zero_point_only",
       "when": ["not present.a_zero_point", "present.b_zero_point", "(rank3ByRank3ShapeOk or rank3ByRank2ShapeOk)", "scalarBZeroOk"],
-      "derive": { "hasAZero": false, "hasBZero": true, "bZeroPerColumn": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
-          "bindings": ["a_2", "b", "b_zero_point", "y", "params_5"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2068,13 +1944,13 @@
     {
       "id": "rank3_by_rank3_b_zero_point_per_column_batched",
       "when": ["present.a_zero_point", "present.b_zero_point", "rank3ByRank3ShapeOk", "scalarAZeroOk and ranks.b_zero_point == 3 and dim(shapes.b_zero_point, 0) == dim(shapes.y, 0) and dim(shapes.b_zero_point, 1) == 1 and dim(shapes.b_zero_point, 2) == dim(shapes.y, 2)"],
-      "derive": { "hasAZero": true, "hasBZero": true, "bZeroPerColumn": false, "bZeroBatchedPerColumn": true },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3PerColumn",
           "shader": "matmul-integer-batched.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "b_zero_point", "y", "params_5"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2086,13 +1962,13 @@
     {
       "id": "rank3_by_rank3_b_zero_point_per_column_batched_no_azp",
       "when": ["not present.a_zero_point", "present.b_zero_point", "rank3ByRank3ShapeOk", "ranks.b_zero_point == 3 and dim(shapes.b_zero_point, 0) == dim(shapes.y, 0) and dim(shapes.b_zero_point, 1) == 1 and dim(shapes.b_zero_point, 2) == dim(shapes.y, 2)"],
-      "derive": { "hasAZero": false, "hasBZero": true, "bZeroPerColumn": false, "bZeroBatchedPerColumn": true },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3PerColumn",
           "shader": "matmul-integer-batched.wgsl.jinja",
-          "bindings": ["a_2", "b", "b_zero_point", "y", "params_5"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2104,13 +1980,13 @@
     {
       "id": "rank2_by_rank3_scalar_zero_points",
       "when": ["present.a_zero_point", "present.b_zero_point", "ranks.a == 2 and ranks.b == 3 and ranks.y == 3 and dim(shapes.a, 1) == dim(shapes.b, 1) and dim(shapes.y, 0) == dim(shapes.b, 0) and dim(shapes.y, 1) == dim(shapes.a, 0) and dim(shapes.y, 2) == dim(shapes.b, 2)", "scalarAZeroOk and scalarBZeroOk"],
-      "derive": { "hasAZero": true, "hasBZero": true, "bZeroPerColumn": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched2x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "b_zero_point", "y", "params_5"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2122,13 +1998,13 @@
     {
       "id": "rank2_by_rank3_a_zero_point_only",
       "when": ["present.a_zero_point", "not present.b_zero_point", "ranks.a == 2 and ranks.b == 3 and ranks.y == 3 and dim(shapes.a, 1) == dim(shapes.b, 1) and dim(shapes.y, 0) == dim(shapes.b, 0) and dim(shapes.y, 1) == dim(shapes.a, 0) and dim(shapes.y, 2) == dim(shapes.b, 2)", "scalarAZeroOk"],
-      "derive": { "hasAZero": true, "hasBZero": false, "bZeroPerColumn": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched2x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
-          "bindings": ["a_2", "b", "a_zero_point", "y", "params_5"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2140,13 +2016,13 @@
     {
       "id": "rank2_by_rank3_b_zero_point_only",
       "when": ["not present.a_zero_point", "present.b_zero_point", "ranks.a == 2 and ranks.b == 3 and ranks.y == 3 and dim(shapes.a, 1) == dim(shapes.b, 1) and dim(shapes.y, 0) == dim(shapes.b, 0) and dim(shapes.y, 1) == dim(shapes.a, 0) and dim(shapes.y, 2) == dim(shapes.b, 2)", "scalarBZeroOk"],
-      "derive": { "hasAZero": false, "hasBZero": true, "bZeroPerColumn": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched2x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
-          "bindings": ["a_2", "b", "b_zero_point", "y", "params_5"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2158,13 +2034,13 @@
     {
       "id": "rank2_by_rank3_no_zero_points",
       "when": ["not present.a_zero_point", "not present.b_zero_point", "ranks.a == 2 and ranks.b == 3 and ranks.y == 3 and dim(shapes.a, 1) == dim(shapes.b, 1) and dim(shapes.y, 0) == dim(shapes.b, 0) and dim(shapes.y, 1) == dim(shapes.a, 0) and dim(shapes.y, 2) == dim(shapes.b, 2)", "true"],
-      "derive": { "hasAZero": false, "hasBZero": false, "bZeroPerColumn": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched2x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
-          "bindings": ["a_2", "b", "y", "params_5"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2183,7 +2059,7 @@
           "name": "MatMulInteger.Batched4",
           "shader": "quant-matmul-accumulate-rank4.wgsl.jinja",
           "derive": { "outputToScratch": false },
-          "bindings": ["a_2", "b", "a_zero_point", "b_zero_point", "y", "params_6"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2202,7 +2078,7 @@
           "name": "MatMulInteger.Batched4",
           "shader": "quant-matmul-accumulate-rank4.wgsl.jinja",
           "derive": { "outputToScratch": false },
-          "bindings": ["a_2", "b", "a_zero_point", "y", "params_6"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2221,7 +2097,7 @@
           "name": "MatMulInteger.Batched4",
           "shader": "quant-matmul-accumulate-rank4.wgsl.jinja",
           "derive": { "outputToScratch": false },
-          "bindings": ["a_2", "b", "b_zero_point", "y", "params_6"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2240,7 +2116,7 @@
           "name": "MatMulInteger.Batched4",
           "shader": "quant-matmul-accumulate-rank4.wgsl.jinja",
           "derive": { "outputToScratch": false },
-          "bindings": ["a_2", "b", "y", "params_6"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
@@ -2258,8 +2134,6 @@
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
         "tileM": 8,
         "tileN": 64,
         "ktWords": 32
@@ -2269,7 +2143,7 @@
           "id": "main",
           "name": "MatMulInteger.Dp4aBatched3x3",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
-          "bindings": ["a_3", "b", "a_zero_point", "b_zero_point", "y", "params_7"],
           "dispatch": { "x": "ceil(dim(shapes.y, 2) / 64)", "y": "ceil(dim(shapes.y, 1) / 8)", "z": "dim(shapes.y, 0)" }
         }
       ]
@@ -2283,8 +2157,6 @@
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
         "tileM": 8,
         "tileN": 64,
         "ktWords": 32
@@ -2294,7 +2166,7 @@
           "id": "main",
           "name": "MatMulInteger.Dp4aBatched4Scalar",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
-          "bindings": ["a_3", "b", "a_zero_point", "b_zero_point", "y", "params_8"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / 64)",
             "y": "ceil(dim(shapes.y, 2) / 8)",
@@ -2312,8 +2184,6 @@
         "hasAZero": false,
         "hasBZero": false,
         "bZeroPerColumn": false,
-        "requant": "\"none\"",
-        "scaleScalar": "\"f32\"",
         "tileM": 8,
         "tileN": 64,
         "ktWords": 32
@@ -2323,7 +2193,7 @@
           "id": "main",
           "name": "MatMulInteger.Dp4aBatched4NoZero",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
-          "bindings": ["a_3", "b", "y", "params_8"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / 64)",
             "y": "ceil(dim(shapes.y, 2) / 8)",

     "EXACT_MATRIX_MIN_INVOCATIONS": { "default": 16384 }
   },
   "derive": {
+    "precastRows": "dim(shapes.a, 0)",
+    "precastContraction": "dim(shapes.a, 1)",
+    "precastCols": "dim(shapes.b, 1)",
+    "precastBatches": 1,
+    "sgmatTileRows": "dim(shapes.a, 0)",
+    "aUnsigned": "tensorDtypes.a == \"uint8\"",
+    "bUnsigned": "tensorDtypes.b == \"uint8\"",
+    "aZeroPerRow": false,
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "wave32Adapter": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 32 and device.adapterInfo.subgroupMaxSize == 32",
     "canPinSubgroupSize32": "device.features.has(\"subgroups\") and device.features.has(\"subgroup-size-control\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize <= 32 and device.adapterInfo.subgroupMaxSize >= 32",
     "pinSubgroupSize32": "canPinSubgroupSize32 and not wave32Adapter",
     "wave32Effective": "wave32Adapter or pinSubgroupSize32",
+    "variableSubgroup16To32": "device.features.has(\"subgroups\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 16 and device.adapterInfo.subgroupMaxSize == 32",
     "sgmatTileN": "64",
     "sgmatTileK": "32",
     "sgmatColumnGroups": "2",
     "sgmatStorageRowGroups": "floor((device.limits.maxComputeWorkgroupStorageSize - sgmatTileN * sgmatTileK * 4) / (16 * sgmatTileK * 4))",
+    "sgmatRowGroups": "max(1, min(4, floor(deviceWorkgroupCap / (32 * sgmatColumnGroups)), sgmatStorageRowGroups, max(floor(sgmatStorageRowGroups / 2), floor(sgmatTileRows / 16))))",
     "sgmatTileM": "16 * sgmatRowGroups",
     "sgmatSubgroups": "sgmatRowGroups * sgmatColumnGroups",
     "sgmatWorkgroup": "32 * sgmatSubgroups",
     "sgmatReuseStaging": "sgmatStagingBytes + sgmatScratchBytes > device.limits.maxComputeWorkgroupStorageSize",
     "sgmatSharedBytes": "sgmatStagingBytes + (0 if sgmatReuseStaging else sgmatScratchBytes)",
     "sgmatDeviceOk": "sgmatWorkgroup <= deviceWorkgroupCap and sgmatSharedBytes <= device.limits.maxComputeWorkgroupStorageSize",
+    "precastSubgroups": "max(1, min(4, ceilDiv(precastRows, 16), floor(deviceWorkgroupCap / 32), floor(device.limits.maxComputeWorkgroupStorageSize / (8 * 64 * 4))))",
     "precastTileM": "16 * precastSubgroups",
     "precastTileN": "4 * 8",
     "precastTileK": "4 * 8",
     "precastWorkgroup": "32 * precastSubgroups",
+    "preparedK": "ceilDiv(precastContraction, precastTileK) * precastTileK",
+    "preparedN": "ceilDiv(precastCols, precastTileN) * precastTileN",
+    "preparedElements": "ceilDiv(precastRows, precastTileM) * precastTileM * preparedK",
+    "preparedBElements": "precastBatches * preparedK * preparedN",
     "prepareWorkgroup": "min(256, deviceWorkgroupCap)",
     "prepareTotalElements": "preparedElements + preparedBElements",
     "prepareGroups": "ceilDiv(prepareTotalElements, 4 * prepareWorkgroup)",
     "prepareDispatchElements": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535)) * min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535)) * prepareWorkgroup * 4",
     "precastResourcesOk": "deviceWorkgroupCap >= precastWorkgroup and precastSubgroups * 8 * 64 * 4 <= device.limits.maxComputeWorkgroupStorageSize",
+    "precastDispatchOk": "prepareTotalElements <= 4294967295 and prepareDispatchElements <= 4294967296 and ceilDiv(precastCols, precastTileN) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and ceilDiv(precastRows, precastTileM) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and prepareGroups <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "precastF32BuffersOk": "preparedElements * 4 <= device.limits.maxStorageBufferBindingSize and preparedElements * 4 <= device.limits.maxBufferSize and preparedBElements * 4 <= device.limits.maxStorageBufferBindingSize and preparedBElements * 4 <= device.limits.maxBufferSize",
     "precastF16BuffersOk": "preparedElements * 2 <= device.limits.maxStorageBufferBindingSize and preparedElements * 2 <= device.limits.maxBufferSize and preparedBElements * 2 <= device.limits.maxStorageBufferBindingSize and preparedBElements * 2 <= device.limits.maxBufferSize",
     "precastWorthwhile": "dim(shapes.a, 0) >= 16 and dim(shapes.b, 1) >= 32 and dim(shapes.a, 1) >= 32 and floor(dim(shapes.a, 0) / 16) * floor(dim(shapes.b, 1) / 32) * 32 >= tunables.EXACT_MATRIX_MIN_INVOCATIONS",
     "portableWorkgroupCols": 16,
+    "portableWorkgroupRows": "pow2ceil(min(16, max(4, pow2ceil(ceilDiv(precastRows, 4))), floor(deviceWorkgroupCap / portableWorkgroupCols), device.limits.maxComputeWorkgroupSizeY) + 1) / 2",
     "portableTileM": "4 * portableWorkgroupRows",
     "portableTileN": "4 * portableWorkgroupCols",
     "portableTileK": 32,
     "portablePartialK": "floor(16777216 / (255 * 255) / portableTileK) * portableTileK",
+    "portableGeometryOk": "portableWorkgroupRows >= 4 and portableWorkgroupCols <= device.limits.maxComputeWorkgroupSizeX and portablePartialK >= portableTileK and ceilDiv(precastRows, portableTileM) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "rank2MatMulShapeOk": "ranks.a == 2 and ranks.b == 2 and ranks.y == 2 and dim(shapes.a, 1) == dim(shapes.b, 0) and dim(shapes.y, 0) == dim(shapes.a, 0) and dim(shapes.y, 1) == dim(shapes.b, 1)",
     "rank3ByRank3ShapeOk": "ranks.a == 3 and ranks.b == 3 and ranks.y == 3 and dim(shapes.a, 2) == dim(shapes.b, 1) and dim(shapes.y, 0) == max(dim(shapes.a, 0), dim(shapes.b, 0)) and (dim(shapes.a, 0) == dim(shapes.y, 0) or dim(shapes.a, 0) == 1) and (dim(shapes.b, 0) == dim(shapes.y, 0) or dim(shapes.b, 0) == 1) and dim(shapes.y, 1) == dim(shapes.a, 1) and dim(shapes.y, 2) == dim(shapes.b, 2)",
     "rank3ByRank2ShapeOk": "ranks.a == 3 and ranks.b == 2 and ranks.y == 3 and dim(shapes.a, 2) == dim(shapes.b, 0) and dim(shapes.y, 0) == dim(shapes.a, 0) and dim(shapes.y, 1) == dim(shapes.a, 1) and dim(shapes.y, 2) == dim(shapes.b, 1)",
     "requiredScalarBZeroOk": "present.b_zero_point and scalarBZeroOk",
     "rank2PerColumnBZeroOk": "ranks.b_zero_point == 1 and dim(shapes.b_zero_point, 0) == dim(shapes.y, 1)",
     "rank2Dp4aLargePrefill": "ranks.y == 2 and ranks.a == 2 and dim(shapes.y, 0) >= tunables.DP4A_LARGE_MIN_M and dim(shapes.y, 1) >= tunables.DP4A_LARGE_MIN_N and dim(shapes.a, 1) >= tunables.DP4A_LARGE_MIN_K",
+    "perColumnDp4aUseNarrowTile": "variableSubgroup16To32 and rank2Dp4aLargePrefill and not present.a_zero_point",
     "rank2Dp4aWideShortK": "ranks.y == 2 and ranks.a == 2 and dim(shapes.y, 0) >= tunables.DP4A_SHORT_K_MIN_MN and dim(shapes.y, 1) >= tunables.DP4A_SHORT_K_MIN_MN and dim(shapes.a, 1) <= tunables.DP4A_SHORT_K_MAX",
     "rank2Dp4aMediumK": "ranks.y == 2 and ranks.a == 2 and dim(shapes.y, 0) >= tunables.DP4A_MEDIUM_K_MIN_MN and dim(shapes.y, 1) >= tunables.DP4A_MEDIUM_K_MIN_MN and dim(shapes.a, 1) > tunables.DP4A_SHORT_K_MAX and dim(shapes.a, 1) <= tunables.DP4A_MEDIUM_K_MAX",
     "rank2Dp4aTileM": "1 if ranks.y < 2 or dim(shapes.y, 0) == 1 else (tunables.DP4A_SHORT_K_TILE_M if rank2Dp4aWideShortK else (tunables.DP4A_MEDIUM_K_TILE_M if rank2Dp4aMediumK else (tunables.DP4A_LARGE_TILE_M if rank2Dp4aLargePrefill else tunables.DP4A_TILE_M)))",
     "rank4YBatchStride": "dim(shapes.y, 2) * dim(shapes.y, 3) if ranks.y >= 4 else 1",
     "aScalar": "\"u32\" if tensorDtypes.a == \"uint8\" else \"i32\"",
     "bScalar": "\"u32\" if tensorDtypes.b == \"uint8\" else \"i32\"",
     "sgmatM": "matrixM",
     "sgmatN": "matrixN",
     "sgmatK": "matrixK",
     "sgmatDispatchOk": "ceilDiv(sgmatN, sgmatTileN) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and ceilDiv(sgmatM, sgmatTileM) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "sgmatParallelInvocations": "floor(sgmatM / sgmatTileM) * floor(sgmatN / sgmatTileN) * sgmatWorkgroup",
+    "sgmatWorthwhile": "sgmatM >= sgmatTileM and sgmatN >= sgmatTileN / 2 and sgmatK >= sgmatTileK and sgmatParallelInvocations >= tunables.EXACT_MATRIX_MIN_INVOCATIONS",
+    "preferPackedDot": "variableSubgroup16To32 and device.wgslLanguageFeatures.has(\"packed_4x8_integer_dot_product\") and rank2MatMulShapeOk and rank2Dp4aDispatchOk and device.limits.maxComputeInvocationsPerWorkgroup >= 128 and device.limits.maxComputeWorkgroupSizeX >= 128 and device.limits.maxComputeWorkgroupStorageSize >= 8448 and matrixM >= tunables.DP4A_LARGE_MIN_M and matrixN >= tunables.DP4A_LARGE_MIN_N and matrixK > tunables.DP4A_MEDIUM_K_MAX / 2 and matrixK % 4 == 0 and matrixM * matrixN * matrixK >= 4 * tunables.DP4A_LARGE_MIN_M * tunables.DP4A_LARGE_MIN_N * tunables.DP4A_MEDIUM_K_MAX",
+    "tileM": "rank2Dp4aTileM",
+    "tileN": "rank2Dp4aTileN",
+    "ktWords": "rank2Dp4aKtWords"
   },
   "bindings": {
+    "a": { "elementType": "$aStorage" },
+    "b": { "elementType": "$bScalar" },
+    "a_zero_point": { "elementType": "$aScalar", "length": 1 },
+    "b_zero_point": { "elementType": "$bScalar" },
+    "y": { "elementType": "i32" },
     "params": {
       "struct": [
         { "name": "M", "type": "u32", "value": "matrixM" },
         { "name": "N", "type": "u32", "value": "matrixN" },
         { "name": "yBatchStride", "type": "u32", "value": 0 }
       ]
     },
+    "a_main": { "name": "a", "elementType": "$aScalar" },
+    "a_a_vec4": { "name": "a", "elementType": "$aVec4" },
+    "params_main": {
       "name": "params",
       "struct": [
         { "name": "M", "type": "u32", "value": "matrixM" },
         { "name": "N", "type": "u32", "value": "matrixN" },
         { "name": "K", "type": "u32", "value": "matrixK" }
       ]
     },
+    "params_m_n_k": {
       "name": "params",
       "struct": [
         { "name": "M", "type": "u32", "value": 1 },
         { "name": "N", "type": "u32", "value": 1 },
         { "name": "K", "type": "u32", "value": "dim(shapes.a, 0)" }
       ]
     },
+    "params__uniform": {
       "name": "params",
       "struct": [
         { "name": "M", "type": "u32", "value": "batchM" },
         { "name": "N", "type": "u32", "value": "batchN" },
         { "name": "batchCount", "type": "u32", "value": "batchCount" }
       ]
     },
+    "params_rank4_no_zero_points": {
       "name": "params",
       "struct": [
         { "name": "M", "type": "u32", "value": "rank4M" },
         { "name": "N", "type": "u32", "value": "rank4N" },
         }
       ]
     },
+    "params_quant_dp4a_matmul": {
       "name": "params",
       "struct": [
         { "name": "M", "type": "u32", "value": "batchM" },
         { "name": "N", "type": "u32", "value": "batchN" },
         { "name": "yBatchStride", "type": "u32", "value": "batchYStride" }
       ]
     },
+    "params_dp4a_rank4_no_zero_points": {
       "name": "params",
       "struct": [
         { "name": "M", "type": "u32", "value": "rank4M" },
         { "name": "N", "type": "u32", "value": "rank4N" },
     "precast_f32_bFloat": { "name": "bFloat", "buffer": "storage", "elementType": "f32" },
     "precast_prepare_params": {
       "name": "params",
       "struct": [
         { "name": "M", "type": "u32", "value": "sgmatM" },
         { "name": "K", "type": "u32", "value": "sgmatK" },
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f16\"",
         "hasAZero": false,
         "hasBZero": false,
+        "bZeroPerColumn": false
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f16_aFloat", "b", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f16_a", "precast_f16_b", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a0_bnone_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "not present.b_zero_point", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
         "hasAZero": false,
         "hasBZero": false,
         "bZeroPerColumn": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f16_aFloat", "b", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f32\"",
         "hasAZero": false,
         "hasBZero": false,
+        "bZeroPerColumn": false
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f32_aFloat", "b", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f32_a", "precast_f32_b", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a0_bnone",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "not present.b_zero_point", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
         "hasAZero": false,
         "hasBZero": false,
         "bZeroPerColumn": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f32_aFloat", "b", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f16\"",
         "hasAZero": false,
         "hasBZero": true,
+        "bZeroPerColumn": false
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f16_a", "precast_f16_b", "b_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a0_bscalar_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "requiredScalarBZeroOk", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f32\"",
         "hasAZero": false,
         "hasBZero": true,
+        "bZeroPerColumn": false
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f32_a", "precast_f32_b", "b_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a0_bscalar",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "requiredScalarBZeroOk", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f16\"",
         "hasAZero": false,
         "hasBZero": true,
+        "bZeroPerColumn": true
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f16_a", "precast_f16_b", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a0_bcolumn_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "present.b_zero_point and rank2PerColumnBZeroOk", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": true,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f32\"",
         "hasAZero": false,
         "hasBZero": true,
+        "bZeroPerColumn": true
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f32_a", "precast_f32_b", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a0_bcolumn",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "not present.a_zero_point", "present.b_zero_point and rank2PerColumnBZeroOk", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": true,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f16\"",
         "hasAZero": true,
         "hasBZero": false,
+        "bZeroPerColumn": false
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f16_aFloat", "b", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f16_a", "precast_f16_b", "a_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a1_bnone_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "not present.b_zero_point", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
         "hasAZero": true,
         "hasBZero": false,
         "bZeroPerColumn": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f16_aFloat", "b", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f32\"",
         "hasAZero": true,
         "hasBZero": false,
+        "bZeroPerColumn": false
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f32_aFloat", "b", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f32_a", "precast_f32_b", "a_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a1_bnone",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "not present.b_zero_point", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
         "hasAZero": true,
         "hasBZero": false,
         "bZeroPerColumn": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f32_aFloat", "b", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f16\"",
         "hasAZero": true,
         "hasBZero": true,
+        "bZeroPerColumn": false
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f16_a", "precast_f16_b", "a_zero_point", "b_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a1_bscalar_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "requiredScalarBZeroOk", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f32\"",
         "hasAZero": true,
         "hasBZero": true,
+        "bZeroPerColumn": false
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f32_a", "precast_f32_b", "a_zero_point", "b_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a1_bscalar",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "requiredScalarBZeroOk", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f16\"",
         "hasAZero": true,
         "hasBZero": true,
+        "bZeroPerColumn": true
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f16_a", "precast_f16_b", "a_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a1_bcolumn_f16",
       "priority": 39,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "present.b_zero_point and rank2PerColumnBZeroOk", "precastF16BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float16\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f16\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": ["shader-f16"] },
       "derive": {
         "precastScalar": "\"f16\"",
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": true,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f16_aFloat", "b", "b_zero_point", "precast_f16_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f16_a", "portable_f16_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "sgmatRowGroups": "precastSubgroups",
         "sgmatSubgroups": "precastSubgroups",
         "sgmatWorkgroup": "precastWorkgroup",
         "precastScalar": "\"f32\"",
         "hasAZero": true,
         "hasBZero": true,
+        "bZeroPerColumn": true
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactMultiply",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["precast_f32_a", "precast_f32_b", "a_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(sgmatN, sgmatTileN)", "y": "ceilDiv(sgmatM, sgmatTileM)" }
         }
       ],
       "id": "portable_precast_a1_bcolumn",
       "priority": 38,
       "when": ["rank2MatMulShapeOk", "requiredScalarAZeroOk", "present.b_zero_point and rank2PerColumnBZeroOk", "precastF32BuffersOk", "precastDispatchOk", "precastWorthwhile", "portableGeometryOk", "(portableTileM + portableTileN) * portableTileK * dtypeBytes(\"float32\") <= device.limits.maxComputeWorkgroupStorageSize", "\"f32\" == \"f16\" or sgmatM >= portableTileN"],
+      "demoteWhen": ["preferPackedDot"],
       "requires": { "features": [] },
       "derive": {
         "precastScalar": "\"f32\"",
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": true,
         "M": "sgmatM",
         "N": "sgmatN",
         "K": "preparedK",
         "workgroupRows": "portableWorkgroupRows",
         "blockRows": "portableTileM",
         "blockCols": "portableTileN",
+        "integerPartialK": "portablePartialK"
       },
       "passes": [
         {
           "id": "prepare",
+          "name": "MatMulInteger.ExactPrepareOperands",
           "shader": "quant-exact-prepare.wgsl.jinja",
+          "bindings": ["a_main", "a_zero_point", "precast_f32_aFloat", "b", "b_zero_point", "precast_f32_bFloat", "precast_prepare_params"],
           "dispatch": {
             "x": "min(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
             "y": "ceilDiv(prepareGroups, min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
         },
         {
           "id": "main",
+          "name": "MatMulInteger.ExactPortableMultiply",
           "shader": "quant-exact-portable.wgsl.jinja",
           "bindings": ["portable_f32_a", "portable_f32_b", "y"],
           "dispatch": { "x": "ceilDiv(sgmatN, portableTileN)", "y": "ceilDiv(sgmatM, portableTileM)" }
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
+      "derive": { "hasAZero": false, "hasBZero": false, "bZeroPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["a_main", "b", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
+      "derive": { "hasAZero": false, "hasBZero": true, "bZeroPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["a_main", "b", "b_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
+      "derive": { "hasAZero": false, "hasBZero": true, "bZeroPerColumn": true },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["a_main", "b", "b_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
+      "derive": { "hasAZero": true, "hasBZero": false, "bZeroPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
+      "derive": { "hasAZero": true, "hasBZero": true, "bZeroPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "b_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [{ "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }]
       },
+      "derive": { "hasAZero": true, "hasBZero": true, "bZeroPerColumn": true },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.ExactMatrix",
           "shader": "quant-exact-matrix.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "b_zero_point", "y", "params_main"],
           "dispatch": { "x": "ceilDiv(matrixN, sgmatTileN)", "y": "ceilDiv(matrixM, sgmatTileM)" }
         }
       ]
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
+        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL"
       },
       "passes": [
         {
         "hasAZero": false,
         "hasBZero": false,
         "bZeroPerColumn": false,
+        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL"
       },
       "passes": [
         {
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
         "tileM": 8,
         "tileN": 64,
         "ktWords": 32
           "name": "MatMulInteger.Dp4aBatched3x2Tail",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
           "bindings": [
+            "a_main",
             "b",
             "a_zero_point",
             "b_zero_point",
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
+        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL"
       },
       "passes": [
         {
         "hasAZero": true,
         "hasBZero": false,
         "bZeroPerColumn": false,
+        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL"
       },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Dp4a",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_a_vec4", "b", "a_zero_point", "y", "params"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.y, 1), rank2Dp4aTileN)",
             "y": "ceilDiv(dim(shapes.y, 0), rank2Dp4aTileM)"
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": false,
+        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL"
       },
       "passes": [
         {
         "hasAZero": false,
         "hasBZero": false,
         "bZeroPerColumn": false,
+        "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL"
       },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Dp4a",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_a_vec4", "b", "y", "params"],
           "dispatch": {
             "x": "ceilDiv(dim(shapes.y, 1), rank2Dp4aTileN)",
             "y": "ceilDiv(dim(shapes.y, 0), rank2Dp4aTileM)"
         "hasAZero": false,
         "hasBZero": true,
         "bZeroPerColumn": true,
         "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL",
+        "tileM": "rank2Dp4aTileM if dim(shapes.y, 0) == 1 else (16 if perColumnDp4aUseNarrowTile else tunables.TILE_M)",
+        "tileN": "rank2Dp4aTileN if dim(shapes.y, 0) == 1 else (2 * tunables.TILE_N if variableSubgroup16To32 and rank2Dp4aLargePrefill and 2 * tunables.TILE_N <= deviceWorkgroupCap else tunables.TILE_N)",
         "ktWords": "rank2Dp4aKtWords if dim(shapes.y, 0) == 1 else tunables.KT_WORDS"
       },
       "passes": [
           "name": "MatMulInteger.Dp4aBPerColumn",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
           "bindings": ["a", "b", "b_zero_point", "y", "params"],
+          "dispatch": { "x": "ceilDiv(dim(shapes.y, 1), tileN)", "y": "ceilDiv(dim(shapes.y, 0), tileM)" }
         }
       ]
     },
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": true,
         "gemvKUnroll": "tunables.DP4A_GEMV_K_UNROLL",
+        "tileM": "rank2Dp4aTileM if dim(shapes.y, 0) == 1 else (16 if perColumnDp4aUseNarrowTile else tunables.TILE_M)",
+        "tileN": "rank2Dp4aTileN if dim(shapes.y, 0) == 1 else (2 * tunables.TILE_N if variableSubgroup16To32 and rank2Dp4aLargePrefill and 2 * tunables.TILE_N <= deviceWorkgroupCap else tunables.TILE_N)",
         "ktWords": "rank2Dp4aKtWords if dim(shapes.y, 0) == 1 else tunables.KT_WORDS"
       },
       "passes": [
           "name": "MatMulInteger.Dp4aAScalarBPerColumn",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
           "bindings": ["a", "b", "a_zero_point", "b_zero_point", "y", "params"],
+          "dispatch": { "x": "ceilDiv(dim(shapes.y, 1), tileN)", "y": "ceilDiv(dim(shapes.y, 0), tileM)" }
         }
       ]
     },
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "b_zero_point", "y", "params_main"],
           "dispatch": {
             "x": "ceil(matrixN / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((matrixM), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
+            "z": 1
           }
         }
       ]
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "y", "params_main"],
           "dispatch": {
             "x": "ceil(matrixN / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((matrixM), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
+            "z": 1
           }
         }
       ]
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
+          "bindings": ["a_main", "b", "b_zero_point", "y", "params_main"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
+            "z": 1
           }
         }
       ]
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
+          "bindings": ["a_main", "b", "b_zero_point", "y", "params_m_n_k"],
           "dispatch": { "x": 1 }
         }
       ]
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
+          "bindings": ["a_main", "b", "y", "params_m_n_k"],
           "dispatch": { "x": 1 }
         }
       ]
           "id": "main",
           "name": "MatMulInteger",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
+          "bindings": ["a_main", "b", "y", "params_main"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
+            "z": 1
           }
         }
       ]
           "id": "main",
           "name": "MatMulInteger.PerColumn",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "b_zero_point", "y", "params_main"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
+            "z": 1
           }
         }
       ]
           "id": "main",
           "name": "MatMulInteger.PerColumnNoAZero",
           "shader": "quant-matmul-accumulate-rank2.wgsl.jinja",
+          "bindings": ["a_main", "b", "b_zero_point", "y", "params_main"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceilDiv((dim(shapes.y, 0)), (tunables.WORKGROUP_SIZE_Y)), DISPATCH_FOLD_WIDTH)",
+            "z": 1
           }
         }
       ]
     {
       "id": "rank3_by_rank3_scalar_zero_points",
       "when": ["present.a_zero_point", "present.b_zero_point", "(rank3ByRank3ShapeOk or rank3ByRank2ShapeOk)", "scalarAZeroOk and scalarBZeroOk"],
+      "derive": { "hasAZero": true, "hasBZero": true, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "b_zero_point", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
     {
       "id": "rank3_by_rank3_no_zero_points",
       "when": ["not present.a_zero_point", "not present.b_zero_point", "(rank3ByRank3ShapeOk or rank3ByRank2ShapeOk)", "true"],
+      "derive": { "hasAZero": false, "hasBZero": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
+          "bindings": ["a_main", "b", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
     {
       "id": "rank3_by_rank3_a_zero_point_only",
       "when": ["present.a_zero_point", "not present.b_zero_point", "(rank3ByRank3ShapeOk or rank3ByRank2ShapeOk)", "scalarAZeroOk"],
+      "derive": { "hasAZero": true, "hasBZero": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
     {
       "id": "rank3_by_rank3_b_zero_point_only",
       "when": ["not present.a_zero_point", "present.b_zero_point", "(rank3ByRank3ShapeOk or rank3ByRank2ShapeOk)", "scalarBZeroOk"],
+      "derive": { "hasAZero": false, "hasBZero": true, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
+          "bindings": ["a_main", "b", "b_zero_point", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
     {
       "id": "rank3_by_rank3_b_zero_point_per_column_batched",
       "when": ["present.a_zero_point", "present.b_zero_point", "rank3ByRank3ShapeOk", "scalarAZeroOk and ranks.b_zero_point == 3 and dim(shapes.b_zero_point, 0) == dim(shapes.y, 0) and dim(shapes.b_zero_point, 1) == 1 and dim(shapes.b_zero_point, 2) == dim(shapes.y, 2)"],
+      "derive": { "hasAZero": true, "hasBZero": true, "bZeroBatchedPerColumn": true },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3PerColumn",
           "shader": "matmul-integer-batched.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "b_zero_point", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
     {
       "id": "rank3_by_rank3_b_zero_point_per_column_batched_no_azp",
       "when": ["not present.a_zero_point", "present.b_zero_point", "rank3ByRank3ShapeOk", "ranks.b_zero_point == 3 and dim(shapes.b_zero_point, 0) == dim(shapes.y, 0) and dim(shapes.b_zero_point, 1) == 1 and dim(shapes.b_zero_point, 2) == dim(shapes.y, 2)"],
+      "derive": { "hasAZero": false, "hasBZero": true, "bZeroBatchedPerColumn": true },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched3x3PerColumn",
           "shader": "matmul-integer-batched.wgsl.jinja",
+          "bindings": ["a_main", "b", "b_zero_point", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
     {
       "id": "rank2_by_rank3_scalar_zero_points",
       "when": ["present.a_zero_point", "present.b_zero_point", "ranks.a == 2 and ranks.b == 3 and ranks.y == 3 and dim(shapes.a, 1) == dim(shapes.b, 1) and dim(shapes.y, 0) == dim(shapes.b, 0) and dim(shapes.y, 1) == dim(shapes.a, 0) and dim(shapes.y, 2) == dim(shapes.b, 2)", "scalarAZeroOk and scalarBZeroOk"],
+      "derive": { "hasAZero": true, "hasBZero": true, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched2x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "b_zero_point", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
     {
       "id": "rank2_by_rank3_a_zero_point_only",
       "when": ["present.a_zero_point", "not present.b_zero_point", "ranks.a == 2 and ranks.b == 3 and ranks.y == 3 and dim(shapes.a, 1) == dim(shapes.b, 1) and dim(shapes.y, 0) == dim(shapes.b, 0) and dim(shapes.y, 1) == dim(shapes.a, 0) and dim(shapes.y, 2) == dim(shapes.b, 2)", "scalarAZeroOk"],
+      "derive": { "hasAZero": true, "hasBZero": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched2x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
+          "bindings": ["a_main", "b", "a_zero_point", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
     {
       "id": "rank2_by_rank3_b_zero_point_only",
       "when": ["not present.a_zero_point", "present.b_zero_point", "ranks.a == 2 and ranks.b == 3 and ranks.y == 3 and dim(shapes.a, 1) == dim(shapes.b, 1) and dim(shapes.y, 0) == dim(shapes.b, 0) and dim(shapes.y, 1) == dim(shapes.a, 0) and dim(shapes.y, 2) == dim(shapes.b, 2)", "scalarBZeroOk"],
+      "derive": { "hasAZero": false, "hasBZero": true, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched2x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
+          "bindings": ["a_main", "b", "b_zero_point", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
     {
       "id": "rank2_by_rank3_no_zero_points",
       "when": ["not present.a_zero_point", "not present.b_zero_point", "ranks.a == 2 and ranks.b == 3 and ranks.y == 3 and dim(shapes.a, 1) == dim(shapes.b, 1) and dim(shapes.y, 0) == dim(shapes.b, 0) and dim(shapes.y, 1) == dim(shapes.a, 0) and dim(shapes.y, 2) == dim(shapes.b, 2)", "true"],
+      "derive": { "hasAZero": false, "hasBZero": false, "bZeroBatchedPerColumn": false },
       "passes": [
         {
           "id": "main",
           "name": "MatMulInteger.Batched2x3",
           "shader": "matmul-integer-batched.wgsl.jinja",
+          "bindings": ["a_main", "b", "y", "params__uniform"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 1) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
           "name": "MatMulInteger.Batched4",
           "shader": "quant-matmul-accumulate-rank4.wgsl.jinja",
           "derive": { "outputToScratch": false },
+          "bindings": ["a_main", "b", "a_zero_point", "b_zero_point", "y", "params_rank4_no_zero_points"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
           "name": "MatMulInteger.Batched4",
           "shader": "quant-matmul-accumulate-rank4.wgsl.jinja",
           "derive": { "outputToScratch": false },
+          "bindings": ["a_main", "b", "a_zero_point", "y", "params_rank4_no_zero_points"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
           "name": "MatMulInteger.Batched4",
           "shader": "quant-matmul-accumulate-rank4.wgsl.jinja",
           "derive": { "outputToScratch": false },
+          "bindings": ["a_main", "b", "b_zero_point", "y", "params_rank4_no_zero_points"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
           "name": "MatMulInteger.Batched4",
           "shader": "quant-matmul-accumulate-rank4.wgsl.jinja",
           "derive": { "outputToScratch": false },
+          "bindings": ["a_main", "b", "y", "params_rank4_no_zero_points"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / tunables.WORKGROUP_SIZE_X)",
             "y": "min(ceil(dim(shapes.y, 2) / tunables.WORKGROUP_SIZE_Y), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
         "tileM": 8,
         "tileN": 64,
         "ktWords": 32
           "id": "main",
           "name": "MatMulInteger.Dp4aBatched3x3",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_a_vec4", "b", "a_zero_point", "b_zero_point", "y", "params_quant_dp4a_matmul"],
           "dispatch": { "x": "ceil(dim(shapes.y, 2) / 64)", "y": "ceil(dim(shapes.y, 1) / 8)", "z": "dim(shapes.y, 0)" }
         }
       ]
         "hasAZero": true,
         "hasBZero": true,
         "bZeroPerColumn": false,
         "tileM": 8,
         "tileN": 64,
         "ktWords": 32
           "id": "main",
           "name": "MatMulInteger.Dp4aBatched4Scalar",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_a_vec4", "b", "a_zero_point", "b_zero_point", "y", "params_dp4a_rank4_no_zero_points"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / 64)",
             "y": "ceil(dim(shapes.y, 2) / 8)",
         "hasAZero": false,
         "hasBZero": false,
         "bZeroPerColumn": false,
         "tileM": 8,
         "tileN": 64,
         "ktWords": 32
           "id": "main",
           "name": "MatMulInteger.Dp4aBatched4NoZero",
           "shader": "quant-dp4a-matmul.wgsl.jinja",
+          "bindings": ["a_a_vec4", "b", "y", "params_dp4a_rank4_no_zero_points"],
           "dispatch": {
             "x": "ceil(dim(shapes.y, 3) / 64)",
             "y": "ceil(dim(shapes.y, 2) / 8)",

build/webgpu/matmul-integer-batched.wgsl.jinja CHANGED Viewed

@@ -26,9 +26,8 @@ fn {{ fn_name }}() -> i32 {
   {% else %}
   return 0;
   {% endif %}
-}
-{%- endmacro %}
-{%- macro quant_read_batched_column_zero(fn_name, buffer_name, is_unsigned, present=true, batched_per_column=false) %}
 fn {{ fn_name }}({% if batched_per_column %}batch: u32, col: u32{% endif %}) -> i32 {
   {% if present %}
   {% if batched_per_column %}
@@ -47,9 +46,7 @@ fn {{ fn_name }}({% if batched_per_column %}batch: u32, col: u32{% endif %}) ->
   {% else %}
   return 0;
   {% endif %}
-}
-{%- endmacro %}
 // b_zero_point lookup. Per-column-per-batch (b_zp shape [batch,1,N]) indexes
 // batch*N + col; per-column (shape [N]) indexes col; otherwise scalar [0].

   {% else %}
   return 0;
   {% endif %}
+}{% endmacro %}
+{% macro quant_read_batched_column_zero(fn_name, buffer_name, is_unsigned, present=true, batched_per_column=false) %}
 fn {{ fn_name }}({% if batched_per_column %}batch: u32, col: u32{% endif %}) -> i32 {
   {% if present %}
   {% if batched_per_column %}
   {% else %}
   return 0;
   {% endif %}
+}{% endmacro %}
 // b_zero_point lookup. Per-column-per-batch (b_zp shape [batch,1,N]) indexes
 // batch*N + col; per-column (shape [N]) indexes col; otherwise scalar [0].

build/webgpu/metadata.json CHANGED Viewed

@@ -1,27 +1,27 @@
 {
   "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"],

 {
   "name": "ai.onnx.MatMulInteger",
+  "id": "_ai_onnx_matmulinteger_webgpu_ffd4ac7",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "tGrApvNDS+q1wEoCX8hzU2NZ/wLXNvgC7s/8RnC42bg=",
+      "manifest.json": "YjPF3+WTy+o5h++04mLYvGAGSCb8BRQyuQVbZwo+lwg=",
+      "matmul-integer-batched.wgsl.jinja": "+21TfArUYYejHfn+7LTkSuf3dDTKm9cvvFdVqxq9Tac=",
+      "quant-dp4a-matmul.wgsl.jinja": "TQnfQK5DnJq8MI2GLS2mTKjA/0//3+ffrNddL8exb74=",
+      "quant-exact-matrix.wgsl.jinja": "cUkyPJwYCVavMtTYH2bXieCoV5zVWXfMAkiQyEXh66g=",
+      "quant-exact-portable.wgsl.jinja": "YhuWQbbMqQyiADZIkF1Sv7fbpQEgw55btuB5lPDAq34=",
+      "quant-exact-prepare.wgsl.jinja": "qiO6Zu0/94h+r9sLOFAJYm8dNQL9rv6TiJh/O8DYYq0=",
+      "quant-matmul-accumulate-rank2.wgsl.jinja": "3iUfINfSvyh5QmhBJGtyqbfBxFIsvMsytuXs2xNrxEg=",
+      "quant-matmul-accumulate-rank4.wgsl.jinja": "sj5qOkUomLu0WuRoeRBVcZ1YzLkL5uqHKxFCnM+AsIE=",
+      "test.json": "9MIc0cTdv/jJKQ2uUB0ugz11YwJn+/JkmK1wrfh3eQQ="
     }
   },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
   "webgpu": {
+    "manifestSpec": "2.1",
     "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"],

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

@@ -1,3 +1,19 @@
 {% 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
@@ -19,7 +35,6 @@
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const TILE_M: u32 = {{ tileM }}u;
 const TILE_N: u32 = {{ tileN }}u;
 const KT_WORDS: u32 = {{ ktWords }}u;
@@ -70,7 +85,6 @@ fn pack_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScala
 fn dot4_packed(a_word: u32, b_word: u32) -> i32 {
   return dot4I8Packed(a_word, b_word);
 }
 {% if needsBSum %}
 fn sum4_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScalar }}) -> i32 {
@@ -83,31 +97,11 @@ fn sum4_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScala
 {% endif %}
 {% if hasAZero %}
-fn read_a_zero() -> i32 {
-{% if aUnsigned %}
-  return i32(a_zero_point[0]);
-{% else %}
-  return a_zero_point[0];
-{% endif %}
-}
 {% endif %}
 {% if hasBZero %}
-fn read_b_zero({% if bZeroPerColumn %}col: u32{% endif %}) -> i32 {
-{% if bZeroPerColumn %}
-{% if bUnsigned %}
-  return i32(b_zero_point[col]);
-{% else %}
-  return b_zero_point[col];
-{% endif %}
-{% else %}
-{% if bUnsigned %}
-  return i32(b_zero_point[0]);
-{% else %}
-  return b_zero_point[0];
-{% endif %}
-{% endif %}
-}
 {% endif %}
 @compute @workgroup_size(TILE_N, 1, 1)

+{% macro quant_read_column_zero(fn_name, buffer_name, is_unsigned, present=true, per_column=false, param_name="col", return_type="i32") %}
+fn {{ fn_name }}({% if per_column %}{{ param_name }}: u32{% endif %}) -> {{ return_type }} {
+  {% if per_column %}
+  {% if is_unsigned %}
+  return i32({{ buffer_name }}[{{ param_name }}]);
+  {% else %}
+  return {{ buffer_name }}[{{ param_name }}];
+  {% endif %}
+  {% else %}
+  {% if is_unsigned %}
+  return i32({{ buffer_name }}[0]);
+  {% else %}
+  return {{ buffer_name }}[0];
+  {% endif %}
+  {% endif %}
+}{% endmacro %}
 {% 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
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const TILE_M: u32 = {{ tileM }}u;
 const TILE_N: u32 = {{ tileN }}u;
 const KT_WORDS: u32 = {{ ktWords }}u;
 fn dot4_packed(a_word: u32, b_word: u32) -> i32 {
   return dot4I8Packed(a_word, b_word);
 }
 {% if needsBSum %}
 fn sum4_b(b0: {{ bScalar }}, b1: {{ bScalar }}, b2: {{ bScalar }}, b3: {{ bScalar }}) -> i32 {
 {% endif %}
 {% if hasAZero %}
+{{ quant_read_column_zero("read_a_zero", "a_zero_point", aUnsigned, per_column=aZeroPerRow, param_name="row") }}
 {% endif %}
 {% if hasBZero %}
+{{ quant_read_column_zero("read_b_zero", "b_zero_point", bUnsigned, per_column=bZeroPerColumn) }}
 {% endif %}
 @compute @workgroup_size(TILE_N, 1, 1)

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

@@ -1,13 +1,12 @@
 {% 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
@@ -32,7 +31,7 @@ var<workgroup> scratch: array<array<array<f32, 64>, 4>, {{ sgmatSubgroups }}>;
 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;
@@ -56,7 +55,7 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>,
         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;
       }

+{% set matrixZeroPerRow = false %}
 {% set precastOperands = precastOperands | default(false) %}
 {% set precastScalar = precastScalar | default("f32") %}
 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
 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 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[{{ "m" if matrixZeroPerRow else "0" }}]){% endif %});
         }
         tile_A[index] = value;
       }

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

@@ -1,5 +1,3 @@
-{% 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.
@@ -7,14 +5,14 @@ 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 %}
@@ -23,7 +21,6 @@ fn store_integer_result(value: i32, output_index: u32) {
 {% 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].
@@ -59,8 +56,8 @@ fn main(
   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) %}

 // Both integer-output and requantizing consumers share the vector GEMM body.
 // Its callback stores scalar outputs so padded operands also cover column tails.
   y[output_index] = value;
 }
+{% set exactIntegerPartialsValue = true %}
 {% 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 = "f32" %}
 {% set workgroupColsValue = workgroupCols if workgroupCols is defined else 16 %}
 {% set workgroupRowsValue = workgroupRows if workgroupRows is defined else 16 %}
 {% set laneCount = workgroupColsValue * workgroupRowsValue %}
 {% set threadCols = threadColVectors * 4 %}
 {% set aVectorCount = (blockRowsValue * kTileValue / (4 * laneCount))|int %}
 {% set bVectorCount = (kTileValue * blockColsValue / (4 * laneCount))|int %}
 // High-intensity matrix multiplication for the large, untransposed layout:
 //   Y[M,N] = alpha * A[M,K] @ B[K,N].
   let li = lid.y * {{ workgroupColsValue }}u + lid.x;
   let mBase = wg.y * BM;
   let nBase = wg.x * BN;
+{% set aOff = "" %}
+{% set bOff = "" %}
 {% for i in range(threadRows) %}
 {% for j in range(threadColVectors) %}

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

@@ -1,4 +1,5 @@
-{% 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) {
@@ -11,15 +12,15 @@ fn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index)
     {
       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;

+{% set prepareBatchB = false %}
+{% set prepareZeroPerRow = false %}
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ prepareWorkgroup }})
 fn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) tid: u32) {
     {
       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[{{ "m" if prepareZeroPerRow else "0" }}]){% endif %});
       }
       aFloat[offset + {{ h }}u] = {% if precastScalar == "f16" %}f16(value){% else %}value{% endif %};
     }
 {% endfor %}
   } else {
     let bOffset = offset - params.aCount;
+    let k = {{ "localOffset" if prepareBatchB else "bOffset" }} / params.paddedN;
+    let n = {{ "localOffset" if prepareBatchB else "bOffset" }} % params.paddedN;
 {% for h in range(4) %}
     {
       var value = 0.0;

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

@@ -26,16 +26,15 @@ fn {{ fn_name }}() -> i32 {
   {% else %}
   return 0;
   {% endif %}
-}
-{%- endmacro %}
-{%- macro quant_read_column_zero(fn_name, buffer_name, is_unsigned, present=true, per_column=false) %}
-fn {{ fn_name }}({% if per_column %}col: u32{% endif %}) -> i32 {
   {% if present %}
   {% if per_column %}
   {% if is_unsigned %}
-  return i32({{ buffer_name }}[col]);
   {% else %}
-  return {{ buffer_name }}[col];
   {% endif %}
   {% else %}
   {% if is_unsigned %}
@@ -47,8 +46,8 @@ fn {{ fn_name }}({% if per_column %}col: u32{% endif %}) -> i32 {
   {% else %}
   return 0;
   {% endif %}
-}
-{%- endmacro %}{{ quant_read_scalar_zero("read_a_zero", "a_zero_point", aUnsigned, hasAZero) }}
 {{ quant_read_column_zero("read_b_zero", "b_zero_point", bUnsigned, hasBZero, bZeroPerColumn) }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE_X }}, {{ tunables.WORKGROUP_SIZE_Y }})

   {% else %}
   return 0;
   {% endif %}
+}{% endmacro %}
+{% macro quant_read_column_zero(fn_name, buffer_name, is_unsigned, present=true, per_column=false, param_name="col", return_type="i32") %}
+fn {{ fn_name }}({% if per_column %}{{ param_name }}: u32{% endif %}) -> {{ return_type }} {
   {% if present %}
   {% if per_column %}
   {% if is_unsigned %}
+  return i32({{ buffer_name }}[{{ param_name }}]);
   {% else %}
+  return {{ buffer_name }}[{{ param_name }}];
   {% endif %}
   {% else %}
   {% if is_unsigned %}
   {% else %}
   return 0;
   {% endif %}
+}{% endmacro %}
+{{ quant_read_scalar_zero("read_a_zero", "a_zero_point", aUnsigned, hasAZero) }}
 {{ quant_read_column_zero("read_b_zero", "b_zero_point", bUnsigned, hasBZero, bZeroPerColumn) }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE_X }}, {{ tunables.WORKGROUP_SIZE_Y }})

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

@@ -26,9 +26,7 @@ fn {{ fn_name }}() -> i32 {
   {% else %}
   return 0;
   {% endif %}
-}
-{%- 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) }}

   {% else %}
   return 0;
   {% endif %}
+}{% 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) }}

build/webgpu/test.json CHANGED Viewed

@@ -1040,7 +1040,7 @@
       "name": "percol_b_zero_point_no_azp_below_dp4a_floor",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/matmul_integer_test.cc",
-        "notes": "Per-column weight zero-point with NO a_zero_point below the dp4a floor: K=8 misses the K >= 64 guard and numel(Y)=12 misses the 8192-element guard, so dp4a_rank2_b_zero_point_per_column is ineligible and the naive per-column kernel runs. b_zero_point rank-1 size N=3 == dim(Y,1); the absent a_zero_point defaults to 0."
       },
       "inputs": {
         "a": {
@@ -1083,10 +1083,41 @@
       "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
       "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
     {
       "name": "k258_unaligned_k_scalar_zp_u8u8",
       "provenance": {
-        "notes": "K=258 is intentionally not divisible by four. The large output would otherwise admit packed accumulation, so this case verifies that the scalar route handles the unaligned reduction tail and applies scalar zero points exactly."
       },
       "inputs": {
         "a": {
@@ -1108,7 +1139,7 @@
     {
       "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": {
@@ -1304,7 +1335,7 @@
     {
       "name": "dp4a_tail_k130_rank2_u8u8_scalar_zp_exact",
       "provenance": {
-        "notes": "Exactness lock for the shared DP4A two-lane final word; full words must remain branch-free while raw-bias padding keeps the missing lanes neutral."
       },
       "inputs": {
         "a": {
@@ -1332,7 +1363,7 @@
     {
       "name": "dp4a_tail_k131_rank2_s8s8_scalar_zp_exact",
       "provenance": {
-        "notes": "Exactness lock for the shared signed DP4A three-lane final word and its single raw-zero padding lane."
       },
       "inputs": {
         "a": {
@@ -1356,9 +1387,7 @@
     },
     {
       "name": "dp4a_tail_k129_batched3_by_rank2_exact",
-      "provenance": {
-        "notes": "Compact correctness lock for realistic adapter widths with a one-element K tail; sized to select the packed tail kernel."
-      },
       "inputs": {
         "a": {
           "dtype": "uint8",
@@ -1512,7 +1541,7 @@
     {
       "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": {
@@ -1692,7 +1721,7 @@
     {
       "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": {
@@ -1712,7 +1741,7 @@
     {
       "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": {
@@ -1733,7 +1762,7 @@
     {
       "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": {
@@ -1754,7 +1783,7 @@
     {
       "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": {
@@ -1775,7 +1804,7 @@
     {
       "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": {
@@ -1797,7 +1826,7 @@
     {
       "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": {
@@ -1819,7 +1848,7 @@
     {
       "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": {
@@ -1839,7 +1868,7 @@
     {
       "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": {
@@ -1860,7 +1889,7 @@
     {
       "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": {
@@ -1881,7 +1910,7 @@
     {
       "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": {
@@ -1902,7 +1931,7 @@
     {
       "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": {
@@ -1924,7 +1953,7 @@
     {
       "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": {
@@ -2266,6 +2295,52 @@
       },
       "outputs": { "y": { "dtype": "int32", "shape": [65, 35], "tolerance": 0 } },
       "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     }
   ]
 }

       "name": "percol_b_zero_point_no_azp_below_dp4a_floor",
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/math/matmul_integer_test.cc",
+        "notes": "Per-column B zero points ([1, -2, 2]) with no A zero point (defaults to 0) check a compact M=4, K=8, N=3 integer matrix product."
       },
       "inputs": {
         "a": {
       "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
       "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
     },
+    {
+      "name": "dp4a_percol_no_azp_narrow_m_tile_variable_subgroup",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/matmul_integer_test.cc",
+        "notes": "Per-column B zero points without an A zero point, with M=64 rows, K=256, and N=128 columns, check exact integer output at a narrow row count."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "uint8",
+          "shape": [64, 256],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/dp4a_u8_pattern" } }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [256, 128],
+          "data": { "kind": "cycle", "values": { "$ref": "#/fixtureArrays/dp4a_s8_pattern" } }
+        },
+        "b_zero_point": {
+          "dtype": "int8",
+          "shape": [128],
+          "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,
+        "DP4A_LARGE_MIN_M": 64,
+        "DP4A_LARGE_MIN_N": 128,
+        "DP4A_LARGE_MIN_K": 256
+      }
+    },
     {
       "name": "k258_unaligned_k_scalar_zp_u8u8",
       "provenance": {
+        "notes": "K=258 is not divisible by four; scalar zero points (a=128, b=127) on uint8 by uint8 operands check that the unaligned reduction tail is included exactly."
       },
       "inputs": {
         "a": {
     {
       "name": "large_accumulator_dp4a_k1024_u8u8_maxmag",
       "provenance": {
+        "notes": "M=64, K=1024, and N=128 with uint8 operands at magnitude 255 produce large but int32-representable sums; nonzero zero points (3 and 7) verify the exact quantized result."
       },
       "inputs": {
         "a": {
     {
       "name": "dp4a_tail_k130_rank2_u8u8_scalar_zp_exact",
       "provenance": {
+        "notes": "Reduction length 130 leaves two elements after complete four-element groups; scalar zero points must produce the exact integer result."
       },
       "inputs": {
         "a": {
     {
       "name": "dp4a_tail_k131_rank2_s8s8_scalar_zp_exact",
       "provenance": {
+        "notes": "Reduction length 131 leaves three elements after complete four-element groups; signed operands and scalar zero points must produce the exact integer result."
       },
       "inputs": {
         "a": {
     },
     {
       "name": "dp4a_tail_k129_batched3_by_rank2_exact",
+      "provenance": { "notes": "Compact integer matrix product with a one-element reduction tail and exact output." },
       "inputs": {
         "a": {
           "dtype": "uint8",
     {
       "name": "dp4a_gemv_no_zero_s8s8_m1_k64_n512",
       "provenance": {
+        "notes": "A signed single-row matrix product with K=64 and N=512 omits both zero points; its result must match direct integer accumulation."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
     {
       "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. Integer partials and requantization follow the exact integer contract."
       },
       "inputs": {
         "a": {
       },
       "outputs": { "y": { "dtype": "int32", "shape": [65, 35], "tolerance": 0 } },
       "tunables": { "EXACT_MATRIX_MIN_INVOCATIONS": 0 }
+    },
+    {
+      "name": "variable_subgroup_packed_crossover_no_zero_points",
+      "provenance": {
+        "notes": "A compact wide integer projection with M=64, K=128, N=128 and no zero points on either operand checks exact int8 output at a tile-aligned shape."
+      },
+      "inputs": {
+        "a": {
+          "dtype": "int8",
+          "shape": [64, 128],
+          "data": { "kind": "cycle", "values": [-128, 127, -37, 0, 19, 101] }
+        },
+        "b": {
+          "dtype": "int8",
+          "shape": [128, 128],
+          "data": { "kind": "cycle", "values": [127, -128, 25, -63, 0, 47] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [64, 128], "tolerance": 0 } },
+      "tunables": {
+        "EXACT_MATRIX_MIN_INVOCATIONS": 0,
+        "DP4A_LARGE_MIN_M": 32,
+        "DP4A_LARGE_MIN_N": 64,
+        "DP4A_MEDIUM_K_MAX": 64
+      }
+    },
+    {
+      "name": "variable_subgroup_packed_crossover_scalar_zero_points",
+      "provenance": { "notes": "Checks exact integer output with both activation and weight zero points." },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [64, 128], "data": { "kind": "cycle", "values": [0, 255, 17, 239, 128, 1] } },
+        "b": {
+          "dtype": "int8",
+          "shape": [128, 128],
+          "data": { "kind": "cycle", "values": [-128, 127, -23, 51, 0, 91] }
+        },
+        "a_zero_point": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [127] } },
+        "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,
+        "DP4A_LARGE_MIN_M": 32,
+        "DP4A_LARGE_MIN_N": 64,
+        "DP4A_MEDIUM_K_MAX": 64
+      }
     }
   ]
 }