#1217 gfx1150 (Radeon 890M / Strix Point) works: report and a one-line CMake patch

Description

## Summary

Strata v0.1.40.1 builds and runs on a **Ryzen AI 9 HX PRO 370 (Radeon 890M, gfx1150, Strix Point)**. It serves
Qwen3.8-Flash-Next IQ3_XXS through the HIP backend, with ROCm from AMD's TheRock gfx1150 tarball.

- Today the build stops with a CMake `FATAL_ERROR`, because gfx1150 is in none of the three architecture lists in
  `cmake/hip_backend.cmake`. Nothing in the code, commits or issues points to a known gfx1150 problem
- The device code already covers gfx1150:
  - the WMMA guards (`__gfx1150__`) and `gfx_arch_is_gfx11_wmma()` include it (41c89b1708, 1bb3279d74)
  - `intrinsics.hpp` gives it sudot4
- With gfx1150 added to the "unvalidated" list (patch below):
  - the HIP test suite passes, except tests that need files a public checkout does not have
  - the model serves chat and tool calls
  - decode runs at 17-19 tokens/s and prompt reading at about 125-135 tokens/s

We would like gfx1150 to be admitted as an unvalidated (or community-validated) target. Some notes for setup and for
unified-memory APUs in general follow.

## Environment

| Item | Value |
|---|---|
| Machine | MINISFORUM N5 PRO, Ryzen AI 9 HX PRO 370 (12C/24T), Radeon 890M (gfx1150, PCI `1002:150e`, 16 CUs) |
| Memory | 96 GB DDR5-5600 (93.4 GiB usable), BIOS UMA carve-out 512 MB |
| GTT | Default 46.7 GiB (half the RAM). Raised to 64 GiB with `ttm.pages_limit=16777216 ttm.page_pool_size=16777216` |
| Host OS | Proxmox VE 9.2 (Debian 13), kernel `7.0.14-20-pve` (in-tree amdgpu / amdkfd) |
| Runtime environment | Unprivileged LXC container (Debian 13.7). `/dev/kfd` and `/dev/dri/renderD128` are passed with the container's render GID, and `lxc.prlimit.memlock: unlimited` is set |
| Toolchain | GCC 14.2, CMake 3.31.6, Ninja 1.12.1, Python 3.13.5 |
| ROCm | TheRock 7.14.1 for gfx1150 (`libhipblaslt.so.1.4`). See below |
| Strata | `v0.1.40.1` (82f46a8c). The engine is the same as v0.1.40 |

### ROCm (TheRock 7.14.1, gfx1150)

```sh
curl -LO https://repo.amd.com/rocm/tarball-multi-arch/therock-dist-linux-gfx1150-7.14.1.tar.gz
# 1,689,736,748 bytes. AMD publishes no checksum; ours was
# da8335c9bc230a8b0e3e43b21dedfb3cfc2456f3ecfecfaa49db40096f08bac5
mkdir -p 7.14.1 && tar -xzf therock-dist-linux-gfx1150-7.14.1.tar.gz -C 7.14.1   # 8.3 GB extracted
```

`rocminfo` reports the agent as `gfx1150`, "AMD Radeon 890M Graphics", 16 compute units. Its GPU memory pool is the
GTT size (46.7 GiB by default, 64 GiB after the change above).

## The patch

```diff
diff --git a/cmake/hip_backend.cmake b/cmake/hip_backend.cmake
--- a/cmake/hip_backend.cmake
+++ b/cmake/hip_backend.cmake
@@ -12,7 +12,7 @@ endif()
 # same wave32 / 64 KiB LDS / sudot4 family as gfx1100, on a unified-memory APU; experimental.
 set(_strata_hip_validated gfx1100 gfx1201)
 set(_strata_hip_community gfx1101 gfx1200)
-set(_strata_hip_unvalidated gfx1012 gfx1102 gfx1030 gfx1031 gfx1034 gfx1151)
+set(_strata_hip_unvalidated gfx1012 gfx1102 gfx1030 gfx1031 gfx1034 gfx1150 gfx1151)
 # CMake hands HIP a ';' list, but a -DCMAKE_HIP_ARCHITECTURES typed by hand (or ROCm's own Windows tooling) may use
 # spaces, which foreach(IN LISTS) would otherwise treat as one element.
 string(REPLACE " " ";" _strata_hip_norm "${CMAKE_HIP_ARCHITECTURES}")
```

The comment above the lists and `AMD_HIP.md` / `STRIX_HALO.md` would also need a line about gfx1150.

## Build

These are the same steps as `docs/STRIX_HALO.md` with the architecture changed. The build takes about 1.5 minutes
with 6 jobs.

```sh
SDK=$PWD/7.14.1
export ROCM_PATH=$SDK HIP_PATH=$SDK HIP_PLATFORM=amd PATH=$SDK/bin:$SDK/lib/llvm/bin:$PATH \
  LD_LIBRARY_PATH=$SDK/lib:$SDK/lib/rocm_sysdeps/lib:$SDK/lib/llvm/lib
cmake -S . -B build-gfx1150 -G Ninja -DCMAKE_BUILD_TYPE=Release \
  -DSTRATA_ENABLE_HIP=ON -DSTRATA_ENABLE_CUDA=OFF -DSTRATA_PREFILL_MMQ=ON -DSTRATA_PARITY_PROMPT_ATTN=ON \
  -DCMAKE_HIP_ARCHITECTURES=gfx1150 \
  -DCMAKE_HIP_COMPILER=$SDK/lib/llvm/bin/clang++ -DCMAKE_HIP_COMPILER_ROCM_ROOT=$SDK \
  "-DCMAKE_PREFIX_PATH=$SDK;$SDK/lib/rocm_sysdeps;$SDK/lib/llvm" \
  "-DCMAKE_HIP_FLAGS=--rocm-path=$SDK --rocm-device-lib-path=$SDK/lib/llvm/amdgcn/bitcode"
cmake --build build-gfx1150 -j 6
```

The only warning specific to this target is the expected one: "gfx1150 builds, but it is not validated on a real card
yet". hipBLASLt is found.

## Tests (`ctest`, 60 registered)

**55 passed, 4 skipped, 1 failed. None of the skips or the failure is a gfx1150 problem:**

| Test | Result | Why |
|---|---|---|
| WMMA, fused MoE / IQ prompt experts, prompt attention, GDN, QSA, KV (q8 / q4 / stream / hybrid), sampler, router, PLE reader and the other HIP and parity tests | Passed | Including `hip_prefill_wmma_gemm_parity`, `hip_prompt_attn_wmma`, `hip_prefill_fused_moe`, `hip_prefill_fused_iq`, `hip_prefill_mmq_parity`, `hip_intrinsics` and `hip_device_selftest` |
| `hip_prefill_hcd_exact_parity`, `hip_prefill_hipblaslt_gemm` | Skipped | There is no hipBLASLt tuning table for gfx1150 (as expected) |
| `iq_parity_fixtures`, `iq_parity` | Skipped | See below. Run by hand, all formats pass |
| `ple_parity` | Failed | It needs data that a public checkout does not have. See below |

**`iq_parity` by hand.** The fixture generator needs llama.cpp's `gguf-py`, which comes with setup's llama.cpp
download, and PyYAML. In a manual build, the copy CMake fetches works: it is the same pinned commit `3cf03257`.

```sh
PYTHONPATH=build-gfx1150/_deps/strata_llamacpp-src/gguf-py python3 tools/iq_fixture.py --out build-gfx1150/iq_fixture
build-gfx1150/iq_parity build-gfx1150/iq_fixture
```

The result was `dequant rel 0.00e+00` and `mmvq rel 4.6e-03 .. 5.7e-03 (ncols 1..8) ok` for all 10 formats:
IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S, IQ1_M, IQ4_NL, IQ4_XS, Q2_0 and Q3_K.

**`ple_parity`.**
- It needs a Q2_0 shard 2 and the dense pack (`pack/full/dense.bin`). With `--gguf <IQ3_XXS shard 2> --pack <the IQ3_XXS
  pack>`, it then stops at `bench/micro/ple_in.bin / ple_out.bin`, which a public checkout does not ship
- So in a public checkout it always fails. It could be registered only when those fixtures exist, as with the other
  `bench/micro` tests

The GPU tests also ran while another process (llama.cpp on Vulkan) used 44 GiB of the same GTT pool. The kernel
logged no GPU faults or resets.

## Running Qwen3.8-Flash-Next IQ3_XXS

`setup.sh` rejects gfx1150, so we prepared the model by hand, following what `setup.py` does.

- Model: `ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-GGUF` at the pinned revision `ed59f920`. We used both IQ3_XXS shards
  and checked their SHA-256 against the Hugging Face LFS ids
- `tools/iq_pack.py --gguf <shard 1> --out <pack>`, then again with `--experts-bin` (42.9 GB)
- MTP: `tools/mtp_fetch.py fetch`, `tools/mtp_pack.py --experts q2_0`, `tools/mtp_rt.py`, then `data/draft_vocab.bin` (cjk)
  copied into `rt/`
- Server: `serve/server.py --engine strata --config <json>`. The JSON has the same keys as setup writes, and
  `STRATA_API_KEY` is set

Engine arguments:

```
--pack <pack> --native <shard 1> --ple-gguf <shard 2> --expert-profile data/expert-profile.bin
--expert-cache 24576 --mmap-experts --prefill auto --spec 4 --spec-min-p 0.5 --mtp <mtp>/rt
--max-context 65536 --kv int8
```

Startup from a cold page cache took about 45 s:

```
[strata] filling the GPU's expert cache (24576 experts, 39.97 GiB of VRAM) ...
strata generate: pre-filled 24576 of 24576 slots from the profile in 23.9 s (1794 MB/s); slot 0 verified
strata generate: token graph hit path: 24576 resident experts, decided on the device
strata serve: 18438 MiB of VRAM free with everything loaded
```

- GTT in use was about 46 GiB in total
- Chat in Japanese and tool calls worked. We tested with Hermes Agent: a web search tool call, then the answer.
  The second turn reused the 4.6K-token prefix from the prompt cache
- Japanese output was fluent, and we saw no Chinese mixed in

## Speed

Each run is one request at a time, with server-reported timings. The engine was restarted for each column.

| Prompt | Default (no switches) | gfx1151 exact switches set by hand (see below) | Exact + rounding-level prompt switches |
|---|---|---|---|
| 51 tokens, 180 out: first token / decode | 0.25 s / 16.7-16.9 tok/s | 0.24 s / 17.1 | 0.24 s / 16.9 |
| 3.6K: prefill / decode | 126 / 17.0 | 124 / 18.3 | **83** / 18.4 |
| 7K: prefill / decode | 135 / 16.9 | 133 / 18.5 | **86** / 16.9 |
| 28K: prefill / decode | 129 / 18.7 | — | — |

- MTP drafts: 81 of 110 accepted on the short chat
- `STRATA_PREFILL_MMQ=1` made no difference (124 / 130 tok/s prefill)

**Prefill is GPU-bound.** This is `STRATA_PREFILL_TIMING=1` for 3,556 tokens with the exact switches:

```
GPU timeline 28362 ms, wall 28469 ms, host staging 3695 ms: hc read 4745 (16.7%) gdn 6889 (24.3%)
qsa proj 2992 (10.6%) qsa attn 2742 (9.7%) router+shared 1334 (4.7%) gather 291 (1.0%) dequant 988 (3.5%)
gemm gate/up 2815 (9.9%) gemm down 1501 (5.3%) combine 535 (1.9%) ple 217 (0.8%) gdn conv+gates 237 (0.8%)
gdn recurrence 432 (1.5%) gdn out proj 2533 (8.9%)
host: chunk setup 2 ms, waiting for each chunk 108 ms, after each chunk 105 ms, PLE 0 ms
```

With 16 CUs, the GDN, hyper-connection and projection work sets the ceiling. A gfx1150 hipBLASLt table
(`tune_hipblaslt`) could help the GEMM parts, about 15% of the time. We have not tried that yet.

## Notes and suggestions

1. **Admit gfx1150** (the patch above). gfx1152 and gfx1103 have no WMMA code (1bb3279d74), but gfx1150 does.
2. **setup.py and gfx1150.** The UMA path (`amd_apply_uma`, the build note, `STRIX_HALO_*`) is keyed on gfx1151 only.
   - gfx1150 is the same kind of unified-memory APU (PCI `0x150E` is already in `AMD_IGPU_PCI`)
   - Running the Strix Halo path for `gfx1150` too (build for gfx1150, count the GTT pool as GPU memory) would let
     setup install it
3. **Unified memory doubles the experts in the default mode.**
   - The default mode copies every expert into a pinned RAM arena and also fills the GPU cache. On an APU both
     live in the same RAM: 2 x 42.9 GB for IQ3_XXS does not fit a 96 GB machine
   - `--mmap-experts` with an expert cache large enough for every expert (24,576 slots) keeps one copy: GTT holds
     the experts, and the file pages are reclaimable after the fill
   - Setup could pick this on an APU whose RAM is below about 2x the arena
4. **Containers and shared GPUs see the wrong free memory on UMA.**
   - `device_free_bytes()` uses `/proc/meminfo` MemAvailable less 6 GiB. Inside an LXC container, LXCFS reports a
     container-scoped MemAvailable, while GTT allocations are not charged to the memory cgroup
   - Also, `hipMemGetInfo` reported the whole 46.7 GiB pool as free while a Vulkan process held 44 GiB of it
   - We set `--expert-cache` to a fixed number. A note in the docs (or a `STRATA_UMA_FREE_GIB`-style override)
     would help
5. **The gfx1151 exact switches also help gfx1150.** We set `arch_default_env`'s table by hand, minus
   `STRATA_HCD_EXACT`, which needs the gfx1151 hipBLASLt table.
   - Decode was about 8% faster (17.0 -> 18.3, 16.9 -> 18.5 tok/s), and prefill did not change
   - The answers to our short Japanese prompt began with the same text. We did not run a bitwise comparison
   - Applying the table to gfx1150 as well, or adding a gfx1150 table, seems worthwhile
   - The rounding-level prompt switches of `STRIX_HALO.md` section 5 (`PF_FUSED`, `PF_GEMM`, `HC_UPMIX`,
     `PA_FAST`, `HIP_WMMA`, `SELECT_WMMA`, `HC_Q8`) were **slower** on gfx1150: prefill went from 124 to 83 tok/s
6. **Tests in a public checkout.** `ple_parity` always fails there, and `iq_parity` is skipped unless gguf-py and
   PyYAML are installed (see above). The CMake-fetched `_deps/strata_llamacpp-src/gguf-py` could be used as a
   fallback in `tools/_paths.py`.

**Not tested:**
- Windows, and the vision encoder
- `"parallel"` slots, and contexts above 64K
- Stability over days (so far a few hours, including a host reboot and service restarts)