#2766 Windows x64 HIP 0.1.44.1 on Strix Halo (gfx1151): default resident-arena start dies with 'the embedding buffer failed'; --mmap-experts works

Descripción

# Windows x64 HIP (0.1.44.1) on Strix Halo / gfx1151: default start exits with "the embedding buffer failed" — `--mmap-experts` runs

**TL;DR** — On a Ryzen AI Max+ 395 (Radeon 8060S, gfx1151) under Windows 11, the ready-made
`strata-windows-x64-hip.zip` (0.1.44.1) **starts, loads the model, prints `session is up`, and then dies**
with `strata generate: the embedding buffer failed` (that is a **10 KB** `cudaMalloc` in `generate.cpp:6018`).
The expert cache is sized from **5.41 GiB of "free" VRAM → 1087 of 24576 slots** in that mode.
Adding `--mmap-experts` (plus `--vram-reserve-mib 512`) makes the same install run correctly, and in that
mode the process happily uses **~57.5 GB of dedicated VRAM**. Details, logs and a hypothesis below.

---

## 1. Environment

| | |
|---|---|
| Machine | Six United AXB35-02 mini-PC |
| APU | AMD Ryzen AI Max+ 395, Radeon 8060S, **gfx1151**, 20 CUs, wave32 |
| Memory | 128 GB LPDDR5X-8532 (8 × 16 GB); **BIOS "UMA Frame Buffer" = 64 GB** → Windows sees **63.6 GB** |
| OS | Windows 11 Pro, build **26200** (25H2) |
| GPU driver | **32.0.31021.5001** (2026-06-28) |
| Radeon software | Adrenalin 26.6.4 |
| Engine | `strata-windows-x64-hip.zip`, v0.1.44.1, sha256 `e40bce7a2cfe47ea42d3f2af1e0adb92d14b8c02a29107c5a309b2cb518f5f52` (matches `SHA256.txt`) |
| `engine/BUILD.json` | `source=prebuilt, backend=hip, platform=windows-x64, version=0.1.44.1, rocm=10.2.0a20260930, hipblaslt_version=100500`, archs incl. `gfx1151`, `vision=none` |
| Python | 3.14.4 (venv created by `START-HERE.bat`) |
| Model | `ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-GGUF`, rev `ed59f920`, **IQ3_S**: shards `54,817,524,224` + `28,800,138,432` B, `mmproj` `907,543,008` B (all sha256-verified) |

`engine\strata-device.exe --list-devices` / `--selftest` (with `engine\rocm\bin` on PATH) are healthy in both modes:

```
device 0: AMD Radeon(TM) 8060S Graphics
  HIP arch            gfx1151 wave32 (compiled for gfx1100,gfx1101,gfx1102,gfx1200,gfx1201,gfx1030,gfx1151)
  multiprocessors     20
  VRAM total / free   107.647 GiB (115585384448 B) / 107.496 GiB (115422916608 B)
  driver / runtime    71726391 / 71726391
  HIP runtime         E:\Strata-win\engine\amdhip64_7.dll
  memory plan  --max-context 20480 … plan + KV vs free   FITS
selftest: over-allocation refused as required
strata-device selftest OK
```

---

## 2. Install path that works end-to-end

1. **Layout.** The three Windows zips on the release page are **engines only** — `START-HERE.bat`,
   `setup.py`, `serve/`, `tools/` are *not* in them. I took the scripts from the source zip of `main`
   (`codeload.github.com/Niko1221/Strata/zip/refs/heads/main`, 21.6 MB) and unzipped
   `strata-windows-x64-hip.zip` into `engine\` (contents: `strata.exe`, `strata-device.exe`, `rocm\`,
   `amdhip64_7.dll`, `amd_comgr.dll`, `rocm_kpack.dll`, `msvcp140.dll`, `vcruntime140*.dll`, `BUILD.json`).
   → **Suggestion:** mention on the release page that the HIP zip is the engine and where the scripts come from.
2. `START-HERE.bat` → venv (Python 3.14.4) → `setup.py`.
3. ```
   setup.py --family qwen --model IQ3_S \
     --gguf-dir <folder with the two IQ3_S shards + the mmproj> \
     --context 131072 --vision no --backend hip --yes
   ```
   * engine: `[ok] ready-made AMD engine already installed`
   * model: `[ok] model files present` (no re-download), pack → `Strata-data\packs\iq3_s`
     (`302` canonical tensors served natively, arena 1.43 GiB; experts are read from the GGUF in place,
     so `Strata-data` is only ~7.9 GB on disk)
   * MTP: 31 tensors / 5.214 GB fetched → `mtp-q2_0.gguf` (0.889 GB) → runtime dir
   * start script written with `--expert-cache auto --prefill auto --spec 4 --spec-min-p 0.5 --mtp … --max-context 131072 --kv int8`

> Note (not a bug, but it costs an hour): the venv `pip install` needs a reachable index. With the
> Tsinghua mirror one pinned wheel 403'd (`numpy==2.5.3` `cp314-win_amd64`); Tencent/aliyun/USTC/pypi.org
> all served it. Setup's `--quiet` makes that step look frozen.

---

## 3. Bug: default (resident expert arena) start dies

Command (exactly what setup generated, no extra flags):

```
engine\strata.exe --serve --pack E:\Strata-data\packs\iq3_s \
  --native …-00001-of-00002.gguf --ple-gguf …-00002-of-00002.gguf \
  --expert-profile E:\Strata-win\data\expert-profile.bin \
  --expert-cache auto --prefill auto --spec 4 --spec-min-p 0.5 \
  --mtp E:\Strata-data\mtp\rt --max-context 131072 --kv int8
```

Engine log (verbatim, one line per entry):

```
strata: gfx1151 (Strix Halo): 18 exact speed switches on by default (…): GDN_HEAD GDN_PP GDN_CONVL2 GDN_NOY
        CVEC_FUSE HCD_EXACT PF_PAD Q8_PACKED Q6_PACKED MMVF_ROWS ATTN_LANECELL EXPERT_V2 TSUM LFUSE GDN_SPLIT QFUSE PLE_BATCH SH_STREAM
strata generate: memory guard: limit 63.6 GiB (source: meminfo), current 16.1 GiB, reclaimable cache credited 0.0 GiB;
        available 47.5 GiB (MemAvailable 47.5 GiB), floor 2560 MiB
strata generate: native pack: … experts (largest blob 2.66 MB), token embedding IQ4_XS in mapped host memory (322 MiB, 0.3 s)
strata generate: 300 native projection matrices, 2018.88 MiB of weights, in 2.7 s
strata generate: PLE on, table 320001536 rows of …-00002-of-00002.gguf
strata mtp: draft layer loaded, 950 MiB of VRAM (experts 675, dense 111), files read in 0.21 s
strata: STRATA_Q6_PACKED=1 packed the output head (248320 x 2560, +497.3 MiB)
strata generate: GPU 0: AMD Radeon(TM) 8060S Graphics (gfx1151)
strata generate: HIP runtime E:\Strata-win\engine\amdhip64_7.dll
strata generate: expert arena read unbuffered (0 of 16 probe reads from the file cache; 0.6 GiB available, 51.1 GiB of files)
strata generate: expert arena: cudaHostRegister PORTABLE ok; large pages refused for 50295996416 B
        (GetLargePageMinimum=2097152, VirtualAlloc error 1314); using 4 KB pages
strata generate: loaded 46.84 GiB at 2.03 GiB/s
strata generate: expert cache auto: Windows keeps 2560 MiB free after the cache on a card this size …
strata generate: expert cache auto: 5.41 GiB free, 2560 MiB reserved (+218 MiB for the draft head) -> 1087 slots
strata generate: pre-filled 0 of 0 slots from the profile in 0.0 s (0 MB/s); slot 0 verified
strata generate: session is up (engine 0.1.44.1)
strata generate: the embedding buffer failed          ← exit code 1
```

`the embedding buffer failed` is the **10 KB** allocation at `src/program/generate.cpp:6017-6021`:

```cpp
float* d_emb = nullptr;
if (cudaMalloc(&d_emb, (size_t) g.n_embd * 4) != cudaSuccess) {   // g.n_embd = 2560 → 10 KB
    std::fprintf(stderr, "strata generate: the embedding buffer failed\n");
    return 1;
}
```

### Reproduced twice, and it is **not** a system-RAM shortage

| run | `memory guard` reading | expert-cache line | outcome |
|---|---|---|---|
| #1 | `current 16.1 GiB`, `available 47.5 GiB` | `5.41 GiB free, 2560 MiB reserved -> 1087 slots` | `the embedding buffer failed` |
| #2 | `current 11.3 GiB`, `available 52.4 GiB` | `5.41 GiB free, 2560 MiB reserved -> 1087 slots` — **identical** | `the embedding buffer failed` |

Freeing 4.8 GB of system RAM (closing background apps) changed the guard reading but produced the
**same 5.41 GiB / 1087 slots** and the same failure, so the number the cache is sized from is not
"free system RAM" on Windows.

---

## 4. Workaround that works

```
--mmap-experts --vram-reserve-mib 512
```

```
[strata] mapping the experts from the model files (about 55 GB, --mmap-experts): they are not loaded into RAM …
[strata] filling the GPU's expert cache (24576 experts, 46.84 GiB of VRAM) …
ready: http://127.0.0.1:8080/v1  (OpenAI: /v1/chat/completions, Anthropic: /v1/messages, context 131072 tokens)
```

In this mode Windows Task Manager reports **57,521 MB dedicated** + ~1.7 GB shared for the process,
and `strata verify: capturing the … window (51573 MiB of VRAM free)`.
So the card is not out of memory — the failure above is specific to the **resident-arena** path.

---

## 5. Hypothesis (please sanity-check — I have not built the engine)

* `src/core/expert_cache.cpp:28 device_free_bytes()` returns `cudaMemGetInfo(&free,&total)` and adds a UMA
  correction **only under `#if defined(__linux__)`**: `/proc/meminfo` `MemAvailable` minus
  `STRATA_UMA_HEADROOM_GIB` (default 6). On Windows that correction is compiled out, so the expert-cache
  sizing uses HIP's free number as-is.
* `src/platform/memory.cpp:112 gpu_shared_memory_budget()` (DXGI
  `QueryVideoMemoryInfo(0, DXGI_MEMORY_SEGMENT_GROUP_NON_LOCAL)`) exists, but a repo-wide search finds its
  **only caller in the SYCL backend** (`sycl/src/core/pinned.dp.cpp:328`) — the Windows HIP path does not
  use it.
* On WDDM, host memory the GPU can reach is charged to the process's **non-local (shared)** budget, while
  the 64 GB BIOS carve-out is a separate **dedicated** segment. The resident arena is registered with
  `cudaHostRegister(…, PORTABLE)` (46.84 GiB). If HIP's `cudaMemGetInfo` free on a UMA adapter tracks that
  shared budget, then after registration the reported "free" collapses to ~5.4 GiB although the dedicated
  segment is nearly empty — which would explain both `1087 slots` and the later 10 KB allocation failure.
* Possible directions (yours to judge): exclude the registered arena from the sizing reading on
  Windows/UMA; use the DXGI non-local budget (or a `STRATA_UMA_HEADROOM_GIB`-equivalent) there; or
  allocate the arena from the dedicated segment; or make the automatic choice fall back to
  `--mmap-experts` on Windows when it would otherwise fail.

---

## 6. Speed, same machine, same model — and the Linux comparison

Windows (this install, `--mmap-experts --vram-reserve-mib 512`, IQ3_S, context 131072, KV int8),
from the engine's own `strata serve: prompt …` lines:

| prompt tok | prefill | generated | decode | MTP drafts accepted |
|---|---|---|---|---|
| 53 | 660 ms (80.3 tok/s) | 40 | 36.5 tok/s | 22 / 34 |
| 57 | 780 ms (73.1) | **1898** | **40.7 tok/s** | 1050 / 1531 |
| 58 | 705 ms (82.3) | 97 | 34.0 | 45 / 83 |
| 82 | 993 ms (82.6) | 600 | 38.0 | 312 / 515 |
| 588 | 1758 ms (334.5) | 748 | 30.7 | 269 / 551 |
| 2221 | 3772 ms (588.8) | 24 | 38.1 | 15 / 24 |
| **7261** | **11032 ms (658.2)** | 16 | 46.6 (too short to be meaningful) | 11 / 12 |

`strata serve: expert tiers: GPU 0 hits this request; since the start RAM 0 blobs, files 0 blobs 50294.7 MB read
(the GGUF in place)` — the byte counter only crept up (50,294.7 → 50,604.7 MB over several requests), i.e. the
cache looks like it is hitting, not thrashing.

Linux reference on the **same box and model** (Ubuntu, kernel 7.0.0, engine built locally for gfx1151,
default placement, same context): decode **47.2 / 47.4 / 47.5 / 48.9 / 55.2 tok/s**
(`done: 736 tokens in 17 s (48.9 tok/s)` … `done: 33 tokens in 1 s (55.2 tok/s)`).
`docs/STRIX_HALO.md` lists IQ3_S at 8K: **1,242 prompt / 59.7 decode**.

So on Windows I measure roughly **53 % of the documented prefill** (658 vs 1,242 tok/s) and
**~30–41 tok/s decode vs 47–60 on Linux/docs**. Question: is that the expected state of the Windows HIP
path on gfx1151 (driver/runtime 10.2.0a20260930, kernel shapes, the mmap expert path), or is something
about this configuration off?

---

## 7. Anything else I can run

Happy to test diagnostics/flags on this machine — including with the BIOS carve-out changed
(e.g. 512 MB / 16 GB instead of 64 GB, which `setup` itself says "leaves the RAM to the model"),
`--mmap-experts` vs arena, different `--spec`/`--expert-cache`/`--prefill` values, or a debug build, if that
helps narrow Bug 1 or the Windows/Linux gap.

**Attachments in this report:** the full engine log lines above come from
`strata-iq3_s.log` (engine) and the complete setup transcripts
(`strata-win-2026-10-16-122313.log`, 41 KB) — I can paste them in full or upload them if useful.