From be840b48819a230d2d1e87c2518718dfa1d56eac Mon Sep 17 00:00:00 2001 From: Frank Date: Mon, 5 Oct 2026 17:35:05 -0700 Subject: [PATCH] frank-large: halogen container cutover, rootless rlimit law, and the DFlash2 drafter shape --- AGENTS.md | 17 +++++++- scripts/serve-frank-large.sh | 11 +++-- serverAppConfigs-docs/LLM/VLLM Overview.md | 40 +++++++++++++++++++ .../LLM/frank-large-halogen27b-dflash2.md | 14 +++++-- 4 files changed, 73 insertions(+), 9 deletions(-) diff --git a/AGENTS.md b/AGENTS.md index c93233b..2d87145 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -25,7 +25,7 @@ merge absorbed the former `vllmScripts` repo, history preserved). | Target | Service | Runtime | Model | Port | |--------|---------|---------|-------|------| | Dankserver | `frank-primary` | tcclaviger RDNA4 vLLM fork (`docker.io/tcclaviger/vllm@sha256:da4029666…`, digest-pinned, plain `podman run`, root systemd unit) | `orcarouter/Qwen3.8-Flash-Next-Uncensored-FP8` (~172.8 GiB, 131 shards), TP=2, expert offload automatic (fork plans it), prefetch OFF, **fp8 KV** (`--kv-cache-dtype fp8`), **no speculative decode**, MNBT 2048 (arena-free: the registry prices overhead flat up to its nbt 4096 anchor), MoE gather 32×64, **`--expert-precise-memory`** (fitted 5.25 GiB overhead, not the 7.60 GiB default), **`VLLM_CUDAGRAPH_MEM_BASE_MB=0` / `_PER_GRAPH_MB=12`** (the plan's CUDA-graph term priced at the measured capture — 0.31 GiB then a NEGATIVE −0.06 GiB delta, so the 256 MiB analytic base was fiction → graph term 0.07 GiB), **`GPU_MAX_HW_QUEUES=2`**, **mounted cache dir** at `/root/.cache` (the image ships no `/root/.cache` at all; unmounted, every boot re-pays `torch.compile took 181.20 s`), **3 slots × 200704 ctx** at KV_BYTES 3,166,500,000/rank (~400,700 KV tokens at the engine's 7,891.8 B/token here) — the **pool, not the slot count, caps resident conversations** (measured live: `GPU KV cache size: 401,408 tokens, Maximum concurrency for 200,704 tokens per request: 2.00×`; the 3.15e9 pool below it measured 398,658/1.99×), and shorter requests use the spare blocks. **`--max-num-seqs` must stay 3: SEQS=2 cannot boot this model** — torch.fx `ConstraintViolationError` on `query_start_loc`, measured at BOTH MNBT 1024 and 2048. **util stays 0.982 BY LAW**: at `le=1` the worker init refuses `budget == total` with `ValueError: Free memory on device cuda:N (31.29/31.86 GiB)` on both ranks before a single weight loads — measured crash loop 2026-10-03. Arena at shape G: **17.77 GiB / 7,762 slots measured**, KV 401,408 tokens / 2.00× concurrency, 0.89 GiB of real free VRAM per card held through a live 4,096-token generation (the ~3.5 GiB of driver-reported free VRAM at shape E was estimator slack — the registry's 9190 B/token KV reserve vs the real 7,892, and the 5.25 GiB overhead vs far less — NOT allocatable), and the client contract wired: `--reasoning-parser … -| Dankclaw | `frank-large` | **halogen container** (`ghcr.io/peonist-ai/halogen@sha256:dc0a39a0…`, digest-pinned, plain `podman run` via `scripts/serve-frank-large.sh`, root systemd unit) | `peonist-ai/halogen-qwen3.8-27b` **`qwen3.8-27b-p1w4d-d2.hgn`** (33.4 GiB; the `.hgn` format loads in NO other engine — llama.cpp, vLLM and transformers all refuse it), **DFlash2 spec decode (`HALOGEN_DRAFTER=2`)**, **1 slot × 200,700 ctx** — `HALOGEN_KV_SLOTS=1` is load-bearing: batching XOR speculation, slots>1 silently disables the drafter (the sibling engine measured 41 → 14 t/s), d2 KV (2× fp8, 64 KiB/token → 12.2 GiB pool, engine sizes it and prints a boot ledger naming any clamp), **TEXT-ONLY** (the engine ignores the vision encoder — vision is frank-small's job), **censored base** (the ROCmFPX AGENT rollback is the refusal escape hatch), **no server-side sampler: absent temperature is GREEDY**, so the card's thinking recipe (temp 1.0 / top_p 0.95) rides the wire in `models.yml`. Publisher card 31.7 t/s mean across eight text workloads vs the retired ROCmFPX endpoint's measured 27.6–28.1. Prompt over the slot is a 400 naming the limit, never a truncation; `max_tokens` over `HALOGEN_MAX_TOKENS_CAP` is a 413. Cutover 2026-10-05|7770| +| Dankclaw | `frank-large` | **halogen container** (`ghcr.io/peonist-ai/halogen@sha256:dc0a39a0…`, digest-pinned, plain `podman run` via `scripts/serve-frank-large.sh`, root systemd unit) | `peonist-ai/halogen-qwen3.8-27b` **`qwen3.8-27b-p1w4d-d2.hgn`** (35,865,565,184 B; the `.hgn` format loads in NO other engine — llama.cpp, vLLM and transformers all refuse it), **DFlash2 spec decode (`HALOGEN_DRAFTER=2`)**, **1 slot × 200,700 ctx** — `HALOGEN_KV_SLOTS=1` is load-bearing: batching XOR speculation, slots>1 silently disables the drafter (measured: 8 slots at serial ~6 t/s/stream), d2 KV (2× fp8, 64 KiB/token → 12.2 GiB pool, engine sizes it and prints a boot ledger naming any clamp), **TEXT-ONLY** (the engine ignores the vision encoder — vision is frank-small's job), **censored base** (the ROCmFPX AGENT rollback is the refusal escape hatch), **no server-side sampler: absent temperature is GREEDY**, so the card's thinking recipe (temp 1.0 / top_p 0.95) rides the wire in `models.yml`. Publisher card 31.7 t/s mean across eight text workloads vs the retired ROCmFPX endpoint's measured 27.6–28.1. Prompt over the slot AND `max_tokens` over `HALOGEN_MAX_TOKENS_CAP` are both 400s naming the limit, never truncations. Serve passes NO `--ulimit nofile` (rootless runc cannot raise Dankclaw's 524288 hard limit — see Key Lessons). Cutover 2026-10-05|7770| | Dankclaw | `frank-small` | llama.cpp master, CUDA sm_86 build (root systemd unit, relinked) | bartowski/MiMo-V2.6-Distill-Qwen-9B-GGUF **Q6_K** (9B dense qwen3_5 hybrid, 7.26 GiB GGUF, no MTP weights so no spec decode) + the repo's f16 mmproj for **vision**, on the RTX 3090, **2 × 200,704** (pool 401,408 — NOT the native window) with **fully unquantized KV (f16 K + f16 V)** under flash attention, **batch 4096 / ubatch 2048** (doubled 2026-10-05: measured +4.6 % prefill at 8,928 tokens for 616 MiB), 21,846 used / **2,279 free** of 24,576 MiB, 94.4 tok/s decode. Resized 2026-10-03 (3 × 180,224 + q8_0 V → 2 × 200,704 + f16 V → + Q6_K rung: the freed slots paid for exact values AND the higher quant; the rung costs a measured −10 % decode — weight-bandwidth-bound); rollback `MODEL_FILE=…/MiMo-Q5KM/…Q5_K_M.gguf` (buys the 10 % back) or `CTX_SIZE=540672 PARALLEL=3 CACHE_TYPE_V=q8_0` (Q5 only) | 7771 | Retired endpoints (`frank-npu`, and every prior runtime) are in @@ -338,7 +338,10 @@ drafter off silently. `HALOGEN_SLOT_CTX=200700` — the per-conversation HARD ceiling (400 naming the limit, never truncation); `models.yml` `contextWindow: 200700` must match or the harness discovers the wall as a 400. `HALOGEN_MAX_TOKENS_CAP=65536` + `HALOGEN_QUEUE_TIMEOUT=7200` are a coupled -pair; the harness `maxTokens: 57344` stays under the cap (over-cap is a 413). +pair; the harness `maxTokens: 57344` stays under the cap (over-cap is a 400 +naming the limit — README, never a 413). Thinking tokens COUNT against the +budget: exhausted mid-thought is `finish_reason: "length"` with empty +`content` and the reasoning in `reasoning_content`. `HALOGEN_DRAFTER=2` = DFlash2. **No server-side sampler**: an absent temperature is greedy, so the thinking recipe exists only in `models.yml` `extraBody`. `response_format` is unsupported (no constrained decoding). @@ -729,6 +732,16 @@ linked runbooks. **Runtime-specific guards** +- **Rootless podman cannot RAISE rlimits, and `--ulimit nofile=1048576` + dies at container create, not at runtime.** runc reports + `error setting rlimit type 7: operation not permitted` (type 7 = + RLIMIT_NOFILE) when the requested hard limit exceeds the host process's + hard limit — on Dankclaw that ceiling is 524288. Inherited-equal limits + are fine, so `memlock=-1:-1` works iff the unit grants + `LimitMEMLOCK=infinity` (systemd is privileged; the value flows through + podman to runc without any raise). Upstream halogen's own run lines pass + NO `--ulimit` at all — copy vendor run lines before adding flags + (measured 2026-10-05 on the frank-large cutover). - **Kairic Edge** (if ever re-deployed): the pinned CK commit needs `third_party/composable_kernel/include/ck/config.h` generated manually (untracked; `git clean` removes it), and its greedy-argmax fastpath rejects diff --git a/scripts/serve-frank-large.sh b/scripts/serve-frank-large.sh index a341492..06f43cd 100755 --- a/scripts/serve-frank-large.sh +++ b/scripts/serve-frank-large.sh @@ -110,14 +110,19 @@ done LOG="${FRANK_LARGE_LOG:-/home/dank/logs/frank_large.log}" mkdir -p "$(dirname "$LOG")" # for the udev-granted supplementary group, seccomp unconfined for the ROCm -# ioctl surface, and unbounded memlock/nofile (the engine mmaps 33.4 GiB of -# weights and wants a big fd budget). No comment lines inside the backslash +# ioctl surface, and unbounded memlock (the unit's LimitMEMLOCK=infinity is +# inherited, so runc never has to RAISE it). NO --ulimit nofile: Dankclaw's +# hard limit is 524288, rootless runc cannot raise a hard limit and dies +# with "error setting rlimit type 7: operation not permitted" (measured +# 2026-10-05 cutover), and 524288 fds is already ~100x what the engine +# wants; the model card and both sibling serve scripts pass no nofile +# either. No comment lines inside the backslash # chain — a mid-chain # terminates the continuation (the 2026-10-05 # exit-127 crash loop on the sibling serve script). podman run -d --name "$FRANK_LARGE_CONTAINER" \ --publish "$FRANK_LARGE_PUBLISH" \ --ipc=host \ - --ulimit memlock=-1:-1 --ulimit nofile=1048576:1048576 \ + --ulimit memlock=-1:-1 \ --device "$DEV_KFD" --device "$DEV_DRI" \ --group-add keep-groups \ --security-opt seccomp=unconfined \ diff --git a/serverAppConfigs-docs/LLM/VLLM Overview.md b/serverAppConfigs-docs/LLM/VLLM Overview.md index e238a6c..dbccf8b 100644 --- a/serverAppConfigs-docs/LLM/VLLM Overview.md +++ b/serverAppConfigs-docs/LLM/VLLM Overview.md @@ -237,6 +237,46 @@ vLLM supports multiple parallelism strategies: - Expert Parallelism (EP) distributes expert networks in MoE models - Context Parallelism (CP) splits context across GPUs + +## Chunked Prefill — the flag surface, and what it actually does + +**YES, the runtime exposes chunked-prefill options, and the running shape +already uses them.** The tcclaviger fork is a vLLM fork, so the standard +scheduler flags are present in the shipped image (verified by reading the +`vllm.engine.arg_utils` and `vllm.config.scheduler` source in the live +container `frank-primary-qwen`, not from a stub): + +| Flag | Default | What it controls | +|---|---|---| +| `--enable-chunked-prefill` | `True` in `SchedulerConfig`; the CLI arg is an auto-flag (`"default": None`) | whether prefill may be split into chunks | +| `--max-num-batched-tokens` (MNBT) | 2048 (`DEFAULT_MAX_NUM_BATCHED_TOKENS`) | tokens processed per prefill iteration | +| `--prefill-schedule-interval` | 0 | gap between prefill iterations | + +The running shape's argv carries `--max-num-batched-tokens 2048` alongside +`--max-model-len 200704`, so chunked prefill is **ON and already active**: a +200704-token prompt is prefilled in ~98 iterations of 2048 tokens, not in one +shot. + +**It is a TTFT lever, not a throughput lever.** MEASURED 2026-09-29 on the +ad hoc qwen-fp8 container (the same fork, same payload, same TP=2): + +- MNBT 8192 vs the default 2048: **neutral** on decode at every depth. +- Small values re-chunk prefill only. They hurt TTFT (each chunk is its own + iteration, so the first token waits on more iterations) and they never touch + decode: a decode step is `max_num_seqs x (spec+1)` tokens, independent of + MNBT. + +**There is no `--ubatch-size` equivalent in vLLM.** `--ubatch-size` is a +llama.cpp compute-buffer concept (`--batch-size`/`--ubatch-size`); the tcclaviger +fork has no ubatch knob at all. The llama.cpp compute-buffer pair that frank-small +was doubled on (2026-10-05, +4.6 % prefill at 8,928 tokens for 616 MiB) has no +counterpart here, and the two backends should not be reasoned about by analogy. + +**MNBT is not a lever on this shape either.** The deployed 2048 is the fork's +default, so a change is a deliberate departure from the shipped value, not a +tune. The one thing MNBT does cost: every change re-enters the compile RANGE +path, so `FRANK_PRIMARY_CACHE_DIR` pays one recompilation. + ## See Also - [[VLLM Index]] diff --git a/serverAppConfigs-docs/LLM/frank-large-halogen27b-dflash2.md b/serverAppConfigs-docs/LLM/frank-large-halogen27b-dflash2.md index 581f237..1d998cb 100644 --- a/serverAppConfigs-docs/LLM/frank-large-halogen27b-dflash2.md +++ b/serverAppConfigs-docs/LLM/frank-large-halogen27b-dflash2.md @@ -48,7 +48,7 @@ bearing numbers: - **`HALOGEN_MAX_TOKENS_CAP=65536` + `HALOGEN_QUEUE_TIMEOUT=7200`** — coupled pair (FLAGS.md table): a request at the cap must not outlive the timeout. 65536 > the harness's 57,344 `maxTokens` so an eager request never meets - the engine's 413. + the engine's 400. - **KV math:** 1 slot x 200,700 x 64 KiB/token (d2 KV = 2x fp8) = 12.2 GiB pool, inside the ~65 GiB left after 33.4 GiB mmap'd weights + engine working set + host baseline on the 121.2 GiB box. The engine sizes the pool itself @@ -56,17 +56,23 @@ bearing numbers: that ledger, not a formula. - **No server-side sampler:** an absent temperature is GREEDY. The card's thinking recipe (temp 1.0 / top_p 0.95) lives in `models.yml` `extraBody` - — on the wire only. + — on the wire only. The template defaults `reasoning_effort` to **xhigh** + when a request sends none, and **thinking tokens consume the + `max_tokens` budget**: a budget exhausted mid-thought returns + `finish_reason: "length"` with EMPTY `content` and the partial reasoning + in `reasoning_content` — clients must read both fields. ## Client contract deltas vs llama.cpp |Behavior|halogen| |---|---| |Prompt over slot|400 naming the limit (no truncation)| -|`max_tokens` over cap|413, max_tokens over queue timeout|503s for the queue| +|`max_tokens` over cap|400 naming the limit (never silent truncation)| +|full-length request outliving `HALOGEN_QUEUE_TIMEOUT`|503s for everyone queued| +|`timings`|every response carries llama-server-shape `timings` (`draft_n`, `draft_n_accepted`) — the drafter check reads the RESPONSE, and it proves speculation fired| |`response_format`|NOT supported (no constrained decoding)| |thinking field|`reasoning_content` (same as before)| -|`reasoning_effort`|same Qwen3.8 template; outside {xhigh,medium,low} rejected — models.yml fold stays| +|`reasoning_effort`|same Qwen3.8 template; outside {xhigh,medium,low} rejected — models.yml fold stays; unset defaults to xhigh, `none` disables thinking| |tools|server-side tool calls advertised — round-trip MUST be re-probed after deploy| |image input|ignored/error — route to frank-small|