hsa_rsrc_factory sync

Change-Id: Idecbc0cdad6068eae5259cb043bdf5746b430aec


[ROCm/rocprofiler commit: 5615ade977]

projects/rocprofiler/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);
     }
   }

projects/rocprofiler/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);
   }

projects/rocprofiler/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;

projects/rocprofiler/test/util/hsa_rsrc_factory.cpp

@@ -144,6 +144,11 @@ HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa) : initialize_hsa_(initialize
   CHECK_STATUS("HSA timer allocation failed",
     (timer_ == NULL) ? HSA_STATUS_ERROR : HSA_STATUS_SUCCESS);
 
+  // Time correlation
+  const uint32_t corr_iters = 1000;
+  CorrelateTime(HsaTimer::TIME_ID_CLOCK_REALTIME, corr_iters);
+  CorrelateTime(HsaTimer::TIME_ID_CLOCK_MONOTONIC, corr_iters);
+
   // System timeout
   timeout_ = (timeout_ns_ == HsaTimer::TIMESTAMP_MAX) ? timeout_ns_ : timer_->ns_to_sysclock(timeout_ns_);
 }
@@ -512,21 +517,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 {
-      CHECK_STATUS("hsa_signal_wait_scacquire()", HSA_STATUS_ERROR);
+    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);
 }
 
@@ -539,7 +548,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);
   }

projects/rocprofiler/test/util/hsa_rsrc_factory.h

@@ -35,6 +35,7 @@ POSSIBILITY OF SUCH DAMAGE.
 #include <stdio.h>
 #include <stdlib.h>
 #include <string.h>
+#include <time.h>
 
 #include <atomic>
 #include <iostream>
@@ -175,6 +176,12 @@ class HsaTimer {
   static const timestamp_t TIMESTAMP_MAX = UINT64_MAX;
   typedef long double freq_t;
 
+  enum time_id_t {
+    TIME_ID_CLOCK_REALTIME = 0,
+    TIME_ID_CLOCK_MONOTONIC = 1,
+    TIME_ID_NUMBER
+  };
+
   HsaTimer(const hsa_pfn_t* hsa_api) : hsa_api_(hsa_api) {
     timestamp_t sysclock_hz = 0;
     hsa_status_t status = hsa_api_->hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz);
@@ -190,6 +197,11 @@ class HsaTimer {
     return timestamp_t((freq_t)time / sysclock_factor_);
   }
 
+  // Method for timespec/ns conversion
+  timestamp_t timespec_to_ns(const timespec& time) const {
+    return ((timestamp_t)time.tv_sec * 1000000000) + time.tv_nsec;
+  }
+
   // Return timestamp in 'ns'
   timestamp_t timestamp_ns() const {
     timestamp_t sysclock;
@@ -198,6 +210,54 @@ class HsaTimer {
     return sysclock_to_ns(sysclock);
   }
 
+  // Return time in 'ns'
+  timestamp_t clocktime_ns(clockid_t clock_id) const {
+    timespec time;
+    clock_gettime(clock_id, &time);
+    return timespec_to_ns(time);
+  }
+
+  // Return pair of correlated values of profiling timestamp and time with
+  // correlation error for a given time ID and number of iterations
+  void correlated_pair_ns(time_id_t time_id, uint32_t iters,
+                          timestamp_t* timestamp_v, timestamp_t* time_v, timestamp_t* error_v) {
+    clockid_t clock_id = 0;
+    switch (clock_id) {
+      case TIME_ID_CLOCK_REALTIME:
+        clock_id = CLOCK_REALTIME;
+        break;
+      case TIME_ID_CLOCK_MONOTONIC:
+        clock_id = CLOCK_MONOTONIC;
+        break;
+      default:
+        CHECK_STATUS("internal error: invalid time_id", HSA_STATUS_ERROR);
+    }
+
+    std::vector<timestamp_t> ts_vec(iters);
+    std::vector<timespec> tm_vec(iters);
+    const uint32_t steps = iters - 1;
+
+    for (uint32_t i = 0; i < iters; ++i) {
+      hsa_api_->hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP, &ts_vec[i]);
+      clock_gettime(clock_id, &tm_vec[i]);
+    }
+
+    const timestamp_t ts_base = sysclock_to_ns(ts_vec.front());
+    const timestamp_t tm_base = timespec_to_ns(tm_vec.front());
+    const timestamp_t error = (ts_vec.back() - ts_vec.front()) / (2 * steps);
+
+    timestamp_t ts_accum = 0;
+    timestamp_t tm_accum = 0;
+    for (uint32_t i = 0; i < iters; ++i) {
+      ts_accum += (ts_vec[i] - ts_base);
+      tm_accum += (timespec_to_ns(tm_vec[i]) - tm_base);
+    }
+
+    *timestamp_v = (ts_accum / iters) + ts_base + error;
+    *time_v = (tm_accum / iters) + tm_base;
+    *error_v = error;
+  }
+
  private:
   // Timestamp frequency factor
   freq_t sysclock_factor_;
@@ -299,7 +359,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;
@@ -357,6 +417,21 @@ class HsaRsrcFactory {
     if (instance_ != NULL) Instance().timeout_ = Instance().timer_->ns_to_sysclock(time);
   }
 
+  void CorrelateTime(HsaTimer::time_id_t time_id, uint32_t iters) {
+    timestamp_t timestamp_v = 0;
+    timestamp_t time_v = 0;
+    timestamp_t error_v = 0;
+    timer_->correlated_pair_ns(time_id, iters, &timestamp_v, &time_v, &error_v);
+    time_shift_[time_id] = time_v - timestamp_v;
+    time_error_[time_id] = error_v;
+  }
+
+  hsa_status_t GetTime(uint32_t time_id, uint64_t value, uint64_t* time) {
+    if (time_id >= HsaTimer::TIME_ID_NUMBER) return HSA_STATUS_ERROR;
+    *time = value + time_shift_[time_id];
+    return HSA_STATUS_SUCCESS;
+  }
+
  private:
   // System agents iterating callback
   static hsa_status_t GetHsaAgentsCallback(hsa_agent_t agent, void* data);
@@ -421,6 +496,10 @@ class HsaRsrcFactory {
   // HSA timer
   HsaTimer* timer_;
 
+  // Time shift array to support time conversion
+  timestamp_t time_shift_[HsaTimer::TIME_ID_NUMBER];
+  timestamp_t time_error_[HsaTimer::TIME_ID_NUMBER];
+
   // CPU/kern-arg memory pools
   hsa_amd_memory_pool_t *cpu_pool_;
   hsa_amd_memory_pool_t *kern_arg_pool_;