#3204 [Engine] 0.2.0-beta: prompt reading ~25% slower on a two-GPU HIP layer split (gfx1151 + gfx1201) since f3be9743; STRATA_PREFILL_WARM=0 restores it

Description

### Area

[Engine]  (speed, kernels, models, the server and API)

### Desktop app, Terminal or source?

Source (START-HERE.bat, setup.sh or a build)

### What happened

On a layer split across a Radeon 8060S (gfx1151, Strix Halo iGPU) and a Radeon AI PRO R9700 (gfx1201), 0.2.0-beta reads
long prompts about 25% slower than 0.1.45: 1,250 against 1,710 tok/s at 64K. That is about the speed of the 8060S
alone, so the second GPU stops helping. I expected 0.1.45's speed.

A `git bisect` between v0.1.45 and v0.2.0-beta lands on **f3be9743**, "prefill: warm the first BF16 and FP16 hipBLAS
products on a thread at start (STRATA_PREFILL_WARM=0 skips)". Setting **`STRATA_PREFILL_WARM=0`** on 0.2.0-beta brings
prompt reading back to 0.1.45's speed. Decode is not affected.

The 8060S alone (R9700 hidden with `ROCR_VISIBLE_DEVICES`) is not affected either: 0.1.45 and 0.2.0-beta both read 64K
at about 1,220-1,230 tok/s there. So this looks specific to the multi-GPU split.

Prompt reading in tok/s, two runs each, one request at a time:

| Version (two-GPU split) | 64K | 128K | 211K | Decode (tok/s) |
|---|---|---|---|---|
| 0.1.43.1 | 1,694 / 1,682 | 1,604 / 1,563 | 1,484 / 1,463 | 45-52 |
| 0.1.45 | 1,718 / 1,705 | | | 47-50 |
| 0.2.0-beta | 1,262 / 1,244 | 1,215 / 1,208 | 1,162 / 1,163 | 44-50 |
| 0.2.0-beta + `STRATA_PREFILL_WARM=0` | 1,720 / 1,709 | 1,658 / 1,618 | 1,559 / 1,548 | 43-51 |

With the warm-up off, 0.2.0-beta is a little faster than 0.1.43.1 at 128K and 211K (+4-5%). A second 64K pair on the
same engine, with prompts it had not seen, read at 1,709 / 1,677.

Ruled out on 0.2.0-beta:

- **The new Linux file-tier defaults.** With `STRATA_COLD_OVERLAP=0 STRATA_IO_PREFETCH=0 STRATA_IO_PF_GATE=0`, 64K
  averaged 1,261 tok/s against 1,244 with them on (two fresh engines each).
- **The known "R9700 + MTP: about half of fresh engines read 10-23% slower" limit.** Six fresh 0.2.0-beta engines all
  read 64K at 1,193-1,275 tok/s, so the drop is not random from engine to engine.

Workaround: set `STRATA_PREFILL_WARM=0`. The only cost is the one-time hipBLAS load (1-2.5 s) on the first request.

Not tested:

- Windows.
- Other two-GPU combinations (two gfx12 cards, or NVIDIA).
- Whether the slowdown depends on the order of the GPUs (here the 8060S is GPU 0).

### Smallest steps or command that reproduce it

1. Build one engine for both GPUs, out of tree:
   `-DSTRATA_ENABLE_HIP=ON -DSTRATA_PREFILL_MMQ=ON -DCMAKE_HIP_ARCHITECTURES="gfx1151;gfx1201"`, HIP flags
   `-DSTRATA_W_LB=8`, ggml at the pinned 3cf03257.
2. Start `serve/server.py` with a two-GPU layer split: 8060S first, `layer_split` 30 (layers 0-29 on the 8060S, 30-47 on
   the R9700), 4 slots. The engine arguments:

   ```
   --pack <packs>/unsloth-ud-iq4_xs --native <models>/Qwen3.8-Flash-Next-UD-IQ4_XS-00001-of-00003.gguf
   --expert-profile data/expert-profile.bin --expert-cache auto --spec 4 --spec-min-p 0.5 --mtp <mtp>/rt
   --max-context 262144 --kv int8 --kv-resident 32768 --mmap-experts --mtp-window 8192 --no-prefill-borrow
   --prefill 4096 --conversation-cache-mib 24576 --conversation-cache-slots 24 --trim-stage-weights
   --batch-groups 1 --batch-mtp --vision
   ```

   The environment:

   ```
   STRATA_HIPBLASLT_TUNING=tools/hip/gfx1151-hipblaslt-100401.txt
   STRATA_PF_FUSED=1  STRATA_PF_GEMM=1  STRATA_PA_FAST=1  STRATA_HIP_WMMA=1
   STRATA_SELECT_WMMA=1  STRATA_PF_FUSED_KQ=1  STRATA_PREFILL_STREAM_MIN=64
   ```

   The experts are 100% resident: the engine reports "layer split: 100% of the experts resident".
3. Send a 16K warm-up request, then a 64K prompt: Python standard-library source cut to 64K tokens, plus a short question.
   Temperature 0, 600 tokens out, one request at a time.
4. Compare the prompt reading speed with and without `STRATA_PREFILL_WARM=0`, each on a fresh engine. Prompt tok/s is
   prompt tokens divided by the time to first token, measured on the client. These figures match the engine's own
   "strata serve: prompt N tokens ... read in X ms" lines.

Every build in this report carries one small local patch, needed for this split on every version tested. Without it the
gfx1201 stage exits on the first prompt with "prefill gemm: a padded X (ldx 6208, K 6144) needs STRATA_PF_PAD's path".
The cause: `STRATA_PF_PAD` is a gfx1151 default, applied process-wide from GPU 0, but the padded GEMM exists only for
gfx11. The patch asks the current device instead (8 lines, 4 files: `Gemm::native`, the two `ld_y` / `ld_a` sites in
`Prefill::run_impl`, and `strata_pf_pad_device()` in `wmma_gemm`). Without the patch, `STRATA_PF_PAD=0` is the
workaround, at 2-3% slower prompts. Happy to send the diff.

**Bisect.** Each step: build the commit (with the PF_PAD patch above), start a fresh engine with the launch above, then
read 64K twice. A step is good if its best run is at least 1,450 tok/s. There were no ambiguous steps and none were
skipped.

| Commit | Date | Subject | 64K (tok/s) | Verdict |
|---|---|---|---|---|
| 417cbab4 | 10-17 | packaging: the engine package carries strata-vision when the build has it | 1,696 / 1,617 | good |
| 83cde2ca | 10-18 | Merge f-151-b70 (c765c5f87) into f-split-int2 | 1,266 / 1,239 | bad |
| 80e6d91a | 10-17 | Merge commit 'a6759787c' into f-split-int | 1,722 / 1,666 | good |
| 83aa59fd | 10-18 | setup: say why the page waits while the bundled engine unpacks | 1,714 / 1,619 | good |
| af3770c6 | 10-18 | Merge f-148-amd (2053388dc, without the IQ2 fused default) into f-split-int2 | 1,722 / 1,567 | good |
| 0e58a652 | 10-18 | tests: iq_multi_parity under STRATA_KQ_SPLIT=1 ... | 1,269 / 1,238 | bad |
| 3980a5ce | 10-18 | HIP: virtual memory ranges ... so --kv-grow can run on ROCm 6+ | 1,273 / 1,245 | bad |
| aef32e22 | 10-18 | prefill timing: the hyper-connection write-back on its own line | 1,260 / 1,253 | bad |
| f3be9743 | 10-18 | prefill: warm the first BF16 and FP16 hipBLAS products on a thread at start | 1,248 / 1,237 | **bad (first bad)** |
| 79ac3164 | 10-18 | gfx12: STRATA_HIP_WMMA=auto ... | 1,719 / 1,708 | good (last good) |

### Reproduction status

Reproduced locally

### Strata version

Strata 0.2.0-beta (tag v0.2.0-beta, engine 0.1.45.3), built from source. First bad commit: f3be9743. Last good: 79ac3164.

### Graphics card, memory and system

Radeon 8060S (gfx1151, AMD Ryzen AI MAX+ 395, 1 GB carve-out + 120 GB GTT) + Radeon AI PRO R9700 32 GB (gfx1201, as an eGPU; the engine's PCIe probe: 7.2 GB/s host to device) · 128 GB RAM · Ubuntu 26.04.1, kernel 7.0.0-38, ROCm 7.14.1 built for both architectures, hipBLASLt 1.4.1 (version 100401)

### Model, size and settings

Qwen3.8-Flash-Next UD-IQ4_XS (unsloth, native pack) · 262K context · int8 KV · MTP drafter · 4 slots · pictures on (CPU reader)

### Logs

With the warm-up on, every start prints this sequence:

```
prefill gemm: hipBLASLt tuning enabled (90 rows, gfx1151, version 100401)
strata prefill: warmed the BF16 and FP16 GEMM libraries on GPU 0 in 2396 ms (STRATA_PREFILL_WARM=0 skips)
prefill gemm: tuning architecture mismatch: file=gfx1151 runtime=gfx1201; using hipBLASEx
prefill gemm: tuning architecture mismatch: file=gfx1151 runtime=gfx1201; using hipBLASEx
strata prefill: warmed the BF16 and FP16 GEMM libraries on GPU 1 in 1048 ms (STRATA_PREFILL_WARM=0 skips)
prefill gemm: hipBLASLt tuning enabled (90 rows, gfx1151, version 100401)
```

`STRATA_PREFILL_TIMING=1` at 64K shows where the time goes. The two stages report differently: one reads in 4,096-token
chunks, the other in blocks of about 32K tokens (two per 64K prompt). These figures compare 0.1.43.1 with 0.2.0-beta, so
other changes between those versions are in them too:

| Per ~32K block | 0.1.43.1 | 0.2.0-beta (warm-up on) |
|---|---|---|
| GPU timeline | 17.2-17.5 s | 24.5 s |
| hc read | 2.7 s | 9.0 s |
| hc write (its own line since aef32e22) | n/a | 1.1 s |
| gemm gate/up | 5.3 s | 5.3 s |
| gdn | 1.7 s | 2.0 s |

With 0.2.0-beta, the 4,096-token chunk lines dropped from about 1.55 s to 0.49 s each. The ~32K-block stage became the
bottleneck instead, and most of its growth is in `hc read`.

On the 0.2.0-beta engine run with timing, the first 64K request after the start also spent 72.8 s in "ple" (525 tok/s);
its second request read at 1,207 tok/s. The other fresh 0.2.0-beta engines read their first 64K at their usual ~1,250.
Not checked whether this is related.

I did not find the mechanism. As far as I can see, the hipBLASLt tuning state is per `Gemm` (`create_hipblaslt_state`),
so the obvious shared-state suspect does not seem to apply. Where the warm-up is called:

- `warm_gemm_libraries()` in `src/prefill/prefill.cpp`.
- A `std::thread` in `src/program/generate.cpp`, started before the expert arena loads.
- The end of `Prefill::init` (once per device).