WIP: gather-MMVQ pilot (review surface, not complete)

[ddvnguyen]

DRAFT review surface for the compact-gather pilot, exactly as it stands. Not a completion claim. Closes nothing.

Measurement context (ARMs)

• ARM 000 OFF single-GPU 18.19 tok/s
• ARM 001 OFF default split 14.39 tok/s
• ARM 002 OFF single-GPU graphs-disabled 13.13 tok/s (denominator for any syncing build)
• ARM 003 naive residency 0/4/8 layers = 18.19/19.23/20.55 tok/s, slope +0.295/layer
• ARM 004 pending (VRAM-matched naive control ~3.4GB, stock llama.cpp)

Scope

• ggml/src/ggml-cuda/ggml-cuda.cu: dispatch-time compact-gather pilot (+347)
• src/llama-context.cpp / .h: pin plumbing (+189/+2)
• src/llama-graph.cpp: hook site (+1)

Known pending: Finding-1 rework (device-side slot_of, in-kernel row-address select; no per-call host readback/sync). Code fixes come after review, not before.

[ddvnguyen]

HEADLINE FINDING (architect ruling, banked §19 / commit c5e078d97): the gather hook is installed on the MMVQ dispatch branch, but decode dispatches to MMQ (engagement counters: mmvq=0, mmq=2192, hit=0, miss=0 — zero gather engagements; fused path bypasses the thunk hook). The mechanism as wired is a decode-time no-op; armed tok/s 17.80/17.85 vs 17.7476 same-session anchor confirms stock behaviour. Any review of this surface should treat hook-dispatch placement (MMQ vs MMVQ) as the primary defect. ARM 006 armed-traffic measurement on this build is VOID (no-op), and the related PCIe premise test has been withdrawn as archaeology on a superseded mechanism — see track bank §18/§19 for the CPU-bound pivot (compute/backend split per expert, not weight-address selection).

[ddvnguyen]

Architect review — gather-MMVQ pilot

Reviewed at a744d8019. Verdict: do not continue this direction. The defects below are real and
fixable, but fixing them does not make the design win, so the recommendation is to close this PR in favour
of per-expert backend selection. Details and the arithmetic behind that call are in
docs/leader-handoff-state.md §22.

The number that governs this review

The expert-ranking prize is 5.23 ms/token (12.44 ms/token per 100% of expert work relocated CPU→GPU,
× h=0.4207). The mechanism fires 144 times per token (48 layers × 3 active expert-matmul sites).

36.3 µs per invocation is break-even. 12.1 µs for a 3× margin.

Every finding below is stated against that budget.

---

B1 — Design: misses are shipped to the GPU instead of computed where their weights live (BLOCKING)

hydra_gather_mmvq, ggml-cuda.cu:2108-2113: every non-pinned expert is copied host→device on every
token, into a ggml_cuda_pool_alloc that dies with the call — no reuse across tokens.

Cost, using the measured 0.597 MiB per expert per site:

10 used × (1 − h=0.4207) = 5.79 misses × 0.597 MiB × 3 sites × 48 layers ≈ 498 MiB/token
one 0.597 MiB H2D at 8 GB/s ≈ 78 µs  → 2.1× over budget for a SINGLE expert
5.79 misses per invocation  ≈ 450 µs → 12× over budget

The design pays for GPU compute by shipping weights across the bus, and on this rig the bus is the
expensive thing. This is the same failure the dual-construction had (2k experts instead of k), reduced by
h rather than cured by it.

Solution — not a patch, an architecture change. Never move the weights; move the activations. Hits
compute on the GPU against the resident slab; misses compute on the CPU, where their weights already
are; combine the two partial outputs. This needs a fork-native op on both backends that skips experts it
does not own — CUDA side compacting the promoted subset on-device (prefix-sum over the pin mask, no host
readback), CPU side looping over the miss subset only. That also removes B2 at the root rather than
patching it. Tracked as ARM 008 (one-site spike, measured against the 36.3 µs bar) before any epic.

B2 — Per-call full stream sync (BLOCKING)

ggml-cuda.cu:2070-2071:

CUDA_CHECK(cudaMemcpyAsync(ids_host.data(), ids->data, n_ids*sizeof(int32_t), cudaMemcpyDeviceToHost, stream));
CUDA_CHECK(cudaStreamSynchronize(stream));

~144 host syncs per token, at 10–50 µs each — at or over break-even before any useful work happens.
Credit where due: ggml_cuda_mul_mat_id_needs_sync was correctly extended at :2163-2175 to veto CUDA
graphs on engaged nodes, so this is not a correctness bug. It is a cost bug, and it is fatal to the budget.

Solution: determine the used-expert set on-device. The ids are already on the GPU; a small
prefix-sum/compaction kernel produces the compact mapping without ever touching the host. Then the graph
veto at :2163-2175 can be dropped too, restoring graph capture.

M1 — The sync is uncounted and untimed on the path that actually runs (MEASUREMENT)

hydra_n_sync / hydra_sync_us are incremented only at :2288 and :2316, both in the old fallback
path. The gather's own sync at :2071 is invisible. We therefore have no instrumentation for the
per-invocation budget on the path under test — which is precisely the number that decides whether expert
ranking is viable at all.

Solution: time every invocation of the engaged path and report p50/p99 µs per site, not just totals.

M2 — Two different mechanisms share one set of hit counters (MEASUREMENT)

Two hydra paths coexist with different semantics:

• gather (:2046) — copies misses H2D, counts hydra_hit_exp once per unique expert in the compacted set
• fallback (:2339-2370) — swaps w_ptr per expert slice, no copies, counts per expert slice with tokens_per_expert > 0

Both increment the same counters, so any run engaging both reports an h that is a mixture of two
incompatible definitions. Neither matches the paper's h, which is over lookups (k=10 per token), not
over unique rows. Both also gate on ne12 == 1, so the reported h is decode-only — correct for the thesis,
but it must be stated rather than inferred.

Solution: one counter set per path, tagged; count lookups not unique rows; label the decode-only scope
in the dump.

M3 — Pin-file loading is silent on every failure (MEASUREMENT, most important for trust)

hydra_init_once_fn, :1938-1970:

• fopen failure → silent. A typo in HYDRA_PIN_FILE yields an unarmed run that looks armed.
• Any line failing sscanf(line, "L %d", &il) → continue, silent. A ranker emitting a different format
  loads zero pins, hydra_active() returns false, and the mechanism does nothing — with no error.
• No summary of what was loaded: no layer count, no experts/layer, no total.

This makes "the mechanism ran and did nothing" indistinguishable from "the mechanism never engaged," which
is exactly the failure mode this track has already been burned by.

Solution: hard-fail on fopen error and on any unparseable line; emit a one-line load summary
(layers, experts/layer, total pins, bytes that will be allocated) at init; and expose an engagement counter
that a test can assert is non-zero before any effect is measured.

R1 — The VRAM slab is invisible to the fit-check and never freed (RESOURCE)

:2022 allocates with raw cudaMalloc outside ggml's buffer accounting, and no cudaFree exists anywhere
for hydra_pin_store::dev_ptr.

At N=38: 38 × 0.597 MiB × 48 layers × 3 sites ≈ 3.2 GB allocated where the fit calculation cannot see it.
That directly compounds #139 — the engine computes a fit, boots, and then dies when the real allocation
lands.

Solution: allocate through a ggml backend buffer so the memory is accounted, and free it with the
backend context.

R2 — Duplicate and out-of-range pins over-allocate silently (RESOURCE)

:2020 sizes the slab as n_slot = hydra_pins[il].size(), but the fill loop at :2024-2027 skips
duplicates and ids >= n_expert. Those slots are allocated and never used.

Solution: deduplicate and bounds-check at load time, and size the slab from the validated set.

N1 — hydra_active() runs std::call_once on the hot path (NIT)

Called at :2217, :2235, :2243, :2251 and from needs_sync, i.e. several times per site per token.
Cheap after completion, but not free, and this is a 36 µs budget.

Solution: resolve once at backend init into a plain bool.

---

Recommendation

Close this PR. File B1/B2 as superseded-by-design rather than to-fix (already open as #140, #141), and land
M1–M3 + R1–R2 separately regardless of which mechanism wins — they are measurement and resource
correctness, and without M3 in particular no number this track produces is trustworthy.

Next step is ARM 008: a one-site spike of per-expert backend selection, instrumented to report wall time
per invocation against the 36.3 µs break-even and the 12.1 µs 3× bar. Above 36.3 µs, the conversion claim
is refuted on this hardware and I will say so.