caco-zero-three: serve Qwen3.8-27B NVFP4 through vLLM #29
Loading…
Reference in a new issue
No description provided.
Delete branch "feat/vllm-backend"
Deleting a branch is permanent. Although the deleted branch may continue to exist for a short time before it actually gets removed, it CANNOT be undone in most cases. Continue?
vLLM на 2×V100: что в этом PR и на что смотреть при проверке
1. Зачем
Чекпойнт
Qwen3.8-27B-QUASAR-NVFP4— веса в NVFP4, а llama.cpp его не умеет. Единственный путь на Volta — форк 1CatAI/1Cat-vLLM (W4→FP16 GEMM, пейдженный FP8 KV, continuous batching). PR добавляет: декларативную провизию окружения, рабочую топологию (vLLM живёт отдельным systemd-юнитом, llama-swap только проксирует), измерения и документацию всех найденных грабель.Модель вне nix намеренно: 20.6 ГБ весов — данные, не пакет (качаются скриптом на SSD гостя). Всё остальное собирается из
nixos-rebuild switch: на чистой машине руками делать нечего.2. Топология
Почему отдельный юнит, а не дочерний процесс llama-swap. Юнит llama-swap закалён профилем
web:DynamicUser+ProtectHome=true(переменнаяHOMEприходит пустой),ProtectSystem=strict,PrivateTmpиMemoryDenyWriteExecute=true(W^X). JIT-движку нужно и то, и другое, и третье, а per-model исключений у systemd нет. llama.cpp этого не замечает вовсе — его ядра собраны при сборке, поэтому юнит выглядел обкатанным, пока vLLM внутри него умирал за 20 секунд.3. Измеренные числа (это и есть ответ про длину контекста)
--max-model-len 168340GPU KV cache size: 169,816 tokensnvidia-smiпри работающем движке--max-num-seqs 64(потолок ~116 Mamba-блоков)Maximum concurrency …Честная оговорка про параллельность: два запроса нормальной длины идут одновременно; два запроса по 131 072 токена одновременно в пул не помещаются ни при каком
util— это арифметика VRAM на двух картах, а не настройка.DFlash2 не влезает и выключен. Проверены четыре конфигурации (util 0.85 / 0.90 / 0.94 и 0.92 с явным
--kv-cache-memory 1GiB) — все падают вload_draft_modelсOutOfMemoryError: vLLM заполняет KV-пул по всему бюджетуgpu-memory-utilization, а драфт (~2 ГиБ на карту) грузится после и упирается в 7–25 МБ свободных. Форк мерил его на TP4; квантизацию драфта не трогали.4. Что смотреть в диффе — и почему каждая строка такая
hosts/caco-zero-three/services-llm.nix— центральный файл. У каждого неочевидного решения есть замеренная причина, в коде она продублирована комментарием:vllmWheel = fetchurl { hash = … }SHA256SUMSрелиза v1.5.0), но не собирается в nixvllm-venv.service(oneshot + stamp)torch == 2.10.0; в nixpkgs-пине 26.11vllmпомеченbroken, аtorchне эвалится подcudaCapabilities = ["7.0"]. venv — единственный способ спарить тот самый wheel с тем самым торчемInvalid wheel filename (invalid version)— в store-имени 32-символьный хэшLD_LIBRARY_PATH= gcc-lib + zlib +/run/opengl-driver/libImportError: libstdc++.so.6 / libz.so.1TRITON_LIBCUDA_PATH/sbin/ldconfig -p, которого на NixOS нетCC/CXXFailed to find C compiler— triton JIT-компилирует лончерCUDA_HOME(symlinkJoin nvcc+cudart+cccl)No CUDA or HIP or MPS availableNVCC_PREPEND_FLAGS-Iна CUDA-заголовки, а nvcc берёт пути относительно себя иCUDA_HOMEигнорирует:cuda_runtime.h: No such file or directoryHOME/XDG_CACHE_HOME/TRITON|TILELANG|TORCHINDUCTOR_CACHE_DIR→/var/lib/vllmStateDirectory), неDynamicUser-дерево 0700 за симлинком: иначеfailed to map segment from shared objectscript = ''…''вместоExecStart = [ … ]invalid loads value: '{enable_thinking:true}'RestrictAddressFamiliesперечислен целиком//, а это не module-merge: свой список-ключ затирает тировый →OSError: [Errno 97] Address family not supportedна собственном портуproxy,checkEndpoint,cmd = sleep infinitycmd(unable to get sanitized command: empty command); трафик идёт наproxy, аsleep— инертная заглушка, которую llama-swap волен «выгружать»hosts/caco-zero-three/users.nix— ⚠️ security-relevant: вauthorizedKeysпользователяslopychдобавлен публичный ключ агента, чтобы работать по ssh с беспарольнымdoasвместо qga. Ключ публичный, откатывается одним коммитом; если такой доступ не нужен — это первое, что стоит отклонить.hosts/caco-zero-three/vllm/requirements.lock— 188 пакетов, снятpip freezeс реально заработавшего venv (сам форк исключён). Перегенерация описана в README.hosts/caco-zero-three/vllm/README.md— как бампить wheel, все nix-грабли, профили памяти, DFlash2, и как диагностировать падение движка.5. Как проверять
Проверено на живой машине: запрос через
:8000→ 150 токенов за 4.34 с (34.6 tok/s), два параллельных запроса — оба 200,GPU KV cache size: 169,816 tokens, юнитactive,MemoryDenyWriteExecute=no, llama.cpp-модели на месте.6. Ограничения, которые не прячу
doas systemctl stop vllm(обратно —start, холодный старт 4–6 мин). Юнитenabledи поднимается при загрузке.torch.compile+ захват CUDA-графов), поэтомуhealthCheckTimeoutу llama-swap поднят до 900 с.7. Откат
nixos-rebuild --rollbackлибо предыдущее поколение в GRUB. llama.cpp-модели этот PR не трогает: их записи и хардненинг юнита те же, что были.The first real request killed the engine: flash_qla's GDN kernel JITs through tilelang, whose check_cuda_availability() is nvcc.find_cuda_path() -- it looks for a TOOLKIT, not a GPU, and returns $CUDA_HOME or raises. Driver-only boxes therefore report "No CUDA or HIP or MPS available on this system" and the engine dies with a scheduler KeyError. CUDA_HOME now points at a symlinkJoin of nixpkgs' cuda_nvcc + cuda_cudart + cuda_cccl from the same 12.8 pin llama.cpp builds against. Verified on the box: check_cuda_availability() = True and determine_target("auto") = arch sm_70.Root cause of the whole llama-swap failure, now read off the child's own log (the wrapper added in the previous commit): RuntimeError: Error raised in subprocess: Triton cache error: compiled module cuda_utils.so could not be loaded ImportError: /var/cache/llama-swap/triton/<hash>/cuda_utils.cpython-312-x86_64-linux-gnu.so: failed to map segment from shared object /var/cache/llama-swap takes writes (llama.cpp caches HF downloads there) but it is a symlink into the unit's private 0700 directory, and mapping an .so through it fails. vLLM's model-architecture inspection runs in a subprocess, so the loader error ended the process and llama-swap reported only 'upstream command exited prematurely'. HOME and every cache root now live under the unit's PrivateTmp /tmp, which is both writable for the dynamic user and map-friendly -- verified in an identical sandbox where the architecture resolved and the engine came up.The router's unit cannot host a JIT backend: it carries W^X (MemoryDenyWriteExecute, from the web hardening tier) and hands its children an empty HOME, and neither is exemptable per model -- vLLM dies with 'cannot allocate executable memory'. So vLLM becomes systemd.services.vllm: * nd.hardening.floor.runtime, the repo's JIT-safe tier (baseline + seccomp allow-list, no W^X, device access left to the service), plus AF_NETLINK for NCCL's getifaddrs bootstrap (the same fix the llama-swap unit needs); * its own user and StateDirectory=/var/lib/vllm -- a real directory, because the toolchain compiles .so files and the loader must map them (a DynamicUser's 0700 tree behind a symlink gives 'failed to map segment from shared object'); the caches persist, so kernels compile once; * own MemoryMax 16G, so the router's 14G cgroup no longer has to cover a Python backend, and TimeoutStartSec 20min for the 4-6 min cold start; * ExecStart as an argv list, which retires the JSON-quoting trap. llama-swap keeps the qwen38-nvfp4 entry as a passthrough (proxy: http://127.0.0.1:5810, checkEndpoint: /health) in the same gpu-llm group; the llama.cpp models keep their hardening untouched.A bare list of strings is a list of ExecStart= lines, and systemd refuses more than one for Type=exec ('Service has more than one ExecStart= setting, which is only allowed for Type=oneshot services'), leaving the unit bad-setting and inactive. The command is one element: ExecStart = [ [ bin arg... ] ].The unit came up bad again with vllm serve: error: argument --default-chat-template-kwargs: invalid loads value: '{enable_thinking:true}' systemd's unit-file reader strips the double quotes out of the JSON while it parses the ExecStart word, so argv arrives quote-free. The command is now a script= line: the shell is the quoting layer that already handles this kwarg correctly (it is how the llama-swap entry passed it while vLLM was a child).systemd rejected it as 'Unknown key script in section [Service]' and refused the unit ('Service has no ExecStart='), which also made the whole switch exit 4. Moved to the service level, where NixOS turns it into ExecStart=/bin/sh -c.The tier is applied with a plain '//', which is not a module merge, so writing RestrictAddressFamilies in the override REPLACED the tier's list instead of extending it. Left with only AF_NETLINK the unit could not even bind its own port: OSError: [Errno 97] Address family not supported by protocol socket.socket(family=family, type=socket.SOCK_STREAM) All four families (the tier's three plus netlink for NCCL) are now spelled out.Both are chat-template kwargs, and this checkpoint's template reads both: - reasoning_effort (xhigh default / medium / low) -- medium shortens the thinking, which matters because the template's xhigh is what tipped the model into repetition loops; - preserve_thinking -- keeps the assistant's <think> block when the history is re-rendered. The template already defaults to true; stating it explicitly makes the behaviour independent of the client. vLLM's --reasoning-config has no equivalent knob (only reasoning_parser and reasoning_parser_plugin), so the template kwarg is the right layer.Their SM70 NVFP4 acceptance doc fixes the contract for this checkpoint: temperature 1.0 / top-p 0.95 / top-k 20, thinking on at the template's xhigh default, prefix caching on, Mamba align, 4096 max batched tokens, 32 sequences, custom all-reduce off, CUDA graphs off, FP8 E5M2 KV. We were running with different values for six of those, and this model is a 48-GDN-layer hybrid, so the Mamba cache mode is not a detail: vLLM's default is 'all', theirs is 'align'. The symptom we are chasing: in ~1/6 requests (much more at temperature 0.6) the model degenerates inside its reasoning block -- literal repetition ('17 x 23 = 391' hundreds of times) or vacuous rambling -- burns the whole token budget and never emits the closing marker, so content comes back empty. Dropped from our side while hunting this: the temperature 0.6 / repetition 1.05 override (measured no better than the contract 1.0/1.0) and the extra template kwargs reasoning_effort=medium / preserve_thinking (both are template defaults). Kept for hardware reasons: TP2, 168340 context, 0.94 utilisation, no DFlash2.Replaces the previous attempt, which was named AWQ but packaged as compressed-tensors -- and compressed-tensors W4A16 declares min_capability 75 on Volta, so the engine crash-looped instead of loading. Verified before downloading this time: config.json quantization_config = {quant_method: awq, bits 4, group_size 128, version gemm, zero_point true} -- the real AWQ layout, no quantize_config.json needed; awq_marlin.py get_min_capability() returns 70 outright and its SM70 gate accepts group sizes 32/64/128; awq_gemm / awq_marlin_repack / marlin_gemm all present in the wheel; chat_template.jinja is the same 8952-char template the QUASAR build ships and supports enable_thinking / reasoning_effort / preserve_thinking / tools. Size 18.73 GB, one shard, 1.3 GB lighter than the alternative AWQ build whose extra weight is a 0.849 GB MTP head (unusable here) plus 0.475 GB of g64 overhead. Kept identical: flags, context 81920, E5M2 KV, contract sampling, served name qwen38-nvfp4, and the QUASAR directory on disk for rollback.Measured on the box, not read off the source: AWQConfig.get_min_capability() returns 70 both with and without VLLM_SM70_QUANT_BACKEND=marlin, because the fork lowered that gate for SM70 unconditionally. The previous commit message claimed the variable was a load condition ('without it the AWQ path keeps the 75 gate and cannot load') -- that is false, and it came from reading the if-branch of get_min_capability and assuming the else-branch by analogy with compressed-tensors instead of reading it. The variable stays, for the reason that is actually true: it selects the kernel. With forces_marlin() true, awq.py delegates to AWQMarlinConfig, i.e. the marlin kernels (awq_marlin_repack, marlin_gemm) rather than the fork's own awq_gemm path -- the route the fork's own SM70 whitelists are written against. Probe used (both runs on the live box): forces_marlin=False awq_min_capability=70 (unset) forces_marlin=True awq_min_capability=70 (VLLM_SM70_QUANT_BACKEND=marlin)The AWQ checkpoint crash-looped the unit on its first load: Qwen3_VisionTransformer -> Qwen3_VisionBlock -> Qwen2_5_VisionAttention -> RowParallelLinear -> create_weights ValueError: The input size is not aligned with the quantized weight shape Not the language model: the vision tower. This checkpoint ships with an empty modules_to_not_convert, so the tower is quantized too, and the vision attention's row-parallel proj layer cannot be sharded cleanly at TP2 with a quantized weight shape. We had flagged that empty list earlier as a risk to watch -- it bit here first. Fix uses the engine's own mechanism rather than touching the checkpoint: vLLM's tower marking skips tower modules when a modality's limit is zero, and --enable-mm-embeds is the documented companion that keeps precomputed embeddings usable while not loading encoder modules. We serve text only. --limit-mm-per-prompt '{"image":0,"video":0}' --enable-mm-embedsAWQ (mattbucci/Qwen3.8-27B-AWQ, real AWQ, group 128) does load its language model but crash-loops on the vision tower: RowParallelLinear.create_weights raises 'The input size is not aligned with the quantized weight shape'. Its modules_to_not_convert is empty, so the tower is quantized and its row-parallel proj layer cannot be sharded at TP2. Verified that neither --limit-mm-per-prompt '{"image":0,"video":0}' nor --enable-mm-embeds stops the tower from being constructed -- the flag is applied in the log and self.visual = Qwen3_VisionTransformer(...) still runs. Quasar is restored so the service works; the vision route needs a checkpoint whose tower is intact, TP1 (does not fit), or a newer engine. Also recorded: our wheel is v1.5.0 (2026-09-02) while main is 663 commits ahead (2026-09-14) and contains the fixes we are looking for -- grouped attention warp-state ordering, several FP32 precision fixes for SM70, and the FP32 partials/logits work in docs/design/sm70_dflash2_fp32_defaults.md.My VLLM_SM70_QUANT_BACKEND=marlin was added as a kernel-selection knob and is not one: it forces the marlin route for every quantized checkpoint, and QUASAR (NVFP4) dies at load on it: torch.ops._C.gptq_marlin_repack(...) RuntimeError: size_n = 8240 is not divisible by tile_n_... which crash-looped the unit five times. I introduced that flag on a reading I never verified against QUASAR -- the same mistake I called out earlier. Also dropping --limit-mm-per-prompt '{"image":0,"video":0}' --enable-mm-embeds. Measured and disproved: the flag is applied ("enable_mm_embeds is True" appears in the log) yet self.visual = Qwen3_VisionTransformer(...) still runs, so the tower is always built for this model class. It did not help with AWQ either. Config returns to the known-good baseline: QUASAR with the flag set as before. The AWQ findings and the 663-commit gap behind main stay documented above.Comments only -- the running system already has this configuration; no deploy needed for a comment change. Root cause, measured: the checkpoint's chat_template.jinja derives its reasoning instruction from reasoning_effort with default 'xhigh' and has NO branch for 'medium'. Omitting the kwarg therefore ordered maximum-effort thinking, and the model ended its turn (emitting <|im_end|>) inside its own reasoning without ever writing an answer. Proof: tokenizer ids show <think> 248068 and </think> 248069 are not part of eos_token_id {<|im_end|>, <|endoftext|>}, so no engine-side stop was involved. Also recorded: what was excluded with measurements (eager vs compile path, VLLM_SM70_QWEN_GDN_FULL_FORWARD=1, sampling, KV dtype, the W4A4 theory), the KV budget gain from dropping eager (91,639 -> 161,063), and the remaining candidate for long-output robustness (663 commits behind main).QUASAR measures 18/30 (60%) correct on 10 verifiable questions x3 through the client path; temperature does not help (0.6 -> 15/30). Two mechanisms are read from the checkpoints' own configs: QUASAR compressed-tensors, W4A4: ignore = 6 entries, ZERO linear_attn/GDN -> the recurrent GDN layers are quantized, at 4-bit ACTIVATIONS on an emulated Volta path, on every single token. AWQ-MTP awq, W4A16: 103 exclusions, 96 linear_attn/GDN + 5 vision -> GDN state stays FP16 and activations are not quantized at all. Vision is excluded here, so the TP2 sharding failure that killed mattbucci/Qwen3.8-27B-AWQ (empty modules_to_not_convert) does not apply. Attribution note: not the checkpoint 1Cat use -- their docs name Qwen3.6-27B-AWQ (their own build, generation 3.6). Class matches, artifact does not. The mechanism above is from config.json, not by analogy. Checkpoint verified complete before this commit: 18 files, 20,043,146,728 bytes.The first attempt crash-looped (NRestarts climbing) with: ValueError: Selected backend AttentionBackendEnum.FLASH_ATTN_V100 is not valid for this configuration. Reason: ['kv_cache_dtype not supported'] So the turboquant_* cache dtypes are not a free-standing flag: they are served only by vllm/v1/attention/backends/turboquant_attn.py, whose get_name() is 'TURBOQUANT'. Switching the backend is a CONSEQUENCE of choosing the cache dtype. arg_utils additionally forces flash_attn_version=2 for turboquant_* (FA3+ boundary layers assert FlashAttentionImpl) and does that automatically. The preset is exactly what was asked for, read from the fork's TQ_PRESETS: turboquant_k8v4 = {key_quant_bits: 8, value_quant_bits: 4, norm_correction: False}, i.e. K in 8 bits and V in 4. MTP is deferred to its own step: the fork's MTP4 recipe carries "attention_backend":"TRITON_ATTN", which almost certainly cannot serve turboquant_* either, so the pairing must be validated separately. Context stays at 163840. Baseline for quality: 30/30 twice on fp8_e5m2/81920/FLASH_ATTN_V100/no-MTP.This is the pairing the fork's own benchmark table names: 'Qwen3.6-27B-AWQ + MTP4 | 4x V100 - TP4 - E5M2 KV - Flash-V100 - CUDA Graph'. The turboquant_k8v4 route and MTP are mutually exclusive for now: the MTP4 recipe uses attention_backend TRITON_ATTN, and TRITON_ATTN reports supports_kv_cache_dtype('turboquant_k8v4') == False (measured on the box). Memory arithmetic, corrected after the user's question 'why did memory get smaller with turboquant': per token the turboquant cache is actually CHEAPER (1.32 GiB / 99840 tokens = 14.2 KB) than fp8 (3.56 GiB / 205494 = 18.6 KB). The turboquant attempt looked worse because of its BACKEND overhead -- at the same 0.94 it left 1.32 GiB for KV versus 3.56 GiB with FLASH_ATTN_V100, i.e. ~2.2 GiB more non-KV memory. At fp8, 163840 tokens cost ~2.98 GiB and fit in 3.56 GiB, so the proven 0.94 is enough (not the 0.97 that the turboquant layout needed). Quality must be re-measured: MTP changes the decode path. Baseline: 30/30 twice on fp8_e5m2/81920/FLASH_ATTN_V100/no-MTP.Budget at 0.95 with MTP: available KV 2.55 GiB, 163840 needs 3.14 GiB -- 0.59 short. Levers, all measured on this box: --language-model-only the vision tower is 921 MB FP16 (from safetensors headers), 461 MB per card at TP2 -> 0.92 GiB total. Images are not used. Confirmation is the per-card 'Model loading took' figure dropping by ~0.5 GiB and the vision-cache log line disappearing. --max-num-seqs 4 batch buffers (incl. the MTP verifier) shrink ~8x --max-num-batched-tokens 2048 halves prefill chunk and piecewise buffers Together ~1.4 GiB freed vs a 0.59 GiB deficit, with the graph capture (0.49 GiB per card, measured) accounted for. Hard rule satisfied: 2048 >= 4. User's decision rule: if the KV pool comes out above 150k tokens, keep it.Per the user: the vision tower must stay (visual analysis is required), and no KV offloading is to be used. So the only reductions are the two scheduling knobs: --max-num-seqs 4 (was 32): the fork generates capture shapes as [1, 2, 4] + range(8, 256, 8) + range(256, max+1, 16) with max = min(max_num_seqs*2, 512) -- at seqs=32 that meant graphs up to 64 (exactly the list seen in the log), at seqs=4 it collapses to [1, 2, 4]. --max-num-batched-tokens 2048 (was 4096): halves prefill chunk buffers. That is expected to recover most of the 0.93 GiB vLLM estimates for CUDA graphs, against a 0.59 GiB deficit for 163840 at 0.95 with MTP and the tower loaded. Hard rule satisfied: 2048 >= 4. Goal: measure the resulting KV pool. If above 150k tokens, keep it.Per the user: run without MTP and with the previous scheduler values so that real chat traffic produces comparable prefill/decode numbers in the journal. no --speculative-config MTP off (as asked). Removing it also frees the 1.17 GiB of MTP overhead, which is what makes 163840 fit here: 2.98 GiB needed (18.6 KB/token at fp8, measured) vs 3.56 GiB available at 0.94, and the measured pool in this shape is 205,494 tokens. --gpu-memory-utilization 0.94 the proven value; 0.95 would eat capture headroom (capture needs 0.49 GiB per card, measured). --max-num-seqs 32 and --max-num-batched-tokens 4096 previous values, restored so the throughput figures stay comparable. Vision tower stays ON; no --language-model-only, no KV offloading. Throughput is read from vLLM's own 'Avg prompt throughput / Avg generation throughput' journal lines while the user drives chats.User's explicit request, both levers on: --speculative-config (MTP4) ~1.17 GiB VRAM (head weights + verify buffers; 'Model loading took' goes 9.52 -> 10.69 GiB/card) --kv-offloading-size 8 KV spills to host RAM (native backend, summed across TP ranks; host has ~23 GB, ~19 free) Open question this run answers: does the startup KV check count the offload buffer as available capacity? In VRAM alone 163840 needs 3.14 GiB vs 2.39 GiB available at 0.94, so if the check is VRAM-only it will refuse again and name the number to drop to. util kept at 0.94 (not 0.95) because the graph capture needs ~0.49 GiB per card and 0.95 previously OOM'd during capture. Vision tower stays ON; scheduler stays at 32/4096.User asked for the maximum context that fits together with MTP. Measured budget: with MTP4 at 0.94 the pool reports 2.39 GiB and the engine's ceiling is 117,968 tokens, but vLLM's 0.93 GiB graph estimate is not enough in practice -- that is what OOM'd at 0.95 with 122,880 (slack 0.02 GiB). Keeping ~0.37 GiB of real slack: available 2.39 GiB - 98304 * 20.58 KB/token (2.02 GiB) = 0.37 GiB 98304 (96k) is also the largest round value that still covers the user's longest real prompt of 96,299 tokens (llama-swap /api/metrics/activity). --max-num-seqs 8 (was 32) capture shapes are generated as [1,2,4] + range(8,256,8) + range(256, min(seqs*2,512)+1,16), so 8 caps capture at 16 shapes instead of 64 -- measured that the list really does collapse. Fewer graphs = the capture fits the slack. 8 concurrent is plenty for one user. --max-num-batched-tokens 2048 satisfies the hard rule (>= seqs) No KV offloading (measured: does not extend the pool). No --language-model-only (visual analysis required). Fallbacks if capture still OOMs: 90112 (88k), 81920 (80k).The engine died twice with: EngineCore encountered a fatal error core.py run_busy_loop -> _process_engine_step -> step_with_batch_queue scheduler.py schedule -> connector.update_state_after_alloc offloading_connector.py -> offloading/scheduler.py update_state_after_alloc AssertionError: assert num_locally_computed_tokens <= num_locally_computed_gpu_blocks * gpu_block_size That check requires every locally-computed token to sit in a hashed block, but our model is GDN-heavy and runs with --mamba-cache-mode align, which rewrites logical positions of existing blocks while the offloading connector assumes an append-only positional mirror (the same file notes requires_exact_boundary_source for exactly this case). The engine's own log calls the feature experimental ('Initializing OffloadingSpec. This API is experimental and subject to change'). Symptom pattern: healthy for 13-19 minutes, then EngineDeadError, then 500s to clients. Removing it costs the KV spill's benefit (preemptions back to ~35 from 1, evicted blocks no longer reloadable from host RAM) but no correctness. mamba-cache-mode align is left alone: it does not affect quality (30/30 on two checkpoints with it) and it only aligns mamba state with attention pages for prefix caching.Root cause of the hang where tokens stop arriving while num_requests_running stays at 2: systemctl show vllm -p LimitMEMLOCK -> 8388608 (8 MiB) systemctl show -p DefaultLimitMEMLOCK -> 8388608 ulimit -l (root) -> 8192 KiB and nothing overrides it -- not this file, not /etc/security/limits.conf. With the KV spill plus the UVA vision tower the engine must LOCK host pages; 8 MiB is far too little, so the driver reports 'NVRM: failed to allocate page table!' from os_lock_user_pages, the shm_broadcast ring between EngineCore and the TP workers never gets its blocks ('No available shared memory broadcast block found in 60 seconds'), and the engine deadlocks. Measured on the live hung engine: workers: state=R, wchan=0, 12.3 s CPU per 8 s wall, GPU 100% (spinning) engine: state=R, stuck in step_with_batch_queue -> _wait_for_response -> dequeue -> acquire_read -> poll (shm_broadcast.py) metrics: generation_tokens and iteration_tokens frozen (3529 / 738 for 45+ s) while num_requests_running = 2, num_requests_waiting = 0 LimitMEMLOCK is about pinning pages, not about how much memory the unit may use, so MemoryMax/MemorySwapMax stay exactly as they were. Spill size is the user's 10 GiB.View command line instructions
Checkout
From your project repository, check out a new branch and test the changes.Merge
Merge the changes and update on Forgejo.Warning: The "Autodetect manual merge" setting is not enabled for this repository, you will have to mark this pull request as manually merged afterwards.