--- type: research title: Dual R9700 Cooperation description: Why frank-medium's two Radeon AI PRO R9700s could not run tensor-parallel on the old B650-E board (PCIe atomic routing), what the 2026-09-20 ProArt B850-Creator swap unblocked, and what still has to be probed. tags: - frank-medium - rocm - rccl - pcie-atomics - tensor-parallelism - dankserver timestamp: 2026-09-16 --- # Dual R9700 Cooperation Dankserver's two R9700s run **one tensor-parallel engine across both cards** — today that is `frank-archon` under radiance 1.0.13, TP=2 ([[frank-archon]]). When this page was written the endpoint was `frank-medium`: first llama.cpp's layer split ([[frank-medium-llamacpp-dflash2]], 2026-09-16 → 09-20), then from 2026-09-20 one vLLM-radiance process at `tensor_parallel_size 2` ([[frank-medium-radiance-mxfp4]], 2026-09-20 → 09-26). This page records why true tensor parallelism was impossible on the previous board (TUF GAMING B650-E WIFI), and what the 2026-09-20 board swap changed. ## The short answer **The wall is gone.** On the B650-E, card #2 sat behind the chipset, whose bridges could not 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 could be dispatched on that card — see the [measured matrix](#measured-matrix) and the [2026-09-16 table](#the-same-gate-as-this-board-reports-it-2026-09-16) below. On 2026-09-20 the board became an **ASUS ProArt B850-Creator WiFi NEO**, which bifurcates the CPU's GPP ports **x8/x8**: the board-swap path this page recommended, at B850 rather than X870E money. Both cards now hang off CPU root ports (`00:01.1`, `00:01.3`), every bridge on both chains reads `AtomicOpsCap: Routing+`, and **both endpoints read `AtomicOpsCtl: ReqEn+`** — the exact precondition this page's own test plan named, and it is now [measured](#measured-verdict-2026-09-20) rather than assumed. What that bought: - **Both cooperation modes now work, and they were both run.** vLLM TP=2 booted and served (its R4D allreduce needs no RCCL collective at all), and llama.cpp `--split-mode tensor` booted and served against an RCCL-backed build — so a collective HAS been dispatched on this host since the swap. - **TP=2 is the production shape.** It wins on every axis measured, by 2.1x on decode and 3.6x at 4-way concurrency, and it replaced the layer split on the day of measurement. The payload has moved since: the MXFP4 `frank-medium` shape is retired 2026-09-26, and what runs today is the iq4r publisher under radiance 1.0.13 ([[frank-archon]]). - **The drafter, not the collective, is what llama.cpp lost.** Its tensor mode could not carry the DFlash2 drafter, which is the honest caveat on the comparison below. - **The bandwidth story moved too.** Both cards ride Gen5 x8 (~31.5 GB/s each): card 0 effectively unchanged (was Gen4 x16), card 1 up ~16x (was Gen4 x1, ~1.97 GB/s). The M.2→PCIe riser plan is dead — `PCIEX16_2` holds card #2 now and M.2_2 is free. ## The gate, in source ```cpp void Device::checkAtomicSupport() { // rocclr/device/rocm/rocdevice.cpp:181 std::vector 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. ### The same gate, as this board reports it (2026-09-16) `lspci -vv` on dankserver. Card 0's chain is root port `00:01.1` → `01:00.0` → `02:00.0` → `03:00.0`; card 1's is `00:02.1` → `05:00.0` → `06:00.0` → `06:09.0` → `09:00.0` → `0a:00.0` → `0b:00.0`: | Device | Role | `AtomicOpsCap` | `AtomicOpsCtl` | Link | |---|---|---|---|---| | `00:01.1` | CPU root port → PCIEX16_1 | Routing+ | — | x16 | | `03:00.0` | **card 0** | 32bit+ 64bit+ | **ReqEn+** | x16 | | `00:02.1` | CPU root port → chipset | **Routing+** | — | x4 | | `05:00.0`, `06:00.0`, `06:09.0` | B650 bridges | **Routing−** | — | x1 at `06:09.0` | | `09:00.0`, `0a:00.0` | PCIe switch ports | Routing+ | — | x1 (downgraded) | | `0b:00.0` | **card 1** | 32bit+ 64bit+ | **ReqEn−** | x16 | Three things fall out of that table, none of them in the source alone: - **The CPU side is fine** — `00:02.1` is `Routing+`. The wall is inside the chipset, not the platform, which is why a slot on CPU lanes is the fix. - **The endpoint's own bit differs between the cards, and that is the kernel's doing.** `pci_enable_atomic_ops_to_root()` walks the bridge chain and refuses to set Requester Enable when any bridge cannot route atomics. ROCclr's `pcie_atomics_` and the endpoint's `ReqEn` are two faces of one cause, and it makes `lspci -vv -s | grep AtomicOpsCtl` a one-command diagnosis: on this host card 0 reads `ReqEn+` and card 1 `ReqEn−`. - **Card 1's effective link is x1**, not the x16 its own capability block advertises: `06:09.0` trained at x1, so ~2 GB/s is the real ceiling. That is a bandwidth fact rather than an atomics one, and it only matters for collectives — a 27B TP=2 allreduce is ≈5.6 ms/token at x1 against ≈1.4 ms at x4. The layer split's boundary copy does not care. ## Re-measured 2026-09-20 (ProArt B850-Creator WiFi NEO) `lspci -vv` after the board swap. Card 0's chain is `00:01.1` → `01:00.0` → `02:00.0` → `03:00.0`; card 1's is `00:01.3` → `05:00.0` → `06:00.0` → `07:00.0`. **No chipset bridge is in either path** — `00:02.1` now carries only LAN/SATA/WiFi: | Device | Role | `AtomicOpsCap` | `AtomicOpsCtl` | Link | |---|---|---|---|---| | `00:01.1` | CPU GPP → PCIEX16_1 | Routing+ | — | Gen5 x8 | | `01:00.0` | card 0 switch upstream | Routing+ | — | Gen5 x8 | | `02:00.0` | card 0 switch downstream | Routing+ | — | Gen5 x16 | | `03:00.0` | **card 0** | 32bit+ 64bit+ | **ReqEn+** | Gen5 x16 (on-card) | | `00:01.3` | CPU GPP → PCIEX16_2 | Routing+ | — | Gen5 x8 | | `05:00.0` | card 1 switch upstream | Routing+ | — | Gen5 x8 | | `06:00.0` | card 1 switch downstream | Routing+ | — | Gen5 x16 | | `07:00.0` | **card 1** | 32bit+ 64bit+ | **ReqEn+** | Gen5 x16 (on-card) | - Both endpoints read `CAP_EXP+28.w = 0x1440` — bit 6 set, which is the kernel's `pci_enable_atomic_ops_to_root()` having walked a chain with no `Routing-` bridge. Same one-command diagnosis, opposite answer. - `dmesg` agrees on bandwidth: `252.056 Gb/s available PCIe bandwidth, limited by 32.0 GT/s PCIe x8 link at 0000:00:01.1` (and `.13`) — ~31.5 GB/s per card. DMI confirms the bifurcation: `PCIEX16(G5)_2`, "8x or x8". - `rocm-smi --showtopo` reported `PCIE` between the cards (Navi 48 has no XGMI), they were in separate IOMMU groups (15 and 20 — re-read 2026-10-18: still 15 and 20), and both root ports had ACS `ReqRedir+`, so this page inferred P2P would hairpin at the root complex. **That inference is superseded:** the live engine reports the two cards reaching each other's memory directly, which [[frank-archon]] and [[Dankserver System Specs]] record. Plain P2P traffic measured 45.27 GB/s bidirectional on the *old* board — DMA was never the gate, atomics were. - Kernel 7.0.0-31-generic with `iommu=pt` then; this host runs 7.0.0-38-generic with `iommu=pt` today (`/proc/cmdline`, read 2026-10-18). Both were ruled out as causes ([above](#ruled-out-empirically-do-not-re-derive)) and this host's IOMMU posture is unchanged. **Still unproven (as of 2026-09-20).** The gate that matters is ROCclr's `checkAtomicSupport()` reading the `kLinkAtomicSupport` link attribute; the ReqEn bit is the same cause seen from the kernel side, but no RCCL kernel had been launched since the swap. Until the 2-rank probe runs — or a `--split-mode row` boot that does *not* log `internal AllReduce init failed ... falling back to meta-backend butterfly` — "collectives work here" was a strong expectation, not a measurement. **RCCL never got that measurement**; what settled the operational question instead is production traffic — radiance's TP=2 wire carries cross-rank all-reduces over the direct P2P path and prints its own precision, and the measured `wht6` banner with its exact-step bound is on [[frank-archon]]. ## 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 the last one is now moot Every row below was written under the old board. The last row — the only door that was ever open — stopped mattering on 2026-09-20, because the hardware gate it existed to route around no longer applies to this host. | 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), and under 0.4 ms at the Gen5 x8 ~31.5 GB/s each card now has. TP is compute-bound here, not interconnect-bound. That is why the cheap option (the M.2 riser) and the expensive one (the board swap) were both considered adequate on bandwidth alone — and why the board that landed, at x8/x8, clears the arithmetic with room to spare. ## 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. ## Measured verdict (2026-09-20) Both cooperation modes were run on this host on 2026-09-20 — the same salted 10-prompt battery (`scripts/bench-conversation.py`), the service stopped and masked between arms. The arm table is [[frank-medium-radiance-mxfp4]]'s bench of record (arms R = radiance TP=2 + DFlash2 as deployed, I = the llama.cpp layer split that preceded it, L = llama.cpp `--split-mode tensor` against RCCL); this page keeps only the verdict. **TP=2 won on every axis — roughly 2.1× decode and 3.6× at 4-way concurrency over the layer split — and replaced the layer split the day of measurement.** Arm L carries a caveat no flag can remove: llama.cpp's tensor mode could not load the DFlash2 drafter (it aborts at `ggml-backend.cpp:941` — a draft weight pre-allocated on the non-primary device has no backend once `n_copies == 1`), so it decoded **without** speculation while the other arms decoded with it; and part of R's margin over I is the 4-bit MXFP4 checkpoint, not the parallelism. Two numbers elsewhere on this page are therefore not interchangeable: - The 272 t/s recorded 2026-09-16 is **prefill**, on a **butterfly** allreduce, with RCCL off, on the old board. It was never a TP decode number, and its paired decode figure is in "The third path" below. - The butterfly tensor arm and the RCCL tensor arm a day later are the same shape family on different collective paths; the measured pair lives with [[frank-medium-radiance-mxfp4]]'s arms. RCCL needs PCIe atomics on *every* participating GPU — the gate the old board failed. vLLM's R4D one-shot allreduce dispatches no collective at all, which is why TP=2 runs today rather than merely "being unblocked": the image ships TheRock RCCL, and TheRock carries no gfx1201 device code (`invalid device function`), so radiance could not have used RCCL on this host even with atomics routed. Ubuntu's RCCL 2.27 does carry gfx1201 kernels, and that is what arm L linked against. **Rule for any future multi-GPU collective experiment: stop and mask the production service first**, per the benchmarking policy, or run it on a host that is not serving users. The production service is `frank-archon`, a **root** unit, so on Dankserver that is `sudo systemctl stop frank-archon && sudo systemctl mask frank-archon` (`scripts/frank-archon.service`; never `disable` — that deletes the unit file, and `mask` is what the bench policy prescribes). The 2026-09-20 wording below — `systemctl --user stop frank-medium && systemctl --user mask frank-medium`, frank-medium being a rootless quadlet — is retired with that service, which today reports `not-found` (read 2026-10-18). This is not theoretical: 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 (production in the layer-split era): 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` ran from 2026-09-16 to 09-20 ([[frank-medium-llamacpp-dflash2]]): one process, one model, both cards, 262144 context. Its measured layer-split rates are that page's; this host's tensor arm under the butterfly fallback measured 272 t/s prefill / 26.3 tok/s decode. **That tensor number is not a prediction for an RCCL-backed build** — it is what tensor split costs when the allreduce degrades to a meta-backend butterfly, and it is the number the board-swap path exists to retire. The ranking therefore changed twice. On 2026-09-16: software already gets both cards onto one model, so the hardware fix stopped being a capability question and became a performance one — *tensor* parallelism, ~40 % decode, plus one pooled KV instead of two independent 122k-token pools. On 2026-09-20 the hardware half of that question was answered by the board swap; the llama.cpp measurement half never ran, and `GGML_HIP_RCCL` still defaults `OFF` in `scripts/build-llama-rdna4.sh:41` (re-read 2026-10-18) for exactly the reason this page documents — moot while no llama.cpp shape is deployed on this host. ## What would change the answer - ~~An RCCL build whose kernels stop declaring a hostcall buffer~~ — **moot since 2026-09-20.** That door only mattered while the hardware gate could not be passed; the board swap passes it. The `-mprintf-kind=buffered` hypothesis and its counter-evidence stay in the escape table for the record. - The symmetric-kernel path inside RCCL, if a release ever ships the device API (`ncclDevComm*`) that torch's symmetric-memory path needs. The shipped TheRock nightly (`~/.local/share/rocm-7.14.0a20260624/lib/librccl.so.1`) still does not. - ~~Firmware or an AGESA update that makes the chipset bridges advertise atomic routing~~ — arrived as new silicon instead. - 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