The two replicas behind :7770 were on plain round-robin, which throws away prefix-cache hits: an agent loop re-sends its system prompt and file context every turn, and alternating replicas means every other turn pays a cold prefill. Measured with 90k-token contexts, warm-turn TTFT was 2676 ms. Pinning everything to one replica is not the answer either — a cold batch (the shape a benchmark or a burst of parallel subagents produces) then loses a card's worth of prefill: 16 requests took 30.1 s pinned against 18.6 s spread. So the policy routes by request shape. A continuation (more than one message — the agent-loop signature) is pinned to the replica that served its prefix; anything else spreads across the cards by counter. Measured: 19.97 s for the burst and 1095 ms warm-turn TTFT — the best of both, without a regression on either workload. The policy lives in proxy/routing.py, separate from the HTTP plumbing in proxy/affinity-proxy.py, so it can be tested without aiohttp, a network or a GPU. tests/frank_medium_proxy.bats covers it, and it earned its keep: the first version hashed fresh prefixes to pick a card, which sends every member of a same-prefix burst to one card (29.8 s, measured), and a stale index in the pin path returned 500 on every continuation. Both are now assertions. ROUTING=roundrobin and ROUTING=affinity stay available to re-measure the choice on this host; the default is adaptive.
14 KiB
type, title, description, tags, timestamp
| type | title | description | tags | timestamp | ||||||||||
|---|---|---|---|---|---|---|---|---|---|---|---|---|---|---|
| configuration | frank-medium vllm-radiance MXFP4 Runbook | vllm-radiance MXFP4 serving just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq as frank-medium on Dankserver (2x AMD Radeon AI PRO R9700, gfx1201), as two single-GPU replicas behind an adaptive-routing proxy — and the PCIe-atomics reason TP=2 cannot run on this host. |
|
2026-09-16 |
frank-medium vllm-radiance MXFP4 Runbook
The ACTIVE frank-medium deployment (since 2026-09-15/16): the
stilldeadcode/vllm-radiance:0.9.3 image plus the ggz14/radiance-vllm-mxfp4
patches, serving just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq on the two
Radeon AI PRO R9700 (gfx1201) cards. It replaced the charlie12345/ROCmFPX
llama.cpp deployment (frank-medium-rocmfpx).
Deployed shape: two single-GPU replicas behind an adaptive-routing proxy. The intended shape — one tensor-parallel process across both cards (TP=2) — cannot launch a collective on this host. The reason is a hardware property of the motherboard's PCIe topology, not the image, the RCCL build, or the kernel version. The evidence is in Why TP=2 cannot run here.
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
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 (rocclr/device/rocm/rocvirtual.cpp:3963), and the launch aborts.
hipErrorIllegalState is that refusal, not a driver fault.
Chain of causation
- RCCL's collective kernels carry a hidden hostcall buffer parameter.
- ROCr passes that buffer only when
pcie_atomics_is true for the device; otherwise it returns "not supported" and the dispatch fails. pcie_atomics_is set once, at device init, from the HSA link info between the CPU agent and the GPU (rocclr/device/rocm/rocdevice.cpp:186,Device::checkAtomicSupport). It is not re-read per process.- 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 advertiseAtomicOp Routingcapability (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, NCCL_P2P_DISABLE, NCCL_SHM_DISABLE, NCCL_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.
- Until then, two replicas are how both cards serve.
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-cachinghas 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 small 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.14 available, and 131072 needs 5.58 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/262144 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) |
GPU_UTIL |
0.98 | upstream default; the card is dedicated |
SPEC_METHOD / SPEC |
dflash / 7 | the DFlash2-FP8 block-diffusion drafter, upstream's measured production depth |
SNAP / DRAFTER |
/models/Qwen3.8-27B-Uncensored-MXFP4-awq / /models/Qwen3.8-27B-DFlash2-FP8 |
the served checkpoint and its drafter |
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 (:7771/:7772),
the proxy (:7770) and the TP=2 target.
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.
Benchmarks
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.
Citations
ROCm/clrrocclr/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 itsPCI_EXP_DEVCAP2_ATOMIC_ROUTErequirements on upstream bridges. - Codeberg
StillDeadcode/vllm-radianceissue #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.