command-center/agent/managed-skills/llamacpp-layer-split-rdna4/SKILL.md

name, description

name	description
llamacpp-layer-split-rdna4	Serve one model across two (or more) AMD RDNA4/ROCm GPUs on a host where RCCL cannot dispatch collectives — mainline llama.cpp built for gfx1201 with a bundled ROCm runtime, --split-mode layer, DFlash2/MTP speculative decoding, and the verification that proves each layer works

llama.cpp layer-split serving on RDNA4 (gfx1201)

Use when one model must use two AMD GPUs whose interconnect cannot carry a collective (no PCIe atomic routing on one card's path), or when a host has no usable RCCL. The shape: ONE llama-server process, --split-mode layer across both cards, spec decode from a drafter GGUF or the model's own MTP head.

When to reach for this

• vLLM cannot help: TP needs a per-layer allreduce and PP's stage handoff is a device-backend send/recv (vllm/distributed/parallel_state.py, send_tensor_dict → tensor handles "on the device backend"), i.e. the same RCCL primitive that fails on an atomics-less card. Verify with sudo lspci -vv -s <gpu> | grep AtomicOpsCtl (want ReqEn+) and lspci -tv (a Routing- bridge is a hardware wall).
• Confirming the wall: any RCCL kernel launch on that card returns hipErrorIllegalState with AMD_LOG_LEVEL=3 printing Pcie atomics not enabled, hostcall not supported → AQL dispatch failed!. A single-card test proves nothing (NCCL short-circuits at world_size=1).
• Never run a multi-GPU collective probe beside a live service on these hosts: stop and mask it first, or use a host that serves nobody.

Build (pin the ref, bundle the runtime)

git clone https://github.com/ggml-org/llama.cpp && cd llama.cpp
git checkout --detach <ref>          # pin; do not float
# ROCm toolchain: system ROCm or a TheRock nightly tarball, e.g.
#   https://rocm.nightlies.amd.com/tarball-multi-arch/therock-dist-linux-gfx120X-all-<ver>.tar.gz
ROCM=$HOME/.local/share/rocm-<ver>
cmake -S . -B build-rdna4 -G Ninja -DCMAKE_BUILD_TYPE=Release \
  -DCMAKE_C_COMPILER=$ROCM/llvm/bin/clang -DCMAKE_CXX_COMPILER=$ROCM/llvm/bin/clang++ \
  -DCMAKE_CXX_FLAGS="-I$ROCM/include" -DCMAKE_POSITION_INDEPENDENT_CODE=ON \
  -DGGML_HIP=ON -DGGML_HIP_RCCL=OFF -DGGML_CUDA=OFF -DGGML_VULKAN=OFF \
  -DCMAKE_HIP_ARCHITECTURES=gfx1201 -DGPU_TARGETS=gfx1201 \
  -DLLAMA_BUILD_SERVER=ON -DLLAMA_BUILD_WEBUI=OFF -DBUILD_SHARED_LIBS=OFF \
  -DCMAKE_PREFIX_PATH=$ROCM -DROCM_DIR=$ROCM
cmake --build build-rdna4 -j $(nproc) --target llama-server llama-cli llama-bench
patchelf --set-rpath '$ORIGIN/../lib' build-rdna4/bin/llama-*

• -DGGML_HIP_RCCL=OFF is deliberate: with RCCL off, only --split-mode tensor degrades (butterfly fallback); layer split is unaffected.
• Bundle device libraries at the exact depth the runtime looks for. Copy $ROCM/lib/rocblas/library/ → build/lib/rocblas/library/ and $ROCM/lib/hipblaslt/library/ → build/lib/hipblaslt/library/ with mkdir -p DEST && cp -r SRC/. DEST/. cp -r SRC DEST/ with a missing DEST flattens one level and the failure appears at REQUEST time, not load time (Cannot read .../lib/rocblas/library/TensileLibrary.dat).
• Build only the arch you own. A sibling-arch binary (gfx1200 on gfx1201) links, loads the model, then segfaults.
• Verify: build-rdna4/bin/llama-server --version and --list-devices (expect both cards; an iGPU can enumerate as a third).

Serve

HIP_VISIBLE_DEVICES=0,1 llama-server \
  --model <target.gguf> --alias <fleet-name> --host 0.0.0.0 --port 7770 \
  --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 \
  --reasoning on --reasoning-format deepseek --reasoning-budget -1 \
  --spec-type draft-dflash --spec-draft-model <dflash2-q8_0.gguf> --spec-draft-n-max 7

• --fit off + an explicit --ctx-size: auto-fit must not resize a context you set on purpose. --parallel N DIVIDES --ctx-size; it does not multiply.
• Spec decode choices: --spec-type draft-dflash with a DFlash2 GGUF (the drafter's own metadata selects the DFlash2 path — the init line prints block_size=8, n_extract=5, sample_from_anchor=true); n_max <= block_size-1. Or draft-mtp for a model whose GGUF carries blk.<n>.nextn.* (no draft file; the head loads from the main weights). Types can be chained (draft-mtp,draft-dflash). With draft-dflash, the target's nextn tensors log as "unused tensor ... ignoring" — that is correct, not a broken download (check bytes with the fetch script's --verify-only).
• Layer split puts the drafter on the card holding the target's last layers (--spec-draft-device ROCm1); under tensor mode an unpinned draft joins the meta-device KV allocation and aborts the KV build.

Verify (each layer separately)

1. Load: n_ctx_slot in the log matches --ctx-size; /health answers.
2. Both cards: rocm-smi --showmeminfo vram — each card holds roughly half the weights; the busier one also carries the drafter. Compare against model_bytes/2 + drafter + ctx*64KiB/2 + ~2 GiB (qwen35 hybrid: only its full-attention layers take KV, 64 KiB/token at f16).
3. Spec decode is live: curl /metrics | grep spec_decode_num_drafts_total and ..._accepted_tokens_total (their ratio is acceptance); a request's timings.draft_n / draft_n_accepted gives it per call. Drafts stuck at 0 with tokens flowing means the drafter is not running.
4. Real completion through the fleet alias, then a deep prompt (a 140k-token prompt proves a raised window is real, not just declared).
5. Prefix cache: repeat the same prompt — timings.prompt_n should collapse to a handful of tokens.

Pitfalls that cost real time

• A healthy-looking build that dies on the first request is almost always the bundled-library depth above.
• Serve scripts that source an env file AFTER defaults cannot be overridden from the command line; override the env-file PATH instead.
• Qwen3.8 chat templates RAISE on reasoning_effort=high (500). Probe low/medium/high/xhigh before declaring an endpoint ready, and map the harness's levels onto the supported set.
• One slot at a large context is the concurrency limit: concurrent agent loops queue, and first-turn TTFT then includes the wait.