From aba6a1d986c468090b685c521786d71469d73768 Mon Sep 17 00:00:00 2001 From: Giovanni LB Date: Wed, 15 Nov 2023 12:18:51 -0300 Subject: [PATCH] SWDEV-432445: ATT continuous mode Change-Id: I52732fc1dba41df3373ef6a19a428b00d729bf04 --- include/rocprofiler/v2/rocprofiler.h | 9 +- plugin/att/att.cpp | 97 ++-- plugin/att/att.py | 53 +- plugin/att/code_printing.cpp | 208 +++++++- plugin/att/code_printing.hpp | 51 +- plugin/att/disassembly.cpp | 87 ++-- plugin/att/disassembly.hpp | 2 +- plugin/att/service.py | 95 ++++ plugin/att/stitch.py | 210 +++++--- plugin/exportmap | 6 + src/api/CMakeLists.txt | 4 +- src/core/hsa/hsa_support.cpp | 140 ++++-- src/core/hsa/hsa_support.h | 1 + src/core/hsa/packets/packets_generator.cpp | 113 ++++- src/core/hsa/packets/packets_generator.h | 17 + src/core/hsa/queues/queue.cpp | 559 +++------------------ src/core/hsa/queues/queue.h | 6 + src/core/isa_capture/code_object_track.cpp | 27 +- src/core/isa_capture/code_object_track.hpp | 11 +- src/core/proxy_queue.h | 4 +- src/core/session/att/att.cpp | 291 ++++++++++- src/core/session/att/att.h | 135 ++++- src/core/session/att/att_header.h | 38 ++ src/core/session/att/continuous.cpp | 195 +++++++ src/core/session/att/singlekernel.cpp | 233 +++++++++ src/core/session/filter.cpp | 7 +- src/core/session/filter.h | 10 +- src/tools/tool.cpp | 78 ++- src/util/hsa_rsrc_factory.cpp | 5 +- 29 files changed, 1841 insertions(+), 851 deletions(-) create mode 100644 plugin/att/service.py create mode 100644 src/core/session/att/att_header.h create mode 100644 src/core/session/att/continuous.cpp create mode 100644 src/core/session/att/singlekernel.cpp diff --git a/include/rocprofiler/v2/rocprofiler.h b/include/rocprofiler/v2/rocprofiler.h index 95b8f2b277..ed51b99f71 100644 --- a/include/rocprofiler/v2/rocprofiler.h +++ b/include/rocprofiler/v2/rocprofiler.h @@ -1206,6 +1206,10 @@ typedef struct { * If the obj is still loaded by the time the record was generated, this value is 0 */ rocprofiler_timestamp_t clock_end; + /** + * Identifier for code object loading. This is the marker ID sent to the ATT buffer. + */ + uint32_t att_marker_id; } rocprofiler_intercepted_codeobj_t; /** @@ -1869,7 +1873,10 @@ typedef struct { rocprofiler_hip_function_name_t* hip_functions_names; rocprofiler_hsa_function_name_t* hsa_functions_names; uint32_t range[2]; - uint64_t* dispatch_ids; + struct { + uint64_t start; + uint64_t end; + }* dispatch_ids; }; /** * Data array count diff --git a/plugin/att/att.cpp b/plugin/att/att.cpp index b686e1f79d..07df60a61a 100644 --- a/plugin/att/att.cpp +++ b/plugin/att/att.cpp @@ -44,13 +44,11 @@ #include "rocprofiler.h" #include "rocprofiler_plugin.h" #include "../utils.h" -#include "code_printing.hpp" -#include "../../src/core/session/att/att.h" +#include "../../src/core/session/att/att_header.h" #define ATT_FILENAME_MAXBYTES 90 #define TEST_INVALID_KERNEL size_t(-1) -namespace { class att_plugin_t { public: att_plugin_t(void* data) { @@ -65,22 +63,16 @@ class att_plugin_t { header.raw = reinterpret_cast(data); header.reserved = 0x11; - - isa_mode = static_cast(header.isadumpmode); - header.isadumpmode = 0; } bool MPI_ENABLE = false; int MPI_RANK = 0; - std::mutex writing_lock; + static std::mutex writing_lock; bool is_valid_{true}; rocprofiler::att_header_packet_t header{.raw = 0}; - rocprofiler::rocprofiler_att_isa_dump_mode isa_mode = rocprofiler::ISA_MODE_DUMP_ALL; bool CheckAddrMatches(uint64_t kernel_addr, uint64_t base_address, uint64_t size) { - if (isa_mode == rocprofiler::ISA_MODE_DUMP_ALL) - return true; return (kernel_addr >= base_address) && (kernel_addr < base_address + size); } @@ -93,7 +85,6 @@ class att_plugin_t { int FlushATTRecord(const rocprofiler_record_att_tracer_t* att_tracer_record, rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { - std::lock_guard lock(writing_lock); if (!att_tracer_record) return ROCPROFILER_STATUS_ERROR; @@ -122,11 +113,11 @@ class att_plugin_t { if (name_demangled.size() > ATT_FILENAME_MAXBYTES) // Limit filename size name_demangled = name_demangled.substr(0, ATT_FILENAME_MAXBYTES); - std::string outfilepath = "."; - if (const char* env = getenv("OUTPUT_PATH")) outfilepath = std::string(env); + std::string output_dir = "."; + if (const char* env = getenv("OUTPUT_PATH")) output_dir = std::string(env); - outfilepath.reserve(outfilepath.size() + 128); // Max filename size - outfilepath += '/' + name_demangled; + std::string outfilepath = output_dir + '/' + name_demangled; + outfilepath.reserve(output_dir.size() + 128); // Max filename size if (MPI_ENABLE) outfilepath += "_rank" + std::to_string(MPI_RANK); outfilepath += "_v"; @@ -139,14 +130,16 @@ class att_plugin_t { auto writer_id = att_tracer_record->writer_id; std::string fname = outfilepath + "_kernel.txt"; - std::ofstream(fname.c_str()) << name_demangled << " dispatch[" << writer_id << "] GPU[" - << att_tracer_record->gpu_id.handle << "]: " << kernel_name_mangled - << '\n'; + std::ofstream kernel_txt_file((outfilepath + "_kernel.txt").c_str()); + kernel_txt_file << name_demangled << " dispatch[" << writer_id << "] GPU[" + << att_tracer_record->gpu_id.handle << "]: " << kernel_name_mangled + << '\n'; // iterate over each shader engine att trace header.navi = !att_tracer_record->intercept_list.userdata & 0x1; int se_num = att_tracer_record->shader_engine_data_count; - for (int i = 0; i < se_num; i++) { + for (int i = 0; i < se_num; i++) + { if (!att_tracer_record->shader_engine_data || !att_tracer_record->shader_engine_data[i].buffer_ptr) continue; @@ -166,50 +159,39 @@ class att_plugin_t { out.write(data_buffer_ptr, se_att_trace->buffer_size); } - if (isa_mode == rocprofiler::ISA_MODE_DUMP_NONE) - return 0; + for (size_t i = 0; i < att_tracer_record->intercept_list.count; i++) + { + const auto& symbol = att_tracer_record->intercept_list.symbols[i]; + if (!symbol.filepath) continue; - uint64_t kernel_addr = att_tracer_record->intercept_list.userdata >> 1; + std::string sfilepath(symbol.filepath); + bool bCopiedData = symbol.data && symbol.data_size; - std::ofstream isafile(outfilepath + "_isa.s"); - if (!isafile.is_open()) { - std::cerr << "Could not open ISA file: " << outfilepath << "_isa.s" << std::endl; - return ROCPROFILER_STATUS_ERROR; - } - isafile << " " << kernel_name_mangled << '\n'; + if (bCopiedData) + { + auto pos = sfilepath.find("://"); + auto rpos = sfilepath.rfind('/'); - for (size_t i = 0; i < att_tracer_record->intercept_list.count; i++) { - const rocprofiler_intercepted_codeobj_t& symbol = - att_tracer_record->intercept_list.symbols[i]; + if (pos == std::string::npos || pos+3 >= sfilepath.size()) continue; - if (!CheckAddrMatches(kernel_addr, symbol.base_address, symbol.mem_size)) continue; - - std::unique_ptr binary; - std::unique_ptr decoder; - - if (symbol.data && symbol.data_size) { - decoder = std::make_unique(symbol.data, symbol.data_size); - } else if (std::string(symbol.filepath).find("file://") != std::string::npos) { - binary = std::make_unique(symbol.filepath); - decoder = - std::make_unique(binary->buffer.data(), binary->buffer.size()); - } else { - continue; + std::string type(sfilepath.begin(), sfilepath.begin()+pos); + std::string cut(sfilepath.begin()+rpos+1, sfilepath.end()); + sfilepath = type + cut + ".out"; } - if (isa_mode == rocprofiler::ISA_MODE_DUMP_KERNEL) - decoder->disassemble_single_kernel(kernel_addr-symbol.base_address); - else - decoder->disassemble_kernels(); + kernel_txt_file << std::hex << "0x" << symbol.base_address << " 0x" << symbol.mem_size + << ' ' << std::dec << symbol.att_marker_id << ' ' << sfilepath << '\n'; - for (auto& instance : decoder->instructions) { - uint64_t addr = instance.address + symbol.base_address; + sfilepath = output_dir + '/' + sfilepath; + if (!bCopiedData || att_file_exists(sfilepath)) continue; - if (decoder->m_symbol_map.find(instance.address) != decoder->m_symbol_map.end()) - isafile << "; Begin " << decoder->m_symbol_map[instance.address].name << '\n'; - if (instance.cpp_reference) isafile << "; " << instance.cpp_reference << '\n'; - isafile << instance.instruction << " // " << std::hex << addr << '\n'; + std::ofstream isafile(sfilepath, std::ios::binary); + if (!isafile.is_open()) { + std::cerr << "Could not open file: " << sfilepath << std::endl; + return ROCPROFILER_STATUS_ERROR; } + + isafile.write(symbol.data, symbol.data_size); } return ROCPROFILER_STATUS_SUCCESS; @@ -248,8 +230,7 @@ class att_plugin_t { }; att_plugin_t* att_plugin = nullptr; - -} // namespace +std::mutex att_plugin_t::writing_lock; ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version, uint32_t rocprofiler_minor_version, @@ -258,6 +239,7 @@ ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_ rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR) return ROCPROFILER_STATUS_ERROR; + std::lock_guard lock(att_plugin_t::writing_lock); if (att_plugin != nullptr) return ROCPROFILER_STATUS_ERROR; att_plugin = new att_plugin_t(data); @@ -270,6 +252,7 @@ ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_ } ROCPROFILER_EXPORT void rocprofiler_plugin_finalize() { + std::lock_guard lock(att_plugin_t::writing_lock); if (!att_plugin) return; delete att_plugin; att_plugin = nullptr; @@ -278,11 +261,13 @@ ROCPROFILER_EXPORT void rocprofiler_plugin_finalize() { ROCPROFILER_EXPORT int rocprofiler_plugin_write_buffer_records( const rocprofiler_record_header_t* begin, const rocprofiler_record_header_t* end, rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { + std::lock_guard lock(att_plugin_t::writing_lock); if (!att_plugin || !att_plugin->IsValid()) return ROCPROFILER_STATUS_ERROR; return att_plugin->WriteBufferRecords(begin, end, session_id, buffer_id); } ROCPROFILER_EXPORT int rocprofiler_plugin_write_record(rocprofiler_record_tracer_t record) { + std::lock_guard lock(att_plugin_t::writing_lock); if (!att_plugin || !att_plugin->IsValid()) return ROCPROFILER_STATUS_ERROR; if (record.header.id.handle == 0) return ROCPROFILER_STATUS_SUCCESS; return ROCPROFILER_STATUS_SUCCESS; diff --git a/plugin/att/att.py b/plugin/att/att.py index 569c504137..2c42dbdf5f 100755 --- a/plugin/att/att.py +++ b/plugin/att/att.py @@ -17,6 +17,7 @@ import numpy as np from stitch import stitch import gc from collections import defaultdict +from service import CodeobjService ATT_VERSION = 2 @@ -189,15 +190,10 @@ def parse_binary(filename, kernel=None): info = SO.wrapped_parse_binary(str(filename).encode("utf-8"), kernel) code = [] - kernel_addr = defaultdict(lambda : "Unknown") - last_known_function = "Unknown" for k in range(info.code_len): code_entry = info.code[k] line = deepcopy(code_entry.line.decode("utf-8")) - if "; Begin " in line: - last_known_function = line.split("; Begin ")[1] - loc = deepcopy(code_entry.loc.decode("utf-8")) to_line = int(code_entry.to_line) if (code_entry.to_line >= 0) else None @@ -207,14 +203,11 @@ def parse_binary(filename, kernel=None): code.append([line, int(code_entry.value), to_line, loc, int(code_entry.index), int(code_entry.line_num), int(code_entry.addr), 0, 0]) - if code[-1][-3] != 0 and len(code) > 1: - kernel_addr[code[-1][-3]] = last_known_function - jumps = {} for k in range(info.jumps_len): jumps[info.jumps[k].key] = info.jumps[k].value - return code, jumps, kernel_addr + return code, jumps def getWaves_binary(name): @@ -248,13 +241,17 @@ def getWaves_binary(name): return (traces_python, waves_python, events, occupancy, flags, kernel_addr, info.flags & 0x4) -def getWaves_stitch(traces, code, jumps, flags, latency_map, hitcount_map, bIsAuto): +def getWaves_stitch(traces, code, jumps, flags, latency_map, hitcount_map, bIsAuto, codeobjservice): for id in traces.keys(): - traces[id].instructions = stitch(traces[id].instructions, code, jumps, flags, bIsAuto) + traces[id].instructions = stitch(traces[id].instructions, code, jumps, flags, bIsAuto, codeobjservice) + if len(code) > hitcount_map.size: + hitcount_map = np.pad(hitcount_map, [0,len(code)-hitcount_map.size]) + latency_map = np.pad(latency_map, [0,len(code)-latency_map.size]) if traces[id].instructions is not None: for inst in traces[id].instructions[0]: hitcount_map[inst.asmline] += inst.num_waves latency_map[inst.asmline] += inst.cycles + return hitcount_map, latency_map def persist(trace_file, SIMD, traces): @@ -301,7 +298,7 @@ def persist(trace_file, SIMD, traces): skips = traces[wave.traceid].instructions[-1] try: for v in traces[wave.traceid].instructions[0]: - if cc in skips: + while cc in skips: cc += 1 t = wave.instructions[cc] insts.append((t[0], v.type, 0, t[1], v.asmline)) @@ -434,7 +431,7 @@ if __name__ == "__main__": "--trace_file", help="Filter for trace files", default=None, type=str ) parser.add_argument( - "--att_kernel", help="Kernel file", type=str, default=pathenv + "/*_kernel.txt" + "--att_kernel", help="Kernel file", type=str, default=os.path.join(pathenv, "*_kernel.txt") ) parser.add_argument("--ports", help="Server and websocket ports, default: 8000,18000") parser.add_argument( @@ -480,17 +477,28 @@ if __name__ == "__main__": for att_kernel in att_kernel_list: print('Parsing:', att_kernel) + att_kernel_f = [] + with open(att_kernel, 'r') as f: + for line in f: + att_kernel_f.append(line.split('\n')[0]) + + get_path_loc = lambda x: np.sum([len(m) for m in x.split(' ')[:3]])+3 + add_pathnv = lambda x: x[:get_path_loc(x)] + os.path.join(pathenv, x[get_path_loc(x):]) + + att_kernel_f = [add_pathnv(p) if '.out' == p[-4:] else p for p in att_kernel_f[1:]] assembly_code = deepcopy(args.assembly_code) # Assembly parsing bIsAuto = False if assembly_code.lower().strip() == 'auto': - assembly_code = att_kernel.split('_kernel.txt')[0]+'_isa.s' bIsAuto = True - path = Path(assembly_code) - if not path.is_file(): + code = [['; Begin ATT ASM', 100, 0, '', 0, 0, 0, 0, 0]] + jumps = [] + elif not Path(assembly_code).is_file(): print("Invalid assembly_code('{0}')!".format(assembly_code)) sys.exit(1) + else: + code, jumps = parse_binary(assembly_code, att_kernel) # Trace Parsing trace_instance_name = att_kernel.split("_kernel.txt")[0] @@ -504,7 +512,6 @@ if __name__ == "__main__": continue print('Att kernel:', att_kernel) - code, jumps, kern_addr = parse_binary(assembly_code, None if bIsAuto else att_kernel) DBFILES = [] EVENTS = [] @@ -512,12 +519,16 @@ if __name__ == "__main__": GFXV = [] analysed_filenames = [] occupancy_filenames = [] - dispatch_kernel_names = {} + kernel_addr = {} latency_map = np.zeros((len(code)), dtype=np.int64) hitcount_map = np.zeros((len(code)), dtype=np.int32) gc.collect() + if bIsAuto: + codeservice = CodeobjService(att_kernel_f, SO.classify_asm_line) + else: + codeservice = None for name in filenames: traces, waves, perfevents, occupancy, gfxv, addrs, ftrace = getWaves_binary(name) @@ -528,7 +539,7 @@ if __name__ == "__main__": assert False for id, addr in enumerate(addrs): - dispatch_kernel_names[id] = kern_addr[addr] + kernel_addr[id] = addr if len(occupancy) > 16: OCCUPANCY.append( occupancy ) occupancy_filenames.append(name) @@ -537,7 +548,7 @@ if __name__ == "__main__": print("No traces from", name) continue - getWaves_stitch(traces, code, jumps, gfxv, latency_map, hitcount_map, bIsAuto) + hitcount_map, latency_map = getWaves_stitch(traces, code, jumps, gfxv, latency_map, hitcount_map, bIsAuto, codeservice) analysed_filenames.append(name) EVENTS.append(perfevents) @@ -559,7 +570,7 @@ if __name__ == "__main__": "EVENT_NAMES": EVENT_NAMES, "OCCUPANCY": OCCUPANCY, "ShaderNames": occupancy_filenames, - "DispatchNames": dispatch_kernel_names, + "DispatchNames": {id: codeservice.getSymbolName(addr) for id, addr in kernel_addr.items()} } view_trace( args, diff --git a/plugin/att/code_printing.cpp b/plugin/att/code_printing.cpp index 4bb509f890..dd089d2837 100644 --- a/plugin/att/code_printing.cpp +++ b/plugin/att/code_printing.cpp @@ -49,6 +49,26 @@ #include #include +#include + +#define C_API_BEGIN try { + +#define C_API_END(returndata) \ +} catch (std::exception& e) \ +{ \ + std::cerr << "Error: " << e.what() << std::endl; \ + return returndata; \ +} \ +catch (std::string& s) \ +{ \ + std::cerr << "Error: " << s << std::endl; \ + return returndata; \ +} \ +catch (...) \ +{ \ + return returndata; \ +} + code_object_decoder_t::code_object_decoder_t(const char* codeobj_data, uint64_t codeobj_size) { buffer = std::vector{}; buffer.resize(codeobj_size); @@ -110,8 +130,16 @@ code_object_decoder_t::code_object_decoder_t(const char* codeobj_data, uint64_t // load_symbol_map(); } - disassembly = std::make_unique(*this); - m_symbol_map = disassembly->GetKernelMap(); + try { + disassembly = std::make_unique(*this); // Can throw + } catch(std::exception& e) { + return; + } + try { + m_symbol_map = disassembly->GetKernelMap(); // Can throw + } catch(std::exception& e) { + return; + } //disassemble_kernels(); } @@ -143,7 +171,29 @@ std::optional code_object_decoder_t::find_symbol(uint64_t vaddr) { return SymbolInfo{symbol_name, symbol.faddr, symbol.mem_size}; } -void code_object_decoder_t::disassemble_kernel(uint64_t faddr, uint64_t vaddr) { +std::pair +code_object_decoder_t::disassemble_instruction(uint64_t faddr, uint64_t vaddr) +{ + if (!disassembly) + throw std::exception(); + + char* cpp_line = nullptr; + auto it = m_line_number_map.find(vaddr); + if (it != m_line_number_map.end()) { + const std::string& file_name = it->second.first; + size_t line_number = it->second.second; + + std::string cpp = file_name + ':' + std::to_string(line_number); + cpp_line = (char*)calloc(cpp.size() + 4, sizeof(char)); + std::memcpy(cpp_line, cpp.data(), cpp.size() * sizeof(char)); + } + size_t size = disassembly->ReadInstruction(faddr, vaddr, cpp_line); + return {disassembly->last_instruction, size}; +} + +void code_object_decoder_t::disassemble_kernel(uint64_t faddr, uint64_t vaddr) +{ + if (!disassembly) return; auto symbol = find_symbol(vaddr); if (!symbol) @@ -155,19 +205,12 @@ void code_object_decoder_t::disassemble_kernel(uint64_t faddr, uint64_t vaddr) { std::cout << "Dumping ISA for " << symbol->name << std::endl; uint64_t end_addr = faddr + symbol->mem_size; - while (faddr < end_addr) { - char* cpp_line = nullptr; - auto it = m_line_number_map.find(vaddr); - if (it != m_line_number_map.end()) { - const std::string& file_name = it->second.first; - size_t line_number = it->second.second; - - std::string cpp = file_name + ':' + std::to_string(line_number); - cpp_line = (char*)calloc(cpp.size() + 4, sizeof(char)); - std::memcpy(cpp_line, cpp.data(), cpp.size() * sizeof(char)); - } - - size_t size = disassembly->ReadInstruction(faddr, vaddr, cpp_line); + while (faddr < end_addr) + { + size_t size; + instruction_instance_t inst; + std::tie(inst, size) = this->disassemble_instruction(faddr, vaddr); + instructions.push_back(inst); faddr += size; vaddr += size; } @@ -182,3 +225,136 @@ void code_object_decoder_t::disassemble_single_kernel(uint64_t kaddr) { if (kaddr >= vaddr && kaddr < vaddr + v.mem_size) disassemble_kernel(v.faddr, vaddr); } + +CodeobjService::CodeobjService(const char* filepath, uint64_t load_base): load_base(load_base) +{ + if (!filepath) throw "Empty filepath."; + + std::string_view fpath(filepath); + + if (fpath.rfind(".out") + 4 == fpath.size()) + { + std::ifstream file(filepath, std::ios::in | std::ios::binary); + + if (!file.is_open()) + throw "Invalid filename " + std::string(filepath); + + std::vector buffer; + file.seekg(0, file.end); + buffer.resize(file.tellg()); + file.seekg(0, file.beg); + file.read(buffer.data(), buffer.size()); + + decoder = std::make_unique(buffer.data(), buffer.size()); + } + else + { + std::unique_ptr binary = std::make_unique(filepath); + decoder = std::make_unique(binary->buffer.data(), binary->buffer.size()); + } +} + +bool CodeobjService::decode_single(uint64_t vaddr, uint64_t faddr) +{ + if (!decoder->disassembly) return false; + + try + { + decoded_map[vaddr] = decoder->disassemble_instruction(faddr, vaddr-load_base); + } + catch(std::exception& e) + { + return false; + } + return true; +} + +std::pair& CodeobjService::getDecoded(uint64_t addr) +{ + if (decoded_map.find(addr) != decoded_map.end()) + return decoded_map[addr]; + + std::optional faddr{}; + + if (!bNotElfFILE) + { + faddr = DisassemblyInstance::va2fo(decoder->buffer.data(), addr-load_base); + if (!faddr) + bNotElfFILE = true; + } + + if (bNotElfFILE && decoder->buffer.size() > 0x100) { + uint64_t f_offset = *reinterpret_cast(decoder->buffer.data()+0xb8); + uint64_t v_offset = *reinterpret_cast(decoder->buffer.data()+0xc8); + + faddr = addr+f_offset-load_base-v_offset; + } + + if (!faddr || !decode_single(addr, *faddr)) + { + std::cerr << "Invalid addr: " << std::hex << addr << std::dec << std::endl; + throw std::exception(); + } + + return decoded_map[addr]; +} + +std::unordered_map> services{}; +std::atomic shandles{1}; + +#define PUBLIC_API __attribute__((visibility("default"))) + +extern "C" +{ + PUBLIC_API uint64_t createService(const char* filename, uint64_t load_base) + { + C_API_BEGIN + + uint64_t handle = shandles.fetch_add(1); + services[handle] = std::make_unique(filename, load_base); + return handle; + + C_API_END(0) + } + PUBLIC_API int deleteService(uint64_t handle) + { + return services.erase(handle); + } + PUBLIC_API const char* getInstruction(uint64_t handle, uint64_t addr) + { + C_API_BEGIN + + return services.at(handle)->getInstruction(addr); + + C_API_END(nullptr) + } + PUBLIC_API const char* getCppref(uint64_t handle, uint64_t addr) + { + C_API_BEGIN + + return services.at(handle)->getCppref(addr); + + C_API_END(nullptr) + } + PUBLIC_API size_t getInstSize(uint64_t handle, uint64_t addr) + { + C_API_BEGIN + + return services.at(handle)->getSize(addr); + + C_API_END(0) + } + PUBLIC_API const char* getSymbolName(uint64_t addr) + { + C_API_BEGIN + + for (auto& [handle, service] : services) + { + if (!service->inrange(addr)) continue; + return service->getSymbolName(addr); + } + return nullptr; + + C_API_END(nullptr) + } +} \ No newline at end of file diff --git a/plugin/att/code_printing.hpp b/plugin/att/code_printing.hpp index 02df1deee9..17c9a7463f 100644 --- a/plugin/att/code_printing.hpp +++ b/plugin/att/code_printing.hpp @@ -27,28 +27,67 @@ #include #include #include +#include #include "disassembly.hpp" class code_object_decoder_t { - public: - // void load_symbol_map(); +public: std::optional find_symbol(uint64_t address); code_object_decoder_t(const char* codeobj_data, uint64_t codeobj_size); ~code_object_decoder_t(); + std::pair + disassemble_instruction(uint64_t faddr, uint64_t vaddr); void disassemble_kernel(uint64_t faddr, uint64_t vaddr); void disassemble_single_kernel(uint64_t kaddr); void disassemble_kernels(); int m_fd; - std::map> m_line_number_map; - std::map m_symbol_map; + std::map> m_line_number_map{}; + std::map m_symbol_map{}; std::string m_uri; std::vector buffer; - std::vector instructions; - std::unique_ptr disassembly; + std::vector instructions{}; + std::unique_ptr disassembly{}; +}; + +class CodeobjService +{ +public: + CodeobjService(const char* filepath, uint64_t load_base); + bool decode_single(uint64_t vaddr, uint64_t faddr); + + std::pair& getDecoded(uint64_t addr); + const char* getInstruction(uint64_t addr) { return getDecoded(addr).first.instruction; } + const char* getCppref(uint64_t addr) { return getDecoded(addr).first.cpp_reference; } + size_t getSize(uint64_t addr) { return getDecoded(addr).second; } + + uint64_t size() const { + if (!decoder) return 0; + return decoder->buffer.size(); + } + uint64_t begin() const { return load_base; }; + uint64_t end() const { return begin() + size(); } + bool inrange(uint64_t addr) const { return addr >= begin() && addr < end(); } + + const char* getSymbolName(uint64_t addr) const { + if (!decoder) return nullptr; + + auto it = decoder->m_symbol_map.find(addr-load_base); + if (it != decoder->m_symbol_map.end()) + return it->second.name.data(); + + return nullptr; + } +private: + const uint64_t load_base; + + std::unordered_map> decoded_map; + std::unique_ptr decoder{nullptr}; + + bool bNotElfFILE = false; }; diff --git a/plugin/att/disassembly.cpp b/plugin/att/disassembly.cpp index 0bb358821e..b7be35de96 100644 --- a/plugin/att/disassembly.cpp +++ b/plugin/att/disassembly.cpp @@ -52,14 +52,23 @@ #include "code_printing.hpp" #include -#define CHECK_COMGR(call) \ - if (amd_comgr_status_s status = call) { \ - const char* reason = ""; \ - amd_comgr_status_string(status, &reason); \ - std::cerr << __LINE__ << " code: " << status << std::endl; \ - std::cerr << __LINE__ << " failed: " << reason << std::endl; \ - exit(1); \ - } +#define THROW_COMGR(call) \ +if (amd_comgr_status_s status = call) { \ + const char* reason = ""; \ + amd_comgr_status_string(status, &reason); \ + std::cerr << __FILE__ << ':' << __LINE__ << " code: " \ + << status << " failed: " << reason << std::endl; \ + throw std::exception(); \ +} + +#define RETURN_COMGR(call) \ +if (amd_comgr_status_s status = call) { \ + const char* reason = ""; \ + amd_comgr_status_string(status, &reason); \ + std::cerr << __FILE__ << ':' << __LINE__ << " code: " \ + << status << " failed: " << reason << std::endl; \ + return AMD_COMGR_STATUS_ERROR; \ +} CodeObjectBinary::CodeObjectBinary(const std::string& uri) : m_uri(uri) { const std::string protocol_delim{"://"}; @@ -83,12 +92,14 @@ CodeObjectBinary::CodeObjectBinary(const std::string& uri) : m_uri(uri) { std::string decoded_path; decoded_path.reserve(path.length()); for (size_t i = 0; i < path.length(); ++i) + { if (path[i] == '%' && std::isxdigit(path[i + 1]) && std::isxdigit(path[i + 2])) { decoded_path += std::stoi(path.substr(i + 1, 2), 0, 16); i += 2; } else { decoded_path += path[i]; } + } /* Tokenize the query/fragment. */ std::vector tokens; @@ -122,26 +133,19 @@ CodeObjectBinary::CodeObjectBinary(const std::string& uri) : m_uri(uri) { if (!(size = std::stoul(size_it->second, nullptr, 0))) return; } - if (protocol != "file") { - printf("\"%s\" protocol not supported\n", protocol.c_str()); - return; - } + if (protocol != "file") throw protocol + " protocol not supported!"; std::ifstream file(decoded_path, std::ios::in | std::ios::binary); - if (!file) { - printf("could not open `%s'\n", decoded_path.c_str()); - return; - } + if (!file || !file.is_open()) throw "could not open " + decoded_path; if (!size) { file.ignore(std::numeric_limits::max()); size_t bytes = file.gcount(); file.clear(); - if (bytes < offset) { - printf("invalid uri `%s' (file size < offset)\n", decoded_path.c_str()); - return; - } + if (bytes < offset) + throw "invalid uri " + decoded_path + " (file size < offset)"; + size = bytes - offset; } @@ -154,17 +158,20 @@ CodeObjectBinary::CodeObjectBinary(const std::string& uri) : m_uri(uri) { DisassemblyInstance::DisassemblyInstance(code_object_decoder_t& decoder) : buffer(reinterpret_cast(decoder.buffer.data())), - size(decoder.buffer.size()), - instructions(decoder.instructions) { - CHECK_COMGR(amd_comgr_create_data(AMD_COMGR_DATA_KIND_EXECUTABLE, &data)); - CHECK_COMGR(amd_comgr_set_data(data, size, decoder.buffer.data())); + size(decoder.buffer.size()) +{ + THROW_COMGR(amd_comgr_create_data(AMD_COMGR_DATA_KIND_EXECUTABLE, &data)); + THROW_COMGR(amd_comgr_set_data(data, size, decoder.buffer.data())); + /*std::cout << "checking isa" << std::endl; char isa_name[128]; size_t isa_size = sizeof(isa_name); CHECK_COMGR(amd_comgr_get_data_isa_name(data, &isa_size, isa_name)); + std::cout << isa_name << std::endl; */ + const char* isa_name = "amdgcn-amd-amdhsa--gfx1100"; - CHECK_COMGR(amd_comgr_create_disassembly_info( - isa_name, //"amdgcn-amd-amdhsa--gfx1100", + THROW_COMGR(amd_comgr_create_disassembly_info( + isa_name, &DisassemblyInstance::memory_callback, &DisassemblyInstance::inst_callback, [](uint64_t address, void* user_data) {}, &info)); } @@ -183,7 +190,7 @@ static bool IsKernelType(amd_comgr_symbol_type_t type) amd_comgr_status_t DisassemblyInstance::symbol_callback(amd_comgr_symbol_t symbol, void* user_data) { amd_comgr_symbol_type_t type; - CHECK_COMGR(amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_TYPE, &type)); + RETURN_COMGR(amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_TYPE, &type)); if (!IsKernelType(type)) return AMD_COMGR_STATUS_SUCCESS; @@ -191,14 +198,14 @@ amd_comgr_status_t DisassemblyInstance::symbol_callback(amd_comgr_symbol_t symbo uint64_t vaddr; uint64_t mem_size; uint64_t name_size; - CHECK_COMGR(amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_VALUE, &vaddr)); - CHECK_COMGR(amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_SIZE, &mem_size)); - CHECK_COMGR(amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_NAME_LENGTH, &name_size)); + RETURN_COMGR(amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_VALUE, &vaddr)); + RETURN_COMGR(amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_SIZE, &mem_size)); + RETURN_COMGR(amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_NAME_LENGTH, &name_size)); std::string name; name.resize(name_size); - CHECK_COMGR(amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_NAME, name.data())); + RETURN_COMGR(amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_NAME, name.data())); DisassemblyInstance& instance = *static_cast(user_data); std::optional faddr = va2fo(instance.buffer, vaddr); @@ -210,23 +217,25 @@ amd_comgr_status_t DisassemblyInstance::symbol_callback(amd_comgr_symbol_t symbo std::map& DisassemblyInstance::GetKernelMap() { symbol_map = {}; - CHECK_COMGR(amd_comgr_iterate_symbols(data, &DisassemblyInstance::symbol_callback, this)); + THROW_COMGR(amd_comgr_iterate_symbols(data, &DisassemblyInstance::symbol_callback, this)); + return symbol_map; } DisassemblyInstance::~DisassemblyInstance() { - CHECK_COMGR(amd_comgr_release_data(data)); - CHECK_COMGR(amd_comgr_destroy_disassembly_info(info)); + amd_comgr_release_data(data); + amd_comgr_destroy_disassembly_info(info); } uint64_t DisassemblyInstance::ReadInstruction(uint64_t faddr, uint64_t vaddr, const char* cpp_line) { uint64_t size_read; uint64_t addr_in_buffer = reinterpret_cast(buffer) + faddr; - CHECK_COMGR(amd_comgr_disassemble_instruction(info, addr_in_buffer, (void*)this, &size_read)); - assert(instructions.size() != 0); - instructions.back().address = vaddr; - instructions.back().cpp_reference = cpp_line; + + THROW_COMGR(amd_comgr_disassemble_instruction(info, addr_in_buffer, (void*)this, &size_read)); + + last_instruction.address = vaddr; + last_instruction.cpp_reference = cpp_line; return size_read; } @@ -241,7 +250,7 @@ uint64_t DisassemblyInstance::memory_callback(uint64_t from, char* to, uint64_t void DisassemblyInstance::inst_callback(const char* instruction, void* user_data) { DisassemblyInstance& instance = *static_cast(user_data); - instance.instructions.push_back({strdup(instruction), nullptr, 0}); + instance.last_instruction.instruction = strdup(instruction); } #define CHECK_VA2FO(x, msg) if (!(x)) { \ diff --git a/plugin/att/disassembly.hpp b/plugin/att/disassembly.hpp index d686c3dca2..98019c0420 100644 --- a/plugin/att/disassembly.hpp +++ b/plugin/att/disassembly.hpp @@ -61,7 +61,7 @@ class DisassemblyInstance { void* buffer; int64_t size; - std::vector& instructions; + instruction_instance_t last_instruction; amd_comgr_disassembly_info_t info; amd_comgr_data_t data; std::map symbol_map; diff --git a/plugin/att/service.py b/plugin/att/service.py new file mode 100644 index 0000000000..bafc19b73b --- /dev/null +++ b/plugin/att/service.py @@ -0,0 +1,95 @@ +import ctypes +from ctypes import * + +pluginpath = '/home/giovanni/Desktop/rocprofiler/build/lib/rocprofiler/libatt_plugin.so' + +attplugin = ctypes.CDLL(pluginpath) + +attplugin.createService.restype = ctypes.c_uint64 +attplugin.createService.argtypes = [ctypes.c_char_p, ctypes.c_uint64] +attplugin.deleteService.restype = ctypes.c_int +attplugin.deleteService.argtypes = [ctypes.c_uint64] +attplugin.getInstruction.restype = ctypes.c_char_p +attplugin.getInstruction.argtypes = [ctypes.c_uint64, ctypes.c_uint64] +attplugin.getCppref.restype = ctypes.c_char_p +attplugin.getCppref.argtypes = [ctypes.c_uint64, ctypes.c_uint64] +attplugin.getInstSize.restype = ctypes.c_size_t +attplugin.getInstSize.argtypes = [ctypes.c_uint64, ctypes.c_uint64] +attplugin.getSymbolName.restype = ctypes.c_char_p +attplugin.getSymbolName.argtypes = [ctypes.c_uint64] + + +class CodeobjInstance: + def __init__(self, line, classification_func): + tokens = line.split(' ') + self.load_base = int(tokens[0], 16) + self.load_end = self.load_base + int(tokens[1], 16) + self.att_marker = int(tokens[2]) + self.fpath = tokens[3] + + self.handle = attplugin.createService(self.fpath.encode('utf-8'), self.load_base) + self.classifier = classification_func + + if self.handle == 0: + print('Warning: Could not open', line) + raise + + def release(self): + attplugin.deleteService(self.handle) + + def inrange(self, addr): + return addr >= self.load_base and addr < self.load_end+0x1000 + + def GetInstruction(self, addr): + inst = attplugin.getInstruction(self.handle, addr) + if inst is None: + return None + inst = inst.decode() + while len(inst) and (inst[0] == '\t' or inst[0] == ' '): + inst = inst[1:] + while len(inst) and (inst[-1] == '\t' or inst[-1] == ' '): + inst = inst[:-1] + cpp = attplugin.getCppref(self.handle, addr) + if cpp: + cpp = cpp.decode() + size = attplugin.getInstSize(self.handle, addr) + if size and inst: + return (self.classifier(inst.encode('utf-8'), len(inst)), inst, cpp, size) + return None + + +class CodeobjService: + def __init__(self, att_kernel_txt, cfunc) -> None: + cfunc.restype = ctypes.c_int + cfunc.argtypes = [ctypes.c_char_p, ctypes.c_size_t] + + self.last_instance = None + self.services = [] + for line in att_kernel_txt: + try: + if 'memory://' == line[0:len('memory://')]: + continue + self.services.append(CodeobjInstance(line, cfunc)) + except: + pass + + def release(self): + for _, _, instance in self.services: + instance.release() + + def GetInstruction(self, addr): + if self.last_instance and self.last_instance.inrange(addr): + return self.last_instance.GetInstruction(addr) + + for instance in self.services: + if instance.inrange(addr): + self.last_instance = instance + return instance.GetInstruction(addr) + + return None + + def getSymbolName(self, addr): + name = attplugin.getSymbolName(addr) + if name: + return name.decode() + return "Addr #"+hex(addr) diff --git a/plugin/att/stitch.py b/plugin/att/stitch.py index 48715a4ad8..c70337cbb4 100644 --- a/plugin/att/stitch.py +++ b/plugin/att/stitch.py @@ -53,12 +53,25 @@ WaveInstCategory = { # Keeps track of register states for hipcc-generated assembly class RegisterWatchList: - def __init__(self, labels): + def __init__(self, labels, code, jump_map, insts): self.registers = {"v" + str(k): [[] for m in range(64)] for k in range(64)} for k in range(128): self.registers["s" + str(k)] = [] self.labels = labels + self.code = code + self.jump_map = jump_map + self.insts = insts + + def jump(self, as_line): + return self.jump_map[as_line[2]] + + def getcode(self, line): + return self.code[line], 1 + + def getincrement(self, line): + return 1 + def try_translate(self, tok): if tok[0] in ["s"]: return self.registers[self.range(tok)[0]] @@ -153,12 +166,64 @@ class RegisterWatchList: except: pass + # Matches tokens in reverse order + def try_match_swapped(self, i, line, increment): + return self.insts[i + 1].type == self.code[line][1] and self.insts[i].type == self.code[line + 1][1] + # Translates PC values to instructions, for auto captured ISA class PCTranslator: - def __init__(self, code, insts): + def __init__(self, insts, code, raw_code, reverse_map, codeservice): + self.codeservice = codeservice + + self.insts = insts + self.addrmap = {c[-3] : (c, self.codeservice.GetInstruction(c[-3])[3]) for c in code if c[-3] > 0} + self.code = code - self.insts = insts[1:] - self.addrmap = {code[m][-3] : m for m in range(len(code))} + self.raw_code = raw_code + self.reverse_map = reverse_map + self.jump_map = {c[-3] : self.getjump_loc(c) for c in code if c[1] == BRANCH} + + def jump(self, as_line): + return self.jump_map[as_line[-3]] + + def getcode(self, addr): + try: + return self.addrmap[addr] + except Exception as ex: + new_inst = self.codeservice.GetInstruction(addr) + if new_inst and new_inst[3]: # Check returned size > 0 + last_line = self.raw_code[-1] + newline = [new_inst[1], new_inst[0], len(self.raw_code), new_inst[2], last_line[4]+1, last_line[5]+1, addr, 0, 0] + if new_inst[0] == BRANCH: + self.jump_map[addr] = self.getjump_loc(newline) + self.addrmap[addr] = (newline, new_inst[3]) + + next = len(self.code) + self.reverse_map[addr] = len(self.raw_code) + self.raw_code.append(newline) + self.code.append(newline) + return newline, new_inst[3] + else: + raise ex + + def jump(self, asm_line): + try: + return self.jump_map[asm_line[-3]] + except: + loc = self.getjump_loc(asm_line) + self.jump_map[asm_line[-3]] = loc + return loc + + def getjump_loc(self, asm_line): + try: + dest = int(asm_line[0].split(' ')[-1]) + if dest >= 32768: dest -= 65536 + return asm_line[-3] + 4*dest+4 + except: + return -1 + + def getincrement(self, addr): + return self.getcode(addr)[1] def try_translate(self, tok): pass @@ -170,15 +235,13 @@ class PCTranslator: pass def swappc(self, line, line_num, inst_index): try: - loc = self.addrmap[self.insts[inst_index+1].cycles] - return loc + return self.getcode(self.insts[inst_index+1].cycles)[0][-3] except: print('SWAPPC warning: Could not find addr', hex(self.insts[inst_index+1].cycles), 'for', inst_index, line) return -1 def setpc(self, line, inst_index): try: - loc = self.addrmap[self.insts[inst_index+1].cycles] - return loc + return self.getcode(self.insts[inst_index+1].cycles)[0][-3] except: print('SETPC warning: Could not find addr', hex(self.insts[inst_index+1].cycles), 'for', inst_index, line) return -1 @@ -189,15 +252,19 @@ class PCTranslator: def updatelane(self, line): pass -# Matches tokens in reverse order -def try_match_swapped(insts, code, i, line): - return insts[i + 1].type == code[line][1] and insts[i].type == code[line + 1][1] + # Matches tokens in reverse order + def try_match_swapped(self, i, addr, increment): + try: + return self.insts[i + 1].type == self.getcode(addr)[0][1] and \ + self.insts[i].type == self.getcode(addr + increment)[0][1] + except Exception as e: + return False -def stitch(insts, raw_code, jumps, gfxv, bIsAuto): +def stitch(insts, raw_code, jumps, gfxv, bIsAuto, codeservice): bGFX9 = gfxv == 'vega' - result, i, line, loopCount = [], 0, 0, defaultdict(int) + result, i, loopCount = [], 0, defaultdict(int) SMEM_INST = [] # scalar memory VLMEM_INST = [] # vector memory load @@ -228,8 +295,15 @@ def stitch(insts, raw_code, jumps, gfxv, bIsAuto): labels[c[0].split(":")[0]] = len(code) reverse_map = {} - for k, v in enumerate(jump_map): - reverse_map[v] = k + if bIsAuto: + for k, v in enumerate(jump_map): + try: + reverse_map[code[v][-3]] = k + except: + pass + else: + for k, v in enumerate(jump_map): + reverse_map[v] = k jumps = {jump_map[j] + 1: j for j in jumps} @@ -242,42 +316,51 @@ def stitch(insts, raw_code, jumps, gfxv, bIsAuto): loops = 0 maxline = 0 + pcskip = [] if bIsAuto: try: - if insts[0].type != PCINFO: + firstinst = insts[0] + insts = insts[1:] + + if firstinst.type != PCINFO: print('Warning: Waves without PCINFO') return None - elif insts[0].cycles == 0: + elif firstinst.cycles == 0: print('Info: Some waves started before the trace') return None - watchlist = PCTranslator(code, insts) - line = watchlist.addrmap[insts[0].cycles] - except Exception as e: - print(e) + + watchlist = PCTranslator(insts, code, raw_code, reverse_map, codeservice) + line = firstinst.cycles + lineincrement = watchlist.getincrement(line) + pcskip.append(0) + except KeyError as e: + print('Auto error invalid addr', hex(e.args[0])) + return None + except Exception as e: + print('Auto error', e) return None - insts = insts[1:] else: - watchlist = RegisterWatchList(labels=labels) + line = 0 + lineincrement = 1 + watchlist = RegisterWatchList(labels=labels, code=code, jump_map=jump_map) N = len(insts) - pcskip = [] - while i < N: + while i < N and line >= 0 and loops < MAX_STITCHED_TOKENS: if insts[i].type == PCINFO: i += 1 + pcskip.append(i) continue - loops += 1 - if line >= len(code) or loops > MAX_STITCHED_TOKENS \ - or num_failed_stitches > MAX_FAILED_STITCHES: + + inst = insts[i] + try: + as_line, lineincrement = watchlist.getcode(line) + except: break - maxline = max(reverse_map[line], maxline) - inst = insts[i] - as_line = code[line] - matched = True - next = line + 1 + next = line + lineincrement if not bIsAuto: if '_mov_' in as_line[0]: @@ -287,7 +370,7 @@ def stitch(insts, raw_code, jumps, gfxv, bIsAuto): if as_line[1] == GETPC: try: - watchlist.getpc(as_line[0], code[line+1][0]) + watchlist.getpc(as_line[0], watchlist.getcode(next)[0]) matched = inst.type in [SALU, JUMP] except: matched = False @@ -298,31 +381,31 @@ def stitch(insts, raw_code, jumps, gfxv, bIsAuto): next = watchlist.setpc(as_line[0], i) matched = inst.type in [SALU, JUMP] if bIsAuto: - pcskip.append(i) i += 1 + pcskip.append(i+1) while next < 0 and i+1 < len(insts): - if insts[i+1].type == PCINFO: - next = watchlist.setpc(as_line[0], i) - pcskip.append(i) - else: - inst.cycles += insts[i+1].cycles i += 1 + if insts[i].type == PCINFO: + next = watchlist.setpc(as_line[0], i-1) + pcskip.append(i+1) + else: + inst.cycles += insts[i].cycles if next < 0: print('Jump to unknown location in line', as_line[0]) break elif as_line[1] == SWAPPC: - next = watchlist.swappc(as_line[0], line, i) matched = inst.type in [SALU, JUMP] + next = watchlist.swappc(as_line[0], line, i) if bIsAuto: - pcskip.append(i) i += 1 + pcskip.append(i+1) while next < 0 and i+1 < len(insts): - if insts[i+1].type == PCINFO: - next = watchlist.swappc(as_line[0], line, i) - pcskip.append(i) - else: - inst.cycles += insts[i+1].cycles i += 1 + if insts[i].type == PCINFO: + next = watchlist.swappc(as_line[0], line, i-1) + pcskip.append(i+1) + else: + inst.cycles += insts[i].cycles if next < 0: print('Jump to unknown location in line', as_line[0]) break @@ -421,28 +504,26 @@ def stitch(insts, raw_code, jumps, gfxv, bIsAuto): num_inflight = NUM_FLAT + NUM_SMEM + NUM_VLMEM + NUM_VSMEM elif inst.type == JUMP and as_line[1] == BRANCH: - next = jump_map[as_line[2]] + next = watchlist.jump(as_line) if next is None or next == 0: print("Jump to unknown location!", as_line) break elif inst.type == NEXT and as_line[1] == BRANCH: - next = line + 1 + pass else: matched = False - next = line + 1 - if i + 1 < N and line + 1 < len(code): - if try_match_swapped(insts, code, i, line): - temp = insts[i] - insts[i] = insts[i + 1] - insts[i + 1] = temp - next = line - elif "s_waitcnt " in as_line[0] or "_load_" in as_line[0]: - if skipped_immed > 0 and "s_waitcnt " in as_line[0]: - matched = True - skipped_immed -= 1 - elif 'scratch_' not in as_line[0]: - print('Parsing terminated at:', as_line) - break + if watchlist.try_match_swapped(i, line, lineincrement): + temp = insts[i] + insts[i] = insts[i + 1] + insts[i + 1] = temp + next = line + elif "s_waitcnt " in as_line[0] or "_load_" in as_line[0]: + if skipped_immed > 0 and "s_waitcnt " in as_line[0]: + matched = True + skipped_immed -= 1 + elif 'scratch_' not in as_line[0]: + print('Parsing terminated at:', as_line) + break if matched: inst.asmline = reverse_map[line] @@ -457,6 +538,8 @@ def stitch(insts, raw_code, jumps, gfxv, bIsAuto): i += 1 else: num_failed_stitches += 1 + + maxline = max(reverse_map[line], maxline) line = next N = max(N, 1) @@ -475,5 +558,6 @@ def stitch(insts, raw_code, jumps, gfxv, bIsAuto): ) break line += 1 + print('Sucessfuly parsed', i, 'tokens') return result, loopCount, mem_unroll, flight_count, maxline, len(result), pcskip diff --git a/plugin/exportmap b/plugin/exportmap index 3eb22197f7..1a8c2d0863 100644 --- a/plugin/exportmap +++ b/plugin/exportmap @@ -3,5 +3,11 @@ global: rocprofiler_plugin_initialize; rocprofiler_plugin_finalize; rocprofiler_plugin_write_buffer_records; rocprofiler_plugin_write_record; + createService; + deleteService; + getInstruction; + getCppref; + getInstSize; + getSymbolName; local: *; }; \ No newline at end of file diff --git a/src/api/CMakeLists.txt b/src/api/CMakeLists.txt index 360d8e503e..72db79591c 100644 --- a/src/api/CMakeLists.txt +++ b/src/api/CMakeLists.txt @@ -184,8 +184,8 @@ file(GLOB ROCPROFILER_TRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/*.cpp) file(GLOB ROCPROFILER_ROCTRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/src/*.cpp) -set(ROCPROFILER_ATT_SRC_FILES - ${PROJECT_SOURCE_DIR}/src/core/session/att/att.cpp) +file(GLOB ROCPROFILER_ATT_SRC_FILES + ${PROJECT_SOURCE_DIR}/src/core/session/att/*.cpp) file(GLOB ROCPROFILER_CLASS_SRC_FILES ${CMAKE_CURRENT_SOURCE_DIR}/rocprofiler_singleton.cpp) file(GLOB ROCPROFILER_SPM_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/spm/spm.cpp) diff --git a/src/core/hsa/hsa_support.cpp b/src/core/hsa/hsa_support.cpp index dc8c7517ee..fd0a828bf6 100644 --- a/src/core/hsa/hsa_support.cpp +++ b/src/core/hsa/hsa_support.cpp @@ -417,67 +417,84 @@ struct CodeObjectCallbackArg { bool unload; }; -hsa_status_t CodeObjectCallback(hsa_executable_t executable, - hsa_loaded_code_object_t loaded_code_object, void* arg) { - hsa_evt_data_t data{}; - rocprofiler::HSASupport_Singleton& hsasupport_singleton = - rocprofiler::HSASupport_Singleton::GetInstance(); - if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE, - &data.codeobj.storage_type) != HSA_STATUS_SUCCESS) +#define CheckInfo(x) if ((x) != HSA_STATUS_SUCCESS) \ rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed"); - if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_FILE) { - if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_FILE, - &data.codeobj.storage_file) != HSA_STATUS_SUCCESS || - data.codeobj.storage_file == -1) +hsa_status_t CodeObjectCallback(hsa_executable_t executable, + hsa_loaded_code_object_t loaded_code_object, void* arg) { + using namespace std::placeholders; + hsa_evt_data_t data{}; + + auto codeobj_info_func = rocprofiler::HSASupport_Singleton::GetInstance() + .GetHSALoaderApi() + .hsa_ven_amd_loader_loaded_code_object_get_info; + auto codeobj_bound = std::bind(codeobj_info_func, loaded_code_object, _1, _2); + + CheckInfo(codeobj_bound( + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE, + &data.codeobj.storage_type + )); + + if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_FILE) + { + CheckInfo(codeobj_bound( + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_FILE, + &data.codeobj.storage_file + )); + if (data.codeobj.storage_file == -1) rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed"); + data.codeobj.memory_base = data.codeobj.memory_size = 0; - } else if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY) { - if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, - HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_BASE, - &data.codeobj.memory_base) != HSA_STATUS_SUCCESS || - hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, - HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_SIZE, - &data.codeobj.memory_size) != HSA_STATUS_SUCCESS) - rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed"); + } + else if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY) + { + CheckInfo(codeobj_bound( + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_BASE, + &data.codeobj.memory_base + )); + CheckInfo(codeobj_bound( + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_SIZE, + &data.codeobj.memory_size + )); data.codeobj.storage_file = -1; - } else if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_NONE) { - return HSA_STATUS_SUCCESS; // FIXME: do we really not care about these - // code objects? - } else { + } + else if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_NONE) + { + return HSA_STATUS_SUCCESS; // FIXME: do we really not care about these code objects? + } + else + { rocprofiler::fatal("unknown code object storage type: %d", data.codeobj.storage_type); } - if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE, - &data.codeobj.load_base) != HSA_STATUS_SUCCESS || - hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE, - &data.codeobj.load_size) != HSA_STATUS_SUCCESS || - hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA, - &data.codeobj.load_delta) != HSA_STATUS_SUCCESS) - rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed"); + CheckInfo(codeobj_bound( + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE, + &data.codeobj.load_base + )); + CheckInfo(codeobj_bound( + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE, + &data.codeobj.load_size + )); + CheckInfo(codeobj_bound( + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA, + &data.codeobj.load_delta + )); - if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH, - &data.codeobj.uri_length) != HSA_STATUS_SUCCESS) - rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed"); + CheckInfo(codeobj_bound( + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH, + &data.codeobj.uri_length + )); std::string uri_str(data.codeobj.uri_length, '\0'); - if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI, uri_str.data()) != - HSA_STATUS_SUCCESS) - rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed"); - if (hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info( - loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_AGENT, - &data.codeobj.agent) != HSA_STATUS_SUCCESS) - rocprofiler::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed"); + CheckInfo(codeobj_bound( + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI, + uri_str.data() + )); + CheckInfo(codeobj_bound( + HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_AGENT, + &data.codeobj.agent + )); data.codeobj.uri = uri_str.c_str(); data.codeobj.unload = *static_cast(arg) ? 1 : 0; @@ -486,8 +503,8 @@ hsa_status_t CodeObjectCallback(hsa_executable_t executable, if (data.codeobj.unload) codeobj_capture_instance::Unload(data.codeobj.load_base); else - codeobj_capture_instance::Load(data.codeobj.load_base, uri_str, data.codeobj.memory_base, - data.codeobj.memory_size); + codeobj_capture_instance::Load(data.codeobj.load_base, + uri_str, data.codeobj.memory_base, data.codeobj.memory_size); hsa_executable_iterate_agent_symbols(executable, data.codeobj.agent, hsa_executable_iteration_callback, &(data.codeobj.unload)); @@ -783,6 +800,7 @@ const Agent::DeviceInfo& HSAAgentInfo::GetDeviceInfo() const { return device_info_; } assert("Attempting to read deviceInfo for a CPU agent"); + return device_info_; } uint64_t HSAAgentInfo::getHandle() const { return agent_.handle; } @@ -1037,7 +1055,6 @@ void HSASupport_Singleton::HSAInitialize(HsaApiTable* table) { // is currently doing as well as the roctracer compatibility // code earlier in this file. uint32_t gpu_id = 0; - hsasupport_singleton.GetCoreApiTable().hsa_agent_get_info_fn( agent, (hsa_agent_info_t)(HSA_AMD_AGENT_INFO_DRIVER_UID), &gpu_id); const Agent::DeviceInfo& device_info = rocprofiler_instance.GetDeviceInfo(gpu_id); @@ -1046,6 +1063,7 @@ void HSASupport_Singleton::HSAInitialize(HsaApiTable* table) { agent, (hsa_agent_info_t)(HSA_AMD_AGENT_INFO_NEAREST_CPU), &nearCpuAgent); rocprofiler::HSAAgentInfo agent_info(agent, device_type); agent_info.SetNearCpuAgent(nearCpuAgent); + agent_info.SetDeviceInfo(device_info); Packet::InitializeGPUPool(agent, &agent_info); hsasupport_singleton.SetHSAAgentInfo(agent, agent_info); @@ -1121,5 +1139,23 @@ void HSASupport_Singleton::HSAFinalize() { FinitKsymbols(); } +#define CHECK_HSA_STATUS(msg, status) \ +if ((status) != HSA_STATUS_SUCCESS && (status) != HSA_STATUS_INFO_BREAK) { \ + try { \ + const char* emsg = nullptr; \ + hsa_status_string(status, &emsg); \ + if (!emsg) emsg = ""; \ + std::cerr << msg << std::endl; \ + std::cerr << emsg << std::endl; \ + } catch (std::exception & e) { \ + } \ + abort(); \ +} + +void HSASupport_Singleton::CreateSignal(uint32_t attribute, hsa_signal_t* signal) { + auto status = GetAmdExtTable().hsa_amd_signal_create_fn(1, 0, nullptr, attribute, signal); + CHECK_HSA_STATUS("Error: hsa_amd_signal_create failed", status); +} + } // namespace rocprofiler diff --git a/src/core/hsa/hsa_support.h b/src/core/hsa/hsa_support.h index 219e074939..4cf2302a8c 100644 --- a/src/core/hsa/hsa_support.h +++ b/src/core/hsa/hsa_support.h @@ -169,6 +169,7 @@ class HSASupport_Singleton { void HSAFinalize(); void InitKsymbols(); void FinitKsymbols(); + void CreateSignal(uint32_t attribute, hsa_signal_t* signal); HSASupport_Singleton(const HSASupport_Singleton&) = delete; HSASupport_Singleton& operator=(const HSASupport_Singleton&) = delete; }; diff --git a/src/core/hsa/packets/packets_generator.cpp b/src/core/hsa/packets/packets_generator.cpp index 8571215392..c724eb3b79 100644 --- a/src/core/hsa/packets/packets_generator.cpp +++ b/src/core/hsa/packets/packets_generator.cpp @@ -18,6 +18,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +//#define HSA_ATT_MARKER_ENABLE + #include "packets_generator.h" #include "src/api/rocprofiler_singleton.h" @@ -235,14 +237,12 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, rocprofiler::Metric* metric; if (std::find(counter_names.begin(), counter_names.end(), result->name) != counter_names.end()) { - // std::cout << "Counter from Result List: " << result->name << std::endl; counters_taken.insert(result->name); metric = const_cast(metricsDict[gpu_agent.handle]->Get(result->name)); if (metric == nullptr) std::cout << result->name << " not found in metricsDict\n"; context->metrics_list.push_back(metric); } else { metrics_counters_taken.insert(result->name); - // std::cout << "Counter Added: " << result->name << std::endl; } } @@ -283,7 +283,6 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, } } if (flag) { - // std::cout << "Counter from Result Map: " << metric_name << std::endl; counters_taken.insert(metric_name); rocprofiler::Metric* metric = const_cast(metricsDict[gpu_agent.handle]->Get(metric_name)); @@ -586,10 +585,12 @@ hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets( #pragma GCC diagnostic push #pragma GCC diagnostic ignored "-Wconversion-null" // Preparing the profile structure to get the packets + auto* params = new hsa_ven_amd_aqlprofile_parameter_t[att_params.size()+1]; + memcpy(params, att_params.data(), att_params.size()*sizeof(params[0])); hsa_ven_amd_aqlprofile_profile_t* profile = new hsa_ven_amd_aqlprofile_profile_t{gpu_agent, HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_TRACE, nullptr, 0, - &att_params[0], (uint32_t)att_params.size(), + params, (uint32_t)att_params.size(), NULL, NULL}; #pragma GCC diagnostic pop @@ -614,4 +615,108 @@ hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets( return profile; } + +// Generate ATT tracer marker packets. Also generate and return +// the descriptor object which has the PM4 buffer for inserting data +hsa_ven_amd_aqlprofile_descriptor_t +GenerateATTMarkerPackets(hsa_agent_t gpu_agent, packet_t& marker_packet, uint32_t data) +{ +#ifdef HSA_ATT_MARKER_ENABLE + // Preparing the profile structure to get the packets + auto pool = rocprofiler::HSASupport_Singleton::GetInstance() + .GetHSAAgentInfo(gpu_agent.handle) + .cpu_pool_; + + hsa_ven_amd_aqlprofile_descriptor_t desc{AllocateSysMemory(gpu_agent, 1024, &pool), 1024}; + hsa_ven_amd_aqlprofile_profile_t profile{ + gpu_agent, HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_TRACE, + nullptr, 0, + nullptr, 0, + {}, desc + }; + + hsa_ven_amd_aqlprofile_att_marker_channel_t channel = HSA_VEN_AMD_AQLPROFILE_ATT_CHANNEL_2; + hsa_status_t status = hsa_ven_amd_aqlprofile_att_marker(&profile, &marker_packet, data, channel); + if (status != HSA_STATUS_SUCCESS) + { + rocprofiler::warning("Could not create ATT Marker Packets."); + desc.size = 0; + desc.ptr = nullptr; + } + + return desc; +#else + return {nullptr,0}; +#endif +} + +void AddVendorSpecificPacket(const packet_t* packet, + std::vector* transformed_packets, + const hsa_signal_t& packet_completion_signal) { + transformed_packets->emplace_back(*packet).completion_signal = packet_completion_signal; +} + +/* + Function name: CreateBarrierPacket. + Argument : The list of transformed packets to add the + barrier packet to. Pointer to the completion signal + and the input signal of the barrier packet to be created. + Description : This packet creates the barrier packet with the given + completion signal and dependency signal. It then adds to + the transformed packets list. +*/ +void CreateBarrierPacket(std::vector* transformed_packets, + const hsa_signal_t* packet_dependency_signal, + const hsa_signal_t* packet_completion_signal + ) { + hsa_barrier_and_packet_t barrier{0}; + barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE | + (1 << HSA_PACKET_HEADER_BARRIER) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); + if (packet_completion_signal != nullptr) barrier.completion_signal = *packet_completion_signal; + if (packet_dependency_signal != nullptr) barrier.dep_signal[0] = *packet_dependency_signal; + void* barrier_ptr = &barrier; + transformed_packets->emplace_back(*reinterpret_cast(barrier_ptr)); +} + +template constexpr Integral bit_mask(int first, int last) { + assert(last >= first && "Error: hsa_support::bit_mask -> invalid argument"); + size_t num_bits = last - first + 1; + return ((num_bits >= sizeof(Integral) * 8) ? ~Integral{0} + /* num_bits exceed the size of Integral */ + : ((Integral{1} << num_bits) - 1)) + << first; +} + +template constexpr Integral bit_extract(Integral x, int first, int last) { + return (x >> first) & bit_mask(0, last - first); +} + +bool IsDispatchPacket(const hsa_barrier_and_packet_t& packet) +{ + return bit_extract( + packet.header, + HSA_PACKET_HEADER_TYPE, + HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1 + ) == HSA_PACKET_TYPE_KERNEL_DISPATCH; +} + +// Returns a list of pointers to dispatch packets. +std::vector ExtractDispatchPackets( + const void* packets, + int pkt_count +) { + std::vector ret; + for (int i = 0; i < pkt_count; ++i) { + auto& original_packet = static_cast(packets)[i]; + + // Skip packets other than kernel dispatch packets. + if (!IsDispatchPacket(original_packet)) continue; + + ret.push_back(static_cast(packets)+i); + } + return ret; +} + } // namespace Packet diff --git a/src/core/hsa/packets/packets_generator.h b/src/core/hsa/packets/packets_generator.h index 866fef3544..205e1c525a 100644 --- a/src/core/hsa/packets/packets_generator.h +++ b/src/core/hsa/packets/packets_generator.h @@ -60,6 +60,8 @@ hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets( std::vector& att_params, packet_t* start_packet, packet_t* stop_packet, size_t att_buffer_size); +hsa_ven_amd_aqlprofile_descriptor_t +GenerateATTMarkerPackets(hsa_agent_t gpu_agent, packet_t& marker_packet, uint32_t data); uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool); @@ -74,6 +76,21 @@ typedef struct { att_memory_pools_t* GetAttMemPools(hsa_agent_t gpu_agent); +void AddVendorSpecificPacket(const packet_t* packet, + std::vector* transformed_packets, + const hsa_signal_t& packet_completion_signal); + +void CreateBarrierPacket(std::vector* transformed_packets, + const hsa_signal_t* packet_dependency_signal, + const hsa_signal_t* packet_completion_signal); + +bool IsDispatchPacket(const hsa_barrier_and_packet_t& packet); + +// Returns a list of pointers to dispatch packets. +std::vector ExtractDispatchPackets( + const void* packets, + int pkt_count +); } // namespace Packet #endif // SRC_CORE_HSA_PACKETS_PACKETS_GENERATOR_H_ diff --git a/src/core/hsa/queues/queue.cpp b/src/core/hsa/queues/queue.cpp index 7897d959e1..b02e238f4f 100644 --- a/src/core/hsa/queues/queue.cpp +++ b/src/core/hsa/queues/queue.cpp @@ -29,6 +29,7 @@ #include #include #include +#include #include "rocprofiler.h" #include "src/api/rocprofiler_singleton.h" @@ -55,8 +56,6 @@ #define __NR_gettid 186 -#define DEFAULT_ATT_BUFFER_SIZE 0x40000000 - std::mutex sessions_pending_signal_lock; namespace rocprofiler { @@ -274,17 +273,6 @@ hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type, return status; } -hsa_status_t attTraceDataCallback(hsa_ven_amd_aqlprofile_info_type_t info_type, - hsa_ven_amd_aqlprofile_info_data_t* info_data, void* data) { - hsa_status_t status = HSA_STATUS_SUCCESS; - att_trace_callback_data_t* passed_data = reinterpret_cast(data); - passed_data->push_back(*info_data); - // TODO: clear output buffers after copying - // either copy here or in ::AddAttRecord - - return status; -} - void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t* pending) { record->counters_count = rocprofiler_record_counters_instances_count_t{pending->context->metrics_list.size()}; @@ -323,52 +311,6 @@ void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_sign } } -void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent, - att_pending_signal_t& pending) { - HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance(); - HSAAgentInfo agent_info = hsasupport_singleton.GetHSAAgentInfo(gpu_agent.handle); - att_trace_callback_data_t data; - hsa_status_t status = - hsa_ven_amd_aqlprofile_iterate_data(pending.profile, attTraceDataCallback, &data); - - if ((status & HSA_STATUS_ERROR_OUT_OF_RESOURCES) == HSA_STATUS_ERROR_OUT_OF_RESOURCES) - rocprofiler::warning("Warning: ATT buffer full!\n"); - if ((status & HSA_STATUS_ERROR_EXCEPTION) == HSA_STATUS_ERROR_EXCEPTION) - rocprofiler::warning("Warning: ATT received a UTC memory error!\n"); - if (status == HSA_STATUS_ERROR) fatal("Thread Trace Error!"); - - // Allocate memory for shader_engine_data - record->shader_engine_data = static_cast( - calloc(data.size(), sizeof(rocprofiler_record_se_att_data_t))); - - att_trace_callback_data_t::iterator trace_data_it; - - uint32_t se_index = 0; - // iterate over the trace data collected from each shader engine - for (trace_data_it = data.begin(); trace_data_it != data.end(); trace_data_it++) { - const void* data_ptr = trace_data_it->trace_data.ptr; - const uint32_t data_size = trace_data_it->trace_data.size; - - void* buffer = NULL; - if (data_size != 0) { - // Allocate buffer on CPU to copy out trace data - buffer = Packet::AllocateSysMemory(gpu_agent, data_size, &agent_info.cpu_pool_); - if (buffer == NULL) fatal("Trace data buffer allocation failed"); - - auto status = - hsasupport_singleton.GetCoreApiTable().hsa_memory_copy_fn(buffer, data_ptr, data_size); - if (status != HSA_STATUS_SUCCESS) fatal("Trace data memcopy to host failed"); - - record->shader_engine_data[se_index].buffer_ptr = buffer; - record->shader_engine_data[se_index].buffer_size = data_size; - ++se_index; - - // TODO: clear output buffers after copying - } - } - record->shader_engine_data_count = data.size(); -} - /* Function name: enable_dispatch Argument : pointer to the the Queue class object @@ -574,104 +516,6 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) { return false; } -bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) { - auto queue_info_session = static_cast(data); - rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton = - rocprofiler::ROCProfiler_Singleton::GetInstance(); - rocprofiler::HSASupport_Singleton& hsasupport_singleton = - rocprofiler::HSASupport_Singleton::GetInstance(); - if (!queue_info_session || !rocprofiler_singleton.GetSession(queue_info_session->session_id) || - !rocprofiler_singleton.GetSession(queue_info_session->session_id)->GetAttTracer()) - return true; - rocprofiler::Session* session = rocprofiler_singleton.GetSession(queue_info_session->session_id); - rocprofiler::att::AttTracer* att_tracer = session->GetAttTracer(); - if (!session->GetAttTracer()) return true; - - std::vector& pending_signals = - const_cast&>( - att_tracer->GetPendingSignals(queue_info_session->writer_id)); - - if (!pending_signals.empty()) { - for (auto it = pending_signals.begin(); it != pending_signals.end(); - it = pending_signals.erase(it)) { - auto& pending = *it; - std::lock_guard lock(session->GetSessionLock()); - if (hsasupport_singleton.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}; - record.kernel_properties = pending.kernel_properties; - record.thread_id = rocprofiler_thread_id_t{pending.thread_id}; - record.queue_idx = rocprofiler_queue_index_t{pending.queue_index}; - record.queue_id = rocprofiler_queue_id_t{queue_info_session->queue_id}; - record.writer_id = queue_info_session->writer_id; - - if (/*pending.counters_count > 0 && */ pending.profile) { - AddAttRecord(&record, queue_info_session->agent, pending); - } - // July/01/2023 -> Changed this to queue_info_session->writer_id - // so we can correlate to dispatches. kernel_id already has the descriptor. - record.header = {ROCPROFILER_ATT_TRACER_RECORD, - rocprofiler_record_id_t{pending.kernel_descriptor}}; - - record.intercept_list = codeobj_record::get_capture(record.header.id); - std::atomic_thread_fence(std::memory_order_release); - - if (pending.session_id.handle == 0) { - pending.session_id = rocprofiler_singleton.GetCurrentSessionId(); - } - if (session->FindBuffer(pending.buffer_id)) { - Memory::GenericBuffer* buffer = session->GetBuffer(pending.buffer_id); - buffer->AddRecord(record); - buffer->Flush(); - } - codeobj_record::free_capture(record.header.id); - - hsa_status_t status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn( - (pending.profile->output_buffer.ptr)); - CHECK_HSA_STATUS("Error: Couldn't free output buffer memory", status); - status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn( - (pending.profile->command_buffer.ptr)); - CHECK_HSA_STATUS("Error: Couldn't free command buffer memory", status); - delete pending.profile; - } - } - delete queue_info_session; - - std::atomic_thread_fence(std::memory_order_seq_cst); - return false; -} - -/* - Function name: CreateBarrierPacket. - Argument : The list of transformed packets to add the - barrier packet to. Pointer to the completion signal - and the input signal of the barrier packet to be created. - Description : This packet creates the barrier packet with the given - completion signal and dependency signal. It then adds to - the transformed packets list. -*/ -void CreateBarrierPacket(std::vector* transformed_packets, - const hsa_signal_t* packet_dependency_signal, - const hsa_signal_t* packet_completion_signal - ) { - hsa_barrier_and_packet_t barrier{0}; - barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE | - (1 << HSA_PACKET_HEADER_BARRIER) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); - if (packet_completion_signal != nullptr) barrier.completion_signal = *packet_completion_signal; - if (packet_dependency_signal != nullptr) barrier.dep_signal[0] = *packet_dependency_signal; - void* barrier_ptr = &barrier; - transformed_packets->emplace_back(*reinterpret_cast(barrier_ptr)); -} - -void AddVendorSpecificPacket(const Packet::packet_t* packet, - std::vector* transformed_packets, - const hsa_signal_t& packet_completion_signal) { - transformed_packets->emplace_back(*packet).completion_signal = packet_completion_signal; -} void SignalAsyncHandler(const hsa_signal_t& signal, void* data) { hsa_status_t status = @@ -680,32 +524,8 @@ void SignalAsyncHandler(const hsa_signal_t& signal, void* data) { CHECK_HSA_STATUS("Error: hsa_amd_signal_async_handler failed", status); } -void signalAsyncHandlerATT(const hsa_signal_t& signal, void* data) { - hsa_status_t status = - HSASupport_Singleton::GetInstance().GetAmdExtTable().hsa_amd_signal_async_handler_fn( - signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalHandlerATT, data); - CHECK_HSA_STATUS("Error: hsa_amd_signal_async_handler for ATT failed", status); -} - void CreateSignal(uint32_t attribute, hsa_signal_t* signal) { - hsa_status_t status = - HSASupport_Singleton::GetInstance().GetAmdExtTable().hsa_amd_signal_create_fn( - 1, 0, nullptr, attribute, signal); - CHECK_HSA_STATUS("Error: hsa_amd_signal_create failed", status); -} - -template constexpr Integral bit_mask(int first, int last) { - assert(last >= first && "Error: hsa_support::bit_mask -> invalid argument"); - size_t num_bits = last - first + 1; - return ((num_bits >= sizeof(Integral) * 8) ? ~Integral{0} - /* num_bits exceed the size of Integral */ - : ((Integral{1} << num_bits) - 1)) - << first; -} - -/* Extract bits [last:first] from t. */ -template constexpr Integral bit_extract(Integral x, int first, int last) { - return (x >> first) & bit_mask(0, last - first); + HSASupport_Singleton::GetInstance().CreateSignal(attribute, signal); } rocprofiler_session_id_t session_id = rocprofiler_session_id_t{0}; @@ -720,11 +540,8 @@ bool is_counter_collection_mode = false; bool is_timestamp_collection_mode = false; bool is_att_collection_mode = false; bool is_pc_sampling_collection_mode = false; -std::vector att_parameters_data; uint32_t replay_mode_count = 0; -std::vector kernel_profile_names; -std::vector kernel_profile_dispatch_ids; -std::vector att_counters_names; + rocprofiler::Session* session = nullptr; @@ -737,162 +554,52 @@ void CheckNeededProfileConfigs() { rocprofiler::ROCProfiler_Singleton::GetInstance(); internal_session_id = rocprofiler_singleton.GetCurrentSessionId(); + if (session_id.handle > 0 && internal_session_id.handle == session_id.handle) return; + if (internal_session_id.handle == 0) return; + session_id = internal_session_id; - if (session_id.handle == 0 || internal_session_id.handle != session_id.handle) { - session_id = internal_session_id; - // Getting Counters count from the Session - if (session_id.handle > 0) { - session = rocprofiler_singleton.GetSession(session_id); - if (session && session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION)) { - rocprofiler_filter_id_t filter_id = - session->GetFilterIdWithKind(ROCPROFILER_COUNTERS_COLLECTION); - rocprofiler::Filter* filter = session->GetFilter(filter_id); - session_data = filter->GetCounterData(); - is_counter_collection_mode = true; - session_data_count = session_data.size(); - buffer_id = filter->GetBufferId(); - } else if (session && - session->FindFilterWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION)) { - is_timestamp_collection_mode = true; - rocprofiler_filter_id_t filter_id = - session->GetFilterIdWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION); - rocprofiler::Filter* filter = session->GetFilter(filter_id); - buffer_id = filter->GetBufferId(); - } else if (session && session->FindFilterWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) { - rocprofiler_filter_id_t filter_id = - session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION); - rocprofiler::Filter* filter = session->GetFilter(filter_id); - att_parameters_data = filter->GetAttParametersData(); - is_att_collection_mode = true; - buffer_id = - session->GetFilter(session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) - ->GetBufferId(); + // Getting Counters count from the Session + session = rocprofiler_singleton.GetSession(session_id); - att_counters_names = filter->GetCounterData(); - kernel_profile_names = std::get>( - filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES)); - kernel_profile_dispatch_ids = - std::get>(filter->GetProperty(ROCPROFILER_FILTER_DISPATCH_IDS)); - } else if (session && session->FindFilterWithKind(ROCPROFILER_PC_SAMPLING_COLLECTION)) { - is_pc_sampling_collection_mode = true; - } - } + if (session && session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION)) { + rocprofiler_filter_id_t filter_id = + session->GetFilterIdWithKind(ROCPROFILER_COUNTERS_COLLECTION); + rocprofiler::Filter* filter = session->GetFilter(filter_id); + session_data = filter->GetCounterData(); + is_counter_collection_mode = true; + session_data_count = session_data.size(); + buffer_id = filter->GetBufferId(); + } else if (session && + session->FindFilterWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION)) { + is_timestamp_collection_mode = true; + rocprofiler_filter_id_t filter_id = + session->GetFilterIdWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION); + rocprofiler::Filter* filter = session->GetFilter(filter_id); + buffer_id = filter->GetBufferId(); + } else if (session && session->FindFilterWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) { + rocprofiler_filter_id_t filter_id = + session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION); + rocprofiler::Filter* filter = session->GetFilter(filter_id); + + auto* att_tracer = session->GetAttTracer(); + att_tracer->SetParameters(filter->GetAttParametersData()); + is_att_collection_mode = true; + buffer_id = session->GetFilter(session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) + ->GetBufferId(); + att_tracer->SetCountersNames(filter->GetCounterData()); + att_tracer->SetKernelsNames(std::get>( + filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES) + )); + att_tracer->SetDispatchIds(std::get>>( + filter->GetProperty(ROCPROFILER_FILTER_DISPATCH_IDS) + )); + } else if (session && session->FindFilterWithKind(ROCPROFILER_PC_SAMPLING_COLLECTION)) { + is_pc_sampling_collection_mode = true; } } -static int KernelInterceptCount = 0; std::atomic WRITER_ID{0}; -std::pair, bool> GetAllowedProfilesList(const void* packets, int pkt_count) { - std::vector can_profile_packet; - bool b_can_profile_anypacket = false; - can_profile_packet.reserve(pkt_count); - rocprofiler::HSASupport_Singleton& hsasupport_singleton = - rocprofiler::HSASupport_Singleton::GetInstance(); - std::lock_guard lock(hsasupport_singleton.ksymbol_map_lock); - assert(hsasupport_singleton.ksymbols); - - uint32_t current_writer_id = WRITER_ID.load(std::memory_order_relaxed); - - for (int i = 0; i < pkt_count; ++i) { - auto& original_packet = static_cast(packets)[i]; - bool b_profile_this_object = false; - - // Skip packets other than kernel dispatch packets. - if (bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE, - HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) == - HSA_PACKET_TYPE_KERNEL_DISPATCH) { - auto& kdispatch = static_cast(packets)[i]; - - // If Dispatch IDs specified, profile based on dispatch ID - for (auto id : kernel_profile_dispatch_ids) b_profile_this_object |= id == current_writer_id; - try { - // Can throw - const std::string& kernel_name = hsasupport_singleton.ksymbols->at(kdispatch.kernel_object); - - // If no filters specified, auto profile this kernel - if (kernel_profile_names.size() == 0 && kernel_profile_dispatch_ids.size() == 0 && - kernel_name.find("__amd_rocclr_") == std::string::npos) - b_profile_this_object = true; - - // Try to match the mangled kernel name with given matches in input.txt - // We want to initiate att profiling if a match exists - for (const std::string& kernel_matches : kernel_profile_names) - if (kernel_name.find(kernel_matches) != std::string::npos) b_profile_this_object = true; - } catch (...) { - rocprofiler::warning("Warning: Unknown name for object %lu\n", kdispatch.kernel_object); - } - current_writer_id += 1; - } - b_can_profile_anypacket |= b_profile_this_object; - can_profile_packet.push_back(b_profile_this_object); - } - // If we're going to skip all packets, need to update writer ID - if (!b_can_profile_anypacket) WRITER_ID.store(current_writer_id, std::memory_order_release); - return {can_profile_packet, b_can_profile_anypacket}; -} - -std::pair ProcessATTParams( - Packet::packet_t& start_packet, Packet::packet_t& stop_packet, Queue& queue_info, - rocprofiler::HSAAgentInfo& agentInfo) { - std::vector att_params; - int num_att_counters = 0; - uint32_t att_buffer_size = DEFAULT_ATT_BUFFER_SIZE; - rocprofiler_codeobj_capture_mode_t capture_mode = ROCPROFILER_CAPTURE_SYMBOLS_ONLY; - - for (rocprofiler_att_parameter_t& param : att_parameters_data) { - switch (param.parameter_name) { - case ROCPROFILER_ATT_PERFCOUNTER_NAME: - break; - case ROCPROFILER_ATT_CAPTURE_MODE: - capture_mode = static_cast(param.value); - break; - case ROCPROFILER_ATT_BUFFER_SIZE: - att_buffer_size = - std::max(96l << 10l, std::min(int64_t(param.value) << 20l, (1l << 32l) - (3l << 20))); - break; // Clip to [96KB, 4GB) - case ROCPROFILER_ATT_PERFCOUNTER: - num_att_counters += 1; - break; - default: - att_params.push_back( - {static_cast(int(param.parameter_name)), - param.value}); - } - } - - if (att_counters_names.size() > 0) { - MetricsDict* metrics_dict_ = MetricsDict::Create(&agentInfo); - - for (const std::string& counter_name : att_counters_names) { - const Metric* metric = metrics_dict_->Get(counter_name); - const BaseMetric* base = dynamic_cast(metric); - if (!base) rocprofiler::fatal("Invalid base metric value: %s\n", counter_name.c_str()); - - std::vector counters; - base->GetCounters(counters); - hsa_ven_amd_aqlprofile_event_t event = counters[0]->event; - if (event.block_name != HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ) - rocprofiler::fatal("Only events from the SQ block can be selected for ATT.\n"); - - att_params.push_back( - {static_cast(int(ROCPROFILER_ATT_PERFCOUNTER)), - event.counter_id | (event.counter_id ? (0xF << 24) : 0)}); - num_att_counters += 1; - } - - hsa_ven_amd_aqlprofile_parameter_t zero_perf = { - static_cast(int(ROCPROFILER_ATT_PERFCOUNTER)), 0}; - - // Fill other perfcounters with 0's - for (; num_att_counters < 16; num_att_counters++) att_params.push_back(zero_perf); - } - // Get the PM4 Packets using packets_generator - return {Packet::GenerateATTPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(), att_params, - &start_packet, &stop_packet, att_buffer_size), - capture_mode}; -} - /** * @brief This function is a queue write interceptor. It intercepts the * packet write function. Creates an instance of packet class with the raw @@ -902,23 +609,19 @@ std::pair */ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index, void* data, hsa_amd_queue_intercept_packet_writer writer) { - static const char* env_MAX_ATT_PROFILES = getenv("ROCPROFILER_MAX_ATT_PROFILES"); - static int MAX_ATT_PROFILES = env_MAX_ATT_PROFILES ? atoi(env_MAX_ATT_PROFILES) : 1; - const Packet::packet_t* packets_arr = reinterpret_cast(packets); std::vector transformed_packets; CheckNeededProfileConfigs(); rocprofiler_session_id_t session_id_snapshot = session_id; + auto& queue_info = *reinterpret_cast(data); + std::lock_guard lk(queue_info.qw_mutex); + if (session_id_snapshot.handle > 0 && pkt_count > 0 && (is_counter_collection_mode || is_timestamp_collection_mode || is_pc_sampling_collection_mode) && session) { - // Getting Queue Data and Information - auto& queue_info = *reinterpret_cast(data); - std::lock_guard lk(queue_info.qw_mutex); - // hsa_ven_amd_aqlprofile_profile_t* profile; std::vector> @@ -929,10 +632,7 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u // +Skip kernel dispatch IDs not wanted // Skip packets other than kernel dispatch packets. - if (session_id_snapshot.handle == 0 || - bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE, - HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) != - HSA_PACKET_TYPE_KERNEL_DISPATCH) { + if (session_id_snapshot.handle == 0 || !Packet::IsDispatchPacket(original_packet)) { transformed_packets.emplace_back(packets_arr[i]); continue; } @@ -959,13 +659,13 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u Creates a barrier packet with its completion signal as the queue's ready signal. */ - CreateBarrierPacket(&transformed_packets, nullptr, &ready_signal); + Packet::CreateBarrierPacket(&transformed_packets, nullptr, &ready_signal); /* Creates a barrier packet with queue's blocksignal as its input and completion signal.This will ensure it is no longer 0 so a later barrier packet waiting on it to be 0 will be blocked */ - CreateBarrierPacket(&transformed_packets, &block_signal, &block_signal); + Packet::CreateBarrierPacket(&transformed_packets, &block_signal, &block_signal); } @@ -978,10 +678,13 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u dummy_signal.handle = 0; profile.first->start_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - AddVendorSpecificPacket(profile.first->start_packet, &transformed_packets, dummy_signal); + Packet::AddVendorSpecificPacket(profile.first->start_packet, &transformed_packets, dummy_signal); - CreateBarrierPacket(&transformed_packets, &profile.first->start_packet->completion_signal, - nullptr); + Packet::CreateBarrierPacket( + &transformed_packets, + &profile.first->start_packet->completion_signal, + nullptr + ); } auto& packet = transformed_packets.emplace_back(packets_arr[i]); @@ -1043,18 +746,18 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u hsa_signal_t dummy_signal{}; profile.first->stop_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - AddVendorSpecificPacket(profile.first->stop_packet, &transformed_packets, dummy_signal); + Packet::AddVendorSpecificPacket(profile.first->stop_packet, &transformed_packets, dummy_signal); profile.first->read_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - AddVendorSpecificPacket(profile.first->read_packet, &transformed_packets, interrupt_signal); + Packet::AddVendorSpecificPacket(profile.first->read_packet, &transformed_packets, interrupt_signal); // Added Interrupt Signal with barrier and provided handler for it - CreateBarrierPacket( &transformed_packets, &interrupt_signal, nullptr); + Packet::CreateBarrierPacket( &transformed_packets, &interrupt_signal, nullptr); } else - CreateBarrierPacket( &transformed_packets, nullptr, &interrupt_signal); + Packet::CreateBarrierPacket( &transformed_packets, nullptr, &interrupt_signal); rocprofiler::HSAAgentInfo& agentInfo = rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo( queue_info.GetGPUAgent().handle); @@ -1073,142 +776,15 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u } /* Write the transformed packets to the hardware queue. */ writer(&transformed_packets[0], transformed_packets.size()); - } else if (session_id_snapshot.handle > 0 && pkt_count > 0 && is_att_collection_mode && session && - KernelInterceptCount < MAX_ATT_PROFILES) { - // att start - // Getting Queue Data and Information - auto& queue_info = *static_cast(data); - std::lock_guard lk(queue_info.qw_mutex); - rocprofiler::HSAAgentInfo& agentInfo = - rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo( - queue_info.GetGPUAgent().handle); - - bool can_profile_anypacket = false; - std::vector can_profile_packet; - std::tie(can_profile_packet, can_profile_anypacket) = - GetAllowedProfilesList(packets, pkt_count); - - if (!can_profile_anypacket) { - /* Write the original packets to the hardware if no patch will be profiled */ - writer(packets, pkt_count); - return; - } - - // Preparing att Packets - Packet::packet_t start_packet{}; - Packet::packet_t stop_packet{}; - hsa_ven_amd_aqlprofile_profile_t* profile = nullptr; - rocprofiler_codeobj_capture_mode_t capture_mode = ROCPROFILER_CAPTURE_SYMBOLS_ONLY; - - if (att_parameters_data.size() > 0) { - std::tie(profile, capture_mode) = - ProcessATTParams(start_packet, stop_packet, queue_info, agentInfo); - } - - // Searching across all the packets given during this write - for (size_t i = 0; i < pkt_count; ++i) { - auto& original_packet = static_cast(packets)[i]; - uint32_t writer_id = 0; - - // Skip all packets marked with !can_profile - if (i >= can_profile_packet.size() || can_profile_packet[i] == false) { - transformed_packets.emplace_back(packets_arr[i]); - - // increment writer ID for every packet - if (bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE, - HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) == - HSA_PACKET_TYPE_KERNEL_DISPATCH) - writer_id = WRITER_ID.fetch_add(1, std::memory_order_release); - - continue; - } - KernelInterceptCount += 1; - writer_id = WRITER_ID.fetch_add(1, std::memory_order_release); - - if (!att_parameters_data.empty() && profile) { - // Adding start packet and its barrier with a dummy signal - hsa_signal_t dummy_signal{}; - dummy_signal.handle = 0; - start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - AddVendorSpecificPacket(&start_packet, &transformed_packets, dummy_signal); - CreateBarrierPacket(&transformed_packets, &start_packet.completion_signal, nullptr) ; - } - - auto& packet = transformed_packets.emplace_back(packets_arr[i]); - auto& dispatch_packet = reinterpret_cast(packet); - - CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal); - - // Adding the dispatch packet newly created signal to the pending signals - // 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 = rocprofiler::ROCProfiler_Singleton::GetInstance().GetUniqueRecordId(); - AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object), - record_id); - - session->GetAttTracer()->AddPendingSignals( - writer_id, record_id, original_packet.completion_signal, packet.completion_signal, - session_id_snapshot, buffer_id, profile, kernel_properties, - (uint32_t)syscall(__NR_gettid), user_pkt_index); - - uint64_t off = dispatch_packet.kernel_object + - GetKernelCode(dispatch_packet.kernel_object)->kernel_code_entry_byte_offset; - codeobj_record::make_capture(rocprofiler_record_id_t{record_id}, capture_mode, off); - - uint64_t IsGFX9 = HSASupport_Singleton::GetInstance() - .GetHSAAgentInfo(queue_info.GetGPUAgent().handle) - .GetDeviceInfo() - .getName() - .find("gfx9") != std::string::npos; - codeobj_record::make_capture(rocprofiler_record_id_t{record_id}, capture_mode, IsGFX9 | (off<<1)); - codeobj_record::start_capture(rocprofiler_record_id_t{record_id}); - codeobj_record::stop_capture(rocprofiler_record_id_t{record_id}); - - // Make a copy of the original packet, adding its signal to a barrier packet - if (original_packet.completion_signal.handle != 0U) { - hsa_barrier_and_packet_t barrier{}; - barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; - barrier.dep_signal[0] = packet.completion_signal; - Packet::packet_t* __attribute__((__may_alias__)) pkt = - (reinterpret_cast(&barrier)); - transformed_packets.emplace_back(*pkt).completion_signal = - original_packet.completion_signal; - } - - // Adding a barrier packet with the original packet's completion signal. - hsa_signal_t interrupt_signal; - CreateSignal(0, &interrupt_signal); - - // Adding Stop PM4 Packets - if (att_parameters_data.size() > 0 && is_att_collection_mode && profile) { - stop_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - AddVendorSpecificPacket(&stop_packet, &transformed_packets, interrupt_signal); - - // Added Interrupt Signal with barrier and provided handler for it - CreateBarrierPacket(&transformed_packets, &interrupt_signal, nullptr); - } else { - hsa_barrier_and_packet_t barrier{}; - barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; - barrier.completion_signal = interrupt_signal; - Packet::packet_t* __attribute__((__may_alias__)) pkt = - (reinterpret_cast(&barrier)); - transformed_packets.emplace_back(*pkt); - } - - // Creating Async Handler to be called every time the interrupt signal is - // marked complete - signalAsyncHandlerATT( - interrupt_signal, - new queue_info_session_t{queue_info.GetGPUAgent(), session_id_snapshot, - queue_info.GetQueueID(), writer_id, interrupt_signal}); - } - /* Write the transformed packets to the hardware queue. */ - writer(&transformed_packets[0], transformed_packets.size()); - // ATT end - } else { - /* Write the original packets to the hardware queue if no profiling session - * is active */ + } else if (!is_att_collection_mode || !session->GetAttTracer()->ATTWriteInterceptor( + packets, + pkt_count, + user_pkt_index, + *static_cast(data), + writer, + buffer_id + )) { + /* Write the original packets to the hardware queue if no profiling session is active */ writer(packets, pkt_count); } } @@ -1227,7 +803,6 @@ Queue::~Queue() { { profiler_serializer_t& serializer = rocprofiler::ROCProfiler_Singleton::GetInstance().GetSerializer(); - // std::cout << GetROCMToolObj()-> std::lock_guard serializer_lock(serializer.serializer_mutex); for (auto it = serializer.dispatch_ready.begin(); it != serializer.dispatch_ready.end();) { if ((*it)->GetQueueID() == GetQueueID()) { diff --git a/src/core/hsa/queues/queue.h b/src/core/hsa/queues/queue.h index c837871ee7..d49fa6c80e 100644 --- a/src/core/hsa/queues/queue.h +++ b/src/core/hsa/queues/queue.h @@ -70,6 +70,12 @@ class Queue { static void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index, void* data, hsa_amd_queue_intercept_packet_writer writer); + static bool ATTWriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index, + void* data, hsa_amd_queue_intercept_packet_writer writer); + static bool ATTSingleWriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index, + void* data, hsa_amd_queue_intercept_packet_writer writer); + static bool ATTContiguousWriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index, + void* data, hsa_amd_queue_intercept_packet_writer writer); hsa_queue_t* GetCurrentInterceptQueue(); hsa_agent_t GetGPUAgent(); hsa_agent_t GetCPUAgent(); diff --git a/src/core/isa_capture/code_object_track.cpp b/src/core/isa_capture/code_object_track.cpp index bf5fba741e..1fa1e57b57 100644 --- a/src/core/isa_capture/code_object_track.cpp +++ b/src/core/isa_capture/code_object_track.cpp @@ -48,14 +48,11 @@ std::mutex codeobj_record::mutex; std::unordered_map codeobj_record::codeobjs{}; std::unordered_map codeobj_record::record_id_map{}; std::unordered_set codeobj_record::listeners; +std::atomic codeobj_capture_instance::loadcount{0}; // Codeobj Record codeobj_record::codeobj_record(rocprofiler_codeobj_capture_mode_t mode) : capture_mode(mode){}; -static uint64_t timestamp() { - return rocprofiler::ROCProfiler_Singleton::GetInstance().timestamp_ns().value; -} - void codeobj_record::start_capture() { listeners.insert(this); for (auto& [addr, capture] : codeobjs) this->addcapture(capture); @@ -72,10 +69,18 @@ void codeobj_record::stop_capture() { } // Codeobj Capture -void codeobj_capture_instance::Load(uint64_t addr, const std::string& URI, uint64_t mem_addr, - uint64_t size) { +void codeobj_capture_instance::Load( + uint64_t addr, + const std::string& URI, + uint64_t mem_addr, + uint64_t size +) { + uint32_t id = loadcount.fetch_add(1, std::memory_order_relaxed)+1; + auto time = rocprofiler::ROCProfiler_Singleton::GetInstance().timestamp_ns().value; + std::lock_guard lock(codeobj_record::mutex); - auto inst = std::make_shared(addr, URI, mem_addr, size, timestamp()); + + auto inst = std::make_shared(addr, URI, mem_addr, size, time, id); codeobj_record::codeobjs[addr] = inst; for (auto* listen : codeobj_record::listeners) listen->addcapture(inst); } @@ -85,7 +90,8 @@ void codeobj_capture_instance::Unload(uint64_t addr) { if (codeobj_record::codeobjs.find(addr) == codeobj_record::codeobjs.end()) return; - codeobj_record::codeobjs.at(addr)->end_time = timestamp(); + auto time = rocprofiler::ROCProfiler_Singleton::GetInstance().timestamp_ns().value; + codeobj_record::codeobjs.at(addr)->end_time = time; codeobj_record::codeobjs.erase(addr); } @@ -179,11 +185,6 @@ std::pair codeobj_capture_instance::parse_uri() { return {offset, size}; } -codeobj_capture_instance::codeobj_capture_instance(uint64_t _addr, const std::string& _uri, - uint64_t mem_addr, uint64_t mem_size, - uint64_t start_time) - : addr(_addr), start_time(start_time), URI(_uri), mem_addr(mem_addr), mem_size(mem_size) {}; - void codeobj_capture_instance::setmode(rocprofiler_codeobj_capture_mode_t mode) { // Only reset when needed & check if codeobj was not unloaded if (static_cast(mode) > capture_mode) reset(mode); diff --git a/src/core/isa_capture/code_object_track.hpp b/src/core/isa_capture/code_object_track.hpp index 96add163b8..6e5de1ed3c 100644 --- a/src/core/isa_capture/code_object_track.hpp +++ b/src/core/isa_capture/code_object_track.hpp @@ -35,22 +35,29 @@ class codeobj_capture_instance { public: codeobj_capture_instance(uint64_t _addr, const std::string& _uri, uint64_t mem_addr, - uint64_t mem_size, uint64_t start_time); + uint64_t mem_size, uint64_t start_time, uint32_t id) + : addr(_addr), start_time(start_time), URI(_uri), + mem_addr(mem_addr), mem_size(mem_size), load_id(id) {}; void setmode(rocprofiler_codeobj_capture_mode_t mode); rocprofiler_intercepted_codeobj_t get() const { const char* buf_ptr = buffer.size() ? buffer.data() : nullptr; - return {URI.c_str(), addr, mem_size, buf_ptr, buffer.size(), start_time, end_time}; + return {URI.c_str(), addr, mem_size, buf_ptr, buffer.size(), start_time, end_time, load_id}; }; const uint64_t addr; const uint64_t start_time; + const uint32_t load_id; static void Load(uint64_t addr, const std::string& URI, uint64_t mem_addr, uint64_t mem_size); static void Unload(uint64_t addr); + static uint32_t GetLoadCount() { return loadcount.load(std::memory_order_relaxed); } private: + //! 32 bits ID because this is the natural channel width for ATT Markers. + //! There is no world in which 4 billions markers can be sent anyway. + static std::atomic loadcount; void reset(rocprofiler_codeobj_capture_mode_t mode); std::pair parse_uri(); diff --git a/src/core/proxy_queue.h b/src/core/proxy_queue.h index b0a1ed6653..3374b81674 100644 --- a/src/core/proxy_queue.h +++ b/src/core/proxy_queue.h @@ -29,8 +29,8 @@ THE SOFTWARE. #include #include -#include "core/queue.h" -#include "core/types.h" +#include "src/core/queue.h" +#include "src/core/types.h" struct HsaApiTable; diff --git a/src/core/session/att/att.cpp b/src/core/session/att/att.cpp index 6b4aeafe2d..2c63c6c299 100644 --- a/src/core/session/att/att.cpp +++ b/src/core/session/att/att.cpp @@ -22,6 +22,10 @@ #include #include +#include "src/core/hsa/packets/packets_generator.h" +#include "src/api/rocprofiler_singleton.h" +#include "src/core/isa_capture/code_object_track.hpp" + namespace rocprofiler { namespace att { @@ -31,16 +35,33 @@ AttTracer::AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t : buffer_id_(buffer_id), filter_id_(filter_id), session_id_(session_id) {} void AttTracer::AddPendingSignals( - 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, hsa_ven_amd_aqlprofile_profile_t* profile, - rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id, uint64_t queue_index) { + 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, + hsa_ven_amd_aqlprofile_profile_t* profile, + rocprofiler_kernel_properties_t kernel_properties, + uint32_t thread_id, uint64_t queue_index +) { std::lock_guard lock(sessions_pending_signals_lock_); - if (sessions_pending_signals_.find(writer_id) == sessions_pending_signals_.end()) - sessions_pending_signals_.emplace(writer_id, std::vector()); - sessions_pending_signals_.at(writer_id).emplace_back(att_pending_signal_t{ - kernel_object, original_completion_signal, new_completion_signal, session_id_, buffer_id, - profile, kernel_properties, thread_id, queue_index}); + + auto pending = sessions_pending_signals_.find(writer_id); + if (pending == sessions_pending_signals_.end()) + pending = sessions_pending_signals_.emplace(writer_id, std::vector()).first; + + pending->second.emplace_back(att_pending_signal_t{ + kernel_object, + original_completion_signal, + new_completion_signal, + session_id_, + buffer_id, + profile, + kernel_properties, + thread_id, + queue_index + }); } const std::vector& AttTracer::GetPendingSignals(uint32_t writer_id) { @@ -50,6 +71,258 @@ const std::vector& AttTracer::GetPendingSignals(uint32_t w return sessions_pending_signals_.at(writer_id); } +#define DEFAULT_ATT_BUFFER_SIZE 0x40000000 + +std::pair +AttTracer::ProcessATTParams( + hsa_ext_amd_aql_pm4_packet_t& start_packet, + hsa_ext_amd_aql_pm4_packet_t& stop_packet, + queue::Queue& queue_info, + rocprofiler::HSAAgentInfo& agentInfo +) { + std::vector att_params; + int num_att_counters = 0; + uint32_t att_buffer_size = DEFAULT_ATT_BUFFER_SIZE; + rocprofiler_codeobj_capture_mode_t capture_mode = ROCPROFILER_CAPTURE_SYMBOLS_ONLY; + + for (rocprofiler_att_parameter_t& param : att_parameters_data) { + switch (param.parameter_name) { + case ROCPROFILER_ATT_PERFCOUNTER_NAME: + break; + case ROCPROFILER_ATT_CAPTURE_MODE: + capture_mode = static_cast(param.value); + break; + case ROCPROFILER_ATT_BUFFER_SIZE: + att_buffer_size = + std::max(96l << 10l, std::min(int64_t(param.value) << 20l, (1l << 32l) - (3l << 20))); + break; // Clip to [96KB, 4GB) + case ROCPROFILER_ATT_PERFCOUNTER: + num_att_counters += 1; + break; + default: + att_params.push_back( + {static_cast(int(param.parameter_name)), + param.value}); + } + } + + if (att_counters_names.size() > 0) { + MetricsDict* metrics_dict_ = MetricsDict::Create(&agentInfo); + + for (const std::string& counter_name : att_counters_names) { + const Metric* metric = metrics_dict_->Get(counter_name); + const BaseMetric* base = dynamic_cast(metric); + if (!base) rocprofiler::fatal("Invalid base metric value: %s\n", counter_name.c_str()); + + std::vector counters; + base->GetCounters(counters); + hsa_ven_amd_aqlprofile_event_t event = counters[0]->event; + if (event.block_name != HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ) + rocprofiler::fatal("Only events from the SQ block can be selected for ATT.\n"); + + att_params.push_back( + {static_cast(int(ROCPROFILER_ATT_PERFCOUNTER)), + event.counter_id | (event.counter_id ? (0xF << 24) : 0)}); + num_att_counters += 1; + } + + hsa_ven_amd_aqlprofile_parameter_t zero_perf = { + static_cast(int(ROCPROFILER_ATT_PERFCOUNTER)), 0}; + + // Fill other perfcounters with 0's + for (; num_att_counters < 16; num_att_counters++) att_params.push_back(zero_perf); + } + // Get the PM4 Packets using packets_generator + return {Packet::GenerateATTPackets( + queue_info.GetCPUAgent(), + queue_info.GetGPUAgent(), + att_params, + &start_packet, + &stop_packet, + att_buffer_size + ), + capture_mode}; + +} + +bool AttTracer::ATTWriteInterceptor( + const void* packets, + uint64_t pkt_count, + uint64_t user_pkt_index, + queue::Queue& queue, + hsa_amd_queue_intercept_packet_writer writer, + rocprofiler_buffer_id_t buffer_id +) { + bool IsSingleDispatchMode = kernel_profile_dispatch_ids.size() == 0; + + if (session_id_.handle == 0 || + pkt_count == 0 || + att_parameters_data.size() == 0 + ) return false; + + if (IsSingleDispatchMode) + return ATTSingleWriteInterceptor(packets, pkt_count, user_pkt_index, queue, writer, buffer_id); + else + return ATTContiguousWriteInterceptor(packets, pkt_count, queue, writer, buffer_id); +} + +void AttTracer::signalAsyncHandlerATT(const hsa_signal_t& signal, void* data) { + hsa_status_t status = + HSASupport_Singleton::GetInstance().GetAmdExtTable().hsa_amd_signal_async_handler_fn( + signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalHandlerATT, data); + if (status != HSA_STATUS_SUCCESS) + rocprofiler::fatal("Error: hsa_amd_signal_async_handler for ATT failed"); +} + +bool AttTracer::AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) { + auto queue_info_session = static_cast(data); + rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton = + rocprofiler::ROCProfiler_Singleton::GetInstance(); + + rocprofiler::HSASupport_Singleton& hsasupport_singleton = + rocprofiler::HSASupport_Singleton::GetInstance(); + + if (!queue_info_session || !rocprofiler_singleton.GetSession(queue_info_session->session_id) || + !rocprofiler_singleton.GetSession(queue_info_session->session_id)->GetAttTracer()) + return true; + + rocprofiler::Session* session = rocprofiler_singleton.GetSession(queue_info_session->session_id); + std::lock_guard lock(session->GetSessionLock()); + rocprofiler::att::AttTracer* att_tracer = session->GetAttTracer(); + + if (!session->GetAttTracer()) return true; + + std::vector& pending_signals = + const_cast&>( + att_tracer->GetPendingSignals(queue_info_session->writer_id)); + + if (!pending_signals.empty()) { + for (auto it = pending_signals.begin(); it != pending_signals.end(); + it = pending_signals.erase(it)) { + + auto& pending = *it; + //if (hsasupport_singleton.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}; + record.kernel_properties = pending.kernel_properties; + record.thread_id = rocprofiler_thread_id_t{pending.thread_id}; + record.queue_idx = rocprofiler_queue_index_t{pending.queue_index}; + record.queue_id = rocprofiler_queue_id_t{queue_info_session->queue_id}; + record.writer_id = queue_info_session->writer_id; + + if (/*pending.counters_count > 0 && */ pending.profile) { + AddAttRecord(&record, queue_info_session->agent, pending); + } + + // July/01/2023 -> Changed this to queue_info_session->writer_id + // so we can correlate to dispatches. kernel_id already has the descriptor. + record.header = {ROCPROFILER_ATT_TRACER_RECORD, + rocprofiler_record_id_t{pending.kernel_descriptor}}; + + record.intercept_list = codeobj_record::get_capture(record.header.id); + std::atomic_thread_fence(std::memory_order_release); + + if (pending.session_id.handle == 0) { + pending.session_id = rocprofiler_singleton.GetCurrentSessionId(); + } + + if (session->FindBuffer(pending.buffer_id)) { + Memory::GenericBuffer* buffer = session->GetBuffer(pending.buffer_id); + buffer->AddRecord(record); + buffer->Flush(); + } + codeobj_record::free_capture(record.header.id); + + hsa_status_t status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn( + (pending.profile->output_buffer.ptr)); + if (status != HSA_STATUS_SUCCESS) + rocprofiler::warning("Error: Couldn't free output buffer memory"); + + status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn( + (pending.profile->command_buffer.ptr)); + if (status != HSA_STATUS_SUCCESS) + rocprofiler::warning("Error: Couldn't free command buffer memory"); + + if (pending.profile->parameters) + delete[] pending.profile->parameters; + delete pending.profile; + } + } + delete queue_info_session; + + std::atomic_thread_fence(std::memory_order_seq_cst); + return false; +} + +void AttTracer::AddAttRecord( + rocprofiler_record_att_tracer_t* record, + hsa_agent_t gpu_agent, + att_pending_signal_t& pending +) { + HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance(); + HSAAgentInfo agent_info = hsasupport_singleton.GetHSAAgentInfo(gpu_agent.handle); + std::vector data; + hsa_status_t status = + hsa_ven_amd_aqlprofile_iterate_data(pending.profile, attTraceDataCallback, &data); + + if ((status & HSA_STATUS_ERROR_OUT_OF_RESOURCES) == HSA_STATUS_ERROR_OUT_OF_RESOURCES) + rocprofiler::warning("Warning: ATT buffer full!\n"); + if ((status & HSA_STATUS_ERROR_EXCEPTION) == HSA_STATUS_ERROR_EXCEPTION) + rocprofiler::warning("Warning: ATT received a UTC memory error!\n"); + if (status == HSA_STATUS_ERROR) fatal("Thread Trace Error!"); + + // Allocate memory for shader_engine_data + record->shader_engine_data = static_cast( + calloc(data.size(), sizeof(rocprofiler_record_se_att_data_t))); + + std::vector::iterator trace_data_it; + + uint32_t se_index = 0; + // iterate over the trace data collected from each shader engine + for (trace_data_it = data.begin(); trace_data_it != data.end(); trace_data_it++) { + const void* data_ptr = trace_data_it->trace_data.ptr; + const uint32_t data_size = trace_data_it->trace_data.size; + + void* buffer = NULL; + if (data_size != 0) { + // Allocate buffer on CPU to copy out trace data + buffer = Packet::AllocateSysMemory(gpu_agent, data_size, &agent_info.cpu_pool_); + if (buffer == NULL) fatal("Trace data buffer allocation failed"); + + auto status = + hsasupport_singleton.GetCoreApiTable().hsa_memory_copy_fn(buffer, data_ptr, data_size); + if (status != HSA_STATUS_SUCCESS) fatal("Trace data memcopy to host failed"); + + record->shader_engine_data[se_index].buffer_ptr = buffer; + record->shader_engine_data[se_index].buffer_size = data_size; + ++se_index; + + // TODO: clear output buffers after copying + } + } + record->shader_engine_data_count = data.size(); +} + +hsa_status_t AttTracer::attTraceDataCallback( + hsa_ven_amd_aqlprofile_info_type_t info_type, + hsa_ven_amd_aqlprofile_info_data_t* info_data, + void* data +) { + hsa_status_t status = HSA_STATUS_SUCCESS; + auto* passed_data = reinterpret_cast*>(data); + passed_data->push_back(*info_data); + // TODO: clear output buffers after copying + // either copy here or in ::AddAttRecord + + return status; +} + +std::unordered_map AttTracer::pending_stop_packets; +std::mutex AttTracer::att_enable_disable_mutex; + + } // namespace att } // namespace rocprofiler diff --git a/src/core/session/att/att.h b/src/core/session/att/att.h index 4da130dd92..0a2397beba 100644 --- a/src/core/session/att/att.h +++ b/src/core/session/att/att.h @@ -27,8 +27,15 @@ #include #include #include +#include +#include #include "rocprofiler.h" +#include "src/utils/helper.h" +#include "src/core/proxy_queue.h" +#include "src/core/hsa/hsa_support.h" +#include "src/core/hsa/queues/queue.h" +#include "att_header.h" namespace rocprofiler { @@ -44,33 +51,25 @@ typedef struct { uint64_t queue_index; } att_pending_signal_t; -enum rocprofiler_att_isa_dump_mode { - ISA_MODE_DUMP_ALL=0, - ISA_MODE_DUMP_OBJ, - ISA_MODE_DUMP_KERNEL, - ISA_MODE_DUMP_NONE -}; - -union att_header_packet_t { - struct { - uint64_t reserved : 14; - uint64_t navi : 1; - uint64_t enable : 1; - uint64_t DSIMDM : 4; - uint64_t DCU : 5; - uint64_t DSA : 1; - uint64_t SEID : 6; - uint64_t isadumpmode : 3; - }; - uint64_t raw; -}; namespace att { +struct ATTRecordSignal +{ + size_t record_id; + size_t writer_id; + size_t last_kernel_exec; + rocprofiler_session_id_t session_id_snapshot; + hsa_ext_amd_aql_pm4_packet_t stop_packet; +}; + class AttTracer { - public: - AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t filter_id, - rocprofiler_session_id_t session_id); +public: + AttTracer( + rocprofiler_buffer_id_t buffer_id, + rocprofiler_filter_id_t filter_id, + rocprofiler_session_id_t session_id + ); void AddPendingSignals(uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& original_completion_signal, @@ -82,10 +81,98 @@ class AttTracer { const std::vector& GetPendingSignals(uint32_t writer_id); - private: + bool ATTWriteInterceptor( + const void* packets, + uint64_t pkt_count, + uint64_t user_pkt_index, + queue::Queue& queue_info, + hsa_amd_queue_intercept_packet_writer writer, + rocprofiler_buffer_id_t buffer_id + ); + + void InsertMarker( + std::vector& transformed_packets, + hsa_agent_t agent, + uint32_t data + ); + + void SetParameters(const std::vector& params) { + att_parameters_data = params; + } + void SetDispatchIds(const std::vector>& ids) { + kernel_profile_dispatch_ids = ids; + } + void SetCountersNames(const std::vector& names) { + att_counters_names = names; + } + void SetKernelsNames(const std::vector& names) { + kernel_profile_names = names; + } + std::optional> RequiresStartPacket(size_t rstart, size_t size); + + static void signalAsyncHandlerATT(const hsa_signal_t& signal, void* data); + static bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data); + + static hsa_status_t attTraceDataCallback( + hsa_ven_amd_aqlprofile_info_type_t info_type, + hsa_ven_amd_aqlprofile_info_data_t* info_data, + void* data + ); + +protected: + using packet_t = hsa_ext_amd_aql_pm4_packet_t; + static std::unordered_map pending_stop_packets; + static std::mutex att_enable_disable_mutex; + +private: + uint32_t codeobj_load_cnt = 0; + + static void AddAttRecord( + rocprofiler_record_att_tracer_t* record, + hsa_agent_t gpu_agent, + att_pending_signal_t& pending + ); + + std::pair + ProcessATTParams( + hsa_ext_amd_aql_pm4_packet_t& start_packet, + hsa_ext_amd_aql_pm4_packet_t& stop_packet, + queue::Queue& queue_info, + rocprofiler::HSAAgentInfo& agentInfo + ); + + bool ATTSingleWriteInterceptor( + const void* packets, + uint64_t pkt_count, + uint64_t user_pkt_index, + queue::Queue& queue_info, + hsa_amd_queue_intercept_packet_writer writer, + rocprofiler_buffer_id_t buffer_id + ); + + bool ATTContiguousWriteInterceptor( + const void* packets, + uint64_t pkt_count, + queue::Queue& queue_info, + hsa_amd_queue_intercept_packet_writer writer, + rocprofiler_buffer_id_t buffer_id + ); + + static void CreateSignal(uint32_t attribute, hsa_signal_t* signal) { + HSASupport_Singleton::GetInstance().CreateSignal(attribute, signal); + } + + std::pair, bool> GetAllowedProfilesList(const void* packets, int pkt_count); + rocprofiler_buffer_id_t buffer_id_; rocprofiler_filter_id_t filter_id_; rocprofiler_session_id_t session_id_; + std::atomic WRITER_ID{1}; + + std::vector kernel_profile_names; + std::vector> kernel_profile_dispatch_ids; + std::vector att_counters_names; + std::vector att_parameters_data; std::mutex sessions_pending_signals_lock_; std::map> sessions_pending_signals_; diff --git a/src/core/session/att/att_header.h b/src/core/session/att/att_header.h new file mode 100644 index 0000000000..6d0fb73275 --- /dev/null +++ b/src/core/session/att/att_header.h @@ -0,0 +1,38 @@ +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#pragma once + +namespace rocprofiler { + +union att_header_packet_t { + struct { + uint64_t reserved : 14; + uint64_t navi : 1; + uint64_t enable : 1; + uint64_t DSIMDM : 4; + uint64_t DCU : 5; + uint64_t DSA : 1; + uint64_t SEID : 6; + }; + uint64_t raw; +}; + +} \ No newline at end of file diff --git a/src/core/session/att/continuous.cpp b/src/core/session/att/continuous.cpp new file mode 100644 index 0000000000..8a4c3dac29 --- /dev/null +++ b/src/core/session/att/continuous.cpp @@ -0,0 +1,195 @@ +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#include "att.h" +#include +#include +#include "src/core/hsa/packets/packets_generator.h" +#include "src/api/rocprofiler_singleton.h" +#include "src/core/isa_capture/code_object_track.hpp" + +#define __NR_gettid 186 + +namespace rocprofiler { + +namespace att { + +void AttTracer::InsertMarker( + std::vector& transformed_packets, + hsa_agent_t agent, + uint32_t data +) { + packet_t marker_packet{}; + auto desc = Packet::GenerateATTMarkerPackets(agent, marker_packet, data); + if (desc.ptr && desc.size) + Packet::AddVendorSpecificPacket(&marker_packet, &transformed_packets, hsa_signal_t{.handle = 0}); +} + + +std::optional> AttTracer::RequiresStartPacket(size_t rstart, size_t size) +{ + for (auto& r : kernel_profile_dispatch_ids) + if (rstart <= r.first && rstart+size > r.first) + return r; + return {}; +} + +bool AttTracer::ATTContiguousWriteInterceptor( + const void* packets, + uint64_t pkt_count, + queue::Queue& queue_info, + hsa_amd_queue_intercept_packet_writer writer, + rocprofiler_buffer_id_t buffer_id +) { + const packet_t* packets_arr = reinterpret_cast(packets); + std::vector transformed_packets; + + // att start + // Getting Queue Data and Information + rocprofiler::HSAAgentInfo& agentInfo = rocprofiler::HSASupport_Singleton::GetInstance() + .GetHSAAgentInfo(queue_info.GetGPUAgent().handle); + + auto dispatchPackets = Packet::ExtractDispatchPackets(packets, pkt_count); + if (dispatchPackets.size() == 0) return false; + + size_t writer_id = WRITER_ID.fetch_add(dispatchPackets.size(), std::memory_order_relaxed); + uint32_t new_load_cnt = codeobj_capture_instance::GetLoadCount(); + + auto bInsertStart = RequiresStartPacket(writer_id, dispatchPackets.size()); + { + std::lock_guard lk(att_enable_disable_mutex); + // If att_start already exists, don't start again + auto agent_pending_packets = pending_stop_packets.find(queue_info.GetGPUAgent().handle); + if (agent_pending_packets != pending_stop_packets.end()) + bInsertStart = {}; + + // If nothing will be added or removed, return + if (!bInsertStart && codeobj_load_cnt == new_load_cnt) + { + if ( + agent_pending_packets == pending_stop_packets.end() || + agent_pending_packets->second.last_kernel_exec > writer_id + dispatchPackets.size() + ) + return false; + } + } + + if (bInsertStart) + { + // Preparing att Packets + packet_t start_packet{}; + packet_t stop_packet{}; + hsa_ven_amd_aqlprofile_profile_t* profile = nullptr; + rocprofiler_codeobj_capture_mode_t capturem = ROCPROFILER_CAPTURE_SYMBOLS_ONLY; + std::tie(profile, capturem) = ProcessATTParams(start_packet, stop_packet, queue_info, agentInfo); + + if (!profile) + { + rocprofiler::warning("Failed to create profile from queue!"); + return false; + } + + uint64_t IsGFX9 = HSASupport_Singleton::GetInstance() + .GetHSAAgentInfo(queue_info.GetGPUAgent().handle) + .GetDeviceInfo() + .getName() + .find("gfx9") != std::string::npos; + + hsa_signal_t dummy_signal{}; + dummy_signal.handle = 0; + start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + Packet::AddVendorSpecificPacket(&start_packet, &transformed_packets, dummy_signal); + Packet::CreateBarrierPacket(&transformed_packets, &start_packet.completion_signal, nullptr); + + uint64_t record_id = rocprofiler::ROCProfiler_Singleton::GetInstance().GetUniqueRecordId(); + AddKernelNameWithDispatchID("ATT_Contiguous", record_id); + + this->AddPendingSignals( + writer_id, + record_id, + start_packet.completion_signal, + start_packet.completion_signal, + session_id_, + buffer_id, + profile, + {0}, + (uint32_t)syscall(__NR_gettid), + 0 + ); + + codeobj_record::make_capture(rocprofiler_record_id_t{record_id}, capturem, IsGFX9); + codeobj_record::start_capture(rocprofiler_record_id_t{record_id}); + + stop_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + std::lock_guard lk(att_enable_disable_mutex); + pending_stop_packets[queue_info.GetGPUAgent().handle] + = {record_id, writer_id, bInsertStart->second, session_id_, stop_packet}; + } + + if (codeobj_load_cnt != new_load_cnt) + { + codeobj_load_cnt = new_load_cnt; + InsertMarker(transformed_packets, queue_info.GetGPUAgent(), new_load_cnt); + } + + // Searching across all the packets given during this write + for (size_t i = 0; i < pkt_count; ++i) + transformed_packets.emplace_back(packets_arr[i]); + + { + std::lock_guard lk(att_enable_disable_mutex); + auto agent_pending_packets = pending_stop_packets.find(queue_info.GetGPUAgent().handle); + + if (agent_pending_packets != pending_stop_packets.end() && + agent_pending_packets->second.last_kernel_exec <= writer_id + dispatchPackets.size() + ) { + const ATTRecordSignal& rsignal = agent_pending_packets->second; + // Adding a barrier packet with the original packet's completion signal. + hsa_signal_t interrupt_signal; + CreateSignal(0, &interrupt_signal); + + // Adding Stop PM4 Packets + Packet::AddVendorSpecificPacket(&rsignal.stop_packet, &transformed_packets, interrupt_signal); + + // Added Interrupt Signal with barrier and provided handler for it + Packet::CreateBarrierPacket(&transformed_packets, &interrupt_signal, nullptr); + + // Creating Async Handler to be called every time the interrupt signal is marked complete + signalAsyncHandlerATT(interrupt_signal, new queue::queue_info_session_t{ + queue_info.GetGPUAgent(), + rsignal.session_id_snapshot, + queue_info.GetQueueID(), + rsignal.writer_id, + interrupt_signal + }); + + codeobj_record::stop_capture(rocprofiler_record_id_t{rsignal.record_id}); + pending_stop_packets.erase(queue_info.GetGPUAgent().handle); + } + } + + /* Write the transformed packets to the hardware queue. */ + writer(&transformed_packets[0], transformed_packets.size()); + return true; +} + +} // namespace att + +} // namespace rocprofiler diff --git a/src/core/session/att/singlekernel.cpp b/src/core/session/att/singlekernel.cpp new file mode 100644 index 0000000000..ff7603e3a5 --- /dev/null +++ b/src/core/session/att/singlekernel.cpp @@ -0,0 +1,233 @@ +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#include "att.h" +#include +#include +#include "src/core/hsa/packets/packets_generator.h" +#include "src/api/rocprofiler_singleton.h" +#include "src/core/isa_capture/code_object_track.hpp" + +#define __NR_gettid 186 + +namespace rocprofiler { + +namespace att { + +std::pair, bool> +AttTracer::GetAllowedProfilesList(const void* packets, int pkt_count) { + + std::vector can_profile_packet; + bool b_can_profile_anypacket = false; + can_profile_packet.reserve(pkt_count); + + rocprofiler::HSASupport_Singleton& hsasupport_singleton = + rocprofiler::HSASupport_Singleton::GetInstance(); + + std::lock_guard lock(hsasupport_singleton.ksymbol_map_lock); + assert(hsasupport_singleton.ksymbols); + + uint32_t current_writer_id = WRITER_ID.load(std::memory_order_relaxed); + + for (int i = 0; i < pkt_count; ++i) { + auto& original_packet = static_cast(packets)[i]; + bool b_profile_this_object = false; + + // Skip packets other than kernel dispatch packets. + if (Packet::IsDispatchPacket(original_packet)) { + auto& kdispatch = static_cast(packets)[i]; + try { + // Can throw + const std::string& kernel_name = hsasupport_singleton.ksymbols->at(kdispatch.kernel_object); + + // If no filters specified, auto profile this kernel + if (kernel_profile_names.size() == 0 && + kernel_name.find("__amd_rocclr_") == std::string::npos + ) { + b_profile_this_object = true; + } else { + // Try to match the mangled kernel name with given matches in input.txt + // We want to initiate att profiling if a match exists + for (const std::string& kernel_matches : kernel_profile_names) + if (kernel_name.find(kernel_matches) != std::string::npos) b_profile_this_object = true; + } + } catch (...) { + rocprofiler::warning("Warning: Unknown name for object %lu\n", kdispatch.kernel_object); + } + current_writer_id += 1; + } + b_can_profile_anypacket |= b_profile_this_object; + can_profile_packet.push_back(b_profile_this_object); + } + // If we're going to skip all packets, need to update writer ID + if (!b_can_profile_anypacket) + WRITER_ID.store(current_writer_id, std::memory_order_release); + + return {can_profile_packet, b_can_profile_anypacket}; +} + +bool AttTracer::ATTSingleWriteInterceptor( + const void* packets, + uint64_t pkt_count, + uint64_t user_pkt_index, + queue::Queue& queue_info, + hsa_amd_queue_intercept_packet_writer writer, + rocprofiler_buffer_id_t buffer_id +) { + static int KernelInterceptCount = 0; + static const char* env_MAX_ATT_PROFILES = getenv("ROCPROFILER_MAX_ATT_PROFILES"); + static int MAX_ATT_PROFILES = env_MAX_ATT_PROFILES ? atoi(env_MAX_ATT_PROFILES) : 1; + + if (KernelInterceptCount >= MAX_ATT_PROFILES) return false; + + const packet_t* packets_arr = reinterpret_cast(packets); + std::vector transformed_packets; + + // att start + // Getting Queue Data and Information + rocprofiler::HSAAgentInfo& agentInfo = + rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo( + queue_info.GetGPUAgent().handle); + + bool can_profile_anypacket = false; + std::vector can_profile_packet; + std::tie(can_profile_packet, can_profile_anypacket) = GetAllowedProfilesList(packets, pkt_count); + + if (!can_profile_anypacket) return false; + + // Preparing att Packets + packet_t start_packet{}; + packet_t stop_packet{}; + hsa_ven_amd_aqlprofile_profile_t* profile = nullptr; + rocprofiler_codeobj_capture_mode_t capturem = ROCPROFILER_CAPTURE_SYMBOLS_ONLY; + + std::tie(profile, capturem) = ProcessATTParams(start_packet, stop_packet, queue_info, agentInfo); + + // Searching across all the packets given during this write + for (size_t i = 0; i < pkt_count; ++i) { + auto& original_packet = static_cast(packets)[i]; + uint32_t writer_id = 0; + + // Skip all packets marked with !can_profile + if (i >= can_profile_packet.size() || can_profile_packet[i] == false) { + transformed_packets.emplace_back(packets_arr[i]); + + // increment writer ID for every packet + if (Packet::IsDispatchPacket(original_packet)) + writer_id = WRITER_ID.fetch_add(1, std::memory_order_release); + + continue; + } + KernelInterceptCount += 1; + writer_id = WRITER_ID.fetch_add(1, std::memory_order_release); + + if (profile) { + // Adding start packet and its barrier with a dummy signal + hsa_signal_t dummy_signal{}; + dummy_signal.handle = 0; + start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + Packet::AddVendorSpecificPacket(&start_packet, &transformed_packets, dummy_signal); + Packet::CreateBarrierPacket(&transformed_packets, &start_packet.completion_signal, nullptr) ; + } + + auto& packet = transformed_packets.emplace_back(packets_arr[i]); + auto& dispatch_packet = reinterpret_cast(packet); + + CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal); + + // Adding the dispatch packet newly created signal to the pending signals + // list to be processed by the signal interrupt + uint64_t record_id = rocprofiler::ROCProfiler_Singleton::GetInstance().GetUniqueRecordId(); + AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object), record_id); + + this->AddPendingSignals( + writer_id, + record_id, + original_packet.completion_signal, + packet.completion_signal, + session_id_, + buffer_id, + profile, + {0}, + (uint32_t)syscall(__NR_gettid), + user_pkt_index + ); + + uint64_t IsGFX9 = HSASupport_Singleton::GetInstance() + .GetHSAAgentInfo(queue_info.GetGPUAgent().handle) + .GetDeviceInfo() + .getName() + .find("gfx9") != std::string::npos; + codeobj_record::make_capture(rocprofiler_record_id_t{record_id}, capturem, IsGFX9); + codeobj_record::start_capture(rocprofiler_record_id_t{record_id}); + codeobj_record::stop_capture(rocprofiler_record_id_t{record_id}); + + // Make a copy of the original packet, adding its signal to a barrier packet + if (original_packet.completion_signal.handle != 0U) { + hsa_barrier_and_packet_t barrier{}; + barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; + barrier.dep_signal[0] = packet.completion_signal; + packet_t* __attribute__((__may_alias__)) pkt = + (reinterpret_cast(&barrier)); + transformed_packets.emplace_back(*pkt).completion_signal = + original_packet.completion_signal; + } + + // Adding a barrier packet with the original packet's completion signal. + hsa_signal_t interrupt_signal; + CreateSignal(0, &interrupt_signal); + + // Adding Stop PM4 Packets + if (profile) { + stop_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + Packet::AddVendorSpecificPacket(&stop_packet, &transformed_packets, interrupt_signal); + + // Added Interrupt Signal with barrier and provided handler for it + Packet::CreateBarrierPacket(&transformed_packets, &interrupt_signal, nullptr); + } else { + hsa_barrier_and_packet_t barrier{}; + barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; + barrier.completion_signal = interrupt_signal; + packet_t* __attribute__((__may_alias__)) pkt = + (reinterpret_cast(&barrier)); + transformed_packets.emplace_back(*pkt); + } + + // Creating Async Handler to be called every time the interrupt signal is + // marked complete + signalAsyncHandlerATT(interrupt_signal, new queue::queue_info_session_t{ + queue_info.GetGPUAgent(), + session_id_, + queue_info.GetQueueID(), + writer_id, + interrupt_signal + }); + } + /* Write the transformed packets to the hardware queue. */ + writer(&transformed_packets[0], transformed_packets.size()); + return true; +} + + + + +} // namespace att + +} // namespace rocprofiler diff --git a/src/core/session/filter.cpp b/src/core/session/filter.cpp index 116cb07f05..7b94772e47 100644 --- a/src/core/session/filter.cpp +++ b/src/core/session/filter.cpp @@ -195,7 +195,7 @@ void Filter::SetProperty(rocprofiler_filter_property_t property) { case ROCPROFILER_FILTER_DISPATCH_IDS: dispatch_id_filter_.clear(); for (uint32_t j = 0; j < property.data_count; j++) - dispatch_id_filter_.emplace_back(property.dispatch_ids[j]); + dispatch_id_filter_.push_back({property.dispatch_ids[j].start, property.dispatch_ids[j].end}); break; default: break; @@ -205,9 +205,8 @@ void Filter::SetProperty(rocprofiler_filter_property_t property) { // "profiler mode!\n"); } } -std::variant, uint32_t*, std::vector> Filter::GetProperty( - rocprofiler_filter_property_kind_t kind) { - std::variant, uint32_t*, std::vector> property; +Filter::filter_property_variant_t Filter::GetProperty(rocprofiler_filter_property_kind_t kind) { + filter_property_variant_t property; switch (kind) { case ROCPROFILER_FILTER_GPU_NAME: { property = agent_names_; diff --git a/src/core/session/filter.h b/src/core/session/filter.h index a926bdba1f..4ffb0f0e8d 100644 --- a/src/core/session/filter.h +++ b/src/core/session/filter.h @@ -33,6 +33,11 @@ namespace rocprofiler { class Filter { public: + typedef std::variant< + std::vector, + uint32_t*, + std::vector> + > filter_property_variant_t; Filter(rocprofiler_filter_id_t id, rocprofiler_filter_kind_t filter_kind, rocprofiler_filter_data_t filter_data, uint64_t data_count); ~Filter(); @@ -53,8 +58,7 @@ class Filter { bool HasCallback(); void SetProperty(rocprofiler_filter_property_t property); - std::variant, uint32_t*, std::vector > GetProperty( - rocprofiler_filter_property_kind_t kind); + filter_property_variant_t GetProperty(rocprofiler_filter_property_kind_t kind); size_t GetPropertiesCount(rocprofiler_filter_property_kind_t kind); rocprofiler_spm_parameter_t* GetSpmParameterData(); @@ -77,7 +81,7 @@ class Filter { std::vector att_parameters_; // ATT Parameters rocprofiler_counters_sampler_parameters_t counters_sampler_parameters_; // sampled counters parameters - std::vector dispatch_id_filter_; + std::vector> dispatch_id_filter_; bool has_sync_callback_{false}; rocprofiler_sync_callback_t callback_; diff --git a/src/tools/tool.cpp b/src/tools/tool.cpp index f29e0f42c9..5bef396963 100644 --- a/src/tools/tool.cpp +++ b/src/tools/tool.cpp @@ -266,9 +266,14 @@ std::vector GetCounterNames() { return counters; } -typedef std::tuple>, - std::vector, std::vector, std::vector, uint64_t> - att_parsed_input_t; +struct att_parsed_input_t +{ + std::vector> params{}; + std::vector kernel_names{}; + std::vector counters_names{}; + std::vector> dispatch_ids{}; + rocprofiler::att_header_packet_t header{.raw = 0}; +}; static int GetMpRank() { std::vector mpivars = {"MPI_RANK", "OMPI_COMM_WORLD_RANK", "MV2_COMM_WORLD_RANK"}; @@ -278,12 +283,9 @@ static int GetMpRank() { } att_parsed_input_t GetATTParams() { - std::vector> parameters; - std::vector kernel_names; - std::vector counters_names; - std::vector dispatch_ids; + att_parsed_input_t ret{}; const char* path = getenv("COUNTERS_PATH"); - if (!path) return {{}, {}, {}, {}, 0}; + if (!path) return ret; // List of parameters the user can set. Maxvalue is unused. std::unordered_map ATT_PARAM_NAMES{}; @@ -317,12 +319,11 @@ att_parsed_input_t GetATTParams() { std::ifstream trace_file(path); if (!trace_file.is_open()) { std::cout << "Unable to open att trace file." << std::endl; - return {{}, {}, {}, {}, 0}; + return ret; } - rocprofiler::att_header_packet_t header{.raw = 0}; - header.enable = 1; - header.DSIMDM = default_params["SIMD_SELECT"]; + ret.header.enable = 1; + ret.header.DSIMDM = default_params["SIMD_SELECT"]; int MPI_RANK = GetMpRank(); bool started_att_counters = false; @@ -350,13 +351,13 @@ att_parsed_input_t GetATTParams() { if (!started_att_counters) continue; if (param_name == "REMOVE_HEADER_PACKET") { - header.enable = 0; + ret.header.enable = 0; continue; } else if (param_name == "KERNEL") { - kernel_names.push_back(line); + ret.kernel_names.push_back(line); continue; } else if (param_name == "PERFCOUNTER") { - counters_names.push_back(line); + ret.counters_names.push_back(line); continue; } else if (param_name == "DISPATCH") { size_t comma = line.find(','); @@ -364,7 +365,14 @@ att_parsed_input_t GetATTParams() { int rank = (comma < line.size() - 1) ? stoi(line.substr(comma + 1)) : 0; if (MPI_RANK < 0 || rank == MPI_RANK) // Only add ID if rank matches the one in input.txt - dispatch_ids.push_back(std::max(id - 1, 0)); // off by 1 in relation to kernel-trace + ret.dispatch_ids.push_back({id, id}); + continue; + } else if (param_name == "DISPATCH_RANGE") { + size_t comma = line.find(','); + int start = stoi(line.substr(0, comma)); + int end = (comma < line.size() - 1) ? stoi(line.substr(comma + 1)) : start; + + ret.dispatch_ids.push_back({start, end}); continue; } // param_value is a number @@ -380,13 +388,8 @@ att_parsed_input_t GetATTParams() { continue; } - if (param_name.find("ISA_DUMP_MODE") != std::string::npos) { - header.isadumpmode = param_value; - continue; - } - if (ATT_PARAM_NAMES.find(param_name) != ATT_PARAM_NAMES.end()) { - parameters.push_back(std::make_pair(ATT_PARAM_NAMES[param_name], param_value)); + ret.params.push_back(std::make_pair(ATT_PARAM_NAMES[param_name], param_value)); try { default_params.erase(param_name); } catch (...) { @@ -398,17 +401,17 @@ att_parsed_input_t GetATTParams() { } if (param_name.find("TARGET_CU") != std::string::npos) - header.DCU = param_value; + ret.header.DCU = param_value; else if (param_name.find("SIMD_SELECT") != std::string::npos) - header.DSIMDM = param_value; + ret.header.DSIMDM = param_value; } - if (!started_att_counters) return {{}, {}, {}, {}, 0}; + if (!started_att_counters) return {}; for (auto& param : default_params) - parameters.push_back(std::make_pair(ATT_PARAM_NAMES[param.first], param.second)); + ret.params.push_back(std::make_pair(ATT_PARAM_NAMES[param.first], param.second)); - return {parameters, kernel_names, counters_names, dispatch_ids, header.raw}; + return ret; } void finish() { @@ -733,20 +736,15 @@ ROCPROFILER_EXPORT bool OnLoad(void* table, uint64_t runtime_version, uint64_t f if (getenv("ROCPROFILER_ROCTX_TRACE")) apis_requested.emplace_back(ACTIVITY_DOMAIN_ROCTX); // ATT Parameters + att_parsed_input_t att_params = GetATTParams(); std::vector parameters; - std::vector> params; - std::vector kernel_names; - std::vector att_counters_names; - std::vector dispatch_ids; - uint64_t attheader; - std::tie(params, kernel_names, att_counters_names, dispatch_ids, attheader) = GetATTParams(); // load the plugins - plugins_load(params.size() ? (void*)attheader : (void*)&counter_names); + plugins_load(att_params.params.size() ? (void*)att_params.header.raw : (void*)&counter_names); - for (auto& kv_pair : params) + for (auto& kv_pair : att_params.params) parameters.emplace_back(rocprofiler_att_parameter_t{kv_pair.first, kv_pair.second}); - for (std::string& name : att_counters_names) { + for (std::string& name : att_params.counters_names) { rocprofiler_att_parameter_t param; param.parameter_name = ROCPROFILER_ATT_PERFCOUNTER_NAME; param.counter_name = name.c_str(); @@ -852,12 +850,12 @@ ROCPROFILER_EXPORT bool OnLoad(void* table, uint64_t runtime_version, uint64_t f rocprofiler_filter_property_t property = {}; std::vector kernel_names_c; - if (dispatch_ids.size()) { // Correlation ID filter + if (att_params.dispatch_ids.size()) { // Correlation ID filter property.kind = ROCPROFILER_FILTER_DISPATCH_IDS; - property.data_count = dispatch_ids.size(); - property.dispatch_ids = dispatch_ids.data(); + property.data_count = att_params.dispatch_ids.size(); + property.dispatch_ids = reinterpret_cast(att_params.dispatch_ids.data()); } else { // Kernel names filter - for (auto& name : kernel_names) kernel_names_c.push_back(name.data()); + for (auto& name : att_params.kernel_names) kernel_names_c.push_back(name.data()); property.kind = ROCPROFILER_FILTER_KERNEL_NAMES; property.data_count = kernel_names_c.size(); diff --git a/src/util/hsa_rsrc_factory.cpp b/src/util/hsa_rsrc_factory.cpp index 8689b92208..4fba4bb439 100644 --- a/src/util/hsa_rsrc_factory.cpp +++ b/src/util/hsa_rsrc_factory.cpp @@ -316,7 +316,10 @@ hsa_status_t HsaRsrcFactory::LoadAqlProfileLib(aqlprofile_pfn_t* api) { api->hsa_ven_amd_aqlprofile_iterate_data = (decltype(::hsa_ven_amd_aqlprofile_iterate_data)*)dlsym( handle, "hsa_ven_amd_aqlprofile_iterate_data"); - +#ifdef HSA_ATT_MARKER_ENABLE + api->hsa_ven_amd_aqlprofile_att_marker = (decltype(::hsa_ven_amd_aqlprofile_att_marker)*) + dlsym(handle, "hsa_ven_amd_aqlprofile_att_marker"); +#endif return HSA_STATUS_SUCCESS; }