docs: vault link integrity restored after rogue-session doc WIP landed
The parallel session's stash pop left GGUF Compatibility, Multi-Device Splits, Serving Sizing, PCIe Link and Atomics, RDNA4 Container Serving, vLLM Porting triage and RDNA4 MoE Budget unindexed and carrying 9 unresolved wikilinks (LLM Bench Windows / Model Swap / AGENTS targets that never existed, a [[nodiscard]] C++ attribute parsed as a wiki link, and two multi-line-wrapped link tokens). All repointed to real pages; six rows added to the vault index; seam green.
This commit is contained in:
Frank committed
1 parent
380aafa1cc
commit
a6ad9af708
10 files changed
+2013
-1
No files matched your search
@@ -0,0 +1,274 @@
|
||||
---
|
||||
type: reference
|
||||
title: Llama Cpp Multi-Device Splits
|
||||
description: What llama.cpp can and cannot express across devices, the no-collective layer-split path, split balancing by measurement, drafter KV sizing from GGUF headers, and heterogeneous placement triage.
|
||||
tags: [llama.cpp, multi-gpu, layer-split, pcie, sizing, speculative]
|
||||
timestamp: 2026-10-05T00:00:00-07:00
|
||||
---
|
||||
|
||||
# Llama Cpp Multi-Device Splits
|
||||
|
||||
## 1. Expressibility law — classify the arrangement before running anything
|
||||
|
||||
llama.cpp expresses exactly two things across devices, and they are not
|
||||
interchangeable:
|
||||
|
||||
| arrangement | how | expressible? |
|
||||
|---|---|---|
|
||||
| **Layer split** — layers 0..n on dev A, rest on dev B | `-sm layer -ts <a>/<b>` | **yes**, no collectives needed; works cross-vendor (CUDA + Vulkan measured 2026-09-16) |
|
||||
| **Per-tensor placement** — attention+KV on A, experts on B | `-ot <regex>=<dev>` | **NO.** `--override-tensor` sets `has_tensor_overrides()`, a REQUIRED term of the scheduler's `pipeline_parallel` flag (only read site: `llama-context.cpp`, in the `pipeline_parallel = …` expression). That flag sets `n_copies = parallel ? MAX : 1` in `ggml_backend_sched_new`; with `n_copies == 1` a pre-allocated weight on a non-primary device has no backend and the graph aborts at `ggml-backend.cpp:941` (`pre-allocated tensor (blk.0.ffn_down_exps.weight) in a buffer (Vulkan1) that cannot run the operation (NONE)` — `(NONE)` is the leaf's own op name). |
|
||||
|
||||
Failure modes observed (mainline `37b53fd`): per-tensor `-ot exps=Vulkan1`
|
||||
with `-dev CUDA0,Vulkan1` → scheduler abort; layer-scoped `-ot` with a
|
||||
cross-backend pair → model load fails with only `failed to load model`;
|
||||
`-ot x=CPU` **works** because the CPU backend is last in the list — **the CPU
|
||||
is the only legitimate partner in a heterogeneous split.** On a CUDA+AMD host
|
||||
a card can be the trunk OR the offload host, never the complement of an
|
||||
AMD-resident peer. Pure Vulkan changes nothing here — the gate is the
|
||||
override, not the driver mix.
|
||||
|
||||
If the plan needs tensors on a specific device: stop, the answer is no for
|
||||
any vendor mix. Expressible alternatives, in order: (1) two independent
|
||||
endpoints, one model per device; (2) single-vendor layer split; (3) device +
|
||||
CPU per-tensor offload (mainline pinning guard: `GGML_CUDA_NO_PINNED`; ik fork
|
||||
was `GGML_CUDA_NO_PINNED_WEIGHTS` — without it a CPU-overridden tensor can
|
||||
land in a pinned `CUDA_Host` buffer and page-lock ~108 GiB).
|
||||
|
||||
Layer-split arithmetic to check before promising: a whole layer carries its
|
||||
experts (MoE) — compute GiB/layer, not per param-group; only some layers carry
|
||||
KV (every 4th on the hybrid), so a share of layers is only a share of KV; big
|
||||
non-repeating tensors (PLE tables) land on the main device and cannot move —
|
||||
they cap how small the primary can get.
|
||||
|
||||
## 2. No collectives? The atomics wall and the layer-split answer
|
||||
|
||||
Decide with one line before designing:
|
||||
|
||||
```bash
|
||||
sudo lspci -vv -s <card2> | grep AtomicOpsCtl # ReqEn- ⇒ no collectives, ever
|
||||
lspci -tv # a Routing- bridge is a hardware wall
|
||||
```
|
||||
|
||||
`ReqEn-` kills vLLM entirely — TP needs a per-layer allreduce AND PP needs a
|
||||
device-backend `send`/`recv` of the hidden state
|
||||
(`vllm/distributed/parallel_state.py`) — i.e. the same RCCL primitive that
|
||||
fails on an atomics-less card. llama.cpp's `--split-mode layer` needs no
|
||||
collective: layers distribute in order, only the boundary activation crosses
|
||||
PCIe once per forward pass. Full probe/repair detail:
|
||||
[[PCIe Link and Atomics]].
|
||||
|
||||
## 3. Build: pin by feature, bundle the runtime at the exact depth
|
||||
|
||||
- Pin the ref by FEATURE list (e.g. DFlash2 + qwen35 MTP on 2026-09-16 →
|
||||
`ggml-org/llama.cpp#27342`); never float. Verify a fork build the same way
|
||||
(`ls <venv>/…/spec_decode/ # dflash/ vs dflash2/`).
|
||||
- `-DCMAKE_HIP_ARCHITECTURES=<gfx only>` — a sibling arch links, loads, then
|
||||
segfaults. `-DGGML_HIP_RCCL=OFF` when there is no atomics path (only tensor
|
||||
mode degrades to the butterfly fallback; layer split unaffected).
|
||||
- **TRAP — `cp -r SRC DST/` flattens when `DST` does not exist** and drops one
|
||||
directory level. rocBLAS/hipBLASLt look for device code at a FIXED depth
|
||||
(`../lib/rocblas/library`, `../lib/hipblaslt/library/<arch>`); flattened,
|
||||
the engine loads every GiB, answers `/health`, and dies on the FIRST request
|
||||
with `Cannot read .../TensileLibrary.dat: No such file or directory for GPU
|
||||
arch`. `mkdir -p DST && cp -r SRC/. DST/`, then
|
||||
`patchelf --set-rpath '$ORIGIN/../lib' build/bin/llama-*`.
|
||||
- Dual-backend build (CUDA + Vulkan is the practical pair; HIP+CUDA do not
|
||||
coexist): `-DGGML_CUDA=ON -DGGML_VULKAN=ON` in one tree; Vulkan needs
|
||||
`glslc` + libvulkan headers. **Device names are not what you expect:**
|
||||
`--list-devices` lists the discrete card under BOTH backends (`CUDA0` AND
|
||||
`Vulkan0`) and the APU LAST (`Vulkan1`) — a wrong `-ot` name makes
|
||||
llama.cpp print the valid set and exit (self-diagnosing).
|
||||
- What the existing build can do: `GPU_BACKEND_*=OFF` lines in its
|
||||
CMakeCache.
|
||||
|
||||
## 4. Serve flags that matter on a split
|
||||
|
||||
```bash
|
||||
HIP_VISIBLE_DEVICES=0,1 llama-server \
|
||||
--device ROCm0,ROCm1 --split-mode layer --tensor-split 1,1 \
|
||||
--gpu-layers 999 --fit off --flash-attn on \
|
||||
--ctx-size <n> --parallel 1 --cache-type-k f16 --cache-type-v f16 \
|
||||
--jinja --metrics --cache-prompt \
|
||||
--spec-type draft-dflash --spec-draft-model <dflash2.gguf> \
|
||||
--spec-draft-device ROCm1 --spec-draft-n-max 7
|
||||
```
|
||||
|
||||
- `--fit off` + explicit `--ctx-size`: auto-fit must not resize a context set
|
||||
on purpose. **`--parallel N` DIVIDES `--ctx-size`** — one slot at large ctx
|
||||
is the concurrency policy, concurrent agent loops queue.
|
||||
- Pin the drafter to the card holding the target's LAST layers
|
||||
(`--spec-draft-device`); under tensor mode an unpinned draft joins the
|
||||
meta-device KV allocation and aborts the KV build. `--spec-draft-n-max` must
|
||||
be ≤ the drafter's `block_size - 1` (DFlash2: 8 ⇒ 7).
|
||||
- Qwen3.8-family templates RAISE on `reasoning_effort=high` (500) — probe
|
||||
low/medium/high/xhigh and map harness levels before declaring ready.
|
||||
|
||||
## 5. Split balancing: measure, never interpolate
|
||||
|
||||
**A `-ts` step is ONE LAYER, and hybrid stacks do not have uniform layers.**
|
||||
qwen35 is 16 × (3 Gated DeltaNet + 1 Gated Attention): one step ≈ 0.43 GiB on
|
||||
a DeltaNet boundary, ≈ 1.43 GiB on a full-attention one (weights + that
|
||||
layer's own f16 KV) — measured on frank-medium. The gap can jump 2.9 GiB in a
|
||||
single step; the best integer split can sit 0.2 GiB from balance while its
|
||||
neighbours are 1+ GiB off. An affine fit reproduces the deployed point and
|
||||
overstates gaps in between (`tests/frank_medium_unit.bats` encodes the fit as
|
||||
a gate; treat a NEGATIVE fitted constant as a unit bug).
|
||||
|
||||
Knobs that move weight and REQUIRE a re-measure: `--ubatch-size` (compute
|
||||
buffers are **~2x heavier on the output/drafter card** — 512→1536 added
|
||||
~3.3 GiB there vs ~1.7 GiB elsewhere), a drafter swap (weights + draft KV),
|
||||
`--ctx-size`/`--parallel` (KV). Total free VRAM is invariant across splits —
|
||||
only the distribution moves; use that to sanity-check a suspicious row.
|
||||
|
||||
**Procedure (one split per load):**
|
||||
|
||||
```bash
|
||||
WORK=/tmp/split.env; cp scripts/serve-<svc>.env $WORK
|
||||
export FRANK_MEDIUM_ENV_FILE=$WORK
|
||||
for ts in 39,26 38,27 37,28 36,29 35,30; do
|
||||
sed -i "s|^TENSOR_SPLIT=.*|TENSOR_SPLIT=$ts|" $WORK
|
||||
nohup bash scripts/serve-<svc>.sh >/tmp/split.log 2>&1 &
|
||||
# poll /health (or process exit), then:
|
||||
rocm-smi --showmeminfo vram | grep 'VRAM Total Used' # sample ~6 s after health
|
||||
pkill -x llama-server; sleep 6
|
||||
done
|
||||
```
|
||||
|
||||
Load → sample → kill, NO requests sent (allocation only). Run under `nohup`
|
||||
and poll from a separate call (a load is 30–90 s). Pick the smallest
|
||||
inter-card gap that clears the per-card floor with room to spare; one `-ts`
|
||||
step ≈ 1.5 GiB at ~6k batch / 262k ctx on 30 GiB cards — enough to bracket,
|
||||
not to solve. Prefer balance: the busy card's headroom bounds the context
|
||||
(`ctx_extra ≈ (floor − measured_free) / per-position cost on the busier
|
||||
card`). A faulted split (`SDMA0 UTCL2 page fault` → `illegal memory access` →
|
||||
`ggml_abort`) is OFF-LIMITS — re-measure, never re-tune; more headroom does
|
||||
not prevent a boundary-copy fault.
|
||||
|
||||
## 6. Drafter KV: sized from the TARGET, bounded by its window policy
|
||||
|
||||
`common/speculative.cpp` builds the draft context with
|
||||
`cparams.n_ctx = llama_n_ctx(ctx_tgt)` — **no flag caps it.** So:
|
||||
|
||||
```
|
||||
draft KV = ctx_tgt × block_count × 2 (K,V) × kv_heads × key_length × dtype
|
||||
```
|
||||
|
||||
…unless the drafter declares a sliding window. Measured at 262,144 target
|
||||
ctx: DFlash2 (`sliding_window = 2048`) ≈ 40 MiB; a same-shape 5-layer drafter
|
||||
with NO window = 20 KiB/position = **5.0 GiB** — a four-layer swing in the
|
||||
balance point, and swapping them invalidated an already-balanced split.
|
||||
|
||||
**Failure signature:** the engine loads the whole model, answers `/health`,
|
||||
and dies on the FIRST draft decode with `ROCm error: out of memory` /
|
||||
`cudaMalloc failed` inside `common_speculative_impl_draft_dflash::process` —
|
||||
a sizing failure, not a broken drafter; under a self-healing unit it presents
|
||||
as a crash-loop. Read `sliding_window`, `block_count`, `head_count_kv`,
|
||||
`key_length` from the drafter's GGUF header first. A compact reader (the
|
||||
three format traps: no KV alignment — v3 alignment applies only to the tensor
|
||||
section; arrays are `type(9) → element_type(u32) → count(u64)`, element type
|
||||
FIRST; the enum is `7=BOOL 9=ARRAY`, 64-bit = 10/11/12, and 8=STRING/9=ARRAY
|
||||
is the v2/v3 file order, NOT gguf.h's 2/3):
|
||||
|
||||
```python
|
||||
import struct, sys
|
||||
def rd(path):
|
||||
f = open(path, 'rb'); assert f.read(4) == b'GGUF'
|
||||
struct.unpack('<I', f.read(4)); struct.unpack('<Q', f.read(8))
|
||||
nk = struct.unpack('<Q', f.read(8))[0]
|
||||
def s():
|
||||
n = struct.unpack('<Q', f.read(8))[0]; return f.read(n).decode('utf-8', 'replace')
|
||||
SCALAR = {0:1,1:1,2:2,3:2,4:4,5:4,6:4,7:1,10:8,11:8,12:8}
|
||||
FMT = {0:'<B',1:'<b',2:'<H',3:'<h',4:'<I',5:'<i',6:'<f',7:'<?',10:'<Q',11:'<q',12:'<d'}
|
||||
keys = {}
|
||||
for _ in range(nk):
|
||||
k = s(); t = struct.unpack('<I', f.read(4))[0]
|
||||
if t == 8: keys[k] = s()
|
||||
elif t in SCALAR: keys[k] = struct.unpack(FMT[t], f.read(SCALAR[t]))[0]
|
||||
else:
|
||||
et = struct.unpack('<I', f.read(4))[0]; n = struct.unpack('<Q', f.read(8))[0]
|
||||
if et == 8: keys[k] = [s() for _ in range(n)]
|
||||
elif et in SCALAR: keys[k] = ('arr', et, n); f.read(n * SCALAR[et])
|
||||
else: raise Exception('unsupported arr elem type %d' % et)
|
||||
return keys
|
||||
```
|
||||
|
||||
## 7. Heterogeneous placement triage (go/no-go with arithmetic)
|
||||
|
||||
- **Link:** `LnkCap` vs `LnkSta` at every hop; expect **idle-downshift** (a
|
||||
Gen4 x4 dock idles at Gen1 x4, ~1 GB/s) — measure under load. Per-hop cost
|
||||
vs hidden state (5120-dim f16 = 10 KiB/hop ≈ 1.3 µs at Gen4 x4): bytes are
|
||||
almost never the constraint; **hop count is** — every boundary is a host
|
||||
sync, and a per-op-split runtime can hit 100 splits/token at ~5 % SM. Read
|
||||
`graph splits` from the loader log.
|
||||
- **KV is usually the bottleneck:** `KV/token = kv_bytes_per_position ×
|
||||
context` (attention reads the whole cache every token). Highest-leverage
|
||||
placement: **KV on the fastest-memory device together with its attention
|
||||
compute** — never KV alone over the link (GBs over 7.88 GB/s = hundreds of
|
||||
ms/token). Check fit: `kv_positions × slots × bytes + attention_trunk ≤
|
||||
VRAM`.
|
||||
- **What placement never buys:** prefill is compute-bound and expert weights
|
||||
are read once per batch — moving them buys nothing and adds a serial stage.
|
||||
Calibration: Strix Halo + 3090, moving 12.5 % of expert mass to the GPU
|
||||
gave decode +12 % but prefill 192 → 32 tok/s. The harmonic
|
||||
`1/(f_big/BW_slow + f_fast/BW_fast)` reproduces the decode gain within ~2
|
||||
points (assumes sequential stages — true at batch 1).
|
||||
- **UMA check before assuming additivity:** `mem_info_vram_total` (BIOS
|
||||
carve-out) vs `mem_info_gtt_total` (aperture onto SYSTEM RAM). A big GTT
|
||||
window means the "GPU memory" IS the same LPDDR: adding a stage there adds
|
||||
a compute engine, not capacity/bandwidth. Only discrete VRAM is additive.
|
||||
- **Speculation as the "cheap" split:** bandwidth-free (token ids per round)
|
||||
but needs the exact tokenizer for a plain drafter or a head trained for
|
||||
that target; a plain drafter costs K sequential forwards (break-even needs
|
||||
p ≈ 0.7+ at r≈0.4–0.5, K=3); compare against the INCUMBENT (a target
|
||||
already committing ~1.7 tok/round on its own MTP head is a high bar); and
|
||||
check the runtime accepts external drafters at all (closed engines 400
|
||||
unknown drafter names by design). Details: [[DFlash DSpark Speculative Decoding]].
|
||||
- **Hazards:** keep paged/sidecar tables where they are read from; keep mmap
|
||||
ON for huge-embedding models (anonymous memory → cgroup kill); CUDA graph
|
||||
capture can abort on unusual graph shapes; make the FIRST measurement a
|
||||
single-device baseline of the runtime on the new device — a split's
|
||||
arithmetic is worthless if the baseline loses to the incumbent; a split
|
||||
consumes the fast device entirely — check what it was serving.
|
||||
|
||||
**Verdict ranking:** (1) two independent endpoints — often the real answer;
|
||||
(2) one process, weights-resident, KV co-located with attention on the fast
|
||||
pool; (3) hot-expert cache only if `misses × bytes × tok/s` fits the link and
|
||||
overlaps — no mainstream runtime implements it; (4) name what does NOT help:
|
||||
TP across unequal devices, KV on the far device, loading over the link,
|
||||
prefill-motivated placements.
|
||||
|
||||
## 8. Harness syntax and verification
|
||||
|
||||
`llama-bench` separates a LIST with `/` and SWEEP values with `,`:
|
||||
`-dev CUDA0,Vulkan1` is TWO single-device runs, never one two-device run;
|
||||
`-ts 0.85,0.15` is two sweep values each becoming `[0.85, 0, …]` — device 0
|
||||
takes everything. Write `-dev A/B -ts x/y`. Read the table's own `dev` column;
|
||||
any placement-dependent number needs an independent memory check
|
||||
(`nvidia-smi` sampling while loaded). (This artifact produced a wrong fleet
|
||||
conclusion 2026-09-16 — "CUDA + Vulkan cannot serve one model" — with two
|
||||
bogus OOMs that were really all-111-GiB-on-the-24-GiB-card.)
|
||||
|
||||
Verification battery: `/health`; both cards' memory ≈ half the weights each
|
||||
(busier card also carries the drafter; compare `model_bytes/2 + drafter +
|
||||
ctx×KV/2 + buffers`); raise ctx by N and watch BOTH cards rise by
|
||||
N·bytes/2 — neither card's total may equal the model file (that arithmetic
|
||||
rules out replication); `spec_decode_num_drafts_total` > 0 and a real
|
||||
completion; a deep prompt (140k) proves a raised window is REAL, not
|
||||
declared; repeat prompt → `timings.prompt_n` collapses (prefix cache live);
|
||||
`/slots` shows expected `(id, n_ctx)` pairs; two concurrent requests with
|
||||
`journalctl -o cat | grep 'slot print_timing'` proves drafting on BOTH slots
|
||||
(default journal output truncates these lines).
|
||||
|
||||
Ops traps: a server started over ssh outlives stopping the runner (keeps the
|
||||
port — `pkill -x <binary>` first); `systemctl start` of a root unit needs
|
||||
sudo (polkit "Access denied" otherwise); old crash `ROCm error` lines live in
|
||||
the journal forever — check `ActiveEnterTimestamp`/`NRestarts` before
|
||||
treating a grep count as a live fault.
|
||||
|
||||
## Related
|
||||
|
||||
- [[PCIe Link and Atomics]] — the atomics wall, link probing, repair
|
||||
- [[Llama Cpp Serving Sizing]] — slot-pool and KV pricing
|
||||
- [[DFlash DSpark Speculative Decoding]] — spec types, chaining, A/B
|
||||
- [[frank-medium-llamacpp-dflash2]] — the layer-split era runbook
|
||||
- [[Dual-R9700-Cooperation]] / [[x4-pair-experiment]] — the measured history
|
||||
Reference in new issue
Block a user