serverAppConfigs/serverAppConfigs-docs/LLM/frank-medium-radiance-mxfp4.md

type, title, description, tags, timestamp

type	title	description	tags	timestamp
reference	frank-medium vllm-radiance MXFP6 Runbook	Retired era: the tensor-parallel radiance shape on two cards, its quant-class swap, and the on-disk rollback payload.	frank-medium dankserver vllm vllm-radiance mxfp6 mxfp4 gfx1201 r9700 rccl pcie-atomics quadlet retired-endpoint	2026-09-21

frank-medium vllm-radiance MXFP6 Runbook

RETIRED era (2026-09-21 → 10-10; the endpoint is frank-archon today and this shape is one row in its Eras table.) The deployment was the ggz14/vllm-radiance-mxfp4:0.13.0-e060895 image (built from PR #51's head, tagged locally as localhost/radiance-vllm:0.13.0-e060895) plus the radiance-vllm-mxfp4 patches at e060895, serving hugypufy/Swift-Qwen3.8-27B-PARO-MXFP6 (ParoQuant MXFP6 W4A8, 25.16 GiB, 18 shards) as ONE tensor-parallel process across both Radeon AI PRO R9700 (gfx1201). It replaced the mainline-llama.cpp layer-split deployment (frank-medium-llamacpp-dflash2, SUPERSEDED 2026-09-20); the same image and model had earlier run here as two single-GPU replicas, which is the shape most of this page describes.

Serving shape since 2026-09-21: ParoQuant MXFP6 with unquantized (bf16) KV. This endpoint cut to the MXFP6 checkpoint above (PR #51's launch6_p kernel, quant_method=paroquant_mxfp6, bits=6, group_size=128, krot=8, rotation on, e4m3fn scale) with --kv-cache-dtype auto (unquantized bf16 KV) and --gpu-memory-utilization 0.98. The Swift variant is a finetune of ukisai/Swift-Qwen3.8-27b (a different base from Qwen/Qwen3.8-27B), so the DFlash2-FP8 drafter's hidden-state distribution is a mismatch — acceptance may drop from the 4.24/7 that the DFlash2 arm measured on the FP8 target (the A/B below ran both arms on FP8, not on MXFP4). The MXFP4 checkpoint is the rollback, and as of the 2026-10-19 host census it is no longer on disk: /home/dank/models/ holds StillDeadcode, rad and tpl.jinja only, so a rollback re-fetches /models/Qwen3.8-27B-Uncensored-MXFP4-awq (Hub just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq) and flips SNAP and the image back to localhost/radiance-vllm:0.9.3 (re-pull docker.io/stilldeadcode/vllm-radiance:0.9.3 if the local tag is gone; the MXFP4 shape auto-sized to 486,464 tokens with bf16 KV and held the full 262,144 window).

Deployed shape since 2026-09-20: one rootless Quadlet, --tensor-parallel-size 2. The ProArt B850-Creator WiFi NEO board swap put both R9700s on CPU GPP root ports (00:01.1 → 03:00.0 and 00:01.3 → 07:00.0), both Gen5 x8, every upstream bridge AtomicOpsCap: Routing+ and both endpoints ReqEn+ — the wall documented in Why TP=2 cannot run here is gone on this board. The collective that actually runs is vLLM's R4D one-shot allreduce, which does not use RCCL at all (TheRock RCCL has no gfx1201 device code; see Dual-R9700-Cooperation).

Deployed flags (2026-09-21)

From dankserver/frank-medium/quadlet/frank-medium.container:

Setting	Value
Model	hugypufy/Swift-Qwen3.8-27B-PARO-MXFP6 (ParoQuant MXFP6 W4A8, 25.16 GiB, 18 shards; a finetune of ukisai/Swift-Qwen3.8-27b)
Speculative	--speculative-config {"method":"dflash","model":"/models/Qwen3.8-27B-DFlash2-FP8","num_speculative_tokens":7} (the z-lab DFlash2 external drafter; measured on the Swift finetune: mean acceptance length 7.00 / 85.7% on easy text, 5.08 / 58.3% on varied — the drafter survives the base-model change. A/B tested against MTP self-speculation — DFlash2 wins on decode, see below)
Parallelism	--tensor-parallel-size 2 (ROCR_VISIBLE_DEVICES=0,1)
Context	--max-model-len 262144
Concurrency	--max-num-seqs 8, chunk 8192
KV	auto = unquantized bf16 (--kv-cache-dtype auto), auto-sized — pool 425,047 tokens (measured at 0.98 fill: 14.42 GiB per card available for KV, 1.62x the 262,144 window); rollback: KV_DTYPE=fp8
Prompt cache	--enable-prefix-caching
GPU fill	--gpu-memory-utilization 0.98 (the 32 GB ceiling — 0.99 fails the startup check)
Parsers	--enable-auto-tool-choice --tool-call-parser qwen3_coder, --reasoning-parser qwen3
Sampler floor	--override-sampling-effects (temp 0.7 / top_p 0.95 / top_k 20)
Multimodal	--limit-mm-per-prompt '{"image": 8}'
Chat template	qwen-sharp-v22.5.0.jinja (peculiar-ragdoll Qwen-Sharp v22.5.0, repo-tracked at dankserver/frank-medium/, mounted at /templates/; terse system prompt by default — opt out per-request with chat_template_kwargs {"terse": false}; rollback: flip CHAT_TEMPLATE back to /templates/qwen-sharp-v22.4.0.jinja)

The KV pool is auto-sized, not pinned. Production ships no KV_MEM pin (KV_MEM=0): the pool comes up at whatever VRAM the weights + drafter leave. MXFP6's weights are 25.16 GiB (~12.58 GiB per card at TP=2); at 0.98 fill (31.22 GiB per card) that leaves 14.42 GiB per card for KV — a measured 425,047-token pool, 1.62x the 262,144 window (boot 2026-09-22). The DFlash2-FP8 drafter costs 2.1 GiB of VRAM on top. 0.98 is the ceiling, not a preference: the R9700 reports 31.86 GiB total but only 31.23 GiB free at device-init (the ROCm/HIP context + NCCL + the radiance custom all-reduce buffers consume the 0.63 GiB gap), so a gpu_memory_utilization above 0.98 (0.99 = 31.54 GiB) fails vLLM's startup check before the weights even load. For reference the MXFP4 checkpoint (the rollback) auto-sized to 486,464 tokens at the same 0.98 with bf16 KV (1.86x the window); MXFP6's larger weights give a smaller pool, still comfortably above the window. After the first healthy start, recalibrate with calibrate-kv.sh.

DFlash2 vs MTP — A/B on the FP8 checkpoint (2026-09-20)

The FP8 checkpoint carries its own single-layer MTP head, so it can self- speculate (SPEC_METHOD=mtp) or use the z-lab DFlash2 external drafter (SPEC_METHOD=dflash, a separate quantization-agnostic 2.1 GiB DFlash2DraftModel). Both were deployed and measured on the same FP8 target (scripts/bench-conversation.py, 10 real prompts, max_tokens=512, greedy, salted, --burst 4):

Arm	Mean accept	Decode (geomean)	TTFT (mean)	Prefill (geomean)	Burst-4 aggregate	KV pool
DFlash2 (depth 7)	4.24/7	100.0 t/s	0.19 s	923.0 t/s	287.1 t/s	384,238 (1.47×)
MTP self-spec (depth 8)	5.66/8	85.9 t/s	0.06 s	1096.9 t/s	179.7 t/s	461,540 (1.76×)

DFlash2 wins on decode (+16 %) and burst-4 (+60 %); MTP wins on TTFT and prefill. The counterintuitive part: DFlash2 accepts fewer tokens per step (4.24 vs 5.66) yet decodes faster. Accept length is a proxy, not the metric — throughput is (accept + 1) / (draft time + verify time). The DFlash2 drafter is a small separate model that steps fast; the MTP head is the target's own layer applied autoregressively 8×, so its draft is slower. The faster step more than makes up the lower accept length. FP8 + DFlash2 (100.0 t/s) also beats the original MXFP4 + DFlash2 (92.4 t/s) — the higher-precision FP8 target verifies more of the drafter's proposals. The drafter costs 2.1 GiB of VRAM (the KV pool is 1.47× vs MTP's 1.76×), an acceptable trade for the decode win. DFlash2 is the deployed shape.

Prefix caching — measured granularity (2026-09-20)

--enable-prefix-caching is on, and radiance says so at boot: "Prefix caching in Mamba cache 'align' mode is currently enabled. Its support for Mamba layers is experimental." The attention block size is 1,648 tokens (boot log: "Setting attention block size to 1648 tokens …"), padded to the mamba/GDN page size, and reuse happens at that granularity — not per token.

Measured on the deployed endpoint, not inferred:

Probe	Cold	Warm	vllm:prefix_cache_hits_total delta
identical 13,021-token prompt	2.30 s	0.68 s (3.4×)	+9,888 of 13,021 queried (76 %)
identical 2,521-token prompt	0.51 s	0.51 s	+0 (under two blocks)

Consequence: a shared prefix shorter than about two blocks (≈3.3k tokens) re-prefills in full, so the benefit starts at ~3.3k tokens of shared prefix. The API returns prompt_tokens_details: null in this build, so the vllm:prefix_cache_hits_total / vllm:prefix_cache_queries_total counters (both token-counted) are the only instrument — see Observability.

Battery results — why this shape won (2026-09-20)

scripts/bench-conversation.py, 10 real-workload prompts, max_tokens=512, greedy, salted per run and per arm, service stopped between arms (scripts/bench-conversation.py --host http://127.0.0.1:7770 --burst 4):

Arm	Shape	Decode (geomean)	TTFT (mean)	Prefill (geomean)	Burst-4 aggregate
I	llama.cpp layer split 32:33 + DFlash2 Q4_K_M drafter	45.7 t/s	0.72 s	~530 t/s	71.7 t/s (26.7/stream)
L	llama.cpp --split-mode tensor 1:1, drafterless	30.5 t/s	0.71 s	—	54.3 t/s
R	vllm-radiance TP=2 + DFlash2-FP8	95.8 t/s	0.21 s	1493.6 t/s	259.4 t/s (65.6/stream)

Re-measured against the production unit after deploy: geomean decode 92.4 t/s, mean TTFT 0.22 s, burst-4 aggregate 248.9 t/s — within noise of arm R.

Honest caveat: arm R carries MXFP4 (4-bit) where the llama arms carry Q8_0 (8-bit), so part of the win is quantization, not tensor parallelism alone. Arm L exists to show that llama.cpp's own TP is not the answer here regardless — it is slower than its own layer split, and cannot carry the drafter at all.

Verified capabilities (2026-09-20, against the live endpoint)

• /v1/models → id: frank-medium; /health → 200.
• Tool calls: finish_reason: tool_calls via qwen3_coder. The harness drives a bash tool through this endpoint end to end.
• Reasoning: reasoning_effort accepts none/low/medium/high natively (vLLM's ReasoningEffort); thinking arrives in the reasoning field, NOT llama.cpp's reasoning_content. max_tokens counts visible content only; with no reasoning budget the model still spends it, so a tiny budget returns content: null with finish_reason: length.
• Vision works despite language_model_only: true in config.json — the Gemma4 merge path answers image requests (~83 prompt tokens for a small synthetic), correctly identifying shape and colour on held-out images. models.yml therefore lists this endpoint as input: [text, image].
• Drafter is live: accept length 2.9–4.3 tokens per draft, acceptance ~0.60.

Rollback

The llama.cpp layer-split trio is archived in scripts/.AGENTTRASH/ 20260920-llama-layersplit/ (frank-medium.service, frank-medium.env, serve-frank-medium.sh) and the runtime + Q8_0 payload remain on disk:

1. git revert the cutover commit.
2. deploy.sh dankserver --service frank-medium (re-installs and re-enables the root unit).
3. Or by hand: restore the three files, move the unit symlink back into /etc/systemd/system/, sudo systemctl daemon-reload && sudo systemctl unmask frank-medium && sudo systemctl start frank-medium.
4. Gate: curl -s http://192.168.0.235:7770/health → 200.

R4D vs RCCL

vLLM's R4D one-shot allreduce bypasses RCCL entirely because TheRock RCCL builds lack gfx1201 device code (invalid device function). The Ubuntu RCCL 2.27 / librccl1 7.1.0 build does load gfx1201 kernels — that is what scripts/build-llama-rdna4.sh with GGML_HIP_RCCL=ON links (build-rdna4-rocm714-rccl/) — so llama.cpp's tensor mode has a real collective too. It is simply the slower shape.

The sections from here down describe the 2026-09-16 → 2026-09-19 era: two single-GPU replicas behind an adaptive-routing proxy, on the TUF GAMING B650-E. The deployed shape is above; this is the evidence trail that made replicas necessary then and what the 2026-09-20 board swap changed.

Architecture

graph LR
    C[clients: Hermes, harness] --> P["frank-medium-proxy<br/>:7770 adaptive routing"]
    P --> A["frank-medium-a<br/>R9700 #1, TP=1, :7771"]
    P --> B["frank-medium-b<br/>R9700 #2, TP=1, :7772"]

Three rootless Quadlet units, all WantedBy=default.target + Restart=on-failure:

Unit	Container	GPU	Port	Role
frank-medium-a	frank-medium-a	HIP_VISIBLE_DEVICES=0	7771	replica
frank-medium-b	frank-medium-b	HIP_VISIBLE_DEVICES=1	7772	replica
frank-medium-proxy	frank-medium-proxy	—	7770	routing

The fleet contract is unchanged: one endpoint on :7770, served model name frank-medium, OpenAI /v1, streaming included. The vLLM OpenAI API is stateless per request, so no session affinity is needed for correctness. Routing is described in Routing.

Only HIP_VISIBLE_DEVICES is set per replica, never a matching ROCR_VISIBLE_DEVICES: ROCR filters first and HIP then indexes into the filtered list, so the single-valued pair (ROCR=1,HIP=1) leaves an empty device set and the engine dies with hipErrorNoDevice.

frank-medium.container (TP=2, KV_MEM=18563072000, MAXLEN=262144) stays in the repo as the target shape for a host whose second card sits on an atomic-capable root port. deploy-remote.sh deliberately does not install it and moves any stale copy into ~/.AGENTTRASH/, because a TP=2 unit on this host would restart forever without ever becoming healthy.

Why TP=2 cannot run here

Resolved 2026-09-20. Everything below is true of the TUF GAMING B650-E WIFI, whose second card sat behind chipset bridges advertising AtomicOpsCap: Routing-. The ProArt B850-Creator WiFi NEO swap moved both cards onto CPU GPP root ports (00:01.1 → 03:00.0, 00:01.3 → 07:00.0, Gen5 x8, every hop Routing+, both endpoints ReqEn+), and TP=2 now boots and serves. The measurement verdict against the alternatives is in Dual-R9700-Cooperation.

Symptom

Any two-rank collective dies at launch:

[FATAL ERROR]: HIP failure: 'the operation cannot be performed in the present state'
  group.cc:337 -> group.cc:656 -> group.cc:801 -> enqueue.cc:3111

With the runtime's own API log (AMD_LOG_LEVEL=3) the failing call is explicit:

[ShaderName : ncclDevKernel_Generic_4(ncclDevKernelArgsStorage<4096ul>)]
[ :1:rocvirtual.cpp :4208:  Pcie atomics not enabled, hostcall not supported]
[ :1:rocvirtual.cpp :4636:  AQL dispatch failed!]
[ :3:hip_module.cpp :893 :  hipLaunchKernel: Returned hipErrorIllegalState]

ROCclr refuses to hand the kernel a hostcall buffer when the device has no PCIe atomics (pcie_atomics_ is true for the device init, and ROCr passes that buffer only when pcie_atomics_ is true; otherwise it returns "not supported" and the dispatch fails) — and the launch aborts. hipErrorIllegalState is that refusal, not a driver fault.

Chain of causation

1. RCCL's collective kernels carry a hidden hostcall buffer parameter.
2. ROCr passes that buffer only when pcie_atomics_ is true for the device; otherwise it returns "not supported" and the dispatch fails.
3. pcie_atomics_ is set once, at device init, from the HSA link info between the CPU agent and the GPU (Device::checkAtomicSupport). It is not re-read per process.
4. The KFD driver populates that link info from pci_enable_atomic_ops_to_root(), which — per the PCIe spec and the kernel's implementation — requires every upstream bridge between the device and its root port to advertise AtomicOp Routing capability (drivers/pci/pci.c).

The hardware difference between the two cards:

Card	PCIe chain (endpoint → root port)	AtomicOpsCap on every hop	AtomicOpsCtl
R9700 #1 (03:00.0)	02:00.0 → 01:00.0 → 00:01.1	Routing+ / Routing+ / Routing+	ReqEn+ (atomics ON)
R9700 #2 (0b:00.0)	0a:00.0 → 09:00.0 → 06:09.0 → 05:00.0 → 00:02.1	Routing+ / Routing+ / Routing− / Routing− / Routing+	ReqEn− (atomics OFF)

Card #1 is on a CPU root port; card #2 is behind the chipset bridges 05:00.0/06:09.0, which advertise AtomicOpsCap: Routing-. That capability bit is read-only — no software can enable atomic routing through those bridges — so the driver's atomic enablement call fails for card #2 and the KFD marks the CPU↔GPU2 link NO_ATOMICS.

Writing the AtomicOp Requester bit into the endpoint's Device Control 2 by hand (setpci -s 0b:00.0 CAP_EXP+28.w) does not help: the driver reads the chain at probe time, and ROCr consumes the driver's cached answer. Only a re-probe (a boot with the write already in place) would change it — and even then the two Routing- bridges would still make pci_enable_atomic_ops_to_root() fail.

What was ruled out

Suspect	Result
Kernel version	7.0.0-28, -30 and -31 fail identically (full apt full-upgrade to the newest available kernel changed nothing)
IOMMU	identical before and after iommu=pt
P2P hardware	fine: 45.3 GB/s bidirectional copy between the two cards; both cards pass rocm-bandwidth-test
The shipped librccl	a from-source rebuild of ROCm/rccl develop (with local fixes for that branch's dangling references) reproduces the same failure at the same call path — the library is not at fault
vLLM/pynccl	a bare two-rank torch.distributed allreduce with no vLLM involved fails the same way; single-rank collectives succeed on both cards (NCCL short-circuits at world_size=1, which is why a one-card test proves nothing)
Transport knobs	NCCL_PROTO=Simple, P2P_DISABLE, SHM_DISABLE, CUMEM_ENABLE=0, MSCCL off, channel counts, launch modes, HSA_ENABLE_SDMA=0 — all identical
Image/support matrix	the same image on a different host with the same kernel (Proxmox VM, Codeberg issue #22) fails byte-identically

What would fix it

• Move card #2 to a slot whose path is atomic-capable (a CPU root port, as card #1 has). This is the only software-free fix and the one that restores true TP=2. On this AM5 platform the CPU provides the x16 slot plus x4 lanes, so a second x16-class card necessarily lands on the chipset path — a platform with more CPU lanes (Threadripper/EPYC) is the structural answer.
• Done 2026-09-20 — this is exactly what happened. The ProArt B850-Creator WiFi NEO routes both slots to CPU root ports, the atomics condition is satisfied, and TP=2 is the deployed shape. The Threadripper/EPYC argument was correct about the old board and is now moot.

Routing

The proxy decides per request, from proxy/routing.py (policy) and proxy/affinity-proxy.py (HTTP plumbing):

• A continuation — a request carrying more than one message, i.e. an agent loop re-sending its context — is pinned to the replica that served the prefix before, because --enable-prefix-caching has that prefix warm there.
• Anything else — a one-shot question, a batch member, a burst of parallel subagents — spreads across the cards by counter, so both prefill in parallel.

Measured on this host, 90k-token contexts and 15k-token prompts, vllm bench plus a sequential-turn client:

strategy	cold 16-request burst	warm-turn TTFT (4 turns × 2 families)
roundrobin (counter)	18.6 s	2676 ms — every other turn pays a cold prefill
affinity (hash of prefix)	30.1 s — one card does all the prefilling	1096 ms
adaptive (deployed)	19.97 s	1095 ms

Two failures found by measuring rather than reasoning, both now covered by tests/frank_medium_proxy.bats:

• Hashing a fresh prefix to pick a card sends every member of a same-prefix burst to one card (29.8 s for the burst above). Only continuations may be pinned; the spread has to come from a counter.
• The identity of a conversation is its first 4 KiB (AFFINITY_BYTES). A prefix shorter than that window is indistinguishable from its own tail, so turns of the same conversation hash differently and never pin.

ROUTING=roundrobin and ROUTING=affinity remain available to re-measure the choice; the default is adaptive.

Deployed configuration

Both replicas, uniform except GPU index and port:

Knob	Value	Why
TP	1	one card per replica
MAXLEN	114688	the largest shape that boots on one card: 262144 needs 9.61 GiB of KV against 5.21 GiB available, and 131072 needs 5.58 GiB against 5.22
KV_MEM	0	passed through literally as --kv-cache-memory 0, which the image's patched KV sizing reads as "profile this shape"; the upstream pin (2x7551-32624) is a TP=2 measurement and is invalid here. Do not blank the variable instead of passing 0: with the flag omitted vLLM sizes the cache from GPU_UTIL alone and the KV allocation dies with torch.OutOfMemoryError (measured: 2.16 GiB requested, 116 MiB free, 29.34 GiB already held)
MAXSEQS	4	matches the per-replica share of the fleet's concurrency
CHUNK	8192	prefill chunk (upstream default), coupled to the context: the prefill arena grows with the chunk size and shares the device budget with the KV cache. Measured: CHUNK=16384 leaves 4.22 GiB of KV against the 5.21 GiB that 114688 context needs, and the engine refuses to start (To serve at least one request with the model's max seq len (114688), 5.21 GiB KV cache is needed, which is larger than the available KV cache memory (4.22 GiB)). Raising CHUNK therefore means cutting MAXLEN; single-card prefill measured ≈2.5k tok/s at 8192 (~15k-token prompt, 5992 ms TTFT)
GPU_UTIL	0.98	upstream default; the card is dedicated
SPEC_METHOD / SPEC	dflash / 7	the z-lab DFlash2-FP8 block-diffusion drafter. Depth 7 is upstream's production default and was re-measured here on a code corpus built from this repo (see Speculative depth); depth 5 wins batch throughput, 7 wins single-stream latency, and this fleet is interactive
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-Qwen3.8-27B-Uncensored-f16.gguf (beside the target)	vision enabled; this build has NO autoload, so the flag is explicit and the serve script refuses to start without the file

Mounts are the TP=2 unit's minus nothing: /models, the HF cache, the shared compile cache, /patches (the pinned radiance clone; its patches are applied at every container start), /r4d (the patched libr4d — the image's shipped r4d.so predates the gated-delta-net overflow fix and NaNs this model), and the entrypoint/healthcheck/proxy scripts mounted read-only out of /home/dank/serverAppConfigs.

The healthcheck reads PORT, so one probe serves the replicas (:7770), the proxy (:7770) and the TP=2 target.

Boot persistence, verified by rebooting

The trio was reboot-tested end to end (2026-09-16), because this repository has been burned before by services that ran only from compose.yaml and stayed dark after a reboot:

sudo reboot                 # unattended; ~2 min 14 s to graphical.target
# ... no manual step at any point ...
systemctl --user is-active frank-medium-a frank-medium-b frank-medium-proxy
  -> active active active            (containers up 59 s after boot)
# ...
curl http://127.0.0.1:7770/health        # both replicas UP ~3.5 min after boot

Why it comes back: each unit's [Install] WantedBy=default.target — the generator wires all three into default.target's wants (they report generated, which is the healthy state on this host's v1 quadlet generator and cannot be enabled) — plus Linger=yes for the deploy user, so the user manager runs without a login session. The replicas need no warm cache: the compile cache lives on disk (/cache) and the model loads cold in about three minutes.

Deploy and verify

./tests/run.sh                       # unit/config seams
./scripts/deploy.sh dankserver --service frank-medium

The handler re-tags docker.io/stilldeadcode/vllm-radiance:0.9.3 as localhost/radiance-vllm:0.9.3 (this host's quadlet path cannot resolve registry-prefixed refs from local storage at create time), installs the three units, retires any stale TP=2 unit into the trash, restarts the trio and waits for http://127.0.0.1:7770/health.

Verification that means something:

curl -s http://192.168.0.235:7770/health        # proxy: names any replica that is down
curl -s http://192.168.0.235:7770/v1/models     # served name must be frank-medium
# and a real completion through the proxy:
curl -s http://192.168.0.235:7770/v1/chat/completions -H 'Content-Type: application/json' \
  -d '{"model":"frank-medium","messages":[{"role":"user","content":"Say OK"}],"max_tokens":400}'

The same call works against either replica directly (:7771, :7772) — when one of them is down the other still serves, which is why the proxy reports the individual replica states instead of a single boolean.

A replica that dies mid-flight does not cost a client a request: a failed connection demotes that replica immediately and the proxy retries on the live one (vLLM's endpoints are stateless, so re-sending a request that never reached an engine is safe) and logs the failover. Verify with card 1 stopped:

systemctl --user stop frank-medium-a.service
for i in 1 2 3 4 5 6; do
  curl -s -m 30 -o /dev/null -w "%{http_code} " -H 'Content-Type: application/json' \
    -d '{"model":"frank-medium","messages":[{"role":"user","content":"probe"}],"max_tokens":8}' \
    http://127.0.0.1:7770/v1/chat/completions
done; echo
podman logs --since 60s frank-medium-proxy | grep -i failover
systemctl --user start frank-medium-a.service

Six 200s and one proxy: failover — … unreachable line is the pass condition. Restarting a stopped replica takes about 150 s (warm compile cache).

Benchmarks

The harnesses live in the repository — the numbers below are reproducible with scripts/bench-vllm-corpus.py (build the code corpus), scripts/bench-vllm-serve.sh (concurrency sweep), scripts/bench-turn-latency.py (per-turn TTFT for an agent-loop shape, which is what the routing decision optimizes) and scripts/bench-llm-soak.py (sustained multi-family load). Each carries the traps in its header: the tokenizer argument that makes a vllm bench run report nothing, the salted prompts that keep a run from inheriting a warm cache, and why the corpus is code-shaped rather than random.

vllm bench serve against the proxy (both cards serving), random dataset, --ignore-eos, 2026-09-16:

Case	Output t/s	TTFT p50	TPOT p50	Acceptance
conc 1, in 1k / out 512	42.8	382 ms	23.6 ms	43.1 %
conc 4, in 1k / out 512	147.6	471 ms	25.3 ms	34.2 %
conc 8, in 1k / out 512	230.8	913 ms	28.8 ms	39.0 %
conc 1, in 32k / out 32	2.5	11837 ms	29.6 ms	25.0 %

The 32k row is a prefill measurement: 32768 prompt tokens in 11.8 s is ≈2.8k tok/s of prefill on a single card. The replicas split the fleet's traffic, they do not pool capacity for one request, against the upstream reference's 4585 t/s at 32k on TP=2 — one card doing the whole prompt lands where you would expect.

Read these against the upstream repo's own reference table (2x R9700, TP=2, BetterBench corpus): step time 22.8-23.2 ms per update, combined decode 224 t/s, prefill 4585 t/s at 32k. Step time (TPOT) is the comparable number — tokens/s swings with draft acceptance, and synthetic random prompts accept far worse than the code/prose mix upstream measures (43 % here against 5.8 tokens/update there), so a lower token rate at the same step time is the dataset, not the deployment.

Bench caveat: the client must be given --tokenizer /models/Qwen3.8-27B-Uncensored-MXFP4-awq with /models mounted. Passing only --model frank-medium makes the client try to resolve that alias on HuggingFace and die before sending a request — a green-looking run that produced no numbers at all.

Speculative depth

Upstream ships depth 7 for dflash and notes the choice is content-dependent; it was measured for TP=2 on their corpus. Re-measured here on one card against a code corpus built from this repository's own sources (16 prompts, 512 output tokens each, two runs per setting):

Depth	conc 1	conc 4
SPEC=7 (deployed)	46.2 tok/s, TPOT 20.8 ms, 3.86 tok/update	118.7 / 127.0 tok/s, TPOT 25.1-29.3 ms
SPEC=5	41.5 tok/s, TPOT 23.1 ms, 3.41 tok/update	136.1 / 129.4 tok/s, TPOT 25.3-26.6 ms

The step time is nearly identical at both depths (≈79 ms against ≈80 ms at conc 1) — depth 7 wins single-stream because it lands more tokens per step (3.86 against 3.41), and depth 5 wins at conc 4 by ≈8 % because more, narrower verify steps pack better. Depth 7 stays: this fleet serves about three people in interactive agent loops, where per-turn latency is the number that is felt, and it is also the better default on code at low concurrency. Switch a replica to SPEC=5 if batch throughput ever matters more than turn latency — it is one environment line in the unit, at the cost of a recompile (the spec depth is part of the graph cache key).

Capacity: how many big conversations fit

Each replica's own ledger at boot states the ceiling:

GPU KV cache size: 122,636 tokens
Maximum concurrency for 114,688 tokens per request: 1.07x
[radiance] kv cache groups: size 8, 9 groups, 202 blocks/request (upstream would pick size 5)
Add 3 padding layers, may waste at most 60.00% KV cache memory

So one card holds ~122k tokens of KV — one full-context conversation plus a little, or about two 30k-token agent loops with their cached prefixes. A 3-family soak (three concurrent loops, each re-sending a ~30k-token document every turn with growing history) shows what happens at that limit: turns alternate between ~1.1 s when the prefix is resident and 12-22 s when it has been evicted and must be prefilled again, because two families per card plus their cached prefixes exceed the pool. Nothing failed, and per-family pinning kept the fast turns on the replica that holds the prefix.

Read that as a capacity budget, not a defect: two concurrent large-context sessions per card are comfortable, three or more of that size trade turns against each other. The soak ran 155 requests, 155 answers, no errors, no failover and both replicas steady — the cost shows up as latency, not failure.

CHUNK is not a lever for this. The arena/KV budget is not monotonic in the chunk size, and 8192 is a sweet spot rather than a compromise — measured available KV per card: 5.44 GiB at CHUNK=8192 (122,636 tokens), 4.22 GiB at 16384, and 4.16 GiB at 4096, where the engine refuses to start at all (5.0 GiB KV cache is needed, which is larger than the available KV cache memory (4.16 GiB)). Both directions from the deployed value cost capacity, so the real lever for more concurrent large contexts is another card.

Placement does not rescue an overloaded working set

If families compete for cache space, the obvious idea is to place a new conversation on the card carrying less context instead of alternating. That was implemented (decayed prompt-token estimate per replica, new families to the lighter card, continuations still pinned) and measured against the counter on the workload it should help: one ~80k-token family plus two ~50k-token families, each turn re-sending the same document.

placement	heavy	medium 1	medium 2
counter (13 min)	13 turns, median 25.4 s, 3 cached (23 %)	14 turns, median 34.7 s, 2 cached (14 %)	13 turns, median 45.2 s, 2 cached (15 %)
balanced by load (~8 min)	5 turns, median 46.7 s, 1 cached (20 %)	7 turns, median 23.4 s, 0 cached	6 turns, median 34.3 s, 0 cached

No improvement — the load arm shows the same near-total loss of cache reuse (note its smaller sample: the arm was stopped early once the pattern was clear), so the change was reverted. routing.py keeps the counter and the policy surface stays as small as it was.

The reason is arithmetic rather than policy: a running request's live sequence and its cached prefix draw on the same pool, so a card that is running two 50k-token generations has already exceeded its 122,636 tokens before any cache entry is considered. Placement can decide which card is over, not whether one is. Three concurrent large sessions do not fit in two cards' capacity, and the levers are therefore scheduling (fewer large contexts in flight, which is what the harness's context compaction does) or a third card.

Caching itself is healthy; the sum is not

Before blaming the cache, it was measured directly: one card, an unloaded engine, a ~50k-token prompt sent three times — 22.3 s cold, then 1.4 s, then 1.4 s — and the same result with a varying tail between turns (22.4 s cold, 1.41 s warm). So long-prefix reuse works exactly as the pinning policy assumes, and the engine's lifetime "Prefix cache hit rate: 29 %" line is an average over every benchmark prompt a replica has ever seen, not a statement about any one request.

Serialising large requests does not fix the overload either. Admission was implemented — at most one >16k-token request per card at a time (LARGE_TOKENS / LARGE_LIMIT), small requests untouched — and run against the same three-family workload:

	heavy	medium 1	medium 2
no admission	13 turns, median 25.4 s, 23 % cached	34.7 s, 14 %	45.2 s, 15 %
one large per card	5 turns, median 48.2 s, 20 % cached	22.7 s, 14 %	22.6 s, 33 %

No clear win on small samples, so it was reverted too. The reason is now visible in the arithmetic: with three families on two cards, one card must hold two families, and its pool has to carry 2 × 50k of cached prefix plus the 50k of the sequence currently generating — 150k against 122,636. Those two co-tenants evict each other turn by turn whatever the concurrency limit is. A card can hold two large families' caches, or one cache plus a live generation, but not both. That is the wall, and it is capacity, not scheduling.

Watching it instead of measuring it by hand

Everything above was diagnosed with ad-hoc probes. The same signals are now scraped — Prometheus has frank-medium and frank-dragoon jobs (shared/prometheus/prometheus.yml), so the numbers that matter have history: vllm:kv_cache_usage_perc (per-card pool fill), vllm:num_requests_waiting (queue depth), rate(vllm:prefix_cache_hits_total[5m]) / rate(vllm:prefix_cache_queries_total[5m]) (whether pinned conversations are served warm), and the request-latency histograms. See Observability, and one Grafana dashboard reads them (LLM fleet — capacity and reuse).

Five rules watch this service (shared/prometheus/rules/llm-fleet.yml): pool saturation above 90 %, requests queueing for 10 minutes, prefix reuse collapsing below 20 % with traffic, the scrape target down, and p95 latency above 60 s — each threshold tied to a measurement on this page rather than to a guess. They evaluate and show up on the dashboard's Alerts firing panel, but nothing is notified: this stack has no Alertmanager, so an alert reaches a person only if that person is looking. That is the honest state of "alerting" here.

Citations

• vLLM upstream: https://github.com/vllm-project/vllm — the R4D one-shot allreduce and the tensor-parallel path this deployment runs.
• ROCm/clr rocclr/device/rocm/rocvirtual.cpp:3963 — the hostcall-buffer refusal; rocclr/device/rocm/rocdevice.cpp:186 — checkAtomicSupport.
• Linux drivers/pci/pci.c — pci_enable_atomic_ops_to_root() and its PCI_EXP_DEVCAP2_ATOMIC_ROUTE requirements on upstream bridges.
• Codeberg StillDeadcode/vllm-radiance issue #22 — an independent host with the same image and kernel 7.0.0-30 failing identically.
• ggz14/radiance-vllm-mxfp4 — the patch set and the reference benchmark table this page compares against.

---