[Feature]: SYCL: opt-in native graph recording (STRATA_NATIVE_GRAPH=1); use_root_sync blocks it (+2-3 % decode on an Arc Pro B70, same tokens) #2416

What do you want to do?

Let the SYCL engine use the native recording mode of sycl_ext_oneapi_graph (oneAPI 2026.1), which gave +2-3 % decode on my Arc Pro B70 with identical output tokens. I have a patch ready and would send it as a pull request if you want it, or you can apply it directly.

Summary

On the SYCL engine every kernel is launched with sycl::ext::oneapi::experimental::use_root_sync (394 sites in 49 files; a leftover of the CUDA→SYCL translation). That property requests a cooperative launch. A kernel that carries it cannot be captured by the new native recording mode of sycl_ext_oneapi_graph (oneAPI 2026.1, property::graph::enable_native_recording): finalize() fails with Failed to instantiate native UR executable graph. No kernel in sycl/ uses a root-group barrier, so the property does nothing for us.

The patch

1. removes use_root_sync from all launch sites (mechanical, 51 files, properties{} or the remaining properties such as grf_size<256> stay), and
2. adds an opt-in switch in dpct/graph.hpp: STRATA_NATIVE_GRAPH=1 creates the graphs with enable_native_recording. If the runtime does not support it the constructor throws; we catch it, print one line and use the normal recording, so switching it on cannot break a start. Default is off, so nothing changes for anyone who does not set it.

Measured on one B70: +2.1 % to +2.6 % decode with the switch on, byte-identical output tokens.

Root cause (minimal repro, 25 lines)

repro/rootsync.cpp records one kernel into a natively recorded graph, once without and once with use_root_sync:

OK    plain nd_range kernel (no properties)
FAIL  same kernel launched with use_root_sync : Failed to instantiate native UR executable graph
OK    use_root_sync kernel outside a graph

(oneAPI 2026.1.1, Level Zero loader 1.34.0, driver compute-runtime 26.35.) The same mini programs also show that native recording accepts everything else Strata's decode graph uses: plain and nd_range kernels with SLM, sub-group-size attributes, atomics, ext_oneapi_submit_barrier, event dependencies inside the in-order queue, memcpy in all directions, pinned and zeMemAllocHost(BIAS_UNCACHED) host memory, spin kernels on host-visible flags, 2500 kernels in one graph, an AOT build for bmg-g31, and the dpct recording path. What it refuses (per the extension text) is cross-queue event dependencies (Failed to end native UR graph capture), host tasks, update() and async malloc; none of them is recorded by the decode path.

Measurements

Machine: Arc Pro B70 32 GB (BMG-G31), Core i7-14700F, Linux 7.0, compute-runtime 26.35.39758, oneAPI 2026.1.1 (20260724), Level Zero loader 1.34.0. Model: Qwen3.8-Flash-Next GSQ-RCO IQ3_S (GGUF + vision mmproj), --spec 4 --spec-min-p 0.5 --kv int8 --kv-resident 20480 --pool-workers 8 --batch 2 --batch-mtp, MTP draft, STRATA_GU_LANES=32 STRATA_DOWN_LANES=8. The benchmark feeds a fixed token continuation (--spec-follow), so no chat template, thinking mode or sampling is involved; the server soak used the OpenAI-style chat endpoint with temperature 0 and no special thinking settings. Method: the same fixed reference continuation (--spec-follow, identical acceptance), variants interleaved in rounds, first round dropped, 640 tokens.

• CPU time per run 128.4 s → 126.8 s.

• The generated tokens are identical in all variants and rounds (same md5 of the output ids), including unpatched 0.1.43 and 0.1.44.

• Prefill is unchanged: a 35-token pass takes 423-430 ms in every variant; an 18,349-token prompt takes 13.1 s (~1400 tok/s) through the server with the switch on.

• Server checks with the switch on (batch 2 + MTP): short chat, tool call, 16k prefill, two concurrent clients (single 42.3 tok/s, two clients 52.1 tok/s together), no Failed to instantiate, aborted, DEVICE_LOST or exception lines in the log.

• Final PR tree (opt-in switch, STRATA_NATIVE_GRAPH), 0.1.44, 3 interleaved rounds: switch on 69.39 tok/s vs off 67.70 (+2.5 %), same output tokens (md5). With the switch on but the 2026.0 runtime the fallback message is printed and decode runs at the normal-recording speed (67.65), no crash.

• Server soak with the switch on (40 min, 74 iterations of short chat, tool call, 16k prompt, two concurrent streams): 0 errors; vision request, 6 client aborts mid-stream and an abort next to a second running stream all fine; log scan for Failed to instantiate, aborted, DEVICE_LOST, timed out, exceptions: 0 hits. 16k prefill median 11.9 s.

Raw runs: data/raw_runs.txt; summaries: data/ab_*.txt.

Where the gain comes from (what is and is not established)

• Recording a graph the normal way puts one event reset per kernel node into the command lists: with the Level Zero tracing layer I counted 13,879 zeCommandListAppendEventReset for 14,144 kernel appends while capturing the decode windows, and 0 of both with native recording (data/zetrace_*.txt).
• That alone does not explain the speed-up. A synthetic graph of 2,400 kernels replays in the same time in both modes (1.526 vs 1.527 ms; with a barrier every 8 nodes 1.526 vs 1.526 ms; with a 4 KiB device copy every 4 nodes 2.009-2.012 vs 2.011-2.015 ms), so the resets cost no GPU time (repro/graphbench.cpp, data/graphbench.txt).
• In the real decode the compute engine is busy 99.6 % of the time in both modes (xe fdinfo, data/gpubusy.txt), so there is no idle gap that native recording closes. The measured gain in that run was +2.7 % (62.54 vs 64.26 tok/s, 4 rounds each). This is a separate run with the engine under fdinfo sampling and a shorter prompt, hence lower absolute speeds than in the table above; the gain is the same size (+2.1 to +2.7 % across all runs).
• So I can only report the measured gain; I have not isolated its mechanism. If a maintainer has an idea which node type in the real graph (spin/wait kernels on host flags, pinned host-to-device copies) behaves differently, I am happy to test it.

What is not tested (please help)

• The *_parity.cpp programs and the identity gates with the property removed. Decode output was identical on the runs above, but that is one prompt.
• Alchemist (A750/A770): the property removal should be neutral, but I have no A-series card; if someone with one can run a before/after pair, that would help.
• Windows, Arc iGPUs, two-card splits, and multi-hour runs under real agent load (the soak above was synthetic and repetitive, 40 min).
• Requirements for the switch: oneAPI 2026.1 or newer and a Level Zero loader of 1.34 or newer and a driver that exposes ZE_extension_record_replay_graph. With the 2026.0.0 runtime the constructor throws ("The property list contains property unsupported for the current object") and the fallback is used. With the 2026.1 runtime on loader 1.28.2 I saw crashes and a hang at the first graph, so please do not turn the switch on there; that is why it stays opt-in.

Reproduce

icpx -fsycl -O1 repro/rootsync.cpp -o rootsync
ZE_AFFINITY_MASK=0 ./rootsync
icpx -fsycl -O2 repro/graphbench.cpp -o graphbench && ZE_AFFINITY_MASK=0 ./graphbench   # replay time, normal vs native, 4 graph shapes
# decode A/B: same binary, STRATA_NATIVE_GRAPH=0 and =1, --spec-follow with a fixed reference continuation

Follow-ups (not in the patch)

• Let setup_intel.py recommend the switch once a loader ≥ 1.34 is detected.
• Fusing small decode kernels (about 2400 graph nodes per verify window) is not worth it on top of this: two fusions (doorbell, QSA-o) kept byte-identical output but gave 0 % decode change on the B70.