dank/serverAppConfigs
1
0
Fork 0
Code Issues 6 Pull Requests Actions Packages Projects Releases Wiki Activity

frank-medium: deploy two single-GPU replicas — TP=2 is blocked by PCIe atomics

Browse Source 
The radiance MXFP4 stack cannot run TP=2 on this host. RCCL's collective
kernels carry a hidden hostcall-buffer parameter, and ROCr refuses to
dispatch them on a device without PCIe atomics: the AQL dispatch fails with
hipErrorIllegalState ("the operation cannot be performed in the present
state"), which is what every two-rank collective hit. The second R9700 sits
behind the chipset bridges 05:00.0/06:09.0, which advertise
`AtomicOpsCap: Routing-` — a read-only capability, so no software can enable
atomic routing on that path and the driver's
pci_enable_atomic_ops_to_root() fails for that card. The first card, on CPU
root port 00:01.1, has atomics enabled.

Ruled out: kernel version (7.0.0-28/-30/-31 identical, after a full upgrade),
iommu=pt, P2P hardware (45.3 GB/s between the cards), the shipped librccl (a
from-source rebuild reproduces the failure at the same call path), vLLM and
pynccl (a bare two-rank torch.distributed allreduce fails the same way), all
NCCL transport knobs, and single-rank tests (NCCL short-circuits at
world_size=1, so one-card success proves nothing).

Deployment: three rootless Quadlets — frank-medium-a (:7771, card 1),
frank-medium-b (:7772, card 2), both TP=1 at 114688 context (262144 and
131072 do not fit one card's KV) with vLLM-profiled KV — and
frank-medium-proxy (:7770), a round-robin that preserves the fleet contract
of one endpoint and one served model name. frank-medium.container stays as
the TP=2 target for a host whose second card is atomic-capable; the deploy
handler does not install it and trashes any stale copy.

Measured through the proxy (vllm bench serve, random dataset, --ignore-eos):
conc 1 42.8 tok/s, conc 4 147.6 tok/s, conc 8 230.8 tok/s, 32k prefill at
~2.8k tok/s per card, TPOT p50 23.6-28.8 ms.

Also: healthcheck reads PORT so one probe serves all three units, the hermes
example and the harness models.yml drop to the real 114688 context, and the
runbook records the full evidence chain.
This commit is contained in:
Frank
2026-09-16 03:47:11 -07:00
parent ad7a9c42f9
commit 4fca739bf8
15 changed files with 817 additions and 105 deletions
Download Patch File Download Diff File Expand all files Collapse all files
serverAppConfigs-docs/LLM/frank-medium-radiance-mxfp4.md
+236
View File
@@ -0,0 +1,236 @@
---
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 a round-robin 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 a round-robin 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 round-robin"]
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 | round-robin |
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 round-robin needs no session affinity.
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.
## 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 | vLLM profiles this shape; the upstream pin (2x7551-32624) is a TP=2/262144 measurement and is invalid here |
| `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/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.
serverAppConfigs-docs/LLM/frank-medium-rocmfpx.md
+6
View File
@@ -13,6 +13,12 @@ tags:
timestamp: 2026-09-15
---
> **PRIOR deployment (retired 2026-09-16).** This page documents the
> llama.cpp era of `frank-medium`, which served the endpoint for part of
> 2026-09-15 before the vllm-radiance MXFP4 stack replaced it. The ACTIVE
> deployment is [[frank-medium-radiance-mxfp4]]; this page is kept as the
> record of the runtime it replaced.
# frank-medium ROCmFPX Runbook
`frank-medium` is the Dankserver LLM endpoint (port 7770) since 2026-09-15,
serverAppConfigs-docs/LLM/index.md
+11 -4
View File
@@ -48,10 +48,17 @@ i1-Q4_K_M GGUF, 262k ctx / 1 slot, vision]
## Active Deployment
- [[frank-medium-rocmfpx]] — the ACTIVE frank-medium runbook:
charlie12345/ROCmFPX llama.cpp serving kingjones777/
Qwen3.8-27B-ROCmFPX-Q8_0-AGENT-GGUF (MTP draft on) on Dankserver's
2x Radeon AI PRO R9700 (since 2026-09-15; replaced frank-small)
- [[frank-medium-radiance-mxfp4]] — the ACTIVE frank-medium runbook: the
vllm-radiance MXFP4 stack (stilldeadcode/vllm-radiance:0.9.3 +
ggz14/radiance-vllm-mxfp4 patches) serving
just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq on Dankserver's 2x Radeon
AI PRO R9700 as two single-GPU replicas behind a round-robin proxy, with
the PCIe-atomics evidence that rules out TP=2 on this host (since
2026-09-15/16)
- [[frank-medium-rocmfpx]] — PRIOR frank-medium deployment (llama.cpp era,
2026-09-15): charlie12345/ROCmFPX serving
kingjones777/Qwen3.8-27B-ROCmFPX-Q8_0-AGENT-GGUF; kept as the record of the
runtime it replaced
- [[freetoken-frank-small]] — RETIRED 2026-09-15 with the RTX 3090: the
FreeToken v0.1.2 frank-small runbook (batsclamp Huihui-Qwen3.6-35B-A3B-
abliterated-FP8), kept as the retirement record