sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +79 -0
build/webgpu/bench.json +69 -0
build/webgpu/fused-gemm.wgsl.jinja +186 -0
build/webgpu/gemm-subgroup-matrix.wgsl.jinja +269 -0
build/webgpu/manifest.json +594 -0
build/webgpu/metadata.json +19 -0
build/webgpu/test.json +662 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.FusedGemm
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 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.
+See the [ONNX Runtime `FusedGemm` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.FusedGemm) for the reference semantics.
+## 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
+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
+| Variable | Allowed dtypes |
+| --- | --- |
+| `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.
+```js
+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" },
+});
+```

build/webgpu/bench.json ADDED Viewed

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

+{
+  "op": "com.microsoft.FusedGemm",
+  "cases": [
+    {
+      "name": "fusedgemm-relu-m3072-k768-n3072",
+      "preset": "smoke",
+      "vars": { "dtype": "float32" },
+      "attrs": { "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [3072, 768], "dtype": "float32", "dist": "normal", "seed": 5300, "scale": 1 },
+        "B": { "shape": [768, 3072], "dtype": "float32", "dist": "normal", "seed": 5301, "scale": 1 },
+        "C": { "shape": [3072], "dtype": "float32", "dist": "normal", "seed": 5302, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [3072, 3072], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 3072 * 768 * 3072" }] }
+    },
+    {
+      "name": "fusedgemm-transb-relu-m3072-k768-n3072",
+      "preset": "smoke",
+      "vars": { "dtype": "float32" },
+      "attrs": { "transB": 1, "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [3072, 768], "dtype": "float32", "dist": "normal", "seed": 5310, "scale": 1 },
+        "B": { "shape": [3072, 768], "dtype": "float32", "dist": "normal", "seed": 5311, "scale": 1 },
+        "C": { "shape": [3072], "dtype": "float32", "dist": "normal", "seed": 5312, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [3072, 3072], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 3072 * 768 * 3072" }] }
+    },
+    {
+      "name": "fusedgemm-f16-relu-m3072-k768-n3072",
+      "preset": "model",
+      "vars": { "dtype": "float16" },
+      "attrs": { "activation": "Relu" },
+      "inputs": {
+        "A": { "shape": [3072, 768], "dtype": "float16", "dist": "normal", "seed": 5320, "scale": 1 },
+        "B": { "shape": [768, 3072], "dtype": "float16", "dist": "normal", "seed": 5321, "scale": 1 },
+        "C": { "shape": [3072], "dtype": "float16", "dist": "normal", "seed": 5322, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [3072, 3072], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 3072 * 768 * 3072" }] }
+    },
+    {
+      "name": "fusedgemm-nobias-tanh-m2048-k1024-n2048",
+      "preset": "model",
+      "vars": { "dtype": "float32" },
+      "attrs": { "activation": "Tanh" },
+      "inputs": {
+        "A": { "shape": [2048, 1024], "dtype": "float32", "dist": "normal", "seed": 5330, "scale": 1 },
+        "B": { "shape": [1024, 2048], "dtype": "float32", "dist": "normal", "seed": 5331, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [2048, 2048], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 2048 * 1024 * 2048" }] }
+    },
+    {
+      "name": "fusedgemm-noactivation-m2048-k1024-n2048",
+      "preset": "model",
+      "vars": { "dtype": "float32" },
+      "attrs": {},
+      "inputs": {
+        "A": { "shape": [2048, 1024], "dtype": "float32", "dist": "normal", "seed": 5340, "scale": 1 },
+        "B": { "shape": [1024, 2048], "dtype": "float32", "dist": "normal", "seed": 5341, "scale": 1 },
+        "C": { "shape": [2048], "dtype": "float32", "dist": "normal", "seed": 5342, "scale": 1 }
+      },
+      "outputs": { "Y": { "shape": [2048, 2048], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "2 * 2048 * 1024 * 2048" }] }
+    }
+  ]
+}

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

	@@ -0,0 +1,186 @@

+{% if usesF16 is defined and usesF16 %}enable f16;
+{% endif %}{% set OUT = "f16" if (usesF16 is defined and usesF16) else "f32" %}
+{{ env.wgsl.resourceDeclarations }}
+// Shape-driven, register-blocked f32 GEMM for all four transpose combinations.
+// Runtime M/N/K share one compiled pipeline. A 4x4 thread micro-tile forms a
+// 64x64 output tile with BK=16; transpose modes change only load strides and
+// keep the stored-contiguous axis coalesced. M/N/K tails are bounds-checked.
+// Both tiles are indexed by their own output axis and group four K values per
+// vector word, so the micro-tile accumulates through dot() and one step reads
+// TM + TN words instead of 4 * (TM + TN) scalars.
+{% set gemmEpi = gemmEpilogue if gemmEpilogue is defined else "none" %}
+{% if gemmEpi == "fastgelu" %}
+fn tanh_safe(x: f32) -> f32 {
+  if (x > 10.0) { return 1.0; }
+  if (x < -10.0) { return -1.0; }
+  return tanh(x);
+}
+fn gelu_tanh(v: f32) -> f32 {
+  return 0.5 * v * (1.0 + tanh_safe(0.7978845608028654 * (v + 0.044715 * v * v * v)));
+}
+{% elif gemmEpi == "activation" %}
+{% 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" %}
+  return tanh(v);
+{% endif %}
+{% if actMode == "HardSigmoid" %}
+  return clamp(f32({{ actA }}) * v + f32({{ actB }}), 0.0, 1.0);
+{% endif %}
+}
+{% endif %}
+{% set epiOpen = "gelu_tanh(" if gemmEpi == "fastgelu" else ("fused_act(" if gemmEpi == "activation" else "") %}
+{% set epiClose = ")" if gemmEpi != "none" else "" %}
+const BK: u32 = 16u;
+const BM: u32 = 64u;
+const BN: u32 = 64u;
+const TM: u32 = 4u;   // per-thread micro-tile rows
+const TN: u32 = 4u;   // per-thread micro-tile cols
+// Stage tiles at the operands' storage width. For f16 inputs this is
+// bit-identical to widening on shared load, keeps f32 FMA registers, and halves
+// workgroup storage.
+{% set tileT = OUT %}
+const K_VECS: u32 = BK / 4u;
+var<workgroup> tileA: array<array<vec4<{{ tileT }}>, K_VECS>, BM>;   // [64][4]  logical A'[m][k/4]
+var<workgroup> tileB: array<array<vec4<{{ tileT }}>, K_VECS>, BN>;   // [64][4]  logical B'[n][k/4]
+@compute @workgroup_size(16, 16, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  let li = lid.y * 16u + lid.x;
+  let M = params.M;
+  let N = params.N;
+  let K = params.K;
+  let n_tiles = (N + BN - 1u) / BN;
+  let m_tiles = (M + BM - 1u) / BM;
+  for (var nt: u32 = wg.x; nt < n_tiles; nt = nt + nwg.x) {
+    let nBase = nt * BN;
+    for (var mt: u32 = wg.y; mt < m_tiles; mt = mt + nwg.y) {
+      let mBase = mt * BM;
+      var acc: array<f32, TM * TN>;
+      for (var i: u32 = 0u; i < TM * TN; i = i + 1u) { acc[i] = 0.0; }
+      let numTiles = (K + BK - 1u) / BK;
+      for (var kt: u32 = 0u; kt < numTiles; kt = kt + 1u) {
+        let kBase = kt * BK;
+        // Cooperative load: one vector word per lane per pass. The lane grid is
+        // laid out so the axis the operand stores contiguously stays coalesced.
+        for (var idx: u32 = li; idx < BM * K_VECS; idx = idx + 256u) {
+{% if transA %}
+          // A stored [K, M]: logical A'[m,k] = a[k*M + m]. Coalesce on m (contiguous).
+          let ar = idx % BM;
+          let ac4 = idx / BM;
+{% else %}
+          // A stored [M, K]: logical A[m,k] = a[m*K + k]. Coalesce on k (contiguous).
+          let ar = idx / K_VECS;
+          let ac4 = idx % K_VECS;
+{% endif %}
+          let am = mBase + ar;
+          let ak = kBase + ac4 * 4u;
+          var aWord = vec4<{{ tileT }}>({{ tileT }}(0.0));
+          if (am < M) {
+{% for component in range(4) %}
+            if (ak + {{ component }}u < K) {
+{% if transA %}
+              aWord[{{ component }}u] = {{ tileT }}(a[(ak + {{ component }}u) * M + am]);
+{% else %}
+              aWord[{{ component }}u] = {{ tileT }}(a[am * K + ak + {{ component }}u]);
+{% endif %}
+            }
+{% endfor %}
+          }
+          tileA[ar][ac4] = aWord;
+        }
+        for (var idx: u32 = li; idx < BN * K_VECS; idx = idx + 256u) {
+{% if transB is defined and transB %}
+          // B stored [N, K]: logical B[k,n] = b[n*K + k]. Coalesce on k (contiguous).
+          let bc = idx / K_VECS;
+          let br4 = idx % K_VECS;
+{% else %}
+          // B stored [K, N]: logical B[k,n] = b[k*N + n]. Coalesce on n (contiguous).
+          let bc = idx % BN;
+          let br4 = idx / BN;
+{% endif %}
+          let bn = nBase + bc;
+          let bk = kBase + br4 * 4u;
+          var bWord = vec4<{{ tileT }}>({{ tileT }}(0.0));
+          if (bn < N) {
+{% for component in range(4) %}
+            if (bk + {{ component }}u < K) {
+{% if transB is defined and transB %}
+              bWord[{{ component }}u] = {{ tileT }}(b[bn * K + bk + {{ component }}u]);
+{% else %}
+              bWord[{{ component }}u] = {{ tileT }}(b[(bk + {{ component }}u) * N + bn]);
+{% endif %}
+            }
+{% endfor %}
+          }
+          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();
+      }
+      let m0 = mBase + lid.y * TM;
+      let n0 = nBase + lid.x * TN;
+      for (var ti: u32 = 0u; ti < TM; ti = ti + 1u) {
+        let m = m0 + ti;
+        if (m >= M) { continue; }
+        for (var tj: u32 = 0u; tj < TN; tj = tj + 1u) {
+          let n = n0 + tj;
+          if (n < N) {
+            // alpha == 0 nulls A*B exactly and avoids Inf*0 becoming NaN.
+            let product = select(acc[ti * TN + tj] * params.alpha, 0.0, params.alpha == 0.0);
+{% if hasBias is defined and hasBias %}
+{% if rowBias is defined and rowBias %}
+            let biasValue = f32(bias[n]);
+{% elif scalarBias is defined and scalarBias %}
+            let biasValue = f32(bias[0u]);
+{% elif columnBias is defined and columnBias %}
+            let biasValue = f32(bias[m]);
+{% else %}
+            let biasValue = f32(bias[m * N + n]);
+{% endif %}
+            let biasTerm = select(biasValue * params.beta, 0.0, params.beta == 0.0);
+            y[m * N + n] = {{ OUT }}({{ epiOpen }}product + biasTerm{{ epiClose }});
+{% else %}
+            y[m * N + n] = {{ OUT }}({{ epiOpen }}product{{ epiClose }});
+{% endif %}
+          }
+        }
+      }
+    }
+  }
+}

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

	@@ -0,0 +1,269 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+enable subgroups;
+{% if pinSubgroupSize32 %}
+enable subgroup_size_control;
+{% endif %}
+enable chromium_experimental_subgroup_matrix;
+diagnostic(off, chromium.subgroup_matrix_uniformity);
+{{ env.wgsl.resourceDeclarations }}
+{% set operandScalar = mmaScalar %}
+{% set accScalar = "f32" %}
+{% set tileMValue = tileM if tileM is defined else 32 %}
+{% set tileNValue = tileN if tileN is defined else 64 %}
+{% set tileKValue = tileK if tileK 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 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; }
+  if (x < -10.0) { return -1.0; }
+  return tanh(x);
+}
+fn gelu_tanh(v: f32) -> f32 {
+  return 0.5 * v * (1.0 + tanh_safe(0.7978845608028654 * (v + 0.044715 * v * v * v)));
+}
+{% elif gemmEpi == "activation" %}
+{% 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" %}
+  return tanh(v);
+{% endif %}
+{% if actMode == "HardSigmoid" %}
+  return clamp(f32({{ actA }}) * v + f32({{ actB }}), 0.0, 1.0);
+{% endif %}
+}
+{% endif %}
+const IN_F:       u32 = {{ inFeatures }}u;
+const OUT_F:      u32 = {{ outFeatures }}u;
+{% set kLoopEndValue = kLoopEnd if kLoopEnd is defined else ("K_LOOP" if tailSafe is defined and tailSafe else "IN_F") %}
+const TILE_COLS:  u32 = {{ tileNValue }}u;
+const TILE_ROWS:  u32 = {{ tileMValue }}u;
+const TILE_K:     u32 = {{ tileKValue }}u;
+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
+  // contiguous as [K,tileM], then transpose in subgroupMatrixLoad.
+  let local_idx = row * {{ (tileKValue / 8)|int }}u + c_idx;
+  let vectors_per_k = TILE_ROWS / 4u;
+  let total_vectors = TILE_K * vectors_per_k;
+  for (var vector_idx = local_idx; vector_idx < total_vectors; vector_idx += 128u) {
+    let k_local = vector_idx / vectors_per_k;
+    let m4 = (vector_idx - k_local * vectors_per_k) * 4u;
+    for (var component = 0u; component < 4u; component++) {
+      let m_local = m4 + component;
+      let global_m = tile_base + m_local;
+      let global_k = k_idx + k_local;
+      tile_A[k_local * TILE_ROWS + m_local] = x[global_k * params.M + global_m];
+    }
+  }
+{% else %}
+  // Clamp the partial-M tail to the last valid row (M-1) instead of zero-filling
+  // it. Some subgroup-matrix implementations let a zero-padded left tile corrupt
+  // the adjacent valid row. Replicating row M-1 avoids that boundary; the
+  // replicated rows are discarded by storeOutput's row_limit guard, so every valid
+  // output row stays bit-identical. M >= 1 is guaranteed by the when-clause args.M > 0.
+  let col: u32 = c_idx * 8u;
+  for (var row_offset: u32 = 0u; row_offset < TILE_ROWS; row_offset += {{ (1024 / tileKValue)|int }}u) {
+  let r: u32 = row + row_offset;
+  let a_global: u32 = min(tile_base + r, params.M - 1u);
+  for (var col_offset: u32 = 0u; col_offset < 8u; col_offset++) {
+    let k: u32 = k_idx + col + col_offset;
+    tile_A[r * TILE_K + col + col_offset] =
+      x[a_global * IN_F + k];
+  }
+  }
+{% 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
+  // the subgroup right operand without a transpose.
+  let local_idx = row * {{ (tileKValue / 16)|int }}u + c_idx;
+  let vectors_per_k = TILE_COLS / 4u;
+  let total_vectors = TILE_K * vectors_per_k;
+  for (var vector_idx = local_idx; vector_idx < total_vectors; vector_idx += 128u) {
+    let k_local = vector_idx / vectors_per_k;
+    let n4 = (vector_idx - k_local * vectors_per_k) * 4u;
+    for (var component = 0u; component < 4u; component++) {
+      let n_local = n4 + component;
+      let global_n = tile_base + n_local;
+      let global_k = k_idx + k_local;
+      tile_B[bank_offset + k_local * TILE_COLS + n_local] = w[global_k * OUT_F + global_n];
+    }
+  }
+{% else %}
+  let col: u32 = c_idx * 16u;
+  for (var row_offset: u32 = 0u; row_offset < TILE_COLS; row_offset += {{ (2048 / tileKValue)|int }}u) {
+  let b_row: u32 = row + row_offset;
+  // Non-power-of-two tuning widths (48/96) use only a prefix of lanes on the
+  // final row band. Keep those inactive lanes out of tile_B and global memory.
+  if (b_row < TILE_COLS) {
+    let w_global: u32 = tile_base + b_row;
+  for (var i: u32 = 0u; i < 16u; i++) {
+    let k: u32 = k_idx + col + i;
+    tile_B[bank_offset + b_row * TILE_K + col + i] =
+      w[w_global * IN_F + k];
+  }
+  }
+  }
+{% endif %}
+}
+{% 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;
+    y[offset + row * OUT_F + col]  = {{ yScalar }}({{ epiOpen }}scratch[src_slot][row * 8u + col]{{ BIAS_1 }}{{ epiClose }});
+    y[offset + row * OUT_F + col2] = {{ yScalar }}({{ epiOpen }}scratch[src_slot][row * 8u + col2]{{ BIAS_2 }}{{ epiClose }});
+  }
+}
+@compute @workgroup_size(128, 1, 1){{ " @subgroup_size(32)" if pinSubgroupSize32 else "" }}
+fn main(
+  @builtin(workgroup_id) workgroup_id: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>,
+  @builtin(local_invocation_index) local_idx: u32,
+  @builtin(subgroup_invocation_id) sg_id: u32,
+  @builtin(subgroup_size) sg_size: u32
+) {
+  // 2D-folded M-tile (row) base: workgroup_id.z carries the high bits when the
+  // row tile exceeds the per-dimension dispatch limit.
+  // Reduces to workgroup_id.y when nwg.z == 1; the row_limit guard in storeOutput
+  // (i32(M) - i32(a_global_base + ...)) drops the over-dispatched tail.
+  let a_global_base: u32 = (ROW_TILE_BASE + workgroup_id.y + workgroup_id.z * nwg.y) * TILE_ROWS;
+  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>>(&tile_A, matrix_a_offset + {{ m * 8 }}u, true, TILE_ROWS);
+{% else %}
+      var matA{{ m }}: subgroup_matrix_left<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_left<{{ operandScalar }}, 8, 8>>(&tile_A, matrix_a_offset{% if m > 0 %} + {{ m * 8 }}u * TILE_K{% endif %}, false, TILE_K);
+{% endif %}
+{% 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 %}
+      var matB{{ n }}: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>>(&tile_B, matrix_b_offset + {{ n * 8 }}u, false, TILE_COLS);
+{% else %}
+      var matB{{ n }}: subgroup_matrix_right<{{ operandScalar }}, 8, 8> = subgroupMatrixLoad<subgroup_matrix_right<{{ operandScalar }}, 8, 8>>(&tile_B, matrix_b_offset + {{ n * 8 }}u * TILE_K, true, TILE_K);
+{% endif %}
+{% endfor %}
+{% 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
+  // retain this guarded scalar epilogue: subgroupMatrixStore scatters values
+  // 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(&scratch[bank + {{ m * bTilesPerSubgroup + n }}u], 0u, matC{{ m }}{{ n }}, false, 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) %}
+  storeOutput(
+    matrix_c_offset + {{ m * 8 }}u * OUT_F + {{ n * 8 }}u{% if needsColBase %},
+    w_global_base + base_B + {{ n * 8 }}u,
+{% else %},
+{% endif %}
+    row,
+    col,
+    bank + {{ m * bTilesPerSubgroup + n }}u,
+    row_limit_{{ m }}
+  );
+{% endfor %}
+{% endfor %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,594 @@

+{
+  "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",
+    "canPinSubgroupSize32": "device.features.has(\"subgroups\") and device.features.has(\"subgroup-size-control\") and has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize <= 32 and device.adapterInfo.subgroupMaxSize >= 32",
+    "pinSubgroupSize32": "canPinSubgroupSize32 and not wave32Adapter",
+    "wave32Effective": "wave32Adapter or pinSubgroupSize32",
+    "transAFlag": "has(attrs, \"transA\") and attrs.transA != 0",
+    "transBFlag": "has(attrs, \"transB\") and attrs.transB != 0",
+    "gemmM": "dim(shapes.A, 1) if transAFlag else dim(shapes.A, 0)",
+    "gemmK": "dim(shapes.A, 0) if transAFlag else dim(shapes.A, 1)",
+    "gemmN": "dim(shapes.B, 0) if transBFlag else dim(shapes.B, 1)",
+    "gemmKB": "dim(shapes.B, 1) if transBFlag else dim(shapes.B, 0)",
+    "shapeOk": "ranks.A == 2 and ranks.B == 2 and ranks.Y == 2 and gemmK == gemmKB and dim(shapes.Y, 0) == gemmM and dim(shapes.Y, 1) == gemmN",
+    "dtypeOk": "(tensorDtypes.A == \"float32\" or tensorDtypes.A == \"float16\") and tensorDtypes.B == tensorDtypes.A and tensorDtypes.Y == tensorDtypes.A and f16Ok(dtypes.T)",
+    "rowBiasOk": "present.C and ((ranks.C == 1 and dim(shapes.C, 0) == gemmN) or (ranks.C == 2 and dim(shapes.C, 0) == 1 and dim(shapes.C, 1) == gemmN))",
+    "scalarBiasOk": "present.C and (ranks.C == 0 or (ranks.C == 1 and dim(shapes.C, 0) == 1) or (ranks.C == 2 and dim(shapes.C, 0) == 1 and dim(shapes.C, 1) == 1))",
+    "columnBiasOk": "present.C and ranks.C == 2 and dim(shapes.C, 0) == gemmM and dim(shapes.C, 1) == 1",
+    "matrixBiasOk": "present.C and ranks.C == 2 and dim(shapes.C, 0) == gemmM and dim(shapes.C, 1) == gemmN",
+    "cOk": "tensorDtypes.C == tensorDtypes.A and (rowBiasOk or scalarBiasOk or columnBiasOk or matrixBiasOk)",
+    "hasActivation": "has(attrs, \"activation\")",
+    "activationName": "attrs.activation if hasActivation else \"\"",
+    "activationOk": "not hasActivation or activationName == \"Relu\" or activationName == \"LeakyRelu\" or activationName == \"Sigmoid\" or activationName == \"Tanh\" or activationName == \"HardSigmoid\"",
+    "activationAlpha": "attrs.activation_alpha if has(attrs, \"activation_alpha\") else (0.2 if activationName == \"HardSigmoid\" else (0.01 if activationName == \"LeakyRelu\" else 0.0))",
+    "activationBeta": "attrs.activation_beta if has(attrs, \"activation_beta\") else (0.5 if activationName == \"HardSigmoid\" else 0.0)",
+    "baseContract": "shapeOk and dtypeOk and activationOk",
+    "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": [
+          { "componentType": "f16", "M": 8, "N": 8, "K": 8 },
+          { "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": [
+          { "componentType": "f16", "M": 8, "N": 8, "K": 8 },
+          { "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": [
+          { "componentType": "f16", "M": 8, "N": 8, "K": 8 },
+          { "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": [
+          { "componentType": "f16", "M": 8, "N": 8, "K": 8 },
+          { "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": [
+          { "componentType": "f16", "M": 8, "N": 8, "K": 8 },
+          { "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": [
+          { "componentType": "f16", "M": 8, "N": 8, "K": 8 },
+          { "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": [
+          { "componentType": "f16", "M": 8, "N": 8, "K": 8 },
+          { "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": [
+          { "componentType": "f16", "M": 8, "N": 8, "K": 8 },
+          { "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)" }
+        }
+      ]
+    },
+    {
+      "id": "notrans_bias",
+      "priority": 0,
+      "when": ["biasContract", "transAFlag == false", "transBFlag == false"],
+      "constants": {
+        "transA": false,
+        "transB": false,
+        "hasBias": true,
+        "rowBias": "rowBiasOk",
+        "scalarBias": "scalarBiasOk",
+        "columnBias": "columnBiasOk"
+      },
+      "passes": [
+        {
+          "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)"
+          }
+        }
+      ]
+    },
+    {
+      "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)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "transA_bias",
+      "priority": 0,
+      "when": ["biasContract", "transAFlag == true", "transBFlag == false"],
+      "constants": {
+        "transA": true,
+        "transB": false,
+        "hasBias": true,
+        "rowBias": "rowBiasOk",
+        "scalarBias": "scalarBiasOk",
+        "columnBias": "columnBiasOk"
+      },
+      "passes": [
+        {
+          "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)"
+          }
+        }
+      ]
+    },
+    {
+      "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)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "transB_bias",
+      "priority": 0,
+      "when": ["biasContract", "transAFlag == false", "transBFlag == true"],
+      "constants": {
+        "transA": false,
+        "transB": true,
+        "hasBias": true,
+        "rowBias": "rowBiasOk",
+        "scalarBias": "scalarBiasOk",
+        "columnBias": "columnBiasOk"
+      },
+      "passes": [
+        {
+          "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)"
+          }
+        }
+      ]
+    },
+    {
+      "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)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "transAB_bias",
+      "priority": 0,
+      "when": ["biasContract", "transAFlag == true", "transBFlag == true"],
+      "constants": {
+        "transA": true,
+        "transB": true,
+        "hasBias": true,
+        "rowBias": "rowBiasOk",
+        "scalarBias": "scalarBiasOk",
+        "columnBias": "columnBiasOk"
+      },
+      "passes": [
+        {
+          "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)"
+          }
+        }
+      ]
+    },
+    {
+      "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)"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,19 @@

+{
+  "name": "com.microsoft.FusedGemm",
+  "id": "_com_microsoft_fusedgemm_webgpu_afa4467",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "5+Pgon1ED+99xPL6smxxz2G47C5gDnfB0kJ+Zj2wZLc=",
+      "fused-gemm.wgsl.jinja": "IkB53JqVNn9Xezw0YcKYDvm9IwDvoZNnSeA/UY4Eci4=",
+      "gemm-subgroup-matrix.wgsl.jinja": "gk0C1ZhDuTerG8rDtHYjQp74zieMlGP6QeFm0XqRZSw=",
+      "manifest.json": "s7cSq5so+VQS0XTbiHalmSHGnQQA2ylTu524o3V815U=",
+      "test.json": "xyQ7rg3TZtQ7zZA5W25GP1jcfMvCLIWiBlk6ASHNmi0="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.FusedGemm" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,662 @@

+{
+  "op": "com.microsoft.FusedGemm",
+  "cases": [
+    {
+      "name": "notrans_rowbias_relu",
+      "provenance": { "notes": "No transpose, a row bias, and Relu." },
+      "attrs": { "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [96, 80],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.18000000000000002, "cosStep": 0.32 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [80, 72],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.12, "cosStep": 0.24000000000000002 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [72],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.3, "cosStep": 0.14 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [96, 72], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "transa_rowbias_relu",
+      "provenance": { "notes": "transA: A stored (K, M), which changes only the load index." },
+      "attrs": { "transA": 1, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [80, 96],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.19, "cosStep": 0.33 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [80, 72],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.13, "cosStep": 0.25 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [72],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.31, "cosStep": 0.15 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [96, 72], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "transb_rowbias_leaky",
+      "provenance": { "notes": "transB with LeakyRelu at its default slope of 0.01." },
+      "attrs": { "transB": 1, "activation": "LeakyRelu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [96, 80],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.2, "cosStep": 0.33999999999999997 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [72, 80],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.14, "cosStep": 0.26 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [72],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.31999999999999995, "cosStep": 0.16 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [96, 72], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "transab_rowbias_tanh",
+      "provenance": { "notes": "Both operands transposed, with Tanh." },
+      "attrs": { "transA": 1, "transB": 1, "activation": "Tanh" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [80, 96],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.21000000000000002, "cosStep": 0.35 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [72, 80],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.15, "cosStep": 0.27 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [72],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.32999999999999996, "cosStep": 0.17 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [96, 72], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "notrans_nobias_sigmoid",
+      "provenance": { "notes": "No bias at all -- the other binding set -- with Sigmoid." },
+      "attrs": { "activation": "Sigmoid" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [96, 80],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.22000000000000003, "cosStep": 0.36 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [80, 72],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.16, "cosStep": 0.28 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [96, 72], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "transa_nobias_none",
+      "provenance": { "notes": "transA with no bias and no activation, which is plain Gemm." },
+      "attrs": { "transA": 1 },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [48, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.23, "cosStep": 0.37 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [48, 40],
+          "data": {
+            "kind": "fillFloat32",
+            "scale": 0.6,
+            "sinStep": 0.16999999999999998,
+            "cosStep": 0.29000000000000004
+          }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 40], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "transab_nobias_hardsigmoid",
+      "provenance": { "notes": "Both transposed, HardSigmoid at its 0.2/0.5 defaults." },
+      "attrs": { "transA": 1, "transB": 1, "activation": "HardSigmoid" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [48, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.25, "cosStep": 0.39 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [40, 48],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.19, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 40], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "scalar_bias_relu",
+      "provenance": { "notes": "A scalar C, broadcast to every element." },
+      "attrs": { "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [48, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.26, "cosStep": 0.4 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [32, 24],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.2, "cosStep": 0.32 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [1],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.38, "cosStep": 0.22 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [48, 24], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "column_bias_relu",
+      "provenance": { "notes": "A column C of shape (M, 1)." },
+      "attrs": { "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [48, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.27, "cosStep": 0.41000000000000003 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [32, 24],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.21000000000000002, "cosStep": 0.33 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [48, 1],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.39, "cosStep": 0.23 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [48, 24], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "matrix_bias_relu",
+      "provenance": { "notes": "A full (M, N) C." },
+      "attrs": { "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [48, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.28, "cosStep": 0.42 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [32, 24],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.22, "cosStep": 0.34 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [48, 24],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.39999999999999997, "cosStep": 0.24 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [48, 24], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "row2d_bias_tanh",
+      "provenance": { "notes": "A (1, N) C, the rank-2 spelling of a row bias." },
+      "attrs": { "activation": "Tanh" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [48, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.29000000000000004, "cosStep": 0.43 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [32, 24],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.22999999999999998, "cosStep": 0.35 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [1, 24],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.41, "cosStep": 0.25 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [48, 24], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "alpha_beta_leaky_explicit",
+      "provenance": { "notes": "Explicit alpha, beta and LeakyRelu slope, none of them defaults." },
+      "attrs": { "alpha": 0.5, "beta": 2, "activation": "LeakyRelu", "activation_alpha": 0.25 },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [48, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.30000000000000004, "cosStep": 0.44 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [32, 24],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.24, "cosStep": 0.36 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [24],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.42, "cosStep": 0.26 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [48, 24], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "alpha_zero_relu",
+      "provenance": {
+        "notes": "alpha = 0 nulls the product exactly rather than multiplying it, so an infinite accumulator cannot become NaN."
+      },
+      "attrs": { "alpha": 0, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [32, 24],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.32, "cosStep": 0.45999999999999996 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [24, 16],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.26, "cosStep": 0.38 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.43999999999999995, "cosStep": 0.28 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [32, 16], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "beta_zero_relu",
+      "provenance": { "notes": "beta = 0 nulls the C term the same way." },
+      "attrs": { "beta": 0, "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [32, 24],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.33, "cosStep": 0.47 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [24, 16],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.27, "cosStep": 0.39 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": {
+            "kind": "fillFloat32",
+            "scale": 0.4,
+            "sinStep": 0.44999999999999996,
+            "cosStep": 0.29000000000000004
+          }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [32, 16], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "small_unaligned_relu",
+      "provenance": { "notes": "M, K and N all below one tile and none a multiple of the micro-tile." },
+      "attrs": { "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [7, 13],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.34, "cosStep": 0.48 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [13, 11],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.28, "cosStep": 0.4 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [11],
+          "data": {
+            "kind": "fillFloat32",
+            "scale": 0.4,
+            "sinStep": 0.45999999999999996,
+            "cosStep": 0.30000000000000004
+          }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [7, 11], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "f16_notrans_bias_relu",
+      "provenance": { "notes": "float16 operands, accumulating in f32." },
+      "attrs": { "activation": "Relu" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [64, 48],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.35, "cosStep": 0.49 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [48, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.41000000000000003 }
+        },
+        "C": {
+          "dtype": "float16",
+          "shape": [32],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.47, "cosStep": 0.31 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [64, 32], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "f16_transb_nobias_tanh",
+      "provenance": { "notes": "float16 with transB and no bias." },
+      "attrs": { "transB": 1, "activation": "Tanh" },
+      "inputs": {
+        "A": {
+          "dtype": "float16",
+          "shape": [64, 48],
+          "data": { "kind": "fillFloat32", "scale": 0.9, "sinStep": 0.36, "cosStep": 0.5 }
+        },
+        "B": {
+          "dtype": "float16",
+          "shape": [32, 48],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.3, "cosStep": 0.42000000000000004 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [64, 32], "tolerance": 0.006, "relTolerance": 0.006 } }
+    },
+    {
+      "name": "pinned_transab_leaky_alpha_beta",
+      "provenance": {
+        "notes": "Hand-computed at float64 from the schema's own definition -- alpha * A' * B' + beta * C, then LeakyRelu -- with both operands transposed and every scalar explicit, so nothing in the pin comes from a default."
+      },
+      "attrs": {
+        "transA": 1,
+        "transB": 1,
+        "alpha": 0.75,
+        "beta": 1.5,
+        "activation": "LeakyRelu",
+        "activation_alpha": 0.125
+      },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [10, 12],
+          "data": {
+            "kind": "values",
+            "values": [0.482, 0.233, -0.17, 1.167, -0.64, 1.418, 0.317, -0.314, 0.628, 0.761, 0.092, -1.224, 1.017, 0.692, -0.239, -1.141, 0.507, 1.027, 0.417, 1.008, -1.377, -1.169, 0.818, 0.125, -0.798, 0.491, -0.397, -0.028, 1.262, 0.159, -1.326, 0.973, -0.57, 0.244, 1.381, -0.037, -1.135, -0.158, -0.767, 0.452, 1.127, -0.526, 1.431, -0.255, -0.505, 0.154, -0.336, 0.414, -0.21, 1.258, -0.143, -0.146, -1.313, 0.676, -0.406, 1.497, -1.321, 1.005, 0.274, -1.266, -1.498, -0.927, 0.328, 0.949, -0.043, 1.29, -0.088, -0.031, 0.975, 0.247, 0.213, -0.993, 0.848, 0.355, 0.745, -0.006, 0.934, -0.819, 0.242, 1.171, -0.907, -1.153, 0.361, -0.612, -0.006, -1.408, 0.302, -0.531, 1.072, 0.002, 1.211, -0.427, 0.901, -0.552, -1.033, 0.857, -1.401, -0.719, -1.259, 0.575, -1.481, 0.209, 0.107, 0.975, 0.802, 1.009, 0.77, 1.493, 1.138, -1.32, 1.026, -0.854, -0.22, 1.076, -1.45, -0.272, 0.24, 1.368, -1.149, 0.598]
+          }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [8, 10],
+          "data": {
+            "kind": "values",
+            "values": [-0.756, 1.16, -0.81, -1.402, -0.32, -0.036, 0.068, -1.181, -1.451, -1.096, -1.089, -0.777, 1.363, 0.017, -1.26, 1.485, 0.309, 1.036, -0.894, 1.06, 0.977, -1.123, -1.369, -1.108, 0.193, 0.923, -0.324, 0.673, 1.35, -1.017, 1.247, -0.487, 0.898, 0.209, 0.78, -1.249, 1.274, -0.422, -1.282, 0.118, 0.51, 0.864, -0.868, 1.06, -1.42, -1.341, -0.608, 0.888, 1.183, 0.199, 1.423, 0.737, -0.285, 0.903, 0.487, -0.826, 0.934, -1.365, -1.034, -1.222, 1.233, -0.428, -0.705, 1.431, -1.459, 0.817, -1.361, 1.45, 0.173, -1.252, -1.415, 1.357, -0.751, 0.894, -0.602, 0.879, 1.197, 0.168, -0.407, 0.107]
+          }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.235, -0.223, 0.474, -0.529, 0.922, 0.252, -0.834, 0.806] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [12, 8],
+          "data": {
+            "kind": "values",
+            "values": [3.370844, -0.198133, -0.199107, 2.110806, 2.109943, 2.351428, -0.592314, 1.807759, 3.546125, -0.528558, -0.131459, 2.441212, -0.108338, 5.288458, -0.547352, 0.359243, 1.611335, 2.095594, -0.005854, 0.27388, -0.058539, -0.047104, -0.393263, 2.309827, -0.150555, -0.116139, 3.783087, -0.084227, 0.70041, 1.53985, 1.740793, -0.062086, 0.411694, 4.257617, -0.456835, 0.679208, 1.923927, 0.92273, 0.173925, 4.421494, -0.056756, -0.076294, 1.740675, -0.189957, 0.772669, -0.025928, -0.126539, 0.139902, -0.006079, -0.264435, 2.411458, -0.232052, 4.870933, 2.281323, 3.98095, 3.349453, 0.326566, -0.256703, -0.033174, 0.35729, -0.059635, 1.535964, -0.698491, 1.830789, -0.243493, 1.843096, 5.123723, -0.469259, 2.88782, -0.427406, 3.358449, -0.070443, -0.430807, -0.122892, 2.170899, -0.110731, 0.818956, -0.183576, -0.207642, -0.302042, 1.494748, -0.206163, 0.184732, -0.042274, -0.008104, 2.140054, -0.375211, 0.824137, -0.229193, 0.648641, 0.168956, -0.499118, 5.959787, -0.428765, 0.006611, 2.00139]
+          },
+          "tolerance": 0.00002,
+          "relTolerance": 0.0002
+        }
+      }
+    },
+    {
+      "name": "pinned_column_bias_hardsigmoid",
+      "provenance": { "notes": "Independently pinned column-bias broadcast with HardSigmoid at explicit parameters." },
+      "attrs": { "activation": "HardSigmoid", "activation_alpha": 0.3, "activation_beta": 0.4 },
+      "inputs": {
+        "A": {
+          "dtype": "float32",
+          "shape": [9, 7],
+          "data": {
+            "kind": "values",
+            "values": [-1.053, 0.354, -0.86, -0.3, 0.047, -0.117, -0.287, -0.319, -1.275, -0.229, 1.483, 0.499, 0.027, -0.762, -0.484, 1.402, -0.067, -1.282, 1.004, -0.316, 0.827, -0.547, 1.007, 0.853, 1.198, 0.01, -0.728, -0.59, 0.576, -0.694, 0.916, -1.297, -1.274, -0.773, -0.042, -0.708, -0.676, -0.463, -0.763, -0.803, 0.241, -0.813, 1.473, 0.728, 0.481, 0.679, -0.263, 0.948, 0.487, -0.141, 0.651, -1.067, 0.751, 1.368, 0.358, 0.552, 0.481, -1.401, 1.166, -0.338, -0.392, 0.853, 1.361]
+          }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [7, 6],
+          "data": {
+            "kind": "values",
+            "values": [-0.006, -1.252, 0.157, -1.417, -1.042, 1.004, -0.995, 0.459, -1.464, 0.63, -0.926, -1.197, -0.321, -0.703, 0.213, 1.042, -0.752, -1.147, -0.595, -0.542, 1.233, 0.919, -1.13, 0.523, -1.206, -1.173, -0.179, 0.925, 0.323, -1.273, 0.567, 0.742, 0.435, 0.361, -1.469, -1.195, -1.211, -0.674, 0.828, 0.501, 0.515, -0.094]
+          }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [9, 1],
+          "data": { "kind": "values", "values": [-0.957, 0.099, -0.259, -0.825, -0.224, -0.876, 0.837, -0.309, -0.723] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [9, 6],
+          "data": {
+            "kind": "values",
+            "values": [0.212855, 0.802754, 0.0, 0.233126, 0.651217, 0.0, 0.669088, 0.165598, 1.0, 0.688478, 0.351265, 0.924192, 0.0, 0.328861, 0.0, 0.78732, 0.898156, 0.0, 0.0, 0.075651, 0.0, 1.0, 0.0, 0.0, 1.0, 0.323342, 0.200611, 0.0, 0.78929, 1.0, 1.0, 1.0, 0.0, 0.0, 0.599796, 0.449499, 0.3431, 0.191342, 0.941447, 0.602963, 0.0, 0.521214, 0.0, 0.039497, 0.334728, 0.865142, 0.216728, 0.0, 0.340949, 0.0, 1.0, 0.173301, 0.069842, 0.182348]
+          },
+          "tolerance": 0.00002,
+          "relTolerance": 0.0002
+        }
+      }
+    },
+    {
+      "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": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.8, "sinStep": 0.11, "cosStep": 0.27 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [64, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.19, "cosStep": 0.13 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.35, "sinStep": 0.23, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "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": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.8, "sinStep": 0.15, "cosStep": 0.27 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [64, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.19, "cosStep": 0.17 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "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": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.8, "sinStep": 0.12, "cosStep": 0.27 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [64, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.19, "cosStep": 0.14 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.35, "sinStep": 0.23, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "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": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.8, "sinStep": 0.16, "cosStep": 0.27 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [64, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.19, "cosStep": 0.18 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "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": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.8, "sinStep": 0.13, "cosStep": 0.27 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.19, "cosStep": 0.15 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.35, "sinStep": 0.23, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "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": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.8, "sinStep": 0.16999999999999998, "cosStep": 0.27 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.19, "cosStep": 0.19 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "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": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.8, "sinStep": 0.14, "cosStep": 0.27 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.19, "cosStep": 0.16 }
+        },
+        "C": {
+          "dtype": "float32",
+          "shape": [128],
+          "data": { "kind": "fillFloat32", "scale": 0.35, "sinStep": 0.23, "cosStep": 0.41 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    },
+    {
+      "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": {
+        "A": {
+          "dtype": "float32",
+          "shape": [64, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.8, "sinStep": 0.18, "cosStep": 0.27 }
+        },
+        "B": {
+          "dtype": "float32",
+          "shape": [128, 64],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.19, "cosStep": 0.2 }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [64, 128], "tolerance": 0.00002, "relTolerance": 0.0002 } }
+    }
+  ]
+}