SWDEV-396239: Automatic ISA dump from ATT

Change-Id: Ia66346c0048b779487157961ead08d7567c9153a

CHANGELOG.md

@@ -299,3 +299,5 @@ Example for file plugin output:
 ### Added
 
 - Updated supported GPU architectures in README with profiler versions
+- Automatic ISA dumping for ATT
+- CSV mode for ATT

README.md

@@ -306,56 +306,32 @@ Usage:
 - ATT (Advanced thread tracer) plugin: advanced hardware traces data in binary format. Please refer ATT section.
 Tool used to collect fine-grained hardware metrics. Provides ISA-level instruction hotspot analysis via hardware tracing.
 
-***
-Note: ATT(Advanced Thread Trace) needs some preparation before running.
-***
+  - Install plugin package. See Plugin Support section for installation
+  - Run the following to view the trace. Att-specific options must come right after the assembly file
 
-- Make sure to generate the assembly file for application by executing the following before compiling your HIP Application. This can be achieved globally by following environment variable
+    ```bash
+    rocprofv2 -i input.txt --plugin att <app_assembly_file> --mode network <app_relative_path>
+    ```
 
-  ```bash
-  export HIPCC_COMPILE_FLAGS_APPEND="--save-temps -g"
-  ```
-
-  Similarly, the --save-temps -g flags can be added per file for better ISA generation control.
-
-- Install plugin package. See Plugin Support section for installation
-- Run the following to view the trace. Att-specific options must come right after the assembly file
-
-  ```bash
-  rocprofv2 -i input.txt --plugin att <app_assembly_file> --mode network <app_relative_path>
-  ```
-
-- Example for vectoradd on navi31.
-
-  ***
-  Note: Special attention to gfx1100.s==navi31 in the ISA file name.
-  Use gfx1030 for navi21, gfx90a for MI200 and gfx940 for MI300
-  ***
-
-  ```bash
-  hipcc -g --save-temps vectoradd_hip.cpp -o vectoradd_hip.exe
-  rocprofv2 -i input.txt --plugin att vectoradd_hip-hip-amdgcn-amd-amdhsa-gfx1100.s --mode network ./vectoradd_hip.exe
-  ```
-
-  Then open the browser at `http://localhost:8000`
-  The ISA can also be obtained from llvm/roc objdump, however, annotations will be different
-
-  - app_assembly_file_relative_path
-      AMDGCN ISA file with .s extension generated in 1st step
+  - app_assembly_file: 
+    On ROCm 6.0, ATT enables automatic capture of the ISA during kernel execution, and does not require recompiling. It is recommeneded to leave at "auto".
   - app_relative_path
-      Path for the running application
+    Path for the running application
   - ATT plugin optional parameters
     - --depth [n]: How many waves per slot to parse (maximum).
     - --mpi [proc]: Parse with this many mpi processes, for greater analysis speed. Does not change results. Requires mpi4py.
     - --att_kernel "filename": Kernel filename to use (instead of ATT asking which one to use).
     - --trace_file "files": glob (wildcards allowed) of traces files to parse. Requires quotes for use with wildcards.
-    - --mode [network, file, off (default)]
+    - --mode [network, file, csv, off (default)]
       - network
           Opens the server with the browser UI.
           att needs 2 ports available (e.g. 8000, 18000). There is an option (default: --ports "8000,18000") to change these.
           In case rocprofv2 is running on a different machine, use port forwarding `ssh -L 8000:localhost:8000 <user@IP>` so the browser can be used locally. For docker, use --network=host --ipc=host -p8000:8000 -p18000:18000
       - file
           Dumps the analyzed json files to disk for vieweing at a later time. Run python3 httpserver.py from within the generated ui/ folder to view the trace, similarly to network mode. The folder can be copied to another machine, and will run without rocm.
+      - csv
+          Dumps the analyzed assembly into a CSV format, with the hitcount and total cycles cost.
+          Use rocprofv2's -o option to specify output file name (default "att_output.csv").
       - off
           Runs trace collection but not analysis, so it can be analyzed at a later time. Run rocprofv2 ATT [network, file] with the same parameters, removing the application binary, to analyze previously generated traces.
   - input.txt
@@ -385,23 +361,47 @@ Note: ATT(Advanced Thread Trace) needs some preparation before running.
     - PERFCOUNTERS_COL_PERIOD=0x3 // Multiplier period for counter collection [0~31]. 0=fastest (usually once every 16 cycles). GFX9 only. Counters will be shown in a graph over time in the browser UI.
     - PERFCOUNTER=counter_name // Add a SQ counter to be collected with ATT; period defined by PERFCOUNTERS_COL_PERIOD. GFX9 only.
     - BUFFER_SIZE=[size] // Sets size of the ATT buffer collection, per dispatch, in megabytes (shared among all shader engines).
+    - ISA_CAPTURE_MODE=[0,1,2] // Set capture mode during kernel dispatch.
+        - 0 = capture symbols only.
+        - 1 = capture symbols for file:// and make a copy of memory://
+        - 2 = Copy file:// and memory://
+
+  - Example for vectoradd.
+
+    ```bash
+    # -g add debugging symbols to the binary. Required only for tracking disassembly back to c++.
+    hipcc -g --save-temps vectoradd_hip.cpp -o vectoradd_hip.exe
+    # "auto" means to use the automatically captured ISA. "csv" dumps result to "my.csv".
+    rocprofv2 -i input.txt -o my.csv --plugin att auto --mode csv ./vectoradd_hip.exe
+    ```
+    Instruction latencies will be in my.csv
+
+  - Example with symbolic ISA (as in ROCm 5.7 or previous).
+    ```bash
+    hipcc -g --save-temps vectoradd_hip.cpp -o vectoradd_hip.exe
+    # A custom ISA can be used such as vectoradd_hip-hip-amdgcn-amd-amdhsa-gfx1100.s
+    # Special attention to the correct architecture for the ISA, such as "gfx1100" (navi31).
+    rocprofv2 -i input.txt --plugin att vectoradd_hip-hip-amdgcn-amd-amdhsa-gfx1100.s --mode network ./vectoradd_hip.exe
+    ```
+
+    Then open the browser at `http://localhost:8000`
 
   ***
   Note: For MPI or long running applications, we recommend to run collection, and later run the parser with already collected data:
-   Run only collection: The assembly file is not used. Use mpirun [...] rocprofv2 [...] if needed.
+  Run only collection: The assembly file is not used. Use mpirun [...] rocprofv2 [...] if needed.
 
   ```bash
-   rocprofv2 -i input.txt --plugin att none ./vectoradd_hip.exe
+  # Run only data collection, not the parser
+  rocprofv2 -i input.txt --plugin att auto --mode off ./vectoradd_hip.exe
   ```
-
-  Remove the binary/application: Only runs the parser.
-  
+  Remove the binary/application from the command line.
   ```bash
-  rocprofv2 -i input.txt --plugin att vectoradd_hip-hip-amdgcn-amd-amdhsa-gfx1100.s --mode network
+  # Only runs the parser on previously collected data.
+  rocprofv2 -i input.txt --plugin att auto --mode network
   ```
 
   ***
-  
+
 ### Flush Interval
 
 Flush interval can be used to control the interval time in milliseconds between the buffers flush for the tool. However, if the buffers are full the flush will be called on its own. This can be used as in the next example:

include/rocprofiler/v2/rocprofiler.h

@@ -1173,6 +1173,75 @@ typedef struct {
   uint32_t buffer_size;
 } rocprofiler_record_se_att_data_t;
 
+/**
+ * struct to store the filepaths and their addresses for intercepted code objects
+*/
+typedef struct {
+  /**
+   * File path (file://, memory://) of the code object
+  */
+  const char* filepath;
+  /**
+   * Addr where codeobj is loaded
+  */
+  uint64_t base_address;
+  /**
+   * If a copy of the codeobj is made, contains the data. Nullptr otherwise.
+  */
+  const char* data;
+  /**
+   * If a copy of the codeobj is made, contains the size of the data. 0 otherwise.
+  */
+  uint64_t size;
+  /**
+   * Timestamp for the time point this codeobj was loaded.
+  */
+  rocprofiler_timestamp_t clock_start;
+  /**
+   * Timestamp for the time point this codeobj was unloaded.
+   * If the obj is still loaded by the time the record was generated, this value is 0
+  */
+  rocprofiler_timestamp_t clock_end;
+} rocprofiler_intercepted_codeobj_t;
+
+/**
+ * Enum defines how code object is captured for ATT and PC Sampling
+*/
+typedef enum {
+  /**
+   * Capture file and memory paths for the loaded code object
+  */
+  ROCPROFILER_CAPTURE_SYMBOLS_ONLY = 0,
+  /**
+   * Capture symbols for file:// and memory:// type objects,
+   * and generate a copy of all kernel code for objects under memory://
+  */
+  ROCPROFILER_CAPTURE_COPY_MEMORY = 1,
+  /**
+   * Capture symbols and all kernel code for file:// and memory:// type objects
+  */
+  ROCPROFILER_CAPTURE_COPY_FILE_AND_MEMORY = 2
+} rocprofiler_codeobj_capture_mode_t;
+
+/**
+ * struct to store the filepaths and their addresses for intercepted code objects
+*/
+typedef struct {
+  /**
+   * List of symbols
+  */
+  const rocprofiler_intercepted_codeobj_t* symbols;
+  /**
+   * Number of symbols
+  */
+  uint64_t count;
+  /**
+   * Userdata space for custom capture.
+   * For ATT records, it is the address of the kernel being launched.
+  */
+  uint64_t userdata;
+} rocprofiler_codeobj_symbols_t;
+
 /**
  * ATT tracing record structure.
  * This will represent all the information reported by the
@@ -1214,6 +1283,10 @@ typedef struct {
    * Queue Index - packet index in the queue
    */
   rocprofiler_queue_index_t queue_idx;
+  /**
+   * Writer ID for counting how many kernels
+   */
+  uint64_t writer_id;
   /**
    * ATT data output from each shader engine.
    */
@@ -1222,6 +1295,10 @@ typedef struct {
    * The count of the shader engine ATT data
    */
   uint64_t shader_engine_data_count;
+  /**
+   * Filepaths for the intercepted code objects at the time of kernel dispatch
+  */
+  rocprofiler_codeobj_symbols_t intercept_list;
 } rocprofiler_record_att_tracer_t;
 
 
@@ -1650,7 +1727,7 @@ typedef enum {
    */
   ROCPROFILER_PC_SAMPLING_COLLECTION = 3,
   /**
-   * ATT Tracing. (Not Yet Supported)
+   * ATT Tracing.
    */
   ROCPROFILER_ATT_TRACE_COLLECTION = 4,
   /**
@@ -1707,16 +1784,50 @@ typedef const char* rocprofiler_hsa_function_name_t;
  * ATT parameters to be used by for collection
  */
 typedef enum {
-  ROCPROFILER_ATT_COMPUTE_UNIT = 0,   //! Select the target compute unit (wgp) for profiling.
-  ROCPROFILER_ATT_VMID_MASK = 1,      //! VMID Mask.
-  ROCPROFILER_ATT_SE_MASK = 5,        //! Shader engine mask for selection.
-  ROCPROFILER_ATT_SIMD_SELECT = 8,    //! Set SIMD Mask (GFX9) or SIMD ID for collection (Navi)
-  ROCPROFILER_ATT_OCCUPANCY = 9,      //! Set true for occupancy collection only.
-  ROCPROFILER_ATT_BUFFER_SIZE = 10,   //! ATT collection max data size, in MB. Shared among shader engines.
-  ROCPROFILER_ATT_PERF_MASK = 240,    //! Mask of which compute units to generate perfcounters. GFX9 only.
-  ROCPROFILER_ATT_PERF_CTRL = 241,    //! Select collection period for perfcounters. GFX9 only.
-  ROCPROFILER_ATT_PERFCOUNTER = 242,  //! Select perfcounter ID (SQ block) for collection. GFX9 only.
-  ROCPROFILER_ATT_PERFCOUNTER_NAME = 243, //! Select perfcounter name (SQ block) for collection. GFX9 only.
+  /**
+   * Select the target compute unit (wgp) for profiling.
+  */
+  ROCPROFILER_ATT_COMPUTE_UNIT = 0,
+  /**
+   * VMID Mask
+  */
+  ROCPROFILER_ATT_VMID_MASK = 1,
+  /**
+   * Shader engine mask for selection.
+  */
+  ROCPROFILER_ATT_SE_MASK = 5,
+  /**
+   * Set SIMD Mask (GFX9) or SIMD ID for collection (Navi)
+  */
+  ROCPROFILER_ATT_SIMD_SELECT = 8,
+  /**
+   * Set true for occupancy collection only.
+  */
+  ROCPROFILER_ATT_OCCUPANCY = 9,
+  /**
+   * ATT collection max data size, in MB. Shared among shader engines.
+  */
+  ROCPROFILER_ATT_BUFFER_SIZE = 10,
+  /**
+   * Set ISA capture during ATT collection (rocprofiler_codeobj_capture_mode_t)
+  */
+  ROCPROFILER_ATT_CAPTURE_MODE = 11,
+  /**
+   * Mask of which compute units to generate perfcounters. GFX9 only.
+  */
+  ROCPROFILER_ATT_PERF_MASK = 240,
+  /**
+   * Select collection period for perfcounters. GFX9 only.
+  */
+  ROCPROFILER_ATT_PERF_CTRL = 241,
+  /**
+   * Select perfcounter ID (SQ block) for collection. GFX9 only.
+  */
+  ROCPROFILER_ATT_PERFCOUNTER = 242,
+  /**
+   * Select perfcounter name (SQ block) for collection. GFX9 only.
+  */
+  ROCPROFILER_ATT_PERFCOUNTER_NAME = 243,
   ROCPROFILER_ATT_MAXVALUE,
 
   ROCPROFILER_ATT_MASK = 2,           //! Deprecated
@@ -2230,6 +2341,62 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_device_profiling_session_stop(
 ROCPROFILER_API rocprofiler_status_t rocprofiler_device_profiling_session_destroy(
     rocprofiler_session_id_t session_id) ROCPROFILER_VERSION_9_0;
 
+/**
+ * Creates a codeobj capture record, returned in ID.
+ * \param[out] id contains a handle for the created record.
+ * \param[in] mode Set to capture symbols only,
+ *                  make a copy of codeobj under memory://
+ *                  or copy all codeobj.
+ * \param[in] userdata userdata to be returned in the record. For ATT records, is the kernel addr.
+ * \retval ::ROCPROFILER_STATUS_SUCCESS The function has been executed successfully.
+ */
+ROCPROFILER_API rocprofiler_status_t
+rocprofiler_codeobj_capture_create(
+  rocprofiler_record_id_t* id,
+  rocprofiler_codeobj_capture_mode_t mode,
+  uint64_t userdata
+);
+
+/**
+ * API to get the captured codeobj.
+ * Each call invalidates the previous pointer for the same ID.
+ * \param[in] id record handle.
+ * \param[out] capture captured code objects.
+ * \retval ::ROCPROFILER_STATUS_SUCCESS The function has been executed successfully.
+ * \retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS invalid ID.
+ */
+ROCPROFILER_API rocprofiler_status_t
+rocprofiler_codeobj_capture_get(rocprofiler_record_id_t id,
+                                rocprofiler_codeobj_symbols_t* capture);
+
+/**
+ * API to delete a record.
+ * Invalidates the pointer returned from rocprofiler_codeobj_capture_get.
+ * \param[in] id record handle.
+ * \retval ::ROCPROFILER_STATUS_SUCCESS The function has been executed successfully.
+ */
+ROCPROFILER_API rocprofiler_status_t
+rocprofiler_codeobj_capture_free(rocprofiler_record_id_t id);
+
+/**
+ * Records the current loaded codeobjs and any following loads until stop() is called.
+ * \param[in] id record handle.
+ * \retval ::ROCPROFILER_STATUS_SUCCESS The function has been executed successfully.
+ * \retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS invalid ID.
+ */
+ROCPROFILER_API rocprofiler_status_t
+rocprofiler_codeobj_capture_start(rocprofiler_record_id_t id);
+
+/**
+ * Stops recording of future codeobjs, until start() is called again.
+ * Calling stop() immediately after a start() snapshots the current state of loaded codeobjs.
+ * \param[in] id record handle.
+ * \retval ::ROCPROFILER_STATUS_SUCCESS The function has been executed successfully.
+ * \retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS invalid ID.
+ */
+ROCPROFILER_API rocprofiler_status_t
+rocprofiler_codeobj_capture_stop(rocprofiler_record_id_t id);
+
 /** @} */
 
 #ifdef __cplusplus

plugin/att/CMakeLists.txt

@@ -26,7 +26,7 @@ set(ENV{ROCPROFV2_ATT_LIB_PATH} $ROCPROFV2_ATT)
 
 # Building att plugin library
 file(GLOB ROCPROFILER_UTIL_SRC_FILES ${PROJECT_SOURCE_DIR}/src/utils/helper.cpp)
-file(GLOB FILE_SOURCES att.cpp)
+file(GLOB FILE_SOURCES att.cpp disassembly.cpp code_printing.cpp)
 add_library(att_plugin SHARED ${FILE_SOURCES} ${ROCPROFILER_UTIL_SRC_FILES})
 
 set_target_properties(
@@ -44,8 +44,7 @@ target_include_directories(att_plugin PRIVATE ${PROJECT_SOURCE_DIR}
 target_link_options(
     att_plugin PRIVATE -Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/../exportmap
     -Wl,--no-undefined)
-target_link_libraries(att_plugin PRIVATE rocprofiler-v2 hsa-runtime64::hsa-runtime64
-                                         stdc++fs)
+target_link_libraries(att_plugin PRIVATE rocprofiler-v2 hsa-runtime64::hsa-runtime64 stdc++fs dw elf amd_comgr)
 
 install(TARGETS att_plugin LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR}/${PROJECT_NAME}
                                    COMPONENT asan)
@@ -56,6 +55,7 @@ configure_file(att.py att/att.py COPYONLY)
 configure_file(trace_view.py att/trace_view.py COPYONLY)
 configure_file(stitch.py att/stitch.py COPYONLY)
 configure_file(drawing.py att/drawing.py COPYONLY)
+configure_file(att_to_csv.py att/att_to_csv.py COPYONLY)
 configure_file(ui/index.html att/ui/index.html COPYONLY)
 configure_file(ui/logo.svg att/ui/logo.svg COPYONLY)
 configure_file(ui/styles.css att/ui/styles.css COPYONLY)

plugin/att/att.cpp

@@ -44,8 +44,11 @@
 #include "rocprofiler.h"
 #include "rocprofiler_plugin.h"
 #include "../utils.h"
+#include "code_printing.hpp"
+
 
 #define ATT_FILENAME_MAXBYTES 90
+#define TEST_INVALID_KERNEL size_t(-1)
 
 namespace {
 
@@ -74,24 +77,33 @@ class att_plugin_t {
 
   bool IsValid() const { return is_valid_; }
 
-  void FlushATTRecord(const rocprofiler_record_att_tracer_t* att_tracer_record,
-                      rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) {
+  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) {
-      printf("No att data buffer received\n");
-      return;
+    if (!att_tracer_record) return ROCPROFILER_STATUS_ERROR;
+
+    std::string kernel_name_mangled{};
+    // Found problem with rocprofiler API for invalid kernel_ids;
+    if (att_tracer_record->kernel_id.handle != TEST_INVALID_KERNEL) {
+      size_t name_length;
+      CHECK_ROCPROFILER(rocprofiler_query_kernel_info_size(
+          ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, &name_length));
+      const char* kernel_name_c = nullptr;
+      CHECK_ROCPROFILER(rocprofiler_query_kernel_info(
+          ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, &kernel_name_c));
+
+      assert(kernel_name_c && "Rocprofv2 returned an invalid kernel name");
+
+      kernel_name_mangled = std::string(kernel_name_c);
+      free(const_cast<char*>(kernel_name_c));
+    } else {  // Temporary. Adding a valid string.
+      kernel_name_mangled = "test_kernel";
     }
 
-    size_t name_length;
-    CHECK_ROCPROFILER(rocprofiler_query_kernel_info_size(
-        ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, &name_length));
-    const char* kernel_name_c = static_cast<const char*>(malloc(name_length * sizeof(char)));
-    CHECK_ROCPROFILER(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME,
-                                                    att_tracer_record->kernel_id, &kernel_name_c));
 
     std::string name_demangled =
-        rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c));
+        rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_mangled));
 
     if (name_demangled.size() > ATT_FILENAME_MAXBYTES)  // Limit filename size
       name_demangled = name_demangled.substr(0, ATT_FILENAME_MAXBYTES);
@@ -110,11 +122,11 @@ class att_plugin_t {
       file_iteration += 1;
 
     outfilepath += std::to_string(file_iteration);
-    auto dispatch_id = att_tracer_record->header.id.handle;
+    auto writer_id = att_tracer_record->writer_id;
 
     std::string fname = outfilepath + "_kernel.txt";
-    std::ofstream(fname.c_str()) << name_demangled << " dispatch[" << dispatch_id << "] GPU["
-                                 << att_tracer_record->gpu_id.handle << "]: " << kernel_name_c
+    std::ofstream(fname.c_str()) << name_demangled << " dispatch[" << writer_id << "] GPU["
+                                 << att_tracer_record->gpu_id.handle << "]: " << kernel_name_mangled
                                  << '\n';
 
     // iterate over each shader engine att trace
@@ -125,29 +137,69 @@ class att_plugin_t {
         continue;
       printf("--------------collecting data for shader_engine %d---------------\n", i);
       rocprofiler_record_se_att_data_t* se_att_trace = &att_tracer_record->shader_engine_data[i];
-      const char* data_buffer_ptr = reinterpret_cast<char*>(se_att_trace->buffer_ptr);
+      char* data_buffer_ptr = reinterpret_cast<char*>(se_att_trace->buffer_ptr);
 
       // dump data in binary format
       std::ofstream out(outfilepath + "_se" + std::to_string(i) + ".att", std::ios::binary);
-      if (out.is_open())
-        out.write((char*)data_buffer_ptr, se_att_trace->buffer_size);
-      else
+      if (!out.is_open()) {
         std::cerr << "ATT Failed to open file: " << outfilepath << "_se" << i << ".att\n";
+        return ROCPROFILER_STATUS_ERROR;
+      }
+      out.write(data_buffer_ptr, se_att_trace->buffer_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;
+    }
+    uint64_t kernel_begin_addr = att_tracer_record->intercept_list.userdata;
+    isafile << "<Kernel> " << kernel_name_mangled << '\n';
+
+    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];
+
+      std::unique_ptr<CodeObjectBinary> binary;
+      std::unique_ptr<code_object_decoder_t> decoder;
+
+      if (symbol.data && symbol.size) {
+        decoder = std::make_unique<code_object_decoder_t>(symbol.data, symbol.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;
+      }
+
+      for (auto& instance : decoder->instructions) {
+        uint64_t addr = instance.address + symbol.base_address;
+
+        if (kernel_begin_addr == addr)
+          isafile << "; Begin <Kernel> " << kernel_name_mangled << '\n';
+        else if (decoder->m_symbol_map.find(instance.address) != decoder->m_symbol_map.end())
+          isafile << "; Begin " << decoder->m_symbol_map[instance.address].first << '\n';
+        if (instance.cpp_reference) isafile << "; " << instance.cpp_reference << '\n';
+        isafile << instance.instruction << " // " << std::hex << addr << '\n';
+      }
+    }
+
+    return ROCPROFILER_STATUS_SUCCESS;
   }
 
   int WriteBufferRecords(const rocprofiler_record_header_t* begin,
                          const rocprofiler_record_header_t* end,
                          rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) {
     while (begin < end) {
-      if (!begin) return 0;
+      if (!begin) return ROCPROFILER_STATUS_ERROR;
       switch (begin->kind) {
         case ROCPROFILER_PROFILER_RECORD:
         case ROCPROFILER_TRACER_RECORD:
         case ROCPROFILER_PC_SAMPLING_RECORD:
         case ROCPROFILER_SPM_RECORD:
         case ROCPROFILER_COUNTERS_SAMPLER_RECORD:
-          printf("Invalid record Kind: %d", begin->kind);
+          rocprofiler::warning("Invalid record Kind: %d\n", begin->kind);
           break;
 
         case ROCPROFILER_ATT_TRACER_RECORD: {
@@ -158,10 +210,11 @@ class att_plugin_t {
           break;
         }
       }
-      rocprofiler_next_record(begin, &begin, session_id, buffer_id);
+      int status = rocprofiler_next_record(begin, &begin, session_id, buffer_id);
+      if (status != ROCPROFILER_STATUS_SUCCESS) return status;
     }
 
-    return 0;
+    return ROCPROFILER_STATUS_SUCCESS;
   }
 
  private:
@@ -176,17 +229,17 @@ ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_
                                                      void* data) {
   if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR ||
       rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR)
-    return -1;
+    return ROCPROFILER_STATUS_ERROR;
 
-  if (att_plugin != nullptr) return -1;
+  if (att_plugin != nullptr) return ROCPROFILER_STATUS_ERROR;
 
   att_plugin = new att_plugin_t();
-  if (att_plugin->IsValid()) return 0;
+  if (att_plugin->IsValid()) return ROCPROFILER_STATUS_SUCCESS;
 
   // The plugin failed to initialied, destroy it and return an error.
   delete att_plugin;
   att_plugin = nullptr;
-  return -1;
+  return ROCPROFILER_STATUS_ERROR;
 }
 
 ROCPROFILER_EXPORT void rocprofiler_plugin_finalize() {
@@ -198,12 +251,12 @@ 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) {
-  if (!att_plugin || !att_plugin->IsValid()) return -1;
+  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) {
-  if (!att_plugin || !att_plugin->IsValid()) return -1;
-  if (record.header.id.handle == 0) return 0;
-  return 0;
+  if (!att_plugin || !att_plugin->IsValid()) return ROCPROFILER_STATUS_ERROR;
+  if (record.header.id.handle == 0) return ROCPROFILER_STATUS_SUCCESS;
+  return ROCPROFILER_STATUS_SUCCESS;
 }

plugin/att/att.py

@@ -40,25 +40,26 @@ class PerfEvent(ctypes.Structure):
 class CodeWrapped(ctypes.Structure):
     """ Matches CodeWrapped on the python side """
     _fields_ = [('line', ctypes.c_char_p),
-                            ('loc', ctypes.c_char_p),
-                            ('value', ctypes.c_int),
-                            ('to_line', ctypes.c_int),
-                            ('index', ctypes.c_int),
-                            ('line_num', ctypes.c_int)]
+                ('loc', ctypes.c_char_p),
+                ('to_line', ctypes.c_int),
+                ('value', ctypes.c_int),
+                ('index', ctypes.c_int),
+                ('line_num', ctypes.c_int),
+                ('addr', ctypes.c_int64)]
 
 
 class KvPair(ctypes.Structure):
     """ Matches pair<int, int> = (key, value) on the python side """
     _fields_ = [('key', ctypes.c_int),
-                            ('value', ctypes.c_int)]
+               ('value', ctypes.c_int)]
 
 
 class ReturnAssemblyInfo(ctypes.Structure):
     """ Matches ReturnAssemblyInfo on the python side """
     _fields_ = [('code', POINTER(CodeWrapped)),
-                            ('jumps', POINTER(KvPair)),
-                            ('code_len', ctypes.c_int),
-                            ('jumps_len', ctypes.c_int)]
+                ('jumps', POINTER(KvPair)),
+                ('code_len', ctypes.c_int),
+                ('jumps_len', ctypes.c_int)]
 
 
 class Wave(ctypes.Structure):
@@ -155,8 +156,9 @@ def parse_binary(filename, kernel=None):
         to_line = int(code_entry.to_line) if (code_entry.to_line >= 0) else None
         loc = loc if len(loc) > 0 else None
 
-        code.append([line, int(code_entry.value), to_line, loc,
-                    int(code_entry.index), int(code_entry.line_num), 0, 0]) # hitcount + cycles
+        # asm, inst_type, addr, loc, index, line_num, hitcount, cycles
+        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])
 
     jumps = {}
     for k in range(info.jumps_len):
@@ -187,9 +189,9 @@ def getWaves_binary(name, shader_engine_data_dict, target_cu, depth):
     shader_engine_data_dict[name] = (waves_python, events, occupancy, flags)
 
 
-def getWaves_stitch(SIMD, code, jumps, flags, latency_map, hitcount_map):
+def getWaves_stitch(SIMD, code, jumps, flags, latency_map, hitcount_map, bIsAuto):
     for pwave in SIMD:
-        pwave.instructions = stitch(pwave.instructions, code, jumps, flags)
+        pwave.instructions = stitch(pwave.instructions, code, jumps, flags, bIsAuto)
 
         for inst in pwave.instructions[0]:
             hitcount_map[inst[-1]] += 1
@@ -363,7 +365,10 @@ if __name__ == "__main__":
                         network: Open att server over the network.''', type=str, default="off")
     args = parser.parse_args()
 
-    if args.mode.lower() == 'file':
+    CSV_MODE = False
+    if args.mode.lower() == 'csv':
+        CSV_MODE = True
+    elif args.mode.lower() == 'file':
         args.dumpfiles = True
     elif args.mode.lower() == 'network':
         args.dumpfiles = False
@@ -387,12 +392,6 @@ if __name__ == "__main__":
     if args.target_cu is None:
         args.target_cu = 1
 
-    # Assembly parsing
-    path = Path(args.assembly_code)
-    if not path.is_file():
-        print("Invalid assembly_code('{0}')!".format(args.assembly_code))
-        sys.exit(1)
-
     att_kernel = glob.glob(args.att_kernel)
 
     if len(att_kernel) == 0:
@@ -418,6 +417,16 @@ if __name__ == "__main__":
     else:
         args.att_kernel = att_kernel[0]
 
+    # Assembly parsing
+    bIsAuto = False
+    if args.assembly_code.lower().strip() == 'auto':
+        args.assembly_code = args.att_kernel.split('_kernel.txt')[0]+'_isa.s'
+        bIsAuto = True
+    path = Path(args.assembly_code)
+    if not path.is_file():
+        print("Invalid assembly_code('{0}')!".format(args.assembly_code))
+        sys.exit(1)
+
     # Trace Parsing
     if args.trace_file is None:
         filenames = glob.glob(args.att_kernel.split('_kernel.txt')[0]+'_*.att')
@@ -431,7 +440,7 @@ if __name__ == "__main__":
     code = jumps = None
     if mpi_root:
         print('Att kernel:', args.att_kernel)
-        code, jumps = parse_binary(args.assembly_code, args.att_kernel)
+        code, jumps = parse_binary(args.assembly_code, None if bIsAuto else args.att_kernel)
 
     DBFILES = []
     TIMELINES = [np.zeros(int(1E4),dtype=np.int16) for k in range(5)]
@@ -460,7 +469,7 @@ if __name__ == "__main__":
         if np.sum([0]+[len(s.instructions) for s in SIMD]) == 0:
             print("No waves from", name)
             continue
-        getWaves_stitch(SIMD, code, jumps, gfxv, latency_map, hitcount_map)
+        getWaves_stitch(SIMD, code, jumps, gfxv, latency_map, hitcount_map, bIsAuto)
 
         analysed_filenames.append(name)
         EVENTS.append(perfevents)
@@ -524,6 +533,12 @@ if __name__ == "__main__":
             code[k][-2] = int(hitcount_map[k])
             code[k][-1] = int(latency_map[k])
 
+    if CSV_MODE:
+        if mpi_root:
+            from att_to_csv import dump_csv
+            dump_csv(code)
+        quit()
+
     gc.collect()
     print("Min time:", min_event_time)
 

plugin/att/att_to_csv.py

@@ -0,0 +1,18 @@
+#!/usr/bin/env python3
+
+import numpy as np
+import csv
+import os
+
+def dump_csv(code):
+    outpath = os.getenv("OUT_FILE_NAME")
+    if outpath is None:
+        outpath = "att_output.csv"
+    if ".csv" not in outpath:
+        outpath += ".csv"
+
+    with open(outpath, 'w') as f:
+        writer = csv.writer(f)
+        writer.writerow(['Line', 'Instruction', 'Hitcount', 'Cycles', 'Addr', 'C++ Reference'])
+        [writer.writerow([m[5], m[0], m[7], m[8], hex(m[6]), m[3]]) for m in code]
+        #[writer.writerow(m) for m in code]

plugin/att/code_printing.cpp

@@ -0,0 +1,226 @@
+/* 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 "code_printing.hpp"
+
+#include <algorithm>
+#include <fstream>
+#include <iomanip>
+#include <iostream>
+#include <map>
+#include <memory>
+#include <mutex>
+#include <optional>
+#include <string>
+#include <type_traits>
+#include <unordered_map>
+#include <vector>
+
+#include <cstdarg>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+#include <sys/types.h>
+#include <sys/stat.h>
+#include <fcntl.h>
+#include <unistd.h>
+
+#include <hsa/amd_hsa_elf.h>
+#include "../utils.h"
+#include <cxxabi.h>
+#include <elfutils/libdw.h>
+#include <sys/mman.h>
+
+code_object_decoder_t::code_object_decoder_t(const char* codeobj_data, uint64_t codeobj_size) {
+  buffer = std::vector<char>{};
+  buffer.resize(codeobj_size);
+  std::memcpy(buffer.data(), codeobj_data, codeobj_size);
+
+  m_fd = -1;
+#if defined(_GNU_SOURCE) && defined(MFD_ALLOW_SEALING) && defined(MFD_CLOEXEC)
+  m_fd = ::memfd_create(m_uri.c_str(), MFD_ALLOW_SEALING | MFD_CLOEXEC);
+#endif
+  if (m_fd == -1) // If fail, attempt under /tmp
+    m_fd = ::open("/tmp", O_TMPFILE | O_RDWR, 0666);
+
+  if (m_fd == -1) {
+    printf("could not create a temporary file for code object\n");
+    return;
+  }
+
+  if (size_t size = ::write(m_fd, buffer.data(), buffer.size()); size != buffer.size()) {
+    printf("could not write to the temporary file\n");
+    return;
+  }
+  ::lseek(m_fd, 0, SEEK_SET);
+  fsync(m_fd);
+
+  m_line_number_map = {};
+
+  std::unique_ptr<Dwarf, void (*)(Dwarf*)> dbg(dwarf_begin(m_fd, DWARF_C_READ),
+                                               [](Dwarf* dbg) { dwarf_end(dbg); });
+
+  /*if (!dbg) {
+    rocprofiler::warning("Error opening Dwarf!\n");
+    return;
+  } */
+
+  if (dbg) {
+    Dwarf_Off cu_offset{0}, next_offset;
+    size_t header_size;
+
+    while (!dwarf_nextcu(dbg.get(), cu_offset, &next_offset, &header_size, nullptr, nullptr,
+                         nullptr)) {
+      Dwarf_Die die;
+      if (!dwarf_offdie(dbg.get(), cu_offset + header_size, &die)) continue;
+
+      Dwarf_Lines* lines;
+      size_t line_count;
+      if (dwarf_getsrclines(&die, &lines, &line_count)) continue;
+
+      for (size_t i = 0; i < line_count; ++i) {
+        Dwarf_Addr addr;
+        int line_number;
+
+        if (Dwarf_Line* line = dwarf_onesrcline(lines, i))
+          if (!dwarf_lineaddr(line, &addr) && !dwarf_lineno(line, &line_number) && line_number) {
+            m_line_number_map.emplace(
+                addr, std::make_pair(dwarf_linesrc(line, nullptr, nullptr), line_number));
+          }
+      }
+      cu_offset = next_offset;
+    }
+    // load_symbol_map();
+  }
+  disassemble_kernels();
+}
+
+
+code_object_decoder_t::~code_object_decoder_t() {
+  if (m_fd) ::close(m_fd);
+}
+
+std::optional<code_object_decoder_t::symbol_info_t> code_object_decoder_t::find_symbol(
+    uint64_t address) {
+  /* Load the symbol table.  */
+  if (auto it = m_symbol_map.upper_bound(address); it != m_symbol_map.begin()) {
+    if (auto&& [symbol_value, symbol] = *std::prev(it); address < (symbol_value + symbol.second)) {
+      std::string symbol_name = symbol.first;
+
+      if (int status; auto* demangled_name =
+                          abi::__cxa_demangle(symbol_name.c_str(), nullptr, nullptr, &status)) {
+        symbol_name = demangled_name;
+        free(demangled_name);
+      }
+      return symbol_info_t{std::move(symbol_name), symbol_value, symbol.second};
+    }
+  }
+  return {};
+}
+
+/*
+void code_object_decoder_t::load_symbol_map() {
+  std::unique_ptr<Elf, void (*)(Elf *)> elf (
+      elf_begin(m_fd, ELF_C_READ, nullptr),
+      [](Elf *elf){ elf_end(elf); });
+
+  if (!elf) {
+    rocprofiler::warning("Error opening ELF!\n");
+    return;
+  }
+
+  Elf64_Ehdr *ehdr = elf64_getehdr(elf.get());
+  if (!ehdr) {
+    printf("elf64_getehdr failed\n");
+    return;
+  }
+
+  // Slurp the symbol table.
+  Elf_Scn *scn = nullptr;
+  while ((scn = elf_nextscn(elf.get(), scn)) != nullptr) {
+    GElf_Shdr shdr_mem;
+    GElf_Shdr *shdr = gelf_getshdr(scn, &shdr_mem);
+    if (shdr->sh_type != SHT_SYMTAB && shdr->sh_type != SHT_DYNSYM) {
+      continue;
+    }
+
+    Elf_Data *data = elf_getdata(scn, nullptr);
+    if (!data) continue;
+
+    size_t symbol_count = data->d_size / gelf_fsize(elf.get(), ELF_T_SYM, 1, EV_CURRENT);
+    for (size_t j = 0; j < symbol_count; ++j) {
+      GElf_Sym sym_mem;
+      GElf_Sym *sym = gelf_getsym(data, j, &sym_mem);
+
+      if (GELF_ST_TYPE(sym->st_info) != STT_FUNC || sym->st_shndx == SHN_UNDEF) continue;
+
+      std::string symbol_name{ elf_strptr(elf.get(), shdr->sh_link, sym->st_name) };
+      auto symbol_pair = std::make_pair(symbol_name, sym->st_size);
+
+      auto [it, success] = m_symbol_map.emplace(sym->st_value, symbol_pair);
+
+      // If there already was a symbol defined at this address, but this
+      // new symbol covers a larger address range, replace the old symbol
+      //         with this new one.
+      if (!success && sym->st_size > it->second.second) it->second = symbol_pair;
+    }
+  }
+} */
+
+void code_object_decoder_t::disassemble_kernel(uint64_t addr) {
+  auto symbol = find_symbol(addr);
+
+  if (!symbol) {
+    std::cerr << "No symbol found at address 0x" << std::hex << addr << std::endl;
+    return;
+  }
+
+  // if (symbol->m_name.find("__amd_rocclr_") == 0)
+  //  return;
+
+  std::cout << "Dumping ISA for " << symbol->m_name << std::endl;
+
+  uint64_t end_addr = addr + symbol->m_size;
+  while (addr < end_addr) {
+    char* cpp_line = nullptr;
+    auto it = m_line_number_map.find(addr);
+    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(addr, cpp_line);
+    addr += size;
+  }
+}
+
+void code_object_decoder_t::disassemble_kernels() {
+  disassembly = std::make_unique<DisassemblyInstance>(*this);
+
+  // if (m_symbol_map.begin() == m_symbol_map.end())
+  m_symbol_map = disassembly->GetKernelMap();
+
+  for (auto& [k, v] : m_symbol_map) disassemble_kernel(k);
+}

plugin/att/code_printing.hpp

@@ -0,0 +1,56 @@
+/* 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
+
+#include <map>
+#include <optional>
+#include <string>
+#include <vector>
+#include "rocprofiler.h"
+#include <memory>
+#include "disassembly.hpp"
+
+class code_object_decoder_t {
+  struct symbol_info_t {
+    const std::string m_name;
+    uint64_t m_value;
+    uint64_t m_size;
+  };
+
+ public:
+  // void load_symbol_map();
+  std::optional<symbol_info_t> find_symbol(uint64_t address);
+
+  code_object_decoder_t(const char* codeobj_data, uint64_t codeobj_size);
+  ~code_object_decoder_t();
+
+  void disassemble_kernel(uint64_t addr);
+  void disassemble_kernels();
+  int m_fd;
+
+  std::map<uint64_t, std::pair<std::string, size_t>> m_line_number_map;
+  std::map<uint64_t, std::pair<std::string, uint64_t>> m_symbol_map;
+
+  std::string m_uri;
+  std::vector<char> buffer;
+  std::vector<instruction_instance_t> instructions;
+  std::unique_ptr<DisassemblyInstance> disassembly;
+};

plugin/att/disassembly.cpp

@@ -0,0 +1,228 @@
+/* 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. */
+
+#if !defined(_GNU_SOURCE) || !defined(_XOPEN_SOURCE)
+#define _XOPEN_SOURCE 700
+#endif
+
+#include "code_printing.hpp"
+
+#include <algorithm>
+#include <fstream>
+#include <iomanip>
+#include <iostream>
+#include <map>
+#include <memory>
+#include <mutex>
+#include <optional>
+#include <string>
+#include <type_traits>
+#include <unordered_map>
+#include <vector>
+
+#include <cstdarg>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+#include <sys/mman.h>
+#include <sys/types.h>
+#include <sys/stat.h>
+#include <fcntl.h>
+#include <unistd.h>
+
+#include <hsa/amd_hsa_elf.h>
+#include "../utils.h"
+#include <cxxabi.h>
+#include <elfutils/libdw.h>
+
+#define CHECK_COMGR(call)                                                                          \
+  if (amd_comgr_status_s status = call) {                                                          \
+    const char* reason = "";                                                                       \
+    amd_comgr_status_string(status, &reason);                                                      \
+    std::cerr << __LINE__ << " failed: " << reason << std::endl;                                   \
+    return;                                                                                        \
+  }
+
+CodeObjectBinary::CodeObjectBinary(const std::string& uri) : m_uri(uri) {
+  const std::string protocol_delim{"://"};
+
+  size_t protocol_end = m_uri.find(protocol_delim);
+  std::string protocol = m_uri.substr(0, protocol_end);
+  protocol_end += protocol_delim.length();
+
+  std::transform(protocol.begin(), protocol.end(), protocol.begin(),
+                 [](unsigned char c) { return std::tolower(c); });
+
+  std::string path;
+  size_t path_end = m_uri.find_first_of("#?", protocol_end);
+  if (path_end != std::string::npos) {
+    path = m_uri.substr(protocol_end, path_end++ - protocol_end);
+  } else {
+    path = m_uri.substr(protocol_end);
+  }
+
+  /* %-decode the string.  */
+  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;
+  size_t pos, last = path_end;
+  while ((pos = m_uri.find('&', last)) != std::string::npos) {
+    tokens.emplace_back(m_uri.substr(last, pos - last));
+    last = pos + 1;
+  }
+  if (last != std::string::npos) {
+    tokens.emplace_back(m_uri.substr(last));
+  }
+
+  /* Create a tag-value map from the tokenized query/fragment.  */
+  std::unordered_map<std::string, std::string> params;
+  std::for_each(tokens.begin(), tokens.end(), [&](std::string& token) {
+    size_t delim = token.find('=');
+    if (delim != std::string::npos) {
+      params.emplace(token.substr(0, delim), token.substr(delim + 1));
+    }
+  });
+
+  buffer = std::vector<char>{};
+  try {
+    size_t offset{0}, size{0};
+
+    if (auto offset_it = params.find("offset"); offset_it != params.end()) {
+      offset = std::stoul(offset_it->second, nullptr, 0);
+    }
+
+    if (auto size_it = params.find("size"); size_it != params.end()) {
+      if (!(size = std::stoul(size_it->second, nullptr, 0))) return;
+    }
+
+    if (protocol != "file") {
+      printf("\"%s\" protocol not supported\n", protocol.c_str());
+      return;
+    }
+
+    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 (!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;
+      }
+      size = bytes - offset;
+    }
+
+    file.seekg(offset, std::ios_base::beg);
+    buffer.resize(size);
+    file.read(&buffer[0], size);
+  } catch (...) {
+  }
+}
+
+DisassemblyInstance::DisassemblyInstance(code_object_decoder_t& decoder)
+    : buffer(reinterpret_cast<int64_t>(decoder.buffer.data())),
+      size(decoder.buffer.size()),
+      instructions(decoder.instructions) {
+  amd_comgr_create_data(AMD_COMGR_DATA_KIND_RELOCATABLE, &data);
+  amd_comgr_set_data(data, size, decoder.buffer.data());
+
+  char isa_name[128];
+  size_t isa_size = sizeof(isa_name);
+  amd_comgr_get_data_isa_name(data, &isa_size, isa_name);
+
+  CHECK_COMGR(amd_comgr_create_disassembly_info(
+      isa_name,  //"amdgcn-amd-amdhsa--gfx1100",
+      &DisassemblyInstance::memory_callback, &DisassemblyInstance::inst_callback,
+      [](uint64_t address, void* user_data) {}, &info));
+}
+
+amd_comgr_status_t DisassemblyInstance::symbol_callback(amd_comgr_symbol_t symbol,
+                                                        void* user_data) {
+  amd_comgr_symbol_type_t type;
+  amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_TYPE, &type);
+
+  if (type != AMD_COMGR_SYMBOL_TYPE_FUNC && type != AMD_COMGR_SYMBOL_TYPE_AMDGPU_HSA_KERNEL)
+    return AMD_COMGR_STATUS_SUCCESS;
+
+  uint64_t addr;
+  amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_VALUE, &addr);
+
+  uint64_t mem_size;
+  amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_SIZE, &mem_size);
+
+  uint64_t name_size;
+  amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_NAME_LENGTH, &name_size);
+
+  std::string name;
+  name.resize(name_size);
+
+  amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_NAME, name.data());
+
+  static_cast<DisassemblyInstance*>(user_data)->symbol_map[addr] = {name, mem_size};
+  return AMD_COMGR_STATUS_SUCCESS;
+}
+
+std::map<uint64_t, std::pair<std::string, uint64_t>>& DisassemblyInstance::GetKernelMap() {
+  symbol_map = std::map<uint64_t, std::pair<std::string, uint64_t>>{};
+  amd_comgr_iterate_symbols(data, &DisassemblyInstance::symbol_callback, this);
+  return symbol_map;
+}
+
+DisassemblyInstance::~DisassemblyInstance() {
+  amd_comgr_release_data(data);
+  CHECK_COMGR(amd_comgr_destroy_disassembly_info(info));
+}
+
+uint64_t DisassemblyInstance::ReadInstruction(uint64_t addr, const char* cpp_line) {
+  uint64_t size_read;
+  amd_comgr_disassemble_instruction(info, buffer + addr, (void*)this, &size_read);
+  instructions.back().address = addr;
+  instructions.back().cpp_reference = cpp_line;
+  return size_read;
+}
+
+uint64_t DisassemblyInstance::memory_callback(uint64_t from, char* to, uint64_t size,
+                                              void* user_data) {
+  DisassemblyInstance& instance = *static_cast<DisassemblyInstance*>(user_data);
+  size_t copysize = std::min((int64_t)size, instance.buffer + instance.size - (int64_t)from);
+  std::memcpy(to, (char*)from, copysize);
+  return copysize;
+}
+
+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});
+}

plugin/att/disassembly.hpp

@@ -0,0 +1,59 @@
+/* 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
+
+#include <string>
+#include <vector>
+#include <amd_comgr/amd_comgr.h>
+#include <memory>
+
+typedef struct {
+  const char* instruction;
+  const char* cpp_reference;
+  uint64_t address;
+} instruction_instance_t;
+
+class CodeObjectBinary {
+ public:
+  CodeObjectBinary(const std::string& uri);
+  std::string m_uri;
+  std::vector<char> buffer;
+};
+
+class DisassemblyInstance {
+ public:
+  DisassemblyInstance(class code_object_decoder_t& decoder);
+  ~DisassemblyInstance();
+
+  uint64_t ReadInstruction(uint64_t addr, const char* cpp_line);
+  std::map<uint64_t, std::pair<std::string, uint64_t>>& GetKernelMap();
+
+  static uint64_t memory_callback(uint64_t from, char* to, uint64_t size, void* user_data);
+  static void inst_callback(const char* instruction, void* user_data);
+  static amd_comgr_status_t symbol_callback(amd_comgr_symbol_t symbol, void* user_data);
+
+  int64_t buffer;
+  int64_t size;
+  std::vector<instruction_instance_t>& instructions;
+  amd_comgr_disassembly_info_t info;
+  amd_comgr_data_t data;
+  std::map<uint64_t, std::pair<std::string, uint64_t>> symbol_map;
+};

plugin/att/drawing.py

@@ -120,6 +120,7 @@ def draw_wave_states(selections, normalize, TIMELINES):
     plt.figure(figsize=(15,4))
 
     maxtime = max([np.max((TIMELINES[k]!=0)*np.arange(0,TIMELINES[k].size)) for k in plot_indices])
+    maxtime = max(maxtime, 1)
     timelines = [deepcopy(TIMELINES[k][:maxtime]) for k in plot_indices]
     timelines = [np.pad(t, [0, maxtime-t.size]) for t in timelines]
 
@@ -161,6 +162,7 @@ def draw_occupancy(selections, normalize, OCCUPANCY, shadernames):
         OCCUPANCY = [occ for occ in OCCUPANCY if len(occ) > 0]
 
     maxtime = 1
+    delta = 1
     for name, occ in zip(shadernames, OCCUPANCY):
         occ_values = [0]
         occ_times = [0]

plugin/att/stitch.py

@@ -6,7 +6,7 @@ if sys.version_info[0] < 3:
 from collections import defaultdict
 from copy import deepcopy
 
-MAX_STITCHED_TOKENS = 10000000
+MAX_STITCHED_TOKENS = 100000000
 MAX_FAILED_STITCHES = 256
 STACK_SIZE_LIMIT = 64
 
@@ -25,6 +25,7 @@ GETPC = 11
 SETPC = 12
 SWAPPC = 13
 LANEIO = 14
+PCINFO = 15
 DONT_KNOW = 100
 
 WaveInstCategory = {
@@ -46,10 +47,11 @@ WaveInstCategory = {
     SETPC: "SETPC",
     SWAPPC: "SWAPPC",
     LANEIO: "LANEIO",
+    PCINFO: "PCINFO",
     DONT_KNOW: "DONT_KNOW",
 }
 
-
+# Keeps track of register states for hipcc-generated assembly
 class RegisterWatchList:
     def __init__(self, labels):
         self.registers = {'v'+str(k): [[] for m in range(64)] for k in range(64)}
@@ -84,7 +86,7 @@ class RegisterWatchList:
         except:
             pass
 
-    def swappc(self, line, line_num):
+    def swappc(self, line, line_num, inst_num):
         try:
             tokens = self.tokenize(line)
             dst = tokens[1]
@@ -97,10 +99,9 @@ class RegisterWatchList:
         except:
             return 0
 
-    def setpc(self, line):
+    def setpc(self, line, inst_num):
         try:
             src = line.split(' ')[1].strip()
-            #print('Going to:', self.registers[self.range(src)[0]], src)
             popped = self.registers[self.range(src)[0]][-1]
             self.registers[self.range(src)[0]] = self.registers[self.range(src)[0]][:-1]
             return popped
@@ -140,16 +141,57 @@ class RegisterWatchList:
         except Exception as e:
             pass
 
+# Translates PC values to instructions, for auto captured ISA
+class PCTranslator:
+    def __init__(self, code, insts):
+        self.code = code
+        self.insts = insts
+        self.addrmap = {code[m][-3] : m for m in range(len(code))}
+
+    def try_translate(self, tok):
+        pass
+    def range(self, r):
+        pass
+    def tokenize(self, line):
+        pass
+    def getpc(self, line, next_line):
+        pass
+    def swappc(self, line, line_num, inst_index):
+        try:
+            loc = self.addrmap[self.insts[inst_index+1][2]]
+            #print('Jumping to:', loc, self.code[loc])
+            return loc
+        except:
+            print('SWAPPC: Could not find addr', self.insts[inst_index+1][2], 'for', line)
+            return -1
+    def setpc(self, line, inst_index):
+        try:
+            loc = self.addrmap[self.insts[inst_index+1][2]]
+            #print('Jumping to:', loc, self.code[loc])
+            return loc
+        except:
+            print('SETPC: Could not find addr', self.insts[inst_index+1][2], 'for', line)
+            return -1
+    def scratch(self, line):
+        pass
+    def move(self, line):
+        pass
+    def updatelane(self, line):
+        pass
+
+# Matches tokens in reverse order
 def try_match_swapped(insts, code, i, line):
     return insts[i+1][1] == code[line][1] and insts[i][1] == code[line+1][1]
 
 FORK_NAMES = 1
+# A successful parsed instruction
 class CachedInst:
     def __init__(self, inst, as_line):
         self.inst_type = inst
         self.as_line = as_line
         self.forks = None
 
+# A branch of the parsing tree
 class Fork:
     def __init__(self):
         global FORK_NAMES
@@ -159,19 +201,21 @@ class Fork:
         FORK_NAMES += 1
         #print('Created new fork: ', self.name)
 
-def move_down_fork(fork, insts, i): #def move_down_fork(fork : Fork, insts : list, i : int):
+# Try to match sequence "insts" with the branch "fork", starting at position "i"
+def move_down_fork(fork, insts, i): #(fork : Fork, insts : list, i : int):
     N = min(len(insts), len(fork.insts))
 
     while i < N:
         if insts[i][1] == fork.insts[i].inst_type:
             i += 1
-        elif i<N-1 and insts[i+1][1] == fork.insts[i].inst_type and insts[i][1] == fork.insts[i+1].inst_type:
+        elif i<N-1  and insts[i+1][1] == fork.insts[i].inst_type \
+                    and insts[i][1] == fork.insts[i+1].inst_type:
             i += 2
         else:
             #print('Failed at', i, insts[i])
             return False, i
 
-    if len(fork.insts) < len(insts):
+    if len(fork.insts) != len(insts):
         #print('Failed at the end at', i, insts[i])
         return False, i
 
@@ -180,11 +224,11 @@ def move_down_fork(fork, insts, i): #def move_down_fork(fork : Fork, insts : lis
 
 FORK_TREE = Fork()
 
+# Check if there exists a previous wave with the same sequence of instructions executed
 def fromDict(insts):
     i = 0
     N = len(insts)
     cur_fork = FORK_TREE
-    #print('Getting from dict')
     while i < N:
         tillEnd, final_pos = move_down_fork(cur_fork, insts, i)
         if tillEnd:
@@ -192,7 +236,6 @@ def fromDict(insts):
             return True, cur_fork
 
         i += final_pos
-        #print('Got fpos:', i, 'of', len(insts))
 
         if i >= len(cur_fork.insts):
             return False, cur_fork
@@ -204,7 +247,6 @@ def fromDict(insts):
         bMatchFork = False
         for fork in last_inst.forks:
             if fork.insts[0].inst_type == insts[0][1]:
-                #print('Found match fork', fork.name)
                 cur_fork = fork
                 bMatchFork = True
                 break
@@ -217,10 +259,30 @@ def fromDict(insts):
     return False, cur_fork
 
 
-def stitch(insts, raw_code, jumps, gfxv):
+def stitch(insts, raw_code, jumps, gfxv, bIsAuto):
     bGFX9 = gfxv == 'vega'
+
+    # Try from cached result from a previous wave that have already been parsed
+    dict_sucess, current_fork = fromDict(insts)
+    if dict_sucess:
+        result, loopCount, mem_unroll, flight_count, maxline, pcsequence = current_fork.data
+        # Check if the sequence of measured PC values are equal for cached and new wave
+        if len(pcsequence) > 0:
+            pcs = [r[2] for r in insts if r[1] == PCINFO]
+            if len(pcs) != len(pcsequence):
+                dict_sucess = False
+            for pc1, pc2 in zip(pcs, pcsequence):
+                if pc1 != pc2:
+                    dict_sucess = False
+
+    # If successful, use resulting assembly from cache
+    if dict_sucess:
+        result = [r+(asm[-1],) for r, asm in zip(insts, result)]
+        return result, loopCount, mem_unroll, flight_count, maxline, len(result)
+
     result, i, line, loopCount, N = [], 0, 0, defaultdict(int), len(insts)
 
+
     SMEM_INST = []  # scalar memory
     VLMEM_INST = [] # vector memory load
     VSMEM_INST = [] # vector memory store
@@ -236,8 +298,14 @@ def stitch(insts, raw_code, jumps, gfxv):
 
     labels = {}
     jump_map = [0]
+
+    # Clean the code and remove comments
     code = [raw_code[0]]
     for c in raw_code[1:]:
+        if bIsAuto and '; Begin ' == c[0][:len('; Begin ')]:
+            if '; Begin <Kernel>' in c[0]:
+                line = len(code)
+                print('Begin at:', line, c)
         c = list(c)
         c[0] = c[0].split(';')[0].split('//')[0].strip()
 
@@ -254,25 +322,22 @@ def stitch(insts, raw_code, jumps, gfxv):
 
     jumps = {jump_map[j]+1: j for j in jumps}
 
+    # Checks if we have guaranteed ordering in memory operations
     smem_ordering = 0
     vlmem_ordering = 0
     vsmem_ordering = 0
 
-    watchlist = RegisterWatchList(labels=labels)
-
     num_failed_stitches = 0
     loops = 0
     maxline = 0
 
-    dict_sucess, current_fork = fromDict(insts)
-    if dict_sucess:
-        result, loopCount, mem_unroll, flight_count, maxline = current_fork.data
-        result = [r+(asm[-1],) for r, asm in zip(insts, result)]
-        return result, loopCount, mem_unroll, flight_count, maxline, len(insts)
+    watchlist = RegisterWatchList(labels=labels) if not bIsAuto else PCTranslator(code, insts)
 
+    pcsequence = []
     while i < N:
         loops += 1
-        if line >= len(code) or loops > MAX_STITCHED_TOKENS or num_failed_stitches > MAX_FAILED_STITCHES:
+        if line >= len(code) or loops > MAX_STITCHED_TOKENS \
+            or num_failed_stitches > MAX_FAILED_STITCHES:
             break
 
         maxline = max(reverse_map[line], maxline)
@@ -282,23 +347,37 @@ def stitch(insts, raw_code, jumps, gfxv):
         matched = True
         next = line+1
 
-        if '_mov_' in as_line[0]:
-            watchlist.move(as_line[0])
-        elif 'scratch_' in as_line[0]:
-            watchlist.scratch(as_line[0])
+        if not bIsAuto:
+            if '_mov_' in as_line[0]:
+                watchlist.move(as_line[0])
+            elif 'scratch_' in as_line[0]:
+                watchlist.scratch(as_line[0])
 
         if as_line[1] == GETPC:
-            watchlist.getpc(as_line[0], code[line+1][0])
-            matched = inst[1] in [SALU, JUMP]
+            try:
+                watchlist.getpc(as_line[0], code[line+1][0])
+                matched = inst[1] in [SALU, JUMP]
+            except:
+                matched = False
         elif as_line[1] == LANEIO:
             watchlist.updatelane(as_line[0])
             matched = inst[1] == VALU
         elif as_line[1] == SETPC:
-            next = watchlist.setpc(as_line[0])
+            next = watchlist.setpc(as_line[0], i)
             matched = inst[1] in [SALU, JUMP]
+            if bIsAuto:
+                matched = next >= 0
+                i += 1
+                result.append((insts[i][0], PCINFO, 0, 0, 0))
+                pcsequence.append(insts[i][2])
         elif as_line[1] == SWAPPC:
-            next = watchlist.swappc(as_line[0], line)
+            next = watchlist.swappc(as_line[0], line, i)
             matched = inst[1] in [SALU, JUMP]
+            if bIsAuto:
+                matched = next >= 0
+                i += 1
+                result.append((insts[i][0], PCINFO, 0, 0, 0))
+                pcsequence.append(insts[i][2])
         elif inst[1] == as_line[1]:
             if line in jumps:
                 loopCount[jumps[line]-1] += 1
@@ -366,10 +445,10 @@ def stitch(insts, raw_code, jumps, gfxv):
 
                 if 'vscnt' in as_line[0] or (bGFX9 and 'vmcnt' in as_line[0]):
                     try:
-                        wait_N = int(as_line[0].split('vmcnt(')[1].split(')')[0])
+                        wait_N = int(as_line[0].split('vscnt(')[1].split(')')[0])
                     except:
                         try:
-                            wait_N = int(as_line[0].split('vscnt(')[1].split(')')[0])
+                            wait_N = int(as_line[0].split('vmcnt(')[1].split(')')[0])
                         except:
                             wait_N = 0
                     flight_count.append([as_line[5], num_inflight, wait_N])
@@ -407,10 +486,13 @@ def stitch(insts, raw_code, jumps, gfxv):
                     if skipped_immed > 0 and 's_waitcnt ' in as_line[0]:
                         matched = True
                         skipped_immed -= 1
-                    else:
+                    elif 'scratch_' not in as_line[0]:
                         print('Parsing terminated at:', as_line)
                         break
 
+        #print(matched, as_line)
+        #print([WaveInstCategory[insts[i+k][1]] for k in range(10) if i+k < len(insts)])
+
         if matched:
             result.append(inst + (reverse_map[line],))
             i += 1
@@ -425,8 +507,8 @@ def stitch(insts, raw_code, jumps, gfxv):
         line = next
 
     N = max(N, 1)
-    if len(result) != N:
-        print('Warning - Stitching rate: '+str(len(result) * 100 / N)+'% matched')
+    if i != N:
+        print('Warning - Stitching rate: '+str(i * 100 / N)+'% matched')
         print('Leftovers:', [WaveInstCategory[insts[i+k][1]] for k in range(20) if i+k < len(insts)])
         try:
             print(line, code[line])
@@ -440,5 +522,6 @@ def stitch(insts, raw_code, jumps, gfxv):
             line += 1
 
     current_fork.insts = [CachedInst(inst[1], inst[-1]) for inst in result]
-    current_fork.data = result, loopCount, mem_unroll, flight_count, maxline
-    return result, loopCount, mem_unroll, flight_count, maxline, len(insts)
+    current_fork.data = result, loopCount, mem_unroll, flight_count, maxline, pcsequence
+    result = [r for r in result if r[1] != PCINFO]
+    return result, loopCount, mem_unroll, flight_count, maxline, len(result) if i == N else N

plugin/att/trace_view.py

@@ -278,7 +278,8 @@ def view_trace(args, code, dbnames, att_filenames, bReturnLoc, OCCUPANCY, bDumpO
             simd_wave_filenames[se_number] = wv_filenames
 
     if mpi_root:
-        JSON_GLOBAL_DICTIONARY['code.json'] = Readable({"code": code[:allse_maxline+16], "top_n": get_top_n(code[:allse_maxline+16])})
+        code_sel = [c[:-3]+c[-2:] for c in code[:allse_maxline+16]]
+        JSON_GLOBAL_DICTIONARY['code.json'] = Readable({"code": code_sel, "top_n": get_top_n(code_sel)})
 
     for key in simd_wave_filenames.keys():
         wv_array = [[

src/api/CMakeLists.txt

@@ -184,11 +184,15 @@ 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)
-file(GLOB ROCPROFILER_ATT_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/att/att.cpp)
+set(ROCPROFILER_ATT_SRC_FILES
+    ${PROJECT_SOURCE_DIR}/src/core/session/att/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)
 
+set(CORE_ISA_CAPTURE_DIR ${PROJECT_SOURCE_DIR}/src/core/isa_capture)
+file(GLOB CORE_ISA_CAPTURE_SRC_FILES ${CORE_ISA_CAPTURE_DIR}/*.cpp)
+
 set(CORE_HARDWARE_DIR ${PROJECT_SOURCE_DIR}/src/core/hardware)
 file(GLOB CORE_HARDWARE_SRC_FILES ${CORE_HARDWARE_DIR}/*.cpp)
 
@@ -265,6 +269,7 @@ add_library(
     ${ROCPROFILER_CLASS_SRC_FILES}
     ${ROCPROFILER_PROFILER_SRC_FILES}
     ${ROCPROFILER_ATT_SRC_FILES}
+    ${CORE_ISA_CAPTURE_SRC_FILES}
     ${CORE_HARDWARE_SRC_FILES}
     ${CORE_HSA_SRC_FILES}
     ${ROCPROFILER_SPM_SRC_FILES}
@@ -326,6 +331,8 @@ if(ASAN)
                 stdc++
                 stdc++fs
                 amd_comgr
+                dw
+                elf
                 ${PCIACCESS_LIBRARIES})
 else()
     target_link_options(
@@ -343,6 +350,8 @@ else()
                 stdc++
                 stdc++fs
                 amd_comgr
+                dw
+                elf
                 ${PCIACCESS_LIBRARIES})
 endif()
 

src/api/exportmap

@@ -38,5 +38,10 @@ global: OnLoad;
         rocprofiler_device_profiling_session_poll;
         rocprofiler_device_profiling_session_stop;
         rocprofiler_device_profiling_session_destroy;
+        rocprofiler_codeobj_capture_get;
+        rocprofiler_codeobj_capture_create;
+        rocprofiler_codeobj_capture_free;
+        rocprofiler_codeobj_capture_start;
+        rocprofiler_codeobj_capture_stop;
 local:  *;
 };

src/api/rocprofilerv2.cpp

@@ -8,6 +8,7 @@
 #include "src/core/hsa/hsa_support.h"
 #include "src/api/rocprofiler_singleton.h"
 #include "src/utils/helper.h"
+#include "src/core/isa_capture/code_object_track.hpp"
 
 // TODO(aelwazir): change that to adapt with our own Exception
 // What about outside exceptions and callbacks exceptions!!
@@ -543,6 +544,59 @@ rocprofiler_device_profiling_session_destroy(rocprofiler_session_id_t session_id
   API_METHOD_SUFFIX
 }
 
+ROCPROFILER_API rocprofiler_status_t
+rocprofiler_codeobj_capture_get(rocprofiler_record_id_t id,
+                                rocprofiler_codeobj_symbols_t* symbols) {
+  API_METHOD_PREFIX
+  try {
+    *symbols = codeobj_record::get_capture(id);
+  } catch(const std::out_of_range& e) {
+    return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS;
+  }
+  API_METHOD_SUFFIX
+}
+
+ROCPROFILER_API rocprofiler_status_t
+rocprofiler_codeobj_capture_create(
+  rocprofiler_record_id_t* id,
+  rocprofiler_codeobj_capture_mode_t mode,
+  uint64_t userdata
+) {
+  API_METHOD_PREFIX
+  id->handle = rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId();
+  codeobj_record::make_capture(*id, mode, userdata);
+  API_METHOD_SUFFIX
+}
+
+ROCPROFILER_API rocprofiler_status_t
+rocprofiler_codeobj_capture_free(rocprofiler_record_id_t id) {
+  API_METHOD_PREFIX
+  codeobj_record::free_capture(id);
+  API_METHOD_SUFFIX
+}
+
+ROCPROFILER_API rocprofiler_status_t
+rocprofiler_codeobj_capture_start(rocprofiler_record_id_t id) {
+  API_METHOD_PREFIX
+  try {
+    codeobj_record::start_capture(id);
+  } catch(const std::out_of_range& e) {
+    return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS;
+  }
+  API_METHOD_SUFFIX
+}
+
+ROCPROFILER_API rocprofiler_status_t
+rocprofiler_codeobj_capture_stop(rocprofiler_record_id_t id) {
+  API_METHOD_PREFIX
+  try {
+    codeobj_record::stop_capture(id);
+  } catch(const std::out_of_range& e) {
+    return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS;
+  }
+  API_METHOD_SUFFIX
+}
+
 
 static bool started{false};
 

src/core/hsa/hsa_support.cpp

@@ -50,6 +50,7 @@
 
 #include "src/core/hsa/queues/queue.h"
 #include "src/api/rocprofiler_singleton.h"
+#include "src/core/isa_capture/code_object_track.hpp"
 
 #include <experimental/filesystem>
 namespace fs = std::experimental::filesystem;
@@ -464,6 +465,16 @@ hsa_status_t CodeObjectCallback(hsa_executable_t executable,
   data.codeobj.unload = *static_cast<bool*>(arg) ? 1 : 0;
   ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ, &data);
 
+  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
+    );
+
   hsa_executable_iterate_agent_symbols(executable, data.codeobj.agent,
                                        hsa_executable_iteration_callback, &(data.codeobj.unload));
 

src/core/hsa/queues/queue.cpp

@@ -33,6 +33,7 @@
 #include "src/core/hsa/packets/packets_generator.h"
 #include "src/core/hsa/hsa_support.h"
 #include "src/utils/helper.h"
+#include "src/core/isa_capture/code_object_track.hpp"
 
 #define CHECK_HSA_STATUS(msg, status)                                                              \
   do {                                                                                             \
@@ -493,17 +494,20 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
 }
 
 bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
-  // TODO: finish implementation to iterate trace data and add it to rocprofiler record
-  // and generic buffer
 
   auto queue_info_session = static_cast<queue_info_session_t*>(data);
-  if (!queue_info_session || !GetROCProfilerSingleton() ||
-      !GetROCProfilerSingleton()->GetSession(queue_info_session->session_id) ||
-      !GetROCProfilerSingleton()->GetSession(queue_info_session->session_id)->GetAttTracer())
+  if (!queue_info_session || !GetROCProfilerSingleton())
     return true;
+
   rocprofiler::Session* session =
       GetROCProfilerSingleton()->GetSession(queue_info_session->session_id);
+  if (!session) return true;
+
+  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));
@@ -512,9 +516,9 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
     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 (hsa_support::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};
@@ -522,13 +526,18 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
       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 writer ID so we can correlate to dispatches
-      // kernel_id already has the descriptor.
+      // 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{queue_info_session->writer_id}};
+                       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 = GetROCProfilerSingleton()->GetCurrentSessionId();
@@ -536,7 +545,10 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
       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 = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
           (pending.profile->output_buffer.ptr));
       CHECK_HSA_STATUS("Error: Couldn't free output buffer memory", status);
@@ -548,6 +560,7 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
   }
   delete queue_info_session;
 
+  std::atomic_thread_fence(std::memory_order_seq_cst);
   return false;
 }
 
@@ -721,17 +734,25 @@ std::pair<std::vector<bool>, bool> GetAllowedProfilesList(const void* packets, i
   return {can_profile_packet, b_can_profile_anypacket};
 }
 
-hsa_ven_amd_aqlprofile_profile_t* ProcessATTParams(Packet::packet_t& start_packet,
-                                                   Packet::packet_t& stop_packet, Queue& queue_info,
-                                                   Agent::AgentInfo& agentInfo) {
+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,
+  Agent::AgentInfo& 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)));
@@ -773,8 +794,8 @@ hsa_ven_amd_aqlprofile_profile_t* ProcessATTParams(Packet::packet_t& start_packe
     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);
+  return {Packet::GenerateATTPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(),
+          att_params, &start_packet, &stop_packet, att_buffer_size), capture_mode};
 }
 
 /**
@@ -786,6 +807,7 @@ hsa_ven_amd_aqlprofile_profile_t* ProcessATTParams(Packet::packet_t& start_packe
  */
 void 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;
 
@@ -952,9 +974,15 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
     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 && is_att_collection_mode)
-      profile = ProcessATTParams(start_packet, stop_packet, queue_info, agentInfo);
+    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) {
@@ -976,7 +1004,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
       KernelInterceptCount += 1;
       writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
 
-      if (att_parameters_data.size() > 0 && is_att_collection_mode && profile) {
+      if (att_parameters_data.size() > 0 && profile) {
         // Adding start packet and its barrier with a dummy signal
         hsa_signal_t dummy_signal{};
         dummy_signal.handle = 0;
@@ -996,17 +1024,17 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
       uint64_t record_id = GetROCProfilerSingleton()->GetUniqueRecordId();
       AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object),
                                   record_id);
-      if (session && profile) {
-        session->GetAttTracer()->AddPendingSignals(
-            writer_id, record_id, original_packet.completion_signal,
-            dispatch_packet.completion_signal, session_id_snapshot, buffer_id, profile,
-            kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index);
-      } else {
-        session->GetAttTracer()->AddPendingSignals(
-            writer_id, record_id, original_packet.completion_signal,
-            dispatch_packet.completion_signal, session_id_snapshot, buffer_id, nullptr,
-            kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index);
-      }
+
+      session->GetAttTracer()->AddPendingSignals(
+          writer_id, record_id, original_packet.completion_signal,
+          dispatch_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);
+      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) {

src/core/isa_capture/code_object_track.cpp

@@ -0,0 +1,248 @@
+/* Copyright (c) 2023 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 <algorithm>
+#include <atomic>
+#include <functional>
+#include <map>
+#include <mutex>
+#include <optional>
+#include <string>
+#include <type_traits>
+#include <unordered_map>
+#include <vector>
+#include <cassert>
+#include <cinttypes>
+#include <cstdint>
+#include <cstdio>
+#include <cstring>
+#include <sys/mman.h>
+#include <hsa/hsa.h>
+#include <amd-dbgapi/amd-dbgapi.h>
+#include <hsa/amd_hsa_kernel_code.h>
+#include <hsa/hsa_ven_amd_loader.h>
+#include <iostream>
+#include "src/utils/helper.h"
+#include "src/api/rocprofiler_singleton.h"
+#include "src/core/isa_capture/code_object_track.hpp"
+#include <amd_comgr/amd_comgr.h>
+
+std::mutex codeobj_capture_instance::mutex;
+std::mutex codeobj_record::mutex;
+
+std::unordered_map<uint64_t, CodeobjPtr> codeobj_capture_instance::codeobjs{};
+std::unordered_map<uint64_t, codeobj_record::RecordInstance> codeobj_record::record_id_map{};
+std::unordered_set<codeobj_record*> codeobj_record::listeners;
+
+// Codeobj Record
+codeobj_record::codeobj_record(rocprofiler_codeobj_capture_mode_t mode) : capture_mode(mode){};
+
+void codeobj_record::start_capture() {
+  listeners.insert(this);
+  for (auto& [addr, capture] : codeobj_capture_instance::codeobjs) this->addcapture(capture);
+}
+
+void codeobj_record::addcapture(CodeobjPtr& capture) {
+  if (captures.find(capture) != captures.end()) return;
+  capture->setmode(capture_mode);
+  captures.insert(capture);
+}
+
+void codeobj_record::stop_capture() {
+  try {
+    listeners.erase(this);
+  } catch (...) {
+  };
+}
+
+// Codeobj Capture
+void codeobj_capture_instance::Load(uint64_t addr, const std::string& URI, uint64_t mem_addr,
+                                    uint64_t mem_size) {
+  std::lock_guard<std::mutex> lock(mutex);
+  codeobjs[addr] = std::make_shared<codeobj_capture_instance>(
+      addr, URI, mem_addr, mem_size, rocprofiler::GetCurrentTimestamp().value);
+  std::atomic_thread_fence(std::memory_order_release);  // Fencing the state of the map
+  {
+    std::lock_guard<std::mutex> lock(codeobj_record::mutex);
+    for (auto* listen : codeobj_record::listeners) listen->addcapture(codeobjs.at(addr));
+  }
+}
+
+void codeobj_capture_instance::Unload(uint64_t addr) {
+  std::lock_guard<std::mutex> lock(mutex);
+  codeobjs.at(addr)->end_time = rocprofiler::GetCurrentTimestamp().value;
+  codeobjs.erase(addr);
+}
+
+void codeobj_capture_instance::copyCodeobjFromFile(uint64_t offset, uint64_t size,
+                                                   const std::string& decoded_path) {
+  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 (!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;
+    }
+    size = bytes - offset;
+  }
+
+  file.seekg(offset, std::ios_base::beg);
+  buffer.resize(size);
+  file.read(&buffer[0], size);
+}
+
+void codeobj_capture_instance::copyCodeobjFromMemory(uint64_t mem_addr, uint64_t mem_size) {
+  buffer.resize(mem_size);
+  std::memcpy(buffer.data(), (uint64_t*)mem_addr, mem_size);
+}
+
+std::pair<size_t, size_t> codeobj_capture_instance::parse_uri() {
+  const std::string protocol_delim{"://"};
+
+  size_t protocol_end = URI.find(protocol_delim);
+  protocol = URI.substr(0, protocol_end);
+  protocol_end += protocol_delim.length();
+
+  std::transform(protocol.begin(), protocol.end(), protocol.begin(),
+                 [](unsigned char c) { return std::tolower(c); });
+
+  std::string path;
+  size_t path_end = URI.find_first_of("#?", protocol_end);
+  if (path_end != std::string::npos) {
+    path = URI.substr(protocol_end, path_end++ - protocol_end);
+  } else {
+    path = URI.substr(protocol_end);
+  }
+
+  /* %-decode the string.  */
+  decoded_path = std::string{};
+  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;
+  size_t pos, last = path_end;
+  while ((pos = URI.find('&', last)) != std::string::npos) {
+    tokens.emplace_back(URI.substr(last, pos - last));
+    last = pos + 1;
+  }
+  if (last != std::string::npos) tokens.emplace_back(URI.substr(last));
+
+  /* Create a tag-value map from the tokenized query/fragment.  */
+  std::unordered_map<std::string, std::string> params;
+  std::for_each(tokens.begin(), tokens.end(), [&](std::string& token) {
+    size_t delim = token.find('=');
+    if (delim != std::string::npos) {
+      params.emplace(token.substr(0, delim), token.substr(delim + 1));
+    }
+  });
+
+  size_t offset{0}, size{0};
+
+  if (auto offset_it = params.find("offset"); offset_it != params.end())
+    offset = std::stoul(offset_it->second, nullptr, 0);
+
+  if (auto size_it = params.find("size"); size_it != params.end()) {
+    if (!(size = std::stoul(size_it->second, nullptr, 0))) throw std::exception();
+  }
+
+  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) {
+  reset(ROCPROFILER_CAPTURE_SYMBOLS_ONLY);
+};
+
+void codeobj_capture_instance::setmode(rocprofiler_codeobj_capture_mode_t mode) {
+  // Only reset when needed & check if codeobj was not unloaded
+  if (end_time == 0 && static_cast<int>(mode) > static_cast<int>(capture_mode))
+    reset(mode);
+}
+
+void codeobj_capture_instance::reset(rocprofiler_codeobj_capture_mode_t mode) {
+  capture_mode = mode;
+
+  size_t offset, size;
+  try {
+    std::tie(offset, size) = parse_uri();
+  } catch (...) {
+    return;
+  }
+
+  buffer = std::vector<char>{};
+
+  if (mode == ROCPROFILER_CAPTURE_SYMBOLS_ONLY) return;
+
+  if (protocol == "file") {
+    if (mode == ROCPROFILER_CAPTURE_COPY_FILE_AND_MEMORY)
+      copyCodeobjFromFile(offset, size, decoded_path);
+  } else if (protocol == "memory") {
+    copyCodeobjFromMemory(mem_addr, mem_size);
+  } else {
+    printf("\"%s\" protocol not supported\n", protocol.c_str());
+    return;
+  }
+}
+
+// Public static funcs
+void codeobj_record::make_capture(rocprofiler_record_id_t id,
+                                  rocprofiler_codeobj_capture_mode_t mode, uint64_t userdata) {
+  std::lock_guard<std::mutex> lock(mutex);
+  record_id_map[id.handle] = {userdata, std::unique_ptr<codeobj_record>{new codeobj_record(mode)}};
+}
+
+void codeobj_record::free_capture(rocprofiler_record_id_t id) {
+  std::lock_guard<std::mutex> lock(mutex);
+  record_id_map.erase(id.handle);
+}
+
+void codeobj_record::start_capture(rocprofiler_record_id_t id) {
+  std::lock_guard<std::mutex> lock(mutex);
+  record_id_map.at(id.handle).second->start_capture();
+}
+
+void codeobj_record::stop_capture(rocprofiler_record_id_t id) {
+  std::lock_guard<std::mutex> lock(mutex);
+  record_id_map.at(id.handle).second->stop_capture();
+}
+
+rocprofiler_codeobj_symbols_t codeobj_record::get_capture(rocprofiler_record_id_t id) {
+  std::atomic_thread_fence(std::memory_order_acquire);  // Fencing the state of the map
+  auto& pair = record_id_map.at(id.handle);
+  return pair.second->get(pair.first);
+}

src/core/isa_capture/code_object_track.hpp

@@ -0,0 +1,133 @@
+/* Copyright (c) 2023 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
+
+#include "src/utils/helper.h"
+#include <unordered_map>
+#include <unordered_set>
+#include <memory>
+#include <mutex>
+#include <string>
+#include <fstream>
+
+/**
+ * A class to keep track of currently loaded code objects.
+ * Only the public static methods are thread-safe and expected to be used.
+ */
+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);
+
+  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, buf_ptr, buffer.size(), start_time, end_time};
+  };
+
+  const uint64_t addr;
+  const uint64_t start_time;
+
+  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 std::unordered_map<uint64_t, std::shared_ptr<codeobj_capture_instance>> codeobjs;
+
+ private:
+  void reset(rocprofiler_codeobj_capture_mode_t mode);
+
+  std::pair<size_t, size_t> parse_uri();
+  void DecodePath();
+  void copyCodeobjFromFile(uint64_t offset, uint64_t size, const std::string& decoded_path);
+  void copyCodeobjFromMemory(uint64_t mem_addr, uint64_t mem_size);
+
+  std::string URI;
+  std::string decoded_path;
+  std::string protocol;
+  std::vector<char> buffer;
+
+  uint64_t mem_addr;
+  uint64_t mem_size;
+  uint64_t end_time = 0;
+  rocprofiler_codeobj_capture_mode_t capture_mode;
+
+  // Address -> codeobj
+  static std::mutex mutex;
+};
+
+typedef std::shared_ptr<codeobj_capture_instance> CodeobjPtr;
+
+template <> struct std::hash<CodeobjPtr> {
+  // addr is typically 2^12-byte aligned. Taking last 44 bits of time == cycle time of many hours.
+  uint64_t operator()(const CodeobjPtr& p) const {
+    return (p->addr >> 12) ^ (p->start_time << 20);
+  };
+};
+
+template <> struct std::equal_to<CodeobjPtr> {
+  bool operator()(const CodeobjPtr& a, const CodeobjPtr& b) const {
+    return (a->addr == b->addr) & (a->start_time == b->start_time);
+  };
+};
+
+/**
+ * A class to keep track of the history of loaded code objets.
+ * Only the public static methods are thread-safe and expected to be used.
+ */
+class codeobj_record {
+ public:
+  codeobj_record(rocprofiler_codeobj_capture_mode_t mode);
+  ~codeobj_record() {
+    if (listeners.find(this) != listeners.end()) stop_capture();
+  };
+
+  void addcapture(CodeobjPtr& capture);
+
+ public:
+  static void make_capture(rocprofiler_record_id_t id, rocprofiler_codeobj_capture_mode_t mode,
+                           uint64_t userdata);
+  static void free_capture(rocprofiler_record_id_t id);
+  static void start_capture(rocprofiler_record_id_t id);
+  static void stop_capture(rocprofiler_record_id_t id);
+  static rocprofiler_codeobj_symbols_t get_capture(rocprofiler_record_id_t id);
+
+  static std::unordered_set<codeobj_record*> listeners;
+  static std::mutex mutex;
+
+ private:
+  rocprofiler_codeobj_symbols_t get(uint64_t userdata) {
+    persist.clear();
+    for (auto& capt : captures) persist.push_back(capt->get());
+    return rocprofiler_codeobj_symbols_t{persist.data(), persist.size(), userdata};
+  };
+
+  void start_capture();
+  void stop_capture();
+
+  rocprofiler_codeobj_capture_mode_t capture_mode;
+  std::vector<rocprofiler_intercepted_codeobj_t> persist;
+  std::unordered_set<CodeobjPtr> captures;
+
+  // Record_id -> codeobj
+  using RecordInstance = std::pair<uint64_t, std::unique_ptr<codeobj_record>>;
+  static std::unordered_map<uint64_t, RecordInstance> record_id_map;
+};

src/core/session/att/att.cpp

@@ -41,11 +41,9 @@ void AttTracer::AddPendingSignals(
   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});
-  std::atomic_thread_fence(std::memory_order_release);
 }
 
 const std::vector<att_pending_signal_t>& AttTracer::GetPendingSignals(uint32_t writer_id) {
-  std::atomic_thread_fence(std::memory_order_acquire);
   std::lock_guard<std::mutex> lock(sessions_pending_signals_lock_);
   assert(sessions_pending_signals_.find(writer_id) != sessions_pending_signals_.end() &&
          "writer_id is not found in the pending_signals");

src/tools/tool.cpp

@@ -287,12 +287,14 @@ att_parsed_input_t GetATTParams() {
 
   ATT_PARAM_NAMES["KERNEL"] = ROCPROFILER_ATT_MAXVALUE;
   ATT_PARAM_NAMES["BUFFER_SIZE"] = ROCPROFILER_ATT_BUFFER_SIZE;
+  ATT_PARAM_NAMES["ISA_CAPTURE_MODE"] = ROCPROFILER_ATT_CAPTURE_MODE;
 
   // Default values used for token generation.
   std::unordered_map<std::string, uint32_t> default_params = {
     {"SE_MASK", 0x111111}, // One every 4 SEs, by default
     {"SIMD_SELECT", 0x3}, // 0x3 works for both gfx9 and Navi
-    {"BUFFER_SIZE", 0x40000000} // 2^30 == 1GB
+    {"BUFFER_SIZE", 0x40000000}, // 2^30 == 1GB
+    {"ISA_CAPTURE_MODE", static_cast<uint32_t>(ROCPROFILER_CAPTURE_SYMBOLS_ONLY)}
   };
 
   std::ifstream trace_file(path);

tests-v2/featuretests/profiler/profiler_gtest.cpp

@@ -1150,6 +1150,146 @@ TEST(ProfilerMPTest, WhenRunningMultiProcessTestItPasses) {
     ASSERT_TRUE(1);
   }
 }
+
+/*
+ * ###################################################
+ * ############ Codeobj capture tests ################
+ * ###################################################
+ */
+
+class CodeobjTest : public ::testing::Test {
+public:
+  virtual void SetUp(const char* app_name) {};
+  virtual void TearDown(){};
+  static void FlushCallback(const rocprofiler_record_header_t* record,
+                            const rocprofiler_record_header_t* end_record,
+                            rocprofiler_session_id_t session_id,
+                            rocprofiler_buffer_id_t buffer_id) {};
+
+  void SetupRocprofiler() {
+    int result = ROCPROFILER_STATUS_ERROR;
+
+    result = rocprofiler_initialize();
+    EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+    result = rocprofiler_create_session(ROCPROFILER_NONE_REPLAY_MODE, &session_id);
+    EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+    result = rocprofiler_codeobj_capture_create(&id, ROCPROFILER_CAPTURE_SYMBOLS_ONLY, 0);
+    EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+  }
+
+  void TearDownRocprofiler() {
+    int result = ROCPROFILER_STATUS_ERROR;
+
+    result = rocprofiler_codeobj_capture_free(id);
+    EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+    result = rocprofiler_destroy_session(session_id);
+    EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+    result = rocprofiler_finalize();
+    EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+  }
+
+  rocprofiler_session_id_t session_id;
+  rocprofiler_record_id_t id;
+};
+
+TEST_F(CodeobjTest, WhenRunningProfilerWithCodeobjCapture) {
+  int result = ROCPROFILER_STATUS_ERROR;
+  SetupRocprofiler();
+
+  result = rocprofiler_codeobj_capture_start(id);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  // Launch a kernel
+  LaunchVectorAddKernel();
+
+  result = rocprofiler_codeobj_capture_stop(id);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  rocprofiler_codeobj_symbols_t capture;
+
+  rocprofiler_codeobj_capture_get(id, &capture);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  EXPECT_GE(capture.count, 1);
+  bool bCaptured_itself = false;
+
+  for (int i=0; i<(int)capture.count; i++) {
+    const char* path = capture.symbols[i].filepath;
+    if (!path) continue;
+    std::string fpath(path);
+    size_t pos = fpath.find("runFeatureTests#offset=");
+    if (pos != std::string::npos && fpath.find("&size=", pos) != std::string::npos)
+      bCaptured_itself = true;
+  }
+  EXPECT_EQ(bCaptured_itself, true);
+
+  TearDownRocprofiler();
+}
+
+
+TEST_F(CodeobjTest, WhenRunningProfilerWithMultipleCaptureAndCopy) {
+  int result = ROCPROFILER_STATUS_ERROR;
+
+  SetupRocprofiler();
+
+  rocprofiler_record_id_t id2;
+  rocprofiler_record_id_t id3;
+
+  result = rocprofiler_codeobj_capture_create(&id2, ROCPROFILER_CAPTURE_COPY_FILE_AND_MEMORY, 0);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  result = rocprofiler_codeobj_capture_start(id);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  result = rocprofiler_codeobj_capture_start(id2);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  result = rocprofiler_codeobj_capture_create(&id3, ROCPROFILER_CAPTURE_COPY_MEMORY, 0);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  result = rocprofiler_codeobj_capture_start(id3);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  // Launch a kernel
+  LaunchVectorAddKernel();
+
+  result = rocprofiler_codeobj_capture_stop(id2);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  result = rocprofiler_codeobj_capture_stop(id3);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  result = rocprofiler_codeobj_capture_free(id3);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  rocprofiler_codeobj_symbols_t capture;
+  rocprofiler_codeobj_capture_get(id2, &capture);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  EXPECT_GE(capture.count, 1);
+
+  for (int i=0; i<(int)capture.count; i++) {
+    EXPECT_NE(capture.symbols[i].base_address, 0);
+    EXPECT_NE(capture.symbols[i].clock_start.value, 0);
+    EXPECT_NE(capture.symbols[i].data, nullptr);
+    EXPECT_NE(capture.symbols[i].size, 0);
+  }
+
+  result = rocprofiler_codeobj_capture_stop(id);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  result = rocprofiler_codeobj_capture_free(id2);
+  EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result);
+
+  // Expect to return error when running get() after free()
+  result = rocprofiler_codeobj_capture_get(id3, &capture);
+  EXPECT_NE(ROCPROFILER_STATUS_SUCCESS, result);
+
+  TearDownRocprofiler();
+}
+
 /*
  * ###################################################
  * ############ Plugin tests ################

tests-v2/unittests/CMakeLists.txt

@@ -5,6 +5,8 @@ find_library(PCIACCESS_LIBRARIES pciaccess REQUIRED)
 enable_testing()
 find_package(GTest REQUIRED)
 
+find_library(GDB rocm-dbgapi PATHS ${ROCM_PATH} REQUIRED)
+
 # Getting Source files for ROCProfiler, Hardware, HSA, Memory, Session, Counters, Utils
 set(CORE_MEMORY_DIR ${PROJECT_SOURCE_DIR}/src/core/memory)
 file(GLOB CORE_MEMORY_SRC_FILES ${CORE_MEMORY_DIR}/*.cpp)
@@ -35,14 +37,15 @@ 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)
-file(GLOB 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_SRC_CLASS_FILES
      ${CMAKE_CURRENT_SOURCE_DIR}/rocprofiler_singleton.cpp)
+file(GLOB ROCPROFILER_ISA_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/isa_capture/*.cpp)
 file(GLOB ROCPROFILER_SPM_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/spm/spm.cpp)
 file(GLOB ROCPROFILER_SRC_API_FILES ${PROJECT_SOURCE_DIR}/src/api/*.cpp)
 
-set(ROCPROFILER_SRC_FILES ${ROCPROFILER_SRC_API_FILES} ${ROCPROFILER_SRC_CLASS_FILES}
-                          ${ROCPROFILER_SRC_PROFILER_FILES} ${ROCPROFILER_ATT_SRC_FILES})
+set(ROCPROFILER_SRC_FILES ${ROCPROFILER_SRC_API_FILES} ${ROCPROFILER_ATT_SRC_FILES}
+     ${ROCPROFILER_ISA_SRC_FILES} ${ROCPROFILER_SRC_PROFILER_FILES} ${ROCPROFILER_ATT_SRC_FILES})
 
 set(CORE_HSA_DIR ${PROJECT_SOURCE_DIR}/src/core/hsa)
 file(GLOB CORE_HSA_SRC_FILES ${CORE_HSA_DIR}/*.cpp)
@@ -98,7 +101,7 @@ target_compile_definitions(
 
 target_link_libraries(
     runUnitTests PRIVATE rocprofiler_tool ${AQLPROFILE_LIB} hsa-runtime64::hsa-runtime64
-                         GTest::gtest GTest::gtest_main stdc++fs ${PCIACCESS_LIBRARIES})
+          GTest::gtest GTest::gtest_main stdc++fs ${PCIACCESS_LIBRARIES} ${GDB} dw elf c dl)
 
 add_dependencies(tests runUnitTests)
 install(TARGETS runUnitTests