Updating to load_delta. Fixing perfetto plugin.

Change-Id: If893f84b0ff108cfa0ccdcf717ee8592aa621032


[ROCm/rocprofiler commit: 4dd21807c0]

projects/rocprofiler/plugin/att/disassembly.cpp

@@ -276,7 +276,7 @@ void DisassemblyInstance::inst_callback(const char* instruction, void* user_data
 // return file offset, if found
 std::optional<uint64_t> DisassemblyInstance::va2fo(uint64_t va)
 {
-  /*CHECK_VA2FO(buffer.size(), "buffer is not large enough");
+  CHECK_VA2FO(buffer.size(), "buffer is not large enough");
 
   uint8_t *e_ident = (uint8_t*)buffer.data();
   CHECK_VA2FO(e_ident, "e_ident is nullptr");
@@ -297,13 +297,11 @@ std::optional<uint64_t> DisassemblyInstance::va2fo(uint64_t va)
     e_ident[EI_ABIVERSION] == 3, "unexpected ei_abiversion"); // ELFABIVERSION_AMDGPU_HSA_V5
 
   Elf64_Ehdr *ehdr = (Elf64_Ehdr*)buffer.data();
-  CHECK_VA2FO(buffer.size() > ehdr->e_phoff + sizeof(Elf64_Ehdr), "buffer is not large enough");
   CHECK_VA2FO(ehdr, "ehdr is nullptr");
   CHECK_VA2FO(ehdr->e_type == ET_DYN, "unexpected e_type");
-  CHECK_VA2FO(ehdr->e_machine == ELF::EM_AMDGPU, "unexpected e_machine"); */
+  CHECK_VA2FO(ehdr->e_machine == ELF::EM_AMDGPU, "unexpected e_machine");
 
   CHECK_VA2FO(buffer.size() > sizeof(Elf64_Ehdr), "buffer is not large enough");
-  Elf64_Ehdr *ehdr = (Elf64_Ehdr*)buffer.data();
   CHECK_VA2FO(ehdr->e_phoff != 0, "unexpected e_phoff");
 
   CHECK_VA2FO(buffer.size() > ehdr->e_phoff + sizeof(Elf64_Phdr), "buffer is not large enough");
@@ -331,7 +329,7 @@ std::optional<uint64_t> DisassemblyInstance::va2fo(uint64_t va)
 
 std::vector<std::pair<uint64_t, uint64_t>> DisassemblyInstance::getSegments()
 {
-  /*CHECK_VA2FO(buffer.size(), "buffer is not large enough");
+  CHECK_VA2FO(buffer.size(), "buffer is not large enough");
 
   uint8_t *e_ident = (uint8_t*)buffer.data();
   CHECK_VA2FO(e_ident, "e_ident is nullptr");
@@ -352,13 +350,11 @@ std::vector<std::pair<uint64_t, uint64_t>> DisassemblyInstance::getSegments()
     e_ident[EI_ABIVERSION] == 3, "unexpected ei_abiversion"); // ELFABIVERSION_AMDGPU_HSA_V5
 
   Elf64_Ehdr *ehdr = (Elf64_Ehdr*)buffer.data();
-  CHECK_VA2FO(buffer.size() > ehdr->e_phoff + sizeof(Elf64_Ehdr), "buffer is not large enough");
   CHECK_VA2FO(ehdr, "ehdr is nullptr");
   CHECK_VA2FO(ehdr->e_type == ET_DYN, "unexpected e_type");
-  CHECK_VA2FO(ehdr->e_machine == ELF::EM_AMDGPU, "unexpected e_machine"); */
+  CHECK_VA2FO(ehdr->e_machine == ELF::EM_AMDGPU, "unexpected e_machine");
 
   CHECK_VA2FO(buffer.size() > sizeof(Elf64_Ehdr), "buffer is not large enough");
-  Elf64_Ehdr *ehdr = (Elf64_Ehdr*)buffer.data();
   CHECK_VA2FO(ehdr->e_phoff != 0, "unexpected e_phoff");
 
   CHECK_VA2FO(buffer.size() > ehdr->e_phoff + sizeof(Elf64_Phdr), "buffer is not large enough");

projects/rocprofiler/plugin/att/stitch.py

@@ -347,10 +347,12 @@ def stitch(insts, raw_code, jumps, gfxv, bIsAuto, codeservice):
             line = firstinst.cycles
             lineincrement = watchlist.getincrement(line)
         except KeyError as e:
-            print('Auto error invalid addr', hex(e.args[0]))
+            print('Warning: Waves from addr', hex(e.args[0]), 'have no codeobj info.')
-            return None
+            for i in range(len(insts)):
+                insts[i].asmline = 0
+            return [i for k, i in enumerate(insts) if i.type != PCINFO], [], [], [], 1, 0, [k for k, i in enumerate(insts) if i.type == PCINFO]
         except Exception as e:
-            print('Auto error', e)
+            print('Unknown error', e)
             return None
     else:
         line = 0

projects/rocprofiler/plugin/perfetto/perfetto.cpp

@@ -298,53 +298,34 @@ class perfetto_plugin_t {
 
     TRACE_EVENT_END("KERNELS", queue_track, profiler_record.timestamps.end.value);
 
-    auto get_counter_track_fn = [&](std::string counter_name) {
+    auto get_counter_track_fn = [&](size_t i, rocprofiler_counter_id_t counter_handler)
-      std ::string counter_track_id = hostname_ + std::to_string(GetPid()) + counter_name;
+    {
-      std::pair<int, std::string> gpu_counter_track_id = std::make_pair(device_id, counter_name);
+      auto& ctrack = counter_tracks_[device_id];
-      std::unordered_map<std::string, perfetto::CounterTrack>::iterator counters_track_it;
-      {
-        counters_track_it = counter_tracks_.find(gpu_counter_track_id.second);
-        if (counters_track_it == counter_tracks_.end()) {
-          /* Create a new perfetto::Track */
-          counters_track_it =
-              counter_tracks_
-                  .emplace(gpu_counter_track_id.second,
-                           perfetto::CounterTrack(counter_track_id.c_str(), gpu_track))
-                  .first;
 
-          auto counter_track_desc = counters_track_it->second.Serialize();
+      if (i<ctrack.size()) return;
-          std::string counter_track_str = "Counter " + gpu_counter_track_id.second;
+
-          counter_track_desc.set_name(counter_track_str);
+      const char* name_c = nullptr;
-          perfetto::TrackEvent::SetTrackDescriptor(counters_track_it->second, counter_track_desc);
+      CHECK_ROCPROFILER(rocprofiler_query_counter_info(session_id, ROCPROFILER_COUNTER_NAME, counter_handler, &name_c));
-        }
+
-      }
+      ctrack.push_back(perfetto::CounterTrack(name_c, gpu_track));
-      return counters_track_it->second;
+      auto counter_track_desc = ctrack.back().Serialize();
+      counter_track_desc.set_name("Counter " + std::string(name_c));
+      perfetto::TrackEvent::SetTrackDescriptor(ctrack.back(), counter_track_desc);
     };
 
     // For Counters
     if (!profiler_record.counters) return 0;
 
-    for (uint64_t i = 0; i < profiler_record.counters_count.value; i++) {
+    for (uint64_t i = 0; i < profiler_record.counters_count.value; i++)
+    {
       if (profiler_record.counters[i].counter_handler.handle == 0) continue;
 
-      size_t name_length = 0;
+      get_counter_track_fn(i, profiler_record.counters[i].counter_handler);
-      CHECK_ROCPROFILER(rocprofiler_query_counter_info_size(
+      auto& counters_track = counter_tracks_.at(device_id).at(i);
-          session_id, ROCPROFILER_COUNTER_NAME, profiler_record.counters[i].counter_handler,
-          &name_length));
-
-      if (name_length <= 1)
-        continue;
-
-      const char* name_c = nullptr;
-      CHECK_ROCPROFILER(rocprofiler_query_counter_info(
-          session_id, ROCPROFILER_COUNTER_NAME, profiler_record.counters[i].counter_handler,
-          &name_c));
-
-      perfetto::CounterTrack counters_track = get_counter_track_fn(std::string(name_c));
       TRACE_COUNTER("COUNTERS", counters_track, profiler_record.timestamps.begin.value,
                     profiler_record.counters[i].value.value);
       // Added an extra zero event for maintaining start-end of the counter
-      TRACE_COUNTER("COUNTERS", counters_track, profiler_record.timestamps.end.value, 0.001);
+      TRACE_COUNTER("COUNTERS", counters_track, profiler_record.timestamps.end.value, 0);
     }
 
     return 0;
@@ -634,7 +615,7 @@ class perfetto_plugin_t {
   // Activity Tracks
   std::unordered_map<uint64_t, perfetto::Track> queue_tracks_;
 
-  std::unordered_map<std::string, perfetto::CounterTrack> counter_tracks_;
+  std::unordered_map<uint64_t, std::vector<perfetto::CounterTrack>> counter_tracks_;
 
   std::atomic<uint64_t> track_counter_{GetPid()};
   std::vector<uint64_t> track_ids_used_;

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

@@ -501,9 +501,9 @@ hsa_status_t CodeObjectCallback(hsa_executable_t executable,
   ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ, &data);
 
   if (data.codeobj.unload)
-    codeobj_capture_instance::Unload(data.codeobj.load_base);
+    codeobj_capture_instance::Unload(data.codeobj.load_delta);
   else
-    codeobj_capture_instance::Load(data.codeobj.load_base, data.codeobj.load_size,
+    codeobj_capture_instance::Load(data.codeobj.load_delta, data.codeobj.load_size,
             uri_str, data.codeobj.memory_base, data.codeobj.memory_size);
 
   hsa_executable_iterate_agent_symbols(executable, data.codeobj.agent,