--cache-ram caps the total bytes of saved prompt states in HOST RAM, and alloc() refuses any single state larger than the whole cap. That refusal was firing in production: 137 times in the two days from 2026-09-16, at 11.0-11.4 GiB per state, so every conversation past ~170k tokens was uncacheable and re-prefilled from scratch on its next turn. The state law at this shape is 68 KiB/token (16 full-attention layers x 4 kv heads x 256 dim x 2 B x 2 for K/V, plus the MTP draft context's 4096 B), so 8192 MiB held ~120k tokens and 16384 holds ~240k: a FULL 196608-token slot state (~12.8 GiB) now fits with ~3 GiB spare. Host RAM only, allocated on demand and never pinned, so the VRAM fit is untouched. New seam: the cap must clear one full slot's state and stay under half the host's RAM (it fails on the old 8192 and passes on 16384). Also corrects the adjacent claim that 0 is the build's default for --ctx-checkpoints -- this binary's --help says 32.
708 lines
39 KiB
Markdown
708 lines
39 KiB
Markdown
---
|
||
type: configuration
|
||
title: frank-medium llama.cpp Runbook (MTP drafter)
|
||
description: The ACTIVE frank-medium deployment since 2026-09-16 — mainline llama.cpp on Dankserver serving ajgazin/Swift-Qwen3.8-27B-Uncensored-Dynamic-MTP at UD-Q8_K_XL with the model's own MTP head (the z-lab DFlash2 sidecar kept on disk, one env change away), as ONE layer-split process across both Radeon AI PRO R9700 cards at a 35:30 split (skewed toward the x16-slot card; see the 2026-09-17 part 3 section), a 393216-position KV pool (196608 per conversation over two slots), vision via the BF16 mmproj.
|
||
tags:
|
||
- frank-medium
|
||
- dankserver
|
||
- llama-cpp
|
||
- dflash2
|
||
- dspark
|
||
- qwen35
|
||
- gfx1201
|
||
- r9700
|
||
- layer-split
|
||
timestamp: 2026-09-17
|
||
---
|
||
|
||
# frank-medium — llama.cpp (DFlash2 drafter)
|
||
|
||
The ACTIVE frank-medium deployment since 2026-09-16: **mainline llama.cpp**
|
||
serving `ajgazin/Swift-Qwen3.8-27B-Uncensored-Dynamic-MTP-GGUF` at `UD-Q8_K_XL`
|
||
with the **z-lab DFlash2** block-diffusion drafter, as **one process across both
|
||
Radeon AI PRO R9700 cards**, layer-split. DSpark was tried in the drafter's
|
||
place for a few hours the same day and reverted — see *The DSpark excursion*.
|
||
It replaced the vllm-radiance MXFP4 replica pair + routing proxy
|
||
([[frank-medium-radiance-mxfp4]], the record and the PCIe-atomics evidence) in
|
||
the same day the pair went live. It is also the box's **only** LLM runtime: the
|
||
replicas, the proxy and their unit files are retired.
|
||
|
||
## Why one process and layer split
|
||
|
||
The endpoint's whole history is one constraint: **the two cards cannot take
|
||
part in an RCCL collective.** RCCL's kernels declare a hidden hostcall buffer,
|
||
and ROCr refuses to supply it to a device without PCIe atomics; the second
|
||
R9700 hangs off the B650 chipset (`AtomicOpsCap: Routing-`, read-only), so
|
||
every collective that touches it dies with `hipErrorIllegalState`. That rules
|
||
out tensor parallelism — and it rules out **vLLM's pipeline parallelism too**,
|
||
whose stage handoff is a `send`/`recv` of the hidden state on the device
|
||
backend (`vllm/distributed/parallel_state.py`), the very primitive that fails.
|
||
The full matrix is in [[Dual-R9700-Cooperation]].
|
||
|
||
llama.cpp's `--split-mode layer` needs **no collective at all**: layers are
|
||
distributed across the cards in order and the boundary activation is a plain
|
||
per-backend copy, once per forward pass. Measured on this host with the
|
||
predecessor runtime (llama-bench, pp4096 / tg128, f16 KV, flash-attn on):
|
||
|
||
| Split | pp4096 | tg128 |
|
||
|---|---:|---:|
|
||
| one R9700 (`-sm none`) | 1160.9 | 19.42 |
|
||
| **two cards, layer (`-sm layer`)** | **1991.2** | **18.76** |
|
||
| two cards, tensor (`-sm tensor`, no RCCL: butterfly fallback) | 271.8 | 26.26 |
|
||
|
||
Tensor split buys ~40 % decode and loses 7× prefill. The workload here is
|
||
prefill-heavy (agentic, long prompts, reasoning traces), so layer split is the
|
||
shape — prefill nearly doubles and one model serves from both cards.
|
||
|
||
## Hardware
|
||
|
||
| Property | Value |
|
||
|---|---|
|
||
| GPUs | 2x AMD Radeon AI PRO R9700 (DID 0x7551), gfx1201 (Navi 48, RDNA4) |
|
||
| VRAM | 34,208,743,424 B (31.86 GiB) each, as `rocm-smi` reports it |
|
||
| Link | PCIe, no XGMI; card #2 behind the chipset, **no atomic routing** |
|
||
| iGPU | Raphael (gfx1036) enumerates as HIP device 2 — hidden via `HIP_VISIBLE_DEVICES=0,1` |
|
||
| CPU / RAM | Ryzen 7 7700X (8c/16t), 60 GiB RAM, 63 GiB swap |
|
||
|
||
## Runtime
|
||
|
||
Mainline llama.cpp at **`fb27a525`** (2026-09-16), built for **gfx1201 only** by
|
||
`scripts/build-llama-rdna4.sh` against a private ROCm `7.14.0a20260624`
|
||
toolchain (TheRock nightly) that is **bundled next to the binary** —
|
||
`RPATH=$ORIGIN/../lib`, no `LD_LIBRARY_PATH`, nothing from the host's ROCm 7.1
|
||
packages.
|
||
|
||
Two features pin that ref, and neither is in the ROCmFPX fork the endpoint used
|
||
earlier the same day (its tip, `c49ebdb`, predates both):
|
||
|
||
- **DFlash2** (`ggml-org/llama.cpp#27342`): grouped dynamic depthwise
|
||
convolution plus the candidate selector, detected from the draft GGUF
|
||
(`selector_top_k > 0`) and served through the ordinary
|
||
`--spec-type draft-dflash` path.
|
||
- **qwen35 MTP**: the `draft-mtp` implementation covers this architecture's
|
||
single trained head.
|
||
|
||
Build with:
|
||
|
||
```bash
|
||
bash scripts/build-llama-rdna4.sh # ~25 min on this box, JOBS=$(nproc)
|
||
```
|
||
|
||
> **Bundling trap, measured.** rocBLAS and hipBLASLt look for their device code
|
||
> at a fixed depth under the binary's lib dir (`../lib/rocblas/library`,
|
||
> `../lib/hipblaslt/library/<arch>`). `cp -r SRC DST/` with a missing `DST`
|
||
> creates `DST` as a copy of `SRC` and **flattens it one level**, and the
|
||
> failure is invisible at startup: the engine loads 29 GiB of weights, answers
|
||
> `/health`, and dies on the **first request** with
|
||
> `rocBLAS error: Cannot read .../lib/rocblas/library/TensileLibrary.dat: No
|
||
> such file or directory for GPU arch : gfx1201`. The build script now declares
|
||
> the destination and copies contents.
|
||
|
||
## Model and drafter
|
||
|
||
| File | Bytes | Role |
|
||
|---|---:|---|
|
||
| `Swift-Qwen3.8-27B-Uncensored-Dynamic-MTP-UD-Q8_K_XL.gguf` | 31,457,990,112 | target, 29.29 GiB |
|
||
| `Qwen3.8-27B-DFlash2-Q8_0.gguf` | 2,056,414,816 | DFlash2 drafter, 1.92 GiB |
|
||
|
||
Fetched by `scripts/download-frank-medium-gguf.sh` (resumable, byte-verified —
|
||
a truncated 31 GB GGUF is otherwise a load-time crash).
|
||
|
||
The target is Unsloth Dynamic 3.0 (per-tensor types from Unsloth's own quant of
|
||
the same size, on Unsloth's importance matrix) of an abliterated Swift
|
||
fine-tune: orcarouter's single refusal direction projected out of 131 tensors,
|
||
vision tower included. Vision has been ON since 2026-09-17
|
||
(`--mmproj`); this is one of the fleet's two image endpoints.
|
||
|
||
**The file also carries the model's MTP head** (`blk.64.nextn.*`), and since
|
||
2026-09-17 that is the deployed drafter: `--spec-type draft-mtp` builds the
|
||
draft context against the main model's own weights (`shares_model`), so there
|
||
is no second copy of the model and no `--model-draft` argument — the cost is
|
||
one nextn layer's KV (4096 B/position) plus small buffers on the card that
|
||
holds the last layers (ROCm1). The DFlash2 sidecar stays on disk and in the
|
||
fetch script: `SPEC_TYPE=draft-dflash` (with the sidecar back on the command
|
||
line) brings the measured-faster chain back, at the price of ~1.1 GiB on ROCm1.
|
||
|
||
## Deployment
|
||
|
||
| File | Role |
|
||
|---|---|
|
||
| `scripts/frank-medium.env` | every tunable, with the sizing arithmetic in comments |
|
||
| `scripts/serve-frank-medium.sh` | env-driven llama-server launcher; no arguments |
|
||
| `scripts/frank-medium.service` | ROOT systemd unit, `User=dank`, `WantedBy=multi-user.target` |
|
||
|
||
```bash
|
||
./scripts/deploy.sh dankserver --service frank-medium
|
||
```
|
||
|
||
The handler verifies the runtime and payload exist (and refuses otherwise),
|
||
retires the three vLLM quadlets into `~/.AGENTTRASH/20260916-vllm-radiance-retire/`,
|
||
waits for **port 7770 to be released** (the retired proxy held it on the host
|
||
network), relinks and restarts the unit, and gates on
|
||
`http://127.0.0.1:7770/health`.
|
||
|
||
Deployed flags:
|
||
|
||
| Flag | Value | Why |
|
||
|---|---|---|
|
||
| `--split-mode` / `--device` / `--tensor-split` | `layer` / `ROCm0,ROCm1` / `35,30` | both cards, one model, no collective; the split is skewed because ROCm1 carries the drafter, and it is one LAYER of granularity — 0.43 GiB on a DeltaNet boundary, 1.43 GiB on an attention one. 35:30 since 2026-09-17 (part 3): ROCm0 is CPU-root-port Gen4 x16, ROCm1 rides the chipset at Gen4 x1, and the drafter-bound card on the x1 path ran 98% against the x16 card's 40%. 36:29 fails both the 1 GiB floor (card 0 free 0.45) and the fit slack (gap 2.22) — this is the farthest safe shift |
|
||
| `--gpu-layers` / `--fit` | `999` / `off` | everything on the cards; auto-fit must not resize a context we set |
|
||
| `--ctx-size` / `--parallel` | `393216` / `2` | the pool over two slots of 196608; `--parallel` **divides** the context. 1.5x the trained 262144 window, and the GGUF has no rope scaling — see the 2026-09-17 section |
|
||
| `--batch-size` / `--ubatch-size` | `6144` / `512` | measured: ubatch 512 is the prefill optimum *and* the cheapest point on the curve; the logical batch is throughput- and VRAM-neutral. See [[frank-medium-batch-ubatch]] |
|
||
| `--ctx-checkpoints` / `--cache-ram` | `0` / `16384` | no per-slot KV snapshots (the build's default is **32** — this is a deliberate deviation, not an inherited default); the RAM-side prompt cache is capped at 16 GiB, which is what lets a full 196608-token slot state be cached at all (2026-09-18, see below) |
|
||
| `--flash-attn` / cache types | `on` / `f16` K and V | an agentic checkpoint is the wrong place to spend KV precision |
|
||
| `--spec-type` / `--spec-draft-n-max` | `draft-mtp` / `3` | the model's own head, `shares_model` — no sidecar file, no second model copy; n_max 3 is the measured MTP winner (32.4 tok/s vs 26.7 at n_max 7). The DFlash2 sidecar stays on disk: `SPEC_TYPE=draft-dflash` brings the 43.5 tok/s chain back at ~1.1 GiB on ROCm1 |
|
||
| `--spec-draft-device` | `ROCm1` | the drafter rides the card that ends up holding the target's last layers |
|
||
| `--temp` / `--top-p` / `--top-k` | `0.7` / `0.95` / `20` | the values the vLLM predecessor served on this endpoint |
|
||
| `--reasoning-format` / `--reasoning-budget` | `deepseek` / `-1` | trace into `reasoning_content`, unlimited |
|
||
| `--mmproj` | `mmproj-BF16.gguf` (888 MiB, beside the target) | vision enabled since 2026-09-17; this build has NO autoload, so the flag is explicit and the serve script refuses to start without the file |
|
||
|
||
## The DSpark excursion (2026-09-16, tried and reverted)
|
||
|
||
The endpoint went live on **z-lab DFlash2**, was switched to
|
||
**magnitudedev/Qwen3.8-27B-DSpark** for a few hours the same day, and was put
|
||
back. Both are 5-layer block-diffusion sidecars, and their differences are what
|
||
decided it:
|
||
|
||
| | DFlash2 (deployed) | DSpark (tried, reverted) |
|
||
|---|---|---|
|
||
| file | 2,056,414,816 B — 1.92 GiB | 1,455,376,576 B — 1.36 GiB |
|
||
| draft depth | `block_size 8` -> `n_max 7` | `block_size 7` -> `n_max 6` (llama.cpp asserts `n_max <= block_size - 1`) |
|
||
| attention | `sliding_window 2048`, `causal false`, candidate selector (rank 256, top-k 16) | **no sliding window**; a Markov head in place of the selector |
|
||
| draft KV at 262144 | ~40 MiB — the window bounds it | **5.0 GiB** — the whole context |
|
||
|
||
> **A fetch-script bug this threw up.** The DSpark file's size was recorded as
|
||
> 1,359,284,737 when it was fetched — 96 MB short of the Hub's 1,455,376,576 —
|
||
> and that wrong number became the fetch script's declared size, so
|
||
> `--verify-only` failed on a payload that was in fact byte-correct. Both
|
||
> drafters' sha256 have since been checked against the Hub's LFS oids and match.
|
||
|
||
**The draft context is sized from the target's.** `cparams.n_ctx =
|
||
llama_n_ctx(ctx_tgt)` in `common/speculative.cpp`, with no flag to cap it
|
||
independently, so a drafter that declares no sliding window caches the entire
|
||
target window on whichever device `--spec-draft-device` names. DSpark's KV is
|
||
5 layers x 2 x 8 kv heads x 128 dim x 2 B = 20 KiB per position.
|
||
|
||
That cost has to be paid somewhere, and the failure was instructive: at a 1:1
|
||
split the model **loaded fine, answered `/health`, and then died on the first
|
||
draft decode** with `ROCm error: out of memory` inside
|
||
`common_speculative_impl_draft_dflash::process` — ROCm1 sat at 33.9 GB of
|
||
34.2 GB. A unit that crash-loops with an allocation failure *after* a
|
||
successful load is a sizing failure, not a broken drafter.
|
||
|
||
Read `sliding_window` out of a drafter's GGUF header before sizing a
|
||
spec-decode deployment — 40 lines of `struct` parsing is enough, and the
|
||
flag-level configuration looks identical either way.
|
||
|
||
### The split moved twice: once for the batch, once for the drafter
|
||
|
||
**Under DSpark.** Tripling the batch (2048/512 -> 6144/1536) made the two cards
|
||
diverge: at 262144 the bigger `--ubatch-size` cost **ROCm1 ~3.3 GiB against
|
||
ROCm0's ~1.7 GiB** — ROCm1 holds the output layers *and* the drafter's graphs —
|
||
which left the draft-carrying card at 0.08 GiB free against ROCm0's 3.78 under
|
||
the 38:27 split that had been balanced at the smaller batch:
|
||
|
||
| `-ts` | GPU 0 free | GPU 1 free |
|
||
|---|---:|---:|
|
||
| 38,27 | 3.78 GiB | 0.08 GiB |
|
||
| 41,24 | 1.47 GiB | 2.38 GiB |
|
||
| 40,25 | 1.88 GiB | 1.88 GiB |
|
||
|
||
**Then DSpark came off**, which frees ~4.3 GiB on ROCm1 (5.0 GiB of draft KV
|
||
out, 0.65 GiB of heavier DFlash2 weights in) and moves the balance the other
|
||
way. Re-measured at the same batch, the same two slots and the same window:
|
||
|
||
| `-ts` | GPU 0 free | GPU 1 free | gap |
|
||
|---|---:|---:|---:|
|
||
| 39,26 | 2.42 GiB | 7.20 GiB | 4.78 |
|
||
| 38,27 | 3.85 GiB | 5.77 GiB | 1.92 |
|
||
| 37,28 | 4.28 GiB | 5.34 GiB | 1.06 |
|
||
| **36,29 — deployed** | **4.71 GiB** | **4.91 GiB** | **0.20** |
|
||
| 35,30 | 5.14 GiB | 4.48 GiB | 0.66 |
|
||
| 34,31 | 6.57 GiB | 3.05 GiB | 3.52 |
|
||
|
||
**One `-ts` step is one layer, and this stack's layers are not the same size.**
|
||
qwen35 is 16 x (3 Gated DeltaNet + 1 Gated Attention), so a step is worth
|
||
**~0.43 GiB** on a DeltaNet boundary and **~1.43 GiB** on a full-attention one
|
||
(0.45 GiB of weights plus ~1.0 GiB of its own f16 KV at 262144). Both fell out
|
||
of that table — 37,28 -> 38,27 moves a DeltaNet layer, 38,27 -> 39,26 an
|
||
attention layer. Balance by measurement, never by interpolating.
|
||
|
||
**Then the batch moved, and the split followed it** (2026-09-16, the same day
|
||
— see [[frank-medium-batch-ubatch]]). 36:29 was the balance at ubatch **1536**;
|
||
the sweep showed 512 is both faster and 3.71 GiB cheaper, and — because the
|
||
compute buffers cost ROCm1 nearly twice what they cost ROCm0 — shrinking them
|
||
handed ROCm1 1.31 GiB more than it handed ROCm0:
|
||
|
||
| `-ts` | ubatch | GPU 0 free | GPU 1 free | gap |
|
||
|---|---:|---:|---:|---:|
|
||
| 36,29 | 1536 | 4.71 GiB | 4.91 GiB | 0.20 |
|
||
| 36,29 | 512 | 6.01 GiB | 7.32 GiB | 1.31 |
|
||
| **35,30 — deployed** | **512** | **6.44 GiB** | **6.90 GiB** | **0.46** |
|
||
|
||
35:30's 0.46 GiB gap is one DeltaNet layer's worth, which is the floor here: no
|
||
split can be tuned finer than a layer, so the 0.20 GiB of the old configuration
|
||
is not reachable at this batch. It clears the operator's 1.5 GiB-per-card floor
|
||
by ~5 GiB, so no headroom is bought out of `CTX_SIZE` and the model's full
|
||
262144-token window stays.
|
||
|
||
A batch of 8192/2048 failed here under DSpark — its 2020 MiB pp buffer could not
|
||
be allocated on ROCm1, which had 0.08 GiB free. That was a *contiguity* failure,
|
||
not a total-bytes one, and the sweep has since removed the reason to care: with
|
||
the batch geometry measured, there is nothing to gain by making either knob
|
||
larger, so the question is closed rather than pending.
|
||
|
||
> **A fault to watch.** A 39:26 split, at 1x batch and one slot, died on its
|
||
> eighth request: the amdgpu driver logged SDMA0 UTCL2 page faults and the
|
||
> engine surfaced `ROCm error: an illegal memory access was encountered` ->
|
||
> `ggml_abort` — a fault in the cross-card boundary copy, which is this layer
|
||
> split's hot path. It has not reproduced (36:29 served the whole verification
|
||
> battery and the e2e gate, and 35:30 at ubatch 512 is the deployed shape now;
|
||
> 39:26 itself was only *loaded* at the deployed batch, not exercised), and it
|
||
> is **not** an OOM: no amount of headroom would prevent it.
|
||
|
||
## 2026-09-17: the pool grows to 327680, the split moves to 36:29, vision comes on
|
||
|
||
The window grew from 262144 to a **327680-position pool** (163840 per
|
||
conversation over the two slots), the split moved one layer to ROCm0
|
||
(35:30 → 36:29), and the **BF16 mmproj** (888 MiB) came on.
|
||
|
||
**Why 36:29, not the balance.** At this context the fit model puts the balance
|
||
point at 34:31 (0.11 GiB gap). The operator moved a layer to ROCm0 to buy
|
||
headroom on ROCm1 — the card that carries the drafter, its KV, the output
|
||
layers and the heavier compute buffers — accepting a 3.15 GiB gap by
|
||
construction:
|
||
|
||
| Card | Fit | Measured at load | Headroom |
|
||
|---|---:|---:|---:|
|
||
| ROCm0 (36 layers + mmproj) | 29.19 GiB | 28.40 GiB | 3.46 GiB |
|
||
| ROCm1 (29 layers + drafter) | 26.04 GiB | 26.66 GiB | 5.20 GiB |
|
||
|
||
Both clear the 1.5 GiB floor with room; the fit lands within 0.8 GiB on both
|
||
cards, which is the accuracy it has always had (at the 35:30 point it
|
||
reproduced ROCm1 to 0.00 and ROCm0 to 0.86). The seam's slack is the 3.15 GiB
|
||
gap rounded up, with the trade written in its comment.
|
||
|
||
**The rope caveat.** The GGUF carries no `context_length` and no rope-scaling
|
||
keys (freq_base 1e7, `dimension_sections [11,11,10,0]`); the checkpoint's
|
||
trained window is 262144. Positions beyond that — the top 65k of the pool —
|
||
extrapolate on the base rope. It runs, and 1.25x is a modest ask of a rope,
|
||
but it is untrained territory, and the seam's ctx cap encodes 327680 with
|
||
that note.
|
||
|
||
**Vision: this build has no mmproj autoload.** The first 327680/36:29 deploy
|
||
came up healthy and *text-only*: image requests answered 500 "image input is
|
||
not supported — hint: … provide the mmproj" and the load log had no mmproj
|
||
line, so `llama-server` fb27a525 does not pick up `mmproj*.gguf` from the
|
||
model directory. The serve script now passes `--mmproj` explicitly (the
|
||
frank-small pattern), refuses to start without the file when vision is on, and
|
||
the deploy handler's preflight checks for it. Verified end-to-end after the
|
||
fix: a generated 64x32 test image was described from its pixels — with
|
||
thinking on the model reasoned about "the three horizontal bars" (its reading
|
||
of the crude glyph) and with thinking off answered the letter question in one
|
||
token.
|
||
|
||
The e2e crawl spec flaked once on the first 327680 deploy (1 failed / 31
|
||
passed) and passed 32/32 on the redeploy; the service itself was healthy
|
||
throughout (health gate green, `NRestarts=0`).
|
||
|
||
## 2026-09-17 (part 2): the drafter switches to MTP, the pool grows to 393216
|
||
|
||
The operator asked for **2x196k contexts**, which the DFlash2 sidecar's
|
||
~1.1 GiB on ROCm1 was eating. The drafter switched to the model's own
|
||
**MTP head** (`SPEC_TYPE=draft-mtp`, `n_max 3` — the measured MTP winner from
|
||
the same-day A/B: 32.4 tok/s at n_max 3 vs 26.7 at n_max 7, against DFlash2's
|
||
43.5), the pool grew 327680 → **393216** (196608 per conversation), and the
|
||
split moved one layer back to **34:31**, the balance point at the new pool.
|
||
|
||
| Card | Fit | Measured at load | Headroom |
|
||
|---|---:|---:|---:|
|
||
| ROCm0 (34 layers + mmproj) | 29.77 GiB | 28.91 GiB | 2.95 GiB |
|
||
| ROCm1 (31 layers + MTP buffers) | 28.32 GiB | 30.78 GiB | 1.08 GiB |
|
||
|
||
The fit predicted 28.32 on ROCm1; the load came in **2.46 GiB hotter** than
|
||
predicted. The MTP charge is one nextn layer's KV (4096 B/position — the
|
||
draft context is sized from the target, same law as the DFlash2 one) plus
|
||
draft buffers the fit only estimated. The seam was recalibrated to the
|
||
measured value: the MTP buffer constant is now 2,901,231,913 B (the 4.20 GiB
|
||
total MTP charge on card 1), and the floor dropped 1.5 → **1.0 GiB** for this
|
||
shape — ROCm1's 1.08 GiB of headroom is real. 35:30 failed card 0's floor
|
||
only at the OLD 1.5 GiB bar (it clears 1.0 with 1.27 GiB to spare); 36:29
|
||
fails both gates. Under the recalibrated 1.0 GiB floor 35:30 became
|
||
reachable — and the split went there the same day (part 3).
|
||
|
||
The rope caveat deepens: the top 131k of a conversation (196608 vs the
|
||
trained 262144… i.e. 393216/2 = 196608 per slot, so the top 131,072 positions
|
||
of a FULL conversation) extrapolates on the base rope. 1.5x is a bigger ask
|
||
than 1.25x; the seam's ctx cap encodes 393216 with that note.
|
||
|
||
Verified after deploy: `adding speculative implementation 'draft-mtp'` with
|
||
`shares_model` in the load log, no `--spec-draft-model` flag, `--mmproj`
|
||
present; 207 drafts / 108 accepted in a text run (52 % acceptance, mean len
|
||
~1.7); vision, tools and `reasoning_effort` all pass; e2e 32/32; `NRestarts=0`.
|
||
|
||
## 2026-09-17 (part 3): the PCIe topology, and the split leans to 35:30
|
||
|
||
`lspci -vv` on dankserver answers why the split now leans toward ROCm0:
|
||
the two cards do not sit on equal roads.
|
||
|
||
| Card | Path | Bottleneck link |
|
||
|---|---|---|
|
||
| ROCm0 `03:00.0` | CPU root port `00:01.1` (Gen4 x16) -> switch -> card | Gen4 **x16** (~31.5 GB/s) |
|
||
| ROCm1 `0b:00.0` | CPU root port `00:02.1` (Gen4 x4) -> chipset -> `06:09.0` -> `09:00.0` -> card | Gen4 **x1** (~1.97 GB/s) |
|
||
|
||
With the MTP drafter, the output layers and the heavier compute buffers
|
||
all pinned to the last-layers card, live `rocm-smi` under traffic read
|
||
**98% busy on the x1 card vs 40% on the x16 card**. The boundary crossing
|
||
itself is cheap (~5.2 MB of activation per 512-token ubatch rides the
|
||
wire — a few ms at 1.97 GB/s against a ~430 ms ubatch), but the
|
||
saturated card should not be the one behind the narrowest link. The
|
||
operator moved one layer to ROCm0: **34:31 -> 35:30**. The moved layer
|
||
is 34 — the 4k+2 position of block 8, a DeltaNet — which is what keeps
|
||
the shift cheap (~0.43 GiB).
|
||
|
||
Fit at the new shape: card 0 30.59 GiB (free 1.27), card 1 30.01 GiB
|
||
(free 1.85), gap 0.58 — inside the 1.0 GiB floor and the 1.483 GiB
|
||
slack. The next step, 36:29, crosses the full-attention layer at 35 and
|
||
fails BOTH gates (card 0 free 0.45 GiB, gap 2.22 GiB), so the
|
||
layer-toward-ROCm0 lever is spent: further rebalancing means a ctx cut,
|
||
not a layer. The seam's balance test re-derives all of this from the env
|
||
file on every run.
|
||
|
||
Verified after deploy: `rocm-smi` reads **30.84 GiB used on ROCm0
|
||
(headroom 1.02) and 28.86 GiB on ROCm1 (headroom 3.00)** — the shift
|
||
landed 1.80 / 1.97 GiB, more than the fit's 0.58 GiB gap predicted, and
|
||
the x1 card's headroom tripled. Utilization under live traffic: 60% on
|
||
the x16 card (was 40%), 98% on the x1 card — the drafter, the output
|
||
layers and the heavier compute buffers are structurally pinned to the
|
||
last-layers card, so one DeltaNet layer was all the split lever had left
|
||
to give. A chat completion answers in 1.4 s end-to-end; MTP acceptance
|
||
runs 0.51–0.91 (mean len 2.7–3.7) on the live traffic.
|
||
|
||
|
||
## 2026-09-18: the prompt cache doubles to 16 GiB
|
||
|
||
`--cache-ram` is the one cache tunable on this service that is **not VRAM**: it
|
||
caps the total bytes of saved prompt states in host RAM (`server_prompt_cache`,
|
||
`tools/server/server-task.cpp`). `update()` evicts oldest-first past the cap,
|
||
and `alloc()` **refuses outright** any single state larger than the whole cap —
|
||
with one `WRN` line and nothing else:
|
||
|
||
```
|
||
- prompt state size 11388.990 MiB exceeds cache size limit 8192.000 MiB, skipping
|
||
```
|
||
|
||
That line appeared **137 times** in the two days from 2026-09-16. At 8192 MiB
|
||
every conversation past ~170k tokens was uncacheable — the state was refused,
|
||
the slot's KV went with it, and the next turn re-prefilled the whole thing. The
|
||
cap is a token budget in disguise, because a state is the serialized KV of a
|
||
whole sequence plus its constant-size DeltaNet state plus the MTP draft
|
||
context's own KV:
|
||
|
||
| Component | Per token |
|
||
|---|---|
|
||
| 16 full-attention layers x 4 kv heads x 256 dim x 2 B x 2 (K,V) | 64 KiB |
|
||
| MTP draft context (one nextn layer, same arithmetic) | 4096 B |
|
||
| **state law** | **68 KiB** |
|
||
|
||
| Cap | Tokens storable | Fits one 196608-token slot (~12.8 GiB)? |
|
||
|---|---|---|
|
||
| 8192 MiB (until 2026-09-18) | ~120k | no — refused, then re-prefilled |
|
||
| **16384 MiB (deployed)** | ~240k | **yes**, with ~3 GiB of budget spare |
|
||
|
||
The cost is host RAM, allocated on demand and never pinned, so no VRAM figure
|
||
above moves. dankserver has 62 GiB beside the media-stack co-tenants, which is
|
||
why the cache is bounded rather than set to `-1` (unlimited): the seam test in
|
||
`tests/frank_medium_unit.bats` gates both ends — the cap must clear one full
|
||
slot's state, and it may not claim more than half the host's RAM.
|
||
|
||
## Verified serving behaviour (2026-09-16)
|
||
|
||
- Load: `model loaded` ~25 s after start with a warm page cache, then
|
||
`listening on http://0.0.0.0:7770`. The drafter adds ~2.5 s.
|
||
- Slot shape: `n_slots = 2, n_ctx_slot = 196608`; `/props` reports
|
||
`modalities: {vision: true}` and the sampler defaults above.
|
||
- The drafter is engaged — the init lines are the proof, not the flags.
|
||
Current (MTP, since 2026-09-17 part 2):
|
||
|
||
```
|
||
common_speculative_impl_draft_mtp: adding speculative implementation 'draft-mtp'
|
||
common_speculative_impl_draft_mtp: - n_max=3, n_min=0, p_min=0.00
|
||
... (draft context shares the main model's weights: `shares_model`)
|
||
```
|
||
|
||
Under the DFlash2 era (2026-09-16 through 2026-09-17 part 1) the block was:
|
||
|
||
```
|
||
common_speculative_init_result: loading draft model '/home/dank/models/Qwen3.8-27B-DFlash2-GGUF/Qwen3.8-27B-DFlash2-Q8_0.gguf'
|
||
common_speculative_impl_draft_dflash: adding speculative implementation 'draft-dflash'
|
||
common_speculative_impl_draft_dflash: - n_max=7, n_min=0, p_min=0.00
|
||
common_speculative_impl_draft_dflash: - block_size=8, mask_token_id=248070, n_extract=5, sample_from_anchor=true
|
||
```
|
||
|
||
`block_size=8` and `n_extract=5` are DFlash2's own metadata; the candidate
|
||
selector is what distinguishes it from DFlash 1. DSpark logs under the same
|
||
`draft_dflash` symbol with `block_size=7` and a Markov head, so read the
|
||
numbers, not the implementation name.
|
||
- Both cards hold the model. The measurement below is the **34:31-era**
|
||
record of the part 2 cutover; the deployed 35:30 shape is measured in
|
||
the part 3 section (ROCm0 30.84 GiB / ROCm1 28.86 GiB, headroom
|
||
1.02 / 3.00 GiB, both above the 1.0 GiB floor):
|
||
at 393216 with the deployed batch and the 34:31 split, `rocm-smi`
|
||
reported **28.91 GiB used on ROCm0 and 30.78 GiB on ROCm1** of 31.86
|
||
GiB each — the fit predicted 29.77 / 28.32, so ROCm1 ran 2.46 GiB
|
||
hotter than predicted and the seam's MTP constant was recalibrated to
|
||
the measured value (see the 2026-09-17 part 2 section). ROCm1's
|
||
1.08 GiB of headroom cleared the 1.0 GiB floor that shape required.
|
||
- Drafting is live on **both slots** (the 2026-09-16 DFlash2-era numbers below
|
||
are the record; MTP-era acceptance is in the part 2 section), which is what
|
||
a 2-slot deployment has to
|
||
prove: `slot print_timing` reported `draft acceptance = 0.62857 (44 accepted /
|
||
70 generated), mean len = 5.40` and `0.48214 (27 / 56)` on slot 1, and
|
||
`0.38776 (57 / 147), mean len = 3.71` on slot 0, from one concurrent pair.
|
||
Acceptance moves with the text; the point is that neither slot is idle.
|
||
- **140,069 prompt tokens prefilled in 134 s** (1043 tok/s) was measured
|
||
through a *single* slot, at the previous configuration: the raised window is
|
||
real rather than declared, and it is 22 % beyond the 114688 the vLLM
|
||
predecessor could offer. The window is now split over two slots.
|
||
|
||
- **Nothing is replicated.** The loader does not print a per-device ledger at
|
||
this verbosity, so this is measured instead of read: raising CTX_SIZE from
|
||
131072 to 262144 added **4.83 GB to GPU 0 and 4.84 GB to GPU 1** — both cards
|
||
carry the KV for their own attention layers, about half each — and neither
|
||
card's *total* usage (24.94 / 29.42 GB at 262144) could hold the 31.46 GB
|
||
file, so no card carries a full copy of the weights. Compare the shape this
|
||
replaced: two vLLM replicas, i.e. two complete copies of the model, one per
|
||
card, which is exactly why a 29.8 GB FP8 checkpoint could not be served there
|
||
at any usable context.
|
||
|
||
### Measured throughput (2026-09-16, deployed service, `temperature 0`)
|
||
|
||
Same filler corpus for every shape, so these are shapes rather than a
|
||
benchmark: a large prompt with a tiny budget isolates prefill, a small prompt
|
||
with a long budget isolates decode. `draft`/`acc` are per request; acceptance
|
||
moves with the text (filler prose accepts far worse than code or a real agent
|
||
loop).
|
||
|
||
| Shape | Prefill | Decode | Acceptance |
|
||
|---|---|---|---|
|
||
| 4k prompt / 32 out | 3994 tok in 3.25 s — 1230 tok/s | 59.8 tok/s | 71 % |
|
||
| 1k prompt / 512 out | 998 tok in 1.41 s — 707 tok/s | **43.8 tok/s** | 41 % |
|
||
| 16k prompt / 16 out | 12,528 tok in 8.94 s — **1402 tok/s** | — | 100 % (13/13) |
|
||
| 64k prompt / 16 out | 48,508 tok in 39.8 s — 1219 tok/s | — | 53 % |
|
||
| 140k prompt / 16 out | 140,069 tok in 134.3 s — 1043 tok/s | — | — |
|
||
| same 1k prompt again | **4 tokens prefilled** | 43.8 tok/s | 41 % |
|
||
|
||
The last row is the prefix cache doing its job: the second request of the same
|
||
conversation re-prefilled 4 of 998 tokens. Prefill holds ~1.0-1.4k tok/s from
|
||
4k to 140k, and decode sits at ~44 tok/s on a long generation — against the
|
||
vLLM predecessor's 42.8 tok/s single-stream (two replicas) and the ROCmFPX
|
||
llama.cpp measured here at 35.7-46.2 tok/s. One process, one model, both cards,
|
||
no proxy.
|
||
|
||
### What the drafter is worth — and why MTP is not the alternative
|
||
|
||
Measured 2026-09-16 with `scripts/bench-spec-ab.py` on the **DFlash2** drafter —
|
||
the one deployed now: four configurations of the same binary, identical prompts
|
||
at `temperature 0`, unit stopped between runs (one host, one copy of a 29 GiB
|
||
model). 1k-token prompt, 512-token budget. These are the numbers that selected a
|
||
trained multi-token drafter over the model's own MTP head.
|
||
|
||
| Configuration | Decode | Drafting | VRAM (card 1) |
|
||
|---|---:|---|---:|
|
||
| `none` — autoregressive | 17.7 tok/s | — | 26.07 GB |
|
||
| `draft-mtp` n_max 3 (the model's own head) | 32.4 tok/s | 58 % of 72 | 29.35 GB |
|
||
| `draft-mtp` n_max 7 | 26.7 tok/s | 33 % of 147 | 29.65 GB |
|
||
| **`draft-dflash` n_max 7 — deployed** | **43.5 tok/s** | 41 % of 119 | 29.42 GB |
|
||
|
||
- **DFlash2 is worth 2.46× over autoregressive and 34 % over the best MTP.**
|
||
Deeper MTP drafting costs more than it returns: at depth 7 only a third of
|
||
the proposals survive, at depth 3 more than half do.
|
||
- **All four produced identical greedy text** (`2, 3, 5, 7, 11`): speculation
|
||
buys speed, never a different answer. The harness warns loudly if a future
|
||
configuration ever breaks that.
|
||
- **MTP is not the cheaper option.** Selecting it loads `blk.64.nextn.*` from
|
||
the target and adds ~3.3 GB to the card that holds it, against 2.06 GB for
|
||
the drafter's file — which is also why the load log says
|
||
`unused tensor blk.64.nextn... ignoring` under `draft-dflash`: correct
|
||
behaviour, not a truncated download.
|
||
- **They cannot be chained.** `--spec-type draft-mtp,draft-dflash` dies at load
|
||
with `context type MTP requested but model doesn't contain MTP layers`. A
|
||
chain has ONE draft context, and MTP builds against whatever model that
|
||
context holds — here the DFlash2 file. MTP's head is loaded only when MTP *is*
|
||
the draft model (`creating MTP draft context against the target model ...`).
|
||
That is structural, not a tuning question.
|
||
|
||
### Chaining the draftless n-gram matcher (deployed)
|
||
|
||
`ngram-mod` **can** ride beside a context-backed drafter, and it is enabled.
|
||
The reason is in `common/speculative.cpp`: the n-gram types are added to the config vector
|
||
unconditionally while every model-backed drafter is gated on
|
||
`ctx_dft != nullptr` — so n-gram needs no draft context and cannot collide with
|
||
the drafter over one. (MTP is the opposite: it *is* context-backed, which is why
|
||
`draft-mtp,draft-dflash` cannot load.) The vector's fixed order — ngram types
|
||
first, then the drafters — also means CLI order is irrelevant.
|
||
|
||
The measurements below are DFlash2's — the deployed configuration. (During the
|
||
DSpark excursion the chain was `draft-dspark,ngram-mod`; the matcher's behaviour
|
||
is a property of the matcher, not of the drafter beside it.)
|
||
|
||
Measured with `scripts/bench-spec-ab.py`, decode tok/s, three shapes that a
|
||
proposer can and cannot match:
|
||
|
||
| Configuration | prose 1k/512 | echo 1.5k/400 (copy 0.94) | structured/400 | identity |
|
||
|---|---:|---:|---:|---:|
|
||
| `draft-dflash@7` | 43.77 | 76.82 | 91.08 | 70.37 |
|
||
| `ngram-mod` alone | 17.72 *(0 drafts)* | 102.50 | 17.80 *(0 drafts)* | 17.77 |
|
||
| **`draft-dflash,ngram-mod@7`** | **43.65** | **200.58** | 91.22 | 70.40 |
|
||
|
||
- **On prompt-echo traffic the chain is 2.6× the drafter alone** (200.6 vs 76.8)
|
||
and twice `ngram-mod` by itself, at 100 % acceptance.
|
||
- **On prose it is indistinguishable from DFlash2 alone** — the same 119 drafts
|
||
and 49 accepts, the same 43.7 tok/s. That is the property that makes it safe
|
||
as a default: where the matcher cannot match, it costs nothing.
|
||
- `ngram-mod` alone is a *negative control*: zero drafts on prose and identity
|
||
(17.7 tok/s = the autoregressive baseline), 100 % acceptance on echo. Nothing
|
||
else could have proposed those tokens, which is how its drafts are identified
|
||
at all — see the trap below.
|
||
- VRAM is unchanged (29.42 GB on card 1 with or without it): the matcher is a
|
||
64 MiB host-RAM hash table.
|
||
- `structured` (25 records with unique ids) rewards no proposer: its lines
|
||
differ, so there is no long span to match. It is a negative control for
|
||
*self-repetition*, not evidence about it.
|
||
|
||
**The window matters, and shorter is not better.** The chain at three windows:
|
||
|
||
| `SPEC_NGRAM_N_MIN`/`N_MAX` | prose 1k/512 | echo 1.5k/400 |
|
||
|---|---:|---:|
|
||
| **48/64 — deployed** | **43.9** | **200.6** |
|
||
| 16/32 | 36.0 | 161.0 |
|
||
| 8/16 | 36.2 | 113.8 |
|
||
|
||
A short window makes the matcher propose on ordinary prose fragments, where the
|
||
proposals are rejected and the verification is still paid — the "ngram hurts"
|
||
result this fleet measured twice before on other payloads. At 48/64 the chain
|
||
keeps DFlash2's exact numbers on prose and wins on echo.
|
||
|
||
> **Operational trap: `ngram-mod` prints no startup banner in this build.**
|
||
> `journalctl -u frank-medium` shows `adding speculative implementation
|
||
> 'draft-dflash'` and nothing for the matcher — the only evidence it is running
|
||
> is behavioural (`draft_n` > 0 where text repeats) or the per-task line
|
||
> `slot print_timing: ... draft acceptance = 1.00000 (256 accepted / 256
|
||
> generated), mean len = 65.00`. Do not conclude from the load log that the
|
||
> chain is not engaged.
|
||
|
||
### The agent-loop shape (fleet harness, 2 families × 2 turns)
|
||
|
||
`scripts/bench-turn-latency.py http://127.0.0.1:7770 2 2 LAYERSPLIT` — two
|
||
agent loops re-sending a ~30k-token document concurrently:
|
||
|
||
| | turn 1 | turn 2 |
|
||
|---|---:|---:|
|
||
| family 1 | 24,799 ms | 30,697 ms |
|
||
| family 0 | 51,424 ms | 8,015 ms |
|
||
|
||
Read that against the prefill table, not as a latency regression: with **one**
|
||
slot the second family's first turn waited for the first (25 s of cold 30k
|
||
prefill + the wait ≈ 51 s), and the last number — 8.0 s — was a re-prefill the
|
||
cache could only partly serve. That was the 1-slot configuration. The
|
||
deployment now runs **two slots of 131072**, so a second agent loop has a slot
|
||
of its own instead of queueing, at the cost of halving the window each one
|
||
gets. `--parallel` splits the context and the KV between slots; it does not add
|
||
throughput, so the trade is concurrency against conversation length.
|
||
|
||
### Benchmarking this endpoint
|
||
|
||
`scripts/bench-batch-ubatch.sh` sweeps `--batch-size` against `--ubatch-size`
|
||
with `llama-bench` and records, per cell, the prefill rate, the decode rate and
|
||
each card's **peak** VRAM. Run it from a transient unit rather than a shell:
|
||
|
||
```bash
|
||
sudo systemd-run --unit=fm-bench --collect --no-block --property=Type=oneshot \
|
||
--property=LimitMEMLOCK=infinity --property=User=dank \
|
||
scripts/bench-batch-ubatch.sh
|
||
```
|
||
|
||
Properties of `llama-bench` shape that harness, and three of them have already
|
||
produced wrong numbers on this fleet:
|
||
|
||
- **It separates devices differently from `llama-server`.** `/` splits devices
|
||
*within one configuration*, `,` splits *sweep values* — so this service's
|
||
`--tensor-split 36,29` is `-ts 36/29` to llama-bench, and the comma form
|
||
silently benches each card on its own instead. Likewise `--device
|
||
ROCm0,ROCm1` → `-dev ROCm0/ROCm1`.
|
||
- **There is no `-c`.** The context is a consequence of the test shape:
|
||
`n_ctx = n_prompt + n_gen + n_depth`. No bench can therefore prove that a
|
||
configuration's activation buffers fit beside the *deployed* KV — that needs
|
||
the live unit at 327680 across two slots.
|
||
- **A test that cannot allocate calls `exit(1)`**, taking the whole process with
|
||
it, so the harness runs one invocation per cell and a failure costs one TSV
|
||
line instead of the rest of the sweep.
|
||
- **A shell's memlock limit is not the unit's.** The unit sets
|
||
`LimitMEMLOCK=infinity` for the ROCm staging buffers; an SSH shell here gets
|
||
8 MiB, which fails the load for a reason that has nothing to do with the
|
||
model being benched.
|
||
- **No `--spec-*` and no `--parallel`.** The drafter cannot take part, so a
|
||
`llama-bench` decode figure is the *drafterless* baseline — 15.9 tok/s at 64k
|
||
depth against the deployed 43.8 — and slots are never exercised. `-d <n>` is
|
||
what puts KV in front of the measurement at all: a `tg` instance runs with
|
||
`n_prompt = 0`, i.e. at ~zero depth unless `-d` says otherwise.
|
||
|
||
## Observability
|
||
|
||
The endpoint exports Prometheus metrics under the **`llamacpp:`** namespace
|
||
(`--metrics`), which is what the fleet's scrape and dashboard read — the
|
||
`vllm:` series the predecessor exported no longer exist:
|
||
|
||
| Need | Series |
|
||
|---|---|
|
||
| Throughput | `llamacpp:prompt_tokens_seconds`, `llamacpp:predicted_tokens_seconds` |
|
||
| Queue | `llamacpp:requests_processing`, `llamacpp:requests_deferred` |
|
||
| Cache reuse | `llamacpp:prompt_tokens_cached_total` vs `llamacpp:prompt_tokens_total` |
|
||
| **Speculative decoding is live** | `llamacpp:spec_decode_num_drafts_total` and `llamacpp:spec_decode_num_accepted_tokens_total` (their ratio is the acceptance rate) |
|
||
|
||
`spec_decode_num_drafts_total` staying at 0 is the one-line answer to "is the
|
||
drafter actually running": the flags can all be right while the draft model
|
||
failed to load.
|
||
|
||
## Client contract
|
||
|
||
OpenAI-compatible `/v1` on **7770**, alias **`frank-medium`**, streaming and
|
||
`/health`, `/props`, `/metrics` as llama.cpp provides them. Reasoning traces
|
||
are projected into `reasoning_content`. A request that sends no sampler fields
|
||
gets the server defaults above; a client that sends its own wins. Text only.
|
||
|
||
## Troubleshooting
|
||
|
||
- **Loads, then dies on the first request with a rocBLAS Tensile error** — the
|
||
bundled device-library depth (see the bundling trap above); check
|
||
`build-rdna4-rocm714/lib/rocblas/library/` exists and is populated.
|
||
- **`hipErrorNoDevice` / a card missing** — `HIP_VISIBLE_DEVICES` must be
|
||
`0,1`; the iGPU enumerates as device 2 and must stay hidden.
|
||
- **Drafting not engaged** — read the journal for
|
||
`common_speculative_impl_draft_dflash` (DFlash2 and DSpark both log under that
|
||
symbol); a missing drafter file makes the serve script exit before
|
||
llama-server starts, and a drafter whose GGUF the build does not recognise
|
||
fails the load instead.
|
||
- **Why is `blk.64.*` "unused"?** That is the model's own MTP head and it is
|
||
correct for `draft-dflash`. It is not a truncated download:
|
||
`bash scripts/download-frank-medium-gguf.sh --verify-only` proves the bytes.
|
||
- **A rebuild segfaults at load** — the build is gfx1201-only on purpose; a
|
||
gfx1200 sibling binary links, loads the model and then segfaults on these
|
||
cards.
|
||
|
||
## See Also
|
||
|
||
- [[frank-medium-radiance-mxfp4]] — the retired vllm-radiance MXFP4 deployment
|
||
(replicas + proxy) and the PCIe-atomics evidence
|
||
- [[frank-medium-rocmfpx]] — the llama.cpp/ROCmFPX deployment earlier on
|
||
2026-09-15, whose `--split-mode layer` measurements this service reuses
|
||
- [[Dual-R9700-Cooperation]] — why the two cards cannot take part in a
|
||
collective, and every software escape that is closed
|
||
- [[frank-small-qwen38-gsq-rco]] — the fleet's other llama.cpp endpoint (RTX
|
||
3090, MTP, the image path)
|
||
- [[Dankserver System Specs]] — hardware inventory
|
||
- [[frank-medium-batch-ubatch]] — the batch/ubatch sweep that set the deployed
|
||
geometry and moved this split from 36:29 to 35:30: the ubatch optimum at 512,
|
||
why the logical batch is free, and the harness traps
|