sync 2e7068faf55e

Browse files

Files changed (6) hide show

README.md +74 -0
build/webgpu/bench.json +25 -0
build/webgpu/manifest.json +74 -0
build/webgpu/metadata.json +18 -0
build/webgpu/test.json +217 -0
build/webgpu/window-cosine-sum.wgsl.jinja +129 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.HannWindow
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 17
+## Description
+Generates a Hann window of a given length using the formula `0.5 - 0.5 * cos(2π * n / N)`, where `N` is `size` for periodic mode or `size - 1` for symmetric mode. The output is a 1-D tensor of length `size`.
+See the [ONNX `HannWindow` spec](https://onnx.ai/onnx/operators/onnx__HannWindow.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `size` | `size` | `T1` | `0` | — | Scalar indicating the number of elements in the output window. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `y` | `T2` | `1` | — | 1-D Hann window tensor of shape `[size]`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `output_datatype` | `1` | Data type of the output tensor, encoded as a DataType enum value from TensorProto (default `1` = FLOAT). |
+| `periodic` | `1` | If `1` (default), returns a periodic window of length `size`; if `0`, returns a symmetric window of length `size`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T1` | `int32` |
+| `T2` | `float32`, `float16`, `uint32`, `int32`, `uint8`, `int8`, `int16` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`window-cosine-sum.wgsl.jinja`](build/webgpu/window-cosine-sum.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader automatically allocates outputs whose metadata it can derive from the manifest contract and this call.
+The explicit `outputs` entries provide shape and logical dtype metadata for the results listed below:
+- `y`
+Each entry either requests an optional result or supplies metadata that cannot be inferred from the inputs.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.HannWindow", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { y } = await kernel({ size: { data: sizeData, shape: [] } }, {
+  outputs: { y: { shape: [1], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

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

+{
+  "op": "ai.onnx.HannWindow",
+  "cases": [
+    {
+      "name": "hann-f32-4096",
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [4096] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [4096] } }
+    },
+    {
+      "name": "hann-f16-1m",
+      "preset": "edge",
+      "attrs": { "output_datatype": 10 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1048576] } } },
+      "outputs": { "y": { "dtype": "float16", "shape": [1048576], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2 * numel(shapes.y)", "name": "write-BW" }] }
+    },
+    {
+      "name": "hann-f32-1m",
+      "preset": "smoke",
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1048576] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1048576], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.y) * 4", "name": "write-BW" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,74 @@

+{
+  "domain": "ai.onnx",
+  "name": "HannWindow",
+  "sinceVersion": 17,
+  "description": "Generates a Hann window of a given length using the formula `0.5 - 0.5 * cos(2π * n / N)`, where `N` is `size` for periodic mode or `size - 1` for symmetric mode. The output is a 1-D tensor of length `size`.",
+  "inputs": [
+    {
+      "role": "size",
+      "dtype": "T1",
+      "description": "Scalar indicating the number of elements in the output window.",
+      "rank": 0
+    }
+  ],
+  "outputs": [
+    { "role": "output", "dtype": "T2", "description": "1-D Hann window tensor of shape `[size]`.", "rank": 1 }
+  ],
+  "attributes": { "output_datatype": 1, "periodic": 1 },
+  "attributeDescriptions": {
+    "output_datatype": "Data type of the output tensor, encoded as a DataType enum value from TensorProto (default `1` = FLOAT).",
+    "periodic": "If `1` (default), returns a periodic window of length `size`; if `0`, returns a symmetric window of length `size`."
+  },
+  "attributeConstraints": { "periodic": { "values": [0, 1] } },
+  "typeConstraints": { "T1": ["int32"], "T2": ["float32", "float16", "uint32", "int32", "uint8", "int8", "int16"] },
+  "args": {
+    "size": { "kind": "tensor", "semantic": "size", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {},
+  "variants": [
+    {
+      "id": "generic",
+      "when": "ranks.output == 1 and attrs.output_datatype == onnxDtypeCode(logicalDtypes.T2) and f16Ok(tensorDtypes.output)",
+      "constants": {
+        "usesF16": "tensorDtypes.output == \"float16\"",
+        "integerOutput": "tensorDtypes.output != \"float32\" and tensorDtypes.output != \"float16\"",
+        "outScalar": "dtypes.T2"
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "HannWindow",
+          "source": { "shader": "window-cosine-sum.wgsl.jinja", "inputs": { "a0": "0.5", "a1": "0.5", "a2": "0" } },
+          "bindings": [
+            {
+              "name": "y",
+              "arg": "y",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$outScalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "count", "type": "u32", "value": "numel(shapes.output)" },
+                  {
+                    "name": "denom",
+                    "type": "f32",
+                    "value": "dim(shapes.output, 0) if attrs.periodic else dim(shapes.output, 0) - 1"
+                  }
+                ]
+              }
+            }
+          ],
+          "dispatch": { "gridStride": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,18 @@

+{
+  "name": "ai.onnx.HannWindow",
+  "id": "_ai_onnx_hannwindow_webgpu_d631c24",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "PEIGhBPxuOnWnbnYZ7QAl796F7fkVaYB3agt7stTl7c=",
+      "manifest.json": "pLxvdWH8iSFyF9krpRrWQGiCCJ/bAW8+ksjxaaY7D1U=",
+      "test.json": "f/qlcQHxoUOnaLhNSSG6jsdaOpn929gU3yLsEAi5XyY=",
+      "window-cosine-sum.wgsl.jinja": "T17RHnl56gmwE9BtwVPMeL6Gb+bO/QFqpqKZ9q6EiBE="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.HannWindow" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,217 @@

+{
+  "op": "ai.onnx.HannWindow",
+  "cases": [
+    {
+      "name": "periodic_f32_size16",
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [16], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "symmetric_f32_size16",
+      "attrs": { "periodic": 0 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [16], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "periodic_f32_size10",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hannwindow" },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [10] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [10], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "symmetric_f32_size10_onnx_backend",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_hannwindow_symmetric" },
+      "attrs": { "periodic": 0 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [10] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [10], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "periodic_f32_size0_empty",
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [0] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "periodic_f32_size1_boundary",
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "symmetric_f32_size1_nan_boundary",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.HannWindowFloat",
+        "notes": "Additional ORT-validated edge: symmetric size 1 uses denominator length - 1 = 0 and produces NaN."
+      },
+      "attrs": { "periodic": 0 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } } },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [1],
+          "tolerance": 0,
+          "allowNaN": true,
+          "data": { "kind": "values", "values": ["NaN"] }
+        }
+      }
+    },
+    {
+      "name": "symmetric_f32_odd_size7",
+      "attrs": { "periodic": 0 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [7] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [7], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "periodic_f16_size32",
+      "attrs": { "output_datatype": 10 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [32] } } },
+      "outputs": { "y": { "dtype": "float16", "shape": [32], "tolerance": 0.002 } }
+    },
+    {
+      "name": "ort_periodic_uint32_output_size8",
+      "provenance": {
+        "source": "ONNX Runtime CPUExecutionProvider",
+        "notes": "ONNX HannWindow computes in float and applies output_datatype=UINT32 as a final cast."
+      },
+      "attrs": { "output_datatype": 12 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [8] } } },
+      "outputs": {
+        "y": {
+          "dtype": "uint32",
+          "shape": [8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 1, 0, 0, 0] }
+        }
+      }
+    },
+    {
+      "name": "ort_symmetric_int32_output_size1_zero_cast",
+      "provenance": {
+        "source": "ONNX Runtime CPUExecutionProvider",
+        "notes": "The symmetric size-one formula produces NaN in float; ONNX Cast to INT32 maps that value to zero."
+      },
+      "attrs": { "output_datatype": 6, "periodic": 0 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [1] } } },
+      "outputs": {
+        "y": { "dtype": "int32", "shape": [1], "tolerance": 0, "data": { "kind": "values", "values": [0] } }
+      }
+    },
+    {
+      "name": "onnx_standard_periodic_int8_output_size8",
+      "attrs": { "output_datatype": 3 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [8] } } },
+      "outputs": {
+        "y": {
+          "dtype": "int8",
+          "shape": [8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 1, 0, 0, 0] }
+        }
+      }
+    },
+    {
+      "name": "onnx_standard_periodic_uint8_output_size8",
+      "attrs": { "output_datatype": 2 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [8] } } },
+      "outputs": {
+        "y": {
+          "dtype": "uint8",
+          "shape": [8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 1, 0, 0, 0] }
+        }
+      }
+    },
+    {
+      "name": "onnx_standard_periodic_int16_output_size8",
+      "attrs": { "output_datatype": 5 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [8] } } },
+      "outputs": {
+        "y": {
+          "dtype": "int16",
+          "shape": [8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 0, 0, 1, 0, 0, 0] }
+        }
+      }
+    },
+    {
+      "name": "periodic_f32_size257_phase_accuracy_gpu_gap",
+      "skipGpu": {
+        "category": "todo",
+        "reason": "The current f32 cosine and range-reduction path crosses the final rounding boundary for this periodic size-257 window. A correctly rounded approximation or software-extended precision could close this implementable gap."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.HannWindowFloat",
+        "notes": "Medium odd-sized periodic window stresses f32 phase formation across many non-cardinal cosine positions; CPU and ORT agree more tightly than the current WGSL f32 cosine-sum path."
+      },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [257] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [257], "tolerance": 6e-8, "relTolerance": 0 } }
+    },
+    {
+      "name": "symmetric_f32_size257_phase_accuracy_gpu_gap",
+      "skipGpu": {
+        "category": "todo",
+        "reason": "The current f32 cosine and range-reduction path crosses the final rounding boundary for this symmetric size-257 window. A correctly rounded approximation or software-extended precision could close this implementable gap."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.HannWindowFloat",
+        "notes": "Symmetric odd-sized companion to the periodic phase edge; the length-1 denominator shifts every non-cardinal cosine sample and exposes f32 phase drift in the WGSL path."
+      },
+      "attrs": { "periodic": 0 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [257] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [257], "tolerance": 6e-8, "relTolerance": 0 } }
+    },
+    {
+      "name": "ort_periodic_f32_size32",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/signal/signal_ops_test.cc",
+        "test": "SignalOpsTest.HannWindowFloat"
+      },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [32] } } },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [32],
+          "tolerance": 0.00001,
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.009607, 0.03806, 0.084265, 0.146447, 0.222215, 0.308658, 0.402455, 0.5, 0.597545, 0.691342, 0.777785, 0.853553, 0.915735, 0.96194, 0.990393, 1.0, 0.990393, 0.96194, 0.915735, 0.853553, 0.777785, 0.691342, 0.597545, 0.5, 0.402455, 0.308658, 0.222215, 0.146447, 0.084265, 0.03806, 0.009607]
+          }
+        }
+      }
+    },
+    {
+      "name": "periodic_int32_size8",
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [8] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "symmetric_f32_size2_denom1",
+      "attrs": { "periodic": 0 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [2] } } },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [2],
+          "tolerance": 0.00001,
+          "data": { "kind": "values", "values": [0.0, 0.0] }
+        }
+      }
+    },
+    {
+      "name": "symmetric_f32_size4_cardinal_verify",
+      "attrs": { "periodic": 0 },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [4] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "symmetric_f16_size16",
+      "attrs": { "output_datatype": 10, "periodic": 0 },
+      "requires": { "features": ["shader-f16"] },
+      "inputs": { "size": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16] } } },
+      "outputs": { "y": { "dtype": "float16", "shape": [16], "tolerance": 0.002 } }
+    }
+  ]
+}

build/webgpu/window-cosine-sum.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,129 @@

+// Cosine-sum window:
+//   w = a0 - a1*cos(phase) + a2*cos(2*phase), phase = 2*pi*n/N
+// Compiled coefficients select Hann, Hamming, or Blackman. N (params.denom) is
+// the window length for a periodic window, or length - 1 for a symmetric window.
+//
+// Matching the f64-rounded target within roughly 0.5 ULP requires two
+// details beyond a naive `a0 - a1*cos(2*pi*n/N)`: keep n/N as a rational turn and
+// use the backend-stable unit-circle polynomial, and form a0 - a1*c with one
+// rounding via fma(-a1, c, a0). WGSL builtin transcendental accuracy is
+// implementation-defined, and a rounded multiply followed by subtraction would
+// double-round. Grid strides cover large windows with a clamped dispatch.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// Backend-stable unit-circle twiddles for DFT kernels. Retain the phase as a
+// rational turn and route all polynomial evaluation through the shared precise
+// trig core so FFT, direct DFT, STFT, and cosine windows agree exactly.
+//
+// phase is in [0, period). sign is -1 for a forward transform and +1 for an
+// inverse transform. The result is (cos, sin).
+// Backend-stable f32 sine/cosine core.
+//
+// Shader transcendental accuracy is implementation-defined, and some portable
+// backends are only accurate to roughly 1e-4. Each path retains the most accurate
+// available phase representation, reduces it to [-pi, pi], then uses these
+// polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
+// truncation error is well below one f32 ULP over its documented interval.
+const PRECISE_TRIG_PI: f32 = 3.141592653589793;
+const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
+const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
+// Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
+fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
+  let x2 = x * x;
+  var sinPolynomial = 1.6059043836821613e-10;
+  sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
+  sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
+  sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
+  sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
+  sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
+  let sine = x * fma(sinPolynomial, x2, 1.0);
+  var cosPolynomial = 2.08767569878681e-9;
+  cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
+  cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
+  cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
+  cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
+  cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
+  let cosine = fma(cosPolynomial, x2, 1.0);
+  return vec2<f32>(cosine, sine);
+}
+// Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
+fn precise_sincos_centered(x: f32) -> vec2<f32> {
+  var folded = x;
+  var cosineSign = 1.0;
+  if (folded > PRECISE_TRIG_HALF_PI) {
+    folded = PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  } else if (folded < -PRECISE_TRIG_HALF_PI) {
+    folded = -PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  }
+  let value = precise_sincos_half_pi(folded);
+  return vec2<f32>(cosineSign * value.x, value.y);
+}
+fn dft_precise_twiddle(phase: u32, period: u32, sign: f32) -> vec2<f32> {
+  let reducedPhase = phase % period;
+  if (reducedPhase == 0u) {
+    return vec2<f32>(1.0, sign * 0.0);
+  }
+  if (period % 4u == 0u) {
+    let quarter = period / 4u;
+    if (reducedPhase == quarter) {
+      return vec2<f32>(0.0, sign);
+    }
+    if (reducedPhase == 2u * quarter) {
+      return vec2<f32>(-1.0, sign * 0.0);
+    }
+    if (reducedPhase == 3u * quarter) {
+      return vec2<f32>(0.0, -sign);
+    }
+  }
+  // Center the rational turn before multiplying by 2*pi. This avoids the
+  // cancellation in `2*pi*(N-1)/N - 2*pi` for bins near the end of a row.
+  let turns = f32(reducedPhase) / f32(period);
+  let centeredTurns = turns - floor(turns + 0.5);
+  return precise_sincos_centered(sign * PRECISE_TRIG_TWO_PI * centeredTurns);
+}
+const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
+fn cos_turns(phase: u32, period: u32) -> f32 {
+  return dft_precise_twiddle(phase % period, period, 1.0).x;
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(
+  @builtin(global_invocation_id) gid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  let stride = nwg.x * WG;
+  for (var i = gid.x; i < params.count; i += stride) {
+    if (params.denom == 0.0) {
+{% if integerOutput %}
+      // The window is evaluated in f32 before the output cast. Spell the
+      // symmetric-size-one integer result directly instead of converting NaN.
+      y[i] = {{ outScalar }}(0);
+{% else %}
+      // Do not synthesize NaN with 0/0: fast-math can fold identical operands to
+      // 1. Incorporating the runtime uniform also keeps
+      // WGSL validators from rejecting a compile-time non-finite constant.
+      let nan_bits = bitcast<u32>(params.denom) | 0x7fc00000u;
+      y[i] = {{ outScalar }}(bitcast<f32>(nan_bits));
+{% endif %}
+      continue;
+    }
+    let period = u32(params.denom);
+    // w = a0 - a1*cos(2*pi*turns) [+ a2*cos(4*pi*turns)], fused single-rounding.
+    var w = fma(-({{ source.a1 }}), cos_turns(i, period), {{ source.a0 }});
+    y[i] = {{ outScalar }}(w);
+  }
+}