SWDEV-439581 - hipEventBlockingSync flag for hip events

Change-Id: I0d7785a568f8007f82f999776a7ad23d0acc81b7


[ROCm/clr commit: a5a4b78606]

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) {

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;
   }

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;
   };
 

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<ProfilingSignal*>(hw_event)->signal_, ActiveWait());
+    constexpr int kHipEventBlockingSync = 0x1;
+    bool active_wait = !(hip_event_flags & kHipEventBlockingSync) && ActiveWait();
+    return WaitForSignal(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_, active_wait);
   }
   return (hsa_signal_load_relaxed(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_) == 0);
 }

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;