diff --git a/projects/rocprofiler-systems/.github/workflows/formatting.yml b/projects/rocprofiler-systems/.github/workflows/formatting.yml
index e16f796a77..01a166ced9 100644
--- a/projects/rocprofiler-systems/.github/workflows/formatting.yml
+++ b/projects/rocprofiler-systems/.github/workflows/formatting.yml
@@ -74,7 +74,7 @@ jobs:
     - name: clang-format
       run: |
         set +e
-        FILES=$(find include src examples -type f | egrep '\.hpp$|\.cpp$|\.cpp\.in$')
+        FILES=$(find source examples -type f | egrep '\.hpp$|\.cpp$|\.cpp\.in$')
         FORMAT_OUT=$(clang-format-11 -output-replacements-xml ${FILES})
         RET=$(echo ${FORMAT_OUT} | grep -c '<replacement ')
         if [ "${RET}" -ne 0 ]; then
@@ -83,3 +83,21 @@ jobs:
           git diff
           exit ${RET}
         fi
+
+  includes:
+    runs-on: ubuntu-20.04
+
+    steps:
+    - uses: actions/checkout@v2
+    - name: check-includes
+      run: |
+        set +e
+        FILES=$(find source examples -type f | egrep '\.hpp$|\.cpp$|\.cpp\.in$')
+        MATCHES=$(egrep 'include "timemory|include <bits' ${FILES})
+        if [ -n "${MATCHES}" ]; then
+          echo -e "\nError! Included timemory header with quotes or bits folder included\n"
+          echo -e "### MATCHES: ###"
+          echo -e "${MATCHES}"
+          echo -e "################"
+          exit 1
+        fi
diff --git a/projects/rocprofiler-systems/CMakeLists.txt b/projects/rocprofiler-systems/CMakeLists.txt
index 5c014697ae..9091751127 100644
--- a/projects/rocprofiler-systems/CMakeLists.txt
+++ b/projects/rocprofiler-systems/CMakeLists.txt
@@ -27,10 +27,28 @@ project(
     DESCRIPTION "CPU/GPU Application tracing with static/dynamic binary instrumentation"
     HOMEPAGE_URL "https://github.com/AMDResearch/omnitrace")
 
+find_package(Git)
+
+if(Git_FOUND AND EXISTS "${PROJECT_SOURCE_DIR}/.git")
+    execute_process(
+        COMMAND ${GIT_EXECUTABLE} describe --tags
+        OUTPUT_VARIABLE OMNITRACE_GIT_DESCRIBE
+        OUTPUT_STRIP_TRAILING_WHITESPACE)
+    execute_process(
+        COMMAND ${GIT_EXECUTABLE} rev-parse HEAD
+        OUTPUT_VARIABLE OMNITRACE_GIT_REVISION
+        OUTPUT_STRIP_TRAILING_WHITESPACE)
+else()
+    set(OMNITRACE_GIT_DESCRIBE "v${OMNITRACE_VERSION}")
+    set(OMNITRACE_GIT_REVISION "")
+endif()
+
 message(
     STATUS
-        "[${PROJECT_NAME}] version ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH}"
+        "[${PROJECT_NAME}] version ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH} (${FULL_VERSION_STRING})"
     )
+message(STATUS "[${PROJECT_NAME}] git revision: ${OMNITRACE_GIT_REVISION}")
+message(STATUS "[${PROJECT_NAME}] git describe: ${OMNITRACE_GIT_DESCRIBE}")
 set(CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake ${PROJECT_SOURCE_DIR}/cmake/Modules
                       ${PROJECT_SOURCE_DIR}/source/python/cmake ${CMAKE_MODULE_PATH})
 set(BUILD_SHARED_LIBS
@@ -103,6 +121,8 @@ omnitrace_add_option(OMNITRACE_USE_HIP "Enable HIP support" ON)
 omnitrace_add_option(OMNITRACE_USE_PAPI "Enable HW counter support via PAPI" ON)
 omnitrace_add_option(OMNITRACE_USE_ROCTRACER "Enable roctracer support"
                      ${OMNITRACE_USE_HIP})
+omnitrace_add_option(OMNITRACE_USE_ROCPROFILER "Enable rocprofiler support"
+                     ${OMNITRACE_USE_HIP})
 omnitrace_add_option(
     OMNITRACE_USE_ROCM_SMI "Enable rocm-smi support for power/temp/etc. sampling"
     ${OMNITRACE_USE_HIP})
@@ -149,16 +169,20 @@ if(NOT OMNITRACE_USE_HIP)
     set(OMNITRACE_USE_ROCTRACER
         OFF
         CACHE BOOL "Disabled via OMNITRACE_USE_HIP=OFF" FORCE)
+    set(OMNITRACE_USE_ROCPROFILER
+        OFF
+        CACHE BOOL "Disabled via OMNITRACE_USE_HIP=OFF" FORCE)
     set(OMNITRACE_USE_ROCM_SMI
         OFF
         CACHE BOOL "Disabled via OMNITRACE_USE_HIP=OFF" FORCE)
 elseif(
     OMNITRACE_USE_HIP
     AND NOT OMNITRACE_USE_ROCTRACER
+    AND NOT OMNITRACE_USE_ROCPROFILER
     AND NOT OMNITRACE_USE_ROCM_SMI)
     omnitrace_message(
         AUTHOR_WARNING
-        "Setting OMNITRACE_USE_HIP=OFF because roctracer and rocm-smi options are disabled"
+        "Setting OMNITRACE_USE_HIP=OFF because roctracer, rocprofiler, and rocm-smi options are disabled"
         )
     set(OMNITRACE_USE_HIP OFF)
 endif()
diff --git a/projects/rocprofiler-systems/VERSION b/projects/rocprofiler-systems/VERSION
index 26aaba0e86..e0f931ac87 100644
--- a/projects/rocprofiler-systems/VERSION
+++ b/projects/rocprofiler-systems/VERSION
@@ -1 +1 @@
-1.2.0
+1.3.0dev0
diff --git a/projects/rocprofiler-systems/cmake/ConfigCPack.cmake b/projects/rocprofiler-systems/cmake/ConfigCPack.cmake
index c8a1d83bed..15f16e742a 100644
--- a/projects/rocprofiler-systems/cmake/ConfigCPack.cmake
+++ b/projects/rocprofiler-systems/cmake/ConfigCPack.cmake
@@ -66,9 +66,9 @@ endif()
 
 if(OMNITRACE_USE_MPI)
     set(VALID_MPI_IMPLS "mpich" "openmpi")
-    if("${MPI_C_COMPILER_INCLUDE_DIRS}" MATCHES "openmpi")
+    if("${MPI_C_COMPILER_INCLUDE_DIRS};${MPI_C_HEADER_DIR}" MATCHES "openmpi")
         set(OMNITRACE_MPI_IMPL "openmpi")
-    elseif("${MPI_C_COMPILER_INCLUDE_DIRS}" MATCHES "mpich")
+    elseif("${MPI_C_COMPILER_INCLUDE_DIRS};${MPI_C_HEADER_DIR}" MATCHES "mpich")
         set(OMNITRACE_MPI_IMPL "mpich")
     else()
         message(
@@ -90,6 +90,8 @@ if(OMNITRACE_USE_MPI)
         set(OMNITRACE_MPI_IMPL_UPPER "OpenMPI")
     elseif("${OMNITRACE_MPI_IMPL}" STREQUAL "mpich")
         set(OMNITRACE_MPI_IMPL_UPPER "MPICH")
+    else()
+        set(OMNITRACE_MPI_IMPL_UPPER "MPI")
     endif()
     set(OMNITRACE_CPACK_PACKAGE_SUFFIX
         "${OMNITRACE_CPACK_PACKAGE_SUFFIX}-${OMNITRACE_MPI_IMPL_UPPER}")
diff --git a/projects/rocprofiler-systems/cmake/MacroUtilities.cmake b/projects/rocprofiler-systems/cmake/MacroUtilities.cmake
index 40bc75a661..7596baf6a5 100644
--- a/projects/rocprofiler-systems/cmake/MacroUtilities.cmake
+++ b/projects/rocprofiler-systems/cmake/MacroUtilities.cmake
@@ -509,10 +509,10 @@ function(omnitrace_custom_compilation)
     cmake_parse_arguments(COMP "GLOBAL;PROJECT" "COMPILER" "DIRECTORY;TARGET;SOURCE"
                           ${ARGN})
 
-    # find omnitrace_launch_compiler
+    # find omnitrace-launch-compiler
     find_program(
         OMNITRACE_COMPILE_LAUNCHER
-        NAMES omnitrace_launch_compiler
+        NAMES omnitrace-launch-compiler
         HINTS ${PROJECT_SOURCE_DIR} ${CMAKE_SOURCE_DIR}
         PATHS ${PROJECT_SOURCE_DIR} ${CMAKE_SOURCE_DIR}
         PATH_SUFFIXES scripts bin)
@@ -524,7 +524,7 @@ function(omnitrace_custom_compilation)
     if(NOT OMNITRACE_COMPILE_LAUNCHER)
         message(
             FATAL_ERROR
-                "omnitrace could not find 'omnitrace_launch_compiler'. Please set '-DOMNITRACE_COMPILE_LAUNCHER=/path/to/launcher'"
+                "omnitrace could not find 'omnitrace-launch-compiler'. Please set '-DOMNITRACE_COMPILE_LAUNCHER=/path/to/launcher'"
             )
     endif()
 
@@ -797,7 +797,7 @@ function(OMNITRACE_INSTALL_TPL _TPL_TARGET _NEW_NAME _BUILD_TREE_DIR)
     # build tree symbolic links
     add_custom_target(
         ${_NEW_NAME}-library ALL
-        ${CMAKE_COMMAND} -E create_symlink $<TARGET_FILE:${_TPL_TARGET}>
+        ${CMAKE_COMMAND} -E copy_if_different $<TARGET_FILE:${_TPL_TARGET}>
         ${_TPL_PREFIX}${_NEW_NAME}${_TPL_SUFFIX}.${_TPL_VERSION}
         COMMAND
             ${CMAKE_COMMAND} -E create_symlink
@@ -809,7 +809,7 @@ function(OMNITRACE_INSTALL_TPL _TPL_TARGET _NEW_NAME _BUILD_TREE_DIR)
             ${_BUILD_TREE_DIR}/${_TPL_PREFIX}${_NEW_NAME}${_TPL_SUFFIX}
         WORKING_DIRECTORY ${_BUILD_TREE_DIR}
         DEPENDS ${_TPL_TARGET}
-        COMMENT "Creating ${_NEW_NAME} symbolic links to ${_TPL_TARGET}...")
+        COMMENT "Creating ${_NEW_NAME} from ${_TPL_TARGET}...")
 
     install(
         FILES $<TARGET_FILE:${_TPL_TARGET}>
diff --git a/projects/rocprofiler-systems/cmake/Modules/FindROCmVersion.cmake b/projects/rocprofiler-systems/cmake/Modules/FindROCmVersion.cmake
index 58526ab559..9efb27b762 100644
--- a/projects/rocprofiler-systems/cmake/Modules/FindROCmVersion.cmake
+++ b/projects/rocprofiler-systems/cmake/Modules/FindROCmVersion.cmake
@@ -196,7 +196,7 @@ function(ROCM_VERSION_PARSE_VERSION_FILES)
         if(EXISTS "${ROCmVersion_VERSION_FILE}" AND IS_ABSOLUTE
                                                     "${ROCmVersion_VERSION_FILE}")
             get_filename_component(_VERSION_DIR "${ROCmVersion_VERSION_FILE}" PATH)
-            get_filename_component(_VERSION_DIR "${_VERSION_DIR}/.." ABSOLUTE)
+            get_filename_component(_VERSION_DIR "${_VERSION_DIR}/.." REALPATH)
             set(ROCmVersion_DIR
                 "${_VERSION_DIR}"
                 CACHE PATH "Root path to ROCm's .info/${ROCmVersion_VERSION_FILE}"
@@ -227,8 +227,14 @@ function(ROCM_VERSION_PARSE_VERSION_FILES)
     if(ROCmVersion_DIR)
         set(_PATHS ${ROCmVersion_DIR})
     else()
-        set(_PATHS ${ROCmVersion_DIR} ${ROCmVersion_ROOT} ${ROCmVersion_ROOT_DIR}
-                   $ENV{CMAKE_PREFIX_PATH} ${CMAKE_PREFIX_PATH} ${ROCM_PATH} /opt/rocm)
+        set(_PATHS)
+        foreach(_DIR ${ROCmVersion_DIR} ${ROCmVersion_ROOT} ${ROCmVersion_ROOT_DIR}
+                     $ENV{CMAKE_PREFIX_PATH} ${CMAKE_PREFIX_PATH} ${ROCM_PATH} /opt/rocm)
+            if(EXISTS ${_DIR})
+                get_filename_component(_ABS_DIR "${_DIR}" REALPATH)
+                list(APPEND _PATHS ${_ABS_DIR})
+            endif()
+        endforeach()
         rocm_version_message(STATUS "ROCmVersion search paths: ${_PATHS}")
     endif()
 
diff --git a/projects/rocprofiler-systems/cmake/Modules/Findrocm-smi.cmake b/projects/rocprofiler-systems/cmake/Modules/Findrocm-smi.cmake
index 034891d970..1136f6f9d5 100644
--- a/projects/rocprofiler-systems/cmake/Modules/Findrocm-smi.cmake
+++ b/projects/rocprofiler-systems/cmake/Modules/Findrocm-smi.cmake
@@ -5,15 +5,24 @@ include(FindPackageHandleStandardArgs)
 
 # ----------------------------------------------------------------------------------------#
 
-# set(_ROCM_PATHS $ENV{ROCM_HOME} /opt/rocm)
+if(NOT ROCM_PATH AND NOT "$ENV{ROCM_PATH}" STREQUAL "")
+    set(ROCM_PATH "$ENV{ROCM_PATH}")
+endif()
+
+foreach(_DIR ${ROCM_PATH} /opt/rocm /opt/rocm/rocm_smi)
+    if(EXISTS ${_DIR})
+        get_filename_component(_ABS_DIR "${_DIR}" REALPATH)
+        list(APPEND _ROCM_SMI_PATHS ${_ABS_DIR})
+    endif()
+endforeach()
 
 # ----------------------------------------------------------------------------------------#
 
 find_path(
     rocm-smi_ROOT_DIR
     NAMES include/rocm_smi/rocm_smi.h
-    HINTS ${_ROCM_PATHS}
-    PATHS ${_ROCM_PATHS}
+    HINTS ${_ROCM_SMI_PATHS}
+    PATHS ${_ROCM_SMI_PATHS}
     PATH_SUFFIXES rocm_smi)
 
 mark_as_advanced(rocm-smi_ROOT_DIR)
@@ -23,9 +32,9 @@ mark_as_advanced(rocm-smi_ROOT_DIR)
 find_path(
     rocm-smi_INCLUDE_DIR
     NAMES rocm_smi/rocm_smi.h
-    HINTS ${rocm-smi_ROOT_DIR} ${_ROCM_PATHS}
-    PATHS ${rocm-smi_ROOT_DIR} ${_ROCM_PATHS}
-    PATH_SUFFIXES rocm_smi/include rocm_smi)
+    HINTS ${rocm-smi_ROOT_DIR} ${_ROCM_SMI_PATHS}
+    PATHS ${rocm-smi_ROOT_DIR} ${_ROCM_SMI_PATHS}
+    PATH_SUFFIXES include rocm_smi/include)
 
 mark_as_advanced(rocm-smi_INCLUDE_DIR)
 
@@ -34,8 +43,8 @@ mark_as_advanced(rocm-smi_INCLUDE_DIR)
 find_library(
     rocm-smi_LIBRARY
     NAMES rocm_smi64 rocm_smi
-    HINTS ${rocm-smi_ROOT_DIR} ${_ROCM_PATHS}
-    PATHS ${rocm-smi_ROOT_DIR} ${_ROCM_PATHS}
+    HINTS ${rocm-smi_ROOT_DIR} ${_ROCM_SMI_PATHS}
+    PATHS ${rocm-smi_ROOT_DIR} ${_ROCM_SMI_PATHS}
     PATH_SUFFIXES rocm_smi/lib rocm_smi/lib64 lib lib64)
 
 if(rocm-smi_LIBRARY)
@@ -65,6 +74,6 @@ endif()
 
 # ------------------------------------------------------------------------------#
 
-unset(_ROCM_PATHS)
+unset(_ROCM_SMI_PATHS)
 
 # ------------------------------------------------------------------------------#
diff --git a/projects/rocprofiler-systems/cmake/Modules/Findrocprofiler.cmake b/projects/rocprofiler-systems/cmake/Modules/Findrocprofiler.cmake
new file mode 100644
index 0000000000..c24079f6f7
--- /dev/null
+++ b/projects/rocprofiler-systems/cmake/Modules/Findrocprofiler.cmake
@@ -0,0 +1,100 @@
+# Distributed under the OSI-approved BSD 3-Clause License.  See accompanying file
+# Copyright.txt or https://cmake.org/licensing for details.
+
+include(FindPackageHandleStandardArgs)
+
+# ----------------------------------------------------------------------------------------#
+
+if(NOT ROCM_PATH AND NOT "$ENV{ROCM_PATH}" STREQUAL "")
+    set(ROCM_PATH "$ENV{ROCM_PATH}")
+endif()
+
+foreach(_DIR ${ROCM_PATH} /opt/rocm /opt/rocm/rocprofiler)
+    if(EXISTS ${_DIR})
+        get_filename_component(_ABS_DIR "${_DIR}" REALPATH)
+        list(APPEND _ROCM_ROCPROFILER_PATHS ${_ABS_DIR})
+    endif()
+endforeach()
+
+# ----------------------------------------------------------------------------------------#
+
+find_path(
+    rocprofiler_ROOT_DIR
+    NAMES include/rocprofiler/rocprofiler.h include/rocprofiler.h
+    HINTS ${_ROCM_ROCPROFILER_PATHS}
+    PATHS ${_ROCM_ROCPROFILER_PATHS}
+    PATH_SUFFIXES rocprofiler)
+
+mark_as_advanced(rocprofiler_ROOT_DIR)
+
+# ----------------------------------------------------------------------------------------#
+
+find_path(
+    rocprofiler_INCLUDE_DIR
+    NAMES rocprofiler.h
+    HINTS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
+    PATHS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
+    PATH_SUFFIXES include include/rocprofiler rocprofiler/include)
+
+mark_as_advanced(rocprofiler_INCLUDE_DIR)
+
+find_path(
+    rocprofiler_hsa_INCLUDE_DIR
+    NAMES hsa.h
+    HINTS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
+    PATHS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
+    PATH_SUFFIXES include include/hsa)
+
+mark_as_advanced(rocprofiler_hsa_INCLUDE_DIR)
+
+# ----------------------------------------------------------------------------------------#
+
+find_library(
+    rocprofiler_LIBRARY
+    NAMES rocprofiler64 rocprofiler
+    HINTS ${rocprofiler_ROOT_DIR}/rocprofiler ${rocprofiler_ROOT_DIR}
+          ${_ROCM_ROCPROFILER_PATHS}
+    PATHS ${rocprofiler_ROOT_DIR}/rocprofiler ${rocprofiler_ROOT_DIR}
+          ${_ROCM_ROCPROFILER_PATHS}
+    PATH_SUFFIXES lib lib64
+    NO_DEFAULT_PATH)
+
+find_library(
+    rocprofiler_hsa-runtime_LIBRARY
+    NAMES hsa-runtime64 hsa-runtime
+    HINTS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
+    PATHS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
+    PATH_SUFFIXES lib lib64)
+
+if(rocprofiler_LIBRARY)
+    get_filename_component(rocprofiler_LIBRARY_DIR "${rocprofiler_LIBRARY}" PATH CACHE)
+endif()
+
+mark_as_advanced(rocprofiler_LIBRARY rocprofiler_hsa-runtime_LIBRARY)
+
+# ----------------------------------------------------------------------------------------#
+find_package_handle_standard_args(
+    rocprofiler DEFAULT_MSG rocprofiler_ROOT_DIR rocprofiler_INCLUDE_DIR
+    rocprofiler_hsa_INCLUDE_DIR rocprofiler_LIBRARY rocprofiler_hsa-runtime_LIBRARY)
+
+# ------------------------------------------------------------------------------#
+
+if(rocprofiler_FOUND)
+    add_library(rocprofiler::rocprofiler INTERFACE IMPORTED)
+    add_library(rocprofiler::roctx INTERFACE IMPORTED)
+    set(rocprofiler_INCLUDE_DIRS ${rocprofiler_INCLUDE_DIR}
+                                 ${rocprofiler_hsa_INCLUDE_DIR})
+    set(rocprofiler_LIBRARIES ${rocprofiler_LIBRARY} ${rocprofiler_hsa-runtime_LIBRARY})
+    set(rocprofiler_LIBRARY_DIRS ${rocprofiler_LIBRARY_DIR})
+
+    target_include_directories(
+        rocprofiler::rocprofiler INTERFACE ${rocprofiler_INCLUDE_DIR}
+                                           ${rocprofiler_hsa_INCLUDE_DIR})
+
+    target_link_libraries(rocprofiler::rocprofiler INTERFACE ${rocprofiler_LIBRARIES})
+endif()
+# ------------------------------------------------------------------------------#
+
+unset(_ROCM_ROCPROFILER_PATHS)
+
+# ------------------------------------------------------------------------------#
diff --git a/projects/rocprofiler-systems/cmake/Modules/Findroctracer.cmake b/projects/rocprofiler-systems/cmake/Modules/Findroctracer.cmake
index 2ef0281909..8dbebf5682 100644
--- a/projects/rocprofiler-systems/cmake/Modules/Findroctracer.cmake
+++ b/projects/rocprofiler-systems/cmake/Modules/Findroctracer.cmake
@@ -5,15 +5,24 @@ include(FindPackageHandleStandardArgs)
 
 # ----------------------------------------------------------------------------------------#
 
-set(_ROCM_PATHS $ENV{ROCM_HOME} /opt/rocm /opt/rocm/roctracer)
+if(NOT ROCM_PATH AND NOT "$ENV{ROCM_PATH}" STREQUAL "")
+    set(ROCM_PATH "$ENV{ROCM_PATH}")
+endif()
+
+foreach(_DIR ${ROCM_PATH} /opt/rocm /opt/rocm/roctracer)
+    if(EXISTS ${_DIR})
+        get_filename_component(_ABS_DIR "${_DIR}" REALPATH)
+        list(APPEND _ROCM_ROCTRACER_PATHS ${_ABS_DIR})
+    endif()
+endforeach()
 
 # ----------------------------------------------------------------------------------------#
 
 find_path(
     roctracer_ROOT_DIR
-    NAMES include/roctracer.h
-    HINTS ${_ROCM_PATHS}
-    PATHS ${_ROCM_PATHS}
+    NAMES include/roctracer/roctracer.h include/roctracer.h
+    HINTS ${_ROCM_ROCTRACER_PATHS}
+    PATHS ${_ROCM_ROCTRACER_PATHS}
     PATH_SUFFIXES roctracer)
 
 mark_as_advanced(roctracer_ROOT_DIR)
@@ -23,17 +32,17 @@ mark_as_advanced(roctracer_ROOT_DIR)
 find_path(
     roctracer_INCLUDE_DIR
     NAMES roctracer.h
-    HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
-    PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
-    PATH_SUFFIXES roctracer/include include)
+    HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
+    PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
+    PATH_SUFFIXES include include/roctracer roctracer/include)
 
 mark_as_advanced(roctracer_INCLUDE_DIR)
 
 find_path(
     roctracer_hsa_INCLUDE_DIR
     NAMES hsa.h
-    HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
-    PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
+    HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
+    PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
     PATH_SUFFIXES include include/hsa)
 
 mark_as_advanced(roctracer_hsa_INCLUDE_DIR)
@@ -43,22 +52,22 @@ mark_as_advanced(roctracer_hsa_INCLUDE_DIR)
 find_library(
     roctracer_LIBRARY
     NAMES roctracer64 roctracer
-    HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
-    PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
+    HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
+    PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
     PATH_SUFFIXES lib lib64)
 
 find_library(
     roctracer_roctx_LIBRARY
     NAMES roctx64 roctx
-    HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
-    PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
+    HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
+    PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
     PATH_SUFFIXES lib lib64)
 
 find_library(
     roctracer_kfdwrapper_LIBRARY
     NAMES kfdwrapper64 kfdwrapper
-    HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
-    PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
+    HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
+    PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
     PATH_SUFFIXES lib lib64)
 
 find_package(hsakmt)
@@ -70,8 +79,8 @@ else()
     find_library(
         roctracer_hsakmt_LIBRARY
         NAMES hsakmt
-        HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
-        PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
+        HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
+        PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
         PATH_SUFFIXES lib lib64)
 endif()
 
@@ -122,6 +131,6 @@ endif()
 
 # ------------------------------------------------------------------------------#
 
-unset(_ROCM_PATHS)
+unset(_ROCM_ROCTRACER_PATHS)
 
 # ------------------------------------------------------------------------------#
diff --git a/projects/rocprofiler-systems/cmake/Packages.cmake b/projects/rocprofiler-systems/cmake/Packages.cmake
index 4179880100..e0e44f610f 100644
--- a/projects/rocprofiler-systems/cmake/Packages.cmake
+++ b/projects/rocprofiler-systems/cmake/Packages.cmake
@@ -16,6 +16,8 @@ omnitrace_add_interface_library(
 omnitrace_add_interface_library(omnitrace-hip "Provides flags and libraries for HIP")
 omnitrace_add_interface_library(omnitrace-roctracer
                                 "Provides flags and libraries for roctracer")
+omnitrace_add_interface_library(omnitrace-rocprofiler
+                                "Provides flags and libraries for rocprofiler")
 omnitrace_add_interface_library(omnitrace-rocm-smi
                                 "Provides flags and libraries for rocm-smi")
 omnitrace_add_interface_library(omnitrace-mpi "Provides MPI or MPI headers")
@@ -31,9 +33,15 @@ omnitrace_add_interface_library(omnitrace-compile-definitions "Compile definitio
 
 # libraries with relevant compile definitions
 set(OMNITRACE_EXTENSION_LIBRARIES
-    omnitrace::omnitrace-hip omnitrace::omnitrace-roctracer omnitrace::omnitrace-rocm-smi
-    omnitrace::omnitrace-mpi omnitrace::omnitrace-ptl omnitrace::omnitrace-ompt
-    omnitrace::omnitrace-papi omnitrace::omnitrace-perfetto)
+    omnitrace::omnitrace-hip
+    omnitrace::omnitrace-roctracer
+    omnitrace::omnitrace-rocprofiler
+    omnitrace::omnitrace-rocm-smi
+    omnitrace::omnitrace-mpi
+    omnitrace::omnitrace-ptl
+    omnitrace::omnitrace-ompt
+    omnitrace::omnitrace-papi
+    omnitrace::omnitrace-perfetto)
 
 target_include_directories(
     omnitrace-headers INTERFACE ${PROJECT_SOURCE_DIR}/source/lib/omnitrace
@@ -45,6 +53,14 @@ target_link_libraries(omnitrace-headers INTERFACE omnitrace::omnitrace-threading
 # ensure the env overrides the appending /opt/rocm later
 string(REPLACE ":" ";" CMAKE_PREFIX_PATH "$ENV{CMAKE_PREFIX_PATH};${CMAKE_PREFIX_PATH}")
 
+set(OMNITRACE_DEFAULT_ROCM_PATH
+    /opt/rocm
+    CACHE PATH "Default search path for ROCM")
+if(EXISTS ${OMNITRACE_DEFAULT_ROCM_PATH})
+    get_filename_component(OMNITRACE_DEFAULT_ROCM_PATH "${OMNITRACE_DEFAULT_ROCM_PATH}"
+                           REALPATH)
+endif()
+
 # ----------------------------------------------------------------------------------------#
 #
 # Threading
@@ -84,7 +100,7 @@ endforeach()
 # ----------------------------------------------------------------------------------------#
 
 if(OMNITRACE_USE_HIP)
-    list(APPEND CMAKE_PREFIX_PATH /opt/rocm)
+    list(APPEND CMAKE_PREFIX_PATH ${OMNITRACE_DEFAULT_ROCM_PATH})
     find_package(hip ${omnitrace_FIND_QUIETLY} REQUIRED)
     omnitrace_target_compile_definitions(omnitrace-hip INTERFACE OMNITRACE_USE_HIP)
     target_link_libraries(omnitrace-hip INTERFACE hip::host)
@@ -97,7 +113,7 @@ endif()
 # ----------------------------------------------------------------------------------------#
 
 if(OMNITRACE_USE_ROCTRACER)
-    list(APPEND CMAKE_PREFIX_PATH /opt/rocm)
+    list(APPEND CMAKE_PREFIX_PATH ${OMNITRACE_DEFAULT_ROCM_PATH})
     find_package(roctracer ${omnitrace_FIND_QUIETLY} REQUIRED)
     omnitrace_target_compile_definitions(omnitrace-roctracer
                                          INTERFACE OMNITRACE_USE_ROCTRACER)
@@ -106,6 +122,20 @@ if(OMNITRACE_USE_ROCTRACER)
     set(CMAKE_INSTALL_RPATH "${CMAKE_INSTALL_RPATH}:${roctracer_LIBRARY_DIRS}")
 endif()
 
+# ----------------------------------------------------------------------------------------#
+#
+# rocprofiler
+#
+# ----------------------------------------------------------------------------------------#
+if(OMNITRACE_USE_ROCPROFILER)
+    list(APPEND CMAKE_PREFIX_PATH ${OMNITRACE_DEFAULT_ROCM_PATH})
+    find_package(rocprofiler ${omnitrace_FIND_QUIETLY} REQUIRED)
+    omnitrace_target_compile_definitions(omnitrace-rocprofiler
+                                         INTERFACE OMNITRACE_USE_ROCPROFILER)
+    target_link_libraries(omnitrace-rocprofiler INTERFACE rocprofiler::rocprofiler)
+    set(CMAKE_INSTALL_RPATH "${CMAKE_INSTALL_RPATH}:${rocprofiler_LIBRARY_DIRS}")
+endif()
+
 # ----------------------------------------------------------------------------------------#
 #
 # rocm-smi
@@ -113,7 +143,7 @@ endif()
 # ----------------------------------------------------------------------------------------#
 
 if(OMNITRACE_USE_ROCM_SMI)
-    list(APPEND CMAKE_PREFIX_PATH /opt/rocm)
+    list(APPEND CMAKE_PREFIX_PATH ${OMNITRACE_DEFAULT_ROCM_PATH})
     find_package(rocm-smi ${omnitrace_FIND_QUIETLY} REQUIRED)
     omnitrace_target_compile_definitions(omnitrace-rocm-smi
                                          INTERFACE OMNITRACE_USE_ROCM_SMI)
@@ -545,7 +575,7 @@ omnitrace_checkout_git_submodule(
     RELATIVE_PATH external/timemory
     WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
     REPO_URL https://github.com/NERSC/timemory.git
-    REPO_BRANCH gpu-kernel-instrumentation)
+    REPO_BRANCH omnitrace)
 
 omnitrace_save_variables(
     BUILD_CONFIG VARIABLES BUILD_SHARED_LIBS BUILD_STATIC_LIBS
diff --git a/projects/rocprofiler-systems/examples/lulesh/CMakeLists.txt b/projects/rocprofiler-systems/examples/lulesh/CMakeLists.txt
index c4899d0a32..df4b2fa0c9 100644
--- a/projects/rocprofiler-systems/examples/lulesh/CMakeLists.txt
+++ b/projects/rocprofiler-systems/examples/lulesh/CMakeLists.txt
@@ -2,6 +2,8 @@ cmake_minimum_required(VERSION 3.16 FATAL_ERROR)
 
 project(lulesh LANGUAGES C CXX)
 
+set(CMAKE_BUILD_TYPE "RelWithDebInfo")
+
 list(INSERT CMAKE_MODULE_PATH 0 ${PROJECT_SOURCE_DIR}/cmake/Modules)
 
 option(LULESH_BUILD_KOKKOS "Build Kokkos from submodule" ON)
diff --git a/projects/rocprofiler-systems/examples/openmp/CMakeLists.txt b/projects/rocprofiler-systems/examples/openmp/CMakeLists.txt
index 534bc3698a..ea6a7a6d6c 100644
--- a/projects/rocprofiler-systems/examples/openmp/CMakeLists.txt
+++ b/projects/rocprofiler-systems/examples/openmp/CMakeLists.txt
@@ -23,10 +23,7 @@ else()
        AND LIBOMP_LIBRARY
        AND COMMAND omnitrace_custom_compilation)
         target_compile_options(openmp-common PUBLIC -W -Wall -fopenmp=libomp)
-        target_compile_options(openmp-cg PRIVATE -W -Wall -fopenmp=libomp)
-        target_link_libraries(openmp-cg PRIVATE ${LIBOMP_LIBRARY})
-        target_compile_options(openmp-lu PRIVATE -W -Wall -fopenmp=libomp)
-        target_link_libraries(openmp-lu PRIVATE ${LIBOMP_LIBRARY})
+        target_link_libraries(openmp-common PUBLIC ${LIBOMP_LIBRARY})
         omnitrace_custom_compilation(COMPILER ${CLANGXX_EXECUTABLE} TARGET openmp-common)
         omnitrace_custom_compilation(COMPILER ${CLANGXX_EXECUTABLE} TARGET openmp-cg)
         omnitrace_custom_compilation(COMPILER ${CLANGXX_EXECUTABLE} TARGET openmp-lu)
@@ -36,6 +33,9 @@ else()
     endif()
 endif()
 
+target_link_libraries(openmp-cg PRIVATE openmp-common)
+target_link_libraries(openmp-lu PRIVATE openmp-common)
+
 if(OMNITRACE_INSTALL_EXAMPLES)
     install(
         TARGETS openmp-cg openmp-lu
diff --git a/projects/rocprofiler-systems/examples/transpose/transpose.cpp b/projects/rocprofiler-systems/examples/transpose/transpose.cpp
index c19a0fd20e..3ca092603e 100644
--- a/projects/rocprofiler-systems/examples/transpose/transpose.cpp
+++ b/projects/rocprofiler-systems/examples/transpose/transpose.cpp
@@ -93,13 +93,18 @@ transpose_a(int* in, int* out, int M, int N)
     out[idx] = tile[threadIdx.x][threadIdx.y];
 }
 
+namespace
+{
+size_t nthreads = 2;
+size_t nitr     = 500;
+size_t nsync    = 10;
+}  // namespace
+
 void
 run(int rank, int tid, hipStream_t stream, int argc, char** argv)
 {
-    size_t       nitr  = 500;
-    size_t       nsync = 10;
-    unsigned int M     = 4960 * 2;
-    unsigned int N     = 4960 * 2;
+    unsigned int M = 4960 * 2;
+    unsigned int N = 4960 * 2;
     if(argc > 2) nitr = atoll(argv[2]);
     if(argc > 3) nsync = atoll(argv[3]);
 
@@ -186,29 +191,26 @@ do_a2a(int rank)
 int
 main(int argc, char** argv)
 {
-    int    rank     = 0;
-    int    size     = 1;
-    int    nthreads = 2;
-    int    nitr     = 5000;
-    size_t nsync    = 10;
+    int rank = 0;
+    int size = 1;
     for(int i = 1; i < argc; ++i)
     {
         auto _arg = std::string{ argv[i] };
         if(_arg == "?" || _arg == "-h" || _arg == "--help")
         {
             fprintf(stderr,
-                    "usage: transpose [NUM_THREADS (%i)] [NUM_ITERATION (%i)] "
+                    "usage: transpose [NUM_THREADS (%zu)] [NUM_ITERATION (%zu)] "
                     "[SYNC_EVERY_N_ITERATIONS (%zu)]\n",
                     nthreads, nitr, nsync);
             exit(EXIT_SUCCESS);
         }
     }
-    if(argc > 1) nthreads = atoi(argv[1]);
-    if(argc > 2) nitr = atoi(argv[2]);
+    if(argc > 1) nthreads = atoll(argv[1]);
+    if(argc > 2) nitr = atoll(argv[2]);
     if(argc > 3) nsync = atoll(argv[3]);
 
-    printf("[transpose] Number of threads: %i\n", nthreads);
-    printf("[transpose] Number of iterations: %i\n", nitr);
+    printf("[transpose] Number of threads: %zu\n", nthreads);
+    printf("[transpose] Number of iterations: %zu\n", nitr);
     printf("[transpose] Syncing every %zu iterations\n", nsync);
 
 #if defined(USE_MPI)
@@ -233,14 +235,14 @@ main(int argc, char** argv)
     {
         std::vector<std::thread> _threads{};
         std::vector<hipStream_t> _streams(nthreads);
-        for(int i = 0; i < nthreads; ++i)
+        for(size_t i = 0; i < nthreads; ++i)
             HIP_API_CALL(hipStreamCreate(&_streams.at(i)));
-        for(int i = 1; i < nthreads; ++i)
+        for(size_t i = 1; i < nthreads; ++i)
             _threads.emplace_back(run, rank, i, _streams.at(i), argc, argv);
         run(rank, 0, _streams.at(0), argc, argv);
         for(auto& itr : _threads)
             itr.join();
-        for(int i = 0; i < nthreads; ++i)
+        for(size_t i = 0; i < nthreads; ++i)
             HIP_API_CALL(hipStreamDestroy(_streams.at(i)));
     }
     HIP_API_CALL(hipDeviceSynchronize());
diff --git a/projects/rocprofiler-systems/external/timemory b/projects/rocprofiler-systems/external/timemory
index 33ef109c08..1c03bec283 160000
--- a/projects/rocprofiler-systems/external/timemory
+++ b/projects/rocprofiler-systems/external/timemory
@@ -1 +1 @@
-Subproject commit 33ef109c0839fa694ba05f39a74106694aaeec15
+Subproject commit 1c03bec2837823a3d574b83a688a2f61b515d862
diff --git a/projects/rocprofiler-systems/scripts/gperftools-cpu-profile.sh b/projects/rocprofiler-systems/scripts/gperftools-cpu-profile.sh
index 06624c4d4e..fabc56737f 100755
--- a/projects/rocprofiler-systems/scripts/gperftools-cpu-profile.sh
+++ b/projects/rocprofiler-systems/scripts/gperftools-cpu-profile.sh
@@ -1,16 +1,18 @@
 #!/bin/bash
 
 EXE=$(basename ${1})
-DIR=cpu.prof.${EXE}
+DIR=gperftools-output
 mkdir -p ${DIR}
 
-# gperf settings
+# gperftools settings
 : ${N:=0}
-: ${GPERF_PROFILE:=""}
-: ${GPERF_PROFILE_BASE:=${DIR}/gperf}
+: ${GPERFTOOLS_PROFILE:=""}
+: ${GPERFTOOLS_PROFILE_BASE:=${DIR}/prof.${EXE}}
 : ${MALLOCSTATS:=1}
 : ${CPUPROFILE_FREQUENCY:=250}
 : ${CPUPROFILE_REALTIME:=1}
+: ${PPROF:=$(which google-pprof)}
+: ${PPROF:=$(which pprof)}
 
 # rendering settings
 : ${INTERACTIVE:=0}
@@ -24,17 +26,18 @@ if [ "$(uname)" = "Darwin" ]; then
         IMG_FORMAT="jpg"
     fi
 fi
+
 run-verbose()
 {
-    echo "${@}" 1>&2
+    echo "### ${@} ###" 1>&2
     eval ${@}
 }
 
-while [ -z "${GPERF_PROFILE}" ]
+while [ -z "${GPERFTOOLS_PROFILE}" ]
 do
-    TEST_FILE=${GPERF_PROFILE_BASE}.${N}
+    TEST_FILE=${GPERFTOOLS_PROFILE_BASE}.${N}
     if [ ! -f "${TEST_FILE}" ]; then
-        GPERF_PROFILE=${TEST_FILE}
+        GPERFTOOLS_PROFILE=${TEST_FILE}
     fi
     N=$((${N}+1))
 done
@@ -43,10 +46,10 @@ export MALLOCSTATS
 export CPUPROFILE_FREQUENCY
 export CPUPROFILE_REALTIME
 
-echo -e "\n\t--> Outputting profile to '${GPERF_PROFILE}'...\n"
+echo -e "\n\t--> Outputting profile to '${GPERFTOOLS_PROFILE}'...\n"
 
 # remove profile file if unsucessful execution
-cleanup-failure() { set +v ; echo "failure"; rm -f ${GPERF_PROFILE}; exit 1; }
+cleanup-failure() { set +v ; echo "failure"; rm -f ${GPERFTOOLS_PROFILE}; exit 1; }
 trap cleanup-failure SIGHUP SIGINT SIGQUIT SIGILL SIGABRT SIGKILL
 
 ADD_LIBS()
@@ -73,66 +76,44 @@ ADD_PRELOAD()
     done
 }
 
+run-pprof()
+{
+    if [ -n "${PPROF}" ]; then
+        run-verbose ${PPROF} ${ADD_LIB_LIST} ${PPROF_ARGS} "${@}"
+    else
+        echo -e "neither google-pprof nor pprof were found!"
+        exit 1
+    fi
+}
+
 # configure pre-loading of profiler library
-PROJECT_LIBRARIES="$(find ${PWD} -type f | egrep 'libtimemory|libctimemory' | egrep -v '\.a$|\.dSYM' | egrep '\.so$|\.dylib$')"
-run-verbose ADD_LIBS ${PROJECT_LIBRARIES}
-if [ "$(uname)" = "Darwin" ]; then
-    run-verbose ADD_PRELOAD $(otool -L ${1} | egrep 'profiler' | awk '{print $1}')
-    LIBS=$(echo ${LIBS} | sed 's/^://g')
-    if [ -n "${LIBS}" ]; then
-        export DYLD_FORCE_FLAT_NAMESPACE=1
-        echo "DYLD_INSERT_LIBRARIES=${LIBS}"
-    fi
-else
-    run-verbose ADD_PRELOAD $(ldd ${1} | egrep 'profiler' | awk '{print $(NF-1)}') /usr/lib/x86_64-linux-gnu/libprofiler.so
-    LIBS=$(echo ${LIBS} | sed 's/^://g')
-    if [ -n "${LIB}" ]; then
-        echo "LD_PRELOAD=${LIBS}"
-    fi
-fi
+for i in $(find ${PWD} -type f | egrep 'libomnitrace' | egrep -v '\.a$' | egrep '\.so$') $(ldd ${1} | awk '{print $(NF-1)}')
+do
+    if [ -f "${i}" ]; then run-verbose ADD_LIBS "${i}"; fi
+done
+run-verbose ADD_PRELOAD $(ldd ${1} | egrep 'profiler' | awk '{print $(NF-1)}') /usr/lib/$(uname -m)-linux-gnu/libprofiler.so
+LIBS=$(echo ${LIBS} | sed 's/^://g')
 
 set -e
 # run the application
-if [ "$(uname)" = "Darwin" ]; then
-    eval DYLD_INSERT_LIBRARIES=${LIBS} CPUPROFILE_FREQUENCY=${CPUPROFILE_FREQUENCY} CPUPROFILE=${GPERF_PROFILE} $@ | tee ${GPERF_PROFILE}.log
-else
-    eval LD_PRELOAD=${LIBS} CPUPROFILE_FREQUENCY=${CPUPROFILE_FREQUENCY} CPUPROFILE=${GPERF_PROFILE} $@ | tee ${GPERF_PROFILE}.log
-fi
+LD_PRELOAD=${LIBS} CPUPROFILE_FREQUENCY=${CPUPROFILE_FREQUENCY} CPUPROFILE=${GPERFTOOLS_PROFILE} ${@} | tee ${GPERFTOOLS_PROFILE}.log
 set +e
 
-echo-dart-measurement()
-{
-    local _NAME=${1}
-    local _TYPE=${2}
-    local _PATH=${3}
-    echo "<DartMeasurementFile name=\"${_NAME}\" type=\"image/${_TYPE}\">${_PATH}</DartMeasurementFile>"
-}
-
 # generate the results
 EXT=so
-if [ "$(uname)" = "Darwin" ]; then EXT=dylib; fi
-if [ -f "${GPERF_PROFILE}" ]; then
-    : ${PPROF:=$(which google-pprof)}
-    : ${PPROF:=$(which pprof)}
-    if [ -n "${PPROF}" ]; then
-        run-verbose ${PPROF} --text ${ADD_LIB_LIST} ${PPROF_ARGS} ${1} ${GPERF_PROFILE} 1> ${GPERF_PROFILE}.txt.tmp
-        run-verbose cat ${GPERF_PROFILE}.txt.tmp | c++filt -n -t &> ${GPERF_PROFILE}.txt
-        run-verbose ${PPROF} --text --cum ${ADD_LIB_LIST} ${PPROF_ARGS} ${1} ${GPERF_PROFILE} 1> ${GPERF_PROFILE}.cum.txt.tmp
-        run-verbose cat ${GPERF_PROFILE}.cum.txt.tmp | c++filt -n -t &> ${GPERF_PROFILE}.cum.txt
-        rm -f *.txt.tmp
-        # if dot is available
-        if [ -n "$(which dot)" ]; then
-            run-verbose ${PPROF} --dot ${ADD_LIB_LIST} ${PPROF_ARGS} ${1} ${GPERF_PROFILE} 1> ${GPERF_PROFILE}.dot
-            run-verbose dot ${DOT_ARGS} -T${IMG_FORMAT} ${GPERF_PROFILE}.dot -o ${GPERF_PROFILE}.${IMG_FORMAT}
-            echo-dart-measurement ${GPERF_PROFILE}.${IMG_FORMAT} ${IMG_FORMAT} ${PWD}/${GPERF_PROFILE}.${IMG_FORMAT}
-        fi
-        if [ "${INTERACTIVE}" -gt 0 ]; then
-            run-verbose ${PPROF} ${ADD_LIB_LIST} ${PPROF_ARGS} ${1} ${GPERF_PROFILE}
-        fi
-    else
-        echo -e "google-pprof/pprof not found!"
+if [ -f "${GPERFTOOLS_PROFILE}" ]; then
+    run-pprof --text ${1} ${GPERFTOOLS_PROFILE} | c++filt -n -t 1> ${GPERFTOOLS_PROFILE}.txt
+    run-pprof ${PPROF} --text --cum ${1} ${GPERFTOOLS_PROFILE} | c++filt -n -t 1> ${GPERFTOOLS_PROFILE}.cum.txt
+    # if dot is available
+    if [ -n "$(which dot)" ]; then
+        run-pprof ${PPROF} --dot ${1} ${GPERFTOOLS_PROFILE} 1> ${GPERFTOOLS_PROFILE}.dot
+        run-verbose $(which dot) ${DOT_ARGS} -T${IMG_FORMAT} ${GPERFTOOLS_PROFILE}.dot -o ${GPERFTOOLS_PROFILE}.${IMG_FORMAT}
+    fi
+    if [ "${INTERACTIVE}" -gt 0 ]; then
+        run-pprof ${PPROF} ${1} ${GPERFTOOLS_PROFILE}
     fi
 else
-    echo -e "profile file \"${GPERF_PROFILE}\" not found!"
+    echo -e "profile file \"${GPERFTOOLS_PROFILE}\" not found!"
     ls -la
+    exit 1
 fi
diff --git a/projects/rocprofiler-systems/scripts/omnitrace_launch_compiler b/projects/rocprofiler-systems/scripts/omnitrace-launch-compiler
similarity index 100%
rename from projects/rocprofiler-systems/scripts/omnitrace_launch_compiler
rename to projects/rocprofiler-systems/scripts/omnitrace-launch-compiler
diff --git a/projects/rocprofiler-systems/source/bin/omnitrace-avail/avail.cpp b/projects/rocprofiler-systems/source/bin/omnitrace-avail/avail.cpp
index 939cad710a..0f3706a7ca 100644
--- a/projects/rocprofiler-systems/source/bin/omnitrace-avail/avail.cpp
+++ b/projects/rocprofiler-systems/source/bin/omnitrace-avail/avail.cpp
@@ -30,6 +30,8 @@
 #include "info_type.hpp"
 
 #include "library/config.hpp"
+#include "library/gpu.hpp"
+#include "library/rocprofiler.hpp"
 
 #include <timemory/components.hpp>
 #include <timemory/components/definition.hpp>
@@ -57,6 +59,12 @@
 #include <utility>
 #include <vector>
 
+#if defined(OMNITRACE_USE_HIP) && OMNITRACE_USE_HIP > 0
+#    include <hip/hip_runtime.h>
+#elif !defined(OMNITRACE_USE_HIP)
+#    define OMNITRACE_USE_HIP 0
+#endif
+
 #if defined(TIMEMORY_UNIX)
 #    include <sys/ioctl.h>  // ioctl() and TIOCGWINSZ
 #    include <unistd.h>     // for STDOUT_FILENO
@@ -98,6 +106,8 @@ void
 write_hw_counter_info(std::ostream&, const array_t<bool, N>& = {},
                       const array_t<bool, N>& = {}, const array_t<string_t, N>& = {});
 
+int gpu_count = 0;
+
 //--------------------------------------------------------------------------------------//
 
 int
@@ -125,6 +135,8 @@ main(int argc, char** argv)
             }
         }
     }
+    _category_options.emplace("hw_counters::CPU");
+    _category_options.emplace("hw_counters::GPU");
 
     array_t<bool, TOTAL> options    = { false, false, false, false, false, false, false };
     array_t<string_t, TOTAL> fields = {};
@@ -164,6 +176,11 @@ main(int argc, char** argv)
     parser.add_argument({ "--debug" }, "Enable debug messages")
         .max_count(1)
         .action([](parser_t& p) { debug_msg = p.get<bool>("debug"); });
+    parser.add_argument({ "--verbose" }, "Enable informational messages")
+        .max_count(1)
+        .action([](parser_t& p) {
+            verbose_level = (p.get_count("verbose") == 0) ? 1 : p.get<int>("verbose");
+        });
     parser.add_argument({ "-a", "--all" }, "Print all available info")
         .max_count(1)
         .action([&](parser_t& p) {
@@ -297,7 +314,7 @@ main(int argc, char** argv)
                       "Display the output filename for the component")
         .max_count(1);
     parser
-        .add_argument({ "--categories" },
+        .add_argument({ "-c", "--categories" },
                       "Display the category information (use --list-categories to see "
                       "the available categories)")
         .dtype("string")
@@ -399,6 +416,20 @@ main(int argc, char** argv)
         return EXIT_FAILURE;
     }
 
+#if OMNITRACE_USE_HIP > 0
+    // initialize HIP and call rocm_metrics() which add choices to OMNITRACE_ROCM_EVENTS
+    // setting
+    auto _status = hipGetDeviceCount(&gpu_count);
+    if(gpu_count > 0 && _status == hipSuccess)
+    {
+        (void) omnitrace::rocprofiler::rocm_metrics();
+    }
+    else
+    {
+        verbprintf(0, "No HIP devices found. GPU HW counters will not be available\n");
+    }
+#endif
+
     auto _parser_set_if_exists = [&parser](auto& _var, const std::string& _opt) {
         using Tp = decay_t<decltype(_var)>;
         if(parser.exists(_opt)) _var = parser.get<Tp>(_opt);
@@ -923,10 +954,13 @@ write_hw_counter_info(std::ostream& os, const array_t<bool, N>& options,
     static_assert(N >= num_hw_counter_options,
                   "Error! Too few hw counter options + fields");
 
-    using width_type = array_t<int64_t, N>;
-    using width_bool = array_t<bool, N>;
+    using width_type       = array_t<int64_t, N>;
+    using width_bool       = array_t<bool, N>;
+    using hwcounter_info_t = std::vector<tim::hardware_counters::info>;
 
     auto _papi_events = tim::papi::available_events_info();
+    auto _rocm_events =
+        (gpu_count > 0) ? omnitrace::rocprofiler::rocm_metrics() : hwcounter_info_t{};
 
     auto _process_counters = [](auto& _events, int32_t _offset) {
         for(auto& itr : _events)
@@ -939,13 +973,25 @@ write_hw_counter_info(std::ostream& os, const array_t<bool, N>& options,
 
     int32_t _offset = 0;
     _offset += _process_counters(_papi_events, _offset);
+    _offset += _process_counters(_rocm_events, _offset);
 
-    using hwcounter_info_t             = std::vector<tim::hardware_counters::info>;
-    auto                 fields        = std::vector<hwcounter_info_t>{ _papi_events };
-    auto                 subcategories = std::vector<std::string>{ "CPU", "GPU", "" };
-    array_t<string_t, N> _labels       = { "HARDWARE COUNTER", "AVAILABLE", "SUMMARY",
+    auto fields        = std::vector<hwcounter_info_t>{ _papi_events, _rocm_events };
+    auto subcategories = std::vector<std::string>{ "CPU", "GPU", "" };
+    array_t<string_t, N> _labels = { "HARDWARE COUNTER", "AVAILABLE", "SUMMARY",
                                      "DESCRIPTION" };
-    array_t<bool, N>     _center       = { false, true, false, false };
+    array_t<bool, N>     _center = { false, true, false, false };
+
+    for(size_t i = 0; i < subcategories.size(); ++i)
+    {
+        if(i >= fields.size()) break;
+        if(!category_view.empty() && category_view.count(subcategories.at(i)) == 0 &&
+           category_view.count(std::string{ "hw_counters::" } + subcategories.at(i)) == 0)
+            fields.at(i).clear();
+        if(!is_category_selected(subcategories.at(i)) &&
+           !is_category_selected(std::string{ "hw_counters::" } + subcategories.at(i)))
+            fields.at(i).clear();
+        if(fields.at(i).empty()) subcategories.at(i).clear();
+    }
 
     width_type _widths;
     width_bool _wusing;
@@ -986,14 +1032,15 @@ write_hw_counter_info(std::ostream& os, const array_t<bool, N>& options,
     os << "\n" << banner(_widths, _wusing, '-');
 
     size_t nitr = 0;
+    size_t nout = 0;
     for(const auto& fitr : fields)
     {
         auto idx = nitr++;
 
         if(idx < subcategories.size())
         {
-            if(!markdown && idx != 0) os << banner(_widths, _wusing, '-');
-            if(subcategories.at(idx).length() > 0)
+            if(!markdown && nout != 0) os << banner(_widths, _wusing, '-');
+            if(!subcategories.at(idx).empty())
             {
                 os << global_delim;
                 if(options[0])
@@ -1008,6 +1055,7 @@ write_hw_counter_info(std::ostream& os, const array_t<bool, N>& options,
                 }
                 os << "\n";
                 if(!markdown) os << banner(_widths, _wusing, '-');
+                ++nout;
             }
         }
         else
diff --git a/projects/rocprofiler-systems/source/bin/omnitrace-avail/common.cpp b/projects/rocprofiler-systems/source/bin/omnitrace-avail/common.cpp
index 61bcb0f06b..e61df01b7a 100644
--- a/projects/rocprofiler-systems/source/bin/omnitrace-avail/common.cpp
+++ b/projects/rocprofiler-systems/source/bin/omnitrace-avail/common.cpp
@@ -38,7 +38,6 @@ bool              alphabetical        = false;
 bool              available_only      = false;
 bool              all_info            = false;
 bool              force_brief         = false;
-bool              debug_msg           = false;
 bool              case_insensitive    = false;
 bool              regex_hl            = false;
 bool              expand_keys         = false;
@@ -50,7 +49,11 @@ int32_t           padding             = 4;
 str_vec_t         regex_keys          = {};
 str_vec_t         category_regex_keys = {};
 str_set_t         category_view       = {};
-std::stringstream lerr                = {};
+std::stringstream lerr{};
+
+bool    debug_msg = tim::get_env<bool>("OMNITRACE_DEBUG_AVAIL", settings::debug());
+int32_t verbose_level =
+    tim::get_env<int32_t>("OMNITRACE_VERBOSE_AVAIL", settings::verbose());
 
 // explicit setting names to exclude
 std::set<std::string> settings_exclude = {
@@ -321,7 +324,8 @@ process_categories(parser_t& p, const str_set_t& _category_options)
 
         if(_category_options.count(itr) == 0)
         {
-            if(!_is_shorthand("component") && !_is_shorthand("settings"))
+            if(!_is_shorthand("component") && !_is_shorthand("settings") &&
+               !_is_shorthand("hw_counters"))
                 throw std::runtime_error(
                     itr + " is not a valid category. Use --list-categories to view "
                           "valid categories");
diff --git a/projects/rocprofiler-systems/source/bin/omnitrace-avail/common.hpp b/projects/rocprofiler-systems/source/bin/omnitrace-avail/common.hpp
index b62130bc4e..ca31b28d53 100644
--- a/projects/rocprofiler-systems/source/bin/omnitrace-avail/common.hpp
+++ b/projects/rocprofiler-systems/source/bin/omnitrace-avail/common.hpp
@@ -99,6 +99,7 @@ extern int32_t           max_width;
 extern int32_t           num_cols;
 extern int32_t           min_width;
 extern int32_t           padding;
+extern int32_t           verbose_level;
 extern str_vec_t         regex_keys;
 extern str_vec_t         category_regex_keys;
 extern str_set_t         category_view;
@@ -143,3 +144,36 @@ remove(std::string inp, const std::set<std::string>& entries);
 
 bool
 file_exists(const std::string&);
+
+// control debug printf statements
+#define errprintf(LEVEL, ...)                                                            \
+    {                                                                                    \
+        if(werror || LEVEL < 0)                                                          \
+        {                                                                                \
+            if(debug_msg || verbose_level >= LEVEL)                                      \
+                fprintf(stderr, "[omnitrace][avail] Error! " __VA_ARGS__);               \
+            char _buff[FUNCNAMELEN];                                                     \
+            sprintf(_buff, "[omnitrace][avail] Error! " __VA_ARGS__);                    \
+            throw std::runtime_error(std::string{ _buff });                              \
+        }                                                                                \
+        else                                                                             \
+        {                                                                                \
+            if(debug_msg || verbose_level >= LEVEL)                                      \
+                fprintf(stderr, "[omnitrace][avail] Warning! " __VA_ARGS__);             \
+        }                                                                                \
+        fflush(stderr);                                                                  \
+    }
+
+// control verbose printf statements
+#define verbprintf(LEVEL, ...)                                                           \
+    {                                                                                    \
+        if(debug_msg || verbose_level >= LEVEL)                                          \
+            fprintf(stderr, "[omnitrace][avail] " __VA_ARGS__);                          \
+        fflush(stderr);                                                                  \
+    }
+
+#define verbprintf_bare(LEVEL, ...)                                                      \
+    {                                                                                    \
+        if(debug_msg || verbose_level >= LEVEL) fprintf(stderr, __VA_ARGS__);            \
+        fflush(stderr);                                                                  \
+    }
diff --git a/projects/rocprofiler-systems/source/bin/omnitrace-avail/generate_config.cpp b/projects/rocprofiler-systems/source/bin/omnitrace-avail/generate_config.cpp
index e887905762..a07a528107 100644
--- a/projects/rocprofiler-systems/source/bin/omnitrace-avail/generate_config.cpp
+++ b/projects/rocprofiler-systems/source/bin/omnitrace-avail/generate_config.cpp
@@ -230,7 +230,7 @@ generate_config(std::string _config_file, const std::set<std::string>& _config_f
         _ar->startNode();
         (*_ar)(cereal::make_nvp("version", std::string{ OMNITRACE_VERSION_STRING }));
         (*_ar)(cereal::make_nvp("date", tim::get_local_datetime("%F_%H.%M", &_time)));
-        settings::serialize_settings(*_ar);
+        settings::serialize_settings(*_ar, *_settings);
         _ar->finishNode();
     };
 
@@ -311,6 +311,7 @@ generate_config(std::string _config_file, const std::set<std::string>& _config_f
             });
         else
         {
+            _settings->ordering();
             std::sort(_data.begin(), _data.end(), [](auto _lhs, auto _rhs) {
                 auto _lomni = _lhs->get_categories().count("omnitrace") > 0;
                 auto _romni = _rhs->get_categories().count("omnitrace") > 0;
@@ -419,7 +420,7 @@ update_choices(std::shared_ptr<settings> _settings)
 {
     std::vector<info_type> _info = get_component_info<TIMEMORY_NATIVE_COMPONENTS_END>();
 
-    if(settings::verbose() >= 2 || settings::debug())
+    if(_settings->get_verbose() >= 2 || _settings->get_debug())
         printf("[omnitrace-avail] # of component found: %zu\n", _info.size());
 
     _info.erase(std::remove_if(_info.begin(), _info.end(),
@@ -442,7 +443,7 @@ update_choices(std::shared_ptr<settings> _settings)
     _component_choices.reserve(_info.size());
     for(const auto& itr : _info)
         _component_choices.emplace_back(itr.id_type());
-    if(settings::verbose() >= 2 || settings::debug())
+    if(_settings->get_verbose() >= 2 || _settings->get_debug())
         printf("[omnitrace-avail] # of component choices: %zu\n",
                _component_choices.size());
     _settings->find("OMNITRACE_TIMEMORY_COMPONENTS")
diff --git a/projects/rocprofiler-systems/source/bin/omnitrace-avail/info_type.cpp b/projects/rocprofiler-systems/source/bin/omnitrace-avail/info_type.cpp
index a478e71462..07d6830a2f 100644
--- a/projects/rocprofiler-systems/source/bin/omnitrace-avail/info_type.cpp
+++ b/projects/rocprofiler-systems/source/bin/omnitrace-avail/info_type.cpp
@@ -30,6 +30,7 @@
 #include "library/components/mpi_gotcha.hpp"
 #include "library/components/omnitrace.hpp"
 #include "library/components/pthread_gotcha.hpp"
+#include "library/components/rocprofiler.hpp"
 #include "library/components/roctracer.hpp"
 #include "library/components/user_region.hpp"
 
diff --git a/projects/rocprofiler-systems/source/lib/common/defines.h b/projects/rocprofiler-systems/source/lib/common/defines.h
index a3705a911b..80b6945d9c 100644
--- a/projects/rocprofiler-systems/source/lib/common/defines.h
+++ b/projects/rocprofiler-systems/source/lib/common/defines.h
@@ -31,6 +31,7 @@
 #define OMNITRACE_HOT               OMNITRACE_ATTRIBUTE(hot)
 #define OMNITRACE_CONST             OMNITRACE_ATTRIBUTE(const)
 #define OMNITRACE_PURE              OMNITRACE_ATTRIBUTE(pure)
+#define OMNITRACE_WEAK              OMNITRACE_ATTRIBUTE(weak)
 #define OMNITRACE_PACKED            OMNITRACE_ATTRIBUTE(__packed__)
 #define OMNITRACE_PACKED_ALIGN(VAL) OMNITRACE_PACKED OMNITRACE_ATTRIBUTE(__aligned__(VAL))
 
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace-dl/dl.cpp b/projects/rocprofiler-systems/source/lib/omnitrace-dl/dl.cpp
index ac7a90429d..8f191fa354 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace-dl/dl.cpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace-dl/dl.cpp
@@ -157,6 +157,25 @@ struct OMNITRACE_HIDDEN_API indirect
         setenv("HSA_TOOLS_LIB", _omni_hsa_lib.c_str(), 0);
 #endif
 
+#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
+        auto _rocm_path    = get_env("ROCM_PATH", "/opt/rocm");
+        auto _rocp_metrics = common::join('/', _rocm_path, "rocprofiler/lib/metrics.xml");
+        setenv("HSA_TOOLS_LIB", m_omnilib.c_str(), 0);
+        setenv("ROCP_TOOL_LIB", m_omnilib.c_str(), 0);
+        setenv("ROCPROFILER_LOG", "1", 0);
+        setenv("ROCP_HSA_INTERCEPT", "1", 0);
+        setenv("ROCP_METRICS", _rocp_metrics.c_str(), 0);
+        setenv("HSA_TOOLS_REPORT_LOAD_FAILURE", "1", 0);
+        if(getenv("ROCM_PATH"))
+        {
+            setenv("OMNITRACE_ROCPROFILER_LIBRARY",
+                   common::join('/', getenv("ROCM_PATH"),
+                                "rocprofiler/lib/librocprofiler64.so")
+                       .c_str(),
+                   0);
+        }
+#endif
+
 #if OMNITRACE_USE_OMPT > 0
         if(get_env("OMNITRACE_USE_OMPT", true))
         {
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/CMakeLists.txt b/projects/rocprofiler-systems/source/lib/omnitrace/CMakeLists.txt
index b31267a5cc..b9fcaa78e3 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/CMakeLists.txt
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/CMakeLists.txt
@@ -32,6 +32,7 @@ target_link_libraries(
         $<BUILD_INTERFACE:omnitrace::omnitrace-ptl>
         $<BUILD_INTERFACE:omnitrace::omnitrace-hip>
         $<BUILD_INTERFACE:omnitrace::omnitrace-roctracer>
+        $<BUILD_INTERFACE:omnitrace::omnitrace-rocprofiler>
         $<BUILD_INTERFACE:omnitrace::omnitrace-rocm-smi>
         $<BUILD_INTERFACE:$<IF:$<BOOL:${OMNITRACE_BUILD_LTO}>,omnitrace::omnitrace-lto,>>
         $<BUILD_INTERFACE:$<IF:$<BOOL:${OMNITRACE_BUILD_STATIC_LIBGCC}>,omnitrace::omnitrace-static-libgcc,>>
@@ -98,6 +99,9 @@ set(library_headers
     ${CMAKE_CURRENT_LIST_DIR}/library/perfetto.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/process_sampler.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/ptl.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/library/rocm.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/library/roctracer.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/runtime.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/sampling.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/state.hpp
@@ -113,8 +117,8 @@ set(library_headers
     ${CMAKE_CURRENT_LIST_DIR}/library/components/mpi_gotcha.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/components/omnitrace.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/components/rocm_smi.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/library/components/rocprofiler.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/components/roctracer.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/library/components/roctracer_callbacks.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/components/pthread_gotcha.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/components/pthread_create_gotcha.hpp
     ${CMAKE_CURRENT_LIST_DIR}/library/components/pthread_mutex_gotcha.hpp
@@ -122,11 +126,28 @@ set(library_headers
 
 target_sources(omnitrace-object-library PRIVATE ${library_sources} ${library_headers})
 
+if(OMNITRACE_USE_ROCTRACER OR OMNITRACE_USE_ROCPROFILER)
+    target_sources(
+        omnitrace-object-library
+        PRIVATE ${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler.cpp
+                ${CMAKE_CURRENT_LIST_DIR}/library/rocm.cpp
+                ${CMAKE_CURRENT_LIST_DIR}/library/components/rocprofiler.cpp)
+endif()
+
 if(OMNITRACE_USE_ROCTRACER)
     target_sources(
         omnitrace-object-library
         PRIVATE ${CMAKE_CURRENT_LIST_DIR}/library/components/roctracer.cpp
-                ${CMAKE_CURRENT_LIST_DIR}/library/components/roctracer_callbacks.cpp)
+                ${CMAKE_CURRENT_LIST_DIR}/library/roctracer.cpp)
+endif()
+
+if(OMNITRACE_USE_ROCPROFILER)
+    target_sources(
+        omnitrace-object-library
+        PRIVATE ${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler.cpp
+                ${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler.hpp
+                ${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler/hsa_rsrc_factory.hpp
+                ${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler/hsa_rsrc_factory.cpp)
 endif()
 
 if(OMNITRACE_USE_ROCM_SMI)
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library.cpp
index ae0d0a7686..5c2a7d7043 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library.cpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library.cpp
@@ -29,6 +29,7 @@
 #include "library/components/pthread_create_gotcha.hpp"
 #include "library/components/pthread_gotcha.hpp"
 #include "library/components/pthread_mutex_gotcha.hpp"
+#include "library/components/rocprofiler.hpp"
 #include "library/config.hpp"
 #include "library/coverage.hpp"
 #include "library/critical_trace.hpp"
@@ -38,6 +39,7 @@
 #include "library/ompt.hpp"
 #include "library/process_sampler.hpp"
 #include "library/ptl.hpp"
+#include "library/rocprofiler.hpp"
 #include "library/sampling.hpp"
 #include "library/thread_data.hpp"
 #include "library/timemory.hpp"
@@ -91,6 +93,27 @@ ensure_finalization(bool _static_init = false)
         tim::set_env("HSA_ENABLE_INTERRUPT", "0", 0);
 #if defined(OMNITRACE_USE_ROCTRACER) && OMNITRACE_USE_ROCTRACER > 0
         tim::set_env("HSA_TOOLS_LIB", "libomnitrace.so", 0);
+#endif
+#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
+        auto _default_rocm_path =
+            JOIN("", "/opt/rocm-", OMNITRACE_HIP_VERSION_MAJOR, '.',
+                 OMNITRACE_HIP_VERSION_MINOR, '.', OMNITRACE_HIP_VERSION_PATCH);
+        auto _rocm_path    = tim::get_env("OMNITRACE_ROCM_PATH",
+                                       tim::get_env("ROCM_PATH", _default_rocm_path));
+        auto _rocp_metrics = JOIN('/', _rocm_path, "rocprofiler/lib/metrics.xml");
+        tim::set_env("HSA_TOOLS_LIB", "libomnitrace.so", 0);
+        tim::set_env("ROCP_TOOL_LIB", "libomnitrace.so", 0);
+        tim::set_env("ROCPROFILER_LOG", "1", 0);
+        tim::set_env("ROCP_HSA_INTERCEPT", "1", 0);
+        tim::set_env("ROCP_METRICS", _rocp_metrics, 0);
+        tim::set_env("HSA_TOOLS_REPORT_LOAD_FAILURE", "1", 0);
+        if(getenv("ROCM_PATH"))
+        {
+            tim::set_env("OMNITRACE_ROCPROFILER_LIBRARY",
+                         JOIN('/', tim::get_env<std::string>("ROCM_PATH"),
+                              "rocprofiler/lib/librocprofiler64.so"),
+                         0);
+        }
 #endif
     }
     return scope::destructor{ []() { omnitrace_finalize_hidden(); } };
@@ -872,6 +895,13 @@ omnitrace_finalize_hidden(void)
         tasking::join();
     }
 
+    if(get_use_rocprofiler())
+    {
+        OMNITRACE_VERBOSE_F(1, "Shutting down rocprofiler...\n");
+        rocprofiler::post_process();
+        rocprofiler::rocm_cleanup();
+    }
+
     if(dmp::rank() == 0) fprintf(stderr, "\n");
 
     OMNITRACE_DEBUG_F("Stopping main bundle...\n");
@@ -1108,8 +1138,7 @@ omnitrace_finalize_hidden(void)
                       _push_count, "vs. popped:", _pop_count)
             .c_str());
 
-    OMNITRACE_DEBUG_F("Disabling signal handling...\n");
-    tim::disable_signal_detection();
+    config::finalize();
 
     OMNITRACE_VERBOSE_F(0, "Finalized\n");
 }
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/common.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/common.hpp
index d87e45b11b..64b13e818b 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/common.hpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/common.hpp
@@ -28,6 +28,8 @@
 #include <timemory/api.hpp>
 #include <timemory/backends/dmp.hpp>
 #include <timemory/backends/process.hpp>
+#include <timemory/utility/demangle.hpp>
+#include <timemory/utility/filepath.hpp>
 
 #include <cassert>
 #include <cstdint>
@@ -48,8 +50,13 @@ TIMEMORY_DEFINE_NS_API(api, rocm_smi)
 
 namespace omnitrace
 {
-namespace api      = tim::api;       // NOLINT
-namespace category = tim::category;  // NOLINT
+namespace api      = ::tim::api;       // NOLINT
+namespace category = ::tim::category;  // NOLINT
+namespace filepath = ::tim::filepath;  // NOLINT
+
+using ::tim::demangle;      // NOLINT
+using ::tim::get_env;       // NOLINT
+using ::tim::try_demangle;  // NOLINT
 }  // namespace omnitrace
 
 // same sort of functionality as python's " ".join([...])
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/fwd.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/fwd.hpp
index c39113940f..836370eff3 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/fwd.hpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/fwd.hpp
@@ -39,6 +39,7 @@ TIMEMORY_DEFINE_NS_API(project, omnitrace)
 TIMEMORY_DEFINE_NS_API(category, process_sampling)
 
 TIMEMORY_DECLARE_COMPONENT(roctracer)
+TIMEMORY_DECLARE_COMPONENT(rocprofiler)
 
 /// \struct tim::trait::name
 /// \brief provides a constexpr string in ::value
@@ -129,6 +130,7 @@ using sampling_gpu_temp   = data_tracker<double, backtrace_gpu_temp>;
 using sampling_gpu_power  = data_tracker<double, backtrace_gpu_power>;
 using sampling_gpu_memory = data_tracker<double, backtrace_gpu_memory>;
 using roctracer           = tim::component::roctracer;
+using rocprofiler         = tim::component::rocprofiler;
 
 template <typename ApiT, typename StartFuncT = default_functor_t,
           typename StopFuncT = default_functor_t>
@@ -140,6 +142,10 @@ struct functors;
 TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, component::roctracer, false_type)
 #endif
 
+#if !defined(OMNITRACE_USE_ROCPROFILER)
+TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, component::rocprofiler, false_type)
+#endif
+
 #if !defined(TIMEMORY_USE_LIBUNWIND)
 TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, omnitrace::api::sampling, false_type)
 TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, omnitrace::component::backtrace, false_type)
@@ -169,6 +175,9 @@ TIMEMORY_SET_COMPONENT_API(omnitrace::component::user_region, project::omnitrace
 TIMEMORY_SET_COMPONENT_API(omnitrace::component::roctracer, project::omnitrace,
                            tpls::rocm, device::gpu, os::supports_linux,
                            category::external)
+TIMEMORY_SET_COMPONENT_API(omnitrace::component::rocprofiler, project::omnitrace,
+                           tpls::rocm, device::gpu, os::supports_linux,
+                           category::external, category::hardware_counter)
 TIMEMORY_SET_COMPONENT_API(omnitrace::component::sampling_wall_clock, project::omnitrace,
                            category::timing, os::supports_unix, category::sampling,
                            category::interrupt_sampling)
@@ -198,6 +207,8 @@ TIMEMORY_PROPERTY_SPECIALIZATION(omnitrace::component::user_region, OMNITRACE_US
                                  "user_region", "omnitrace_user_region")
 TIMEMORY_PROPERTY_SPECIALIZATION(omnitrace::component::roctracer, OMNITRACE_ROCTRACER,
                                  "roctracer", "omnitrace_roctracer")
+TIMEMORY_PROPERTY_SPECIALIZATION(omnitrace::component::rocprofiler, OMNITRACE_ROCPROFILER,
+                                 "rocprofiler", "omnitrace_rocprofiler")
 TIMEMORY_PROPERTY_SPECIALIZATION(omnitrace::component::sampling_wall_clock,
                                  OMNITRACE_SAMPLING_WALL_CLOCK, "sampling_wall_clock", "")
 TIMEMORY_PROPERTY_SPECIALIZATION(omnitrace::component::sampling_cpu_clock,
@@ -226,6 +237,8 @@ TIMEMORY_METADATA_SPECIALIZATION(
     "Used by OMPT")
 TIMEMORY_METADATA_SPECIALIZATION(omnitrace::component::roctracer, "roctracer",
                                  "High-precision ROCm API and kernel tracing", "")
+TIMEMORY_METADATA_SPECIALIZATION(omnitrace::component::rocprofiler, "rocprofiler",
+                                 "ROCm kernel hardware counters", "")
 TIMEMORY_METADATA_SPECIALIZATION(omnitrace::component::sampling_wall_clock,
                                  "sampling_wall_clock", "Wall-clock timing",
                                  "Derived from statistical sampling")
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/pthread_create_gotcha.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/pthread_create_gotcha.cpp
index 0274742a9e..db69d0017b 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/pthread_create_gotcha.cpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/pthread_create_gotcha.cpp
@@ -154,7 +154,8 @@ pthread_create_gotcha::wrapper::operator()() const
         if(!thread_bundle_data_t::instances().at(_tid))
         {
             thread_data<omnitrace_thread_bundle_t>::construct(
-                TIMEMORY_JOIN("", get_exe_name(), "/thread-", threading::get_id()),
+                TIMEMORY_JOIN('/', "omnitrace/process", process::get_id(), "thread",
+                              threading::get_id()),
                 quirk::config<quirk::auto_start>{});
             thread_bundle_data_t::instances().at(_tid)->start();
         }
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/pthread_mutex_gotcha.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/pthread_mutex_gotcha.cpp
index 872eb67896..6c87ee3de5 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/pthread_mutex_gotcha.cpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/pthread_mutex_gotcha.cpp
@@ -29,8 +29,8 @@
 #include "library/runtime.hpp"
 #include "library/sampling.hpp"
 #include "library/utility.hpp"
-#include "timemory/backends/threading.hpp"
 
+#include <timemory/backends/threading.hpp>
 #include <timemory/utility/signals.hpp>
 #include <timemory/utility/types.hpp>
 
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/rocprofiler.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/rocprofiler.cpp
new file mode 100644
index 0000000000..47c33eda91
--- /dev/null
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/rocprofiler.cpp
@@ -0,0 +1,344 @@
+// MIT License
+//
+// Copyright (c) 2022 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 "library/components/rocprofiler.hpp"
+#include "library/common.hpp"
+#include "library/components/pthread_create_gotcha.hpp"
+#include "library/components/pthread_gotcha.hpp"
+#include "library/config.hpp"
+#include "library/debug.hpp"
+#include "library/defines.hpp"
+#include "library/dynamic_library.hpp"
+#include "library/perfetto.hpp"
+#include "library/redirect.hpp"
+#include "library/rocprofiler.hpp"
+#include "library/sampling.hpp"
+#include "library/thread_data.hpp"
+
+#include <timemory/storage/types.hpp>
+#include <timemory/utility/types.hpp>
+#include <timemory/variadic/functional.hpp>
+#include <timemory/variadic/lightweight_tuple.hpp>
+
+#include <rocprofiler.h>
+
+#include <cstdint>
+#include <string_view>
+#include <type_traits>
+
+TIMEMORY_STATISTICS_TYPE(component::rocm_data_tracker, component::rocm_feature_value)
+TIMEMORY_DEFINE_CONCRETE_TRAIT(report_units, component::rocm_data_tracker, false_type)
+
+namespace tim
+{
+namespace component
+{
+namespace
+{
+auto&
+rocprofiler_activity_count()
+{
+    static std::atomic<int64_t> _v{ 0 };
+    return _v;
+}
+}  // namespace
+
+omnitrace::unique_ptr_t<rocm_data_t>&
+rocm_data(int64_t _tid)
+{
+    using thread_data_t = omnitrace::thread_data<rocm_data_t, rocm_event>;
+    static auto& _v     = thread_data_t::instances(thread_data_t::construct_on_init{});
+    return _v.at(_tid);
+}
+
+rocm_event::rocm_event(uint32_t _dev, uint32_t _thr, uint32_t _queue,
+                       std::string _event_name, rocm_metric_type _begin,
+                       rocm_metric_type _end, uint32_t _feature_count, void* _features_v)
+: device_id{ _dev }
+, thread_id{ _thr }
+, queue_id{ _queue }
+, entry{ _begin }
+, exit{ _end }
+, name(std::move(_event_name))
+{
+    feature_values.reserve(_feature_count);
+    feature_names.reserve(_feature_count);
+    auto* _features = static_cast<rocprofiler_feature_t*>(_features_v);
+    for(uint32_t i = 0; i < _feature_count; ++i)
+    {
+        const rocprofiler_feature_t* p = &_features[i];
+        feature_names.emplace_back(p->name);
+        switch(p->data.kind)
+        {
+            // Output metrics results
+            case ROCPROFILER_DATA_KIND_UNINIT: break;
+            case ROCPROFILER_DATA_KIND_BYTES:
+                feature_values.emplace_back(
+                    rocm_feature_value{ p->data.result_bytes.size });
+                break;
+            case ROCPROFILER_DATA_KIND_INT32:
+                feature_values.emplace_back(rocm_feature_value{ p->data.result_int32 });
+                break;
+            case ROCPROFILER_DATA_KIND_FLOAT:
+                feature_values.emplace_back(rocm_feature_value{ p->data.result_float });
+                break;
+            case ROCPROFILER_DATA_KIND_DOUBLE:
+                feature_values.emplace_back(rocm_feature_value{ p->data.result_double });
+                break;
+            case ROCPROFILER_DATA_KIND_INT64:
+                feature_values.emplace_back(rocm_feature_value{ p->data.result_int64 });
+                break;
+        }
+    }
+}
+
+std::string
+rocm_event::as_string() const
+{
+    std::stringstream _ss{};
+    _ss << name << ", device: " << device_id << ", queue: " << queue_id
+        << ", thread: " << thread_id << ", entry: " << entry << ", exit = " << exit;
+    _ss.precision(3);
+    _ss << std::fixed;
+    for(size_t i = 0; i < feature_names.size(); ++i)
+    {
+        _ss << ", " << feature_names.at(i) << " = ";
+        auto _as_string = [&_ss](auto&& itr) { _ss << std::setw(4) << itr; };
+        std::visit(_as_string, feature_values.at(i));
+    }
+    return _ss.str();
+}
+
+void
+rocprofiler::preinit()
+{
+    rocprofiler_data::label()       = "rocprofiler";
+    rocprofiler_data::description() = "ROCm hardware counters";
+}
+
+void
+rocprofiler::start()
+{
+    if(tracker_type::start() == 0) setup();
+}
+
+void
+rocprofiler::stop()
+{
+    if(tracker_type::stop() == 0) shutdown();
+}
+
+bool
+rocprofiler::is_setup()
+{
+    return omnitrace::rocprofiler::is_setup();
+}
+
+void
+rocprofiler::add_setup(const std::string&, std::function<void()>&&)
+{}
+
+void
+rocprofiler::add_shutdown(const std::string&, std::function<void()>&&)
+{}
+
+void
+rocprofiler::remove_setup(const std::string&)
+{}
+
+void
+rocprofiler::remove_shutdown(const std::string&)
+{}
+
+void
+rocprofiler::setup()
+{
+    OMNITRACE_VERBOSE_F(1, "rocprofiler is setup\n");
+}
+
+void
+rocprofiler::shutdown()
+{
+    omnitrace::rocprofiler::post_process();
+    omnitrace::rocprofiler::rocm_cleanup();
+    /*
+    using storage_type = typename rocprofiler_data::storage_type;
+    using bundle_t     = rocprofiler_data;
+    using tag_t        = api::omnitrace;
+
+    auto    _data   = omnitrace::rocprofiler::get_data();
+    auto    _labels = omnitrace::rocprofiler::get_data_labels();
+    auto    _info   = omnitrace::rocprofiler::rocm_metrics();
+    int64_t _idx    = 0;
+    auto    _scope  = tim::scope::get_default();
+
+    auto _get_metric_desc = [_info](std::string_view _v) {
+        for(auto itr : _info)
+        {
+            if(itr.symbol().find(_v) == 0 || itr.short_description().find(_v) == 0)
+                return std::make_pair(itr.short_description(), itr.long_description());
+        }
+        return std::make_pair(std::string{}, std::string{});
+    };
+
+    auto _debug       = settings::debug();
+    settings::debug() = true;
+
+    struct hw_counters
+    {};
+
+    using rocm_counter = omnitrace::rocprofiler::rocm_counter;
+
+    struct perfetto_rocm_event
+    {
+        rocm_counter      entry = {};
+        rocm_counter      exit  = {};
+        rocprofiler_value value = {};
+
+        bool operator<(const perfetto_rocm_event& _v) const
+        {
+            return (entry.at(0) == _v.entry.at(0)) ? exit.at(0) < _v.exit.at(0)
+                                                   : entry.at(0) < _v.entry.at(0);
+        }
+    };
+
+    // contains the necessary info for export to perfetto
+    auto _perfetto_raw_data =
+        std::map<int64_t, std::map<int64_t, std::vector<perfetto_rocm_event>>>{};
+    // contains the time-stamp regions for the counter tracks
+    auto _perfetto_time_regions =
+        std::map<int64_t, std::map<int64_t, std::set<uint64_t>>>{};
+
+    // create a layout compatible for exporting to perfetto
+    for(const auto& itr : _labels)
+    {
+        auto _dev_id   = itr.first;
+        auto _dev_name = JOIN("", '[', _dev_id, ']');
+
+        for(size_t i = 0; i < itr.second.size(); ++i)
+        {
+            auto _metric_name = itr.second.at(i);
+            auto _idx         = perfetto_counter_track<hw_counters>::emplace(
+                _dev_id, JOIN(' ', "Device", _metric_name, _dev_name));
+            auto& _raw = _perfetto_raw_data[_dev_id][_idx];
+            auto& _reg = _perfetto_time_regions[_dev_id][_idx];
+            for(const auto& ditr : _data)
+            {
+                _raw.emplace_back(
+                    perfetto_rocm_event{ ditr.entry, ditr.exit, ditr.data.at(i) });
+            }
+            std::sort(_raw.begin(), _raw.end());
+            for(auto ritr : _raw)
+            {
+                if(pthread_create_gotcha::is_valid_execution_time(0, ritr.entry.at(0)))
+                    _reg.emplace(ritr.entry.at(0));
+                if(pthread_create_gotcha::is_valid_execution_time(0, ritr.exit.at(0)))
+                    _reg.emplace(ritr.exit.at(0));
+            }
+        }
+    }
+
+    for(auto& ditr : _perfetto_time_regions)
+        for(auto& citr : ditr.second)
+        {
+            for(auto _ts = citr.second.begin(); _ts != citr.second.end(); ++_ts)
+            {
+                rocprofiler_value _v    = {};
+                auto              _curr = _ts;
+                auto              _next = std::next(_ts);
+                if(_next == citr.second.end()) continue;
+                auto _min_ts = *_curr;
+                auto _max_ts = (_next == citr.second.end()) ? *_curr : *_next;
+                for(auto itr : _perfetto_raw_data[ditr.first][citr.first])
+                {
+                    if(itr.entry[0] >= _min_ts && itr.exit[0] <= _max_ts)
+                    {
+                        using namespace tim::stl;
+                        _v += itr.value;
+                    }
+                }
+
+                auto _write_counter = [&](auto _v) {
+                    if(_min_ts == _max_ts)
+                    {
+                        using value_type = std::remove_reference_t<
+                            std::remove_cv_t<decay_t<decltype(_v)>>>;
+                        _v = static_cast<value_type>(0);
+                    }
+
+                    TRACE_COUNTER(
+                        "hardware_counter",
+                        perfetto_counter_track<hw_counters>::at(ditr.first, citr.first),
+                        _min_ts, _v);
+                };
+                std::visit(_write_counter, _v);
+            }
+        }
+
+    for(const auto& itr : _labels)
+    {
+        for(size_t i = 0; i < itr.second.size(); ++i)
+        {
+            auto _metric_name         = itr.second.at(i);
+            auto _metric_desc         = _get_metric_desc(_metric_name).second;
+            rocprofiler_data::label() = _metric_name;
+            if(!_metric_desc.empty())
+                rocprofiler_data::description() = JOIN(" - ", "rocprof", _metric_desc);
+            auto _dev_id = itr.first;
+            auto _label  = JOIN('-', "rocprofiler", _metric_name, "device", _dev_id);
+            storage_type          _storage{ standalone_storage{}, ++_idx, _label };
+            std::vector<bundle_t> _bundles = {};
+            _bundles.reserve(_data.size());
+            for(const auto& ditr : _data)
+            {
+                auto _hash = add_hash_id(ditr.name);
+                auto _v    = ditr.data.at(i);
+                auto _obj  = std::tie(_bundles.emplace_back(bundle_t{}));
+                invoke::reset<tag_t>(_obj);
+                invoke::push<tag_t>(_obj, _scope, _hash, &_storage, _dev_id);
+                invoke::start<tag_t>(_obj);
+                invoke::store<tag_t>(_obj, _v);
+                invoke::stop<tag_t>(_obj);
+                invoke::pop<tag_t>(_obj, &_storage, _dev_id);
+            }
+
+            _storage.write(_label);
+        }
+    }
+    settings::debug() = _debug;
+    */
+
+    OMNITRACE_VERBOSE_F(1, "rocprofiler is shutdown\n");
+}
+
+scope::transient_destructor
+rocprofiler::protect_flush_activity()
+{
+    return scope::transient_destructor([]() { --rocprofiler_activity_count(); },
+                                       []() { ++rocprofiler_activity_count(); });
+}
+}  // namespace component
+}  // namespace tim
+
+TIMEMORY_INSTANTIATE_EXTERN_COMPONENT(rocprofiler, false, void)
+TIMEMORY_INSTANTIATE_EXTERN_COMPONENT(rocprofiler_data, true,
+                                      tim::component::rocprofiler_value)
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/rocprofiler.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/rocprofiler.hpp
new file mode 100644
index 0000000000..3c76db0150
--- /dev/null
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/rocprofiler.hpp
@@ -0,0 +1,236 @@
+// MIT License
+//
+// Copyright (c) 2022 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 "library/components/fwd.hpp"
+#include "library/defines.hpp"
+#include "library/thread_data.hpp"
+
+#include <timemory/api.hpp>
+#include <timemory/backends/hardware_counters.hpp>
+#include <timemory/components/base.hpp>
+#include <timemory/components/data_tracker/components.hpp>
+#include <timemory/components/macros.hpp>
+#include <timemory/enum.h>
+#include <timemory/macros.hpp>
+#include <timemory/macros/os.hpp>
+#include <timemory/mpl/concepts.hpp>
+#include <timemory/mpl/macros.hpp>
+#include <timemory/mpl/type_traits.hpp>
+#include <timemory/mpl/types.hpp>
+#include <timemory/utility/transient_function.hpp>
+
+#include <array>
+#include <cstdint>
+#include <string>
+#include <string_view>
+#include <variant>
+#include <vector>
+
+#if !defined(OMNITRACE_MAX_COUNTERS)
+#    define OMNITRACE_MAX_COUNTERS 25
+#endif
+
+#if !defined(OMNITRACE_ROCM_LOOK_AHEAD)
+#    define OMNITRACE_ROCM_LOOK_AHEAD 128
+#endif
+
+#if !defined(OMNITRACE_MAX_ROCM_QUEUES)
+#    define OMNITRACE_MAX_ROCM_QUEUES OMNITRACE_MAX_THREADS
+#endif
+
+namespace tim
+{
+namespace component
+{
+using rocm_metric_type   = unsigned long long;
+using rocm_info_entry    = ::tim::hardware_counters::info;
+using rocm_feature_value = std::variant<uint32_t, float, uint64_t, double>;
+
+struct rocm_counter
+{
+    std::array<rocm_metric_type, OMNITRACE_MAX_COUNTERS> counters;
+};
+
+struct rocm_event
+{
+    using value_type = rocm_feature_value;
+
+    uint32_t                        device_id      = 0;
+    uint32_t                        thread_id      = 0;
+    uint32_t                        queue_id       = 0;
+    rocm_metric_type                entry          = 0;
+    rocm_metric_type                exit           = 0;
+    std::string                     name           = {};
+    std::vector<std::string_view>   feature_names  = {};
+    std::vector<rocm_feature_value> feature_values = {};
+
+    rocm_event() = default;
+    rocm_event(uint32_t _dev, uint32_t _thr, uint32_t _queue, std::string _event_name,
+               rocm_metric_type begin, rocm_metric_type end, uint32_t _feature_count,
+               void* _features);
+
+    std::string as_string() const;
+
+    friend std::ostream& operator<<(std::ostream& _os, const rocm_event& _v)
+    {
+        return (_os << _v.as_string());
+    }
+
+    friend bool operator<(const rocm_event& _lhs, const rocm_event& _rhs)
+    {
+        return std::tie(_lhs.device_id, _lhs.queue_id, _lhs.entry, _lhs.thread_id) <
+               std::tie(_rhs.device_id, _rhs.queue_id, _rhs.entry, _rhs.thread_id);
+    }
+};
+
+using rocm_data_t       = std::vector<rocm_event>;
+using rocm_data_tracker = data_tracker<rocm_feature_value, rocm_event>;
+
+omnitrace::unique_ptr_t<rocm_data_t>&
+rocm_data(int64_t _tid = threading::get_id());
+
+using rocprofiler_value = typename rocm_event::value_type;
+using rocprofiler_data  = data_tracker<rocprofiler_value, rocprofiler>;
+
+struct rocprofiler
+: base<rocprofiler, void>
+, private policy::instance_tracker<rocprofiler, false>
+{
+    using value_type   = void;
+    using base_type    = base<rocprofiler, void>;
+    using tracker_type = policy::instance_tracker<rocprofiler, false>;
+
+    TIMEMORY_DEFAULT_OBJECT(rocprofiler)
+
+    static void preinit();
+    static void global_init() { setup(); }
+    static void global_finalize() { shutdown(); }
+
+    static bool is_setup();
+    static void setup();
+    static void shutdown();
+    static void add_setup(const std::string&, std::function<void()>&&);
+    static void add_shutdown(const std::string&, std::function<void()>&&);
+    static void remove_setup(const std::string&);
+    static void remove_shutdown(const std::string&);
+
+    void start();
+    void stop();
+
+    // this function protects rocprofiler_flush_activty from being called
+    // when omnitrace exits during a callback
+    [[nodiscard]] static scope::transient_destructor protect_flush_activity();
+};
+
+#if !defined(OMNITRACE_USE_ROCTRACER)
+inline void
+rocprofiler::setup()
+{}
+
+inline void
+rocprofiler::shutdown()
+{}
+
+inline bool
+rocprofiler::is_setup()
+{
+    return false;
+}
+#endif
+}  // namespace component
+
+namespace operation
+{
+template <>
+struct set_storage<component::rocm_data_tracker>
+{
+    using T                             = component::rocm_data_tracker;
+    static constexpr size_t max_threads = 4096;
+    using type                          = T;
+    using storage_array_t               = std::array<storage<type>*, max_threads>;
+    friend struct get_storage<component::rocm_data_tracker>;
+
+    TIMEMORY_DEFAULT_OBJECT(set_storage)
+
+    auto operator()(storage<type>*, size_t) const {}
+    auto operator()(type&, size_t) const {}
+    auto operator()(storage<type>* _v) const { get().fill(_v); }
+
+private:
+    static storage_array_t& get()
+    {
+        static storage_array_t _v = { nullptr };
+        return _v;
+    }
+};
+
+template <>
+struct get_storage<component::rocm_data_tracker>
+{
+    using type = component::rocm_data_tracker;
+
+    TIMEMORY_DEFAULT_OBJECT(get_storage)
+
+    auto operator()(const type&) const
+    {
+        return operation::set_storage<type>::get().at(0);
+    }
+
+    auto operator()() const
+    {
+        type _obj{};
+        return (*this)(_obj);
+    }
+
+    auto operator()(size_t _idx) const
+    {
+        return operation::set_storage<type>::get().at(_idx);
+    }
+
+    auto operator()(type&, size_t _idx) const { return (*this)(_idx); }
+};
+}  // namespace operation
+}  // namespace tim
+
+#if !defined(OMNITRACE_USE_ROCTRACER)
+TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, component::rocprofiler_data, false_type)
+#endif
+
+TIMEMORY_SET_COMPONENT_API(component::rocprofiler_data, project::timemory,
+                           category::timing, os::supports_unix)
+TIMEMORY_DEFINE_CONCRETE_TRAIT(is_timing_category, component::rocprofiler_data,
+                               false_type)
+TIMEMORY_DEFINE_CONCRETE_TRAIT(uses_timing_units, component::rocprofiler_data, false_type)
+TIMEMORY_DEFINE_CONCRETE_TRAIT(report_units, component::rocprofiler_data, false_type)
+TIMEMORY_STATISTICS_TYPE(component::rocprofiler_data, component::rocprofiler_value)
+
+#if !defined(OMNITRACE_EXTERN_COMPONENTS) ||                                             \
+    (defined(OMNITRACE_EXTERN_COMPONENTS) && OMNITRACE_EXTERN_COMPONENTS > 0)
+
+#    include <timemory/operations.hpp>
+
+TIMEMORY_DECLARE_EXTERN_COMPONENT(rocprofiler, false, void)
+TIMEMORY_DECLARE_EXTERN_COMPONENT(rocprofiler_data, true, double)
+
+#endif
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer.cpp
index 73cf3af4ad..bd8f62fbb2 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer.cpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer.cpp
@@ -22,12 +22,12 @@
 
 #include "library/components/roctracer.hpp"
 #include "library/components/pthread_gotcha.hpp"
-#include "library/components/roctracer_callbacks.hpp"
 #include "library/config.hpp"
 #include "library/debug.hpp"
 #include "library/defines.hpp"
 #include "library/dynamic_library.hpp"
 #include "library/redirect.hpp"
+#include "library/roctracer.hpp"
 #include "library/sampling.hpp"
 #include "library/thread_data.hpp"
 
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/config.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/config.cpp
index 6f686e2222..86d4fa971e 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/config.cpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/config.cpp
@@ -130,6 +130,14 @@ _settings_are_configured()
 }
 }  // namespace
 
+void
+finalize()
+{
+    OMNITRACE_DEBUG("[omnitrace_finalize] Disabling signal handling...\n");
+    tim::disable_signal_detection();
+    _settings_are_configured() = false;
+}
+
 bool
 settings_are_configured()
 {
@@ -142,7 +150,6 @@ configure_settings(bool _init)
 {
     volatile bool _v = _settings_are_configured();
     if(_v) return;
-    _settings_are_configured() = true;
 
     static bool _once = false;
     if(_once) return;
@@ -161,6 +168,8 @@ configure_settings(bool _init)
     tim::manager::add_metadata("OMNITRACE_VERSION_MAJOR", OMNITRACE_VERSION_MAJOR);
     tim::manager::add_metadata("OMNITRACE_VERSION_MINOR", OMNITRACE_VERSION_MINOR);
     tim::manager::add_metadata("OMNITRACE_VERSION_PATCH", OMNITRACE_VERSION_PATCH);
+    tim::manager::add_metadata("OMNITRACE_GIT_DESCRIBE", OMNITRACE_GIT_DESCRIBE);
+    tim::manager::add_metadata("OMNITRACE_GIT_REVISION", OMNITRACE_GIT_REVISION);
 
 #if OMNITRACE_HIP_VERSION > 0
     tim::manager::add_metadata("OMNITRACE_HIP_VERSION", OMNITRACE_HIP_VERSION_STRING);
@@ -206,8 +215,13 @@ configure_settings(bool _init)
                              !_config->get<bool>("OMNITRACE_USE_PERFETTO"), "backend",
                              "timemory");
 
-    OMNITRACE_CONFIG_SETTING(bool, "OMNITRACE_USE_ROCTRACER", "Enable ROCM tracing", true,
-                             "backend", "roctracer", "rocm");
+    OMNITRACE_CONFIG_SETTING(bool, "OMNITRACE_USE_ROCTRACER",
+                             "Enable ROCm API and kernel tracing", true, "backend",
+                             "roctracer", "rocm");
+
+    OMNITRACE_CONFIG_SETTING(bool, "OMNITRACE_USE_ROCPROFILER",
+                             "Enable ROCm hardware counters", true, "backend",
+                             "rocprofiler", "rocm");
 
     OMNITRACE_CONFIG_SETTING(
         bool, "OMNITRACE_USE_ROCM_SMI",
@@ -343,6 +357,13 @@ configure_settings(bool _init)
     OMNITRACE_CONFIG_SETTING(std::string, "OMNITRACE_ROCTRACER_HSA_API_TYPES",
                              "HSA API type to collect", "", "roctracer", "rocm");
 
+    OMNITRACE_CONFIG_SETTING(
+        std::string, "OMNITRACE_ROCM_EVENTS",
+        "ROCm hardware counters. Use ':device=N' syntax to specify collection on device "
+        "number N, e.g. ':device=0'. If no device specification is provided, the event "
+        "is collected on every available device",
+        "", "rocprofiler", "rocm", "hardware_counters");
+
     OMNITRACE_CONFIG_SETTING(bool, "OMNITRACE_CRITICAL_TRACE_DEBUG",
                              "Enable debugging for critical trace", _omnitrace_debug,
                              "debugging", "critical_trace");
@@ -553,6 +574,8 @@ configure_settings(bool _init)
     configure_signal_handler();
     configure_disabled_settings();
 
+    _settings_are_configured() = true;
+
     OMNITRACE_CONDITIONAL_BASIC_PRINT(get_verbose_env() > 0, "configuration complete\n");
 }
 
@@ -582,6 +605,7 @@ configure_mode_settings()
         _set("OMNITRACE_USE_TIMEMORY", false);
         _set("OMNITRACE_USE_ROCM_SMI", false);
         _set("OMNITRACE_USE_ROCTRACER", false);
+        _set("OMNITRACE_USE_ROCPROFILER", false);
         _set("OMNITRACE_USE_KOKKOSP", false);
         _set("OMNITRACE_USE_OMPT", false);
         _set("OMNITRACE_USE_SAMPLING", false);
@@ -599,8 +623,9 @@ configure_mode_settings()
     {
         OMNITRACE_VERBOSE_F(
             1, "No HIP devices were found: disabling roctracer and rocm_smi...\n");
-        get_use_roctracer() = false;
-        get_use_rocm_smi()  = false;
+        _set("OMNITRACE_USE_ROCTRACER", false);
+        _set("OMNITRACE_USE_ROCPROFILER", false);
+        _set("OMNITRACE_USE_ROCM_SMI", false);
     }
 
     get_instrumentation_interval() = std::max<size_t>(get_instrumentation_interval(), 1);
@@ -624,6 +649,7 @@ configure_mode_settings()
         _set("OMNITRACE_USE_TIMEMORY", false);
         _set("OMNITRACE_USE_ROCM_SMI", false);
         _set("OMNITRACE_USE_ROCTRACER", false);
+        _set("OMNITRACE_USE_ROCPROFILER", false);
         _set("OMNITRACE_USE_KOKKOSP", false);
         _set("OMNITRACE_USE_OMPT", false);
         _set("OMNITRACE_USE_SAMPLING", false);
@@ -723,6 +749,7 @@ configure_disabled_settings()
     _handle_use_option("OMNITRACE_USE_OMPT", "ompt");
     _handle_use_option("OMNITRACE_USE_ROCM_SMI", "rocm_smi");
     _handle_use_option("OMNITRACE_USE_ROCTRACER", "roctracer");
+    _handle_use_option("OMNITRACE_USE_ROCPROFILER", "rocprofiler");
     _handle_use_option("OMNITRACE_CRITICAL_TRACE", "critical_trace");
 
 #if !defined(OMNITRACE_USE_ROCTRACER) || OMNITRACE_USE_ROCTRACER == 0
@@ -731,6 +758,12 @@ configure_disabled_settings()
         _config->find(itr)->second->set_hidden(true);
 #endif
 
+#if !defined(OMNITRACE_USE_ROCPROFILER) || OMNITRACE_USE_ROCPROFILER == 0
+    _config->find("OMNITRACE_USE_ROCPROFILER")->second->set_hidden(true);
+    for(const auto& itr : _config->disable_category("rocprofiler"))
+        _config->find(itr)->second->set_hidden(true);
+#endif
+
 #if !defined(OMNITRACE_USE_ROCM_SMI) || OMNITRACE_USE_ROCM_SMI == 0
     _config->find("OMNITRACE_USE_ROCM_SMI")->second->set_hidden(true);
     for(const auto& itr : _config->disable_category("rocm_smi"))
@@ -1064,7 +1097,8 @@ is_binary_rewrite()
 bool
 get_debug_env()
 {
-    return tim::get_env<bool>("OMNITRACE_DEBUG", false);
+    return (settings_are_configured()) ? get_debug()
+                                       : tim::get_env<bool>("OMNITRACE_DEBUG", false);
 }
 
 bool
@@ -1106,7 +1140,8 @@ get_debug_sampling()
 int
 get_verbose_env()
 {
-    return tim::get_env<int>("OMNITRACE_VERBOSE", 0);
+    return (settings_are_configured()) ? get_verbose()
+                                       : tim::get_env<int>("OMNITRACE_VERBOSE", 0);
 }
 
 int
@@ -1130,27 +1165,36 @@ get_use_timemory()
     return static_cast<tim::tsettings<bool>&>(*_v->second).get();
 }
 
-bool&
+bool
 get_use_roctracer()
 {
 #if defined(OMNITRACE_USE_ROCTRACER) && OMNITRACE_USE_ROCTRACER > 0
     static auto _v = get_config()->find("OMNITRACE_USE_ROCTRACER");
     return static_cast<tim::tsettings<bool>&>(*_v->second).get();
 #else
-    static auto _v            = false;
-    return _v;
+    return false;
 #endif
 }
 
-bool&
+bool
+get_use_rocprofiler()
+{
+#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
+    static auto _v = get_config()->find("OMNITRACE_USE_ROCPROFILER");
+    return static_cast<tim::tsettings<bool>&>(*_v->second).get();
+#else
+    return false;
+#endif
+}
+
+bool
 get_use_rocm_smi()
 {
 #if defined(OMNITRACE_USE_ROCM_SMI) && OMNITRACE_USE_ROCM_SMI > 0
     static auto _v = get_config()->find("OMNITRACE_USE_ROCM_SMI");
     return static_cast<tim::tsettings<bool>&>(*_v->second).get();
 #else
-    static auto _v = false;
-    return _v;
+    return false;
 #endif
 }
 
@@ -1442,6 +1486,13 @@ get_trace_thread_locks()
     return static_cast<tim::tsettings<bool>&>(*_v->second).get();
 }
 
+std::string
+get_rocm_events()
+{
+    static auto _v = get_config()->find("OMNITRACE_ROCM_EVENTS");
+    return static_cast<tim::tsettings<std::string>&>(*_v->second).get();
+}
+
 bool
 get_trace_thread_rwlocks()
 {
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/config.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/config.hpp
index c33a331aca..49835430c8 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/config.hpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/config.hpp
@@ -57,6 +57,9 @@ configure_signal_handler();
 void
 configure_disabled_settings();
 
+void
+finalize();
+
 void
 handle_deprecated_setting(const std::string& _old, const std::string& _new,
                           int _verbose = 0);
@@ -162,10 +165,13 @@ get_use_perfetto() OMNITRACE_HOT;
 bool&
 get_use_timemory() OMNITRACE_HOT;
 
-bool&
+bool
 get_use_roctracer() OMNITRACE_HOT;
 
-bool&
+bool
+get_use_rocprofiler() OMNITRACE_HOT;
+
+bool
 get_use_rocm_smi() OMNITRACE_HOT;
 
 bool&
@@ -276,6 +282,9 @@ get_trace_thread_locks();
 
 bool
 get_trace_thread_rwlocks();
+
+std::string
+get_rocm_events();
 }  // namespace config
 
 //
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/defines.hpp.in b/projects/rocprofiler-systems/source/lib/omnitrace/library/defines.hpp.in
index c9d4c07207..05602ccb22 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/defines.hpp.in
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/defines.hpp.in
@@ -27,6 +27,8 @@
 #define OMNITRACE_VERSION_MAJOR @PROJECT_VERSION_MAJOR@
 #define OMNITRACE_VERSION_MINOR @PROJECT_VERSION_MINOR@
 #define OMNITRACE_VERSION_PATCH @PROJECT_VERSION_PATCH@
+#define OMNITRACE_GIT_DESCRIBE "@OMNITRACE_GIT_DESCRIBE@"
+#define OMNITRACE_GIT_REVISION "@OMNITRACE_GIT_REVISION@"
 
 #define OMNITRACE_HIP_VERSION_STRING "@OMNITRACE_HIP_VERSION@"
 #define OMNITRACE_HIP_VERSION_MAJOR @OMNITRACE_HIP_VERSION_MAJOR@
@@ -40,20 +42,25 @@
 #endif
 // clang-format on
 
+#define OMNITRACE_VERSION                                                                \
+    ((10000 * OMNITRACE_VERSION_MAJOR) + (100 * OMNITRACE_VERSION_MINOR) +               \
+     OMNITRACE_VERSION_PATCH)
+
 #define OMNITRACE_HIP_VERSION                                                            \
     ((10000 * OMNITRACE_HIP_VERSION_MAJOR) + (100 * OMNITRACE_HIP_VERSION_MINOR) +       \
      OMNITRACE_HIP_VERSION_PATCH)
 
 #define TIMEMORY_USER_COMPONENT_ENUM                                                     \
     OMNITRACE_COMPONENT_idx, OMNITRACE_USER_REGION_idx, OMNITRACE_ROCTRACER_idx,         \
-        OMNITRACE_SAMPLING_WALL_CLOCK_idx, OMNITRACE_SAMPLING_CPU_CLOCK_idx,             \
-        OMNITRACE_SAMPLING_PERCENT_idx, OMNITRACE_SAMPLING_GPU_POWER_idx,                \
-        OMNITRACE_SAMPLING_GPU_TEMP_idx, OMNITRACE_SAMPLING_GPU_BUSY_idx,                \
-        OMNITRACE_SAMPLING_GPU_MEMORY_USAGE_idx,
+        OMNITRACE_ROCPROFILER_idx, OMNITRACE_SAMPLING_WALL_CLOCK_idx,                    \
+        OMNITRACE_SAMPLING_CPU_CLOCK_idx, OMNITRACE_SAMPLING_PERCENT_idx,                \
+        OMNITRACE_SAMPLING_GPU_POWER_idx, OMNITRACE_SAMPLING_GPU_TEMP_idx,               \
+        OMNITRACE_SAMPLING_GPU_BUSY_idx, OMNITRACE_SAMPLING_GPU_MEMORY_USAGE_idx,
 
 #define OMNITRACE_COMPONENT                 OMNITRACE_COMPONENT_idx
 #define OMNITRACE_USER_REGION               OMNITRACE_USER_REGION_idx
 #define OMNITRACE_ROCTRACER                 OMNITRACE_ROCTRACER_idx
+#define OMNITRACE_ROCPROFILER               OMNITRACE_ROCPROFILER_idx
 #define OMNITRACE_SAMPLING_WALL_CLOCK       OMNITRACE_SAMPLING_WALL_CLOCK_idx
 #define OMNITRACE_SAMPLING_CPU_CLOCK        OMNITRACE_SAMPLING_CPU_CLOCK_idx
 #define OMNITRACE_SAMPLING_PERCENT          OMNITRACE_SAMPLING_PERCENT_idx
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/dynamic_library.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/dynamic_library.hpp
index f27220c94a..8e79d5f931 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/dynamic_library.hpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/dynamic_library.hpp
@@ -26,6 +26,7 @@
 
 #include <dlfcn.h>
 #include <string>
+#include <unistd.h>
 
 namespace omnitrace
 {
@@ -45,6 +46,26 @@ struct dynamic_library
     bool open();
     int  close() const;
 
+    template <typename RetT, typename... Args>
+    RetT invoke(std::string_view _name, RetT (*&_func)(Args...), Args... _args)
+    {
+        if(!handle) open();
+        if(handle)
+        {
+            *(void**) (&_func) = dlsym(handle, _name.data());
+            if(_func)
+            {
+                return (*_func)(_args...);
+            }
+            else
+            {
+                fprintf(stderr, "[omnitrace][pid=%i]> %s :: %s\n", getpid(), _name.data(),
+                        dlerror());
+            }
+        }
+        return RetT{};
+    }
+
     std::string envname  = {};
     std::string filename = {};
     int         flags    = 0;
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/gpu.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/gpu.cpp
index ee040cc941..a778f67e6f 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/gpu.cpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/gpu.cpp
@@ -24,25 +24,41 @@
 
 #if defined(OMNITRACE_USE_ROCM_SMI) && OMNITRACE_USE_ROCM_SMI > 0
 #    include "library/components/rocm_smi.hpp"
-#elif defined(OMNITRACE_USE_HIP) && OMNITRACE_USE_HIP > 0
+#elif !defined(OMNITRACE_USE_ROCM_SMI)
+#    define OMNITRACE_USE_ROCM_SMI 0
+#endif
+
+#if defined(OMNITRACE_USE_HIP) && OMNITRACE_USE_HIP > 0
 #    if !defined(TIMEMORY_USE_HIP)
 #        define TIMEMORY_USE_HIP 1
 #    endif
 #    include <timemory/components/hip/backends.hpp>
+#elif !defined(OMNITRACE_USE_HIP)
+#    define OMNITRACE_USE_HIP 0
 #endif
 
 namespace omnitrace
 {
 namespace gpu
 {
+int
+hip_device_count()
+{
+#if OMNITRACE_USE_HIP > 0
+    return ::tim::hip::device_count();
+#else
+    return 0;
+#endif
+}
+
 int
 device_count()
 {
-#if defined(OMNITRACE_USE_ROCM_SMI) && OMNITRACE_USE_ROCM_SMI > 0
+#if OMNITRACE_USE_ROCM_SMI > 0
     // store as static since calls after rsmi_shutdown will return zero
     static auto _v = rocm_smi::device_count();
     return _v;
-#elif defined(OMNITRACE_USE_HIP) && OMNITRACE_USE_HIP > 0
+#elif OMNITRACE_USE_HIP > 0
     return ::tim::hip::device_count();
 #else
     return 0;
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/gpu.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/gpu.hpp
index 3ef15ae569..324c7c4180 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/gpu.hpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/gpu.hpp
@@ -28,5 +28,8 @@ namespace gpu
 {
 int
 device_count();
-}
+
+int
+hip_device_count();
+}  // namespace gpu
 }  // namespace omnitrace
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/ompt.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/ompt.cpp
index fb702ea044..f41952c33f 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/ompt.cpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/ompt.cpp
@@ -68,8 +68,8 @@ setup()
     comp::user_ompt_bundle::reset();
     tim::auto_lock_t lk{ tim::type_mutex<ompt_handle_t>() };
     comp::user_ompt_bundle::configure<omnitrace::component::user_region>();
-    f_bundle =
-        std::make_unique<ompt_bundle_t>("ompt", quirk::config<quirk::auto_start>{});
+    f_bundle = std::make_unique<ompt_bundle_t>("omnitrace/ompt",
+                                               quirk::config<quirk::auto_start>{});
 }
 
 void
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/perfetto.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/perfetto.hpp
index 13d9d329e2..62de2252be 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/perfetto.hpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/perfetto.hpp
@@ -173,7 +173,8 @@ struct perfetto_counter_track
                 _missing.emplace_back(std::make_tuple(*itr, itr->c_str(), false));
             }
         }
-        auto&       _name = _name_data.emplace_back(std::make_unique<std::string>(_v));
+        auto        _index = _track_data.size();
+        auto&       _name  = _name_data.emplace_back(std::make_unique<std::string>(_v));
         const char* _unit_name = (_units && strlen(_units) > 0) ? _units : nullptr;
         _track_data.emplace_back(perfetto::CounterTrack{ _name->c_str() }
                                      .set_unit_name(_unit_name)
@@ -217,6 +218,7 @@ struct perfetto_counter_track
                 }
             }
         }
+        return _index;
     }
 
     static auto& at(size_t _idx, size_t _n) { return get_data().second.at(_idx).at(_n); }
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/rocm.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocm.cpp
new file mode 100644
index 0000000000..161242ccca
--- /dev/null
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocm.cpp
@@ -0,0 +1,264 @@
+// MIT License
+//
+// Copyright (c) 2022 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 "library/rocm.hpp"
+#include "library.hpp"
+#include "library/components/rocm_smi.hpp"
+#include "library/components/rocprofiler.hpp"
+#include "library/components/roctracer.hpp"
+#include "library/config.hpp"
+#include "library/critical_trace.hpp"
+#include "library/debug.hpp"
+#include "library/rocprofiler.hpp"
+#include "library/rocprofiler/hsa_rsrc_factory.hpp"
+#include "library/roctracer.hpp"
+#include "library/runtime.hpp"
+#include "library/sampling.hpp"
+#include "library/thread_data.hpp"
+#include "library/tracing.hpp"
+
+#include <timemory/backends/cpu.hpp>
+#include <timemory/backends/threading.hpp>
+#include <timemory/utility/types.hpp>
+
+#include <atomic>
+#include <chrono>
+#include <cstdint>
+#include <mutex>
+#include <tuple>
+
+#define HIP_PROF_HIP_API_STRING 1
+
+#include <roctracer_ext.h>
+#include <roctracer_hcc.h>
+#include <roctracer_hip.h>
+
+#define AMD_INTERNAL_BUILD 1
+#include <roctracer_hsa.h>
+
+#if __has_include(<hip/amd_detail/hip_prof_str.h>) || (defined(OMNITRACE_USE_HIP) && OMNITRACE_USE_HIP > 0)
+#    include <hip/amd_detail/hip_prof_str.h>
+#    define OMNITRACE_HIP_API_ARGS 1
+#else
+#    define OMNITRACE_HIP_API_ARGS 0
+#endif
+
+#include <rocprofiler.h>
+
+using namespace omnitrace;
+
+namespace omnitrace
+{
+namespace rocm
+{
+std::mutex rocm_mutex = {};
+bool       is_loaded  = false;
+}  // namespace rocm
+}  // namespace omnitrace
+
+// HSA-runtime tool on-load method
+extern "C"
+{
+#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
+    void OnUnloadTool()
+    {
+        OMNITRACE_BASIC_VERBOSE(2, "Inside %s\n", __FUNCTION__);
+
+        rocm::lock_t _lk{ rocm::rocm_mutex, std::defer_lock };
+        if(!_lk.owns_lock()) _lk.lock();
+
+        if(!rocm::is_loaded) return;
+        rocm::is_loaded = false;
+
+        _lk.unlock();
+
+        // stop_top_level_timer_if_necessary();
+        // Final resources cleanup
+        omnitrace::rocprofiler::rocm_cleanup();
+    }
+
+    void OnLoadToolProp(rocprofiler_settings_t* settings)
+    {
+        OMNITRACE_BASIC_VERBOSE(2, "Inside %s\n", __FUNCTION__);
+
+        rocm::lock_t _lk{ rocm::rocm_mutex, std::defer_lock };
+        if(!_lk.owns_lock()) _lk.lock();
+
+        if(rocm::is_loaded) return;
+        rocm::is_loaded = true;
+
+        _lk.unlock();
+
+        // Enable timestamping
+        settings->timestamp_on = 1u;
+
+        // Initialize profiling
+        omnitrace::rocprofiler::rocm_initialize();
+        HsaRsrcFactory::Instance().PrintGpuAgents("ROCm");
+    }
+#endif
+
+    bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
+                const char* const* failed_tool_names)
+    {
+        OMNITRACE_BASIC_VERBOSE(2, "Inside %s\n", __FUNCTION__);
+
+        if(!tim::get_env("OMNITRACE_INIT_TOOLING", true)) return true;
+        if(!tim::settings::enabled()) return true;
+
+        roctracer_is_init() = true;
+        pthread_gotcha::push_enable_sampling_on_child_threads(false);
+        OMNITRACE_BASIC_VERBOSE_F(1, "\n");
+
+        tim::consume_parameters(table, runtime_version, failed_tool_count,
+                                failed_tool_names);
+
+        if(!config::settings_are_configured() && get_state() < State::Active)
+            omnitrace_init_tooling_hidden();
+
+        OMNITRACE_SCOPED_THREAD_STATE(ThreadState::Internal);
+
+        static auto _setup = [=]() {
+            try
+            {
+                OMNITRACE_VERBOSE(1, "[OnLoad] setting up HSA...\n");
+
+                // const char* output_prefix = getenv("ROCP_OUTPUT_DIR");
+                const char* output_prefix = nullptr;
+
+                bool trace_hsa_api = get_trace_hsa_api();
+
+                // Enable HSA API callbacks/activity
+                if(trace_hsa_api)
+                {
+                    std::vector<std::string> hsa_api_vec =
+                        tim::delimit(get_trace_hsa_api_types());
+
+                    // initialize HSA tracing
+                    roctracer_set_properties(ACTIVITY_DOMAIN_HSA_API, (void*) table);
+
+                    OMNITRACE_VERBOSE(1, "    HSA-trace(");
+                    if(!hsa_api_vec.empty())
+                    {
+                        for(const auto& itr : hsa_api_vec)
+                        {
+                            uint32_t    cid = HSA_API_ID_NUMBER;
+                            const char* api = itr.c_str();
+                            ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_HSA_API, api,
+                                                             &cid, nullptr));
+                            ROCTRACER_CALL(roctracer_enable_op_callback(
+                                ACTIVITY_DOMAIN_HSA_API, cid, hsa_api_callback, nullptr));
+
+                            OMNITRACE_VERBOSE(1, " %s", api);
+                        }
+                    }
+                    else
+                    {
+                        ROCTRACER_CALL(roctracer_enable_domain_callback(
+                            ACTIVITY_DOMAIN_HSA_API, hsa_api_callback, nullptr));
+                    }
+                    OMNITRACE_VERBOSE(1, " )\n");
+                }
+
+                bool trace_hsa_activity = get_trace_hsa_activity();
+                // Enable HSA GPU activity
+                if(trace_hsa_activity)
+                {
+                    // initialize HSA tracing
+                    ::roctracer::hsa_ops_properties_t ops_properties{
+                        table,
+                        reinterpret_cast<activity_async_callback_t>(
+                            hsa_activity_callback),
+                        nullptr, output_prefix
+                    };
+                    roctracer_set_properties(ACTIVITY_DOMAIN_HSA_OPS, &ops_properties);
+
+                    OMNITRACE_VERBOSE(1, "    HSA-activity-trace()\n");
+                    ROCTRACER_CALL(roctracer_enable_op_activity(ACTIVITY_DOMAIN_HSA_OPS,
+                                                                HSA_OP_ID_COPY));
+                }
+            } catch(std::exception& _e)
+            {
+                OMNITRACE_BASIC_PRINT("Exception was thrown in HSA setup: %s\n",
+                                      _e.what());
+            }
+        };
+
+        static auto _shutdown = []() {
+            OMNITRACE_DEBUG_F("roctracer_disable_domain_callback\n");
+            ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HSA_API));
+
+            OMNITRACE_DEBUG_F("roctracer_disable_op_activity\n");
+            ROCTRACER_CALL(
+                roctracer_disable_op_activity(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY));
+        };
+
+        (void) omnitrace::get_clock_skew();
+
+        comp::roctracer::add_setup("hsa", _setup);
+        comp::roctracer::add_shutdown("hsa", _shutdown);
+
+        rocm_smi::set_state(State::Active);
+        comp::roctracer::setup();
+
+#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
+        bool _force_rocprofiler_init =
+            tim::get_env("OMNITRACE_FORCE_ROCPROFILE_INIT", false, false);
+#else
+        bool _force_rocprofiler_init = false;
+#endif
+
+        bool _success = true;
+        bool _is_empty =
+            (config::settings_are_configured() && config::get_rocm_events().empty());
+        if(_force_rocprofiler_init || (get_use_rocprofiler() && !_is_empty))
+        {
+            auto _rocprof =
+                dynamic_library{ "OMNITRACE_ROCPROFILER_LIBRARY", "librocprofiler64.so",
+                                 (RTLD_LAZY | RTLD_GLOBAL), true };
+
+            on_load_t _rocprof_load = nullptr;
+            _success = _rocprof.invoke("OnLoad", _rocprof_load, table, runtime_version,
+                                       failed_tool_count, failed_tool_names);
+            OMNITRACE_CONDITIONAL_PRINT_F(!_success,
+                                          "Warning! Invoking rocprofiler's OnLoad "
+                                          "failed! OMNITRACE_ROCPROFILER_LIBRARY=%s\n",
+                                          _rocprof.filename.c_str());
+            OMNITRACE_CI_THROW(!_success,
+                               "Warning! Invoking rocprofiler's OnLoad "
+                               "failed! OMNITRACE_ROCPROFILER_LIBRARY=%s\n",
+                               _rocprof.filename.c_str());
+        }
+        pthread_gotcha::pop_enable_sampling_on_child_threads();
+        return _success;
+    }
+
+    // HSA-runtime on-unload method
+    void OnUnload()
+    {
+        OMNITRACE_BASIC_VERBOSE(2, "Inside %s\n", __FUNCTION__);
+        rocm_smi::set_state(State::Finalized);
+        comp::roctracer::shutdown();
+        comp::rocprofiler::shutdown();
+        omnitrace_finalize_hidden();
+    }
+}
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/rocm.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocm.hpp
new file mode 100644
index 0000000000..94f982ce48
--- /dev/null
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocm.hpp
@@ -0,0 +1,58 @@
+// MIT License
+//
+// Copyright (c) 2022 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 "library/defines.hpp"
+
+#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
+#    include <rocprofiler.h>
+#endif
+
+#include <cstdint>
+#include <mutex>
+
+namespace omnitrace
+{
+namespace rocm
+{
+using lock_t = std::unique_lock<std::mutex>;
+
+extern std::mutex rocm_mutex;
+extern bool       is_loaded;
+}  // namespace rocm
+}  // namespace omnitrace
+
+extern "C"
+{
+    struct HsaApiTable;
+    using on_load_t = bool (*)(HsaApiTable*, uint64_t, uint64_t, const char* const*);
+
+    bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
+                const char* const* failed_tool_names) OMNITRACE_PUBLIC_API;
+    void OnUnload() OMNITRACE_PUBLIC_API;
+
+#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
+    void OnLoadToolProp(rocprofiler_settings_t* settings) OMNITRACE_PUBLIC_API;
+    void OnUnloadTool() OMNITRACE_PUBLIC_API;
+#endif
+}
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler.cpp
new file mode 100644
index 0000000000..c663ae5c13
--- /dev/null
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler.cpp
@@ -0,0 +1,791 @@
+// MIT License
+//
+// Copyright (c) 2022 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 "library/rocprofiler.hpp"
+#include "library/common.hpp"
+#include "library/config.hpp"
+#include "library/debug.hpp"
+#include "library/gpu.hpp"
+#include "library/perfetto.hpp"
+#include "library/rocm.hpp"
+#include "library/rocprofiler/hsa_rsrc_factory.hpp"
+
+#include <timemory/backends/hardware_counters.hpp>
+#include <timemory/manager.hpp>
+#include <timemory/mpl/concepts.hpp>
+#include <timemory/storage/types.hpp>
+#include <timemory/utility/types.hpp>
+
+#include <rocprofiler.h>
+
+#include <atomic>
+#include <cstdlib>
+#include <dlfcn.h>
+#include <hsa.h>
+#include <iostream>
+#include <mutex>
+#include <sstream>
+#include <string.h>
+#include <string_view>
+#include <type_traits>
+#include <unistd.h>
+#include <vector>
+
+namespace omnitrace
+{
+namespace rocprofiler
+{
+namespace
+{
+auto&
+get_event_names()
+{
+    static auto _v = std::map<uint32_t, std::vector<rocprofiler_feature_t>>{};
+    return _v;
+}
+}  // namespace
+
+// Error handler
+void
+fatal(const std::string& msg)
+{
+    OMNITRACE_PRINT_F("\n");
+    OMNITRACE_PRINT_F("%s\n", msg.c_str());
+    abort();
+}
+
+// Check returned HSA API status
+const char*
+rocm_error_string(hsa_status_t _status)
+{
+    const char* _err_string = nullptr;
+    if(_status != HSA_STATUS_SUCCESS) rocprofiler_error_string(&_err_string);
+    return _err_string;
+}
+
+// Check returned HSA API status
+bool
+rocm_check_status(hsa_status_t _status, const std::set<hsa_status_t>& _nonfatal = {})
+{
+    if(_status != HSA_STATUS_SUCCESS)
+    {
+        if(_nonfatal.count(_status) == 0)
+            fatal(JOIN(" :: ", "ERROR", rocm_error_string(_status)));
+
+        OMNITRACE_PRINT_F("Warning! %s\n", rocm_error_string(_status));
+        return false;
+    }
+    return true;
+}
+
+// Context stored entry type
+struct context_entry_t
+{
+    bool                        valid;
+    hsa_agent_t                 agent;
+    rocprofiler_group_t         group;
+    rocprofiler_callback_data_t data;
+};
+
+// Context callback arg
+struct callbacks_arg_t
+{
+    rocprofiler_pool_t** pools;
+};
+
+// Handler callback arg
+struct handler_arg_t
+{
+    rocprofiler_feature_t* features;
+    unsigned               feature_count;
+};
+
+bool&
+is_setup()
+{
+    static bool _v = false;
+    return _v;
+}
+
+std::map<uint32_t, std::vector<std::string_view>>
+get_data_labels()
+{
+    auto _v = std::map<uint32_t, std::vector<std::string_view>>{};
+    for(const auto& itr : get_event_names())
+    {
+        _v[itr.first] = {};
+        for(auto vitr : itr.second)
+            _v[itr.first].emplace_back(std::string_view{ vitr.name });
+    }
+    return _v;
+}
+
+// Dump stored context entry
+void
+rocm_dump_context_entry(context_entry_t* entry, rocprofiler_feature_t* features,
+                        unsigned feature_count)
+{
+    // static rocm_metric_type last_timestamp = get_last_timestamp_ns();
+
+    volatile std::atomic<bool>* valid =
+        reinterpret_cast<std::atomic<bool>*>(&entry->valid);
+    while(valid->load() == false)
+        sched_yield();
+
+    const rocprofiler_dispatch_record_t* record = entry->data.record;
+
+    if(!record) return;  // there is nothing to do here.
+
+    auto _queue_id    = entry->data.queue_id;
+    auto _thread_id   = entry->data.thread_id;
+    auto _dev_id      = HsaRsrcFactory::Instance().GetAgentInfo(entry->agent)->dev_index;
+    auto _kernel_name = std::string{ entry->data.kernel_name };
+    auto _pos         = _kernel_name.find_last_of(')');
+    if(_pos != std::string::npos) _kernel_name = _kernel_name.substr(0, _pos + 1);
+
+    rocprofiler_group_t& group = entry->group;
+    if(group.context == nullptr)
+    {
+        fatal("context is nullptr\n");
+    }
+
+    if(feature_count > 0)
+    {
+        rocm_check_status(rocprofiler_group_get_data(&group));
+        rocm_check_status(rocprofiler_get_metrics(group.context));
+    }
+
+    auto _evt = comp::rocm_event{ _dev_id,       _thread_id,  _queue_id,     _kernel_name,
+                                  record->begin, record->end, feature_count, features };
+
+    comp::rocm_data()->emplace_back(_evt);
+}
+
+// Profiling completion handler
+// Dump and delete the context entry
+// Return true if the context was dumped successfully
+bool
+rocm_context_handler(const rocprofiler_pool_entry_t* entry, void* arg)
+{
+    // Context entry
+    context_entry_t* ctx_entry   = reinterpret_cast<context_entry_t*>(entry->payload);
+    handler_arg_t*   handler_arg = reinterpret_cast<handler_arg_t*>(arg);
+
+    rocm::lock_t _lk{ rocm::rocm_mutex, std::defer_lock };
+    if(!_lk.owns_lock()) _lk.lock();
+
+    rocm_dump_context_entry(ctx_entry, handler_arg->features, handler_arg->feature_count);
+
+    return true;
+}
+
+// Kernel disoatch callback
+hsa_status_t
+rocm_dispatch_callback(const rocprofiler_callback_data_t* callback_data, void* arg,
+                       rocprofiler_group_t* group)
+{
+    // Passed tool data
+    hsa_agent_t agent = callback_data->agent;
+
+    // Open profiling context
+    const unsigned   gpu_id = HsaRsrcFactory::Instance().GetAgentInfo(agent)->dev_index;
+    callbacks_arg_t* callbacks_arg = reinterpret_cast<callbacks_arg_t*>(arg);
+    rocprofiler_pool_t*      pool  = callbacks_arg->pools[gpu_id];
+    rocprofiler_pool_entry_t pool_entry{};
+    rocm_check_status(rocprofiler_pool_fetch(pool, &pool_entry));
+    // Profiling context entry
+    rocprofiler_t*   context = pool_entry.context;
+    context_entry_t* entry   = reinterpret_cast<context_entry_t*>(pool_entry.payload);
+
+    // Get group[0]
+    rocm_check_status(rocprofiler_get_group(context, 0, group));
+
+    // Fill profiling context entry
+    entry->agent            = agent;
+    entry->group            = *group;
+    entry->data             = *callback_data;
+    entry->data.kernel_name = strdup(callback_data->kernel_name);
+    reinterpret_cast<std::atomic<bool>*>(&entry->valid)->store(true);
+
+    return HSA_STATUS_SUCCESS;
+}
+
+unsigned
+metrics_input(unsigned _device, rocprofiler_feature_t** ret)
+{
+    // OMNITRACE_THROW("%s\n", __FUNCTION__);
+    // Profiling feature objects
+    auto                     _events = tim::delimit(config::get_rocm_events(), ", ;\t\n");
+    std::vector<std::string> _features    = {};
+    auto                     _this_device = JOIN("", ":device=", _device);
+    for(auto itr : _events)
+    {
+        OMNITRACE_VERBOSE_F(3, "Processing feature '%s' for device %u...\n", itr.c_str(),
+                            _device);
+        auto _pos = itr.find(":device=");
+        if(_pos != std::string::npos)
+        {
+            if(itr.find(_this_device) != std::string::npos)
+            {
+                _features.emplace_back(itr.substr(0, _pos));
+            }
+        }
+        else
+        {
+            _features.emplace_back(itr);
+        }
+    }
+    const unsigned         feature_count = _features.size();
+    rocprofiler_feature_t* features      = new rocprofiler_feature_t[feature_count];
+    memset(features, 0, feature_count * sizeof(rocprofiler_feature_t));
+
+    // PMC events
+    for(unsigned i = 0; i < feature_count; ++i)
+    {
+        OMNITRACE_VERBOSE_F(3, "Adding feature '%s' for device %u...\n",
+                            _features.at(i).c_str(), _device);
+        features[i].kind            = ROCPROFILER_FEATURE_KIND_METRIC;
+        features[i].name            = strdup(_features.at(i).c_str());
+        features[i].parameters      = nullptr;
+        features[i].parameter_count = 0;
+    }
+
+    *ret = features;
+    return feature_count;
+}
+
+struct info_data
+{
+    const AgentInfo*                    agent = nullptr;
+    std::vector<comp::rocm_info_entry>* data  = nullptr;
+};
+
+hsa_status_t
+info_data_callback(const rocprofiler_info_data_t info, void* arg)
+{
+    using qualifier_t     = tim::hardware_counters::qualifier;
+    using qualifier_vec_t = std::vector<qualifier_t>;
+    auto*       _arg      = static_cast<info_data*>(arg);
+    const auto* _agent    = _arg->agent;
+    auto*       _data     = _arg->data;
+
+    switch(info.kind)
+    {
+        case ROCPROFILER_INFO_KIND_METRIC:
+        {
+            auto _device_qualifier_sym = JOIN("", ":device=", _agent->dev_index);
+            auto _device_qualifier     = tim::hardware_counters::qualifier{
+                true, static_cast<int>(_agent->dev_index), _device_qualifier_sym,
+                JOIN(" ", "Device", _agent->dev_index)
+            };
+            auto _long_desc = std::string{ info.metric.description };
+            auto _units     = std::string{};
+            auto _pysym     = std::string{};
+            if(info.metric.expr != nullptr)
+            {
+                auto _sym        = JOIN("", info.metric.name, _device_qualifier_sym);
+                auto _short_desc = JOIN("", "Derived counter: ", info.metric.expr);
+                _data->emplace_back(comp::rocm_info_entry(
+                    true, tim::hardware_counters::api::rocm, _data->size(), 0, _sym,
+                    _pysym, _short_desc, _long_desc, _units,
+                    qualifier_vec_t{ _device_qualifier }));
+            }
+            else
+            {
+                if(info.metric.instances == 1)
+                {
+                    auto _sym = JOIN("", info.metric.name, _device_qualifier_sym);
+                    auto _short_desc =
+                        JOIN("", info.metric.name, " on device ", _agent->dev_index);
+                    _data->emplace_back(comp::rocm_info_entry(
+                        true, tim::hardware_counters::api::rocm, _data->size(), 0, _sym,
+                        _pysym, _short_desc, _long_desc, _units,
+                        qualifier_vec_t{ _device_qualifier }));
+                }
+                else
+                {
+                    for(uint32_t i = 0; i < info.metric.instances; ++i)
+                    {
+                        auto _instance_qualifier_sym = JOIN("", '[', i, ']');
+                        auto _instance_qualifier =
+                            tim::hardware_counters::qualifier{ true, static_cast<int>(i),
+                                                               _instance_qualifier_sym,
+                                                               JOIN(" ", "Instance", i) };
+                        auto _sym = JOIN("", info.metric.name, _instance_qualifier_sym,
+                                         _device_qualifier_sym);
+                        auto _short_desc = JOIN("", info.metric.name, " instance ", i,
+                                                " on device ", _agent->dev_index);
+                        _data->emplace_back(comp::rocm_info_entry(
+                            true, tim::hardware_counters::api::rocm, _data->size(), 0,
+                            _sym, _pysym, _short_desc, _long_desc, _units,
+                            qualifier_vec_t{ _device_qualifier, _instance_qualifier }));
+                    }
+                }
+            }
+            break;
+        }
+        default: printf("wrong info kind %u\n", info.kind); return HSA_STATUS_ERROR;
+    }
+    return HSA_STATUS_SUCCESS;
+}
+
+std::vector<comp::rocm_info_entry>
+rocm_metrics()
+{
+    std::vector<comp::rocm_info_entry> _data = {};
+    try
+    {
+        (void) HsaRsrcFactory::Instance();
+    } catch(std::runtime_error& _e)
+    {
+        OMNITRACE_VERBOSE_F(0, "%s\n", _e.what());
+        return _data;
+    }
+
+    // Available GPU agents
+    const unsigned gpu_count = HsaRsrcFactory::Instance().GetCountOfGpuAgents();
+
+    std::vector<AgentInfo*> _gpu_agents(gpu_count, nullptr);
+    for(unsigned i = 0; i < gpu_count; ++i)
+    {
+        const AgentInfo*  _agent   = _gpu_agents[i];
+        const AgentInfo** _agent_p = &_agent;
+        HsaRsrcFactory::Instance().GetGpuAgentInfo(i, _agent_p);
+
+        auto _v = info_data{ _agent, &_data };
+        if(!rocm_check_status(
+               rocprofiler_iterate_info(&_agent->dev_id, ROCPROFILER_INFO_KIND_METRIC,
+                                        info_data_callback, reinterpret_cast<void*>(&_v)),
+               { HSA_STATUS_ERROR_NOT_INITIALIZED }))
+            return _data;
+    }
+
+    auto _settings = tim::settings::shared_instance();
+    if(_settings)
+    {
+        auto ritr = _settings->find("OMNITRACE_ROCM_EVENTS");
+        if(ritr != _settings->end())
+        {
+            auto _rocm_events = ritr->second;
+            if(_rocm_events->get_choices().empty())
+            {
+                std::vector<std::string> _choices = {};
+                _choices.reserve(_data.size());
+                for(auto itr : _data)
+                {
+                    if(!itr.symbol().empty()) _choices.emplace_back(itr.symbol());
+                }
+                _rocm_events->set_choices(_choices);
+            }
+        }
+    }
+
+    return _data;
+}
+
+void
+rocm_initialize()
+{
+    // Available GPU agents
+    const unsigned gpu_count = HsaRsrcFactory::Instance().GetCountOfGpuAgents();
+
+    (void) rocm_metrics();
+
+    // Adding dispatch observer
+    callbacks_arg_t* callbacks_arg = new callbacks_arg_t{};
+    callbacks_arg->pools           = new rocprofiler_pool_t*[gpu_count];
+    for(unsigned gpu_id = 0; gpu_id < gpu_count; gpu_id++)
+    {
+        // Getting profiling features
+        rocprofiler_feature_t* features      = nullptr;
+        unsigned               feature_count = metrics_input(gpu_id, &features);
+
+        if(features)
+        {
+            get_event_names()[gpu_id].clear();
+            get_event_names()[gpu_id].reserve(feature_count);
+            for(unsigned i = 0; i < feature_count; ++i)
+                get_event_names().at(gpu_id).emplace_back(features[i]);
+        }
+
+        // Handler arg
+        handler_arg_t* handler_arg = new handler_arg_t{};
+        handler_arg->features      = features;
+        handler_arg->feature_count = feature_count;
+
+        // Context properties
+        rocprofiler_pool_properties_t properties{};
+        properties.num_entries   = 100;
+        properties.payload_bytes = sizeof(context_entry_t);
+        properties.handler       = rocm_context_handler;
+        properties.handler_arg   = handler_arg;
+
+        // Getting GPU device info
+        const AgentInfo* agent_info = nullptr;
+        if(HsaRsrcFactory::Instance().GetGpuAgentInfo(gpu_id, &agent_info) == false)
+        {
+            fprintf(stderr, "GetGpuAgentInfo failed\n");
+            abort();
+        }
+
+        // Open profiling pool
+        rocprofiler_pool_t* pool = nullptr;
+        uint32_t            mode = 0;  // ROCPROFILER_MODE_SINGLEGROUP
+        rocm_check_status(rocprofiler_pool_open(agent_info->dev_id, features,
+                                                feature_count, &pool, mode, &properties));
+        callbacks_arg->pools[gpu_id] = pool;
+    }
+
+    rocprofiler_queue_callbacks_t callbacks_ptrs{};
+    callbacks_ptrs.dispatch = rocm_dispatch_callback;
+    int err = rocprofiler_set_queue_callbacks(callbacks_ptrs, callbacks_arg);
+    OMNITRACE_VERBOSE_F(3, "err=%d, rocprofiler_set_queue_callbacks\n", err);
+
+    is_setup() = true;
+}
+
+void
+rocm_cleanup()
+{
+    // Unregister dispatch callback
+    rocm_check_status(rocprofiler_remove_queue_callbacks());
+    // close profiling pool
+    // rocm_check_status(rocprofiler_pool_flush(pool));
+    // rocm_check_status(rocprofiler_pool_close(pool));
+}
+
+namespace
+{
+using rocm_event         = comp::rocm_event;
+using rocm_data_t        = comp::rocm_data_t;
+using rocm_metric_type   = comp::rocm_metric_type;
+using rocm_feature_value = comp::rocm_feature_value;
+using rocm_data_tracker  = comp::rocm_data_tracker;
+
+void
+post_process_perfetto()
+{
+    using counter_track = perfetto_counter_track<rocm_event>;
+
+    static bool _once = false;
+    if(_once) return;
+
+    auto _data          = rocm_data_t{};
+    auto _device_data   = std::map<uint32_t, std::vector<rocm_event*>>{};
+    auto _device_fields = std::map<uint32_t, std::vector<std::string_view>>{};
+    auto _device_range  = std::map<uint32_t, std::set<rocm_metric_type>>{};
+
+    for(size_t i = 0; i < OMNITRACE_MAX_THREADS; ++i)
+    {
+        auto& _v = comp::rocm_data(i);
+        if(_v)
+        {
+            _data.reserve(_data.size() + _v->size());
+            for(auto& itr : *_v)
+                _data.emplace_back(itr);
+        }
+    }
+
+    if(_data.empty()) return;
+    _once = true;
+
+    std::sort(_data.begin(), _data.end());
+
+    auto _get_events = [](std::vector<rocm_event*>& _inp, rocm_metric_type _ts) {
+        auto _v = std::vector<rocm_event*>{};
+        for(const auto& itr : _inp)
+        {
+            if(_ts >= itr->entry && _ts <= itr->exit) _v.emplace_back(itr);
+        }
+        return _v;
+    };
+
+    {
+        auto _device_time = std::map<uint32_t, std::set<rocm_metric_type>>{};
+        for(auto& itr : _data)
+        {
+            _device_data[itr.device_id].emplace_back(&itr);
+            _device_time[itr.device_id].emplace(itr.entry);
+            _device_time[itr.device_id].emplace(itr.exit);
+            auto _dev_id = itr.device_id;
+            if(get_use_perfetto() && !counter_track::exists(_dev_id))
+            {
+                auto addendum = [&](auto&& _v) {
+                    return JOIN(" ", "Device", _v, JOIN("", '[', _dev_id, ']'));
+                };
+                for(auto nitr : itr.feature_names)
+                    counter_track::emplace(_dev_id, addendum(nitr));
+            }
+        }
+
+        for(auto& ditr : _device_time)
+        {
+            for(auto itr = ditr.second.begin(); itr != ditr.second.end(); ++itr)
+            {
+                auto _next = std::next(itr);
+                if(_next == ditr.second.end()) continue;
+                _device_range[ditr.first].emplace(((*_next / 2) + (*itr / 2)));
+            }
+        }
+    }
+
+    for(auto& ditr : _device_range)
+    {
+        auto _dev_id = ditr.first;
+        auto _values = std::vector<rocm_feature_value>{};
+        for(const auto& itr : ditr.second)
+        {
+            auto     _v  = _get_events(_device_data[_dev_id], itr);
+            uint64_t _ts = itr;
+            for(auto* vitr : _v)
+            {
+                size_t _n = vitr->feature_values.size();
+                if(_values.empty())
+                {
+                    _values.reserve(_n);
+                    for(size_t i = 0; i < _n; ++i)
+                    {
+                        _values.emplace_back(vitr->feature_values.at(i));
+                    }
+                }
+                else
+                {
+                    for(size_t i = 0; i < _n; ++i)
+                    {
+                        auto _plus = [](auto& _lhs, auto&& _rhs) { _lhs += _rhs; };
+                        std::visit(_plus, _values.at(i), vitr->feature_values.at(i));
+                    }
+                }
+            }
+
+            for(size_t i = 0; i < _values.size(); ++i)
+            {
+                auto _trace_counter = [_dev_id, i, _ts](auto&& _v) {
+                    TRACE_COUNTER("hardware_counter", counter_track::at(_dev_id, i), _ts,
+                                  _v);
+                };
+                std::visit(_trace_counter, _values.at(i));
+            }
+        }
+    }
+}
+
+void
+post_process_timemory()
+{
+    static bool _once = false;
+    if(_once) return;
+
+    auto _data          = rocm_data_t{};
+    auto _device_data   = std::map<uint32_t, std::vector<rocm_event*>>{};
+    auto _device_fields = std::map<uint32_t, std::vector<std::string_view>>{};
+    auto _device_range  = std::map<uint32_t, std::set<rocm_metric_type>>{};
+
+    for(size_t i = 0; i < OMNITRACE_MAX_THREADS; ++i)
+    {
+        auto& _v = comp::rocm_data(i);
+        if(_v)
+        {
+            _data.reserve(_data.size() + _v->size());
+            for(auto& itr : *_v)
+                _data.emplace_back(itr);
+        }
+    }
+
+    if(_data.empty()) return;
+    _once = true;
+
+    std::sort(_data.begin(), _data.end());
+
+    for(auto& itr : _data)
+    {
+        _device_data[itr.device_id].emplace_back(&itr);
+    }
+
+    using storage_type = typename rocm_data_tracker::storage_type;
+    using bundle_type  = tim::lightweight_tuple<rocm_data_tracker>;
+
+    auto        _info            = rocm_metrics();
+    static auto _get_description = [&_info](std::string_view _v) {
+        for(auto& itr : _info)
+        {
+            if(itr.symbol().find(_v) == 0 || itr.short_description().find(_v) == 0)
+            {
+                return itr.long_description();
+            }
+        }
+        return std::string{};
+    };
+
+    struct local_event
+    {
+        rocm_event*              parent   = nullptr;
+        std::vector<local_event> children = {};
+
+        TIMEMORY_DEFAULT_OBJECT(local_event)
+
+        explicit local_event(rocm_event* _v)
+        : parent{ _v }
+        {}
+
+        bool operator()(rocm_event* _v)
+        {
+            OMNITRACE_CI_THROW(!parent, "Error! '%s' has nullptr", __PRETTY_FUNCTION__);
+
+            if(_v->device_id != parent->device_id) return false;
+            if(_v->entry > parent->entry && _v->exit <= parent->exit)
+            {
+                for(auto& itr : children)
+                {
+                    if(itr(_v)) return true;
+                }
+                children.emplace_back(_v);
+                std::sort(children.begin(), children.end());
+                return true;
+            }
+            return false;
+        }
+
+        bool operator<(const local_event& _v) const
+        {
+            OMNITRACE_CI_THROW(!parent, "Error! '%s' has nullptr", __PRETTY_FUNCTION__);
+            OMNITRACE_CI_THROW(!_v.parent, "Error! '%s' passed nullptr",
+                               __PRETTY_FUNCTION__);
+
+            return *parent < *_v.parent;
+        }
+
+        void operator()(int64_t _index, scope::config _scope) const
+        {
+            OMNITRACE_CI_THROW(!parent, "Error! '%s' has nullptr", __PRETTY_FUNCTION__);
+
+            bundle_type _bundle{ parent->name, _scope };
+            _bundle.push(parent->queue_id)
+                .start()
+                .store(parent->feature_values.at(_index));
+
+            for(const auto& itr : children)
+                itr(_index, _scope);
+
+            _bundle.stop().pop(parent->queue_id);
+        }
+    };
+
+    struct local_storage
+    {
+        int64_t                       index              = 0;
+        std::string                   metric_name        = {};
+        std::string                   metric_description = {};
+        std::unique_ptr<storage_type> storage            = {};
+
+        local_storage(uint32_t _devid, size_t _idx, std::string_view _name)
+        : index{ static_cast<int64_t>(_idx) }
+        , metric_name{ _name }
+        , metric_description{ _get_description(metric_name) }
+        {
+            auto _metric_name = std::string{ _name };
+            _metric_name      = std::regex_replace(
+                _metric_name, std::regex{ "(.*)\\[([0-9]+)\\]" }, "$1_$2");
+            storage = std::make_unique<storage_type>(
+                tim::standalone_storage{}, index,
+                JOIN('-', "rocprof", "device", _devid, _metric_name));
+        }
+
+        void operator()(const local_event& _event, scope::config _scope) const
+        {
+            operation::set_storage<rocm_data_tracker>{}(storage.get());
+            _event(index, _scope);
+        }
+
+        void write() const
+        {
+            rocm_data_tracker::label()       = metric_name;
+            rocm_data_tracker::description() = metric_description;
+            storage->write();
+        }
+    };
+
+    auto _local_data = std::map<uint32_t, std::vector<local_event>>{};
+    auto _scope      = scope::get_default();
+
+    for(auto& ditr : _device_data)
+    {
+        auto _storage = std::vector<local_storage>{};
+        for(auto& itr : ditr.second)
+        {
+            auto _n = itr->feature_names.size();
+            if(_n > _storage.size())
+            {
+                _storage.reserve(_n);
+                for(size_t i = _storage.size(); i < _n; ++i)
+                    _storage.emplace_back(ditr.first, i, itr->feature_names.at(i));
+            }
+        }
+
+        auto& _local = _local_data[ditr.first];
+        for(auto& itr : ditr.second)
+        {
+            for(auto& litr : _local)
+            {
+                if(litr(itr))
+                {
+                    goto _bypass_insert;
+                }
+            }
+            _local.emplace_back(itr);
+        _bypass_insert:;
+        }
+
+        for(auto& sitr : _storage)
+        {
+            for(auto& itr : _local_data[ditr.first])
+                sitr(itr, _scope);
+        }
+
+        for(auto& itr : _storage)
+            itr.write();
+    }
+
+    tim::trait::runtime_enabled<omnitrace::rocprofiler::rocm_data_tracker>::set(false);
+}
+}  // namespace
+
+void
+post_process()
+{
+    if(get_use_perfetto()) post_process_perfetto();
+
+    if(get_use_timemory())
+    {
+        auto _manager = tim::manager::master_instance();
+        if(_manager)
+        {
+            _manager->add_cleanup("rocprofiler", &post_process_timemory);
+        }
+        else
+        {
+            post_process_timemory();
+        }
+    }
+}
+}  // namespace rocprofiler
+}  // namespace omnitrace
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler.hpp
new file mode 100644
index 0000000000..3c0e75f1f3
--- /dev/null
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler.hpp
@@ -0,0 +1,88 @@
+// MIT License
+//
+// Copyright (c) 2022 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 "library/components/rocprofiler.hpp"
+#include "library/defines.hpp"
+#include "library/timemory.hpp"
+
+#include <timemory/backends/hardware_counters.hpp>
+#include <timemory/macros.hpp>
+#include <timemory/mpl/concepts.hpp>
+#include <timemory/mpl/macros.hpp>
+
+#include <array>
+#include <atomic>
+#include <cstring>
+#include <dlfcn.h>
+#include <iostream>
+#include <list>
+#include <map>
+#include <string>
+#include <string_view>
+#include <tuple>
+#include <unistd.h>
+#include <utility>
+#include <variant>
+#include <vector>
+
+namespace omnitrace
+{
+namespace rocprofiler
+{
+std::map<uint32_t, std::vector<std::string_view>>
+get_data_labels();
+
+void
+rocm_initialize();
+
+void
+rocm_cleanup();
+
+bool&
+is_setup();
+
+void
+post_process();
+
+std::vector<comp::rocm_info_entry>
+rocm_metrics();
+
+#if !defined(OMNITRACE_USE_ROCPROFILER) || OMNITRACE_USE_ROCPROFILER == 0
+inline void
+post_process()
+{}
+
+inline void
+rocm_cleanup()
+{}
+
+inline std::vector<comp::rocm_info_entry>
+rocm_metrics()
+{
+    return std::vector<comp::rocm_info_entry>{};
+}
+#endif
+
+}  // namespace rocprofiler
+}  // namespace omnitrace
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler/hsa_rsrc_factory.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler/hsa_rsrc_factory.cpp
new file mode 100644
index 0000000000..ed805344cd
--- /dev/null
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler/hsa_rsrc_factory.cpp
@@ -0,0 +1,976 @@
+/******************************************************************************
+Copyright (c) 2018 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 "library/rocprofiler/hsa_rsrc_factory.hpp"
+#include "library/debug.hpp"
+
+#include <timemory/manager.hpp>
+
+#include <rocprofiler.h>
+
+#define PUBLIC_API      __attribute__((visibility("default")))
+#define CONSTRUCTOR_API __attribute__((constructor))
+#define DESTRUCTOR_API  __attribute__((destructor))
+
+#include <atomic>
+#include <cassert>
+#include <dlfcn.h>
+#include <fcntl.h>
+#include <fstream>
+#include <hsa.h>
+#include <hsa_ext_amd.h>
+#include <hsa_ext_finalize.h>
+#include <iostream>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <string>
+#include <sys/mman.h>
+#include <sys/stat.h>
+#include <sys/types.h>
+#include <vector>
+
+// Callback function to get available in the system agents
+hsa_status_t
+HsaRsrcFactory::GetHsaAgentsCallback(hsa_agent_t agent, void* data)
+{
+    hsa_status_t     status     = HSA_STATUS_ERROR;
+    HsaRsrcFactory*  hsa_rsrc   = reinterpret_cast<HsaRsrcFactory*>(data);
+    const AgentInfo* agent_info = hsa_rsrc->AddAgentInfo(agent);
+    if(agent_info != nullptr) status = HSA_STATUS_SUCCESS;
+    return status;
+}
+
+// This function checks to see if the provided
+// pool has the HSA_AMD_SEGMENT_GLOBAL property. If the kern_arg flag is true,
+// the function adds an additional requirement that the pool have the
+// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT property. If kern_arg is false,
+// pools must NOT have this property.
+// Upon finding a pool that meets these conditions, HSA_STATUS_INFO_BREAK is
+// returned. HSA_STATUS_SUCCESS is returned if no errors were encountered, but
+// no pool was found meeting the requirements. If an error is encountered, we
+// return that error.
+static hsa_status_t
+FindGlobalPool(hsa_amd_memory_pool_t pool, void* data, bool kern_arg)
+{
+    hsa_status_t      err;
+    hsa_amd_segment_t segment;
+    uint32_t          flag;
+
+    if(nullptr == data)
+    {
+        return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+    }
+
+    err = HsaRsrcFactory::HsaApi()->hsa_amd_memory_pool_get_info(
+        pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment);
+    CHECK_STATUS("hsa_amd_memory_pool_get_info", err);
+    if(HSA_AMD_SEGMENT_GLOBAL != segment)
+    {
+        return HSA_STATUS_SUCCESS;
+    }
+
+    err = HsaRsrcFactory::HsaApi()->hsa_amd_memory_pool_get_info(
+        pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag);
+    CHECK_STATUS("hsa_amd_memory_pool_get_info", err);
+
+    uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT;
+
+    if((karg_st == 0 && kern_arg) || (karg_st != 0 && !kern_arg))
+    {
+        return HSA_STATUS_SUCCESS;
+    }
+
+    *(reinterpret_cast<hsa_amd_memory_pool_t*>(data)) = pool;
+    return HSA_STATUS_INFO_BREAK;
+}
+
+// This is the call-back function for hsa_amd_agent_iterate_memory_pools() that
+// finds a pool with the properties of HSA_AMD_SEGMENT_GLOBAL and that is NOT
+// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT
+hsa_status_t
+FindStandardPool(hsa_amd_memory_pool_t pool, void* data)
+{
+    return FindGlobalPool(pool, data, false);
+}
+
+// This is the call-back function for hsa_amd_agent_iterate_memory_pools() that
+// finds a pool with the properties of HSA_AMD_SEGMENT_GLOBAL and that IS
+// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT
+hsa_status_t
+FindKernArgPool(hsa_amd_memory_pool_t pool, void* data)
+{
+    return FindGlobalPool(pool, data, true);
+}
+
+// Constructor of the class
+HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa)
+: initialize_hsa_(initialize_hsa)
+{
+    hsa_status_t status;
+
+    cpu_pool_      = nullptr;
+    kern_arg_pool_ = nullptr;
+
+    InitHsaApiTable(nullptr);
+
+    // Initialize the Hsa Runtime
+    if(initialize_hsa_)
+    {
+        status = hsa_api_.hsa_init();
+        CHECK_STATUS("Error in hsa_init", status);
+    }
+
+    // Discover the set of Gpu devices available on the platform
+    status = hsa_api_.hsa_iterate_agents(GetHsaAgentsCallback, this);
+    CHECK_STATUS("Error Calling hsa_iterate_agents", status);
+    if(cpu_pool_ == nullptr)
+        CHECK_STATUS("CPU memory pool is not found", HSA_STATUS_ERROR);
+    if(kern_arg_pool_ == nullptr)
+        CHECK_STATUS("Kern-arg memory pool is not found", HSA_STATUS_ERROR);
+
+    // Get AqlProfile API table
+    aqlprofile_api_ = { nullptr };
+#ifdef ROCP_LD_AQLPROFILE
+    status = LoadAqlProfileLib(&aqlprofile_api_);
+#else
+    status = hsa_api_.hsa_system_get_major_extension_table(
+        HSA_EXTENSION_AMD_AQLPROFILE, hsa_ven_amd_aqlprofile_VERSION_MAJOR,
+        sizeof(aqlprofile_api_), &aqlprofile_api_);
+#endif
+    CHECK_STATUS("aqlprofile API table load failed", status);
+
+    // Get Loader API table
+    loader_api_ = { nullptr };
+    status      = hsa_api_.hsa_system_get_major_extension_table(
+        HSA_EXTENSION_AMD_LOADER, 1, sizeof(loader_api_), &loader_api_);
+    CHECK_STATUS("loader API table query failed", status);
+
+    // Instantiate HSA timer
+    timer_ = new HsaTimer(&hsa_api_);
+    CHECK_STATUS("HSA timer allocation failed",
+                 (timer_ == nullptr) ? HSA_STATUS_ERROR : HSA_STATUS_SUCCESS);
+
+    // System timeout
+    timeout_ = (timeout_ns_ == HsaTimer::TIMESTAMP_MAX)
+                   ? timeout_ns_
+                   : timer_->ns_to_sysclock(timeout_ns_);
+}
+
+// Destructor of the class
+HsaRsrcFactory::~HsaRsrcFactory()
+{
+    delete timer_;
+    for(const auto* p : cpu_list_)
+        delete p;
+    for(const auto* p : gpu_list_)
+        delete p;
+    if(initialize_hsa_)
+    {
+        hsa_status_t status = hsa_api_.hsa_shut_down();
+        try
+        {
+            CHECK_STATUS("Error in hsa_shut_down", status);
+        } catch(std::runtime_error& _e)
+        {
+            fflush(stderr);
+            fprintf(stderr, "%s\n", _e.what());
+            fflush(stderr);
+            abort();
+        }
+    }
+}
+
+void
+HsaRsrcFactory::InitHsaApiTable(HsaApiTable* table)
+{
+    std::lock_guard<mutex_t> lck(mutex_);
+
+    if(hsa_api_.hsa_init == nullptr)
+    {
+        if(table != nullptr)
+        {
+            hsa_api_.hsa_init           = table->core_->hsa_init_fn;
+            hsa_api_.hsa_shut_down      = table->core_->hsa_shut_down_fn;
+            hsa_api_.hsa_agent_get_info = table->core_->hsa_agent_get_info_fn;
+            hsa_api_.hsa_iterate_agents = table->core_->hsa_iterate_agents_fn;
+
+            hsa_api_.hsa_queue_create  = table->core_->hsa_queue_create_fn;
+            hsa_api_.hsa_queue_destroy = table->core_->hsa_queue_destroy_fn;
+            hsa_api_.hsa_queue_load_write_index_relaxed =
+                table->core_->hsa_queue_load_write_index_relaxed_fn;
+            hsa_api_.hsa_queue_store_write_index_relaxed =
+                table->core_->hsa_queue_store_write_index_relaxed_fn;
+            hsa_api_.hsa_queue_load_read_index_relaxed =
+                table->core_->hsa_queue_load_read_index_relaxed_fn;
+
+            hsa_api_.hsa_signal_create        = table->core_->hsa_signal_create_fn;
+            hsa_api_.hsa_signal_destroy       = table->core_->hsa_signal_destroy_fn;
+            hsa_api_.hsa_signal_load_relaxed  = table->core_->hsa_signal_load_relaxed_fn;
+            hsa_api_.hsa_signal_store_relaxed = table->core_->hsa_signal_store_relaxed_fn;
+            hsa_api_.hsa_signal_wait_scacquire =
+                table->core_->hsa_signal_wait_scacquire_fn;
+            hsa_api_.hsa_signal_store_screlease =
+                table->core_->hsa_signal_store_screlease_fn;
+
+            hsa_api_.hsa_code_object_reader_create_from_file =
+                table->core_->hsa_code_object_reader_create_from_file_fn;
+            hsa_api_.hsa_executable_create_alt =
+                table->core_->hsa_executable_create_alt_fn;
+            hsa_api_.hsa_executable_load_agent_code_object =
+                table->core_->hsa_executable_load_agent_code_object_fn;
+            hsa_api_.hsa_executable_freeze = table->core_->hsa_executable_freeze_fn;
+            hsa_api_.hsa_executable_get_symbol =
+                table->core_->hsa_executable_get_symbol_fn;
+            hsa_api_.hsa_executable_symbol_get_info =
+                table->core_->hsa_executable_symbol_get_info_fn;
+            hsa_api_.hsa_executable_iterate_symbols =
+                table->core_->hsa_executable_iterate_symbols_fn;
+
+            hsa_api_.hsa_system_get_info = table->core_->hsa_system_get_info_fn;
+            hsa_api_.hsa_system_get_major_extension_table =
+                table->core_->hsa_system_get_major_extension_table_fn;
+
+            hsa_api_.hsa_amd_agent_iterate_memory_pools =
+                table->amd_ext_->hsa_amd_agent_iterate_memory_pools_fn;
+            hsa_api_.hsa_amd_memory_pool_get_info =
+                table->amd_ext_->hsa_amd_memory_pool_get_info_fn;
+            hsa_api_.hsa_amd_memory_pool_allocate =
+                table->amd_ext_->hsa_amd_memory_pool_allocate_fn;
+            hsa_api_.hsa_amd_agents_allow_access =
+                table->amd_ext_->hsa_amd_agents_allow_access_fn;
+            hsa_api_.hsa_amd_memory_async_copy =
+                table->amd_ext_->hsa_amd_memory_async_copy_fn;
+
+            hsa_api_.hsa_amd_signal_async_handler =
+                table->amd_ext_->hsa_amd_signal_async_handler_fn;
+            hsa_api_.hsa_amd_profiling_set_profiler_enabled =
+                table->amd_ext_->hsa_amd_profiling_set_profiler_enabled_fn;
+            hsa_api_.hsa_amd_profiling_get_async_copy_time =
+                table->amd_ext_->hsa_amd_profiling_get_async_copy_time_fn;
+            hsa_api_.hsa_amd_profiling_get_dispatch_time =
+                table->amd_ext_->hsa_amd_profiling_get_dispatch_time_fn;
+        }
+        else
+        {
+            hsa_api_.hsa_init           = hsa_init;
+            hsa_api_.hsa_shut_down      = hsa_shut_down;
+            hsa_api_.hsa_agent_get_info = hsa_agent_get_info;
+            hsa_api_.hsa_iterate_agents = hsa_iterate_agents;
+
+            hsa_api_.hsa_queue_create  = hsa_queue_create;
+            hsa_api_.hsa_queue_destroy = hsa_queue_destroy;
+            hsa_api_.hsa_queue_load_write_index_relaxed =
+                hsa_queue_load_write_index_relaxed;
+            hsa_api_.hsa_queue_store_write_index_relaxed =
+                hsa_queue_store_write_index_relaxed;
+            hsa_api_.hsa_queue_load_read_index_relaxed =
+                hsa_queue_load_read_index_relaxed;
+
+            hsa_api_.hsa_signal_create          = hsa_signal_create;
+            hsa_api_.hsa_signal_destroy         = hsa_signal_destroy;
+            hsa_api_.hsa_signal_load_relaxed    = hsa_signal_load_relaxed;
+            hsa_api_.hsa_signal_store_relaxed   = hsa_signal_store_relaxed;
+            hsa_api_.hsa_signal_wait_scacquire  = hsa_signal_wait_scacquire;
+            hsa_api_.hsa_signal_store_screlease = hsa_signal_store_screlease;
+
+            hsa_api_.hsa_code_object_reader_create_from_file =
+                hsa_code_object_reader_create_from_file;
+            hsa_api_.hsa_executable_create_alt = hsa_executable_create_alt;
+            hsa_api_.hsa_executable_load_agent_code_object =
+                hsa_executable_load_agent_code_object;
+            hsa_api_.hsa_executable_freeze          = hsa_executable_freeze;
+            hsa_api_.hsa_executable_get_symbol      = hsa_executable_get_symbol;
+            hsa_api_.hsa_executable_symbol_get_info = hsa_executable_symbol_get_info;
+            hsa_api_.hsa_executable_iterate_symbols = hsa_executable_iterate_symbols;
+
+            hsa_api_.hsa_system_get_info = hsa_system_get_info;
+            hsa_api_.hsa_system_get_major_extension_table =
+                hsa_system_get_major_extension_table;
+
+            hsa_api_.hsa_amd_agent_iterate_memory_pools =
+                hsa_amd_agent_iterate_memory_pools;
+            hsa_api_.hsa_amd_memory_pool_get_info = hsa_amd_memory_pool_get_info;
+            hsa_api_.hsa_amd_memory_pool_allocate = hsa_amd_memory_pool_allocate;
+            hsa_api_.hsa_amd_agents_allow_access  = hsa_amd_agents_allow_access;
+            hsa_api_.hsa_amd_memory_async_copy    = hsa_amd_memory_async_copy;
+
+            hsa_api_.hsa_amd_signal_async_handler = hsa_amd_signal_async_handler;
+            hsa_api_.hsa_amd_profiling_set_profiler_enabled =
+                hsa_amd_profiling_set_profiler_enabled;
+            hsa_api_.hsa_amd_profiling_get_async_copy_time =
+                hsa_amd_profiling_get_async_copy_time;
+            hsa_api_.hsa_amd_profiling_get_dispatch_time =
+                hsa_amd_profiling_get_dispatch_time;
+        }
+    }
+}
+
+hsa_status_t
+HsaRsrcFactory::LoadAqlProfileLib(aqlprofile_pfn_t* api)
+{
+    void* handle = dlopen(kAqlProfileLib, RTLD_NOW);
+    if(handle == nullptr)
+    {
+        fprintf(stderr, "Loading '%s' failed, %s\n", kAqlProfileLib, dlerror());
+        return HSA_STATUS_ERROR;
+    }
+    dlerror(); /* Clear any existing error */
+
+    api->hsa_ven_amd_aqlprofile_error_string =
+        (decltype(::hsa_ven_amd_aqlprofile_error_string)*) dlsym(
+            handle, "hsa_ven_amd_aqlprofile_error_string");
+    api->hsa_ven_amd_aqlprofile_validate_event =
+        (decltype(::hsa_ven_amd_aqlprofile_validate_event)*) dlsym(
+            handle, "hsa_ven_amd_aqlprofile_validate_event");
+    api->hsa_ven_amd_aqlprofile_start = (decltype(::hsa_ven_amd_aqlprofile_start)*) dlsym(
+        handle, "hsa_ven_amd_aqlprofile_start");
+    api->hsa_ven_amd_aqlprofile_stop = (decltype(::hsa_ven_amd_aqlprofile_stop)*) dlsym(
+        handle, "hsa_ven_amd_aqlprofile_stop");
+#ifdef AQLPROF_NEW_API
+    api->hsa_ven_amd_aqlprofile_read = (decltype(::hsa_ven_amd_aqlprofile_read)*) dlsym(
+        handle, "hsa_ven_amd_aqlprofile_read");
+#endif
+    api->hsa_ven_amd_aqlprofile_legacy_get_pm4 =
+        (decltype(::hsa_ven_amd_aqlprofile_legacy_get_pm4)*) dlsym(
+            handle, "hsa_ven_amd_aqlprofile_legacy_get_pm4");
+    api->hsa_ven_amd_aqlprofile_get_info =
+        (decltype(::hsa_ven_amd_aqlprofile_get_info)*) dlsym(
+            handle, "hsa_ven_amd_aqlprofile_get_info");
+    api->hsa_ven_amd_aqlprofile_iterate_data =
+        (decltype(::hsa_ven_amd_aqlprofile_iterate_data)*) dlsym(
+            handle, "hsa_ven_amd_aqlprofile_iterate_data");
+
+    return HSA_STATUS_SUCCESS;
+}
+
+// Add system agent info
+const AgentInfo*
+HsaRsrcFactory::AddAgentInfo(const hsa_agent_t agent)
+{
+    // Determine if device is a Gpu agent
+    hsa_status_t status;
+    AgentInfo*   agent_info = nullptr;
+
+    hsa_device_type_t type;
+    status = hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type);
+    CHECK_STATUS("Error Calling hsa_agent_get_info", status);
+
+    if(type == HSA_DEVICE_TYPE_CPU)
+    {
+        agent_info            = new AgentInfo{};
+        agent_info->dev_id    = agent;
+        agent_info->dev_type  = HSA_DEVICE_TYPE_CPU;
+        agent_info->dev_index = cpu_list_.size();
+
+        status = hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindStandardPool,
+                                                             &agent_info->cpu_pool);
+        if((status == HSA_STATUS_INFO_BREAK) && (cpu_pool_ == nullptr))
+            cpu_pool_ = &agent_info->cpu_pool;
+        status = hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindKernArgPool,
+                                                             &agent_info->kern_arg_pool);
+        if((status == HSA_STATUS_INFO_BREAK) && (kern_arg_pool_ == nullptr))
+            kern_arg_pool_ = &agent_info->kern_arg_pool;
+        agent_info->gpu_pool = {};
+
+        cpu_list_.push_back(agent_info);
+        cpu_agents_.push_back(agent);
+    }
+
+    if(type == HSA_DEVICE_TYPE_GPU)
+    {
+        agent_info           = new AgentInfo{};
+        agent_info->dev_id   = agent;
+        agent_info->dev_type = HSA_DEVICE_TYPE_GPU;
+        hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, agent_info->name);
+        strncpy(agent_info->gfxip, agent_info->name, 4);
+        agent_info->gfxip[4] = '\0';
+        hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_WAVEFRONT_SIZE,
+                                    &agent_info->max_wave_size);
+        hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_QUEUE_MAX_SIZE,
+                                    &agent_info->max_queue_size);
+        hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_info->profile);
+        agent_info->is_apu = (agent_info->profile == HSA_PROFILE_FULL) ? true : false;
+        hsa_api_.hsa_agent_get_info(
+            agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT),
+            &agent_info->cu_num);
+        hsa_api_.hsa_agent_get_info(
+            agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU),
+            &agent_info->waves_per_cu);
+        hsa_api_.hsa_agent_get_info(
+            agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU),
+            &agent_info->simds_per_cu);
+        hsa_api_.hsa_agent_get_info(
+            agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ENGINES),
+            &agent_info->se_num);
+        hsa_api_.hsa_agent_get_info(
+            agent,
+            static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ARRAYS_PER_SE),
+            &agent_info->shader_arrays_per_se);
+
+        agent_info->cpu_pool      = {};
+        agent_info->kern_arg_pool = {};
+        status = hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindStandardPool,
+                                                             &agent_info->gpu_pool);
+        CHECK_ITER_STATUS("hsa_amd_agent_iterate_memory_pools(gpu pool)", status);
+
+        // GFX8 and GFX9 SGPR/VGPR block sizes
+        agent_info->sgpr_block_dflt = (strcmp(agent_info->gfxip, "gfx8") == 0) ? 1 : 2;
+        agent_info->sgpr_block_size = 8;
+        agent_info->vgpr_block_size = 4;
+
+        // Set GPU index
+        agent_info->dev_index = gpu_list_.size();
+        gpu_list_.push_back(agent_info);
+        gpu_agents_.push_back(agent);
+    }
+
+    if(agent_info) agent_map_[agent.handle] = agent_info;
+
+    return agent_info;
+}
+
+// Return systen agent info
+const AgentInfo*
+HsaRsrcFactory::GetAgentInfo(const hsa_agent_t agent)
+{
+    const AgentInfo* agent_info = nullptr;
+    auto             it         = agent_map_.find(agent.handle);
+    if(it != agent_map_.end())
+    {
+        agent_info = it->second;
+    }
+    return agent_info;
+}
+
+// Get the count of Hsa Gpu Agents available on the platform
+//
+// @return uint32_t Number of Gpu agents on platform
+//
+uint32_t
+HsaRsrcFactory::GetCountOfGpuAgents()
+{
+    return uint32_t(gpu_list_.size());
+}
+
+// Get the count of Hsa Cpu Agents available on the platform
+//
+// @return uint32_t Number of Cpu agents on platform
+//
+uint32_t
+HsaRsrcFactory::GetCountOfCpuAgents()
+{
+    return uint32_t(cpu_list_.size());
+}
+
+// Get the AgentInfo handle of a Gpu device
+//
+// @param idx Gpu Agent at specified index
+//
+// @param agent_info Output parameter updated with AgentInfo
+//
+// @return bool true if successful, false otherwise
+//
+bool
+HsaRsrcFactory::GetGpuAgentInfo(uint32_t idx, const AgentInfo** agent_info)
+{
+    // Determine if request is valid
+    uint32_t size = uint32_t(gpu_list_.size());
+    if(idx >= size)
+    {
+        return false;
+    }
+
+    // Copy AgentInfo from specified index
+    *agent_info = gpu_list_[idx];
+
+    return true;
+}
+
+// Get the AgentInfo handle of a Cpu device
+//
+// @param idx Cpu Agent at specified index
+//
+// @param agent_info Output parameter updated with AgentInfo
+//
+// @return bool true if successful, false otherwise
+//
+bool
+HsaRsrcFactory::GetCpuAgentInfo(uint32_t idx, const AgentInfo** agent_info)
+{
+    // Determine if request is valid
+    uint32_t size = uint32_t(cpu_list_.size());
+    if(idx >= size)
+    {
+        return false;
+    }
+
+    // Copy AgentInfo from specified index
+    *agent_info = cpu_list_[idx];
+    return true;
+}
+
+// Create a Queue object and return its handle. The queue object is expected
+// to support user requested number of Aql dispatch packets.
+//
+// @param agent_info Gpu Agent on which to create a queue object
+//
+// @param num_Pkts Number of packets to be held by queue
+//
+// @param queue Output parameter updated with handle of queue object
+//
+// @return bool true if successful, false otherwise
+//
+bool
+HsaRsrcFactory::CreateQueue(const AgentInfo* agent_info, uint32_t num_pkts,
+                            hsa_queue_t** queue)
+{
+    hsa_status_t status;
+    status = hsa_api_.hsa_queue_create(agent_info->dev_id, num_pkts, HSA_QUEUE_TYPE_MULTI,
+                                       nullptr, nullptr, UINT32_MAX, UINT32_MAX, queue);
+    return (status == HSA_STATUS_SUCCESS);
+}
+
+// Create a Signal object and return its handle.
+// @param value Initial value of signal object
+// @param signal Output parameter updated with handle of signal object
+// @return bool true if successful, false otherwise
+bool
+HsaRsrcFactory::CreateSignal(uint32_t value, hsa_signal_t* signal)
+{
+    hsa_status_t status;
+    status = hsa_api_.hsa_signal_create(value, 0, nullptr, signal);
+    return (status == HSA_STATUS_SUCCESS);
+}
+
+// Allocate memory for use by a kernel of specified size in specified
+// agent's memory region.
+// @param agent_info Agent from whose memory region to allocate
+// @param size Size of memory in terms of bytes
+// @return uint8_t* Pointer to buffer, null if allocation fails.
+uint8_t*
+HsaRsrcFactory::AllocateLocalMemory(const AgentInfo* agent_info, size_t size)
+{
+    hsa_status_t status = HSA_STATUS_ERROR;
+    uint8_t*     buffer = nullptr;
+    size                = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
+    status       = hsa_api_.hsa_amd_memory_pool_allocate(agent_info->gpu_pool, size, 0,
+                                                   reinterpret_cast<void**>(&buffer));
+    uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : nullptr;
+    return ptr;
+}
+
+// Allocate memory to pass kernel parameters.
+// Memory is alocated accessible for all CPU agents and for GPU given by AgentInfo
+// parameter.
+// @param agent_info Agent from whose memory region to allocate
+// @param size Size of memory in terms of bytes
+// @return uint8_t* Pointer to buffer, null if allocation fails.
+uint8_t*
+HsaRsrcFactory::AllocateKernArgMemory(const AgentInfo* agent_info, size_t size)
+{
+    hsa_status_t status = HSA_STATUS_ERROR;
+    uint8_t*     buffer = nullptr;
+    if(!cpu_agents_.empty())
+    {
+        size   = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
+        status = hsa_api_.hsa_amd_memory_pool_allocate(*kern_arg_pool_, size, 0,
+                                                       reinterpret_cast<void**>(&buffer));
+        // Both the CPU and GPU can access the kernel arguments
+        if(status == HSA_STATUS_SUCCESS)
+        {
+            hsa_agent_t ag_list[1] = { agent_info->dev_id };
+            status = hsa_api_.hsa_amd_agents_allow_access(1, ag_list, nullptr, buffer);
+        }
+    }
+    uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : nullptr;
+    return ptr;
+}
+
+// Allocate system memory accessible by both CPU and GPU
+// @param agent_info Agent from whose memory region to allocate
+// @param size Size of memory in terms of bytes
+// @return uint8_t* Pointer to buffer, null if allocation fails.
+uint8_t*
+HsaRsrcFactory::AllocateSysMemory(const AgentInfo* agent_info, size_t size)
+{
+    hsa_status_t status = HSA_STATUS_ERROR;
+    uint8_t*     buffer = nullptr;
+    size                = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
+    if(!cpu_agents_.empty())
+    {
+        status = hsa_api_.hsa_amd_memory_pool_allocate(*cpu_pool_, size, 0,
+                                                       reinterpret_cast<void**>(&buffer));
+        // Both the CPU and GPU can access the memory
+        if(status == HSA_STATUS_SUCCESS)
+        {
+            hsa_agent_t ag_list[1] = { agent_info->dev_id };
+            status = hsa_api_.hsa_amd_agents_allow_access(1, ag_list, nullptr, buffer);
+        }
+    }
+    uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : nullptr;
+    return ptr;
+}
+
+// Allocate memory for command buffer.
+// @param agent_info Agent from whose memory region to allocate
+// @param size Size of memory in terms of bytes
+// @return uint8_t* Pointer to buffer, null if allocation fails.
+uint8_t*
+HsaRsrcFactory::AllocateCmdMemory(const AgentInfo* agent_info, size_t size)
+{
+    size         = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
+    uint8_t* ptr = (agent_info->is_apu && CMD_MEMORY_MMAP)
+                       ? reinterpret_cast<uint8_t*>(
+                             mmap(nullptr, size, PROT_READ | PROT_WRITE | PROT_EXEC,
+                                  MAP_SHARED | MAP_ANONYMOUS, 0, 0))
+                       : AllocateSysMemory(agent_info, size);
+    return ptr;
+}
+
+// Wait signal
+void
+HsaRsrcFactory::SignalWait(const hsa_signal_t& signal) const
+{
+    while(1)
+    {
+        const hsa_signal_value_t signal_value = hsa_api_.hsa_signal_wait_scacquire(
+            signal, HSA_SIGNAL_CONDITION_LT, 1, timeout_, HSA_WAIT_STATE_BLOCKED);
+        if(signal_value == 0)
+        {
+            break;
+        }
+        else
+        {
+            CHECK_STATUS("hsa_signal_wait_scacquire()", HSA_STATUS_ERROR);
+        }
+    }
+}
+
+// Wait signal with signal value restore
+void
+HsaRsrcFactory::SignalWaitRestore(const hsa_signal_t&       signal,
+                                  const hsa_signal_value_t& signal_value) const
+{
+    SignalWait(signal);
+    hsa_api_.hsa_signal_store_relaxed(const_cast<hsa_signal_t&>(signal), signal_value);
+}
+
+// Copy data from GPU to host memory
+bool
+HsaRsrcFactory::Memcpy(const hsa_agent_t& agent, void* dst, const void* src, size_t size)
+{
+    hsa_status_t status = HSA_STATUS_ERROR;
+    if(!cpu_agents_.empty())
+    {
+        hsa_signal_t s = {};
+        status         = hsa_api_.hsa_signal_create(1, 0, nullptr, &s);
+        CHECK_STATUS("hsa_signal_create()", status);
+        status = hsa_api_.hsa_amd_memory_async_copy(dst, cpu_agents_[0], src, agent, size,
+                                                    0, nullptr, s);
+        CHECK_STATUS("hsa_amd_memory_async_copy()", status);
+        SignalWait(s);
+        status = hsa_api_.hsa_signal_destroy(s);
+        CHECK_STATUS("hsa_signal_destroy()", status);
+    }
+    return (status == HSA_STATUS_SUCCESS);
+}
+
+bool
+HsaRsrcFactory::Memcpy(const AgentInfo* agent_info, void* dst, const void* src,
+                       size_t size)
+{
+    return Memcpy(agent_info->dev_id, dst, src, size);
+}
+
+// Memory free method
+bool
+HsaRsrcFactory::FreeMemory(void* ptr)
+{
+    const hsa_status_t status = hsa_memory_free(ptr);
+    CHECK_STATUS("hsa_memory_free", status);
+    return (status == HSA_STATUS_SUCCESS);
+}
+
+// Loads an Assembled Brig file and Finalizes it into Device Isa
+// @param agent_info Gpu device for which to finalize
+// @param brig_path File path of the Assembled Brig file
+// @param kernel_name Name of the kernel to finalize
+// @param code_desc Handle of finalized Code Descriptor that could
+// be used to submit for execution
+// @return bool true if successful, false otherwise
+bool
+HsaRsrcFactory::LoadAndFinalize(const AgentInfo* agent_info, const char* brig_path,
+                                const char* kernel_name, hsa_executable_t* executable,
+                                hsa_executable_symbol_t* code_desc)
+{
+    hsa_status_t status = HSA_STATUS_ERROR;
+
+    // Build the code object filename
+    std::string filename(brig_path);
+    std::clog << "Code object filename: " << filename << std::endl;
+
+    // Open the file containing code object
+    hsa_file_t file_handle = open(filename.c_str(), O_RDONLY);
+    if(file_handle == -1)
+    {
+        std::cerr << "Error: failed to load '" << filename << "'" << std::endl;
+        assert(false);
+        return false;
+    }
+
+    // Create code object reader
+    hsa_code_object_reader_t code_obj_rdr = { 0 };
+    status = hsa_api_.hsa_code_object_reader_create_from_file(file_handle, &code_obj_rdr);
+    if(status != HSA_STATUS_SUCCESS)
+    {
+        std::cerr << "Failed to create code object reader '" << filename << "'"
+                  << std::endl;
+        return false;
+    }
+
+    // Create executable.
+    status = hsa_api_.hsa_executable_create_alt(
+        HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, executable);
+    CHECK_STATUS("Error in creating executable object", status);
+
+    // Load code object.
+    status = hsa_api_.hsa_executable_load_agent_code_object(
+        *executable, agent_info->dev_id, code_obj_rdr, nullptr, nullptr);
+    CHECK_STATUS("Error in loading executable object", status);
+
+    // Freeze executable.
+    status = hsa_api_.hsa_executable_freeze(*executable, "");
+    CHECK_STATUS("Error in freezing executable object", status);
+
+    // Get symbol handle.
+    hsa_executable_symbol_t kernelSymbol;
+    status = hsa_api_.hsa_executable_get_symbol(*executable, nullptr, kernel_name,
+                                                agent_info->dev_id, 0, &kernelSymbol);
+    CHECK_STATUS("Error in looking up kernel symbol", status);
+
+    // Update output parameter
+    *code_desc = kernelSymbol;
+    return true;
+}
+
+// Print the various fields of Hsa Gpu Agents
+bool
+HsaRsrcFactory::PrintGpuAgents(const std::string&)
+{
+    std::cout << std::flush;
+    // std::clog << header << " :" << std::endl;
+
+    char key[1024], value[1024];
+
+    const AgentInfo* agent_info;
+    int              size = uint32_t(gpu_list_.size());
+    for(int idx = 0; idx < size; idx++)
+    {
+        agent_info = gpu_list_[idx];
+
+        /*  std::clog << "> agent[" << idx << "] :" << std::endl;
+            std::clog << ">> Name : " << agent_info->name << std::endl;
+            std::clog << ">> APU : " << agent_info->is_apu << std::endl;
+            std::clog << ">> HSAIL profile : " << agent_info->profile << std::endl;
+            std::clog << ">> Max Wave Size : " << agent_info->max_wave_size << std::endl;
+            std::clog << ">> Max Queue Size : " << agent_info->max_queue_size <<
+           std::endl; std::clog << ">> CU number : " << agent_info->cu_num << std::endl;
+            std::clog << ">> Waves per CU : " << agent_info->waves_per_cu << std::endl;
+            std::clog << ">> SIMDs per CU : " << agent_info->simds_per_cu << std::endl;
+            std::clog << ">> SE number : " << agent_info->se_num << std::endl;
+            std::clog << ">> Shader Arrays per SE : " << agent_info->shader_arrays_per_se
+           << std::endl;
+        */
+
+#define OMNITRACE_METADATA(...) tim::manager::add_metadata(__VA_ARGS__)
+
+        sprintf(key, "ROCM_AGENT_%d_NAME", idx);
+        sprintf(value, "%s", agent_info->name);
+        OMNITRACE_METADATA(key, value);
+
+        sprintf(key, "ROCM_AGENT_%d_IS_APU", idx);
+        sprintf(value, "%d", static_cast<int>(agent_info->is_apu));
+        OMNITRACE_METADATA(key, value);
+
+        sprintf(key, "ROCM_AGENT_%d_HSA_PROFILE", idx);
+        sprintf(value, "%d", agent_info->profile);
+        OMNITRACE_METADATA(key, value);
+
+        sprintf(key, "ROCM_AGENT_%d_MAX_WAVE_SIZE", idx);
+        sprintf(value, "%d", agent_info->max_wave_size);
+        OMNITRACE_METADATA(key, value);
+
+        sprintf(key, "ROCM_AGENT_%d_MAX_QUEUE_SIZE", idx);
+        sprintf(value, "%d", agent_info->max_queue_size);
+        OMNITRACE_METADATA(key, value);
+
+        sprintf(key, "ROCM_AGENT_%d_CU_NUMBER", idx);
+        sprintf(value, "%d", agent_info->cu_num);
+        OMNITRACE_METADATA(key, value);
+
+        sprintf(key, "ROCM_AGENT_%d_WAVES_PER_CU", idx);
+        sprintf(value, "%d", agent_info->waves_per_cu);
+        OMNITRACE_METADATA(key, value);
+
+        sprintf(key, "ROCM_AGENT_%d_SIMDs_PER_CU", idx);
+        sprintf(value, "%d", agent_info->simds_per_cu);
+        OMNITRACE_METADATA(key, value);
+
+        sprintf(key, "ROCM_AGENT_%d_SE_NUMBER", idx);
+        sprintf(value, "%d", agent_info->se_num);
+        OMNITRACE_METADATA(key, value);
+
+        sprintf(key, "ROCM_AGENT_%d_SHADER_ARRAYS_PER_SE", idx);
+        sprintf(value, "%d", agent_info->shader_arrays_per_se);
+        OMNITRACE_METADATA(key, value);
+    }
+    return true;
+}
+
+uint64_t
+HsaRsrcFactory::Submit(hsa_queue_t* queue, const void* packet)
+{
+    const uint32_t slot_size_b = CMD_SLOT_SIZE_B;
+
+    // adevance command queue
+    const uint64_t write_idx = hsa_api_.hsa_queue_load_write_index_relaxed(queue);
+    hsa_api_.hsa_queue_store_write_index_relaxed(queue, write_idx + 1);
+    while((write_idx - hsa_api_.hsa_queue_load_read_index_relaxed(queue)) >= queue->size)
+    {
+        sched_yield();
+    }
+
+    uint32_t  slot_idx   = (uint32_t)(write_idx % queue->size);
+    uint32_t* queue_slot = reinterpret_cast<uint32_t*>((uintptr_t)(queue->base_address) +
+                                                       (slot_idx * slot_size_b));
+    const uint32_t* slot_data = reinterpret_cast<const uint32_t*>(packet);
+
+    // Copy buffered commands into the queue slot.
+    // Overwrite the AQL invalid header (first dword) last.
+    // This prevents the slot from being read until it's fully written.
+    memcpy(&queue_slot[1], &slot_data[1], slot_size_b - sizeof(uint32_t));
+    std::atomic<uint32_t>* header_atomic_ptr =
+        reinterpret_cast<std::atomic<uint32_t>*>(&queue_slot[0]);
+    header_atomic_ptr->store(slot_data[0], std::memory_order_release);
+
+    // ringdoor bell
+    hsa_api_.hsa_signal_store_relaxed(queue->doorbell_signal, write_idx);
+
+    return write_idx;
+}
+
+uint64_t
+HsaRsrcFactory::Submit(hsa_queue_t* queue, const void* packet, size_t size_bytes)
+{
+    const uint32_t slot_size_b = CMD_SLOT_SIZE_B;
+    if((size_bytes & (slot_size_b - 1)) != 0)
+    {
+        fprintf(stderr, "HsaRsrcFactory::Submit: Bad packet size %zx\n", size_bytes);
+        abort();
+    }
+
+    const char* begin     = reinterpret_cast<const char*>(packet);
+    const char* end       = begin + size_bytes;
+    uint64_t    write_idx = 0;
+    for(const char* ptr = begin; ptr < end; ptr += slot_size_b)
+    {
+        write_idx = Submit(queue, ptr);
+    }
+
+    return write_idx;
+}
+
+const char*
+HsaRsrcFactory::GetKernelName(uint64_t addr)
+{
+    std::lock_guard<mutex_t> lck(mutex_);
+    const auto               it = symbols_map_->find(addr);
+    if(it == symbols_map_->end())
+    {
+        fprintf(stderr, "HsaRsrcFactory::kernel addr (0x%lx) is not found\n", addr);
+        abort();
+    }
+    return strdup(it->second);
+}
+
+void
+HsaRsrcFactory::EnableExecutableTracking(HsaApiTable* table)
+{
+    std::lock_guard<mutex_t> lck(mutex_);
+    executable_tracking_on_                = true;
+    table->core_->hsa_executable_freeze_fn = hsa_executable_freeze_interceptor;
+}
+
+hsa_status_t
+HsaRsrcFactory::executable_symbols_cb(hsa_executable_t /*exec*/,
+                                      hsa_executable_symbol_t symbol, void* /*data*/)
+{
+    hsa_symbol_kind_t value  = (hsa_symbol_kind_t) 0;
+    hsa_status_t      status = hsa_api_.hsa_executable_symbol_get_info(
+        symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &value);
+    CHECK_STATUS("Error in getting symbol info", status);
+    if(value == HSA_SYMBOL_KIND_KERNEL)
+    {
+        uint64_t addr = 0;
+        uint32_t len  = 0;
+        status        = hsa_api_.hsa_executable_symbol_get_info(
+            symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &addr);
+        CHECK_STATUS("Error in getting kernel object", status);
+        status = hsa_api_.hsa_executable_symbol_get_info(
+            symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &len);
+        CHECK_STATUS("Error in getting name len", status);
+        char* name = new char[len + 1];
+        status     = hsa_api_.hsa_executable_symbol_get_info(
+            symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
+        CHECK_STATUS("Error in getting kernel name", status);
+        name[len] = 0;
+        auto ret  = symbols_map_->insert({ addr, name });
+        if(ret.second == false)
+        {
+            delete[] ret.first->second;
+            ret.first->second = name;
+        }
+    }
+    return HSA_STATUS_SUCCESS;
+}
+
+hsa_status_t
+HsaRsrcFactory::hsa_executable_freeze_interceptor(hsa_executable_t executable,
+                                                  const char*      options)
+{
+    std::lock_guard<mutex_t> lck(mutex_);
+    if(symbols_map_ == nullptr) symbols_map_ = new symbols_map_t;
+    hsa_status_t status = hsa_api_.hsa_executable_iterate_symbols(
+        executable, executable_symbols_cb, nullptr);
+    CHECK_STATUS("Error in iterating executable symbols", status);
+    return hsa_api_.hsa_executable_freeze(executable, options);
+    ;
+}
+
+std::atomic<HsaRsrcFactory*>   HsaRsrcFactory::instance_{};
+HsaRsrcFactory::mutex_t        HsaRsrcFactory::mutex_;
+HsaRsrcFactory::timestamp_t    HsaRsrcFactory::timeout_ns_ = HsaTimer::TIMESTAMP_MAX;
+hsa_pfn_t                      HsaRsrcFactory::hsa_api_{};
+bool                           HsaRsrcFactory::executable_tracking_on_ = false;
+HsaRsrcFactory::symbols_map_t* HsaRsrcFactory::symbols_map_            = nullptr;
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler/hsa_rsrc_factory.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler/hsa_rsrc_factory.hpp
new file mode 100644
index 0000000000..5c2a54270f
--- /dev/null
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler/hsa_rsrc_factory.hpp
@@ -0,0 +1,470 @@
+// MIT License
+//
+// Copyright (c) 2022 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
+
+#define AMD_INTERNAL_BUILD 1
+
+#include <hsa.h>
+#include <hsa_api_trace.h>
+#include <hsa_ext_amd.h>
+#include <hsa_ext_finalize.h>
+#include <hsa_ven_amd_aqlprofile.h>
+#include <hsa_ven_amd_loader.h>
+
+#include <atomic>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <iostream>
+#include <map>
+#include <mutex>
+#include <string>
+#include <vector>
+
+#define HSA_ARGUMENT_ALIGN_BYTES 16
+#define HSA_QUEUE_ALIGN_BYTES    64
+#define HSA_PACKET_ALIGN_BYTES   64
+#define HSA_MESSAGE_LENGTH       4096
+
+#define CHECK_STATUS(msg, status)                                                        \
+    do                                                                                   \
+    {                                                                                    \
+        if((status) != HSA_STATUS_SUCCESS)                                               \
+        {                                                                                \
+            const char* emsg = 0;                                                        \
+            hsa_status_string(status, &emsg);                                            \
+            char _buffer[HSA_MESSAGE_LENGTH];                                            \
+            snprintf(_buffer, HSA_MESSAGE_LENGTH - 1, "%s: %s", msg,                     \
+                     emsg ? emsg : "<unknown error>");                                   \
+            throw std::runtime_error(_buffer);                                           \
+        }                                                                                \
+    } while(0)
+
+#define CHECK_ITER_STATUS(msg, status)                                                   \
+    do                                                                                   \
+    {                                                                                    \
+        if((status) != HSA_STATUS_INFO_BREAK)                                            \
+        {                                                                                \
+            const char* emsg = 0;                                                        \
+            hsa_status_string(status, &emsg);                                            \
+            char _buffer[HSA_MESSAGE_LENGTH];                                            \
+            snprintf(_buffer, HSA_MESSAGE_LENGTH - 1, "%s: %s", msg,                     \
+                     emsg ? emsg : "<unknown error>");                                   \
+            throw std::runtime_error(_buffer);                                           \
+        }                                                                                \
+    } while(0)
+
+static const size_t                   MEM_PAGE_BYTES = 0x1000;
+static const size_t                   MEM_PAGE_MASK  = MEM_PAGE_BYTES - 1;
+typedef decltype(hsa_agent_t::handle) hsa_agent_handle_t;
+
+struct hsa_pfn_t
+{
+    decltype(::hsa_init)*           hsa_init;
+    decltype(::hsa_shut_down)*      hsa_shut_down;
+    decltype(::hsa_agent_get_info)* hsa_agent_get_info;
+    decltype(::hsa_iterate_agents)* hsa_iterate_agents;
+
+    decltype(::hsa_queue_create)*                    hsa_queue_create;
+    decltype(::hsa_queue_destroy)*                   hsa_queue_destroy;
+    decltype(::hsa_queue_load_write_index_relaxed)*  hsa_queue_load_write_index_relaxed;
+    decltype(::hsa_queue_store_write_index_relaxed)* hsa_queue_store_write_index_relaxed;
+    decltype(::hsa_queue_load_read_index_relaxed)*   hsa_queue_load_read_index_relaxed;
+
+    decltype(::hsa_signal_create)*          hsa_signal_create;
+    decltype(::hsa_signal_destroy)*         hsa_signal_destroy;
+    decltype(::hsa_signal_load_relaxed)*    hsa_signal_load_relaxed;
+    decltype(::hsa_signal_store_relaxed)*   hsa_signal_store_relaxed;
+    decltype(::hsa_signal_wait_scacquire)*  hsa_signal_wait_scacquire;
+    decltype(::hsa_signal_store_screlease)* hsa_signal_store_screlease;
+
+    decltype(::hsa_code_object_reader_create_from_file)*
+                                           hsa_code_object_reader_create_from_file;
+    decltype(::hsa_executable_create_alt)* hsa_executable_create_alt;
+    decltype(
+        ::hsa_executable_load_agent_code_object)* hsa_executable_load_agent_code_object;
+    decltype(::hsa_executable_freeze)*            hsa_executable_freeze;
+    decltype(::hsa_executable_get_symbol)*        hsa_executable_get_symbol;
+    decltype(::hsa_executable_symbol_get_info)*   hsa_executable_symbol_get_info;
+    decltype(::hsa_executable_iterate_symbols)*   hsa_executable_iterate_symbols;
+
+    decltype(::hsa_system_get_info)* hsa_system_get_info;
+    decltype(
+        ::hsa_system_get_major_extension_table)* hsa_system_get_major_extension_table;
+
+    decltype(::hsa_amd_agent_iterate_memory_pools)* hsa_amd_agent_iterate_memory_pools;
+    decltype(::hsa_amd_memory_pool_get_info)*       hsa_amd_memory_pool_get_info;
+    decltype(::hsa_amd_memory_pool_allocate)*       hsa_amd_memory_pool_allocate;
+    decltype(::hsa_amd_agents_allow_access)*        hsa_amd_agents_allow_access;
+    decltype(::hsa_amd_memory_async_copy)*          hsa_amd_memory_async_copy;
+
+    decltype(::hsa_amd_signal_async_handler)* hsa_amd_signal_async_handler;
+    decltype(
+        ::hsa_amd_profiling_set_profiler_enabled)* hsa_amd_profiling_set_profiler_enabled;
+    decltype(
+        ::hsa_amd_profiling_get_async_copy_time)* hsa_amd_profiling_get_async_copy_time;
+    decltype(::hsa_amd_profiling_get_dispatch_time)* hsa_amd_profiling_get_dispatch_time;
+};
+
+// Encapsulates information about a Hsa Agent such as its
+// handle, name, max queue size, max wavefront size, etc.
+struct AgentInfo
+{
+    // Handle of Agent
+    hsa_agent_t dev_id;
+
+    // Agent type - Cpu = 0, Gpu = 1 or Dsp = 2
+    uint32_t dev_type;
+
+    // APU flag
+    bool is_apu;
+
+    // Agent system index
+    uint32_t dev_index;
+
+    // GFXIP name
+    char gfxip[64];
+
+    // Name of Agent whose length is less than 64
+    char name[64];
+
+    // Max size of Wavefront size
+    uint32_t max_wave_size;
+
+    // Max size of Queue buffer
+    uint32_t max_queue_size;
+
+    // Hsail profile supported by agent
+    hsa_profile_t profile;
+
+    // CPU/GPU/kern-arg memory pools
+    hsa_amd_memory_pool_t cpu_pool;
+    hsa_amd_memory_pool_t gpu_pool;
+    hsa_amd_memory_pool_t kern_arg_pool;
+
+    // The number of compute unit available in the agent.
+    uint32_t cu_num;
+
+    // Maximum number of waves possible in a Compute Unit.
+    uint32_t waves_per_cu;
+
+    // Number of SIMD's per compute unit CU
+    uint32_t simds_per_cu;
+
+    // Number of Shader Engines (SE) in Gpu
+    uint32_t se_num;
+
+    // Number of Shader Arrays Per Shader Engines in Gpu
+    uint32_t shader_arrays_per_se;
+
+    // SGPR/VGPR block sizes
+    uint32_t sgpr_block_dflt;
+    uint32_t sgpr_block_size;
+    uint32_t vgpr_block_size;
+};
+
+// HSA timer class
+// Provides current HSA timestampa and system-clock/ns conversion API
+class HsaTimer
+{
+public:
+    typedef uint64_t         timestamp_t;
+    static const timestamp_t TIMESTAMP_MAX = UINT64_MAX;
+    typedef long double      freq_t;
+
+    HsaTimer(const hsa_pfn_t* hsa_api)
+    : hsa_api_(hsa_api)
+    {
+        timestamp_t  sysclock_hz = 0;
+        hsa_status_t status      = hsa_api_->hsa_system_get_info(
+            HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz);
+        CHECK_STATUS("hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY)", status);
+        sysclock_factor_ = (freq_t) 1000000000 / (freq_t) sysclock_hz;
+    }
+
+    // Methods for system-clock/ns conversion
+    timestamp_t sysclock_to_ns(const timestamp_t& sysclock) const
+    {
+        return timestamp_t((freq_t) sysclock * sysclock_factor_);
+    }
+    timestamp_t ns_to_sysclock(const timestamp_t& time) const
+    {
+        return timestamp_t((freq_t) time / sysclock_factor_);
+    }
+
+    // Return timestamp in 'ns'
+    timestamp_t timestamp_ns() const
+    {
+        timestamp_t  sysclock;
+        hsa_status_t status =
+            hsa_api_->hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP, &sysclock);
+        CHECK_STATUS("hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP)", status);
+        return sysclock_to_ns(sysclock);
+    }
+
+private:
+    // Timestamp frequency factor
+    freq_t sysclock_factor_;
+    // HSA API table
+    const hsa_pfn_t* const hsa_api_;
+};
+
+class HsaRsrcFactory
+{
+public:
+    static const size_t           CMD_SLOT_SIZE_B = 0x40;
+    typedef std::recursive_mutex  mutex_t;
+    typedef HsaTimer::timestamp_t timestamp_t;
+
+    static HsaRsrcFactory* Create(bool initialize_hsa = true)
+    {
+        std::lock_guard<mutex_t> lck(mutex_);
+        HsaRsrcFactory*          obj = instance_.load(std::memory_order_relaxed);
+        if(obj == nullptr)
+        {
+            obj = new HsaRsrcFactory(initialize_hsa);
+            instance_.store(obj, std::memory_order_release);
+        }
+        return obj;
+    }
+
+    static HsaRsrcFactory& Instance()
+    {
+        HsaRsrcFactory* obj = instance_.load(std::memory_order_acquire);
+        if(obj == nullptr) obj = Create(false);
+        hsa_status_t status = (obj != nullptr) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
+        CHECK_STATUS("HsaRsrcFactory::Instance() failed", status);
+        return *obj;
+    }
+
+    static void Destroy()
+    {
+        std::lock_guard<mutex_t> lck(mutex_);
+        if(instance_) delete instance_.load();
+        instance_ = nullptr;
+    }
+
+    // Return system agent info
+    const AgentInfo* GetAgentInfo(const hsa_agent_t agent);
+
+    // Get the count of Hsa Gpu Agents available on the platform
+    // @return uint32_t Number of Gpu agents on platform
+    uint32_t GetCountOfGpuAgents();
+
+    // Get the count of Hsa Cpu Agents available on the platform
+    // @return uint32_t Number of Cpu agents on platform
+    uint32_t GetCountOfCpuAgents();
+
+    // Get the AgentInfo handle of a Gpu device
+    // @param idx Gpu Agent at specified index
+    // @param agent_info Output parameter updated with AgentInfo
+    // @return bool true if successful, false otherwise
+    bool GetGpuAgentInfo(uint32_t idx, const AgentInfo** agent_info);
+
+    // Get the AgentInfo handle of a Cpu device
+    // @param idx Cpu Agent at specified index
+    // @param agent_info Output parameter updated with AgentInfo
+    // @return bool true if successful, false otherwise
+    bool GetCpuAgentInfo(uint32_t idx, const AgentInfo** agent_info);
+
+    // Create a Queue object and return its handle. The queue object is expected
+    // to support user requested number of Aql dispatch packets.
+    // @param agent_info Gpu Agent on which to create a queue object
+    // @param num_Pkts Number of packets to be held by queue
+    // @param queue Output parameter updated with handle of queue object
+    // @return bool true if successful, false otherwise
+    bool CreateQueue(const AgentInfo* agent_info, uint32_t num_pkts, hsa_queue_t** queue);
+
+    // Create a Signal object and return its handle.
+    // @param value Initial value of signal object
+    // @param signal Output parameter updated with handle of signal object
+    // @return bool true if successful, false otherwise
+    bool CreateSignal(uint32_t value, hsa_signal_t* signal);
+
+    // Allocate local GPU memory
+    // @param agent_info Agent from whose memory region to allocate
+    // @param size Size of memory in terms of bytes
+    // @return uint8_t* Pointer to buffer, null if allocation fails.
+    uint8_t* AllocateLocalMemory(const AgentInfo* agent_info, size_t size);
+
+    // Allocate memory tp pass kernel parameters
+    // Memory is alocated accessible for all CPU agents and for GPU given by AgentInfo
+    // parameter.
+    // @param agent_info Agent from whose memory region to allocate
+    // @param size Size of memory in terms of bytes
+    // @return uint8_t* Pointer to buffer, null if allocation fails.
+    uint8_t* AllocateKernArgMemory(const AgentInfo* agent_info, size_t size);
+
+    // Allocate system memory accessible from both CPU and GPU
+    // Memory is alocated accessible to all CPU agents and AgentInfo parameter is ignored.
+    // @param agent_info Agent from whose memory region to allocate
+    // @param size Size of memory in terms of bytes
+    // @return uint8_t* Pointer to buffer, null if allocation fails.
+    uint8_t* AllocateSysMemory(const AgentInfo* agent_info, size_t size);
+
+    // Allocate memory for command buffer.
+    // @param agent_info Agent from whose memory region to allocate
+    // @param size Size of memory in terms of bytes
+    // @return uint8_t* Pointer to buffer, null if allocation fails.
+    uint8_t* AllocateCmdMemory(const AgentInfo* agent_info, size_t size);
+
+    // Wait signal
+    void SignalWait(const hsa_signal_t& signal) const;
+
+    // Wait signal with signal value restore
+    void SignalWaitRestore(const hsa_signal_t&       signal,
+                           const hsa_signal_value_t& signal_value) const;
+
+    // Copy data from GPU to host memory
+    bool Memcpy(const hsa_agent_t& agent, void* dst, const void* src, size_t size);
+    bool Memcpy(const AgentInfo* agent_info, void* dst, const void* src, size_t size);
+
+    // Memory free method
+    static bool FreeMemory(void* ptr);
+
+    // Loads an Assembled Brig file and Finalizes it into Device Isa
+    // @param agent_info Gpu device for which to finalize
+    // @param brig_path File path of the Assembled Brig file
+    // @param kernel_name Name of the kernel to finalize
+    // @param code_desc Handle of finalized Code Descriptor that could
+    // be used to submit for execution
+    // @return true if successful, false otherwise
+    bool LoadAndFinalize(const AgentInfo* agent_info, const char* brig_path,
+                         const char* kernel_name, hsa_executable_t* hsa_exec,
+                         hsa_executable_symbol_t* code_desc);
+
+    // Print the various fields of Hsa Gpu Agents
+    bool PrintGpuAgents(const std::string& header);
+
+    // Submit AQL packet to given queue
+    static uint64_t Submit(hsa_queue_t* queue, const void* packet);
+    static uint64_t Submit(hsa_queue_t* queue, const void* packet, size_t size_bytes);
+
+    // Enable executables loading tracking
+    static bool        IsExecutableTracking() { return executable_tracking_on_; }
+    static void        EnableExecutableTracking(HsaApiTable* table);
+    static const char* GetKernelName(uint64_t addr);
+
+    // Initialize HSA API table
+    void static InitHsaApiTable(HsaApiTable* table);
+    static const hsa_pfn_t* HsaApi() { return &hsa_api_; }
+
+    // Return AqlProfile API table
+    typedef hsa_ven_amd_aqlprofile_pfn_t aqlprofile_pfn_t;
+    const aqlprofile_pfn_t* AqlProfileApi() const { return &aqlprofile_api_; }
+
+    // Return Loader API table
+    const hsa_ven_amd_loader_1_00_pfn_t* LoaderApi() const { return &loader_api_; }
+
+    // Methods for system-clock/ns conversion and timestamp in 'ns'
+    timestamp_t SysclockToNs(const timestamp_t& sysclock) const
+    {
+        return timer_->sysclock_to_ns(sysclock);
+    }
+    timestamp_t NsToSysclock(const timestamp_t& time) const
+    {
+        return timer_->ns_to_sysclock(time);
+    }
+    timestamp_t TimestampNs() const { return timer_->timestamp_ns(); }
+
+    timestamp_t        GetSysTimeout() const { return timeout_; }
+    static timestamp_t GetTimeoutNs() { return timeout_ns_; }
+    static void        SetTimeoutNs(const timestamp_t& time)
+    {
+        std::lock_guard<mutex_t> lck(mutex_);
+        timeout_ns_ = time;
+        if(instance_ != nullptr)
+            Instance().timeout_ = Instance().timer_->ns_to_sysclock(time);
+    }
+
+private:
+    // System agents iterating callback
+    static hsa_status_t GetHsaAgentsCallback(hsa_agent_t agent, void* data);
+
+    // Callback function to find and bind kernarg region of an agent
+    static hsa_status_t FindMemRegionsCallback(hsa_region_t region, void* data);
+
+    // Load AQL profile HSA extension library directly
+    static hsa_status_t LoadAqlProfileLib(aqlprofile_pfn_t* api);
+
+    // Constructor of the class. Will initialize the Hsa Runtime and
+    // query the system topology to get the list of Cpu and Gpu devices
+    explicit HsaRsrcFactory(bool initialize_hsa);
+
+    // Destructor of the class
+    ~HsaRsrcFactory();
+
+    // Add an instance of AgentInfo representing a Hsa Gpu agent
+    const AgentInfo* AddAgentInfo(const hsa_agent_t agent);
+
+    // To mmap command buffer memory
+    static const bool CMD_MEMORY_MMAP = false;
+
+    // HSA was initialized
+    const bool initialize_hsa_;
+
+    static std::atomic<HsaRsrcFactory*> instance_;
+    static mutex_t                      mutex_;
+
+    // Used to maintain a list of Hsa Gpu Agent Info
+    std::vector<const AgentInfo*> gpu_list_;
+    std::vector<hsa_agent_t>      gpu_agents_;
+
+    // Used to maintain a list of Hsa Cpu Agent Info
+    std::vector<const AgentInfo*> cpu_list_;
+    std::vector<hsa_agent_t>      cpu_agents_;
+
+    // System agents map
+    std::map<hsa_agent_handle_t, const AgentInfo*> agent_map_;
+
+    // Executables loading tracking
+    typedef std::map<uint64_t, const char*> symbols_map_t;
+    static symbols_map_t*                   symbols_map_;
+    static bool                             executable_tracking_on_;
+    static hsa_status_t hsa_executable_freeze_interceptor(hsa_executable_t executable,
+                                                          const char*      options);
+    static hsa_status_t executable_symbols_cb(hsa_executable_t        exec,
+                                              hsa_executable_symbol_t symbol, void* data);
+
+    // HSA runtime API table
+    static hsa_pfn_t hsa_api_;
+
+    // AqlProfile API table
+    aqlprofile_pfn_t aqlprofile_api_;
+
+    // Loader API table
+    hsa_ven_amd_loader_1_00_pfn_t loader_api_;
+
+    // System timeout, ns
+    static timestamp_t timeout_ns_;
+    // System timeout, sysclock
+    timestamp_t timeout_;
+
+    // HSA timer
+    HsaTimer* timer_;
+
+    // CPU/kern-arg memory pools
+    hsa_amd_memory_pool_t* cpu_pool_;
+    hsa_amd_memory_pool_t* kern_arg_pool_;
+};
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer_callbacks.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/roctracer.cpp
similarity index 85%
rename from projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer_callbacks.cpp
rename to projects/rocprofiler-systems/source/lib/omnitrace/library/roctracer.cpp
index dc836b65b7..1a02ddeb2f 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer_callbacks.cpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/roctracer.cpp
@@ -20,7 +20,7 @@
 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 // SOFTWARE.
 
-#include "library/components/roctracer_callbacks.hpp"
+#include "library/roctracer.hpp"
 #include "library.hpp"
 #include "library/config.hpp"
 #include "library/critical_trace.hpp"
@@ -77,65 +77,6 @@ hip_api_string(hip_api_id_t id, const hip_api_data_t* data)
     tim::consume_parameters(id, data);
 #endif
 }
-//
-int64_t
-get_clock_skew()
-{
-    static auto _use = tim::get_env("OMNITRACE_USE_ROCTRACER_CLOCK_SKEW", true);
-    static auto _v   = []() {
-        namespace cpu = tim::cpu;
-        // synchronize timestamps
-        // We'll take a CPU timestamp before and after taking a GPU timestmp, then
-        // take the average of those two, hoping that it's roughly at the same time
-        // as the GPU timestamp.
-        static auto _cpu_now = []() {
-            cpu::fence();
-            return comp::wall_clock::record();
-        };
-
-        static auto _gpu_now = []() {
-            cpu::fence();
-            uint64_t _v = 0;
-            ROCTRACER_CALL(roctracer_get_timestamp(&_v));
-            return _v;
-        };
-
-        do
-        {
-            // warm up cache and allow for any static initialization
-            (void) _cpu_now();
-            (void) _gpu_now();
-        } while(false);
-
-        auto _compute = [](volatile uint64_t& _cpu_ts, volatile uint64_t& _gpu_ts) {
-            _cpu_ts = 0;
-            _gpu_ts = 0;
-            _cpu_ts += _cpu_now() / 2;
-            _gpu_ts += _gpu_now() / 1;
-            _cpu_ts += _cpu_now() / 2;
-            return static_cast<int64_t>(_cpu_ts) - static_cast<int64_t>(_gpu_ts);
-        };
-        constexpr int64_t _n       = 10;
-        int64_t           _cpu_ave = 0;
-        int64_t           _gpu_ave = 0;
-        int64_t           _diff    = 0;
-        for(int64_t i = 0; i < _n; ++i)
-        {
-            volatile uint64_t _cpu_ts = 0;
-            volatile uint64_t _gpu_ts = 0;
-            _diff += _compute(_cpu_ts, _gpu_ts);
-            _cpu_ave += _cpu_ts / _n;
-            _gpu_ave += _gpu_ts / _n;
-        }
-        OMNITRACE_BASIC_VERBOSE(2, "CPU timestamp: %li\n", _cpu_ave);
-        OMNITRACE_BASIC_VERBOSE(2, "HIP timestamp: %li\n", _gpu_ave);
-        OMNITRACE_BASIC_VERBOSE(1, "CPU/HIP timestamp skew: %li (used: %s)\n", _diff,
-                                _use ? "yes" : "no");
-        _diff /= _n;
-        return _diff;
-    }();
-    return (_use) ? _v : 0;
-}
 
 int&
 get_current_device()
@@ -218,6 +159,66 @@ get_hip_activity_mutex(int64_t _tid = threading::get_id())
 }
 }  // namespace
 
+//
+int64_t
+get_clock_skew()
+{
+    static auto _use = tim::get_env("OMNITRACE_USE_ROCTRACER_CLOCK_SKEW", true);
+    static auto _v   = []() {
+        namespace cpu = tim::cpu;
+        // synchronize timestamps
+        // We'll take a CPU timestamp before and after taking a GPU timestmp, then
+        // take the average of those two, hoping that it's roughly at the same time
+        // as the GPU timestamp.
+        static auto _cpu_now = []() {
+            cpu::fence();
+            return comp::wall_clock::record();
+        };
+
+        static auto _gpu_now = []() {
+            cpu::fence();
+            uint64_t _v = 0;
+            ROCTRACER_CALL(roctracer_get_timestamp(&_v));
+            return _v;
+        };
+
+        do
+        {
+            // warm up cache and allow for any static initialization
+            (void) _cpu_now();
+            (void) _gpu_now();
+        } while(false);
+
+        auto _compute = [](volatile uint64_t& _cpu_ts, volatile uint64_t& _gpu_ts) {
+            _cpu_ts = 0;
+            _gpu_ts = 0;
+            _cpu_ts += _cpu_now() / 2;
+            _gpu_ts += _gpu_now() / 1;
+            _cpu_ts += _cpu_now() / 2;
+            return static_cast<int64_t>(_cpu_ts) - static_cast<int64_t>(_gpu_ts);
+        };
+        constexpr int64_t _n       = 10;
+        int64_t           _cpu_ave = 0;
+        int64_t           _gpu_ave = 0;
+        int64_t           _diff    = 0;
+        for(int64_t i = 0; i < _n; ++i)
+        {
+            volatile uint64_t _cpu_ts = 0;
+            volatile uint64_t _gpu_ts = 0;
+            _diff += _compute(_cpu_ts, _gpu_ts);
+            _cpu_ave += _cpu_ts / _n;
+            _gpu_ave += _gpu_ts / _n;
+        }
+        OMNITRACE_BASIC_VERBOSE(2, "CPU timestamp: %li\n", _cpu_ave);
+        OMNITRACE_BASIC_VERBOSE(2, "HIP timestamp: %li\n", _gpu_ave);
+        OMNITRACE_BASIC_VERBOSE(1, "CPU/HIP timestamp skew: %li (used: %s)\n", _diff,
+                                _use ? "yes" : "no");
+        _diff /= _n;
+        return _diff;
+    }();
+    return (_use) ? _v : 0;
+}
+
 // HSA API callback function
 void
 hsa_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg)
@@ -227,6 +228,15 @@ hsa_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void*
 
     OMNITRACE_SCOPED_THREAD_STATE(ThreadState::Internal);
 
+    static thread_local std::once_flag _once{};
+    std::call_once(_once, []() {
+        if(threading::get_id() != 0)
+        {
+            sampling::block_signals();
+            threading::set_thread_name("roctracer.hsa");
+        }
+    });
+
     (void) arg;
     const hsa_api_data_t* data = reinterpret_cast<const hsa_api_data_t*>(callback_data);
     OMNITRACE_CONDITIONAL_PRINT_F(
@@ -343,7 +353,7 @@ hsa_activity_callback(uint32_t op, activity_record_t* record, void* arg)
     static thread_local std::once_flag _once{};
     std::call_once(_once, []() {
         sampling::block_signals();
-        threading::set_thread_name("omni.roctracer");
+        threading::set_thread_name("roctracer.hsa");
     });
 
     auto&& _protect = comp::roctracer::protect_flush_activity();
@@ -699,7 +709,7 @@ hip_activity_callback(const char* begin, const char* end, void*)
     static thread_local std::once_flag _once{};
     std::call_once(_once, []() {
         sampling::block_signals();
-        threading::set_thread_name("omni.roctracer");
+        threading::set_thread_name("roctracer.hip");
     });
 
     auto&& _protect = comp::roctracer::protect_flush_activity();
@@ -883,134 +893,3 @@ roctracer_shutdown_routines()
     return _v;
 }
 }  // namespace omnitrace
-
-#include "library/components/rocm_smi.hpp"
-
-using namespace omnitrace;
-
-// HSA-runtime tool on-load method
-extern "C"
-{
-    bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
-                const char* const* failed_tool_names) OMNITRACE_VISIBILITY("default");
-    void OnUnload() OMNITRACE_VISIBILITY("default");
-
-    bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
-                const char* const* failed_tool_names)
-    {
-        if(!tim::get_env("OMNITRACE_INIT_TOOLING", true)) return true;
-        if(!tim::settings::enabled()) return true;
-
-        roctracer_is_init() = true;
-        pthread_gotcha::push_enable_sampling_on_child_threads(false);
-        OMNITRACE_CONDITIONAL_BASIC_PRINT_F(get_debug_env() || get_verbose_env() > 0,
-                                            "\n");
-        tim::consume_parameters(table, runtime_version, failed_tool_count,
-                                failed_tool_names);
-
-        if(!config::settings_are_configured() && get_state() < State::Active)
-            omnitrace_init_tooling_hidden();
-
-        OMNITRACE_SCOPED_THREAD_STATE(ThreadState::Internal);
-
-        static auto _setup = [=]() {
-            try
-            {
-                OMNITRACE_CONDITIONAL_BASIC_PRINT_F(get_debug() || get_verbose() > 1,
-                                                    "setting up HSA...\n");
-
-                // const char* output_prefix = getenv("ROCP_OUTPUT_DIR");
-                const char* output_prefix = nullptr;
-
-                bool trace_hsa_api = get_trace_hsa_api();
-
-                // Enable HSA API callbacks/activity
-                if(trace_hsa_api)
-                {
-                    std::vector<std::string> hsa_api_vec =
-                        tim::delimit(get_trace_hsa_api_types());
-
-                    // initialize HSA tracing
-                    roctracer_set_properties(ACTIVITY_DOMAIN_HSA_API, (void*) table);
-
-                    OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug() || get_verbose() > 1,
-                                                      "    HSA-trace(");
-                    if(!hsa_api_vec.empty())
-                    {
-                        for(const auto& itr : hsa_api_vec)
-                        {
-                            uint32_t    cid = HSA_API_ID_NUMBER;
-                            const char* api = itr.c_str();
-                            ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_HSA_API, api,
-                                                             &cid, nullptr));
-                            ROCTRACER_CALL(roctracer_enable_op_callback(
-                                ACTIVITY_DOMAIN_HSA_API, cid, hsa_api_callback, nullptr));
-
-                            OMNITRACE_CONDITIONAL_BASIC_PRINT(
-                                get_debug() || get_verbose() > 1, " %s", api);
-                        }
-                    }
-                    else
-                    {
-                        ROCTRACER_CALL(roctracer_enable_domain_callback(
-                            ACTIVITY_DOMAIN_HSA_API, hsa_api_callback, nullptr));
-                    }
-                    OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug() || get_verbose() > 1,
-                                                      "\n");
-                }
-
-                bool trace_hsa_activity = get_trace_hsa_activity();
-                // Enable HSA GPU activity
-                if(trace_hsa_activity)
-                {
-                    // initialize HSA tracing
-                    ::roctracer::hsa_ops_properties_t ops_properties{
-                        table,
-                        reinterpret_cast<activity_async_callback_t>(
-                            hsa_activity_callback),
-                        nullptr, output_prefix
-                    };
-                    roctracer_set_properties(ACTIVITY_DOMAIN_HSA_OPS, &ops_properties);
-
-                    OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug() || get_verbose() > 1,
-                                                      "    HSA-activity-trace()\n");
-                    ROCTRACER_CALL(roctracer_enable_op_activity(ACTIVITY_DOMAIN_HSA_OPS,
-                                                                HSA_OP_ID_COPY));
-                }
-            } catch(std::exception& _e)
-            {
-                OMNITRACE_BASIC_PRINT("Exception was thrown in HSA setup: %s\n",
-                                      _e.what());
-            }
-        };
-
-        static auto _shutdown = []() {
-            OMNITRACE_DEBUG_F("roctracer_disable_domain_callback\n");
-            ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HSA_API));
-
-            OMNITRACE_DEBUG_F("roctracer_disable_op_activity\n");
-            ROCTRACER_CALL(
-                roctracer_disable_op_activity(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY));
-        };
-
-        (void) get_clock_skew();
-
-        comp::roctracer::add_setup("hsa", _setup);
-        comp::roctracer::add_shutdown("hsa", _shutdown);
-
-        rocm_smi::set_state(State::Active);
-        comp::roctracer::setup();
-
-        pthread_gotcha::pop_enable_sampling_on_child_threads();
-        return true;
-    }
-
-    // HSA-runtime on-unload method
-    void OnUnload()
-    {
-        OMNITRACE_DEBUG_F("\n");
-        rocm_smi::set_state(State::Finalized);
-        comp::roctracer::shutdown();
-        omnitrace_finalize_hidden();
-    }
-}
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer_callbacks.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/roctracer.hpp
similarity index 99%
rename from projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer_callbacks.hpp
rename to projects/rocprofiler-systems/source/lib/omnitrace/library/roctracer.hpp
index bc598bdb1c..044f0fb40c 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer_callbacks.hpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/roctracer.hpp
@@ -77,6 +77,9 @@ roctracer_is_init();
 bool&
 roctracer_is_setup();
 
+int64_t
+get_clock_skew();
+
 roctracer_functions_t&
 roctracer_setup_routines();
 
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/runtime.cpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/runtime.cpp
index 6fb74019d5..9c06b3327d 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/runtime.cpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/runtime.cpp
@@ -157,7 +157,8 @@ std::unique_ptr<main_bundle_t>&
 get_main_bundle()
 {
     static auto _v =
-        std::make_unique<main_bundle_t>("omnitrace", quirk::config<quirk::auto_start>{});
+        std::make_unique<main_bundle_t>(JOIN('/', "omnitrace/process", process::get_id()),
+                                        quirk::config<quirk::auto_start>{});
     return _v;
 }
 
@@ -166,7 +167,8 @@ get_gotcha_bundle()
 {
     static auto _v =
         (setup_gotchas(), std::make_unique<gotcha_bundle_t>(
-                              "omnitrace", quirk::config<quirk::auto_start>{}));
+                              JOIN('/', "omnitrace/process", process::get_id()),
+                              quirk::config<quirk::auto_start>{}));
     return _v;
 }
 
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/timemory.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/timemory.hpp
index e078cab872..b6d4028146 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/timemory.hpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/timemory.hpp
@@ -44,15 +44,17 @@
 
 namespace omnitrace
 {
-namespace audit     = tim::audit;      // NOLINT
-namespace comp      = tim::component;  // NOLINT
-namespace quirk     = tim::quirk;      // NOLINT
-namespace threading = tim::threading;  // NOLINT
-namespace scope     = tim::scope;      // NOLINT
-namespace dmp       = tim::dmp;        // NOLINT
-namespace process   = tim::process;    // NOLINT
-namespace units     = tim::units;      // NOLINT
-namespace trait     = tim::trait;      // NOLINT
-namespace api       = tim::api;        // NOLINT
-namespace operation = tim::operation;  // NOLINT
+namespace audit     = ::tim::audit;      // NOLINT
+namespace comp      = ::tim::component;  // NOLINT
+namespace quirk     = ::tim::quirk;      // NOLINT
+namespace threading = ::tim::threading;  // NOLINT
+namespace scope     = ::tim::scope;      // NOLINT
+namespace dmp       = ::tim::dmp;        // NOLINT
+namespace process   = ::tim::process;    // NOLINT
+namespace units     = ::tim::units;      // NOLINT
+namespace trait     = ::tim::trait;      // NOLINT
+namespace api       = ::tim::api;        // NOLINT
+namespace operation = ::tim::operation;  // NOLINT
+
+using settings = ::tim::settings;  // NOLINT
 }  // namespace omnitrace
diff --git a/projects/rocprofiler-systems/source/lib/omnitrace/library/tracing.hpp b/projects/rocprofiler-systems/source/lib/omnitrace/library/tracing.hpp
index 3a7d59ee09..c5ccf5b7ea 100644
--- a/projects/rocprofiler-systems/source/lib/omnitrace/library/tracing.hpp
+++ b/projects/rocprofiler-systems/source/lib/omnitrace/library/tracing.hpp
@@ -91,11 +91,10 @@ inline void
 thread_init()
 {
     static thread_local auto _thread_setup = []() {
-        auto _exe = get_exe_name();
         if(threading::get_id() > 0)
             threading::set_thread_name(JOIN(" ", "Thread", threading::get_id()).c_str());
         thread_data<omnitrace_thread_bundle_t>::construct(
-            JOIN("", _exe, "/thread-", threading::get_id()),
+            JOIN('/', "omnitrace", process::get_id(), "thread", threading::get_id()),
             quirk::config<quirk::auto_start>{});
         get_interval_data()->reserve(512);
         // save the hash maps