#497 0.1.35: Windows AMD report: the ready-made zip on an RX 6800 (gfx1030)

Beschreibung

As i am testing also on my other Machine With AMD, here a Report:

The ready-made zip on a discrete card, as docs/AMD_HIP.md asks: 0.1.35's `strata-windows-x64-hip.zip` (and 0.1.34's before it) runs on an RX 6800 (gfx1030) and serves Qwen3.8-Flash-Next IQ3_XXS at 131072 context. Two things showed up that concern the zip on Windows:

1. **0.1.35's zip with setup's `--expert-cache auto` decodes at 31.4 tok/s; 0.1.34's zip (the same kernels and ROCm) with a fixed `--expert-cache 3072` at 42.0 tok/s.** `auto` sizes the cache from `hipMemGetInfo`, which on Windows does not subtract what the desktop holds. The cache then fills the card past what Windows lets the process keep in VRAM, and Windows keeps about 1.1 GiB more of the engine in system memory than with a budget-aware build. #380 fixes this in the engine, so a zip built with it gets the fix too.
2. **The PCIe probe reads 22,093.5 GB/s** (18,970.7 GB/s on 0.1.34; the local build with #377, timed on the host, reads 27.2). That is harmless on this x16 link, since the default 0.55 stays from 20 GB/s up, but the zip would not notice a slow link. #377 fixes this.

### The PC

RX 6800 16 GB (gfx1030, drives the desktop), driver 32.0.21045.9003 (AMD Software: Adrenalin Edition 26.9.1), Ryzen 7 5800X3D, 64 GB RAM, PCIe 4.0 x16, Windows 11 (build 26300).

System32 holds the driver's `amdhip64_6.dll` and `amd_comgr_2.dll` and no `amdhip64_7.dll`, so #468/#461 did not apply here. 0.1.34's zip loaded its runtime from `rocm\bin` (the process's module list), and 0.1.35's loads it from beside `strata.exe`.

### strata-device (0.1.35, `rocm\bin` first on the PATH)

```
device 0: AMD Radeon RX 6800
  arch gfx1030, 16.0 GiB, wave32
```

```
device 0: AMD Radeon RX 6800
  HIP arch            gfx1030 wave32 (compiled for gfx1100,gfx1101,gfx1102,gfx1200,gfx1201,gfx1030)
  multiprocessors     30
  VRAM total / free   15.984 GiB (17163091968 B) / 15.847 GiB (17015111680 B)
  driver / runtime    71726391 / 71726391
  HIP runtime         D:\AI\Strata\engine-0135\amdhip64_7.dll

memory plan  --max-context 20480
  KV + indexer keys   0.267387 GB
  recurrent state     0.000000 GB  (GDN recurrence + conv history, NOT evictable)
  expert cache        5.675613 GB  (4105 slots of 1,382,400 B)
  VRAM pool           5.943000 GB
  expert arena (DRAM) 0.000000 GB
  card free           15.847 GiB (17015111680 B)
  plan + KV vs free   FITS

selftest: arena 0.062 GiB (67108864 B), used 0.002 GiB (2097152 B) after two 1 MiB allocations
selftest: over-allocation refused as required
strata-device selftest OK
```

0.1.34's output was the same, without the `HIP runtime` line.

### The engine

The zips were unpacked into `engine-0134\` and `engine-0135\` beside a self-built `engine\`. The zip's `strata.exe`, with its `rocm\bin` first on the PATH, was driven through `--serve` by a script rather than the server. Model and configuration are this PC's: Qwen3.8-Flash-Next IQ3_XXS (native pack, MTP draft head) with `--prefill auto --spec 4 --spec-min-p 0.5 --max-context 131072 --kv int8 --kv-resident 32768 --pcie-frac 0.2`, and the cache size as given. Each variant was started twice. Each start ran a first answer, then `tools/calibrate.py`'s three prompts twice (128 tokens each), then a 32K-token prompt.

The end of the startup (0.1.35 with `--expert-cache auto`, from a start without `--pcie-frac` so that the probe runs; with `--pcie-frac 0.2` the last line reads 581 MiB):

```
strata generate: PCIe probe: 22093.5 GB/s host->device -> pcie_frac 0.55 (default 0.55)
...
strata generate: GPU 0: AMD Radeon RX 6800 (gfx1030)
strata generate: HIP runtime D:\AI\Strata\engine-0135\amdhip64_7.dll
strata generate: expert arena: cudaHostRegister PORTABLE ok; large pages refused for 42916118528 B (GetLargePageMinimum=2097152, VirtualAlloc error 1450); using 4 KB pages
strata generate: loaded 39.97 GiB at 2.97 GiB/s
strata generate: experimental native Q5_K head, 437043200 bytes
strata generate: expert cache auto: 10.69 GiB free, 700 MiB reserved (+184 MiB for the draft head) -> 4530 slots
strata generate: expert cache 6072 slots, 9.83 GiB of VRAM; policy is
...
strata generate: pre-filled 6072 of 6072 slots from the profile; slot 0 verified
strata generate: R4 hit path ON - resident experts are computed on the GPU
strata generate: 7 expert-pool workers + the host thread
strata generate: session is up (engine 0.1.35)
strata generate: token graph hit path: 6072 resident experts, decided on the device
strata serve: prompt chunk auto: 8192 tokens
strata serve: the prompt path borrows 2770 CUDA0 cache slots (4.44 GiB)
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 verify: window up to 6 tokens, 69.6 MiB of device buffers
strata mtp: draft head over 106299 tokens (178.4 MiB)
strata serve: 503 MiB of VRAM free with everything loaded
```

The speed line of a first answer (a 34-token prompt, 32 tokens):

```
0.1.35, --expert-cache auto:  prompt 34 tokens = 0 reused + 34 read in 759 ms (44.8 tok/s), 32 generated in 1204 ms (26.6 tok/s), drafts accepted 24 of 30, 1 checkpoints
0.1.34, --expert-cache 3072:  prompt 34 tokens = 0 reused + 34 read in 637 ms (53.4 tok/s), 32 generated in 805 ms (39.8 tok/s), drafts accepted 24 of 30, 1 checkpoints
```

| | expert cache | decode, median of 12 requests | 32K prompt, mean of 2 |
|---|---|---|---|
| zip 0.1.34, `--expert-cache 3072` | 4090 slots, 6.66 GiB | 42.0 tok/s | 310 tok/s |
| zip 0.1.35, `--expert-cache auto` | 6072 slots, 9.83 GiB | 31.4 tok/s | 285 tok/s |
| local build with #380, `auto` | 4960 slots, 8.05 GiB | 42.8 tok/s | 317 tok/s |
| the same build, `auto`, `STRATA_WDDM_BUDGET=0` | 6072 slots, 9.83 GiB | 32.4 tok/s | 298 tok/s |
| the same build, `--expert-cache 3072` | 4090 slots, 6.66 GiB | 40.5 tok/s | 319 tok/s |

The local build is 0.1.32 with #356, #373, #376, #377 and #380, plus the verify-window scratch reserve described on #356 (branch `win-scratch` on my fork, not proposed), compiled on this PC with the HIP SDK 7.2. 0.1.34 and 0.1.35 have the same kernels and ROCm; 0.1.35's engine changes are the runtime log line, #467's resident RAM mode (not used here) and #448's chunk choice for a layer split. With `--expert-cache 3072` (4090 slots in both), 0.1.34's zip decodes 2.6-3.7 % faster than the local build and reads a 32K prompt 2.6-2.7 % slower. `STRATA_PREFILL_TIMING` puts the prompt difference in hc read, gdn recurrence and gdn out proj.

### `--expert-cache auto` on Windows

- **Cache size and memory:** without the budget, both engines start the sizing from 10.69 GiB free and keep all 6072 slots. Windows' own GPU memory counters (`GPU Process Memory` and `GPU Adapter Memory`, read after the start quoted above) show:
  - without the budget: the engine has 14.65-14.66 GiB in VRAM, and the card has 15.72-15.74 of its 15.98 GiB in use;
  - with #380 (Windows budgets 15569 MiB for this process): 4960 slots, the engine 13.98 GiB in VRAM, the card 15.05 GiB in use.
- **The overflow:** without the budget the cache is 1.78 GiB larger, but the engine holds only 0.67-0.68 GiB more in VRAM; the other 1.11-1.12 GiB is in system memory. Windows keeps the overflow out of VRAM. Forty seconds into the 32K prompt of the second start, the budget-off variants had 0.47 and 0.48 GiB less in VRAM (0.48 GiB more in system memory) than at the same point of the first start. The budget-on variant read the same in both starts.
- **Where decode loses** (`STRATA_DECODE_TIMING`, mean per verify window, the local build): 79.6 ms without the budget, 59.5 ms with it (the zip with `auto`: 81.0 ms).
  - The draft step goes from 5.6 to 23.1 ms (the zip: 22.6 ms).
  - The wait for the GPU goes from 19.8 to 25.1 ms (the zip: 25.9 ms).
  - The larger cache does take work off the CPU: 7.6 instead of 8.8 CPU experts (10.3 instead of 11.8 entries) per layer and window, and the host's per-layer work drops from 27.3 to 24.5 ms per window.
  - A smaller part of the decode gap is the drafts: 2.47 instead of 2.56 tokens per window (about 3.5 %). A different cache size changes which experts the GPU computes, and the GPU rounds differently from the CPU, so the drafts are accepted a little differently.
  - The draft step's slowdown fits Windows having moved the draft layer (835 MiB) to system memory. I have not checked which allocations moved.
- **The zip with `auto`** matches the local build with the budget off: the same 6072 slots, the same memory figures and nearly the same slowdown (31.4 against 32.4 tok/s). On this card, setup's default costs the zip about a quarter of its decode speed.

### The verify-window hang from 0.1.30 (mentioned on #356)

It did not come back. The zip has no scratch reserve and ran all 33 of its requests (both versions) without a hang. The local build with the reserve switched off also answered normally. I am dropping the reserve.

Not checked: the content of the answers (only that the drafts are accepted as often as with the local build at the same cache size), the zip installed through START-HERE, and images.