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

type, title, description, tags, timestamp

type	title	description	tags	timestamp
research	Dual R9700 Cooperation	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.	frank-medium rocm rccl pcie-atomics tensor-parallelism dankserver	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 and the 2026-09-16 table 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 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

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) 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 — two-machine research (the PP variant of this problem)
• Dankserver System Specs — hardware inventory
• Observability — the scrape/dashboard/alert stack around these endpoints