serverAppConfigs/serverAppConfigs-docs/LLM/PCIe Link and Atomics.md

type, title, description, tags, timestamp

type	title	description	tags	timestamp
reference	PCIe Link and Atomics	Link probing (idle downshift, path ceilings, CUDA/HIP bandwidth probes), the streaming-vs-compute ceiling arithmetic, and the PCIe-atomics wall that decides whether AMD multi-GPU collectives can exist at all.	pcie rccl atomics multi-gpu bandwidth amd nvidia	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). 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.

// 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:

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:

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