Raise errors for prompts that are too long - #436
Merged
Merged
Conversation
merrymercy
changed the title
timethink
pushed a commit
to timethink/sglang
that referenced
this pull request
Mar 9, 2025
efschu
added a commit
to efschu/htsglang
that referenced
this pull request
Aug 2, 2026
…r on this rig The transfer_kv_all_layer_direct_lf_pf segfault (cu12/cu13 cudaMemcpyBatchAsync ABI split) that blocked page_first_direct host-tier transfers is fixed by the cu13 sgl_kernel rebuild merged in chore/436-kernel-cu13.
efschu
added a commit
to efschu/htsglang
that referenced
this pull request
Aug 3, 2026
…ash, not installed The kernels are only reachable through a rebuilt sgl-kernel, and the wheel is pinned separately from the source. This records the artifact so the install is a documented step rather than a rediscovery: /spinning/wt-398-wheel/sglang_kernel-0.4.4-cp310-abi3-linux_x86_64.whl sha256 67f03cfa755efa01498c7732bd6ae015ec5673feffe9a51452fefdbe0dcd4664 16 638 372 B, 39 sgl_kernel files, dist name sglang-kernel (the fork dist -- the sgl-project#384 two-dist situation is unchanged and sgl-kernel 0.3.21 must stay uninstalled), objdump shows libcudart/libcublas/libcublasLt .so.13 only. Build knobs copied verbatim from sgl-project#436 (cu13 nvcc, archs 86;120, SKIP_SM90_VARIANT=ON, ENABLE_FA3=OFF, COMPILE_THREADS=1, MAX_JOBS=4), empty dedicated ccache, CPU-only (CUDA_VISIBLE_DEVICES=99). Two runs: the second is an incremental rebuild of the gguf TU after a late comment edit, so the wheel corresponds to the committed source rather than to an intermediate state. DELIBERATELY NOT INSTALLED: another session's pytest mapped sgl_kernel in /proc/*/maps throughout, and this section's own rule is to swap those files only while the venv is quiet. Until it is installed the tree behaves exactly as before sgl-project#398 -- the capability probe answers False, MXFP4 stays out of the type sets, the Q5_0 repack carries it. Marker op verified in the binary the only way a stripped .so allows: strings finds _Z17ggml_mxfp4_nativev, the schema "ggml_mxfp4_native() -> int" and the plain name, 3 occurrences, the same count as the ggml_mmvq_kq_tuned control. Also fixes the IdentityMap snippet in TICKET_398 to the real API (registry.nvml.identity_map, not IdentityMap().entries()).
efschu
pushed a commit
to efschu/htsglang
that referenced
this pull request
Aug 4, 2026
…che observability gap Records the boot recipe decisions (env-driven storage directory, explicit 24 GB host L2 pin over the derived 2.0 ratio, page_first rather than the sgl-project#436-blocked page_first_direct) and the validation plan. Notes the gap that forced the indirect activation proof: the hierarchical cache emits NO Prometheus series whatsoever -- no hit, miss, write or eviction counters and no tier-bytes gauge, only two log lines. L3 activity is therefore not separable from the radix-level cache_hit_rate, and disk capacity and eviction are unmeasured. The kvso spill tier already exposes spill_tier_used_bytes / spill_tier_total_bytes, so the pattern to copy is in tree. Follow-up task candidate. Also records that cgroup memory.stat is the only trustworthy RAM source here: lxcfs reports 120 GB against a real 98 GB, so any guard reading psutil.available is blind in this container.
efschu
added a commit
to efschu/htsglang
that referenced
this pull request
Aug 17, 2026
…riant, refuse both guard flips (a) ROOT, at file:line. transfer_item_warp moves bytes with 64-bit PTX -- ld.global.nc.b64 / st.global.cg.b64 (transfer.cu:29-30) -- which faults on a misaligned address. The only alignment guard in the path is TORCH_CHECK(item_size % 8 == 0) at transfer.cu:285. That guard is sufficient for the lf_pf sibling and NOT for lf_ph, and the difference is entirely in the offset arithmetic. I compared the two launcher calls line by line (:496-513 against :532-551): identical empty tensors, identical tables in identical slots, identical layout dims. But get_global_offset_pf builds offsets only as multiples of item_size and layout_dim, so alignment is inherited from the checked guard, while get_global_offset_ph (transfer.cu:106-119) SUBDIVIDES by head_num in three of its four terms and the page-head kernel copies head_size_bytes = item_size / head_num (transfer.cu:145). Nothing requires that quotient to be 8-byte aligned. MISSING INVARIANT: item_size % (8 * head_num) == 0. Smallest faulting shape head_num=2 head_dim=2 fp16 -- item_size=8 passes the guard, head_size=4 does not, 24 enumerated offsets misalign. SCOPE, STATED RATHER THAN STRETCHED: this does NOT explain the reported segfault. test_minimax_sparse_pool_host_unit's shapes (page_size=4, float32, head_num=1, head_dim=2) give head_size=8 and zero misaligned offsets, and so do this rig's serving shapes (head_num=4, head_dim=256, fp8 -> head_size=256). The defect is real and latent; no production shape here is exposed. Both facts are PINNED so the attribution cannot quietly widen. Three hypotheses checked and closed on the way: the layer pointer table is correctly a device uint64 tensor (mha.py:135-144); layout_dim is the right per-token all-layer stride for the ph formula (mha.py:182); the k-only pool cannot reach lf_ph, it raises for any non-page_first layout under the kernel backend (mha.py:1113-1118). Attribution needs metal, so the falsifier is FILED NOT RUN: tools/441/falsify_lf_ph_441.py, three arms (repro / alignment / bisect), one arm per invocation because each is expected to kill the process, to be run under a gpu-arb claim. (b) BOTH GUARDS: DO NOT FLIP. The brief's premise was that the sgl-project#436 rebuild made them obsolete. Checking them at the code says otherwise, for two reasons, neither of which is the ABI issue. test_kvcacheio.py:20-24 module skip claims "segfaults on CUDA 13.x". But the lf_ph segfault reproduces on BOTH wheels -- that is this ticket's own premise and why it is not sgl-project#436. So the reason is wrong or incomplete, and flipping would expose a live cu12 crash. It must stay until (a) is fixed; what should change now is the REASON STRING, because a wrong reason sends the next reader to the wrong wheel. _DIRECT_PF_BATCHCOPY_BROKEN_CUDA13 never guarded a wheel issue at all. Per the measured matrix above it, it guards a TEST-SHAPE violation of the cudaMemcpyBatchAsync contract: pageable+default stream fails, pinned+side stream is OK, and production pins (mha.py:97) and uses a side stream (cache_controller.py:276,:742-749). The file's own comment already says unskipping alone leaves it red. The rebuild landing does not change a contract the test violates by construction; the fix is a test-shape change, and verifying it is GPU work. PIN TESTS NOT WRITTEN, deliberately: a pin that "goes red if the breakage returns" presupposes the breakage is gone, and for both guards it is not. Writing them now would encode a green state that does not exist. NOT DELIVERED: the sgl-project#261-Gate short-run script. Preparing it turnkey needs the gate's shim details verified at the code, which I have not done, and shipping a script I cannot stand behind would be worse than saying so. Tests, hermetic (CUDA_VISIBLE_DEVICES="", no CUDA context, no launch, no GPU): test_page_head_offset_alignment_441.py 6 passed + 6 subtests in 0.07 s pure arithmetic transcription of get_global_offset_ph; states the proposed guard positively and fails if that predicate is not the right one test/registered/unit/mem_cache/ 940 failed / 796 passed -- the 940 are PRE-EXISTING ON THIS BASE (b786858 lacks the sgl-project#585 conftest that turns no-accelerator errors into skips). Proven by re-running with my file ignored: 940 failed / 790 passed, i.e. identical failures and exactly my 6 passes added. ruff clean; falsifier byte-compiles without executing.
efschu
added a commit
to efschu/htsglang
that referenced
this pull request
Aug 17, 2026
…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.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.