sync 2e7068faf55e

Browse files

Files changed (13) hide show

README.md +83 -0
build/webgpu/bench.json +252 -0
build/webgpu/dft-tiled-real.wgsl.jinja +140 -0
build/webgpu/fft-radix2-dit-shared-inplace.wgsl.jinja +149 -0
build/webgpu/fft-radix2-dit-storage-pass.wgsl.jinja +215 -0
build/webgpu/fft-radix2-dit-storage.wgsl.jinja +84 -0
build/webgpu/fft-real-packed-stockham-shared.wgsl.jinja +156 -0
build/webgpu/fft-stockham-shared.wgsl.jinja +292 -0
build/webgpu/fft-stockham-twiddles.wgsl.jinja +125 -0
build/webgpu/manifest.json +644 -0
build/webgpu/metadata.json +25 -0
build/webgpu/stft-general.wgsl.jinja +164 -0
build/webgpu/test.json +847 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.STFT
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 17
+## Description
+Computes the short-time Fourier transform (STFT) of a rank-3 real or complex signal by sliding a complete frame over the input and applying the DFT. The `frameStep` scalar represents the required positive `frame_step` input. At least one of `window` or `frameLength` is required; when both are supplied, their lengths must match. The output contains the real and imaginary components of every frequency bin for each complete frame.
+See the [ONNX `STFT` spec](https://onnx.ai/onnx/operators/onnx__STFT.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `signal` | `signal` | `T` | `3` | — | Real input signal with shape `(batch_size, signal_length, 1)`, or complex input with shape `(batch_size, signal_length, 2)`. | required |
+| `window` | `window` | `T` | `1` | — | Optional 1-D window tensor with shape `(window_length)` to multiply by each frame before the DFT. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | `4` | derived; see description | STFT result with shape `(batch_size, frames, dft_unique_bins, 2)` in one-sided mode or `(batch_size, frames, frame_length, 2)` in two-sided mode, with real and imaginary parts in the last dimension. `frames = floor((signal_length - frame_length) / frame_step) + 1`; partial trailing frames are not emitted. | required |
+## Runtime arguments
+| Name | Kind | Semantic | Description | Presence |
+| --- | --- | --- | --- | --- |
+| `frameStep` | `i32` | `frame_step` | Positive number of signal positions between the starts of consecutive frames. | required |
+| `frameLength` | `i32` | `frame_length` | Optional number of signal values in each frame; must equal the window length when both are supplied. | optional |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `onesided` | `1` | When 1 (default), returns `floor(n_fft / 2) + 1` non-redundant frequency bins with indices `[0, ..., floor(n_fft / 2)]` for a real-valued input (RFFT); set to 0 to return the full spectrum. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`dft-tiled-real.wgsl.jinja`](build/webgpu/dft-tiled-real.wgsl.jinja)
+- [`fft-radix2-dit-shared-inplace.wgsl.jinja`](build/webgpu/fft-radix2-dit-shared-inplace.wgsl.jinja)
+- [`fft-radix2-dit-storage-pass.wgsl.jinja`](build/webgpu/fft-radix2-dit-storage-pass.wgsl.jinja)
+- [`fft-radix2-dit-storage.wgsl.jinja`](build/webgpu/fft-radix2-dit-storage.wgsl.jinja)
+- [`fft-real-packed-stockham-shared.wgsl.jinja`](build/webgpu/fft-real-packed-stockham-shared.wgsl.jinja)
+- [`fft-stockham-shared.wgsl.jinja`](build/webgpu/fft-stockham-shared.wgsl.jinja)
+- [`fft-stockham-twiddles.wgsl.jinja`](build/webgpu/fft-stockham-twiddles.wgsl.jinja)
+- [`stft-general.wgsl.jinja`](build/webgpu/stft-general.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.STFT", { version: 1 });
+const { output } = await kernel({
+  signal: { data: signalData, shape: [1, 4, 1] },
+  window: { data: windowData, shape: [4] },
+  frameStep: 4,
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,252 @@

+{
+  "op": "ai.onnx.STFT",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "batch4_signal1024_frame256_step128",
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [4, 1024, 1] },
+        "window": { "dtype": "float32", "shape": [256] }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 7, 129, 2] } },
+      "args": { "frameStep": 128 }
+    },
+    {
+      "name": "fft_frame256_signal2048_step128_batch4",
+      "preset": "smoke",
+      "vars": { "frames": 60, "n": 256, "log2n": 8 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [4, 2048, 1], "dist": "normal", "seed": 911 },
+        "window": { "dtype": "float32", "shape": [256], "dist": "uniform", "seed": 912 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 15, 129, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.frames * args.n * args.log2n" }] },
+      "args": { "frameStep": 128 }
+    },
+    {
+      "name": "fft_frame512_signal4096_step256_batch2",
+      "preset": "smoke",
+      "vars": { "frames": 30, "n": 512, "log2n": 9 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [2, 4096, 1], "dist": "normal", "seed": 913 },
+        "window": { "dtype": "float32", "shape": [512], "dist": "uniform", "seed": 914 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 15, 257, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.frames * args.n * args.log2n" }] },
+      "args": { "frameStep": 256 }
+    },
+    {
+      "name": "fft_frame1024_signal131072_step256_audio",
+      "preset": "smoke",
+      "vars": { "frames": 509, "n": 1024, "log2n": 10 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [1, 131072, 1], "dist": "normal", "seed": 915 },
+        "window": { "dtype": "float32", "shape": [1024], "dist": "uniform", "seed": 916 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 509, 513, 2] } },
+      "bench": { "primary": true, "metrics": [{ "type": "gflops", "value": "5 * args.frames * args.n * args.log2n" }] },
+      "args": { "frameStep": 256 }
+    },
+    {
+      "name": "pow2_gate_FFT_win512_signal8192_step160",
+      "preset": "smoke",
+      "vars": { "batch": 2, "frames": 49, "bins": 257, "n": 512, "log2n": 9 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [2, 8192, 1], "dist": "normal", "seed": 2001 },
+        "window": { "dtype": "float32", "shape": [512], "dist": "uniform", "seed": 2002 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 49, 257, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.batch * args.frames * args.n * args.log2n" }] },
+      "args": { "frameStep": 160 }
+    },
+    {
+      "name": "mixed_radix_win480_signal8192_step160_tail15",
+      "preset": "smoke",
+      "vars": { "batch": 2, "frames": 49, "bins": 241, "n": 480, "log2radix2": 5, "tail": 15 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [2, 8192, 1], "dist": "normal", "seed": 2001 },
+        "window": { "dtype": "float32", "shape": [480], "dist": "uniform", "seed": 2003 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 49, 241, 2] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "args.batch * args.frames * (5 * args.n * args.log2radix2 + 8 * args.bins * args.tail)"
+          }
+        ]
+      },
+      "args": { "frameStep": 160 }
+    },
+    {
+      "name": "mixed_radix_win400_signal16000_step160_batch4_tail25",
+      "preset": "smoke",
+      "vars": { "batch": 4, "frames": 98, "bins": 201, "n": 400, "log2radix2": 4, "tail": 25 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [4, 16000, 1], "dist": "normal", "seed": 2010 },
+        "window": { "dtype": "float32", "shape": [400], "dist": "uniform", "seed": 2011 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 98, 201, 2] } },
+      "bench": {
+        "metrics": [
+          {
+            "type": "gflops",
+            "value": "args.batch * args.frames * (5 * args.n * args.log2radix2 + 8 * args.bins * args.tail)"
+          }
+        ]
+      },
+      "args": { "frameStep": 160 }
+    },
+    {
+      "name": "radix2_win512_signal16000_step160_batch4",
+      "preset": "smoke",
+      "vars": { "batch": 4, "frames": 97, "bins": 257, "n": 512, "log2n": 9 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [4, 16000, 1], "dist": "normal", "seed": 2010 },
+        "window": { "dtype": "float32", "shape": [512], "dist": "uniform", "seed": 2012 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 97, 257, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.batch * args.frames * args.n * args.log2n" }] },
+      "args": { "frameStep": 160 }
+    },
+    {
+      "name": "dispatch_cliff_NAIVE_batch48_frames1397_win512_step160",
+      "preset": "stress",
+      "vars": { "batch": 48, "frames": 1397, "bins": 257, "n": 512, "log2n": 9, "signalLength": 224000 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [48, 224000, 1], "dist": "normal", "seed": 2020 },
+        "window": { "dtype": "float32", "shape": [512], "dist": "uniform", "seed": 2021 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [48, 1397, 257, 2] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "5 * args.batch * args.frames * args.n * args.log2n" },
+          {
+            "type": "bandwidth",
+            "value": "4 * (args.batch * args.signalLength + args.n + args.batch * args.frames * args.bins * 2)"
+          }
+        ]
+      },
+      "args": { "frameStep": 160 }
+    },
+    {
+      "name": "dispatch_cliff_sibling_FFT_batch32_frames1397_win512_step160",
+      "preset": "stress",
+      "vars": { "batch": 32, "frames": 1397, "bins": 257, "n": 512, "log2n": 9, "signalLength": 224000 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [32, 224000, 1], "dist": "normal", "seed": 2020 },
+        "window": { "dtype": "float32", "shape": [512], "dist": "uniform", "seed": 2022 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [32, 1397, 257, 2] } },
+      "bench": {
+        "metrics": [
+          { "type": "gflops", "value": "5 * args.batch * args.frames * args.n * args.log2n" },
+          {
+            "type": "bandwidth",
+            "value": "4 * (args.batch * args.signalLength + args.n + args.batch * args.frames * args.bins * 2)"
+          }
+        ]
+      },
+      "args": { "frameStep": 160 }
+    },
+    {
+      "name": "large_window_FFT_win4096_signal131072_step1024",
+      "preset": "smoke",
+      "vars": { "batch": 1, "frames": 125, "bins": 2049, "n": 4096, "log2n": 12 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [1, 131072, 1], "dist": "normal", "seed": 2030 },
+        "window": { "dtype": "float32", "shape": [4096], "dist": "uniform", "seed": 2031 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 125, 2049, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.batch * args.frames * args.n * args.log2n" }] },
+      "args": { "frameStep": 1024 }
+    },
+    {
+      "name": "large_window_FFT_win8192_signal131072_step2048",
+      "preset": "stress",
+      "vars": { "batch": 1, "frames": 61, "bins": 4097, "n": 8192, "log2n": 13 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [1, 131072, 1], "dist": "normal", "seed": 2030 },
+        "window": { "dtype": "float32", "shape": [8192], "dist": "uniform", "seed": 2033 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 61, 4097, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.batch * args.frames * args.n * args.log2n" }] },
+      "args": { "frameStep": 2048 }
+    },
+    {
+      "name": "large_window_sibling_FFT_win2048_signal131072_step1024",
+      "preset": "smoke",
+      "vars": { "batch": 1, "frames": 127, "bins": 1025, "n": 2048, "log2n": 11 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [1, 131072, 1], "dist": "normal", "seed": 2030 },
+        "window": { "dtype": "float32", "shape": [2048], "dist": "uniform", "seed": 2032 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 127, 1025, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.batch * args.frames * args.n * args.log2n" }] },
+      "args": { "frameStep": 1024 }
+    },
+    {
+      "name": "mixed_radix_win1000_tail125_batch2_step250_twosided",
+      "preset": "stress",
+      "vars": { "batch": 2, "frames": 13, "bins": 1000, "n": 1000 },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [2, 4000, 1], "dist": "normal", "seed": 3120 },
+        "window": { "dtype": "float32", "shape": [1000], "dist": "uniform", "seed": 3121 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 13, 1000, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.batch * args.frames * args.bins * args.n" }] },
+      "args": { "frameStep": 250 }
+    },
+    {
+      "name": "tiled_win2047_nonpow2_boundary_batch2_step256",
+      "preset": "stress",
+      "vars": { "batch": 2, "frames": 30, "bins": 1024, "n": 2047 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [2, 9471, 1], "dist": "normal", "seed": 3101 },
+        "window": { "dtype": "float32", "shape": [2047], "dist": "uniform", "seed": 3102 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 30, 1024, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.batch * args.frames * args.bins * args.n" }] },
+      "args": { "frameStep": 256 }
+    },
+    {
+      "name": "mixed_radix_win1536_batch4_step384",
+      "preset": "stress",
+      "vars": { "batch": 4, "frames": 13, "bins": 769, "n": 1536 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [4, 6144, 1], "dist": "normal", "seed": 3110 },
+        "window": { "dtype": "float32", "shape": [1536], "dist": "uniform", "seed": 3111 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 13, 769, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.batch * args.frames * args.bins * args.n" }] },
+      "args": { "frameStep": 384 }
+    },
+    {
+      "name": "tiled_win513_nonpow2_boundary_batch2_signal8193_step160",
+      "preset": "stress",
+      "vars": { "batch": 2, "frames": 49, "bins": 257, "n": 513 },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [2, 8193, 1], "dist": "normal", "seed": 3120 },
+        "window": { "dtype": "float32", "shape": [513], "dist": "uniform", "seed": 3121 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 49, 257, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.batch * args.frames * args.bins * args.n" }] },
+      "args": { "frameStep": 160 }
+    },
+    {
+      "name": "audio-no-window-frame512-batch8-signal160000-pathology",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Large no-window transform that exercises rank3_no_window; the windowed FFT control isolates the optional-window route gap."
+      },
+      "vars": { "batch": 8, "frames": 997, "bins": 257, "n": 512, "log2n": 9 },
+      "attrs": { "onesided": 1 },
+      "inputs": { "signal": { "dtype": "float32", "shape": [8, 160000, 1], "dist": "normal", "seed": 3130 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [8, 997, 257, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.batch * args.frames * args.n * args.log2n" }] },
+      "args": { "frameStep": 160, "frameLength": 512 }
+    }
+  ]
+}

build/webgpu/dft-tiled-real.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,140 @@

+// Portable direct transform for large real signals. A 256-lane workgroup owns
+// 256 adjacent frequency bins and stages one input row or windowed STFT frame
+// once, avoiding a global reload per (frequency, sample) pair. Twiddles advance
+// by complex multiplication and reset from sin/cos every 64 samples to bound
+// phase drift without transcendental functions in the inner loop.
+{{ env.wgsl.resourceDeclarations }}
+// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
+// rational turn and route all polynomial evaluation through the shared precise
+// trig core so FFT, direct DFT, STFT, and cosine windows agree exactly.
+//
+// phase is in [0, period). sign is -1 for a forward transform and +1 for an
+// inverse transform. The result is (cos, sin).
+// Backend-stable f32 sine/cosine core.
+//
+// Shader transcendental accuracy is implementation-defined, and some portable
+// backends are only accurate to roughly 1e-4. Each path retains the most accurate
+// available phase representation, reduces it to [-pi, pi], then uses these
+// polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
+// truncation error is well below one f32 ULP over its documented interval.
+const PRECISE_TRIG_PI: f32 = 3.141592653589793;
+const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
+const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
+// Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
+fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
+  let x2 = x * x;
+  var sinPolynomial = 1.6059043836821613e-10;
+  sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
+  sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
+  sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
+  sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
+  sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
+  let sine = x * fma(sinPolynomial, x2, 1.0);
+  var cosPolynomial = 2.08767569878681e-9;
+  cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
+  cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
+  cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
+  cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
+  cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
+  let cosine = fma(cosPolynomial, x2, 1.0);
+  return vec2<f32>(cosine, sine);
+}
+// Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
+fn precise_sincos_centered(x: f32) -> vec2<f32> {
+  var folded = x;
+  var cosineSign = 1.0;
+  if (folded > PRECISE_TRIG_HALF_PI) {
+    folded = PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  } else if (folded < -PRECISE_TRIG_HALF_PI) {
+    folded = -PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  }
+  let value = precise_sincos_half_pi(folded);
+  return vec2<f32>(cosineSign * value.x, value.y);
+}
+fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
+  let reducedPhase = phase % period;
+  if (reducedPhase == 0u) {
+    return vec2<f32>(1.0, sign * 0.0);
+  }
+  if (period % 4u == 0u) {
+    let quarter = period / 4u;
+    if (reducedPhase == quarter) {
+      return vec2<f32>(0.0, sign);
+    }
+    if (reducedPhase == 2u * quarter) {
+      return vec2<f32>(-1.0, sign * 0.0);
+    }
+    if (reducedPhase == 3u * quarter) {
+      return vec2<f32>(0.0, -sign);
+    }
+  }
+  // Center the rational turn before multiplying by 2*pi. This avoids the
+  // cancellation in `2*pi*(N-1)/N - 2*pi` for bins near the end of a row.
+  let turns = f32(reducedPhase) / f32(period);
+  let centeredTurns = turns - floor(turns + 0.5);
+  return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
+}
+const N: u32 = {{ source.transformLength }}u;
+const NUM_BINS: u32 = {{ source.numBins }}u;
+const WG: u32 = {{ source.workgroupSize }}u;
+const TWIDDLE_RESET: u32 = 64u;
+var<workgroup> signalTile: array<f32, N>;
+@compute @workgroup_size(WG, 1, 1)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>,
+) {
+  let lane = lid.x;
+  let frame = wg.y;
+  let batch = wg.z;
+  let start = frame * params.frameStep;
+  for (var n = lane; n < N; n = n + WG) {
+    let signalIndex = start + n;
+    signalTile[n] = signal[batch * params.signalLength + signalIndex] * window[n];
+  }
+  workgroupBarrier();
+  let freq = wg.x * WG + lane;
+  if (freq >= NUM_BINS) {
+    return;
+  }
+  let stepTwiddle = dft_precise_twiddle(freq, N, -1.0);
+  var real = 0.0;
+  var imag = 0.0;
+  for (var base = 0u; base < N; base = base + TWIDDLE_RESET) {
+    // Reduce the reset phase in integer space first. Computing angleStep*base
+    // directly loses enough f32 precision at high bins/long signals to create
+    // visible drift even though each 64-sample recurrence remains stable.
+    let phase = (freq * base) % N;
+    var twiddle = dft_precise_twiddle(phase, N, -1.0);
+    let end = min(N, base + TWIDDLE_RESET);
+    for (var n = base; n < end; n = n + 1u) {
+      let value = signalTile[n];
+      real = real + value * twiddle.x;
+      imag = imag + value * twiddle.y;
+      twiddle = vec2<f32>(
+        twiddle.x * stepTwiddle.x - twiddle.y * stepTwiddle.y,
+        twiddle.x * stepTwiddle.y + twiddle.y * stepTwiddle.x,
+      );
+    }
+  }
+  let outputBase = ((batch * params.numFrames + frame) * NUM_BINS + freq) * 2u;
+  output[outputBase] = real;
+  output[outputBase + 1u] = imag;
+}

build/webgpu/fft-radix2-dit-shared-inplace.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,149 @@

+// Workgroup-resident in-place radix-2 DIT FFT.
+//
+// This route fills the gap between the double-buffered Stockham kernel
+// (16 bytes/element) and the global-scratch DIT fallback. A single complex
+// array needs only 8 bytes/element, extending the resident FFT range one
+// power-of-two rung beyond the 16-byte/element Stockham kernel on every device.
+// Each workgroup owns one transform; all stages remain local and use uniform
+// barriers, so only the input and requested spectrum touch global memory.
+// STFT is forward-only, so the transform sign is fixed at -1.
+{{ env.wgsl.resourceDeclarations }}
+// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
+// rational turn and route all polynomial evaluation through the shared precise
+// trig core so FFT, direct DFT, STFT, and cosine windows agree exactly.
+//
+// phase is in [0, period). sign is -1 for a forward transform and +1 for an
+// inverse transform. The result is (cos, sin).
+// Backend-stable f32 sine/cosine core.
+//
+// Shader transcendental accuracy is implementation-defined, and some portable
+// backends are only accurate to roughly 1e-4. Each path retains the most accurate
+// available phase representation, reduces it to [-pi, pi], then uses these
+// polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
+// truncation error is well below one f32 ULP over its documented interval.
+const PRECISE_TRIG_PI: f32 = 3.141592653589793;
+const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
+const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
+// Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
+fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
+  let x2 = x * x;
+  var sinPolynomial = 1.6059043836821613e-10;
+  sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
+  sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
+  sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
+  sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
+  sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
+  let sine = x * fma(sinPolynomial, x2, 1.0);
+  var cosPolynomial = 2.08767569878681e-9;
+  cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
+  cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
+  cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
+  cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
+  cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
+  let cosine = fma(cosPolynomial, x2, 1.0);
+  return vec2<f32>(cosine, sine);
+}
+// Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
+fn precise_sincos_centered(x: f32) -> vec2<f32> {
+  var folded = x;
+  var cosineSign = 1.0;
+  if (folded > PRECISE_TRIG_HALF_PI) {
+    folded = PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  } else if (folded < -PRECISE_TRIG_HALF_PI) {
+    folded = -PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  }
+  let value = precise_sincos_half_pi(folded);
+  return vec2<f32>(cosineSign * value.x, value.y);
+}
+fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
+  let reducedPhase = phase % period;
+  if (reducedPhase == 0u) {
+    return vec2<f32>(1.0, sign * 0.0);
+  }
+  if (period % 4u == 0u) {
+    let quarter = period / 4u;
+    if (reducedPhase == quarter) {
+      return vec2<f32>(0.0, sign);
+    }
+    if (reducedPhase == 2u * quarter) {
+      return vec2<f32>(-1.0, sign * 0.0);
+    }
+    if (reducedPhase == 3u * quarter) {
+      return vec2<f32>(0.0, -sign);
+    }
+  }
+  // Center the rational turn before multiplying by 2*pi. This avoids the
+  // cancellation in `2*pi*(N-1)/N - 2*pi` for bins near the end of a row.
+  let turns = f32(reducedPhase) / f32(period);
+  let centeredTurns = turns - floor(turns + 0.5);
+  return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
+}
+const N: u32 = {{ source.fftLength }}u;
+const N_HALF: u32 = N / 2u;
+const NUM_BINS: u32 = {{ source.numBins }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const SIGN: f32 = -1.0;
+var<workgroup> fftData: array<vec2<f32>, {{ source.fftLength }}>;
+@compute @workgroup_size(WG, 1, 1)
+fn main(@builtin(workgroup_id) wgid: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let transform = wgid.x + wgid.y * nwg.x;
+  let batch = transform / params.numFrames;
+  let frame = transform - batch * params.numFrames;
+  let start = frame * params.frameStep;
+  for (var position = lid.x; position < N; position = position + WG) {
+    let n = reverseBits(position) >> (32u - {{ source.log2Length }}u);
+    let signalIndex = start + n;
+    let inputBase = (batch * params.signalLength + signalIndex) * {{ source.inComp }}u;
+    let w = window[n];
+    var value = vec2<f32>(signal[inputBase] * w, 0.0);
+{% if source.inComp == 2 %}
+    value.y = signal[inputBase + 1u] * w;
+{% endif %}
+    fftData[position] = value;
+  }
+  workgroupBarrier();
+{% for stage in range(1, source.log2Length + 1) %}
+  {
+    const SPAN: u32 = 1u << {{ stage }}u;
+    const HALF: u32 = SPAN >> 1u;
+    for (var butterfly = lid.x; butterfly < N_HALF; butterfly = butterfly + WG) {
+      let group = butterfly / HALF;
+      let j = butterfly - group * HALF;
+      let lo = group * SPAN + j;
+      let hi = lo + HALF;
+      let a = fftData[lo];
+      let b = fftData[hi];
+      let twiddle = dft_precise_twiddle(j, SPAN, SIGN);
+      let t = vec2<f32>(
+        b.x * twiddle.x - b.y * twiddle.y,
+        b.x * twiddle.y + b.y * twiddle.x,
+      );
+      fftData[lo] = a + t;
+      fftData[hi] = a - t;
+    }
+    workgroupBarrier();
+  }
+{% endfor %}
+  for (var bin = lid.x; bin < NUM_BINS; bin = bin + WG) {
+    let dst = (transform * NUM_BINS + bin) * 2u;
+    output[dst] = fftData[bin].x;
+    output[dst + 1u] = fftData[bin].y;
+  }
+}

build/webgpu/fft-radix2-dit-storage-pass.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,215 @@

+// Globally staged radix-2 DIT FFT. A single-workgroup storage FFT is efficient
+// when a device can keep a whole transform resident, but architectures that
+// serialize storage barriers or spill a large workgroup can become severely
+// under-occupied. This route exposes every transform element or butterfly to
+// the global scheduler. Separate dispatches provide the required stage barriers
+// while preserving the same bit-reversed DIT algorithm and scratch layout.
+// STFT is forward-only, so the transform sign is fixed at -1.
+{{ env.wgsl.resourceDeclarations }}
+{% if source.phase == "local_prefix" or source.phase == "stage" %}
+// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
+// rational turn and route all polynomial evaluation through the shared precise
+// trig core so FFT, direct DFT, STFT, and cosine windows agree exactly.
+//
+// phase is in [0, period). sign is -1 for a forward transform and +1 for an
+// inverse transform. The result is (cos, sin).
+{% set needPreciseTrigCentered = true %}
+{% set needPreciseTrigTwoPi = true %}
+// Backend-stable f32 sine/cosine core.
+//
+// Shader transcendental accuracy is implementation-defined, and some portable
+// backends are only accurate to roughly 1e-4. Each path retains the most accurate
+// available phase representation, reduces it to [-pi, pi], then uses these
+// polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
+// truncation error is well below one f32 ULP over its documented interval.
+{% if needPreciseTrigCentered %}
+const PRECISE_TRIG_PI: f32 = 3.141592653589793;
+{% endif %}
+{% if needPreciseTrigTwoPi %}
+const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
+{% endif %}
+{% if needPreciseTrigCentered %}
+const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
+{% endif %}
+// Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
+fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
+  let x2 = x * x;
+  var sinPolynomial = 1.6059043836821613e-10;
+  sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
+  sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
+  sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
+  sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
+  sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
+  let sine = x * fma(sinPolynomial, x2, 1.0);
+  var cosPolynomial = 2.08767569878681e-9;
+  cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
+  cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
+  cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
+  cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
+  cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
+  let cosine = fma(cosPolynomial, x2, 1.0);
+  return vec2<f32>(cosine, sine);
+}
+{% if needPreciseTrigCentered %}
+// Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
+fn precise_sincos_centered(x: f32) -> vec2<f32> {
+  var folded = x;
+  var cosineSign = 1.0;
+  if (folded > PRECISE_TRIG_HALF_PI) {
+    folded = PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  } else if (folded < -PRECISE_TRIG_HALF_PI) {
+    folded = -PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  }
+  let value = precise_sincos_half_pi(folded);
+  return vec2<f32>(cosineSign * value.x, value.y);
+}
+{% endif %}
+fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
+  let reducedPhase = phase % period;
+  if (reducedPhase == 0u) {
+    return vec2<f32>(1.0, sign * 0.0);
+  }
+  if (period % 4u == 0u) {
+    let quarter = period / 4u;
+    if (reducedPhase == quarter) {
+      return vec2<f32>(0.0, sign);
+    }
+    if (reducedPhase == 2u * quarter) {
+      return vec2<f32>(-1.0, sign * 0.0);
+    }
+    if (reducedPhase == 3u * quarter) {
+      return vec2<f32>(0.0, -sign);
+    }
+  }
+  // Center the rational turn before multiplying by 2*pi. This avoids the
+  // cancellation in `2*pi*(N-1)/N - 2*pi` for bins near the end of a row.
+  let turns = f32(reducedPhase) / f32(period);
+  let centeredTurns = turns - floor(turns + 0.5);
+  return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
+}
+{% endif %}
+const N: u32 = {{ source.fftLength }}u;
+{% if source.phase == "stage" %}
+const N_HALF: u32 = N / 2u;
+{% endif %}
+{% if source.phase == "finish" %}
+const NUM_BINS: u32 = {{ source.numBins }}u;
+{% endif %}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+{% if source.phase == "stage" %}
+const SPAN: u32 = 1u << {{ source.stage }}u;
+const HALF: u32 = SPAN >> 1u;
+const SIGN: f32 = -1.0;
+{% elif source.phase == "local_prefix" %}
+const LOCAL_N: u32 = {{ source.localLength }}u;
+const BLOCKS_PER_TRANSFORM: u32 = N / LOCAL_N;
+const SIGN: f32 = -1.0;
+var<workgroup> localData: array<vec2<f32>, {{ source.localLength }}>;
+{% endif %}
+{% if source.phase != "local_prefix" %}fn folded_global_index(gid: vec3<u32>, nwg: vec3<u32>) -> u32 {
+  return gid.x + gid.y * nwg.x * WG;
+}{% endif %}
+@compute @workgroup_size(WG, 1, 1)
+fn main({% if source.phase != "local_prefix" %}@builtin(global_invocation_id) gid: vec3<u32>,
+        {% else %}@builtin(local_invocation_id) lid: vec3<u32>,
+        @builtin(workgroup_id) wgid: vec3<u32>,
+        {% endif %}
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+{% if source.phase != "local_prefix" %}
+  let flat = folded_global_index(gid, nwg);
+{% endif %}
+  let transforms = arrayLength(&scratch) / (N * 2u);
+{% if source.phase == "local_prefix" %}
+  let flatGroup = wgid.x + wgid.y * nwg.x;
+  let transform = flatGroup / BLOCKS_PER_TRANSFORM;
+  if (transform >= transforms) { return; }
+  let block = flatGroup - transform * BLOCKS_PER_TRANSFORM;
+  let position = block * LOCAL_N + lid.x;
+  let n = reverseBits(position) >> (32u - {{ source.log2Length }}u);
+  let batch = transform / params.numFrames;
+  let frame = transform - batch * params.numFrames;
+  let start = frame * params.frameStep;
+  let signalIndex = start + n;
+  let inputBase = (batch * params.signalLength + signalIndex) * {{ source.inComp }}u;
+  let w = window[n];
+  var value = vec2<f32>(signal[inputBase] * w, 0.0);
+{% if source.inComp == 2 %}
+  value.y = signal[inputBase + 1u] * w;
+{% endif %}
+  localData[lid.x] = value;
+  workgroupBarrier();
+{% for stage in range(1, source.localLog2Length + 1) %}
+  if (lid.x < LOCAL_N / 2u) {
+    const LOCAL_SPAN_{{ stage }}: u32 = 1u << {{ stage }}u;
+    const LOCAL_HALF_{{ stage }}: u32 = LOCAL_SPAN_{{ stage }} >> 1u;
+    let group = lid.x / LOCAL_HALF_{{ stage }};
+    let j = lid.x - group * LOCAL_HALF_{{ stage }};
+    let lo = group * LOCAL_SPAN_{{ stage }} + j;
+    let hi = lo + LOCAL_HALF_{{ stage }};
+    let a = localData[lo];
+    let b = localData[hi];
+    let twiddle = dft_precise_twiddle(j, LOCAL_SPAN_{{ stage }}, SIGN);
+    let t = vec2<f32>(
+      b.x * twiddle.x - b.y * twiddle.y,
+      b.x * twiddle.y + b.y * twiddle.x,
+    );
+    localData[lo] = a + t;
+    localData[hi] = a - t;
+  }
+  workgroupBarrier();
+{% endfor %}
+  let dst = (transform * N + position) * 2u;
+  scratch[dst] = localData[lid.x].x;
+  scratch[dst + 1u] = localData[lid.x].y;
+{% elif source.phase == "stage" %}
+  let count = transforms * N_HALF;
+  if (flat >= count) { return; }
+  let transform = flat / N_HALF;
+  let butterfly = flat - transform * N_HALF;
+  let group = butterfly / HALF;
+  let j = butterfly - group * HALF;
+  let lo = (transform * N + group * SPAN + j) * 2u;
+  let hi = lo + HALF * 2u;
+  let a = vec2<f32>(scratch[lo], scratch[lo + 1u]);
+  let b = vec2<f32>(scratch[hi], scratch[hi + 1u]);
+  let twiddle = dft_precise_twiddle(j, SPAN, SIGN);
+  let t = vec2<f32>(
+    b.x * twiddle.x - b.y * twiddle.y,
+    b.x * twiddle.y + b.y * twiddle.x,
+  );
+  let sum = a + t;
+  let difference = a - t;
+  scratch[lo] = sum.x;
+  scratch[lo + 1u] = sum.y;
+  scratch[hi] = difference.x;
+  scratch[hi + 1u] = difference.y;
+{% else %}
+  let count = transforms * NUM_BINS;
+  if (flat >= count) { return; }
+  let transform = flat / NUM_BINS;
+  let bin = flat - transform * NUM_BINS;
+  let src = (transform * N + bin) * 2u;
+  let dst = (transform * NUM_BINS + bin) * 2u;
+  output[dst] = scratch[src];
+  output[dst + 1u] = scratch[src + 1u];
+{% endif %}
+}

build/webgpu/fft-radix2-dit-storage.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,84 @@

+// Portable single-dispatch radix-2 DIT FFT for transforms too large for the
+// guaranteed workgroup-storage budget. One workgroup owns each transform and
+// uses either its full-spectrum output or a private global scratch row as an
+// in-place complex buffer. Uniform storageBarrier calls replace the global
+// barrier that would otherwise require one dispatch per FFT stage.
+//
+// Direct output is used only when each transform owns a contiguous full-spectrum
+// row. Scratch-backed DFT supports one-sided output. STFT accepts real or
+// complex windowed frames and copies NUM_BINS after the final stage.
+{{ env.wgsl.resourceDeclarations }}
+{% set STFT = source.mode == "stft" %}
+const PI: f32 = 3.141592653589793;
+const N: u32 = {{ source.fftLength }}u;
+const N_HALF: u32 = N / 2u;
+const LOG2_N: u32 = {{ source.log2Length }}u;
+const NUM_BINS: u32 = {{ source.numBins }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const SIGN: f32 = {{ "1.0" if source.inverse else "-1.0" }};
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(workgroup_id) wgid: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>,
+) {
+  let tid = lid.x;
+  let transform = wgid.x + wgid.y * nwg.x;
+  let transformBase = transform * N * 2u;
+  let batch = transform / params.numFrames;
+  let frame = transform - batch * params.numFrames;
+  let start = frame * params.frameStep;
+  for (var n = tid; n < N; n = n + WG) {
+    var value = vec2<f32>(0.0);
+    let signalIndex = start + n;
+    let inputBase = (batch * params.signalLength + signalIndex) * {{ source.inComp }}u;
+    let w = window[n];
+    value.x = signal[inputBase] * w;
+{% if source.inComp == 2 %}
+    value.y = signal[inputBase + 1u] * w;
+{% endif %}
+    let reversed = reverseBits(n) >> (32u - LOG2_N);
+    let dst = transformBase + reversed * 2u;
+    scratch[dst] = value.x;
+    scratch[dst + 1u] = value.y;
+  }
+  storageBarrier();
+{% for stage in range(1, source.log2Length + 1) %}
+  {
+    const SPAN: u32 = 1u << {{ stage }}u;
+    const HALF: u32 = SPAN >> 1u;
+    for (var butterfly = tid; butterfly < N_HALF; butterfly = butterfly + WG) {
+      let group = butterfly / HALF;
+      let j = butterfly - group * HALF;
+      let lo = transformBase + (group * SPAN + j) * 2u;
+      let hi = lo + HALF * 2u;
+      let a = vec2<f32>(scratch[lo], scratch[lo + 1u]);
+      let b = vec2<f32>(scratch[hi], scratch[hi + 1u]);
+      let angle = SIGN * 2.0 * PI * f32(j) / f32(SPAN);
+      let c = cos(angle);
+      let sn = sin(angle);
+      let t = vec2<f32>(b.x * c - b.y * sn, b.x * sn + b.y * c);
+      let sum = a + t;
+      let difference = a - t;
+      scratch[lo] = sum.x;
+      scratch[lo + 1u] = sum.y;
+      scratch[hi] = difference.x;
+      scratch[hi + 1u] = difference.y;
+    }
+    storageBarrier();
+  }
+{% endfor %}
+  const SCALE: f32 = 1.0;
+  for (var k = tid; k < NUM_BINS; k = k + WG) {
+    let src = transformBase + k * 2u;
+    let dst = (transform * NUM_BINS + k) * 2u;
+    output[dst] = scratch[src] * SCALE;
+    output[dst + 1u] = scratch[src + 1u] * SCALE;
+  }
+}

build/webgpu/fft-real-packed-stockham-shared.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,156 @@

+{{ env.wgsl.resourceDeclarations }}
+// Real-input RFFT specialization. Pack even/odd samples into one N/2-point
+// complex transform, then reconstruct the N/2+1 non-redundant bins. This cuts
+// the Stockham butterfly count almost in half while preserving the same fused
+// window load as the full complex STFT kernel.
+const N: u32 = {{ source.fftLength }}u;
+const PACKED_N: u32 = N / 2u;
+const PACKED_HALF: u32 = PACKED_N / 2u;
+const NUM_BINS: u32 = PACKED_N + 1u;
+const FFT_LANES: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+const TRANSFORMS_PER_WG: u32 = {{ source.transformsPerWorkgroup | default(1) }}u;
+const WG: u32 = FFT_LANES * TRANSFORMS_PER_WG;
+const STAGE_TWIDDLES: u32 = PACKED_N - 1u;
+const SHARED_STRIDE: u32 = 2u * PACKED_N;
+var<workgroup> shRe: array<f32, SHARED_STRIDE * TRANSFORMS_PER_WG>;
+var<workgroup> shIm: array<f32, SHARED_STRIDE * TRANSFORMS_PER_WG>;
+fn loadPacked(frameBase: u32, packed: u32, validTransform: bool) -> vec2<f32> {
+  let even = 2u * packed;
+  let odd = even + 1u;
+  var value = vec2<f32>(0.0);
+  if (validTransform) {
+    value.x = signal[frameBase + even] * window[even];
+    value.y = signal[frameBase + odd] * window[odd];
+  }
+  return value;
+}
+@compute @workgroup_size(WG)
+fn main(@builtin(workgroup_id) wgid: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
+  let transform = lid.x / FFT_LANES;
+  let tid = lid.x - transform * FFT_LANES;
+  let batch = wgid.y;
+  let frame = wgid.x * TRANSFORMS_PER_WG + transform;
+  let validTransform = frame < params.numFrames;
+  let start = frame * params.frameStep;
+  let frameBase = batch * params.signalLength + start;
+  let sharedBase = transform * SHARED_STRIDE;
+  // Fuse pairs of radix-2 Stockham stages into radix-4 butterflies. An odd
+  // packed log2 length starts with one radix-2 stage, after which every stage
+  // can be paired. The second stage's root B is sufficient: A=B^2 and C=B^3.
+  // The first butterfly loads windowed samples directly, avoiding a separate
+  // shared-memory load phase and its barrier.
+  var srcOff = sharedBase + PACKED_N;
+  var dstOff = sharedBase;
+  var radix = 1u;
+  if ((countTrailingZeros(PACKED_N) & 1u) != 0u) {
+    for (var j = tid; j < PACKED_HALF; j = j + FFT_LANES) {
+      let a = loadPacked(frameBase, j, validTransform);
+      let b = loadPacked(frameBase, PACKED_HALF + j, validTransform);
+      shRe[srcOff + 2u * j] = a.x + b.x;
+      shIm[srcOff + 2u * j] = a.y + b.y;
+      shRe[srcOff + 2u * j + 1u] = a.x - b.x;
+      shIm[srcOff + 2u * j + 1u] = a.y - b.y;
+    }
+    workgroupBarrier();
+    radix = 2u;
+  } else {
+    for (var j = tid; j < PACKED_N / 4u; j = j + FFT_LANES) {
+      let x0 = loadPacked(frameBase, j, validTransform);
+      let x1 = loadPacked(frameBase, PACKED_N / 4u + j, validTransform);
+      let x2 = loadPacked(frameBase, PACKED_N / 2u + j, validTransform);
+      let x3 = loadPacked(frameBase, 3u * PACKED_N / 4u + j, validTransform);
+      let evenSum = x0 + x2;
+      let evenDiff = x0 - x2;
+      let oddSum = x1 + x3;
+      let oddDiff = x1 - x3;
+      shRe[srcOff + 4u * j] = evenSum.x + oddSum.x;
+      shIm[srcOff + 4u * j] = evenSum.y + oddSum.y;
+      shRe[srcOff + 4u * j + 1u] = evenDiff.x + oddDiff.y;
+      shIm[srcOff + 4u * j + 1u] = evenDiff.y - oddDiff.x;
+      shRe[srcOff + 4u * j + 2u] = evenSum.x - oddSum.x;
+      shIm[srcOff + 4u * j + 2u] = evenSum.y - oddSum.y;
+      shRe[srcOff + 4u * j + 3u] = evenDiff.x - oddDiff.y;
+      shIm[srcOff + 4u * j + 3u] = evenDiff.y + oddDiff.x;
+    }
+    workgroupBarrier();
+    radix = 4u;
+  }
+  for (; radix < PACKED_N; radix = radix * 4u) {
+    for (var j = tid; j < PACKED_N / 4u; j = j + FFT_LANES) {
+      let jm = j & (radix - 1u);
+      let outBase = ((j / radix) * (4u * radix)) + jm;
+      let x0r = shRe[srcOff + j];
+      let x0i = shIm[srcOff + j];
+      let x1r = shRe[srcOff + PACKED_N / 4u + j];
+      let x1i = shIm[srcOff + PACKED_N / 4u + j];
+      let x2r = shRe[srcOff + PACKED_N / 2u + j];
+      let x2i = shIm[srcOff + PACKED_N / 2u + j];
+      let x3r = shRe[srcOff + 3u * PACKED_N / 4u + j];
+      let x3i = shIm[srcOff + 3u * PACKED_N / 4u + j];
+      let twiddleIndex = 2u * radix - 1u + jm;
+      let br = twiddles[2u * twiddleIndex];
+      let bi = twiddles[2u * twiddleIndex + 1u];
+      let ar = br * br - bi * bi;
+      let ai = 2.0 * br * bi;
+      let cr = ar * br - ai * bi;
+      let ci = ar * bi + ai * br;
+      let u1r = x1r * br - x1i * bi;
+      let u1i = x1r * bi + x1i * br;
+      let u2r = x2r * ar - x2i * ai;
+      let u2i = x2r * ai + x2i * ar;
+      let u3r = x3r * cr - x3i * ci;
+      let u3i = x3r * ci + x3i * cr;
+      let evenSumR = x0r + u2r;
+      let evenSumI = x0i + u2i;
+      let evenDiffR = x0r - u2r;
+      let evenDiffI = x0i - u2i;
+      let oddSumR = u1r + u3r;
+      let oddSumI = u1i + u3i;
+      let oddDiffR = u1r - u3r;
+      let oddDiffI = u1i - u3i;
+      shRe[dstOff + outBase] = evenSumR + oddSumR;
+      shIm[dstOff + outBase] = evenSumI + oddSumI;
+      shRe[dstOff + outBase + radix] = evenDiffR + oddDiffI;
+      shIm[dstOff + outBase + radix] = evenDiffI - oddDiffR;
+      shRe[dstOff + outBase + 2u * radix] = evenSumR - oddSumR;
+      shIm[dstOff + outBase + 2u * radix] = evenSumI - oddSumI;
+      shRe[dstOff + outBase + 3u * radix] = evenDiffR - oddDiffI;
+      shIm[dstOff + outBase + 3u * radix] = evenDiffI + oddDiffR;
+    }
+    workgroupBarrier();
+    let swap = srcOff;
+    srcOff = dstOff;
+    dstOff = swap;
+  }
+  let outRow = (batch * params.numFrames + frame) * NUM_BINS;
+  for (var k = tid; validTransform && k < NUM_BINS; k = k + FFT_LANES) {
+    let packedK = k & (PACKED_N - 1u);
+    let mirrorK = (PACKED_N - packedK) & (PACKED_N - 1u);
+    let zr = shRe[srcOff + packedK];
+    let zi = shIm[srcOff + packedK];
+    let mr = shRe[srcOff + mirrorK];
+    let mi = shIm[srcOff + mirrorK];
+    let sumRe = zr + mr;
+    let sumIm = zi - mi;
+    let diffRe = zr - mr;
+    let diffIm = zi + mi;
+    let rootIndex = STAGE_TWIDDLES + k;
+    let c = twiddles[2u * rootIndex];
+    let sn = twiddles[2u * rootIndex + 1u];
+    let outBase = (outRow + k) * 2u;
+    output[outBase] = 0.5 * (sumRe + c * diffIm + sn * diffRe);
+    output[outBase + 1u] = 0.5 * (sumIm - c * diffRe + sn * diffIm);
+  }
+}

build/webgpu/fft-stockham-shared.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,292 @@

+{{ env.wgsl.resourceDeclarations }}
+{% set MIXED_TAIL = source.mixedRadixTail if source.mixedRadixTail is defined else false %}
+{% set USE_TWIDDLE_TABLE = source.useTwiddleTable if source.useTwiddleTable is defined else false %}
+{% set NEED_PRECISE_TWIDDLE = not USE_TWIDDLE_TABLE or MIXED_TAIL %}
+{% if NEED_PRECISE_TWIDDLE %}
+// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
+// rational turn and route all polynomial evaluation through the shared precise
+// trig core so FFT, direct DFT, STFT, and cosine windows agree exactly.
+//
+// phase is in [0, period). sign is -1 for a forward transform and +1 for an
+// inverse transform. The result is (cos, sin).
+{% set needPreciseTrigCentered = true %}
+{% set needPreciseTrigTwoPi = true %}
+// Backend-stable f32 sine/cosine core.
+//
+// Shader transcendental accuracy is implementation-defined, and some portable
+// backends are only accurate to roughly 1e-4. Each path retains the most accurate
+// available phase representation, reduces it to [-pi, pi], then uses these
+// polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
+// truncation error is well below one f32 ULP over its documented interval.
+{% if needPreciseTrigCentered %}
+const PRECISE_TRIG_PI: f32 = 3.141592653589793;
+{% endif %}
+{% if needPreciseTrigTwoPi %}
+const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
+{% endif %}
+{% if needPreciseTrigCentered %}
+const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
+{% endif %}
+// Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
+fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
+  let x2 = x * x;
+  var sinPolynomial = 1.6059043836821613e-10;
+  sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
+  sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
+  sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
+  sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
+  sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
+  let sine = x * fma(sinPolynomial, x2, 1.0);
+  var cosPolynomial = 2.08767569878681e-9;
+  cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
+  cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
+  cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
+  cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
+  cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
+  let cosine = fma(cosPolynomial, x2, 1.0);
+  return vec2<f32>(cosine, sine);
+}
+{% if needPreciseTrigCentered %}
+// Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
+fn precise_sincos_centered(x: f32) -> vec2<f32> {
+  var folded = x;
+  var cosineSign = 1.0;
+  if (folded > PRECISE_TRIG_HALF_PI) {
+    folded = PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  } else if (folded < -PRECISE_TRIG_HALF_PI) {
+    folded = -PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  }
+  let value = precise_sincos_half_pi(folded);
+  return vec2<f32>(cosineSign * value.x, value.y);
+}
+{% endif %}
+fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
+  let reducedPhase = phase % period;
+  if (reducedPhase == 0u) {
+    return vec2<f32>(1.0, sign * 0.0);
+  }
+  if (period % 4u == 0u) {
+    let quarter = period / 4u;
+    if (reducedPhase == quarter) {
+      return vec2<f32>(0.0, sign);
+    }
+    if (reducedPhase == 2u * quarter) {
+      return vec2<f32>(-1.0, sign * 0.0);
+    }
+    if (reducedPhase == 3u * quarter) {
+      return vec2<f32>(0.0, -sign);
+    }
+  }
+  // Center the rational turn before multiplying by 2*pi. This avoids the
+  // cancellation in `2*pi*(N-1)/N - 2*pi` for bins near the end of a row.
+  let turns = f32(reducedPhase) / f32(period);
+  let centeredTurns = turns - floor(turns + 0.5);
+  return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
+}
+{% endif %}
+{% set TAIL_ANCHOR = 32 %}
+{% set REANCHOR_TAIL = MIXED_TAIL
+    and source.fftLength / source.radix2Length > TAIL_ANCHOR %}
+// Radix-2 Stockham autosort FFT: one transform per workgroup, entirely in
+// shared memory. log2(N) ping-pong stages alternate between the two halves of
+// a 2*N scratch pair (no bit reversal; the autosort write pattern leaves the
+// final stage in natural order). Threads stride over the N/2 butterflies of
+// each stage with a workgroupBarrier between stages. Large transform batches
+// may bind a compact precomputed twiddle table; smaller batches evaluate the
+// same roots inline and avoid the extra dispatch.
+// Mode "dft" views the input as [outer, N, inner, inComp] and the output as
+// [outer, numBins, inner, 2]; params.inner carries the (runtime) inner count
+// and dispatch.x = outer * inner workgroups. Mode "stft" runs one workgroup
+// per (batch, frame): the window multiply is fused into the frame load. STFT
+// selection admits only complete frames, so every frame load is in bounds. The
+// inverse+onesided (IRFFT) combination is outside this specialization.
+// Mode "logmel" runs one workgroup per spectrogram frame of
+// a real mono signal: the frame load is centered (start - N/2) with reflection
+// below zero and zeros at/after params.audioLength, and the epilogue folds the
+// one-sided power spectrum through packed triangular mel bands into log10
+// mel-major output instead of storing complex bins.
+const N: u32 = {{ source.fftLength }}u;
+const N_HALF: u32 = N / 2u;
+{% if MIXED_TAIL %}
+const RADIX2_N: u32 = {{ source.radix2Length }}u;
+const RADIX2_HALF: u32 = RADIX2_N / 2u;
+const TAIL_RADIX: u32 = N / RADIX2_N;
+{% if REANCHOR_TAIL %}
+const TAIL_ANCHOR: u32 = {{ TAIL_ANCHOR }}u;
+{% endif %}
+{% endif %}
+const NUM_BINS: u32 = {{ source.numBins }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+{% if NEED_PRECISE_TWIDDLE %}
+// Twiddle = exp(SIGN * i * 2*pi * k / span); SIGN flips for the inverse
+// transform, exactly like the naive kernels' `sign` select.
+const SIGN: f32 = {{ "1.0" if source.inverse else "-1.0" }};
+{% endif %}
+{% if MIXED_TAIL %}
+// Walk consecutive powers of one unit root in each mixed-radix tail. A single
+// backend-stable root evaluation per output bin replaces one precise
+// polynomial evaluation per tail residue for every supported factorization.
+fn complex_mul(a: vec2<f32>, b: vec2<f32>) -> vec2<f32> {
+  return vec2<f32>(a.x * b.x - a.y * b.y, a.x * b.y + a.y * b.x);
+}
+{% endif %}
+// Ping-pong scratch: stage reads [srcOff, srcOff+N), writes [dstOff, dstOff+N).
+var<workgroup> shRe: array<f32, 2u * N>;
+var<workgroup> shIm: array<f32, 2u * N>;
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(workgroup_id) wgid: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>,
+        @builtin(num_workgroups) nwg: vec3<u32>) {
+  let tid = lid.x;
+  // Flat workgroup index, 2D-folded: dispatches that would exceed the
+  // maxComputeWorkgroupsPerDimension limit on a single axis split the
+  // one-workgroup-per-(batch,frame) [or per-(outer,inner)] grid across x/y and
+  // recover the flat index here. Collapses to wgid.x when nwg.y == 1 (every DFT
+  // dispatch and any STFT dispatch under the per-axis limit).
+  let flat = wgid.x + wgid.y * nwg.x;
+  let batch = flat / params.numFrames;
+  let frame = flat - batch * params.numFrames;
+  let start = frame * params.frameStep;
+  // Windowed frame load (window multiply fused into the first read). The mixed
+  // path transposes n = n2 + TAIL_RADIX*n1 into contiguous n2 Stockham rows.
+  for (var n = tid; n < N; n = n + WG) {
+{% if MIXED_TAIL %}
+    let sample_n = (n / RADIX2_N) + TAIL_RADIX * (n % RADIX2_N);
+{% else %}
+    let sample_n = n;
+{% endif %}
+    let signal_index = start + sample_n;
+    let base = (batch * params.signalLength + signal_index) * {{ source.inComp }}u;
+{% if source.hasWindow is defined and not source.hasWindow %}
+    let w = 1.0;
+{% else %}
+    let w = window[sample_n];
+{% endif %}
+    let xr = signal[base] * w;
+{% if source.inComp == 2 %}
+    let xi = signal[base + 1u] * w;
+{% else %}
+    let xi = 0.0;
+{% endif %}
+    shRe[n] = xr;
+    shIm[n] = xi;
+  }
+  workgroupBarrier();
+  // Stockham stages. Stage s (ns = 2^s) reads a = src[j], b = src[j + RADIX2_N/2],
+  // applies w = exp(SIGN*i*pi*(j mod ns)/ns) to b, and writes
+  //   dst[(j div ns)*2*ns + (j mod ns)]      = a + w*b
+  //   dst[(j div ns)*2*ns + (j mod ns) + ns] = a - w*b
+  // |angle| <= pi, so twiddle accuracy never depends on large-argument
+  // sin/cos reduction (unlike the naive kernels' k*n angles).
+  var srcOff = 0u;
+  var dstOff = N;
+{% if MIXED_TAIL %}
+  for (var s = 0u; (1u << s) < RADIX2_N; s = s + 1u) {
+    let ns = 1u << s;
+    for (var j = tid; j < N_HALF; j = j + WG) {
+      let sequence = j / RADIX2_HALF;
+      let local_j = j - sequence * RADIX2_HALF;
+      let sequence_base = sequence * RADIX2_N;
+      let jm = local_j & (ns - 1u);
+      let outBase = sequence_base + ((local_j >> s) << (s + 1u)) + jm;
+      let ar = shRe[srcOff + sequence_base + local_j];
+      let ai = shIm[srcOff + sequence_base + local_j];
+      let br = shRe[srcOff + sequence_base + RADIX2_HALF + local_j];
+      let bi = shIm[srcOff + sequence_base + RADIX2_HALF + local_j];
+      let twiddle = dft_precise_twiddle(jm, 2u * ns, SIGN);
+      let c = twiddle.x;
+      let sn = twiddle.y;
+      let tr = br * c - bi * sn;
+      let ti = br * sn + bi * c;
+      shRe[dstOff + outBase] = ar + tr;
+      shIm[dstOff + outBase] = ai + ti;
+      shRe[dstOff + outBase + ns] = ar - tr;
+      shIm[dstOff + outBase + ns] = ai - ti;
+    }
+{% else %}
+  for (var s = 0u; (1u << s) < N; s = s + 1u) {
+    let ns = 1u << s;
+    for (var j = tid; j < N_HALF; j = j + WG) {
+      let jm = j & (ns - 1u);
+      let outBase = ((j >> s) << (s + 1u)) + jm;
+      let ar = shRe[srcOff + j];
+      let ai = shIm[srcOff + j];
+      let br = shRe[srcOff + N_HALF + j];
+      let bi = shIm[srcOff + N_HALF + j];
+{% if USE_TWIDDLE_TABLE %}
+      let twiddleIndex = ns - 1u + jm;
+      let c = twiddles[2u * twiddleIndex];
+      let sn = twiddles[2u * twiddleIndex + 1u];
+{% else %}
+      let twiddle = dft_precise_twiddle(jm, 2u * ns, SIGN);
+      let c = twiddle.x;
+      let sn = twiddle.y;
+{% endif %}
+      let tr = br * c - bi * sn;
+      let ti = br * sn + bi * c;
+      shRe[dstOff + outBase] = ar + tr;
+      shIm[dstOff + outBase] = ai + ti;
+      shRe[dstOff + outBase + ns] = ar - tr;
+      shIm[dstOff + outBase + ns] = ai - ti;
+    }
+{% endif %}
+    workgroupBarrier();
+    let swap = srcOff;
+    srcOff = dstOff;
+    dstOff = swap;
+  }
+  // The final swap leaves the natural-order result at srcOff. The 1/N inverse
+  // scale is a power-of-two reciprocal, so the multiply is exact and matches
+  // the naive kernels' division.
+  let outRow = (batch * params.numFrames + frame) * NUM_BINS;
+  for (var k = tid; k < NUM_BINS; k = k + WG) {
+    let outBase = (outRow + k) * 2u;
+{% if MIXED_TAIL %}
+    let k1 = k % RADIX2_N;
+    var sum_re = 0.0;
+    var sum_im = 0.0;
+    let tail_step = dft_precise_twiddle(k, N, SIGN);
+    var tail_twiddle = vec2<f32>(1.0, SIGN * 0.0);
+    for (var n2 = 0u; n2 < TAIL_RADIX; n2 = n2 + 1u) {
+{% if REANCHOR_TAIL %}
+      // Restart the walk from a precise root every TAIL_ANCHOR steps. At
+      // n2 == 0 this reproduces the initializer exactly (phase 0 returns the
+      // literal unit), so the anchored and unanchored entries agree bit
+      // for bit and only the chain length changes.
+      if (n2 % TAIL_ANCHOR == 0u) {
+        tail_twiddle = dft_precise_twiddle(k * n2, N, SIGN);
+      }
+{% endif %}
+      let v_index = srcOff + n2 * RADIX2_N + k1;
+      let vr = shRe[v_index];
+      let vi = shIm[v_index];
+      sum_re = sum_re + vr * tail_twiddle.x - vi * tail_twiddle.y;
+      sum_im = sum_im + vr * tail_twiddle.y + vi * tail_twiddle.x;
+      tail_twiddle = complex_mul(tail_twiddle, tail_step);
+    }
+    output[outBase] = sum_re;
+    output[outBase + 1u] = sum_im;
+{% else %}
+    output[outBase] = shRe[srcOff + k];
+    output[outBase + 1u] = shIm[srcOff + k];
+{% endif %}
+  }
+}

build/webgpu/fft-stockham-twiddles.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,125 @@

+{{ env.wgsl.resourceDeclarations }}
+// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
+// rational turn and route all polynomial evaluation through the shared precise
+// trig core so FFT, direct DFT, STFT, and cosine windows agree exactly.
+//
+// phase is in [0, period). sign is -1 for a forward transform and +1 for an
+// inverse transform. The result is (cos, sin).
+// Backend-stable f32 sine/cosine core.
+//
+// Shader transcendental accuracy is implementation-defined, and some portable
+// backends are only accurate to roughly 1e-4. Each path retains the most accurate
+// available phase representation, reduces it to [-pi, pi], then uses these
+// polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
+// truncation error is well below one f32 ULP over its documented interval.
+const PRECISE_TRIG_PI: f32 = 3.141592653589793;
+const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
+const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
+// Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
+fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
+  let x2 = x * x;
+  var sinPolynomial = 1.6059043836821613e-10;
+  sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
+  sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
+  sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
+  sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
+  sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
+  let sine = x * fma(sinPolynomial, x2, 1.0);
+  var cosPolynomial = 2.08767569878681e-9;
+  cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
+  cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
+  cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
+  cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
+  cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
+  let cosine = fma(cosPolynomial, x2, 1.0);
+  return vec2<f32>(cosine, sine);
+}
+// Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
+fn precise_sincos_centered(x: f32) -> vec2<f32> {
+  var folded = x;
+  var cosineSign = 1.0;
+  if (folded > PRECISE_TRIG_HALF_PI) {
+    folded = PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  } else if (folded < -PRECISE_TRIG_HALF_PI) {
+    folded = -PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  }
+  let value = precise_sincos_half_pi(folded);
+  return vec2<f32>(cosineSign * value.x, value.y);
+}
+fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
+  let reducedPhase = phase % period;
+  if (reducedPhase == 0u) {
+    return vec2<f32>(1.0, sign * 0.0);
+  }
+  if (period % 4u == 0u) {
+    let quarter = period / 4u;
+    if (reducedPhase == quarter) {
+      return vec2<f32>(0.0, sign);
+    }
+    if (reducedPhase == 2u * quarter) {
+      return vec2<f32>(-1.0, sign * 0.0);
+    }
+    if (reducedPhase == 3u * quarter) {
+      return vec2<f32>(0.0, -sign);
+    }
+  }
+  // Center the rational turn before multiplying by 2*pi. This avoids the
+  // cancellation in `2*pi*(N-1)/N - 2*pi` for bins near the end of a row.
+  let turns = f32(reducedPhase) / f32(period);
+  let centeredTurns = turns - floor(turns + 0.5);
+  return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
+}
+// One compact twiddle table shared by every transform in a large Stockham
+// dispatch. Stage s uses ns=2^s distinct roots at table offset ns-1, matching
+// the exact angle expression in the shared Stockham FFT path. Precomputing
+// each root once avoids repeating sin/cos for every frame while preserving the
+// same f32 twiddle values and butterfly arithmetic.
+// STFT is forward-only, so the transform sign is fixed at -1.
+const N: u32 = {{ source.fftLength }}u;
+const SIGN: f32 = -1.0;
+{% if source.realPacked %}
+const PACKED_N: u32 = N / 2u;
+const STAGE_TWIDDLES: u32 = PACKED_N - 1u;
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let tableIndex = gid.x;
+{% if source.realPacked %}
+  if (tableIndex >= N) {
+    return;
+  }
+  if (tableIndex < STAGE_TWIDDLES) {
+    let ns = 1u << (31u - countLeadingZeros(tableIndex + 1u));
+    let jm = tableIndex - (ns - 1u);
+    let twiddle = dft_precise_twiddle(jm, 2u * ns, SIGN);
+    twiddles[2u * tableIndex] = twiddle.x;
+    twiddles[2u * tableIndex + 1u] = twiddle.y;
+    return;
+  }
+  let k = tableIndex - STAGE_TWIDDLES;
+  let twiddle = dft_precise_twiddle(k, N, SIGN);
+  twiddles[2u * tableIndex] = twiddle.x;
+  twiddles[2u * tableIndex + 1u] = twiddle.y;
+{% else %}
+  if (tableIndex >= N - 1u) {
+    return;
+  }
+  let ns = 1u << (31u - countLeadingZeros(tableIndex + 1u));
+  let jm = tableIndex - (ns - 1u);
+  let twiddle = dft_precise_twiddle(jm, 2u * ns, SIGN);
+  twiddles[2u * tableIndex] = twiddle.x;
+  twiddles[2u * tableIndex + 1u] = twiddle.y;
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,644 @@

+{
+  "domain": "ai.onnx",
+  "name": "STFT",
+  "sinceVersion": 17,
+  "description": "Computes the short-time Fourier transform (STFT) of a rank-3 real or complex signal by sliding a complete frame over the input and applying the DFT. The `frameStep` scalar represents the required positive `frame_step` input. At least one of `window` or `frameLength` is required; when both are supplied, their lengths must match. The output contains the real and imaginary components of every frequency bin for each complete frame.",
+  "inputs": [
+    {
+      "role": "signal",
+      "dtype": "T",
+      "rank": 3,
+      "description": "Real input signal with shape `(batch_size, signal_length, 1)`, or complex input with shape `(batch_size, signal_length, 2)`."
+    },
+    {
+      "role": "window",
+      "dtype": "T",
+      "rank": 1,
+      "optional": true,
+      "description": "Optional 1-D window tensor with shape `(window_length)` to multiply by each frame before the DFT."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": 4,
+      "description": "STFT result with shape `(batch_size, frames, dft_unique_bins, 2)` in one-sided mode or `(batch_size, frames, frame_length, 2)` in two-sided mode, with real and imaginary parts in the last dimension. `frames = floor((signal_length - frame_length) / frame_step) + 1`; partial trailing frames are not emitted.",
+      "shape": ["dim(shapes.signal, 0)", "expectedFrames", "floor(frameLength / 2) + 1 if attrs.onesided == 1 else frameLength", 2]
+    }
+  ],
+  "attributes": { "onesided": 1 },
+  "attributeDescriptions": {
+    "onesided": "When 1 (default), returns `floor(n_fft / 2) + 1` non-redundant frequency bins with indices `[0, ..., floor(n_fft / 2)]` for a real-valued input (RFFT); set to 0 to return the full spectrum."
+  },
+  "attributeConstraints": { "onesided": { "values": [0, 1] } },
+  "typeConstraints": { "T": ["float32"] },
+  "args": {
+    "signal": { "kind": "tensor", "semantic": "signal", "role": "input" },
+    "window": { "kind": "tensor", "semantic": "window", "role": "input", "required": false },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" },
+    "frameStep": {
+      "kind": "i32",
+      "semantic": "frame_step",
+      "role": "attribute",
+      "description": "Positive number of signal positions between the starts of consecutive frames."
+    },
+    "frameLength": {
+      "kind": "i32",
+      "semantic": "frame_length",
+      "role": "attribute",
+      "required": false,
+      "description": "Optional number of signal values in each frame; must equal the window length when both are supplied."
+    }
+  },
+  "tunables": {
+    "WORKGROUP_SIZE": 128,
+    "TWIDDLE_TABLE_MIN_TRANSFORMS": 4096,
+    "TILED_REAL_MIN_LENGTH": 512,
+    "GLOBAL_STAGE_LOCAL_LENGTH": 256,
+    "GLOBAL_STAGE_MAX_DEVICE_STORAGE_BYTES": 16384
+  },
+  "derive": {
+    "windowLength": "dim(shapes.window, 0) if (present.window and ranks.window == 1) else 0",
+    "frameLength": "windowLength if (present.window and ranks.window == 1) else (args.frameLength if args.frameLength is defined else 0)",
+    "frameLengthMatchesWindow": "args.frameLength is undefined or not present.window or args.frameLength == windowLength",
+    "frameControlsValid": "args.frameStep > 0 and frameLength > 0 and frameLength <= dim(shapes.signal, 1) and frameLengthMatchesWindow",
+    "expectedFrames": "floor((dim(shapes.signal, 1) - frameLength) / args.frameStep) + 1 if frameControlsValid else 0",
+    "rank3OutputLayout": "ranks.signal == 3 and ranks.output == 4 and dim(shapes.output, 0) == dim(shapes.signal, 0) and dim(shapes.output, 1) == expectedFrames and dim(shapes.signal, 2) >= 1 and dim(shapes.signal, 2) <= 2 and dim(shapes.output, 3) == 2",
+    "rank3WindowLayout": "present.window and ranks.window == 1 and rank3OutputLayout and frameControlsValid",
+    "rank3WindowSpectrum": "false if not present.window else ((attrs.onesided == 0 and dim(shapes.output, 2) == dim(shapes.window, 0)) or (attrs.onesided != 0 and dim(shapes.signal, 2) == 1 and dim(shapes.output, 2) == floor(dim(shapes.window, 0) / 2) + 1))",
+    "rank3WindowDispatchFits": "dim(shapes.output, 1) <= device.limits.maxComputeWorkgroupsPerDimension and dim(shapes.output, 0) <= device.limits.maxComputeWorkgroupsPerDimension",
+    "rank3WindowSharedFftFits": "false if not present.window else (rank3WindowDispatchFits and pow2ceil(dim(shapes.window, 0)) == dim(shapes.window, 0) and dim(shapes.window, 0) >= 4 and dim(shapes.window, 0) * 16 <= device.limits.maxComputeWorkgroupStorageSize)",
+    "rank3WindowFftRadix2": "1 if not present.window else pick([[dim(shapes.window, 0) % 256 == 0, 256], [dim(shapes.window, 0) % 128 == 0, 128], [dim(shapes.window, 0) % 64 == 0, 64], [dim(shapes.window, 0) % 32 == 0, 32], [dim(shapes.window, 0) % 16 == 0, 16], [dim(shapes.window, 0) % 8 == 0, 8], [dim(shapes.window, 0) % 4 == 0, 4]], 1)",
+    "rank3WindowFftTailRadix": "1 if not present.window else dim(shapes.window, 0) / rank3WindowFftRadix2",
+    "rank3WindowMixedFftFits": "false if not present.window else (rank3WindowDispatchFits and pow2ceil(dim(shapes.window, 0)) != dim(shapes.window, 0) and rank3WindowFftRadix2 >= 4 and rank3WindowFftTailRadix > 1 and rank3WindowFftTailRadix <= 128 and dim(shapes.window, 0) * 16 <= device.limits.maxComputeWorkgroupStorageSize)",
+    "rank3WindowFftLog2": "0 if not present.window else log2ceil(dim(shapes.window, 0))",
+    "rank3WindowRadix2ScratchFits": "false if not present.window else (rank3WindowDispatchFits and pow2ceil(dim(shapes.window, 0)) == dim(shapes.window, 0) and dim(shapes.window, 0) >= 4 and dim(shapes.window, 0) * 16 > device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 0) * dim(shapes.output, 1) > 0 and dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) * 2 * 4 <= device.limits.maxStorageBufferBindingSize and dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) * 2 * 4 <= device.limits.maxBufferSize)",
+    "rank3WindowInplaceFftFits": "false if not present.window else (rank3WindowDispatchFits and pow2ceil(dim(shapes.window, 0)) == dim(shapes.window, 0) and dim(shapes.window, 0) >= 4 and dim(shapes.window, 0) * 16 > device.limits.maxComputeWorkgroupStorageSize and dim(shapes.window, 0) * 8 <= device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 0) * dim(shapes.output, 1) > 0)",
+    "rank3WindowGlobalFftFits": "false if not present.window else (rank3WindowRadix2ScratchFits and dim(shapes.window, 0) >= 2 * tunables.GLOBAL_STAGE_LOCAL_LENGTH and dim(shapes.window, 0) * 8 > device.limits.maxComputeWorkgroupStorageSize)",
+    "rank3WindowTiledRealFits": "false if not present.window else (rank3WindowDispatchFits and ranks.signal == 3 and dim(shapes.signal, 2) == 1 and dim(shapes.window, 0) >= tunables.TILED_REAL_MIN_LENGTH and dim(shapes.window, 0) * 4 <= device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 0) > 0 and dim(shapes.output, 1) > 0 and ceilDiv(dim(shapes.output, 2), 256) <= device.limits.maxComputeWorkgroupsPerDimension)",
+    "packedFourWayFits": "false if not present.window else (device.limits.maxComputeInvocationsPerWorkgroup >= 256 and device.limits.maxComputeWorkgroupSizeX >= 256 and dim(shapes.window, 0) * 32 <= device.limits.maxComputeWorkgroupStorageSize)",
+    "packedTransformsPerWorkgroup": "4 if packedFourWayFits else 1",
+    "rank3NoWindowLayout": "not present.window and args.frameLength is defined and rank3OutputLayout and frameControlsValid",
+    "rank3NoWindowSpectrum": "(attrs.onesided == 0 and dim(shapes.output, 2) == frameLength) or (attrs.onesided != 0 and dim(shapes.signal, 2) == 1 and dim(shapes.output, 2) == floor(frameLength / 2) + 1)",
+    "rank3NoWindowSharedFftFits": "false if ranks.signal != 3 else ((attrs.onesided == 0 or (dim(shapes.signal, 2) == 1 and dim(shapes.output, 2) >= 3)) and pow2ceil(frameLength) == frameLength and frameLength >= 4 and frameLength * 16 <= device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 1) <= device.limits.maxComputeWorkgroupsPerDimension and dim(shapes.output, 0) <= device.limits.maxComputeWorkgroupsPerDimension)"
+  },
+  "bindingSets": {
+    "stftWindowFft": [
+      {
+        "name": "signal",
+        "arg": "signal",
+        "semantic": "signal",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "window",
+        "arg": "window",
+        "semantic": "window",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
+            { "name": "numFrames", "type": "u32", "value": "dim(shapes.output, 1)" },
+            { "name": "frameStep", "type": "u32", "value": "args.frameStep" }
+          ]
+        }
+      }
+    ],
+    "stftTwiddleWrite": [
+      { "name": "twiddles", "semantic": "fft_twiddles", "buffer": { "type": "storage" }, "elementType": "f32" }
+    ],
+    "stftWindowFftTwiddles": [
+      {
+        "name": "signal",
+        "arg": "signal",
+        "semantic": "signal",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "window",
+        "arg": "window",
+        "semantic": "window",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "twiddles",
+        "semantic": "fft_twiddles",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
+            { "name": "numFrames", "type": "u32", "value": "dim(shapes.output, 1)" },
+            { "name": "frameStep", "type": "u32", "value": "args.frameStep" }
+          ]
+        }
+      }
+    ],
+    "stftNoWindowFft": [
+      {
+        "name": "signal",
+        "arg": "signal",
+        "semantic": "signal",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
+            { "name": "numFrames", "type": "u32", "value": "dim(shapes.output, 1)" },
+            { "name": "frameStep", "type": "u32", "value": "args.frameStep" }
+          ]
+        }
+      }
+    ],
+    "stftWindowScratch": [
+      {
+        "name": "signal",
+        "arg": "signal",
+        "semantic": "signal",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "window",
+        "arg": "window",
+        "semantic": "window",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      { "name": "scratch", "semantic": "fft_scratch", "buffer": { "type": "storage" }, "elementType": "f32" },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
+            { "name": "numFrames", "type": "u32", "value": "dim(shapes.output, 1)" },
+            { "name": "frameStep", "type": "u32", "value": "args.frameStep" }
+          ]
+        }
+      }
+    ],
+    "stftWindowScratchLocalPrefix": [
+      {
+        "name": "signal",
+        "arg": "signal",
+        "semantic": "signal",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "window",
+        "arg": "window",
+        "semantic": "window",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      { "name": "scratch", "semantic": "fft_scratch", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
+            { "name": "numFrames", "type": "u32", "value": "dim(shapes.output, 1)" },
+            { "name": "frameStep", "type": "u32", "value": "args.frameStep" }
+          ]
+        }
+      }
+    ],
+    "stftWindowScratchStage": [
+      { "name": "scratch", "semantic": "fft_scratch", "buffer": { "type": "storage" }, "elementType": "f32" }
+    ],
+    "stftWindowScratchFinish": [
+      { "name": "scratch", "semantic": "fft_scratch", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" }
+    ],
+    "stftWindowGeneral": [
+      {
+        "name": "signal",
+        "arg": "signal",
+        "semantic": "signal",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "window",
+        "arg": "window",
+        "semantic": "window",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.output) / 2" },
+            { "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
+            { "name": "signalComponents", "type": "u32", "value": "dim(shapes.signal, 2)" },
+            { "name": "frameLength", "type": "u32", "value": "dim(shapes.window, 0)" },
+            { "name": "frames", "type": "u32", "value": "dim(shapes.output, 1)" },
+            { "name": "bins", "type": "u32", "value": "dim(shapes.output, 2)" },
+            { "name": "frameStep", "type": "u32", "value": "args.frameStep" }
+          ]
+        }
+      }
+    ],
+    "stftNoWindowGeneral": [
+      {
+        "name": "signal",
+        "arg": "signal",
+        "semantic": "signal",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.output) / 2" },
+            { "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
+            { "name": "signalComponents", "type": "u32", "value": "dim(shapes.signal, 2)" },
+            { "name": "frameLength", "type": "u32", "value": "args.frameLength" },
+            { "name": "frames", "type": "u32", "value": "dim(shapes.output, 1)" },
+            { "name": "bins", "type": "u32", "value": "dim(shapes.output, 2)" },
+            { "name": "frameStep", "type": "u32", "value": "args.frameStep" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "rank3_no_window_fft",
+      "priority": 20,
+      "when": ["rank3NoWindowLayout", "rank3NoWindowSpectrum", "rank3NoWindowSharedFftFits"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "STFT.NoWindowFftShared",
+          "source": {
+            "shader": "fft-stockham-shared.wgsl.jinja",
+            "inputs": {
+              "mode": "\"stft\"",
+              "fftLength": "frameLength",
+              "numBins": "dim(shapes.output, 2)",
+              "inComp": "dim(shapes.signal, 2)",
+              "inverse": false,
+              "hasWindow": false
+            }
+          },
+          "bindings": "stftNoWindowFft",
+          "dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
+        }
+      ]
+    },
+    {
+      "id": "rank3_window_fft_shared_inplace",
+      "priority": 32,
+      "requires": { "limits": { "maxComputeInvocationsPerWorkgroup": 256, "maxComputeWorkgroupSizeX": 256 } },
+      "tunables": { "WORKGROUP_SIZE": 256 },
+      "when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowInplaceFftFits"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "STFT.FftSharedInplace",
+          "source": {
+            "shader": "fft-radix2-dit-shared-inplace.wgsl.jinja",
+            "inputs": {
+              "fftLength": "dim(shapes.window, 0)",
+              "log2Length": "rank3WindowFftLog2",
+              "numBins": "dim(shapes.output, 2)",
+              "inComp": "dim(shapes.signal, 2)"
+            }
+          },
+          "bindings": "stftWindowFft",
+          "dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
+        }
+      ]
+    },
+    {
+      "id": "rank3_window_fft_global_stages",
+      "priority": 31,
+      "requires": { "limits": { "maxComputeInvocationsPerWorkgroup": 256, "maxComputeWorkgroupSizeX": 256 } },
+      "tunables": { "WORKGROUP_SIZE": 256 },
+      "demoteWhen": ["device.limits.maxComputeWorkgroupStorageSize > tunables.GLOBAL_STAGE_MAX_DEVICE_STORAGE_BYTES"],
+      "when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowGlobalFftFits"],
+      "intermediates": [
+        {
+          "id": "fft_scratch",
+          "dtype": "float32",
+          "shape": "[dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) * 2]"
+        }
+      ],
+      "passes": [
+        {
+          "id": "local_prefix",
+          "name": "STFT.FftLocalPrefix",
+          "source": {
+            "shader": "fft-radix2-dit-storage-pass.wgsl.jinja",
+            "inputs": {
+              "phase": "\"local_prefix\"",
+              "fftLength": "dim(shapes.window, 0)",
+              "log2Length": "rank3WindowFftLog2",
+              "localLength": "tunables.GLOBAL_STAGE_LOCAL_LENGTH",
+              "localLog2Length": "log2ceil(tunables.GLOBAL_STAGE_LOCAL_LENGTH)",
+              "numBins": "dim(shapes.output, 2)",
+              "inComp": "dim(shapes.signal, 2)"
+            }
+          },
+          "bindings": "stftWindowScratchLocalPrefix",
+          "dispatch": {
+            "workgroups": "dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) / tunables.GLOBAL_STAGE_LOCAL_LENGTH"
+          }
+        },
+        {
+          "id": "stages",
+          "repeat": { "count": "rank3WindowFftLog2 - log2ceil(tunables.GLOBAL_STAGE_LOCAL_LENGTH)", "index": "stage" },
+          "passes": [
+            {
+              "id": "butterflies",
+              "name": "STFT.FftGlobalStage",
+              "source": {
+                "shader": "fft-radix2-dit-storage-pass.wgsl.jinja",
+                "inputs": {
+                  "phase": "\"stage\"",
+                  "stage": "repeat.stage + log2ceil(tunables.GLOBAL_STAGE_LOCAL_LENGTH) + 1",
+                  "fftLength": "dim(shapes.window, 0)",
+                  "log2Length": "rank3WindowFftLog2",
+                  "numBins": "dim(shapes.output, 2)",
+                  "inComp": "dim(shapes.signal, 2)"
+                }
+              },
+              "bindings": "stftWindowScratchStage",
+              "dispatch": {
+                "threads": "dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) / 2",
+                "workgroupSize": "tunables.WORKGROUP_SIZE"
+              }
+            }
+          ]
+        },
+        {
+          "id": "finish",
+          "name": "STFT.FftGlobalFinish",
+          "source": {
+            "shader": "fft-radix2-dit-storage-pass.wgsl.jinja",
+            "inputs": {
+              "phase": "\"finish\"",
+              "fftLength": "dim(shapes.window, 0)",
+              "log2Length": "rank3WindowFftLog2",
+              "numBins": "dim(shapes.output, 2)",
+              "inComp": "dim(shapes.signal, 2)"
+            }
+          },
+          "bindings": "stftWindowScratchFinish",
+          "dispatch": {
+            "threads": "dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.output, 2)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "rank3_window_fft_storage_inplace",
+      "priority": 30,
+      "when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowRadix2ScratchFits"],
+      "intermediates": [
+        {
+          "id": "fft_scratch",
+          "dtype": "float32",
+          "shape": "[dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) * 2]"
+        }
+      ],
+      "passes": [
+        {
+          "id": "main",
+          "name": "STFT.FftStorageInplace",
+          "source": {
+            "shader": "fft-radix2-dit-storage.wgsl.jinja",
+            "inputs": {
+              "mode": "\"stft\"",
+              "fftLength": "dim(shapes.window, 0)",
+              "log2Length": "rank3WindowFftLog2",
+              "numBins": "dim(shapes.output, 2)",
+              "inComp": "dim(shapes.signal, 2)",
+              "inverse": false
+            }
+          },
+          "bindings": "stftWindowScratch",
+          "dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
+        }
+      ]
+    },
+    {
+      "id": "rank3_window_tiled_real",
+      "priority": 19,
+      "requires": { "limits": { "maxComputeInvocationsPerWorkgroup": 256, "maxComputeWorkgroupSizeX": 256 } },
+      "tunables": { "WORKGROUP_SIZE": 256 },
+      "when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowTiledRealFits"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "STFT.TiledReal",
+          "source": {
+            "shader": "dft-tiled-real.wgsl.jinja",
+            "inputs": {
+              "mode": "\"stft\"",
+              "transformLength": "dim(shapes.window, 0)",
+              "numBins": "dim(shapes.output, 2)",
+              "workgroupSize": "tunables.WORKGROUP_SIZE"
+            }
+          },
+          "bindings": "stftWindowFft",
+          "dispatch": {
+            "x": "ceilDiv(dim(shapes.output, 2), 256)",
+            "y": "dim(shapes.output, 1)",
+            "z": "dim(shapes.output, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "rank3_window_rfft_packed",
+      "priority": 22,
+      "tunables": { "WORKGROUP_SIZE": 64 },
+      "when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowSharedFftFits", "attrs.onesided != 0", "dim(shapes.signal, 2) == 1", "dim(shapes.window, 0) >= 64", "dim(shapes.output, 0) * dim(shapes.output, 1) >= tunables.TWIDDLE_TABLE_MIN_TRANSFORMS"],
+      "intermediates": [{ "id": "fft_twiddles", "dtype": "float32", "shape": "[2 * dim(shapes.window, 0)]" }],
+      "passes": [
+        {
+          "id": "twiddles",
+          "name": "STFT.RealPackedTwiddleTable",
+          "source": {
+            "shader": "fft-stockham-twiddles.wgsl.jinja",
+            "inputs": { "fftLength": "dim(shapes.window, 0)", "realPacked": 1 }
+          },
+          "bindings": "stftTwiddleWrite",
+          "dispatch": { "threads": "dim(shapes.window, 0)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "main",
+          "name": "STFT.RealPackedFftShared",
+          "source": {
+            "shader": "fft-real-packed-stockham-shared.wgsl.jinja",
+            "inputs": { "fftLength": "dim(shapes.window, 0)", "transformsPerWorkgroup": "packedTransformsPerWorkgroup" }
+          },
+          "bindings": "stftWindowFftTwiddles",
+          "dispatch": {
+            "x": "ceilDiv(dim(shapes.output, 1), packedTransformsPerWorkgroup)",
+            "y": "dim(shapes.output, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "rank3_window_fft_twiddles",
+      "priority": 21,
+      "when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowSharedFftFits", "dim(shapes.window, 0) >= 64", "dim(shapes.output, 0) * dim(shapes.output, 1) >= tunables.TWIDDLE_TABLE_MIN_TRANSFORMS"],
+      "intermediates": [{ "id": "fft_twiddles", "dtype": "float32", "shape": "[2 * (dim(shapes.window, 0) - 1)]" }],
+      "passes": [
+        {
+          "id": "twiddles",
+          "name": "STFT.FftTwiddleTable",
+          "source": {
+            "shader": "fft-stockham-twiddles.wgsl.jinja",
+            "inputs": { "fftLength": "dim(shapes.window, 0)", "realPacked": 0 }
+          },
+          "bindings": "stftTwiddleWrite",
+          "dispatch": { "threads": "dim(shapes.window, 0) - 1", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "main",
+          "name": "STFT.FftSharedTwiddles",
+          "source": {
+            "shader": "fft-stockham-shared.wgsl.jinja",
+            "inputs": {
+              "mode": "\"stft\"",
+              "fftLength": "dim(shapes.window, 0)",
+              "numBins": "dim(shapes.output, 2)",
+              "inComp": "dim(shapes.signal, 2)",
+              "inverse": false,
+              "useTwiddleTable": true
+            }
+          },
+          "bindings": "stftWindowFftTwiddles",
+          "dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
+        }
+      ]
+    },
+    {
+      "id": "rank3_window_fft_mixed",
+      "priority": 25,
+      "when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowMixedFftFits"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "STFT.FftMixedRecurrent",
+          "source": {
+            "shader": "fft-stockham-shared.wgsl.jinja",
+            "inputs": {
+              "mode": "\"stft\"",
+              "fftLength": "dim(shapes.window, 0)",
+              "radix2Length": "rank3WindowFftRadix2",
+              "mixedRadixTail": true,
+              "numBins": "dim(shapes.output, 2)",
+              "inComp": "dim(shapes.signal, 2)",
+              "inverse": false
+            }
+          },
+          "bindings": "stftWindowFft",
+          "dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
+        }
+      ]
+    },
+    {
+      "id": "rank3_window_fft",
+      "priority": 20,
+      "when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowSharedFftFits"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "STFT.FftShared",
+          "source": {
+            "shader": "fft-stockham-shared.wgsl.jinja",
+            "inputs": {
+              "mode": "\"stft\"",
+              "fftLength": "dim(shapes.window, 0)",
+              "numBins": "dim(shapes.output, 2)",
+              "inComp": "dim(shapes.signal, 2)",
+              "inverse": false
+            }
+          },
+          "bindings": "stftWindowFft",
+          "dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
+        }
+      ]
+    },
+    {
+      "id": "rank3_window",
+      "priority": 0,
+      "supersededBy": ["rank3_window_fft_shared_inplace", "rank3_window_fft_global_stages", "rank3_window_fft_storage_inplace", "rank3_window_fft_mixed", "rank3_window_rfft_packed", "rank3_window_fft_twiddles", "rank3_window_fft", "rank3_window_tiled_real"],
+      "when": ["rank3WindowLayout", "rank3WindowSpectrum"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "STFT",
+          "source": { "shader": "stft-general.wgsl.jinja", "inputs": { "hasWindowInput": true } },
+          "bindings": "stftWindowGeneral",
+          "dispatch": { "threads": "numel(shapes.output) / 2", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "rank3_no_window",
+      "priority": 10,
+      "when": ["rank3NoWindowLayout", "rank3NoWindowSpectrum"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "STFT.GeneralNoWindow",
+          "source": { "shader": "stft-general.wgsl.jinja", "inputs": { "hasWindowInput": false } },
+          "bindings": "stftNoWindowGeneral",
+          "dispatch": { "threads": "numel(shapes.output) / 2", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,25 @@

+{
+  "name": "ai.onnx.STFT",
+  "id": "_ai_onnx_stft_webgpu_285ab84",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "OKXSfrXKgDJBKuWa61pzqlE0q6XeDp3+Ac78kZJzvEs=",
+      "dft-tiled-real.wgsl.jinja": "jWXC4IENyJwpE7YFeiMBTtrmtncyfXddBe6JQBB3+h4=",
+      "fft-radix2-dit-shared-inplace.wgsl.jinja": "IGz/c29DBNlDXBpYQk9d44DHXtmORm86rJSyJxDQ4oU=",
+      "fft-radix2-dit-storage-pass.wgsl.jinja": "kDmHLr6NfPBnO2OPHYAkBgCA/dUzsHq9yIJS/v4eIpk=",
+      "fft-radix2-dit-storage.wgsl.jinja": "rlD1QCYyRVzQf4VEZTseGFYvUzjKNRneHUatxpz5Y9k=",
+      "fft-real-packed-stockham-shared.wgsl.jinja": "w9XapiIdKLjNy/3BwXrzmnBzUYWCzSvOoHE6lMQQJok=",
+      "fft-stockham-shared.wgsl.jinja": "ihNoaimYQY8F/oomzVbEh9g/g08yolKLYVQn79cZG1o=",
+      "fft-stockham-twiddles.wgsl.jinja": "6/+Ugc8XhOPl7iNTPc0Lwd0Moh8AFw+RVQZSWCxJbKY=",
+      "manifest.json": "stekrpAUb0OTNiTqln6FHfTBeHeZmISRKX/dP8eEkVU=",
+      "stft-general.wgsl.jinja": "UGbzUabSlzivMPxP2LJ2YRNK5Z6msagBugJvn3Tpu0s=",
+      "test.json": "Y05Ud3usGRQ2tsWy3twV3mndN+3FjHo3LjOwTA24Nog="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.STFT" }
+}

build/webgpu/stft-general.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,164 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{{ env.wgsl.resourceDeclarations }}
+// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
+// rational turn and route all polynomial evaluation through the shared precise
+// trig core so FFT, direct DFT, STFT, and cosine windows agree exactly.
+//
+// phase is in [0, period). sign is -1 for a forward transform and +1 for an
+// inverse transform. The result is (cos, sin).
+// Backend-stable f32 sine/cosine core.
+//
+// Shader transcendental accuracy is implementation-defined, and some portable
+// backends are only accurate to roughly 1e-4. Each path retains the most accurate
+// available phase representation, reduces it to [-pi, pi], then uses these
+// polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
+// truncation error is well below one f32 ULP over its documented interval.
+const PRECISE_TRIG_PI: f32 = 3.141592653589793;
+const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
+const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
+// Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
+fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
+  let x2 = x * x;
+  var sinPolynomial = 1.6059043836821613e-10;
+  sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
+  sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
+  sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
+  sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
+  sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
+  let sine = x * fma(sinPolynomial, x2, 1.0);
+  var cosPolynomial = 2.08767569878681e-9;
+  cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
+  cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
+  cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
+  cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
+  cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
+  let cosine = fma(cosPolynomial, x2, 1.0);
+  return vec2<f32>(cosine, sine);
+}
+// Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
+fn precise_sincos_centered(x: f32) -> vec2<f32> {
+  var folded = x;
+  var cosineSign = 1.0;
+  if (folded > PRECISE_TRIG_HALF_PI) {
+    folded = PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  } else if (folded < -PRECISE_TRIG_HALF_PI) {
+    folded = -PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  }
+  let value = precise_sincos_half_pi(folded);
+  return vec2<f32>(cosineSign * value.x, value.y);
+}
+fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
+  let reducedPhase = phase % period;
+  if (reducedPhase == 0u) {
+    return vec2<f32>(1.0, sign * 0.0);
+  }
+  if (period % 4u == 0u) {
+    let quarter = period / 4u;
+    if (reducedPhase == quarter) {
+      return vec2<f32>(0.0, sign);
+    }
+    if (reducedPhase == 2u * quarter) {
+      return vec2<f32>(-1.0, sign * 0.0);
+    }
+    if (reducedPhase == 3u * quarter) {
+      return vec2<f32>(0.0, -sign);
+    }
+  }
+  // Center the rational turn before multiplying by 2*pi. This avoids the
+  // cancellation in `2*pi*(N-1)/N - 2*pi` for bins near the end of a row.
+  let turns = f32(reducedPhase) / f32(period);
+  let centeredTurns = turns - floor(turns + 0.5);
+  return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
+}
+// General direct STFT for rank-3 real or complex signals. The frame length is
+// the window length when a window is bound, otherwise the adapted frame_length
+// scalar. Window samples default to 1.0 when no window input is present.
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("index", guardInline=true, note="") }}
+  let bin = index % params.bins;
+  var t = index / params.bins;
+  let frame = t % params.frames;
+  let batch = t / params.frames;
+  let start = frame * params.frameStep;
+  let frame_length = params.frameLength;
+  var real = 0.0;
+  var imag = 0.0;
+  let phasePeriod = max(frame_length, 1u);
+  let phaseStep = bin % phasePeriod;
+  var phase = 0u;
+  for (var n = 0u; n < frame_length; n = n + 1u) {
+    let signal_index = start + n;
+    let signal_base = (batch * params.signalLength + signal_index) * params.signalComponents;
+{% if source.hasWindowInput %}
+    let w = window[n];
+{% else %}
+    let w = 1.0;
+{% endif %}
+    let xr = signal[signal_base] * w;
+    var xi = 0.0;
+    if (params.signalComponents == 2u) {
+      xi = signal[signal_base + 1u] * w;
+    }
+    // Reduce the integer phase before conversion. The stable unit-circle
+    // approximation avoids backend-dependent residuals at exact DFT roots.
+    let twiddle = dft_precise_twiddle(phase, frame_length, -1.0);
+    let c = twiddle.x;
+    let s = twiddle.y;
+    real = real + xr * c - xi * s;
+    imag = imag + xr * s + xi * c;
+    // Overflow-safe modular increment. A direct `bin * n` wraps u32 for long
+    // transforms before the modulus is applied.
+    let distanceToWrap = phasePeriod - phaseStep;
+    phase = select(phase + phaseStep, phase - distanceToWrap, phase >= distanceToWrap);
+  }
+  let output_base = index * 2u;
+  output[output_base] = real;
+  output[output_base + 1u] = imag;
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,847 @@

+{
+  "op": "ai.onnx.STFT",
+  "fixtureArrays": {
+    "onnx_backend_arange128_cos_window16_input_window": [1, 0.961939766, 0.853553391, 0.691341716, 0.5, 0.308658284, 0.146446609, 0.038060234, 0, 0.038060234, 0.146446609, 0.308658284, 0.5, 0.691341716, 0.853553391, 0.961939766],
+    "ort_complex_input_batched_twosided_frame_length_input_signal": [1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0]
+  },
+  "cases": [
+    {
+      "name": "dispatch_cliff_rank3_over_8m_elements",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [1, 2200000, 1] },
+        "window": { "dtype": "float32", "shape": [6], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2199995, 4, 2], "tolerance": 0.001 } },
+      "args": { "frameStep": 1 }
+    },
+    {
+      "name": "many_frames_fft_twiddle_table_win64",
+      "provenance": {
+        "notes": "Exercises the shared twiddle-table FFT selected for thousands of independent transforms without making correctness coverage depend on the much larger dispatch-cliff benchmark allocation."
+      },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [1, 4159, 1] },
+        "window": { "dtype": "float32", "shape": [64] }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4096, 33, 2], "tolerance": 0.0001 } },
+      "args": { "frameStep": 1 }
+    },
+    {
+      "name": "many_frames_fft_twiddle_table_win64_twosided",
+      "provenance": {
+        "notes": "Twosided sibling of the twiddle-table selector above. The packed real-FFT kernel that otherwise wins at this transform count requires onesided != 0, so a full-spectrum output is what leaves the shared twiddle-table FFT as the selected route."
+      },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 4159, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.029, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [64],
+          "data": { "kind": "fillFloat32", "sinStep": 0.007, "cosStep": 0.013, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 4096, 64, 2], "tolerance": 0.001, "relTolerance": 0.0001 }
+      },
+      "args": { "frameStep": 1 }
+    },
+    {
+      "name": "many_frames_packed_rfft_win512_radix4",
+      "provenance": {
+        "notes": "Selector and correctness lock for the packed real N=512 path. The 4097 transforms cross the twiddle-table threshold and leave one partial four-transform workgroup on capable devices."
+      },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 4608, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.029, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.007, "cosStep": 0.013, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 4097, 257, 2], "tolerance": 0.01, "relTolerance": 0.0001 }
+      },
+      "args": { "frameStep": 1 }
+    },
+    {
+      "name": "real_window_onesided",
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 6, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "window": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, 1.0, 1.0, 0.5] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 3, 2], "tolerance": 0.00002 } },
+      "args": { "frameStep": 2 }
+    },
+    {
+      "name": "real_subnormal_impulse_window_onesided_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal FTZ: subnormal 1e-40 impulse flushes to zero on GPU, every one-sided bin reads 0 instead of the subnormal value (tolerance 0)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.STFTFloat",
+        "notes": "Focused amplitude edge: a one-frame rectangular STFT of a subnormal impulse should preserve that value in every real one-sided bin."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 4, 1],
+          "data": { "kind": "values", "values": [1e-40, 0.0, 0.0, 0.0] }
+        },
+        "window": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 3, 2], "tolerance": 0 } },
+      "args": { "frameStep": 1 }
+    },
+    {
+      "name": "complex_twosided",
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 4, 2],
+          "data": { "kind": "values", "values": [1.0, 0.0, 2.0, 1.0, 3.0, -1.0, 4.0, 0.0] }
+        },
+        "window": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 0.5] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 2, 2], "tolerance": 0.00002 } },
+      "args": { "frameStep": 2 }
+    },
+    {
+      "name": "real_three_frames_window",
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 5, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "window": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.25, 1.0, 0.25] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 2, 2], "tolerance": 0.00002 } },
+      "args": { "frameStep": 1 }
+    },
+    {
+      "name": "ort_scalar_frame_step_windowed",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.STFTFloat",
+        "notes": "Compact scalar-frame-step projection of ORT's windowed STFT test."
+      },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 6, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "window": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 3, 2], "tolerance": 0.00002 } },
+      "args": { "frameStep": 2 }
+    },
+    {
+      "name": "real_twosided_window",
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 6, 1],
+          "data": { "kind": "values", "values": [0.0, 1.0, 0.0, -1.0, 0.0, 1.0] }
+        },
+        "window": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 0.5, 0.5, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 4, 2], "tolerance": 0.00002 } },
+      "args": { "frameStep": 2 }
+    },
+    {
+      "name": "batch2_onesided_stride2_window3",
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [2, 5, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, -1.0, 0.0, 1.0, 0.0, -1.0] }
+        },
+        "window": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.5, 1.0, 0.25] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 2, 2], "tolerance": 0.00002 } },
+      "args": { "frameStep": 2 }
+    },
+    {
+      "name": "single_frame_full_length_onesided",
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 4, 1],
+          "data": { "kind": "values", "values": [1.0, 0.0, -1.0, 0.0] }
+        },
+        "window": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 3, 2], "tolerance": 0.00002 } },
+      "args": { "frameStep": 4 }
+    },
+    {
+      "name": "ort_constant_64_window16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.STFTFloat"
+      },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [1, 64, 1], "data": { "kind": "constant", "value": 1.0 } },
+        "window": { "dtype": "float32", "shape": [16], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 7, 9, 2],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": [16.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 16.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 16.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 16.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 16.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 16.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 16.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0]
+          }
+        }
+      },
+      "args": { "frameStep": 8 }
+    },
+    {
+      "name": "onnx_backend_arange128_cos_window16",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_stft_with_window" },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 128, 1],
+          "data": { "kind": "linspace", "start": 0.0, "end": 127.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_arange128_cos_window16_input_window" }
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 15, 9, 2], "tolerance": 0.00025 } },
+      "args": { "frameStep": 8 }
+    },
+    {
+      "name": "real_batch2_twosided_stride2",
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [2, 5, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, -1.0, 0.0, 1.0, 0.0, -1.0] }
+        },
+        "window": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 0.5, 0.25] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 3, 2], "tolerance": 0.00002 } },
+      "args": { "frameStep": 2 }
+    },
+    {
+      "name": "onnx_backend_stft_rectangular_window",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_stft",
+        "notes": "Official rectangular-window fixture: window is omitted and frame_length supplies the DFT size."
+      },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 128, 1],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0, 30.0, 31.0, 32.0, 33.0, 34.0, 35.0, 36.0, 37.0, 38.0, 39.0, 40.0, 41.0, 42.0, 43.0, 44.0, 45.0, 46.0, 47.0, 48.0, 49.0, 50.0, 51.0, 52.0, 53.0, 54.0, 55.0, 56.0, 57.0, 58.0, 59.0, 60.0, 61.0, 62.0, 63.0, 64.0, 65.0, 66.0, 67.0, 68.0, 69.0, 70.0, 71.0, 72.0, 73.0, 74.0, 75.0, 76.0, 77.0, 78.0, 79.0, 80.0, 81.0, 82.0, 83.0, 84.0, 85.0, 86.0, 87.0, 88.0, 89.0, 90.0, 91.0, 92.0, 93.0, 94.0, 95.0, 96.0, 97.0, 98.0, 99.0, 100.0, 101.0, 102.0, 103.0, 104.0, 105.0, 106.0, 107.0, 108.0, 109.0, 110.0, 111.0, 112.0, 113.0, 114.0, 115.0, 116.0, 117.0, 118.0, 119.0, 120.0, 121.0, 122.0, 123.0, 124.0, 125.0, 126.0, 127.0]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 15, 9, 2], "tolerance": 0.00025 } },
+      "args": { "frameStep": 8, "frameLength": 16 }
+    },
+    {
+      "name": "ort_constant_64_window16_twosided",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.STFTFloat",
+        "notes": "Two-sided projection of ORT's constant-window STFT setup; validates the full spectrum instead of the default one-sided bins."
+      },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [1, 64, 1], "data": { "kind": "constant", "value": 1.0 } },
+        "window": { "dtype": "float32", "shape": [16], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 7, 16, 2], "tolerance": 0.00025 } },
+      "args": { "frameStep": 8 }
+    },
+    {
+      "name": "ort_frame_length_without_window",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.STFTFloat",
+        "notes": "Covers the ONNX optional-input form where frame_length is provided and the window input is omitted."
+      },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 6, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 2, 3, 2],
+          "tolerance": 0.00002,
+          "data": { "kind": "values", "values": [10.0, 0.0, -2.0, 2.0, -2.0, 0.0, 18.0, 0.0, -2.0, 2.0, -2.0, 0.0] }
+        }
+      },
+      "args": { "frameStep": 2, "frameLength": 4 }
+    },
+    {
+      "name": "ort_complex_input_batched_twosided_frame_length",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.STFTFloatComplexInputBatched",
+        "notes": "Compact optional-window projection of ORT's batched complex-input STFT regression."
+      },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [2, 8, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_complex_input_batched_twosided_frame_length_input_signal" }
+          }
+        }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [2, 3, 4, 2],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": [4.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 4.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 4.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 396.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 396.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 396.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0]
+          }
+        }
+      },
+      "args": { "frameStep": 2, "frameLength": 4 }
+    },
+    {
+      "name": "fft_window256_signal1024_batch2_onesided",
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [2, 1024, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.029, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [256],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.011, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [2, 4, 129, 2], "tolerance": 0.001, "relTolerance": 0.0001 }
+      },
+      "args": { "frameStep": 256 }
+    },
+    {
+      "name": "fft_window512_complex_twosided_step256",
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 2048, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.023, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.009, "cosStep": 0.007, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 7, 512, 2], "tolerance": 0.01, "relTolerance": 0.0001 }
+      },
+      "args": { "frameStep": 256 }
+    },
+    {
+      "name": "ort_complex_batched_no_window_frame_length",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.STFTFloatComplexInputBatched",
+        "notes": "ORT's optional-window path uses frame_length as the rectangular-window length. This case locks the optimized no-window FFT route for batched complex input."
+      },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [2, 8, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_complex_input_batched_twosided_frame_length_input_signal" }
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 4, 2], "tolerance": 0.00005 } },
+      "args": { "frameStep": 2, "frameLength": 4 }
+    },
+    {
+      "name": "ort_complex_batched_frame_offsets_twosided",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.STFTFloatComplexInputBatched",
+        "notes": "Compact projection of ORT's batched complex-input frame-offset regression."
+      },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [2, 8, 2],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_complex_input_batched_twosided_frame_length_input_signal" }
+          }
+        },
+        "window": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 1.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 4, 2], "tolerance": 0.00005 } },
+      "args": { "frameStep": 2 }
+    },
+    {
+      "name": "fft_window512_real_twosided_step256",
+      "provenance": {
+        "notes": "Real/full-spectrum coverage for the power-of-two FFT path; complements complex/twosided N=512 and real/onesided N=1024 fixtures."
+      },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 1024, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.007, "cosStep": 0.013, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 3, 512, 2], "tolerance": 0.01, "relTolerance": 0.0001 }
+      },
+      "args": { "frameStep": 256 }
+    },
+    {
+      "name": "fft_window1024_onesided_step512",
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 4096, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.027, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "fillFloat32", "sinStep": 0.005, "cosStep": 0.003, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 7, 513, 2], "tolerance": 0.004, "relTolerance": 0.0001 }
+      },
+      "args": { "frameStep": 512 }
+    },
+    {
+      "name": "fft_window4096_radix2_dit_onesided",
+      "provenance": {
+        "notes": "Locks the capability-selected 4096-point radix-2 DIT routes: workgroup-resident on devices with 32 KiB of shared storage, and globally staged/storage-backed on lower-limit devices."
+      },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 4096, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.027, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.005, "cosStep": 0.003, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 1, 2049, 2], "tolerance": 0.04, "relTolerance": 0.0002 }
+      },
+      "args": { "frameStep": 4096 }
+    },
+    {
+      "name": "fft_window2048_radix2_dit_onesided_min_storage",
+      "provenance": {
+        "notes": "The in-place radix-2 rung is selected inside a BAND — rank3WindowInplaceFftFits wants window*16 > maxComputeWorkgroupStorageSize (the two-buffer form does not fit) and window*8 <= it (the in-place one does) — so the band moves with the device and no single window covers every tier. window 4096 is the band at 32768 (the adapter here); 2048 is the band at the 16384 WebGPU guaranteed minimum, which is what the min-limits tier reports. Without this sibling the rung is rendered on every low-storage tier and executed on none of them. Deliberately no requires.limits: outside its band the case is still valid, it just selects the staged rung."
+      },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 2048, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.027, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [2048],
+          "data": { "kind": "fillFloat32", "sinStep": 0.005, "cosStep": 0.003, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 1, 1025, 2], "tolerance": 0.04, "relTolerance": 0.0002 }
+      },
+      "args": { "frameStep": 2048 }
+    },
+    {
+      "name": "fft_window4096_radix2_dit_complex_twosided",
+      "provenance": {
+        "notes": "Exercises the complex load and full-spectrum store arms shared by the capability-selected 4096-point DIT routes."
+      },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 4096, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.027, "scale": 0.2 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [4096],
+          "data": { "kind": "fillFloat32", "sinStep": 0.005, "cosStep": 0.003, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 1, 4096, 2], "tolerance": 0.05, "relTolerance": 0.0003 }
+      },
+      "args": { "frameStep": 4096 }
+    },
+    {
+      "name": "empty_zero_dim",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [0, 6, 1], "data": { "kind": "values", "values": [] } },
+        "window": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, 1.0, 1.0, 0.5] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 2, 3, 2], "tolerance": 0 } },
+      "args": { "frameStep": 2 }
+    },
+    {
+      "name": "mixed_stockham_n400_real_onesided_tail25",
+      "provenance": {
+        "notes": "Real onesided 400-point mixed-radix path: 16-point shared Stockham transforms over 25 residue classes, followed by the direct radix-25 tail."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 880, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.021, "cosStep": 0.013, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [400],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0079, "cosStep": 0.0079, "scale": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 201, 2], "tolerance": 0.01, "relTolerance": 0.001 } },
+      "args": { "frameStep": 160 }
+    },
+    {
+      "name": "mixed_stockham_n288_real_onesided_tail9",
+      "provenance": {
+        "notes": "Generic mixed-radix recurrence with a 32-point shared Stockham factor and radix-9 tail."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 432, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.017, "scale": 0.7 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [288],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.007, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 145, 2], "tolerance": 0.02, "relTolerance": 0.001 } },
+      "args": { "frameStep": 144 }
+    },
+    {
+      "name": "mixed_stockham_n400_complex_twosided",
+      "provenance": { "notes": "Complex/full-spectrum branch of the 16x25 mixed-radix decomposition." },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 400, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.031, "scale": 0.3 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [400],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0079, "cosStep": 0.0079, "scale": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 400, 2], "tolerance": 0.02, "relTolerance": 0.001 } },
+      "args": { "frameStep": 400 }
+    },
+    {
+      "name": "mixed_stockham_n480_real_twosided_tail15",
+      "provenance": {
+        "notes": "Real full-spectrum 480-point mixed-radix path: 32-point shared Stockham transforms over 15 residue classes, followed by the direct radix-15 tail."
+      },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 960, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.029, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [480],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0065, "cosStep": 0.0065, "scale": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 480, 2], "tolerance": 0.02, "relTolerance": 0.001 } },
+      "args": { "frameStep": 240 }
+    },
+    {
+      "name": "mixed_stockham_n496_real_onesided_tail31",
+      "provenance": {
+        "notes": "Generic mixed-radix recurrence at the supported tail bound: a 16-point shared Stockham factor and radix-31 tail."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 744, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.031, "scale": 0.6 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [496],
+          "data": { "kind": "fillFloat32", "sinStep": 0.009, "cosStep": 0.005, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 249, 2], "tolerance": 0.02, "relTolerance": 0.001 } },
+      "args": { "frameStep": 248 }
+    },
+    {
+      "name": "mixed_stockham_n132_real_onesided_tail33",
+      "provenance": {
+        "notes": "First tail past the unanchored bound: a 4-point shared Stockham factor and radix-33 tail, so the tail walk re-anchors exactly once."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 198, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.031, "scale": 0.6 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [132],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.017, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 67, 2], "tolerance": 0.02, "relTolerance": 0.001 } },
+      "args": { "frameStep": 66 }
+    },
+    {
+      "name": "window_and_frame_length_copresent_fft",
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [1, 64, 1], "data": { "kind": "constant", "value": 1.0 } },
+        "window": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/onnx_backend_arange128_cos_window16_input_window" }
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 7, 9, 2], "tolerance": 0.0002 } },
+      "args": { "frameStep": 8, "frameLength": 16 }
+    },
+    {
+      "name": "mixed_stockham_n1000_real_twosided_tail125",
+      "provenance": {
+        "notes": "Window length 1000 factors as 8 x 125 and exercises the shared Stockham mixed-radix route after the tail-bound transition. Odd lengths 513 and 777 retain independent tiled-real coverage."
+      },
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 2000, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0091, "cosStep": 0.0071, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [1000],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0031, "cosStep": 0.0031, "scale": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 2, 1000, 2], "tolerance": 0.05, "relTolerance": 0.002 }
+      },
+      "args": { "frameStep": 1000 }
+    },
+    {
+      "name": "fft_pow2_win8_drops_partial_last_frame",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 12, 1],
+          "data": { "kind": "values", "values": [1.0, -2.0, 3.0, -4.0, 5.0, -6.0, 7.0, -8.0, 9.0, -10.0, 11.0, -12.0] }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, 0.9, 0.8, 0.7, 0.6, 0.5, 0.4, 0.3] }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 1, 5, 2], "tolerance": 0.0002, "relTolerance": 0.0001 }
+      },
+      "args": { "frameStep": 8 },
+      "provenance": {
+        "notes": "The final four samples cannot fill an eight-sample frame, so ONNX floor framing emits only the complete frame at offset zero."
+      }
+    },
+    {
+      "name": "fft_pow2_win4_complex_twosided_drops_partial_last_frame",
+      "attrs": { "onesided": 0 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 6, 2],
+          "data": { "kind": "values", "values": [1.0, 0.5, 2.0, -0.5, 3.0, 1.0, -4.0, 0.0, 5.0, -2.0, 6.0, 0.25] }
+        },
+        "window": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 0.75, 0.5, 0.25] } }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 1, 4, 2], "tolerance": 0.0002, "relTolerance": 0.0001 }
+      },
+      "args": { "frameStep": 4 },
+      "provenance": {
+        "notes": "The frame at offset four would extend past the six-sample signal, so ONNX floor framing emits only the complete frame at offset zero."
+      }
+    },
+    {
+      "name": "fft_pow2_win16_window_with_interior_zeros",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": { "dtype": "float32", "shape": [1, 32, 1], "data": { "kind": "linspace", "start": 0.0, "end": 31.0 } },
+        "window": {
+          "dtype": "float32",
+          "shape": [16],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.25, 0.5, 0.75, 1.0, 0.75, 0.5, 0.0, 0.0, 0.5, 0.75, 1.0, 0.75, 0.5, 0.25, 0.0]
+          }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 3, 9, 2], "tolerance": 0.0002, "relTolerance": 0.0001 }
+      },
+      "args": { "frameStep": 8 }
+    },
+    {
+      "name": "dispatch_fold_fft_batch2_frames32800_win4_onesided",
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [2, 32803, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.029, "scale": 1.0 }
+        },
+        "window": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.5, 1.0, 1.0, 0.5] } }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [2, 32800, 3, 2], "tolerance": 0.001, "relTolerance": 0.0001 }
+      },
+      "args": { "frameStep": 1 }
+    },
+    {
+      "name": "rank3_window_win513_nonpow2_onesided_compact",
+      "provenance": {
+        "notes": "Compact selector and correctness lock for the first non-power-of-two window above 512. Three frames exercise frame addressing; 257 bins exercise the partial second 256-frequency block."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 1025, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0091, "cosStep": 0.0071, "scale": 1.0 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [513],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0061, "cosStep": 0.0061, "scale": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 257, 2], "tolerance": 0.02, "relTolerance": 0.001 } },
+      "args": { "frameStep": 256 }
+    },
+    {
+      "name": "rank3_window_win777_tiled_real_generic",
+      "provenance": { "notes": "Non-513 correctness lock for the storage-bounded tiled real STFT path." },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 777, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0067, "cosStep": 0.0091, "scale": 0.5 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [777],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0041, "cosStep": 0.0051, "scale": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 389, 2], "tolerance": 0.04, "relTolerance": 0.002 } },
+      "args": { "frameStep": 777 }
+    },
+    {
+      "name": "rank3_window_fft8192_storage_generic",
+      "provenance": {
+        "notes": "Non-4096 correctness lock for the geometry-derived large radix-2 FFT; device storage capacity selects the single-dispatch or globally staged implementation."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "signal": {
+          "dtype": "float32",
+          "shape": [1, 8192, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0045, "cosStep": 0.0085, "scale": 0.25 }
+        },
+        "window": {
+          "dtype": "float32",
+          "shape": [8192],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0025, "cosStep": 0.0015, "scale": 0.5, "offset": 0.5 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [1, 1, 4097, 2], "tolerance": 0.08, "relTolerance": 0.0005 }
+      },
+      "args": { "frameStep": 8192 }
+    }
+  ]
+}