Refitting a BLAS beneath wgpu

Why the BLAS must follow the bend

A shadow ray in Praxis walks the TLAS down to a BLAS and meets the triangles it holds. A BLAS is built from the vertices its Mesh has when it is built. When the bend pass rewrites a figure's vertices, nothing changes in its BLAS, so the figure would cast the shadow of the shape it was built in.

There are two ways to bring a BLAS to new vertices. A rebuild makes the tree again from the triangles, as if they were new. A refit, called an update in Vulkan and wgpu, keeps the tree: which triangles share a leaf and how the leaves group stay as they were, and only the boxes are recomputed, from the leaves up, to hold the triangles where they are now.

The world-design pre-mortem priced the rebuild: "A hundred 5,120-triangle BLAS take 7.8 to 9.3 ms of GPU to rebuild (measured [...]), 65 to 78% of ADR 0013's 12 ms". The owner's Budget for the crowd's acceleration-structure update is 1.0 ms of GPU and 0.5 ms of CPU per frame (the decision on map #17, item 14). A refit does less work than a rebuild, which is why M3 times it. What it gives up:

• The same triangles. A refit repeats its build's description: in the code's words, "the same buffers, offsets, counts, stride and formats". Only the vertex positions in those ranges may change.
• Built to allow it. A refitted Mesh's BLAS is built with ALLOW_UPDATE, which wgpu-hal hands Metal as MTLAccelerationStructureUsage::Refit (wgpu-hal/src/metal/conv.rs:486-487). Apple's documentation warns that Metal then builds the structure "more conservatively, which can reduce its intersection performance".
• Never compacted. Every other BLAS in Praxis is built with ALLOW_COMPACTION and compacted once built (#202). A refitted one stays as built.
• Its index stream kept. #208 loaded the bent figures with no index stream, since only a BLAS build reads one: its hand-off counts 6.2 MB of the 16 MiB index Arena saved. A refit reads the indices again in every frame, so a refitted figure keeps its stream.
• A tree chosen for another shape. The grouping was picked for the shape the BLAS was built in. As the triangles drift from it, the recomputed boxes grow and overlap, and a ray tests more of them. The crowd's figures are built upright and sway either way; the demo below shows the effect on a model.
• Scratch memory. A refit works in a scratch region of the size Metal asks for its update (refitScratchBufferSize, read at wgpu-hal/src/metal/device.rs:2088). On the M4 Pro, Metal asks for none for these BLASes (chapter 7).

A refit costs less to make than a rebuild and traces worse the further its triangles have drifted. The crowd sways back and forth through the upright shape its BLASes are built in, so the drift never grows past the sway's widest lean, a curvature of 0.35 rad/m, where the model's refitted tree stays within 7% of a rebuilt one at every level. No figure isolates what the drift costs the shadow rays on the M4 Pro: part 2 reports no shadow mask time for the rebuild form, and the 50 µs more that the shadow mask pass took in the refit form than with the proxies compares 100 BLASes of 5,120 triangles with 1,600 proxies (chapter 8). What part 2 did measure is that the refits alone cost more than the whole Budget.

Three crowd forms

Which way the figures reach the TLAS is a Render switch, RenderSwitches::crowd_form, with three values. The default, Proxies, is #208's form, unchanged. Refit is the new form. Rebuild is the same setup with each BLAS rebuilt through wgpu instead: it runs on any backend and gives the shapes a refit must match.

A fourth way to load a Mesh

#208 had three MeshUse values: drawn and traced, drawn alone, traced alone. Part 1 adds Refitted, "Drawn and traced, a Mesh the GPU bends whose BLAS is updated each frame from its bent shape: built to allow updates, never compacted". Being traced, it keeps its index stream. The load builds it with flags kept in one constant, which every refit repeats:

The figures become Moving Instances

The static TLAS is rebuilt only when its static Instances change, so it would keep each figure's box as loaded while the figure swayed out of it. In the update forms each figure therefore carries PreviousWorldTransform, which makes its Instance a Moving Instance: it joins the dynamic TLAS, which is rebuilt every frame over each BLAS's current bounds.

Refused off Metal

State::new checks the switch before it makes the Renderer, and fails with the reason, so a capture or a bench asked for the refit form on Vulkan or DX12 stops with "Refit is unavailable: [...]". BlasRefit::new checks again and logs; the review kept that second check as defence in depth.

The default form draws and traces as before; the PR's bench step found every counter in every scene equal to its Baseline. It changes one thing for Proxies: the dynamic TLAS's command buffer is now recorded after the bend pass instead of before it (chapter 5).

Three layers of wgpu

wgpu is layered. Praxis calls its safe API. Beneath it, wgpu-core validates every call, tracks each resource's state and lifetime and records the barriers between uses. Beneath that, wgpu-hal is one thin, unsafe implementation per backend. Metal's wgpu-hal can refit, as Vulkan's and DX12's can; wgpu-core never asks any of them to.

Why wgpu-core cannot refit

wgpu has an update mode in its API, AccelerationStructureUpdateMode::PreferUpdate, but wgpu-core 30.0.1 treats it as a request it cannot meet. For a BLAS it logs and builds:

wgpu-hal's Metal backend implements both modes. For Update it calls Metal's refit, refitAccelerationStructure:descriptor:destination:scratchBuffer:scratchBufferOffset:, from the source structure into the destination, which for Praxis are the same BLAS (wgpu-hal/src/metal/command.rs:1864-1873). ADR 0007 allows exactly this: "using wgpu-hal only for a subsystem that exceeds wgpu's safe API".

The raw command encoder

To record wgpu-hal commands, Praxis asks a wgpu command encoder, which records a command buffer, for the wgpu-hal encoder beneath it, with CommandEncoder::as_hal_mut, which lends a &mut hal::metal::CommandEncoder to a closure. wgpu's docs set its terms: the raw encoder "must not be manually destroyed", "You must not end the command buffer; wgpu will do it when you call finish", and "The wgpu command encoder must not be interacted with in any way while recording is happening to the wgpu_hal or backend command encoder". wgpu-core enforces the last: "Mixing the wgpu encoding API with access to the raw hal encoder via as_hal_mut is not supported", and it tracks which API an encoder uses "to detect and reject invalid usage" (wgpu-core/src/command/mod.rs:625-627).

So the refits get an encoder of their own, labelled "BLAS Refit Encoder", which only ever sees wgpu-hal. wgpu finishes it into a command buffer like any other, and Praxis submits it right after the frame's first commands (chapter 5).

Guards and the snatch lock

A refit needs several wgpu-hal objects at once: the vertex Arena, the index Arena, the scratch buffer and every refitted BLAS. Each wgpu resource lends its wgpu-hal object through as_hal, which returns an as_hal guard. The guard of a buffer or a BLAS holds an Arc of the resource and a read of its device's snatch lock, the lock that stops a resource's raw object from being destroyed ("snatched") while someone reads it (wgpu-core/src/as_hal.rs:59-81). A thread may not hold it twice. In debug builds wgpu-core checks, and panics:

So the refit cannot hold four guards at once. Instead, unguarded reads the object's address through each guard, then lets the guard go, lock and all, and keeps the reference:

Letting the guard go also drops its Arc, so the guard no longer keeps the resource alive. Praxis does: the Mesh pool holds the Arenas, BlasRefit holds the scratch buffer, and the refit clones every BLAS handle into a local handles list for the length of the recording.

Who keeps the resources alive on the GPU

wgpu-core keeps a resource alive until the GPU has finished every command buffer that uses it, but it learns what a command buffer uses only from the commands recorded through it. The refit's command buffer was recorded through wgpu-hal, so wgpu-core knows nothing of it. What holds its resources is Metal itself: wgpu-hal's Metal backend makes its command buffers with retained references by default (retain_command_buffer_references: true, wgpu-hal/src/metal/mod.rs:383-393), and a Metal command buffer made that way retains every resource its commands name until it completes. The review caught the first SAFETY text crediting wgpu-core with this, and the PR corrected it.

The unsafe block, invariant by invariant

An unsafe block tells the compiler "trust me"; its SAFETY comment says why. The ticket asked that the block state "every invariant it relies on". The PR's description lists seven; the merged comment has eight, adding the rule about the raw encoder from chapter 3. They are taken here in the code's order.

1. Every handle is Metal's

   // - Every handle is Metal's: `new` refuses the refit off Metal.

   The block reinterprets wgpu objects as hal::metal types. BlasRefit::new returns None for Refit on any other backend, so the block never runs there. Each as_hal::<Metal>() also returns None for a foreign object, and the block then refuses and logs once instead of recording.

2. Every wgpu-hal object outlives its use past its guard

   // - Every wgpu-hal object outlives its use past its guard
   //   (`metal::unguarded`): the Mesh pool, this refit and `blases`
   //   hold their handles throughout, nothing destroys them, and only
   //   the main thread calls into wgpu meanwhile (ADR 0016).

   This is the promise unguarded asks for (chapter 3). Its last clause matters because the snatch lock's job was to stop another thread destroying a resource mid-read; with the guards gone, the promise holds because only the main thread touches wgpu while the refit records.

3. The encoder is wgpu-hal's alone

   // - The encoder is used through wgpu-hal alone, never ended here,
   //   and finished by wgpu; wgpu-core panics on a mix of the two.

   wgpu's own terms for as_hal_mut, from chapter 3. The refit's encoder records nothing through wgpu, and encoder.finish() ends it.

4. The GPU's use ends before any object is freed

   // - The GPU's use ends before any object is freed: wgpu-core never
   //   sees what this command buffer uses, but wgpu-hal's Metal command
   //   buffers retain every resource their commands name until they
   //   complete (its `retain_command_buffer_references`, on by
   //   default).

   The CPU-side promise (invariant 2) covers recording. This one covers the GPU running the commands after the CPU has moved on: if a BLAS were dropped next frame while a refit of it still ran, Metal's retained reference would keep it.

5. Each refit repeats its BLAS's build

   // - Each refit's descriptor is the one wgpu-core built the BLAS from
   //   (`MeshPool::load_for`), with its flags, which allow updates: the
   //   same buffers, offsets, counts, stride and formats.

   A refit is only defined over the triangles the structure was built from. metal::triangles rebuilds wgpu-hal's description of them from MeshBlas, the record load_for kept: the same vertex and index Arenas, first vertex, vertex count and format, the 28-byte CookedVertex stride, u32 indices from first_index, and the geometry's flags. REFITTED_BLAS_FLAGS is the one constant both the build and the refit read.

6. Each refit has its own scratch region

   // - Each refit has its own scratch region, of the size Metal asks
   //   for its update, in a buffer any usage makes on Metal. The next
   //   frame's refits reuse the regions, and the frames before them may
   //   still run: Metal tracks hazards on wgpu-hal's buffers, scratch
   //   included, so they wait.

   All the refits are in one call and may run side by side, so no two share scratch. metal::scratch asks wgpu-hal for each BLAS's update_scratch_size, rounds it up to SCRATCH_ALIGNMENT, 256 bytes, "a margin over the one-byte alignment wgpu-hal reports for Metal", and lays the regions end to end in one buffer made at the first refit. A frame in flight may still be refitting when the next frame's refits are recorded into the same regions. Metal's hazard tracking sees the scratch buffer named in both and orders them. The buffer joins the Renderer's GPU memory through BlasRefit::gpu_memory_bytes.

7. The bend comes first; the build waits

   // - The bend pass's writes to the vertex Arena come first in the
   //   queue, and Metal tracks hazards on wgpu-hal's buffers, so each
   //   refit reads the bent shapes; the TLAS build that reads the
   //   BLASes waits for this command buffer through `metal::Order`, or
   //   on the CPU.

   Two orders, of two kinds. Before the refit, Metal's hazard tracking is enough: the bend pass writes the vertex Arena and the refit names it. After the refit it is not: the TLAS build reads the BLASes without naming them. That half is chapter 5.

8. wgpu-core never learns of the refit, and need not

   // - wgpu-core never learns of the refit, and needs not: it changes no
   //   BLAS's size, handle or build index, so the TLAS built next still
   //   validates against it, and wgpu-core's barriers are no-ops on
   //   Metal.

   wgpu-core keeps, per BLAS, the index of the build that last wrote it. A TLAS build fails if one of its BLASes was never built, and a pass that uses a TLAS fails if one of its BLASes was built after it (wgpu-core/src/command/ray_tracing.rs:495-534). A refit in place leaves each BLAS as wgpu-core last saw it, built at load, so the dynamic TLAS built after it still validates. And the barriers wgpu-core would have recorded around a write it knew of do nothing on Metal: wgpu-hal's place_acceleration_structure_barrier for Metal is an empty function (wgpu-hal/src/metal/command.rs:1877-1881). The module's own summary: "a refit changes only what the BLAS holds".

Ordering on Metal

Metal orders GPU work that writes and reads the same resource, but only resources a command names. wgpu-hal 30.0.1 builds a TLAS from instance descriptors that name each BLAS by its resource ID, so Metal never sees the build read them, and nothing orders a TLAS build after a BLAS write. #244 found this, for BLAS builds and compaction at load (gfx-rs/wgpu#9215, "Metal BLAS -> TLAS builds need synchronisation", open). A refit each frame followed by the dynamic TLAS's build is the same pair.

#244 measured what a missing order does: "a TLAS built right after a BLAS build or compaction copy, with no wait between them, missed the BLAS in 3 to 23 of 80 iterations while other devices loaded the GPU" (its heads-up on #209). GPU Ordering Explained, volume 13, tells that story. The heads-up gave #209 two ways to order the pair:

• A CPU wait. Submit the refit, then poll its submission. That is what compact_blases does at load, but every frame "it would stall the main thread on the GPU, against criterion 8's 0.5 ms of CPU".
• wgpu-hal events. Metal's Queue::add_signal_event and add_wait_event can signal after the refit's command buffer and wait before the build's. In #9215's thread, one application had dropped this route because staging is queue-wide: whatever submit comes next takes a staged event, whether another thread makes it or wgpu-core makes one of its own, as it can when a buffer with pending writes is mapped. In Praxis only the main thread submits (ADR 0016), and Order::stage's SAFETY comment, below, states that nothing submits in between.

Part 1 takes the events, and keeps the CPU wait as a fallback.

The frame in two submissions

Until #289 the dynamic TLAS's build went in a command buffer submitted before the frame's, and the bend pass was recorded in the frame's encoder, after it. In the Proxies form the build's command buffer is still submitted first; only its recording now follows the bend pass. A refit must sit between the bend and the build, so in the update forms the frame is now cut after its bend pass. Everything recorded so far, the "head", goes in a submission of its own with the refits.

BlasRefit::submit does the cut's work: in the rebuild form it records the rebuilds into the head itself, through wgpu; in the refit form it submits the head and the refit command buffer together; either way it stages the signal before and the wait after:

A shared event, staged for the next submit

A shared event is a Metal counter on the GPU: one command buffer signals it to a value as it completes, another waits until it reaches that value. metal::Order gets one from a wgpu-hal fence (raw_shared_event) and counts the values it has signalled, one more each frame.

wgpu-hal does not take an event per command buffer. add_signal_event and add_wait_event stage it for the queue, and "any add_wait_event call is consumed by whichever Queue::submit runs next on this queue" (wgpu-hal/src/metal/mod.rs:551-556). That is why submit stages the signal right before submission 1 and the wait right after it, and why Order::stage's SAFETY says "only the main thread submits (ADR 0016), and Renderer::render submits nothing between the update's submission and the dynamic TLAS's build". The signal is drained onto the submission's last command buffer, the refits' (since part 2, when the refit is timed, the marker's, which follows them; chapter 7), "after the submit's own completion signal, so a foreign API waiting on (event, value) observes the wgpu work as done".

Why strict event sync

A staged wait alone is not an order. By default wgpu-hal encodes it on an internal command buffer of its own, committed ahead of the submission's, and its docs say what that leaves open:

The dynamic TLAS's build shares no Metal-tracked resource with that internal command buffer, so Order::new turns on strict event sync. Then every command buffer, as it starts encoding, first waits on an internal relay event, at the value the next submit will release; and each submit's internal command buffer releases the relay only after the staged waits it carries (wgpu-hal/src/metal/mod.rs:507-514, and begin_encoding in command.rs). The build's command buffer is encoded after submission 1, so it is held until value n is signalled, which happens when the refits complete.

Strict mode has a price and no off switch: "one extra encodeWaitForEvent per command buffer plus one extra internal command buffer per submit", and "once enabled, the queue stays in strict mode for its lifetime". So BlasRefit makes an Order, and turns strict mode on, only in the update forms on Metal. Part 2 measured the price: nothing measurable (chapter 8).

The fallback: wait on the CPU

If Metal makes no shared event, Order::new returns None and BlasRefit::new logs "no shared event: each frame waits on the CPU for its BLAS updates". Each frame then polls submission 1 to completion before recording the build, as compact_blases does at load. It is correct and slow: the main thread waits on the GPU in every frame.

How the test knows

The ticket's acceptance criterion: "a capture test shows shadows from refitted BLASes matching those from a full rebuild". The rebuild form is the oracle: wgpu-core builds those BLASes from scratch, through the safe API, from the same bent vertices. If the refit reads the bent shapes and the build waits for it, the two forms trace the same triangles and cast the same shadows, pixel for pixel.

Why four devices at once

A missing order shows only sometimes, "and more often under load (#244)", in the comment on DEVICES, which is 4. Each thread makes its own wgpu device and captures the refit form, so the four contend for one GPU, as the other devices did in #244's 80-iteration loop.

Red first

A test earns trust by failing when the thing it checks is broken. The PR reports five experiments:

Passing is not a guarantee

Strict event sync with no event wait passed 10 runs of 10. It holds the build until its own submission's internal command buffer runs, and that usually happens to be after the refits; wgpu-hal's docs above say single-stream work "often" serializes so. The PR's words: "It happens to serialize the work but guarantees nothing, so this PR relies on the event wait."

Ten passes say little about a race. If a race lost one run in 20, ten clean runs in a row would still come up about 60% of the time (0.95¹⁰ ≈ 0.60, this guide's arithmetic). A test can show that an order is missing, as the 19 red runs did; it cannot show that one is present. What makes the merged form correct is the API's promise: the build's command buffer waits on the GPU until the event reaches a value only the refits' completion signals. The 20 green runs at 8 devices are the check that the promise is kept, not the reason to believe it.

The review

The local review found 24 items: 15 fixed, 4 declined, 5 deferred to part 2. Among the fixes: the SAFETY text that credited wgpu-core with keeping the refit's resources alive (chapter 3), the scratch-reuse and staged-events invariants, the fence leak in Order::new, and a refused refit that now logs once, with its real reason.

Timing the refit with markers

Part 1 made the refit correct and timed only its CPU stage, blas_refit, which no bench figure counted. On the GPU, in the update forms the build was not bracketed, and the frame's span ran from the bend pass, holding the build and the gap between submissions. Part 2 times the refit on both processors, so it can be held to the Budget's 1.0 ms of GPU and 0.5 ms of CPU. The GPU half needs a marker, and the marker has to be built with care.

Where Metal reads the clock

Praxis asks wgpu for a timestamp at the start or the end of a pass. wgpu can also write one between commands, behind its native TIMESTAMP_QUERY_INSIDE_ENCODERS and TIMESTAMP_QUERY_INSIDE_PASSES features, but wgpu 30 documents those on Metal for AMD and Intel GPUs, not Apple's, so on the M4 Pro a pass's ends are all there is. On Metal, wgpu-hal 30.0.1 offers timestamps only where the device samples counters at stage boundaries (MTLCounterSamplingPoint::AtStageBoundary, wgpu-hal/src/metal/adapter.rs:676-681), and attaches the query set to the compute pass's encoder as a sample buffer with a start-of-encoder and an end-of-encoder index (command.rs:1700-1745). The clock is read where the encoder starts and ends on the GPU. #204, timing the dynamic TLAS's build, found what that implies (docs/performance.md, "Bracketing on Metal"):

• An empty pass times nothing. "Empty compute passes write no timestamps". A marker has to dispatch real work.
• An unrelated pass brackets nothing. "passes that touch no TLAS do not bracket the build, which Metal runs alongside them; markers that query the TLAS do." Metal orders work through hazard tracking, so a marker has to name a resource the timed work writes, or it may run beside that work rather than after it.

The dynamic TLAS's markers each run one ray query against the TLAS, on either side of its build. The refit needed a marker of its own.

The refit's marker reads the scratch

The refits are recorded through wgpu-hal alone, on an encoder that takes no wgpu call once wgpu-hal has recorded into it (chapter 3), so no wgpu pass, and no timestamp write, can go in their command buffer. #290 puts the marker in a third command buffer of the refit's submission: one invocation, mark_blas_refit, that reads the first word of the refits' scratch buffer and stores it where nothing reads it. Every refit names that buffer as its scratch, so Metal runs the marker once the refits are done, and the timestamp at the marker's start ends their time; the three-times check below shows it waiting. The regions are in fact empty on the M4 Pro: Metal asks for no scratch to refit these BLASes, so the buffer is its 256-byte minimum (this guide's check of the sizes Metal reports, on macOS 27.0.1).

The marker writes only its start: marker_after gives it the refit's second query slot as beginning_of_pass_write_index and no end (src/render/readback.rs:463-469). Nothing writes a timestamp where the refits begin, so their time runs from the end of the bend pass, the last timed pass in the head command buffer before them, to the marker's start. That span includes what Metal spends between command buffers: docs/performance.md puts it at "some 15 to 25 µs of fixed cost".

Checking the timer: three times the work

A timer that measures the wrong thing still returns plausible numbers. So part 2 tested it with a local patch that recorded every refit three times. blas_refit read 3,332 and 3,333 µs against 1,127 and 1,131 once, "and the GPU frame grew by the same 2.2 ms: the marker waits for the refits". Three times the work read 2.96 times the time (3,332 / 1,127, this guide's arithmetic), and the frame grew by about the difference, 2,205 µs. A marker that ran beside the refits would have read the same small number whatever they did.

The frame's GPU time, submission by submission

In the refit form a frame's GPU work arrives in three submissions: the head, the refits and their marker; the dynamic TLAS's build, bracketed by its markers; and the rest of the frame. Between them the GPU can sit idle while the CPU records. frame_times, now a pure function with a unit test, sums the spans and leaves the gaps out:

• The build counts once. In a frame whose refit went untimed, the span from the bend pass holds the build's submission, so the build is not added again: the review's fix, commit 5201f26, carried by FrameShape (src/render/renderer.rs:676).
• The oracle stays untimed. The rebuild form times neither its rebuild nor its build: its rebuilds through wgpu "would fail the first marker's query", since wgpu-core checks that query against the BLASes of the last build.
• The CPU side. The blas_refit stage, which records the refits and submits them with the head, now counts in dynamic_tlas_cpu_microseconds beside the build's recording and submit (src/bench/mod.rs). The head's submit is in it too, which the proxies form makes with the rest of the frame, so the refit form's CPU "reads a little high".
• One Baseline change everywhere. The new timed pass adds two queries a frame, so the readback's buffers, and every scene's gpu_memory_bytes, grow by 112 B; PR #290 re-recorded every Baseline.

The result, and M7's form

The bench: Apple M4 Pro, wgpu 30.0.1 on Metal, release builds of part 2 on 66b5b57, the crowd scene at 2560×1440 with 4x MSAA, 248 warm-up and 1,800 measured frames, four rounds alternating the forms. docs/performance.md says "GPU medians agree within 1% across rounds"; in the table below they agree within 1.1%, apart from form (a)'s build, 253.2 to 259.2 µs, 2.4% apart (this guide's arithmetic). The update is the dynamic TLAS's build in both forms, plus the refits in form (a); the bend pass, which both forms run, is outside it, at 21 to 27 µs.

• Proxies: within both. The build takes 597 to 603 µs, 60% of the GPU Budget; the CPU 319 to 326 µs at the median, its p95s 397 to 466.
• Refit: over the GPU Budget. The refits alone take 1,114 to 1,123 µs for 100 BLASes of 5,120 triangles, 512,000 in all: "about 2.2 ns a triangle with some 15 to 25 µs of fixed cost" (the three-times check gives 2.15 ns a triangle: 2,205 µs more for 1,024,000 more triangles, this guide's arithmetic). The update takes 1.37 to 1.38 ms, over the Budget by 0.37 to 0.38 ms. The CPU is within it at the median, 393 to 451 µs, and over it at p95, 504 to 628.
• A smaller build. Over 100 entries rather than 1,600 proxies, the build takes 253 to 259 µs. Part of the gap is strict event sync's doing, below, so form (b)'s build figure is an upper bound.

Strict event sync, measured

Chapter 5 quoted wgpu-hal's price for strict event sync: a wait per command buffer and a command buffer per submit. Part 2 forced it on in the proxies form, with a local patch, three rounds alternating. Every figure moved within noise or fell:

• The update's CPU: 297–315 µs to 287–292.
• The CPU frame: 727–756 µs to 711–723.
• The GPU frame: 3,301–3,308 µs to 3,181–3,257.
• The build: 593–610 µs to 483–520, because "its first marker no longer starts while the frame before still runs".

Beyond the update

• CPU frame. Form (a) places no proxies, so its CPU frame is lower, by 90 to 200 µs across the rounds.
• Rays. They traverse 100 BLASes of 5,120 triangles rather than 1,600 proxies: the shadow mask pass takes 50 µs more.
• Memory. Form (a)'s GPU memory reads 47.4 MB more. Its 100 BLASes, which form (b) does not have, are uncompacted (chapter 1), and Praxis counts them at the 94 B a triangle measured on the M4 Pro: 48.1 MB (512,000 × 94 B). Form (b)'s per-frame staging, which grows with the 1,600 Instances it writes each frame against form (a)'s 100, holds 0.7 MB more (both this guide's arithmetic), and the refits' scratch is one 256-byte buffer (chapter 7). The index streams form (a) keeps add nothing to this figure, which counts each Arena's whole buffer in both forms.
• Shapes. At 1280×720 the two forms' images differ in 5,398 of 921,600 pixels, all along the shadows' outlines: the proxies' shadows are a little thinner.
• The oracle. The rebuild form's rebuild and build are not timed; in one run its GPU frame spanned 9,303 µs, waits between submissions included, and its update's CPU took 709 µs.

The decision

PR #290 ended by asking the owner to pick M7's form, recommending the proxies. The owner picked them, as recorded on map #17:

Why, in the recommendation's terms:

• Proxies fit both Budgets with 40% of the GPU's to spare, run on every backend and need no unsafe. Their costs: the thinner shadows, 16 of the dynamic TLAS's 10,000 entries per figure, and the heap allocations wgpu-core makes per entry.
• Refit is 0.37 ms over the GPU Budget, runs on Metal alone, and needs wgpu-hal plus strict event sync.
• The refit stays. Behind its Render switch it is the reference for the shapes. tests/crowd.rs keeps its unsafe path and its order tested on Metal, and praxis bench <dir> --crowd-form refit still times it, though no Baseline holds its figures: PR #290 dropped the one it recorded, since verify benches each Baseline's name as a scene.

Reading list

Grouped by source, each with what it is good for here. Repo links are pinned to 66b5b57, the commit that merged PR #289, or to 0534eb6, PR #290's head when this guide was written, whose src/ and docs/ it merged unchanged; wgpu's to its v30.0.1 tag, the version the lockfile pins.

Glossary

Where CONTEXT.md defines a term, the definition here and in the tooltips is its own at 66b5b57, word for word; PR #290 changes none of them.