SWDEV-391589: Fixing File plugin output text

Change-Id: I5e1d46431ca2b93661772b062996ed62574c36f9

plugin/CMakeLists.txt

@@ -24,3 +24,4 @@ add_subdirectory(file)
 add_subdirectory(perfetto)
 add_subdirectory(ctf)
 add_subdirectory(att)
+add_subdirectory(cli)

plugin/att/att.cpp

@@ -79,9 +79,10 @@ class att_plugin_t {
     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));
+    std::string name_demangled =
+        rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c));
 
-    if (name_demangled.size() > ATT_FILENAME_MAXBYTES) { // Limit filename size
+    if (name_demangled.size() > ATT_FILENAME_MAXBYTES) {  // Limit filename size
       name_demangled = name_demangled.substr(0, ATT_FILENAME_MAXBYTES);
     }
 
@@ -167,7 +168,8 @@ att_plugin_t* att_plugin = nullptr;
 }  // namespace
 
 ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version,
-                                                     uint32_t rocprofiler_minor_version) {
+                                                     uint32_t rocprofiler_minor_version,
+                                                     void* data) {
   if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR ||
       rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR)
     return -1;

plugin/cli/CMakeLists.txt

@@ -0,0 +1,47 @@
+# ###############################################################################
+# # 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.
+# ###############################################################################
+
+file(GLOB ROCPROFILER_UTIL_SRC_FILES ${PROJECT_SOURCE_DIR}/src/utils/helper.cpp)
+
+file(GLOB CLI_SOURCES "*.cpp")
+add_library(cli_plugin SHARED ${CLI_SOURCES} ${ROCPROFILER_UTIL_SRC_FILES})
+
+set_target_properties(cli_plugin PROPERTIES
+  CXX_VISIBILITY_PRESET hidden
+  LINK_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/../exportmap
+  LIBRARY_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR})
+
+target_compile_definitions(cli_plugin
+  PRIVATE HIP_PROF_HIP_API_STRING=1 __HIP_PLATFORM_HCC__=1)
+
+target_include_directories(cli_plugin PRIVATE ${PROJECT_SOURCE_DIR})
+
+target_link_options(cli_plugin PRIVATE -Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/../exportmap -Wl,--no-undefined)
+
+target_link_libraries(cli_plugin PRIVATE rocprofiler-v2 hsa-runtime64::hsa-runtime64 stdc++fs atomic amd_comgr dl)
+
+install(TARGETS cli_plugin LIBRARY
+  DESTINATION ${CMAKE_INSTALL_LIBDIR}/${PROJECT_NAME}
+  COMPONENT asan)
+install(TARGETS cli_plugin LIBRARY
+  DESTINATION ${CMAKE_INSTALL_LIBDIR}/${PROJECT_NAME}
+  COMPONENT runtime)

plugin/cli/cli.cpp

@@ -0,0 +1,484 @@
+/* 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 <cxxabi.h>
+#include <stdarg.h>
+#include <stdio.h>
+#include <string.h>
+#include <sys/syscall.h>
+#include <sys/types.h>
+#include <unistd.h>
+
+#include <atomic>
+#include <cassert>
+#include <cstddef>
+#include <cstdint>
+#include <experimental/filesystem>
+#include <fstream>
+#include <iostream>
+#include <memory>
+#include <optional>
+#include <ostream>
+#include <sstream>
+#include <string>
+#include <hsa/hsa.h>
+#include <mutex>
+#include <unordered_map>
+
+#include "rocprofiler.h"
+#include "rocprofiler_plugin.h"
+#include "../utils.h"
+
+namespace fs = std::experimental::filesystem;
+
+namespace {
+
+static std::string output_file_name;
+class file_plugin_t {
+ private:
+  enum class output_type_t { COUNTER, TRACER, PC_SAMPLING };
+
+  class output_file_t {
+   public:
+    output_file_t(std::string name) : name_(std::move(name)) {}
+
+    std::string name() const { return name_; }
+
+    template <typename T> std::ostream& operator<<(T&& value) {
+      if (!is_open()) open();
+      return stream_ << std::forward<T>(value);
+    }
+
+    std::ostream& operator<<(std::ostream& (*func)(std::ostream&)) {
+      if (!is_open()) open();
+      return stream_ << func;
+    }
+
+    void open() {
+      // If the stream is already in the failed state, there's no need to try
+      // to open the file.
+      if (fail()) return;
+
+      const char* output_dir = getenv("OUTPUT_PATH");
+      output_file_name = getenv("OUT_FILE_NAME") ? std::string(getenv("OUT_FILE_NAME")) + "_" : "";
+
+      if (output_dir == nullptr && getenv("OUT_FILE_NAME") == nullptr) {
+        stream_.copyfmt(std::cout);
+        stream_.clear(std::cout.rdstate());
+        stream_.basic_ios<char>::rdbuf(std::cout.rdbuf());
+        return;
+      }
+      if (output_dir == nullptr) output_dir = "./";
+
+      fs::path output_prefix(output_dir);
+      if (!fs::is_directory(fs::status(output_prefix))) {
+        if (!stream_.fail()) rocprofiler::warning("Cannot open output directory '%s'", output_dir);
+        stream_.setstate(std::ios_base::failbit);
+        return;
+      }
+
+      std::stringstream ss;
+      output_file_name = replace_MPI_macros(output_file_name);
+
+      ss << output_file_name << GetPid() << "_" << name_;
+      stream_.open(output_prefix / ss.str());
+    }
+
+    bool is_open() const { return stream_.is_open(); }
+    bool fail() const { return stream_.fail(); }
+
+    // Returns a string with the MPI %macro replaced with the corresponding envvar
+    std::string replace_MPI_macros(std::string output_file_name) {
+      std::unordered_map<const char*, const char*> MPI_BUILTINS = {
+          {"MPI_RANK", "%rank"},
+          {"OMPI_COMM_WORLD_RANK", "%rank"},
+          {"MV2_COMM_WORLD_RANK", "%rank"}};
+
+      for (const auto& [envvar, key] : MPI_BUILTINS) {
+        size_t key_find = output_file_name.rfind(key);
+        if (key_find == std::string::npos) continue;  // Does not contain a %?rank var
+
+        const char* env_var_set = getenv(envvar);
+        if (env_var_set == nullptr) continue;  // MPI_COMM_WORLD_x var is does not exist
+
+        int rank = atoi(env_var_set);
+        output_file_name = output_file_name.substr(0, key_find) + std::to_string(rank) +
+            output_file_name.substr(key_find + std::string(key).size());
+      }
+
+      return output_file_name;
+    }
+
+   private:
+    const std::string name_;
+    std::ofstream stream_;
+  };
+
+  output_file_t* get_output_file(output_type_t output_type, uint32_t domain = 0) {
+    switch (output_type) {
+      case output_type_t::COUNTER:
+        return &output_file_;
+      case output_type_t::TRACER:
+        switch (domain) {
+          case ACTIVITY_DOMAIN_ROCTX:
+            return &roctx_file_;
+          case ACTIVITY_DOMAIN_HSA_API:
+            return &hsa_api_file_;
+          case ACTIVITY_DOMAIN_HIP_API:
+            return &hip_api_file_;
+          case ACTIVITY_DOMAIN_HIP_OPS:
+            return &hip_activity_file_;
+          case ACTIVITY_DOMAIN_HSA_OPS:
+            return &hsa_async_copy_file_;
+          default:
+            assert(!"domain/op not supported!");
+            break;
+        }
+        break;
+      case output_type_t::PC_SAMPLING:
+        return &pc_sample_file_;
+    }
+    return nullptr;
+  }
+
+ public:
+  file_plugin_t() {
+    valid_ = true;
+  }
+
+  std::mutex writing_lock;
+
+  const char* GetDomainName(rocprofiler_tracer_activity_domain_t domain) {
+    switch (domain) {
+      case ACTIVITY_DOMAIN_ROCTX:
+        return "ROCTX_DOMAIN";
+        break;
+      case ACTIVITY_DOMAIN_HIP_API:
+        return "HIP_API_DOMAIN";
+        break;
+      case ACTIVITY_DOMAIN_HIP_OPS:
+        return "HIP_OPS_DOMAIN";
+        break;
+      case ACTIVITY_DOMAIN_HSA_API:
+        return "HSA_API_DOMAIN";
+        break;
+      case ACTIVITY_DOMAIN_HSA_OPS:
+        return "HSA_OPS_DOMAIN";
+        break;
+      case ACTIVITY_DOMAIN_HSA_EVT:
+        return "HSA_EVT_DOMAIN";
+        break;
+      default:
+        return "";
+    }
+  }
+
+  void FlushTracerRecord(rocprofiler_record_tracer_t tracer_record,
+                         rocprofiler_session_id_t session_id,
+                         rocprofiler_buffer_id_t buffer_id = rocprofiler_buffer_id_t{0}) {
+    std::lock_guard<std::mutex> lock(writing_lock);
+    if (tracer_record.timestamps.end.value <= 0 && tracer_record.domain != ACTIVITY_DOMAIN_ROCTX)
+      return;
+    std::string function_name;
+    std::string kernel_name;
+    std::string roctx_message;
+    uint64_t roctx_id;
+    if ((tracer_record.operation_id.id == 0 && tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS)) {
+      if (tracer_record.name) {
+        kernel_name = rocprofiler::cxx_demangle(tracer_record.name);
+        std::string key = "\"";
+        std::size_t found = kernel_name.rfind(key);
+        while (found != std::string::npos) {
+          kernel_name.replace(found, key.length(), "'");
+          found = kernel_name.rfind(key, found - 1);
+        }
+      }
+    }
+    size_t function_name_size = 0;
+    char* function_name_c = nullptr;
+    if (tracer_record.domain == ACTIVITY_DOMAIN_HSA_API) {
+      CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info_size(
+          rocprofiler_session_id_t{0}, ROCPROFILER_HSA_FUNCTION_NAME, tracer_record.api_data_handle,
+          tracer_record.operation_id, &function_name_size));
+      function_name_c = new char[function_name_size];
+      if (function_name_size > 1) {
+        CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info(
+            rocprofiler_session_id_t{0}, ROCPROFILER_HSA_FUNCTION_NAME,
+            tracer_record.api_data_handle, tracer_record.operation_id, &function_name_c));
+        if (function_name_c) function_name = std::string(function_name_c);
+      }
+    }
+    if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_API) {
+      CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info_size(
+          rocprofiler_session_id_t{0}, ROCPROFILER_HIP_FUNCTION_NAME, tracer_record.api_data_handle,
+          tracer_record.operation_id, &function_name_size));
+      if (function_name_size > 1) {
+        CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info(
+            session_id, ROCPROFILER_HIP_FUNCTION_NAME, tracer_record.api_data_handle,
+            tracer_record.operation_id, &function_name_c));
+        if (function_name_c) function_name = std::string(function_name_c);
+      }
+      if (tracer_record.name) {
+        kernel_name = rocprofiler::cxx_demangle(std::string(tracer_record.name));
+        std::string key = "\"";
+        std::size_t found = kernel_name.rfind(key);
+        while (found != std::string::npos) {
+          kernel_name.replace(found, key.length(), "'");
+          found = kernel_name.rfind(key, found - 1);
+        }
+        // TODO: Change how this API returns a string.
+      }
+    }
+    if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX) {
+      if (tracer_record.name) roctx_message = rocprofiler::cxx_demangle(tracer_record.name);
+      roctx_id = tracer_record.operation_id.id;
+    }
+    char* activity_name = nullptr;
+    if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS) {
+      if (tracer_record.api_data_handle.handle) {
+        kernel_name = rocprofiler::cxx_demangle(
+            const_cast<char*>(reinterpret_cast<const char*>(tracer_record.api_data_handle.handle)));
+      }
+      size_t activity_name_size = 0;
+      CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info_size(
+          session_id, ROCPROFILER_HIP_ACTIVITY_NAME, tracer_record.api_data_handle,
+          tracer_record.operation_id, &activity_name_size));
+      if (activity_name_size > 1) {
+        activity_name = nullptr;
+        CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info(
+            session_id, ROCPROFILER_HIP_ACTIVITY_NAME, tracer_record.api_data_handle,
+            tracer_record.operation_id, &activity_name));
+      }
+    }
+    if (tracer_record.domain == ACTIVITY_DOMAIN_HSA_OPS) {
+      size_t activity_name_size = 0;
+      CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info_size(
+          session_id, ROCPROFILER_HSA_ACTIVITY_NAME, tracer_record.api_data_handle,
+          tracer_record.operation_id, &activity_name_size));
+      if (activity_name_size > 1) {
+        activity_name = nullptr;
+        CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info(
+            session_id, ROCPROFILER_HSA_ACTIVITY_NAME, tracer_record.api_data_handle,
+            tracer_record.operation_id, &activity_name));
+      }
+    }
+
+    output_file_t* output_file = get_output_file(output_type_t::TRACER, tracer_record.domain);
+    *output_file << "Record_ID(" << tracer_record.header.id.handle << "), "
+                 << "Domain(" << GetDomainName(tracer_record.domain) << "), ";
+    if (function_name.size() > 1) *output_file << "Function(" << function_name << "), ";
+    if (activity_name) *output_file << "Operation_Name(" << activity_name << "), ";
+    if (kernel_name.size() > 1) *output_file << "Kernel_Name(" << kernel_name.c_str() << "), ";
+    if (tracer_record.domain != ACTIVITY_DOMAIN_ROCTX) {
+      *output_file << "Start_Timestamp(" << tracer_record.timestamps.begin.value << "), "
+                   << "End_Timestamp(" << tracer_record.timestamps.end.value << "), "
+                   << "Correlation_ID(" << tracer_record.correlation_id.value << ")";
+    } else {
+      *output_file << "Timestamp(" << tracer_record.timestamps.begin.value << "), ";
+    }
+    if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX && roctx_id >= 0)
+      *output_file << "ROCTX_ID(" << roctx_id << "), ";
+    if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX && roctx_message.size() > 1)
+      *output_file << "ROCTX_Message(" << roctx_message << ")";
+    *output_file << std::endl;
+  }
+
+  void FlushProfilerRecord(const rocprofiler_record_profiler_t* profiler_record,
+                           rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) {
+    std::lock_guard<std::mutex> lock(writing_lock);
+    size_t name_length = 0;
+    output_file_t* output_file{nullptr};
+    output_file = get_output_file(output_type_t::COUNTER);
+    CHECK_ROCPROFILER(rocprofiler_query_kernel_info_size(ROCPROFILER_KERNEL_NAME,
+                                                         profiler_record->kernel_id, &name_length));
+    // Taken from rocprofiler: The size hasn't changed in  recent past
+    static const uint32_t lds_block_size = 128 * 4;
+    const char* kernel_name_c = nullptr;
+    if (name_length > 1) {
+      CHECK_ROCPROFILER(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME,
+                                                      profiler_record->kernel_id, &kernel_name_c));
+    }
+    *output_file << "Record_ID(" << std::to_string(profiler_record->header.id.handle) << "), "
+                 << "GPU_ID(" << std::to_string(profiler_record->gpu_id.handle) << "), "
+                 << "Queue_ID(" << std::to_string(profiler_record->queue_id.handle) << "), "
+                 << "Queue_Index(" << std::to_string(profiler_record->queue_idx.value) << "), "
+                 << "Process_ID(" << std::to_string(GetPid()) << "), "
+                 << "Thread_ID(" << std::to_string(profiler_record->thread_id.value) << "), "
+                 << "Grid_Size(" << std::to_string(profiler_record->kernel_properties.grid_size)
+                 << "), "
+                 << "Workgroup_Size("
+                 << std::to_string(profiler_record->kernel_properties.workgroup_size) << "), "
+                 << "LDS("
+                 << std::to_string(
+                        ((profiler_record->kernel_properties.lds_size + (lds_block_size - 1)) &
+                         ~(lds_block_size - 1)))
+                 << "), "
+                 << "Scratch_Size("
+                 << std::to_string(profiler_record->kernel_properties.scratch_size) << "), "
+                 << "Arch_VGPR("
+                 << std::to_string(profiler_record->kernel_properties.arch_vgpr_count) << "), "
+                 << "Accumulative_VGPR("
+                 << std::to_string(profiler_record->kernel_properties.accum_vgpr_count) << "), "
+                 << "SGPR(" << std::to_string(profiler_record->kernel_properties.sgpr_count)
+                 << "), "
+                 << "Wave_Size(" << std::to_string(profiler_record->kernel_properties.wave_size);
+    std::string kernel_name = "";
+    if (name_length > 1) {
+      kernel_name = rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c));
+      std::string key = "\"";
+      std::size_t found = kernel_name.rfind(key);
+      while (found != std::string::npos) {
+        kernel_name.replace(found, key.length(), "'");
+        found = kernel_name.rfind(key, found - 1);
+      }
+    }
+    *output_file << "), "
+                 << "Kernel_Name(\"" << kernel_name << "\"), "
+                 << "Begin_Timestamp(" << std::to_string(profiler_record->timestamps.begin.value)
+                 << "), "
+                 << "End_Timestamp(" << std::to_string(profiler_record->timestamps.end.value)
+                 << ")";
+
+    // For Counters
+    if (profiler_record->counters) {
+      *output_file << ", ";
+      for (uint64_t i = 0; i < profiler_record->counters_count.value; i++) {
+        if (profiler_record->counters[i].counter_handler.handle > 0) {
+          size_t counter_name_length = 0;
+          CHECK_ROCPROFILER(rocprofiler_query_counter_info_size(
+              session_id, ROCPROFILER_COUNTER_NAME, profiler_record->counters[i].counter_handler,
+              &counter_name_length));
+          if (counter_name_length > 1) {
+            const char* name_c = nullptr;
+            CHECK_ROCPROFILER(rocprofiler_query_counter_info(
+                session_id, ROCPROFILER_COUNTER_NAME, profiler_record->counters[i].counter_handler,
+                &name_c));
+            *output_file << name_c << "("
+                         << std::to_string(profiler_record->counters[i].value.value) << ")";
+            if (i < profiler_record->counters_count.value - 1) *output_file << ", ";
+          }
+        }
+      }
+    }
+    *output_file << std::endl;
+    if (kernel_name_c) {
+      free(const_cast<char*>(kernel_name_c));
+    }
+  }
+
+  void FlushPCSamplingRecord(const rocprofiler_record_pc_sample_t* pc_sampling_record) {
+    output_file_t* output_file{nullptr};
+    output_file = get_output_file(output_type_t::PC_SAMPLING);
+    const auto& sample = pc_sampling_record->pc_sample;
+    *output_file << "dispatch[" << sample.dispatch_id.value << "], "
+                 << "timestamp(" << sample.timestamp.value << "), "
+                 << "gpu_id(" << sample.gpu_id.handle << "), "
+                 << "pc-sample(" << std::hex << std::showbase << sample.pc << "), "
+                 << "se(" << sample.se << ')' << std::endl;
+  }
+  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;
+      switch (begin->kind) {
+        case ROCPROFILER_PROFILER_RECORD: {
+          const rocprofiler_record_profiler_t* profiler_record =
+              reinterpret_cast<const rocprofiler_record_profiler_t*>(begin);
+          FlushProfilerRecord(profiler_record, session_id, buffer_id);
+          break;
+        }
+        case ROCPROFILER_TRACER_RECORD: {
+          rocprofiler_record_tracer_t* tracer_record = const_cast<rocprofiler_record_tracer_t*>(
+              reinterpret_cast<const rocprofiler_record_tracer_t*>(begin));
+          FlushTracerRecord(*tracer_record, session_id, buffer_id);
+          break;
+        }
+        case ROCPROFILER_ATT_TRACER_RECORD: {
+          break;
+        }
+        case ROCPROFILER_PC_SAMPLING_RECORD: {
+          const rocprofiler_record_pc_sample_t* pc_sampling_record =
+              reinterpret_cast<const rocprofiler_record_pc_sample_t*>(begin);
+          FlushPCSamplingRecord(pc_sampling_record);
+          break;
+        }
+        default:
+          break;
+      }
+      rocprofiler_next_record(begin, &begin, session_id, buffer_id);
+    }
+    return 0;
+  }
+
+  bool is_valid() const { return valid_; }
+
+ private:
+  bool valid_{false};
+  std::atomic<bool> tracer_header_written_{false};
+  std::atomic<bool> profiler_header_written_{false};
+
+  output_file_t roctx_file_{"roctx_trace.txt"}, hsa_api_file_{"hsa_api_trace.txt"},
+      hip_api_file_{"hip_api_trace.txt"}, hip_activity_file_{"hcc_ops_trace.txt"},
+      hsa_async_copy_file_{"async_copy_trace.txt"}, pc_sample_file_{"pcs_trace.txt"},
+      output_file_{"results.txt"};
+};
+
+file_plugin_t* file_plugin = nullptr;
+
+}  // namespace
+
+ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version,
+                                                     uint32_t rocprofiler_minor_version,
+                                                     void* data) {
+  if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR ||
+      rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR)
+    return -1;
+
+  if (file_plugin != nullptr) return -1;
+
+  file_plugin = new file_plugin_t();
+  if (file_plugin->is_valid()) return 0;
+
+  // The plugin failed to initialized, destroy it and return an error.
+  delete file_plugin;
+  file_plugin = nullptr;
+  return -1;
+}
+
+ROCPROFILER_EXPORT void rocprofiler_plugin_finalize() {
+  if (!file_plugin) return;
+  delete file_plugin;
+  file_plugin = nullptr;
+}
+
+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 (!file_plugin || !file_plugin->is_valid()) return -1;
+  return file_plugin->WriteBufferRecords(begin, end, session_id, buffer_id);
+}
+
+ROCPROFILER_EXPORT int rocprofiler_plugin_write_record(rocprofiler_record_tracer_t record) {
+  if (!file_plugin || !file_plugin->is_valid()) return -1;
+  if (record.header.id.handle == 0) return 0;
+  file_plugin->FlushTracerRecord(record, rocprofiler_session_id_t{0}, rocprofiler_buffer_id_t{0});
+  return 0;
+}

plugin/ctf/ctf.cpp

@@ -38,7 +38,8 @@ rocm_ctf::Plugin* the_plugin = nullptr;
 }  // namespace
 
 ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(const uint32_t rocprofiler_major_version,
-                                                     const uint32_t rocprofiler_minor_version) {
+                                                     const uint32_t rocprofiler_minor_version,
+                                                     void* data) {
   if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR ||
       rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR) {
     return -1;

plugin/file/file.cpp

@@ -26,6 +26,7 @@
 #include <sys/types.h>
 #include <unistd.h>
 
+#include <atomic>
 #include <cassert>
 #include <cstddef>
 #include <cstdint>
@@ -49,6 +50,36 @@ namespace fs = std::experimental::filesystem;
 
 namespace {
 
+std::vector<std::string> GetCounterNames() {
+  std::vector<std::string> counters;
+  const char* line_c_str = getenv("ROCPROFILER_COUNTERS");
+  if (line_c_str) {
+    std::string line = line_c_str;
+    // skip commented lines
+    auto found = line.find_first_not_of(" \t");
+    if (found != std::string::npos) {
+      if (line[found] == '#') return {};
+    }
+    if (line.find("pmc") == std::string::npos) return counters;
+    char seperator = ' ';
+    std::string::size_type prev_pos = 0, pos = line.find(seperator, prev_pos);
+    prev_pos = ++pos;
+    if (pos != std::string::npos) {
+      while ((pos = line.find(seperator, pos)) != std::string::npos) {
+        std::string substring(line.substr(prev_pos, pos - prev_pos));
+        if (substring.length() > 0 && substring != ":") {
+          counters.push_back(substring);
+        }
+        prev_pos = ++pos;
+      }
+      if (!line.substr(prev_pos, pos - prev_pos).empty()) {
+        counters.push_back(line.substr(prev_pos, pos - prev_pos));
+      }
+    }
+  }
+  return counters;
+}
+
 static std::string output_file_name;
 class file_plugin_t {
  private:
@@ -78,7 +109,7 @@ class file_plugin_t {
       if (fail()) return;
 
       const char* output_dir = getenv("OUTPUT_PATH");
-      output_file_name = getenv("OUT_FILE_NAME") ? std::string(getenv("OUT_FILE_NAME")) + "_" : "";
+      output_file_name = getenv("OUT_FILE_NAME") ? std::string(getenv("OUT_FILE_NAME")) : "";
 
       if (output_dir == nullptr && getenv("OUT_FILE_NAME") == nullptr) {
         stream_.copyfmt(std::cout);
@@ -96,10 +127,12 @@ class file_plugin_t {
         return;
       }
 
-      std::stringstream ss;
       output_file_name = replace_MPI_macros(output_file_name);
 
-      ss << output_file_name << GetPid() << "_" << name_;
+      std::stringstream ss;
+      ss << name_ << "_" << ((output_file_name.empty()) ? std::to_string(GetPid()) : "")
+         << output_file_name << ".csv";
+      std::cout << "Results File: " << output_prefix / ss.str() << std::endl;
       stream_.open(output_prefix / ss.str());
     }
 
@@ -163,45 +196,90 @@ class file_plugin_t {
   }
 
  public:
-  file_plugin_t() {
-    output_file_t hsa_handles("hsa_handles.txt", true);
-
-    [[maybe_unused]] hsa_status_t status = hsa_iterate_agents(
-        [](hsa_agent_t agent, void* user_data) {
-          auto* file = static_cast<decltype(hsa_handles)*>(user_data);
-          hsa_device_type_t type;
-
-          if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type) != HSA_STATUS_SUCCESS)
-            return HSA_STATUS_ERROR;
-
-          if (!file->isStdOut())
-            *file << std::hex << std::showbase << agent.handle << " agent "
-                  << ((type == HSA_DEVICE_TYPE_CPU) ? "cpu" : "gpu") << std::endl;
-          return HSA_STATUS_SUCCESS;
-        },
-        &hsa_handles);
-    assert(status == HSA_STATUS_SUCCESS && "failed to iterate HSA agents");
-    if (hsa_handles.fail()) {
-      rocprofiler::warning("Cannot write to '%s'", hsa_handles.name().c_str());
-      return;
-    }
-
-    // App begin timestamp begin_ts_file.txt
-    output_file_t begin_ts("begin_ts_file.txt", true);
-
-    [[maybe_unused]] rocprofiler_timestamp_t app_begin_timestamp = {};
-    CHECK_ROCPROFILER(rocprofiler_get_timestamp(&app_begin_timestamp));
-
-    if (!begin_ts.isStdOut()) begin_ts << std::dec << app_begin_timestamp.value << std::endl;
-
-    if (begin_ts.fail()) {
-      rocprofiler::warning("Cannot write to '%s'", begin_ts.name().c_str());
-      return;
-    }
+  file_plugin_t(void* data) {
+    if (data) counter_names_ = GetCounterNames();
 
     valid_ = true;
   }
 
+  void WriteHeader(output_type_t type, rocprofiler_tracer_activity_domain_t domain) {
+    output_file_t* output_file;
+    switch (domain) {
+      case ACTIVITY_DOMAIN_HSA_API: {
+        if (hsa_api_header_written_.load(std::memory_order_relaxed)) return;
+        output_file = get_output_file(output_type_t::TRACER, ACTIVITY_DOMAIN_HSA_API);
+        *output_file << "Record_ID,Domain,Function,Start_Timestamp,End_Timestamp,Correlation_ID"
+                     << std::endl;
+        *output_file << std::endl;
+        hsa_api_header_written_.exchange(true, std::memory_order_release);
+        return;
+      }
+      case ACTIVITY_DOMAIN_HIP_API: {
+        if (hip_api_header_written_.load(std::memory_order_relaxed)) return;
+        output_file = get_output_file(output_type_t::TRACER, ACTIVITY_DOMAIN_HIP_API);
+        *output_file
+            << "Record_ID,Domain,Function,Kernel_Name,Start_Timestamp,End_Timestamp,Correlation_ID"
+            << std::endl;
+        *output_file << std::endl;
+        hip_api_header_written_.exchange(true, std::memory_order_release);
+        return;
+      }
+      case ACTIVITY_DOMAIN_ROCTX: {
+        if (roctx_header_written_.load(std::memory_order_relaxed)) return;
+        output_file = get_output_file(output_type_t::TRACER, ACTIVITY_DOMAIN_ROCTX);
+        *output_file << "Record_ID,Domain,ROCTX_ID,Message,Timestamp" << std::endl;
+        *output_file << std::endl;
+        roctx_header_written_.exchange(true, std::memory_order_release);
+        return;
+      }
+      case ACTIVITY_DOMAIN_HSA_OPS: {
+        if (hsa_async_copy_header_written_.load(std::memory_order_relaxed)) return;
+        output_file = get_output_file(output_type_t::TRACER, ACTIVITY_DOMAIN_HSA_OPS);
+        *output_file << "Record_ID,Domain,Operation,Start_Timestamp,Stop_Timestamp,Correlation_ID"
+                     << std::endl;
+        *output_file << std::endl;
+        hsa_async_copy_header_written_.exchange(true, std::memory_order_release);
+        return;
+      }
+      case ACTIVITY_DOMAIN_HIP_OPS: {
+        if (hip_activity_header_written_.load(std::memory_order_relaxed)) return;
+        output_file = get_output_file(output_type_t::TRACER, ACTIVITY_DOMAIN_HIP_OPS);
+        *output_file << "Record_ID,Domain,Operation,Kernel_Name,Start_Timestamp,Stop_Timestamp,"
+                        "Correlation_ID"
+                     << std::endl;
+        *output_file << std::endl;
+        hip_activity_header_written_.exchange(true, std::memory_order_release);
+        return;
+      }
+      default: {
+        if (type == output_type_t::COUNTER) {
+          if (kernel_dispatches_header_written_.load(std::memory_order_relaxed)) return;
+          output_file = get_output_file(output_type_t::COUNTER);
+
+          *output_file
+              << "Dispatch_ID,GPU_ID,Queue_ID,Queue_Index,PID,TID,GRD,WGR,LDS,SCR,Arch_VGPR,"
+                 "ACCUM_VGPR,SGPR,Wave_Size,SIG,OBJ,Kernel_Name,Start_Timestamp,End_Timestamp";
+          if (counter_names_.size() > 0) {
+            for (uint32_t i = 0; i < counter_names_.size(); i++)
+              *output_file << "," << counter_names_[i];
+          }
+          *output_file << std::endl;
+          *output_file << std::endl;
+          kernel_dispatches_header_written_.exchange(true, std::memory_order_release);
+          return;
+        } else if (type == output_type_t::PC_SAMPLING) {
+          if (pc_sample_header_written_.load(std::memory_order_relaxed)) return;
+          output_file = get_output_file(output_type_t::PC_SAMPLING);
+          *output_file << "Dispatch_ID,Timestamp,GPU_ID,PC_Sample,Shader_Engines" << std::endl;
+          *output_file << std::endl;
+          pc_sample_header_written_.exchange(true, std::memory_order_release);
+          return;
+        }
+        return;
+      }
+    }
+  }
+
   std::mutex writing_lock;
 
   const char* GetDomainName(rocprofiler_tracer_activity_domain_t domain) {
@@ -235,13 +313,21 @@ class file_plugin_t {
     std::lock_guard<std::mutex> lock(writing_lock);
     if (tracer_record.timestamps.end.value <= 0 && tracer_record.domain != ACTIVITY_DOMAIN_ROCTX)
       return;
+    WriteHeader(output_type_t::TRACER, tracer_record.domain);
+    std::string function_name;
     std::string kernel_name;
     std::string roctx_message;
     uint64_t roctx_id;
-    if (tracer_record.name) {
-      if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_API)
+    if ((tracer_record.operation_id.id == 0 && tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS)) {
+      if (tracer_record.name) {
         kernel_name = rocprofiler::cxx_demangle(tracer_record.name);
-      if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX) roctx_message = tracer_record.name;
+        std::string key = "\"";
+        std::size_t found = kernel_name.rfind(key);
+        while (found != std::string::npos) {
+          kernel_name.replace(found, key.length(), "'");
+          found = kernel_name.rfind(key, found - 1);
+        }
+      }
     }
     size_t function_name_size = 0;
     char* function_name_c = nullptr;
@@ -253,6 +339,7 @@ class file_plugin_t {
         CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info(
             rocprofiler_session_id_t{0}, ROCPROFILER_HSA_FUNCTION_NAME,
             tracer_record.api_data_handle, tracer_record.operation_id, &function_name_c));
+        if (function_name_c) function_name = std::string(function_name_c);
       }
     }
     if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_API) {
@@ -261,51 +348,87 @@ class file_plugin_t {
           tracer_record.operation_id, &function_name_size));
       if (function_name_size > 1) {
         CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info(
-            rocprofiler_session_id_t{0}, ROCPROFILER_HIP_FUNCTION_NAME,
-            tracer_record.api_data_handle, tracer_record.operation_id, &function_name_c));
+            session_id, ROCPROFILER_HIP_FUNCTION_NAME, tracer_record.api_data_handle,
+            tracer_record.operation_id, &function_name_c));
+        if (function_name_c) function_name = std::string(function_name_c);
       }
-    }
-    output_file_t* output_file = get_output_file(output_type_t::TRACER, tracer_record.domain);
-    *output_file << "Record(" << tracer_record.header.id.handle << "), Domain("
-                 << GetDomainName(tracer_record.domain) << "),";
-    if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX && roctx_id >= 0)
-      *output_file << " ROCTX_ID(" << tracer_record.operation_id.id << "),";
-    if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX && tracer_record.name)
-      *output_file << " ROCTX_Message(" << reinterpret_cast<const char*>(tracer_record.name)
-                   << "),";
-    if (function_name_c) *output_file << " Function(" << function_name_c << "),";
-    if (kernel_name.size() > 1) *output_file << " Kernel_Name(" << kernel_name.c_str() << "),";
-    if (tracer_record.domain == ACTIVITY_DOMAIN_HSA_OPS ||
-        tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS) {
-      switch (tracer_record.operation_id.id) {
-        case 0:
-          *output_file << " Operation(DISPATCH_OP),";
-          break;
-        case 1:
-          *output_file << " Operation(COPY_OP),";
-          break;
-        case 2:
-          *output_file << " Operation(BARRIER_OP),";
-          break;
-        default:
-          break;
+      if (tracer_record.name) {
+        kernel_name = rocprofiler::cxx_demangle(std::string(tracer_record.name));
+        std::string key = "\"";
+        std::size_t found = kernel_name.rfind(key);
+        while (found != std::string::npos) {
+          kernel_name.replace(found, key.length(), "'");
+          found = kernel_name.rfind(key, found - 1);
+        }
+        // TODO: Change how this API returns a string.
       }
     }
     if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX) {
-      *output_file << " timestamp(" << tracer_record.timestamps.begin.value << ")";
-    } else if (tracer_record.phase == ROCPROFILER_PHASE_EXIT ||
-               tracer_record.phase == ROCPROFILER_PHASE_NONE) {
-      *output_file << " Begin(" << tracer_record.timestamps.begin.value << "), End("
-                   << tracer_record.timestamps.end.value << ")";
+      if (tracer_record.name) roctx_message = rocprofiler::cxx_demangle(tracer_record.name);
+      roctx_id = tracer_record.operation_id.id;
     }
-    if (tracer_record.domain != ACTIVITY_DOMAIN_ROCTX)
-      *output_file << ", Correlation_ID(" << tracer_record.correlation_id.value << ")";
-    *output_file << '\n';
+    char* activity_name = nullptr;
+    if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS) {
+      if (tracer_record.api_data_handle.handle) {
+        kernel_name = rocprofiler::cxx_demangle(
+            const_cast<char*>(reinterpret_cast<const char*>(tracer_record.api_data_handle.handle)));
+      }
+      size_t activity_name_size = 0;
+      CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info_size(
+          session_id, ROCPROFILER_HIP_ACTIVITY_NAME, tracer_record.api_data_handle,
+          tracer_record.operation_id, &activity_name_size));
+      if (activity_name_size > 1) {
+        activity_name = nullptr;
+        CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info(
+            session_id, ROCPROFILER_HIP_ACTIVITY_NAME, tracer_record.api_data_handle,
+            tracer_record.operation_id, &activity_name));
+      }
+    }
+    if (tracer_record.domain == ACTIVITY_DOMAIN_HSA_OPS) {
+      size_t activity_name_size = 0;
+      CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info_size(
+          session_id, ROCPROFILER_HSA_ACTIVITY_NAME, tracer_record.api_data_handle,
+          tracer_record.operation_id, &activity_name_size));
+      if (activity_name_size > 1) {
+        activity_name = nullptr;
+        CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info(
+            session_id, ROCPROFILER_HSA_ACTIVITY_NAME, tracer_record.api_data_handle,
+            tracer_record.operation_id, &activity_name));
+      }
+    }
+    // return;
+    output_file_t* output_file = get_output_file(output_type_t::TRACER, tracer_record.domain);
+    *output_file << "" << tracer_record.header.id.handle << ","
+                 << GetDomainName(tracer_record.domain);
+    if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX && roctx_id >= 0)
+      *output_file << "," << roctx_id;
+    if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX) {
+      if (roctx_message.size() > 1)
+        *output_file << ",\"" << roctx_message << "\"";
+      else
+        *output_file << ",";
+    }
+    if (function_name.size() > 1) *output_file << ",\"" << function_name << "\"";
+    if (activity_name) *output_file << ",\"" << activity_name << "\"";
+    if (kernel_name.size() > 1)
+      *output_file << ",\"" << kernel_name.c_str() << "\"";
+    else if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_API ||
+             tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS)
+      *output_file << ",";
+    if (tracer_record.domain != ACTIVITY_DOMAIN_ROCTX) {
+      *output_file << "," << tracer_record.timestamps.begin.value << ","
+                   << tracer_record.timestamps.end.value;
+      *output_file << "," << tracer_record.correlation_id.value;
+    } else {
+      *output_file << "," << tracer_record.timestamps.begin.value;
+    }
+    *output_file << std::endl;
   }
 
   void FlushProfilerRecord(const rocprofiler_record_profiler_t* profiler_record,
                            rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) {
     std::lock_guard<std::mutex> lock(writing_lock);
+    WriteHeader(output_type_t::COUNTER, ACTIVITY_DOMAIN_NUMBER);
     size_t name_length = 0;
     output_file_t* output_file{nullptr};
     output_file = get_output_file(output_type_t::COUNTER);
@@ -318,62 +441,42 @@ class file_plugin_t {
       CHECK_ROCPROFILER(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME,
                                                       profiler_record->kernel_id, &kernel_name_c));
     }
-    *output_file << std::string("dispatch[") << std::to_string(profiler_record->header.id.handle)
-                 << "], " << std::string("gpu_id(")
-                 << std::to_string(profiler_record->gpu_id.handle) << "), "
-                 << std::string("queue_id(") << std::to_string(profiler_record->queue_id.handle)
-                 << "), " << std::string("queue_index(")
-                 << std::to_string(profiler_record->queue_idx.value) << "), " << std::string("pid(")
-                 << std::to_string(GetPid()) << "), " << std::string("tid(")
-                 << std::to_string(profiler_record->thread_id.value) << ")";
-    *output_file << ", " << std::string("grd(")
-                 << std::to_string(profiler_record->kernel_properties.grid_size) << "), "
-                 << std::string("wgr(")
-                 << std::to_string(profiler_record->kernel_properties.workgroup_size) << "), "
-                 << std::string("lds(")
+    *output_file << std::to_string(profiler_record->header.id.handle) << ","
+                 << std::to_string(profiler_record->gpu_id.handle) << ","
+                 << std::to_string(profiler_record->queue_id.handle) << ","
+                 << std::to_string(profiler_record->queue_idx.value) << ","
+                 << std::to_string(GetPid()) << ","
+                 << std::to_string(profiler_record->thread_id.value);
+    *output_file << "," << std::to_string(profiler_record->kernel_properties.grid_size) << ","
+                 << std::to_string(profiler_record->kernel_properties.workgroup_size) << ","
                  << std::to_string(
                         ((profiler_record->kernel_properties.lds_size + (lds_block_size - 1)) &
                          ~(lds_block_size - 1)))
-                 << "), " << std::string("scr(")
-                 << std::to_string(profiler_record->kernel_properties.scratch_size) << "), "
-                 << std::string("arch_vgpr(")
-                 << std::to_string(profiler_record->kernel_properties.arch_vgpr_count) << "), "
-                 << std::string("accum_vgpr(")
-                 << std::to_string(profiler_record->kernel_properties.accum_vgpr_count) << "), "
-                 << std::string("sgpr(")
-                 << std::to_string(profiler_record->kernel_properties.sgpr_count) << "), "
-                 << std::string("wave_size(")
-                 << std::to_string(profiler_record->kernel_properties.wave_size) << "), "
-                 << std::string("sig(")
+                 << "," << std::to_string(profiler_record->kernel_properties.scratch_size) << ","
+                 << std::to_string(profiler_record->kernel_properties.arch_vgpr_count) << ","
+                 << std::to_string(profiler_record->kernel_properties.accum_vgpr_count) << ","
+                 << std::to_string(profiler_record->kernel_properties.sgpr_count) << ","
+                 << std::to_string(profiler_record->kernel_properties.wave_size) << ","
                  << std::to_string(profiler_record->kernel_properties.signal_handle);
     std::string kernel_name = "";
     if (name_length > 1) {
-      kernel_name = rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c));
+      kernel_name = rocprofiler::cxx_demangle(kernel_name_c);
+      std::string key = "\"";
+      std::size_t found = kernel_name.rfind(key);
+      while (found != std::string::npos) {
+        kernel_name.replace(found, key.length(), "'");
+        found = kernel_name.rfind(key, found - 1);
+      }
     }
-    *output_file << "), " << std::string("obj(")
-                 << std::to_string(profiler_record->kernel_id.handle) << "), "
-                 << std::string("kernel-name(\"") << kernel_name << "\")"
-                 << std::string(", start_time(")
-                 << std::to_string(profiler_record->timestamps.begin.value) << ")"
-                 << std::string(", end_time(")
-                 << std::to_string(profiler_record->timestamps.end.value) << ")";
+    *output_file << "," << std::to_string(profiler_record->kernel_id.handle) << ",\"" << kernel_name
+                 << "\"," << std::to_string(profiler_record->timestamps.begin.value) << ","
+                 << std::to_string(profiler_record->timestamps.end.value);
 
     // For Counters
     if (profiler_record->counters) {
       for (uint64_t i = 0; i < profiler_record->counters_count.value; i++) {
         if (profiler_record->counters[i].counter_handler.handle > 0) {
-          size_t counter_name_length = 0;
-          CHECK_ROCPROFILER(rocprofiler_query_counter_info_size(
-              session_id, ROCPROFILER_COUNTER_NAME, profiler_record->counters[i].counter_handler,
-              &counter_name_length));
-          if (counter_name_length > 1) {
-            const char* name_c = nullptr;
-            CHECK_ROCPROFILER(rocprofiler_query_counter_info(
-                session_id, ROCPROFILER_COUNTER_NAME, profiler_record->counters[i].counter_handler,
-                &name_c));
-            *output_file << ", " << name_c << " ("
-                         << std::to_string(profiler_record->counters[i].value.value) << ')';
-          }
+          *output_file << "," << std::to_string(profiler_record->counters[i].value.value);
         }
       }
     }
@@ -384,14 +487,13 @@ class file_plugin_t {
   }
 
   void FlushPCSamplingRecord(const rocprofiler_record_pc_sample_t* pc_sampling_record) {
+    WriteHeader(output_type_t::PC_SAMPLING, ACTIVITY_DOMAIN_NUMBER);
     output_file_t* output_file{nullptr};
     output_file = get_output_file(output_type_t::PC_SAMPLING);
     const auto& sample = pc_sampling_record->pc_sample;
-    *output_file << "dispatch[" << sample.dispatch_id.value << "], "
-                 << "timestamp(" << sample.timestamp.value << "), "
-                 << "gpu_id(" << sample.gpu_id.handle << "), "
-                 << "pc-sample(" << std::hex << std::showbase << sample.pc << "), "
-                 << "se(" << sample.se << ')' << std::endl;
+    *output_file << sample.dispatch_id.value << "," << sample.timestamp.value << ","
+                 << sample.gpu_id.handle << "," << std::hex << std::showbase << sample.pc << ","
+                 << sample.se << std::endl;
   }
   int WriteBufferRecords(const rocprofiler_record_header_t* begin,
                          const rocprofiler_record_header_t* end,
@@ -432,11 +534,17 @@ class file_plugin_t {
 
  private:
   bool valid_{false};
+  std::vector<std::string> counter_names_;
 
-  output_file_t roctx_file_{"roctx_trace.txt"}, hsa_api_file_{"hsa_api_trace.txt"},
-      hip_api_file_{"hip_api_trace.txt"}, hip_activity_file_{"hcc_ops_trace.txt"},
-      hsa_async_copy_file_{"async_copy_trace.txt"}, pc_sample_file_{"pcs_trace.txt"},
-      output_file_{"results.txt"};
+  std::atomic<bool> roctx_header_written_{false}, hsa_api_header_written_{false},
+      hip_api_header_written_{false}, hip_activity_header_written_{false},
+      hsa_async_copy_header_written_{false}, pc_sample_header_written_{false},
+      kernel_dispatches_header_written_{false};
+
+  output_file_t roctx_file_{"roctx_trace"}, hsa_api_file_{"hsa_api_trace"},
+      hip_api_file_{"hip_api_trace"}, hip_activity_file_{"hcc_ops_trace"},
+      hsa_async_copy_file_{"async_copy_trace"}, pc_sample_file_{"pcs_trace"},
+      output_file_{"results"};
 };
 
 file_plugin_t* file_plugin = nullptr;
@@ -444,14 +552,15 @@ file_plugin_t* file_plugin = nullptr;
 }  // namespace
 
 ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version,
-                                                     uint32_t rocprofiler_minor_version) {
+                                                     uint32_t rocprofiler_minor_version,
+                                                     void* data) {
   if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR ||
       rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR)
     return -1;
 
   if (file_plugin != nullptr) return -1;
 
-  file_plugin = new file_plugin_t();
+  file_plugin = new file_plugin_t(data);
   if (file_plugin->is_valid()) return 0;
 
   // The plugin failed to initialized, destroy it and return an error.

plugin/perfetto/perfetto.cpp

@@ -768,7 +768,7 @@ perfetto_plugin_t* perfetto_plugin = nullptr;
 }  // namespace
 
 int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version,
-                                  uint32_t rocprofiler_minor_version) {
+                                  uint32_t rocprofiler_minor_version, void* data) {
   if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR ||
       rocprofiler_minor_version > ROCPROFILER_VERSION_MINOR)
     return -1;