#2416 [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)

Descrição

## 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.

| tree | normal recording | native recording | change |
|---|---|---|---|
| 0.1.43 + this change (n=4) | 67.39 tok/s (67.05-67.57) | 68.78 (68.27-69.03) | **+2.1 %** |
| 0.1.44 + this change (n=3) | 66.54 (66.53-67.42) | 68.25 (68.13-68.41) | **+2.6 %** |

- 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.