Measuring Compute Pass Duration with Timestamp Query

The specific question here is how long one spatial compute dispatch actually runs on the GPU — the scatter that bins ten million GPS pings, or the kernel that culls a vector-tile set — not how long submit() took to return on the CPU. performance.now() around the dispatch measures only command recording and returns before the GPU has touched the buffers, so it reports single-digit microseconds for a pass that costs eight milliseconds of silicon time. The correct instrument is a two-slot timestamp query that stamps the GPU’s nanosecond counter at the pass’s begin and end, resolved and read back after submission. This page is a single self-contained class you can drop around any compute pass to get a stable millisecond reading; the broader lifecycle and per-pass attribution live in the parent reference on frame profiling with WebGPU timestamp queries.

Runnable reference implementation

The ComputePassTimer below owns a two-query set, a resolve buffer, and a mappable readback buffer, all allocated once. timeDispatch wraps a caller-supplied record callback in a timestampWrites-instrumented compute pass, resolves, submits, and reads back asynchronously. It exponentially smooths the reading so a single disjoint or quantized sample does not jump the number, and it guards every WebGPU precondition — feature presence, buffer map state, disjoint deltas — so it degrades cleanly instead of throwing.

typescript
export interface TimerSample {
  lastMs: number;   // most recent valid GPU duration
  emaMs: number;    // exponential moving average, resistant to outliers
  samples: number;  // count of valid samples folded in
}

export class ComputePassTimer {
  private readonly querySet: GPUQuerySet;
  private readonly resolveBuffer: GPUBuffer;
  private readonly readbackBuffer: GPUBuffer;
  private readonly supported: boolean;
  private ema = 0;
  private count = 0;
  private last = 0;

  // Discard any delta outside [0, MAX_NS]; a disjoint counter reads negative
  // (as bigint underflow) or absurdly large across a GPU clock change.
  private static readonly MAX_NS = 100_000_000n; // 100 ms ceiling
  private static readonly ALPHA = 0.2;           // EMA weight on the newest sample

  constructor(private readonly device: GPUDevice) {
    this.supported = device.features.has("timestamp-query");
    if (!this.supported) {
      // Build inert placeholders so callers need no branching; timeDispatch no-ops.
      this.querySet = undefined as unknown as GPUQuerySet;
      this.resolveBuffer = undefined as unknown as GPUBuffer;
      this.readbackBuffer = undefined as unknown as GPUBuffer;
      return;
    }
    this.querySet = device.createQuerySet({ type: "timestamp", count: 2 });
    this.resolveBuffer = device.createBuffer({
      size: 2 * 8, // two BigUint64 counters
      usage: GPUBufferUsage.QUERY_RESOLVE | GPUBufferUsage.COPY_SRC,
    });
    this.readbackBuffer = device.createBuffer({
      size: 2 * 8,
      usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
    });
  }

  /**
   * Record and submit `body` inside a timed compute pass. The callback receives
   * the instrumented pass; it must set its pipeline, bind groups, and dispatch,
   * but must NOT call pass.end() — the timer owns the pass lifecycle.
   */
  timeDispatch(body: (pass: GPUComputePassEncoder) => void): void {
    const encoder = this.device.createCommandEncoder();

    if (!this.supported) {
      // Fallback: run the work untimed so the frame still renders.
      const pass = encoder.beginComputePass();
      body(pass);
      pass.end();
      this.device.queue.submit([encoder.finish()]);
      return;
    }

    const pass = encoder.beginComputePass({
      timestampWrites: {
        querySet: this.querySet,
        beginningOfPassWriteIndex: 0,
        endOfPassWriteIndex: 1,
      },
    });
    body(pass);
    pass.end();

    // Resolve the opaque counters into a readable buffer.
    encoder.resolveQuerySet(this.querySet, 0, 2, this.resolveBuffer, 0);

    // Only stage to the readback buffer when it is free; skipping a frame is fine.
    const canRead = this.readbackBuffer.mapState === "unmapped";
    if (canRead) {
      encoder.copyBufferToBuffer(this.resolveBuffer, 0, this.readbackBuffer, 0, 2 * 8);
    }

    this.device.queue.submit([encoder.finish()]);
    if (canRead) void this.readBack();
  }

  private async readBack(): Promise<void> {
    try {
      await this.readbackBuffer.mapAsync(GPUMapMode.READ);
    } catch {
      return; // device lost or buffer destroyed mid-flight; ignore this sample
    }
    // Copy out before unmap detaches the backing memory.
    const ns = new BigUint64Array(this.readbackBuffer.getMappedRange().slice(0));
    this.readbackBuffer.unmap();

    const delta = ns[1] - ns[0];
    if (delta <= 0n || delta > ComputePassTimer.MAX_NS) return; // disjoint / quantized

    const ms = Number(delta) / 1e6;
    this.last = ms;
    this.ema = this.count === 0 ? ms
      : ComputePassTimer.ALPHA * ms + (1 - ComputePassTimer.ALPHA) * this.ema;
    this.count++;
  }

  read(): TimerSample {
    return { lastMs: this.last, emaMs: this.ema, samples: this.count };
  }

  destroy(): void {
    if (!this.supported) return;
    this.querySet.destroy();
    this.resolveBuffer.destroy();
    this.readbackBuffer.destroy();
  }
}

Wiring it around a real spatial scatter pass is three lines. The timer never touches your pipeline or bind groups — it only owns the pass boundary where the stamps land:

typescript
const timer = new ComputePassTimer(device);

function frame(pointCount: number): void {
  timer.timeDispatch((pass) => {
    pass.setPipeline(scatterPipeline);
    pass.setBindGroup(0, scatterBindGroup);
    pass.dispatchWorkgroups(Math.ceil(pointCount / 256)); // one invocation per point
  });

  const { emaMs, lastMs, samples } = timer.read();
  if (samples > 0) hud.set(`scatter ${emaMs.toFixed(2)} ms (last ${lastMs.toFixed(2)})`);
  requestAnimationFrame(() => frame(pointCount));
}

Because readBack trails the dispatch it measured by a submission or two, the HUD shows a rolling average rather than the in-flight frame — the right behaviour for catching an intermittent spike when a pan pushes the point count up. The exponential average keeps that number readable while lastMs still exposes the raw spike.

Interpreting the reading

A timestamp delta measures GPU execution, not the wall time you feel as jank, and the two diverge in ways worth knowing before you act on the number. The first sample after a pipeline is created is almost always an outlier: the driver compiles and warms the shader on first use, and the initial dispatch pays that one-off cost inside the timed span. Discard the first few readings, or warm the pipeline with a throwaway dispatch before the timer goes live, so the EMA reflects steady-state cost rather than compilation.

A single-pass delta also hides why a pass is slow. A scatter that reads 8 ms could be arithmetic-bound, memory-bandwidth-bound on the ten-million-point payload, or serialized on atomic contention into a few hot cells — the timer reports the total, not the cause. Treat it as the signal that tells you which pass to open, then reach for workgroup occupancy optimization for spatial kernels to find the mechanism. Conversely, if the GPU delta is small but the frame still stutters, the cost is on the CPU recording side or in readback stalls, not in the pass this timer measures — a performance.now() bracket around submit will confirm it. Comparing the GPU delta against that CPU bracket is the fastest way to classify a frame as GPU-bound or record-bound.

Parameter reference

Every tunable value in the implementation, with guidance for spatial compute workloads. These tables scroll horizontally on narrow viewports.

Parameter Value Guidance
querySet.count 2 Exactly one begin + one end slot for a single pass; do not oversize.
Resolve buffer usage QUERY_RESOLVE | COPY_SRC resolveQuerySet requires the first flag; COPY_SRC stages to the mappable buffer.
Readback buffer usage COPY_DST | MAP_READ The only combination that permits mapAsync; a resolve buffer cannot be mapped.
ALPHA (EMA weight) 0.2 Lower (0.1) for a steadier HUD, higher (0.4) to track fast regressions during zoom.
MAX_NS ceiling 100_000_000n (100 ms) Rejects disjoint counters; set just above your worst legitimate pass, not higher.
Delta lower bound > 0n Rejects quantized-to-zero and non-monotonic samples in one check.
Dispatch divisor 256 Must match @workgroup_size(256) in the WGSL; a mismatch under-dispatches and skews timing.
Sample cadence every frame, read when free The mapState guard auto-throttles readback; never await on the render path.

The two guards — mapState === "unmapped" before the copy and the [0, MAX_NS] window on the delta — are what make the timer safe to leave running: the first prevents a validation error when the previous readback is still in flight, the second rejects the disjoint and quantized samples the spec explicitly permits. For authoritative wording on resolveQuerySet and buffer mapping, see the W3C WebGPU specification and the MDN GPUBuffer.mapAsync reference.

Failure modes

  • Timer reads 0.00 ms forever. The document is not cross-origin isolated, so the counter’s low bits are zeroed and a fast pass differences to 0n, which the > 0n guard discards — leaving samples at zero. Detection: read() returns samples: 0 despite the pass clearly running. Fix: serve with COOP/COEP headers for fine resolution, or, if profiling a genuinely tiny pass, batch several dispatches inside one timed pass so the measured span clears the quantization bucket.
  • GPUValidationError at copyBufferToBuffer. Removing the mapState guard lets the copy target a buffer still mapped from the previous frame’s mapAsync. Detection: a synchronous validation error every few frames, correlated with readback latency. Fix: keep the canRead gate so the copy is recorded only when the readback buffer is "unmapped".
  • HUD frozen after a device loss. If the device is lost mid-pan — an oversized dispatch tripping the driver watchdog — mapAsync rejects and, unguarded, would leave an unhandled rejection. Detection: mapAsync throws and samples stop advancing. Fix: the try/catch in readBack swallows the rejection; recreate the device and timer through the browser support fallback routing strategies.
  • Number jumps to a negative or multi-second value. A disjoint counter across a GPU clock or power-state change makes ns[1] - ns[0] underflow or spike. Detection: without the ceiling, the EMA lurches wildly for one sample. Fix: the delta <= 0n || delta > MAX_NS window drops it; never fold a rejected sample into the average.
  • First reading is wildly high, then settles. The initial dispatch after pipeline creation pays a one-off shader-compilation cost inside the timed span, so the first EMA value can be many times the steady state. Detection: emaMs starts high and decays over the first second of profiling. Fix: warm the pipeline with a throwaway dispatch before wiring the timer, or ignore samples until read().samples clears a small threshold.

Backend / Python interop note

When the same spatial pass runs headless under wgpu-py on a Python rendering server, the timing surface is identical: the bindings expose create_query_set, resolve_query_set, and buffer mapping with the same two-slot pattern, so a ComputePassTimer equivalent ports directly. Feed the server-side profiler the same fixed workload — a pinned GeoParquet tile of known point count — on every run so the millisecond figures are comparable across deploys, and log the EMA rather than raw samples so a single disjoint counter on a shared cloud GPU does not poison a regression alert.

Up: Frame Profiling with WebGPU Timestamp Queries