isGraphCaptureActive/queryGraphCaptureActive calls through an unresolved ZeGraphExt.zeCommandListIsGraphCaptureEnabledExp pointer - crashes on first SYCL queue submit on Arc Pro B70 (Battlemage)

[Kinetic69]

Target repository

intel/llvm — the crashing code lives at
unified-runtime/source/adapters/level_zero/v2/command_list_manager.cpp inside this monorepo
(the oneapi-src/unified-runtime standalone mirror is largely inactive now; intel/llvm is where
unified-runtime development actually happens per its commit history).

graphcapture-nullcheck.patch

sycl-fix-snapshot.zip

Summary

libur_adapter_level_zero_v2.so segfaults on the first kernel submission to any SYCL queue on
an Intel Arc Pro B70 (Battlemage, PCI ID 8086:e223). Root-caused via gdb + objdump against the
actual shipped binary (which ships with full DWARF debug info, not stripped) down to a specific
instruction: ur_command_list_manager::isGraphCaptureActive() calls through
ZeGraphExt.zeCommandListIsGraphCaptureEnabledExp without checking whether the pointer actually
resolved — call *0x278(%r15) with a null value stored at that offset. ZeGraphExt.Supported
(the gate this function trusts) is apparently true even though this one entry point never got
populated for this device/driver combination.

This function runs on every single SYCL queue submission, not only when a caller explicitly
uses SYCL graph capture — queue_impl::isNativeRecording() calls it unconditionally as part of
handler::finalize(). So it's not an edge case: it's the very first thing that happens on the very
first kernel launch.

Root cause

Traced through intel/llvm history: PR #21145 — "Fix urQueueIsGraphCaptureEnabledExp when a fork
is created in a graph"
(merged 2026-02-02) replaced a safe, purely-internal state read —

*pResult = graphCapture.isActive();

— with a raw, unguarded driver call:

ze_result_t ZeResult =
    ZE_CALL_NOCHECK(hContext.get()
                        ->getPlatform()
                        ->ZeGraphExt.zeCommandListIsGraphCaptureEnabledExp,
                    (getZeCommandList()));
*pResult = (ZeResult == ZE_RESULT_QUERY_TRUE);

ZE_CALL_NOCHECK does not null-check the function pointer before invoking it. This is exactly the
code our shipped 2026.1.1 adapter still runs (confirmed by symbol name — the shipped binary still
calls this function isGraphCaptureActive, matching this commit; a later commit renamed it to
queryGraphCaptureActive, so the packaged 2026.1.1 build is pinned to an internal snapshot from
around this point, despite its later build date).

checkGraphExtensionSupport() / ZeGraphExt.Supported is supposed to gate this — platform.cpp
computes it by ANDing (*funcAddr != nullptr) over every entry in the graph extension's
mandatory-function table, zeCommandListIsGraphCaptureEnabledExp included. On this device/driver
combination that computation evidently doesn't hold: Supported reads true at call time, but the
specific pointer is null in memory (confirmed directly via objdump on the installed .so, not
inferred from source).

Environment

• GPU: 2× Intel Arc Pro B70, 32 GB VRAM each (Battlemage, PCI ID 8086:e223, "BMG G31")
• Board/CPU: Gigabyte X299 AORUS Gaming 7, Intel Core i7-7800X (Skylake-X)
• OS/kernel: Ubuntu-based, kernel 7.0.x, xe DRM driver
• oneAPI: DPC++/C++ Compiler 2026.1.1 (2026.1.1.20260724), oneMKL 2026.1.0
• Level Zero loader: 1.28.2
• libze-intel-gpu1 (compute-runtime): 26.27.39122.11
• llama.cpp: reproduced on both current master (876a432, ggml 0.18.0) and tag b8935
  (f454bd7, ggml 0.10.0) — i.e. across ~months of unrelated ggml-sycl changes, with identical
  crash signature both times.
• GPU BARs: full 32 GiB resizable BAR on both cards, confirmed via lspci (Region 2: Memory ... [size=32G]) and zeInit/zeDeviceGet succeeding cleanly outside of any compute call.

Repro

source /opt/intel/oneapi/setvars.sh --force
llama-bench -m <any-gguf> -ngl 99 -dev SYCL0 -p 32 -n 16 -r 1

Crashes identically regardless of model architecture, quantization, or prompt/generation length,
100% of the time, on the very first compute call.

Backtrace (gdb, full symbols — binary ships with DWARF debug info)

Thread 1 "llama-bench" received signal SIGSEGV, Segmentation fault.
0x0000000000000000 in ?? ()

#0  0x0000000000000000 in ?? ()
#1  0x00007fffe8d32414 in ur_command_list_manager::isGraphCaptureActive(bool*) ()
    from libur_adapter_level_zero_v2.so.0
#2  0x00007fffe8d95314 in v2::ur_queue_immediate_in_order_t::queueIsGraphCapteEnabledExp(bool*) ()
    from libur_adapter_level_zero_v2.so.0
#3  0x00007fffe8d6d331 in ur::level_zero::urQueueIsGraphCaptureEnabledExp(ur_queue_handle_t_*, bool*) ()
    from libur_adapter_level_zero_v2.so.0
#4  0x00007fffe39704cb in urQueueIsGraphCaptureEnabledExp () from libur_loader.so.0
#5  0x00007fffefeb4085 in sycl::_V1::detail::queue_impl::isNativeRecording() const () from libsycl.so.9
#6  0x00007fffeff2fea5 in sycl::_V1::handler::finalize() () from libsycl.so.9
#7  ... queue_impl::finalizeHandlerInOrderNoEventsUnlocked<sycl::_V1::handler>(...) () from libsycl.so.9
#8  ... queue_impl::submit_impl(...) () from libsycl.so.9
#9  ... sycl::_V1::queue::submit_with_event_impl(...) () from libsycl.so.9
#10 oneapi::mkl::gpu::launch_kernel_3D(...) () from libmkl_sycl_blas.so.6
#11 oneapi::mkl::gpu::run_check(...) () from libmkl_sycl_blas.so.6
#12 oneapi::mkl::gpu::have_binary_kernels(int*, sycl::_V1::queue*) () from libmkl_sycl_blas.so.6
#13 oneapi::mkl::gpu::mkl_blas_gpu_sgemm_driver_sycl(...) () from libmkl_sycl_blas.so.6
#14 oneapi::mkl::gpu::sgemm_sycl_internal(...) () from libmkl_sycl_blas.so.6
#15 oneapi::mkl::gpu::sgemm_sycl(...) () from libmkl_sycl_blas.so.6
#16 oneapi::mkl::blas::sgemm(...) () from libmkl_sycl_blas.so.6
#17 oneapi::mkl::blas::column_major::gemm(...) () from libmkl_sycl_blas.so.6
#18 ggml_sycl_op_mul_mat_sycl(...) () from libggml-sycl.so.0        <- last llama.cpp frame
#19 ggml_sycl_op_mul_mat<no_quantize_q8_1>(...) () from libggml-sycl.so.0
#20 ggml_sycl_mul_mat(...) () from libggml-sycl.so.0
#21 ggml_backend_sycl_graph_compute_impl(...) () from libggml-sycl.so.0
...

Disassembly of the actual crashing function in the shipped .so (objdump -d,
ur_command_list_manager::isGraphCaptureActive(bool*) at file offset 0xed3a0):

ed3aa: sub    $0x18,%rsp
...
ed3c2: mov    (%rdi),%rdi
ed3c5: call   f9a60 <ur_context_handle_t_::getPlatform() const>
ed3ca: cmpb   $0x0,0x230(%rax)                 ; ZeGraphExt.Supported
ed3d1: je     ed46f                            ; skip if !Supported (not taken here)
ed3d7: mov    (%r14),%rdi
ed3da: call   f9a60 <ur_context_handle_t_::getPlatform() const>
ed3df: mov    %rax,%r15
...
ed409: mov    0x68(%r14),%rdi                  ; getZeCommandList()
ed40d: call   *0x278(%r15)                     ; <-- ZeGraphExt.zeCommandListIsGraphCaptureEnabledExp, NULL

Supported (platform+0x230) reads true; the call target stored at platform+0x278
(ZeGraphExt.zeCommandListIsGraphCaptureEnabledExp) is 0. rip after the call is exactly 0x0,
consistent with a null value at that memory location, not a bad this or stack corruption.

Fix (patch attached, tested working)

Null-check the resolved pointer at each of the five call sites in command_list_manager.cpp that
trust ZeGraphExt.Supported to guarantee a specific pointer is non-null, and fall back to
UR_RESULT_ERROR_UNSUPPORTED_FEATURE (or, for the capture-active query specifically, to
*pResult = false, since that query runs unconditionally on the hot path and a hard error there
would be a behavior change SYCL callers don't expect from what has always been a cheap boolean
check). Full diff: graphcapture-nullcheck.patch (attached).

Built and tested locally: cloned intel/llvm at commit fc6ef33cb092 (the exact commit whose code
the shipped 2026.1.1 binary still runs, confirmed by matching symbol names), applied this patch,
built just the ur_adapter_level_zero_v2 CMake target against the system Level Zero loader
(1.28.2) and oneMKL 2026.1.0, and loaded the patched .so via UR_ADAPTERS_FORCE_LOAD in place of
the shipped one — no other changes.

Result: llama-bench (llama.cpp master, 876a432) runs to completion on both GPUs, repeatedly,
with correct output, where the unpatched adapter crashed 100% of the time on the first call.

Performance (patched v2 adapter vs. v1-adapter workaround vs. Vulkan)

Same model, same machine, same settings — deepseek2 30B.A3B Q4_K_M (GLM-4.7-Flash, MoE), single
Arc Pro B70:

Backend	pp512 (t/s)	tg128 (t/s)
Vulkan	246.8	70.1
SYCL, v1 adapter (UR_ADAPTERS_FORCE_LOAD workaround)	810.0	37.2
SYCL, v2 adapter with this patch	877.2	46.9

The patched v2 adapter is faster than both alternatives on prompt processing, and meaningfully
faster than the v1-adapter workaround on token generation too (+26%), though Vulkan still leads on
tg128 for this model. Confirmed stable across repeated runs and on both installed GPUs (SYCL0 and
SYCL1 independently: pp512 877.2 / 877.5 t/s, tg128 46.9 / 48.8 t/s).

This is a separate finding from the crash itself — noted for context since it was measured while
validating the fix, and because it changes the shape of the general "SYCL slower than Vulkan on
Battlemage" observation (see #22413): at least for MoE prompt processing, the v2 adapter is
actually ahead of Vulkan once it isn't crashing.

Ask

Happy to open this as a PR directly if that's preferred over a patch-attached issue — the change is
small, contained to one file, and already validated against real hardware. Also happy to add a
regression test to test/conformance/exp_graph/ exercising a driver where
zeCommandListIsGraphCaptureEnabledExp resolves to null but Supported is otherwise true, if
that's a useful signal to add given Supported's computation didn't catch this case in practice.

[duaned]

Confirming the same crash on the same GPU model, Arc Pro B70 (8086:e223), in a different environment: oneAPI 2026.1.1 apt, compute-runtime 26.27.39122.12, libze 1.28.6, kernel 7.0.14 (Proxmox), Ubuntu 26.04 in an unprivileged LXC.

Any kernel submit crashes, including a 20 line vector-add compiled with icpx. Backtrace matches: call through null pointer, frame 1 ur_command_list_manager::isGraphCaptureActive, libur_adapter_level_zero_v2.so.0. Crash reproduces with both immediate and batched command lists and with SYCL_CACHE_PERSISTENT=0, so it is independent of #21972.

Workaround that fully restores function here: SYCL_UR_USE_LEVEL_ZERO_V2=0 (also set UR_LOADER_USE_LEVEL_ZERO_V2=0). llama.cpp SYCL then runs at expected Battlemage performance, so this is adapter side, not driver capability.

OpenCL backend on the same driver stack works, which also points at the v2 adapter rather than compute-runtime, though presumably a driver shipping ZeGraphExt would mask it. The null check patch attached above looks like the right fix.

[kswiecicki]

Workaround that fully restores function here: SYCL_UR_USE_LEVEL_ZERO_V2=0 (also set UR_LOADER_USE_LEVEL_ZERO_V2=0). llama.cpp SYCL then runs at expected Battlemage performance, so this is adapter side, not driver capability.

L0v1 adapter doesn't implement the graph record & replay extension, perhaps there's a fallback in llama.cpp for that case.

Since the performance with null-check overhead is still satisfactory I think it's good idea to apply your fix to resolve the immediate issue.
I see there's a PR already, thanks @kweronsx.