diff --git a/AGENTS.md b/AGENTS.md index b76af0a..8691240 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -200,6 +200,17 @@ Scratch ports are not used on LLM hosts — a test process cannot share the single GPU anyway; validate against the service itself. +This covers **any GPU experiment, not just benchmarks** — including a bare +2-rank RCCL/`torch.distributed` probe. On 2026-09-16 one such probe was run +against the live service "because it only needs 200 MB of VRAM"; the host went +hard down ~20 minutes later (no ARP reply, 100% ICMP loss from two subnets, no +recovery — a power cycle was required). A hung AQL dispatch wedging the GPU +driver is the plausible mechanism, and this host has precedent for going +unreachable until a power cycle. If an experiment needs GPUs, stop and mask +first, or use a host that serves nobody. Telling an off machine from a wedged +one: `ip neigh show ` from a peer — FAILED means off/NIC dead, while a +merely wedged kernel usually still answers ARP. + 1. Stop and mask (self-healing would restart it): - Dankclaw `frank-large` (root unit since the 2026-09-15 ik cutover): `sudo systemctl stop frank-large && sudo systemctl mask frank-large` @@ -276,9 +287,10 @@ a device without PCIe atomics; the second R9700 sits behind the chipset bridges `05:00.0`/`06:09.0`, which advertise `AtomicOpsCap: Routing-` — a read-only capability, so no software can enable atomics on that path. Not the kernel version (7.0.0-28/-30/-31 identical), not the image, not the RCCL -build (a from-source rebuild fails identically). A slot on a CPU root port — -there is none free on this AM5 platform for a second x16 card — or a platform -with more CPU lanes restores TP=2. `frank-medium.container` stays in the repo +build (a from-source rebuild fails identically). Putting both cards on CPU root +ports restores TP=2: on this board that means an M.2→PCIe x4 riser into the +free, CPU-attached M.2_2 slot, or a board whose CPU x16 bifurcates to x8/x8. +`frank-medium.container` stays in the repo as the TP=2 target; the deploy handler does not install it. Runbook: `serverAppConfigs-docs/LLM/frank-medium-radiance-mxfp4.md`. @@ -304,6 +316,16 @@ linked runbooks. and the rest behind the chipset (usually not) — that asymmetry is what blocked frank-medium's TP=2. A one-card collective proves nothing: NCCL short-circuits at `world_size=1`. +- **The fix is physical — and cheap on this board.** Two paths put both cards + on CPU root ports, which restores the atomics routing and therefore TP=2: an + M.2→PCIe x4 riser into a *CPU-attached* M.2 slot (on dankserver's TUF + GAMING B650-E WIFI that is M.2_2, which is free — per ASUS's spec page; + `PCIEX16_2`, where card #2 sits today, is chipset and x1-mode), or a board + whose CPU x16 bifurcates to x8/x8. Bandwidth is not the constraint: a 27B + TP=2 allreduce at batch 4 is ≈10 MB per token-step, ≈1.4 ms at Gen4 x4. + RCCL's P2P `send`/`recv` fails on such a card too, so pipeline parallelism is + not an escape either. Full evidence matrix, the open rebuild hypothesis, and + the test plan: [[Dual-R9700-Cooperation]]. - **When a GPU collective dies, read the runtime, not the framework.** The message that identified the cause came from `AMD_LOG_LEVEL=3` (`Pcie atomics not enabled, hostcall not supported` → `AQL dispatch diff --git a/serverAppConfigs-docs/LLM/Dual-R9700-Cooperation.md b/serverAppConfigs-docs/LLM/Dual-R9700-Cooperation.md new file mode 100644 index 0000000..5af528d --- /dev/null +++ b/serverAppConfigs-docs/LLM/Dual-R9700-Cooperation.md @@ -0,0 +1,158 @@ +--- +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 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 | 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 used for the replicas +(`scripts/bench-vllm-corpus.py` / `bench-turn-latency.py`) 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. + +## 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 diff --git a/serverAppConfigs-docs/index.md b/serverAppConfigs-docs/index.md index 6767800..640c8d5 100644 --- a/serverAppConfigs-docs/index.md +++ b/serverAppConfigs-docs/index.md @@ -120,6 +120,12 @@ The LLM inference fleet on the two hosts. Full index: AI PRO R9700 as two single-GPU replicas behind a round-robin proxy, plus the PCIe-atomics evidence that rules out TP=2 on this host (since 2026-09-15/16) +- [[Dual-R9700-Cooperation]] — why the two R9700s cannot run tensor-parallel + together: the PCIe-atomics gate quoted from ROCclr, the measured RCCL matrix + (collectives *and* P2P `send`/`recv` both fail on card #2, which rules out + pipeline parallelism too), what is closed and what is still open, and the two + hardware paths — a ~$30 M.2→PCIe x4 riser into the free, CPU-attached M.2_2 + slot, or a board whose CPU x16 bifurcates to x8/x8 - [[frank-medium-rocmfpx]] — PRIOR frank-medium deployment (llama.cpp era): charlie12345/ROCmFPX serving kingjones777/ Qwen3.8-27B-ROCmFPX-Q8_0-AGENT-GGUF; kept as the record of the runtime it