---
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
---

> **SUPERSEDED 2026-09-16.** This page is the record of the retired
> vllm-radiance MXFP4 deployment. The ACTIVE frank-medium is mainline
> llama.cpp with the DFlash2 drafter ([[frank-medium-llamacpp-dflash2]]);
> the PCIe-atomics evidence in this page remains the reason the runtime is
> layer-split llama.cpp rather than tensor-parallel vLLM.

# 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.

### 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 `enable`d) —
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 — ~120 requests, no errors, no failover,
both replicas steady — 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-large` 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

- `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.