[SYCL][HIP] add missing get_native free function, updates to use SYCL 2020 …

sycl/include/sycl/backend.hpp

@@ -25,8 +25,12 @@
 #endif
 #endif
 #if SYCL_EXT_ONEAPI_BACKEND_HIP
+#ifdef SYCL_EXT_ONEAPI_BACKEND_HIP_EXPERIMENTAL
+#include <sycl/ext/oneapi/experimental/backend/backend_traits_hip.hpp>
+#else
 #include <sycl/detail/backend_traits_hip.hpp>
 #endif
+#endif
 #if SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO
 #include <sycl/detail/backend_traits_level_zero.hpp>
 #endif
@@ -110,10 +114,8 @@ auto get_native_buffer(const buffer<DataT, Dimensions, AllocatorT, void> &Obj)
   // No check for backend mismatch because buffer can be allocated on different
   // backends
   if (BackendName == backend::ext_oneapi_level_zero)
-    throw sycl::runtime_error(
-        errc::feature_not_supported,
-        "Buffer interop is not supported by level zero yet",
-        PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(sycl::errc::feature_not_supported,
+                          "Buffer interop is not supported by level zero yet");
   return Obj.template getNative<BackendName>();
 }
 #endif
@@ -122,10 +124,8 @@ auto get_native_buffer(const buffer<DataT, Dimensions, AllocatorT, void> &Obj)
 template <backend BackendName, class SyclObjectT>
 auto get_native(const SyclObjectT &Obj)
     -> backend_return_t<BackendName, SyclObjectT> {
-  // TODO use SYCL 2020 exception when implemented
   if (Obj.get_backend() != BackendName) {
-    throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
-                              PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(sycl::errc::backend_mismatch, "Backends mismatch");
   }
   return reinterpret_cast<backend_return_t<BackendName, SyclObjectT>>(
       Obj.getNative());
@@ -134,10 +134,8 @@ auto get_native(const SyclObjectT &Obj)
 template <backend BackendName, bundle_state State>
 auto get_native(const kernel_bundle<State> &Obj)
     -> backend_return_t<BackendName, kernel_bundle<State>> {
-  // TODO use SYCL 2020 exception when implemented
   if (Obj.get_backend() != BackendName) {
-    throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
-                              PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(sycl::errc::backend_mismatch, "Backends mismatch");
   }
   return Obj.template getNative<BackendName>();
 }
@@ -153,10 +151,8 @@ auto get_native(const buffer<DataT, Dimensions, AllocatorT> &Obj)
 template <>
 inline backend_return_t<backend::opencl, event>
 get_native<backend::opencl, event>(const event &Obj) {
-  // TODO use SYCL 2020 exception when implemented
   if (Obj.get_backend() != backend::opencl) {
-    throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
-                              PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(sycl::errc::backend_mismatch, "Backends mismatch");
   }
   backend_return_t<backend::opencl, event> ReturnValue;
   for (auto const &element : Obj.getNativeVector()) {
@@ -173,10 +169,8 @@ get_native<backend::opencl, event>(const event &Obj) {
 template <>
 inline backend_return_t<backend::ext_oneapi_cuda, device>
 get_native<backend::ext_oneapi_cuda, device>(const device &Obj) {
-  // TODO use SYCL 2020 exception when implemented
   if (Obj.get_backend() != backend::ext_oneapi_cuda) {
-    throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
-                              PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(sycl::errc::backend_mismatch, "Backends mismatch");
   }
   // CUDA uses a 32-bit int instead of an opaque pointer like other backends,
   // so we need a specialization with static_cast instead of reinterpret_cast.
@@ -185,6 +179,20 @@ get_native<backend::ext_oneapi_cuda, device>(const device &Obj) {
 }
 #endif
 
+#if SYCL_EXT_ONEAPI_BACKEND_HIP
+template <>
+inline backend_return_t<backend::ext_oneapi_hip, device>
+get_native<backend::ext_oneapi_hip, device>(const device &Obj) {
+  if (Obj.get_backend() != backend::ext_oneapi_hip) {
+    throw sycl::exception(sycl::errc::backend_mismatch, "Backends mismatch");
+  }
+  // HIP uses a 32-bit int instead of an opaque pointer like other backends,
+  // so we need a specialization with static_cast instead of reinterpret_cast.
+  return static_cast<backend_return_t<backend::ext_oneapi_hip, device>>(
+      Obj.getNative());
+}
+#endif
+
 // Native handle of an accessor should be accessed through interop_handler
 template <backend BackendName, typename DataT, int Dimensions,
           access::mode AccessMode, access::target AccessTarget,

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

@@ -0,0 +1,127 @@
+//===------- backend_traits_hip.hpp - Backend traits for HIP ---*-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 HIP backend but there is no sycl::detail::InteropFeatureSupportMap
+// specialization for the HIP backend.
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/accessor.hpp>
+#include <sycl/context.hpp>
+#include <sycl/detail/backend_traits.hpp>
+#include <sycl/device.hpp>
+#include <sycl/event.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/queue.hpp>
+
+#include <vector>
+
+typedef int HIPdevice;
+typedef struct HIPctx_st *HIPcontext;
+typedef struct HIPstream_st *HIPstream;
+typedef struct HIPevent_st *HIPevent;
+typedef struct HIPmod_st *HIPmodule;
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+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_hip, context> {
+  using type = HIPcontext;
+};
+
+template <> struct interop<backend::ext_oneapi_hip, device> {
+  using type = HIPdevice;
+};
+
+template <> struct interop<backend::ext_oneapi_hip, event> {
+  using type = HIPevent;
+};
+
+template <> struct interop<backend::ext_oneapi_hip, queue> {
+  using type = HIPstream;
+};
+
+template <> struct interop<backend::ext_oneapi_hip, platform> {
+  using type = std::vector<HIPdevice>;
+};
+
+template <typename DataT, int Dimensions, typename AllocatorT>
+struct BackendInput<backend::ext_oneapi_hip,
+                    buffer<DataT, Dimensions, AllocatorT>> {
+  using type = DataT *;
+};
+
+template <typename DataT, int Dimensions, typename AllocatorT>
+struct BackendReturn<backend::ext_oneapi_hip,
+                     buffer<DataT, Dimensions, AllocatorT>> {
+  using type = DataT *;
+};
+
+template <> struct BackendInput<backend::ext_oneapi_hip, context> {
+  using type = HIPcontext;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_hip, context> {
+  using type = std::vector<HIPcontext>;
+};
+
+template <> struct BackendInput<backend::ext_oneapi_hip, device> {
+  using type = HIPdevice;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_hip, device> {
+  using type = HIPdevice;
+};
+
+template <> struct BackendInput<backend::ext_oneapi_hip, event> {
+  using type = HIPevent;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_hip, event> {
+  using type = HIPevent;
+};
+
+template <> struct BackendInput<backend::ext_oneapi_hip, queue> {
+  using type = HIPstream;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_hip, queue> {
+  using type = HIPstream;
+};
+
+template <> struct BackendInput<backend::ext_oneapi_hip, platform> {
+  using type = std::vector<HIPdevice>;
+};
+
+template <> struct BackendReturn<backend::ext_oneapi_hip, platform> {
+  using type = std::vector<HIPdevice>;
+};
+
+template <> struct InteropFeatureSupportMap<backend::ext_oneapi_hip> {
+  static constexpr bool MakePlatform = false;
+  static constexpr bool MakeDevice = true;
+  static constexpr bool MakeContext = true;
+  static constexpr bool MakeQueue = true;
+  static constexpr bool MakeEvent = true;
+  static constexpr bool MakeBuffer = false;
+  static constexpr bool MakeKernel = false;
+  static constexpr bool MakeKernelBundle = false;
+};
+
+} // namespace detail
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl

sycl/include/sycl/ext/oneapi/experimental/backend/hip.cpp

@@ -0,0 +1,97 @@
+//==--------- hip.hpp - SYCL HIP 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 <sycl/backend.hpp>
+#include <sycl/context.hpp>
+#include <sycl/ext/oneapi/experimental/backend/backend_traits_hip.hpp>
+
+#include <vector>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace ext {
+namespace oneapi {
+namespace hip {
+
+// Implementation of ext_oneapi_hip::make<device>
+inline __SYCL_EXPORT device make_device(pi_native_handle NativeHandle) {
+  return sycl::detail::make_device(NativeHandle, backend::ext_oneapi_hip);
+}
+
+// Implementation of hip::has_native_event
+inline __SYCL_EXPORT bool has_native_event(event sycl_event) {
+  if (sycl_event.get_backend() == backend::ext_oneapi_hip)
+    return get_native<backend::ext_oneapi_hip>(sycl_event) != nullptr;
+
+  return false;
+}
+
+} // namespace hip
+} // namespace oneapi
+} // namespace ext
+
+// HIP context specialization
+template <>
+inline auto get_native<backend::ext_oneapi_hip, context>(const context &C)
+    -> backend_return_t<backend::ext_oneapi_hip, context> {
+  // create a vector to be returned
+  backend_return_t<backend::ext_oneapi_hip, context> ret;
+
+  // get the native HIP context from the SYCL object
+  auto native = reinterpret_cast<
+      backend_return_t<backend::ext_oneapi_hip, context>::value_type>(
+      C.getNative());
+  ret.push_back(native);
+
+  return ret;
+}
+
+// Specialisation of interop_handles get_native_context
+template <>
+inline backend_return_t<backend::ext_oneapi_hip, context>
+interop_handle::get_native_context<backend::ext_oneapi_hip>() 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
+}
+
+// HIP device specialization
+template <>
+inline device make_device<backend::ext_oneapi_hip>(
+    const backend_input_t<backend::ext_oneapi_hip, device> &BackendObject) {
+  pi_native_handle NativeHandle = static_cast<pi_native_handle>(BackendObject);
+  return ext::oneapi::hip::make_device(NativeHandle);
+}
+
+// HIP event specialization
+template <>
+inline event make_event<backend::ext_oneapi_hip>(
+    const backend_input_t<backend::ext_oneapi_hip, event> &BackendObject,
+    const context &TargetContext) {
+  return detail::make_event(detail::pi::cast<pi_native_handle>(BackendObject),
+                            TargetContext, true,
+                            /*Backend*/ backend::ext_oneapi_hip);
+}
+
+// HIP queue specialization
+template <>
+inline queue make_queue<backend::ext_oneapi_hip>(
+    const backend_input_t<backend::ext_oneapi_hip, queue> &BackendObject,
+    const context &TargetContext, const async_handler Handler) {
+  return detail::make_queue(detail::pi::cast<pi_native_handle>(BackendObject),
+                            TargetContext, nullptr, true, Handler,
+                            /*Backend*/ backend::ext_oneapi_hip);
+}
+
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl

sycl/plugins/hip/pi_hip.hpp

@@ -77,15 +77,15 @@ struct _pi_device {
 private:
   using native_type = hipDevice_t;
 
-  native_type cuDevice_;
+  native_type hipDevice_;
   std::atomic_uint32_t refCount_;
   pi_platform platform_;
 
 public:
-  _pi_device(native_type cuDevice, pi_platform platform)
-      : cuDevice_(cuDevice), refCount_{1}, platform_(platform) {}
+  _pi_device(native_type hipDevice, pi_platform platform)
+      : hipDevice_(hipDevice), refCount_{1}, platform_(platform) {}
 
-  native_type get() const noexcept { return cuDevice_; };
+  native_type get() const noexcept { return hipDevice_; };
 
   pi_uint32 get_reference_count() const noexcept { return refCount_; }