hsa_rsrc_factory sync

Change-Id: Idecbc0cdad6068eae5259cb043bdf5746b430aec

src/core/tracker.h

@@ -167,7 +167,12 @@ class Tracker {
     auto end = sig_list_.end();
     while (it != end) {
       auto cur = it++;
-      hsa_rsrc_->SignalWait((*cur)->signal);
+// The wait should be optiona as there possible some inter kernel dependencies and it possible to wait for
+// the kernels will never be lunched as the application was finished by some reason.
+#if 0
+      // FIXME: currently the signal value for tracking signals are taken from original application signal
+      hsa_rsrc_->SignalWait((*cur)->signal, 1);
+#endif
       Erase(cur);
     }
   }

src/util/hsa_rsrc_factory.cpp

@@ -44,9 +44,6 @@ POSSIBILITY OF SUCH DAMAGE.
 #include <string>
 #include <vector>
 
-#include "util/exception.h"
-#include "util/logger.h"
-
 namespace rocprofiler {
 namespace util {
 
@@ -523,22 +520,25 @@ uint8_t* HsaRsrcFactory::AllocateCmdMemory(const AgentInfo* agent_info, size_t s
 }
 
 // Wait signal
-void HsaRsrcFactory::SignalWait(const hsa_signal_t& signal) const {
+hsa_signal_value_t HsaRsrcFactory::SignalWait(const hsa_signal_t& signal, const hsa_signal_value_t& signal_value) const {
+  const hsa_signal_value_t exp_value = signal_value - 1;
+  hsa_signal_value_t ret_value = signal_value;
   while (1) {
-    const hsa_signal_value_t signal_value =
-      hsa_api_.hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, 1, timeout_, HSA_WAIT_STATE_BLOCKED);
-    if (signal_value == 0) {
-      break;
-    } else {
-      if (signal_value == 1) WARN_LOGGING("signal waiting...");
-      else EXC_RAISING(HSA_STATUS_ERROR, "hsa_signal_wait_scacquire (" << signal_value << ")");
+    ret_value =
+      hsa_api_.hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, signal_value, timeout_, HSA_WAIT_STATE_BLOCKED);
+    if (ret_value == exp_value) break;
+    if (ret_value != signal_value) {
+      std::cerr << "Error: HsaRsrcFactory::SignalWait: signal_value(" << signal_value
+                << "), ret_value(" << ret_value << ")" << std::endl << std::flush;
+      abort();
     }
   }
+  return ret_value;
 }
 
 // Wait signal with signal value restore
 void HsaRsrcFactory::SignalWaitRestore(const hsa_signal_t& signal, const hsa_signal_value_t& signal_value) const {
-  SignalWait(signal);
+  SignalWait(signal, signal_value);
   hsa_api_.hsa_signal_store_relaxed(const_cast<hsa_signal_t&>(signal), signal_value);
 }
 
@@ -551,7 +551,7 @@ bool HsaRsrcFactory::Memcpy(const hsa_agent_t& agent, void* dst, const void* src
     CHECK_STATUS("hsa_signal_create()", status);
     status = hsa_api_.hsa_amd_memory_async_copy(dst, cpu_agents_[0], src, agent, size, 0, NULL, s);
     CHECK_STATUS("hsa_amd_memory_async_copy()", status);
-    SignalWait(s);
+    SignalWait(s, 1);
     status = hsa_api_.hsa_signal_destroy(s);
     CHECK_STATUS("hsa_signal_destroy()", status);
   }

src/util/hsa_rsrc_factory.h

@@ -361,7 +361,7 @@ class HsaRsrcFactory {
   uint8_t* AllocateCmdMemory(const AgentInfo* agent_info, size_t size);
 
   // Wait signal
-  void SignalWait(const hsa_signal_t& signal) const;
+  hsa_signal_value_t SignalWait(const hsa_signal_t& signal, const hsa_signal_value_t& signal_value) const;
 
   // Wait signal with signal value restore
   void SignalWaitRestore(const hsa_signal_t& signal, const hsa_signal_value_t& signal_value) const;