Skip to content

Intel Arc Pro B70: carry the 0.1.39 SYCL build and copy-engine fixes - #1

Open
ai-anant wants to merge 9 commits into
mainfrom
intel-arc-b70
Open

Intel Arc Pro B70: carry the 0.1.39 SYCL build and copy-engine fixes#1
ai-anant wants to merge 9 commits into
mainfrom
intel-arc-b70

Conversation

@ai-anant

@ai-anant ai-anant commented Oct 4, 2026

Copy link
Copy Markdown
Owner

Carry of upstream Niko1221#809 onto this fork, unchanged.

The Arc Pro B70 32 GB is the card that port was written for, but Strata 0.1.39's sycl/ tree does not build: the port's copies of session and native_dense were not updated for Niko1221#626 (ThreadAffinity) and Niko1221#559 (NativeDense::load layer range). This branch is exactly the nine commits maxfridbe measured on a B70 (tip b56c45e). No file outside sycl/ and the Intel docs is touched. Nothing was pushed to Niko1221/Strata.

What the author measured (Arc Pro B70 32 GB, Ubuntu 24.04, oneAPI 2026.1, AOT bmg-g31)

  • Build: 0 errors
  • Kernel tests on the card: 22 of 25 pass. ple_parity needs bench captures that are not in the repo. s2_expert_grouped_parity is bitwise after the FMA fix. iq_parity passes with generated fixtures.
  • A host-only SYCL memcpy (pinned host to pageable host) hung the Arc copy engine and could wedge the card. Those copies now go through the CPU.
  • Coder IQ1_M, 256K context: about 73-75 tok/s decode. Prompt reading is faster than the unfixed 0.1.39 port (626 tok/s vs 485 tok/s at 2,185 tokens).

Build on the B70

See docs/INTEL_ARC.md. Short form:

source /opt/intel/oneapi/setvars.sh
# AOT for the Pro B70
cmake -S sycl -B build-sycl-aot -G Ninja -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DSTRATA_SYCL_AOT=bmg-g31
cmake --build build-sycl-aot --target strata
./setup.sh --backend sycl

Reference: Niko1221#809

maxfridbe and others added 9 commits October 4, 2026 12:34
maxfridbe @claude
sycl: the 0.1.39 port builds - Niko1221#626's thread affinity and Nik…
0ff8e57 
…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
maxfridbe @claude
sycl: benchy v1 (the standard Intel bench), the probe programs and th…
1473c0b 
…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
maxfridbe @claude
sycl: prompt-path speedups - the lend mirror, QSA block scores as GEM…
0dcd75b 
…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
maxfridbe @claude
sycl: several Intel cards - setup --check sizes models across every c…
58ef487 
…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
maxfridbe @claude
docs: INTEL.md as the port's reference; every measured speed in INTEL…
7080003 
…_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
maxfridbe @claude
sycl: Niko1221#559's conversation parking in the port's conversation_…
dec0097 
…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
maxfridbe @claude
sycl: conversation_snapshot_test frees its pinned buffers with sycl::…
cdb290a 
…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
maxfridbe @claude
sycl: host-only copies go through the CPU - the Arc's copy engine hun…
5780652 
…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
maxfridbe @claude
sycl: the kernel tests all pass on the B70 - iq/ple_parity registered…
b56c45e 
… 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
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

1 participant

@ai-anant