[SYCL][PI][CUDA] Implements get_native interoperability

Implements get_native for CUDA allowing queries for native handles on
SYCL objects; queue, event, context, and device.

Signed-off-by: Steffen Larsen <[email protected]>

Author: Steffen Larsen

sycl/include/CL/sycl.hpp

@@ -10,6 +10,7 @@
 
 #include <CL/sycl/accessor.hpp>
 #include <CL/sycl/atomic.hpp>
+#include <CL/sycl/backend.hpp>
 #include <CL/sycl/buffer.hpp>
 #include <CL/sycl/builtins.hpp>
 #include <CL/sycl/context.hpp>

sycl/include/CL/sycl/backend.hpp

@@ -0,0 +1,34 @@
+//==---------------- backend.hpp - SYCL PI backends ------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/accessor.hpp>
+#include <CL/sycl/backend_types.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+
+template <backend BackendName, class SyclObjectT>
+auto get_native(const SyclObjectT &Obj) ->
+    typename interop<BackendName, SyclObjectT>::type {
+  return Obj.template get_native<BackendName>();
+}
+
+// Native handle of an accessor should be accessed through interop_handler
+template <backend BackendName, typename DataT, int Dimensions,
+          access::mode AccessMode, access::target AccessTarget,
+          access::placeholder IsPlaceholder>
+auto get_native(const accessor<DataT, Dimensions, AccessMode, AccessTarget,
+                               IsPlaceholder> &Obj) ->
+    typename interop<BackendName, accessor<DataT, Dimensions, AccessMode,
+                                           AccessTarget, IsPlaceholder>>::type =
+    delete;
+
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/backend/cuda.hpp

@@ -1,3 +1,4 @@
+
 //==---------------- cuda.hpp - SYCL CUDA backend --------------------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
@@ -6,29 +7,45 @@
 //
 //===----------------------------------------------------------------------===//
 
+#pragma once
+
+#include <CL/sycl/accessor.hpp>
+#include <CL/sycl/backend_types.hpp>
+#include <CL/sycl/context.hpp>
 #include <CL/sycl/detail/defines.hpp>
+#include <CL/sycl/device.hpp>
+#include <CL/sycl/event.hpp>
+#include <CL/sycl/queue.hpp>
+
+typedef int CUdevice;
+typedef struct CUctx_st *CUcontext;
+typedef struct CUstream_st *CUstream;
+typedef struct CUevent_st *CUevent;
+
+// As defined in the CUDA 10.1 header file. This requires CUDA version > 3.2
+#if defined(_WIN64) || defined(__LP64__)
+typedef unsigned long long CUdeviceptr;
+#else
+typedef unsigned int CUdeviceptr;
+#endif
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-namespace backend {
-namespace cuda {
-
-// CUDA backend specific options
-// TODO: Use values that won't overlap with others
-
-// Mem Object info: Retrieve the raw CUDA pointer from a cl_mem
-#define PI_CUDA_RAW_POINTER (0xFF01)
-// Context creation: Use a primary CUDA context instead of a custom one by
-//                   providing a property value of PI_TRUE for the following
-//                   property ID.
-#define PI_CONTEXT_PROPERTIES_CUDA_PRIMARY (0xFF02)
-
-// PI Command Queue using Default stream
-#define PI_CUDA_USE_DEFAULT_STREAM (0xFF03)
-// PI Command queue will sync with default stream
-#define PI_CUDA_SYNC_WITH_DEFAULT (0xFF04)
-
-} // namespace cuda
-} // namespace backend
+
+template <> struct interop<backend::cuda, device> { using type = CUdevice; };
+
+template <> struct interop<backend::cuda, context> { using type = CUcontext; };
+
+template <> struct interop<backend::cuda, queue> { using type = CUstream; };
+
+template <> struct interop<backend::cuda, event> { using type = CUevent; };
+
+template <typename DataT, int Dimensions, access::mode AccessMode>
+struct interop<backend::cuda, accessor<DataT, Dimensions, AccessMode,
+                                       access::target::global_buffer,
+                                       access::placeholder::false_t>> {
+  using type = CUdeviceptr;
+};
+
 } // namespace sycl
 } // namespace cl

sycl/include/CL/sycl/backend_types.hpp

@@ -0,0 +1,19 @@
+//==-------------- backend_types.hpp - SYCL backend types ------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+namespace cl {
+namespace sycl {
+
+enum class backend { host, opencl, cuda };
+
+template <backend name, typename SYCLObjectT> struct interop;
+
+} // namespace sycl
+} // namespace cl

sycl/include/CL/sycl/context.hpp

@@ -7,6 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 #pragma once
+
+#include <CL/sycl/backend_types.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/export.hpp>
 #include <CL/sycl/exception_list.hpp>
@@ -135,10 +137,21 @@ class __SYCL_EXPORT context {
   /// \return a vector of valid SYCL device instances.
   vector_class<device> get_devices() const;
 
+  /// Gets the native handle of the SYCL context.
+  ///
+  /// \return a native handle, the type of which defined by the backend.
+  template <backend BackendName>
+  auto get_native() const -> typename interop<BackendName, context>::type {
+    return reinterpret_cast<typename interop<BackendName, context>::type>(
+        getNative());
+  }
+
 private:
   /// Constructs a SYCL context object from a valid context_impl instance.
   context(shared_ptr_class<detail::context_impl> Impl);
 
+  pi_native_handle getNative() const;
+
   shared_ptr_class<detail::context_impl> impl;
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);

sycl/include/CL/sycl/detail/buffer_impl.hpp

@@ -44,6 +44,8 @@ class __SYCL_EXPORT buffer_impl final : public SYCLMemObjT {
   using typename BaseT::MemObjType;
 
 public:
+  static constexpr pi_handle_type PIHandleType = PI_NATIVE_HANDLE_MEM;
+
   buffer_impl(size_t SizeInBytes, size_t RequiredAlign,
               const property_list &Props,
               unique_ptr_class<SYCLMemObjAllocator> Allocator)

sycl/include/CL/sycl/detail/cuda_definitions.hpp

@@ -0,0 +1,24 @@
+//==------------ cuda_definitions.hpp - SYCL CUDA backend ------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+// CUDA backend specific options
+// TODO: Use values that won't overlap with others
+
+// Mem Object info: Retrieve the raw CUDA pointer from a cl_mem
+#define PI_CUDA_RAW_POINTER (0xFF01)
+// Context creation: Use a primary CUDA context instead of a custom one by
+//                   providing a property value of PI_TRUE for the following
+//                   property ID.
+#define PI_CONTEXT_PROPERTIES_CUDA_PRIMARY (0xFF02)
+
+// PI Command Queue using Default stream
+#define PI_CUDA_USE_DEFAULT_STREAM (0xFF03)
+// PI Command queue will sync with default stream
+#define PI_CUDA_SYNC_WITH_DEFAULT (0xFF04)

sycl/include/CL/sycl/detail/pi.def

@@ -113,4 +113,7 @@ _PI_API(piextUSMGetMemAllocInfo)
 
 _PI_API(piextKernelSetArgMemObj)
 
+// Interop
+_PI_API(piGetNativeHandle)
+
 #undef _PI_API

sycl/include/CL/sycl/detail/pi.h

@@ -59,6 +59,16 @@ using pi_uint32 = uint32_t;
 using pi_uint64 = uint64_t;
 using pi_bool = pi_uint32;
 using pi_bitfield = pi_uint64;
+using pi_native_handle = uintptr_t;
+
+enum pi_handle_type {
+  PI_NATIVE_HANDLE_CONTEXT,
+  PI_NATIVE_HANDLE_DEVICE,
+  PI_NATIVE_HANDLE_QUEUE,
+  PI_NATIVE_HANDLE_EVENT,
+  PI_NATIVE_HANDLE_MEM,
+  PI_NATIVE_HANDLE_PLATFORM
+};
 
 //
 // NOTE: prefer to map 1:1 to OpenCL so that no translation is needed
@@ -1346,6 +1356,14 @@ pi_result piextUSMGetMemAllocInfo(pi_context context, const void *ptr,
                                   size_t param_value_size, void *param_value,
                                   size_t *param_value_size_ret);
 
+/// Gets the native handle of a SYCL object.
+///
+/// \param handleType is a representation of the native handle.
+/// \param piObject is the PI object to get the native handle of.
+/// \param nativeHandle is the native handle of piObject.
+pi_result piGetNativeHandle(pi_handle_type handleType, void *piObject,
+                            pi_native_handle *nativeHandle);
+
 struct _pi_plugin {
   // PI version supported by host passed to the plugin. The Plugin
   // checks and writes the appropriate Function Pointers in

sycl/include/CL/sycl/device.hpp

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <CL/sycl/backend_types.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/export.hpp>
 #include <CL/sycl/info/info_desc.hpp>
@@ -165,10 +166,21 @@ class __SYCL_EXPORT device {
   static vector_class<device>
   get_devices(info::device_type deviceType = info::device_type::all);
 
+  /// Gets the native handle of the SYCL device.
+  ///
+  /// \return a native handle, the type of which defined by the backend.
+  template <backend BackendName>
+  auto get_native() const -> typename interop<BackendName, device>::type {
+    return static_cast<typename interop<BackendName, device>::type>(
+        getNative());
+  }
+
 private:
   shared_ptr_class<detail::device_impl> impl;
   device(shared_ptr_class<detail::device_impl> impl) : impl(impl) {}
 
+  pi_native_handle getNative() const;
+
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
 

sycl/include/CL/sycl/event.hpp

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <CL/sycl/backend_types.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/export.hpp>
 #include <CL/sycl/info/info_desc.hpp>
@@ -113,9 +114,19 @@ class __SYCL_EXPORT event {
   typename info::param_traits<info::event_profiling, param>::return_type
   get_profiling_info() const;
 
+  /// Gets the native handle of the SYCL event.
+  ///
+  /// \return a native handle, the type of which defined by the backend.
+  template <backend BackendName>
+  auto get_native() const -> typename interop<BackendName, event>::type {
+    return static_cast<typename interop<BackendName, event>::type>(getNative());
+  }
+
 private:
   event(shared_ptr_class<detail::event_impl> EventImpl);
 
+  pi_native_handle getNative() const;
+
   shared_ptr_class<detail::event_impl> impl;
 
   template <class Obj>

sycl/include/CL/sycl/queue.hpp

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <CL/sycl/backend_types.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/export.hpp>
 #include <CL/sycl/device.hpp>
@@ -638,7 +639,18 @@ class __SYCL_EXPORT queue {
   /// Equivalent to has_property<property::queue::in_order>()
   bool is_in_order() const;
 
+  /// Gets the native handle of the SYCL queue.
+  ///
+  /// \return a native handle, the type of which defined by the backend.
+  template <backend BackendName>
+  auto get_native() const -> typename interop<BackendName, queue>::type {
+    return reinterpret_cast<typename interop<BackendName, queue>::type>(
+        getNative());
+  }
+
 private:
+  pi_native_handle getNative() const;
+
   shared_ptr_class<detail::queue_impl> impl;
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);

sycl/plugins/cuda/pi_cuda.cpp

@@ -11,7 +11,7 @@
 ///
 /// \ingroup sycl_pi_cuda
 
-#include <CL/sycl/backend/cuda.hpp>
+#include <CL/sycl/detail/cuda_definitions.hpp>
 #include <CL/sycl/detail/pi.hpp>
 #include <pi_cuda.hpp>
 
@@ -3554,6 +3554,46 @@ pi_result cuda_piextUSMGetMemAllocInfo(pi_context context, const void *ptr,
   return result;
 }
 
+// Native interop
+
+pi_result cuda_piGetNativeHandle(pi_handle_type handleType, void *piObject,
+                                 pi_native_handle *nativeHandle) {
+  switch (handleType) {
+  case pi_handle_type::PI_NATIVE_HANDLE_CONTEXT: {
+    pi_context context = static_cast<pi_context>(piObject);
+    *nativeHandle = reinterpret_cast<pi_native_handle>(context->get());
+    return PI_SUCCESS;
+  }
+  case pi_handle_type::PI_NATIVE_HANDLE_DEVICE: {
+    pi_device device = static_cast<pi_device>(piObject);
+    *nativeHandle = static_cast<pi_native_handle>(device->get());
+    return PI_SUCCESS;
+  }
+  case pi_handle_type::PI_NATIVE_HANDLE_QUEUE: {
+    pi_queue queue = static_cast<pi_queue>(piObject);
+    *nativeHandle = reinterpret_cast<pi_native_handle>(queue->get());
+    return PI_SUCCESS;
+  }
+  case pi_handle_type::PI_NATIVE_HANDLE_EVENT: {
+    pi_event event = static_cast<pi_event>(piObject);
+    if (event->is_user_event()) {
+      return PI_INVALID_EVENT;
+    }
+    *nativeHandle = reinterpret_cast<pi_native_handle>(event->get());
+    return PI_SUCCESS;
+  }
+  case pi_handle_type::PI_NATIVE_HANDLE_MEM: {
+    pi_mem mem = static_cast<pi_mem>(piObject);
+    *nativeHandle = static_cast<pi_native_handle>(mem->get());
+    return PI_SUCCESS;
+  }
+  default:
+    PI_HANDLE_UNKNOWN_PARAM_NAME(handleType);
+  }
+  cl::sycl::detail::pi::die("Native handle request not implemented");
+  return {};
+}
+
 const char SupportedVersion[] = _PI_H_VERSION_STRING;
 
 pi_result piPluginInit(pi_plugin *PluginInit) {
@@ -3673,6 +3713,8 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piextUSMGetMemAllocInfo, cuda_piextUSMGetMemAllocInfo)
 
   _PI_CL(piextKernelSetArgMemObj, cuda_piextKernelSetArgMemObj)
+  // Interop
+  _PI_CL(piGetNativeHandle, cuda_piGetNativeHandle)
 
 #undef _PI_CL