[SYCL][Graph] Support for native-command

Support [sycl_ext_codeplay_enqueue_native_command](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc) with SYCL-Graph.

Introduces `interop_handle::ext_codeplay_get_native_graph<backend>()` to
give the user access to the native graph object which native commands
can be appended to.

To use CUDA as an example, code using `ext_codeplay_enqueue_native_command`
eagerly can be updated from:

```cpp
 CGH.ext_codeplay_enqueue_native_command([=](interop_handle IH) {
       auto NativeStream = IH.get_native_queue<cuda>();
       myNativeLibraryCall(NativeStream);
}
```

To

```cpp
CGH.ext_codeplay_enqueue_native_command([=](interop_handle IH) {
     if (IH.ext_codeplay_has_graph())  {
       auto NativeGraph = IH.ext_codeplay_get_native_graph<cuda>();
       auto NativeStream = IH.get_native_queue<cuda>();

       // Start capture stream calls into graph
       cuStreamBeginCaptureToGraph(NativeStream, NativeGraph, nullptr,
                                             nullptr, 0,
                                             CU_STREAM_CAPTURE_MODE_GLOBAL);

       myNativeLibraryCall(NativeStream);

       // Stop capturing stream calls into graph
       cuStreamEndCapture(NativeStream, &NativeGraph);
     } else {
       auto NativeStream = IH.get_native_queue<cuda>();
       myNativeLibraryCall(NativeStream );
    }
}
```

Example of how this integration could work in GROMACS https://gitlab.com/gromacs/gromacs/-/merge_requests/4954

Author: EwanC

sycl/doc/extensions/experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc

@@ -173,13 +173,99 @@ dependencies are satisfied.
 The SYCL command described above completes once all of the native asynchronous
 tasks it contains have completed.
 
+TODO - reword
 The call to `interopCallable` must not submit any synchronous tasks to the
-native backend object, and it must not block waiting for any tasks to complete.
-The call also must not add tasks to backend objects that underly any other
-queue, aside from the queue that is associated with this handler. If it does
-any of these things, the behavior is undefined.
+native backend object. The call must also not block
+waiting for any tasks to complete. Additionally, the call must not add tasks to
+backend objects that underly any other queue, aside from the queue that is
+associated with this handler. If it does any of these things, the behavior is undefined.
 
-== Example
+=== SYCL-Graph Interaction
+
+This section defines the interaction with the
+link:../experimental/sycl_ext_oneapi_graph.asciidoc[sycl_ext_oneapi_graph]
+extension.
+
+The `interopCallable` object will be invoked during `command_graph::finalize()`
+when the backend object for the graph is available to give to the user as a
+handle. The user then may add nodes using native APIs to the backend graph
+object queried with `interop_handle::ext_codeplay_get_native_graph()`. The
+runtime will schedule the dependencies of the user added nodes such
+that they respect the graph node edges.
+
+==== Interop Handle Class Modifications
+
+```c++
+using graph = ext::oneapi::experimental::command_graph<
+      ext::oneapi::experimental::graph_state::executable>;
+
+class interop_handle {
+  bool ext_codeplay_has_graph() const;
+
+  template <backend Backend>
+  backend_return_t<Backend, graph> ext_codeplay_get_native_graph() const;
+
+};
+```
+
+Table {counter: tableNumber}. Native types for
+`template <backend Backend, class T> backend_return_t<Backend, T>` for `T` as
+`command_graph<graph_state::executable`.
+
+[cols="2a,a"]
+|===
+|Backend|Native graph type
+
+| `backend::opencl`
+| `cl_command_buffer_khr`
+
+| `backend::ext_oneapi_level_zero`
+| `ze_command_list_handle_t`
+
+| `backend::ext_oneapi_cuda`
+| `CUGraph`
+
+| `backend::ext_oneapi_hip`
+| `hipGraph_t`
+
+|===
+
+==== New Interop Handle Member Functions
+
+Table {counter: tableNumber}. Additional member functions of the `sycl::interop_handle` class.
+[cols="2a,a"]
+|===
+|Member function|Description
+
+|
+[source,c++]
+----
+bool interop_handle::ext_codeplay_has_graph() const;
+----
+
+| Query if the `interop_handle` object has a native graph object available.
+
+|
+[source,c++]
+----
+template <backend Backend>
+backend_return_t<Backend, graph>
+interop_handle::ext_codeplay_get_native_graph() const;
+----
+
+| Return the native graph object associated with the `interop_handle`.
+
+Exceptions:
+
+* Throws with error code `invalid` if there is no native graph object
+  associated with the interop handle.
+
+|===
+
+
+== Examples
+
+=== HIP Native Task
 
 This example demonstrates how to use this extension to enqueue asynchronous
 native tasks on the HIP backend.
@@ -207,11 +293,29 @@ q.submit([&](sycl::handler &cgh) {
 q.wait();
 ```
 
-== Issues
-
-=== sycl_ext_oneapi_graph
+=== CUDA Stream Record Native Task.
 
-`ext_codeplay_enqueue_native_command`
-cannot be used in graph nodes. A synchronous exception will be thrown with error
-code `invalid` if a user tries to add them to a graph.
+This example demonstrates how to use this extension to add stream recorded
+native nodes to a SYCL-Graph object on the CUDA backend.
 
+```
+CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
+     if (IH.ext_codeplay_has_graph())  {
+       auto NativeGraph = IH.ext_codeplay_get_native_graph<cuda>();
+       auto NativeStream = IH.get_native_queue<cuda>();
+
+       // Start capture stream calls into graph
+       cuStreamBeginCaptureToGraph(NativeStream, NativeGraph, nullptr,
+                                   nullptr, 0,
+                                   CU_STREAM_CAPTURE_MODE_GLOBAL);
+
+       myNativeLibraryCall(NativeStream);
+
+       // Stop capturing stream calls into graph
+       cuStreamEndCapture(NativeStream, &NativeGraph);
+     } else {
+       auto NativeStream = IH.get_native_queue<cuda>();
+       myNativeLibraryCall(NativeStream );
+    }
+}
+```

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -2119,13 +2119,9 @@ extensions.
 
 ==== sycl_ext_codeplay_enqueue_native_command
 
-`ext_codeplay_enqueue_native_command`, defined in
+`ext_codeplay_enqueue_native_command` commands, defined in
 link:../experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc[sycl_ext_codeplay_enqueue_native_command]
-cannot be used in graph nodes. A synchronous exception will be thrown with error
-code `invalid` if a user tries to add them to a graph.
-
-Removing this restriction is something we may look at for future revisions of
-`sycl_ext_oneapi_graph`.
+can be used in graph nodes. See the section on SYCL-Graph interaction.
 
 ==== sycl_ext_intel_queue_index
 

sycl/include/sycl/detail/backend_traits_cuda.hpp

@@ -17,13 +17,15 @@
 #include <sycl/detail/backend_traits.hpp>
 #include <sycl/device.hpp>
 #include <sycl/event.hpp>
+#include <sycl/ext/oneapi/experimental/graph.hpp>
 #include <sycl/queue.hpp>
 
 typedef int CUdevice;
 typedef struct CUctx_st *CUcontext;
 typedef struct CUstream_st *CUstream;
 typedef struct CUevent_st *CUevent;
 typedef struct CUmod_st *CUmodule;
+typedef struct CUgraph_st *CUgraph;
 
 // As defined in the CUDA 10.1 header file. This requires CUDA version > 3.2
 #if defined(_WIN64) || defined(__LP64__)
@@ -102,6 +104,12 @@ template <> struct BackendReturn<backend::ext_oneapi_cuda, queue> {
   using type = CUstream;
 };
 
+using graph = ext::oneapi::experimental::command_graph<
+    ext::oneapi::experimental::graph_state::executable>;
+template <> struct BackendReturn<backend::ext_oneapi_cuda, graph> {
+  using type = CUgraph;
+};
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl

sycl/include/sycl/detail/backend_traits_hip.hpp

@@ -17,6 +17,7 @@
 #include <sycl/detail/backend_traits.hpp>
 #include <sycl/device.hpp>
 #include <sycl/event.hpp>
+#include <sycl/ext/oneapi/experimental/graph.hpp>
 #include <sycl/queue.hpp>
 
 typedef int HIPdevice;
@@ -25,6 +26,8 @@ typedef struct ihipStream_t *HIPstream;
 typedef struct ihipEvent_t *HIPevent;
 typedef struct ihipModule_t *HIPmodule;
 typedef void *HIPdeviceptr;
+typedef struct ihipGraph *HIPGraph;
+typedef struct hipGraphNode *HIPGraphNode;
 
 namespace sycl {
 inline namespace _V1 {
@@ -96,6 +99,12 @@ template <> struct BackendReturn<backend::ext_oneapi_hip, queue> {
   using type = HIPstream;
 };
 
+using graph = ext::oneapi::experimental::command_graph<
+    ext::oneapi::experimental::graph_state::executable>;
+template <> struct BackendReturn<backend::ext_oneapi_hip, graph> {
+  using type = HIPGraph;
+};
+
 template <> struct InteropFeatureSupportMap<backend::ext_oneapi_hip> {
   static constexpr bool MakePlatform = false;
   static constexpr bool MakeDevice = true;

sycl/include/sycl/detail/backend_traits_level_zero.hpp

@@ -20,13 +20,14 @@
 #include <sycl/device.hpp>                                  // for device
 #include <sycl/event.hpp>                                   // for event
 #include <sycl/ext/oneapi/backend/level_zero_ownership.hpp> // for ownership
-#include <sycl/image.hpp>                                   // for image
-#include <sycl/kernel.hpp>                                  // for kernel
-#include <sycl/kernel_bundle.hpp>                           // for kernel_b...
-#include <sycl/kernel_bundle_enums.hpp>                     // for bundle_s...
-#include <sycl/platform.hpp>                                // for platform
-#include <sycl/property_list.hpp>                           // for property...
-#include <sycl/range.hpp>                                   // for range
+#include <sycl/ext/oneapi/experimental/graph.hpp>
+#include <sycl/image.hpp>               // for image
+#include <sycl/kernel.hpp>              // for kernel
+#include <sycl/kernel_bundle.hpp>       // for kernel_b...
+#include <sycl/kernel_bundle_enums.hpp> // for bundle_s...
+#include <sycl/platform.hpp>            // for platform
+#include <sycl/property_list.hpp>       // for property...
+#include <sycl/range.hpp>               // for range
 
 #include <variant> // for variant
 #include <vector>  // for vector
@@ -207,6 +208,12 @@ template <> struct BackendReturn<backend::ext_oneapi_level_zero, kernel> {
   using type = ze_kernel_handle_t;
 };
 
+using graph = ext::oneapi::experimental::command_graph<
+    ext::oneapi::experimental::graph_state::executable>;
+template <> struct BackendReturn<backend::ext_oneapi_level_zero, graph> {
+  using type = ze_command_list_handle_t;
+};
+
 template <> struct InteropFeatureSupportMap<backend::ext_oneapi_level_zero> {
   static constexpr bool MakePlatform = true;
   static constexpr bool MakeDevice = true;

sycl/include/sycl/detail/backend_traits_opencl.hpp

@@ -21,9 +21,10 @@
 #include <sycl/detail/ur.hpp>             // for assertion and ur handles
 #include <sycl/device.hpp>                // for device
 #include <sycl/event.hpp>                 // for event
-#include <sycl/kernel.hpp>                // for kernel
-#include <sycl/kernel_bundle_enums.hpp>   // for bundle_state
-#include <sycl/platform.hpp>              // for platform
+#include <sycl/ext/oneapi/experimental/graph.hpp>
+#include <sycl/kernel.hpp>              // for kernel
+#include <sycl/kernel_bundle_enums.hpp> // for bundle_state
+#include <sycl/platform.hpp>            // for platform
 
 #include <vector> // for vector
 
@@ -132,6 +133,12 @@ template <> struct BackendReturn<backend::opencl, kernel> {
   using type = cl_kernel;
 };
 
+using graph = ext::oneapi::experimental::command_graph<
+    ext::oneapi::experimental::graph_state::executable>;
+template <> struct BackendReturn<backend::opencl, graph> {
+  using type = cl_command_buffer_khr;
+};
+
 template <> struct InteropFeatureSupportMap<backend::opencl> {
   static constexpr bool MakePlatform = true;
   static constexpr bool MakeDevice = true;

sycl/include/sycl/ext/oneapi/experimental/graph.hpp

@@ -114,7 +114,8 @@ enum class node_type {
   prefetch = 6,
   memadvise = 7,
   ext_oneapi_barrier = 8,
-  host_task = 9
+  host_task = 9,
+  native_command = 10
 };
 
 /// Class representing a node in the graph, returned by command_graph::add().

sycl/include/sycl/handler.hpp

@@ -1934,9 +1934,6 @@ class __SYCL_EXPORT handler {
                                               void(interop_handle)>::value>
   ext_codeplay_enqueue_native_command([[maybe_unused]] FuncT &&Func) {
 #ifndef __SYCL_DEVICE_ONLY__
-    throwIfGraphAssociated<
-        ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
-            sycl_ext_codeplay_enqueue_native_command>();
     ext_codeplay_enqueue_native_command_impl(Func);
 #endif
   }

sycl/include/sycl/interop_handle.hpp

@@ -16,6 +16,7 @@
 #include <sycl/detail/impl_utils.hpp> // for getSyclObjImpl
 #include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/accessor_property_list.hpp> // for accessor_property_list
+#include <sycl/ext/oneapi/experimental/graph.hpp>
 #include <sycl/image.hpp>                             // for image
 #include <ur_api.h> // for ur_mem_handle_t, ur...
 
@@ -49,6 +50,9 @@ class interop_handle {
   /// interop_handle.
   __SYCL_EXPORT backend get_backend() const noexcept;
 
+  /// Returns true if command-group is being added to a graph as a node
+  __SYCL_EXPORT bool ext_codeplay_has_graph() const noexcept;
+
   /// Receives a SYCL accessor that has been defined as a requirement for the
   /// command group, and returns the underlying OpenCL memory object that is
   /// used by the SYCL runtime. If the accessor passed as parameter is not part
@@ -134,6 +138,26 @@ class interop_handle {
 #endif
   }
 
+  using graph = ext::oneapi::experimental::command_graph<
+      ext::oneapi::experimental::graph_state::executable>;
+  template <backend Backend = backend::opencl>
+  backend_return_t<Backend, graph> ext_codeplay_get_native_graph() const {
+#ifndef __SYCL_DEVICE_ONLY__
+    // TODO: replace the exception thrown below with the SYCL 2020 exception
+    // with the error code 'errc::backend_mismatch' when those new exceptions
+    // are ready to be used.
+    if (Backend != get_backend())
+      throw exception(make_error_code(errc::invalid),
+                      "Incorrect backend argument was passed");
+
+    // C-style cast required to allow various native types
+    return (backend_return_t<Backend, graph>)getNativeGraph();
+#else
+    // we believe this won't be ever called on device side
+    return 0;
+#endif
+  }
+
   /// Returns the SYCL application interoperability native backend object
   /// associated with the device associated with the SYCL queue that the host
   /// task was submitted to. The native backend object returned must be in
@@ -186,8 +210,9 @@ class interop_handle {
   interop_handle(std::vector<ReqToMem> MemObjs,
                  const std::shared_ptr<detail::queue_impl> &Queue,
                  const std::shared_ptr<detail::device_impl> &Device,
-                 const std::shared_ptr<detail::context_impl> &Context)
-      : MQueue(Queue), MDevice(Device), MContext(Context),
+                 const std::shared_ptr<detail::context_impl> &Context,
+                 const ur_exp_command_buffer_handle_t &Graph)
+      : MQueue(Queue), MDevice(Device), MContext(Context), MGraph(Graph),
         MMemObjs(std::move(MemObjs)) {}
 
   template <backend Backend, typename DataT, int Dims>
@@ -211,10 +236,12 @@ class interop_handle {
   getNativeQueue(int32_t &NativeHandleDesc) const;
   __SYCL_EXPORT ur_native_handle_t getNativeDevice() const;
   __SYCL_EXPORT ur_native_handle_t getNativeContext() const;
+  __SYCL_EXPORT ur_native_handle_t getNativeGraph() const;
 
   std::shared_ptr<detail::queue_impl> MQueue;
   std::shared_ptr<detail::device_impl> MDevice;
   std::shared_ptr<detail::context_impl> MContext;
+  ur_exp_command_buffer_handle_t MGraph;
 
   std::vector<ReqToMem> MMemObjs;
 };

sycl/source/detail/graph_impl.cpp

@@ -828,6 +828,8 @@ ur_exp_command_buffer_sync_point_t exec_graph_impl::enqueueNode(
     std::shared_ptr<node_impl> Node) {
 
   // Queue which will be used for allocation operations for accessors.
+  // Will also be used in native commands to return to the user in
+  // `interop_handler::get_native_queue()` calls
   auto AllocaQueue = std::make_shared<sycl::detail::queue_impl>(
       DeviceImpl, sycl::detail::getSyclObjImpl(Ctx), sycl::async_handler{},
       sycl::property_list{});

sycl/source/detail/graph_impl.hpp

@@ -71,6 +71,8 @@ inline node_type getNodeTypeFromCG(sycl::detail::CGType CGType) {
     return node_type::host_task;
   case sycl::detail::CGType::ExecCommandBuffer:
     return node_type::subgraph;
+  case sycl::detail::CGType::EnqueueNativeCommand:
+    return node_type::native_command;
   default:
     assert(false && "Invalid Graph Node Type");
     return node_type::empty;
@@ -704,6 +706,9 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
     case sycl::detail::CGType::ExecCommandBuffer:
       Stream << "CGExecCommandBuffer \\n";
       break;
+    case sycl::detail::CGType::EnqueueNativeCommand:
+      Stream << "CGNativeCommand \\n";
+      break;
     default:
       Stream << "Other \\n";
       break;