[SYCL][HIP] Implemented supported make_* interop functions. (#10526)

This PR adds missing functions in the hip backend to allow for
interoperability in programs that create sycl objects from native hip
objects. The new function implementations are:

- `make_device`
- `make_queue`
- `make_event`

Note that it would really make sense for
#10491 to be merged first because this
PR makes the same code change in pi2ur, for a fix that is attributed to
#10491.

---------

Signed-off-by: Jack Kirk <[email protected]>

Author: JackAKirk

sycl/include/sycl/backend.hpp

@@ -298,9 +298,11 @@ std::enable_if_t<detail::InteropFeatureSupportMap<Backend>::MakeQueue == true,
 make_queue(const typename backend_traits<Backend>::template input_type<queue>
                &BackendObject,
            const context &TargetContext, const async_handler Handler = {}) {
+  auto KeepOwnership =
+      Backend == backend::ext_oneapi_cuda || Backend == backend::ext_oneapi_hip;
   return detail::make_queue(detail::pi::cast<pi_native_handle>(BackendObject),
-                            false, TargetContext, nullptr, false, {}, Handler,
-                            Backend);
+                            false, TargetContext, nullptr, KeepOwnership, {},
+                            Handler, Backend);
 }
 
 template <backend Backend>

sycl/include/sycl/detail/backend_traits_hip.hpp

@@ -95,6 +95,18 @@ template <> struct BackendReturn<backend::ext_oneapi_hip, queue> {
   using type = HIPstream;
 };
 
+template <> struct InteropFeatureSupportMap<backend::ext_oneapi_hip> {
+  static constexpr bool MakePlatform = false;
+  static constexpr bool MakeDevice = true;
+  static constexpr bool MakeContext = false;
+  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;
+  static constexpr bool MakeImage = false;
+};
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl

sycl/include/sycl/ext/oneapi/backend/hip.hpp

@@ -16,7 +16,7 @@ inline namespace _V1 {
 template <>
 inline backend_return_t<backend::ext_oneapi_hip, device>
 get_native<backend::ext_oneapi_hip, device>(const device &Obj) {
-  // TODO use SYCL 2020 exception when implemented
+  // TODO swap with SYCL 2020 exception when in ABI-break window
   if (Obj.get_backend() != backend::ext_oneapi_hip) {
     throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
                               PI_ERROR_INVALID_OPERATION);
@@ -27,5 +27,23 @@ get_native<backend::ext_oneapi_hip, device>(const device &Obj) {
       Obj.getNative());
 }
 
+template <>
+inline device make_device<backend::ext_oneapi_hip>(
+    const backend_input_t<backend::ext_oneapi_hip, device> &BackendObject) {
+  auto devs = device::get_devices(info::device_type::gpu);
+  for (auto &dev : devs) {
+    if (dev.get_backend() == backend::ext_oneapi_hip &&
+        BackendObject == get_native<backend::ext_oneapi_hip>(dev)) {
+      return dev;
+    }
+  }
+  // The ext_oneapi_hip platform(s) adds all n available devices where n
+  // is returned from call to `hipGetDeviceCount`.
+  // Hence if this code is reached then the requested device ordinal must
+  // not be visible to the driver.
+  throw sycl::exception(make_error_code(errc::invalid),
+                        "Native device has an invalid ordinal.");
+}
+
 } // namespace _V1
 } // namespace sycl

sycl/plugins/unified_runtime/ur/adapters/cuda/queue.cpp

@@ -242,7 +242,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueCreateWithNativeHandle(
     ur_native_handle_t hNativeQueue, ur_context_handle_t hContext,
     ur_device_handle_t hDevice, const ur_queue_native_properties_t *pProperties,
     ur_queue_handle_t *phQueue) {
-  (void)pProperties;
   (void)hDevice;
 
   unsigned int CuFlags;
@@ -263,13 +262,14 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueCreateWithNativeHandle(
 
   // Create queue and set num_compute_streams to 1, as computeCuStreams has
   // valid stream
-  *phQueue = new ur_queue_handle_t_{std::move(ComputeCuStreams),
-                                    std::move(TransferCuStreams),
-                                    hContext,
-                                    hContext->getDevice(),
-                                    CuFlags,
-                                    Flags,
-                                    /*backend_owns*/ false};
+  *phQueue =
+      new ur_queue_handle_t_{std::move(ComputeCuStreams),
+                             std::move(TransferCuStreams),
+                             hContext,
+                             hContext->getDevice(),
+                             CuFlags,
+                             Flags,
+                             /*backend_owns*/ pProperties->isNativeHandleOwned};
   (*phQueue)->NumComputeStreams = 1;
 
   return Return;

sycl/plugins/unified_runtime/ur/adapters/hip/event.cpp

@@ -15,10 +15,10 @@ ur_event_handle_t_::ur_event_handle_t_(ur_command_t Type,
                                        ur_context_handle_t Context,
                                        ur_queue_handle_t Queue,
                                        hipStream_t Stream, uint32_t StreamToken)
-    : CommandType{Type}, RefCount{1}, HasBeenWaitedOn{false}, IsRecorded{false},
-      IsStarted{false}, StreamToken{StreamToken}, EvEnd{nullptr},
-      EvStart{nullptr}, EvQueued{nullptr}, Queue{Queue}, Stream{Stream},
-      Context{Context} {
+    : CommandType{Type}, RefCount{1}, HasOwnership{true},
+      HasBeenWaitedOn{false}, IsRecorded{false}, IsStarted{false},
+      StreamToken{StreamToken}, EvEnd{nullptr}, EvStart{nullptr},
+      EvQueued{nullptr}, Queue{Queue}, Stream{Stream}, Context{Context} {
 
   bool ProfilingEnabled = Queue->URFlags & UR_QUEUE_FLAG_PROFILING_ENABLE;
 
@@ -36,6 +36,15 @@ ur_event_handle_t_::ur_event_handle_t_(ur_command_t Type,
   urContextRetain(Context);
 }
 
+ur_event_handle_t_::ur_event_handle_t_(ur_context_handle_t Context,
+                                       hipEvent_t EventNative)
+    : CommandType{UR_COMMAND_EVENTS_WAIT}, RefCount{1}, HasOwnership{false},
+      HasBeenWaitedOn{false}, IsRecorded{false}, IsStarted{false},
+      StreamToken{std::numeric_limits<uint32_t>::max()}, EvEnd{EventNative},
+      EvStart{nullptr}, EvQueued{nullptr}, Queue{nullptr}, Context{Context} {
+  urContextRetain(Context);
+}
+
 ur_event_handle_t_::~ur_event_handle_t_() {
   if (Queue != nullptr) {
     urQueueRelease(Queue);
@@ -160,6 +169,9 @@ ur_result_t ur_event_handle_t_::wait() {
 }
 
 ur_result_t ur_event_handle_t_::release() {
+  if (!backendHasOwnership())
+    return UR_RESULT_SUCCESS;
+
   assert(Queue != nullptr);
   UR_CHECK_ERROR(hipEventDestroy(EvEnd));
 
@@ -302,15 +314,18 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventGetNativeHandle(
 }
 
 /// Created a UR event object from a HIP event handle.
-/// TODO: Implement this.
-/// NOTE: The created UR object takes ownership of the native handle.
+/// NOTE: The created UR object doesn't take ownership of the native handle.
 ///
 /// \param[in] hNativeEvent The native handle to create UR event object from.
 /// \param[out] phEvent Set to the UR event object created from native handle.
-///
-/// \return UR_RESULT_ERROR_UNSUPPORTED_FEATURE
 UR_APIEXPORT ur_result_t UR_APICALL urEventCreateWithNativeHandle(
-    ur_native_handle_t, ur_context_handle_t,
-    const ur_event_native_properties_t *, ur_event_handle_t *) {
-  return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
+    ur_native_handle_t hNativeEvent, ur_context_handle_t hContext,
+    const ur_event_native_properties_t *pProperties,
+    ur_event_handle_t *phEvent) {
+  std::ignore = pProperties;
+
+  *phEvent = ur_event_handle_t_::makeWithNative(
+      hContext, reinterpret_cast<hipEvent_t>(hNativeEvent));
+
+  return UR_RESULT_SUCCESS;
 }

sycl/plugins/unified_runtime/ur/adapters/hip/event.hpp

@@ -60,6 +60,8 @@ struct ur_event_handle_t_ {
 
   uint32_t getEventId() const noexcept { return EventId; }
 
+  bool backendHasOwnership() const noexcept { return HasOwnership; }
+
   // Returns the counter time when the associated command(s) were enqueued
   uint64_t getQueuedTime() const;
 
@@ -77,6 +79,11 @@ struct ur_event_handle_t_ {
                                   StreamToken);
   }
 
+  static ur_event_handle_t makeWithNative(ur_context_handle_t context,
+                                          hipEvent_t eventNative) {
+    return new ur_event_handle_t_(context, eventNative);
+  }
+
   ur_result_t release();
 
   ~ur_event_handle_t_();
@@ -88,10 +95,16 @@ struct ur_event_handle_t_ {
                      ur_queue_handle_t Queue, hipStream_t Stream,
                      uint32_t StreamToken);
 
+  // This constructor is private to force programmers to use the
+  // makeWithNative for event interop
+  ur_event_handle_t_(ur_context_handle_t Context, hipEvent_t EventNative);
+
   ur_command_t CommandType; // The type of command associated with event.
 
   std::atomic_uint32_t RefCount; // Event reference count.
 
+  bool HasOwnership; // 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/plugins/unified_runtime/ur/adapters/hip/queue.cpp

@@ -193,6 +193,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueRelease(ur_queue_handle_t hQueue) {
   try {
     std::unique_ptr<ur_queue_handle_t_> QueueImpl(hQueue);
 
+    if (!hQueue->backendHasOwnership())
+      return UR_RESULT_SUCCESS;
+
     ScopedContext Active(hQueue->getContext()->getDevice());
 
     hQueue->forEachStream([](hipStream_t S) {
@@ -252,19 +255,44 @@ urQueueGetNativeHandle(ur_queue_handle_t hQueue, ur_queue_native_desc_t *,
 }
 
 /// Created a UR queue object from a HIP queue handle.
-/// TODO: Implement this.
-/// NOTE: The created UR object takes ownership of the native handle.
+/// NOTE: The created UR object doesn't takes ownership of the native handle.
 ///
 /// \param[in] hNativeQueue The native handle to create UR queue object from.
 /// \param[in] hContext is the UR context of the queue.
 /// \param[out] phQueue Set to the UR queue object created from native handle.
-/// \param pProperties->isNativeHandleOwned tells if SYCL RT should assume the
-/// ownership of
-///        the native handle, if it can.
-///
-/// \return UR_RESULT_ERROR_UNSUPPORTED_FEATURE
 UR_APIEXPORT ur_result_t UR_APICALL urQueueCreateWithNativeHandle(
-    ur_native_handle_t, ur_context_handle_t, ur_device_handle_t,
-    const ur_queue_native_properties_t *, ur_queue_handle_t *) {
-  return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
+    ur_native_handle_t hNativeQueue, ur_context_handle_t hContext,
+    ur_device_handle_t hDevice, const ur_queue_native_properties_t *pProperties,
+    ur_queue_handle_t *phQueue) {
+  (void)hDevice;
+
+  unsigned int HIPFlags;
+  hipStream_t HIPStream = reinterpret_cast<hipStream_t>(hNativeQueue);
+
+  auto Return = UR_CHECK_ERROR(hipStreamGetFlags(HIPStream, &HIPFlags));
+
+  ur_queue_flags_t Flags = 0;
+  if (HIPFlags == hipStreamDefault)
+    Flags = UR_QUEUE_FLAG_USE_DEFAULT_STREAM;
+  else if (HIPFlags == hipStreamNonBlocking)
+    Flags = UR_QUEUE_FLAG_SYNC_WITH_DEFAULT_STREAM;
+  else
+    detail::ur::die("Unknown hip stream");
+
+  std::vector<hipStream_t> ComputeHIPStreams(1, HIPStream);
+  std::vector<hipStream_t> TransferHIPStreams(0);
+
+  // Create queue and set num_compute_streams to 1, as computeHIPStreams has
+  // valid stream
+  *phQueue =
+      new ur_queue_handle_t_{std::move(ComputeHIPStreams),
+                             std::move(TransferHIPStreams),
+                             hContext,
+                             hContext->getDevice(),
+                             HIPFlags,
+                             Flags,
+                             /*backend_owns*/ pProperties->isNativeHandleOwned};
+  (*phQueue)->NumComputeStreams = 1;
+
+  return Return;
 }

sycl/plugins/unified_runtime/ur/adapters/hip/queue.hpp

@@ -49,11 +49,13 @@ struct ur_queue_handle_t_ {
   std::mutex ComputeStreamMutex;
   std::mutex TransferStreamMutex;
   std::mutex BarrierMutex;
+  bool HasOwnership;
 
   ur_queue_handle_t_(std::vector<native_type> &&ComputeStreams,
                      std::vector<native_type> &&TransferStreams,
                      ur_context_handle_t Context, ur_device_handle_t Device,
-                     unsigned int Flags, ur_queue_flags_t URFlags)
+                     unsigned int Flags, ur_queue_flags_t URFlags,
+                     bool BackendOwns = true)
       : ComputeStreams{std::move(ComputeStreams)},
         TransferStreams{std::move(TransferStreams)},
         DelayCompute(this->ComputeStreams.size(), false),
@@ -62,7 +64,7 @@ struct ur_queue_handle_t_ {
         Device{Device}, RefCount{1}, EventCount{0}, ComputeStreamIdx{0},
         TransferStreamIdx{0}, NumComputeStreams{0}, NumTransferStreams{0},
         LastSyncComputeStreams{0}, LastSyncTransferStreams{0}, Flags(Flags),
-        URFlags(URFlags) {
+        URFlags(URFlags), HasOwnership{BackendOwns} {
     urContextRetain(Context);
     urDeviceRetain(Device);
   }
@@ -235,4 +237,6 @@ struct ur_queue_handle_t_ {
   uint32_t getReferenceCount() const noexcept { return RefCount; }
 
   uint32_t getNextEventId() noexcept { return ++EventCount; }
+
+  bool backendHasOwnership() const noexcept { return HasOwnership; }
 };

sycl/source/backend.cpp

@@ -37,6 +37,8 @@ static const PluginPtr &getPlugin(backend Backend) {
     return pi::getPlugin<backend::ext_oneapi_level_zero>();
   case backend::ext_oneapi_cuda:
     return pi::getPlugin<backend::ext_oneapi_cuda>();
+  case backend::ext_oneapi_hip:
+    return pi::getPlugin<backend::ext_oneapi_hip>();
   default:
     throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
                           "getPlugin: Unsupported backend " +

sycl/test/basic_tests/interop-hip.cpp

@@ -56,5 +56,9 @@ int main() {
   hip_event = get_native<backend::ext_oneapi_hip>(Event);
   hip_queue = get_native<backend::ext_oneapi_hip>(Queue);
 
+  device InteropDevice = make_device<backend::ext_oneapi_hip>(hip_device);
+  event InteropEvent = make_event<backend::ext_oneapi_hip>(hip_event, Context);
+  queue InteropQueue = make_queue<backend::ext_oneapi_hip>(hip_queue, Context);
+
   return 0;
 }