Chapter 35 — Ordering, hazards, and the dispatch gap

status: polished · path: Muse Glimmer, pinned Muser tree

Prerequisites: Ch 2 (command buffers, encoders, dispatch — this chapter uses “closure,” “encoder,” and “commit” fluently), Ch 3 (unified memory and the buffer substrate), Ch 10 (the one-token graph whose dispatches we will count), Ch 34 (who submits). No GPU-synchronization background is assumed; the hazard taxonomy is built from zero.

Ch 34 ended with a submission: the batcher packs rows, a permit is acquired, one concurrent encoder records a whole token, one commit, one wait. That chapter treated the recorded tape as if pressing “record” and “play” were all there was to correctness. They are not. The moment more than one kernel writes GPU memory inside one command buffer — and Muser’s token graph runs hundreds of dispatch groups — you must answer a question the hardware will not answer for you on every path: when do this kernel’s writes become visible to that kernel? Get it wrong in one direction and you read stale bytes; get it wrong in the other and you stall the GPU for no reason.

This chapter does three things. It teaches the hazard taxonomy — RAW, WAW, WAR — with timelines small enough to check by eye. It walks Muser’s actual ordering tools, verified in source: one command buffer per token, tracked buffers by default, two barrier forms, a serial/concurrent encoder switch, and queue-ordered submissions instead of fences. Then it spends that vocabulary on the campaign’s most instructive measurement: the bounded one-token diagnosis that reconciled a +196 dispatch-closure gap into four named families and rejected the largest of them for changing logprobs beyond contract. That last part is a story about wanting something and not being allowed to have it, and we tell it that way, because the recurring question of this book — what may be moved without breaking the exactness contract? — has never had a sharper answer than “not these 104 groups.”

35.1 What one queue buys — and what it does not

Start with the question every ordering bug is an answer to: what has the platform already promised us, and what do we still owe it ourselves? Getting that boundary wrong in either direction is expensive, so it is worth drawing carefully before any Muser code appears.

Recall the Metal execution model of Ch 2. You record kernel dispatches onto a MTLComputeCommandEncoder, end the encoding, commit the command buffer, and the queue executes it. Two guarantees come free:

1. Within one serial encoder, dispatches run in recorded order. Whatever ordering problems exist, they are not scheduling problems: kernel B recorded after kernel A does not start before A on a serial encoder.
2. Command buffers on one queue complete in commit order. The engine leans on this for its teacher-forced lane — the comment at the submission site: “The queue serializes GPU work so token i+1 cannot race token i’s residual/KV, while the host encodes i+1 during i’s GPU interval” [crates/muser-engine/src/decode.rs:2151-2155].

So where is the problem? In memory visibility. A dispatch “completing” does not mean its writes are visible to the next dispatch — caches, and the freedom a concurrent encoder grants its dispatch closures to overlap, sit between the two. Ch 10 already showed the serving token graph uses one concurrent encoder with explicit barriers precisely because the four independent Q/K/V/gate matvecs are allowed to overlap [crates/muser-engine/src/decode.rs:5448-5458]. Ordering within overlap is the subject of this chapter. And even on a serial encoder, the fine print matters: Metal orders dispatch execution, but Muser’s own diagnostic notes teach that relying on implicit visibility instead of saying what you mean is how Ch 34’s clean design acquires invisible state.

35.2 The three hazards, from zero

Before the definitions, the stake. What does it cost to classify one of these memory relationships wrongly? Not a crash, usually. The failure is a number that is merely different — different enough to fail a bit-exactness diff weeks later, on someone else’s machine, with nothing in the log to point at. That is why the taxonomy below is worth memorizing rather than looking up.

A hazard is a pair of memory operations whose order changes the result. Three kinds can hurt you; classify every buffer relationship into exactly one of four buckets (Figure 35.1). Two kernels A (first) and B (second), one buffer X:

  RAW — Read After Write   (true dependency: B reads what A wrote)

     kernel A          ████ write X
     kernel B                              ▒▒▒▒ read X   ← must see A's bytes
                       ────────── time ──────────→
     Rule: A's writes must be visible to B. THE decode hazard: the token
     graph is a producer→consumer chain (norm → matvec → gate → …).

  WAW — Write After Write  (order of writers decides final bytes)

     kernel A          ████ write X = 1
     kernel B                    ████ write X = 2   ← final value must be 2
                       ────────── time ──────────→
     Rule: writers serialize. Rare on the token path (buffers have one
     writer per dispatch group) but lives in workspaces reused across
     dispatch groups and in-place kernels.

  WAR — Write After Read  (B overwrites bytes A still needs)

     kernel A          ▒▒▒▒ read X (old)
     kernel B                    ████ write X (new) ← must wait for A's read
                       ────────── time ──────────→
     Rule: B waits. Appears when a ring buffer's producer laps a slow
     consumer — Ch 15's ring rotations are the in-engine near-miss.

  RAR — Read After Read   (both readers: order irrelevant)
     NOT A HAZARD. No barrier, ever.

Figure 35.1: the hazard taxonomy as timelines. ████ = write, ▒▒▒▒ = read. The rule column is the whole discipline: name the bucket, then place (or omit) a barrier accordingly.

Worked micro-example, one buffer x, two dispatch closures on a concurrent encoder. Recorded order A then B, but the encoder may overlap them:

  A: x = x + 1        (read-modify-write — contains a RAW against itself)
  B: y = x * 2        (reads x)

  If B's read overlaps A's write:  y may see the old or the new x.
  This is a RAW hazard from A to B → a barrier between the closures is
  REQUIRED. Without one, the graph is correct only by scheduling luck.

Now the same pair with B: x = x - 1 instead: both closures write x, a WAW hazard — the final value of x depends on which write lands last, so the closures must serialize even though neither reads the other. And if A reads x while B writes it, you have WAR — on the decode path, think “the KV ring’s next row is also a row someone is still attending from.” Ch 15’s explicit origin_logical/origin_physical bookkeeping exists so that the writer and the readers name different rows and the WAR never fires.

One more non-hazard worth naming because the engine exploits it: reads of immutable data never need ordering. All four packed decode rows of Ch 34 read the same weights concurrently — RAR against the weight arena, zero barriers, by construction.

35.3 Muser’s actual ordering tools

So what do we actually reach for when a dependency has to be stated? It is worth asking suspiciously, because the ancestor answered the same question with a great deal of machinery. The Ferrite engine solved this problem with a compiled Vec<Op> program and a frozen per-model barrier plan (§35.8). Muser has no such machinery — there is no Op enum, no compiler pass, no plan. What it has instead is five explicit devices, each verifiable in source, and the discipline to place them by hand.

Tool 1 — one command buffer per token, one queue

The serving token graph is one command buffer: encoder opened, whole 52-layer graph recorded, end_encoding, commit, bounded wait [crates/muser-engine/src/decode.rs:5448-5460]. The packed group graph of Ch 34 is the same shape with four rows [crates/muser-engine/src/decode.rs:4920-4937]. Cross-token ordering is the queue’s commit-order guarantee (decode.rs:2151-2155); cross-graph ordering is the scheduler’s single-owner rule of [Ch 34 §34.2]. There are no MTLFence or MTLSharedEvent objects anywhere in the engine — a grep for fence construction across crates/muser-engine/src/ finds the word only in a comment [crates/muser-engine/src/metal/buffer.rs:10]. Where a boundary must be crossed, Muser ends the command buffer and waits — and even the wait is bounded: wait_for_completion parks the blocking call on a watcher thread with a condvar so “a hang becomes a logged Deadline error, not a frozen box” [crates/muser-engine/src/metal/context.rs:153-175].

Tool 2 — tracked buffers by default

Metal buffers can be created tracked (the driver inserts the visibility synchronization your dispatches imply) or untracked (it does not; you owe every dependency an explicit fence or barrier — see [Metal-PG]). Muser’s allocation path routes every buffer through one function:

#![allow(unused)]
fn main() {
// crates/muser-engine/src/metal/buffer.rs:7-14
fn shared_tracked() -> MTLResourceOptions {
    // Several accepted Muse paths still cross compute encoders (notably
    // target-hidden prefill/capture). Untracked resources are only valid when
    // every such dependency has an explicit fence/barrier. b9678d4 enabled
    // untracked mode globally before that contract existed and empirically
    // changed DFlash conditioning while leaving final greedy IDs unchanged.
    MTLResourceOptions::StorageModeShared
}
}

That comment is a war story compressed into a code block, and it is worth unpacking, because the bet it records is one we lost. The fork is a tempting one: tracking is not free — the driver re-derives, on every dispatch, the dependencies you already know — and switching it off is the standard way to buy that cost back. So an earlier revision (b9678d4) flipped the engine to untracked globally, before any explicit-dependency contract existed. What we expected was a cheaper encode with identical output, on the reasoning that the graph’s real dependencies were already written down as barriers anyway.

What we got was quieter than a crash and worse. The final greedy token IDs stayed identical — decode the fixture, read the text, and nothing at all looks wrong — but the DFlash draft’s conditioning, the hidden states it reads, changed. The lesson is the one this chapter keeps circling: a knob that sits nowhere near the arithmetic can still perturb numerics, here through scheduling overlap, and in an engine whose contract is reproducible bits “the tokens still match” is not a passing grade. So the decision was to revert and stay reverted: tracked is the default, and untracked mode is not re-openable until every crossing dependency carries an explicit fence or barrier.

One documented exception survives, and it is worth separating from the story above so it is not misread as a crack in it: multi-gigabyte KV planes may allocate uninitialized — and that is an initialization contract, not a tracking one; debug builds poison the bytes (0xDEAD) so a premature read is conspicuous, and the ring metadata guarantees every read row was written first [crates/muser-engine/src/metal/buffer.rs:246-277].

Tool 3 — two barrier forms

With tracking on, why barriers at all? Because the concurrent encoder still must be told where dispatch closures depend on each other, and because the engine wants dependencies to be auditable in source rather than implied. Form one, broad scope, is inserted automatically between closures:

#![allow(unused)]
fn main() {
// crates/muser-engine/src/decode.rs:6256-6266
impl EncodeTarget for GraphEncoder<'_> {
    fn before_dispatch(&self) {
        if self.concurrent && self.has_dispatch.replace(true) {
            // Broad buffer scope exactly matches llama.cpp's dependency reset.
            // Independent kernels are deliberately grouped into one dispatch
            // closure, so every closure boundary is a real graph dependency.
            unsafe {
                let _: () = objc::msg_send![self.encoder, memoryBarrierWithScope: 1u64];
            }
        }
    }
}

Every dispatch(command, |encoder| { … }) call routes through before_dispatch [crates/muser-engine/src/decode.rs:6273-6279]: on a concurrent encoder, each closure after the first is preceded by a whole-buffer-scope memory barrier — deliberately the same reset llama.cpp performs between its own kernel groups [crates/muser-engine/src/decode.rs:6259]. The cost of the broad form is that it also orders unrelated reads; which is why form two exists:

Form two is the targeted resource barrier, naming exactly the buffers a dependency flows through. The clearest example is on the decode attention route — KV store then attention, the RAW hazard made visible as two lines:

#![allow(unused)]
fn main() {
// crates/muser-engine/src/decode.rs:5660-5671
dispatch(command, |encoder| {
    self.kernels.encode_kv_store_f16(
        encoder,
        &self.activations.k,
        &self.activations.v,
        &plane.key,
        &plane.value,
        write_physical,
    );
    let kv: [&metal::ResourceRef; 2] = [plane.key.metal(), plane.value.metal()];
    encoder.memory_barrier_with_resources(&kv);
    self.kernels.encode_llama_flash_attn_decode_vec_f16(
}

Store writes K/V; barrier names K/V only; attention reads them. The scoped form has its own didactic comment on the splitk route, where the producer’s partials scratch (not the whole world) must reach the reducer: “Scope the dependency to that allocation instead of stalling every buffer used by the 52-layer command buffer” [crates/muser-engine/src/metal/encode/attn.rs:771-775]. Targeted barriers appear at each genuine RAW boundary — mask-before-attention and pad-before-vec on the llama routes, staged-shadow-before-attention on the SWA route, partials-before-reducer on the split routes [attn.rs:247, 256, 293, 316, 428, 558, 614, 632, 747, 783; decode.rs:4305].

Tool 4 — serial versus concurrent encoder, as a flag

Serial or concurrent — which should the prefill encoder be? The engine declines to answer permanently. It ships a default and keeps the loser reachable behind an environment variable, with the A/B history recorded in the comment beside it:

#![allow(unused)]
fn main() {
// crates/muser-engine/src/decode.rs:975-979
    // Prefill-only concurrent Q/K/V/gate and FFN gate+up. Decode already
    // groups those projections; serial prefill paid a launch tax on PP128.
    // `MUSER_SERIAL_PREFILL_DISPATCH` restores the previous encoder for A/B.
    concurrent_prefill_dispatch: bool,
}

Default: concurrent grouping of the independent projection GEMMs (the MUSER_SERIAL_PREFILL_DISPATCH flag at decode.rs:1331-1333 restores the serial encoder for exact A/B comparisons). The retained flag matters for the same reason the receipts do — a performance default you cannot switch off is a default nobody can re-measure, and the comment names the measurement that decided this one: the launch tax the serial encoder paid on the prefill benchmark. The same select-by-graph structure appears in new_prefill_graph_encoder, which also documents the command buffer’s reference contract — unretained references, “This matches pinned llama.cpp’s commandBufferWithUnretainedReferences contract” [crates/muser-engine/src/decode.rs:6102-6122].

Tool 5 — discipline as documentation

Count the tools and you notice what is missing: no barrier planner, no analysis pass, no proof type. The contract is maintained the way Ch 34’s isolation contract is — in the structure of the code and its comments. Independent kernels are grouped into one closure so that “every closure boundary is a real graph dependency” (decode.rs:6260-6261); the barriers you find are therefore exactly the hazard statements, no more. When a new hazard class appeared — the SWA staging shadow needing its staged bytes visible to the vec kernel — the fix was three named resources and a comment [crates/muser-engine/src/decode.rs:4300-4305].

That is the whole ordering model. Now spend it on the measurement.

35.4 The instrument, corrected first

On 2026-08-15 the campaign asked: where does the one-token decode deficit live? The instrument was MUSER_METAL_PHASE_PROFILE and its PhaseProfiler — and the first finding of the investigation was not about decode at all. The instrument itself was wrong in two ways. We went looking for the deficit and found the ruler bent first, which is why the note corrects the ruler before it presents any number [docs/decode-dispatch-gap-20260815.md].

What a “closure” is. The profiler counts calls to the Rust dispatch closure, each submitted as its own command buffer with a synchronous wait — it does not count raw dispatch_thread_groups calls. One qkvg closure contains four kernel dispatches but contributes one count. Production serving encodes the graph into one shared encoder with a single wait, so “diagnostic wall time minus reported GPU time” mostly measures the diagnostic’s own hundreds of submit/wait cycles. Any statement like “the gap is N dispatches of overhead” is therefore a category error; the honest unit is profiling closures. Put plainly, the profiler measures a graph that production never runs, and its unit is a boundary in our own Rust code rather than a launch on the GPU. Every count in the next two sections is in that unit: read them as boundaries we can name, never as kernels the hardware issued.

Two label defects had shifted timings. The tell was arithmetic that could not be true: more labels printed than samples taken. Production labels omitted lm_head (its time silently attributed to the following softcap label), and the legacy schedule declared a separate SWA kv_store in all 39 SWA layers although the implementation at position 2,048 combines KV publication and attention in one closure — 603 labels printed for 564 samples, “every legacy per-label timing after the first extra label was shifted.” The fix went past the two defects to the class: the profiler now derives labels from post-append ring state and aborts outright on a label/sample count mismatch, so the next such bug stops the run instead of quietly re-attributing time [docs/decode-dispatch-gap-20260815.md].

This is the book’s measurement culture in miniature, and it recurs in Ch 38: before a number can mean anything, the instrument’s own error model must be on the table.

35.5 The +196 reconciliation

With a trustworthy instrument we could finally ask the counting question and believe the answer: how much bigger is the production graph than the legacy schedule, and is the difference waste? The bounded one-token diagnostic counted 760 profiling closures for the production graph against 564 for the legacy schedule — a difference of +196. The fixture behind that count is worth stating once, because every number in this section and the next shares it: the pinned 16,756,681,056-byte target, a 2,048-token fixture, one teacher token.

The interesting part is not the size of the gap. It is that the gap reconciled exactly, with no unexplained remainder, into four families [docs/decode-dispatch-gap-20260815.md §Corrected closure-count diff]:

Table 35.1: the +196-closure reconciliation at position 2,048 [docs/decode-dispatch-gap-20260815.md]. Closures are Rust profiling closures, not raw Metal dispatches (§35.4).

Two structural reads of this table matter more than the arithmetic. First, the 52 splits and 39 staging groups are the price of the ordering and exactness disciplines you just learned: the store-then-barrier-then-attend sequence of §35.3 is what the 52 splits are, and the staging groups are Ch 16’s decision to reproduce llama’s reduction lanes rather than read a wrapped ring “mathematically equivalently.” Second — and the note says this in bold — “No repeated closure performing provably identical arithmetic was found” [docs/decode-dispatch-gap-20260815.md]. The gap is not waste shaped like duplicated work. It is boundaries, each of which exists for a reason.

35.6 What was landed, what was rejected

Naming four families is a diagnosis, not a cure. The question we actually cared about is narrower and harsher: which of them can be removed without changing a single bit of the output? So the note ran the candidates as experiments rather than as a plan — each on the fixture just described (pinned target, 2,048-token fixture, one teacher token, bounded phase diagnostic), and each judged by one gate, the SHA-256 of the full logit row [docs/decode-dispatch-gap-20260815.md §Landed and rejected reductions]:

Table 35.2: the reduction table. GPU times are single-run diagnostics; the SHA-256 is over the full logit row — the exactness gate [docs/decode-dispatch-gap-20260815.md].

The one landed change is the +1 family: remove the last-row copy. It is bit-exact by construction — no arithmetic is touched, one 6,656-element f32 copy simply stops happening — so it never had to argue with the gate at all. Its single-run GPU delta is −0.136 ms (−0.34 %). Its wall sample went up 4.380 ms, and that number we do not believe: it is the diagnostic’s own submit/wait noise, which is precisely why the instrument had to be characterized before this table could be read, and no wall claim was made [docs/decode-dispatch-gap-20260815.md]. A lesson in miniature: the only guaranteed-free lunch in the whole table was a bookkeeping copy, and it was worth a third of a percent.

The 104-group fusion is rejected on exactness, not on speed. This is the family we wanted most — the largest one, and the one whose removal looks like pure structure rather than arithmetic — so we took the swing. The available dual-norm fusion does collapse the separated norm boundaries into 655 groups, exactly as advertised. It also changes the logit bytes (Table 35.2, row 3), which ends the argument before performance is even discussed. The gate is not “faster.” The gate is “identical.”

The second swing was the careful one. Replace the four-SIMD-group rsqrt reduction with the pinned llama 32-SIMD-group reduction twice, preserve the intervening f32 device-memory publication and reread, and the bits survive (row 4). We expected that to be the answer: the same structural win as before, merely performed correctly this time. It was not. Its five-sample streamed serving result was 26.714 tok/s median (CV 0.079 %) versus llama’s 33.428, and the note’s own verdict is the sentence to carry away: “the exact reduction is useful but does not by itself close Stage A” [docs/decode-dispatch-gap-20260815.md]. The one-query GQA specialization said the same thing from another direction, measuring 28.290 tok/s median (CV 0.193 %), ratio 0.8463×, prefill 1.0366× — Stage A still 13.37 points below its 98 % decode bar at that moment [docs/decode-dispatch-gap-20260815.md]. What the two swings taught together is that closure count is not the currency the gap is denominated in: removing a boundary buys less than counting boundaries suggests it should. (How Stage A eventually closed — by changing the anchor itself, J0/J1 — is Ch 38’s story; the gap families survived it.)

The hybrid postmortem. Before either of those disciplined attempts there was a bolder one, and it is the one worth re-telling, because it failed in the most instructive way on offer. The idea was to take the whole prize at once: reuse the legacy retained-activation schedule and select the fast fused boundaries wherever they existed. The first look said it worked — the greedy token was preserved, so you could decode the fixture, print the text, and see nothing at all. Then we looked underneath the token, at the full distribution, and found that it had failed the public numerical contract. Every number in that postmortem is worth quoting [docs/decode-dispatch-gap-20260815.md §Rejected hybrid postmortem]:

• full-logit maximum absolute error: 4.6300888e-4; mean absolute error 1.5170774150033564e-4;
• normalized logprob maximum absolute error: 3.197146176834309e-4 — above the 1e-4 contract;
• 201,970 of 202,048 logits differed;
• the first KV difference: layer 1, value plane element 524,115, f16 bits 39,892 versus 39,893 — a single ULP flip in one value element, one layer in, propagating to a hundred thousand logits.

Read that last bullet slowly, because it is this book’s precision thesis in one measurement. A one-ULP difference in one f16 value — the seventeenth-bit wobble of a rounding-order change in the fused residual/norm chain — is invisible in the sampled token and fatal to the contract. Put the other way round: the argmax that picks a token is a lossy hash of the distribution, so a check written against the token cannot see a change in the thing the token was computed from. That is why the attempt “was removed rather than hidden behind a tolerance or shipped as an alternate route” [docs/decode-dispatch-gap-20260815.md], and why we kept the run that convicted it — evidence retained under muser-receipt://pinned-token-parity-20260814-v{3,4}/.

35.7 What the gap reveals about what an engine is for

Step back and state the tradeoff plainly. The available 104-group norm fusion is the single largest closure family. Fusing it changes normalized logprobs to 3.2e-4 against a 1e-4 contract. Therefore Muser keeps the boundaries. The ~3 % of GPU time that fusion-class changes might recover (40.330 → 39.274 ms for the exactable subset, Table 35.2) is not purchasable at that price — and the campaign’s own ranked list records the principle as its first item: “Implement an exact boundary fusion only if it reproduces the standalone reduction and store order bit for bit. The current 104-group fusion is a negative fixture, not a candidate” [docs/decode-dispatch-gap-20260815.md §Ranked remaining exact work].

This is the recurring question’s hardest answer. What does one token cost? Some of the cost is not removable — not because the techniques are unknown, but because the engine is for something: producing tokens whose full distributions can be diffed against a pinned reference, byte for byte, so that every other claim in this book (parity, handoffs, warm reuse, sessions) has a ground truth to stand on. An engine that quietly rounded differently under load could not host Ch 32’s bounded-drift gates — it would have already spent its credibility on its own scheduler. Bit-exactness beats throughput, and not sentimentally: it is the enabling asset for every measured claim Parts IV–VI made.

35.8 Ancestor contrast: the compiled barrier plan that was not ported

One question is still owed an answer, and it is the one a skeptical reader has been holding since the tools section: is hand-placing every barrier a choice, or just what happens when nobody gets around to building the tool? The ancestor is the control experiment, because it did build the tool. That is why this contrast belongs in the chapter rather than in an appendix — it decides how you should read every barrier you have just been shown, as deliberate design or as absence. The ancestor Ferrite engine ordered its decode differently, and the contrast is instructive enough to box.

The Ferrite design (lineage, not Muser). Ferrite compiled its decode graph into a Vec<Op> program at load time — routing decisions resolved once, then a pipeline of passes fused, reordered, and emitted a frozen per-model barrier plan of byte-range hazards, replayed per token; a sealed OverlapProof type made concurrency certification constructible only by the analyzer [ferrite-book Ch 21; ferrite-book Ch 22]. It ran its buffers untracked under that plan, buying a measured −2 to −4 % tracking cost back on the ancestor’s A18 Pro hardware [ferrite-book Ch 22]. Why none of it ported. The audit that planned this book verified the divergence in the Muser tree: there is no Op enum and no VM program — decode is hand-written encode methods, the shape of Ferrite’s legacy route, chosen for Muser’s fixed single-model, pinned-kernel discipline where route reasoning does not benefit from a program representation (weights and routes are constant after load). And the ancestor’s own corrections register later established that the comparison motivating its untracked mode was false — “llama.cpp uses untracked hazard tracking by default” is marked FALSE, the fabricated-claim finding that forced the ancestor book’s largest correction pass [ferrite-book CORRECTIONS-2026-07 §3a]. Muser keeps the lesson and not the machinery: hazard-tracking is default-on (§35.3, the b9678d4 reversal), and barriers are hand-placed statements rather than compiler output. What survives as genuine lineage is the taxonomy (§35.2), the scope-versus-per-resource barrier insight, and the measurement instinct that barrier overhead must be counted, never assumed [ferrite-book Ch 22].

The gentle summary: Ferrite’s design was a bet that a compiler could own ordering so the runtime could be free with it. Muser’s design is a bet that a small, fixed, audited graph does not need a compiler — only a taxonomy, five tools, and the discipline to write every dependency down.

35.9 Tradeoffs

Four choices in this chapter were live enough to have an alternative someone could reasonably have picked. Here they are with the evidence attached, so that the next person to reopen one starts from what was measured rather than from what seems obvious.

Tracked-by-default vs untracked-plus-plan. Tracking costs driver-side dependency analysis on every dispatch; the ancestor measured −2 to −4 % for turning it off [A18-neo, ferrite-book Ch 22] — Ferrite-lineage hardware, never measured on Muser’s M3 Ultra [unverified]. Muser pays the tracking cost because the one attempt at untracked mode “empirically changed DFlash conditioning while leaving final greedy IDs unchanged” [crates/muser-engine/src/metal/buffer.rs:10-12] — a silent-numerics perturbation, the exact failure class this engine exists to exclude.

The 104 groups: paid in boundaries, refunded in trust. The separated norm-boundary closures cost GPU time relative to a fused schedule (Table 35.2: 40.330 vs 39.274 ms for the bit-exact variant on the diagnostic fixture), and the hybrid that fused them aggressively breached contract at 3.197e-4 normalized-logprob error [docs/decode-dispatch-gap-20260815.md]. The refund is every exactness-gated claim downstream — most concretely, the parity matrices of Ch 38 that could not have been run against an engine with drifting boundaries.

Broad-scope closure barriers vs targeted resource barriers. The broad form is simpler and matches llama’s own dependency reset (decode.rs:6259); the targeted form exists where a broad stall would serialize unrelated work — “instead of stalling every buffer used by the 52-layer command buffer” [crates/muser-engine/src/metal/encode/attn.rs:771-773]. No measurement separates the two forms’ costs [unverified]; the split is driven by the shape of the dependencies (closure-granular on the token graph, allocation-granular inside attention producer/reducer pairs).

One-CB-per-token vs fences. Ending the command buffer and waiting is coarser than an MTLFence would be — the teacher-forced lane’s comment shows the trade was considered and taken for host-encode overlap instead (decode.rs:2151-2155). The compensating control is the bounded wait (context.rs:153-175): the coarser primitive cannot wedge a serving thread indefinitely.

35.10 What comes next

Ordering and hazards were the last invisible ingredient in the token’s journey — and the +196 reconciliation named one family, the 39 SWA staging groups, that belongs to a graph this book has only glimpsed: the prefill route, where whole prompts flow through batch-shaped kernels, wrapped rings force a staging shadow, and chunk boundaries yield to waiting decoders. Decode is a matvec story; prefill is a GEMM story; the same weights star in both. Ch 36 walks the second graph.

References

• crates/muser-engine/src/metal/buffer.rs:7-14, 246-277 — shared_tracked (quoted, with the b9678d4 incident) and the uninitialized-KV contract.
• crates/muser-engine/src/decode.rs:5448-5460 — one command buffer per token, concurrent encoder, explicit barriers.
• crates/muser-engine/src/decode.rs:6102-6122, 6256-6279 — unretained- references prefill command buffer; serial/concurrent GraphEncoder; the broad-scope closure barrier (quoted); the dispatch trampoline.
• crates/muser-engine/src/decode.rs:975-983, 1331-1333 — encoder-mode and fusion defaults with their A/B flags.
• crates/muser-engine/src/decode.rs:2151-2155, 5660-5671 — queue serialization comment; the store→barrier→attend RAW (quoted).
• crates/muser-engine/src/decode.rs:4297-4305 — the staged-shadow resource barrier on the SWA route.
• crates/muser-engine/src/metal/encode/attn.rs:247, 256, 293, 316, 428, 558, 614, 632, 745-783 — the targeted-barrier inventory; the scoped- dependency comment (quoted).
• crates/muser-engine/src/metal/context.rs:142-184 — bounded completion waits.
• [docs/decode-dispatch-gap-20260815.md] — read in full for this chapter: the instrumentation correction, the 760/564/+196 reconciliation (Table 35.1), the landed/rejected reduction table (Table 35.2), the hybrid postmortem numbers, and the ranked remaining work.
• crates/muser-engine/src/decode.rs:6143-6149, 6224-6236 — PhaseProfiler as the closure-counting instrument behind §35.4.
• [ferrite-book Ch 21], [ferrite-book Ch 22], [ferrite-book CORRECTIONS-2026-07 §3a] — the ancestor’s compiled-program/barrier-plan design and its corrections register (lineage/contrast only; A18 Pro numbers are ancestor context, never Muser results).
• Ch 16 §16.9 — the staging and split families from the attention side; Ch 38 — how the anchor change (J0/J1) closed what the fusions could not; Ch 40 — the norm-boundary fusion as a catalogued rejection.