#1890 sycl: Windows/OpenCL support for the merged port, and a live free-VRAM figure for the expert cache (#1549)

Description

## What this is





`main` now carries the SYCL port (`a1c47bc9`), built and measured on Linux with Level Zero, where oneMKL


and command graphs both work. This is the **Windows / OpenCL half** ΓÇö the pieces that backend needs and


that a Linux build never exercises ΓÇö plus the expert-cache sizing fix that came out of building it.





Single card, all measurements: **Arc Pro B70, 32 GB, Windows 11, driver 32.0.101.8976, OpenCL backend,


eager replay, engine 0.1.42** on this branch.





## What it changes





**1. The Windows build layer** (`sycl/CMakeLists.txt`, `sycl/tools/build.bat`,


`sycl/tools/onemath_config.hpp`, `sycl/serve/strata-sycl.bat`)





Compiler-ID gating (`/fp:precise` and `/W3` for MSVC, `-fp-model=precise` elsewhere ΓÇö the precise model


is load-bearing, `quantize_act_parity` goes from byte-exact to 303k mismatches without it),


`STRATA_SYCL_JIT_STATELESS` (the IGC >4 GiB surface-offset bug: an Arc A770, Coder IQ1_M, 4.46 GiB cache,


every expert in the last 0.46 GiB dequantized to NaN and every prompt answered `!!!!`), the oneMKL stub


for an installer-free conda install, and the DLLs that must sit beside `strata.exe`.





**2. Portability fixes in the migrated copies.** Upstream's migration was done on Linux, so several of


its files only build there. Each is a small, local change:





| file | fix |


|---|---|


| `sycl/src/core/pinned.dp.cpp` | restore the `STRATA_USE_CUDA` guard around the `cudaDeviceProp` block |


| `sycl/src/kernels/cuda/verify_kernels.dp.cpp` | rename `template <bool OUT>` → `WRITE_OUT` (a Windows SDK macro; `docs/INTEL.md` records it as "renamed, not worked around") |


| `sycl/src/core/gguf_expert_source.cpp` | Windows file I/O ΓÇö the `#ifdef _WIN32` guard and `portable_open` / `portable_pread` / `portable_close`, **on top of** upstream's new chunked mirror, which is kept |


| `sycl/src/program/generate.cpp` | rename the `small` identifier (also an SDK macro) and replace `memmem` with `std::search`. Identifier only — the earlier `fixups.py` did a global replace that also rewrote comments and strings (`"A small header binds"` → `"A small_chunk header binds"`), which is a bug in that tool and is deliberately not repeated |


| `sycl/src/core/verify.cpp` | restore the `STRATA_VERIFY_EAGER` guard on `capture_commit` ΓÇö without it an OpenCL run reaches `begin_recording` and dies on `"opencl backend is not supported by SYCL Graph extension"` |





**3. The plain-SYCL fallback GEMM, merged with upstream's staged GEMM** (`sycl/src/prefill/gemm.dp.cpp`,


+247 lines). This is the one that matters. Upstream has **no fallback at all**: on Linux/Level Zero


oneMKL works, so there is nothing to fall back to. On OpenCL oneMKL is not usable, so without this the


prompt path has no expert GEMM and the engine dies on the first request:





```


The oneAPI Math Kernel Library (oneMKL) Interfaces Project does not support this API.


prefill gemm: cublasGemmEx (staged): cuBLAS status 999


```





So this keeps upstream's `staged_gemm` (the oneMKL >4 GiB operand fix, measured on an A770) and gates the


oneMKL attempt on `onemkl_usable()`, with `fallback_gemm_tiled` behind it. A Level Zero build is


unchanged; an OpenCL build gets the fallback instead of a crash.





**4. A live free-VRAM figure** (`sycl/src/core/expert_cache.cpp`, +49). On this backend there is no


free-VRAM query at all ΓÇö `ext_intel_free_memory` needs Level Zero sysman, which the driver does not


expose ΓÇö so `dpct`'s `get_memory_info` answers from `STRATA_DEVICE_FREE_MIB`, a setup-time constant


`(vram_gb - 1.0) * 1024` = 31,744 MiB written by `setup_intel.py` from the registry. It is the same


number before and after the weights, the head and the draft layer are resident, because those are


allocated *before* the cache is sized. `own_gpu_local()` asks DXGI for this process's dedicated VRAM and


the OS's budget for it ΓÇö both live, and the budget already excludes whatever else holds the card ΓÇö and


`device_free_bytes()` prefers that figure. `sycl/probe/vram_pressure.cpp` measures both: the process's


figure tracks an allocation exactly (0.07 → 28.07 GiB for a 28 GiB fill) where `dpct`'s gives 0.00 and an


error. `STRATA_DXGI_FREE=0` takes the old path; that is the A/B.





**5. `STRATA_EXPERT_CACHE_HEADROOM=<fraction>`** (`include/strata/core/expert_cache.hpp`,


`src/program/generate.cpp`) ΓÇö the opt-in A/B knob, **default 0 everywhere**, so no build's default


sizing changes.





## What it leaves alone





Every CUDA and HIP default: the headroom is 0, and `onemkl_usable()` is true where oneMKL is, so cuBLAS


still takes the GEMM and the fallback is never reached. A Level Zero build takes `staged_gemm` exactly as


`main` does. `--expert-cache N` is never cut. Upstream's chunked expert mirror, the AVX1 CPU sources, the


iq/ple test registration and the `router_fast` / `payload_pairs` probes are all kept as they are on


`main`.





## How I checked it





**Build:** `cmake --build build-sycl --target strata`, conda `dpcpp_win-64` 2026.1.1 + VS18 Build Tools,


Ninja. Clean, `BUILD EXIT 0`.





**The sizing fix, before and after, same binary, three interleaved reps with a 30-second settle,**


Coder IQ1_M, 2,048-token prompt then 128 greedy tokens:





| | slots | prompt | decode |


|---|---|---|---|


| the constant (`STRATA_DXGI_FREE=0`) | 12,288 | 73.8 tok/s | **13.0 tok/s** |


| the DXGI figure | 11,348 | **113.1 tok/s** | **56.3 tok/s** |





The arithmetic says why: `auto` was told 31.00 GiB, sized a 28.82 GiB cache on a 31.16 GiB card and then


allocated 6.8 GiB of weights, KV, draft head and buffers on top of it ΓÇö 35.6 GiB of demand against a


31.16 GiB card, so the driver evicted and every kernel touching an evicted page crawled. Under


`STRATA_PREFILL_SYNC=1` the slowdown is spread across ordinary kernels rather than the expert path:


`gdn` 55x, the QSA projections 42x, the GDN output projection 53x, the two expert GEMMs 39x, the whole


prompt 36x. With the live figure `auto` is told 23.84 GiB and sizes 11,348 slots / 21.63 GiB, and both


paths are faster.





**The fallback:** with the figure live the engine fills its profile mirror, borrows its prompt slots and


runs the prompt path through the fallback with oneMKL bypassed ΓÇö the `oneMKL ... does not support this


API` crash above is gone. Prompt reading for the Coder at 2,048 tokens is 113 tok/s, decode 56.3.





## What is not done, and it is a real gap

**The MTP eager-replay path needs a bigger VRAM reserve on this card, and I have not made that automatic.**

Commit `b8280a90` ports the eager-replay refactor (`record_*` bodies with replaying `capture_*`
wrappers, the same shape `verify.cpp` uses). It builds, and a run with `--spec 4 --mtp` on the B70
completes at prompt 107.0 tok/s / decode 55.7 tok/s (2,048-token prompt, 128 tokens, DXGI figure sizing
11,348 slots / 21.63 GiB). The `mtp: begin capture` failure this pull request was opened with is gone.

What is not solved is the crash that follows when the card is nearly full. With a `--max-context 262144`
config and `--vram-reserve-mib 2048` the engine logs its own warning first:

```
strata serve: 0 MiB of VRAM free with everything loaded - LOW: requests may stall;
              add --vram-reserve-mib 2560 to the config's args (or lower --max-context)
```

and then dies with `0xC0000005` in `sycl::platform::platform(const sycl::platform&)` on a worker
thread - which is queue creation, i.e. an allocation that cannot be satisfied, not a logic error. Taking
the engine's advice (`--vram-reserve-mib 4096`) makes the same 262,144-context server come up and
answer, so the failure mode is understood well enough to avoid, just not prevented.

Why I am flagging it rather than fixing it: the DXGI figure sizes the cache from live memory, so it
*should* already account for the KV this config reserves, but the serve-side check that prints the LOW
warning still reads the old constant (`get_memory_info: ext_intel_free_memory is not supported`). So
there are now two free-VRAM figures in the port and they disagree, and the one guarding the request is
the wrong one. Making them agree is a change I would rather do deliberately than smuggle in here.

Reproduced on: Arc Pro B70, Windows, OpenCL, driver 32.0.101.8976, Coder IQ1_M, `--spec 4 --mtp`
262,144-context config; disappears with `--vram-reserve-mib 4096`. An instrumented build (fprintf inside
`record_round`) also escaped it, which is consistent with a timing-dependent allocation failure rather
than a logic one.

Also untested here: any NVIDIA or AMD card, a Level Zero build on Windows with this driver (it exposes no
user-mode adapter on my machine), and whether the cliff exists at all on another backend — an RTX 5070 Ti
A/B of `STRATA_EXPERT_CACHE_HEADROOM` found no difference there, which is consistent with `cudaMemGetInfo`
being live and the post-write check working on CUDA.

Closes the sizing half of #1549. This supersedes #1841, which is the same helper on a base that has since
been merged upstream.

Closes the sizing half of #1549. This supersedes #1841, which is the same helper on a base that has since

been merged upstream.



Closes the sizing half of #1549. This supersedes #1841, which is the same helper on a base that has since


been merged upstream.





Closes the sizing half of #1549. This supersedes #1841, which is the same helper on a base that has since


been merged upstream.