command-center/agent/managed-skills/pcie-link-capability-probe/SKILL.md

name, description

name	description
pcie-link-capability-probe	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)

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:

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