You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
{{ message }}
Repository navigation
Misc. bug: Missing Lightning Indexer and HC Fused Ops with DS4F on Intel B70 (SYCL) #26549
resolve_fused_ops: Lightning Indexer not supported, set to disabled
resolve_fused_ops: fused DeepSeek V4 HC pre not supported, set to disabled
resolve_fused_ops: fused DeepSeek V4 HC comb not supported, set to disabled
resolve_fused_ops: fused DeepSeek V4 HC post not supported, set to disabled
ggml-sycl supports_op gates lacking cases for these four ops. The scheduler falls back to unfused ggml primitives on CPU — correct, but slower and (for the Lightning Indexer) at significantly higher VRAM usage (the unfused indexer materializes a [n_kv x ubatch] score tensor per layer).
Once the Lightning Indexer is implemented, 1M context is easily achieved on B70.
HC ops (DSV4_HC_PRE/COMB/POST): CUDA port (CUDA ref: dsv4-hc.cu, 294
lines). Mechanical: element-wise and small-matrix math. No tensor cores.
Lightning Indexer: CUDA port (lightning-indexer.cu, 588 lines, uses
NVIDIA WMMA). Our kernel uses plain SYCL (no DPAS/WMMA). Deliberate
16-lane subgroup adaptation (CUDA/Metal hard-assume WARP_SIZE=32 with
explicit TODOs). Flat-index store with an explanatory comment addresses a
known IGC Xe2 miscompile (strided stores via computed base pointers);
AMD/Metal reviewers should not "tidy" it.
Confirmed working on dual Arc Pro B70 (Ubuntu 26.04, Linux 7.0) with PR #26515 + #26568 applied to master.
All four resolve_fused_ops ... not supported warnings are gone, no SIGSEGV at first decode, and output is coherent. Same GGUF segfaults at first decode on unpatched master.
~10 tok/s generation, UD-Q2_K_XL (91 GB), -ngl all --n-cpu-moe 30 -sm layer, 64 GB VRAM + 64 GB system RAM.
Note for anyone else on limited RAM: --load-mode none was worth ~100x over the default mmap here (0.1 → 10 tok/s). With weights larger than system RAM, mmap pages from disk on every token.
Config differs from the original report: Q2_K_XL rather than IQ3_XXS, and two GPUs with layer split.
Confirmed working on Arc Pro B65 (BMG-G31, 32GB) — official ghcr.io/ggml-org/llama.cpp:server-intel image (build b10573-d775b8967, oneAPI 2025.3 userspace), Unraid, kernel 6.18.38, xe KMD, single device via ONEAPI_DEVICE_SELECTOR=level_zero:0.
All four resolve_fused_ops ... not supported warnings gone, no SIGSEGV at first decode, coherent output. DSv4F-0731 UD-IQ3_XXS (~104 GB) hybrid: -ngl 99 --n-cpu-moe 35 -c 65536 -ub 1024 -sm none on 32 GB VRAM + 96 GB DDR4 → 7.2–7.4 t/s sustained over a 14.7k-token generation.
Gotcha for anyone still seeing a silent first-decode crash with these fixes applied: check SYCL_CACHE_PERSISTENT. With SYCL_CACHE_PERSISTENT=1 set, we get a SIGSEGV at first decode that presents identically to the pre-#26515 crash — silent death right after llama threadpool init, dmesg showing a GPF in libc followed by an xe bcs engine reset (the reset is cleanup fallout, not the cause). gdb:
#0 __strcmp_avx2
#1 sycl::_V1::detail::getSortedImages(...)
#2 sycl::_V1::detail::PersistentDeviceCodeCache::getItemFromDisc(...)
#3 sycl::_V1::detail::ProgramManager::getOrCreateURProgram(...)
...
#17 ggml_sycl_op_bin_bcast<op_repeat> ('hc_init' — first op of the DSv4F graph)
#18 ggml_sycl_repeat
Reproduces even with a freshly created empty cache directory, so it's the disk-read path itself, not cache corruption. Unsetting the variable (in-memory JIT only) makes the identical command run. Looks like an oneAPI 2025.3 runtime bug (will file against intel/llvm) rather than llama.cpp, but since this issue is where people land when DSv4F dies at first decode on SYCL, noting it here. None of the usual mitigations touch it (GGML_SYCL_DISABLE_OPT/GRAPH, UR_L0_V2_FORCE_DISABLE_COPY_OFFLOAD=1, legacy L0 adapter, NEO EnableBlitterOperationsSupport=0 — all confirmed applied, crash unchanged).
Two smaller datapoints for hybrid (CPU-expert) configs:
DSpark spec decode (draft-dspark, Q2K-Q8 drafter, n_max clamps to 5) measured net-negative here: 5.3–5.6 t/s @ 0.35 acceptance vs 6.4 t/s drafterless, same config at 32k. Batched verification multiplies DDR4 expert reads; with the fused ops the backend is faster without the drafter.
Activity
mobilinkd commentedon Aug 4, 2026
dsv4-hc-sycl.patch
lightning-indexer-sycl.patch
@arthw -- two patches to address missing Lightning Indexer and HC fused ops.
do not depend on [SYCL] enhance OP set_rows to support all missed data types #26515 for compilation; [SYCL] enhance OP set_rows to support all missed data types #26515 is needed for
end-to-end verification (DSv4F on SYCL crashes at first decode on master
without the set_rows F16 fix).
lines). Mechanical: element-wise and small-matrix math. No tensor cores.
NVIDIA WMMA). Our kernel uses plain SYCL (no DPAS/WMMA). Deliberate
16-lane subgroup adaptation (CUDA/Metal hard-assume WARP_SIZE=32 with
explicit TODOs). Flat-index store with an explanatory comment addresses a
known IGC Xe2 miscompile (strided stores via computed base pointers);
AMD/Metal reviewers should not "tidy" it.
arthw commentedon Aug 4, 2026
@mobilinkd
Got it!
I will check it soon!
Thank you!
MarkC11235 commentedon Aug 10, 2026
Confirmed working on dual Arc Pro B70 (Ubuntu 26.04, Linux 7.0) with PR #26515 + #26568 applied to master.
All four resolve_fused_ops ... not supported warnings are gone, no SIGSEGV at first decode, and output is coherent. Same GGUF segfaults at first decode on unpatched master.
~10 tok/s generation, UD-Q2_K_XL (91 GB), -ngl all --n-cpu-moe 30 -sm layer, 64 GB VRAM + 64 GB system RAM.
Note for anyone else on limited RAM: --load-mode none was worth ~100x over the default mmap here (0.1 → 10 tok/s). With weights larger than system RAM, mmap pages from disk on every token.
Config differs from the original report: Q2_K_XL rather than IQ3_XXS, and two GPUs with layer split.
cfsill commentedon Aug 22, 2026
Confirmed working on Arc Pro B65 (BMG-G31, 32GB) — official
ghcr.io/ggml-org/llama.cpp:server-intelimage (buildb10573-d775b8967, oneAPI 2025.3 userspace), Unraid, kernel 6.18.38, xe KMD, single device viaONEAPI_DEVICE_SELECTOR=level_zero:0.All four
resolve_fused_ops ... not supportedwarnings gone, no SIGSEGV at first decode, coherent output. DSv4F-0731 UD-IQ3_XXS (~104 GB) hybrid:-ngl 99 --n-cpu-moe 35 -c 65536 -ub 1024 -sm noneon 32 GB VRAM + 96 GB DDR4 → 7.2–7.4 t/s sustained over a 14.7k-token generation.Gotcha for anyone still seeing a silent first-decode crash with these fixes applied: check
SYCL_CACHE_PERSISTENT. WithSYCL_CACHE_PERSISTENT=1set, we get a SIGSEGV at first decode that presents identically to the pre-#26515 crash — silent death right afterllama threadpool init, dmesg showing a GPF in libc followed by an xebcsengine reset (the reset is cleanup fallout, not the cause). gdb:Reproduces even with a freshly created empty cache directory, so it's the disk-read path itself, not cache corruption. Unsetting the variable (in-memory JIT only) makes the identical command run. Looks like an oneAPI 2025.3 runtime bug (will file against intel/llvm) rather than llama.cpp, but since this issue is where people land when DSv4F dies at first decode on SYCL, noting it here. None of the usual mitigations touch it (
GGML_SYCL_DISABLE_OPT/GRAPH,UR_L0_V2_FORCE_DISABLE_COPY_OFFLOAD=1, legacy L0 adapter, NEOEnableBlitterOperationsSupport=0— all confirmed applied, crash unchanged).Two smaller datapoints for hybrid (CPU-expert) configs:
draft-dspark, Q2K-Q8 drafter, n_max clamps to 5) measured net-negative here: 5.3–5.6 t/s @ 0.35 acceptance vs 6.4 t/s drafterless, same config at 32k. Batched verification multiplies DDR4 expert reads; with the fused ops the backend is faster without the drafter.UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORYwhenever the L0 context spans both devices, any model — that's [GSD-12641] USM device allocation fails with UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY when Level Zero context spans 2 discrete Arc A770 GPUs (regression from 25.40 to 26.09) intel/compute-runtime#916 (drm/xeXE_VM_BINDbug, kernel fix in 7.1-rc1), unrelated to these ops. Selector-pinning one device per process avoids it.rriggs commentedon Sep 24, 2026
#26568 closes this issue