sync 91d990483a17
Browse files- README.md +15 -11
- build/webgpu/bench.json +0 -1
- build/webgpu/dft-tiled-real.wgsl.jinja +9 -10
- build/webgpu/fft-radix2-dit-shared-inplace.wgsl.jinja +15 -16
- build/webgpu/fft-radix2-dit-storage-pass.wgsl.jinja +35 -38
- build/webgpu/fft-radix2-dit-storage.wgsl.jinja +8 -8
- build/webgpu/fft-real-packed-stockham-shared.wgsl.jinja +2 -2
- build/webgpu/fft-stockham-shared.wgsl.jinja +32 -44
- build/webgpu/fft-stockham-twiddles.wgsl.jinja +9 -10
- build/webgpu/manifest.json +199 -415
- build/webgpu/metadata.json +29 -14
- build/webgpu/stft-general.wgsl.jinja +22 -23
- build/webgpu/test.json +6 -7
README.md
CHANGED
|
@@ -18,20 +18,20 @@ See the [ONNX `STFT` spec](https://onnx.ai/onnx/operators/onnx__STFT.html) for t
|
|
| 18 |
|
| 19 |
## Inputs
|
| 20 |
|
| 21 |
-
| Name |
|
| 22 |
-
| --- | --- | --- | --- | --- | --- |
|
| 23 |
-
| `signal` | `
|
| 24 |
-
| `window` | `
|
| 25 |
|
| 26 |
## Outputs
|
| 27 |
|
| 28 |
-
| Name |
|
| 29 |
-
| --- | --- | --- | --- | --- | --- |
|
| 30 |
-
| `output` | `
|
| 31 |
|
| 32 |
## Runtime arguments
|
| 33 |
|
| 34 |
-
| Name | Kind |
|
| 35 |
| --- | --- | --- | --- | --- |
|
| 36 |
| `frameStep` | `i32` | `frame_step` | Positive number of signal positions between the starts of consecutive frames. | required |
|
| 37 |
| `frameLength` | `i32` | `frame_length` | Optional number of signal values in each frame; must equal the window length when both are supplied. | optional |
|
|
@@ -52,7 +52,7 @@ Default values (overridable per request):
|
|
| 52 |
|
| 53 |
## Files
|
| 54 |
|
| 55 |
-
- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
|
| 56 |
- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
|
| 57 |
- [`test.json`](build/webgpu/test.json) — correctness cases
|
| 58 |
- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
|
|
@@ -67,10 +67,14 @@ Default values (overridable per request):
|
|
| 67 |
|
| 68 |
## Use with `@huggingface/kernels`
|
| 69 |
|
| 70 |
-
|
| 71 |
-
|
|
|
|
|
|
|
|
|
|
| 72 |
|
| 73 |
The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
|
|
|
|
| 74 |
|
| 75 |
Replace each `*Data` placeholder with a typed array containing the corresponding input data.
|
| 76 |
|
|
|
|
| 18 |
|
| 19 |
## Inputs
|
| 20 |
|
| 21 |
+
| Name | Logical dtype | Rank | Shape | Description | Presence |
|
| 22 |
+
| --- | --- | --- | --- | --- | --- |
|
| 23 |
+
| `signal` | `T` | `3` | — | Real input signal with shape `(batch_size, signal_length, 1)`, or complex input with shape `(batch_size, signal_length, 2)`. | required |
|
| 24 |
+
| `window` | `T` | `1` | — | Optional 1-D window tensor with shape `(window_length)` to multiply by each frame before the DFT. | optional |
|
| 25 |
|
| 26 |
## Outputs
|
| 27 |
|
| 28 |
+
| Name | Logical dtype | Rank | Shape | Description | Presence |
|
| 29 |
+
| --- | --- | --- | --- | --- | --- |
|
| 30 |
+
| `output` | `T` | `4` | derived | STFT result with shape `(batch_size, frames, dft_unique_bins, 2)` in one-sided mode or `(batch_size, frames, frame_length, 2)` in two-sided mode, with real and imaginary parts in the last dimension. `frames = floor((signal_length - frame_length) / frame_step) + 1`; partial trailing frames are not emitted. | required |
|
| 31 |
|
| 32 |
## Runtime arguments
|
| 33 |
|
| 34 |
+
| Name | Kind | Upstream attribute | Description | Presence |
|
| 35 |
| --- | --- | --- | --- | --- |
|
| 36 |
| `frameStep` | `i32` | `frame_step` | Positive number of signal positions between the starts of consecutive frames. | required |
|
| 37 |
| `frameLength` | `i32` | `frame_length` | Optional number of signal values in each frame; must equal the window length when both are supplied. | optional |
|
|
|
|
| 52 |
|
| 53 |
## Files
|
| 54 |
|
| 55 |
+
- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
|
| 56 |
- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
|
| 57 |
- [`test.json`](build/webgpu/test.json) — correctness cases
|
| 58 |
- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
|
|
|
|
| 67 |
|
| 68 |
## Use with `@huggingface/kernels`
|
| 69 |
|
| 70 |
+
```sh
|
| 71 |
+
npm install --save-exact @huggingface/kernels@0.0.1-preview.2
|
| 72 |
+
```
|
| 73 |
+
|
| 74 |
+
Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
|
| 75 |
|
| 76 |
The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
|
| 77 |
+
It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
|
| 78 |
|
| 79 |
Replace each `*Data` placeholder with a typed array containing the corresponding input data.
|
| 80 |
|
build/webgpu/bench.json
CHANGED
|
@@ -1,5 +1,4 @@
|
|
| 1 |
{
|
| 2 |
-
"op": "ai.onnx.STFT",
|
| 3 |
"tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
|
| 4 |
"cases": [
|
| 5 |
{
|
|
|
|
| 1 |
{
|
|
|
|
| 2 |
"tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
|
| 3 |
"cases": [
|
| 4 |
{
|
build/webgpu/dft-tiled-real.wgsl.jinja
CHANGED
|
@@ -5,18 +5,17 @@
|
|
| 5 |
// phase drift without transcendental functions in the inner loop.
|
| 6 |
{{ env.wgsl.resourceDeclarations }}
|
| 7 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 8 |
-
// rational turn and route
|
| 9 |
-
//
|
| 10 |
//
|
| 11 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 12 |
// inverse transform. The result is (cos, sin).
|
| 13 |
// Backend-stable f32 sine/cosine core.
|
| 14 |
//
|
| 15 |
-
// Shader transcendental accuracy is implementation-defined
|
| 16 |
-
//
|
| 17 |
-
//
|
| 18 |
-
//
|
| 19 |
-
// truncation error is well below one f32 ULP over its documented interval.
|
| 20 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 21 |
const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
|
| 22 |
const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
|
|
@@ -84,9 +83,9 @@ fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
|
|
| 84 |
return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
|
| 85 |
}
|
| 86 |
|
| 87 |
-
const N: u32 = {{
|
| 88 |
-
const NUM_BINS: u32 = {{
|
| 89 |
-
const WG: u32 = {{
|
| 90 |
const TWIDDLE_RESET: u32 = 64u;
|
| 91 |
|
| 92 |
var<workgroup> signalTile: array<f32, N>;
|
|
|
|
| 5 |
// phase drift without transcendental functions in the inner loop.
|
| 6 |
{{ env.wgsl.resourceDeclarations }}
|
| 7 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 8 |
+
// rational turn and route polynomial evaluation through one trigonometric core
|
| 9 |
+
// for FFT, direct DFT, STFT, and cosine-window calculations.
|
| 10 |
//
|
| 11 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 12 |
// inverse transform. The result is (cos, sin).
|
| 13 |
// Backend-stable f32 sine/cosine core.
|
| 14 |
//
|
| 15 |
+
// Shader transcendental accuracy is implementation-defined. Each path retains
|
| 16 |
+
// its available phase representation, reduces it to [-pi, pi], then evaluates
|
| 17 |
+
// explicit polynomials in a fixed order. The half-pi core is degree 13 for sine
|
| 18 |
+
// and degree 12 for cosine.
|
|
|
|
| 19 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 20 |
const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
|
| 21 |
const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
|
|
|
|
| 83 |
return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
|
| 84 |
}
|
| 85 |
|
| 86 |
+
const N: u32 = {{ transformLength }}u;
|
| 87 |
+
const NUM_BINS: u32 = {{ numBins }}u;
|
| 88 |
+
const WG: u32 = {{ workgroupSizeSpec }}u;
|
| 89 |
const TWIDDLE_RESET: u32 = 64u;
|
| 90 |
|
| 91 |
var<workgroup> signalTile: array<f32, N>;
|
build/webgpu/fft-radix2-dit-shared-inplace.wgsl.jinja
CHANGED
|
@@ -2,25 +2,24 @@
|
|
| 2 |
//
|
| 3 |
// This route fills the gap between the double-buffered Stockham kernel
|
| 4 |
// (16 bytes/element) and the global-scratch DIT fallback. A single complex
|
| 5 |
-
// array needs only 8 bytes/element
|
| 6 |
-
// power-of-two
|
| 7 |
// Each workgroup owns one transform; all stages remain local and use uniform
|
| 8 |
// barriers, so only the input and requested spectrum touch global memory.
|
| 9 |
// STFT is forward-only, so the transform sign is fixed at -1.
|
| 10 |
{{ env.wgsl.resourceDeclarations }}
|
| 11 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 12 |
-
// rational turn and route
|
| 13 |
-
//
|
| 14 |
//
|
| 15 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 16 |
// inverse transform. The result is (cos, sin).
|
| 17 |
// Backend-stable f32 sine/cosine core.
|
| 18 |
//
|
| 19 |
-
// Shader transcendental accuracy is implementation-defined
|
| 20 |
-
//
|
| 21 |
-
//
|
| 22 |
-
//
|
| 23 |
-
// truncation error is well below one f32 ULP over its documented interval.
|
| 24 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 25 |
const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
|
| 26 |
const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
|
|
@@ -89,12 +88,12 @@ fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
|
|
| 89 |
}
|
| 90 |
|
| 91 |
|
| 92 |
-
const N: u32 = {{
|
| 93 |
const N_HALF: u32 = N / 2u;
|
| 94 |
-
const NUM_BINS: u32 = {{
|
| 95 |
const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
|
| 96 |
const SIGN: f32 = -1.0;
|
| 97 |
-
var<workgroup> fftData: array<vec2<f32>, {{
|
| 98 |
|
| 99 |
@compute @workgroup_size(WG, 1, 1)
|
| 100 |
fn main(@builtin(workgroup_id) wgid: vec3<u32>,
|
|
@@ -106,19 +105,19 @@ fn main(@builtin(workgroup_id) wgid: vec3<u32>,
|
|
| 106 |
let start = frame * params.frameStep;
|
| 107 |
|
| 108 |
for (var position = lid.x; position < N; position = position + WG) {
|
| 109 |
-
let n = reverseBits(position) >> (32u - {{
|
| 110 |
let signalIndex = start + n;
|
| 111 |
-
let inputBase = (batch * params.signalLength + signalIndex) * {{
|
| 112 |
let w = window[n];
|
| 113 |
var value = vec2<f32>(signal[inputBase] * w, 0.0);
|
| 114 |
-
{% if
|
| 115 |
value.y = signal[inputBase + 1u] * w;
|
| 116 |
{% endif %}
|
| 117 |
fftData[position] = value;
|
| 118 |
}
|
| 119 |
workgroupBarrier();
|
| 120 |
|
| 121 |
-
{% for stage in range(1,
|
| 122 |
{
|
| 123 |
const SPAN: u32 = 1u << {{ stage }}u;
|
| 124 |
const HALF: u32 = SPAN >> 1u;
|
|
|
|
| 2 |
//
|
| 3 |
// This route fills the gap between the double-buffered Stockham kernel
|
| 4 |
// (16 bytes/element) and the global-scratch DIT fallback. A single complex
|
| 5 |
+
// array needs only 8 bytes/element rather than Stockham's 16, extending the
|
| 6 |
+
// resident range by one power-of-two step for a fixed workgroup-storage limit.
|
| 7 |
// Each workgroup owns one transform; all stages remain local and use uniform
|
| 8 |
// barriers, so only the input and requested spectrum touch global memory.
|
| 9 |
// STFT is forward-only, so the transform sign is fixed at -1.
|
| 10 |
{{ env.wgsl.resourceDeclarations }}
|
| 11 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 12 |
+
// rational turn and route polynomial evaluation through one trigonometric core
|
| 13 |
+
// for FFT, direct DFT, STFT, and cosine-window calculations.
|
| 14 |
//
|
| 15 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 16 |
// inverse transform. The result is (cos, sin).
|
| 17 |
// Backend-stable f32 sine/cosine core.
|
| 18 |
//
|
| 19 |
+
// Shader transcendental accuracy is implementation-defined. Each path retains
|
| 20 |
+
// its available phase representation, reduces it to [-pi, pi], then evaluates
|
| 21 |
+
// explicit polynomials in a fixed order. The half-pi core is degree 13 for sine
|
| 22 |
+
// and degree 12 for cosine.
|
|
|
|
| 23 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 24 |
const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
|
| 25 |
const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
|
|
|
|
| 88 |
}
|
| 89 |
|
| 90 |
|
| 91 |
+
const N: u32 = {{ fftLength }}u;
|
| 92 |
const N_HALF: u32 = N / 2u;
|
| 93 |
+
const NUM_BINS: u32 = {{ numBins }}u;
|
| 94 |
const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
|
| 95 |
const SIGN: f32 = -1.0;
|
| 96 |
+
var<workgroup> fftData: array<vec2<f32>, {{ fftLength }}>;
|
| 97 |
|
| 98 |
@compute @workgroup_size(WG, 1, 1)
|
| 99 |
fn main(@builtin(workgroup_id) wgid: vec3<u32>,
|
|
|
|
| 105 |
let start = frame * params.frameStep;
|
| 106 |
|
| 107 |
for (var position = lid.x; position < N; position = position + WG) {
|
| 108 |
+
let n = reverseBits(position) >> (32u - {{ log2Length }}u);
|
| 109 |
let signalIndex = start + n;
|
| 110 |
+
let inputBase = (batch * params.signalLength + signalIndex) * {{ inComp }}u;
|
| 111 |
let w = window[n];
|
| 112 |
var value = vec2<f32>(signal[inputBase] * w, 0.0);
|
| 113 |
+
{% if inComp == 2 %}
|
| 114 |
value.y = signal[inputBase + 1u] * w;
|
| 115 |
{% endif %}
|
| 116 |
fftData[position] = value;
|
| 117 |
}
|
| 118 |
workgroupBarrier();
|
| 119 |
|
| 120 |
+
{% for stage in range(1, log2Length + 1) %}
|
| 121 |
{
|
| 122 |
const SPAN: u32 = 1u << {{ stage }}u;
|
| 123 |
const HALF: u32 = SPAN >> 1u;
|
build/webgpu/fft-radix2-dit-storage-pass.wgsl.jinja
CHANGED
|
@@ -1,15 +1,14 @@
|
|
| 1 |
-
// Globally staged radix-2 DIT FFT
|
| 2 |
-
//
|
| 3 |
-
//
|
| 4 |
-
//
|
| 5 |
-
//
|
| 6 |
-
// while preserving the same bit-reversed DIT algorithm and scratch layout.
|
| 7 |
// STFT is forward-only, so the transform sign is fixed at -1.
|
| 8 |
{{ env.wgsl.resourceDeclarations }}
|
| 9 |
-
{% if
|
| 10 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 11 |
-
// rational turn and route
|
| 12 |
-
//
|
| 13 |
//
|
| 14 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 15 |
// inverse transform. The result is (cos, sin).
|
|
@@ -17,11 +16,10 @@
|
|
| 17 |
{% set needPreciseTrigTwoPi = true %}
|
| 18 |
// Backend-stable f32 sine/cosine core.
|
| 19 |
//
|
| 20 |
-
// Shader transcendental accuracy is implementation-defined
|
| 21 |
-
//
|
| 22 |
-
//
|
| 23 |
-
//
|
| 24 |
-
// truncation error is well below one f32 ULP over its documented interval.
|
| 25 |
{% if needPreciseTrigCentered %}
|
| 26 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 27 |
{% endif %}
|
|
@@ -99,62 +97,61 @@ fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
|
|
| 99 |
|
| 100 |
|
| 101 |
{% endif %}
|
| 102 |
-
const N: u32 = {{
|
| 103 |
-
{% if
|
| 104 |
const N_HALF: u32 = N / 2u;
|
| 105 |
{% endif %}
|
| 106 |
-
{% if
|
| 107 |
-
const NUM_BINS: u32 = {{
|
| 108 |
{% endif %}
|
| 109 |
const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
|
| 110 |
-
{% if
|
| 111 |
-
const SPAN: u32 = 1u << {{
|
| 112 |
const HALF: u32 = SPAN >> 1u;
|
| 113 |
const SIGN: f32 = -1.0;
|
| 114 |
-
{% elif
|
| 115 |
-
const LOCAL_N: u32 = {{
|
| 116 |
const BLOCKS_PER_TRANSFORM: u32 = N / LOCAL_N;
|
| 117 |
const SIGN: f32 = -1.0;
|
| 118 |
-
var<workgroup> localData: array<vec2<f32>, {{
|
| 119 |
{% endif %}
|
| 120 |
|
| 121 |
-
{% if
|
| 122 |
-
return gid.x + gid.y *
|
| 123 |
}{% endif %}
|
| 124 |
|
| 125 |
|
| 126 |
@compute @workgroup_size(WG, 1, 1)
|
| 127 |
-
fn main({% if
|
| 128 |
{% else %}@builtin(local_invocation_id) lid: vec3<u32>,
|
| 129 |
-
@builtin(workgroup_id) wgid: vec3<u32>
|
| 130 |
-
{% endif %}
|
| 131 |
-
|
| 132 |
-
|
| 133 |
-
let flat = folded_global_index(gid, nwg);
|
| 134 |
{% endif %}
|
| 135 |
let transforms = arrayLength(&scratch) / (N * 2u);
|
| 136 |
|
| 137 |
-
{% if
|
| 138 |
-
let flatGroup = wgid.x + wgid.y *
|
| 139 |
let transform = flatGroup / BLOCKS_PER_TRANSFORM;
|
| 140 |
if (transform >= transforms) { return; }
|
| 141 |
let block = flatGroup - transform * BLOCKS_PER_TRANSFORM;
|
| 142 |
let position = block * LOCAL_N + lid.x;
|
| 143 |
-
let n = reverseBits(position) >> (32u - {{
|
| 144 |
let batch = transform / params.numFrames;
|
| 145 |
let frame = transform - batch * params.numFrames;
|
| 146 |
let start = frame * params.frameStep;
|
| 147 |
let signalIndex = start + n;
|
| 148 |
-
let inputBase = (batch * params.signalLength + signalIndex) * {{
|
| 149 |
let w = window[n];
|
| 150 |
var value = vec2<f32>(signal[inputBase] * w, 0.0);
|
| 151 |
-
{% if
|
| 152 |
value.y = signal[inputBase + 1u] * w;
|
| 153 |
{% endif %}
|
| 154 |
localData[lid.x] = value;
|
| 155 |
workgroupBarrier();
|
| 156 |
|
| 157 |
-
{% for stage in range(1,
|
| 158 |
if (lid.x < LOCAL_N / 2u) {
|
| 159 |
const LOCAL_SPAN_{{ stage }}: u32 = 1u << {{ stage }}u;
|
| 160 |
const LOCAL_HALF_{{ stage }}: u32 = LOCAL_SPAN_{{ stage }} >> 1u;
|
|
@@ -179,7 +176,7 @@ fn main({% if source.phase != "local_prefix" %}@builtin(global_invocation_id) gi
|
|
| 179 |
scratch[dst] = localData[lid.x].x;
|
| 180 |
scratch[dst + 1u] = localData[lid.x].y;
|
| 181 |
|
| 182 |
-
{% elif
|
| 183 |
let count = transforms * N_HALF;
|
| 184 |
if (flat >= count) { return; }
|
| 185 |
let transform = flat / N_HALF;
|
|
|
|
| 1 |
+
// Globally staged radix-2 DIT FFT for transforms not admitted to the
|
| 2 |
+
// workgroup-resident route. Each dispatch exposes every transform element or
|
| 3 |
+
// butterfly as independent work. Dispatch boundaries provide the required
|
| 4 |
+
// stage barriers while preserving the bit-reversed DIT algorithm and scratch
|
| 5 |
+
// layout.
|
|
|
|
| 6 |
// STFT is forward-only, so the transform sign is fixed at -1.
|
| 7 |
{{ env.wgsl.resourceDeclarations }}
|
| 8 |
+
{% if phase == "local_prefix" or phase == "stage" %}
|
| 9 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 10 |
+
// rational turn and route polynomial evaluation through one trigonometric core
|
| 11 |
+
// for FFT, direct DFT, STFT, and cosine-window calculations.
|
| 12 |
//
|
| 13 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 14 |
// inverse transform. The result is (cos, sin).
|
|
|
|
| 16 |
{% set needPreciseTrigTwoPi = true %}
|
| 17 |
// Backend-stable f32 sine/cosine core.
|
| 18 |
//
|
| 19 |
+
// Shader transcendental accuracy is implementation-defined. Each path retains
|
| 20 |
+
// its available phase representation, reduces it to [-pi, pi], then evaluates
|
| 21 |
+
// explicit polynomials in a fixed order. The half-pi core is degree 13 for sine
|
| 22 |
+
// and degree 12 for cosine.
|
|
|
|
| 23 |
{% if needPreciseTrigCentered %}
|
| 24 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 25 |
{% endif %}
|
|
|
|
| 97 |
|
| 98 |
|
| 99 |
{% endif %}
|
| 100 |
+
const N: u32 = {{ fftLength }}u;
|
| 101 |
+
{% if phase == "stage" %}
|
| 102 |
const N_HALF: u32 = N / 2u;
|
| 103 |
{% endif %}
|
| 104 |
+
{% if phase == "finish" %}
|
| 105 |
+
const NUM_BINS: u32 = {{ numBins }}u;
|
| 106 |
{% endif %}
|
| 107 |
const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
|
| 108 |
+
{% if phase == "stage" %}
|
| 109 |
+
const SPAN: u32 = 1u << {{ stage }}u;
|
| 110 |
const HALF: u32 = SPAN >> 1u;
|
| 111 |
const SIGN: f32 = -1.0;
|
| 112 |
+
{% elif phase == "local_prefix" %}
|
| 113 |
+
const LOCAL_N: u32 = {{ localLength }}u;
|
| 114 |
const BLOCKS_PER_TRANSFORM: u32 = N / LOCAL_N;
|
| 115 |
const SIGN: f32 = -1.0;
|
| 116 |
+
var<workgroup> localData: array<vec2<f32>, {{ localLength }}>;
|
| 117 |
{% endif %}
|
| 118 |
|
| 119 |
+
{% if phase != "local_prefix" %}fn folded_global_index(gid: vec3<u32>) -> u32 {
|
| 120 |
+
return gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
|
| 121 |
}{% endif %}
|
| 122 |
|
| 123 |
|
| 124 |
@compute @workgroup_size(WG, 1, 1)
|
| 125 |
+
fn main({% if phase != "local_prefix" %}@builtin(global_invocation_id) gid: vec3<u32>
|
| 126 |
{% else %}@builtin(local_invocation_id) lid: vec3<u32>,
|
| 127 |
+
@builtin(workgroup_id) wgid: vec3<u32>
|
| 128 |
+
{% endif %}) {
|
| 129 |
+
{% if phase != "local_prefix" %}
|
| 130 |
+
let flat = folded_global_index(gid);
|
|
|
|
| 131 |
{% endif %}
|
| 132 |
let transforms = arrayLength(&scratch) / (N * 2u);
|
| 133 |
|
| 134 |
+
{% if phase == "local_prefix" %}
|
| 135 |
+
let flatGroup = wgid.x + wgid.y * {{ DISPATCH_FOLD_WIDTH }}u;
|
| 136 |
let transform = flatGroup / BLOCKS_PER_TRANSFORM;
|
| 137 |
if (transform >= transforms) { return; }
|
| 138 |
let block = flatGroup - transform * BLOCKS_PER_TRANSFORM;
|
| 139 |
let position = block * LOCAL_N + lid.x;
|
| 140 |
+
let n = reverseBits(position) >> (32u - {{ log2Length }}u);
|
| 141 |
let batch = transform / params.numFrames;
|
| 142 |
let frame = transform - batch * params.numFrames;
|
| 143 |
let start = frame * params.frameStep;
|
| 144 |
let signalIndex = start + n;
|
| 145 |
+
let inputBase = (batch * params.signalLength + signalIndex) * {{ inComp }}u;
|
| 146 |
let w = window[n];
|
| 147 |
var value = vec2<f32>(signal[inputBase] * w, 0.0);
|
| 148 |
+
{% if inComp == 2 %}
|
| 149 |
value.y = signal[inputBase + 1u] * w;
|
| 150 |
{% endif %}
|
| 151 |
localData[lid.x] = value;
|
| 152 |
workgroupBarrier();
|
| 153 |
|
| 154 |
+
{% for stage in range(1, localLog2Length + 1) %}
|
| 155 |
if (lid.x < LOCAL_N / 2u) {
|
| 156 |
const LOCAL_SPAN_{{ stage }}: u32 = 1u << {{ stage }}u;
|
| 157 |
const LOCAL_HALF_{{ stage }}: u32 = LOCAL_SPAN_{{ stage }} >> 1u;
|
|
|
|
| 176 |
scratch[dst] = localData[lid.x].x;
|
| 177 |
scratch[dst + 1u] = localData[lid.x].y;
|
| 178 |
|
| 179 |
+
{% elif phase == "stage" %}
|
| 180 |
let count = transforms * N_HALF;
|
| 181 |
if (flat >= count) { return; }
|
| 182 |
let transform = flat / N_HALF;
|
build/webgpu/fft-radix2-dit-storage.wgsl.jinja
CHANGED
|
@@ -8,15 +8,15 @@
|
|
| 8 |
// row. Scratch-backed DFT supports one-sided output. STFT accepts real or
|
| 9 |
// complex windowed frames and copies NUM_BINS after the final stage.
|
| 10 |
{{ env.wgsl.resourceDeclarations }}
|
| 11 |
-
{% set STFT =
|
| 12 |
|
| 13 |
const PI: f32 = 3.141592653589793;
|
| 14 |
-
const N: u32 = {{
|
| 15 |
const N_HALF: u32 = N / 2u;
|
| 16 |
-
const LOG2_N: u32 = {{
|
| 17 |
-
const NUM_BINS: u32 = {{
|
| 18 |
const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
|
| 19 |
-
const SIGN: f32 = {{ "1.0" if
|
| 20 |
|
| 21 |
@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
|
| 22 |
fn main(
|
|
@@ -35,10 +35,10 @@ fn main(
|
|
| 35 |
for (var n = tid; n < N; n = n + WG) {
|
| 36 |
var value = vec2<f32>(0.0);
|
| 37 |
let signalIndex = start + n;
|
| 38 |
-
let inputBase = (batch * params.signalLength + signalIndex) * {{
|
| 39 |
let w = window[n];
|
| 40 |
value.x = signal[inputBase] * w;
|
| 41 |
-
{% if
|
| 42 |
value.y = signal[inputBase + 1u] * w;
|
| 43 |
{% endif %}
|
| 44 |
let reversed = reverseBits(n) >> (32u - LOG2_N);
|
|
@@ -48,7 +48,7 @@ fn main(
|
|
| 48 |
}
|
| 49 |
storageBarrier();
|
| 50 |
|
| 51 |
-
{% for stage in range(1,
|
| 52 |
{
|
| 53 |
const SPAN: u32 = 1u << {{ stage }}u;
|
| 54 |
const HALF: u32 = SPAN >> 1u;
|
|
|
|
| 8 |
// row. Scratch-backed DFT supports one-sided output. STFT accepts real or
|
| 9 |
// complex windowed frames and copies NUM_BINS after the final stage.
|
| 10 |
{{ env.wgsl.resourceDeclarations }}
|
| 11 |
+
{% set STFT = modeSpec == "stft" %}
|
| 12 |
|
| 13 |
const PI: f32 = 3.141592653589793;
|
| 14 |
+
const N: u32 = {{ fftLength }}u;
|
| 15 |
const N_HALF: u32 = N / 2u;
|
| 16 |
+
const LOG2_N: u32 = {{ log2Length }}u;
|
| 17 |
+
const NUM_BINS: u32 = {{ numBins }}u;
|
| 18 |
const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
|
| 19 |
+
const SIGN: f32 = {{ "1.0" if inverse else "-1.0" }};
|
| 20 |
|
| 21 |
@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
|
| 22 |
fn main(
|
|
|
|
| 35 |
for (var n = tid; n < N; n = n + WG) {
|
| 36 |
var value = vec2<f32>(0.0);
|
| 37 |
let signalIndex = start + n;
|
| 38 |
+
let inputBase = (batch * params.signalLength + signalIndex) * {{ inComp }}u;
|
| 39 |
let w = window[n];
|
| 40 |
value.x = signal[inputBase] * w;
|
| 41 |
+
{% if inComp == 2 %}
|
| 42 |
value.y = signal[inputBase + 1u] * w;
|
| 43 |
{% endif %}
|
| 44 |
let reversed = reverseBits(n) >> (32u - LOG2_N);
|
|
|
|
| 48 |
}
|
| 49 |
storageBarrier();
|
| 50 |
|
| 51 |
+
{% for stage in range(1, log2Length + 1) %}
|
| 52 |
{
|
| 53 |
const SPAN: u32 = 1u << {{ stage }}u;
|
| 54 |
const HALF: u32 = SPAN >> 1u;
|
build/webgpu/fft-real-packed-stockham-shared.wgsl.jinja
CHANGED
|
@@ -4,12 +4,12 @@
|
|
| 4 |
// complex transform, then reconstruct the N/2+1 non-redundant bins. This cuts
|
| 5 |
// the Stockham butterfly count almost in half while preserving the same fused
|
| 6 |
// window load as the full complex STFT kernel.
|
| 7 |
-
const N: u32 = {{
|
| 8 |
const PACKED_N: u32 = N / 2u;
|
| 9 |
const PACKED_HALF: u32 = PACKED_N / 2u;
|
| 10 |
const NUM_BINS: u32 = PACKED_N + 1u;
|
| 11 |
const FFT_LANES: u32 = {{ tunables.WORKGROUP_SIZE }}u;
|
| 12 |
-
const TRANSFORMS_PER_WG: u32 = {{
|
| 13 |
const WG: u32 = FFT_LANES * TRANSFORMS_PER_WG;
|
| 14 |
const STAGE_TWIDDLES: u32 = PACKED_N - 1u;
|
| 15 |
const SHARED_STRIDE: u32 = 2u * PACKED_N;
|
|
|
|
| 4 |
// complex transform, then reconstruct the N/2+1 non-redundant bins. This cuts
|
| 5 |
// the Stockham butterfly count almost in half while preserving the same fused
|
| 6 |
// window load as the full complex STFT kernel.
|
| 7 |
+
const N: u32 = {{ fftLength }}u;
|
| 8 |
const PACKED_N: u32 = N / 2u;
|
| 9 |
const PACKED_HALF: u32 = PACKED_N / 2u;
|
| 10 |
const NUM_BINS: u32 = PACKED_N + 1u;
|
| 11 |
const FFT_LANES: u32 = {{ tunables.WORKGROUP_SIZE }}u;
|
| 12 |
+
const TRANSFORMS_PER_WG: u32 = {{ transformsPerWorkgroup | default(1) }}u;
|
| 13 |
const WG: u32 = FFT_LANES * TRANSFORMS_PER_WG;
|
| 14 |
const STAGE_TWIDDLES: u32 = PACKED_N - 1u;
|
| 15 |
const SHARED_STRIDE: u32 = 2u * PACKED_N;
|
build/webgpu/fft-stockham-shared.wgsl.jinja
CHANGED
|
@@ -1,11 +1,12 @@
|
|
| 1 |
{{ env.wgsl.resourceDeclarations }}
|
| 2 |
-
{% set MIXED_TAIL =
|
| 3 |
-
{% set USE_TWIDDLE_TABLE =
|
|
|
|
| 4 |
{% set NEED_PRECISE_TWIDDLE = not USE_TWIDDLE_TABLE or MIXED_TAIL %}
|
| 5 |
{% if NEED_PRECISE_TWIDDLE %}
|
| 6 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 7 |
-
// rational turn and route
|
| 8 |
-
//
|
| 9 |
//
|
| 10 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 11 |
// inverse transform. The result is (cos, sin).
|
|
@@ -13,11 +14,10 @@
|
|
| 13 |
{% set needPreciseTrigTwoPi = true %}
|
| 14 |
// Backend-stable f32 sine/cosine core.
|
| 15 |
//
|
| 16 |
-
// Shader transcendental accuracy is implementation-defined
|
| 17 |
-
//
|
| 18 |
-
//
|
| 19 |
-
//
|
| 20 |
-
// truncation error is well below one f32 ULP over its documented interval.
|
| 21 |
{% if needPreciseTrigCentered %}
|
| 22 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 23 |
{% endif %}
|
|
@@ -96,49 +96,42 @@ fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
|
|
| 96 |
{% endif %}
|
| 97 |
{% set TAIL_ANCHOR = 32 %}
|
| 98 |
{% set REANCHOR_TAIL = MIXED_TAIL
|
| 99 |
-
and
|
| 100 |
|
| 101 |
// Radix-2 Stockham autosort FFT: one transform per workgroup, entirely in
|
| 102 |
// shared memory. log2(N) ping-pong stages alternate between the two halves of
|
| 103 |
// a 2*N scratch pair (no bit reversal; the autosort write pattern leaves the
|
| 104 |
// final stage in natural order). Threads stride over the N/2 butterflies of
|
| 105 |
-
// each stage with a workgroupBarrier between stages.
|
| 106 |
-
|
| 107 |
-
//
|
| 108 |
-
|
| 109 |
-
//
|
| 110 |
-
|
| 111 |
-
//
|
| 112 |
-
//
|
| 113 |
-
// inverse+onesided (IRFFT) combination is outside this specialization.
|
| 114 |
-
// Mode "logmel" runs one workgroup per spectrogram frame of
|
| 115 |
-
// a real mono signal: the frame load is centered (start - N/2) with reflection
|
| 116 |
-
// below zero and zeros at/after params.audioLength, and the epilogue folds the
|
| 117 |
-
// one-sided power spectrum through packed triangular mel bands into log10
|
| 118 |
-
// mel-major output instead of storing complex bins.
|
| 119 |
|
| 120 |
-
const N: u32 = {{
|
| 121 |
const N_HALF: u32 = N / 2u;
|
| 122 |
{% if MIXED_TAIL %}
|
| 123 |
-
const RADIX2_N: u32 = {{
|
| 124 |
const RADIX2_HALF: u32 = RADIX2_N / 2u;
|
| 125 |
const TAIL_RADIX: u32 = N / RADIX2_N;
|
| 126 |
{% if REANCHOR_TAIL %}
|
| 127 |
const TAIL_ANCHOR: u32 = {{ TAIL_ANCHOR }}u;
|
| 128 |
{% endif %}
|
| 129 |
{% endif %}
|
| 130 |
-
const NUM_BINS: u32 = {{
|
| 131 |
const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
|
| 132 |
{% if NEED_PRECISE_TWIDDLE %}
|
| 133 |
// Twiddle = exp(SIGN * i * 2*pi * k / span); SIGN flips for the inverse
|
| 134 |
-
// transform
|
| 135 |
-
const SIGN: f32 = {{ "1.0" if
|
| 136 |
{% endif %}
|
| 137 |
|
| 138 |
{% if MIXED_TAIL %}
|
| 139 |
-
// Walk consecutive powers of one unit root in each mixed-radix tail.
|
| 140 |
-
//
|
| 141 |
-
// polynomial evaluation per tail residue for every supported factorization.
|
| 142 |
fn complex_mul(a: vec2<f32>, b: vec2<f32>) -> vec2<f32> {
|
| 143 |
return vec2<f32>(a.x * b.x - a.y * b.y, a.x * b.y + a.y * b.x);
|
| 144 |
}
|
|
@@ -153,11 +146,8 @@ fn main(@builtin(workgroup_id) wgid: vec3<u32>,
|
|
| 153 |
@builtin(local_invocation_id) lid: vec3<u32>,
|
| 154 |
@builtin(num_workgroups) nwg: vec3<u32>) {
|
| 155 |
let tid = lid.x;
|
| 156 |
-
// Flat workgroup index
|
| 157 |
-
//
|
| 158 |
-
// one-workgroup-per-(batch,frame) [or per-(outer,inner)] grid across x/y and
|
| 159 |
-
// recover the flat index here. Collapses to wgid.x when nwg.y == 1 (every DFT
|
| 160 |
-
// dispatch and any STFT dispatch under the per-axis limit).
|
| 161 |
let flat = wgid.x + wgid.y * nwg.x;
|
| 162 |
|
| 163 |
let batch = flat / params.numFrames;
|
|
@@ -172,14 +162,14 @@ fn main(@builtin(workgroup_id) wgid: vec3<u32>,
|
|
| 172 |
let sample_n = n;
|
| 173 |
{% endif %}
|
| 174 |
let signal_index = start + sample_n;
|
| 175 |
-
let base = (batch * params.signalLength + signal_index) * {{
|
| 176 |
-
{% if
|
| 177 |
let w = 1.0;
|
| 178 |
{% else %}
|
| 179 |
let w = window[sample_n];
|
| 180 |
{% endif %}
|
| 181 |
let xr = signal[base] * w;
|
| 182 |
-
{% if
|
| 183 |
let xi = signal[base + 1u] * w;
|
| 184 |
{% else %}
|
| 185 |
let xi = 0.0;
|
|
@@ -193,8 +183,7 @@ fn main(@builtin(workgroup_id) wgid: vec3<u32>,
|
|
| 193 |
// applies w = exp(SIGN*i*pi*(j mod ns)/ns) to b, and writes
|
| 194 |
// dst[(j div ns)*2*ns + (j mod ns)] = a + w*b
|
| 195 |
// dst[(j div ns)*2*ns + (j mod ns) + ns] = a - w*b
|
| 196 |
-
// |angle| <= pi,
|
| 197 |
-
// sin/cos reduction (unlike the naive kernels' k*n angles).
|
| 198 |
var srcOff = 0u;
|
| 199 |
var dstOff = N;
|
| 200 |
{% if MIXED_TAIL %}
|
|
@@ -254,8 +243,7 @@ fn main(@builtin(workgroup_id) wgid: vec3<u32>,
|
|
| 254 |
}
|
| 255 |
|
| 256 |
// The final swap leaves the natural-order result at srcOff. The 1/N inverse
|
| 257 |
-
// scale is
|
| 258 |
-
// the naive kernels' division.
|
| 259 |
let outRow = (batch * params.numFrames + frame) * NUM_BINS;
|
| 260 |
for (var k = tid; k < NUM_BINS; k = k + WG) {
|
| 261 |
let outBase = (outRow + k) * 2u;
|
|
|
|
| 1 |
{{ env.wgsl.resourceDeclarations }}
|
| 2 |
+
{% set MIXED_TAIL = mixedRadixTail if mixedRadixTail is defined else false %}
|
| 3 |
+
{% set USE_TWIDDLE_TABLE = useTwiddleTable if useTwiddleTable is defined else false %}
|
| 4 |
+
{% set radix2Length = radix2Length | default(fftLength) %}
|
| 5 |
{% set NEED_PRECISE_TWIDDLE = not USE_TWIDDLE_TABLE or MIXED_TAIL %}
|
| 6 |
{% if NEED_PRECISE_TWIDDLE %}
|
| 7 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 8 |
+
// rational turn and route polynomial evaluation through one trigonometric core
|
| 9 |
+
// for FFT, direct DFT, STFT, and cosine-window calculations.
|
| 10 |
//
|
| 11 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 12 |
// inverse transform. The result is (cos, sin).
|
|
|
|
| 14 |
{% set needPreciseTrigTwoPi = true %}
|
| 15 |
// Backend-stable f32 sine/cosine core.
|
| 16 |
//
|
| 17 |
+
// Shader transcendental accuracy is implementation-defined. Each path retains
|
| 18 |
+
// its available phase representation, reduces it to [-pi, pi], then evaluates
|
| 19 |
+
// explicit polynomials in a fixed order. The half-pi core is degree 13 for sine
|
| 20 |
+
// and degree 12 for cosine.
|
|
|
|
| 21 |
{% if needPreciseTrigCentered %}
|
| 22 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 23 |
{% endif %}
|
|
|
|
| 96 |
{% endif %}
|
| 97 |
{% set TAIL_ANCHOR = 32 %}
|
| 98 |
{% set REANCHOR_TAIL = MIXED_TAIL
|
| 99 |
+
and fftLength / radix2Length > TAIL_ANCHOR %}
|
| 100 |
|
| 101 |
// Radix-2 Stockham autosort FFT: one transform per workgroup, entirely in
|
| 102 |
// shared memory. log2(N) ping-pong stages alternate between the two halves of
|
| 103 |
// a 2*N scratch pair (no bit reversal; the autosort write pattern leaves the
|
| 104 |
// final stage in natural order). Threads stride over the N/2 butterflies of
|
| 105 |
+
// each stage with a workgroupBarrier between stages.
|
| 106 |
+
{% if USE_TWIDDLE_TABLE %}
|
| 107 |
+
// This specialization reads roots from a compact precomputed twiddle table.
|
| 108 |
+
{% else %}
|
| 109 |
+
// This specialization evaluates roots inline.
|
| 110 |
+
{% endif %}
|
| 111 |
+
// One workgroup owns each (batch, frame). The window multiplication is fused
|
| 112 |
+
// into an in-bounds frame load, and the output stores numBins complex values.
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 113 |
|
| 114 |
+
const N: u32 = {{ fftLength }}u;
|
| 115 |
const N_HALF: u32 = N / 2u;
|
| 116 |
{% if MIXED_TAIL %}
|
| 117 |
+
const RADIX2_N: u32 = {{ radix2Length }}u;
|
| 118 |
const RADIX2_HALF: u32 = RADIX2_N / 2u;
|
| 119 |
const TAIL_RADIX: u32 = N / RADIX2_N;
|
| 120 |
{% if REANCHOR_TAIL %}
|
| 121 |
const TAIL_ANCHOR: u32 = {{ TAIL_ANCHOR }}u;
|
| 122 |
{% endif %}
|
| 123 |
{% endif %}
|
| 124 |
+
const NUM_BINS: u32 = {{ numBins }}u;
|
| 125 |
const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
|
| 126 |
{% if NEED_PRECISE_TWIDDLE %}
|
| 127 |
// Twiddle = exp(SIGN * i * 2*pi * k / span); SIGN flips for the inverse
|
| 128 |
+
// transform.
|
| 129 |
+
const SIGN: f32 = {{ "1.0" if inverse else "-1.0" }};
|
| 130 |
{% endif %}
|
| 131 |
|
| 132 |
{% if MIXED_TAIL %}
|
| 133 |
+
// Walk consecutive powers of one unit root in each mixed-radix tail. One
|
| 134 |
+
// precise root evaluation per output bin anchors the recurrence.
|
|
|
|
| 135 |
fn complex_mul(a: vec2<f32>, b: vec2<f32>) -> vec2<f32> {
|
| 136 |
return vec2<f32>(a.x * b.x - a.y * b.y, a.x * b.y + a.y * b.x);
|
| 137 |
}
|
|
|
|
| 146 |
@builtin(local_invocation_id) lid: vec3<u32>,
|
| 147 |
@builtin(num_workgroups) nwg: vec3<u32>) {
|
| 148 |
let tid = lid.x;
|
| 149 |
+
// Flat workgroup index for a grid folded across x and y. This reduces to
|
| 150 |
+
// wgid.x when the dispatch uses one row.
|
|
|
|
|
|
|
|
|
|
| 151 |
let flat = wgid.x + wgid.y * nwg.x;
|
| 152 |
|
| 153 |
let batch = flat / params.numFrames;
|
|
|
|
| 162 |
let sample_n = n;
|
| 163 |
{% endif %}
|
| 164 |
let signal_index = start + sample_n;
|
| 165 |
+
let base = (batch * params.signalLength + signal_index) * {{ inComp }}u;
|
| 166 |
+
{% if hasWindow is defined and not hasWindow %}
|
| 167 |
let w = 1.0;
|
| 168 |
{% else %}
|
| 169 |
let w = window[sample_n];
|
| 170 |
{% endif %}
|
| 171 |
let xr = signal[base] * w;
|
| 172 |
+
{% if inComp == 2 %}
|
| 173 |
let xi = signal[base + 1u] * w;
|
| 174 |
{% else %}
|
| 175 |
let xi = 0.0;
|
|
|
|
| 183 |
// applies w = exp(SIGN*i*pi*(j mod ns)/ns) to b, and writes
|
| 184 |
// dst[(j div ns)*2*ns + (j mod ns)] = a + w*b
|
| 185 |
// dst[(j div ns)*2*ns + (j mod ns) + ns] = a - w*b
|
| 186 |
+
// |angle| <= pi, which bounds the argument supplied to sin and cos.
|
|
|
|
| 187 |
var srcOff = 0u;
|
| 188 |
var dstOff = N;
|
| 189 |
{% if MIXED_TAIL %}
|
|
|
|
| 243 |
}
|
| 244 |
|
| 245 |
// The final swap leaves the natural-order result at srcOff. The 1/N inverse
|
| 246 |
+
// scale is an exactly representable power-of-two reciprocal applied once.
|
|
|
|
| 247 |
let outRow = (batch * params.numFrames + frame) * NUM_BINS;
|
| 248 |
for (var k = tid; k < NUM_BINS; k = k + WG) {
|
| 249 |
let outBase = (outRow + k) * 2u;
|
build/webgpu/fft-stockham-twiddles.wgsl.jinja
CHANGED
|
@@ -1,17 +1,16 @@
|
|
| 1 |
{{ env.wgsl.resourceDeclarations }}
|
| 2 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 3 |
-
// rational turn and route
|
| 4 |
-
//
|
| 5 |
//
|
| 6 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 7 |
// inverse transform. The result is (cos, sin).
|
| 8 |
// Backend-stable f32 sine/cosine core.
|
| 9 |
//
|
| 10 |
-
// Shader transcendental accuracy is implementation-defined
|
| 11 |
-
//
|
| 12 |
-
//
|
| 13 |
-
//
|
| 14 |
-
// truncation error is well below one f32 ULP over its documented interval.
|
| 15 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 16 |
const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
|
| 17 |
const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
|
|
@@ -86,9 +85,9 @@ fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
|
|
| 86 |
// each root once avoids repeating sin/cos for every frame while preserving the
|
| 87 |
// same f32 twiddle values and butterfly arithmetic.
|
| 88 |
// STFT is forward-only, so the transform sign is fixed at -1.
|
| 89 |
-
const N: u32 = {{
|
| 90 |
const SIGN: f32 = -1.0;
|
| 91 |
-
{% if
|
| 92 |
const PACKED_N: u32 = N / 2u;
|
| 93 |
const STAGE_TWIDDLES: u32 = PACKED_N - 1u;
|
| 94 |
|
|
@@ -96,7 +95,7 @@ const STAGE_TWIDDLES: u32 = PACKED_N - 1u;
|
|
| 96 |
@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
|
| 97 |
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 98 |
let tableIndex = gid.x;
|
| 99 |
-
{% if
|
| 100 |
if (tableIndex >= N) {
|
| 101 |
return;
|
| 102 |
}
|
|
|
|
| 1 |
{{ env.wgsl.resourceDeclarations }}
|
| 2 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 3 |
+
// rational turn and route polynomial evaluation through one trigonometric core
|
| 4 |
+
// for FFT, direct DFT, STFT, and cosine-window calculations.
|
| 5 |
//
|
| 6 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 7 |
// inverse transform. The result is (cos, sin).
|
| 8 |
// Backend-stable f32 sine/cosine core.
|
| 9 |
//
|
| 10 |
+
// Shader transcendental accuracy is implementation-defined. Each path retains
|
| 11 |
+
// its available phase representation, reduces it to [-pi, pi], then evaluates
|
| 12 |
+
// explicit polynomials in a fixed order. The half-pi core is degree 13 for sine
|
| 13 |
+
// and degree 12 for cosine.
|
|
|
|
| 14 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 15 |
const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
|
| 16 |
const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
|
|
|
|
| 85 |
// each root once avoids repeating sin/cos for every frame while preserving the
|
| 86 |
// same f32 twiddle values and butterfly arithmetic.
|
| 87 |
// STFT is forward-only, so the transform sign is fixed at -1.
|
| 88 |
+
const N: u32 = {{ fftLength }}u;
|
| 89 |
const SIGN: f32 = -1.0;
|
| 90 |
+
{% if realPacked %}
|
| 91 |
const PACKED_N: u32 = N / 2u;
|
| 92 |
const STAGE_TWIDDLES: u32 = PACKED_N - 1u;
|
| 93 |
|
|
|
|
| 95 |
@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
|
| 96 |
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 97 |
let tableIndex = gid.x;
|
| 98 |
+
{% if realPacked %}
|
| 99 |
if (tableIndex >= N) {
|
| 100 |
return;
|
| 101 |
}
|
build/webgpu/manifest.json
CHANGED
|
@@ -2,61 +2,27 @@
|
|
| 2 |
"domain": "ai.onnx",
|
| 3 |
"name": "STFT",
|
| 4 |
"sinceVersion": 17,
|
| 5 |
-
"
|
| 6 |
-
"
|
| 7 |
-
{
|
| 8 |
-
"role": "signal",
|
| 9 |
-
"dtype": "T",
|
| 10 |
-
"rank": 3,
|
| 11 |
-
"description": "Real input signal with shape `(batch_size, signal_length, 1)`, or complex input with shape `(batch_size, signal_length, 2)`."
|
| 12 |
-
},
|
| 13 |
-
{
|
| 14 |
-
"role": "window",
|
| 15 |
-
"dtype": "T",
|
| 16 |
-
"rank": 1,
|
| 17 |
-
"optional": true,
|
| 18 |
-
"description": "Optional 1-D window tensor with shape `(window_length)` to multiply by each frame before the DFT."
|
| 19 |
-
}
|
| 20 |
-
],
|
| 21 |
-
"outputs": [
|
| 22 |
-
{
|
| 23 |
-
"role": "output",
|
| 24 |
"dtype": "T",
|
| 25 |
"rank": 4,
|
| 26 |
-
"description": "STFT result with shape `(batch_size, frames, dft_unique_bins, 2)` in one-sided mode or `(batch_size, frames, frame_length, 2)` in two-sided mode, with real and imaginary parts in the last dimension. `frames = floor((signal_length - frame_length) / frame_step) + 1`; partial trailing frames are not emitted.",
|
| 27 |
"shape": ["dim(shapes.signal, 0)", "expectedFrames", "floor(frameLength / 2) + 1 if attrs.onesided == 1 else frameLength", 2]
|
| 28 |
}
|
| 29 |
-
],
|
| 30 |
-
"attributes": { "onesided": 1 },
|
| 31 |
-
"attributeDescriptions": {
|
| 32 |
-
"onesided": "When 1 (default), returns `floor(n_fft / 2) + 1` non-redundant frequency bins with indices `[0, ..., floor(n_fft / 2)]` for a real-valued input (RFFT); set to 0 to return the full spectrum."
|
| 33 |
},
|
| 34 |
-
"attributeConstraints": { "onesided": { "values": [0, 1] } },
|
| 35 |
-
"typeConstraints": { "T": ["float32"] },
|
| 36 |
"args": {
|
| 37 |
-
"
|
| 38 |
-
"
|
| 39 |
-
"output": { "kind": "tensor", "semantic": "output", "role": "output" },
|
| 40 |
-
"frameStep": {
|
| 41 |
-
"kind": "i32",
|
| 42 |
-
"semantic": "frame_step",
|
| 43 |
-
"role": "attribute",
|
| 44 |
-
"description": "Positive number of signal positions between the starts of consecutive frames."
|
| 45 |
-
},
|
| 46 |
-
"frameLength": {
|
| 47 |
-
"kind": "i32",
|
| 48 |
-
"semantic": "frame_length",
|
| 49 |
-
"role": "attribute",
|
| 50 |
-
"required": false,
|
| 51 |
-
"description": "Optional number of signal values in each frame; must equal the window length when both are supplied."
|
| 52 |
-
}
|
| 53 |
},
|
|
|
|
|
|
|
|
|
|
| 54 |
"tunables": {
|
| 55 |
-
"WORKGROUP_SIZE": 128,
|
| 56 |
-
"TWIDDLE_TABLE_MIN_TRANSFORMS": 4096,
|
| 57 |
-
"TILED_REAL_MIN_LENGTH": 512,
|
| 58 |
-
"GLOBAL_STAGE_LOCAL_LENGTH": 256,
|
| 59 |
-
"GLOBAL_STAGE_MAX_DEVICE_STORAGE_BYTES": 16384
|
| 60 |
},
|
| 61 |
"derive": {
|
| 62 |
"windowLength": "dim(shapes.window, 0) if (present.window and ranks.window == 1) else 0",
|
|
@@ -67,7 +33,7 @@
|
|
| 67 |
"rank3OutputLayout": "ranks.signal == 3 and ranks.output == 4 and dim(shapes.output, 0) == dim(shapes.signal, 0) and dim(shapes.output, 1) == expectedFrames and dim(shapes.signal, 2) >= 1 and dim(shapes.signal, 2) <= 2 and dim(shapes.output, 3) == 2",
|
| 68 |
"rank3WindowLayout": "present.window and ranks.window == 1 and rank3OutputLayout and frameControlsValid",
|
| 69 |
"rank3WindowSpectrum": "false if not present.window else ((attrs.onesided == 0 and dim(shapes.output, 2) == dim(shapes.window, 0)) or (attrs.onesided != 0 and dim(shapes.signal, 2) == 1 and dim(shapes.output, 2) == floor(dim(shapes.window, 0) / 2) + 1))",
|
| 70 |
-
"rank3WindowDispatchFits": "dim(shapes.output, 1) <= device.limits.maxComputeWorkgroupsPerDimension and dim(shapes.output, 0) <= device.limits.maxComputeWorkgroupsPerDimension",
|
| 71 |
"rank3WindowSharedFftFits": "false if not present.window else (rank3WindowDispatchFits and pow2ceil(dim(shapes.window, 0)) == dim(shapes.window, 0) and dim(shapes.window, 0) >= 4 and dim(shapes.window, 0) * 16 <= device.limits.maxComputeWorkgroupStorageSize)",
|
| 72 |
"rank3WindowFftRadix2": "1 if not present.window else pick([[dim(shapes.window, 0) % 256 == 0, 256], [dim(shapes.window, 0) % 128 == 0, 128], [dim(shapes.window, 0) % 64 == 0, 64], [dim(shapes.window, 0) % 32 == 0, 32], [dim(shapes.window, 0) % 16 == 0, 16], [dim(shapes.window, 0) % 8 == 0, 8], [dim(shapes.window, 0) % 4 == 0, 4]], 1)",
|
| 73 |
"rank3WindowFftTailRadix": "1 if not present.window else dim(shapes.window, 0) / rank3WindowFftRadix2",
|
|
@@ -76,235 +42,54 @@
|
|
| 76 |
"rank3WindowRadix2ScratchFits": "false if not present.window else (rank3WindowDispatchFits and pow2ceil(dim(shapes.window, 0)) == dim(shapes.window, 0) and dim(shapes.window, 0) >= 4 and dim(shapes.window, 0) * 16 > device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 0) * dim(shapes.output, 1) > 0 and dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) * 2 * 4 <= device.limits.maxStorageBufferBindingSize and dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) * 2 * 4 <= device.limits.maxBufferSize)",
|
| 77 |
"rank3WindowInplaceFftFits": "false if not present.window else (rank3WindowDispatchFits and pow2ceil(dim(shapes.window, 0)) == dim(shapes.window, 0) and dim(shapes.window, 0) >= 4 and dim(shapes.window, 0) * 16 > device.limits.maxComputeWorkgroupStorageSize and dim(shapes.window, 0) * 8 <= device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 0) * dim(shapes.output, 1) > 0)",
|
| 78 |
"rank3WindowGlobalFftFits": "false if not present.window else (rank3WindowRadix2ScratchFits and dim(shapes.window, 0) >= 2 * tunables.GLOBAL_STAGE_LOCAL_LENGTH and dim(shapes.window, 0) * 8 > device.limits.maxComputeWorkgroupStorageSize)",
|
| 79 |
-
"rank3WindowTiledRealFits": "false if not present.window else (rank3WindowDispatchFits and ranks.signal == 3 and dim(shapes.signal, 2) == 1 and dim(shapes.window, 0) >= tunables.TILED_REAL_MIN_LENGTH and dim(shapes.window, 0) * 4 <= device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 0) > 0 and dim(shapes.output, 1) > 0 and ceilDiv(dim(shapes.output, 2), 256) <= device.limits.maxComputeWorkgroupsPerDimension)",
|
| 80 |
"packedFourWayFits": "false if not present.window else (device.limits.maxComputeInvocationsPerWorkgroup >= 256 and device.limits.maxComputeWorkgroupSizeX >= 256 and dim(shapes.window, 0) * 32 <= device.limits.maxComputeWorkgroupStorageSize)",
|
| 81 |
"packedTransformsPerWorkgroup": "4 if packedFourWayFits else 1",
|
| 82 |
"rank3NoWindowLayout": "not present.window and args.frameLength is defined and rank3OutputLayout and frameControlsValid",
|
| 83 |
"rank3NoWindowSpectrum": "(attrs.onesided == 0 and dim(shapes.output, 2) == frameLength) or (attrs.onesided != 0 and dim(shapes.signal, 2) == 1 and dim(shapes.output, 2) == floor(frameLength / 2) + 1)",
|
| 84 |
-
"rank3NoWindowSharedFftFits": "false if ranks.signal != 3 else ((attrs.onesided == 0 or (dim(shapes.signal, 2) == 1 and dim(shapes.output, 2) >= 3)) and pow2ceil(frameLength) == frameLength and frameLength >= 4 and frameLength * 16 <= device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 1) <= device.limits.maxComputeWorkgroupsPerDimension and dim(shapes.output, 0) <= device.limits.maxComputeWorkgroupsPerDimension)"
|
| 85 |
},
|
| 86 |
-
"
|
| 87 |
-
"
|
| 88 |
-
|
| 89 |
-
|
| 90 |
-
|
| 91 |
-
|
| 92 |
-
"
|
| 93 |
-
"
|
| 94 |
-
|
| 95 |
-
|
| 96 |
-
|
| 97 |
-
|
| 98 |
-
|
| 99 |
-
|
| 100 |
-
|
| 101 |
-
|
| 102 |
-
|
| 103 |
-
|
| 104 |
-
|
| 105 |
-
"
|
| 106 |
-
"
|
| 107 |
-
"
|
| 108 |
-
|
| 109 |
-
|
| 110 |
-
|
| 111 |
-
|
| 112 |
-
|
| 113 |
-
|
| 114 |
-
|
| 115 |
-
|
| 116 |
-
|
| 117 |
-
|
| 118 |
-
|
| 119 |
-
|
| 120 |
-
|
| 121 |
-
|
| 122 |
-
"name": "
|
| 123 |
-
"
|
| 124 |
-
"
|
| 125 |
-
|
| 126 |
-
|
| 127 |
-
},
|
| 128 |
-
{
|
| 129 |
-
"name": "window",
|
| 130 |
-
"arg": "window",
|
| 131 |
-
"semantic": "window",
|
| 132 |
-
"buffer": { "type": "read-only-storage" },
|
| 133 |
-
"elementType": "f32"
|
| 134 |
-
},
|
| 135 |
-
{
|
| 136 |
-
"name": "twiddles",
|
| 137 |
-
"semantic": "fft_twiddles",
|
| 138 |
-
"buffer": { "type": "read-only-storage" },
|
| 139 |
-
"elementType": "f32"
|
| 140 |
-
},
|
| 141 |
-
{ "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
|
| 142 |
-
{
|
| 143 |
-
"name": "params",
|
| 144 |
-
"semantic": "kernel.params",
|
| 145 |
-
"buffer": { "type": "uniform" },
|
| 146 |
-
"struct": {
|
| 147 |
-
"name": "Params",
|
| 148 |
-
"fields": [
|
| 149 |
-
{ "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
|
| 150 |
-
{ "name": "numFrames", "type": "u32", "value": "dim(shapes.output, 1)" },
|
| 151 |
-
{ "name": "frameStep", "type": "u32", "value": "args.frameStep" }
|
| 152 |
-
]
|
| 153 |
-
}
|
| 154 |
-
}
|
| 155 |
-
],
|
| 156 |
-
"stftNoWindowFft": [
|
| 157 |
-
{
|
| 158 |
-
"name": "signal",
|
| 159 |
-
"arg": "signal",
|
| 160 |
-
"semantic": "signal",
|
| 161 |
-
"buffer": { "type": "read-only-storage" },
|
| 162 |
-
"elementType": "f32"
|
| 163 |
-
},
|
| 164 |
-
{ "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
|
| 165 |
-
{
|
| 166 |
-
"name": "params",
|
| 167 |
-
"semantic": "kernel.params",
|
| 168 |
-
"buffer": { "type": "uniform" },
|
| 169 |
-
"struct": {
|
| 170 |
-
"name": "Params",
|
| 171 |
-
"fields": [
|
| 172 |
-
{ "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
|
| 173 |
-
{ "name": "numFrames", "type": "u32", "value": "dim(shapes.output, 1)" },
|
| 174 |
-
{ "name": "frameStep", "type": "u32", "value": "args.frameStep" }
|
| 175 |
-
]
|
| 176 |
-
}
|
| 177 |
-
}
|
| 178 |
-
],
|
| 179 |
-
"stftWindowScratch": [
|
| 180 |
-
{
|
| 181 |
-
"name": "signal",
|
| 182 |
-
"arg": "signal",
|
| 183 |
-
"semantic": "signal",
|
| 184 |
-
"buffer": { "type": "read-only-storage" },
|
| 185 |
-
"elementType": "f32"
|
| 186 |
-
},
|
| 187 |
-
{
|
| 188 |
-
"name": "window",
|
| 189 |
-
"arg": "window",
|
| 190 |
-
"semantic": "window",
|
| 191 |
-
"buffer": { "type": "read-only-storage" },
|
| 192 |
-
"elementType": "f32"
|
| 193 |
-
},
|
| 194 |
-
{ "name": "scratch", "semantic": "fft_scratch", "buffer": { "type": "storage" }, "elementType": "f32" },
|
| 195 |
-
{ "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
|
| 196 |
-
{
|
| 197 |
-
"name": "params",
|
| 198 |
-
"semantic": "kernel.params",
|
| 199 |
-
"buffer": { "type": "uniform" },
|
| 200 |
-
"struct": {
|
| 201 |
-
"name": "Params",
|
| 202 |
-
"fields": [
|
| 203 |
-
{ "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
|
| 204 |
-
{ "name": "numFrames", "type": "u32", "value": "dim(shapes.output, 1)" },
|
| 205 |
-
{ "name": "frameStep", "type": "u32", "value": "args.frameStep" }
|
| 206 |
-
]
|
| 207 |
-
}
|
| 208 |
-
}
|
| 209 |
-
],
|
| 210 |
-
"stftWindowScratchLocalPrefix": [
|
| 211 |
-
{
|
| 212 |
-
"name": "signal",
|
| 213 |
-
"arg": "signal",
|
| 214 |
-
"semantic": "signal",
|
| 215 |
-
"buffer": { "type": "read-only-storage" },
|
| 216 |
-
"elementType": "f32"
|
| 217 |
-
},
|
| 218 |
-
{
|
| 219 |
-
"name": "window",
|
| 220 |
-
"arg": "window",
|
| 221 |
-
"semantic": "window",
|
| 222 |
-
"buffer": { "type": "read-only-storage" },
|
| 223 |
-
"elementType": "f32"
|
| 224 |
-
},
|
| 225 |
-
{ "name": "scratch", "semantic": "fft_scratch", "buffer": { "type": "storage" }, "elementType": "f32" },
|
| 226 |
-
{
|
| 227 |
-
"name": "params",
|
| 228 |
-
"semantic": "kernel.params",
|
| 229 |
-
"buffer": { "type": "uniform" },
|
| 230 |
-
"struct": {
|
| 231 |
-
"name": "Params",
|
| 232 |
-
"fields": [
|
| 233 |
-
{ "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
|
| 234 |
-
{ "name": "numFrames", "type": "u32", "value": "dim(shapes.output, 1)" },
|
| 235 |
-
{ "name": "frameStep", "type": "u32", "value": "args.frameStep" }
|
| 236 |
-
]
|
| 237 |
-
}
|
| 238 |
-
}
|
| 239 |
-
],
|
| 240 |
-
"stftWindowScratchStage": [
|
| 241 |
-
{ "name": "scratch", "semantic": "fft_scratch", "buffer": { "type": "storage" }, "elementType": "f32" }
|
| 242 |
-
],
|
| 243 |
-
"stftWindowScratchFinish": [
|
| 244 |
-
{ "name": "scratch", "semantic": "fft_scratch", "buffer": { "type": "read-only-storage" }, "elementType": "f32" },
|
| 245 |
-
{ "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" }
|
| 246 |
-
],
|
| 247 |
-
"stftWindowGeneral": [
|
| 248 |
-
{
|
| 249 |
-
"name": "signal",
|
| 250 |
-
"arg": "signal",
|
| 251 |
-
"semantic": "signal",
|
| 252 |
-
"buffer": { "type": "read-only-storage" },
|
| 253 |
-
"elementType": "f32"
|
| 254 |
-
},
|
| 255 |
-
{
|
| 256 |
-
"name": "window",
|
| 257 |
-
"arg": "window",
|
| 258 |
-
"semantic": "window",
|
| 259 |
-
"buffer": { "type": "read-only-storage" },
|
| 260 |
-
"elementType": "f32"
|
| 261 |
-
},
|
| 262 |
-
{ "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
|
| 263 |
-
{
|
| 264 |
-
"name": "params",
|
| 265 |
-
"semantic": "kernel.params",
|
| 266 |
-
"buffer": { "type": "uniform" },
|
| 267 |
-
"struct": {
|
| 268 |
-
"name": "Params",
|
| 269 |
-
"fields": [
|
| 270 |
-
{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 2" },
|
| 271 |
-
{ "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
|
| 272 |
-
{ "name": "signalComponents", "type": "u32", "value": "dim(shapes.signal, 2)" },
|
| 273 |
-
{ "name": "frameLength", "type": "u32", "value": "dim(shapes.window, 0)" },
|
| 274 |
-
{ "name": "frames", "type": "u32", "value": "dim(shapes.output, 1)" },
|
| 275 |
-
{ "name": "bins", "type": "u32", "value": "dim(shapes.output, 2)" },
|
| 276 |
-
{ "name": "frameStep", "type": "u32", "value": "args.frameStep" }
|
| 277 |
-
]
|
| 278 |
-
}
|
| 279 |
-
}
|
| 280 |
-
],
|
| 281 |
-
"stftNoWindowGeneral": [
|
| 282 |
-
{
|
| 283 |
-
"name": "signal",
|
| 284 |
-
"arg": "signal",
|
| 285 |
-
"semantic": "signal",
|
| 286 |
-
"buffer": { "type": "read-only-storage" },
|
| 287 |
-
"elementType": "f32"
|
| 288 |
-
},
|
| 289 |
-
{ "name": "output", "arg": "output", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "f32" },
|
| 290 |
-
{
|
| 291 |
-
"name": "params",
|
| 292 |
-
"semantic": "kernel.params",
|
| 293 |
-
"buffer": { "type": "uniform" },
|
| 294 |
-
"struct": {
|
| 295 |
-
"name": "Params",
|
| 296 |
-
"fields": [
|
| 297 |
-
{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 2" },
|
| 298 |
-
{ "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
|
| 299 |
-
{ "name": "signalComponents", "type": "u32", "value": "dim(shapes.signal, 2)" },
|
| 300 |
-
{ "name": "frameLength", "type": "u32", "value": "args.frameLength" },
|
| 301 |
-
{ "name": "frames", "type": "u32", "value": "dim(shapes.output, 1)" },
|
| 302 |
-
{ "name": "bins", "type": "u32", "value": "dim(shapes.output, 2)" },
|
| 303 |
-
{ "name": "frameStep", "type": "u32", "value": "args.frameStep" }
|
| 304 |
-
]
|
| 305 |
-
}
|
| 306 |
-
}
|
| 307 |
-
]
|
| 308 |
},
|
| 309 |
"variants": [
|
| 310 |
{
|
|
@@ -315,18 +100,16 @@
|
|
| 315 |
{
|
| 316 |
"id": "main",
|
| 317 |
"name": "STFT.NoWindowFftShared",
|
| 318 |
-
"
|
| 319 |
-
|
| 320 |
-
"
|
| 321 |
-
|
| 322 |
-
|
| 323 |
-
|
| 324 |
-
|
| 325 |
-
|
| 326 |
-
"hasWindow": false
|
| 327 |
-
}
|
| 328 |
},
|
| 329 |
-
"bindings": "
|
| 330 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 331 |
}
|
| 332 |
]
|
|
@@ -334,23 +117,21 @@
|
|
| 334 |
{
|
| 335 |
"id": "rank3_window_fft_shared_inplace",
|
| 336 |
"priority": 32,
|
| 337 |
-
"requires": { "limits": { "maxComputeInvocationsPerWorkgroup": 256, "maxComputeWorkgroupSizeX": 256 } },
|
| 338 |
-
"tunables": { "WORKGROUP_SIZE": 256 },
|
| 339 |
"when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowInplaceFftFits"],
|
|
|
|
|
|
|
| 340 |
"passes": [
|
| 341 |
{
|
| 342 |
"id": "main",
|
| 343 |
"name": "STFT.FftSharedInplace",
|
| 344 |
-
"
|
| 345 |
-
|
| 346 |
-
"
|
| 347 |
-
|
| 348 |
-
|
| 349 |
-
|
| 350 |
-
"inComp": "dim(shapes.signal, 2)"
|
| 351 |
-
}
|
| 352 |
},
|
| 353 |
-
"bindings": "
|
| 354 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 355 |
}
|
| 356 |
]
|
|
@@ -358,10 +139,10 @@
|
|
| 358 |
{
|
| 359 |
"id": "rank3_window_fft_global_stages",
|
| 360 |
"priority": 31,
|
| 361 |
-
"requires": { "limits": { "maxComputeInvocationsPerWorkgroup": 256, "maxComputeWorkgroupSizeX": 256 } },
|
| 362 |
-
"tunables": { "WORKGROUP_SIZE": 256 },
|
| 363 |
-
"demoteWhen": ["device.limits.maxComputeWorkgroupStorageSize > tunables.GLOBAL_STAGE_MAX_DEVICE_STORAGE_BYTES"],
|
| 364 |
"when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowGlobalFftFits"],
|
|
|
|
|
|
|
|
|
|
| 365 |
"intermediates": [
|
| 366 |
{
|
| 367 |
"id": "fft_scratch",
|
|
@@ -373,21 +154,21 @@
|
|
| 373 |
{
|
| 374 |
"id": "local_prefix",
|
| 375 |
"name": "STFT.FftLocalPrefix",
|
| 376 |
-
"
|
| 377 |
-
|
| 378 |
-
"
|
| 379 |
-
|
| 380 |
-
|
| 381 |
-
|
| 382 |
-
|
| 383 |
-
|
| 384 |
-
|
| 385 |
-
"inComp": "dim(shapes.signal, 2)"
|
| 386 |
-
}
|
| 387 |
},
|
| 388 |
-
"bindings": "
|
| 389 |
"dispatch": {
|
| 390 |
-
"
|
|
|
|
|
|
|
| 391 |
}
|
| 392 |
},
|
| 393 |
{
|
|
@@ -397,21 +178,20 @@
|
|
| 397 |
{
|
| 398 |
"id": "butterflies",
|
| 399 |
"name": "STFT.FftGlobalStage",
|
| 400 |
-
"
|
| 401 |
-
|
| 402 |
-
"
|
| 403 |
-
|
| 404 |
-
|
| 405 |
-
|
| 406 |
-
|
| 407 |
-
|
| 408 |
-
"inComp": "dim(shapes.signal, 2)"
|
| 409 |
-
}
|
| 410 |
},
|
| 411 |
-
"bindings": "
|
| 412 |
"dispatch": {
|
| 413 |
-
"
|
| 414 |
-
"
|
|
|
|
| 415 |
}
|
| 416 |
}
|
| 417 |
]
|
|
@@ -419,20 +199,22 @@
|
|
| 419 |
{
|
| 420 |
"id": "finish",
|
| 421 |
"name": "STFT.FftGlobalFinish",
|
| 422 |
-
"
|
| 423 |
-
|
| 424 |
-
"
|
| 425 |
-
|
| 426 |
-
|
| 427 |
-
|
| 428 |
-
|
| 429 |
-
"inComp": "dim(shapes.signal, 2)"
|
| 430 |
-
}
|
| 431 |
},
|
| 432 |
-
"bindings":
|
|
|
|
|
|
|
|
|
|
| 433 |
"dispatch": {
|
| 434 |
-
"
|
| 435 |
-
"
|
|
|
|
| 436 |
}
|
| 437 |
}
|
| 438 |
]
|
|
@@ -452,18 +234,16 @@
|
|
| 452 |
{
|
| 453 |
"id": "main",
|
| 454 |
"name": "STFT.FftStorageInplace",
|
| 455 |
-
"
|
| 456 |
-
|
| 457 |
-
"
|
| 458 |
-
|
| 459 |
-
|
| 460 |
-
|
| 461 |
-
|
| 462 |
-
|
| 463 |
-
"inverse": false
|
| 464 |
-
}
|
| 465 |
},
|
| 466 |
-
"bindings": "
|
| 467 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 468 |
}
|
| 469 |
]
|
|
@@ -471,23 +251,21 @@
|
|
| 471 |
{
|
| 472 |
"id": "rank3_window_tiled_real",
|
| 473 |
"priority": 19,
|
| 474 |
-
"requires": { "limits": { "maxComputeInvocationsPerWorkgroup": 256, "maxComputeWorkgroupSizeX": 256 } },
|
| 475 |
-
"tunables": { "WORKGROUP_SIZE": 256 },
|
| 476 |
"when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowTiledRealFits"],
|
|
|
|
|
|
|
| 477 |
"passes": [
|
| 478 |
{
|
| 479 |
"id": "main",
|
| 480 |
"name": "STFT.TiledReal",
|
| 481 |
-
"
|
| 482 |
-
|
| 483 |
-
"
|
| 484 |
-
|
| 485 |
-
|
| 486 |
-
|
| 487 |
-
"workgroupSize": "tunables.WORKGROUP_SIZE"
|
| 488 |
-
}
|
| 489 |
},
|
| 490 |
-
"bindings": "
|
| 491 |
"dispatch": {
|
| 492 |
"x": "ceilDiv(dim(shapes.output, 2), 256)",
|
| 493 |
"y": "dim(shapes.output, 1)",
|
|
@@ -499,28 +277,28 @@
|
|
| 499 |
{
|
| 500 |
"id": "rank3_window_rfft_packed",
|
| 501 |
"priority": 22,
|
| 502 |
-
"tunables": { "WORKGROUP_SIZE": 64 },
|
| 503 |
"when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowSharedFftFits", "attrs.onesided != 0", "dim(shapes.signal, 2) == 1", "dim(shapes.window, 0) >= 64", "dim(shapes.output, 0) * dim(shapes.output, 1) >= tunables.TWIDDLE_TABLE_MIN_TRANSFORMS"],
|
|
|
|
| 504 |
"intermediates": [{ "id": "fft_twiddles", "dtype": "float32", "shape": "[2 * dim(shapes.window, 0)]" }],
|
| 505 |
"passes": [
|
| 506 |
{
|
| 507 |
"id": "twiddles",
|
| 508 |
"name": "STFT.RealPackedTwiddleTable",
|
| 509 |
-
"
|
| 510 |
-
|
| 511 |
-
|
| 512 |
-
|
| 513 |
-
|
| 514 |
-
|
|
|
|
|
|
|
| 515 |
},
|
| 516 |
{
|
| 517 |
"id": "main",
|
| 518 |
"name": "STFT.RealPackedFftShared",
|
| 519 |
-
"
|
| 520 |
-
|
| 521 |
-
|
| 522 |
-
},
|
| 523 |
-
"bindings": "stftWindowFftTwiddles",
|
| 524 |
"dispatch": {
|
| 525 |
"x": "ceilDiv(dim(shapes.output, 1), packedTransformsPerWorkgroup)",
|
| 526 |
"y": "dim(shapes.output, 0)"
|
|
@@ -537,28 +315,28 @@
|
|
| 537 |
{
|
| 538 |
"id": "twiddles",
|
| 539 |
"name": "STFT.FftTwiddleTable",
|
| 540 |
-
"
|
| 541 |
-
|
| 542 |
-
|
| 543 |
-
|
| 544 |
-
|
| 545 |
-
|
|
|
|
|
|
|
| 546 |
},
|
| 547 |
{
|
| 548 |
"id": "main",
|
| 549 |
"name": "STFT.FftSharedTwiddles",
|
| 550 |
-
"
|
| 551 |
-
|
| 552 |
-
"
|
| 553 |
-
|
| 554 |
-
|
| 555 |
-
|
| 556 |
-
|
| 557 |
-
|
| 558 |
-
"useTwiddleTable": true
|
| 559 |
-
}
|
| 560 |
},
|
| 561 |
-
"bindings": "
|
| 562 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 563 |
}
|
| 564 |
]
|
|
@@ -571,19 +349,17 @@
|
|
| 571 |
{
|
| 572 |
"id": "main",
|
| 573 |
"name": "STFT.FftMixedRecurrent",
|
| 574 |
-
"
|
| 575 |
-
|
| 576 |
-
"
|
| 577 |
-
|
| 578 |
-
|
| 579 |
-
|
| 580 |
-
|
| 581 |
-
|
| 582 |
-
|
| 583 |
-
"inverse": false
|
| 584 |
-
}
|
| 585 |
},
|
| 586 |
-
"bindings": "
|
| 587 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 588 |
}
|
| 589 |
]
|
|
@@ -596,17 +372,15 @@
|
|
| 596 |
{
|
| 597 |
"id": "main",
|
| 598 |
"name": "STFT.FftShared",
|
| 599 |
-
"
|
| 600 |
-
|
| 601 |
-
"
|
| 602 |
-
|
| 603 |
-
|
| 604 |
-
|
| 605 |
-
|
| 606 |
-
"inverse": false
|
| 607 |
-
}
|
| 608 |
},
|
| 609 |
-
"bindings": "
|
| 610 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 611 |
}
|
| 612 |
]
|
|
@@ -614,15 +388,20 @@
|
|
| 614 |
{
|
| 615 |
"id": "rank3_window",
|
| 616 |
"priority": 0,
|
| 617 |
-
"supersededBy": ["rank3_window_fft_shared_inplace", "rank3_window_fft_global_stages", "rank3_window_fft_storage_inplace", "rank3_window_fft_mixed", "rank3_window_rfft_packed", "rank3_window_fft_twiddles", "rank3_window_fft", "rank3_window_tiled_real"],
|
| 618 |
"when": ["rank3WindowLayout", "rank3WindowSpectrum"],
|
|
|
|
| 619 |
"passes": [
|
| 620 |
{
|
| 621 |
"id": "main",
|
| 622 |
"name": "STFT",
|
| 623 |
-
"
|
| 624 |
-
"
|
| 625 |
-
"
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 626 |
}
|
| 627 |
]
|
| 628 |
},
|
|
@@ -634,9 +413,14 @@
|
|
| 634 |
{
|
| 635 |
"id": "main",
|
| 636 |
"name": "STFT.GeneralNoWindow",
|
| 637 |
-
"
|
| 638 |
-
"
|
| 639 |
-
"
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 640 |
}
|
| 641 |
]
|
| 642 |
}
|
|
|
|
| 2 |
"domain": "ai.onnx",
|
| 3 |
"name": "STFT",
|
| 4 |
"sinceVersion": 17,
|
| 5 |
+
"inputs": { "signal": { "dtype": "T", "rank": 3 }, "window": { "dtype": "T", "rank": 1, "optional": true } },
|
| 6 |
+
"outputs": {
|
| 7 |
+
"output": {
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 8 |
"dtype": "T",
|
| 9 |
"rank": 4,
|
|
|
|
| 10 |
"shape": ["dim(shapes.signal, 0)", "expectedFrames", "floor(frameLength / 2) + 1 if attrs.onesided == 1 else frameLength", 2]
|
| 11 |
}
|
|
|
|
|
|
|
|
|
|
|
|
|
| 12 |
},
|
|
|
|
|
|
|
| 13 |
"args": {
|
| 14 |
+
"frameStep": { "kind": "i32", "onnx": "frame_step" },
|
| 15 |
+
"frameLength": { "kind": "i32", "onnx": "frame_length", "required": false }
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 16 |
},
|
| 17 |
+
"attributes": { "onesided": { "default": 1 } },
|
| 18 |
+
"attributeConstraints": { "onesided": { "values": [0, 1] } },
|
| 19 |
+
"typeConstraints": { "T": ["float32"] },
|
| 20 |
"tunables": {
|
| 21 |
+
"WORKGROUP_SIZE": { "default": 128 },
|
| 22 |
+
"TWIDDLE_TABLE_MIN_TRANSFORMS": { "default": 4096 },
|
| 23 |
+
"TILED_REAL_MIN_LENGTH": { "default": 512 },
|
| 24 |
+
"GLOBAL_STAGE_LOCAL_LENGTH": { "default": 256 },
|
| 25 |
+
"GLOBAL_STAGE_MAX_DEVICE_STORAGE_BYTES": { "default": 16384 }
|
| 26 |
},
|
| 27 |
"derive": {
|
| 28 |
"windowLength": "dim(shapes.window, 0) if (present.window and ranks.window == 1) else 0",
|
|
|
|
| 33 |
"rank3OutputLayout": "ranks.signal == 3 and ranks.output == 4 and dim(shapes.output, 0) == dim(shapes.signal, 0) and dim(shapes.output, 1) == expectedFrames and dim(shapes.signal, 2) >= 1 and dim(shapes.signal, 2) <= 2 and dim(shapes.output, 3) == 2",
|
| 34 |
"rank3WindowLayout": "present.window and ranks.window == 1 and rank3OutputLayout and frameControlsValid",
|
| 35 |
"rank3WindowSpectrum": "false if not present.window else ((attrs.onesided == 0 and dim(shapes.output, 2) == dim(shapes.window, 0)) or (attrs.onesided != 0 and dim(shapes.signal, 2) == 1 and dim(shapes.output, 2) == floor(dim(shapes.window, 0) / 2) + 1))",
|
| 36 |
+
"rank3WindowDispatchFits": "dim(shapes.output, 1) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and dim(shapes.output, 0) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535)",
|
| 37 |
"rank3WindowSharedFftFits": "false if not present.window else (rank3WindowDispatchFits and pow2ceil(dim(shapes.window, 0)) == dim(shapes.window, 0) and dim(shapes.window, 0) >= 4 and dim(shapes.window, 0) * 16 <= device.limits.maxComputeWorkgroupStorageSize)",
|
| 38 |
"rank3WindowFftRadix2": "1 if not present.window else pick([[dim(shapes.window, 0) % 256 == 0, 256], [dim(shapes.window, 0) % 128 == 0, 128], [dim(shapes.window, 0) % 64 == 0, 64], [dim(shapes.window, 0) % 32 == 0, 32], [dim(shapes.window, 0) % 16 == 0, 16], [dim(shapes.window, 0) % 8 == 0, 8], [dim(shapes.window, 0) % 4 == 0, 4]], 1)",
|
| 39 |
"rank3WindowFftTailRadix": "1 if not present.window else dim(shapes.window, 0) / rank3WindowFftRadix2",
|
|
|
|
| 42 |
"rank3WindowRadix2ScratchFits": "false if not present.window else (rank3WindowDispatchFits and pow2ceil(dim(shapes.window, 0)) == dim(shapes.window, 0) and dim(shapes.window, 0) >= 4 and dim(shapes.window, 0) * 16 > device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 0) * dim(shapes.output, 1) > 0 and dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) * 2 * 4 <= device.limits.maxStorageBufferBindingSize and dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) * 2 * 4 <= device.limits.maxBufferSize)",
|
| 43 |
"rank3WindowInplaceFftFits": "false if not present.window else (rank3WindowDispatchFits and pow2ceil(dim(shapes.window, 0)) == dim(shapes.window, 0) and dim(shapes.window, 0) >= 4 and dim(shapes.window, 0) * 16 > device.limits.maxComputeWorkgroupStorageSize and dim(shapes.window, 0) * 8 <= device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 0) * dim(shapes.output, 1) > 0)",
|
| 44 |
"rank3WindowGlobalFftFits": "false if not present.window else (rank3WindowRadix2ScratchFits and dim(shapes.window, 0) >= 2 * tunables.GLOBAL_STAGE_LOCAL_LENGTH and dim(shapes.window, 0) * 8 > device.limits.maxComputeWorkgroupStorageSize)",
|
| 45 |
+
"rank3WindowTiledRealFits": "false if not present.window else (rank3WindowDispatchFits and ranks.signal == 3 and dim(shapes.signal, 2) == 1 and dim(shapes.window, 0) >= tunables.TILED_REAL_MIN_LENGTH and dim(shapes.window, 0) * 4 <= device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 0) > 0 and dim(shapes.output, 1) > 0 and ceilDiv(dim(shapes.output, 2), 256) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
|
| 46 |
"packedFourWayFits": "false if not present.window else (device.limits.maxComputeInvocationsPerWorkgroup >= 256 and device.limits.maxComputeWorkgroupSizeX >= 256 and dim(shapes.window, 0) * 32 <= device.limits.maxComputeWorkgroupStorageSize)",
|
| 47 |
"packedTransformsPerWorkgroup": "4 if packedFourWayFits else 1",
|
| 48 |
"rank3NoWindowLayout": "not present.window and args.frameLength is defined and rank3OutputLayout and frameControlsValid",
|
| 49 |
"rank3NoWindowSpectrum": "(attrs.onesided == 0 and dim(shapes.output, 2) == frameLength) or (attrs.onesided != 0 and dim(shapes.signal, 2) == 1 and dim(shapes.output, 2) == floor(frameLength / 2) + 1)",
|
| 50 |
+
"rank3NoWindowSharedFftFits": "false if ranks.signal != 3 else ((attrs.onesided == 0 or (dim(shapes.signal, 2) == 1 and dim(shapes.output, 2) >= 3)) and pow2ceil(frameLength) == frameLength and frameLength >= 4 and frameLength * 16 <= device.limits.maxComputeWorkgroupStorageSize and dim(shapes.output, 1) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) and dim(shapes.output, 0) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535))"
|
| 51 |
},
|
| 52 |
+
"bindings": {
|
| 53 |
+
"signal": { "buffer": "read-only-storage", "elementType": "f32" },
|
| 54 |
+
"output": { "buffer": "storage", "elementType": "f32" },
|
| 55 |
+
"params": {
|
| 56 |
+
"buffer": "uniform",
|
| 57 |
+
"struct": [
|
| 58 |
+
{ "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
|
| 59 |
+
{ "name": "numFrames", "type": "u32", "value": "dim(shapes.output, 1)" },
|
| 60 |
+
{ "name": "frameStep", "type": "u32" }
|
| 61 |
+
]
|
| 62 |
+
},
|
| 63 |
+
"window": { "buffer": "read-only-storage", "elementType": "f32" },
|
| 64 |
+
"scratch": { "scratch": "fft_scratch", "buffer": "storage", "elementType": "f32" },
|
| 65 |
+
"twiddles": { "scratch": "fft_twiddles", "buffer": "storage", "elementType": "f32" },
|
| 66 |
+
"twiddles_2": { "scratch": "fft_twiddles", "name": "twiddles", "buffer": "read-only-storage", "elementType": "f32" },
|
| 67 |
+
"params_2": {
|
| 68 |
+
"name": "params",
|
| 69 |
+
"buffer": "uniform",
|
| 70 |
+
"struct": [
|
| 71 |
+
{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 2" },
|
| 72 |
+
{ "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
|
| 73 |
+
{ "name": "signalComponents", "type": "u32", "value": "dim(shapes.signal, 2)" },
|
| 74 |
+
{ "name": "frameLength", "type": "u32", "value": "dim(shapes.window, 0)" },
|
| 75 |
+
{ "name": "frames", "type": "u32", "value": "dim(shapes.output, 1)" },
|
| 76 |
+
{ "name": "bins", "type": "u32", "value": "dim(shapes.output, 2)" },
|
| 77 |
+
{ "name": "frameStep", "type": "u32" }
|
| 78 |
+
]
|
| 79 |
+
},
|
| 80 |
+
"params_3": {
|
| 81 |
+
"name": "params",
|
| 82 |
+
"buffer": "uniform",
|
| 83 |
+
"struct": [
|
| 84 |
+
{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 2" },
|
| 85 |
+
{ "name": "signalLength", "type": "u32", "value": "dim(shapes.signal, 1)" },
|
| 86 |
+
{ "name": "signalComponents", "type": "u32", "value": "dim(shapes.signal, 2)" },
|
| 87 |
+
{ "name": "frameLength", "type": "u32" },
|
| 88 |
+
{ "name": "frames", "type": "u32", "value": "dim(shapes.output, 1)" },
|
| 89 |
+
{ "name": "bins", "type": "u32", "value": "dim(shapes.output, 2)" },
|
| 90 |
+
{ "name": "frameStep", "type": "u32" }
|
| 91 |
+
]
|
| 92 |
+
}
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 93 |
},
|
| 94 |
"variants": [
|
| 95 |
{
|
|
|
|
| 100 |
{
|
| 101 |
"id": "main",
|
| 102 |
"name": "STFT.NoWindowFftShared",
|
| 103 |
+
"shader": "fft-stockham-shared.wgsl.jinja",
|
| 104 |
+
"derive": {
|
| 105 |
+
"modeSpec": "\"stft\"",
|
| 106 |
+
"fftLength": "frameLength",
|
| 107 |
+
"numBins": "dim(shapes.output, 2)",
|
| 108 |
+
"inComp": "dim(shapes.signal, 2)",
|
| 109 |
+
"inverse": false,
|
| 110 |
+
"hasWindow": false
|
|
|
|
|
|
|
| 111 |
},
|
| 112 |
+
"bindings": ["signal", "output", "params"],
|
| 113 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 114 |
}
|
| 115 |
]
|
|
|
|
| 117 |
{
|
| 118 |
"id": "rank3_window_fft_shared_inplace",
|
| 119 |
"priority": 32,
|
|
|
|
|
|
|
| 120 |
"when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowInplaceFftFits"],
|
| 121 |
+
"requires": { "limits": { "maxComputeInvocationsPerWorkgroup": 256, "maxComputeWorkgroupSizeX": 256 } },
|
| 122 |
+
"tunables": { "WORKGROUP_SIZE": { "default": 256 } },
|
| 123 |
"passes": [
|
| 124 |
{
|
| 125 |
"id": "main",
|
| 126 |
"name": "STFT.FftSharedInplace",
|
| 127 |
+
"shader": "fft-radix2-dit-shared-inplace.wgsl.jinja",
|
| 128 |
+
"derive": {
|
| 129 |
+
"fftLength": "dim(shapes.window, 0)",
|
| 130 |
+
"log2Length": "rank3WindowFftLog2",
|
| 131 |
+
"numBins": "dim(shapes.output, 2)",
|
| 132 |
+
"inComp": "dim(shapes.signal, 2)"
|
|
|
|
|
|
|
| 133 |
},
|
| 134 |
+
"bindings": ["signal", "window", "output", "params"],
|
| 135 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 136 |
}
|
| 137 |
]
|
|
|
|
| 139 |
{
|
| 140 |
"id": "rank3_window_fft_global_stages",
|
| 141 |
"priority": 31,
|
|
|
|
|
|
|
|
|
|
| 142 |
"when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowGlobalFftFits"],
|
| 143 |
+
"demoteWhen": ["device.limits.maxComputeWorkgroupStorageSize > tunables.GLOBAL_STAGE_MAX_DEVICE_STORAGE_BYTES"],
|
| 144 |
+
"requires": { "limits": { "maxComputeInvocationsPerWorkgroup": 256, "maxComputeWorkgroupSizeX": 256 } },
|
| 145 |
+
"tunables": { "WORKGROUP_SIZE": { "default": 256 } },
|
| 146 |
"intermediates": [
|
| 147 |
{
|
| 148 |
"id": "fft_scratch",
|
|
|
|
| 154 |
{
|
| 155 |
"id": "local_prefix",
|
| 156 |
"name": "STFT.FftLocalPrefix",
|
| 157 |
+
"shader": "fft-radix2-dit-storage-pass.wgsl.jinja",
|
| 158 |
+
"derive": {
|
| 159 |
+
"phase": "\"local_prefix\"",
|
| 160 |
+
"fftLength": "dim(shapes.window, 0)",
|
| 161 |
+
"log2Length": "rank3WindowFftLog2",
|
| 162 |
+
"localLength": "tunables.GLOBAL_STAGE_LOCAL_LENGTH",
|
| 163 |
+
"localLog2Length": "log2ceil(tunables.GLOBAL_STAGE_LOCAL_LENGTH)",
|
| 164 |
+
"numBins": "dim(shapes.output, 2)",
|
| 165 |
+
"inComp": "dim(shapes.signal, 2)"
|
|
|
|
|
|
|
| 166 |
},
|
| 167 |
+
"bindings": ["signal", "window", "scratch", "params"],
|
| 168 |
"dispatch": {
|
| 169 |
+
"x": "min(dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) / tunables.GLOBAL_STAGE_LOCAL_LENGTH, 65535)",
|
| 170 |
+
"y": "ceilDiv(dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) / tunables.GLOBAL_STAGE_LOCAL_LENGTH, 65535)",
|
| 171 |
+
"z": 1
|
| 172 |
}
|
| 173 |
},
|
| 174 |
{
|
|
|
|
| 178 |
{
|
| 179 |
"id": "butterflies",
|
| 180 |
"name": "STFT.FftGlobalStage",
|
| 181 |
+
"shader": "fft-radix2-dit-storage-pass.wgsl.jinja",
|
| 182 |
+
"derive": {
|
| 183 |
+
"phase": "\"stage\"",
|
| 184 |
+
"stage": "repeat.stage + log2ceil(tunables.GLOBAL_STAGE_LOCAL_LENGTH) + 1",
|
| 185 |
+
"fftLength": "dim(shapes.window, 0)",
|
| 186 |
+
"log2Length": "rank3WindowFftLog2",
|
| 187 |
+
"numBins": "dim(shapes.output, 2)",
|
| 188 |
+
"inComp": "dim(shapes.signal, 2)"
|
|
|
|
|
|
|
| 189 |
},
|
| 190 |
+
"bindings": ["scratch"],
|
| 191 |
"dispatch": {
|
| 192 |
+
"x": "min(ceilDiv((dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) / 2), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 193 |
+
"y": "ceilDiv(ceilDiv((dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.window, 0) / 2), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 194 |
+
"z": 1
|
| 195 |
}
|
| 196 |
}
|
| 197 |
]
|
|
|
|
| 199 |
{
|
| 200 |
"id": "finish",
|
| 201 |
"name": "STFT.FftGlobalFinish",
|
| 202 |
+
"shader": "fft-radix2-dit-storage-pass.wgsl.jinja",
|
| 203 |
+
"derive": {
|
| 204 |
+
"phase": "\"finish\"",
|
| 205 |
+
"fftLength": "dim(shapes.window, 0)",
|
| 206 |
+
"log2Length": "rank3WindowFftLog2",
|
| 207 |
+
"numBins": "dim(shapes.output, 2)",
|
| 208 |
+
"inComp": "dim(shapes.signal, 2)"
|
|
|
|
|
|
|
| 209 |
},
|
| 210 |
+
"bindings": [
|
| 211 |
+
{ "scratch": "fft_scratch", "name": "scratch", "buffer": "read-only-storage", "elementType": "f32" },
|
| 212 |
+
"output"
|
| 213 |
+
],
|
| 214 |
"dispatch": {
|
| 215 |
+
"x": "min(ceilDiv((dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.output, 2)), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 216 |
+
"y": "ceilDiv(ceilDiv((dim(shapes.output, 0) * dim(shapes.output, 1) * dim(shapes.output, 2)), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 217 |
+
"z": 1
|
| 218 |
}
|
| 219 |
}
|
| 220 |
]
|
|
|
|
| 234 |
{
|
| 235 |
"id": "main",
|
| 236 |
"name": "STFT.FftStorageInplace",
|
| 237 |
+
"shader": "fft-radix2-dit-storage.wgsl.jinja",
|
| 238 |
+
"derive": {
|
| 239 |
+
"modeSpec": "\"stft\"",
|
| 240 |
+
"fftLength": "dim(shapes.window, 0)",
|
| 241 |
+
"log2Length": "rank3WindowFftLog2",
|
| 242 |
+
"numBins": "dim(shapes.output, 2)",
|
| 243 |
+
"inComp": "dim(shapes.signal, 2)",
|
| 244 |
+
"inverse": false
|
|
|
|
|
|
|
| 245 |
},
|
| 246 |
+
"bindings": ["signal", "window", "scratch", "output", "params"],
|
| 247 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 248 |
}
|
| 249 |
]
|
|
|
|
| 251 |
{
|
| 252 |
"id": "rank3_window_tiled_real",
|
| 253 |
"priority": 19,
|
|
|
|
|
|
|
| 254 |
"when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowTiledRealFits"],
|
| 255 |
+
"requires": { "limits": { "maxComputeInvocationsPerWorkgroup": 256, "maxComputeWorkgroupSizeX": 256 } },
|
| 256 |
+
"tunables": { "WORKGROUP_SIZE": { "default": 256 } },
|
| 257 |
"passes": [
|
| 258 |
{
|
| 259 |
"id": "main",
|
| 260 |
"name": "STFT.TiledReal",
|
| 261 |
+
"shader": "dft-tiled-real.wgsl.jinja",
|
| 262 |
+
"derive": {
|
| 263 |
+
"modeSpec": "\"stft\"",
|
| 264 |
+
"transformLength": "dim(shapes.window, 0)",
|
| 265 |
+
"numBins": "dim(shapes.output, 2)",
|
| 266 |
+
"workgroupSizeSpec": "tunables.WORKGROUP_SIZE"
|
|
|
|
|
|
|
| 267 |
},
|
| 268 |
+
"bindings": ["signal", "window", "output", "params"],
|
| 269 |
"dispatch": {
|
| 270 |
"x": "ceilDiv(dim(shapes.output, 2), 256)",
|
| 271 |
"y": "dim(shapes.output, 1)",
|
|
|
|
| 277 |
{
|
| 278 |
"id": "rank3_window_rfft_packed",
|
| 279 |
"priority": 22,
|
|
|
|
| 280 |
"when": ["rank3WindowLayout", "rank3WindowSpectrum", "rank3WindowSharedFftFits", "attrs.onesided != 0", "dim(shapes.signal, 2) == 1", "dim(shapes.window, 0) >= 64", "dim(shapes.output, 0) * dim(shapes.output, 1) >= tunables.TWIDDLE_TABLE_MIN_TRANSFORMS"],
|
| 281 |
+
"tunables": { "WORKGROUP_SIZE": { "default": 64 } },
|
| 282 |
"intermediates": [{ "id": "fft_twiddles", "dtype": "float32", "shape": "[2 * dim(shapes.window, 0)]" }],
|
| 283 |
"passes": [
|
| 284 |
{
|
| 285 |
"id": "twiddles",
|
| 286 |
"name": "STFT.RealPackedTwiddleTable",
|
| 287 |
+
"shader": "fft-stockham-twiddles.wgsl.jinja",
|
| 288 |
+
"derive": { "fftLength": "dim(shapes.window, 0)", "realPacked": 1 },
|
| 289 |
+
"bindings": ["twiddles"],
|
| 290 |
+
"dispatch": {
|
| 291 |
+
"x": "min(ceilDiv((dim(shapes.window, 0)), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 292 |
+
"y": "ceilDiv(ceilDiv((dim(shapes.window, 0)), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 293 |
+
"z": 1
|
| 294 |
+
}
|
| 295 |
},
|
| 296 |
{
|
| 297 |
"id": "main",
|
| 298 |
"name": "STFT.RealPackedFftShared",
|
| 299 |
+
"shader": "fft-real-packed-stockham-shared.wgsl.jinja",
|
| 300 |
+
"derive": { "fftLength": "dim(shapes.window, 0)", "transformsPerWorkgroup": "packedTransformsPerWorkgroup" },
|
| 301 |
+
"bindings": ["signal", "window", "twiddles_2", "output", "params"],
|
|
|
|
|
|
|
| 302 |
"dispatch": {
|
| 303 |
"x": "ceilDiv(dim(shapes.output, 1), packedTransformsPerWorkgroup)",
|
| 304 |
"y": "dim(shapes.output, 0)"
|
|
|
|
| 315 |
{
|
| 316 |
"id": "twiddles",
|
| 317 |
"name": "STFT.FftTwiddleTable",
|
| 318 |
+
"shader": "fft-stockham-twiddles.wgsl.jinja",
|
| 319 |
+
"derive": { "fftLength": "dim(shapes.window, 0)", "realPacked": 0 },
|
| 320 |
+
"bindings": ["twiddles"],
|
| 321 |
+
"dispatch": {
|
| 322 |
+
"x": "min(ceilDiv((dim(shapes.window, 0) - 1), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 323 |
+
"y": "ceilDiv(ceilDiv((dim(shapes.window, 0) - 1), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 324 |
+
"z": 1
|
| 325 |
+
}
|
| 326 |
},
|
| 327 |
{
|
| 328 |
"id": "main",
|
| 329 |
"name": "STFT.FftSharedTwiddles",
|
| 330 |
+
"shader": "fft-stockham-shared.wgsl.jinja",
|
| 331 |
+
"derive": {
|
| 332 |
+
"modeSpec": "\"stft\"",
|
| 333 |
+
"fftLength": "dim(shapes.window, 0)",
|
| 334 |
+
"numBins": "dim(shapes.output, 2)",
|
| 335 |
+
"inComp": "dim(shapes.signal, 2)",
|
| 336 |
+
"inverse": false,
|
| 337 |
+
"useTwiddleTable": true
|
|
|
|
|
|
|
| 338 |
},
|
| 339 |
+
"bindings": ["signal", "window", "twiddles_2", "output", "params"],
|
| 340 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 341 |
}
|
| 342 |
]
|
|
|
|
| 349 |
{
|
| 350 |
"id": "main",
|
| 351 |
"name": "STFT.FftMixedRecurrent",
|
| 352 |
+
"shader": "fft-stockham-shared.wgsl.jinja",
|
| 353 |
+
"derive": {
|
| 354 |
+
"modeSpec": "\"stft\"",
|
| 355 |
+
"fftLength": "dim(shapes.window, 0)",
|
| 356 |
+
"radix2Length": "rank3WindowFftRadix2",
|
| 357 |
+
"mixedRadixTail": true,
|
| 358 |
+
"numBins": "dim(shapes.output, 2)",
|
| 359 |
+
"inComp": "dim(shapes.signal, 2)",
|
| 360 |
+
"inverse": false
|
|
|
|
|
|
|
| 361 |
},
|
| 362 |
+
"bindings": ["signal", "window", "output", "params"],
|
| 363 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 364 |
}
|
| 365 |
]
|
|
|
|
| 372 |
{
|
| 373 |
"id": "main",
|
| 374 |
"name": "STFT.FftShared",
|
| 375 |
+
"shader": "fft-stockham-shared.wgsl.jinja",
|
| 376 |
+
"derive": {
|
| 377 |
+
"modeSpec": "\"stft\"",
|
| 378 |
+
"fftLength": "dim(shapes.window, 0)",
|
| 379 |
+
"numBins": "dim(shapes.output, 2)",
|
| 380 |
+
"inComp": "dim(shapes.signal, 2)",
|
| 381 |
+
"inverse": false
|
|
|
|
|
|
|
| 382 |
},
|
| 383 |
+
"bindings": ["signal", "window", "output", "params"],
|
| 384 |
"dispatch": { "x": "dim(shapes.output, 1)", "y": "dim(shapes.output, 0)" }
|
| 385 |
}
|
| 386 |
]
|
|
|
|
| 388 |
{
|
| 389 |
"id": "rank3_window",
|
| 390 |
"priority": 0,
|
|
|
|
| 391 |
"when": ["rank3WindowLayout", "rank3WindowSpectrum"],
|
| 392 |
+
"supersededBy": ["rank3_window_fft_shared_inplace", "rank3_window_fft_global_stages", "rank3_window_fft_storage_inplace", "rank3_window_fft_mixed", "rank3_window_rfft_packed", "rank3_window_fft_twiddles", "rank3_window_fft", "rank3_window_tiled_real"],
|
| 393 |
"passes": [
|
| 394 |
{
|
| 395 |
"id": "main",
|
| 396 |
"name": "STFT",
|
| 397 |
+
"shader": "stft-general.wgsl.jinja",
|
| 398 |
+
"derive": { "hasWindowInput": true },
|
| 399 |
+
"bindings": ["signal", "window", "output", "params_2"],
|
| 400 |
+
"dispatch": {
|
| 401 |
+
"x": "min(ceilDiv((numel(shapes.output) / 2), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 402 |
+
"y": "ceilDiv(ceilDiv((numel(shapes.output) / 2), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 403 |
+
"z": 1
|
| 404 |
+
}
|
| 405 |
}
|
| 406 |
]
|
| 407 |
},
|
|
|
|
| 413 |
{
|
| 414 |
"id": "main",
|
| 415 |
"name": "STFT.GeneralNoWindow",
|
| 416 |
+
"shader": "stft-general.wgsl.jinja",
|
| 417 |
+
"derive": { "hasWindowInput": false },
|
| 418 |
+
"bindings": ["signal", "output", "params_3"],
|
| 419 |
+
"dispatch": {
|
| 420 |
+
"x": "min(ceilDiv((numel(shapes.output) / 2), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 421 |
+
"y": "ceilDiv(ceilDiv((numel(shapes.output) / 2), (tunables.WORKGROUP_SIZE)), 65535)",
|
| 422 |
+
"z": 1
|
| 423 |
+
}
|
| 424 |
}
|
| 425 |
]
|
| 426 |
}
|
build/webgpu/metadata.json
CHANGED
|
@@ -1,25 +1,40 @@
|
|
| 1 |
{
|
| 2 |
"name": "ai.onnx.STFT",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"backend": { "type": "webgpu" },
|
| 7 |
"digest": {
|
| 8 |
"algorithm": "sha256",
|
| 9 |
"files": {
|
| 10 |
-
"bench.json": "
|
| 11 |
-
"dft-tiled-real.wgsl.jinja": "
|
| 12 |
-
"fft-radix2-dit-shared-inplace.wgsl.jinja": "
|
| 13 |
-
"fft-radix2-dit-storage-pass.wgsl.jinja": "
|
| 14 |
-
"fft-radix2-dit-storage.wgsl.jinja": "
|
| 15 |
-
"fft-real-packed-stockham-shared.wgsl.jinja": "
|
| 16 |
-
"fft-stockham-shared.wgsl.jinja": "
|
| 17 |
-
"fft-stockham-twiddles.wgsl.jinja": "
|
| 18 |
-
"manifest.json": "
|
| 19 |
-
"stft-general.wgsl.jinja": "
|
| 20 |
-
"test.json": "
|
| 21 |
}
|
| 22 |
},
|
| 23 |
-
"provenance": { "kernel": { "sha": "
|
| 24 |
-
"webgpu": {
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 25 |
}
|
|
|
|
| 1 |
{
|
| 2 |
"name": "ai.onnx.STFT",
|
| 3 |
+
"id": "_ai_onnx_stft_webgpu_06cb038",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"backend": { "type": "webgpu" },
|
| 7 |
"digest": {
|
| 8 |
"algorithm": "sha256",
|
| 9 |
"files": {
|
| 10 |
+
"bench.json": "vHcB5tti+CM3DMtdpAb1k6Lb72nQ3dLXXqVWPYHWi+Y=",
|
| 11 |
+
"dft-tiled-real.wgsl.jinja": "nls3dnjWMqMRdmwudzIVekmyeWr13k1DaSsSxfdX+Xo=",
|
| 12 |
+
"fft-radix2-dit-shared-inplace.wgsl.jinja": "uKO/mdFdypgtav2qghpelG0iwgR4WXM/mpgJVy2+MQA=",
|
| 13 |
+
"fft-radix2-dit-storage-pass.wgsl.jinja": "lSFcs2O2JtHaV8WGmyvtfo/QZi6xaflVR71cOip6zJ8=",
|
| 14 |
+
"fft-radix2-dit-storage.wgsl.jinja": "Rc4YVTVGib28SSfikk4IPyLTDL3BlEGCAsH04WyfqbE=",
|
| 15 |
+
"fft-real-packed-stockham-shared.wgsl.jinja": "fPcCFvTUojCPabsRnC9y7GvJVkogRQ7zcQtC48YAR4s=",
|
| 16 |
+
"fft-stockham-shared.wgsl.jinja": "W1qZ/FnWKVSZSpE43qo/06isxY6D+Va6+tzG7/myock=",
|
| 17 |
+
"fft-stockham-twiddles.wgsl.jinja": "SVeaa/kM8Fx9RHH3yVcvx/twJqPjZtrffvi0pGesdyg=",
|
| 18 |
+
"manifest.json": "8J0Y8qFjzAsLgbgEydZN8iD67cDTFYyOs1daSYYUeaA=",
|
| 19 |
+
"stft-general.wgsl.jinja": "R006F8Q2a4U7S+V1FyFsc+nUUMjx5ya4yGOFkcIiW+w=",
|
| 20 |
+
"test.json": "R2V8u8+fDmmwimCB8Cwj+xradfacxn/hrIJ2L/oF49Q="
|
| 21 |
}
|
| 22 |
},
|
| 23 |
+
"provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
|
| 24 |
+
"webgpu": {
|
| 25 |
+
"manifestSpec": "2.0",
|
| 26 |
+
"variants": {
|
| 27 |
+
"rank3_no_window_fft": ["fft-stockham-shared.wgsl.jinja"],
|
| 28 |
+
"rank3_window_fft_shared_inplace": ["fft-radix2-dit-shared-inplace.wgsl.jinja"],
|
| 29 |
+
"rank3_window_fft_global_stages": ["fft-radix2-dit-storage-pass.wgsl.jinja"],
|
| 30 |
+
"rank3_window_fft_storage_inplace": ["fft-radix2-dit-storage.wgsl.jinja"],
|
| 31 |
+
"rank3_window_tiled_real": ["dft-tiled-real.wgsl.jinja"],
|
| 32 |
+
"rank3_window_rfft_packed": ["fft-real-packed-stockham-shared.wgsl.jinja", "fft-stockham-twiddles.wgsl.jinja"],
|
| 33 |
+
"rank3_window_fft_twiddles": ["fft-stockham-shared.wgsl.jinja", "fft-stockham-twiddles.wgsl.jinja"],
|
| 34 |
+
"rank3_window_fft_mixed": ["fft-stockham-shared.wgsl.jinja"],
|
| 35 |
+
"rank3_window_fft": ["fft-stockham-shared.wgsl.jinja"],
|
| 36 |
+
"rank3_window": ["stft-general.wgsl.jinja"],
|
| 37 |
+
"rank3_no_window": ["stft-general.wgsl.jinja"]
|
| 38 |
+
}
|
| 39 |
+
}
|
| 40 |
}
|
build/webgpu/stft-general.wgsl.jinja
CHANGED
|
@@ -1,30 +1,30 @@
|
|
| 1 |
{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
|
| 2 |
{% if note == "dispatch-limit" %}
|
| 3 |
-
// 2D-folded flat index: gid.y carries the high bits past the
|
| 4 |
-
//
|
| 5 |
{% elif note == "limit" %}
|
| 6 |
-
// 2D-folded flat index: gid.y carries the high bits past the
|
| 7 |
-
//
|
| 8 |
{% elif note == "device-axis" %}
|
| 9 |
-
// The flat dispatch is folded across x/y at
|
| 10 |
-
//
|
| 11 |
{% elif note == "vec4-limit" %}
|
| 12 |
-
// 2D-folded flat vec4 index: gid.y carries the high bits past the
|
| 13 |
-
//
|
| 14 |
{% elif note == "element-limit" %}
|
| 15 |
// 2D-folded flat element index: gid.y carries the high bits past the
|
| 16 |
-
//
|
| 17 |
{% elif note == "dispatch" %}
|
| 18 |
-
// 2D-folded flat index: gid.y carries the high bits past the
|
| 19 |
-
//
|
| 20 |
{% endif %}
|
| 21 |
{% if bound == "" %}
|
| 22 |
-
let {{ name }} = gid.x + gid.y *
|
| 23 |
{%- elif guardInline %}
|
| 24 |
-
let {{ name }} = gid.x + gid.y *
|
| 25 |
if ({{ name }} >= {{ bound }}) { return; }
|
| 26 |
{%- else %}
|
| 27 |
-
let {{ name }} = gid.x + gid.y *
|
| 28 |
if ({{ name }} >= {{ bound }}) {
|
| 29 |
return;
|
| 30 |
}
|
|
@@ -33,18 +33,17 @@
|
|
| 33 |
|
| 34 |
{{ env.wgsl.resourceDeclarations }}
|
| 35 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 36 |
-
// rational turn and route
|
| 37 |
-
//
|
| 38 |
//
|
| 39 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 40 |
// inverse transform. The result is (cos, sin).
|
| 41 |
// Backend-stable f32 sine/cosine core.
|
| 42 |
//
|
| 43 |
-
// Shader transcendental accuracy is implementation-defined
|
| 44 |
-
//
|
| 45 |
-
//
|
| 46 |
-
//
|
| 47 |
-
// truncation error is well below one f32 ULP over its documented interval.
|
| 48 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 49 |
const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
|
| 50 |
const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
|
|
@@ -118,7 +117,7 @@ fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
|
|
| 118 |
// scalar. Window samples default to 1.0 when no window input is present.
|
| 119 |
|
| 120 |
@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
|
| 121 |
-
fn main(@builtin(global_invocation_id) gid: vec3<u32>
|
| 122 |
{{ flat_index_2d("index", guardInline=true, note="") }}
|
| 123 |
|
| 124 |
let bin = index % params.bins;
|
|
@@ -136,7 +135,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
|
|
| 136 |
for (var n = 0u; n < frame_length; n = n + 1u) {
|
| 137 |
let signal_index = start + n;
|
| 138 |
let signal_base = (batch * params.signalLength + signal_index) * params.signalComponents;
|
| 139 |
-
{% if
|
| 140 |
let w = window[n];
|
| 141 |
{% else %}
|
| 142 |
let w = 1.0;
|
|
|
|
| 1 |
{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
|
| 2 |
{% if note == "dispatch-limit" %}
|
| 3 |
+
// 2D-folded flat index: gid.y carries the high bits past the dispatch's
|
| 4 |
+
// per-axis workgroup fold width (outputs > 16.7M elements).
|
| 5 |
{% elif note == "limit" %}
|
| 6 |
+
// 2D-folded flat index: gid.y carries the high bits past the dispatch's
|
| 7 |
+
// per-axis workgroup fold width.
|
| 8 |
{% elif note == "device-axis" %}
|
| 9 |
+
// The flat dispatch is folded across x/y at a fixed per-axis workgroup
|
| 10 |
+
// width; gid.y carries the high portion of the output index.
|
| 11 |
{% elif note == "vec4-limit" %}
|
| 12 |
+
// 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
|
| 13 |
+
// per-axis workgroup fold width (the dispatch caps x and spills into y).
|
| 14 |
{% elif note == "element-limit" %}
|
| 15 |
// 2D-folded flat element index: gid.y carries the high bits past the
|
| 16 |
+
// dispatch's per-axis workgroup fold width.
|
| 17 |
{% elif note == "dispatch" %}
|
| 18 |
+
// 2D-folded flat index: gid.y carries the high bits past the dispatch's
|
| 19 |
+
// per-axis workgroup fold width.
|
| 20 |
{% endif %}
|
| 21 |
{% if bound == "" %}
|
| 22 |
+
let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
|
| 23 |
{%- elif guardInline %}
|
| 24 |
+
let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
|
| 25 |
if ({{ name }} >= {{ bound }}) { return; }
|
| 26 |
{%- else %}
|
| 27 |
+
let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
|
| 28 |
if ({{ name }} >= {{ bound }}) {
|
| 29 |
return;
|
| 30 |
}
|
|
|
|
| 33 |
|
| 34 |
{{ env.wgsl.resourceDeclarations }}
|
| 35 |
// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
|
| 36 |
+
// rational turn and route polynomial evaluation through one trigonometric core
|
| 37 |
+
// for FFT, direct DFT, STFT, and cosine-window calculations.
|
| 38 |
//
|
| 39 |
// phase is in [0, period). sign is -1 for a forward transform and +1 for an
|
| 40 |
// inverse transform. The result is (cos, sin).
|
| 41 |
// Backend-stable f32 sine/cosine core.
|
| 42 |
//
|
| 43 |
+
// Shader transcendental accuracy is implementation-defined. Each path retains
|
| 44 |
+
// its available phase representation, reduces it to [-pi, pi], then evaluates
|
| 45 |
+
// explicit polynomials in a fixed order. The half-pi core is degree 13 for sine
|
| 46 |
+
// and degree 12 for cosine.
|
|
|
|
| 47 |
const PRECISE_TRIG_PI: f32 = 3.141592653589793;
|
| 48 |
const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
|
| 49 |
const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
|
|
|
|
| 117 |
// scalar. Window samples default to 1.0 when no window input is present.
|
| 118 |
|
| 119 |
@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
|
| 120 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 121 |
{{ flat_index_2d("index", guardInline=true, note="") }}
|
| 122 |
|
| 123 |
let bin = index % params.bins;
|
|
|
|
| 135 |
for (var n = 0u; n < frame_length; n = n + 1u) {
|
| 136 |
let signal_index = start + n;
|
| 137 |
let signal_base = (batch * params.signalLength + signal_index) * params.signalComponents;
|
| 138 |
+
{% if hasWindowInput %}
|
| 139 |
let w = window[n];
|
| 140 |
{% else %}
|
| 141 |
let w = 1.0;
|
build/webgpu/test.json
CHANGED
|
@@ -1,5 +1,4 @@
|
|
| 1 |
{
|
| 2 |
-
"op": "ai.onnx.STFT",
|
| 3 |
"fixtureArrays": {
|
| 4 |
"onnx_backend_arange128_cos_window16_input_window": [1, 0.961939766, 0.853553391, 0.691341716, 0.5, 0.308658284, 0.146446609, 0.038060234, 0, 0.038060234, 0.146446609, 0.308658284, 0.5, 0.691341716, 0.853553391, 0.961939766],
|
| 5 |
"ort_complex_input_batched_twosided_frame_length_input_signal": [1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0]
|
|
@@ -18,7 +17,7 @@
|
|
| 18 |
{
|
| 19 |
"name": "many_frames_fft_twiddle_table_win64",
|
| 20 |
"provenance": {
|
| 21 |
-
"notes": "
|
| 22 |
},
|
| 23 |
"inputs": {
|
| 24 |
"signal": { "dtype": "float32", "shape": [1, 4159, 1] },
|
|
@@ -30,7 +29,7 @@
|
|
| 30 |
{
|
| 31 |
"name": "many_frames_fft_twiddle_table_win64_twosided",
|
| 32 |
"provenance": {
|
| 33 |
-
"notes": "
|
| 34 |
},
|
| 35 |
"attrs": { "onesided": 0 },
|
| 36 |
"inputs": {
|
|
@@ -53,7 +52,7 @@
|
|
| 53 |
{
|
| 54 |
"name": "many_frames_packed_rfft_win512_radix4",
|
| 55 |
"provenance": {
|
| 56 |
-
"notes": "
|
| 57 |
},
|
| 58 |
"inputs": {
|
| 59 |
"signal": {
|
|
@@ -425,7 +424,7 @@
|
|
| 425 |
{
|
| 426 |
"name": "fft_window512_real_twosided_step256",
|
| 427 |
"provenance": {
|
| 428 |
-
"notes": "
|
| 429 |
},
|
| 430 |
"attrs": { "onesided": 0 },
|
| 431 |
"inputs": {
|
|
@@ -489,7 +488,7 @@
|
|
| 489 |
{
|
| 490 |
"name": "fft_window2048_radix2_dit_onesided_min_storage",
|
| 491 |
"provenance": {
|
| 492 |
-
"notes": "
|
| 493 |
},
|
| 494 |
"inputs": {
|
| 495 |
"signal": {
|
|
@@ -684,7 +683,7 @@
|
|
| 684 |
{
|
| 685 |
"name": "mixed_stockham_n1000_real_twosided_tail125",
|
| 686 |
"provenance": {
|
| 687 |
-
"notes": "Window length
|
| 688 |
},
|
| 689 |
"attrs": { "onesided": 0 },
|
| 690 |
"inputs": {
|
|
|
|
| 1 |
{
|
|
|
|
| 2 |
"fixtureArrays": {
|
| 3 |
"onnx_backend_arange128_cos_window16_input_window": [1, 0.961939766, 0.853553391, 0.691341716, 0.5, 0.308658284, 0.146446609, 0.038060234, 0, 0.038060234, 0.146446609, 0.308658284, 0.5, 0.691341716, 0.853553391, 0.961939766],
|
| 4 |
"ort_complex_input_batched_twosided_frame_length_input_signal": [1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 1, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0, 99, 0]
|
|
|
|
| 17 |
{
|
| 18 |
"name": "many_frames_fft_twiddle_table_win64",
|
| 19 |
"provenance": {
|
| 20 |
+
"notes": "A length-64 window over 4,159 samples produces 4,096 independent transforms and selects the shared twiddle-table FFT geometry."
|
| 21 |
},
|
| 22 |
"inputs": {
|
| 23 |
"signal": { "dtype": "float32", "shape": [1, 4159, 1] },
|
|
|
|
| 29 |
{
|
| 30 |
"name": "many_frames_fft_twiddle_table_win64_twosided",
|
| 31 |
"provenance": {
|
| 32 |
+
"notes": "A 4,096-frame, 64-point transform with `onesided=0` requires the full spectrum and exercises the shared twiddle-table FFT."
|
| 33 |
},
|
| 34 |
"attrs": { "onesided": 0 },
|
| 35 |
"inputs": {
|
|
|
|
| 52 |
{
|
| 53 |
"name": "many_frames_packed_rfft_win512_radix4",
|
| 54 |
"provenance": {
|
| 55 |
+
"notes": "A real 512-point STFT over 4,097 transforms exercises packed processing and leaves one partial four-transform workgroup when that specialization is available."
|
| 56 |
},
|
| 57 |
"inputs": {
|
| 58 |
"signal": {
|
|
|
|
| 424 |
{
|
| 425 |
"name": "fft_window512_real_twosided_step256",
|
| 426 |
"provenance": {
|
| 427 |
+
"notes": "A real length-512 transform with onesided=0 emits the full 512-bin spectrum; frame step 256 produces three frames."
|
| 428 |
},
|
| 429 |
"attrs": { "onesided": 0 },
|
| 430 |
"inputs": {
|
|
|
|
| 488 |
{
|
| 489 |
"name": "fft_window2048_radix2_dit_onesided_min_storage",
|
| 490 |
"provenance": {
|
| 491 |
+
"notes": "A 2048-point window exercises the in-place radix-2 specialization at the 16 KiB WebGPU workgroup-storage minimum: its 8 bytes per element scratch fits while the 16 bytes per element two-buffer scratch does not. On larger storage limits the case remains valid and may use the two-buffer specialization instead, so it requires no explicit device limit."
|
| 492 |
},
|
| 493 |
"inputs": {
|
| 494 |
"signal": {
|
|
|
|
| 683 |
{
|
| 684 |
"name": "mixed_stockham_n1000_real_twosided_tail125",
|
| 685 |
"provenance": {
|
| 686 |
+
"notes": "Window length 1,000 factors as 8 x 125, selecting the shared mixed-radix transform. With onesided=0, each frame emits the full 1,000-bin spectrum."
|
| 687 |
},
|
| 688 |
"attrs": { "onesided": 0 },
|
| 689 |
"inputs": {
|