#2837 Windows Strix Halo (gfx1151, Ryzen AI Max+ 395): hipMemGetInfo reports only the carve-out's free space, and the gfx1151 arch defaults never apply

Description

# Windows Strix Halo (gfx1151, Ryzen AI Max+ 395): hipMemGetInfo reports only the carve-out's free space, and the gfx1151 arch defaults never apply

## Hardware and software

- AMD Ryzen AI Max+ 395 w/ Radeon 8060S (Strix Halo, `gfx1151`, wave32, 40 CUs), Windows 11
  (build 26100), Adrenalin driver 32.0.31041.1004, 96 GB LPDDR5X (BIOS carve-out settable).
- Strata v0.1.44.1 (also reproduced on main 296b59d), ready-made Windows HIP engine
  (`strata-windows-x64-hip.zip`, ROCm 10.2.0a20260930), and a self-built engine from the same tree
  (`tools/hip/build_windows.bat`, `STRATA_HIP_ARCHS=gfx1151`).
- Model: ISTA-DASLab GSQ-RCO IQ3_S, 200K or 32K context, int8 KV, MTP on.

## Finding 1: `device_free_bytes()` is not unified-memory aware on Windows

`cudaMemGetInfo` on this APU reports the free space of the BIOS's dedicated carve-out (what the
desktop has not taken of it), not what the unified memory can host. The Linux-only correction in
`src/core/expert_cache.cpp` (`#if defined(__linux__)` → MemAvailable minus
`STRATA_UMA_HEADROOM_GIB`) does not apply, so the auto expert cache sizes from a tiny figure.

Measured free figure (after the session's allocations), by BIOS carve-out, 32K context unless said:

| carve-out | free figure | what happens |
|---|---|---|
| 0.5 GB | ~0 | auto refuses; forced `--expert-cache N` shrinks to 0 slots |
| 16 GB | 0.37 GiB | auto refuses: "needs at least 241 slots (611 MiB), about 757 MiB more than this card has free" |
| 32 GB | ~1.1 GiB | 32K works: 438 slots, filled from the profile, 96% hit rate, ~29 tok/s decode with MTP |
| 64 GB | 6.69 GiB (200K) | sizes 19339 slots, but the 200K prompt paths crash (below) |

Consequences of the small figure, all reproduced:

- auto: "expert cache auto: 0.00 GiB free → 0 slots" → "no VRAM is left for the expert cache".
- forced `--expert-cache N`: the sizing pass sets `expert_cache = 0` (keep_bytes ≤ 0), so the cache
  opens with no slots and the profile fill reports "pre-filled 0 of 0 slots from the profile"; the
  server then exits because the MTP verify tier needs the profile-filled VRAM expert tier. An IQ
  pack cannot drop MTP (`it needs --native SHARD1, --spec T and --prefill CHUNK`), so there is no
  working fallback: a catch-22.

A patch that applies the same unified-memory correction on Windows (integrated GPU →
`GlobalMemoryStatusEx().ullAvailPhys` minus the same headroom) fixed all of it here: at 16 GB
carve-out and 200K context, the auto cache sized 19339 slots (36.84 GiB), filled 19336 of them from
the profile, and served 40-45 tok/s decode at 96.9% hit rate. I will send it as a PR.

## Finding 2: the gfx1151 arch defaults never apply on Windows HIP

`src/program/generate.cpp` (the `#if defined(STRATA_USE_HIP)` block) applies
`apply_arch_defaults(ad.gcnArchName)` only when `cudaGetDeviceProperties` succeeds, and swallows the
error otherwise. On this machine (Windows HIP 10.2.0a20260930 wheels) it fails, so the 18 exact
speed switches silently stay off — no "strata: gfx1151 (Strix Halo): N exact speed switches" line
at startup. Setting the 18 switches as environment variables directly (same values the table sets)
took decode from 26 tok/s to 40-45 tok/s at 200K, byte-identical answers as documented. A fix would
either retry the property query, log the failure, or apply the defaults when the device query
cannot run.

## Finding 3: at 200K on Windows, the prompt/verify paths crash whenever the cache is large

With the probe reading larger figures (64 GB carve-out, pre-patch; or any large cache), each
configuration died in a different subsystem, all in the `unspecified launch failure` (719) family:

- `--prefill auto` (auto → 6144) + auto cache: `strata serve: cannot create the refill stream` /
  `Hip error 719` inside hipblaslt.cpp:468;
- the same with `STRATA_GFX1151_DEFAULTS=0` (portable path): `strata serve: verify: stream create
  failed`;
- forced cache that lends slots to the prompt path: `prefill: the token id buffer (unspecified
  launch failure)`, and at 32K with a large borrowed region: `prefill: setting the FP16 activation
  image failed`;
- with a small cache that cannot lend (no borrowing): no GPU crash at all, but then verify's tier
  is missing (finding 1) and the server exits.

32K with the stock defaults on the 32 GB carve-out (cache 438 slots, prompt path borrowing only
310 of them) is the one configuration that served stably — which is what makes me suspect the
borrowed-region path breaks once the borrowed region (or the cache) crosses some size on this
driver, though I could not pin the threshold and did not test beyond what is listed.

## Also noted

- The engine's own `--prefill 8192` warning (#1630) matches: with a forced `--prefill 8192` the
  commit charge jumped ~60-70 GB and the process died; `--prefill 6144` and below were stable here
  too. A bare `--prefill auto` stopping at 6144 on Windows is what kept the working 32K setup safe.
- BIOS carve-out interacts with everything above through finding 1; on Windows, unlike the Linux
  guide's advice, a LARGER carve-out helps (it is the only thing that grows the reported figure).

## What I did not test

- Any Linux build; the SYCL port; multi-GPU; images (`--vision cpu` was configured but unused);
  other model sizes; `--kv q4_0` / `k8v4`; the CUDA backend (no NVIDIA GPU).
- Long-running stability at 200K beyond a benchmark session (hours), and 256K/384K/512K contexts.

Happy to test builds or run more configurations on this machine — it is currently serving 200K at
40-45 tok/s decode with the patch and the 18 switches set explicitly.