#2800 [Bug]: SYCL, 2x Arc Pro B70 layer split: --pipeline-windows 2 fails the first verify window (UR_RESULT_ERROR_UNINITIALIZED)

Description

## What happened

With `--pipeline-windows 2` on a two-card layer split, the SYCL engine loads, captures all of its verify window
graphs without an error, reads the prompt, and then fails the first verify window. The request ends with
`verify: error` and the engine stops with a SYCL exception:

```
strata serve: --pipeline-windows 2: two verifiers per stage, the stages overlap across windows (decode too; the GDN snapshots beside the window)
strata verify: captured the 1-token window (upload no error, sync no error)
... (16 captures: 1-4 tokens, two verifiers, two stages, all "no error")
strata trace: read 24 tokens (windows) in 347.1 ms
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)
```

The same build without `--pipeline-windows 2` (the same flags otherwise, at `--max-context 65536` and the default reserve) serves normally (deterministic output across
restarts, coherent replies, per-card health checks pass after each stop). The cards stay healthy after the crash
(per-card health check passes, no reset needed).

Command (the server's engine line, upstream main 296b59d2 built for bmg-g31; the last three flags were appended
to make the experts fit next to the second verifier, see below):

```
strata --serve --pack <ista-gsq-rco-iq3-xxs pack> --native Qwen3.8-Flash-Next-GSQ-RCO-IQ3_XXS-00001-of-00002.gguf
  --ple-gguf Qwen3.8-Flash-Next-GSQ-RCO-IQ3_XXS-00002-of-00002.gguf --ple-io direct --stream-experts
  --expert-profile data/expert-profile-sycl.bin --expert-cache auto --prefill 8192 --no-prefill-borrow --kv int8
  --spec 4 --mtp <mtp rt> --spec-min-p 0 --conversation-cache-mib 32768 --conversation-cache-slots 4
  --suffix-draft 0 --lookup-chain 0 --mtp-max-t 4 --layer-split 30
  --pipeline-windows 2 --max-context 32768 --vram-reserve-mib 512
```

Container environment as in docs/INTEL.md (`STRATA_VERIFY_DEVICE_PLAN=1`, `STRATA_VERIFY_NO_HOST=1`,
`STRATA_STAGER_THREADS=12`, `ONEAPI_DEVICE_SELECTOR=level_zero:*`).

Steps: start the server with the line above, send one chat completion (any short prompt, `max_tokens` 128). Both
starts that loaded failed this way at the first request: this one on main 296b59d2, and one earlier the same day
on a build of main plus an unrelated GDN change (`--max-context 200000`, `--prefill 8192`, same split).

What I expected: decode with the two verifiers overlapping, as the flag describes.

Side note (not the bug): at `--max-context 200000 --vram-reserve-mib 1024` the second verifier's 303 MiB means
909 experts no longer fit, and the engine refuses to start with the clear "REFUSED: ... experts are neither in
VRAM nor mirrored" message; at 32K context and a 512 MiB reserve everything fits. Without
`--pipeline-windows 2`, `--max-context 65536` with the default 1024 MiB reserve fits all 24,576 experts.

Possibly related: #1980 ends in the same `UR_RESULT_ERROR_UNINITIALIZED` abort at the first verify window on 2x B60,
but there it follows a `UR_RESULT_ERROR_DEVICE_LOST` with experts missing from VRAM. Here every expert is resident, no
DEVICE_LOST is logged, and the run without `--pipeline-windows 2` is fine.

Not tested: a single B70, other Intel cards, other models, `--pipeline-windows 1`. I did not look for the failing
call inside the verify path.

## Strata version, GPU, OS

Strata main 296b59d2 (0.1.44.1), SYCL build (bmg-g31 AOT, oneAPI DPC++ 2026.1.0), 2x Intel Arc Pro B70 32 GB,
Ubuntu 24.04 host kernel 7.1.0, Compute Runtime 26.22.38646.4 (libze-intel-gpu1), in the strata-sycl-dev style
container.

## Engine log

The engine log of the failing run, from the start line to the abort (local paths shortened):

```
[strata] 2026-10-16 17:45:49 engine started: strata --serve --pack <pack ista-gsq-rco-iq3-xxs> --native Qwen3.8-Flash-Next-GSQ-RCO-IQ3_XXS-00001-of-00002.gguf --ple-gguf Qwen3.8-Flash-Next-GSQ-RCO-IQ3_XXS-00002-of-00002.gguf --ple-io direct --stream-experts --expert-profile data/expert-profile-sycl.bin --expert-cache auto --vram-reserve-mib 1024 --prefill 8192 --no-prefill-borrow --max-context 200000 --kv int8 --spec 4 --mtp <mtp rt> --spec-min-p 0 --conversation-cache-mib 32768 --conversation-cache-slots 4 --suffix-draft 0 --lookup-chain 0 --mtp-max-t 4 --layer-split 30 --pipeline-windows 2 --max-context 32768 --vram-reserve-mib 512
strata generate: layer split across 2 GPUs: CUDA0, then CUDA1 (split 30)
strata generate: PCIe probe: 14.3 GB/s host->device (best of 14.2 14.2 14.3 14.3) -> pcie_frac 0.39 (default 0.55)
strata generate: this CPU has no AVX-512: the expert kernels run on AVX-2 (multi-token for the i-quant gate/up rows)
strata generate: native pack: <pack ista-gsq-rco-iq3-xxs> experts (largest blob 2.33 MB), token embedding IQ3_S in mapped host memory (260 MiB, 0.2 s)
strata generate: layer split: CUDA0 loads the dense weights of layers 0-29 only
strata generate: 943 MiB of weights loaded from <pack ista-gsq-rco-iq3-xxs> in 0.3 s (590 canonical tensors skipped: served natively)
strata generate: 187 native projection matrices, 1075.89 MiB of weights, in 0.4 s
strata generate: PLE direct read probe: 84334 rows/s (fine)
strata generate: the SSD is kept awake while rows are read: one page of the table after 100 ms without a read, until 60 s after the last request (STRATA_SSD_KEEPALIVE=0 turns it off)
strata generate: layer split: CUDA1 loads the dense weights of layers 30-47 only
strata generate: layer split: CUDA1 PCIe probe 14.2 GB/s (best of 13.8 14.0 14.2 13.8) -> pcie_frac 0.39
strata generate: layer split: CUDA1 holds its weights; 30.48 GiB free (its session follows the split search)
strata generate: profile data/expert-profile-sycl.bin: 24576 ranked pairs, built for 24576 slots
strata generate: PLE on, table 320001536 rows of Qwen3.8-Flash-Next-GSQ-RCO-IQ3_XXS-00002-of-00002.gguf
strata generate: layer split: CUDA1 holds its weights, session [30, 48) and the head; 29.84 GiB free
strata mtp: draft layer loaded, 840 MiB of VRAM (experts 675, dense 111), files read in 0.42 s (1875 MiB/s)
strata generate: GPU 0: Intel(R) Arc(TM) Pro B70 Graphics, compute capability 20.2
strata generate: experts streamed from the GGUF on demand (--stream-experts): no host arena; every expert must fit the VRAM cache
strata generate: 15 pool workers on logical processors 1,2,3,4,5,6,7,8,9,10,11,12,13,14,15, the host thread on 0 (draining too) (--host-core first)
strata trace: 30134 MiB free after the weights, the session and the drafter
strata generate: --pipeline-windows 2: 303 MiB of CUDA0 kept out of the expert cache (the second verifier, the GDN snapshots)
strata generate: expert cache auto: 29.43 GiB free, 512 MiB reserved (+0 MiB for the draft head) -> 10676 slots
strata generate: expert cache 15360 slots, 22.68 GiB of VRAM; policy is
strata trace: 6909 MiB free after opening the expert cache
strata generate: the GPU computes the experts in the cache; it rounds differently from the CPU,
                 so a reply can differ slightly from a run without the cache (same quality:
                 bench/results/2026-09-27-cache-parity).
                 PROFILE, ranked by routing frequency, no eviction.
strata generate: profile fill: 22.68 GiB in 2.4 s (10.02 GB/s, pipelined)
strata trace: 6909 MiB free after the profile fill
strata generate: pre-filled 15360 of 15360 slots from the profile in 2.4 s (10016 MB/s); slot 0 verified
strata generate: layer split, CUDA1: 29.01 GiB free of 31.89, room for experts 22.59 GiB
strata generate: layer split: CUDA1 runs layers 30-47, expert cache 9216 slots (17.29 GiB), 9216 of its 9216 profiled pairs; slot 0 verified
strata generate: layer split: CUDA0 runs layers 0-29
strata trace: 6909 MiB free after the R4 hit path
strata generate: R4 hit path ON - resident experts are computed on the GPU
strata generate: 15 expert-pool workers + the host thread
strata trace: 6909 MiB free after the expert cache and the graphs
strata generate: session is up (engine 0.1.44.1)
strata generate: token graph hit path: 24576 resident experts, decided on the device
strata serve: layer split: 100% 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 trace: 3513 MiB free after the head and the prompt path
strata verify: window up to 4 tokens, 58.1 MiB of device buffers (100% VRAM resident: zero-doorbell graph)
strata serve: layer split: layers 0-29 (CUDA0), 30-47 (CUDA1), one hand-off per window
strata verify: window up to 4 tokens, 58.1 MiB of device buffers (100% VRAM resident: zero-doorbell graph)
strata verify: window up to 4 tokens, 58.1 MiB of device buffers (100% VRAM resident: zero-doorbell graph)
strata verify: window up to 4 tokens, 58.1 MiB of device buffers (100% VRAM resident: zero-doorbell graph)
strata trace: 3387 MiB free after the verifier and the drafter's binding
strata serve: --pipeline-windows 2: two verifiers per stage, the stages overlap across windows (decode too; the GDN snapshots beside the window)
strata serve: 3240 MiB of VRAM free with everything loaded
strata trace: request 31 0
strata trace: prompt start 30 -1
strata verify: captured the 1-token window (upload no error, sync no error)
strata verify: captured the 2-token window (upload no error, sync no error)
strata verify: captured the 3-token window (upload no error, sync no error)
strata verify: captured the 4-token window (upload no error, sync no error)
strata verify: captured the 1-token window (upload no error, sync no error)
strata verify: captured the 2-token window (upload no error, sync no error)
strata verify: captured the 3-token window (upload no error, sync no error)
strata verify: captured the 4-token window (upload no error, sync no error)
strata verify: captured the 1-token window (upload no error, sync no error)
strata verify: captured the 2-token window (upload no error, sync no error)
strata verify: captured the 3-token window (upload no error, sync no error)
strata verify: captured the 4-token window (upload no error, sync no error)
strata verify: captured the 1-token window (upload no error, sync no error)
strata verify: captured the 2-token window (upload no error, sync no error)
strata verify: captured the 3-token window (upload no error, sync no error)
strata verify: captured the 4-token window (upload no error, sync no error)
strata trace: read 24 tokens (windows) in 347.1 ms
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)
```

Prepared with AI assistance (Claude Code): the reproduction, the log and this text were produced by Claude on the submitter's two Arc Pro B70s, at the submitter's request.