From 227e49b3fd963a0ebcb322defa0bebb8a2e1da83 Mon Sep 17 00:00:00 2001 From: Giovanni LB Date: Fri, 1 Mar 2024 20:02:12 -0300 Subject: [PATCH] Updating to load_delta. Fixing perfetto plugin. Change-Id: If893f84b0ff108cfa0ccdcf717ee8592aa621032 [ROCm/rocprofiler commit: 4dd21807c0d5fcbd3d6b23337e54332e879c3f03] --- .../rocprofiler/plugin/att/disassembly.cpp | 12 ++-- projects/rocprofiler/plugin/att/stitch.py | 8 ++- .../rocprofiler/plugin/perfetto/perfetto.cpp | 55 ++++++------------- .../rocprofiler/src/core/hsa/hsa_support.cpp | 4 +- 4 files changed, 29 insertions(+), 50 deletions(-) diff --git a/projects/rocprofiler/plugin/att/disassembly.cpp b/projects/rocprofiler/plugin/att/disassembly.cpp index c89387ae6e..f0b5d23eb2 100644 --- a/projects/rocprofiler/plugin/att/disassembly.cpp +++ b/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 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 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 DisassemblyInstance::va2fo(uint64_t va) std::vector> 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> 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"); diff --git a/projects/rocprofiler/plugin/att/stitch.py b/projects/rocprofiler/plugin/att/stitch.py index bbc2c6f555..47c4450806 100644 --- a/projects/rocprofiler/plugin/att/stitch.py +++ b/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])) - return None + print('Warning: Waves from addr', hex(e.args[0]), 'have no codeobj info.') + 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 diff --git a/projects/rocprofiler/plugin/perfetto/perfetto.cpp b/projects/rocprofiler/plugin/perfetto/perfetto.cpp index fcb0bb79e0..e8a1d4c515 100644 --- a/projects/rocprofiler/plugin/perfetto/perfetto.cpp +++ b/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) { - std ::string counter_track_id = hostname_ + std::to_string(GetPid()) + counter_name; - std::pair gpu_counter_track_id = std::make_pair(device_id, counter_name); - std::unordered_map::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 get_counter_track_fn = [&](size_t i, rocprofiler_counter_id_t counter_handler) + { + auto& ctrack = counter_tracks_[device_id]; - auto counter_track_desc = counters_track_it->second.Serialize(); - std::string counter_track_str = "Counter " + gpu_counter_track_id.second; - counter_track_desc.set_name(counter_track_str); - perfetto::TrackEvent::SetTrackDescriptor(counters_track_it->second, counter_track_desc); - } - } - return counters_track_it->second; + if (i queue_tracks_; - std::unordered_map counter_tracks_; + std::unordered_map> counter_tracks_; std::atomic track_counter_{GetPid()}; std::vector track_ids_used_; diff --git a/projects/rocprofiler/src/core/hsa/hsa_support.cpp b/projects/rocprofiler/src/core/hsa/hsa_support.cpp index ff8324d4d0..7daab9ffbc 100644 --- a/projects/rocprofiler/src/core/hsa/hsa_support.cpp +++ b/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,