Pipelining Compute Dispatches with Double Buffering

The stall this page removes is the serialization between a clustering compute pass and the work that consumes its output. A naive loop dispatches the point-in-polygon clustering kernel, then calls mapAsync on the result buffer to read cluster labels back to the CPU, then renders — and every stage waits for the previous one to finish because they all touch the same buffer. The GPU sits idle during the mapAsync latency, the CPU sits idle during the dispatch, and the frame time becomes the sum of compute, readback, and render instead of their maximum. At ten million points a single clustering pass plus its readback can blow a 16 ms budget on its own, so the map judders on every pan even though no single stage is slow. The fix is to give the pipeline two independent sets of buffers and alternate between them: while the GPU computes into set B for frame N+1, the CPU reads and the renderer draws from set A for frame N. The buffer dependency that forced serialization disappears because the two frames no longer share memory, and steady-state frame time collapses to the cost of the single longest stage.

This is a scheduling technique layered on top of async dispatch patterns for spatial clustering; it assumes the clustering kernel itself — the WGSL kernel for point-in-polygon clustering — already runs correctly, and it treats that kernel as a black box whose only requirement is that its input and output buffers can be duplicated. Because a compute pass is a distinct GPU program that can execute while the renderer paints, the overlap here builds directly on the compute versus render pipeline fundamentals.

Runnable reference implementation

The orchestrator below holds two FrameSet records, each carrying its own input, output, and staging buffers plus its own bind group. Every frame it computes into the back set, then swaps, so the set the renderer and readback read is always the one the GPU finished last frame. The clustering shader is unchanged from a single-buffered pipeline — only the buffers it binds differ per set.

interface FrameSet {
  input: GPUBuffer;    // packed vec2<f32> positions for this frame
  output: GPUBuffer;   // cluster labels written by the kernel (STORAGE | COPY_SRC)
  staging: GPUBuffer;  // MAP_READ | COPY_DST target for readback
  bindGroup: GPUBindGroup;
  reading: boolean;    // true while staging is mapped; do not reuse until cleared
}

class DoubleBufferedClustering {
  private sets: [FrameSet, FrameSet];
  private front = 0;                    // set consumed this frame
  private readonly labelBytes: number;  // pointCount * 4 (one u32 label per point)

  constructor(
    private device: GPUDevice,
    private pipeline: GPUComputePipeline,
    private pointCount: number,
    makeBindGroup: (input: GPUBuffer, output: GPUBuffer) => GPUBindGroup,
  ) {
    this.labelBytes = pointCount * 4;
    const make = (): FrameSet => {
      const input = device.createBuffer({
        size: pointCount * 8,            // vec2<f32> = 8 bytes/point
        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
      });
      const output = device.createBuffer({
        size: this.labelBytes,
        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
      });
      const staging = device.createBuffer({
        size: this.labelBytes,
        usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
      });
      return { input, output, staging, bindGroup: makeBindGroup(input, output), reading: false };
    };
    this.sets = [make(), make()];
  }

  // Advance one frame. Returns the labels resolved from the PREVIOUS frame's
  // back set, so the caller always renders data that is one frame old but never
  // waits on a fresh dispatch.
  async frame(positions: Float32Array): Promise<Uint32Array | null> {
    const back = this.sets[this.front ^ 1];   // the set we compute into now

    // If the back set's staging is still mapped from an earlier frame, skip its
    // upload this turn rather than stall; ringDepth of 2 makes this rare.
    if (back.reading) return null;

    this.device.queue.writeBuffer(back.input, 0, positions);

    const enc = this.device.createCommandEncoder();
    const pass = enc.beginComputePass();
    pass.setPipeline(this.pipeline);
    pass.setBindGroup(0, back.bindGroup);
    pass.dispatchWorkgroups(Math.ceil(this.pointCount / 256));
    pass.end();
    // Copy this frame's fresh labels into its own staging buffer for readback.
    enc.copyBufferToBuffer(back.output, 0, back.staging, 0, this.labelBytes);
    this.device.queue.submit([enc.finish()]);

    // Swap: the set just dispatched becomes the front the NEXT call consumes.
    this.front ^= 1;

    // Resolve the front set — the one computed last frame. Its GPU work has had
    // a full frame to finish, so mapAsync typically resolves without blocking.
    const front = this.sets[this.front];
    front.reading = true;
    await front.staging.mapAsync(GPUMapMode.READ);
    const labels = new Uint32Array(front.staging.getMappedRange().slice(0));
    front.staging.unmap();
    front.reading = false;
    return labels;
  }

  destroy(): void {
    for (const s of this.sets) {
      s.input.destroy();
      s.output.destroy();
      s.staging.destroy();
    }
  }
}

The kernel the orchestrator dispatches is buffer-set agnostic: it reads whichever positions and writes whichever labels the bind group supplies, so the identical pipeline serves both set A and set B. Nothing in the WGSL knows or cares which frame it is computing — the double buffering lives entirely in the host-side bind-group selection.

@group(0) @binding(0) var<storage, read>       positions: array<vec2<f32>>; // this set's input
@group(0) @binding(1) var<storage, read_write> labels:    array<u32>;        // this set's output
@group(0) @binding(2) var<uniform>             grid:      vec4<f32>;          // origin.xy, inv_cell.xy

const GRID_W: u32 = 1024u;

// Assign each point a coarse cluster label by hashing it to a grid cell. The
// same kernel runs for whichever buffer set the host bind group names, so the
// overlap is purely a scheduling decision made in TypeScript.
@compute @workgroup_size(256)
fn cluster(@builtin(global_invocation_id) gid: vec3<u32>) {
    let i = gid.x;
    if (i >= arrayLength(&positions)) { return; }  // guard the ragged tail workgroup
    let p = positions[i];
    let cx = u32((p.x - grid.x) * grid.z);         // grid.z = inverse cell width
    let cy = u32((p.y - grid.y) * grid.w);         // grid.w = inverse cell height
    labels[i] = cy * GRID_W + cx;                  // row-major cluster id, written to this set
}

The load-bearing detail is the this.front ^= 1 swap sitting between the submit and the mapAsync. It hands the just-dispatched set forward to the next frame and resolves the set from the prior frame, whose GPU work has already had a full frame’s wall-clock time to complete. That is what makes mapAsync resolve near-instantly instead of blocking: you never read a buffer the GPU is still writing. If your renderer draws directly from the label buffer on the GPU rather than reading labels to the CPU, drop the staging/mapAsync path entirely and simply bind front.output to the render pass — the same swap keeps the read and write frames on different memory.

Parameter reference

Failure modes

• Torn frame from a shared buffer. Reusing one output buffer for both the in-flight compute and the current render lets the renderer sample labels the kernel is still overwriting, producing flickering cluster assignments. Detection: cluster colors shimmer during continuous pan but settle when the map is still. Fix: give each frame set its own output buffer as shown, and only ever read the front set while writing the back set.
• mapAsync on an already-mapped buffer. Calling mapAsync on a staging buffer whose previous map has not been unmapped raises OperationError: buffer is already mapped. Detection: the promise rejects synchronously on a fast pan when frames overtake the readback. Fix: gate reuse on the reading flag as shown, and either skip the frame or add a third buffer set so a mapped buffer is never re-dispatched.
• Unbounded queue growth. Submitting a new dispatch every frame without ever awaiting completion lets the GPU queue grow until VRAM is exhausted and the device is lost. Detection: memory climbs steadily and device.lost fires after sustained interaction. Fix: keep the await front.staging.mapAsync(...) in the loop — it provides natural backpressure — or gate submissions on queue.onSubmittedWorkDone() when rendering GPU-side without readback.
• Stale labels rendered as fresh. Because the front set is always one frame old, wiring its labels to a hit-test or selection that assumes current-frame coordinates mismatches the pointer by one frame of motion. Detection: click targets lag the cursor by one frame during fast pans. Fix: transform the returned labels by the delta between their capture extent and the current extent, or accept one frame of latency for non-interactive layers.

Related

• Async Dispatch Patterns for Spatial Clustering — the broader scheduling model this double-buffering sits inside.
• Writing a WGSL Kernel for Point-in-Polygon Clustering — the clustering kernel whose buffers are duplicated here.
• Frame Profiling with Timestamp Queries — confirm that the overlap actually collapses frame time to the longest stage.
• Compute vs Render Pipeline Fundamentals — why a compute pass can run while the renderer paints.

Up: Async Dispatch Patterns for Spatial Clustering