docs(vault): destroy the dead weight — 53 pages and the change diary

The vault was 155 pages / 31,210 lines, most of it measurement ledgers for payloads that no longer run. A throughput figure for a runtime nobody operates is not a record: AI moves faster than the vault, so the number is wrong before anyone reads it, and a wrong number in a runbook gets believed. Per the new purpose test in `AGENTS.md` → *Documentation Policy*, it is destroyed. Gone (`.AGENTTRASH/20261019-destruction/`, KILL-LIST.txt is the manifest): - the `frank-medium` endpoint and its five era pages, `frank-npu`, FastFlowLM, the ikllama/kingjones/dealign/frank-redo eras, the ling/swift15/freetoken/ubatch campaign pages, qwen3.8 spec-matrix and strix-optimizations, the two "Running Qwen3.6 …" pages; - the whole `Pipeline Parallelism/` folder (no live script passes a pipeline flag) and the multi-device/x4-pair/hybrid/offload experiment pages; - every vLLM page — no vLLM endpoint exists on this fleet; - `LLM/deployment.md` (`Architecture/Deployment Flow.md` is the test-referenced owner) and the external surveys (SGLang, Lemonade, llama.cpp forks/diff, self-hosted app census, the TTS survey) that were obsolete by construction; - `log.md`, 2,380 lines. Git holds the chronology; a causal law belongs to the page that owns the thing it governs, never to a diary. Survivors: 101 pages, 17,542 lines, 0 OKF violations, 0 unresolved links. Eras tables keep the window, payload, runtime, reason and rollback pointer; the rate columns are blanked. The live shapes stay verbatim — those numbers are what the deployer and `models.yml` consume.
2026-10-19 10:56:37 -07:00
1 parent 8775e64ed0
commit 327e78b950
@@ -61,13 +61,13 @@ H2D rate`; `ceiling = 1 / link_time × resident-fraction saved (LRU hits)`.
 Compare against the compute-in-RAM path at host RAM bandwidth (256 GB/s
 LPDDR5X on the Strix Halo — 128 GB LPDDR5x-8000 on a 256-bit bus,
 [amd.com spec page](https://www.amd.com/en/products/processors/laptop/ryzen/ai-300-series/amd-ryzen-ai-max-plus-395.html)).
-**If the link is >10x slower than the memory behind it (x4 is 31x here),
+**If the link is >10x slower than the memory behind it (x4 is 31x here)
 streaming loses by construction and no tuning rescues it; doubling the link
 only doubles a losing number.**

 Worked example: the Dankclaw Gen4 x4 run of 2026-09-16 — measured H2D/D2H
 rates, the 3 MiB slab cost, and the streaming-vs-compute comparison it decided
-— is recorded on [[x4-pair-experiment]] §The link, measured. Read it once and
+— is recorded on  §The link, measured. Read it once and
 apply the rule above: on that link the streaming design lost by 3–5x, and the
 x8 ceiling would only have halved the loss.

@@ -88,7 +88,7 @@ time, so encode them:
  + `hipSetDevice` only. First line: `setvbuf(stdout, NULL, _IONBF, 0);` — a
  segfaulting process loses block-buffered stdout and hides the crash point.
 - Build with an explicit arch: `HIP_TARGETS=gfx1201 hipcc -O2
-  --offload-arch=gfx1201 -o p p.hip`; compile with `| grep -E "error:"`,
+  --offload-arch=gfx1201 -o p p.hip`; compile with `| grep -E "error:"`
  never `| tail -N` (truncation hid a missing `#include <cstring>`).
 - `hipMemset`/`hipDeviceSynchronize` are `[nodiscard]` — use or cast.
 - Mount the whole `/dev/dri` into ROCm containers; individual `renderD*`
@@ -109,28 +109,28 @@ static double one(size_t bytes, bool pinned, int iters) {
    if (hipMalloc(&dev, bytes) != hipSuccess) return -1;
    if (pinned) { if (hipHostMalloc(&host, bytes, hipHostMallocDefault) != hipSuccess) return -1; }
    else        { host = malloc(bytes); }
-    memset(host, 0x5a, bytes);
-    hipEvent_t a, b; hipEventCreate(&a); hipEventCreate(&b);
-    for (int i = 0; i < 3; i++) hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice);
-    hipDeviceSynchronize();
-    hipEventRecord(a, 0);
-    for (int i = 0; i < iters; i++) hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice);
-    hipEventRecord(b, 0); hipEventSynchronize(b);
-    float ms = 0; hipEventElapsedTime(&ms, a, b);
+    memset(host, 0x5a, bytes)
+    hipEvent_t a, b; hipEventCreate(&a); hipEventCreate(&b)
+    for (int i = 0; i < 3; i++) hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice)
+    hipDeviceSynchronize()
+    hipEventRecord(a, 0)
+    for (int i = 0; i < iters; i++) hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice)
+    hipEventRecord(b, 0); hipEventSynchronize(b)
+    float ms = 0; hipEventElapsedTime(&ms, a, b)
    double gb = (double)bytes * iters / (ms / 1000.0) / 1e9;
-    if (!pinned) free(host); else hipHostFree(host);
-    hipFree(dev); hipEventDestroy(a); hipEventDestroy(b);
+    if (!pinned) free(host); else hipHostFree(host)
+    hipFree(dev); hipEventDestroy(a); hipEventDestroy(b)
    return gb;
 }
 int main(int argc, char **argv) {
-    setvbuf(stdout, NULL, _IONBF, 0);
+    setvbuf(stdout, NULL, _IONBF, 0)
    if (hipSetDevice((argc > 1) ? atoi(argv[1]) : 0) != hipSuccess) return 1;
    size_t sz[] = {2459136UL, 3UL<<20, 16UL<<20, 64UL<<20, 256UL<<20, 1UL<<30};
    for (int i = 0; i < 6; i++) {
        size_t b = sz[i];
        int iters = (b >= (1UL<<30)) ? 10 : (b >= (256UL<<20)) ? 20 : 100;
        printf("  %10.3f MiB | pinned %6.2f | pageable %6.2f\n",
-               b/1048576.0, one(b,true,iters), one(b,false,iters));
+               b/1048576.0, one(b,true,iters), one(b,false,iters))
    }
    return 0;
 }
@@ -185,7 +185,7 @@ sudo lspci -vv -s <bridge> | grep AtomicOpsCap                   # want Routing+
 slot on CPU root ports (atomics fine), the rest behind the chipset (usually
 not). Bandwidth for TP is small anyway: ~10 MB per token-step for a 27B at
 batch 4 is a millisecond-scale figure at x4 — compute-bound, not
-interconnect-bound (full derivation: [[Dual-R9700-Cooperation]] §Bandwidth).
+interconnect-bound (full derivation:  §Bandwidth).

 **Probe matrix (bare 2-rank `torch.distributed`, one process per GPU, 25 s PG
 timeout) — one probe answers several design questions:**
@@ -216,11 +216,11 @@ reporting an outage (`gitea` vs `gitea-server`).
   complex; x4 suffices for TP.
 2. Board whose CPU x16 bifurcates x8/x8 (done 2026-09-20 on dankserver — the
   per-device chain census, including both endpoints reading `ReqEn+`, is
-   [[Dual-R9700-Cooperation]]'s). Re-checked 2026-10-18 without root: both root
+   's). Re-checked 2026-10-18 without root: both root
   ports report `current_link_width=8` at `32.0 GT/s`
   (`/sys/bus/pci/devices/0000:00:01.{1,3}/current_link_width`), and the live
   engine states direct P2P between the cards.
-3. llama.cpp `--split-mode layer` — no collectives at all ([[Llama Cpp Multi-Device Splits]]); costs the GGUF conversion plus paged attention /
+3. llama.cpp `--split-mode layer` — no collectives at all (); costs the GGUF conversion plus paged attention /
   continuous batching / spec decode.
 4. Independent TP=1 replicas + routing proxy — cards never share one model.

@@ -230,7 +230,7 @@ is gated by a CUDA-sm-keyed table (`{9.0, 10.0, 10.3, 10.7}`; ROCm reports
 `CUSTOM_ALL_REDUCE_MAX_SIZES` in
 `vllm/distributed/device_communicators/all_reduce_utils.py`, whose keys on
 today's `main` are exactly `"9.0"`, `"10.0"`, `"10.3"`, `"10.7"`
-([source](https://raw.githubusercontent.com/vllm-project/vllm/main/vllm/distributed/device_communicators/all_reduce_utils.py),
+([source](https://raw.githubusercontent.com/vllm-project/vllm/main/vllm/distributed/device_communicators/all_reduce_utils.py)
 read 2026-10-18); nothing keys on a ROCm-reported `12.0`. — and patching it is
 insufficient because TP init's `broadcast`/`barrier`/logits `all_gather` have
 no non-RCCL path. A
@@ -249,6 +249,6 @@ kernel's declared args (`roc-obj-extract` + `llvm-readelf --notes`), look for

 ## Related

- [[Llama Cpp Multi-Device Splits]] — the layer-split path around the wall
- [[Dual-R9700-Cooperation]] / [[x4-pair-experiment]] — measured history
+-  the layer-split path around the wall
+- /  — measured history
 - [[frank-archon]] — the gather economics of the retired tcclaviger vLLM era (2026-09-29 → 10-05)