rocprofiler-sdk-tool library intermediate binary output (#734)

* Support for binary temporary files * clang formatting * formating ring buffer.hpp * Update source/lib/common/container/ring_buffer.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * fixing bugs * fix loop range * Fix for v3 test failures * bug fix * fix bug * fix memory leaks * destructing agent_info * Update CMakeLists.txt * clang-tidy fixes * Fix data race on destructor of rocprofiler_agent_t map in rocprofiler-sdk-tool library * Create lib/rocproifler-sdk-tool/tmp_file.* - move tmp_file class into separate header/implementation * Agent Info CSV in rocprofiler-sdk-tool - update tests to use agent_info.csv instead of rocminfo * Update lib/rocprofiler-sdk-tool/tool.cpp - use logical_node_id instead of node_id * Adding stats file * Adding tests for stats * Update scratch memory support - convert scratch memory support to use binary output * Tool Update: scratch memory stats + extended statistics - replace generate_*_csv with generate_csv overloads - added generate_csv for scratch memory - enable stats for scratch memory - replace ROCPROF_*_STATS env variables with ROCPROF_STATS env variable * rocprofv3 update - simple --stats option - add scratch memory trace to --sys-trace * Update tests/rocprofv3/tracing-hip-in-libraries - extend validate.py to test stats data - fix conftest.py for memory_copy_stats_data * Code coverage fixes - invoke __gcov_dump to ensure that code coverage is flushed after finalization --------- Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Co-authored-by: Ammar ELWazir <ammar.elwazir@amd.com> Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>
2024-04-09 05:25:28 -05:00
@@ -76,7 +76,7 @@ target_compile_options(
    INTERFACE $<$<COMPILE_LANGUAGE:C>:$<$<C_COMPILER_ID:GNU>:-rdynamic>>
              $<$<COMPILE_LANGUAGE:CXX>:$<$<CXX_COMPILER_ID:GNU>:-rdynamic>>)

-if(NOT APPLE)
+if(NOT APPLE AND NOT ROCPROFILER_ENABLE_CLANG_TIDY)
    target_link_options(rocprofiler-debug-flags INTERFACE
                        $<$<CXX_COMPILER_ID:GNU>:-rdynamic>)
 endif()
@@ -42,7 +42,8 @@ usage() {
  echo -e "${GREEN}--hsa-image-trace ${RESET}            For Collecting HSA API Traces (Image-extenson API)"
  echo -e "${GREEN}--hsa-finalizer-trace ${RESET}        For Collecting HSA API Traces (Finalizer-extension API)"
  echo -e ""
-  echo -e "${GREEN}--sys-trace ${RESET}                  For Collecting HIP,HSA, Memory Copy, (marker)ROCTx and Kernel dispatch traces\n"
+  echo -e "${GREEN}--sys-trace ${RESET}                  For Collecting HIP, HSA, Marker (ROCTx), Memory copy, Scratch memory, and Kernel dispatch traces\n"
+  echo -e "${GREEN}--stats ${RESET}                      For Collecting statistics of enabled tracing types\n"
  echo -e ""
  echo -e "${GREEN}-o   | --output-file ${RESET}         For the output file name"
  echo -e "\t#${GREY} usage (with current dir): rocprofv3 --hsa-trace -o <file_name> <executable>"
@@ -177,6 +178,10 @@ while true; do
    export ROCPROF_MARKER_API_TRACE=1
    export ROCPROF_KERNEL_TRACE=1
    export ROCPROF_MEMORY_COPY_TRACE=1
+    export ROCPROF_SCRATCH_MEMORY_TRACE=1
+    shift
+  elif [ "$1" == "--stats" ]; then
+    export ROCPROF_STATS=1
    shift
  elif [ "$1" == "--" ]; then
    shift
@@ -92,7 +92,7 @@ struct ring_buffer
    /// memory is uninitialized. Typically used by allocators. If Tp is a class type,
    /// be sure to use a placement new instead of a memcpy.
    template <typename Tp>
-    Tp* request();
+    Tp* request(bool wrap = true);

    /// Read class-type data from buffer (uses placement new).
    template <typename Tp>
@@ -210,11 +210,11 @@ ring_buffer::write(Tp* in, std::enable_if_t<!std::is_class<Tp>::value, int>)
 //
 template <typename Tp>
 Tp*
-ring_buffer::request()
+ring_buffer::request(bool wrap)
 {
    if(m_ptr == nullptr) return nullptr;

-    return request(sizeof(Tp));
+    return reinterpret_cast<Tp*>(request(sizeof(Tp), wrap));
 }
 //
 template <typename Tp>
@@ -266,7 +266,7 @@ ring_buffer::retrieve() const
 {
    if(m_ptr == nullptr) return nullptr;

-    return retrieve(sizeof(Tp));
+    return reinterpret_cast<Tp*>(retrieve(sizeof(Tp)));
 }
 //
 }  // namespace base
@@ -320,7 +320,7 @@ struct ring_buffer : private base::ring_buffer
    Tp* read(Tp* _dest) const { return base_type::read<Tp>(_dest).second; }

    /// Get an uninitialized address at tail of buffer.
-    Tp* request() { return base_type::request<Tp>(); }
+    Tp* request(bool wrap = true) { return base_type::request<Tp>(wrap); }

    /// Read data from head of buffer.
    Tp* retrieve() { return base_type::retrieve<Tp>(); }
@@ -337,6 +337,8 @@ struct ring_buffer : private base::ring_buffer
    /// Returns if the buffer is full.
    bool is_full() const { return (base_type::free() < sizeof(Tp)); }

+    bool clear() { return base_type::clear(); }
+
    template <typename... Args>
    auto emplace(Args&&... args)
    {
@@ -2,10 +2,15 @@
 # Tool library used by rocprofiler
 #

+rocprofiler_activate_clang_tidy()
+
+set(TOOL_HEADERS config.hpp csv.hpp generateCSV.hpp helper.hpp output_file.hpp
+                 statistics.hpp tmp_file.hpp)
+set(TOOL_SOURCES config.cpp generateCSV.cpp helper.cpp output_file.cpp tmp_file.cpp
+                 tool.cpp)
+
 add_library(rocprofiler-sdk-tool SHARED)
-target_sources(
-    rocprofiler-sdk-tool PRIVATE config.hpp config.cpp csv.hpp helper.hpp helper.cpp
-                                 output_file.hpp output_file.cpp tool.cpp)
+target_sources(rocprofiler-sdk-tool PRIVATE ${TOOL_SOURCES} ${TOOL_HEADERS})

 add_subdirectory(plugins)

@@ -15,6 +20,11 @@ target_link_libraries(
            rocprofiler::rocprofiler-build-flags rocprofiler::rocprofiler-memcheck
            rocprofiler::rocprofiler-common-library)

+if(ROCPROFILER_BUILD_CODECOV)
+    target_compile_definitions(rocprofiler-sdk-tool PRIVATE ROCPROFILER_CODECOV=1)
+    target_link_libraries(rocprofiler-sdk-tool PRIVATE gcov)
+endif()
+
 set_target_properties(
    rocprofiler-sdk-tool
    PROPERTIES LIBRARY_OUTPUT_DIRECTORY
@@ -48,13 +48,13 @@ namespace tool
 {
 namespace
 {
-auto launch_time = std::make_unique<std::time_t>(std::time_t{std::time(nullptr)});
+auto launch_time = new std::time_t{std::time(nullptr)};

 std::string
 get_local_datetime(const char* dt_format, std::time_t* dt_curr)
 {
    char mbstr[512];
-    if(!dt_curr) dt_curr = launch_time.get();
+    if(!dt_curr) dt_curr = launch_time;

    if(std::strftime(mbstr, sizeof(mbstr), dt_format, std::localtime(dt_curr)) != 0)
        return std::string{mbstr};
@@ -267,7 +267,7 @@ output_keys(std::string _tag)
        }
    }

-    auto* _launch_time = launch_time.get();
+    auto* _launch_time = launch_time;
    auto  _time_format = get_env<std::string>("ROCP_TIME_FORMAT", "%F_%H.%M");

    auto _mpi_size = get_mpi_size();
@@ -72,10 +72,13 @@ struct config
    bool        hip_compiler_api_trace      = get_env("ROCPROF_HIP_COMPILER_API_TRACE", false);
    bool        list_metrics                = get_env("ROCPROF_LIST_METRICS", false);
    bool        list_metrics_output_file    = get_env("ROCPROF_OUTPUT_LIST_METRICS_FILE", false);
+    bool        stats                       = get_env("ROCPROF_STATS", false);
    int         mpi_size                    = get_mpi_size();
    int         mpi_rank                    = get_mpi_rank();
-    std::string output_path = get_env("ROCPROF_OUTPUT_PATH", fs::current_path().string());
-    std::string output_file = get_env("ROCPROF_OUTPUT_FILE_NAME", std::to_string(getpid()));
+    std::string output_path   = get_env("ROCPROF_OUTPUT_PATH", fs::current_path().string());
+    std::string output_file   = get_env("ROCPROF_OUTPUT_FILE_NAME", std::to_string(getpid()));
+    std::string output_format = get_env("ROCPROF_OUTPUT_FORMAT", "CSV");
+    std::string tmp_directory = get_env("ROCPROF_TMPDIR", fs::current_path().string());
    std::vector<std::string> kernel_names = {};
    std::set<std::string>    counters     = {};
 };
@@ -86,14 +89,14 @@ get_config()
 {
    if constexpr(ContextT == config_context::global)
    {
-        static auto _v = config{};
-        return _v;
+        static auto* _v = new config{};
+        return *_v;
    }
    else
    {
        // context specific config copied from global config
-        static auto _v = get_config<config_context::global>();
-        return _v;
+        static auto* _v = new config{get_config<config_context::global>()};
+        return *_v;
    }
 }

@@ -26,8 +26,11 @@

 #include <array>
 #include <cstddef>
+#include <iomanip>
+#include <ios>
 #include <ostream>
 #include <string_view>
+#include <type_traits>

 namespace rocprofiler
 {
@@ -40,10 +43,22 @@ std::ostream&
 write_csv_entry(std::ostream& ofs, TupleT&& _data, std::index_sequence<Idx...>)
 {
    auto _write = [&ofs](size_t idx, auto&& _val) {
+        using value_type = common::mpl::unqualified_type_t<decltype(_val)>;
        if(idx > 0) ofs << ",";
-        if constexpr(rocprofiler::common::mpl::is_string_type<decltype(_val)>::value) ofs << "\"";
-        ofs << _val;
-        if constexpr(rocprofiler::common::mpl::is_string_type<decltype(_val)>::value) ofs << "\"";
+        if constexpr(std::is_floating_point<value_type>::value)
+        {
+            constexpr value_type one = 1;
+            if(_val >= one)
+                ofs << std::setprecision(6) << std::fixed << _val;
+            else
+                ofs << std::setprecision(6) << std::scientific << _val;
+        }
+        else
+        {
+            if constexpr(common::mpl::is_string_type<value_type>::value) ofs << "\"";
+            ofs << _val;
+            if constexpr(common::mpl::is_string_type<value_type>::value) ofs << "\"";
+        }
    };

    (_write(Idx, std::get<Idx>(_data)), ...);
@@ -75,6 +90,7 @@ struct csv_encoder
 };

 using api_csv_encoder                  = csv_encoder<7>;
+using agent_info_csv_encoder           = csv_encoder<53>;
 using kernel_trace_csv_encoder         = csv_encoder<16>;
 using counter_collection_csv_encoder   = csv_encoder<15>;
 using memory_copy_csv_encoder          = csv_encoder<7>;
@@ -82,6 +98,7 @@ using marker_csv_encoder               = csv_encoder<7>;
 using list_basic_metrics_csv_encoder   = csv_encoder<5>;
 using list_derived_metrics_csv_encoder = csv_encoder<5>;
 using scratch_memory_encoder           = csv_encoder<8>;
+using stats_csv_encoder                = csv_encoder<8>;
 }  // namespace csv
 }  // namespace tool
 }  // namespace rocprofiler
@@ -0,0 +1,446 @@
+// MIT License
+//
+// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#include "generateCSV.hpp"
+#include "csv.hpp"
+#include "helper.hpp"
+#include "statistics.hpp"
+
+#include <rocprofiler-sdk/fwd.h>
+#include <rocprofiler-sdk/marker/api_id.h>
+
+#include <utility>
+
+namespace rocprofiler
+{
+namespace tool
+{
+namespace
+{
+using float_type   = double;
+using stats_data_t = statistics<uint64_t, float_type>;
+using stats_map_t  = std::map<std::string_view, stats_data_t>;
+
+void
+write_stats(output_file& ofs, timestamps_t* app_timestamps, const stats_map_t& data)
+{
+    auto       _ss          = std::stringstream{};
+    float_type app_duration = (app_timestamps->app_end_time - app_timestamps->app_start_time);
+    for(const auto& [id, value] : data)
+    {
+        auto        duration_ns = value.get_sum();
+        auto        calls       = value.get_count();
+        const auto& name        = id;
+        float_type  avg_ns      = value.get_mean();
+        float_type  percentage  = (duration_ns / app_duration) * static_cast<float_type>(100);
+
+        rocprofiler::tool::csv::stats_csv_encoder::write_row(_ss,
+                                                             name,
+                                                             calls,
+                                                             duration_ns,
+                                                             avg_ns,
+                                                             percentage,
+                                                             value.get_min(),
+                                                             value.get_max(),
+                                                             value.get_stddev());
+    }
+    ofs << _ss.str();
+}
+}  // namespace
+
+void
+generate_csv(tool_table* tool_functions, std::vector<rocprofiler_agent_v0_t>& data)
+{
+    std::sort(data.begin(), data.end(), [](rocprofiler_agent_v0_t lhs, rocprofiler_agent_v0_t rhs) {
+        return lhs.node_id < rhs.node_id;
+    });
+
+    for(auto& itr : data)
+    {
+        auto _type = std::string_view{};
+        if(itr.type == ROCPROFILER_AGENT_TYPE_CPU)
+            _type = "CPU";
+        else if(itr.type == ROCPROFILER_AGENT_TYPE_GPU)
+            _type = "GPU";
+        else
+            _type = "UNK";
+
+        auto agent_info_ss = std::stringstream{};
+        rocprofiler::tool::csv::agent_info_csv_encoder::write_row(agent_info_ss,
+                                                                  itr.node_id,
+                                                                  itr.logical_node_id,
+                                                                  _type,
+                                                                  itr.cpu_cores_count,
+                                                                  itr.simd_count,
+                                                                  itr.cpu_core_id_base,
+                                                                  itr.simd_id_base,
+                                                                  itr.max_waves_per_simd,
+                                                                  itr.lds_size_in_kb,
+                                                                  itr.gds_size_in_kb,
+                                                                  itr.num_gws,
+                                                                  itr.wave_front_size,
+                                                                  itr.num_xcc,
+                                                                  itr.cu_count,
+                                                                  itr.array_count,
+                                                                  itr.num_shader_banks,
+                                                                  itr.simd_arrays_per_engine,
+                                                                  itr.cu_per_simd_array,
+                                                                  itr.simd_per_cu,
+                                                                  itr.max_slots_scratch_cu,
+                                                                  itr.gfx_target_version,
+                                                                  itr.vendor_id,
+                                                                  itr.device_id,
+                                                                  itr.location_id,
+                                                                  itr.domain,
+                                                                  itr.drm_render_minor,
+                                                                  itr.num_sdma_engines,
+                                                                  itr.num_sdma_xgmi_engines,
+                                                                  itr.num_sdma_queues_per_engine,
+                                                                  itr.num_cp_queues,
+                                                                  itr.max_engine_clk_ccompute,
+                                                                  itr.max_engine_clk_fcompute,
+                                                                  itr.sdma_fw_version.Value,
+                                                                  itr.fw_version.Value,
+                                                                  itr.capability.Value,
+                                                                  itr.cu_per_engine,
+                                                                  itr.max_waves_per_cu,
+                                                                  itr.family_id,
+                                                                  itr.workgroup_max_size,
+                                                                  itr.grid_max_size,
+                                                                  itr.local_mem_size,
+                                                                  itr.hive_id,
+                                                                  itr.gpu_id,
+                                                                  itr.workgroup_max_dim.x,
+                                                                  itr.workgroup_max_dim.y,
+                                                                  itr.workgroup_max_dim.z,
+                                                                  itr.grid_max_dim.x,
+                                                                  itr.grid_max_dim.y,
+                                                                  itr.grid_max_dim.z,
+                                                                  itr.name,
+                                                                  itr.vendor_name,
+                                                                  itr.product_name,
+                                                                  itr.model_name);
+        tool_functions->tool_get_agent_info_file_fn() << agent_info_ss.str();
+    }
+}
+
+void
+generate_csv(tool_table* tool_functions, std::vector<kernel_dispatch_ring_buffer_t>& data)
+{
+    auto kernel_stats = stats_map_t{};
+    for(auto& buf : data)
+    {
+        while(true)
+        {
+            auto kernel_trace_ss                                        = std::stringstream{};
+            rocprofiler_buffer_tracing_kernel_dispatch_record_t* record = buf.retrieve();
+            if(record == nullptr) break;
+            auto kernel_name = tool_functions->tool_get_kernel_name_fn(record->kernel_id);
+            rocprofiler::tool::csv::kernel_trace_csv_encoder::write_row(
+                kernel_trace_ss,
+                tool_functions->tool_get_domain_name_fn(record->kind),
+                tool_functions->tool_get_agent_node_id_fn(record->agent_id),
+                record->queue_id.handle,
+                record->kernel_id,
+                kernel_name,
+                record->correlation_id.internal,
+                record->start_timestamp,
+                record->end_timestamp,
+                record->private_segment_size,
+                record->group_segment_size,
+                record->workgroup_size.x,
+                record->workgroup_size.y,
+                record->workgroup_size.z,
+                record->grid_size.x,
+                record->grid_size.y,
+                record->grid_size.z);
+
+            if(tool::get_config().stats)
+                kernel_stats[kernel_name] += (record->end_timestamp - record->start_timestamp);
+
+            tool_functions->tool_get_kernel_trace_file_fn() << kernel_trace_ss.str();
+        }
+    }
+
+    if(tool::get_config().stats)
+        write_stats(tool_functions->tool_get_kernel_stats_file_fn(),
+                    tool_functions->tool_get_app_timestamps_fn(),
+                    kernel_stats);
+}
+
+void
+generate_csv(tool_table* tool_functions, std::vector<hip_ring_buffer_t>& data)
+{
+    auto hip_stats = stats_map_t{};
+    for(auto& buf : data)
+    {
+        while(true)
+        {
+            auto                                         hip_trace_ss = std::stringstream{};
+            rocprofiler_buffer_tracing_hip_api_record_t* record       = buf.retrieve();
+            if(record == nullptr) break;
+            auto api_name =
+                tool_functions->tool_get_operation_name_fn(record->kind, record->operation);
+            rocprofiler::tool::csv::api_csv_encoder::write_row(
+                hip_trace_ss,
+                tool_functions->tool_get_domain_name_fn(record->kind),
+                api_name,
+                getpid(),
+                record->thread_id,
+                record->correlation_id.internal,
+                record->start_timestamp,
+                record->end_timestamp);
+
+            if(tool::get_config().stats)
+                hip_stats[api_name] += (record->end_timestamp - record->start_timestamp);
+
+            tool_functions->tool_get_hip_api_trace_file_fn() << hip_trace_ss.str();
+        }
+    }
+
+    if(tool::get_config().stats)
+    {
+        write_stats(tool_functions->tool_get_hip_stats_file_fn(),
+                    tool_functions->tool_get_app_timestamps_fn(),
+                    hip_stats);
+    }
+}
+
+void
+generate_csv(tool_table* tool_functions, std::vector<hsa_ring_buffer_t>& data)
+{
+    auto hsa_stats = stats_map_t{};
+    for(auto& buf : data)
+    {
+        while(true)
+        {
+            auto                                         hsa_trace_ss = std::stringstream{};
+            rocprofiler_buffer_tracing_hsa_api_record_t* record       = buf.retrieve();
+            if(record == nullptr) break;
+            auto api_name =
+                tool_functions->tool_get_operation_name_fn(record->kind, record->operation);
+            rocprofiler::tool::csv::api_csv_encoder::write_row(
+                hsa_trace_ss,
+                tool_functions->tool_get_domain_name_fn(record->kind),
+                api_name,
+                getpid(),
+                record->thread_id,
+                record->correlation_id.internal,
+                record->start_timestamp,
+                record->end_timestamp);
+
+            if(tool::get_config().stats)
+                hsa_stats[api_name] += (record->end_timestamp - record->start_timestamp);
+
+            tool_functions->tool_get_hsa_api_trace_file_fn() << hsa_trace_ss.str();
+        }
+    }
+
+    if(tool::get_config().stats)
+    {
+        write_stats(tool_functions->tool_get_hsa_stats_file_fn(),
+                    tool_functions->tool_get_app_timestamps_fn(),
+                    hsa_stats);
+    }
+}
+
+void
+generate_csv(tool_table* tool_functions, std::vector<memory_copy_ring_buffer_t>& data)
+{
+    auto memory_copy_stats = stats_map_t{};
+    for(auto& buf : data)
+    {
+        while(true)
+        {
+            auto memory_copy_trace_ss                               = std::stringstream{};
+            rocprofiler_buffer_tracing_memory_copy_record_t* record = buf.retrieve();
+            if(record == nullptr) break;
+            auto api_name =
+                tool_functions->tool_get_operation_name_fn(record->kind, record->operation);
+            rocprofiler::tool::csv::memory_copy_csv_encoder::write_row(
+                memory_copy_trace_ss,
+                tool_functions->tool_get_domain_name_fn(record->kind),
+                api_name,
+                tool_functions->tool_get_agent_node_id_fn(record->src_agent_id),
+                tool_functions->tool_get_agent_node_id_fn(record->dst_agent_id),
+                record->correlation_id.internal,
+                record->start_timestamp,
+                record->end_timestamp);
+
+            if(tool::get_config().stats)
+                memory_copy_stats[api_name] += (record->end_timestamp - record->start_timestamp);
+
+            tool_functions->tool_get_memory_copy_trace_file_fn() << memory_copy_trace_ss.str();
+        }
+    }
+
+    if(tool::get_config().stats)
+    {
+        write_stats(tool_functions->tool_get_memory_copy_stats_file_fn(),
+                    tool_functions->tool_get_app_timestamps_fn(),
+                    memory_copy_stats);
+    }
+}
+
+void
+generate_csv(tool_table* tool_functions, std::vector<marker_api_ring_buffer_t>& data)
+{
+    for(auto& buf : data)
+    {
+        while(true)
+        {
+            auto                              marker_api_trace_ss = std::stringstream{};
+            rocprofiler_tool_marker_record_t* record              = buf.retrieve();
+            if(record == nullptr) break;
+            if(record->kind == ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_API &&
+               ((record->op == ROCPROFILER_MARKER_CORE_API_ID_roctxMarkA &&
+                 record->phase == ROCPROFILER_CALLBACK_PHASE_EXIT) ||
+                (record->op == ROCPROFILER_MARKER_CORE_API_ID_roctxRangePop &&
+                 record->phase == ROCPROFILER_CALLBACK_PHASE_ENTER) ||
+                (record->op == ROCPROFILER_MARKER_CORE_API_ID_roctxRangeStop &&
+                 record->phase == ROCPROFILER_CALLBACK_PHASE_ENTER)))
+            {
+                tool::csv::marker_csv_encoder::write_row(
+                    marker_api_trace_ss,
+                    tool_functions->tool_get_callback_kind_fn(record->kind),
+                    tool_functions->tool_get_roctx_msg_fn(record->cid),
+                    record->pid,
+                    record->tid,
+                    record->cid,
+                    record->start_timestamp,
+                    record->end_timestamp);
+            }
+            else
+            {
+                tool::csv::marker_csv_encoder::write_row(
+                    marker_api_trace_ss,
+                    tool_functions->tool_get_callback_kind_fn(record->kind),
+                    tool_functions->tool_get_callback_op_name_fn(record->kind, record->op),
+                    record->pid,
+                    record->tid,
+                    record->cid,
+                    record->start_timestamp,
+                    record->end_timestamp);
+            }
+            tool_functions->tool_get_marker_api_trace_file_fn() << marker_api_trace_ss.str();
+        }
+    }
+}
+
+void
+generate_csv(tool_table* tool_functions, std::vector<counter_collection_ring_buffer_t>& data)
+{
+    for(auto& buf : data)
+    {
+        while(true)
+        {
+            rocprofiler_tool_counter_collection_record_t* record = buf.retrieve();
+            if(record == nullptr) break;
+            auto kernel_id          = record->dispatch_data.kernel_id;
+            auto counter_name_value = std::map<std::string, uint64_t>{};
+            for(const auto& count : record->profiler_record)
+            {
+                auto        rec          = static_cast<rocprofiler_record_counter_t>(count);
+                std::string counter_name = tool_functions->tool_get_counter_info_name_fn(rec.id);
+                auto        search       = counter_name_value.find(counter_name);
+                if(search == counter_name_value.end())
+                    counter_name_value.emplace(
+                        std::pair<std::string, uint64_t>{counter_name, rec.counter_value});
+                else
+                    search->second = search->second + rec.counter_value;
+            }
+
+            const auto& correlation_id = record->dispatch_data.correlation_id;
+
+            auto magnitude = [](rocprofiler_dim3_t dims) { return (dims.x * dims.y * dims.z); };
+            auto counter_collection_ss = std::stringstream{};
+            for(auto& itr : counter_name_value)
+            {
+                tool::csv::counter_collection_csv_encoder::write_row(
+                    counter_collection_ss,
+                    correlation_id.internal,
+                    record->dispatch_index,
+                    tool_functions->tool_get_agent_node_id_fn(record->dispatch_data.agent_id),
+                    record->dispatch_data.queue_id.handle,
+                    record->pid,
+                    record->thread_id,
+                    magnitude(record->dispatch_data.grid_size),
+                    tool_functions->tool_get_kernel_name_fn(kernel_id),
+                    magnitude(record->dispatch_data.workgroup_size),
+                    record->lds_block_size_v,
+                    record->private_segment_size,
+                    record->arch_vgpr_count,
+                    record->sgpr_count,
+                    itr.first,
+                    itr.second);
+            }
+            tool_functions->tool_get_counter_collection_file_fn() << counter_collection_ss.str();
+        }
+    }
+}
+
+void
+generate_csv(tool_table* tool_functions, std::vector<scratch_memory_ring_buffer_t>& data)
+{
+    auto* ofs = tool_functions->tool_get_scratch_memory_file_fn();
+    if(!ofs) throw std::runtime_error{"error creating scratch memory output file"};
+
+    auto scratch_memory_stats = stats_map_t{};
+    for(auto& buf : data)
+    {
+        while(true)
+        {
+            rocprofiler_buffer_tracing_scratch_memory_record_t* record = buf.retrieve();
+            if(record == nullptr) break;
+
+            auto scratch_memory_trace = std::stringstream{};
+            auto kind_name            = tool_functions->tool_get_domain_name_fn(record->kind);
+            auto op_name =
+                tool_functions->tool_get_operation_name_fn(record->kind, record->operation);
+
+            tool::csv::scratch_memory_encoder::write_row(
+                scratch_memory_trace,
+                kind_name,
+                op_name,
+                tool_functions->tool_get_agent_node_id_fn(record->agent_id),
+                record->queue_id.handle,
+                record->thread_id,
+                record->flags,
+                record->start_timestamp,
+                record->end_timestamp);
+
+            if(tool::get_config().stats)
+                scratch_memory_stats[op_name] += (record->end_timestamp - record->start_timestamp);
+
+            (*ofs) << scratch_memory_trace.str();
+        }
+    }
+
+    if(tool::get_config().stats)
+    {
+        auto* stats_ofs = tool_functions->tool_get_scratch_memory_stats_file_fn();
+        if(!stats_ofs) throw std::runtime_error{"error creating scratch memory stats output file"};
+        write_stats(*stats_ofs, tool_functions->tool_get_app_timestamps_fn(), scratch_memory_stats);
+    }
+}
+}  // namespace tool
+}  // namespace rocprofiler
@@ -0,0 +1,57 @@
+// MIT License
+//
+// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#pragma once
+
+#include "helper.hpp"
+
+#include <rocprofiler-sdk/agent.h>
+
+namespace rocprofiler
+{
+namespace tool
+{
+void
+generate_csv(tool_table* tool_functions, std::vector<rocprofiler_agent_v0_t>& data);
+
+void
+generate_csv(tool_table* tool_functions, std::vector<kernel_dispatch_ring_buffer_t>& data);
+
+void
+generate_csv(tool_table* tool_functions, std::vector<hip_ring_buffer_t>& data);
+
+void
+generate_csv(tool_table* tool_functions, std::vector<hsa_ring_buffer_t>& data);
+
+void
+generate_csv(tool_table* tool_functions, std::vector<memory_copy_ring_buffer_t>& data);
+
+void
+generate_csv(tool_table* tool_functions, std::vector<marker_api_ring_buffer_t>& data);
+
+void
+generate_csv(tool_table* tool_functions, std::vector<counter_collection_ring_buffer_t>& data);
+
+void
+generate_csv(tool_table* tool_functions, std::vector<scratch_memory_ring_buffer_t>& data);
+}  // namespace tool
+}  // namespace rocprofiler
@@ -30,6 +30,7 @@
 #include <atomic>
 #include <iostream>
 #include <mutex>
+#include <tuple>
 #include <unordered_map>
 #include <unordered_set>
 #include <utility>
@@ -22,8 +22,11 @@

 #pragma once

+#include "lib/common/container/ring_buffer.hpp"
+#include "lib/common/container/small_vector.hpp"
 #include "lib/common/defines.hpp"
 #include "lib/common/filesystem.hpp"
+#include "output_file.hpp"

 #include <rocprofiler-sdk/registration.h>
 #include <rocprofiler-sdk/rocprofiler.h>
@@ -60,7 +63,8 @@
                rocprofiler_get_status_string(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__));        \
            std::cerr << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg            \
                      << " failed with error code " << ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) \
-                      << ": " << status_msg << std::endl;                                          \
+                      << ": " << status_msg << "\n"                                                \
+                      << std::flush;                                                               \
            std::stringstream errmsg{};                                                            \
            errmsg << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg " failure ("  \
                   << status_msg << ")";                                                           \
@@ -77,6 +81,9 @@ using rocprofiler_tool_buffer_kind_operation_names_t =
    std::unordered_map<rocprofiler_buffer_tracing_kind_t,
                       std::unordered_map<uint32_t, const char*>>;

+namespace common = ::rocprofiler::common;
+namespace tool   = ::rocprofiler::tool;
+
 struct rocprofiler_tool_buffer_name_info_t
 {
    rocprofiler_tool_buffer_kind_names_t           kind_names      = {};
@@ -100,3 +107,109 @@ get_buffer_id_names();

 rocprofiler_tool_callback_name_info_t
 get_callback_id_names();
+
+struct rocprofiler_tool_marker_record_t
+{
+    rocprofiler_callback_tracing_kind_t kind;
+    uint32_t                            op    = 0;
+    uint32_t                            phase = 0;
+    uint64_t                            pid   = 0;
+    uint64_t                            tid   = 0;
+    uint64_t                            cid   = 0;
+
+    rocprofiler_timestamp_t start_timestamp;
+    rocprofiler_timestamp_t end_timestamp;
+};
+
+struct rocprofiler_tool_counter_collection_record_t
+{
+    rocprofiler_profile_counting_dispatch_data_t                       dispatch_data;
+    common::container::small_vector<rocprofiler_record_counter_t, 128> profiler_record;
+    uint64_t                                                           pid                  = 0;
+    uint64_t                                                           id                   = 0;
+    uint64_t                                                           thread_id            = 0;
+    uint64_t                                                           dispatch_index       = 0;
+    uint64_t                                                           private_segment_size = 0;
+    uint64_t                                                           arch_vgpr_count      = 0;
+    uint64_t                                                           sgpr_count           = 0;
+    uint64_t                                                           lds_block_size_v     = 0;
+};
+
+struct timestamps_t
+{
+    rocprofiler_timestamp_t app_start_time;
+    rocprofiler_timestamp_t app_end_time;
+};
+
+using hip_ring_buffer_t =
+    rocprofiler::common::container::ring_buffer<rocprofiler_buffer_tracing_hip_api_record_t>;
+using hsa_ring_buffer_t =
+    rocprofiler::common::container::ring_buffer<rocprofiler_buffer_tracing_hsa_api_record_t>;
+using kernel_dispatch_ring_buffer_t = rocprofiler::common::container::ring_buffer<
+    rocprofiler_buffer_tracing_kernel_dispatch_record_t>;
+using memory_copy_ring_buffer_t =
+    rocprofiler::common::container::ring_buffer<rocprofiler_buffer_tracing_memory_copy_record_t>;
+using counter_collection_buffer_t =
+    rocprofiler::common::container::ring_buffer<rocprofiler_record_counter_t>;
+using marker_api_ring_buffer_t =
+    rocprofiler::common::container::ring_buffer<rocprofiler_tool_marker_record_t>;
+using counter_collection_ring_buffer_t =
+    rocprofiler::common::container::ring_buffer<rocprofiler_tool_counter_collection_record_t>;
+using scratch_memory_ring_buffer_t =
+    rocprofiler::common::container::ring_buffer<rocprofiler_buffer_tracing_scratch_memory_record_t>;
+
+using tool_get_agent_node_id_fn_t      = uint64_t (*)(rocprofiler_agent_id_t);
+using tool_get_app_timestamps_fn_t     = timestamps_t* (*) ();
+using tool_get_kernel_name_fn_t        = std::string_view (*)(uint64_t);
+using tool_get_domain_name_fn_t        = std::string_view (*)(rocprofiler_buffer_tracing_kind_t);
+using tool_get_operation_name_fn_t     = std::string_view (*)(rocprofiler_buffer_tracing_kind_t,
+                                                          rocprofiler_tracing_operation_t);
+using tool_get_callback_kind_name_fn_t = std::string_view (*)(rocprofiler_callback_tracing_kind_t);
+using tool_get_callback_op_name_fn_t   = std::string_view (*)(rocprofiler_callback_tracing_kind_t,
+                                                            uint32_t);
+using tool_get_roctx_msg_fn_t          = std::string_view (*)(uint64_t);
+using tool_get_counter_info_name_fn_t  = std::string (*)(uint64_t);
+using tool_get_output_file_ref_fn_t    = rocprofiler::tool::output_file& (*) ();
+using tool_get_output_file_ptr_fn_t    = rocprofiler::tool::output_file*& (*) ();
+
+enum class buffer_type_t
+{
+    ROCPROFILER_TOOL_BUFFER_HSA = 0,
+    ROCPROFILER_TOOL_BUFFER_HIP,
+    ROCPROFILER_TOOL_BUFFER_MEMORY_COPY,
+    ROCPROFILER_TOOL_BUFFER_COUNTER_COLLECTION,
+    ROCPROFILER_TOOL_BUFFER_KERNEL_DISPATCH,
+    ROCPROFILER_TOOL_BUFFER_MARKER_API,
+    ROCPROFILER_TOOL_BUFFER_SCRATCH_MEMORY,
+};
+
+struct tool_table
+{
+    // node id
+    tool_get_agent_node_id_fn_t tool_get_agent_node_id_fn = nullptr;
+    // timestamps
+    tool_get_app_timestamps_fn_t tool_get_app_timestamps_fn = nullptr;
+    // names and messages
+    tool_get_kernel_name_fn_t        tool_get_kernel_name_fn       = nullptr;
+    tool_get_domain_name_fn_t        tool_get_domain_name_fn       = nullptr;
+    tool_get_operation_name_fn_t     tool_get_operation_name_fn    = nullptr;
+    tool_get_counter_info_name_fn_t  tool_get_counter_info_name_fn = nullptr;
+    tool_get_callback_kind_name_fn_t tool_get_callback_kind_fn     = nullptr;
+    tool_get_callback_op_name_fn_t   tool_get_callback_op_name_fn  = nullptr;
+    tool_get_roctx_msg_fn_t          tool_get_roctx_msg_fn         = nullptr;
+    // trace files
+    tool_get_output_file_ref_fn_t tool_get_agent_info_file_fn         = nullptr;
+    tool_get_output_file_ref_fn_t tool_get_kernel_trace_file_fn       = nullptr;
+    tool_get_output_file_ref_fn_t tool_get_hsa_api_trace_file_fn      = nullptr;
+    tool_get_output_file_ref_fn_t tool_get_hip_api_trace_file_fn      = nullptr;
+    tool_get_output_file_ref_fn_t tool_get_marker_api_trace_file_fn   = nullptr;
+    tool_get_output_file_ref_fn_t tool_get_counter_collection_file_fn = nullptr;
+    tool_get_output_file_ref_fn_t tool_get_memory_copy_trace_file_fn  = nullptr;
+    tool_get_output_file_ptr_fn_t tool_get_scratch_memory_file_fn     = nullptr;
+    // stats files
+    tool_get_output_file_ref_fn_t tool_get_kernel_stats_file_fn         = nullptr;
+    tool_get_output_file_ref_fn_t tool_get_hip_stats_file_fn            = nullptr;
+    tool_get_output_file_ref_fn_t tool_get_hsa_stats_file_fn            = nullptr;
+    tool_get_output_file_ref_fn_t tool_get_memory_copy_stats_file_fn    = nullptr;
+    tool_get_output_file_ptr_fn_t tool_get_scratch_memory_stats_file_fn = nullptr;
+};
@@ -0,0 +1,231 @@
+// MIT License
+//
+// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#pragma once
+
+#include <cmath>
+#include <cstdint>
+#include <fstream>
+#include <functional>
+#include <iomanip>
+#include <iostream>
+#include <limits>
+#include <type_traits>
+
+namespace rocprofiler
+{
+namespace tool
+{
+/// \struct statistics
+/// \tparam Tp data type for statistical accumulation
+/// \tparam Fp floating point data type to use for division
+/// \brief A generic class for statistical accumulation.
+///
+template <typename Tp, typename Fp = double>
+struct statistics
+{
+public:
+    using value_type = Tp;
+    using float_type = Fp;
+    using this_type  = statistics<Tp, Fp>;
+    static_assert(std::is_arithmetic<Tp>::value, "only supports arithmetic types");
+
+public:
+    inline statistics()                      = default;
+    inline ~statistics()                     = default;
+    inline statistics(const statistics&)     = default;
+    inline statistics(statistics&&) noexcept = default;
+    inline statistics& operator=(const statistics&) = default;
+    inline statistics& operator=(statistics&&) noexcept = default;
+
+    explicit statistics(value_type val)
+    : m_cnt(1)
+    , m_sum(val)
+    , m_sqr(val * val)
+    , m_min(val)
+    , m_max(val)
+    {}
+
+    statistics& operator=(value_type val)
+    {
+        m_cnt = 1;
+        m_sum = val;
+        m_min = val;
+        m_max = val;
+        m_sqr = (val * val);
+        return *this;
+    }
+
+public:
+    // Accumulated values
+    inline int64_t    get_count() const { return m_cnt; }
+    inline value_type get_min() const { return m_min; }
+    inline value_type get_max() const { return m_max; }
+    inline value_type get_sum() const { return m_sum; }
+    inline value_type get_sqr() const { return m_sqr; }
+    inline float_type get_mean() const { return static_cast<float_type>(m_sum) / m_cnt; }
+    inline float_type get_variance() const
+    {
+        if(m_cnt < 2) return (m_sum - m_sum);
+
+        auto _sum_of_squared_samples = m_sqr;
+        auto _sum_squared_mean       = (m_sum * m_sum) / static_cast<float_type>(m_cnt);
+        return (_sum_of_squared_samples - _sum_squared_mean) / static_cast<float_type>(m_cnt - 1);
+    }
+
+    inline float_type get_stddev() const { return ::std::sqrt(::std::abs(get_variance())); }
+
+    // Modifications
+    inline void reset()
+    {
+        m_cnt = 0;
+        m_sum = value_type{};
+        m_sqr = value_type{};
+        m_min = value_type{};
+        m_max = value_type{};
+    }
+
+public:
+    // Operators (value_type)
+    inline statistics& operator+=(value_type val)
+    {
+        if(m_cnt == 0)
+        {
+            m_sum = val;
+            m_sqr = (val * val);
+            m_min = val;
+            m_max = val;
+        }
+        else
+        {
+            m_sum += val;
+            m_sqr += (val * val);
+            m_min = ::std::min(m_min, val);
+            m_max = ::std::max(m_max, val);
+        }
+        ++m_cnt;
+
+        return *this;
+    }
+
+    inline statistics& operator-=(value_type val)
+    {
+        if(m_cnt > 1) --m_cnt;
+        m_sum -= val;
+        m_sqr -= (val * val);
+        m_min -= val;
+        m_max -= val;
+        return *this;
+    }
+
+    inline statistics& operator*=(value_type val)
+    {
+        m_sum *= val;
+        m_sqr *= (val * val);
+        m_min *= val;
+        m_max *= val;
+        return *this;
+    }
+
+    inline statistics& operator/=(value_type val)
+    {
+        m_sum /= val;
+        m_sqr /= (val * val);
+        m_min /= val;
+        m_max /= val;
+        return *this;
+    }
+
+public:
+    // Operators (this_type)
+    inline statistics& operator+=(const statistics& rhs)
+    {
+        if(m_cnt == 0)
+        {
+            m_sum = rhs.m_sum;
+            m_sqr = rhs.m_sqr;
+            m_min = rhs.m_min;
+            m_max = rhs.m_max;
+        }
+        else
+        {
+            m_sum += rhs.m_sum;
+            m_sqr += rhs.m_sqr;
+            m_min = ::std::min(m_min, rhs.m_min);
+            m_max = ::std::max(m_max, rhs.m_max);
+        }
+        m_cnt += rhs.m_cnt;
+        return *this;
+    }
+
+    // Operators (this_type)
+    inline statistics& operator-=(const statistics& rhs)
+    {
+        if(m_cnt > 0)
+        {
+            m_sum -= rhs.m_sum;
+            m_sqr -= rhs.m_sqr;
+            m_min = ::std::min(m_min, rhs.m_min);
+            m_max = ::std::max(m_max, rhs.m_max);
+        }
+        return *this;
+    }
+
+private:
+    // summation of each history^1
+    int64_t    m_cnt = 0;
+    value_type m_sum = value_type{};
+    value_type m_sqr = value_type{};
+    value_type m_min = value_type{};
+    value_type m_max = value_type{};
+
+public:
+    // friend operator for addition
+    friend statistics operator+(const statistics& lhs, const statistics& rhs)
+    {
+        return statistics(lhs) += rhs;
+    }
+
+    friend statistics operator-(const statistics& lhs, const statistics& rhs)
+    {
+        return statistics(lhs) -= rhs;
+    }
+};
+}  // namespace tool
+}  // namespace rocprofiler
+
+namespace std
+{
+template <typename Tp>
+::rocprofiler::tool::statistics<Tp>
+max(::rocprofiler::tool::statistics<Tp> lhs, const Tp& rhs)
+{
+    return lhs.get_max(rhs);
+}
+
+template <typename Tp>
+::rocprofiler::tool::statistics<Tp>
+min(::rocprofiler::tool::statistics<Tp> lhs, const Tp& rhs)
+{
+    return lhs.get_min(rhs);
+}
+}  // namespace std
@@ -0,0 +1,131 @@
+// MIT License
+//
+// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#include "tmp_file.hpp"
+#include "config.hpp"
+
+#include "lib/common/filesystem.hpp"
+
+namespace fs = ::rocprofiler::common::filesystem;
+
+bool
+tmp_file::fopen(const char* _mode)
+{
+    if(!fs::exists(filename))
+    {
+        // if the filepath does not exist, open in out mode to create it
+        std::ofstream _ofs{filename};
+    }
+    file = std::fopen(filename.c_str(), _mode);
+    if(file) fd = ::fileno(file);
+
+    return (file != nullptr && fd > 0);
+}
+
+tmp_file::tmp_file(std::string _filename)
+: filename(std::move(_filename))
+{}
+
+tmp_file::~tmp_file()
+{
+    close();
+    remove();
+}
+
+bool
+tmp_file::flush()
+{
+    if(stream.is_open())
+    {
+        stream.flush();
+    }
+    else if(file != nullptr)
+    {
+        int _ret = fflush(file);
+        int _cnt = 0;
+        while(_ret == EAGAIN || _ret == EINTR)
+        {
+            // std::this_thread::sleep_for(std::chrono::milliseconds{ 100 });
+            _ret = fflush(file);
+            if(++_cnt > 10) break;
+        }
+        return (_ret == 0);
+    }
+
+    return true;
+}
+
+bool
+tmp_file::close()
+{
+    flush();
+
+    if(stream.is_open())
+    {
+        stream.close();
+        return !stream.is_open();
+    }
+    else if(file != nullptr)
+    {
+        auto _ret = fclose(file);
+        if(_ret == 0)
+        {
+            file = nullptr;
+            fd   = -1;
+        }
+        return (_ret == 0);
+    }
+
+    return true;
+}
+
+bool
+tmp_file::open(std::ios::openmode _mode)
+{
+    if(!fs::exists(filename))
+    {
+        // if the filepath does not exist, open in out mode to create it
+        std::ofstream _ofs{};
+        _ofs.open(filename, std::ofstream::binary | std::ofstream::out);
+    }
+
+    stream.open(filename, _mode);
+    return (stream.is_open() && stream.good());
+}
+
+bool
+tmp_file::remove()
+{
+    close();
+    if(fs::exists(filename))
+    {
+        auto _ret = ::remove(filename.c_str());
+        return (_ret == 0);
+    }
+
+    return true;
+}
+
+tmp_file::operator bool() const
+{
+    return (stream.is_open() && stream.good()) || (file != nullptr && fd > 0);
+}
@@ -0,0 +1,51 @@
+// MIT License
+//
+// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#pragma once
+
+#include <fstream>
+#include <ios>
+#include <mutex>
+#include <set>
+#include <string>
+
+struct tmp_file
+{
+    tmp_file(std::string _filename);
+    ~tmp_file();
+
+    bool open(std::ios::openmode = std::ios::binary | std::ios::in | std::ios::out);
+    bool fopen(const char* = "r+");
+    bool flush();
+    bool close();
+    bool remove();
+
+    explicit operator bool() const;
+
+    std::string              filename     = {};
+    std::string              subdirectory = {};
+    std::fstream             stream       = {};
+    FILE*                    file         = nullptr;
+    int                      fd           = -1;
+    std::set<std::streampos> file_pos     = {};
+    std::mutex               file_mutex   = {};
+};
@@ -253,7 +253,7 @@ hip_api_impl<TableIdx, OpIdx>::functor(Args... args)

    if(callback_contexts.empty() && buffered_contexts.empty())
    {
-        auto _ret = exec(info_type::get_table_func(), std::forward<Args>(args)...);
+        [[maybe_unused]] auto _ret = exec(info_type::get_table_func(), std::forward<Args>(args)...);
        if constexpr(!std::is_void<RetT>::value)
            return _ret;
        else
@@ -31,12 +31,13 @@ def test_validate_counter_collection_pmc2(input_dir: pd.DataFrame):
            file_path = os.path.join(subdirectory_path, file_name)
            assert os.path.isfile(file_path), f"{file_path} is not a file."

-            with open(file_path, "r") as file:
-                df = pd.read_csv(file)
-                # check if kernel-name is present
-                assert len(df["Kernel_Name"]) > 0
-                # check if counter value is positive
-                assert len(df["Counter_Value"]) > 0
+            if "agent_info.csv" not in file_path:
+                with open(file_path, "r") as file:
+                    df = pd.read_csv(file)
+                    # check if kernel-name is present
+                    assert len(df["Kernel_Name"]) > 0
+                    # check if counter value is positive
+                    assert len(df["Counter_Value"]) > 0


 if __name__ == "__main__":
@@ -15,7 +15,7 @@ add_test(
    COMMAND
        $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --hip-runtime-trace
        --hip-compiler-trace --hsa-core-trace --hsa-amd-trace --hsa-image-trace
-        --hsa-finalizer-trace --kernel-trace --memory-copy-trace -d
+        --hsa-finalizer-trace --kernel-trace --memory-copy-trace --stats -d
        ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out $<TARGET_FILE:hip-in-libraries>)

 string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV
@@ -51,13 +51,26 @@ add_test(
        --kernel-input
        ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_trace.csv
        --memory-copy-input
-        ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_trace.csv)
+        ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_trace.csv
+        --agent-input
+        ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_agent_info.csv
+        --kernel-stats
+        ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_stats.csv
+        --hip-stats ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_stats.csv
+        --hsa-stats ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_stats.csv
+        --memory-copy-stats
+        ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv)

 set(VALIDATION_FILES
    ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_trace.csv
    ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_api_trace.csv
    ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_api_trace.csv
-    ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_trace.csv)
+    ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_trace.csv
+    ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_agent_info.csv
+    ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_stats.csv
+    ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_stats.csv
+    ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_stats.csv
+    ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv)

 set_tests_properties(
    rocprofv3-test-trace-hip-in-libraries-validate
@@ -6,6 +6,11 @@ import pytest


 def pytest_addoption(parser):
+    parser.addoption(
+        "--agent-input",
+        action="store",
+        help="Path to agent info CSV file.",
+    )
    parser.addoption(
        "--hsa-input",
        action="store",
@@ -31,6 +36,38 @@ def pytest_addoption(parser):
        action="store",
        help="Path to HIP runtime and compiler API tracing CSV file.",
    )
+    parser.addoption(
+        "--hip-stats",
+        action="store",
+        help="Path to HIP stats CSV file.",
+    )
+    parser.addoption(
+        "--hsa-stats",
+        action="store",
+        help="Path to HSA stats CSV file.",
+    )
+    parser.addoption(
+        "--kernel-stats",
+        action="store",
+        help="Path to kernel stats CSV file.",
+    )
+    parser.addoption(
+        "--memory-copy-stats",
+        action="store",
+        help="Path to memory copy stats CSV file.",
+    )
+
+
+@pytest.fixture
+def agent_info_input_data(request):
+    filename = request.config.getoption("--agent-input")
+    data = []
+    with open(filename, "r") as inp:
+        reader = csv.DictReader(inp)
+        for row in reader:
+            data.append(row)
+
+    return data


@pytest.fixture
@@ -92,3 +129,55 @@ def hip_input_data(request):
                data.append(row)

    return data
+
+
+@pytest.fixture
+def hip_stats_data(request):
+    filename = request.config.getoption("--hip-stats")
+    data = []
+    if os.path.exists(filename):
+        with open(filename, "r") as inp:
+            reader = csv.DictReader(inp)
+            for row in reader:
+                data.append(row)
+
+    return data
+
+
+@pytest.fixture
+def hsa_stats_data(request):
+    filename = request.config.getoption("--hsa-stats")
+    data = []
+    if os.path.exists(filename):
+        with open(filename, "r") as inp:
+            reader = csv.DictReader(inp)
+            for row in reader:
+                data.append(row)
+
+    return data
+
+
+@pytest.fixture
+def kernel_stats_data(request):
+    filename = request.config.getoption("--kernel-stats")
+    data = []
+    if os.path.exists(filename):
+        with open(filename, "r") as inp:
+            reader = csv.DictReader(inp)
+            for row in reader:
+                data.append(row)
+
+    return data
+
+
+@pytest.fixture
+def memory_copy_stats_data(request):
+    filename = request.config.getoption("--memory-copy-stats")
+    data = []
+    if os.path.exists(filename):
+        with open(filename, "r") as inp:
+            reader = csv.DictReader(inp)
+            for row in reader:
+                data.append(row)
+
+    return data
@@ -3,7 +3,6 @@
 import re
 import sys
 import pytest
-import subprocess


 class dim3(object):
@@ -16,7 +15,29 @@ class dim3(object):
        return (self.x, self.y, self.z)


-def test_api_trace(hsa_input_data, hip_input_data):
+def validate_average(row):
+    avg_ns = float(row["AverageNs"])
+    tot_ns = int(row["TotalDurationNs"])
+    num_cnt = float(row["Calls"])
+    assert abs(avg_ns - (tot_ns / num_cnt)) < 1.0e-3, f"{row}"
+
+
+def validate_stats(row):
+    min_v = int(row["MinNs"])
+    max_v = int(row["MaxNs"])
+    avg_v = float(row["AverageNs"])
+    cnt_v = int(row["Calls"])
+    stddev_v = float(row["StdDev"])
+
+    assert min_v > 0, f"{row}"
+    assert max_v > 0, f"{row}"
+    assert min_v < max_v if cnt_v > 1 else min_v == max_v, f"{row}"
+    assert min_v < avg_v if cnt_v > 1 else min_v == int(avg_v), f"{row}"
+    assert max_v > avg_v if cnt_v > 1 else max_v == int(avg_v), f"{row}"
+    assert stddev_v > 0.0 if cnt_v > 1 else int(stddev_v) == 0, f"{row}"
+
+
+def test_api_trace(hsa_input_data, hip_input_data, hip_stats_data):
    functions = []
    correlation_ids = []
    for row in hsa_input_data:
@@ -77,8 +98,15 @@ def test_api_trace(hsa_input_data, hip_input_data):
        ):
            assert itr in functions

+    for row in hip_stats_data:
+        assert int(row["TotalDurationNs"]) > 0
+        assert int(row["Calls"]) > 0
+        validate_average(row)
+        assert float(row["Percentage"]) > 0.0
+        validate_stats(row)

-def test_kernel_trace(kernel_input_data):
+
+def test_kernel_trace(kernel_input_data, kernel_stats_data):
    valid_kernel_names = sorted(
        [
            "(anonymous namespace)::transpose(int const*, int*, int, int)",
@@ -125,27 +153,42 @@ def test_kernel_trace(kernel_input_data):
    kernels = sorted(list(set(kernels)))
    assert kernels == valid_kernel_names

+    for row in kernel_stats_data:
+        assert int(row["TotalDurationNs"]) > 0
+        assert int(row["Calls"]) > 0
+        validate_average(row)
+        assert float(row["Percentage"]) > 0.0
+        validate_stats(row)
+
+
+def test_memory_copy_trace(
+    agent_info_input_data,
+    memory_copy_input_data,
+    hsa_input_data,
+    hsa_stats_data,
+    memory_copy_stats_data,
+):
+    def get_agent(node_id):
+        for row in agent_info_input_data:
+            if row["Logical_Node_Id"] == node_id:
+                return row
+        return None

-def test_memory_copy_trace(memory_copy_input_data, hsa_input_data):
    for row in memory_copy_input_data:
        assert row["Kind"] == "MEMORY_COPY"
        assert row["Direction"] in ("HOST_TO_DEVICE", "DEVICE_TO_HOST")
+
+        src_agent = get_agent(row["Source_Agent_Id"])
+        dst_agent = get_agent(row["Destination_Agent_Id"])
+        assert src_agent is not None and dst_agent is not None, f"{agent_info_input_data}"
+
        if row["Direction"] == "HOST_TO_DEVICE":
-            output = subprocess.check_output(
-                'rocminfo | grep "Node: *'
-                + row["Source_Agent_Id"]
-                + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'',
-                shell=True,
-            )
-            assert int(str(output).find("CPU")) >= 0
+            assert src_agent["Agent_Type"] == "CPU"
+            assert dst_agent["Agent_Type"] == "GPU"
        elif row["Direction"] == "DEVICE_TO_HOST":
-            output = subprocess.check_output(
-                'rocminfo | grep "Node: *'
-                + row["Destination_Agent_Id"]
-                + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'',
-                shell=True,
-            )
-            assert int(str(output).find("CPU")) >= 0
+            assert src_agent["Agent_Type"] == "GPU"
+            assert dst_agent["Agent_Type"] == "CPU"
+
        assert int(row["Correlation_Id"]) > 0
        assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"])

@@ -155,6 +198,20 @@ def test_memory_copy_trace(memory_copy_input_data, hsa_input_data):
            valid_length += 1
    assert len(memory_copy_input_data) == valid_length

+    for row in hsa_stats_data:
+        assert int(row["TotalDurationNs"]) > 0
+        assert int(row["Calls"]) > 0
+        validate_average(row)
+        assert float(row["Percentage"]) > 0.0
+        validate_stats(row)
+
+    for row in memory_copy_stats_data:
+        assert int(row["TotalDurationNs"]) > 0
+        assert int(row["Calls"]) > 0
+        validate_average(row)
+        assert float(row["Percentage"]) > 0.0
+        validate_stats(row)
+

 if __name__ == "__main__":
    exit_code = pytest.main(["-x", __file__] + sys.argv[1:])
@@ -31,10 +31,11 @@ def test_validate_counter_collection_plus_tracing(input_dir: pd.DataFrame):
            file_path = os.path.join(subdirectory_path, file_name)
            assert os.path.isfile(file_path), f"{file_path} is not a file."

-            with open(file_path, "r") as file:
-                df = pd.read_csv(file)
-                # check if either kernel-name/FUNCTION is present
-                assert "Kernel_Name" in df.columns or "Function" in df.columns
+            if "agent_info.csv" not in file_path:
+                with open(file_path, "r") as file:
+                    df = pd.read_csv(file)
+                    # check if either kernel-name/FUNCTION is present
+                    assert "Kernel_Name" in df.columns or "Function" in df.columns


 if __name__ == "__main__":
@@ -56,13 +56,16 @@ add_test(
        --memory-copy-input
        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_memory_copy_trace.csv
        --marker-input
-        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_marker_api_trace.csv)
+        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_marker_api_trace.csv
+        --agent-input
+        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_agent_info.csv)

 set(VALIDATION_FILES
    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_memory_copy_trace.csv
    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_hsa_api_trace.csv
    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_kernel_trace.csv
-    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_marker_api_trace.csv)
+    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_marker_api_trace.csv
+    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_agent_info.csv)

 set_tests_properties(
    rocprofv3-test-trace-validate
@@ -111,13 +114,16 @@ add_test(
        --memory-copy-input
        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_memory_copy_trace.csv
        --marker-input
-        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_marker_api_trace.csv)
+        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_marker_api_trace.csv
+        --agent-input
+        ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_agent_info.csv)

 set(SYS_VALIDATION_FILES
    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_memory_copy_trace.csv
    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_hsa_api_trace.csv
    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_kernel_trace.csv
-    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_marker_api_trace.csv)
+    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_marker_api_trace.csv
+    ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_agent_info.csv)

 set_tests_properties(
    rocprofv3-test-systrace-validate
@@ -5,6 +5,11 @@ import pytest


 def pytest_addoption(parser):
+    parser.addoption(
+        "--agent-input",
+        action="store",
+        help="Path to agent info CSV file.",
+    )
    parser.addoption(
        "--hsa-input",
        action="store",
@@ -32,6 +37,18 @@ def pytest_addoption(parser):
    )


+@pytest.fixture
+def agent_info_input_data(request):
+    filename = request.config.getoption("--agent-input")
+    data = []
+    with open(filename, "r") as inp:
+        reader = csv.DictReader(inp)
+        for row in reader:
+            data.append(row)
+
+    return data
+
+
@pytest.fixture
 def hsa_input_data(request):
    filename = request.config.getoption("--hsa-input")
@@ -1,10 +1,27 @@
 #!/usr/bin/env python3

 import sys
-import subprocess
 import pytest


+def test_agent_info(agent_info_input_data):
+    logical_node_id = max([int(itr["Logical_Node_Id"]) for itr in agent_info_input_data])
+
+    assert logical_node_id + 1 == len(agent_info_input_data)
+
+    for row in agent_info_input_data:
+        agent_type = row["Agent_Type"]
+        assert agent_type in ("CPU", "GPU")
+        if agent_type == "CPU":
+            assert int(row["Cpu_Cores_Count"]) > 0
+            assert int(row["Simd_Count"]) == 0
+            assert int(row["Max_Waves_Per_Simd"]) == 0
+        else:
+            assert int(row["Cpu_Cores_Count"]) == 0
+            assert int(row["Simd_Count"]) > 0
+            assert int(row["Max_Waves_Per_Simd"]) > 0
+
+
 def test_hsa_api_trace(hsa_input_data):
    functions = []
    correlation_ids = []
@@ -60,49 +77,37 @@ def test_kernel_trace(kernel_input_data):
        assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"])


-def test_memory_copy_trace(memory_copy_input_data):
+def test_memory_copy_trace(agent_info_input_data, memory_copy_input_data):
+
+    def get_agent(node_id):
+        for row in agent_info_input_data:
+            if row["Logical_Node_Id"] == node_id:
+                return row
+        return None
+
    for row in memory_copy_input_data:
        assert row["Kind"] == "MEMORY_COPY"

    assert len(memory_copy_input_data) == 2

-    row = memory_copy_input_data[0]
-    assert row["Direction"] == "HOST_TO_DEVICE"
-    output = subprocess.check_output(
-        'rocminfo | grep "Node: *'
-        + row["Source_Agent_Id"]
-        + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'',
-        shell=True,
-    )
-    assert int(str(output).find("CPU")) >= 0
-    output = subprocess.check_output(
-        'rocminfo | grep "Node: *'
-        + row["Destination_Agent_Id"]
-        + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'',
-        shell=True,
-    )
-    assert int(str(output).find("GPU")) >= 0
-    assert int(row["Correlation_Id"]) > 0
-    assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"])
+    def test_row(idx, direction):
+        assert direction in ("HOST_TO_DEVICE", "DEVICE_TO_HOST")
+        row = memory_copy_input_data[idx]
+        assert row["Direction"] == direction
+        src_agent = get_agent(row["Source_Agent_Id"])
+        dst_agent = get_agent(row["Destination_Agent_Id"])
+        assert src_agent is not None and dst_agent is not None, f"{agent_info_input_data}"
+        if direction == "HOST_TO_DEVICE":
+            assert src_agent["Agent_Type"] == "CPU"
+            assert dst_agent["Agent_Type"] == "GPU"
+        else:
+            assert src_agent["Agent_Type"] == "GPU"
+            assert dst_agent["Agent_Type"] == "CPU"
+        assert int(row["Correlation_Id"]) > 0
+        assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"])

-    row = memory_copy_input_data[1]
-    assert row["Direction"] == "DEVICE_TO_HOST"
-    output = subprocess.check_output(
-        'rocminfo | grep "Node: *'
-        + row["Source_Agent_Id"]
-        + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'',
-        shell=True,
-    )
-    assert int(str(output).find("GPU")) >= 0
-    output = subprocess.check_output(
-        'rocminfo | grep "Node: *'
-        + row["Destination_Agent_Id"]
-        + '" -A 1 | grep "Device Type" | sed \'s/.*: *//\'',
-        shell=True,
-    )
-    assert int(str(output).find("CPU")) >= 0
-    assert int(row["Correlation_Id"]) > 0
-    assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"])
+    test_row(0, "HOST_TO_DEVICE")
+    test_row(1, "DEVICE_TO_HOST")


 def test_marker_api_trace(marker_input_data):