SYCL: Fix test-backend-ops crashes with SYCL-Graph

Currently on a CUDA backend to SYCL when running
`GGML_SYCL_DISABLE_GRAPH=0 ./bin/test-backend-ops -b SYCL0` I
see crashes from 3 operations:

1) `-o MUL_MAT`: Issue arising from recording of oneMath `ext_codeplay_enqueue_native_command`.
2) `-o CONCAT` : Use of blocking waits on a queue that's being recorded https://github.com/ggml-org/llama.cpp/blob/master/ggml/src/ggml-sycl/concat.cpp#L185-L187, can these wait calls just be removed?
3) `-o MUL_MAT_ID`: Blocking wait on a recording queue for a copy to host memory https://github.com/ggml-org/llama.cpp/blob/master/ggml/src/ggml-sycl/ggml-sycl.cpp#L3072-L3074 , host work could be wrapped in a host-task?

For 1) I have come up with a oneMath fix in uxlfoundation/oneMath#669, I've put a provisional git tag to pull in this PR for testing, but will update to the upstream commit once merged.

For 2 & 3) we've noticed that `ggml-cuda.cu` has the
[check_node_graph_compatibility_and_refresh_copy_ops](https://github.com/ggml-org/llama.cpp/blob/39e73ae0d69f882d7e29cecc6dd8f5052fca6731/ggml/src/ggml-cuda/ggml-cuda.cu#L2458-L2458) method for checking if a graph can be used, even if enabled. I've taken a similar approach in this PR by adding a method to `ggml-sycl.cpp` for checking if a graph can be used for the operations even if a user has asked for it to be enabled.

Author: Ewan Crawford

ggml/src/ggml-sycl/CMakeLists.txt

@@ -135,8 +135,8 @@ else()
         endif()
         FetchContent_Declare(
             ONEMATH
-            GIT_REPOSITORY https://github.com/uxlfoundation/oneMath.git
-            GIT_TAG c255b1b4c41e2ee3059455c1f96a965d6a62568a
+            GIT_REPOSITORY https://github.com/EwanC/oneMath.git
+            GIT_TAG 671d9bcc5aa6cc52cce6b6518c9b8f126c0352f4
         )
         FetchContent_MakeAvailable(ONEMATH)
         # Create alias to match with find_package targets name

ggml/src/ggml-sycl/ggml-sycl.cpp

@@ -3639,11 +3639,37 @@ static void ggml_backend_sycl_graph_compute_impl(ggml_backend_sycl_context * syc
     }
 }
 
+#ifdef GGML_SYCL_GRAPH
+static bool check_node_graph_compatibility(ggml_cgraph * cgraph) {
+    for (int i = 0; i < cgraph->n_nodes; i++) {
+        ggml_tensor * node = cgraph->nodes[i];
+        switch(node->op) {
+            default: break;
+            case GGML_OP_CONCAT:
+            // ggml_sycl_op_concat() does a blocking host wait after memcpy operations,
+            // but wait() can't be called on the events returned by a queue recording
+            // to a graph.
+            [[fallthrough]];
+            case GGML_OP_MUL_MAT_ID:
+            // ggml_sycl_mul_mat_id() does a blocking host wait on the sycl queue after
+            // submitting a memcpy operation, but wait() can't be called on a queue that
+            // is recording to a graph.
+#ifndef NDEBUG
+            GGML_LOG_DEBUG("%s: disabling SYCL graphs due to unsupported node type\n", __func__);
+#endif
+            return false;
+        }
+    }
+    return true;
+}
+#endif
+
 static ggml_status ggml_backend_sycl_graph_compute(ggml_backend_t backend, ggml_cgraph * cgraph) {
     auto * sycl_ctx = static_cast<ggml_backend_sycl_context *>(backend->context);
 
 #ifdef GGML_SYCL_GRAPH
-    if (!g_ggml_sycl_disable_graph) {
+    bool use_sycl_graph = !g_ggml_sycl_disable_graph && check_node_graph_compatibility(cgraph);
+    if (use_sycl_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);