Frank
a5fe045b58
Structural repair: - 13 pages were the same document twice (each exactly 2x its HEAD size), a second frontmatter block whose closing fence had been glued onto a body line by a frontmatter rewrite that re-read its own output. Un-doubled by cutting at the second frontmatter and keeping the copy whose body is a superset; the orphan census reports 0 lost tokens against HEAD. - log.md: one entry per date, bare `## YYYY-MM-DD` headings, newest first — 66 dated headings merged into 34, titles demoted to bold leads, every content line verified present. The file had been descending at the head and ascending at the tail. Ownership moves: - New `LLM/frank-medium.md` is the locator for the retired Dankserver endpoint: era windows quoted from each era page's own banner, no numbers restated, and the two unresolved date conflicts recorded instead of silently resolved. `frank-archon.md` links to it instead of restating the pre-rename window. - The four Pipeline-Parallelism stubs (frontmatter only, no body) are gone; `Multi-Node Runtimes.md` owns runtime construction, launch, mixed-vendor verdicts and cross-split spec decode, and the four links that pointed at the stubs were repointed at it or at the drafter page. - Authentik's three stub pages fold into `Authentik.md`, with the process detail and the outpost reference split out so the parent stays under the length hint. Verified corrections (host census + live probe, 2026-10-19): - The MXFP4 checkpoint two pages called the "on-disk rollback" is gone: /home/dank/models/ holds StillDeadcode, rad and tpl.jinja. Stated once, on the era page that owns that payload; the swap page now links to it. - frank-medium-radiance-mxfp4: the 4.24/7 accept rate belongs to the DFlash2 arm measured on the FP8 target, not to the MXFP4 era; the 3-family soak run census is kept once (155 requests) instead of twice (~120 and 155). - Fleet shapes re-confirmed against the live endpoints before touching any page: 3x200000 fp8 KV, 2x200704, 2x262144. Index rows are regenerated, not hand-carved. Net -1,849 vault lines.
255 lines
13 KiB
Markdown
255 lines
13 KiB
Markdown
---
|
||
type: reference
|
||
title: PCIe Link and Atomics
|
||
description: "Link probing, the streaming-versus-compute ceiling arithmetic, and the atomics wall that decides multi-GPU collectives."
|
||
tags: [pcie, rccl, atomics, multi-gpu, bandwidth, amd, nvidia]
|
||
timestamp: 2026-10-05T00:00:00-07:00
|
||
---
|
||
|
||
# PCIe Link and Atomics
|
||
|
||
## 1. Degraded or just idle?
|
||
|
||
The trap: an idle GPU reports `LnkSta: Speed 2.5GT/s (downgraded)` — the
|
||
driver downshifts to Gen1 for power. That is not a fault and not the slot's
|
||
fault. Never conclude from an idle reading.
|
||
|
||
1. **Path ceiling from the kernel (cheapest, most decisive):**
|
||
`sudo journalctl -k -b --no-pager | grep "available PCIe bandwidth"` →
|
||
`pci 0000:c4:00.0: 63.012 Gb/s available PCIe bandwidth, limited by 16.0
|
||
GT/s PCIe x4 link at 0000:00:03.1 (capable of 252.048 Gb/s with ... x16)`.
|
||
The "limited by" device names the **bottleneck** (usually a root port) and
|
||
its width — that is the number a swap could move; "capable of x16" is the
|
||
endpoint's OWN capability, and the gap between them is the story.
|
||
`dmesg` rotates; `journalctl -k -b` keeps this boot.
|
||
2. **Both sides separately:** `LnkCap`/`LnkSta`/`LnkCtl2` on endpoint AND root
|
||
port. `LnkCtl2: Target Link Speed: 16GT/s` on both ends = the platform
|
||
intends Gen4, so a Gen1 current speed is power management, not
|
||
misconfiguration. A root port whose `LnkCap` width exceeds the trained
|
||
width means the socket/adapter wires fewer lanes — **M.2 connectors are 4
|
||
lanes by pinout; an M.2 riser can never deliver more.**
|
||
3. **Prove the link under load (the step people skip):** on NVIDIA sample NVML
|
||
while the GPU works
|
||
(`nvidia-smi --query-gpu=pcie.link.gen.current,pcie.link.gen.hostmax,pcie.link.width.current,pcie.link.width.max`):
|
||
`gen.current` climbing 1 → `hostmax` = idle downshift, healthy; stuck below
|
||
`hostmax` under load = real degradation; `width.current < width.max` under
|
||
load = **physical lane limit** — no software touches it. On a busy
|
||
production GPU this is a usage smoke test, not a benchmark — never stop a
|
||
live service to read its link.
|
||
4. **Do not plan hardware from `dmidecode -t slot`** — consumer firmware
|
||
labels Ethernet/SD controllers as "M.2 slots" and advertises slots in PCI
|
||
domains that do not exist. Trust `lspci -nn | grep "PCI bridge"` and
|
||
`lspci -tv` (empty root ports still enumerate; an absent port is absent).
|
||
5. **"Would moving the card help?"** Compare the bottleneck device from step 1
|
||
with the target port's `LnkCap` width; two M.2 sockets = both x4 = a
|
||
no-op. Clean `DevSta`/`CESta` with a low trained speed means "wired
|
||
narrow", not "failing". Report ceiling, card capability, what the idle
|
||
reading meant, and whether any physical change moves the number.
|
||
|
||
Gen4 usable payload after 128b/130b: **x4 7.877 GB/s, x8 15.75, x16 31.5**.
|
||
|
||
## 2. The engine shape decides whether the link matters at all
|
||
|
||
| engine shape | what crosses PCIe per token | link matters? |
|
||
|---|---|---|
|
||
| streams weights (expert LRU, paging) | the missed weights (~GB/token) | **yes — decisive** |
|
||
| computes offloaded weights in RAM (llama.cpp `-ot exps=CPU`) | activations only (~KB) | no; host RAM bandwidth decides |
|
||
| everything resident | nothing | no |
|
||
|
||
**Ceiling arithmetic:** `link time/token = active expert bytes / measured
|
||
H2D rate`; `ceiling = 1 / link_time × resident-fraction saved (LRU hits)`.
|
||
Compare against the compute-in-RAM path at host RAM bandwidth (256 GB/s
|
||
LPDDR5X on the Strix Halo — 128 GB LPDDR5x-8000 on a 256-bit bus,
|
||
[amd.com spec page](https://www.amd.com/en/products/processors/laptop/ryzen/ai-300-series/amd-ryzen-ai-max-plus-395.html)).
|
||
**If the link is >10x slower than the memory behind it (x4 is 31x here),
|
||
streaming loses by construction and no tuning rescues it; doubling the link
|
||
only doubles a losing number.**
|
||
|
||
Worked example: the Dankclaw Gen4 x4 run of 2026-09-16 — measured H2D/D2H
|
||
rates, the 3 MiB slab cost, and the streaming-vs-compute comparison it decided
|
||
— is recorded on [[x4-pair-experiment]] §The link, measured. Read it once and
|
||
apply the rule above: on that link the streaming design lost by 3–5x, and the
|
||
x8 ceiling would only have halved the loss.
|
||
|
||
## 3. Measure the link: two probes, same mechanics
|
||
|
||
Stop **all** production services first (fleet bench law) — the 3090's VRAM
|
||
may be nearly full, and the probe contends for the very link it measures.
|
||
|
||
**CUDA probe** (`pcie-link-test.cu`): sweep {3, 16, 64, 256 MiB, 1, 2 GiB}
|
||
with `cudaMalloc` + `cudaHostAlloc`, event-timed; per size measure sync H2D,
|
||
async H2D, D2H, one single transfer, and pageable H2D. `3 MiB` is the row
|
||
that matters for expert streaming; big sizes give the ceiling.
|
||
|
||
**HIP probe for RDNA4** (`pcie_bw.hip`) — the ROCm landmines each cost real
|
||
time, so encode them:
|
||
|
||
- `hipGetDeviceProperties` SEGFAULTS on some builds; use `hipGetDeviceCount`
|
||
+ `hipSetDevice` only. First line: `setvbuf(stdout, NULL, _IONBF, 0);` — a
|
||
segfaulting process loses block-buffered stdout and hides the crash point.
|
||
- Build with an explicit arch: `HIP_TARGETS=gfx1201 hipcc -O2
|
||
--offload-arch=gfx1201 -o p p.hip`; compile with `| grep -E "error:"`,
|
||
never `| tail -N` (truncation hid a missing `#include <cstring>`).
|
||
- `hipMemset`/`hipDeviceSynchronize` are `[nodiscard]` — use or cast.
|
||
- Mount the whole `/dev/dri` into ROCm containers; individual `renderD*`
|
||
nodes enumerate but fail context creation.
|
||
- `rocm-smi --showmetrics`: `pcie_link_width` is trustworthy, but
|
||
`pcie_link_speed` reads the ASPM idle value and `pcie_bandwidth_acc/inst`
|
||
are usually `N/A` on consumer cards — live traffic generally cannot be
|
||
sampled.
|
||
|
||
```cuda
|
||
// pcie_bw.hip — sweep sizes, pinned vs pageable
|
||
#include <cstdio>
|
||
#include <cstring>
|
||
#include <cstdlib>
|
||
#include <hip/hip_runtime.h>
|
||
static double one(size_t bytes, bool pinned, int iters) {
|
||
void *dev = NULL, *host = NULL;
|
||
if (hipMalloc(&dev, bytes) != hipSuccess) return -1;
|
||
if (pinned) { if (hipHostMalloc(&host, bytes, hipHostMallocDefault) != hipSuccess) return -1; }
|
||
else { host = malloc(bytes); }
|
||
memset(host, 0x5a, bytes);
|
||
hipEvent_t a, b; hipEventCreate(&a); hipEventCreate(&b);
|
||
for (int i = 0; i < 3; i++) hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice);
|
||
hipDeviceSynchronize();
|
||
hipEventRecord(a, 0);
|
||
for (int i = 0; i < iters; i++) hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice);
|
||
hipEventRecord(b, 0); hipEventSynchronize(b);
|
||
float ms = 0; hipEventElapsedTime(&ms, a, b);
|
||
double gb = (double)bytes * iters / (ms / 1000.0) / 1e9;
|
||
if (!pinned) free(host); else hipHostFree(host);
|
||
hipFree(dev); hipEventDestroy(a); hipEventDestroy(b);
|
||
return gb;
|
||
}
|
||
int main(int argc, char **argv) {
|
||
setvbuf(stdout, NULL, _IONBF, 0);
|
||
if (hipSetDevice((argc > 1) ? atoi(argv[1]) : 0) != hipSuccess) return 1;
|
||
size_t sz[] = {2459136UL, 3UL<<20, 16UL<<20, 64UL<<20, 256UL<<20, 1UL<<30};
|
||
for (int i = 0; i < 6; i++) {
|
||
size_t b = sz[i];
|
||
int iters = (b >= (1UL<<30)) ? 10 : (b >= (256UL<<20)) ? 20 : 100;
|
||
printf(" %10.3f MiB | pinned %6.2f | pageable %6.2f\n",
|
||
b/1048576.0, one(b,true,iters), one(b,false,iters));
|
||
}
|
||
return 0;
|
||
}
|
||
```
|
||
|
||
**Reading it:**
|
||
|
||
- **Engine-vs-link is the diagnosis, not the raw number:** compare
|
||
`engine_effective_GB_s` to `link_GB_s` at the SAME transfer size. Engine at
|
||
~40 % of the link at the same granularity = the wire is not the problem and
|
||
no hardware change will help.
|
||
- Achieved GB/s is itself the proof of the trained link state (6.7 GB/s only
|
||
exists at Gen4 x4) — no separate under-load `LnkSta` sample needed.
|
||
- **pinned ≈ pageable ⇒ pinning is not a lever.** async ≈ sync ⇒ link-limited
|
||
(the usual "pageable costs 2x" rule does not apply at this link).
|
||
- Small-transfer penalty (2.345 MiB: ~106 µs vs ~78 µs ideal, ~21 %) is
|
||
per-transfer overhead — recoverable only by BATCHING transfers, not by a
|
||
faster link. The engine-vs-wire version of that reading on Dankserver
|
||
(bulk slab rate against the engine's effective gather rate, same
|
||
granularity) is [[frank-archon]]'s tcclaviger-era record.
|
||
- Both GPUs symmetric ⇒ no cross-device rebalance is available; measure both.
|
||
- Don't quote a derived line from inside a probe without hand-checking it —
|
||
a printf format/argument mismatch in one of these printed a nonsense
|
||
ceiling. GiB→GB before dividing.
|
||
|
||
## 4. The PCIe-atomics wall (AMD multi-GPU collectives)
|
||
|
||
**Symptom:** multi-GPU AMD host fails TP/PP; the failure surfaces as
|
||
`hipLaunchKernel: hipErrorIllegalState` for `ncclDevKernel_*` — nothing in
|
||
vLLM/RCCL/PyTorch names atomics. Get the real message:
|
||
|
||
```bash
|
||
AMD_LOG_LEVEL=3 <the failing process>
|
||
# → "Pcie atomics not enabled, hostcall not supported" → "AQL dispatch failed!"
|
||
```
|
||
|
||
**Cause chain:** RCCL kernels carry a hidden hostcall-buffer parameter; ROCr
|
||
refuses to fill it without PCIe atomics, which it reads **once at driver
|
||
probe** (`rocclr/device/rocm/rocdevice.cpp` `checkAtomicSupport()` — no else,
|
||
no override). **No env var, module param, or sysfs toggle exists** —
|
||
`setpci -s <gpu> CAP_EXP+28.w=0440` (DCTL2 bit 6) changes nothing at runtime.
|
||
|
||
**Check the CHAIN, not the card:**
|
||
|
||
```bash
|
||
sudo lspci -vv -s <gpu> | grep -E "AtomicOpsCtl|LnkSta|LnkCap" # want ReqEn+
|
||
lspci -tv # every hop
|
||
sudo lspci -vv -s <bridge> | grep AtomicOpsCap # want Routing+
|
||
```
|
||
|
||
`Routing-` is a read-only hardware capability. Typical AM5 layout: one x16
|
||
slot on CPU root ports (atomics fine), the rest behind the chipset (usually
|
||
not). Bandwidth for TP is small anyway: ~10 MB per token-step for a 27B at
|
||
batch 4 is a millisecond-scale figure at x4 — compute-bound, not
|
||
interconnect-bound (full derivation: [[Dual-R9700-Cooperation]] §Bandwidth).
|
||
|
||
**Probe matrix (bare 2-rank `torch.distributed`, one process per GPU, 25 s PG
|
||
timeout) — one probe answers several design questions:**
|
||
|
||
| Op | CPU-lane card | chipset card | Decides |
|
||
|---|---|---|---|
|
||
| `broadcast`/`all_reduce`/`barrier` | OK | FAIL | TP is dead |
|
||
| `send`/`recv` (P2P), in ISOLATION | hang | FAIL | **PP is dead too** — vLLM's PP handoff IS this primitive |
|
||
|
||
Test `send`/`recv` alone (a prior failed op leaves the group broken). Run the
|
||
probe from a REAL file (`torch.multiprocessing.spawn` re-executes it; piping
|
||
into `python3 -` fails with `FileNotFoundError: '/<stdin>'`). `world_size=1`
|
||
always succeeds — **one-card success proves nothing.**
|
||
|
||
**SAFETY — non-negotiable:** never run a GPU experiment beside a serving
|
||
unit: `sudo systemctl stop <svc> && sudo systemctl mask <svc>` first (mask,
|
||
never disable — disable deletes a relinked unit file), or use a host that
|
||
serves nobody. A hung AQL dispatch can wedge the GPU driver; a probe beside a
|
||
live service on 2026-09-16 preceded a hard host down (no ARP, no ICMP from
|
||
two subnets, power cycle required). Telling off from wedged afterwards: `ip
|
||
neigh show <ip>` from a peer — `FAILED`/`INCOMPLETE` = off/NIC dead; a wedged
|
||
kernel usually still answers ARP. Validate your own probe's naming before
|
||
reporting an outage (`gitea` vs `gitea-server`).
|
||
|
||
**Fixes, ranked:**
|
||
|
||
1. CPU-attached M.2 → PCIe x4 riser (~$30) — puts the card on the CPU root
|
||
complex; x4 suffices for TP.
|
||
2. Board whose CPU x16 bifurcates x8/x8 (done 2026-09-20 on dankserver — the
|
||
per-device chain census, including both endpoints reading `ReqEn+`, is
|
||
[[Dual-R9700-Cooperation]]'s). Re-checked 2026-10-18 without root: both root
|
||
ports report `current_link_width=8` at `32.0 GT/s`
|
||
(`/sys/bus/pci/devices/0000:00:01.{1,3}/current_link_width`), and the live
|
||
engine states direct P2P between the cards.
|
||
3. llama.cpp `--split-mode layer` — no collectives at all ([[Llama Cpp Multi-Device Splits]]); costs the GGUF conversion plus paged attention /
|
||
continuous batching / spec decode.
|
||
4. Independent TP=1 replicas + routing proxy — cards never share one model.
|
||
|
||
**Closed software routes (do not re-litigate):** vLLM's IPC custom allreduce
|
||
is gated by a CUDA-sm-keyed table (`{9.0, 10.0, 10.3, 10.7}`; ROCm reports
|
||
12.0 → off on RDNA4 by code limit) — the table is
|
||
`CUSTOM_ALL_REDUCE_MAX_SIZES` in
|
||
`vllm/distributed/device_communicators/all_reduce_utils.py`, whose keys on
|
||
today's `main` are exactly `"9.0"`, `"10.0"`, `"10.3"`, `"10.7"`
|
||
([source](https://raw.githubusercontent.com/vllm-project/vllm/main/vllm/distributed/device_communicators/all_reduce_utils.py),
|
||
read 2026-10-18); nothing keys on a ROCm-reported `12.0`. — and patching it is
|
||
insufficient because TP init's `broadcast`/`barrier`/logits `all_gather` have
|
||
no non-RCCL path. A
|
||
fork wrapping `.all_reduce` (e.g. radiance `RADIANCE_USE_R4D_AR` one-shot
|
||
P2P, gloo for IPC-handle exchange) does not intercept
|
||
`PyNcclCommunicator.__init__`'s warmup allreduce. NCCL symmetric memory needs
|
||
`ncclDevComm*` (many builds ship the version string without the device API).
|
||
P/D disaggregation needs a KV connector AND coordination collectives. gloo is
|
||
CPU/TCP — unreachable for per-layer TP.
|
||
|
||
**Open, unverified:** rebuilding RCCL so kernels stop declaring a hostcall
|
||
buffer (`-mprintf-kind=buffered` is the usual claim) — counter-evidence: no
|
||
live `printf`/`assert` in RCCL device source. Test path: rebuild, dump a
|
||
kernel's declared args (`roc-obj-extract` + `llvm-readelf --notes`), look for
|
||
`hidden_hostcall_buffer`; if gone, retry the 2-rank probe (service stopped).
|
||
|
||
## Related
|
||
|
||
- [[Llama Cpp Multi-Device Splits]] — the layer-split path around the wall
|
||
- [[Dual-R9700-Cooperation]] / [[x4-pair-experiment]] — measured history
|
||
- [[frank-archon]] — the gather economics of the retired tcclaviger vLLM era (2026-09-29 → 10-05)
|