Tools update (#397)

* Srnagara/tool counters collect (#331)

* Adding counter collection capability to tools

* Adding counter collection feature to tools

* Adding counter collection capability to tools

* Fixing merge down issues

* Small tool fixes for build + prevent profile realloc

* Reproducing the counter name query issue in buffered callback

* Minor fix for init order + sample that directly uses sdk-tool for debug purposes

* Adding a temporary fix to print the counter names

* Fixing the output file name and reverting the changes of caching the profile config

* Fixing SGPR_Count value

* cleaning up debug prints

* Adding header to counter collection file

* Adding kernel filtering support

* Remove threading

* Cleaning up the code

* Removing redundant prints

* Revert "Remove threading"

This reverts commit 05c58fb9de826e92cf8d2e3d1c31d5578525dcb4.

* Revert "Cleaning up the code"

This reverts commit 1d964882bf2396dee8ad020cbb6c83b36e0674e9.

* Changing the tools code to align with init-order fix

* cmake formatting (cmake-format) (#335)

Co-authored-by: SrirakshaNag <SrirakshaNag@users.noreply.github.com>

* source formatting (clang-format v11) (#336)

Co-authored-by: SrirakshaNag <SrirakshaNag@users.noreply.github.com>

* Adding support for async memory copy

* source formatting (clang-format v11) (#391)

Co-authored-by: SrirakshaNag <SrirakshaNag@users.noreply.github.com>

* Fixing header typo

* Fixing tool_fini

* Replaceing the direction and kind fields values with description

* Update lib/rocprofiler-sdk-tool/helper.cpp

- Remove use of VLA

* Update lib/rocprofiler-sdk-tool/tool.cpp

- Formatting

* Migrate common/config.* to rocprofiler-sdk-tool

* Update lib/rocprofiler-sdk-tool/tool.cpp

- fix clang-tidy issues

* source formatting (clang-format v11) (#392)

Co-authored-by: jrmadsen <jrmadsen@users.noreply.github.com>

* Update lib/common/mpl.hpp

- is_string_type / is_string_type_impl for deducing if type is a string type

* Update include/rocprofiler-sdk/fwd.h

- ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_NONE starts at zero

* Update lib/rocprofiler-sdk/hsa/async_copy.*

- functions for operation ids and names

* Update lib/rocprofiler-sdk/buffer_tracing.cpp

- support iterating and getting names for ROCPROFILER_BUFFER_TRACING_MEMORY_COPY

* Update lib/rocprofiler-sdk-tool/config.*

- env ROCPROFILER_ prefix -> ROCPROF_ prefix
- add support for memory copy tracing, counter collection, etc.

* Update lib/rocprofiler-sdk-tool/helper.*

- removed TracerFlushRecord
- removed cxa_demangle (use one in common library)
- removed GetCounterNames (handled in config)
- removed GetKernelNames (handled in config)

* Add lib/rocprofiler-sdk-tool/output_file.*

- separate out get_output_stream function and output_file struct from tool.cpp

* Add lib/rocprofiler-sdk-tool/csv.hpp

- write_csv_entry automatically quotes strings
- csv_encoder struct enforces correct number of columns

* Update lib/rocprofiler-sdk-tool/CMakeLists.txt

- add new files

* Update lib/rocprofiler-sdk-tool/tool.cpp

- update construction of output_file class
- add kernel_symbol_data for serializing kernel trace data
- use config instead of env lookups
- optimize counter collection profile config lookup/creation

* Update bin/rocprofv3

- rocprofv3 --help exits with 0 (as it should)
- command-line arg for memory copy tracing
- command-line arg for mangled kernels
- command-line arg for truncated kernels
- env ROCPROFILER_ prefix -> env ROCPROF_ prefix

* Update tests/async-copy-tracing/validate.py

- update test_async_copy_direction to new enum values

* Update tests/kernel-tracing/validate.py

- update test_async_copy_direction to new enum values

* Update tests/tools/json-tool.cpp

- add ROCPROFILER_BUFFER_TRACING_MEMORY_COPY to supported buffer_name_info

* Update samples/counter_collection/{CMakeLists.txt,main.cpp}

- remove counter-collection-sdk-tool

* Update .github/workflows/docs.yml

- fix paths triggering running the workflow

---------

Co-authored-by: Benjamin Welton <bewelton@amd.com>
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: SrirakshaNag <SrirakshaNag@users.noreply.github.com>
Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>
Co-authored-by: jrmadsen <jrmadsen@users.noreply.github.com>

* adding counter collection support

* Adding counter collection test

* changing directory structure of counter collection tests

* Fixing test path for rocprofv3

* Adding hsa-tracing basic test

* cmake formatting (cmake-format) (#362)

Co-authored-by: bgopesh <bgopesh@users.noreply.github.com>

* counter collection tests drop2

* fixing hsa-trace test for rocprofv3 path

* python formatting (black) (#371)

Co-authored-by: bgopesh <bgopesh@users.noreply.github.com>

* both counter colleciton and tracing should work together

* Fixing rocprofv3 path

* Attempt to fix Segfault with AddressSanitizer

* fixing sanitizer segfault

* Update rocprofv3

* Update lib/rocprofiler-sdk-tool/README.md

- update env variables

* Update lib/rocprofiler-sdk/buffer_tracing.cpp

- return ROCPROFILER_STATUS_BUFFER_NOT_FOUND if buffer tracing service is configured with invalid buffer

* Update lib/rocprofiler-sdk-tool/tool.cpp

- designated hsa API trace buffer

* Update tests/hsa-tracing/CMakeLists.txt

- Fix environment

* Update rocprofv3

- do not override HSA_TOOLS_LIB
- support ROCPROF_PRELOAD
- LD_PRELOAD librocprofiler-sdk.so

* Restructure tests directory

- move all rocprofv3 integration tests into subfolder

* Update cmake/Templates/rocprofiler-sdk/config.cmake.in

- create rocprofiler-sdk::rocprofv3 cmake target

* Update tests/rocprofv3/hsa-tracing

- improve validate.py
- convert input to dict via csv.DictReader

* Update tests/apps/CMakeLists.txt

- fix build rpath for simple-transpose

* Update  cmake/rocprofiler_memcheck.cmake

- prefer libtsan.so.0

* Update tests/rocprofv3/hsa-tracing

- move to tests/rocprofv3/tracing
- include kernel tracing and memory copy tracing

* Update lib/rocprofiler-sdk-tool/tool.cpp

- normalize "_ID" vs. "_Id" in CSV column names (use "_Id")

* Update lib/rocprofiler-sdk/buffer.{hpp,cpp}

- change signature of buffer::get_buffers()
- buffer::get_buffers() uses static_object

* Update lib/rocprofiler-sdk/context/context.cpp

- update usage of buffer::get_buffers()
  - now returns pointer

* Update lib/rocprofiler-sdk/tests/buffer.cpp

- update to change for signature of buffer::get_buffers()

* Update tests/rocprofv3/tracing/CMakeLists.txt

- use %argt% with -d argument

* Update lib/rocprofiler-sdk-tool/tool.cpp

- use atexit for finalization

* Update tests/rocprofv3/tracing/CMakeLists.txt

- tweaked name of tests

* Update lib/rocprofiler-sdk/hsa/async_copy.*

- async_copy_fini + reference counting signals

* Update lib/rocprofiler-sdk/registration.cpp

- invoke hsa::async_copy_fini() to prevent data race on signals

---------

Co-authored-by: SrirakshaNag <104580803+SrirakshaNag@users.noreply.github.com>
Co-authored-by: Benjamin Welton <bewelton@amd.com>
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: SrirakshaNag <SrirakshaNag@users.noreply.github.com>
Co-authored-by: gobhardw <gopesh.bhardwaj@amd.com>
Co-authored-by: bgopesh <bgopesh@users.noreply.github.com>

[ROCm/rocprofiler-sdk commit: 1f4cf1aa39]

projects/rocprofiler-sdk/.github/workflows/docs.yml

@@ -9,7 +9,7 @@ on:
       - 'VERSION'
       - 'source/docs/**'
       - 'source/scripts/update-docs.sh'
-      - 'source/include/rocprofiler/*'
+      - 'source/include/rocprofiler-sdk/**'
       - '.github/workflows/docs.yml'
   pull_request:
     branches: [main]
@@ -18,7 +18,7 @@ on:
       - 'VERSION'
       - 'source/docs/**'
       - 'source/scripts/update-docs.sh'
-      - 'source/include/rocprofiler/*'
+      - 'source/include/rocprofiler-sdk/**'
       - '.github/workflows/docs.yml'
 
 concurrency:

projects/rocprofiler-sdk/cmake/Templates/rocprofiler-sdk/config.cmake.in

@@ -96,8 +96,7 @@ else()
         @hip_DIR@
         @rocm_version_DIR@)
 
-    include(
-        "${@PACKAGE_NAME@_CMAKE_DIR}/@PACKAGE_NAME@-config-nolink-target.cmake")
+    include("${@PACKAGE_NAME@_CMAKE_DIR}/@PACKAGE_NAME@-config-nolink-target.cmake")
 
     add_library(@PACKAGE_NAME@::@PACKAGE_NAME@-external-nolink INTERFACE IMPORTED)
 
@@ -129,11 +128,15 @@ else()
         endforeach()
     endif()
 
-    target_link_libraries(
-        @PROJECT_NAME@::@PROJECT_NAME@
-        INTERFACE @PACKAGE_NAME@::@PACKAGE_NAME@-external-nolink)
+    target_link_libraries(@PROJECT_NAME@::@PROJECT_NAME@
+                          INTERFACE @PACKAGE_NAME@::@PACKAGE_NAME@-external-nolink)
 endif()
 
+add_executable(@PACKAGE_NAME@::rocprofv3 IMPORTED)
+set_property(
+    TARGET @PACKAGE_NAME@::rocprofv3
+    PROPERTY IMPORTED_LOCATION ${@PACKAGE_NAME@_ROOT_DIR}/@CMAKE_INSTALL_BINDIR@/rocprofv3)
+
 include(FindPackageHandleStandardArgs)
 find_package_handle_standard_args(
     @PACKAGE_NAME@

projects/rocprofiler-sdk/cmake/rocprofiler_memcheck.cmake

@@ -25,7 +25,7 @@ endfunction()
 function(rocprofiler_set_memcheck_env _TYPE _LIB_BASE)
     set(_LIBS ${_LIB_BASE})
 
-    foreach(_N 6 5 4 3 2 1 0)
+    foreach(_N ${ARGN} 6 5 4 3 2 1 0)
         list(
             APPEND _LIBS
             ${CMAKE_SHARED_LIBRARY_PREFIX}${_LIB_BASE}${CMAKE_SHARED_LIBRARY_SUFFIX}.${_N}
@@ -34,7 +34,7 @@ function(rocprofiler_set_memcheck_env _TYPE _LIB_BASE)
 
     foreach(_LIB ${_LIBS})
         if(NOT ${_TYPE}_LIBRARY)
-            find_library(${_TYPE}_LIBRARY NAMES ${_LIB} ${ARGN})
+            find_library(${_TYPE}_LIBRARY NAMES ${_LIB})
         endif()
     endforeach()
 
@@ -60,7 +60,7 @@ elseif(ROCPROFILER_MEMCHECK STREQUAL "MemorySanitizer")
     rocprofiler_add_memcheck_flags("memory")
 elseif(ROCPROFILER_MEMCHECK STREQUAL "ThreadSanitizer")
     rocprofiler_add_memcheck_flags("thread")
-    rocprofiler_set_memcheck_env("${ROCPROFILER_MEMCHECK}" "tsan")
+    rocprofiler_set_memcheck_env("${ROCPROFILER_MEMCHECK}" "tsan" 0)
 elseif(ROCPROFILER_MEMCHECK STREQUAL "UndefinedBehaviorSanitizer")
     rocprofiler_add_memcheck_flags("undefined")
     rocprofiler_set_memcheck_env("${ROCPROFILER_MEMCHECK}" "ubsan")

projects/rocprofiler-sdk/source/bin/rocprofv3

@@ -1,11 +1,13 @@
 #!/bin/bash -e
 
-set -eou pipefail
+set -eo pipefail
 
-ROCPROFV2_DIR=$(dirname -- "$(realpath "${BASH_SOURCE[0]}")");
-ROCM_DIR=$( dirname -- "${ROCPROFV2_DIR}"; )
+ROCPROFV3_DIR=$(dirname -- "$(realpath "${BASH_SOURCE[0]}")")
+ROCM_DIR=$(dirname -- "${ROCPROFV3_DIR}")
 
-export HSA_TOOLS_LIB=${ROCM_DIR}/lib/librocprofiler-sdk.so.0
+: ${HSA_TOOLS_LIB:="${ROCM_DIR}/lib/librocprofiler-sdk.so"}
+
+export HSA_TOOLS_LIB
 
 # Define color codes
 GREEN='\033[0;32m'
@@ -13,61 +15,130 @@ GREY='\033[0;90m'
 RESET='\033[0m'
 
 usage() {
+  local EC=${1}
+  if [ -z "${EC}" ]; then EC=1; fi
   echo -e "${RESET}ROCProfilerV3 Run Script Usage:"
   echo -e "${GREEN}-h   | --help ${RESET}                For showing this message"
+  echo -e "${GREEN}-i   | --input ${RESET}               For adding counters file path (every line in the text file represents a counter)"
   echo -e "${GREEN}--hsa-trace ${RESET}                  For Collecting HSA API Traces"
   echo -e "${GREEN}--kernel-trace ${RESET}               For Collecting Kernel Dispatch Traces"
+  echo -e "${GREEN}--memory-copy-trace ${RESET}          For Collecting Memory Copy Traces"
   echo -e "${GREEN}-o   | --output-file ${RESET}         For the output file name"
   echo -e "\t#${GREY} usage e.g:(with current dir): rocprofv3 --hsa-trace -o <file_name> <executable>"
   echo -e "\t#${GREY} usage e.g:(with custom dir):  rocprofv3 --hsa-trace -d <out_dir> -o <file_name> <executable>${RESET}\n"
   echo -e "${GREEN}-d   | --output-directory ${RESET}    For adding output path where the output files will be saved"
-  echo -e "\t#${GREY} usage e.g:(with custom dir):  rocprofv3 --hsa-trace -d <out_dir> <executable>${RESET}\n"
-  exit 1
+  echo -e "\t#${GREY} usage e.g:(with custom dir):  rocprofv3 --hsa-trace -d <out_dir> <executable>${RESET}"
+  echo -e "${GREEN}-M   | --mangled-kernels ${RESET}     Do not demangle the kernel names"
+  echo -e "${GREEN}-T   | --truncate-kernels ${RESET}    Truncate the demangled kernel names"
+  echo -e ""
+  exit ${EC}
 }
 
 if [ -z "$1" ]; then
-  usage
+  usage 1
 fi
 
-: "${ROCPROFILER_OUTPUT_PATH:="."}"
-: "${ROCPROFILER_OUTPUT_FILE_NAME:=""}"
+: ${ROCPROF_OUTPUT_PATH:="."}
+: ${ROCPROF_OUTPUT_PATH_INTERNAL:="."}
+: ${ROCPROF_OUTPUT_FILE_NAME:=""}
+: ${ROCPROF_COUNTERS_PATH:=""}
+: ${ROCPROF_PRELOAD:=""}
 
 while true; do
   if [[ "$1" == "-h" || "$1" == "--help" ]]; then
-    usage
+    usage 0
+  elif [[ "$1" == "-M" || "$1" == "--mangled-kernels" ]]; then
+    export ROCPROF_DEMANGLE_KERNELS=0
+    shift
+  elif [[ "$1" == "-T" || "$1" == "--truncate-kernels" ]]; then
+    export ROCPROF_TRUNCATE_KERNELS=1
+    shift
+  elif [[ "$1" == "-i" || "$1" == "--input" ]]; then
+    if [ "$2" ] && [ -n "$2" ] && [ -r "$2" ]; then
+      export ROCPROF_COUNTERS_PATH=$2
+      export ROCPROF_COUNTER_COLLECTION=1
+    else
+      echo -e "Error: input file \"$2\" doesn't exist!"
+      usage 1
+    fi
+    shift
+    shift
   elif [[ "$1" == "-o" || "$1" == "--output-file-name" ]]; then
     if [ "$2" ]; then
-      export ROCPROFILER_OUTPUT_FILE_NAME=$2
+      export ROCPROF_OUTPUT_FILE_NAME=$2
     else
-      usage
+      usage 1
     fi
     shift
     shift
   elif [[ "$1" == "-d" || "$1" == "--output-directory" ]]; then
     if [ "$2" ]; then
-      ROCPROFILER_OUTPUT_PATH=$2
+      ROCPROF_OUTPUT_PATH_INTERNAL=$2
+      export ROCPROF_OUTPUT_PATH=$ROCPROF_OUTPUT_PATH_INTERNAL
     else
-      usage
+      usage 1
     fi
     shift
     shift
   elif [ "$1" == "--hsa-trace" ]; then
-    export ROCPROFILER_HSA_API_TRACE=1
+    export ROCPROF_HSA_API_TRACE=1
     shift
   elif [ "$1" == "--kernel-trace" ]; then
-    export ROCPROFILER_KERNEL_TRACE=1
+    export ROCPROF_KERNEL_TRACE=1
+    shift
+  elif [ "$1" == "--memory-copy-trace" ]; then
+    export ROCPROF_MEMORY_COPY_TRACE=1
     shift
   elif [ "$1" == "--" ]; then
     shift
     break
   elif [[ "$1" == "-"* || "$1" == "--"* ]]; then
     echo -e "Wrong option \"$1\", Please use the following options:\n"
-    usage
+    usage 1
   else
     break
   fi
 done
 
-export ROCPROFILER_OUTPUT_PATH
+# read input counter file
+PMC_LINES=()
+if [ -n "$ROCPROF_COUNTERS_PATH" ]; then
+  input=$ROCPROF_COUNTERS_PATH
+  while IFS= read -r line || [[ -n "$line" ]]; do
+    #skip empty lines
+    if [[ -z "$line" ]]; then
+      continue
+    fi
+    PMC_LINES+=("$line")
+  done <"$input"
+fi
 
-ROCP_TOOL_LIBRARIES=${ROCM_DIR}/lib/rocprofiler-sdk/librocprofiler-sdk-tool.so "${@}"
+if [ -n "${PMC_LINES:-}" ]; then
+  #for counter collection
+  COUNTER=1
+  for i in "${!PMC_LINES[@]}"; do
+    export ROCPROF_COUNTERS="${PMC_LINES[$i]}"
+    if [[ ! ${PMC_LINES[$i]} =~ "pmc" ]]; then
+      continue
+    fi
+
+    RESULT_PATH="$ROCPROF_OUTPUT_PATH_INTERNAL/pmc_$COUNTER"
+    if [ -n "$ROCPROF_OUTPUT_FILE_NAME" ] || [ -n "$ROCPROF_OUTPUT_PATH" ]; then
+      export ROCPROF_OUTPUT_PATH=$RESULT_PATH
+    fi
+    ((COUNTER++))
+    ROCP_TOOL_LIBRARIES="${ROCM_DIR}/lib/rocprofiler-sdk/librocprofiler-sdk-tool.so" \
+    LD_LIBRARY_PATH=${ROCM_DIR}/lib:${LD_LIBRARY_PATH} \
+    LD_PRELOAD="${ROCPROF_PRELOAD}:${ROCM_DIR}/lib/librocprofiler-sdk.so" \
+    "${@}"
+    if [ -n "$ROCPROF_OUTPUT_PATH" ]; then
+      echo -e "\nThe output path for the following counters: $ROCPROF_OUTPUT_PATH"
+    fi
+  done
+else
+  # for non counter collection. e.g: tracing
+  ROCP_TOOL_LIBRARIES="${ROCM_DIR}/lib/rocprofiler-sdk/librocprofiler-sdk-tool.so" \
+  LD_LIBRARY_PATH=${ROCM_DIR}/lib:${LD_LIBRARY_PATH} \
+  LD_PRELOAD="${ROCPROF_PRELOAD}:${ROCM_DIR}/lib/librocprofiler-sdk.so" \
+  "${@}"
+fi

projects/rocprofiler-sdk/source/bin/tests/CMakeLists.txt

@@ -1 +0,0 @@
-

projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h

@@ -171,7 +171,7 @@ typedef enum  // NOLINT(performance-enum-size)
  */
 typedef enum  // NOLINT(performance-enum-size)
 {
-    ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_NONE = -1,         ///< Unknown memory copy direction
+    ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_NONE = 0,          ///< Unknown memory copy direction
     ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_HOST_TO_HOST,      ///< Memory copy from host to host
     ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_HOST_TO_DEVICE,    ///< Memory copy from host to device
     ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_DEVICE_TO_HOST,    ///< Memory copy from device to host

projects/rocprofiler-sdk/source/lib/common/CMakeLists.txt

@@ -3,10 +3,9 @@
 #
 rocprofiler_activate_clang_tidy()
 
-set(common_sources config.cpp environment.cpp demangle.cpp logging.cpp static_object.cpp
-                   utility.cpp xml.cpp)
+set(common_sources environment.cpp demangle.cpp logging.cpp static_object.cpp utility.cpp
+                   xml.cpp)
 set(common_headers
-    config.hpp
     defines.hpp
     environment.hpp
     demangle.hpp

projects/rocprofiler-sdk/source/lib/common/mpl.hpp

@@ -23,6 +23,8 @@
 #pragma once
 
 #include <cstddef>
+#include <string>
+#include <string_view>
 #include <tuple>
 #include <type_traits>
 
@@ -107,6 +109,18 @@ template <typename Tp>
 struct is_pair : is_pair_impl<std::remove_cv_t<std::remove_reference_t<std::decay_t<Tp>>>>
 {};
 
+template <typename Tp>
+struct is_string_type_impl
+{
+    static constexpr auto value =
+        is_one_of<Tp, type_list<const char*, char*, std::string, std::string_view>>::value;
+};
+
+template <typename Tp>
+struct is_string_type
+: is_string_type_impl<std::remove_cv_t<std::remove_reference_t<std::decay_t<Tp>>>>
+{};
+
 template <typename, typename = void>
 constexpr bool is_type_complete_v = false;  // NOLINT(misc-definitions-in-headers)
 

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt

@@ -3,8 +3,9 @@
 #
 
 add_library(rocprofiler-sdk-tool SHARED)
-target_sources(rocprofiler-sdk-tool PRIVATE helper.hpp helper.cpp tool.cpp
-                                            trace_buffer.hpp)
+target_sources(
+    rocprofiler-sdk-tool PRIVATE config.hpp config.cpp csv.hpp helper.hpp helper.cpp
+                                 output_file.hpp output_file.cpp tool.cpp)
 
 add_subdirectory(plugins)
 

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/README.md

@@ -17,14 +17,14 @@ hsa-api trace.
 The tool uses the following environment variables
 to read the user choices.
 
-- `ROCPROFILER_KERNEL_TRACE=1` to obtain kernel trace
-- `ROCPROFILER_HSA_API_TRACE=1` to obtain hsa api trace
+- `ROCPROF_KERNEL_TRACE=1` to obtain kernel trace
+- `ROCPROF_HSA_API_TRACE=1` to obtain hsa api trace
 
 The user can also specify the output filename and output file path
 to which the traces are written to.
 
-- `ROCPROFILER_OUTPUT_PATH=<directory>` to set the output directory path
-- `ROCPROFILER_OUTPUT_FILE_NAME=<filename-without-extension>` to set the output file name
+- `ROCPROF_OUTPUT_PATH=<directory>` to set the output directory path
+- `ROCPROF_OUTPUT_FILE_NAME=<filename-without-extension>` to set the output file name
 
 ## CHANGELOG
 

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp

@@ -21,7 +21,8 @@
 // THE SOFTWARE.
 //
 
-#include "lib/common/config.hpp"
+#include "config.hpp"
+
 #include "lib/common/defines.hpp"
 #include "lib/common/demangle.hpp"
 #include "lib/common/environment.hpp"
@@ -36,23 +37,24 @@
 #include <ctime>
 #include <fstream>
 #include <regex>
+#include <set>
 #include <sstream>
 #include <string>
 #include <vector>
 
 namespace rocprofiler
 {
-namespace common
+namespace tool
 {
 namespace
 {
-std::time_t* launch_time = new std::time_t{std::time(nullptr)};
+auto launch_time = std::make_unique<std::time_t>(std::time_t{std::time(nullptr)});
 
 std::string
 get_local_datetime(const char* dt_format, std::time_t* dt_curr)
 {
     char mbstr[512];
-    if(!dt_curr) dt_curr = launch_time;
+    if(!dt_curr) dt_curr = launch_time.get();
 
     if(std::strftime(mbstr, sizeof(mbstr), dt_format, std::localtime(dt_curr)) != 0)
         return std::string{mbstr};
@@ -109,6 +111,95 @@ get_num_siblings(pid_t _id = getppid())
 {
     return get_siblings(_id).size();
 }
+
+// replace unsuported specail chars with space
+void
+handle_special_chars(std::string& str)
+{
+    // Iterate over the string and replace any special characters with a space.
+    auto pos = std::string::npos;
+    while((pos = str.find_first_of("!@#$%&(),*+-./;<=>?@{}^`~|:")) != std::string::npos)
+        str.at(pos) = ' ';
+}
+
+bool
+has_kernel_name_format(std::string const& str)
+{
+    return std::find_if(str.begin(), str.end(), [](unsigned char ch) {
+               return (isalnum(ch) != 0 || ch == '_');
+           }) != str.end();
+}
+
+bool
+has_counter_format(std::string const& str)
+{
+    return std::find_if(str.begin(), str.end(), [](unsigned char ch) {
+               return (isalnum(ch) != 0 || ch == '_');
+           }) != str.end();
+}
+
+// validate kernel names
+auto
+parse_kernel_names(const std::string& line)
+{
+    auto kernel_names_v = std::vector<std::string>{};
+    if(line.empty()) return kernel_names_v;
+
+    auto kernel_names = std::set<std::string>{};
+    trim(line);
+    auto input_line  = std::stringstream{line};
+    auto kernel_name = std::string{};
+    while(getline(input_line, kernel_name, ','))
+    {
+        if(has_kernel_name_format(kernel_name))
+        {
+            LOG(INFO) << "kernel name " << kernel_names.size() << ": " << kernel_name;
+            kernel_names.emplace(kernel_name);
+        }
+        else
+        {
+            LOG(ERROR) << "invalid kernel name: " << kernel_name;
+        }
+    }
+
+    kernel_names_v.reserve(kernel_names.size());
+    for(const auto& itr : kernel_names)
+        kernel_names_v.emplace_back(itr);
+
+    return kernel_names_v;
+}
+
+std::set<std::string>
+parse_counters(std::string line)
+{
+    auto counters = std::set<std::string>{};
+
+    if(line.empty()) return counters;
+
+    // trim line for any white spaces
+    trim(line);
+
+    if(!(line[0] == '#' || line.find("pmc") == std::string::npos))
+    {
+        handle_special_chars(line);
+
+        std::stringstream input_line(line);
+        std::string       counter;
+        while(getline(input_line, counter, ' '))
+        {
+            if(counter.substr(0, 3) != "pmc" && has_counter_format(counter))
+            {
+                counters.emplace(counter);
+            }
+            else
+            {
+                LOG(ERROR) << "invalid counter: " << counter;
+            }
+        }
+    }
+
+    return counters;
+}
 }  // namespace
 
 int
@@ -127,12 +218,17 @@ get_mpi_rank()
     return _v;
 }
 
+config::config()
+: kernel_names{parse_kernel_names(get_env("ROCPROF_KERNEL_NAMES", std::string{}))}
+, counters{parse_counters(get_env("ROCPROF_COUNTERS", std::string{}))}
+{}
+
 std::vector<output_key>
 output_keys(std::string _tag)
 {
     using strpair_t = std::pair<std::string, std::string>;
 
-    auto _cmdline = read_command_line(getpid());
+    auto _cmdline = common::read_command_line(getpid());
 
     if(_tag.empty() && !_cmdline.empty()) _tag = ::basename(_cmdline.front().c_str());
 
@@ -175,11 +271,11 @@ output_keys(std::string _tag)
         }
     }
 
-    auto* _launch_time = launch_time;
+    auto* _launch_time = launch_time.get();
     auto  _time_format = get_env<std::string>("ROCP_TIME_FORMAT", "%F_%H.%M");
 
-    auto _mpi_size = get_env<int>("OMPI_COMM_WORLD_SIZE", get_env<int>("MV2_COMM_WORLD_SIZE", 0));
-    auto _mpi_rank = get_env<int>("OMPI_COMM_WORLD_RANK", get_env<int>("MV2_COMM_WORLD_RANK", -1));
+    auto _mpi_size = get_mpi_size();
+    auto _mpi_rank = get_mpi_rank();
 
     auto _dmp_size      = fmt::format("{}", (_mpi_size) > 0 ? _mpi_size : 1);
     auto _dmp_rank      = fmt::format("{}", (_mpi_rank) > 0 ? _mpi_rank : 0);
@@ -312,69 +408,18 @@ format(std::string _fpath, const std::string& _tag)
     return _fpath;
 }
 
-std::string
-compose_filename(const config& _cfg)
-{
-    auto _output_path = _cfg.output_path;
-    auto _output_file = _cfg.output_file;
-    auto _output_ext  = _cfg.output_ext;
-
-    if(_output_path.empty()) _output_path = ".";
-    if(_cfg.mpi_size > 0)
-    {
-        if(_cfg.mpi_rank >= 0)
-        {
-            _output_file = fmt::format("{}.{}", _output_file, _cfg.mpi_rank);
-        }
-        else
-        {
-            _output_file = fmt::format("{}.{}", _output_file, getpid());
-        }
-    }
-    if(!_output_ext.empty())
-    {
-        if(_output_ext.find('.') == std::string::npos) _output_ext.insert(0, ".");
-        if(_output_file.length() < _output_ext.length() ||
-           _output_file.find(_output_ext) != _output_file.length() - _output_ext.length())
-            _output_file += _output_ext;
-    }
-
-    // join <OUTPUT_PATH>/<OUTPUT_FILE> and replace any keys with values
-    auto _prefix = format(common::filesystem::path{_output_path} / _output_file);
-
-    // return on empty
-    if(_prefix.empty()) return std::string{};
-
-    // get the absolute path
-    auto _fname = common::filesystem::absolute(common::filesystem::path{_prefix});
-
-    // create the directory if necessary
-    auto _fname_path = _fname.parent_path();
-    if(!common::filesystem::exists(_fname_path))
-        common::filesystem::create_directories(_fname.parent_path());
-
-    return _fname.string();
-}
-
 std::string
 format_name(std::string_view _name, const config& _cfg)
 {
-    if(_cfg.demangle && _cfg.truncate)
-    {
-        return truncate_name(cxx_demangle(_name));
-    }
+    if(!_cfg.demangle && !_cfg.truncate) return std::string{_name};
 
-    if(_cfg.demangle)
-    {
-        return cxx_demangle(_name);
-    }
+    // truncating requires demangling first so always demangle
+    auto _demangled_name =
+        common::cxx_demangle(std::regex_replace(_name.data(), std::regex{"(\\.kd)$"}, ""));
 
-    if(_cfg.truncate)
-    {
-        return truncate_name(_name);
-    }
+    if(_cfg.truncate) return common::truncate_name(_demangled_name);
 
-    return std::string{_name};
+    return _demangled_name;
 }
 
 void
@@ -388,5 +433,5 @@ output_key::output_key(std::string _key, std::string _val, std::string _desc)
 , value{std::move(_val)}
 , description{std::move(_desc)}
 {}
-}  // namespace common
+}  // namespace tool
 }  // namespace rocprofiler

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp

@@ -24,14 +24,19 @@
 #pragma once
 
 #include "lib/common/environment.hpp"
+#include "lib/common/filesystem.hpp"
 
+#include <set>
 #include <string>
 #include <vector>
 
 namespace rocprofiler
 {
-namespace common
+namespace tool
 {
+namespace fs = common::filesystem;
+using common::get_env;
+
 enum class config_context
 {
     global = 0,
@@ -50,13 +55,21 @@ get_mpi_rank();
 
 struct config
 {
-    bool        demangle    = get_env("ROCP_DEMANGLE_KERNELS", true);
-    bool        truncate    = get_env("ROCP_TRUNCATE_KERNELS", false);
-    int         mpi_size    = get_mpi_size();
-    int         mpi_rank    = get_mpi_rank();
-    std::string output_path = get_env<std::string>("ROCP_OUTPUT_PATH", ".");
-    std::string output_file = get_env<std::string>("ROCP_OUTPUT_FILE", "results");
-    std::string output_ext  = {};
+    config();
+
+    bool        demangle           = get_env("ROCPROF_DEMANGLE_KERNELS", true);
+    bool        truncate           = get_env("ROCPROF_TRUNCATE_KERNELS", false);
+    bool        kernel_trace       = get_env("ROCPROF_KERNEL_TRACE", false);
+    bool        hsa_api_trace      = get_env("ROCPROF_HSA_API_TRACE", false);
+    bool        marker_api_trace   = get_env("ROCPROF_MARKER_API_TRACE", false);
+    bool        memory_copy_trace  = get_env("ROCPROF_MEMORY_COPY_TRACE", false);
+    bool        counter_collection = get_env("ROCPROF_COUNTER_COLLECTION", false);
+    int         mpi_size           = get_mpi_size();
+    int         mpi_rank           = get_mpi_rank();
+    std::string output_path        = get_env("ROCPROF_OUTPUT_PATH", fs::current_path().string());
+    std::string output_file        = get_env("ROCPROF_OUTPUT_FILE_NAME", std::to_string(getpid()));
+    std::vector<std::string> kernel_names = {};
+    std::set<std::string>    counters     = {};
 };
 
 template <config_context ContextT = config_context::global>
@@ -90,9 +103,6 @@ struct output_key
 std::vector<output_key>
 output_keys(std::string _tag = {});
 
-std::string
-compose_filename(const config&);
-
 std::string
 format(std::string _fpath, const std::string& _tag = {});
 
@@ -101,5 +111,5 @@ format_name(std::string_view _name, const config& = get_config<>());
 
 void
 initialize();
-}  // namespace common
+}  // namespace tool
 }  // namespace rocprofiler

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/csv.hpp

@@ -0,0 +1,83 @@
+// MIT License
+//
+// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+//
+// 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 "lib/common/mpl.hpp"
+
+#include <array>
+#include <cstddef>
+#include <ostream>
+#include <string_view>
+
+namespace rocprofiler
+{
+namespace tool
+{
+namespace csv
+{
+template <typename TupleT, size_t... Idx>
+std::ostream&
+write_csv_entry(std::ostream& ofs, TupleT&& _data, std::index_sequence<Idx...>)
+{
+    auto _write = [&ofs](size_t idx, auto&& _val) {
+        if(idx > 0) ofs << ",";
+        if constexpr(rocprofiler::common::mpl::is_string_type<decltype(_val)>::value) ofs << "\"";
+        ofs << _val;
+        if constexpr(rocprofiler::common::mpl::is_string_type<decltype(_val)>::value) ofs << "\"";
+    };
+
+    (_write(Idx, std::get<Idx>(_data)), ...);
+    return (ofs << '\n');
+}
+
+template <size_t NumCols>
+struct csv_encoder
+{
+    static constexpr auto columns = NumCols;
+
+    template <typename... Args,
+              typename Tp                                       = void,
+              std::enable_if_t<sizeof...(Args) == columns, int> = 0>
+    static auto write_row(std::ostream& ofs, Args&&... args)
+    {
+        write_csv_entry(
+            ofs, std::make_tuple(std::forward<Args>(args)...), std::make_index_sequence<columns>{});
+        return csv_encoder<columns>{};
+    }
+
+    template <typename Tp, size_t N>
+    static auto write_row(std::ostream& ofs, const std::array<Tp, N>& arr)
+    {
+        static_assert(N == columns, "Error! too many/few args passed");
+        write_csv_entry(ofs, arr, std::make_index_sequence<columns>{});
+        return csv_encoder<columns>{};
+    }
+};
+
+using hsa_csv_encoder                = csv_encoder<7>;
+using kernel_trace_csv_encoder       = csv_encoder<16>;
+using counter_collection_csv_encoder = csv_encoder<14>;
+using memory_copy_csv_encoder        = csv_encoder<7>;
+}  // namespace csv
+}  // namespace tool
+}  // namespace rocprofiler

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.cpp

@@ -21,6 +21,8 @@
 // SOFTWARE.
 
 #include "helper.hpp"
+#include "config.hpp"
+#include "rocprofiler-sdk/fwd.h"
 
 #include <glog/logging.h>
 
@@ -28,6 +30,8 @@
 #include <iostream>
 #include <mutex>
 #include <unordered_map>
+#include <unordered_set>
+#include <utility>
 
 namespace
 {
@@ -73,7 +77,6 @@ enum amd_kernel_code_property_t
     AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_RESERVED1, 12, 4),
 };
 
-std::mutex                                             kernel_descriptor_name_map_mutex;
 std::unordered_map<rocprofiler_address_t, const char*> kernel_descriptor_name_map;
 
 std::mutex kernel_properties_correlation_mutex;
@@ -120,6 +123,7 @@ sgpr_count(const std::string_view& name, const kernel_descriptor_t& kernel_code)
     // GFX10 and later always allocate 128 sgprs.
 
     // TODO(srnagara): Recheck the extraction of gfxip from gpu name
+
     const char*  name_data       = name.data();
     const size_t gfxip_label_len = std::min(name.size() - 2, size_t{63});
     if(gfxip_label_len > 0 && strnlen(name_data, gfxip_label_len + 1) >= gfxip_label_len)
@@ -169,19 +173,12 @@ GetKernelCode(uint64_t kernel_object)
 }
 }  // namespace
 
-void
-SetKernelDescriptorName(rocprofiler_address_t kernel_descriptor, const char* kernel_name)
-{
-    std::lock_guard<std::mutex> kernel_descriptor_name_map_lock(kernel_descriptor_name_map_mutex);
-    kernel_descriptor_name_map[kernel_descriptor] = kernel_name;
-}
-
 void
 SetKernelProperties(uint64_t correlation_id, rocprofiler_tool_kernel_properties_t kernel_properties)
 {
     std::lock_guard<std::mutex> kernel_properties_correlation_map_lock(
         kernel_properties_correlation_mutex);
-    kernel_properties_correlation_map[correlation_id] = kernel_properties;
+    kernel_properties_correlation_map[correlation_id] = std::move(kernel_properties);
 }
 
 rocprofiler_tool_kernel_properties_t
@@ -197,74 +194,24 @@ GetKernelProperties(uint64_t correlation_id)
     }
     return it->second;
 }
-const char*
-GetKernelDescriptorName(rocprofiler_address_t kernel_descriptor)
-{
-    std::lock_guard<std::mutex> kernel_descriptor_name_map_lock(kernel_descriptor_name_map_mutex);
-    auto                        it = kernel_descriptor_name_map.find(kernel_descriptor);
-    if(it == kernel_descriptor_name_map.end())
-    {
-        std::cout << "kernel name not found" << std::endl;
-        abort();
-    }
-    return it->second;
-}
-
-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;
-}
 
 void
 populate_kernel_properties_data(rocprofiler_tool_kernel_properties_t* kernel_properties,
-                                const hsa_kernel_dispatch_packet_t    dispatch_packet)
+                                const hsa_kernel_dispatch_packet_t*   dispatch_packet)
 {
-    const uint64_t kernel_object = dispatch_packet.kernel_object;
+    const uint64_t kernel_object = dispatch_packet->kernel_object;
 
     const kernel_descriptor_t* kernel_code = GetKernelCode(kernel_object);
     uint64_t                   grid_size =
-        dispatch_packet.grid_size_x * dispatch_packet.grid_size_y * dispatch_packet.grid_size_z;
+        dispatch_packet->grid_size_x * dispatch_packet->grid_size_y * dispatch_packet->grid_size_z;
     if(grid_size > UINT32_MAX) abort();
     kernel_properties->grid_size = grid_size;
-    uint64_t workgroup_size = dispatch_packet.workgroup_size_x * dispatch_packet.workgroup_size_y *
-                              dispatch_packet.workgroup_size_z;
+    uint64_t workgroup_size      = dispatch_packet->workgroup_size_x *
+                              dispatch_packet->workgroup_size_y * dispatch_packet->workgroup_size_z;
     if(workgroup_size > UINT32_MAX) abort();
     kernel_properties->workgroup_size = (uint32_t) workgroup_size;
-    kernel_properties->lds_size       = dispatch_packet.group_segment_size;
-    kernel_properties->scratch_size   = dispatch_packet.private_segment_size;
+    kernel_properties->lds_size       = dispatch_packet->group_segment_size;
+    kernel_properties->scratch_size   = dispatch_packet->private_segment_size;
     kernel_properties->arch_vgpr_count =
         arch_vgpr_count(kernel_properties->gpu_agent.name, *kernel_code);
     kernel_properties->accum_vgpr_count =
@@ -275,79 +222,62 @@ populate_kernel_properties_data(rocprofiler_tool_kernel_properties_t* kernel_pro
                          AMD_KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32)
             ? 32
             : 64;
-    kernel_properties->signal_handle = dispatch_packet.completion_signal.handle;
+    kernel_properties->signal_handle = dispatch_packet->completion_signal.handle;
 }
 
-std::string
-cxa_demangle(std::string_view _mangled_name, int* _status)
+rocprofiler_tool_buffer_name_info_t
+get_buffer_id_names()
 {
-    constexpr size_t buffer_len = 4096;
-    // return the mangled since there is no buffer
-    if(_mangled_name.empty())
-    {
-        *_status = -2;
-        return std::string{};
-    }
+    static auto supported = std::unordered_set<rocprofiler_buffer_tracing_kind_t>{
+        ROCPROFILER_BUFFER_TRACING_HSA_API,
+        ROCPROFILER_BUFFER_TRACING_MEMORY_COPY,
+        ROCPROFILER_BUFFER_TRACING_MARKER_API};
 
-    auto _demangled_name = std::string{_mangled_name};
+    auto cb_name_info = rocprofiler_tool_buffer_name_info_t{};
+    //
+    // callback for each kind operation
+    //
+    static auto tracing_kind_operation_cb =
+        [](rocprofiler_buffer_tracing_kind_t kindv, uint32_t operation, void* data_v) {
+            auto* name_info_v = static_cast<rocprofiler_tool_buffer_name_info_t*>(data_v);
 
-    // PARAMETERS to __cxa_demangle
-    //  mangled_name:
-    //      A nullptr-terminated character string containing the name to be demangled.
-    //  buffer:
-    //      A region of memory, allocated with malloc, of *length bytes, into which the
-    //      demangled name is stored. If output_buffer is not long enough, it is expanded
-    //      using realloc. output_buffer may instead be nullptr; in that case, the demangled
-    //      name is placed in a region of memory allocated with malloc.
-    //  _buflen:
-    //      If length is non-nullptr, the length of the buffer containing the demangled name
-    //      is placed in *length.
-    //  status:
-    //      *status is set to one of the following values
-    size_t _demang_len = 0;
-    char*  _demang = abi::__cxa_demangle(_demangled_name.c_str(), nullptr, &_demang_len, _status);
-    switch(*_status)
-    {
-        //  0 : The demangling operation succeeded.
-        // -1 : A memory allocation failure occurred.
-        // -2 : mangled_name is not a valid name under the C++ ABI mangling rules.
-        // -3 : One of the arguments is invalid.
-        case 0:
+            if(supported.count(kindv) > 0)
+            {
+                const char* name = nullptr;
+                ROCPROFILER_CALL(rocprofiler_query_buffer_tracing_kind_operation_name(
+                                     kindv, operation, &name, nullptr),
+                                 "query buffer failed");
+                if(name) name_info_v->operation_names[kindv][operation] = name;
+            }
+
+            return 0;
+        };
+
+    //
+    //  callback for each kind (i.e. domain)
+    //
+    static auto tracing_kind_cb = [](rocprofiler_buffer_tracing_kind_t kind, void* data) {
+        //  store the buffer kind name
+        auto*       name_info_v = static_cast<rocprofiler_tool_buffer_name_info_t*>(data);
+        const char* name        = nullptr;
+        ROCPROFILER_CALL(rocprofiler_query_buffer_tracing_kind_name(kind, &name, nullptr),
+                         "query buffer failed");
+
+        if(name) name_info_v->kind_names[kind] = name;
+
+        if(supported.count(kind) > 0)
         {
-            if(_demang) _demangled_name = std::string{_demang};
-            break;
+            ROCPROFILER_CALL(rocprofiler_iterate_buffer_tracing_kind_operations(
+                                 kind, tracing_kind_operation_cb, static_cast<void*>(data)),
+                             "query buffer failed");
         }
-        case -1:
-        {
-            char _msg[buffer_len];
-            ::memset(_msg, '\0', buffer_len * sizeof(char));
-            ::snprintf(_msg,
-                       buffer_len,
-                       "memory allocation failure occurred demangling %s",
-                       _demangled_name.c_str());
-            ::perror(_msg);
-            break;
-        }
-        case -2: break;
-        case -3:
-        {
-            char _msg[buffer_len];
-            ::memset(_msg, '\0', buffer_len * sizeof(char));
-            ::snprintf(_msg,
-                       buffer_len,
-                       "Invalid argument in: (\"%s\", nullptr, nullptr, %p)",
-                       _demangled_name.c_str(),
-                       (void*) _status);
-            ::perror(_msg);
-            break;
-        }
-        default: break;
+
+        return 0;
     };
 
-    // if it "demangled" but the length is zero, set the status to -2
-    if(_demang_len == 0 && *_status == 0) *_status = -2;
+    ROCPROFILER_CALL(rocprofiler_iterate_buffer_tracing_kinds(tracing_kind_cb,
+                                                              static_cast<void*>(&cb_name_info)),
+                     "iterate_buffer failed");
 
-    // free allocated buffer
-    ::free(_demang);
-    return _demangled_name;
+    return cb_name_info;
 }

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.hpp

@@ -20,6 +20,22 @@
 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 // SOFTWARE.
 
+#pragma once
+
+#include "lib/common/defines.hpp"
+#include "lib/common/filesystem.hpp"
+
+#include <rocprofiler-sdk/registration.h>
+#include <rocprofiler-sdk/rocprofiler.h>
+
+#include <amd_comgr/amd_comgr.h>
+#include <hsa/amd_hsa_kernel_code.h>
+#include <hsa/hsa.h>
+#include <hsa/hsa_api_trace.h>
+#include <hsa/hsa_ext_amd.h>
+#include <hsa/hsa_ven_amd_aqlprofile.h>
+#include <hsa/hsa_ven_amd_loader.h>
+
 #include <cxxabi.h>
 #include <sys/syscall.h>
 #include <sys/types.h>
@@ -29,32 +45,22 @@
 #include <map>
 #include <ostream>
 #include <regex>
+#include <set>
 #include <sstream>
 #include <string>
 #include <string_view>
 #include <vector>
 
-#include "lib/common/filesystem.hpp"
-
-#include <amd_comgr/amd_comgr.h>
-#include <hsa/amd_hsa_kernel_code.h>
-#include <hsa/hsa.h>
-#include <hsa/hsa_api_trace.h>
-#include <hsa/hsa_ext_amd.h>
-#include <hsa/hsa_ven_amd_aqlprofile.h>
-#include <hsa/hsa_ven_amd_loader.h>
-#include <rocprofiler-sdk/registration.h>
-#include <rocprofiler-sdk/rocprofiler.h>
-
 #define ROCPROFILER_CALL(result, msg)                                                              \
     {                                                                                              \
-        rocprofiler_status_t CHECKSTATUS = result;                                                 \
-        if(CHECKSTATUS != ROCPROFILER_STATUS_SUCCESS)                                              \
+        rocprofiler_status_t ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) = result;                 \
+        if(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) != ROCPROFILER_STATUS_SUCCESS)              \
         {                                                                                          \
-            std::string status_msg = rocprofiler_get_status_string(CHECKSTATUS);                   \
+            std::string status_msg =                                                               \
+                rocprofiler_get_status_string(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__));        \
             std::cerr << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg            \
-                      << " failed with error code " << CHECKSTATUS << ": " << status_msg           \
-                      << std::endl;                                                                \
+                      << " failed with error code " << ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) \
+                      << ": " << status_msg << std::endl;                                          \
             std::stringstream errmsg{};                                                            \
             errmsg << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg " failure ("  \
                    << status_msg << ")";                                                           \
@@ -86,16 +92,12 @@ typedef struct
     uint64_t               signal_handle;
     uint64_t               kernel_object;
     rocprofiler_queue_id_t queue_id;
+    std::string            kernel_name;
     rocprofiler_agent_t    gpu_agent;
+    uint64_t               thread_id;
 
 } rocprofiler_tool_kernel_properties_t;
 
-typedef struct
-{
-    std::vector<rocprofiler_agent_t> gpu_agents_lists;
-
-} rocprofiler_tool_agent_callback_t;
-
 struct kernel_descriptor_t
 {
     uint8_t  reserved0[16];
@@ -108,19 +110,20 @@ struct kernel_descriptor_t
     uint8_t  reserved2[6];
 };
 
-using rocprofiler_tool_callback_kind_names_t =
-    std::map<rocprofiler_callback_tracing_kind_t, const char*>;
-using rocprofiler_tool_callback_kind_operation_names_t =
-    std::map<rocprofiler_callback_tracing_kind_t, std::map<uint32_t, const char*>>;
+using rocprofiler_tool_buffer_kind_names_t =
+    std::unordered_map<rocprofiler_buffer_tracing_kind_t, const char*>;
+using rocprofiler_tool_buffer_kind_operation_names_t =
+    std::unordered_map<rocprofiler_buffer_tracing_kind_t,
+                       std::unordered_map<uint32_t, const char*>>;
 
-struct rocprofiler_tool_callback_name_info_t
+struct rocprofiler_tool_buffer_name_info_t
 {
-    rocprofiler_tool_callback_kind_names_t           kind_names      = {};
-    rocprofiler_tool_callback_kind_operation_names_t operation_names = {};
+    rocprofiler_tool_buffer_kind_names_t           kind_names      = {};
+    rocprofiler_tool_buffer_kind_operation_names_t operation_names = {};
 };
 
-std::vector<std::string>
-GetCounterNames();
+// std::vector<std::string>
+// GetCounterNames();
 
 void
 SetKernelDescriptorName(rocprofiler_address_t kernel_descriptor, const char* name);
@@ -140,10 +143,7 @@ GetKernelDescriptorName(rocprofiler_address_t kernel_descriptor);
 
 void
 populate_kernel_properties_data(rocprofiler_tool_kernel_properties_t* kernel_properties,
-                                const hsa_kernel_dispatch_packet_t    dispatch_packet);
+                                const hsa_kernel_dispatch_packet_t*   dispatch_packet);
 
-void
-TracerFlushRecord(void* data, rocprofiler_callback_tracing_kind_t kind);
-
-std::string
-cxa_demangle(std::string_view _mangled_name, int* _status);
+rocprofiler_tool_buffer_name_info_t
+get_buffer_id_names();

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp

@@ -0,0 +1,62 @@
+// MIT License
+//
+// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+//
+// 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 "output_file.hpp"
+#include "config.hpp"
+
+#include <fmt/format.h>
+
+namespace rocprofiler
+{
+namespace tool
+{
+namespace fs = common::filesystem;
+
+std::pair<std::ostream*, void (*)(std::ostream*&)>
+get_output_stream(const std::string& fname, const std::string& ext)
+{
+    auto output_path      = fs::path{tool::format(tool::get_config().output_path)};
+    auto output_file_name = tool::format(tool::get_config().output_file);
+
+    if(output_path.string().empty()) return {&std::clog, [](auto*&) {}};
+
+    if(fs::exists(output_path) && !fs::is_directory(fs::status(output_path)))
+        throw std::runtime_error{
+            fmt::format("ROCPROFILER_OUTPUT_PATH ({}) already exists and is not a directory",
+                        output_path.string())};
+    if(!fs::exists(output_path)) fs::create_directories(output_path);
+
+    auto  output_file = tool::format(output_path / (output_file_name + "_" + fname + ext));
+    auto* _ofs        = new std::ofstream{output_file};
+    if(!_ofs && !*_ofs)
+        throw std::runtime_error{fmt::format("Failed to open {} for output", output_file)};
+
+    LOG(ERROR) << "Results File: " << output_file;
+
+    return {_ofs, [](std::ostream*& v) {
+                if(v) dynamic_cast<std::ofstream*>(v)->close();
+                delete v;
+                v = nullptr;
+            }};
+}
+}  // namespace tool
+}  // namespace rocprofiler

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp

@@ -0,0 +1,101 @@
+// MIT License
+//
+// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+//
+// 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 "config.hpp"
+#include "csv.hpp"
+
+#include "lib/common/filesystem.hpp"
+#include "lib/rocprofiler-sdk-tool/csv.hpp"
+
+#include <array>
+#include <iostream>
+#include <mutex>
+#include <ostream>
+#include <sstream>
+#include <stdexcept>
+#include <string>
+#include <vector>
+
+namespace rocprofiler
+{
+namespace tool
+{
+std::pair<std::ostream*, void (*)(std::ostream*&)>
+get_output_stream(const std::string& fname, const std::string& ext = ".csv");
+
+struct output_file
+{
+    template <size_t N>
+    output_file(std::string name, csv::csv_encoder<N>, std::array<std::string_view, N>&& header);
+
+    ~output_file() { m_dtor(m_stream); }
+
+    output_file(const output_file&) = delete;
+    output_file& operator=(const output_file&) = delete;
+
+    std::string name() const { return m_name; }
+
+    template <typename T>
+    std::ostream& operator<<(T&& value)
+    {
+        auto _lk = std::unique_lock<std::mutex>{m_mutex};
+        return ((m_stream) ? *m_stream : std::cerr) << std::forward<T>(value);
+    }
+
+    std::ostream& operator<<(std::ostream& (*func)(std::ostream&) )
+    {
+        auto _lk = std::unique_lock<std::mutex>{m_mutex};
+        return ((m_stream) ? *m_stream : std::cerr) << func;
+    }
+
+    operator bool() const { return m_stream != nullptr; }
+
+private:
+    using stream_dtor_t = void (*)(std::ostream*&);
+
+    const std::string m_name   = {};
+    std::mutex        m_mutex  = {};
+    std::ostream*     m_stream = nullptr;
+    stream_dtor_t     m_dtor   = [](std::ostream*&) {};
+};
+
+template <size_t N>
+output_file::output_file(std::string                       name,
+                         csv::csv_encoder<N>               encoder,
+                         std::array<std::string_view, N>&& header)
+: m_name{std::move(name)}
+{
+    std::tie(m_stream, m_dtor) = get_output_stream(m_name);
+
+    for(auto& itr : header)
+    {
+        LOG_IF(FATAL, itr.empty())
+            << "CSV file for " << m_name << " was not provided the correct number of headers";
+    }
+
+    // write the csv header
+    if(m_stream) encoder.write_row(*m_stream, header);
+}
+}  // namespace tool
+}  // namespace rocprofiler

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp

@@ -20,192 +20,186 @@
 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 // SOFTWARE.
 
+#include "config.hpp"
+#include "csv.hpp"
 #include "helper.hpp"
-#include "rocprofiler-sdk/context.h"
-#include "trace_buffer.hpp"
+#include "output_file.hpp"
 
+#include "lib/common/demangle.hpp"
 #include "lib/common/environment.hpp"
 #include "lib/common/filesystem.hpp"
+#include "lib/common/logging.hpp"
+#include "lib/common/synchronized.hpp"
+#include "lib/common/utility.hpp"
+
+#include <rocprofiler-sdk/agent.h>
+#include <rocprofiler-sdk/fwd.h>
+#include <rocprofiler-sdk/rocprofiler.h>
+
+#include <glog/logging.h>
 
 #include <fmt/core.h>
 #include <unistd.h>
-
+#include <cassert>
 #include <fstream>
 #include <iomanip>
 #include <mutex>
+#include <optional>
 #include <shared_mutex>
 #include <unordered_map>
+#include <unordered_set>
 #include <vector>
 
-namespace common = ::rocprofiler::common;
-namespace fs     = common::filesystem;
-
-TRACE_BUFFER_INSTANTIATE();
-
-// static const uint32_t lds_block_size = 128 * 4;
+namespace common                     = ::rocprofiler::common;
+namespace tool                       = ::rocprofiler::tool;
+static const uint32_t lds_block_size = 128 * 4;
 
 namespace
-{
-auto tool_buffer = rocprofiler_buffer_id_t{};
-auto context_id  = rocprofiler_context_id_t{};
-auto output_path =
-    fs::path{common::get_env<std::string>("ROCPROFILER_OUTPUT_PATH", fs::current_path().string())};
-auto output_file_name =
-    common::get_env<std::string>("ROCPROFILER_OUTPUT_FILE_NAME", std::to_string(getpid()));
-
-std::pair<std::ostream*, void (*)(std::ostream*&)>
-get_output_stream(const std::string& fname, const std::string& ext = ".csv")
-{
-    if(output_path.string().empty()) return {&std::clog, [](auto*&) {}};
-
-    if(fs::exists(output_path) && !fs::is_directory(fs::status(output_path)))
-        throw std::runtime_error{
-            fmt::format("ROCPROFILER_OUTPUT_PATH ({}) already exists and is not a directory",
-                        output_path.string())};
-    if(!fs::exists(output_path)) fs::create_directories(output_path);
-
-    auto  output_file = output_path / (output_file_name + "_" + fname + ext);
-    auto* _ofs        = new std::ofstream{output_file};
-    if(!_ofs && !*_ofs)
-        throw std::runtime_error{fmt::format("Failed to open {} for output", output_file.string())};
-    std::cout << "Results File: " << output_file << std::endl;
-    return {_ofs, [](std::ostream*& v) {
-                if(v) dynamic_cast<std::ofstream*>(v)->close();
-                delete v;
-                v = nullptr;
-            }};
-}
-
-template <typename Tp>
-std::string
-as_hex(Tp _v, size_t _width = 16)
-{
-    auto _ss = std::stringstream{};
-    _ss.fill('0');
-    _ss << "0x" << std::hex << std::setw(_width) << _v;
-    return _ss.str();
-}
-}  // namespace
-
-struct output_file
-{
-    output_file(std::string name, std::vector<std::string>&& header)
-    : m_name{std::move(name)}
-    {
-        std::tie(m_stream, m_dtor) = get_output_stream(m_name);
-        auto ss                    = std::stringstream{};
-        for(auto&& itr : header)
-        {
-            ss << "," << itr;
-        }
-
-        // write the csv header
-        if(!ss.str().empty()) *m_stream << ss.str().substr(1) << '\n';
-    }
-
-    ~output_file() { m_dtor(m_stream); }
-
-    output_file(const output_file&) = delete;
-    output_file& operator=(const output_file&) = delete;
-
-    std::string name() const { return m_name; }
-
-    template <typename T>
-    std::ostream& operator<<(T&& value)
-    {
-        return (*m_stream) << std::forward<T>(value);
-    }
-
-    std::ostream& operator<<(std::ostream& (*func)(std::ostream&) ) { return (*m_stream) << func; }
-
-    operator bool() const { return m_stream != nullptr; }
-
-private:
-    using stream_dtor_t = void (*)(std::ostream*&);
-
-    const std::string m_name   = {};
-    std::ostream*     m_stream = nullptr;
-    stream_dtor_t     m_dtor   = [](std::ostream*&) {};
-};
+{}  // namespace
 
 auto&
 get_hsa_api_file()
 {
-    static auto _v =
-        output_file{"hsa_api_trace",
-                    {"DOMAIN", "FUNCTION", "START_TIMESTAMP", "END_TIMESTAMP", "CORRELATION_ID"}};
+    static auto _v = tool::output_file{"hsa_api_trace",
+                                       tool::csv::hsa_csv_encoder{},
+                                       {"Domain",
+                                        "Function",
+                                        "Process_Id",
+                                        "Thread_Id",
+                                        "Correlation_Id",
+                                        "Start_Timestamp",
+                                        "End_Timestamp"}};
     return _v;
 }
 
 auto&
 get_kernel_trace_file()
 {
-    static auto _v = output_file{"kernel_trace",
-                                 {"AGENT_ID",
-                                  "QUEUE_ID",
-                                  "KERNEL_ID",
-                                  "KERNEL_NAME",
-                                  "CONTEXT_ID",
-                                  "BUFFER_ID",
-                                  "CORRELATION_ID",
-                                  "KIND",
-                                  "START_TS",
-                                  "END_TS",
-                                  "PRIVATE_SEGMENT_SIZE",
-                                  "GROUP_SEGMENT_SIZE",
-                                  "WORKGROUP_SIZE_X",
-                                  "WORKGROUP_SIZE_Y",
-                                  "WORKGROUP_SIZE_Z",
-                                  "GRID_SIZE_X",
-                                  "GRID_SIZE_Y",
-                                  "GRID_SIZE_Z"}};
+    static auto _v = tool::output_file{"kernel_trace",
+                                       tool::csv::kernel_trace_csv_encoder{},
+                                       {"Kind",
+                                        "Agent_Id",
+                                        "Queue_Id",
+                                        "Kernel_Id",
+                                        "Kernel_Name",
+                                        "Correlation_Id",
+                                        "Start_Timestamp",
+                                        "End_Timestamp",
+                                        "Private_Segment_Size",
+                                        "Group_Segment_Size",
+                                        "Workgroup_Size_X",
+                                        "Workgroup_Size_Y",
+                                        "Workgroup_Size_Z",
+                                        "Grid_Size_X",
+                                        "Grid_Size_Y",
+                                        "Grid_Size_Z"}};
     return _v;
 }
 
-std::shared_mutex                                        kernel_data_mutex;
-std::unordered_map<rocprofiler_kernel_id_t, std::string> kernel_data;
-
-struct hsa_api_trace_entry_t
+auto&
+get_counter_collection_file()
 {
-    std::atomic<uint32_t>                 valid;
-    rocprofiler_callback_tracing_record_t record;
-    rocprofiler_timestamp_t               begin_timestamp;
-    rocprofiler_timestamp_t               end_timestamp;
-    std::string_view                      api_name;
+    static auto _v = tool::output_file{"counter_collection",
+                                       tool::csv::counter_collection_csv_encoder{},
+                                       {"Counter_Id",
+                                        "Agent_Id",
+                                        "Queue_Id",
+                                        "Process_Id",
+                                        "Thread_Id",
+                                        "Grid_Size",
+                                        "Kernel-Name",
+                                        "Workgroup_Size",
+                                        "LDS_Block_Size",
+                                        "Scratch_Size",
+                                        "VGPR_Count",
+                                        "SGPR_Count",
+                                        "Counter_Name",
+                                        "Counter_Value"}};
+    return _v;
+}
 
-    hsa_api_trace_entry_t(rocprofiler_timestamp_t               begin,
-                          rocprofiler_timestamp_t               end,
-                          rocprofiler_callback_tracing_record_t tracer_record,
-                          std::string_view                      name)
-    : valid(TRACE_ENTRY_INIT)
-    , record(tracer_record)
-    , begin_timestamp(begin)
-    , end_timestamp(end)
-    , api_name(name)
+auto&
+get_memory_copy_trace_file()
+{
+    static auto _v = tool::output_file{"memory_copy_trace",
+                                       tool::csv::memory_copy_csv_encoder{},
+                                       {"Kind",
+                                        "Direction",
+                                        "Source_Agent_Id",
+                                        "Destination_Agent_Id",
+                                        "Correlation_Id",
+                                        "Start_Timestamp",
+                                        "End_Timestamp"}};
+    return _v;
+}
+
+rocprofiler_buffer_id_t&
+get_hsa_api_trace_buffer()
+{
+    static rocprofiler_buffer_id_t hsa_api_buf = {};
+    return hsa_api_buf;
+}
+
+rocprofiler_buffer_id_t&
+get_kernel_trace_buffer()
+{
+    static rocprofiler_buffer_id_t kernel_trace_buf = {};
+    return kernel_trace_buf;
+}
+
+rocprofiler_buffer_id_t&
+get_counter_collection_buffer()
+{
+    static rocprofiler_buffer_id_t counter_collection_buf = {};
+    return counter_collection_buf;
+}
+
+rocprofiler_buffer_id_t&
+get_memory_copy_trace_buffer()
+{
+    static rocprofiler_buffer_id_t memory_copy_buf = {};
+    return memory_copy_buf;
+}
+
+using rocprofiler_kernel_symbol_data_t =
+    rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t;
+
+struct kernel_symbol_data : rocprofiler_kernel_symbol_data_t
+{
+    using base_type = rocprofiler_kernel_symbol_data_t;
+
+    kernel_symbol_data(const base_type& _base)
+    : base_type{_base}
+    , formatted_kernel_name{tool::format_name(CHECK_NOTNULL(_base.kernel_name))}
+    , demangled_kernel_name{common::cxx_demangle(CHECK_NOTNULL(_base.kernel_name))}
+    , truncated_kernel_name{common::truncate_name(demangled_kernel_name)}
     {}
+
+    std::string formatted_kernel_name = {};
+    std::string demangled_kernel_name = {};
+    std::string truncated_kernel_name = {};
 };
 
-TraceBuffer<hsa_api_trace_entry_t> hsa_api_buffer("HSA API",
-                                                  0x200000,
-                                                  [](hsa_api_trace_entry_t* entry) {
-                                                      TracerFlushRecord(
-                                                          entry,
-                                                          ROCPROFILER_CALLBACK_TRACING_HSA_API);
-                                                  });
+using kernel_symbol_data_map_t = std::unordered_map<rocprofiler_kernel_id_t, kernel_symbol_data>;
+auto kernel_data               = common::Synchronized<kernel_symbol_data_map_t, true>{};
+auto name_info                 = get_buffer_id_names();
 
-rocprofiler_tool_callback_name_info_t name_info;
+auto&
+get_client_ctx()
+{
+    static rocprofiler_context_id_t context_id;
+    return context_id;
+}
 
 void
-TracerFlushRecord(void* data, rocprofiler_callback_tracing_kind_t kind)
+flush()
 {
-    if(kind == ROCPROFILER_CALLBACK_TRACING_HSA_API)
+    for(auto itr : {get_memory_copy_trace_buffer(),
+                    get_kernel_trace_buffer(),
+                    get_counter_collection_buffer(),
+                    get_hsa_api_trace_buffer()})
     {
-        auto* entry = reinterpret_cast<hsa_api_trace_entry_t*>(data);
-        get_hsa_api_file() << "ACTIVITY_DOMAIN_HSA_API"
-                           << ","
-                           << "\"" << entry->api_name << "\""
-                           << "," << entry->begin_timestamp << "," << entry->end_timestamp << ","
-                           << entry->record.correlation_id.internal << '\n';
+        if(itr.handle > 0) ROCPROFILER_CALL(rocprofiler_flush_buffer(itr), "buffer flush");
     }
 }
 
@@ -214,35 +208,11 @@ rocprofiler_tracing_callback(rocprofiler_callback_tracing_record_t record,
                              rocprofiler_user_data_t*              user_data,
                              void*                                 data)
 {
-    if(record.kind == ROCPROFILER_CALLBACK_TRACING_HIP_API)
-    {
-        // To be implemented
-        throw std::runtime_error{"not implemented"};
-    }
+    throw std::runtime_error{"not implemented"};
 
-    else if(record.kind == ROCPROFILER_CALLBACK_TRACING_HSA_API)
-    {
-        auto timestamp = rocprofiler_timestamp_t{};
-        ROCPROFILER_CALL(rocprofiler_get_timestamp(&timestamp), "timestamp failed");
-
-        if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER)
-        {
-            user_data->value = timestamp;
-        }
-        else if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT)
-        {
-            const auto* info_name_str = name_info.operation_names[record.kind][record.operation];
-            hsa_api_trace_entry_t& entry =
-                hsa_api_buffer.Emplace(user_data->value, timestamp, record, info_name_str);
-            entry.valid.store(TRACE_ENTRY_COMPLETE, std::memory_order_release);
-        }
-    }
-    else if(record.kind == ROCPROFILER_CALLBACK_TRACING_MARKER_API)
-    {
-        // To be implemented
-        throw std::runtime_error{"not implemented"};
-    }
-    (void) (data);
+    (void) record;
+    (void) user_data;
+    (void) data;
 }
 
 void
@@ -255,38 +225,22 @@ code_object_tracing_callback(rocprofiler_callback_tracing_record_t record,
     {
         if(record.phase == ROCPROFILER_CALLBACK_PHASE_UNLOAD)
         {
-            // flush the buffer to ensure that any lookups for the client kernel names for the code
-            // object are completed
-            auto flush_status = rocprofiler_flush_buffer(tool_buffer);
-            if(flush_status != ROCPROFILER_STATUS_ERROR_BUFFER_BUSY)
-                ROCPROFILER_CALL(flush_status, "buffer flush");
+            flush();
         }
     }
 
     if(record.kind == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT &&
        record.operation == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER)
     {
-        auto* sym_data =
-            static_cast<rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t*>(
-                record.payload);
-
+        auto* sym_data = static_cast<rocprofiler_kernel_symbol_data_t*>(record.payload);
         if(record.phase == ROCPROFILER_CALLBACK_PHASE_LOAD)
         {
-            auto kernel_name =
-                std::regex_replace(sym_data->kernel_name, std::regex{"(\\.kd)$"}, "");
-            int demangle_status = 0;
-            kernel_name         = cxa_demangle(kernel_name, &demangle_status);
-            std::unique_lock<std::shared_mutex> lock(kernel_data_mutex);
-            kernel_data.emplace(sym_data->kernel_id, kernel_name);
+            kernel_data.wlock(
+                [](kernel_symbol_data_map_t& kdata, rocprofiler_kernel_symbol_data_t* sym_data_v) {
+                    kdata.emplace(sym_data_v->kernel_id, kernel_symbol_data{*sym_data_v});
+                },
+                sym_data);
         }
-        // The map entry cannot be erased here
-        // since we are tracing the kernel symbols here not the kernel dispatch
-        // else if(record.phase == ROCPROFILER_CALLBACK_PHASE_UNLOAD)
-        //{
-
-        //  kernel_data.erase(data->kernel_id);
-
-        //}
     }
 
     (void) user_data;
@@ -294,174 +248,369 @@ code_object_tracing_callback(rocprofiler_callback_tracing_record_t record,
 }
 
 void
-kernel_tracing_callback(rocprofiler_context_id_t      context,
-                        rocprofiler_buffer_id_t       buffer_id,
-                        rocprofiler_record_header_t** headers,
-                        size_t                        num_headers,
-                        void*                         user_data,
-                        uint64_t /*drop_count*/)
+buffered_callback(rocprofiler_context_id_t /*context*/,
+                  rocprofiler_buffer_id_t /*buffer_id*/,
+                  rocprofiler_record_header_t** headers,
+                  size_t                        num_headers,
+                  void* /*user_data*/,
+                  uint64_t /*drop_count*/)
 {
+    static auto _sync = std::mutex{};
+    auto        _lk   = std::lock_guard<std::mutex>{_sync};
+
     if(num_headers == 0)
-        throw std::runtime_error{
-            "rocprofiler invoked a buffer callback with no headers. this should never happen"};
+        throw std::runtime_error{"rocprofiler invoked a buffer callback with no headers "
+                                 "this should never happen"};
+
     else if(headers == nullptr)
         throw std::runtime_error{"rocprofiler invoked a buffer callback with a null pointer to the "
                                  "array of headers. this should never happen"};
 
-    auto kernel_trace_ss = std::stringstream{};
     for(size_t i = 0; i < num_headers; ++i)
     {
         auto* header = headers[i];
 
-        if(header == nullptr)
+        if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING)
         {
-            throw std::runtime_error{
-                "rocprofiler provided a null pointer to header. this should never happen"};
-        }
-        else if(header->hash !=
-                rocprofiler_record_header_compute_hash(header->category, header->kind))
-        {
-            throw std::runtime_error{"rocprofiler_record_header_t (category | kind) != hash"};
-        }
-        else if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
-                header->kind == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH)
-        {
-            auto* record =
-                static_cast<rocprofiler_buffer_tracing_kernel_dispatch_record_t*>(header->payload);
-            std::string kernel_name;
+            if(header->kind == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH)
             {
-                std::shared_lock<std::shared_mutex> lock(kernel_data_mutex);
-                kernel_name = kernel_data.at(record->kernel_id);
-            }
+                auto* record = static_cast<rocprofiler_buffer_tracing_kernel_dispatch_record_t*>(
+                    header->payload);
+                std::string kernel_name = kernel_data.rlock(
+                    [](const kernel_symbol_data_map_t&                      kdata,
+                       rocprofiler_buffer_tracing_kernel_dispatch_record_t* record_v) {
+                        return kdata.at(record_v->kernel_id).formatted_kernel_name;
+                    },
+                    record);
 
-            kernel_trace_ss << record->agent_id.handle << "," << record->queue_id.handle << ","
-                            << record->kernel_id << ",\"" << kernel_name << "\"," << context.handle
-                            << "," << buffer_id.handle << "," << record->correlation_id.internal
-                            << "," << record->kind << "," << record->start_timestamp << ","
-                            << record->end_timestamp << "," << record->private_segment_size << ","
-                            << record->group_segment_size << "," << record->workgroup_size.x << ","
-                            << record->workgroup_size.y << "," << record->workgroup_size.z << ","
-                            << record->grid_size.x << "," << record->grid_size.y << ","
-                            << record->grid_size.z << '\n';
+                auto kernel_trace_ss = std::stringstream{};
+                tool::csv::kernel_trace_csv_encoder::write_row(
+                    kernel_trace_ss,
+                    name_info.kind_names.at(record->kind),
+                    record->agent_id.handle,
+                    record->queue_id.handle,
+                    record->kernel_id,
+                    std::move(kernel_name),
+                    record->correlation_id.internal,
+                    record->start_timestamp,
+                    record->end_timestamp,
+                    record->private_segment_size,
+                    record->group_segment_size,
+                    record->workgroup_size.x,
+                    record->workgroup_size.y,
+                    record->workgroup_size.z,
+                    record->grid_size.x,
+                    record->grid_size.y,
+                    record->grid_size.z);
+
+                get_kernel_trace_file() << kernel_trace_ss.str();
+            }
+            else if(header->kind == ROCPROFILER_BUFFER_TRACING_HSA_API)
+            {
+                auto* record =
+                    static_cast<rocprofiler_buffer_tracing_hsa_api_record_t*>(header->payload);
+
+                auto hsa_trace_ss = std::stringstream{};
+                tool::csv::hsa_csv_encoder::write_row(
+                    hsa_trace_ss,
+                    name_info.kind_names.at(record->kind),
+                    name_info.operation_names.at(record->kind).at(record->operation),
+                    getpid(),
+                    record->thread_id,
+                    record->correlation_id.internal,
+                    record->start_timestamp,
+                    record->end_timestamp);
+
+                get_hsa_api_file() << hsa_trace_ss.str();
+            }
+            else if(header->kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY)
+            {
+                auto* record =
+                    static_cast<rocprofiler_buffer_tracing_memory_copy_record_t*>(header->payload);
+
+                auto memory_copy_trace_ss = std::stringstream{};
+                tool::csv::memory_copy_csv_encoder::write_row(
+                    memory_copy_trace_ss,
+                    name_info.kind_names.at(record->kind),
+                    name_info.operation_names.at(record->kind).at(record->operation),
+                    record->src_agent_id.handle,
+                    record->dst_agent_id.handle,
+                    record->correlation_id.internal,
+                    record->start_timestamp,
+                    record->end_timestamp);
+
+                get_memory_copy_trace_file() << memory_copy_trace_ss.str();
+            }
+            else
+            {
+                LOG(FATAL) << fmt::format(
+                    "unsupported category + kind: {} + {}", header->category, header->kind);
+            }
+        }
+
+        if(header->category == ROCPROFILER_BUFFER_CATEGORY_COUNTERS && header->kind == 0)
+        {
+            auto* profiler_record = static_cast<rocprofiler_record_counter_t*>(header->payload);
+            rocprofiler_tool_kernel_properties_t kernel_properties =
+                GetKernelProperties(profiler_record->corr_id.internal);
+            rocprofiler_counter_id_t counter_id;
+            const char*              counter_name;
+            size_t                   size, pos;
+            rocprofiler_query_record_counter_id(profiler_record->id, &counter_id);
+            rocprofiler_query_counter_name(counter_id, &counter_name, &size);
+            rocprofiler_query_record_dimension_position(profiler_record->id, 0, &pos);
+
+            auto counter_collection_ss = std::stringstream{};
+            counter_collection_ss << counter_id.handle << ","
+                                  << kernel_properties.gpu_agent.id.handle << ","
+                                  << kernel_properties.queue_id.handle << "," << getpid() << ","
+                                  << kernel_properties.thread_id << ",";
+            counter_collection_ss << kernel_properties.grid_size << ","
+                                  << kernel_properties.kernel_name << ","
+                                  << kernel_properties.workgroup_size << ","
+                                  << ((kernel_properties.lds_size + (lds_block_size - 1)) &
+                                      ~(lds_block_size - 1))
+                                  << "," << kernel_properties.scratch_size << ","
+                                  << kernel_properties.arch_vgpr_count << ","
+                                  << kernel_properties.sgpr_count << ",";
+            /*
+             Iterate through the N dimensional that is obtained for the counter.
+              given instance id what is the counter id
+              given counter id what is the counter name
+              given instance how many dimension
+              iterate through dimensions
+                what is the dimension id
+                what is the dimension name
+                what pos in the dimension.
+            */
+            // ss << counter_name << "[" << info.name << "," << pos << "]" << ",";
+            // ss << profiler_record->counter_value << "\n";
+            counter_collection_ss << counter_name << "["
+                                  << "," << pos << "]"
+                                  << ",";
+            counter_collection_ss << counter_name << ",";
+            counter_collection_ss << profiler_record->counter_value << "\n";
+            get_counter_collection_file() << counter_collection_ss.str() << "\n";
         }
     }
-
-    static auto _sync = std::mutex{};
-    auto        _lk   = std::unique_lock<std::mutex>{_sync};
-    if(get_kernel_trace_file())
-        get_kernel_trace_file() << kernel_trace_ss.str();
-    else
-        std::cerr << "kernel trace file already closed: " << kernel_trace_ss.str();
-
-    (void) (user_data);
 }
 
-rocprofiler_tool_callback_name_info_t
-get_callback_id_names()
+using counter_vec_t = std::vector<rocprofiler_counter_id_t>;
+using agent_counter_map_t =
+    std::unordered_map<const rocprofiler_agent_t*, std::optional<rocprofiler_profile_config_id_t>>;
+
+// this function creates a rocprofiler profile config on the first entry
+auto
+get_agent_profile(const rocprofiler_agent_t* agent)
 {
-    auto cb_name_info = rocprofiler_tool_callback_name_info_t{};
-    //
-    // callback for each kind operation
-    //
-    static auto tracing_kind_operation_cb =
-        [](rocprofiler_callback_tracing_kind_t kindv, uint32_t operation, void* data_v) {
-            auto* name_info_v = static_cast<rocprofiler_tool_callback_name_info_t*>(data_v);
+    static auto data = common::Synchronized<agent_counter_map_t>{};
 
-            if(kindv == ROCPROFILER_CALLBACK_TRACING_HSA_API)
+    auto profile = std::optional<rocprofiler_profile_config_id_t>{};
+    data.ulock(
+        [agent, &profile](const agent_counter_map_t& data_v) {
+            auto itr = data_v.find(agent);
+            if(itr != data_v.end())
             {
-                const char* name = nullptr;
-                ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_operation_name(
-                                     kindv, operation, &name, nullptr),
-                                 "query callback failed");
-                if(name) name_info_v->operation_names[kindv][operation] = name;
+                profile = itr->second;
+                return true;
+            }
+            return false;
+        },
+        [agent, &profile](agent_counter_map_t& data_v) {
+            auto counters_v = counter_vec_t{};
+            ROCPROFILER_CALL(
+                rocprofiler_iterate_agent_supported_counters(
+                    *agent,
+                    [](rocprofiler_counter_id_t* counters, size_t num_counters, void* user_data) {
+                        auto* vec = static_cast<counter_vec_t*>(user_data);
+                        for(size_t i = 0; i < num_counters; i++)
+                        {
+                            const char* name = nullptr;
+                            size_t      len  = 0;
+
+                            ROCPROFILER_CALL(
+                                rocprofiler_query_counter_name(counters[i], &name, &len),
+                                "Could not query name");
+
+                            if(name && len > 0)
+                            {
+                                if(tool::get_config().counters.count(name) > 0)
+                                    vec->emplace_back(counters[i]);
+                            }
+                        }
+                        return ROCPROFILER_STATUS_SUCCESS;
+                    },
+                    static_cast<void*>(&counters_v)),
+                "iterate agent supported counters");
+
+            if(!counters_v.empty())
+            {
+                auto profile_v = rocprofiler_profile_config_id_t{};
+                ROCPROFILER_CALL(rocprofiler_create_profile_config(
+                                     *agent, counters_v.data(), counters_v.size(), &profile_v),
+                                 "Could not construct profile cfg");
+                profile = profile_v;
             }
 
-            if(kindv == ROCPROFILER_CALLBACK_TRACING_HIP_API)
-            {
-                const char* name = nullptr;
-                ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_operation_name(
-                                     kindv, operation, &name, nullptr),
-                                 "query callback failed");
-                if(name) name_info_v->operation_names[kindv][operation] = name;
-            }
-            return 0;
-        };
+            data_v.emplace(agent, profile);
+            return true;
+        });
 
-    //
-    //  callback for each callback kind (i.e. domain)
-    //
-    static auto tracing_kind_cb = [](rocprofiler_callback_tracing_kind_t kind, void* data) {
-        //  store the callback kind name
-        auto*       name_info_v = static_cast<rocprofiler_tool_callback_name_info_t*>(data);
-        const char* name        = nullptr;
-        ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_name(kind, &name, nullptr),
-                         "query callback failed");
+    return profile;
+}
 
-        if(name) name_info_v->kind_names[kind] = name;
+void
+dispatch_callback(rocprofiler_queue_id_t              queue_id,
+                  const rocprofiler_agent_t*          agent,
+                  rocprofiler_correlation_id_t        correlation_id,
+                  const hsa_kernel_dispatch_packet_t* dispatch_packet,
+                  uint64_t                            kernel_id,
+                  void* /*callback_data_args*/,
+                  rocprofiler_profile_config_id_t* config)
+{
+    rocprofiler_tool_kernel_properties_t kernel_properties;
+    const auto&                          kernel_info =
+        kernel_data.rlock([](const kernel_symbol_data_map_t& kdata,
+                             uint64_t kernel_id_v) { return kdata.at(kernel_id_v); },
+                          kernel_id);
 
-        if(kind == ROCPROFILER_CALLBACK_TRACING_HSA_API)
+    auto is_targeted_kernel = [&kernel_info]() {
+        for(const auto& name : tool::get_config().kernel_names)
         {
-            ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kind_operations(
-                                 kind, tracing_kind_operation_cb, static_cast<void*>(data)),
-                             "query callback failed");
+            if(name == kernel_info.truncated_kernel_name)
+                return true;
+            else
+            {
+                auto dkernel_name = std::string_view{kernel_info.demangled_kernel_name};
+                auto pos          = dkernel_name.find(name);
+                // if the demangled kernel name contains name and the next character is '(' then
+                // mark as found
+                if(pos != std::string::npos && (pos + 1) < dkernel_name.size() &&
+                   dkernel_name.at(pos + 1) == '(')
+                    return true;
+            }
         }
-        return 0;
+        return false;
     };
 
-    ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kinds(tracing_kind_cb,
-                                                                static_cast<void*>(&cb_name_info)),
-                     "iterate_callback failed");
+    if(!is_targeted_kernel()) return;
 
-    return cb_name_info;
+    auto profile = get_agent_profile(agent);
+
+    if(profile)
+    {
+        kernel_properties.kernel_name = kernel_info.formatted_kernel_name;
+        kernel_properties.queue_id    = queue_id;
+        kernel_properties.gpu_agent   = *agent;
+        kernel_properties.thread_id   = common::get_tid();
+        populate_kernel_properties_data(&kernel_properties, dispatch_packet);
+        SetKernelProperties(correlation_id.internal, kernel_properties);
+
+        *config = *profile;
+    }
 }
 
+rocprofiler_client_finalize_t client_finalizer  = nullptr;
+rocprofiler_client_id_t*      client_identifier = nullptr;
+
 int
-tool_init(rocprofiler_client_finalize_t /*fini_func*/, void* tool_data)
+tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
 {
-    name_info = get_callback_id_names();
+    client_finalizer = fini_func;
 
-    ROCPROFILER_CALL(rocprofiler_create_context(&context_id), "create context failed");
+    ROCPROFILER_CALL(rocprofiler_create_context(&get_client_ctx()), "create context failed");
 
-    if(common::get_env("ROCPROFILER_KERNEL_TRACE", false))
+    ROCPROFILER_CALL(
+        rocprofiler_configure_callback_tracing_service(get_client_ctx(),
+                                                       ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT,
+                                                       nullptr,
+                                                       0,
+                                                       code_object_tracing_callback,
+                                                       nullptr),
+        "code object tracing configure failed");
+
+    if(tool::get_config().kernel_trace)
     {
-        ROCPROFILER_CALL(
-            rocprofiler_configure_callback_tracing_service(context_id,
-                                                           ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT,
-                                                           nullptr,
-                                                           0,
-                                                           code_object_tracing_callback,
-                                                           nullptr),
-            "tracing configure failed");
-
-        ROCPROFILER_CALL(rocprofiler_create_buffer(context_id,
+        ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(),
                                                    4096,
                                                    2048,
                                                    ROCPROFILER_BUFFER_POLICY_LOSSLESS,
-                                                   kernel_tracing_callback,
+                                                   buffered_callback,
                                                    tool_data,
-                                                   &tool_buffer),
+                                                   &get_kernel_trace_buffer()),
                          "buffer creation");
 
         ROCPROFILER_CALL(
-            rocprofiler_configure_buffer_tracing_service(
-                context_id, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, nullptr, 0, tool_buffer),
+            rocprofiler_configure_buffer_tracing_service(get_client_ctx(),
+                                                         ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH,
+                                                         nullptr,
+                                                         0,
+                                                         get_kernel_trace_buffer()),
             "buffer tracing service for kernel dispatch configure");
     }
 
-    if(common::get_env("ROCPROFILER_HSA_API_TRACE", false))
+    if(tool::get_config().memory_copy_trace)
     {
-        // Requesting all operations
+        ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(),
+                                                   4096,
+                                                   2048,
+                                                   ROCPROFILER_BUFFER_POLICY_LOSSLESS,
+                                                   buffered_callback,
+                                                   nullptr,
+                                                   &get_memory_copy_trace_buffer()),
+                         "create memory copy buffer");
+
         ROCPROFILER_CALL(
-            rocprofiler_configure_callback_tracing_service(context_id,
-                                                           ROCPROFILER_CALLBACK_TRACING_HSA_API,
-                                                           nullptr,
-                                                           0,
-                                                           rocprofiler_tracing_callback,
-                                                           nullptr),
-            "tracing configure failed");
+            rocprofiler_configure_buffer_tracing_service(get_client_ctx(),
+                                                         ROCPROFILER_BUFFER_TRACING_MEMORY_COPY,
+                                                         nullptr,
+                                                         0,
+                                                         get_memory_copy_trace_buffer()),
+            "buffer tracing service for memory copy configure");
     }
 
-    ROCPROFILER_CALL(rocprofiler_start_context(context_id), "start context failed");
+    if(tool::get_config().hsa_api_trace)
+    {
+        ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(),
+                                                   4096,
+                                                   2048,
+                                                   ROCPROFILER_BUFFER_POLICY_LOSSLESS,
+                                                   buffered_callback,
+                                                   tool_data,
+                                                   &get_hsa_api_trace_buffer()),
+                         "buffer creation");
+
+        ROCPROFILER_CALL(
+            rocprofiler_configure_buffer_tracing_service(get_client_ctx(),
+                                                         ROCPROFILER_BUFFER_TRACING_HSA_API,
+                                                         nullptr,
+                                                         0,
+                                                         get_hsa_api_trace_buffer()),
+            "buffer tracing service for memory copy configure");
+    }
+
+    if(tool::get_config().counter_collection)
+    {
+        ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(),
+                                                   4096,
+                                                   2048,
+                                                   ROCPROFILER_BUFFER_POLICY_LOSSLESS,
+                                                   buffered_callback,
+                                                   nullptr,
+                                                   &get_counter_collection_buffer()),
+                         "buffer creation failed");
+
+        ROCPROFILER_CALL(
+            rocprofiler_configure_buffered_dispatch_profile_counting_service(
+                get_client_ctx(), get_counter_collection_buffer(), dispatch_callback, nullptr),
+            "Could not setup buffered service");
+    }
+
+    ROCPROFILER_CALL(rocprofiler_start_context(get_client_ctx()), "start context failed");
+
+    std::atexit([]() {
+        if(client_finalizer && client_identifier) client_finalizer(*client_identifier);
+    });
 
     return 0;
 }
@@ -469,8 +618,12 @@ tool_init(rocprofiler_client_finalize_t /*fini_func*/, void* tool_data)
 void
 tool_fini(void* tool_data)
 {
-    rocprofiler_flush_buffer(tool_buffer);
-    rocprofiler_stop_context(context_id);
+    client_identifier = nullptr;
+    client_finalizer  = nullptr;
+
+    flush();
+    rocprofiler_stop_context(get_client_ctx());
+
     (void) (tool_data);
 }
 
@@ -480,14 +633,17 @@ rocprofiler_configure(uint32_t /*version*/,
                       uint32_t                 priority,
                       rocprofiler_client_id_t* id)
 {
+    common::init_logging("ROCPROF_LOG_LEVEL");
+    FLAGS_colorlogtostderr = true;
+
     // only activate if main tool
     if(priority > 0) return nullptr;
 
     // set the client name
-    id->name = "rocporfiler-tool";
+    id->name = "rocprofiler-tool";
 
     // store client info
-    // client::client_id = id;
+    client_identifier = id;
 
     // create configure data
     static auto cfg = rocprofiler_tool_configure_result_t{

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.cpp

@@ -26,6 +26,7 @@
 #include <rocprofiler-sdk/rocprofiler.h>
 
 #include "lib/common/container/stable_vector.hpp"
+#include "lib/common/static_object.hpp"
 #include "lib/common/utility.hpp"
 #include "lib/rocprofiler-sdk/context/context.hpp"
 #include "lib/rocprofiler-sdk/context/domain.hpp"
@@ -70,24 +71,26 @@ get_buffer_offset()
 bool
 is_valid_buffer_id(rocprofiler_buffer_id_t id)
 {
-    auto nbuffers = get_buffers().size();
+    if(!get_buffers()) return false;
+    auto nbuffers = get_buffers()->size();
     auto offset   = get_buffer_offset();
     return (id.handle >= offset && id.handle < (offset + nbuffers));
 }
 
-unique_buffer_vec_t&
+unique_buffer_vec_t*
 get_buffers()
 {
-    static auto _v = unique_buffer_vec_t{reserve_size_t{unique_buffer_vec_t::chunk_size}};
+    static auto*& _v = common::static_object<unique_buffer_vec_t>::construct(
+        reserve_size_t{unique_buffer_vec_t::chunk_size});
     return _v;
 }
 
 instance*
 get_buffer(rocprofiler_buffer_id_t buffer_id)
 {
-    if(is_valid_buffer_id(buffer_id))
+    if(is_valid_buffer_id(buffer_id) && get_buffers())
     {
-        for(auto& itr : get_buffers())
+        for(auto& itr : *get_buffers())
         {
             if(itr && itr->buffer_id == buffer_id.handle)
             {
@@ -101,6 +104,8 @@ get_buffer(rocprofiler_buffer_id_t buffer_id)
 std::optional<rocprofiler_buffer_id_t>
 allocate_buffer()
 {
+    if(registration::get_fini_status() > 0) return std::nullopt;
+
     // ensure buffer has thread to handle flushing it
     static auto _init_threads_once = std::once_flag{};
     std::call_once(_init_threads_once, []() { internal_threading::initialize(); });
@@ -109,13 +114,13 @@ allocate_buffer()
     auto _lk = std::unique_lock<std::mutex>{get_buffers_mutex()};
 
     // initial context identifier number
-    auto _idx = get_buffer_offset() + get_buffers().size();
+    auto _idx = get_buffer_offset() + CHECK_NOTNULL(get_buffers())->size();
 
     // make space in registered
-    get_buffers().emplace_back(nullptr);
+    CHECK_NOTNULL(get_buffers())->emplace_back(nullptr);
 
     // create an entry in the registered
-    auto& _cfg_v = get_buffers().back();
+    auto& _cfg_v = CHECK_NOTNULL(get_buffers())->back();
     _cfg_v       = allocator::make_unique_static<buffer::instance>();
     auto* _cfg   = _cfg_v.get();
 
@@ -157,7 +162,7 @@ flush(rocprofiler_buffer_id_t buffer_id, bool wait)
     auto idx = buff->buffer_idx++;
 
     auto _task = [buffer_id, idx, offset]() {
-        auto& buff_v          = get_buffers().at(buffer_id.handle - offset);
+        auto& buff_v          = CHECK_NOTNULL(get_buffers())->at(buffer_id.handle - offset);
         auto& buff_internal_v = buff_v->get_internal_buffer(idx);
 
         if(!buff_internal_v.is_empty())
@@ -232,8 +237,8 @@ rocprofiler_create_buffer(rocprofiler_context_id_t        context,
     if(!opt_buff_id) return ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND;
     buffer_id->handle = opt_buff_id->handle;
 
-    auto& buff = rocprofiler::buffer::get_buffers().at(opt_buff_id->handle -
-                                                       rocprofiler::buffer::get_buffer_offset());
+    auto& buff = CHECK_NOTNULL(rocprofiler::buffer::get_buffers())
+                     ->at(opt_buff_id->handle - rocprofiler::buffer::get_buffer_offset());
 
     // allocate the buffers. if it is lossless, we allocate a second buffer to store data while
     // other buffer is being flushed
@@ -264,8 +269,8 @@ rocprofiler_destroy_buffer(rocprofiler_buffer_id_t buffer_id)
         return ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND;
 
     auto  offset  = rocprofiler::buffer::get_buffer_offset();
-    auto& buffers = rocprofiler::buffer::get_buffers();
-    auto& buff    = buffers.at(buffer_id.handle - offset);
+    auto* buffers = CHECK_NOTNULL(rocprofiler::buffer::get_buffers());
+    auto& buff    = buffers->at(buffer_id.handle - offset);
 
     // buffer is currently being flushed or destroyed
     if(buff->syncer.test_and_set()) return ROCPROFILER_STATUS_ERROR_BUFFER_BUSY;

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp

@@ -71,7 +71,7 @@ is_valid_buffer_id(rocprofiler_buffer_id_t id);
 std::optional<rocprofiler_buffer_id_t>
 allocate_buffer();
 
-unique_buffer_vec_t&
+unique_buffer_vec_t*
 get_buffers();
 
 instance*

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer_tracing.cpp

@@ -25,6 +25,7 @@
 
 #include "lib/rocprofiler-sdk/context/context.hpp"
 #include "lib/rocprofiler-sdk/context/domain.hpp"
+#include "lib/rocprofiler-sdk/hsa/async_copy.hpp"
 #include "lib/rocprofiler-sdk/hsa/hsa.hpp"
 #include "lib/rocprofiler-sdk/marker/marker.hpp"
 #include "lib/rocprofiler-sdk/registration.hpp"
@@ -96,6 +97,8 @@ rocprofiler_configure_buffer_tracing_service(rocprofiler_context_id_t          c
 
     if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND;
 
+    if(buffer_id.handle == 0) return ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND;
+
     constexpr auto invalid_buffer_id =
         rocprofiler_buffer_id_t{std::numeric_limits<uint64_t>::max()};
 
@@ -147,6 +150,8 @@ rocprofiler_query_buffer_tracing_kind_operation_name(rocprofiler_buffer_tracing_
     const char* val = nullptr;
     if(kind == ROCPROFILER_BUFFER_TRACING_HSA_API)
         val = rocprofiler::hsa::name_by_id(operation);
+    else if(kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY)
+        val = rocprofiler::hsa::async_copy::name_by_id(operation);
     else if(kind == ROCPROFILER_BUFFER_TRACING_MARKER_API)
         val = rocprofiler::marker::name_by_id<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi>(operation);
     else
@@ -187,6 +192,8 @@ rocprofiler_iterate_buffer_tracing_kind_operations(
     auto ops = std::vector<uint32_t>{};
     if(kind == ROCPROFILER_BUFFER_TRACING_HSA_API)
         ops = rocprofiler::hsa::get_ids();
+    else if(kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY)
+        ops = rocprofiler::hsa::async_copy::get_ids();
     else if(kind == ROCPROFILER_BUFFER_TRACING_MARKER_API)
         ops = rocprofiler::marker::get_ids<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi>();
     else

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp

@@ -420,11 +420,11 @@ deregister_client_contexts(rocprofiler_client_id_t client_id)
 {
     for(auto& itr : get_registered_contexts_impl())
     {
-        if(itr->client_idx == client_id.handle)
+        if(itr->client_idx == client_id.handle && buffer::get_buffers())
         {
-            for(auto& bitr : buffer::get_buffers())
+            for(auto& bitr : *buffer::get_buffers())
             {
-                if(bitr->context_id == itr->context_idx) bitr.reset();
+                if(bitr && bitr->context_id == itr->context_idx) bitr.reset();
             }
             itr.reset();
         }

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp

@@ -23,6 +23,7 @@
 #include "lib/rocprofiler-sdk/hsa/async_copy.hpp"
 #include "lib/common/defines.hpp"
 #include "lib/common/scope_destructor.hpp"
+#include "lib/common/static_object.hpp"
 #include "lib/common/utility.hpp"
 #include "lib/rocprofiler-sdk/agent.hpp"
 #include "lib/rocprofiler-sdk/buffer.hpp"
@@ -61,12 +62,78 @@ namespace rocprofiler
 {
 namespace hsa
 {
+namespace async_copy
+{
 namespace
 {
 using context_t              = context::context;
 using context_array_t        = common::container::small_vector<const context_t*>;
 using external_corr_id_map_t = std::unordered_map<const context_t*, rocprofiler_user_data_t>;
 
+template <size_t Idx>
+struct async_copy_info;
+
+#define SPECIALIZE_ASYNC_COPY_INFO(DIRECTION)                                                      \
+    template <>                                                                                    \
+    struct async_copy_info<ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_##DIRECTION>                     \
+    {                                                                                              \
+        static constexpr auto operation_idx = ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_##DIRECTION;  \
+        static constexpr auto name          = #DIRECTION;                                          \
+    };
+
+SPECIALIZE_ASYNC_COPY_INFO(NONE)
+SPECIALIZE_ASYNC_COPY_INFO(HOST_TO_HOST)
+SPECIALIZE_ASYNC_COPY_INFO(HOST_TO_DEVICE)
+SPECIALIZE_ASYNC_COPY_INFO(DEVICE_TO_HOST)
+SPECIALIZE_ASYNC_COPY_INFO(DEVICE_TO_DEVICE)
+
+#undef SPECIALIZE_ASYNC_COPY_INFO
+
+template <size_t Idx, size_t... IdxTail>
+const char*
+name_by_id(const uint32_t id, std::index_sequence<Idx, IdxTail...>)
+{
+    if(Idx == id) return async_copy_info<Idx>::name;
+    if constexpr(sizeof...(IdxTail) > 0)
+        return name_by_id(id, std::index_sequence<IdxTail...>{});
+    else
+        return nullptr;
+}
+
+template <size_t Idx, size_t... IdxTail>
+uint32_t
+id_by_name(const char* name, std::index_sequence<Idx, IdxTail...>)
+{
+    if(std::string_view{async_copy_info<Idx>::name} == std::string_view{name})
+        return async_copy_info<Idx>::operation_idx;
+    if constexpr(sizeof...(IdxTail) > 0)
+        return id_by_name(name, std::index_sequence<IdxTail...>{});
+    else
+        return ROCPROFILER_HSA_API_ID_NONE;
+}
+
+template <size_t... Idx>
+void
+get_ids(std::vector<uint32_t>& _id_list, std::index_sequence<Idx...>)
+{
+    auto _emplace = [](auto& _vec, uint32_t _v) {
+        if(_v < ROCPROFILER_HSA_API_ID_LAST) _vec.emplace_back(_v);
+    };
+
+    (_emplace(_id_list, async_copy_info<Idx>::operation_idx), ...);
+}
+
+template <size_t... Idx>
+void
+get_names(std::vector<const char*>& _name_list, std::index_sequence<Idx...>)
+{
+    auto _emplace = [](auto& _vec, const char* _v) {
+        if(_v != nullptr && strnlen(_v, 1) > 0) _vec.emplace_back(_v);
+    };
+
+    (_emplace(_name_list, async_copy_info<Idx>::name), ...);
+}
+
 bool
 context_filter(const context::context* ctx)
 {
@@ -90,6 +157,13 @@ struct async_copy_data
     external_corr_id_map_t              extern_corr_ids = {};
 };
 
+auto*
+get_active_signals()
+{
+    static auto* _v = common::static_object<std::atomic<int64_t>>::construct();
+    return _v;
+}
+
 template <typename Tp, typename Up>
 constexpr Tp*
 convert_hsa_handle(Up _hsa_object)
@@ -197,6 +271,7 @@ async_copy_handler(hsa_signal_value_t signal_value, void* arg)
     {
         ROCP_HSA_TABLE_CALL(ERROR, get_table().core_->hsa_signal_destroy_fn(_data->rocp_signal));
         delete _data;
+        get_active_signals()->fetch_sub(1);
     }
 
     return (signal_value > 0);
@@ -367,6 +442,10 @@ async_copy_impl(Args... args)
             return invoke(
                 get_next_dispatch<Idx>(), std::move(_tied_args), std::make_index_sequence<N>{});
         }
+        else
+        {
+            get_active_signals()->fetch_add(1);
+        }
     }
 
     {
@@ -384,6 +463,8 @@ async_copy_impl(Args... args)
             ROCP_HSA_TABLE_CALL(ERROR, get_table().core_->hsa_signal_destroy_fn(_data->rocp_signal))
                 << ":: failed to destroy signal after async handler failed";
 
+            get_active_signals()->fetch_sub(1);
+
             delete _data;
             return invoke(
                 get_next_dispatch<Idx>(), std::move(_tied_args), std::make_index_sequence<N>{});
@@ -455,20 +536,64 @@ using async_copy_index_seq_t =
     std::index_sequence<async_copy_id, async_copy_on_engine_id, async_copy_rect_id>;
 }  // namespace
 
+// check out the assembly here... this compiles to a switch statement
+const char*
+name_by_id(uint32_t id)
+{
+    return name_by_id(id, std::make_index_sequence<ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST>{});
+}
+
+uint32_t
+id_by_name(const char* name)
+{
+    return id_by_name(name,
+                      std::make_index_sequence<ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST>{});
+}
+
+std::vector<uint32_t>
+get_ids()
+{
+    auto _data = std::vector<uint32_t>{};
+    _data.reserve(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST);
+    get_ids(_data, std::make_index_sequence<ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST>{});
+    return _data;
+}
+
+std::vector<const char*>
+get_names()
+{
+    auto _data = std::vector<const char*>{};
+    _data.reserve(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST);
+    get_names(_data, std::make_index_sequence<ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST>{});
+    return _data;
+}
+}  // namespace async_copy
+
 void
 async_copy_init(hsa_api_table_t* _orig)
 {
     if(_orig)
     {
-        async_copy_save(_orig, async_copy_index_seq_t{});
+        async_copy::async_copy_save(_orig, async_copy::async_copy_index_seq_t{});
 
-        auto ctxs = context::get_registered_contexts(context_filter);
+        auto ctxs = context::get_registered_contexts(async_copy::context_filter);
         if(!ctxs.empty())
         {
             _orig->amd_ext_->hsa_amd_profiling_async_copy_enable_fn(true);
-            async_copy_wrap(_orig, async_copy_index_seq_t{});
+            async_copy::async_copy_wrap(_orig, async_copy::async_copy_index_seq_t{});
         }
     }
 }
+
+void
+async_copy_fini()
+{
+    if(!async_copy::get_active_signals()) return;
+    while(async_copy::get_active_signals()->load() > 0)
+    {
+        std::this_thread::yield();
+        std::this_thread::sleep_for(std::chrono::milliseconds{50});
+    }
+}
 }  // namespace hsa
 }  // namespace rocprofiler

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.hpp

@@ -30,7 +30,25 @@ namespace rocprofiler
 {
 namespace hsa
 {
+namespace async_copy
+{
+const char*
+name_by_id(uint32_t id);
+
+uint32_t
+id_by_name(const char* name);
+
+std::vector<const char*>
+get_names();
+
+std::vector<uint32_t>
+get_ids();
+}  // namespace async_copy
+
 void
 async_copy_init(hsa_api_table_t* _orig);
+
+void
+async_copy_fini();
 }  // namespace hsa
 }  // namespace rocprofiler

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp

@@ -557,6 +557,7 @@ finalize()
     static auto _once = std::once_flag{};
     std::call_once(_once, []() {
         set_fini_status(-1);
+        hsa::async_copy_fini();
         hsa::code_object_shutdown();
         if(get_init_status() > 0)
         {

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/buffer.cpp

@@ -40,14 +40,14 @@ TEST(rocprofiler_lib, buffer)
     namespace buffer = ::rocprofiler::buffer;
     namespace common = ::rocprofiler::common;
 
-    ASSERT_EQ(buffer::get_buffers().size(), 0)
+    ASSERT_EQ(buffer::get_buffers()->size(), 0)
         << "no buffers should have been created at this point";
     auto buffer_id = buffer::allocate_buffer();
 
     EXPECT_TRUE(buffer_id) << "failed to allocate buffer";
     EXPECT_GT(buffer_id->handle, 0);
     EXPECT_TRUE(buffer::is_valid_buffer_id(*buffer_id)) << "id=" << buffer_id->handle;
-    ASSERT_EQ(buffer::get_buffers().size(), 1) << "incorrect number of buffers created";
+    ASSERT_EQ(buffer::get_buffers()->size(), 1) << "incorrect number of buffers created";
 
     // get pointer to buffer
     auto* buffer_v = buffer::get_buffer(*buffer_id);

projects/rocprofiler-sdk/tests/CMakeLists.txt

@@ -42,3 +42,6 @@ add_subdirectory(apps)
 # validation tests
 add_subdirectory(kernel-tracing)
 add_subdirectory(async-copy-tracing)
+
+# rocprofv3 validation tests
+add_subdirectory(rocprofv3)

projects/rocprofiler-sdk/tests/apps/CMakeLists.txt

@@ -5,10 +5,15 @@ cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
 
 project(rocprofiler-test-apps LANGUAGES C CXX)
 
+set(CMAKE_BUILD_RPATH "\$ORIGIN:\$ORIGIN/../lib")
+
+# applications used by integration tests which DO NOT link to rocprofiler-sdk-roctx
+add_subdirectory(simple-transpose)
+
 set(CMAKE_BUILD_RPATH
     "\$ORIGIN:\$ORIGIN/../lib:$<TARGET_FILE_DIR:rocprofiler-sdk-roctx::rocprofiler-sdk-roctx-shared-library>"
     )
 
-# applications used by integration tests
+# applications used by integration tests which DO link to rocprofiler-sdk-roctx
 add_subdirectory(reproducible-runtime)
 add_subdirectory(transpose)

projects/rocprofiler-sdk/tests/apps/simple-transpose/CMakeLists.txt

@@ -0,0 +1,47 @@
+#
+#
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+if(NOT CMAKE_HIP_COMPILER)
+    find_program(
+        amdclangpp_EXECUTABLE
+        NAMES amdclang++
+        HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
+        PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
+        PATH_SUFFIXES bin llvm/bin NO_CACHE)
+    mark_as_advanced(amdclangpp_EXECUTABLE)
+
+    if(amdclangpp_EXECUTABLE)
+        set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}")
+    endif()
+endif()
+
+project(rocprofiler-tool-test-app-transpose LANGUAGES CXX HIP)
+
+foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO)
+    if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "")
+        set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}")
+    endif()
+endforeach()
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_EXTENSIONS OFF)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_HIP_STANDARD 17)
+set(CMAKE_HIP_EXTENSIONS OFF)
+set(CMAKE_HIP_STANDARD_REQUIRED ON)
+
+set_source_files_properties(simple-transpose.cpp PROPERTIES LANGUAGE HIP)
+add_executable(simple-transpose)
+target_sources(simple-transpose PRIVATE simple-transpose.cpp)
+target_compile_options(simple-transpose PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow
+                                                -Werror)
+
+find_package(Threads REQUIRED)
+target_link_libraries(simple-transpose PRIVATE Threads::Threads)
+
+install(
+    TARGETS simple-transpose
+    DESTINATION bin
+    COMPONENT tests)

projects/rocprofiler-sdk/tests/apps/simple-transpose/simple-transpose.cpp

@@ -0,0 +1,162 @@
+/*
+Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+
+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 <iostream>
+#include <mutex>
+
+// hip header file
+#include "hip/hip_runtime.h"
+
+#define WIDTH 1024
+
+#define NUM (WIDTH * WIDTH)
+
+#define THREADS_PER_BLOCK_X 4
+#define THREADS_PER_BLOCK_Y 4
+#define THREADS_PER_BLOCK_Z 1
+
+#define HIP_API_CALL(CALL)                                                                         \
+    {                                                                                              \
+        hipError_t error_ = (CALL);                                                                \
+        if(error_ != hipSuccess)                                                                   \
+        {                                                                                          \
+            auto _hip_api_print_lk = auto_lock_t{print_lock};                                      \
+            fprintf(stderr,                                                                        \
+                    "%s:%d :: HIP error : %s\n",                                                   \
+                    __FILE__,                                                                      \
+                    __LINE__,                                                                      \
+                    hipGetErrorString(error_));                                                    \
+            throw std::runtime_error("hip_api_call");                                              \
+        }                                                                                          \
+    }
+
+namespace
+{
+using auto_lock_t = std::unique_lock<std::mutex>;
+auto print_lock   = std::mutex{};
+}  // namespace
+
+// Device (Kernel) function, it must be void
+__global__ void
+matrixTranspose(float* out, float* in, const int width)
+{
+    int x = blockDim.x * blockIdx.x + threadIdx.x;
+    int y = blockDim.y * blockIdx.y + threadIdx.y;
+
+    out[y * width + x] = in[x * width + y];
+}
+
+// CPU implementation of matrix transpose
+void
+matrixTransposeCPUReference(float* output, float* input, const unsigned int width)
+{
+    for(unsigned int j = 0; j < width; j++)
+    {
+        for(unsigned int i = 0; i < width; i++)
+        {
+            output[i * width + j] = input[j * width + i];
+        }
+    }
+}
+
+int
+main()
+{
+    float* Matrix;
+    float* TransposeMatrix;
+    float* cpuTransposeMatrix;
+
+    float* gpuMatrix;
+    float* gpuTransposeMatrix;
+
+    hipDeviceProp_t devProp;
+    HIP_API_CALL(hipGetDeviceProperties(&devProp, 0));
+
+    std::cout << "Device name " << devProp.name << std::endl;
+
+    int i;
+    int errors;
+
+    Matrix             = (float*) malloc(NUM * sizeof(float));
+    TransposeMatrix    = (float*) malloc(NUM * sizeof(float));
+    cpuTransposeMatrix = (float*) malloc(NUM * sizeof(float));
+
+    // initialize the input data
+    for(i = 0; i < NUM; i++)
+    {
+        Matrix[i] = (float) i * 10.0f;
+    }
+
+    // allocate the memory on the device side
+    HIP_API_CALL(hipMalloc((void**) &gpuMatrix, NUM * sizeof(float)));
+    HIP_API_CALL(hipMalloc((void**) &gpuTransposeMatrix, NUM * sizeof(float)));
+
+    // Memory transfer from host to device
+    HIP_API_CALL(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice));
+
+    // Lauching kernel from host
+    hipLaunchKernelGGL(matrixTranspose,
+                       dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y),
+                       dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                       0,
+                       0,
+                       gpuTransposeMatrix,
+                       gpuMatrix,
+                       WIDTH);
+
+    // Memory transfer from device to host
+    HIP_API_CALL(
+        hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost));
+
+    // CPU MatrixTranspose computation
+    matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
+
+    // verify the results
+    errors     = 0;
+    double eps = 1.0E-6;
+    for(i = 0; i < NUM; i++)
+    {
+        if(std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps)
+        {
+            errors++;
+        }
+    }
+    if(errors != 0)
+    {
+        printf("FAILED: %d errors\n", errors);
+    }
+    else
+    {
+        printf("PASSED!\n");
+    }
+
+    // free the resources on device side
+    HIP_API_CALL(hipFree(gpuMatrix));
+    HIP_API_CALL(hipFree(gpuTransposeMatrix));
+
+    // free the resources on host side
+    free(Matrix);
+    free(TransposeMatrix);
+    free(cpuTransposeMatrix);
+
+    return errors;
+}

projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py

@@ -154,26 +154,26 @@ def test_async_copy_direction(input_data):
     sdk_data = data["rocprofiler-sdk-json-tool"]
 
     # Direction values:
-    #  -1 == ??? (unknown)
-    #   0 == H2H (host to host)
-    #   1 == H2D (host to device)
-    #   2 == D2H (device to host)
-    #   3 == D2D (device to device)
-    async_dir_cnt = dict([(idx, 0) for idx in range(-1, 4)])
+    #   0 == ??? (unknown)
+    #   1 == H2H (host to host)
+    #   2 == H2D (host to device)
+    #   3 == D2H (device to host)
+    #   4 == D2D (device to device)
+    async_dir_cnt = dict([(idx, 0) for idx in range(0, 5)])
     for itr in sdk_data["buffer_records"]["memory_copies"]:
         op_id = itr["operation"]
-        assert op_id > 0
-        assert op_id < 3
+        assert op_id > 1
+        assert op_id < 4
         async_dir_cnt[op_id] += 1
 
     # in the transpose test which generates the input file,
     # two threads and the main thread (so three threads total)
     # each perform one H2D + one D2H memory copy
-    assert async_dir_cnt[-1] == 0
     assert async_dir_cnt[0] == 0
-    assert async_dir_cnt[1] == 6
+    assert async_dir_cnt[1] == 0
     assert async_dir_cnt[2] == 6
-    assert async_dir_cnt[3] == 0
+    assert async_dir_cnt[3] == 6
+    assert async_dir_cnt[4] == 0
 
 
 if __name__ == "__main__":

projects/rocprofiler-sdk/tests/kernel-tracing/validate.py

@@ -156,23 +156,23 @@ def test_async_copy_direction(input_data):
     sdk_data = data["rocprofiler-sdk-json-tool"]
 
     # Direction values:
-    #  -1 == ??? (unknown)
-    #   0 == H2H (host to host)
-    #   1 == H2D (host to device)
-    #   2 == D2H (device to host)
-    #   3 == D2D (device to device)
-    async_dir_cnt = dict([(idx, 0) for idx in range(-1, 4)])
+    #   0 == ??? (unknown)
+    #   1 == H2H (host to host)
+    #   2 == H2D (host to device)
+    #   3 == D2H (device to host)
+    #   4 == D2D (device to device)
+    async_dir_cnt = dict([(idx, 0) for idx in range(0, 5)])
     for itr in sdk_data["buffer_records"]["memory_copies"]:
         op_id = itr["operation"]
         async_dir_cnt[op_id] += 1
 
     # in the reproducible-runtime test which generates the input file,
     # we don't expect any async memory copy operations
-    assert async_dir_cnt[-1] == 0
     assert async_dir_cnt[0] == 0
     assert async_dir_cnt[1] == 0
     assert async_dir_cnt[2] == 0
     assert async_dir_cnt[3] == 0
+    assert async_dir_cnt[4] == 0
 
 
 if __name__ == "__main__":

projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt

@@ -0,0 +1,43 @@
+#
+#
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+project(rocprofiler-rocprofv3-tests LANGUAGES C CXX)
+
+set(CMAKE_INSTALL_DEFAULT_COMPONENT_NAME "tests")
+
+if(NOT CMAKE_BUILD_TYPE)
+    set(CMAKE_BUILD_TYPE
+        "Release"
+        CACHE STRING "" FORCE)
+endif()
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_EXTENSIONS OFF)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_HIP_STANDARD 17)
+set(CMAKE_HIP_EXTENSIONS OFF)
+set(CMAKE_HIP_STANDARD_REQUIRED ON)
+
+enable_testing()
+include(CTest)
+
+# needed for validation
+find_package(Python3 REQUIRED)
+
+execute_process(
+    COMMAND ${Python3_EXECUTABLE} -c "import pandas"
+    RESULT_VARIABLE PYTHON_PANDAS_IMPORT
+    OUTPUT_QUIET)
+
+if(NOT ${PYTHON_PANDAS_IMPORT} EQUAL 0)
+    message(
+        FATAL_ERROR
+            "The 'pandas' Python3 package is not installed. Please install it using the following command: '${Python3_EXECUTABLE} -m pip install pandas'"
+        )
+endif()
+
+add_subdirectory(tracing)
+add_subdirectory(counter-collection)
+add_subdirectory(tracing-plus-cc)

projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/CMakeLists.txt

@@ -0,0 +1,7 @@
+foreach(FILENAME conftest.py pytest.ini)
+    configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME}
+                   ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY)
+endforeach()
+
+add_subdirectory(input1)
+add_subdirectory(input2)

projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/conftest.py

@@ -0,0 +1,26 @@
+#!/usr/bin/env python3
+
+import json
+import pytest
+import pandas as pd
+
+
+def pytest_addoption(parser):
+    parser.addoption("--input", action="store", help="Path to csv file.")
+    parser.addoption("--input-dir", action="store", help="Path to output dir.")
+
+
+@pytest.fixture
+def input_data(request):
+    filename = request.config.getoption("--input")
+    if filename:
+        with open(filename, "r") as inp:
+            return pd.read_csv(filename)
+    else:
+        return None
+
+
+@pytest.fixture
+def input_dir(request):
+    dirname = request.config.getoption("--input-dir")
+    return dirname

projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/CMakeLists.txt

@@ -0,0 +1,51 @@
+#
+#
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+project(
+    rocprofiler-tests-counter-collection
+    LANGUAGES CXX
+    VERSION 0.0.0)
+
+find_package(rocprofiler-sdk REQUIRED)
+
+foreach(FILENAME validate.py input.txt)
+    configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME}
+                   ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY)
+endforeach()
+
+# pmc1
+add_test(
+    NAME rocprofv3-test-counter-collection-pmc1-execute
+    COMMAND
+        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> -i
+        ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_1 -o
+        pmc1 $<TARGET_FILE:simple-transpose>)
+
+set(cc-env-pmc1
+    "${PRELOAD_ENV}"
+    "HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>"
+    "LD_LIBRARY_PATH=$<TARGET_FILE_DIR:rocprofiler::rocprofiler-shared-library>:$ENV{LD_LIBRARY_PATH}"
+    )
+
+set_tests_properties(
+    rocprofv3-test-counter-collection-pmc1-execute
+    PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${cc-env-pmc1}"
+               FAIL_REGULAR_EXPRESSION "threw an exception")
+
+add_test(NAME rocprofv3-test-counter-collection-pmc1-validate
+         COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input
+                 "${CMAKE_CURRENT_BINARY_DIR}/out_cc_1/pmc_1/pmc1_counter_collection.csv")
+
+set_tests_properties(
+    rocprofv3-test-counter-collection-pmc1-validate
+    PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS
+               rocprofv3-test-counter-collection-pmc1-execute FAIL_REGULAR_EXPRESSION
+               "threw an exception")
+
+# Needs to be enabled once counter colelction tool support is merged
+set_tests_properties(rocprofv3-test-counter-collection-pmc1-execute PROPERTIES DISABLED
+                                                                               TRUE)
+set_tests_properties(rocprofv3-test-counter-collection-pmc1-validate PROPERTIES DISABLED
+                                                                                TRUE)

projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/input.txt

@@ -0,0 +1 @@
+pmc: SQ_WAVES

projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/validate.py

@@ -0,0 +1,16 @@
+import pandas as pd
+import sys
+import pytest
+
+
+def test_validate_counter_collection_pmc1(input_data: pd.DataFrame):
+    df = input_data
+
+    assert df.empty == False
+    assert df["agent-id"].map(type).eq(int).all()
+    assert len(df["kernel-name"]) > 0
+
+
+if __name__ == "__main__":
+    exit_code = pytest.main(["-x", __file__] + sys.argv[1:])
+    sys.exit(exit_code)

projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/CMakeLists.txt

@@ -0,0 +1,50 @@
+#
+#
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+project(
+    rocprofiler-tests-counter-collection
+    LANGUAGES CXX
+    VERSION 0.0.0)
+
+find_package(rocprofiler-sdk REQUIRED)
+
+foreach(FILENAME validate.py input.txt)
+    configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME}
+                   ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY)
+endforeach()
+
+# pmc2
+add_test(
+    NAME rocprofv3-test-counter-collection-pmc2-execute
+    COMMAND
+        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> -i
+        ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_2 -o
+        pmc2 $<TARGET_FILE:simple-transpose>)
+set(cc-env-pmc2
+    "${PRELOAD_ENV}"
+    "HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>"
+    "LD_LIBRARY_PATH=$<TARGET_FILE_DIR:rocprofiler::rocprofiler-shared-library>:$ENV{LD_LIBRARY_PATH}"
+    )
+
+set_tests_properties(
+    rocprofv3-test-counter-collection-pmc2-execute
+    PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${cc-env-pmc2}"
+               FAIL_REGULAR_EXPRESSION "threw an exception")
+
+add_test(NAME rocprofv3-test-counter-collection-pmc2-validate
+         COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py
+                 --input-dir "${CMAKE_CURRENT_BINARY_DIR}/out_cc_2")
+
+set_tests_properties(
+    rocprofv3-test-counter-collection-pmc2-validate
+    PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS
+               rocprofv3-test-counter-collection-pmc2-execute FAIL_REGULAR_EXPRESSION
+               "threw an exception")
+
+# Needs to be enabled once counter colelction tool support is merged
+set_tests_properties(rocprofv3-test-counter-collection-pmc2-execute PROPERTIES DISABLED
+                                                                               TRUE)
+set_tests_properties(rocprofv3-test-counter-collection-pmc2-validate PROPERTIES DISABLED
+                                                                                TRUE)

projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/input.txt

@@ -0,0 +1,2 @@
+pmc: SQ_WAVES
+pmc: GRBM_COUNT

projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/validate.py

@@ -0,0 +1,44 @@
+import pandas as pd
+import os
+import sys
+import pytest
+
+
+def test_validate_counter_collection_pmc2(input_dir: pd.DataFrame):
+    directory_path = input_dir
+
+    # Check if the directory is not empty
+    assert os.path.isdir(directory_path), f"{directory_path} is not a directory."
+    assert os.listdir(directory_path), f"{directory_path} is empty."
+
+    # Check if there are 2 subdirectories pmc_1 and pmc_2
+    subdirectories = [
+        d
+        for d in os.listdir(directory_path)
+        if os.path.isdir(os.path.join(directory_path, d))
+    ]
+    assert (
+        len(subdirectories) == 2
+    ), f"Expected 2 subdirectories, found {len(subdirectories)}."
+
+    # Check if each subdirectory has files
+    for subdirectory in subdirectories:
+        subdirectory_path = os.path.join(directory_path, subdirectory)
+        assert os.listdir(subdirectory_path), f"{subdirectory_path} is empty."
+
+        # Check if each file in the subdirectory has some data
+        for file_name in os.listdir(subdirectory_path):
+            file_path = os.path.join(subdirectory_path, file_name)
+            assert os.path.isfile(file_path), f"{file_path} is not a file."
+
+            with open(file_path, "r") as file:
+                df = pd.read_csv(file)
+                # check if kernel-name is present
+                assert len(df["kernel-name"]) > 0
+                # check if counter value is positive
+                assert len(df["counter_value"]) > 0
+
+
+if __name__ == "__main__":
+    exit_code = pytest.main(["-x", __file__] + sys.argv[1:])
+    sys.exit(exit_code)

projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/pytest.ini

@@ -0,0 +1,5 @@
+
+[pytest]
+addopts = --durations=20 -ras -vv
+testpaths = input1/validate.py
+            input2/validate.py

projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt

@@ -0,0 +1,49 @@
+#
+#
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+project(
+    rocprofiler-tests-counter-collection
+    LANGUAGES CXX
+    VERSION 0.0.0)
+
+find_package(rocprofiler-sdk REQUIRED)
+
+foreach(FILENAME validate.py conftest.py pytest.ini input.txt)
+    configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME}
+                   ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY)
+endforeach()
+
+# pmc3
+add_test(
+    NAME rocprofv3-test-tracing-plus-cc-execute
+    COMMAND
+        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --hsa-trace -i
+        ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace
+        -o pmc3 $<TARGET_FILE:simple-transpose>)
+
+set(cc-tracing-env
+    "${PRELOAD_ENV}"
+    "HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>"
+    "LD_LIBRARY_PATH=$<TARGET_FILE_DIR:rocprofiler::rocprofiler-shared-library>:$ENV{LD_LIBRARY_PATH}"
+    )
+
+set_tests_properties(
+    rocprofv3-test-tracing-plus-cc-execute
+    PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${cc-tracing-env}"
+               FAIL_REGULAR_EXPRESSION "threw an exception")
+
+add_test(NAME rocprofv3-test-tracing-plus-cc-validate
+         COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py
+                 --input-dir "${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace")
+
+set_tests_properties(
+    rocprofv3-test-tracing-plus-cc-validate
+    PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS
+               rocprofv3-test-tracing-plus-cc-execute FAIL_REGULAR_EXPRESSION
+               "threw an exception")
+
+# Needs to be enabled once counter colelction tool support is merged
+set_tests_properties(rocprofv3-test-tracing-plus-cc-execute PROPERTIES DISABLED TRUE)
+set_tests_properties(rocprofv3-test-tracing-plus-cc-validate PROPERTIES DISABLED TRUE)

projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/conftest.py

@@ -0,0 +1,15 @@
+#!/usr/bin/env python3
+
+import json
+import pytest
+import pandas as pd
+
+
+def pytest_addoption(parser):
+    parser.addoption("--input-dir", action="store", help="Path to output dir.")
+
+
+@pytest.fixture
+def input_dir(request):
+    dirname = request.config.getoption("--input-dir")
+    return dirname

projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/input.txt

@@ -0,0 +1,6 @@
+# multi block counters
+
+pmc: GRBM_COUNT GRBM_GUI_ACTIVE SQ_CYCLES SQ_BUSY_CYCLES SQ_WAVES 
+pmc: TCC_REQ_sum                            TCC_STREAMING_REQ_sum                 TCC_HIT_sum                            TCC_MISS_sum
+pmc: TCC_EA_WRREQ_sum                       TCC_EA_WRREQ_64B_sum                  TCC_EA_WR_UNCACHED_32B_sum 
+pmc: TCC_TOO_MANY_EA_WRREQS_STALL_sum       TCC_EA_ATOMIC_sum                     TCC_EA_RDREQ_sum                       TCC_EA_RDREQ_32B_sum

projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/pytest.ini

@@ -0,0 +1,4 @@
+
+[pytest]
+addopts = --durations=20 -ras -vv
+testpaths = validate.py

projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/validate.py

@@ -0,0 +1,42 @@
+import pandas as pd
+import os
+import sys
+import pytest
+
+
+def test_validate_counter_collection_plus_tracing(input_dir: pd.DataFrame):
+    directory_path = input_dir
+
+    # Check if the directory is not empty
+    assert os.path.isdir(directory_path), f"{directory_path} is not a directory."
+    assert os.listdir(directory_path), f"{directory_path} is empty."
+
+    # Check if there are 4 subdirectories for pmc's
+    subdirectories = [
+        d
+        for d in os.listdir(directory_path)
+        if os.path.isdir(os.path.join(directory_path, d))
+    ]
+    assert (
+        len(subdirectories) == 4
+    ), f"Expected 4 subdirectories, found {len(subdirectories)}."
+
+    # Check if each subdirectory has files
+    for subdirectory in subdirectories:
+        subdirectory_path = os.path.join(directory_path, subdirectory)
+        assert os.listdir(subdirectory_path), f"{subdirectory_path} is empty."
+
+        # Check if each file in the subdirectory has some data
+        for file_name in os.listdir(subdirectory_path):
+            file_path = os.path.join(subdirectory_path, file_name)
+            assert os.path.isfile(file_path), f"{file_path} is not a file."
+
+            with open(file_path, "r") as file:
+                df = pd.read_csv(file)
+                # check if either kernel-name/FUNCTION is present
+                assert "kernel-name" in df.columns or "FUNCTION" in df.columns
+
+
+if __name__ == "__main__":
+    exit_code = pytest.main(["-x", __file__] + sys.argv[1:])
+    sys.exit(exit_code)

projects/rocprofiler-sdk/tests/rocprofv3/tracing/CMakeLists.txt

@@ -0,0 +1,47 @@
+#
+# rocprofv3 tool test
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+project(
+    rocprofiler-tests-rocprofv3-tracing
+    LANGUAGES CXX
+    VERSION 0.0.0)
+
+find_package(rocprofiler-sdk REQUIRED)
+
+add_test(
+    NAME rocprofv3-test-trace-execute
+    COMMAND
+        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --hsa-trace --kernel-trace
+        --memory-copy-trace -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out
+        $<TARGET_FILE:simple-transpose>)
+
+string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV
+               "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}")
+
+set(tracing-env "${PRELOAD_ENV}"
+                "HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>")
+
+set_tests_properties(rocprofv3-test-trace-execute PROPERTIES LABELS "integration-tests"
+                                                             ENVIRONMENT "${tracing-env}")
+
+foreach(FILENAME validate.py conftest.py)
+    configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME}
+                   ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY)
+endforeach()
+
+add_test(
+    NAME rocprofv3-test-trace-validate
+    COMMAND
+        ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --hsa-input
+        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_hsa_api_trace.csv
+        --kernel-input
+        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_kernel_trace.csv
+        --memory-copy-input
+        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_memory_copy_trace.csv)
+
+set_tests_properties(
+    rocprofv3-test-trace-validate
+    PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS rocprofv3-test-trace-execute
+               FAIL_REGULAR_EXPRESSION "AssertionError")

projects/rocprofiler-sdk/tests/rocprofv3/tracing/conftest.py

@@ -0,0 +1,58 @@
+#!/usr/bin/env python3
+
+import csv
+import pytest
+
+
+def pytest_addoption(parser):
+    parser.addoption(
+        "--hsa-input",
+        action="store",
+        help="Path to HSA API tracing CSV file.",
+    )
+    parser.addoption(
+        "--kernel-input",
+        action="store",
+        help="Path to kernel tracing CSV file.",
+    )
+    parser.addoption(
+        "--memory-copy-input",
+        action="store",
+        help="Path to memory-copy tracing CSV file.",
+    )
+
+
+@pytest.fixture
+def hsa_input_data(request):
+    filename = request.config.getoption("--hsa-input")
+    data = []
+    with open(filename, "r") as inp:
+        reader = csv.DictReader(inp)
+        for row in reader:
+            data.append(row)
+
+    return data
+
+
+@pytest.fixture
+def kernel_input_data(request):
+    filename = request.config.getoption("--kernel-input")
+    data = []
+    with open(filename, "r") as inp:
+        reader = csv.DictReader(inp)
+        for row in reader:
+            data.append(row)
+
+    return data
+
+
+@pytest.fixture
+def memory_copy_input_data(request):
+    filename = request.config.getoption("--memory-copy-input")
+    data = []
+    with open(filename, "r") as inp:
+        reader = csv.DictReader(inp)
+        for row in reader:
+            data.append(row)
+
+    return data

projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py

@@ -0,0 +1,68 @@
+#!/usr/bin/env python3
+
+import sys
+import pytest
+
+
+def test_hsa_api_trace(hsa_input_data):
+    functions = []
+    correlation_ids = []
+    for row in hsa_input_data:
+        assert row["Domain"] == "HSA_API"
+        assert int(row["Process_Id"]) > 0
+        assert int(row["Thread_Id"]) >= int(row["Process_Id"])
+        assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"])
+        functions.append(row["Function"])
+        correlation_ids.append(int(row["Correlation_Id"]))
+
+    correlation_ids = sorted(list(set(correlation_ids)))
+
+    # all correlation ids are unique
+    assert len(correlation_ids) == len(hsa_input_data)
+    # correlation ids are numbered from 1 to N
+    assert correlation_ids[0] == 1
+    assert correlation_ids[-1] == len(correlation_ids)
+
+    functions = list(set(functions))
+    assert "hsa_amd_memory_async_copy_on_engine" in functions
+
+
+def test_kernel_trace(kernel_input_data):
+    for row in kernel_input_data:
+        assert row["Kind"] == "KERNEL_DISPATCH"
+        assert int(row["Agent_Id"]) > 0
+        assert int(row["Queue_Id"]) > 0
+        assert int(row["Kernel_Id"]) > 0
+        assert row["Kernel_Name"] == "matrixTranspose(float*, float*, int)"
+        assert int(row["Correlation_Id"]) > 0
+        assert int(row["Workgroup_Size_X"]) == 4
+        assert int(row["Workgroup_Size_Y"]) == 4
+        assert int(row["Workgroup_Size_Z"]) == 1
+        assert int(row["Grid_Size_X"]) == 1024
+        assert int(row["Grid_Size_Y"]) == 1024
+        assert int(row["Grid_Size_Z"]) == 1
+        assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"])
+
+
+def test_memory_copy_trace(memory_copy_input_data):
+    for row in memory_copy_input_data:
+        assert row["Kind"] == "MEMORY_COPY"
+
+    row = memory_copy_input_data[0]
+    assert row["Direction"] == "HOST_TO_DEVICE"
+    assert int(row["Source_Agent_Id"]) == 0
+    assert int(row["Destination_Agent_Id"]) == 1
+    assert int(row["Correlation_Id"]) > 0
+    assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"])
+
+    row = memory_copy_input_data[1]
+    assert row["Direction"] == "DEVICE_TO_HOST"
+    assert int(row["Source_Agent_Id"]) == 1
+    assert int(row["Destination_Agent_Id"]) == 0
+    assert int(row["Correlation_Id"]) > 0
+    assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"])
+
+
+if __name__ == "__main__":
+    exit_code = pytest.main(["-x", __file__] + sys.argv[1:])
+    sys.exit(exit_code)

projects/rocprofiler-sdk/tests/tools/json-tool.cpp

@@ -181,7 +181,9 @@ buffer_name_info
 get_buffer_tracing_names()
 {
     static const auto supported = std::unordered_set<rocprofiler_buffer_tracing_kind_t>{
-        ROCPROFILER_BUFFER_TRACING_HSA_API, ROCPROFILER_BUFFER_TRACING_MARKER_API};
+        ROCPROFILER_BUFFER_TRACING_HSA_API,
+        ROCPROFILER_BUFFER_TRACING_MARKER_API,
+        ROCPROFILER_BUFFER_TRACING_MEMORY_COPY};
 
     auto cb_name_info = buffer_name_info{};
     //