#2282 [Bug] SYCL on a Meteor Lake iGPU (8086:7d55) won't start, then freezes: 3 small fixes resolves issues

Descrição

[strata-alloc-fail.log](https://github.com/user-attachments/files/33389649/strata-alloc-fail.log)
[strata-window-hang.log](https://github.com/user-attachments/files/33389648/strata-window-hang.log)
[strata-start-log.log](https://github.com/user-attachments/files/33389650/strata-start-log.log)

### What happened

On this integrated GPU the SYCL engine doesn't start, and once that is fixed it freezes after a few requests. There are three separate bugs, each with a small fix. With all three applied it runs steadily (see "What I tested").

| # | What you see | Fix |
|---|---|---|
| 1 | `expert cache auto: 0.00 GiB free ... -> 0 slots` | one line in `sycl/include/dpct/device.hpp` |
| 2 | `ExpertCache: cudaMalloc(33.02 GiB) failed` | `UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1` |
| 3 | engine stalls on the 3rd to 5th request | `STRATA_WINDOW_SYNC=1` |

Two more findings follow the bugs: the engine's free memory at start-up depends on `MemAvailable`, and setup misreads this GPU's memory size.
Related: #515 asked for support for a shared-memory Intel iGPU (Arc B390, on Windows). This report is the Meteor Lake iGPU on Linux.

#### System

ASUS NUC14RVH-B, Core Ultra 5 125H, Intel Arc iGPU (Meteor Lake-P, `8086:7d55`, `xe` driver), 64 GB RAM shared with the GPU, Ubuntu 24.04, kernel 7.0.0-38.
Strata 0.1.43 (commit `7695859a`), JIT build in the `strata-sycl-dev` Docker image (built from `llama-sycl-b70:stable`; oneAPI DPC++/C++ 2026.1.0, NEO 26.35.39758.10). Model: IQ2_XS, 8K context. Retested on v0.1.43.1 (`a0116b27`, 16K context).

#### What I tested

- **Bugs 1 to 3** were found on 0.1.43.
- **v0.1.43.1 with the patch** (engine rebuilt from that source; the log says `session is up (engine 0.1.43.1)`): 20 short requests in a row (about 7 to 12 tokens/s on the 15 to 17 token answers), six 300-token answers (33.7 to 34.8 s each), a 922-token prompt twice (32.5 s cold, 1.8 s repeated), and a request with thinking (151 tokens, 17.6 s). All completed.
- **v0.1.43.1 without the patch:** `./setup.sh --low-ram auto --backend sycl --check` still reports 0.2 GB of VRAM (see "Also"), and `set_integrated()` still has no caller in the source (bug 1).

#### Bug 1: free memory reads as 0

- **Symptom:** `expert cache auto: 0.00 GiB free ... -> 0 slots`, then `no VRAM is left for the expert cache`.
- **Cause:** `device_free_bytes()` reads free RAM only if the GPU is flagged as integrated, and nothing sets that flag. `sycl/include/dpct/device.hpp:201` is the definition of `set_integrated()`, and `grep -rn set_integrated sycl src include` finds no call, on v0.1.43.1 as well.
- **Fix:** in `get_device_info()` in `sycl/include/dpct/device.hpp`, mark the GPU as integrated, then rebuild:

  ```cpp
  prop.set_integrated(dev.has(sycl::aspect::ext_oneapi_is_integrated_gpu));
  ```

  The local patch below does the same for PCI id 7d55 only (`set_integrated(1)`), and that is the version I ran.
- **Note:** after the fix the log still prints `get_memory_info: ext_intel_free_memory is not supported.`, and the serve check can print `0 MiB of VRAM free ... LOW`. Other call sites query the device directly, so these lines are wrong on an iGPU, but they did not stop it from running.

#### Bug 2: the expert cache can't be allocated

- **Symptom:** `ExpertCache: cudaMalloc(33.02 GiB) failed`. The engine retries at 75% of the size each time (eight failures, down to 4.40 GiB), then tries 2,449 slots, and stops with `REFUSED: 21384 experts are neither in VRAM nor mirrored`.
- **Cause:** this GPU allows at most 4 GiB in one allocation (`clinfo`: `Max memory allocation 4294959104 (4GiB)`), and the expert cache is one 33 GiB allocation.
- **Fix:** set `UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1` (a Level Zero setting that lifts the 4 GiB limit). Suggest making it the default for integrated GPUs.

#### Bug 3: freeze after a few requests

- **Symptom:** the engine stalled during the 3rd request in one run and the 5th in the next (two and four requests had finished first), then `no progress for 60 s ... (verify window: waiting for the GPU to finish the window ...)`. The attached stall report shows all 8 expert-pool workers parked (sleeping) when it fires.
- **Cause (likely):** the hang described in the comment at `sycl/src/core/verify.cpp:2031` for the A750. The `STRATA_WINDOW_SYNC` fix for it is on by default only on i915 (line 2035), and this GPU uses xe.
- **Fix:** set `STRATA_WINDOW_SYNC=1`. Suggest making it the default for this GPU, or for xe.

#### How to apply the fixes

1. Apply the bug 1 change and rebuild the engine.
2. Add the two settings to the `env` block of `strata-<model>.json`:

   ```json
   "UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS": "1",
   "STRATA_WINDOW_SYNC": "1"
   ```

3. Before each start, run `sync; echo 2 | sudo tee /proc/sys/vm/drop_caches` (the next section explains why).
4. Start the model.

I also used `--vram-reserve-mib 2048`, `STRATA_UMA_HEADROOM_GIB=12` and `STRATA_MIRROR_MIB=1024`, which are my choices for this machine, not part of the fixes.

#### The engine's free memory at start-up depends on `MemAvailable`

Once the GPU is flagged as integrated, `device_free_bytes()` (`sycl/src/core/expert_cache.cpp`) takes `MemAvailable` less `STRATA_UMA_HEADROOM_GIB` (default 6, I set 12) as the free figure, and the expert cache is sized from it. If `MemAvailable` is low at start, the cache comes out too small and the start is refused.

**A start that was refused** (patched 0.1.43 engine, started by the v0.1.43.1 setup):

```
expert cache auto: 12.60 GiB free, 2048 MiB reserved (+143 MiB for the draft head) -> 6378 slots
...
REFUSED: 17162 experts are neither in VRAM nor mirrored; ... raise STRATA_MIRROR_MIB or the free RAM, or lower --max-context
```

Right after it failed, `free -h` showed 35 GiB used and `MemAvailable` was 27,441,636 kB (26.2 GiB), and `pgrep` listed no Strata process. I have not isolated what held the memory.

**The same config after `sync; echo 2 | sudo tee /proc/sys/vm/drop_caches`:** `MemAvailable` was 59,767,612 kB (57.0 GiB) and the next start read:

```
expert cache auto: 38.92 GiB free, 2048 MiB reserved (+143 MiB for the draft head) -> 24576 slots
expert cache 24576 slots, 33.02 GiB of VRAM
```

On the v0.1.43.1 engine the same command gave `MemAvailable` 59,150,540 kB (56.4 GiB) and a start with 38.17 GiB free and all 24,576 experts cached. That is about 6 GiB less than `MemAvailable` minus my 12 GiB headroom (44.4 GiB); I have not traced the difference.

**Over all starts:** `strata-iq2_xs.log` (attached as `strata-start-log.log`) holds 11 starts of this config. Eight read 37.65 to 39.74 GiB free and loaded all 24,576 experts. Three read 11.29 to 13.22 GiB free, got 5,697 to 7,131 cache slots, and were refused (`REFUSED: 16695`, `18134` and `17162 experts are neither in VRAM nor mirrored`). The last start is the v0.1.43.1 engine; the ten before it print `engine 0.1.43`.

Suggestion: warn at start-up when the free figure is far below `MemTotal` on an integrated GPU. The 0.1.40.2 fix for "host OOM at start on Linux when RAM was mostly file cache" looks related and may be worth extending to this case.

#### How to reproduce

I saw bugs 1 to 3 on 0.1.43. Start with the unmodified code to see bug 1. Add fix 1 to see bug 2. Add fix 2, start, and run this loop from a second terminal to see bug 3. Add fix 3 and the loop completes (it did on v0.1.43.1).

```
ok=1
for L in {A..T}; do
  curl -sf localhost:8085/v1/chat/completions -H 'Content-Type: application/json' \
    -d "{\"model\":\"x\",\"messages\":[{\"role\":\"user\",\"content\":\"Name one animal that starts with the letter $L.\"}],\"max_tokens\":24,\"chat_template_kwargs\":{\"enable_thinking\":false}}" \
    >/dev/null || { echo "failed at $L"; ok=0; break; }
done
[ "$ok" = 1 ] && echo "all 20 done"
```

#### Also: setup misreads this GPU's memory

`sycl/setup_intel.py` reads this GPU's memory size from PCI BAR 2 (0.25 GB), so setup treats it as a card with almost no memory.
Output of `./setup.sh --low-ram auto --backend sycl --check` on v0.1.43.1, unmodified and with the local patch:

| | Unmodified | With the local patch |
|---|---|---|
| GPU line | `Intel Intel GPU 7d55 (xe), 0.2 GB VRAM` | `Intel Arc Graphics (Meteor Lake-P, shared RAM), 45.8 GB VRAM` |
| Warning | "less than 12 GB of VRAM ... most experts stay on the CPU" | none |
| IQ2_XS | "fits with part of its experts mirrored in RAM (slower)" | "fits in VRAM" |
| IQ3_XXS | "fits with part of its experts mirrored in RAM (slower)" | "does not fit" |

With the patch, IQ3_XXS shows "does not fit" because the patch assumes no pinned mirror for this GPU. That is my choice for this machine, not a measured limit.

The patch below is what I run locally. It sets the integrated flag (bug 1), writes the settings for bugs 2 and 3 (plus my headroom and mirror values) into the generated config, and gives setup the right memory size. It only changes anything for PCI id 7d55.

<details><summary>Local patch (2 files, +15 / -3)</summary>

```diff
diff --git a/sycl/include/dpct/device.hpp b/sycl/include/dpct/device.hpp
index b55d2dcf..8fd8ddcc 100644
--- a/sycl/include/dpct/device.hpp
+++ b/sycl/include/dpct/device.hpp
@@ -345,6 +345,9 @@ static void get_device_info(device_info &out, const sycl::device &dev) {
   if (dev.has(sycl::aspect::ext_intel_device_id)) {
     prop.set_device_id(
         dev.get_info<sycl::ext::intel::info::device::device_id>());
+    // Meteor Lake-P iGPU (PCI 7d55) only: nothing else sets the integrated flag, so device_free_bytes()
+    // (core/expert_cache.cpp) never took its unified-memory branch and read 0 bytes free on this GPU.
+    if ((prop.get_device_id() & 0xffffu) == 0x7d55) prop.set_integrated(1);
   }
   if (dev.has(sycl::aspect::ext_intel_device_info_uuid)) {
     prop.set_uuid(dev.get_info<sycl::ext::intel::info::device::uuid>());
diff --git a/sycl/setup_intel.py b/sycl/setup_intel.py
index 8c0ed962..9dc2cc76 100644
--- a/sycl/setup_intel.py
+++ b/sycl/setup_intel.py
@@ -52,6 +52,7 @@ INTEL_ARC = {"e223": ("Arc Pro B70", 32.0), "e221": ("Arc Pro B60", 24.0), "e211
 # pinned mirror (--mmap-experts + STRATA_VERIFY_NO_HOST=0 - docs/INTEL.md, "Arc A750"). On xe (the A770) the mirror is
 # built in chunks past that limit and the A-series uses the ordinary streamed path; see to_sycl().
 ALCHEMIST_IDS = frozenset({"56a0", "56a1", "56a2", "56a5", "56a6", "5690"})
+MTL_IGPU = "7d55"   # Meteor Lake-P iGPU on xe: no VRAM, the GPU shares the RAM. Only this id takes the changes marked MTL.
 
 
 def intel_gpus():
@@ -75,6 +76,8 @@ def intel_gpus():
             if driver != "xe":                          # i915 without a known Arc id is an integrated GPU
                 continue
             name = f"Intel GPU {devid} (xe)"
+        if devid == MTL_IGPU and driver == "xe":        # MTL: BAR 2 is a 256 MB aperture, not memory; the GPU can address ~93% of
+            name, vram = "Arc Graphics (Meteor Lake-P, shared RAM)", max(S.ram_gb() * 0.93 - 12, 0.1)   # the RAM, less 12 for the OS
         if not vram:                                    # resource line 2 (BAR 2) is the VRAM aperture on xe cards
             try:
                 start, end, _ = (dev / "resource").read_text().splitlines()[2].split()
@@ -172,7 +175,8 @@ def to_sycl(cfg: dict, exe: Path, ram: float, keep: dict, vram_gb: float = 0.0,
         args += ["--kv-resident", "32768"]
     asked_reserve = flag(args, "--vram-reserve-mib")    # setup writes it only when the user gave it
     drop(args, "--vram-reserve-mib", True)
-    reserve = asked_reserve or (str(SMALL_RESERVE_MIB) if small_card(vram_gb, driver) else
+    reserve = asked_reserve or ("2048" if devid == MTL_IGPU else   # MTL: what is allocated after the expert cache
+                                str(SMALL_RESERVE_MIB) if small_card(vram_gb, driver) else
                                 "1024" if ctx <= 32768 else "2048")
     alchemist = driver == "i915" or devid in ALCHEMIST_IDS
     if driver == "i915":
@@ -183,7 +187,7 @@ d