Fixing Kernel Names

The lifetime of the kernel descriptor handle is unknown, so we currently uses the record id to identify the kernel name

Change-Id: I8802ce245bb021bcb43dd97cd6334bc88de0faf4


[ROCm/rocprofiler commit: 1322e527a6]

projects/rocprofiler/src/api/rocmtool.cpp

@@ -199,7 +199,7 @@ size_t rocmtool::GetKernelInfoSize(rocprofiler_kernel_info_kind_t kind,
                                    rocprofiler_kernel_id_t kernel_id) {
   switch (kind) {
     case ROCPROFILER_KERNEL_NAME:
-      return GetKernelNameFromKsymbols(kernel_id.handle).size();
+      return GetKernelNameUsingDispatchID(kernel_id.handle).size();
     default:
       warning("The provided Kernel Kind is not yet supported!");
       return 0;
@@ -209,7 +209,7 @@ const char* rocmtool::GetKernelInfo(rocprofiler_kernel_info_kind_t kind,
                                     rocprofiler_kernel_id_t kernel_id) {
   switch (kind) {
     case ROCPROFILER_KERNEL_NAME:
-      return strdup(GetKernelNameFromKsymbols(kernel_id.handle).c_str());
+      return strdup(GetKernelNameUsingDispatchID(kernel_id.handle).c_str());
     default:
       warning("The provided Kernel Kind is not yet supported!");
       return "";

projects/rocprofiler/src/core/hsa/hsa_support.cpp

@@ -52,6 +52,7 @@ namespace {
 
 hsa_status_t hsa_executable_iteration_callback(hsa_executable_t executable, hsa_agent_t agent,
                                                hsa_executable_symbol_t symbol, void* args) {
+  
   hsa_symbol_kind_t type;
   rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
       symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &type);
@@ -62,14 +63,21 @@ hsa_status_t hsa_executable_iteration_callback(hsa_executable_t executable, hsa_
     // TODO(aelwazir): to be removed if the HSA fixed the issue of corrupted
     // names overflowing the length given
     if (name_length > 1) {
-      char name[name_length + 1];
-      uint64_t kernel_object;
-      rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
-          symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
-      rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
-          symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel_object);
-      std::string kernel_name = std::string(name).substr(0, name_length);
-      rocmtools::AddKernelName(kernel_object, kernel_name);
+      if(!(*static_cast<bool*>(args))) {
+        char name[name_length + 1];
+        uint64_t kernel_object;
+        rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
+            symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
+        rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
+            symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel_object);
+        std::string kernel_name = std::string(name).substr(0, name_length);
+        rocmtools::AddKernelName(kernel_object, kernel_name);
+      } else {
+        uint64_t kernel_object;
+        rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
+            symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel_object);
+        rocmtools::RemoveKernelName(kernel_object);
+      }
     }
   }
 
@@ -447,7 +455,7 @@ hsa_status_t CodeObjectCallback(hsa_executable_t executable,
   ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ, &data);
 
   hsa_executable_iterate_agent_symbols(executable, data.codeobj.agent,
-                                       hsa_executable_iteration_callback, nullptr);
+                                       hsa_executable_iteration_callback, &(data.codeobj.unload));
 
   return HSA_STATUS_SUCCESS;
 }

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

@@ -63,6 +63,10 @@ void AddKernelName(uint64_t handle, std::string name) {
   std::lock_guard<std::mutex> lock(ksymbol_map_lock);
   ksymbols->emplace(handle, name);
 }
+void RemoveKernelName(uint64_t handle) {
+  std::lock_guard<std::mutex> lock(ksymbol_map_lock);
+  ksymbols->erase(handle);
+}
 std::string GetKernelNameFromKsymbols(uint64_t handle) {
   std::lock_guard<std::mutex> lock(ksymbol_map_lock);
   if(ksymbols->find(handle)!=ksymbols->end())
@@ -70,11 +74,39 @@ std::string GetKernelNameFromKsymbols(uint64_t handle) {
   else
     return "Unknown Kernel!";
 }
+
+static std::mutex kernel_names_map_lock;
+static std::map<std::string, std::vector<uint64_t>>* kernel_names;
+static std::atomic<bool> kernel_names_flag{true};
+void AddKernelNameWithDispatchID(std::string name, uint64_t id) {
+  std::lock_guard<std::mutex> lock(kernel_names_map_lock);
+  if(kernel_names->find(name) == kernel_names->end())
+    kernel_names->emplace(name, std::vector<uint64_t>());
+  kernel_names->at(name).push_back(id);
+}
+std::string GetKernelNameUsingDispatchID(uint64_t given_id) {
+  std::lock_guard<std::mutex> lock(kernel_names_map_lock);
+  for(auto kernel_name : (*kernel_names)) {
+    for(auto dispatch_id : kernel_name.second) {
+      if(dispatch_id == given_id)
+        return kernel_name.first;
+    }
+  }
+  return "Unknown Kernel!";
+}
+
 void InitKsymbols() {
   if (ksymbols_flag.load(std::memory_order_relaxed)) {
-    std::lock_guard<std::mutex> lock(ksymbol_map_lock);
-    ksymbols = new std::map<uint64_t, std::string>();
-    ksymbols_flag.exchange(false, std::memory_order_release);
+    {
+      std::lock_guard<std::mutex> lock(ksymbol_map_lock);
+      ksymbols = new std::map<uint64_t, std::string>();
+      ksymbols_flag.exchange(false, std::memory_order_release);
+    }
+    {
+      std::lock_guard<std::mutex> lock(kernel_names_map_lock);
+      kernel_names = new std::map<std::string, std::vector<uint64_t>>();
+      kernel_names_flag.exchange(false, std::memory_order_release);
+    }
   }
 }
 void FinitKsymbols() {
@@ -84,8 +116,16 @@ void FinitKsymbols() {
     delete ksymbols;
     ksymbols_flag.exchange(true, std::memory_order_release);
   }
+  if (!kernel_names_flag.load(std::memory_order_relaxed)) {
+    std::lock_guard<std::mutex> lock(kernel_names_map_lock);
+    kernel_names->clear();
+    delete kernel_names;
+    kernel_names_flag.exchange(true, std::memory_order_release);
+  }
 }
 
+
+
 struct kernel_descriptor_t {
   uint8_t reserved0[16];
   int64_t kernel_code_entry_byte_offset;
@@ -416,7 +456,6 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
       hsa_support::GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn(
           queue_info_session->agent, pending.signal, &time);
       rocprofiler_record_profiler_t record{};
-      record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor};
       record.gpu_id = rocprofiler_agent_id_t{
           (uint64_t)hsa_support::GetAgentInfo(queue_info_session->agent.handle).getIndex()};
       record.kernel_properties = pending.kernel_properties;
@@ -429,7 +468,8 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
         AddRecordCounters(&record, pending);
       }
       record.header = {ROCPROFILER_PROFILER_RECORD,
-                       rocprofiler_record_id_t{GetROCMToolObj()->GetUniqueRecordId()}};
+                       rocprofiler_record_id_t{pending.kernel_descriptor}};
+      record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor};
 
       if (pending.session_id.handle == 0) {
         pending.session_id = GetROCMToolObj()->GetCurrentSessionId();
@@ -509,7 +549,7 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
         AddAttRecord(&record, queue_info_session->agent, pending);
       }
       record.header = {ROCPROFILER_ATT_TRACER_RECORD,
-                       rocprofiler_record_id_t{GetROCMToolObj()->GetUniqueRecordId()}};
+                       rocprofiler_record_id_t{pending.kernel_descriptor}};
 
       if (pending.session_id.handle == 0) {
         pending.session_id = GetROCMToolObj()->GetCurrentSessionId();
@@ -740,14 +780,16 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
         rocprofiler_kernel_properties_t kernel_properties =
             set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent());
         if (session) {
+          uint64_t record_id = GetROCMToolObj()->GetUniqueRecordId();
+          AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object), record_id);
           if (profiles && replay_mode_count > 0) {
             session->GetProfiler()->AddPendingSignals(
-                writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal,
+                writer_id, record_id, 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);
           } else {
             session->GetProfiler()->AddPendingSignals(
-                writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal,
+                writer_id, record_id, dispatch_packet.completion_signal,
                 session_id, buffer_id, nullptr, 0, nullptr, kernel_properties,
                 (uint32_t)syscall(__NR_gettid), user_pkt_index);
           }
@@ -929,13 +971,15 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
       // list to be processed by the signal interrupt
       rocprofiler_kernel_properties_t kernel_properties =
           set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent());
+      uint64_t record_id = GetROCMToolObj()->GetUniqueRecordId();
+      AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object), record_id);
       if (session && profile) {
         session->GetAttTracer()->AddPendingSignals(
-            writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal, session_id,
+            writer_id, record_id, 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, dispatch_packet.kernel_object, dispatch_packet.completion_signal, session_id,
+            writer_id, record_id, dispatch_packet.completion_signal, session_id,
             buffer_id, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index);
       }
 

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

@@ -42,6 +42,9 @@ namespace rocmtools {
 void InitKsymbols();
 void FinitKsymbols();
 void AddKernelName(uint64_t handle, std::string kernel_name);
+void RemoveKernelName(uint64_t handle);
+void AddKernelNameWithDispatchID(std::string name, uint64_t id);
+std::string GetKernelNameUsingDispatchID(uint64_t given_id);
 std::string GetKernelNameFromKsymbols(uint64_t handle);
 uint32_t GetCurrentActiveInterruptSignalsCount();