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.
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:
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> 0nguard discards — leavingsamplesat zero. Detection:read()returnssamples: 0despite the pass clearly running. Fix: serve withCOOP/COEPheaders 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. GPUValidationErroratcopyBufferToBuffer. Removing themapStateguard lets the copy target a buffer still mapped from the previous frame’smapAsync. Detection: a synchronous validation error every few frames, correlated with readback latency. Fix: keep thecanReadgate 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 —
mapAsyncrejects and, unguarded, would leave an unhandled rejection. Detection:mapAsyncthrows and samples stop advancing. Fix: thetry/catchinreadBackswallows 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: thedelta <= 0n || delta > MAX_NSwindow 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:
emaMsstarts high and decays over the first second of profiling. Fix: warm the pipeline with a throwaway dispatch before wiring the timer, or ignore samples untilread().samplesclears 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.
Related
- Frame Profiling with WebGPU Timestamp Queries — the full lifecycle and multi-pass attribution this single-pass timer specializes.
- Finding Render Pass Bottlenecks with GPU Timestamps — the same instrument applied across render passes and draw batches.
- Workgroup Occupancy Optimization for Spatial Kernels — what to tune once the timer flags a compute pass as the frame’s cost center.
- Feature-Detecting timestamp-query Support — the guard that decides whether this timer instruments or no-ops.