#1440 [Bug]: [SYCL] 0.1.40.2 on 2x Arc Pro B70: --layer-split fails three ways (#1054's fix from #1111 is not in main, plus two new split regressions)

本文

What happened:
Two Arc Pro B70s, Qwen3.8-Flash-Next **IQ3_XXS** (all 24,576 experts fit across the two cards), `--layer-split auto`, built from
`main` at **`e8ca9af`** (v0.1.40.2) with `-DSTRATA_SYCL_AOT=bmg-g31`, run through `sycl/serve/strata-sycl.sh`. It fails in three
separate ways, one after another as each is worked around. The same model on the same machine **works** with the 0.1.39 port at
`a5682fe` (the commit the 2x B60 community report used): 100% of the experts resident, 4/4 correctness checks, greedy output
repeatable, 56–82 tok/s decode.

### 1. The host mirror still takes the later stages' layers (#1054, fixed in PR #1111 but not in `main`)

`main`'s re-migration (`23a38a0`) does not carry #1111's fix. The mirror's second loop in `sycl/src/program/generate.cpp` still walks
every layer (`for (int64_t l = 0; l < g.n_layers; ++l)`), against stage 0's cache only, before the later stages' caches exist. On
IQ3_XXS it mirrors 11,317 experts / 20.5 GiB of CUDA1's layers. On UD-IQ4_XS it tries to pin 39 GiB in one `sycl::malloc_host`, which
fails, because a single host allocation on this card tops out below ~30 GiB (a probe got 24 GiB and failed at 32 GiB, with or without
`UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1`). With `STRATA_VERIFY_NO_HOST=1` the engine then refuses:

```
strata generate: mirroring the experts missing from VRAM: mirror: no pinned host memory for 39264 MiB (they stay on the SSD)
strata generate: REFUSED: 16983 experts are neither in VRAM nor mirrored; with STRATA_VERIFY_NO_HOST the device plan cannot run them ...
```

The planner on the same run had said `the caches hold 19833 of 24576 profiled pairs`. Limiting the loop to stage 0's layers (as
#1111 does with `mirror_end`) fixes it. The local patch used:

```diff
         for (int64_t l = 0; l < g.n_layers; ++l)                     // pairs the profile does not list at all
             for (int64_t e = 0; e < g.n_expert; ++e)
-                if (xcache.slot_of(l, e) == strata::core::kNotResident &&
+                if ((!multi_gpu || stage_of(l) == 0) &&
+                    xcache.slot_of(l, e) == strata::core::kNotResident &&
```

With it, IQ3_XXS mirrors only stage 0's own 1,589 misses (2.29 GiB).

### 2. `device_free_bytes()` charges the first card's allocations to the second card (`93b129a`)

After fix 1, CUDA1's cache still would not open:

```
strata generate: layer split, CUDA1: 26.97 GiB free of 31.89, room for experts 22.86 GiB
strata generate: layer split, CUDA1 expert cache: ExpertCache: ... = 18.23 GiB, but only 3.81 GiB of VRAM is free (28.08 GiB of 31.89 GiB total). Lower --expert-cache.; trying 8755 of 9728 slots (try 1 of 3)
...
```

The first line, from `get_memory_info` on the current device, is right. The second comes from `ExpertCache::open`, which replaces it
with `device_free_bytes()`. That function subtracts `own_drm_local_bytes()`, which sums `drm-total-local0` over **every** DRM fd of
the process, so on a split it is both cards' allocations together. CUDA0's 21.7 GiB cache is charged to CUDA1. The fdinfo fallback is
only needed when the driver reports the whole card free (the A750/i915 case in `93b129a`). The local patch used:

```diff
-    if (const unsigned long long own = own_drm_local_bytes(); own > 0 && total_b > own && total_b - own < free_b) {
+    if (const unsigned long long own = (free_b + (16ull << 20) >= total_b) ? own_drm_local_bytes() : 0ull;
+        own > 0 && total_b > own && total_b - own < free_b) {
```

A per-device fix would match `drm-pdev:` in fdinfo to the current device's PCI address instead.

With fixes 1 and 2 it loads: `layer split: 100% of the experts resident`, `4567 MiB of VRAM free with everything loaded`.

### 3. The first request kills the engine: graph across contexts

```
strata verify: captured the 6-token window (upload <FIXME: Placeholder>, sync <FIXME: Placeholder>)
Cannot submit to a queue with a dependency from a graph that is associated with a different context.Exception caught at file:/work/Strata/sycl/src/core/verify.cpp, line:1706
```

The server answers `the engine stopped unexpectedly (exit code 1); the next request restarts it`, and every retry repeats it.

**Likely cause, not yet tested:** #1111 hard-wires `sh_stream_on()` to `false` in the SYCL port ("A recorded SYCL graph takes only what
its own queue records"). `main`'s re-migration kept the CUDA default, which is `true` on a non-HIP build. So the shared expert forks
onto a side queue during window capture, and on a split that queue's context is not the stage's. `STRATA_SH_STREAM=0` (passed into the
container) should tell. `strata-sycl.sh` passes only a fixed set of `-e` variables, so it needs adding there.

### What works

`a5682fe` (0.1.39-sycl) runs this setup with two small local changes:
- `--remote-expert-opt` removed from the config, because that engine rejects it.
- The SYCL `resident_plan` definition takes the trailing `uint32_t* plan_err` parameter that `include/strata/kernels/verify_kernels.hpp`
  gained. Without it, linking fails with an undefined reference to `resident_plan(...)`.

It still mirrors CUDA1's layers (18.23 GiB pinned, unused), but with 72 GB of RAM that does no harm.

Results on `a5682fe`, IQ3_XXS, 32K context, INT8 KV, `--spec 4 --mtp`:

| test | decode | prompt |
|---|---|---|
| short prompt, 400 tokens | 56–61 tok/s | – |
| ~720-token prompt, 256 tokens | 82 tok/s | ~290 tok/s |
| ~3.2K-token prompt, 256 tokens | 51 tok/s | ~540 tok/s |

Checks: arithmetic, factual and code-execution prompts all passed. The generated `fib` passes `fib(50) == 12586269025`. Repeated greedy
runs gave identical output.

Also needed in an unprivileged Proxmox LXC: `--group-add 44 --group-add 993` on the `docker run` in `strata-sycl.sh`. Without them,
`sycl-ls` in the container finds no Level Zero devices and the engine stops at `No device of requested type available`. This could be
worth a line in `docs/INTEL.md`.

Strata version, GPU, OS:
0.1.40.2 (e8ca9af, AOT bmg-g31) fails; 0.1.39-sycl (a5682fe + plan_err link fix) works · 2x Intel Arc Pro B70 32 GB (8086:e223), each PCIe Gen4 x16 · Proxmox VE 9.2, kernel 7.0.14-17-pve, xe, in a Debian 13 LXC + Docker; strata-sycl-dev image (oneAPI 2026.1.0, Level Zero V2 1.17.39758); EPYC 7F52, 72 GB RAM to the container

Investigated with Claude Code, can run tests on dual B70's if needed.