sync 91d990483a17

Browse files

Files changed (8) hide show

README.md +37 -16
build/webgpu/bench.json +383 -1
build/webgpu/fused-gemm.wgsl.jinja +33 -29
build/webgpu/gemm-subgroup-matrix.wgsl.jinja +123 -40
build/webgpu/gemv-contiguous.wgsl.jinja +158 -0
build/webgpu/manifest.json +364 -207
build/webgpu/metadata.json +31 -8
build/webgpu/test.json +1950 -13

README.md CHANGED Viewed

@@ -18,17 +18,17 @@ See the [ONNX Runtime `FusedGemm` contrib-operator spec](https://github.com/micr
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `A` | `A` | `T` | `2` | — | Left operand, `(M, K)` when `transA` is 0 and `(K, M)` otherwise. | required |
-| `B` | `B` | `T` | `2` | — | Right operand, `(K, N)` when `transB` is 0 and `(N, K)` otherwise. | required |
-| `C` | `C` | `T` | — | — | Optional additive term, unidirectionally broadcastable to `(M, N)`: a scalar, a row `(N)`, a column `(M, 1)`, or the full matrix. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `Y` | `T` | `2` | derived; see description | `act(alpha * A' * B' + beta * C)` with shape `(M, N)`. | required |
 ## Attributes
@@ -36,13 +36,13 @@ Attributes and default values (overridable per request):
 | Attribute | Default | Description |
 | --- | --- | --- |
 | `alpha` | `1` | Scalar multiplier for the product `A' * B'`; the standard default is 1. |
 | `beta` | `1` | Scalar multiplier for `C`; the standard default is 1. |
 | `transA` | `0` | Whether `A` is stored transposed. The standard default is 0. |
 | `transB` | `0` | Whether `B` is stored transposed. The standard default is 0. |
-| `activation` | — | Optional fused activation name. Supported modes are `Relu`, `LeakyRelu`, `Sigmoid`, `Tanh` and `HardSigmoid`; omission applies none. |
-| `activation_alpha` | — | First activation parameter: the slope for `LeakyRelu` or `alpha` for `HardSigmoid`. |
-| `activation_beta` | — | Second activation parameter: `beta` for `HardSigmoid`. |
 ## Type constraints
@@ -50,25 +50,45 @@ Attributes and default values (overridable per request):
 | --- | --- |
 | `T` | `float32`, `float16` |
 ## Device requirements
 Some implementation variants require `subgroup-matrix` and `subgroups`. These are route-specific capabilities, not package-wide requirements; availability also depends on the request shape and dtype.
 ## Files
-- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 - [`fused-gemm.wgsl.jinja`](build/webgpu/fused-gemm.wgsl.jinja)
 - [`gemm-subgroup-matrix.wgsl.jinja`](build/webgpu/gemm-subgroup-matrix.wgsl.jinja)
 ## Use with `@huggingface/kernels`
-The loader derives every required output's shape and logical dtype from the manifest contract and this call.
-It then allocates the result tensors automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.
@@ -76,7 +96,8 @@ Replace each `*Data` placeholder with a typed array containing the corresponding
 import { getKernel } from "@huggingface/kernels";
 const kernel = await getKernel("webgpu-kernels/com.microsoft.FusedGemm", { version: 1 });
-const { Y } = await kernel({ A: { data: AData, shape: [7, 13] }, B: { data: BData, shape: [13, 11] } }, {
-  attrs: { activation: "Relu" },
 });
 ```

 ## Inputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `A` | `T` | `2` | — | Left operand, `(M, K)` when `transA` is 0 and `(K, M)` otherwise. | required |
+| `B` | `T` | `2` | — | Right operand, `(K, N)` when `transB` is 0 and `(N, K)` otherwise. | required |
+| `C` | `T` | — | — | Optional additive term, unidirectionally broadcastable to `(M, N)`: a scalar, a row `(N)`, a column `(M, 1)`, or the full matrix. | optional |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `Y` | `T` | `2` | derived | `act(alpha * A' * B' + beta * C)` with shape `(M, N)`. | required |
 ## Attributes
 | Attribute | Default | Description |
 | --- | --- | --- |
+| `activation` | — | Optional fused activation name. Supported modes are `Relu`, `LeakyRelu`, `Sigmoid`, `Tanh` and `HardSigmoid`; omission applies none. |
+| `activation_alpha` | — | First activation parameter: the slope for `LeakyRelu` or `alpha` for `HardSigmoid`. |
+| `activation_beta` | — | Second activation parameter: `beta` for `HardSigmoid`. |
 | `alpha` | `1` | Scalar multiplier for the product `A' * B'`; the standard default is 1. |
 | `beta` | `1` | Scalar multiplier for `C`; the standard default is 1. |
 | `transA` | `0` | Whether `A` is stored transposed. The standard default is 0. |
 | `transB` | `0` | Whether `B` is stored transposed. The standard default is 0. |
 ## Type constraints
 | --- | --- |
 | `T` | `float32`, `float16` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `gemv_plain` — Reduces contiguous weight rows for a single output row in f32, applies alpha and activation, and casts once. Four aligned columns share the input vector; other widths use one column. The reduction width is bounded by K and device workgroup limits. Either storage orientation of A is supported.
+- `gemv_bias` — Reduces contiguous weight rows for a single output row in f32, applies alpha, broadcast beta-scaled bias and activation, and casts once. Four aligned columns share the input vector; other widths use one column. The reduction width is bounded by K and device workgroup limits. Either storage orientation of A is supported.
+- `notrans_sgmat_bias` — Uses supported subgroup matrices with f32 accumulation and a fused epilogue. Complete non-transposed f32 A tiles load directly from storage; long reductions on that path merge two-tile sums with f32 compensation. Staged A clamps row tails. Workgroup memory is checked against device limits.
+- `notrans_sgmat` — Uses supported subgroup matrices with f32 accumulation and a fused epilogue. Complete non-transposed f32 A tiles load directly from storage; long reductions on that path merge two-tile sums with f32 compensation. Staged A clamps row tails. Workgroup memory is checked against device limits.
+- `transA_sgmat_bias` — Uses supported subgroup matrices with f32 accumulation and a fused epilogue. Complete non-transposed f32 A tiles load directly from storage; long reductions on that path merge two-tile sums with f32 compensation. Staged A clamps row tails. Workgroup memory is checked against device limits.
+- `transA_sgmat` — Uses supported subgroup matrices with f32 accumulation and a fused epilogue. Complete non-transposed f32 A tiles load directly from storage; long reductions on that path merge two-tile sums with f32 compensation. Staged A clamps row tails. Workgroup memory is checked against device limits.
+- `transB_sgmat_bias` — Uses supported subgroup matrices with f32 accumulation and a fused epilogue. Complete non-transposed f32 A tiles load directly from storage; long reductions on that path merge two-tile sums with f32 compensation. Staged A clamps row tails. Workgroup memory is checked against device limits.
+- `transB_sgmat` — Uses supported subgroup matrices with f32 accumulation and a fused epilogue. Complete non-transposed f32 A tiles load directly from storage; long reductions on that path merge two-tile sums with f32 compensation. Staged A clamps row tails. Workgroup memory is checked against device limits.
+- `transAB_sgmat_bias` — Uses supported subgroup matrices with f32 accumulation and a fused epilogue. Complete non-transposed f32 A tiles load directly from storage; long reductions on that path merge two-tile sums with f32 compensation. Staged A clamps row tails. Workgroup memory is checked against device limits.
+- `transAB_sgmat` — Uses supported subgroup matrices with f32 accumulation and a fused epilogue. Complete non-transposed f32 A tiles load directly from storage; long reductions on that path merge two-tile sums with f32 compensation. Staged A clamps row tails. Workgroup memory is checked against device limits.
 ## Device requirements
 Some implementation variants require `subgroup-matrix` and `subgroups`. These are route-specific capabilities, not package-wide requirements; availability also depends on the request shape and dtype.
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
 - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 - [`fused-gemm.wgsl.jinja`](build/webgpu/fused-gemm.wgsl.jinja)
 - [`gemm-subgroup-matrix.wgsl.jinja`](build/webgpu/gemm-subgroup-matrix.wgsl.jinja)
+- [`gemv-contiguous.wgsl.jinja`](build/webgpu/gemv-contiguous.wgsl.jinja)
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.
 import { getKernel } from "@huggingface/kernels";
 const kernel = await getKernel("webgpu-kernels/com.microsoft.FusedGemm", { version: 1 });
+const { Y } = await kernel({
+  A: { data: AData, shape: [32, 224] },
+  B: { data: BData, shape: [224, 64] },
 });
 ```

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.FusedGemm",
   "cases": [
     {
       "name": "fusedgemm-relu-m3072-k768-n3072",
@@ -64,6 +63,389 @@
       },
       "outputs": { "Y": { "shape": [2048, 2048], "dtype": "float32" } },
       "bench": { "metrics": [{ "type": "gflops", "value": "2 * 2048 * 1024 * 2048" }] }
     }
   ]
 }

 {
   "cases": [
     {
       "name": "fusedgemm-relu-m3072-k768-n3072",
       },
       "outputs": { "Y": { "shape": [2048, 2048], "dtype": "float32" } },
       "bench": { "metrics": [{ "type": "gflops", "value": "2 * 2048 * 1024 * 2048" }] }
+    },
+    {
+      "name": "fusedgemm-f16-transb-relu-m1-k4096-n4096-sgmat-floor-pathology",
+      "preset": "model",
+      "vars": { "dtype": "float16" },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [1, 4096], "dtype": "float16", "dist": "normal", "seed": 5350, "scale": 1 },
+        "B": { "shape": [4096, 4096], "dtype": "float16", "dist": "normal", "seed": 5351, "scale": 1 },
+        "C": { "shape": [4096], "dtype": "float16", "dist": "normal", "seed": 5352, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [1, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * 1 * 4096 * 4096" },
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.A) + numel(shapes.B) + numel(shapes.Y) + numel(shapes.C))",
+            "provenance": "logical"
+          }
+        ]
+      }
+    },
+    {
+      "name": "fusedgemm-f16-transb-relu-m64-k4096-n4096-control",
+      "preset": "model",
+      "vars": { "dtype": "float16" },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [64, 4096], "dtype": "float16", "dist": "normal", "seed": 5380, "scale": 1 },
+        "B": { "shape": [4096, 4096], "dtype": "float16", "dist": "normal", "seed": 5381, "scale": 1 },
+        "C": { "shape": [4096], "dtype": "float16", "dist": "normal", "seed": 5382, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [64, 4096], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 64 * 4096 * 4096" }] }
+    },
+    {
+      "name": "fusedgemm-f16-single-row-k1024-n64-transb1-leakyrelu",
+      "preset": "model",
+      "vars": { "M": 1, "K": 1024, "N": 64, "dtype": "float16" },
+      "attrs": { "transB": 1, "activation": "LeakyRelu" },
+      "inputs": {
+        "A": { "shape": [1, 1024], "dtype": "float16", "dist": "normal", "seed": 5350, "scale": 1 },
+        "B": { "shape": [64, 1024], "dtype": "float16", "dist": "normal", "seed": 5351, "scale": 1 },
+        "C": { "shape": [64], "dtype": "float16", "dist": "normal", "seed": 5352, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [1, 64], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * args.M * args.K * args.N" },
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.A) + numel(shapes.B) + numel(shapes.Y) + numel(shapes.C))",
+            "provenance": "logical"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic benchmark",
+        "notes": "Single-row fused projection spanning reduction length, output width, weight orientation and activation, with f32 accumulation and f16 output."
+      }
+    },
+    {
+      "name": "fusedgemm-f16-single-row-k128-n4096-transb1-hardsigmoid",
+      "preset": "model",
+      "vars": { "M": 1, "K": 128, "N": 4096, "dtype": "float16" },
+      "attrs": { "transB": 1, "activation": "HardSigmoid" },
+      "inputs": {
+        "A": { "shape": [1, 128], "dtype": "float16", "dist": "normal", "seed": 5350, "scale": 1 },
+        "B": { "shape": [4096, 128], "dtype": "float16", "dist": "normal", "seed": 5351, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [1, 4096], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * args.M * args.K * args.N" },
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.A) + numel(shapes.B) + numel(shapes.Y))",
+            "provenance": "logical"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic benchmark",
+        "notes": "Single-row fused projection spanning reduction length, output width, weight orientation and activation, with f32 accumulation and f16 output."
+      }
+    },
+    {
+      "name": "fusedgemm-f16-single-row-k32-n256-transb1-tanh",
+      "preset": "model",
+      "vars": { "M": 1, "K": 32, "N": 256, "dtype": "float16" },
+      "attrs": { "transB": 1, "activation": "Tanh" },
+      "inputs": {
+        "A": { "shape": [1, 32], "dtype": "float16", "dist": "normal", "seed": 5350, "scale": 1 },
+        "B": { "shape": [256, 32], "dtype": "float16", "dist": "normal", "seed": 5351, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [1, 256], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * args.M * args.K * args.N" },
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.A) + numel(shapes.B) + numel(shapes.Y))",
+            "provenance": "logical"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic benchmark",
+        "notes": "Single-row fused projection spanning reduction length, output width, weight orientation and activation, with f32 accumulation and f16 output."
+      }
+    },
+    {
+      "name": "fusedgemm-f16-single-row-k1024-n64-transb0-relu",
+      "preset": "model",
+      "vars": { "M": 1, "K": 1024, "N": 64, "dtype": "float16" },
+      "attrs": { "transB": 0, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [1, 1024], "dtype": "float16", "dist": "normal", "seed": 5350, "scale": 1 },
+        "B": { "shape": [1024, 64], "dtype": "float16", "dist": "normal", "seed": 5351, "scale": 1 },
+        "C": { "shape": [64], "dtype": "float16", "dist": "normal", "seed": 5352, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [1, 64], "dtype": "float16" } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * args.M * args.K * args.N" },
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.A) + numel(shapes.B) + numel(shapes.Y) + numel(shapes.C))",
+            "provenance": "logical"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic benchmark",
+        "notes": "Single-row fused projection spanning reduction length, output width, weight orientation and activation, with f32 accumulation and f16 output."
+      }
+    },
+    {
+      "name": "gemv-float32-k128-n4096-ta0-none-plain",
+      "preset": "model",
+      "vars": { "dtype": "float32" },
+      "attrs": { "transA": 0, "transB": 1 },
+      "inputs": {
+        "A": { "shape": [1, 128], "dtype": "float32", "dist": "normal", "seed": 5350, "scale": 1 },
+        "B": { "shape": [4096, 128], "dtype": "float32", "dist": "normal", "seed": 5351, "scale": 1 }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 4096] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": 1048576 },
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.A) + numel(shapes.B) + numel(shapes.Y))",
+            "provenance": "logical"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic benchmark",
+        "notes": "Single-row fused projection covers dtype, stored A orientation, output alignment, broadcast bias and activation."
+      }
+    },
+    {
+      "name": "gemv-float32-k4096-n4096-ta1-none-Sigmoid",
+      "preset": "model",
+      "vars": { "dtype": "float32" },
+      "attrs": { "transA": 1, "transB": 1, "activation": "Sigmoid" },
+      "inputs": {
+        "A": { "shape": [4096, 1], "dtype": "float32", "dist": "normal", "seed": 5350, "scale": 1 },
+        "B": { "shape": [4096, 4096], "dtype": "float32", "dist": "normal", "seed": 5351, "scale": 1 }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 4096] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": 33554432 },
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.A) + numel(shapes.B) + numel(shapes.Y))",
+            "provenance": "logical"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic benchmark",
+        "notes": "Single-row fused projection covers dtype, stored A orientation, output alignment, broadcast bias and activation."
+      }
+    },
+    {
+      "name": "gemv-float16-k1023-n257-ta0-scalar-LeakyRelu",
+      "preset": "model",
+      "vars": { "dtype": "float16" },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "activation": "LeakyRelu",
+        "alpha": 0.75,
+        "beta": -0.25,
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": { "shape": [1, 1023], "dtype": "float16", "dist": "normal", "seed": 5350, "scale": 1 },
+        "B": { "shape": [257, 1023], "dtype": "float16", "dist": "normal", "seed": 5351, "scale": 1 },
+        "C": { "shape": [], "dtype": "float16", "dist": "normal", "seed": 5352, "scale": 1 }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 257] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": 525822 },
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.A) + numel(shapes.B) + numel(shapes.Y) + numel(shapes.C))",
+            "provenance": "logical"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic benchmark",
+        "notes": "Single-row fused projection covers dtype, stored A orientation, output alignment, broadcast bias and activation."
+      }
+    },
+    {
+      "name": "gemv-float16-k4096-n4096-ta1-scalar-HardSigmoid",
+      "preset": "model",
+      "vars": { "dtype": "float16" },
+      "attrs": { "transA": 1, "transB": 1, "activation": "HardSigmoid" },
+      "inputs": {
+        "A": { "shape": [4096, 1], "dtype": "float16", "dist": "normal", "seed": 5350, "scale": 1 },
+        "B": { "shape": [4096, 4096], "dtype": "float16", "dist": "normal", "seed": 5351, "scale": 1 },
+        "C": { "shape": [], "dtype": "float16", "dist": "normal", "seed": 5352, "scale": 1 }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 4096] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": 33554432 },
+          {
+            "type": "bandwidth",
+            "name": "logical input/output bytes",
+            "value": "dtypeBytes(args.dtype) * (numel(shapes.A) + numel(shapes.B) + numel(shapes.Y) + numel(shapes.C))",
+            "provenance": "logical"
+          }
+        ]
+      },
+      "provenance": {
+        "source": "synthetic benchmark",
+        "notes": "Single-row fused projection covers dtype, stored A orientation, output alignment, broadcast bias and activation."
+      }
+    },
+    {
+      "name": "direct_a_m31_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [31, 96], "dtype": "float32", "dist": "normal", "seed": 5400, "scale": 0.125 },
+        "B": { "shape": [128, 96], "dtype": "float32", "dist": "normal", "seed": 5500, "scale": 0.125 },
+        "C": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 5600, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [31, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "761856" }] }
+    },
+    {
+      "name": "direct_a_m32_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [32, 96], "dtype": "float32", "dist": "normal", "seed": 5401, "scale": 0.125 },
+        "B": { "shape": [128, 96], "dtype": "float32", "dist": "normal", "seed": 5501, "scale": 0.125 },
+        "C": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 5601, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [32, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "786432" }] }
+    },
+    {
+      "name": "direct_a_m33_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [33, 96], "dtype": "float32", "dist": "normal", "seed": 5402, "scale": 0.125 },
+        "B": { "shape": [128, 96], "dtype": "float32", "dist": "normal", "seed": 5502, "scale": 0.125 },
+        "C": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 5602, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [33, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "811008" }] }
+    },
+    {
+      "name": "direct_a_m63_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [63, 96], "dtype": "float32", "dist": "normal", "seed": 5403, "scale": 0.125 },
+        "B": { "shape": [128, 96], "dtype": "float32", "dist": "normal", "seed": 5503, "scale": 0.125 },
+        "C": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 5603, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [63, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "1548288" }] }
+    },
+    {
+      "name": "direct_a_m64_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [64, 96], "dtype": "float32", "dist": "normal", "seed": 5404, "scale": 0.125 },
+        "B": { "shape": [128, 96], "dtype": "float32", "dist": "normal", "seed": 5504, "scale": 0.125 },
+        "C": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 5604, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [64, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "1572864" }] }
+    },
+    {
+      "name": "direct_a_m65_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [65, 96], "dtype": "float32", "dist": "normal", "seed": 5405, "scale": 0.125 },
+        "B": { "shape": [128, 96], "dtype": "float32", "dist": "normal", "seed": 5505, "scale": 0.125 },
+        "C": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 5605, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [65, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "1597440" }] }
+    },
+    {
+      "name": "direct_a_m64_k32_n64_tb0",
+      "attrs": { "transB": 0, "activation": "Tanh" },
+      "inputs": {
+        "A": { "shape": [64, 32], "dtype": "float32", "dist": "normal", "seed": 5406, "scale": 0.125 },
+        "B": { "shape": [32, 64], "dtype": "float32", "dist": "normal", "seed": 5506, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [64, 64], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "262144" }] }
+    },
+    {
+      "name": "direct_a_m32_k224_n64_tb0",
+      "attrs": { "transB": 0 },
+      "inputs": {
+        "A": { "shape": [32, 224], "dtype": "float32", "dist": "normal", "seed": 5408, "scale": 0.125 },
+        "B": { "shape": [224, 64], "dtype": "float32", "dist": "normal", "seed": 5508, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [32, 64], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "917504" }] }
+    },
+    {
+      "name": "direct_a_m32_k256_n64_tb1",
+      "attrs": { "transB": 1 },
+      "inputs": {
+        "A": { "shape": [32, 256], "dtype": "float32", "dist": "normal", "seed": 5409, "scale": 0.125 },
+        "B": { "shape": [64, 256], "dtype": "float32", "dist": "normal", "seed": 5509, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [32, 64], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "1048576" }] }
+    },
+    {
+      "name": "direct_a_m32_k288_n64_tb0",
+      "attrs": { "transB": 0 },
+      "inputs": {
+        "A": { "shape": [32, 288], "dtype": "float32", "dist": "normal", "seed": 5410, "scale": 0.125 },
+        "B": { "shape": [288, 64], "dtype": "float32", "dist": "normal", "seed": 5510, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [32, 64], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "1179648" }] }
+    },
+    {
+      "name": "direct_a_m32_k768_n128_tb0",
+      "attrs": { "transB": 0, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [32, 768], "dtype": "float32", "dist": "normal", "seed": 5411, "scale": 0.125 },
+        "B": { "shape": [768, 128], "dtype": "float32", "dist": "normal", "seed": 5511, "scale": 0.125 },
+        "C": { "shape": [128], "dtype": "float32", "dist": "normal", "seed": 5611, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [32, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "6291456" }] }
+    },
+    {
+      "name": "direct_a_m64_k1024_n128_tb0",
+      "attrs": { "transB": 0, "activation": "Tanh" },
+      "inputs": {
+        "A": { "shape": [64, 1024], "dtype": "float32", "dist": "normal", "seed": 5412, "scale": 0.125 },
+        "B": { "shape": [1024, 128], "dtype": "float32", "dist": "normal", "seed": 5512, "scale": 0.125 }
+      },
+      "outputs": { "Y": { "shape": [64, 128], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "16777216" }] }
     }
   ]
 }

build/webgpu/fused-gemm.wgsl.jinja CHANGED Viewed

@@ -24,26 +24,27 @@ fn gelu_tanh(v: f32) -> f32 {
 {% set actMode = gemmActivation | default("") %}
 {% set actA = gemmActAlpha | default(0.0) %}
 {% set actB = gemmActBeta | default(0.0) %}
-fn fused_act(v: f32) -> f32 {
-{% if actMode == "Relu" %}
   return max(v, 0.0);
-{% endif %}
-{% if actMode == "LeakyRelu" %}
-  return select(v * f32({{ actA }}), v, v >= 0.0);
-{% endif %}
-{% if actMode == "Sigmoid" %}
   return 1.0 / (1.0 + exp(-v));
-{% endif %}
-{% if actMode == "Tanh" %}
-  // Clamp first: the hardware tanh evaluates (e^2v - 1)/(e^2v + 1), which is
-  // Inf/Inf = NaN once e^2v overflows (v > ~44). A GEMM accumulator reaches that
-  // easily because it grows with K. tanh is already +/-1 to full f32 precision by
-  // |v| ~ 9, so the clamp changes no representable result.
   return tanh(clamp(v, -10.0, 10.0));
 {% endif %}
-{% if actMode == "HardSigmoid" %}
-  return clamp(f32({{ actA }}) * v + f32({{ actB }}), 0.0, 1.0);
-{% endif %}
 }
 {% endif %}
@@ -141,19 +142,22 @@ fn main(
           tileB[bc][br4] = bWord;
         }
         workgroupBarrier();
-{% set regIndent = "        " %}{% set regT = "f32" %}{% set regCast = 1 if tileT != "f32" else 0 %}{{ regIndent }}let aRow = lid.y * TM;
-{{ regIndent }}let bCol = lid.x * TN;
-{{ regIndent }}for (var kv: u32 = 0u; kv < BK / 4u; kv = kv + 1u) {
-{{ regIndent }}  var av: array<vec4<{{ regT }}>, TM>;
-{{ regIndent }}  var bv: array<vec4<{{ regT }}>, TN>;
-{{ regIndent }}  for (var i: u32 = 0u; i < TM; i = i + 1u) { av[i] = {% if regCast %}vec4<{{ regT }}>(tileA[aRow + i][kv]){% else %}tileA[aRow + i][kv]{% endif %}; }
-{{ regIndent }}  for (var j: u32 = 0u; j < TN; j = j + 1u) { bv[j] = {% if regCast %}vec4<{{ regT }}>(tileB[bCol + j][kv]){% else %}tileB[bCol + j][kv]{% endif %}; }
-{{ regIndent }}  for (var i: u32 = 0u; i < TM; i = i + 1u) {
-{{ regIndent }}    for (var j: u32 = 0u; j < TN; j = j + 1u) {
-{{ regIndent }}      acc[i * TN + j] = acc[i * TN + j] + dot(av[i], bv[j]);
-{{ regIndent }}    }
-{{ regIndent }}  }
-{{ regIndent }}}
         workgroupBarrier();
       }

 {% set actMode = gemmActivation | default("") %}
 {% set actA = gemmActAlpha | default(0.0) %}
 {% set actB = gemmActBeta | default(0.0) %}
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
   return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
   return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
   return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
+  return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
 {% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(actMode, actA, actB) -}}
 }
 {% endif %}
           tileB[bc][br4] = bWord;
         }
         workgroupBarrier();
+{% set regT = "f32" %}{% set regCast = 1 if tileT != "f32" else 0 %}
+{% filter indent(8, true) %}
+let aRow = lid.y * TM;
+let bCol = lid.x * TN;
+for (var kv: u32 = 0u; kv < BK / 4u; kv = kv + 1u) {
+  var av: array<vec4<{{ regT }}>, TM>;
+  var bv: array<vec4<{{ regT }}>, TN>;
+  for (var i: u32 = 0u; i < TM; i = i + 1u) { av[i] = {% if regCast %}vec4<{{ regT }}>(tileA[aRow + i][kv]){% else %}tileA[aRow + i][kv]{% endif %}; }
+  for (var j: u32 = 0u; j < TN; j = j + 1u) { bv[j] = {% if regCast %}vec4<{{ regT }}>(tileB[bCol + j][kv]){% else %}tileB[bCol + j][kv]{% endif %}; }
+  for (var i: u32 = 0u; i < TM; i = i + 1u) {
+    for (var j: u32 = 0u; j < TN; j = j + 1u) {
+      acc[i * TN + j] = acc[i * TN + j] + dot(av[i], bv[j]);
+    }
+  }
+}
+{% endfilter %}
         workgroupBarrier();
       }

build/webgpu/gemm-subgroup-matrix.wgsl.jinja CHANGED Viewed

@@ -1,6 +1,3 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 enable subgroups;
 {% if pinSubgroupSize32 %}
 enable subgroup_size_control;
@@ -8,10 +5,12 @@ enable subgroup_size_control;
 enable chromium_experimental_subgroup_matrix;
 diagnostic(off, chromium.subgroup_matrix_uniformity);
 {{ env.wgsl.resourceDeclarations }}
 {% set operandScalar = mmaScalar %}
 {% set accScalar = "f32" %}
 {% set OUT = outputBuffer if outputBuffer is defined else "y" %}
 {% set OUT_SCALAR = outScalar if outScalar is defined else yScalar %}
 {% set tileMValue = tileM if tileM is defined else 32 %}
@@ -20,19 +19,46 @@ diagnostic(off, chromium.subgroup_matrix_uniformity);
 {% set gemmEpi = gemmEpilogue if gemmEpilogue is defined else "none" %}
 {% set epiOpen = "gelu_tanh(" if gemmEpi == "fastgelu" else ("fused_act(" if gemmEpi == "activation" else "") %}
 {% set epiClose = ")" if gemmEpi != "none" else "" %}
 {% set useDoubleBufferedB = doubleBufferedB is defined and doubleBufferedB %}
 {% set aTransposed = aTransposedStorage if aTransposedStorage is defined else (matrixInputsTransposed is defined and matrixInputsTransposed) %}
 {% set bTransposed = bTransposedStorage if bTransposedStorage is defined else (matrixInputsTransposed is defined and matrixInputsTransposed) %}
 {% set columnTileBaseValue = columnTileBase if columnTileBase is defined else 0 %}
 {% set rowTileBaseValue = rowTileBase if rowTileBase is defined else 0 %}
 // Four subgroups cover 32xN as a 2x2 grid, or taller tiles as four row bands.
 {% set subtileCols = (tileNValue / 2)|int %}
 {% set subtileRows = 16 %}
 {% set aTilesPerSubgroup = (subtileRows / 8)|int %}
 {% set bTilesPerSubgroup = (subtileCols / 8)|int %}
 {% set subgroupCount = 4 %}
 {% set scratchBanks = subgroupCount * aTilesPerSubgroup * bTilesPerSubgroup %}
 {% if gemmEpi == "fastgelu" %}
 fn tanh_safe(x: f32) -> f32 {
   if (x > 10.0) { return 1.0; }
@@ -47,26 +73,27 @@ fn gelu_tanh(v: f32) -> f32 {
 {% set actMode = gemmActivation | default("") %}
 {% set actA = gemmActAlpha | default(0.0) %}
 {% set actB = gemmActBeta | default(0.0) %}
-fn fused_act(v: f32) -> f32 {
-{% if actMode == "Relu" %}
   return max(v, 0.0);
-{% endif %}
-{% if actMode == "LeakyRelu" %}
-  return select(v * f32({{ actA }}), v, v >= 0.0);
-{% endif %}
-{% if actMode == "Sigmoid" %}
   return 1.0 / (1.0 + exp(-v));
-{% endif %}
-{% if actMode == "Tanh" %}
-  // Clamp first: the hardware tanh evaluates (e^2v - 1)/(e^2v + 1), which is
-  // Inf/Inf = NaN once e^2v overflows (v > ~44). A GEMM accumulator reaches that
-  // easily because it grows with K. tanh is already +/-1 to full f32 precision by
-  // |v| ~ 9, so the clamp changes no representable result.
   return tanh(clamp(v, -10.0, 10.0));
 {% endif %}
-{% if actMode == "HardSigmoid" %}
-  return clamp(f32({{ actA }}) * v + f32({{ actB }}), 0.0, 1.0);
-{% endif %}
 }
 {% endif %}
@@ -80,13 +107,16 @@ const SUB_COLS:   u32 = {{ subtileCols }}u;
 const SUB_ROWS:   u32 = {{ subtileRows }}u;
 const COLUMN_TILE_BASE: u32 = {{ columnTileBaseValue }}u;
 const ROW_TILE_BASE: u32 = {{ rowTileBaseValue }}u;
 var<workgroup> tile_A: array<{{ operandScalar }}, {{ tileMValue }} * {{ tileKValue }}>;
 var<workgroup> tile_B: array<{{ operandScalar }}, {{ (2 if useDoubleBufferedB else 1) * tileNValue }} * {{ tileKValue }}>;
 // Distinct readback banks for every result matrix, so one barrier can publish
 // the full subtile without write-after-read reuse.
 var<workgroup> scratch: array<array<{{ accScalar }}, 64>, {{ scratchBanks }}>;
 fn loadSHMA(tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
 {% if aTransposed %}
   // Generic transA stores A as [K,M]. Keep both the global and workgroup writes
@@ -123,6 +153,7 @@ fn loadSHMA(tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
 {% endif %}
 }
 fn loadSHMB(tile_base: u32, k_idx: u32, row: u32, c_idx: u32, bank_offset: u32) {
 {% if bTransposed %}
   // Generic transB=0 stores B as [K,N]. Stage that native orientation and load
@@ -161,11 +192,13 @@ fn loadSHMB(tile_base: u32, k_idx: u32, row: u32, c_idx: u32, bank_offset: u32)
 {% set needsColBase = hasBias or (tailSafe is defined and tailSafe) %}
 {% set BIAS_1 = (" + " ~ accScalar ~ "(bias[col_base + col])") if hasBias else "" %}
 {% set BIAS_2 = (" + " ~ accScalar ~ "(bias[col_base + col2])") if hasBias else "" %}
 fn storeOutput(offset: u32{% if needsColBase %}, col_base: u32{% endif %}, row: u32, col: u32, src_slot: u32, row_limit: i32) {
   if (row_limit > 0 && row < u32(row_limit)) {
     let col2: u32 = col + 1u;
-    {{ OUT }}[offset + row * OUT_F + col]  = {{ OUT_SCALAR }}({{ epiOpen }}scratch[src_slot][row * 8u + col]{{ BIAS_1 }}{{ epiClose }});
-    {{ OUT }}[offset + row * OUT_F + col2] = {{ OUT_SCALAR }}({{ epiOpen }}scratch[src_slot][row * 8u + col2]{{ BIAS_2 }}{{ epiClose }});
   }
 }
@@ -185,34 +218,59 @@ fn main(
   let w_global_base: u32 = (COLUMN_TILE_BASE + workgroup_id.x) * TILE_COLS;
   let subtile_id: u32 = local_idx / sg_size;
   let subtile_idx: u32 = subtile_id / 2u;
   let subtile_idy: u32 = subtile_id % 2u;
   let base_A: u32 = subtile_idy * SUB_ROWS;
   let base_B: u32 = subtile_idx * SUB_COLS;
-{% for m in range(aTilesPerSubgroup) %}
-{% for n in range(bTilesPerSubgroup) %}
-  var matC{{ m }}{{ n }}: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
-{% endfor %}
-{% endfor %}
   for (var kidx: u32 = 0u; kidx < {{ kLoopEndValue }}; kidx += TILE_K) {
     loadSHMA(a_global_base, kidx, local_idx / {{ (tileKValue / 8)|int }}u, local_idx % {{ (tileKValue / 8)|int }}u);
     loadSHMB(w_global_base, kidx, local_idx / {{ (tileKValue / 16)|int }}u, local_idx % {{ (tileKValue / 16)|int }}u, 0u);
     workgroupBarrier();
     for (var step: u32 = 0u; step < TILE_K; step += 8u) {
 {% set dynamicATiles = aTilesPerSubgroup if aTilesPerSubgroup is defined else 2 %}
 {% set aTransposed = aTransposedStorage if aTransposedStorage is defined else (matrixInputsTransposed is defined and matrixInputsTransposed) %}
 {% set bTransposed = bTransposedStorage if bTransposedStorage is defined else (matrixInputsTransposed is defined and matrixInputsTransposed) %}
 {% set B_BANK = "b_bank_offset + " if doubleBufferedB is defined and doubleBufferedB else "" %}
-{% if aTransposed %}
-      let matrix_a_offset = step * TILE_ROWS + subtile_idy * SUB_ROWS;
 {% else %}
-      let matrix_a_offset = subtile_idy * SUB_ROWS * TILE_K + step;
 {% endif %}
 {% for m in range(dynamicATiles) %}
-{% if aTransposed %}
       var matA{{ m }}: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>, col_major>(&tile_A, matrix_a_offset + {{ m * 8 }}u, TILE_ROWS);
 {% else %}
       var matA{{ m }}: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>, row_major>(&tile_A, matrix_a_offset{% if m > 0 %} + {{ m * 8 }}u * TILE_K{% endif %}, TILE_K);
@@ -220,9 +278,9 @@ fn main(
 {% endfor %}
 {% if bTransposed %}
-      let matrix_b_offset = {{ B_BANK }}step * TILE_COLS + subtile_idx * SUB_COLS;
 {% else %}
-      let matrix_b_offset = {{ B_BANK }}subtile_idx * SUB_COLS * TILE_K + step;
 {% endif %}
 {% for n in range(bTilesPerSubgroup) %}
 {% if bTransposed %}
@@ -234,15 +292,43 @@ fn main(
 {% for m in range(dynamicATiles) %}
 {% for n in range(bTilesPerSubgroup) %}
-      matC{{ m }}{{ n }} = subgroupMatrixMultiplyAccumulate(matA{{ m }}, matB{{ n }}, matC{{ m }}{{ n }});
 {% endfor %}
 {% endfor %}
     }
     workgroupBarrier();
   }
   let matrix_c_offset: u32 = (a_global_base + base_A) * OUT_F + w_global_base + base_B;
   let row: u32 = sg_id / 4u;
   let col: u32 = (sg_id % 4u) * 2u;
   // Stage every result into a distinct bank before one barrier. Tail/bias paths
@@ -250,13 +336,10 @@ fn main(
   // across lanes, so its cross-lane readback must be published before partial-M
   // guards diverge. The epilogue also handles output conversion and bias.
   let bank: u32 = subtile_id * {{ aTilesPerSubgroup * bTilesPerSubgroup }}u;
-{% for m in range(aTilesPerSubgroup) %}
-{% for n in range(bTilesPerSubgroup) %}
-  subgroupMatrixStore<row_major>(&scratch[bank + {{ m * bTilesPerSubgroup + n }}u], 0u, matC{{ m }}{{ n }}, 8u);
-{% endfor %}
-{% endfor %}
   workgroupBarrier();
 {% for m in range(aTilesPerSubgroup) %}
   let row_limit_{{ m }}: i32 = i32(params.M) - i32(a_global_base + base_A + {{ m * 8 }}u);
 {% for n in range(bTilesPerSubgroup) %}
@@ -268,7 +351,7 @@ fn main(
     row,
     col,
     bank + {{ m * bTilesPerSubgroup + n }}u,
-    row_limit_{{ m }}
   );
 {% endfor %}
 {% endfor %}

 enable subgroups;
 {% if pinSubgroupSize32 %}
 enable subgroup_size_control;
 enable chromium_experimental_subgroup_matrix;
 diagnostic(off, chromium.subgroup_matrix_uniformity);
 {{ env.wgsl.resourceDeclarations }}
 {% set operandScalar = mmaScalar %}
 {% set accScalar = "f32" %}
+{% set SPLIT_K = splitK if splitK is defined else 1 %}
 {% set OUT = outputBuffer if outputBuffer is defined else "y" %}
 {% set OUT_SCALAR = outScalar if outScalar is defined else yScalar %}
 {% set tileMValue = tileM if tileM is defined else 32 %}
 {% set gemmEpi = gemmEpilogue if gemmEpilogue is defined else "none" %}
 {% set epiOpen = "gelu_tanh(" if gemmEpi == "fastgelu" else ("fused_act(" if gemmEpi == "activation" else "") %}
 {% set epiClose = ")" if gemmEpi != "none" else "" %}
+{% set useDirectStore = directStore is defined and directStore and gemmEpi == "none" %}
+{% set useDirectAStorage = directAStorage is defined and directAStorage %}
 {% set useDoubleBufferedB = doubleBufferedB is defined and doubleBufferedB %}
+{% set useBalancedAccumulation = balancedMatrixAccumulation is defined and balancedMatrixAccumulation and useDoubleBufferedB and not useDirectStore and operandScalar == "f32" %}
+{% set matrixChunkTiles = compensatedMatrixChunkTiles if compensatedMatrixChunkTiles is defined else 0 %}
+{% set useChunkedAccumulation = matrixChunkTiles > 0 and useDirectAStorage and not useDoubleBufferedB and not useDirectStore and operandScalar == "f32" and SPLIT_K == 1 %}
 {% set aTransposed = aTransposedStorage if aTransposedStorage is defined else (matrixInputsTransposed is defined and matrixInputsTransposed) %}
 {% set bTransposed = bTransposedStorage if bTransposedStorage is defined else (matrixInputsTransposed is defined and matrixInputsTransposed) %}
 {% set columnTileBaseValue = columnTileBase if columnTileBase is defined else 0 %}
 {% set rowTileBaseValue = rowTileBase if rowTileBase is defined else 0 %}
 // Four subgroups cover 32xN as a 2x2 grid, or taller tiles as four row bands.
+{% if tileMValue > 32 %}
+{% set subtileCols = tileNValue %}
+{% set subtileRows = (tileMValue / 4)|int %}
+{% else %}
 {% set subtileCols = (tileNValue / 2)|int %}
 {% set subtileRows = 16 %}
+{% endif %}
 {% set aTilesPerSubgroup = (subtileRows / 8)|int %}
 {% set bTilesPerSubgroup = (subtileCols / 8)|int %}
 {% set subgroupCount = 4 %}
 {% set scratchBanks = subgroupCount * aTilesPerSubgroup * bTilesPerSubgroup %}
+{% macro matrixAccumulators() %}
+{% for m in range(aTilesPerSubgroup) %}
+{% for n in range(bTilesPerSubgroup) %}
+  var matC{{ m }}{{ n }}: subgroup_matrix_result<{{ accScalar }}, 8, 8>;
+{% endfor %}
+{% endfor %}
+{% endmacro %}
+{% macro storeMatrixAccumulators() %}
+{% for m in range(aTilesPerSubgroup) %}
+{% for n in range(bTilesPerSubgroup) %}
+  subgroupMatrixStore<row_major>(&scratch[bank + {{ m * bTilesPerSubgroup + n }}u], 0u, matC{{ m }}{{ n }}, 8u);
+{% endfor %}
+{% endfor %}
+{% endmacro %}
 {% if gemmEpi == "fastgelu" %}
 fn tanh_safe(x: f32) -> f32 {
   if (x > 10.0) { return 1.0; }
 {% set actMode = gemmActivation | default("") %}
 {% set actA = gemmActAlpha | default(0.0) %}
 {% set actB = gemmActBeta | default(0.0) %}
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
   return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
   return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
   return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
+  return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
 {% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(actMode, actA, actB) -}}
 }
 {% endif %}
 const SUB_ROWS:   u32 = {{ subtileRows }}u;
 const COLUMN_TILE_BASE: u32 = {{ columnTileBaseValue }}u;
 const ROW_TILE_BASE: u32 = {{ rowTileBaseValue }}u;
+{% if not useDirectAStorage %}
 var<workgroup> tile_A: array<{{ operandScalar }}, {{ tileMValue }} * {{ tileKValue }}>;
+{% endif %}
 var<workgroup> tile_B: array<{{ operandScalar }}, {{ (2 if useDoubleBufferedB else 1) * tileNValue }} * {{ tileKValue }}>;
 // Distinct readback banks for every result matrix, so one barrier can publish
 // the full subtile without write-after-read reuse.
 var<workgroup> scratch: array<array<{{ accScalar }}, 64>, {{ scratchBanks }}>;
+{% if not useDirectAStorage %}
 fn loadSHMA(tile_base: u32, k_idx: u32, row: u32, c_idx: u32) {
 {% if aTransposed %}
   // Generic transA stores A as [K,M]. Keep both the global and workgroup writes
 {% endif %}
 }
+{% endif %}
 fn loadSHMB(tile_base: u32, k_idx: u32, row: u32, c_idx: u32, bank_offset: u32) {
 {% if bTransposed %}
   // Generic transB=0 stores B as [K,N]. Stage that native orientation and load
 {% set needsColBase = hasBias or (tailSafe is defined and tailSafe) %}
 {% set BIAS_1 = (" + " ~ accScalar ~ "(bias[col_base + col])") if hasBias else "" %}
 {% set BIAS_2 = (" + " ~ accScalar ~ "(bias[col_base + col2])") if hasBias else "" %}
+{% set ACC_1 = "(prior.x + (scratch[src_slot][row * 8u + col] + tile_B[src_slot * 64u + row * 8u + col]))" if useBalancedAccumulation else "scratch[src_slot][row * 8u + col]" %}
+{% set ACC_2 = "(prior.y + (scratch[src_slot][row * 8u + col2] + tile_B[src_slot * 64u + row * 8u + col2]))" if useBalancedAccumulation else "scratch[src_slot][row * 8u + col2]" %}
 fn storeOutput(offset: u32{% if needsColBase %}, col_base: u32{% endif %}, row: u32, col: u32, src_slot: u32, row_limit: i32) {
   if (row_limit > 0 && row < u32(row_limit)) {
     let col2: u32 = col + 1u;
+    {{ OUT }}[offset + row * OUT_F + col]  = {{ OUT_SCALAR }}({{ epiOpen }}{{ ACC_1 }}{{ BIAS_1 }}{{ epiClose }});
+    {{ OUT }}[offset + row * OUT_F + col2] = {{ OUT_SCALAR }}({{ epiOpen }}{{ ACC_2 }}{{ BIAS_2 }}{{ epiClose }});
   }
 }
   let w_global_base: u32 = (COLUMN_TILE_BASE + workgroup_id.x) * TILE_COLS;
   let subtile_id: u32 = local_idx / sg_size;
+{% if tileMValue > 32 %}
+  let subtile_idx: u32 = 0u;
+  let subtile_idy: u32 = subtile_id;
+{% else %}
   let subtile_idx: u32 = subtile_id / 2u;
   let subtile_idy: u32 = subtile_id % 2u;
+{% endif %}
   let base_A: u32 = subtile_idy * SUB_ROWS;
   let base_B: u32 = subtile_idx * SUB_COLS;
+{% if not useChunkedAccumulation %}
+{{ matrixAccumulators() }}
+{% endif %}
+{% if useChunkedAccumulation %}
+  let row = sg_id / 4u;
+  let col = (sg_id % 4u) * 2u;
+  let bank = subtile_id * {{ aTilesPerSubgroup * bTilesPerSubgroup }}u;
+  var chunkSums: array<vec2<f32>, {{ aTilesPerSubgroup * bTilesPerSubgroup }}>;
+  var chunkErrors: array<vec2<f32>, {{ aTilesPerSubgroup * bTilesPerSubgroup }}>;
+  for (var chunk_base = 0u; chunk_base < IN_F; chunk_base += TILE_K * {{ matrixChunkTiles }}u) {
+{{ matrixAccumulators() }}
+    for (var kidx = chunk_base; kidx < min(chunk_base + TILE_K * {{ matrixChunkTiles }}u, IN_F); kidx += TILE_K) {
+{% else %}
   for (var kidx: u32 = 0u; kidx < {{ kLoopEndValue }}; kidx += TILE_K) {
+{% endif %}
+{% if not useDirectAStorage %}
     loadSHMA(a_global_base, kidx, local_idx / {{ (tileKValue / 8)|int }}u, local_idx % {{ (tileKValue / 8)|int }}u);
+{% endif %}
     loadSHMB(w_global_base, kidx, local_idx / {{ (tileKValue / 16)|int }}u, local_idx % {{ (tileKValue / 16)|int }}u, 0u);
     workgroupBarrier();
     for (var step: u32 = 0u; step < TILE_K; step += 8u) {
 {% set dynamicATiles = aTilesPerSubgroup if aTilesPerSubgroup is defined else 2 %}
+{% set matrixAccumulator = matrixAccumulator | default("matC") %}
+{% set matrixStepOffset = matrixStepOffset | default(0) %}
+{% set matrixStep = "step" if matrixStepOffset == 0 else "(step + " ~ matrixStepOffset ~ "u)" %}
 {% set aTransposed = aTransposedStorage if aTransposedStorage is defined else (matrixInputsTransposed is defined and matrixInputsTransposed) %}
 {% set bTransposed = bTransposedStorage if bTransposedStorage is defined else (matrixInputsTransposed is defined and matrixInputsTransposed) %}
+{% set directA = directAStorage is defined and directAStorage %}
 {% set B_BANK = "b_bank_offset + " if doubleBufferedB is defined and doubleBufferedB else "" %}
+{% if directA %}
+      // Direct storage guarantees a full aligned M/K tile, so subgroup-matrix
+      // loads need no guarded workgroup staging.
+      let matrix_a_offset = (a_global_base + base_A) * IN_F + kidx + {{ matrixStep }};
+{% elif aTransposed %}
+      let matrix_a_offset = {{ matrixStep }} * TILE_ROWS + subtile_idy * SUB_ROWS;
 {% else %}
+      let matrix_a_offset = subtile_idy * SUB_ROWS * TILE_K + {{ matrixStep }};
 {% endif %}
 {% for m in range(dynamicATiles) %}
+{% if directA %}
+      var matA{{ m }}: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>, row_major>(&x, matrix_a_offset{% if m > 0 %} + {{ m * 8 }}u * IN_F{% endif %}, IN_F);
+{% elif aTransposed %}
       var matA{{ m }}: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>, col_major>(&tile_A, matrix_a_offset + {{ m * 8 }}u, TILE_ROWS);
 {% else %}
       var matA{{ m }}: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>, row_major>(&tile_A, matrix_a_offset{% if m > 0 %} + {{ m * 8 }}u * TILE_K{% endif %}, TILE_K);
 {% endfor %}
 {% if bTransposed %}
+      let matrix_b_offset = {{ B_BANK }}{{ matrixStep }} * TILE_COLS + subtile_idx * SUB_COLS;
 {% else %}
+      let matrix_b_offset = {{ B_BANK }}subtile_idx * SUB_COLS * TILE_K + {{ matrixStep }};
 {% endif %}
 {% for n in range(bTilesPerSubgroup) %}
 {% if bTransposed %}
 {% for m in range(dynamicATiles) %}
 {% for n in range(bTilesPerSubgroup) %}
+      {{ matrixAccumulator }}{{ m }}{{ n }} = subgroupMatrixMultiplyAccumulate(matA{{ m }}, matB{{ n }}, {{ matrixAccumulator }}{{ m }}{{ n }});
 {% endfor %}
 {% endfor %}
     }
     workgroupBarrier();
   }
+{% if useChunkedAccumulation %}
+{{ storeMatrixAccumulators() }}
+    workgroupBarrier();
+{% for m in range(aTilesPerSubgroup) %}
+{% for n in range(bTilesPerSubgroup) %}
+    {
+      let partial = vec2<f32>(scratch[bank + {{ m * bTilesPerSubgroup + n }}u][row * 8u + col], scratch[bank + {{ m * bTilesPerSubgroup + n }}u][row * 8u + col + 1u]);
+      let corrected = partial - chunkErrors[{{ m * bTilesPerSubgroup + n }}];
+      let next = chunkSums[{{ m * bTilesPerSubgroup + n }}] + corrected;
+      // Infinite totals keep a zero correction: Inf - Inf must not poison
+      // a later finite partial. NaN totals still propagate through the sum.
+      chunkErrors[{{ m * bTilesPerSubgroup + n }}] = select(vec2<f32>(0.0), (next - chunkSums[{{ m * bTilesPerSubgroup + n }}]) - corrected, abs(next) <= vec2<f32>(3.4028234663852886e38));
+      chunkSums[{{ m * bTilesPerSubgroup + n }}] = next;
+    }
+{% endfor %}
+{% endfor %}
+    workgroupBarrier();
+  }
+{% endif %}
   let matrix_c_offset: u32 = (a_global_base + base_A) * OUT_F + w_global_base + base_B;
+{% if useChunkedAccumulation %}
+{% for m in range(aTilesPerSubgroup) %}
+{% for n in range(bTilesPerSubgroup) %}
+  scratch[bank + {{ m * bTilesPerSubgroup + n }}u][row * 8u + col] = chunkSums[{{ m * bTilesPerSubgroup + n }}].x;
+  scratch[bank + {{ m * bTilesPerSubgroup + n }}u][row * 8u + col + 1u] = chunkSums[{{ m * bTilesPerSubgroup + n }}].y;
+{% endfor %}
+{% endfor %}
+  workgroupBarrier();
+{% else %}
   let row: u32 = sg_id / 4u;
   let col: u32 = (sg_id % 4u) * 2u;
   // Stage every result into a distinct bank before one barrier. Tail/bias paths
   // across lanes, so its cross-lane readback must be published before partial-M
   // guards diverge. The epilogue also handles output conversion and bias.
   let bank: u32 = subtile_id * {{ aTilesPerSubgroup * bTilesPerSubgroup }}u;
+{{ storeMatrixAccumulators() }}
   workgroupBarrier();
+{% endif %}
 {% for m in range(aTilesPerSubgroup) %}
   let row_limit_{{ m }}: i32 = i32(params.M) - i32(a_global_base + base_A + {{ m * 8 }}u);
 {% for n in range(bTilesPerSubgroup) %}
     row,
     col,
     bank + {{ m * bTilesPerSubgroup + n }}u,
+    row_limit_{{ m }}{{ ", pairTotals[" ~ (m * bTilesPerSubgroup + n) ~ "]" if useBalancedAccumulation else "" }}
   );
 {% endfor %}
 {% endfor %}

build/webgpu/gemv-contiguous.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,158 @@

+{% set gemvFused = gemvFused | default(false) %}
+{% set COLS = gemvColumns | default(4) %}
+{% set ACC = "f32" if COLS == 1 else "vec4<f32>" %}
+{% macro wgsl_tree_fold_stmt(a, op, idx, svar) %}
+{% if op == "max" %}
+{{ a }}[{{ idx }}] = max({{ a }}[{{ idx }}], {{ a }}[{{ idx }} + {{ svar }}]);
+{%- else %}
+{{ a }}[{{ idx }}] = {{ a }}[{{ idx }}] + {{ a }}[{{ idx }} + {{ svar }}];
+{%- endif %}
+{% endmacro %}
+{% macro wgsl_tree_fold(arrays, op="add", idx="lid", wg="WORKGROUP_SIZE", svar="stride", typed=false, form="tail", breakInline=false, bodyInline=false, barrierFirst=false) %}
+  var {{ svar }}{{ ": u32 " if typed else " " }}= {{ wg }} / 2u;
+  loop {
+{% if form == "head" %}
+{% if breakInline %}
+    if ({{ svar }} == 0u) { break; }
+{% else %}
+    if ({{ svar }} == 0u) {
+      break;
+    }
+{% endif %}
+{% endif %}
+{% if bodyInline %}
+    if ({{ idx }} < {{ svar }}) { {{ wgsl_tree_fold_stmt(arrays[0], op, idx, svar) }} }
+{% else %}
+    if ({{ idx }} < {{ svar }}) {
+{% for a in arrays %}
+      {{ wgsl_tree_fold_stmt(a, op, idx, svar) }}
+{% endfor %}
+    }
+{% endif %}
+{% if form == "head" %}
+{% if barrierFirst %}
+    workgroupBarrier();
+    {{ svar }} = {{ svar }} / 2u;
+{% else %}
+    {{ svar }} = {{ svar }} / 2u;
+    workgroupBarrier();
+{% endif %}
+{% else %}
+    workgroupBarrier();
+    if ({{ svar }} == 1u) {
+      break;
+    }
+    {{ svar }} = {{ svar }} / 2u;
+{% endif %}
+  }
+{%- endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+{% if gemvFused and (hasActivation | default(false)) %}
+{% set actMode = gemmActivation | default("") %}
+{% set actA = gemmActAlpha | default(0.0) %}
+{% set actB = gemmActBeta | default(0.0) %}
+{% macro fused_act_return(mode, alpha, beta) -%}
+{% if mode == "Relu" %}
+  return max(v, 0.0);
+{% elif mode == "Clip" %}
+  return clamp(v, f32({{ alpha }}), f32({{ beta }}));
+{% elif mode == "LeakyRelu" %}
+  return select(v * f32({{ alpha }}), v, v >= 0.0);
+{% elif mode == "Sigmoid" %}
+  return 1.0 / (1.0 + exp(-v));
+{% elif mode == "Tanh" %}
+  // tanh is already +/-1 to f32 precision at the clamp bounds. Clamping before
+  // the builtin preserves that saturated result for larger accumulators.
+  return tanh(clamp(v, -10.0, 10.0));
+{% elif mode == "HardSigmoid" %}
+  return clamp(f32({{ alpha }}) * v + f32({{ beta }}), 0.0, 1.0);
+{% else %}
+  return v * clamp(v * 0.16666666666666666 + 0.5, 0.0, 1.0);
+{% endif %}
+{%- endmacro -%}
+fn fused_act(v: f32) -> f32 {
+{{ fused_act_return(actMode, actA, actB) -}}
+}
+{% endif %}
+const IN_F: u32 = {{ inFeatures }}u;
+const OUT_F: u32 = {{ outFeatures }}u;
+const WG: u32 = {{ workgroupSize }}u;
+var<workgroup> partials: array<{{ ACC }}, {{ workgroupSize }}>;
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>
+) {
+  let tid = lid.x;
+  // wg.z carries column-group bits past the per-dimension dispatch limit.
+  // This reduces to wg.x when the dispatch does not fold. col_base is workgroup-uniform, so
+  // this guard returns the whole over-dispatched workgroup before any barrier.
+  let col_base = (wg.x + wg.z * {{ DISPATCH_FOLD_WIDTH }}u) * {{ COLS }}u;
+  if (col_base >= OUT_F) {
+    return;
+  }
+  var acc = {{ ACC }}(0.0);
+  for (var i = tid; i < IN_F; i = i + WG) {
+{% if xScalar == "f16" %}
+    let xv = f32(x[i]);
+{% else %}
+    let xv = x[i];
+{% endif %}
+{% for j in range(COLS) %}
+    let w{{ j }} = w[(col_base + {{ j }}u) * IN_F + i];
+{% endfor %}
+{% if wScalar == "f16" %}
+    acc = acc + xv * {{ "f32(w0)" if COLS == 1 else "vec4<f32>(f32(w0), f32(w1), f32(w2), f32(w3))" }};
+{% else %}
+    acc = acc + xv * {{ "w0" if COLS == 1 else "vec4<f32>(w0, w1, w2, w3)" }};
+{% endif %}
+  }
+  partials[tid] = acc;
+  workgroupBarrier();
+{{ wgsl_tree_fold(["partials"], idx="tid", wg="WG", form="head") }}
+  if (tid == 0u) {
+{% if gemvFused %}
+{% for j in range(COLS) %}
+    {
+      let n = col_base + {{ j }}u;
+      let product = select(partials[0]{% if COLS > 1 %}[{{ j }}u]{% endif %} * params.alpha, 0.0, params.alpha == 0.0);
+{% if hasBias %}
+      let bias_term = select(f32(bias[{% if scalarBiasOk | default(false) %}0u{% else %}n{% endif %}]) * params.beta, 0.0, params.beta == 0.0);
+{% endif %}
+      y[n] = {{ yScalar }}({% if hasActivation | default(false) %}fused_act({% endif %}product{% if hasBias %} + bias_term{% endif %}{% if hasActivation | default(false) %}){% endif %});
+    }
+{% endfor %}
+{% else %}
+{% if hasBias is defined and hasBias %}
+    // C is a row vector over the output columns, so these four columns read four
+    // consecutive entries. This specialization requires beta = 1, so bias is
+    // added without a per-output multiply.
+    let total = partials[0] + vec4<f32>(
+      f32(bias[col_base + 0u]), f32(bias[col_base + 1u]),
+      f32(bias[col_base + 2u]), f32(bias[col_base + 3u]));
+{% else %}
+    let total = partials[0];
+{% endif %}
+{% if yScalar == "f16" %}
+    y[col_base + 0u] = f16(total.x);
+    y[col_base + 1u] = f16(total.y);
+    y[col_base + 2u] = f16(total.z);
+    y[col_base + 3u] = f16(total.w);
+{% else %}
+    y[col_base + 0u] = total.x;
+    y[col_base + 1u] = total.y;
+    y[col_base + 2u] = total.z;
+    y[col_base + 3u] = total.w;
+{% endif %}
+{% endif %}
+  }
+}

build/webgpu/manifest.json CHANGED Viewed

@@ -2,54 +2,23 @@
   "domain": "com.microsoft",
   "name": "FusedGemm",
   "sinceVersion": 1,
-  "description": "Gemm with a fused activation: `Y = act(alpha * A' * B' + beta * C)`, where `A'` and `B'` are optionally transposed and `C` is broadcastable to `(M, N)`. The activation runs in the f32 accumulator before the single output narrowing. This package supports `Relu`, `LeakyRelu`, `Sigmoid`, `Tanh` and `HardSigmoid`; the other activation strings and numeric types admitted by the open schema are not implemented. Omitting `activation` gives plain Gemm.",
-  "inputs": [
-    {
-      "role": "A",
-      "dtype": "T",
-      "rank": 2,
-      "description": "Left operand, `(M, K)` when `transA` is 0 and `(K, M)` otherwise."
-    },
-    {
-      "role": "B",
-      "dtype": "T",
-      "rank": 2,
-      "description": "Right operand, `(K, N)` when `transB` is 0 and `(N, K)` otherwise."
-    },
-    {
-      "role": "C",
-      "dtype": "T",
-      "optional": true,
-      "description": "Optional additive term, unidirectionally broadcastable to `(M, N)`: a scalar, a row `(N)`, a column `(M, 1)`, or the full matrix."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "Y",
-      "dtype": "T",
-      "rank": 2,
-      "shape": "[gemmM, gemmN]",
-      "description": "`act(alpha * A' * B' + beta * C)` with shape `(M, N)`."
-    }
-  ],
-  "attributes": { "alpha": 1, "beta": 1, "transA": 0, "transB": 0 },
-  "attributeDescriptions": {
-    "alpha": "Scalar multiplier for the product `A' * B'`; the standard default is 1.",
-    "beta": "Scalar multiplier for `C`; the standard default is 1.",
-    "transA": "Whether `A` is stored transposed. The standard default is 0.",
-    "transB": "Whether `B` is stored transposed. The standard default is 0.",
-    "activation": "Optional fused activation name. Supported modes are `Relu`, `LeakyRelu`, `Sigmoid`, `Tanh` and `HardSigmoid`; omission applies none.",
-    "activation_alpha": "First activation parameter: the slope for `LeakyRelu` or `alpha` for `HardSigmoid`.",
-    "activation_beta": "Second activation parameter: `beta` for `HardSigmoid`."
   },
   "attributeConstraints": { "activation": { "values": ["Relu", "LeakyRelu", "Sigmoid", "Tanh", "HardSigmoid"] } },
   "typeConstraints": { "T": ["float32", "float16"] },
-  "args": {
-    "A": { "kind": "tensor", "semantic": "A", "role": "input" },
-    "B": { "kind": "tensor", "semantic": "B", "role": "weights" },
-    "C": { "kind": "tensor", "semantic": "C", "role": "weights", "required": false },
-    "Y": { "kind": "tensor", "semantic": "Y", "role": "output" }
-  },
   "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",
@@ -78,89 +47,114 @@
     "noBiasContract": "baseContract and not present.C",
     "biasContract": "baseContract and present.C and cOk",
     "sgmatOperandBytes": "2 if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else 4",
-    "sgmatStorageBytes": "(32 * 32 + 64 * 32) * sgmatOperandBytes + 32 * 64 * 4",
     "sgmatResourcesFit": "128 <= deviceWorkgroupCap and sgmatStorageBytes <= device.limits.maxComputeWorkgroupStorageSize",
     "sgmatScalingOk": "(attrs.alpha if has(attrs, \"alpha\") else 1) == 1 and (attrs.beta if has(attrs, \"beta\") else 1) == 1",
-    "sgmatLayoutOk": "gemmM > 0 and gemmK % 32 == 0 and gemmN % 64 == 0 and ((tensorDtypes.A == \"float16\" and device.features.has(\"shader-f16\") and gemmM >= 2) or (tensorDtypes.A == \"float32\" and gemmM >= 32)) and ceilDiv(gemmM, 32) <= device.limits.maxComputeWorkgroupsPerDimension and ceilDiv(gemmN, 64) <= device.limits.maxComputeWorkgroupsPerDimension",
     "sgmatContract": "wave32Effective and sgmatLayoutOk and sgmatScalingOk and sgmatResourcesFit",
-    "sgmatRowBiasContract": "sgmatContract and rowBiasOk and ranks.C == 1"
-  },
-  "constants": {
     "usesF16": "tensorDtypes.A == \"float16\"",
     "gemmActivation": "activationName",
     "gemmActAlpha": "activationAlpha",
     "gemmActBeta": "activationBeta",
-    "gemmEpilogue": "\"activation\" if hasActivation else \"none\""
   },
-  "bindingSets": {
-    "noBias": [
-      { "name": "a", "arg": "A", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "b", "arg": "B", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "M", "type": "u32", "value": "gemmM" },
-            { "name": "N", "type": "u32", "value": "gemmN" },
-            { "name": "K", "type": "u32", "value": "gemmK" },
-            { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" }
-          ]
-        }
-      }
-    ],
-    "withBias": [
-      { "name": "a", "arg": "A", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "b", "arg": "B", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "bias", "arg": "C", "semantic": "C", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": {
-          "name": "Params",
-          "fields": [
-            { "name": "M", "type": "u32", "value": "gemmM" },
-            { "name": "N", "type": "u32", "value": "gemmN" },
-            { "name": "K", "type": "u32", "value": "gemmK" },
-            { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" },
-            { "name": "beta", "type": "f32", "value": "attrs.beta if has(attrs, \"beta\") else 1" }
-          ]
-        }
-      }
-    ],
-    "sgmatNoBias": [
-      { "name": "x", "arg": "A", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "w", "arg": "B", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "M", "type": "u32", "value": "gemmM" }] }
-      }
-    ],
-    "sgmatBias": [
-      { "name": "x", "arg": "A", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "w", "arg": "B", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "bias", "arg": "C", "semantic": "C", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "M", "type": "u32", "value": "gemmM" }] }
-      }
-    ]
   },
   "variants": [
     {
       "id": "notrans_sgmat_bias",
       "priority": 100,
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
@@ -168,32 +162,38 @@
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
-      "when": ["sgmatRowBiasContract", "transAFlag == false", "transBFlag == false"],
-      "constants": {
         "hasBias": true,
         "aTransposedStorage": false,
         "bTransposedStorage": "not false",
-        "xScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "wScalar": "\"f16\" if tensorDtypes.B == \"float16\" else \"f32\"",
-        "bScalar": "\"f16\" if tensorDtypes.C == \"float16\" else \"f32\"",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
-        "outFeatures": "gemmN"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrixBias",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
-          "bindings": "sgmatBias",
-          "dispatch": { "workgroups": "ceilDiv(gemmM, 32)", "x": "ceilDiv(gemmN, 64)" }
         }
       ]
     },
     {
       "id": "notrans_sgmat",
       "priority": 100,
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
@@ -201,31 +201,38 @@
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
-      "when": ["noBiasContract and sgmatContract", "transAFlag == false", "transBFlag == false"],
-      "constants": {
         "hasBias": false,
         "aTransposedStorage": false,
         "bTransposedStorage": "not false",
-        "xScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "wScalar": "\"f16\" if tensorDtypes.B == \"float16\" else \"f32\"",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
-        "outFeatures": "gemmN"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrix",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
-          "bindings": "sgmatNoBias",
-          "dispatch": { "workgroups": "ceilDiv(gemmM, 32)", "x": "ceilDiv(gemmN, 64)" }
         }
       ]
     },
     {
       "id": "transA_sgmat_bias",
       "priority": 100,
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
@@ -233,32 +240,38 @@
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
-      "when": ["sgmatRowBiasContract", "transAFlag == true", "transBFlag == false"],
-      "constants": {
         "hasBias": true,
         "aTransposedStorage": true,
         "bTransposedStorage": "not false",
-        "xScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "wScalar": "\"f16\" if tensorDtypes.B == \"float16\" else \"f32\"",
-        "bScalar": "\"f16\" if tensorDtypes.C == \"float16\" else \"f32\"",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
-        "outFeatures": "gemmN"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrixBias",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
-          "bindings": "sgmatBias",
-          "dispatch": { "workgroups": "ceilDiv(gemmM, 32)", "x": "ceilDiv(gemmN, 64)" }
         }
       ]
     },
     {
       "id": "transA_sgmat",
       "priority": 100,
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
@@ -266,31 +279,38 @@
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
-      "when": ["noBiasContract and sgmatContract", "transAFlag == true", "transBFlag == false"],
-      "constants": {
         "hasBias": false,
         "aTransposedStorage": true,
         "bTransposedStorage": "not false",
-        "xScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "wScalar": "\"f16\" if tensorDtypes.B == \"float16\" else \"f32\"",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
-        "outFeatures": "gemmN"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrix",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
-          "bindings": "sgmatNoBias",
-          "dispatch": { "workgroups": "ceilDiv(gemmM, 32)", "x": "ceilDiv(gemmN, 64)" }
         }
       ]
     },
     {
       "id": "transB_sgmat_bias",
       "priority": 100,
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
@@ -298,32 +318,38 @@
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
-      "when": ["sgmatRowBiasContract", "transAFlag == false", "transBFlag == true"],
-      "constants": {
         "hasBias": true,
         "aTransposedStorage": false,
         "bTransposedStorage": "not true",
-        "xScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "wScalar": "\"f16\" if tensorDtypes.B == \"float16\" else \"f32\"",
-        "bScalar": "\"f16\" if tensorDtypes.C == \"float16\" else \"f32\"",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
-        "outFeatures": "gemmN"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrixBias",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
-          "bindings": "sgmatBias",
-          "dispatch": { "workgroups": "ceilDiv(gemmM, 32)", "x": "ceilDiv(gemmN, 64)" }
         }
       ]
     },
     {
       "id": "transB_sgmat",
       "priority": 100,
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
@@ -331,31 +357,38 @@
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
-      "when": ["noBiasContract and sgmatContract", "transAFlag == false", "transBFlag == true"],
-      "constants": {
         "hasBias": false,
         "aTransposedStorage": false,
         "bTransposedStorage": "not true",
-        "xScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "wScalar": "\"f16\" if tensorDtypes.B == \"float16\" else \"f32\"",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
-        "outFeatures": "gemmN"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrix",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
-          "bindings": "sgmatNoBias",
-          "dispatch": { "workgroups": "ceilDiv(gemmM, 32)", "x": "ceilDiv(gemmN, 64)" }
         }
       ]
     },
     {
       "id": "transAB_sgmat_bias",
       "priority": 100,
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
@@ -363,32 +396,38 @@
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
-      "when": ["sgmatRowBiasContract", "transAFlag == true", "transBFlag == true"],
-      "constants": {
         "hasBias": true,
         "aTransposedStorage": true,
         "bTransposedStorage": "not true",
-        "xScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "wScalar": "\"f16\" if tensorDtypes.B == \"float16\" else \"f32\"",
-        "bScalar": "\"f16\" if tensorDtypes.C == \"float16\" else \"f32\"",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
-        "outFeatures": "gemmN"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrixBias",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
-          "bindings": "sgmatBias",
-          "dispatch": { "workgroups": "ceilDiv(gemmM, 32)", "x": "ceilDiv(gemmN, 64)" }
         }
       ]
     },
     {
       "id": "transAB_sgmat",
       "priority": 100,
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
@@ -396,25 +435,31 @@
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
-      "when": ["noBiasContract and sgmatContract", "transAFlag == true", "transBFlag == true"],
-      "constants": {
         "hasBias": false,
         "aTransposedStorage": true,
         "bTransposedStorage": "not true",
-        "xScalar": "\"f16\" if tensorDtypes.A == \"float16\" else \"f32\"",
-        "wScalar": "\"f16\" if tensorDtypes.B == \"float16\" else \"f32\"",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
-        "outFeatures": "gemmN"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrix",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
-          "bindings": "sgmatNoBias",
-          "dispatch": { "workgroups": "ceilDiv(gemmM, 32)", "x": "ceilDiv(gemmN, 64)" }
         }
       ]
     },
@@ -422,7 +467,7 @@
       "id": "notrans_bias",
       "priority": 0,
       "when": ["biasContract", "transAFlag == false", "transBFlag == false"],
-      "constants": {
         "transA": false,
         "transB": false,
         "hasBias": true,
@@ -435,10 +480,25 @@
           "id": "main",
           "name": "FusedGemm.Bias",
           "shader": "fused-gemm.wgsl.jinja",
-          "bindings": "withBias",
           "dispatch": {
-            "x": "min(ceilDiv(gemmN, 64), device.limits.maxComputeWorkgroupsPerDimension)",
-            "y": "min(ceilDiv(gemmM, 64), device.limits.maxComputeWorkgroupsPerDimension)"
           }
         }
       ]
@@ -447,16 +507,29 @@
       "id": "notrans",
       "priority": 0,
       "when": ["noBiasContract", "transAFlag == false", "transBFlag == false"],
-      "constants": { "transA": false, "transB": false, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm",
           "shader": "fused-gemm.wgsl.jinja",
-          "bindings": "noBias",
           "dispatch": {
-            "x": "min(ceilDiv(gemmN, 64), device.limits.maxComputeWorkgroupsPerDimension)",
-            "y": "min(ceilDiv(gemmM, 64), device.limits.maxComputeWorkgroupsPerDimension)"
           }
         }
       ]
@@ -465,7 +538,7 @@
       "id": "transA_bias",
       "priority": 0,
       "when": ["biasContract", "transAFlag == true", "transBFlag == false"],
-      "constants": {
         "transA": true,
         "transB": false,
         "hasBias": true,
@@ -478,10 +551,25 @@
           "id": "main",
           "name": "FusedGemm.Bias",
           "shader": "fused-gemm.wgsl.jinja",
-          "bindings": "withBias",
           "dispatch": {
-            "x": "min(ceilDiv(gemmN, 64), device.limits.maxComputeWorkgroupsPerDimension)",
-            "y": "min(ceilDiv(gemmM, 64), device.limits.maxComputeWorkgroupsPerDimension)"
           }
         }
       ]
@@ -490,16 +578,29 @@
       "id": "transA",
       "priority": 0,
       "when": ["noBiasContract", "transAFlag == true", "transBFlag == false"],
-      "constants": { "transA": true, "transB": false, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm",
           "shader": "fused-gemm.wgsl.jinja",
-          "bindings": "noBias",
           "dispatch": {
-            "x": "min(ceilDiv(gemmN, 64), device.limits.maxComputeWorkgroupsPerDimension)",
-            "y": "min(ceilDiv(gemmM, 64), device.limits.maxComputeWorkgroupsPerDimension)"
           }
         }
       ]
@@ -508,7 +609,7 @@
       "id": "transB_bias",
       "priority": 0,
       "when": ["biasContract", "transAFlag == false", "transBFlag == true"],
-      "constants": {
         "transA": false,
         "transB": true,
         "hasBias": true,
@@ -521,10 +622,25 @@
           "id": "main",
           "name": "FusedGemm.Bias",
           "shader": "fused-gemm.wgsl.jinja",
-          "bindings": "withBias",
           "dispatch": {
-            "x": "min(ceilDiv(gemmN, 64), device.limits.maxComputeWorkgroupsPerDimension)",
-            "y": "min(ceilDiv(gemmM, 64), device.limits.maxComputeWorkgroupsPerDimension)"
           }
         }
       ]
@@ -533,16 +649,29 @@
       "id": "transB",
       "priority": 0,
       "when": ["noBiasContract", "transAFlag == false", "transBFlag == true"],
-      "constants": { "transA": false, "transB": true, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm",
           "shader": "fused-gemm.wgsl.jinja",
-          "bindings": "noBias",
           "dispatch": {
-            "x": "min(ceilDiv(gemmN, 64), device.limits.maxComputeWorkgroupsPerDimension)",
-            "y": "min(ceilDiv(gemmM, 64), device.limits.maxComputeWorkgroupsPerDimension)"
           }
         }
       ]
@@ -551,7 +680,7 @@
       "id": "transAB_bias",
       "priority": 0,
       "when": ["biasContract", "transAFlag == true", "transBFlag == true"],
-      "constants": {
         "transA": true,
         "transB": true,
         "hasBias": true,
@@ -564,10 +693,25 @@
           "id": "main",
           "name": "FusedGemm.Bias",
           "shader": "fused-gemm.wgsl.jinja",
-          "bindings": "withBias",
           "dispatch": {
-            "x": "min(ceilDiv(gemmN, 64), device.limits.maxComputeWorkgroupsPerDimension)",
-            "y": "min(ceilDiv(gemmM, 64), device.limits.maxComputeWorkgroupsPerDimension)"
           }
         }
       ]
@@ -576,16 +720,29 @@
       "id": "transAB",
       "priority": 0,
       "when": ["noBiasContract", "transAFlag == true", "transBFlag == true"],
-      "constants": { "transA": true, "transB": true, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm",
           "shader": "fused-gemm.wgsl.jinja",
-          "bindings": "noBias",
           "dispatch": {
-            "x": "min(ceilDiv(gemmN, 64), device.limits.maxComputeWorkgroupsPerDimension)",
-            "y": "min(ceilDiv(gemmM, 64), device.limits.maxComputeWorkgroupsPerDimension)"
           }
         }
       ]

   "domain": "com.microsoft",
   "name": "FusedGemm",
   "sinceVersion": 1,
+  "inputs": {
+    "A": { "dtype": "T", "rank": 2 },
+    "B": { "dtype": "T", "rank": 2 },
+    "C": { "dtype": "T", "optional": true }
+  },
+  "outputs": { "Y": { "dtype": "T", "rank": 2, "shape": "[gemmM, gemmN]" } },
+  "attributes": {
+    "alpha": { "default": 1 },
+    "beta": { "default": 1 },
+    "transA": { "default": 0 },
+    "transB": { "default": 0 },
+    "activation": {},
+    "activation_alpha": {},
+    "activation_beta": {}
   },
   "attributeConstraints": { "activation": { "values": ["Relu", "LeakyRelu", "Sigmoid", "Tanh", "HardSigmoid"] } },
   "typeConstraints": { "T": ["float32", "float16"] },
   "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",
     "noBiasContract": "baseContract and not present.C",
     "biasContract": "baseContract and present.C and cOk",
     "sgmatOperandBytes": "2 if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else 4",
+    "sgmatTileM": "32",
+    "sgmatTileN": "64",
+    "sgmatTileK": "32",
+    "sgmatDirectA": "tensorDtypes.A == \"float32\" and not transAFlag and gemmK > 0 and gemmM % sgmatTileM == 0",
+    "sgmatStorageBytes": "((0 if sgmatDirectA else sgmatTileM) + sgmatTileN) * sgmatTileK * sgmatOperandBytes + sgmatTileM * sgmatTileN * 4",
     "sgmatResourcesFit": "128 <= deviceWorkgroupCap and sgmatStorageBytes <= device.limits.maxComputeWorkgroupStorageSize",
     "sgmatScalingOk": "(attrs.alpha if has(attrs, \"alpha\") else 1) == 1 and (attrs.beta if has(attrs, \"beta\") else 1) == 1",
+    "sgmatLayoutOk": "gemmM > 0 and gemmK % sgmatTileK == 0 and gemmN % sgmatTileN == 0 and ((tensorDtypes.A == \"float16\" and device.features.has(\"shader-f16\") and (gemmM >= 2 or (not transAFlag and transBFlag))) or (tensorDtypes.A == \"float32\" and gemmM >= sgmatTileM)) and ceilDiv(gemmM, sgmatTileM) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and ceilDiv(gemmN, sgmatTileN) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
     "sgmatContract": "wave32Effective and sgmatLayoutOk and sgmatScalingOk and sgmatResourcesFit",
+    "sgmatRowBiasContract": "sgmatContract and rowBiasOk and ranks.C == 1",
     "usesF16": "tensorDtypes.A == \"float16\"",
     "gemmActivation": "activationName",
     "gemmActAlpha": "activationAlpha",
     "gemmActBeta": "activationBeta",
+    "gemmEpilogue": "\"activation\" if hasActivation else \"none\"",
+    "gemvColumns": "4 if gemmN % 4 == 0 else 1",
+    "gemvWorkgroup": "min(pow(2, log2ceil(deviceWorkgroupCap + 1) - 1), pow2ceil(max(gemmK, 1)), 64)"
   },
+  "bindings": {
+    "x": { "arg": "A", "buffer": "read-only-storage", "elementType": "$T" },
+    "w": { "arg": "B", "buffer": "read-only-storage", "elementType": "$T" },
+    "bias": { "arg": "C", "buffer": "read-only-storage", "elementType": "$T" },
+    "y": { "arg": "Y", "buffer": "storage", "elementType": "$T" },
+    "params": { "buffer": "uniform", "struct": [{ "name": "M", "type": "u32", "value": "gemmM" }] },
+    "a": { "arg": "A", "buffer": "read-only-storage", "elementType": "$T" },
+    "b": { "arg": "B", "buffer": "read-only-storage", "elementType": "$T" }
   },
   "variants": [
+    {
+      "id": "gemv_plain",
+      "priority": 200,
+      "when": ["noBiasContract", "gemmM == 1 and gemmN > 0 and gemmK > 0 and transBFlag", "gemvWorkgroup * gemvColumns * 4 <= device.limits.maxComputeWorkgroupStorageSize"],
+      "derive": {
+        "hasBias": false,
+        "workgroupSize": "gemvWorkgroup",
+        "gemvFused": true,
+        "inFeatures": "gemmK",
+        "outFeatures": "gemmN",
+        "xScalar": "\"f16\" if usesF16 else \"f32\"",
+        "wScalar": "\"f16\" if usesF16 else \"f32\"",
+        "yScalar": "\"f16\" if usesF16 else \"f32\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "FusedGemm.Gemv",
+          "shader": "gemv-contiguous.wgsl.jinja",
+          "bindings": [
+            "x",
+            "w",
+            "y",
+            {
+              "name": "params",
+              "struct": [{ "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" }]
+            }
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv(gemmN, gemvColumns), DISPATCH_FOLD_WIDTH)",
+            "y": "1",
+            "z": "ceilDiv(ceilDiv(gemmN, gemvColumns), DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "gemv_bias",
+      "priority": 200,
+      "when": ["biasContract", "gemmM == 1 and gemmN > 0 and gemmK > 0 and transBFlag", "gemvWorkgroup * gemvColumns * 4 <= device.limits.maxComputeWorkgroupStorageSize"],
+      "derive": {
+        "hasBias": true,
+        "workgroupSize": "gemvWorkgroup",
+        "gemvFused": true,
+        "inFeatures": "gemmK",
+        "outFeatures": "gemmN",
+        "xScalar": "\"f16\" if usesF16 else \"f32\"",
+        "wScalar": "\"f16\" if usesF16 else \"f32\"",
+        "yScalar": "\"f16\" if usesF16 else \"f32\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "FusedGemm.Gemv",
+          "shader": "gemv-contiguous.wgsl.jinja",
+          "bindings": [
+            "x",
+            "w",
+            "bias",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" },
+                { "name": "beta", "type": "f32", "value": "attrs.beta if has(attrs, \"beta\") else 1" }
+              ]
+            }
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv(gemmN, gemvColumns), DISPATCH_FOLD_WIDTH)",
+            "y": "1",
+            "z": "ceilDiv(ceilDiv(gemmN, gemvColumns), DISPATCH_FOLD_WIDTH)"
+          }
+        }
+      ]
+    },
     {
       "id": "notrans_sgmat_bias",
       "priority": 100,
+      "when": ["sgmatRowBiasContract", "transAFlag == false", "transBFlag == false"],
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
+      "derive": {
         "hasBias": true,
         "aTransposedStorage": false,
         "bTransposedStorage": "not false",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
+        "outFeatures": "gemmN",
+        "tileM": "sgmatTileM",
+        "tileN": "sgmatTileN",
+        "tileK": "sgmatTileK",
+        "directAStorage": "sgmatDirectA",
+        "compensatedMatrixChunkTiles": "2 if sgmatDirectA and gemmK >= 8 * sgmatTileK else 0"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrixBias",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
+          "bindings": ["x", "w", "bias", "y", "params"],
+          "dispatch": {
+            "x": "ceilDiv(gemmN, sgmatTileN)",
+            "y": "min(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)",
+            "z": "ceilDiv(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)"
+          }
         }
       ]
     },
     {
       "id": "notrans_sgmat",
       "priority": 100,
+      "when": ["noBiasContract and sgmatContract", "transAFlag == false", "transBFlag == false"],
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
+      "derive": {
         "hasBias": false,
         "aTransposedStorage": false,
         "bTransposedStorage": "not false",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
+        "outFeatures": "gemmN",
+        "tileM": "sgmatTileM",
+        "tileN": "sgmatTileN",
+        "tileK": "sgmatTileK",
+        "directAStorage": "sgmatDirectA",
+        "compensatedMatrixChunkTiles": "2 if sgmatDirectA and gemmK >= 8 * sgmatTileK else 0"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrix",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
+          "bindings": ["x", "w", "y", "params"],
+          "dispatch": {
+            "x": "ceilDiv(gemmN, sgmatTileN)",
+            "y": "min(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)",
+            "z": "ceilDiv(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)"
+          }
         }
       ]
     },
     {
       "id": "transA_sgmat_bias",
       "priority": 100,
+      "when": ["sgmatRowBiasContract", "transAFlag == true", "transBFlag == false"],
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
+      "derive": {
         "hasBias": true,
         "aTransposedStorage": true,
         "bTransposedStorage": "not false",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
+        "outFeatures": "gemmN",
+        "tileM": "sgmatTileM",
+        "tileN": "sgmatTileN",
+        "tileK": "sgmatTileK",
+        "directAStorage": "sgmatDirectA",
+        "compensatedMatrixChunkTiles": "2 if sgmatDirectA and gemmK >= 8 * sgmatTileK else 0"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrixBias",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
+          "bindings": ["x", "w", "bias", "y", "params"],
+          "dispatch": {
+            "x": "ceilDiv(gemmN, sgmatTileN)",
+            "y": "min(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)",
+            "z": "ceilDiv(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)"
+          }
         }
       ]
     },
     {
       "id": "transA_sgmat",
       "priority": 100,
+      "when": ["noBiasContract and sgmatContract", "transAFlag == true", "transBFlag == false"],
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
+      "derive": {
         "hasBias": false,
         "aTransposedStorage": true,
         "bTransposedStorage": "not false",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
+        "outFeatures": "gemmN",
+        "tileM": "sgmatTileM",
+        "tileN": "sgmatTileN",
+        "tileK": "sgmatTileK",
+        "directAStorage": "sgmatDirectA",
+        "compensatedMatrixChunkTiles": "2 if sgmatDirectA and gemmK >= 8 * sgmatTileK else 0"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrix",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
+          "bindings": ["x", "w", "y", "params"],
+          "dispatch": {
+            "x": "ceilDiv(gemmN, sgmatTileN)",
+            "y": "min(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)",
+            "z": "ceilDiv(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)"
+          }
         }
       ]
     },
     {
       "id": "transB_sgmat_bias",
       "priority": 100,
+      "when": ["sgmatRowBiasContract", "transAFlag == false", "transBFlag == true"],
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
+      "derive": {
         "hasBias": true,
         "aTransposedStorage": false,
         "bTransposedStorage": "not true",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
+        "outFeatures": "gemmN",
+        "tileM": "sgmatTileM",
+        "tileN": "sgmatTileN",
+        "tileK": "sgmatTileK",
+        "directAStorage": "sgmatDirectA",
+        "compensatedMatrixChunkTiles": "2 if sgmatDirectA and gemmK >= 8 * sgmatTileK else 0"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrixBias",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
+          "bindings": ["x", "w", "bias", "y", "params"],
+          "dispatch": {
+            "x": "ceilDiv(gemmN, sgmatTileN)",
+            "y": "min(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)",
+            "z": "ceilDiv(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)"
+          }
         }
       ]
     },
     {
       "id": "transB_sgmat",
       "priority": 100,
+      "when": ["noBiasContract and sgmatContract", "transAFlag == false", "transBFlag == true"],
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
+      "derive": {
         "hasBias": false,
         "aTransposedStorage": false,
         "bTransposedStorage": "not true",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
+        "outFeatures": "gemmN",
+        "tileM": "sgmatTileM",
+        "tileN": "sgmatTileN",
+        "tileK": "sgmatTileK",
+        "directAStorage": "sgmatDirectA",
+        "compensatedMatrixChunkTiles": "2 if sgmatDirectA and gemmK >= 8 * sgmatTileK else 0"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrix",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
+          "bindings": ["x", "w", "y", "params"],
+          "dispatch": {
+            "x": "ceilDiv(gemmN, sgmatTileN)",
+            "y": "min(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)",
+            "z": "ceilDiv(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)"
+          }
         }
       ]
     },
     {
       "id": "transAB_sgmat_bias",
       "priority": 100,
+      "when": ["sgmatRowBiasContract", "transAFlag == true", "transBFlag == true"],
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
+      "derive": {
         "hasBias": true,
         "aTransposedStorage": true,
         "bTransposedStorage": "not true",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
+        "outFeatures": "gemmN",
+        "tileM": "sgmatTileM",
+        "tileN": "sgmatTileN",
+        "tileK": "sgmatTileK",
+        "directAStorage": "sgmatDirectA",
+        "compensatedMatrixChunkTiles": "2 if sgmatDirectA and gemmK >= 8 * sgmatTileK else 0"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrixBias",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
+          "bindings": ["x", "w", "bias", "y", "params"],
+          "dispatch": {
+            "x": "ceilDiv(gemmN, sgmatTileN)",
+            "y": "min(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)",
+            "z": "ceilDiv(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)"
+          }
         }
       ]
     },
     {
       "id": "transAB_sgmat",
       "priority": 100,
+      "when": ["noBiasContract and sgmatContract", "transAFlag == true", "transBFlag == true"],
       "requires": {
         "features": ["subgroups", "chromium-experimental-subgroup-matrix"],
         "subgroupMatrixConfigs": [
           { "componentType": "f32", "resultComponentType": "f32", "M": 8, "N": 8, "K": 8 }
         ]
       },
+      "derive": {
         "hasBias": false,
         "aTransposedStorage": true,
         "bTransposedStorage": "not true",
         "yScalar": "\"f16\" if tensorDtypes.Y == \"float16\" else \"f32\"",
         "mmaScalar": "\"f16\" if tensorDtypes.A == \"float16\" and tensorDtypes.B == \"float16\" else \"f32\"",
         "inFeatures": "gemmK",
+        "outFeatures": "gemmN",
+        "tileM": "sgmatTileM",
+        "tileN": "sgmatTileN",
+        "tileK": "sgmatTileK",
+        "directAStorage": "sgmatDirectA",
+        "compensatedMatrixChunkTiles": "2 if sgmatDirectA and gemmK >= 8 * sgmatTileK else 0"
       },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm.SubgroupMatrix",
           "shader": "gemm-subgroup-matrix.wgsl.jinja",
+          "bindings": ["x", "w", "y", "params"],
+          "dispatch": {
+            "x": "ceilDiv(gemmN, sgmatTileN)",
+            "y": "min(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)",
+            "z": "ceilDiv(ceilDiv(gemmM, sgmatTileM), DISPATCH_FOLD_WIDTH)"
+          }
         }
       ]
     },
       "id": "notrans_bias",
       "priority": 0,
       "when": ["biasContract", "transAFlag == false", "transBFlag == false"],
+      "derive": {
         "transA": false,
         "transB": false,
         "hasBias": true,
           "id": "main",
           "name": "FusedGemm.Bias",
           "shader": "fused-gemm.wgsl.jinja",
+          "bindings": [
+            "a",
+            "b",
+            "bias",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "name": "M", "type": "u32", "value": "gemmM" },
+                { "name": "N", "type": "u32", "value": "gemmN" },
+                { "name": "K", "type": "u32", "value": "gemmK" },
+                { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" },
+                { "name": "beta", "type": "f32", "value": "attrs.beta if has(attrs, \"beta\") else 1" }
+              ]
+            }
+          ],
           "dispatch": {
+            "x": "min(ceilDiv(gemmN, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": "min(ceilDiv(gemmM, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
           }
         }
       ]
       "id": "notrans",
       "priority": 0,
       "when": ["noBiasContract", "transAFlag == false", "transBFlag == false"],
+      "derive": { "transA": false, "transB": false, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm",
           "shader": "fused-gemm.wgsl.jinja",
+          "bindings": [
+            "a",
+            "b",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "name": "M", "type": "u32", "value": "gemmM" },
+                { "name": "N", "type": "u32", "value": "gemmN" },
+                { "name": "K", "type": "u32", "value": "gemmK" },
+                { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" }
+              ]
+            }
+          ],
           "dispatch": {
+            "x": "min(ceilDiv(gemmN, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": "min(ceilDiv(gemmM, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
           }
         }
       ]
       "id": "transA_bias",
       "priority": 0,
       "when": ["biasContract", "transAFlag == true", "transBFlag == false"],
+      "derive": {
         "transA": true,
         "transB": false,
         "hasBias": true,
           "id": "main",
           "name": "FusedGemm.Bias",
           "shader": "fused-gemm.wgsl.jinja",
+          "bindings": [
+            "a",
+            "b",
+            "bias",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "name": "M", "type": "u32", "value": "gemmM" },
+                { "name": "N", "type": "u32", "value": "gemmN" },
+                { "name": "K", "type": "u32", "value": "gemmK" },
+                { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" },
+                { "name": "beta", "type": "f32", "value": "attrs.beta if has(attrs, \"beta\") else 1" }
+              ]
+            }
+          ],
           "dispatch": {
+            "x": "min(ceilDiv(gemmN, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": "min(ceilDiv(gemmM, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
           }
         }
       ]
       "id": "transA",
       "priority": 0,
       "when": ["noBiasContract", "transAFlag == true", "transBFlag == false"],
+      "derive": { "transA": true, "transB": false, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm",
           "shader": "fused-gemm.wgsl.jinja",
+          "bindings": [
+            "a",
+            "b",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "name": "M", "type": "u32", "value": "gemmM" },
+                { "name": "N", "type": "u32", "value": "gemmN" },
+                { "name": "K", "type": "u32", "value": "gemmK" },
+                { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" }
+              ]
+            }
+          ],
           "dispatch": {
+            "x": "min(ceilDiv(gemmN, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": "min(ceilDiv(gemmM, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
           }
         }
       ]
       "id": "transB_bias",
       "priority": 0,
       "when": ["biasContract", "transAFlag == false", "transBFlag == true"],
+      "derive": {
         "transA": false,
         "transB": true,
         "hasBias": true,
           "id": "main",
           "name": "FusedGemm.Bias",
           "shader": "fused-gemm.wgsl.jinja",
+          "bindings": [
+            "a",
+            "b",
+            "bias",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "name": "M", "type": "u32", "value": "gemmM" },
+                { "name": "N", "type": "u32", "value": "gemmN" },
+                { "name": "K", "type": "u32", "value": "gemmK" },
+                { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" },
+                { "name": "beta", "type": "f32", "value": "attrs.beta if has(attrs, \"beta\") else 1" }
+              ]
+            }
+          ],
           "dispatch": {
+            "x": "min(ceilDiv(gemmN, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": "min(ceilDiv(gemmM, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
           }
         }
       ]
       "id": "transB",
       "priority": 0,
       "when": ["noBiasContract", "transAFlag == false", "transBFlag == true"],
+      "derive": { "transA": false, "transB": true, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm",
           "shader": "fused-gemm.wgsl.jinja",
+          "bindings": [
+            "a",
+            "b",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "name": "M", "type": "u32", "value": "gemmM" },
+                { "name": "N", "type": "u32", "value": "gemmN" },
+                { "name": "K", "type": "u32", "value": "gemmK" },
+                { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" }
+              ]
+            }
+          ],
           "dispatch": {
+            "x": "min(ceilDiv(gemmN, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": "min(ceilDiv(gemmM, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
           }
         }
       ]
       "id": "transAB_bias",
       "priority": 0,
       "when": ["biasContract", "transAFlag == true", "transBFlag == true"],
+      "derive": {
         "transA": true,
         "transB": true,
         "hasBias": true,
           "id": "main",
           "name": "FusedGemm.Bias",
           "shader": "fused-gemm.wgsl.jinja",
+          "bindings": [
+            "a",
+            "b",
+            "bias",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "name": "M", "type": "u32", "value": "gemmM" },
+                { "name": "N", "type": "u32", "value": "gemmN" },
+                { "name": "K", "type": "u32", "value": "gemmK" },
+                { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" },
+                { "name": "beta", "type": "f32", "value": "attrs.beta if has(attrs, \"beta\") else 1" }
+              ]
+            }
+          ],
           "dispatch": {
+            "x": "min(ceilDiv(gemmN, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": "min(ceilDiv(gemmM, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
           }
         }
       ]
       "id": "transAB",
       "priority": 0,
       "when": ["noBiasContract", "transAFlag == true", "transBFlag == true"],
+      "derive": { "transA": true, "transB": true, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FusedGemm",
           "shader": "fused-gemm.wgsl.jinja",
+          "bindings": [
+            "a",
+            "b",
+            "y",
+            {
+              "name": "params",
+              "struct": [
+                { "name": "M", "type": "u32", "value": "gemmM" },
+                { "name": "N", "type": "u32", "value": "gemmN" },
+                { "name": "K", "type": "u32", "value": "gemmK" },
+                { "name": "alpha", "type": "f32", "value": "attrs.alpha if has(attrs, \"alpha\") else 1" }
+              ]
+            }
+          ],
           "dispatch": {
+            "x": "min(ceilDiv(gemmN, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": "min(ceilDiv(gemmM, 64), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
           }
         }
       ]

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,42 @@
 {
   "name": "com.microsoft.FusedGemm",
-  "id": "_com_microsoft_fusedgemm_webgpu_61fbd6b",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "5+Pgon1ED+99xPL6smxxz2G47C5gDnfB0kJ+Zj2wZLc=",
-      "fused-gemm.wgsl.jinja": "XL5rMA9oUlUJZCMHq2PZAzmORxFZoz/xa5Mo05/aK3k=",
-      "gemm-subgroup-matrix.wgsl.jinja": "PF4rqax3Fvj0sfPRjj+0BeuUkfoJVM2Bt6r7yz4Q2y0=",
-      "manifest.json": "s7cSq5so+VQS0XTbiHalmSHGnQQA2ylTu524o3V815U=",
-      "test.json": "LRhFk3xI+8tnvhAwEvFnept0OjRWCvXhjpjQ31ZIK0E="
     }
   },
-  "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.FusedGemm" }
 }

 {
   "name": "com.microsoft.FusedGemm",
+  "id": "_com_microsoft_fusedgemm_webgpu_abf1d0b",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "wb8UHX4mLph1fuImZI56hUkFg3rsazdLNuDl69XrKdM=",
+      "fused-gemm.wgsl.jinja": "A7FcGNQgqc57hkI4skKQutwPNJzfif1nI3C2f36hZcc=",
+      "gemm-subgroup-matrix.wgsl.jinja": "PSvZd1vPhlsbysyRrWuBMC4iHBhQq6VuOjGmFWsF2K8=",
+      "gemv-contiguous.wgsl.jinja": "sBwHBkgw/X6QKltFcxbi4TSPvrCh5xM8QAGZIt8nMzg=",
+      "manifest.json": "JSX1Yq1ypgxc4wVUAoDBbUKhyOn5NiuF5pzNihs85SY=",
+      "test.json": "j+1eYKc0EE8Rt4YEY/oS5mEuBucLF/pRQCkqnb0zB6Q="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "gemv_plain": ["gemv-contiguous.wgsl.jinja"],
+      "gemv_bias": ["gemv-contiguous.wgsl.jinja"],
+      "notrans_sgmat_bias": ["gemm-subgroup-matrix.wgsl.jinja"],
+      "notrans_sgmat": ["gemm-subgroup-matrix.wgsl.jinja"],
+      "transA_sgmat_bias": ["gemm-subgroup-matrix.wgsl.jinja"],
+      "transA_sgmat": ["gemm-subgroup-matrix.wgsl.jinja"],
+      "transB_sgmat_bias": ["gemm-subgroup-matrix.wgsl.jinja"],
+      "transB_sgmat": ["gemm-subgroup-matrix.wgsl.jinja"],
+      "transAB_sgmat_bias": ["gemm-subgroup-matrix.wgsl.jinja"],
+      "transAB_sgmat": ["gemm-subgroup-matrix.wgsl.jinja"],
+      "notrans_bias": ["fused-gemm.wgsl.jinja"],
+      "notrans": ["fused-gemm.wgsl.jinja"],
+      "transA_bias": ["fused-gemm.wgsl.jinja"],
+      "transA": ["fused-gemm.wgsl.jinja"],
+      "transB_bias": ["fused-gemm.wgsl.jinja"],
+      "transB": ["fused-gemm.wgsl.jinja"],
+      "transAB_bias": ["fused-gemm.wgsl.jinja"],
+      "transAB": ["fused-gemm.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.FusedGemm",
   "cases": [
     {
       "name": "notrans_rowbias_relu",
@@ -95,7 +94,7 @@
     },
     {
       "name": "notrans_nobias_sigmoid",
-      "provenance": { "notes": "No bias at all -- the other binding set -- with Sigmoid." },
       "attrs": { "activation": "Sigmoid" },
       "inputs": {
         "A": {
@@ -294,7 +293,7 @@
     },
     {
       "name": "beta_zero_relu",
-      "provenance": { "notes": "beta = 0 nulls the C term the same way." },
       "attrs": { "beta": 0, "activation": "Relu" },
       "inputs": {
         "A": {
@@ -481,7 +480,7 @@
     {
       "name": "sgmat_notrans_rowbias_relu",
       "provenance": {
-        "notes": "Aligned for the subgroup-matrix gate (M >= 32, K % 32 == 0, N % 64 == 0), so this selects notrans_sgmat_bias where the unaligned fixtures above stay on the tiled kernel. Stages A as [M,K] and B as [K,N], which is the operand orientation this transpose pair asks the matrix kernel for, and fuses Relu in the accumulator domain over a rank-1 row bias."
       },
       "attrs": { "activation": "Relu" },
       "inputs": {
@@ -506,7 +505,7 @@
     {
       "name": "sgmat_notrans_nobias_tanh",
       "provenance": {
-        "notes": "The no-bias arm of the same gate and layout: the matrix kernel renders a separate set of store sites without the bias term."
       },
       "attrs": { "activation": "Tanh" },
       "inputs": {
@@ -526,7 +525,7 @@
     {
       "name": "notrans_rowbias_tanh_saturating_preactivation",
       "provenance": {
-        "notes": "Saturating pre-activations: the row bias is large enough that most |alpha*A*B + beta*C| land past 44, where the hardware tanh's internal e^2v overflows and the quotient becomes Inf/Inf = NaN. The bias carries the magnitude rather than the operands so the accumulator stays small and well conditioned — large operands make the near-cancelling dot products disagree between accumulation orders, which is a property of the fixture rather than of the activation. Unaligned for the subgroup-matrix gate, so this covers the tiled store path."
       },
       "attrs": { "activation": "Tanh" },
       "inputs": {
@@ -551,7 +550,7 @@
     {
       "name": "sgmat_notrans_rowbias_tanh_saturating_preactivation",
       "provenance": {
-        "notes": "The aligned twin of the saturating fixture above (M >= 32, K % 32 == 0, N % 64 == 0), so the subgroup-matrix route evaluates the same overflowing pre-activations through its own store sites."
       },
       "attrs": { "activation": "Tanh" },
       "inputs": {
@@ -576,7 +575,7 @@
     {
       "name": "sgmat_transa_rowbias_tanh",
       "provenance": {
-        "notes": "Aligned for the subgroup-matrix gate (M >= 32, K % 32 == 0, N % 64 == 0), so this selects transa_sgmat_bias where the unaligned fixtures above stay on the tiled kernel. Stages A as [K,M] and B as [K,N], which is the operand orientation this transpose pair asks the matrix kernel for, and fuses Tanh in the accumulator domain over a rank-1 row bias."
       },
       "attrs": { "transA": 1, "activation": "Tanh" },
       "inputs": {
@@ -601,7 +600,7 @@
     {
       "name": "sgmat_transa_nobias_sigmoid",
       "provenance": {
-        "notes": "The no-bias arm of the same gate and layout: the matrix kernel renders a separate set of store sites without the bias term."
       },
       "attrs": { "transA": 1, "activation": "Sigmoid" },
       "inputs": {
@@ -621,7 +620,7 @@
     {
       "name": "sgmat_transb_rowbias_sigmoid",
       "provenance": {
-        "notes": "Aligned for the subgroup-matrix gate (M >= 32, K % 32 == 0, N % 64 == 0), so this selects transb_sgmat_bias where the unaligned fixtures above stay on the tiled kernel. Stages A as [M,K] and B as [N,K], which is the operand orientation this transpose pair asks the matrix kernel for, and fuses Sigmoid in the accumulator domain over a rank-1 row bias."
       },
       "attrs": { "transB": 1, "activation": "Sigmoid" },
       "inputs": {
@@ -646,7 +645,7 @@
     {
       "name": "sgmat_transb_nobias_leakyrelu",
       "provenance": {
-        "notes": "The no-bias arm of the same gate and layout: the matrix kernel renders a separate set of store sites without the bias term."
       },
       "attrs": { "transB": 1, "activation": "LeakyRelu" },
       "inputs": {
@@ -666,7 +665,7 @@
     {
       "name": "sgmat_transab_rowbias_leakyrelu",
       "provenance": {
-        "notes": "Aligned for the subgroup-matrix gate (M >= 32, K % 32 == 0, N % 64 == 0), so this selects transab_sgmat_bias where the unaligned fixtures above stay on the tiled kernel. Stages A as [K,M] and B as [N,K], which is the operand orientation this transpose pair asks the matrix kernel for, and fuses LeakyRelu in the accumulator domain over a rank-1 row bias."
       },
       "attrs": { "transA": 1, "transB": 1, "activation": "LeakyRelu" },
       "inputs": {
@@ -691,7 +690,7 @@
     {
       "name": "sgmat_transab_nobias_relu",
       "provenance": {
-        "notes": "The no-bias arm of the same gate and layout: the matrix kernel renders a separate set of store sites without the bias term."
       },
       "attrs": { "transA": 1, "transB": 1, "activation": "Relu" },
       "inputs": {
@@ -707,6 +706,1944 @@
         }
       },
       "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
     }
   ]
 }

 {
   "cases": [
     {
       "name": "notrans_rowbias_relu",
     },
     {
       "name": "notrans_nobias_sigmoid",
+      "provenance": { "notes": "No bias, no transposes, and a fused Sigmoid activation." },
       "attrs": { "activation": "Sigmoid" },
       "inputs": {
         "A": {
     },
     {
       "name": "beta_zero_relu",
+      "provenance": { "notes": "`beta = 0` suppresses the C term exactly." },
       "attrs": { "beta": 0, "activation": "Relu" },
       "inputs": {
         "A": {
     {
       "name": "sgmat_notrans_rowbias_relu",
       "provenance": {
+        "notes": "Aligned `M >= 32`, `K % 32 = 0`, and `N % 64 = 0` select the subgroup-matrix route. A is stored as [M, K], B as [K, N], and Relu is fused over a rank-1 row bias."
       },
       "attrs": { "activation": "Relu" },
       "inputs": {
     {
       "name": "sgmat_notrans_nobias_tanh",
       "provenance": {
+        "notes": "Aligned non-transposed operands exercise the no-bias subgroup-matrix store with fused Tanh."
       },
       "attrs": { "activation": "Tanh" },
       "inputs": {
     {
       "name": "notrans_rowbias_tanh_saturating_preactivation",
       "provenance": {
+        "notes": "A large row bias drives most pre-activations beyond magnitude 44 while the dot products remain small and well conditioned. Unaligned dimensions exercise guarded Tanh saturation on the tiled store path."
       },
       "attrs": { "activation": "Tanh" },
       "inputs": {
     {
       "name": "sgmat_notrans_rowbias_tanh_saturating_preactivation",
       "provenance": {
+        "notes": "Aligned dimensions and a large row bias exercise guarded Tanh saturation on the subgroup-matrix store path."
       },
       "attrs": { "activation": "Tanh" },
       "inputs": {
     {
       "name": "sgmat_transa_rowbias_tanh",
       "provenance": {
+        "notes": "Aligned dimensions select the subgroup-matrix route with A stored as [K, M], B as [K, N], a rank-1 row bias, and fused Tanh."
       },
       "attrs": { "transA": 1, "activation": "Tanh" },
       "inputs": {
     {
       "name": "sgmat_transa_nobias_sigmoid",
       "provenance": {
+        "notes": "Aligned dimensions exercise the transposed-A, no-bias subgroup-matrix store with fused Sigmoid."
       },
       "attrs": { "transA": 1, "activation": "Sigmoid" },
       "inputs": {
     {
       "name": "sgmat_transb_rowbias_sigmoid",
       "provenance": {
+        "notes": "Aligned dimensions select the subgroup-matrix route with A stored as [M, K], B as [N, K], a rank-1 row bias, and fused Sigmoid."
       },
       "attrs": { "transB": 1, "activation": "Sigmoid" },
       "inputs": {
     {
       "name": "sgmat_transb_nobias_leakyrelu",
       "provenance": {
+        "notes": "Aligned dimensions exercise the transposed-B, no-bias subgroup-matrix store with fused LeakyRelu."
       },
       "attrs": { "transB": 1, "activation": "LeakyRelu" },
       "inputs": {
     {
       "name": "sgmat_transab_rowbias_leakyrelu",
       "provenance": {
+        "notes": "Aligned dimensions select the subgroup-matrix route with A stored as [K, M], B as [N, K], a rank-1 row bias, and fused LeakyRelu."
       },
       "attrs": { "transA": 1, "transB": 1, "activation": "LeakyRelu" },
       "inputs": {
     {
       "name": "sgmat_transab_nobias_relu",
       "provenance": {
+        "notes": "Aligned dimensions exercise the doubly transposed, no-bias subgroup-matrix store with fused Relu."
       },
       "attrs": { "transA": 1, "transB": 1, "activation": "Relu" },
       "inputs": {
         }
       },
       "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "f16_transb_rowbias_relu_m1_sgmat_floor_pathology",
+      "provenance": {
+        "notes": "A single float16 decode row with aligned K and N remains below the subgroup-matrix row floor and exercises the generic register tile."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_transb_rowbias_relu_m2_sgmat_floor_control",
+      "provenance": {
+        "notes": "Two float16 rows with aligned K and N meet the subgroup-matrix row floor and exercise that route."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [2, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [2, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_nobias_identity",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1 },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_nobias_Relu",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_nobias_LeakyRelu",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1, "activation": "LeakyRelu", "activation_alpha": 0.125 },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_nobias_Sigmoid",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1, "activation": "Sigmoid" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_nobias_Tanh",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1, "activation": "Tanh" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_nobias_HardSigmoid",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1, "activation": "HardSigmoid", "activation_alpha": 0.15, "activation_beta": 0.4 },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_empty_reduction_nobias",
+      "provenance": { "notes": "An empty reduction leaves the additive term, or zero without C, before applying Relu." },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "dtype": "float16", "shape": [1, 0], "data": { "kind": "constant", "value": 0.0 } },
+        "B": { "dtype": "float16", "shape": [128, 0], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_bias_identity",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1 },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_bias_LeakyRelu",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1, "activation": "LeakyRelu", "activation_alpha": 0.125 },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_bias_Sigmoid",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1, "activation": "Sigmoid" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_bias_Tanh",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1, "activation": "Tanh" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_bias_HardSigmoid",
+      "provenance": {
+        "notes": "A single output row with contiguous reduction dimensions in both operands. The activation is applied to the f32 accumulator before the final f16 store."
+      },
+      "attrs": { "transB": 1, "activation": "HardSigmoid", "activation_alpha": 0.15, "activation_beta": 0.4 },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_single_row_empty_reduction_bias",
+      "provenance": { "notes": "An empty reduction leaves the additive term, or zero without C, before applying Relu." },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "dtype": "float16", "shape": [1, 0], "data": { "kind": "constant", "value": 0.0 } },
+        "B": { "dtype": "float16", "shape": [128, 0], "data": { "kind": "constant", "value": 0.0 } },
+        "C": {
+          "dtype": "float16",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 128], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float32-k64-n9-none-identity-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0,
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [9, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 9], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k63-n9-scalar-Relu-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0,
+        "activation": "Relu",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [63, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [9, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 9], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k64-n9-scalar1d-LeakyRelu-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0,
+        "activation": "LeakyRelu",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [9, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [1],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 9], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k63-n9-scalar2d-Sigmoid-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0,
+        "activation": "Sigmoid",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [9, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 9], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k64-n9-row-Tanh-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0,
+        "activation": "Tanh",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [9, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [9],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 9], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k63-n9-row2d-HardSigmoid-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0,
+        "activation": "HardSigmoid",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [63, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [9, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [1, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 9], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k64-n12-none-Tanh-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "Tanh",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [12, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 12], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k63-n12-scalar-HardSigmoid-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "HardSigmoid",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [63, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [12, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 12], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k64-n12-scalar1d-identity-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [12, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [1],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 12], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k63-n12-scalar2d-Relu-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "Relu",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [63, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [12, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 12], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k64-n12-row-LeakyRelu-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "LeakyRelu",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [12, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 12], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float32-k63-n12-row2d-Sigmoid-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "Sigmoid",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [12, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [1, 12],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 12], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-probe-float16-k63-n9-none-HardSigmoid-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0,
+        "activation": "HardSigmoid",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [9, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 9], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k64-n9-scalar-identity-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [64, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [9, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 9], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k63-n9-scalar1d-Relu-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "Relu",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [9, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [1],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 9], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k64-n9-scalar2d-LeakyRelu-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "LeakyRelu",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [9, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 9], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k63-n9-row-Sigmoid-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "Sigmoid",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [63, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [9, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [9],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 9], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k64-n9-row2d-Tanh-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "Tanh",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [64, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [9, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [1, 9],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 9], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k63-n12-none-Sigmoid-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "Sigmoid",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [63, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [12, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 12], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k64-n12-scalar-Tanh-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "Tanh",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [64, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [12, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 12], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k63-n12-scalar1d-HardSigmoid-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0.25,
+        "activation": "HardSigmoid",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [12, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [1],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 12], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k64-n12-scalar2d-identity-ta1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0,
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [64, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [12, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 12], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k63-n12-row-Relu-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0,
+        "activation": "Relu",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [12, 63],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [12],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 12], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-probe-float16-k64-n12-row2d-LeakyRelu-ta0",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": {
+        "transA": 0,
+        "transB": 1,
+        "alpha": -0.75,
+        "beta": 0,
+        "activation": "LeakyRelu",
+        "activation_alpha": 0.125,
+        "activation_beta": 0.375
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [12, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [1, 12],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 12], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-float32-alpha-zero-nonfinite-input",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": { "transB": 1, "alpha": 0, "beta": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "dtype": "float32", "shape": [1, 8], "data": { "kind": "constant", "value": 1e+39 } },
+        "B": { "dtype": "float32", "shape": [12, 8], "data": { "kind": "constant", "value": 1.0 } },
+        "C": { "dtype": "float32", "shape": [12], "data": { "kind": "constant", "value": 2.0 } }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 12], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-float32-beta-zero-nonfinite-input",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": { "transB": 1, "alpha": 1, "beta": 0, "activation": "Relu" },
+      "inputs": {
+        "A": { "dtype": "float32", "shape": [1, 8], "data": { "kind": "constant", "value": 1.0 } },
+        "B": { "dtype": "float32", "shape": [12, 8], "data": { "kind": "constant", "value": 1.0 } },
+        "C": {
+          "dtype": "float32",
+          "shape": [12],
+          "data": { "kind": "constant", "value": 2.0, "nanStart": 0.0, "nanCount": 12.0 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 12], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-float16-alpha-zero-nonfinite-input",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": { "transB": 1, "alpha": 0, "beta": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "dtype": "float16", "shape": [1, 8], "data": { "kind": "constant", "value": 1e+39 } },
+        "B": { "dtype": "float16", "shape": [12, 8], "data": { "kind": "constant", "value": 1.0 } },
+        "C": { "dtype": "float16", "shape": [12], "data": { "kind": "constant", "value": 2.0 } }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 12], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-float16-beta-zero-nonfinite-input",
+      "provenance": {
+        "notes": "Single-row contiguous-weight reduction with an odd reduction or output tail, optional broadcast bias, scaling, and fused activation."
+      },
+      "attrs": { "transB": 1, "alpha": 1, "beta": 0, "activation": "Relu" },
+      "inputs": {
+        "A": { "dtype": "float16", "shape": [1, 8], "data": { "kind": "constant", "value": 1.0 } },
+        "B": { "dtype": "float16", "shape": [12, 8], "data": { "kind": "constant", "value": 1.0 } },
+        "C": {
+          "dtype": "float16",
+          "shape": [12],
+          "data": { "kind": "constant", "value": 2.0, "nanStart": 0.0, "nanCount": 12.0 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 12], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-boundary-float32-k1-n1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 1], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-boundary-float32-k7-n5",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [5, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 5], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-boundary-float32-k31-n33",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 31],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [33, 31],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 33], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-boundary-float32-k32-n32",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [32, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 32], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-boundary-float32-k33-n33",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 33],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [33, 33],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 33], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-boundary-float32-k129-n16",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 129],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [16, 129],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 16], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-boundary-float32-k1023-n257",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 1023],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [257, 1023],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 257], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-boundary-float32-k4096-n64",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 4096],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [64, 4096],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 64], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-boundary-float32-k1-n65537",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [65537, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 65537], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-boundary-float32-k1-n262148",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [262148, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 262148], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-boundary-float16-k1-n1",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 1], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-boundary-float16-k7-n5",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [5, 7],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 5], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-boundary-float16-k31-n33",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 31],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [33, 31],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 33], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-boundary-float16-k32-n32",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [32, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 32], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-boundary-float16-k33-n33",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 33],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [33, 33],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 33], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-boundary-float16-k129-n16",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 129],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [16, 129],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 16], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-boundary-float16-k1023-n257",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 1023],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [257, 1023],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 257], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-boundary-float16-k4096-n64",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 4096],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [64, 4096],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 64], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-boundary-float16-k1-n65537",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [65537, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 65537], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-boundary-float16-k1-n262148",
+      "provenance": {
+        "notes": "Single-row contiguous-weight projection covers reduction-tree widths, scalar column tails and folded output-column dispatch."
+      },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [262148, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 262148], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-empty-output-float32-m0-n8",
+      "provenance": { "notes": "Empty output with a contiguous stored weight reduction." },
+      "attrs": { "transB": 1 },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [0, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [8, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [0, 8], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-empty-output-float32-m1-n0",
+      "provenance": { "notes": "Empty output with a contiguous stored weight reduction." },
+      "attrs": { "transB": 1 },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [0, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 0], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "gemv-empty-output-float16-m0-n8",
+      "provenance": { "notes": "Empty output with a contiguous stored weight reduction." },
+      "attrs": { "transB": 1 },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [0, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [8, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [0, 8], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "gemv-empty-output-float16-m1-n0",
+      "provenance": { "notes": "Empty output with a contiguous stored weight reduction." },
+      "attrs": { "transB": 1 },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [1, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [0, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [1, 0], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "direct_a_m31_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [31, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.39, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [31, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m32_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [32, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.39, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [32, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m33_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [33, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.39, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [33, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m63_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [63, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.39, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [63, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m64_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.39, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m65_k96_n128_tb1",
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [65, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.39, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [65, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m64_k32_n64_tb0",
+      "attrs": { "transB": 0, "activation": "Tanh" },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 32],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 64], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m64_k0_n64_tb1",
+      "attrs": { "transB": 1, "activation": "HardSigmoid" },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 0],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [64, 0],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.39, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 64], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m32_k224_n64_tb0",
+      "attrs": { "transB": 0 },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [32, 224],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [224, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m32_k256_n64_tb1",
+      "attrs": { "transB": 1 },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [32, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [64, 256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m32_k288_n64_tb0",
+      "attrs": { "transB": 0 },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [32, 288],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [288, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m32_k768_n128_tb0",
+      "attrs": { "transB": 0, "activation": "Relu" },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [32, 768],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [768, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.39, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [32, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "direct_a_m64_k1024_n128_tb0",
+      "attrs": { "transB": 0, "activation": "Tanh" },
+      "provenance": {
+        "notes": "Complete A tiles and adjacent row/reduction boundaries exercise direct storage loads, compensated reduction blocks, and guarded staging with signed nonuniform data."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.037, "scale": 0.2 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [1024, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.043, "scale": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "compensated_negative_infinity",
+      "provenance": {
+        "notes": "A 257-element cycle spans 256-element rows: one row stays finite with distinct column results while others contain negative infinity. This detects misplaced outputs and requires infinite partial sums to keep zero compensation before later finite partials."
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [32, 256],
+          "data": {
+            "kind": "cycle",
+            "values": [-3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, 2.0, 3.0, -3.0, -2.0, -1.0, 0.0, 1.0, "-Infinity", 3.0, -3.0, -2.0, -1.0, 0.0, 1.0]
+          }
+        },
+        "B": { "dtype": "float32", "shape": [256, 64], "data": { "kind": "cycle", "values": [1.0, 2.0, 3.0] } }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [32, 64], "tolerance": 0 } }
     }
   ]
 }