#866 sycl: keep cudaStreamQuery's answer in the per-layer ring waits

本文

The SYCL port's per-layer ring wait gives up too early. With experts mirrored in host RAM, a layer whose GPU time passes 2 ms is reported as `layer N never rang (graph finished)` while the graph is still running. On one Arc Pro B60 the Coder IQ1_M run stopped at layer 15 every time.

## Where it broke

The CUDA code asks the stream whether it is still running, and only gives up when it is not (`src/core/verify.cpp:1458`):

```cpp
const cudaError_t q = cudaStreamQuery(cs_);
if (q != cudaErrorNotReady && *seq < want) { ...fail "never rang"... }
```

The migration turned the query into `const dpct::err0 q = DPCT_CHECK_ERROR(cs_->ext_oneapi_empty())`. `DPCT_CHECK_ERROR` evaluates to 0 whenever the expression does not throw, so `q` is always 0, `q != 1` is always true, and any ring that takes longer than 2 ms is a failure. The same pattern is in `session_run_token` (`sycl/src/core/session.cpp`) and `Verifier::run` (`sycl/src/core/verify.cpp`).

With every expert in VRAM the host does not wait layer by layer (`STRATA_VERIFY_NO_HOST` waits for the whole window), so the paths measured on the B70 never reached it. With part of the experts mirrored in RAM the host does wait, and the first layer slow enough to trip it was layer 15 of the Coder.

## The change

`q` is now 0 when the queue is empty and 1 while it is running, which is what the guard and the trace comment ("1 = still running") already expect:

```cpp
const dpct::err0 q = cs_->ext_oneapi_empty() ? 0 : 1;
```

Two sites, 8 lines added and 3 removed in the engine files, plus a `fixups.py` entry (#8) that produces the same code from a fresh dpct run, so a re-migration keeps it.

## Verified

One Arc Pro B60 24 GB, Coder IQ1_M, 4,042 of 12,288 experts mirrored in pinned RAM, `--max-context 8192 --kv int8 --spec 4`, 128 greedy tokens, AOT `bmg-g21`, Level Zero V2, NEO 26.09.37435.12, oneAPI 2026.1.1, Ryzen 5 5600 with 64 GB RAM:

- before: exit 1 at layer 15, `verify: layer 15 never rang (graph finished)`, 5 of 5 tries;
- after: exit 0 in 7 of 7 runs of builds that contain it, 11.2 to 11.9 tok/s, coherent Python.

The all-resident path is unchanged (two B60s, `--layer-split`, every expert in VRAM: Coder 56.6 tok/s, IQ2_XS 58.6 to 61.3, same as before the change). Setup and numbers are in the results report, #867.

## Notes

- The same bug was found independently in `verify.cpp` by maxious's fork (`maxious/Strata_SYCL`, branch `b70`, "`DPCT_CHECK_ERROR` discards the value of its expression", fixed there as an explicit `idle = cs_->ext_oneapi_empty()`). The `session.cpp` copy (`session_run_token`) is still unfixed in that fork and in `maxfridbe/Strata_B70` at the time of writing. This PR covers both sites on upstream's tree, with a reproduction on a card.
- #667 found that a spin over host memory may need a system fence. I tried it: `sycl::atomic_fence(seq_cst, system)` in `strata_spin_pause()`. Before this fix it fails at layer 15 the same way, so it is not the cause of this failure. After this fix the run with the fence did 11.4 tok/s and without it 11.2 to 11.9, so it made no difference here. It is not part of this PR.
- Decode on that one-card path is slow (11.9 tok/s). The engine's stage table shows about 3.9 ms per layer in `wait for rings`, and it falls with the device-side spin bound `kSpinMax` (25.1 tok/s at 2,000, 28.4 at 200). That looks like a separate problem, the GPU not seeing the host's flag store, and I have not changed it.
- `src/core/verify.cpp` has two more `q != cudaErrorNotReady` waits (`:1983`, `:2098`) with no SYCL copy yet. They need the same treatment when ported.
- The tree this was checked on also has the compile fix (#784, #809): at `6f32ec0` `sycl/` does not build without it.