Improve benchmark scripts by Ying1123 · Pull Request #717 · sgl-project/sglang

Ying1123 · 2024-07-24T21:43:18Z

No description provided.

* Update pr-test-npu-debug.yml * debug testcase * debug testcase * debug testcase * debug testcase * debug testcase * debug testcase * debug testcase * debug testcase * debug testcase * debug testcase * Update pr-test-npu-debug.yml * Update pr-test-npu-debug.yml * Update llama4.py * Update pr-test-npu-innersource.yml * Update pr-test-npu-innersource.yml

…ead them the same _resident_ceiling returns -1 for two OPPOSITE states, and its own docstring says so: "Highest row a RESIDENT REQUEST pins, or -1 when none/unknown." build_flip_live_slots_fn sets req_max to -1 when it has no request parts, so an IDLE box is encoded exactly like an UNREADABLE split. The two demand opposite behaviour: * unreadable -> evict NOTHING. Unmapping a row a live request is reading is the one unrecoverable error, so the conservative branch is right here. * none resident -> the entire live set is tree-only and MAXIMALLY evictable. _evict_floor_rows took the conservative branch for both, which means it refused to price an eviction precisely when it had the most to win. On the specimen: BOTH BLOCKED ... 0 req resident KV rung: current=137216 rows, floor=398471, slack=0 397,958 max_live against a 137,216 cap -- 2.9x over -- so slack pins to 0, the sgl-project#688 evict-rung funding path is structurally dead, and every flip falls back on the raw seam budget. That is why three FLIP ABANDONED events turned on a mere 55 MiB shortfall. ROOT DIFFERS FROM THE FILED HYPOTHESIS, and the distinction matters for the neighbouring tasks: the ids are NOT stale pre-shrink survivors needing a rebase. 397,958 is a perfectly valid id in the ~437k id space; the pool simply has only 137,216 rows BACKED. The floor is max(live id) + reserve, a SPARSE high-water mark, so ~137k live rows scattered across the id space demand 398,471 rows of backing. Nothing is un-rebased; the mark is just the wrong statistic when the live set is sparse. A rebase would not have fixed this. The two states are separable from data already recorded -- the live-set split carries req_rows beside req_max -- so no new enumeration is needed. Tests, hermetic (CUDA_VISIBLE_DEVICES=""): test_the_specimen_shape_reproduces ......... the plain floor IS 398471, and it exceeds the 137216 cap, so slack pins to 0 test_nothing_resident_prices_the_eviction .. the fix: floor collapses under the cap and an eviction is priced test_unreadable_split_still_refuses ........ CAN-FAIL: the conservative branch must survive for its real case. A fix that treated -1 as "evict everything" passes the fix test and fails this one -- and that mistake unmaps live rows. test_resident_rows_present_is_unchanged, test_nothing_resident_predicate_is_exact -> 5 passed + 3 subtests; 76 passed + 5 subtests across 717 + flip-runtime + 715 suites. ruff clean in the changed region. CAN-FAIL PROVEN BY MUTATION: restoring the conflation fails 1 of 5 (the fix test), reverted and re-verified 5 passed. Desk only -- not yet observed on metal. The live proof is a flip funded by the evict rung on an idle box, which the next boot can show.

…ndition, not a cause Follow-up to b9f4580. F4-r4 rooted sgl-project#717 (c4e5579) and DISPROVED the framing I had put in that message. I wrote that a floor above the cap means "live row ids outliving a pool shrink/reshard". That is wrong. 397,958 is a VALID id in the ~437k id space while only 137,216 rows are backed: the live set is SPARSE, so a high-water ID above the backed-row count is normal, not evidence of staleness. max_live is an ID, not a count, and I read it as a count. The actual root was _resident_ceiling encoding "none" and "unknown" as the same -1 sentinel, so an idle box read as an unreadable split and eviction was never priced -- a sentinel-collision defect, nothing to do with rebasing. The guard itself stands and is unchanged in behaviour: whatever the cause, a floor above the cap pins slack to 0 and the rung cannot fund, which is worth naming and was invisible before. Only the diagnosis is removed. The message now states the condition and says explicitly that max_live is a high-water ID over a sparse set -- "state the fact, do not infer the cause" -- so the next reader is not sent down the road I went down. Tests unchanged, 5 passed: they assert the condition and the gap, never the cause, which is why the correction needed no test edit.

…seam refuses Scope (c) was framed wrong and the correction is the headline. The attention KV ops do NOT route through any process group -- get_kv_buffer/set_kv_buffer are LOCAL memory operations, and a repo-wide grep for group accessors inside forward_extend/forward_decode/set_kv_buffer/get_kv_buffer, radix_attention.py, base_attn_backend.py and forward_context.py returns nothing. The group enters one level up at the cross-rank LSE merge: cp_lse_ag_out_ar_mha_uneven (layers/dcp/comm.py:231) takes cp_group as a PARAMETER, and its callers supply it as `group = get_parallel().dcp_group` (flashinfer_backend.py:5634, :2478, :5684, :5747, :5863; triton_backend.py:2311). So B1 needed no new call site. runtime_context.py:392-394 resolves dcp_group through _v() (:240-242), which calls the getter on EVERY access rather than caching, and get_dcp_group() already carries the B1 branch (:2744). What this slice owed was a TEST of that resolution, not a wiring change -- and the test pins the uncached property specifically, because if dcp_group ever starts caching, arming would silently stop taking effect after the first read. There is also no choke point: ~15 attention-backend files call the KV buffers inline. Ownership is expressed in the POOL SHAPE and the MERGE GROUP, never at the access sites, which is why that sprawl does not have to be touched. R6's build anchor was wrong. It cited model_runner_kv_cache_mixin.py:2496-2500, which is real but lives in _init_unified_mamba_pools (:2467) -- reached only under --enable-unified-memory (:3194-3200), a path this deployment does not take. The filter governing our hybrid-mamba config is the classic-path one at :3929-3934 inside the HybridLinearKVPool construction (:3913). Same logic, wrong citation, so the byte-identity claim had never been stated against the path it would be tested on. Corrected in the docstring. R6 had also DUPLICATED a shipped sizing rule with the wrong rounding. plan_decoupled computed round(share * C); the shipped rule is dcp_compact_pool_rows (layers/dcp/owner.py:155-181), already called by this same build path at :3765-3767, whose docstring says it is written once because "a sizing rule whose off-by-one has already cost a debugging round must not exist twice". I had made it exist twice AND rounded the wrong way: the shipped rule CEILS to a whole owner block, and flooring is what let trailing partial block slots scatter out of bounds (async illegal memory access, found by the kv-session-offload S1 test at --max-total-tokens 3000 on cp_S=64). A planner handing the build a floored row count re-creates that bug at the seam. plan_decoupled now delegates to dcp_compact_pool_rows. Consequences taken, not papered over: the plan carries the period it ceiled against, and world conservation DERIVES the ceil slack from it (at most `period` rows across the world) rather than being handed a slack that could be widened until it passed; a zero-row pool became unreachable rather than caught. Arming point: phase_flip_runtime.py step 1, alongside set_phase_flip_tp_active (:1196) inside _cutover (:1160) -- after the group-wide quiesce (on_round consensus :3066, hold :3121-3126, execute :3128) and after the KV wave-move, before any group handle is re-derived (:1224). B1 is a PP-PREFILL mechanism, so arm_for_cutover arms into PP and DISARMS into TP; it does not lean on the flip route out-ranking B1 in get_dcp_group, because precedence is the net and disarming is the contract. Relying on branch order is how a later reorder becomes silent capture. The red-first case: armed routing over a stage-local pool. A rank may be asked for ANY attention layer, and a stage-local pool has no rows for those it does not own, so _transfer_full_attention_id (memory_pool.py:3855-3859) misses the mapping or returns another layer's rows -- wrong output, no error. record_pool_plan records what the build ACTUALLY built (an intention cannot be compared against reality), and arm_decoupled_kv refuses unless that plan is DECOUPLED and covers the full attention set. Disarming is never refused: requiring a healthy pool in order to STOP using it turns the recovery path into a second failure. set_decoupled_kv_active now refuses to arm without an initialized group, following set_phase_flip_tp_active (parallel_state.py:2670-2687) instead of the silently-no-oping set_dcp_spill_active. That precedent names the failure class: "a silent no-op all-reduce, the exact corruption class this routing exists to prevent". The get_dcp_group fall-through stays as defence in depth for a state the setter now makes unreachable. Test results (hermetic, CUDA_VISIBLE_DEVICES="", interpreter /spinning/htsglang-gpu/.venv/bin/python3): test_decoupled_kv_arming_704b.py (new) + test_decoupled_kv_pool_plan_704b.py + test_decoupled_kv_group_b1.py 40 -> 58 passed can-fail by breakage: neutering assert_pool_supports_arming and the set_decoupled_kv_active group check turns EXACTLY 6 tests red; restoring returns 58. No collateral. regression measured against a REMOVED-DIFF BASELINE, not assumed: baseline mem_cache 944 failed / 835 passed, distributed 21 / 2706 with diff mem_cache 944 failed / 853 passed, distributed 21 / 2706 => same failures, +18 passes, zero new reds. planner 2805 passed. model_executor clean. The 965 pre-existing reds are the hermetic-mode consequence ("No accelerator (CUDA, XPU, HPU, NPU, MUSA, MPS) or platform plugin is available", utils/common.py:1322), not a repo defect and not caused here. A hermetic run of these directories can only be read as a delta against that baseline. BOOT DEPENDENCIES, nothing below is proven without a reviewed boot: 1. phase_flip_runtime.py is NOT wired. arm_for_cutover exists and is tested, but no call was inserted at :1196, because this worktree's copy (6666 lines, [sgl-project#690]) is BEHIND the deploy tree's (6984 lines, [sgl-project#717], ~355 diff lines apart) and the deploy tree has no B1 symbols at all. A call against the stale revision would likely be discarded on merge. The insertion point is documented in the function docstring. 2. SGLANG_DECOUPLED_KV=1 has never been booted; no pool allocated from it. 3. The share vector at boot is unverified -- whether get_cp_token_ratios() and cp_token_split_factor(dcp_size) are populated when _init_pools runs is a boot fact, not a desk fact. 4. The LSE merge has never run over B1; only the resolution chain is pinned. 5. Nothing yet cross-checks that every rank recorded a plan with the same period and layer set. B1's manifest covers the group, not the pool. Deploy tree /spinning/wt-678-deploy untouched, as required.

…e rung read them the same" This reverts commit c4e5579.

…ferable vs lazy-fundable The floor was measured once and frozen, never designed, and that is visible in the source rather than a matter of opinion: arming_floor = corridor band floor + seam_entry_reserve, where the first half is a stated policy with a stated tolerance (1024 - 20% = 819 MiB) and the second is one scalar whose own docstring calls it "the shipped allowance". Per-rank it decomposes exactly: rank 0 (5090) 1728 MiB = 819 + 909 seam draw rank 1 (3080 x4) 1825 MiB = 819 + 1006 rank 2 (3080 x8) 2467 MiB = 819 + 1648 Rank 2 holds 81% more than rank 0 with no recorded reason. The defect is not the size of the number, it is that it is ONE number: a monolithic holdback cannot be traded, because trading requires knowing which part buys what. STAGING, AND WHAT ACTUALLY FORCES IT. Under barlink a cross-card transfer is one PCIe crossing, so a host bounce (two) is strictly worse and is not offered there. But the reshard stages in VRAM on both sides regardless of transport, and rollback semantics are NOT why: _dist_exchange allocates a uint8 receive buffer per peer and sends a gathered contiguous buffer (kv_reshard.py:939-995), because the wire format is a flat byte stream per peer while the destination rows are scattered ids. The obstacle to BAR1 landing in-place is the SCATTER, not rollback and not write ordering -- so the cross-card staging term goes to ~zero only under a layout co-design that makes a peer's destination extent contiguous. That is the largest structural reduction available here. RESTORE, NEVER REBUILD adopted as a named invariant: the flip may pay copy time, never build time. Consequence: capture-moment workspace is a BOOT-time component, not a per-flip floor term (both layouts captured once, graph state parked to host and restored per flip, 40-85 ms band, sgl-project#464 coalescing pending). The weights refill already complies. TIME PRICE of host-bouncing, from the H2D rates I measured in sgl-project#690 rather than nameplate: 125 / 199 / 181 ms for the whole seam draw on ranks 0/1/2, i.e. 4-6% of a ~3.1 s flip to return 909-1648 MiB per card permanently. Stated as a FLOOR on the cost, not a wall-clock delta: the refill is already 41-52% of the flip and shares the same link, so bounces queue behind it -- worst on rank 1's x4 card, where reslotting remains the cheapest single intervention. NAMED AS UNMEASURED: the seam draw cannot currently be split into components; nothing records it per component. Until one instrument attributes the peak instant to (send buffers, receive buffers, graph state, allocator transient), any per-component trade is arithmetic on an undivided number -- the same error class as pricing a flip on intention rather than completion. Also records the two honest limits on the rebuilt evict rung (sgl-project#717): it now delivers less than it prices by design, so repricing must key on the DELIVERED amount; and the ~413 MiB it leaves on the table is recoverable only by evicting deeper than the cap. Desk analysis, no code change. Feeds sgl-project#702 repricing and sgl-project#677 economics.

…ason strings, file the host-tier gate 1. THE GUARD. check_page_head_alignment in csrc/kvcacheio/transfer.cu, requiring item_size % (8 * head_num) == 0, called from BOTH page-head entries: transfer_kv_all_layer_lf_ph and transfer_kv_per_layer_ph_lf. Guarding only the reported one would have left the same asymmetry sgl-project#717 was reverted for -- a defect that looks fixed from one side. The message names the two facts that EXPLAIN the rule, not just the rule: the .b64 PTX width (ld.global.nc.b64 / st.global.cg.b64 fault on a misaligned address) and the subdividing offset formula (get_global_offset_ph divides three of its four terms by head_num). A reader who learns only the rule works around it by changing head_num. Pinned: the predicate refuses the faulting shape (head_num=2 head_dim=2 fp16, item_size=8 passes the OLD guard, head_size=4 does not), does NOT falsely refuse the reported-crash shapes or this rig's shapes, and AGREES WITH ACTUAL OFFSET ALIGNMENT on every shape -- so a future edit cannot keep the tests green by tuning the constant. A source pin holds it wired at both entries. MUTATION PROOF: unguarding one entry and weakening the predicate back to % 8 turns 2 pins red. 2. BOTH REASON STRINGS. test_kvcacheio.py's skip said "segfaults on CUDA 13.x". It now states the evidence: reproduces on BOTH wheels, root NOT yet attributed, the alignment defect found here is real but is NOT this crash's cause (these shapes are aligned), and the metal falsifier is at tools/441/. The >= 13 condition is deliberately left as-is with the reason recorded -- widening it would hide the cu12 signal the falsifier needs, removing it would hand CI a known segfault. test_minimax_sparse_pool_host_unit.py's batchcopy comment now LEADS with the reading it previously only implied: the flag guards a CONTRACT VIOLATION COMMITTED BY THE TEST, not a wheel, ABI or CUDA-13 bug, and the measured matrix proves it because pinned+side-stream passes on the same wheel that fails the other three. It says explicitly that the sgl-project#436 rebuild cannot make it obsolete, since a rebuild cannot change a contract the test breaks by construction. 3. sgl-project#261 GATE: THE "SHIM" IS LOAD-BEARING. Verified at the code, which is what was owed. Step 4 is `python -m sglang.srt.mem_cache.hicache_migrate` with --target-tp-size / --target-ratios / GDN layout arguments. The live gate runs source A at TP=1 and destination B at TP=N (live_handover_gate.sh:37), and stored KV pages and GDN state blobs carry the SOURCE's shard geometry. Without the conversion B is not reading the host tier the hard way -- it is handed bytes in a layout it cannot interpret, and the run fails for a reason unrelated to the host tier. So a no-shim 1->N run is structurally meaningless, which is the stop condition the brief named; the cross-geometry handover IS the umsharder. What IS meaningful is holding the geometry EQUAL and dropping the migration: no conversion is needed by construction and the host tier is what remains under test -- write-through, park/export, prefetch, resume from cache, byte-identical continuation. Weaker than the 1->N gate, and real. Filed as scripts/handover/hostier_gate_noshim.sh, syntax-checked and refusing cleanly on missing env, NOT RUN. Its preflight REFUSES when the two servers' geometry/checkpoint differ rather than producing a red a reader could mistake for a host-tier finding -- that refusal is the script's safety property. It keeps the original gate's two load-bearing pieces: the A-vs-A floor FIRST (a cross-server byte claim on a rig whose own repeat is not byte-identical is not a claim about handover) and --expect-cached (without it a destination that re-prefilled reproduces the same greedy tokens and the comparison passes while proving nothing -- observed for real on the sgl-project#261 gate before it was added). Tests, hermetic (CUDA_VISIBLE_DEVICES="", no CUDA context, no launch, no GPU, no server contacted): test_page_head_offset_alignment_441.py 13 passed + 14 subtests in 0.07 s test/registered/unit/mem_cache/ 940 failed / 803 passed -- the 940 are the unchanged PRE-EXISTING baseline on this base (b786858 lacks the sgl-project#585 conftest); passes moved 796 -> 803, exactly the 7 new pins. ruff + codespell clean; the shell script passes bash -n; edited Python files byte-compile. The C++ guard is NOT compiled here (no CUDA build in a hermetic run). Its predicate is pinned arithmetically and its wiring by source inspection; a compile rides the next kernel build.

…aim: the arena tails are measured and recorded The sgl-project#702 repricing under the cold-spill doctrine found the number I said did not exist. NOTE_677 section 8 stated "no 1456 MiB figure appears anywhere in the records or source". WRONG: managers/phase_flip_seam_reserve.py, record_path docstring, records the measured per-rank arena tails -- "1436 MiB on rank2 against 466 MiB on rank1 and 0 on rank0". 1436 rather than 1456, but plainly the figure the earlier candidate meant. I missed it by grepping the docs tree and the ticket number instead of the module that owns the quantity. The sgl-project#685 ticket attribution was still wrong (it is an UnboundLocalError); the NUMBER was real. It also settles, without a boot, what NOTE_677 left as a candidate. Against the standing pool reduction per rank -- 704 / 801 / 1443 MiB, i.e. arming_floor_subtrahend_bytes = floor - max(corridor law, already_reserved) -- the measured tail is 99.5% of rank 2's, 58% of rank 1's, and 0% of rank 0's. Rank 2's floor excess IS its arena tail. The repricing itself lands in the evidence tree (not a git repo): /spinning/evidence-665-f1/NOTE_702_CUT_TABLE.md sections D1-D6, referenced from an append-only block in PLAN_PERF_PIPELINE_2026-08-16.md (head-298 sha256 verified byte-identical, 81560e94b7628e6c). Headline: crediting the standing reduction back, [31,17,16] recovers to 530,381 -- EXACTLY [31,18,15]'s spill-funded figure, because once floors stop binding both bind PP0 on the same bytes. The pool advantage that made [31,18,15] interesting was an artefact of VRAM-resident floors, as the user said; with it gone [31,17,16] wins outright on speed and is the only one of the two above the 14.1% noise floor. Build NOT undertaken and NOT small: nothing credits ON-DEMAND capacity against the standing floor, and crediting the evict rung's PRICED capacity would be pricing on intention rather than completion -- sgl-project#717's defect one layer up, with an OOM at arm time as its failure mode. Ticket filed at D6 with one live _record_seam_peak reading as prerequisite. Desk only, no boot, no GPU, no serving contact.

…s what the watchdog polls Root for the family isolated in ANALYSE_734 (S2 CUDA illegal memory access 2026-08-17; S3 Bar1CollectiveAborted 2026-08-05, different model, tp3/pp1). THE DEFECT. barlink_bar1.py Bar1Transport.close() (:5563) releases the device memory the watchdog is still polling and disarms the watchdog FOURTEEN LINES LATER: :5643 self._cuda.vmm_free(*w) <- memory released :5646 self._ctl_dev = None :5657 self._abort_poll_active = False <- watchdog disarmed The code's own comment two lines above the disarm states the correct intent -- "the watchdog poll ... reads _ctl_dev directly -- it has to stand down with the word". It does not stand down WITH the word, it stands down AFTER it. The watchdog runs on its own thread (poll_status_word, :4807), takes no lock, and is not joined before the free. ONE DEFECT, TWO PRESENTATIONS, chosen by the allocator rather than by anything in the failure: - block unmapped -> the poll faults -> CUDA illegal memory access (S2) - block recycled -> the poll reads garbage -> non-zero status word and the second arm closes because THE ABORT CODE IS NEVER VALIDATED: there is no known-code table anywhere in the module, so check_aborted raises "a spin kernel took its abort path" for any non-zero word. S3 therefore reports a spin-kernel abort that no spin kernel performed. That is the unification. WHY THE LOG ORDERING DOES NOT INCRIMINATE THE POLL, now by mechanism rather than caution: barlink_abort_gate.poll_status_words() (:327) swallows the fault under "except Exception # a watchdog must not die". A CUDA IMA is STICKY -- it poisons the context, so every later CUDA call re-raises it. The gate keeps polling and re-reports the same fault, which is exactly S2's "four repeated" errors, and the serving path only learns of it at its next CUDA touch, which happened to be the kv_reshard checksum. Gate = first reporter, payload path = second reporter, neither is the perpetrator. Same reporter-vs-perpetrator shape as sgl-project#722, recurring one layer down. THREE CONDITIONS keep it alive independently, and one of them widens the verdict's reach: there is NO record_stream() anywhere in the module (searched, zero hits), while the poll copies from _ctl_dev on a PRIVATE stream with non_blocking=True. That is the documented caching-allocator hazard on its own, so the family does not strictly require a close() -- a plain recycle under memory churn suffices. Which is consistent with S3 happening on a boot with no phase flip at all. Also: pause_polling() exists but has exactly ONE real caller (parallel_state.py:2888, for graph capture, and its docstring says why); teardown and arena remap never got the same treatment despite the same cross-thread hazard. WHY S2 SURFACED ON THE FLIP PATH: its server_args read tp_size=1, pp_size=3 with enable_phase_flip=True and phase_flip_tp_vector='32,16,16'. At tp_size=1 there is no steady-state TP collective for barlink to serve -- the transport exists in that boot FOR THE FLIP's TP layout, so the flip is what builds and tears these down while the server is live. FIX SHAPE filed (5.1 disarm-then-quiesce before free; 5.2 atomic guard; 5.3 record_stream, required for the S3 arm; 5.4 pause across flip teardown reusing the existing context manager; 5.5 narrow the gate's except so a sticky CUDA error is not swallowed; 5.6 validate the abort code so garbage becomes a named refusal instead of a false accusation). Not built in this commit. Open and named: lineage (which commit each specimen ran, sgl-project#634's exact closure and reopen wording, sgl-project#717/sgl-project#722 dates) is being pulled separately and the sgl-project#634 reopen decision is NOT made here; and I have not traced the flip's call into close() line by line -- the ordering defect stands on its own reading and condition 2 means the family does not depend on that link, but the S2 attribution does. No boot; source and the two logs only.

…flip, and sgl-project#722/sgl-project#634 stay closed I reproduced a RETRACTED verdict. My previous commit (1b6b989) rooted the barlink-BAR1 family in a teardown-ordering race in Bar1Transport.close(). That is task sgl-project#722, which was already filed with the same mechanism -- OPERATOR-STATE :9053-9057 records it almost word for word ("poll_status_word liest stale _ctl_dev-Mapping ... Guard prueft is-None = IDENTITAET statt VALIDITAET") -- and then RETRACTED on live evidence: the abort-poll is REPORTER not perpetrator, the real root was sgl-project#717's backing mismatch, and the revert was proven over 216+33 flips with 0 faults. WHY I MISSED IT: I gated the prior-art search on the repository. Both sgl-project#722 and sgl-project#634 live in OPERATOR-STATE.md, outside it. The subagent's "NOT FOUND in the repo" was correct and the repo was the wrong place to look. The gate is only as wide as the corpus you point it at, and I pointed it at one. THE REAL ROOT, from S2's own log: 48743 19:30:31 PHASE-FLIP armed: pp_to_tp 48749 19:30:31 FLIP EXTENT PROBE: seqlen=51311 kv_allocated_len=51310 48740 19:30:31 CENSUS at-arm: cached=127182 cur_slot_reqs=4, backing=309464, highest live row=183998 48931 19:30:40 KV-BACKING EVICTED 127731 recomputable rows to bring the high-water mark below 61303 (resident ceiling -1) 48935 19:30:40 backing 61303 instead of 116736, highest live row 0 48955 19:30:4x CUDA illegal memory access Twenty-four lines from eviction to fault. _nothing_resident() returns True when _last_live_split["req_rows"] == 0. During a pp_to_tp flip the requests are PARKED for the consensus round, so req_rows reads 0 while 127182 rows are still cached and the flip's own probe has just measured a request at seqlen=51311. The rung takes the nothing-resident branch, evicts 127731 rows, shrinks backing 309464 -> 61303 -- and the flip then reads the rows it was about to move (phase_flip_runtime:6926 _execute -> _pack_outgoing -> kv_reshard.py:359 _checksum). Those rows are above the new cap. THE SAFETY NET CANNOT COVER THIS BY CONSTRUCTION. Site A's comment says the optimistic pricing "is only safe BECAUSE" _shrink_to re-reads the live set afterwards. It re-reads the live REQUEST set, and the flip's pending pack is not a live request -- it is parked, which is the very state that made req_rows zero. Check and hazard share one blind spot. This is b786858's revert text ("backing shrank to 69,054 rows under a highest live row of 233,289, and the next access above the cap was an illegal address") recurring through the REBUILD 675793c, on a path the rebuild did not consider: not a live request above the cap, but the phase flip. LINEAGE, by git merge-base --is-ancestor against serving commit 9255258: c4e5579 (attempt), b786858 (revert) AND 675793c (rebuild) are all IN. So S2 is post-revert AND post-rebuild -- and that distinction is the answer. Post-revert alone would suggest a barlink reopen; post-rebuild says the reverted defect returned through its own replacement. S3 (2026-08-05) predates sgl-project#622 (08-08) and sgl-project#717 (08-17) entirely: old world, not a trigger. VERDICT: sgl-project#722 and sgl-project#634 do NOT reopen. The trigger is a new post-revert specimen with a BARLINK perpetrator; S2's perpetrator is the KV rung and the abort-poll is again the reporter -- so this note is a second independent confirmation of sgl-project#722's retraction, not a challenge to it. sgl-project#741 reclassifies as a sgl-project#717-rebuild x phase-flip regression: live, rooted, unowned. The barlink stand-down/re-arm change stays UNBUILT per the standing order; my section 5 is superseded. Fix shape for the real root (not built, three options with a preference): make the flip's pending extent visible to the split so req_rows counts parked-for-flip rows (preferred, smallest, closest to the defect); or gate the rung off while a flip is armed (blunt, costs the idle-box win, cannot be wrong); or extend _shrink_to's re-read to the flip's pending set so the net covers what its comment claims. Sections 1-4 of the note are KEPT, marked retracted, as the record of what was claimed and why it is wrong.

Fix for the root filed in ANALYSE_741. On 2026-08-17 19:30:40 the KV backing rung evicted 127,731 rows that the phase flip was about to pack, and the next access above the new cap was an illegal memory access 24 log lines later. It killed that day's serving. WHY BOTH HALVES WERE BLIND AT ONCE, which is what shapes the fix. _nothing_resident() asks the live split for req_rows, and _shrink_to's safety net re-measures through _max_live_row() -- which calls the SAME _live_slots_fn. A flip quiesces its requests before packing them, and a quiesced request sits in none of the batch structures _live_reqs enumerates, so that one enumeration reported zero to both. Fixing only the predicate would have left the net equally blind to any other caller that shrinks during a park, so both are fed from ONE side channel. LINE 1 -- the parked extent is visible to predicate AND net. _nothing_resident() consults _flip_pending() first; _max_live_row() folds the parked max row id in, which is the half a predicate-only fix would have missed. (-1, -1) means UNKNOWN and blocks in both -- never reads as empty. LINE 2 -- the rung refuses outright while a flip is armed, on the pricing and the collecting site both, because the failure is a SILENT eviction followed by a DELAYED fault and one mechanism is not enough. Gated on ARMED ONLY: outside a flip the rung stays fully live, which sgl-project#688's evict-rung funding path depends on. WIRED, not merely defined. Both channels are passed at the production factory against scheduler.phase_flip_runtime, and an AST pin asserts the factory actually passes them -- parsed rather than grepped, because the module's own prose names both kwargs while explaining them. A fix wired into nothing is the defect class this repo keeps finding; it is not shipping as one here. TEST RESULTS (repo rule: documented before commit) New suite test_evict_rung_flip_park_744.py: 15 tests + 2 subtests, all green. Can-fail proofs cut BOTH ways on purpose -- removing the extent visibility must let the eviction happen again, and the gate must NOT make the rung dead outside flips. Mutation matrix, 6 mutants, ALL KILLED (run against the new suite plus both sgl-project#717 suites): M1 predicate ignores parked extent -> 3 failed M2 net ignores parked extent -> 1 failed M3 gate removed from pricing site -> 1 failed M4 gate removed from collecting site -> 2 failed M5 gate ALWAYS on (rung dead outside flips) -> 3 failed <- the sgl-project#688 guard M6 unknown extent treated as empty -> 1 failed M5 is the one that matters most: it proves the gate is not permitted to disable the rung outside a flip. Pinned suites green: test_evict_rung_nothing_resident_717, test_evict_rung_floor_invariant_717, test_kv_backing_cap_agreement_656, test_kv_backing_recovery_clamp_684, test_kv_backing_exhaustion_662_f4 -- 88 passed with the new suite. Full managers regression, compared PER SUITE against HEAD in a scratch worktree (per the truncated-baseline lesson -- a total can hide a swap): baseline HEAD : 29 failed, 4 errors, 2522 passed with sgl-project#744 : 28 failed, 4 errors, 2540 passed NEW failures : ZERO The single baseline failure absent from my run (test_phase_flip_mover_streaming_631 :: test_the_waved_price_is_short_of_the_ measured_live_set) is NOT claimed as a fix: it passes 3/3 in isolation on BOTH trees, so it is order-dependent in the full run and unrelated to this change. Ruff and codespell clean. Two things stated plainly rather than dressed up. The fix was built BEFORE the falsifier, contrary to red-first; the mutation matrix is the equivalent evidence and the ordering was mine. And the sgl-project#717 suites caught a real error at once -- they invoke the rung's methods UNBOUND against stubs, my first version used direct attribute access, and all four went red; fixed by matching the file's existing getattr(self, ..., None) idiom. KNOWN LIMITATION, filed as its own follow-up: _flip_pending() returns the last enumeration that saw requests, consulted only while armed. That makes line 1 CONSERVATIVE rather than exact -- a flip that arms before any enumeration has seen requests answers UNKNOWN and blocks. Safe, but the exact arm-time snapshot is not shipped here. No boot. Boot proof goes on F4-r4's window list.

…- unstrangle the funder My sgl-project#744 fix has a live regression and this is the refinement. Specimen 21:46:32 on 72696b0aec: tp_to_pp refused 35x, IDLE-LOCK with 407,622 tokens pending and 0 resident, the guard reporting "no KV provider is registered" because the armed-only gate refused the rung at BOTH sites and left the provider list empty. Seam staging (PP1: needs 1269 MiB, spendable 883, rung wanted SHRINK to 126506 with slack 45526 = a 1455 MiB deficit that WAS coverable) could never be funded. THE ERROR WAS MINE AND IT WAS A DESIGN ERROR, not a slip. Seam-funding eviction of recomputable prefix rows is REQUESTED BY the flip machinery -- it is what "KV capacity is the funder" means -- so refusing the rung for the duration of a flip protects against the 21:18 crash by disabling the thing the flip is waiting on. A wholesale gate cannot be right when the protected party and the requesting party are the same machinery. THERE WERE TWO STRANGLE POINTS, not one. The armed gate was the obvious half. The second was quieter: _nothing_resident() returning False on a parked extent dropped through to the unknown-refuse branch, so even with the gate removed the rung would still have declined. Both are fixed. THE REFINEMENT. The parked extent already carries exactly the information needed to be selective, which is why this is clean rather than a compromise: rows INSIDE it are what the flip is about to pack and may not be touched; every row ABOVE it is recomputable prefix and is precisely what the funding wants. So the extent PINS THE CEILING instead of closing the rung. _parked_ceiling() returns the highest parked row id, -1 when nothing is parked, and -2 for the one case that still refuses wholesale: an UNKNOWN extent while a flip is armed, where there is no boundary to name. Both call sites take req_max = max(req_max, parked), so the evictor is handed a ceiling that cannot reach into the extent. UNKNOWN, DECIDED AND DOCUMENTED as asked: it refuses ONLY while a flip is armed. Outside a flip there is nothing parked to protect, and closing the rung there is exactly the strangle this commit removes. sgl-project#746 (the exact arm-time snapshot) is what removes this last wholesale case; until then it is one narrow, named condition rather than the whole flip window. TEST RESULTS test_evict_rung_flip_park_744.py refined to the new semantics, 19 tests + 2 subtests green. The two tests that asserted the wholesale gate are REPLACED rather than deleted -- they now assert the funding path delivers -- and the crash protection gets its own explicit test that the evictor's ceiling never reaches into the parked extent. Mutation matrix, 4 mutants, ALL KILLED, covering both directions: N1 exclusion removed (parked ignored) -> 1 failed CRASH protection dies N2 exclusion covers everything (gate back) -> 1 failed FUNDING dies N3 collecting site ignores parked ceiling -> 2 failed N4 unknown-while-armed no longer refuses -> 3 failed N1 and N2 are the (c) pair: neither over- nor under-covering survives. Pinned suites green together, 142 passed + 87 subtests: sgl-project#744, sgl-project#731 + the #731x#744 interaction pin, both sgl-project#717 suites, sgl-project#656, sgl-project#684, sgl-project#662-f4, sgl-project#713 admission intake, sgl-project#739 prefill progress. Ruff and codespell clean. PP2 SECONDARY, checked and NOT claimed as fixed. The 0 MiB shrink against an unmoved driver-free column is an already-instrumented condition (kv_backing_relief.py:1915-1927): the code detects it, names retained handles (SGLANG_FLIP_SEAM_RETAIN_HANDLES) or a missing commit chunk as the cause, keeps the cap on and returns 0. The gate refinement lets the rung PRICE again; it cannot make a retained-handle arena PAY. Necessary, not sufficient -- the retain setting needs verifying at the boot, and this commit does not pretend to close it. No boot.

…ce in one harvest plan, with a proven runner Consolidated from the comp4 gate run (progress.662-F4-r5), WINDOW_TICKET_745/755, NOTE_747 par.8-9, NOTE_738/755, TICKET_727 and the operator ledger. Two hard incompatibilities shape the plan: WINDOW_TICKET_745 Arm 1 excludes the checkpoint interval that sgl-project#758's anchor-cadence observable requires (-> ARM I hicache harvest, ARM II = one flag more), and sgl-project#713's TTFT<3s needs a quiet router while every other loaded gate needs the soak backlog (-> sgl-project#713 is the idle sub-phase BEFORE the backlog, not a separate boot; the 06:44Z soak driver is the load source per the Lastprobe rule). ARM I phases: load-time (sgl-project#738 no-99G-plateau, file-backed-image reclaim), idle-quiet (sgl-project#713, health, corridor), loaded (Gates A/B/C, - sgl-project#757 race-holds, sgl-project#748 all three shapes, sgl-project#744/sgl-project#717 rung-funded flip, - sgl-project#690 refill census, sgl-project#758-2 mamba host resume, WT_745's three lines, corridor minima), teardown (image reclaim). ARM II adds --mamba-checkpoint-interval 8192 for sgl-project#758-1 anchor cadence + NOTE_747 par.8.1-8.3. SEPARATE windows named with reasons: sgl-project#727 four-boot A/B, WT_755 slots A/B (pool-geometry confound), sgl-project#755 metal retraction (mechanism not built -- nothing to measure), sgl-project#709, sgl-project#735 Step-2. sgl-project#602 and sgl-project#536/sgl-project#537 carried as HONEST unresolved slots (owner-held detail / not found with acceptance shape) rather than invented readouts. Runner run_window_ladder.sh: PASS/FAIL/UNOBS table from boot log + live server; never boots, never kills, never touches the soak driver; soak-tolerant by construction (loaded checks are log observations, the one latency check runs only in --phase idle, and choosing that phase IS the operator's quiet-router assertion). Missing-emitter cases (anchor cadence) report UNOBS, never FAIL -- absence of an instrument is not absence of the property. Mock-smoked per the desk rule, both directions: a fixture built from the comp4 specimen lines reproduces the real run's verdicts exactly (GATE-C crash, both sgl-project#748 shapes + vacuous relief, sgl-project#757 sentence = 5 FAIL; sgl-project#744/sgl-project#690 PASS; exit 1) and a clean fixture goes fully green incl. the ARM II cadence line (exit 0). bash -n clean, codespell clean. Nothing was booted.

…s a release granule The floor clamp silently undoes the granularity round-up, and nothing re-checked it: rows_wanted = max(rows_wanted, self._min_release_rows()) target = max(floor, current - rows_wanted) The round-up exists because a shrink smaller than one commit chunk per buffer clears no extent anywhere (the 2026-08-11 measurement in that comment). When the eviction floor binds, the clamp hands back a target whose distance from current is below one granule again -- and the shrink was attempted regardless. The cap engaged, decommit_range cleared no extent, and the rank lost capacity in exchange for nothing. MEASURED, boot_798_0822_0737.log: 15 occurrences of "reported 0 MiB but the driver's free column did not move". PP2's shape: current=126976, floor=88945, granule=229376. The round-up asks 229376; the clamp yields 88945; the real distance is 38031 rows, one sixth of a granule. This is the grant-vs-payment gap. The group's decision was already correct (51 GRANTED / 0 DECLINED); what failed was the payment. Note the granule can EXCEED the whole pool (229376 > 126976 here), so on such a rank no shrink can pay at all at --flip-seam-chunk-mib 8. That is a sizing question and is deliberately NOT papered over here: this guard's job is to stop paying a cap for a release that cannot happen, not to pretend it can. The sizing lever is filed separately. DIRECTION OF SAFETY: the guard only ever turns a shrink into NO shrink. It never deepens one, so it cannot pull backing below the highest live row -- the sgl-project#717 fault that reverted c4e5579 and killed boots. The can-fail test pins that direction explicitly. Correcting my own earlier exclusion: I ruled granularity out by comparing an ask of 107049 ROWS against "8 MiB" without converting between them. They are one axis -- _min_release_rows is ceil(chunk_bytes * buffers / bytes_per_row). The exclusion was apples to oranges and it was wrong. Tests: test_floor_clamp_defeats_granule_796.py, 3 tests, red first. Pre-existing failures in test_evict_rung_floor_invariant_717.py and test_kv_arena_*_631.py are untouched by this diff (verified: it touches neither _buffers nor _min_release_rows).

… a correctness invariant for a funding win THE PREVIOUS COMMIT SHIPPED A REGRESSION AND THE TREE CAUGHT IT. _floor_rows clamped the floor DOWN to this rank's backed rows whenever the floor exceeded them. It looked like the ordered fix -- "derive the per-rank floor from the per-rank cap; a floor >100% of its own cap is a defect by construction" -- and the arithmetic is indeed a defect. The clamp is not its remedy. The floor IS `live set + 1 + margin + admission reserve`. Lowering it to a smaller cap does not make an under-backed rank able to shrink; it authorises a cap BELOW rows that are still in use. Measured, not argued: test_residency_cap_flip_levelling_792 ::TheLevellingMustNotCapBelowTheLiveSet ::test_the_recovery_levelling_leaves_every_rank_able_to_pay with the clamp 1 failed without the clamp 12 passed The invariant is older than this ticket and the tree already carried it. My ID-level regression is what surfaced it, and it is exactly the danger direction I had written into my own mutant rationale one commit earlier -- and then implemented anyway. WHAT STAYS. The detection: floor_exceeds_local_cap() splits the two cases _floor_ppm collapsed -- floor == cap is a HEALTHY full pool, floor > cap is an under-backed rank -- and _floor_rows now LOGS that condition by name instead of letting it vanish into a group-wide freeze. The floor itself is untouched. WHAT THIS MEANS FOR THE VETO. It is not fixed here, and the honest statement is that it cannot be fixed locally at all. An under-backed rank genuinely has nothing to give; the repair is on the BACKING side (grow it) or in the agreement shape (stop agreeing a proportion of each rank's own cap). Both are enforcement and both need metal. Ticket with acceptance criteria in the COORD. MUTANT MA1 IS NOW THE CLAMP ITSELF, so the withdrawn version cannot creep back: reintroducing it turns the suite red. Added test_the_floor_never_drops_below_the_live_set so this file cannot drift toward it either. Also ran black over the touched files (the sgl-project#656 formatting ratchet). The ratchet's own two failures are PRE-EXISTING and not mine: the files it flags are phase_flip_boot.py, phase_flip_seam_reserve.py, kv_vmm_backing.py and two test files, none of which this branch touches. TEST RESULTS (hermetic, CUDA_VISIBLE_DEVICES="") test_funding_authority_770.py + test_floor_local_cap_812.py + test_residency_cap_flip_levelling_792.py 76 passed test/srt/mutants_770.sh 11/11 KILLED test/srt/mutants_812.sh 4/4 KILLED (+1 equivalent) ruff --select=F401,F821,UP037 All checks passed codespell clean PRE-EXISTING FAILURES CONFIRMED NOT MINE, each by running the same file against the base module 481c732: test_evict_rung_floor_invariant_717.py 3 failed on base AND on mine (AttributeError '_buffers' -- the sgl-project#717 stubs, unrelated to this branch) test_r1_private_constant_gate_584.py flags VERDICT_BAR_MIB, NEAR_OOM_MIB, BYTES_PER_GB, MIN_STAGING_GB -- all in files this branch never touches test_scheduler_chunked_req_gate.py 3, unrelated test_pp_proxy_cross_epoch_mispair_795.py 1, unrelated

Ying1123 added 2 commits July 24, 2024 07:13

fix bugs

55ba630

fix

7a258cc

Ying1123 merged commit 30d8e13 into main Jul 24, 2024

Ying1123 deleted the fix-bugs branch July 24, 2024 21:44

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

Improve benchmark scripts (sgl-project#717)

1db4fd3

efschu added a commit to efschu/htsglang that referenced this pull request Aug 17, 2026

Revert "[sgl-project#717] "Nothing resident" is not "unknown", and th…

b786858

…e rung read them the same" This reverts commit c4e5579.

Navigation Menu

Uh oh!

Improve benchmark scripts - #717

Improve benchmark scripts#717
Ying1123 merged 2 commits into
mainfrom
fix-bugs

Ying1123 commented Jul 24, 2024

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Conversation

Ying1123 commented Jul 24, 2024

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant