Fixing Kernel Dispatches timestamps issue

Change-Id: I4946ea13e2796b018dc3a9114c232b0fa5f483dc


[ROCm/rocprofiler commit: 941a0373e6]

projects/rocprofiler/src/core/hsa/queues/queue.cpp

@@ -397,10 +397,10 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
     for (auto it = pending_signals.begin(); it != pending_signals.end();
          it = pending_signals.erase(it)) {
       auto& pending = *it;
-      if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->signal)) return true;
+      if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->new_signal)) return true;
       hsa_amd_profiling_dispatch_time_t time;
       hsa_support::GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn(
-          queue_info_session->agent, pending->signal, &time);
+          queue_info_session->agent, pending->original_signal, &time);
       uint32_t record_count = 1;
       bool is_individual_xcc_mode = false;
       uint32_t xcc_count = queue_info_session->xcc_count;
@@ -469,8 +469,8 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
         }
         delete pending->context;
       }
-      if (pending->signal.handle)
-        hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(pending->signal);
+      if (pending->new_signal.handle)
+        hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(pending->new_signal);
       if (queue_info_session->interrupt_signal.handle)
         hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(queue_info_session->interrupt_signal);
     }
@@ -500,7 +500,7 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
          it = pending_signals.erase(it)) {
       auto& pending = *it;
       std::lock_guard<std::mutex> lock(session->GetSessionLock());
-      if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.signal)) return true;
+      if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.new_signal)) return true;
       rocprofiler_record_att_tracer_t record{};
       record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor};
       record.gpu_id = rocprofiler_agent_id_t{(uint64_t)queue_info_session->gpu_index};
@@ -725,13 +725,6 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
 
       auto& packet = transformed_packets.emplace_back(packets_arr[i]);
       auto& dispatch_packet = reinterpret_cast<hsa_kernel_dispatch_packet_t&>(packet);
-
-      /*
-       * Only PC sampling relies on this right now, so it would be better to
-       * only generate an ID if PC sampling is active to conserve IDs, but it's
-       * unlikely 64 bits' worth of identifiers will be exhausted during the
-       * lifetime of the ROCProfiler_Singleton.
-       */
       uint64_t correlation_id = dispatch_packet.reserved2;
 
       CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal);
@@ -745,12 +738,12 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
                                     record_id);
         if (profiles.size() > 0 && replay_mode_count > 0) {
           session->GetProfiler()->AddPendingSignals(
-              writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id,
+              writer_id, record_id, original_packet.completion_signal, dispatch_packet.completion_signal, session_id, buffer_id,
               profile.first, profile.first->metrics_list.size(), profile.second, kernel_properties,
               (uint32_t)syscall(__NR_gettid), user_pkt_index, correlation_id);
         } else {
           session->GetProfiler()->AddPendingSignals(
-              writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id,
+              writer_id, record_id, original_packet.completion_signal, dispatch_packet.completion_signal, session_id, buffer_id,
               nullptr, 0, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid),
               user_pkt_index, correlation_id);
         }
@@ -946,11 +939,11 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
                                   record_id);
       if (session && profile) {
         session->GetAttTracer()->AddPendingSignals(
-            writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id, profile,
+            writer_id, record_id, original_packet.completion_signal, dispatch_packet.completion_signal, session_id, buffer_id, profile,
             kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index);
       } else {
         session->GetAttTracer()->AddPendingSignals(
-            writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id, nullptr,
+            writer_id, record_id, original_packet.completion_signal, dispatch_packet.completion_signal, session_id, buffer_id, nullptr,
             kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index);
       }
 

projects/rocprofiler/src/core/session/att/att.cpp

@@ -31,7 +31,7 @@ AttTracer::AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t
 AttTracer::~AttTracer() {}
 
 void AttTracer::AddPendingSignals(uint32_t writer_id, uint64_t kernel_object,
-                                   const hsa_signal_t& completion_signal,
+                                   const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal,
                                    rocprofiler_session_id_t session_id,
                                    rocprofiler_buffer_id_t buffer_id,
                                    hsa_ven_amd_aqlprofile_profile_t* profile,
@@ -41,7 +41,7 @@ void AttTracer::AddPendingSignals(uint32_t writer_id, uint64_t kernel_object,
   if (sessions_pending_signals_.find(writer_id) == sessions_pending_signals_.end())
     sessions_pending_signals_.emplace(writer_id, std::vector<att_pending_signal_t>());
   sessions_pending_signals_.at(writer_id).emplace_back(
-      att_pending_signal_t{kernel_object, completion_signal, session_id_, buffer_id, profile,
+      att_pending_signal_t{kernel_object, original_completion_signal, new_completion_signal, session_id_, buffer_id, profile,
                             kernel_properties, thread_id, queue_index});
 }
 

projects/rocprofiler/src/core/session/att/att.h

@@ -34,7 +34,8 @@ namespace rocprofiler {
 
 typedef struct {
   uint64_t kernel_descriptor;
-  hsa_signal_t signal;
+  hsa_signal_t original_signal;
+  hsa_signal_t new_signal;
   rocprofiler_session_id_t session_id;
   rocprofiler_buffer_id_t buffer_id;
   hsa_ven_amd_aqlprofile_profile_t* profile;
@@ -52,7 +53,7 @@ class AttTracer {
   ~AttTracer();
 
   void AddPendingSignals(uint32_t writer_id, uint64_t kernel_object,
-                         const hsa_signal_t& completion_signal, rocprofiler_session_id_t session_id,
+                         const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id,
                          rocprofiler_buffer_id_t buffer_id, hsa_ven_amd_aqlprofile_profile_t* profile,
                          rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id,
                          uint64_t queue_index);

projects/rocprofiler/src/core/session/profiler/profiler.cpp

@@ -125,7 +125,7 @@ bool Profiler::HasActivePass() {
 }
 
 void Profiler::AddPendingSignals(
-    uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& completion_signal,
+    uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal,
     rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id,
     rocprofiler::profiling_context_t* context, uint64_t session_data_count,
     hsa_ven_amd_aqlprofile_profile_t* profile, rocprofiler_kernel_properties_t kernel_properties,
@@ -134,7 +134,7 @@ void Profiler::AddPendingSignals(
   if (sessions_pending_signals_->find(writer_id) == sessions_pending_signals_->end())
     sessions_pending_signals_->emplace(writer_id, std::vector<pending_signal_t*>());
   sessions_pending_signals_->at(writer_id).emplace_back(new pending_signal_t{
-      kernel_object, completion_signal, session_id_, buffer_id, context, session_data_count,
+      kernel_object, original_completion_signal, new_completion_signal, session_id_, buffer_id, context, session_data_count,
       profile, kernel_properties, thread_id, queue_index, correlation_id});
 }
 

projects/rocprofiler/src/core/session/profiler/profiler.h

@@ -44,7 +44,8 @@ namespace rocprofiler {
 
 typedef struct {
   uint64_t kernel_descriptor;
-  hsa_signal_t signal;
+  hsa_signal_t original_signal;
+  hsa_signal_t new_signal;
   rocprofiler_session_id_t session_id;
   rocprofiler_buffer_id_t buffer_id;
   rocprofiler::profiling_context_t* context;
@@ -67,7 +68,7 @@ class Profiler {
   ~Profiler();
 
   void AddPendingSignals(uint32_t writer_id, uint64_t kernel_object,
-                           const hsa_signal_t& completion_signal, rocprofiler_session_id_t session_id,
+                           const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id,
                            rocprofiler_buffer_id_t buffer_id,
                            rocprofiler::profiling_context_t* context, uint64_t session_data_count,
                            hsa_ven_amd_aqlprofile_profile_t* profile,