#2394 setup: Alchemist on xe gets UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1 (A770 expert cache 1,906 -> 4,540 slots, decode +45%)

本文

## Title
Issue: #1545, #1606 (Arc A770)

## Summary
On an Arc A770 16 GB under `xe`, Level Zero refuses a single device allocation above DG2's `max_mem_alloc_size` (~4 GiB) unless `UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1`. `--expert-cache auto` then falls back to what fits under it: **1,906 slots (3.64 GiB) instead of 4,540 (8.64 GiB)**, and decode drops by a third. Nothing in setup, `strata-sycl.sh` or the docs sets the variable today. This makes setup write it into an Alchemist card's config on `xe` (next to the FP64-emulation variables). The JIT build's `-ze-opt-greater-than-4GB-buffer-required` already keeps the kernels correct past 4 GiB, which is what makes the bigger arena safe.

## What changed
- `sycl/setup_intel.py` (`to_sycl`): `env["UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS"] = "1"` for an Alchemist card on `xe`. i915 (the A750) and the B-series are unchanged.
- `tools/test_setup_sycl.py`: Alchemist on `xe` gets it; a B-series card and an A750 on i915 do not. `python tools/test_setup_sycl.py`: 19 tests OK.
- `docs/INTEL.md`: one sentence in the ">4 GiB" bullet.

## How it was checked
Arc A770 16 GB, `xe`, Linux 7.0, oneAPI 2026.1 (JIT), Ryzen 7 3800X, 30 GB RAM, 2026-10-14. Same binary (`v0.1.44.1` + #2207 + #2215), Coder IQ1_M, `--resident-experts --expert-cache auto --max-context 131072 --kv int8 --kv-resident 20480 --vram-reserve-mib 1200 --prefill-borrow --prefill auto --spec 4 --mtp`, `STRATA_VERIFY_STEPPED=1 STRATA_PF_XMX=2`, without and with the variable:

| | without | with |
|---|---|---|
| expert cache | 1,906 slots, 3.64 GiB | 4,540 slots, 8.64 GiB |
| 8K prompt | 237.9, 363.9 tok/s (2 chunks) | 509.2, 547.1, 559.8 tok/s (1 chunk) |
| 8K decode | 16.51, 17.46 tok/s (hits 0.52) | 23.97, 26.75, 27.62 tok/s (hits 0.76) |
| 2K decode | 14.67, 14.71 tok/s (hits 0.37) | 23.20, 25.95 tok/s (hits 0.61) |

The generated text is a correct answer in both (decoded and read). The token ids differ between the two because the cache holds different experts (the CPU-vs-GPU rounding). For contrast, our 2026-10-08 fork build, from before #1602 landed, answered `!!!!` (token 0) with the variable set: the bigger arena needs the 64-bit-address build.

With setup's own `--stream-experts` config this card hit the same limit on Monday (`ExpertCache: cudaMalloc(4.86 GiB) failed ... trying a smaller expert cache: 1891 slots`).

Not tested: an A770 on i915 (#1606), which likely has the same limit but is left alone here; a Docker (`strata-sycl.sh`) run (the wrapper passes `UR_*` through).

🤖 Generated with [Claude Code](https://claude.com/claude-code)