feat(experimental) Add reusable WebGPU 2D FFT foundation (#2831)

Author: ibgreen-openai

docs/api-reference/experimental/README.md

@@ -57,7 +57,8 @@ OIT resolve pipelines compose into one render stack.
 
 The [GPU Primitives and Command Graphs guide](/docs/api-reference/experimental/gpu-primitives)
 introduces explicit command scheduling, typed table-backed graph views, hierarchical scan, stable
-compaction, stable key/value sorting, and GPU-written indirect draw commands.
+compaction, stable key/value sorting, bounded two-dimensional complex FFTs, and GPU-written
+indirect draw commands.
 
 ## GPU Simulations
 

docs/api-reference/experimental/gpu-primitives/README.md

@@ -58,8 +58,8 @@ later render pass consume that count without a CPU synchronization point.
 The implementation consists of `GPUCommandGraph`, typed graph data views, `GPUScan`,
 `GPUCompaction`, `GPUMask`, `GPUVisibilityWorkflow`, `GPUHierarchyLayout`, `GPUGraphTraversal`,
 `GPUAncestorProjection`, `GPUSort`, `GPUBatchSort`, `GPUReduction`, `GPUHistogram`, `GPUGridBinning`,
-`GPUGridAggregation`, `GPUGroupAggregation`, `GPUIndexPickingTarget`, `GPUReadbackRing`, and
-`DrawCommandBuffer`. The accompanying hierarchical trace viewer applies these primitives to
+`GPUGridAggregation`, `GPUGroupAggregation`, `GPUFFT2D`, `GPUIndexPickingTarget`,
+`GPUReadbackRing`, and `DrawCommandBuffer`. The accompanying hierarchical trace viewer applies these primitives to
 process and thread collapse, source and topology filtering, dependency focusing, visible-parent
 projection, GPU picking, activity histograms, and indirect span and edge rendering over up to
 four million spans. The sort and data-analysis examples demonstrate independent composable
@@ -1585,6 +1585,7 @@ close enough to WebGPU that developers can reason about cost, ordering, and owne
 - [`GPUGraphTraversal`](/docs/api-reference/experimental/gpu-primitives/gpu-graph-traversal)
 - [`GPUAncestorProjection`](/docs/api-reference/experimental/gpu-primitives/gpu-ancestor-projection)
 - [`GPUSort`](/docs/api-reference/experimental/gpu-primitives/gpu-sort)
+- [`GPUFFT2D`](/docs/api-reference/experimental/gpu-primitives/gpu-fft2d)
 - [`GPUReduction`](/docs/api-reference/experimental/gpu-primitives/gpu-reduction)
 - [`GPUHistogram`](/docs/api-reference/experimental/gpu-primitives/gpu-histogram)
 - [`GPUGridBinning`](/docs/api-reference/experimental/gpu-primitives/gpu-grid-binning)

docs/api-reference/experimental/gpu-primitives/gpu-fft2d.md

@@ -0,0 +1,162 @@
+import {GPUPrimitivesDocsTabs} from '@site/src/components/docs/gpu-primitives-docs-tabs';
+
+# GPUFFT2D
+
+<GPUPrimitivesDocsTabs active="fft2d" />
+
+## Overview
+
+`GPUFFT2D` records a bounded, out-of-place two-dimensional complex fast Fourier transform on
+WebGPU. It accepts caller-owned row-major storage buffers, owns one equally sized scratch buffer,
+and records every bit-reversal and radix-2 butterfly pass onto the application's
+`CommandEncoder`. It never submits commands or reads values back to the CPU.
+
+The initial implementation targets reusable simulation and signal-processing foundations such as
+spectral oceans, frequency-domain filters, convolution, and procedural fields. It deliberately
+does not own textures, convert real-valued inputs, select padding dimensions, or hide command
+submission.
+
+## Usage
+
+Each complex value occupies two consecutive `float32` components: real followed by imaginary.
+Values are row-major, so the complete field contains `width * height * 2` floats.
+
+```ts
+import {Buffer} from '@luma.gl/core';
+import {GPUFFT2D} from '@luma.gl/experimental';
+
+const width = 256;
+const height = 256;
+const complexByteLength = width * height * 2 * Float32Array.BYTES_PER_ELEMENT;
+
+const inputBuffer = device.createBuffer({
+  data: initialComplexValues,
+  usage: Buffer.STORAGE | Buffer.COPY_DST
+});
+const frequencyBuffer = device.createBuffer({
+  byteLength: complexByteLength,
+  usage: Buffer.STORAGE
+});
+const reconstructedBuffer = device.createBuffer({
+  byteLength: complexByteLength,
+  usage: Buffer.STORAGE
+});
+const transform = new GPUFFT2D(device, {width, height});
+
+const commandEncoder = device.createCommandEncoder({id: 'spectral-step'});
+transform.encode(commandEncoder, {
+  inputBuffer,
+  outputBuffer: frequencyBuffer,
+  direction: 'forward'
+});
+transform.encode(commandEncoder, {
+  inputBuffer: frequencyBuffer,
+  outputBuffer: reconstructedBuffer,
+  direction: 'inverse'
+});
+device.submit(commandEncoder.finish());
+```
+
+The two calls above compose in one command buffer. The second transform observes the first
+transform's output through ordinary WebGPU command ordering; no intermediate submission or CPU
+synchronization is required.
+
+## Constructor
+
+### `new GPUFFT2D(device, props)`
+
+```ts
+type GPUFFT2DProps = {
+  id?: string;
+  width: number;
+  height: number;
+};
+```
+
+`width` and `height` must each be powers of two from 2 through 2048. Rectangular transforms are
+supported. The bound keeps allocation and dispatch costs predictable: the maximum field contains
+4,194,304 complex values and occupies 32 MiB per complex buffer.
+
+Construction allocates one field-sized scratch buffer, one compute pipeline, and two immutable
+32-byte parameter buffers per pass so forward and inverse encodings never race through rewritten
+uniforms. Input and output storage remain caller-owned.
+
+## Encoding
+
+### `encode(commandEncoder, options): Buffer`
+
+```ts
+type GPUFFT2DEncodeOptions = {
+  inputBuffer: Buffer;
+  outputBuffer: Buffer;
+  direction?: 'forward' | 'inverse';
+};
+```
+
+Both buffers must belong to the transform's device, declare `Buffer.STORAGE`, and contain at least
+`stats.complexBufferByteLength` bytes. They must be separate allocations; the source is never
+modified. `encode()` returns `outputBuffer` for convenient downstream binding.
+
+The normalization convention is:
+
+- `forward`: negative complex exponent and no normalization;
+- `inverse`: positive complex exponent and division by `width * height` on the final pass.
+
+The transform first bit-reverses and evaluates every row, then does the same for every column.
+Passes ping-pong between the class-owned scratch field and the caller's output so the final pass
+always lands in `outputBuffer`.
+
+## Support query
+
+### `getGPUFFT2DSupport(device, props): GPUFFT2DSupport`
+
+The support query validates dimensions before allocation and reports WebGPU compute, workgroup,
+dispatch, storage-binding, and buffer-size limits. A valid plan is included in `stats` even when a
+device limit prevents execution.
+
+```ts
+const support = getGPUFFT2DSupport(device, {width: 512, height: 256});
+if (!support.supported) {
+  console.warn(support.reason);
+}
+```
+
+## Statistics
+
+`transform.stats` is an immutable `GPUFFT2DStats` object:
+
+| Field | Meaning |
+| --- | --- |
+| `width`, `height`, `elementCount` | Logical complex-field dimensions and value count. |
+| `complexBufferByteLength` | Minimum byte length of input, output, and scratch fields. |
+| `horizontalStageCount`, `verticalStageCount` | Radix-2 butterfly stages per axis. |
+| `passCount`, `dispatchCountPerEncode` | Two bit-reversal passes plus all butterfly stages. |
+| `workgroupSize`, `workgroupCount` | Fixed 8-by-8 invocation tile and dispatch grid. |
+| `scratchBufferByteLength` | Class-owned transform scratch. |
+| `parameterBufferCount`, `parameterBufferByteLength` | Immutable forward/inverse pass metadata. |
+
+`makeGPUFFT2DStats(width, height)` computes the same plan without a device or GPU allocation.
+
+## Ownership and lifecycle
+
+`GPUFFT2D` owns only its compute pipeline, scratch buffer, and parameter buffers. The caller owns
+input buffers, output buffers, command encoders, submission, and any optional readback. Destroying
+the transform releases only class-owned resources and is idempotent. Previously supplied caller
+buffers remain valid.
+
+The same instance may encode more than one ordered transform into a command encoder. Do not encode
+the same instance concurrently into command buffers that may execute simultaneously because those
+encodings share its scratch field. Use separate instances when independent queues or overlapping
+submissions need the same dimensions.
+
+## Current limits
+
+- WebGPU only; no WebGL fallback.
+- Power-of-two dimensions from 2 through 2048.
+- Packed row-major complex `float32` fields only.
+- Out-of-place input and output only.
+- One shared scratch field per instance; no concurrent execution contract.
+- No hidden padding, real-to-complex packing, texture conversion, submission, or readback.
+
+These constraints keep the primitive small and predictable while allowing higher-level systems to
+define spectrum generation, texture outputs, cascade policy, and scheduling separately.

docs/table-of-contents.json

@@ -220,6 +220,7 @@
               "api-reference/experimental/gpu-primitives/gpu-graph-traversal",
               "api-reference/experimental/gpu-primitives/gpu-ancestor-projection",
               "api-reference/experimental/gpu-primitives/gpu-sort",
+              "api-reference/experimental/gpu-primitives/gpu-fft2d",
               "api-reference/experimental/gpu-primitives/gpu-reduction",
               "api-reference/experimental/gpu-primitives/gpu-histogram",
               "api-reference/experimental/gpu-primitives/gpu-grid-binning",
@@ -315,6 +316,7 @@
               "api-reference/experimental/gpu-primitives/gpu-graph-traversal",
               "api-reference/experimental/gpu-primitives/gpu-ancestor-projection",
               "api-reference/experimental/gpu-primitives/gpu-sort",
+              "api-reference/experimental/gpu-primitives/gpu-fft2d",
               "api-reference/experimental/gpu-primitives/gpu-reduction",
               "api-reference/experimental/gpu-primitives/gpu-histogram",
               "api-reference/experimental/gpu-primitives/gpu-grid-binning",

docs/whats-new.md

@@ -34,6 +34,7 @@ Target Release Date: Q3, 2026
 - **GPU scan, compaction, and indirect drawing** - Typed graph views compose hierarchical `uint32` scan, stable ID compaction, and GPU-written `DrawCommandBuffer` instance counts. Scan and compaction accept fixed-width `GPUVector` imports as one logical sequence while preserving chunk topology. The [GPU Trace Viewer](/examples/experimental/gpu-trace-viewer) demonstrates the path over up to four million spans, while [GPU Frustum Culling](/examples/experimental/gpu-frustum-culling) applies it to indexed indirect rendering of a 3D instance field.
 - **GPU trace manipulation primitives** - `GPUMask` composes chunk-preserving selection predicates; `GPUHierarchyLayout` computes scan-based process and thread expansion; `GPUGraphTraversal` expands bounded, cycle-safe CSR dependency frontiers; and `GPUAncestorProjection` reconnects hidden spans to their nearest visible canonical parent. The [GPU Hierarchical Trace Viewer](/examples/experimental/gpu-trace-viewer) applies all four to live hierarchy controls, topology filters, dependency focusing, GPU picking, projected indirect edges, and collapsed-process activity.
 - **Graph-native GPU sort** - `GPUSort` stably orders one paired packed `uint32` domain, while `GPUBatchSort` independently orders aligned GPU vector chunks without hidden packing or lost batch boundaries. Bitonic or binary LSD radix selection occurs per work unit. The [GPU Sort example](/examples/experimental/gpu-sort) contrasts packed global order with preserved Arrow batches and exposes graph compilation and transient reuse.
+- **Reusable 2D GPU FFT** - [`GPUFFT2D`](/docs/api-reference/experimental/gpu-primitives/gpu-fft2d) records bounded power-of-two complex transforms into caller-owned WebGPU command encoders. Forward and normalized inverse passes share one explicit scratch field without hidden submission or readback, providing a reusable spectral-simulation and signal-processing foundation.
 - **Graph-native GPU data analysis** - `GPUReduction`, `GPUHistogram`, `GPUGridBinning`, `GPUGridAggregation`, and `GPUGroupAggregation` add deterministic scalar aggregates, equal-width or irregular-edge histogram counts, filtered categorical counts and floating-point statistics, row-major spatial counts, and weighted floating-point sum/min/max/mean cell statistics. GPU-resident histogram edges and group-selection masks can change between encodings without CPU readback or graph recompilation. Analysis operations initialize once and accumulate fixed-width vector chunks without packing. The [GPU Data Analysis example](/examples/experimental/gpu-data-analysis) composes the operations without hidden submission or readback.
 - **Semantic G-buffer targets** - `GBuffer` owns WebGPU MRT scene color, normal-roughness, velocity, and depth targets plus named extra channels, then exposes the standard depth, normal, and velocity bindings consumed by screen-space effect pipelines.
 - **Composable deferred lighting** - `deferredLighting` resolves Cook-Torrance opaque lighting from G-buffer material channels, reconstructed depth, one directional light, and a fixed-capacity WebGPU point-light storage buffer. The [Deferred Illumination Lab](/examples/experimental/deferred-rendering) exposes the material channels and animated lights live.

modules/experimental/src/gpu-primitives/gpu-fft2d-shaders.ts

@@ -0,0 +1,97 @@
+// luma.gl
+// SPDX-License-Identifier: MIT
+// Copyright (c) vis.gl contributors
+
+/** Number of invocations along each dimension of one GPUFFT2D workgroup. */
+export const GPU_FFT2D_WORKGROUP_DIMENSION = 8;
+
+/** Byte length of the uniform block consumed by every GPUFFT2D pass. */
+export const GPU_FFT2D_PARAMETER_BYTE_LENGTH = 32;
+
+/** Shared bit-reversal and butterfly kernel used by every GPUFFT2D pass. */
+export const GPU_FFT2D_SHADER = /* wgsl */ `\
+struct GPUFFT2DParameters {
+  width: u32,
+  height: u32,
+  axis: u32,
+  passKind: u32,
+  transformSize: u32,
+  stage: u32,
+  directionSign: f32,
+  normalizationScale: f32,
+};
+
+@group(0) @binding(0) var<storage, read> inputValues: array<vec2f>;
+@group(0) @binding(1) var<storage, read_write> outputValues: array<vec2f>;
+@group(0) @binding(2) var<uniform> parameters: GPUFFT2DParameters;
+
+const GPU_FFT2D_PI: f32 = 3.14159265358979323846;
+
+fn reverseLowBits(value: u32, bitCount: u32) -> u32 {
+  var source = value;
+  var reversed = 0u;
+  for (var bitIndex = 0u; bitIndex < bitCount; bitIndex++) {
+    reversed = (reversed << 1u) | (source & 1u);
+    source = source >> 1u;
+  }
+  return reversed;
+}
+
+fn getLinearIndex(xCoordinate: u32, yCoordinate: u32) -> u32 {
+  return yCoordinate * parameters.width + xCoordinate;
+}
+
+fn multiplyComplex(left: vec2f, right: vec2f) -> vec2f {
+  return vec2f(
+    left.x * right.x - left.y * right.y,
+    left.x * right.y + left.y * right.x
+  );
+}
+
+@compute @workgroup_size(${GPU_FFT2D_WORKGROUP_DIMENSION}, ${GPU_FFT2D_WORKGROUP_DIMENSION}, 1)
+fn main(@builtin(global_invocation_id) globalIdentifier: vec3u) {
+  if (globalIdentifier.x >= parameters.width || globalIdentifier.y >= parameters.height) {
+    return;
+  }
+
+  let horizontal = parameters.axis == 0u;
+  let coordinate = select(globalIdentifier.y, globalIdentifier.x, horizontal);
+  var sourceCoordinate = coordinate;
+
+  if (parameters.passKind == 0u) {
+    sourceCoordinate = reverseLowBits(coordinate, parameters.stage);
+  } else {
+    let butterflySpan = 1u << parameters.stage;
+    let butterflyHalfSpan = butterflySpan >> 1u;
+    let butterflyOffset = coordinate & (butterflySpan - 1u);
+    let twiddleIndex = butterflyOffset & (butterflyHalfSpan - 1u);
+    let butterflyStart = coordinate - butterflyOffset;
+    let firstCoordinate = butterflyStart + twiddleIndex;
+    let secondCoordinate = firstCoordinate + butterflyHalfSpan;
+
+    let firstX = select(globalIdentifier.x, firstCoordinate, horizontal);
+    let firstY = select(firstCoordinate, globalIdentifier.y, horizontal);
+    let secondX = select(globalIdentifier.x, secondCoordinate, horizontal);
+    let secondY = select(secondCoordinate, globalIdentifier.y, horizontal);
+    let firstValue = inputValues[getLinearIndex(firstX, firstY)];
+    let secondValue = inputValues[getLinearIndex(secondX, secondY)];
+    let angle = parameters.directionSign * 2.0 * GPU_FFT2D_PI *
+      f32(twiddleIndex) / f32(butterflySpan);
+    let twiddle = vec2f(cos(angle), sin(angle));
+    let rotatedSecondValue = multiplyComplex(secondValue, twiddle);
+    let butterflyValue = select(
+      firstValue + rotatedSecondValue,
+      firstValue - rotatedSecondValue,
+      butterflyOffset >= butterflyHalfSpan
+    );
+    outputValues[getLinearIndex(globalIdentifier.x, globalIdentifier.y)] =
+      butterflyValue * parameters.normalizationScale;
+    return;
+  }
+
+  let sourceX = select(globalIdentifier.x, sourceCoordinate, horizontal);
+  let sourceY = select(sourceCoordinate, globalIdentifier.y, horizontal);
+  outputValues[getLinearIndex(globalIdentifier.x, globalIdentifier.y)] =
+    inputValues[getLinearIndex(sourceX, sourceY)] * parameters.normalizationScale;
+}
+`;