#1798 sycl: --expert-parallel-device N, two Arc cards compute each layer's experts together

Beschreibung

## Title
Issue: none.

## Summary
An opt-in way for the SYCL port to use two Arc cards for each layer's experts: `--expert-parallel-device N`. Card 0 runs every layer, card N holds the experts card 0 does not, and for each layer of a decode window both cards compute their own experts at the same time. Without the option nothing changes.

Needs #1589 (the Verifier's null-stream sentinel, issue #1555), now in main: without it card N's verifier takes card 0's queue and its first capture fails.

### Measured
2x Arc Pro B70 (32 GB each, xe, Linux), Flash-Next IQ3_XXS, Ryzen 9 9950X3D, 60 GB RAM, through the server (`sycl/serve/server_intel.py`, 131,072 context, INT8 KV with `--kv-resident 32768`, `--vram-reserve-mib 2048`, the default drafting), greedy, thinking off, the same request 3 times (the 2nd and 3rd; the 1st reads the prompt):

| | `--layer-split auto` | `--expert-parallel-device 1` |
|---|---|---|
| where the experts live | layers 0-28 on card 0, 29-47 on card 1, every expert in VRAM | card 0: 283 per layer, card 1: the other 229, every expert in VRAM |
| decode, a BST class (600 tokens) | 94.3 tok/s | 107.9 tok/s (+14.4%) |
| decode, a 300-word explanation (384 tokens) | 64.2 tok/s | 72.6 tok/s (+13.1%) |
| output | | byte-identical to the layer split, all 6 requests of each |


### Why not 2x
Only the experts run on both cards. Stage times of a 4-token window (`STRATA_VERIFY_EAGER=1 STRATA_VERIFY_PROFILE=1`): the experts take 11.8 ms with the layer split and 8.2 ms with expert parallelism, the rest 41.0 and 40.7 ms. These run without the graph (the port's in-graph timestamps read 0), every kernel launched and timed from the host, so they are larger than with the graph and their ratio (52.8 / 48.9 = 1.08x) is not the decode's gain; they show which part got shorter. A unitrace of 96 decoded tokens shows what the handshake still costs card 0 per layer: its wait for card N's rows 26 us (card N's inbox copy, plan and row send are on its path), the send 3.5 us. In the same trace each card's expert kernels take 160-170 ms of the trace's 1.1 s.

## What changed
- `sycl/src/program/generate.cpp`: the option and its checks (a second card, `--stream-experts`; not with `--layer-split`, `--peer-device` or `--pipeline-windows`). Each card's set of experts is fixed at the start: each layer's experts are dealt to the two cards hottest first, in pairs alternating 0,1 / 1,0, until card 0's per-layer quota is full, so both cards serve about the same share of each layer's routing. Card N gets an `ExpertCache` of its own, filled from the GGUF. Per-layer slots, no prompt loans and no adaptive swaps, since either would leave an expert on neither card mid-request.
- `sycl/src/core/verify.cpp`, `sycl/include/strata/core/verify.hpp`: `Verifier::set_ep`. Card N's half of each window is a graph of its own, recorded beside card 0's and launched just before it; the same for the `--batch` graphs. Per layer, card 0 pushes the routed ids and the q8_1 MoE input into card N's inbox; card N plans and computes its entries and writes their rows straight into card 0's `parts_`; card 0 waits for that before its combine, which is the single-card one. After each window, `ep_check` reports a wait that ran to its bound or an expert on neither card.
- `sycl/src/kernels/cuda/ep_kernels.dp.cpp`, `sycl/include/strata/kernels/ep_kernels.hpp` (new): the handshake kernels. The flags hold an epoch each card's graph advances on the device (a host value would freeze at capture), and every spin is bounded. Two orderings mattered, each found by a greedy answer that changed between identical requests: data the other card wrote is read past the caches (an inbound PCIe write does not invalidate the reader's caches), and card N writes its rows and raises its flag in one kernel (`ep_send_rows`); with the flag in a later kernel, 2 of 6 requests read a row that was not all there yet.
- `sycl/include/strata/core/ep_mem.hpp` (new): the regions one card writes and the other reads. dpct gives each device its own `sycl::context`, so each region is a Level Zero allocation of its own on its owner, opened on the other card through an IPC handle in the same process.
- `sycl/src/kernels/cuda/verify_kernels.dp.cpp`: `resident_plan_ep`, the device plan for one card of the pair. It reuses `resident_plan_kernel` with two new parameters; with them, entries the other card owns are left out of the groups instead of being a plan error. Their defaults keep the old behaviour, and the existing calls do not pass them, so `resident_plan` is unchanged.
- `sycl/src/kernels/resident_plan_ep_parity.cpp` (new, in the parity tests): both cards' plans against a host replay.
- `sycl/CMakeLists.txt`: the new kernel file and the parity test.
- `docs/INTEL.md`: a section with the table above, how it works, what it needs and the checks.

## How it was checked
- `resident_plan_ep_parity`: 420 cases, 0 failures; `resident_plan_parity`: 641 cases, 0 failures (GPU, synthetic, no model). It checks that the two plans together cover every entry once, an expert on neither card sets the plan error, and with every expert on card 0 the plan is `resident_plan`'s word for word.
- `STRATA_EP_VERIFY=1` (with `STRATA_EP_MIRROR=1`) has card 0 compute every row itself and compare card N's rows with them word for word: no row differed in the 448 windows the three runs reported. `STRATA_EP_SELF=1` runs the window on card 0's own rows only; its output is byte-identical to the expert-parallel one.
- Output with the option: byte-identical to `--layer-split` for both prompts, all 6 requests of each (table above).
- `--batch 4` with the option (server, 128K context, this branch's build): 4 requests at once (512 tokens each) gave the same output as one request alone.
- Without the option: the same tokens as main's binary, one card with the RAM mirror, `--tokens` with a 40-token prompt, greedy, 32K context: T=1 (`--spec 2 --mtp-max-t 1 --suffix-draft 0`, 256 tokens) and MTP (`--spec 4 --spec-min-p 0.5 --mtp`, 512 tokens), identical output hashes.
- SYCL build (AOT `bmg-g31`): all targets including the parity tests, and `strata` again after the last change (a comment and the docs).

## Extra Notes
- The whole diff was reviewed by a human (the author).
- Only the SYCL port is touched; CUDA and HIP are unchanged.
- The prompt path reads card N's experts from the GGUF on card 0 (a 1,996-token prompt: 719 tok/s against 768 on one card with the RAM mirror). `STRATA_EP_MIRROR=1` also mirrors them in pinned RAM (18 GB for IQ3_XXS).
- A run without `--mtp` ends with an error at exit after its output, as on main; #1642 fixes that and is separate.
- Tested only on 2x B70 under xe on Linux; not on Arc A-series cards (i915), on Windows, or with cards of two different kinds.

Assisted-by: Claude Code with Opus 5.5 medium