Compute Filtering

Filtering the Catalog with a Compute Predicate and Scan

Stream compaction keeps the items that pass a test, packed in order: compute 0/1 flags, scan them, and write each kept item to the slot its scan gives. The test here is "in stock and under $25", on catalog.js's prices as WGSL constants, and out[0] receives the count a later drawIndirect() could use:

Filtering BookNest's six books with flags, a scan and a scatterJavaScript
// out: 7 x u32
@group(0) @binding(0) var<storage, read_write> out: array<u32>;   // count, then ids
const price = array(14.99, 39.50, 24.00, 18.75, 16.20, 21.30, 0, 0);
const inStock = array(1u, 1, 0, 1, 1, 1, 0, 0);    // Salt and Saffron (3) is sold out
var<workgroup> s: array<u32, 8>;
@compute @workgroup_size(8) fn main(@builtin(local_invocation_index) l: u32) {
  let keep = u32(inStock[l] == 1 && price[l] < 25.0);   // 1. the predicate
  s[l] = keep;
  workgroupBarrier();
  for (var d = 1u; d < 8; d <<= 1) {               // 2. inclusive scan of the flags
    let left = select(0u, s[l - d], l >= d);
    workgroupBarrier();
    s[l] += left;
    workgroupBarrier();
  }
  if (keep == 1) { out[s[l]] = l + 1; }            // 3. scatter: slot = inclusive sum
  if (l == 7) { out[0] = s[7]; }                   // the total is the last inclusive sum
}
Output
out: 4, 1, 4, 5, 6, 0, 0
gpu: 0.00 ms (median of 5)

Books 1, 4, 5 and 6 pass, in catalog order, exactly as BOOKS.filter((b) => b.inStock && b.price < 25) returns them in JavaScript; Salt and Saffron fails on stock alone. Unlike the atomicAdd() append of Culled Shelf, the order is stable and reproducible, which a sorted shelf or a paginated list needs.

Stream compaction of BookNest's six books: a predicate, a scan of the flags, and a scatter that packs the survivors in orderHTMLLive
<!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 titles = ['The Quiet Harbor', 'Small Steps', 'Salt and Saffron', "Clockmaker's", 'Night Trains', 'Paper Gardens'];
const prices = [14.99, 39.50, 24.00, 18.75, 16.20, 21.30];
const code = /* wgsl */ `
@group(0) @binding(0) var<storage, read_write> out: array<u32>;   // count, then ids
const price = array(14.99, 39.50, 24.00, 18.75, 16.20, 21.30, 0, 0);
const inStock = array(1u, 1, 0, 1, 1, 1, 0, 0);    // Salt and Saffron (3) is sold out
var<workgroup> s: array<u32, 8>;
@compute @workgroup_size(8) fn main(@builtin(local_invocation_index) l: u32) {
  let keep = u32(inStock[l] == 1 && price[l] < 25.0);   // 1. the predicate
  s[l] = keep;
  workgroupBarrier();
  for (var d = 1u; d < 8; d <<= 1) {               // 2. inclusive scan of the flags
    let left = select(0u, s[l - d], l >= d);
    workgroupBarrier();
    s[l] += left;
    workgroupBarrier();
  }
  if (keep == 1) { out[s[l]] = l + 1; }            // 3. scatter: slot = inclusive sum
  if (l == 7) { out[0] = s[7]; }                   // the total is the last inclusive sum
}`;
// Row 1: the catalog (kept books bright). Row 2: the packed output, placed by out[1..count].
const view = /* wgsl */ `
@group(0) @binding(0) var<storage> out: array<u32>;
struct Out { @builtin(position) pos: vec4f, @location(0) color: vec3f }
const tints = array(vec3f(.12, .37, .55), vec3f(.36, .25, .60), vec3f(.71, .27, .18), vec3f(.88, .60, .06), vec3f(.25, .49, .23), vec3f(.16, .62, .56));
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
  let q = vec2f(f32(v & 1), f32(v >> 1));
  if (i < 6) {                                              // the whole catalog
    var kept = false;
    for (var k = 1u; k <= out[0]; k++) { kept = kept || out[k] == i + 1; }
    let c = select(mix(tints[i], vec3f(0.95, 0.94, 0.91), 0.75), tints[i], kept);
    return Out(vec4f(-0.9 + f32(i) * 0.3 + q.x * 0.24, 0.2 + q.y * 0.5, 0, 1), c);
  }
  let slot = i - 6;                                         // the compacted list
  if (slot >= out[0]) { return Out(vec4f(0, 0, 0, 1), vec3f(0)); }   // degenerate: nothing drawn
  let book = out[slot + 1] - 1;
  return Out(vec4f(-0.9 + f32(slot) * 0.3 + q.x * 0.24, -0.72 + q.y * 0.5, 0, 1), tints[book]);
}
@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 out = device.createBuffer({ size: 28, usage: B.STORAGE | B.COPY_SRC });
  const read = device.createBuffer({ size: 28, usage: B.COPY_DST | B.MAP_READ });
  const compute = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code }) } });
  const module = device.createShaderModule({ code: view });
  const render = device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
    vertex: { module }, fragment: { module, targets: [{ format }] } });
  const encoder = device.createCommandEncoder();
  const cp = encoder.beginComputePass();
  cp.setPipeline(compute);
  cp.setBindGroup(0, device.createBindGroup({ layout: compute.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: out } }] }));
  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, device.createBindGroup({ layout: render.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: out } }] }));
  pass.draw(4, 12);
  pass.end();
  encoder.copyBufferToBuffer(out, 0, read, 0, 28);
  device.queue.submit([encoder.finish()]);
  await read.mapAsync(GPUMapMode.READ);
  const o = new Uint32Array(read.getMappedRange().slice(0));
  read.unmap();

  ink.font = 'bold 13px system-ui, sans-serif'; ink.fillStyle = '#222';
  ink.fillText('Catalog: keep "in stock and under $25"', 14, 24);
  ink.fillText(`Packed by the scan: out = [${[...o.slice(0, o[0] + 1)].join(', ')}]  (count first, then 1-based ids)`, 14, 186);
  ink.font = '10.5px system-ui, sans-serif'; ink.textAlign = 'center';
  titles.forEach((t, i) => { ink.fillStyle = '#333'; ink.fillText(t, 66 + i * 90, 148); ink.fillText(i === 2 ? 'sold out' : `$${prices[i].toFixed(2)}`, 66 + i * 90, 162); });
  for (let s = 0; s < o[0]; s++) ink.fillText(titles[o[s + 1] - 1], 66 + s * 90, 310);
  ink.textAlign = 'left'; ink.fillStyle = '#1e6b3a'; ink.font = '11.5px system-ui, sans-serif';
  ink.fillText('Same order as BOOKS.filter(...) in JavaScript: stable and reproducible, unlike an atomic append.', 14, 326);
}
main();
</script>