Indirect Draw Lists

Building an Indirect Draw List from a Compute Pass

GPU-Built Arguments appended survivors with atomicAdd() on the instanceCount of a DrawArgs struct. BookNest's aisle does the same with the frustum test of GPU-Driven Frustum Culling, and the shader file also holds the drawing entry points, so the camera matrix and the procedural book layout (six shelves of 200, heights varied by slot) exist once:

demos/ch04/shelf-cull.wgsl: frustum culling into drawIndirect() argumentsWGSL
// shelf-cull.wgsl: frustum-cull 1,200 books into an indirect draw, then draw the survivors
struct Camera { eye: vec3f, yaw: f32, aspect: f32 }
struct DrawArgs { vertexCount: u32, instanceCount: atomic<u32>, firstVertex: u32,
                  firstInstance: u32 }
@group(0) @binding(0) var<uniform> cam: Camera;
@group(0) @binding(1) var<storage, read_write> args: DrawArgs;
@group(0) @binding(2) var<storage, read_write> visible: array<u32>;
const BOOKS = 1200u;                                  // 6 shelves of 200 books
fn viewProj() -> mat4x4f {                            // perspective (fov 1 rad) * view
  let f = 1 / tan(0.5);
  let s = sin(cam.yaw);
  let c = cos(cam.yaw);
  let fwd = vec4f(s, 0, -c, -dot(vec3f(s, 0, -c), cam.eye));    // distance along the view
  let right = vec4f(c, 0, s, -dot(vec3f(c, 0, s), cam.eye));
  let up = vec4f(0, 1, 0, -cam.eye.y);
  let z = fwd * 50 / 49.9 - vec4f(0, 0, 0, 5 / 49.9);   // near 0.1, far 50 -> depth 0..1
  return transpose(mat4x4f(right * f / cam.aspect, up * f, z, fwd));   // rows -> matrix
}
fn book(i: u32) -> vec4f {                            // left x, bottom y, height, width
  let slot = i % 200;
  return vec4f(f32(slot) * 0.12 - 12, f32(i / 200) * 0.5 - 1.3,
               0.3 + 0.025 * f32((slot * 7) % 5), 0.1);
}
@compute @workgroup_size(64) fn cull(@builtin(global_invocation_id) id: vec3u) {
  if (id.x >= BOOKS) { return; }
  let b = book(id.x);
  let center = vec4f(b.x + b.w / 2, b.y + b.z / 2, 0, 1);
  let radius = length(b.zw) / 2;                      // bounding sphere
  let m = transpose(viewProj());                      // m[k] = row k of the matrix
  let planes = array(m[3] + m[0], m[3] - m[0], m[3] + m[1], m[3] - m[1], m[2], m[3] - m[2]);
  for (var k = 0; k < 6; k++) {                       // left, right, bottom, top, near, far
    if (dot(planes[k], center) < -radius * length(planes[k].xyz)) { return; }   // outside
  }
  visible[atomicAdd(&args.instanceCount, 1)] = id.x;  // inside or crossing: keep
}
@group(0) @binding(2) var<storage> kept: array<u32>;  // the same list, read-only
struct Out { @builtin(position) pos: vec4f, @location(0) color: vec3f }
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
  let n = kept[i];                                    // instance i draws survivor n
  let b = book(n);
  let q = array(vec2f(0, 0), vec2f(1, 0), vec2f(0, 1), vec2f(1, 1))[v];
  let tint = array(vec3f(.12, .37, .55), vec3f(.36, .25, .60), vec3f(.88, .60, .06),
                   vec3f(.25, .49, .23), vec3f(.71, .27, .18), vec3f(.16, .62, .56))[n % 6];
  let world = vec4f(b.x + q.x * b.w, b.y + q.y * b.z, 0, 1);
  return Out(viewProj() * world, tint * (0.75 + 0.25 * q.x));   // a lit spine edge
}
@fragment fn fs(in: Out) -> @location(0) vec4f { return vec4f(in.color, 1); }

viewProj() builds the matrix from four rows (right, up, depth and forward, each with its translation in w), so the planes are sums and differences of m[0] to m[3] after the transpose. Appended ids arrive in scheduling order; use Compute Filtering's scan when order matters. Several meshes need one DrawArgs each at different offsets and one drawIndirect() per mesh: core WebGPU has no multi-draw indirect, and this adapter lists no such feature.

BookNest's 1,200-book aisle frustum-culled into drawIndirect() arguments every frame as the camera turnsHTMLLive
<!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="340"></canvas>
  <canvas id="labels" width="600" height="340"></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);
}

// shelf-cull.wgsl: frustum-cull 1,200 books into an indirect draw, then draw the survivors
const code = /* wgsl */ `
struct Camera { eye: vec3f, yaw: f32, aspect: f32 }
struct DrawArgs { vertexCount: u32, instanceCount: atomic<u32>, firstVertex: u32, firstInstance: u32 }
@group(0) @binding(0) var<uniform> cam: Camera;
@group(0) @binding(1) var<storage, read_write> args: DrawArgs;
@group(0) @binding(2) var<storage, read_write> visible: array<u32>;
const BOOKS = 1200u;                                  // 6 shelves of 200 books
fn viewProj() -> mat4x4f {                            // perspective (fov 1 rad) * view
  let f = 1 / tan(0.5);
  let s = sin(cam.yaw);
  let c = cos(cam.yaw);
  let fwd = vec4f(s, 0, -c, -dot(vec3f(s, 0, -c), cam.eye));    // distance along the view
  let right = vec4f(c, 0, s, -dot(vec3f(c, 0, s), cam.eye));
  let up = vec4f(0, 1, 0, -cam.eye.y);
  let z = fwd * 50 / 49.9 - vec4f(0, 0, 0, 5 / 49.9);   // near 0.1, far 50 -> depth 0..1
  return transpose(mat4x4f(right * f / cam.aspect, up * f, z, fwd));   // rows -> matrix
}
fn book(i: u32) -> vec4f {                            // left x, bottom y, height, width
  let slot = i % 200;
  return vec4f(f32(slot) * 0.12 - 12, f32(i / 200) * 0.5 - 1.3,
               0.3 + 0.025 * f32((slot * 7) % 5), 0.1);
}
@compute @workgroup_size(64) fn cull(@builtin(global_invocation_id) id: vec3u) {
  if (id.x >= BOOKS) { return; }
  let b = book(id.x);
  let center = vec4f(b.x + b.w / 2, b.y + b.z / 2, 0, 1);
  let radius = length(b.zw) / 2;                      // bounding sphere
  let m = transpose(viewProj());                      // m[k] = row k of the matrix
  let planes = array(m[3] + m[0], m[3] - m[0], m[3] + m[1], m[3] - m[1], m[2], m[3] - m[2]);
  for (var k = 0; k < 6; k++) {                       // left, right, bottom, top, near, far
    if (dot(planes[k], center) < -radius * length(planes[k].xyz)) { return; }   // outside
  }
  visible[atomicAdd(&args.instanceCount, 1)] = id.x;  // inside or crossing: keep
}
@group(0) @binding(2) var<storage> kept: array<u32>;  // the same list, read-only
struct Out { @builtin(position) pos: vec4f, @location(0) color: vec3f }
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
  let n = kept[i];                                    // instance i draws survivor n
  let b = book(n);
  let q = array(vec2f(0, 0), vec2f(1, 0), vec2f(0, 1), vec2f(1, 1))[v];
  let tint = array(vec3f(.12, .37, .55), vec3f(.36, .25, .60), vec3f(.88, .60, .06),
                   vec3f(.25, .49, .23), vec3f(.71, .27, .18), vec3f(.16, .62, .56))[n % 6];
  let world = vec4f(b.x + q.x * b.w, b.y + q.y * b.z, 0, 1);
  return Out(viewProj() * world, tint * (0.75 + 0.25 * q.x));   // a lit spine edge
}
@fragment fn fs(in: Out) -> @location(0) vec4f { return vec4f(in.color, 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 camera = device.createBuffer({ size: 32, usage: B.UNIFORM | B.COPY_DST });
  const args = device.createBuffer({ size: 16, usage: B.INDIRECT | B.STORAGE | B.COPY_DST | B.COPY_SRC });
  const visible = device.createBuffer({ size: 1200 * 4, usage: B.STORAGE });
  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 = device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
    vertex: { module, entryPoint: 'vs' }, fragment: { module, entryPoint: 'fs', targets: [{ format }] } });
  const cullGroup = device.createBindGroup({ layout: cull.getBindGroupLayout(0), entries: [
    { binding: 0, resource: { buffer: camera } }, { binding: 1, resource: { buffer: args } }, { binding: 2, resource: { buffer: visible } }] });
  const drawGroup = device.createBindGroup({ layout: draw.getBindGroupLayout(0), entries: [
    { binding: 0, resource: { buffer: camera } }, { binding: 2, resource: { buffer: visible } }] });

  let count = '?';
  function frame(now) {
    const yaw = Math.sin(now / 2600) * 0.9;
    device.queue.writeBuffer(camera, 0, new Float32Array([Math.sin(now / 4000) * 4, 0, 4, yaw, 600 / 340]));
    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(1200 / 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(draw); pass.setBindGroup(0, drawGroup);
    pass.drawIndirect(args, 0);                         // only the survivors
    pass.end();
    const staging = readback.find((b) => b.mapState === 'unmapped');
    if (staging) encoder.copyBufferToBuffer(args, 0, staging, 0, 16);
    device.queue.submit([encoder.finish()]);
    if (staging) staging.mapAsync(GPUMapMode.READ).then(() => { count = new Uint32Array(staging.getMappedRange())[1]; staging.unmap(); });
    ink.clearRect(0, 0, 600, 340);
    ink.fillStyle = 'rgba(247,244,238,0.9)'; ink.fillRect(0, 0, 600, 26);
    ink.font = '12.5px ui-monospace, monospace'; ink.fillStyle = '#222';
    ink.fillText(`drawIndirect(args): instanceCount = ${count} of 1,200 books   yaw ${yaw.toFixed(2)} rad`, 10, 18);
    requestAnimationFrame(frame);
  }
  requestAnimationFrame(frame);
}
main();
</script>