diff --git a/sycl/include/sycl/backend.hpp b/sycl/include/sycl/backend.hpp index 445a44f59acda..c47f077b4319a 100644 --- a/sycl/include/sycl/backend.hpp +++ b/sycl/include/sycl/backend.hpp @@ -298,9 +298,11 @@ std::enable_if_t::MakeQueue == true, make_queue(const typename backend_traits::template input_type &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(BackendObject), - false, TargetContext, nullptr, false, {}, Handler, - Backend); + false, TargetContext, nullptr, KeepOwnership, {}, + Handler, Backend); } template diff --git a/sycl/include/sycl/detail/backend_traits_hip.hpp b/sycl/include/sycl/detail/backend_traits_hip.hpp index aa480f8a099d2..eef21a7f413b2 100644 --- a/sycl/include/sycl/detail/backend_traits_hip.hpp +++ b/sycl/include/sycl/detail/backend_traits_hip.hpp @@ -95,6 +95,18 @@ template <> struct BackendReturn { using type = HIPstream; }; +template <> struct InteropFeatureSupportMap { + 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 diff --git a/sycl/include/sycl/ext/oneapi/backend/hip.hpp b/sycl/include/sycl/ext/oneapi/backend/hip.hpp index 0ffd01337a65c..20526932fcbd6 100644 --- a/sycl/include/sycl/ext/oneapi/backend/hip.hpp +++ b/sycl/include/sycl/ext/oneapi/backend/hip.hpp @@ -16,7 +16,7 @@ inline namespace _V1 { template <> inline backend_return_t get_native(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(const device &Obj) { Obj.getNative()); } +template <> +inline device make_device( + const backend_input_t &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(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 diff --git a/sycl/plugins/unified_runtime/ur/adapters/cuda/queue.cpp b/sycl/plugins/unified_runtime/ur/adapters/cuda/queue.cpp index 40cd2228c26f4..acb8e2d3b9fe5 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/cuda/queue.cpp +++ b/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; diff --git a/sycl/plugins/unified_runtime/ur/adapters/hip/event.cpp b/sycl/plugins/unified_runtime/ur/adapters/hip/event.cpp index d1ad57e196bf1..f4ccb200525e6 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/hip/event.cpp +++ b/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::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(hNativeEvent)); + + return UR_RESULT_SUCCESS; } diff --git a/sycl/plugins/unified_runtime/ur/adapters/hip/event.hpp b/sycl/plugins/unified_runtime/ur/adapters/hip/event.hpp index cfa9adaa8df4b..1e418519a8fde 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/hip/event.hpp +++ b/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. diff --git a/sycl/plugins/unified_runtime/ur/adapters/hip/queue.cpp b/sycl/plugins/unified_runtime/ur/adapters/hip/queue.cpp index a752f871049f1..bbbba8f17b61e 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/hip/queue.cpp +++ b/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 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(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 ComputeHIPStreams(1, HIPStream); + std::vector 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; } diff --git a/sycl/plugins/unified_runtime/ur/adapters/hip/queue.hpp b/sycl/plugins/unified_runtime/ur/adapters/hip/queue.hpp index 55c27962b980a..8b3c049638623 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/hip/queue.hpp +++ b/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 &&ComputeStreams, std::vector &&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; } }; diff --git a/sycl/source/backend.cpp b/sycl/source/backend.cpp index 5f138dcd7981c..b2cfa4e0465b1 100644 --- a/sycl/source/backend.cpp +++ b/sycl/source/backend.cpp @@ -37,6 +37,8 @@ static const PluginPtr &getPlugin(backend Backend) { return pi::getPlugin(); case backend::ext_oneapi_cuda: return pi::getPlugin(); + case backend::ext_oneapi_hip: + return pi::getPlugin(); default: throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "getPlugin: Unsupported backend " + diff --git a/sycl/test/basic_tests/interop-hip.cpp b/sycl/test/basic_tests/interop-hip.cpp index 58255c163536d..579139ff4300b 100644 --- a/sycl/test/basic_tests/interop-hip.cpp +++ b/sycl/test/basic_tests/interop-hip.cpp @@ -56,5 +56,9 @@ int main() { hip_event = get_native(Event); hip_queue = get_native(Queue); + device InteropDevice = make_device(hip_device); + event InteropEvent = make_event(hip_event, Context); + queue InteropQueue = make_queue(hip_queue, Context); + return 0; }