research: how to make the two R9700s work together (and one correction to my own claim)
New page: LLM/Dual-R9700-Cooperation.md — registered in the vault index, with the measured evidence rather than the theory. **The answer.** Card #2 sits behind the B650 chipset, whose bridges advertise `AtomicOpsCap: Routing-` (read-only), so ROCr refuses to fill the hostcall buffer RCCL's kernels declare and no RCCL op can dispatch on it. Two paths fix it: an M.2→PCIe x4 riser into the free, CPU-attached M.2_2 slot (~$30; per ASUS's spec page M.2_1 and M.2_2 are CPU-attached, PCIEX16_2 is chipset in x1 mode, and there is no x8/x8 bifurcation on this board), or a board whose CPU x16 bifurcates to x8/x8. Bandwidth is not the constraint: ≈10 MB per token-step ≈ 1.4 ms at Gen4 x4. **New measured results.** A bare 2-rank torch.distributed run: rank 0 (CPU-lane card) passes broadcast/all_reduce/barrier; rank 1 (chipset card) fails all three, and `send`/`recv` fails too — which kills pipeline parallelism as an escape, since vLLM's PP handoff is exactly that primitive. **Correction.** I previously said no software path exists. That overshot. One is open and unverified: recompiling RCCL so its kernels stop declaring a hostcall buffer (`-mprintf-kind=buffered` is the claim). Counter-evidence recorded: RCCL's device source contains no live printf/assert/__assertfail, so the trigger is unexplained, and the claim traces to search summaries, not a primary source. The page carries the two-command test that would settle it. Also corrected: 0.9.3 is radiance's version, not vLLM's — the base is vLLM 0.27.1 / torch 2.11.0 / ROCm 7.14, and its shipped RCCL exposes no `ncclDevComm*`, so torch's symmetric-memory path cannot run. **Rule added to the benchmarking policy**: it covers ANY GPU experiment, not just benchmarks — including a bare 2-rank probe. Running one against the live service on 2026-09-16 preceded the host going hard down (no ARP, no ICMP, power cycle required); a hung AQL dispatch wedging the driver is the plausible mechanism. The policy also now says how to tell off from wedged: `ip neigh` from a peer.
This commit is contained in:
Frank
@@ -0,0 +1,158 @@
|
||||
---
|
||||
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 used for the replicas
|
||||
(`scripts/bench-vllm-corpus.py` / `bench-turn-latency.py`) 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.
|
||||
|
||||
## 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
|
||||
Reference in New Issue
Block a user