sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +73 -0
build/webgpu/bench.json +78 -0
build/webgpu/causal-conv-with-state-tiled.wgsl.jinja +177 -0
build/webgpu/causal-conv-with-state-vec4.wgsl.jinja +112 -0
build/webgpu/causal-conv-with-state.wgsl.jinja +108 -0
build/webgpu/manifest.json +1077 -0
build/webgpu/metadata.json +20 -0
build/webgpu/test.json +358 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.CausalConvWithState
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 27
+## Description
+Stateful causal 1-D depthwise convolution. It left-pads each channel with `past_state` (or zeros), applies a `(channels, 1, kernel_size)` depthwise convolution and optional SiLU/Swish activation, and returns both the output and the last `kernel_size - 1` values as `present_state`. This inference implementation supports float16 and float32 tensors; bfloat16 is not yet implemented.
+See the [ONNX `CausalConvWithState` spec](https://onnx.ai/onnx/operators/onnx__CausalConvWithState.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `inputT` | `T` | `3` | — | Input tensor with shape `(batch_size, channels, length)` in channels-first layout. | required |
+| `weight` | `weightT` | `T` | `3` | — | Depthwise convolution kernel with shape `(channels, 1, kernel_size)`, matching the ONNX Conv weight layout for `group = channels`. | required |
+| `bias` | `biasT` | `T` | `1` | — | Optional per-channel bias with shape `(channels)`. | optional |
+| `past_state` | `pastStateT` | `T` | `3` | — | Carry state from the previous step with shape `(batch_size, channels, kernel_size - 1)`. If omitted, the left 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` | `3` | derived; see description | Updated carry state with shape `(batch_size, channels, kernel_size - 1)`, including prior state or zero-padding when the current input is shorter than the state. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `activation` | `"none"` | Optional fused activation. `silu` and `swish` are aliases; `none` leaves the convolution result unchanged. |
+## 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/ai.onnx.CausalConvWithState", { version: 1 });
+const { outputT, presentStateT } = await kernel({
+  inputT: { data: inputTData, shape: [1, 1, 2] },
+  weightT: { data: weightTData, shape: [1, 1, 5] },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,78 @@

+{
+  "op": "ai.onnx.CausalConvWithState",
+  "tunableSpace": { "workgroupSize": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "decode-f16-b1c1024t1k4",
+      "preset": "smoke",
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": { "shape": [1, 1024, 1], "dtype": "float16", "dist": "normal", "seed": 2751, "scale": 0.2 },
+        "weightT": { "shape": [1024, 1, 4], "dtype": "float16", "dist": "normal", "seed": 2752, "scale": 0.1 },
+        "biasT": { "shape": [1024], "dtype": "float16", "dist": "normal", "seed": 2753, "scale": 0.05 },
+        "pastStateT": { "shape": [1, 1024, 3], "dtype": "float16", "dist": "normal", "seed": 2754, "scale": 0.2 }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1024, 1], "dtype": "float16" },
+        "presentStateT": { "shape": [1, 1024, 3], "dtype": "float16" }
+      },
+      "bench": {
+        "primary": true,
+        "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)) * 2"
+          }
+        ]
+      }
+    },
+    {
+      "name": "prefill-f32-b2c32t256k4",
+      "preset": "smoke",
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": { "shape": [2, 32, 256], "dtype": "float32", "dist": "normal", "seed": 2755, "scale": 0.2 },
+        "weightT": { "shape": [32, 1, 4], "dtype": "float32", "dist": "normal", "seed": 2756, "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, 2)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    },
+    {
+      "name": "large-kernel-f32-b2c1024t512k128",
+      "preset": "stress",
+      "tunableSpace": { "tiledWorkgroupSize": [64, 128, 256] },
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": { "shape": [2, 1024, 512], "dtype": "float32", "dist": "normal", "seed": 2757, "scale": 0.2 },
+        "weightT": { "shape": [1024, 1, 128], "dtype": "float32", "dist": "normal", "seed": 2758, "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, 2)" },
+          {
+            "type": "bandwidth",
+            "value": "(numel(shapes.inputT) + numel(shapes.weightT) + numel(shapes.outputT) + numel(shapes.presentStateT)) * 4"
+          }
+        ]
+      }
+    }
+  ]
+}

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

	@@ -0,0 +1,177 @@

+{% 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 %}
+  // Materialize the convolution result in T before the f32 SiLU epilogue.
+  let convolution_output = f32({{ outputScalar }}(value));
+  let sigmoid = 1.0 / (1.0 + exp(-convolution_output));
+  return convolution_output * sigmoid;
+{% 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,112 @@

+{% 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 %}
+  // Materialize the convolution result in T before the f32 SiLU epilogue.
+  let convolution_output = vec4<f32>({{ outputVec4 }}(value));
+  let sigmoid = vec4<f32>(1.0) / (vec4<f32>(1.0) + exp(-convolution_output));
+  return convolution_output * sigmoid;
+{% 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,108 @@

+{% 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 %}
+  // Materialize the convolution result in T before the f32 SiLU epilogue.
+  let convolution_output = f32({{ outputScalar }}(value));
+  let sigmoid = 1.0 / (1.0 + exp(-convolution_output));
+  return convolution_output * sigmoid;
+{% 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,1077 @@

+{
+  "domain": "ai.onnx",
+  "name": "CausalConvWithState",
+  "sinceVersion": 27,
+  "description": "Stateful causal 1-D depthwise convolution. It left-pads each channel with `past_state` (or zeros), applies a `(channels, 1, kernel_size)` depthwise convolution and optional SiLU/Swish activation, and returns both the output and the last `kernel_size - 1` values as `present_state`. This inference implementation supports float16 and float32 tensors; bfloat16 is not yet implemented.",
+  "inputs": [
+    {
+      "role": "input",
+      "dtype": "T",
+      "rank": 3,
+      "description": "Input tensor with shape `(batch_size, channels, length)` in channels-first layout."
+    },
+    {
+      "role": "weight",
+      "dtype": "T",
+      "rank": 3,
+      "description": "Depthwise convolution kernel with shape `(channels, 1, kernel_size)`, matching the ONNX Conv weight layout for `group = channels`."
+    },
+    {
+      "role": "bias",
+      "dtype": "T",
+      "rank": 1,
+      "optional": true,
+      "description": "Optional per-channel bias with shape `(channels)`."
+    },
+    {
+      "role": "past_state",
+      "dtype": "T",
+      "rank": 3,
+      "optional": true,
+      "description": "Carry state from the previous step with shape `(batch_size, channels, kernel_size - 1)`. If omitted, the left 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,
+      "shape": "[dim(shapes.input, 0), dim(shapes.input, 1), dim(shapes.weight, 2) - 1]",
+      "description": "Updated carry state with shape `(batch_size, channels, kernel_size - 1)`, including prior state or zero-padding when the current input is shorter than the state."
+    }
+  ],
+  "attributes": { "activation": "none" },
+  "attributeConstraints": { "activation": { "values": ["none", "silu", "swish"] } },
+  "attributeDescriptions": {
+    "activation": "Optional fused activation. `silu` and `swish` are aliases; `none` leaves the convolution result unchanged."
+  },
+  "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": {
+    "kernelSize": "dim(shapes.weightT, 2)",
+    "kernelSizePadded": "ceilDiv(kernelSize, 4) * 4",
+    "stateLength": "kernelSize - 1",
+    "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",
+    "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",
+    "commonContract": "ranks.inputT == 3 and ranks.weightT == 3 and ranks.outputT == 3 and kernelSize >= 1 and dim(shapes.weightT, 0) == dim(shapes.inputT, 1) and dim(shapes.weightT, 1) == 1 and sameShape(shapes.outputT, shapes.inputT) and presentStateOk and f16Ok(dtypes.T)",
+    "zeroStateContract": "commonContract and not present.pastStateT and not present.biasT",
+    "biasNoStateContract": "commonContract and not present.pastStateT and present.biasT and ranks.biasT == 1 and dim(shapes.biasT, 0) == dim(shapes.inputT, 1)",
+    "stateNoBiasContract": "commonContract and present.pastStateT and not present.biasT and pastStateShapeOk",
+    "stateBiasContract": "commonContract and present.pastStateT and present.biasT and ranks.biasT == 1 and dim(shapes.biasT, 0) == dim(shapes.inputT, 1) and pastStateShapeOk"
+  },
+  "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" }
+          ]
+        }
+      }
+    ],
+    "biasNoStateScalar": [
+      {
+        "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" }
+          ]
+        }
+      }
+    ],
+    "stateNoBiasScalar": [
+      {
+        "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" }
+          ]
+        }
+      }
+    ],
+    "stateBiasScalar": [
+      {
+        "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" }
+          ]
+        }
+      }
+    ],
+    "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)" }
+          ]
+        }
+      }
+    ],
+    "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)" }
+          ]
+        }
+      }
+    ],
+    "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)" }
+          ]
+        }
+      }
+    ],
+    "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)" }
+          ]
+        }
+      }
+    ],
+    "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)" }
+          ]
+        }
+      }
+    ],
+    "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)" }
+          ]
+        }
+      }
+    ],
+    "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)" }
+          ]
+        }
+      }
+    ],
+    "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)" }
+          ]
+        }
+      }
+    ]
+  },
+  "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": false,
+        "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": true }
+          },
+          "bindings": "zeroVec4",
+          "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": false,
+        "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": true }
+          },
+          "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": false,
+        "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": true }
+          },
+          "bindings": "stateNoBiasVec4",
+          "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": false,
+        "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": true }
+          },
+          "bindings": "stateBiasVec4",
+          "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": false,
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.TiledLargeKernel",
+          "source": {
+            "shader": "causal-conv-with-state-tiled.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": true }
+          },
+          "bindings": "zeroTiled",
+          "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": false,
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.TiledLargeKernel",
+          "source": {
+            "shader": "causal-conv-with-state-tiled.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": true }
+          },
+          "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": false,
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.TiledLargeKernel",
+          "source": {
+            "shader": "causal-conv-with-state-tiled.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": true }
+          },
+          "bindings": "stateNoBiasTiled",
+          "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": false,
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState.TiledLargeKernel",
+          "source": {
+            "shader": "causal-conv-with-state-tiled.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": true }
+          },
+          "bindings": "stateBiasTiled",
+          "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": false,
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState",
+          "source": {
+            "shader": "causal-conv-with-state.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": true }
+          },
+          "bindings": "zeroScalar",
+          "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": false,
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState",
+          "source": {
+            "shader": "causal-conv-with-state.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": true }
+          },
+          "bindings": "biasNoStateScalar",
+          "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": false,
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState",
+          "source": {
+            "shader": "causal-conv-with-state.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": true }
+          },
+          "bindings": "stateNoBiasScalar",
+          "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": false,
+        "usesF16": "tensorDtypes.inputT == \"float16\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CausalConvWithState",
+          "source": {
+            "shader": "causal-conv-with-state.wgsl.jinja",
+            "inputs": { "materializeConvBeforeActivation": true }
+          },
+          "bindings": "stateBiasScalar",
+          "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": "ai.onnx.CausalConvWithState",
+  "id": "_ai_onnx_causalconvwithstate_webgpu_0d9e334",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "fCrdgbJNyH7TGMSVJ0qqqadgttfP8Aab9dv/haLYRZM=",
+      "causal-conv-with-state-tiled.wgsl.jinja": "X/uUVtlr1/PhjaCIQI7Z6EyqZzcldbk6agI1eX2UGqY=",
+      "causal-conv-with-state-vec4.wgsl.jinja": "BMP+FWcwzs8ig4DN9gh7PdEe/jTVZKSZah48B9WvTZM=",
+      "causal-conv-with-state.wgsl.jinja": "jf0H+njpBJ0uBPp3hWfdn+dyWkjSTIeM5H357XyeaN0=",
+      "manifest.json": "Y+eFrTwuYgJYldZPqxz/xEgCwYLiOg3A/X/Ni6/aj/E=",
+      "test.json": "yhv3pfQudb0pD6fFnYnV7jLeNKkxCRkzvg90AXG7Cy4="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.CausalConvWithState" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,358 @@

+{
+  "op": "ai.onnx.CausalConvWithState",
+  "cases": [
+    {
+      "name": "scalar_zero_state",
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "shape": [1, 2, 3],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [2, 1, 3],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.11, "scale": 0.2 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 2, 3], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 2, 2], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "scalar_bias_no_state",
+      "attrs": { "activation": "none" },
+      "inputs": {
+        "inputT": {
+          "shape": [1, 2, 5],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.31, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [2, 1, 5],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.27, "scale": 0.2 }
+        },
+        "biasT": {
+          "shape": [2],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.07, "scale": 0.1 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 2, 5], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 2, 4], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "scalar_state_no_bias_swish",
+      "attrs": { "activation": "swish" },
+      "inputs": {
+        "inputT": {
+          "shape": [1, 2, 1],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.33, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [2, 1, 4],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.29, "cosStep": 0.17, "scale": 0.2 }
+        },
+        "pastStateT": {
+          "shape": [1, 2, 3],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.21, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 2, 1], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 2, 3], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "scalar_f16_state_bias_silu_cast_boundary",
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "shape": [1, 1, 3],
+          "dtype": "float16",
+          "data": { "kind": "values", "values": [1.001, -0.731, 0.287] }
+        },
+        "weightT": {
+          "shape": [1, 1, 3],
+          "dtype": "float16",
+          "data": { "kind": "values", "values": [0.619, -1.127, 0.443] }
+        },
+        "biasT": { "shape": [1], "dtype": "float16", "data": { "kind": "values", "values": [0.037] } },
+        "pastStateT": {
+          "shape": [1, 1, 2],
+          "dtype": "float16",
+          "data": { "kind": "values", "values": [-0.219, 0.887] }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 3], "dtype": "float16" },
+        "presentStateT": { "shape": [1, 1, 2], "dtype": "float16" }
+      }
+    },
+    {
+      "name": "kernel_size_one",
+      "inputs": {
+        "inputT": {
+          "shape": [2, 2, 3],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [2, 1, 1],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.09, "scale": 0.2 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [2, 2, 3], "dtype": "float32" },
+        "presentStateT": { "shape": [2, 2, 0], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "input_shorter_than_state_zero_padding",
+      "inputs": {
+        "inputT": { "shape": [1, 1, 2], "dtype": "float32", "data": { "kind": "values", "values": [2.0, 3.0] } },
+        "weightT": {
+          "shape": [1, 1, 5],
+          "dtype": "float32",
+          "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0, 1.0] }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 2], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 1, 4], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "zero_length_reemits_past_state",
+      "inputs": {
+        "inputT": { "shape": [1, 1, 0], "dtype": "float32", "data": { "kind": "values", "values": [] } },
+        "weightT": { "shape": [1, 1, 3], "dtype": "float32", "data": { "kind": "values", "values": [1.0, 1.0, 1.0] } },
+        "pastStateT": { "shape": [1, 1, 2], "dtype": "float32", "data": { "kind": "values", "values": [4.0, 5.0] } }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 0], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 1, 2], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "vec4_zero_state_f16_silu",
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "shape": [2, 3, 8],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [3, 1, 4],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.27, "cosStep": 0.11, "scale": 0.2 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [2, 3, 8], "dtype": "float16" },
+        "presentStateT": { "shape": [2, 3, 3], "dtype": "float16" }
+      }
+    },
+    {
+      "name": "vec4_bias_no_state_swish",
+      "attrs": { "activation": "swish" },
+      "inputs": {
+        "inputT": {
+          "shape": [1, 2, 8],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.13, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [2, 1, 3],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.31, "scale": 0.2 }
+        },
+        "biasT": {
+          "shape": [2],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.33, "cosStep": 0.07, "scale": 0.1 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 2, 8], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 2, 2], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "vec4_state_no_bias",
+      "inputs": {
+        "inputT": {
+          "shape": [1, 2, 4],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.15, "cosStep": 0.25, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [2, 1, 2],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.35, "cosStep": 0.05, "scale": 0.2 }
+        },
+        "pastStateT": {
+          "shape": [1, 2, 1],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.29, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 2, 4], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 2, 1], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "vec4_state_bias_silu",
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "shape": [1, 2, 8],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.27, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [2, 1, 4],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.25, "cosStep": 0.13, "scale": 0.2 }
+        },
+        "biasT": {
+          "shape": [2],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.39, "cosStep": 0.05, "scale": 0.1 }
+        },
+        "pastStateT": {
+          "shape": [1, 2, 3],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.33, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 2, 8], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 2, 3], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "tiled_zero_state",
+      "inputs": {
+        "inputT": {
+          "shape": [1, 1, 256],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.03, "cosStep": 0.05, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [1, 1, 32],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.09, "scale": 0.1 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 256], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 1, 31], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "tiled_bias_no_state_silu",
+      "attrs": { "activation": "silu" },
+      "inputs": {
+        "inputT": {
+          "shape": [1, 1, 256],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.04, "cosStep": 0.06, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [1, 1, 32],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.08, "cosStep": 0.1, "scale": 0.1 }
+        },
+        "biasT": { "shape": [1], "dtype": "float32", "data": { "kind": "constant", "value": 0.05 } }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 256], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 1, 31], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "tiled_state_no_bias",
+      "inputs": {
+        "inputT": {
+          "shape": [1, 1, 256],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.07, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [1, 1, 32],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.11, "scale": 0.1 }
+        },
+        "pastStateT": {
+          "shape": [1, 1, 31],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.17, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 256], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 1, 31], "dtype": "float32" }
+      }
+    },
+    {
+      "name": "tiled_state_bias_f16_swish",
+      "attrs": { "activation": "swish" },
+      "inputs": {
+        "inputT": {
+          "shape": [1, 1, 256],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.06, "cosStep": 0.08, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [1, 1, 32],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.1, "cosStep": 0.12, "scale": 0.1 }
+        },
+        "biasT": { "shape": [1], "dtype": "float16", "data": { "kind": "constant", "value": -0.05 } },
+        "pastStateT": {
+          "shape": [1, 1, 31],
+          "dtype": "float16",
+          "data": { "kind": "fillFloat32", "sinStep": 0.14, "cosStep": 0.18, "scale": 0.25 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 256], "dtype": "float16" },
+        "presentStateT": { "shape": [1, 1, 31], "dtype": "float16" }
+      }
+    },
+    {
+      "name": "tiled_zero_state_unaligned_k35_weight_tile_pad",
+      "inputs": {
+        "inputT": {
+          "shape": [1, 1, 256],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.03, "cosStep": 0.05, "scale": 0.25 }
+        },
+        "weightT": {
+          "shape": [1, 1, 35],
+          "dtype": "float32",
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.09, "scale": 0.1 }
+        }
+      },
+      "outputs": {
+        "outputT": { "shape": [1, 1, 256], "dtype": "float32" },
+        "presentStateT": { "shape": [1, 1, 34], "dtype": "float32" }
+      },
+      "provenance": {
+        "notes": "Kernel length 3 mod 4. The shared tiled kernel consumes four taps per iteration, so this shape reaches it only through the zero-padded weight tile; ai.onnx carries its own copy of that gate, so it needs its own lock."
+      }
+    }
+  ]
+}