Repository navigation
Conversation
The Strata engine is CUDA. On a GPU without it - written for and measured on an Intel Arc Pro
B70 - the same GGUF weights run through llama.cpp's SYCL backend. serve/engine_llama.py puts
llama-server behind Engine.generate(ids, max_new, sampling, cancel), so the OpenAI and
Anthropic endpoints, streaming, tool calls, MCP and the web app are unchanged above it.
It sends the prompt as token ids (the server's own tokenizer, from the same GGUF) with
return_tokens, streams ids back with llama-server's timings mapped onto Strata's `last`,
heartbeats while a long prompt is read, and stops the server by closing the connection when a
client goes away. Attach to a running server ("llama": {"url"}) or spawn one ("llama": {"exe"}).
Not yet: images (Strata's Vision produces embeddings for the CUDA engine; llama-server takes
the image itself) - start without "vision" and image requests are refused cleanly.
Tests: a scripted llama-server (the wire format observed on b29c606), plus a live round trip
when STRATA_LLAMA_URL is set. Measured on the B70 with Coder IQ1_M: 23-25 tok/s decode,
~150 tok/s prefill, correct code, GPU-bound.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…erver itself Without nvidia-smi, setup now looks for an Intel discrete GPU in sysfs (vendor 0x8086 under xe or i915, named by PCI id - lspci's database does not know an Arc Pro B70 yet) and takes a path with no compiler and no CUDA toolkit: download the model, export the tokenizer, find a llama-server (--llama-server PATH, one on PATH, or the B70 container image via docker), write a config with a "llama" block and a start-llama.sh, and a run script that uses --engine llama. The CUDA engine's rules that do not apply are stepped around rather than deleted: the compute capability and driver-version checks, the experts-in-RAM gates (on llama.cpp the fit is VRAM, shard 1 = the download minus the 28.8 GB lookup table), the MTP draft download (5 GB the llama engine cannot use), and the engine's own argument list. engine_from_config runs the config's start script when nothing answers at the url, so the run script is one click and closing the server stops the container. Verified on an Arc Pro B70: setup runs steps 1-7 clean against the downloaded Coder IQ1_M. docs/INTEL.md has the whole picture and the measurements. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…webui here) Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
About 20 KiB of KV per token (only 12 of 48 layers are full attention). 32k, 64k, 96k and 128k all load with the whole model on the card, 128k with 1.6 GB to spare; 160k is under the safety margin and 256k does not fit - asking for it spilled VRAM into host RAM and took the machine down until its watchdog reset it, so the doc says so in bold. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…4 tok/s Prompt reading scales up with prompt size (150 tok/s at 2.7k, 424 tok/s at 105k), so the full window costs about five minutes, not fifteen. 30.3 GB in use with 1.5 GB free at 131,072. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
… kernel parity tests pass on the Arc Pro B70 The port lives in sycl/ and changes no upstream file: sycl/src and sycl/include hold only the files the migration touched, sycl/CMakeLists.txt takes everything else from the original tree. How it was made, reproducibly (sycl/tools/): a dev image (the llama.cpp SYCL image + SYCLomatic 2025.3 + CUDA 12.8 headers from the pip wheels), migrate.sh (compilation database + dpct over 86 translation units), fixups.py (every hand fix with its reason, idempotent), build.sh (icpx, ninja). What dpct got wrong and the script repairs: CUDA's null stream became a null sycl::queue* (strata::q_of maps it to the default queue); __ldg((const float*) p) lost its cast and read one byte; __fadd_rn lost parentheses; ggml's lookup tables were threaded through kernel parameters with the wrong table per template; helper headers renamed since 2025.3; cudaGraphUpload, graph introspection and %globaltimer have no equivalents. Two compiler flags are load-bearing: -fp-model=precise, and the device compiler's -cl-fp32-correctly-rounded-divide-sqrt - the Arc's fp32 divide is not correctly rounded by default, and quantize_act_parity went from 303k mismatches to byte-exact with it. Not ported: the three kernels with inline PTX (mma.sync, ldmatrix, cp.async) take their existing pre-sm_80 fallback; the launchers refuse the device. The XMX joint_matrix versions are next. The engine itself starts on the card, loads the pack, the projections, the PLE table and the 23.4 GiB expert arena, and is then stopped by the test rig's memory guard: Strata keeps the experts as a resident host copy, and the rig has 23 GiB of RAM (upstream asks for 32). A streaming fill of the VRAM cache from the GGUF is the next piece; the whole model fits in the 32 GB card. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
… Pro B70 --stream-experts (GgufExpertSource): every expert goes straight from the GGUF into the VRAM cache through a small ring of staging buffers; no resident host arena, so the engine runs on a 23 GiB machine. What it took to get from "loads" to "<think>\nThe user wants a Python function": - the doorbell handshake through host-mapped memory: volatile device loads do not bypass the caches on Intel; system-scope atomics (strata::sys_load/sys_store) do. A device-side spin is bounded (kSpinMax): an unbounded one wedged the GT twice when its process died. - STRATA_VERIFY_NO_HOST=1 + STRATA_VERIFY_DEVICE_PLAN=1: with every expert resident the GPU plans each layer itself and the host waits for the whole window graph; per-layer ring visibility inside a graph is not reliable on this platform (measured with sycl/probe). - -fsycl-default-sub-group-size=32: the kernels are written for 32-lane warps; dpct pinned only 134 of 289. - cudaMemcpy is synchronous; dpct's default-queue memcpy had no wait, and the streaming ring reuses its buffers - the expert cache filled from overwritten memory (non-deterministic residuals, NaN by layer 5). - native_expert_parity hand-ported (dpct cannot parse it without ggml-cpu.h): the GPU native expert kernel matches the float reference on real IQ1_M rows. - verify.cpp debug readbacks (STRATA_VERIFY_DEBUG): the per-layer residual ladder that found all of it. Not yet: speed (0.17 tok/s - the spin budgets), the segfault at exit, the tensor-core kernels. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…, eager stage profiler, docs Measured on the B70: 64 tokens in 52.0 s, of which ~47 s is the first window (the runtime JIT-compiles every kernel on first use: 6 tokens take 47.5 s, 64 take 52.0 s). The output is the model's real answer. Debug tooling kept in the port: STRATA_VERIFY_EAGER=1 runs the window on the queue instead of as a graph and turns the stage stamps into host clocks, so STRATA_VERIFY_PROFILE=1 prints ms per stage without a device timer; STRATA_VERIFY_DEBUG=1 prints the residual after every layer and the round's phase times. build.sh takes BUILD_DIR and AOT (an ahead-of-time bmg-g31 build removes the JIT cost). Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…st token in 0.85 s, no JIT The -Xs device flag was mis-quoted; fixed. Measured on the B70 with the same 19-token prompt, 32 tokens: decode 15.17 tok/s, prefill 28.79 tok/s, a 4-token window 73 ms on the card. The persistent JIT cache (SYCL_CACHE_PERSISTENT=1) is the alternative for the JIT build: 15 MB, second start 15.6 tok/s. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…er on iq3_xxs layers A window routes the same expert for up to 6 tokens; the grouped kernels recomputed the dequantised weights (grid lookups, sign unpacking, the byte-wise compare/xor/subtract emulations) once per token. Multi<TY> splits each of the five formats this model uses (iq2_s, iq3_xxs, iq3_s gate/up; iq4_nl, q2_0 down) into prep (once) and acts + finish (per token); the kernels take entries four at a time. Lanes per row is a compile-time constant (STRATA_EXPERT_LANES, 8): 32 lanes on an 80-call row was mostly reduction. Measured with the new NATIVE_BENCH=1 mode of native_expert_parity (10 experts x 4 entries, one layer): iq3_xxs/iq4_nl 1.467 -> 0.649 ms, iq2_s/q2_0 ~0.68 -> 0.579 ms, iq3_s/iq4_nl 0.750 ms; parity unchanged. strata::dp4a wraps dpct::dp4a (sycl::ext::oneapi::dot_acc is the same emulation with a header that breaks the link across translation units). Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…tok/s decode route<NE> templated on the expert count; native_router_top10_multi_ne dispatches 256/512. Both call sites (the verify window and the per-token moe_route) take it for the Coder. 64 tokens: 16.25 -> 16.95 tok/s, same output. Also: Dockerfile.unitrace (per-kernel Level Zero profiling image). Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…: the mmvq kernels run at 110-130 GB/s warm - exit segfault: dpct's static global_memory/constant_memory objects destructed after the default queue (queue::get_device in the core); heap-allocated, never freed. Exit code 0; unitrace can write its table. - Dockerfile.unitrace (pti-gpu unitrace in the dev image) + sycl/rank_kernels.py: per-kernel device time. - Q6KTraits: the column-independent unpack (shifts, masks, saturating byte subtract) moved into load(); acts()/dot() split; a row-blocked kernel (native_mmvq_rowwarp_kernel, STRATA_MMVQ_RPW) kept but off: with a 300 ms warm-up the shared kernel already does 2560x2560x4 in 48 us (110 GB/s) and the variants tie. - mmvq_bench: the dense kernel timing harness (warm-up + 400 iterations; a 5 ms run measures the clock ramp). Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…a thread-safe 512-slot keyed cache The prompt path collects hundreds of blob pointers per layer and copies from them on other threads later; the 8-slot ring handed back pointers that were overwritten by then, and its index was not thread-safe. Slots are now keyed by (layer, expert), reused oldest-first among 512 (~1.4 GB), under a mutex. Long prompts also hung in the prompt path's streamed-expert handshake (lent cache slots make experts non-resident, the host supplies them behind a flag the GPU wait kernel does not see). --no-prefill-borrow keeps every expert resident and removes the handshake: 2,000 tokens 693 tok/s, 2,185 tokens 560 tok/s, 300 tokens 180 tok/s (llama.cpp on the same card: 150 short, 424 at 105k). Rig: watchdog reports and kills a hung run at the expected time (EXPECT_S/QUIET_S); --spec 2 default (20.75 tok/s decode vs 16.95 at spec 4: the suffix drafter accepts 9%). Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…rrectness check; bandwidth probe bw probe (incompressible data): 4-byte loads stream at 395 GB/s, 16-byte at 596. The Q6_K matvec sits at 130-160 GB/s in every variant tried (lanes per row, rows per warp, unroll, 16-byte loads, software pipelining): latency-bound on a dependent load->unpack->dp4a chain, not ALU or bandwidth. The 16-byte kernel is kept for n_out >= 4096 (+13% there, slower on small shapes); rel 5e-8 against the shared one. Decode 20.80 tok/s (spec 2). The drafter is the lever now: MTP fetch + pack running. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
… vs 138-319 tok/s, decode 20.8 vs 24-26 Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…cepted) - 1.8x llama.cpp on the same card Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…75 tok/s in the same run) Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
# Conflicts: # setup.py
…), AOT divide rounding, hybrid K8V4 KV The merge of upstream main (a790805) left 28 migrated copies behind. They are refreshed by re-migration: dpct over the old and the merged tree, sycl/tools/normalize.sh (dpct's names, unchanged files, message serials, fixups.py) on both, then a 3-way merge per file. Copies dpct never produced (verify.cpp, mtp.cpp) take upstream's K8V4 diff by hand. The PTX gate fixup follows upstream's new __HIPCC__ guard (#elif). The AOT build never received -cl-fp32-correctly-rounded-divide-sqrt (it went to the spir64 backend, which an AOT build does not have): quantize_act_parity had its 303k mismatches back. The option now reaches ocloc through -device "... -options ...". 19 of 22 parity tests pass again; kv_hybrid_parity (new) passes all but its tensor-core step and is built, not registered. B70, merged engine, AOT, --stream-experts --no-prefill-borrow --spec 4 --mtp: 2,184-token prompt 566.5 tok/s (was 475), decode 38.3 tok/s at that context (was 42.7; 77% draft acceptance on this text), MTP prompt pass 38.6 ms (E-9 batched). Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…k/s decode, 566 tok/s prompt) Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…r-store dequant, XMX kernels (opt-in), profiles Measured on the B70, outputs identical before and after each change (docs/INTEL.md "Speed work"): - every window graph and the drafter's graphs captured at load (STRATA_WARM_GRAPHS=0 restores first-use capture) - the commit graph is left running while the drafter's round runs on its own queue: decode 38.3 -> 43.1 tok/s at the 2,184-token context, 46.5 -> 49.7 on a short prompt - the draft layer's batched prompt pass takes prompts under 64 rows (106 -> 16 ms on 19 tokens) - one device module per kernel: the first launch in the prompt path 245 -> 1 ms - the expert dequant writes each thread's run as one vector store: 0.085 -> 0.030 ms per expert; the 2,184-token prompt 566 -> 682 tok/s, 8,000 tokens 720 -> 841 - SWAR sign compare/subtract in the expert dots, a local-memory resident-plan kernel, a split-K fused down kernel (bit-identical to the single-token kernel, gr_parity): parity-clean, no measurable decode change XMX: oneMKL's FP16 GEMMs already run on the XMX units (30-60 TFLOP/s). Two joint_matrix kernels are in the tree, correct and opt-in because both lose to the paths they would replace on this card: xmx_gemm_iq (fused dequant + GEMM from the quantized rows; xmx_gemm_bench) and qsa_prompt_attn_xmx (the mma.sync prompt attention's port; qsa_prompt_attn_parity passes at 1e-6 of scale; STRATA_PROMPT_ATTN_XMX=1). Profiling: unitrace -q needs Intel's metrics libraries (Dockerfile.metrics) and dev.xe.observation_paranoid=0; the per-kernel table is in INTEL.md. A trap recorded there: --prefill-until with a native pack drops the prompt tail (the token loop is skipped), so a "windows for short prompts" routing was invalid and is not in this commit. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…t timers around the prompt path's per-layer grouping Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…(2,184 tokens 711 -> 765 tok/s, 8,000 857 -> 986) The rows of a 2,184-token prompt are 27k random 4 KB O_DIRECT reads (114 MB for ~7 MB of rows), all before the first kernel: 466 ms with the reader's 16 blocking threads. 64 threads (the port's default now, a copy of src/platform/direct_file.cpp; STRATA_IO_THREADS still overrides) give 330 ms and 128 / 256 the same - the drive tops out near 85k IOPS. A 256-token first chunk (STRATA_PREFILL_FIRST, 0: off) lets the GPU start while the rest of the rows are read behind it: 2,184 tokens 3,075 -> 2,860 ms, time to first token ~150 ms less; 8,000 tokens 986 tok/s. The chunk boundary moves rounding in the expert GEMMs, so a long greedy continuation can diverge late (token 45 of 64 on the test prompt); short prompts are bit-identical. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…streaming per chunk, where the 77 s go Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…he SSD is the floor of the streamed experts Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
# Conflicts: # setup.py
…mpt path (stale ring pointers, stream-all hang) - Re-migration merge of upstream a790805..d6708a4 (GDN recurrence pipelining, block scores read once per window, the GPU-wide sampler, the draft head's VRAM reserve, correctness fixes); verify.cpp/mtp.cpp ported by hand; two new fixups (a cudaStream_t cast in the sampler, a const sycl::free in native_head). Parity unchanged (19 of 21). B70: 2,184-token prompt 765 -> 792 tok/s, decode at that context 43.1 -> 45.2 tok/s, outputs identical. - The prompt path's stream plan kept GgufExpertSource::blob() pointers for ~1,900 experts per chunk against a 512-slot ring, so blobs were overwritten before they were copied (the first 80k measurement was partly wrong experts). The stager threads now read (layer, expert) themselves via GgufExpertSource::read_into. - The stream-all walk hangs in its first large chunk on this card (copy engine stuck on a barrier, any ring size, either issuer; the old read path too). The port defaults to the per-layer routed-only walk (upstream cf0d12c's STRATA_PREFILL_RING=8, included); STRATA_PREFILL_STREAM_ALL=1 restores it for debugging. - The PLE reader copy keeps its 64-deep queue. 80,000 tokens now complete: 101 s = 790 tok/s. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…0 tokens at 1,062 tok/s, decode 35-40 after Without borrowing, the VRAM reserve for the KV state and the chunk buffers evicts ~1,900 experts, the prompt streams them (790 tok/s at 80,000 tokens) and decode after it runs at 2 tok/s. Borrowing lends ~940 slots and refills them in ~1 s: 1,062 tok/s and 35-40 tok/s decode. It costs ~1 s per prompt at short contexts (2,184 tokens 610 vs 792 tok/s), so the default is by context: borrow above 32K. --prefill-borrow / --no-prefill-borrow decide explicitly; layer splits keep their own buffers as before. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
This was referenced Oct 4, 2026
|
Dang. Closed the PR without merge during the time that I was running the benchmark. Here's the results anyways. benchy v1 (2026-10-04)
System and build:
|
|
@Niko1221 You should at least bring the |
Author
|
good news everyone massive speedups pushed Update: prompt-path speedups on the Arc Pro B70, and a standard bench (benchy v1) Three prompt-path changes in
The last two change summation order, so they are FP32-level but not bitwise, like the CUDA build's tensor-core paths: greedy outputs follow the same text and part ways at a near-tie after tens of tokens. benchy v1, Arc Pro B70 32 GB, each model on its served config, cold page cache, 256 greedy tokens per run:
Against the same bench yesterday: Coder 256K TTFT 328.8 -> 235.1 s, 128K 140.9 -> 111.2 s; IQ2_XS 128K 191.2 -> 155.5 s. Resources (from the same run):
System and build:
If you post numbers from your card, please run the same bench: |
Author
|
Update: conversation_snapshot_test is fixed in 5780652 and passes on the B70 (2,169 checks, no copy-engine resets). A queue memcpy with no device end (pinned host ↔ pageable host) hangs the Arc's copy engine (Engine reset: engine_class=bcs; repeated runs wedged the card). Host-only copies now go through the CPU. The engine's conversation_kv_save/restore had the same copy for --kv-resident sessions, used by the opt-in conversation cache (#596). Now 22 of 25 kernel tests pass. |
Contributor
samsonnguyen
pushed a commit
to samsonnguyen/Strata_B70
that referenced
this pull request
Oct 5, 2026
…_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
Niko1221
added a commit
that referenced
this pull request
Oct 6, 2026
Niko1221
added a commit
that referenced
this pull request
Oct 6, 2026
…s/INTEL.md) The SYCL port of PR #423 tracked 0.1.38; 13 of its migrated copies were behind 0.1.39's sources. Re-migrated 0.1.38 and 0.1.39 with sycl/tools/migrate.sh, normalized, merged with sycl/tools/merge_upstream.py: - clean merges: expert_cache, expert_source, native_dense, pinned, elementwise, kv_stream, iq_multi_parity, prefill, verify.hpp; q8_1_finite.hpp new (#606) - verify.cpp: the #649 trace kept with the port's queue handling (STRATA_VERIFY_TRACE; the HIP-only coherent/doorbell switches stay HIP-only) - generate.cpp: the port's --prefill-borrow beside --vram-elastic/--vram-segment-mib - iq_kernels / native_mmvq: the port's kernels kept, #606's finite q8_1 scale and sum applied by hand (dpct's conflicts there were kernel-name hashes) - expert_cache: --vram-elastic (#533) is CUDA-only on the port too (no driver VMM on Level Zero through dpct); it reports so instead of failing to build - version 0.1.39 Compiled here with icpx 2026.1.1 (SPIR-V, no Intel GPU on this PC): strata builds. Not run on an Arc.
Niko1221
pushed a commit
that referenced
this pull request
Oct 6, 2026
…none) dpct left the DPCT1106 note but dropped cudaMemGetInfo in generate.cpp's stage_room, so every later stage of --layer-split sized its expert cache from 0 bytes free and the second card crashed with 'expert cache: no room'. Reported with the fix by tmking01 in #423 (2x Arc Pro B70); also a fixups.py entry so a re-migration keeps it.
Niko1221
added a commit
that referenced
this pull request
Oct 6, 2026
The image pins level_zero:0, so a two-card layer split could not see the second Arc (#423, tmking01). Set ONEAPI_DEVICE_SELECTOR=level_zero:* for it.
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
…s/INTEL.md) The SYCL port of PR Niko1221#423 tracked 0.1.38; 13 of its migrated copies were behind 0.1.39's sources. Re-migrated 0.1.38 and 0.1.39 with sycl/tools/migrate.sh, normalized, merged with sycl/tools/merge_upstream.py: - clean merges: expert_cache, expert_source, native_dense, pinned, elementwise, kv_stream, iq_multi_parity, prefill, verify.hpp; q8_1_finite.hpp new (Niko1221#606) - verify.cpp: the Niko1221#649 trace kept with the port's queue handling (STRATA_VERIFY_TRACE; the HIP-only coherent/doorbell switches stay HIP-only) - generate.cpp: the port's --prefill-borrow beside --vram-elastic/--vram-segment-mib - iq_kernels / native_mmvq: the port's kernels kept, Niko1221#606's finite q8_1 scale and sum applied by hand (dpct's conflicts there were kernel-name hashes) - expert_cache: --vram-elastic (Niko1221#533) is CUDA-only on the port too (no driver VMM on Level Zero through dpct); it reports so instead of failing to build - version 0.1.39 Compiled here with icpx 2026.1.1 (SPIR-V, no Intel GPU on this PC): strata builds. Not run on an Arc.
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
…none) dpct left the DPCT1106 note but dropped cudaMemGetInfo in generate.cpp's stage_room, so every later stage of --layer-split sized its expert cache from 0 bytes free and the second card crashed with 'expert cache: no room'. Reported with the fix by tmking01 in Niko1221#423 (2x Arc Pro B70); also a fixups.py entry so a re-migration keeps it.
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
The image pins level_zero:0, so a two-card layer split could not see the second Arc (Niko1221#423, tmking01). Set ONEAPI_DEVICE_SELECTOR=level_zero:* for it.
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
… experimental Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
…s/INTEL.md) The SYCL port of PR Niko1221#423 tracked 0.1.38; 13 of its migrated copies were behind 0.1.39's sources. Re-migrated 0.1.38 and 0.1.39 with sycl/tools/migrate.sh, normalized, merged with sycl/tools/merge_upstream.py: - clean merges: expert_cache, expert_source, native_dense, pinned, elementwise, kv_stream, iq_multi_parity, prefill, verify.hpp; q8_1_finite.hpp new (Niko1221#606) - verify.cpp: the Niko1221#649 trace kept with the port's queue handling (STRATA_VERIFY_TRACE; the HIP-only coherent/doorbell switches stay HIP-only) - generate.cpp: the port's --prefill-borrow beside --vram-elastic/--vram-segment-mib - iq_kernels / native_mmvq: the port's kernels kept, Niko1221#606's finite q8_1 scale and sum applied by hand (dpct's conflicts there were kernel-name hashes) - expert_cache: --vram-elastic (Niko1221#533) is CUDA-only on the port too (no driver VMM on Level Zero through dpct); it reports so instead of failing to build - version 0.1.39 Compiled here with icpx 2026.1.1 (SPIR-V, no Intel GPU on this PC): strata builds. Not run on an Arc. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
…none) dpct left the DPCT1106 note but dropped cudaMemGetInfo in generate.cpp's stage_room, so every later stage of --layer-split sized its expert cache from 0 bytes free and the second card crashed with 'expert cache: no room'. Reported with the fix by tmking01 in Niko1221#423 (2x Arc Pro B70); also a fixups.py entry so a re-migration keeps it. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
The image pins level_zero:0, so a two-card layer split could not see the second Arc (Niko1221#423, tmking01). Set ONEAPI_DEVICE_SELECTOR=level_zero:* for it. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Syakyr
pushed a commit
to Syakyr/Strata
that referenced
this pull request
Oct 7, 2026
…_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
Niko1221
added a commit
that referenced
this pull request
Oct 7, 2026
…HED; null-allocation checks On the Arc A750 a kernel never saw the host's doorbell and result-row stores to ordinary host USM (sycl/probe/doorbell.cpp: only the uncached allocation completes the handshake), so every ring wait hit its bound and the window read stale rows: NaN logits. verify's mapped staging and the pool's y_miss rows now use host_malloc_polled (sycl_queue.hpp; needs libze-dev's header, else sycl::malloc_host; STRATA_HOST_UNCACHED=0 turns it off). IQ3_XXS on the A750 now answers 12*12 = 144. Also: the expert cache and the MTP draft head check Level Zero's null return for a too-large allocation (and an Arc A750 needs UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1 for a single allocation over 4 GiB); gather_rows no longer asks for a cooperative launch it does not use; the expert-cache zeroing in generate.cpp is chunked too. Based on the probe in PR #423 (maxfridbe).
0xrjman
pushed a commit
to 0xrjman/Strata
that referenced
this pull request
Oct 7, 2026
…HED; null-allocation checks On the Arc A750 a kernel never saw the host's doorbell and result-row stores to ordinary host USM (sycl/probe/doorbell.cpp: only the uncached allocation completes the handshake), so every ring wait hit its bound and the window read stale rows: NaN logits. verify's mapped staging and the pool's y_miss rows now use host_malloc_polled (sycl_queue.hpp; needs libze-dev's header, else sycl::malloc_host; STRATA_HOST_UNCACHED=0 turns it off). IQ3_XXS on the A750 now answers 12*12 = 144. Also: the expert cache and the MTP draft head check Level Zero's null return for a too-large allocation (and an Arc A750 needs UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1 for a single allocation over 4 GiB); gather_rows no longer asks for a cooperative launch it does not use; the expert-cache zeroing in generate.cpp is chunked too. Based on the probe in PR Niko1221#423 (maxfridbe).
|
@maxfridbe thanks for the amazing work on B70. Hi @Niko1221, I'm Juan from Intel Technical team. Great work on the B70 SYCL port in this PR. I may be able to help with Arc Pro B70 hardware for your testing. What's the best way to reach you privately to discuss this? Thanks! |
Closed
You should help out @maxfridbe too. Look at his repos. He's already putting in work on SYCL support for various things. |
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
9 participants
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.
Hello,
I've burned through my fable-opus credits to get to the point of this pr (not required) but i thought I'd share. I tried to make a pr that is non-invasive to your codebase. Hopefully this will help intel people out there. Basically your kernels in sycl form with lots of me testing to fit the b70 to make sure it works at all. Ill keep updating as your branch updates with optimizations. I just thought I'd share.