RX 7900 XTX (gfx1100): 925 -> 2,503 tok/s prompt read (+171%) - the method, and the 26% cliff below full lend coverage

Start with docs/AMD_HIP.md and the FileExpertSource startup message described in the issue; the requested work is to warn when resident coverage is incomplete and document STRATA_RESIDENT_HEADROOM_GIB. The issue also proposes making resident-set sizing deterministic or retrying it, which needs investigation of the existing pacing in #1250. Check the startup warning against full and incomplete coverage, and verify the documentation explains ROCm-version retuning.

gfx1100 (RX 7900 XTX): 100% lend coverage is most of the prefill story, and STRATA_RESIDENT_HEADROOM_GIB is undocumented

Measurements of the prompt path on a single RX 7900 XTX, with what helped, and the dead ends so nobody repeats
them. No request attached: one observation you may want to act on, plus the method behind a ~2.5K tok/s prompt read.

Everything below was measured on v0.1.40.2, which needs no local patch on either ROCm — #1180 covers
both of the clang-22 problems I had worked around on 0.1.40 (NW_LB=1 for clang >= 22, and the missing
__syncthreads() after put(raw, (s + 1) & 1) in native_w11_kernel, now at moe_fused_iq.cu:828 with the
comment "the writes of every thread are in LDS before the next stage reads them (#1180)").

Setup

Ubuntu 24.04 in a privileged LXC (8 vCPU, 50 GB RAM, no swap), RX 7900 XTX 24 GB (gfx1100). Engine v0.1.40.2
built with -DSTRATA_ENABLE_HIP=ON -DSTRATA_PREFILL_MMQ=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100. Model: Swift 1.5
IQ3_XXS (native pack, 48 layers, IQ2_XXS/IQ2_XS/IQ3_XXS/IQ3_S/IQ2_S gate-up + Q2_0/IQ4_NL down).

Engine arguments: --expert-cache auto --prefill auto:32768 --spec 4 --mtp <rt> --max-context 524288 --kv int8 --kv-resident 65536 --resident-experts --vision --vram-reserve-mib 700, with
STRATA_RESIDENT_HEADROOM_GIB set, and STRATA_HIPBLASLT_TUNING pointing at the table that matches the
installed hipBLASLt.

ROCm versions: the method is not version-specific, the numbers are

I report both because they differ, and because I expect most readers are on the wheels setup.py pins:

	ROCm	hipBLASLt / table	25K	74K	150K
setup.py default	7.10.0a20251120	1.2.0 / gfx1100-hipblaslt-100200.txt	2,409	2,158	1,914
what I run now	10.2.0a20261003 (TheRock nightly)	1.5.0 / gfx1100-hipblaslt-100500.txt	2,744	2,559	2,339

• The subject of this issue — the lend coverage, wait copy, the file reads — is engine-side and reproduces
  on both. The good/bad tables below are from the 10.2 run, but the 7.10 run swings the same way (I saw 1,518 /
  1,646 / 1,668 tok/s at 74K on restarts that landed short).
• The 7.10 → 10.2 move needed STRATA_RESIDENT_HEADROOM_GIB 8 → 7, because the newer runtime keeps roughly
  1 GiB more host RAM. That is exactly the kind of thing item 3 below is about.
• One caveat if you are on 7.10: the 10.2 long-prompt numbers above came with the matching 100500 table. On
  any ROCm the engine refuses a table whose version does not match the library and prefills on plain hipBLAS,
  which cost about 45% here (926 vs 1,687 tok/s on a 131K prompt) — keep the pairing.

For reference, the same box/model with the 0.1.40 defaults (no --resident-experts, no fused gate/up,
--prefill auto, pageable expert cache) read a 97,488-token prompt at 925 tok/s.

The method, by measured contribution

lever	effect
fused int8 WMMA gate/up + down (STRATA_PF_FUSED=1, STRATA_PF_GEMM=1, STRATA_PF_SWITCH_MIN_T=4096)	925 → 1,092 tok/s; removes the 16-expert host grouping that was 37.3% of the stage sum
--resident-experts with a page-locked complement	1,092 → 2,503 tok/s; wait copy 45,604 ms (52.2%) → 24 ms (0.1%)
--prefill auto:32768	+23% at 74K (fewer chunks, less expert re-streaming)
the matching hipBLASLt table	see the caveat above: without it the engine drops to plain hipBLAS

The observation: the resident set is decided by MemAvailable at start, and below 100% lend coverage the prompt read silently drops ~26%

FileExpertSource prints the number that decides everything:

FileExpertSource: 6201 of the prompt path's 6201 lendable slots keep their experts in RAM too
FileExpertSource: 5093 of the prompt path's 6202 lendable slots keep their experts in RAM too
                 (the others are read from the file when lent: not enough RAM for them)

When that line is not N of N, the uncovered experts are read from the 76 GB pack during the prompt, and
wait copy — which otherwise sits at 0.1% — swallows the run. Same 74K prompt, same binary, only the startup
differed:

	resident	coverage	file blobs	wait copy	74K prompt
good start	33.24 GiB	6201 / 6201	0	134–485 ms (0.2%)	2,583 tok/s
bad start	31.47 GiB	5093 / 6202	11,090	37,019 ms (36.8%)	1,810 tok/s

The stage breakdown of the bad run is unambiguous — every other stage is unchanged:

stage (ms)	good	bad
wait copy	485 (1.7%)	11,552 (26.2%)
gemm gate/up	6,741	7,692
combine	2,553	5,494
qsa attn / select / gdn / hc read	3,244 / 2,988 / 2,460 / 2,588	3,402 / 3,126 / 2,458 / 2,688
KV streaming	79.00% hit VRAM, 83.4 MiB from RAM	78.97%, 83.5 MiB — identical

Why it is length-dependent: a short prompt only routes through part of the experts, so the uncovered slots are
rarely touched (25K reads the same either way: 2,443 vs 2,409 tok/s). A 74K+ prompt routes through nearly all of
them, so the misses repeat. That is the whole "short prompts look fine, long prompts collapse" pattern.

Why it varies at all: the complement is sized from MemAvailable at start. On this 50 GB box the resident set
landed anywhere between 31.5 and 33.2 GiB across restarts, i.e. 75–100% coverage, with nothing else changed.
The only knob is STRATA_RESIDENT_HEADROOM_GIB (8 for the 7.10 wheels, 7 after moving to 10.2). It is not in
--help, not in docs/, and searching the issues for HEADROOM_GIB returns 0 hits, so I expect most people
never find it: they just get a slower long prompt.

What I would ask for

1. Warn loudly when the line is not N of N — one WARNING with the measured cost ("this prompt path will
   read K experts from the file; expect roughly X% less prompt speed") would have saved me a day of looking at
   wait copy. Right now it is one INFO line among ~40 at start.
2. Make the resident set deterministic, or retry when it lands below full coverage. #1250 already paces
   the pinning; the same pacing could re-run the sizing (or hold back a little RAM) until the complement covers
   the prompt path, since the outcome swings ±26% on a coin flip.
3. Document STRATA_RESIDENT_HEADROOM_GIB in docs/AMD_HIP.md, including that it needs re-tuning when the
   ROCm version changes (7.10 → 10.2 needed 8 → 7 here). A one-line note "check the FileExpertSource line
   after a restart" would cover it.

Dead ends (measured, so nobody repeats them)

• KV placement does not matter. --kv-resident 20,480 / 65,536 / 131,072 with 100% coverage: 25K
  2,481 / 2,453 / 2,454, 74K 2,553 / 2,577 / 2,567, 150K 2,365 / 2,351 / 2,347 tok/s, and the KV streaming
  stats are identical (97.7–98.0% hit VRAM, ~250 MiB from RAM). Moving the KV slice out of VRAM buys +331
  expert slots (0.53 GiB) and no speed. Note --kv-resident 0 is the sentinel for all cells in VRAM, not
  "none" — it costs ~6 GiB of free VRAM and 11% of decode.
• A bigger prefill chunk does not help. --prefill 65536 / 131072 vs auto:32768: 25K 2,454 / 2,460 vs
  2,465, 74K 2,571 / 2,568 vs 2,585, 150K 2,351 / 2,348 vs 2,363. (Also, generate.cpp:1672 caps auto at
  32,768, so auto:32768 is already the maximum.)
• A self-calibrated hipBLASLt 1.5.0 small-T table does nothing (only relevant if you are on a 1.5.0 build).
  I took the geometry set of the new gfx1100-hipblaslt-100401.txt (45 rows, small-T buckets) and re-ran
  tune_hipblaslt against hipBLASLt 1.5.0 on this card: 45 rows, 12 of them small-T, header gfx1100 100500.
  A/B with a warm-up: 512 tokens 794 vs 696 (the 696 had a 605 outlier; the other sample was 787), 1K 1,140 vs
  1,146, 2K 1,842 vs 1,843, 74K 2,528 vs 2,539 — i.e. flat. The reason is visible in the diff: the tuner picks
  the same solution ids the heuristic already used (bf16 10240 2560 10240 is 1036 at T=128…6144, exactly
  what the old table has at 4096/8192). So the +32% you measured for the 100401 table looks specific to that
  hipBLASLt version's old table, and does not transfer to 1.5.0. Short prompts here are also dominated by fixed
  per-request cost, not the GEMMs: 512 tokens reads at ~700–790 tok/s, 4.2K at ~2,400.
• --spec above 4 is flat. With full coverage, spec 3/4/5/6: window 22.5/24.9/27.0/27.9 ms, tokens per
  window 2.00/2.26/2.29/2.37, decode 87.9/91.4/89.8/91.4 tok/s. More drafts cost per step and buy tokens per
  step at about the same rate, so it saturates at 4.
• STRATA_SPEC_GUMBEL=1 under greedy decoding: 70.1 vs 72.2 tok/s (no gain; it is a sampled-decoding
  feature, and it does change the output).
• pin=N / strata_prefix works and is worth documenting harder. With {"strata_prefix": {"messages": 1}}
  on a 100K-token document, tools/research_run.py goes from 101 s (median first token 14.8 s) to 19 s
  (0.81 s) — with no --conversation-cache-mib at all. It does stop working the moment another request is
  interleaved (even a 31-token one), which is where the conversation cache is needed.

Reproduction

# build (either ROCm; no patch needed on 0.1.40.2)
cmake -S . -B build-hip -G Ninja -DCMAKE_BUILD_TYPE=Release -DSTRATA_ENABLE_HIP=ON \
  -DSTRATA_PREFILL_MMQ=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100
ninja -C build-hip strata

# the check that matters, after the first request of a start
grep -o "FileExpertSource: .* of the prompt path's .*" strata-*.log | tail -1
# 6201 of the prompt path's 6201   -> good; anything else -> the prompt path reads the pack

Then read the same 74K prompt twice, once with the coverage full and once with it short, and diff the
wait copy share. On this box the difference is 0.2% vs 36.8%, and 2,583 vs 1,810 tok/s.