feat(monitoring,harness): follow frank-medium onto vllm:* metrics and fix red-by-construction targets
The radiance TP=2 cutover changed the metric namespace frank-medium exports (llamacpp:* -> vllm:*) and left two scrape targets addressing host processes by an address a container cannot reach. - shared/prometheus/rules/llm-fleet.yml: two namespaces side by side. Added LlmRequestsWaiting, LlmKvCacheSaturated (vllm-only gauges llama.cpp never had), the job:llm_vllm_prefix_cache_reuse:rate5m and job:llm_spec_accept_length:rate5m recording rules, and re-cut LlmPrefixCacheCollapse / LlmSpecDecodeOff to OR both families with per-family traffic floors so an idle endpoint cannot fire them. Every expression validated against live Prometheus. - shared/prometheus/prometheus.yml: frank-small -> host.containers.internal:7771 and frank-large scraped at host.containers.internal:7770. Both are HOST processes on Dankclaw and netavark DNATs published container ports only, so the LAN-IP targets were red by construction; frank-large's CIRU llama.cpp build does export the llamacpp:* family, contra the old comment. - tests/observability_quadlet.bats: the two seams now encode the new contract (per-endpoint reachability by process type; every named series must be on the measured list for its engine; engine-independent rules must name both). - homepage docs.html/services.json: frank-medium sees pictures, frank-large is CIRU Orca with a single seat (was halogen w4b, two seats, 200k). - docs: radiance runbook is ACTIVE (deployed flags, three-arm battery, verified capabilities, rollback, replica-era sections labelled); llama.cpp runbook carries the SUPERSEDED banner; Dual-R9700-Cooperation states the measured verdict; Observability records the namespace split and the DNAT lesson; AGENTS.md topology/frank-medium/Model Config/Benchmarking Policy/runbook index flipped; both indexes updated; work-log entry added. Also repaired agent-pass damage: prose prepended above frontmatter openers, duplicated runbook bodies, a literal tool-call token pasted into a runbook, and an index destroyed by diff artifacts (restored from HEAD, flips re-applied). Vault link check: 128 pages, 0 unresolved. ./tests/run.sh: 325 pass.
This commit is contained in:
Frank committed
1 parent
c2116658f8
commit
08f1c7ad09
13 files changed
+687
-286
No files matched your search
@@ -14,12 +14,11 @@ timestamp: 2026-09-16
|
||||
|
||||
# Dual R9700 Cooperation
|
||||
|
||||
Dankserver's two R9700s serve `frank-medium` as **one llama.cpp process with a
|
||||
layer split** ([[frank-medium-llamacpp-dflash2]]). This page records why true
|
||||
tensor parallelism was impossible on the previous board (TUF GAMING B650-E
|
||||
WIFI), and what the 2026-09-20 board swap changed. See
|
||||
[[frank-medium-radiance-mxfp4]] for the replica-pair deployment that ran under
|
||||
the old constraint.
|
||||
Dankserver's two R9700s serve `frank-medium` as **one vLLM-radiance process at
|
||||
tensor_parallel_size 2** ([[frank-medium-radiance-mxfp4]]). This page records why
|
||||
true tensor parallelism was impossible on the previous board (TUF GAMING B650-E
|
||||
WIFI), and what the 2026-09-20 board swap changed. The llama.cpp layer split that
|
||||
served the endpoint until that swap is [[frank-medium-llamacpp-dflash2]].
|
||||
|
||||
## The short answer
|
||||
|
||||
@@ -35,18 +34,21 @@ bifurcates the CPU's GPP ports **x8/x8**: the board-swap path this page
|
||||
recommended, at B850 rather than X870E money. Both cards now hang off CPU root
|
||||
ports (`00:01.1`, `00:01.3`), every bridge on both chains reads
|
||||
`AtomicOpsCap: Routing+`, and **both endpoints read `AtomicOpsCtl: ReqEn+`** —
|
||||
the exact precondition this page's own [test plan](#test-plan) names.
|
||||
the exact precondition this page's own test plan named, and it is now
|
||||
[measured](#measured-verdict-2026-09-20) rather than assumed.
|
||||
|
||||
What that buys, and what it does not:
|
||||
What that bought:
|
||||
|
||||
- **Unblocked in principle:** vLLM TP=2, vLLM PP=2, and llama.cpp
|
||||
`--split-mode row`. All three died on the same hostcall gate.
|
||||
- **Not yet measured.** Nothing has dispatched a collective on this host since
|
||||
the swap. The 2-rank probe is the remaining gate, and policy is stop-and-mask
|
||||
production first.
|
||||
- **Not needed for capability.** The layer split already puts both cards on one
|
||||
model. This is a performance question: an RCCL-backed tensor split, and the
|
||||
single pooled KV that comes with it.
|
||||
- **Both cooperation modes now work, and they were both run.** vLLM TP=2 booted
|
||||
and served (its R4D allreduce needs no RCCL collective at all), and llama.cpp
|
||||
`--split-mode tensor` booted and served against an RCCL-backed build — so a
|
||||
collective HAS been dispatched on this host since the swap.
|
||||
- **One of them is production.** TP=2 wins on every axis measured, by 2.1x on
|
||||
decode and 3.6x at 4-way concurrency, so it replaced the layer split on the day
|
||||
of measurement. See [[frank-medium-radiance-mxfp4]] for the deployed shape.
|
||||
- **The drafter, not the collective, is what llama.cpp lost.** Its tensor mode
|
||||
could not carry the DFlash2 drafter, which is the honest caveat on the
|
||||
comparison below.
|
||||
- **The bandwidth story moved too.** Both cards ride Gen5 x8 (~31.5 GB/s each):
|
||||
card 0 effectively unchanged (was Gen4 x16), card 1 up ~16x (was Gen4 x1,
|
||||
~1.97 GB/s). The M.2→PCIe riser plan is dead — `PCIEX16_2` holds card #2 now
|
||||
@@ -211,28 +213,47 @@ landed, at x8/x8, clears the arithmetic with room to spare.
|
||||
- Independent corroboration: StillDeadcode/vllm-radiance Codeberg issue #22 —
|
||||
different host, same image, kernel 7.0.0-30, same `enqueue.cc` failure.
|
||||
|
||||
## Test plan (when the hardware changes)
|
||||
## Measured verdict (2026-09-20)
|
||||
|
||||
**Precondition met 2026-09-20:** the line below reads `ReqEn+` on both cards, so
|
||||
this is now a standing plan rather than a hypothesis.
|
||||
Both cooperation modes were run on this host on 2026-09-20 — the same salted
|
||||
10-prompt battery (`scripts/bench-conversation.py`), the service stopped and
|
||||
masked between arms. Decode/prefill are geometric means over the ten rows;
|
||||
burst-4 is four concurrent copies of one prompt, reported as aggregate tokens/s.
|
||||
|
||||
The verdict to look for is one line of `lspci`, before touching any model:
|
||||
| Shape | Cross-device path | Decode | TTFT | Prefill | burst-4 |
|
||||
|---|---|---|---|---|---|
|
||||
| **radiance TP=2 + DFlash2 — deployed** | R4D one-shot allreduce | **95.8 t/s** | **0.21 s** | **1493.6 t/s** | **259.4 t/s (65.6/stream)** |
|
||||
| llama.cpp layer split 32:33 + DFlash2 — predecessor | none, boundary activation copy | 45.7 t/s | 0.72 s | ~530 t/s | 71.7 t/s |
|
||||
| llama.cpp `--split-mode tensor` 1:1 | RCCL allreduce | 30.5 t/s | 0.71 s | n/r | 54.3 t/s |
|
||||
|
||||
```bash
|
||||
sudo lspci -vv -s <card2> | grep AtomicOpsCtl # want: ReqEn+
|
||||
```
|
||||
Row 3 carries a caveat that no flag can remove: llama.cpp's tensor mode could not
|
||||
load the DFlash2 drafter (it aborts at `ggml-backend.cpp:941` — a draft weight
|
||||
pre-allocated on the non-primary device has no backend once `n_copies == 1`), so
|
||||
it decoded **without** speculation while the other two rows decoded with it. The
|
||||
comparison that decides anything is row 1 against row 2, and some of that gap is
|
||||
the 4-bit MXFP4 checkpoint rather than the parallelism.
|
||||
|
||||
If that reads `ReqEn+`, the rest is mechanical: `--tensor-parallel-size 2` with
|
||||
`HIP_VISIBLE_DEVICES=0,1`, then the same benchmark set
|
||||
(`scripts/bench-turn-latency.py`, and the endpoint's own `timings` block) so the
|
||||
two architectures are compared on identical workloads. If it still reads
|
||||
`ReqEn-`, the card is still on a non-atomic path — check `lspci -tv` for the
|
||||
bridge chain rather than assuming the move worked.
|
||||
Two numbers elsewhere on this page are therefore not interchangeable:
|
||||
|
||||
- The 272 t/s recorded 2026-09-16 is **prefill**, on a **butterfly** allreduce,
|
||||
with RCCL off, on the old board. It was never a TP decode number.
|
||||
- 26.3 t/s from the same table is that butterfly tensor's decode. Row 3's 30.5 is
|
||||
an RCCL tensor decode a day later: same shape family, different collective path.
|
||||
|
||||
RCCL needs PCIe atomics on *every* participating GPU — the gate the old board
|
||||
failed. vLLM's R4D one-shot allreduce dispatches no collective at all, which is
|
||||
why TP=2 runs today rather than merely "being unblocked": the image ships TheRock
|
||||
RCCL, and TheRock carries no gfx1201 device code (`invalid device function`), so
|
||||
radiance could not have used RCCL on this host even with atomics routed. Ubuntu's
|
||||
RCCL 2.27 does carry gfx1201 kernels, and that is what row 3 linked against.
|
||||
|
||||
**Rule for any future multi-GPU collective experiment: stop and mask the
|
||||
production service first** (`sudo systemctl stop frank-medium && sudo systemctl
|
||||
mask frank-medium`, per the benchmarking policy), or run it on a host that is
|
||||
not serving users. On 2026-09-16 a bare RCCL P2P probe was run against the live
|
||||
production service first**, per the benchmarking policy, or run it on a host that
|
||||
is not serving users. Since 2026-09-20 `frank-medium` is a rootless quadlet, so
|
||||
that is `systemctl --user stop frank-medium && systemctl --user mask frank-medium`
|
||||
on Dankserver.
|
||||
|
||||
This is not theoretical: on 2026-09-16 a bare RCCL P2P probe was run against the live
|
||||
service; the host went hard down ~20 minutes later (no ARP, no ICMP, required a
|
||||
power cycle). The mechanism is not proven, but a hung AQL dispatch wedging the
|
||||
GPU driver is a plausible cause and this host has precedent for going
|
||||
|
||||
@@ -1,7 +1,7 @@
|
||||
---
|
||||
type: configuration
|
||||
title: frank-medium llama.cpp Runbook (DFlash2 drafter)
|
||||
description: The ACTIVE frank-medium deployment since 2026-09-16 (model swapped 2026-09-19) — mainline llama.cpp on Dankserver serving orcarouter/Qwen3.8-27B-Uncensored-GGUF at Q8_0 with the z-lab DFlash2 Q4_K_M drafter (the model's own MTP head stays in the file, one SPEC_TYPE change away), as ONE layer-split process across both Radeon AI PRO R9700 cards at a balanced 32:33 split, a 327680-position KV pool (163840 per conversation over two slots), f16 KV plus a 16 GiB host-RAM prompt cache, vision via the BF16 mmproj.
|
||||
description: The former ACTIVE frank-medium deployment since 2026-09-16 (model swapped 2026-09-19) — mainline llama.cpp on Dankserver serving orcarouter/Qwen3.8-27B-Uncensored-GGUF at Q8_0 with the z-lab DFlash2 Q4_K_M drafter (the model's own MTP head stays in the file, one SPEC_TYPE change away), as ONE layer-split process across both Radeon AI PRO R9700 cards at a balanced 32:33 split, a 327680-position KV pool (163840 per conversation over two slots), f16 KV plus a 16 GiB host-RAM prompt cache, vision via the BF16 mmproj.
|
||||
tags:
|
||||
- frank-medium
|
||||
- dankserver
|
||||
@@ -12,12 +12,17 @@ tags:
|
||||
- gfx1201
|
||||
- r9700
|
||||
- layer-split
|
||||
timestamp: 2026-09-19
|
||||
timestamp: 2026-09-20
|
||||
---
|
||||
|
||||
> **SUPERSEDED 2026-09-20.** Replaced by [[frank-medium-radiance-mxfp4]]
|
||||
> (radiance TP=2, measured faster on every axis). The layer-split shape below is
|
||||
> the rollback target, and the llama.cpp measurements in this page stay valid for
|
||||
> it — but nothing here is the deployed shape any more.
|
||||
|
||||
# frank-medium — llama.cpp (DFlash2 drafter)
|
||||
|
||||
The ACTIVE frank-medium deployment since 2026-09-16 (model swapped to
|
||||
The frank-medium deployment from 2026-09-16 to 2026-09-20 (model swapped to
|
||||
orcarouter Q8_0 on 2026-09-19): **mainline llama.cpp**
|
||||
serving `orcarouter/Qwen3.8-27B-Uncensored-GGUF` at `Q8_0`
|
||||
with the **z-lab DFlash2** `Q4_K_M` block-diffusion drafter, as **one process across both
|
||||
@@ -26,9 +31,11 @@ model (ajgazin Swift-Qwen3.8-27B UD-Q8_K_XL, 35:30 skew, 393216 pool) is the
|
||||
calibration history in the sizing tables below. DSpark was tried in the drafter's
|
||||
place on 2026-09-16 and reverted — see *The DSpark excursion*.
|
||||
It replaced the vllm-radiance MXFP4 replica pair + routing proxy
|
||||
([[frank-medium-radiance-mxfp4]], the record and the PCIe-atomics evidence) in
|
||||
the same day the pair went live. It is also the box's **only** LLM runtime: the
|
||||
replicas, the proxy and their unit files are retired.
|
||||
([[frank-medium-radiance-mxfp4]] holds that record, and the PCIe-atomics
|
||||
evidence) the same day the pair went live — and was itself replaced by
|
||||
radiance TP=2 on 2026-09-20, once the board swap made a collective legal.
|
||||
During its window it was the box's **only** LLM runtime: the replicas, the
|
||||
proxy and their unit files were retired.
|
||||
|
||||
## Why one process and layer split
|
||||
|
||||
|
||||
@@ -1,7 +1,7 @@
|
||||
---
|
||||
type: configuration
|
||||
title: frank-medium vllm-radiance MXFP4 Runbook
|
||||
description: vllm-radiance MXFP4 serving just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq as frank-medium on Dankserver (2x AMD Radeon AI PRO R9700, gfx1201), as two single-GPU replicas behind an adaptive-routing proxy — and the PCIe-atomics reason TP=2 cannot run on this host.
|
||||
description: ACTIVE frank-medium runbook since 2026-09-20 — vllm-radiance MXFP4 serving just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq as ONE tensor-parallel process (TP=2) on Dankserver's two AMD Radeon AI PRO R9700 (gfx1201), with a DFlash2-FP8 drafter. Vision is served by the loaded checkpoint (verified on synthetic images 2026-09-20). The sections below describe the 2026-09-16 two-replica era.
|
||||
tags:
|
||||
- frank-medium
|
||||
- dankserver
|
||||
@@ -13,28 +13,121 @@ tags:
|
||||
- rccl
|
||||
- pcie-atomics
|
||||
- quadlet
|
||||
timestamp: 2026-09-16
|
||||
timestamp: 2026-09-20
|
||||
---
|
||||
|
||||
> **SUPERSEDED 2026-09-16.** This page is the record of the retired
|
||||
> vllm-radiance MXFP4 deployment. The ACTIVE frank-medium is mainline
|
||||
> llama.cpp with the DFlash2 drafter ([[frank-medium-llamacpp-dflash2]]);
|
||||
> the PCIe-atomics evidence in this page remains the reason the runtime is
|
||||
> layer-split llama.cpp rather than tensor-parallel vLLM.
|
||||
|
||||
# frank-medium vllm-radiance MXFP4 Runbook
|
||||
|
||||
The ACTIVE frank-medium deployment (since 2026-09-15/16): the
|
||||
`stilldeadcode/vllm-radiance:0.9.3` image plus the `ggz14/radiance-vllm-mxfp4`
|
||||
patches, serving `just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq` on the two
|
||||
Radeon AI PRO R9700 (gfx1201) cards. It replaced the charlie12345/ROCmFPX
|
||||
llama.cpp deployment ([[frank-medium-rocmfpx]]).
|
||||
The ACTIVE frank-medium deployment: the `stilldeadcode/vllm-radiance:0.9.3`
|
||||
image (tagged locally as `localhost/radiance-vllm:0.9.3`) plus the
|
||||
`ggz14/radiance-vllm-mxfp4` patches, serving
|
||||
`just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq` as **ONE
|
||||
tensor-parallel process across both Radeon AI PRO R9700 (gfx1201)**.
|
||||
It replaced the mainline-llama.cpp layer-split deployment
|
||||
([[frank-medium-llamacpp-dflash2]], SUPERSEDED 2026-09-20); the same image and
|
||||
model had earlier run here as two single-GPU replicas, which is the shape most
|
||||
of this page describes.
|
||||
|
||||
**Deployed shape: two single-GPU replicas behind an adaptive-routing proxy.** The
|
||||
intended shape — one tensor-parallel process across both cards (TP=2) — cannot
|
||||
launch a collective on this host. The reason is a hardware property of the
|
||||
motherboard's PCIe topology, not the image, the RCCL build, or the kernel
|
||||
version. The evidence is in [Why TP=2 cannot run here](#why-tp2-cannot-run-here).
|
||||
**Deployed shape since 2026-09-20: one rootless Quadlet, `--tensor-parallel-size 2`.**
|
||||
The ProArt B850-Creator WiFi NEO board swap put both R9700s on CPU GPP root
|
||||
ports (`00:01.1 → 03:00.0` and `00:01.3 → 07:00.0`), both Gen5 x8, every
|
||||
upstream bridge `AtomicOpsCap: Routing+` and both endpoints `ReqEn+` — the wall
|
||||
documented in [Why TP=2 cannot run here](#why-tp2-cannot-run-here) is gone on this
|
||||
board. The collective that actually runs is vLLM's R4D one-shot allreduce, which
|
||||
does not use RCCL at all (TheRock RCCL has no gfx1201 device code; see
|
||||
[[Dual-R9700-Cooperation]]).
|
||||
|
||||
## Deployed flags (2026-09-20)
|
||||
|
||||
From `dankserver/frank-medium/quadlet/frank-medium.container`:
|
||||
|
||||
| Setting | Value |
|
||||
|---|---|
|
||||
| Model | `just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq` |
|
||||
| Speculative | `--speculative-config {"method":"dflash","model":"/models/draft","num_speculative_tokens":7}` (z-lab DFlash2-FP8) |
|
||||
| Parallelism | `--tensor-parallel-size 2` (`ROCR_VISIBLE_DEVICES=0,1`) |
|
||||
| Context | `--max-model-len 262144` |
|
||||
| Concurrency | `--max-num-seqs 8`, chunk 8192 |
|
||||
| KV | fp8 (`--kv-cache-dtype fp8`), **auto-sized — no `KV_CACHE_MEMORY_MB`** |
|
||||
| Prompt cache | `--enable-prefix-caching` |
|
||||
| GPU fill | `--gpu-memory-utilization 0.98` |
|
||||
| Parsers | `--enable-auto-tool-choice --tool-call-parser qwen3_coder`, `--reasoning-parser qwen3` |
|
||||
| Sampler floor | `--override-sampling-effects` (temp 0.7 / top_p 0.95 / top_k 20) |
|
||||
| Multimodal | `--limit-mm-per-prompt '{"image": 8}'` |
|
||||
| Chat template | `qwen-fixed-v22.3.jinja` (patch dir bind-mounted, `MM_MERGE_ENGINE=AXROW_SEPARATE` for Gemma4 vision) |
|
||||
|
||||
**The KV pool is auto-sized, not pinned.** The arm-R benchmark run used
|
||||
`KV_CACHE_MEMORY_MB=18563072000`, which bought exactly one 262,144-token
|
||||
sequence (262,720 tokens at 70,656 B/token — a profile inflated by the
|
||||
8192-chunk activation peak). Production ships no such pin: the pool came up at
|
||||
**943,581 tokens = 3.59× the full window**, and the boot line read
|
||||
`Available KV cache memory: 18.17 GiB → GPU KV cache size: 943,581 tokens`.
|
||||
Per-token cost in steady state is 19,673 B — fp8 KV across the model's 16
|
||||
full-attention layers plus the drafter's. Container RSS measured 28.4 GiB
|
||||
against the host's 62 GiB and stayed flat through a 4-way burst at 512-token
|
||||
max; host-wide swap 3.3 GiB, unchanged. The 2026-09-16 swap storm was the
|
||||
two-replica vision shape, not this one.
|
||||
|
||||
## Battery results — why this shape won (2026-09-20)
|
||||
|
||||
`scripts/bench-conversation.py`, 10 real-workload prompts, `max_tokens=512`,
|
||||
greedy, salted per run and per arm, service stopped between arms
|
||||
(`scripts/bench-conversation.py --host http://127.0.0.1:7770 --burst 4`):
|
||||
|
||||
| Arm | Shape | Decode (geomean) | TTFT (mean) | Prefill (geomean) | Burst-4 aggregate |
|
||||
|---|---|---|---|---|---|
|
||||
| I | llama.cpp layer split 32:33 + DFlash2 Q4_K_M drafter | 45.7 t/s | 0.72 s | ~530 t/s | 71.7 t/s (26.7/stream) |
|
||||
| L | llama.cpp `--split-mode tensor` 1:1, drafterless | 30.5 t/s | 0.71 s | — | 54.3 t/s |
|
||||
| **R** | **vllm-radiance TP=2 + DFlash2-FP8** | **95.8 t/s** | **0.21 s** | **1493.6 t/s** | **259.4 t/s (65.6/stream)** |
|
||||
|
||||
Re-measured against the production unit after deploy: geomean decode 92.4 t/s,
|
||||
mean TTFT 0.22 s, burst-4 aggregate 248.9 t/s — within noise of arm R.
|
||||
|
||||
Honest caveat: arm R carries MXFP4 (4-bit) where the llama arms carry Q8_0
|
||||
(8-bit), so part of the win is quantization, not tensor parallelism alone. Arm
|
||||
L exists to show that llama.cpp's own TP is not the answer here regardless —
|
||||
it is slower than its own layer split, and cannot carry the drafter at all.
|
||||
|
||||
## Verified capabilities (2026-09-20, against the live endpoint)
|
||||
|
||||
- `/v1/models` → `id: frank-medium`; `/health` → 200.
|
||||
- **Tool calls**: `finish_reason: tool_calls` via `qwen3_coder`. The harness
|
||||
drives a bash tool through this endpoint end to end.
|
||||
- **Reasoning**: `reasoning_effort` accepts `none`/`low`/`medium`/`high`
|
||||
natively (vLLM's `ReasoningEffort`); thinking arrives in the `reasoning`
|
||||
field, NOT llama.cpp's `reasoning_content`. `max_tokens` counts visible
|
||||
content only; with no reasoning budget the model still spends it, so a tiny
|
||||
budget returns `content: null` with `finish_reason: length`.
|
||||
- **Vision works** despite `language_model_only: true` in `config.json` — the
|
||||
Gemma4 merge path answers image requests (~83 prompt tokens for a small
|
||||
synthetic), correctly identifying shape and colour on held-out images.
|
||||
`models.yml` therefore lists this endpoint as `input: [text, image]`.
|
||||
- Drafter is live: accept length 2.9–4.3 tokens per draft, acceptance ~0.60.
|
||||
|
||||
## Rollback
|
||||
|
||||
The llama.cpp layer-split trio is archived in `scripts/.AGENTTRASH/
|
||||
20260920-llama-layersplit/` (`frank-medium.service`, `frank-medium.env`,
|
||||
`serve-frank-medium.sh`) and the runtime + Q8_0 payload remain on disk:
|
||||
|
||||
1. `git revert` the cutover commit.
|
||||
2. `deploy.sh dankserver --service frank-medium` (re-installs and re-enables the root unit).
|
||||
3. Or by hand: restore the three files, move the unit symlink back into `/etc/systemd/system/`, `sudo systemctl daemon-reload && sudo systemctl unmask frank-medium && sudo systemctl start frank-medium`.
|
||||
4. Gate: `curl -s http://192.168.0.235:7770/health` → 200.
|
||||
|
||||
## R4D vs RCCL
|
||||
|
||||
vLLM's R4D one-shot allreduce bypasses RCCL entirely because TheRock RCCL
|
||||
builds lack gfx1201 device code (`invalid device function`). The Ubuntu
|
||||
RCCL 2.27 / `librccl1` 7.1.0 build does load gfx1201 kernels — that is what
|
||||
`scripts/build-llama-rdna4.sh` with `GGML_HIP_RCCL=ON` links
|
||||
(`build-rdna4-rocm714-rccl/`) — so llama.cpp's tensor mode has a real
|
||||
collective too. It is simply the slower shape.
|
||||
|
||||
> **The sections from here down describe the 2026-09-16 → 2026-09-19 era: two
|
||||
> single-GPU replicas behind an adaptive-routing proxy, on the TUF GAMING
|
||||
> B650-E.** The deployed shape is above; this is the evidence trail that made
|
||||
> replicas necessary then and what the 2026-09-20 board swap changed.
|
||||
|
||||
## Architecture
|
||||
|
||||
@@ -71,6 +164,14 @@ host would restart forever without ever becoming healthy.
|
||||
|
||||
## Why TP=2 cannot run here
|
||||
|
||||
> **Resolved 2026-09-20.** Everything below is true of the **TUF GAMING
|
||||
> B650-E WIFI**, whose second card sat behind chipset bridges advertising
|
||||
> `AtomicOpsCap: Routing-`. The ProArt B850-Creator WiFi NEO swap moved both
|
||||
> cards onto CPU GPP root ports (`00:01.1 → 03:00.0`, `00:01.3 → 07:00.0`, Gen5
|
||||
> x8, every hop `Routing+`, both endpoints `ReqEn+`), and TP=2 now boots and
|
||||
> serves. The measurement verdict against the alternatives is in
|
||||
> [[Dual-R9700-Cooperation]].
|
||||
|
||||
### Symptom
|
||||
|
||||
Any two-rank collective dies at launch:
|
||||
@@ -83,14 +184,16 @@ Any two-rank collective dies at launch:
|
||||
With the runtime's own API log (`AMD_LOG_LEVEL=3`) the failing call is explicit:
|
||||
|
||||
```
|
||||
ShaderName : ncclDevKernel_Generic_4(ncclDevKernelArgsStorage<4096ul>)
|
||||
:1:rocvirtual.cpp :4208: Pcie atomics not enabled, hostcall not supported
|
||||
:1:rocvirtual.cpp :4636: AQL dispatch failed!
|
||||
:3:hip_module.cpp :893 : hipLaunchKernel: Returned hipErrorIllegalState
|
||||
[ShaderName : ncclDevKernel_Generic_4(ncclDevKernelArgsStorage<4096ul>)]
|
||||
[ :1:rocvirtual.cpp :4208: Pcie atomics not enabled, hostcall not supported]
|
||||
[ :1:rocvirtual.cpp :4636: AQL dispatch failed!]
|
||||
[ :3:hip_module.cpp :893 : hipLaunchKernel: Returned hipErrorIllegalState]
|
||||
```
|
||||
|
||||
ROCclr refuses to hand the kernel a hostcall buffer when the device has no
|
||||
PCIe atomics (`rocclr/device/rocm/rocvirtual.cpp:3963`), and the launch aborts.
|
||||
PCIe atomics (`pcie_atomics_` is true for the device init, and ROCr passes
|
||||
that buffer only when `pcie_atomics_` is true; otherwise it returns "not
|
||||
supported" and the dispatch fails) — and the launch aborts.
|
||||
`hipErrorIllegalState` is that refusal, not a driver fault.
|
||||
|
||||
### Chain of causation
|
||||
@@ -99,15 +202,15 @@ PCIe atomics (`rocclr/device/rocm/rocvirtual.cpp:3963`), and the launch aborts.
|
||||
2. ROCr passes that buffer only when `pcie_atomics_` is true for the device;
|
||||
otherwise it returns "not supported" and the dispatch fails.
|
||||
3. `pcie_atomics_` is set once, at device init, from the HSA link info between
|
||||
the CPU agent and the GPU (`rocclr/device/rocm/rocdevice.cpp:186`,
|
||||
`Device::checkAtomicSupport`). It is not re-read per process.
|
||||
the CPU agent and the GPU (`Device::checkAtomicSupport`). It is not re-read
|
||||
per process.
|
||||
4. The KFD driver populates that link info from
|
||||
`pci_enable_atomic_ops_to_root()`, which — per the PCIe spec and the
|
||||
kernel's implementation — requires **every upstream bridge** between the
|
||||
device and its root port to advertise `AtomicOp Routing` capability
|
||||
(`drivers/pci/pci.c`).
|
||||
|
||||
### The hardware difference between the two cards
|
||||
The hardware difference between the two cards:
|
||||
|
||||
| Card | PCIe chain (endpoint → root port) | `AtomicOpsCap` on every hop | AtomicOpsCtl |
|
||||
|---|---|---|---|
|
||||
@@ -135,7 +238,7 @@ boot with the write already in place) would change it — and even then the two
|
||||
| P2P hardware | fine: 45.3 GB/s bidirectional copy between the two cards; both cards pass `rocm-bandwidth-test` |
|
||||
| The shipped librccl | a from-source rebuild of `ROCm/rccl` develop (with local fixes for that branch's dangling references) reproduces the same failure at the same call path — the library is not at fault |
|
||||
| vLLM/pynccl | a bare two-rank `torch.distributed` allreduce with no vLLM involved fails the same way; single-rank collectives succeed on both cards (NCCL short-circuits at `world_size=1`, which is why a one-card test proves nothing) |
|
||||
| Transport knobs | `NCCL_PROTO=Simple`, `NCCL_P2P_DISABLE`, `NCCL_SHM_DISABLE`, `NCCL_CUMEM_ENABLE=0`, MSCCL off, channel counts, launch modes, `HSA_ENABLE_SDMA=0` — all identical |
|
||||
| Transport knobs | `NCCL_PROTO=Simple`, `P2P_DISABLE`, `SHM_DISABLE`, `CUMEM_ENABLE=0`, MSCCL off, channel counts, launch modes, `HSA_ENABLE_SDMA=0` — all identical |
|
||||
| Image/support matrix | the same image on a different host with the same kernel (Proxmox VM, Codeberg issue #22) fails byte-identically |
|
||||
|
||||
### What would fix it
|
||||
@@ -145,7 +248,10 @@ boot with the write already in place) would change it — and even then the two
|
||||
true TP=2. On this AM5 platform the CPU provides the x16 slot plus x4 lanes,
|
||||
so a second x16-class card necessarily lands on the chipset path — a
|
||||
platform with more CPU lanes (Threadripper/EPYC) is the structural answer.
|
||||
- Until then, **two replicas** are how both cards serve.
|
||||
- **Done 2026-09-20 — this is exactly what happened.** The ProArt
|
||||
B850-Creator WiFi NEO routes both slots to CPU root ports, the atomics
|
||||
condition is satisfied, and TP=2 is the deployed shape. The Threadripper/EPYC
|
||||
argument was correct about the old board and is now moot.
|
||||
|
||||
## Routing
|
||||
|
||||
@@ -159,7 +265,7 @@ The proxy decides per request, from `proxy/routing.py` (policy) and
|
||||
subagents — spreads across the cards by counter, so both prefill in parallel.
|
||||
|
||||
Measured on this host, 90k-token contexts and 15k-token prompts, `vllm bench`
|
||||
plus a small sequential-turn client:
|
||||
plus a sequential-turn client:
|
||||
|
||||
| strategy | cold 16-request burst | warm-turn TTFT (4 turns × 2 families) |
|
||||
|---|--:|--:|
|
||||
@@ -187,13 +293,16 @@ Both replicas, uniform except GPU index and port:
|
||||
| Knob | Value | Why |
|
||||
|---|---|---|
|
||||
| `TP` | 1 | one card per replica |
|
||||
| `MAXLEN` | 114688 | the largest shape that boots on one card: 262144 needs 9.61 GiB of KV against 5.14 available, and 131072 needs 5.58 against 5.22 |
|
||||
| `KV_MEM` | 0 | passed through literally as `--kv-cache-memory 0`, which the image's patched KV sizing reads as "profile this shape"; the upstream pin (2x7551-32624) is a TP=2/262144 measurement and is invalid here. **Do not blank the variable instead of passing 0**: with the flag omitted vLLM sizes the cache from `GPU_UTIL` alone and the KV allocation dies with `torch.OutOfMemoryError` (measured: 2.16 GiB requested, 116 MiB free, 29.34 GiB already held) |
|
||||
| `MAXLEN` | 114688 | the largest shape that boots on one card: 262144 needs 9.61 GiB of KV against 5.21 GiB available, and 131072 needs 5.58 GiB against 5.22 |
|
||||
| `KV_MEM` | 0 | passed through literally as `--kv-cache-memory 0`, which the image's patched KV sizing reads as "profile this shape"; the upstream pin (2x7551-32624) is a TP=2 measurement and is invalid here. **Do not blank the variable instead of passing 0**: with the flag omitted vLLM sizes the cache from `GPU_UTIL` alone and the KV allocation dies with `torch.OutOfMemoryError` (measured: 2.16 GiB requested, 116 MiB free, 29.34 GiB already held) |
|
||||
| `MAXSEQS` | 4 | matches the per-replica share of the fleet's concurrency |
|
||||
| `CHUNK` | 8192 | prefill chunk (upstream default), **coupled to the context**: the prefill arena grows with the chunk size and shares the device budget with the KV cache. Measured: `CHUNK=16384` leaves 4.22 GiB of KV against the 5.21 GiB that 114688 context needs, and the engine refuses to start (`To serve at least one request with the model's max seq len (114688), 5.21 GiB KV cache is needed, which is larger than the available KV cache memory (4.22 GiB)`). Raising CHUNK therefore means cutting `MAXLEN`; single-card prefill measured ≈2.5k tok/s at 8192 (~15k-token prompt, 5992 ms TTFT), so the trade was not worth taking |
|
||||
| `CHUNK` | 8192 | prefill chunk (upstream default), **coupled to the context**: the prefill arena grows with the chunk size and shares the device budget with the KV cache. Measured: `CHUNK=16384` leaves 4.22 GiB of KV against the 5.21 GiB that 114688 context needs, and the engine refuses to start (`To serve at least one request with the model's max seq len (114688), 5.21 GiB KV cache is needed, which is larger than the available KV cache memory (4.22 GiB)`). Raising CHUNK therefore means cutting `MAXLEN`; single-card prefill measured ≈2.5k tok/s at 8192 (~15k-token prompt, 5992 ms TTFT) |
|
||||
| `GPU_UTIL` | 0.98 | upstream default; the card is dedicated |
|
||||
| `SPEC_METHOD` / `SPEC` | dflash / 7 | the DFlash2-FP8 block-diffusion drafter. Depth 7 is upstream's production default and was re-measured here on a code corpus built from this repo (see [Speculative depth](#speculative-depth)); depth 5 wins batch throughput, 7 wins single-stream latency, and this fleet is interactive |
|
||||
| `SNAP` / `DRAFTER` | `/models/Qwen3.8-27B-Uncensored-MXFP4-awq` / `/models/Qwen3.8-27B-DFlash2-FP8` | the served checkpoint and its drafter |
|
||||
| `SPEC_METHOD` / `SPEC` | dflash / 7 | the z-lab DFlash2-FP8 block-diffusion drafter. Depth 7 is upstream's production default and was re-measured here on a code corpus built from this repo (see [Speculative depth](#speculative-depth)); depth 5 wins batch throughput, 7 wins single-stream latency, and this fleet is interactive |
|
||||
| `spec-draft-device` | `ROCm1` | the drafter rides the card that ends up holding the target's last layers |
|
||||
| `temp` / `top-p` / `top-k` | `0.7` / `0.95` / `20` | the values the vLLM predecessor served on this endpoint |
|
||||
| `reasoning-format` / `reasoning-budget` | `deepseek` / `-1` | trace into `reasoning_content`, unlimited |
|
||||
| `mmproj` | `mmproj-Qwen3.8-27B-Uncensored-f16.gguf` (beside the target) | vision enabled; this build has NO autoload, so the flag is explicit and the serve script refuses to start without the file |
|
||||
|
||||
Mounts are the TP=2 unit's minus nothing: `/models`, the HF cache, the shared
|
||||
compile cache, `/patches` (the pinned radiance clone; its patches are applied
|
||||
@@ -202,7 +311,7 @@ at every container start), `/r4d` (the patched libr4d — the image's shipped
|
||||
the entrypoint/healthcheck/proxy scripts mounted read-only out of
|
||||
`/home/dank/serverAppConfigs`.
|
||||
|
||||
The healthcheck reads `PORT`, so one probe serves the replicas (:7771/:7772),
|
||||
The healthcheck reads `PORT`, so one probe serves the replicas (:7770),
|
||||
the proxy (:7770) and the TP=2 target.
|
||||
|
||||
### Boot persistence, verified by rebooting
|
||||
@@ -216,7 +325,8 @@ sudo reboot # unattended; ~2 min 14 s to graphical.target
|
||||
# ... no manual step at any point ...
|
||||
systemctl --user is-active frank-medium-a frank-medium-b frank-medium-proxy
|
||||
-> active active active (containers up 59 s after boot)
|
||||
curl http://127.0.0.1:7770/health -> both replicas UP ~3.5 min after boot
|
||||
# ...
|
||||
curl http://127.0.0.1:7770/health # both replicas UP ~3.5 min after boot
|
||||
```
|
||||
|
||||
Why it comes back: each unit's `[Install] WantedBy=default.target` — the generator
|
||||
|
||||
@@ -1,13 +1,13 @@
|
||||
# LLM Fleet — Knowledge Base
|
||||
|
||||
Documentation for the LLM inference fleet on the two hosts (2026-09-19):
|
||||
**mainline llama.cpp on Dankserver** serving frank-medium as one layer-split
|
||||
process across both Radeon AI PRO R9700 cards (orcarouter Qwen3.8-27B-
|
||||
Uncensored Q8_0 + the z-lab DFlash2 Q4_K_M drafter, balanced 32:33), and TWO
|
||||
endpoints on Dankclaw — the CIRU v4.4.1 runtime (frank-large, jcbtc
|
||||
Qwen3.8-Flash-CIRU-STRIX-Orca w4b on the gfx1151 iGPU, port 7770) and mainline
|
||||
llama.cpp (frank-small, HauhauCS Gemma4-12B QAT Q4_K_M with MTP and vision on
|
||||
the RTX 3090, port 7771).
|
||||
Documentation for the LLM inference fleet on the two hosts (2026-09-20):
|
||||
**vllm-radiance on Dankserver** serving frank-medium as ONE tensor-parallel
|
||||
process across both Radeon AI PRO R9700 cards (just1moremodel
|
||||
Qwen3.8-27B-Uncensored MXFP4-awq + the z-lab DFlash2-FP8 drafter, TP=2 over the
|
||||
R4D one-shot allreduce), and TWO endpoints on Dankclaw — the CIRU v4.4.1
|
||||
runtime (frank-large, jcbtc Qwen3.8-Flash-CIRU-STRIX-Orca w4b on the gfx1151
|
||||
iGPU, port 7770) and mainline llama.cpp (frank-small, HauhauCS Gemma4-12B QAT
|
||||
Q4_K_M with MTP and vision on the RTX 3090, port 7771).
|
||||
|
||||
> 2026-08-22: the former `vllmScripts` repository is merged into the
|
||||
> serverAppConfigs repository. This knowledge base moved from the old repo's
|
||||
@@ -21,10 +21,10 @@ graph LR
|
||||
subgraph Dankserver[AMD Dankserver - Ubuntu 26.04]
|
||||
A[2x Radeon AI PRO R9700
|
||||
gfx1201, 32 GiB each]
|
||||
B2[llama.cpp, layer split
|
||||
B2[vllm-radiance 0.9.3, TP=2
|
||||
frank-medium, port 7770]
|
||||
C[orcarouter Qwen3.8-27B Q8_0 + DFlash2 Q4_K_M
|
||||
ONE process on both cards, 2x163840 ctx, balanced 32:33]
|
||||
C[Qwen3.8-27B MXFP4-awq + DFlash2-FP8
|
||||
ONE process on both cards, 262144 ctx, up to 8 seqs]
|
||||
A --> B2 --> C
|
||||
end
|
||||
|
||||
@@ -46,6 +46,7 @@ frank-small, port 7771]
|
||||
CLI[LLM Clients / API] -->|HTTP :7770| B2
|
||||
CLI -->|HTTP :7770| E
|
||||
CLI -->|HTTP :7771| G
|
||||
```
|
||||
|
||||
## Servers
|
||||
|
||||
@@ -56,21 +57,25 @@ frank-small, port 7771]
|
||||
|
||||
## Active Deployment
|
||||
|
||||
- [[frank-medium-llamacpp-dflash2]] — the ACTIVE frank-medium runbook since
|
||||
2026-09-16 (model swapped 2026-09-19): mainline llama.cpp serving
|
||||
orcarouter/Qwen3.8-27B-Uncensored-GGUF at Q8_0 with the z-lab DFlash2 Q4_K_M
|
||||
drafter on Dankserver's 2x Radeon AI PRO R9700, as ONE layer-split process
|
||||
behind :7770 (tensor parallelism was impossible on the old board's second
|
||||
card; the 2026-09-20 swap cleared that gate but the runtime still ships
|
||||
RCCL-less and the collective is unprobed), balanced 32:33 over two 163,840-token
|
||||
slots
|
||||
- [[frank-medium-radiance-mxfp4]] — RETIRED 2026-09-16, one day after it was
|
||||
deployed: the vllm-radiance MXFP4 stack (stilldeadcode/vllm-radiance:0.9.3 +
|
||||
- [[frank-medium-radiance-mxfp4]] — the ACTIVE frank-medium runbook since
|
||||
2026-09-20: vllm-radiance MXFP4 (stilldeadcode/vllm-radiance:0.9.3 +
|
||||
ggz14/radiance-vllm-mxfp4 patches) serving
|
||||
just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq as two single-GPU replicas
|
||||
behind a round-robin proxy; kept as the record of the deployment it
|
||||
succeeded, plus the PCIe-atomics evidence that ruled out TP=2 on the previous
|
||||
board — that gate cleared on 2026-09-20, see [[Dual-R9700-Cooperation]]
|
||||
just1moremodel/Qwen3.8-27B-Uncensored-MXFP4-awq with the z-lab DFlash2-FP8
|
||||
drafter as ONE tensor-parallel process (TP=2) behind :7770 — the three-arm
|
||||
battery that chose it, the auto-sized 943,581-token KV pool, verified
|
||||
capabilities (tools, reasoning, vision), and the record of the 2026-09-16
|
||||
two-replica era
|
||||
- [[frank-medium-llamacpp-dflash2]] — RETIRED 2026-09-16 → 2026-09-20: mainline
|
||||
llama.cpp serving orcarouter/Qwen3.8-27B-Uncensored-GGUF at Q8_0 with the
|
||||
DFlash2 Q4_K_M drafter as ONE layer-split process (32:33 over two
|
||||
163,840-token slots); superseded by the radiance TP=2 cutover, and the
|
||||
rollback target. Keeps the batch/ubatch optimum curve and the per-layer split
|
||||
arithmetic
|
||||
- [[Dual-R9700-Cooperation]] — whether the two R9700s can serve ONE model: the
|
||||
RCCL/PCIe-atomics wall on the old TUF GAMING B650-E, the 2026-09-20 ProArt
|
||||
B850-Creator x8/x8 swap that cleared it, and the measured three-way verdict
|
||||
(radiance TP=2 95.8 t/s decode vs layer split 45.7 vs llama tensor mode 30.5)
|
||||
that put TP=2 in production
|
||||
- [[frank-medium-rocmfpx]] — PRIOR frank-medium deployment (llama.cpp era,
|
||||
2026-09-15): charlie12345/ROCmFPX serving
|
||||
kingjones777/Qwen3.8-27B-ROCmFPX-Q8_0-AGENT-GGUF; kept as the record of the
|
||||
|
||||
Reference in new issue
Block a user