---
name: pcie-link-capability-probe
description: "Measure a GPU-to-host PCIe link the way a weight-streaming engine experiences it, and convert the result into a tokens/s ceiling (streamed expert weights) versus the compute-in-RAM alternative — before recommending or rejecting an offload design."
---

# PCIe link capability probe

Use when someone asks "is the link fast enough to run model X with engine Y (on one GPU)?"
and Y keeps some weights in host RAM — MoE expert offload with an LRU GPU cache, layer
paging, weight streaming. The answer is arithmetic once the link is measured, and the
*shape* of the engine 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 on the GPU | nothing | no |

## 1. Establish the hardware ceiling (never trust the card alone)

```bash
lspci -tv | grep -iA2 vga                       # which bridge the GPU hangs off
d=$(lspci -D | grep -i nvidia | awk '{print $1}' | head -1)
sudo lspci -vv -s "$d" | grep -E 'LnkCap:|LnkSta:'    # the CARD's capability + trained state
sudo lspci -vv -s 00:03.1 | grep -E 'LnkCap:|LnkSta:'  # the PLATFORM PORT (this is the real limit)
```

- `LnkCap` of the **upstream port** = the electrical ceiling. A card with `Width x16` under
  an `x8` port cannot exceed x8; an x4 riser/cable caps both.
- `LnkSta` at idle usually shows a downgraded speed (PCIe ASPM downshift, e.g.
  `2.5GT/s (downgraded)`) — that is NOT the loaded rate. Trust the measurement in step 2.
- Gen4 usable payload: x4 **7.877 GB/s**, x8 15.75, x16 31.5 (128b/130b encoding).

## 2. Measure it (the program that did the job)

Stop **all** production services on the host first (fleet policy), then compile and run:

```cuda
// pcie-link-test.cu : nvcc -O2 -o /tmp/pcie-link-test /tmp/pcie-link-test.cu
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cuda_runtime.h>
#define CK(x) do{ cudaError_t e=(x); if(e!=cudaSuccess){fprintf(stderr,"CUDA %s (%s:%d)\n",cudaGetErrorString(e),__FILE__,__LINE__); return 1;} }while(0)
static double ms_between(cudaEvent_t a, cudaEvent_t b){float ms=0; cudaEventElapsedTime(&ms,a,b); return (double)ms;}
static int bench_size(size_t bytes){
    const int iters = bytes>=(1UL<<30)?10:(bytes>=(256UL<<20)?20:100);
    void *dev=NULL,*host=NULL,*pageable=NULL; cudaStream_t s=NULL;
    CK(cudaMalloc(&dev,bytes)); CK(cudaHostAlloc(&host,bytes,cudaHostAllocDefault));
    memset(host,0x5a,bytes); CK(cudaStreamCreate(&s));
    cudaEvent_t a,b; CK(cudaEventCreate(&a)); CK(cudaEventCreate(&b));
    for(int i=0;i<3;i++) CK(cudaMemcpy(dev,host,bytes,cudaMemcpyHostToDevice));
    CK(cudaDeviceSynchronize());
    CK(cudaEventRecord(a)); for(int i=0;i<iters;i++) CK(cudaMemcpy(dev,host,bytes,cudaMemcpyHostToDevice));
    CK(cudaEventRecord(b)); CK(cudaEventSynchronize(b)); double h2d=(double)bytes*iters/(ms_between(a,b)/1000.0)/1e9;
    CK(cudaEventRecord(a)); for(int i=0;i<iters;i++) CK(cudaMemcpyAsync(dev,host,bytes,cudaMemcpyHostToDevice,s));
    CK(cudaEventRecord(b)); CK(cudaEventSynchronize(b)); double h2da=(double)bytes*iters/(ms_between(a,b)/1000.0)/1e9;
    CK(cudaEventRecord(a)); for(int i=0;i<iters;i++) CK(cudaMemcpy(host,dev,bytes,cudaMemcpyDeviceToHost));
    CK(cudaEventRecord(b)); CK(cudaEventSynchronize(b)); double d2h=(double)bytes*iters/(ms_between(a,b)/1000.0)/1e9;
    CK(cudaMemcpy(dev,host,bytes,cudaMemcpyHostToDevice));
    CK(cudaEventRecord(a)); CK(cudaMemcpy(dev,host,bytes,cudaMemcpyHostToDevice)); CK(cudaEventRecord(b)); CK(cudaEventSynchronize(b));
    double one_us=ms_between(a,b)*1000.0;
    pageable=malloc(bytes); memset(pageable,0x5a,bytes);
    for(int i=0;i<2;i++) CK(cudaMemcpy(dev,pageable,bytes,cudaMemcpyHostToDevice));
    CK(cudaEventRecord(a)); for(int i=0;i<iters;i++) CK(cudaMemcpy(dev,pageable,bytes,cudaMemcpyHostToDevice));
    CK(cudaEventRecord(b)); CK(cudaEventSynchronize(b)); double pag=(double)bytes*iters/(ms_between(a,b)/1000.0)/1e9;
    printf("%8.1f MiB | %6.2f | %6.2f | %6.2f | %8.1f us | %6.2f\n", bytes/1048576.0,h2d,h2da,d2h,one_us,pag);
    free(pageable); cudaFreeHost(host); cudaFree(dev); cudaEventDestroy(a); cudaEventDestroy(b); cudaStreamDestroy(s); return 0;
}
int main(void){ cudaDeviceProp p; CK(cudaGetDeviceProperties(&p,0));
    printf("device %s | CUDA %d.%d\n", p.name, p.major, p.minor);
    printf("   MiB  | H2D    | async  | D2H    | 1x xfer    | pageable\n");
    size_t sz[]={3UL<<20,16UL<<20,64UL<<20,256UL<<20,1UL<<30,2UL<<30};
    for(int i=0;i<6;i++) if(bench_size(sz[i])) return 1; return 0; }
```

Size choice matters: **3 MiB** is the granularity an expert-streaming engine moves, so
its single-transfer time is the number that matters; the big sizes give the ceiling.

## 3. Read it

- **Achieved GB/s is itself the proof of the trained link state** — 6.7 GB/s is only
  possible at Gen4 x4 (Gen1 x4 ≈ 1, Gen2 ≈ 2, Gen3 ≈ 3.94), so no separate `LnkSta`
  sample under load is needed.
- **async ≈ sync and pageable ≈ pinned ⇒ link-limited, not staging-limited.** If async is
  much faster than sync, the engine's fill pipeline (batching copies on one stream) is
  where the win is; if pageable is much slower, pinning is a real lever. Measure, don't
  assume.
- Per-slab launch overhead = single-xfer time − (slab bytes ÷ rate). Small next to
  bandwidth cost at 3 MiB.

## 4. Convert to a tokens/s ceiling

```
active expert bytes/token  = active_params × bytes/param   (A3B @ Q4_K_M ≈ 1.70 GiB)
link time/token            = active expert bytes / measured H2D rate
ceiling (tok/s)            = 1 / link time/token      × resident-fraction saved (LRU hits)
```

Then compare against the alternatives on the same host — and against the
compute-in-RAM path computed at host RAM bandwidth (256 GB/s LPDDR5X on this fleet's
Strix Halo). If the link is >10× slower than the memory behind it (x4 is 31× here),
streaming loses by construction and no tuning rescues it; doubling the link (x4 → x8)
only doubles a losing number.

## Worked example (Dankclaw 2026-09-16)

Gen4 x4, measured H2D 6.72 / D2H 6.60 GB/s (85 % of spec), 3 MiB slab = 474 µs.
A3B MoE, Q4_K_M (75 GiB experts, ~1.70 GiB active/token): 272 ms/token ⇒ **3.7 tok/s**
(4.6 with 20 % resident). Same host, same model, measured: **20.46 tok/s** with the
experts computed on the CPU, 25 (halogen), 26.5 (mainline on the iGPU). Verdict: the
streaming design loses 3-5×, and the engine in question could not load the GGUF anyway.

## Pitfalls

- Stop production services before running this (fleet policy) — and remember the 3090's
  VRAM may be almost fully occupied, so a 2 GiB test buffer needs the service stopped.
- Don't quote a derived line from inside the program without checking it: a printf
  format/argument mismatch in this very probe printed a nonsense ceiling. Recompute by hand.
- GiB vs GB: convert before dividing (1.70 GiB = 1.825 GB).
- The idle `LnkSta` (2.5 GT/s, downgraded) is not the loaded link — quoting it makes a
  working x4 look like a broken x1.