dispatch vs profiling packets signals race fixed

Change-Id: I7f250d01f3d45f3f8ec2539b2187ab99503850aa

src/core/context.h

@@ -6,6 +6,7 @@
 #include <hsa.h>
 #include <hsa_ext_amd.h>
 #include <map>
+#include <mutex>
 #include <vector>
 
 #include "core/metrics.h"
@@ -125,8 +126,7 @@ class Group {
   }
   void ResetRefs() { refs_ = n_profiles_; }
   uint32_t DecrRefs() {
-    --refs_;
-    return refs_;
+    return (refs_ > 0) ? --refs_ : 0;
   }
 
  private:
@@ -395,7 +395,7 @@ class Context {
     uint32_t r = group->DecrRefs();
     group->GetContext()->mutex_.unlock();
     if (r == 0) {
-      group->GetContext()->handler_(group->GetGroup(), group->GetContext()->handler_arg_);
+      return group->GetContext()->handler_(group->GetGroup(), group->GetContext()->handler_arg_);
     }
     return false;
   }

src/core/intercept_queue.cpp

@@ -14,4 +14,5 @@ InterceptQueue::obj_map_t* InterceptQueue::obj_map_ = NULL;
 const char* InterceptQueue::kernel_none_ = "";
 uint64_t InterceptQueue::timeout_ = UINT64_MAX;
 Tracker* InterceptQueue::tracker_ = NULL;
+bool InterceptQueue::tracker_on_ = true;
 }  // namespace rocprofiler

src/core/intercept_queue.h

@@ -41,7 +41,7 @@ class InterceptQueue {
                                            group_segment_size, queue, &status);
     if (status != HSA_STATUS_SUCCESS) abort();
 
-    if (!tracker_) tracker_ = new Tracker(timeout_);
+    if (tracker_on_ && (tracker_ == NULL)) tracker_ = new Tracker(timeout_);
     status = hsa_amd_profiling_set_profiler_enabled(*queue, true);
     if (status != HSA_STATUS_SUCCESS) abort();
 
@@ -89,15 +89,19 @@ class InterceptQueue {
         const hsa_kernel_dispatch_packet_t* dispatch_packet =
             reinterpret_cast<const hsa_kernel_dispatch_packet_t*>(packet);
         const char* kernel_name = GetKernelName(dispatch_packet);
-        const auto* entry = tracker_->Add(obj->agent_info_->dev_id, dispatch_packet->completion_signal);
-        const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal = entry->signal;
+        const rocprofiler_dispatch_record_t* record = NULL;
+        if (tracker_ != NULL) {
+          const auto* entry = tracker_->Add(obj->agent_info_->dev_id, dispatch_packet->completion_signal);
+          const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal = entry->signal;
+          record = entry->record;
+        }
         rocprofiler_callback_data_t data = {obj->agent_info_->dev_id,
                                             obj->agent_info_->dev_index,
                                             obj->queue_,
                                             user_que_idx,
                                             dispatch_packet->kernel_object,
                                             kernel_name,
-                                            entry->record};
+                                            record};
         hsa_status_t status = dispatch_callback_(&data, callback_data_, &group);
         free(const_cast<char*>(kernel_name));
         if ((status == HSA_STATUS_SUCCESS) && (group.context != NULL)) {
@@ -137,6 +141,7 @@ class InterceptQueue {
   }
 
   static void SetTimeout(uint64_t timeout) { timeout_ = timeout; }
+  static void TrackerOn(bool on) { tracker_on_ = on; }
 
  private:
   InterceptQueue(const hsa_agent_t& agent, hsa_queue_t* const queue, ProxyQueue* proxy) :
@@ -188,7 +193,7 @@ class InterceptQueue {
   static const char* kernel_none_;
   static uint64_t timeout_;
   static Tracker* tracker_;
-  static const bool tracker_on_ = true;
+  static bool tracker_on_;
 
   hsa_queue_t* const queue_;
   ProxyQueue* const proxy_;

src/core/rocprofiler.cpp

@@ -148,6 +148,11 @@ CONSTRUCTOR_API void constructor() {
     Context::SetTimeout(timeout_val);
     InterceptQueue::SetTimeout(timeout_val);
   }
+  const char* tracker_on_str = getenv("ROCP_TRACKER_ON");
+  if (tracker_on_str != NULL) {
+    if (strncmp(tracker_on_str, "true", 4) == 0) InterceptQueue::TrackerOn(true);
+    if (strncmp(tracker_on_str, "false", 4) == 0) InterceptQueue::TrackerOn(false);
+  }
 }
 
 DESTRUCTOR_API void destructor() {
@@ -173,6 +178,7 @@ const MetricsDict* GetMetrics(const hsa_agent_t& agent) {
   return metrics;
 }
 
+rocprofiler::Tracker::mutex_t rocprofiler::Tracker::mutex_;
 util::Logger::mutex_t util::Logger::mutex_;
 util::Logger* util::Logger::instance_ = NULL;
 uint64_t Context::timeout_ = UINT64_MAX;

src/core/tracker.h

@@ -6,6 +6,7 @@
 #include <hsa_ext_amd.h>
 
 #include <list>
+#include <mutex>
 
 #include "inc/rocprofiler.h"
 #include "util/exception.h"
@@ -15,6 +16,7 @@ namespace rocprofiler {
 
 class Tracker {
   public:
+  typedef std::mutex mutex_t;
   typedef rocprofiler_dispatch_record_t record_t;
   struct entry_t;
   typedef std::list<entry_t*> sig_list_t;
@@ -29,6 +31,7 @@ class Tracker {
 
   Tracker(uint64_t timeout = UINT64_MAX) : timeout_(timeout) {}
   ~Tracker() {
+    mutex_.lock();
     for (entry_t* entry : sig_list_) {
       assert(entry != NULL);
       while (1) {
@@ -43,6 +46,7 @@ class Tracker {
       }
       Del(entry);
     }
+    mutex_.unlock();
   }
 
   // Add tracker entry
@@ -50,7 +54,9 @@ class Tracker {
     entry_t* entry = new entry_t{};
     assert(entry);
     entry->tracker = this;
+    mutex_.lock();
     entry->it = sig_list_.insert(sig_list_.begin(), entry);
+    mutex_.unlock();
 
     entry->agent = agent;
     entry->orig = orig;
@@ -73,7 +79,9 @@ class Tracker {
   // Delete tracker entry
   void Del(entry_t* entry) {
     hsa_signal_destroy(entry->signal);
+    mutex_.lock();
     sig_list_.erase(entry->it);
+    mutex_.unlock();
     delete entry;
   } 
 
@@ -84,6 +92,7 @@ class Tracker {
 
     hsa_status_t status = hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP, &record->complete);
     if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP)");
+    if (record->complete == 0) EXC_RAISING(status, "hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP), time is zero");
 
     hsa_amd_profiling_dispatch_time_t dispatch_time{};
     status = hsa_amd_profiling_get_dispatch_time(entry->agent, entry->signal, &dispatch_time);
@@ -106,6 +115,8 @@ class Tracker {
   uint64_t timeout_;
   // Tracked signals list
   sig_list_t sig_list_;
+  // Inter-thread synchronization
+  static mutex_t mutex_;
 };
 
 } // namespace rocprofiler