GPU-Driven Rendering

Chapter 21 showed the CPU cost of draw calls, and instancing as the first cure. GPU-driven rendering takes the idea to its conclusion. The scene's objects, transforms and materials live in GPU buffers; compute shaders cull them and build the draw commands; and the CPU's job shrinks to a handful of dispatches per frame, whether the scene holds a thousand objects or a million. This is how today's most detailed real-time scenes are drawn.

An ordinary pass.draw(vertexCount, instanceCount) takes its numbers from JavaScript. Its twin, pass.drawIndirect(buffer, offset), reads the same four numbers (vertex count, instance count, first vertex, first instance) from a GPU buffer, when the GPU gets to that command. So a compute shader that runs earlier in the frame can decide them. The pattern, used in the hero:

1. Every object's data (position, size, material) sits in a storage buffer.
2. A culling compute shader runs one thread per object. Each visible object claims a slot in a "visible list" with an atomic add on the draw command's instance count (Chapter 18), and writes its index there.
3. One drawIndirect draws "instance count" instances; each instance's vertex shader looks up its object through the visible list.

The CPU never learns how many objects were visible. It doesn't need to. Larger renderers go further with multi-draw indirect (many draw commands in one buffer, one per mesh or material), which some graphics APIs offer and WebGPU is adding.

Frustum culling (Chapter 9) is the easy half: is the object's bounding sphere entirely outside any of the view's six planes? One thread per object, six dot products. Occlusion culling is the valuable half: in a city street, most of the city is hidden behind the nearest buildings. The GPU test uses the same Hi-Z pyramid as Chapter 33's reflections, but keeping the furthest depth in each block: a coarse texel says "everything in this block is at least this close". To test an object:

1. Project its bounding box onto the screen: a rectangle, and the nearest depth of any of its corners.
2. Pick the pyramid level where that rectangle covers about two texels across, and read those few texels.
3. If the object's nearest point is further away than the furthest depth in all of them, something closer covers every pixel it could touch. It's hidden: don't draw it.

Where does the depth pyramid come from, before the frame is drawn? Three common answers: draw a few big, simple occluders first (the hero does this with the buildings); reuse last frame's depth, reprojected; or draw in two phases, first the objects that were visible last frame, then build the pyramid from their depth and test everything else against it, drawing what turns out to be newly visible. The two-phase method is exact and needs no hand-picked occluders, which is why it's become the standard.

Culling whole objects is coarse: a huge rock that's half hidden is drawn in full. Modern pipelines split every mesh into meshlets (or clusters): small groups of about 64 to 128 neighboring triangles. Each stores a bounding sphere and a normal cone, the range of directions its triangles face. Then every meshlet can be culled on its own: outside the frustum, hidden behind something, or facing entirely away from the camera (if every triangle's normal points away, the whole meshlet is back-facing). Graphics APIs with mesh shaders process meshlets directly; elsewhere, a compute shader culls meshlets and writes their triangles into an index buffer for an indirect draw.

Level of detail by screen error

Meshlets also enable fine-grained level of detail (LOD). A distant rock doesn't need 100,000 triangles to cover 40 pixels. The test for "detailed enough" is simple: how far, in pixels, would the simplified version's surface deviate from the full one? If under a pixel, nobody can tell. Classic LOD swaps whole objects between a few precomputed versions, with visible pops. Part 4 does it per cluster, so every region of a surface picks its own level.

Put the pieces together and you can stop worrying about triangle counts. Artists bring in film-quality meshes, millions of triangles each, with no hand-made levels of detail; the renderer draws, every frame, about as many triangles as the screen has pixels, wherever they're needed, whatever the scene holds. It's called virtualized geometry, by analogy with virtual memory and virtual textures: the full data exists, but only the parts at the detail the current view needs are loaded and drawn. The idea of a hierarchy of triangle patches goes back to batched multi-triangulation (Paolo Cignoni and colleagues, 2005); Brian Karis and colleagues made it practical at film detail in 2021, and this part follows their design.

Building the hierarchy

The hierarchy is built once, offline, from the full mesh:

1. Split the mesh into clusters of up to 128 neighboring triangles.
2. Gather neighboring clusters into groups of 8 to 32.
3. Simplify each group, as one piece, to half its triangles, but with every vertex on the group's border held fixed. The simplified group meets its neighbors along exactly the same edges as before, whatever detail they're drawn at.
4. Split the simplified group back into clusters, half as many as went in. These are the next level up.
5. Repeat from step 2 with the new clusters, until one cluster is left. The groups come out different each time, so the borders that had to stay fixed at one level are inside a group at the next, and get simplified then.

Because a group goes in as several clusters and comes out as several, the result isn't a tree: each cluster has several parents, and each parent several children. It's a directed acyclic graph. Each cluster stores two numbers with their bounding spheres: the error of the simplification that made it (zero for the full detail), and the error of the group it was simplified in, which its parents have as their own. The builder makes sure errors only grow going up (a group's error is at least its children's), and so do the spheres.

Choosing a cut

Each frame, the renderer needs a cut through the graph: a set of clusters that covers the mesh exactly once, fine where the camera is close, coarse where it isn't. The rule for one cluster is short. Project its own error onto the screen (the error divided by the distance to its sphere, times the focal length in pixels), and its parents' error the same way. Draw the cluster if its own error is under the allowed amount, a pixel say, and its parents' isn't:

let selfOk = pixels(selfError, selfSphere) <= allowed;
let parentOk = pixels(parentError, parentSphere) <= allowed;
if (selfOk && !parentOk) {
  append(cluster);   // this is the right level here
}

No cluster needs to know what any other decided, so every cluster of every object is tested in parallel, with no tree to walk. It works because errors only grow going up: along any path from the full detail to the top, the test flips from "fine enough" to "not fine enough" exactly once, so exactly one cluster on each path is drawn. And there are no cracks: all the clusters of a group share the same parents, so they all make the same decision, and a group is either drawn whole or replaced whole by its parents, which meet the neighbors along the same locked border.

The figure's hierarchy is built in JavaScript when the page loads, in about a second: the 81,920 triangles become about 3,000 clusters in five or six levels, grouped 8 at a time and simplified by collapsing edges, cheapest first, by quadric error (Michael Garland and Paul Heckbert, 1997). The compute shader runs one thread for each cluster of each asteroid: the level test above, then the frustum and the normal cone, then an atomic append. The draw is one indirect call with an instance per visible cluster, each instance 128 triangles long; a cluster with fewer triangles collapses its extra ones to nothing. A production renderer adds the occlusion culling of Part 2 (drawing last frame's visible clusters first, building the depth pyramid from them, and testing the rest against it), and culls whole groups at a time through a bounding volume hierarchy, so it doesn't have to test every cluster of every object.

Triangles the size of pixels

With the error allowed at a pixel, most triangles end up a pixel or two across, and the hardware rasterizer is bad at those. It sets up triangles at a fixed rate, and shades pixels in 2 × 2 blocks, so a one-pixel triangle costs as much as four. Renderers built this way rasterize the small triangles themselves, in a compute shader: one thread per triangle loops over the few pixels it covers and, for each, keeps the nearest with a 64-bit atomic maximum on a value packing depth (in the high bits) with the cluster and triangle number, which is exactly a visibility buffer (Part 5). Big triangles still go to the hardware. Core WebGPU has only 32-bit atomics, so in a browser the depth and the triangle number have to share 32 bits, or be resolved in two passes.

Streaming

The last piece is memory. The hierarchy is stored on disk in pages of a few groups each. Each frame, the cut says which clusters are wanted; pages that aren't loaded yet are requested, and until they arrive, their coarser parents are drawn instead, so detail sharpens a moment later instead of leaving holes. The top levels are always kept loaded, so there's always something to fall back on. Memory use then follows the screen's resolution and what's in view, not the size of the scene: a single mesh can be bigger than the GPU's memory.

Deferred shading (Chapter 21) writes a fat G-buffer: color, normal, roughness and more, for every pixel. With tiny triangles, much of that writing is wasted: triangles a few pixels big are shaded in 2 × 2 pixel blocks (so the GPU can compute texture mip levels), and at their edges up to three of every four shaded pixels are discarded. The visibility buffer (Christopher Burns and Warren Hunt, 2013) stores almost nothing during rasterization: per pixel, just the instance and triangle that's visible, packed into one 32-bit or 64-bit number. A later compute pass reads that number, fetches the triangle's three vertices, recomputes the exact position and texture coordinates for the pixel, and evaluates the material once per pixel, never for pixels that end up hidden.

Two numbers name a visible triangle: which instance it belongs to, and its number within that instance's mesh. Here they share one 32-bit integer: the instance in the top 12 bits (up to 4,096 instances) and the triangle in the bottom 20 (up to about a million triangles per mesh). The fragment shader writes that single integer into an r32uint target, cleared beforehand to all ones to mean "nothing here". The vertex shader passes both numbers along unchanged; drawing without an index buffer, so that every three vertices make one triangle, the triangle's number is simply vertex_index / 3:

const TRI_BITS = 20u;                  // the triangle's share of the 32 bits

// The whole fragment shader of the visibility pass. "flat": the same
// value for the whole triangle, not blended between its corners.
@fragment
fn visFS(@location(0) @interpolate(flat) instance: u32,
         @location(1) @interpolate(flat) tri: u32) -> @location(0) u32 {
  return (instance << TRI_BITS) | tri;
}

How the bits are split is a choice. A renderer with more instances uses a 64-bit target, or, with Part 4's clusters, stores the slot in the visible list (which knows the instance and the cluster) plus the triangle's number within its cluster, which never needs more than 7 bits.

The material pass runs one compute thread per pixel. It takes the number apart, fetches the triangle's three corners and projects them to the screen, just as the vertex shader did. Then it works out the pixel's barycentric coordinates: three weights saying how close the pixel is to each corner. Each one is the area of the small triangle formed by the pixel and the other two corners, as a share of the whole triangle's area. Blending the corners' texture coordinates (or normals, or anything else) with those weights gives the values for this pixel.

@compute @workgroup_size(8, 8)
fn resolve(@builtin(global_invocation_id) id: vec3u) {
  let packed = textureLoad(visTex, id.xy, 0).r;
  if (packed == 0xffffffffu) {
    return;                             // no triangle covers this pixel
  }
  let instance = packed >> TRI_BITS;
  let tri = packed & ((1u << TRI_BITS) - 1u);
  var clip: array<vec4f, 3>;
  var px: array<vec2f, 3>;
  var uv: array<vec2f, 3>;
  for (var k = 0u; k < 3u; k++) {
    let v = tri * 3u + k;               // every 3 vertices are one triangle
    clip[k] = u.viewProj * models[instance] * vec4f(positions[v], 1.0);
    px[k] = (clip[k].xy / clip[k].w * vec2f(0.5, -0.5) + 0.5) * u.screen;
    uv[k] = uvs[v];
  }
  // each corner's weight: the area across from it, over the whole area
  let p = vec2f(id.xy) + 0.5;           // this pixel's center
  let area = cross2(px[1] - px[0], px[2] - px[0]);
  var b = vec3f(cross2(px[1] - p, px[2] - p),
                cross2(px[2] - p, px[0] - p),
                cross2(px[0] - p, px[1] - p)) / area;
  // undo perspective: weight each corner by 1 / w, then rescale to sum 1
  b /= vec3f(clip[0].w, clip[1].w, clip[2].w);
  b /= b.x + b.y + b.z;
  let texCoord = b.x * uv[0] + b.y * uv[1] + b.z * uv[2];
  textureStore(outTex, id.xy, vec4f(shade(instance, texCoord), 1.0));
}

And what it reads and the two small functions it calls. shade stands in for a real material: it paints each instance's color as a checkerboard, enough to show that the texture coordinates come out right.

struct Uniforms { viewProj: mat4x4f, screen: vec2f };
@group(0) @binding(0) var<uniform> u: Uniforms;
@group(0) @binding(1) var visTex: texture_2d<u32>;
@group(0) @binding(2) var outTex: texture_storage_2d<rgba8unorm, write>;
// the mesh's vertices, three per triangle (one mesh, shared by every
// instance here), and each instance's placement and color
@group(0) @binding(3) var<storage, read> positions: array<vec3f>;
@group(0) @binding(4) var<storage, read> uvs: array<vec2f>;
@group(0) @binding(5) var<storage, read> models: array<mat4x4f>;
@group(0) @binding(6) var<storage, read> colors: array<vec4f>;

// twice the signed area of the triangle with sides a and b
fn cross2(a: vec2f, b: vec2f) -> f32 {
  return a.x * b.y - a.y * b.x;
}
// a stand-in material: the instance's color, in an 8 x 8 checkerboard
fn shade(instance: u32, uv: vec2f) -> vec3f {
  let cell = vec2u(floor(uv * 8.0));
  let dark = ((cell.x + cell.y) & 1u) == 0u;
  return colors[instance].rgb * select(1.0, 0.6, dark);
}

The two lines after the areas matter. On screen, the far half of a tilted triangle is squeezed smaller than the near half, so weights measured in pixels would slide the texture toward the far end. Dividing by each corner's w, its depth in front of the camera, and then making the weights add up to 1 again, is the same correction the rasterizer applies to every interpolated value. Real renderers also compute how fast the texture coordinates change from one pixel to the next, which the 2 × 2 blocks used to give for free, so texture sampling can still pick a mip level.

There is a catch in "evaluates the material". One compute pass now shades the pixels of every material in the frame, so one shader has to be able to read every material's textures. Part 6 is about how.

The hero draws 65,536 objects with one call, and it gets away with that because they all share one material. A real scene has hundreds of materials, each with its own handful of textures, and here the binding model of Chapter 17 gets in the way. A shader can only reach the textures in the bind groups that were set when the draw was recorded, and by default no more than 16 sampled textures per shader stage. So a GPU-built draw list has to be split by material after all: one indirect draw per material, with a setBindGroup in front of each, recorded by the CPU, which has to know in advance which materials exist and where each one's range of the list starts. The visibility buffer's material pass is worse off: it shades the pixels of every material in a single dispatch, so there is no point at which to switch bind groups at all.

Renderers have long worked around this. A texture array packs many textures into one binding, but every layer must have the same size and format. An atlas packs them into one big texture, but filtering bleeds across neighbors, mip levels blur them together, and space is wasted. Both turn adding a texture into a repacking job.

Indices instead of bindings

Early GPUs had a fixed set of hardware slots for textures, and setting a bind group, or its equivalent, filled those slots. Hardware moved on long ago. A modern GPU names a texture by a pointer or an index into a table in memory, so nothing stops a shader from choosing which texture to read while it runs. The native graphics APIs expose this as bindless resources: descriptor heaps in Direct3D 12 (2015), argument buffers in Metal 2 (2017), and descriptor indexing in Vulkan, core since version 1.2 (2020). Every texture the scene uses goes into one large table, each material stores the index numbers of its textures, and the shader looks them up:

1. The table is filled once, and changed only when textures are loaded or unloaded.
2. The per-object data the culling shader already reads (Part 1) carries a material number; the material buffer holds that material's table indices.
3. The fragment shader, or the visibility buffer's material pass, reads the indices and samples the textures from the table. One draw, or one dispatch, covers every material.

How WebGPU plans to do it

The proposal adds one new object, the resource table, and two optional device features. With "sampling-resource-table", a table can hold sampled textures and samplers, which covers materials. "heterogeneous-resource-table" builds on it and allows almost any kind of resource, including storage textures and, later, storage buffers. They are separate because many devices that can index textures freely can't do the same for everything else.

// Proposed API: a draft, which may change before it ships.
const device = await adapter.requestDevice({ requiredFeatures: ["sampling-resource-table"] });
const table = device.createResourceTable({ size: 4096 });    // up to 65,536 slots

// insert() puts a resource in a free slot and returns the slot's number.
const linear = table.insert(device.createSampler({ magFilter: "linear", minFilter: "linear", mipmapFilter: "linear" }));
for (const m of materials) {
  m.albedoSlot = table.insert(m.albedo.createView());
  m.normalSlot = table.insert(m.normal.createView());
}
writeMaterialBuffer(materials, linear);   // each material's slot numbers, for the shader

// Pipelines that read the table say so in their layout ("auto" layouts work it out).
const layout = device.createPipelineLayout({ bindGroupLayouts: [sceneLayout], usesResourceTable: true });

// A render pass gets its table when it begins; a compute pass with setResourceTable(table).
const pass = encoder.beginRenderPass({ colorAttachments, depthStencilAttachment, resourceTable: table });
pass.setPipeline(pipeline);
pass.setBindGroup(0, sceneGroup);   // objects, materials, camera: the same for every draw
pass.drawIndirect(args, 0);          // every visible object, every material, one draw
pass.end();

In WGSL, an enable resource_table; directive turns on two built-in functions. The table has no name and no @binding: it's simply whichever table the pass was given. getResource<T>(i) returns the resource in slot i as type T, and hasResource<T>(i) reports whether that slot really holds one.

enable resource_table;

struct Material { albedo: u32, normal: u32, sampler: u32, baseColor: u32 };
@group(0) @binding(1) var<storage, read> materials: array<Material>;

@fragment
fn fs(in: VOut) -> @location(0) vec4f {
  let m = materials[in.material];
  let s = getResource<sampler>(m.sampler);
  let texel = textureSample(getResource<texture_2d<f32>>(m.albedo), s, in.uv);
  // A texture that's still streaming in falls back to the material's flat color.
  let ready = hasResource<texture_2d<f32>>(m.albedo);
  let albedo = select(unpack4x8unorm(m.baseColor).rgb, texel.rgb, ready);
  return vec4f(shade(albedo, in), 1.0);
}

A bad index can't crash anything, or read another page's memory. If slot i is out of range, empty, or holds a different type of resource, getResource returns a harmless default of the requested type instead. That's why the example samples first and decides with select after: sampling is always safe, and the sample stays outside any branch. WGSL wants textureSample there, because it picks a mip level by comparing a 2 × 2 block of neighboring pixels (Chapter 11), and inside a branch some of those neighbors might not be running.

What it costs

Freedom moves work from the driver to the application, and the proposal is mostly about keeping that safe:

• Lifetimes. The GPU may still be reading a slot from frames already submitted. table.remove(slot) hides it at once, but the slot is only handed out again by insert after that work has finished, so a streaming renderer can unload and load textures every frame without a race. table.update(slot, resource) places a resource in a chosen slot instead, under the same rule.
• Reads and writes. WebGPU guarantees that no shader reads a texture while another part of the same pass writes it. With bind groups, it checks this from what's bound. With a table it can't know which slots a shader will touch, so any slot whose texture is being written in the current pass is hidden, and reads as the default. Each resource also has a usage in the table, "readonly" (the default), "writable" or "none", set with setResourceUsage. Marking shadow maps "none" while they're drawn, then "readonly" for the main pass, saves the browser from guarding them after every pass.
• Divergent indices. Neighboring pixels in a quad may use different materials. Some hardware can only index the table with the same value across a group of threads, and handles different values by looping over them, which is slower where materials meet. Sorting pixels by material, as the visibility buffer does, keeps that rare.
• Memory. Bind groups tell the driver exactly which textures a draw needs. A table doesn't, so a scene that overflows GPU memory can no longer rely on the driver to swap the right textures in; the application has to keep its working set in budget.
• One table at a time. A pass sees a single table. Samplers have lower limits than textures on some hardware, so a scene keeps a handful of shared samplers rather than one per material.

Together, cluster culling, cluster LOD, a visibility buffer, GPU-built draw lists and bindless resources form the architecture of the most detailed real-time renderers today: the CPU barely touches the geometry, or the materials, at all.

• Indirect draws take their counts from a GPU buffer, so compute shaders can build the draw list; atomics append visible objects to it.
• GPU culling tests every object in parallel: against the frustum's six planes, and against a Hi-Z pyramid of furthest depths for occlusion. Two-phase culling builds the pyramid from what was visible last frame.
• Meshlets of ~100 triangles are culled individually, including by their normal cones.
• Virtualized geometry builds a graph of clusters by simplifying groups with locked borders; each frame every cluster decides alone, from two projected errors, whether it's in the cut, so the triangles drawn follow the pixels, not the scene.
• The visibility buffer stores only which triangle covers each pixel, and shades materials afterward, exactly once per pixel.
• Bindless resources put every texture in one table that shaders index by number, so one draw or dispatch can cover every material. WebGPU's version, the resource table, is a draft that hasn't shipped yet.

Chapter 35 fills the space between surfaces: fog, light shafts, the sky and clouds.

1. A frustum of your own. In the playground, change RULE so the visible dots form a slowly turning wedge, like a camera's view from the center.

   Hint

   Compare the angle of p with the view direction: dot(normalize(p), vec2f(cos(t), sin(t))) > 0.8. The test is the 2D version of a frustum's side planes.

2. The cost of culling. In the hero, freeze the culling camera, then drive away and turn back. Where are objects missing, and why is that fine when the cameras agree?

   Hint

   Everything outside the frozen view, or behind its buildings, was culled for the frozen camera. When culling and viewing use the same camera, exactly those objects couldn't be seen anyway, except right at the edges of the bounding volumes, which is why the tests are made conservative.

3. Why the furthest depth? The occlusion pyramid stores the furthest depth of each block, while Chapter 33's reflection pyramid stored the nearest. Why the difference?

   Hint

   Occlusion needs a guarantee that every pixel in the block has something in front of the object: the furthest depth is the weakest of them, so beating it beats them all. Ray marching needs to know where the ray might first hit something: the nearest depth.

4. Counting commands. In the bindless figure, set 24 materials. How many commands does the CPU record with bind groups, and with the table? What else does the CPU need to know, with bind groups, before it can record them?

   Hint

   With bind groups it's two per material in the list, a setBindGroup and a draw, and each draw needs to know where its material's run of the list starts and how long it is, numbers only the GPU's culling pass knows. A second compute pass can count and sort objects by material and write one indirect draw per material, but the CPU still has to record a draw for every material that might be visible. With the table it's two commands, whatever the scene.