andersson82 · Oct 12, 2026 · Oct 12, 2026
diff --git a/docs/INTEL.md b/docs/INTEL.md
@@ -69,6 +69,8 @@ questions are setup's own. What changes:
   64K up when the RAM holds the KV. A `model_switcher` or `sampling` block from an earlier config is kept.
   `run-<model>.sh` starts `sycl/serve/server_intel.py`.
 - **Arc A-series (`i915`, e.g. the A750 with 8 GB):** a different config, see "Arc A750 and the other Alchemist cards" below.
+- **An integrated Arc (`xe`, e.g. the Arc 140T):** no Docker, the engine runs on the PC's own oneAPI; see "Integrated
+  Arc" below.
 - **Engine settings:** the config's `"env"` block reaches the engine. `strata-sycl.sh` forwards every variable starting with
   `STRATA_`, `ONEAPI_`, `UR_`, `IGC_`, `SYCL_` or `ZES_` into the container and sets `STRATA_VERIFY_DEVICE_PLAN=1`,
   `STRATA_VERIFY_NO_HOST=1` and `STRATA_STAGER_THREADS=12` unless the environment says otherwise (for example
@@ -794,6 +796,87 @@ against 1,240), and on the B70 without `--stream-experts` the prompt is faster (
 regression test is `sycl/tools/test_prompt_deadlock.sh` (the 2,744-token prompt; it times out on the old build and passes on
 the fix).
 
+## Integrated Arc: the Arc 140T (`xe`, 2026-10-12, re-run 2026-10-22)
+
+Measured on the Arc 140T (PCI id 7d51, Xe-LPG+) of a Core Ultra 9 285H with 62 GiB of RAM, Fedora 44 (kernel 7.2.9,
+`xe` driver), oneAPI 2026.0 installed on the PC, with the Coder IQ1_M, engine 0.1.43 and again with engine 0.1.47
+(Strata 0.2.2) on 2026-10-22. The GPU's memory is the system
+RAM: it reports 58 GiB of global memory and no free-memory query.
+
+| | Arc Pro B70 (`xe`, 32 GB) | Arc 140T (`xe`, integrated) |
+|---|---|---|
+| where the engine runs | in the `strata-sycl-dev` image (`sycl/serve/strata-sycl.sh`) | on the PC's own oneAPI (`sycl/serve/strata-sycl-native.sh`), no Docker |
+| build | AOT `bmg-g31` | JIT (`BUILD_DIR=build-sycl bash sycl/tools/build.sh`) |
+| experts | `--stream-experts`: from the GGUF into VRAM, the rest in a pinned RAM mirror | the same flags; the cache and the mirror are both in the RAM |
+| how big the cache is | the card's free VRAM | `MemAvailable` from `/proc/meminfo` (`STRATA_SYCL_FREE_MIB` overrides it) |
+| prompt-slot borrowing | from 32K up | off |
+| `SYCL_CACHE_PERSISTENT` | 0 (it crashed on Xe2) | 1 works: the start takes 7 s instead of 79 s from the second start on (0.1.40) |
+
+- **The driver.** Only `xe` was tested. On a kernel that binds the iGPU to `i915`, `i915.force_probe=!7d51 xe.force_probe=7d51`
+  on the kernel command line moves it to `xe`.
+- **Build.** Install the oneAPI Base Toolkit (2026.0 here, Intel's installer, into `/opt/intel/oneapi`) and the GPU compute runtime of the
+  distribution (Fedora 44: `intel-compute-runtime`, `intel-level-zero`, `oneapi-level-zero`, and `cmake`, `ninja-build`). Then, in the checkout,
+  `BUILD_DIR=build-sycl bash sycl/tools/build.sh` (it sources `setvars.sh` itself). The build ends with `BUILD EXIT 0`.
+- **Free memory.** The 140T has no `ext_intel_free_memory`, also with `ZES_ENABLE_SYSMAN=1`, so the engine got 0 and the
+  expert cache refused to start. On a device with `host_unified_memory` it now reads `MemAvailable`. The `xe` driver keeps the
+  memory a GPU program freed in its own pool, which `MemAvailable` does not count until a reboot: on the first start after a
+  boot all 12,288 experts were cached, on later starts 10,496 to 10,870 (the rest mirrored). `STRATA_SYCL_FREE_MIB` sets the
+  figure by hand.
+- **No prompt-slot borrowing.** With it, the answer after a prompt that borrowed the cache's slots was token 0 (every expert
+  activation NaN from the second prompt chunk; 0.1.40, any context). The prompt now gets its own buffers on a device with
+  `host_unified_memory`. Borrowing saved about 7% of prompt time in a test without the bug (112 vs 105 tok/s, 0.1.40).
+- **The first decode window stalled** on every request until `STRATA_WINDOW_SYNC=1` (#2110): 10 of 10 in 0.1.43,
+  6 of 6 in 0.1.47.
+- **What fits.** The experts' file (shard 1) has to fit in the RAM beside everything else. The Coder IQ1_M and the IQ2_XS
+  (shard 1 36.5 GiB) ran with 10 to 13 GiB of RAM left; the IQ3_XXS (43.8 GiB) left 3 to 4 GiB and did not run well (0.1.40).
+  Programs that use the GPU's memory (another model server) take from the same RAM: stop them first.
+- **The launcher.** `sycl/serve/strata-sycl-native.sh` is the config's `exe`. It sources oneAPI, sets the port's defaults
+  (`STRATA_VERIFY_DEVICE_PLAN=1`, `STRATA_VERIFY_NO_HOST=1`, `STRATA_STAGER_THREADS=12`, `STRATA_NO_MULTI_GR=1`), each one
+  overridden by the config's `"env"` (`0` turns one off), and gives the engine an OOM score of 1000, so that a PC out of
+  memory loses the engine and not the desktop. `STRATA_SYCL_BIN` names another binary (relative to the checkout or absolute).
+
+**With setup.** `./setup.sh --backend sycl` finds the integrated GPU in sysfs (an Intel GPU at PCI `00:02.0` on `xe`) and needs
+oneAPI in `/opt/intel/oneapi` (or `ONEAPI_ROOT`) and the engine in `build-sycl` (or `build-sycl-aot`); Docker is not used. What
+changes from a card:
+
+- **Which size fits.** A model needs its shard 1 plus 20 GiB of RAM (the 14 GiB the measured runs used beside shard 1, and
+  6 GiB more). The size menu marks the sizes that need more, recommends the first one that fits, and asks before it installs
+  one that does not (`--model X --yes` installs it anyway). `--check` says the same per model.
+- **The config.** As the one below: the launcher as `exe`, the PC's own paths, `--stream-experts`, no KV streaming (the KV is in
+  the same RAM either way), the VRAM reserve as on a card (1,024 MiB up to 32K, 2,048 above), `--draft-vocab en`, and
+  `SYCL_CACHE_PERSISTENT=1` on the 140T only (it crashed on Xe2). An AOT build adds `STRATA_SYCL_BIN`.
+- **On `i915`** setup stops and names the kernel command line that moves the GPU to `xe`.
+
+On the 140T, `./setup.sh --backend sycl --family coder --model IQ1_M --context 98304 --no-start --yes` wrote the config below
+(less `"fit_max_tokens"`) and `run-coder-iq1_m.sh`.
+
+**Running it by hand.** A config like this, started with
+`.venv/bin/python sycl/serve/server_intel.py --engine strata --config strata-coder-iq1_m.json --port 8080`:
+
+    {
+     "exe": "<checkout>/sycl/serve/strata-sycl-native.sh",
+     "args": ["--pack", "<data>/packs/coder-iq1_m", "--native", "<data>/models/coder-IQ1_M/<shard 1>.gguf",
+              "--ple-gguf", "<data>/models/coder-IQ1_M/<shard 2>.gguf", "--expert-profile", "<checkout>/data/expert-profile-coder.bin",
+              "--expert-cache", "auto", "--prefill", "auto", "--spec", "4", "--spec-min-p", "0.5", "--mtp", "<data>/mtp/rt",
+              "--max-context", "98304", "--kv", "int8", "--stream-experts", "--vram-reserve-mib", "2048"],
+     "cwd": "<checkout>", "tokenizer": "<data>/packs/coder-iq1_m/tokenizer", "model_name": "qwen3.8-flash-next-coder-iq1_m",
+     "log": "<checkout>/strata-coder-iq1_m.log", "port": 8080, "backend": "sycl", "draft_vocab": "en", "fit_max_tokens": true,
+     "env": {"SYCL_CACHE_PERSISTENT": "1"}
+    }
+
+**Speed** (through the server, greedy, thinking off, with #2110):
+
+| | |
+|---|---|
+| decode, 100-token answers to short questions, 0.1.43 | 12.0 to 17.4 tok/s all experts cached, 11.5 to 17.2 with 1,792 of them mirrored |
+| the same, 0.1.47 (four runs of 10, 1,915 to 2,660 experts mirrored) | 8.4 to 15.2 tok/s; two runs of the same config gave 8.5-13.3 and 10.6-15.2 |
+| a 72,648-token prompt, 0.1.43 | read at 156.5 tok/s (7 min 44 s), the fact at its start recalled |
+| a 74,939-token prompt, 0.1.47 | read at 149.9 tok/s (8 min 20 s), the fact at its start recalled |
+| a 24-token prompt | 1.2 s (0.1.43), 2.3 s (0.1.47) |
+
+The 140T reads 90.5 GB/s from device memory and 76.1 GB/s from host memory (a read probe, 2026-10-05). Where a decode round's
+time goes is not profiled yet.
+
 ## Not done
 
 - Images on the llama.cpp config: not wired. On the SYCL port they are wired, see "Pictures on an Arc" below. Open there: a

diff --git a/docs/INTEL_ARC.md b/docs/INTEL_ARC.md
@@ -50,7 +50,8 @@ Full flags, per-request timings and engine logs are in the report linked in the
   The HIP build is not affected (the option is off by default).
 
 **Not tested by anyone yet:** Windows (no native build path; see below), Arc on WSL2 for the 0.1.39 port, integrated Arc
-GPUs (the B390 / Panther Lake in #515; the shared-memory planning does not exist yet), and pictures on a Battlemage card
+GPUs other than the Arc 140T on Linux (the B390 / Panther Lake on Windows in #515; the shared-memory planning does not
+exist yet; the 140T: [INTEL.md](INTEL.md), "Integrated Arc"), and pictures on a Battlemage card
 (the first picture on the Arc Pro B70 reset the GPU; they are off by default there, [INTEL.md](INTEL.md#pictures-on-an-arc-2026-10-19-branch-f-021-arcvis)).
 Pictures work on the A750 (on by default, read on the CPU). The Alchemist A-series was tested later, on an A750 (rows above and
 [INTEL.md](INTEL.md), "Arc A750 and the other Alchemist cards").

diff --git a/sycl/include/dpct/device.hpp b/sycl/include/dpct/device.hpp
@@ -12,7 +12,9 @@
 #include <algorithm>
 #include <array>
 #include <climits>
+#include <cstdlib>
 #include <cstring>
+#include <fstream>
 #include <functional>
 #include <iostream>
 #include <map>
@@ -533,8 +535,28 @@ class device_ext : public sycl::device {
   void get_memory_info(size_t &free_memory, size_t &total_memory) {
 #if (defined(__SYCL_COMPILER_VERSION) && __SYCL_COMPILER_VERSION >= 20221105)
     if (!has(sycl::aspect::ext_intel_free_memory)) {
-      std::cerr << "get_memory_info: ext_intel_free_memory is not supported." << std::endl;
+      // An integrated GPU (the Arc 140T measured here) has no free-memory query, not even with
+      // ZES_ENABLE_SYSMAN=1, and allocates from system RAM: what is free for it is what the system can still
+      // give, MemAvailable. Without this the figure is 0 and the expert cache refuses to start.
       free_memory = 0;
+#if defined(__linux__)
+      // STRATA_SYCL_FREE_MIB overrides it: the xe driver keeps freed GPU pages in a pool that MemAvailable does
+      // not count but the next GPU allocation reuses (measured: 6 GiB allocated, MemAvailable unchanged).
+      const char* free_mib = std::getenv("STRATA_SYCL_FREE_MIB");
+      if (free_mib != nullptr && get_info<sycl::info::device::host_unified_memory>()) {
+        free_memory = (size_t)std::strtoull(free_mib, nullptr, 10) << 20;
+      } else if (get_info<sycl::info::device::host_unified_memory>()) {
+        std::ifstream meminfo("/proc/meminfo");
+        std::string key;
+        size_t kib = 0;
+        while (meminfo >> key >> kib) {
+          if (key == "MemAvailable:") { free_memory = kib * 1024; break; }
+          meminfo.ignore(256, '\n');
+        }
+      }
+#endif
+      if (free_memory == 0)
+        std::cerr << "get_memory_info: ext_intel_free_memory is not supported." << std::endl;
     } else {
       free_memory = get_info<sycl::ext::intel::info::device::free_memory>();
     }

diff --git a/sycl/serve/strata-sycl-native.sh b/sycl/serve/strata-sycl-native.sh
@@ -0,0 +1,28 @@
+#!/usr/bin/env bash
+# strata-sycl-native.sh: the SYCL engine as the `exe` for serve/server.py, run on the host's own oneAPI (no Docker).
+# Made for an integrated Arc (measured on the Arc 140T, Core Ultra 9 285H, Fedora 44, oneAPI 2026.0). The server
+# runs it as `strata-sycl-native.sh --serve <args>` and talks to it over stdin/stdout; stderr goes to its log.
+#   STRATA_SYCL_BIN    the engine binary, relative to the repo or absolute (default: build-sycl/strata)
+#   ONEAPI_ROOT        the oneAPI install         (default: /opt/intel/oneapi)
+# The port's run-time switches below are defaults: the config's "env" block overrides them (a value of 0 turns one off).
+set -eo pipefail
+here=$(cd "$(dirname "$0")/../.." && pwd)                # the repo
+# setvars.sh reads unset variables, so -u only after it; its output would corrupt the serve protocol on stdout
+source "${ONEAPI_ROOT:-/opt/intel/oneapi}/setvars.sh" >/dev/null 2>&1 || true
+set -u
+export ONEAPI_DEVICE_SELECTOR=${ONEAPI_DEVICE_SELECTOR:-level_zero:0}
+export SYCL_CACHE_PERSISTENT=${SYCL_CACHE_PERSISTENT:-0}         # docs/INTEL.md: the persistent JIT cache crashed on Xe2
+export UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1                  # the expert cache is one 23 GiB allocation (> 4 GiB)
+# as strata-sycl.sh: the device-built verify plan and no host handshakes in it (xe); STRATA_NO_MULTI_GR: the generic
+# GR kernels (the multi-token ones decoded 15.1 vs 18.3 tok/s on the 140T in 0.1.40, same output)
+for kv in STRATA_VERIFY_DEVICE_PLAN=1 STRATA_VERIFY_NO_HOST=1 STRATA_STAGER_THREADS=12 STRATA_NO_MULTI_GR=1; do
+    k=${kv%%=*}
+    [ -n "${!k+x}" ] || export "$kv"
+    # the engine tests some switches by presence: 0 or empty means "not set"
+    case "${!k}" in 0|"") unset "$k" ;; esac
+done
+# The iGPU's memory is system RAM: if the machine runs out, the kernel should kill the engine, not the desktop
+echo 1000 > /proc/self/oom_score_adj 2>/dev/null || true
+bin=${STRATA_SYCL_BIN:-build-sycl/strata}
+case "$bin" in /*) ;; *) bin=$here/$bin ;; esac
+exec "$bin" "$@"