[Engine] SYCL: +10-35% prompt speed on Arc Pro B70/B65, copy-engine hang fix, int8 crossover cache, failing tests fixed by maxfridbe · Pull Request #3248 · Niko1221/Strata

maxfridbe · 2026-10-19T19:47:37Z

Sorry it took so long i have been working on nextsycl i had forgotten to bring in the enhancements here, hopefully this will get approved, Improvements all around!

Branch: maxfridbe/Strata_B70 intel-arc-0.2.0 (head d119210), based on main f57c0ad; merges cleanly into main e88ee39. Replaces #1111 (closed in the 0.2.0 tracker reset), cut down to what main does not have yet. Only sycl/ and docs/INTEL*.md change.
Problem and result

On Arc Pro B70 and B65 (Battlemage, 32 GB), main’s SYCL engine has one hang, two failing tests and a slow first prompt, and its long prompts are slower than they need to be:

The copy engine hangs (the card ends up wedged). A queue copy with no device end (pinned host <-> pageable host) goes to the Arc’s copy engine, which hangs on it (Engine reset: engine_class=bcs; repeated runs wedge the device until a reboot). conversation_snapshot_test does this on every run, and so would saving or restoring a --kv-resident conversation. After: host-only copies run on the CPU; the test passes with 0 resets.
The int8 crossover is re-measured at every start. strata-sycl.sh runs the engine in a --rm container, so STRATA_PREFILL_MMQ=auto’s cache in ~/.cache/strata is lost. The first 2,185-token prompt (B70, Coder IQ1_M) runs at 761 tok/s instead of 1,450. After: the cache is kept under the data root (STRATA_CACHE_DIR).
Two tests fail on SYCL. s2_expert_grouped_parity (4 failures: the transposed grouped kernels get FMA-contracted under -fp-model=precise, the reference kernels do not) and gr_parity (its QFUSE check expects a q8_1 image the SYCL fused read never writes; the engine is unaffected). After: both pass, and wrong bytes still fail.
Long prompts are slower than they could be. Two prompt-path kernels change, each with an off switch:
    prompt attention scores one work-item per cell, 2.0-2.2x the warp kernel (STRATA_ATTN_PERCELL=0 restores it);
    from 8K cells, QSA block selection scores with oneMKL GEMM tiles, so selection at 40K goes from 2.1 s to 0.7 s (STRATA_SELECT_GEMM=0 restores it).

Prompt speed, interleaved A/B on the same B70 binary: +10% at 2K, +21% at 40K and +35% at 128K. Generation speed is unchanged.

Also included, as small tooling fixes for multi-card and Podman hosts:

STRATA_SYCL_RUNNER=docker|podman|native|distrobox:<name> for strata-sycl.sh.
gpustat.py reads two cards correctly. Before, it showed 163% busy and both cards’ VRAM summed as one.
setup_intel.py --check sizes a model across all cards.
setup_intel.py knows the B65’s PCI id (e222).

Related discussion

#1111 (closed by the 0.2.0 reset; this is its rebase), #423, #809. The hang fix, the cache and the test fixes need no scope agreement. The two prompt kernels change a SYCL-only prompt-path default; both have off switches. Happy to split this into three PRs (hang and test fixes / prompt kernels / tooling) if you prefer.
Approach

Everything stays in sycl/. The new kernels are declared locally in prefill.cpp, not in the shared headers.
The GEMM selection only turns on from 8K cells, because oneMKL’s launch cost loses below that (B70: -2% at 2K, even at 8K, 3x at 40K). It reuses the attention scratch, so no VRAM is added.
Measured and left out, because none of them won on the B70 (identical tokens in each case):
    STRATA_EXP_MODE=8: 82.5 vs 82.5 tok/s decode.
    Queues without profiling: +0.25% decode.
    A 66-wide register top-k for 262K contexts: 2.6x faster in isolation, but it spills registers, so decode was -1.5% and the 128K prompt -6.6%.

Validation

Environment: Arc Pro B70 (bmg-g31) + Arc Pro B65, 32 GB each, xe driver, PCIe Gen5 x8, Ryzen 9 9950X, 61 GiB RAM, Fedora Silverblue 44 (kernel 7.1.5), strata-sycl-dev image under rootless Podman (oneAPI 2026.1), AOT=bmg-g31. Revision: intel-arc-0.2.0 c44fb63 (d119210 only adds the B65 name), against main f57c0ad built the same way.

Passed:

Build: AOT=bmg-g31 JOBS=8 bash sycl/tools/build.sh gives 0 errors on both main and the branch.
ctest on the B70, and on the B65 (ONEAPI_DEVICE_SELECTOR=level_zero:1): 34 of 35 pass on each. That includes conversation_snapshot_test, s2_expert_grouped_parity, gr_parity and the newly registered qsa_prompt_attn_parity, which passes on int8, FP16 and Q4_0 pools. For comparison, main passes 30 of 33, with the snapshot test excluded because it wedges the card. No Engine reset or wedged appeared in the kernel log after any run.
Token parity. Settings: --pcie-frac 0 --adapt-every 0 STRATA_PREFILL_MMQ=1, IQ2_XS, 2,185-token prompt, 256 greedy tokens. Two runs of main are identical, and the branch with both new kernels off is identical to main.
    With the new kernels on, the first-token KL against the old kernels is 3.9e-10 nats at 8K and 1.3e-11 at 40K, with the same top-1. That is the same size as main’s own int8-vs-FP16 switch (8.4e-10 and 5.1e-12).
    The 256 greedy tokens after a 128K prompt are identical.
    Note: with the default STRATA_PREFILL_MMQ=auto, two runs of main differ, because the per-group choice comes from a timing.
setup_intel.py --check on the two-card host. The output is in commit 7b70671a.
Benchmark. Qwen3.8-Flash-Next Coder IQ1_M and IQ2_XS with the served flags (--prefill 4096 --spec 4 --spec-min-p 0.5 --mtp --kv int8 --stream-experts --vram-reserve-mib 2048 --kv-resident 32768, MMQ auto with a warm crossover cache for both builds). Cold page cache, 256 greedy tokens, prefill / decode in tok/s:

card model prompt main f57c0ad intel-arc-0.2.0
B70 Coder IQ1_M 2,185 1,312 / 87.6 1,449 / 88.1
B70 Coder IQ1_M 8,000 1,620 / 84.9 1,826 / 88.5
B70 Coder IQ1_M 40,000 1,580 / 77.6 1,909 / 78.7
B70 Coder IQ1_M 128,000 1,351 / 78.5 1,821 / 76.4
B70 IQ2_XS 2,185 1,053 / 80.1 1,355 / 77.4
B70 IQ2_XS 40,000 1,376 / 74.1 1,603 / 77.3
B65 Coder IQ1_M 2,185 948 / 60.1 1,018 / 62.0
B65 Coder IQ1_M 40,000 1,039 / 54.7 1,315 / 54.8
B65 Coder IQ1_M 128,000 861 / 53.6 1,246 / 54.4
B65 IQ2_XS 40,000 942 / 53.2 1,162 / 50.1

Decode moves with the text: the prompt kernels change the greedy output after a few dozen tokens, and MTP acceptance changes with it. With identical tokens, decode is unchanged.

Interleaved A/B on the B70 (same binary, --pcie-frac 0 --adapt-every 0 STRATA_PREFILL_MMQ=1, 2 rounds, prefill tok/s, old kernels → new):

Coder at 2,185 tokens: 1,329 → 1,468 (+10%).
Coder at 40,000: 1,562 → 1,885 (+21%).
Coder at 128,000: 1,336 → 1,798 (+35%), with identical tokens.
IQ2_XS at 40,000: 1,373 → 1,613 (+18%).

Failed:

ple_parity: the Q2_0 shard-2 fixture is not on this machine. It fails the same way on main.

Not run:

CUDA, HIP and Metal builds: nothing outside sycl/ and docs changed.
Arc A-series, B580 and B60: no such cards here. The new kernels are plain SYCL (sub-group 32) plus oneMKL, so they build for every target, but their speed there is unmeasured. Both have off switches.
A layer split across the B70 + B65 end to end: the engine runs above used one card each.
Sampled (non-greedy) quality, beyond the first-token KL above.

Limits

The two prompt kernels are not bit-identical to the warp kernels (FP32 summation order). Their switches restore the old kernels.
The 8K threshold for the GEMM selection was measured on a B70 only.
STRATA_SYCL_RUNNER=podman was exercised through strata-sycl.sh itself. native and distrobox: were not re-run on this branch.

…gs on pinned-host <-> pageable copies A queue memcpy/memset with neither end in device memory (USM host pinned <-> pageable host) runs on the Arc's copy engine, which hangs on it on a B70: dmesg "Engine reset: engine_class=bcs", and after repeated runs the device is declared wedged (needs a reboot). conversation_snapshot_test's fill() into its pinned KV pools hit it on 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. 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.

…nsposed grouped kernels The transposed grouped kernels (gu/down_grouped_t_kernel) 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; on the B70 the last bit differed and s2_expert_grouped_parity reported 4 failures (moe_grouped_s2 / moe_group_resident, both within 1e-7 of the host reference). `#pragma clang fp contract(off)` in the two kernels, kept by fixups.py (9d) across re-migrations; 9c (the snapshot test's sycl::free) is back in fixups.py for the same reason.

…GR read does not write q8_1 fused_gr_read_multi on SYCL never writes the q8_1 image of `mixed` (fused_gr.dp.cpp returns false: the caller quantizes), so gr_parity's S26 check failed on every SYCL device with "QFUSE: T=1 rep 0: not written". The engine's path is fine (verify.cpp quantizes when the read returns false); the test now says "skipped" when the read writes nothing at T=1, and a read that writes different bytes still fails.

…distrobox (STRATA_SYCL_RUNNER) The serve wrapper hard-coded `docker`. Hosts that keep the toolchain in containers (Fedora Silverblue: rootless Podman, toolbox/distrobox with oneAPI) could not serve the SYCL engine without a Docker daemon. STRATA_SYCL_RUNNER=docker (default, unchanged) | podman (the same image and flags, plus label=disable for a data disk without SELinux labels) | native (already inside a toolbox with oneAPI, data root at /work) | distrobox:<name> (enters that distrobox from the host). The settings the wrapper forwards (STRATA_*, ONEAPI_*, ...) reach the engine the same way in every mode.

With two Intel cards every process that opens both lists both in its fdinfo, so the sampler added the other card's clients: busy read 163% of one card and vram_used_mb was both cards' sum (server_intel's VRAM pill and setup's free-memory reading use this file). Clients are filtered by drm-pdev to the sampled card and keyed by (card, client id), since client ids count per card. Also the Arc Pro B65 (e222) by name, and the lspci fallback reads `-vmm`'s Device line (`-mm`'s quoted fields do not split when an unquoted -p00 sits among them).

With several cards, --check judged each model by the first card only: on two 32 GB cards the IQ2_XS, IQ3_XXS and IQ3_S read "fits with part of its experts mirrored in RAM (slower)" although a --layer-split holds them in VRAM. A model that does not fit one card but fits the cards together (each less the same 1.5 GB of room) now reads "fits in VRAM across the N cards (--layer-split)"; one card and the A-series wording are unchanged.

…rp kernel, prompt +10% at 2K) On a B70 the QSA prompt attention is 15% of a 2K prompt and 28% of a 40K one (STRATA_PREFILL_TIMING). The warp kernel splits each cell's 256 dimensions over a sub-group and reduces them with a 5-step shuffle tree per query head, 60 shuffles a cell. qsa_prompt_attn_percell gives work-item c cell c's key row (16-byte loads) and all 12 heads' dot products, reading the query heads from local memory, in 128-cell chunks; the value pass and the merge are the warp kernel's. Each 16-dimension block is summed in four short chains before it joins the cell's sum, so its FP32 error stays near the warp kernel's. Prompt path only (prefill.cpp); decode and the draft layer keep qsa_decode_attn_batch. STRATA_ATTN_PERCELL=0 brings the warp kernel back. Not bitwise: FP32 in another summation order (as the XMX prompt kernel). qsa_prompt_attn_parity now checks it too (against FP64 and the warp kernel, int8 / FP16 / Q4_0 pools, 0 failures on a B70: error vs FP64 1.3-1.6x the warp kernel's, 1.2-2.2x faster) and is registered with ctest; the XMX kernel's refusal of Q4_0 pools no longer ends the test. Arc Pro B70, Coder IQ1_M, 256K config, `--pcie-frac 0 --adapt-every 0 STRATA_PREFILL_MMQ=1`, 256 greedy tokens: 2,185-token prompt 1,330 -> 1,481 tok/s (+11%, 2 rounds each). First-token logits against the warp kernel at 8K: KL 3.9e-10 nats, top-1 equal (upstream's int8 vs FP16 prompt experts on the same prompt: KL 8.4e-10).

…3x faster at 40K, prompt +21% at 40K, +35% at 128K) The prompt path scores every pooled key block against every query's indexer heads; on SYCL that is the warp kernel (the CUDA build's tensor-core path is not ported), 8% of a 40K prompt on a B70 and growing with the context. As a dense product it is one oneMKL FP32 GEMM per tile (pooled keys x the batch's 4 x 256 indexer queries) and a small kernel that applies the relu per head, sums in head order, and gives block n_bid the dead key and its +1e9 rule (qsa_block_scores_reduce, the warp kernel's semantics). The tile is carved from the attention scratch, free until that layer's attention on the same queue, so no VRAM is added. Below ~8K cells oneMKL's launch cost loses (-2% at 2K, even at 8K), so by default only a batch reaching 2,048 blocks takes it; STRATA_SELECT_GEMM=0 never, =1 always. FP32-level, not bitwise (as the CUDA 3xTF32 path). Arc Pro B70, `--pcie-frac 0 --adapt-every 0 STRATA_PREFILL_MMQ=1`, with the per-cell attention, old kernels -> new: Coder IQ1_M 40K 1,562 -> 1,885 tok/s, 128K 1,336 -> 1,798 (the 256 greedy tokens identical); IQ2_XS 40K 1,373 -> 1,613. First-token KL against the old kernels at 40K 1.3e-11 nats, top-1 equal. docs/INTEL.md has the table.

…he int8 crossover was re-measured at every start STRATA_PREFILL_MMQ=auto (the Battlemage default) times the int8 and FP16 expert chains at the first prompt of each expert format and caches the result in ~/.cache/strata/i8-crossover.txt. strata-sycl.sh runs the engine in a `--rm` container, so that home - and the cache - is gone after every run: each server start paid the measurement again inside its first prompt. Arc Pro B70, Coder IQ1_M (six format pairs), 2,185-token prompt: 761 tok/s on the first run without the cache against 1,450 once it is read back (int8 forced: 1,465-1,469); the same choice is made either way (500 of 500 groups int8). The wrapper now sets STRATA_CACHE_DIR=/work/.strata-cache (the mounted data root) unless the environment names another.

On a B70 + B65 host setup listed the second card as "Intel Intel GPU e222 (xe)": the id was missing from the name table (lspci knows it as Battlemage G31 [Arc Pro B65], 32 GB). setup --check now reads "GPU 1: Intel Arc Pro B65, 32 GB VRAM - can be used".

deriknel · 2026-10-20T06:47:34Z

Tested on 2x Arc Pro B70 (CachyOS, kernel 7.2.11, xe) with Swift 1.5 Flash-Next IQ2_XS: layer split 26, every expert in VRAM, int8 KV, MTP spec 4, prefill 8192, 262144 context. The image is v0.2.0-beta (3d1fb48) plus our SYCL patches, including #2808 (STRATA_PROMPT_ATTN_UNION=1). The baseline arm turned this PR off with STRATA_ATTN_PERCELL=0 STRATA_SELECT_GEMM=0. Each arm ran on a fresh engine, two rounds. Prompt reads (tok/s, after a warm-up bench):

prompt	off (2 rounds)	on (2 rounds)	change
~1.2K tokens	916 / 898	920 / 915	+1%
~4.5K	1,434 / 1,419	1,438 / 1,425	0%
~33.6K	2,275 / 2,267	2,344 / 2,342	+3%
~148.8K	2,275 / 2,275	2,719 / 2,714	+19%

The ~148.8K read went from 66 s to 56 s. The needle was found in every read. Decode didn't change (123.1 against 123.0 tok/s mean over 5 greedy prompts), and the greedy text was identical on all 5 prompts in both rounds.

Kernel suite: on the same image without this PR, gr_parity and s2_expert_grouped_parity fail. With it both pass, leaving ple_parity (needs a Q2_0 model file we don't have) as the only failure out of 39. No xe resets or page faults across the run. We didn't hit the copy-engine hang this fixes, so we can't say anything about that part.

Together with #2802 and #2810, the ~148.8K read stays at 2,724 tok/s. We carry this as a local patch until it merges.

We investigated (myself and Claude Opus 5.5 in Claude Code) and I reviewed the findings and ran every test on my two Arc Pro B70s on my CachyOS machine.

SleepinDevil · 2026-10-20T14:30:40Z

This is impressive, great work! Eager to see this merged into the project!

…gine hang fix, int8 crossover cache) — integ batch 19 [gpustat: integ's multi-card design + PR's B65 name; CMake registers PR's args-form parity test beside integ's union test]

…hole request)

Mouzon1775 · 2026-11-02T19:27:51Z

Thanks @maxfridbe for this PR, and @hotschmoe and @Niko1221 for the work it builds on. I cherry-picked d95a9c4 and 8cc654f (authorship kept) onto hotschmoe’s 1c36052 and published the result, with our launch settings: https://github.com/Mouzon1775/Strata/tree/f3-4xb70

Hardware: 4× Intel Arc Pro B70 32 GB, Qwen3.8-Flash-Next UD-Q4_K_XL, layer split across all four cards.

No GPU faults (xe fault/reset/wedge) in about 4.75 hours of testing.
100K-token prompt, 3 runs each, compared with plain 1c36052: first token 30.7 s → 27.2 s (−11%), whole request 37.3 s → 34.0 s (−8.8%).
Generated text was byte-identical to plain 1c36052 on all 14 prompts we could compare.

One request from the 4-card side: on SYCL, --pipeline-windows currently works only with a two-stage layer split. Multi-stage support, as the CUDA/HIP engine has, would help 4×B70 setups a lot, and I’m happy to test branches on this machine.

): host-only copies on the CPU, s2 grouped kernels without FMA contraction, gr_parity QFUSE skip, crossover cache dir Ported from maxfridbe's pull request Niko1221#3248 (commits d95a9c4, 5fbef34, 344e43e, c44fb63 of that branch), cherry-picked without the prompt-kernel commits. - conversation_snapshot: a copy with neither end in device memory (pinned host <-> pageable host) is a memcpy on the CPU; through the queue it runs on the Arc's copy engine, which the PR reports hanging (Engine reset: engine_class=bcs). - s2_expert_grouped: #pragma clang fp contract(off) in the two transposed grouped kernels (fixups.py 9d keeps it across re-migrations; 9c the snapshot test's sycl::free). - gr_parity: the STRATA_QFUSE read check is skipped when the SYCL fused GR read writes no q8_1. - strata-sycl.sh: STRATA_CACHE_DIR defaults to the mounted data root so the int8 crossover cache survives the --rm container.

…: per-cell prompt attention fallback, QSA block scores as oneMKL GEMM tiles Ported from maxfridbe's pull request Niko1221#3248 (commits e55a11c and 8cc654f of that branch). - STRATA_ATTN_PERCELL (default on): qsa_prompt_attn_percell scores one work-item per cell in the prompt path's fallback attention (reached when the Battlemage ESIMD kernel refuses the pools, e.g. FP16 or Q4_0 KV, or STRATA_PROMPT_ATTN_ESIMD=0). =0 keeps the warp kernel. Not bitwise: FP32 in another summation order. - STRATA_SELECT_GEMM (default: from 2048 blocks): QSA block scores as oneMKL FP32 GEMM tiles plus qsa_block_scores_reduce; =0 the warp kernel, =1 always. Not bitwise. - qsa_prompt_attn_parity also checks the per-cell kernel and is registered with ctest. Conflict resolution against the tree: qsa_select.dp.cpp keeps qsa_block_scores_rows beside the new reduce kernel; docs/INTEL.md keeps this tree's text (the PR's table is not copied; measurements are in the Intel notes).

…RATA_<name>=0 turns each off), STRATA_SYCL_IQ_DECODE_FAST default ON in CMake and passed by the SYCL build scripts; STRATA_SELECT_GEMM_MIN sets the block threshold of the Niko1221#3248 GEMM block scores Controls of glorysdj's pull request Niko1221#4185: unset means on for a Battlemage device and off for every other device (strata::bm_default_on); a set value still decides. Measured results are in phaseA-tests/pr-priorities/RESULTS-INTEL.md.

…RATA_<name>=0 turns each off), STRATA_SYCL_IQ_DECODE_FAST default ON in CMake and passed by the SYCL build scripts; STRATA_SELECT_GEMM_MIN sets the block threshold of the Niko1221#3248 GEMM block scores Controls of glorysdj's pull request Niko1221#4185: unset means on for a Battlemage device and off for every other device (strata::bm_default_on); a set value still decides. Measured results are in phaseA-tests/pr-priorities/RESULTS-INTEL.md. (cherry picked from commit 8fffe81) b70 port: the controls read through strata::bm_default_on next to apply_b70_defaults in sycl_queue.hpp. Where this tree has its own kernel for the same call, its order from the Niko1221#4185 port stays: STRATA_IQ_DECODE_FAST before the row-reuse / SLM expert kernels, attn_merge_v2 before the active merge (attn_merge_mode now takes the queue), KREG before the Q6_K shape GRF. STRATA_GR_REDUCE_ILP keeps this tree's 8-loads-in-flight default (=0 off), as the plain reduce has it. STRATA_RESIDENT_PLAN_COMPACT stays unwired (see the Niko1221#4185 port).

maxfridbe added 10 commits October 19, 2026 12:06

deriknel mentioned this pull request Oct 20, 2026

[Engine] sycl: cheaper IQ decodes in the grouped expert kernels (STRATA_IQ_DECODE_FAST=1, opt-in, bitwise) #2802

Open

deriknel mentioned this pull request Oct 20, 2026

[Engine] sycl: STRATA_HC_Q8=1 reads the Q8_0 hyper-connection projections (it loaded them and never read them) #2810

Open

Merge branch 'main' into intel-arc-0.2.0

fce03f9

Mouzon1775 added a commit to Mouzon1775/Strata that referenced this pull request Nov 2, 2026

README: correct the PR Niko1221#3248 timing figures (first token vs w…

845e6cd

…hole request)

Navigation Menu

Repository navigation

[Engine] SYCL: +10-35% prompt speed on Arc Pro B70/B65, copy-engine hang fix, int8 crossover cache, failing tests fixed - #3248

[Engine] SYCL: +10-35% prompt speed on Arc Pro B70/B65, copy-engine hang fix, int8 crossover cache, failing tests fixed#3248
maxfridbe wants to merge 11 commits into
Niko1221:mainfrom
maxfridbe:intel-arc-0.2.0

maxfridbe commented Oct 19, 2026

Uh oh!

deriknel commented Oct 20, 2026

Uh oh!

SleepinDevil commented Oct 20, 2026

Uh oh!

Mouzon1775 commented Nov 2, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

maxfridbe commented Oct 19, 2026

Uh oh!

deriknel commented Oct 20, 2026

Uh oh!

SleepinDevil commented Oct 20, 2026

Uh oh!

Mouzon1775 commented Nov 2, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants