Mid-generation self-healing: recovery decision logic by dddabtc · Pull Request #2 · leyten/shard

dddabtc · 2026-06-20T21:48:51Z

What

First increment toward mid-generation fault tolerance. Today, if one stage node dies during a request, the whole request fails (coordinate* raises TransportError, node_kv head os._exits). Detection is already fast; this adds the recovery decision logic so a single node death becomes recoverable instead of fatal — important because the swarm targets churning consumer GPUs.

Recovery is stop, swap, replay, continue: localize the dead stage → swap in a warm spare holding the same layer block → reset live KV and replay the committed prefix → resume. Output stays token-identical because greedy decoding resumes from the exact committed tokens.

Scope of this PR

Only the pure decision module shard/heal.py — no torch, no sockets, no model — so it runs and is unit-tested anywhere:

locate_failure(reports) — the coordinator only sits on the stage-0 and tail edges, so a middle death surfaces as a result that never returns. Each stage reports reachability + highest verify-chunk received/forwarded; the failed node is the first unreachable stage or the first gap (predecessor forwarded a chunk the successor never received). This also catches a frozen, not-dead node.
plan_replay(prompt_ids, committed_out) — the replay prefix is the prompt plus all committed output except the current driver token, since re-feeding it must greedily reproduce that token.
run_recovery(ops, ...) — the ordered state machine over an injected SwarmOps interface (the engine implements the real node ops later), and it aborts loudly if the replay does not reproduce the known driver token rather than silently diverging. That check is the greedy-identity guarantee.

tests/test_heal.py — 13 stdlib unittest cases (localization incl. gap/frozen, replay planning edge cases, and recovery call-order / no-spare / replay-mismatch / head-death paths). All pass.

docs/research/self-healing.md records the design, the scope, and the deferred items.

Deliberately deferred (named, not hidden)

Spare coverage — this increment assumes one warm spare for the protected block; a spare-per-block (any single death survivable) or cold-load-from-pool is scheduler work.
Long-context replay, multiple simultaneous deaths, coordinator death, Byzantine nodes — out of scope here.

Follow-ups (need a multi-node GPU rig to validate end to end)

Engine wiring in phase0/specpipe.py: a per-stage control channel beside the data socket; re-dialing the downstream socket after startup (stages today only re-accept upstream); and turning the coordinator's "timeout → fail" into "timeout → run_recovery".
sidecar/main.go: a small local admin API to swap a forward route to the spare at runtime (no weight reload, no restart).

Acceptance for the engine PR: fixed prompt/K/depth, record output_ids; rerun killing the protected middle stage after a committed chunk; assert the final output_ids match exactly (ids, not text) with no weight reload and no no-failure slowdown.

I don't have a GPU swarm to run the end-to-end path, so this PR is scoped to the part that's fully testable in isolation. Happy to adjust the API or scope if it doesn't fit how you'd like the engine side to call in.

…in-region coordinator) Pipelined async-draft coordinator (specpipe.coordinate_pipe) takes gpt-oss-120B from a latency-bound ~18 to ~40 tok/s (peak ~42), greedy/exact, over WAN on 3 scattered RTX 4090s + a coordinator. Two structural wins beyond pipelining: a 3-stage 12-layer ring (4 WAN hops not 5, fits a 24GB card) and placing the layer-less coordinator in-region (cut the ring 174->102ms). Also lands a fixed-topology tree-verify path (fastverify.tree_decode, validated bit-exact) kept as groundwork. launch_oss.py = the swarm launcher. Verifiable receipt: docs/receipts/gpt-oss-120b-wan-20260619.json (4 distinct GPU UUIDs / IPs / US states, WAN edge RTTs, output token hash, sync==pipelined token match). This is the build target for the permissionless Phases 3-6.

…JOIN/step 1) Replace the shared-SHARD_PSK TCP wire with a libp2p sidecar: a Go go-libp2p daemon that runs next to the Python engine as a transparent TCP-over-libp2p tunnel. Each node holds its own ed25519 key (PeerId) — no shared secret anywhere. The engine swaps transport by one import (wire -> shard.transport, the PSK-free message codec); its logic is untouched. Proven: gpt-oss-120B split across 4 scattered boxes (UT/CA/NV/WA) served over libp2p produces bit-identical greedy tokens to the trusted-wire receipt — same 4 physical GPUs, sha f646e0db...3f70. Receipt: docs/receipts/gpt-oss-120b-libp2p-20260619.json. - sidecar/: go-libp2p tunnel + per-node identity (tunnel + self-test modes) - shard/transport.py: PSK-free send_msg/recv_msg, drop-in for wire - phase0/specpipe.py: --dump on the sync coordinator path - docs/INTEGRATION.md: full done-right architecture (5-verb build); STATE.md: status - docs/NETWORK.md: the network narrative (c0mpute = network, Shard = engine) Next (step 2): NAT traversal (DCUtR + relay; today used mapped public ports), identity<->cwt_ account binding, and re-enabling the pipelined perf path (~40 tok/s) over the sidecar.

…better Re-enable the pipelined perf path over the sidecar. The earlier failure was a latent race in serve_tail_fast: it identified the coordinator's direct-return channel by arrival order (first-readable), but over the sidecar the fire-forwarded reset can reach the predecessor connection before hello_return reaches the return connection, so the tail mis-read a reset as the hello and tore both down. Now identify the return channel by content (hello_return), tolerant of either order — transport-agnostic, no regression on TCP. Result on the 4-box 120B ring over libp2p (plain TCP, no QUIC): PIPE K=4 depth=2 warm 44.79 tok/s, bit-identical greedy output (tokens_match_sync=True, same sha as the wire receipt). vs trusted-wire 39.8 — parity-or-better this window. Receipt updated to the perf path.

…d (step 2.1/2.2) Add the NAT-traversal stack to the sidecar so home GPUs behind a router can join: QUIC, DCUtR hole-punching, circuit-relay-v2 (service + client), AutoNAT, -announce (advertise the real public addr behind a container's port mapping), explicit client.Reserve (deterministic + observable, vs autorelay's black-box finder), and a connection monitor that logs RELAY vs DIRECT. Proven on real boxes (UT relay, CA public, a genuinely NAT-blocked node): the NAT'd node reserves a relay slot and data crosses the relay both ways — i.e. a bedroom GPU can join. The direct hole-punched line (DCUtR) is wired but can't be demonstrated in a 3-node Docker test: DCUtR waits for the node's confirmed public address, which needs several observer peers / AutoNAT, and a Docker box is likely symmetric-NAT. Not a code bug — it activates on the real multi-node network with a real (cone) home NAT. Next: 2.3 identity <-> cwt_ binding (unblocked); validate the direct line on the real network or by standing up AutoNAT/relay helper nodes.

…T lab (step 2.2) Bump go-libp2p v0.33.2 -> v0.48.0 (Go 1.25.11). The old version isn't the issue, but latest is the right place for production (hardened DCUtR/AutoNAT). Code compiled clean against v0.48 (no API breaks). Proved the direct line end-to-end in a controlled lab: two Linux network namespaces, each behind its own NAT (full-cone, UPnP-style), a relay in the middle. AutoNAT confirmed each node's public address, DCUtR upgraded the relay rendezvous to a DIRECT QUIC connection between the two NAT'd nodes, and 100 KB transferred byte-identical over it (circuit-relay caps ~2 KB, so 100 KB proves it went direct, not relayed). Gotchas worth recording: host ufw silently drops the lab's UDP (bypass for the bridge); rp_filter drops NAT'd returns (disable); and libp2p refuses to dial TEST-NET ranges (203.0.113.x) as "blocked observed address" -> use a routable range (11.0.0.x). So: every node joins; full-cone home routers get a direct line via DCUtR, restricted/ symmetric fall back to the proven relay. Same model IPFS runs at scale.

A node proves it controls its libp2p PeerId by signing a challenge with its node key: sidecar -prove <nonce> -> PEERID + base64 SIG sidecar -verify peerid,nonce,b64sig -> OK/FAIL (reference for the c0mpute verifier) The node-agent sends {PeerId, sig} to c0mpute alongside its cwt_ token; c0mpute verifies and records PeerId <-> account (that half lives in the c0mpute repo, per the boundary law). Tested: correct proof verifies, tampered nonce + impersonated PeerId both fail; and cross-language (Go signs, c0mpute's TS verifier confirms). Step 2 (JOIN) complete.

…N piece

Context was silently capped at 2048. The fast verify static KV cache was a fixed 2048 with no bounds check and corrupted past it, and the vLLM draft was started with max_model_len 2048. Both are configurable now (--max-ctx on the stages, --max-len on the draft) and an overflow raises a clean ContextOverflow instead of a CUDA out of bounds, so the node logs it and stays up. Long prompts also need memory efficient prefill, which eager attention can't give us: gpt-oss attention sinks block sdpa and the flash sink kernel is Hopper only, so eager was the only option and it OOMs around 8k. Moved the stages onto flex_attention (Ada native, handles sinks and the sliding window, validated at cosine 0.9996 vs eager) and added chunked prefill so per chunk activations stay bounded while the KV cache grows. A 95,690 token prompt now runs through an N=4 split on four 4090s with no OOM, prefill around 207 tok/s, and the model answers it correctly. Decode at long context is still slow on the simple path; graphed flex decode is the next step. Separately, settled the spec decode question: output is deterministic at a fixed K and matches the trusted wire receipt. The differences across different K are floating point non associativity in the batched verify, not a quality regression. Numbers and repro details are in docs/receipts.

…ndowed-draft spec-decode Long-context decode was unusably slow, about 0.5 tok/s at 95k, because flex recompiled every decode step as the KV grew. Two changes fix it. 1. Split prefill and decode. Prefill stays on flex (memory efficient, the only thing that survives a big query on Ada). Decode flips the layers back to eager and runs the existing CUDA graphed path: eager attention with q=K+1 is cheap even over 100k keys and the graph is a fixed-shape replay, so nothing recompiles. That alone took plain decode to 3.5 tok/s at 95k. 2. Spec-decode with a windowed draft. A full-context draft is ~800ms/round at 95k and kills spec-decode. The draft only needs recent context to predict the next tokens, so the draft query is windowed to the last few thousand tokens (--draft-ctx) while the full swarm still verifies with the complete context. That drops the draft to ~100ms and hides it behind the verify. With draft_ctx=2048, depth 4, spec-decode runs at 7.6 tok/s at 95k with correct output, about 15x the naive path. Mechanics: FastVerify does a chunked flex prefill into its static cache (mask via _compile=True so it doesn't materialise the dense mask and OOM); the fast serve loop routes prefill chunks vs decode by a flag and forwards it down the ring (else a middle stage treats a prefill chunk as a decode and tries to allocate a 49GB eager score matrix); coordinate_pipe chunks the prefill and windows the draft. The verify round-trip (about 250ms over four WAN-separated stages) is the floor now. Next lever for 10+ is sliding window KV: half the layers only need a 128 key window but currently read the whole buffer. Numbers in docs/receipts.

When a stage node dies during a request the swarm currently fails the whole request. This adds the decision logic to recover instead: localize the failed stage, swap in a warm spare holding the same layer block, replay the committed prefix, and resume. Output stays token-identical because greedy decoding resumes from the exact committed tokens, and recovery aborts loudly if the replay does not reproduce the known driver token. shard/heal.py is pure decision logic (no torch, no sockets), so it is unit tested with the standard library (tests/test_heal.py, 13 cases). Engine wiring (phase0/specpipe.py control channel + reconnectable downstream socket) and the libp2p sidecar runtime forward-swap are follow-ups that need a multi-node GPU rig; docs/research/self-healing.md records the design, scope, and those steps.

dddabtc · 2026-06-20T23:24:50Z

Quick heads-up: I should have a real multi-node GPU swarm available in a couple of days. Once it's up I can take the deferred pieces end to end on actual hardware rather than in isolation:

the engine wiring in phase0/specpipe.py (per-stage control channel, reconnectable downstream socket, and turning the coordinator's timeout path into run_recovery),
the sidecar/main.go runtime forward-swap, and
the real acceptance test — kill the protected middle stage mid-generation and assert the final output_ids match an un-interrupted run exactly, with no weight reload and no no-failure slowdown.

That also lets me do the more realistic latency/throughput tuning the recovery path needs (replay cost, localization timing) instead of guessing. So this PR is the testable-in-isolation core; the hardware-dependent follow-ups can land properly after that.

If you have a preference for how the engine should call into this (the SwarmOps surface here is just a first cut), happy to shape it to fit before I wire it up.

Both broke the first residential join (LAUNCH.md P0-#2): - node_kv's flat `import transport` (the pushed-flat box layout) now falls back to shard.transport, so a repo checkout needs no hand-set PYTHONPATH. - m25_pull_range read /root/.hf_token unconditionally and died on any non-vast box: now env HF_TOKEN wins, else ~/.hf_token, else huggingface_hub's own login chain (public/cached repos need none). Pinned by tests/test_shard_stage.py (lands with the shard.stage entrypoint): the clean-env no-PYTHONPATH subprocess gate + the token-resolution unit.

…emon Promotes the operator SSH launch string (m25_scatter_pipe.stage_cmd) into the CLI the Leg-7 node daemon execs (c0mpute NODE_DAEMON.md §4): the ring assignment arrives as flags, the engine env is derived in-process, and the stage speaks a machine-readable stdout contract a supervisor can wait on (SHARD_STAGE_OK preflight / SHARD_STAGE_READY from serve() / SHARD_STAGE_FATAL + nonzero exit). Secrets stay env-only, never argv. --check gives the daemon a no-network preflight (engine import + model-dir sanity). Also: the receipt node-key default is ~-relative (identical on root boxes, works elsewhere). tests/test_shard_stage.py: clean-env subprocess --check (PYTHONPATH stripped — the P0-#2 regression gate), FATAL contract on bad dir/missing assignment, the READY line + a frame flowing through the real serve() loop on CPU, HF-token resolution order. Full suite 554 passed / 1 skipped.

shard #103/#104 + c0mpute #27/#28 landed: shard.stage CLI (ready/fatal contract), both P0-#2 landmines dead, auto-update removed (leyten call), the --mode shard daemon skeleton merged. Next-ranked: forward-leg peer addressing (the blocker), runtime artifact, challenge sketch, warm re-join.

warmring-20260718: coordinate served a real ring job, receipts pinned-verified, nonce threading proven. P1-#2 residue = assignment-EPOCH + pay-model.

…watchdog (#118) P1-#1 done (map live, sim fallback); P1-#2 code-complete (epoch settlement #37). Next: EAGLE P0-#5 mitigation-first.

…, next=EAGLE (#119) Sync LAUNCH.md (was stamped 07-16) + RESUME-HERE to today's true state: - remaining-blockers-at-a-glance header so the list is scannable - P1-#2 CODE-COMPLETE: pay-model built (c0mpute PR #41, per-worker cut after flat-by-layers split, gated off, deploys at launch) — no longer a stub; residue = flip-flag-at-launch + Phase-2 billing - P0-#6: the self-heal MECHANISMS all landed (#28/#34/#36/#37); residue = the churn-survival PROOF demo - economics section: DECIDED (USDC via existing revenue-share, farming unprofitable by the platform cut) + built - RESUME-HERE: the prod-safety operating lesson (c0mpute services exec tsx from the working dir; develop in the c0mpute-dev clone, never prod; nothing deploys before the PoC is finished). Next action unchanged: EAGLE watchdog.

c0mpute PR #41 was merged to master this session, gated off until launch; LAUNCH.md P1-#2 and the M25_ENGINE LATEST-2 block still read OPEN.

Both broke the first residential join (LAUNCH.md P0-#2): - node_kv's flat `import transport` (the pushed-flat box layout) now falls back to shard.transport, so a repo checkout needs no hand-set PYTHONPATH. - m25_pull_range read /root/.hf_token unconditionally and died on any non-vast box: now env HF_TOKEN wins, else ~/.hf_token, else huggingface_hub's own login chain (public/cached repos need none). Pinned by tests/test_shard_stage.py (lands with the shard.stage entrypoint): the clean-env no-PYTHONPATH subprocess gate + the token-resolution unit.

…emon Promotes the operator SSH launch string (m25_scatter_pipe.stage_cmd) into the CLI the Leg-7 node daemon execs (c0mpute NODE_DAEMON.md §4): the ring assignment arrives as flags, the engine env is derived in-process, and the stage speaks a machine-readable stdout contract a supervisor can wait on (SHARD_STAGE_OK preflight / SHARD_STAGE_READY from serve() / SHARD_STAGE_FATAL + nonzero exit). Secrets stay env-only, never argv. --check gives the daemon a no-network preflight (engine import + model-dir sanity). Also: the receipt node-key default is ~-relative (identical on root boxes, works elsewhere). tests/test_shard_stage.py: clean-env subprocess --check (PYTHONPATH stripped — the P0-#2 regression gate), FATAL contract on bad dir/missing assignment, the READY line + a frame flowing through the real serve() loop on CPU, HF-token resolution order. Full suite 554 passed / 1 skipped.

shard #103/#104 + c0mpute #27/#28 landed: shard.stage CLI (ready/fatal contract), both P0-#2 landmines dead, auto-update removed (leyten call), the --mode shard daemon skeleton merged. Next-ranked: forward-leg peer addressing (the blocker), runtime artifact, challenge sketch, warm re-join.

warmring-20260718: coordinate served a real ring job, receipts pinned-verified, nonce threading proven. P1-#2 residue = assignment-EPOCH + pay-model.

…watchdog (#118) P1-#1 done (map live, sim fallback); P1-#2 code-complete (epoch settlement #37). Next: EAGLE P0-#5 mitigation-first.

…, next=EAGLE (#119) Sync LAUNCH.md (was stamped 07-16) + RESUME-HERE to today's true state: - remaining-blockers-at-a-glance header so the list is scannable - P1-#2 CODE-COMPLETE: pay-model built (c0mpute PR #41, per-worker cut after flat-by-layers split, gated off, deploys at launch) — no longer a stub; residue = flip-flag-at-launch + Phase-2 billing - P0-#6: the self-heal MECHANISMS all landed (#28/#34/#36/#37); residue = the churn-survival PROOF demo - economics section: DECIDED (USDC via existing revenue-share, farming unprofitable by the platform cut) + built - RESUME-HERE: the prod-safety operating lesson (c0mpute services exec tsx from the working dir; develop in the c0mpute-dev clone, never prod; nothing deploys before the PoC is finished). Next action unchanged: EAGLE watchdog.

c0mpute PR #41 was merged to master this session, gated off until launch; LAUNCH.md P1-#2 and the M25_ENGINE LATEST-2 block still read OPEN.

Settlement credited every stage on a real ring for the first time: 48 tokens, 6 receipts tiling [0:62), no rejections, and no 'stage' key on the wire, which is #145 doing its job. A tail SIGKILLed mid-serve degraded in 3s, re-formed in 33s, and served and settled again with nobody paid for the interrupted job. A warm node re-joined 12.2s after assignment, verifying its range from disk. The ring-level heal still takes ~30 min because a re-form re-tiles every range and makes untouched survivors re-download ~25GB. That is the honest number and the next engineering win; the launch list is otherwise leyten's to execute.

…ed EU ring (#163) * feat(v4): opt-in partial CUDA graphs over the decode-step islands A V4 decode layer dispatches ~240 kernel launches and profiles at ~8.75ms CPU-dispatch vs ~2.4ms GPU on a 5090 under co-tenant contention -- the wall varies 3-24ms across boxes purely on how starved the launch thread is. K3 answered this by capturing the whole layer block, but V4 cannot: the attention core writes a rotating ring slot (kv_cache[:, start_pos % win]), reads a compressed region whose width grows with position, and the Indexer scores over a growing kv slice; the Compressor branches on (start_pos+1) % ratio; and the MoE routes per token through a host-syncing .tolist(). Each freezes wrong into a graph. So this captures only the position- and data-INDEPENDENT islands the reference exposes as pure Block methods -- the two hc_pre (mix + Sinkhorn), the two hc_post, and the attn/ffn RMSNorms -- and leaves attn and ffn eager between them, fed through static input buffers. V4_CUDA_GRAPH=1 arms it; default OFF keeps the path byte-identical and the CPU suite untouched. A stage refuses (stays eager, says why) without CUDA, and a failed/over-budget capture falls back to a whole-Block eager call with no KV state double-advanced. Measured on the real 158GiB checkpoint, layers [40:43) (I,C,I), sm_120: bit-exact -- graphed h, logits, and every KV/compressor buffer torch.equal the eager reference over 24 decode steps (a graph replays the same kernels, so it is the same bytes, not reassociation-drifted). 9 graphs/stage, 50 fewer launches per layer (20% of ~240; cpu_ms -32%), 1.11-1.19x wall from idle to 2x CPU oversubscription. The ceiling is real and not this lever's to move: the MoE is half the launches AND half the GPU time and is un-graphable as routed, and the attention core is position-dependent -- so graphs cap dispatch contention ~20%, they do not make the layer contention-immune. Batch-invariant/grouped-expert MoE is the next lever. tests/test_v4_stage_graph.py proves bit-exact graphed==eager on GPU (skipped without CUDA) and the CPU refusal; existing v4 suites stay green. * feat(v4): W-deep speculative rollback ring for pipelined speculation Pipelined speculation (PipeInfer/FlowSpec) streams s=1 speculative frames back-to-back so the 5-of-6 idle pipeline stages fill, which needs a stage that can roll back more than one chunk: a rejection arriving W frames downstream must rewind across every compression boundary those frames crossed. Generalize the one-chunk `_spec_ckpt` to `_spec_ckpts`, a position-keyed deque of the last W pre-frame snapshots (maxlen caps the depth). `_seek` now restores the newest checkpoint covering the target, replays the accepted prefix, and spends that checkpoint plus every newer one; `commit(pos)` drops checkpoints the ring has settled past. One-chunk behavior is a special case and every existing rollback test still passes. The correctness gate this was blocked on: the compressed regions are still NOT snapshotted, on the argument that a slot poisoned by a rejected frame is rewritten before its first read. That argument is DEPTH-INVARIANT — the read set [0,(P+1)//ratio) is a function of position P alone, so slot j enters it exactly at the position q=(j+1)*ratio-1 that writes it, and a rewind to r=(committed+1) re-processes every position >= r in order, so every poisoned slot (all at q>=r) is rewritten at its own q before that q reads it, however many ratio-4 (overlap) and ratio-8 (plain) boundaries the W frames crossed. What the snapshot must carry fully for this to hold is the whole window ring plus BOTH accumulators of every compressor incl. the Indexer's (the overlap compressor mutates them at each boundary via the kv_state[:ratio]=kv_state[ratio:] shift). Proven on the CPU oracle: test_multi_deep_rollback_across_boundaries streams W frames, rewinds up to W deep across several boundaries, NaN-poisons the stale compressed region, and matches sequential decode bit-for-bit — incl. a long-prompt case where the Indexer discriminates and a rewind deeper than the window ring. test_multi_deep_rollback_mutation_check proves the green is not vacuous: dropping any one snapshotted region (window / kv_state / score_state / the Indexer's accumulators) makes the rollback diverge. Bounded-W refusal and commit-drop are pinned too, and the split chain matches the single stage. * docs(v4): pipelined-speculation design + throughput projection The async coordinate_dspark redesign that turns the serial one-chunk DSpark round into a streamed one: send the drafted block as B+1 separate s=1 frames back-to-back so the D=6 pipeline stages fill instead of 5/6 idling, and the per-token replay penalty (why DSpark ties greedy) disappears. Records the coordinator state model (sender + receiver threads, committed vs speculative frontier kept <=W in flight, epoch-fenced early inference cancellation), the wire/protocol deltas, and the throughput projection tok/s=(R+1)/((R+D)tau), R=a/(1-a): 10-12.5 tok/s at graphed tau, 15-19 at grouped-MoE+graph tau, for a=0.7-0.8, D=6 — the 10-20 tok/s single-stream target, vs ~1.3 today. Verdict GO: the correctness gate is proven green (multi-deep rollback is bit-exact across compression boundaries, mutation-checked non-vacuous), the MTP drafter needs no rollback by construction (test_cache_never_speculative), and the coordinator rewrite is moderate transport plumbing with the risk retired. The one thing to measure first on a warm ring is whether compute (tau) dominates per-hop WAN latency; the projection is a compute-bound ceiling. * feat(v4): opt-in partial CUDA graphs over the decode-step islands A V4 decode layer dispatches ~240 kernel launches and profiles at ~8.75ms CPU-dispatch vs ~2.4ms GPU on a 5090 under co-tenant contention -- the wall varies 3-24ms across boxes purely on how starved the launch thread is. K3 answered this by capturing the whole layer block, but V4 cannot: the attention core writes a rotating ring slot (kv_cache[:, start_pos % win]), reads a compressed region whose width grows with position, and the Indexer scores over a growing kv slice; the Compressor branches on (start_pos+1) % ratio; and the MoE routes per token through a host-syncing .tolist(). Each freezes wrong into a graph. So this captures only the position- and data-INDEPENDENT islands the reference exposes as pure Block methods -- the two hc_pre (mix + Sinkhorn), the two hc_post, and the attn/ffn RMSNorms -- and leaves attn and ffn eager between them, fed through static input buffers. V4_CUDA_GRAPH=1 arms it; default OFF keeps the path byte-identical and the CPU suite untouched. A stage refuses (stays eager, says why) without CUDA, and a failed/over-budget capture falls back to a whole-Block eager call with no KV state double-advanced. Measured on the real 158GiB checkpoint, layers [40:43) (I,C,I), sm_120: bit-exact -- graphed h, logits, and every KV/compressor buffer torch.equal the eager reference over 24 decode steps (a graph replays the same kernels, so it is the same bytes, not reassociation-drifted). 9 graphs/stage, 50 fewer launches per layer (20% of ~240; cpu_ms -32%), 1.11-1.19x wall from idle to 2x CPU oversubscription. The ceiling is real and not this lever's to move: the MoE is half the launches AND half the GPU time and is un-graphable as routed, and the attention core is position-dependent -- so graphs cap dispatch contention ~20%, they do not make the layer contention-immune. Batch-invariant/grouped-expert MoE is the next lever. tests/test_v4_stage_graph.py proves bit-exact graphed==eager on GPU (skipped without CUDA) and the CPU refusal; existing v4 suites stay green. * feat(v4): chunked verify path, one pass per layer, opt-in A speculative verify chunk of s tokens costs s times a single-token traversal today: Stage.forward replays it position by position, because the reference's decode branch is a hard seqlen == 1 (kv_cache[:, start_pos % win] = kv.squeeze(1), and the same in Compressor). Measured at 3.93x for s=6 on the live 7x5090 ring, which is why g~4 speculation nets ~1.0x end to end. V4_FAST_VERIFY=1 runs the chunk in ONE pass per layer instead. The mechanics the reference has no branch for are written here, driving its own weights and its own Block/hc/MoE/Compressor: the chunk's kv goes to a scratch region appended to kv_cache so the window ring keeps its pre-chunk contents through the attention (chunk position p+i lands on exactly the slot holding the oldest token of p's own window, so a pre-written ring answers p with tokens from p's future), the compressor advances one position at a time through the reference's own decode branch, and the indexer is split into the contiguous runs that share an end_pos // ratio so each scores and picks top-k at exactly the width it would have had alone. Measured on one real V4 layer on a 5090, ctx 1024, s=6: attention + HC goes 5.7x -> 1.4x a single token; the whole layer 5.9x -> 3.9x, i.e. 1.48x, because a 256-expert MoE at 6 tokens touches ~6x the experts through the reference's per-expert python loop. That is the next lever and it is orthogonal to this one. Not bit-identical to the loop, and cannot be: torch picks kernels by tensor size, so a batched pass reassociates (MKL's M=1 sgemm, cuBLAS's gemv, bf16 rsqrt's vectorized path). What is proved instead, exactly, is that nothing else differs -- every position attends the same places in the same order and those places hold the same bytes, and the ring, compressed region and both fp32 accumulators come out torch.equal per layer. On the GPU the whole chunked attention is torch.equal to the loop against the real fp8/fp4 checkpoint; the drift there is entirely in the rented hc_pre/MoE GEMMs. Payload drift <= 24 bf16 ulps and an identical greedy stream over 90/90 swept chunks at s <= 6; 9/18 at s=8, which is why the stream tests stop at the widths a verify chunk actually has and why the flag is off by default. Four mutations of the new code (naive ring write, untruncated compressed rows, one batched indexer group, a skipped compressor step) all fail the suite. v4 suite 193 passed; pipe selftest ALL PASS with V4_FAST_VERIFY=1 forced. * test(v4): batch-invariant harness + full window-wrap coverage for chunked verify CI failed three chunked-verify cases at the window-wrap seam (window_size 16, chunk at 15 spanning 15,16 where 16%16 wraps the ring slot to 0). Root cause is NOT a structural wrap bug: an 81-case sweep over every window-straddling position (15,16,31,32,63,127 x s in 2,4,6, all layer kinds) is bit-exact once float reassociation is removed. The mechanics — which kv each position attends, which slot holds what, what the ring and compressor accumulators end up as — are provably identical to the per-token loop at the wrap. Instrumentation confirmed it: at the failing case every index and every gathered kv byte matches the loop; the only difference is sparse_attn's internal einsum reassociating its d-contraction at s>1 vs s=1 (and the linear GEMMs likewise). This box happened to reassociate the small toy identically at s<=6, so the tests passed here; CI's build did not. The tests were the bug: they asserted torch.equal on raw arithmetic, which is only well-posed once size-dependent reassociation is removed — exactly what the design docstring already claimed ("bit-exact under a batch-invariant GEMM") but never applied. Fix: - add a batch_invariant() context manager that swaps every linear and every einsum (sparse_attn's two, the o-projection, the indexer score) for a broadcast-multiply-then-fp32-sum, whose reduction order is fixed per output element and independent of s. Under it both paths compute identical numbers, so a surviving difference is a real mechanical error. - run attends_exactly / leaves_the_same_state / rows_are_the_greedy_rows under it, and parametrize them over MECH_CASES = the compression-boundary geometries plus WRAP_CASES (15,16,31,32 x s in 2,4,6) — every window multiple, permanently. - rows_are_the_greedy_rows now proves the ALGORITHM's losslessness exactly under the harness (all widths, s=8 included), with the real-arithmetic near-tie cost documented and bounded by drift_is_reassociation_sized (kept on real arithmetic as the structural tripwire). Engine unchanged — there was nothing to fix in v4_stage.py. Five mutations (naive ring write, untruncated compressed rows, one batched indexer group, skipped compressor step, and a wrap-specific contiguous ring commit) all fail the updated suite; control clean. Full v4 suite green; pipe selftest ALL PASS with V4_FAST_VERIFY=1. * feat(v4): default the ring launch to CUDA graphs ON The V4 ring is CPU-launch-bound on a serial pipe, so the partial island graphs (v4_stage._BlockGraphs — the position/data-independent hc_pre, hc_post and norm islands, bit-exact per test_v4_stage_graph.py) are a real ~+12% steady-state single-stream win by collapsing ~68 of ~240 kernel launches per layer. They shipped but were launched OFF, so the first live ring never saw the gain. stage_launch_cmd now sets V4_CUDA_GRAPH=1 by default (new cuda_graph kwarg, opt-out preserved; the env sits before extra_env so an explicit override still wins). The module default stays OFF, so a bare import, the CPU parity suite and in-process rings are unchanged. Capture cost is documented on the launcher: the ~533s first-token cascade is tilelang JIT autotune+compile of the sparse/fp8/fp4 kernels at V4 shapes, not the graph capture itself — tilelang memoises those to its on-disk cache, so the compile half is paid once per box and reused on a re-warm; only a fresh box pays it again. It is a one-time front- loaded tax amortised over the whole generation. * fix(v4): survive a coordinator disconnect without cascading the ring The ring died whenever the coordinator/bench exited: the head reads only from the coordinator and the tail answers only it, so a closed coordinator socket raised in the head's recv (cascading every stage) and left the tail unreachable. Every bench cost a full ~25min re-warm and iteration was blocked on it. Now a coordinator exit/restart is survivable: * the head re-accepts a reconnecting coordinator on its predecessor leg (reaccept closure, head only) — the forward leg to the rest of the ring stays warm and _fwd_open heals it if it idled out; a middle stage still cascades on a genuine upstream death. * the tail's return channel is a swappable _RetChannel; a background thread re-accepts a reconnecting coordinator's hello_return and swaps the socket in FIRST, then acks, so by the time connect_ring returns the tail is already answering on the new socket. A send to a departed coordinator is dropped, not fatal. A reset opens every job, so nothing partial survives the gap. _coord_cli documents the persistent-coordinator pattern this enables: a re-run bench reconnects to the SAME warm ring instead of re-warming. Test: test_ring_survives_a_coordinator_disconnect drops both coordinator sockets with no stop op, re-dials the same head/tail, and decodes the same stream twice more over the still-running stages. selftest + selftest-relay still ALL PASS. * feat(v4): confidence-gated adaptive send-length for the dspark path The DSpark drafter emits a per-position confidence with every block and RingDrafter already ships it as `conf`, but coordinate_dspark ignored it and always offered the whole block. On a round the drafter itself expects to lose early, those extra positions are chunk work the ring computes and throws away; on a round it expects a long run, the whole block banks the most tokens per traversal. coordinate_dspark can now truncate the OFFERED block to the confidence- predicted survival prefix (keep the leading run whose conf stays >= a threshold, floored at conf_min). It is lossless by construction, not by tuning: the tail verifies exactly the chunk it receives and both ends run plan_verify_round over the same sent drafts, so the committed stream is byte-identical whatever the send-length — a pure throughput knob that can cost acceptance but never a token. Default OFF (V4_DSPARK_CONF_GATE / per-job confGate) and the reason is in the code: measured on the CPU reference, `conf` is a RAW logit near 0 that goes negative, so there is no universal cutoff and the raw-score floor must be calibrated to the real model first. A `conf_probe` hook and the new `sent`/`send_hist` stats are the calibration path: run ungated on a live ring, record (conf, accepted) per round, set the floor. Research puts the win at ~+15-25%. Tests: _conf_send_len survival-prefix unit test, and a real drafted ring run three ways (ungated / trim-to-1 / whole-block) all emitting the exact greedy stream while `sent` proves the knob moved. * feat(v4): grouped fp4 MoE decode kernel — 3.2x MoE, CUDA-graph-capturable The V4 MoE decode path is CPU-launch-bound: a bincount().tolist()/nonzero host sync plus ~120 tiny fp4 expert launches per layer per token, for microseconds of actual fp4 arithmetic. v4_moe_grouped collapses the six routed experts (w1/w3/w2) into three grouped fp4 GEMM launches over a gathered [G,N,K] expert bank and drops the last host sync, so the whole MoE.forward is sync-free. The grouped GEMM is kernel.fp4_gemm_kernel with a batch (expert-slot) axis on the grid; the gathered weight/scale are indexed by grid position. Every output element sees the vendored kernel's exact arithmetic (block_K=32 = the fp4 scale group, same FP4->FP8 cast, same per-32 x per-128 scales, same fp32 accumulate), and the ascending-expert-id fold reproduces the reference loop order, so it is torch.equal to the reference MoE at s=1 — verified bit-exact over 8 draws at V4's shipped dims on a real 5090, three-way against both the reference and v4_moe_decode. Measured (RTX 5090, sm_120, single-token decode, real MoE dims): reference MoE.forward 4.15 ms v4_moe_decode 3.39 ms (1.22x) grouped 1.30 ms (3.19x vs ref, 2.61x vs decode) grouped + CUDA graph 0.38 ms (dropping the .tolist() lets the whole MoE forward capture; replay bit-exact) Opt-in, default OFF (V4_MOE_GROUPED=1); with the env unset install() is a no-op and the decode path is byte-identical. Falls back to the captured reference for s>1 (grouped MoE is not token-count invariant), world_size>1, and hash-routed layers — the same envelope as v4_moe_decode. sm_120 note: a device-side W[eids[g]] dereference (a data-dependent per-block index into the packed-fp4 bank) mis-addresses on this tilelang build — uniform eids work, distinct eids collapse every slot onto one weight — so the slot->expert map lives in a torch gather (device-side, no host sync) and the kernel indexes the gathered bank by grid position. The FP4->FP8 MMA JITs clean on sm_120, no tcgen05/TMEM. * perf(v4): collapse the DSpark drafter's wasted intermediate forwards advance_and_draft commits n positions per round (n = accepted + 1, up to block_size + 1) by running the reference forward_spec ONCE PER position and keeping only the last block. Every intermediate call runs a full MoE stack, the vocab-wide head and the Markov loop and discards all of it to leave one byte behind: DSparkAttention writes kv_cache[start_pos % win] = main_kv and nothing else (main_kv derives from main_hidden alone; the draft block's own K/V is dropped). So ~5 of every 6 draft forwards are wasted, and the round is CPU-dispatch-bound with the GPU idle. v4_dspark_fast rebinds advance_and_draft (the v4_moe_decode install pattern; opt-in V4_DSPARK_FAST, default OFF, wired from ring_drafter) to run only that slot write for the n-1 intermediate positions and the full forward once for the kept block. Bit-exact vs the reference loop -- drafts, logits, confidence and every mtp KV buffer torch.equal -- on the CPU suite and on a real-dim GPU drafter. Per position, not batched: batching the intermediate KV GEMM is bit-exact on CPU but an M=k matmul reassociates its reduction differently from k separate M=1 ones on a GPU (the fp32 confidence head diverged at k>=2), and this is the spec-decode verify path where an accepted draft is committed output. Measured (RTX 5090, synthetic bf16 drafter at real dims, ~20.5 ms/forward): current scales n x forward (n=6: 123.7 ms), cache-advance stays ~flat (n=6: 27.4 ms, 4.5x); at n=1 there is no change, as there are no intermediate positions. On the real fp4 drafter (~37 ms/forward) this is the 224 -> ~40 ms the round predicts. A second opt-in lever (V4_DSPARK_GRAPH, CUDA only, requires lever 1) CUDA-graphs the kept forward's forward_head -- the one fixed-shape, position-independent, host-sync-free slice of forward_spec. Bit-exact (graph replay == eager) but only ~1.05x here: the MoE-bearing block bodies dominate the forward and cannot be captured, because the reference expert dispatch drains the device (indices[0].tolist()) and branches on the result in Python. The head is the capturable ceiling, not the whole forward. * feat(v4): reference-compute slim overrides — indexer skip + QAT-sim drop Two removable blocks of per-layer decode work for the DeepSeek-V4-Flash ring, each an install-override that rebinds the vendored reference after model.py is executed (v4_moe_decode's pattern), each behind a default-OFF env flag so the reference path is byte-identical with both off. Item 1 (V4_REF_SLIM): skip the Indexer while context is short. Every ratio-4 layer runs an fp8 GEMM + Hadamard + fp4 quant + its own compressor + score einsum + top-k (~15-22 launches) to pick top-`index_topk` compressed slots — but while end_pos//ratio <= index_topk the top-k selects EVERY slot, so the attended set is exactly the fixed get_compress_topk_idxs index the ratio-128 layers already use (shipped crossover end_pos 2051). Substituting it is set-identical (order-only diff => the same gather-order bf16 ULP the sm120 retile ships; 0 at CPU-oracle scale). Correctness gate: a job that can cross the crossover keeps the indexer's compressor advanced (skip only scoring) so re-engage is bit-exact; the compressor is skipped entirely only when set_job_max_pos guarantees short. Item 2 (V4_REF_SLIM_NOQAT): no-op the inplace act_quant/fp4_act_quant QAT round-trip (quantize->dequantize back to bf16) that every layer runs to match an fp8/fp4 KV deployment. A bf16-KV deployment wants the full bf16 KV; the real non-inplace GEMM quantization is delegated untouched. Approximate (removes a precision reduction); gated separately, keep OFF for an fp8-KV cache. CPU-oracle proof in tests/test_v4_ref_slim.py (12 tests): env-off no-op + idempotent install; select-all set-identity; short-ctx logits within 2 bf16 ULP (0 measured); re-engage bit-exact past the crossover when the compressor is kept; mutation-check that skipping it diverges; item-2 no-op/real-quant split + bounded logit move. Full v4 CPU suite green (137 total). * feat(v4): carry the job horizon to every stage for the ref-slim indexer skip v4_ref_slim's indexer skip can also drop the Indexer's Compressor, but only for a job that provably never leaves the select-all regime — and the Compressor is STATE, so getting that wrong short means the indexer re-engages past index_topk*ratio against a half-filled cache and picks the wrong keys, silently, in plausible numbers. The module shipped set_job_max_pos() and left the call site to the serve path; this is that call site. The horizon rides the reset frame that already opens a job and is propagated down the ring unchanged, so every stage sets it before the first step lands. All three coordinators declare an UPPER BOUND, never an estimate: greedy is exactly prompt+max_new, coordinate_spec adds K+1 for the draft chunk a round puts on the wire above the committed length, and coordinate_dspark — which cannot know the tail's MTP block size at reset — takes a deliberately fat fixed margin. The asymmetry is the whole design: over-declaring keeps the compressor advanced (correct at any length, two cheap GEMMs), under-declaring is wrong, and an absent key lands as None, which v4_ref_slim already reads as the safe answer. So a reset frame from an older coordinator degrades to unoptimised, never to incorrect. serve_stage clears the horizon on teardown. In a real ring that is cosmetic (one process per stage); in the in-process shape the selftest and tests use, the stages are threads beside the coordinator and the oracle and the global is shared, so a dead stage must not leave one job's value behind it. Tests: the declared max_pos upper-bounds every position actually driven, on all three coordinators; an absent key is None; and on a real 2-stage ring the horizon reaches the stages, does not change the answer, and is gone after teardown. * feat(v4): install the grouped fp4 MoE kernel in the serve path, after the decode one The grouped kernel shipped with an install() nothing called, so V4_MOE_GROUPED=1 was a flag with no effect on a ring. load_ref() is where the MoE overrides go — after model.py is executed, in the same window v4_moe_decode already uses — and it is now wired there. The ORDER of the two calls is the precedence, not a style choice. Each install captures whatever MoE.forward is bound at that moment as its own fallback, so installing grouped SECOND builds the chain grouped -> decode -> reference: grouped claims the single-token score-routed decode step and hands back what it declines (s>1, world_size>1, hash-routed layers), decode handles those, and the reference gets what neither claims. Reversed, decode would sit on top and claim that same decode step, and the grouped kernel would be installed and unreachable. Both flags off leaves the reference byte-identical, and grouped additionally refuses without CUDA, so a CPU box importing this pays nothing (tilelang stays deferred inside the kernel builder). Tests pin the precedence hermetically on a stub module — rebinding the real dsv4_model.MoE.forward would follow every later test in the run — plus the both-off no-op and that load_ref actually makes the call in that order. * test(v4): close the ratio and batch gaps in the W-deep rollback proof An adversarial pass over the multi-deep rollback could not break it (36 cases: interleaved push/spend/re-push, boundary-exact targets, deep multi-boundary rewinds with the overlap shift in the rejected tail, differing correction tokens, commit off-by-one, eviction, long-context Indexer, coarse-checkpoint replay, fuzzer) but flagged three untested surfaces. One (multi-stage split) was already covered; the other two are closed here, because both are load-bearing for the claim that a toy oracle licenses a statement about the shipped model: larger ratio the shipped config compresses at 4 and 128 while cpu_args uses 4 and 8, so "does a big ratio differ" was untested. The depth-invariance argument is ratio-agnostic by construction (the read set [0,(P+1)//ratio) has the same shape at every ratio); ratio 16 with a rejected tail crossing p=31 checks it. batch b=2 _snapshot clones whole buffers rather than [:bsz], so a snapshot silently covering only row 0 would pass every other test in this file. Also pin the threads in the run line: at cpu_args' toy shape torch's intra-op threading is pure contention, 1.78 s/decode-step at 4 threads vs 0.079 s at 1, which is a 22x slowdown over a suite that decodes thousands of single positions (41 passed in 4m16s pinned, vs 32m unpinned). No numerics change — every comparison is stage-vs-oracle in one process. * feat(v4): whole-layer decode CUDA graphs — capture-safe attention core The island graphs (V4_CUDA_GRAPH=1) only fold in the hc_pre/hc_post/norm islands and leave attn and MoE eager, so they cap the decode-step dispatch at ~1.2x. This makes the ATTENTION CORE capture-safe so the whole layer captures as one graph — the paged-decode-graph technique M2.5 got from vLLM and K3 punted on for MLA. Three things baked position/data into a graph; each is removed by a second transcription of the reference's decode branch that calls the reference's OWN kernels and parameters (v4_whole_layer_graph.py), so a correct capture replays the same kernels on the same bytes — bit-exact, not reassociation: #1 device-side position: start_pos enters as a length-1 device buffer, and every kv_cache[:, p%win] / kv_state[:, p%ratio] / kv_cache[:, p//ratio] store becomes index_copy_ on a device-derived slot chosen at replay. #2 fixed-width masked read: the Indexer scores q against the WHOLE kv_cache (fixed max width), masks columns past end_pos//ratio to -inf, takes a fixed topk, and maps future-slot picks to -1 — which sparse_attn already treats as a no-op, so the fixed-width topk is byte-identical to the reference's growing one. #3 two graphs by position: should_compress is host-known, so the block is captured twice (compress / no-compress) and the replay picks by position, keeping zero data-dependent control flow inside a graph. V4_CUDA_GRAPH=whole arms it, default OFF keeps every path byte-identical. Because the routed MoE host-syncs (bincount().tolist()), the stage runs it EAGER between two graphs (moe_eager) — real-serving-safe and bit-exact to the reference today; a graph-safe grouped-fp4 MoE folds the FFN in for the full win. moe_stub (a fixed expert set) proves + measures the whole-layer ceiling. Proven bit-exact: the capture-safe block == model.py Block.forward on CPU over 40 decode steps that wrap the window and cross the ratio-4/8 compression boundaries (no GPU needed); on a 5090 the whole-layer graph == eager (stub), the moe_eager graph == the reference block, and a V4_CUDA_GRAPH=whole Stage == the eager Stage — hidden, logits, and every KV/compressor buffer. Measured on a 5090 (sm_120), synthetic weights at real dims (dim 4096, 64x512 heads, window 128, index_topk 512), per decode layer, dispatch-bound box: island vs eager (real MoE): 1.20x moe_eager vs eager (real MoE): 2.11x (deployable, bit-exact) whole vs eager (stub MoE): 8.02x wall / 13.2x cpu (ceiling) Receipt: docs/receipts/v4-whole-layer-graph-20260801.json Rollback safety is unchanged: the graph writes are position-parameterized, and a speculative rewind (_seek) still replays the accepted prefix through the eager per-token path (the _replaying guard), so it rides on the existing _snapshot/_seek proof. tests/test_v4_whole_layer.py: the CPU bit-exactness (runs anywhere) and the GPU stage parity (skipped without CUDA). Existing v4 suites stay green. * fix(v4): make the grouped MoE decline a bank that would not fit, not OOM the stage Composing the four perf branches turned up a hazard none of them could see alone. _expert_bank stacks a layer's routed experts into a contiguous bank and caches it, while the reference's per-expert nn.Parameters stay alive — so at the shipped dims it is ~3.2 GiB of DUPLICATE weights per layer beside the ~3.7 GiB already there. The module doc always said the real fix is a load-time layout choice (store the experts as a bank, drop the per-expert tensors); that lives in the loader and is not on this branch. Wiring install() into the serve path made the flag reachable without it. The failure was not graceful. The bank is built lazily, on the first decode token of a layer, after the graph pools are pinned and the KV cache is allocated, and nothing upstream catches an OOM out of ffn — _BlockGraphs guards only the graph capture, Stage.forward and _forward_loop guard nothing. A stage that tried and failed would die and cascade the ring on the first job. The GPU bench builds ONE layer, where the copy fits, so it could not have found this; the CPU suite never installs at all. So the bank build now measures free VRAM first and declines — once per layer, cached, with a line saying why — and the layer falls back to the decode path. The lever is safe to leave in a ring's env: it either engages or says it did not. Also closes a cycle this branch's install ORDER made possible: grouped captures decode_forward as its fallback, so a second v4_moe_decode.install would capture grouped_forward and the two would call each other until the stack blew on the first prefill. load_ref runs the pair once, so nothing reaches it today — decode now refuses to install over grouped rather than leaving the trap armed. * docs(v4): the composed launch recipe for the perf round 2 ring What to set, what to leave off, and the two things the composition found that no branch could see alone. The headline is that V4_FAST_VERIFY is a COMPETING strategy, not an additive lever. Every other lever gates on a single-token shape — graphs on h.shape[1]==1, grouped and decode MoE on xv.size(0)!=1, and the slim indexer only exists on Indexer.forward, which the chunked path bypasses entirely via its own _chunk_indexer. So the chunk that fast-verify claims is exactly the chunk where the other three switch off. On a drafted ring that is nearly all decode work (dspark_block_size 5 => s=6, which _chunk_ok accepts), so stacking it on top does not add its win to theirs, it replaces them. The A/B ladder here runs it last and alone rather than third, which is what makes a step-3 regression attributable. Second: V4_MOE_GROUPED will decline on a full stage until the loader stores the experts as a bank, and the 11x figure came from a whole-MoE graph capture the serve path cannot do (_BlockGraphs leaves ffn eager by construction). 3.19x eager is the ceiling, where the bank fits at all. Also records what "lossless" does and does not cover: V4_REF_SLIM is ULP-level, not bit-exact, so a token-identity gate can in principle trip on a near-tie. * fix(v4): pin the Indexer's top-k tie-break, and correct an overclaimed bar The previous commit claimed the capture-safe attention core is bit-exact to model.py's Block.forward over a full decode run. That was WRONG. It passed under default CPU threading and fails at OMP_NUM_THREADS=1, where the hidden state diverges by ~6e-3 at pos 43. The cause is not reassociation and not the masking or the compress split — the scores over the valid columns are bit-identical to the reference's narrow einsum (measured max|narrow - wide| = 0.0). It is the TIE-BREAK. index_score is full of exact ties, because relu_ floors every negative score to a hard 0.0, and torch.topk resolves them by an artifact of its partition that depends on the array WIDTH and on the CPU THREAD COUNT. So a fixed-width read breaks ties differently than the reference's growing-width one — and, worse, the reference does not agree with ITSELF across thread counts. Selection is now a STABLE sort, (score DESC, column index ASC): deterministic on every width, thread count and device, agreeing with the reference on every tie-free selection and never picking a lower-scoring column. That is a property the eager reference path does not have, and it is a latent ring hazard on its own — two boxes running the same layer can otherwise select different compressed slots, and a speculative verify can reject its own drafter for no modelling reason. Recorded as the top risk in the receipt. The bars are re-scoped to what is provable, and every one of them is now green at OMP_NUM_THREADS=1: * block == reference, bit-exact, wherever the selection is unambiguous — 15 steps at index_topk 8, and the FULL 40-step run at index_topk 16 (wrapping the window 2.5x, crossing ten ratio-4 and five ratio-8 boundaries); * the selection rewrite itself, tested directly: identical selected SET whenever the boundary is tie-free, identical selected SCORE MULTISET always, plus the -1 padding, the offset mapping and width-independence; * graphed == eager capture-safe, unconditional, unchanged. A dropped claim, too: buffers do NOT stay equal under selection pressure — a diverged pick in an early indexer layer changes a later layer's input. That test failed and was removed rather than weakened. The CPU block-parity tests now skip when the process bound the tilelang kernels, instead of failing on a device mismatch. Re-measured after the fix (the sort costs a little over topk): island 1.18x, moe_eager 2.08x, whole 7.26x wall / 11.9x cpu. * feat(v4): incremental accept rule on the tail for streamed s=1 frames Pipelined speculation streams the drafted block as separate one-token frames instead of one chunk, so the tail can no longer advance its drafter "over the round's committed prefix" — there is no round, and when a frame is forwarded nobody yet knows whether its token will survive. RingDrafter gains a `pipelined` mode (armed per job by the reset) that applies the SAME accept rule one position at a time. Two scalars are the whole of it. Frames reach the tail in the order they were injected, so the frame at q is on the committed path exactly when q == cfront+1 and its token equals mfront, the greedy the tail produced at cfront. A rejected frame does not move the frontier, so every frame streamed behind it fails the same test — the poisoning of a speculative tail falls out of the rule instead of needing a flag, including the sharp case where a later draft happens to equal the greedy token the REJECTED frame produced. The coordinator's correction lands back at the frontier carrying mfront and re-opens it. The drafter advances one position per committed frame, off that frame's own tap, and the block it returns proposes q+2..q+B+1 — which is the serial chunk [cur]+drafts decomposed. That decomposition is bit-identical, not merely close: advance_and_draft already loops per position internally, and if the two cadences produced different mtp state then an acceptance rate measured on one path would say nothing about the other. Pinned on two independent real drafters prefilled from the same tensors. Nothing here rolls back — the mtp cache only ever records committed positions (test_cache_never_speculative), which is why only the layer stage needed the W-deep rewind. 29 dspark tests green (6 new), CPU, OMP_NUM_THREADS=1. * feat(v4): async pipelined DSpark coordinator, epoch-fenced and lossless coordinate_dspark sends [cur]+drafts as ONE frame and waits, so exactly one chunk is ever in flight (5 of 6 stages idle) and every stage replays it position by position — which is why the drafted path merely ties greedy. coordinate_dspark_ pipelined streams the same block as separate s=1 frames back-to-back: stage k works on token i while stage k-1 works on token i+1, and no frame is ever longer than one position, so both costs go at once. Opt-in via V4_PIPELINED_SPEC or a job's `pipelined` flag; the serial path is untouched and byte-identical with the flag off (no epoch => no fence, no cpos => no commit, no extra reply keys). Losslessness is not an argument about speculation. Every emitted token is m_p, the tail's own greedy token from a frame whose entire history is committed, taken at M=b through ParallelHead exactly as greedy decode takes it. Accepting a draft commits m_p, which merely happens to equal it. So on the accepted path the ring sees precisely the frame sequence greedy decode would send, plus rejected frames that Stage._seek undoes before the next accepted frame — same frames, same shapes, same order. THE EPOCH FENCE, and what it honestly is not. Every frame carries a generation id the coordinator bumps on each cancel. The receiver is where it is load-bearing: replies for discarded frames are still in flight when the correction is streamed, and without the epoch they would be read as judgments of the NEW frames at the same positions. The sender drops queued frames of a dead epoch before the wire, with the check and the owed-a-reply bookkeeping under one lock so a drop can never strand the receiver. The stages pass a stale frame on without computing it — a GUARD, not PipeInfer's early cancellation: on a single in-band FIFO ring a cancel cannot overtake the frames it cancels, so shrinking the bubble would need an out-of-band path to every stage that this topology does not have. What it buys is that a replaying relay or a coordinator bug cannot _seek a stage onto a discarded future. Fenced frames are excluded from the receipt chain on EVERY stage (the marker propagates), so out_root == in_root still holds hop for hop. The cancel costs less than the design doc assumed: the reply that reveals a rejection is a COMMITTED frame's, so the tail already drafted a fresh block off the correction token — [correction]+block streams immediately and the pipeline refills as it drains. `cpos = c-1` rides each frame as the commit watermark, bounding the W-ring's clone memory without a control frame. Proven on CPU, bit-identical to the reference Transformer's greedy stream and to the serial dspark path, in every regime that could move the answer: zero accept (every block rejected, every reply rewinds W deep), full accept (pipeline fills to block+1 frames, no cancel), partial accept, EOS and max_new landing mid-block with frames still in the ring, and a cancel landing exactly on a compression boundary (positions 7 and 11 of cpu_args' ratio-4/8 mix — the zero-margin case). Receipts settle with check_chain on all of them. The stale-frame test carries its own anti-vacuity half: with _fenced stubbed out the same injected frame corrupts the stream. 64 pipe tests green (16 new), plus both selftests. * perf(v4): bucket the indexer read, and grade the graph on m25's two-tier bar Two corrections, both from precedent already in this repo. BUCKET THE READ, DON'T MASK A MAX WIDTH. The indexer read was the whole max_seq_len/ratio cache at every position, masked down — which paid full-context cost from token one, and put a top-k downstream of a masked wide array. m25_stage's DECODE_BUCKETS is the right shape and V4 already had the warning in its own tree: _chunk_indexer says "Masking a common-width score to fake the shorter reads would change which slots a topk near-tie picks". So the read is now the smallest rung >= end_pos//ratio, one graph per rung, captured lazily on a crossing, capped by V4_GRAPH_MAX with over-budget layers staying eager, counted and logged. Bounded a priori by the ladder, not by the sequence. The stable-sort selection is what makes bucketing legal: with topk the answer would depend on WHICH BUCKET a position landed in, i.e. on capture history. Selecting by (score DESC, index ASC) is width-independent, so a rung crossing is invisible. That is now stated where it is relied on. GRADE TWO-TIER, m25's bar (research/graph_aux_check.py), the one its CUDA graphs shipped +74% under: TIER 1, hard torch.equal — graphed == the eager twin (_eager): same math, same bucket decisions, no capture. This is the real capture-correctness proof and it is unconditional. TIER 2, named and bounded — vs the vendored reference. Token-identity there is not a bar the reference can meet: its topk tie-break is not reproducible against ITSELF across array widths or OMP_NUM_THREADS. Measured at real dims: 40/40 steps bit-exact, max|graph - reference| = 0.0. FRESHNESS GATE — the same input at a DIFFERENT position must MOVE the output, across a bucket rung. The cheapest insurance against position baked in at capture, which is the worst silent failure here and the one m25 actually shipped once (stale EAGLE aux, e8d2c82). Also noted at the source: V4's DSpark tap is safe only because it sits OUTSIDE the captured region. The moment a graph spans more than one layer it becomes e8d2c82 verbatim — a Python side effect inside a graph is recorded once at capture and skipped by every replay. Comment + gate added. Bucketing paid for the stable sort and then some — the CPU suite went 32s to ~5s, and at real dims (dim 4096, 64x512 heads, window 128, index_topk 512): island 1.22x, moe_eager 2.13x (deployable), whole 7.66x wall / 12.6x cpu. Fixed on the way: the stub capture path had lost its _feed_capture, so a compress variant captured at pos 0 read freqs_cis[1 - ratio] and tripped a device-side assert. * fix(v4): width-invariant Indexer top-k in the capture-safe decode The whole-layer capture read the Indexer's scores at a fixed max width and masked the tail to -inf, on the claim that a masked wide topk is byte-identical to the reference's narrow one. It is not. Tensor.topk does not define its tie order -- it is whatever the selection algorithm leaves behind, it differs between the CPU and CUDA backends, and it is not invariant to array length -- and index_score is bf16 behind a relu_() that floors negatives to a hard 0.0, so ties at the k-th rank are routine. At decode pos 43 the wide read picked compressed slot 1 where the reference picked slot 4 (scores tied at exactly -0.0751953), and the hidden state diverged by ~2e-3. A different attention support set, not a rounding drift. Select by (value DESC, index ASC) instead: a total order, so the top-k is unique and cannot depend on how many -inf columns are padded on the end. All fixed-shape device-side ops, no host sync. That also makes BUCKETING the read (m25_stage.DECODE_BUCKETS' shape) a pure cost lever -- bucketing alone does not fix this, since a bucket is still wider than the position needs. Second cause, found once the first was neutralised: v4_kernels_cpu.sparse_attn reduced over the true topk width in one flat pass, so padding a 31-wide index list to 48 with -1 regrouped its pairwise tree and changed the last bits (~1% of calls; it bit a 200-step run at step 122). kernel.py walks topk in FIXED 64-wide blocks, where an extra all-masked block rescales by exp(0)==1 and adds 0 -- padding is bitwise free on the GPU. Emulate that blocking so the GPU-less oracle is faithful to the kernel it stands in for. Grade in two tiers, because token-identity against an arbitrary tie order was never the right bar. Tier 1, hard: the read width is a cost knob (max, bucketed and exact widths give torch.equal hidden/logits/KV -- 4 seeds x 200 steps x 3 widths), plus a freshness gate that the same input at a different position must move the output. Tier 2, named and bounded: bit-identical to the vendored reference until its own top-k stops being well defined, with the tie census reported (23 steps clean, first tie at pos 43, 23/120 steps tied). Steps raised 40 -> 120: the 40-step bar walked straight past the reduction width bug and passed on seed 23 while failing on 7 and 11. pytest tests/test_v4_whole_layer.py: 3 passed, 1 skipped. Full V4 suite green (126 passed, 2 skipped). * fix(v4): reconcile with e1db515 — kernel-faithful oracle, sorted lane order Takes the two fixes from the independent triage on v4/whole-layer-graph-fix and keeps this branch's bucketing, budget and moe_eager path on top. TAKEN: the block-faithful v4_kernels_cpu.sparse_attn. The CPU oracle reduced in one flat pass over the true topk width, so padding a list with -1 regrouped its pairwise tree and moved last bits on ~1% of calls — while the real kernel walks fixed 64-wide blocks where an all-masked block rescales by exp(0)==1 and adds 0, making padding bitwise free. A GPU-less oracle that is not blocked the same way fails parity the GPU passes. Verified here: 0 mismatches in 200 trials x 3 pad widths, and the indexer einsum is width-invariant on the valid columns (0/200). TAKEN: the width-invariant selection contract, and the two-tier framing. IMPROVED, and this branch ships the difference. Both selections are width-invariant, but the tie-admission form (k-th value + cumsum) emits picks in ascending INDEX order, while topk hands sparse_attn its picks in DESCENDING SCORE order — and sparse_attn's per-block reduction is order-sensitive, so re-laning the same support set regroups the sum. Measured at real dims, moe_eager vs the vendored reference over 40 steps: ascending-index lanes 2/40 bit-exact (max 9.4e-2); a STABLE DESCENDING SORT — the same (value DESC, index ASC) total order, but keeping the reference's lane order — 40/40 bit-exact (0.0). The sort costs a little more and removes a whole class of Tier-2 noise. Bucketing keeps its floor at index_topk, as that branch's note requires: the read has to be wide enough to hold the k picks. Tests re-cut around what is actually provable. The long-run bar is no longer "equals the reference" — it cannot be, because neither the reference's tie order nor its lane order is part of its contract — but TIER 1: the read width is a COST KNOB, max vs bucketed giving torch.equal hidden/logits/KV over 120 steps on seeds 7, 11 and 23. That bar needs no reference and would have caught the original bug on its own. Steps 40 -> 120 and the seed sweep are that branch's point: the first cut failed on 7 and 11 at different positions and PASSED on 23. pytest tests/test_v4_whole_layer.py: 6 passed, 2 skipped at OMP_NUM_THREADS=1 (CPU) and 6 passed, 2 skipped on the 5090 (tilelang). Multipliers hold with the width-invariant selection in place — the sort is within run-to-run noise: island 1.22x, moe_eager 2.15x, whole 7.31x wall / 12.0x cpu. * perf(v4): make the routed experts the grouped-MoE bank, not a copy of it The grouped fp4 MoE kernel measured 3.19x on a one-layer bench and then never fired on a real stage. Its per-step gather needs a contiguous [E, N, K] expert bank, and the only way to get one was to stack the reference's per-expert Parameters on the first decode token -- a second copy of the layer's weights, ~3.2 GiB per layer at the shipped dims. A seven-layer stage on a 32 GiB 5090 cannot hold both, so the bank build checked free VRAM and declined. Measured on the card: of a seven-layer stage, exactly ONE layer got a bank and six declined. So build the bank at load instead of budgeting for it. Stage.__init__ now calls v4_moe_grouped.bank_layout() between constructing the Blocks and loading them: it allocates each layer's bank and repoints every expert Linear's parameter at its slice (p.data = bank[j], zero-copy -- a slice of a contiguous bank is contiguous), freeing the tensor the constructor allocated. load_state_dict then writes THROUGH those views, so the checkpoint lands in the bank and nowhere else. One copy on the card; the bank costs nothing. Keeping the per-expert Linears as views rather than deleting them is what solves the s > 1 edge. Prefill, a verify chunk, a hash-routed layer and any world_size > 1 rank all still take the reference MoE.forward, which reaches the weights through experts[i].w1.weight -- same objects, same dtypes, same contiguous bytes, now addressing bank memory. That path runs unchanged and stays bit-exact by construction, rather than needing an s > 1 grouped kernel that would not be bit-exact anyway (grouped-and-padded MoE is not token-count invariant). empty_cache per weight kind is load-bearing, and only shows up on hardware: the freed per-expert blocks are ~4 MiB, the bank is one ~1 GiB request, and the caching allocator cannot serve the second out of the first. Without it memory_allocated stays flat while memory_reserved -- what the driver sees -- climbs to 19.75 GiB on a stage holding 10.17 GiB of weights. Measured, RTX 5090, converted 43-layer checkpoint, real weights, stage[40:43): 10.168 GiB allocated with the layout and 10.168 GiB without it (against 19.730 by stacking), 3/3 layers banked, none declined. At shipped dims, stage[0:7): 23.641 GiB either way, 7/7 banked, where stacking took 26.829 GiB for 1/7. Decode step 26.902 -> 18.599 ms eager and 23.627 -> 15.387 ms under V4_CUDA_GRAPH=1; MoE.forward alone 4.75 -> 2.02 ms, the honest multi-layer 2.35x. 41 decoded steps torch.equal to the same run with the lever off, both graph settings; the GPU parity harness now also checks decode AND s > 1 over the bank layout. Default OFF is untouched: with V4_MOE_GROUPED unset nothing is allocated and nothing is repointed. * docs: V4-Flash engine living state There was no cross-session anchor for the V4 engine the way M25_ENGINE.md anchors M2.5, so each session was re-deriving the same facts from branches and logs. Records what is measured rather than what is hoped for. The headline: pipelined speculative decode runs 3.68 tok/s on a 6-box scattered ring, bit-identical to greedy, receipts verified out-of-process against a tamper test. That is 1.94x greedy and 3.5x serial. Also records the three things blocking 20 tok/s, so they are not rediscovered: the ring is ~11-15% pipeline-efficient (bubbles, not compute, are the prize); six 32GB cards hold ~45 layers against the model's 43, which makes the ring VRAM-bound and caps any re-tiling at 1.15x; and acceptance regressed 4.0 -> 3.05 under two levers, which cancelled most of their gain. Includes the ring-ops lessons that cost real time this session — screen boxes on CPU load twice to separate a provisioning burst from a co-tenant, place by role because the tail runs the drafter, and dedup harder than machine_id after two "distinct" hosts reported identical load triples. * fix(v4): release the experts before the grouped bank asks for them The bank layout was measured with `memory_allocated`, which is flat across it to the byte -- the per-expert storages really are freed, and the banks really are the only copy. What was never measured is the interval. `_relayout_moe` allocated each [256, N, K] bank BEFORE walking the experts it replaces, so every bank was a request for memory the layer was still holding: +1024 MiB of live bytes per weight kind at the shipped dims, and the blocks that came back behind it were 256 scattered 4 MiB per-expert allocations sharing large-pool segments with the kinds not yet relaid -- the wrong shape to satisfy the request that freed them, and not wholly-free segments, so `empty_cache` could not return them either. `memory_reserved` therefore climbed by every bank built so far and only fell at the END of the layer, making the real peak steady + a whole layer's experts (3.19 GiB), not the +1.07 GiB the allocated high-water mark suggested. That is why seven layers measured clean and eight died. Against a 30.76 GiB budget (31.36 usable less the ~0.6 GiB CUDA context) a head stage[0:7) is 24.62 GiB steady and peaked at 27.81; stage[0:8) is 27.98 and peaked at 31.17. The reported OOM reads back as exactly that: 27.97 GiB allocated (26.998 of layers + 0.986 of embedding, flat), 2.13 GiB reserved but unallocated (= 1024 + 64 + 1024 + 64 MiB, the w1/w1_s/w3/w3_s banks already built in _BANK_KINDS order), 669 MiB free, and a 1024.00 MiB request -- the w2 weight bank, fifth of six -- that cannot be met. So invert the order. `_relayout_moe(moe, preserve=False)` rebinds every routed-expert parameter of the layer to a void tensor first, hands the now wholly-free run back with one `empty_cache`, and only then allocates the six banks, into the 3.19 GiB the layer just vacated. Peak allocated and peak reserved both equal the steady state. `preserve` defaults to True and keeps the copy for the caller that has already written its weights (the GPU parity harness's standalone layer); `Stage.__init__` is the only caller that passes False, and it may because the Blocks are two statements old, every routed-expert byte is still the constructor's uninitialised `torch.empty`, and `Stage.load` then load_state_dicts every one of them strictly through the views. Gated on CPU at the REAL expert count. `_Meter` puts a weakref finalizer on each UntypedStorage's PyObject, so a storage is counted out the instant its last reference dies, and runs the ledger the failure was about: bytes requested from the allocator against bytes released to it. On the stage's path peak == steady == baseline and the debt is 0; with the old ordering the same meter reports the layout holding 1.00 of a whole bank twice, which is the test that fails on 24633c0 and passes here. A companion test pins that the toy dims are the shipped layer scaled (fp4 packs 2/byte, e8m0 is 1 per 32, so each weight bank is exactly 16/51 of the layer whatever dim is), and a new CPU parity test runs two MoEs from one state dict -- one plain, one laid out release-first before loading -- through the reference MoE.forward on real fp4 weights: torch.equal at s=1, 3, 9. tests/test_v4_moe_grouped.py 26 passed 2 skipped; test_v4_pipe.py 57 passed; phase0/v4_pipe.py selftest ALL PASS (10/10). phase0/deepseek_v4_ref/ untouched. * test(v4): prove every lever in the composed stack actually fires The stack is now seven opt-in flags deep and the expensive failure is not a wrong answer, it is a meaningless one: m25 banked "CUDA graphs don't help" off a ring where a forgotten env made every job run eager. Nothing in that run could distinguish "the lever fired and did nothing" from "the lever never fired". So each flag is armed in a FRESH INTERPRETER and the process reports the state it reached. Subprocesses are not fussiness -- every one of these is read at module import, so an in-process monkeypatch proves nothing about a stage that imported the module first. The assertions read Stage.__repr__, which is the same surface an operator reads on the ring, so a green test and a green ring log are the same evidence. 21 reachability proofs, plus two that cost 22s each and are worth it: the ring recipe serves exactly what the default serves. ALL PASS is NOT this bar and cannot be: the selftest compares each config's ring against THAT CONFIG'S reference, so a lever that moves both together passes in-config while changing what the ring serves. Only a cross-config comparison catches it. and it caught one. V4_REF_SLIM_NOQAT moves the tokens from step 3 -- it removes the reference's deliberate fp8/fp4 QAT precision simulation, so the run is strictly MORE precise and still a different answer. Documented APPROXIMATE in v4_ref_slim, invisible to the selftest, and disqualifying for an arm whose headline claim is "bit-identical to greedy". It is asserted DIVERGENT rather than merely left out, so it cannot drift into the recipe later without someone deciding to. Also removes a merge artifact in v4_stage: the islands branch and the whole-layer branch each added the V4_CUDA_GRAPH block, git kept both, and the bool copy sat above the mode copy as dead code that read as live. Exactly the shape of thing this file exists to catch. 23 passed. Full v4 suite 392 passed, 5 skipped. * feat(v4): measure what one V4 ring hop actually puts on the wire V4's inter-stage payload is h [b, s, hc_mult, dim] — the hyper-connections keep four residual streams per token the whole way down the stack and only collapse at hc_head — so a hop moves 4x what an ordinary pipeline-parallel model would. That makes the wire the thing to argue about on a scattered WAN ring, and the argument was being made from shape arithmetic on paper. This builds real frames with the engine's own _make_step_frame and measures them through the real transport codec, so the JSON header and the length prefixes are counted where they fall. At the shipped config (dim 4096, hc_mult 4): frame s bf16 B fp8 B x prefill, whole prompt 8192 268,501,190 134,349,081 1.999 prefill, 2048 prompt 2048 67,125,446 33,587,481 1.999 dspark verify chunk 6 196,848 98,672 1.995 decode, s=1 1 32,968 16,672 1.977 Two things fall out that were not obvious. The framing overhead does NOT eat the saving on the small frame — 192 B of a 32,968 B s=1 decode frame, 0.58% — because a single V4 token is already 32 KiB of hyper-connections; on a one-stream model that frame would be 8 KiB and the case would be much weaker. And a bf16 prefill at V4_MAX_SEQ's own default of 8192 is 268,501,190 B against transport's 268,435,456 B MAX_FRAME: the ring cannot prefill its advertised context in one frame at all, and recv_msg refuses the length before allocating. The ceiling is 8,189 positions bf16 against 16,367 fp8. * docs(v4): the composed full-stack launch recipe, and a launcher that can reach it docs/V4_FULL_STACK.md is the branch's operating document: the exact env for the ring, which levers compose with pipelining and which are exclusive (with the dispatch reason, not an assertion), the merge decisions, and the full test matrix. Two corrections to docs/V4_PERF_ROUND_2.md, which is now marked superseded for launch but kept as the per-lever source of record: - the grouped MoE no longer declines on a full stage. The bank layout landed, so the real multi-layer numbers replace the one-layer bench's 3.19x: MoE.forward 2.35x, whole stage 1.45x eager / 1.54x graphed. - V4_CUDA_GRAPH is a mode, not a bool. And the launcher can now SAY the mode. stage_launch_cmd only ever emitted V4_CUDA_GRAPH =1, i.e. island, so reaching whole meant remembering an extra_env override -- one forgotten string between "measured island" and "reported whole". It takes the mode as an argument and RAISES on a value the stage would resolve to off, because a ring that silently runs eager while the launch line says otherwise costs a provisioning cycle to discover and produces a number nobody can interpret. GRAPH_MODE_VALUES is duplicated out of v4_stage._graph_mode rather than imported -- stage_launch_cmd is a pure builder and tests/test_v4_pipe.py runs it with no model on the box. A test asserts the two lists agree in both directions, which is what makes the duplication safe. * test(v4): pin what binds pipelining to the graphs' Tier-1 bar A pipelined cancel calls Stage._replay to rebuild the window ring and both compressor accumulators over the accepted prefix. _replay sets _replaying=True and _run's graph gate excludes it -- so a rollback rebuilds that state EAGER while the frames that originally wrote it went through the GRAPH. If graphed and eager differ by one bit, every cancel leaves the stage holding state the un-cancelled run would not have had, silently, behind valid receipts. Pipelining cancels often: the CPU selftest shows 3 in 4 cycles. So this is the common path. Which means the whole-layer graph's Tier-1 bar (graphed == its eager twin, torch.equal, incl. across bucket crossings) is the CORRECTNESS PRECONDITION for composing V4_CUDA_GRAPH with V4_PIPELINED_SPEC, not a quality nicety -- and Tier 2 alone would not do, because an approximate-but-defensible graph still poisons every rollback. Nothing stated this, and a CPU box cannot measure it, so it is pinned structurally: the replay flags itself, the graph gate excludes it, and the replay drives _run per position. Noted and NOT taken: _replay is s=1 per position and discards its outputs, so it could use the graph -- faster, and it would dissolve the dependency. That is a behaviour change to a separately verified module and wants its own measurement. * perf(v4): scale the fp8 wire per position, not per tensor V4_FP8_WIRE halves the bytes on every hop, which is the right trade on a scattered ring — but it shipped with one scale for the whole tensor, and that quietly costs acceptance on the DSpark path. The reduction axis is a CORRECTNESS property here, not an accuracy one. e4m3 is a float: it carries its own 4-bit exponent, 2^14.8 of normal dynamic range, so a coarser scale buys almost nothing in error — measured over the boundary tensors of a real V4 stack, a per-tensor scale is 1.03x worse than a per-128-block one even when the chunk spans 2^-20 of dynamic range, while block scales cost 4-32x more sidecar bytes. Granularity is simply not where the accuracy lives. Where it decides something is speculation. A scale reduced over the SEQUENCE axis makes a position's transmitted bytes depend on which other positions shared its frame, and DSpark deliberately sends the same position in chunks of different lengths: the first round of a generation is a bare [cur] chunk (s=1), every round after carries a whole block (s=dspark_block_size+1). So the drafted stream stops being bit-identical to the greedy stream — the property coordinate_dspark's accept rule is built on — for a reason that has nothing to do with the model, and the verifier starts rejecting drafts the drafter derived from the same weights. Measured on the CPU ring with a drafter proposing that ring's own greedy continuation, so a self-consistent wire must accept everything (K=3, ceiling 4): bf16 wire g = 4.000, 15 rounds, stream == greedy fp8, per (position,stream) g = 4.000, 15 rounds, stream == greedy fp8, per tensor g = 1.714, 35 rounds, stream != greedy Reducing over `dim` alone makes position p's packing a function of position p, and costs 2 B per (position, stream) — 8 B against 32 KiB per token, 0.024%. Three numerics fixes ride along, each with a test that fails without it. The amax is taken in fp32, because clamp(min=1e-8) coerces its bound to the tensor dtype and in float16 1e-8 underflows to 0.0, so an all-zero row divided by zero. The quotient is clamped to +-448 (the vendored act_quant's own contract), because torch's e4m3fn cast NaNs past ~464 rather than saturating. And the scale is clamped to bf16 max so a single inf cannot poison the values sharing its scale: an infinite amax sends every finite neighbour to 0 and the inf itself to NaN, turning one bad slot into a whole bad row. _recv_hids now cross-checks h.dtype against the presen…

leyten and others added 10 commits June 19, 2026 09:59

docs: STATE — step 3 (content-addressed weights) is the remaining JOI…

e246973

…N piece

leyten force-pushed the master branch 2 times, most recently from fd32138 to df9792d Compare June 21, 2026 14:11

leyten mentioned this pull request Jul 15, 2026

feat: python -m shard.stage + P0-#2 portability fixes #104

Merged

leyten mentioned this pull request Jul 15, 2026

docs: living-state resume + LAUNCH progress for the Leg-7 session #105

Merged

leyten mentioned this pull request Jul 18, 2026

docs: tick P1-#2 live-ring validation #117

Merged

leyten added a commit that referenced this pull request Jul 18, 2026

docs: tick P1-#2 live-ring validation (#117)

6248e70

warmring-20260718: coordinate served a real ring job, receipts pinned-verified, nonce threading proven. P1-#2 residue = assignment-EPOCH + pay-model.

leyten mentioned this pull request Jul 20, 2026

docs: session 07-20 — epoch fix landed, map wired live; next = EAGLE watchdog #118

Merged

leyten added a commit that referenced this pull request Jul 20, 2026

docs: session 07-20 — epoch fix landed, map wired live; next = EAGLE …

2259736

…watchdog (#118) P1-#1 done (map live, sim fallback); P1-#2 code-complete (epoch settlement #37). Next: EAGLE P0-#5 mitigation-first.

leyten mentioned this pull request Jul 20, 2026

docs: audit launch list to 07-20 — pay-model built, dev-safety lesson, next=EAGLE #119

Merged

leyten mentioned this pull request Jul 20, 2026

docs: mark c0mpute pay-model #41 as merged (was stale OPEN) #124

Merged

leyten added a commit that referenced this pull request Jul 20, 2026

docs: mark c0mpute pay-model #41 as merged (was stale OPEN) (#124)

b9def16

c0mpute PR #41 was merged to master this session, gated off until launch; LAUNCH.md P1-#2 and the M25_ENGINE LATEST-2 block still read OPEN.

leyten force-pushed the master branch from 318db57 to 600f4ff Compare July 22, 2026 21:58

leyten added a commit that referenced this pull request Jul 22, 2026

docs: tick P1-#2 live-ring validation (#117)

935c572

warmring-20260718: coordinate served a real ring job, receipts pinned-verified, nonce threading proven. P1-#2 residue = assignment-EPOCH + pay-model.

leyten added a commit that referenced this pull request Jul 22, 2026

docs: session 07-20 — epoch fix landed, map wired live; next = EAGLE …

d566a2f

…watchdog (#118) P1-#1 done (map live, sim fallback); P1-#2 code-complete (epoch settlement #37). Next: EAGLE P0-#5 mitigation-first.

leyten added a commit that referenced this pull request Jul 22, 2026

docs: mark c0mpute pay-model #41 as merged (was stale OPEN) (#124)

c17dcb0

c0mpute PR #41 was merged to master this session, gated off until launch; LAUNCH.md P1-#2 and the M25_ENGINE LATEST-2 block still read OPEN.

Navigation Menu

Mid-generation self-healing: recovery decision logic - #2

Mid-generation self-healing: recovery decision logic#2
dddabtc wants to merge 10 commits into
leyten:masterfrom
dddabtc:feat/self-heal-recovery-core

dddabtc commented Jun 20, 2026

Uh oh!

dddabtc commented Jun 20, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

dddabtc commented Jun 20, 2026

What

Scope of this PR

Deliberately deferred (named, not hidden)

Follow-ups (need a multi-node GPU rig to validate end to end)

Uh oh!

dddabtc commented Jun 20, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants