sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +75 -0
build/webgpu/bench.json +402 -0
build/webgpu/causal-conv-with-state-tiled.wgsl.jinja +174 -0
build/webgpu/causal-conv-with-state-vec4.wgsl.jinja +109 -0
build/webgpu/causal-conv-with-state.wgsl.jinja +105 -0
build/webgpu/manifest.json +1257 -0
build/webgpu/metadata.json +20 -0
build/webgpu/test.json +1283 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.CausalConvWithState
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Legacy Microsoft contrib form of stateful 1-D causal depthwise convolution. Each channel uses its own `(channels, 1, kernel)` weight over current and past positions, with optional activation and `past_state`/`present_state` tensors for incremental decoding. The contrib-only `state_window` attribute may retain several rollback states. This inference implementation preserves the existing contrib ABI with `ndim = 1`, float16 or float32 tensors, and float32 accumulation; spatial ranks 2 and 3 and bfloat16 are not implemented.
+See the [ONNX Runtime `CausalConvWithState` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.CausalConvWithState) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `inputT` | `T` | `3` | — | Channels-first input tensor with shape `(batch_size, channels, sequence_length)` for the supported 1-D mode. | required |
+| `weight` | `weightT` | `T` | `3` | — | Depthwise convolution kernel with shape `(channels, 1, kernel_size)` for the supported 1-D mode. | required |
+| `bias` | `biasT` | `T` | `1` | — | Optional per-channel bias with shape `(channels,)`. | optional |
+| `past_state` | `pastStateT` | `T` | derived | — | Carry state from the previous step; shape `(batch_size, channels, k_1 - 1)`, or `(W, batch_size, channels, k_1 - 1)` when `state_window = W > 0`, in which case only slot `W - 1` is read. If absent, the left-side padding is zero. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `outputT` | `T` | `3` | same as `input` | Convolution output with the same shape as `input`. | required |
+| `present_state` | `presentStateT` | `T` | derived | derived; see description | Updated carry state; shape `(batch_size, channels, k_1 - 1)`, or `(W, batch_size, channels, k_1 - 1)` when `state_window = W > 0`. Slot `W - 1` holds the last `k - 1` values along the causal axis; slot `j` holds the same for the prefix ending at position `seq_len - W + j`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `activation` | `"none"` | Activation applied after convolution and bias. Defaults to `none`; `swish` is an alias of SiLU. |
+| `ndim` | `1` | Number of spatial dimensions. This implementation supports the contrib 1D mode (`ndim = 1`). |
+| `state_window` | `0` | Contrib extension selecting the number of rollback state slots to retain, in the range 0 through 8. Defaults to 0. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## 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
+- [`causal-conv-with-state-tiled.wgsl.jinja`](build/webgpu/causal-conv-with-state-tiled.wgsl.jinja)
+- [`causal-conv-with-state-vec4.wgsl.jinja`](build/webgpu/causal-conv-with-state-vec4.wgsl.jinja)
+- [`causal-conv-with-state.wgsl.jinja`](build/webgpu/causal-conv-with-state.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/com.microsoft.CausalConvWithState", { version: 1 });
+const { outputT, presentStateT } = await kernel({
+  inputT: { data: inputTData, shape: [1, 1, 5] },
+  weightT: { data: weightTData, shape: [1, 1, 4] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,402 @@

+{
+  "op": "com.microsoft.CausalConvWithState",
+  "tunableSpace": { "workgroupSize": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "causal-conv-f32-b2c32t256k4",
+      "preset": "smoke",
+      "vars": { "batch": 2, "channels": 32, "length": 256, "kernel": 4 },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": { "shape": [2, 32, 256], "dtype": "float32", "dist": "normal", "seed": 205, "scale": 0.2 },
+        "weightT": { "shape": [32, 1, 4], "dtype": "float32", "dist": "normal", "seed": 206, "scale": 0.1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [2, 32, 256], "dtype": "float32" },
+        "presentStateT": { "shape": [2, 32, 3], "dtype": "float32" }
+      },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 1)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-f32-state-bias-silu-b2c32t256k4",
+      "preset": "smoke",
+      "vars": { "batch": 2, "channels": 32, "length": 256, "kernel": 4 },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "shape": [2, 32, 256], "dtype": "float32", "dist": "normal", "seed": 206, "scale": 0.2 },
+        "weightT": { "shape": [32, 1, 4], "dtype": "float32", "dist": "normal", "seed": 207, "scale": 0.1 },
+        "biasT": { "shape": [32], "dtype": "float32", "dist": "normal", "seed": 208, "scale": 0.05 },
+        "pastStateT": { "shape": [2, 32, 3], "dtype": "float32", "dist": "normal", "seed": 209, "scale": 0.2 }
+      },
+      "outputs": {
+        "outputT": { "shape": [2, 32, 256], "dtype": "float32" },
+        "presentStateT": { "shape": [2, 32, 3], "dtype": "float32" }
+      },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 1)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.biasT) + numel(shapes.pastStateT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-f32-large-kernel-prefill-b2c1024t512k128",
+      "preset": "stress",
+      "tunableSpace": { "tiledWorkgroupSize": [64, 128, 256] },
+      "vars": { "batch": 2, "channels": 1024, "length": 512, "kernel": 128 },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "normal", "seed": 301, "scale": 0.2 },
+        "weightT": { "shape": [1024, 1, 128], "dtype": "float32", "dist": "normal", "seed": 302, "scale": 0.05 }
+      },
+      "outputs": {
+        "outputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "empty" },
+        "presentStateT": { "shape": [2, 1024, 127], "dtype": "float32", "dist": "empty" }
+      },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 1)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-f32-large-kernel127-prefill-b2c1024t512-alignment-pathology",
+      "preset": "stress",
+      "provenance": {
+        "notes": "A realistic odd-width causal kernel adjacent to the K=128 tiled control. K=127 is valid but fails the tiled kernel's K % 4 contract and falls back to the scalar zero-state implementation."
+      },
+      "vars": { "batch": 2, "channels": 1024, "length": 512, "kernel": 127 },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "normal", "seed": 310, "scale": 0.2 },
+        "weightT": { "shape": [1024, 1, 127], "dtype": "float32", "dist": "normal", "seed": 311, "scale": 0.05 }
+      },
+      "outputs": {
+        "outputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "empty" },
+        "presentStateT": { "shape": [2, 1024, 126], "dtype": "float32", "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 1)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-f32-bias-large-kernel-prefill-b2c1024t512k128-pathology",
+      "preset": "stress",
+      "tunableSpace": { "tiledWorkgroupSize": [64, 128, 256] },
+      "provenance": {
+        "notes": "Regression lock for the historical large-kernel bias cliff: this common per-channel bias path previously excluded the tiled specialization and selected the scalar kernel."
+      },
+      "vars": { "batch": 2, "channels": 1024, "length": 512, "kernel": 128 },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "normal", "seed": 307, "scale": 0.2 },
+        "weightT": { "shape": [1024, 1, 128], "dtype": "float32", "dist": "normal", "seed": 308, "scale": 0.05 },
+        "biasT": { "shape": [1024], "dtype": "float32", "dist": "normal", "seed": 309, "scale": 0.05 }
+      },
+      "outputs": {
+        "outputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "empty" },
+        "presentStateT": { "shape": [2, 1024, 127], "dtype": "float32", "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 1)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.biasT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-f32-state-large-kernel-prefill-b2c1024t512k128-pathology",
+      "preset": "stress",
+      "tunableSpace": { "tiledWorkgroupSize": [64, 128, 256] },
+      "provenance": {
+        "notes": "Regression lock for the historical continuation cliff: valid carry state previously excluded the tiled specialization even without bias or activation."
+      },
+      "vars": { "batch": 2, "channels": 1024, "length": 512, "kernel": 128 },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "normal", "seed": 312, "scale": 0.2 },
+        "weightT": { "shape": [1024, 1, 128], "dtype": "float32", "dist": "normal", "seed": 313, "scale": 0.05 },
+        "pastStateT": { "shape": [2, 1024, 127], "dtype": "float32", "dist": "normal", "seed": 314, "scale": 0.2 }
+      },
+      "outputs": {
+        "outputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "empty" },
+        "presentStateT": { "shape": [2, 1024, 127], "dtype": "float32", "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 1)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.pastStateT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-f32-state-bias-large-kernel-prefill-b2c1024t512k128-pathology",
+      "preset": "stress",
+      "tunableSpace": { "tiledWorkgroupSize": [64, 128, 256] },
+      "provenance": {
+        "notes": "A realistic chunked continuation and regression lock: valid carry state plus bias previously excluded the tiled specialization and selected the scalar state+bias implementation."
+      },
+      "vars": { "batch": 2, "channels": 1024, "length": 512, "kernel": 128 },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "normal", "seed": 303, "scale": 0.2 },
+        "weightT": { "shape": [1024, 1, 128], "dtype": "float32", "dist": "normal", "seed": 304, "scale": 0.05 },
+        "biasT": { "shape": [1024], "dtype": "float32", "dist": "normal", "seed": 305, "scale": 0.05 },
+        "pastStateT": { "shape": [2, 1024, 127], "dtype": "float32", "dist": "normal", "seed": 306, "scale": 0.2 }
+      },
+      "outputs": {
+        "outputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "empty" },
+        "presentStateT": { "shape": [2, 1024, 127], "dtype": "float32", "dist": "empty" }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 1)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.biasT) + numel(shapes.pastStateT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-f32-huge-numel-fold-b2c4096t4096k4",
+      "preset": "stress",
+      "vars": { "batch": 2, "channels": 4096, "length": 4096, "kernel": 4 },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": { "shape": [2, 4096, 4096], "dtype": "float32", "dist": "normal", "seed": 311, "scale": 0.2 },
+        "weightT": { "shape": [4096, 1, 4], "dtype": "float32", "dist": "normal", "seed": 312, "scale": 0.1 }
+      },
+      "outputs": {
+        "outputT": { "shape": [2, 4096, 4096], "dtype": "float32", "dist": "empty" },
+        "presentStateT": { "shape": [2, 4096, 3], "dtype": "float32", "dist": "empty" }
+      },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 1)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "nemo-conv-decode-9728",
+      "preset": "smoke",
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [1, 9728, 1] },
+        "weightT": { "dtype": "float32", "shape": [9728, 1, 4] },
+        "biasT": { "dtype": "float32", "shape": [9728] },
+        "pastStateT": { "dtype": "float32", "shape": [1, 9728, 3] }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 9728, 1] },
+        "presentStateT": { "dtype": "float32", "shape": [1, 9728, 3] }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 1)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.biasT) + numel(shapes.pastStateT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-qwen3next-decode-b1-c8192-t1-k4",
+      "preset": "model",
+      "provenance": {
+        "notes": "Qwen3-Next class defaults (linear_conv_kernel_dim 4) over the concatenated q/k/v channels of the linear-attention branch (2*16*128 + 32*128 = 8192) at a decode step."
+      },
+      "vars": { "batch": 1, "channels": 8192, "length": 1, "kernel": 4 },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "shape": [1, 8192, 1], "dtype": "float32", "dist": "normal", "seed": 8700, "scale": 0.5 },
+        "weightT": { "shape": [8192, 1, 4], "dtype": "float32", "dist": "normal", "seed": 8701, "scale": 0.3 },
+        "biasT": { "shape": [8192], "dtype": "float32", "dist": "normal", "seed": 8702, "scale": 0.1 },
+        "pastStateT": { "shape": [1, 8192, 3], "dtype": "float32", "dist": "normal", "seed": 8703, "scale": 0.2 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 8192, 1], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 8192, 3], "dtype": "float32" }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 2)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.biasT) + numel(shapes.pastStateT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-qwen3next-prefill-b1-c8192-t2048-k4",
+      "preset": "model",
+      "provenance": { "notes": "Qwen3-Next class defaults over a 2048-token prefill chunk." },
+      "vars": { "batch": 1, "channels": 8192, "length": 2048, "kernel": 4 },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "shape": [1, 8192, 2048], "dtype": "float32", "dist": "normal", "seed": 8800, "scale": 0.5 },
+        "weightT": { "shape": [8192, 1, 4], "dtype": "float32", "dist": "normal", "seed": 8801, "scale": 0.3 },
+        "biasT": { "shape": [8192], "dtype": "float32", "dist": "normal", "seed": 8802, "scale": 0.1 },
+        "pastStateT": { "shape": [1, 8192, 3], "dtype": "float32", "dist": "normal", "seed": 8803, "scale": 0.2 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 8192, 2048], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 8192, 3], "dtype": "float32" }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 2)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.biasT) + numel(shapes.pastStateT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-mamba2-decode-b1-c10240-t1-k4",
+      "preset": "model",
+      "provenance": {
+        "notes": "Mamba2 class defaults (hidden_size 4096, expand 2, n_groups 8, state_size 128, conv_kernel 4): the mixer conv spans 2*4096 + 2*8*128 = 10240 channels."
+      },
+      "vars": { "batch": 1, "channels": 10240, "length": 1, "kernel": 4 },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "shape": [1, 10240, 1], "dtype": "float32", "dist": "normal", "seed": 8900, "scale": 0.5 },
+        "weightT": { "shape": [10240, 1, 4], "dtype": "float32", "dist": "normal", "seed": 8901, "scale": 0.3 },
+        "biasT": { "shape": [10240], "dtype": "float32", "dist": "normal", "seed": 8902, "scale": 0.1 },
+        "pastStateT": { "shape": [1, 10240, 3], "dtype": "float32", "dist": "normal", "seed": 8903, "scale": 0.2 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 10240, 1], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 10240, 3], "dtype": "float32" }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 2)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.biasT) + numel(shapes.pastStateT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-mamba2-prefill-b1-c10240-t2048-k4",
+      "preset": "model",
+      "provenance": { "notes": "Mamba2 class defaults over a 2048-token prefill chunk." },
+      "vars": { "batch": 1, "channels": 10240, "length": 2048, "kernel": 4 },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "shape": [1, 10240, 2048], "dtype": "float32", "dist": "normal", "seed": 9000, "scale": 0.5 },
+        "weightT": { "shape": [10240, 1, 4], "dtype": "float32", "dist": "normal", "seed": 9001, "scale": 0.3 },
+        "biasT": { "shape": [10240], "dtype": "float32", "dist": "normal", "seed": 9002, "scale": 0.1 },
+        "pastStateT": { "shape": [1, 10240, 3], "dtype": "float32", "dist": "normal", "seed": 9003, "scale": 0.2 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 10240, 2048], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 10240, 3], "dtype": "float32" }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 2)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.biasT) + numel(shapes.pastStateT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-lfm2-decode-b1-c2560-t1-k3",
+      "preset": "model",
+      "provenance": {
+        "notes": "LFM2 class defaults (hidden_size 2560, conv_L_cache 3): the short-conv block at a decode step, a kernel one tap narrower than the Mamba family."
+      },
+      "vars": { "batch": 1, "channels": 2560, "length": 1, "kernel": 3 },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "shape": [1, 2560, 1], "dtype": "float32", "dist": "normal", "seed": 9100, "scale": 0.5 },
+        "weightT": { "shape": [2560, 1, 3], "dtype": "float32", "dist": "normal", "seed": 9101, "scale": 0.3 },
+        "biasT": { "shape": [2560], "dtype": "float32", "dist": "normal", "seed": 9102, "scale": 0.1 },
+        "pastStateT": { "shape": [1, 2560, 2], "dtype": "float32", "dist": "normal", "seed": 9103, "scale": 0.2 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 2560, 1], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 2560, 2], "dtype": "float32" }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 2)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.biasT) + numel(shapes.pastStateT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "causal-conv-lfm2-prefill-b1-c2560-t2048-k3",
+      "preset": "model",
+      "provenance": { "notes": "LFM2 class defaults over a 2048-token prefill chunk." },
+      "vars": { "batch": 1, "channels": 2560, "length": 2048, "kernel": 3 },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "shape": [1, 2560, 2048], "dtype": "float32", "dist": "normal", "seed": 9200, "scale": 0.5 },
+        "weightT": { "shape": [2560, 1, 3], "dtype": "float32", "dist": "normal", "seed": 9201, "scale": 0.3 },
+        "biasT": { "shape": [2560], "dtype": "float32", "dist": "normal", "seed": 9202, "scale": 0.1 },
+        "pastStateT": { "shape": [1, 2560, 2], "dtype": "float32", "dist": "normal", "seed": 9203, "scale": 0.2 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 2560, 2048], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 2560, 2], "dtype": "float32" }
+      },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "2 * numel(shapes.outputT) * dim(shapes.weightT, 2)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.biasT) + numel(shapes.pastStateT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/causal-conv-with-state-tiled.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,174 @@

+{% set pastSlot = "(params.stateWindow - 1u) * params.stateSlotStride + " if hasStateWindow else "" %}
+{% if usesF16 %}
+enable f16;
+{% endif -%}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+const TILE: u32 = {{ tileSize }}u;
+const OUTPUTS_PER_THREAD: u32 = 8u;
+const KERNEL_SIZE: u32 = {{ kernelSize }}u;
+// The tap loop consumes four weights per iteration. Rounding the weight tile up
+// to a multiple of four and zero-filling the tail lets a kernel of any length
+// use it: the phantom taps contribute nothing. The input tile grows to match,
+// because the last lane reads TILE + KERNEL_PADDED - 2.
+const KERNEL_PADDED: u32 = {{ kernelSizePadded }}u;
+const STATE_LENGTH: u32 = KERNEL_SIZE - 1u;
+const INPUT_TILE_SIZE: u32 = TILE + KERNEL_PADDED - 1u;
+var<workgroup> input_tile: array<f32, {{ inputTileSize }}>;
+var<workgroup> weight_tile: array<f32, {{ kernelSizePadded }}>;
+fn activate(value: f32) -> f32 {
+{% if useSilu %}
+  return value / (1.0 + exp(-value));
+{% else %}
+  return value;
+{% endif %}
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(local_invocation_id) lid3: vec3<u32>,
+        @builtin(workgroup_id) wid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let lane = lid3.x;
+  let tiles_per_row = (params.length + TILE - 1u) / TILE;
+  // Recover the logical workgroup index after an oversized grid folds into y.
+  let flat_wg = wid.x + wid.y * nwg.x;
+  let total_wg = params.batchSize * params.channels * tiles_per_row;
+  if (flat_wg >= total_wg) {
+    return;
+  }
+  let tile_index = flat_wg % tiles_per_row;
+  let bc = flat_wg / tiles_per_row;
+  let channel = bc % params.channels;
+  let tile_start = tile_index * TILE;
+  var i = lane;
+  while (i < INPUT_TILE_SIZE) {
+    // input_tile[i] is the concatenated state/input sample needed at output
+    // tile_start + i and kernel tap zero.
+    let virtual_pos = tile_start + i;
+    var value = 0.0;
+    if (virtual_pos >= STATE_LENGTH) {
+      let input_pos = virtual_pos - STATE_LENGTH;
+      if (input_pos < params.length) {
+        value = f32(input[bc * params.length + input_pos]);
+      }
+    } else {
+{% if hasState %}
+      value = f32(past_state[{{ pastSlot }}bc * STATE_LENGTH + virtual_pos]);
+{% endif %}
+    }
+    input_tile[i] = value;
+    i = i + WG;
+  }
+  i = lane;
+  while (i < KERNEL_PADDED) {
+    var weight_value = 0.0;
+    if (i < KERNEL_SIZE) {
+      weight_value = f32(weight[channel * KERNEL_SIZE + i]);
+    }
+    weight_tile[i] = weight_value;
+    i = i + WG;
+  }
+  workgroupBarrier();
+  let pos = tile_start + lane * OUTPUTS_PER_THREAD;
+  if (pos < params.length) {
+{% if hasBias %}
+    var acc0 = vec4<f32>(f32(bias[channel]));
+    var acc1 = acc0;
+{% else %}
+    var acc0 = vec4<f32>(0.0);
+    var acc1 = vec4<f32>(0.0);
+{% endif %}
+    // Each lane computes eight adjacent outputs and reuses a weight across them.
+    var k = 0u;
+    while (k < KERNEL_PADDED) {
+      let w = vec4<f32>(
+        weight_tile[k],
+        weight_tile[k + 1u],
+        weight_tile[k + 2u],
+        weight_tile[k + 3u]
+      );
+      let base0 = lane * OUTPUTS_PER_THREAD + k;
+      let base1 = base0 + 4u;
+      let x00 = vec4<f32>(input_tile[base0], input_tile[base0 + 1u], input_tile[base0 + 2u], input_tile[base0 + 3u]);
+      let x01 = vec4<f32>(input_tile[base0 + 1u], input_tile[base0 + 2u], input_tile[base0 + 3u], input_tile[base0 + 4u]);
+      let x02 = vec4<f32>(input_tile[base0 + 2u], input_tile[base0 + 3u], input_tile[base0 + 4u], input_tile[base0 + 5u]);
+      let x03 = vec4<f32>(input_tile[base0 + 3u], input_tile[base0 + 4u], input_tile[base0 + 5u], input_tile[base0 + 6u]);
+      let x10 = vec4<f32>(input_tile[base1], input_tile[base1 + 1u], input_tile[base1 + 2u], input_tile[base1 + 3u]);
+      let x11 = vec4<f32>(input_tile[base1 + 1u], input_tile[base1 + 2u], input_tile[base1 + 3u], input_tile[base1 + 4u]);
+      let x12 = vec4<f32>(input_tile[base1 + 2u], input_tile[base1 + 3u], input_tile[base1 + 4u], input_tile[base1 + 5u]);
+      let x13 = vec4<f32>(input_tile[base1 + 3u], input_tile[base1 + 4u], input_tile[base1 + 5u], input_tile[base1 + 6u]);
+      acc0 = fma(x00, vec4<f32>(w.x), acc0);
+      acc0 = fma(x01, vec4<f32>(w.y), acc0);
+      acc0 = fma(x02, vec4<f32>(w.z), acc0);
+      acc0 = fma(x03, vec4<f32>(w.w), acc0);
+      acc1 = fma(x10, vec4<f32>(w.x), acc1);
+      acc1 = fma(x11, vec4<f32>(w.y), acc1);
+      acc1 = fma(x12, vec4<f32>(w.z), acc1);
+      acc1 = fma(x13, vec4<f32>(w.w), acc1);
+      k = k + 4u;
+    }
+    let output_base = bc * params.length + pos;
+    output[output_base] = {{ outputScalar }}(activate(acc0.x));
+    output[output_base + 1u] = {{ outputScalar }}(activate(acc0.y));
+    output[output_base + 2u] = {{ outputScalar }}(activate(acc0.z));
+    output[output_base + 3u] = {{ outputScalar }}(activate(acc0.w));
+    output[output_base + 4u] = {{ outputScalar }}(activate(acc1.x));
+    output[output_base + 5u] = {{ outputScalar }}(activate(acc1.y));
+    output[output_base + 6u] = {{ outputScalar }}(activate(acc1.z));
+    output[output_base + 7u] = {{ outputScalar }}(activate(acc1.w));
+  }
+  // Exactly one tile per (batch, channel) updates the carry state, with all
+  // lanes sharing the copy.
+  if (tile_index == 0u) {
+{% if hasStateWindow %}
+    // Windowed state slot j follows position (length - stateWindow + j), so the
+    // final slot is the after-last state. A slot whose prefix
+    // ends before position 0 holds no position from this call and stays zero. Lanes stride over
+    // the flattened (slot, element) grid so a narrow workgroup still covers every slot.
+    var e = lane;
+    let state_elems = params.stateWindow * STATE_LENGTH;
+    while (e < state_elems) {
+      let slot = e / STATE_LENGTH;
+      let s = e % STATE_LENGTH;
+      var state_value = 0.0;
+      if (slot + params.length >= params.stateWindow) {
+        let virtual_pos = params.length + slot + 1u - params.stateWindow + s;
+        if (virtual_pos >= STATE_LENGTH) {
+          let input_pos = virtual_pos - STATE_LENGTH;
+          state_value = f32(input[bc * params.length + input_pos]);
+        } else {
+{% if hasState %}
+          state_value = f32(past_state[{{ pastSlot }}bc * STATE_LENGTH + virtual_pos]);
+{% endif %}
+        }
+      }
+      present_state[slot * params.stateSlotStride + bc * STATE_LENGTH + s] = {{ outputScalar }}(state_value);
+      e = e + WG;
+    }
+{% else %}
+    var s = lane;
+    while (s < STATE_LENGTH) {
+      let virtual_pos = params.length + s;
+      var state_value = 0.0;
+      if (virtual_pos >= STATE_LENGTH) {
+        let input_pos = virtual_pos - STATE_LENGTH;
+        state_value = f32(input[bc * params.length + input_pos]);
+      } else {
+{% if hasState %}
+        state_value = f32(past_state[bc * STATE_LENGTH + virtual_pos]);
+{% endif %}
+      }
+      present_state[bc * STATE_LENGTH + s] = {{ outputScalar }}(state_value);
+      s = s + WG;
+    }
+{% endif %}
+  }
+}

build/webgpu/causal-conv-with-state-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,109 @@

+{% set pastSlot = "(params.stateWindow - 1u) * params.stateSlotStride + " if hasStateWindow else "" %}
+{% if usesF16 %}
+enable f16;
+{% endif -%}
+{{ env.wgsl.resourceDeclarations }}
+// Four outputs per invocation for a kernel narrow enough that every tap of a
+// four-wide output lands in this vector or the one before it -- kernel size 2 to
+// 4. The taps are unrolled swizzles of two loaded vectors rather than a
+// runtime-bounded loop.
+const WG: u32 = {{ workgroupSize }}u;
+{% set S = kernelSize - 1 %}
+const STATE_LENGTH: u32 = {{ S }}u;
+fn activate4(value: vec4<f32>) -> vec4<f32> {
+{% if useSilu %}
+  return value / (vec4<f32>(1.0) + exp(-value));
+{% else %}
+  return value;
+{% endif %}
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let row_vecs = params.length / 4u;
+  let work_size = params.batchSize * params.channels * row_vecs;
+  // Recover the logical 1D index after an oversized dispatch is folded into y.
+  let index = gid.x + gid.y * nwg.x * WG;
+  if (index >= work_size) {
+    return;
+  }
+  let pos4 = index % row_vecs;
+  let bc = index / row_vecs;
+  let channel = bc % params.channels;
+  // Widen on load and narrow on store so every tap and accumulation stays f32,
+  // independent of the tensor type.
+  let current = vec4<f32>(input[index]);
+  var previous = vec4<f32>(0.0);
+  if (pos4 != 0u) {
+    previous = vec4<f32>(input[index - 1u]);
+{% if hasState %}
+  } else {
+    // The first vector of a row continues the previous call: its taps come from
+    // the carried state, in the same trailing lanes the interior case reads.
+    let sb = {{ pastSlot }}bc * STATE_LENGTH;
+    previous = vec4<f32>({% for i in range(4 - S) %}0.0, {% endfor %}{% for i in range(S) %}f32(past_state[sb{{ " + " ~ i ~ "u" if i else "" }}]){{ ", " if not loop.last else "" }}{% endfor %});
+{% endif %}
+  }
+{% if kernelSize == 4 %}
+  let w = vec4<f32>(weight[channel]);
+{% else %}
+{% for k in range(kernelSize) %}
+  let w{{ k }} = f32(weight[channel * {{ kernelSize }}u + {{ k }}u]);
+{% endfor %}
+{% endif %}
+{% for k in range(kernelSize) %}
+{% set wk = ("w." ~ ["x", "y", "z", "w"][k]) if kernelSize == 4 else ("w" ~ k) %}
+{% if loop.first and hasBias %}
+  var value = vec4<f32>(f32(bias[channel])) + {{ wk }} * vec4<f32>(
+{%- for j in range(4) %}{{ ("current." ~ ["x", "y", "z", "w"][j - S + k]) if j - S + k >= 0 else ("previous." ~ ["x", "y", "z", "w"][4 + j - S + k]) }}{{ ", " if not loop.last else "" }}{% endfor -%}
+);
+{% elif loop.first %}
+  var value = {{ wk }} * vec4<f32>(
+{%- for j in range(4) %}{{ ("current." ~ ["x", "y", "z", "w"][j - S + k]) if j - S + k >= 0 else ("previous." ~ ["x", "y", "z", "w"][4 + j - S + k]) }}{{ ", " if not loop.last else "" }}{% endfor -%}
+);
+{% else %}
+  value = value + {{ wk }} * vec4<f32>(
+{%- for j in range(4) %}{{ ("current." ~ ["x", "y", "z", "w"][j - S + k]) if j - S + k >= 0 else ("previous." ~ ["x", "y", "z", "w"][4 + j - S + k]) }}{{ ", " if not loop.last else "" }}{% endfor -%}
+);
+{% endif %}
+{% endfor %}
+  output[index] = {{ outputVec4 }}(activate4(value));
+  if (pos4 == 0u) {
+{% if hasStateWindow %}
+    // Windowed state: slot j holds the carry state after position (length - stateWindow + j).
+    // A slot whose prefix ends before position 0 holds no position from this call and stays
+    // zero; a slot that reaches back before position 0 continues the carried state.
+    // present_state is scalar-typed here while input is vec4, hence the lane split.
+    for (var slot = 0u; slot < params.stateWindow; slot = slot + 1u) {
+      let in_window = slot + params.length >= params.stateWindow;
+      for (var s = 0u; s < STATE_LENGTH; s = s + 1u) {
+        var state_value = 0.0;
+        if (in_window) {
+          let virtual_pos = params.length + slot + 1u - params.stateWindow + s;
+          if (virtual_pos >= STATE_LENGTH) {
+            let input_pos = virtual_pos - STATE_LENGTH;
+            state_value = f32(input[bc * row_vecs + input_pos / 4u][input_pos % 4u]);
+          } else {
+{% if hasState %}
+            state_value = f32(past_state[{{ pastSlot }}bc * STATE_LENGTH + virtual_pos]);
+{% endif %}
+          }
+        }
+        present_state[slot * params.stateSlotStride + bc * STATE_LENGTH + s] = {{ outputScalar }}(state_value);
+      }
+    }
+{% else %}
+    let tail = vec4<f32>(input[bc * row_vecs + row_vecs - 1u]);
+    let state_base = bc * STATE_LENGTH;
+{% for i in range(S) %}
+    present_state[state_base{{ " + " ~ i ~ "u" if i else "" }}] = {{ outputScalar }}(tail.{{ "xyzw"[4 - S + i] }});
+{% endfor %}
+{% endif %}
+  }
+}

build/webgpu/causal-conv-with-state.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,105 @@

+{% set pastSlot = "(params.stateWindow - 1u) * params.stateSlotStride + " if hasStateWindow else "" %}
+{% if usesF16 %}
+enable f16;
+{% endif -%}
+{{ env.wgsl.resourceDeclarations }}
+const WG: u32 = {{ workgroupSize }}u;
+fn activate(value: f32) -> f32 {
+{% if useSilu %}
+  return value / (1.0 + exp(-value));
+{% else %}
+  return value;
+{% endif %}
+}
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  // length == 0 still writes the present_state carryover. The nonzero grid keeps
+  // one state-writing thread per (batch, channel) even when there is no output.
+  let len_nz = max(1u, params.length);
+  let work_size = params.batchSize * params.channels * len_nz;
+  // gid.y carries the high bits past the per-dimension dispatch limit.
+  let index = gid.x + gid.y * nwg.x * WG;
+  if (index >= work_size) {
+    return;
+  }
+  let pos = index % len_nz;
+  let bc = index / len_nz;
+  let batch = bc / params.channels;
+  let channel = bc % params.channels;
+  let state_length = params.kernelSize - 1u;
+  // Output exists only for real positions. At length zero this thread performs
+  // only the present_state carryover below.
+  if (pos < params.length) {
+    var acc = 0.0;
+{% if hasBias %}
+    acc = acc + f32(bias[channel]);
+{% endif %}
+    for (var k: u32 = 0u; k < params.kernelSize; k = k + 1u) {
+      var value = 0.0;
+      let virtual_pos = pos + k;
+      if (virtual_pos >= state_length) {
+        let input_pos = virtual_pos - state_length;
+        let input_index = (batch * params.channels + channel) * params.length + input_pos;
+        value = f32(input[input_index]);
+      } else {
+{% if hasState %}
+        let state_index = {{ pastSlot }}(batch * params.channels + channel) * state_length + virtual_pos;
+        value = f32(past_state[state_index]);
+{% endif %}
+      }
+      acc = acc + value * f32(weight[channel * params.kernelSize + k]);
+    }
+    output[index] = {{ outputScalar }}(activate(acc));
+  }
+  if (pos == 0u) {
+{% if hasStateWindow %}
+    // Windowed state slot j follows position (length - stateWindow + j), so the
+    // final slot is the after-last state. A slot whose prefix
+    // ends before position 0 holds no position from this call and stays zero.
+    for (var slot: u32 = 0u; slot < params.stateWindow; slot = slot + 1u) {
+      let in_window = slot + params.length >= params.stateWindow;
+      for (var s: u32 = 0u; s < state_length; s = s + 1u) {
+        var state_value = 0.0;
+        if (in_window) {
+          let virtual_pos = params.length + slot + 1u - params.stateWindow + s;
+          if (virtual_pos >= state_length) {
+            let input_pos = virtual_pos - state_length;
+            let input_index = (batch * params.channels + channel) * params.length + input_pos;
+            state_value = f32(input[input_index]);
+          } else {
+{% if hasState %}
+            let state_index = {{ pastSlot }}(batch * params.channels + channel) * state_length + virtual_pos;
+            state_value = f32(past_state[state_index]);
+{% endif %}
+          }
+        }
+        present_state[slot * params.stateSlotStride + (batch * params.channels + channel) * state_length + s] = {{ outputScalar }}(state_value);
+      }
+    }
+{% else %}
+    for (var s: u32 = 0u; s < state_length; s = s + 1u) {
+      var state_value = 0.0;
+      let virtual_pos = params.length + s;
+      if (virtual_pos >= state_length) {
+        let input_pos = virtual_pos - state_length;
+        let input_index = (batch * params.channels + channel) * params.length + input_pos;
+        state_value = f32(input[input_index]);
+      } else {
+{% if hasState %}
+        let state_index = (batch * params.channels + channel) * state_length + virtual_pos;
+        state_value = f32(past_state[state_index]);
+{% endif %}
+      }
+      present_state[(batch * params.channels + channel) * state_length + s] = {{ outputScalar }}(state_value);
+    }
+{% endif %}
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,1257 @@

+{
+  "domain": "com.microsoft",
+  "name": "CausalConvWithState",
+  "sinceVersion": 1,
+  "description": "Legacy Microsoft contrib form of stateful 1-D causal depthwise convolution. Each channel uses its own `(channels, 1, kernel)` weight over current and past positions, with optional activation and `past_state`/`present_state` tensors for incremental decoding. The contrib-only `state_window` attribute may retain several rollback states. This inference implementation preserves the existing contrib ABI with `ndim = 1`, float16 or float32 tensors, and float32 accumulation; spatial ranks 2 and 3 and bfloat16 are not implemented.",
+  "inputs": [
+    {
+      "role": "input",
+      "dtype": "T",
+      "rank": 3,
+      "description": "Channels-first input tensor with shape `(batch_size, channels, sequence_length)` for the supported 1-D mode."
+    },
+    {
+      "role": "weight",
+      "dtype": "T",
+      "rank": 3,
+      "description": "Depthwise convolution kernel with shape `(channels, 1, kernel_size)` for the supported 1-D mode."
+    },
+    {
+      "role": "bias",
+      "dtype": "T",
+      "rank": 1,
+      "optional": true,
+      "description": "Optional per-channel bias with shape `(channels,)`."
+    },
+    {
+      "role": "past_state",
+      "dtype": "T",
+      "rank": "3 if attrs.state_window == 0 else 4",
+      "optional": true,
+      "description": "Carry state from the previous step; shape `(batch_size, channels, k_1 - 1)`, or `(W, batch_size, channels, k_1 - 1)` when `state_window = W > 0`, in which case only slot `W - 1` is read. If absent, the left-side padding is zero."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": 3,
+      "shape": "shapes.input",
+      "description": "Convolution output with the same shape as `input`."
+    },
+    {
+      "role": "present_state",
+      "dtype": "T",
+      "rank": "3 if attrs.state_window == 0 else 4",
+      "shape": "[dim(shapes.input, 0), dim(shapes.input, 1), dim(shapes.weight, 2) - 1] if attrs.state_window == 0 else [attrs.state_window, dim(shapes.input, 0), dim(shapes.input, 1), dim(shapes.weight, 2) - 1]",
+      "description": "Updated carry state; shape `(batch_size, channels, k_1 - 1)`, or `(W, batch_size, channels, k_1 - 1)` when `state_window = W > 0`. Slot `W - 1` holds the last `k - 1` values along the causal axis; slot `j` holds the same for the prefix ending at position `seq_len - W + j`."
+    }
+  ],
+  "attributes": { "activation": "none", "ndim": 1, "state_window": 0 },
+  "attributeConstraints": { "activation": { "values": ["none", "silu", "swish"] }, "ndim": { "values": [1] } },
+  "attributeDescriptions": {
+    "activation": "Activation applied after convolution and bias. Defaults to `none`; `swish` is an alias of SiLU.",
+    "ndim": "Number of spatial dimensions. This implementation supports the contrib 1D mode (`ndim = 1`).",
+    "state_window": "Contrib extension selecting the number of rollback state slots to retain, in the range 0 through 8. Defaults to 0."
+  },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "inputT": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "weightT": { "kind": "tensor", "semantic": "weight", "role": "input" },
+    "biasT": { "kind": "tensor", "semantic": "bias", "role": "input", "required": false },
+    "pastStateT": { "kind": "tensor", "semantic": "past_state", "role": "input", "required": false },
+    "outputT": { "kind": "tensor", "semantic": "output", "role": "output" },
+    "presentStateT": { "kind": "tensor", "semantic": "present_state", "role": "output" }
+  },
+  "tunables": { "workgroupSize": 256, "tiledWorkgroupSize": 128 },
+  "derive": {
+    "stateWindow": "attrs.state_window",
+    "windowed": "stateWindow > 0",
+    "stateWindowOk": "stateWindow >= 0 and stateWindow <= 8",
+    "kernelSize": "dim(shapes.weightT, ranks.weightT - 1)",
+    "kernelSizePadded": "ceilDiv(kernelSize, 4) * 4",
+    "weightRankOk": "ranks.weightT == 3 and dim(shapes.weightT, 1) == 1",
+    "stateLength": "kernelSize - 1",
+    "stateSlotStride": "dim(shapes.inputT, 0) * dim(shapes.inputT, 1) * stateLength",
+    "windowedLengthOk": "not windowed or dim(shapes.inputT, 2) > 0",
+    "presentStateOk": "(ranks.presentStateT == 3 and dim(shapes.presentStateT, 0) == dim(shapes.inputT, 0) and dim(shapes.presentStateT, 1) == dim(shapes.inputT, 1) and dim(shapes.presentStateT, 2) == stateLength) if not windowed else (ranks.presentStateT == 4 and dim(shapes.presentStateT, 0) == stateWindow and dim(shapes.presentStateT, 1) == dim(shapes.inputT, 0) and dim(shapes.presentStateT, 2) == dim(shapes.inputT, 1) and dim(shapes.presentStateT, 3) == stateLength)",
+    "pastStateShapeOk": "present.pastStateT and ((ranks.pastStateT == 3 and dim(shapes.pastStateT, 0) == dim(shapes.inputT, 0) and dim(shapes.pastStateT, 1) == dim(shapes.inputT, 1) and dim(shapes.pastStateT, 2) == stateLength) if not windowed else (ranks.pastStateT == 4 and dim(shapes.pastStateT, 0) == stateWindow and dim(shapes.pastStateT, 1) == dim(shapes.inputT, 0) and dim(shapes.pastStateT, 2) == dim(shapes.inputT, 1) and dim(shapes.pastStateT, 3) == stateLength))",
+    "commonContract": "ranks.inputT == 3 and weightRankOk and ranks.outputT == 3 and (tensorDtypes.inputT == \"float32\" or tensorDtypes.inputT == \"float16\") and tensorDtypes.weightT == tensorDtypes.inputT and tensorDtypes.outputT == tensorDtypes.inputT and tensorDtypes.presentStateT == tensorDtypes.inputT and f16Ok(dtypes.T) and dim(shapes.inputT, 1) == dim(shapes.weightT, 0) and dim(shapes.outputT, 0) == dim(shapes.inputT, 0) and dim(shapes.outputT, 1) == dim(shapes.inputT, 1) and dim(shapes.outputT, 2) == dim(shapes.inputT, 2) and stateWindowOk and windowedLengthOk and presentStateOk",
+    "zeroStateContract": "commonContract and not present.pastStateT and not present.biasT",
+    "biasNoStateContract": "commonContract and not present.pastStateT and present.biasT and ranks.biasT == 1 and tensorDtypes.biasT == tensorDtypes.inputT and dim(shapes.biasT, 0) == dim(shapes.inputT, 1)",
+    "stateNoBiasContract": "commonContract and present.pastStateT and not present.biasT and tensorDtypes.pastStateT == tensorDtypes.inputT and pastStateShapeOk",
+    "stateBiasContract": "commonContract and present.pastStateT and present.biasT and ranks.biasT == 1 and tensorDtypes.pastStateT == tensorDtypes.inputT and tensorDtypes.biasT == tensorDtypes.inputT and pastStateShapeOk and dim(shapes.biasT, 0) == dim(shapes.inputT, 1)"
+  },
+  "bindingSets": {
+    "zeroScalar": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "kernelSize", "type": "u32", "value": "kernelSize" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "zeroVec4": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputVec4"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$weightElem"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputVec4"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "biasNoState": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "kernelSize", "type": "u32", "value": "kernelSize" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "stateNoBias": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "past_state",
+        "arg": "pastStateT",
+        "semantic": "past_state",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "kernelSize", "type": "u32", "value": "kernelSize" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "stateBias": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "past_state",
+        "arg": "pastStateT",
+        "semantic": "past_state",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "kernelSize", "type": "u32", "value": "kernelSize" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "zeroScalarIo": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      }
+    ],
+    "biasNoStateIo": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      }
+    ],
+    "stateNoBiasIo": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "past_state",
+        "arg": "pastStateT",
+        "semantic": "past_state",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      }
+    ],
+    "stateBiasIo": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "past_state",
+        "arg": "pastStateT",
+        "semantic": "past_state",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      }
+    ],
+    "zeroTiled": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "biasNoStateTiled": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "stateNoBiasTiled": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "past_state",
+        "arg": "pastStateT",
+        "semantic": "past_state",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "stateBiasTiled": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "past_state",
+        "arg": "pastStateT",
+        "semantic": "past_state",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "biasNoStateVec4": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputVec4"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$weightElem"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputVec4"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "stateNoBiasVec4": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputVec4"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$weightElem"
+      },
+      {
+        "name": "past_state",
+        "arg": "pastStateT",
+        "semantic": "past_state",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputVec4"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ],
+    "stateBiasVec4": [
+      {
+        "name": "input",
+        "arg": "inputT",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputVec4"
+      },
+      {
+        "name": "weight",
+        "arg": "weightT",
+        "semantic": "weight",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$weightElem"
+      },
+      {
+        "name": "bias",
+        "arg": "biasT",
+        "semantic": "bias",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "past_state",
+        "arg": "pastStateT",
+        "semantic": "past_state",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inputScalar"
+      },
+      {
+        "name": "output",
+        "arg": "outputT",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputVec4"
+      },
+      {
+        "name": "present_state",
+        "arg": "presentStateT",
+        "semantic": "present_state",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "batchSize", "type": "u32", "value": "dim(shapes.inputT, 0)" },
+            { "name": "channels", "type": "u32", "value": "dim(shapes.inputT, 1)" },
+            { "name": "length", "type": "u32", "value": "dim(shapes.inputT, 2)" },
+            { "name": "stateWindow", "type": "u32", "value": "stateWindow" },
+            { "name": "stateSlotStride", "type": "u32", "value": "stateSlotStride" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "zero_state_vec4",
+      "priority": 20,
+      "when": ["zeroStateContract", "kernelSize >= 2", "kernelSize <= 4", "dim(shapes.inputT, 2) >= 4", "dim(shapes.inputT, 2) % 4 == 0"],
+      "constants": {
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "workgroupSize": 256,
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "inputVec4": "\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\"",
+        "outputVec4": "\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\"",
+        "hasBias": false,
+        "hasState": false,
+        "kernelSize": "kernelSize",
+        "kernelSizePadded": "kernelSizePadded",
+        "weightElem": "(\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\") if kernelSize == 4 else dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.Vec4",
+          "source": {
+            "shader": "causal-conv-with-state-vec4.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "zeroVec4",
+          "dispatch": {
+            "threads": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * (dim(shapes.outputT, 2) / 4)",
+            "workgroupSize": "constants.workgroupSize"
+          }
+        }
+      ]
+    },
+    {
+      "id": "state_bias_vec4",
+      "priority": 20,
+      "when": ["stateBiasContract", "kernelSize >= 2", "kernelSize <= 4", "dim(shapes.inputT, 2) >= 4", "dim(shapes.inputT, 2) % 4 == 0"],
+      "constants": {
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "workgroupSize": 256,
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "inputVec4": "\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\"",
+        "outputVec4": "\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\"",
+        "hasBias": true,
+        "hasState": true,
+        "kernelSize": "kernelSize",
+        "kernelSizePadded": "kernelSizePadded",
+        "weightElem": "(\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\") if kernelSize == 4 else dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.Vec4",
+          "source": {
+            "shader": "causal-conv-with-state-vec4.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "stateBiasVec4",
+          "dispatch": {
+            "threads": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * (dim(shapes.outputT, 2) / 4)",
+            "workgroupSize": "constants.workgroupSize"
+          }
+        }
+      ]
+    },
+    {
+      "id": "bias_no_state_vec4",
+      "priority": 20,
+      "when": ["biasNoStateContract", "kernelSize >= 2", "kernelSize <= 4", "dim(shapes.inputT, 2) >= 4", "dim(shapes.inputT, 2) % 4 == 0"],
+      "constants": {
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "workgroupSize": 256,
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "inputVec4": "\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\"",
+        "outputVec4": "\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\"",
+        "hasBias": true,
+        "hasState": false,
+        "kernelSize": "kernelSize",
+        "kernelSizePadded": "kernelSizePadded",
+        "weightElem": "(\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\") if kernelSize == 4 else dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.Vec4",
+          "source": {
+            "shader": "causal-conv-with-state-vec4.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "biasNoStateVec4",
+          "dispatch": {
+            "threads": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * (dim(shapes.outputT, 2) / 4)",
+            "workgroupSize": "constants.workgroupSize"
+          }
+        }
+      ]
+    },
+    {
+      "id": "state_no_bias_vec4",
+      "priority": 20,
+      "when": ["stateNoBiasContract", "kernelSize >= 2", "kernelSize <= 4", "dim(shapes.inputT, 2) >= 4", "dim(shapes.inputT, 2) % 4 == 0"],
+      "constants": {
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "workgroupSize": 256,
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "inputVec4": "\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\"",
+        "outputVec4": "\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\"",
+        "hasBias": false,
+        "hasState": true,
+        "kernelSize": "kernelSize",
+        "kernelSizePadded": "kernelSizePadded",
+        "weightElem": "(\"vec4<f16>\" if tensorDtypes.inputT == \"float16\" else \"vec4<f32>\") if kernelSize == 4 else dtypes.T"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.Vec4",
+          "source": {
+            "shader": "causal-conv-with-state-vec4.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "stateNoBiasVec4",
+          "dispatch": {
+            "threads": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * (dim(shapes.outputT, 2) / 4)",
+            "workgroupSize": "constants.workgroupSize"
+          }
+        }
+      ]
+    },
+    {
+      "id": "zero_state_tiled_large_kernel",
+      "priority": 10,
+      "when": ["zeroStateContract", "kernelSize >= 32", "dim(shapes.inputT, 2) >= 256", "dim(shapes.inputT, 2) % 8 == 0", "tunables.tiledWorkgroupSize >= 1", "floor(tunables.tiledWorkgroupSize) == tunables.tiledWorkgroupSize", "tunables.tiledWorkgroupSize <= device.limits.maxComputeInvocationsPerWorkgroup", "tunables.tiledWorkgroupSize <= device.limits.maxComputeWorkgroupSizeX", "(tunables.tiledWorkgroupSize * 8 + 2 * kernelSizePadded - 1) * 4 <= device.limits.maxComputeWorkgroupStorageSize"],
+      "constants": {
+        "hasBias": false,
+        "hasState": false,
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "workgroupSize": "tunables.tiledWorkgroupSize",
+        "tileSize": "tunables.tiledWorkgroupSize * 8",
+        "kernelSize": "kernelSize",
+        "kernelSizePadded": "kernelSizePadded",
+        "inputTileSize": "tunables.tiledWorkgroupSize * 8 + kernelSizePadded - 1",
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.TiledLargeKernel",
+          "source": {
+            "shader": "causal-conv-with-state-tiled.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "zeroTiled",
+          "dispatch": {
+            "workgroups": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * ceilDiv(dim(shapes.outputT, 2), constants.tileSize)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "state_bias_tiled_large_kernel",
+      "priority": 10,
+      "when": ["stateBiasContract", "kernelSize >= 32", "dim(shapes.inputT, 2) >= 256", "dim(shapes.inputT, 2) % 8 == 0", "tunables.tiledWorkgroupSize >= 1", "floor(tunables.tiledWorkgroupSize) == tunables.tiledWorkgroupSize", "tunables.tiledWorkgroupSize <= device.limits.maxComputeInvocationsPerWorkgroup", "tunables.tiledWorkgroupSize <= device.limits.maxComputeWorkgroupSizeX", "(tunables.tiledWorkgroupSize * 8 + 2 * kernelSizePadded - 1) * 4 <= device.limits.maxComputeWorkgroupStorageSize"],
+      "constants": {
+        "hasBias": true,
+        "hasState": true,
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "workgroupSize": "tunables.tiledWorkgroupSize",
+        "tileSize": "tunables.tiledWorkgroupSize * 8",
+        "kernelSize": "kernelSize",
+        "kernelSizePadded": "kernelSizePadded",
+        "inputTileSize": "tunables.tiledWorkgroupSize * 8 + kernelSizePadded - 1",
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.TiledLargeKernel",
+          "source": {
+            "shader": "causal-conv-with-state-tiled.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "stateBiasTiled",
+          "dispatch": {
+            "workgroups": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * ceilDiv(dim(shapes.outputT, 2), constants.tileSize)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "bias_no_state_tiled_large_kernel",
+      "priority": 10,
+      "when": ["biasNoStateContract", "kernelSize >= 32", "dim(shapes.inputT, 2) >= 256", "dim(shapes.inputT, 2) % 8 == 0", "tunables.tiledWorkgroupSize >= 1", "floor(tunables.tiledWorkgroupSize) == tunables.tiledWorkgroupSize", "tunables.tiledWorkgroupSize <= device.limits.maxComputeInvocationsPerWorkgroup", "tunables.tiledWorkgroupSize <= device.limits.maxComputeWorkgroupSizeX", "(tunables.tiledWorkgroupSize * 8 + 2 * kernelSizePadded - 1) * 4 <= device.limits.maxComputeWorkgroupStorageSize"],
+      "constants": {
+        "hasBias": true,
+        "hasState": false,
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "workgroupSize": "tunables.tiledWorkgroupSize",
+        "tileSize": "tunables.tiledWorkgroupSize * 8",
+        "kernelSize": "kernelSize",
+        "kernelSizePadded": "kernelSizePadded",
+        "inputTileSize": "tunables.tiledWorkgroupSize * 8 + kernelSizePadded - 1",
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.TiledLargeKernel",
+          "source": {
+            "shader": "causal-conv-with-state-tiled.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "biasNoStateTiled",
+          "dispatch": {
+            "workgroups": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * ceilDiv(dim(shapes.outputT, 2), constants.tileSize)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "state_no_bias_tiled_large_kernel",
+      "priority": 10,
+      "when": ["stateNoBiasContract", "kernelSize >= 32", "dim(shapes.inputT, 2) >= 256", "dim(shapes.inputT, 2) % 8 == 0", "tunables.tiledWorkgroupSize >= 1", "floor(tunables.tiledWorkgroupSize) == tunables.tiledWorkgroupSize", "tunables.tiledWorkgroupSize <= device.limits.maxComputeInvocationsPerWorkgroup", "tunables.tiledWorkgroupSize <= device.limits.maxComputeWorkgroupSizeX", "(tunables.tiledWorkgroupSize * 8 + 2 * kernelSizePadded - 1) * 4 <= device.limits.maxComputeWorkgroupStorageSize"],
+      "constants": {
+        "hasBias": false,
+        "hasState": true,
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "workgroupSize": "tunables.tiledWorkgroupSize",
+        "tileSize": "tunables.tiledWorkgroupSize * 8",
+        "kernelSize": "kernelSize",
+        "kernelSizePadded": "kernelSizePadded",
+        "inputTileSize": "tunables.tiledWorkgroupSize * 8 + kernelSizePadded - 1",
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.TiledLargeKernel",
+          "source": {
+            "shader": "causal-conv-with-state-tiled.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "stateNoBiasTiled",
+          "dispatch": {
+            "workgroups": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * ceilDiv(dim(shapes.outputT, 2), constants.tileSize)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "zero_state",
+      "priority": 0,
+      "when": ["zeroStateContract"],
+      "constants": {
+        "hasBias": false,
+        "hasState": false,
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "workgroupSize": "tunables.workgroupSize",
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState",
+          "source": {
+            "shader": "causal-conv-with-state.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "zeroScalar",
+          "dispatch": {
+            "threads": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * max(1, dim(shapes.outputT, 2))",
+            "workgroupSize": "constants.workgroupSize"
+          }
+        }
+      ]
+    },
+    {
+      "id": "state_bias",
+      "priority": 0,
+      "when": ["stateBiasContract"],
+      "constants": {
+        "hasBias": true,
+        "hasState": true,
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "workgroupSize": "tunables.workgroupSize",
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState",
+          "source": {
+            "shader": "causal-conv-with-state.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "stateBias",
+          "dispatch": {
+            "threads": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * max(1, dim(shapes.outputT, 2))",
+            "workgroupSize": "constants.workgroupSize"
+          }
+        }
+      ]
+    },
+    {
+      "id": "bias_no_state",
+      "priority": 0,
+      "when": ["biasNoStateContract"],
+      "constants": {
+        "hasBias": true,
+        "hasState": false,
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "workgroupSize": "tunables.workgroupSize",
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState",
+          "source": {
+            "shader": "causal-conv-with-state.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "biasNoState",
+          "dispatch": {
+            "threads": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * max(1, dim(shapes.outputT, 2))",
+            "workgroupSize": "constants.workgroupSize"
+          }
+        }
+      ]
+    },
+    {
+      "id": "state_no_bias",
+      "priority": 0,
+      "when": ["stateNoBiasContract"],
+      "constants": {
+        "hasBias": false,
+        "hasState": true,
+        "useSilu": "attrs.activation == \"silu\" or attrs.activation == \"swish\"",
+        "inputScalar": "dtypes.T",
+        "outputScalar": "dtypes.T",
+        "workgroupSize": "tunables.workgroupSize",
+        "hasStateWindow": "windowed",
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState",
+          "source": {
+            "shader": "causal-conv-with-state.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": false }
+          },
+          "bindings": "stateNoBias",
+          "dispatch": {
+            "threads": "dim(shapes.outputT, 0) * dim(shapes.outputT, 1) * max(1, dim(shapes.outputT, 2))",
+            "workgroupSize": "constants.workgroupSize"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,20 @@

+{
+  "name": "com.microsoft.CausalConvWithState",
+  "id": "_com_microsoft_causalconvwithstate_webgpu_9e6e59f",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "JZU8nd2+4ByxWEyHozqDzF2xGDNwA4gaZ1jk71H/kF0=",
+      "causal-conv-with-state-tiled.wgsl.jinja": "3AcQURBK9/4LFE/TCfWxDMIAMsjLIVcfsZ9dK/qlzGk=",
+      "causal-conv-with-state-vec4.wgsl.jinja": "jjKAFL4nTGXXrRLF+TyBiJO4fNy0bEF1gK0KfIGeRoc=",
+      "causal-conv-with-state.wgsl.jinja": "B03ROsgmw6YXaTQ6373ibBbKyt7lxTBThsR7ThfyXVU=",
+      "manifest.json": "sR2aWi0R4vk7io6t9uwLdcVyP8jdQ5PIA0QHMHKQkaU=",
+      "test.json": "oC4aXSiNHpXHdk6nE+nXhAZvkGulY8eJOcfiXZ0jXdw="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.CausalConvWithState" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1283 @@

+{
+  "op": "com.microsoft.CausalConvWithState",
+  "cases": [
+    {
+      "name": "ort_kernel1_zero_size_state",
+      "provenance": {
+        "source": "onnxruntime/test/python/transformers/test_parity_linear_attention_causal_conv.py",
+        "test": "TestLinearAttentionCausalConvCPUParity.test_causal_conv_with_state_cpu_kernel_1",
+        "notes": "Direct standard rank-3 weight fixture for the ORT kernel=1 zero-size state edge case."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 4, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [4, 1, 1],
+          "data": { "kind": "values", "values": [0.5, -1.0, 1.5, -0.25] }
+        },
+        "biasT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.25, -0.5, 0.75, -1.0] } },
+        "pastStateT": { "dtype": "float32", "shape": [2, 4, 0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 4, 5], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 4, 0], "tolerance": 0 }
+      }
+    },
+    {
+      "name": "ort_basic_no_state_no_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.BasicNoStateNoBias",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 0.5, 1.5, 2.5, 3.5] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_silu_with_bias_and_state",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.SiluActivationWithBiasAndState",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 0.5, 1.5, 2.5, 3.5] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        },
+        "biasT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.1, -0.2] } },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [-1.0, 0.5, 0.3, -0.7] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_basic_with_bias",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.BasicWithBias",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 0.5, 1.5, 2.5, 3.5] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        },
+        "biasT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.1, -0.2] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_basic_with_state",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.BasicWithState",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 0.5, 1.5, 2.5] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [-1.0, 0.5, 0.3, -0.7] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 3], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_with_state_and_bias_none",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.WithStateAndBias",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 0.5, 1.5, 2.5] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        },
+        "biasT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.1, -0.2] } },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [-1.0, 0.5, 0.3, -0.7] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 3], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_silu_no_state",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.SiluActivationNoState",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 0.5, 1.5, 2.5, 3.5] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_silu_with_state",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.SiluActivationWithState",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 0.5, 1.5, 2.5] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [-1.0, 0.5, 0.3, -0.7] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 3], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_kernel_size2_state_silu",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.KernelSize2",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 0.5, 1.5, 2.5, 3.5] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 2],
+          "data": { "kind": "values", "values": [0.3, 0.7, 0.4, 0.6] }
+        },
+        "pastStateT": { "dtype": "float32", "shape": [1, 2, 1], "data": { "kind": "values", "values": [0.5, -0.3] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 1], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_kernel_size4_state_none",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.KernelSize4",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 5],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 4],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4] }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [-1.0, 0.0, 0.5] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 5], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 1, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_multi_batch_state_bias_silu",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.MultiBatch",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 0.5, 1.5, 2.5, -1.0, 0.0, 1.0, 0.2, 0.4, 0.6] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        },
+        "biasT": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.1, -0.1] } },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [-0.5, 0.5, 0.3, -0.3, 0.1, -0.1, 0.7, 0.8] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_single_token_decode_state_bias_silu",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.SingleTokenDecode",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 4, 1],
+          "data": { "kind": "values", "values": [0.5, -0.3, 1.2, 0.8] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [4, 1, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, -0.1, -0.2, 0.1, 0.2, 0.3, 0.3, 0.3, 0.3]
+          }
+        },
+        "biasT": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 0.1, -0.1, 0.0] } },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 4, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, -1.0, 0.0, 1.0, 0.5, 0.5, 0.5, -0.2, 0.4, -0.6] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 4, 1], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 4, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "ort_single_token_decode_multi_batch_silu",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.SingleTokenDecodeMultiBatch",
+        "notes": "Direct ORT depthwise weight shape [D,1,K]."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 2, 1],
+          "data": { "kind": "values", "values": [0.5, -0.3, 1.2, 0.8] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, -1.0, 0.0, 0.5, 0.5, -0.2, 0.4] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 2, 1], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "zero_state",
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 3, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [3, 1, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 3, 5], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 3, 2], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "scalar_bias_no_state_odd_length",
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 3, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [3, 1, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 3, 5], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 3, 2], "tolerance": 0.000001 }
+      },
+      "provenance": {
+        "notes": "Odd sequence length keeps the bias/no-state scalar fallback covered when the vec4 route is ineligible."
+      }
+    },
+    {
+      "name": "state_bias_silu",
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.41 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 2, 4], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "vec4_bias_no_state_silu",
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.24, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.18, "cosStep": 0.23 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.14, "cosStep": 0.41 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 2, 4], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.000001 }
+      },
+      "provenance": {
+        "notes": "Kernel 4 over a length that divides into vec4 lanes, with a bias and no carried state: the vectorized arm where the first lane's taps are the zero prefix rather than past_state."
+      }
+    },
+    {
+      "name": "vec4_state_no_bias_silu",
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 2, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.28, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.22, "cosStep": 0.23 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [2, 2, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.3, "cosStep": 0.13 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 2, 4], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.000001 }
+      },
+      "provenance": {
+        "notes": "Kernel 4 over a length that divides into vec4 lanes, with carried state and no bias: the vectorized arm that reads past_state into the first lane's taps but adds no bias term."
+      }
+    },
+    {
+      "name": "ort_larger_dimensions_state_bias_silu",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/causal_conv_with_state_op_test.cc",
+        "test": "CausalConvWithStateTest.LargerDimensions",
+        "notes": "Compact deterministic projection of ORT's larger-dimension state+bias SiLU stress case."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 8, 16],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.1, "cosStep": 0.0 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [8, 1, 4],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.0, "cosStep": 0.2 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0.0, 0.01, 0.02, 0.03, 0.04, 0.05, 0.06, 0.07] }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [2, 8, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.5, "sinStep": 0.3, "cosStep": 0.0 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 8, 16], "tolerance": 0.00002 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 8, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "zero_length_present_state_carryover_dropped",
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [1, 2, 0], "data": { "kind": "values", "values": [] } },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 3],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, 0.6] }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] }
+        }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 0],
+          "data": { "kind": "values", "values": [] },
+          "tolerance": 0
+        },
+        "presentStateT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "length_shorter_than_state_with_past_silu",
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [1.0, -2.0, 0.5, 3.0] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 5],
+          "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4, 0.5, -0.1, -0.2, 0.15, 0.25, 0.35] }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "values", "values": [-1.0, 0.5, 0.3, -0.7, 0.2, -0.4, 0.6, -0.8] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "length_shorter_than_state_no_state_zero_pad",
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 2],
+          "data": { "kind": "values", "values": [2.0, -1.0, 0.5, 4.0] }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": { "kind": "values", "values": [0.25, 0.5, -0.5, 1.0, 0.1, 0.2, 0.3, 0.4] }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 2], "tolerance": 0.00001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "vec4_zero_state_silu_compact",
+      "provenance": {
+        "notes": "Compact correctness lock for the aligned K=4 vec4 prefill path, including causal zero padding, SiLU, multi-batch rows, and present-state tails."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 3, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [3, 1, 4],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.11, "cosStep": 0.23 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 8], "tolerance": 0.00002 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 3, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_zero_state_compact",
+      "provenance": {
+        "notes": "Compact correctness lock for the workgroup-tiled large-kernel prefill path and its cooperative present-state update."
+      },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 31], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_bias_no_state_compact",
+      "provenance": {
+        "notes": "Compact correctness lock for the bias-only specialization of the workgroup-tiled large-kernel prefill path."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.125] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 1, 31], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_state_no_bias_compact",
+      "provenance": {
+        "notes": "Compact correctness lock for the carry-state specialization of the workgroup-tiled large-kernel prefill path."
+      },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 1, 31],
+          "data": { "kind": "fillFloat32", "scale": 0.15, "sinStep": 0.027, "cosStep": 0.019 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 1, 31], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_state_bias_silu_compact",
+      "provenance": {
+        "notes": "Compact correctness lock for the carry-state, bias, and SiLU specialization used by the production-shape fixture."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.075] } },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 1, 31],
+          "data": { "kind": "fillFloat32", "scale": 0.15, "sinStep": 0.027, "cosStep": 0.019 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 1, 31], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_state_bias_k128_wg64_multitile",
+      "provenance": {
+        "notes": "Smallest swept workgroup at the production kernel size; length 520 forces a partial second output tile."
+      },
+      "attrs": { "activation": "silu" },
+      "tunables": { "tiledWorkgroupSize": 64 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 520],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.075] } },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 1, 127],
+          "data": { "kind": "fillFloat32", "scale": 0.15, "sinStep": 0.027, "cosStep": 0.019 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 520], "tolerance": 0.0001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 1, 127], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_state_bias_k128_wg256",
+      "provenance": { "notes": "Largest swept workgroup at the production kernel and sequence sizes." },
+      "attrs": { "activation": "silu" },
+      "tunables": { "tiledWorkgroupSize": 256 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 512],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 128],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.075] } },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 1, 127],
+          "data": { "kind": "fillFloat32", "scale": 0.15, "sinStep": 0.027, "cosStep": 0.019 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 512], "tolerance": 0.0001 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 1, 127], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "state_window2_pinned",
+      "provenance": {
+        "notes": "Hand-computed from the ONNX Runtime state_window contract (onnxruntime/core/graph/contrib_ops/bert_defs.cc, the CausalConvWithState schema). Upstream's own state_window cases are CUDA-only and compare against a replayed reference rather than pinned numbers, so the expected values here were worked out by hand instead of ported. Slot 0 is the carry state after position 1 and slot 1 after position 2, so slot 1 repeats what the unwindowed op writes."
+      },
+      "attrs": { "activation": "none", "state_window": 2 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [1.0, 10.0, 100.0] }
+        }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [100.0, 210.0, 321.0] },
+          "tolerance": 0
+        },
+        "presentStateT": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 2.0, 3.0] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "state_window4_longer_than_sequence",
+      "provenance": {
+        "notes": "Hand-computed from the ONNX Runtime state_window contract (onnxruntime/core/graph/contrib_ops/bert_defs.cc, the CausalConvWithState schema). Upstream's own state_window cases are CUDA-only and compare against a replayed reference rather than pinned numbers, so the expected values here were worked out by hand instead of ported. W exceeds the sequence length, so the leading W - T slots must be zero rather than uninitialized."
+      },
+      "attrs": { "activation": "none", "state_window": 4 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [2.0, -1.0, 4.0] } },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [1.0, 10.0, 100.0] }
+        }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [200.0, -80.0, 392.0] },
+          "tolerance": 0
+        },
+        "presentStateT": {
+          "dtype": "float32",
+          "shape": [4, 1, 1, 2],
+          "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 2.0, 2.0, -1.0, -1.0, 4.0] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "state_window2_past_slot_pinned",
+      "provenance": {
+        "notes": "Hand-computed from the ONNX Runtime state_window contract (onnxruntime/core/graph/contrib_ops/bert_defs.cc, the CausalConvWithState schema). Upstream's own state_window cases are CUDA-only and compare against a replayed reference rather than pinned numbers, so the expected values here were worked out by hand instead of ported. past_state slot 0 is poisoned with large negatives that no correct read touches; only slot W-1 carries the previous call's state."
+      },
+      "attrs": { "activation": "none", "state_window": 2 },
+      "inputs": {
+        "inputT": { "dtype": "float32", "shape": [1, 1, 2], "data": { "kind": "values", "values": [1.0, 2.0] } },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 3],
+          "data": { "kind": "values", "values": [1.0, 10.0, 100.0] }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 2],
+          "data": { "kind": "values", "values": [-1000.0, -2000.0, 5.0, 7.0] }
+        }
+      },
+      "outputs": {
+        "outputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 2],
+          "data": { "kind": "values", "values": [175.0, 217.0] },
+          "tolerance": 0
+        },
+        "presentStateT": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 2],
+          "data": { "kind": "values", "values": [7.0, 1.0, 1.0, 2.0] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "vec4_state_window3",
+      "provenance": {
+        "notes": "Gives the aligned K=4 vec4 prefill path a windowed present_state; its scalar-typed state output has to be gathered lane by lane out of the vec4 input row."
+      },
+      "attrs": { "activation": "silu", "state_window": 3 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.11, "cosStep": 0.23 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 8], "tolerance": 0.00002 },
+        "presentStateT": { "dtype": "float32", "shape": [3, 1, 2, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_zero_state_window2",
+      "provenance": {
+        "notes": "Windowed present_state on the large-kernel tiled path with no past state; the state-writing tile now strides over a (slot, element) grid instead of a single slot."
+      },
+      "attrs": { "activation": "none", "state_window": 2 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 1, 1, 31], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_bias_no_state_window2",
+      "provenance": {
+        "notes": "Exercises windowed present-state publication on the large-kernel tiled route when bias is present but past state is absent."
+      },
+      "attrs": { "activation": "silu", "state_window": 2 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.075] } }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 1, 1, 31], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_state_no_bias_window2",
+      "provenance": {
+        "notes": "Exercises windowed past-state reads and present-state publication on the large-kernel tiled route without bias."
+      },
+      "attrs": { "activation": "none", "state_window": 2 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 31],
+          "data": { "kind": "fillFloat32", "scale": 0.15, "sinStep": 0.027, "cosStep": 0.019 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 1, 1, 31], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_state_bias_window2",
+      "provenance": {
+        "notes": "Windowed present_state on the large-kernel tiled path with a windowed past_state and bias; the earliest slot still reaches back into the carried state."
+      },
+      "attrs": { "activation": "silu", "state_window": 2 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 1, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [1, 1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "biasT": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [-0.075] } },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [2, 1, 1, 31],
+          "data": { "kind": "fillFloat32", "scale": 0.15, "sinStep": 0.027, "cosStep": 0.019 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 1, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 1, 1, 31], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "vec4_state_window6_longer_than_sequence",
+      "provenance": {
+        "notes": "W = 6 exceeds the four-position input, so the vec4 path's two leading window slots hold no position from this call and must be zero. Only this variant can reach that branch with a window: the tiled path demands at least 256 positions, which no legal window exceeds."
+      },
+      "attrs": { "activation": "none", "state_window": 6 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.19, "cosStep": 0.37 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.00002 },
+        "presentStateT": { "dtype": "float32", "shape": [6, 1, 2, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "vec4_state_window_past_state_prefix",
+      "provenance": {
+        "notes": "A windowed state whose window reaches back further than this call is long, WITH a past state: the early slots carry positions from before this call, so they have to come from past_state rather than from the input row."
+      },
+      "attrs": { "activation": "silu", "state_window": 6 },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 4],
+          "data": { "kind": "fillFloat32", "scale": 0.6, "sinStep": 0.29, "cosStep": 0.13 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 4],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.19, "cosStep": 0.37 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [6, 1, 2, 3],
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.37, "scale": 0.5 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 4], "tolerance": 0.00002 },
+        "presentStateT": { "dtype": "float32", "shape": [6, 1, 2, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "f16_scalar_state_bias_silu",
+      "provenance": {
+        "notes": "float16 tensors on the scalar kernel. ONNX Runtime registers this operator for the whole supported float set; this port pinned float32. Every tap and accumulation still runs in f32 and only the store narrows, which is what the kernel already did for float32."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [1, 2, 4],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float16",
+          "shape": [2, 1, 3],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "biasT": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.1, -0.2] } },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [1, 2, 2],
+          "data": { "kind": "fillFloat32", "scale": 0.3, "sinStep": 0.29, "cosStep": 0.13 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 2, 4], "tolerance": 0.005 },
+        "presentStateT": { "dtype": "float16", "shape": [1, 2, 2], "tolerance": 0.005 }
+      }
+    },
+    {
+      "name": "f16_k4_vec4_zero_state_silu",
+      "provenance": {
+        "notes": "float16 on the four-tap vectorized kernel, which read the bound element type directly and so was the only one of the three actually pinned to float32."
+      },
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [2, 3, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float16",
+          "shape": [3, 1, 4],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.11, "cosStep": 0.23 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [2, 3, 8], "tolerance": 0.005 },
+        "presentStateT": { "dtype": "float16", "shape": [2, 3, 3], "tolerance": 0.005 }
+      }
+    },
+    {
+      "name": "f16_large_kernel_tiled_state_bias",
+      "provenance": {
+        "notes": "float16 on the tiled large-kernel path, which stages the weight and the virtual input in float32 workgroup memory regardless of the tensor type."
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float16",
+          "shape": [1, 1, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float16",
+          "shape": [1, 1, 32],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "biasT": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [-0.075] } },
+        "pastStateT": {
+          "dtype": "float16",
+          "shape": [1, 1, 31],
+          "data": { "kind": "fillFloat32", "scale": 0.15, "sinStep": 0.027, "cosStep": 0.019 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float16", "shape": [1, 1, 256], "tolerance": 0.01 },
+        "presentStateT": { "dtype": "float16", "shape": [1, 1, 31], "tolerance": 0.005 }
+      }
+    },
+    {
+      "name": "weight_rank3_k4_vec4_zero_state",
+      "provenance": {
+        "notes": "A rank-3 weight on the four-tap vectorized kernel, where the kernel extent is read as a vec4 rather than element by element."
+      },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [2, 3, 8],
+          "data": { "kind": "fillFloat32", "scale": 0.4, "sinStep": 0.17, "cosStep": 0.31 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [3, 1, 4],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.11, "cosStep": 0.23 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [2, 3, 8], "tolerance": 0.00002 },
+        "presentStateT": { "dtype": "float32", "shape": [2, 3, 3], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_unaligned_k33_weight_tile_pad",
+      "provenance": {
+        "notes": "Kernel length 1 mod 4. The tap loop consumes four weights per iteration, so this shape reaches the tiled path only via the zero-padded weight tile; before that it fell to the untiled kernel."
+      },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 33],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 32], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_unaligned_k34_weight_tile_pad",
+      "provenance": { "notes": "Kernel length 2 mod 4 -- the other half of the padded-tail arithmetic." },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 34],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 33], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_unaligned_k35_bias_weight_tile_pad",
+      "provenance": {
+        "notes": "Kernel length 3 mod 4, the largest pad, with a bias so the padded tail is exercised on the bias arm of the family too."
+      },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 35],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "biasT": {
+          "dtype": "float32",
+          "shape": [2],
+          "data": { "kind": "fillFloat32", "scale": 0.1, "sinStep": 0.07, "cosStep": 0.03 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 34], "tolerance": 0.000001 }
+      }
+    },
+    {
+      "name": "large_kernel_tiled_unaligned_k37_state_weight_tile_pad",
+      "provenance": {
+        "notes": "Kernel length 1 mod 4 carrying past state, so the padded weight tile is covered on the stateful arm where STATE_LENGTH stays the true kernel-1."
+      },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "dtype": "float32",
+          "shape": [1, 2, 256],
+          "data": { "kind": "fillFloat32", "scale": 0.2, "sinStep": 0.013, "cosStep": 0.021 }
+        },
+        "weightT": {
+          "dtype": "float32",
+          "shape": [2, 1, 37],
+          "data": { "kind": "fillFloat32", "scale": 0.05, "sinStep": 0.031, "cosStep": 0.017 }
+        },
+        "pastStateT": {
+          "dtype": "float32",
+          "shape": [1, 2, 36],
+          "data": { "kind": "fillFloat32", "scale": 0.15, "sinStep": 0.011, "cosStep": 0.029 }
+        }
+      },
+      "outputs": {
+        "outputT": { "dtype": "float32", "shape": [1, 2, 256], "tolerance": 0.00005 },
+        "presentStateT": { "dtype": "float32", "shape": [1, 2, 36], "tolerance": 0.000001 }
+      }
+    }
+  ]
+}