SWDEV-432445: ATT continuous mode

Change-Id: I52732fc1dba41df3373ef6a19a428b00d729bf04


[ROCm/rocprofiler commit: aba6a1d986]

projects/rocprofiler/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

projects/rocprofiler/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<uint64_t>(data);
     header.reserved = 0x11;
-
-    isa_mode = static_cast<decltype(isa_mode)>(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<std::mutex> 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> " << 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<CodeObjectBinary> binary;
-      std::unique_ptr<code_object_decoder_t> decoder;
-
-      if (symbol.data && symbol.data_size) {
-        decoder = std::make_unique<code_object_decoder_t>(symbol.data, symbol.data_size);
-      } else if (std::string(symbol.filepath).find("file://") != std::string::npos) {
-        binary = std::make_unique<CodeObjectBinary>(symbol.filepath);
-        decoder =
-            std::make_unique<code_object_decoder_t>(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<std::mutex> 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<std::mutex> 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<std::mutex> 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<std::mutex> 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;

projects/rocprofiler/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,

projects/rocprofiler/plugin/att/code_printing.cpp

@@ -49,6 +49,26 @@
 #include <elfutils/libdw.h>
 #include <sys/mman.h>
 
+#include <atomic>
+
+#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<char>{};
   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<DisassemblyInstance>(*this);
-  m_symbol_map = disassembly->GetKernelMap();
+  try {
+    disassembly = std::make_unique<DisassemblyInstance>(*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<SymbolInfo> 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<instruction_instance_t, size_t>
+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<char> 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<code_object_decoder_t>(buffer.data(), buffer.size());
+  }
+  else
+  {
+    std::unique_ptr<CodeObjectBinary> binary = std::make_unique<CodeObjectBinary>(filepath);
+    decoder = std::make_unique<code_object_decoder_t>(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<instruction_instance_t, size_t>& CodeobjService::getDecoded(uint64_t addr)
+{
+  if (decoded_map.find(addr) != decoded_map.end())
+    return decoded_map[addr];
+
+  std::optional<uint64_t> 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<uint32_t*>(decoder->buffer.data()+0xb8);
+    uint64_t v_offset = *reinterpret_cast<uint32_t*>(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<uint64_t, std::unique_ptr<CodeobjService>> services{};
+std::atomic<uint64_t> 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<CodeobjService>(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)
+  }
+}

projects/rocprofiler/plugin/att/code_printing.hpp

@@ -27,28 +27,67 @@
 #include <string>
 #include <vector>
 #include <memory>
+#include <unordered_map>
 
 #include "disassembly.hpp"
 
 class code_object_decoder_t {
- public:
-  // void load_symbol_map();
+public:
   std::optional<SymbolInfo> find_symbol(uint64_t address);
 
   code_object_decoder_t(const char* codeobj_data, uint64_t codeobj_size);
   ~code_object_decoder_t();
 
+  std::pair<instruction_instance_t, size_t>
+  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<uint64_t, std::pair<std::string, size_t>> m_line_number_map;
-  std::map<uint64_t, SymbolInfo> m_symbol_map;
+  std::map<uint64_t, std::pair<std::string, size_t>> m_line_number_map{};
+  std::map<uint64_t, SymbolInfo> m_symbol_map{};
 
   std::string m_uri;
   std::vector<char> buffer;
-  std::vector<instruction_instance_t> instructions;
-  std::unique_ptr<DisassemblyInstance> disassembly;
+  std::vector<instruction_instance_t> instructions{};
+  std::unique_ptr<DisassemblyInstance> disassembly{};
+};
+
+class CodeobjService
+{
+public:
+  CodeobjService(const char* filepath, uint64_t load_base);
+  bool decode_single(uint64_t vaddr, uint64_t faddr);
+
+  std::pair<instruction_instance_t, size_t>& 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<uint64_t, std::pair<instruction_instance_t, size_t>> decoded_map;
+  std::unique_ptr<code_object_decoder_t> decoder{nullptr};
+
+  bool bNotElfFILE = false;
 };

projects/rocprofiler/plugin/att/disassembly.cpp

@@ -52,14 +52,23 @@
 #include "code_printing.hpp"
 #include <hsa/amd_hsa_elf.h>
 
-#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<std::string> 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<std::streamsize>::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<void*>(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<DisassemblyInstance*>(user_data);
   std::optional<uint64_t> faddr = va2fo(instance.buffer, vaddr);
@@ -210,23 +217,25 @@ amd_comgr_status_t DisassemblyInstance::symbol_callback(amd_comgr_symbol_t symbo
 
 std::map<uint64_t, SymbolInfo>& 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<uint64_t>(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<DisassemblyInstance*>(user_data);
-  instance.instructions.push_back({strdup(instruction), nullptr, 0});
+  instance.last_instruction.instruction = strdup(instruction);
 }
 
 #define CHECK_VA2FO(x, msg) if (!(x)) {                                \

projects/rocprofiler/plugin/att/disassembly.hpp

@@ -61,7 +61,7 @@ class DisassemblyInstance {
 
   void* buffer;
   int64_t size;
-  std::vector<instruction_instance_t>& instructions;
+  instruction_instance_t last_instruction;
   amd_comgr_disassembly_info_t info;
   amd_comgr_data_t data;
   std::map<uint64_t, SymbolInfo> symbol_map;

projects/rocprofiler/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)

projects/rocprofiler/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

projects/rocprofiler/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:  *;
 };

projects/rocprofiler/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)

projects/rocprofiler/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<bool*>(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 = "<Unknown HSA Error>";                                                   \
+    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

projects/rocprofiler/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;
 };

projects/rocprofiler/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<rocprofiler::Metric*>(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<rocprofiler::Metric*>(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<packet_t>* 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<packet_t>* 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<packet_t*>(barrier_ptr));
+}
+
+template <typename Integral = uint64_t> 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 <typename Integral> constexpr Integral bit_extract(Integral x, int first, int last) {
+  return (x >> first) & bit_mask<Integral>(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<const hsa_kernel_dispatch_packet_s*> ExtractDispatchPackets(
+  const void* packets,
+  int pkt_count
+) {
+  std::vector<const hsa_kernel_dispatch_packet_s*> ret;
+  for (int i = 0; i < pkt_count; ++i) {
+    auto& original_packet = static_cast<const hsa_barrier_and_packet_t*>(packets)[i];
+
+    // Skip packets other than kernel dispatch packets.
+    if (!IsDispatchPacket(original_packet)) continue;
+
+    ret.push_back(static_cast<const hsa_kernel_dispatch_packet_s*>(packets)+i);
+  }
+  return ret;
+}
+
 }  // namespace Packet

projects/rocprofiler/src/core/hsa/packets/packets_generator.h

@@ -60,6 +60,8 @@ hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets(
     std::vector<hsa_ven_amd_aqlprofile_parameter_t>& 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<packet_t>* transformed_packets,
+                             const hsa_signal_t& packet_completion_signal);
+
+void CreateBarrierPacket(std::vector<packet_t>* 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<const hsa_kernel_dispatch_packet_s*> ExtractDispatchPackets(
+  const void* packets,
+  int pkt_count
+);
 
 }  // namespace Packet
 #endif  // SRC_CORE_HSA_PACKETS_PACKETS_GENERATOR_H_

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

@@ -29,6 +29,7 @@
 #include <utility>
 #include <algorithm>
 #include <numa.h>
+#include <unordered_map>
 
 #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<att_trace_callback_data_t*>(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<rocprofiler_record_se_att_data_t*>(
-      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<queue_info_session_t*>(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<att_pending_signal_t>& pending_signals =
-      const_cast<std::vector<att_pending_signal_t>&>(
-          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<std::mutex> 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<Packet::packet_t>* 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<Packet::packet_t*>(barrier_ptr));
-}
-
-void AddVendorSpecificPacket(const Packet::packet_t* packet,
-                             std::vector<Packet::packet_t>* 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 <typename Integral = uint64_t> 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 <typename Integral> constexpr Integral bit_extract(Integral x, int first, int last) {
-  return (x >> first) & bit_mask<Integral>(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<rocprofiler_att_parameter_t> att_parameters_data;
 uint32_t replay_mode_count = 0;
-std::vector<std::string> kernel_profile_names;
-std::vector<uint64_t> kernel_profile_dispatch_ids;
-std::vector<std::string> 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<std::vector<std::string>>(
-            filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES));
-        kernel_profile_dispatch_ids =
-            std::get<std::vector<uint64_t>>(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<std::vector<std::string>>(
+      filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES)
+    ));
+    att_tracer->SetDispatchIds(std::get<std::vector<std::pair<uint64_t,uint64_t>>>(
+      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<uint32_t> WRITER_ID{0};
 
-std::pair<std::vector<bool>, bool> GetAllowedProfilesList(const void* packets, int pkt_count) {
-  std::vector<bool> 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<std::mutex> 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<const hsa_barrier_and_packet_t*>(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<const hsa_kernel_dispatch_packet_s*>(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<hsa_ven_amd_aqlprofile_profile_t*, rocprofiler_codeobj_capture_mode_t> ProcessATTParams(
-    Packet::packet_t& start_packet, Packet::packet_t& stop_packet, Queue& queue_info,
-    rocprofiler::HSAAgentInfo& agentInfo) {
-  std::vector<hsa_ven_amd_aqlprofile_parameter_t> 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<rocprofiler_codeobj_capture_mode_t>(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<hsa_ven_amd_aqlprofile_parameter_name_t>(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<const BaseMetric*>(metric);
-      if (!base) rocprofiler::fatal("Invalid base metric value: %s\n", counter_name.c_str());
-
-      std::vector<const counter_t*> 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<hsa_ven_amd_aqlprofile_parameter_name_t>(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<hsa_ven_amd_aqlprofile_parameter_name_t>(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<hsa_ven_amd_aqlprofile_profile_t*, rocprofiler_codeobj_capture_mode_t>
  */
 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<const Packet::packet_t*>(packets);
   std::vector<Packet::packet_t> transformed_packets;
 
   CheckNeededProfileConfigs();
   rocprofiler_session_id_t session_id_snapshot = session_id;
 
+  auto& queue_info = *reinterpret_cast<Queue*>(data);
+  std::lock_guard<std::mutex> 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<Queue*>(data);
-    std::lock_guard<std::mutex> lk(queue_info.qw_mutex);
-
 
     // hsa_ven_amd_aqlprofile_profile_t* profile;
     std::vector<std::pair<rocprofiler::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
@@ -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<Queue*>(data);
-    std::lock_guard<std::mutex> lk(queue_info.qw_mutex);
-    rocprofiler::HSAAgentInfo& agentInfo =
-        rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo(
-            queue_info.GetGPUAgent().handle);
-
-    bool can_profile_anypacket = false;
-    std::vector<bool> 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<const hsa_barrier_and_packet_t*>(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<hsa_kernel_dispatch_packet_t&>(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<Packet::packet_t*>(&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<Packet::packet_t*>(&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<Queue*>(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<std::mutex> serializer_lock(serializer.serializer_mutex);
     for (auto it = serializer.dispatch_ready.begin(); it != serializer.dispatch_ready.end();) {
       if ((*it)->GetQueueID() == GetQueueID()) {

projects/rocprofiler/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();

projects/rocprofiler/src/core/isa_capture/code_object_track.cpp

@@ -48,14 +48,11 @@ std::mutex codeobj_record::mutex;
 std::unordered_map<uint64_t, CodeobjPtr> codeobj_record::codeobjs{};
 std::unordered_map<uint64_t, codeobj_record::RecordInstance> codeobj_record::record_id_map{};
 std::unordered_set<codeobj_record*> codeobj_record::listeners;
+std::atomic<uint32_t> 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<std::mutex> lock(codeobj_record::mutex);
-  auto inst = std::make_shared<codeobj_capture_instance>(addr, URI, mem_addr, size, timestamp());
+
+  auto inst = std::make_shared<codeobj_capture_instance>(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<size_t, size_t> 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<int>(mode) > capture_mode) reset(mode);

projects/rocprofiler/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<uint32_t> loadcount;
   void reset(rocprofiler_codeobj_capture_mode_t mode);
 
   std::pair<size_t, size_t> parse_uri();

projects/rocprofiler/src/core/proxy_queue.h

@@ -29,8 +29,8 @@ THE SOFTWARE.
 #include <map>
 #include <mutex>
 
-#include "core/queue.h"
-#include "core/types.h"
+#include "src/core/queue.h"
+#include "src/core/types.h"
 
 struct HsaApiTable;
 

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

@@ -22,6 +22,10 @@
 #include <cassert>
 #include <atomic>
 
+#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<std::mutex> lock(sessions_pending_signals_lock_);
-  if (sessions_pending_signals_.find(writer_id) == sessions_pending_signals_.end())
-    sessions_pending_signals_.emplace(writer_id, std::vector<att_pending_signal_t>());
-  sessions_pending_signals_.at(writer_id).emplace_back(att_pending_signal_t{
-      kernel_object, 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<att_pending_signal_t>()).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<att_pending_signal_t>& AttTracer::GetPendingSignals(uint32_t writer_id) {
@@ -50,6 +71,258 @@ const std::vector<att_pending_signal_t>& AttTracer::GetPendingSignals(uint32_t w
   return sessions_pending_signals_.at(writer_id);
 }
 
+#define DEFAULT_ATT_BUFFER_SIZE 0x40000000
+
+std::pair<hsa_ven_amd_aqlprofile_profile_t*, rocprofiler_codeobj_capture_mode_t>
+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<hsa_ven_amd_aqlprofile_parameter_t> 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<rocprofiler_codeobj_capture_mode_t>(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<hsa_ven_amd_aqlprofile_parameter_name_t>(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<const BaseMetric*>(metric);
+      if (!base) rocprofiler::fatal("Invalid base metric value: %s\n", counter_name.c_str());
+
+      std::vector<const counter_t*> 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<hsa_ven_amd_aqlprofile_parameter_name_t>(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<hsa_ven_amd_aqlprofile_parameter_name_t>(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<queue::queue_info_session_t*>(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<std::mutex> lock(session->GetSessionLock());
+  rocprofiler::att::AttTracer* att_tracer = session->GetAttTracer();
+
+  if (!session->GetAttTracer()) return true;
+
+  std::vector<att_pending_signal_t>& pending_signals =
+      const_cast<std::vector<att_pending_signal_t>&>(
+          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<hsa_ven_amd_aqlprofile_info_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<rocprofiler_record_se_att_data_t*>(
+      calloc(data.size(), sizeof(rocprofiler_record_se_att_data_t)));
+
+  std::vector<hsa_ven_amd_aqlprofile_info_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();
+}
+
+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<std::vector<hsa_ven_amd_aqlprofile_info_data_t>*>(data);
+  passed_data->push_back(*info_data);
+  // TODO: clear output buffers after copying
+  // either copy here or in ::AddAttRecord
+
+  return status;
+}
+
+std::unordered_map<uint64_t, ATTRecordSignal> AttTracer::pending_stop_packets;
+std::mutex AttTracer::att_enable_disable_mutex;
+
+
 }  // namespace att
 
 }  // namespace rocprofiler

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

@@ -27,8 +27,15 @@
 #include <mutex>
 #include <string>
 #include <vector>
+#include <unordered_map>
+#include <atomic>
 
 #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<att_pending_signal_t>& 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<packet_t>& transformed_packets,
+    hsa_agent_t agent,
+    uint32_t data
+  );
+
+  void SetParameters(const std::vector<rocprofiler_att_parameter_t>& params) {
+    att_parameters_data = params;
+  }
+  void SetDispatchIds(const std::vector<std::pair<uint64_t,uint64_t>>& ids) {
+    kernel_profile_dispatch_ids = ids;
+  }
+  void SetCountersNames(const std::vector<std::string>& names) {
+    att_counters_names = names;
+  }
+  void SetKernelsNames(const std::vector<std::string>& names) {
+    kernel_profile_names = names;
+  }
+  std::optional<std::pair<size_t, size_t>> 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<uint64_t, ATTRecordSignal> 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<hsa_ven_amd_aqlprofile_profile_t*, rocprofiler_codeobj_capture_mode_t>
+  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<std::vector<bool>, 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<uint32_t> WRITER_ID{1};
+
+  std::vector<std::string> kernel_profile_names;
+  std::vector<std::pair<uint64_t,uint64_t>> kernel_profile_dispatch_ids;
+  std::vector<std::string> att_counters_names;
+  std::vector<rocprofiler_att_parameter_t> att_parameters_data;
 
   std::mutex sessions_pending_signals_lock_;
   std::map<uint32_t, std::vector<att_pending_signal_t>> sessions_pending_signals_;

projects/rocprofiler/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;
+};
+
+}

projects/rocprofiler/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 <cassert>
+#include <atomic>
+#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<packet_t>& 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<std::pair<size_t, size_t>> 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<const packet_t*>(packets);
+  std::vector<packet_t> 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<std::mutex> 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<std::mutex> 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<std::mutex> 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

projects/rocprofiler/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 <cassert>
+#include <atomic>
+#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<std::vector<bool>, bool>
+AttTracer::GetAllowedProfilesList(const void* packets, int pkt_count) {
+
+  std::vector<bool> 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<std::mutex> 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<const hsa_barrier_and_packet_t*>(packets)[i];
+    bool b_profile_this_object = false;
+
+    // Skip packets other than kernel dispatch packets.
+    if (Packet::IsDispatchPacket(original_packet)) {
+      auto& kdispatch = static_cast<const hsa_kernel_dispatch_packet_s*>(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<const packet_t*>(packets);
+  std::vector<packet_t> 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<bool> 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<const hsa_barrier_and_packet_t*>(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<hsa_kernel_dispatch_packet_t&>(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<packet_t*>(&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<packet_t*>(&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

projects/rocprofiler/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<std::vector<std::string>, uint32_t*, std::vector<uint64_t>> Filter::GetProperty(
-    rocprofiler_filter_property_kind_t kind) {
-  std::variant<std::vector<std::string>, uint32_t*, std::vector<uint64_t>> 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_;

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

@@ -33,6 +33,11 @@ namespace rocprofiler {
 
 class Filter {
  public:
+  typedef std::variant<
+    std::vector<std::string>,
+    uint32_t*,
+    std::vector<std::pair<uint64_t,uint64_t>>
+  > 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<std::vector<std::string>, uint32_t*, std::vector<uint64_t> > 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<rocprofiler_att_parameter_t> att_parameters_;        // ATT Parameters
   rocprofiler_counters_sampler_parameters_t
       counters_sampler_parameters_;  // sampled counters parameters
-  std::vector<uint64_t> dispatch_id_filter_;
+  std::vector<std::pair<uint64_t,uint64_t>> dispatch_id_filter_;
 
   bool has_sync_callback_{false};
   rocprofiler_sync_callback_t callback_;

projects/rocprofiler/src/tools/tool.cpp

@@ -266,9 +266,14 @@ std::vector<std::string> GetCounterNames() {
   return counters;
 }
 
-typedef std::tuple<std::vector<std::pair<rocprofiler_att_parameter_name_t, uint32_t>>,
-        std::vector<std::string>, std::vector<std::string>, std::vector<uint64_t>, uint64_t>
-    att_parsed_input_t;
+struct att_parsed_input_t
+{
+  std::vector<std::pair<rocprofiler_att_parameter_name_t, uint32_t>> params{};
+  std::vector<std::string> kernel_names{};
+  std::vector<std::string> counters_names{};
+  std::vector<std::pair<uint64_t, uint64_t>> dispatch_ids{};
+  rocprofiler::att_header_packet_t header{.raw = 0};
+};
 
 static int GetMpRank() {
   std::vector<const char*> 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<std::pair<rocprofiler_att_parameter_name_t, uint32_t>> parameters;
-  std::vector<std::string> kernel_names;
-  std::vector<std::string> counters_names;
-  std::vector<uint64_t> 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<std::string, rocprofiler_att_parameter_name_t> 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<rocprofiler_att_parameter_t> parameters;
-  std::vector<std::pair<rocprofiler_att_parameter_name_t, uint32_t>> params;
-  std::vector<std::string> kernel_names;
-  std::vector<std::string> att_counters_names;
-  std::vector<uint64_t> 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<const char*> 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<decltype(property.dispatch_ids)>(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();

projects/rocprofiler/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;
 }