[SYCL][CUDA] Add experimental cuda interop with event (#6288)

This PR is adds part of the CUDA-backend spec interop proposed in KhronosGroup/SYCL-Docs#197. The changes work with the CUDA CTS interop checks KhronosGroup/SYCL-CTS#336.

This PR just adds the event interop. 

llvm-test-suite: intel/llvm-test-suite#1053

Author: AidanBeltonS

sycl/include/sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp

@@ -132,7 +132,7 @@ template <> struct InteropFeatureSupportMap<backend::ext_oneapi_cuda> {
   static constexpr bool MakeDevice = true;
   static constexpr bool MakeContext = true;
   static constexpr bool MakeQueue = true;
-  static constexpr bool MakeEvent = false;
+  static constexpr bool MakeEvent = true;
   static constexpr bool MakeBuffer = false;
   static constexpr bool MakeKernel = false;
   static constexpr bool MakeKernelBundle = false;

sycl/include/sycl/ext/oneapi/experimental/backend/cuda.hpp

@@ -24,6 +24,14 @@ inline __SYCL_EXPORT device make_device(pi_native_handle NativeHandle) {
   return sycl::detail::make_device(NativeHandle, backend::ext_oneapi_cuda);
 }
 
+// Implementation of cuda::has_native_event
+inline __SYCL_EXPORT bool has_native_event(event sycl_event) {
+  if (sycl_event.get_backend() == backend::ext_oneapi_cuda)
+    return get_native<backend::ext_oneapi_cuda>(sycl_event) != nullptr;
+
+  return false;
+}
+
 } // namespace cuda
 } // namespace oneapi
 } // namespace ext
@@ -71,6 +79,16 @@ inline device make_device<backend::ext_oneapi_cuda>(
   return ext::oneapi::cuda::make_device(NativeHandle);
 }
 
+// CUDA event specialization
+template <>
+inline event make_event<backend::ext_oneapi_cuda>(
+    const backend_input_t<backend::ext_oneapi_cuda, event> &BackendObject,
+    const context &TargetContext) {
+  return detail::make_event(detail::pi::cast<pi_native_handle>(BackendObject),
+                            TargetContext, true,
+                            /*Backend*/ backend::ext_oneapi_cuda);
+}
+
 // CUDA queue specialization
 template <>
 inline queue make_queue<backend::ext_oneapi_cuda>(

sycl/plugins/cuda/pi_cuda.cpp

@@ -450,8 +450,8 @@ CUstream _pi_queue::get_next_transfer_stream() {
 
 _pi_event::_pi_event(pi_command_type type, pi_context context, pi_queue queue,
                      CUstream stream, pi_uint32 stream_token)
-    : commandType_{type}, refCount_{1}, hasBeenWaitedOn_{false},
-      isRecorded_{false}, isStarted_{false},
+    : commandType_{type}, refCount_{1}, has_ownership_{true},
+      hasBeenWaitedOn_{false}, isRecorded_{false}, isStarted_{false},
       streamToken_{stream_token}, evEnd_{nullptr}, evStart_{nullptr},
       evQueued_{nullptr}, queue_{queue}, stream_{stream}, context_{context} {
 
@@ -471,6 +471,13 @@ _pi_event::_pi_event(pi_command_type type, pi_context context, pi_queue queue,
   cuda_piContextRetain(context_);
 }
 
+_pi_event::_pi_event(pi_context context, CUevent eventNative)
+    : commandType_{PI_COMMAND_TYPE_USER}, refCount_{1}, has_ownership_{false},
+      hasBeenWaitedOn_{false}, isRecorded_{false}, isStarted_{false},
+      streamToken_{std::numeric_limits<pi_uint32>::max()}, evEnd_{eventNative},
+      evStart_{nullptr}, evQueued_{nullptr}, queue_{nullptr}, context_{
+                                                                  context} {}
+
 _pi_event::~_pi_event() {
   if (queue_ != nullptr) {
     cuda_piQueueRelease(queue_);
@@ -583,7 +590,11 @@ pi_result _pi_event::wait() {
 }
 
 pi_result _pi_event::release() {
+  if (!backend_has_ownership())
+    return PI_SUCCESS;
+
   assert(queue_ != nullptr);
+
   PI_CHECK_ERROR(cuEventDestroy(evEnd_));
 
   if (queue_->properties_ & PI_QUEUE_PROFILING_ENABLE) {
@@ -3910,11 +3921,19 @@ pi_result cuda_piextEventGetNativeHandle(pi_event event,
 /// \param[out] event Set to the PI event object created from native handle.
 ///
 /// \return TBD
-pi_result cuda_piextEventCreateWithNativeHandle(pi_native_handle, pi_context,
-                                                bool, pi_event *) {
-  cl::sycl::detail::pi::die(
-      "Creation of PI event from native handle not implemented");
-  return {};
+pi_result cuda_piextEventCreateWithNativeHandle(pi_native_handle nativeHandle,
+                                                pi_context context,
+                                                bool ownNativeHandle,
+                                                pi_event *event) {
+  (void)ownNativeHandle;
+  assert(!ownNativeHandle);
+
+  std::unique_ptr<_pi_event> event_ptr{nullptr};
+
+  *event = _pi_event::make_with_native(context,
+                                       reinterpret_cast<CUevent>(nativeHandle));
+
+  return PI_SUCCESS;
 }
 
 /// Creates a PI sampler object

sycl/plugins/cuda/pi_cuda.hpp

@@ -635,6 +635,8 @@ struct _pi_event {
 
   pi_uint32 get_event_id() const noexcept { return eventId_; }
 
+  bool backend_has_ownership() const noexcept { return has_ownership_; }
+
   // Returns the counter time when the associated command(s) were enqueued
   //
   pi_uint64 get_queued_time() const;
@@ -655,6 +657,10 @@ struct _pi_event {
                          stream_token);
   }
 
+  static pi_event make_with_native(pi_context context, CUevent eventNative) {
+    return new _pi_event(context, eventNative);
+  }
+
   pi_result release();
 
   ~_pi_event();
@@ -665,10 +671,16 @@ struct _pi_event {
   _pi_event(pi_command_type type, pi_context context, pi_queue queue,
             CUstream stream, pi_uint32 stream_token);
 
+  // This constructor is private to force programmers to use the
+  // make_with_native for event introp
+  _pi_event(pi_context context, CUevent eventNative);
+
   pi_command_type commandType_; // The type of command associated with event.
 
   std::atomic_uint32_t refCount_; // Event reference count.
 
+  bool has_ownership_; // Signifies if event owns the native type.
+
   bool hasBeenWaitedOn_; // Signifies whether the event has been waited
                          // on through a call to wait(), which implies
                          // that it has completed.

sycl/test/basic_tests/interop-cuda-experimental.cpp

@@ -28,6 +28,7 @@ int main() {
   device Device;
   context Context(Device);
   queue Queue(Device);
+  event Event;
 
   // 4.5.1.1 For each SYCL runtime class T which supports SYCL application
   // interoperability with the SYCL backend, a specialization of return_type
@@ -40,6 +41,7 @@ int main() {
 
   backend_traits<backend::ext_oneapi_cuda>::return_type<device> cu_device;
   backend_traits<backend::ext_oneapi_cuda>::return_type<context> cu_context;
+  backend_traits<backend::ext_oneapi_cuda>::return_type<event> cu_event;
   backend_traits<backend::ext_oneapi_cuda>::return_type<queue> cu_queue;
 
   // 4.5.1.2 For each SYCL runtime class T which supports SYCL application
@@ -51,6 +53,7 @@ int main() {
 
   cu_device = get_native<backend::ext_oneapi_cuda>(Device);
   cu_context = get_native<backend::ext_oneapi_cuda>(Context);
+  cu_event = get_native<backend::ext_oneapi_cuda>(Event);
   cu_queue = get_native<backend::ext_oneapi_cuda>(Queue);
 
   // Check deprecated
@@ -62,6 +65,9 @@ int main() {
   cu_context = Context.get_native<backend::ext_oneapi_cuda>();
   // expected-warning@+2 {{'get_native' is deprecated: Use SYCL 2020 sycl::get_native free function}}
   // expected-warning@+1 {{'get_native<sycl::backend::ext_oneapi_cuda>' is deprecated: Use SYCL 2020 sycl::get_native free function}}
+  cu_event = Event.get_native<backend::ext_oneapi_cuda>();
+  // expected-warning@+2 {{'get_native' is deprecated: Use SYCL 2020 sycl::get_native free function}}
+  // expected-warning@+1 {{'get_native<sycl::backend::ext_oneapi_cuda>' is deprecated: Use SYCL 2020 sycl::get_native free function}}
   cu_queue = Queue.get_native<backend::ext_oneapi_cuda>();
 
   // 4.5.1.1 For each SYCL runtime class T which supports SYCL application
@@ -89,6 +95,7 @@ int main() {
       cu_context[0]};
   context InteropContext =
       make_context<backend::ext_oneapi_cuda>(InteropContextInput);
+  event InteropEvent = make_event<backend::ext_oneapi_cuda>(cu_event, Context);
 
   queue InteropQueue = make_queue<backend::ext_oneapi_cuda>(cu_queue, Context);