#1602 sycl: make the Intel Arc A-series (Alchemist) work — prompt handshake, mirror, and the >4 GiB kernel-address bug

本文

## Title
sycl: make the Intel Arc A-series (Alchemist) work — prompt handshake, mirror, and the >4 GiB kernel-address bug

Issue: Resolves # (none)

## Summary
The SYCL base engine on an **Intel Arc A770** answered short prompts but produced garbage (`"!!!!!!!!!!!!!!!!"`)
for prompts longer than a few dozen tokens. This series fixes the A-series path end to end: the doorbell the host
polls, the expert mirror on Alchemist, the prompt path's own buffers, and — the actual cause of the garbage — a JIT
kernel-address bug where IGC turned a kernel's pointer arguments into 32-bit offsets from the allocation's base, so
kernels misread any bytes more than 4 GiB into an allocation while a memcpy of the same bytes was right.

Measured on an Arc A770 (16 GB, `xe`, PCI `56a0`), Qwen3.8-Flash-Next Coder IQ1_M, `--stream-experts` with a 4.46 GiB
VRAM expert cache: before, a 542-token prompt generated 16 tokens of `!` (every prompt answered token 0); after,
prompts of 73-542 tokens have 0 non-finite values (`STRATA_DBG_NAN=1`) and the 542-token prompt answers right
(`17 + 24 = 41`). Short prompts were correct throughout.

## What changed
Commits (oldest first):

- **`8f69d94` sycl: publish the doorbell with the uncached write hint** — a plain atomic store to host USM reaches the
  host only when the kernel ends on `xe`; the doorbell now stores through `write_hint<cache_mode::uncached>` so the
  host sees the sequence number mid-kernel (probe: 230 ms vs 2038 ms). Adds `sycl/probe/doorbell_store.cpp`.
- **`33064c3` sycl: chunk the expert mirror and raise memlock in the wrapper** — Alchemist refuses a single
  `sycl::malloc_host` past ~3 GiB; the mirror is built in chunks so a 17.6 GiB mirror loads on the A770.
- **`039992e` sycl: keep Alchemist's prompt path on its own buffers** — the prompt path no longer carves its buffers
  from borrowed expert-cache slots on A-series; this removes a prefill crash (`routed id out of range`, exit 139).
- **`c756d4b` sycl: add the doorbell payload checksum** — the publish kernels write a ring tag and a checksum of the
  payload next to the sequence word; the host waits for the whole payload before the CPU pool reads it. Robustness
  for the CPU-expert mode (not the A770 cure).
- **`aa71188` sycl: stage oneMKL GEMM operands that lie more than 4 GiB into their allocation** — oneMKL's own
  precompiled GEMM reads the last partial tile of an operand wrongly past 4 GiB (its kernels do not get the build
  option below). A small allocation registry (`device_offset_end`) plus staging in `Gemm::bf16`/`f16`;
  `STRATA_SYCL_GEMM_NO_STAGE=1` is the negative control.
- **`b40f69f` sycl: build Alchemist kernels for allocations past 4 GiB** — **the fix.** The JIT (non-AOT) build now
  passes `-ze-opt-greater-than-4GB-buffer-required` to the `spir64` backend (in the same argument as the
  divide/sqrt option; CMake drops a repeated `-Xsycl-target-backend=spir64`). AOT build unchanged.
- **`ebe06f9` sycl: setup picks the Intel card by PCI id and configures Alchemist** — `setup_intel.py` carries the
  PCI device id through so an A-series on `xe` is told from a Battlemage; every Alchemist id gets the driver's FP64
  emulation. On `i915` the experts still load into a CPU-computed RAM arena; on `xe` the A-series takes the ordinary
  streamed path (mirror chunked + the build option above).

Files: `sycl/CMakeLists.txt`, `sycl/include/strata/sycl_doorbell.hpp`, `sycl/include/strata/sycl_queue.hpp`,
`sycl/src/core/{expert_cache,layer,session,verify}.cpp`, `sycl/src/prefill/gemm.dp.cpp`,
`sycl/src/kernels/cuda/elementwise.dp.cpp`, `sycl/src/program/generate.cpp`, `include/strata/kernels/elementwise.hpp`,
`sycl/setup_intel.py`, `docs/INTEL.md`.

## Extra Notes
- The oneMKL staging (`aa71188`) did **not** fix the A770 garbage by itself: instrumentation showed no operand ever
  lies >4 GiB into its allocation in the reproducing config. It is kept as a defensive fix for the case the build
  option cannot cover (oneMKL's precompiled kernels), matching what the CUDA/HIP shims do. The `b40f69f` build option
  is the actual cure.
- The change is A-series-specific in effect but safe for other Intel cards: the backend option keeps kernels on
  64-bit addresses, and the doorbell/mirror/setup changes are gated by driver or device id.
- `tools/test_setup_sycl.py` passes (14 tests). Speed cost of the build option is not measured.
- Build: `docker run ... strata-sycl-dev "bash sycl/tools/build.sh strata"` (JIT). No AOT change.