sync 2e7068faf55e

Browse files

Files changed (11) hide show

README.md +82 -0
build/webgpu/bench.json +147 -0
build/webgpu/dft-contiguous-naive.wgsl.jinja +170 -0
build/webgpu/dft-general.wgsl.jinja +221 -0
build/webgpu/dft-rank4.wgsl.jinja +205 -0
build/webgpu/dft-tiled-real.wgsl.jinja +136 -0
build/webgpu/fft-radix2-dit-storage.wgsl.jinja +233 -0
build/webgpu/fft-stockham-shared.wgsl.jinja +193 -0
build/webgpu/manifest.json +816 -0
build/webgpu/metadata.json +23 -0
build/webgpu/test.json +1374 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.DFT
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 20
+## Description
+Computes the discrete Fourier Transform (DFT) of the input along a specified axis. For a signal of length N, output bin k is `sum_{n=0}^{N-1} exp(-2*pi*j*k*n/N) * x[n]`; the inverse divides by N and negates the exponent sign. Supports forward/inverse, real-to-complex (RFFT), and complex-to-real (IRFFT) modes via the `onesided` and `inverse` attributes.
+See the [ONNX `DFT` spec](https://onnx.ai/onnx/operators/onnx__DFT.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `input` | `input` | `T` | same as logical dtype | — | — | Input signal tensor; the last dimension is 1 for real values or 2 for complex `(real, imaginary)` pairs. | required |
+| `dft_length` | `dft_length` | `L` | same as logical dtype | `0` | — | Optional int32 scalar controlling the signal length used for the transform; input is zero-padded or truncated to this length. | optional |
+| `axis` | `axis` | `I` | `int32` | `0` | — | Optional logical int64 scalar specifying the dimension over which to compute the DFT (the last axis is reserved for the real/imaginary component and is excluded); signed axes use int32 WebGPU storage and default to `-2` when omitted. | optional |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | same as `input` | — | DFT result tensor; last dimension is 2 (complex) for forward DFT and RFFT, or 1 (real) for IRFFT. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `inverse` | `0` | When set to 1, computes the inverse DFT (IDFT/IRFFT) instead of the forward transform; default is 0 (forward). |
+| `onesided` | `0` | When set to 1, exploits conjugate symmetry to return only the non-redundant half of the spectrum (RFFT for forward, IRFFT for inverse); default is 0 (full spectrum). |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32` |
+| `L` | `int32` |
+| `I` | `int64` |
+## 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-contiguous-naive.wgsl.jinja`](build/webgpu/dft-contiguous-naive.wgsl.jinja)
+- [`dft-general.wgsl.jinja`](build/webgpu/dft-general.wgsl.jinja)
+- [`dft-rank4.wgsl.jinja`](build/webgpu/dft-rank4.wgsl.jinja)
+- [`dft-tiled-real.wgsl.jinja`](build/webgpu/dft-tiled-real.wgsl.jinja)
+- [`fft-radix2-dit-storage.wgsl.jinja`](build/webgpu/fft-radix2-dit-storage.wgsl.jinja)
+- [`fft-stockham-shared.wgsl.jinja`](build/webgpu/fft-stockham-shared.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader automatically allocates outputs whose metadata it can derive from the manifest contract and this call.
+The explicit `outputs` entries provide shape and logical dtype metadata for the results listed below:
+- `output`
+Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
+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.DFT", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { output } = await kernel({ input: { data: inputData, shape: [1, 1, 1, 1, 1, 1, 2, 1] } }, {
+  outputs: { output: { shape: [1, 1, 1, 1, 1, 1, 2, 2], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,147 @@

+{
+  "op": "ai.onnx.DFT",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "batch4_len64_complex",
+      "inputs": { "input": { "dtype": "float32", "shape": [4, 64, 2] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 64, 2] } }
+    },
+    {
+      "name": "fft_batch_axis1_len1024_batch64_real_full",
+      "preset": "smoke",
+      "vars": { "transforms": 64, "n": 1024, "log2n": 10 },
+      "inputs": { "input": { "dtype": "float32", "shape": [64, 1024, 1], "dist": "normal", "seed": 901 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [64, 1024, 2] } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "gflops", "value": "5 * args.transforms * args.n * args.log2n" }]
+      }
+    },
+    {
+      "name": "fft_batch_axis1_len2048_batch32_real_onesided",
+      "preset": "smoke",
+      "attrs": { "onesided": 1 },
+      "vars": { "transforms": 32, "n": 2048, "log2n": 11 },
+      "inputs": { "input": { "dtype": "float32", "shape": [32, 2048, 1], "dist": "normal", "seed": 902 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [32, 1025, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.transforms * args.n * args.log2n" }] }
+    },
+    {
+      "name": "fft_rank4_axis1_len1024_batch2_inner8",
+      "preset": "smoke",
+      "attrs": {},
+      "vars": { "transforms": 16, "n": 1024, "log2n": 10 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": { "dtype": "float32", "shape": [2, 1024, 8, 1], "dist": "normal", "seed": 903 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 1024, 8, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.transforms * args.n * args.log2n" }] }
+    },
+    {
+      "name": "fft_rank2_axis0_len512_single",
+      "preset": "smoke",
+      "attrs": {},
+      "vars": { "transforms": 1, "n": 512, "log2n": 9 },
+      "inputs": { "input": { "dtype": "float32", "shape": [512, 1], "dist": "normal", "seed": 904 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [512, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.transforms * args.n * args.log2n" }] }
+    },
+    {
+      "name": "tiled_batch_axis1_len1000_batch64_real_full",
+      "preset": "stress",
+      "vars": { "transforms": 64, "n": 1000 },
+      "inputs": { "input": { "dtype": "float32", "shape": [64, 1000, 1], "dist": "normal", "seed": 905 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [64, 1000, 2] } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "8 * args.transforms * args.n * args.n" }] }
+    },
+    {
+      "name": "fft_storage_inplace_batch_axis1_len4096_batch8_real_full",
+      "preset": "stress",
+      "vars": { "transforms": 8, "n": 4096, "log2n": 12 },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": { "input": { "dtype": "float32", "shape": [8, 4096, 1], "dist": "normal", "seed": 921, "scale": 2 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [8, 4096, 2], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "gflops", "value": "5 * args.transforms * args.n * args.log2n" }]
+      }
+    },
+    {
+      "name": "tiled_batch_axis1_len4095_batch8_fft_boundary_minus1",
+      "preset": "stress",
+      "vars": { "transforms": 8, "n": 4095 },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": { "input": { "dtype": "float32", "shape": [8, 4095, 1], "dist": "normal", "seed": 925, "scale": 2 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [8, 4095, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "8 * args.transforms * args.n * args.n" }] }
+    },
+    {
+      "name": "fft_storage_scratch_batch_axis1_len4096_batch8_real_onesided",
+      "preset": "stress",
+      "vars": { "transforms": 8, "n": 4096, "log2n": 12 },
+      "attrs": { "inverse": 0, "onesided": 1 },
+      "inputs": { "input": { "dtype": "float32", "shape": [8, 4096, 1], "dist": "normal", "seed": 924, "scale": 2 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [8, 2049, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.transforms * args.n * args.log2n" }] }
+    },
+    {
+      "name": "general_len_pow2_1024_batch32_dftlength_demotes_fft",
+      "preset": "stress",
+      "vars": { "transforms": 32, "n": 1024, "log2n": 10 },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [32, 1024, 1], "dist": "normal", "seed": 922, "scale": 2 },
+        "dft_length": { "dtype": "int32", "shape": [], "dist": "constant", "value": 1024 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [32, 1024, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.transforms * args.n * args.log2n" }] }
+    },
+    {
+      "name": "audio-rfft-explicit-length2048-batch32-pathology",
+      "preset": "model",
+      "attrs": { "inverse": 0, "onesided": 1 },
+      "vars": { "transforms": 32, "n": 2048, "log2n": 11 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [32, 2048, 1], "dist": "normal", "seed": 923, "scale": 2 },
+        "dft_length": { "dtype": "int32", "shape": [], "dist": "constant", "value": 2048 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [32, 1025, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.transforms * args.n * args.log2n" }] }
+    },
+    {
+      "name": "audio-fft-explicit-axis2048-batch32-pathology",
+      "preset": "stress",
+      "provenance": {
+        "notes": "Explicit-axis FFT paired with an equal-element-count implicit-axis control to expose the optional-axis dispatch boundary."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "vars": { "transforms": 32, "n": 2048, "log2n": 11 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [32, 2048, 1], "dist": "normal", "seed": 926, "scale": 2 },
+        "axis": { "dtype": "int32", "shape": [], "dist": "constant", "value": 1 }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [32, 2048, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.transforms * args.n * args.log2n" }] }
+    },
+    {
+      "name": "fft_storage_generic_batch_axis1_len8192_batch4_real_onesided",
+      "preset": "stress",
+      "vars": { "transforms": 4, "n": 8192, "log2n": 13 },
+      "attrs": { "inverse": 0, "onesided": 1 },
+      "inputs": { "input": { "dtype": "float32", "shape": [4, 8192, 1], "dist": "normal", "seed": 927, "scale": 2 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4097, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "5 * args.transforms * args.n * args.log2n" }] }
+    },
+    {
+      "name": "tiled_real_nonpow2_len4097_batch2_onesided",
+      "preset": "stress",
+      "vars": { "transforms": 2, "n": 4097, "bins": 2049 },
+      "attrs": { "inverse": 0, "onesided": 1 },
+      "inputs": { "input": { "dtype": "float32", "shape": [2, 4097, 1], "dist": "normal", "seed": 928, "scale": 2 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2049, 2], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "gflops", "value": "8 * args.transforms * args.bins * args.n" }] }
+    }
+  ]
+}

build/webgpu/dft-contiguous-naive.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,170 @@

+{{ 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);
+}
+fn is_nyquist(freq: u32, signal_length: u32) -> bool {
+  return signal_length % 2u == 0u && freq * 2u == signal_length;
+}
+{% macro dft_naive_accumulate(transform_index, input_axis_length, output_axis_length, batch_index, other_index, include_batch=true, include_other=true) %}
+  var real = 0.0;
+  var imag = 0.0;
+  let inverse = params.inverse != 0u;
+  let onesided_inverse = inverse && params.onesided != 0u;
+  let signal_length = select({{ input_axis_length }}, {{ output_axis_length }}, onesided_inverse);
+  if (onesided_inverse) {
+    for (var freq = 0u; freq < {{ input_axis_length }}; freq = freq + 1u) {
+      let input_base = transform_input_base(freq{% if include_batch %}, {{ batch_index }}{% endif %}{% if include_other %}, {{ other_index }}{% endif %});
+      let xr = input[input_base];
+      var xi = 0.0;
+      if (params.inputComponents == 2u) {
+        xi = input[input_base + 1u];
+      }
+      let phase = ({{ transform_index }} * freq) % signal_length;
+      let twiddle = dft_precise_twiddle(phase, signal_length, 1.0);
+      let c = twiddle.x;
+      let s = twiddle.y;
+      var weight = 1.0;
+      if (freq != 0u && !is_nyquist(freq, signal_length)) {
+        weight = 2.0;
+      }
+      real = real + weight * (xr * c - xi * s);
+      imag = imag + weight * (xr * s + xi * c);
+    }
+    real = real / f32(signal_length);
+    imag = imag / f32(signal_length);
+  } else {
+    let sign = select(-1.0, 1.0, inverse);
+    for (var n = 0u; n < {{ input_axis_length }}; n = n + 1u) {
+      let input_base = transform_input_base(n{% if include_batch %}, {{ batch_index }}{% endif %}{% if include_other %}, {{ other_index }}{% endif %});
+      let xr = input[input_base];
+      var xi = 0.0;
+      if (params.inputComponents == 2u) {
+        xi = input[input_base + 1u];
+      }
+      let phase = ({{ transform_index }} * n) % signal_length;
+      let twiddle = dft_precise_twiddle(phase, signal_length, sign);
+      let c = twiddle.x;
+      let s = twiddle.y;
+      real = real + xr * c - xi * s;
+      imag = imag + xr * s + xi * c;
+    }
+    if (inverse) {
+      real = real / f32(signal_length);
+      imag = imag / f32(signal_length);
+    }
+  }
+{% endmacro %}
+fn transform_input_base(sample: u32{% if source.batched %}, batch: u32{% endif %}) -> u32 {
+{% if source.batched %}
+  return (batch * params.signalLength + sample) * params.inputComponents;
+{% else %}
+  return sample * params.inputComponents;
+{% endif %}
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // Fold a long 1D output grid into y at the device dispatch limit.
+  let index = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (index >= params.count) {
+    return;
+  }
+{% if source.batched %}
+  let transform_index = index % params.outputLength;
+  let batch = index / params.outputLength;
+{{ dft_naive_accumulate("transform_index", "params.signalLength", "params.outputLength", "batch", "0u", true, false) }}
+{% else %}
+  let transform_index = index;
+{{ dft_naive_accumulate("transform_index", "params.signalLength", "params.outputLength", "0u", "0u", false, false) }}
+{% endif %}
+  let output_base = index * params.outputComponents;
+  output[output_base] = real;
+  if (params.outputComponents == 2u) {
+    output[output_base + 1u] = imag;
+  }
+}

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

	@@ -0,0 +1,221 @@

+{{ 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 naive DFT with an optional runtime transform axis and an optional
+// runtime dft_length (opset-20 `axis` / `dft_length` tensor inputs). The output
+// shape already encodes the per-axis transform length, so we read every output
+// element's logical coords from the output dims, locate the transform axis, and
+// sum over the (zero-padded / truncated) signal of `signal_length` samples.
+//
+// Layout: the last tensor dimension is the complex-component axis (1 = real,
+// 2 = complex). The dims passed in params are the logical dims (tensor dims
+// minus the component axis) packed into 7 slots; leading unused slots hold 1.
+// `pad = 7 - n_logical` maps a logical axis to its padded slot, so the row-major
+// stride math is uniform.
+//
+// Packing depends on whether the axis input is present. The default -2 axis
+// uses a right-aligned (outer, axis, inner) shape. A runtime axis instead keeps
+// every logical dimension right-aligned in the same seven slots (rank <= 8).
+fn is_nyquist(freq: u32, signal_length: u32) -> bool {
+  return signal_length % 2u == 0u && freq * 2u == signal_length;
+}
+fn resolve_axis() -> u32 {
+{% if source.hasAxisInput %}
+  let raw = axis_in[0];
+  let rank = i32(params.rank);
+  let norm = select(raw, raw + rank, raw < 0);
+  return u32(norm);
+{% else %}
+  return params.axisAttr;
+{% endif %}
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded flat index over output elements (one per complex sample).
+  let index = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (index >= params.count) { return; }
+  let in_dims = array<u32, 7>(params.inD0, params.inD1, params.inD2, params.inD3, params.inD4, params.inD5, params.inD6);
+  let out_dims = array<u32, 7>(params.outD0, params.outD1, params.outD2, params.outD3, params.outD4, params.outD5, params.outD6);
+  // pad = number of leading 1-slots; logical dims live in [pad .. 6].
+  let n_logical = params.rank - 1u;
+  let pad = 7u - n_logical;
+  let resolved_axis = resolve_axis();
+{% if source.skipRuntimeAxis1 %}
+  // Complementary to the guarded shared-memory FFT pass in the dynamic-axis
+  // specialization. Axis 1 is already complete, so this dispatch is a uniform
+  // no-op for that case.
+  if (resolved_axis == 1u) { return; }
+{% endif %}
+  let transform_slot = resolved_axis + pad;
+  // Decode the right-aligned logical output coords from the flat index.
+  var coords = array<u32, 7>(0u, 0u, 0u, 0u, 0u, 0u, 0u);
+  var rem = index;
+  for (var slot = 7u; slot > pad; slot = slot - 1u) {
+    let dim = out_dims[slot - 1u];
+    coords[slot - 1u] = rem % dim;
+    rem = rem / dim;
+  }
+  let transform_index = coords[transform_slot];
+  let input_axis_length = in_dims[transform_slot];
+  let output_axis_length = out_dims[transform_slot];
+  let inverse = params.inverse != 0u;
+  let onesided_inverse = inverse && params.onesided != 0u;
+  var signal_length = input_axis_length;
+{% if source.hasDftLengthInput %}
+  if (dft_len_in[0] > 0) { signal_length = u32(dft_len_in[0]); }
+{% else %}
+  if (params.dftLengthAttr > 0u) { signal_length = params.dftLengthAttr; }
+{% endif %}
+  if (onesided_inverse) { signal_length = output_axis_length; }
+  var real = 0.0;
+  var imag = 0.0;
+  if (onesided_inverse) {
+    for (var freq = 0u; freq < input_axis_length; freq = freq + 1u) {
+      var in_coords = coords;
+      in_coords[transform_slot] = freq;
+      var in_linear = 0u;
+      for (var slot = 0u; slot < 7u; slot = slot + 1u) {
+        in_linear = in_linear * in_dims[slot] + in_coords[slot];
+      }
+      let in_base = in_linear * params.inputComponents;
+      let xr = input[in_base];
+      var xi = 0.0;
+      if (params.inputComponents == 2u) { xi = input[in_base + 1u]; }
+      // Range-reduce the phase: (k*n mod N) keeps |angle| < 2*PI so cos/sin stay
+      // accurate. f32(k)*f32(n) for k,n up to ~2047 reaches ~12867 rad where f32
+      // ULP (~1e-3 rad) corrupts the twiddle; k*n is exact in u32 here.
+      let kn_mod = (transform_index * freq) % signal_length;
+      let twiddle = dft_precise_twiddle(kn_mod, signal_length, 1.0);
+      let c = twiddle.x;
+      let s = twiddle.y;
+      var weight = 1.0;
+      if (freq != 0u && !is_nyquist(freq, signal_length)) { weight = 2.0; }
+      real = real + weight * (xr * c - xi * s);
+      imag = imag + weight * (xr * s + xi * c);
+    }
+    real = real / f32(signal_length);
+    imag = imag / f32(signal_length);
+  } else {
+    let sign = select(-1.0, 1.0, inverse);
+    for (var n = 0u; n < signal_length; n = n + 1u) {
+      if (n >= input_axis_length) { continue; }
+      var in_coords = coords;
+      in_coords[transform_slot] = n;
+      var in_linear = 0u;
+      for (var slot = 0u; slot < 7u; slot = slot + 1u) {
+        in_linear = in_linear * in_dims[slot] + in_coords[slot];
+      }
+      let in_base = in_linear * params.inputComponents;
+      let xr = input[in_base];
+      var xi = 0.0;
+      if (params.inputComponents == 2u) { xi = input[in_base + 1u]; }
+      // Range-reduce the phase (see onesided_inverse branch): (k*n mod N) keeps
+      // |angle| < 2*PI; k*n is exact in u32 for the naive path's index range.
+      let kn_mod = (transform_index * n) % signal_length;
+      let twiddle = dft_precise_twiddle(kn_mod, signal_length, sign);
+      let c = twiddle.x;
+      let s = twiddle.y;
+      real = real + xr * c - xi * s;
+      imag = imag + xr * s + xi * c;
+    }
+    if (inverse) {
+      real = real / f32(signal_length);
+      imag = imag / f32(signal_length);
+    }
+  }
+  let output_base = index * params.outputComponents;
+  output[output_base] = real;
+  if (params.outputComponents == 2u) {
+    output[output_base + 1u] = imag;
+  }
+}

build/webgpu/dft-rank4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,205 @@

+{% 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);
+}
+fn is_nyquist(freq: u32, signal_length: u32) -> bool {
+  return signal_length % 2u == 0u && freq * 2u == signal_length;
+}
+{% macro dft_naive_accumulate(transform_index, input_axis_length, output_axis_length, batch_index, other_index, include_batch=true, include_other=true) %}
+  var real = 0.0;
+  var imag = 0.0;
+  let inverse = params.inverse != 0u;
+  let onesided_inverse = inverse && params.onesided != 0u;
+  let signal_length = select({{ input_axis_length }}, {{ output_axis_length }}, onesided_inverse);
+  if (onesided_inverse) {
+    for (var freq = 0u; freq < {{ input_axis_length }}; freq = freq + 1u) {
+      let input_base = transform_input_base(freq{% if include_batch %}, {{ batch_index }}{% endif %}{% if include_other %}, {{ other_index }}{% endif %});
+      let xr = input[input_base];
+      var xi = 0.0;
+      if (params.inputComponents == 2u) {
+        xi = input[input_base + 1u];
+      }
+      let phase = ({{ transform_index }} * freq) % signal_length;
+      let twiddle = dft_precise_twiddle(phase, signal_length, 1.0);
+      let c = twiddle.x;
+      let s = twiddle.y;
+      var weight = 1.0;
+      if (freq != 0u && !is_nyquist(freq, signal_length)) {
+        weight = 2.0;
+      }
+      real = real + weight * (xr * c - xi * s);
+      imag = imag + weight * (xr * s + xi * c);
+    }
+    real = real / f32(signal_length);
+    imag = imag / f32(signal_length);
+  } else {
+    let sign = select(-1.0, 1.0, inverse);
+    for (var n = 0u; n < {{ input_axis_length }}; n = n + 1u) {
+      let input_base = transform_input_base(n{% if include_batch %}, {{ batch_index }}{% endif %}{% if include_other %}, {{ other_index }}{% endif %});
+      let xr = input[input_base];
+      var xi = 0.0;
+      if (params.inputComponents == 2u) {
+        xi = input[input_base + 1u];
+      }
+      let phase = ({{ transform_index }} * n) % signal_length;
+      let twiddle = dft_precise_twiddle(phase, signal_length, sign);
+      let c = twiddle.x;
+      let s = twiddle.y;
+      real = real + xr * c - xi * s;
+      imag = imag + xr * s + xi * c;
+    }
+    if (inverse) {
+      real = real / f32(signal_length);
+      imag = imag / f32(signal_length);
+    }
+  }
+{% endmacro %}
+fn input_index(batch: u32, d1: u32, d2: u32) -> u32 {
+  return ((batch * params.inputD1 + d1) * params.inputD2 + d2) * params.inputComponents;
+}
+fn output_index(batch: u32, d1: u32, d2: u32) -> u32 {
+  return ((batch * params.outputD1 + d1) * params.outputD2 + d2) * params.outputComponents;
+}
+fn transform_input_base(sample: u32, batch: u32, other: u32) -> u32 {
+  return input_index(batch, other, sample);
+}
+@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) }}
+  let out_d2 = index % params.outputD2;
+  var t = index / params.outputD2;
+  let out_d1 = t % params.outputD1;
+  let batch = t / params.outputD1;
+  let transform_index = out_d2;
+  let other_index = out_d1;
+  let input_axis_length = params.inputD2;
+  let output_axis_length = params.outputD2;
+{{ dft_naive_accumulate("transform_index", "input_axis_length", "output_axis_length", "batch", "other_index") }}
+  let out_base = output_index(batch, out_d1, out_d2);
+  output[out_base] = real;
+  if (params.outputComponents == 2u) {
+    output[out_base + 1u] = imag;
+  }
+}

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

	@@ -0,0 +1,136 @@

+// 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 batch = wg.y;
+  for (var n = lane; n < N; n = n + WG) {
+    signalTile[n] = input[(batch * N + n) * params.inner];
+  }
+  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 * NUM_BINS + freq) * 2u;
+  output[outputBase] = real;
+  output[outputBase + 1u] = imag;
+}

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

	@@ -0,0 +1,233 @@

+// 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" %}
+{% set USE_SCRATCH = source.mode != "dft" %}
+{% if not STFT %}
+// 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 %}
+{% if STFT %}
+const PI: f32 = 3.141592653589793;
+{% endif %}
+const N: u32 = {{ source.fftLength }}u;
+const N_HALF: u32 = N / 2u;
+const LOG2_N: u32 = {{ source.log2Length }}u;
+{% if USE_SCRATCH %}
+const NUM_BINS: u32 = {{ source.numBins }}u;
+{% endif %}
+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;
+{% if STFT %}
+  let batch = transform / params.numFrames;
+  let frame = transform - batch * params.numFrames;
+  let start = frame * params.frameStep;
+{% else %}
+  let outerIndex = transform / params.inner;
+  let innerIndex = transform - outerIndex * params.inner;
+{% endif %}
+  for (var n = tid; n < N; n = n + WG) {
+    var value = vec2<f32>(0.0);
+{% if STFT %}
+    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 %}
+{% else %}
+    let inputBase = ((outerIndex * N + n) * params.inner + innerIndex) * {{ source.inComp }}u;
+    value.x = input[inputBase];
+{% if source.inComp == 2 %}
+    value.y = input[inputBase + 1u];
+{% endif %}
+{% endif %}
+    let reversed = reverseBits(n) >> (32u - LOG2_N);
+    let dst = transformBase + reversed * 2u;
+{% if USE_SCRATCH %}
+    scratch[dst] = value.x;
+    scratch[dst + 1u] = value.y;
+{% else %}
+    output[dst] = value.x;
+    output[dst + 1u] = value.y;
+{% endif %}
+  }
+  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;
+{% if USE_SCRATCH %}
+      let a = vec2<f32>(scratch[lo], scratch[lo + 1u]);
+      let b = vec2<f32>(scratch[hi], scratch[hi + 1u]);
+{% else %}
+      let a = vec2<f32>(output[lo], output[lo + 1u]);
+      let b = vec2<f32>(output[hi], output[hi + 1u]);
+{% endif %}
+{% if STFT %}
+      let angle = SIGN * 2.0 * PI * f32(j) / f32(SPAN);
+      let c = cos(angle);
+      let sn = sin(angle);
+{% else %}
+      let twiddle = dft_precise_twiddle(j, SPAN, SIGN);
+      let c = twiddle.x;
+      let sn = twiddle.y;
+{% endif %}
+      let t = vec2<f32>(b.x * c - b.y * sn, b.x * sn + b.y * c);
+      let sum = a + t;
+      let difference = a - t;
+{% if USE_SCRATCH %}
+      scratch[lo] = sum.x;
+      scratch[lo + 1u] = sum.y;
+      scratch[hi] = difference.x;
+      scratch[hi + 1u] = difference.y;
+{% else %}
+      output[lo] = sum.x;
+      output[lo + 1u] = sum.y;
+      output[hi] = difference.x;
+      output[hi + 1u] = difference.y;
+{% endif %}
+    }
+    storageBarrier();
+  }
+{% endfor %}
+{% if source.inverse %}
+  const SCALE: f32 = 1.0 / f32(N);
+{% elif USE_SCRATCH %}
+  const SCALE: f32 = 1.0;
+{% endif %}
+{% if USE_SCRATCH %}
+  for (var k = tid; k < NUM_BINS; k = k + WG) {
+    let src = transformBase + k * 2u;
+{% if STFT %}
+    let dst = (transform * NUM_BINS + k) * 2u;
+{% else %}
+    let dst = ((outerIndex * NUM_BINS + k) * params.inner + innerIndex) * 2u;
+{% endif %}
+    output[dst] = scratch[src] * SCALE;
+    output[dst + 1u] = scratch[src + 1u] * SCALE;
+  }
+{% elif source.inverse %}
+  for (var k = tid; k < N; k = k + WG) {
+    let index = transformBase + k * 2u;
+    output[index] = output[index] * SCALE;
+    output[index + 1u] = output[index + 1u] * SCALE;
+  }
+{% endif %}
+}

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

	@@ -0,0 +1,193 @@

+{{ 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);
+}
+// 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;
+const NUM_BINS: u32 = {{ source.numBins }}u;
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+// 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" }};
+// 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;
+{% if source.runtimeAxis is defined and source.runtimeAxis %}
+  // Dynamic-axis DFT pairs this axis-specialized pass with a general pass.
+  // The uniform guard ensures exactly one performs work for each invocation.
+  let rawAxis = axis_in[0];
+  let normalizedAxis = select(rawAxis, rawAxis + {{ source.axisRank }}, rawAxis < 0);
+  if (normalizedAxis != {{ source.expectedAxis }}) { return; }
+{% endif %}
+  // 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 outerIndex = flat / params.inner;
+  let innerIndex = flat - outerIndex * params.inner;
+  for (var n = tid; n < N; n = n + WG) {
+    let base = ((outerIndex * N + n) * params.inner + innerIndex) * {{ source.inComp }}u;
+    shRe[n] = input[base];
+{% if source.inComp == 2 %}
+    shIm[n] = input[base + 1u];
+{% else %}
+    shIm[n] = 0.0;
+{% endif %}
+  }
+  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;
+  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];
+      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;
+    }
+    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.
+{% if source.inverse %}
+  const SCALE: f32 = 1.0 / f32(N);
+{% else %}
+  const SCALE: f32 = 1.0;
+{% endif %}
+  for (var k = tid; k < NUM_BINS; k = k + WG) {
+    let outBase = ((outerIndex * NUM_BINS + k) * params.inner + innerIndex) * 2u;
+    output[outBase] = shRe[srcOff + k] * SCALE;
+    output[outBase + 1u] = shIm[srcOff + k] * SCALE;
+  }
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,816 @@

+{
+  "domain": "ai.onnx",
+  "name": "DFT",
+  "sinceVersion": 20,
+  "description": "Computes the discrete Fourier Transform (DFT) of the input along a specified axis. For a signal of length N, output bin k is `sum_{n=0}^{N-1} exp(-2*pi*j*k*n/N) * x[n]`; the inverse divides by N and negates the exponent sign. Supports forward/inverse, real-to-complex (RFFT), and complex-to-real (IRFFT) modes via the `onesided` and `inverse` attributes.",
+  "inputs": [
+    {
+      "role": "input",
+      "dtype": "T",
+      "description": "Input signal tensor; the last dimension is 1 for real values or 2 for complex `(real, imaginary)` pairs."
+    },
+    {
+      "role": "dft_length",
+      "dtype": "L",
+      "optional": true,
+      "description": "Optional int32 scalar controlling the signal length used for the transform; input is zero-padded or truncated to this length.",
+      "rank": 0
+    },
+    {
+      "role": "axis",
+      "dtype": "I",
+      "optional": true,
+      "description": "Optional logical int64 scalar specifying the dimension over which to compute the DFT (the last axis is reserved for the real/imaginary component and is excluded); signed axes use int32 WebGPU storage and default to `-2` when omitted.",
+      "rank": 0
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "description": "DFT result tensor; last dimension is 2 (complex) for forward DFT and RFFT, or 1 (real) for IRFFT.",
+      "rank": "ranks.input"
+    }
+  ],
+  "attributes": { "inverse": 0, "onesided": 0 },
+  "attributeDescriptions": {
+    "inverse": "When set to 1, computes the inverse DFT (IDFT/IRFFT) instead of the forward transform; default is 0 (forward).",
+    "onesided": "When set to 1, exploits conjugate symmetry to return only the non-redundant half of the spectrum (RFFT for forward, IRFFT for inverse); default is 0 (full spectrum)."
+  },
+  "attributeConstraints": { "inverse": { "values": [0, 1] }, "onesided": { "values": [0, 1] } },
+  "typeConstraints": { "T": ["float32"], "L": ["int32"], "I": ["int64"] },
+  "args": {
+    "input": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "dft_length": { "kind": "tensor", "semantic": "dft_length", "role": "dft_length", "required": false },
+    "axis": {
+      "kind": "tensor",
+      "semantic": "axis",
+      "role": "axis",
+      "dtype": "int32",
+      "narrowing": "checked",
+      "required": false
+    },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 128, "TILED_REAL_WORKGROUP_SIZE": 256, "TILED_REAL_MIN_LENGTH": 512 },
+  "derive": {
+    "dftAxisAttr": "ranks.input - 2",
+    "dftAttrOuter": "outer(shapes.input, dftAxisAttr)",
+    "dftAttrInputInner": "inner(shapes.input, dftAxisAttr) / dim(shapes.input, ranks.input - 1)",
+    "dftAttrOutputOuter": "outer(shapes.output, dftAxisAttr)",
+    "dftAttrOutputInner": "inner(shapes.output, dftAxisAttr) / dim(shapes.output, ranks.output - 1)",
+    "rank3BatchDispatchFits": "dim(shapes.input, 0) <= device.limits.maxComputeWorkgroupsPerDimension",
+    "rank3Axis1SharedFftFits": "rank3BatchDispatchFits and pow2ceil(dim(shapes.input, 1)) == dim(shapes.input, 1) and dim(shapes.input, 1) >= 4 and dim(shapes.input, 1) * 16 <= device.limits.maxComputeWorkgroupStorageSize",
+    "rank3Axis1StorageFftFits": "rank3BatchDispatchFits and pow2ceil(dim(shapes.input, 1)) == dim(shapes.input, 1) and dim(shapes.input, 1) >= 4 and dim(shapes.input, 1) * 16 > device.limits.maxComputeWorkgroupStorageSize",
+    "rank3Axis1FftLog2": "log2ceil(dim(shapes.input, 1))",
+    "rank3Axis1ScratchFits": "dim(shapes.input, 0) * dim(shapes.input, 1) * 2 * 4 <= device.limits.maxStorageBufferBindingSize and dim(shapes.input, 0) * dim(shapes.input, 1) * 2 * 4 <= device.limits.maxBufferSize"
+  },
+  "bindingSets": {
+    "contiguousTransform": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "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": "inner", "type": "u32", "value": 1 }] }
+      }
+    ],
+    "dftScratch": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "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": "inner", "type": "u32", "value": 1 }] }
+      }
+    ],
+    "fftLen": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "dft_len_in",
+        "arg": "dft_length",
+        "semantic": "dft_length",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32",
+        "length": 1
+      },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "inner", "type": "u32", "value": 1 }] }
+      }
+    ],
+    "fftLenStatic": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "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": "inner", "type": "u32", "value": 1 }] }
+      }
+    ],
+    "fftAxis": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      {
+        "name": "axis_in",
+        "arg": "axis",
+        "semantic": "axis",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32",
+        "length": 1
+      },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "inner", "type": "u32", "value": 1 }] }
+      }
+    ],
+    "sharedFftTransform": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "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": "inner",
+              "type": "u32",
+              "value": "inner(shapes.input, dftAxisAttr) / dim(shapes.input, ranks.input - 1)"
+            }
+          ]
+        }
+      }
+    ],
+    "rank2Naive": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "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) / dim(shapes.output, 1)" },
+            { "name": "signalLength", "type": "u32", "value": "dim(shapes.input, 0)" },
+            { "name": "inputComponents", "type": "u32", "value": "dim(shapes.input, 1)" },
+            { "name": "outputComponents", "type": "u32", "value": "dim(shapes.output, 1)" },
+            { "name": "outputLength", "type": "u32", "value": "dim(shapes.output, 0)" },
+            { "name": "inverse", "type": "u32", "value": "attrs.inverse" },
+            { "name": "onesided", "type": "u32", "value": "attrs.onesided" }
+          ]
+        }
+      }
+    ],
+    "rank3Naive": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "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) / dim(shapes.output, 2)" },
+            { "name": "signalLength", "type": "u32", "value": "dim(shapes.input, 1)" },
+            { "name": "inputComponents", "type": "u32", "value": "dim(shapes.input, 2)" },
+            { "name": "outputComponents", "type": "u32", "value": "dim(shapes.output, 2)" },
+            { "name": "outputLength", "type": "u32", "value": "dim(shapes.output, 1)" },
+            { "name": "inverse", "type": "u32", "value": "attrs.inverse" },
+            { "name": "onesided", "type": "u32", "value": "attrs.onesided" }
+          ]
+        }
+      }
+    ],
+    "rank4Naive": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "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) / dim(shapes.output, 3)" },
+            { "name": "inputD1", "type": "u32", "value": "dim(shapes.input, 1)" },
+            { "name": "inputD2", "type": "u32", "value": "dim(shapes.input, 2)" },
+            { "name": "outputD1", "type": "u32", "value": "dim(shapes.output, 1)" },
+            { "name": "outputD2", "type": "u32", "value": "dim(shapes.output, 2)" },
+            { "name": "inputComponents", "type": "u32", "value": "dim(shapes.input, 3)" },
+            { "name": "outputComponents", "type": "u32", "value": "dim(shapes.output, 3)" },
+            { "name": "inverse", "type": "u32", "value": "attrs.inverse" },
+            { "name": "onesided", "type": "u32", "value": "attrs.onesided" }
+          ]
+        }
+      }
+    ],
+    "genNone": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "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) / dim(shapes.output, ranks.output - 1)" },
+            { "name": "rank", "type": "u32", "value": "4" },
+            { "name": "inputComponents", "type": "u32", "value": "dim(shapes.input, ranks.input - 1)" },
+            { "name": "outputComponents", "type": "u32", "value": "dim(shapes.output, ranks.output - 1)" },
+            { "name": "inverse", "type": "u32", "value": "attrs.inverse" },
+            { "name": "onesided", "type": "u32", "value": "attrs.onesided" },
+            { "name": "inD0", "type": "u32", "value": 1 },
+            { "name": "inD1", "type": "u32", "value": 1 },
+            { "name": "inD2", "type": "u32", "value": 1 },
+            { "name": "inD3", "type": "u32", "value": 1 },
+            { "name": "inD4", "type": "u32", "value": "dftAttrOuter" },
+            { "name": "inD5", "type": "u32", "value": "dim(shapes.input, dftAxisAttr)" },
+            { "name": "inD6", "type": "u32", "value": "dftAttrInputInner" },
+            { "name": "outD0", "type": "u32", "value": 1 },
+            { "name": "outD1", "type": "u32", "value": 1 },
+            { "name": "outD2", "type": "u32", "value": 1 },
+            { "name": "outD3", "type": "u32", "value": 1 },
+            { "name": "outD4", "type": "u32", "value": "dftAttrOutputOuter" },
+            { "name": "outD5", "type": "u32", "value": "dim(shapes.output, dftAxisAttr)" },
+            { "name": "outD6", "type": "u32", "value": "dftAttrOutputInner" },
+            { "name": "axisAttr", "type": "u32", "value": "1" },
+            { "name": "dftLengthAttr", "type": "u32", "value": 0 }
+          ]
+        }
+      }
+    ],
+    "genAxis": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "axis_in",
+        "arg": "axis",
+        "semantic": "axis",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32",
+        "length": 1
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.output) / dim(shapes.output, ranks.output - 1)" },
+            { "name": "rank", "type": "u32", "value": "ranks.input" },
+            { "name": "inputComponents", "type": "u32", "value": "dim(shapes.input, ranks.input - 1)" },
+            { "name": "outputComponents", "type": "u32", "value": "dim(shapes.output, ranks.output - 1)" },
+            { "name": "inverse", "type": "u32", "value": "attrs.inverse" },
+            { "name": "onesided", "type": "u32", "value": "attrs.onesided" },
+            { "name": "inD0", "type": "u32", "value": "dim(shapes.input, ranks.input - 8) if ranks.input >= 8 else 1" },
+            { "name": "inD1", "type": "u32", "value": "dim(shapes.input, ranks.input - 7) if ranks.input >= 7 else 1" },
+            { "name": "inD2", "type": "u32", "value": "dim(shapes.input, ranks.input - 6) if ranks.input >= 6 else 1" },
+            { "name": "inD3", "type": "u32", "value": "dim(shapes.input, ranks.input - 5) if ranks.input >= 5 else 1" },
+            { "name": "inD4", "type": "u32", "value": "dim(shapes.input, ranks.input - 4) if ranks.input >= 4 else 1" },
+            { "name": "inD5", "type": "u32", "value": "dim(shapes.input, ranks.input - 3) if ranks.input >= 3 else 1" },
+            { "name": "inD6", "type": "u32", "value": "dim(shapes.input, ranks.input - 2)" },
+            {
+              "name": "outD0",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 8) if ranks.output >= 8 else 1"
+            },
+            {
+              "name": "outD1",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 7) if ranks.output >= 7 else 1"
+            },
+            {
+              "name": "outD2",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 6) if ranks.output >= 6 else 1"
+            },
+            {
+              "name": "outD3",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 5) if ranks.output >= 5 else 1"
+            },
+            {
+              "name": "outD4",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 4) if ranks.output >= 4 else 1"
+            },
+            {
+              "name": "outD5",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 3) if ranks.output >= 3 else 1"
+            },
+            { "name": "outD6", "type": "u32", "value": "dim(shapes.output, ranks.output - 2)" },
+            { "name": "dftLengthAttr", "type": "u32", "value": 0 }
+          ]
+        }
+      }
+    ],
+    "genLen": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "dft_len_in",
+        "arg": "dft_length",
+        "semantic": "dft_length",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32",
+        "length": 1
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.output) / dim(shapes.output, ranks.output - 1)" },
+            { "name": "rank", "type": "u32", "value": "4" },
+            { "name": "inputComponents", "type": "u32", "value": "dim(shapes.input, ranks.input - 1)" },
+            { "name": "outputComponents", "type": "u32", "value": "dim(shapes.output, ranks.output - 1)" },
+            { "name": "inverse", "type": "u32", "value": "attrs.inverse" },
+            { "name": "onesided", "type": "u32", "value": "attrs.onesided" },
+            { "name": "inD0", "type": "u32", "value": 1 },
+            { "name": "inD1", "type": "u32", "value": 1 },
+            { "name": "inD2", "type": "u32", "value": 1 },
+            { "name": "inD3", "type": "u32", "value": 1 },
+            { "name": "inD4", "type": "u32", "value": "dftAttrOuter" },
+            { "name": "inD5", "type": "u32", "value": "dim(shapes.input, dftAxisAttr)" },
+            { "name": "inD6", "type": "u32", "value": "dftAttrInputInner" },
+            { "name": "outD0", "type": "u32", "value": 1 },
+            { "name": "outD1", "type": "u32", "value": 1 },
+            { "name": "outD2", "type": "u32", "value": 1 },
+            { "name": "outD3", "type": "u32", "value": 1 },
+            { "name": "outD4", "type": "u32", "value": "dftAttrOutputOuter" },
+            { "name": "outD5", "type": "u32", "value": "dim(shapes.output, dftAxisAttr)" },
+            { "name": "outD6", "type": "u32", "value": "dftAttrOutputInner" },
+            { "name": "axisAttr", "type": "u32", "value": "1" }
+          ]
+        }
+      }
+    ],
+    "genAxisLen": [
+      {
+        "name": "input",
+        "arg": "input",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "f32"
+      },
+      { "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
+      {
+        "name": "dft_len_in",
+        "arg": "dft_length",
+        "semantic": "dft_length",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32",
+        "length": 1
+      },
+      {
+        "name": "axis_in",
+        "arg": "axis",
+        "semantic": "axis",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "i32",
+        "length": 1
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [
+            { "name": "count", "type": "u32", "value": "numel(shapes.output) / dim(shapes.output, ranks.output - 1)" },
+            { "name": "rank", "type": "u32", "value": "ranks.input" },
+            { "name": "inputComponents", "type": "u32", "value": "dim(shapes.input, ranks.input - 1)" },
+            { "name": "outputComponents", "type": "u32", "value": "dim(shapes.output, ranks.output - 1)" },
+            { "name": "inverse", "type": "u32", "value": "attrs.inverse" },
+            { "name": "onesided", "type": "u32", "value": "attrs.onesided" },
+            { "name": "inD0", "type": "u32", "value": "dim(shapes.input, ranks.input - 8) if ranks.input >= 8 else 1" },
+            { "name": "inD1", "type": "u32", "value": "dim(shapes.input, ranks.input - 7) if ranks.input >= 7 else 1" },
+            { "name": "inD2", "type": "u32", "value": "dim(shapes.input, ranks.input - 6) if ranks.input >= 6 else 1" },
+            { "name": "inD3", "type": "u32", "value": "dim(shapes.input, ranks.input - 5) if ranks.input >= 5 else 1" },
+            { "name": "inD4", "type": "u32", "value": "dim(shapes.input, ranks.input - 4) if ranks.input >= 4 else 1" },
+            { "name": "inD5", "type": "u32", "value": "dim(shapes.input, ranks.input - 3) if ranks.input >= 3 else 1" },
+            { "name": "inD6", "type": "u32", "value": "dim(shapes.input, ranks.input - 2)" },
+            {
+              "name": "outD0",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 8) if ranks.output >= 8 else 1"
+            },
+            {
+              "name": "outD1",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 7) if ranks.output >= 7 else 1"
+            },
+            {
+              "name": "outD2",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 6) if ranks.output >= 6 else 1"
+            },
+            {
+              "name": "outD3",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 5) if ranks.output >= 5 else 1"
+            },
+            {
+              "name": "outD4",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 4) if ranks.output >= 4 else 1"
+            },
+            {
+              "name": "outD5",
+              "type": "u32",
+              "value": "dim(shapes.output, ranks.output - 3) if ranks.output >= 3 else 1"
+            },
+            { "name": "outD6", "type": "u32", "value": "dim(shapes.output, ranks.output - 2)" }
+          ]
+        }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "batch_runtime_axis1_fft",
+      "priority": 30,
+      "when": ["present.axis", "numel(shapes.axis) == 1", "not present.dft_length", "ranks.input == 3", "ranks.output == 3", "attrs.onesided == 0", "dim(shapes.input, 0) == dim(shapes.output, 0)", "dim(shapes.input, 1) == dim(shapes.output, 1)", "rank3Axis1SharedFftFits", "dim(shapes.input, 2) >= 1", "dim(shapes.input, 2) <= 2", "dim(shapes.output, 2) == 2"],
+      "passes": [
+        {
+          "id": "axis1_fft",
+          "name": "DFT.BatchRuntimeAxis1FftShared",
+          "source": {
+            "shader": "fft-stockham-shared.wgsl.jinja",
+            "inputs": {
+              "mode": "\"dft\"",
+              "fftLength": "dim(shapes.input, 1)",
+              "numBins": "dim(shapes.output, 1)",
+              "inComp": "dim(shapes.input, 2)",
+              "inverse": "attrs.inverse != 0",
+              "runtimeAxis": true,
+              "axisRank": 3,
+              "expectedAxis": 1
+            }
+          },
+          "bindings": "fftAxis",
+          "dispatch": { "x": "dim(shapes.input, 0)" }
+        },
+        {
+          "id": "other_axis_general",
+          "name": "DFT.BatchRuntimeOtherAxisGeneral",
+          "source": {
+            "shader": "dft-general.wgsl.jinja",
+            "inputs": { "hasAxisInput": true, "hasDftLengthInput": false, "skipRuntimeAxis1": true }
+          },
+          "bindings": "genAxis",
+          "dispatch": {
+            "threads": "numel(shapes.output) / dim(shapes.output, ranks.output - 1)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "batch_axis1_fft_storage_onesided",
+      "priority": 32,
+      "when": ["not present.axis", "not present.dft_length", "ranks.input == 3", "ranks.output == 3", "dftAxisAttr == 1", "attrs.onesided != 0", "attrs.inverse == 0", "dim(shapes.input, 0) == dim(shapes.output, 0)", "dim(shapes.input, 0) > 0", "rank3BatchDispatchFits", "rank3Axis1StorageFftFits", "rank3Axis1ScratchFits", "dim(shapes.output, 1) == floor(dim(shapes.input, 1) / 2) + 1", "dim(shapes.input, 2) == 1", "dim(shapes.output, 2) == 2"],
+      "intermediates": [
+        { "id": "fft_scratch", "dtype": "float32", "shape": "[dim(shapes.input, 0) * dim(shapes.input, 1) * 2]" }
+      ],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.BatchAxis1FftStorageOnesided",
+          "source": {
+            "shader": "fft-radix2-dit-storage.wgsl.jinja",
+            "inputs": {
+              "mode": "\"dft_scratch\"",
+              "fftLength": "dim(shapes.input, 1)",
+              "log2Length": "rank3Axis1FftLog2",
+              "numBins": "dim(shapes.output, 1)",
+              "inComp": 1,
+              "inverse": false
+            }
+          },
+          "bindings": "dftScratch",
+          "dispatch": { "x": "dim(shapes.input, 0)" }
+        }
+      ]
+    },
+    {
+      "id": "batch_axis1_fft_storage_inplace",
+      "priority": 31,
+      "when": ["not present.axis", "not present.dft_length", "ranks.input == 3", "ranks.output == 3", "dftAxisAttr == 1", "attrs.onesided == 0", "dim(shapes.input, 0) == dim(shapes.output, 0)", "rank3BatchDispatchFits", "rank3Axis1StorageFftFits", "dim(shapes.output, 1) == dim(shapes.input, 1)", "dim(shapes.input, 2) >= 1", "dim(shapes.input, 2) <= 2", "dim(shapes.output, 2) == 2"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.BatchAxis1FftStorageInplace",
+          "source": {
+            "shader": "fft-radix2-dit-storage.wgsl.jinja",
+            "inputs": {
+              "mode": "\"dft\"",
+              "fftLength": "dim(shapes.input, 1)",
+              "log2Length": "rank3Axis1FftLog2",
+              "numBins": "dim(shapes.output, 1)",
+              "inComp": "dim(shapes.input, 2)",
+              "inverse": "attrs.inverse != 0"
+            }
+          },
+          "bindings": "contiguousTransform",
+          "dispatch": { "x": "dim(shapes.input, 0)" }
+        }
+      ]
+    },
+    {
+      "id": "batch_axis1_tiled_real_nonpow2",
+      "priority": 29,
+      "requires": { "limits": { "maxComputeInvocationsPerWorkgroup": 256, "maxComputeWorkgroupSizeX": 256 } },
+      "when": ["not present.axis", "not present.dft_length", "ranks.input == 3", "ranks.output == 3", "dftAxisAttr == 1", "attrs.inverse == 0", "dim(shapes.input, 0) == dim(shapes.output, 0)", "dim(shapes.input, 0) > 0", "rank3BatchDispatchFits", "dim(shapes.input, 1) >= tunables.TILED_REAL_MIN_LENGTH", "pow2ceil(dim(shapes.input, 1)) != dim(shapes.input, 1)", "dim(shapes.input, 1) * 4 <= device.limits.maxComputeWorkgroupStorageSize", "dim(shapes.input, 2) == 1", "dim(shapes.output, 2) == 2", "((attrs.onesided == 0 and dim(shapes.output, 1) == dim(shapes.input, 1)) or (attrs.onesided != 0 and dim(shapes.output, 1) == floor(dim(shapes.input, 1) / 2) + 1))", "ceilDiv(dim(shapes.output, 1), tunables.TILED_REAL_WORKGROUP_SIZE) <= device.limits.maxComputeWorkgroupsPerDimension"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.BatchAxis1TiledRealNonPowerOfTwo",
+          "source": {
+            "shader": "dft-tiled-real.wgsl.jinja",
+            "inputs": {
+              "mode": "\"dft\"",
+              "transformLength": "dim(shapes.input, 1)",
+              "numBins": "dim(shapes.output, 1)",
+              "workgroupSize": "tunables.TILED_REAL_WORKGROUP_SIZE"
+            }
+          },
+          "bindings": "contiguousTransform",
+          "dispatch": {
+            "x": "ceilDiv(dim(shapes.output, 1), tunables.TILED_REAL_WORKGROUP_SIZE)",
+            "y": "dim(shapes.input, 0)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "axis_fft_shared",
+      "priority": 30,
+      "when": ["not present.axis", "not present.dft_length", "ranks.input == ranks.output", "((ranks.input == 2 and dftAxisAttr == 0) or (ranks.input == 3 and dftAxisAttr == 1 and dim(shapes.input, 0) == dim(shapes.output, 0)) or (ranks.input == 4 and dftAxisAttr == 2 and dim(shapes.input, 0) == dim(shapes.output, 0) and dim(shapes.input, 1) == dim(shapes.output, 1)))", "pow2ceil(dim(shapes.input, dftAxisAttr)) == dim(shapes.input, dftAxisAttr)", "dim(shapes.input, dftAxisAttr) >= 4", "dim(shapes.input, dftAxisAttr) * 16 <= device.limits.maxComputeWorkgroupStorageSize", "outer(shapes.input, dftAxisAttr) * inner(shapes.input, dftAxisAttr) / dim(shapes.input, ranks.input - 1) <= device.limits.maxComputeWorkgroupsPerDimension", "dim(shapes.input, ranks.input - 1) >= 1", "dim(shapes.input, ranks.input - 1) <= 2", "((attrs.onesided == 0 and dim(shapes.output, dftAxisAttr) == dim(shapes.input, dftAxisAttr) and dim(shapes.output, ranks.input - 1) == 2) or (attrs.onesided != 0 and attrs.inverse == 0 and dim(shapes.input, ranks.input - 1) == 1 and dim(shapes.output, dftAxisAttr) == floor(dim(shapes.input, dftAxisAttr) / 2) + 1 and dim(shapes.output, ranks.input - 1) == 2))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.AxisFftShared",
+          "source": {
+            "shader": "fft-stockham-shared.wgsl.jinja",
+            "inputs": {
+              "mode": "\"dft\"",
+              "fftLength": "dim(shapes.input, dftAxisAttr)",
+              "numBins": "dim(shapes.output, dftAxisAttr)",
+              "inComp": "dim(shapes.input, ranks.input - 1)",
+              "inverse": "attrs.inverse != 0"
+            }
+          },
+          "bindings": "sharedFftTransform",
+          "dispatch": {
+            "x": "outer(shapes.input, dftAxisAttr) * inner(shapes.input, dftAxisAttr) / dim(shapes.input, ranks.input - 1)"
+          }
+        }
+      ]
+    },
+    {
+      "id": "batch_axis1_fft_explicit_equal_length",
+      "priority": 31,
+      "when": ["not present.axis", "present.dft_length", "numel(shapes.dft_length) == 1", "ranks.input == 3", "ranks.output == 3", "dftAxisAttr == 1", "dim(shapes.input, 0) == dim(shapes.output, 0)", "rank3Axis1SharedFftFits", "dim(shapes.input, 2) >= 1", "dim(shapes.input, 2) <= 2", "((attrs.onesided == 0 and dim(shapes.output, 1) == dim(shapes.input, 1) and dim(shapes.output, 2) == 2) or (attrs.onesided != 0 and attrs.inverse == 0 and dim(shapes.input, 2) == 1 and dim(shapes.output, 1) == floor(dim(shapes.input, 1) / 2) + 1 and dim(shapes.output, 2) == 2))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.BatchAxis1FftSharedExplicitLength",
+          "source": {
+            "shader": "fft-stockham-shared.wgsl.jinja",
+            "inputs": {
+              "mode": "\"dft\"",
+              "fftLength": "dim(shapes.input, 1)",
+              "numBins": "dim(shapes.output, 1)",
+              "inComp": "dim(shapes.input, 2)",
+              "inverse": "attrs.inverse != 0"
+            }
+          },
+          "bindings": "fftLenStatic",
+          "dispatch": { "x": "dim(shapes.input, 0)" }
+        }
+      ]
+    },
+    {
+      "id": "rank2_axis0",
+      "priority": 20,
+      "when": ["not present.axis", "not present.dft_length", "ranks.input == 2", "ranks.output == 2", "dftAxisAttr == 0", "dim(shapes.input, 1) >= 1", "dim(shapes.input, 1) <= 2", "((attrs.onesided == 0 and dim(shapes.output, 0) == dim(shapes.input, 0) and dim(shapes.output, 1) == 2) or (attrs.onesided != 0 and attrs.inverse == 0 and dim(shapes.input, 1) == 1 and dim(shapes.output, 0) == floor(dim(shapes.input, 0) / 2) + 1 and dim(shapes.output, 1) == 2) or (attrs.onesided != 0 and attrs.inverse != 0 and dim(shapes.input, 1) == 2 and dim(shapes.input, 0) == floor(dim(shapes.output, 0) / 2) + 1 and dim(shapes.output, 1) == 1))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.Rank2Axis0",
+          "source": { "shader": "dft-contiguous-naive.wgsl.jinja", "inputs": { "batched": false } },
+          "bindings": "rank2Naive",
+          "dispatch": {
+            "threads": "numel(shapes.output) / dim(shapes.output, 1)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "batch_axis1",
+      "priority": 10,
+      "when": ["not present.axis", "not present.dft_length", "ranks.input == 3", "ranks.output == 3", "dftAxisAttr == 1", "dim(shapes.input, 0) == dim(shapes.output, 0)", "dim(shapes.input, 2) >= 1", "dim(shapes.input, 2) <= 2", "((attrs.onesided == 0 and dim(shapes.output, 1) == dim(shapes.input, 1) and dim(shapes.output, 2) == 2) or (attrs.onesided != 0 and attrs.inverse == 0 and dim(shapes.input, 2) == 1 and dim(shapes.output, 1) == floor(dim(shapes.input, 1) / 2) + 1 and dim(shapes.output, 2) == 2) or (attrs.onesided != 0 and attrs.inverse != 0 and dim(shapes.input, 2) == 2 and dim(shapes.input, 1) == floor(dim(shapes.output, 1) / 2) + 1 and dim(shapes.output, 2) == 1))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT",
+          "source": { "shader": "dft-contiguous-naive.wgsl.jinja", "inputs": { "batched": true } },
+          "bindings": "rank3Naive",
+          "dispatch": {
+            "threads": "numel(shapes.output) / dim(shapes.output, 2)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "rank4_axis2",
+      "priority": 10,
+      "when": ["not present.axis", "not present.dft_length", "ranks.input == 4", "ranks.output == 4", "dftAxisAttr == 2", "dim(shapes.input, 0) == dim(shapes.output, 0)", "dim(shapes.input, 1) == dim(shapes.output, 1)", "dim(shapes.input, 3) >= 1", "dim(shapes.input, 3) <= 2", "((attrs.onesided == 0 and dim(shapes.output, 2) == dim(shapes.input, 2) and dim(shapes.output, 3) == 2) or (attrs.onesided != 0 and attrs.inverse == 0 and dim(shapes.input, 3) == 1 and dim(shapes.output, 2) == floor(dim(shapes.input, 2) / 2) + 1 and dim(shapes.output, 3) == 2) or (attrs.onesided != 0 and attrs.inverse != 0 and dim(shapes.input, 3) == 2 and dim(shapes.input, 2) == floor(dim(shapes.output, 2) / 2) + 1 and dim(shapes.output, 3) == 1))"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.Rank4Axis2",
+          "shader": "dft-rank4.wgsl.jinja",
+          "bindings": "rank4Naive",
+          "dispatch": {
+            "threads": "numel(shapes.output) / dim(shapes.output, 3)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "general_attr",
+      "priority": 5,
+      "when": ["not present.axis", "not present.dft_length", "ranks.input >= 2", "dftAxisAttr <= ranks.input - 2", "ranks.input == ranks.output", "dim(shapes.input, ranks.input - 1) >= 1", "dim(shapes.input, ranks.input - 1) <= 2", "dim(shapes.output, ranks.output - 1) >= 1", "dim(shapes.output, ranks.output - 1) <= 2"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.General",
+          "source": {
+            "shader": "dft-general.wgsl.jinja",
+            "inputs": { "hasAxisInput": false, "hasDftLengthInput": false, "skipRuntimeAxis1": false }
+          },
+          "bindings": "genNone",
+          "dispatch": {
+            "threads": "numel(shapes.output) / dim(shapes.output, ranks.output - 1)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "general_axis_len",
+      "priority": 5,
+      "when": ["present.axis and numel(shapes.axis) == 1", "present.dft_length and numel(shapes.dft_length) == 1", "ranks.input >= 2", "ranks.input <= 8", "ranks.input == ranks.output", "dim(shapes.input, ranks.input - 1) >= 1", "dim(shapes.input, ranks.input - 1) <= 2", "dim(shapes.output, ranks.output - 1) >= 1", "dim(shapes.output, ranks.output - 1) <= 2"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.General",
+          "source": {
+            "shader": "dft-general.wgsl.jinja",
+            "inputs": { "hasAxisInput": true, "hasDftLengthInput": true, "skipRuntimeAxis1": false }
+          },
+          "bindings": "genAxisLen",
+          "dispatch": {
+            "threads": "numel(shapes.output) / dim(shapes.output, ranks.output - 1)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "general_axis",
+      "priority": 5,
+      "when": ["present.axis and numel(shapes.axis) == 1", "not present.dft_length", "ranks.input >= 2", "ranks.input <= 8", "ranks.input == ranks.output", "dim(shapes.input, ranks.input - 1) >= 1", "dim(shapes.input, ranks.input - 1) <= 2", "dim(shapes.output, ranks.output - 1) >= 1", "dim(shapes.output, ranks.output - 1) <= 2"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.General",
+          "source": {
+            "shader": "dft-general.wgsl.jinja",
+            "inputs": { "hasAxisInput": true, "hasDftLengthInput": false, "skipRuntimeAxis1": false }
+          },
+          "bindings": "genAxis",
+          "dispatch": {
+            "threads": "numel(shapes.output) / dim(shapes.output, ranks.output - 1)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "general_len",
+      "priority": 5,
+      "when": ["not present.axis", "present.dft_length and numel(shapes.dft_length) == 1", "ranks.input >= 2", "dftAxisAttr <= ranks.input - 2", "ranks.input == ranks.output", "dim(shapes.input, ranks.input - 1) >= 1", "dim(shapes.input, ranks.input - 1) <= 2", "dim(shapes.output, ranks.output - 1) >= 1", "dim(shapes.output, ranks.output - 1) <= 2"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "DFT.General",
+          "source": {
+            "shader": "dft-general.wgsl.jinja",
+            "inputs": { "hasAxisInput": false, "hasDftLengthInput": true, "skipRuntimeAxis1": false }
+          },
+          "bindings": "genLen",
+          "dispatch": {
+            "threads": "numel(shapes.output) / dim(shapes.output, ranks.output - 1)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,23 @@

+{
+  "name": "ai.onnx.DFT",
+  "id": "_ai_onnx_dft_webgpu_dbca2e5",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "jbPgs835f9m0aqqbSqY7zJJ+sRdHAdkjUzasXhjUEMM=",
+      "dft-contiguous-naive.wgsl.jinja": "HizdEBeuURKgl0Zxsz3xs62LJA6qmXvTYeYgC/XW5FE=",
+      "dft-general.wgsl.jinja": "fRvFH2YPz8DQLaTWSwXfScJErAFyeim5yqU39h+DYIc=",
+      "dft-rank4.wgsl.jinja": "QYtIZb6gpPc4fPIglkSt50Nr3KgbmuqoZk3EsuCgS3w=",
+      "dft-tiled-real.wgsl.jinja": "hbc3jiXVFIDZzBLo4O5hKQDqPIrNFdNktE85DhjWAzY=",
+      "fft-radix2-dit-storage.wgsl.jinja": "xLm5AhxAvxBdznovw1gQsQsaJGxusSmdtcaxUs00i2w=",
+      "fft-stockham-shared.wgsl.jinja": "7J4RLO2Q8puehryHRRHVhrrNxjNik5718G/qu6zoF6I=",
+      "manifest.json": "9IvgHvMofuhCPOTfi61UqnxP5X0IdE11o+JDI1V+dSs=",
+      "test.json": "pmZjTxGLBuze2IB7hjXRl0ks8BfzGeJ1OWCsaNDm/ao="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.DFT" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1374 @@

+{
+  "op": "ai.onnx.DFT",
+  "fixtureArrays": {
+    "ort_rank2_complex_axis0_output_output": [36, 0, -4, 9.65685, -4, 4, -4, 1.65685, -4, 0, -4, -1.65685, -4, -4, -4, -9.65685]
+  },
+  "cases": [
+    {
+      "name": "dispatch_cliff_batch_axis1_over_8m_elements",
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": { "input": { "dtype": "float32", "shape": [2099200, 4, 1] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [2099200, 4, 2], "tolerance": 0.001 } }
+    },
+    {
+      "name": "dispatch_cliff_rank4_axis1_over_8m_elements",
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1050000, 4, 2, 1] },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1050000, 4, 2, 2], "tolerance": 0.001 } }
+    },
+    {
+      "name": "dispatch_cliff_rank4_axis2_over_8m_elements",
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": { "input": { "dtype": "float32", "shape": [1050000, 2, 4, 1] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1050000, 2, 4, 2], "tolerance": 0.001 } }
+    },
+    {
+      "name": "real_full_default_axis",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 0.0, -1.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "default_axis_minus2_rank2",
+      "provenance": {
+        "source": "ONNX DFT-20 schema",
+        "notes": "Omitting the optional axis input must use -2. For rank 2 that normalizes to axis 0, the only signal axis."
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4, 1], "data": { "kind": "values", "values": [1.0, 0.0, 0.0, 0.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "tolerance": 0.00001,
+          "data": { "kind": "values", "values": [1.0, 0.0, 1.0, 0.0, 1.0, 0.0, 1.0, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "default_axis_minus2_rank4",
+      "provenance": {
+        "source": "ONNX DFT-20 schema",
+        "notes": "Omitting the optional axis input must use -2. For rank 4 that is axis 2, independently transforming each [batch, inner] signal."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 4, 1],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, 0.0, 2.0, 0.0, 0.0, 0.0] }
+        }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 2, 4, 2],
+          "tolerance": 0.00001,
+          "data": {
+            "kind": "values",
+            "values": [1.0, 0.0, 1.0, 0.0, 1.0, 0.0, 1.0, 0.0, 2.0, 0.0, 2.0, 0.0, 2.0, 0.0, 2.0, 0.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "real_subnormal_impulse_preserves_all_bins_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 input flushes to zero on GPU, every output bin reads 0 instead of the subnormal value (tolerance 0)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_radix2",
+        "notes": "Focused amplitude edge: a real impulse of 1e-40 has a DFT with every real frequency bin equal to that subnormal value."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4, 1],
+          "data": { "kind": "values", "values": [1e-40, 0.0, 0.0, 0.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "complex_subnormal_impulse_preserves_all_bins_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 input flushes to zero on GPU, every output bin reads 0 instead of the subnormal value (tolerance 0)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_radix2",
+        "notes": "Complex companion to the real impulse edge: a subnormal complex impulse at sample zero should appear unchanged in every frequency bin."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4, 2],
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "complex_inverse_subnormal_dc_preserves_time_signal_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 input flushes to zero on GPU, every output bin reads 0 instead of the subnormal value (tolerance 0)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_inverse",
+        "notes": "Inverse companion: a subnormal DC bin scaled by 1/N should produce a valid finite subnormal constant signal."
+      },
+      "attrs": { "inverse": 1, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4, 2],
+          "data": { "kind": "values", "values": [4e-40, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "rank2_axis0_subnormal_impulse_preserves_all_bins_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 input flushes to zero on GPU, every output bin reads 0 instead of the subnormal value (tolerance 0)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_radix2",
+        "notes": "Rank-2 axis-0 FFT companion: a real subnormal impulse should preserve the same subnormal real value in every bin."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [4, 1], "data": { "kind": "values", "values": [1e-40, 0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "rank4_axis1_subnormal_impulse_preserves_all_bins_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 input flushes to zero on GPU, every output bin reads 0 instead of the subnormal value (tolerance 0)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_radix2",
+        "notes": "Rank-4 axis-1 FFT companion: the subnormal impulse is along the first signal axis with one inner position."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4, 1, 1],
+          "data": { "kind": "values", "values": [1e-40, 0.0, 0.0, 0.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "rank4_axis2_subnormal_impulse_preserves_all_bins_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 input flushes to zero on GPU, every output bin reads 0 instead of the subnormal value (tolerance 0)."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_radix2",
+        "notes": "Rank-4 axis-2 FFT companion: the subnormal impulse is along the second signal axis with one outer channel."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 1],
+          "data": { "kind": "values", "values": [1e-40, 0.0, 0.0, 0.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 4, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "negative_axis_minus2_rank3",
+      "provenance": {
+        "source": "https://github.com/onnx/onnx/blob/main/docs/Operators.md#DFT",
+        "notes": "DFT axis accepts negative values in [-r, -2]; for rank 3, axis=-2 is the last signal axis and matches axis=1."
+      },
+      "attrs": {},
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 0.0, -1.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_axis_input_overrides_default_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_invertible_real",
+        "notes": "Opset-20 axis is an input tensor. This same-shape axis=0 case is a regression lock ensuring the runtime axis overrides the default axis attribute."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4, 3, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.21, "scale": 4.0 }
+        },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 3, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_axis_input_overrides_default_rank4_axis2",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_invertible_real",
+        "notes": "Rank-4 projection of the opset-20 axis-input path. The output shape also matches the default axis, so this locks correct runtime-axis decoding rather than allowing a silent axis-1 transform."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4, 4, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.17, "scale": 3.0 }
+        },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 4, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "real_onesided_even",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 0.0, -1.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "complex_inverse",
+      "attrs": { "inverse": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4, 2],
+          "data": { "kind": "values", "values": [2.0, 0.0, 1.0, -1.0, 0.0, 0.0, 1.0, 1.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "real_onesided_odd_length",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 5, 1],
+          "data": { "kind": "values", "values": [0.0, 1.0, 2.0, 3.0, 4.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "real_full_batch2",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 4, 1],
+          "data": { "kind": "values", "values": [1.0, 0.0, -1.0, 0.0, 0.0, 2.0, 0.0, -2.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 4, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "complex_inverse_length3",
+      "attrs": { "inverse": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 3, 2],
+          "data": { "kind": "values", "values": [6.0, 0.0, -3.0, 1.7320508, -3.0, -1.7320508] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "real_radix2_onesided_length8_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_radix2_onesided"
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 8, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 5, 2], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "fft_batch_axis1_len256_batch4_real_onesided_perf_compact",
+      "provenance": {
+        "notes": "Compact correctness sibling for the batch-axis onesided FFT benchmark; preserves real input, axis=1, power-of-two length, and onesided output without benchmark-scale tensors."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": { "input": { "dtype": "float32", "shape": [4, 256, 1] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 129, 2], "tolerance": 0.0002 } }
+    },
+    {
+      "name": "complex_full_batch2_length3",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 2],
+          "data": { "kind": "values", "values": [1.0, 0.0, 0.0, 1.0, -1.0, 0.0, 2.0, 0.0, 0.0, -1.0, 1.0, 1.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "real_naive_length5_full_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_naive"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 5, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 5, 2], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "real_radix2_length8_full_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_radix2"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 8, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 8, 2], "tolerance": 0.00004 } }
+    },
+    {
+      "name": "real_naive_length5_onesided_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_naive_onesided"
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 5, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 2], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "ort_dft_length_truncates_forward",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_Float_naive",
+        "notes": "Valid opset-20 dft_length input that truncates the transform length below the input extent."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 5, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 3, 2],
+          "tolerance": 0.00002,
+          "data": { "kind": "values", "values": [6.0, 0.0, -1.5, 0.866025, -1.5, -0.866025] }
+        }
+      }
+    },
+    {
+      "name": "ort_dft_length_zero_pads_forward",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_Float_naive",
+        "notes": "Valid opset-20 dft_length input that zero-pads the transform length beyond the input extent."
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [5] } },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 5, 2],
+          "tolerance": 0.00002,
+          "data": {
+            "kind": "values",
+            "values": [6.0, 0.0, -0.809017, -3.665469, 0.309017, 1.677599, 0.309017, -1.677599, -0.809017, 3.665469]
+          }
+        }
+      }
+    },
+    {
+      "name": "complex_inverse_length5_ort",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT17_Float_inverse"
+      },
+      "attrs": { "inverse": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 5, 2],
+          "data": {
+            "kind": "values",
+            "values": [15.0, 0.0, -2.499999, 3.440955, -2.5, 0.8123, -2.499999, -0.812299, -2.500003, -3.440953]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 5, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_irfft_radix2_even_length",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_IRFFT_radix2"
+      },
+      "attrs": { "onesided": 1, "inverse": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 5, 2],
+          "data": { "kind": "values", "values": [36.0, 0.0, -4.0, 9.65685, -4.0, 4.0, -4.0, 1.65685, -4.0, 0.0] }
+        }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 8, 1],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_irfft_radix2_rank3_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_IRFFT_radix2",
+        "notes": "Axis-0 rank-3 projection of ORT's even-length IRFFT case."
+      },
+      "attrs": { "onesided": 1, "inverse": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [5, 1, 2],
+          "data": { "kind": "values", "values": [36.0, 0.0, -4.0, 9.65685, -4.0, 4.0, -4.0, 1.65685, -4.0, 0.0] }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [8] } },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [8, 1, 1],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_irfft_naive_odd_length",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_IRFFT_naive",
+        "notes": "Odd-length IRFFT needs explicit dft_length to disambiguate from the even default."
+      },
+      "attrs": { "onesided": 1, "inverse": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 3, 2],
+          "data": { "kind": "values", "values": [15.0, 0.0, -2.499999, 3.440955, -2.5, 0.8123] }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [5] } },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 5, 1],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_invertible_real_rank4_axis1_projection",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_invertible_real",
+        "notes": "Single forward DFT projection from ORT's rank-4 real invertibility sweep."
+      },
+      "attrs": {},
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 2, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19, "scale": 8.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 2, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_invertible_complex_rank4_axis2_projection",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_invertible_complex",
+        "notes": "Single forward DFT projection from ORT's rank-4 complex invertibility sweep."
+      },
+      "attrs": {},
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 2, 4, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 5.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 4, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_invertible_complex_rank4_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_invertible_complex",
+        "notes": "Single forward DFT projection from ORT's rank-4 complex invertibility sweep."
+      },
+      "attrs": {},
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [4, 2, 3, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.17, "scale": 3.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 2, 3, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_invertible_real_rank3_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_invertible_real",
+        "notes": "Single forward DFT projection from ORT's axis-0 invertibility sweep."
+      },
+      "attrs": {},
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [4, 3, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.09, "cosStep": 0.21, "scale": 4.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 3, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_rank2_complex_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_2D_complex"
+      },
+      "attrs": {},
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 0.0, 2.0, 0.0, 3.0, 0.0, 4.0, 0.0, 5.0, 0.0, 6.0, 0.0, 7.0, 0.0, 8.0, 0.0]
+          }
+        }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_rank2_complex_axis0_output_output" } }
+        }
+      }
+    },
+    {
+      "name": "ort_rank2_real_onesided_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_2D_real"
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [8, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [5, 2],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [36.0, 0.0, -4.0, 9.65685, -4.0, 4.0, -4.0, 1.65685, -4.0, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_rank2_real_full_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_2D_real"
+      },
+      "attrs": {},
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [8, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [8] } },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [8, 2],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/ort_rank2_complex_axis0_output_output" } }
+        }
+      }
+    },
+    {
+      "name": "ort_rank2_irfft_even_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_2D_complex_onesided_inverse"
+      },
+      "attrs": { "onesided": 1, "inverse": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [5, 2],
+          "data": { "kind": "values", "values": [36.0, 0.0, -4.0, 9.65685, -4.0, 4.0, -4.0, 1.65685, -4.0, 0.0] }
+        }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [8, 1],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_rank2_irfft_odd_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_IRFFT_naive",
+        "notes": "Rank-2 axis-0 projection of ORT's odd-length IRFFT case; explicit output length disambiguates odd from even."
+      },
+      "attrs": { "onesided": 1, "inverse": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [15.0, 0.0, -2.499999, 3.440955, -2.5, 0.8123] }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [5] } },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [5, 1],
+          "tolerance": 0.0001,
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0] }
+        }
+      }
+    },
+    {
+      "name": "onnx_backend_dft_axis1_rank4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dft",
+        "notes": "Uses the official shape and attributes with deterministic generated data to avoid committing the large random tensor literal."
+      },
+      "attrs": {},
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 10, 10, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.031, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 10, 10, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_dft_axis2_rank4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dft_axis",
+        "notes": "Uses the official shape and attributes with deterministic generated data to avoid committing the large random tensor literal."
+      },
+      "attrs": {},
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 10, 10, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.019, "cosStep": 0.031, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 10, 10, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_dft_inverse_axis1_rank4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dft_inverse",
+        "notes": "Uses the official shape and attributes with deterministic generated data to avoid committing the large random tensor literal."
+      },
+      "attrs": { "inverse": 1 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 10, 10, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.029, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 10, 10, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_dft_irfft_axis1_rank4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dft_irfft",
+        "notes": "Uses the official shape and attributes with deterministic generated data to avoid committing the large random tensor literal."
+      },
+      "attrs": { "inverse": 1, "onesided": 1 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 6, 10, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.037, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 10, 10, 1], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_dft_rfft_axis1_rank4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dft_rfft",
+        "notes": "Uses the official shape and attributes with deterministic generated data to avoid committing the large random tensor literal."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 10, 10, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.021, "cosStep": 0.034, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 6, 10, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_dft_opset19_axis1_rank4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dft_opset19",
+        "notes": "Uses the official shape and attributes with deterministic generated data to avoid committing the large random tensor literal."
+      },
+      "attrs": {},
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 10, 10, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.026, "cosStep": 0.041, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 10, 10, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_dft_axis_opset19_axis2_rank4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dft_axis_opset19",
+        "notes": "Uses the official shape and attributes with deterministic generated data to avoid committing the large random tensor literal."
+      },
+      "attrs": {},
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 10, 10, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.026, "cosStep": 0.041, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 10, 10, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_dft_inverse_opset19_axis1_rank4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dft_inverse_opset19",
+        "notes": "Uses the official shape and attributes with deterministic generated data to avoid committing the large random tensor literal."
+      },
+      "attrs": { "inverse": 1 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 10, 10, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.027, "cosStep": 0.039, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 10, 10, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_dft_irfft_opset19_axis1_rank4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dft_irfft_opset19",
+        "notes": "Uses the official shape and attributes with deterministic generated data to avoid committing the large random tensor literal."
+      },
+      "attrs": { "inverse": 1, "onesided": 1 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 6, 10, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.025, "cosStep": 0.043, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 10, 10, 1], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "onnx_backend_dft_rfft_opset19_axis1_rank4",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_dft_rfft_opset19",
+        "notes": "Uses the official shape and attributes with deterministic generated data to avoid committing the large random tensor literal."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 10, 10, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.046, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 6, 10, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "fft_rank2_axis0_real_onesided_len256",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [256, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.021, "cosStep": 0.031, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [129, 2], "tolerance": 0.001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "fft_rank2_axis0_complex_full_len512",
+      "attrs": {},
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [512, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.029, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [512, 2], "tolerance": 0.02, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "fft_batch_axis1_real_full_len1024_batch2",
+      "attrs": {},
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 1024, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.041, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 1024, 2], "tolerance": 0.1, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "tiled_real_nonpow2_len513_full",
+      "provenance": {
+        "notes": "Compact correctness lock for the large non-power-of-two tiled direct DFT, including its 64-lane tunable dispatch, final partial frequency workgroup, and twiddle resets."
+      },
+      "tunables": { "WORKGROUP_SIZE": 64 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 513, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.027, "scale": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 513, 2], "tolerance": 0.03, "relTolerance": 0.0002 } }
+    },
+    {
+      "name": "tiled_real_nonpow2_len4095_full_boundary",
+      "provenance": {
+        "notes": "Boundary-length accuracy lock for the non-power-of-two tiled DFT. Exercises all shared-memory slots and many twiddle resets at the largest supported signal length."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4095, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0013, "cosStep": 0.0027, "scale": 0.05 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4095, 2], "tolerance": 0.1, "relTolerance": 0.0005 } }
+    },
+    {
+      "name": "fft_batch_axis1_complex_inverse_len512",
+      "attrs": { "inverse": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 512, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.037, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 512, 2], "tolerance": 0.0001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "fft_batch_axis1_real_onesided_len2048",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2048, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.019, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1025, 2], "tolerance": 0.06, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "fft_batch_axis1_real_full_len4096_storage_inplace",
+      "provenance": {
+        "notes": "Locks the portable single-dispatch storage-buffer radix-2 route above the shared Stockham size ceiling."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4096, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.009, "cosStep": 0.017, "scale": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4096, 2], "tolerance": 0.35, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "fft_batch_axis1_complex_inverse_len4096_storage_inplace",
+      "provenance": {
+        "notes": "The storage-buffer radix-2 route only takes over above the shared-memory ceiling at length 2048, and every case that long was a forward transform of a real signal. This one is the inverse of a complex signal, which is the route's other three knobs at once: the twiddle sign, the two-component load, and the 1/N scaling."
+      },
+      "attrs": { "inverse": 1, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4096, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.009, "cosStep": 0.017, "scale": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4096, 2], "tolerance": 0.0001, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "fft_batch_axis1_real_onesided_len4096_storage_scratch",
+      "provenance": {
+        "notes": "Covers the one-dispatch storage-buffer DIT route when a onesided output is too small to serve as its own FFT scratch."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4096, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.009, "cosStep": 0.017, "scale": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2049, 2], "tolerance": 0.35, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "fft_rank4_axis1_real_full_len256_inner3",
+      "attrs": {},
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 256, 3, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.027, "cosStep": 0.033, "scale": 1.0 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [2, 256, 3, 2], "tolerance": 0.015, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "fft_rank4_axis1_real_onesided_len16_inner5",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 16, 5, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.043, "cosStep": 0.021, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 9, 5, 2], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "fft_rank4_axis2_real_onesided_len512",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 3, 512, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.039, "scale": 1.0 }
+        }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [2, 3, 257, 2], "tolerance": 0.002, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "empty_zero_dim",
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": { "input": { "dtype": "float32", "shape": [0, 4, 1], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 4, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_caseB_empty",
+      "attrs": { "onesided": 1 },
+      "inputs": { "input": { "dtype": "float32", "shape": [1, 0, 1], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 1, 2],
+          "data": { "kind": "values", "values": [0.0, 0.0] },
+          "tolerance": 0.001
+        }
+      }
+    },
+    {
+      "name": "irfft_naive_axis1_over_128_dispatch_guard",
+      "provenance": {
+        "source": "synthetic",
+        "test": "naive batch_axis1 IRFFT, numel(output) > 128 dispatch coverage",
+        "notes": "Odd output length 255 (non-power-of-2) forces the naive batch_axis1 variant (FFT radix-2 paths don't apply). dim(output,2)==1 on the onesided-inverse path, so the dispatch numerator must be numel(output)/dim(output,2)=255, not numel(output)/2=127: the latter only launches ceil(127/WG) workgroups and silently drops outputs >= ~128. numel(output)=255 > WG(128) so the drop is observable."
+      },
+      "attrs": { "inverse": 1, "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 128, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.019, "scale": 1.0 }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [255] } },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 255, 1], "tolerance": 0.01, "relTolerance": 0.01 } }
+    },
+    {
+      "name": "dft_length_default_axis_no_axis_input",
+      "provenance": {
+        "notes": "dft_length input present, axis input omitted (default transform axis 1). Exercises the general_len variant: forward DFT of a real rank-3 signal where dft_length matches the signal extent."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 0.0, -1.0] }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [4] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "general_axis_input_len2048_naive_vs_fft_batch32",
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [32, 2048, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0007, "cosStep": 0.0013, "scale": 1.0 }
+        },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [32, 2048, 2], "tolerance": 0.05, "relTolerance": 0.01 } }
+    },
+    {
+      "name": "general_axis_input_gidy_fold_over_8m_rows",
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2099200, 4, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.013, "cosStep": 0.019, "scale": 1.0 }
+        },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [2099200, 4, 2], "tolerance": 0.0001, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "whisper_nfft400_naive_batch_axis1_onesided",
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [200, 400, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.011, "cosStep": 0.017, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [200, 201, 2], "tolerance": 0.02, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "irfft_batch_axis1_len512_reconstruct_real",
+      "attrs": { "inverse": 1, "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [8, 257, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.023, "cosStep": 0.031, "scale": 1.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [8, 512, 1], "tolerance": 0.02, "relTolerance": 0.01 } }
+    },
+    {
+      "name": "general_len_rank4_axis1_dftlength_zeropad_len6",
+      "provenance": {
+        "notes": "dft_length=6 zero-pads a rank-4 axis-1 signal of length 4 up to 6 samples, forcing the general_len variant (rank-4 coordinate decode + zero-pad guard n>=input_axis_length). Untested rank-4 dft_length path."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 4, 3, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.037, "cosStep": 0.061, "scale": 2.0 }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [6] } }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [2, 6, 3, 2], "tolerance": 0.0002, "relTolerance": 0.0001 }
+      }
+    },
+    {
+      "name": "general_irfft_rank4_axis2_odd_len5_no_nyquist",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.DFT20_IRFFT_naive",
+        "notes": "Rank-4 axis-2 odd-length IRFFT via general_axis_len: odd output length 5 has no exact Nyquist bin, so no weight-2 doubling should occur. Exercises the general kernel onesided_inverse branch + rank-4 coordinate decode."
+      },
+      "attrs": { "inverse": 1, "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 2, 3, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.029, "cosStep": 0.047, "scale": 3.0 }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [5] } },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 5, 1], "tolerance": 0.0005, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "general_axis_len_negative_axis_rank3_dftlength_truncate",
+      "provenance": {
+        "notes": "Negative axis input (-2 -> axis 1 for rank 3) combined with a truncating dft_length (6 -> 4) forces general_axis_len and exercises runtime resolve_axis() negative normalization together with signal truncation. Untested combination."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 6, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.033, "cosStep": 0.052, "scale": 2.0 }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4, 2], "tolerance": 0.0002, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "general_len_rank2_axis0_dftlength_zeropad_odd",
+      "provenance": {
+        "notes": "Rank-2 axis-0 with dft_length=5 zero-padding a length-3 signal forces general_len; validates the rank-2 (n_logical=1, pad=2) coordinate decode of dft-general under zero-padding."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [5] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5, 2], "tolerance": 0.00002 } }
+    },
+    {
+      "name": "fft_explicit_equal_length128_batch4_onesided",
+      "provenance": {
+        "notes": "Correctness lock for the optimized explicit-dft_length route when the runtime length equals the power-of-two signal length."
+      },
+      "attrs": { "inverse": 0, "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [4, 128, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.033, "cosStep": 0.052, "scale": 2.0 }
+        },
+        "dft_length": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [128] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 65, 2], "tolerance": 0.002, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "fft_runtime_axis1_full_len16_batch2",
+      "provenance": {
+        "notes": "A scalar runtime axis selects axis 1, checking the guarded Stockham FFT and the complementary general-axis pass on the same input."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 16, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.031, "cosStep": 0.047, "scale": 1.5 }
+        },
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [2, 16, 2], "tolerance": 0.00002, "relTolerance": 0.00001 }
+      }
+    },
+    {
+      "name": "fft_batch_axis1_real_full_len8192_storage_inplace_generic",
+      "provenance": { "notes": "Non-4096 correctness lock for the geometry-derived storage-buffer radix-2 route." },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 8192, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0045, "cosStep": 0.0085, "scale": 0.25 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 8192, 2], "tolerance": 0.7, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "fft_batch_axis1_real_onesided_len8192_storage_scratch_generic",
+      "provenance": { "notes": "Non-4096 correctness lock for the geometry-derived one-sided radix-2 scratch route." },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 8192, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0045, "cosStep": 0.0085, "scale": 0.25 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 4097, 2], "tolerance": 0.7, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "tiled_real_nonpow2_len4097_onesided_above_old_boundary",
+      "provenance": {
+        "notes": "Proves the tiled direct transform is bounded by device workgroup storage rather than the former 4095 shape cutoff."
+      },
+      "attrs": { "onesided": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 4097, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.0013, "cosStep": 0.0027, "scale": 0.05 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2049, 2], "tolerance": 0.15, "relTolerance": 0.001 } }
+    },
+    {
+      "name": "rank5_batched_axis3",
+      "provenance": {
+        "notes": "ONNX DFT permits arbitrary batch rank around the transformed axis. With the axis given as an attribute the general kernel needs only (outer, axis, inner), so the batch rank never reaches it."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 4, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.29, "scale": 0.5 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 1, 2, 4, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "rank6_axis2_batch_and_inner",
+      "provenance": {
+        "notes": "A rank the packing has never seen, with batch dims on both sides of the transform axis (outer 4, inner 2). If any rank ceiling had merely been raised rather than removed, this is the case that finds it."
+      },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "axis": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 2, 3, 1, 2, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.23, "cosStep": 0.41, "scale": 0.6 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 3, 1, 2, 2], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "rank8_default_axis_contract_ceiling",
+      "provenance": { "notes": "Exercises the manifest's advertised rank-8 ceiling with the ONNX default axis (-2)." },
+      "attrs": { "inverse": 0, "onesided": 0 },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 1, 1, 2, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0] }
+        }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 1, 1, 2, 2],
+          "tolerance": 0.00001,
+          "data": { "kind": "values", "values": [3.0, 0.0, -1.0, 0.0] }
+        }
+      }
+    }
+  ]
+}