sync 91d990483a17

Browse files

Files changed (12) hide show

README.md +41 -15
build/webgpu/bench.json +646 -7
build/webgpu/manifest.json +0 -0
build/webgpu/metadata.json +65 -12
build/webgpu/moe-ffn-gemv.wgsl.jinja +1 -1
build/webgpu/moe-ffn-grouped.wgsl.jinja +4 -9
build/webgpu/moe-ffn-stage.wgsl.jinja +5 -6
build/webgpu/moe-grouped-sgmat.wgsl.jinja +166 -0
build/webgpu/moe-mix-stage.wgsl.jinja +4 -5
build/webgpu/moe-output-stage.wgsl.jinja +4 -5
build/webgpu/moe-route-stage.wgsl.jinja +2 -3
build/webgpu/test.json +695 -24

README.md CHANGED Viewed

@@ -18,22 +18,22 @@ See the [ONNX Runtime `MoE` contrib-operator spec](https://github.com/microsoft/
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `input` | `inputT` | `T` | — | — | Token activations, either 2D `(num_tokens, hidden_size)` or 3D `(batch_size, sequence_length, hidden_size)`. | required |
-| `router_probs` | `routerT` | `T` | `2` | — | 2D router logits of shape `(num_tokens, num_experts)`, where `num_tokens` is the product of every leading dimension of `input`. Despite the historical port name, the operator applies a full softmax before top-k selection. | required |
-| `fc1_experts_weights` | `fc1T` | `T` | `3` | — | 3D first-layer expert weights of shape `(num_experts, fusion_size * inter_size, hidden_size)`, where `fusion_size` is 2 for fused SwiGLU (`swiglu_fusion` 1 or 2) and 1 otherwise. | required |
-| `fc1_experts_bias` | `fc1BiasT` | `T` | `2` | — | Optional 2D FC1 bias of shape `(num_experts, fusion_size * inter_size)`. | optional |
-| `fc2_experts_weights` | `fc2T` | `T` | `3` | — | 3D second-layer expert weights of shape `(num_experts, hidden_size, inter_size)`. | required |
-| `fc2_experts_bias` | `fc2BiasT` | `T` | `2` | — | Optional 2D FC2 bias of shape `(num_experts, hidden_size)`, added per expert before that expert's routing weight is applied. | optional |
-| `fc3_experts_weights` | `fc3T` | `T` | `3` | — | Optional 3D third-layer expert weights of shape `(num_experts, inter_size, hidden_size)`. It supplies the separate linear operand for SwiGLU when `swiglu_fusion` is 0, or the multiplicative linear projection for SiLU gating. Other activations do not consume FC3. | optional |
-| `fc3_experts_bias` | `fc3BiasT` | `T` | `2` | — | Optional 2D FC3 bias of shape `(num_experts, inter_size)`. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `output` | `outputT` | `T` | same as `input` | same as `input` | Routed expert output with the same shape as `input`. | required |
 ## Attributes
@@ -47,8 +47,8 @@ Attributes and default values (overridable per request):
 | `k` | `1` | Number of experts selected per token; the schema default is 1. |
 | `normalize_routing_weights` | `0` | Whether to normalize the selected routing weights; the schema default is 0. |
 | `swiglu_fusion` | `0` | 0 keeps the SwiGLU operands in separate FC1/FC3 GEMMs, 1 interleaves them in one FC1 row, and 2 concatenates them. The schema default is 0. |
-| `use_sparse_mixer` | `0` | Whether to use sparse-mixer routing. The standard default and only supported value is 0. |
 | `swiglu_limit` | — | Optional SwiGLU clamp limit; omission means no clamp. |
 ## Type constraints
@@ -56,9 +56,30 @@ Attributes and default values (overridable per request):
 | --- | --- |
 | `T` | `float32` |
 ## 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
@@ -66,6 +87,7 @@ Attributes and default values (overridable per request):
 - [`moe-ffn-gemv.wgsl.jinja`](build/webgpu/moe-ffn-gemv.wgsl.jinja)
 - [`moe-ffn-grouped.wgsl.jinja`](build/webgpu/moe-ffn-grouped.wgsl.jinja)
 - [`moe-ffn-stage.wgsl.jinja`](build/webgpu/moe-ffn-stage.wgsl.jinja)
 - [`moe-mix-stage.wgsl.jinja`](build/webgpu/moe-mix-stage.wgsl.jinja)
 - [`moe-output-gemv.wgsl.jinja`](build/webgpu/moe-output-gemv.wgsl.jinja)
 - [`moe-output-grouped.wgsl.jinja`](build/webgpu/moe-output-grouped.wgsl.jinja)
@@ -74,10 +96,14 @@ Attributes and default values (overridable per request):
 ## Use with `@huggingface/kernels`
-The loader derives every required output's shape and logical dtype from the manifest contract and this call.
-It then allocates the result tensors automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

 ## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `inputT` | `input` | `T` | — | — | Token activations, either 2D `(num_tokens, hidden_size)` or 3D `(batch_size, sequence_length, hidden_size)`. | required |
+| `routerT` | `router_probs` | `T` | `2` | — | 2D router logits of shape `(num_tokens, num_experts)`, where `num_tokens` is the product of every leading dimension of `input`. A full softmax is applied before top-k selection. | required |
+| `fc1T` | `fc1_experts_weights` | `T` | `3` | — | 3D first-layer expert weights of shape `(num_experts, fusion_size * inter_size, hidden_size)`, where `fusion_size` is 2 for fused SwiGLU (`swiglu_fusion` 1 or 2) and 1 otherwise. | required |
+| `fc1BiasT` | `fc1_experts_bias` | `T` | `2` | — | Optional 2D FC1 bias of shape `(num_experts, fusion_size * inter_size)`. | optional |
+| `fc2T` | `fc2_experts_weights` | `T` | `3` | — | 3D second-layer expert weights of shape `(num_experts, hidden_size, inter_size)`. | required |
+| `fc2BiasT` | `fc2_experts_bias` | `T` | `2` | — | Optional 2D FC2 bias of shape `(num_experts, hidden_size)`, added per expert before that expert's routing weight is applied. | optional |
+| `fc3T` | `fc3_experts_weights` | `T` | `3` | — | Optional 3D third-layer expert weights of shape `(num_experts, inter_size, hidden_size)`. It supplies the separate linear operand for SwiGLU when `swiglu_fusion` is 0, or the multiplicative linear projection for SiLU gating. Other activations do not consume FC3. | optional |
+| `fc3BiasT` | `fc3_experts_bias` | `T` | `2` | — | Optional 2D FC3 bias of shape `(num_experts, inter_size)`. | optional |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `outputT` | `output` | `T` | same as `inputT` | same as `inputT` | Routed expert output with the same shape as `input`. | required |
 ## Attributes
 | `k` | `1` | Number of experts selected per token; the schema default is 1. |
 | `normalize_routing_weights` | `0` | Whether to normalize the selected routing weights; the schema default is 0. |
 | `swiglu_fusion` | `0` | 0 keeps the SwiGLU operands in separate FC1/FC3 GEMMs, 1 interleaves them in one FC1 row, and 2 concatenates them. The schema default is 0. |
 | `swiglu_limit` | — | Optional SwiGLU clamp limit; omission means no clamp. |
+| `use_sparse_mixer` | `0` | Whether to use sparse-mixer routing. The standard default and only supported value is 0. |
 ## Type constraints
 | --- | --- |
 | `T` | `float32` |
+## Implementation variants
+One implementation is selected per call from the device capabilities, the request shapes and the dtypes; these notes say what each one covers.
+- `sgmat_grouped_routed_fc1plain_fc3none_fc2plain` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1plain_fc3none_fc2bias` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1plain_fc3plain_fc2plain` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1plain_fc3plain_fc2bias` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1plain_fc3biased_fc2plain` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1plain_fc3biased_fc2bias` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1bias_fc3none_fc2plain` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1bias_fc3none_fc2bias` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1bias_fc3plain_fc2plain` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1bias_fc3plain_fc2bias` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1bias_fc3biased_fc2plain` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+- `sgmat_grouped_routed_fc1bias_fc3biased_fc2bias` — Expert-grouped f32 matrix projections load public weight layouts directly. Two statically interleaved accumulation chains limit rounding growth; the input tile is reused for result publication. Requires compatible subgroups, f32 fragments, workgroup limits, and complete weight tiles.
+## 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
 - [`moe-ffn-gemv.wgsl.jinja`](build/webgpu/moe-ffn-gemv.wgsl.jinja)
 - [`moe-ffn-grouped.wgsl.jinja`](build/webgpu/moe-ffn-grouped.wgsl.jinja)
 - [`moe-ffn-stage.wgsl.jinja`](build/webgpu/moe-ffn-stage.wgsl.jinja)
+- [`moe-grouped-sgmat.wgsl.jinja`](build/webgpu/moe-grouped-sgmat.wgsl.jinja)
 - [`moe-mix-stage.wgsl.jinja`](build/webgpu/moe-mix-stage.wgsl.jinja)
 - [`moe-output-gemv.wgsl.jinja`](build/webgpu/moe-output-gemv.wgsl.jinja)
 - [`moe-output-grouped.wgsl.jinja`](build/webgpu/moe-output-grouped.wgsl.jinja)
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.MoE",
   "tunableSpace": {
     "workgroupSize": [64, 128],
     "decodeLanes": [32],
@@ -72,7 +71,7 @@
       "outputs": { "outputT": { "dtype": "float32", "shape": [1, 1024] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "args.topK * 3 * args.hidden * args.inter * 4" }] },
       "provenance": {
-        "notes": "Scored as bandwidth, not gflops: a decode step reads whole expert weight matrices to produce one output row, so memory is its roofline and the compute peak is the wrong denominator. The prefill cases stay on gflops."
       }
     },
     {
@@ -105,7 +104,7 @@
       "outputs": { "outputT": { "dtype": "float32", "shape": [1, 1024] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "args.topK * 2 * args.hidden * args.inter * 4" }] },
       "provenance": {
-        "notes": "Scored as bandwidth, not gflops: a decode step reads whole expert weight matrices to produce one output row, so memory is its roofline and the compute peak is the wrong denominator. The prefill cases stay on gflops."
       }
     },
     {
@@ -138,7 +137,7 @@
       "outputs": { "outputT": { "dtype": "float32", "shape": [8, 1024] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "args.topK * 3 * args.hidden * args.inter * 4" }] },
       "provenance": {
-        "notes": "Scored as bandwidth, not gflops: a decode step reads whole expert weight matrices to produce one output row, so memory is its roofline and the compute peak is the wrong denominator. The prefill cases stay on gflops."
       }
     },
     {
@@ -369,7 +368,7 @@
       "name": "moe-f32-qwen3-moe-decode-t1-h2048-e32-k8-i768",
       "preset": "model",
       "provenance": {
-        "notes": "Qwen3-MoE class defaults (hidden_size 2048, moe_intermediate_size 768, num_experts_per_tok 8); expert count capped at 32 of the config's 128 to keep each weight tensor under 512 MB. Scored as bandwidth, not gflops: a decode step reads whole expert weight matrices to produce one output row, so memory is its roofline and the compute peak is the wrong denominator. The prefill cases stay on gflops."
       },
       "vars": { "tokens": 1, "hidden": 2048, "experts": 32, "inter": 768, "topK": 8 },
       "attrs": { "k": 8, "activation_type": "swiglu", "swiglu_fusion": 2, "normalize_routing_weights": 1 },
@@ -405,7 +404,7 @@
       "name": "moe-f32-olmoe-decode-t1-h2048-e16-k8-i2048",
       "preset": "model",
       "provenance": {
-        "notes": "OLMoE class defaults (hidden_size 2048, intermediate_size 2048, num_experts_per_tok 8); expert count capped at 16 of the config's 64. Scored as bandwidth, not gflops: a decode step reads whole expert weight matrices to produce one output row, so memory is its roofline and the compute peak is the wrong denominator. The prefill cases stay on gflops."
       },
       "vars": { "tokens": 1, "hidden": 2048, "experts": 16, "inter": 2048, "topK": 8 },
       "attrs": { "k": 8, "activation_type": "swiglu", "swiglu_fusion": 2, "normalize_routing_weights": 1 },
@@ -439,7 +438,7 @@
       "name": "moe-f32-qwen2-moe-decode-t1-h2048-e20-k4-i1408",
       "preset": "model",
       "provenance": {
-        "notes": "Qwen2-MoE class defaults (hidden_size 2048, moe_intermediate_size 1408, num_experts_per_tok 4); expert count capped at 20 of the config's 60. Scored as bandwidth, not gflops: a decode step reads whole expert weight matrices to produce one output row, so memory is its roofline and the compute peak is the wrong denominator. The prefill cases stay on gflops."
       },
       "vars": { "tokens": 1, "hidden": 2048, "experts": 20, "inter": 1408, "topK": 4 },
       "attrs": { "k": 4, "activation_type": "swiglu", "swiglu_fusion": 2, "normalize_routing_weights": 1 },
@@ -468,6 +467,646 @@
       "bench": {
         "metrics": [{ "type": "gflops", "value": "2 * args.tokens * args.topK * args.hidden * args.inter * 3" }]
       }
     }
   ]
 }

 {
   "tunableSpace": {
     "workgroupSize": [64, 128],
     "decodeLanes": [32],
       "outputs": { "outputT": { "dtype": "float32", "shape": [1, 1024] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "args.topK * 3 * args.hidden * args.inter * 4" }] },
       "provenance": {
+        "notes": "The bandwidth metric counts one float32 read of the selected experts' three gated-MLP weight matrices."
       }
     },
     {
       "outputs": { "outputT": { "dtype": "float32", "shape": [1, 1024] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "args.topK * 2 * args.hidden * args.inter * 4" }] },
       "provenance": {
+        "notes": "The bandwidth metric counts one float32 read of the selected expert's two ReLU-MLP weight matrices."
       }
     },
     {
       "outputs": { "outputT": { "dtype": "float32", "shape": [8, 1024] } },
       "bench": { "metrics": [{ "type": "bandwidth", "value": "args.topK * 3 * args.hidden * args.inter * 4" }] },
       "provenance": {
+        "notes": "The bandwidth metric counts one float32 read of the selected experts' three gated-MLP weight matrices."
       }
     },
     {
       "name": "moe-f32-qwen3-moe-decode-t1-h2048-e32-k8-i768",
       "preset": "model",
       "provenance": {
+        "notes": "Qwen3-MoE geometry uses hidden size 2,048, intermediate size 768, top-k 8, and 32 instantiated experts to keep each weight tensor below 512 MB. The bandwidth metric counts the selected experts' three float32 weight matrices."
       },
       "vars": { "tokens": 1, "hidden": 2048, "experts": 32, "inter": 768, "topK": 8 },
       "attrs": { "k": 8, "activation_type": "swiglu", "swiglu_fusion": 2, "normalize_routing_weights": 1 },
       "name": "moe-f32-olmoe-decode-t1-h2048-e16-k8-i2048",
       "preset": "model",
       "provenance": {
+        "notes": "OLMoE geometry uses hidden and intermediate sizes of 2,048, top-k 8, and 16 instantiated experts. The bandwidth metric counts the selected experts' three float32 weight matrices."
       },
       "vars": { "tokens": 1, "hidden": 2048, "experts": 16, "inter": 2048, "topK": 8 },
       "attrs": { "k": 8, "activation_type": "swiglu", "swiglu_fusion": 2, "normalize_routing_weights": 1 },
       "name": "moe-f32-qwen2-moe-decode-t1-h2048-e20-k4-i1408",
       "preset": "model",
       "provenance": {
+        "notes": "Qwen2-MoE geometry uses hidden size 2,048, intermediate size 1,408, top-k 4, and 20 instantiated experts. The bandwidth metric counts the selected experts' three float32 weight matrices."
       },
       "vars": { "tokens": 1, "hidden": 2048, "experts": 20, "inter": 1408, "topK": 4 },
       "attrs": { "k": 4, "activation_type": "swiglu", "swiglu_fusion": 2, "normalize_routing_weights": 1 },
       "bench": {
         "metrics": [{ "type": "gflops", "value": "2 * args.tokens * args.topK * args.hidden * args.inter * 3" }]
       }
+    },
+    {
+      "name": "boundary-matrix_silu_gate_fc3_bias_only",
+      "attrs": { "k": 1, "activation_type": "silu", "normalize_routing_weights": 0 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.5 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.25, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.18 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.16 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.23, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 64] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1plain_fc3none_fc2plain_relu",
+      "attrs": { "k": 2, "activation_type": "relu", "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [48, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.29, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.3, "cosStep": 0.24, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.12, "cosStep": 0.36, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.14, "cosStep": 0.26, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [48, 64] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1plain_fc3none_fc2bias_swiglu2",
+      "attrs": {
+        "activation_type": "swiglu",
+        "activation_alpha": 1.702,
+        "activation_beta": 0.05,
+        "swiglu_fusion": 2,
+        "normalize_routing_weights": 1
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.107, "cosStep": 0.281, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.317, "cosStep": 0.231, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 192, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.137, "cosStep": 0.351, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.157, "cosStep": 0.251, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.277, "cosStep": 0.321, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 128] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1bias_fc3none_fc2plain_swiglu1",
+      "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 1, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 24, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.124, "cosStep": 0.272, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.334, "cosStep": 0.222, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 192, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.154, "cosStep": 0.342, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.234, "cosStep": 0.422, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.174, "cosStep": 0.242, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 24, 64] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1bias_fc3none_fc2bias_swiglu2",
+      "attrs": { "activation_type": "swiglu", "swiglu_fusion": 2, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.141, "cosStep": 0.263, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.351, "cosStep": 0.213, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 192, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.171, "cosStep": 0.333, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.251, "cosStep": 0.413, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.191, "cosStep": 0.233, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.311, "cosStep": 0.303, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 128] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_partial_tiles_gelu",
+      "attrs": { "activation_type": "gelu", "normalize_routing_weights": 1, "k": 2 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.173, "cosStep": 0.229, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [32, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.347, "cosStep": 0.163, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.199, "cosStep": 0.271, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.251, "cosStep": 0.413, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.241, "cosStep": 0.179, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.311, "cosStep": 0.303, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [32, 64] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1plain_fc3plain_fc2plain_swiglu0",
+      "attrs": { "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 48, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.158, "cosStep": 0.254, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.368, "cosStep": 0.204, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.188, "cosStep": 0.324, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.208, "cosStep": 0.224, "scale": 0.15 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.228, "cosStep": 0.364, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 48, 128] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1plain_fc3plain_fc2bias_swiglu0",
+      "attrs": {
+        "k": 2,
+        "activation_type": "swiglu",
+        "activation_alpha": 1.702,
+        "activation_beta": 0.05,
+        "swiglu_fusion": 0,
+        "normalize_routing_weights": 1
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [48, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.175, "cosStep": 0.245, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.385, "cosStep": 0.195, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.205, "cosStep": 0.315, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.225, "cosStep": 0.215, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.345, "cosStep": 0.285, "scale": 0.2 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.245, "cosStep": 0.355, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [48, 64] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1plain_fc3biased_fc2plain_swiglu0",
+      "attrs": {
+        "activation_type": "swiglu",
+        "activation_alpha": 1.702,
+        "activation_beta": 0.05,
+        "swiglu_fusion": 0,
+        "normalize_routing_weights": 1
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.192, "cosStep": 0.236, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.402, "cosStep": 0.186, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.222, "cosStep": 0.306, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.242, "cosStep": 0.206, "scale": 0.15 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.262, "cosStep": 0.346, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.322, "cosStep": 0.176, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 128] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1plain_fc3biased_fc2bias_swiglu0",
+      "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 24, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.209, "cosStep": 0.227, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.419, "cosStep": 0.177, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.239, "cosStep": 0.297, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.259, "cosStep": 0.197, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.379, "cosStep": 0.267, "scale": 0.2 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.279, "cosStep": 0.337, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.339, "cosStep": 0.167, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 24, 64] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1bias_fc3plain_fc2plain_swiglu0",
+      "attrs": {
+        "k": 2,
+        "activation_type": "swiglu",
+        "activation_alpha": 1.702,
+        "activation_beta": 0.05,
+        "swiglu_fusion": 0,
+        "normalize_routing_weights": 1
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [48, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.226, "cosStep": 0.218, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.436, "cosStep": 0.168, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.256, "cosStep": 0.288, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.336, "cosStep": 0.368, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.276, "cosStep": 0.188, "scale": 0.15 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.296, "cosStep": 0.328, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [48, 128] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1bias_fc3plain_fc2bias_swiglu0",
+      "attrs": { "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.243, "cosStep": 0.209, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.453, "cosStep": 0.159, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.273, "cosStep": 0.279, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.353, "cosStep": 0.359, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.293, "cosStep": 0.179, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.413, "cosStep": 0.249, "scale": 0.2 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.313, "cosStep": 0.319, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 64] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1bias_fc3biased_fc2plain_swiglu0",
+      "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 24, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.26, "cosStep": 0.2, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.47, "cosStep": 0.15, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.27, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.35, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.17, "scale": 0.15 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.33, "cosStep": 0.31, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.39, "cosStep": 0.14, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 24, 128] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_fc1bias_fc3biased_fc2bias_swiglu0",
+      "attrs": {
+        "activation_type": "swiglu",
+        "activation_alpha": 1.702,
+        "activation_beta": 0.05,
+        "swiglu_fusion": 0,
+        "normalize_routing_weights": 1
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.277, "cosStep": 0.191, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.487, "cosStep": 0.141, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.307, "cosStep": 0.261, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.387, "cosStep": 0.341, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.327, "cosStep": 0.161, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.447, "cosStep": 0.231, "scale": 0.2 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.347, "cosStep": 0.301, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.407, "cosStep": 0.131, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 64] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_identity_no_fc3",
+      "attrs": { "k": 2, "activation_type": "identity", "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [48, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.181, "cosStep": 0.217, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.391, "cosStep": 0.127, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 64, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.211, "cosStep": 0.287, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.231, "cosStep": 0.187, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [48, 128] } },
+      "preset": "model"
+    },
+    {
+      "name": "boundary-matrix_unaligned_hidden_inter",
+      "attrs": { "activation_type": "silu", "normalize_routing_weights": 1, "k": 2 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 48, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.163, "cosStep": 0.239, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.337, "cosStep": 0.173, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.209, "cosStep": 0.281, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.261, "cosStep": 0.423, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.251, "cosStep": 0.189, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.321, "cosStep": 0.313, "scale": 0.2 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.227, "cosStep": 0.197, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.283, "cosStep": 0.359, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 48, 64] } },
+      "preset": "model"
     }
   ]
 }

build/webgpu/manifest.json CHANGED Viewed

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

build/webgpu/metadata.json CHANGED Viewed

@@ -1,26 +1,79 @@
 {
   "name": "com.microsoft.MoE",
-  "id": "_com_microsoft_moe_webgpu_9c7f8e1",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "Oq7rf+LMUf5McimJEVeHF3PM3CEHu7X/yCrhSJdoZq4=",
       "expert-group-slots.wgsl.jinja": "Ta+3H2FA1qRRzksdkwgKLZiO+AokmuM+B9JVN8cv1EQ=",
-      "manifest.json": "1u7xRTHJz3qXke+74gJ0e7GZQdJ9SCJw/8gpaOkB9U4=",
-      "moe-ffn-gemv.wgsl.jinja": "loPQLZYp+Lkz/jq0ML5cYbMFP+IWHlpULSiZcaOx33s=",
-      "moe-ffn-grouped.wgsl.jinja": "m5rGC8aq9oW2TsQ2P3B2uV7Hm7u90DFORehlxFnSw5c=",
-      "moe-ffn-stage.wgsl.jinja": "uIz5vVs9A8s5K5B2ZDj4fYDX3BGsWUfR8LxThIR95Tk=",
-      "moe-mix-stage.wgsl.jinja": "+deuFyuBn04UeIX8Hk2/ntCHNfM1vgkNDTcsh3uXVdI=",
       "moe-output-gemv.wgsl.jinja": "5yj+BoTsRHiHW92Q8fVvtojdkY7r1jHupMX+FPVQMBk=",
       "moe-output-grouped.wgsl.jinja": "qzP5G/WtwF4zOBrtySCLTCBflHuvKlU8dk4/CRIzxe8=",
-      "moe-output-stage.wgsl.jinja": "wtB+CtWxl8XChH4Y4/8LEBdHOznrdsUlAo65tmGEuWw=",
-      "moe-route-stage.wgsl.jinja": "6uVRcLbpzp0Y1a8+jRTHOfvBhfCKpCoVju0dBDCVlgo=",
-      "test.json": "T5bB5QuC2mzb2c0i+BQMaTbJjyyvFaL3ikEc4mFJBaw="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.MoE" }
 }

 {
   "name": "com.microsoft.MoE",
+  "id": "_com_microsoft_moe_webgpu_7f9cfff",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "dk7Y+0dMicwCWAn3BjOAjL0VyY0CqCabC1DuWCMDMis=",
       "expert-group-slots.wgsl.jinja": "Ta+3H2FA1qRRzksdkwgKLZiO+AokmuM+B9JVN8cv1EQ=",
+      "manifest.json": "CRRrdafIg5G627yy63AxYRy9pW6V5I93qKobNmgV6vs=",
+      "moe-ffn-gemv.wgsl.jinja": "2LCF3pr7u1cfGz0eUHwkSXr5JQbFgynU8IZS7wFnAuU=",
+      "moe-ffn-grouped.wgsl.jinja": "s1GymR5z3Kj0z7/NOdw6jFgA6D5LM4yfYISrcqmoy+Y=",
+      "moe-ffn-stage.wgsl.jinja": "NNU6+WY7JCjo8itC+pYuTQIwPbjX2PZb1P2WXNS3z/Y=",
+      "moe-grouped-sgmat.wgsl.jinja": "+bkbK0RGxiTY2JlmPLkynnUctqoUcRkdu5OMEl+d9nA=",
+      "moe-mix-stage.wgsl.jinja": "+xEc6boXboFbx+4AJ2/KWElM04LRuufmiIKfEK51/oo=",
       "moe-output-gemv.wgsl.jinja": "5yj+BoTsRHiHW92Q8fVvtojdkY7r1jHupMX+FPVQMBk=",
       "moe-output-grouped.wgsl.jinja": "qzP5G/WtwF4zOBrtySCLTCBflHuvKlU8dk4/CRIzxe8=",
+      "moe-output-stage.wgsl.jinja": "DqLA91ZM2ELUfyDxELZMxHqb0qgmDh/V6Q9k1wVkQH8=",
+      "moe-route-stage.wgsl.jinja": "aj6lm39GKmPH+aAkK7f4i5JpGdb47wPNURH4JQW/A1w=",
+      "test.json": "5rXwo178i5IeETKSYI6HhcwOn6rOVmXBaiy8J12coIk="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "split_routed_fc1plain_fc3none_fc2plain": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1plain_fc3none_fc2plain": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1plain_fc3none_fc2bias": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1plain_fc3none_fc2bias": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1plain_fc3plain_fc2plain": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1plain_fc3plain_fc2plain": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1plain_fc3plain_fc2bias": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1plain_fc3plain_fc2bias": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1plain_fc3biased_fc2plain": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1plain_fc3biased_fc2plain": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1plain_fc3biased_fc2bias": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1plain_fc3biased_fc2bias": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1bias_fc3none_fc2plain": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1bias_fc3none_fc2plain": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1bias_fc3none_fc2bias": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1bias_fc3none_fc2bias": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1bias_fc3plain_fc2plain": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1bias_fc3plain_fc2plain": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1bias_fc3plain_fc2bias": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1bias_fc3plain_fc2bias": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1bias_fc3biased_fc2plain": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1bias_fc3biased_fc2plain": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "split_routed_fc1bias_fc3biased_fc2bias": ["moe-ffn-stage.wgsl.jinja", "moe-output-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "gemv_routed_fc1bias_fc3biased_fc2bias": ["moe-ffn-gemv.wgsl.jinja", "moe-output-gemv.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1plain_fc3none_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1plain_fc3none_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1plain_fc3none_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1plain_fc3none_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1plain_fc3plain_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1plain_fc3plain_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1plain_fc3plain_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1plain_fc3plain_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1plain_fc3biased_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1plain_fc3biased_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1plain_fc3biased_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1plain_fc3biased_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1bias_fc3none_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1bias_fc3none_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1bias_fc3none_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1bias_fc3none_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1bias_fc3plain_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1bias_fc3plain_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1bias_fc3plain_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1bias_fc3plain_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1bias_fc3biased_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1bias_fc3biased_fc2plain": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "sgmat_grouped_routed_fc1bias_fc3biased_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-grouped-sgmat.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-route-stage.wgsl.jinja"],
+      "grouped_routed_fc1bias_fc3biased_fc2bias": ["expert-group-slots.wgsl.jinja", "moe-ffn-grouped.wgsl.jinja", "moe-mix-stage.wgsl.jinja", "moe-output-grouped.wgsl.jinja", "moe-route-stage.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/moe-ffn-gemv.wgsl.jinja CHANGED Viewed

@@ -41,7 +41,7 @@ fn swiglu(gate_in: f32, up_in: f32) -> f32 {
   let gate = gate_in;
   let up = up_in;
 {% endif %}
-  return gate / (1.0 + exp(-params.activationAlpha * gate)) * (up + params.activationBeta);
 }
 {% endif %}

   let gate = gate_in;
   let up = up_in;
 {% endif %}
+  return gate / (1.0 + exp(-{{ matrixActivationAlpha if matrixActivationAlpha is defined else "params.activationAlpha" }} * gate)) * (up + {{ matrixActivationBeta if matrixActivationBeta is defined else "params.activationBeta" }});
 }
 {% endif %}

build/webgpu/moe-ffn-grouped.wgsl.jinja CHANGED Viewed

@@ -1,13 +1,8 @@
 {{ env.wgsl.resourceDeclarations }}
-// FC1 (and, where the schema splits them, FC3) projection plus the activation, on the grouped
-// schedule. The other two schedules read a whole expert weight matrix once per routed slot,
-// which at prefill is the entire cost: for hidden 1024 / inter 512 that is 2 MB of weights per
-// slot, re-read for every slot the expert serves.
-//
-// Here the group stage has already sorted the chunk's slots by expert, so a tile of MTILE rows
-// shares ONE expert -- and therefore one weight tile, staged in workgroup memory and reused by
-// every row.
 const HIDDEN: u32 = {{ hidden }}u;
 const INTER: u32 = {{ inter }}u;
 const FC1_ROWS: u32 = {{ fc1Rows }}u;
@@ -52,7 +47,7 @@ fn swiglu(gate_in: f32, up_in: f32) -> f32 {
   let gate = gate_in;
   let up = up_in;
 {% endif %}
-  return gate / (1.0 + exp(-params.activationAlpha * gate)) * (up + params.activationBeta);
 }
 {% endif %}

 {{ env.wgsl.resourceDeclarations }}
+// Grouped FC1 projection (plus FC3 when present) and activation. Routed slots are
+// sorted by expert before this pass, so every row in an MTILE tile uses the same
+// expert and reuses its staged weight tile.
 const HIDDEN: u32 = {{ hidden }}u;
 const INTER: u32 = {{ inter }}u;
 const FC1_ROWS: u32 = {{ fc1Rows }}u;
   let gate = gate_in;
   let up = up_in;
 {% endif %}
+  return gate / (1.0 + exp(-{{ matrixActivationAlpha if matrixActivationAlpha is defined else "params.activationAlpha" }} * gate)) * (up + {{ matrixActivationBeta if matrixActivationBeta is defined else "params.activationBeta" }});
 }
 {% endif %}

build/webgpu/moe-ffn-stage.wgsl.jinja CHANGED Viewed

@@ -64,17 +64,16 @@ fn swiglu(gate_in: f32, up_in: f32) -> f32 {
   let gate = gate_in;
   let up = up_in;
 {% endif %}
-  return gate / (1.0 + exp(-params.activationAlpha * gate)) * (up + params.activationBeta);
 }
 {% endif %}
 @compute @workgroup_size(WG, 1, 1)
-fn main(@builtin(global_invocation_id) gid: vec3<u32>,
-        @builtin(num_workgroups) nwg: vec3<u32>) {
-  // 2D-folded flat index: gid.y carries the high bits past the maxComputeWorkgroupsPerDimension
-  // workgroup-per-dimension dispatch limit. Reduces to gid.x when nwg.y == 1.
-  let index = gid.x + gid.y * nwg.x * WG;
   let total = params.tokenCount * TOP_K * INTER;
   if (index >= total) {
     return;

   let gate = gate_in;
   let up = up_in;
 {% endif %}
+  return gate / (1.0 + exp(-{{ matrixActivationAlpha if matrixActivationAlpha is defined else "params.activationAlpha" }} * gate)) * (up + {{ matrixActivationBeta if matrixActivationBeta is defined else "params.activationBeta" }});
 }
 {% endif %}
 @compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits past the per-axis dispatch fold width.
+  // Reduces to gid.x when the dispatch does not fold.
+  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
   let total = params.tokenCount * TOP_K * INTER;
   if (index >= total) {
     return;

build/webgpu/moe-grouped-sgmat.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,166 @@

+enable subgroups;
+{% if pinSubgroupSize32 %}
+enable subgroup_size_control;
+{% endif %}
+enable chromium_experimental_subgroup_matrix;
+diagnostic(off, chromium.subgroup_matrix_uniformity);
+{{ env.wgsl.resourceDeclarations }}
+{% set ffn = matrixStage == "ffn" %}
+{% set second = ffn and (hasFc3 or activation == "swiglu") %}
+{% set reduction = hidden if ffn else inter %}
+{% set columns = inter if ffn else hidden %}
+{% if ffn and activation == "gelu" %}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)));
+}{% endif %}
+{% if ffn %}{% if activation == "swiglu" %}
+fn swiglu(gate_in: f32, up_in: f32) -> f32 {
+{% if hasSwigluLimit %}
+  // swiglu_limit clamps the gate operand from above and the linear operand to [-limit, limit]
+  // before the product; when the attribute is absent no clamp is applied.
+  let gate = min(gate_in, {{ swigluLimit }});
+  let up = clamp(up_in, -({{ swigluLimit }}), {{ swigluLimit }});
+{% else %}
+  let gate = gate_in;
+  let up = up_in;
+{% endif %}
+  return gate / (1.0 + exp(-{{ matrixActivationAlpha if matrixActivationAlpha is defined else "params.activationAlpha" }} * gate)) * (up + {{ matrixActivationBeta if matrixActivationBeta is defined else "params.activationBeta" }});
+}
+{% endif %}
+{% endif %}
+const HIDDEN: u32 = {{ hidden }}u;
+const INTER: u32 = {{ inter }}u;
+{% if ffn %}
+const TOP_K: u32 = {{ topK }}u;
+const FC1_ROWS: u32 = {{ fc1Rows }}u;
+{% endif %}
+const TILE_K: u32 = 32u;
+const SUB_ROWS: u32 = 16u;
+const SUB_COLS: u32 = {{ 16 if second else 32 }}u;
+// Two independent f32 chains reduce rounding growth over long reductions.
+// Jinja assigns alternating 8-wide steps to the chains at compile time.
+// Each subgroup publishes four result banks per chain after A is dead.
+// The two row groups reuse those banks with barriers on both sides.
+var<workgroup> tile_A: array<f32, {{ (groupedSgmatSharedBytes / 4) | int }}>;
+{% macro weight_offset(c) %}
+{% if ffn %}
+{% set column = "n_base + subtile_idx * SUB_COLS + " ~ ((c % 2) * 8 if second else c * 8) ~ "u" %}
+{% if second and c >= 2 and hasFc3 %}
+(expert * INTER + {{ column }}) * HIDDEN + kidx + step
+{% else %}
+(expert * FC1_ROWS + {% if second and not hasFc3 and swigluFusion == 1 %}({{ column }}) * 2u{% if c >= 2 %} + 1u{% endif %}{% else %}{{ column }}{% if second and c >= 2 %} + INTER{% endif %}{% endif %}) * HIDDEN + kidx + step
+{% endif %}
+{% else %}
+(expert * HIDDEN + n_base + subtile_idx * SUB_COLS + {{ c * 8 }}u) * INTER + kidx + step
+{% endif %}
+{% endmacro %}
+@compute @workgroup_size(128){{ " @subgroup_size(32)" if pinSubgroupSize32 else "" }}
+fn main(@builtin(workgroup_id) wid: vec3<u32>,
+        @builtin(local_invocation_index) local_idx: u32,
+        @builtin(subgroup_invocation_id) lane: u32,
+        @builtin(subgroup_size) sg_size: u32) {
+  if (wid.x >= tile_meta[0]) { return; }
+  let expert = tile_meta[1u + wid.x * 3u];
+  let slice_base = tile_meta[2u + wid.x * 3u];
+  let rows = tile_meta[3u + wid.x * 3u];
+  let a_row = local_idx / 4u;
+  let a_slot = select(0u, slot_list[slice_base + a_row], a_row < rows);
+  let subtile_id = local_idx / sg_size;
+  let subtile_idy = subtile_id % 2u;
+  let subtile_idx = subtile_id / 2u;
+  let n_base = wid.y * {{ 32 if second else 64 }}u;
+{% for r in range(2) %}{% for c in range(4) %}{% for chain in range(2) %}
+  var mat{{ ["C","D","E","F"][chain] }}{{ r }}{{ c }}: subgroup_matrix_result<f32, 8, 8>;
+{% endfor %}{% endfor %}{% endfor %}
+  for (var kidx = 0u; kidx < {{ reduction }}u; kidx += TILE_K) {
+    for (var i = 0u; i < 8u; i++) {
+      let k = kidx + (local_idx % 4u) * 8u + i;
+      var v = 0.0;
+      if (a_row < rows) {
+{% if ffn %}
+        v = input[(params.tokenOffset + a_slot / TOP_K) * HIDDEN + k];
+{% else %}
+        v = hidden_act[a_slot * INTER + k];
+{% endif %}
+      }
+      tile_A[a_row * TILE_K + (local_idx % 4u) * 8u + i] = v;
+    }
+    workgroupBarrier();
+{% for stepIndex in range(4) %}
+    {
+      let step = {{ stepIndex * 8 }}u;
+{% for r in range(2) %}
+      let matA{{ r }} = subgroupMatrixLoad<subgroup_matrix_left<f32, 8, 8>, row_major>(&tile_A, (subtile_idy * SUB_ROWS + {{ r * 8 }}u) * TILE_K + step, TILE_K);
+{% endfor %}
+{% for c in range(4) %}
+{% set source = ("fc3_experts_weights" if second and c >= 2 and hasFc3 else "fc1_experts_weights") if ffn else "fc2_experts_weights" %}
+{% set stride = hidden * (2 if second and not hasFc3 and swigluFusion == 1 else 1) if ffn else inter %}
+      let matB{{ c }} = subgroupMatrixLoad<subgroup_matrix_right<f32, 8, 8>, col_major>(&{{ source }}, {{ weight_offset(c) }}, {{ stride }}u);
+{% endfor %}
+{% for r in range(2) %}{% for c in range(4) %}
+      mat{{ ["C","D","E","F"][stepIndex % 2] }}{{ r }}{{ c }} = subgroupMatrixMultiplyAccumulate(matA{{ r }}, matB{{ c }}, mat{{ ["C","D","E","F"][stepIndex % 2] }}{{ r }}{{ c }});
+{% endfor %}{% endfor %}
+    }
+{% endfor %}
+    workgroupBarrier();
+  }
+  let row = lane / 4u;
+  let lane_col = (lane % 4u) * 2u;
+{% for r in range(2) %}
+{% for c in range(4) %}
+{% for chain in range(2) %}
+  subgroupMatrixStore<row_major>(&tile_A, {{ chain * 1024 }}u + (subtile_id * 4u + {{ c }}u) * 64u, mat{{ ["C","D","E","F"][chain] }}{{ r }}{{ c }}, 8u);
+{% endfor %}
+{% endfor %}
+  workgroupBarrier();
+  {
+    let m = subtile_idy * SUB_ROWS + {{ r * 8 }}u + row;
+    if (m < rows) {
+      let out_slot = slot_list[slice_base + m];
+{% for c in range(2 if second else 4) %}{% for half in range(2) %}
+      {
+        let n = n_base + subtile_idx * SUB_COLS + {{ c * 8 + half }}u + lane_col;
+        let col = n;
+        if (col < {{ columns }}u) {
+          let bank = (subtile_id * 4u + {{ c }}u) * 64u + row * 8u + lane_col + {{ half }}u;
+{% if ffn %}
+          let a = ((tile_A[bank] + tile_A[1024u + bank])){% if hasFc1Bias %} + fc1_experts_bias[expert * FC1_ROWS + {{ "col * 2u" if second and not hasFc3 and swigluFusion == 1 else "col" }}]{% endif %};
+{% if second %}
+          let b = ((tile_A[bank + 2u * 64u] + tile_A[1024u + bank + 2u * 64u])){% if hasFc3 and hasFc3Bias %} + fc3_experts_bias[expert * INTER + col]{% elif not hasFc3 and hasFc1Bias %} + fc1_experts_bias[expert * FC1_ROWS + {{ "col * 2u + 1u" if swigluFusion == 1 else "INTER + col" }}]{% endif %};
+{% endif %}
+{% if activation == "swiglu" %}
+          let value = swiglu(a, b);
+{% else %}
+{% if activation == "relu" %}
+          let activated = max(a, 0.0);
+{% elif activation == "gelu" %}
+          let activated = gelu_tanh(a);
+{% elif activation == "silu" %}
+          let activated = a / (1.0 + exp(-a));
+{% else %}
+          let activated = a;
+{% endif %}
+          let value = activated{% if hasFc3 %} * b{% endif %};
+{% endif %}
+          hidden_act[out_slot * INTER + col] = value;
+{% else %}
+          slot_out[out_slot * HIDDEN + col] = ((tile_A[bank] + tile_A[1024u + bank])){% if hasFc2Bias %} + fc2_experts_bias[expert * HIDDEN + col]{% endif %};
+{% endif %}
+        }
+      }
+{% endfor %}{% endfor %}
+    }
+  }
+{% if r == 0 %}  workgroupBarrier();{% endif %}
+{% endfor %}
+}

build/webgpu/moe-mix-stage.wgsl.jinja CHANGED Viewed

@@ -10,11 +10,10 @@ const TOP_K: u32 = {{ topK }}u;
 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
-fn main(@builtin(global_invocation_id) gid: vec3<u32>,
-        @builtin(num_workgroups) nwg: vec3<u32>) {
-  // 2D-folded flat index: gid.y carries the high bits past the maxComputeWorkgroupsPerDimension
-  // workgroup-per-dimension dispatch limit. Reduces to gid.x when nwg.y == 1.
-  let index = gid.x + gid.y * nwg.x * WG;
   let total = params.tokenCount * HIDDEN;
   if (index >= total) {
     return;

 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits past the per-axis dispatch fold width.
+  // Reduces to gid.x when the dispatch does not fold.
+  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
   let total = params.tokenCount * HIDDEN;
   if (index >= total) {
     return;

build/webgpu/moe-output-stage.wgsl.jinja CHANGED Viewed

@@ -8,11 +8,10 @@ const TOP_K: u32 = {{ topK }}u;
 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
-fn main(@builtin(global_invocation_id) gid: vec3<u32>,
-        @builtin(num_workgroups) nwg: vec3<u32>) {
-  // 2D-folded flat index: gid.y carries the high bits past the maxComputeWorkgroupsPerDimension
-  // workgroup-per-dimension dispatch limit. Reduces to gid.x when nwg.y == 1.
-  let index = gid.x + gid.y * nwg.x * WG;
   let total = params.tokenCount * HIDDEN;
   if (index >= total) {
     return;

 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits past the per-axis dispatch fold width.
+  // Reduces to gid.x when the dispatch does not fold.
+  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
   let total = params.tokenCount * HIDDEN;
   if (index >= total) {
     return;

build/webgpu/moe-route-stage.wgsl.jinja CHANGED Viewed

@@ -10,10 +10,9 @@ const TOP_K: u32 = {{ topK }}u;
 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
-fn main(@builtin(global_invocation_id) gid: vec3<u32>,
-        @builtin(num_workgroups) nwg: vec3<u32>) {
   // gid.y carries the high bits past the per-dimension dispatch limit.
-  let token = gid.x + gid.y * nwg.x * WG;
   if (token >= TOKENS) {
     return;
   }

 const WG: u32 = {{ workgroupSize }}u;
 @compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // gid.y carries the high bits past the per-dimension dispatch limit.
+  let token = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
   if (token >= TOKENS) {
     return;
   }

build/webgpu/test.json CHANGED Viewed

@@ -1,12 +1,11 @@
 {
-  "op": "com.microsoft.MoE",
   "cases": [
     {
       "name": "relu_top1_no_bias",
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "default activation_type (relu), one expert per token",
-        "notes": "The schema's default activation and nothing optional: no biases, no FC3. Every other case is a delta from this one."
       },
       "attrs": { "normalize_routing_weights": 1 },
       "inputs": {
@@ -38,7 +37,7 @@
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "activation_type gelu with an FC1 bias",
-        "notes": "The tanh-approximate GELU used by ONNX Runtime's MoE kernel, over an FC1 that carries its own per-expert bias. The bias joins the projection before the activation, which is what separates it from a bias on the output side."
       },
       "attrs": { "k": 1, "activation_type": "gelu", "normalize_routing_weights": 0 },
       "inputs": {
@@ -149,7 +148,7 @@
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "all three optional biases present at once",
-        "notes": "FC1, FC2 and FC3 biases together on the provider-defined SiLU gated path."
       },
       "attrs": { "k": 2, "activation_type": "silu", "normalize_routing_weights": 1 },
       "inputs": {
@@ -461,7 +460,7 @@
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "com.microsoft.MoE optional-input combination",
-        "notes": "Each optional input changes the binding layout, so every combination of FC1 bias, FC3 (with or without its own bias) and FC2 bias is a distinct variant. This case is the one that selects fc1bias / no FC3 / fc2bias; without it that variant would ship unexercised."
       },
       "attrs": { "k": 1, "activation_type": "gelu", "normalize_routing_weights": 0 },
       "inputs": {
@@ -587,7 +586,7 @@
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "com.microsoft.MoE optional-input combination",
-        "notes": "Each optional input changes the binding layout, so every combination of FC1 bias, FC3 (with or without its own bias) and FC2 bias is a distinct variant. This case is the one that selects no FC1 bias / FC3 with bias / fc2bias; without it that variant would ship unexercised."
       },
       "attrs": { "k": 2, "activation_type": "silu", "normalize_routing_weights": 1 },
       "inputs": {
@@ -634,7 +633,7 @@
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "com.microsoft.MoE optional-input combination",
-        "notes": "Each optional input changes the binding layout, so every combination of FC1 bias, FC3 (with or without its own bias) and FC2 bias is a distinct variant. This case is the one that selects fc1bias / FC3 without bias / no FC2 bias; without it that variant would ship unexercised."
       },
       "attrs": { "k": 1, "activation_type": "silu", "normalize_routing_weights": 0 },
       "inputs": {
@@ -676,7 +675,7 @@
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "com.microsoft.MoE optional-input combination",
-        "notes": "Each optional input changes the binding layout, so every combination of FC1 bias, FC3 (with or without its own bias) and FC2 bias is a distinct variant. This case is the one that selects fc1bias / FC3 without bias / fc2bias; without it that variant would ship unexercised."
       },
       "attrs": { "k": 2, "activation_type": "silu", "normalize_routing_weights": 1 },
       "inputs": {
@@ -723,7 +722,7 @@
       "provenance": {
         "source": "onnxruntime/contrib_ops/cpu/moe/moe_cpu.cc",
         "test": "full-softmax routing and equal-probability pair ordering",
-        "notes": "Equal zero logits become probabilities [0.5, 0.5]. ONNX Runtime selects the higher expert index on the tie; that expert emits 6, and the non-normalized route probability scales the result to exactly 3. This also exercises the schema defaults k=1, activation_type=relu, and normalize_routing_weights=0."
       },
       "inputs": {
         "inputT": { "dtype": "float32", "shape": [1, 1], "data": { "kind": "values", "values": [1.0] } },
@@ -745,7 +744,7 @@
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "com.microsoft.MoE optional-input combination",
-        "notes": "Each optional input changes the binding layout, so every combination of FC1 bias, FC3 (with or without its own bias) and FC2 bias is a distinct variant. This case is the one that selects fc1bias / FC3 with bias / no FC2 bias; without it that variant would ship unexercised."
       },
       "attrs": { "k": 1, "activation_type": "silu", "normalize_routing_weights": 0 },
       "inputs": {
@@ -790,7 +789,7 @@
     {
       "name": "deep_reduction_fc1plain_fc3none_fc2plain_identity",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": { "activation_type": "identity", "normalize_routing_weights": 1 },
       "inputs": {
@@ -820,7 +819,7 @@
     {
       "name": "deep_reduction_fc1plain_fc3none_fc2bias_swiglu2",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": {
         "k": 2,
@@ -862,7 +861,7 @@
     {
       "name": "deep_reduction_fc1bias_fc3none_fc2plain_swiglu1",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 1, "normalize_routing_weights": 1 },
       "inputs": {
@@ -897,7 +896,7 @@
     {
       "name": "deep_reduction_fc1bias_fc3none_fc2bias_swiglu2",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": { "activation_type": "swiglu", "swiglu_fusion": 2, "normalize_routing_weights": 1 },
       "inputs": {
@@ -937,7 +936,7 @@
     {
       "name": "deep_reduction_fc1plain_fc3plain_fc2plain_swiglu0",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": { "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
       "inputs": {
@@ -972,7 +971,7 @@
     {
       "name": "deep_reduction_fc1plain_fc3plain_fc2bias_swiglu0",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": {
         "k": 2,
@@ -1019,7 +1018,7 @@
     {
       "name": "deep_reduction_fc1plain_fc3biased_fc2plain_swiglu0",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": {
         "activation_type": "swiglu",
@@ -1065,7 +1064,7 @@
     {
       "name": "deep_reduction_fc1plain_fc3biased_fc2bias_swiglu0",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
       "inputs": {
@@ -1110,7 +1109,7 @@
     {
       "name": "deep_reduction_fc1bias_fc3plain_fc2plain_swiglu0",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": {
         "k": 2,
@@ -1157,7 +1156,7 @@
     {
       "name": "deep_reduction_fc1bias_fc3plain_fc2bias_swiglu0",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": { "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
       "inputs": {
@@ -1202,7 +1201,7 @@
     {
       "name": "deep_reduction_fc1bias_fc3biased_fc2plain_swiglu0",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
       "inputs": {
@@ -1247,7 +1246,7 @@
     {
       "name": "deep_reduction_fc1bias_fc3biased_fc2bias_swiglu0",
       "provenance": {
-        "notes": "A 128-long reduction on both projections, deep enough for the cooperative schedule; the 4-wide cases above stay on the one-thread-per-column schedule."
       },
       "attrs": {
         "activation_type": "swiglu",
@@ -1303,7 +1302,7 @@
     {
       "name": "identity_activation_top1",
       "provenance": {
-        "notes": "activation_type identity passes the FC1 projection through unchanged, which is the only declared activation with no other case."
       },
       "attrs": { "activation_type": "identity" },
       "inputs": {
@@ -1514,7 +1513,7 @@
     {
       "name": "grouped_prefill_partial_tiles_gelu",
       "provenance": {
-        "notes": "Half the routed slots the grouped schedule used to demand, so most of every expert's tile is padding: the tiles are 32 rows and 3 experts share only 64 slots. Covers the band the slot threshold opened, where the store guard drops more rows than it keeps."
       },
       "attrs": { "activation_type": "gelu", "normalize_routing_weights": 1, "k": 2 },
       "inputs": {
@@ -1996,6 +1995,678 @@
         }
       },
       "outputs": { "outputT": { "dtype": "float32", "shape": [96, 34], "tolerance": 0.00002 } }
     }
   ]
 }

 {
   "cases": [
     {
       "name": "relu_top1_no_bias",
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "default activation_type (relu), one expert per token",
+        "notes": "Exercises the schema's default activation with no optional biases or FC3 input."
       },
       "attrs": { "normalize_routing_weights": 1 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "activation_type gelu with an FC1 bias",
+        "notes": "Tanh-approximate GELU after the FC1 projection and its per-expert bias."
       },
       "attrs": { "k": 1, "activation_type": "gelu", "normalize_routing_weights": 0 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "all three optional biases present at once",
+        "notes": "The SiLU gated path applies FC1, FC2, and FC3 biases in one invocation."
       },
       "attrs": { "k": 2, "activation_type": "silu", "normalize_routing_weights": 1 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "com.microsoft.MoE optional-input combination",
+        "notes": "Uses FC1 and FC2 biases without an FC3 projection, exercising GELU on the FC1 projection before the biased FC2 projection."
       },
       "attrs": { "k": 1, "activation_type": "gelu", "normalize_routing_weights": 0 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "com.microsoft.MoE optional-input combination",
+        "notes": "Uses the gated SiLU branch with an FC3 bias and an FC2 bias while omitting the FC1 bias; routing weights are normalized across the selected experts."
       },
       "attrs": { "k": 2, "activation_type": "silu", "normalize_routing_weights": 1 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "com.microsoft.MoE optional-input combination",
+        "notes": "Uses the gated SiLU branch with an FC1 bias, an unbiased FC3 projection, and no FC2 bias."
       },
       "attrs": { "k": 1, "activation_type": "silu", "normalize_routing_weights": 0 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "com.microsoft.MoE optional-input combination",
+        "notes": "Uses normalized top-2 routing with the gated SiLU branch, an FC1 bias, an unbiased FC3 projection, and an FC2 bias."
       },
       "attrs": { "k": 2, "activation_type": "silu", "normalize_routing_weights": 1 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/contrib_ops/cpu/moe/moe_cpu.cc",
         "test": "full-softmax routing and equal-probability pair ordering",
+        "notes": "Equal zero logits become probabilities [0.5, 0.5], with ties selecting the higher expert index. That expert emits 6, and its non-normalized route probability scales the result to exactly 3. Also exercises the defaults k=1, activation_type=relu, and normalize_routing_weights=0."
       },
       "inputs": {
         "inputT": { "dtype": "float32", "shape": [1, 1], "data": { "kind": "values", "values": [1.0] } },
       "provenance": {
         "source": "onnxruntime/docs/ContribOperators.md#com.microsoft.MoE",
         "test": "com.microsoft.MoE optional-input combination",
+        "notes": "Uses the gated SiLU branch with FC1 and FC3 biases and no FC2 bias."
       },
       "attrs": { "k": 1, "activation_type": "silu", "normalize_routing_weights": 0 },
       "inputs": {
     {
       "name": "deep_reduction_fc1plain_fc3none_fc2plain_identity",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with plain fc1, no fc3, plain fc2, and identity activation."
       },
       "attrs": { "activation_type": "identity", "normalize_routing_weights": 1 },
       "inputs": {
     {
       "name": "deep_reduction_fc1plain_fc3none_fc2bias_swiglu2",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with plain fc1, no fc3, biased fc2, and SwiGLU formula 2."
       },
       "attrs": {
         "k": 2,
     {
       "name": "deep_reduction_fc1bias_fc3none_fc2plain_swiglu1",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with biased fc1, no fc3, plain fc2, and SwiGLU formula 1."
       },
       "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 1, "normalize_routing_weights": 1 },
       "inputs": {
     {
       "name": "deep_reduction_fc1bias_fc3none_fc2bias_swiglu2",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with biased fc1, no fc3, biased fc2, and SwiGLU formula 2."
       },
       "attrs": { "activation_type": "swiglu", "swiglu_fusion": 2, "normalize_routing_weights": 1 },
       "inputs": {
     {
       "name": "deep_reduction_fc1plain_fc3plain_fc2plain_swiglu0",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with plain fc1, plain fc3, plain fc2, and formula 0."
       },
       "attrs": { "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
       "inputs": {
     {
       "name": "deep_reduction_fc1plain_fc3plain_fc2bias_swiglu0",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with plain fc1, plain fc3, biased fc2, and formula 0."
       },
       "attrs": {
         "k": 2,
     {
       "name": "deep_reduction_fc1plain_fc3biased_fc2plain_swiglu0",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with plain fc1, biased fc3, plain fc2, and formula 0."
       },
       "attrs": {
         "activation_type": "swiglu",
     {
       "name": "deep_reduction_fc1plain_fc3biased_fc2bias_swiglu0",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with plain fc1, biased fc3, biased fc2, and formula 0."
       },
       "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
       "inputs": {
     {
       "name": "deep_reduction_fc1bias_fc3plain_fc2plain_swiglu0",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with biased fc1, plain fc3, plain fc2, and formula 0."
       },
       "attrs": {
         "k": 2,
     {
       "name": "deep_reduction_fc1bias_fc3plain_fc2bias_swiglu0",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with biased fc1, plain fc3, biased fc2, and formula 0."
       },
       "attrs": { "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
       "inputs": {
     {
       "name": "deep_reduction_fc1bias_fc3biased_fc2plain_swiglu0",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with biased fc1, biased fc3, plain fc2, and formula 0."
       },
       "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
       "inputs": {
     {
       "name": "deep_reduction_fc1bias_fc3biased_fc2bias_swiglu0",
       "provenance": {
+        "notes": "A 128-element reduction on both projections exercises the cooperative schedule with biased fc1, biased fc3, biased fc2, and formula 0."
       },
       "attrs": {
         "activation_type": "swiglu",
     {
       "name": "identity_activation_top1",
       "provenance": {
+        "notes": "With `activation_type=identity`, the FC1 projection passes through unchanged before the output projection."
       },
       "attrs": { "activation_type": "identity" },
       "inputs": {
     {
       "name": "grouped_prefill_partial_tiles_gelu",
       "provenance": {
+        "notes": "Three experts share 64 routed slots across 32-row tiles, so most lanes address padding and the store guard discards more rows than it writes."
       },
       "attrs": { "activation_type": "gelu", "normalize_routing_weights": 1, "k": 2 },
       "inputs": {
         }
       },
       "outputs": { "outputT": { "dtype": "float32", "shape": [96, 34], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_silu_gate_fc3_bias_only",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_silu_gate_fc3_bias_only. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "k": 1, "activation_type": "silu", "normalize_routing_weights": 0 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.31, "scale": 0.5 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.25, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.18 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41, "scale": 0.16 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.23, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 64], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1plain_fc3none_fc2plain_relu",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1plain_fc3none_fc2plain_relu. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "k": 2, "activation_type": "relu", "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [48, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.29, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.3, "cosStep": 0.24, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.12, "cosStep": 0.36, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.14, "cosStep": 0.26, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [48, 64], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1plain_fc3none_fc2bias_swiglu2",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1plain_fc3none_fc2bias_swiglu2. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": {
+        "activation_type": "swiglu",
+        "activation_alpha": 1.702,
+        "activation_beta": 0.05,
+        "swiglu_fusion": 2,
+        "normalize_routing_weights": 1
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.107, "cosStep": 0.281, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.317, "cosStep": 0.231, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 192, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.137, "cosStep": 0.351, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.157, "cosStep": 0.251, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.277, "cosStep": 0.321, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 128], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1bias_fc3none_fc2plain_swiglu1",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1bias_fc3none_fc2plain_swiglu1. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 1, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 24, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.124, "cosStep": 0.272, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.334, "cosStep": 0.222, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 192, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.154, "cosStep": 0.342, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.234, "cosStep": 0.422, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.174, "cosStep": 0.242, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 24, 64], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1bias_fc3none_fc2bias_swiglu2",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1bias_fc3none_fc2bias_swiglu2. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "activation_type": "swiglu", "swiglu_fusion": 2, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.141, "cosStep": 0.263, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.351, "cosStep": 0.213, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 192, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.171, "cosStep": 0.333, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.251, "cosStep": 0.413, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.191, "cosStep": 0.233, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.311, "cosStep": 0.303, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 128], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_partial_tiles_gelu",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_partial_tiles_gelu. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "activation_type": "gelu", "normalize_routing_weights": 1, "k": 2 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [32, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.173, "cosStep": 0.229, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [32, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.347, "cosStep": 0.163, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.199, "cosStep": 0.271, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.251, "cosStep": 0.413, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.241, "cosStep": 0.179, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.311, "cosStep": 0.303, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [32, 64], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1plain_fc3plain_fc2plain_swiglu0",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1plain_fc3plain_fc2plain_swiglu0. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 48, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.158, "cosStep": 0.254, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.368, "cosStep": 0.204, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.188, "cosStep": 0.324, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.208, "cosStep": 0.224, "scale": 0.15 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.228, "cosStep": 0.364, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 48, 128], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1plain_fc3plain_fc2bias_swiglu0",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1plain_fc3plain_fc2bias_swiglu0. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": {
+        "k": 2,
+        "activation_type": "swiglu",
+        "activation_alpha": 1.702,
+        "activation_beta": 0.05,
+        "swiglu_fusion": 0,
+        "normalize_routing_weights": 1
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [48, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.175, "cosStep": 0.245, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.385, "cosStep": 0.195, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.205, "cosStep": 0.315, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.225, "cosStep": 0.215, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.345, "cosStep": 0.285, "scale": 0.2 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.245, "cosStep": 0.355, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [48, 64], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1plain_fc3biased_fc2plain_swiglu0",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1plain_fc3biased_fc2plain_swiglu0. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": {
+        "activation_type": "swiglu",
+        "activation_alpha": 1.702,
+        "activation_beta": 0.05,
+        "swiglu_fusion": 0,
+        "normalize_routing_weights": 1
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.192, "cosStep": 0.236, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.402, "cosStep": 0.186, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.222, "cosStep": 0.306, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.242, "cosStep": 0.206, "scale": 0.15 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.262, "cosStep": 0.346, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.322, "cosStep": 0.176, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 128], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1plain_fc3biased_fc2bias_swiglu0",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1plain_fc3biased_fc2bias_swiglu0. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 24, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.209, "cosStep": 0.227, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.419, "cosStep": 0.177, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.239, "cosStep": 0.297, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.259, "cosStep": 0.197, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.379, "cosStep": 0.267, "scale": 0.2 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.279, "cosStep": 0.337, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.339, "cosStep": 0.167, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 24, 64], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1bias_fc3plain_fc2plain_swiglu0",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1bias_fc3plain_fc2plain_swiglu0. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": {
+        "k": 2,
+        "activation_type": "swiglu",
+        "activation_alpha": 1.702,
+        "activation_beta": 0.05,
+        "swiglu_fusion": 0,
+        "normalize_routing_weights": 1
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [48, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.226, "cosStep": 0.218, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.436, "cosStep": 0.168, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.256, "cosStep": 0.288, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.336, "cosStep": 0.368, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.276, "cosStep": 0.188, "scale": 0.15 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.296, "cosStep": 0.328, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [48, 128], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1bias_fc3plain_fc2bias_swiglu0",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1bias_fc3plain_fc2bias_swiglu0. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.243, "cosStep": 0.209, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.453, "cosStep": 0.159, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.273, "cosStep": 0.279, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.353, "cosStep": 0.359, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.293, "cosStep": 0.179, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.413, "cosStep": 0.249, "scale": 0.2 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.313, "cosStep": 0.319, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 64], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1bias_fc3biased_fc2plain_swiglu0",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1bias_fc3biased_fc2plain_swiglu0. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "k": 2, "activation_type": "swiglu", "swiglu_fusion": 0, "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 24, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.26, "cosStep": 0.2, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.47, "cosStep": 0.15, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.27, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.35, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.17, "scale": 0.15 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.33, "cosStep": 0.31, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.39, "cosStep": 0.14, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 24, 128], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_fc1bias_fc3biased_fc2bias_swiglu0",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_fc1bias_fc3biased_fc2bias_swiglu0. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": {
+        "activation_type": "swiglu",
+        "activation_alpha": 1.702,
+        "activation_beta": 0.05,
+        "swiglu_fusion": 0,
+        "normalize_routing_weights": 1
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.277, "cosStep": 0.191, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.487, "cosStep": 0.141, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.307, "cosStep": 0.261, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.387, "cosStep": 0.341, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.327, "cosStep": 0.161, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.447, "cosStep": 0.231, "scale": 0.2 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.347, "cosStep": 0.301, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.407, "cosStep": 0.131, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [96, 64], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_identity_no_fc3",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_identity_no_fc3. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "k": 2, "activation_type": "identity", "normalize_routing_weights": 1 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [48, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.181, "cosStep": 0.217, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [48, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.391, "cosStep": 0.127, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 64, 128],
+          "data": { "kind": "fillFloat32", "sinStep": 0.211, "cosStep": 0.287, "scale": 0.15 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 128, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.231, "cosStep": 0.187, "scale": 0.15 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [48, 128], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "matrix_unaligned_hidden_inter",
+      "provenance": {
+        "notes": "Direct-weight matrix sibling of grouped_prefill_unaligned_hidden_inter. Preserves the activation, biases, input recipes, and tolerance; varies aligned reduction extents and rank while routed expert slices may end in partial row tiles."
+      },
+      "attrs": { "activation_type": "silu", "normalize_routing_weights": 1, "k": 2 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 48, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.163, "cosStep": 0.239, "scale": 0.4 }
+        },
+        "routerT": {
+          "dtype": "float32",
+          "shape": [96, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.337, "cosStep": 0.173, "scale": 0.7 }
+        },
+        "fc1T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.209, "cosStep": 0.281, "scale": 0.15 }
+        },
+        "fc1BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.261, "cosStep": 0.423, "scale": 0.2 }
+        },
+        "fc2T": {
+          "dtype": "float32",
+          "shape": [3, 64, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.251, "cosStep": 0.189, "scale": 0.15 }
+        },
+        "fc2BiasT": {
+          "dtype": "float32",
+          "shape": [3, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.321, "cosStep": 0.313, "scale": 0.2 }
+        },
+        "fc3T": {
+          "dtype": "float32",
+          "shape": [3, 96, 64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.227, "cosStep": 0.197, "scale": 0.15 }
+        },
+        "fc3BiasT": {
+          "dtype": "float32",
+          "shape": [3, 96],
+          "data": { "kind": "fillFloat32", "sinStep": 0.283, "cosStep": 0.359, "scale": 0.2 }
+        }
+      },
+      "outputs": { "outputT": { "dtype": "float32", "shape": [2, 48, 64], "tolerance": 0.00002 } }
     }
   ]
 }