GPU Price Sort

To sort records, sort keys that carry them: each book's key is its price in cents times 8 plus its index, so the low three bits name the book, and 0xffffffff fills the two spare slots of the 8-key network. One workgroup runs Bitonic Sort on the GPU's network in workgroup memory; a render pass then draws a spine per slot, as tall as its price and in its book's color, from the buffer the sort left behind. The read-back only checks the result:

demos/ch04/sorted-shelf.html: a bitonic sort by price feeds the draw callHTMLLive

<canvas id="shelf" width="600" height="120"></canvas>
<script type="module">
const device = await (await navigator.gpu.requestAdapter()).requestDevice();
const { STORAGE, COPY_SRC, COPY_DST, MAP_READ } = GPUBufferUsage;
const prices = [14.99, 39.50, 24.00, 18.75, 16.20, 21.30];      // BookNest's six books
const keys = new Uint32Array(8).fill(0xffffffff);               // padding sorts last
prices.forEach((p, book) => { keys[book] = Math.round(p * 100) * 8 + book; });  // cents|id
const order = device.createBuffer({ size: 32, usage: STORAGE | COPY_SRC | COPY_DST });
device.queue.writeBuffer(order, 0, keys);
const module = device.createShaderModule({ code: /* wgsl */ `
  @group(0) @binding(0) var<storage, read_write> keys: array<u32, 8>;
  var<workgroup> k8: array<u32, 8>;
  @compute @workgroup_size(8) fn sort(@builtin(local_invocation_index) i: u32) {
    k8[i] = keys[i];
    workgroupBarrier();
    for (var k = 2u; k <= 8; k <<= 1) {                          // Section 4.17.7
      for (var j = k >> 1; j > 0; j >>= 1) {
        let p = i ^ j;
        if (p > i && (k8[i] > k8[p]) == ((i & k) == 0)) {
          let t = k8[i]; k8[i] = k8[p]; k8[p] = t;
        }
        workgroupBarrier();
      }
    }
    keys[i] = k8[i];
  }
  @group(0) @binding(0) var<storage> sorted: array<u32, 8>;
  const 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));
  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 q = array(vec2f(0, 0), vec2f(1, 0), vec2f(0, 1), vec2f(1, 1))[v];
    let cents = f32(sorted[i] >> 3);                             // slot i's price...
    let p = vec2f(-0.95 + f32(i) * 0.32 + q.x * 0.28, q.y * cents / 2300 - 0.9);
    return Out(vec4f(p, 0, 1), tint[sorted[i] & 7]);             // ...and its book
  }
  @fragment fn fs(in: Out) -> @location(0) vec4f { return vec4f(in.color, 1); }` });
const format = navigator.gpu.getPreferredCanvasFormat(), context = shelf.getContext('webgpu');
context.configure({ device, format });
const sort = device.createComputePipeline({ layout: 'auto', compute: { module } });
const draw = device.createRenderPipeline({ layout: 'auto', vertex: { module },
  fragment: { module, targets: [{ format }] }, primitive: { topology: 'triangle-strip' } });
const bind = (pipeline) => device.createBindGroup({ layout: pipeline.getBindGroupLayout(0),
  entries: [{ binding: 0, resource: order }] });
const read = device.createBuffer({ size: 32, usage: COPY_DST | MAP_READ });
const encoder = device.createCommandEncoder(), compute = encoder.beginComputePass();
compute.setPipeline(sort), compute.setBindGroup(0, bind(sort));
compute.dispatchWorkgroups(1), compute.end();
const pass = encoder.beginRenderPass({ colorAttachments: [{ loadOp: 'clear', storeOp: 'store',
  view: context.getCurrentTexture().createView(), clearValue: [0.96, 0.94, 0.90, 1] }] });
pass.setPipeline(draw), pass.setBindGroup(0, bind(draw));
pass.draw(4, 6), pass.end();                           // six spines, cheapest first
encoder.copyBufferToBuffer(order, read);
device.queue.submit([encoder.finish()]), await read.mapAsync(GPUMapMode.READ);
const gpu = [...new Uint32Array(read.getMappedRange()).slice(0, 6)].map((key) => key & 7);
const js = [0, 1, 2, 3, 4, 5].sort((a, b) => prices[a] - prices[b]);
console.log(`GPU order: ${gpu.map((b) => `${b + 1} $${prices[b].toFixed(2)}`).join(', ')}`);
console.log(`matches JavaScript's sort: ${gpu.join() === js.join()}`);
window.__done = true;
</script>

Browser output of 94

Output of 94

GPU order: 1 $14.99, 5 $16.20, 4 $18.75, 6 $21.30, 3 $24.00, 2 $39.50
matches JavaScript's sort: true

The spines rise from The Quiet Harbor (blue, $14.99) to Patterns of the Deep Web (purple, $39.50), in their cover colors. Packing works while cents times 8 fits in 32 bits; for wider keys, sort vec2u pairs.

MENU