Files

T

Frank 4fca739bf8 frank-medium: deploy two single-GPU replicas — TP=2 is blocked by PCIe atomics

The radiance MXFP4 stack cannot run TP=2 on this host. RCCL's collective
kernels carry a hidden hostcall-buffer parameter, and ROCr refuses to
dispatch them on a device without PCIe atomics: the AQL dispatch fails with
hipErrorIllegalState ("the operation cannot be performed in the present
state"), which is what every two-rank collective hit. The second R9700 sits
behind the chipset bridges 05:00.0/06:09.0, which advertise
`AtomicOpsCap: Routing-` — a read-only capability, so no software can enable
atomic routing on that path and the driver's
pci_enable_atomic_ops_to_root() fails for that card. The first card, on CPU
root port 00:01.1, has atomics enabled.

Ruled out: kernel version (7.0.0-28/-30/-31 identical, after a full upgrade),
iommu=pt, P2P hardware (45.3 GB/s between the cards), the shipped librccl (a
from-source rebuild reproduces the failure at the same call path), vLLM and
pynccl (a bare two-rank torch.distributed allreduce fails the same way), all
NCCL transport knobs, and single-rank tests (NCCL short-circuits at
world_size=1, so one-card success proves nothing).

Deployment: three rootless Quadlets — frank-medium-a (:7771, card 1),
frank-medium-b (:7772, card 2), both TP=1 at 114688 context (262144 and
131072 do not fit one card's KV) with vLLM-profiled KV — and
frank-medium-proxy (:7770), a round-robin that preserves the fleet contract
of one endpoint and one served model name. frank-medium.container stays as
the TP=2 target for a host whose second card is atomic-capable; the deploy
handler does not install it and trashes any stale copy.

Measured through the proxy (vllm bench serve, random dataset, --ignore-eos):
conc 1 42.8 tok/s, conc 4 147.6 tok/s, conc 8 230.8 tok/s, 32k prefill at
~2.8k tok/s per card, TPOT p50 23.6-28.8 ms.

Also: healthcheck reads PORT so one probe serves all three units, the hermes
example and the harness models.yml drop to the real 114688 context, and the
runbook records the full evidence chain.

2026-09-16 03:47:11 -07:00

10 KiB

Raw Blame History

type, title, description, tags, timestamp

type

title

description

frank-medium ROCmFPX Runbook

frank-medium is the Dankserver LLM endpoint (port 7770) since 2026-09-15, when the RTX 3090 left the box and two AMD Radeon AI PRO R9700 cards (Navi 48, RDNA4, gfx1201, 32,624 MiB each, PCIe 5.0 x16) took its place. It replaced the retired FreeToken/CUDA frank-small (freetoken-frank-small), which cannot run on this hardware at all.

Stack: charlie12345/ROCmFPX llama.cpp (main, c49ebdb, 2026-08-22) serving kingjones777/Qwen3.8-27B-ROCmFPX-Q8_0-AGENT-GGUF with the companion MTP draft head. Stock llama.cpp cannot load this GGUF — Q8_0_ROCMFPX_AGENT is ftype 115 / GGML_TYPE_Q8_0_ROCMFPX (type 103) and exists only in the fork.

Hardware

Property	Value
GPUs	2x AMD Radeon AI PRO R9700 (DID 0x7551), rev c0
Architecture	gfx1201 (Navi 48, RDNA4)
VRAM	32,624 MiB each (rocm-smi)
Link	PCIe, no XGMI; cards sit in different IOMMU groups
iGPU	Raphael (gfx1036) enumerates as ROCm device 2 — hidden from HIP
CPU / RAM	Ryzen 7 7700X (8c/8t), 60 GiB RAM, 63 GiB swap
Kernel driver	in-tree `amdgpu`; no new driver needed

HIP_VISIBLE_DEVICES=0,1 hides the iGPU so nothing can drift onto shared system RAM. Devices are addressed as ROCm0,ROCm1 (llama-server) or ROCm0/ROCm1 (llama-bench multi-device syntax — see the bench gotcha below).

Runtime build (self-contained ROCm 7.14)

The host ROCm install is quarantined rather than repaired: the build uses the fork's self-contained toolchain, so the binary carries its own libamdhip64/libhsa-runtime64/rocblas/hipblaslt with RPATH=$ORIGIN/../lib and never touches system ROCm.

cd /home/dank/GitPackages/charlie-ROCmFPX
env JOBS=8 ./scripts/build-rocmfp4-rocm714-local.sh
# -> build-rdna4-rocm714/bin/llama-server

The script downloads the gfx120X-all TheRock ROCm 7.14 nightly tarball into ~/.local/share/rocm-7.14.0a20260624 (~17 GiB extracted; tarball kept next to the build tree) and builds gfx1200;gfx1201.

gfx1201 is not interchangeable with gfx1200 — the fork's docs/BUILD-AMD-ARCHITECTURES.md warns a wrong-sibling build links, loads the model, then segfaults. scripts/rocmfpx-hip-arch.sh auto-detects and refuses.

Model

File	Size	Role
`Qwen3.8-27B-Q8_0_ROCMFPX_AGENT.gguf`	28,214,425,952 B (26.28 GiB, 8.39 bpw)	target
`mtp-Qwen3.8-27B-Q4_0.gguf`	1,680,271,648 B	MTP draft head
`mmproj-Qwen3.8-27B-BF16.gguf`	931,145,888 B	vision tower — unused (`--no-mmproj`)

The AGENT routing keeps 13 of the 17 *.output.weight tensors at full Q8_0 (12 more than the plain Q8_0_ROCMFPX build), which raises MTP draft acceptance (0.953 vs 0.911 on the author's gfx1151 box). Take this variant only when serving with MTP. Fetch with scripts/download-q8agent.sh (resumable, size-verified; run --verify-only to re-check).

The model card documents 131,072 context; llama.cpp reports n_ctx_train (262144) from the GGUF, but the author never tested beyond 131k, so the deployment uses 131,072.

Parallelism: layer split, not tensor (measured)

llama-bench on the host (pp4096 / tg128, same at depth 32768, f16 KV, flash-attn on, 3 reps):

Split	pp4096	tg128	pp4096 @d32768	tg128 @d32768
1x R9700 (`-sm none`)	1160.9	19.42	~512	~18.4
2x layer (`-sm layer`)	1991.2	18.76	859.1	17.64
2x tensor (`-sm tensor`)	271.8	26.26	228.6	25.09

Tensor parallel wins decode by ~40 % but collapses prefill 7x: this build ships without RCCL, so llama.cpp logs internal AllReduce init failed (n_devices != 2?); falling back to meta-backend butterfly and every layer pays a slow cross-card reduction. The workload is prefill-heavy (agentic, long prompts, reasoning traces), so the operator switched to layer split on 2026-09-15. Revisit if the fork ships a working RCCL tensor path (-DGGML_HIP_RCCL=ON).

Served-path comparison with MTP on (real completions): layer 35.7–46.2 tok/s vs tensor 28.7 tok/s — layer wins decode too once the draft head is in play.

llama-bench multi-device gotcha

llama-bench and llama-server disagree on device syntax. In llama-bench, -dev ROCm0,ROCm1 means two separate single-card benches (the value is split on , into separate device sets, one model load each); a single split bench needs -dev ROCm0/ROCm1. The invalid form silently produces single-card numbers in a table that looks like a split test.

Tensor-mode draft abort (if tensor is ever revisited)

Under --split-mode tensor the MTP draft head must be pinned to exactly one device (--spec-draft-device ROCm0); with it empty the draft joins the meta-device KV allocation and the server aborts while building the KV cache:

ggml/src/ggml-backend.cpp:119: GGML_ASSERT(buffer) failed
  <- ggml_backend_meta_alloc_ctx_tensors_from_buft <- llama_kv_cache ctor

Deployment

File	Role
`scripts/frank-medium.env`	all tunables (sampler, split, draft, context)
`scripts/serve-frank-medium.sh`	env-driven llama-server launcher
`scripts/frank-medium.service`	root systemd unit (relinked into /etc/systemd/system)
`scripts/download-q8agent.sh`	model fetch + size verification

./scripts/deploy.sh dankserver --service frank-medium

First deploy (2026-09-15) also ran retire_root_llm_small() in scripts/deploy-remote.sh: the CUDA-era frank-small unit was stopped, disabled and moved to ~/.AGENTTRASH/20260915-frank-small-retire/host-units (the host's systemctl disable has been observed to delete unit files, 2026-08-26 incident, so both outcomes are tolerated).

Deployed flags (see scripts/frank-medium.env for the full set): --split-mode layer --device ROCm0,ROCm1 --gpu-layers 999 --fit off, --ctx-size 131072 --parallel 1, f16 K/V, flash-attn on, --spec-type draft-mtp (n_max 6, p_min 0.6) with the draft pinned to ROCm1 — the card holding the target model's last layers, where the final hidden state is produced. Reasoning on, deepseek format, unlimited budget; sampler defaults are the checkpoint's instruct table (temp 0.7, top-p 0.8, top-k 20, min-p 0, presence 1.5). Text only — vision stays on frank-large (Dankclaw).

Verified serving behavior (2026-09-15)

Health {"status":"ok"} ~18 s after start; model load ~17 s (warm page cache).
MTP engaged: draft acceptance 60–69 % across prompts (timings.draft_n / draft_n_accepted in every chat response).
Decode 35.7–46.2 tok/s on a 200-token budget; cold prompt 949 tokens prefilled in 2036 ms (~466 tok/s at ubatch 512).
Prefix caching works with the draft head: the model card's mtp-prompt-cache-fix.patch defect (draft head disables prefix caching) is already fixed in c49ebdb — the speculative boundary rides inside the context checkpoint (common_prompt_checkpoint::data_spec). Measured: turn 2 of a conversation reprocessed 22 tokens (cache_n 945) instead of 949.
Quirk: the request-level "speculative": {"n_max": 0} override does not disable drafting (still drafts and accepts). To serve without MTP, restart with SPEC_TYPE=none.

Client contract

OpenAI-compatible /v1 on port 7770, alias frank-medium. Reasoning traces are projected into reasoning_content (deepseek format); requests that send no sampler fields get the server's CLI defaults, not greedy. Text only. Harness routing after the cutover: blankpad smol/tiny/commit/default → frank-medium, Hermes title_generation (and its other dankserver-small routes) → frank-medium, vision stays frank-large.

Host ROCm posture (post-cleanup)

Historically the host had a broken mix: AMD-repo ROCm 6.3.3 in /opt/rocm-6.3.3 hijacked libhsa-runtime64 via /etc/ld.so.conf.d/10-rocm-opencl.conf, breaking Ubuntu's ROCm 7.1 userspace (undefined reference to 'hsa_amd_memory_get_preferred_copy_engine'). On 2026-09-15 the stale 6.3.3 packages (~24.04-suffixed, 167 pkgs) were purged, the hijack conf and the repo.radeon.com apt pin moved to ~/.AGENTTRASH/20260915-rocm633-cleanup/, and ldconfig now resolves only Ubuntu's ROCm 7.1 libs. rocm-smi is Ubuntu's own 7.1.1 package. The self-contained toolchain in ~/.local/share/rocm-7.14.0a20260624 does not depend on any of it.

Troubleshooting

Wrong arch build: a gfx1200 binary on this card links, loads, then segfaults. Rebuild with the pinned script and check the "Build complete for gfx1200;gfx1201" line plus the emitted code objects.
Crash at load: read the journal (sudo journalctl -u frank-medium -n 100), then run the serve script directly; the bundled libs make the script and the raw binary behave identically. The known abort shape is the tensor-mode draft KV abort above.
MTP "not engaged": check timings.draft_n in a response; zero means the draft head is not loaded (look for common_speculative_init: adding speculative implementation 'draft-mtp' in the journal).
Model files missing/truncated: bash scripts/download-q8agent.sh --verify-only.

10 KiB Raw Blame History Unescape Escape