The point of an indirect buffer is that a compute shader fills it. The usual pattern declares the arguments as a struct with an atomic count, resets the count before each frame, and lets every invocation that keeps an object append it to a list:
struct DrawArgs { vertexCount: u32, instanceCount: atomic<u32>, firstVertex: u32,
firstInstance: u32 }
@group(0) @binding(0) var<storage, read_write> args: DrawArgs; // INDIRECT | STORAGE
@group(0) @binding(1) var<storage, read_write> visible: array<u32>;
fn keep(object: u32) {
visible[atomicAdd(&args.instanceCount, 1)] = object; // atomicAdd returns the old count
}The vertex shader then reads visible[instance_index], so instance k draws the k-th survivor. The struct's layout matches the four u32s of drawIndirect, since atomic<u32> has the size and alignment of a u32. Survivors arrive in whatever order the invocations ran, so sort (GPU Price Sort) if order matters, and reset instanceCount with writeBuffer() or a tiny compute pass before the culling pass, never after it.
<!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 COUNT = 96;
const code = /* wgsl */ `
struct DrawArgs { vertexCount: u32, instanceCount: atomic<u32>, firstVertex: u32, firstInstance: u32 }
@group(0) @binding(0) var<storage, read_write> args: DrawArgs; // INDIRECT | STORAGE
@group(0) @binding(1) var<storage, read_write> visible: array<u32>;
@group(0) @binding(2) var<storage> prices: array<f32>;
@group(0) @binding(3) var<uniform> limit: f32;
fn keep(object: u32) {
visible[atomicAdd(&args.instanceCount, 1)] = object; // atomicAdd returns the old count
}
@compute @workgroup_size(64) fn cull(@builtin(global_invocation_id) id: vec3u) {
if (id.x < arrayLength(&prices) && prices[id.x] <= limit) { keep(id.x); }
}
@group(0) @binding(1) var<storage> kept: array<u32>; // the same list, read-only
fn place(book: u32, q: vec2f) -> vec2f {
let col = f32(book % 24); let row = f32(book / 24);
return vec2f(-0.95 + col * 0.08 + q.x * 0.065, 0.36 - row * 0.4 + q.y * 0.3 * (0.5 + prices[book] / 60));
}
struct Out { @builtin(position) pos: vec4f, @location(0) color: vec4f }
@vertex fn allBooks(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
return Out(vec4f(place(i, vec2f(f32(v & 1), f32(v >> 1))), 0, 1), vec4f(0.85, 0.83, 0.80, 1)); // faint: everything
}
@vertex fn survivors(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
let book = kept[i]; // instance i draws survivor i
return Out(vec4f(place(book, vec2f(f32(v & 1), f32(v >> 1))), 0, 1), vec4f(0.08, 0.40 + 0.3 * f32(book % 3) / 2, 0.75, 1));
}
@fragment fn fs(in: Out) -> @location(0) vec4f { return in.color; }`;
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 priceList = new Float32Array(COUNT).map((_, i) => 8 + ((i * 7919) % 97) / 97 * 42);
const prices = device.createBuffer({ size: priceList.byteLength, usage: B.STORAGE | B.COPY_DST });
device.queue.writeBuffer(prices, 0, priceList);
const args = device.createBuffer({ size: 16, usage: B.INDIRECT | B.STORAGE | B.COPY_DST | B.COPY_SRC });
const visible = device.createBuffer({ size: COUNT * 4, usage: B.STORAGE });
const limit = device.createBuffer({ size: 4, usage: B.UNIFORM | B.COPY_DST });
const readback = [0, 1, 2].map(() => device.createBuffer({ size: 16, usage: B.COPY_DST | B.MAP_READ }));
const module = device.createShaderModule({ code });
const cull = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'cull' } });
const draw = (entryPoint) => device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
vertex: { module, entryPoint }, fragment: { module, entryPoint: 'fs', targets: [{ format }] } });
const allPipeline = draw('allBooks'), survivorPipeline = draw('survivors');
const cullGroup = device.createBindGroup({ layout: cull.getBindGroupLayout(0), entries: [
{ binding: 0, resource: { buffer: args } }, { binding: 1, resource: { buffer: visible } },
{ binding: 2, resource: { buffer: prices } }, { binding: 3, resource: { buffer: limit } }] });
const allGroup = device.createBindGroup({ layout: allPipeline.getBindGroupLayout(0), entries: [{ binding: 2, resource: { buffer: prices } }] });
const survivorGroup = device.createBindGroup({ layout: survivorPipeline.getBindGroupLayout(0), entries: [
{ binding: 1, resource: { buffer: visible } }, { binding: 2, resource: { buffer: prices } }] });
let shownCount = '?', shownLimit = 0;
function frame(now) {
const maxPrice = 29 + 21 * Math.sin(now / 1400);
device.queue.writeBuffer(limit, 0, new Float32Array([maxPrice]));
device.queue.writeBuffer(args, 0, new Uint32Array([4, 0, 0, 0])); // reset the count BEFORE culling
const encoder = device.createCommandEncoder();
const cp = encoder.beginComputePass();
cp.setPipeline(cull); cp.setBindGroup(0, cullGroup); cp.dispatchWorkgroups(Math.ceil(COUNT / 64)); 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(allPipeline); pass.setBindGroup(0, allGroup); pass.draw(4, COUNT);
pass.setPipeline(survivorPipeline); pass.setBindGroup(0, survivorGroup);
pass.drawIndirect(args, 0); // the GPU decides how many
pass.end();
const staging = readback.find((b) => b.mapState === 'unmapped'); // a count for the label, never awaited
if (staging) encoder.copyBufferToBuffer(args, 0, staging, 0, 16);
device.queue.submit([encoder.finish()]);
if (staging) {
const at = maxPrice;
staging.mapAsync(GPUMapMode.READ).then(() => { shownCount = new Uint32Array(staging.getMappedRange())[1]; shownLimit = at; staging.unmap(); });
}
ink.clearRect(0, 0, 600, 330);
ink.font = '12.5px ui-monospace, monospace'; ink.fillStyle = '#222';
ink.fillText(`keep books with price <= $${maxPrice.toFixed(2)} (atomicAdd appends each survivor)`, 12, 22);
ink.fillStyle = '#1f4f8a';
ink.fillText(`drawIndirect(args): instanceCount = ${shownCount} of ${COUNT} (read back a frame late, at $${shownLimit.toFixed(2)})`, 12, 320);
requestAnimationFrame(frame);
}
requestAnimationFrame(frame);
}
main();
</script>