#1980 [Bug]: 2x Arc Pro B60 layer split: DEVICE_LOST at the first verify window whenever experts are missing from VRAM; a stock setup_intel.py config lands at 96-97% resident and always hits it

本文

## What happened

A two-card SYCL layer split on 0.1.43 loses the device at the first speculative verify window every time the load is less than 100% resident with `STRATA_VERIFY_NO_HOST=1`. The same build, box and quant at 100% resident runs clean. The 2x B60 report (`bench/results/2026-10-04-community-2x-arc-pro-b60/README.md`, "Not done") lists "a run with `STRATA_VERIFY_NO_HOST=1` and experts outside VRAM" as untested. This is that cell, run seven times.

The stock two-card config from `sycl/setup_intel.py` lands in that cell. It writes neither `STRATA_STAGE_TRIM` nor an explicit `--layer-split`, so each card keeps a full copy of every layer's dense weights, 1.4 GiB per card on Flash-Next IQ2_XS, and the load stops at 96-97% resident. The report's clean run used a hand-written `iq2b.json` with `STRATA_STAGE_TRIM=1` and `--layer-split 25`.

Runs 2-6 on 0.1.43 (`7695859`), run 1 on 0.1.42 (`61b3fb5`); same two cards, same box, nothing else on the GPUs:

| Run | Quant | Context / reserve / KV | Resident | Result |
|---|---|---|---|---|
| 1 (0.1.42) | IQ3_XXS | 32768 / 1024 / int8 | 79% (19,481 of 24,576) | DEVICE_LOST at first verify window, 2 CAT faults |
| 2 | IQ2_XS | 32768 / 1024 / int8 | 97% (23,950 of 24,576) | DEVICE_LOST at first verify window, 5 retries, 5 CAT faults |
| 3 | IQ2_XS | 8192 / 1536 / int8 | 96% (23,487) | load only, no request sent |
| 4 | IQ2_XS | report's exact `iq2b.json` args and env | 100% | clean: 0 DEVICE_LOST, 0 CAT faults, 46.5 t/s median decode (n=9, 512-token replies, drafts accepted 73-77%) |
| 5 | IQ2_XS | run 2 plus `STRATA_STAGE_TRIM=1` only | 100% | load only; the later card freed 1.41 GiB, the first card did not move, `auto` moved the split to 24 |
| 6 | IQ2_XS | run 2 plus `STRATA_STAGE_TRIM=1` and `--layer-split 25` | 100% | load only; both cards freed ~1.4 GiB, 3,284 MiB free at load end |

Kernel log on every crash (7 of 7), always on the second-stage card (`0000:2f:00.0`, `level_zero:1`, layers 25-47 and the head):

```
[61412.776888] xe 0000:2f:00.0: [drm] Tile0: GT0: Fault response: Unsuccessful -ENOENT
[61412.776909] xe 0000:2f:00.0: [drm] Tile0: GT0: Fault response: Unsuccessful -ENOENT
[61412.776971] xe 0000:2f:00.0: [drm] Tile0: GT0: Engine memory CAT error [18]: class=ccs, logical_mask: 0x1, guc_id=2
[61412.777019] xe 0000:2f:00.0: [drm] Tile0: GT0: Engine memory CAT error [18]: class=ccs, logical_mask: 0x1, guc_id=2
[61412.778804] xe 0000:2f:00.0: [drm] Tile0: GT0: Engine reset: engine_class=ccs, logical_mask: 0x1, guc_id=2, state=0x249
[61412.778848] xe 0000:2f:00.0: [drm] Tile0: GT0: Timedout job: seqno=4294967184, lrc_seqno=4294967184, guc_id=2, flags=0x20 in strata [46941]
```

(Run 1, second attempt. Run 2's five faults print the same block with different seqnos.) No `CAT error` or engine reset on `0000:23:00.0` in the kernel log since boot.

Where the missing experts sat, from the engine log: the first card's shortfall is mirrored and the second card's is not logged as mirrored.

| Run | level_zero:0 short | its mirror line | level_zero:1 short | its mirror line |
|---|---|---|---|---|
| 1 | 1,976 (11,336 of 13,312) | `1976 of 1976 experts missing from VRAM mirrored in pinned host memory (2.83 GiB)` | 3,119 (8,145 of 11,264) | none |
| 2 | 115 (12,685 of 12,800) | `115 of 115 experts missing from VRAM ... mirrored in pinned host memory (0.15 GiB)` | 511 (11,265 of 11,776) | none |

Every fault is on the card whose missing experts have no mirror line. The source says why there is no line. `sycl/src/program/generate.cpp` builds the host mirror for the first stage only: `mirror_end = multi_gpu && !split_at.empty() ? split_at[0] : g.n_layers` (comment: "#1054, #1440: with a layer split the mirror holds only the first GPU's layers ... the later stages' caches do not exist yet"). `unmirrored_misses` is counted over the same `mirror_end` range, so the `REFUSED: N experts are neither in VRAM nor mirrored; with STRATA_VERIFY_NO_HOST ...` gate (the lines after `const bool no_host = ...`) sees 0 for a second stage that is short by 511 or 3,119 experts, and the run starts. The mirror table is attached to the device plan once, relative to the first stage's residency array (`resident_plan_set_mirror(thits.d_res, mirror_table_d)` in generate.cpp; `mir = g_mirror_table + (res_layer - g_mirror_res)` in `sycl/src/kernels/cuda/verify_kernels.dp.cpp`), so the second stage's plan has no mirror entry for its missing experts, and under `STRATA_VERIFY_NO_HOST=1` there is no host fallback for them either. Runs 1 and 2 are the only runs where the second stage was short; both faulted on that card at the first verify window. Run 4, with the second stage at 100%, ran clean.

A start-time check that counted the later stages' non-resident experts in `unmirrored_misses` would have refused runs 1 to 3 at start, with the existing message, instead of losing the device.

The engine log at the crash shows the verify window captured ("upload no error, sync no error") and then the device lost. The box recovered on its own each time; no reset was needed. Miss volume did not move the crash point: 5,095 missing experts (run 1) and 626 (run 2) both died at the first window.

Why the stock load is short: `sycl/src/program/generate.cpp` (`stage_trim = trim_asked && multi_gpu && !split_auto`) trims only when `STRATA_STAGE_TRIM=1` is set, and trims the first card only with an explicit `--layer-split`. `setup_intel.py` sets neither for a `gpu` list.

Memory lines, per card, at the "holds its weights" point (GiB free), runs 2 / 5 / 6 / report:

| Card | run 2 | run 5 | run 6 | report |
|---|---|---|---|---|
| level_zero:1 | 20.88 | 22.29 | 22.35 | 22.36 |
| level_zero:0 (cache auto free) | 20.57 | 20.57 | 21.93 | 22.11 |

Both cards report 23.91 GiB usable (`usable size exclude stolen 0x5fa000000`), the same as the report's cards.

## Suggested changes (not tested as a patch; no pull request opened)

1. Count the later stages' non-resident experts in `unmirrored_misses`, so a short second stage is refused at start with the existing `REFUSED` message instead of losing the device.
2. In `sycl/setup_intel.py` `write_run_script`, write `STRATA_STAGE_TRIM=1` for a two-card layer split, next to the existing `ONEAPI_DEVICE_SELECTOR` default:

```diff
+        if split:
+            cfg.setdefault("env", {}).setdefault("STRATA_STAGE_TRIM", "1")
```

Run 5 above is the load this produces on our box: 100% resident under `--layer-split auto`. It removes the trigger for a stock two-card install but not the gap in item 1.

## Strata version, GPU, OS

- Strata 0.1.43, commit `7695859`, SYCL AOT build, `STRATA_SYCL_AOT=bmg-g21` (CMake cache of `build-sycl-aot`), source build, no Docker
- 2x Intel Arc Pro B60 24 GB (`8086:e211`), PCI `0000:23:00.0` and `0000:2f:00.0`
- Ubuntu 24.04.5 LTS, kernel 7.0.0-38-generic (`xe`), NEO compute runtime 26.31.39395.14 (`intel-opencl-icd`, `libze-intel-gpu1`), Level Zero loader 1.32.0 (`libze1`), oneAPI DPC++ 2026.1.1 (2026.1.1.20260724)
- HP Z6 G4 Workstation, Xeon Platinum 8173M; the engine's own probe reads 6.8 GB/s host to device on each card (`PCIe probe 6.8 GB/s`)
- Report's box for comparison: Ubuntu 24.04, kernel 6.17, NEO 26.09.37435.12, oneAPI 2026.1.1 (README line 8)

## Engine log

Paths rewritten: `<data root>` is the models/packs/MTP directory, `<strata checkout>` the Strata source tree, `<path>` other local paths.

<details><summary>Engine log excerpts and both cfg.json files</summary>

#### Run 2 engine log (`iq2xs-2card-1011T1753`), residency lines through the DEVICE_LOST

```
strata generate: 115 of 115 experts missing from VRAM (+ 0 in the prompt loan's slots) mirrored in pinned host memory (0.15 GiB, 0.1 s); the GPU reads them over PCIe
strata generate: adaptive swaps off: with STRATA_VERIFY_NO_HOST an expert they move out of VRAM would be neither there nor mirrored
strata generate: layer split, CUDA1: 19.45 GiB free of 23.91, room for experts 15.41 GiB
strata generate: layer split: CUDA1 runs layers 25-47, expert cache 11265 slots (15.41 GiB), 11265 of its 11776 profiled pairs; slot 0 verified
strata generate: layer split: CUDA0 runs layers 0-24
strata generate: R4 hit path ON - resident experts are computed on the GPU
strata generate: 27 expert-pool workers + the host thread
strata generate: session is up (engine 0.1.43)
strata generate: token graph hit path: 23950 resident experts, decided on the device
strata serve: layer split: 97% of the experts resident, the prompt path's streamed ring 96 slots
strata serve: the prompt path allocates its own buffers (too few cache slots to borrow)
strata: copy engine: the prompt path's expert copies run on the GPU's copy engine (queue group 1), not the compute engine (STRATA_COPY_ENGINE=0 turns it off)
strata verify: window up to 6 tokens, 56.9 MiB of device buffers
strata serve: layer split: layers 0-24 (CUDA0), 25-47 (CUDA1), one hand-off per window
strata verify: window up to 6 tokens, 56.9 MiB of device buffers
strata mtp: draft head over 106299 tokens (137.9 MiB)
strata serve: 2050 MiB of VRAM free with everything loaded
strata verify: captured the 6-token window (upload no error, sync no error)
strata verify: captured the 6-token window (upload no error, sync no error)
level_zero backend failed with error: 20 (UR_RESULT_ERROR_DEVICE_LOST)
strata serve: verify: error
terminate called after throwing an instance of 'sycl::_V1::exception'
  what():  level_zero backend failed with error: 37 (UR_RESULT_ERROR_UNINITIALIZED)
```

#### Run 4 engine log (`iq2xs-2card-upexact-1011T1802`), residency lines

```
strata generate: layer split: CUDA1 holds its weights; 22.35 GiB free (its session follows the split search)
strata generate: layer split: CUDA1 holds its weights, session [25, 48) and the head; 21.92 GiB free
strata generate: expert cache auto: 22.09 GiB free, 1536 MiB reserved (+0 MiB for the draft head) -> 12800 slots
strata generate: layer split, CUDA1: 21.13 GiB free of 23.91, room for experts 17.92 GiB
strata generate: layer split: CUDA1 runs layers 25-47, expert cache 11776 slots (16.11 GiB), 11776 of its 11776 profiled pairs; slot 0 verified
strata generate: layer split: CUDA0 runs layers 0-24
strata generate: session is up (engine 0.1.43)
strata serve: layer split: 100% of the experts resident, the prompt path's streamed ring 96 slots
strata serve: 4222 MiB of VRAM free with everything loaded
```

#### Run 2 cfg.json

```json
{
 "exe": "<path>",
 "args": [
  "--pack",
  "<data root>/packs/iq2_xs",
  "--native",
  "<data root>/models/IQ2_XS/Qwen3.8-Flash-Next-GSQ-RCO-IQ2_XS-00001-of-00002.gguf",
  "--ple-gguf",
  "<data root>/models/IQ2_XS/Qwen3.8-Flash-Next-GSQ-RCO-IQ2_XS-00002-of-00002.gguf",
  "--expert-profile",
  "<strata checkout>/data/expert-profile-sycl.bin",
  "--expert-cache",
  "auto",
  "--spec",
  "4",
  "--mtp",
  "<data root>/mtp/rt",
  "--max-context",
  "32768",
  "--kv",
  "int8",
  "--remote-expert-opt",
  "--spec-min-p",
  "0.70",
  "--stream-experts",
  "--vram-reserve-mib",
  "1024",
  "--prefill",
  "4096"
 ],
 "cwd": "<strata checkout>",
 "tokenizer": "<data root>/packs/iq2_xs/tokenizer",
 "model_name": "x",
 "log": "<path>",
 "port": 8090,
 "gpu": [
  0,
  1
 ],
 "gpus_asked": true,
 "layer_split": "auto",
 "host": "127.0.0.1",
 "open_browser": false
}
```

#### Run 4 cfg.json

```json
{
 "exe": "<path>",
 "args": [
  "--pack",
  "<data root>/packs/iq2_xs",
  "--native",
  "<data root>/models/IQ2_XS/Qwen3.8-Flash-Next-GSQ-RCO-IQ2_XS-00001-of-00002.gguf",
  "--ple-gguf",
  "<data root>/models/IQ2_XS/Qwen3.8-Flash-Next-GSQ-RCO-IQ2_XS-00002-of-00002.gguf",
  "--expert-profile",
  "data/expert-profile.bin",
  "--expert-cache",
  "auto",
  "--stream-experts",
  "--prefill",
  "auto",
  "--spec",
  "4",
  "--spec-min-p",
  "0.5",
  "--mtp",
  "<data root>/mtp/rt",
  "--max-context",
  "8192",
  "--kv",
  "int8",
  "--vram-reserve-mib",
  "1536",
  "--no-prefill-borrow",
  "--layer-split",
  "25"
 ],
 "cwd": "<strata checkout>",
 "tokenizer": "<data root