#1545 Intel A770 xe driver

本文

[strata-a770-xe-alchemist.patch](https://github.com/user-attachments/files/33200827/strata-a770-xe-alchemist.patch)
## Arc A770 16 GB — Alchemist on `xe`: three fixes, with numbers

I got the SYCL port running on an **Arc A770 16 GB on the `xe` driver**, which is the gap in the two recipes the docs cover: every Alchemist result here is an A750 on `i915`, and every `xe` result is a Battlemage card. That matters because the port's Alchemist defaults are keyed off the driver name (`intel_i915_gpu()`), so an Alchemist card on `xe` silently inherits the Battlemage values.

Nothing here changes behaviour on an NVIDIA, AMD, `i915` or Battlemage-on-`xe` card — two of the three are diagnostic/reporting paths and the third only fires when a driver reports no free memory at all.

**Host:** Minisforum BD895i SE (Ryzen 9 8945HX), 29 GB usable RAM, CachyOS kernel 7.2.9, `xe`, NEO 26.35.39758, oneAPI 2026.0, native build (no Docker), `level-zero-loader` 1.32.0.
**Model:** Coder IQ1_M (23.42 GiB of experts), 64K context, int8 KV, MTP `--spec 4 --spec-min-p 0.5`, `--mmap-experts`.

### Results

| | |
|---|---|
| decode, 300-token greedy (5 runs) | **6.99–7.54 tok/s**, all five byte-identical in content |
| prompt | 24–32 tok/s |
| through the OpenAI API, 128-token answers | 17.5–18.2 tok/s, MTP drafts **72–80% accepted** |
| VRAM free with everything loaded | 2,144 MiB of 15.11 GiB |
| parity tests | **22 of 25 pass on the card** (upstream's CPU-device baseline is 14 of 25); no kernel hangs |

Decode is below the Pro B50's ~23 tok/s from #423 and the reason is placement, not the port: that model's experts all fit VRAM, this one's cannot (a 4 GiB pinned-host ceiling, below), so the CPU computes 7–15 distinct experts per layer and the GPU waits for them.

### 1. The uncached doorbell needs `<level_zero/ze_api.h>` at build time — without it, every token is 0

`sycl_queue.hpp` allocates the doorbell/ring buffers with `ZE_HOST_MEM_ALLOC_FLAG_BIAS_UNCACHED`, but only when `STRATA_HAVE_ZE` is defined, and that needs the header **at build time**. Distributions split the Level Zero loader from its headers (Arch: `level-zero-loader` vs `level-zero-headers`; Ubuntu `libze1` vs `libze-dev`), so a box whose Level Zero runtime works fine builds the port without them. `host_malloc_polled` then falls back to `sycl::malloc_host`, a running kernel never sees the host's stores, every ring wait runs to `kSpinMax`, and the window goes on with the CPU experts' rows missing: NaN logits and token 0 for every token.

The failure is completely silent — clean build, engine loads, prefill reports a sane rate, parity tests pass. Measured on one config, before and after:

| | header absent | header present |
|---|---|---|
| `wait for rings` | **42,140 ms/round** | **98.7 ms/round** |
| `pool` (the CPU experts, unchanged) | 26.8 ms/round | 34.1 ms/round |
| decode | 0.08 tok/s | 12.27 tok/s |
| output | **32 of 32 tokens were 0** | **0 of 1,500 were** (5 runs × 300) |

The host was never slow — it served its experts in 27 ms while the GPU waited 42 s for a store it could not see.

The patch adds a `check_include_file_cxx` at configure time that warns when the header is missing (naming both failure and package), plus `-DSTRATA_LZ_INCLUDE=<dir>` for a header installed elsewhere, and links `ze_loader`, which the doorbell calls directly. It does **not** vendor the header. `nm -D strata | grep zeMemAllocHost` tells you whether a build has it.

### 2. `xe` reports no free VRAM, so every cache sized to 0 slots

This driver does not expose `sycl::aspect::ext_intel_free_memory` (`ZES_ENABLE_SYSMAN=1` does not change it), so dpct returns 0 bytes free and `--spec` refuses to start with `it needs --spec T (T >= 2)`.

`device_free_bytes()` already has a DRM-fdinfo fallback, but its guard is:

```cpp
own > 0 && total_b > own && total_b - own < free_b
```

which can never be true for `free_b == 0` (unsigned) — so the fallback is unreachable exactly in the case that needs it. Arming it on `free_b == 0` as well is the whole fix; `xe` does report `drm-total-vram0` in `/proc/self/fdinfo/*`. Expert cache: **0 slots → 1,934** (3.72 GiB), nothing else changed. The `i915` case the guard was written for is untouched.

### 3. The serve-time VRAM check read a bare 0

`generate.cpp`'s "N MiB of VRAM free with everything loaded" check called `get_memory_info()` directly instead of `device_free_bytes()` like the four other free-VRAM readings in the same file. On `xe` it therefore printed `0 MiB free ... add --vram-reserve-mib N` on every start **no matter how high the reserve was raised** — advice that chases its own tail, and a card that could never satisfy it. Routed through `device_free_bytes()` it reports the true 2,144 MiB.

### Two runtime knobs (not in the patch)

**`STRATA_WINDOW_SYNC=1` is required here.** `verify.cpp` defaults it to 1 on `i915` and 0 elsewhere, for the documented Alchemist behaviour that a window launched behind other in-flight GPU work never starts. This card reports `xe`, so it gets 0, and long runs then stall with `verify: timed out at layer 0` — with the host having *already served* that layer, so it is not the handshake of #1. Measured on 300-token runs with the adaptive expert tier on (its async swaps are exactly the "other GPU work"):

| lever | runs | completed |
|---|---|---|
| `WINDOW_SYNC=0` (the inherited default) | 3 | **0** |
| `WINDOW_SYNC=1` | 4 | **4** |
| `WINDOW_SYNC=0` + `--adapt-every 0` | 1 | 1 |

Both levers work, so the stall needs the adaptive tier *and* a window launched behind it. I deliberately left the default alone: keying it correctly means detecting Alchemist by *architecture* rather than driver name, and 8 runs on one machine is thin evidence for a change that would cost every Battlemage-on-`xe` card a queue drain per window. Your call — happy to send that as a follow-up if you'd rather it were automatic.

**`--mmap-experts` is required on a small-RAM host.** Without it the engine pins the whole 23.42 GiB arena (`cudaHostRegister PORTABLE ok`), leaving ~3 GB reclaimable on 29 GB: long runs were OOM-killed (exit 137) mid-decode with 338k major faults and 7.3 GiB swapped. mmap costs decode (**12.27 → 7.0–7.5 tok/s**) but keeps the box alive. Prefill is unaffected — 3.7 tok/s on the first run after boot, 24–25 once the file cache holds the pack, in either mode.

Related, and why the B70's `--stream-experts` recipe is not available here: its mirror is one pinned host allocation, and a `malloc_host` ladder on this card returns **null at 4 GiB** (2 GiB succeeds, held or freed), with `ulimit -l` at 8 MiB and unraisable without root.

### Recipe

```sh
OverrideDefaultFP64Settings=1 IGC_EnableDPEmulation=1 NEOReadDebugKeys=1   # Alchemist: no FP64 hardware
ONEAPI_DEVICE_SELECTOR=level_zero:gpu ZE_AFFINITY_MASK=0                    # not the iGPU
SYCL_CACHE_PERSISTENT=0 UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1
SYCL_PROGRAM_APPEND_COMPILE_OPTIONS=-ze-intel-greater-than-4GB-buffer-required
ZES_ENABLE_SYSMAN=1 STRATA_WINDOW_SYNC=1
STRATA_VERIFY_DEVICE_PLAN=1 STRATA_VERIFY_NO_HOST=0                         # Alchemist, not the xe/B70 value

strata --pack <pack> --native <shard1> --ple-gguf <shard2> --expert-profile <profile> \
       --mmap-experts --expert-cache auto --no-prefill-borrow --prefill 512 \
       --spec 4 --spec-min-p 0.5 --mtp <mtp-dir> \
       --max-context 65536 --kv int8 --vram-reserve-mib 2048
```

The pairing that is easy to get wrong: `STRATA_VERIFY_NO_HOST=0` (the Alchemist value, since the CPU computes experts outside VRAM here) together with `STRATA_WINDOW_SYNC=1` (also Alchemist) on a card that reports `xe`. `NO_HOST=1` — the B70's — hangs the window here; `WINDOW_SYNC=0` — the Battlemage-on-`xe` default — stalls long runs.

The A750's documented teardown crash (a CPU-pool worker segfaulting on exit) also happens here, as does `UR_RESULT_ERROR_UNINITIALIZED` at exit — both **after** a completed generation, so a run is judged by its `decode`/`output` lines, not its exit code.

The patch also adds all of this to `docs/INTEL_ARC.md`, including a row in the "what has been run" table, the header as a build dependency, and a correction to "not tested by anyone yet" (the Alchemist A-series is now tested — on `xe`, not `i915`).

*Attached: `strata-a770-xe-alchemist.patch` (4 files, +187/−7). Verified with `git apply --check` and `git am` against a fresh `main` clone at `6674a00`, then built with icpx.*

I personally didn't get this working I used Qwen3.8max