Repository: ggml-org/llama.cpp
Pull request: SYCL: Support sycl_ext_oneapi_limited_graph

The current usage of the SYCL-Graph extension checks for the `sycl_ext_oneapi_graph` device aspect. However, it is also possible to support `sycl_ext_oneapi_limied_graph` devices that don't support update. This is primarily OpenCL and Level-Zero backends to DPC++, as CUDA and HIP backends always have full graph support.


Tested Using `./bin/test-backend-ops -b SYCL0 -o RWKV_WKV7` and using SYCL_UR_TRACE to verify changes by usage of UR entry-points for implementation SYCL-Graph.


CUDA backend to DPC++  supports the full graph aspect, and so can do update. This is update behavior is consistent both before and after the change, with usage of entry-points for both graph update and graph creation:
```
$ SYCL_UR_TRACE=2 GGML_SYCL_DISABLE_GRAPH=0  ./bin/test-backend-ops -b SYCL0 -o RWKV_WKV7 2> /dev/null | grep -c "urCommandBufferUpdateKernelLaunchExp"
8

$ SYCL_UR_TRACE=2 GGML_SYCL_DISABLE_GRAPH=0  ./bin/test-backend-ops -b SYCL0 -o RWKV_WKV7 2> /dev/null | grep -c "urCommandBufferCreateExp"
24
```

Level-Zero backend to DPC++ by default doesn't support the full graph aspect. Before this change, the code was falling back to the non-graph path.
```
$ SYCL_UR_TRACE=2 GGML_SYCL_DISABLE_GRAPH=0  ./bin/test-backend-ops -b SYCL0 -o RWKV_WKV7 2> /dev/null | grep -c "urCommandBufferCreateExp"
0
```

With this change, I observe the graph path being used on Level-Zero
```
$ SYCL_UR_TRACE=2 GGML_SYCL_DISABLE_GRAPH=0  ./bin/test-backend-ops -b SYCL0 -o RWKV_WKV7 2> /dev/null | grep -c "urCommandBufferCreateExp"
96
```

Changed files:
filename: ggml/src/ggml-sycl/ggml-sycl.cpp
patch: @@ -3653,7 +3653,8 @@ static ggml_status ggml_backend_sycl_graph_compute(ggml_backend_t backend, ggml_
 
 #ifdef GGML_SYCL_GRAPH
     if (!g_ggml_sycl_disable_graph) {
-        if (!sycl_ctx->exec_graph && !dpct::get_device(sycl_ctx->device).has(sycl::aspect::ext_oneapi_graph)) {
+        const bool graph_support = dpct::get_device(sycl_ctx->device).has(sycl::aspect::ext_oneapi_limited_graph);
+        if (!graph_support) {
             GGML_SYCL_DEBUG("[SYCL-GRAPH] can not use graphs on device:%d\n", sycl_ctx->device);
             ggml_backend_sycl_graph_compute_impl(sycl_ctx, cgraph);
             return GGML_STATUS_SUCCESS;
@@ -3664,8 +3665,10 @@ static ggml_status ggml_backend_sycl_graph_compute(ggml_backend_t backend, ggml_
         ggml_backend_sycl_graph_compute_impl(sycl_ctx, cgraph);
         model_sycl_graph.end_recording();
 
-        if (!sycl_ctx->exec_graph) {
-            auto exec_graph = model_sycl_graph.finalize({sycl_ex::property::graph::updatable{}});
+        const bool graph_update_support = dpct::get_device(sycl_ctx->device).has(sycl::aspect::ext_oneapi_graph);
+        if (!sycl_ctx->exec_graph || !graph_update_support) {
+            auto exec_graph = graph_update_support ? model_sycl_graph.finalize(sycl_ex::property::graph::updatable{}) :
+                                                     model_sycl_graph.finalize();
             sycl_ctx->exec_graph = std::make_unique<
                 sycl_ex::command_graph<sycl_ex::graph_state::executable>>(exec_graph);
         } else {