Skip to content

Update version to 0.1.15 - #431

Merged
merrymercy merged 1 commit into
mainfrom
update
May 12, 2024
Merged

merrymercy merged 1 commit into
mainfrom
update

Conversation

@merrymercy

Copy link
Copy Markdown
Contributor

No description provided.

@merrymercy
merrymercy merged commit 72bb344 into main May 12, 2024
@merrymercy
merrymercy deleted the update branch May 12, 2024 21:22
timethink pushed a commit to timethink/sglang that referenced this pull request Mar 9, 2025
Garrybest pushed a commit to Garrybest/sglang that referenced this pull request Jan 9, 2026
* add grok2 launching guide
HanHan009527 pushed a commit to HanHan009527/sglang that referenced this pull request Apr 7, 2026
EdwardXuy pushed a commit to shun8686/sglang that referenced this pull request May 9, 2026
efschu added a commit to efschu/htsglang that referenced this pull request Aug 3, 2026
… correct the §1.3 model, and make the sgl-project#431 guard non-blocking

Two commits:

test(sgl-project#483,sgl-project#459): sgl-project#483 ("check_regressions prices prefill on stale rates")
is REFUTED -- check_regressions already prices on the RESOLVED rates
(landed with sgl-project#475). What was actually missing was the PIN: the only test
calling the function is checkpoint-gated and skips in every hermetic run,
so a revert would be invisible. test_check_regressions_pricing_483.py is
the hermetic falsifier. Also makes the s12 harness honest: an absent
measurement is None, never 0.0 (a spec-off tick has no accept length, and
0.0 silently survives every arithmetic it enters); and any procedural
property a verdict rests on (warm-up discarded, draws back-to-back, gap
between draws) is recorded as a measured number in the artifact, not taken
on the harness's word -- the s12 warm-up draw ran for months and was
discarded silently, so a warmed-up point and a cold one wrote the same
file.

perf(sgl-project#517): names the actual decode-loop BAR1 collectives and corrects the
catalog's round model -- a decode round is NOT free of host-path
collectives (the belief the old model rested on). Per NEXTN round there
are 3 replay boundaries (the draft chain is one captured graph, not one
per step) and 5 host-path BAR1 broadcasts, all sgl-project#50 speculative
rank-agreement syncs, reaching barlink because a barlink boot never
constructs pynccl. Measured cost: the sgl-project#431 blocking status-word read costs
-9.22% decode TPS against the same-tree NCCL baseline; removing it
recovers +2.68%. Fix: the abort-guard status word is now read
asynchronously (non-blocking D2H + cudaEventQuery, returning what an
earlier check staged), trading reporting LATENCY for zero stream
synchronization per check -- detection is unaffected because ctlStatus is
sticky. SGLANG_BARLINK_BAR1_ABORT_MAX_LAG (default 4) bounds the
deferral with a can-fail proof; ..._ABORT_DEFER=0 restores the exact
pre-sgl-project#517 blocking behavior. NOTE_517_bar1_guard_desk.md has the full
model, desk-only (no GPU arm run yet -- ticketed).

Two real conflicts, both resolved by combining rather than choosing a
side:
- FEATURE_CATALOG.md §12 (same insertion-point pattern as every merge
  today): HEAD's "Verified carve-outs (sgl-project#514)" paragraph kept, this
  branch's two sgl-project#459 paragraphs ("absent measurement is None" /
  "procedural property must be in the artifact") appended after. Two
  OTHER hunks (the sgl-project#483-REFUTED note near §2, the "Guard cost" note near
  §7) landed with no conflict at all -- neither region had been touched
  since this branch's base.
- test/registered/unit/distributed/_bar1_marker_source.py: both HEAD (via
  sgl-project#514's earlier edits) and this branch carry DIFFERENT LINE_* pins into
  barlink.py/barlink_bar1.py, because each was written against its own
  version of those files. Neither side's numbers are correct once both
  sets of source edits are combined. Resolved by reading the ACTUAL
  post-merge source: grepped the real logger.info/logger.warning/raise
  RuntimeError call-start lines in the merged barlink.py and
  barlink_bar1.py (2236/2250/3207/810) and verified all four renderers
  execute correctly against them before committing -- not a textual
  merge of the two guesses.

All 17 catalog sections and every other additive paragraph verified
present.

Tests (CUDA_VISIBLE_DEVICES=99, PYTHONPATH pinned to this worktree):
- The five directly touched/new test files (_bar1_marker_source.py's
  renderers verified live against the real merged source, plus
  test_barlink_bar1_abort_431.py, test_barlink_bar1_abort_deferred_517.py,
  test_s12_harness_459.py, test_check_regressions_pricing_483.py): 67 + 7
  passed, 0 failed.
- test_s12_log_analyse.py specifically (s12_log_analyse.py is touched):
  same 4 pre-existing failures, unchanged (fixture-data-dependent, not
  source-code-dependent).
- test/registered/unit/distributed/ (full): 24 failed, 2206 passed,
  8 skipped, 723 subtests -- same 24 failing IDs as the established
  baseline, byte-identical.
- ruff check --select=F401,F821,UP037: clean.
- codespell: 7 hits, all confirmed pre-existing (a German prose sample
  string already in scripts/gpu_battery/s12_prefill_kurve.py and
  s14_decode_punkt.py at this branch's base, untouched by its diff).
efschu pushed a commit to efschu/htsglang that referenced this pull request Aug 8, 2026
…e, never __trap

The last live instance of the sgl-project#583 context-killing mechanism: wait_ge in
the host-staged transport still called __trap() on spin expiry, which
destroys the CUDA context and surfaces as sticky 'unspecified launch
failure' at whatever unrelated call site runs next. Unreachable in every
production configuration (host is explicitly-selected, no-fallback,
bar1 pinned) — but a landmine for any host-transport measurement boot.

Ported the device transport's sgl-project#583 fix honestly, not minimally:
- wait_ge returns bool and writes a per-site code into the abort word on
  expiry (put reuse-guard=1, reduce peer-wait=2, copyout wait=3); the
  BARLINK_HOST_TRAP macro and its hipify plumbing are deleted.
- Block-uniform exit: shared abortS written by the spinning thread 0,
  __syncthreads(), all threads return together before any payload loop,
  fence, or last_block barrier — verified at all three sites. Terminal-
  abort contract documented: cross-block skew and stale blk_ctr are
  acceptable residue of a process the host is about to kill by raise.
- The transport's THREE sequence-counter families (collective, per-peer
  send, per-peer recv) are unified into one zero-initialized backing
  tensor with 2-word rows — the spec's naive grow-to-2 would have made
  the p2p kernels' abort store corrupt a NEIGHBOR's sequence number, and
  the unified shape makes check_aborted a single read over every abort
  word the transport can produce.
- HostCollectiveAborted mirrors DeviceCollectiveAborted (rank/world/
  code/where, garbage-results contract) and appends the sgl-project#650 peer
  statement. check_aborted is a plain in-line read (documented choice:
  this is a measurement vehicle, not the bs=1 hot path) wired twice:
  the duck-typed _after_transport seam (sgl-project#431) for eager collectives,
  and abort-gate registration for replay-only collectives — host is
  capturable, so a graph-resident collective never reaches the seam
  (the same gap sgl-project#583 closed for the device transport).
- The load-bearing reachability claim is pinned: 'host' is in the
  no-fallback set via the real parallel_state lookup, and the dispatch
  seam is proven transport-stable mid-run — an explicit
  SGLANG_BARLINK_TRANSPORT=host boot is the only way this code runs.

Falsifiers, red-first (executed against the pre-fix source inside the
test): old wait_ge provably trapped and returned void, old seq tensor
had no abort-word room; fixed source has zero traps, bool returns, all
three consuming sites, distinct codes. nvcc compile smoke rc=0 (sm_86).
GPU-gated on-card falsifier written and documented (put-kernel forced
timeout must set the word, leave the context ALIVE, and raise named)
— skips at the desk by design.

Counting basis: 35 passed + 1 GPU-skip + 19 subtests in the new file;
full battery with abort/stub-drift/sentinel regression suites 81 passed
+ 1 skipped + 19 subtests (97.9 s, my rerun); ruff and imports clean.
efschu added a commit to efschu/htsglang that referenced this pull request Aug 9, 2026
…trol message

Reverts the request-injection vehicle entirely. Injection put a SECOND
blocking channel -- a control message on the PP point-to-point chain --
alongside the flip's group reduction, with no global order between them,
and it deadlocked twice on metal:

  async forward, arm in the same pass:
    rank0 on_round -> bounded_collective, ranks1-2 -> _pull_raw_reqs
  synchronous forward (the "fix"):
    rank0 -> point_to_point_pyobj send, ranks1-2 -> hidden-states recv

Both are instances of the fork's standing collective family (sgl-project#431, sgl-project#616,
sgl-project#639): two independent blocking channels with no order between them.
Patching the ordering case by case (defer-by-one was the third attempt)
fights the pattern; removing the channel ends it by construction. A
control message riding the PP data chain can always be scheduled against
a rank that is currently blocked in the other channel.

The channel was never needed, because THE SIGNAL IS ALREADY REPLICATED.
Under PP every stage receives the same request stream over the chain and
runs the same scheduling decisions on it, so waiting_queue and
running_batch hold the same values on every rank -- visible directly in
the scheduler's log, where the three stages print identical
#running-req, #queue-req and #pending-token at the same timestamp. decide
is a pure function of those values, so every rank computes an identical
verdict and arms ITSELF via the same arm_phase_flip the RPC path calls.
Zero messages, no delivery, nothing to order.

The existing machinery then does what it was built for: ranks arm at
slightly different rounds, an armed rank enters the reduction only once
locally PARKED (owing no pipeline send), and peers converge on their own
arm+drain -- "MIN-skew is legal" (PhaseFlipRuntime.on_round). Skew is
absorbed; only DISAGREEMENT is fatal, and disagreement cannot arise from
identical inputs. The reduction's equality family (epoch, config_fp,
direction once armed) is the standing loud check on that assumption.

Evaluated from _phase_flip_on_round, immediately before the consensus
round that carries the arm, on every rank and in both loop families.
A refused arm releases the dwell clock so the policy can re-decide
rather than sit out a window it never used.

request_receiver.py and scheduler_pp_mixin.py are byte-identical to
their pre-policy state again; a test asserts that by source inspection,
so neither can regrow a policy hook or an internal-arm special case.

Tests: 313 passed in the sgl-project#631 flip family, CPU-only. New: three ranks
with identical replicated inputs produce identical arms; the policy
returns no request; a refusal releases the dwell.

Operational acceptance under load NOT yet achieved -- boot with this
design is running.
efschu pushed a commit to efschu/htsglang that referenced this pull request Aug 17, 2026
…c, cross-ref sgl-project#726

Docs only; no production file touched. Closes the three items my own retro-sweep
left open.

## 1. EVERY=32 (former section 19.3): NOT DEAD, but path-dependent

Traced the constant to its consumers rather than trusting either ledger entry.

  check_every() returns 1 when the env var is unset (barlink_abort_gate.py:
    189-197). The default is ONE, not 32 -- "=32" is only meaningful if
    somebody sets it.
  barlink_bar1.py:5177 reaches it UNCONDITIONALLY, after the
    abort_check_enabled / pending-launch gates. LIVE on bar1.
  barlink_device.py:1562 sits in the in-line fallback path BELOW sgl-project#517's
    _abort_poll_active early return, so on the device transport it is bypassed
    entirely whenever the watchdog feeds the flag.

BOTH LEDGER READINGS WERE PARTIALLY RIGHT AND NEITHER WAS COMPLETE. "Dead since
sgl-project#517" is true of the DEVICE path only -- that is exactly what phase 2 made
cheap -- while bar1 never got that treatment and still consults the knob.
NOTE_517 itself treated ..._EVERY as arm B3, expected to buy nothing on top of
the staged read; "buys nothing measurable" is not "the consumer was removed",
and that elision is how the two readings drifted apart. The sgl-project#431 utilisation
figure (72-84% -> 86-91%) predates sgl-project#517 and was never re-measured after it.

THE RECIPE QUESTION, answered: the production boot recipe does NOT set it --
absent from startkommandos-rig.md (grep -c = 0). So there is no dead knob in the
recipe misleading the next operator, which was the specific worry.

Worth recording: this is the SAME bar1-vs-device asymmetry I found independently
in sgl-project#673, where bar1's check_aborted never consults _abort_poll_active while the
device path short-circuits on it. Two investigations converged on one structural
fact about these two transports.

Section 19.3 is now empty and points at the settled 19.1 entry.

## 2. Stale spill-matrix doc (section 19.4): corrected at both sites

ANALYSE_spill_matrix_20260804.md claimed kvso x HiCache is mutually exclusive
and that "kvso cannot run on the production recipe at all". sgl-project#550 superseded
that: the combination is OPT-IN, gated on KVSO_ALLOW_HICACHE
(server_args.py:7502, `if os.environ.get("KVSO_ALLOW_HICACHE", "0") != "1":
raise`). Per the refusal's own text the two host pools are independent objects
with DISJOINT KEY SPACES and their pinned host RAM is summed by ONE JOINT BUDGET
GUARD instead of each validating alone; what remains is a measurement
(spill-copy vs prefetch contention), not a mechanism.

Fixed in section S1 and in the H15 matrix row, with the original text kept
inline so the change is legible rather than silently rewritten.

ANCHOR CORRECTION: the gate is at server_args.py:7502, NOT the :7385-7395 given
in the brief. Cited by symbol per the anchor rule -- the sgl-project#621 sweep found audit
line numbers had drifted, two of three into the wrong file.

## 3. sgl-project#726 cross-reference added to section 19.1

The INT8-KV IMMA-QK builder must CREATE the dtype surface, not find it: choices
at server_args.py:1026 carry no int8, the fp8 KV scale path is hard-coded
per-tensor, and the one per-group-scale precedent (MHATokenToKVPoolFP4)
dequantises eagerly before any backend sees it -- copying it buys VRAM savings
but NOT the bandwidth savings an int8-KV lane exists for. Budget the plumbing as
new work. The anchor is VERIFIED; the FP4-precedent detail is marked
UNVERIFIED-LEDGER (relayed from ANALYSE_726, not re-read here).

PRIOR-ART GATE honoured: grepped FEATURE_CATALOG, docs/dev and git log --all
--grep before each change; both absence claims carry the file:line of the
refusing gate (KVSO_ALLOW_HICACHE at :7502, kv_cache_dtype choices at :1026).
efschu pushed a commit to efschu/htsglang that referenced this pull request Aug 25, 2026
…urvive

`poll_status_words` caught every exception so the watchdog thread could not
die. That is right for a poll that hiccups and WRONG for a CUDA illegal memory
access: once a context takes one it is unusable, and every later CUDA call in
the process raises the SAME error at whatever site happens to run next.
Swallowing it did not keep serving alive. It only decided that the crash would
be reported somewhere innocent.

THE SPECIMEN. W40, boot_w40_857strict_0825_2113.log, the CUDA_LAUNCH_BLOCKING=1
arm. First CUDA fault in the whole log, 21:17:13, immediately after
"SEAM DRAIN tp_to_pp: device-tier streams quiesced in 0.0 ms at the no-return
point":

    barlink_abort_gate.py:351 poll_status_words -> barlink_bar1.py:4937/4948
      self._abort_poll_dst.copy_(self._ctl_dev[0:1], non_blocking=True)
    torch.AcceleratorError: CUDA error: an illegal memory access was encountered

Logged as "barlink-BAR1 status poll failed" and continued past. Downstream, in
the same second, the scheduler died in `get_cpu_copy` inside the seam CAPTURE;
the boot before it (..._2107.log, no CLB) died in `load_cpu_copy` inside the
seam RESTORE. Three sites, one fault, two of them innocent -- and that cost this
shift two wrong roots before CUDA_LAUNCH_BLOCKING separated them.

THE CLASS: an exception handler that assumes its failure is survivable. It could
not tell a hiccup from a context kill because it caught bare `Exception`.

THIS MODULE ALREADY OWNED THE CLASS, which is why the fix belongs here. Its own
docstring opens on sgl-project#431: a run that tripped the cap on essentially every
collective produced an abort file with ZERO matching lines, so "nothing tripped"
and "everything tripped" were indistinguishable. The handler below it had made
"a poll hiccuped" and "the CUDA context is dead" indistinguishable in exactly
the same way, inside the module written to end that.

WHAT CHANGED
  is_poison_error()   classifies by MESSAGE, not by class, and that is the whole
    point: torch raises AcceleratorError for a recoverable OOM and for an
    illegal access alike, so isinstance would over-refuse. Markers are the
    context-killing cudaErrors: illegal memory access, illegal instruction,
    misaligned address, unspecified launch failure, device-side assert.
  record_poison()/poison_record()  FIRST wins, not latest. The record exists to
    name the ORIGIN, and in a poisoned process every later call reports the same
    error -- a record that tracked the newest would name the innocent site,
    which is the defect itself.
  poll_status_words   on a poison-class fault: record, log ERROR saying in words
    that the context is unusable and that THIS is the origin rather than the
    traceback that lands next, and BREAK out of the round. The watchdog still
    does not raise -- raising kills the thread and loses the record -- but it
    stops reading devices, because every further read returns the same error and
    multiplies the misattribution.

FOURTH INSTANCE OF ONE PATTERN, and the pattern is now worth more than any
single fix: WHERE IT SURFACED IS NOT WHERE IT ORIGINATED. (1) W40b blamed
restore_seam_state, actually the offload side. (2) the seam emitter sat AFTER
the dangerous call, so the one event it existed to explain was the one it could
not witness (#783b). (3) W40 arm 1 without CLB blamed the trailing
synchronize(). (4) this. CUDA_LAUNCH_BLOCKING=1 in one standing acceptance arm
is what separates them; without it every IMA is misattributed by construction.

PRIOR ART, CHECKED BEFORE BUILDING
  pause_polling() (:309) EXISTS and excludes the watchdog's device reads for a
    scoped window; one caller, parallel_state.graph_capture:2888, scoped by its
    docstring to CUDA-graph capture. THE PHASE FLIP DOES NOT TAKE IT. Filed, not
    used -- see below.
  hicache_phase_binding's generation stamping (BindingState.advance,
    host_pool_at, write_back_stamp_is_current:527) is the reusable shape for a
    phase-generation stamp IF the mapping question resolves that way. NOT built
    here: a stamp must not ship on a guess about which mapping moves.
  poison-class handling: DOES NOT EXIST anywhere in
    srt/distributed/device_communicators/. The only matches for "illegal memory
    access" are four comments in pynccl.py:137,191,238,263 about avoiding one.

STILL OPEN, DELIBERATELY, and nothing here assumes an answer:
  * WHICH POINTER IS FREED UNDERNEATH. `_ctl_dev` is allocated in `_build_up`
    (barlink_bar1.py:2596) and nulled in `__init__` (:1765) and `close()`
    (:5794), where vmm_free runs on _own/_own_flag BEFORE the null -- but
    `poll_status_word` early-returns when `_ctl_dev is None`, so a COMPLETED
    close() cannot produce this fault. Either a race inside close(), or a
    different mapping entirely. Not established.
  * WHETHER THE POLL IS THE ORIGIN OR MERELY THE FIRST REPORTER. Earliest in
    time is evidence for origin, not proof. Falsifiable: if it is only
    reporting, pausing polling across the cutover MOVES the crash rather than
    removing it. That is the next arm, and it is exactly why pause_polling was
    filed instead of applied -- applying it now would silence a messenger whose
    guilt is unproven.
  * W40a's nine successful restores versus W40b's first failing one. Narrowed,
    not closed: with #783b's range check never firing (guard_refusals=0 on
    metal) and the earliest fault in barlink, the restore was probably never the
    origin in either boot. "Probably" is not a finding.

TESTS  test_abort_gate_poison_867, 10 passed. Hermetic, no CUDA: the poison is a
real torch.AcceleratorError carrying the specimen's message, raised by a stub
transport. A/B on this tip: 5 failed / 5 passed WITHOUT the change, 10 passed
WITH it -- and the 5 that pass either way are the controls, deliberately routed
through getattr shims so they RUN on unpatched code instead of erroring. An
all-red file would have proved nothing about which behaviour this change owns.
Both directions of every judgement can fail: the classifier must match the
specimen text AND must NOT match a recoverable OOM; polling must stop after a
poison AND must continue after a hiccup.

Hermetic suites, run SEPARATELY per sgl-project#749 against fc24204 + 088d5d6, count
gate (^FAILED + ^SUBFAILED + ^ERROR == summary) PASSED on all three before any
comparison:

  distributed  95 failed / 3027 passed / 1170 subtests   gate 92+3+0 == 95
               baseline 3017 passed; +10 is exactly this commit's tests.
               Summary extracted BY PATTERN, not tail -1: these arms print UCX
               teardown warnings after pytest's summary.
  mem_cache    2 failed / 1774 passed / 361 subtests      gate 2 == 2
  managers     15 failed / 4343 passed / 356 subtests     gate 15 == 15

Zero new failures on any arm; every failure identical to baseline by count and
by name (arena_high_water_631 x7, phase_flip_rotation_wiring_809 x4,
restore_never_rebuild_677 x4, acceptance_emitters_758 x2 -- all
"RuntimeError: No CUDA GPUs are available" in a hermetic run).

ruff: barlink_abort_gate.py 0 before and 0 after; the new test carries 4 x E402
from the house register_cpu_ci()-between-imports convention. codespell clean.
No line over 88 columns. black/isort deliberately NOT run over these files: the
venv's black is older than the one this tree was formatted with.

NOT PROVEN HERE: that this removes the W40 crash. It removes the MISATTRIBUTION.
The acceptance boot is what says whether the fault itself is still there, and
now its traceback will point at the origin instead of at the next innocent
caller.
efschu pushed a commit to efschu/htsglang that referenced this pull request Aug 31, 2026
… while the peers spin

BOOT-PROVEN (boot 22, boot_855_1033b_0840f82601_0831_131955, 13:24-13:26).
The first TP forward after a pp_to_tp cutover touches a shape specialisation
never loaded in this process. Rank 0's crash-time stack (log line 53248 ff.):

  _dcp_write_scatter (flashinfer_backend.py:2574)
    run (triton/runtime/jit.py:743) -> _init_handles (compiler.py:466)
      loadBinary -> cuModuleLoadData (libcuda.so) + 11 libcuda frames

NOT the compiler -- cuModuleLoadData. The module is built; LOADING it into the
CUDA context blocks, because the load needs the device and the device is
saturated by the peers' barlink BAR1 spin kernels, which are waiting in
all_gather for this very rank. The cycle closes and the spin deadline fires:
Bar1CollectiveAborted (ranks 1/3 and 2/3, group flip_dcp:0) -> SIGQUIT.

WHY THE EXISTING MECHANISM MISSED IT, measured: 352 build-window lines in that
boot, ZERO in 13:24-13:26, all 352 carrying 'full cuda-graph capture warmup'.
cold_build_window had exactly three production callers (barlink BAR1 build,
sampling warmup #603b, capture warmup) and none is on the path a cutover
re-dispatches into. sgl-project#640 on a path sgl-project#615 never saw. The fix makes that path a
caller rather than building a second mechanism beside it.

A STALE COMMENT NEARLY REFUTED THIS FIX, and the correction ships with it.
sampler_warmup.py said wrapping a lazy build in cold_build_window "does NOT
work ... the window is PROCESS-LOCAL". True when written (8bddb93,
2026-08-06); falsified ONE DAY LATER by sgl-project#615 (38ec4fb, 2026-08-07), which
hooked publication into cold_build_window so "every existing call site
therefore becomes group-visible without moving". Never revised, both ancestors
of this pin, verified with git log -S rather than assumed. The paragraph is
annotated in place rather than deleted, because it is still right about ITS
module (warm-at-boot + barrier REMOVES the race; a window only EXTENDS it).

WHY A WINDOW OVER THE REAL FORWARDS, NOT AN ENUMERATED WARM SET. The warm set
is not knowable by inspection -- it varies with direction, spec-decode, the
sgl-project#887 one-chunk grant and any backend swap -- and missing one member reproduces
the wedge exactly, at the next first-loader. Whatever loads, loads under the
window. Residual stated in the code: this EXTENDS deadlines (900 s cap) rather
than removing the race; boot 22's stall was ~150 s, so the cap is not binding.

FIRST-LOADER CENSUS, as a standing table: devtools/CENSUS_1033c_first_loaders.md
-- every first-loader site x triggering state change x covering window x
collective proximity. Rows 1-6 verified at file:line (row 5 on metal); rows 7
(runtime recapture / drafter switch) and 8 (rung change sgl-project#704, resume-restore
sgl-project#89) are written UNVERIFIED, so the table is a lower bound on the covered set
and never a proof that nothing else is exposed. It also names the inverse
column (armed windows that may cover paths that no longer first-load -- the
352 are a count of WARMUPS, not of BUILDS, and the window carries no
modules-loaded counter, so live and dead coverage are indistinguishable from
the log) and answers the upstream-minimal question: there are not four window
mechanisms but two plus a front door, with one site (barlink_device.py:865)
bypassing the front door. Named, not rebuilt.

#1033d, same boot, INDEPENDENT DEFECT, minimal hardening only: the prefetch IO
aux thread caught only Empty, so a page whose geometry did not match the
incoming binding (mha.py:556 reshape, '[2,16,1,4,256]'=32768 against a 16384
page) ended the thread three seconds before the cutover -- killing no process,
setting no exit code, appearing in no health probe, and leaving storage
prefetch dead for the rest of the boot. It now fails the OPERATION (host slots
released, so the requester gets a refusal instead of an unreachable
completion), logs loudly with a counter, and stays alive. The broad except is
defensible only because the alternative is silent thread death with no
supervisor above the loop; the underlying two-geometry host pool across a flip
(sgl-project#718/sgl-project#719/sgl-project#875 family, and the fork's own 'sgl-project#939 RE-HOME VIA RE-READ ... source
page 16384 elems vs destination page 32768 elems' line names the same mismatch
two lines earlier while handling it correctly) is NOT fixed here and is its own
posten.

DESK CHECKS, matched to each edit's failure class:
* devtools/check_1033c_cutover_window.py -- 19 cases, hermetic, driving the real
  run_batch wrapper: RED-FIRST arm reproducing the boot-22 condition (counter
  unarmed -> no window -> nothing published), first forward covered AND
  published, window closes (the sgl-project#431 open-without-close shape), budget finite
  and steady state byte-identical, window closed and budget spent even when the
  forward raises, and an AST check that the arming sits under no `if` (the naive
  grep form of that assertion failed on a COMMENT -- fixed to ask the AST).
* AST check on the aux thread: broad handler present, logs at ERROR, releases
  the failed operation's host slots, and continues.
* Method-split check: wrapper delegates, window only on the armed path, lazy
  import off the fast path, inner body intact.
ruff F401/F821/UP037 clean.

DESK-PROVEN. Metal proof is boot 23.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant