Address editorial feedback from Gordon

Author: EwanC

sycl/doc/extensions/experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc

@@ -177,15 +177,15 @@ dependencies are satisfied.
 The SYCL command described above completes once all of the native asynchronous
 tasks it contains have completed.
 
-The call to `interopCallable` is done by the host thread, users should
-therefore not perform blocking or synchronization tasks inside
-`interopCallable`, as it will defer the host thread returning to SYCL user
-code.
-
 The call to `interopCallable` must not add tasks to backend objects that underly
 any other queue, aside from the queue that is associated with this handler,
 otherwise, the behavior is undefined.
 
+[_Note:_ The function object `interopCallable` is invoked to enqueue commands to a
+native queue or graph and therefore, APIs which block or synchronize could
+prolong or interfere with other commands being enqueued to the backend.
+_{endnote}_]
+
 === SYCL-Graph Interaction
 
 This section defines the interaction with the
@@ -202,6 +202,8 @@ that they respect the graph node edges.
 ==== Interop Handle Class Modifications
 
 ```c++
+// Alias is for editorial brevity in the ext_codeplay_get_native_graph
+// definition, and is non-normative.
 using graph = ext::oneapi::experimental::command_graph<
       ext::oneapi::experimental::graph_state::executable>;
 
@@ -213,28 +215,6 @@ class interop_handle {
 };
 ```
 
-Table {counter: tableNumber}. Native types for
-`template <backend Backend, class T> backend_return_t<Backend, T>` where `T` is
-instantiated 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.
@@ -248,11 +228,16 @@ Table {counter: tableNumber}. Additional member functions of the `sycl::interop_
 bool interop_handle::ext_codeplay_has_graph() const;
 ----
 
-| Query if the `interop_handle` object has a native graph object available.
-  Note that host-task nodes in a `command_graph` will return `false` for this,
-  as host-task commands are executed by the SYCL runtime rather than the
-  backend device, so there is no native graph object associated with the
-  command.
+|
+_Returns_: True if the `interop_handle object` was constructed and passed to
+an enqueue native command function object by `ext_codeplay_enqueue_native_command`,
+that was invoked when adding a graph node, either explicitly or implicitly
+via graph record.
+
+[_Note:_ that host-task nodes in a `command_graph` will return `false` from this
+query, as the host-task callable is invoked during graph execution rather than
+graph finalization.
+_{endnote}_]
 
 |
 [source,c++]
@@ -262,12 +247,42 @@ backend_return_t<Backend, graph>
 interop_handle::ext_codeplay_get_native_graph() const;
 ----
 
-| Return the native graph object associated with the `interop_handle`.
+|
+_Returns_: The native graph object associated with the `interop_handle`.
+
+_Throws_: An exception with the `errc::invalid` error code if
+`ext_codeplay_has_graph()` returns `false`.
+
+|===
+
+== Implementation Notes
+
+When `interop_handle::get_native_queue()` is invoked in a native command
+function object on graph finalize, the queue that is returned to the user is an
+internal queue created by the SYCL runtime, as there is no user provided queue
+at the point of graph finalization. This queue has the same device and context
+as the graph was created with. The only valid usage of this queue is to perform
+stream capture to a graph for backend APIs that provide this functionality.
+
+Table {counter: tableNumber}. Native types for
+`template <backend Backend, class T> backend_return_t<Backend, T>` where `T` is
+instantiated 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`
 
-Exceptions:
+| `backend::ext_oneapi_cuda`
+| `CUGraph`
 
-* Throws with error code `invalid` if `ext_codeplay_has_graph()` returns
-  `false`.
+| `backend::ext_oneapi_hip`
+| `hipGraph_t`
 
 |===
 

sycl/doc/syclgraph/SYCLGraphUsageGuide.md

@@ -171,6 +171,57 @@ In this case it may be necessary to first manually trigger the warmup by calling
 `Graph.begin_recording(Queue)` to prevent the warmup from being captured in a
 graph when recording.
 
+#### ext_codeplay_enqueue_native_command
+
+The SYCL-Graph extension is compatible with the
+[ext_codeplay_enqueue_native_command](../extensions/experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc)
+extension that can be used to capture asynchronous library commands as graph
+nodes. However, existing `ext_codeplay_enqueue_native_command` user code will
+need modifications to work correctly for submission to a sycl queue that can be
+in either the executable or recording state.
+
+Using the CUDA backend as an example, existing code which uses a
+native-command to invoke a library call:
+
+```c++
+q.submit([&](sycl::handler &CGH) {
+    CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
+       auto NativeStream = IH.get_native_queue<cuda>();
+       myNativeLibraryCall(NativeStream);
+    });
+});
+```
+
+Can be ported as below to work with SYCL-Graph, where the queue may be in
+a recording state. If the code is not ported but the queue is in a recording
+state, then asynchronous work in `myNativeLibraryCall` will be scheduled
+immediately as part of graph finalize, rather than being added to the graph as
+a node, which is unlikely to be the desired user behavior.
+
+```c++
+q.submit([&](sycl::handler &CGH) {
+    CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
+        auto NativeStream = h.get_native_queue<cuda>();
+        if (IH.ext_codeplay_has_graph())  {
+            auto NativeGraph =
+              IH.ext_codeplay_get_native_graph<sycl::backend::ext_oneapi_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 {
+            myNativeLibraryCall(NativeStream);
+        }
+    });
+});
+```
+
 ## Code Examples
 
 The examples below demonstrate intended usage of the extension, but may not be

sycl/include/sycl/interop_handle.hpp

@@ -144,9 +144,6 @@ class interop_handle {
   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");