fix: set ulimit -n 65535 - #647
Merged
Merged
Conversation
Ying1123
requested changes
Jul 18, 2024
Contributor
Author
|
I'll merge the main branch. |
Contributor
Author
|
done cc @Ying1123 |
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 8, 2026
…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.
efschu
added a commit
to efschu/htsglang
that referenced
this pull request
Aug 8, 2026
…recting rev1's verdict
Rev1 concluded the bring-up "cannot run on this laptop". That was right about
THIS TREE and wrong about the MACHINE, because it was written before anyone
looked at the laptop. A read-only sweep of it found a complete working setup
left by a prior session in /root/lh/, unrecorded anywhere.
WHAT ACTUALLY RUNS: Qwen3.6-35B-A3B Q4_K_M GGUF serves on the Radeon 780M iGPU
through htsglang at ~12.3-12.8 tok/s decode and ~157-172 tok/s prefill, bs=1,
ctx 8192, TP=1, no speculation. Eight logs reach "server is fired up"; q4_d/e/f
ran clean and answered POST /generate 200 OK, ended only by operator SIGTERM.
Weight load 74.33 s, 21.38 GB resident, 2.86 GB left for KV at ctx 8192.
That measurement VALIDATES the rev1 budget model: predicted room under the
25600-MiB ceiling was 3137 MiB, measured free-after-weights 3.38 GB. Section 4
is retitled accordingly and can now be trusted for planning.
HOW, given section 2 is still entirely correct: the gap was closed OUTSIDE the
repository. torch 2.10.0+rocm7.0 (hip 7.0.51831, cuda None), sglang editable
from /root/lh/sglang_src, no sgl_kernel at all, and a purpose-built standalone
sglang_gguf_rocm extension built for gfx1100 exporting every op the GGUF path
needs plus ggml_mmvq_kq_tuned and ggml_mxfp4_native, which have no in-tree
counterpart. The modified gguf.py imports it at :99 behind an `elif _is_hip:`
arm at :75 -- precisely the binding rev1 section 7 listed as work to be done.
HSA_OVERRIDE_GFX_VERSION=11.0.0 is load-bearing and verified: without it a
64x64 matmul fails hipErrorInvalidDeviceFunction.
The decisive schema check was run and its OBSERVED output now replaces rev1's
"Expect False" in section 2.6: sgl_kernel is not installed (it cannot be -- the
gfx gate), so the in-tree question is moot on this machine; the ops come from
the out-of-tree extension instead.
THE TWO REAL PROBLEMS NOW:
1. Speculation hard-faults the HIP context. NEXTN boots, serves a batch, then
dies with "unspecified launch failure" within 10-40 s -- twice, at ctx 2048
and 8192, against a same-config no-spec control four minutes apart that ran
clean. The crash site (scheduler.py:5206, a .to("cpu")) is only where an
async fault surfaces; AMD_SERIALIZE_KERNEL=3 was never run to localize it.
Prime suspect: the draft path reusing a GGUF op validated only on the target
path. Two earlier spec bugs were already fixed:
--speculative-draft-model-quantization gguf is REQUIRED (without it the sgl-project#647
pair of section 9.1 shows up live as 2 unloaded MTP parameters), and at ctx
8192 the spec arm cannot fit its mamba cache (max_mamba_cache_size=0).
2. The working code is unversioned. /root/lh/sglang_src is NOT a git repo and
has no baked commit id. The port that makes this machine run exists in one
unversioned copy on one laptop. This is now the largest risk in the ticket.
NEW CORRECTNESS FINDING, section 1.5.4: Q6_K dequant is non-deterministically
WRONG on gfx1103 -- eight runs on one fixed input differ, worst max|d| 5.8e-01
vs the numpy reference, up to 75 non-finite of 262144, wrong on the first call
in a fresh process, while Q4_K and Q5_K are byte-identical across the same runs.
Eight hypotheses recorded as falsified, including that it is an HSA_OVERRIDE
artefact (a native gfx1103 build is worse -- Q5_K becomes affected too).
Contained, not fixed: ggml_mul_mat_a8 is validated correct for Q6_K on gfx1103
(max|d| 5.5e-04) so Q6_K is pinned to MMQ, and the load path is rescued by an
exact one-time CPU dequantise. Narrow -- 4 tensors of 753 -- but one is the
lm_head this checkpoint ships as Q6_K.
Section 7 is rewritten as a real work queue: localize the spec fault, save a
bandwidth floor and bench output (none survives -- the 12.5 tok/s figure has no
denominator and bench_spec.py's accept-length instrument has never executed),
free headroom for the spec arm, then recover sglang_src into version control and
land the kernel build + HIP binding in-tree.
Also corrected: the marker preamble claimed the model had never generated a
token -- it has; what remains UNVERIFIED is coherence, since no content-checked
probe.py run survives. And rev1's own falsification of "no smaller fallback
without a new download" is kept -- Q2_K_XL is on the laptop and is the right
first bring-up vehicle.
Not evidence, recorded so nobody miscites it: boot_pd631.sh runs a safetensors
2B model, not the GGUF, so it proves PD plumbing only.
Nothing on the laptop was started, stopped or modified; the sweep was read-only.
No rig GPU was used.
efschu
pushed a commit
to efschu/htsglang
that referenced
this pull request
Aug 12, 2026
…it is on the card The GGUF loader's page-cache half is already handled: GGUFReader mmaps, and ConsumedPageDropper (sgl-project#391, a8a2f7b) madvises each tensor's extent away behind the stream. What madvise cannot touch is ANON memory, and that is where the model was sitting twice. fused_moe_triton/layer.py:1318-1324 stores every expert's CPU tensor in TWO aliasing holders (param.expert_data_map and param.data_container). The default materialize branch then did: stacked = torch.stack([get(i) for i in range(count)], dim=0) param.materialize(stacked.shape, dtype=stacked.dtype) param.data.copy_(stacked) continue -- torch.stack building a second full host copy of the layer's expert set on top of the loaded one, and `continue` leaving BOTH holders populated for the process's lifetime, after the experts were already resident in the parameter. Every other terminal branch already releases them: the offload branch at :2772/:2780-2781 and the #391c streaming drain at :1649-1650. The default branch was the only one that did not, so a plain GGUF MoE boot at resident fraction 1.0 took the leaking path. This box was OOM-killed twice by host pressure; a ~120-GiB checkpoint makes the second copy concrete. Fix, one hunk: clear data_container BEFORE the fill (with both holders aliasing, releasing one frees nothing, so this is what lets drop() return each expert's bytes as it lands), materialize from the row_shape and dtype the source already reports, then copy expert-by-expert with del + drop per expert, and clear expert_data_map after. The torch.stack second copy is gone; shape and dtype are provably the same values torch.stack derived, since _gguf_expert_source computes them from the same first expert. One semantic guard added deliberately: torch.stack REJECTED a ragged or mixed-dtype expert set, while copy_ would broadcast or cast it silently. The rejection is kept as an explicit ValueError naming the parameter and the offending expert index, so the refactor cannot turn a hard failure into silent corruption. Cost: per-expert H2D instead of one bulk copy -- for a 128-expert layer, 128 pageable copies of order 10 us against one full host stack allocation plus memcpy per layer. Expected net win; NOT measured, and not claimed. Red first, verified independently of the implementing agent by reverting the hunk and re-running: 3 failures on the pre-fix tree from three independent instruments -- - object level: data_container still holds 8 live expert tensors - peak level: two whole-set torch.stack calls, one per parameter - RSS level: 8728576 bytes beyond the parameter survived materialization of a 16 MiB expert set The object-level guard is the one that pins the fix; the RSS assertion is corroborating only and is documented as such in the test, because it is below the allocator noise floor for small payloads (an earlier 2-MiB-payload draft PASSED on broken code). Not changed, and why: - linear.py:979-980/1014-1015/1780-1781 append .narrow() views that pin full host storages, but they are cleared during process_weights_after_loading (gguf.py:1295/:1668) -- a transient load-time peak, not permanent residency. Bounding it needs a .clone() that trades peak for a copy and is a net loss at TP=1. - the `len(data_container) > 1` clear gates at gguf.py:1260/:1651 are correct, not leaks: create_weights sets the container on every linear while only the merged/QKV sites append, so len is 0 or >= 2. len == 1 would leave qweight unmaterialized and break at forward time first. Reasoned from the append sites, not from an exhaustive enumeration -- recorded as unproven. - tests/moe_offload/test_gguf_moe_offload.py:285 is now misnamed (test_materialize_default_path_still_builds_the_full_stack); it still passes because it asserts shape and byte identity, which this change preserves, but the name asserts something no longer true. Left alone to keep the diff minimal; rename when next touched. Test: test/registered/unit/model_loader/test_gguf_host_residency_644.py 6 passed (3 failed pre-fix) unit/model_loader/ + unit/quantization/ 467 passed, 29 skipped, 94 subtests (base 455 + 6 from sgl-project#647 + 6 from this = 467, no failures, skips and subtests unchanged) tests/moe_offload/test_gguf_moe_offload.py 20 passed unit/layers/moe/ 214 passed, 395 subtests ruff --select=F401,F821,UP037 clean CPU-only; no GPU boot this shift, so the H2D path is exercised CPU-to-CPU. What is proven is that the source references no longer outlive the copy, plus byte identity against real Q4_K/Q6_K payloads through the existing moe_offload suite.
efschu
pushed a commit
to efschu/htsglang
that referenced
this pull request
Aug 12, 2026
… bundle, errors first Per-bug verdict (fixed / refused / recipe-gated), the evidence each rests on, and what was deliberately NOT done. The recipes are the honest part: the two-instance PD boot for a general non-divisible split, the GGUF corpus sweep for sgl-project#647's one-entry table, and the single GPU boot that sgl-project#647 and sgl-project#644 both still need. Records three things the next shift would otherwise re-derive: that a parse-time gate for sgl-project#643 cannot exist (a PD arm's ServerArgs never names the peer's TP size), that compute_head_slice_params cannot be repaired in place (its return type is the wall), and that a mincore test cannot see sgl-project#644 (it measures file pages; the leak is anon). Also pins a pre-existing red so it is not later misread as a regression from this branch: unit/distributed/test_uneven_dcp_pool_geometry.py and test_uneven_tp_memory.py are 8 failed / 40 passed on BOTH the pristine base aca5037 (measured in a scratch worktree) and this branch.
efschu
pushed a commit
to efschu/htsglang
that referenced
this pull request
Aug 12, 2026
…e-existing distributed red with both-tree numbers The unit/distributed reds are now measured on BOTH trees rather than asserted: whole-directory unit/server_args + unit/distributed gives 22 failed / 3303 passed / 12 skipped / 1048 subtests on the pristine base aca5037 (scratch worktree, since git stash is not available here) and the identical figure on this branch. Narrowed to the two owning files, 8 failed / 40 passed on both. Recorded so a later reader cannot mistake them for a regression from this bundle, and flagged as something to file rather than tolerate -- a red in a directory people run trains readers to ignore red. Also adds the single-boot validation recipe sgl-project#647 and sgl-project#644 both still need (load-time gate assertion, behaviour diff for sgl-project#647 where a DIFFERENCE is the pass condition, host-RSS plateau comparison for sgl-project#644 where byte IDENTITY is the pass condition, and the ordering that keeps the two from being confused for each other).
efschu
added a commit
to efschu/htsglang
that referenced
this pull request
Aug 12, 2026
…rt-gate completion fix The dense-GGUF suffix table had one entry, '.gate.weight'. A MoE shared-expert gate's HF name ends in '_gate.weight', so a BF16 ffn_gate_inp_shexp took the .qweight rename, matched no parameter, and was dropped -- which the draft loader's own guard turned into a HARD BOOT REFUSAL for every GGUF MoE checkpoint carrying an MTP block. Carries docs/dev/631/HANDOFF_VAL_R4.md.
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
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.
Motivation
cc @merrymercy @Ying1123 @hnyls2002
When the number of
num-promptsis relatively large, for example,10,000, the defaultulimit -non my development machine is1024. This causes most requests to fail. This fix ensures that the server automatically sets ulimit upon startup and similarly requiresulimitduring client benchmark.hardware: A100 80G
client: https://github.com/vllm-project/vllm/blob/main/benchmarks/benchmark_serving.py
Note: disable streaming
ref: https://www.ibm.com/support/pages/recommended-values-os-ulimit-feature
Modification
as titled
Checklist