dispatch vs profiling packets signals race fixed

Change-Id: I7f250d01f3d45f3f8ec2539b2187ab99503850aa


[ROCm/rocprofiler commit: caa15d7186]

projects/rocprofiler/inc/rocprofiler.h

@@ -169,7 +169,7 @@ typedef enum {
 } rocprofiler_mode_t;
 
 // Profiling handler, calling on profiling completion
-typedef void (*rocprofiler_handler_t)(rocprofiler_group_t group, void* arg);
+typedef bool (*rocprofiler_handler_t)(rocprofiler_group_t group, void* arg);
 
 // Profiling preperties
 typedef struct {

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

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

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

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

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

projects/rocprofiler/test/ctrl/tool.cpp

@@ -8,6 +8,7 @@
 #include <dirent.h>
 #include <hsa.h>
 #include <pthread.h>
+#include <sched.h>
 #include <stdio.h>
 #include <stdlib.h>
 #include <sys/types.h>
@@ -243,10 +244,15 @@ void output_group(FILE* file, const rocprofiler_group_t* group, const char* str)
 }
 
 // Dump stored context profiling output data
-void dump_context(context_entry_t* entry) {
+bool dump_context(context_entry_t* entry) {
   hsa_status_t status = HSA_STATUS_ERROR;
 
   if (entry->valid) {
+    const rocprofiler_dispatch_record_t* record = entry->data.record;
+    if (record) {
+      if (record->complete == 0) return false;
+    }
+
     ++context_collected;
     entry->valid = 0;
     const uint32_t index = entry->index;
@@ -254,18 +260,22 @@ void dump_context(context_entry_t* entry) {
     const rocprofiler_feature_t* features = entry->features;
     const unsigned feature_count = entry->feature_count;
 
-    fprintf(file_handle, "dispatch[%u], queue_index(%lu), kernel_name(\"%s\"), time(%lu,%lu,%lu,%lu)\n",
+    fprintf(file_handle, "dispatch[%u], queue_index(%lu), kernel_name(\"%s\")",
       index,
       entry->data.queue_index,
-      entry->data.kernel_name,
-      entry->data.record->dispatch,
-      entry->data.record->begin,
-      entry->data.record->end,
-      entry->data.record->complete);
+      entry->data.kernel_name);
+    if (record) fprintf(file_handle, ", time(%lu,%lu,%lu,%lu)",
+      record->dispatch,
+      record->begin,
+      record->end,
+      record->complete);
+    fprintf(file_handle, "\n");
     fflush(file_handle);
 
-    delete entry->data.record;
-    entry->data.record = NULL;
+    if (record) {
+      delete record;
+      entry->data.record = NULL;
+    }
 
     rocprofiler_group_t& group = entry->group;
     if (group.context != NULL) {
@@ -285,6 +295,8 @@ void dump_context(context_entry_t* entry) {
       rocprofiler_close(group.context);
     }
   }
+
+  return true;
 }
 
 // Dump all stored contexts profiling output data
@@ -303,8 +315,9 @@ void dump_context_array() {
 }
 
 // Profiling completion handler
-void handler(rocprofiler_group_t group, void* arg) {
+bool handler(rocprofiler_group_t group, void* arg) {
   context_entry_t* entry = reinterpret_cast<context_entry_t*>(arg);
+  bool ret = false;
 
   if (pthread_mutex_lock(&mutex) != 0) {
     perror("pthread_mutex_lock");
@@ -312,14 +325,16 @@ void handler(rocprofiler_group_t group, void* arg) {
   }
 
   if (context_array->find(entry->index) != context_array->end()) {
-    dump_context(entry);
-    dealloc_context_entry(entry);
+    if (dump_context(entry)) dealloc_context_entry(entry);
+    else ret = true;
   }
 
   if (pthread_mutex_unlock(&mutex) != 0) {
     perror("pthread_mutex_unlock");
     exit(1);
   }
+
+  return ret;
 }
 
 // Kernel disoatch callback

projects/rocprofiler/test/gfx_metrics.xml

@@ -40,7 +40,6 @@
   <metric name="TCC_MC_WRREQ_STALL" block=TCC event=28 descr="Number of cycles a write request was stalled."></metric>
 
   <metric name="TCP_TA_DATA_STALL_CYCLES" block=TCP event=3 descr="TCP stalls TA data interface. Now Windowed."></metric>
-  <metric name="TCP_TCP_TA_DATA_STALL_CYCLES" block=TCP event=3 descr="TCP stalls TA data interface. Not Windowed."></metric>
 
   <metric name=CPC_ALWAYS_COUNT block=CPC event=0 ></metric>
   <metric name=CPC_ME1_STALL_WAIT_ON_RCIU_READ block=CPC event=8 ></metric>
@@ -86,7 +85,7 @@
   <metric name=TCC_EA_RDREQ block=TCC event=41 ></metric>
   <metric name=TCC_EA_RDREQ_32B block=TCC event=42 ></metric>
 
-  <metric name=TCP_TA_DATA_STALL_CYCLES block=TCP event=6 descr="TCP stalls TA data interface. Now Windowed."></metric>
+  <metric name="TCP_TA_DATA_STALL_CYCLES" block=TCP event=6 descr="TCP stalls TA data interface. Now Windowed."></metric>
 
   <metric name=CPC_ALWAYS_COUNT block=CPC event=0 ></metric>
   <metric name=CPC_ME1_STALL_WAIT_ON_RCIU_READ block=CPC event=8 ></metric>

projects/rocprofiler/test/metrics.xml

@@ -178,7 +178,7 @@
   <metric
     name="MemUnitStalled"
     descr="The percentage of GPUTime the memory unit is stalled. Try reducing the number or size of fetches and writes if possible. Value range: 0% (optimal) to 100% (bad)."
-    expr=100*max(TCP_TCP_TA_DATA_STALL_CYCLES,16)/GRBM_GUI_ACTIVE/SE_NUM
+    expr=100*max(TCP_TA_DATA_STALL_CYCLES,16)/GRBM_GUI_ACTIVE/SE_NUM
   ></metric>
 
   # WriteUnitStalled The percentage of GPUTime the Write unit is stalled. Value range: 0% to 100% (bad).