Repository navigation
Conversation
…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
The Intel port's existing runtime is docker-in-docker: sycl/serve/strata-sycl.sh starts the engine inside the strata-sycl-dev image with the folder above the checkout mounted at /work. That is the right shape for a host install where only the engine needs oneAPI, and the wrong shape for a deployable image - you cannot nest it. This adds the second shape. sycl/setup_intel.py: STRATA_SYCL_NATIVE=1 runs the engine in this process' own environment. It skips the docker checks, uses paths as they are instead of remapping them to /work, and makes the config's exe the native wrapper. STRATA_INTEL_GPUS names the cards without sysfs, for a container that does not show the host's /sys/class/drm and for CI, which has no Arc at all. sycl/serve/strata-native.sh: the exe for that mode. Sets the port's runtime defaults (docs/INTEL.md) and execs STRATA_SYCL_BIN. Dockerfile.arc: two stages. The builder installs the oneAPI compiler, MKL and ocloc and builds only the `strata` target with AOT for bmg-g21 (the Arc Pro B60's target), which needs no GPU - ocloc cross-compiles on the CPU. The runtime installs the oneAPI runtime and the Arc compute runtime and copies the engine in. The compiler packages come from Intel's oneAPI apt repo; ocloc and the compute runtime are not in it and come from the kobuk-team/intel-graphics PPA. Neither oneAPI runtime package writes an ld.so.conf.d entry, so the image writes one from the trees that actually got installed, and fails the build if any shared object the engine needs does not resolve. docker-entrypoint-arc.sh: setup on first start, serve after. It relinks the persisted config on every start rather than only when /opt/strata has none - the same flaw as upstream Niko1221#1244 in docker-entrypoint.sh, where a regular file left by an earlier setup pass survives `docker restart` and wins over the edited copy on the volume. sycl/tools/test_arc_image.py: the checks a machine with no Intel GPU can answer, run in CI and runnable on the card before trusting the image. docs/ARC_DOCKER.md: the host requirements (the kernel driver is the one thing an image cannot ship), the env reference, the AOT choice, and what is verified against the repositories versus what has never been run on a card.
…d failing main The first CI run failed at the tag, before any real build work: invalid tag "ghcr.io/Syakyr/strata-arc:0.1.40-bmg-g21": repository name must be lowercase github.repository_owner keeps its case and GHCR will not take it, so IMAGE is now computed in a step with the owner lowercased rather than declared in the job env. Verified locally: Syakyr -> ghcr.io/syakyr/strata-arc, and the key stays IMAGE (lowercasing the whole echo line would have broken that). The second failure was our own guard doing its job on a push to main and showing up as a red X there. A dispatch against a ref with no Dockerfile.arc is a mistake worth reporting; a routine push there is not. The check is now a separate guard job that fails only on workflow_dispatch, and build runs on needs: guard when the file is present - so a main push goes grey instead of red. A step cannot stop the rest of its job, which is why it had to be a job.
The runtime stage died at Dockerfile.arc:136 with exit 1 and no useful message.
It was not a missing library. The step was:
for d in A B C; do [ -d "$d" && echo "$d" >> conf; done && ldconfig
When the last glob in the list misses, that $d && list returns 1, the for loop
returns 1, and the outer && ldconfig never runs. set -e does not fire (the
failing test is inside a && list, which is exempt), so the only symptom is the
RUN exiting 1 right after the last [ -d ] - which reads like a missing .so and
is really shell exit-status.
The trace showed exactly that: the last line executed was the final [ -d ]
against an unexpanded glob, and ldconfig never appeared.
Now the trees are discovered from what is actually installed (find -type d
-name lib, keeping only those holding .so files) with the if form, which
returns 0 when the condition is false. The verification is a per-library
ok/MISS report and a diagnostic dump of /opt/intel on failure, so the next
failure tells us where things are instead of just dying.
Reproduced and verified against a fake tree with the real && chaining:
old form -> exit 1, never reaches ldconfig
new form -> reaches ldconfig, both lib dirs found, exit 0
The SYCL build stops at the compile with: sycl/src/kernels/cuda/fused_gr.dp.cpp:1607:6: error: functions that differ only in their return type cannot be overloaded The header declares it bool: include/strata/kernels/fused_gr.hpp:59 bool fused_gr_read_multi(...) "Returns true when it also wrote the q8_1 images" and the CUDA original agrees (src/kernels/cuda/fused_gr.cu:1621, bool). The DPCT-translated SYCL copy still had void from before that return type was introduced - the header changed in bc6b6f8 (the Aurora gfx1151 merge onto 0.1.39) and the Intel port was not carried along with it. This port does not implement the q8_1 image write at all (no q8 handling anywhere in the function body), so it reports false on every exit path. The bare `return;` at 1755 became `return false;` and the function's end got `return false;` too. Checked that nothing in the SYCL tree reads the value: no `= fused_gr_read_multi` and no `return fused_gr_read_multi` under sycl/, so the constant false changes no behaviour there. Also scanned the rest of sycl/**/*.dp.cpp for the same class of mismatch against include/. Three name collisions came up (launch, valid, finish) and all three are file-local free functions vs class members in unrelated scopes with different signatures - not conflicts.
sycl/src/core/expert_source.cpp:2938:14: error: call to 'load_experts_gguf' is ambiguous
header include/strata/core/expert_source.hpp:767 6 params, last defaulted
port sycl/src/core/expert_source.cpp:2632 5 params
A five-argument call matched both. Same class as the fused_gr_read_multi fix:
the header moved, the DPCT-translated port did not. Not touched by this branch.
Ported the parameter rather than stubbing it, mirroring the CUDA original
(src/core/expert_source.cpp:2985), including its wait_ready spin gate:
auto wait_ready = [ready](int64_t l) {
if (ready != nullptr)
while (ready->load(std::memory_order_acquire) <= l + 1) std::this_thread::yield();
};
called after a layer is claimed in each of the two worker loops, matching
CUDA's two sites. The only call site in this tree passes no 'ready', so the
gate never fires here today - but the port now honours the contract instead
of silently dropping it. <atomic>/<thread>/<chrono> already included.
…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.
…e's exit code
Same mistake a third time. The fallback test ran the wrapper with no
STRATA_SYCL_BIN and asserted rc == 0 - but that resolves the REAL engine,
which builds a SYCL context before doing anything and aborts with
"No device of requested type available" on any host without a card:
FAILED: 1 - the wrapper falls back to $STRATA_HOME/engine-arc/strata
The assertion was about the wrong thing. What we want to know is that the
wrapper looks in $STRATA_HOME/engine-arc/strata, not that the engine can run
without a GPU.
Now the test builds a stub at $TMP/engine-arc/strata that echoes its own path,
and asserts the wrapper execs exactly that. Stronger than before: it pins the
resolved path instead of just requiring a zero exit.
Verified by running the real wrapper against the real stub:
STRATA_HOME=/tmp/arcfbs, no STRATA_SYCL_BIN
-> stub engine resolved at /tmp/arcfbs/engine-arc/strata rc=0
STRATA_SYCL_BIN=/nonexistent/strata
-> STRATA_SYCL_BIN=/nonexistent/strata is not an executable file rc=1
Note for whoever hits this next: three of our checks have now assumed the engine
is inert before it opens a device. It is not - SYCL context init comes first, so
any device-free gate has to go through a stub or accept the no-device abort.
…f the image The image BUILDS. `Build` is green - the engine compiles, links, resolves every shared object, and the in-build test_arc_image.py passes. This failure was the next step, and it was the test's own guard doing its job: No /dev/dri in this container, and STRATA_INTEL_GPUS is not set. Start it with the GPU: docker run --device /dev/dri ... Or name the cards: -e STRATA_INTEL_GPUS="Arc Pro B60:24,Arc Pro B60:24" The Dockerfile sets that env internally, so the in-build run passed; the workflow's separate `docker run` did not carry it. The smoke step now names the cards - which is the supported way to exercise this without a card, and the same flag a container that cannot see the host's /sys/class/drm needs anyway. Also: .github was not in .dockerignore, so a workflow-only commit invalidated the builder's `COPY . .` layer and cost a full ~15 minute engine recompile before the change could even be exercised. CI files do not belong in the image; excluding them makes workflow iteration cheap instead of punishing.
…them
The image has ENTRYPOINT ["./docker-entrypoint-arc.sh"] and no CMD, and the
entrypoint is purely env-driven - so it never looked at its arguments. Anything
passed to `docker run` was discarded and the container fell through to the
setup-and-serve flow.
What that cost in CI: the smoke step ran
docker run --rm -e STRATA_INTEL_GPUS=... IMAGE .venv/bin/python sycl/tools/test_arc_image.py
and instead of running the test, the entrypoint downloaded two GGUF shards
(~29 GB), prepared the pack, built the MTP draft layer, wrote run-iq2_xs.sh,
started serving and then died with
RuntimeError: the engine exited before it was ready
(see /opt/strata/strata-iq2_xs.log)
Ten minutes and 29 GB of egress doing the wrong thing. The same footgun faces
any user: `docker run strata-arc bash` would have started a model download
rather than giving a shell.
Now a command given to `docker run` is exec'd as-is. Placed BEFORE the GPU
guard deliberately - a shell inside the container is exactly what you want when
the container cannot see a card, and the smoke step runs with no card by design.
Verified both branches with a script file (not `sh -c`, whose $0 swallows the
first argument and made my first check meaningless):
with args -> exec target: [.venv/bin/python] [sycl/tools/test_arc_image.py]
no args -> falls through to the env-driven flow, GPU guard next
The second smoke step ("the entrypoint refuses to start without a GPU") passes
no args, so it still exercises the guard as intended.
… 26.04 is fine Two corrections to how the host requirement was framed. First, it read like a caveat specific to this image. It is not. A container ships userspace and nothing else, so the kernel driver always comes from the host - for arcint, for kyuz0's toolboxes, for every GPU image there is. If another Xe image already runs on the machine, the requirement is already met and there is nothing to do. Strata opens the card over the same Level Zero path those do. Second, the page only covered Ubuntu 24.04, where Battlemage needs the HWE kernel. On Ubuntu 26.04 LTS that is a non-issue: 26.04 ships Linux 7.0, which carries a mature xe KMD, so there is no HWE step to take and it is the better of the two. (i915 still does not drive BMG, so xe is the driver either way.) Also notes the cross-release case explicitly: this image is 24.04 userspace on a 26.04 host, which is fine - the only interface is the kernel ABI on /dev/dri/renderD*, and the compute runtime in the image (26.31) is newer than the host distribution rather than older, so there is nothing to catch up to.
…ner fetch
GGUF_DIR maps to setup's --gguf-dir. Without it the entrypoint always lets setup
download the model, so the "put the GGUFs in /home/.../Models yourself" workflow
cannot work - the container would re-fetch ~68 GB on first start.
The llama-swap config for the B60 box lives in scratch/b60-setup/ (gitignored by
the operator's global .gitignore, so it stays local and out of the repo). It adds
a strata-base macro plus the qwen3.8-flash-next-iq2xs-strata entry alongside the
five existing models, and a commented single-card Coder IQ1_M variant.
Two deliberate differences from the existing podman-base, both explained inline:
- No ZE_AFFINITY_MASK. Strata picks cards itself via GPUS/LAYER_SPLIT; a mask
pinned to 0 would hide the second card from it. The single-card variant sets
it back.
- STRATA_INTEL_GPUS names the cards outright, because podman does not reliably
expose the host's /sys/class/drm, which is how setup normally counts and
sizes them.
compose.yml needs no change: llama-swap launches model containers through the
podman socket, so their volumes come from the `podman run` line, not compose.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What this is
A deployable Intel Arc container image, and the one thing the Intel port needed to make that possible: a way to run
the SYCL engine without a Docker daemon inside the container.
The port's existing runtime is docker-in-docker.
sycl/serve/strata-sycl.shstarts the engine withdocker runinside the
strata-sycl-devimage, with the folder above the checkout mounted at/work. That is the right shapefor a host install where only the engine needs oneAPI. It is the wrong shape for an image you ship: you cannot nest
it. This branch adds the other shape — engine, oneAPI runtime and Arc compute runtime in one image, the GPU coming
in through
--device /dev/dri.Target hardware: Intel Arc Pro B60 (24 GB), single or dual, per Niko1221#870 and
bench/results/2026-10-04-community-2x-arc-pro-b60.This branch includes someone else's work - please read before reviewing
The SYCL port at 0.1.40 does not build as shipped. This PR did not introduce that; it is reported upstream:
shipped (3 compile errors, 12 undefined symbols)
Upstream
mainhas no CI at all (no.github/workflowsdirectory), so the port is never compiled before releaseand has drifted from the shared headers every cycle - Niko1221#784 did this same job for 0.1.39.
So this branch merges PR #1111 "Intel arc 0.1.40" by
@maxfridbe (merge commit
e4db92e, 19 commits). That PR re-migrates the port with dpct and 3-way merges itagainst
5047172, the commit it was migrated from: 175 conflicts resolved, three files hand-ported(
verify.cpp,mtp.cpp,native_expert_parity.cpp), and dpct's mistranslations documented indocs/INTEL.md.The port refresh is maxfridbe's, not ours. We are the ones who needed it, not the ones who made it. If you are
reviewing this PR: the port-correctness review belongs on Niko1221#1111. What is ours is the Docker packaging, the native
run mode, the workflow and the docs.
Conflict resolution in that merge: the only two conflicts were
sycl/src/core/expert_source.cppandsycl/src/kernels/cuda/fused_gr.dp.cpp- the two files this branch had hand-patched to get past the first errorsbefore Niko1221#1111 was found. We took theirs in both, because their re-migration supersedes our stopgap signature
edits. Verified after the merge:
fused_gr_read_multiisbool(matchesfused_gr.hpp:59) andload_experts_gguftakes the header's 6thconst std::atomic<int>* readyparameter.sycl/setup_intel.pyauto-merged with this branch's native-mode changes intact.
Changes
sycl/setup_intel.py— native mode.STRATA_SYCL_NATIVE=1runs the engine in the current process' ownenvironment: it skips the
docker/image checks, uses paths as they are instead of remapping them to/work, andmakes the config's
exethe native wrapper.STRATA_INTEL_GPUS="Arc Pro B60:24,Arc Pro B60:24"names the cardswithout sysfs — for a container that does not show the host's
/sys/class/drm, and for CI, which has no Arc atall. The default (docker) path is unchanged; that is regression-tested.
sycl/serve/strata-native.sh— theexefor that mode. Sets the port's runtime defaults fromdocs/INTEL.md(SYCL_CACHE_PERSISTENT=0,ZES_ENABLE_SYSMAN=1,ONEAPI_DEVICE_SELECTOR=level_zero:*) withoutclobbering anything the config set, and execs
STRATA_SYCL_BIN.Dockerfile.arc— two stages.ocloc, builds only thestratatarget (STRATA_SYCL_PARITY=OFF— thekernel tests run on a card and the card is not there), AOT for
bmg-g21. No GPU needed:ocloccross-compiles on the CPU.
The compiler packages come from Intel's oneAPI apt repo.
oclocand the compute runtime are not in that repoand come from the
kobuk-team/intel-graphicsPPA that Intel's client-GPU guide points at. Neither oneAPI runtime package writes an
ld.so.conf.dentry, so theimage writes one from the trees that actually got installed, and then fails the build if any shared object the
engine needs does not resolve.
docker-entrypoint-arc.sh— setup on first start, serve after. It relinks the persisted config on everystart rather than only when
/opt/stratahas none. See the note on Niko1221#1244 below.sycl/tools/test_arc_image.py— the checks a machine with no Intel GPU can answer. Runs in CI, and can be runon the card before trusting the image:
docs/ARC_DOCKER.md— host requirements, the env reference, the AOT choice, troubleshooting. Plus pointersfrom
docs/INTEL_ARC.md.Relation to Niko1221#870
Niko1221#870 is the first B60 report, and its four fixes are already in 0.1.40 — checked against this checkout:
e211→("Arc Pro B60", 24.0)inINTEL_ARCsycl/setup_intel.py:58)ThreadAffinityforSessionLoopScratch::pinned_core(Niko1221#626 drift)NativeDense::loadlayer range (Niko1221#559 drift)write_run_scripttakes 4 argssycl/setup_intel.py:278)So this branch adds nothing to those. What it does take from Niko1221#870 is the device-selector gotcha: that report's
box had a B60 and a B580, and
setup's config"gpu": 1never reaches the SYCL path, so the engine had to bepinned by hand with
ONEAPI_DEVICE_SELECTOR=level_zero:1. The image pinsONEAPI_DEVICE_SELECTOR=level_zero:*instead, so a two-card host sees both cards without the user knowing the bus order.
Two other things worth carrying over from that report for B60 owners: 31 tok/s at setup's defaults, 33–36 tok/s
with the expert-cache/reserve tuning, and roughly 2 tok/s per GiB of expert set removed from the PCIe mirror —
sharper than the B70 numbers in
docs/INTEL.mdsuggest, and the reason a 24 GB card is worth tuning for.And Niko1221#1244, which this branch avoids rather than fixes
Niko1221#1244 is a live bug in the existing
docker-entrypoint.sh: the persisted config is linked in only when/opt/strata/strata-$tag.jsonis absent, so a regular file left there by an earlier setup pass survivesdocker restart, the test is true, and edits to/data/config/strata-*.jsonare silently ignored.I wrote that same line in the ARC entrypoint at first. It is fixed here — the ARC entrypoint relinks every start, and
adopts a pre-
/data/configconfig if the volume copy is empty. The NVIDIA entrypoint still has the bug upstream;that fix is split out separately so it can go in on its own.
What is verified, and what is not
Verified against the live repositories (2026-10-07): the oneAPI compiler/runtime/MKL versions; that
oclocand thecompute runtime are absent from the oneAPI repo and present in the PPA; that the PPA's
intel-oclocinstallsocloc-26.31.1and links/usr/bin/oclocthroughupdate-alternativesin postinst (soicpx's PATH lookupworks); that the SYCL runtime lands
libsycl.so.9+ the UR Level Zero adapters in/opt/intel/oneapi/compiler/2026.1/lib, the MKL runtime landslibmkl_sycl_blas.so.6in/opt/intel/oneapi/redist/lib, and neither writes anld.so.conf.dentry.Verified by running: the native-mode config transform (drops
--resident-experts/--mmap-experts/--kv-resident,emits the native wrapper and
STRATA_SYCL_BIN, 128K context →--prefill 4096and a 2048 MiB reserve); thedefault docker path unchanged (regression test); the wrapper execs the engine and fails loudly on a bad binary;
the card list resolves from
STRATA_INTEL_GPUS; the workflow YAML parses; shell syntax on both scripts.Not verified: this image has never run on an Arc. There is no Docker in the environment
it was written in and no Intel GPU anywhere near it. The first build is the CI run on this branch; the first run on
a card is the next step. Build status: green, and published.
ghcr.io/syakyr/strata-arc:0.1.40-bmg-g21— 3.11 GB, manifestdigest
sha256:d151e31c3cc59522f50c2537386d493a71847c9c47824db24c99b3c110ac1194, package visibilitypublic. Pushed via
workflow_dispatchwithpush=true(run 37596190347, 24m43s - the extra time over the12 min PR build is the image upload). Every check passes: engine compiles and links, every shared object
resolves, the device-free suite reports
all checks passed, and the entrypoint's no-GPU guard fires correctly.Note the
latesttag is deliberately not applied - the workflow only adds it on a tag ref, not a branchdispatch. Pull the version string explicitly.
Seven failures stood between the first attempt and that, and all seven were ours, not the port's:
github.repository_ownerkeeps its case; GHCR rejects itmainredld.so.confstep died beforeldconfig[ -d "$d" ] && ...in aforreturns 1 on a missed glob, killing the&&chain--helpassumed device-freeSTRATA_INTEL_GPUSENTRYPOINTwith noCMD, sodocker run IMAGE <cmd>discarded<cmd>The last one cost ten minutes and a 29 GB model download before it was noticed: the smoke step asked for
test_arc_image.pyand instead got the full setup-and-serve flow. The entrypoint now execs passed args, placedbefore the GPU guard so a shell is available in a card-less container - which is exactly what you need when
debugging one.
The performance numbers quoted above are the community's, measured without Docker — a
container adds a
--devicepassthrough and nothing else on the GPU path, but nobody has put a stopwatch on that.Build / run
The host kernel is the one thing the image cannot ship: Battlemage needs the kernel's
xedriver, which on Ubuntu24.04 means the HWE kernel. The image carries the userspace compute runtime and nothing else about the driver.
Note the image deliberately does not set
STRATA_VERIFY_NO_HOST=1: it is only valid when every expert isVRAM-resident, and on a card that mirrors part of them in RAM that is the path that has hung.