[SYCL] [CUDA]Add experimental cuda interop

sycl/include/CL/sycl/accessor.hpp

@@ -2218,6 +2218,12 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
   bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
 };
 
+// local_accessor
+// Same semantics and restrictions as accessor with target::local
+template <typename DataT, int Dimensions = 1>
+using local_accessor =
+    accessor<DataT, Dimensions, access_mode::read_write, access::target::local>;
+
 /// Image accessors.
 ///
 /// Available only when accessTarget == access::target::image.

sycl/include/CL/sycl/backend.hpp

@@ -18,8 +18,12 @@
 #include <CL/sycl/detail/backend_traits_opencl.hpp>
 #endif
 #if SYCL_EXT_ONEAPI_BACKEND_CUDA
+#ifdef SYCL_EXT_ONEAPI_BACKEND_CUDA_EXPERIMENTAL
+#include <sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp>
+#else
 #include <CL/sycl/detail/backend_traits_cuda.hpp>
 #endif
+#endif
 #if SYCL_EXT_ONEAPI_BACKEND_HIP
 #include <CL/sycl/detail/backend_traits_hip.hpp>
 #endif
@@ -187,15 +191,25 @@ inline backend_return_t<backend::opencl, event> get_native<
 }
 #endif
 
-// Native handle of an accessor should be accessed through interop_handler
+// Native accessor handle for kernel function interop
 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 detail::interop<
-        BackendName, accessor<DataT, Dimensions, AccessMode, AccessTarget,
-                              IsPlaceholder>>::type = delete;
+    typename detail::interop<BackendName,
+                             accessor<DataT, Dimensions, AccessMode,
+                                      AccessTarget, IsPlaceholder>>::type {
+#ifdef __SYCL_DEVICE_ONLY__
+  return reinterpret_cast<typename detail::interop<
+      BackendName, accessor<DataT, Dimensions, AccessMode, AccessTarget,
+                            IsPlaceholder>>::type>(Obj.get_pointer().get());
+
+#else
+  throw runtime_error("Get native accessor is not support on host.",
+                      PI_INVALID_VALUE);
+#endif
+}
 
 namespace detail {
 // Forward declaration

sycl/include/CL/sycl/context.hpp

@@ -27,6 +27,8 @@ class platform;
 namespace detail {
 class context_impl;
 }
+template <backend Backend, class SyclT>
+auto get_native(const SyclT &Obj) -> backend_return_t<Backend, SyclT>;
 
 /// The context class represents a SYCL context on which kernel functions may
 /// be executed.
@@ -230,6 +232,10 @@ class __SYCL_EXPORT context {
   pi_native_handle getNative() const;
 
   std::shared_ptr<detail::context_impl> impl;
+
+  template <backend Backend, class SyclT>
+  friend auto get_native(const SyclT &Obj) -> backend_return_t<Backend, SyclT>;
+
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
 

sycl/include/CL/sycl/platform.hpp

@@ -27,6 +27,8 @@ class device;
 namespace detail {
 class platform_impl;
 }
+template <backend Backend, class SyclT>
+auto get_native(const SyclT &Obj) -> backend_return_t<Backend, SyclT>;
 
 /// Encapsulates a SYCL platform on which kernels may be executed.
 ///
@@ -147,6 +149,9 @@ class __SYCL_EXPORT platform {
   std::shared_ptr<detail::platform_impl> impl;
   platform(std::shared_ptr<detail::platform_impl> impl) : impl(impl) {}
 
+  template <backend Backend, class SyclT>
+  friend auto get_native(const SyclT &Obj) -> backend_return_t<Backend, SyclT>;
+
   template <class T>
   friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
   template <class Obj>

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

@@ -0,0 +1,165 @@
+//===------- backend_traits_cuda.hpp - Backend traits for CUDA ---*-C++ -*-===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines the specializations of the sycl::detail::interop,
+// sycl::detail::BackendInput and sycl::detail::BackendReturn class templates
+// for the CUDA backend but there is no sycl::detail::InteropFeatureSupportMap
+// specialization for the CUDA backend.
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/accessor.hpp>
+#include <CL/sycl/context.hpp>
+#include <CL/sycl/detail/backend_traits.hpp>
+#include <CL/sycl/device.hpp>
+#include <CL/sycl/event.hpp>
+#include <CL/sycl/kernel_bundle.hpp>
+#include <CL/sycl/queue.hpp>
+
+#include <vector>
+
+typedef int CUdevice;
+typedef struct CUctx_st *CUcontext;
+typedef struct CUstream_st *CUstream;
+typedef struct CUevent_st *CUevent;
+typedef struct CUmod_st *CUmodule;
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+// TODO the interops for context, device, event, platform and program
+// may be removed after removing the deprecated 'get_native()' methods
+// from the corresponding classes. The interop<backend, queue> specialization
+// is also used in the get_queue() method of the deprecated class
+// interop_handler and also can be removed after API cleanup.
+template <> struct interop<backend::ext_oneapi_cuda, context> {
+  using type = CUcontext;
+};
+
+template <> struct interop<backend::ext_oneapi_cuda, device> {
+  using type = CUdevice;
+};
+
+template <> struct interop<backend::ext_oneapi_cuda, event> {
+  using type = CUevent;
+};
+
+template <> struct interop<backend::ext_oneapi_cuda, queue> {
+  using type = CUstream;
+};
+
+template <> struct interop<backend::ext_oneapi_cuda, platform> {
+  using type = std::vector<CUdevice>;
+};
+
+#ifdef __SYCL_INTERNAL_API
+template <> struct interop<backend::ext_oneapi_cuda, program> {
+  using type = CUmodule;
+};
+#endif
+
+template <typename DataT, int Dimensions, access::mode AccessMode>
+struct interop<backend::ext_oneapi_cuda,
+               accessor<DataT, Dimensions, AccessMode, access::target::device,
+                        access::placeholder::false_t>> {
+  using type = DataT *;
+};
+
+template <typename DataT, int Dimensions, access::mode AccessMode>
+struct interop<
+    backend::ext_oneapi_cuda,
+    accessor<DataT, Dimensions, AccessMode, access::target::constant_buffer,
+             access::placeholder::false_t>> {
+  using type = DataT *;
+};
+
+template <typename DataT, int Dimensions, access::mode AccessMode>
+struct interop<backend::ext_oneapi_cuda,
+               accessor<DataT, Dimensions, AccessMode, access::target::local,
+                        access::placeholder::false_t>> {
+  using type = DataT *;
+};
+
+template <typename DataT, int Dimensions, typename AllocatorT>
+struct BackendInput<backend::ext_oneapi_cuda,
+                    buffer<DataT, Dimensions, AllocatorT>> {
+  using type = DataT *;
+};
+
+template <typename DataT, int Dimensions, typename AllocatorT>
+struct BackendReturn<backend::ext_oneapi_cuda,
+                     buffer<DataT, Dimensions, AllocatorT>> {
+  using type = DataT *;
+};
+
+template <> struct BackendInput<backend::ext_oneapi_cuda, context> {
+  using type = CUcontext;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_cuda, context> {
+  using type = std::vector<CUcontext>;
+};
+
+template <> struct BackendInput<backend::ext_oneapi_cuda, device> {
+  using type = CUdevice;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_cuda, device> {
+  using type = CUdevice;
+};
+
+template <> struct BackendInput<backend::ext_oneapi_cuda, event> {
+  using type = CUevent;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_cuda, event> {
+  using type = CUevent;
+};
+
+template <> struct BackendInput<backend::ext_oneapi_cuda, queue> {
+  using type = CUstream;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_cuda, queue> {
+  using type = CUstream;
+};
+
+template <> struct BackendInput<backend::ext_oneapi_cuda, platform> {
+  using type = std::vector<CUdevice>;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_cuda, platform> {
+  using type = std::vector<CUdevice>;
+};
+
+#ifdef __SYCL_INTERNAL_API
+template <> struct BackendInput<backend::ext_oneapi_cuda, program> {
+  using type = CUmodule;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_cuda, program> {
+  using type = CUmodule;
+};
+#endif
+
+template <> struct InteropFeatureSupportMap<backend::ext_oneapi_cuda> {
+  static constexpr bool MakePlatform = true;
+  static constexpr bool MakeDevice = true;
+  static constexpr bool MakeContext = true;
+  static constexpr bool MakeQueue = true;
+  static constexpr bool MakeEvent = true;
+  static constexpr bool MakeBuffer = true;
+  static constexpr bool MakeKernel = true;
+  static constexpr bool MakeKernelBundle = true;
+};
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

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

@@ -0,0 +1,120 @@
+//==--------- cuda.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
+
+#include <CL/sycl/backend.hpp>
+#include <CL/sycl/context.hpp>
+#include <CL/sycl/interop_handle.hpp>
+
+#include <vector>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace oneapi {
+namespace cuda {
+
+// Implementation of cuda::make<device>
+inline __SYCL_EXPORT device make_device(pi_native_handle NativeHandle) {
+  return sycl::detail::make_device(NativeHandle, backend::cuda);
+}
+
+// Implementation of cuda::make<platform>
+inline __SYCL_EXPORT platform make_platform(pi_native_handle NativeHandle) {
+  return sycl::detail::make_platform(NativeHandle, backend::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
+
+// CUDA context specialization
+template <>
+inline auto get_native<backend::ext_oneapi_cuda, context>(const context &C)
+    -> backend_return_t<backend::ext_oneapi_cuda, context> {
+  // create a vector to be returned
+  backend_return_t<backend::ext_oneapi_cuda, context> ret;
+
+  // get the native CUDA context from the SYCL object
+  auto native = reinterpret_cast<
+      backend_return_t<backend::ext_oneapi_cuda, context>::value_type>(
+      C.getNative());
+  ret.push_back(native);
+
+  return ret;
+}
+
+// Specialisation of non-free context get_native
+template <>
+inline backend_return_t<backend::ext_oneapi_cuda, context>
+context::get_native<backend::ext_oneapi_cuda>() const {
+  return sycl::get_native<backend::ext_oneapi_cuda, context>(*this);
+}
+
+// CUDA device specialization
+template <>
+inline device make_device<backend::ext_oneapi_cuda>(
+    const backend_input_t<backend::ext_oneapi_cuda, device> &BackendObject) {
+  pi_native_handle NativeHandle = static_cast<pi_native_handle>(BackendObject);
+  return ext::oneapi::cuda::make_device(NativeHandle);
+}
+
+// CUDA platform specialization
+template <>
+inline auto get_native<backend::ext_oneapi_cuda, platform>(const platform &C)
+    -> backend_return_t<backend::ext_oneapi_cuda, platform> {
+  // get list of platform devices, and transfer to native platform type
+  std::vector<device> platform_devices = C.get_devices();
+  std::vector<CUdevice> native_devices(platform_devices.size());
+
+  // Get the native CUdevice type for each device in platform
+  for (unsigned int i = 0; i < platform_devices.size(); ++i)
+    native_devices[i] =
+        get_native<backend::ext_oneapi_cuda>(platform_devices[i]);
+
+  return native_devices;
+}
+
+// Specialisation of non-free platform get_native
+template <>
+inline backend_return_t<backend::ext_oneapi_cuda, platform>
+platform::get_native<backend::ext_oneapi_cuda>() const {
+  return sycl::get_native<backend::ext_oneapi_cuda, platform>(*this);
+}
+
+template <>
+inline platform make_platform<backend::ext_oneapi_cuda>(
+    const backend_input_t<backend::ext_oneapi_cuda, platform> &BackendObject) {
+  pi_native_handle NativeHandle =
+      detail::pi::cast<pi_native_handle>(&BackendObject);
+  return ext::oneapi::cuda::make_platform(NativeHandle);
+}
+
+// Specialisation of interop_handles get_native_context
+template <>
+inline backend_return_t<backend::ext_oneapi_cuda, context>
+interop_handle::get_native_context<backend::ext_oneapi_cuda>() const {
+#ifndef __SYCL_DEVICE_ONLY__
+  return std::vector{reinterpret_cast<CUcontext>(getNativeContext())};
+#else
+  // we believe this won't be ever called on device side
+  return {};
+#endif
+}
+
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)