Critical trace updates (#24)

* Source code restructuring * Critical trace updates following restructuring * thread_sampler, timestamps - thread_sampler - CPU frequency managed via thread_sampler - rocm-smi managed via thread_sampler - Use consistent timestamps for perfetto - removed hsa_timer_t in favor of wall_clock::record() - disable KokkosP by default - re-enable critical-trace testing * cmake-format * Fix for defines.hpp.in * Remove OMNITRACE_ROCM_SMI_FREQ - thread_sampler freq is set via OMNITRACE_SAMPLING_FREQ w/ max of 1000 * Increase CI Install Dyninst timeout * Debug macros + omnitrace_init_tooling + config - new debug macros - extern "C" omnitrace_init_tooling - guard get_rocm_smi_devices * Miscellaneous tweaks - tweak to transpose - critical_trace::Device::ANY - perfetto "critical-trace" category - OMNITRACE_VERBOSE usage * Disable key and tid data for HIP API calls - non-kernels are ignored in activity callback * critical-trace exe updates - fix perfetto generation - improved logging - improved readability * timemory submodule update - lulesh example cmake tweaks
2022-02-19 02:00:59 -06:00
parent 39f17ae8b8
commit b016c8929f
70 changed files with 2369 additions and 1434 deletions
@@ -216,7 +216,7 @@ jobs:
        rm -rf elfutils*

    - name: Install Dyninst
-      timeout-minutes: 20
+      timeout-minutes: 25
      run:
        cmake --version &&
        git submodule update --init external/dyninst &&
@@ -310,7 +310,7 @@ jobs:
        echo "LD_LIBRARY_PATH=/opt/omnitrace/lib:/opt/dyninst/lib:${LD_LIBRARY_PATH}" >> $GITHUB_ENV

    - name: Install Dyninst
-      timeout-minutes: 20
+      timeout-minutes: 25
      run:
        cmake --version &&
        git submodule update --init external/dyninst &&
@@ -429,7 +429,7 @@ jobs:
        rm -rf elfutils*

    - name: Install Dyninst
-      timeout-minutes: 20
+      timeout-minutes: 25
      run:
        cmake --version &&
        git submodule update --init external/dyninst &&
@@ -130,186 +130,22 @@ else()
    set(HIP_VERSION_PATCH 0)
 endif()

-configure_file(${PROJECT_SOURCE_DIR}/include/library/defines.hpp.in
-               ${PROJECT_BINARY_DIR}/include/library/defines.hpp @ONLY)
-
 omnitrace_activate_clang_tidy()

 # custom visibility settings
 if(OMNITRACE_BUILD_HIDDEN_VISIBILITY)
-    set(CMAKE_C_VISIBILITY_PRESET "hidden")
-    set(CMAKE_CXX_VISIBILITY_PRESET "hidden")
+    set(CMAKE_C_VISIBILITY_PRESET "internal")
+    set(CMAKE_CXX_VISIBILITY_PRESET "internal")
    set(CMAKE_VISIBILITY_INLINES_HIDDEN ON)
 endif()

-if(NOT DEFINED CMAKE_INTERPROCEDURAL_OPTIMIZATION)
-    set(CMAKE_INTERPROCEDURAL_OPTIMIZATION OFF)
-endif()
-
-if(OMNITRACE_BUILD_LTO)
-    omnitrace_save_variables(LTO VARIABLES CMAKE_INTERPROCEDURAL_OPTIMIZATION)
-    set(CMAKE_INTERPROCEDURAL_OPTIMIZATION ON)
-endif()
-
 # ------------------------------------------------------------------------------#
 #
-# omnitrace-library target
+# library and executables
 #
 # ------------------------------------------------------------------------------#

-set(library_sources
-    ${CMAKE_CURRENT_LIST_DIR}/src/library.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/config.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/critical_trace.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/kokkosp.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/gpu.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/perfetto.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/ptl.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/sampling.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/state.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/thread_data.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/timemory.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/components/backtrace.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/components/fork_gotcha.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/components/mpi_gotcha.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/components/omnitrace.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/library/components/pthread_gotcha.cpp
-    ${perfetto_DIR}/sdk/perfetto.cc)
-
-set(library_headers
-    ${CMAKE_CURRENT_LIST_DIR}/include/library.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/api.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/config.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/common.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/critical_trace.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/debug.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/gpu.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/perfetto.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/ptl.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/sampling.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/state.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/thread_data.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/timemory.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/fwd.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/backtrace.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/fork_gotcha.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/mpi_gotcha.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/omnitrace.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/rocm_smi.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/roctracer.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/roctracer_callbacks.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/pthread_gotcha.hpp
-    ${perfetto_DIR}/sdk/perfetto.h)
-
-add_library(omnitrace-library SHARED ${library_sources} ${library_headers})
-
-if(OMNITRACE_USE_ROCTRACER)
-    target_sources(
-        omnitrace-library
-        PRIVATE ${CMAKE_CURRENT_LIST_DIR}/src/library/components/roctracer.cpp
-                ${CMAKE_CURRENT_LIST_DIR}/src/library/components/roctracer_callbacks.cpp)
-endif()
-
-if(OMNITRACE_USE_ROCM_SMI)
-    target_sources(omnitrace-library
-                   PRIVATE ${CMAKE_CURRENT_LIST_DIR}/src/library/components/rocm_smi.cpp)
-endif()
-
-target_include_directories(omnitrace-library SYSTEM PRIVATE ${perfetto_DIR}/sdk)
-
-target_compile_definitions(
-    omnitrace-library
-    PRIVATE OMNITRACE_MAX_THREADS=${OMNITRACE_MAX_THREADS}
-            $<IF:$<BOOL:${OMNITRACE_CUSTOM_DATA_SOURCE}>,CUSTOM_DATA_SOURCE,>)
-
-target_link_libraries(
-    omnitrace-library
-    PUBLIC $<BUILD_INTERFACE:omnitrace::omnitrace-headers>
-           $<BUILD_INTERFACE:omnitrace::omnitrace-threading>
-           $<BUILD_INTERFACE:omnitrace::omnitrace-compile-options>
-           $<BUILD_INTERFACE:omnitrace::omnitrace-hip>
-           $<BUILD_INTERFACE:omnitrace::omnitrace-roctracer>
-           $<BUILD_INTERFACE:omnitrace::omnitrace-rocm-smi>
-           $<BUILD_INTERFACE:omnitrace::omnitrace-mpi>
-           $<BUILD_INTERFACE:omnitrace::omnitrace-ptl>
-           $<BUILD_INTERFACE:timemory::timemory-headers>
-           $<BUILD_INTERFACE:timemory::timemory-gotcha>
-           $<BUILD_INTERFACE:timemory::timemory-cxx-shared>
-           $<IF:$<BOOL:${OMNITRACE_USE_SANITIZER}>,omnitrace::omnitrace-sanitizer,>)
-
-if(OMNITRACE_DYNINST_API_RT)
-    get_filename_component(OMNITRACE_DYNINST_API_RT_DIR "${OMNITRACE_DYNINST_API_RT}"
-                           DIRECTORY)
-endif()
-
-set_target_properties(
-    omnitrace-library
-    PROPERTIES OUTPUT_NAME omnitrace
-               INSTALL_RPATH
-               "\$ORIGIN:\$ORIGIN/timemory/libunwind:\$ORIGIN/dyninst-tpls/libs")
-
-install(
-    TARGETS omnitrace-library
-    DESTINATION ${CMAKE_INSTALL_LIBDIR}
-    OPTIONAL)
-
-# ------------------------------------------------------------------------------#
-#
-# omnitrace-avail target
-#
-# ------------------------------------------------------------------------------#
-
-add_executable(omnitrace-avail ${CMAKE_CURRENT_LIST_DIR}/src/avail.cpp
-                               ${CMAKE_CURRENT_LIST_DIR}/include/avail.hpp)
-
-target_include_directories(omnitrace-avail PRIVATE ${CMAKE_CURRENT_LIST_DIR}/include)
-target_compile_definitions(omnitrace-avail PRIVATE OMNITRACE_EXTERN_COMPONENTS=0)
-target_link_libraries(omnitrace-avail PRIVATE omnitrace-library)
-set_target_properties(omnitrace-avail PROPERTIES INSTALL_RPATH_USE_LINK_PATH ON)
-
-install(
-    TARGETS omnitrace-avail
-    DESTINATION bin
-    OPTIONAL)
-
-# ------------------------------------------------------------------------------#
-#
-# omnitrace-exe target
-#
-# ------------------------------------------------------------------------------#
-
-add_executable(
-    omnitrace-exe
-    ${_EXCLUDE} ${CMAKE_CURRENT_LIST_DIR}/src/omnitrace.cpp
-    ${CMAKE_CURRENT_LIST_DIR}/include/omnitrace.hpp
-    ${CMAKE_CURRENT_LIST_DIR}/src/omnitrace/details.cpp)
-
-target_link_libraries(
-    omnitrace-exe
-    PRIVATE omnitrace::omnitrace-headers
-            omnitrace::omnitrace-dyninst
-            omnitrace::omnitrace-compile-options
-            $<BUILD_INTERFACE:timemory::timemory-headers>
-            $<IF:$<BOOL:${OMNITRACE_USE_SANITIZER}>,omnitrace::omnitrace-sanitizer,>)
-
-set_target_properties(
-    omnitrace-exe
-    PROPERTIES
-        OUTPUT_NAME omnitrace
-        INSTALL_RPATH_USE_LINK_PATH ON
-        INSTALL_RPATH
-        "\$ORIGIN/../${CMAKE_INSTALL_LIBDIR}:\$ORIGIN/../${CMAKE_INSTALL_LIBDIR}/timemory/libunwind:\$ORIGIN/../${CMAKE_INSTALL_LIBDIR}/dyninst-tpls/lib"
-    )
-
-if(CMAKE_BUILD_TYPE MATCHES "^(DEBUG|Debug)")
-    string(REPLACE " " ";" _FLAGS "${CMAKE_CXX_FLAGS_RELWITHDEBINFO}")
-    target_compile_options(omnitrace-exe PRIVATE ${_FLAGS})
-endif()
-
-install(
-    TARGETS omnitrace-exe
-    DESTINATION ${CMAKE_INSTALL_BINDIR}
-    OPTIONAL)
+add_subdirectory(source)

 # ------------------------------------------------------------------------------#
 #
@@ -41,9 +41,9 @@ find_program(OMNITRACE_CLANG_FORMAT_EXE NAMES clang-format-11 clang-format-mp-11
                                              clang-format)

 if(OMNITRACE_CLANG_FORMAT_EXE)
-    file(GLOB_RECURSE sources ${PROJECT_SOURCE_DIR}/src/*.cpp)
-    file(GLOB_RECURSE headers ${PROJECT_SOURCE_DIR}/include/*.hpp
-         ${PROJECT_SOURCE_DIR}/include/*.hpp.in)
+    file(GLOB_RECURSE sources ${PROJECT_SOURCE_DIR}/source/*.cpp)
+    file(GLOB_RECURSE headers ${PROJECT_SOURCE_DIR}/source/*.hpp
+         ${PROJECT_SOURCE_DIR}/source/*.hpp.in)
    file(GLOB_RECURSE examples ${PROJECT_SOURCE_DIR}/examples/*.cpp
         ${PROJECT_SOURCE_DIR}/examples/*.hpp)
    file(GLOB_RECURSE external ${PROJECT_SOURCE_DIR}/examples/lulesh/external/*.cpp
@@ -8,9 +8,8 @@ include(CMakeDependentOption)
 include(CMakeParseArguments)

 # -----------------------------------------------------------------------
-# function - capitalize - make a string capitalized (first letter is capital) usage:
-# capitalize("SHARED" CShared) message(STATUS "-- CShared is \"${CShared}\"") $ -- CShared
-# is "Shared"
+# function - capitalize - make a string capitalized (first letter is capital)
+#
 function(CAPITALIZE str var)
    # make string lower
    string(TOLOWER "${str}" str)
@@ -24,7 +23,7 @@ function(CAPITALIZE str var)
 endfunction()

 # ----------------------------------------------------------------------------------------#
-# macro CHECKOUT_GIT_SUBMODULE()
+# function CHECKOUT_GIT_SUBMODULE()
 #
 # Run "git submodule update" if a file in a submodule does not exist
 #
@@ -45,7 +44,7 @@ function(CHECKOUT_GIT_SUBMODULE)
    endif()

    if(NOT CHECKOUT_TEST_FILE)
-        set(CHECKOUT_TEST_FILE "Makefile")
+        set(CHECKOUT_TEST_FILE "CMakeLists.txt")
    endif()

    # default assumption
@@ -156,160 +155,3 @@ function(CHECKOUT_GIT_SUBMODULE)
    endif()

 endfunction()
-
-# ----------------------------------------------------------------------------------------#
-# require variable
-#
-function(CHECK_REQUIRED VAR)
-    if(NOT DEFINED ${VAR} OR "${${VAR}}" STREQUAL "")
-        message(FATAL_ERROR "Variable '${VAR}' must be defined and not empty")
-    endif()
-endfunction()
-
-# -----------------------------------------------------------------------
-# function add_feature(<NAME> <DOCSTRING>) Add a project feature, whose activation is
-# specified by the existence of the variable <NAME>, to the list of enabled/disabled
-# features, plus a docstring describing the feature
-#
-function(ADD_FEATURE _var _description)
-    set(EXTRA_DESC "")
-    foreach(currentArg ${ARGN})
-        if(NOT "${currentArg}" STREQUAL "${_var}" AND NOT "${currentArg}" STREQUAL
-                                                      "${_description}")
-            set(EXTRA_DESC "${EXTA_DESC}${currentArg}")
-        endif()
-    endforeach()
-
-    set_property(GLOBAL APPEND PROPERTY ${PROJECT_NAME}_FEATURES ${_var})
-    set_property(GLOBAL PROPERTY ${_var}_DESCRIPTION "${_description}${EXTRA_DESC}")
-
-    if("CMAKE_DEFINE" IN_LIST ARGN)
-        set_property(GLOBAL APPEND PROPERTY ${PROJECT_NAME}_CMAKE_DEFINES
-                                            "${_var} @${_var}@")
-    endif()
-endfunction()
-
-# ----------------------------------------------------------------------------------------#
-# function add_option(<OPTION_NAME> <DOCSRING> <DEFAULT_SETTING> [NO_FEATURE]) Add an
-# option and add as a feature if NO_FEATURE is not provided
-#
-function(ADD_OPTION _NAME _MESSAGE _DEFAULT)
-    option(${_NAME} "${_MESSAGE}" ${_DEFAULT})
-    if("NO_FEATURE" IN_LIST ARGN)
-        mark_as_advanced(${_NAME})
-    else()
-        add_feature(${_NAME} "${_MESSAGE}")
-    endif()
-    if("ADVANCED" IN_LIST ARGN)
-        mark_as_advanced(${_NAME})
-    endif()
-endfunction()
-
-# ----------------------------------------------------------------------------------------#
-# function print_enabled_features() Print enabled  features plus their docstrings.
-#
-function(PRINT_ENABLED_FEATURES)
-    set(_basemsg "The following features are defined/enabled (+):")
-    set(_currentFeatureText "${_basemsg}")
-    get_property(_features GLOBAL PROPERTY ${PROJECT_NAME}_FEATURES)
-    if(NOT "${_features}" STREQUAL "")
-        list(REMOVE_DUPLICATES _features)
-        list(SORT _features)
-    endif()
-    foreach(_feature ${_features})
-        if(${_feature})
-            # add feature to text
-            set(_currentFeatureText "${_currentFeatureText}\n     ${_feature}")
-            # get description
-            get_property(_desc GLOBAL PROPERTY ${_feature}_DESCRIPTION)
-            # print description, if not standard ON/OFF, print what is set to
-            if(_desc)
-                if(NOT "${${_feature}}" STREQUAL "ON" AND NOT "${${_feature}}" STREQUAL
-                                                          "TRUE")
-                    set(_currentFeatureText
-                        "${_currentFeatureText}: ${_desc} -- [\"${${_feature}}\"]")
-                else()
-                    string(REGEX REPLACE "^${PROJECT_NAME}_USE_" "" _feature_tmp
-                                         "${_feature}")
-                    string(TOLOWER "${_feature_tmp}" _feature_tmp_l)
-                    capitalize("${_feature_tmp}" _feature_tmp_c)
-                    foreach(_var _feature _feature_tmp _feature_tmp_l _feature_tmp_c)
-                        set(_ver "${${${_var}}_VERSION}")
-                        if(NOT "${_ver}" STREQUAL "")
-                            set(_desc "${_desc} -- [found version ${_ver}]")
-                            break()
-                        endif()
-                        unset(_ver)
-                    endforeach()
-                    set(_currentFeatureText "${_currentFeatureText}: ${_desc}")
-                endif()
-                set(_desc NOTFOUND)
-            endif()
-        endif()
-    endforeach()
-
-    if(NOT "${_currentFeatureText}" STREQUAL "${_basemsg}")
-        message(STATUS "${_currentFeatureText}\n")
-    endif()
-endfunction()
-
-# ----------------------------------------------------------------------------------------#
-# function print_disabled_features() Print disabled features plus their docstrings.
-#
-function(PRINT_DISABLED_FEATURES)
-    set(_basemsg "The following features are NOT defined/enabled (-):")
-    set(_currentFeatureText "${_basemsg}")
-    get_property(_features GLOBAL PROPERTY ${PROJECT_NAME}_FEATURES)
-    if(NOT "${_features}" STREQUAL "")
-        list(REMOVE_DUPLICATES _features)
-        list(SORT _features)
-    endif()
-    foreach(_feature ${_features})
-        if(NOT ${_feature})
-            set(_currentFeatureText "${_currentFeatureText}\n     ${_feature}")
-            get_property(_desc GLOBAL PROPERTY ${_feature}_DESCRIPTION)
-            if(_desc)
-                set(_currentFeatureText "${_currentFeatureText}: ${_desc}")
-                set(_desc NOTFOUND)
-            endif(_desc)
-        endif()
-    endforeach(_feature)
-
-    if(NOT "${_currentFeatureText}" STREQUAL "${_basemsg}")
-        message(STATUS "${_currentFeatureText}\n")
-    endif()
-endfunction()
-
-# ----------------------------------------------------------------------------------------#
-# function print_features() Print all features plus their docstrings.
-#
-function(PRINT_FEATURES)
-    message(STATUS "")
-    print_enabled_features()
-    print_disabled_features()
-endfunction()
-
-# ----------------------------------------------------------------------------------------#
-# macro ADD_SUBPROJECT() Does a git submodule update + add_subdirectory
-#
-macro(ADD_SUBPROJECT PACKAGE_NAME)
-    # parse args
-    cmake_parse_arguments(PACKAGE "SUBMODULE" "DIRECTORY" "" ${ARGN})
-    if(NOT PACKAGE_DIRECTORY)
-        set(PACKAGE_DIRECTORY ${PACKAGE_NAME})
-    endif()
-    # if specified in options
-    if("${PACKAGE_NAME}" IN_LIST PROJECTS)
-        if(PACKAGE_SUBMODULE)
-            checkout_git_submodule(RECURSIVE RELATIVE_PATH ${PACKAGE_DIRECTORY})
-        endif()
-        if(NOT EXISTS "${PROJECT_SOURCE_DIR}/${PACKAGE_DIRECTORY}/CMakeLists.txt")
-            message(
-                STATUS
-                    "Warning! '${PROJECT_SOURCE_DIR}/${PACKAGE_DIRECTORY}/CMakeLists.txt' does not exist!"
-                )
-        else()
-            add_subdirectory(${PACKAGE_DIRECTORY})
-        endif()
-    endif()
-endmacro()
@@ -1,4 +1,4 @@
-# include(Utilities)
+include(Utilities)

 option(LULESH_USE_CUDA "Enable Kokkos CUDA backend for lulesh" OFF)
 option(LULESH_USE_HIP "Enable Kokkos HIP backend for lulesh" OFF)
@@ -120,9 +120,6 @@ run(int rank, int tid, hipStream_t stream, int argc, char** argv)
    HIP_API_CALL(hipMemcpy(in, matrix, size, hipMemcpyHostToDevice));
    HIP_API_CALL(hipDeviceSynchronize());

-    hipDeviceProp_t props;
-    HIP_API_CALL(hipGetDeviceProperties(&props, 0));
-
    dim3 grid(M / 32, N / 32, 1);
    dim3 block(32, 32, 1);  // transpose_a

@@ -189,6 +186,9 @@ main(int argc, char** argv)
    if(argc > 1) nthreads = atoi(argv[1]);
    if(argc > 2) nitr = atoi(argv[2]);

+    printf("[transpose] Number of threads: %i\n", nthreads);
+    printf("[transpose] Number of iterations: %i\n", nitr);
+
 #if defined(USE_MPI)
    MPI_Init(&argc, &argv);
    MPI_Comm_rank(MPI_COMM_WORLD, &rank);
@@ -200,10 +200,12 @@ main(int argc, char** argv)
    int ndevice = 0;
    int devid   = rank;
    HIP_API_CALL(hipGetDeviceCount(&ndevice));
+    printf("[transpose] Number of devices found: %i\n", ndevice);
    if(ndevice > 0)
    {
        devid = rank % ndevice;
        HIP_API_CALL(hipSetDevice(devid));
+        printf("[transpose] Rank %i assigned to device %i\n", rank, devid);
    }
    if(rank == devid && rank < ndevice)
    {
@@ -1,140 +0,0 @@
-// 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"
-
-#include <timemory/api.hpp>
-#include <timemory/backends/dmp.hpp>
-#include <timemory/backends/process.hpp>
-#include <timemory/utility/utility.hpp>
-
-#include <cstdio>
-
-namespace omnitrace
-{
-inline namespace config
-{
-bool
-get_debug();
-
-bool
-get_debug_tid();
-
-bool
-get_debug_pid();
-
-bool
-get_critical_trace_debug();
-}  // namespace config
-}  // namespace omnitrace
-
-#if defined(TIMEMORY_USE_MPI)
-#    define OMNITRACE_CONDITIONAL_PRINT(COND, ...)                                       \
-        if((COND) && get_debug_tid() && get_debug_pid())                                 \
-        {                                                                                \
-            fflush(stderr);                                                              \
-            tim::auto_lock_t _lk{ tim::type_mutex<decltype(std::cerr)>() };              \
-            fprintf(stderr, "[omnitrace][%i][%li] ", static_cast<int>(tim::dmp::rank()), \
-                    tim::threading::get_id());                                           \
-            fprintf(stderr, __VA_ARGS__);                                                \
-            fflush(stderr);                                                              \
-        }
-#else
-#    define OMNITRACE_CONDITIONAL_PRINT(COND, ...)                                       \
-        if((COND) && get_debug_tid() && get_debug_pid())                                 \
-        {                                                                                \
-            fflush(stderr);                                                              \
-            tim::auto_lock_t _lk{ tim::type_mutex<decltype(std::cerr)>() };              \
-            fprintf(stderr, "[omnitrace][%i][%li] ",                                     \
-                    static_cast<int>(tim::process::get_id()), tim::threading::get_id()); \
-            fprintf(stderr, __VA_ARGS__);                                                \
-            fflush(stderr);                                                              \
-        }
-#endif
-
-#define OMNITRACE_CONDITIONAL_BASIC_PRINT(COND, ...)                                     \
-    if((COND) && get_debug_tid() && get_debug_pid())                                     \
-    {                                                                                    \
-        fflush(stderr);                                                                  \
-        tim::auto_lock_t _lk{ tim::type_mutex<decltype(std::cerr)>() };                  \
-        fprintf(stderr, "[omnitrace] ");                                                 \
-        fprintf(stderr, __VA_ARGS__);                                                    \
-        fflush(stderr);                                                                  \
-    }
-
-#if defined(TIMEMORY_USE_MPI)
-#    define OMNITRACE_CONDITIONAL_THROW(COND, ...)                                       \
-        if(COND)                                                                         \
-        {                                                                                \
-            char _msg_buffer[2048];                                                      \
-            snprintf(_msg_buffer, 2048, "[omnitrace][%i][%li] ",                         \
-                     static_cast<int>(tim::dmp::rank()), tim::threading::get_id());      \
-            auto len = strlen(_msg_buffer);                                              \
-            snprintf(_msg_buffer + len, 2048 - len, __VA_ARGS__);                        \
-            throw std::runtime_error(_msg_buffer);                                       \
-        }
-#else
-#    define OMNITRACE_CONDITIONAL_THROW(COND, ...)                                       \
-        if(COND)                                                                         \
-        {                                                                                \
-            char _msg_buffer[2048];                                                      \
-            snprintf(_msg_buffer, 2048, "[omnitrace][%i][%li] ",                         \
-                     static_cast<int>(tim::process::get_id()),                           \
-                     tim::threading::get_id());                                          \
-            auto len = strlen(_msg_buffer);                                              \
-            snprintf(_msg_buffer + len, 2048 - len, __VA_ARGS__);                        \
-            throw std::runtime_error(_msg_buffer);                                       \
-        }
-#endif
-
-#define OMNITRACE_CONDITIONAL_BASIC_THROW(COND, ...)                                     \
-    if(COND)                                                                             \
-    {                                                                                    \
-        char _msg_buffer[2048];                                                          \
-        snprintf(_msg_buffer, 2048, "[omnitrace] " __VA_ARGS__);                         \
-        auto len = strlen(_msg_buffer);                                                  \
-        snprintf(_msg_buffer + len, 2048 - len, __VA_ARGS__);                            \
-        throw std::runtime_error(_msg_buffer);                                           \
-    }
-
-#define OMNITRACE_DEBUG(...)                                                             \
-    OMNITRACE_CONDITIONAL_PRINT(::omnitrace::get_debug(), __VA_ARGS__)
-#define OMNITRACE_PRINT(...) OMNITRACE_CONDITIONAL_PRINT(true, __VA_ARGS__)
-#define OMNITRACE_CT_DEBUG(...)                                                          \
-    OMNITRACE_CONDITIONAL_PRINT(::omnitrace::get_critical_trace_debug(), __VA_ARGS__)
-#define OMNITRACE_BASIC_PRINT(...) OMNITRACE_CONDITIONAL_BASIC_PRINT(true, __VA_ARGS__)
-
-#define OMNITRACE_THROW(...)       OMNITRACE_CONDITIONAL_THROW(true, __VA_ARGS__)
-#define OMNITRACE_BASIC_THROW(...) OMNITRACE_CONDITIONAL_BASIC_THROW(true, __VA_ARGS__)
-
-#include <string>
-
-namespace std
-{
-inline std::string
-to_string(bool _v)
-{
-    return (_v) ? "true" : "false";
-}
-}  // namespace std
@@ -0,0 +1,14 @@
+if(NOT DEFINED CMAKE_INTERPROCEDURAL_OPTIMIZATION)
+    set(CMAKE_INTERPROCEDURAL_OPTIMIZATION OFF)
+endif()
+
+if(OMNITRACE_BUILD_LTO)
+    set(CMAKE_INTERPROCEDURAL_OPTIMIZATION ON)
+endif()
+
+set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR})
+set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR})
+set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR})
+
+add_subdirectory(lib)
+add_subdirectory(bin)
@@ -0,0 +1,3 @@
+add_subdirectory(omnitrace-avail)
+add_subdirectory(omnitrace-critical-trace)
+add_subdirectory(omnitrace)
@@ -0,0 +1,20 @@
+# ------------------------------------------------------------------------------#
+#
+# omnitrace-avail target
+#
+# ------------------------------------------------------------------------------#
+
+add_executable(
+    omnitrace-avail
+    ${CMAKE_CURRENT_LIST_DIR}/avail.cpp ${CMAKE_CURRENT_LIST_DIR}/avail.hpp
+    $<TARGET_OBJECTS:omnitrace::omnitrace-object-library>)
+
+target_include_directories(omnitrace-avail PRIVATE ${CMAKE_CURRENT_LIST_DIR}/include)
+target_compile_definitions(omnitrace-avail PRIVATE OMNITRACE_EXTERN_COMPONENTS=0)
+target_link_libraries(omnitrace-avail PRIVATE omnitrace::omnitrace-interface-library)
+set_target_properties(omnitrace-avail PROPERTIES INSTALL_RPATH_USE_LINK_PATH ON)
+
+install(
+    TARGETS omnitrace-avail
+    DESTINATION bin
+    OPTIONAL)
@@ -0,0 +1,23 @@
+# ------------------------------------------------------------------------------#
+#
+# omnitrace-critical-trace target
+#
+# ------------------------------------------------------------------------------#
+
+add_executable(
+    omnitrace-critical-trace
+    ${CMAKE_CURRENT_LIST_DIR}/critical-trace.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/critical-trace.hpp
+    $<TARGET_OBJECTS:omnitrace::omnitrace-object-library>)
+
+target_include_directories(omnitrace-critical-trace
+                           PRIVATE ${CMAKE_CURRENT_LIST_DIR}/include)
+target_compile_definitions(omnitrace-critical-trace PRIVATE OMNITRACE_EXTERN_COMPONENTS=0)
+target_link_libraries(omnitrace-critical-trace
+                      PRIVATE omnitrace::omnitrace-interface-library)
+set_target_properties(omnitrace-critical-trace PROPERTIES INSTALL_RPATH_USE_LINK_PATH ON)
+
+install(
+    TARGETS omnitrace-critical-trace
+    DESTINATION bin
+    OPTIONAL)
@@ -0,0 +1,952 @@
+// 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 "critical-trace.hpp"
+
+#include "library/api.hpp"
+#include "library/perfetto.hpp"
+
+#include <timemory/hash/types.hpp>
+
+#include <cstdlib>
+
+namespace config         = omnitrace::config;
+namespace critical_trace = omnitrace::critical_trace;
+
+int
+main(int argc, char** argv)
+{
+    omnitrace_init_library();
+
+    config::set_setting_value("OMNITRACE_CRITICAL_TRACE", true);
+    // config::set_setting_value("OMNITRACE_CRITICAL_TRACE_DEBUG", true);
+    config::set_setting_value<int64_t>("OMNITRACE_CRITICAL_TRACE_COUNT", 500);
+    config::set_setting_value<int64_t>("OMNITRACE_CRITICAL_TRACE_PER_ROW", 100);
+    config::set_setting_value<int64_t>("OMNITRACE_CRITICAL_TRACE_NUM_THREADS",
+                                       std::thread::hardware_concurrency());
+    config::set_setting_value("OMNITRACE_CRITICAL_TRACE_SERIALIZE_NAMES", true);
+
+    perfetto::TracingInitArgs               args{};
+    perfetto::TraceConfig                   cfg{};
+    perfetto::protos::gen::TrackEventConfig track_event_cfg{};
+
+    auto  shmem_size_hint = config::get_perfetto_shmem_size_hint();
+    auto  buffer_size     = config::get_perfetto_buffer_size();
+    auto* buffer_config   = cfg.add_buffers();
+
+    buffer_config->set_size_kb(buffer_size);
+    buffer_config->set_fill_policy(
+        perfetto::protos::gen::TraceConfig_BufferConfig_FillPolicy_DISCARD);
+
+    auto* ds_cfg = cfg.add_data_sources()->mutable_config();
+    ds_cfg->set_name("track_event");
+    ds_cfg->set_track_event_config_raw(track_event_cfg.SerializeAsString());
+
+    args.backends |= perfetto::kInProcessBackend;
+    args.shmem_size_hint_kb = shmem_size_hint;
+
+    perfetto::Tracing::Initialize(args);
+    perfetto::TrackEvent::Register();
+
+    auto tracing_session = perfetto::Tracing::NewTrace();
+    tracing_session->Setup(cfg);
+    tracing_session->StartBlocking();
+
+    for(int i = 1; i < argc; ++i)
+    {
+        critical_trace::complete_call_chain = {};
+        OMNITRACE_BASIC_PRINT_F("Loading call-chain %s...\n", argv[i]);
+        critical_trace::load_call_chain(argv[i], "call_chain",
+                                        critical_trace::complete_call_chain);
+        for(const auto& itr : *tim::get_hash_ids())
+            critical_trace::complete_hash_ids.emplace(itr.second);
+        OMNITRACE_BASIC_PRINT_F("Computing critical trace for %s...\n", argv[i]);
+        critical_trace::compute_critical_trace();
+    }
+
+    // Make sure the last event is closed for this example.
+    perfetto::TrackEvent::Flush();
+
+    OMNITRACE_DEBUG_F("Stopping the blocking perfetto trace sessions...\n");
+    tracing_session->StopBlocking();
+
+    OMNITRACE_DEBUG_F("Getting the trace data...\n");
+    std::vector<char> trace_data{ tracing_session->ReadTraceBlocking() };
+
+    if(trace_data.empty())
+    {
+        OMNITRACE_BASIC_PRINT_F(
+            "> trace data is empty. File '%s' will not be written...\n",
+            config::get_perfetto_output_filename().c_str());
+    }
+    else
+    {
+        // Write the trace into a file.
+        OMNITRACE_CONDITIONAL_BASIC_PRINT_F(
+            config::get_verbose() >= 0,
+            "> Outputting '%s' (%.2f KB / %.2f MB / %.2f GB)... ",
+            config::get_perfetto_output_filename().c_str(),
+            static_cast<double>(trace_data.size()) / tim::units::KB,
+            static_cast<double>(trace_data.size()) / tim::units::MB,
+            static_cast<double>(trace_data.size()) / tim::units::GB);
+
+        std::ofstream ofs{};
+        if(!tim::filepath::open(ofs, config::get_perfetto_output_filename(),
+                                std::ios::out | std::ios::binary))
+        {
+            OMNITRACE_BASIC_PRINT_F("> Error opening '%s'...\n",
+                                    config::get_perfetto_output_filename().c_str());
+            return EXIT_FAILURE;
+        }
+        else
+        {
+            // Write the trace into a file.
+            if(config::get_verbose() >= 0) fprintf(stderr, "Done\n");
+            ofs.write(&trace_data[0], trace_data.size());
+        }
+        ofs.close();
+    }
+}
+
+namespace omnitrace
+{
+namespace critical_trace
+{
+namespace
+{
+//--------------------------------------------------------------------------------------//
+
+std::string
+get_perf_name(std::string _func)
+{
+    const auto _npos = std::string::npos;
+    auto       _pos  = std::string::npos;
+    while((_pos = _func.find('_')) != _npos)
+        _func = _func.replace(_pos, 1, " ");
+    if(_func.length() > 0) _func.at(0) = std::toupper(_func.at(0));
+    return _func;
+}
+
+//--------------------------------------------------------------------------------------//
+
+void
+save_call_graph(const std::string& _fname, const std::string& _label,
+                const call_graph_t& _call_graph, bool _msg = false,
+                std::string _func = {})
+{
+    OMNITRACE_CT_DEBUG("\n");
+
+    using perfstats_t =
+        tim::lightweight_tuple<comp::wall_clock, comp::peak_rss, comp::page_rss>;
+    perfstats_t _perf{ get_perf_name(__FUNCTION__) };
+    _perf.start();
+
+    std::stringstream oss{};
+    {
+        namespace cereal = tim::cereal;
+        auto ar = tim::policy::output_archive<cereal::MinimalJSONOutputArchive>::get(oss);
+
+        auto _hash_map = *tim::hash::get_hash_ids();
+        for(auto& itr : _hash_map)
+            itr.second = tim::demangle(itr.second);
+        ar->setNextName("omnitrace");
+        ar->startNode();
+        (*ar)(cereal::make_nvp("hash_map", _hash_map));
+        ar->setNextName(_label.c_str());
+        ar->startNode();
+        serialize_graph(*ar, _call_graph);
+        ar->finishNode();
+        ar->finishNode();
+    }
+
+    std::ofstream ofs{};
+    if(tim::filepath::open(ofs, _fname))
+    {
+        if(_msg)
+        {
+            if(_func.empty()) _func = __FUNCTION__;
+            OMNITRACE_CONDITIONAL_BASIC_PRINT(get_verbose() >= 0,
+                                              "[%s] Outputting '%s'...\n", _func.c_str(),
+                                              _fname.c_str());
+        }
+        ofs << oss.str() << std::endl;
+    }
+
+    _perf.stop();
+    if(_msg)
+    {
+        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).substr(4).c_str());
+    }
+}
+
+void
+save_critical_trace(const std::string& _fname, const std::string& _label,
+                    const std::vector<call_chain>& _cchain, bool _msg = false,
+                    std::string _func = {})
+{
+    OMNITRACE_CT_DEBUG("\n");
+
+    using perfstats_t =
+        tim::lightweight_tuple<comp::wall_clock, comp::peak_rss, comp::page_rss>;
+    perfstats_t _perf{ get_perf_name(__FUNCTION__) };
+    _perf.start();
+
+    auto _save = [&](std::ostream& _os) {
+        namespace cereal = tim::cereal;
+        auto ar = tim::policy::output_archive<cereal::MinimalJSONOutputArchive>::get(_os);
+
+        auto _hash_map = *tim::hash::get_hash_ids();
+        for(auto& itr : _hash_map)
+            itr.second = tim::demangle(itr.second);
+        ar->setNextName("omnitrace");
+        ar->startNode();
+        (*ar)(cereal::make_nvp("hash_map", _hash_map),
+              cereal::make_nvp(_label.c_str(), _cchain));
+        ar->finishNode();
+    };
+
+    std::ofstream ofs{};
+    if(tim::filepath::open(ofs, _fname))
+    {
+        if(_msg)
+        {
+            if(_func.empty()) _func = __FUNCTION__;
+            OMNITRACE_CONDITIONAL_BASIC_PRINT(get_verbose() >= 0,
+                                              "[%s] Outputting '%s'...\n", _func.c_str(),
+                                              _fname.c_str());
+        }
+        std::stringstream oss{};
+        if(_cchain.size() > 1000)
+        {
+            _save(ofs);
+        }
+        else
+        {
+            _save(oss);
+            ofs << oss.str() << std::endl;
+        }
+    }
+
+    _perf.stop();
+    if(_msg)
+    {
+        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).substr(4).c_str());
+    }
+}
+
+void
+save_call_chain_text(const std::string& _fname, const call_chain& _call_chain,
+                     bool _msg = false, std::string _func = {})
+{
+    OMNITRACE_CT_DEBUG("\n");
+
+    using perfstats_t =
+        tim::lightweight_tuple<comp::wall_clock, comp::peak_rss, comp::page_rss>;
+    perfstats_t _perf{ get_perf_name(__FUNCTION__) };
+    _perf.start();
+
+    std::ofstream ofs{};
+    if(tim::filepath::open(ofs, _fname))
+    {
+        if(_msg)
+        {
+            if(_func.empty()) _func = __FUNCTION__;
+            OMNITRACE_CONDITIONAL_BASIC_PRINT(get_verbose() >= 0,
+                                              "[%s] Outputting '%s'...\n", _func.c_str(),
+                                              _fname.c_str());
+        }
+        ofs << _call_chain << "\n";
+    }
+
+    _perf.stop();
+    if(_msg)
+    {
+        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).substr(4).c_str());
+    }
+}
+
+void
+save_call_chain_json(const std::string& _fname, const std::string& _label,
+                     const call_chain& _call_chain, bool _msg = false,
+                     std::string _func = {})
+{
+    OMNITRACE_CT_DEBUG("\n");
+
+    using perfstats_t =
+        tim::lightweight_tuple<comp::wall_clock, comp::peak_rss, comp::page_rss>;
+    perfstats_t _perf{ get_perf_name(__FUNCTION__) };
+    _perf.start();
+
+    auto _save = [&](std::ostream& _os) {
+        namespace cereal = tim::cereal;
+        auto ar = tim::policy::output_archive<cereal::MinimalJSONOutputArchive>::get(_os);
+
+        auto _hash_map = *tim::hash::get_hash_ids();
+        for(auto& itr : _hash_map)
+            itr.second = tim::demangle(itr.second);
+        ar->setNextName("omnitrace");
+        ar->startNode();
+        (*ar)(cereal::make_nvp("hash_map", _hash_map),
+              cereal::make_nvp(_label.c_str(), _call_chain));
+        ar->finishNode();
+    };
+
+    std::ofstream ofs{};
+    if(tim::filepath::open(ofs, _fname))
+    {
+        if(_msg)
+        {
+            if(_func.empty()) _func = __FUNCTION__;
+            OMNITRACE_CONDITIONAL_BASIC_PRINT(get_verbose() >= 0,
+                                              "[%s] Outputting '%s'...\n", _func.c_str(),
+                                              _fname.c_str());
+        }
+        std::stringstream oss{};
+        if(_call_chain.size() > 100000)
+        {
+            _save(ofs);
+        }
+        else
+        {
+            _save(oss);
+            ofs << oss.str() << std::endl;
+        }
+    }
+
+    _perf.stop();
+    if(_msg)
+    {
+        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).substr(4).c_str());
+    }
+}
+
+void
+load_call_chain(const std::string& _fname, const std::string& _label,
+                call_chain& _call_chain)
+{
+    std::ifstream ifs{};
+    ifs.open(_fname);
+    if(ifs && ifs.is_open())
+    {
+        namespace cereal = tim::cereal;
+        auto ar          = tim::policy::input_archive<cereal::JSONInputArchive>::get(ifs);
+
+        ar->setNextName("omnitrace");
+        ar->startNode();
+        (*ar)(cereal::make_nvp(_label.c_str(), _call_chain));
+        ar->finishNode();
+    }
+}
+
+auto
+get_indexed(const call_chain& _chain)
+{
+    OMNITRACE_CT_DEBUG("\n");
+    std::map<int64_t, std::vector<entry>> _indexed = {};
+
+    // allocate for all cpu correlation ids
+    for(const auto& itr : _chain)
+    {
+        _indexed.emplace(static_cast<int64_t>(itr.cpu_cid), std::vector<entry>{});
+        _indexed.emplace(static_cast<int64_t>(itr.parent_cid), std::vector<entry>{});
+    }
+
+    // index based on parent correlation id
+    for(const auto& itr : _chain)
+    {
+        if(itr.depth < 1 && itr.phase == Phase::BEGIN) continue;
+        _indexed[static_cast<int64_t>(itr.parent_cid)].emplace_back(itr);
+    }
+
+    for(auto& itr : _indexed)
+        std::sort(itr.second.begin(), itr.second.end(),
+                  [](const entry& lhs, const entry& rhs) {
+                      // return lhs.cpu_cid < rhs.cpu_cid;
+                      return lhs.begin_ns < rhs.begin_ns;
+                  });
+
+    return _indexed;
+}
+
+void
+find_children(PTL::ThreadPool& _tp, call_graph_t& _graph, const call_chain& _chain)
+{
+    OMNITRACE_CT_DEBUG("\n");
+
+    using iterator_t      = call_graph_sibling_itr_t;
+    using itr_entry_vec_t = std::vector<std::pair<iterator_t, entry>>;
+    using task_group_t    = PTL::TaskGroup<void>;
+
+    auto                                _indexed = get_indexed(_chain);
+    std::map<entry, std::vector<entry>> _entry_map{};
+
+    // allocate all entries
+    OMNITRACE_CT_DEBUG_F("Allocating...\n");
+    for(const auto& itr : _chain)
+    {
+        auto _ins = _entry_map.emplace(itr, std::vector<entry>{});
+        if(!_ins.second)
+        {
+            auto _existing = _ins.first->first;
+            OMNITRACE_BASIC_PRINT("Warning! Duplicate entry for [%s] :: [%s]\n",
+                                  JOIN("", _existing).c_str(), JOIN("", itr).c_str());
+        }
+    }
+
+    task_group_t _tg{ &_tp };
+    OMNITRACE_CT_DEBUG_F("Parallel mapping...\n");
+    for(const auto& itr : _chain)
+    {
+        _tg.run([&]() { _entry_map[itr] = _indexed.at(itr.cpu_cid); });
+    }
+    _tg.join();
+
+    std::function<void(iterator_t, const entry&)> _recursive_func;
+    _recursive_func = [&](iterator_t itr, const entry& _v) {
+        auto _child    = _graph.append_child(itr, _v);
+        auto _children = std::move(_entry_map[_v]);
+        _entry_map[_v].clear();
+        for(auto&& vitr : _children)
+        {
+            _recursive_func(_child, vitr);
+        }
+    };
+
+    // the recursive version of _func + _loop_func has a tendency to overflow the stack
+    auto _func = [&](iterator_t itr, const entry& _v) {
+        auto _child    = _graph.append_child(itr, _v);
+        auto _children = std::move(_entry_map[_v]);
+        _entry_map[_v].clear();
+        itr_entry_vec_t _data{};
+        for(auto&& vitr : _children)
+            _data.emplace_back(_child, vitr);
+        return _data;
+    };
+
+    auto _loop_func = [&_func](itr_entry_vec_t& _data) {
+        auto _inp = _data;
+        _data.clear();
+        for(auto itr : _inp)
+        {
+            for(auto&& fitr : _func(itr.first, itr.second))
+                _data.emplace_back(std::move(fitr));
+        }
+        // if data is empty return false so we can break out of while loop
+        return !_data.empty();
+    };
+
+    if(!_indexed.at(-1).empty())
+    {
+        OMNITRACE_CT_DEBUG_F("Setting root (line %i)...\n", __LINE__);
+        _graph.set_head(_indexed.at(-1).front());
+    }
+    else
+    {
+        OMNITRACE_CT_DEBUG_F("Setting root (line %i)...\n", __LINE__);
+        auto  _depth = static_cast<uint16_t>(-1);
+        entry _root{ 0, Device::NONE, Phase::NONE, _depth, 0, 0, 0, 0, 0, 0, 0 };
+        _graph.set_head(_root);
+    }
+
+    iterator_t _root = _graph.begin();
+    for(auto&& itr : _entry_map)
+    {
+        if(itr.first.depth == _root->depth + 1)
+        {
+            OMNITRACE_CT_DEBUG_F("Generating call-graph...\n");
+            // _recursive_func(_root, itr.first);
+            itr_entry_vec_t _data = _func(_root, itr.first);
+            while(_loop_func(_data))
+            {}
+        }
+    }
+}
+
+void
+find_sequences(PTL::ThreadPool& _tp, call_graph_t& _graph,
+               std::vector<call_chain>& _chain)
+{
+    OMNITRACE_CT_DEBUG("\n");
+    /*
+    using sibling_itr_t = call_graph_sibling_itr_t;
+    using sibling_vec_t = std::vector<sibling_itr_t>;
+    using sibling_map_t = std::map<int64_t, sibling_vec_t>;
+
+    std::function<void(sibling_map_t & _v, sibling_itr_t root)> _no_overlap{};
+    _no_overlap = [&](sibling_map_t& _v, sibling_itr_t root) {
+        sibling_map_t _l{};
+        int64_t       n = _graph.number_of_children(root);
+        if(n == 0) return;
+
+        //_graph.sort(sibling_itr_t{ root },
+        //            [](auto lhs, auto rhs) { return lhs.get_cost() > rhs.get_cost(); });
+
+        for(int64_t i = 0; i < n; ++i)
+        {
+            if(_l.empty())
+            {
+                auto itr = _graph.child(root, i);
+                _l[itr->tid].emplace_back(itr);
+            }
+            else
+            {
+                auto itr       = _graph.child(root, i);
+                bool _overlaps = false;
+                for(auto& litr : _l[itr->tid])
+                {
+                    if(litr->device == itr->device && litr->get_overlap(*itr) > 0)
+                    {
+                        _overlaps = true;
+                        break;
+                    }
+                }
+                if(!_overlaps) _l[itr->tid].emplace_back(itr);
+            }
+        }
+        for(auto& iitr : _l)
+        {
+            for(auto itr : iitr.second)
+            {
+                _v[iitr.first].emplace_back(itr);
+                _no_overlap(_v, itr);
+            }
+        }
+    };
+
+    std::map<int64_t, sibling_vec_t> _tot{};
+    for(sibling_itr_t itr = _graph.begin(); itr != _graph.end(); ++itr)
+    {
+        _no_overlap(_tot, itr);
+    }
+
+    for(const auto& iitr : _tot)
+    {
+        call_chain _cc{};
+        _cc.emplace_back(*_graph.begin());
+        for(const auto& itr : iitr.second)
+            _cc.emplace_back(*itr);
+        _chain.emplace_back(_cc);
+    }
+
+    (void) _tp;
+    */
+
+    using iterator_t = call_graph_preorder_itr_t;
+    std::vector<iterator_t> _end_nodes{};
+    size_t                  _n = 0;
+    for(iterator_t itr = _graph.begin(); itr != _graph.end(); ++itr, ++_n)
+    {
+        auto _nchild = _graph.number_of_children(itr);
+        if(_nchild > 0)
+        {
+            OMNITRACE_CT_DEBUG("Skipping node #%zu with %u children :: %s\n", _n, _nchild,
+                               JOIN("", *itr).c_str());
+            continue;
+        }
+        _end_nodes.emplace_back(itr);
+    }
+    OMNITRACE_CT_DEBUG("Number of end nodes: %zu\n", _end_nodes.size());
+    _chain.resize(_end_nodes.size());
+
+    auto _construct = [&](size_t i) {
+        auto itr = _end_nodes.at(i);
+        while(itr != nullptr && _graph.is_valid(itr))
+        {
+            _chain.at(i).emplace_back(*itr);
+            itr = _graph.parent(itr);
+        }
+        std::reverse(_chain.at(i).begin(), _chain.at(i).end());
+        std::sort(
+            _chain.at(i).begin(), _chain.at(i).end(),
+            [](const entry& lhs, const entry& rhs) { return lhs.begin_ns > rhs.end_ns; });
+    };
+
+    PTL::TaskGroup<void> _tg{ &_tp };
+    for(size_t i = 0; i < _end_nodes.size(); ++i)
+        _tg.run(_construct, i);
+    _tg.join();
+
+    std::sort(_chain.begin(), _chain.end(),
+              [](const call_chain& lhs, const call_chain& rhs) {
+                  return lhs.get_cost() > rhs.get_cost();
+              });
+
+    /*
+    std::vector<call_chain> _new_chain{};
+    for(auto& itr : _chain)
+    {
+        if(itr.empty()) continue;
+        if(_new_chain.empty())
+        {
+            _new_chain.emplace_back(std::move(itr));
+            continue;
+        }
+        std::sort(itr.begin(), itr.end(), [](const entry& lhs, const entry& rhs) {
+            return lhs.get_cost() > rhs.get_cost();
+        });
+
+        call_chain* _append_chain = nullptr;
+        for(auto& nitr : _new_chain)
+        {
+            if(nitr.at(0).tid == itr.at(0).tid && nitr.at(0).get_overlap(itr.at(0)) <= 0)
+            {
+                _append_chain = &nitr;
+                break;
+            }
+        }
+
+        if(_append_chain)
+        {
+            for(auto& oitr : itr)
+                _append_chain->emplace_back(oitr);
+            std::sort(_append_chain->begin(), _append_chain->end(),
+                      [](const entry& lhs, const entry& rhs) {
+                          return lhs.get_cost() > rhs.get_cost();
+                      });
+        }
+        else
+        {
+            _new_chain.emplace_back(std::move(itr));
+        }
+        itr.clear();
+    }
+
+    _chain = _new_chain;*/
+}
+
+template <typename ArchiveT, typename T, typename AllocatorT>
+void
+serialize_graph(ArchiveT& ar, const tim::graph<T, AllocatorT>& t)
+{
+    OMNITRACE_CT_DEBUG("\n");
+
+    namespace cereal = tim::cereal;
+    using iterator_t = typename tim::graph<T, AllocatorT>::sibling_iterator;
+
+    ar(cereal::make_nvp("graph_nodes", t.size()));
+    ar.setNextName("graph");
+    ar.startNode();
+    ar.makeArray();
+    for(iterator_t itr = t.begin(); itr != t.end(); ++itr)
+        serialize_subgraph(ar, t, itr);
+    ar.finishNode();
+}
+
+template <typename ArchiveT, typename T, typename AllocatorT>
+void
+serialize_subgraph(ArchiveT& ar, const tim::graph<T, AllocatorT>& _graph,
+                   typename tim::graph<T, AllocatorT>::iterator _root)
+{
+    using iterator_t = typename tim::graph<T, AllocatorT>::sibling_iterator;
+
+    if(_graph.empty()) return;
+
+    ar.setNextName("node");
+    ar.startNode();
+    ar(*_root);
+    {
+        ar.setNextName("children");
+        ar.startNode();
+        ar.makeArray();
+        for(iterator_t itr = _graph.begin(_root); itr != _graph.end(_root); ++itr)
+            serialize_subgraph(ar, _graph, itr);
+        ar.finishNode();
+    }
+    ar.finishNode();
+}
+
+template <Device DevT>
+std::vector<call_chain>
+get_top(const std::vector<call_chain>& _chain, size_t _count)
+{
+    OMNITRACE_CT_DEBUG("\n");
+    std::vector<call_chain> _data{};
+    _data.reserve(_count);
+    for(const auto& itr : _chain)
+    {
+        if(_data.size() >= _count) break;
+        if(itr.query<>([](const entry& _v) {
+               return (DevT == Device::ANY) ? true : (_v.device == DevT);
+           }))
+        {
+            _data.emplace_back(itr);
+        }
+    }
+    return _data;
+}
+
+template <Device DevT>
+void
+generate_perfetto(const std::vector<call_chain>& _data)
+{
+    OMNITRACE_CT_DEBUG("\n");
+
+    auto _nrows = std::min<size_t>(get_critical_trace_per_row(), _data.size());
+
+    // run in separate thread(s) so that it ends up in unique row
+    if(_nrows < 1) _nrows = _data.size();
+
+    std::string _dev    = (DevT == Device::NONE)  ? ""
+                          : (DevT == Device::ANY) ? "CPU + GPU "
+                          : (DevT == Device::CPU) ? "CPU "
+                                                  : "GPU ";
+    std::string _cpname = _dev + "CritPath";
+    auto        _func   = [&](size_t _idx, size_t _beg, size_t _end) {
+        if(DevT != Device::NONE)
+        {
+            if(_nrows != 1)
+                threading::set_thread_name(TIMEMORY_JOIN(" ", _cpname, _idx).c_str());
+            else
+                threading::set_thread_name(_cpname.c_str());
+        }
+        // ensure all hash ids exist
+        copy_hash_ids();
+        std::set<entry> _used{};
+        for(size_t i = _beg; i < _end; ++i)
+        {
+            if(i >= _data.size()) break;
+            _data.at(i).generate_perfetto<DevT>(_used);
+        }
+    };
+
+    for(size_t i = 0; i < _data.size(); i += _nrows)
+    {
+        if(DevT == Device::NONE)
+            _func(i, i, i + _nrows);
+        else
+            std::thread{ _func, i, i, i + _nrows }.join();
+    }
+}
+
+template <typename Tp, template <typename...> class ContainerT, typename... Args,
+          typename FuncT = bool (*)(const Tp&, const Tp&)>
+inline Tp*
+find(
+    const Tp& _v, ContainerT<Tp, Args...>& _vec,
+    FuncT&& _func = [](const Tp& _lhs, const Tp& _rhs) { return (_lhs == _rhs); })
+{
+    for(auto& itr : _vec)
+    {
+        if(std::forward<FuncT>(_func)(_v, itr)) return &itr;
+    }
+    return nullptr;
+};
+
+template <typename FuncT = bool (*)(const entry&, const entry&)>
+inline entry*
+find(
+    const entry& _v, call_chain& _vec,
+    FuncT&& _func = [](const entry& _lhs, const entry& _rhs) { return (_lhs == _rhs); })
+{
+    return find(_v, reinterpret_cast<std::vector<entry>&>(_vec),
+                std::forward<FuncT>(_func));
+}
+
+void
+squash_critical_path(call_chain& _targ)
+{
+    OMNITRACE_CT_DEBUG("\n");
+    static auto _strict_equal = [](const entry& _lhs, const entry& _rhs) {
+        auto _same_phase  = (_lhs.phase == _rhs.phase);
+        bool _phase_check = true;
+        if(_same_phase) _phase_check = (_lhs.get_timestamp() == _rhs.get_timestamp());
+        return (_lhs == _rhs && _lhs.parent_cid == _rhs.parent_cid && _phase_check);
+    };
+
+    std::sort(_targ.begin(), _targ.end());
+
+    call_chain _squashed{};
+    for(auto& itr : _targ)
+    {
+        if(itr.phase == Phase::DELTA)
+        {
+            _squashed.emplace_back(itr);
+        }
+        else if(itr.phase == Phase::BEGIN)
+        {
+            if(!find(itr, _squashed, _strict_equal)) _squashed.emplace_back(itr);
+        }
+        else
+        {
+            entry* _match = nullptr;
+            if((_match = find(itr, _squashed)) != nullptr)
+                *_match += itr;
+            else
+                _squashed.emplace_back(itr);
+        }
+    }
+
+    std::swap(_targ, _squashed);
+    std::sort(_targ.begin(), _targ.end());
+}
+
+void
+compute_critical_trace()
+{
+    OMNITRACE_CT_DEBUG_F("Generating critical trace...\n");
+
+    // ensure all hash ids exist
+    copy_hash_ids();
+
+    using perfstats_t =
+        tim::lightweight_tuple<comp::wall_clock, comp::cpu_clock, comp::cpu_util,
+                               comp::peak_rss, comp::page_rss>;
+
+    perfstats_t _ct_perf{};
+    _ct_perf.start();
+
+    auto _report_perf = [](auto& _perf, const char* _func, const std::string& _label) {
+        _perf.stop().rekey(_label);
+        OMNITRACE_BASIC_PRINT("[%s] %s\n", _func, JOIN("", _perf).substr(5).c_str());
+        OMNITRACE_BASIC_PRINT("\n");
+        _perf.reset().start();
+    };
+
+    OMNITRACE_BASIC_PRINT("\n");
+
+    try
+    {
+        PTL::ThreadPool _tp{ get_critical_trace_num_threads(), []() { copy_hash_ids(); },
+                             []() {} };
+        _tp.set_verbose(-1);
+        PTL::TaskGroup<void> _tg{ &_tp };
+
+        perfstats_t _perf{};
+        _perf.start();
+
+        OMNITRACE_BASIC_PRINT_F("sorting %zu call chain entries\n",
+                                complete_call_chain.size());
+
+        // sort the complete call chain
+        std::sort(complete_call_chain.begin(), complete_call_chain.end());
+        _report_perf(_perf, __FUNCTION__, "sorting call chain");
+
+        OMNITRACE_BASIC_PRINT_F("squashing call chain...\n");
+
+        // squash the critical path (combine start/stop into delta)
+        squash_critical_path(complete_call_chain);
+        _report_perf(_perf, __FUNCTION__, "squashing critical path");
+
+        // generate the perfetto
+        if(config::get_use_perfetto())
+        {
+            OMNITRACE_BASIC_PRINT_F("generating perfetto for call chain...\n");
+            generate_perfetto<Device::NONE>({ complete_call_chain });
+            generate_perfetto<Device::CPU>({ complete_call_chain });
+            generate_perfetto<Device::GPU>({ complete_call_chain });
+            _report_perf(_perf, __FUNCTION__, "perfetto generation");
+        }
+
+        OMNITRACE_BASIC_PRINT_F("finding children...\n");
+        call_graph_t _graph{};
+        find_children(_tp, _graph, complete_call_chain);
+        _report_perf(_perf, __FUNCTION__, "finding children");
+
+        // sort the call-graph based on cost
+        OMNITRACE_BASIC_PRINT_F("sorting %zu call-graph entries...\n", _graph.size() - 1);
+        _graph.sort([](auto lhs, auto rhs) { return lhs.get_cost() > rhs.get_cost(); },
+                    [&_tg](auto _f) { _tg.run(_f); }, [&_tg]() { _tg.join(); });
+        _report_perf(_perf, __FUNCTION__, "call-graph sort");
+
+        OMNITRACE_BASIC_PRINT_F("saving call-graph...\n");
+        save_call_graph(tim::settings::compose_output_filename("call-graph", ".json"),
+                        "call_graph", _graph, true, __FUNCTION__);
+        _report_perf(_perf, __FUNCTION__, "saving call-graph");
+
+        OMNITRACE_BASIC_PRINT_F("finding sequences...\n");
+        std::vector<call_chain> _top{};
+        find_sequences(_tp, _graph, _top);
+        _report_perf(_perf, __FUNCTION__, "call-graph sequence search");
+
+        OMNITRACE_BASIC_PRINT_F("number of sequences found: %zu (%zu)...\n", _top.size(),
+                                (_top.empty()) ? 0 : _top.at(0).size());
+
+        if(get_critical_trace_count() == 0)
+        {
+            OMNITRACE_CT_DEBUG_F("saving critical trace...\n");
+            save_critical_trace(
+                tim::settings::compose_output_filename("critical-trace", ".json"),
+                "critical_trace", _top, true, __FUNCTION__);
+        }
+        else
+        {
+            // get the top CPU critical traces
+            OMNITRACE_BASIC_PRINT_F("getting top CPU functions...\n");
+            auto _top_cpu = get_top<Device::CPU>(_top, get_critical_trace_count());
+
+            // get the top GPU critical traces
+            OMNITRACE_BASIC_PRINT_F("getting top GPU functions...\n");
+            auto _top_gpu = get_top<Device::GPU>(_top, get_critical_trace_count());
+
+            // get the top CPU + GPU critical traces
+            OMNITRACE_BASIC_PRINT_F("getting top CPU + GPU functions...\n");
+            auto _top_any = get_top<Device::ANY>(_top, get_critical_trace_count());
+
+            if(!_top_cpu.empty())
+            {
+                OMNITRACE_BASIC_PRINT_F(
+                    "generating %zu perfetto CPU critical traces...\n", _top_cpu.size());
+                if(config::get_use_perfetto()) generate_perfetto<Device::CPU>(_top_cpu);
+                OMNITRACE_CT_DEBUG_F("saving CPU critical traces...\n");
+                save_critical_trace(
+                    tim::settings::compose_output_filename("critical-trace-cpu", ".json"),
+                    "critical_trace", _top_cpu, true, __FUNCTION__);
+            }
+
+            if(!_top_gpu.empty())
+            {
+                OMNITRACE_BASIC_PRINT_F(
+                    "generating %zu perfetto GPU critical traces...\n", _top_gpu.size());
+                if(config::get_use_perfetto()) generate_perfetto<Device::GPU>(_top_gpu);
+                OMNITRACE_CT_DEBUG_F("saving GPU critical traces...\n");
+                save_critical_trace(
+                    tim::settings::compose_output_filename("critical-trace-gpu", ".json"),
+                    "critical_trace", _top_gpu, true, __FUNCTION__);
+            }
+
+            if(!_top_any.empty())
+            {
+                OMNITRACE_BASIC_PRINT_F(
+                    "generating %zu perfetto CPU + GPU critical traces...\n",
+                    _top_gpu.size());
+                if(config::get_use_perfetto()) generate_perfetto<Device::ANY>(_top_gpu);
+                OMNITRACE_CT_DEBUG_F("saving CPU + GPU critical traces...\n");
+                save_critical_trace(
+                    tim::settings::compose_output_filename("critical-trace-any", ".json"),
+                    "critical_trace", _top_any, true, __FUNCTION__);
+            }
+        }
+
+        _tg.join();
+        _tp.destroy_threadpool();
+    } catch(std::exception& e)
+    {
+        OMNITRACE_BASIC_PRINT("Thread exited '%s' with exception: %s\n", __FUNCTION__,
+                              e.what());
+        TIMEMORY_CONDITIONAL_DEMANGLED_BACKTRACE(true, 32);
+    }
+
+    _report_perf(_ct_perf, __FUNCTION__, "critical trace computation");
+}
+}  // namespace
+}  // namespace critical_trace
+}  // namespace omnitrace
@@ -0,0 +1,113 @@
+// 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/config.hpp"
+#include "library/critical_trace.hpp"
+#include "library/debug.hpp"
+#include "library/defines.hpp"
+#include "library/perfetto.hpp"
+#include "library/ptl.hpp"
+
+#include <PTL/ThreadPool.hh>
+#include <timemory/backends/dmp.hpp>
+#include <timemory/backends/threading.hpp>
+#include <timemory/hash/types.hpp>
+#include <timemory/tpls/cereal/cereal/archives/json.hpp>
+#include <timemory/tpls/cereal/cereal/cereal.hpp>
+#include <timemory/utility/macros.hpp>
+#include <timemory/utility/types.hpp>
+#include <timemory/utility/utility.hpp>
+
+#include <cctype>
+#include <cstdint>
+#include <exception>
+#include <iomanip>
+#include <sstream>
+#include <stdexcept>
+#include <utility>
+
+namespace omnitrace
+{
+namespace critical_trace
+{
+namespace
+{
+using call_graph_t              = tim::graph<entry>;
+using call_graph_itr_t          = typename call_graph_t::iterator;
+using call_graph_sibling_itr_t  = typename call_graph_t::sibling_iterator;
+using call_graph_preorder_itr_t = typename call_graph_t::pre_order_iterator;
+
+hash_ids   complete_hash_ids{};
+call_chain complete_call_chain{};
+std::mutex complete_call_mutex{};
+
+void
+update_critical_path(call_chain _chain, int64_t _tid);
+
+void
+load_call_chain(const std::string& _fname, const std::string& _label,
+                call_chain& _call_chain);
+
+void
+compute_critical_trace();
+
+void
+find_children(PTL::ThreadPool& _tp, call_graph_t& _graph, const call_chain& _chain);
+
+void
+find_sequences(PTL::ThreadPool& _tp, call_graph_t& _graph,
+               std::vector<call_chain>& _chain);
+
+void
+find_sequences(PTL::ThreadPool& _tp, call_graph_t& _graph, call_graph_itr_t _root,
+               std::vector<call_chain>& _chain);
+
+template <typename ArchiveT, typename T, typename AllocatorT>
+void
+serialize_graph(ArchiveT& ar, const tim::graph<T, AllocatorT>& _graph);
+
+template <typename ArchiveT, typename T, typename AllocatorT>
+void
+serialize_subgraph(ArchiveT& ar, const tim::graph<T, AllocatorT>& _graph,
+                   typename tim::graph<T, AllocatorT>::iterator _root);
+
+void
+compute_critical_trace();
+
+template <Device DevT>
+void
+generate_perfetto(const std::vector<call_chain>& _data);
+
+inline void
+copy_hash_ids()
+{
+    // make copy to avoid parallel iteration issues
+    auto _hash_ids = complete_hash_ids;
+    // ensure all hash ids exist
+    for(const auto& itr : _hash_ids)
+        tim::hash::add_hash_id(itr);
+}
+}  // namespace
+}  // namespace critical_trace
+}  // namespace omnitrace
@@ -0,0 +1,37 @@
+# ------------------------------------------------------------------------------#
+#
+# omnitrace-exe target
+#
+# ------------------------------------------------------------------------------#
+
+add_executable(
+    omnitrace-exe
+    ${_EXCLUDE} ${CMAKE_CURRENT_LIST_DIR}/omnitrace.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/omnitrace.hpp ${CMAKE_CURRENT_LIST_DIR}/details.cpp)
+
+target_link_libraries(
+    omnitrace-exe
+    PRIVATE omnitrace::omnitrace-headers
+            omnitrace::omnitrace-dyninst
+            omnitrace::omnitrace-compile-options
+            $<BUILD_INTERFACE:timemory::timemory-headers>
+            $<IF:$<BOOL:${OMNITRACE_USE_SANITIZER}>,omnitrace::omnitrace-sanitizer,>)
+
+set_target_properties(
+    omnitrace-exe
+    PROPERTIES
+        OUTPUT_NAME omnitrace
+        INSTALL_RPATH_USE_LINK_PATH ON
+        INSTALL_RPATH
+        "\$ORIGIN/../${CMAKE_INSTALL_LIBDIR}:\$ORIGIN/../${CMAKE_INSTALL_LIBDIR}/timemory/libunwind:\$ORIGIN/../${CMAKE_INSTALL_LIBDIR}/dyninst-tpls/lib"
+    )
+
+if(CMAKE_BUILD_TYPE MATCHES "^(DEBUG|Debug)")
+    string(REPLACE " " ";" _FLAGS "${CMAKE_CXX_FLAGS_RELWITHDEBINFO}")
+    target_compile_options(omnitrace-exe PRIVATE ${_FLAGS})
+endif()
+
+install(
+    TARGETS omnitrace-exe
+    DESTINATION ${CMAKE_INSTALL_BINDIR}
+    OPTIONAL)
@@ -0,0 +1,139 @@
+# ------------------------------------------------------------------------------#
+#
+# omnitrace interface library
+#
+# ------------------------------------------------------------------------------#
+
+add_library(omnitrace-interface-library INTERFACE)
+add_library(omnitrace::omnitrace-interface-library ALIAS omnitrace-interface-library)
+
+target_include_directories(
+    omnitrace-interface-library INTERFACE ${CMAKE_CURRENT_SOURCE_DIR}/include
+                                          ${CMAKE_CURRENT_BINARY_DIR}/include)
+target_include_directories(omnitrace-interface-library SYSTEM
+                           INTERFACE ${perfetto_DIR}/sdk)
+
+target_compile_definitions(
+    omnitrace-interface-library
+    INTERFACE OMNITRACE_MAX_THREADS=${OMNITRACE_MAX_THREADS}
+              $<IF:$<BOOL:${OMNITRACE_CUSTOM_DATA_SOURCE}>,CUSTOM_DATA_SOURCE,>)
+
+target_link_libraries(
+    omnitrace-interface-library
+    INTERFACE $<BUILD_INTERFACE:omnitrace::omnitrace-headers>
+              $<BUILD_INTERFACE:omnitrace::omnitrace-threading>
+              $<BUILD_INTERFACE:omnitrace::omnitrace-compile-options>
+              $<BUILD_INTERFACE:omnitrace::omnitrace-hip>
+              $<BUILD_INTERFACE:omnitrace::omnitrace-roctracer>
+              $<BUILD_INTERFACE:omnitrace::omnitrace-rocm-smi>
+              $<BUILD_INTERFACE:omnitrace::omnitrace-mpi>
+              $<BUILD_INTERFACE:omnitrace::omnitrace-ptl>
+              $<BUILD_INTERFACE:timemory::timemory-headers>
+              $<BUILD_INTERFACE:timemory::timemory-gotcha>
+              $<BUILD_INTERFACE:timemory::timemory-cxx-shared>
+              $<IF:$<BOOL:${OMNITRACE_USE_SANITIZER}>,omnitrace::omnitrace-sanitizer,>)
+
+# ------------------------------------------------------------------------------#
+#
+# omnitrace object library
+#
+# ------------------------------------------------------------------------------#
+
+add_library(omnitrace-object-library OBJECT)
+add_library(omnitrace::omnitrace-object-library ALIAS omnitrace-object-library)
+
+configure_file(${CMAKE_CURRENT_SOURCE_DIR}/include/library/defines.hpp.in
+               ${CMAKE_CURRENT_BINARY_DIR}/include/library/defines.hpp @ONLY)
+
+set(library_sources
+    ${CMAKE_CURRENT_LIST_DIR}/src/library.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/config.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/cpu_freq.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/critical_trace.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/kokkosp.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/gpu.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/perfetto.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/ptl.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/sampling.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/state.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/thread_data.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/thread_sampler.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/timemory.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/components/backtrace.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/components/fork_gotcha.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/components/mpi_gotcha.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/components/omnitrace.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/src/library/components/pthread_gotcha.cpp
+    ${perfetto_DIR}/sdk/perfetto.cc)
+
+set(library_headers
+    ${CMAKE_CURRENT_LIST_DIR}/include/library.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/api.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/config.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/common.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/cpu_freq.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/critical_trace.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/debug.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/gpu.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/perfetto.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/ptl.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/sampling.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/state.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/thread_data.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/thread_sampler.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/timemory.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/fwd.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/backtrace.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/fork_gotcha.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/mpi_gotcha.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/omnitrace.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/rocm_smi.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/roctracer.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/roctracer_callbacks.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/include/library/components/pthread_gotcha.hpp
+    ${perfetto_DIR}/sdk/perfetto.h)
+
+target_sources(omnitrace-object-library PRIVATE ${library_sources} ${library_headers})
+
+if(OMNITRACE_USE_ROCTRACER)
+    target_sources(
+        omnitrace-object-library
+        PRIVATE ${CMAKE_CURRENT_LIST_DIR}/src/library/components/roctracer.cpp
+                ${CMAKE_CURRENT_LIST_DIR}/src/library/components/roctracer_callbacks.cpp)
+endif()
+
+if(OMNITRACE_USE_ROCM_SMI)
+    target_sources(omnitrace-object-library
+                   PRIVATE ${CMAKE_CURRENT_LIST_DIR}/src/library/components/rocm_smi.cpp)
+endif()
+
+target_link_libraries(omnitrace-object-library PRIVATE omnitrace-interface-library)
+
+if(OMNITRACE_DYNINST_API_RT)
+    get_filename_component(OMNITRACE_DYNINST_API_RT_DIR "${OMNITRACE_DYNINST_API_RT}"
+                           DIRECTORY)
+endif()
+
+# ------------------------------------------------------------------------------#
+#
+# omnitrace shared library
+#
+# ------------------------------------------------------------------------------#
+
+add_library(omnitrace-library SHARED $<TARGET_OBJECTS:omnitrace-object-library>)
+add_library(omnitrace::omnitrace-library ALIAS omnitrace-library)
+
+target_link_libraries(omnitrace-library PRIVATE omnitrace-interface-library)
+
+set_target_properties(
+    omnitrace-library
+    PROPERTIES OUTPUT_NAME omnitrace
+               VERSION ${PROJECT_VERSION}
+               SOVERSION ${PROJECT_VERSION_MAJOR}
+               INSTALL_RPATH
+               "\$ORIGIN:\$ORIGIN/timemory/libunwind:\$ORIGIN/dyninst-tpls/libs")
+
+install(
+    TARGETS omnitrace-library
+    DESTINATION ${CMAKE_INSTALL_LIBDIR}
+    OPTIONAL)
@@ -50,4 +50,7 @@ extern "C"

    /// stops an instrumentation region
    void omnitrace_pop_trace(const char* name) TIMEMORY_VISIBILITY("default");
+
+    /// used by omnitrace-critical-trace
+    bool omnitrace_init_tooling() TIMEMORY_VISIBILITY("hidden");
 }
@@ -51,9 +51,20 @@ namespace rocm_smi
 void
 setup();

+void
+config();
+
+void
+sample();
+
 void
 shutdown();

+void
+post_process();
+
+void set_state(State);
+
 uint32_t
 device_count();

@@ -77,12 +88,6 @@ struct data
    void sample(uint32_t _dev_id);
    void print(std::ostream& _os) const;

-    template <typename Tp                                                      = nsec_t,
-              std::enable_if_t<!std::is_same_v<std::decay_t<Tp>, nsec_t>, int> = 0>
-    static void poll(std::atomic<State>* _state, Tp&& _interval, promise_t*);
-
-    static void set_state(State);
-    static void poll(std::atomic<State>* _state, nsec_t _interval, promise_t*);
    static void post_process(uint32_t _dev_id);

    uint32_t    m_dev_id    = std::numeric_limits<uint32_t>::max();
@@ -100,7 +105,10 @@ struct data

 private:
    friend void omnitrace::rocm_smi::setup();
+    friend void omnitrace::rocm_smi::config();
+    friend void omnitrace::rocm_smi::sample();
    friend void omnitrace::rocm_smi::shutdown();
+    friend void omnitrace::rocm_smi::post_process();

    static size_t                        device_count;
    static std::set<uint32_t>            device_list;
@@ -111,28 +119,28 @@ private:
    static bool                          shutdown();
 };

-template <
-    typename Tp,
-    std::enable_if_t<!std::is_same_v<std::decay_t<Tp>, std::chrono::nanoseconds>, int>>
-void
-data::poll(std::atomic<State>* _state, Tp&& _interval, promise_t* _prom)
-{
-    poll(_state, std::chrono::duration_cast<nsec_t>(_interval), _prom);
-}
-
-using bundle_t          = std::deque<data>;
-using sampler_instances = thread_data<bundle_t, api::rocm_smi>;
-
 #if !defined(OMNITRACE_USE_ROCM_SMI)
 inline void
 setup()
 {}

+inline void
+config()
+{}
+
+inline void
+sample()
+{}
+
 inline void
 shutdown()
 {}

-inline void data::post_process(uint32_t) {}
+inline void
+post_process()
+{}
+
+inline void set_state(State) {}
 #endif
 }  // namespace rocm_smi
 }  // namespace omnitrace
@@ -54,17 +54,12 @@

 namespace omnitrace
 {
-using hsa_timer_t = hsa_rt_utils::Timer;
-using timestamp_t = hsa_timer_t::timestamp_t;
 using roctracer_bundle_t =
    tim::component_bundle<api::omnitrace, comp::roctracer_data, comp::wall_clock>;
 using roctracer_hsa_bundle_t =
    tim::component_bundle<api::omnitrace, comp::roctracer_data>;
 using roctracer_functions_t = std::vector<std::pair<std::string, std::function<void()>>>;

-std::unique_ptr<hsa_timer_t>&
-get_hsa_timer();
-
 // HSA API callback function
 void
 hsa_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg);
@@ -212,7 +212,7 @@ get_backend();

 // make this visible so omnitrace-avail can call it
 std::string&
-get_perfetto_output_filename() TIMEMORY_VISIBILITY("default");
+get_perfetto_output_filename();

 int64_t
 get_critical_trace_count();
@@ -227,7 +227,7 @@ double&
 get_sampling_delay();

 double&
-get_rocm_smi_freq();
+get_thread_sampling_freq();

 std::string
 get_rocm_smi_devices();
@@ -0,0 +1,44 @@
+// 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
+
+namespace omnitrace
+{
+namespace cpu_freq
+{
+void
+setup();
+
+void
+config();
+
+void
+sample();
+
+void
+shutdown();
+
+void
+post_process();
+}  // namespace cpu_freq
+}  // namespace omnitrace
@@ -26,7 +26,9 @@
 #include "library/defines.hpp"
 #include "library/thread_data.hpp"

-#include <timemory/tpls/cereal/cereal/cereal.hpp>
+#include <timemory/hash/types.hpp>
+#include <timemory/tpls/cereal/cereal.hpp>
+#include <timemory/utility/demangle.hpp>

 #include <cstdint>
 #include <cstdlib>
@@ -43,6 +45,7 @@ enum class Device : short
    NONE = 0,
    CPU,
    GPU,
+    ANY,
 };

 enum class Phase : short
@@ -106,27 +109,44 @@ struct entry
        return _os;
    }
    template <typename Archive>
-    void serialize(Archive& ar, unsigned int);
+    void save(Archive& ar, unsigned int) const;
+
+    template <typename Archive>
+    void load(Archive& ar, unsigned int);
 };

 template <typename Archive>
 void
-entry::serialize(Archive& ar, unsigned int)
+entry::save(Archive& ar, unsigned int) const
 {
    namespace cereal = tim::cereal;
+    std::string _name{};
+    if(hash > 0) _name = tim::get_hash_identifier(hash);
    ar(cereal::make_nvp("priority", priority), cereal::make_nvp("device", device),
       cereal::make_nvp("phase", phase), cereal::make_nvp("depth", depth),
       cereal::make_nvp("tid", tid), cereal::make_nvp("cpu_cid", cpu_cid),
       cereal::make_nvp("gpu_cid", gpu_cid), cereal::make_nvp("parent_cid", parent_cid),
       cereal::make_nvp("begin_ns", begin_ns), cereal::make_nvp("end_ns", end_ns),
-       cereal::make_nvp("hash", hash));
+       cereal::make_nvp("hash", hash), cereal::make_nvp("name", _name),
+       cereal::make_nvp("demangled_name", tim::demangle(_name)));
+}

-    if(get_critical_trace_serialize_names())
-    {
-        std::string _name{};
-        if(hash > 0) _name = tim::demangle(tim::get_hash_identifier(hash));
-        ar(cereal::make_nvp("name", _name));
-    }
+template <typename Archive>
+void
+entry::load(Archive& ar, unsigned int)
+{
+    namespace cereal = tim::cereal;
+    std::string _name{};
+    std::string _demangled_name{};
+    ar(cereal::make_nvp("priority", priority), cereal::make_nvp("device", device),
+       cereal::make_nvp("phase", phase), cereal::make_nvp("depth", depth),
+       cereal::make_nvp("tid", tid), cereal::make_nvp("cpu_cid", cpu_cid),
+       cereal::make_nvp("gpu_cid", gpu_cid), cereal::make_nvp("parent_cid", parent_cid),
+       cereal::make_nvp("begin_ns", begin_ns), cereal::make_nvp("end_ns", end_ns),
+       cereal::make_nvp("hash", hash), cereal::make_nvp("name", _name),
+       cereal::make_nvp("demangled_name", _demangled_name));
+
+    tim::get_hash_ids()->emplace(hash, _name);
 }

 struct call_chain : private std::vector<entry>
@@ -175,12 +195,23 @@ struct call_chain : private std::vector<entry>
    }

    template <Device DevT>
-    void generate_perfetto(std::set<entry>& _used, bool _basic = false) const;
+    void generate_perfetto(std::set<entry>& _used) const;

    template <bool BoolV = true, typename FuncT>
    bool query(FuncT&&) const;
 };

+template <bool BoolV, typename FuncT>
+bool
+call_chain::query(FuncT&& _func) const
+{
+    for(const auto& itr : *this)
+    {
+        if(std::forward<FuncT>(_func)(itr)) return BoolV;
+    }
+    return !BoolV;
+}
+
 using hash_ids = std::unordered_set<std::string>;

 uint64_t
@@ -0,0 +1,218 @@
+// 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"
+
+#include <timemory/api.hpp>
+#include <timemory/backends/dmp.hpp>
+#include <timemory/backends/process.hpp>
+#include <timemory/utility/utility.hpp>
+
+#include <cstdio>
+
+namespace omnitrace
+{
+inline namespace config
+{
+bool
+get_debug();
+
+bool
+get_debug_tid();
+
+bool
+get_debug_pid();
+
+bool
+get_critical_trace_debug();
+}  // namespace config
+}  // namespace omnitrace
+
+#if defined(TIMEMORY_USE_MPI)
+#    define OMNITRACE_PROCESS_IDENTIFIER static_cast<int>(::tim::dmp::rank())
+#elif defined(TIMEMORY_USE_MPI_HEADERS)
+#    define OMNITRACE_PROCESS_IDENTIFIER                                                 \
+        (::tim::dmp::is_initialized()) ? static_cast<int>(::tim::dmp::rank())            \
+                                       : static_cast<int>(::tim::process::get_id())
+#else
+#    define OMNITRACE_PROCESS_IDENTIFIER static_cast<int>(::tim::process::get_id())
+#endif
+
+#define OMNITRACE_CONDITIONAL_PRINT(COND, ...)                                           \
+    if((COND) && ::omnitrace::config::get_debug_tid() &&                                 \
+       ::omnitrace::config::get_debug_pid())                                             \
+    {                                                                                    \
+        fflush(stderr);                                                                  \
+        tim::auto_lock_t _lk{ tim::type_mutex<decltype(std::cerr)>() };                  \
+        fprintf(stderr, "[omnitrace][%i][%li] ", OMNITRACE_PROCESS_IDENTIFIER,           \
+                tim::threading::get_id());                                               \
+        fprintf(stderr, __VA_ARGS__);                                                    \
+        fflush(stderr);                                                                  \
+    }
+
+#define OMNITRACE_CONDITIONAL_BASIC_PRINT(COND, ...)                                     \
+    if((COND) && ::omnitrace::config::get_debug_tid() &&                                 \
+       ::omnitrace::config::get_debug_pid())                                             \
+    {                                                                                    \
+        fflush(stderr);                                                                  \
+        tim::auto_lock_t _lk{ tim::type_mutex<decltype(std::cerr)>() };                  \
+        fprintf(stderr, "[omnitrace] ");                                                 \
+        fprintf(stderr, __VA_ARGS__);                                                    \
+        fflush(stderr);                                                                  \
+    }
+
+#define OMNITRACE_CONDITIONAL_PRINT_F(COND, ...)                                         \
+    if((COND) && ::omnitrace::config::get_debug_tid() &&                                 \
+       ::omnitrace::config::get_debug_pid())                                             \
+    {                                                                                    \
+        fflush(stderr);                                                                  \
+        tim::auto_lock_t _lk{ tim::type_mutex<decltype(std::cerr)>() };                  \
+        fprintf(stderr, "[omnitrace][%i][%li][%s] ", OMNITRACE_PROCESS_IDENTIFIER,       \
+                tim::threading::get_id(), __FUNCTION__);                                 \
+        fprintf(stderr, __VA_ARGS__);                                                    \
+        fflush(stderr);                                                                  \
+    }
+
+#define OMNITRACE_CONDITIONAL_BASIC_PRINT_F(COND, ...)                                   \
+    if((COND) && ::omnitrace::config::get_debug_tid() &&                                 \
+       ::omnitrace::config::get_debug_pid())                                             \
+    {                                                                                    \
+        fflush(stderr);                                                                  \
+        tim::auto_lock_t _lk{ tim::type_mutex<decltype(std::cerr)>() };                  \
+        fprintf(stderr, "[omnitrace][%s] ", __FUNCTION__);                               \
+        fprintf(stderr, __VA_ARGS__);                                                    \
+        fflush(stderr);                                                                  \
+    }
+
+#define OMNITRACE_CONDITIONAL_THROW(COND, ...)                                           \
+    if(COND)                                                                             \
+    {                                                                                    \
+        char _msg_buffer[2048];                                                          \
+        snprintf(_msg_buffer, 2048, "[omnitrace][%i][%li][%s] ",                         \
+                 OMNITRACE_PROCESS_IDENTIFIER, tim::threading::get_id(), __FUNCTION__);  \
+        auto len = strlen(_msg_buffer);                                                  \
+        snprintf(_msg_buffer + len, 2048 - len, __VA_ARGS__);                            \
+        throw std::runtime_error(_msg_buffer);                                           \
+    }
+
+#define OMNITRACE_CONDITIONAL_BASIC_THROW(COND, ...)                                     \
+    if(COND)                                                                             \
+    {                                                                                    \
+        char _msg_buffer[2048];                                                          \
+        snprintf(_msg_buffer, 2048, "[omnitrace][%s] ", __FUNCTION__);                   \
+        auto len = strlen(_msg_buffer);                                                  \
+        snprintf(_msg_buffer + len, 2048 - len, __VA_ARGS__);                            \
+        throw std::runtime_error(_msg_buffer);                                           \
+    }
+
+#define OMNITRACE_STRINGIZE(...) #__VA_ARGS__
+#define OMNITRACE_ESC(...)       __VA_ARGS__
+
+//--------------------------------------------------------------------------------------//
+//
+//  Debug macros
+//
+//--------------------------------------------------------------------------------------//
+#define OMNITRACE_DEBUG(...)                                                             \
+    OMNITRACE_CONDITIONAL_PRINT(::omnitrace::get_debug(), __VA_ARGS__)
+
+#define OMNITRACE_BASIC_DEBUG(...)                                                       \
+    OMNITRACE_CONDITIONAL_BASIC_PRINT(::omnitrace::get_debug_env(), __VA_ARGS__)
+
+#define OMNITRACE_DEBUG_F(...)                                                           \
+    OMNITRACE_CONDITIONAL_PRINT_F(::omnitrace::get_debug(), __VA_ARGS__)
+
+#define OMNITRACE_BASIC_DEBUG_F(...)                                                     \
+    OMNITRACE_CONDITIONAL_BASIC_PRINT_F(::omnitrace::get_debug_env(), __VA_ARGS__)
+
+#define OMNITRACE_CT_DEBUG(...)                                                          \
+    OMNITRACE_CONDITIONAL_PRINT(::omnitrace::get_critical_trace_debug(), __VA_ARGS__)
+
+#define OMNITRACE_CT_DEBUG_F(...)                                                        \
+    OMNITRACE_CONDITIONAL_PRINT_F(::omnitrace::get_critical_trace_debug(), __VA_ARGS__)
+
+//--------------------------------------------------------------------------------------//
+//
+//  Verbose macros
+//
+//--------------------------------------------------------------------------------------//
+#define OMNITRACE_VERBOSE(LEVEL, ...)                                                    \
+    OMNITRACE_CONDITIONAL_PRINT(                                                         \
+        ::omnitrace::get_debug() || ::omnitrace::get_verbose() >= LEVEL, __VA_ARGS__)
+
+#define OMNITRACE_BASIC_VERBOSE(LEVEL, ...)                                              \
+    OMNITRACE_CONDITIONAL_BASIC_PRINT(::omnitrace::get_debug_env() ||                    \
+                                          ::omnitrace::get_verbose_env() >= LEVEL,       \
+                                      __VA_ARGS__)
+
+#define OMNITRACE_VERBOSE_F(LEVEL, ...)                                                  \
+    OMNITRACE_CONDITIONAL_PRINT_F(                                                       \
+        ::omnitrace::get_debug() || ::omnitrace::get_verbose() >= LEVEL, __VA_ARGS__)
+
+#define OMNITRACE_BASIC_VERBOSE_F(LEVEL, ...)                                            \
+    OMNITRACE_CONDITIONAL_BASIC_PRINT_F(::omnitrace::get_debug_env() ||                  \
+                                            ::omnitrace::get_verbose_env() >= LEVEL,     \
+                                        __VA_ARGS__)
+
+//--------------------------------------------------------------------------------------//
+//
+//  Basic print macros (basic means it will not provide PID/RANK or TID) and will not
+//  initialize the settings.
+//
+//--------------------------------------------------------------------------------------//
+
+#define OMNITRACE_BASIC_PRINT(...) OMNITRACE_CONDITIONAL_BASIC_PRINT(true, __VA_ARGS__)
+
+#define OMNITRACE_BASIC_PRINT_F(...) OMNITRACE_CONDITIONAL_BASIC_PRINT(true, __VA_ARGS__)
+
+//--------------------------------------------------------------------------------------//
+//
+//  Print macros. Will provide PID/RANK and TID (will initialize settings)
+//
+//--------------------------------------------------------------------------------------//
+
+#define OMNITRACE_PRINT(...) OMNITRACE_CONDITIONAL_PRINT(true, __VA_ARGS__)
+
+#define OMNITRACE_PRINT_F(...) OMNITRACE_CONDITIONAL_PRINT_F(true, __VA_ARGS__)
+
+//--------------------------------------------------------------------------------------//
+//
+//  Throw macros
+//
+//--------------------------------------------------------------------------------------//
+
+#define OMNITRACE_THROW(...) OMNITRACE_CONDITIONAL_THROW(true, __VA_ARGS__)
+
+#define OMNITRACE_BASIC_THROW(...) OMNITRACE_CONDITIONAL_BASIC_THROW(true, __VA_ARGS__)
+
+#include <string>
+
+namespace std
+{
+inline std::string
+to_string(bool _v)
+{
+    return (_v) ? "true" : "false";
+}
+}  // namespace std
@@ -36,6 +36,8 @@
            perfetto::Category("rocm_smi").SetDescription("Device-level metrics"),       \
            perfetto::Category("sampling")                                               \
                .SetDescription("Metrics derived from sampling"),                        \
+            perfetto::Category("critical-trace")                                         \
+                .SetDescription("Combined critical traces"),                             \
            perfetto::Category("host-critical-trace")                                    \
                .SetDescription("Host-side critical traces"),                            \
            perfetto::Category("device-critical-trace")                                  \
@@ -47,6 +49,8 @@
            perfetto::Category("rocm_smi").SetDescription("Device-level metrics"),       \
            perfetto::Category("sampling")                                               \
                .SetDescription("Metrics derived from sampling"),                        \
+            perfetto::Category("critical-trace")                                         \
+                .SetDescription("Combined critical traces"),                             \
            perfetto::Category("host-critical-trace")                                    \
                .SetDescription("Host-side critical traces"),                            \
            perfetto::Category("device-critical-trace")                                  \
@@ -0,0 +1,103 @@
+// 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/common.hpp"
+#include "library/components/fwd.hpp"
+#include "library/defines.hpp"
+#include "library/state.hpp"
+#include "library/thread_data.hpp"
+
+#include <atomic>
+#include <chrono>
+#include <cstdint>
+#include <functional>
+#include <future>
+#include <memory>
+#include <type_traits>
+#include <vector>
+
+namespace omnitrace
+{
+namespace thread_sampler
+{
+struct instance
+{
+    std::function<void()> setup        = []() {};
+    std::function<void()> shutdown     = []() {};
+    std::function<void()> config       = []() {};
+    std::function<void()> sample       = []() {};
+    std::function<void()> post_process = []() {};
+};
+//
+struct sampler
+{
+    using msec_t    = std::chrono::milliseconds;
+    using usec_t    = std::chrono::microseconds;
+    using nsec_t    = std::chrono::nanoseconds;
+    using promise_t = std::promise<void>;
+    using future_t  = std::future<void>;
+    using state_t   = State;
+
+    using timestamp_t = int64_t;
+
+    template <typename Tp                                                      = nsec_t,
+              std::enable_if_t<!std::is_same_v<std::decay_t<Tp>, nsec_t>, int> = 0>
+    static void poll(std::atomic<state_t>* _state, Tp&& _interval, promise_t*);
+
+    static void setup();
+    static void shutdown();
+    static void post_process();
+    static void set_state(state_t);
+    static void poll(std::atomic<state_t>* _state, nsec_t _interval, promise_t*);
+};
+//
+template <
+    typename Tp,
+    std::enable_if_t<!std::is_same_v<std::decay_t<Tp>, std::chrono::nanoseconds>, int>>
+void
+sampler::poll(std::atomic<state_t>* _state, Tp&& _interval, promise_t* _prom)
+{
+    poll(_state, std::chrono::duration_cast<nsec_t>(_interval), _prom);
+}
+//
+inline void
+setup()
+{
+    sampler::setup();
+}
+
+inline void
+shutdown()
+{
+    sampler::shutdown();
+}
+
+inline void
+post_process()
+{
+    sampler::post_process();
+}
+//
+}  // namespace thread_sampler
+}  // namespace omnitrace
@@ -23,7 +23,6 @@
 #include "library.hpp"
 #include "library/components/fork_gotcha.hpp"
 #include "library/components/mpi_gotcha.hpp"
-#include "library/components/rocm_smi.hpp"
 #include "library/config.hpp"
 #include "library/critical_trace.hpp"
 #include "library/debug.hpp"
@@ -31,6 +30,7 @@
 #include "library/gpu.hpp"
 #include "library/sampling.hpp"
 #include "library/thread_data.hpp"
+#include "library/thread_sampler.hpp"
 #include "library/timemory.hpp"
 #include "timemory/mpl/type_traits.hpp"

@@ -112,9 +112,6 @@ get_cpu_cid_parents()

 using Device = critical_trace::Device;
 using Phase  = critical_trace::Phase;
-
-bool
-omnitrace_init_tooling();
 }  // namespace

 //======================================================================================//
@@ -133,9 +130,8 @@ omnitrace_push_trace(const char* name)
    if(get_state() != State::Active && !omnitrace_init_tooling())
    {
        static auto _debug = get_debug_env();
-        OMNITRACE_CONDITIONAL_BASIC_PRINT(
-            _debug, "[%s] %s :: not active and perfetto not initialized\n", __FUNCTION__,
-            name);
+        OMNITRACE_CONDITIONAL_BASIC_PRINT_F(_debug, "%s :: not active. state = %s\n",
+                                            name, std::to_string(get_state()).c_str());
        return;
    }

@@ -342,9 +338,7 @@ omnitrace_init_library()

 //======================================================================================//

-namespace
-{
-bool
+extern "C" bool
 omnitrace_init_tooling()
 {
    static bool _once = false;
@@ -366,22 +360,10 @@ omnitrace_init_tooling()
    OMNITRACE_DEBUG("[%s]\n", __FUNCTION__);

    auto _dtor = scope::destructor{ []() {
-        // if roctracer is not enabled, we cannot rely on OnLoad(...) via HSA tools
-        // to setup rocm-smi
-        if constexpr(!trait::is_available<comp::roctracer>::value)
-        {
-            try
-            {
-                rocm_smi::setup();
-            } catch(std::exception& _e)
-            {
-                OMNITRACE_PRINT("Exception: %s\n", _e.what());
-            }
-        }
-
        if(get_use_sampling())
        {
            pthread_gotcha::enable_sampling_on_child_threads() = false;
+            thread_sampler::setup();
            sampling::setup();
            pthread_gotcha::enable_sampling_on_child_threads() = true;
            sampling::unblock_signals();
@@ -500,7 +482,8 @@ omnitrace_init_tooling()

    static auto _push_perfetto = [](const char* name) {
        _thread_init();
-        TRACE_EVENT_BEGIN("host", perfetto::StaticString(name));
+        uint64_t _ts = comp::wall_clock::record();
+        TRACE_EVENT_BEGIN("host", perfetto::StaticString(name), _ts);
    };

    static auto _pop_timemory = [](const char* name) {
@@ -517,7 +500,10 @@ omnitrace_init_tooling()
        _data.bundles.pop_back();
    };

-    static auto _pop_perfetto = [](const char*) { TRACE_EVENT_END("host"); };
+    static auto _pop_perfetto = [](const char*) {
+        uint64_t _ts = comp::wall_clock::record();
+        TRACE_EVENT_END("host", _ts);
+    };

    if(get_use_perfetto() && get_use_timemory())
    {
@@ -572,7 +558,6 @@ omnitrace_init_tooling()

    return true;
 }
-}  // namespace

 //======================================================================================//

@@ -590,6 +575,10 @@ omnitrace_init(const char* _mode, bool _is_binary_rewrite, const char* _argv0)
    tim::set_env("OMNITRACE_MODE", _mode, 0);
    config::is_binary_rewrite() = _is_binary_rewrite;

+    // default to KokkosP enabled when sampling, otherwise default to off
+    tim::set_env("OMNITRACE_USE_KOKKOSP", (get_mode() == Mode::Sampling) ? "ON" : "OFF",
+                 0);
+
    if(!_set_mpi_called)
    {
        _start_gotcha_callback = []() { get_gotcha_bundle()->start(); };
@@ -606,7 +595,12 @@ extern "C" void
 omnitrace_finalize(void)
 {
    // return if not active
-    if(get_state() != State::Active) return;
+    if(get_state() != State::Active)
+    {
+        OMNITRACE_DEBUG_F("State = %s. Finalization skipped\n",
+                          std::to_string(get_state()).c_str());
+        return;
+    }

    pthread_gotcha::enable_sampling_on_child_threads() = false;

@@ -643,12 +637,7 @@ omnitrace_finalize(void)
    }

    pthread_gotcha::shutdown();
-
-    if(get_use_rocm_smi())
-    {
-        OMNITRACE_DEBUG("[%s] Shutting down rocm-smi sampling...\n", __FUNCTION__);
-        rocm_smi::shutdown();
-    }
+    thread_sampler::shutdown();

    OMNITRACE_DEBUG("[%s] Shutting down roctracer...\n", __FUNCTION__);
    // ensure that threads running roctracer callbacks shutdown
@@ -746,20 +735,7 @@ omnitrace_finalize(void)
        }
    }

-    // ensure that all the MT instances are flushed
-    if(get_use_rocm_smi())
-    {
-        size_t _ndevice = gpu::device_count();
-        OMNITRACE_CONDITIONAL_PRINT(
-            get_debug() || get_verbose() > 0,
-            "[%s] Post-processing rocm-smi samples for %zu devices...\n", __FUNCTION__,
-            _ndevice);
-        for(size_t i = 0; i < _ndevice; ++i)
-        {
-            rocm_smi::data::post_process(i);
-            rocm_smi::sampler_instances::instances().at(i).reset();
-        }
-    }
+    thread_sampler::post_process();

    if(get_use_critical_trace())
    {
@@ -60,7 +60,9 @@ namespace omnitrace
 namespace rocm_smi
 {
 using tim::type_mutex;
-using auto_lock_t = tim::auto_lock_t;
+using auto_lock_t       = tim::auto_lock_t;
+using bundle_t          = std::deque<data>;
+using sampler_instances = thread_data<bundle_t, api::rocm_smi>;

 namespace
 {
@@ -125,17 +127,51 @@ data::print(std::ostream& _os) const

 namespace
 {
-struct cpu_freq
-{};
-using freq_pair_t                                     = std::pair<size_t, double>;
-std::vector<std::vector<freq_pair_t>> cpu_frequencies = {};
+std::vector<std::unique_ptr<bundle_t>*> _bundle_data{};
+}

-struct cpu_mem
-{};
-using cpu_mem_usage_pair_t                      = std::pair<size_t, int64_t>;
-std::vector<cpu_mem_usage_pair_t> cpu_mem_usage = {};
-}  // namespace
+void
+config()
+{
+    _bundle_data.resize(data::device_count, nullptr);
+    for(size_t i = 0; i < data::device_count; ++i)
+    {
+        if(data::device_list.count(i) > 0)
+        {
+            _bundle_data.at(i) = &sampler_instances::instances().at(i);
+            if(!*_bundle_data.at(i)) *_bundle_data.at(i) = std::make_unique<bundle_t>();
+        }
+    }

+    data::get_initial().resize(data::device_count);
+    for(auto itr : data::device_list)
+        data::get_initial().at(itr).sample(itr);
+}
+
+void
+sample()
+{
+    if(get_rocm_smi_state() != State::Active) return;
+
+    for(auto itr : data::device_list)
+    {
+        OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug(),
+                                          "Polling rocm-smi for device %u...\n", itr);
+        auto& _data = *_bundle_data.at(itr);
+        if(!_data) continue;
+        _data->emplace_back(data{ itr });
+        OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug(), "    %s\n",
+                                          TIMEMORY_JOIN("", _data->back()).c_str());
+    }
+}
+
+void
+set_state(State _v)
+{
+    get_rocm_smi_state().store(_v);
+}
+
+/*
 void
 data::poll(std::atomic<State>* _state, nsec_t _interval, promise_t* _ready)
 {
@@ -155,71 +191,6 @@ data::poll(std::atomic<State>* _state, nsec_t _interval, promise_t* _ready)
        }
    }

-    auto                _ncpu = threading::affinity::hw_concurrency();
-    std::vector<size_t> _cpu_mhz_pos{};
-    std::ifstream       _ifs{ "/proc/cpuinfo" };
-    if(_ifs)
-    {
-        for(size_t i = 0; i < _ncpu; ++i)
-        {
-            short       _n = 0;
-            std::string _st{};
-            while(_ifs && _ifs.good())
-            {
-                std::string _s{};
-                _ifs >> _s;
-                if(!_ifs.good() || !_ifs) break;
-
-                if(_s == "cpu" || _s == "MHz" || _s == ":")
-                {
-                    ++_n;
-                    _st += _s + " ";
-                }
-                else
-                {
-                    _n  = 0;
-                    _st = {};
-                }
-
-                if(_n == 3)
-                {
-                    size_t _pos = _ifs.tellg();
-                    _cpu_mhz_pos.emplace_back(_pos + 1);
-                    _ifs >> _s;
-                    if(!_ifs.good() || !_ifs) break;
-                    OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug() || get_verbose() > 1,
-                                                      "[%zu] %s %s (pos = %zu)\n", i,
-                                                      _st.c_str(), _s.c_str(), _pos + 1);
-                    break;
-                }
-            }
-        }
-    }
-
-    cpu_frequencies.resize(_ncpu);
-
-    _ifs                = std::ifstream{ "/proc/cpuinfo", std::ifstream::binary };
-    auto _read_cpu_freq = [&_cpu_mhz_pos, &_ifs](size_t _idx) {
-        double _freq = 0;
-        _ifs.seekg(_cpu_mhz_pos.at(_idx), _ifs.beg);
-        _ifs >> _freq;
-        return _freq;
-    };
-
-    auto _read_cpu_mem_usage = []() {
-        cpu_mem_usage.emplace_back(tim::get_clock_real_now<size_t, std::nano>(),
-                                   tim::get_page_rss());
-    };
-
-    auto _read_cpu_freqs = [&_read_cpu_freq, _ncpu]() {
-        auto _ts = tim::get_clock_real_now<size_t, std::nano>();
-        for(size_t i = 0; i < _ncpu; ++i)
-        {
-            auto _freq = _read_cpu_freq(i);
-            cpu_frequencies.at(i).emplace_back(_ts, _freq);
-        }
-    };
-
    OMNITRACE_CONDITIONAL_BASIC_PRINT(
        get_verbose() > 0 || get_debug(),
        "Polling rocm-smi for %zu device(s) at an interval of %f seconds...\n",
@@ -235,8 +206,6 @@ data::poll(std::atomic<State>* _state, nsec_t _interval, promise_t* _ready)
    {
        std::this_thread::sleep_until(_now);
        if(_state->load() != State::Active) continue;
-        _read_cpu_mem_usage();
-        _read_cpu_freqs();
        for(auto itr : device_list)
        {
            OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug(),
@@ -254,6 +223,7 @@ data::poll(std::atomic<State>* _state, nsec_t _interval, promise_t* _ready)

    if(polling_finished) polling_finished->set_value();
 }
+*/

 std::vector<data>&
 data::get_initial()
@@ -262,73 +232,20 @@ data::get_initial()
    return _v;
 }

-std::unique_ptr<std::thread>&
-data::get_thread()
-{
-    static std::unique_ptr<std::thread> _v;
-    return _v;
-}
-
-void
-data::set_state(State _state)
-{
-    get_rocm_smi_state().store(_state);
-}
-
 bool
 data::setup()
 {
-    perfetto_counter_track<cpu_freq>::init();
-    perfetto_counter_track<cpu_mem>::init();
    perfetto_counter_track<data>::init();
-
-    // shutdown if already running
-    shutdown();
-
-    OMNITRACE_DEBUG("Configuring rocm-smi...\n");
-
-    auto     _freq      = get_rocm_smi_freq();
-    uint64_t _msec_freq = (1.0 / _freq) * 1.0e3;
-
-    promise_t _prom{};
-    auto      _fut   = _prom.get_future();
-    polling_finished = std::make_unique<promise_t>();
-
    set_state(State::PreInit);
-    get_thread() = std::make_unique<std::thread>(
-        &data::poll<msec_t>, &get_rocm_smi_state(), msec_t{ _msec_freq }, &_prom);
-
-    _fut.wait();
    return true;
 }

 bool
 data::shutdown()
 {
-    auto& _thread = get_thread();
-    if(_thread)
-    {
-        OMNITRACE_DEBUG("Shutting down rocm-smi...\n");
-        set_state(State::Finalized);
-        if(polling_finished)
-        {
-            auto     _fut  = polling_finished->get_future();
-            uint64_t _freq = (1.0 / get_rocm_smi_freq()) * 1.0e3;
-            _fut.wait_for(msec_t{ 5 * _freq });
-            _thread->join();
-        }
-        else
-        {
-            uint64_t _freq = (1.0 / get_rocm_smi_freq()) * 1.0e3;
-            std::this_thread::sleep_for(msec_t{ 5 * _freq });
-            pthread_cancel(_thread->native_handle());
-            _thread->detach();
-        }
-        _thread          = std::unique_ptr<std::thread>{};
-        polling_finished = std::unique_ptr<promise_t>{};
-        return true;
-    }
-    return false;
+    OMNITRACE_DEBUG("Shutting down rocm-smi...\n");
+    set_state(State::Finalized);
+    return true;
 }

 #define GPU_METRIC(COMPONENT, ...)                                                       \
@@ -355,50 +272,6 @@ data::post_process(uint32_t _dev_id)
    using bundle_t = tim::lightweight_tuple<sampling_gpu_busy, sampling_gpu_temp,
                                            sampling_gpu_power, sampling_gpu_memory>;

-    auto _process_frequencies = [](size_t _idx) {
-        using counter_track = perfetto_counter_track<cpu_freq>;
-        if(!counter_track::exists(_idx))
-        {
-            auto _devname = TIMEMORY_JOIN("", "[CPU ", _idx, "] ");
-            auto addendum = [&](const char* _v) { return _devname + std::string{ _v }; };
-            counter_track::emplace(_idx, addendum("Frequency (S)"), "MHz");
-        }
-
-        for(auto& itr : cpu_frequencies.at(_idx))
-        {
-            uint64_t _ts   = itr.first;
-            double   _freq = itr.second;
-            TRACE_COUNTER("sampling", counter_track::at(_idx, 0), _ts, _freq);
-        }
-    };
-
-    auto _process_cpu_mem_usage = []() {
-        using counter_track = perfetto_counter_track<cpu_mem>;
-        if(!counter_track::exists(0))
-        {
-            auto _devname = TIMEMORY_JOIN("", "[CPU] ");
-            auto addendum = [&](const char* _v) { return _devname + std::string{ _v }; };
-            counter_track::emplace(0, addendum("Memory Usage (S)"), "MB");
-        }
-
-        for(auto& itr : cpu_mem_usage)
-        {
-            uint64_t _ts        = itr.first;
-            double   _mem_usage = itr.second;
-            TRACE_COUNTER("sampling", counter_track::at(0, 0), _ts,
-                          _mem_usage / units::megabyte);
-        }
-    };
-
-    static bool _once = false;
-    if(!_once)
-    {
-        _once = true;
-        _process_cpu_mem_usage();
-        for(size_t i = 0; i < cpu_frequencies.size(); ++i)
-            _process_frequencies(i);
-    }
-
    if(device_count < _dev_id) return;

    auto& _rocm_smi_v = sampler_instances::instances().at(_dev_id);
@@ -474,6 +347,7 @@ setup()

    if(is_initialized() || !get_use_rocm_smi()) return;

+    auto _enable_samp = pthread_gotcha::enable_sampling_on_child_threads();
    pthread_gotcha::enable_sampling_on_child_threads() = false;

    // assign the data value to determined by rocm-smi
@@ -513,8 +387,7 @@ setup()

    data::setup();

-    pthread_gotcha::enable_sampling_on_child_threads() = true;
-    omnitrace::rocm_smi::data::set_state(State::Active);
+    pthread_gotcha::enable_sampling_on_child_threads() = _enable_samp;
 }

 void
@@ -532,6 +405,13 @@ shutdown()
    is_initialized() = false;
 }

+void
+post_process()
+{
+    for(auto itr : data::device_list)
+        data::post_process(itr);
+}
+
 uint32_t
 device_count()
 {
@@ -234,9 +234,6 @@ extern "C"
                                                  "[%s] setting up HSA...\n",
                                                  __FUNCTION__);

-                get_hsa_timer() =
-                    std::make_unique<hsa_timer_t>(table->core_->hsa_system_get_info_fn);
-
                // const char* output_prefix = getenv("ROCP_OUTPUT_DIR");
                const char* output_prefix = nullptr;

@@ -314,7 +311,7 @@ extern "C"
        comp::roctracer::add_setup("hsa", std::move(_setup));
        comp::roctracer::add_shutdown("hsa", std::move(_shutdown));

-        rocm_smi::setup();
+        rocm_smi::set_state(State::Active);
        comp::roctracer::setup();

        pthread_gotcha::enable_sampling_on_child_threads() = true;
@@ -325,7 +322,7 @@ extern "C"
    void OnUnload()
    {
        OMNITRACE_DEBUG("[%s]\n", __FUNCTION__);
-        rocm_smi::shutdown();
+        rocm_smi::set_state(State::Finalized);
        comp::roctracer::shutdown();
    }
 }
@@ -82,13 +82,6 @@ get_hip_activity_callbacks(int64_t _tid = threading::get_id())
    return _v.at(_tid);
 }

-std::unique_ptr<hsa_timer_t>&
-get_hsa_timer()
-{
-    static auto _v = std::unique_ptr<hsa_timer_t>{};
-    return _v;
-}
-
 using hip_activity_mutex_t = std::decay_t<decltype(get_hip_activity_callbacks())>;
 using key_data_mutex_t     = std::decay_t<decltype(get_roctracer_key_data())>;
 using hip_data_mutex_t     = std::decay_t<decltype(get_roctracer_hip_data())>;
@@ -114,17 +107,14 @@ hsa_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void*
                    roctracer_op_string(domain, cid, 0), cid, data->correlation_id,
                    (data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit");

-    static thread_local timestamp_t begin_timestamp = 0;
-    static auto&                    timer           = get_hsa_timer();
-    static auto                     _scope          = []() {
+    static thread_local int64_t begin_timestamp = 0;
+    static auto                 _scope          = []() {
        auto _v = scope::config{};
        if(get_roctracer_timeline_profile()) _v += scope::timeline{};
        if(get_roctracer_flat_profile()) _v += scope::flat{};
        return _v;
    }();

-    if(!timer) return;
-
    switch(cid)
    {
        case HSA_API_ID_hsa_init:
@@ -172,22 +162,22 @@ hsa_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void*
        {
            if(data->phase == ACTIVITY_API_PHASE_ENTER)
            {
-                begin_timestamp = timer->timestamp_fn_ns();
+                begin_timestamp = comp::wall_clock::record();
            }
            else
            {
-                const auto*       _name         = roctracer_op_string(domain, cid, 0);
-                const timestamp_t end_timestamp = (cid == HSA_API_ID_hsa_shut_down)
-                                                      ? begin_timestamp
-                                                      : timer->timestamp_fn_ns();
+                const auto* _name         = roctracer_op_string(domain, cid, 0);
+                const auto  end_timestamp = (cid == HSA_API_ID_hsa_shut_down)
+                                                ? begin_timestamp
+                                                : comp::wall_clock::record();

                if(begin_timestamp > end_timestamp) return;

                if(get_use_perfetto())
                {
                    TRACE_EVENT_BEGIN("device", perfetto::StaticString{ _name },
-                                      begin_timestamp);
-                    TRACE_EVENT_END("device", end_timestamp);
+                                      static_cast<uint64_t>(begin_timestamp));
+                    TRACE_EVENT_END("device", static_cast<uint64_t>(end_timestamp));
                }

                if(get_use_timemory())
@@ -247,8 +237,9 @@ hsa_activity_callback(uint32_t op, activity_record_t* record, void* arg)
    auto _func = [_begin_ns, _end_ns, _name]() {
        if(get_use_perfetto())
        {
-            TRACE_EVENT_BEGIN("device", perfetto::StaticString{ *_name }, _begin_ns);
-            TRACE_EVENT_END("device", _end_ns);
+            TRACE_EVENT_BEGIN("device", perfetto::StaticString{ *_name },
+                              static_cast<uint64_t>(_begin_ns));
+            TRACE_EVENT_END("device", static_cast<uint64_t>(_end_ns));
        }
        if(get_use_timemory())
        {
@@ -312,6 +303,8 @@ hip_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void*
        default: break;
    }

+    int64_t _ts = comp::wall_clock::record();
+
    if(data->phase == ACTIVITY_API_PHASE_ENTER)
    {
        switch(cid)
@@ -349,21 +342,13 @@ hip_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void*
                }
                break;
            }
-            default:
-            {
-                if(get_use_perfetto() || get_use_timemory() || get_use_rocm_smi())
-                {
-                    tim::auto_lock_t _lk{ tim::type_mutex<key_data_mutex_t>() };
-                    get_roctracer_key_data().emplace(data->correlation_id, op_name);
-                    get_roctracer_tid_data().emplace(data->correlation_id,
-                                                     threading::get_id());
-                }
-            }
+            default: break;
        }

        if(get_use_perfetto())
        {
-            TRACE_EVENT_BEGIN("device", perfetto::StaticString{ op_name });
+            TRACE_EVENT_BEGIN("device", perfetto::StaticString{ op_name },
+                              static_cast<uint64_t>(_ts));
        }
        if(get_use_timemory())
        {
@@ -388,8 +373,7 @@ hip_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void*
            auto     _parent_cid = (get_cpu_cid_stack()->empty())
                                       ? get_cpu_cid_stack(0)->back()
                                       : get_cpu_cid_stack()->back();
-            int64_t  _ts         = comp::wall_clock::record();
-            add_critical_trace<Device::GPU, Phase::BEGIN>(
+            add_critical_trace<Device::CPU, Phase::BEGIN>(
                threading::get_id(), _cid, data->correlation_id, _parent_cid, _ts, 0,
                critical_trace::add_hash_id(op_name), _depth);
            tim::auto_lock_t _lk{ tim::type_mutex<cid_data_mutex_t>() };
@@ -405,7 +389,7 @@ hip_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void*

        if(get_use_perfetto())
        {
-            TRACE_EVENT_END("device");
+            TRACE_EVENT_END("device", static_cast<uint64_t>(_ts));
        }
        if(get_use_timemory())
        {
@@ -438,8 +422,7 @@ hip_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void*
                std::tie(_cid, _parent_cid, _depth) =
                    get_roctracer_cid_data().at(data->correlation_id);
            }
-            int64_t _ts = comp::wall_clock::record();
-            add_critical_trace<Device::GPU, Phase::END>(
+            add_critical_trace<Device::CPU, Phase::END>(
                threading::get_id(), _cid, data->correlation_id, _parent_cid, _ts, _ts,
                critical_trace::add_hash_id(op_name), _depth);
        }
@@ -471,20 +454,50 @@ hip_activity_callback(const char* begin, const char* end, void*)

    OMNITRACE_DEBUG("Activity records:\n");

+    auto&& _advance_record = [&record]() {
+        ROCTRACER_CALL(roctracer_next_record(record, &record));
+    };
+
    while(record < end_record)
    {
+        // make sure every iteration advances regardless of where return point happens
+        scope::destructor _next_dtor{ _advance_record };
+
        const char* op_name =
            roctracer_op_string(record->domain, record->correlation_id, 0);
        if(op_name == nullptr) op_name = hip_api_name(record->correlation_id);

+        switch(record->kind)
+        {
+            case HIP_API_ID_hipLaunchKernel:
+            case HIP_API_ID_hipLaunchCooperativeKernel:
+            case HIP_API_ID_hipModuleLaunchKernel: break;
+            case HIP_API_ID_hipGetLastError: continue;
+            default:
+            {
+                if(op_name != nullptr && strcmp(op_name, "unknown") != 0 &&
+                   strcmp(op_name, "InternalMarker") != 0)
+                {
+                    OMNITRACE_BASIC_VERBOSE_F(2, "[%s] ignoring callback for %s\n",
+                                              __FUNCTION__, op_name);
+                    continue;
+                }
+                break;
+            }
+        }
+
+        auto _dev_id = record->device_id;
+        auto _thr_id = record->thread_id;
+        auto _prc_id = record->process_id;
+        auto _que_id = record->queue_id;
+
        if(op_name != nullptr)
        {
-            OMNITRACE_DEBUG("\t%-30s\tcorrelation_id(%6lu) time_ns(%12lu:%12lu) "
-                            "delta_ns(%12lu) device_id(%d) "
-                            "stream_id(%lu)\n",
-                            op_name, record->correlation_id, record->begin_ns,
-                            record->end_ns, (record->end_ns - record->begin_ns),
-                            record->device_id, record->queue_id);
+            OMNITRACE_DEBUG(
+                "\t%-30s\tcorrelation_id(%6lu) time_ns(%12lu:%12lu) "
+                "delta_ns(%12lu) device_id(%d) stream_id(%lu) proc_id(%u) thr_id(%u)\n",
+                op_name, record->correlation_id, record->begin_ns, record->end_ns,
+                (record->end_ns - record->begin_ns), _dev_id, _que_id, _prc_id, _thr_id);
        }

        auto        _begin_ns = record->begin_ns;
@@ -543,8 +556,8 @@ hip_activity_callback(const char* begin, const char* end, void*)
                    TRACE_EVENT_BEGIN(
                        "device",
                        perfetto::StaticString{ _kernel_names.at(_name).c_str() },
-                        _begin_ns);
-                    TRACE_EVENT_END("device", _end_ns);
+                        static_cast<uint64_t>(_begin_ns));
+                    TRACE_EVENT_END("device", static_cast<uint64_t>(_end_ns));
                }
                if(get_use_timemory())
                {
@@ -589,8 +602,6 @@ hip_activity_callback(const char* begin, const char* end, void*)
            tim::auto_lock_t _lk{ get_hip_activity_mutex(_tid) };
            _async_ops->emplace_back(std::move(_func));
        }
-
-        ROCTRACER_CALL(roctracer_next_record(record, &record));
    }
 }

@@ -127,11 +127,6 @@ configure_settings()
    OMNITRACE_CONFIG_SETTING(std::string, "OMNITRACE_ROCM_SMI_DEVICES",
                             "Devices to query when OMNITRACE_USE_ROCM_SMI=ON", "all",
                             "backend", "rocm-smi");
-
-    OMNITRACE_CONFIG_SETTING(
-        double, "OMNITRACE_ROCM_SMI_FREQ",
-        "Number of rocm-smi samples per second when OMNITTRACE_USE_ROCM_SMI=ON", 5.0,
-        "backend", "rocm-smi");
 #endif

    OMNITRACE_CONFIG_SETTING(bool, "OMNITRACE_USE_SAMPLING",
@@ -144,7 +139,7 @@ configure_settings()
        "io");

    OMNITRACE_CONFIG_SETTING(bool, "OMNITRACE_USE_KOKKOSP",
-                             "Enable support for Kokkos Tools", true, "kokkos");
+                             "Enable support for Kokkos Tools", false, "kokkos");

    OMNITRACE_CONFIG_SETTING(size_t, "OMNITRACE_INSTRUMENTATION_INTERVAL",
                             "Instrumentation only takes measurements once every N "
@@ -335,16 +330,6 @@ configure_settings()
    if(tim::mpi::is_initialized()) settings::default_process_suffix() = tim::mpi::rank();
 #endif
    OMNITRACE_CONDITIONAL_BASIC_PRINT(get_verbose_env() > 0, "configuration complete\n");
-
-#if defined(OMNITRACE_USE_ROCM_SMI)
-    auto _rsmi_freq = _config->find("OMNITRACE_ROCM_SMI_FREQ");
-    if(_rsmi_freq != _config->end())
-    {
-        double& _rsmi_freq_v =
-            static_cast<tim::tsettings<double>&>(*_rsmi_freq->second).get();
-        if(_rsmi_freq_v > 1000) _rsmi_freq_v = 1000.;
-    }
-#endif
 }

 void
@@ -846,17 +831,21 @@ get_critical_trace_count()
 }

 double&
-get_rocm_smi_freq()
+get_thread_sampling_freq()
 {
-    static auto _v = get_config()->find("OMNITRACE_ROCM_SMI_FREQ");
-    return static_cast<tim::tsettings<double>&>(*_v->second).get();
+    static auto _v = std::min<double>(get_sampling_freq(), 1000.0);
+    return _v;
 }

 std::string
 get_rocm_smi_devices()
 {
+#if defined(OMNITRACE_USE_ROCM_SMI)
    static auto _v = get_config()->find("OMNITRACE_ROCM_SMI_DEVICES");
    return static_cast<tim::tsettings<std::string>&>(*_v->second).get();
+#else
+    return std::string{};
+#endif
 }

 bool
@@ -0,0 +1,184 @@
+// 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/cpu_freq.hpp"
+#include "library/common.hpp"
+#include "library/config.hpp"
+#include "library/debug.hpp"
+#include "library/defines.hpp"
+#include "library/perfetto.hpp"
+#include "library/timemory.hpp"
+
+#include <cstdlib>
+#include <utility>
+#include <vector>
+
+namespace omnitrace
+{
+namespace cpu_freq
+{
+namespace
+{
+struct cpu_freq
+{};
+using freq_pair_t                                    = std::pair<size_t, double>;
+std::vector<std::deque<freq_pair_t>> cpu_frequencies = {};
+
+struct cpu_mem
+{};
+using cpu_mem_usage_pair_t                     = std::pair<size_t, int64_t>;
+std::deque<cpu_mem_usage_pair_t> cpu_mem_usage = {};
+
+int64_t                        ncpu        = threading::affinity::hw_concurrency();
+std::unique_ptr<std::ifstream> ifs         = {};
+std::vector<size_t>            cpu_mhz_pos = {};
+
+}  // namespace
+
+void
+setup()
+{
+    perfetto_counter_track<cpu_freq>::init();
+    perfetto_counter_track<cpu_mem>::init();
+}
+
+void
+config()
+{
+    auto                _ncpu = threading::affinity::hw_concurrency();
+    std::vector<size_t> _cpu_mhz_pos{};
+    std::ifstream       _ifs{ "/proc/cpuinfo" };
+    if(_ifs)
+    {
+        for(size_t i = 0; i < _ncpu; ++i)
+        {
+            short       _n = 0;
+            std::string _st{};
+            while(_ifs && _ifs.good())
+            {
+                std::string _s{};
+                _ifs >> _s;
+                if(!_ifs.good() || !_ifs) break;
+
+                if(_s == "cpu" || _s == "MHz" || _s == ":")
+                {
+                    ++_n;
+                    _st += _s + " ";
+                }
+                else
+                {
+                    _n  = 0;
+                    _st = {};
+                }
+
+                if(_n == 3)
+                {
+                    size_t _pos = _ifs.tellg();
+                    _cpu_mhz_pos.emplace_back(_pos + 1);
+                    _ifs >> _s;
+                    if(!_ifs.good() || !_ifs) break;
+                    OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug() || get_verbose() > 1,
+                                                      "[%zu] %s %s (pos = %zu)\n", i,
+                                                      _st.c_str(), _s.c_str(), _pos + 1);
+                    break;
+                }
+            }
+        }
+    }
+
+    _ifs.close();
+
+    cpu_frequencies.resize(_ncpu);
+    cpu_mhz_pos = _cpu_mhz_pos;
+    ifs         = std::make_unique<std::ifstream>("/proc/cpuinfo", std::ifstream::binary);
+}
+
+void
+sample()
+{
+    cpu_mem_usage.emplace_back(tim::get_clock_real_now<size_t, std::nano>(),
+                               tim::get_page_rss());
+
+    if(!ifs) return;
+
+    auto _read_cpu_freq = [](size_t _idx) {
+        double _freq = 0;
+        ifs->seekg(cpu_mhz_pos.at(_idx), ifs->beg);
+        (*ifs) >> _freq;
+        return _freq;
+    };
+
+    auto _ts = tim::get_clock_real_now<size_t, std::nano>();
+    for(int64_t i = 0; i < ncpu; ++i)
+        cpu_frequencies.at(i).emplace_back(_ts, _read_cpu_freq(i));
+}
+
+void
+shutdown()
+{}
+
+void
+post_process()
+{
+    OMNITRACE_PRINT("Post-processing %zu cpu freqs and %zu memory usage entries\n",
+                    cpu_frequencies.size(), cpu_mem_usage.size());
+    auto _process_frequencies = [](size_t _idx) {
+        using counter_track = perfetto_counter_track<cpu_freq>;
+        if(!counter_track::exists(_idx))
+        {
+            auto _devname = TIMEMORY_JOIN("", "[CPU ", _idx, "] ");
+            auto addendum = [&](const char* _v) { return _devname + std::string{ _v }; };
+            counter_track::emplace(_idx, addendum("Frequency (S)"), "MHz");
+        }
+
+        for(auto& itr : cpu_frequencies.at(_idx))
+        {
+            uint64_t _ts   = itr.first;
+            double   _freq = itr.second;
+            TRACE_COUNTER("sampling", counter_track::at(_idx, 0), _ts, _freq);
+        }
+    };
+
+    auto _process_cpu_mem_usage = []() {
+        using counter_track = perfetto_counter_track<cpu_mem>;
+        if(!counter_track::exists(0))
+        {
+            auto _devname = TIMEMORY_JOIN("", "[CPU] ");
+            auto addendum = [&](const char* _v) { return _devname + std::string{ _v }; };
+            counter_track::emplace(0, addendum("Memory Usage (S)"), "MB");
+        }
+
+        for(auto& itr : cpu_mem_usage)
+        {
+            uint64_t _ts        = itr.first;
+            double   _mem_usage = itr.second;
+            TRACE_COUNTER("sampling", counter_track::at(0, 0), _ts,
+                          _mem_usage / units::megabyte);
+        }
+    };
+
+    _process_cpu_mem_usage();
+    for(size_t i = 0; i < cpu_frequencies.size(); ++i)
+        _process_frequencies(i);
+}
+}  // namespace cpu_freq
+}  // namespace omnitrace
@@ -66,26 +66,6 @@ update_critical_path(call_chain _chain, int64_t _tid);
 void
 compute_critical_trace();

-void
-find_children(PTL::ThreadPool& _tp, call_graph_t& _graph, const call_chain& _chain);
-
-void
-find_sequences(PTL::ThreadPool& _tp, call_graph_t& _graph,
-               std::vector<call_chain>& _chain);
-
-void
-find_sequences(PTL::ThreadPool& _tp, call_graph_t& _graph, call_graph_itr_t _root,
-               std::vector<call_chain>& _chain);
-
-template <typename ArchiveT, typename T, typename AllocatorT>
-void
-serialize_graph(ArchiveT& ar, const tim::graph<T, AllocatorT>& _graph);
-
-template <typename ArchiveT, typename T, typename AllocatorT>
-void
-serialize_subgraph(ArchiveT& ar, const tim::graph<T, AllocatorT>& _graph,
-                   typename tim::graph<T, AllocatorT>::iterator _root);
-
 void
 copy_hash_ids()
 {
@@ -417,93 +397,85 @@ call_chain::get_top_chains()
    return _v;
 }

-template <bool BoolV, typename FuncT>
-bool
-call_chain::query(FuncT&& _func) const
+template <>
+void
+call_chain::generate_perfetto<Device::NONE>(std::set<entry>& _used) const
 {
    for(const auto& itr : *this)
    {
-        if(std::forward<FuncT>(_func)(itr)) return BoolV;
-    }
-    return !BoolV;
-}
-
-template <>
-void
-call_chain::generate_perfetto<Device::NONE>(std::set<entry>&, bool) const
-{}
-
-template <>
-void
-call_chain::generate_perfetto<Device::CPU>(std::set<entry>& _used, bool _basic) const
-{
-    static std::set<std::string> _static_strings{};
-    static std::mutex            _static_mutex{};
-    for(const auto& itr : *this)
-    {
-        if(_basic)
+        if(!_used.emplace(itr).second) continue;
+        if(itr.device == Device::CPU)
        {
-            if(itr.device == Device::CPU)
-            {
-                TRACE_EVENT_BEGIN("device-critical-trace", "CPU",
-                                  static_cast<uint64_t>(itr.begin_ns));
-            }
-            else if(itr.device == Device::GPU)
-            {
-                TRACE_EVENT_BEGIN("device-critical-trace", "GPU",
-                                  static_cast<uint64_t>(itr.begin_ns));
-            }
-            TRACE_EVENT_END("device-critical-trace", static_cast<uint64_t>(itr.end_ns));
-        }
-        else
-        {
-            if(!_used.emplace(itr).second) continue;
-            std::string _name = tim::demangle(tim::get_hash_identifier(itr.hash));
-            _static_mutex.lock();
-            auto sitr = _static_strings.emplace(_name);
-            _static_mutex.unlock();
-            TRACE_EVENT_BEGIN("host-critical-trace",
-                              perfetto::StaticString{ sitr.first->c_str() },
+            TRACE_EVENT_BEGIN("device-critical-trace", "CPU",
                              static_cast<uint64_t>(itr.begin_ns));
-            TRACE_EVENT_END("host-critical-trace", static_cast<uint64_t>(itr.end_ns));
        }
+        else if(itr.device == Device::GPU)
+        {
+            TRACE_EVENT_BEGIN("device-critical-trace", "GPU",
+                              static_cast<uint64_t>(itr.begin_ns));
+        }
+        TRACE_EVENT_END("device-critical-trace", static_cast<uint64_t>(itr.end_ns));
    }
 }

 template <>
 void
-call_chain::generate_perfetto<Device::GPU>(std::set<entry>& _used, bool _basic) const
+call_chain::generate_perfetto<Device::CPU>(std::set<entry>& _used) const
 {
    static std::set<std::string> _static_strings{};
    static std::mutex            _static_mutex{};
    for(const auto& itr : *this)
    {
-        if(_basic)
-        {
-            if(itr.device == Device::CPU)
-            {
-                TRACE_EVENT_BEGIN("device-critical-trace", "CPU",
-                                  static_cast<uint64_t>(itr.begin_ns));
-            }
-            else if(itr.device == Device::GPU)
-            {
-                TRACE_EVENT_BEGIN("device-critical-trace", "GPU",
-                                  static_cast<uint64_t>(itr.begin_ns));
-            }
-            TRACE_EVENT_END("device-critical-trace", static_cast<uint64_t>(itr.end_ns));
-        }
-        else
-        {
-            if(!_used.emplace(itr).second) continue;
-            std::string _name = tim::demangle(tim::get_hash_identifier(itr.hash));
-            _static_mutex.lock();
-            auto sitr = _static_strings.emplace(_name);
-            _static_mutex.unlock();
-            TRACE_EVENT_BEGIN("device-critical-trace",
-                              perfetto::StaticString{ sitr.first->c_str() },
-                              static_cast<uint64_t>(itr.begin_ns));
-            TRACE_EVENT_END("device-critical-trace", static_cast<uint64_t>(itr.end_ns));
-        }
+        if(!_used.emplace(itr).second) continue;
+        if(itr.device != Device::CPU) continue;
+        std::string _name = tim::demangle(tim::get_hash_identifier(itr.hash));
+        _static_mutex.lock();
+        auto sitr = _static_strings.emplace(_name);
+        _static_mutex.unlock();
+        TRACE_EVENT_BEGIN("host-critical-trace",
+                          perfetto::StaticString{ sitr.first->c_str() },
+                          static_cast<uint64_t>(itr.begin_ns));
+        TRACE_EVENT_END("host-critical-trace", static_cast<uint64_t>(itr.end_ns));
+    }
+}
+
+template <>
+void
+call_chain::generate_perfetto<Device::GPU>(std::set<entry>& _used) const
+{
+    static std::set<std::string> _static_strings{};
+    static std::mutex            _static_mutex{};
+    for(const auto& itr : *this)
+    {
+        if(!_used.emplace(itr).second) continue;
+        if(itr.device != Device::GPU) continue;
+        std::string _name = tim::demangle(tim::get_hash_identifier(itr.hash));
+        _static_mutex.lock();
+        auto sitr = _static_strings.emplace(_name);
+        _static_mutex.unlock();
+        TRACE_EVENT_BEGIN("device-critical-trace",
+                          perfetto::StaticString{ sitr.first->c_str() },
+                          static_cast<uint64_t>(itr.begin_ns));
+        TRACE_EVENT_END("device-critical-trace", static_cast<uint64_t>(itr.end_ns));
+    }
+}
+
+template <>
+void
+call_chain::generate_perfetto<Device::ANY>(std::set<entry>& _used) const
+{
+    static std::set<std::string> _static_strings{};
+    static std::mutex            _static_mutex{};
+    for(const auto& itr : *this)
+    {
+        if(!_used.emplace(itr).second) continue;
+        std::string _name = tim::demangle(tim::get_hash_identifier(itr.hash));
+        _static_mutex.lock();
+        auto sitr = _static_strings.emplace(_name);
+        _static_mutex.unlock();
+        TRACE_EVENT_BEGIN("critical-trace", perfetto::StaticString{ sitr.first->c_str() },
+                          static_cast<uint64_t>(itr.begin_ns));
+        TRACE_EVENT_END("critical-trace", static_cast<uint64_t>(itr.end_ns));
    }
 }

@@ -590,8 +562,6 @@ compute(int64_t _tid)

 namespace
 {
-//--------------------------------------------------------------------------------------//
-
 std::string
 get_perf_name(std::string _func)
 {
@@ -603,141 +573,6 @@ get_perf_name(std::string _func)
    return _func;
 }

-//--------------------------------------------------------------------------------------//
-
-void
-save_call_graph(const std::string& _fname, const std::string& _label,
-                const call_graph_t& _call_graph, bool _msg = false,
-                std::string _func = {})
-{
-    OMNITRACE_CT_DEBUG("[%s]\n", __FUNCTION__);
-
-    using perfstats_t =
-        tim::lightweight_tuple<comp::wall_clock, comp::peak_rss, comp::page_rss>;
-    perfstats_t _perf{ get_perf_name(__FUNCTION__) };
-    _perf.start();
-
-    std::stringstream oss{};
-    {
-        namespace cereal = tim::cereal;
-        auto ar = tim::policy::output_archive<cereal::MinimalJSONOutputArchive>::get(oss);
-
-        auto _hash_map = *tim::hash::get_hash_ids();
-        for(auto& itr : _hash_map)
-            itr.second = tim::demangle(itr.second);
-        ar->setNextName("omnitrace");
-        ar->startNode();
-        (*ar)(cereal::make_nvp("hash_map", _hash_map));
-        ar->setNextName(_label.c_str());
-        ar->startNode();
-        serialize_graph(*ar, _call_graph);
-        ar->finishNode();
-        ar->finishNode();
-    }
-
-    std::ofstream ofs{};
-    if(tim::filepath::open(ofs, _fname))
-    {
-        if(_msg)
-        {
-            if(_func.empty()) _func = __FUNCTION__;
-            OMNITRACE_CONDITIONAL_PRINT(get_verbose() >= 0, "[%s] Outputting '%s'...\n",
-                                        _func.c_str(), _fname.c_str());
-        }
-        ofs << oss.str() << std::endl;
-    }
-
-    _perf.stop();
-    if(_msg)
-    {
-        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).c_str());
-    }
-}
-
-void
-save_critical_trace(const std::string& _fname, const std::string& _label,
-                    const std::vector<call_chain>& _cchain, bool _msg = false,
-                    std::string _func = {})
-{
-    OMNITRACE_CT_DEBUG("[%s]\n", __FUNCTION__);
-
-    using perfstats_t =
-        tim::lightweight_tuple<comp::wall_clock, comp::peak_rss, comp::page_rss>;
-    perfstats_t _perf{ get_perf_name(__FUNCTION__) };
-    _perf.start();
-
-    auto _save = [&](std::ostream& _os) {
-        namespace cereal = tim::cereal;
-        auto ar = tim::policy::output_archive<cereal::MinimalJSONOutputArchive>::get(_os);
-
-        auto _hash_map = *tim::hash::get_hash_ids();
-        for(auto& itr : _hash_map)
-            itr.second = tim::demangle(itr.second);
-        ar->setNextName("omnitrace");
-        ar->startNode();
-        (*ar)(cereal::make_nvp("hash_map", _hash_map),
-              cereal::make_nvp(_label.c_str(), _cchain));
-        ar->finishNode();
-    };
-
-    std::ofstream ofs{};
-    if(tim::filepath::open(ofs, _fname))
-    {
-        if(_msg)
-        {
-            if(_func.empty()) _func = __FUNCTION__;
-            OMNITRACE_CONDITIONAL_PRINT(get_verbose() >= 0, "[%s] Outputting '%s'...\n",
-                                        _func.c_str(), _fname.c_str());
-        }
-        std::stringstream oss{};
-        if(_cchain.size() > 1000)
-        {
-            _save(ofs);
-        }
-        else
-        {
-            _save(oss);
-            ofs << oss.str() << std::endl;
-        }
-    }
-
-    _perf.stop();
-    if(_msg)
-    {
-        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).c_str());
-    }
-}
-
-void
-save_call_chain_text(const std::string& _fname, const call_chain& _call_chain,
-                     bool _msg = false, std::string _func = {})
-{
-    OMNITRACE_CT_DEBUG("[%s]\n", __FUNCTION__);
-
-    using perfstats_t =
-        tim::lightweight_tuple<comp::wall_clock, comp::peak_rss, comp::page_rss>;
-    perfstats_t _perf{ get_perf_name(__FUNCTION__) };
-    _perf.start();
-
-    std::ofstream ofs{};
-    if(tim::filepath::open(ofs, _fname))
-    {
-        if(_msg)
-        {
-            if(_func.empty()) _func = __FUNCTION__;
-            OMNITRACE_CONDITIONAL_PRINT(get_verbose() >= 0, "[%s] Outputting '%s'...\n",
-                                        _func.c_str(), _fname.c_str());
-        }
-        ofs << _call_chain << "\n";
-    }
-
-    _perf.stop();
-    if(_msg)
-    {
-        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).c_str());
-    }
-}
-
 void
 save_call_chain_json(const std::string& _fname, const std::string& _label,
                     const call_chain& _call_chain, bool _msg = false,
@@ -792,24 +627,6 @@ save_call_chain_json(const std::string& _fname, const std::string& _label,
    }
 }

-void
-load_call_chain(const std::string& _fname, const std::string& _label,
-                call_chain& _call_chain)
-{
-    std::ifstream ifs{};
-    ifs.open(_fname);
-    if(ifs && ifs.is_open())
-    {
-        namespace cereal = tim::cereal;
-        auto ar          = tim::policy::input_archive<cereal::JSONInputArchive>::get(ifs);
-
-        ar->setNextName("omnitrace");
-        ar->startNode();
-        (*ar)(cereal::make_nvp(_label.c_str(), _call_chain));
-        ar->finishNode();
-    }
-}
-
 template <typename Tp, template <typename...> class ContainerT, typename... Args,
          typename FuncT = bool (*)(const Tp&, const Tp&)>
 inline Tp*
@@ -912,224 +729,16 @@ combine_critical_path(call_chain& _targ, call_chain _chain)
    // squash_critical_path(_targ);
 }

-auto
-get_indexed(const call_chain& _chain)
-{
-    OMNITRACE_CT_DEBUG("[%s]\n", __FUNCTION__);
-    std::map<int64_t, std::vector<entry>> _indexed = {};
-
-    // allocate for all cpu correlation ids
-    for(const auto& itr : _chain)
-    {
-        _indexed.emplace(static_cast<int64_t>(itr.cpu_cid), std::vector<entry>{});
-        _indexed.emplace(static_cast<int64_t>(itr.parent_cid), std::vector<entry>{});
-    }
-
-    // index based on parent correlation id
-    for(const auto& itr : _chain)
-    {
-        if(itr.depth < 1 && itr.phase == Phase::BEGIN) continue;
-        _indexed[static_cast<int64_t>(itr.parent_cid)].emplace_back(itr);
-    }
-
-    for(auto& itr : _indexed)
-        std::sort(itr.second.begin(), itr.second.end(),
-                  [](const entry& lhs, const entry& rhs) {
-                      // return lhs.cpu_cid < rhs.cpu_cid;
-                      return lhs.begin_ns < rhs.begin_ns;
-                  });
-
-    return _indexed;
-}
-
 void
-find_children(PTL::ThreadPool& _tp, call_graph_t& _graph, const call_chain& _chain)
-{
-    OMNITRACE_CT_DEBUG("[%s]\n", __FUNCTION__);
-
-    using iterator_t      = call_graph_sibling_itr_t;
-    using itr_entry_vec_t = std::vector<std::pair<iterator_t, entry>>;
-    using task_group_t    = PTL::TaskGroup<void>;
-
-    auto                                _indexed = get_indexed(_chain);
-    std::map<entry, std::vector<entry>> _entry_map{};
-
-    // allocate all entries
-    OMNITRACE_CT_DEBUG("[%s] Allocating...\n", __FUNCTION__);
-    for(const auto& itr : _chain)
-    {
-        auto _ins = _entry_map.emplace(itr, std::vector<entry>{});
-        if(!_ins.second)
-        {
-            auto _existing = _ins.first->first;
-            OMNITRACE_PRINT("Warning! Duplicate entry for [%s] :: [%s]\n",
-                            JOIN("", _existing).c_str(), JOIN("", itr).c_str());
-        }
-    }
-
-    task_group_t _tg{ &_tp };
-    OMNITRACE_CT_DEBUG("[%s] Parallel mapping...\n", __FUNCTION__);
-    for(const auto& itr : _chain)
-    {
-        _tg.run([&]() { _entry_map[itr] = _indexed.at(itr.cpu_cid); });
-    }
-    _tg.join();
-
-    std::function<void(iterator_t, const entry&)> _recursive_func;
-    _recursive_func = [&](iterator_t itr, const entry& _v) {
-        auto _child    = _graph.append_child(itr, _v);
-        auto _children = std::move(_entry_map[_v]);
-        _entry_map[_v].clear();
-        for(auto&& vitr : _children)
-        {
-            _recursive_func(_child, vitr);
-        }
-    };
-
-    // the recursive version of _func + _loop_func has a tendency to overflow the stack
-    auto _func = [&](iterator_t itr, const entry& _v) {
-        auto _child    = _graph.append_child(itr, _v);
-        auto _children = std::move(_entry_map[_v]);
-        _entry_map[_v].clear();
-        itr_entry_vec_t _data{};
-        for(auto&& vitr : _children)
-            _data.emplace_back(_child, vitr);
-        return _data;
-    };
-
-    auto _loop_func = [&_func](itr_entry_vec_t& _data) {
-        auto _inp = _data;
-        _data.clear();
-        for(auto itr : _inp)
-        {
-            for(auto&& fitr : _func(itr.first, itr.second))
-                _data.emplace_back(std::move(fitr));
-        }
-        // if data is empty return false so we can break out of while loop
-        return !_data.empty();
-    };
-
-    if(!_indexed.at(-1).empty())
-    {
-        OMNITRACE_CT_DEBUG("[%s] Setting root (line %i)...\n", __FUNCTION__, __LINE__);
-        _graph.set_head(_indexed.at(-1).front());
-    }
-    else
-    {
-        OMNITRACE_CT_DEBUG("[%s] Setting root (line %i)...\n", __FUNCTION__, __LINE__);
-        auto  _depth = static_cast<uint16_t>(-1);
-        entry _root{ 0, Device::NONE, Phase::NONE, _depth, 0, 0, 0, 0, 0, 0, 0 };
-        _graph.set_head(_root);
-    }
-
-    iterator_t _root = _graph.begin();
-    for(auto&& itr : _entry_map)
-    {
-        if(itr.first.depth == _root->depth + 1)
-        {
-            OMNITRACE_CT_DEBUG("[%s] Generating call-graph...\n", __FUNCTION__);
-            // _recursive_func(_root, itr.first);
-            itr_entry_vec_t _data = _func(_root, itr.first);
-            while(_loop_func(_data))
-            {}
-        }
-    }
-}
-
-void
-find_sequences(PTL::ThreadPool& _tp, call_graph_t& _graph,
-               std::vector<call_chain>& _chain)
-{
-    OMNITRACE_CT_DEBUG("[%s]\n", __FUNCTION__);
-    using iterator_t = call_graph_preorder_itr_t;
-    std::vector<iterator_t> _end_nodes{};
-    size_t                  _n = 0;
-    for(iterator_t itr = _graph.begin(); itr != _graph.end(); ++itr, ++_n)
-    {
-        auto _nchild = _graph.number_of_children(itr);
-        if(_nchild > 0)
-        {
-            OMNITRACE_CT_DEBUG("Skipping node #%zu with %u children :: %s\n", _n, _nchild,
-                               JOIN("", *itr).c_str());
-            continue;
-        }
-        _end_nodes.emplace_back(itr);
-    }
-    OMNITRACE_CT_DEBUG("Number of end nodes: %zu\n", _end_nodes.size());
-    _chain.resize(_end_nodes.size());
-
-    auto _construct = [&](size_t i) {
-        auto itr = _end_nodes.at(i);
-        while(itr != nullptr && _graph.is_valid(itr))
-        {
-            _chain.at(i).emplace_back(*itr);
-            itr = _graph.parent(itr);
-        }
-        std::reverse(_chain.at(i).begin(), _chain.at(i).end());
-    };
-
-    PTL::TaskGroup<void> _tg{ &_tp };
-    for(size_t i = 0; i < _end_nodes.size(); ++i)
-        _tg.run(_construct, i);
-    _tg.join();
-
-    std::sort(_chain.begin(), _chain.end(),
-              [](const call_chain& lhs, const call_chain& rhs) {
-                  return lhs.get_cost() > rhs.get_cost();
-              });
-}
-
-template <typename ArchiveT, typename T, typename AllocatorT>
-void
-serialize_graph(ArchiveT& ar, const tim::graph<T, AllocatorT>& t)
-{
-    OMNITRACE_CT_DEBUG("[%s]\n", __FUNCTION__);
-
-    namespace cereal = tim::cereal;
-    using iterator_t = typename tim::graph<T, AllocatorT>::sibling_iterator;
-
-    ar(cereal::make_nvp("graph_nodes", t.size()));
-    ar.setNextName("graph");
-    ar.startNode();
-    ar.makeArray();
-    for(iterator_t itr = t.begin(); itr != t.end(); ++itr)
-        serialize_subgraph(ar, t, itr);
-    ar.finishNode();
-}
-
-template <typename ArchiveT, typename T, typename AllocatorT>
-void
-serialize_subgraph(ArchiveT& ar, const tim::graph<T, AllocatorT>& _graph,
-                   typename tim::graph<T, AllocatorT>::iterator _root)
-{
-    using iterator_t = typename tim::graph<T, AllocatorT>::sibling_iterator;
-
-    if(_graph.empty()) return;
-
-    ar.setNextName("node");
-    ar.startNode();
-    ar(*_root);
-    {
-        ar.setNextName("children");
-        ar.startNode();
-        ar.makeArray();
-        for(iterator_t itr = _graph.begin(_root); itr != _graph.end(_root); ++itr)
-            serialize_subgraph(ar, _graph, itr);
-        ar.finishNode();
-    }
-    ar.finishNode();
-}
-
-void
-update_critical_path(call_chain _chain, int64_t _tid)
+update_critical_path(call_chain _chain, int64_t)
 {
    OMNITRACE_CT_DEBUG("[%s]\n", __FUNCTION__);
    try
    {
        // remove any data not
-        auto _diff_tid = [_tid](const entry& _v) { return _v.tid != _tid; };
-        _chain.erase(std::remove_if(_chain.begin(), _chain.end(), _diff_tid),
-                     _chain.end());
+        // auto _diff_tid = [_tid](const entry& _v) { return _v.tid != _tid; };
+        //_chain.erase(std::remove_if(_chain.begin(), _chain.end(), _diff_tid),
+        //             _chain.end());
        combine_critical_path(complete_call_chain, std::move(_chain));
    } catch(const std::exception& e)
    {
@@ -1138,56 +747,6 @@ update_critical_path(call_chain _chain, int64_t _tid)
    }
 }

-template <Device DevT>
-std::vector<call_chain>
-get_top(const std::vector<call_chain>& _chain, size_t _count)
-{
-    OMNITRACE_CT_DEBUG("[%s]\n", __FUNCTION__);
-    std::vector<call_chain> _data{};
-    _data.reserve(_count);
-    for(const auto& itr : _chain)
-    {
-        if(_data.size() >= _count) break;
-        if(itr.query<>([](const entry& _v) { return (_v.device == DevT); }))
-        {
-            _data.emplace_back(itr);
-        }
-    }
-    return _data;
-}
-
-template <Device DevT>
-void
-generate_perfetto(const std::vector<call_chain>& _data, bool _basic = false)
-{
-    OMNITRACE_CT_DEBUG("[%s]\n", __FUNCTION__);
-
-    auto _nrows = std::min<size_t>(get_critical_trace_per_row(), _data.size());
-
-    // run in separate thread(s) so that it ends up in unique row
-    if(_nrows < 1) _nrows = _data.size();
-
-    auto _func = [&](size_t _idx, size_t _beg, size_t _end) {
-        if(_nrows != 1)
-            threading::set_thread_name(TIMEMORY_JOIN(" ", "CriticalPath", _idx).c_str());
-        else
-            threading::set_thread_name("CritialPath");
-        // ensure all hash ids exist
-        copy_hash_ids();
-        std::set<entry> _used{};
-        for(size_t i = _beg; i < _end; ++i)
-        {
-            if(i >= _data.size()) break;
-            _data.at(i).generate_perfetto<DevT>(_used, _basic);
-        }
-    };
-
-    for(size_t i = 0; i < _data.size(); i += _nrows)
-    {
-        std::thread{ _func, i, i, i + _nrows }.join();
-    }
-}
-
 void
 compute_critical_trace()
 {
@@ -1196,7 +755,8 @@ compute_critical_trace()

    if(_computed) return;

-    OMNITRACE_CT_DEBUG("[%s] Generating critical trace...\n", __FUNCTION__);
+    OMNITRACE_CONDITIONAL_PRINT(get_critical_trace_debug() || get_verbose() >= 0,
+                                "[%s] Generating critical trace...\n", __FUNCTION__);

    // ensure all hash ids exist
    copy_hash_ids();
@@ -1209,11 +769,6 @@ compute_critical_trace()

    try
    {
-        PTL::ThreadPool _tp{ get_critical_trace_num_threads(), []() { copy_hash_ids(); },
-                             []() {} };
-        _tp.set_verbose(-1);
-        PTL::TaskGroup<void> _tg{ &_tp };
-
        OMNITRACE_CT_DEBUG("[%s] initial call chain: %zu entries\n", __FUNCTION__,
                           complete_call_chain.size());

@@ -1225,110 +780,14 @@ compute_critical_trace()
        _perf.stop().rekey("Sorting critical trace");
        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).c_str());

-        _perf.start();
-
-        squash_critical_path(complete_call_chain);
-
-        OMNITRACE_CT_DEBUG("[%s] complete call chain: %zu entries\n", __FUNCTION__,
-                           complete_call_chain.size());
-
-        _perf.stop().rekey("Squash critical path");
-        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).c_str());
-
-        _tg.run(
-            [](call_chain _chain, std::string _func) {  // NOLINT
-                save_call_chain_json(tim::settings::compose_output_filename(
-                                         "call-chain", ".json", get_use_pid(),
-                                         (tim::dmp::is_initialized())
-                                             ? tim::dmp::rank()
-                                             : process::get_id()),
-                                     "call_chain", _chain, true, std::move(_func));
-            },
-            complete_call_chain, __FUNCTION__);
-
        _perf.reset().start();
+        save_call_chain_json(
+            tim::settings::compose_output_filename("call-chain", ".json"), "call_chain",
+            complete_call_chain, true, __FUNCTION__);

-        OMNITRACE_CT_DEBUG("[%s] Finding children...\n", __FUNCTION__);
-        call_graph_t _graph{};
-        find_children(_tp, _graph, complete_call_chain);
-
-        OMNITRACE_CT_DEBUG("[%s] complete call graph: %zu entries\n", __FUNCTION__,
-                           _graph.size() - 1);
-
-        _perf.stop().rekey("Finding children");
+        _perf.stop().rekey("Save call-chain");
        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).c_str());

-        _tg.run(
-            [&](std::string _func) {
-                save_call_graph(tim::settings::compose_output_filename(
-                                    "call-graph", ".json", get_use_pid(),
-                                    (tim::dmp::is_initialized()) ? tim::dmp::rank()
-                                                                 : process::get_id()),
-                                "call_graph", _graph, true, std::move(_func));
-            },
-            __FUNCTION__);
-
-        _tg.join();
-
-        _perf.reset().start();
-
-        generate_perfetto<Device::GPU>({ complete_call_chain }, true);
-
-        OMNITRACE_CT_DEBUG("[%s] Finding sequences...\n", __FUNCTION__);
-        // find the sequences
-        std::vector<call_chain> _top{};
-        find_sequences(_tp, _graph, _top);
-
-        _perf.stop().rekey("Finding sequences");
-        OMNITRACE_CT_DEBUG("%s\n", JOIN("", _perf).c_str());
-
-        if(get_critical_trace_count() == 0)
-        {
-            OMNITRACE_CT_DEBUG("[%s] Saving critical trace...\n", __FUNCTION__);
-            save_critical_trace(
-                tim::settings::compose_output_filename(
-                    "critical-trace", ".json", get_use_pid(),
-                    (tim::dmp::is_initialized()) ? tim::dmp::rank() : process::get_id()),
-                "critical_trace", _top, true, __FUNCTION__);
-        }
-        else
-        {
-            OMNITRACE_CT_DEBUG("[%s] Getting top CPU functions...\n", __FUNCTION__);
-            // get the top CPU critical traces
-            auto _top_cpu = get_top<Device::CPU>(_top, get_critical_trace_count());
-            if(!_top_cpu.empty())
-            {
-                OMNITRACE_CT_DEBUG("[%s] Generating perfetto CPU critical traces...\n",
-                                   __FUNCTION__);
-                generate_perfetto<Device::CPU>(_top_cpu);
-                OMNITRACE_CT_DEBUG("[%s] Saving CPU critical traces...\n", __FUNCTION__);
-                save_critical_trace(tim::settings::compose_output_filename(
-                                        "critical-trace-cpu", ".json", get_use_pid(),
-                                        (tim::dmp::is_initialized()) ? tim::dmp::rank()
-                                                                     : process::get_id()),
-                                    "critical_trace", _top_cpu, true, __FUNCTION__);
-            }
-
-            OMNITRACE_CT_DEBUG("[%s] Getting top GPU functions...\n", __FUNCTION__);
-            // get the top GPU critical traces
-            auto _top_gpu = get_top<Device::GPU>(_top, get_critical_trace_count());
-            if(!_top_gpu.empty())
-            {
-                OMNITRACE_CT_DEBUG("[%s] Generating perfetto GPU critical traces...\n",
-                                   __FUNCTION__);
-                generate_perfetto<Device::GPU>(_top_gpu);
-                OMNITRACE_CT_DEBUG("[%s] Saving GPU critical traces...\n", __FUNCTION__);
-                save_critical_trace(tim::settings::compose_output_filename(
-                                        "critical-trace-gpu", ".json", get_use_pid(),
-                                        (tim::dmp::is_initialized()) ? tim::dmp::rank()
-                                                                     : process::get_id()),
-                                    "critical_trace", _top_gpu, true, __FUNCTION__);
-            }
-        }
-
-        _tg.join();
-        _tp.destroy_threadpool();
-        _computed = true;
    } catch(std::exception& e)
    {
        OMNITRACE_PRINT("Thread exited '%s' with exception: %s\n", __FUNCTION__,
@@ -319,8 +319,4 @@ extern "C"
    //----------------------------------------------------------------------------------//
 }

-//--------------------------------------------------------------------------------------//
-
 TIMEMORY_INITIALIZE_STORAGE(kokkosp::memory_tracker)
-
-//--------------------------------------------------------------------------------------//
@@ -0,0 +1,196 @@
+// 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/thread_sampler.hpp"
+#include "library/components/rocm_smi.hpp"
+#include "library/config.hpp"
+#include "library/cpu_freq.hpp"
+#include "library/debug.hpp"
+
+#include <memory>
+#include <vector>
+
+namespace omnitrace
+{
+namespace thread_sampler
+{
+namespace
+{
+using tim::type_mutex;
+using auto_lock_t                                       = tim::auto_lock_t;
+using promise_t                                         = std::promise<void>;
+std::unique_ptr<promise_t>             polling_finished = {};
+std::vector<std::unique_ptr<instance>> instances        = {};
+
+bool&
+is_initialized()
+{
+    static bool _v = false;
+    return _v;
+}
+
+std::unique_ptr<std::thread>&
+get_thread()
+{
+    static std::unique_ptr<std::thread> _v;
+    return _v;
+}
+
+std::atomic<State>&
+get_sampler_state()
+{
+    static std::atomic<State> _v{ State::PreInit };
+    return _v;
+}
+}  // namespace
+
+void
+sampler::poll(std::atomic<State>* _state, nsec_t _interval, promise_t* _ready)
+{
+    threading::set_thread_name("omni.sampler");
+
+    // notify thread started
+    if(_ready) _ready->set_value();
+
+    for(auto& itr : instances)
+        itr->config();
+
+    OMNITRACE_CONDITIONAL_BASIC_PRINT(
+        get_verbose() > 0 || get_debug(),
+        "Thread sampler polling at an interval of %f seconds...\n",
+        std::chrono::duration_cast<std::chrono::duration<double>>(_interval).count());
+
+    auto _now = std::chrono::steady_clock::now();
+    while(_state && _state->load() != State::Finalized && get_state() != State::Finalized)
+    {
+        std::this_thread::sleep_until(_now);
+        if(_state->load() != State::Active) continue;
+        for(auto& itr : instances)
+            itr->sample();
+        while(_now < std::chrono::steady_clock::now())
+            _now += _interval;
+    }
+
+    OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug(),
+                                      "Thread sampler polling completed...\n");
+
+    if(polling_finished) polling_finished->set_value();
+}
+
+void
+sampler::setup()
+{
+    OMNITRACE_VERBOSE(1, "Setting up background sampler...\n");
+
+    // shutdown if already running
+    shutdown();
+
+    auto _enable_samp = pthread_gotcha::enable_sampling_on_child_threads();
+    pthread_gotcha::enable_sampling_on_child_threads() = false;
+
+    if(get_use_rocm_smi())
+    {
+        auto& _rocm_smi         = instances.emplace_back(std::make_unique<instance>());
+        _rocm_smi->setup        = []() { rocm_smi::setup(); };
+        _rocm_smi->shutdown     = []() { rocm_smi::shutdown(); };
+        _rocm_smi->post_process = []() { rocm_smi::post_process(); };
+        _rocm_smi->config       = []() { rocm_smi::config(); };
+        _rocm_smi->sample       = []() { rocm_smi::sample(); };
+    }
+
+    auto& _cpu_freq         = instances.emplace_back(std::make_unique<instance>());
+    _cpu_freq->setup        = []() { cpu_freq::setup(); };
+    _cpu_freq->shutdown     = []() { cpu_freq::shutdown(); };
+    _cpu_freq->post_process = []() { cpu_freq::post_process(); };
+    _cpu_freq->config       = []() { cpu_freq::config(); };
+    _cpu_freq->sample       = []() { cpu_freq::sample(); };
+
+    for(auto& itr : instances)
+        itr->setup();
+
+    polling_finished = std::make_unique<promise_t>();
+
+    auto     _freq      = get_thread_sampling_freq();
+    uint64_t _msec_freq = (1.0 / _freq) * 1.0e3;
+
+    promise_t _prom{};
+    auto      _fut   = _prom.get_future();
+    polling_finished = std::make_unique<promise_t>();
+
+    set_state(State::PreInit);
+    get_thread() = std::make_unique<std::thread>(&poll<msec_t>, &get_sampler_state(),
+                                                 msec_t{ _msec_freq }, &_prom);
+
+    _fut.wait();
+
+    pthread_gotcha::enable_sampling_on_child_threads() = _enable_samp;
+    set_state(State::Active);
+}
+
+void
+sampler::shutdown()
+{
+    for(auto& itr : instances)
+        itr->shutdown();
+
+    auto& _thread = get_thread();
+    if(_thread)
+    {
+        OMNITRACE_VERBOSE(1, "Shutting down background sampler...\n");
+        set_state(State::Finalized);
+        if(polling_finished)
+        {
+            auto     _fut  = polling_finished->get_future();
+            uint64_t _freq = (1.0 / get_thread_sampling_freq()) * 1.0e3;
+            _fut.wait_for(msec_t{ 5 * _freq });
+            _thread->join();
+        }
+        else
+        {
+            uint64_t _freq = (1.0 / get_thread_sampling_freq()) * 1.0e3;
+            std::this_thread::sleep_for(msec_t{ 5 * _freq });
+            pthread_cancel(_thread->native_handle());
+            _thread->detach();
+        }
+        _thread          = std::unique_ptr<std::thread>{};
+        polling_finished = std::unique_ptr<promise_t>{};
+    }
+
+    is_initialized() = false;
+}
+
+void
+sampler::post_process()
+{
+    for(auto& itr : instances)
+        itr->post_process();
+
+    instances.clear();
+}
+
+void
+sampler::set_state(state_t _state)
+{
+    get_sampler_state().store(_state);
+}
+}  // namespace thread_sampler
+}  // namespace omnitrace
@@ -214,7 +214,7 @@ omnitrace_add_test(
    REWRITE_ARGS -e -v 2
    RUNTIME_ARGS -e -v 1 --label file line return args
    RUN_ARGS ""
-    ENVIRONMENT "${_base_environment};OMNITRACE_CRITICAL_TRACE=OFF")
+    ENVIRONMENT "${_base_environment};OMNITRACE_CRITICAL_TRACE=ON")

 omnitrace_add_test(
    NAME transpose-no-save-fpr