diff --git a/docs/INTEL.md b/docs/INTEL.md index beaa6ad79b9..1ca27e6a174 100644 --- 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-.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": "/sycl/serve/strata-sycl-native.sh", + "args": ["--pack", "/packs/coder-iq1_m", "--native", "/models/coder-IQ1_M/.gguf", + "--ple-gguf", "/models/coder-IQ1_M/.gguf", "--expert-profile", "/data/expert-profile-coder.bin", + "--expert-cache", "auto", "--prefill", "auto", "--spec", "4", "--spec-min-p", "0.5", "--mtp", "/mtp/rt", + "--max-context", "98304", "--kv", "int8", "--stream-experts", "--vram-reserve-mib", "2048"], + "cwd": "", "tokenizer": "/packs/coder-iq1_m/tokenizer", "model_name": "qwen3.8-flash-next-coder-iq1_m", + "log": "/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 index 34ffd184421..64d86aa03fc 100644 --- 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 index b55d2dcf08b..3d48547816f 100644 --- a/sycl/include/dpct/device.hpp +++ b/sycl/include/dpct/device.hpp @@ -12,7 +12,9 @@ #include #include #include +#include #include +#include #include #include #include @@ -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()) { + free_memory = (size_t)std::strtoull(free_mib, nullptr, 10) << 20; + } else if (get_info()) { + 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(); } diff --git a/sycl/serve/strata-sycl-native.sh b/sycl/serve/strata-sycl-native.sh new file mode 100755 index 00000000000..ce0b3d027a0 --- /dev/null +++ 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 ` 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" "$@" diff --git a/sycl/setup_intel.py b/sycl/setup_intel.py index 617ea2771a3..7281a966a9c 100644 --- a/sycl/setup_intel.py +++ b/sycl/setup_intel.py @@ -56,23 +56,77 @@ # built in chunks past that limit and the A-series uses the ordinary streamed path; see to_sycl(). ALCHEMIST_IDS = frozenset({"56a0", "56a1", "56a2", "56a5", "56a6", "5690"}) +# Integrated Arc GPUs (docs/INTEL.md, "Integrated Arc"): Intel's integrated GPU sits at PCI 00:02.0 (a discrete card is +# behind a PCIe port, never there). Its memory is the system RAM. Only the xe driver is supported: the port was +# measured there (an Arc 140T, Core Ultra 9 285H, 62 GiB); on i915 setup says how to boot with xe instead. The engine +# runs on the host's oneAPI (sycl/serve/strata-sycl-native.sh), not in the Docker image. +IGPU_SLOT = "0000:00:02.0" +INTEL_IGPU = {"7d51": "Arc 140T"} # the measured one; another integrated GPU is named by its PCI id +SYCL_NATIVE = ROOT / "sycl" / "serve" / "strata-sycl-native.sh" +ONEAPI = Path(os.environ.get("ONEAPI_ROOT") or "/opt/intel/oneapi") +# What a model takes on an integrated GPU beside the experts' file (shard 1): the KV, the prompt buffers, the engine, +# the OS and the desktop. Measured on the Arc 140T (62 GiB): the Coder IQ1_M and IQ2_XS (shard 1 36.5 GiB) ran with +# 10-13 GiB of RAM left, IQ3_XXS (43.8 GiB) left 3-4 GiB and did not run well, so shard 1 + 14 GiB is in use and +# 6 GiB more is the room asked for. +IGPU_REST_GIB = 20 + + +def igpu_need_gib(model: str) -> float: + """The RAM a model needs on an integrated GPU: its shard 1 (all but the 28.8 GB per-layer lookup table, which is + read from the SSD) plus IGPU_REST_GIB.""" + d = S.MODELS.get(model, {}) + return (d.get("download_gb", 0) - 28.8) * 1e9 / 2**30 + IGPU_REST_GIB + + +def igpu_fits(model: str, ram: float) -> bool: + return igpu_need_gib(model) <= ram + + +def _pci(card: Path): + """(vendor, device id, driver, PCI slot) of a /sys/class/drm card, None when unreadable.""" + dev = card / "device" + try: + vendor = (dev / "vendor").read_text().strip().lower() + devid = (dev / "device").read_text().strip().lower().replace("0x", "") + driver = os.path.basename(os.path.realpath(dev / "driver")) if (dev / "driver").exists() else "" + except OSError: + return None + return vendor, devid, driver, os.path.basename(os.path.realpath(dev)) + + +def i915_igpus(drm="/sys/class/drm"): + """The PCI ids of Intel integrated GPUs that run on i915 (which the port does not support on them).""" + out = [] + for card in sorted(Path(drm).glob("card[0-9]*")): + p = None if "-" in card.name else _pci(card) + if p and p[0] == "0x8086" and p[2] == "i915" and p[3] == IGPU_SLOT and p[1] not in INTEL_ARC: + out.append(p[1]) + return out -def intel_gpus(): - """Intel discrete GPUs from sysfs: vendor 0x8086 under the xe or i915 driver, named by PCI device id. VRAM comes - from the id table, else from the size of the card's VRAM BAR.""" + +def intel_gpus(drm="/sys/class/drm"): + """Intel GPUs from sysfs: vendor 0x8086 under the xe or i915 driver, named by PCI device id. A discrete card's VRAM + comes from the id table, else from the size of its VRAM BAR. An integrated GPU on xe is listed with "igpu" and + setup's unified-memory fields (uma: its memory is the RAM, less what the OS keeps).""" found = [] - for card in sorted(Path("/sys/class/drm").glob("card[0-9]*")): + for card in sorted(Path(drm).glob("card[0-9]*")): if "-" in card.name: # connectors (card0-DP-1) share the card's device continue - dev = card / "device" - try: - vendor = (dev / "vendor").read_text().strip().lower() - devid = (dev / "device").read_text().strip().lower().replace("0x", "") - driver = os.path.basename(os.path.realpath(dev / "driver")) if (dev / "driver").exists() else "" - except OSError: + p = _pci(card) + if p is None: continue + vendor, devid, driver, slot = p + dev = card / "device" if vendor != "0x8086" or driver not in ("xe", "i915"): continue + if slot == IGPU_SLOT and devid not in INTEL_ARC: + if driver == "xe": + ram = S.ram_gb() + shared = max(0.0, ram - S.UMA_OS_LEFT_GB) + found.append({"index": len(found), "name": f"Intel {INTEL_IGPU.get(devid, f'integrated GPU {devid}')}", + "vram_gb": shared, "arch": driver, "driver": driver, "devid": devid, "igpu": True, + "uma": True, "dedicated_gb": 0.0, "shared_gb": shared}) + continue # i915: install() says how to boot with xe name, vram = INTEL_ARC.get(devid, (None, 0.0)) if name is None: if driver != "xe": # i915 without a known Arc id is an integrated GPU @@ -95,12 +149,18 @@ def real_mem_have(mem_have, real_ram: float, fake_ram: float): return lambda model, ram, vram=0.0: mem_have(model, real_ram if ram >= fake_ram else ram, vram) -def sycl_engine(): - """The SYCL build: (binary, why-not).""" +def sycl_engine(igpu: bool = False): + """The SYCL build: (binary, why-not). An integrated GPU runs it on the host's oneAPI, so that must be installed.""" if S.WIN: return None, "the SYCL port runs on Linux only" exe = next((b for b in (ROOT / "build-sycl-aot" / "strata", ROOT / "build-sycl" / "strata") if b.exists()), None) + if igpu and not (ONEAPI / "setvars.sh").is_file(): + return None, (f"an integrated GPU runs the engine on this PC's Intel oneAPI, and {ONEAPI}/setvars.sh is missing " + "(docs/INTEL.md, \"Integrated Arc\"; ONEAPI_ROOT names another install)") if exe is None: + if igpu: + return None, ("it is not built (docs/INTEL.md, \"Integrated Arc\": BUILD_DIR=build-sycl bash " + "sycl/tools/build.sh, with oneAPI on this PC)") return None, "it is not built (docs/INTEL.md, \"How to build it\": Docker, then sycl/tools/build.sh in the image)" return exe, None @@ -388,15 +448,62 @@ def to_sycl(cfg: dict, exe: Path, ram: float, keep: dict, vram_gb: float = 0.0, return out +def to_sycl_igpu(cfg: dict, exe: Path, keep: dict, devid: str = "") -> dict: + """setup's config -> the one for an integrated GPU on xe, run by sycl/serve/strata-sycl-native.sh on the host's + oneAPI: the host's own paths (no container), the experts streamed from the GGUF (--stream-experts; the ones the + cache does not hold are mirrored in host memory, which is the same RAM), no KV streaming (the KV is in that RAM + either way), setup's prompt chunks. The engine sizes its cache from MemAvailable (the iGPU has no free-memory + query) and does not lend the prompt the cache's slots there (a prompt that borrowed decoded only token 0 on the + Arc 140T). A reserve the user asked for (--vram-reserve-mib) is kept.""" + args = list(cfg["args"]) + for f in ("--resident-experts", "--mmap-experts"): # setup's low-RAM mode is the CUDA engine's + drop(args, f) + drop(args, "--kv-resident", True) + if flag(args, "--ple-io") == "ram" and "--ple-io" not in (cfg.get("user_args") or []): + # setup's rule for it read the 1024 GB stand-in; here the 29 GB n-gram table would take that much of the GPU's + # own memory (the expert cache's). One the user wrote (user_args) is kept + drop(args, "--ple-io", True) + if "--expert-profile" in args and (ROOT / "data" / SYCL_PROFILE).is_file(): # as on a card on xe (to_sycl) + i = args.index("--expert-profile") + 1 + if args[i].replace("\\", "/").rsplit("/", 1)[-1] == "expert-profile.bin": + args[i] = args[i][:-len("expert-profile.bin")] + SYCL_PROFILE + ctx = int(flag(args, "--max-context") or 32768) + asked_reserve = flag(args, "--vram-reserve-mib") + drop(args, "--vram-reserve-mib", True) + drop(args, "--stream-experts") # a hand-written one (the docs' config has it): written once + args += ["--stream-experts", "--vram-reserve-mib", asked_reserve or ("1024" if ctx <= 32768 else "2048")] + out = {k: v for k, v in cfg.items() if k not in ("lib_dirs", "env", "vision", "gpus")} + out.update({"backend": "sycl", "exe": str(SYCL_NATIVE), "args": args}) + env = {} + if devid in INTEL_IGPU: + # the persistent JIT cache: the Arc 140T's start 79 s -> 7 s from the second start on. Not on Xe2 (it crashed + # there in the first compile, "Things that matter on this GPU"), so only on the measured GPU + env["SYCL_CACHE_PERSISTENT"] = "1" + if exe != ROOT / "build-sycl" / "strata": # strata-sycl-native.sh's default + env["STRATA_SYCL_BIN"] = exe.relative_to(ROOT).as_posix() + if env: + out["env"] = env + out.update(keep) + return out + + def install(argv) -> None: intel = [] if S.WIN else intel_gpus() if not intel: + on_i915 = [] if S.WIN else i915_igpus() + if on_i915: + d = on_i915[0] + S.fail(f"this PC's integrated Intel GPU ({d}) runs on the i915 driver; Strata's SYCL port runs an integrated " + "GPU on the xe driver only", + f"boot with xe for it: add i915.force_probe=!{d} xe.force_probe={d} to the kernel command line, " + "reboot and run this again (docs/INTEL.md, \"Integrated Arc\")") S.fail("no Intel Arc found (an xe or i915 card in /sys/class/drm)", "on an NVIDIA or AMD card, run ./setup.sh") - exe, why = sycl_engine() + igpu_only = all(g.get("igpu") for g in intel) # otherwise setup picks the discrete card (amd_rank) + exe, why = sycl_engine(igpu_only) if exe is None: S.fail(f"Strata's SYCL engine cannot be used: {why}", "docs/INTEL.md: build it, then run this again") for name in ("gpus", "amd_gpus", "amd_problem", "hip_vision", "build_engine_hip", "hipblaslt_table", "ram_gb", - "write_run_script", "start", "say", "main"): + "write_run_script", "start", "say", "main", "ask", "confirm_risk"): if not callable(getattr(S, name, None)): S.fail(f"setup.py has no {name}() any more: sycl/setup_intel.py needs updating for this setup.py") @@ -421,6 +528,7 @@ def install(argv) -> None: say = S.say fake_ram = max(real_ram, 1024.0) + sizes = {} # the size menu's numbers -> models (an integrated GPU only) alch = intel[0]["driver"] == "i915" or intel[0].get("devid") in ALCHEMIST_IDS def say_intel(msg=""): @@ -429,10 +537,35 @@ def say_intel(msg=""): msg = msg.replace("Your AMD GPUs:", "Your Intel GPUs:").replace("just run ./setup.sh", "just run ./setup.sh --backend sycl") msg = re.sub(r"\b(xe|i915) \(AMD: docs/AMD_HIP\.md\)", r"\1 driver (Intel Arc: docs/INTEL.md)", msg) where = ("the experts the card does not hold are computed on the CPU" if alch + else "the integrated GPU's memory is this RAM" if igpu_only else "the experts are streamed into VRAM") msg = msg.replace(f"RAM: {fake_ram:.0f} GB", f"RAM: {real_ram:.0f} GB ({where})") msg = msg.replace(f"({fake_ram:.0f} GB RAM +", f"({real_ram:.0f} GB RAM +").replace(f"you have {fake_ram:.0f} GB RAM", f"you have {real_ram:.0f} GB RAM") - if re.match(r"\s+\S+\s+needs ~\d+ GB RAM:", msg): # --check's CUDA verdicts: replaced by the Intel one + # setup's unified-memory words are an AMD APU's (a BIOS carve-out); an Intel iGPU has none worth naming + msg = re.sub(r"unified memory: 0\.0 GB BIOS carve-out \+ (\d+) GB shared with the RAM, \d+ GB usable by the GPU", + r"shared memory: \1 GB of the RAM usable by the GPU", msg) + if igpu_only and re.search(r"\[ok\]\s+KV streaming on:", msg): # to_sycl_igpu drops it: the KV is in the RAM + msg = re.sub(r"KV streaming on:.*", "KV streaming: not used on an integrated GPU (the KV cache is in the " + "same RAM either way)", msg) + if igpu_only and re.search(r"\[ok\]\s+\d+ GB of RAM: the \d+ GB n-gram table is kept in RAM", msg): # to_sycl_igpu drops it + msg = (" [ok] n-gram table: read from the SSD (on an integrated GPU the RAM is the GPU's memory; --ple-io ram " + "in the config's args keeps it in RAM)") + if igpu_only: + msg = msg.replace("If this AMD card also drives your desktop and the desktop", "If the desktop") + menu = re.match(r" (\d+)\) (\S+) ", msg) if igpu_only else None + if menu and menu.group(2) in S.MODELS: # the size menu: which sizes an integrated GPU holds + sizes[menu.group(1)] = m = menu.group(2) + msg = re.sub(r"\s+<- .*$", "", msg) # setup's notes are for its RAM rule (1024 GB here) + if not igpu_fits(m, real_ram): + msg += f" <- needs ~{igpu_need_gib(m):.0f} GB RAM on an integrated GPU, you have {real_ram:.0f}" + if re.match(r"\s+\S+\s+needs ~[\d.]+ GB of VRAM \+ RAM:", msg) and igpu_only: # --check, a Gyro size + msg = (f" {msg.split()[0]:8s} not tested on an integrated GPU (the SYCL engine's GPU path does not run " + "the Gyro types)") + elif re.match(r"\s+\S+\s+needs ~\d+ GB RAM:", msg) and igpu_only: # --check: the integrated GPU's verdict + m = msg.split()[0] + msg = (f" {m:8s} needs ~{igpu_need_gib(m):.0f} GB RAM on an integrated GPU: " + + ("fits" if igpu_fits(m, real_ram) else "does not fit")) + elif re.match(r"\s+\S+\s+needs ~\d+ GB RAM:", msg): # --check's CUDA verdicts: replaced by the Intel one m = msg.split()[0] d = S.MODELS.get(m, {}) shard1 = d.get("download_gb", 0) - 28.8 # all but the per-layer lookup table (read from disk) @@ -466,6 +599,28 @@ def build_engine_hip(gpu, llama, vision="none"): # RAM together, you have ...", never a refusal) therefore reads the real RAM, not the 1024 GB above. S.mem_have = real_mem_have(S.mem_have, real_ram, fake_ram) + def igpu_risk(model, explicit, yes): + S.confirm_risk(f"{model} needs about {igpu_need_gib(model):.0f} GB of RAM on an integrated GPU (its experts and " + f"everything else share the RAM) and this PC has {real_ram:.0f} GB: it would start with little " + "room left, or not at all", explicit, yes, f"{model} does not fit this integrated GPU's memory", + "choose a smaller size (setup's menu marks the ones that fit), or --model " + f"{model} --yes to install it anyway", " Install it anyway?") + + ask = S.ask + + def ask_intel(question, choices, default, yes): + """On an integrated GPU the size question recommends the size setup would (IQ3_XXS, else the first) only where + it fits, else the first that fits; a size that does not fit is asked again.""" + if not (igpu_only and question == "Which size?" and sizes): + return ask(question, choices, default, yes) + if not igpu_fits(sizes.get(default, ""), real_ram): + default = next((i for i, m in sizes.items() if igpu_fits(m, real_ram)), default) + answer = ask(question, choices, default, yes) + if answer in sizes and not igpu_fits(sizes[answer], real_ram): + igpu_risk(sizes[answer], False, yes) + return answer + S.ask = ask_intel + bench_tips = S.bench_tips def bench_tips_intel(args, env, ram, *a, **k): @@ -485,12 +640,20 @@ def write_run_script(model, cfg_path, port, open_browser=True): # setup.write_ split = isinstance(gsel, list) idx = gsel if isinstance(gsel, int) and 0 <= gsel < len(intel) else 0 sel = intel[idx] - fit_notes = [] - cfg = to_sycl(cfg, exe, real_ram, keep.get(Path(cfg_path).name, {}), sel["vram_gb"], sel["driver"], - sel.get("devid", ""), fit_notes=fit_notes, model=model, ctx_asked=asked_ctx, - auto_reserve=not asked_reserve, fit=not split) - for line in fit_notes: # one plain line: what was chosen and why (never a refusal) - S.say(" " + line) + if sel.get("igpu"): + if split: + S.fail("an integrated GPU cannot share a model with another card", "choose one card (--gpu N)") + iexe, iwhy = sycl_engine(True) + if iexe is None: + S.fail(f"Strata's SYCL engine cannot be used: {iwhy}", "docs/INTEL.md: build it, then run this again") + cfg = to_sycl_igpu(cfg, iexe, keep.get(Path(cfg_path).name, {}), sel.get("devid", "")) + else: + fit_notes = [] + cfg = to_sycl(cfg, exe, real_ram, keep.get(Path(cfg_path).name, {}), sel["vram_gb"], sel["driver"], + sel.get("devid", ""), fit_notes=fit_notes, model=model, ctx_asked=asked_ctx, + auto_reserve=not asked_reserve, fit=not split) + for line in fit_notes: # one plain line: what was chosen and why (never a refusal) + S.say(" " + line) if len(intel) > 1: # pin the SYCL device(s) # the Level Zero order followed the sysfs card order on the tested machines; a machine where it does not # needs an explicit ONEAPI_DEVICE_SELECTOR (kept if the user set one) @@ -544,10 +707,21 @@ def start_sycl(cfg_path, *a, **k): argv += ["--draft-vocab", "en"] # ~110 MiB less VRAM than the default CJK subset if "--vram-reserve-mib" not in argv: argv += ["--vram-reserve-mib", str(SMALL_RESERVE_MIB)] + if igpu_only: + if "--draft-vocab" not in argv: # the English subset: what the Arc 140T was measured with + argv += ["--draft-vocab", "en"] + model = next((argv[i + 1] for i, x in enumerate(argv[:-1]) if x == "--model"), + next((x.split("=", 1)[1] for x in argv if x.startswith("--model=")), None)) + if model in S.MODELS and not igpu_fits(model, real_ram) and "--check" not in argv: # before the download + igpu_risk(model, True, "--yes" in argv) sys.argv = [str(ROOT / "setup.py"), *argv] if "--backend" not in argv: sys.argv += ["--backend", "hip"] - if is_b_series(intel[0]["driver"], intel[0].get("devid", "")): + if igpu_only: # is_b_series counts any xe id that is not an Alchemist one, the 140T too: not a B-series card + if "--vision" not in argv: + sys.argv += ["--vision", "none"] + S.say(" Pictures are off on an integrated GPU: nobody has run them on one yet (--vision yes turns them on)") + elif is_b_series(intel[0]["driver"], intel[0].get("devid", "")): S.say(" " + B_SERIES_LINE) if "--low-ram" not in argv: sys.argv += ["--low-ram", "off"] diff --git a/sycl/src/program/generate.cpp b/sycl/src/program/generate.cpp index 2a4db4d163b..739fede4ea9 100644 --- a/sycl/src/program/generate.cpp +++ b/sycl/src/program/generate.cpp @@ -2311,9 +2311,19 @@ int main(int argc, char **argv) try { } } catch (...) { return false; } }(); + // Not on an integrated GPU either (the Arc 140T, Xe-LPG+, 2026-10-06): its "VRAM" is the system RAM, so the + // experts the cache cannot hold are mirrored in host memory the GPU reads and borrowing saves nothing - and + // there a prompt that borrowed decoded only token 0 after it (all expert activations NaN from the 2nd chunk; + // likely the Alchemist GEMM bug above). Its own buffers: the reference output at a 96K context. + static const bool igpu = [] { + try { return dpct::get_current_device().get_info(); } + catch (...) { return false; } + }(); if (!multi_gpu && alchemist) { o.no_prefill_borrow = true; if (o.prefill_auto && o.prefill_chunk > 2048) o.prefill_chunk = o.prefill_auto_max = 2048; + } else if (!multi_gpu && igpu) { + o.no_prefill_borrow = true; } else { o.no_prefill_borrow = o.max_context <= 32768; } diff --git a/tools/test_setup_sycl.py b/tools/test_setup_sycl.py index 98dac7aa73c..f813fffe1e8 100644 --- a/tools/test_setup_sycl.py +++ b/tools/test_setup_sycl.py @@ -278,7 +278,10 @@ def test_pictures_are_on_by_default_on_every_arc_and_a_battlemage_is_told_it_is_ self.assertIn("issues", m.B_SERIES_LINE) src = (ROOT / "sycl" / "setup_intel.py").read_text(encoding="utf-8") self.assertIn("B_SERIES_LINE", src) - self.assertNotIn('sys.argv += ["--vision", "none"]', src) # no B-series card has pictures switched off for it + install = src[src.index("def install("):] + self.assertEqual(install.count('sys.argv += ["--vision", "none"]'), 1) # no B-series card has pictures switched off for it, + self.assertLess(install.index("if igpu_only: #"), install.index('sys.argv += ["--vision", "none"]')) # only an integrated GPU + self.assertLess(install.index('sys.argv += ["--vision", "none"]'), install.index("elif is_b_series(")) self.assertIn("S.hip_vision = intel_vision", src) self.assertIn("S.build_vision_cpu(", src) # the CPU reader is compiled by setup's own step self.assertIn("build_vision_sycl(", src) # --vision gpu: the SYCL reader, when the oneAPI compiler is there @@ -317,5 +320,100 @@ def test_a_picture_that_cannot_reach_the_container_is_said_once(self): self.assertEqual(out.getvalue().count("/dev/shm has no room"), 1) +class IntegratedGpu(unittest.TestCase): + """An integrated Arc on xe (measured: the Arc 140T, PCI 8086:7d51, 62 GiB): found at PCI 00:02.0 with its memory + as the RAM, run natively on the host's oneAPI with the host's paths, sized by the RAM it shares. On i915 it is not + offered (setup says how to boot with xe). A discrete card at another slot is listed as before.""" + + mod = staticmethod(ToSycl.mod) + + @staticmethod + def drm(root: Path, cards): + """A fake /sys/class/drm: cards = [(slot, devid, driver)].""" + for i, (slot, devid, driver) in enumerate(cards): + dev = root / "devices" / slot + (dev).mkdir(parents=True) + (dev / "vendor").write_text("0x8086\n") + (dev / "device").write_text(f"0x{devid}\n") + (root / "drivers" / driver).mkdir(parents=True, exist_ok=True) + (dev / "driver").symlink_to(root / "drivers" / driver) + (dev / "resource").write_text("0x0 0x0 0x0\n0x0 0x0 0x0\n0x4000000000 0x43ffffffff 0x0\n") + (root / "drm" / f"card{i}").mkdir(parents=True) + (root / "drm" / f"card{i}" / "device").symlink_to(dev) + return str(root / "drm") + + def test_the_140t_on_xe_is_integrated_and_a_discrete_card_is_unchanged(self): + m = self.mod() + with tempfile.TemporaryDirectory() as d, mock.patch.object(m.S, "ram_gb", return_value=62.0): + found = m.intel_gpus(self.drm(Path(d), [("0000:00:02.0", "7d51", "xe"), ("0000:03:00.0", "e223", "xe")])) + igpu, card = found + self.assertEqual(igpu["name"], "Intel Arc 140T") + self.assertTrue(igpu["igpu"] and igpu["uma"]) + self.assertEqual(igpu["vram_gb"], 62.0 - m.S.UMA_OS_LEFT_GB) + self.assertEqual(card, {"index": 1, "name": "Intel Arc Pro B70", "vram_gb": 32.0, "arch": "xe", "driver": "xe", + "devid": "e223"}) + self.assertLess(m.S.amd_rank(card), m.S.amd_rank(igpu)) # setup picks the discrete card first + + def test_an_integrated_gpu_on_i915_is_not_offered_but_named(self): + m = self.mod() + with tempfile.TemporaryDirectory() as d: + drm = self.drm(Path(d), [("0000:00:02.0", "7d51", "i915")]) + self.assertEqual(m.intel_gpus(drm), []) + self.assertEqual(m.i915_igpus(drm), ["7d51"]) + + def test_the_fit_rule_follows_the_140t(self): + m = self.mod() + for model in ("IQ1_M", "IQ2_XS"): # ran with 10-13 GiB left at 62 GiB + self.assertTrue(m.igpu_fits(model, 62.2), model) + self.assertFalse(m.igpu_fits("IQ3_XXS", 62.2)) # 3-4 GiB left: does not run well + self.assertTrue(m.igpu_fits("IQ3_XXS", 96.0)) + + def cfg(self, extra=()): + return {"exe": "x", "cwd": "/home/u/Strata", "lib_dirs": [], "env": {"X": "1"}, + "args": ["--pack", "/home/u/Strata-data/packs/p", "--native", "/home/u/Strata-data/models/a.gguf", + "--expert-cache", "auto", "--prefill", "auto", "--max-context", "98304", "--kv", "int8", + "--kv-resident", "32768", *extra]} + + def test_config_runs_natively_with_the_host_paths(self): + m = self.mod() + c = m.to_sycl_igpu(self.cfg(), m.ROOT / "build-sycl" / "strata", {}, "7d51") + self.assertEqual(c["exe"], str(m.ROOT / "sycl" / "serve" / "strata-sycl-native.sh")) + self.assertNotIn("sycl_root", c) + self.assertEqual(c["args"][c["args"].index("--pack") + 1], "/home/u/Strata-data/packs/p") + self.assertIn("--stream-experts", c["args"]) + self.assertNotIn("--kv-resident", c["args"]) + self.assertEqual(c["args"][c["args"].index("--prefill") + 1], "auto") + self.assertEqual(c["args"][c["args"].index("--vram-reserve-mib") + 1], "2048") + self.assertEqual(c["env"], {"SYCL_CACHE_PERSISTENT": "1"}) + self.assertEqual(c["backend"], "sycl") + + def test_a_hand_written_stream_experts_is_written_once(self): + m = self.mod() + c = m.to_sycl_igpu(self.cfg(["--stream-experts", "--vram-reserve-mib", "2048"]), m.ROOT / "build-sycl" / "strata", {}, "7d51") + self.assertEqual(c["args"].count("--stream-experts"), 1) + self.assertEqual(c["args"].count("--vram-reserve-mib"), 1) + + def test_setups_ple_io_ram_is_dropped_and_a_users_kept(self): + m = self.mod() + c = m.to_sycl_igpu(self.cfg(["--ple-io", "ram"]), m.ROOT / "build-sycl" / "strata", {}, "7d51") + self.assertNotIn("--ple-io", c["args"]) # setup's RAM rule read the 1024 GB stand-in + cfg = self.cfg(["--ple-io", "ram"]) + cfg["user_args"] = ["--ple-io", "ram"] + c = m.to_sycl_igpu(cfg, m.ROOT / "build-sycl" / "strata", {}, "7d51") + self.assertEqual(c["args"][c["args"].index("--ple-io") + 1], "ram") + + def test_the_persistent_cache_only_on_the_measured_gpu_and_an_aot_build_named(self): + m = self.mod() + c = m.to_sycl_igpu(self.cfg(["--vram-reserve-mib", "500"]), m.ROOT / "build-sycl-aot" / "strata", {}, "64a0") + self.assertEqual(c["env"], {"STRATA_SYCL_BIN": "build-sycl-aot/strata"}) + self.assertEqual(c["args"].count("--vram-reserve-mib"), 1) + self.assertEqual(c["args"][c["args"].index("--vram-reserve-mib") + 1], "500") + + def test_the_native_launcher_resolves_the_binary_in_the_repo(self): + src = (ROOT / "sycl" / "serve" / "strata-sycl-native.sh").read_text(encoding="utf-8") + self.assertIn('bin=${STRATA_SYCL_BIN:-build-sycl/strata}', src) + self.assertIn('case "$bin" in /*) ;; *) bin=$here/$bin ;; esac', src) + + if __name__ == "__main__": unittest.main()