SAVEPOINT HOT PATH — where the +3.59% went, and the route that removes it
=========================================================================

THE INSTRUMENT WAS ALREADY RIGHT, AND NOBODY READ IT
----------------------------------------------------
  savepoint_spill_fence_total            24
  savepoint_spill_fence_not_ready_total  24
  savepoint_spill_fence_residual_ms      238.65 ms / 24  =  9.94 ms per fence
  savepoint_write_ms                     285.02 ms / 24  = 11.88 ms per write

Eight cold turns of the receipt shape (docs/receipts/savepoint-four-
integration-2026-09-09.txt: --max-batch 8, c=1, 6300-token prompt, 512
tokens ignore-eos, fresh nonce per rep), ARBI_SAVEPOINT_ENABLED=1, box 206
GPU 1. Every fence the run enqueued found the savepoint's device→host copy
still in flight, and the compute stream sat on it for ~10 ms each time. Three
fences per cold turn is ~30 ms of a 1.25 s TTFT — 2.4% — and the histogram
that says so has been in the tree since the spill stream landed. The previous
receipt measured the subsystem eight ways and read none of them; this one
starts from the counter.

THE LINE-ITEM TABLE (ON minus OFF, per cold 6300-token turn)
------------------------------------------------------------
Measured on the integration tip (22e32f6ec) before any change: 4 warm-up +
8 measured reps per arm, ARBI_STEP_PHASE_PROBE=1 in both arms, one Kineto
trace of a cold turn per arm, one 30 s py-spy per arm.

  #  item                                              instrument                       ms/turn
  1  whole-row 154.9 MB snapshot at each governor      spill_fence_* histograms above;   25.1 (trace)
     crossing (2048/4096/6144), fenced at the next     trace: gaps of 9.34/9.27/6.44 ms  29.8 (hist)
     step's first kernel                               each after the previous step's
                                                       last DtoD, before the next
                                                       step's first pinned HtoD;
                                                       spill stream 288 copies /
                                                       464.8 MB / 35.5 ms
  2  inter-step host work at the crossings: issuing    probe commit_snapshot_us 911 us;   +6.8
     the snapshot (96 copy launches, ring hand-out,    seed_slate_built->seed_drains_done
     chain digest, put) + the seed-path drains         +2.59 ms/step; trace inter-step
                                                       gaps 11.05 vs 4.25 ms
  3  emit-specialised chunk kernel on every launch     _h_emit 22.25 vs _h 21.55 ms       +0.7
                                                       over 192 launches
  4  unconditional per-layer conv staging              py-spy: 5.1% of the launch          ~0 visible
     (stage_conv_window)                               thread's busy samples; mid-forward  (host only)
                                                       gaps 6.1 (ON) vs 7.6 (OFF) ms
  5  promoter prefault holds the GIL (mmap.write)      py-spy: 187 of 292 promoter         ~0 in mean;
                                                       samples; CPU: one 155 MB region     spread (sd
                                                       stalls a Python thread 88.5 ms in   6.2 vs 2.7)
                                                       17 slices <= 5.6 ms; ctypes.memset
                                                       0 stalls
  6  admission (page hashes, best_match, probe_reach)  not visible in py-spy at 6.3k       < 0.1
  ----------------------------------------------------------------------------------------------
     sum of 1+2+3 (trace)                                                                  32.6
     wall-clock turn-0 TTFT, ON 1.2526 s vs OFF 1.2228 s (n=8)             +29.9 ms, +2.44% +-0.39%

Cold-turn window of the two traces (first 1300 ms): ON idle 83.6 ms, OFF
53.9 ms; the difference is items 1 and 2 and nothing else. Decode over the
same boots: +0.43% +-4.1%, unresolved at that n, as before. Turn 1 (resume):
0.176 vs 1.294 s, -86.4%.

The receipt's +44 ms is 14 ms more than this boot pair shows, and no line
item here accounts for it. The A/B/A/B below is the number that stands; the
difference between it and the receipt is boot-to-boot spread the receipt's
own interval did not contain.

THE OBVIOUS OPTIMISATION, AND WHY IT WAS WRONG
----------------------------------------------
The steer was to stop writing at the governor crossings and write once, at
the prompt's end — three whole-row copies per cold turn become one, and the
reasoning was that the ring's in-place reclaim (arbicity/arbi-serve#2296)
destroys each intermediate entry the moment the next write lands, so they
buy nothing. They do not get destroyed. promote_total read 32 of 32 on the
profile boot: the promoter copies every crossing entry to pageable RAM
within milliseconds, before the reclaim, and it survives as a
dominated-but-resident entry that best_match_hashes still returns. That is
what a SHARED SYSTEM PROMPT resumes from — a conversation sharing only the
first 4096 tokens of a 6300-token prompt matches the 4096 entry, and matches
nothing if only prompt ends are written — and it is what the retention
harness's own shared_prefix scenario depends on. The crossing writes are
retention with real coverage; what was wrong with them was their ROUTE (a
whole row copied after the step, fenced in front of the next one), and the
route is what this change replaces. Policy unchanged, nothing traded.

THE BUSY-RETRY LOOP NOBODY WOULD HAVE FOUND BY MEASURING HARDER
---------------------------------------------------------------
With the disk tier on, the flush band's worker ran its idle cadence — walk
the band, take the store lock, try a flush — twenty times a second for the
whole life of the boot, and the store's byte cap (1 GiB, ~7 entries) meant
the band always had entries it could not flush: 4987 attempts deferred for
budget and 2546 races lost in one six-minute arm, all Python under the GIL
beside the launch thread. The promoter had the same shape at 200 Hz. Both
are event-driven now (a write wakes them; an empty pass doubles the sleep
to a cap), and their removal is what took decode from a resolved +0.50% to
no effect detected.

WHERE THE FRAMING WAS WRONG
---------------------------
The brief assumed the cost was paid DISCOVERING ABSENCE — a walk, a hash, a
probe on a turn with nothing to find. It is paid WRITING: three whole-row
snapshots per cold turn, each 11.9 ms of PCIe, each fenced in front of the
next prefill step. The admission side is under 0.1 ms at this prompt size.
Two items on the brief's list (encode_tokens_u64, probe_reach) are real at
128k and noise here; they are on perf/admission-obvious and claimed there.

THE IDLE-GAP INVENTORY, PRE-EXISTING AND NEW (gaps >= 50 us)
------------------------------------------------------------
prefill window (~1.25 s), OFF -> ON:
  A  savepoint fence (prev step's last DtoD -> next input HtoD)   0.4 -> 25.1 ms   NEW
  B  step boundary host (commit/schedule/build -> input HtoD)     4.5 -> 11.5 ms   +6.9 NEW
  Z  small bubbles inside the compiled forward (<= 0.46 ms)       7.7 -> 10.6 ms   +3.0 NEW
     (the new ones sit between _h_emit and chunk_fwd_kernel_o:
      the per-layer conv staging on unarmed steps)
  D  after input HtoD, before the first kernel                    3.5 ->  3.0 ms   pre-existing
  C  first decode steps' page-meta boundary                       2.3 ->  2.3 ms   pre-existing
  E  int8 prefill a_prep -> exl3_i8_gemm launch                   1.25 ms          pre-existing
  total                                                          20.1 -> 53.2 ms
Also pre-existing, in both arms, and on TTFT's critical path: the probe's
fused_forward_done -> fused_prefill_committed phase (the seed tail that
emits the first token: sampling, drafter seed, detok/SSE) is 12.5 ms of host
time per turn in OFF and ON alike, of which the trace shows ~3.3 ms as GPU
idle at the prefill->decode transition (elementwise -> HtoD 1.68 ms,
unrolled -> page-meta 1.36 ms). Not this subsystem; listed for the
follow-up.
decode window (~4.5 s, ~150 MTP steps): 191.6 (OFF) vs 189.7 (ON) ms — none
of it moves with ARBI_SAVEPOINT_ENABLED. Eager glue between elementwise
kernels 48.5 ms (561 gaps), small other 48.4 (740), per-step input-HtoD
boundary 39.7 (536 x ~74 us), DtoD -> fused RMSNorm 21.6 (160), around exl3
GEMM launches 14.7 (208), drain/plan/launch before the page-meta kernel 12.8
(21), batch-build tail 5.8. That 4.2% of decode wall is the MTP verify/draft
path's host-bound sections and is the next pot; it is not this subsystem and
is not touched here.

WHAT CHANGED
------------
1. Every prefill savepoint is staged INSIDE the forward. The in-forward emit
   route (kernel emits the boundary state into the per-layer device slot,
   the spill stream copies it while later layers compute) existed for a
   boundary strictly inside a step. It now takes every position:
     * inside the step — unchanged;
     * at the step's END — the kernel emits at emit_chunk == NT, the state
       after the row's last chunk (the same registers its final store
       writes; one more gated store after the loop in both the GDN and the
       Mamba SSD emit kernels);
     * at the step's START — emit at chunk 0 (the state the row entered
       with) and the conv half from the conv op's own prior_conv input.
   The scheduler arms all three; the commit asks the staged route first
   wherever the step ended and takes the whole-row copy only as a fallback
   for a step that ended on the grid and staged nothing. Both routes are
   counted (savepoint_snapshot_route_total{route=in_forward|whole_row}).
   A completing step whose start state the step before already wrote is
   refused by one predicate on both sides (boundary_already_stored) and
   counted as snapshot_skipped_total{reason=start_already_stored}.
2. The per-layer conv staging runs only under an armed plan on an eager
   launch; a captured launch keeps it unconditionally (a replayed graph must
   hold it).
3. The promoter's prefault of a cold pageable region goes through
   ctypes.memset, which releases the GIL, instead of an mmap.write loop.
3b. The per-layer host cost of the in-forward route, which the first cut
   exposed as launch-bound bubbles on every armed step (+10 ms per cold
   turn against OFF): the device slot is one byte region laid out as the
   host arena lays a layer (state, 64-byte alignment, conv window), so a
   layer spills as ONE device→host copy; the conv window is gathered along
   the transposed input straight into the slot (one kernel), and the row's
   start and end need no gather — they are the conv op's own prior_conv[row]
   and final_buf[row]; the side stream forks with wait_stream; the slot-wait
   timing pairs are read once per step rather than per layer.
4. NOT changed, and why: every launch still carries the emit table. The
   plain and the emitting kernel are two Triton specializations that
   autotune independently, and on a 2048-token row their final states
   differ at accumulation-order level; a prompt that took one when armed
   and the other when not would hold a state no single-kernel prefill
   produces. Item 3's 0.7 ms stays; it is 0.06%.
5. NOT changed: the governor policy. The intermediate crossing entries are
   promoted to pageable RAM before the next write reclaims their slot
   (promote_total 32 of 32 snapshots on the profile boot), survive as
   dominated-but-resident entries, and are what a shared system prompt
   resumes from: a conversation sharing only the first 4096 tokens of a
   6300-token prompt matches the 4096 entry and matches nothing if only
   prompt ends are written. Their coverage is unchanged; only their route
   was wrong.

KERNEL PARITY, ON THE CARD (tools/savepoint/check_end_emit.py)
--------------------------------------------------------------
  lens=[256]          row=0 end   chunk= 4  emit==ht  True
  lens=[256]          row=0 start chunk= 0  emit==h0  True
  lens=[156]          row=0 end   chunk= 3  emit==ht  True
  lens=[1,2048,156]   row=1 end   chunk=32  emit==ht  True
  lens=[1,2048,156]   row=2 start chunk= 0  emit==h0  True
  lens=[1,2048,156]   row=2 end   chunk= 3  emit==ht  True
torch.equal, not allclose: the same registers. tests/test_savepoint_fold_
split_gpu.py and tests/test_mamba2_fold_emit_gpu.py (-m gpu): 9 passed.

A harness note worth keeping: the first version of this check fed the kernel
unnormalised random inputs, the delta-rule state diverged to NaN past ~20
chunks, and every comparison read NaN == NaN as False — which looked exactly
like "the emit is inert on rows longer than 1344 tokens". It was the
harness. The tool now shapes its inputs as the block does and asserts
finiteness before it compares.

WHAT THE FIRST CUT LEFT, AND HOW IT WAS FOUND
---------------------------------------------
With every write in-forward (v1: 32/32 in_forward, 0 fences, 5.4 us mean
slot wait) the cold turn read +1.53% +-0.53% (n=8) against OFF — the fence
was gone and ~19 ms remained. Two further 8-rep boots of the SAME v1 code
(the clone's pull had been refused by a locally modified tool file, which
the absent gdn_chunk_offsets_* counters exposed — an instrument that could
not fire is how the wrong tree was caught) read +1.57% +-0.60% and +1.32%
+-0.20%: that is the boot-to-boot spread of an 8-rep cell, and it is why
the number that stands is the four-boot A/B/A/B below and not any of these.
The v1 trace's cold window put +10.7 ms of the residual in SMALL gaps
(50-200 us, 168 against OFF's 93) spread across the whole GDN layer launch
sequence (l2norm -> cumsum, elementwise -> l2norm, ... -> a_prep), the
signature of the launch thread running closer to the card through an
armed layer. The py-spy leaf frames under the chunk orchestrator named the
cost:

                                    OFF            v2 (armed)
  cdiv (triton/__init__.py)          54  5.5%      296  22.8%
  prepare_chunk_offsets (fla index) 156 16.0%      230  17.7%
  __floordiv__ (torch/_tensor.py)    15  1.5%      138  10.6%
  pad (torch/nn/functional.py)       40  4.1%      118   9.1%

prepare_chunk_offsets — a prepare_lens, a cdiv, a pad and a cumsum, four
device launches and their Python — runs on EVERY layer's launch in both
arms: FLA's identity-keyed tensor cache is disabled process-wide for
capture safety (arbi_serve/_fla_persistent_cache.py), and only
prepare_chunk_indices had been given a per-step memo in its place. The
emitting kernel's wrapper adds its own call per layer, which is the 3x.
So the residual was pre-existing per-layer index work that the armed
route made worse, not the staging code. The offsets table now gets the
same per-step memo as the indices (resolved once from the host twin in the
metadata build, keyed on the device tensor's identity, bypassed under
capture), primed at the same seam, counted as gdn_chunk_offsets_primed —
for both arms.

THE NEW CODE, SAME SHAPE, SAME BOX (ON, 8 reps, tree 44fc0f45)
---------------------------------------------------------------
  savepoint_snapshot_route_total{route="in_forward"}   48   (all of them)
  savepoint_spill_fence_total / _not_ready_total        0 / 0
  savepoint_snapshot_skipped_total{start_already_stored} 8 (the completing
    step of every cold turn, whose start the step before wrote)
  savepoint_fold_emit_slot_wait_ms   228 samples, 1.26 ms  = 5.5 us each
  gdn_chunk_offsets_step_memo        hit and miss both fire; primed every
    served prefill step (the refusals are boot forwards ahead of the
    install seam)
  turn 0 TTFT   OFF 1.2228 s   ON 1.2352 s   +1.02% +-0.67%   (n=8)
  turn 1 TTFT   OFF 1.2936 s   ON 0.1718 s   -86.7%
  turn 0 decode                             +0.84% +-3.9%    unresolved at n=8

Against the OFF boot of the OLD tree: the offsets memo speeds both arms, so
the matched-tree A/B/A/B below is the comparison, not this one. The cold
window of this arm's trace holds 29.7 ms of idle against OFF's 20.1 (v1:
34.1): the fence class is 0.38 ms, step boundaries +2.4 ms (commit 0.58 ms
+ arming per step, both in the probe), and +5.2 ms of small gaps spread
through the armed layers' launch sequences — what an armed layer still
costs the launch thread after the trims (a slot wait, a gather or copy, a
fork, one copy, one fence per layer).

WHAT REMAINS, STRUCTURALLY
--------------------------
An armed step still issues ~6 CUDA calls per GDN layer on the launch
thread, and the GDN section is launch-bound even in OFF (its own 7.7 ms of
small gaps per cold turn). Two levers remain, in order of cost: grouped
spills (2N contiguous slots, one copy per N layers, double-buffered so the
reuse wait lands N layers later; N=4 is ~25 MB of booked VRAM and ~2 calls
per layer), and the commit's 0.58 ms (chain digest, put, retention
bookkeeping). Neither was needed for the number below.

A/B/A/B, ON vs ABSENT, n=48 PER ARM — THE FIRST MATCHED NUMBER (tree 44fc0f45)
-----------------------------------------------------------------------------
Tree 44fc0f45 in BOTH arms, box 206 GPU 1, the receipt's recipe and the
receipt's arms (ON = ARBI_SAVEPOINT_ENABLED=1 with ARBI_SAVEPOINT_DISK_DIR
set and ARBI_PREFIX_TIER=1; OFF = ARBI_SAVEPOINT_ENABLED=0). Four boots,
OFF/ON/OFF/ON, 4 warm-up reps discarded and 24 measured per boot. Turn 0 is
the pure-cost cell (fresh nonce, nothing to resume).

  cell                 OFF        ON         ON-OFF    95% CI            MDE
  turn 0 TTFT          1.2239 s   1.2413 s   +1.42%    [+1.20%, +1.65%]  0.32%
  turn 1 TTFT (reuse)  1.2924 s   0.1716 s   -86.72%   [-86.78, -86.66]  0.08%
  turn 0 decode, ITL   27.61 ms   27.74 ms   +0.50%    [+0.41%, +0.60%]  0.13%
    paired by token position (tools/savepoint/itl_ab.py, which lives on
    branch origin/lcmix-conformance, not on main),
    173 positions, 7632 ON / 15716 OFF gaps

Per-boot turn-0 means: OFF 1.2228 / 1.2250, ON 1.2422 / 1.2405 — the two
replicate pairs read +1.59% and +1.27%, and the arms do not drift into each
other.

VERDICT, STATED THE WAY THE BAR IS STATED
  TTFT: +1.42% [+1.20%, +1.65%]. The whole interval lies above 1%; a
  one-sided test of "effect < 1%" gives p = 1.000. This FAILS the bar. It
  is 2.5x smaller than the +3.59% [+2.88%, +4.31%] the integration receipt
  measured; the host-side commits that followed are measured by the second
  A/B/A/B below, which is the PR's number.
  DECODE: +0.50% [+0.41%, +0.60%], resolved at an MDE of 0.13% — the first
  resolved decode reading this subsystem has had (the per-request rate
  could resolve 2.8-4.2% and reported -0.98% as if it meant something).
  Under the 1% bar, and NOT parity: the machinery costs decode half a
  percent, and its spread is wider (ON rsd 0.74% vs OFF 0.21% on TTFT; the
  ON arm's slowest four cold turns sit 30-47 ms above its median).

WHAT THE ON ARM IS PAYING FOR THAT THE DEFAULT DEPLOYMENT DOES NOT
  The receipt's ON arm sets ARBI_SAVEPOINT_DISK_DIR; the deployed compose
  does not. With it set, the flush band's worker thread ran the whole
  arm at its idle cadence with nothing it could flush: 4987 attempts
  deferred for budget and 2546 races lost in six minutes (Python, under
  the GIL, beside the launch thread), while the store evicted 87 entries by
  LRU under a 1 GiB byte cap. The worker now backs off (cb342f537), and the
  cells below measure the tree with that fix, once with the disk tier as
  the receipt had it and once at the deployed default.

THE FINAL TREE (cb342f53), 8-rep ON cells against the A/B's pooled OFF (n=48)
  disk tier ON (the receipt's arm):
    turn 0 TTFT   1.2325 s   +0.70%  [+0.50%, +0.91%]  MDE 0.29%
      p(effect < 1%) = 0.002; the ON arm's spread fell to rsd 0.28%
      (1226..1236 ms; the four 30-47 ms tails of the A/B are gone)
    turn 0 ITL    +0.34%  [+0.19%, +0.50%]  MDE 0.22%, 178 positions
    flushes_deferred_no_budget 75 over the cell (the A/B arm: 4987),
    32/32 writes in_forward, slot wait 4 us mean
  disk tier OFF (the deployed default), 24 reps, same tree:
    turn 0 TTFT   1.2364 s   +1.02%  [+0.76%, +1.29%]  MDE 0.38%
      p(effect < 1%) = 0.567 — straddles the bar; median 1235 ms (+0.9%),
      five of 24 turns 8-30 ms above it
    turn 0 ITL    +0.39%  [+0.28%, +0.49%]  MDE 0.15%, 178 positions
    96/96 writes in_forward, 24/24 start_already_stored, no fences
  The two cells overlap: at this level boot-to-boot spread is as large as
  the disk tier's difference, so the deployed default is not measurably
  cheaper than the receipt's arm once the flush worker backs off. What
  the tails share is the promoter polling the store lock at 200 Hz through
  every decode step; it is woken by writes from b745de04a on, which is the
  tree the matched A/B/A/B below measures.

A/B/A/B ON THE FINAL TREE (b745de04a), n=48 PER ARM — THE PR'S NUMBER
--------------------------------------------------------------------
Same recipe, same arms (disk tier on, as the receipt had it), four boots.

  cell                 OFF        ON         ON-OFF    95% CI            MDE
  turn 0 TTFT          1.2264 s   1.2422 s   +1.28%    [+1.05%, +1.52%]  0.33%
  turn 1 TTFT (reuse)  1.2931 s   0.1719 s   -86.70%   [-86.74, -86.66]  0.06%
  turn 0 decode, ITL   27.64 ms   27.67 ms   +0.10%    [-0.02%, +0.21%]  0.17%
    179 positions paired, 7689 ON / 15786 OFF gaps

Per-boot turn-0 means: OFF 1.2262 / 1.2267, ON 1.2424 / 1.2420 — replicate
pairs +1.32% and +1.25%; no drift.

VERDICT
  TTFT: +1.28% [+1.05%, +1.52%]. The lower bound is above 1% by 0.05
  points; a one-sided test of "effect < 1%" gives p = 0.992. This FAILS
  the bar as stated, by about a third of a point — from +3.59% [+2.88%,
  +4.31%] at the integration tip. The 8-rep cells that read +0.70% and
  +1.02% along the way were single boot pairs and are superseded by this
  four-boot number.
  DECODE: +0.10% [-0.02%, +0.21%] at an MDE of 0.17% — no effect detected,
  and the run could have seen 1%. The +0.50% the first A/B resolved was
  the flush worker's 20 Hz retry and the promoter's 200 Hz poll under the
  GIL through every decode step; both are gone.
  The ON arm's spread (rsd 0.70% vs OFF 0.41%) remains the open item; see
  "what is left to attack" for the two unmeasured candidates.

THE TWO CONFIGURATIONS, SIDE BY SIDE — DO NOT CONFLATE THEM
-----------------------------------------------------------
The receipt's ON arm turns the disk tier ON (ARBI_SAVEPOINT_DISK_DIR set,
ARBI_PREFIX_TIER=1). The deployed compose sets neither: what ships is
ARBI_SAVEPOINT_ENABLED=1 with prefix_tier at its default (False) and no
disk tier. The bar applies to what ships. Both are honest numbers for the
deployment they describe, and their difference prices the disk tier's
hot-path cost, which nobody had isolated. Same tree (b745de04a), same box,
same recipe, four boots and n=48 per arm each.

  configuration                      turn-0 TTFT   95% CI              decode ITL
  disk tier ON, prefix tier ON       +1.28%        [+1.05%, +1.52%]    +0.10% [-0.02, +0.21]
    (the receipt's arm)
  DEPLOYED DEFAULT (tier off,        +1.16%        [+0.99%, +1.33%]    +0.04% [-0.08, +0.15]
    prefix tier default, nothing       MDE 0.24%   p(effect<1%)=0.967   MDE 0.16%
    else set)

  Deployed arm detail: OFF 1.2248 s (rsd 0.24%), ON 1.2390 s (rsd 0.55%);
  per-boot OFF 1.2246 / 1.2250, ON 1.2388 / 1.2392 — replicate pairs +1.16%
  and +1.16%; turn 1 -86.75% [-86.80, -86.71]; ITL over 187 positions, 7595
  ON / 15949 OFF gaps; 96/96 writes in_forward per ON arm, 24/24
  start_already_stored, no fences.

VERDICT FOR WHAT SHIPS
  +1.16% [+0.99%, +1.33%]. The interval contains 1% — 3% of it lies
  below the bar, and a one-sided test of "effect < 1%" gives p = 0.967.
  This is NEITHER a pass nor a clean fail: it straddles, and it is not to
  be read as a pass. The point estimate is a sixth of a point over the
  bar, down from +3.59% at the integration tip (the receipt's arm, which
  turns the disk tier on, reads +1.28% [+1.05, +1.52] on the same tree;
  the tier's own hot-path cost is therefore about a tenth of a point,
  inside both intervals). Decode: no effect detected at an MDE of 0.16%.
  What remains structurally is in "what is left to attack" below; the
  ON arm's spread (rsd 0.55% vs 0.24%) is the first of those questions.

THE TWO TAIL CANDIDATES, MEASURED (deployed ON, 24 reps, tree bb7c545e3)
------------------------------------------------------------------------
Python's cyclic collector, read by gc_probe over the whole cell (six
minutes of serving, 24 cold turns and 24 resumes):
  gen0: n=258, 47.6 ms total, max 0.77 ms
  gen1: n=23,  15.4 ms total, max 1.44 ms
  gen2: n=0
  pauses >= 2 ms kept: 0
No gen-2 collection ran while serving, and no pause reached two
milliseconds. The collector is NOT the tail, and gc.freeze() has nothing to
move (ARBI_GC_FREEZE_AFTER_BOOT exists, measured, and stays OFF).
The region unmap: one drop in the cell, 5.5 ms on the promoter thread with
the pages released through madvise with the GIL down (region_unmap_ms).
Not the tail either — once in six minutes.
And the cell itself: turn 0 TTFT +1.01% [+0.81%, +1.22%] against the
deployed OFF arms (n=48), median 1236 ms, ONE turn above 1245 ms of 24 (rsd
0.48%, from 0.55-0.74% on the earlier boots); ITL +0.02% [-0.07, +0.10]. The
tails the flush worker and the promoter used to leave are gone, and what
remains is a STEADY ~12 ms per cold turn.

THE FREEZE, AND THE PROBED OFF ARM (24 reps each, same tree)
  ARBI_GC_FREEZE_AFTER_BOOT=1 froze 1,092,560 objects at the ready seam;
  gen2 still ran zero times; turn 0 TTFT +1.07% [+0.91%, +1.24%] — no
  change, as the zero gen-2 count said it would be. The flag stays OFF.
  The OFF arm with the probe: gen0 n=221, gen1 n=20, gen2 n=0; the
  collector behaves the same in both arms and is not what separates them.
  Its TTFT against the deployed OFF arms: +0.05% [-0.06%, +0.15%] — the
  null control's null control, and it reads zero.

THE SECOND DEPLOYED A/B/A/B — THE TREE THAT LANDS (bb7c545e3)
-------------------------------------------------------------
Same deployed arms, four boots, n=48 per arm, the probe armed and the
freeze off (the default):

  cell                 OFF        ON         ON-OFF    95% CI            MDE
  turn 0 TTFT          1.2253 s   1.2402 s   +1.22%    [+1.07%, +1.37%]  0.22%
  turn 1 TTFT (reuse)  1.2926 s   0.1716 s   -86.72%   [-86.76, -86.68]  0.06%
  turn 0 decode, ITL   27.59 ms   27.64 ms   +0.19%    [+0.09%, +0.29%]  0.14%
    179 positions, 7558 ON / 15678 OFF gaps

Per-boot turn-0 means: OFF 1.2249 / 1.2257, ON 1.2399 / 1.2405 (rsd 0.16%
vs 0.51%). One-sided p(effect < 1%) = 0.998.

WHERE THIS STOPS, STATED
  Two independent four-boot A/B/A/Bs of the deployed configuration read
  +1.16% [+0.99%, +1.33%] and +1.22% [+1.07%, +1.37%]. Taken together the
  subsystem costs the shipped configuration about 1.2% of a cold 6300-token
  turn's TTFT — 14-15 ms — and the bar is 1%. This FAILS the bar by a fifth
  of a point, from +3.59% [+2.88%, +4.31%] at the integration tip. Decode is
  within 0.2% (resolved). The 14-15 ms is STEADY, not tail: the tails were
  the background threads and are gone; the collector and the unmap were
  measured and are not it. What is left is the launch thread's ~6 CUDA calls
  per GDN layer on the three armed steps of a cold turn (~9-10 ms) and the
  commit's ~0.6 ms per write (~2 ms). The zero-VRAM levers still on the
  table are each under a tenth of a point; the one lever that would take a
  real bite — grouped spills — spends 5.5 KV pages and is priced below for
  the owner's decision. Stated here rather than chased: another decimal
  place of the same number is not a result.

THE EIGHT CALLS PER GDN LAYER ON AN ARMED STEP, AND WHAT EACH IS WORTH
----------------------------------------------------------------------
Read off the code, per layer, in order:
  1  cudaStreamIsCapturing        conv op         route: capture vs plan-gated eager
  2  cudaStreamWaitEvent(done)    _wait_slot_free the slot's previous D2H must have read it
  3  memcpy D2D / gather kernel   stage_conv*     the conv half into the slot (payload)
  4  cudaStreamIsCapturing        delta-rule op   the same routing decision, other op
     (the emit kernel: real work, stays, and carries the table in both arms)
  5  cudaEventRecord(written)     spill_layer     the emit kernel is done
  6  cudaStreamWaitEvent(written) spill_layer     the side stream waits for it
  7  cudaMemcpyAsync D2H (1)      spill_layer     the layer's region to the ring slot
  8  cudaEventRecord(done)        spill_layer     what 2 waits on, next layer
Each is 1-2 us of CUDA and 5-10 us of pybind and Python; 8 x 48 layers x 3
armed steps is ~1150 calls and 10-14 ms of launch-thread time per cold turn,
of which the traces showed 5-10 ms surfacing as bubbles. That is the steady
remainder.

  FREE: 1 and 4. The staging carries a capture flag (begin_capture /
  join_capture); once every capture admin brackets its recording forward
  with it, the two queries go. ~0.06%.
  STAYS: 3. It is the data.
  IRREDUCIBLE AT ONE SLOT: 2, 5, 6, 7, 8. They exist because layer L+1's
  conv and emit overwrite the slot: the side stream must wait for THIS
  layer's kernel, the copy must be issued, the next writer must wait for
  the copy. "Group in time, not in space" is not available: grouping copies
  requires the K layers' bytes to coexist, and with one slot layer L+1
  destroys layer L. _foreach_copy_ does not apply either: K adjacent slots
  and K adjacent host regions are both contiguous, so a group's copy is
  already ONE memcpy.

GROUPED SPILLS, PRICED SINGLE-BUFFERED
  Over K layers the ordering costs 5/K calls per layer plus the conv copy.
  Double-buffering (2K slots) is NOT needed at this chunk width: a group's
  copy is issued after its last layer's kernel and its slots are next
  written K layers later — after that layer's out-proj, MLP and the
  interleaved attention layers, several ms — while a K-slot copy takes
  K x 0.25 ms. fold_emit_slot_wait_ms is the instrument that says, per
  boot, whether that held (4-5 us mean at K=1 today; a stall reads there).

    K   slots   VRAM        KV pages (of 933)   ordering calls/layer   est. launch ms/turn
    1   1       3.08 MB     0.7 (booked today)  5  (+2 queries, +1 conv)  ~10-14
    2   2       +3.08 MB    +0.7                2.5 (+1 conv)             ~6
    4   4       +9.2 MB     +2.1                1.25 (+1 conv)            ~4
    8   8       +21.6 MB    +4.8                0.6 (+1 conv)             ~3

  K=4 is the smallest step that plausibly moves the deployed number under
  the bar; K=2 is borderline. Either spends VRAM the owner has said is
  zero: the owner's call, with both numbers in front of him.

  REJECTED, for the record: for a boundary at a step's END (every governor
  crossing) both halves of the state are in the slab rows after the step
  and could be copied from there with no slot at all — but the next step
  overwrites those rows layer by layer, so the copies need either the
  whole-row fence (the original defect) or one wait per layer in the NEXT
  step: the same call count moved one step later.

THE TWO FREE CALLS, TAKEN (28c51b275)
  #1 and #4 — the two cudaStreamIsCapturing queries per GDN layer — now
  read the staging's own flag. It is set by the site that records a
  forward: the prefill capture bracketed its record already, the mixed
  capture now does (FoldEmitStaging.capture(stream) sets the flag, joins
  the staging stream's captured copies, verifies coverage, and CLEARS the
  flag when the forward raises, so an aborted capture cannot route every
  later eager prefill down the capture path). The piecewise layer sweep
  carries no bracket, deliberately: capture_layer refuses an unsplit
  prefill and capture_layer_split_gdn keeps the GDN scan eager between its
  pre/post graphs, so no per-layer graph ever holds a GDN prefill op and a
  bracket there would declare a recording that is not happening.
  tests/test_savepoint_fold_emit.py pins the bracket's reset and that no
  op module names the driver query.

GROUPED SPILLS, BUILT AS K, DEFAULT 1 (8b8ae0a07, 52673af8e)
  ARBI_SAVEPOINT_FOLD_EMIT_GROUP=K. The staging holds K slots in ONE device
  buffer at the host arena's stride (align(rec) + align(conv), the same
  stride the snapshot walk packs a kind's layers at), single-buffered. The
  kind's emitting layers take slots in ORDINAL order (layer index is not
  usable: Qwen3.5's GDN layers are 3 of every 4), and the layer that fills
  slot K-1 — or the kind's last layer, for a partial final group — spills
  the group as ONE copy with one fork event, one join event and one fence;
  the snapshot marks the group's layers staged together, when they are.
  K=1 is one copy per layer, the behaviour before the parameter existed;
  tests/test_savepoint_fold_emit_group.py pins that first, then the
  mapping, the one-copy path and its per-layer fallback (a walk that puts
  another attribute between two layers — a PLE context — copies per
  layer, same bytes), a cycle nothing closed (dropped BY NAME at the next
  arm, the seam refuses the snapshot as fold_split_incomplete), and the
  instrument. Mamba2's scan emits BEFORE its conv window is staged, so it
  now claims the group's fence ahead of its kernel — at K=1 that wait used
  to land after the emit had already written the slot.

  THE INSTRUMENT, AND WHAT THE FIRST VERSION GOT WRONG. The coordinator
  asked that a boot at K>1 state whether single-buffering held and that a
  violated assumption be loud. The first cut queried the fence at enqueue
  (was the copy retired when the next writer was launched?) and the very
  first K=4 boot answered "in flight" for a wait the compute stream sat
  0.002 ms on: the launch thread is ahead of the card whenever the card is
  fed, so an enqueue-time query measures the thread's lead, not the
  buffer. The shipped instrument reads the DEVICE's own order: the flush
  that knows its reuse is the next sample records a timing event as the
  fence, the sampled wait is bracketed by a timing pair on the compute
  stream, and the seam reads copy-finished against stream-arrived from the
  three timestamps. fold_emit_slot_reuse_total{copy=retired|in_flight} is
  the order, fold_emit_slot_wait_ms{copy} the time sat, and the first
  in_flight sample of a boot is a WARNING naming K, the copy's lateness
  and the time sat. The sample also moved from the step's first wait (the
  previous step's copy, long retired — every K=1 reading so far was of
  that wait) to its second, the first reuse INSIDE the step, which at K=4
  is the only reuse the single buffer is actually exposed on.

  NULL CONTROL, K=4, deployed configuration, 4 cold turns (smoke_k4):
    savepoint fold emit: 4 device staging slot(s), 12.31 MiB attached,
    K=4 layers per copy, single-buffered
    snapshot_route_total{in_forward} 16 / 16; no fold_split_incomplete
    fold_emit_slot_reuse_total{copy="retired"} 37, {in_flight} 0
    fold_emit_slot_wait_ms{retired}: 37 samples, 0.048 ms total (1.3 us
    each — the fence's own overhead; the copy had finished every time)
  Single-buffering HELD at K=4 on this recipe: the group's 12 MB copy
  finishes inside the four layers (plus interleaved attention) before its
  slots are next written. Whether it holds at other widths is what the
  counter is for; it can fire, and its first version did.

K=4 ON THE CARD: IT DOES NOT CLEAR THE BAR, AND THE PRICE TABLE ABOVE WAS
WRONG ABOUT WHAT SLOTS BUY
--------------------------------------------------------------------------
Deployed configuration (ARBI_SAVEPOINT_ENABLED=1 vs =0, nothing else set
but ARBI_SAVEPOINT_FOLD_EMIT_GROUP=4 on the ON arm), tree 52673af8e, four
boots OFF/ON/OFF/ON, 4 warm-up + 24 measured reps per boot, n=48 per arm,
206 GPU 1 (ab_k4d):

  cell                 OFF        ON (K=4)   ON-OFF    95% CI            MDE
  turn 0 TTFT          1.2249 s   1.2396 s   +1.21%    [+1.05%, +1.37%]  0.23%
  turn 1 TTFT (reuse)  1.2920 s   0.1721 s   -86.7%
  turn 0 decode, ITL   27.59 ms   27.64 ms   +0.18%    [+0.09%, +0.27%]  0.13%
    180 positions, 7603 ON / 15892 OFF gaps

Per-boot turn-0 means: OFF 1.2245 / 1.2252, ON 1.2395 / 1.2397. Share of
the interval below 1%: 0%; one-sided p(effect < 1%) = 0.994. FAILS the bar,
and is INDISTINGUISHABLE from K=1's two readings (+1.16% [+0.99, +1.33] and
+1.22% [+1.07, +1.37]).

The change took effect: the instrument read 133 sampled reuses per boot,
266/266 retired (the copy had finished before the compute stream reached
the slots), 1.3 us each on the fence; 112/112 snapshots per boot in-forward;
no in_flight sample, no warning. The slots genuinely grouped, the launch
thread genuinely issued a quarter of the copies, events and waits — and
TTFT did not move.

SO THE PER-LAYER ORDERING CALLS WERE NOT THE REMAINING COST. Two people
priced the wrong thing: the eight-calls table and the "est. launch
ms/turn" column above convert call counts into milliseconds as if the
launch thread were the critical path on an armed step, and this
measurement says it is not. A menu of slots-for-milliseconds is only
useful if the items are real; that column is not, and it is left standing
above only so this section can say so beside it. Every call-count
reduction on this branch moved the number less than its arithmetic
predicted; this one moved it by nothing.

THE TWO NUMBERS THE OWNER ASKED TO SEE SIDE BY SIDE
  K=1 (shipped default):  933 servable pages, TTFT +1.16% / +1.22%
  K=4 (candidate):        931 servable pages (+9.2 MB), TTFT +1.21%
The collision he expected — VRAM against the bar — does not exist: K=4
spends two pages and buys nothing. The default stays 1.

THE CURRENT HYPOTHESIS: THE BYTES, NOT THE CALLS
  The spill stream carries 288 copies / 464.8 MB / 35.5 ms per cold turn
  (three savepoints of 154.9 MB each, on a side stream, against a ~1.23 s
  turn). Perfectly overlapped that DMA is free; the residual is ~14 ms,
  about 40% of the copy time, and a copy engine sharing PCIe with each
  step's own pinned HtoD at the step start is where imperfect overlap would
  show. That reframes the lever from "how many launches" to "how many bytes,
  and do they overlap": fewer or smaller savepoints, or better overlap.
  Two cheap tests before anyone builds anything, neither run yet:
    * the same calls, the same ordering, the payload copied to a
      device-side sink instead of the pinned slot — if the residual
      persists it is not the D2H bytes; if it collapses, it is;
    * the copy engine's own occupancy read off the Kineto trace against
      the compute stream's gaps, rather than overlap inferred from wall
      clock.

WHAT IS LEFT TO ATTACK, IN THIS SUBSYSTEM AND OUTSIDE IT
--------------------------------------------------------
In the ON arm (this subsystem):
  * NOT the launch thread's per-layer calls: K=4 cut them by three quarters
    and moved TTFT by nothing (the K=4 section). The steady ~14 ms is
    unexplained by call count; the current hypothesis is the ~465 MB of
    D2H per cold turn and how well it overlaps compute, with the two cheap
    tests named there;
  * the commit's ~0.58 ms per write (chain digest, put, retention
    bookkeeping) and the arming's ~0.2 ms per step;
  * the tails are gone: the background threads' spins were them (the
    flush worker's 20 Hz retry, the promoter's 200 Hz poll); the collector
    and the unmap were measured above and are not it. What is left is
    steady.
Outside it (both arms, not moved by ARBI_SAVEPOINT_ENABLED):
  * prefill, ~20 ms of idle per cold 6300-token turn: the ~0.85 ms host
    gap at every chunk boundary, ~0.35 ms of compiled-graph entry per step,
    a 2 ms admission-to-first-launch tail on a request's first step, the
    ~3.3 ms of idle at the prefill->decode transition inside a 12.5 ms host
    phase (seed tail: sample, drafter seed, detok/SSE);
  * decode, ~190 ms of idle per 512-token reply (4.2% of decode wall): the
    MTP verify/draft path's eager sections, the per-step input HtoD
    boundary, the drafter's layer boundaries — the largest pot on the card
    and a different task with its own null control.

CPU LANE
--------
  tests/ -n 4 in registry.arbi.work/arbi-serve:test-latest, CUDA hidden, on
  the tree that lands bb7c545e3: 19444 passed, 386 skipped, 0 failed; on
  the K tree 52673af8e: 19460 passed, 386 skipped, 0 failed. ruff 0.15.21
  check + format clean. Kernel parity (check_end_emit.py) and the -m gpu
  fold-emit suites (9) re-run on the card at every tree.
