diff --git a/serverAppConfigs-docs/LLM/DFlash DSpark Speculative Decoding.md b/serverAppConfigs-docs/LLM/DFlash DSpark Speculative Decoding.md index ca12461..bb0b2b6 100644 --- a/serverAppConfigs-docs/LLM/DFlash DSpark Speculative Decoding.md +++ b/serverAppConfigs-docs/LLM/DFlash DSpark Speculative Decoding.md @@ -95,6 +95,51 @@ llama-server --hf-repo \ The DFlash drafter proposes, then ngram-mod runs as an additional implementation in the chain. Verified against upstream `common/speculative.cpp` (b10133+). +### Chaining laws (measured on this fleet, 2026-09 → 10) + +`--spec-type` takes a comma list, and the gate is `ctx_dft` in +`common/speculative.cpp` (`common_speculative_init_result`): **draftless** +types (the ngram family — simple, map-k, map-k4v, mod, cache) are added +unconditionally, while **context-backed** types (`draft-mtp`, `draft-dflash`, +`draft-dspark`, `draft-eagle3`) all require a draft CONTEXT. Consequences: + +- A chain may hold ONE context-backed drafter plus any number of ngram types. + `draft-dflash,ngram-mod` is legal and measured: 2.6x the drafter alone on + prompt-echo traffic, byte-identical to it on prose at the fleet's 48/64 + window. +- Two context-backed types cannot share one process: `draft-mtp,draft-dflash` + dies at load with `context type MTP requested but model doesn't contain MTP + layers` — MTP builds its context against whatever the draft context holds, + so it met the DRAFTER's file, not a missing head. Pin that in a seam so + nobody "fixes the flag" later. +- Internal order is fixed (ngram types first, then drafters); CLI order is + irrelevant. +- **A legal chain can still lose on a recurrent model.** `draft-mtp,ngram-mod` + on the qwen35 hybrid (frank-large era) measured −15 % prose / −7 % + block-copy (2026-10-03): adding the ngram layer logs `recurrent state + rollback is not compatible with 'ngram-mod' - disabling rollback support`, + and MTP verification depends on that rollback, so acceptance drops on ALL + traffic even where the matcher fires. Measure chains per architecture, + never by analogy. +- ngram-mod prints NO startup banner — the load log shows only the drafter's + `adding speculative implementation` line. Never conclude "chain not + engaged" from the load log; identify behaviourally (`draft_n` > 0 only + where text repeats) or by the per-task `draft acceptance = ... mean len` + line. It costs ~64 MiB of host-RAM hash table. + +### Drafter KV: read the window policy from the header + +The draft context is built with `cparams.n_ctx = llama_n_ctx(ctx_tgt)` +(`common/speculative.cpp`) — there is no flag to cap it — so a drafter that +declares NO sliding window caches the whole target window on whatever device +`--spec-draft-device` names. DFlash declares `sliding_window = 2048` (~40 MiB +of draft KV); DSpark declares none: 5 layers x 2 x 8 kv heads x 128 dim x 2 B += 20 KiB per position, i.e. **5.0 GiB at 262144** — enough to kill the first +draft graph capture with `ROCm error: out of memory` AFTER the model loaded +and `/health` answered (measured 2026-09-16). Read `sliding_window` out of +the drafter's GGUF header before sizing any spec-decode deployment, and +prefer a skewed `-ts` over a context cut when the full window matters. + DFlash on llama.cpp: ``` @@ -209,6 +254,151 @@ docker run --rm --name vllm-qwen36-dflash \ History: before `--spec-default`, the box ran DSpark (`draft-dspark`, retired 2026-08-15 after a ROCm illegal-memory-access crash under two concurrent 16k prompts) and before that bundled MTP (`draft-mtp`, retired same day). DFlash was retired on 2026-08-14 (premature-EOS under parallel load). The old DFlash draft GGUF (`Qwen3.6-27B-DFlash-Q8_0.gguf`), the MTP-only drafter (`Qwen3.6-27B-MTP-ONLY-Q8_0.gguf`), and the DSpark draft (`Qwen3.8-27B-DSpark-Q8_0.gguf`) remain in `/home/dank/models/` as reference only. +## A/B protocol on a live llama.cpp endpoint + +Verdicts are **traffic-shaped**: a proposer can only draft text it has seen, +so a single prompt shape answers the wrong question. Use fixed shapes: + +| Shape | Prompt | Who it can favour | +|---|---|---| +| prose | 1k filler + "reply with one sentence" | model-backed drafters only | +| echo | ~1.5k passage of UNIQUE lines + "repeat verbatim" | ngram (output IS the prompt) | +| structured | repetitive output, no prompt overlap | ngram, if spans really repeat | +| identity | short deterministic question | controls for lossless verification | + +Design the echo passage so no long span repeats INSIDE it — otherwise you +measure the matcher against itself. For `structured`, give each line a unique +id: repeated lines make it a negative control, not a repetition test. + +Protocol: + +1. Stop the unit (`sudo systemctl stop `), `pkill -x llama-server`; wait + for the port (a retired predecessor's published-port proxy leaves the + endpoint port in TIME_WAIT — the bind fails with address-in-use long after + the unit stopped). +2. Per config: write an env file with the overrides, start the serve script + with `FRANK_*_ENV_FILE`, poll `/health`, measure. Ready-made harness: + `scripts/bench-spec-ab.py `. +3. **Restore the unit in a `finally`** and shout if it does not come back + healthy. A crashed harness that leaves production down is worse than no + harness. +4. **Prove losslessness**: one fixed prompt at temp 0 to every config, diff + the greedy text. Speculation only buys speed; differing texts are a + correctness bug. Empty answers (`content: null` when reasoning ate the + budget) make the comparison vacuous — bench clients must fall back to + `reasoning_content`, and echo/structured prompts should disable thinking + per request: `chat_template_kwargs: {"enable_thinking": false}`. +5. **Sweep the depth, not just the mode** — draft depth is model-specific and + single-valued; deeper often loses (measured: at n=7 a third of proposals + survived, at n=3 more than half). +6. **Sweep the ngram window in the same run as the configs**: too short a + `--spec-ngram-mod-n-min/-n-max` proposes on ordinary prose fragments — + rejected, verification still paid (measured 43.9 → 36.0 tok/s). + +Read the right counters — per request `timings.draft_n` / `draft_n_accepted` / +`predicted_per_second`; cumulative `/metrics` +`llamacpp:spec_decode_num_drafts_total` / `..._accepted_tokens_total`; +per-task `draft acceptance = ... (N accepted / M generated)`. `draft_n == 0` +while tokens flow = the proposer is NOT engaged, whatever the flags say. A +zero-draft config is a useful negative control: it proves the other config's +drafts came from the proposer under test. + +Harness hygiene: absolute script paths (a relative path plus `cd` to its own +dir double-counts the prefix); `pkill -x`, never `-f` (the pattern matches +your own command line); nested f-strings with backslash escapes are a +SyntaxError on the fleet's Python (<3.12 semantics) — build log lines with +concatenation. Report `decode tok/s` per shape plus acceptance and mean +accepted length; a config that wins on echo and loses on prose is a traffic +decision, not a winner. Wire the winner into the deployed env file with the +measured table as its justification comment; record losing numbers with "do +not switch without re-measuring". + +The A/B vehicle for ANY single-GPU llama.cpp knob (batch/ubatch, KV type, +slot split, quant rung — not just spec flags) is the transient systemd +drop-in procedure in [[frank-large-benchmark]] and [[frank-small-benchmark]]. + +## Is it worth it? (arithmetic before deployment) + +Answer against the INCUMBENT, not against zero. + +**Step 0 — check whether the target already speculates.** Many engines ship a +drafter of their own; then your baseline is `A_incumbent / t_s`, not `1 / t_s`. +halogen-flash-server: MTP head ON by default — read its own accounting, e.g. +`serve_api: mtp 63 tok in 1.31s = 47.97 t/s | 33 rounds, commit 1.88/round`. +llama.cpp: `/metrics` counters and per-request `draft acceptance` lines. + +**Step 1 — measure the two step times, never guess.** Target step `t_s` = 1 / +(client-visible tok/s ÷ commit-per-round); ask for the rate the operator +actually SEES (engine-internal decode-only numbers run ~1.7–2x optimistic). +Drafter step `d_s` from llama.cpp's per-task `eval time` lines when it is a +plain model on the same box. + +**Step 2 — arithmetic:** + +``` +round_time = K * d_s + 1.1 * t_s # K drafts proposed, K+1 verified in ONE target pass +A = (1 - p^(K+1)) / (1 - p) # expected accepted tokens, p = per-position acceptance +rate = A / round_time # compare to the incumbent +break-even : A > rate_incumbent * (K*d_s + 1.1*t_s) +``` + +Worked example that decided a real case: `t_s = 67 ms` (25 tok/s, 1.67 +commit/round), `d_s = 12 ms` (27B IQ3_S drafter on the 3090) → break-even +`p ≈ 0.70` at K=3, `≈0.68` at K=5; a realistic cross-finetune `p = 0.4–0.6` +gives 17–20 tok/s, i.e. a LOSS. A drafter only ~3x faster per token cannot +pay for K sequential forwards. + +**Step 3 — classify the drafter by per-round cost:** + +| drafter | cost per round | needs | +|---|---|---| +| plain model (`-md`/`draft-simple`) | **K sequential forwards** | 10–50x faster than the target | +| n-gram / prompt-lookup (`ngram-*`) | ~0, no model | repeated context (agent edit loops) | +| trained head (`draft-mtp`, EAGLE/DFlash) | **~1 pass** | the head must exist FOR THAT CHECKPOINT | + +Trained heads are target-specific: EAGLE/DFlash consume the TARGET's hidden +states, so a dense-27B DFlash2 is useless for a 512-expert MoE trunk. + +**Step 4 — acceptance reality checks.** Fleet datum: a trained z-lab DFlash2 +at depth 7 accepted ~41 % on its matched target (`draft-dflash 7 → 43.5 +tok/s`); treat 40–60 % as the honest band for a matched pair. Cheapest proxy +for `p` at temp 0 when both models are greedy and share a tokenizer: greedy +continuation agreement on identical prompts, two API calls, no weights moved. +llama.cpp can also synthesize the projection without any draft artifact: +`--spec-synth-rates P0,P1,...` / `--spec-synth-len L`. Speculation pays IN +PROPORTION TO HOW SLOW THE TARGET IS — the same 41 % gave 2.3x on an +18.8 tok/s layer-split target and a loss on a 25 tok/s one; fast targets are +usually sparse MoE, the worst case for drafting. + +**Step 5 — check the attachment surface BEFORE the arithmetic.** Closed +engines refuse external drafters by design: halogen-flash-server's +`DRAFTER_CAP` gates `dflash2` on the loaded checkpoint and an unavailable +name is a deliberate 400 (a silent downgrade would make comparisons lie); +there is no device routing in its loopback token protocol, so a drafter can +never sit on another card. Pure llama.cpp DOES route the draft: +`--spec-draft-device/-devd` (default follows `--device`), `-ngld`, +`-otd`, `-ctkd/-ctvd`, `--spec-draft-n-max` (default 3), `--spec-draft-p-min`. +One process sees only the devices its backends were built with — CUDA+Vulkan +is the combination that shows a CUDA card AND an AMD iGPU (HIP+CUDA cannot +coexist); a CUDA-only build has no `Vulkan0` to draft with. + +**Step 6 — cross-model pairs need a byte-identical tokenizer.** Speculation +compares token IDs, so draft and target must map id → same string. Verify +cheaply from the GGUF headers (first ~8 MB over HTTP, no full download): +parse `tokenizer.ggml.tokens`, print vocab size, merges, `pre`, and sha256 +over the first 2000 tokens — + +``` +vocab=248320 merges=247587 model=gpt2 pre=qwen35 +sha256(first 2000 tokens)=8619c68fe4466d3a # identical => ids are interchangeable +``` + +Identity is necessary, NOT sufficient — the drafter still has to be cheap and +accepted. Report in this order: does the target already draft; break-even +`p` at K=1,3,7 versus the incumbent; the realistic band and resulting rate; +whether the drafter fits (VRAM) and the attachment surface exists; the one +measurement that would settle it. + ## Notes - Speculative decoding helps low-QPS, memory-bound, latency-focused workloads. It helps less at high QPS. @@ -225,4 +415,8 @@ History: before `--spec-default`, the box ran DSpark (`draft-dspark`, retired 20 - [[Dankclaw System Specs]] - [[Dankserver System Specs]] - [[VLLM Overview]] -- [[deployment]] \ No newline at end of file +- [[deployment]] +- [[GGUF Compatibility]] — header parsing and custom tensor-type lock-in +- [[Llama Cpp Serving Sizing]] — slot-pool and drafter KV sizing arithmetic +- [[Llama Cpp Multi-Device Splits]] — where the drafter's KV actually lands +- [[frank-large-benchmark]] and [[frank-small-benchmark]] — the single-GPU env drop-in A/B vehicle \ No newline at end of file diff --git a/serverAppConfigs-docs/LLM/GGUF Compatibility.md b/serverAppConfigs-docs/LLM/GGUF Compatibility.md new file mode 100644 index 0000000..bda69a9 --- /dev/null +++ b/serverAppConfigs-docs/LLM/GGUF Compatibility.md @@ -0,0 +1,123 @@ +--- +type: reference +title: GGUF Compatibility and Runtime Lock-in +description: How to answer "will engine X load this GGUF?" from the file's header alone, why a custom tensor type forces the runtime, and the discipline for pinning a fork build that can actually serve the file. +tags: [gguf, llama.cpp, forks, quantization, speculative] +timestamp: 2026-10-05T00:00:00-07:00 +--- + +# GGUF Compatibility and Runtime Lock-in + +**The law: a GGUF's tensor-type histogram and ftype pick the runtime — read the +header before promising a file to any loader.** Stock llama.cpp, Ollama and LM +Studio refuse custom types at load (`unknown tensor type`); a fork-private type +means exactly one engine can serve the file. Measured on this fleet repeatedly: +`Qwen3.8-27B-Q8_0_ROCMFPX_AGENT.gguf` is ftype 115 with 311 tensors of type 103 +(`GGML_TYPE_Q8_0_ROCMFPX`) — only charlie12345/ROCmFPX reads it. The full type +tables live in [[Llama Cpp Runtimes]] and [[Llama Cpp Code Diff]] +(ggml_type 100–108, ftype families 100–106 / 110–113 / 118). + +## 1. Read the header, do not guess + +Parse the GGUF KV section and histogram the tensor types plus +`general.architecture.file_type` — one python read over the first few MiB, no +model load. Counting types out of the KV section settles a runtime argument in +minutes; repo names and model cards do not. + +## 2. Compare routing variants by tensor mix, not by name + +Two files of the same quant can differ in how many tensors stay at plain Q8_0 +(measured: 187 vs 2 output-head tensors in an `_AGENT` variant of the same +size). That difference is what moves spec-decode acceptance (0.953 vs 0.911) +— the "same size" file is not the same file. + +## 3. Pin the build by commit SHA + +A fork that can read the file today can still crash mid-load: the 2026-08-25 +ROCmFPX builds read the ftype-115 file and then SIGSEGV in +`ggml_compute_forward_set_rows`. Record the SHA in the build script +(`scripts/build-rocmfpx.sh` pins `36e9acd`) and check it out explicitly — never +a branch. + +## 4. Header diff method (candidate vs known-good) + +Answer "will engine X load this GGUF?" by diffing the candidate's header +against a file the engine **already loads** — never by filenames or model cards; +filenames lie about quantization and about MTP/vision heads. + +1. **Get the bytes**: both files, from byte 0, at least 64 MiB. Tokenizer KVs + (151k token strings + merges) routinely eat 5–8 MiB, so a small range + desynchronises the parser. True size from `curl -sIL` → `x-linked-size` + (a bare HEAD `Content-Length` can lie): + + ```sh + curl -sL -r 0-67108863 -o cand.bin \ + "https://huggingface.co///resolve/main/.gguf" + ssh 'head -c 67108864 /path/to/known-good.gguf' > good.bin + ``` + +2. **Parse — and lock the parser against the known-good file first.** Layout + (v2/v3): `magic[4]` `version u32` `n_tensors u64` `n_kv u64`, then KVs, then + tensor infos. Strings are u64 length + bytes. **The value-type enum is NOT + llama.cpp's `gguf.h`:** in the files seen (2026-09-28, unsloth + ajgazin UD + GGUFs) type **8 = STRING, 9 = ARRAY**; llama.cpp's header has STRING=2, + ARRAY=3 and desynchronises immediately. Working table: + + ``` + 0 u8 1 i8 2 u16 3 i16 4 u32 5 i32 6 f32 + 7 bool 8 STR 9 ARRAY(u32 et, u64 n) 10 u64 11 i64 12 f64 + ``` + + Validate before trusting: the parse is right iff the tensor-info block ends + somewhere sane and names look real (`blk.0.attn_q.weight`, + `token_embd.weight`, `output_norm.weight`). EOF throw = wrong enum or + offset — do not work around it. + +3. **Diff what engines actually bind on**: `general.architecture`; tensor + count; tensor name list in order; per-tensor `ggml_type` and dims; geometry + KVs (`.block_count`, `.embedding_length`, `.feed_forward_length`, + `.attention.head_count[_kv]`, `.rope.*`, `.ssm.*`, `.context_length`, + `.nextn_predict_layers`); MTP tensors (`nextn` in the name); vision tensors + (`visual`, `v_` prefixes). A same-family fine-tune quantized by the same + tool usually comes out byte-identical in every one of these — that is the + answer. + +4. **Then read the engine's source before calling a delta a blocker.** A + metadata difference only matters if the engine reads that key + (`grep -rn "file_type\|recurrent_layers\|full_attention_interval" /`). + Worked example (gufo v0.1.1): `general.file_type` is deliberately not + consulted (UD-*_K_XL ships a stale ftype 15; the tensor table rules), and + `qwen35.attention.recurrent_layers` is never read (layer kinds come from + `full_attention_interval`) — both deltas were non-issues the header diff + alone would have flagged wrongly. + +5. **Say what you did not do.** Header analysis is strong evidence, not a + load. State byte sizes, check free disk before proposing a fetch, and do + not download tens of GB unasked — pointing the engine at the file is the + definitive test. + +**Report shape:** verdict first (loads / will not / loads-with-caveats), then +the side-by-side diff table, then only the load-bearing caveats: text-only vs +vision tensors, chat-template differences (matters for llama.cpp, not for +engines with compiled-in Jinja), license, disk headroom, and whether the base +model is abliterated/uncensored. + +## 5. A/B a fork's spec decode by stripping the whole flag set + +`--model-draft` without `--spec-type` segfaults at load (measured 2026-10-02, +core dump, no diagnostic). A baseline "no-draft" arm must remove `--model-draft` +AND every `--spec-draft-*` AND `--spec-type` together — the flags move as a set. + +## 6. Before deleting a retired payload, read the rollback tag's env file + +A rollback may need a small companion dir: a drafter living in a directory +named after an old quant looks like trash but is load-bearing. Same class as +the gufo payload removal of 2026-10-03, which turned that rollback into a +re-download. + +## Related + +- [[Llama Cpp Runtimes]] — which fork ships which custom types +- [[Llama Cpp Code Diff]] — the enum tables and code-level differences +- [[frank-large-rocmfpx-q8agent-mtp]] — the ftype-115 lock-in in production +- [[Llama Cpp Serving Sizing]] — pricing what a payload costs once it loads diff --git a/serverAppConfigs-docs/LLM/Llama Cpp Multi-Device Splits.md b/serverAppConfigs-docs/LLM/Llama Cpp Multi-Device Splits.md new file mode 100644 index 0000000..eb13c17 --- /dev/null +++ b/serverAppConfigs-docs/LLM/Llama Cpp Multi-Device Splits.md @@ -0,0 +1,274 @@ +--- +type: reference +title: Llama Cpp Multi-Device Splits +description: What llama.cpp can and cannot express across devices, the no-collective layer-split path, split balancing by measurement, drafter KV sizing from GGUF headers, and heterogeneous placement triage. +tags: [llama.cpp, multi-gpu, layer-split, pcie, sizing, speculative] +timestamp: 2026-10-05T00:00:00-07:00 +--- + +# Llama Cpp Multi-Device Splits + +## 1. Expressibility law — classify the arrangement before running anything + +llama.cpp expresses exactly two things across devices, and they are not +interchangeable: + +| arrangement | how | expressible? | +|---|---|---| +| **Layer split** — layers 0..n on dev A, rest on dev B | `-sm layer -ts /` | **yes**, no collectives needed; works cross-vendor (CUDA + Vulkan measured 2026-09-16) | +| **Per-tensor placement** — attention+KV on A, experts on B | `-ot =` | **NO.** `--override-tensor` sets `has_tensor_overrides()`, a REQUIRED term of the scheduler's `pipeline_parallel` flag (only read site: `llama-context.cpp`, in the `pipeline_parallel = …` expression). That flag sets `n_copies = parallel ? MAX : 1` in `ggml_backend_sched_new`; with `n_copies == 1` a pre-allocated weight on a non-primary device has no backend and the graph aborts at `ggml-backend.cpp:941` (`pre-allocated tensor (blk.0.ffn_down_exps.weight) in a buffer (Vulkan1) that cannot run the operation (NONE)` — `(NONE)` is the leaf's own op name). | + +Failure modes observed (mainline `37b53fd`): per-tensor `-ot exps=Vulkan1` +with `-dev CUDA0,Vulkan1` → scheduler abort; layer-scoped `-ot` with a +cross-backend pair → model load fails with only `failed to load model`; +`-ot x=CPU` **works** because the CPU backend is last in the list — **the CPU +is the only legitimate partner in a heterogeneous split.** On a CUDA+AMD host +a card can be the trunk OR the offload host, never the complement of an +AMD-resident peer. Pure Vulkan changes nothing here — the gate is the +override, not the driver mix. + +If the plan needs tensors on a specific device: stop, the answer is no for +any vendor mix. Expressible alternatives, in order: (1) two independent +endpoints, one model per device; (2) single-vendor layer split; (3) device + +CPU per-tensor offload (mainline pinning guard: `GGML_CUDA_NO_PINNED`; ik fork +was `GGML_CUDA_NO_PINNED_WEIGHTS` — without it a CPU-overridden tensor can +land in a pinned `CUDA_Host` buffer and page-lock ~108 GiB). + +Layer-split arithmetic to check before promising: a whole layer carries its +experts (MoE) — compute GiB/layer, not per param-group; only some layers carry +KV (every 4th on the hybrid), so a share of layers is only a share of KV; big +non-repeating tensors (PLE tables) land on the main device and cannot move — +they cap how small the primary can get. + +## 2. No collectives? The atomics wall and the layer-split answer + +Decide with one line before designing: + +```bash +sudo lspci -vv -s | grep AtomicOpsCtl # ReqEn- ⇒ no collectives, ever +lspci -tv # a Routing- bridge is a hardware wall +``` + +`ReqEn-` kills vLLM entirely — TP needs a per-layer allreduce AND PP needs a +device-backend `send`/`recv` of the hidden state +(`vllm/distributed/parallel_state.py`) — i.e. the same RCCL primitive that +fails on an atomics-less card. llama.cpp's `--split-mode layer` needs no +collective: layers distribute in order, only the boundary activation crosses +PCIe once per forward pass. Full probe/repair detail: +[[PCIe Link and Atomics]]. + +## 3. Build: pin by feature, bundle the runtime at the exact depth + +- Pin the ref by FEATURE list (e.g. DFlash2 + qwen35 MTP on 2026-09-16 → + `ggml-org/llama.cpp#27342`); never float. Verify a fork build the same way + (`ls /…/spec_decode/ # dflash/ vs dflash2/`). +- `-DCMAKE_HIP_ARCHITECTURES=` — a sibling arch links, loads, then + segfaults. `-DGGML_HIP_RCCL=OFF` when there is no atomics path (only tensor + mode degrades to the butterfly fallback; layer split unaffected). +- **TRAP — `cp -r SRC DST/` flattens when `DST` does not exist** and drops one + directory level. rocBLAS/hipBLASLt look for device code at a FIXED depth + (`../lib/rocblas/library`, `../lib/hipblaslt/library/`); flattened, + the engine loads every GiB, answers `/health`, and dies on the FIRST request + with `Cannot read .../TensileLibrary.dat: No such file or directory for GPU + arch`. `mkdir -p DST && cp -r SRC/. DST/`, then + `patchelf --set-rpath '$ORIGIN/../lib' build/bin/llama-*`. +- Dual-backend build (CUDA + Vulkan is the practical pair; HIP+CUDA do not + coexist): `-DGGML_CUDA=ON -DGGML_VULKAN=ON` in one tree; Vulkan needs + `glslc` + libvulkan headers. **Device names are not what you expect:** + `--list-devices` lists the discrete card under BOTH backends (`CUDA0` AND + `Vulkan0`) and the APU LAST (`Vulkan1`) — a wrong `-ot` name makes + llama.cpp print the valid set and exit (self-diagnosing). +- What the existing build can do: `GPU_BACKEND_*=OFF` lines in its + CMakeCache. + +## 4. Serve flags that matter on a split + +```bash +HIP_VISIBLE_DEVICES=0,1 llama-server \ + --device ROCm0,ROCm1 --split-mode layer --tensor-split 1,1 \ + --gpu-layers 999 --fit off --flash-attn on \ + --ctx-size --parallel 1 --cache-type-k f16 --cache-type-v f16 \ + --jinja --metrics --cache-prompt \ + --spec-type draft-dflash --spec-draft-model \ + --spec-draft-device ROCm1 --spec-draft-n-max 7 +``` + +- `--fit off` + explicit `--ctx-size`: auto-fit must not resize a context set + on purpose. **`--parallel N` DIVIDES `--ctx-size`** — one slot at large ctx + is the concurrency policy, concurrent agent loops queue. +- Pin the drafter to the card holding the target's LAST layers + (`--spec-draft-device`); under tensor mode an unpinned draft joins the + meta-device KV allocation and aborts the KV build. `--spec-draft-n-max` must + be ≤ the drafter's `block_size - 1` (DFlash2: 8 ⇒ 7). +- Qwen3.8-family templates RAISE on `reasoning_effort=high` (500) — probe + low/medium/high/xhigh and map harness levels before declaring ready. + +## 5. Split balancing: measure, never interpolate + +**A `-ts` step is ONE LAYER, and hybrid stacks do not have uniform layers.** +qwen35 is 16 × (3 Gated DeltaNet + 1 Gated Attention): one step ≈ 0.43 GiB on +a DeltaNet boundary, ≈ 1.43 GiB on a full-attention one (weights + that +layer's own f16 KV) — measured on frank-medium. The gap can jump 2.9 GiB in a +single step; the best integer split can sit 0.2 GiB from balance while its +neighbours are 1+ GiB off. An affine fit reproduces the deployed point and +overstates gaps in between (`tests/frank_medium_unit.bats` encodes the fit as +a gate; treat a NEGATIVE fitted constant as a unit bug). + +Knobs that move weight and REQUIRE a re-measure: `--ubatch-size` (compute +buffers are **~2x heavier on the output/drafter card** — 512→1536 added +~3.3 GiB there vs ~1.7 GiB elsewhere), a drafter swap (weights + draft KV), +`--ctx-size`/`--parallel` (KV). Total free VRAM is invariant across splits — +only the distribution moves; use that to sanity-check a suspicious row. + +**Procedure (one split per load):** + +```bash +WORK=/tmp/split.env; cp scripts/serve-.env $WORK +export FRANK_MEDIUM_ENV_FILE=$WORK +for ts in 39,26 38,27 37,28 36,29 35,30; do + sed -i "s|^TENSOR_SPLIT=.*|TENSOR_SPLIT=$ts|" $WORK + nohup bash scripts/serve-.sh >/tmp/split.log 2>&1 & + # poll /health (or process exit), then: + rocm-smi --showmeminfo vram | grep 'VRAM Total Used' # sample ~6 s after health + pkill -x llama-server; sleep 6 +done +``` + +Load → sample → kill, NO requests sent (allocation only). Run under `nohup` +and poll from a separate call (a load is 30–90 s). Pick the smallest +inter-card gap that clears the per-card floor with room to spare; one `-ts` +step ≈ 1.5 GiB at ~6k batch / 262k ctx on 30 GiB cards — enough to bracket, +not to solve. Prefer balance: the busy card's headroom bounds the context +(`ctx_extra ≈ (floor − measured_free) / per-position cost on the busier +card`). A faulted split (`SDMA0 UTCL2 page fault` → `illegal memory access` → +`ggml_abort`) is OFF-LIMITS — re-measure, never re-tune; more headroom does +not prevent a boundary-copy fault. + +## 6. Drafter KV: sized from the TARGET, bounded by its window policy + +`common/speculative.cpp` builds the draft context with +`cparams.n_ctx = llama_n_ctx(ctx_tgt)` — **no flag caps it.** So: + +``` +draft KV = ctx_tgt × block_count × 2 (K,V) × kv_heads × key_length × dtype +``` + +…unless the drafter declares a sliding window. Measured at 262,144 target +ctx: DFlash2 (`sliding_window = 2048`) ≈ 40 MiB; a same-shape 5-layer drafter +with NO window = 20 KiB/position = **5.0 GiB** — a four-layer swing in the +balance point, and swapping them invalidated an already-balanced split. + +**Failure signature:** the engine loads the whole model, answers `/health`, +and dies on the FIRST draft decode with `ROCm error: out of memory` / +`cudaMalloc failed` inside `common_speculative_impl_draft_dflash::process` — +a sizing failure, not a broken drafter; under a self-healing unit it presents +as a crash-loop. Read `sliding_window`, `block_count`, `head_count_kv`, +`key_length` from the drafter's GGUF header first. A compact reader (the +three format traps: no KV alignment — v3 alignment applies only to the tensor +section; arrays are `type(9) → element_type(u32) → count(u64)`, element type +FIRST; the enum is `7=BOOL 9=ARRAY`, 64-bit = 10/11/12, and 8=STRING/9=ARRAY +is the v2/v3 file order, NOT gguf.h's 2/3): + +```python +import struct, sys +def rd(path): + f = open(path, 'rb'); assert f.read(4) == b'GGUF' + struct.unpack('