mapAsync() showed the mechanics: copy into a MAP_READ | COPY_DST buffer, await mapAsync(), copy the data out, unmap(). For compute work that runs every frame, three habits keep read-backs from stalling the page:
Never await in the frame loop. Handle the data in mapAsync().then(), a frame or more later.
Keep a small ring of staging buffers. A buffer is busy from mapAsync() until unmap(), so take one whose mapState is 'unmapped' and skip the read when none is free.
Read only what JavaScript needs, such as a count, and let render pipelines consume the rest on the GPU, as Culled Shelf did with its indirect arguments.
A test page read a GPU frame counter for 120 animation frames through a ring of three staging buffers. In two runs every read was correct and none skipped; the median latency was 0.6 ms, but the worst were 17.9 and 31.3 ms, which is why the loop must not wait. mapAsync(mode, offset, size) and getMappedRange(offset, size) handle part of a buffer, and texture read-backs pad bytesPerRow to a multiple of 256 (Buffer-Texture Copies).
<!doctype html>
<style>
body { margin: 0; background: #f7f4ee; font: 14px system-ui, sans-serif; }
.stage { position: relative; width: 100%; max-width: 600px; }
.stage canvas { display: block; width: 100%; }
.stage canvas + canvas { position: absolute; inset: 0; pointer-events: none; }
</style>
<div class="stage">
<canvas id="view" width="600" height="330"></canvas>
<canvas id="labels" width="600" height="330"></canvas>
</div>
<script>
const canvas = document.getElementById('view');
const ink = document.getElementById('labels').getContext('2d');
function showMessage(text) { // 2D fallback when WebGPU is missing
const ctx = canvas.getContext('2d');
ctx.fillStyle = '#fbeaea'; ctx.fillRect(0, 0, canvas.width, canvas.height);
ctx.fillStyle = '#8a2b2b'; ctx.font = '18px system-ui, sans-serif'; ctx.textAlign = 'center';
ctx.fillText(text, canvas.width / 2, canvas.height / 2);
}
const code = /* wgsl */ `
@group(0) @binding(0) var<storage, read_write> counter: array<u32, 2>; // frame count, books sold
@compute @workgroup_size(64) fn tick(@builtin(global_invocation_id) id: vec3u) {
if (id.x == 0) { counter[0] += 1; }
if (id.x % 7 == counter[0] % 7) { counter[1] += 1; } // some GPU-side work to report
}
@group(0) @binding(0) var<storage> shown: array<u32, 2>;
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> @builtin(position) vec4f {
let q = vec2f(f32(v & 1), f32(v >> 1));
let a = f32(shown[0]) * 0.05 + f32(i) * 1.047; // the GPU counter spins the books
let c = vec2f(cos(a), sin(a)) * 0.35 + vec2f(-0.5, 0.05);
return vec4f(c + (q - 0.5) * vec2f(0.06, 0.2), 0, 1);
}
@fragment fn fs() -> @location(0) vec4f { return vec4f(0.08, 0.40, 0.75, 1); }`;
async function main() {
const adapter = await navigator.gpu?.requestAdapter();
if (!adapter) return showMessage('WebGPU is not available in this browser');
const device = await adapter.requestDevice();
const context = canvas.getContext('webgpu');
const format = navigator.gpu.getPreferredCanvasFormat();
context.configure({ device, format });
const B = GPUBufferUsage;
const counter = device.createBuffer({ size: 8, usage: B.STORAGE | B.COPY_SRC });
const ring = [0, 1, 2].map(() => device.createBuffer({ size: 8, usage: B.MAP_READ | B.COPY_DST }));
const module = device.createShaderModule({ code });
const compute = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'tick' } });
const render = device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
vertex: { module, entryPoint: 'vs' }, fragment: { module, entryPoint: 'fs', targets: [{ format }] } });
const cGroup = device.createBindGroup({ layout: compute.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: counter } }] });
const rGroup = device.createBindGroup({ layout: render.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: counter } }] });
const latencies = [];
let reads = 0, skipped = 0, last = [0, 0], jsFrames = 0;
function frame() {
jsFrames++;
const encoder = device.createCommandEncoder();
const cp = encoder.beginComputePass();
cp.setPipeline(compute); cp.setBindGroup(0, cGroup); cp.dispatchWorkgroups(1); cp.end();
const pass = encoder.beginRenderPass({ colorAttachments: [{ view: context.getCurrentTexture().createView(),
clearValue: [0.97, 0.96, 0.93, 1], loadOp: 'clear', storeOp: 'store' }] });
pass.setPipeline(render); pass.setBindGroup(0, rGroup); pass.draw(4, 6); pass.end();
const staging = ring.find((b) => b.mapState === 'unmapped'); // busy from mapAsync() until unmap()
if (staging) encoder.copyBufferToBuffer(counter, 0, staging, 0, 8);
else skipped++; // none free: skip this read
device.queue.submit([encoder.finish()]);
if (staging) {
const t0 = performance.now();
staging.mapAsync(GPUMapMode.READ).then(() => { // handled a frame or more later
last = [...new Uint32Array(staging.getMappedRange())];
staging.unmap();
latencies.push(performance.now() - t0); reads++;
if (latencies.length > 60) latencies.shift();
});
}
drawStats();
requestAnimationFrame(frame); // never await in the loop
}
function drawStats() {
ink.clearRect(0, 0, 600, 330);
ink.font = '12.5px system-ui, sans-serif'; ink.fillStyle = '#222';
const lines = [`JavaScript frames: ${jsFrames}`, `GPU counter read back: ${last[0]}`, `books sold (GPU): ${last[1]}`,
`reads completed: ${reads}, skipped: ${skipped}`, `ring: ${ring.map((b) => b.mapState).join(', ')}`];
lines.forEach((l, i) => ink.fillText(l, 330, 60 + i * 22));
ink.fillText('read latency per frame (ms), last 60 reads', 330, 190);
const max = Math.max(20, ...latencies);
latencies.forEach((ms, i) => {
const h = ms / max * 100;
ink.fillStyle = ms > 16 ? '#d08a1c' : '#2a8f4d'; ink.fillRect(330 + i * 4, 300 - h, 3, h);
});
ink.fillStyle = '#555'; ink.font = '11px system-ui, sans-serif';
ink.fillText(`median ${median(latencies).toFixed(1)} ms, worst ${Math.max(0, ...latencies).toFixed(1)} ms (so never await)`, 330, 318);
ink.fillStyle = '#222'; ink.font = '12.5px system-ui, sans-serif';
ink.fillText('The books spin from the GPU counter; the numbers', 12, 290); ink.fillText('on the right arrive a frame or more later.', 12, 308);
}
const median = (a) => a.length ? [...a].sort((x, y) => x - y)[a.length >> 1] : 0;
requestAnimationFrame(frame);
}
main();
</script>