diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 636b72d40a..b7f039399f 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -74,12 +74,12 @@ hipError_t Event::synchronize() { auto hip_device = g_devices[deviceId()]; // Check HW status of the ROCcrl event. Note: not all ROCclr modes support HW status static constexpr bool kWaitCompletion = true; - if (!hip_device->devices()[0]->IsHwEventReady(*event_, kWaitCompletion)) { + if (!hip_device->devices()[0]->IsHwEventReady(*event_, kWaitCompletion, flags)) { if (event_->HwEvent() != nullptr) { amd::Command* command = nullptr; hipError_t status = recordCommand(command, event_->command().queue(), flags); command->enqueue(); - hip_device->devices()[0]->IsHwEventReady(command->event(), kWaitCompletion); + hip_device->devices()[0]->IsHwEventReady(command->event(), kWaitCompletion, flags); command->release(); } else { event_->awaitCompletion(); @@ -93,7 +93,7 @@ bool Event::awaitEventCompletion() { } bool EventDD::awaitEventCompletion() { - return g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_, true); + return g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_, true, flags); } hipError_t Event::elapsedTime(Event& eStop, float& ms) { diff --git a/projects/clr/hipamd/src/hip_event.hpp b/projects/clr/hipamd/src/hip_event.hpp index 0f3ec2113b..e97fa28c2a 100644 --- a/projects/clr/hipamd/src/hip_event.hpp +++ b/projects/clr/hipamd/src/hip_event.hpp @@ -102,7 +102,7 @@ class Event { if (type == Query) { ready = g_devices[deviceId()]->devices()[0]->IsHwEventReadyForcedWait(*event_); } else { - ready = g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_); + ready = g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_, flags); } return ready; } diff --git a/projects/clr/rocclr/device/device.hpp b/projects/clr/rocclr/device/device.hpp index 1a2bb18169..f246856d8e 100644 --- a/projects/clr/rocclr/device/device.hpp +++ b/projects/clr/rocclr/device/device.hpp @@ -1929,7 +1929,8 @@ class Device : public RuntimeObject { // Returns the status of HW event, associated with amd::Event virtual bool IsHwEventReady( const amd::Event& event, //!< AMD event for HW status validation - bool wait = false) const { //!< If true then forces the event completion + bool wait = false, //!< If true then forces the event completion + int hip_event_flags = 0) const { return false; }; diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index e229094138..8384545bef 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -2924,14 +2924,16 @@ bool Device::IsHwEventReadyForcedWait(const amd::Event& event) const { } // ================================================================================================ -bool Device::IsHwEventReady(const amd::Event& event, bool wait) const { +bool Device::IsHwEventReady(const amd::Event& event, bool wait, int hip_event_flags) const { void* hw_event = (event.NotifyEvent() != nullptr) ? event.NotifyEvent()->HwEvent() : event.HwEvent(); if (hw_event == nullptr) { ClPrint(amd::LOG_INFO, amd::LOG_SIG, "No HW event"); return false; } else if (wait) { - return WaitForSignal(reinterpret_cast(hw_event)->signal_, ActiveWait()); + constexpr int kHipEventBlockingSync = 0x1; + bool active_wait = !(hip_event_flags & kHipEventBlockingSync) && ActiveWait(); + return WaitForSignal(reinterpret_cast(hw_event)->signal_, active_wait); } return (hsa_signal_load_relaxed(reinterpret_cast(hw_event)->signal_) == 0); } diff --git a/projects/clr/rocclr/device/rocm/rocdevice.hpp b/projects/clr/rocclr/device/rocm/rocdevice.hpp index 3db4de144a..8a18f896c9 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.hpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.hpp @@ -287,7 +287,11 @@ class NullDevice : public amd::Device { return true; } - bool IsHwEventReady(const amd::Event& event, bool wait = false) const override { return false; } + bool IsHwEventReady(const amd::Event& event, bool wait = false, + int hip_event_flags = 0) const override { + return false; + } + bool IsHwEventReadyForcedWait(const amd::Event& event) const override { return false; } void getHwEventTime(const amd::Event& event, uint64_t* start, uint64_t* end) const override{}; void ReleaseGlobalSignal(void* signal) const override {} @@ -483,7 +487,8 @@ class Device : public NullDevice { virtual bool SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeInput, cl_set_device_clock_mode_output_amd* pSetClockModeOutput); - virtual bool IsHwEventReady(const amd::Event& event, bool wait = false) const; + virtual bool IsHwEventReady(const amd::Event& event, bool wait = false, + int hip_event_flags = 0) const; virtual bool IsHwEventReadyForcedWait(const amd::Event& event) const; virtual void getHwEventTime(const amd::Event& event, uint64_t* start, uint64_t* end) const; virtual void ReleaseGlobalSignal(void* signal) const;