Upstream ships dflash depth 7 as a content-dependent default measured for TP=2
on their corpus. Re-measured on one card here, two runs per setting, 16
repo-derived prompts at 512 output tokens:
SPEC=7 conc1 46.2 tok/s, TPOT 20.8 ms, 3.86 tok/update
conc4 118.7 / 127.0 tok/s
SPEC=5 conc1 41.5 tok/s, TPOT 23.1 ms, 3.41 tok/update
conc4 136.1 / 129.4 tok/s
Step time is nearly identical at both depths (~79 ms vs ~80 ms single stream):
7 wins by landing more tokens per step, 5 wins at conc 4 by ~8% because more
narrow verify steps pack better. Depth 7 stays — this fleet is interactive
agent loops where per-turn latency is what is felt — and the runbook now
records the numbers and the one-line switch for a batch-first workload.
293 lines
15 KiB
Markdown
293 lines
15 KiB
Markdown
---
|
||
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<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](#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.
|