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

Descrição

# 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

```sh
# 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.