#1386 2x RX 7900 XTX (gfx1100, layer split): STRATA_PF_GEMM is worth ~7% prefill and is off by default

Descrição

# 2× RX 7900 XTX (gfx1100, layer split): `STRATA_PF_GEMM` is worth ~7% prefill and is off by default

**TL;DR**

- On gfx1100 with a 125B-class MoE (IQ3_S native experts on a 2-card layer split), 0.1.40 improves two of the three
  directions a lot: **cold start 166 s → 95 s (−43%, #699 read-ahead)** and **decode 46.7 → 56.6 tok/s
  (+21%, shared-expert fork default-off, #826/#816)**. Both claims verified on AMD hardware here.
- But 0.1.40 measured **9–11% slower prefill** than 0.1.39 for me until I found that the new **`STRATA_PF_GEMM` is
  opt-in and off by default**. That switch is worth **+7.4%** over the 0.1.40 default, which brings prefill back to
  **within 3% of 0.1.39**. Worth a release-note line, or a default-on for gfx11 if the numbers hold elsewhere.
- I calibrated `gfx1100-hipblaslt-100201.txt` with `tools/hip/tune_hipblaslt` for this ROCm (hipBLASLt 1.2.1 →
  version number `100201`); the shipped tables are 100100/100200/100202/100401/100500, so 1.2.1 is refused by the
  version check. The table loads, `hip_prefill_hipblaslt_gemm` reports `fallbacks=0`, and it is **end-to-end neutral
  on my model** (my dense GEMMs are WMMA-dominated) although single shapes gain up to **5.46×** over hipBLASEx in
  the tuner. Happy to send it as a PR (like #745) if you want that library version covered.
- **Open question**: with a self-calibrated table, `hip_prefill_hcd_exact_parity` stays skipped with
  `the table sends no tested T to solution 1176 / 1177`, i.e. the HCD-exact kernel never engages. Expected?

## Environment

| | |
|---|---|
| GPUs | 2× Radeon RX 7900 XTX (gfx1100, 24 GB each), **layer split**: layers 0-24 (card 0), 25-47 (card 1), one hand-off per window |
| Host | Dell Precision T7910, 2× Xeon E5-2683 v4 (32 threads), 125 GB RAM, Ubuntu; 4 GPUs total (the other two serve unrelated work) |
| ROCm | **7.2.0** (`7.2.26015-fc0010cf6a`), hipBLASLt **1.2.1 → version number 100201**, ROCM-SMI-LIB 7.8.0 |
| Engine | self-built HIP/gfx1100 from the **v0.1.40.1** tree (`BUILD.json`: version 0.1.40, backend hip, archs gfx1100, `vision: none`), pinned llama.cpp `3cf0325` |
| Server | `serve/` from v0.1.40.1 (the engine is the 0.1.40 one, as your release note says) |
| Model | Qwen3.8-Flash-Next GSQ-RCO **IQ3_S**, 48 layers, pack built with `tools/iq_pack.py` (native experts); the PLE shard is passed as `--ple-gguf` |
| Flags | `--max-context 262144 --kv int8 --kv-resident 32768 --expert-cache auto --prefill auto --spec 4 --spec-min-p 0.5 --layer-split auto` |
| Env | `STRATA_WMMA_GEMM=1` (+ the switches discussed below) |
| Loaded state | 19,084 expert slots, **78% of the experts resident**, 592 MiB VRAM free, prompt chunk auto 8192 / 96-slot ring, MTP draft head 212.9 MiB |

All measurements below are the engine's own `strata serve: prompt …` lines, same pack, same config, same script per
pair; prompts are greedy, `reasoning_effort=none`, generation capped at 8 tokens for the ladder cells (so the numbers
are prefill throughput, not end-to-end).

## 1. Cold start: read-ahead (#699) works on RDNA3

166 s → 95 s (**−43%**) to HTTP-ready, both arms after evicting the page cache by reading an unrelated 80 GB file
(the box has no passwordless sudo, so no `drop_caches`; reading the other model's weights is my substitute, same
method for both arms). Warm restart of the same config: 60–79 s. Nothing here claims your RTX 5090 number
(920 s → 70 s); this is just the AMD-side confirmation.

## 2. Decode: shared-expert fork default-off matches your claim

| arm | decode, 400 generated tokens ×3 (median) |
|---|---:|
| 0.1.39 (fork default on there) | 46.7 tok/s |
| **0.1.40 default (fork off on HIP)** | **56.6 tok/s (+21%)** |
| 0.1.40 with `STRATA_SH_STREAM=1` | 50.2 tok/s |

The fork is **decode-only** on this box: across three arms (default, `STRATA_SH_STREAM=1`, `STRATA_PF_BK64=0`) the
same 152,068-token cold read measured 1,688.4 / 1,688.4 / 1,688.4 tok/s. So defaulting it off on HIP is right here.

## 3. `STRATA_PF_GEMM` (new, opt-in) is worth ~7% prefill on gfx1100 — the main finding

Same pack, same config; the only difference is the env var. Five cold-read cells plus the small-prompt warmup:

| cell (prompt tokens) | 0.1.40 default | **`STRATA_PF_GEMM=1`** | Δ |
|---|---:|---:|---:|
| 8,468 (warmup) | 842.2 | 849.4 | +0.9% |
| 33,668 | 1,453.4 | **1,560.9** | **+7.4%** |
| 67,268 | 1,381.8 | **1,481.6** | **+7.2%** |
| 122,148 | 1,492.9 | **1,603.4** | **+7.4%** |
| 201,668 | 1,451.9 | **1,550.5** | **+6.8%** |
| 152,068 (cold, no reuse) | 1,688.4 | **1,822.7** | **+8.0%** |
| decode 400 tok ×3 (median) | 56.6 | 55.3 | −2.3% (inside our ±7%) |
| cancel latency (134 k, 5 s disconnect) | 13.7 s | 12.3 s | better |

Against 0.1.39 (same pack/config/script) the picture becomes: prefill 1,603.0 / 1,660.8 / 1,891.7 for 33k / 122k / the
152 k cold read, i.e. with `PF_GEMM=1` 0.1.40 is only **−2.6% … −3.6%** behind instead of −9% … −11%.

**Ask:** is off-by-default deliberate on gfx11 (does it regress on some card/quant you tested?), or is gfx1100 simply
untested for it? From `src/prefill/gemm.cu` the S23 path runs *before* the WMMA path in `Gemm::f16` when
`pf_on && T >= pf_min_t` (default 0), and it returns false for shapes it does not take — on our IQ3_S expert shapes it
looks like a straight win. If you want, I can re-run the pair a few more times and post the raw engine lines.

## 4. Negative / neutral results (so nobody re-tests them)

| switch | result on this box |
|---|---|
| `STRATA_SH_STREAM=1` | decode 56.6 → 50.2 (−11%), prefill unchanged |
| `STRATA_PF_BK64=0` | no change (prefill identical cell by cell) |
| `STRATA_IQ3S_MT1=1` | prefill slightly *worse* (p33k 1,560.9 → 1,438.2), decode unchanged |
| `STRATA_PF_SWITCH_MIN_T=4096` | no change (our ladder cells are all ≥ 33 k, so a threshold only affects small T) |
| `STRATA_PF_PAD=1` | +0.1…0.8% on all five cells — inside our ±1% noise; its only large number (+6.2% at an 8.5 k prompt) was falsified by the parallel arm *without* it (909.4 vs 901.7 tok/s), i.e. noise |
| `STRATA_PA_WMMA` (env) | no effect — 0.1.40 turned it into an internal arch macro, so an env gate is ignored (fine, just noting that the old env var is now silent) |

## 5. A calibrated hipBLASLt table for 100201 (offer to contribute)

```sh
cmake --build build-hip --target tune_hipblaslt
CASES=$(awk 'NR>2 {printf " --case %s,%s,%s,%s,%s", $1,$5,$2,$3,$4}' tools/hip/gfx1100-hipblaslt-100200.txt)
./build-hip/tune_hipblaslt $CASES --tuning-out tools/hip/gfx1100-hipblaslt-100201.txt
```

- 26 shapes, **68 s**, 28-line file, header `STRATA_HIPBLASLT_TUNING_V1 gfx1100 100201`.
- `STRATA_HIPBLASLT_TUNING=<file> ./build-hip/hip_prefill_hipblaslt_gemm` → **rc 0, 4 cases PASS, `fallbacks=0`**;
  the engine logs `prefill gemm: hipBLASLt tuning enabled (26 rows, gfx1100, version 100201)`.
- Micro: e.g. `f16 T=8192 N=12288 K=2560`: hipBLASEx 35.43 ms → hipBLASLt solution 1147 **6.48 ms (5.46×)**.
- **End-to-end on my model: no change** (p33k 1,450.6 vs 1,453.4; 152 k cold 1,687.6 vs 1,688.4; decode unchanged) —
  my dense prefill work goes through the WMMA path, so hipBLASLt only catches leftovers.
- `hip_prefill_hcd_exact_parity` stays **SKIP**: `the table sends no tested T to solution 1176 / 1177`. So a
  self-calibrated table silently leaves the HCD-exact kernel off. Is that intended (those two ids only being valid for
  the library build the shipped table was measured on), or should the tuner be steered to prefer 1176/1177 for the
  N 320 / K 10240 shape so the exact kernel can engage?

I can open a PR that adds `tools/hip/gfx1100-hipblaslt-100201.txt` (setup picks `<arch>-hipblaslt-<version>.txt` by
name, so the file is the whole change) and a line in the README's shipped-table list — say the word if that ROCm
version matters to you.

## 6. Variance, in case it is useful for the docs

- Cold reads at a large T (33 k–200 k): **±1%** run-to-run — usable for sub-percent A/B decisions.
- A 8.5 k prompt: **±7%** (two arms of the *same* configuration: 849.4 vs 909.4 tok/s) — cannot judge sub-percent.
- Decode 400 tokens: **±7%** (also seen on 0.1.38 → 0.1.39, where a first-sample "+8.6%" turned into a dead heat on
  A/B/A).
- Recommendation that follows: A/B a switch on the big-T cold-read cells and require the whole row to move, not one
  cell.

## 7. What I did *not* test

Single card, 4-GPU, other quantisations (Q2_0 / IQ2_XS / Q4_0), other architectures, `--vision` (not built here),
`--parallel` / `--batch-mtp`, Windows, and the #848 resident-RAM-mode soak beyond what the default split mode does.
No GPU faults, stalls, errors or 500s in ~3 h of 0.1.40 with repeated 200 k-token prompts, 134 k cancellations and one
long-context recall probe (4/4 markers, pp 1,841.6 tok/s, TTFT 116.1 s).