#3529 [Engine] A770 (i915) dual-card split: stale free-VRAM figure blocks the second stage; --pipeline-windows and xe+--stream-experts crash

Descripción

# Community benchmark: 2x Intel Arc A770 16 GB (i915), Core i5-13600K, 32*2 GB DDR4 RAM

Two consumer Alchemist cards on the `i915` driver, a CPU without AVX-512, and a
**dual-card layer split that works** — which the port's own notes list as broken
(`docs/A770_SPEED_SPEC.md`, "Dual-GPU on the A770 + B580 (peer tier / layer
split): BROKEN in the SYCL port"). It does run here, with two small fixes to the
free-VRAM accounting (attached). Reported per `docs/COMMUNITY_BENCHMARKS.md`.

> A note from the submitter: the point of this machine was to see what two
> 16 GB Alchemist cards can actually do with a 125B MoE, and to find out *why*
> it is slow. The answer turned out to be neither the compute nor the model —
> see "What limits this machine" below. The two fixes attached are what it
> took to get the dual-card path to start at all on i915.

## Hardware and software

| | |
| --- | --- |
| GPU | 2x Intel Arc A770, 16 GB each (DG2/Alchemist, PCI id `56a0`), **i915** driver |
| GPU-GPU | **No usable peer access.** `zeDeviceGetP2PProperties` reports `access=YES` for both directions, but a peer copy fails with `UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY` |
| PCIe | Card A `0000:03:00.0` in the CPU x16 slot: **21.2 GB/s** host->device measured. Card B `0000:0c:00.0` in a **chipset x4** slot: **2.9 GB/s** measured. Both are behind the cards' own PCIe switches (`8086:4fa0` upstream / `8086:4fa4` downstream), and `lspci`/`sysfs` report "2.5 GT/s x1" for both cards — the **measured transfer** is the only trustworthy figure here |
| CPU | Intel Core i5-13600K (6 P + 8 E cores, 20 threads), **no AVX-512** (AVX-VNNI only), so the expert kernels take the AVX-2 path |
| RAM | 62 GB |
| Storage | ZHITAI TiPro5000 NVMe |
| OS | Ubuntu 24.04, kernel `7.0.0-34-generic` |
| oneAPI | 2026.1 (`/opt/intel/oneapi`); Level Zero loader 1.21.9; compute-runtime 25.18.33578 |
| Strata | v0.1.42 source, **native SYCL build** (`icpx -fsycl -O2`, `-fsycl-device-code-split=per_kernel` as shipped), no Docker |

The engine was built natively (oneAPI 2026.1 extracted from the project's own
`strata-sycl-dev` image, since the Intel apt source measured 2.3 kB/s from this
network and no domestic mirror carries oneAPI). Needs `CPATH` and
`LIBRARY_PATH` pointed at `mkl/2026.1/include` and `mkl/2026.1/lib` — see
`FIXES.md`.

## Model and configuration

Qwen3.8-Flash-Next **GSQ-RCO IQ3_S** (3.5 bpw), repository
`ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-GGUF`, two GGUF shards
(54,817,524,224 + 28,800,138,432 bytes = 83.62 GB), downloaded from ModelScope
(`--source modelscope`), packed with `tools/iq_pack.py` into a native pack
(`dense.bin` 1.54 GB, 1079 tensors, 302 served natively). The MTP draft layer is
the stock one from `setup.sh` (`Strata-data/mtp/rt`).

```
--pack Strata-data/packs/iq3_s
--native ...Qwen3.8-Flash-Next-GSQ-RCO-IQ3_S-00001-of-00002.gguf
--expert-profile data/expert-profile.bin
--expert-cache auto --prefill auto --spec 4 --spec-min-p 0.5 --mtp Strata-data/mtp/rt
--max-context 32768 --kv int8 --resident-budget-gib 28 --pool-workers 9
--ple-io direct --vram-reserve-mib 300
--layer-split 23 --split-device 1
```

Environment (the `UR_*` variable is essential, see below):

```
ONEAPI_DEVICE_SELECTOR=level_zero:0,1      # the two A770s; * is not enough, it pulls the iGPU into the split
UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1
STRATA_VERIFY_NO_HOST=0                    # i915 path
IGC_EnableDPEmulation=1 OverrideDefaultFP64Settings=1 NEOReadDebugKeys=1
```

A second model was measured for comparison: Unsloth **UD-Q3_K_XL** (90.0 GB,
three shards, `unsloth/Qwen3.8-Flash-Next-GGUF`), same split and settings.

## Results

Greedy, 250–300 generated tokens per run, warm (second and later requests),
engine-version 0.1.42, one engine session per row unless noted.

| Configuration | Decode (tok/s) | Expert-cache hit | Caches (card A + card B) |
| --- | ---: | ---: | --- |
| **IQ3_S, dual card** (`--layer-split 23`) | **14.3 / 14.4 / 15.2** | 84 – 94 % | 9.34 + 7.75 GiB |
| UD-Q3_K_XL, dual card | 13.4 / 14.1 / 14.3 | 86 – 90 % | 8.70 + 6.93 GiB |
| UD-Q3_K_XL, single card (`level_zero:0`, context 4,096, `--resident-budget-gib 28`) | **5.9** | — | 3.73 GiB |

The single-card row is from the same machine and model but a **4,096-token
context** (a smoke test run before the split was working), so it is not a
matched-configuration comparison; the dual-card rows are 32,768. It is included
because it is the only same-hardware single-card figure available here.

Prompt reading is 20–29 tok/s. The split is `layers 0-22 (CUDA0), 23-47 (CUDA1)`,
one hand-off per window.

Worth noting: **cutting the quantization from ~4.5 effective bits (UD-Q3_K_XL) to
3.5 (IQ3_S) did not change decode speed** — the cache hit rate rose from ~87 % to
~94 % and the caches grew, and the throughput stayed in the same band. The
bottleneck is not the volume of expert data being fed (see below).

## What limits this machine

Measured during decode (`drm-engine-*` counters in the engine process's
`/proc/<pid>/fdinfo`, differenced over the whole request):

- **CPU: 2.0 % busy** (idle 98 %) over a 300-token decode. Not the CPU.
- **Card A 41 % / card B 48 % busy — and they run serially.** Their busy times
  *add up* to the whole wall clock (9.9 s + 11.7 s in a 24.4 s window), i.e. the
  layer split makes the cards take turns rather than overlap. The second card
  buys **expert-cache capacity** (which is real: 5.9 -> 14.5 tok/s), not
  parallelism.
- Disk is not a factor either: one 300-token request reads about 640 MB from the
  NVMe (measured on the block device).
- **A single allocation above 4 GiB fails** (`maxMemAllocSize` = 4 GiB on this
  driver+runtime, `ZE_RESULT_ERROR_UNSUPPORTED_SIZE`) unless
  `UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1` is set; with it, a single 12/13/14 GiB
  allocation succeeds **at 474 GB/s**, i.e. real VRAM at full bandwidth. Your
  source comment already names that variable — it is worth more prominence,
  because without it the A770 cannot hold more than a ~3.7 GiB expert cache.
- Only ~7,400 of the model's 24,576 experts fit in the two cards' VRAM
  (~50 GB of experts, 32 GB of VRAM), so the rest must be fed from host RAM.

The two patches attached make the free-VRAM figure usable on both drivers; with
them the split starts and holds a 17 GiB expert cache across the two cards.

## Things that do not work on this hardware (reproducible)

1. **`--pipeline-windows 2` hangs the SYCL port.** The engine accepts the flag and
   prints `two verifiers per stage, the stages overlap across windows (decode too)`;
   the window graphs capture ("captured the N-token window ... no error"), then the
   first request gives `strata serve: verify: error`, the engine wedges
   (`reading the prompt: 63 tokens, 300 s so far`) and the server ends it after
   301 s of silence. Consistent with the port's single-context command-graph
   capture, as `docs/A770_SPEED_SPEC.md` describes.
2. **`xe` + `--stream-experts`, two cards, segfaults.** Both A770s forced onto
   `xe` (`xe.force_probe=56a0`): the model loads, the caches and the **chunked
   pinned RAM mirror build fine** (`6818 experts missing from VRAM mirrored in
   pinned host memory (12.30 GiB)`), then the first request wedges the verify
   window and the engine dies:
   `segfault ... in libze_intel_gpu.so.1.6.33578`, then a general protection
   fault in `libsycl.so.9.0.0`, exit code -6.
3. **`xe` reports no free VRAM at all.** `get_memory_info:
   ext_intel_free_memory is not supported.` -> `expert cache auto: 0.00 GiB free
   -> 0 slots` -> `no VRAM is left for the expert cache` and the engine refuses
   to start. This is the *opposite* misreport from i915 (which reports the whole
   card as free), and it is the more surprising one, since the port's notes give
   the A-series the ordinary streamed path **on xe**. Fix attached.

## Files

- `strata-iq3s-dual.json` — the exact configuration behind the IQ3_S row.
- `strata-iq3s-xe.json` — the `xe` + `--stream-experts` configuration that loads
  and then crashes (item 2 above).
- `engine-log-iq3s-dual.txt` — the engine log of a dual-card IQ3_S run
  (split lines, cache sizing, per-request timings, cache hit rates).
- `FIXES.md` — the two patches, as unified diffs against `v0.1.42`, and the
  native-build environment.
- `hardware.txt` — `lspci`, `sycl-ls`, the measured PCIe bandwidths and the
  driver/firmware versions.

<details><summary>patch: sycl/src/core/expert_cache.cpp</summary>

```diff
--- sycl/src/core/expert_cache.cpp.bak-20261012	2026-10-11 02:13:10.375218601 +0800
+++ sycl/src/core/expert_cache.cpp	2026-10-14 16:32:49.166887968 +0800
@@ -19,6 +19,7 @@
 #include <algorithm>
 #include <atomic>
 #include <set>
+#include <string>
 #include <cstdio>
 #include <filesystem>
 #include <cstdlib>
@@ -31,7 +32,11 @@
 // Level Zero's free-memory query reads the constant total on an Arc A750 (i915): the engine then sized its caches
 // from the whole card and over-allocated (UR_RESULT_ERROR_OUT_OF_HOST_MEMORY at the first launch). Other processes'
 // memory is not in this number, so it only ever LOWERS the figure the driver gave.  0 = not available.
-static unsigned long long own_drm_local_bytes() {
+// #1440: `pdev` is the CALLING device's PCI address ("0000:0c:00.0"). Every card the process uses keeps its own DRM
+// clients in /proc/self/fdinfo, so summing them all charges the first card's cache to the second and the split's
+// second card looks full: only the clients whose drm-pdev matches count. An empty pdev (no address from the runtime)
+// keeps the old sum, so a one-GPU run is unchanged.
+static unsigned long long own_drm_local_bytes(const std::string& pdev) {
 #if defined(__linux__)
     unsigned long long total = 0;
     std::error_code ec;
@@ -40,13 +45,15 @@
         std::FILE* f = std::fopen(e.path().string().c_str(), "r");
         if (f == nullptr) continue;
         char line[256];
-        bool drm = false;
+        bool drm = false, this_card = pdev.empty();
         unsigned long long client = 0, bytes = 0;
         while (std::fgets(line, sizeof line, f)) {
             unsigned long long v = 0;
             char unit[8] = {0};
+            char card[64] = {0};
             if (std::strncmp(line, "drm-driver:", 11) == 0) drm = true;
             else if (std::sscanf(line, "drm-client-id: %llu", &client) == 1) {}
+            else if (std::sscanf(line, "drm-pdev: %63s", card) == 1) this_card = pdev.empty() || pdev == card;
             else if ((std::sscanf(line, "drm-total-local0: %llu %7s", &v, unit) >= 1 ||
                       std::sscanf(line, "drm-total-vram0: %llu %7s", &v, unit) >= 1) && v > 0) {
                 const unsigned long long mul = unit[0] == 'K' ? 1024ull : unit[0] == 'M' ? 1048576ull
@@ -55,14 +62,25 @@
             }
         }
         std::fclose(f);
-        if (drm && bytes > 0 && seen.insert(client).second) total += bytes;
+        if (drm && this_card && bytes > 0 && seen.insert(client).second) total += bytes;
     }
     return total;
 #else
+    (void) pdev;
     return 0;
 #endif
 }
 
+// The current device's PCI address ("0000:0c:00.0"), the key drm-pdev carries in /proc/self/fdinfo (#1440). Empty
+// when the runtime has no answer, and own_drm_local_bytes() then sums every card as it did before.
+static std::string current_pci_address() {
+    try {
+        return dpct::get_current_device().get_info<sycl::ext::intel::info::device::pci_address>();
+    } catch (...) {
+        return std::string();
+    }
+}
+
 // See the header.  NVIDIA's DGX Spark Porting Guide (section 5.5, "Memory reporting on UMA systems") recommends the same:
 // not to rely on cudaMemGetInfo alone but to count the memory the OS can reclaim.  Swap is not counted here (unlike
 // NVIDIA's reference snippet): an expert cache that pushes the system into swap would be far slower than a smaller one.
@@ -75,15 +93,25 @@
     */
     dpct::get_current_device().get_memory_info(free_b, total_b);
     // #1440: own_drm_local