--- type: configuration title: frank-medium vllm-radiance MXFP4 Runbook description: 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. tags: - frank-medium - dankserver - vllm - vllm-radiance - mxfp4 - gfx1201 - r9700 - rccl - pcie-atomics - quadlet timestamp: 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](#why-tp2-cannot-run-here). ## Architecture ```mermaid graph LR C[clients: Hermes, harness] --> P["frank-medium-proxy
:7770 adaptive routing"] P --> A["frank-medium-a
R9700 #1, TP=1, :7771"] P --> B["frank-medium-b
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](#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 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 (`rocclr/device/rocm/rocdevice.cpp:186`, `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`, `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-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 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), **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), so the trade was not worth taking | | `GPU_UTIL` | 0.98 | upstream default; the card is dedicated | | `SPEC_METHOD` / `SPEC` | dflash / 7 | the 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](#speculative-depth)); depth 5 wins batch throughput, 7 wins single-stream latency, and this fleet is interactive | | `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. ### 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). ## Citations - `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.