Intel arc 0.1.40 by maxfridbe · Pull Request #1111 · Niko1221/Strata

maxfridbe · 2026-10-06T11:25:49Z

Replaces #809, which closed when main's history was rewritten. Rebuilt on the new main: the 18 commits
of #809 cherry-picked, plus one commit that refreshes the port to 0.1.40. Nothing outside sycl/ and
docs/INTEL*.md changes.

Why the refresh was needed

main's SYCL build had stopped compiling: fused_gr, expert_source and prefill/gemm no longer
matched the 0.1.40 headers, because sycl/ still held its 0.1.39 copies.

The port is re-migrated with dpct and 3-way merged against 5047172, the commit it was migrated from.
That leaves 175 conflicts and three hand-ported files (verify.cpp, mtp.cpp,
native_expert_parity.cpp), now resolved. dpct's mistranslations in the new code are listed in
docs/INTEL.md (0.1.40).

Fixed with it

sycl: keep cudaStreamQuery's answer in the per-layer ring waits #866: the per-layer ring waits read DPCT_CHECK_ERROR(q->ext_oneapi_empty()), which is always 0.
So whenever experts were mirrored and STRATA_VERIFY_NO_HOST was off, any layer slower than 2 ms
was reported as "never rang".
[SYCL / Intel Arc] 2x Arc Pro B70 --layer-split auto deadlocks at startup after the host-mirror fill #1054: with a layer split, the host mirror now holds only the first GPU's layers.
- Before, it was built before the later stages' caches existed, in the first GPU's context. Every
  later-stage expert was mirrored (13-18 GiB of RAM), then the later GPU's cache fill copied from
  that pinned memory. Using it across contexts is undefined in SYCL; on two B70s it hung at startup.
- This doesn't reproduce on a B70 + B65, so the fix is untested on the reporter's hardware.

Port-side choices

Shared-expert side stream: STRATA_SH_STREAM is off. A recorded SYCL graph only takes its
own queue's work.
MTP drafter: keeps launch-then-wait instead of host spins on mapped outputs.
Expert kernels: upstream's warp-shaped kernels are opt-in (STRATA_EXPERT_SPLIT=1); the port's
grids stay the default.
VMM: vmm.cpp is upstream's no-VMM branch, so --kv-grow and --vram-elastic stay off.
Not built: gemm_bf16_parity (cuBLAS) and native_expert_bench (raw CUDA).
Version: the engine reports 0.1.40-sycl, which setup.py's MIN_ENGINE requires.

Tested (B70, plus a B70 + B65 split)

ctest: 34 of 34 pass, with no engine resets. mmvq_multi_parity's negative control has no power
on the port's wide kernels, which use one layout at every column count; the test reports that the way
it reports gfx906's wave64 layout.
Output: IQ2_XS greedy output is identical to the 0.1.39 port. A --layer-split 29 run over
B70 + B65 matches one card token for token.
Decode (short chat prompt, 64 tokens): 63.5 -> 68.0 tok/s on one card, 57.9 -> 62.7 tok/s on
the split.
Conversation cache: with streamed K/V, two conversations sharing a 15K-token prefix took turns
(park, mount the shared checkpoint, restore). Every reply was sane, with STRATA_VERIFY_NO_HOST=1
and with per-layer waits. The B60 token-0 report from Intel arc 0.1.39 + perf fixes #809 did not reproduce.

…o1221#559's per-stage dense layers in the port's copies The 0.1.39 re-migration (5047172) predates two upstream changes merged for the release, and the port's copies of the files they touch did not get them, so `sycl/` stopped compiling: - Niko1221#626 (ec35cf0): pin_current_thread returns a ThreadAffinity; sycl/src/core/session.cpp and session.hpp still held a long long core (3-way merged from ec35cf0^ -> ec35cf0). - Niko1221#559 (ba707d8, e50f266): NativeDense::load takes layer_lo / layer_hi, and the layer-range filter; the port's native_dense.cpp kept the old signature (3-way merged from 37fcef2 -> 0.1.39; dpct's function-try-block kept). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…ree kernel benches sycl/benchy.sh runs every set-up model x the v1 prompts (sycl/bench/v1: 20 / 2,185 / 8K / 40K / 128K / 256K tokens) with each model's serve config, from a cold page cache, and writes a report (speed, then VRAM / RAM / SSD / experts offloaded / power) plus each config's args. sycl/probe/: the small programs behind the platform findings in INTEL.md. attn_bench / sel_scores_bench / xmx_int8_bench: the prompt attention, the QSA block scores, an int8 XMX prototype. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…M tiles, per-cell attention scores - The lend mirror: with --stream-experts, the experts in the slots a prompt may lend are kept in pinned RAM, so the prompt path DMAs them per chunk and the refill copies from RAM instead of reading the GGUF twice. Same bytes, identical outputs. STRATA_LEND_MIRROR=0: off. - QSA block scores: oneMKL fp32 GEMM tiles of pooled keys x the batch's indexer queries plus a relu-sum kernel (7-8x the per-pair warp kernel at 40K-256K, relative error 2-4e-7). STRATA_SELECT_GEMM=0: the warp kernel. - Prompt attention: scores one work-item per cell, 128-cell chunks (1.5x). Decode keeps the warp kernel. STRATA_ATTN_PERCELL=0: the warp kernel. The last two are FP32 in another summation order (not bitwise, like the CUDA build's tensor-core paths). Also host timers for the per-layer expert grouping (STRATA_PREFILL_TIMING=1). Measured on the B70 (0.1.38 base): Coder 256K TTFT 328.8 -> 235.1 s, 128K 140.9 -> 111.2 s; IQ2_XS 128K 191.2 -> 155.5 s. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…ard; a --layer-split sees every card setup_intel.py --check judged each model by the first card only (with two the IQ2_XS / IQ3_XXS / IQ3_S fit in VRAM across both). strata-sycl.sh: with no ONEAPI_DEVICE_SELECTOR set, a --layer-split gets level_zero:gpu (the image pins level_zero:0); a set one is passed through, as before. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…_PERFORMANCE.md; sycl/TODO.md INTEL.md explains the port, its build and its switches (the accuracy pass and the 2026-10-04 prompt-path changes); INTEL_PERFORMANCE.md holds the numbers (benchy v1 on the B70, history, other people's cards from Niko1221#423, and how to submit numbers); sycl/TODO.md the speedup list with what each attempt measured. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…state.cpp and its test The port's copies of conversation_state.cpp (synced at 0.1.31) and conversation_snapshot_test.cpp (0.1.32) predate Niko1221#559's parking changes (f67bf51, 537588f), while the port builds against upstream's conversation_snapshot.hpp, which Niko1221#559 changed: conversation_snapshot_test aborted on the B70 (munmap_chunk(): invalid pointer). 3-way merged from e46ccca / 3d4905d to 0.1.39, no conflicts. Not ported yet (opt-in, unchanged): Niko1221#559's batch slots (--batch / "parallel"), whose verify.cpp and verify.hpp parts the port's copies do not have. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…free (dpct wrote libc free) The test's cudaFreeHost(p) on sycl::malloc_host memory came out of dpct as free(p), so the test aborted on the B70 with "munmap_chunk(): invalid pointer". That, not the stale parking code, was the abort: dec0097's Niko1221#559 merge is still needed (the port's copies predated it), but its message names the wrong cause. fixups.py keeps the fix across a re-migration. The engine has no other such free (grep: none). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…g on pinned-host <-> pageable copies A queue memcpy/memset with neither end in device memory (USM host pinned <-> pageable host) runs on the B70's copy engine, which hung on it: dmesg "Engine reset: engine_class=bcs", and after repeated runs the device was declared wedged (needs a reboot). conversation_snapshot_test's fill() into its pinned KV pools hit it every run (the urEventWait "hang" was the driver resetting the engine under the copy), and so would the engine's conversation_kv_save / restore of a --kv-resident session, whose pools are pinned host memory (the opt-in conversation cache, Niko1221#596). transfer() and the test's fill / fill_after now memcpy / memset on the CPU when sycl::get_pointer_type says both ends are host or pageable; copies with a device end stay on the queue. On the B70: conversation_snapshot_test 2169 checks passed, no engine resets before or after. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

… as upstream does; s2 grouped kernels bitwise - iq_parity ran as `iq_parity --selftest`, which read "--selftest" as its fixture directory and failed. Registered as the top-level CMakeLists does: an iq_parity_fixtures setup test (tools/iq_fixture.py, when a python3 with numpy exists) and iq_parity on its output, both a skip on exit 3. On the B70 with generated fixtures: all 10 types ok (dequant rel 0, mmvq rel 4.6-5.7e-3). - ple_parity needs ggml's captured block (bench/micro/ple_in.bin, ple_out.bin), which the repository does not carry: registered only where the captures exist (STRATA_PLE_FIXTURE_DIR), instead of failing everywhere. - s2_expert_grouped_parity: the new transposed grouped kernels write `acc += dw * dx * v` in one expression, which -fp-model=precise contracts to an FMA; the previous kernels compute the product in chunk_dot and add it in the caller, which it does not (nvcc contracts both, so they are bitwise equal on CUDA). The last bit differed (both within 1e-7 of the host reference). No contraction in gu/down_grouped_t_kernel; fixups.py keeps it. Also an S2_DEBUG=1 print of the first differences in the test. ctest on the B70: 23 of 23 pass (iq_parity not registered in the image: no numpy), no engine resets. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

… host, B70 and B65 - build.sh sources setvars.sh with -u off: under -u it ended the script outside the image. - STRATA_SYCL_RUNNER: native (inside the toolbox) or distrobox:<name> (from the host), for sycl/serve/strata-sycl.sh and benchy; docker stays the default. The image's single-card selector (level_zero:0) is kept unless ONEAPI_DEVICE_SELECTOR is set. - benchy records the card(s) each run held its VRAM on ("Ran on"), and says that the power column sums every card's energy counter when there are several. - gpustat: card names from lspci -vmm (-mm's fields did not split), the B65's id. - INTEL.md: the toolbox path, including the Level Zero loader it needs (>= 1.21; Debian 13's 1.20.6 fails the first KV allocation with UR_RESULT_ERROR_UNSUPPORTED_FEATURE). - INTEL_PERFORMANCE.md: benchy v1 on a Ryzen 9 9950X host, B70 and B65 at PCIe Gen5 x8. The B70 against the earlier Gen3 x8 host: IQ2_XS/Swift decode +15-30% and long prompts ~25% faster, Coder prompt +15-40%. The B65 runs at 70-77% of the B70 at ~70% of its power. The earlier host's other tables are kept under their own heading until re-run here. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

- strata-sycl.sh and benchy run the strata-sycl-dev image with podman, SELinux labels off for the data root's mount (an unlabeled data disk is unreadable to a confined container). - benchy runs as root; podman's commands (image check, run, kill) run as the user who ran sudo, whose store holds the rootless image. - INTEL.md: building and running with Podman. Its benchy rows match the toolbox run to within noise (9 rows: decode within 0.1 tok/s, prompt within 0.3%). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…benchy passes a config's env - benchy hands the engine the config's "env" map as serve/server.py does (it dropped it), which is also how its A/B variants set STRATA_PREFILL_STREAM_ALL / STRATA_PREFILL_TIMING. - Re-run on the Ryzen 9 9950X host, B70 at Gen5 x8, the dev image under Podman (2026-10-04): the 32K test config (Coder 2,185 tokens at 980-984 tok/s), long contexts by KV format, stream-all against routed-only, borrowing on/off, the API, start time, VRAM/RAM, and the prompt-path time shares. - Findings: k8v4 is now the fastest Coder KV format at 128K-256K on both counts (1,506 / 65.1 tok/s at 128K against int8's 1,327 / 63.9); stream-all and routed-only read prompts at the same rate since the lend mirror; borrowing keeps decode 2-4% ahead; logprobs cost ~6% (the old "no measurable time" was within the earlier host's noise). - The earlier host's tables move to History. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…EMMs, int8 linears, fused decode kernels) Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

… fp16 wins on large expert groups, int8 1.6-2.4x On the B70 at the prompt path's shapes: dense projections (4,096 rows) oneMKL fp16 128-147 TFLOP/s, oneDNN fp16 1.00-1.14x, oneDNN int8 1.6-2.4x; expert gate/up at 256-512 rows oneMKL 28-32 TFLOP/s against oneDNN fp16's 54-56. TODO: oneDNN fp16 for expert groups >= 128 rows (same precision) before any int8 work. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…1.39 (B60 runs, int8 DPAS benches, spin pricing) Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

….2-1.9x in isolation, not a decode win on 32 GB Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…ault From maxious/Strata_SYCL (branch b70-intel-arc-0139, their experiments 26-30 and 42; the code is theirs): - STRATA_MMVQ_PREUNPACK=1: each Q6_K / Q5_K dense matrix unpacked once to signed bytes, so decode skips the per-token 6-bit unpack (native_mmvq, native_dense); STRATA_MMVQ_LOOP=1: the wide Q6_K kernel's 5+ columns as a loop. Both off by default: the default path is unchanged (B70, Coder k8v4: 916 / 75.2 tok/s at 2,185 tokens against 918 / 75.2, the same drafts accepted; 40K 1,571 / 69.1 against 1,571 / 69.2). - On a 32 GB card the pre-unpacked copies cost ~2.5 GiB of VRAM: 1,313 fewer cache slots, decode 72.0 (-4%), as their own engine A/B found. For cards with VRAM to spare. - q6k_preunpack_parity / q5k_preunpack_parity (pass on the B65), q6k/q5k_preunpack_bench, xmx_mmvq_bench (their int8 DPAS matvec: 1.2-1.9x the shipped kernel at decode widths on a B65; bench only). - dpct::error_string / last_error (DPCT_CHECK_ERROR keeps the exception's message) and the standard sub-group helpers in sycl_math.hpp, which their files use; added beside, nothing replaced. - TODO: their branch's other switches A/B'd on the B70 (none beats this port; their default build drafts once and then decodes without speculation at 40K, and STRATA_PREFILL_INT8 does not fit the 256K config's buffers). Co-Authored-By: maxious <maxious@lambdacomplex.org> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

…3% of one) Every process's fdinfo lists each card it has open; the VRAM and busy sums now keep the clients whose drm-pdev is the card being reported. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

and Niko1221#1054 fixed main's SYCL build had stopped compiling: sycl/ still held its 0.1.39 copies. Re-migrated with dpct and 3-way merged against 5047172, the commit the port was migrated from. 175 conflicts and three hand-ported files are resolved, and dpct's mistranslations in the new code are fixed (docs/INTEL.md, 0.1.40). - Niko1221#866: the per-layer ring waits counted any layer slower than 2 ms as "never rang". - Niko1221#1054: with a layer split, the host mirror holds only the first GPU's layers. The later GPU no longer reads another context's pinned memory, and 13-18 GiB of RAM is no longer spent. - Engine version 0.1.40-sycl. ctest 34/34 on a B70. IQ2_XS greedy output is identical to the 0.1.39 port, and a K=29 split matches one card. Decode 63.5 -> 68.0 tok/s on one card, 57.9 -> 62.7 on the split. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

timnevits · 2026-10-06T14:04:23Z

Thanks, Max. We tested 3f372817 on a single Arc Pro B65 (32 GiB), PCIe Gen4 x16, i5-12600K and 128 GB DDR4. Codex, an OpenAI AI agent, ran these tests and drafted and posted this comment with my approval. The measurements come from my B65 hardware.

Ubuntu 26.04.1, xe/NEO 26.22.38646.7, oneAPI 2026.1.0, native Release/JIT build using the serve pipes. Original full-expert IQ2_XS, 8K context, INT8 KV, prefill 512, 3 GiB reserve, four workers, MTP window 4/min-p 0.5, all missing experts mirrored in pinned RAM and STRATA_VERIFY_NO_HOST=1. Prompt/conversation/adaptation caches off.

Three fresh processes per benchmark, native decode timing:

Workload	Decode tok/s, median(range)
Five tasks, 512 input/640 output	45.27 (44.32–45.28)
Five tasks, 7000 input/640 output	44.70 (43.93–44.72)
Public 20 input/256 output	54.22 (54.19–56.49)
Public 2185 input/256 output	52.29 (52.27–53.02)

The five-task rows summarize each run's task median. Compared with our old public-fixture controls before/after these tests, measured decode was about 8–9.5% higher. Full continuations differ between branches, so this isn't a pure kernel-speed comparison. Repeats within your branch match exactly; its no-MTP, one-draft and current-depth pilot outputs also match.

Prefill 1024 reduced 7K TTFT from 21.95 to 14.86 seconds across three fresh runs each. It changed the long continuation and lowered the story's median decode about 9% (ranges overlap), so we're keeping 512 as the general setting. The other small depth/confidence/worker pilots didn't show an overall win.

34 kernel tests passed, plus existing IQ fixtures and native-expert checks; two IQ fixture tests skipped. Both prefill profiles passed 24 API/lifecycle checks each, including near-8K recall, cancellation and restart. No GPU resets or PCIe errors.

All results include our two retained local safety patches: incomplete-mirror startup refusal and explicit pinned-mirror release before QUIT exit. These change startup/shutdown only; the arithmetic kernels are your branch's. The baseline artifacts/settings and earlier patches are in #955.

maxfridbe · 2026-10-06T15:55:27Z

I was attempting to get glm 5.3 flash to use strata, and hit a massive rewrite that would have to happen to achieve it, I went ahead and created a new repo for it as i dont think that it aligns with the 16gb targets of this one, i did manage to get 20t/s on it however ill leave a link here in case anthing from that repo is valuable to strata https://github.com/maxfridbe/nextsycl

Also i say I and really i mean I have pushed claude opus 5.5 to its limits of testing and tweaking and benching... I'm a developer but not a gpu coder.

djbrettb · 2026-10-06T18:33:51Z

Tested 3f372817 on the reporter's hardware from #1054. The two-B70 layer split now starts and runs. Thanks, Max.

Setup

2× Arc Pro B70 (32 GB), each PCIe Gen5 x8 (x8x8 bifurcation through an MCIO retimer).
EPYC 4464P, 128 GB DDR5 ECC, Proxmox VE 9 / kernel 7.0.14, xe.
Built with -DSTRATA_SYCL_AOT=bmg-g31; it compiled clean with no local patches.
Ran in the dev image (oneAPI 2026.1.1).

Model and launch

Qwen3.8-Flash-Next IQ3_XXS with --layer-split auto, --stream-experts, --max-context 32768, --kv int8, --spec 4 + MTP.
Same config that hung in [SYCL / Intel Arc] 2x Arc Pro B70 --layer-split auto deadlocks at startup after the host-mirror fill #1054.

Startup

Ready in about 25 s. Before the fix it hung at the same point every time.
Split at layer 29: CUDA0 runs layers 0-28, CUDA1 runs 29-47.
100% of the experts are resident across the two cards (26.5 / 24.3 GiB).

Speed

Test	Strata `3f372817`	llama.cpp `-sm layer`, same GGUF, same box
Short code reply (36-token prompt)	96.4 tok/s	24.9 tok/s
2,477-token prompt: time to first token	3.94 s	11.97 s
2,477-token prompt: decode	67.1 tok/s	25.2 tok/s
6-prompt mixed set, no thinking	63.6 tok/s (wall)	23.9 tok/s
6-prompt mixed set, thinking	75.3 tok/s (wall)	25.6 tok/s

Stability

Both cards did work during the bench: each peaked at 100% compute and averaged about 50%, as expected for a layer split.
Before the fix, CUDA1 sat at 0%.
12-minute streaming soak: 170/170 requests OK.
- Decode: median 73.1 tok/s (min 58.3, max 144.6).
- TTFT: median 0.31 s.
Tool calls stream correctly; reasoning answers are correct.
No xe engine resets and no PCIe AER events for the whole run.

As far as I can tell, #1054 can be closed once this lands.

…rc-docker The SYCL port at 0.1.40 does not build as shipped. This is reported upstream: Niko1221#1113 sycl: the v0.1.40.1 port does not build as shipped (3 compile errors, 12 undefined symbols) Niko1221#1248 SYCL build fails on Intel Arc A770 - same two errors, same commit Upstream main has no CI at all (no .github/workflows directory), so the port is never built before release and has drifted from the shared headers every version - Niko1221#784 did the same job for 0.1.39, Niko1221#1111 does it for 0.1.40. We are building on maxfridbe's work rather than hand-patching. PR Niko1221#1111 re-migrates the port with dpct and 3-way merges it against 5047172, the commit it was migrated from: 175 conflicts resolved, three files hand-ported (verify.cpp, mtp.cpp, native_expert_parity.cpp), dpct's mistranslations listed in docs/INTEL.md. Conflict resolution: the only two conflicts were sycl/src/core/expert_source.cpp and sycl/src/kernels/cuda/fused_gr.dp.cpp - the two files this branch had hand-patched to get past the first errors. We took THEIRS in both. Their re-migration is the real fix; our hand-edits were stopgaps at individual signatures and are superseded. Verified after the merge: fused_gr_read_multi -> bool (matches include/strata/kernels/fused_gr.hpp:59) load_experts_gguf -> 6 params with const std::atomic<int>* ready (matches include/strata/core/expert_source.hpp:767) sycl/setup_intel.py auto-merged cleanly: this branch's STRATA_SYCL_NATIVE / STRATA_INTEL_GPUS native run mode is intact alongside their changes. Credit: the port refresh is maxfridbe's, not ours. See Niko1221#1111

The engine COMPILES AND LINKS now - the Niko1221#1111 merge did it. This failure was purely our own verification step being wrong about the engine: + /opt/strata/engine-arc/strata --help terminate called after throwing an instance of 'sycl::_V1::exception' what(): No device of requested type available Aborted (core dumped) The Dockerfile comment claimed --help prints usage "before any device is opened". It does not - the engine builds its SYCL context first, so on a GPU-less builder it aborts. `ldd` was clean, so the runtime package set is complete; only the check was wrong. Both places that ran `--help` and demanded rc 0 now accept the no-device abort as the expected outcome while still failing on real breakage: Dockerfile.arc the build-time gate test_arc_image.py:62 "the engine runs to its usage line" Tested the branch logic against three fake engines: aborts with "No device of requested type" -> PASS (expected, no GPU) prints "usage: ..." -> PASS exits 139 with no message -> FAIL (real breakage still caught) Keying on the message rather than the exit code is deliberate: 134 is also what a genuine abort produces, so the exit code alone cannot tell them apart.

…layers; the fdinfo free-memory fallback only where the driver reports the whole card free - generate.cpp: with a layer split the host mirror walked the later stages' layers too (#1054, fixed in #1111 but lost in the re-migration): 20 GiB of pinned RAM for experts the other card holds, or one 39 GiB pinned allocation that fails. Both loops stop at the first stage's last layer. - expert_cache.cpp: device_free_bytes() subtracted the DRM bytes of every card of the process from the card it asked about, so on a layer split the second card was charged the first one's cache and its expert cache would not open. The fallback is for drivers that report the whole card as free (Arc A750/i915); it now runs only then. Reported by crobe201 with both diffs; MaineTim confirmed the second on 2x B60. Co-authored-by: crobe201 <273473305+crobe201@users.noreply.github.com>

maxfridbe · 2026-10-08T22:36:44Z

alas this took a while and im out of claude credits, will merge in 3days

ClayShoaf · 2026-10-08T23:06:06Z

IQ3_S
Single B70
PCIe Gen 3
64 GB DDR4

"Current" is strata 1.39

Input tokens	Current PP / decode tok/s	PR1111 PP / decode tok/s	nextsycl PP / decode tok/s
20	10.69 / 37.44	32.68 / 36.15	31.00 / 30.95
2185	286.80 / 53.52	433.19 / 55.00	407.70 / 44.78
8000	587.81 / 52.48	681.87 / 53.96	693.90 / 44.41
40000	657.92 / 50.63	743.90 / 47.97	776.10 / 41.07
128000	611.96 / 45.99	736.41 / 48.75	774.40 / 40.32
256000	554.96 / 38.53	721.10 / 43.10	765.30 / 36.85

mpisat · 2026-10-10T13:09:16Z

Have you guys seen https://github.com/MistVVK/XeStrata - someone made a fork for intel series. They claim massive improvements.
I only have intel b50 and first results are:

| Matched workload | Strata v0.1.41 B50 baseline | XeStrata |
|---|---:|---:|
| Cold prefill, 120,000 tokens | 366.5 tok/s | 462.0 tok/s (+26%) |
| Decode, 8,192 generated tokens | 21.5 tok/s | 32.9 tok/s (+53%) |
| Run time | 707.64 s | 508.66 s (28% less) |

timnevits · 2026-10-10T16:33:59Z

@mpisat We tried it too and saw useful gains:

7K input, 640 output	Strata v0.1.41	XeStrata xe0.1.40.2.1
Greedy decode	43.3 tok/s	51.0 tok/s (+18%)
Sampled decode	37.2 tok/s	43.5 tok/s (+17%)
Time to first token	18.6 s	16.0 s

Promising short-context results. We're keeping our current serving setup for now.

Codex, an OpenAI AI agent, ran the tests and drafted and posted this reply at my request.

…stage's layers; the fdinfo free-memory fallback only where the driver reports the whole card free - generate.cpp: with a layer split the host mirror walked the later stages' layers too (Niko1221#1054, fixed in Niko1221#1111 but lost in the re-migration): 20 GiB of pinned RAM for experts the other card holds, or one 39 GiB pinned allocation that fails. Both loops stop at the first stage's last layer. - expert_cache.cpp: device_free_bytes() subtracted the DRM bytes of every card of the process from the card it asked about, so on a layer split the second card was charged the first one's cache and its expert cache would not open. The fallback is for drivers that report the whole card as free (Arc A750/i915); it now runs only then. Reported by crobe201 with both diffs; MaineTim confirmed the second on 2x B60. Co-authored-by: crobe201 <273473305+crobe201@users.noreply.github.com>

Niko1221 · 2026-10-17T07:05:01Z

Thank you. The SYCL port is in main since 0.1.40.2 and has moved a long way since this was written, so this PR no longer applies.

maxfridbe and others added 19 commits October 6, 2026 06:22

TODO: speedups suggested by another SYCL project on the B70 (oneDNN G…

ee282f2

…EMMs, int8 linears, fused decode kernels) Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

TODO: review maxious/Strata_SYCL's experiments on top of intel-arc-0.…

70b3361

…1.39 (B60 runs, int8 DPAS benches, spin pricing) Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

TODO: maxious's DPAS / pre-unpack decode matvec re-run on the B65 - 1…

0be097d

….2-1.9x in isolation, not a decode win on 32 GB Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

maxfridbe force-pushed the intel-arc-0.1.40 branch from 0149b0d to 3f37281 Compare October 6, 2026 11:28

james151br mentioned this pull request Oct 7, 2026

0.1.40 SYCL port on 2x Arc Pro B65: expert-cache fill stalls (1-core spin + xe "Timedout job" + device coredump); 0.1.38-based engine works on the same card #1302

Closed

Syakyr mentioned this pull request Oct 7, 2026

arc: a self-contained Intel Arc (SYCL) image, and the native run mode it needs Syakyr/Strata#1

Open

ditronicos mentioned this pull request Oct 7, 2026

Support Intel Arc Pro B50 (Xe2 / BMG-G21) as a tested SYCL target — 16 GB at 70 W is the card users are actually buying #1410

Closed

8 tasks

ClayShoaf mentioned this pull request Oct 10, 2026

XeStrata Integration? maxfridbe/nextsycl#1

Closed

Navigation Menu

Repository navigation

Intel arc 0.1.40 - #1111

Intel arc 0.1.40#1111
maxfridbe wants to merge 19 commits into
Niko1221:mainfrom
maxfridbe:intel-arc-0.1.40

maxfridbe commented Oct 6, 2026

Uh oh!

timnevits commented Oct 6, 2026 •

edited

Loading

Uh oh!

maxfridbe commented Oct 6, 2026

Uh oh!

djbrettb commented Oct 6, 2026

Uh oh!

maxfridbe commented Oct 8, 2026

Uh oh!

ClayShoaf commented Oct 8, 2026

Uh oh!

mpisat commented Oct 10, 2026 •

edited

Loading

Uh oh!

timnevits commented Oct 10, 2026

Uh oh!

Niko1221 commented Oct 17, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

Conversation

maxfridbe commented Oct 6, 2026

Why the refresh was needed

Fixed with it

Port-side choices

Tested (B70, plus a B70 + B65 split)

Uh oh!

timnevits commented Oct 6, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

maxfridbe commented Oct 6, 2026

Uh oh!

djbrettb commented Oct 6, 2026

Uh oh!

maxfridbe commented Oct 8, 2026

Uh oh!

ClayShoaf commented Oct 8, 2026

Uh oh!

mpisat commented Oct 10, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

timnevits commented Oct 10, 2026

Uh oh!

Niko1221 commented Oct 17, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

timnevits commented Oct 6, 2026 •

edited

Loading

mpisat commented Oct 10, 2026 •

edited

Loading