serverAppConfigs/serverAppConfigs-docs/LLM/Dual-R9700-Cooperation.md

---
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 serve `frank-medium` as **one vLLM-radiance process at
tensor_parallel_size 2** ([[frank-medium-radiance-mxfp4]]). 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 llama.cpp layer split that
served the endpoint until that swap is [[frank-medium-llamacpp-dflash2]].

## 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.
- **One of them is production.** TP=2 wins on every axis measured, by 2.1x on
  decode and 3.6x at 4-way concurrency, so it replaced the layer split on the day
  of measurement. See [[frank-medium-radiance-mxfp4]] for the deployed shape.
- **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<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.

### 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 <bdf> | 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` still reports `PCIE` between the cards (Navi 48 has no
  XGMI), they are still in separate IOMMU groups (15 and 20), and both root
  ports have ACS `ReqRedir+`, so P2P hairpins at the root complex. 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`. Both were ruled out as causes
  ([above](#ruled-out-empirically-do-not-re-derive)) and both are unchanged.

**Still unproven.** 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 has 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" is a strong expectation, not a measurement.

## 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. Decode/prefill are geometric means over the ten rows;
burst-4 is four concurrent copies of one prompt, reported as aggregate tokens/s.

| Shape | Cross-device path | Decode | TTFT | Prefill | burst-4 |
|---|---|---|---|---|---|
| **radiance TP=2 + DFlash2 — deployed** | R4D one-shot allreduce | **95.8 t/s** | **0.21 s** | **1493.6 t/s** | **259.4 t/s (65.6/stream)** |
| llama.cpp layer split 32:33 + DFlash2 — predecessor | none, boundary activation copy | 45.7 t/s | 0.72 s | ~530 t/s | 71.7 t/s |
| llama.cpp `--split-mode tensor` 1:1 | RCCL allreduce | 30.5 t/s | 0.71 s | n/r | 54.3 t/s |

Row 3 carries a caveat that 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 two rows decoded with it. The
comparison that decides anything is row 1 against row 2, and some of that gap is
the 4-bit MXFP4 checkpoint rather than 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.
- 26.3 t/s from the same table is that butterfly tensor's decode. Row 3's 30.5 is
  an RCCL tensor decode a day later: same shape family, different collective path.

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 row 3 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. Since 2026-09-20 `frank-medium` is a rootless quadlet, so
that is `systemctl --user stop frank-medium && systemctl --user mask frank-medium`
on Dankserver.

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, 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). **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 measurement
half has not run yet, and `GGML_HIP_RCCL` is still `OFF` in
`scripts/build-llama-rdna4.sh` for exactly the reason this page documents.

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