agent: repoint roles to frank-medium/dankclaw-3090; sync niri tweaks, memories, managed skills
- config.yml: default/commit -> frank-medium (dankserver); task/tiny/smol/vision -> frank-small (dankclaw-3090); vision xhigh - models.yml: add llamacpp (dankclaw-3090) provider; frank-large preserve_thinking + context notes - niri: focus-ring/border/shadow layout, hot-corners off, re-add noctalia.kdl include; blankpad eDP-1 scale 1.0 - memories: flush rollouts; retire frank_large + marreta_authentik skills; add auth-chain/kv-pool/quadlet-log skills - managed-skills: 40+ new fleet runbook skills - agent/skills: import built-in skill set
This commit is contained in:
Frank committed
1 parent
83c9c92b3f
commit
135bb2194d
103 files changed
+5333
-170
No files matched your search
@@ -0,0 +1,122 @@
|
||||
---
|
||||
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.
|
||||
Reference in new issue
Block a user