Frank
c97e359bc2
The LLM index still opened with vllm-radiance as frank-medium's runtime and drew the replica pair in its topology diagram; the cooperation page's thesis (collectives are impossible, so the answer is hardware) now has to name the path the fleet actually took: llama.cpp's layer split, which launches no RCCL kernel on either card. That reframes the M.2 riser as a means to TENSOR parallelism rather than a means to using both cards at all.
182 lines
10 KiB
Markdown
182 lines
10 KiB
Markdown
---
|
|
type: research
|
|
title: Dual R9700 Cooperation
|
|
description: Why frank-medium's two Radeon AI PRO R9700s run as independent replicas instead of tensor-parallel, which options are ruled out with evidence, and the two hardware paths that restore true TP=2.
|
|
tags:
|
|
- frank-medium
|
|
- rocm
|
|
- rccl
|
|
- pcie-atomics
|
|
- tensor-parallelism
|
|
- dankserver
|
|
timestamp: 2026-09-16
|
|
---
|
|
|
|
# Dual R9700 Cooperation
|
|
|
|
Dankserver's `frank-medium` runs as **two independent single-GPU replicas**
|
|
(`:7771` on card #1, `:7772` on card #2) behind an adaptive-routing proxy on
|
|
`:7770`. This page answers why, and what would let the two cards work on one
|
|
model instead. See [[frank-medium-radiance-mxfp4]] for the deployment itself.
|
|
|
|
## The short answer
|
|
|
|
**Card #2 sits behind the B650 chipset, whose bridges cannot route PCIe atomic
|
|
operations.** RCCL's kernels carry a hidden hostcall-buffer parameter, and ROCr
|
|
refuses to supply it to a device without atomics — so *no* RCCL collective or
|
|
P2P operation can be dispatched on card #2. Two paths restore cooperation:
|
|
|
|
1. **Move card #2 to the free, CPU-attached M.2_2 slot** (~$30 M.2→PCIe x4
|
|
riser). Cheapest, and bandwidth is not a constraint (below).
|
|
2. **A board whose CPU x16 bifurcates to x8/x8** (X670E/X870E class,
|
|
~$250-400). Best performance. This board has no x8/x8 mode.
|
|
|
|
## The gate, in source
|
|
|
|
```cpp
|
|
void Device::checkAtomicSupport() { // rocclr/device/rocm/rocdevice.cpp:181
|
|
std::vector<amd::Device::LinkAttrType> link_attrs;
|
|
link_attrs.push_back(std::make_pair(LinkAttribute::kLinkAtomicSupport, 0));
|
|
if (findLinkInfo(cpu_agent_info_->fine_grain_pool, &link_attrs)) {
|
|
if (link_attrs[0].second == 1) {
|
|
info_.pcie_atomics_ = true; // no else branch
|
|
}
|
|
}
|
|
}
|
|
```
|
|
|
|
The flag is a **queried hardware property**, consumed at
|
|
`rocvirtual.cpp:3954` (the hostcall/AQL dispatch path), in `rocblit.cpp` (SDMA
|
|
copy paths) and exported as `hipDeviceProp.hostNativeAtomicSupported`. There is
|
|
no environment variable, module parameter, or sysfs toggle around it — verified
|
|
by grepping ROCclr for `getenv`/`HSA_*`/`ROCR_*`/`AMD_*` in the gate's
|
|
neighborhood: none. Writing the endpoint's Device Control 2 bit 6 by hand
|
|
(`setpci -s 0b:00.0 CAP_EXP+28.w=0440`) does nothing because the value is
|
|
consumed at driver probe time.
|
|
|
|
What has to be true upstream of it: every PCIe bridge on the GPU's path to the
|
|
CPU must advertise `AtomicOpsCap: Routing+`. On this board card #2's path
|
|
includes `05:00.0` and `06:09.0`, both `Routing-` — a **read-only capability
|
|
bit**, so no software can enable it.
|
|
|
|
## Measured matrix
|
|
|
|
Bare 2-rank `torch.distributed` (`backend=nccl`) on the two R9700s, no vLLM, no
|
|
model. Rank 0 = card #1 (CPU lanes), rank 1 = card #2 (chipset):
|
|
|
|
| Operation | rank 0 (card #1) | rank 1 (card #2) |
|
|
|---|---|---|
|
|
| `broadcast` | OK | **FAIL** `DistBackendError` |
|
|
| `all_reduce` | OK | **FAIL** |
|
|
| `barrier` | OK | **FAIL** |
|
|
| `send` / `recv` (P2P) | blocked, then timeout | **FAIL** |
|
|
| `all_gather` (after failure) | group unusable | — |
|
|
|
|
The `AMD_LOG_LEVEL=3` trace names the cause: `Pcie atomics not enabled, hostcall
|
|
not supported` → `AQL dispatch failed!` → `hipLaunchKernel` returns
|
|
`hipErrorIllegalState` for `ncclDevKernel_Generic_4(...)`. Nothing in vLLM or
|
|
RCCL mentions atomics; the runtime log is the only place it appears.
|
|
|
|
**`send`/`recv` failing is the decisive result**: it kills pipeline parallelism
|
|
as an option, because vLLM's PP stage handoff is exactly that primitive.
|
|
|
|
## Software escapes: closed, and one still open
|
|
|
|
| Idea | Finding |
|
|
|---|---|
|
|
| vLLM's IPC custom allreduce | **Gated off by a table keyed to CUDA sm versions** — `{9.0, 10.0, 10.3, 10.7}`; ROCm reports capability `12.0` for gfx1201, so the gate is `False`. `custom_ar` itself is compiled into the image. |
|
|
| ...and if that table were patched | Still insufficient: `broadcast`, `barrier`, and the TP logits `all_gather` have no non-RCCL path. |
|
|
| The radiance fork's own allreduce | Already present and *on by default* (`RADIANCE_USE_R4D_AR`, size gate `RADIANCE_AR_MAX_KB`): a one-shot P2P push allreduce wrapping `CudaCommunicator.all_reduce`, exchanging IPC handles over the **gloo** group. With the gate raised it covers prefill and decode, and `ncclDevKernel` disappears from the profile. But it does not intercept `PyNcclCommunicator.__init__`'s **warmup allreduce**, nor `pynccl_comm.all_gather` — both direct RCCL calls. |
|
|
| NCCL symmetric memory | Unavailable: the shipped RCCL advertises NCCL 2.30.4 but exposes no `ncclDevComm*` API, so torch's symmetric-memory path cannot run. |
|
|
| Prefill/decode disaggregation | Needs a KV-transfer connector plus coordination collectives; both out. |
|
|
| gloo for GPU work | CPU/TCP collectives — orders of magnitude too slow for per-layer TP. |
|
|
| **Rebuild RCCL so its kernels do not declare a hostcall buffer** | **OPEN — unverified.** The claim is that the hostcall requirement is triggered by device-side `printf`/`assert`, and that `-mprintf-kind=buffered` avoids it. Counter-evidence: RCCL's device source at `src/device/` contains **no live** `printf`, `assert`, or `__assertfail` (checked), so what makes its kernels declare `HiddenHostcallBuffer` is unexplained. The claim traces to search-result summaries, not a primary source. Cheap to settle once a host is available: recompile librccl with `-mprintf-kind=buffered`, dump a kernel's declared arguments (`roc-obj-extract` + `llvm-readelf --notes`) and look for `hidden_hostcall_buffer`; if it is gone, a 2-rank probe should dispatch. **Stop and mask the service first** — see the note below. |
|
|
|
|
## Bandwidth: why x4 is enough
|
|
|
|
A 27B TP=2 allreduce at batch 4 moves roughly
|
|
$2 \times \text{batch} \times \text{hidden} \times 2\ \text{bytes} \approx 160\ \text{KB}$
|
|
per layer per forward pass; across ~62 layers that is **≈10 MB per token-step**,
|
|
or **~1.4 ms at PCIe 4.0 x4** (~7.9 GB/s). TP is compute-bound here, not
|
|
interconnect-bound. That is the whole reason the M.2 riser is the interesting
|
|
option rather than the board swap.
|
|
|
|
## Ruled out empirically (do not re-derive)
|
|
|
|
- Kernel versions 7.0.0-28 / -30 / -31 — byte-identical failure.
|
|
- `iommu=pt`, and `amd_iommu=off` is unrelated (that was a Dankclaw prefill knob).
|
|
- P2P hardware: 45.274 GB/s bidirectional between the cards works fine — plain
|
|
memory traffic is not the problem; *atomic* routing is.
|
|
- A from-source librccl rebuild: identical failure at the identical call path.
|
|
- Every NCCL transport/algorithm knob.
|
|
- A one-card test proves nothing: NCCL short-circuits at `world_size=1`, which
|
|
is exactly why the two replicas work at all.
|
|
- Independent corroboration: StillDeadcode/vllm-radiance Codeberg issue #22 —
|
|
different host, same image, kernel 7.0.0-30, same `enqueue.cc` failure.
|
|
|
|
## Test plan (when the hardware changes)
|
|
|
|
The verdict to look for is one line of `lspci`, before touching any model:
|
|
|
|
```bash
|
|
sudo lspci -vv -s <card2> | grep AtomicOpsCtl # want: ReqEn+
|
|
```
|
|
|
|
If that reads `ReqEn+`, the rest is mechanical: `--tensor-parallel-size 2` with
|
|
`HIP_VISIBLE_DEVICES=0,1`, then the same benchmark set
|
|
(`scripts/bench-turn-latency.py`, and the endpoint's own `timings` block) so the
|
|
two architectures are compared on identical workloads. If it still reads
|
|
`ReqEn-`, the card is still on a non-atomic path — check `lspci -tv` for the
|
|
bridge chain rather than assuming the move worked.
|
|
|
|
**Rule for any future multi-GPU collective experiment: stop and mask the
|
|
production service first** (`sudo systemctl stop frank-medium && sudo systemctl
|
|
mask frank-medium`, per the benchmarking policy), or run it on a host that is
|
|
not serving users. On 2026-09-16 a bare RCCL P2P probe was run against the live
|
|
service; the host went hard down ~20 minutes later (no ARP, no ICMP, required a
|
|
power cycle). The mechanism is not proven, but a hung AQL dispatch wedging the
|
|
GPU driver is a plausible cause and this host has precedent for going
|
|
unreachable until a power cycle.
|
|
|
|
|
|
## The third path, and the one in production: llama.cpp's layer split
|
|
|
|
Everything above is about *collectives*, and vLLM cannot proceed without them
|
|
in any of its parallelisms — TP needs a per-layer allreduce, and PP's stage
|
|
handoff is a p2p `send`/`recv` of the hidden state, the primitive this host
|
|
also fails. **llama.cpp's `--split-mode layer` needs no collective at all**:
|
|
layers are distributed across the cards and the boundary activation is a
|
|
plain per-backend copy, once per forward pass. No RCCL kernel is ever launched
|
|
on the second card, so the atomics gate never comes up.
|
|
|
|
That is what `frank-medium` runs today ([[frank-medium-llamacpp-dflash2]]): one
|
|
process, one model, both cards, 262144 context. Measured against tensor split
|
|
on this host: 1991 t/s prefill / 18.8 tok/s decode layer-split, against 272 /
|
|
26.3 tensor (butterfly fallback, no RCCL). Tensor would still win decode on a
|
|
build with a working RCCL — which is exactly what the M.2 riser or a board swap
|
|
in [the short answer](#the-short-answer) would buy.
|
|
|
|
So the ranking changed on 2026-09-16: software already gets both cards onto one
|
|
model, and the hardware fix is no longer a means to that end — it is a means to
|
|
*tensor* parallelism, which buys ~40 % decode and would let this runtime drop
|
|
layer-split's serial stages entirely.
|
|
|
|
## What would change the answer
|
|
|
|
- An RCCL build whose kernels stop declaring a hostcall buffer — the one
|
|
software route that is *not* closed. See the last row of the escape table for
|
|
the hypothesis, its counter-evidence, and the two-command test.
|
|
- The symmetric-kernel path inside RCCL, if a release ever ships the device API
|
|
(`ncclDevComm*`) that torch's symmetric-memory path needs. The shipped
|
|
release does not.
|
|
- Firmware or an AGESA update that makes the chipset bridges advertise atomic
|
|
routing — a capability bit, so this would show up in `lspci` on either side.
|
|
- A GPU-side workaround: none found in ROCclr, and the gate is per-device,
|
|
queried once at probe.
|
|
|
|
## See Also
|
|
|
|
- [[Pipeline Parallelism/index|Pipeline Parallelism]] — two-machine research (the PP variant of this problem)
|
|
- [[Dankserver System Specs]] — hardware inventory
|
|
- [[Observability]] — the scrape/dashboard/alert stack around these endpoints
|