---
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
(`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). Tensor would still win decode on a
build with a working RCCL — which is exactly what the M.2 riser or a board swap
in [the short answer](#the-short-answer) would buy.

So the ranking changed on 2026-09-16: software already gets both cards onto one
model, and the hardware fix is no longer a means to that end — it is a means to
*tensor* parallelism, which buys ~40 % decode and would let this runtime drop
layer-split's serial stages entirely.

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