#3480 [Engine] gfx1200 (RX 9060 XT): prompt speed depends on the draft side stream, 0.2.1 queue priming (-39% with capped HW queues), the 96-slot prompt ring, and SH_STREAM hangs

Descrição

### Area
Engine (HIP prompt path, gfx1200)

### Desktop app, Terminal or source?
Source build in Docker (ROCm 7.2.3 userspace, hipBLASLt 1.2.2), `serve/server.py` with `strata --serve`.

### What happened
On an RX 9060 XT (gfx1200) four prompt-path settings decide a large part of the prompt speed, and three of the defaults
are the slow side on this card. Every number is a fresh engine per arm, interleaved with the baseline on the same box,
the run config otherwise identical (IQ2_XS native pack, experts in place with `--resident-experts`,
`--kv q4_0 --max-context 524288 --batch 2 --batch-mtp`, MTP on). Answers were unchanged in every arm unless noted.

**1. The draft side stream is needed on gfx1200 (since #1240).** With the HIP draft priority on (the gfx12 default),
the prompt path is ~35% faster when `STRATA_MTP_SIDE_STREAM=1` creates the unused draft side stream. 0.1.44.1, 4K prompt,
median of 3: default 650.6 tok/s, `STRATA_MTP_SIDE_STREAM=1` **874.8**, `STRATA_MTP_HIP_PRIO=0` 596.3, both off 652.1,
`GPU_MAX_HW_QUEUES=4` / `=8` 597.7 / 596.6. A bisect with the priority injected into the older commits gives
`ac52c2c7` (#1240) as the first slow commit (82f46a8c + prio 845.0, 3b499f17 + prio 847.0, ac52c2c7 + prio 632.8, later
commits 632-652). `STRATA_PREFILL_TIMING` puts the difference in the per-chunk "wait copy". #1240's card was a gfx1201.

**2. 0.2.1's queue priming costs 39% here.** We also set `GPU_MAX_HW_QUEUES=2` (neutral before 0.2.1). Three fresh starts each:

| engine | 4K prompt tok/s | 100K-token document tok/s |
|---|---:|---:|
| 0.2.0 + the settings above | 1221-1238 | 1380 |
| 0.2.1 default (priming on) | 681-753 | 994 |
| 0.2.1, `STRATA_HIP_QUEUE_PRIME=0` | 1188-1238 | 1386 |
| 0.2.1, priming on, `GPU_MAX_HW_QUEUES=4` | 899-1234 (one slow start of three) | not run |

Decode, two-slot batching and short chat turns were equal in all arms. Our reading (not verified in the runtime):
the two primed normal-priority queues take both hardware queues, so the high-priority draft stream and the side stream of
point 1 share them. Capping the queues is our own setting, so this may only hit people who do the same; the priming could
skip itself (or prime "nh") when `GPU_MAX_HW_QUEUES` <= 2, or when the draft priority is on.

**3. The 96-slot prompt ring is slow for the fused IQ2_XS experts.** `ring_slots()` (prefill.cpp) returns 96 whenever
`qsa_prompt_attn_matrix_cores_requested()` is true; on gfx12 that is the default even with `--kv q4_0`, where the int8
WMMA kernel never runs. With `STRATA_PF_FUSED=1` on IQ2_XS, 0.2.0:

| `STRATA_PREFILL_RING` | 4K prompt | 100K document | 1.9K-token follow-up at 98K cached |
|---|---:|---:|---:|
| unset (96) | 949-962 | 1169-1179 | 553-567 |
| 384 | 1164-1185 | 1269-1275 | 610-629 |
| 1024 | 1147-1172 | 1275 | 619-629 |

The 96 rule was measured on a 9070 XT with the pinned path; here (fused experts, resident RAM tier) it goes the other way.

**4. `STRATA_SH_STREAM=1` hangs on gfx1200**: `strata serve: verify: timed out at layer 1; its GPU waits were released
and the GPU finished (#267)` on the first decode after start, every time. It is off by default on HIP; `setup.sh --calibrate`
tries it.

Smaller notes: the shipped `tools/hip/gfx1200-hipblaslt-100202.txt` has T buckets 256 / 512 / 2560 only, so 8192-token
chunks use the 2560 rows (e.g. bf16 N=2560 K=12288: 14.2 ms, its row retuned at T=8192: 8.9 ms); a retuned table gave +1-2%
end to end. `STRATA_PA_FAST=1` on gfx1200: kernel 1.12x but its parity check fails (error 1e-3 of the output scale), no
end-to-end gain - fine as opt-in.

The final configuration (side stream on, priming off, ring 384) is measured with `strata bench run` in #3481: 4K 1219.5 tok/s, 16K 1507.8, decode 41.6-45.9, sanity 12/12.

### Reproduction status
Reproduced locally, repeatedly (fresh engines, interleaved arms). Not tested on gfx1201 or on Windows.

### Strata version
0.2.1 (engine 0.1.46.1, upstream 9ae93f7b) for point 2; 0.2.0 for point 3; 0.1.44.1 and the bisect for point 1;
0.2.0 for point 4. Runs with our ten local commits (github.com/shenmintao/Strata branch nas/gfx1200-021); none of them
touches the code paths above.

### Graphics card, memory and system
AMD Radeon RX 9060 XT 16 GB (gfx1200) on PCIe 4.0 x8; i7-12700H (no AVX-512); 64 GB DDR4-3200; fnOS (Debian 12 based,
kernel 6.18) host, Docker with ROCm 7.2.3 userspace.

Related: #3481 (standard benchmark of this configuration), #3482 (draft: opt-in Q4_0-KV WMMA prompt attention).

Happy to run other variants on this card. (Prepared with an AI assistant; every number comes from the logged runs.)