Iterate on spec and tests

Author: EwanC

sycl/doc/extensions/experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc

@@ -258,12 +258,11 @@ interop_handle::ext_codeplay_get_native_graph() const;
 
 Exceptions:
 
-* Throws with error code `invalid` if there is no native graph object
-  associated with the interop handle.
+* Throws with error code `invalid` if `ext_codeplay_has_graph()` returns
+  `false`.
 
 |===
 
-
 == Examples
 
 === HIP Native Task
@@ -284,27 +283,122 @@ q.submit([&](sycl::handler &cgh) {
 
       // Can enqueue arbitrary backend operations. This could also be a kernel
       // launch or call to a library that enqueues operations on the stream etc
-      //
-      // Important: Enqueuing a *synchronous* backend operation results in
-      // undefined behavior.
       hipMemcpyAsync(target_ptr, native_mem, test_size * sizeof(int),
                       hipMemcpyDeviceToHost, stream);
     });
   });
 q.wait();
 ```
 
-=== CUDA Stream Record Native Task
+=== Level-Zero Add Native Graph Node
+
+This example demonstrates how to use this extension to add a native command
+to a SYCL-Graph object on the Level-Zero backend. The command is doing a memory
+copy between two USM pointers.
+
+```c++
+Graph.add([&](sycl::handler &CGH) {
+    CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
+      ze_command_list_handle_t NativeGraph =
+          IH.ext_codeplay_get_native_graph<sycl::backend::ext_oneapi_level_zero>();
+
+      zeCommandListAppendMemoryCopy(
+          NativeGraph, PtrY, PtrX, Size * sizeof(int), nullptr, 0, nullptr);
+    });
+  });
+```
+
+=== OpenCL Add Native Graph Node
+
+This example demonstrates how to use this extension to add a native command to
+a SYCL-Graph object on the OpenCL backend. The command is doing a copy between
+two buffer objects.
+
+```c++
+sycl::queue Queue;
+auto Platform = get_native<sycl::backend::opencl>(Queue.get_context().get_platform());
+clCommandCopyBufferKHR_fn clCommandCopyBufferKHR =
+    reinterpret_cast<clCommandCopyBufferKHR_fn>(
+        clGetExtensionFunctionAddressForPlatform(Platform, "clCommandCopyBufferKHR"));
+
+Graph.add([&](sycl::handler &CGH) {
+    auto AccX = BufX.get_access(CGH);
+    auto AccY = BufY.get_access(CGH);
+    CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
+        cl_command_buffer_khr NativeGraph =
+            IH.ext_codeplay_get_native_graph<sycl::backend::opencl>();
+        auto SrcBuffer = IH.get_native_mem<sycl::backend::opencl>(AccX);
+        auto DstBuffer = IH.get_native_mem<sycl::backend::opencl>(AccY);
+
+        clCommandCopyBufferKHR(
+            NativeGraph, nullptr, nullptr, SrcBuffer[0], DstBuffer[0], 0, 0,
+            Size * sizeof(int), 0, nullptr, nullptr, nullptr);
+    });
+  });
+```
+
+=== CUDA Add Native Graph Node
+
+This example demonstrates how to use this extension to add a native command to
+a SYCL-Graph object on the CUDA backend. The command is doing a memory copy
+between two device USM pointers.
+
+```c++
+Graph.add([&](sycl::handler &CGH) {
+    CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
+      CUgraph NativeGraph =
+          IH.ext_codeplay_get_native_graph<sycl::backend::ext_oneapi_cuda>();
+
+      CUDA_MEMCPY3D Params;
+      std::memset(&Params, 0, sizeof(CUDA_MEMCPY3D));
+      Params.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+      Params.srcDevice = (CUdeviceptr)PtrX;
+      Params.srcHost = nullptr;
+      Params.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+      Params.dstDevice = (CUdeviceptr)PtrY;
+      Params.dstHost = nullptr;
+      Params.WidthInBytes = Size * sizeof(int);
+      Params.Height = 1;
+      Params.Depth = 1;
+
+      CUgraphNode Node;
+      CUcontext Context = IH.get_native_context<sycl::backend::ext_oneapi_cuda>();
+       cuGraphAddMemcpyNode(&Node, NativeGraph, nullptr, 0, &Params, Context);
+    });
+  });
+```
+
+=== HIP Add Native Graph Node
+
+This example demonstrates how to use this extension to add a native command to
+a SYCL-Graph object on the HIP backend. The command is doing a memory copy
+between two device USM pointers.
+
+```c++
+Graph.add([&](sycl::handler &CGH) {
+    CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
+      HIPGraph NativeGraph =
+          IH.ext_codeplay_get_native_graph<sycl::backend::ext_oneapi_hip>();
+
+      HIPGraphNode Node;
+      hipGraphAddMemcpyNode1D(&Node, NativeGraph, nullptr, 0, PtrY, PtrX,
+                              Size * sizeof(int), hipMemcpyDefault);
+    });
+  });
+```
+
+=== CUDA Stream Record Native Graph Nodes
 
 This example demonstrates how to use this extension to add stream recorded
 native nodes to a SYCL-Graph object on the CUDA backend.
 
-```
-q.submit([&](sycl::handler &cgh) {
-    cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle h) {
+```c++
+q.submit([&](sycl::handler &CGH) {
+    CGH.ext_codeplay_enqueue_native_command([=](sycl::interop_handle IH) {
         auto NativeStream = h.get_native_queue<cuda>();
-        if (h.ext_codeplay_has_graph())  {
-            auto NativeGraph = h.ext_codeplay_get_native_graph<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,

sycl/include/sycl/interop_handle.hpp

@@ -208,25 +208,22 @@ class interop_handle {
   friend class detail::DispatchHostTask;
   using ReqToMem = std::pair<detail::AccessorImplHost *, ur_mem_handle_t>;
 
-#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
-  // MGraph should become a member of this class on the next ABI breaking
-  // window
-  // TODO create and link GitHub Issue
   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,
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+                 [[maybe_unused]]
+#endif
                  ur_exp_command_buffer_handle_t Graph = nullptr)
-      : MQueue(Queue), MDevice(Device), MContext(Context), MGraph(Graph),
-        MMemObjs(std::move(MemObjs)) {}
-#else
-  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),
-        MMemObjs(std::move(MemObjs)) {}
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+        // CMPLRLLVM-66082 - MGraph should become a member of this class on the
+        // next ABI breaking window.
+        MGraph(Graph),
 #endif
+        MMemObjs(std::move(MemObjs)) {
+  }
 
   template <backend Backend, typename DataT, int Dims>
   backend_return_t<Backend, buffer<DataT, Dims>>
@@ -255,9 +252,8 @@ class interop_handle {
   std::shared_ptr<detail::device_impl> MDevice;
   std::shared_ptr<detail::context_impl> MContext;
 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
-  // MGraph should become a member of this class on the next ABI breaking
-  // window
-  // TODO link github issue
+  // CMPLRLLVM-66082 - MGraph should become a member of this class on the
+  // next ABI breaking window.
   ur_exp_command_buffer_handle_t MGraph;
 #endif
 

sycl/source/detail/queue_impl.hpp

@@ -754,9 +754,9 @@ class queue_impl {
   }
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // CMPLRLLVM-66082
   // These methods are for accessing a member that should live in the
   // sycl::interop_handle class and will be moved on next ABI breaking window.
-  // TODO create and link github issue
   ur_exp_command_buffer_handle_t getInteropGraph() const {
     return MInteropGraph;
   }
@@ -1014,10 +1014,10 @@ class queue_impl {
   std::weak_ptr<ext::oneapi::experimental::detail::graph_impl> MGraph{};
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // CMPLRLLVM-66082
   // This member should be part of the sycl::interop_handle class, but it
   // in an API breaking change. So member lives here temporarily where it can
   // be accessed through the queue member of the interop_handle
-  // TODO create and link github issue
   ur_exp_command_buffer_handle_t MInteropGraph{};
 #endif
 

sycl/source/detail/scheduler/commands.cpp

@@ -3066,12 +3066,12 @@ ur_result_t ExecCGCommand::enqueueImpCommandBuffer() {
     // available if a user asks for them inside the interop task scope
     std::vector<interop_handle::ReqToMem> ReqToMem;
     const std::vector<Requirement *> &HandlerReq = HostTask->getRequirements();
-    auto ReqToMemConv = [&ReqToMem, HostTask](Requirement *Req) {
+    auto ReqToMemConv = [&ReqToMem, ContextImpl](Requirement *Req) {
       const std::vector<AllocaCommandBase *> &AllocaCmds =
           Req->MSYCLMemObj->MRecord->MAllocaCommands;
 
       for (AllocaCommandBase *AllocaCmd : AllocaCmds)
-        if (getContext(HostTask->MQueue) == getContext(AllocaCmd->getQueue())) {
+        if (ContextImpl == getContext(AllocaCmd->getQueue())) {
           auto MemArg =
               reinterpret_cast<ur_mem_handle_t>(AllocaCmd->getMemAllocation());
           ReqToMem.emplace_back(std::make_pair(Req, MemArg));
@@ -3090,10 +3090,10 @@ ur_result_t ExecCGCommand::enqueueImpCommandBuffer() {
     ur_exp_command_buffer_handle_t InteropCommandBuffer =
         ChildCommandBuffer ? ChildCommandBuffer : MCommandBuffer;
 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+    // CMPLRLLVM-66082
     // The native command-buffer should be a member of the sycl::interop_handle
-    // class, but it is in an API breaking change to add it. So member lives in
+    // class, but it is in an ABI breaking change to add it. So member lives in
     // the queue as a intermediate workaround.
-    // TODO create and link github issue
     interop_handle IH{ReqToMem, MQueue, DeviceImpl, ContextImpl,
                       InteropCommandBuffer};
 #else
@@ -3109,6 +3109,11 @@ ur_result_t ExecCGCommand::enqueueImpCommandBuffer() {
         MSyncPointDeps.empty() ? nullptr : MSyncPointDeps.data(),
         &OutSyncPoint);
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+    // See CMPLRLLVM-66082
+    MQueue->setInteropGraph(nullptr);
+#endif
+
     if (ChildCommandBuffer) {
       ur_result_t Res = Adapter->call_nocheck<
           sycl::detail::UrApiKind::urCommandBufferReleaseExp>(

sycl/source/interop_handle.cpp

@@ -25,7 +25,8 @@ backend interop_handle::get_backend() const noexcept {
 
 bool interop_handle::ext_codeplay_has_graph() const noexcept {
 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
-  // TODO create and link GitHub
+  // CMPLRLLVM-66082 - MGraph should become a member of this class on the
+  // next ABI breaking window.
   return MGraph != nullptr;
 #else
   return MQueue->getInteropGraph() != nullptr;
@@ -64,7 +65,8 @@ interop_handle::getNativeQueue(int32_t &NativeHandleDesc) const {
 
 ur_native_handle_t interop_handle::getNativeGraph() const {
 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
-  // TODO create and link github issue
+  // CMPLRLLVM-66082 - MGraph should become a member of this class on the
+  // next ABI breaking window.
   auto Graph = MGraph;
 #else
   auto Graph = MQueue->getInteropGraph();

sycl/test-e2e/EnqueueNativeCommand/Graph/cuda_explicit_usm.cpp

@@ -0,0 +1,85 @@
+// RUN: %{build} -o %t.out %cuda_options
+// RUN: %{run} %t.out
+// RUN: %if preview-breaking-changes-supported %{ %{build} -fpreview-breaking-changes -o %t2.out %cuda_options %}
+// RUN: %if preview-breaking-changes-supported %{ %{run} %t2.out %}
+// REQUIRES: cuda, cuda_dev_kit
+
+#include <cuda.h>
+#include <sycl/backend.hpp>
+#include <sycl/ext/oneapi/experimental/graph.hpp>
+#include <sycl/interop_handle.hpp>
+#include <sycl/usm.hpp>
+
+namespace exp_ext = sycl::ext::oneapi::experimental;
+using namespace sycl;
+
+int main() {
+  queue Queue;
+
+  const size_t Size = 128;
+  int *PtrX = malloc_device<int>(Size, Queue);
+  int *PtrY = malloc_device<int>(Size, Queue);
+
+  exp_ext::command_graph Graph{Queue};
+
+  auto NodeA = Graph.add([&](handler &CGH) {
+    CGH.single_task([=]() {
+      for (size_t i = 0; i < Size; i++) {
+        PtrX[i] = i;
+        PtrY[i] = 0;
+      }
+    });
+  });
+
+  auto NodeB = Graph.add(
+      [&](handler &CGH) {
+        CGH.ext_codeplay_enqueue_native_command([=](interop_handle IH) {
+          if (!IH.ext_codeplay_has_graph()) {
+            assert(false && "Native Handle should have a graph");
+          }
+          CUgraph NativeGraph =
+              IH.ext_codeplay_get_native_graph<backend::ext_oneapi_cuda>();
+
+          CUDA_MEMCPY3D Params;
+          std::memset(&Params, 0, sizeof(CUDA_MEMCPY3D));
+          Params.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+          Params.srcDevice = (CUdeviceptr)PtrX;
+          Params.srcHost = nullptr;
+          Params.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+          Params.dstDevice = (CUdeviceptr)PtrY;
+          Params.dstHost = nullptr;
+          Params.WidthInBytes = Size * sizeof(int);
+          Params.Height = 1;
+          Params.Depth = 1;
+
+          CUgraphNode Node;
+          CUcontext Context = IH.get_native_context<backend::ext_oneapi_cuda>();
+          auto Res = cuGraphAddMemcpyNode(&Node, NativeGraph, nullptr, 0,
+                                          &Params, Context);
+          assert(Res == CUDA_SUCCESS);
+        });
+      },
+      exp_ext::property::node::depends_on(NodeA));
+
+  Graph.add(
+      [&](handler &CGH) {
+        CGH.parallel_for(range<1>{Size}, [=](id<1> it) { PtrY[it] *= 2; });
+      },
+      exp_ext::property::node::depends_on(NodeB));
+
+  auto ExecGraph = Graph.finalize();
+  Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  std::vector<int> HostData(Size);
+
+  Queue.copy(PtrY, HostData.data(), Size).wait();
+  for (size_t i = 0; i < Size; i++) {
+    const int Ref = i * 2;
+    assert(Ref == HostData[i]);
+  }
+
+  free(PtrX, Queue);
+  free(PtrY, Queue);
+
+  return 0;
+}