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 llama.cpp process with a layer split (frank-medium-llamacpp-dflash2). 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. See frank-medium-radiance-mxfp4 for the replica-pair deployment that ran under the old constraint.

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 names.

What that buys, and what it does not:

• Unblocked in principle: vLLM TP=2, vLLM PP=2, and llama.cpp --split-mode row. All three died on the same hostcall gate.
• Not yet measured. Nothing has dispatched a collective on this host since the swap. The 2-rank probe is the remaining gate, and policy is stop-and-mask production first.
• Not needed for capability. The layer split already puts both cards on one model. This is a performance question: an RCCL-backed tensor split, and the single pooled KV that comes with it.
• 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.

Test plan (when the hardware changes)

Precondition met 2026-09-20: the line below reads ReqEn+ on both cards, so this is now a standing plan rather than a hypothesis.

The verdict to look for is one line of lspci, before touching any model:

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). 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