Disable NCCL_NVLS by default by Ying1123 · Pull Request #631 · sgl-project/sglang

Ying1123 · 2024-07-16T10:23:23Z

No description provided.

…backends are closed The user supplied the real target hardware, which the 2026-08-07 revision never had: a ThinkPad P14s Gen5 AMD -- Ryzen 7 PRO 8840HS, Radeon 780M iGPU (gfx1103, RDNA3), 32 GB DDR5 shared between CPU and iGPU. The old revision assumed a discrete NVIDIA laptop and its conclusions are wrong. The geometry, the sgl-project#647 fix, and the parallelism analysis survive unchanged and are kept. The laptop is REACHABLE from the rig box (root@192.168.0.116, efeu-TP14); the earlier "unreachable" note is stale. Everything in section 1 is measured on it. BOTH BACKENDS ARE CLOSED, each for an independent, evidenced reason. iGPU via ROCm: the GGUF K-quant kernels are not in the ROCm build at all -- not broken, absent. The sources exist (8328 lines under csrc/quantization/gguf/) and CMakeLists.txt:323 compiles them for CUDA, but setup_rocm.py:43-60 is the whole ROCm source list and contains no csrc/quantization/ entry; there is no hipify step. common_extension.cc:428-480 binds every ggml op to torch::kCUDA, while the ROCm build compiles common_extension_rocm.cc (setup_rocm.py:47), which has zero ggml_* symbols -- so the op does not exist as a SCHEMA on ROCm. The omission is deliberate, not a stale shared list: setup_musa.py:91 does list the file. Independently, setup_rocm.py:77-81 exits 1 for any arch other than gfx942/gfx950, so sgl-kernel will not build for gfx1103 at all. Worse, it fails silently. is_cuda() (common.py:146-148) is False on a ROCm torch build and is_hip() is True, and in gguf.py:41-62 the ggml imports AND their None fallbacks are all nested inside `if _is_cuda:`. So on ROCm the names are never bound and the first forward dies with a bare NameError (gguf.py:840, :927, :932, :1010, :1092, :1107) -- after a clean load. The warning at :78-79 is suppressed on exactly this hardware (`if not _is_hip`), and supports_current_device() (:134-152) returns None off CUDA so _enforce_capability_floor abstains. Do not read a clean load, or the absent warning, as support. The tree already documented the conclusion twice and it was missed: quantization.md:39 (gguf ROCm = No) and amd_gpu.md:119. CPU-only: no CPU K-quant kernel exists (get_quant_method, gguf.py:165-192, branches only on _is_npu), so a CPU stage must materialize dense at the measured 3.17x (506.2 -> 1604.2 MiB/layer) = ~67 GiB against 29.5 GiB of MemTotal, a 2.27x overshoot. Dense bf16 is quant-independent, so the Q2 checkpoint lands on the same ~67 GiB and does not rescue it. This is arithmetic, not a tuning problem. CONSTRUCTIVE HALF: the port is mechanically shallow. The gguf kernels contain no CUDA-only hardware intrinsics (zero hits for __ldg, cp.async, asm volatile, mma, wmma, __shfl, __ballot), the AMD shims are already inherited -- ggml-common.h:1019 supplies __vsubss4, __dp4a via __builtin_amdgcn_sdot4 (:1046-1048) and __vcmpeq4 -- moe.cuh carries ~20 USE_ROCM tuning branches, and WARP_SIZE_GGUF 32 (ggml-common.h:6) is already correct for RDNA wave32. So section 7 orders the real work: widen the gfx allow-list, add the source + register the ops, bind an elif _is_hip: arm in gguf.py, then validate numerics on RDNA. MEASURED BUDGET (docs/dev/651/apu_budget.py, new). MemTotal 30211 MiB; BIOS UMA VRAM 1024 MiB, which the user confirms is the BIOS MINIMUM and cannot go lower; GTT 24576 MiB pinned by amdgpu.gttsize=24576. The binding ceiling is therefore GPU-addressable 25600 MiB, not MemTotal, and GTT is backed by the same DDR5 -- it caps what the iGPU may pin, it is not extra memory. Against that ceiling Q4_K_M (21614 MiB) leaves 3137 MiB of slack and Q2_K_XL (11992 MiB) leaves 12758. Context is not the binding constraint -- weights are; even Q4 clears ~82k tokens at fp16 KV. The laptop carries Q4_K_M and Q2_K_XL in /root/lh/models/, NOT the rig's Q4_K_XL, so the earlier "no smaller fallback without a new download" is false there: staging now starts on Q2_K_XL. RESHARD ON SHARED MEMORY (user-confirmed): the PP-prefill/TP-decode flip is logical only -- everything already lives in one RAM, so it is an ownership/view reinterpretation and bytes never move. Route A's hardest problem, that the two layouts want different weight bytes per rank, simply does not exist with one pool and one process. Latent optimization recorded: on unified memory the PD KV transfer could hand over ownership, where today the code copies unconditionally. ACHIEVABLE GOAL, once a backend exists: TP=1/PP=1 + NEXTN speculation, CPU/RAM as the only tier, --cpu-offload-gb for headroom (composition with GGUF unverified -- first test), MoE expert-offload walled for GGUF by sgl-project#123. PP+spec stays mutually exclusive tree-wide (server_args.py:16264 one-server assert; both PD arms reject spec at pd_disaggregation_hook.py:194-229) -- owned by Route A / sgl-project#631, not built here. Carried forward: --pp-layer-ratio sums to backbone depth 40 not block_count 41, max_total_num_tokens is min-reduced across the world group, and the sgl-project#647 fix is 0155ff2. Also noted: docs/dev/651/boot.sh is NVML/CUDA-only (pynvml UUIDs, CUDA_VISIBLE_DEVICES, nvidia/cu13 LD_LIBRARY_PATH) and must be rewritten for ROCm before use; its staging logic still applies. No GPU ran any of this. The model has still never generated a token.

…ialization host-OOM at flip boot payload.to(int64).sum() materializes an int64 copy of the whole payload (8x its size as one transient; torch sum(dtype=int64) over an integral tensor casts the input first, measured on CPU). At the first real-metal INT8-W8A8 flip boot every rank hit this in image_from_tensors during the PP weight snapshot: ~90+ GB of transient host allocations across the three ranks, host OOM killer SIGKILLed rank 2 right after cuda_graph_capture. Reproduced twice; 500-ms memory trace on file (/tmp/route-a-631/boot_mem_trace.csv). Fix: uint8_checksum() walks the payload in 16 MiB chunks with an int64 accumulator per chunk -- transient bounded at 128 MiB, device-agnostic, one host sync. Same value bit-for-bit (equivalence test included). Applied to the whole checksum family: weights_arena (3 sites, the boot killer), gdn_flip_mover, kv_reshard (2 sites), phase_flip_runtime. Falsifier-first: TestChecksumMemory peak-RSS gate fails on the old idiom (3386 MiB peak for a 256 MiB payload) and passes on the fix; family suite 219/219 (run_631_flip_family.sh).

…ild path The hermetic suite never executed this path on hardware; the first INT8-W8A8 PP=3 flip boot walked a ladder of desk-written-never-executed faults, each boot-validated in sequence (attempts 3-9, all logs under /tmp/route-a-631/): 1. scheduler.gpu_id does not exist -- the Scheduler keeps it on its ParallelState (scheduler.ps.gpu_id). 2. --pp-layer-ratio exports SGLANG_PP_LAYER_PARTITION process-wide and the TP stack's pp_size=1 model build dies on the 3-way partition in get_pp_indices. phase_flip_tp_scope now masks the variable for the build and restores it afterwards (the runtime's pp_size=3 layer-map derivations need it back). Pinned by TestTpScopeEnvMask. 3. attention-registry MTP shortcut (full_attn_layers=[0]) mis-classified every GDN layer of the flip TP stack, which rides is_draft_worker for the secondary-runner gates: RadixLinearAttention was routed into flashinfer and died on is_cross_attention. Added the is_phase_flip_tp_stack exemption (same shape as the dual-group lane target exemption above it). 4. same is_draft_worker-ride family in FlashInferAttnBackend: the draft-replicated exemption forced uneven_dcp=False for the flip TP stack, so decode wrote local-head KV (8,1,256) into the full-width token-sharded pool (row_dim 1024) -- store_cache row-count mismatch at graph capture. The store_cache call site now re-raises with the full call geometry (k/v shapes, row_dim, cache shape, index dtype). 5. build_phase_flip_runtime read full_attention_layer_ids off sglang's ModelConfig; for hybrid GDN models it is a property of the HF text config (Qwen3NextConfig) -- read hf_text_config, mirroring the attention registry's mambaish_config access. With these the boot reaches serving (warmup prefill through all three PP ranks, transport barlink) at RANK_MIB 17000,11000,11000; the subsequent kv-consensus-under-PP wedge is a separate fix. Family suite 220 passed.

…icrobatch iteration Measured wedge (first serving flip boot, 2026-08-08 11:06Z, log preserved as server.log.attempt9-WEDGE-kvpressure-consensus): right after the warmup prefill, PP0+PP1 sat in the flip runtime's bounded MIN-reduction while PP2 sat in recv-from-PP1; barlink liveness caught the no-progress after 120 s on every rank and the tree self-terminated cleanly (the sgl-project#622 liveness machinery working as designed). Root: the rank-local-state-feeds-collective family in PP form. The hook ran inside get_next_batch_to_run -- the TOP of the pp iteration, before this rank's sends are issued -- so a rank could enter the blocking world-reduction still owing the send its successor needs to reach ITS reduction. Under event_loop_normal (lockstep TP rounds) the placement is safe and stays; under event_loop_pp the hook is now deferred (_defer_flip_round_to_pp_loop) and runs at the END of each microbatch iteration, after every send of the iteration is flushed -- the reduction becomes the last blocking op of the iteration and no recv/reduction cycle can close. PhaseFlipLoopExit now raises from that quiescent boundary. Extracted _phase_flip_on_round() serves both call sites; the pp loop resets the defer flag on any exit so the post-flip TP loop runs the hook inline again. Falsifier-first: TestPpLoopConsensusOrdering drives the REAL on_round consensus through a bounded barrier channel in the measured composition (last stage recvs mid-iteration before its hook, middle stage's send trails its hook): top placement deadlocks (negative control, broken barrier), end placement completes. Family suite 222 passed. Also books the flip-build VRAM transient (TP originals + arena coexist at arena allocation; measured deficits 4.4/1.7 GiB at the old budgets) into the 3.4a ledger as its own term.

…t on cadence Boot 10 refuted the end-of-iteration placement as sufficient: the ranks' ABSOLUTE round counts diverge under event_loop_pp (pipeline fill, conditional per-slot ops), so PP0 sat in the reduction at its own round boundary while PP1+PP2 waited in the commit send/recv for data that sits behind PP0's reduction -- same wedge, new composition (log preserved). Under the pp loop the reduction is now entered only when this rank is ARMED (there is nothing to agree on otherwise; arming arrives on every rank via the broadcast RPC) AND locally PARKED (ready_fn: drained microbatches, no partial chunk) -- a parked rank owes no pipeline send, so no recv/reduction cycle can close, and MIN-skew across ranks stays legal while peers converge on their own arm+drain. Unarmed serving under PP therefore performs ZERO flip collectives. event_loop_normal (lockstep TP rounds) keeps the periodic consensus unchanged. A flip under continuous load needs a posted-async two-phase consensus -- named follow-up, out of this slice. Pins: TestArmedParkedGate (unarmed and unparked rounds touch the channel zero times; armed+parked enters), TestPpLoopConsensusOrdering kept as the ordering invariant's negative control. Family suite 225 passed.

…not channel totals First armed flip on metal: the honesty-contract refusal fired with 'head shards (22, 12, 12) do not partition 48 heads'. The mover partitioned the value CHANNEL total (6144 x 30/64 = 2880 -> floor 22 heads) with the PP pool's partition_units, which are absent on the tp=1 pool -- while the model splits HEAD COUNTS in whole gdn_tp_units (qwen3_5 local_num_*_heads): (24, 12, 12) under 30,17,17. The uneven-TP-head-geometry class: per-rank geometry comes from the model, never recomputed. gdn_flip_preconditions now takes the hybrid HF text config (linear_num_*_heads + the model-stashed gdn_tp_units) and replicates the model's tp_partition_size-over-heads calls; the legacy channel-total fallback stays for geometry-less callers and its refusal stays reachable. Pins: TestRealConfigHeadGeometry (Qwen3.6-27B constants, 3.4b item 3 -- model-identical (24,12,12) accepted per rank, channel-split negative control still refuses loudly). Family suite green.

…ip-time fix (device-side refill verify) + measured ledger terms POST-FLIP WEDGE (first serving attempt after a completed cutover, 2026-08-08, py-spy on file): the cutover swapped scheduler.ps to a new instance but the components that captured the old object kept PP semantics -- the request receiver relayed requests PP-chain-style (point_to_point from a stage that now runs TP broadcast semantics; one rank wedged there, another in the pool-budget all_reduce) and the output streamer mis-gated the detokenizer send (heartbeat loss). Cutover step 4b now rebuilds request_receiver (ps + tp/attn_tp group handles via dataclasses.replace), output_streamer and load_inquirer against the freshly-routed handles, and the step-9 completeness self-check pins each holder so a future component joining the snapshot list fails loudly at cutover, not as a wedge. The protocol suite's stub scheduler now carries the step-4b holders, so the rebuild list runs under the REAL cutover in the hermetic tests. FLIP-TIME ECONOMICS: the measured 22-33 s flip wall (vs ~2-2.3 s design estimate) decomposes to the HOST-side image checksum in arena_refill -- a single-core uint8 sum at 0.82 GiB/s (measured; 3 concurrent ranks share host bandwidth), not the H2D copy. arena_refill now copies first and verifies on the ARENA's device (device bandwidth); a checksum mismatch restores the CURRENT phase's layout from its creation-time- verified image so the abort stays clean (restore pair threaded through PhaseFlipStacks.refill). Falsifiers updated: corrupted refill restores the active layout's views byte-identically; the no-restore path flags the arena content as undefined. LEDGER (3.4a updated): NCCL two-group sets MEASURED 198/112/112 MiB per card -- the ~500 MiB guess retired; host pinned images booked as a new term (28.1/14.6/17.0 GiB per rank, ~60 GiB total); flip-time decomposition recorded. Family suite 228 passed.

…ge-residency decision RAM basis correction (coordinator note): the box has 98 GiB and no swap; lxcfs's 120 GiB /proc/meminfo is a known lie on this rig (only cgroup memory.stat is truthful). The 59.7 GiB of pinned flip images are therefore >60% of the machine; the remaining host budget for ALL host-RAM consumers together is ~30 GiB and is now stated explicitly. Image residency priced with a measured rig figure (ZFS pool, 869 MB/s O_DIRECT / 1.6 GB/s buffered = 0.6-1.1 ms/MiB, 2-4x the laptop NVMe): both-images-on-disk is rejected (the next-flip target image is on the hot path, +9-15 s per flip); cold-image spill with background promotion (hot image pinned, current phase's restore image on disk, roles swap per flip) saves ~28 GiB steady at the price of a ~7-15 s re-read on the rare abort-restore. Decision: that is the design, deferred as the named follow-up -- the corrected budget suffices for the current slice, and the spill choreography needs its own falsifiers (torn writes, promotion-not-finished-at-flip).

…llectives that have no wire Three defects, all measured on window-2 boot 13 (2026-08-08), all in the same family: rank-local state feeding a group collective. 1. CENSUS DOMAIN. The census counted collectives on world_size==1 groups. Those short-circuit without touching the wire, so their counts measure this rank's local work; comparing them across ranks is a category error. Under PP=3/TP=1 with stage ratio 2,1,1 the size-1 "tp" group reported "tp.all_gather: counts [536, 1096, 1096]" and the detector called a correct configuration a desync. The gate is now ONE invariant computed where world_size is first known -- an inlined copy of it read self.world_size above the block that assigns it, which raised AttributeError inside GroupCoordinator.__init__ and turned every group construction into a retry storm (2 tests red). This excludes non-wire events from a wire census. It does not weaken the check on any group that can actually desync. 2. CENSUS CADENCE. The detector's own comparison is a blocking collective on a group that also carries payload traffic. Its cadence premise -- every rank fires in the same round -- holds in the TP loop and fails in the PP loop, and a flip cutover swaps a size-1 compare group for a real one. Fired at drifted rounds it mispairs the group FIFO: the instrument seeds the exact desync it exists to catch. The round now rides in the payload, a drift stands the periodic comparison down permanently and says so, and the wedge-proof local dump stays armed. Re-zeroing at the group-aligned cutover restores the premise for the post-flip loop. 3. PARK DEADLINE. An armed flip withholds new work so the in-flight state drains, which is what lets it interpose between a request's prefill and its decode rather than only after every stream finishes. Unbounded, that park holds the requests of a rank that never reaches quiescence forever. A rank armed past the deadline now joins the reduction carrying `expired`, and every participating rank abandons the flip on the reduced maximum -- group-agreed, not rank-local. The FLIP is abandoned loudly; the parked requests are never aborted, which is why the abort path returns rather than raises. Falsifiers red first in every case: the drifted-cadence and round-key tests fail with the self-check removed; the never-quiescent-stream test fails with the expiry forced to 0. Tests: scripts/run_631_flip_family.sh 243/243 green (228 before, +9 census domain/cadence, +6 park deadline).

The flip stack had zero speculation plumbing, and pp_size>1 refused speculative_algorithm outright at argument time. Route A's target is one instance that prefills under PP and decodes under TP on the same ranks, with NEXTN live in the decode phase -- so speculation has to follow the TP stack, not the instance. It does that by being a property of the PHASE: - The argument-time rule is not waived. There is no PP-shaped draft worker (the constructors take no pp_rank), so none is built in a PP phase. The assert now enforces that by construction instead of by refusing the flag combination: PP + speculation is still refused unless --enable-phase-flip is set. - The scheduler boots with spec_algorithm NONE and draft_worker None, and keeps the configured algorithm aside as flip_spec_algorithm. Every spec-keyed branch in the PP phase therefore takes exactly the path it takes on an instance without speculation. - phase_flip_boot builds the draft worker on the TP stack, inside the flip scope with the TP server args published -- a draft built against the target's PP geometry would shard its heads for the wrong topology. It shares the TP stack's request pool and KV allocator, exactly as Scheduler.init_memory_pools does at boot, so draft KV is sized rank-locally inside the already-profiled TP budget rather than profiling a second time against memory the two resident pools have already claimed. Its weights are its own model and are NOT arena-backed: there is no second layout for them to flip between. - The cutover arms it with the stack it targets and disarms it on the return trip, and verify_flip_cutover pins both directions. A half-armed cutover -- the algorithm swapped in without its draft worker, or a draft worker left armed against the PP stack it was never built for -- is a loud refusal, not a round that runs on it. NGRAM is refused for this phase: its external corpus manager is wired to the tokenizer channel and is not on the cutover rebuild list. Half-arming it would be worse than not offering it. Tests: scripts/run_631_flip_family.sh 253/253 green (243 before, +4 argument-time gate, +6 cutover arm/disarm including both half-armed can-fail directions). The stack builder itself needs a real model and is covered by the metal boot, not by a desk test.

…king a construction gate what a draft model is Boots 14-16 of the spec-armed flip, each failing one step further in. 1. The flip's V1 blocker list still refused --speculative-algorithm, naming "the TP+NEXTN decode arm" as a follow-up. That arm shipped in the previous commit. The blanket refusal is replaced by the two shapes for which "armed in the TP phase only" is not a complete answer: ngram (its external corpus manager is wired to the tokenizer channel and is not on the cutover rebuild list) and solo draft placement (shadow-rank identity is not modelled across a flip, where every rank changes topology). 2+3. Two more is_draft_worker RIDE bugs, and the reason there keep being more. `is_draft_worker` is a CONSTRUCTION gate -- "build me as a secondary runner: no distributed re-init, no process-global installs" -- and it has THREE producers: a speculative draft worker, the sgl-project#274 dual-group lane, and the sgl-project#631 phase-flip TP stack. Only the first holds draft weights. A site that asks the construction gate when it means draft-NESS gets the wrong answer for the other two. The codebase already made this distinction once, for pools (is_draft_pool_worker, with a docstring telling pool sites never to ask is_draft_worker directly). It is now stated once for the general case as is_draft_model_runner, and the two sites that were newly reachable use it: - the KV pool-config branch demanded a caller-supplied memory_pool_config from the TP stack, which is supposed to RESOLVE one ("Draft worker requires memory_pool_config", all three ranks, boot 15); - the decode graph runner sent the TP stack down the draft branch and raised "This should not happen" (all three ranks, boot 16) -- correct in its judgement, wrong about the runner: it is not a draft one. Both were unreachable before this feature, because speculation and the flip were mutually refused. That is why they had not been found. Also adds scripts/route_a_631_acceptance.py: the acceptance driver. TTFT from a STREAMED request (first token's arrival, not whole-response wall time), decode tok/s excluding the prefill-produced first token, and the accept length read from meta_info["spec_accept_length"] -- not the server's rolling spec_ema_accept_len, which is a different quantity. Tests: scripts/run_631_flip_family.sh 258/258 green (253 before, +5 flip V1 blocker cases; the NEXTN pin in test_v1_blockers_named is retired with a note, since the behaviour it pinned is the thing that shipped).

Boots 17-22 of the spec-armed flip. The sequence now works end to end: PP=3 prefill -> live flip -> TP=3 decode with NEXTN speculating, and back. Four more defects, each found by a boot that got one step further. 1. EMPTY LIVE SET. Flipping with nothing live crashed every rank: the byte view inferred its row width with view(n, -1), which torch refuses at n == 0. An empty flip is not exotic -- it is an idle server, and it is exactly what a caller reaches after flushing the cache to make room for the flip. The width is now passed in. The plan layer had always handled the empty set correctly; only the byte view had not. 2. POOL FIT IS A RUNTIME QUANTITY, AND ITS REFUSAL MUST BE UNANIMOUS. The pre-move bound raised KvReshardError, which climbed into the event loop and killed a server that was serving perfectly well in its current phase. Worse, the reading is RANK-LOCAL -- each rank has its own pool sizes and compact rows -- so a rank that raised while a peer proceeded would leave the group half-flipped. Nothing is mutated at that point, so the verdict is now reduced across the group and, if any rank does not fit, every rank abandons the FLIP and keeps serving. Observed doing exactly that on metal: all three ranks, no bytes moved, serving uninterrupted. Whether the live set fits is not a boot-time property: it grows with the resident prefix cache, and the TP pool shrinks when a draft-KV allocation shares its budget. Both are now named in the message. 3. THE BATCH RESULT PROCESSOR CACHES THE WORKERS. It is on the decode hot path and calls into the spec worker (on_verify_complete_cpu), but it was built at boot, where a phase-flip instance deliberately has no draft worker. The first post-flip decode died with "'TpModelWorker' object has no attribute 'on_verify_complete_cpu'" -- the boot-cached target being asked to behave like the draft stack just armed around it. Added to the cutover rebuild list and to the completeness pin. 4. SPEC METRICS ON A NON-SPEC BATCH. The tokenizer manager indexed spec_verify_ct without the length check its sibling list already had. A spec-configured instance whose PP phase runs without a draft worker emits that field empty, so every PP-phase response raised IndexError. Measured, boot 22 (Qwen3.6-27B INT8-W8A8, 2x3080 + 1x5090, PP stage ratio 2,1,1, flip vector 30,17,17, NEXTN steps 3 / draft tokens 4): prefill (PP=3) 2048 tok 472.5 ms 4335 tok/s 8192 tok 1101.2 ms 7439 tok/s 32768 tok 4621.8 ms 7090 tok/s reshard pp->tp 837 / 1087 / 1563 ms per rank reshard tp->pp 982 / 1242 / 1345 ms per rank (1031 live slots, 6.4-8.5 MiB moved per rank; the KV move itself is ~5 ms -- the second is the weights-arena refill) TP decode+NEXTN TTFT 100.0 ms, 78.3 tok/s (natural prompt, 512 tok) accept length 2.23-3.80, accept rate 0.41-0.93, cuda graph active The accept-length spread tracks output content, as this rig's benchmark notes predict; the number is the scheduler's own spec_num_accept_tokens / spec_num_forward_ct, not the rolling EMA. Tests: scripts/run_631_flip_family.sh 261/261 green (258 before, +3 empty-live-set; the undersized-pool test is rewritten to the abandon contract, with the pools asserted untouched).

PP=3 prefill -> live flip -> TP=3 decode with NEXTN speculating, on the same three ranks, and back. docs/dev/631/FINAL_631.md carries the measurements, the two defect families the work uncovered, the operating notes, and the residuals. prefill (PP=3) 2048 / 8192 / 32768 tok 472.5 / 1101.2 / 4621.8 ms = 4335 / 7439 / 7090 tok/s reshard pp->tp 837 / 1087 / 1563 ms per rank reshard tp->pp 982 / 1242 / 1345 ms per rank, 1031 live slots (KV move ~5 ms; the rest is the weights-arena refill) TP decode+NEXTN TTFT 100.0 ms, 78.3 tok/s, accept length 2.23-3.80 (rate 0.41-0.93), cuda graph active Corridor: the acceptance budget breached the 1024 MiB/card floor at the BOOT peak only (961 MiB on the 5090, three samples). Trimmed to RANK_MIB 16150,10550,10550 and re-verified under load -- minimum free 2338 / 1221 / 1130 MiB, with the flip and speculation unaffected (903/1154/1627 ms, TTFT 101.2 ms, 75.8 tok/s, accept length 2.62). That is the configuration to run. Evidence json copied beside the document; the per-boot server logs stay in /tmp/route-a-631, each named for the defect it found. Tests: scripts/run_631_flip_family.sh 261/261.

…budget The Route A flip built its TP weights arena while the freshly loaded TP weights were still resident, so boot peaked at originals + arena on every rank (measured, torch allocated_peak_bytes): 29.27 GiB on the 5090, 16.64 and 17.81 GiB on the 3080s. The peak never recurs at runtime -- the serving state sits ~7 GiB below it -- but the corridor floor is a CONTINUOUS minimum, so the boot spike alone forced --rank-gpu-memory-mib down to roughly half of each card, and every MiB of that was permanent KV pool given away. Step 2 already solves this for the PP layout (snapshot to host, free the device originals, then load the second layout); step 4 simply did not do it for TP. image_from_tensors exists for exactly this and says so. The arena's host image was already built one line later, so this is a reordering, not a new allocation: peak becomes max(originals, arena) instead of their sum. Second defect, same area: the flip installed ONE vector for both the weight shard plan and the KV token split. Those optimise against different resources -- the weight shard follows compute, the token split follows each rank's memory left AFTER its weights land -- so sizing KV with the compute vector made the most compute-loaded rank bind the allocator's min-reduce and dragged the pool to its unit: rank 0: 12750 tok / ratio 30 = unit 425 <- binds rank 1: 68646 tok / ratio 17 = unit 4038 rank 2: 30515 tok / ratio 17 = unit 1795 -> global max_total_num_tokens 27200, ranks 1 and 2 left idle The server already computes the token-proportional vector and logs it as a restart hint, but nothing could act on it because this line overwrote it. parse_flip_token_vector reads SGLANG_UNEVEN_TOKEN_VECTOR and refuses a wrong length, a non-positive entry or garbage; unset it returns the flip vector, so the default path is byte-identical. Measured on the rig (Qwen3.6-27B-INT8-W8A8, 5090 + 2x 3080, barlink bar1, ctx 65536 -> 262144, --max-running-requests 4): max_total_num_tokens 46422 -> 278104 PP / 318176 TP max_running_requests silently reduced 4 -> 1, now a real 4 NVML free per card 2323 / 2158 / 2075 MiB idle, 2255 / 2146 / 2029 MiB under load (floor 1024) Tests: scripts/run_631_flip_family.sh 266/266 (261 before, +5 new TestFlipTokenVector cases covering unset/override/length/zero/garbage). Prefill ladder in the PP phase 4227.8 / 7183.8 / 6712.8 tok/s at 2048 / 8192 / 32768; decode in the TP phase with NEXTN live, cuda graph on, accept len 2.85-3.57.

… benchmarks Boot recipe, the BAR1 aperture budget that barlink x phase-flip needs, the KV sizing defects and their measured effect, all Nenngroessen read from the server, and the club-3090 prefill/decode numbers with the phase proof each was taken under. Also makes --kv-pressure-ladder opt-in: 'auto' refuses on this rig because it cannot map ranks to cards on a mixed-model node.

…ansition plan The flip kept ONE vector on PhaseFlipStacks and reinstalled it at every cutover. That vector is the WEIGHT shard split. Two consumers are token-space and were reading it: phase_flip_runtime.py set_cp_token_ratios(list(stacks.vector)) phase_flip_runtime.py tp_vector=stacks.vector The second is fed straight to build_phase_flip_transition, whose own docstring calls the argument 'the weighted DCP token vector of the TP layout'. So after a cutover the owner rule and the flip's row-routing plan both split rows under a vector the pools were NOT sized under. That is an out-of-bounds KV slot id, not a slow path -- store_kvcache's guard (SGL_DEVICE_ASSERT(index >= 0 && index < size_limit)) is what catches it, as a device-side assert that takes down every rank. Latent until now: the two vectors were equal by construction, so nothing could diverge. Making the token side overridable is what made this reachable, so it is fixed in the same breath. PhaseFlipStacks now carries both, named for the question each answers, and the two consumers read token_vector. Unset SGLANG_UNEVEN_TOKEN_VECTOR keeps them equal, so the default path is unchanged. Tests: scripts/run_631_flip_family.sh 268/268 (+2: the dataclass carries both vectors, and a source pin on the two consumers -- exercising them live needs a three-rank group).

…and the KV capacity design Adds the post-fix numbers (narrative 79.02 at CV 0.9 %, code 103.21 -- the wrong owner rule cost ~7 % and most of the run-to-run variance), the HiCache-carrier evaluation and why it is rejected, and the shared-arena design with its projected capacities.

…n row at the flip Prerequisite for the cross-phase shared KV arena. The local leg read and wrote per layer in one loop: for f in tr.local_layers: data = src.read_rows(...) dst.write_rows(...) and the comment above it justified the order with 'source and destination are different pools'. That premise is exactly what sharing removes. With one arena per rank sized max(PP, TP) instead of PP + TP, a destination write can land on a source row the transition still owes -- the sgl-project#297 reads-before-writes hazard, unobservable while the pools are disjoint, which is why it survived until sharing was attempted. The peer legs already materialise their payloads before the write region; this makes the local leg agree, at the cost of one list of tensors that were being materialised one at a time anyway. The falsifier builds each rank's PP and TP buffers as overlapping views into ONE arena at different strides, so a premature write corrupts rows, and checks byte-identity against the same flip run with disjoint pools: sharing must change memory economics, never an output byte. Both cases FAIL on the pre-fix loop and pass with the hoist (verified by reverting). GdnFlipMover already satisfies this invariant -- it packs all outgoing payloads, then the local one, then writes -- so the mamba half needs no equivalent change. Tests: scripts/run_631_flip_family.sh 270/270 (+2 aliased-arena falsifiers, proven red on the pre-fix code).

…e read/write seam Groundwork for making the FULL KV budget available to whichever phase is active, instead of splitting it between two permanently resident pools. MHATokenToKVPool gains swappable_backing: allocate on a VA reservation (the existing KvVmmBufferOwner) and back the whole span immediately, then release_backing()/restore_backing() unmap and remap the physical pages behind unchanged addresses. Sizing is untouched -- self.size stays what the constructor computed, already DCP-translated by _dcp_token_sharded_pool_rows, so the sgl-project#592 resize-translation gap on finalize_backing does not apply: the span released and restored is always the pool's own row count. That is deliberately NOT post-capture sizing, which is gated off for this config anyway (dcp_size == 1 and enabled prefill graphs, neither true here). PhaseFlipRuntime gains pre_write_fns, fired at the read/write seam. That seam is the only instant where backing may move: every byte the transition owes has been read (peer legs materialised, local leg hoisted in 93ee1e2) and the destination has not been touched. Empty by default, so this commit changes no behaviour. Tests: scripts/run_631_flip_family.sh 271/271 (+1 seam-ordering pin, per rank because the ranks run concurrently and a global event order would interleave and prove nothing). Not yet wired: the pools are not constructed with swappable_backing and nothing registers a swap -- arming it needs the tree-cache/allocator binding question settled first (they are bound to the PP stack's pools for process life), which is the next step.

…ng, corridor evidence, arena status The 496671-token boot is the finding: matching the KV token vector to each rank's measured capacity equalises the allocator's per-rank units and reaches the shared-arena design number -- the allocator was never the limit, the corridor is. It OOMed only because both phases' pools are resident, so the arena's payoff is now measured rather than modelled. Also records that the corridor must be verified against the 32768-token prefill rung (~600 MiB/card transient) and not decode (~70 MiB): a decode-only sample passed at 1986/2197/1672 while the same config breached at 716/1097/856 under prefill.

…ator's slot-id space The scheduler keeps ONE allocator for process life -- the PP stack's, built by build_kv_cache before the TP stack exists and never swapped at cutover. That is required, not an oversight: the flip identifies a KV row by its GLOBAL slot id across both layouts, so one id space is what makes the transition expressible. The consequence was an unchecked invariant. The TP stack derives its capacity from its own budget and token vector, so it can come out SMALLER than the id space the allocator issues from -- and then the first decode touching a high id writes past the end of the TP KV pool. It does not fail gracefully. It lands in store_kvcache's own bounds guard, SGL_DEVICE_ASSERT(index >= 0 && index < size_limit), as a device-side assert that takes down all three ranks with an async CUDA error pointing at whatever host call synchronised next. That is the crash this rig hit mid-benchmark at PP/allocator C = 46422 against TP C = 27200; it read as 'pool exhaustion' and was neither. Checked at boot where both numbers are on hand, so the refusal names them and says which knob raises the TP side. Both capacities are already min-reduced over the world group by _apply_token_constraints, so the comparison is deterministic per rank and aborts the boot identically everywhere instead of half of it. The shipped configuration satisfies it with room: id space 278104, TP capacity 340474. Tests: scripts/run_631_flip_family.sh 274/274 (+3).

…ving capacity Retracts this file's claim that raising the TP pool 318176 -> 340474 was a +7.0 % capacity gain. It was not a gain at all: the scheduler's single allocator is the PP stack's and is never swapped at cutover, because the flip identifies rows by global slot id and needs one id space. Serving capacity is therefore the PP stack's 278104; the TP pool's 340474 is headroom above the id space. Also records the invariant that fell out of it (TP capacity >= id space), that violating it is what actually caused the store_kvcache device assert previously filed as pool exhaustion, and that the capacity lever is the PP pool -- which needs boot-time exclusive backing, not just the runtime flip-seam swap.

…holds pages The flip ran two KV layouts whose pools were BOTH resident for process life, so each could only be sized against half the per-rank budget. The boot peak (PP pool + TP pool) is what set --rank-gpu-memory-mib, and the corridor floor being a CONTINUOUS minimum meant that one peak capped the budget permanently. Now exactly one layout holds physical pages at a time, on a fixed VA reservation so no address a captured graph baked in ever moves: boot back PP -> RELEASE PP -> build TP pools + capture TP decode graphs -> release TP -> restore PP (the boot phase) flip at the read/write seam, where the source is drained and the destination untouched: release source, restore destination Source-then-destination at the seam is deliberate: the reverse holds both layouts for the width of the swap, and a few milliseconds still counts against a continuous corridor minimum. Releasing first is safe because every row the transition owes is already in the payloads. Measured on the rig at unchanged budgets: PP release hands back 4320 MiB (5090) / 2160 MiB (3080), TP release 5888 / 4480 MiB, and free memory per card goes from 0.28-1.28 GB to 4.44-6.13 GB. That headroom is the point -- it is what the budget can now be raised into. Two things this uncovered, both fixed here: * the TP stack's pool came up UNSWAPPABLE. derive_tp_stack_server_args deliberately clears enable_phase_flip on the TP copy (it describes a TP stack, it does not enable a nested flip), so keying on that flag caught only the PP side. Keyed on is_phase_flip_tp_stack as well. * flush_cache killed every rank with cudaErrorIllegalAddress in the TP phase. _flush_zero_kv_buffers zeroed the scheduler's pool, and the scheduler's pool is the PP stack's -- released while TP serves. It now zeroes every layout that currently HOLDS pages and skips the unbacked one, which also closes the converse gap: zeroing only what the allocator can reach would have left the active TP layout un-zeroed and quietly broken the bit-for-bit-after-flush property. That crash is the answer to 'does anything touch the PP pool during TP decode'. It was found by arming the release and letting it fail loudly, which is why the release was armed before the budget was raised. Validated on metal: full flip round trip clean; short-answer output byte-identical across it with cuda graph True (256-token outputs are NOT a valid instrument here -- three draws in one phase with no flip give three hashes, the documented upstream GDN prefill nondeterminism, so the identity check runs inside the reproducible regime). Tests: scripts/run_631_flip_family.sh 274/274.

…67704 (+32.2 %) Makes the measured configuration the script default and records the pass. The TP pool is capped with --max-total-tokens: it can only ever address ids the scheduler's (PP) allocator issues, so its uncapped 788026 against an id space of 367704 was pure hoarding. Capping moved TP-phase free memory from 663/2676/1117 to 3231/5528/3109 MiB, which is what let the PP budget -- and therefore serving capacity -- grow. Corridor acceptance on the binding 32768-token prefill rung, 100 ms sampling, two flips and the decode bench inside the window: 2198/4033/2292 MiB against a floor of 1024. No regression: prefill within noise, decode within CV, flips slightly faster. Records the honest ceiling too. It is no longer the corridor (1.2-3.0 GiB sit unused above the floor) but the per-rank physical-availability check, which sizes the PP stack against memory the TP stack has not yet claimed and, under exclusive backing, never holds at the same time.

…l acceptance recorded The expected next lever -- teaching _assert_budget_physically_available that the two phases never hold pages together -- does not apply: the PP pool is sized at scheduler.py:1232, before the flip's TP stack is built at :1249, so there is no TP pool resident to discount. The check was never summing the phases. What actually caps rank 0 at ~22.4 GiB is the sgl-project#652 residual: on the 5090 the check sees 8.68 GiB held + 13.26 GiB free of a 32.6 GiB card, ~9.4 GiB accounted as held outside a process that carries one rank. Recorded so the next attempt starts from the evidence. Also fixes the boot script's HICACHE default to 0 -- the flip refuses --enable-hierarchical-cache at argument time, so a default of 1 made the script unbootable without an explicit override. Final acceptance on production: id space 367704, corridor 2286/4001/2380 on the binding 32768-token prefill rung, prefill 4236.8/7245.5/6842.6, decode 76.67/101.99, flips 997/1246/1720 ms, probe clean.

…ls the group (c) as designed is not implementable and the attempt was fatal on metal. Bounding entry to the flip's consensus all_reduce and abandoning from inside fired as intended -- and killed the server: CollectiveTimeoutError: barlink collective 'kv_pressure_ladder.consensus' made no progress for 45s Fatal Python error: Aborted RuntimeError: gloo/transport/tcp/pair.cc:547 Connection closed by peer A rank that has ENTERED an all_reduce owes that all_reduce. The moment this rank walked away, its peers' gloo pairs saw a closed connection and every rank aborted. So a wedge inside the reduction cannot be broken from inside it: any bound must be applied BEFORE entry (do not enter unless the peers are known to be joining), or the reduction must become a non-blocking poll a rank re-enters -- a different design, not a timeout. Withdrawn rather than fixed, because this sits unconditionally in on_round and would therefore have made MANUAL flips fatal too, and the manual path is currently the only working one. A test pins the withdrawal by source inspection so it is not re-added as the "obvious missing safety net" -- it is the opposite. Secondary bug, recorded: the handler caught PeerLostError while the channel raises CollectiveTimeoutError, so the error escaped rather than being handled. Irrelevant now, but it is why the first symptom was a bare abort. DELIVERY-BEFORE-BLOCK is kept and stays green, but it does NOT fix the wedge either: with the arm deferred by one pass, a downstream stage must re-enter the chain recv to reach the pass where it acts, and that recv blocks because upstream is already inside the reduction. The property the manual flip relies on is arm-and-join IN THE SAME ITERATION (process_input_requests at the top, the flip hook at the end of the same microbatch pass); deferral breaks exactly that property. Tests: 316 passed in the sgl-project#631 flip family, CPU-only.

Replaces defer-by-one with the invariant that actually holds, in three parts: (i) a rank receiving or originating a flip arm arms AND reaches the consensus reduction in the SAME pass, never returning to the chain recv in between; (ii) before joining, a rank owing an arm forward has THAT SPECIFIC send committed -- a targeted commit of one work handle, never a blanket synchronous send; (iii) the last stage owes no forward and joins directly. WHY THE BLOCKING REDUCTION IS SAFE, by induction: when rank k enters the reduction, (ii) guarantees rank k+1 already HAS the arm in its recv buffer; rank k+1 processes it at the top of its current or next pass and, by (i)+(ii), arms, completes its own forward and joins within that pass. No rank inside the reduction is owed anything by a rank outside it except the join, which arrives by induction. Worst case is LATENCY (a peer mid-prefill-chunk finishes its pass first), never deadlock. Every clause is a measured failure, not a precaution: * omitting (ii) = variant A: an async forward is progressed by _pp_commit_comm_work at the TOP of the NEXT pass, which never comes once this rank arms and blocks. rank 0 in bounded_collective, ranks 1-2 in _pull_raw_reqs, 0 % GPU. * deferring the arm to satisfy (ii) breaks (i) and is a GUARANTEED miss: the downstream stage must re-enter the chain recv to reach the pass where it acts, and that recv blocks because upstream is already inside the reduction (boot 12: all three armed, cutovers=0, dead at 40 s). * satisfying (ii) with async_send=False = variant B: deadlocks against the HIDDEN-STATES exchange, since a peer need not be in the chain recv at all. Hence "targeted" -- commit the request-chain handle only, every other channel stays async. Applied to the MANUAL flip too. Strictly safer, and it removes manual's latent at-idle deadlock: manual has only ever been exercised UNDER TRAFFIC, where the loop keeps cycling and the commit happens by accident rather than by construction. The property to engineer was never "be like manual" but the stronger invariant manual satisfies only accidentally. Also aligns the exception contract: the channel raises CollectiveTimeoutError where only PeerLostError was caught, which is why the withdrawn (c) surfaced as a bare "Fatal Python error: Aborted" instead of a handled error. Tests: 317 passed in the sgl-project#631 flip family, CPU-only. Can-fail proven on both specimens: reintroducing defer-by-one turns the same-pass, manual and last-stage tests red; skipping the targeted commit turns the commit and manual tests red. Ordinary traffic is asserted to KEEP its uncommitted async forward, so the targeted commit cannot silently serialise the chain. Operational acceptance NOT yet claimed; the boot carrying this is up next.

THE DESIGN LAW, from seven measured corpses: NO RANK MAY BLOCK ON ANY CHANNEL WHILE A PEER MAY BE IN A DIFFERENT BLOCKING CHANNEL. A arm same-pass, async forward rank0 in reduction, peers in chain recv B arm same-pass, sync forward rank0 in send, peers in hidden-states recv B' arm same-pass, targeted commit IDENTICAL to B (boot 13) D defer arm by one pass peers never reach the acting pass (boot 12) -- message-free local decision rank0 alone in the reduction (boot 10) -- bounded chain recv breaks the 1:1 send/consume contract; the SENDERS block -- bounded join (abandon inside) FATAL; closed the gloo pairs, aborted every rank B' is recorded as a falsification of my own reasoning: the targeted-vs-blanket distinction has NO force. Both block rank 0 on the chain send while a peer may sit in the hidden-states channel, and the induction failed at its first premise -- the downstream stage was not "waiting for the arm" at all. THE FIX: make ENTRY to the blocking reduction conditional on knowing every peer is already at that entry. Then entering is safe by CONSTRUCTION, not by argument, because no participant is anywhere else. * managers/phase_flip_presence.py: epoch-stamped, monotone, pollable per-rank markers in /dev/shm (the sgl-project#615 build-window precedent for group-visible state on this single-node topology). Presence is file EXISTENCE, written by atomic rename, so a half-written marker cannot be misread. Flags are NEVER cleared -- a retraction mints a new epoch -- which is what makes a poll safe against a racing writer without any clear-coordination. * _await_group_presence: the armed poll loop. Per iteration it only announces, PUMPS its outstanding arm forward non-blockingly, polls peers' flags, and sleeps. Nothing blocks, so the design law holds by construction. * the entry gate: enter the reduction only on all-ready. * PRE-ENTRY bound (60 s): if the group never assembles, disarm loudly naming the missing ranks and return to cycling; the policy re-arms a new epoch. Legal here precisely because nothing was entered -- the contrast with the withdrawn (c) is the point, and (c)'s fatal pin stays so nobody moves a bound back inside. The blocking targeted commit is REMOVED from the forward path (corpse B'); the arm is delivered by the pump instead, which also fixes corpse A (an async send is otherwise progressed only by the commit at the top of the NEXT pass, which never arrives once this rank is armed and polling). (a) posted-and-polled all_reduce Work is REJECTED on the record: its load-bearing premise -- that such a Work progresses while merely polled, without explicit progress calls -- is an unverified transport assumption of exactly the kind that has already killed designs here. This fork's async SENDS demonstrably do not progress without an explicit commit (corpse A). Betting the gate on the reduction transport differing, unverified, is the same bet. Tests: 326 passed in the sgl-project#631 flip family, CPU-only. Can-fail proven: opening the gate with a peer missing turns three gate tests red; reintroducing a blocking commit in the armed path turns the boot-13 specimen red. Also pinned: pre-entry timeout disarms without raising and leaves the flip retryable; epochs are monotone and a retraction is a new epoch; sweep never drops the epoch in flight; a failing pump cannot break the gate; the manual path takes the same non-blocking route. Operational acceptance NOT yet claimed.

…esync) The presence gate WORKED: all three ranks announced for the epoch (/dev/shm markers e0.r0, e0.r1, e0.r2) and no rank was stuck in any channel. The deadlock class is gone. What killed boot 14 was an interaction between two bounds I left running concurrently. The park deadline (30 s) measures "armed but never reached a quiescent boundary". The presence bound (60 s) measures "the group never assembled". Both were counting from the arm, so a rank whose peers took longer than 30 s to arrive abandoned on the PARK deadline while those peers were still polling for presence. The ranks then disagreed around a gloo collective, which is fatal rather than merely wrong: FLIP ABANDONED: pp_to_tp was armed for 30.0s ... (all three ranks) RuntimeError: gloo/transport/tcp/pair.cc:547 Connection closed by peer Fatal Python error: Aborted Fix is semantic, not a tuning tweak: "armed but never quiescent" is only a meaningful question once the group is ASSEMBLED, so the park clock is re-based when the gate opens. The two bounds now run in SEQUENCE -- presence governs assembly, park governs quiescence -- instead of racing. Tests: 48 in the policy file, full flip family green. The boot-14 specimen is pinned: the park clock must not move before assembly and must be re-based at it, or it expires mid-quiescence and desyncs the group around the collective.

…ot-15) Boot 15's gate opened "after 0.00s" on STALE markers and rank 0 entered the reduction alone -- the gate causing the exact failure it exists to prevent. The instance tag was os.getpid()//100000, which COLLIDES across consecutive boots: 3163115 and 3180590 both give 31. Boot 15 therefore read boot 14's leftover /dev/shm markers as its own quorum, before its peers had armed. The tag has two load-bearing properties and the old one had neither reliably: * UNIQUE PER BOOT, or a later boot inherits an earlier quorum; * IDENTICAL ACROSS RANKS, because the flags are a rendezvous -- a per-process value gives every rank a different quorum and none ever assembles. So it comes from the environment, set ONCE by the boot script (SGLANG_PHASE_FLIP_INSTANCE, timestamp+pid) and inherited by every rank. The in-process fallback is deliberately not process-derived either: it uses the process-group leader's start time, which is shared across the boot's ranks and differs between boots. Plus a sweep of foreign instances at construction, before anything can poll: a stale marker that merely exists is harmless, one that is READ is a false gate. Tests: 329 passed in the sgl-project#631 flip family. The boot-15 specimen is pinned (an earlier boot's full quorum must not open a new boot's gate), as is the rendezvous property (all ranks of one boot agree on the tag).

…t-16) Boot 16 was the first run with a correct rendezvous: fresh per-boot markers, all three ranks announced, the gate opened. Then all three abandoned unanimously (non-fatally, unlike boot 14) with: FLIP ABANDONED: pp_to_tp was armed for 0.0s without the group reaching a quiescent boundary "0.0s" is the tell. The park expiry was computed at the TOP of on_round from the pre-rebase clock, so a rank that had waited out the presence poll was already flagged expired; the gate then opened, re-based _armed_at, and the very next check abandoned the flip it had just assembled. The park deadline asks "armed but never quiescent?" -- meaningful only once the group is ASSEMBLED. So the time spent waiting for assembly must not count toward it, which means the gate has to run BEFORE the expiry is computed, not after. With that order the two bounds are genuinely sequential: presence governs assembly, park governs quiescence, neither consuming the other's budget. Pinned by source inspection (the gate must appear before the expiry in on_round), because this is an ORDERING property that no value-based assertion would catch. Tests: 330 passed in the sgl-project#631 flip family, CPU-only.

…ound (boot-17) Boot 17 reached the healthiest state yet -- fresh rendezvous, all three ranks present, gate opening cleanly -- and still never flipped: PHASE-FLIP group present for epoch 0 after 0.00s (repeatedly, all ranks) arms=1 cutovers=0 abandoned=0, /generate answering nothing The gate opens EVERY round once the group is present, and I re-based the park clock on every opening. That makes the park deadline unreachable: a flip that can never reach quiescence holds for ever, with its requests parked, and the server answers nothing while looking healthy. The re-base is a per-ARM event, not a per-round one -- it exists so the time spent ASSEMBLING does not count toward the quiescence budget, which is a single transition, not a repeating condition. Guarded by _gate_open_epoch so it fires once per epoch and later rounds leave the clock alone; the park deadline can then expire and abandon loudly, which is the designed outcome for a flip that cannot reach quiescence. Tests: 331 passed in the sgl-project#631 flip family, CPU-only. The boot-17 specimen is pinned: the gate re-opening on a later round must NOT move the park clock.

…it rests on The decided fix for the boot-18 wedge was clauses (i)+(ii): announce presence only once a rank owes no chain send, and keep consuming the chain non-blockingly while armed so no peer can block on an armed rank. Both are implemented here. Clause (ii) does not work on this build, and the measurement that shows it also undermines the diagnosis both clauses were derived from, so the pair ships OFF. MEASURED (test_measured_gloo_does_not_progress_a_posted_irecv_by_polling, two processes, real gloo, 8 B and 512 KiB, boot-18 geometry): * a downstream with a posted irecv that only calls is_completed() NEVER completes it -- 4 s of polling at 10 ms, both sizes. A drain built on that predicate absorbs nothing, ever. The pre-existing send-side pump (pp_pump_send_req_work) reaps on the same predicate and is therefore also a no-op; * the upstream's commit returns in 0.00 s with the forward UNCONSUMED. So a single unconsumed forward does not block the upstream at all, and "the downstream stopped consuming" does not by itself explain rank 1 blocking at scheduler_pp_mixin :705 -> :1109. This is the premise phase_flip_presence already rejected for a posted-and-polled all_reduce -- "an unverified transport assumption of exactly the kind that has already killed designs here". It is false for point-to-point too, and is now pinned as an executable measurement so a torch/gloo change that alters it is noticed here rather than in a wedge. Wired live the pair would be strictly worse than the defect: with clause (i) active the forward is re-issued every pass, so send_req_work is never empty when the round hook runs, presence is withheld for ever, and EVERY flip abandons at the presence deadline -- a server that silently stops flipping. Everything is therefore gated behind SGLANG_PP_CHAIN_RECEIVER=1 (default off) and the default path is byte-for-byte unchanged. Boot 18's evidence was also re-examined rather than taken on trust. The py-spy is gone and the serving log was truncated by the next boot; rank 0 (in the reduction) and rank 1 (at :705 -> :1109) are faithful to tree cf478d1, but rank 2's stack was never recorded and rank 2 is the LAST PP stage, for which :705 is structurally unreachable. The corpse table now separates what was observed from what is inferred, and the boot script rotates the serving log instead of truncating it, which is what made the question unanswerable. Kept and tested rather than deleted: PpChainReceiver is a correct state machine for a two-step wire format that cannot be consumed by halves, and is the piece any future design needs whatever drives progress. Tests: scripts/run_631_flip_family.sh 337 passed (was 331; +6 new pins), test_pp_chain_receiver.py 7 passed. Can-fail proven by mutation for all three new clause pins: dropping clause (i) reddens test_can_fail_presence_is_withheld_while_this_rank_owes_a_send and test_can_fail_a_rank_that_never_flushes_abandons_instead_of_wedging; dropping clause (ii) reddens test_can_fail_armed_round_consumes_the_chain_before_it_announces; restoring the blocking commit in the armed path reddens test_armed_forward_path_pumps_and_issues_no_new_forward. A pre-entry deadline bug found by its own test while building it (withholding presence skipped the bound) is fixed. Production on 30030 unaffected and healthy throughout (POLICY=manual, health_generate 200).

…nd always was Standalone finding, separated from the boot-18 work because it is older than that wedge and changes how several earlier results should be read. pp_pump_send_req_work reaps an outstanding chain send when the work handle reports is_completed(). Measured 2026-08-08 on a real gloo pair: that predicate NEVER fires for an isend on this build -- and the control that makes this conclusive is that it does not fire EVEN AFTER THE PEER HAS FULLY CONSUMED THE MESSAGE. The pump has therefore never cleared send_req_work in any circumstance. The only thing that has ever reaped a chain send here is the blocking _pp_commit_comm_work, whose work.wait() clears the list. What this retro-explains: arms have reached downstream stages via those stages' OWN blocking chain recv all along -- the recv side's wait() is what progresses the transfer -- never because an armed rank "pumped the arm forward while it waited". Design notes crediting the pump with delivery were reasoning about a no-op. No one-line repair exists: only wait() progresses a send here, and blocking is precisely what the armed path may not do. The function is kept rather than deleted -- it cannot mutate state, it is where a working predicate would go if the transport gained one, and deleting it would erase the record of what was tried -- but it is now documented as dead and nothing may rely on it. The transport premise is now falsified from three directions and the corpse table says so in one place: a posted all_reduce (rejected on argument), a posted irecv (measured), and a posted isend (measured). On this build ONLY wait() progresses a transfer. Also adds scripts/route_a_631_wedge_capture.sh, the evidence harness for the boot-18 reproduction. It captures ALL THREE ranks to disk -- Python stacks, --locals (send_req_work depth, flip pending state) and --native (what :1109 waits on below Python) -- plus the /dev/shm presence markers, process table and a bounded log slice, on an automatic trigger (presence abandonment in the log, or repeated /health_generate failure), never on manual timing. Boot 18's rank-2 stack was lost because capture was manual and the log was truncated; that is the failure this removes. Tests: test_pp_chain_receiver.py 8 passed (+1, the pump pin), flip family 337 passed. Capture harness dry-run against the live healthy server on 30030: 26 files, all three ranks present with non-empty python/native/ locals stacks and no capture errors. Production untouched by this commit.

…is per-round, its evidence is per-epoch Reproduction on metal, POLICY=auto from 526e53c, wedged 2026-08-08 23:12:38Z. Evidence captured automatically and kept: /spinning/evidence-631/wedge_20260808T231450Z_INSIDE_REDUCTION (all three ranks, python + native + locals, presence markers, log slice). rank 0 bounded_collective -> _reduce -> on_round (in the reduction) rank 2 bounded_collective -> _reduce -> on_round (in the reduction) rank 1 _pp_commit_comm_work <- _pp_forward_and_process_input_requests, blocked in work.wait() on ONE P2PWork Boot 18's geometry is confirmed and rank 2 -- the stack that was lost and that every later design decision had to guess at -- is in the reduction beside rank 0, as inferred. THE CAUSE IS NOT WHAT WAS INFERRED. The log shows all three ranks announcing and the gate opening on all three ("group present for epoch 0 after 0.00s"), so rank 1 passed the gate and completed a consensus round. The deadlock is in the NEXT round: rank 1 must traverse its top-of-pass commit to get back to the hook, that commit blocks because rank 2 has posted no matching irecv (rank 2 is already inside round N+1's reduction), and ranks 0 and 2 entered round N+1 by re-opening the gate INSTANTLY on the epoch-0 flags still up from round N. Flags are never cleared, so the quorum that authorised round N silently authorises N+1 and every round after it. THE GATE'S GUARANTEE IS PER-ROUND; ITS EVIDENCE IS PER-EPOCH. The monotone property that makes a poll safe against a racing writer is the same property that makes the gate a rubber stamp after its first use. Neither deadline intervenes: both are evaluated before entry, and the gate opens in 0.00 s on stale evidence. This also settles the earlier probe that appeared to contradict the diagnosis. An unconsumed forward returns in 0.00 s when the receiver has POSTED an irecv and merely not completed it; when the receiver has posted no irecv at all, the sender's wait() blocks. Both facts are needed to describe the wire and the corpse table now carries both. Consequence for the decided fix, recorded so it is not retried: clauses (i)+(ii) would not have prevented this, for a reason independent of the transport falsification. Clause (i) gates the FIRST announce, and rank 1's flag was already up from round N. Clause (ii) cannot help because the rank that must keep consuming is inside the blocking reduction, not in the poll loop where a drain would run. The next design has to address the INTER-ROUND INTERVAL, not the entry. No behaviour change in this commit: corpse-table finding only. The clause (i)/(ii) build stays parked behind SGLANG_PP_CHAIN_RECEIVER (default off) pending that design decision. Tests: test_phase_policy.py + test_pp_chain_receiver.py 66 passed. Production restored to POLICY=manual after the evidence was on disk.

…atches its guarantee THE RULE: evidence must have the same scope as the guarantee it licenses. The gate guarantees "every participant is at THIS reduction's entry" -- per round. The flags were stamped per EPOCH and are never cleared, so round N's quorum was a standing authorisation for N+1 and every round after it. Reproduced on metal with all three stacks (evidence-631/wedge_20260808T231450Z_INSIDE_REDUCTION): ranks 0 and 2 inside round N+1's reduction, rank 1 blocked at its top-of-pass commit between rounds, gate re-opened in 0.00 s on stale flags. Markers are now stamped (epoch, round). The round is the count of consensus reductions this arm has COMPLETED, and it is incremented in exactly one place: immediately after the collective returns. That is the one instant at which the ranks provably agree, which is what lets the count serve as a shared stamp without ever being exchanged. It is deliberately NOT any rank's local loop counter -- those diverge in absolute value under event_loop_pp, which is the very divergence the gate exists to tolerate, so stamping on them would be circular. THE INDUCTION NOW CLOSES, and round-scoping is what closes it. Announcing for round R is only REACHABLE after this rank completed its round-R top-of-pass commit, because that commit sits on the path to the hook. So a round-stamped flag means "my chain is settled for R": rank k announced R => k's send to k+1 is settled for R all ranks announced R => the chain 0->1->2 is settled for R => no rank inside round R's reduction is owed a chain operation => the blocking reduction is safe, per round, by construction This is the intent of the withdrawn "announce only once you owe no send" clause, obtained at the right granularity and with NO new machinery -- no drain, no progress engine, no predicate this transport cannot honour (corpse F). The remaining failure mode, a rank that never reaches round R because it is busy, is converted by the per-round pre-entry bound into a loud unanimous abandonment and a later retry. Monotonicity is preserved where it mattered: within an (epoch, round) stamp a flag is set once and never unset, and retraction still mints a new epoch rather than clearing. Round-scoping removes only the over-reach -- the flag no longer answers a question it was never evidence for. The consumer-thread alternative stays priced-but-unbuilt, and the parked SGLANG_PP_CHAIN_RECEIVER build stays parked (default off). Tests: flip family 340 passed (+3), test_pp_chain_receiver 8 passed. Can-fail proven by mutation: reverting the gate to epoch-scoped reads reddens test_can_fail_a_completed_round_does_not_open_the_next_one with the reproduction's own message. Also pinned: per-round pre-entry budget, and monotonicity within a stamp. Metal verification against the at-boot zero-traffic wedge follows in the next commit.

…: the gap is between announce and entry Metal verification of b51480f under POLICY=auto, plus the correction of two of my own wrong readings. Evidence: evidence-631/wedge_20260808T233910Z_KVPRESSURE_DIVERGENCE (all three stacks), markers e0.n0.r0/r1/r2. WHAT ROUND-SCOPING FIXED, confirmed: the rubber stamp is gone. On the epoch-scoped build ranks re-opened the gate in 0.00 s on a previous round's flags; on this build the markers show one round only and no rank re-enters a reduction on stale evidence. WHAT IT DID NOT FIX: the flip still wedges, inside a SINGLE round. rank 2 in the reduction for round 0 rank 1 blocked at its top-of-pass commit (:724 -> :1187) rank 0 blocked at its top-of-pass commit (:724 -> :1187) Announcing and ENTERING are not the same instant, and a whole pass can sit between them. A rank announces at the hook; if the quorum is not yet complete it returns, goes around the loop, and the next thing it meets is the top-of-pass commit. So the LAST rank to announce enters immediately while every earlier announcer must traverse a blocking chain commit to get back to the entry -- and that commit blocks precisely because the rank that already entered has stopped consuming. rank 2 in the reduction -> does not recv -> rank 1 stuck committing to rank 2 -> does not recv -> rank 0 stuck committing to rank 1. The flag means "I was at the entry once", not "I am at the entry". The evidence is now correctly scoped in TIME and still wrong in PLACE. TWO CORRECTIONS TO MY OWN EARLIER READINGS, recorded so they are not inherited: 1. I reported this as a separate subsystem diverging ('kv_pressure_ladder.consensus' timing out) and hypothesised a latent local-round-cadence bug in the KV pressure ladder. WRONG. The flip's consensus channel is built from kv_pressure_runtime.default_collective_min, and that helper hardcodes its own module's label, so the flip's OWN reduction reports under another feature's name. There is no second bug. The misleading label is now called out in the corpse table. 2. The earlier probe conclusion that "an unconsumed forward does not block the upstream" held only for a receiver that had POSTED an irecv. These stacks show the blocking case is real when the peer has posted none -- which is what a rank sitting inside a reduction looks like. Round-scoping is kept: it is correct, it is independently pinned, and it removes a real hole. It is simply not the whole fix. Any next step must remove the blocking chain operation from the announce-to-entry interval rather than tighten the stamp again. Tests: flip family 340 passed. No behaviour change in this commit beyond the corpse-table record. Production restored to POLICY=manual.

The flag's meaning changes from 'I was at the entry once' to 'I am at the entry, quiescent, and owe nothing'. Evidence is now correct in TIME (round-scoped, unchanged) and in PLACE. Two clauses compose: round-scoping stops stale evidence ACROSS rounds; quiescent-announce stops it WITHIN a round. (1) A rank announces only once ready_fn holds -- drained microbatches, no admissions, no owed payload -- evaluated at the hook. A non-quiescent rank returns to the pass loop without announcing, which is how it drains. An EXPIRED rank is exempt so it can still reach the reduction and make the abandonment group-agreed. (2) Once announced the rank SPINS at the hook and does not return to the pass loop. That interval is what killed the previous build: the rank met its top-of-pass commit before it could re-check the gate, and that commit blocked behind whichever rank had already entered. The spin touches no channel -- flags and a sleep -- and keeps the existing per-round pre-entry bound, so it is abandonable and cannot hang. Safety, as the completed argument: a quiescent rank owes neither hidden states nor chain payload, so the only message it stops producing is the empty keep-alive forward. Peers are either quiescent and spinning too (need nothing), or mid-drain -- and that case is a BOUNDED RETRY, not a wedge: the spinners' bound expires, they abandon loudly, return to the loop, resume forwarding, the straggler drains, a later epoch retries with everyone genuinely quiescent. At true idle all ranks are quiescent at once and the gate opens on live evidence. Scoped to a wired presence channel: with no presence there is no announce, so the plain consensus path keeps its uniform-hold behaviour exactly. Boot-16's pin is re-expressed as BEHAVIOUR rather than source order (the order legitimately inverts here). Its invariant now holds by construction: a rank is either draining on the park clock or spinning on the per-round presence clock, never both. Also fixes the mislabeled collective that sent a live wedge into the wrong subsystem: default_collective_min took its label from its own module, so the flip's reduction timed out as 'kv_pressure_ladder.consensus'. The label is now a parameter and the flip passes 'phase_flip.consensus'. Tests: flip family 344 passed. Can-fail proven by mutation: announcing without quiescence reddens test_can_fail_a_non_quiescent_rank_does_not_announce. Also pinned: the spin does not return to its caller between polls, and mid-drain skew abandons within the bound leaving the rank disarmed and serving.

…, two new defects exposed Boot POLICY=auto from d547568, 2026-08-09 00:06-00:09Z. The wedge class is GONE and the bounded-retry argument held exactly as written: rank 0 armed, drained first, announced, spun; ranks 1 and 2 never reached the entry; rank 0's per-round bound expired at 60.0s; it abandoned LOUDLY with NOTHING entered, returned to the loop, and ranks 1 and 2 then reached the entry and announced. No rank blocked behind another at a top-of-pass commit. Every step of case (b) behaved as designed. Two new defects underneath it, neither a variant of anything in the table: G. SPINNING STARVES THE DOWNSTREAM. A spinning rank stops issuing the per-pass chain forward, and the downstream stages reach the hook ONLY by returning from their blocking chain recv -- which that forward is what satisfies. The first rank to become quiescent (rank 0, the intake rank, always) therefore prevents every rank behind it from becoming ready. The retry is bounded but NOT convergent: the same rank always drains first, so the starvation reproduces identically each epoch. The safety argument was right that this is not a wedge; it did not predict that the condition recurs. A quiescent spinner evidently must keep EMITTING the keep-alive forward (an isend needs no peer and never blocks) while still not CONSUMING -- a design decision that owes an answer for the sends that then accumulate unconsumed. H. A PRE-ENTRY ABANDONMENT LEAVES A LIVE FLAG. _abandon_no_quorum is rank-local by design and mints a new epoch, but it cannot retract the marker it already wrote. Rank 0 abandoned epoch 0 and re-armed at epoch 1 while ranks 1 and 2, arriving moments later, formed a full epoch-0 quorum USING RANK 0'S STALE FLAG and entered epoch 0's reduction without it. Epochs diverged; the group died on the collective timeout at 00:09:39. 'Retraction mints a new epoch' protects the retracting rank and does nothing for peers still reading the old one. Evidence that a rank has LEFT is as load-bearing as evidence that it arrived, and only one of the two currently exists. Recorded, not fixed: G is a design decision of the class that goes back to main, and H's fix (a publishable withdrawal) touches the monotonicity rule that three earlier corpses depend on. No behaviour change in this commit. Tests: flip family 344 passed. Production restored to POLICY=manual.

…er, with two-phase entry Corpse H, measured 00:07:34Z: a pre-entry abandonment is rank-local and mints a new epoch, but cannot retract the presence marker it already wrote. Rank 0 abandoned epoch 0 and re-armed at epoch 1 while ranks 1 and 2 formed a full epoch-0 quorum ON RANK 0'S STALE FLAG and entered a reduction it would never join; the group died on the collective timeout at 00:09:39Z. Evidence that a rank has LEFT is exactly as load-bearing as evidence that it arrived, and only one of the two existed. MONOTONICITY SURVIVES PER MARKER. presence(e,r,rank) is still write-once and is never mutated or cleared. WITHDRAWN(e,r,rank) and ENTERING(e,r,rank) are their own write-once markers, each with a single writer -- its own rank. Nothing here reintroduces the ordering problem that made flags monotone in the first place, which matters because three earlier corpses depend on that property. Quorum is now: all present AND none withdrawn. THE INVARIANT that makes the race sound: A WITHDRAWAL IS ONLY EFFECTIVE IF NOBODY COMMITTED ON IT. Any commit converts every committed-or-withdrawing rank into an enterer. A rank carrying both markers is an ENTERER -- it discovered a peer had already committed on its presence and follows through. Resolved in one place (PhaseFlipPresence.withdrawn) so no caller has to reason about it. Two-phase entry: publish ENTERING, then re-check quorum; if a WITHDRAWN appeared, wait -- and that wait terminates BY CONSTRUCTION, because this rank's ENTERING forces the withdrawer to follow through, after which it stops counting as withdrawn. Withdrawal: only while no peer is ENTERING; re-check after publishing, because a peer may commit in the window between check and write. If one did, the rank enters instead of stranding it. Tests: flip family 344 passed, phase policy 70 passed (+5). Can-fail proven by mutation on both halves: a quorum that ignores withdrawals reddens the stale-flag specimen; dropping the both-markers tie-break reddens the forced-entry pin. Also pinned: withdrawal is round-scoped, a rank's own ENTERING never blocks its own withdrawal, and the entering side waits out a withdrawal then enters. G (the armed service loop) is NOT in this commit.

…ed G design, and the acceptance program Compact successor brief in docs/dev/631/HANDOFF_656.md. Points at the corpse table in phase_flip_presence as the design document, and carries: - current HEAD and what each of the last eight commits contributes; - THE MEASURED TRANSPORT FACTS as a numbered list, because every design that ignored one of them has died: is_completed() never fires for a posted irecv OR an isend (the latter not even after the peer has fully consumed), only wait() progresses a transfer, the two-sided posted-vs-not-posted wire fact, and the one positive behaviour to build on -- the recv side's wait() drives the transfer; - defect G and the approved ARMED SERVICE LOOP with the send-counter, including the load-bearing ordering constraint (publish the counter strictly AFTER the isend is posted, so the only skew is counter-lags-send: a message may be consumed late, never phantom received), the both-inbound-channels requirement, why this is not the bounded-recv corpse, and the channels-empty-at-entry assert; - the acceptance program verbatim, with the honesty bar spelled out: a pass is a completed PP->TP->PP cycle under load in one unmanned log; - pointers to the three three-rank specimens in /spinning/evidence-631/ and to the capture harness; - the traps that already cost time, including the mislabeled collective; - the standing warnings: POLICY=auto must not boot until G lands, production stays manual and serving, watchdog stays decommissioned, port 30099 untouched. Documentation only; no behaviour change.

…d-counter Defect G, measured 2026-08-09 00:06-00:09Z: a rank that became quiescent and spun at the flip's entry gate stopped issuing its per-pass chain forward. Its downstream reached the hook ONLY by returning from the blocking chain recv that this very forward satisfied, so the first rank to quiesce -- rank 0, the intake rank, always -- prevented every rank behind it from becoming ready. Bounded by the pre-entry deadline, but NOT convergent: the same rank drained first every epoch, so the starvation reproduced identically, epoch after epoch. The obvious fix is wrong on its own terms. Keeping the spinner EMITTING a keep-alive forward owes an answer for the sends that then accumulate unconsumed, which converts a starvation into the bounded-chain-recv corpse. The fix is the other side: stop the downstream from NEEDING the forward. While armed, a rank services its channels each turn and reaches the hook by its own poll, so no rank's readiness depends on a peer's traffic. Consuming without is_completed(), which is the load-bearing part. That predicate never fires on this transport in either direction (corpse F), so readiness cannot come from the transport at all. It comes from a pollable side channel instead: phase_flip_counters, monotone per-message counts on /dev/shm under the presence gate's directory and instance tag. Each sender publishes STRICTLY AFTER posting its isend, so the only skew a peer can observe is counter-lags-send -- a real message seen one poll late. A receiver makes the BLOCKING recv only once the count exceeds its own consumed count, so the message provably exists and the block is bounded by transfer time rather than by peer scheduling. That is deliberate use of the one transport behaviour with positive evidence: the recv side's wait() drives the transfer (arms propagated by exactly this route across boots 14-18). Publishing before the post would invert the skew and send a peer into an unbounded block for a message nobody posted. The ordering is the design, and is pinned by a mutation test. Changes: - phase_flip_counters.py: the counter channel. One file per (channel, kind, rank), single writer, atomic rename, monotone. Two channels: the request chain, and the tensor-dict wire that proxy and output SHARE and must therefore count together. - pp_chain_receiver: consume_up_to(sent_count) -- the greedy counted consume that poll() could never be -- plus a consumed-count callback. The parked receiver is now ON by default under the flip; poll() is kept only as the pinned record of corpse F. - scheduler_pp_mixin: counter bumps at both real send sites and at the dict recv; pp_flip_service (consume, then flush), the counter-gated flush that replaces the dead send-side pump, and pp_flip_channels_empty. - The armed rules are gated on the FLIP, not on the chain receiver. That condition was also a bug: the receiver exists only on ranks with an upstream, so the armed intake rule was off on rank 0 -- the rank that must stop admitting work for the group to reach a quiescent boundary at all, and the rank whose starvation defined G. - phase_flip_runtime: service_fn and channels_empty_fn replace pump_fn and drain_fn in the wiring. Flip-commit hygiene: a rank withholds presence while a channel is not empty (bounded by the pre-entry deadline, so it stays convergent), and re-checks at the instant of entry, abandoning PRE-ENTRY rather than letting a message cross the re-formation and misframe the post-flip stream. Test results (CPU desk run, PYTHONPATH=worktree/python, CUDA_VISIBLE_DEVICES=99): - bash scripts/run_631_flip_family.sh -> 377 passed, 91.75s. Baseline before this change on the same tree: 349 passed. The family script now also carries test_pp_chain_receiver.py (8, previously run separately) and the new test_phase_flip_counters.py (15). - New can-fail pins, each shown to FAIL under its mutation and to terminate rather than hang: * test_can_fail_publishing_before_the_post_wedges_the_receiver * test_can_fail_flushing_without_the_counter_gate_blocks_on_the_peer * test_can_fail_a_non_empty_channel_at_entry_abandons_before_entering - ruff: no new findings on any touched file (the one F841 in scheduler_pp_mixin is present at HEAD unchanged); new files clean. Metal validation of G follows in the next commit -- this one is the build plus its desk evidence.

Boot POLICY=auto 2026-08-09 01:11-01:16Z, three-rank specimen persisted at /spinning/evidence-631/wedge_20260809T011602Z_G_FIXED_PROXY_CHANNEL_EXPOSED. G IS FIXED, MEASURED. All three ranks reach the flip's entry with no traffic driving them there: "group present for epoch 0 after 0.00s/0.01s" on ranks 0, 1 and 2, all three ENTERING markers on disk, and rank 0's stack INSIDE the consensus reduction. Every predecessor boot had at least one rank blocked upstream of the gate; that starvation is gone, and for the predicted reason -- no rank's readiness depends on a peer's traffic. The first attempt failed on a one-line wiring bug whose failure MODE is the real lesson, so both are fixed here: - The consumed-counter callback raised NameError on its first call (a missing import in Scheduler._build_pp_chain_receiver, not in the module under test). It was caught as best-effort and logged, so the whole unit suite passed while the live system published no consumed count at all: req.c1 and req.c2 never appeared on /dev/shm, every upstream withheld presence for ever, and all three epochs abandoned at the 60 s deadline. Pinned now by test_the_receiver_wiring_actually_publishes_the_consumed_ count, which drives the REAL factory's callback through the real state machine rather than constructing one. Verified can-fail: with the import removed it reproduces the exact NameError. - The symptom pointed at the wrong place. The abandonment said ranks 0 and 1 "never reached the flip entry" when they had reached it and declined to announce. A withholding rank is otherwise invisible in the log, so it now says so, with the reason, throttled to a quarter of the presence deadline. - The receiver's publish-failure log is throttled and counted, because a permanent failure previously buried itself at one line per message. WHAT G UNCOVERED (defect I, in the corpse table): quiescence is rank-local, the obligation is pairwise. Three stacks: rank 0 in the reduction, rank 1 spinning at the gate withholding (its downstream is not consuming), rank 2 blocked in _pp_recv_proxy_tensors -- the HIDDEN-STATES channel -- waiting for hidden states from rank 1, which had meanwhile declared itself quiescent and gone to the gate. _pp_microbatches_drained cannot see that a downstream is committed to receiving from this rank. The arm that would have armed rank 2 sat behind that same wedge (req.s1=4441 against req.c2=4440, exactly one unconsumed message), so rank 2 stayed UNARMED while ranks 0 and 1 re-armed and the epochs diverged. This is the SENDER side of the channel the service loop deliberately does not consume. The argument for not consuming it -- a rank with an inbound dict message is by definition not quiescent -- is sound for the receiver and says nothing about a peer that stops sending. Nothing wedged inside a collective, nothing was aborted, no request was touched, and the server kept answering throughout. The bounded pre-entry machinery held. Test results: - bash scripts/run_631_flip_family.sh -> 378 passed, 92.94s (was 377 before the wiring pin; 349 before this defect-G work began). - Metal: POLICY=auto boot reaches the entry on all three ranks (above). The acceptance program is NOT yet runnable -- defect I blocks it. - Production restored on POLICY=manual from this tree and verified generating: GET /health_generate 200, /v1/chat/completions 200.

…lf-contradictory Boot POLICY=auto 2026-08-09 01:37:12Z. The gate assembled 0.01 s after arming and the flip went through on all three ranks: PHASE-FLIP cutover complete: active stack tp, ps tp=3 pp=1 PHASE-FLIP DONE pp_to_tp (epoch 1) in 1038.8 / 1265.0 / 1762.0 ms: 80 live slots, sent 368 cells / 0.72 MiB, received 272 / 0.53 MiB PHASE-FLIP event loop re-dispatch after pp_to_tp (active stack now tp) This is the first policy-driven flip that has ever committed. Log preserved at /spinning/evidence-631/FIRST_AUTO_FLIP_COMMIT_20260809T0137Z.log. WHAT UNBLOCKED IT was not another gate fix. build_flip_quiescence_fn called Scheduler._pp_microbatches_drained -- the FULLY-IDLE predicate, which also requires every running_mbs slot to be empty. running_mbs is the RESIDENT DECODE SET and empties only when requests FINISH, while the policy arms pp_to_tp precisely BECAUSE requests are decoding. The arming condition and the quiescence condition could never hold at the same time, so the group assembled, entered the reduction, agreed and abandoned at the park deadline with ready=0 on every rank, indefinitely. It contradicted the function's own docstring ("does NOT require ... an empty running batch") and the rest of the design: build_flip_live_slots_fn exists to move the KV rows of requests resident at the flip. What must be quiet is the PIPELINE -- no forward in flight, no half-written chunk -- which is what mbs answers. Quiescence now reads mbs only. Diagnosis came from the per-rank quiescence reason added here: "NOT QUIESCENT: PP microbatches not drained (live mb slots [], running_mbs slots [0])" -- nothing in flight, the decode set alone holding the flip. Previously the log carried only "ready=0" and the question had to be answered from py-spy stacks. NEXT DEFECT, exposed by the commit and recorded as J in the corpse table: post-flip pool accounting for the carried decode set. One pass after the cutover, on_idle's invariant checker raised pool memory leak detected! [full] total=367704, available=367623, protected=80, leaked_full_pages={81}, leaked_mamba_pages={2} one full page and one mamba page unaccounted after 80 live slots crossed. The KV move itself reported balanced cells; the allocator-side bookkeeping for the carried rows in the destination stack does not survive. All three ranks raised it and the group went down on SIGQUIT. No previous boot ever got far enough to allocate against the TP stack with a resident decode set, so this could not have been seen before. Test results: - bash scripts/run_631_flip_family.sh -> 379 passed. The quiescence pin that encoded the old predicate is corrected, and test_a_resident_decode_set_alone_does_not_block_the_flip now pins the contract that was measured wrong. - Metal: automatic PP->TP cutover committed on all three ranks (above). - Production restored on POLICY=manual and verified generating: GET /health_generate 200.

The leak the first committed flip exposed is off by exactly one page, and the corpse table now carries that arithmetic rather than only the message: available 367623 + protected 80 = 367703 against total 367704, with the moved set (80 live slots) and the protected set agreeing with each other and disagreeing with the pool by one row. Stated hypothesis, marked as one: build_flip_live_slots_fn enumerates req_to_token[idx, :seqlen], and during decode the row already reserved for the token being generated sits at index seqlen, outside that slice -- so it is allocated, never enumerated, never moved, and owned by nothing in the destination stack. To be verified against the real allocated extent before any change. No behaviour change; documentation only. Suite unchanged at 379 passed.

…sified VERIFY BEFORE FIXING, applied. The hypothesis recorded last commit was that build_flip_live_slots_fn enumerates req_to_token[idx, :seqlen] while the allocator owns kv_allocated_len -- structurally different under sgl-project#486, whose spec reserve is W + L slots ahead of kv_committed_len (W = the draft/verify write footprint, several slots on this rig's NEXTN config, not one). It fits the symptom exactly. It is wrong. Measured with a census bracket around the flip, identical on all three ranks (2026-08-09 02:08:17Z): POOL CENSUS at-arm live_reqs=1 cached=80 unaccounted=1 [81] POOL CENSUS pre-cutover live_reqs=0 cached=80 unaccounted=1 [81] POOL CENSUS post-cutover live_reqs=0 cached=80 unaccounted=1 [81] Page 81 is already unaccounted before the flip moves a byte, and the set is unchanged across the move and the cutover. A no-flip control boot (POLICY=manual, one request served to completion, server idle) stayed clean: leaks=0, flips=0, health 200. WHAT THIS LOCALISES. At arm the row is legitimately held -- live_reqs=1, charged by the checker as uncached = kv_allocated_len - cache_protected_len, so the invariant balances. By pre-cutover the request has FINISHED, its 80 committed rows are in the tree, and its one uncached row was never freed: nothing owns page 81 and no live request remains to charge it to. The defect is in the COMPLETION PATH OF A REQUEST THAT FINISHES WHILE A FLIP IS ARMED, not in the flip's KV move. First suspects are what the armed state defers or suppresses around completion (the abort-deferral window), and the checker's own invariant assert req.kv_committed_freed == req.kv_overallocated_freed. Changes, all diagnostics; no behaviour change to what the flip moves: - _pool_census(when, direction), reproducing the invariant checker's expected - free - cached arithmetic, called at arm and on both sides of the cutover. Read-only and exception-swallowing: a census must never be able to affect the flip it watches. - _probe_allocated_extent, reporting seqlen vs kv_allocated_len vs cache_protected_len per live request, so the sgl-project#486 reserve is visible rather than inferred. - _live_reqs (running_batch + last_batch, the pair the checker walks), used by diagnostics ONLY. - build_flip_live_slots_fn is UNCHANGED, deliberately. I had widened it while adding the probe; the census falsified the reason, so it is reverted. Widening it would have moved rows nobody had shown were missing, in the one place where a wrong guess corrupts a request's context silently and the leak detector would never have said a word. Test results: - bash scripts/run_631_flip_family.sh -> 379 passed. - Metal: automatic flip reproduced twice more (cutover complete on all three ranks); census bracket above; no-flip control clean. - Production restored on POLICY=manual and verified: GET /health_generate 200.

…he resident set Proven by census, and found only because the previous commit's reading ("the request finished while armed") was itself re-tested rather than built on. scheduler.running_batch and scheduler.last_batch are REBOUND to running_mbs[mb_id] / last_mbs[mb_id] at the top of every slot iteration under event_loop_pp. They describe ONE microbatch slot -- whichever slot's iteration is running -- not the rank's resident set. The flip's round hook fires at the END of an arbitrary slot iteration, so build_flip_live_slots_fn sampled a slot that happened to be empty, and an empty slot is indistinguishable from "nothing resident". Census reporting both scopes (2026-08-09 02:21:03Z): at-arm cur_slot_reqs=1 resident_reqs=1 resident_slots=[1] pre-cutover cur_slot_reqs=0 resident_reqs=1 resident_slots=[1] The request was resident in slot 1 throughout. The hook simply ran for a different slot, so the enumeration covered the radix tree only and the request's rows were never moved. THIS IS NOT AN ACCOUNTING BUG. Rows that are not enumerated are not MOVED: the resident request's freshest KV stays in the source pool and is never written into the destination layout, so its context is silently wrong. The leak detector notices the arithmetic; nothing would have noticed the corruption. _live_reqs now enumerates every resident slot (running_mbs), unioned with running_batch/last_batch for the non-PP event loop. J.2, MEASURED AND DELIBERATELY NOT ACTED ON YET. With J.1 fixed the extent probe fires on a real flip for the first time (page_size=1): seqlen=82 kv_allocated_len=81 kv_committed_len=81 cache_protected_len=80 delta_vs_seqlen=-1 seqlen OVER-counts by one against the allocator -- the opposite direction from the hypothesis falsified last commit. Enumerating req_to_token[idx, :seqlen] reads one row BEYOND what the allocator owns and moves it as if it were live KV. The authoritative extent is kv_allocated_len, page-aligned, which is what the invariant checker charges. One measurement on one config is not enough to re-cut an enumeration whose errors are silent; the change is a one-liner once a second flip confirms the sign. J.3, STILL OPEN and the most important. Post-cutover the census reports resident_reqs=0 while the unaccounted page persists, so the resident request appears not to SURVIVE the flip. No flip observed so far has carried a request to the far side, which means the KV-move path has never been exercised end to end with a surviving request. A determined-answer probe on a request decoding ACROSS a cutover is owed before any acceptance claim. Test results: - bash scripts/run_631_flip_family.sh -> 379 passed. - Metal: census bracket above; extent probe now firing; automatic flip still commits on all three ranks. - Production restored on POLICY=manual and verified: GET /health_generate 200.

…ive the cutover The survival oracle, run 2026-08-09 02:36:05-07Z with the idle leak check in WARN mode (SGLANG_ENABLE_STRICT_MEM_CHECK_DURING_IDLE=0) so the accounting crash could not mask the result, and a determined-answer request decoding across the flip: POOL CENSUS at-arm cur_slot_reqs=1 resident_reqs=1 slots=[0] POOL CENSUS pre-cutover cur_slot_reqs=1 resident_reqs=1 slots=[0] POOL CENSUS post-cutover cur_slot_reqs=0 resident_reqs=0 slots=[] cutover complete x3, then AssertionError: x_lru should not be locked when idle, x_lru.full_lock_ref=1, x_lru.id=5 -> Mamba Radix tree sanity check failed -> SIGQUIT Of the three outcomes this experiment could have had, it is the third: the request DIES. It is present and enumerated right up to the cutover and gone immediately after, leaving a stranded KV page AND a stranded mamba lock (full_lock_ref=1 with the tree idle). No content verdict was reachable because nothing continued. This relocates the defect away from the enumeration for the second time. With J.1 fixed, live_slots_fn enumerated the request correctly and the KV move reported balanced cells; the loss is in the CUTOVER, which swaps stacks and scheduler topology without carrying the resident requests across. The pool leak and the mamba lock are two symptoms of that single omission -- fixing either alone would have been treating a shadow. CONSEQUENCE FOR ACCEPTANCE, stated plainly: a flip under load is not merely unproven, it is currently IMPOSSIBLE. Any request resident at the cutover is destroyed. Every flip observed so far committed only because nothing had to survive it, and the acceptance bar (a PP->TP->PP cycle under load in one unmanned log) cannot be met until the cutover carries the resident set. Also recorded: an audit candidate flagged but deliberately not chased -- the false assumption behind J.1 (that scheduler.running_batch names the rank's resident set) is available to any code reading it from a per-slot hook under event_loop_pp, since the attribute is rebound per slot rather than slot-qualified at its use sites. No behaviour change in this commit; findings and the corpse table only. Test results: - bash scripts/run_631_flip_family.sh -> 379 passed. - Metal: oracle run above; cutover completed on all three ranks; leak demoted to warning confirmed the crash was not masking the result. - Production restored on POLICY=manual and verified: GET /health_generate 200.

… J's root cause Supersedes the defect-G and acceptance-readiness sections of the original sgl-project#656 handoff. Records: G fixed and proven on metal; the first policy-driven flip in the feature's history committing at 01:37:12Z; the self-contradictory quiescence predicate that was the real unblocker; and defect J in three parts -- J.1 slot scope (proven, fixed), J.2 row extent (measured, deliberately not cut, with the live-spec-reserve measurement owed before it is), J.3 the cutover not carrying the resident decode set (root cause, and the reason a flip under load is currently impossible). Names the next build precisely -- resident-request carry across the stack/topology swap in build_production_flip_cutover and the scheduler topology snapshot, covering request objects, scheduler bookkeeping and mamba/GDN state AND locks, not just KV cells -- and flags the standing architecture context that design must compose with (sgl-project#635/sgl-project#636 PP dcp_size=1 to TP dcp_size=3 handover and its four silent preconditions; sgl-project#212 store routes truncating GDN state). Also records the two hypotheses of mine that died on the way, so they are not re-derived from a symptom they both fit perfectly, and the SGLANG_ENABLE_STRICT_MEM_CHECK_DURING_IDLE=0 demotion that unmasked J.3. Documentation only; no behaviour change. Suite unchanged at 379 passed.

The blocker: a flip under load was IMPOSSIBLE. Any request resident at the cutover was destroyed, and every flip observed so far committed only because nothing had to survive it. THE DROP SITE IS ONE LINE. Cutover step 6 calls init_pp_loop_state(), which rebinds running_mbs to fresh empty ScheduleBatch objects. Under event_loop_pp running_mbs IS the resident decode set -- running_batch and last_batch are per-slot aliases (J.1) -- so the rebind dropped every resident request: unreachable Req objects whose KV rows stay allocated (the leaked page the idle checker reports) and whose mamba slot locks stay held (x_lru.full_lock_ref=1 -> SIGQUIT). The stranded page and the stranded lock were never two bugs. WHERE THE FIX LIVES IS THE DESIGN. Not at the cutover: init_pp_loop_state has three callers -- boot, the cutover, and event_loop_pp's own entry -- and the TP->PP leg re-dispatches into that loop right after the cutover, so a carry installed only at the cutover would be wiped by the loop it was installed for. The rule is stated at the function that destroys the state: init must never destroy a resident request. Harvest before the rebind, re-seed after it. At boot nothing is resident, so the default path is bit-for-bit unchanged. Once the carry exists the hazard inverts from loss to DUPLICATION: merge_batch extends self.reqs in place and is not idempotent, while init_pp_loop_state is called repeatedly. Harvest therefore dedupes by batch identity (running_batch is normally an alias of a slot) and refuses loudly if one Req is reachable through two distinct batches. Nothing is remapped, and that is a property of the boot rather than an omission: phase_flip_boot step 5a rebinds both stacks' req_to_token and req_index_to_mamba_index_mapping to the SAME tensors, and both layouts key on global slot ids, so every handle a carried Req holds stays valid across the layout swap. SECOND OCCURRENCE OF J.1, FOUND BY AUDIT WHILE BUILDING THIS, and worse than the first because it is silent: gdn_flip_mover enumerated scheduler.running_batch -- one microbatch slot -- so the GDN leg moved the conv/ssm state of whichever slot was current and left behind the linear state of every request resident elsewhere. Since J.1 the KV move carries those requests correctly, so the request would decode on with its linear state truncated at the flip point: sgl-project#212's shape, and nothing raises. Now a named resident_mamba_slots() over the same _live_reqs authority the KV enumeration uses. Pins (the can-fail set is written from the mechanism this build produces, not from one imagined in advance): - the REGRESSION ARM: a cutover whose init drops the set fails loudly with "CUTOVER DROPPED THE RESIDENT DECODE SET", at the cutover, not a pass later with the evidence already stale; - repeated init does not duplicate requests (the merge is not idempotent, the carry has to be); - one Req in two batches is refused; a Req reachable only through last_mbs is refused as a quiescence-predicate bug, not absorbed; - a set left in the other phase's handle is refused by verify_flip_cutover (surviving into a handle the active loop never reads means the requests simply never decode again); - the 3-slot -> 1-slot shrink the TP topology forces; - GDN slots span every resident slot, and the old current-slot-only enumeration is shown returning a strict subset. Also: scripts/route_a_631_survival_oracle.py, the ad-hoc oracle of the J.3 diagnosis turned into a standing harness (determined counting probe, per-chunk stall watchdog, --round-trip), and SPEC=off on the production boot script, which isolates the carry from the draft-state question -- a request that prefills in the PP phase has no draft KV, and that is a second question, not part of this one. Tests: bash scripts/run_631_flip_family.sh -> 399 passed (379 at 46a64ca plus 20 new pins). Desk-only so far; the metal oracle is the next step and this commit does not claim it.

…nd trip works The carry (4c2de6b) made pp_to_tp survivable. Two more defects sat between that and a ROUND TRIP, and neither was visible from the code -- each was found by a leg that could not commit. DEFECT L: tp_to_pp could never reach quiescence under load. Under event_loop_normal the result is processed in the SAME iteration as the forward and ``last_batch = batch`` is set afterwards, so at the hook a non-empty last_batch means "requests are resident", NOT "work is in flight". A decoding request makes it non-empty on every iteration for ever. Measured 03:11:22Z and 03:12:52Z: "NOT QUIESCENT: last_batch is not empty (1 req(s) visible)" on all three ranks, abandoned at the park deadline both times, minutes after pp_to_tp had carried the same request across the other way. This is the SAME CATEGORY ERROR as the _pp_microbatches_drained one that blocked every automatic flip before it: a term that refuses because requests EXIST. Quiescence now asks the narrower question the carry actually needs -- is every live request reachable through the handle the carry harvests? -- which is briefly false right after a prefill and clears itself in one iteration. DEFECT M: the PP chain's ring was read off the LIVE ps. The cutover rewrites ps per phase and the TP phase gets pp_rank=0, pp_size=1, so (pp_rank - 1) % pp_size made UPSTREAM == SELF on every rank. The flip-commit hygiene check then compared a rank's own dict SEND counter against its own dict CONSUME counter -- two different wires -- and rank 0, the first PP stage, sends proxy dicts and consumes none, so its imbalance was permanent and grew with the PP phase's traffic. Measured 03:21:08-03:22:08Z: rank 0 WITHHELD presence for 8889 rounds with "tensor-dict wire has 24 unconsumed message(s) from rank 0" -- itself -- and tp_to_pp abandoned for want of a quorum it could not form. No message was ever unconsumed; the ring was. The counters are built once from the PP topology at boot, so they are the ring's one authority and the helpers now read it from there. METAL RESULT, the thing this feature has never had (SPEC=off boot, one determined-answer request decoding throughout): PHASE-FLIP-CARRY carried 1 resident request(s) ... into the tp phase PHASE-FLIP DONE pp_to_tp (epoch 1) x3 ranks PHASE-FLIP-CARRY carried 1 resident request(s) ... into the pp phase PHASE-FLIP DONE tp_to_pp (epoch 2) x3 ranks oracle: committed legs {'pp_to_tp': 3, 'tp_to_pp': 3} oracle: correct prefix 197 numbers (no break: every token was the next integer) oracle: VERDICT PASS A PP->TP->PP round trip under a live request, both cutovers committed on every rank, and the answer is exactly right. Before the carry the same run destroyed the request. The oracle earned two properties the hard way, both kept: - it reads COMMIT evidence from the serving log, not from the HTTP response. /phase_flip returns 200 for ARMED, and a leg that parks and abandons returns 200 too -- a refused leg once read as a green round trip here; - its verdict is anchored to the FLIP POINT, not to total length. The first probe was a raw completion and the model editorialised about how far to count; the no-flip control drifted EARLIER (32 numbers) than the flip run (43), which is what proved the drift was the model rather than the cutover. Tests: bash scripts/run_631_flip_family.sh -> 402 passed, 0 failed. Docs: PROD_BRINGUP_BENCH.md section 8 records all three defects, the oracle recipe and the SPEC=off knob. NOT claimed here: the acceptance program (POLICY=auto, mixed load, one unmanned log) and the speculating rung -- a request that prefills in the PP phase has no draft KV, which is a separate question and is why these runs are SPEC=off.

…al form of both L and M were fixed in cdcd8d6 but only the commit message carried the reasoning. The corpse table is the design document this feature is navigated by, so both belong in it, together with the general form each one turns out to have: L -- WHAT MUST BE QUIET IS THE MACHINERY, NEVER THE WORKLOAD. Twice now a quiescence term has refused because requests EXIST rather than because work was in flight, and a term with that shape contradicts a feature whose whole purpose is to flip while requests are alive. M -- ANY QUANTITY DERIVED FROM ps IS PHASE-SCOPED, because the cutover rewrites ps. The PP chain's ring must come from something that does not move; the counters are built once from the boot topology and are that thing. This is the flip's own version of the running_batch audit note already in the table.

… is speculation v2 closed with 'a flip under load is not merely unproven, it is currently IMPOSSIBLE'. That is no longer true: a PP->TP->PP round trip under a live request, both cutovers committed on all three ranks, answer exactly right (197 consecutive integers, no break). Records the three defects that stood in the way (K the carry, L the return leg's quiescence, M the PP ring read off a ps the cutover rewrites), the second SILENT occurrence of J.1 in the GDN mover, and -- the thing the next reader most needs -- the WALL, named precisely with its traceback: SPEC=on plus a carried request kills the instance, because a request that prefills in the PP phase has no draft state and the PP phase carries no draft worker by design. That is sgl-project#631's central path, not an edge case, and it is a separate build. Three options are listed with the silent hazard each shares (sglang assumes draft seq_len == target seq_len).

…held the wrong layout REPORTED FROM THE OUTSIDE, which is worth recording: the operator looked at the running rig and said the load was clearly being served by the TP layout and not the PP one. It was, and the reason was a fourth instance of the failure shape this feature keeps producing -- a quantity that does not mean what its name and its comment say. THE METRIC. maybe_arm_phase_policy fed the policy pending_prefill_tokens=sum(len(req.origin_input_ids) for req in self.waiting_queue) while the comment directly above it said "prompt tokens ADMITTED but not yet computed". The waiting queue is the NOT-yet-admitted work. A long prompt under chunked prefill does not sit there: it hangs off scheduler.chunked_req and is filled a chunk per round. So for the whole duration of exactly the work the PP layout exists to do, the policy read 0 pending prefill, and the TP->PP rule (pending > N) could not fire. MEASURED, POLICY=auto, N=4096, 8k/32k prompts arriving every 5s (2026-08-09 03:36-03:39Z): the acceptance ran its ENTIRE 186 s mixed phase in the TP layout on ONE policy decision, and the layout only corrected itself on the idle return to rest. phase at start: tp observed phase timeline (2 transitions): t= 2.0s tp t= 194.1s pp requests: 39 total, 39 ok, 0 ABORTED The fill boundary is extend_range.end, which is what the scheduler's own chunked-prefill code uses (_compute_chunked_req_next_prompt_token), so the remainder behind it is the admitted-and-uncomputed quantity the policy was always meant to weigh. Evaluated on the request-origin rank only, so no cross-rank replication question arises. A DECLINING POLICY WAS ALSO SILENT. Only arming decisions were logged, so "the layout is wrong under load" and "the hook never runs" looked identical from the log -- and the second was my first hypothesis, wrongly. One throttled line now carries the standing hold reason with the two inputs that produced it. Same bet as the withhold reason and the quiescence reason, each of which named a defect in a single boot. ALSO IN THIS COMMIT: speculation is honest by WAITING. A carried request has no draft state and SIGQUITs the draft graph runner (handoff v3 §3), so with speculation armed for the TP phase a pp_to_tp flip is now simply NOT READY while anything is resident -- the flip happens as soon as nothing has to survive it, which is the regime every flip before the carry ran in. Waiting rather than refusing at arm time is deliberate: a rank-local refusal inside arm() would let one rank decline while its peers armed, and diverging epochs is corpse H, which is fatal. Readiness runs through the bounded park/abandon machinery, which is unanimous by construction. Tests: bash scripts/run_631_flip_family.sh -> 404 passed, 0 failed. New pins: a chunked prefill larger than N moves the layout; a nearly finished one does not (the REMAINDER is what counts, not the prompt's size); the speculating TP phase waits for the resident set to drain, and does not wait when there is nothing to carry or when spec is off.

…ays changed Caught on the boot that introduced it, which is the only reason it is not in production: the policy's hold reason carries live quantities ("min dwell: 13.4s since last flip < 15s"), so keying the throttle on the reason STRING made every call look like a new reason and the throttle never engaged. Measured: three identical lines inside one second, and 368 hold lines in two minutes of which 344 were "min dwell" repeats. A 12765-line log flood has already cost this feature a self-kill once (see the origin guard in request_receiver), so this is not cosmetic. The key is now the reason with its digits removed -- the SHAPE of the hold rather than its instantaneous value -- which still logs immediately when the hold changes CHARACTER (min dwell -> pending prefill -> decoding in tp) and otherwise at most every 10 s. Tests: bash scripts/run_631_flip_family.sh -> 404 passed, 0 failed.

init nccl env

b961081

Ying1123 requested a review from merrymercy July 16, 2024 10:23

merrymercy merged commit 0aa189f into main Jul 16, 2024

merrymercy deleted the ying-env branch July 16, 2024 16:05

timethink pushed a commit to timethink/sglang that referenced this pull request Mar 9, 2025

Disable NCCL_NVLS by default (sgl-project#631)

b40ef6d

Navigation Menu

Uh oh!

Disable NCCL_NVLS by default - #631

Disable NCCL_NVLS by default#631
merrymercy merged 1 commit into
mainfrom
ying-env

Ying1123 commented Jul 16, 2024

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

Ying1123 commented Jul 16, 2024

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants