---
name: llamacpp-multigpu-layer-split-cutover
description: "Cut a fleet LLM endpoint over to llama.cpp layer-split across multiple GPUs when cross-GPU collectives are impossible (RCCL/PCIe-atomics wall), including the build recipe, the bundling trap that fails on the first request, DFlash2/MTP wiring, and the measurement battery that proves the split and the drafter are real."
---

# llama.cpp layer-split cutover (no-collective multi-GPU)

Use when one model must serve from TWO+ cards and the box cannot dispatch a
cross-card collective (second card behind a chipset bridge with
`AtomicOpsCap: Routing-`, so ROCr refuses RCCL's hostcall-buffer kernels).
Decide with one line before designing anything:

```bash
sudo lspci -vv -s <card2> | grep AtomicOpsCtl      # ReqEn- ⇒ no collectives, ever
```

`ReqEn-` kills vLLM entirely — tensor parallelism needs an allreduce **and
pipeline parallelism needs an RCCL `send`/`recv` of the hidden state**
(`vllm/distributed/parallel_state.py`: metadata on gloo, tensors on the device
backend). llama.cpp's `--split-mode layer` needs no collective: layers are
distributed in order, only the boundary activation crosses PCIe once per
forward pass.

## 1. Pin the runtime by FEATURE, not by "latest"

Write down which features the deployment needs, then find the ref that has all
of them and pin the 40-hex sha in the build script. On 2026-09-16 that was
DFlash2 (`ggml-org/llama.cpp#27342`, detected from the draft GGUF:
`selector_top_k > 0`, spec type stays `draft-dflash`) and qwen35 MTP. The
ROCmFPX/r9700-toolbox images/fork tips predate both — verify by listing the
built image:

```bash
ls <venv>/lib/python*/site-packages/vllm/v1/worker/gpu/spec_decode/   # dflash/ vs dflash2/
```

## 2. Build self-contained for one arch

- Private ROCm toolchain (TheRock nightly) with `ROCM_PATH`/`HIP_PATH` exported,
  `-DCMAKE_HIP_ARCHITECTURES=<gfx only>` (a sibling arch links, loads and then
  segfaults), `-DGGML_HIP_RCCL=OFF` when there is no atomics path, `-DGGML_NATIVE=OFF`.
- Bundle the runtime into `build/lib` and set `--set-rpath '$ORIGIN/../lib'`.
- **TRAP — `cp -r SRC DST/` flattens when `DST` does not exist.** rocBLAS and
  hipBLASLt look for device code at a FIXED depth (`../lib/rocblas/library`,
  `../lib/hipblaslt/library/<arch>`). Flattened, the engine loads every GiB of
  weights, 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/`.

## 3. Deploy as an env + serve script + unit trio

One env file holds every tunable with the sizing arithmetic in comments; the
serve script defaults each var, refuses to start when the runtime or payload is
missing, and builds an args array (spec flags assembled conditionally so the
same script serves `draft-dflash` with a file and `draft-mtp` in-model). Root
systemd unit: `User=<deploy user>`, `LimitMEMLOCK=infinity`,
`Restart=on-failure`, `WantedBy=multi-user.target`, plus an
`Environment=…_ENV_FILE=<repo path>` so the repo file is the single source.

Deploy handler must: refuse when runtime/payload is missing; retire the previous
runtime's units (into trash, never `rm`); **wait for the port to be released**
(a retired reverse proxy or container holds it); then enable/restart and gate on
`/health`.

DRLaf: `--split-mode layer --device ROCm0,ROCm1 --tensor-split 1,1`,
`--parallel 1`, f16 KV, `--fit off`. Pin the drafter with `--spec-draft-device`
to the card holding the target's last layers. `--spec-draft-n-max` must be
<= the draft's `block_size - 1` (DFlash2: 8 ⇒ 7).

## 4. Prove it, do not assert it

- **The split is real**: raise `--ctx-size` by N tokens and watch `rocm-smi` —
  a layer split adds ~N·(bytes/token)/2 to EACH card (both cards carry the KV
  for their own attention layers), and neither card's TOTAL can equal the model
  file, which is the arithmetic that rules out replication.
- **The drafter is live**: `timings.draft_n`/`draft_n_accepted` per request, and
  the cumulative `llamacpp:spec_decode_num_drafts_total` — the flags can be
  right while the draft model silently failed to load. The loader's init line
  (`block_size=`, `n_extract=`, `mask_token_id=`) identifies the drafter family.
- **Watch for a 500 on `reasoning_effort`**: Qwen3.8-family chat templates
  accept only low/medium/xhigh and RAISE on `high`; map every harness level or
  subagents 500 on their first request.
- **Re-point the monitors in the same change**: a runtime swap changes the
  metric namespace (`vllm:` → `llamacpp:`), and rules/dashboard panels naming
  series that no longer exist draw nothing forever without erroring. Retire
  what cannot be expressed (llama.cpp exports no KV-fill gauge and no latency
  histogram) and name the consequence instead (queue depth, spec-decode-off).
- One slot is a concurrency policy, not a bug: `--parallel` DIVIDES the context.