#3555 [Engine] Windows + RX 7900 XTX (gfx1100): engines 0.1.45.3 / 0.1.46.1 never become ready - "verify: stream create failed", HIP 719 in hipBLASLt init, ~1.9 GiB VRAM free, 0.1.43.1 works

Description

### Desktop app, Terminal or source?

Source - a checkout of the tree, launched as `python serve/server.py --engine strata --config strata-unsloth-ud-iq4_xs.json --port 8084`. Not a git clone, no app and no Terminal; the run config is hand-written. Updating means replacing `engine/` and the `serve/` tree by hand (which is how I tested 0.1.45.3 and 0.1.46.1).

### What happened

0.1.43.1 works. Upgrading to 0.2.1 (**engine 0.1.46.1**, package `strata-engine-0.1.46.1-windows-x86_64-hip-rdna3`) the engine never becomes ready. It builds the whole startup - arena, session, native head, drafter, expert cache, prompt path - and then dies in the verify step:

```
strata trace:  1950 MiB free after the head and the prompt path
strata serve: verify: stream create failed
Hip error: 'unspecified launch failure'(719) at C:/home/runner/_work/rockrel/rockrel/rocm-libraries/projects/hipblaslt/library/src/amd_detail/hipblaslt.cpp:468
```

The same happens with **engine 0.1.45.3** (`strata-engine-0.1.45.3-windows-x86_64-hip-rdna3` from 0.2.0-beta), so the regression is not 0.1.46-only. With 0.1.45.3 the line before the failure differs, the failure does not:

```
prefill gemm: cannot open tuning file; using hipBLASEx
strata trace:  1968 MiB free after the head and the prompt path
strata serve: verify: stream create failed
Hip error: 'unspecified launch failure'(719) at .../hipblaslt.cpp:468
```

### Smallest steps or command that reproduce it

1. Windows source checkout of 0.1.43.1 works with this config: `--max-context 131072 --kv int8 --spec 4 --spec-min-p 0.7 --resident-budget-gib 44 --vram-reserve-mib 1856 --prefill auto --expert-cache auto`.
2. Replace `engine/` with `strata-engine-0.1.46.1-windows-x86_64-hip-rdna3` (or `0.1.45.3`), keep the same config and model.
3. Start via `serve/server.py`.
4. It loads everything (44 GiB pinned complement, expert cache 5702 slots / 12.85 GiB) and then exits with the two lines above.

Nothing else on the GPU; a fresh start each time; 10 attempts, all identical.

### Reproduction status

Reproduced locally - 10 fresh engine starts across the two engine versions (rows below).

### Strata version

Strata 0.2.1 - engine **0.1.46.1** (BUILD.json: hip, `hipblaslt_version 100500`, archs include gfx1100, rocm 10.2.0a20260930).
Also tested: engine **0.1.45.3** from 0.2.0-beta - fails identically.
**Working**: engine **0.1.43.1** with the same model, pack and config.

### Graphics card, memory and system

Radeon RX 7900 XTX 24 GB (**gfx1100**) - driver **32.0.31041.1004** (2026-08-17) - 64 GB DDR5-5600 (CL38) - Windows 11 Pro 25H2 build **26200.9711** - Ryzen 7 9700X - SAPPHIRE PULSE B850MP WIFI.

### Model, size and settings

Unsloth **UD-IQ4_XS** of Qwen3.8-Flash-Next, `--native` GGUF shards + pack (`--pack`), `--expert-profile`, `--mtp`, 131072 context, pictures off, `--resident-budget-gib 44`. Serving is fully resident on 0.1.43.1: `resident 42.6 GiB of experts in RAM` and **0 blob reads** from the file tier.

### Logs

Startup trace with `STRATA_TRACE=1` (engine 0.1.46.1, production config):

```
strata trace: 22169 MiB free after the CUDA context and the weight arena
strata trace: 22169 MiB free after the pack weights
strata trace: 17201 MiB free after the session (KV, GDN state, PLE history)
strata trace: 15633 MiB free after the native head
strata trace: 15620 MiB free after the weights, the session and the drafter
strata trace:  2003 MiB free after opening the expert cache
strata trace:  2010 MiB free after the profile fill
strata trace:  2010 MiB free after the R4 hit path
strata trace:  2010 MiB free after the expert cache and the graphs
strata trace:  2011 MiB free after before the prompt path init
prefill gemm: hipBLASLt tuning enabled (32 rows, gfx1100, version 100500)
strata trace:  1950 MiB free after the head and the prompt path
strata hc: CUDA0: the hyper-connection read runs as staged (the norm per token and stream, the down projection's activations staged ahead by cp.async); checked bit for bit against the plain read on this card (STRATA_HC_SPLIT=1 or 0 for the earlier ones)
strata serve: verify: stream create failed
Hip error: 'unspecified launch failure'(719) at .../hipblaslt.cpp:468
```

With `STRATA_HIPBLASLT_VERBOSE=1`, right before the failure:

```
prefill gemm: hipBLASLt tuning enabled (32 rows, gfx1100, version 100500)
strata serve: verify: stream create failed
prefill gemm: hipBLASLt summary launches=0 fallbacks=0 unique_fallback_shapes=0
Hip error: 'unspecified launch failure'(719) at .../hipblaslt.cpp:468
```

`launches=0 fallbacks=0` - no GEMM was ever launched, so this is not a kernel for the wrong arch, it is the handle/stream creation.

### It is not memory

* `STRATA_TRACE=1` reports **1950-1968 MiB free** at the moment of the failure; the expert cache is sized normally (5702 slots, 12.85 GiB).
* Shrinking the prompt path changes nothing: with `--vram-reserve-mib 4096 --prefill 1024` the cache drops to 5175 slots / 10.63 GiB and the prompt path borrows only 0.86 GiB - same failure, ~1.97 GiB free.
* `HIP_LAUNCH_BLOCKING=1` does not move the failure earlier, so it is not a sticky error left by a previous asynchronous kernel.

### What I tried (each row a fresh engine start, everything else unchanged)

| # | Engine | Change | Result |
|---|---|---|---|
| 1 | 0.1.43.1 | production config | **works** (42.6 GiB resident, 0 blob reads, 50-67 tok/s) |
| 2 | 0.1.46.1 | production config | fail, same call |
| 3 | 0.1.46.1 | `--vram-reserve-mib 3072 --prefill 2048` | fail |
| 4 | 0.1.46.1 | `--vram-reserve-mib 4096 --prefill 1024` (cache 5175 slots / 10.63 GiB, borrow 0.86 GiB) | fail |
| 5 | 0.1.46.1 | `STRATA_PF_GEMM=0` | fail |
| 6 | 0.1.46.1 | `STRATA_HIP_QUEUE_PRIME=0` + `STRATA_PREFILL_STREAM_AHEAD=0` | fail |
| 7 | 0.1.46.1 | `STRATA_ARENA_PIN_GIB=16` (the fix that worked in #2013) | fail |
| 8 | 0.1.46.1 | `STRATA_HIPBLASLT_TUNING` unset | fail |
| 9 | 0.1.45.3 | production config | fail, same call |
| 10 | 0.1.45.3 / 0.1.46.1 | `HIP_LAUNCH_BLOCKING=1` + `STRATA_TRACE=1` | fail, error does not move |

### Why I think it is the prefill GEMM path

* 0.1.43.1 never prints a `prefill gemm:` line with this configuration and starts fine; 0.1.45.3 and 0.1.46.1 print it and die in `hipblaslt.cpp:468` with `launches=0`.
* The outcome is identical with the tuning table present, absent, or the package's own `engine/tuning/gfx1100-hipblaslt-100500.txt`.
* gfx1100 is supported by the shipped library: both engine packages carry `rocm/bin/hipblaslt/library/gfx1100` (98 files, incl. `Kernels.so-000-gfx1100.hsaco` and `hipblasltTransform_gfx1100.hsaco`), a gfx1100 tuning table, and the same `libhipblaslt.dll 1.5.0.0` as 0.1.43.1.

### Possibly related issues

* **#2013** - Windows Strix Halo (gfx1151), same `unspecified launch failure` class during serve init; `STRATA_ARENA_PIN_GIB=16` fixed it there. No effect here (row 7).
* **#3263** - 0.2.0-beta engine 0.1.45.3 on a real RX 7900 XTX (gfx1100) **on Linux** runs, same speed as 0.1.44.1. So gfx1100 itself is fine; the Windows build is where it breaks.
* **#2960**, **#3463** - other Windows + gfx1100 reports.
* 0.1.45 release notes: "Not run: Windows hardware (...), Intel Arc, the RTX 5070."

### Questions

1. Is there an environment switch that turns the hipBLASEx / hipBLASLt prefill path **completely** off? `STRATA_PF_GEMM=0` did not change the failure.
2. Is there a diagnostic build, or `STRATA_TRACE=2`, that prints which hipBLASLt call returns 719 (handle, stream, workspace query)?
3. Does 0.1.46.2+ touch this initialisation on Windows? I can test any engine package you point me at - the machine is free in the evenings and I keep a full rollback of 0.1.43.1.

Happy to run any diagnostic you want on this exact machine.