Culled Shelf

Indirect-Drawing BookNest's Culled Shelf

The shelf now filters on the GPU. A compute shader tests each book against a uniform rule (in stock, at most a price, at least a rating), appends survivors to visible and counts them in the indirect arguments; a bundle recorded once holds a single drawIndirect(). Two frames run with different rules, "under $20" (left canvas) and "rated 4.3 or better, under $40" (right), and the arguments are copied back after each frame:

demos/ch04/culled-shelf.html: compute culling, drawIndirect and one bundle replayedHTML
<canvas id="cheap" width="300" height="80"></canvas>
<canvas id="rated" width="300" height="80"></canvas>
<script src="mipmaps.js"></script>
<script type="module">
const device = await (await navigator.gpu.requestAdapter()).requestDevice();
const { STORAGE, UNIFORM, INDIRECT, COPY_DST, COPY_SRC, MAP_READ } = GPUBufferUsage;
const bitmaps = await Promise.all([1, 2, 3, 4, 5, 6].map(async (n) =>
  createImageBitmap(await (await fetch(`booknest-cover-${n}.png`)).blob())));
const covers = device.createTexture({ size: [256, 384, 6], format: 'rgba8unorm',
  mipLevelCount: 9, usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.COPY_DST |
  GPUTextureUsage.RENDER_ATTACHMENT });
bitmaps.forEach((source, layer) => device.queue.copyExternalImageToTexture(
  { source }, { texture: covers, origin: [0, 0, layer] }, [256, 384]));
generateMipmaps(device, covers);                       // Section 4.13.2
const buffer = (size, usage) => device.createBuffer({ size, usage });
const books = buffer(72, STORAGE | COPY_DST), filter = buffer(8, UNIFORM | COPY_DST);
const args = buffer(16, INDIRECT | STORAGE | COPY_DST | COPY_SRC);   // drawIndirect's 4 u32
const visible = buffer(24, STORAGE), readback = buffer(32, COPY_DST | MAP_READ);
device.queue.writeBuffer(books, 0, new Float32Array([14.99, 4.6, 1, 39.50, 4.3, 1,
  24.00, 4.8, 0, 18.75, 4.1, 1, 16.20, 4.5, 1, 21.30, 4.4, 1]));   // price, rating, stock
const cullCode = /* wgsl */ `
struct Book { price: f32, rating: f32, inStock: f32 }
struct Filter { maxPrice: f32, minRating: f32 }
struct DrawArgs { vertexCount: u32, instanceCount: atomic<u32>, firstVertex: u32,
                  firstInstance: u32 }
@group(0) @binding(0) var<storage> books: array<Book>;
@group(0) @binding(1) var<uniform> rule: Filter;       // "filter" is reserved
@group(0) @binding(2) var<storage, read_write> args: DrawArgs;
@group(0) @binding(3) var<storage, read_write> visible: array<u32>;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) id: vec3u) {
  if (id.x >= arrayLength(&books)) { return; }
  let b = books[id.x];
  if (b.inStock > 0 && b.price <= rule.maxPrice && b.rating >= rule.minRating) {
    visible[atomicAdd(&args.instanceCount, 1)] = id.x;   // append a survivor
  }
}`;
const drawCode = /* wgsl */ `
@group(0) @binding(0) var<storage> visible: array<u32>;
@group(0) @binding(1) var covers: texture_2d_array<f32>;
@group(0) @binding(2) var smp: sampler;
struct Out { @builtin(position) pos: vec4f, @location(0) uv: vec2f,
             @location(1) @interpolate(flat) book: u32 }
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
  let q = array(vec2f(0, 0), vec2f(1, 0), vec2f(0, 1), vec2f(1, 0), vec2f(1, 1),
                vec2f(0, 1))[v];
  let book = visible[i];                               // instance i draws a survivor
  let p = vec2f(-0.95 + f32(book) * 0.32 + q.x * 0.30, q.y * 1.7 - 0.85);
  return Out(vec4f(p, 0, 1), vec2f(q.x, 1 - q.y), book);
}
@fragment fn fs(in: Out) -> @location(0) vec4f {
  return textureSample(covers, smp, in.uv, in.book);
}`;
const cull = device.createComputePipeline({ layout: 'auto',
  compute: { module: device.createShaderModule({ code: cullCode }) } });
const cullGroup = device.createBindGroup({ layout: cull.getBindGroupLayout(0), entries:
  [books, filter, args, visible].map((resource, binding) => ({ binding, resource })) });
const format = navigator.gpu.getPreferredCanvasFormat();
const module = device.createShaderModule({ code: drawCode });
const draw = device.createRenderPipeline({ layout: 'auto', vertex: { module },
  fragment: { module, targets: [{ format }] } });
const sampler = device.createSampler({ magFilter: 'linear', minFilter: 'linear',
  mipmapFilter: 'linear' });
const bundleEncoder = device.createRenderBundleEncoder({ colorFormats: [format] });
bundleEncoder.setPipeline(draw);                       // recorded once...
bundleEncoder.setBindGroup(0, device.createBindGroup({ layout: draw.getBindGroupLayout(0),
  entries: [visible, covers.createView(), sampler].map((resource, binding) =>
    ({ binding, resource })) }));
bundleEncoder.drawIndirect(args, 0);                   // ...with counts read at run time
const bundle = bundleEncoder.finish();
[[cheap, 20, 0], [rated, 40, 4.3]].forEach(([canvas, maxPrice, minRating], frame) => {
  device.queue.writeBuffer(filter, 0, new Float32Array([maxPrice, minRating]));
  device.queue.writeBuffer(args, 0, new Uint32Array([6, 0, 0, 0]));   // reset the count
  const context = canvas.getContext('webgpu');
  context.configure({ device, format });
  const encoder = device.createCommandEncoder();
  const compute = encoder.beginComputePass();
  compute.setPipeline(cull), compute.setBindGroup(0, cullGroup);
  compute.dispatchWorkgroups(1), compute.end();        // 64 invocations, 6 books
  const pass = encoder.beginRenderPass({ colorAttachments: [{ loadOp: 'clear',
    storeOp: 'store', view: context.getCurrentTexture().createView(),
    clearValue: [0.96, 0.94, 0.90, 1] }] });
  pass.executeBundles([bundle]), pass.end();           // the same bundle every frame
  encoder.copyBufferToBuffer(args, 0, readback, frame * 16, 16);
  device.queue.submit([encoder.finish()]);
});
await readback.mapAsync(GPUMapMode.READ);
const counts = new Uint32Array(readback.getMappedRange());
console.log(`under $20: drawIndirect(${counts.slice(0, 4).join(', ')})`);
console.log(`4.3 stars and up: drawIndirect(${counts.slice(4, 8).join(', ')})`);
window.__done = true;
</script>
Browser output of Listing 4.82
Browser output of 82
Output of 82
under $20: drawIndirect(6, 3, 0, 0)
4.3 stars and up: drawIndirect(6, 4, 0, 0)

Each frame's JavaScript is the same: two writeBuffer() calls, one dispatch, one executeBundles(). The GPU decided to draw three covers, then four; out-of-stock Salt and Saffron never appears, and the survivors keep their columns because the vertex shader places each by its book index. One trap met on the way: filter is a reserved word in WGSL, so the uniform is called rule.

BookNest's shelf filtered on the GPU: compute culling fills drawIndirect() arguments for one bundle, replayed under two rulesHTMLLive
<!doctype html>
<style>
  body { margin: 0; background: #f7f4ee; font: 13px system-ui, sans-serif; color: #222; }
  .row { display: flex; gap: 8px; max-width: 600px; }
  .row div { flex: 1; min-width: 0; }
  canvas { display: block; width: 100%; }
  pre { margin: 6px 0; font: 12px ui-monospace, monospace; max-width: 600px; white-space: pre-wrap; }
</style>
<div class="row">
  <div><b>Under $20</b><canvas id="cheap" width="296" height="90"></canvas></div>
  <div><b>Rated 4.3+, under $40</b><canvas id="rated" width="296" height="90"></canvas></div>
</div>
<pre id="out"></pre>
<script>
const cheap = document.getElementById('cheap'), rated = document.getElementById('rated');
function showMessage(text) {                     // 2D fallback when WebGPU is missing
  const ctx = cheap.getContext('2d');
  ctx.fillStyle = '#fbeaea'; ctx.fillRect(0, 0, cheap.width, cheap.height);
  ctx.fillStyle = '#8a2b2b'; ctx.font = '13px system-ui, sans-serif'; ctx.textAlign = 'center';
  ctx.fillText(text, cheap.width / 2, cheap.height / 2);
}

const covers = [['Quiet', 'Harbor', '#1f5f8b'], ['Small', 'Steps', '#5c4099'], ['Salt &', 'Saffron', '#b5452e'],
                ['Clock-', 'maker', '#e0990f'], ['Night', 'Trains', '#2d5d2a'], ['Paper', 'Gardens', '#299e8f']];
const cullCode = /* wgsl */ `
struct Book { price: f32, rating: f32, inStock: f32 }
struct Filter { maxPrice: f32, minRating: f32 }
struct DrawArgs { vertexCount: u32, instanceCount: atomic<u32>, firstVertex: u32, firstInstance: u32 }
@group(0) @binding(0) var<storage> books: array<Book>;
@group(0) @binding(1) var<uniform> rule: Filter;       // "filter" is reserved in WGSL
@group(0) @binding(2) var<storage, read_write> args: DrawArgs;
@group(0) @binding(3) var<storage, read_write> visible: array<u32>;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) id: vec3u) {
  if (id.x >= arrayLength(&books)) { return; }
  let b = books[id.x];
  if (b.inStock > 0 && b.price <= rule.maxPrice && b.rating >= rule.minRating) {
    visible[atomicAdd(&args.instanceCount, 1)] = id.x;   // append a survivor
  }
}`;
const drawCode = /* wgsl */ `
@group(0) @binding(0) var<storage> visible: array<u32>;
@group(0) @binding(1) var covers: texture_2d_array<f32>;
@group(0) @binding(2) var smp: sampler;
struct Out { @builtin(position) pos: vec4f, @location(0) uv: vec2f, @location(1) @interpolate(flat) book: u32 }
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
  let q = array(vec2f(0, 0), vec2f(1, 0), vec2f(0, 1), vec2f(1, 0), vec2f(1, 1), vec2f(0, 1))[v];
  let book = visible[i];                               // instance i draws a survivor
  let p = vec2f(-0.95 + f32(book) * 0.32 + q.x * 0.30, q.y * 1.7 - 0.85);
  return Out(vec4f(p, 0, 1), vec2f(q.x, 1 - q.y), book);
}
@fragment fn fs(in: Out) -> @location(0) vec4f { return textureSample(covers, smp, in.uv, in.book); }`;

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 { STORAGE, UNIFORM, INDIRECT, COPY_DST, COPY_SRC, MAP_READ } = GPUBufferUsage;

  const texture = device.createTexture({ size: [64, 96, 6], format: 'rgba8unorm',
    usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.COPY_DST | GPUTextureUsage.RENDER_ATTACHMENT });
  covers.forEach(([a, b, color], layer) => {            // small covers painted in code
    const art = new OffscreenCanvas(64, 96), g = art.getContext('2d');
    g.fillStyle = color; g.fillRect(0, 0, 64, 96); g.fillStyle = '#f5ecd7'; g.fillRect(0, 28, 64, 34);
    g.fillStyle = '#222'; g.font = 'bold 11px Georgia, serif'; g.textAlign = 'center'; g.fillText(a, 32, 42); g.fillText(b, 32, 56);
    device.queue.copyExternalImageToTexture({ source: art }, { texture, origin: [0, 0, layer] }, [64, 96]);
  });
  const buffer = (size, usage) => device.createBuffer({ size, usage });
  const books = buffer(72, STORAGE | COPY_DST), filter = buffer(8, UNIFORM | COPY_DST);
  const args = buffer(16, INDIRECT | STORAGE | COPY_DST | COPY_SRC);   // drawIndirect's 4 u32
  const visible = buffer(24, STORAGE), readback = buffer(32, COPY_DST | MAP_READ);
  device.queue.writeBuffer(books, 0, new Float32Array([14.99, 4.6, 1, 39.50, 4.3, 1,
    24.00, 4.8, 0, 18.75, 4.1, 1, 16.20, 4.5, 1, 21.30, 4.4, 1]));   // price, rating, stock

  const cull = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code: cullCode }) } });
  const cullGroup = device.createBindGroup({ layout: cull.getBindGroupLayout(0), entries:
    [books, filter, args, visible].map((b, binding) => ({ binding, resource: { buffer: b } })) });
  const format = navigator.gpu.getPreferredCanvasFormat();
  const module = device.createShaderModule({ code: drawCode });
  const draw = device.createRenderPipeline({ layout: 'auto', vertex: { module }, fragment: { module, targets: [{ format }] } });
  const bundleEncoder = device.createRenderBundleEncoder({ colorFormats: [format] });
  bundleEncoder.setPipeline(draw);                     // recorded once...
  bundleEncoder.setBindGroup(0, device.createBindGroup({ layout: draw.getBindGroupLayout(0), entries: [
    { binding: 0, resource: { buffer: visible } }, { binding: 1, resource: texture.createView() },
    { binding: 2, resource: device.createSampler({ magFilter: 'linear', minFilter: 'linear' }) }] }));
  bundleEncoder.drawIndirect(args, 0);                 // ...with counts read at run time
  const bundle = bundleEncoder.finish();

  [[cheap, 20, 0], [rated, 40, 4.3]].forEach(([canvas, maxPrice, minRating], frame) => {
    device.queue.writeBuffer(filter, 0, new Float32Array([maxPrice, minRating]));
    device.queue.writeBuffer(args, 0, new Uint32Array([6, 0, 0, 0]));   // reset the count
    const context = canvas.getContext('webgpu');
    context.configure({ device, format });
    const encoder = device.createCommandEncoder();
    const compute = encoder.beginComputePass();
    compute.setPipeline(cull); compute.setBindGroup(0, cullGroup);
    compute.dispatchWorkgroups(1); compute.end();       // 64 invocations, 6 books
    const pass = encoder.beginRenderPass({ colorAttachments: [{ loadOp: 'clear', storeOp: 'store',
      view: context.getCurrentTexture().createView(), clearValue: [0.96, 0.94, 0.90, 1] }] });
    pass.executeBundles([bundle]); pass.end();          // the same bundle every frame
    encoder.copyBufferToBuffer(args, 0, readback, frame * 16, 16);
    device.queue.submit([encoder.finish()]);
  });
  await readback.mapAsync(GPUMapMode.READ);
  const counts = new Uint32Array(readback.getMappedRange());
  document.getElementById('out').textContent =
    `under $20: drawIndirect(${counts.slice(0, 4).join(', ')})\n4.3 stars and up: drawIndirect(${counts.slice(4, 8).join(', ')})\n` +
    'Salt and Saffron is out of stock and never appears; survivors keep their columns.';
}
main();
</script>