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

# Target repository

**[intel/llvm](https://github.com/intel/llvm/issues/new)** — 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](https://github.com/user-attachments/files/31163501/graphcapture-nullcheck.patch)

[sycl-fix-snapshot.zip](https://github.com/user-attachments/files/31163572/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"](https://github.com/intel/llvm/commit/fc6ef33cb0923dbef1a28a785eeed0913c5fef94)
(merged 2026-02-02) replaced a safe, purely-internal state read —

```cpp
*pResult = graphCapture.isActive();
```

— with a raw, unguarded driver call:

```cpp
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

```bash
source /opt/intel/oneapi/setvars.sh --force
llama-bench -m -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(...) () 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(...) () 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
ed3ca: cmpb $0x0,0x230(%rax) ; ZeGraphExt.Supported
ed3d1: je ed46f ; skip if !Supported (not taken here)
ed3d7: mov (%r14),%rdi
ed3da: call f9a60
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.