Perfetto traces from cached data (#1704)

## Motivation The idea is to unify the way and place where we store our traces. Current implementation uses `trace_cache` for rocpd traces, but perfetto is in lined inside of each module. This change allows us to have a single point in code where we will collect data, process it and store it in the desired format. This means that we can declutter the code further and have single point of responsibility and single point of failure. ## Technical Details New `processor` (perfetto_post_processing.cpp) is added to the `trace_cache` which purpose is to use the cached data to populate perfetto tracks. Cache manager is responsible for keeping the instance of this processor and for its lifetime.
2025-12-01 15:59:16 +01:00
parent b506c75f28
commit 3b11e01716
22 changed files with 1889 additions and 178 deletions
@@ -339,8 +339,8 @@ generate_config(std::string _config_file, const std::set<std::string>& _config_f
                if(_lomni && !_romni) return true;
                if(_romni && !_lomni) return false;
                for(const auto* itr :
-                    { "ROCPROFSYS_CONFIG", "ROCPROFSYS_MODE", "ROCPROFSYS_TRACE",
-                      "ROCPROFSYS_PROFILE", "ROCPROFSYS_USE_SAMPLING",
+                    { "ROCPROFSYS_CONFIG", "ROCPROFSYS_MODE", "ROCPROFSYS_TRACE_CACHED",
+                      "ROCPROFSYS_TRACE", "ROCPROFSYS_PROFILE", "ROCPROFSYS_USE_SAMPLING",
                      "ROCPROFSYS_USE_PROCESS_SAMPLING", "ROCPROFSYS_USE_ROCM",
                      "ROCPROFSYS_USE_AMD_SMI", "ROCPROFSYS_USE_KOKKOSP",
                      "ROCPROFSYS_USE_OMPT", "ROCPROFSYS_USE", "ROCPROFSYS_OUTPUT" })
@@ -185,9 +185,10 @@ get_initial_environment()
    update_env(_env, "ROCPROFSYS_MODE", "causal");
    update_env(_env, "ROCPROFSYS_USE_CAUSAL", true);
    update_env(_env, "ROCPROFSYS_USE_SAMPLING", false);
-    update_env(_env, "ROCPROFSYS_TRACE", false);
+    update_env(_env, "ROCPROFSYS_TRACE_CACHED", false);
    update_env(_env, "ROCPROFSYS_PROFILE", false);
    update_env(_env, "ROCPROFSYS_USE_PROCESS_SAMPLING", false);
+    update_env(_env, "ROCPROFSYS_TRACE", false);
    update_env(_env, "ROCPROFSYS_THREAD_POOL_SIZE",
               get_env<int>("ROCPROFSYS_THREAD_POOL_SIZE", 0));
    update_env(_env, "ROCPROFSYS_LAUNCHER", "rocprof-sys-causal");
@@ -378,6 +378,13 @@ parse_args(int argc, char** argv, std::vector<char*>& _env)
        .action([&](parser_t& p) {
            update_env(_env, "ROCPROFSYS_TRACE", p.get<bool>("trace"));
        });
+    parser
+        .add_argument({ "--trace-cached" },
+                      "Generate a detailed trace (perfetto output) from cached data ")
+        .max_count(1)
+        .action([&](parser_t& p) {
+            update_env(_env, "ROCPROFSYS_TRACE_CACHED", p.get<bool>("trace-cached"));
+        });
    parser
        .add_argument(
            { "-P", "--profile" },
@@ -60,6 +60,7 @@ set(core_headers
    ${CMAKE_CURRENT_LIST_DIR}/argparse.hpp
    ${CMAKE_CURRENT_LIST_DIR}/categories.hpp
    ${CMAKE_CURRENT_LIST_DIR}/common.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/common_types.hpp
    ${CMAKE_CURRENT_LIST_DIR}/concepts.hpp
    ${CMAKE_CURRENT_LIST_DIR}/config.hpp
    ${CMAKE_CURRENT_LIST_DIR}/constraint.hpp
@@ -0,0 +1,100 @@
+// MIT License
+//
+// Copyright (c) 2025 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 <cstdint>
+#include <string>
+#include <vector>
+
+namespace rocprofsys
+{
+/**
+ * @brief Structure for function/region argument information
+ *
+ * Represents metadata about function or region arguments including
+ * their position, type, name, and value information.
+ */
+struct argument_info
+{
+    uint32_t    arg_number = 0;   ///< Argument position/index
+    std::string arg_type   = {};  ///< Argument type (e.g., "int", "float*")
+    std::string arg_name   = {};  ///< Argument name
+    std::string arg_value  = {};  ///< Argument value as string
+};
+
+using function_args_t = std::vector<argument_info>;
+
+inline std::string
+get_args_string(const function_args_t& args)
+{
+    std::string args_str;
+    std::for_each(args.begin(), args.end(), [&args_str](const argument_info& arg) {
+        const auto*       delimiter = ";;";
+        std::stringstream ss;
+        ss << arg.arg_number << delimiter << arg.arg_type << delimiter << arg.arg_name
+           << delimiter << arg.arg_value << delimiter;
+        args_str.append(ss.str());
+    });
+    return args_str;
+}
+
+inline function_args_t
+process_arguments_string(const std::string& arg_str)
+{
+    function_args_t   args;
+    const std::string delimiter = ";;";
+
+    auto split = [](const std::string& str, const std::string& _delimiter) {
+        std::vector<std::string> tokens;
+        size_t                   start = 0;
+        size_t                   end   = str.find(_delimiter);
+
+        while(end != std::string::npos)
+        {
+            tokens.push_back(str.substr(start, end - start));
+            start = end + _delimiter.length();
+            end   = str.find(_delimiter, start);
+        }
+
+        return tokens;
+    };
+
+    auto tokens = split(arg_str, delimiter);
+
+    // Ensure the number of tokens is a multiple of 4
+    if(tokens.size() % 4 != 0)
+    {
+        throw std::invalid_argument("Malformed argument string.");
+    }
+
+    for(auto it = tokens.begin(); it != tokens.end(); it += 4)
+    {
+        argument_info arg = { static_cast<uint32_t>(std::stoi(*it)), *(it + 1), *(it + 2),
+                              *(it + 3) };
+        args.push_back(arg);
+    }
+
+    return args;
+}
+
+}  // namespace rocprofsys
@@ -321,6 +321,10 @@ configure_settings(bool _init)
    ROCPROFSYS_CONFIG_SETTING(bool, "ROCPROFSYS_USE_ROCPD", "Enable rocpd backend", false,
                              "backend", "rocpd");

+    ROCPROFSYS_CONFIG_SETTING(bool, "ROCPROFSYS_TRACE_CACHED",
+                              "Enable perfetto with trace cache", false, "backend",
+                              "perfetto_caching");
+
    ROCPROFSYS_CONFIG_SETTING(bool, "ROCPROFSYS_USE_ROCM",
                              "Enable ROCm API and kernel tracing", true, "backend",
                              "rocm");
@@ -711,6 +715,11 @@ configure_settings(bool _init)
        "written to a file and re-loaded during finalization",
        true, "io", "data", "advanced");

+    ROCPROFSYS_CONFIG_SETTING(bool, "ROCPROFSYS_MERGE_PERFETTO_FILES",
+                              "Merge Perfetto traces. If not explicitly set, it will "
+                              "default to the value of ROCPROFSYS_COLLAPSE_PROCESSES",
+                              false, "perfetto", "data", "advanced");
+
    ROCPROFSYS_CONFIG_SETTING(
        std::string, "ROCPROFSYS_TMPDIR", "Base directory for temporary files",
        get_env<std::string>("TMPDIR", "/tmp"), "io", "data", "advanced");
@@ -2370,6 +2379,13 @@ get_use_tmp_files()
    return static_cast<tim::tsettings<bool>&>(*_v->second).get();
 }

+bool
+get_merge_perfetto_files()
+{
+    static auto _v = get_config()->find("ROCPROFSYS_MERGE_PERFETTO_FILES");
+    return static_cast<tim::tsettings<bool>&>(*_v->second).get();
+}
+
 std::string
 get_tmpdir()
 {
@@ -2404,6 +2420,51 @@ get_database_absolute_path(std::string_view database_name, std::string_view suff
    return _val;
 }

+std::string
+get_perfetto_output_filename_with_suffix(std::string_view suffix)
+{
+    static auto _v   = get_config()->find("ROCPROFSYS_PERFETTO_FILE");
+    auto        _val = static_cast<tim::tsettings<std::string>&>(*_v->second).get();
+
+    // If absolute path is provided, return it as-is
+    if(!_val.empty() && _val.at(0) == '/') return _val;
+
+    auto _pos_dir = _val.find_last_of('/');
+    auto _dir     = std::string{};
+    auto _ext     = std::string{ "proto" };
+
+    if(_pos_dir != std::string::npos)
+    {
+        _dir = _val.substr(0, _pos_dir + 1);
+        _val = _val.substr(_pos_dir + 1);
+    }
+
+    auto _pos_ext = _val.find_last_of('.');
+    if(_pos_ext + 1 < _val.length())
+    {
+        _ext = _val.substr(_pos_ext + 1);
+        _val = _val.substr(0, _pos_ext);
+    }
+
+    // Check if explicitly set via environment OR config file
+    // If explicitly set, don't add suffix; otherwise use provided suffix
+    bool _explicitly_set =
+        (_v->second->get_environ_updated() || _v->second->get_config_updated());
+
+    auto _cfg = settings::compose_filename_config{
+        !_explicitly_set && !suffix.empty(),  // use_suffix only if not explicitly set
+        suffix,                               // suffix value
+        true,                                 // make_dir
+        _dir                                  // explicit_path
+    };
+
+    _val = settings::compose_output_filename(_val, _ext, _cfg);
+    if(!_val.empty() && _val.at(0) != '/')
+        return settings::format(JOIN('/', "%env{PWD}%", _val), get_config()->get_tag());
+
+    return _val;
+}
+
 bool&
 get_use_rocpd()
 {
@@ -2411,6 +2472,13 @@ get_use_rocpd()
    return static_cast<tim::tsettings<bool>&>(*_v).get();
 }

+bool&
+get_caching_perfetto()
+{
+    static auto _v = get_config()->at("ROCPROFSYS_TRACE_CACHED");
+    return static_cast<tim::tsettings<bool>&>(*_v).get();
+}
+
 tmp_file::tmp_file(std::string _v)
 : filename{ std::move(_v) }
 {}
@@ -364,9 +364,18 @@ get_tmpdir();
 std::string
 get_database_absolute_path(std::string_view database_name, std::string_view tag);

+std::string
+get_perfetto_output_filename_with_suffix(std::string_view suffix = "");
+
 bool&
 get_use_rocpd() ROCPROFSYS_HOT;

+bool&
+get_caching_perfetto() ROCPROFSYS_HOT;
+
+bool
+get_merge_perfetto_files();
+
 struct tmp_file
 {
    tmp_file(std::string);
@@ -208,6 +208,8 @@ post_process(tim::manager* _timemory_manager, bool& _perfetto_output_error)
            _data = char_vec_t{ tracing_session->ReadTraceBlocking() };
        }

+        tracing_session.reset();
+
        return _data;
    };

@@ -25,6 +25,7 @@ set(trace_cache_sources
    ${CMAKE_CURRENT_LIST_DIR}/buffer_storage.cpp
    ${CMAKE_CURRENT_LIST_DIR}/metadata_registry.cpp
    ${CMAKE_CURRENT_LIST_DIR}/rocpd_processor.cpp
+    ${CMAKE_CURRENT_LIST_DIR}/perfetto_processor.cpp
 )

 set(trace_cache_headers
@@ -36,6 +37,7 @@ set(trace_cache_headers
    ${CMAKE_CURRENT_LIST_DIR}/cache_type_traits.hpp
    ${CMAKE_CURRENT_LIST_DIR}/metadata_registry.hpp
    ${CMAKE_CURRENT_LIST_DIR}/rocpd_processor.hpp
+    ${CMAKE_CURRENT_LIST_DIR}/perfetto_processor.hpp
    ${CMAKE_CURRENT_LIST_DIR}/sample_processor.hpp
    ${CMAKE_CURRENT_LIST_DIR}/sample_type.hpp
 )
@@ -23,6 +23,7 @@
 #include "cache_manager.hpp"

 #include "core/trace_cache/metadata_registry.hpp"
+#include "core/trace_cache/perfetto_processor.hpp"
 #include "core/trace_cache/rocpd_processor.hpp"
 #include "core/trace_cache/sample_processor.hpp"

@@ -33,7 +34,9 @@
 #include "library/runtime.hpp"

 #include <algorithm>
+#include <cstring>
 #include <memory>
+#include <sstream>
 #include <vector>

 namespace rocprofsys
@@ -50,38 +53,128 @@ struct cache_files_t
    inline bool empty() const { return buff_storage.empty() || metadata.empty(); }
 };

+struct format_t
+{
+    bool        process_parallel;
+    bool        enabled;
+    const char* name;
+};
+
 struct enabled_formats_t
 {
-    bool rocpd = get_use_rocpd();
+    std::vector<format_t> formats = { { true, get_use_rocpd(), "rocpd" },
+                                      { false, get_caching_perfetto(), "perfetto" } };

    void print() const
    {
-        constexpr std::pair<const char*, bool enabled_formats_t::*> formats[] = {
-            { "rocpd", &enabled_formats_t::rocpd },
-        };
+        if(std::none_of(formats.begin(), formats.end(),
+                        [](const auto& f) { return f.enabled; }))
+            return;

-        bool any_enabled = false;
-        for(const auto& fmt : formats)
-            any_enabled |= this->*(fmt.second);
-
-        if(!any_enabled) return;
-
-        bool              first = true;
        std::stringstream ss;
+        bool              first = true;

        for(const auto& fmt : formats)
        {
-            if(this->*(fmt.second))
+            if(fmt.enabled)
            {
-                if(!first && sizeof(formats) > 1) ss << ", ";
-                ss << fmt.first;
+                if(!first) ss << ", ";
+                ss << fmt.name;
                first = false;
            }
        }
+
        ROCPROFSYS_PRINT(
            "Generating [%s] format(s) with collected data from trace cache. This may "
            "take a while..\n",
            ss.str().c_str());
+
+        if(has_parallel_formats())
+        {
+            std::stringstream parallel_ss;
+            bool              first_parallel = true;
+            for(const auto& fmt : formats)
+            {
+                if(fmt.enabled && fmt.process_parallel)
+                {
+                    if(!first_parallel) parallel_ss << ", ";
+                    parallel_ss << fmt.name;
+                    first_parallel = false;
+                }
+            }
+            ROCPROFSYS_PRINT("  - Using parallel processing for: %s\n",
+                             parallel_ss.str().c_str());
+        }
+
+        if(has_sequential_formats())
+        {
+            std::stringstream sequential_ss;
+            bool              first_sequential = true;
+            for(const auto& fmt : formats)
+            {
+                if(fmt.enabled && !fmt.process_parallel)
+                {
+                    if(!first_sequential) sequential_ss << ", ";
+                    sequential_ss << fmt.name;
+                    first_sequential = false;
+                }
+            }
+            ROCPROFSYS_PRINT("  - Using sequential processing for: %s\n",
+                             sequential_ss.str().c_str());
+        }
+    }
+
+    bool has_parallel_formats() const
+    {
+        return std::any_of(formats.begin(), formats.end(),
+                           [](const auto& f) { return f.enabled && f.process_parallel; });
+    }
+
+    bool has_sequential_formats() const
+    {
+        return std::any_of(formats.begin(), formats.end(), [](const auto& f) {
+            return f.enabled && !f.process_parallel;
+        });
+    }
+
+    enabled_formats_t get_parallel_formats() const
+    {
+        enabled_formats_t parallel_formats;
+        parallel_formats.formats.clear();
+        for(const auto& fmt : formats)
+        {
+            parallel_formats.formats.push_back(
+                { true, fmt.enabled && fmt.process_parallel, fmt.name });
+        }
+        return parallel_formats;
+    }
+
+    enabled_formats_t get_sequential_formats() const
+    {
+        enabled_formats_t sequential_formats;
+        sequential_formats.formats.clear();
+        for(const auto& fmt : formats)
+        {
+            sequential_formats.formats.push_back(
+                { false, fmt.enabled && !fmt.process_parallel, fmt.name });
+        }
+        return sequential_formats;
+    }
+
+    bool is_rocpd_enabled() const
+    {
+        auto it = std::find_if(formats.begin(), formats.end(), [](const auto& f) {
+            return std::strcmp(f.name, "rocpd") == 0;
+        });
+        return it != formats.end() && it->enabled;
+    }
+
+    bool is_perfetto_enabled() const
+    {
+        auto it = std::find_if(formats.begin(), formats.end(), [](const auto& f) {
+            return std::strcmp(f.name, "perfetto") == 0;
+        });
+        return it != formats.end() && it->enabled;
    }
 };

@@ -105,7 +198,8 @@ struct processor_config_t

 struct processor_storage_t
 {
-    std::shared_ptr<rocpd_processor_t> rocpd_processor{ nullptr };
+    std::shared_ptr<rocpd_processor_t>    rocpd_processor{ nullptr };
+    std::shared_ptr<perfetto_processor_t> perfetto_processor{ nullptr };
 };

 using directory_files_t    = std::vector<std::string>;
@@ -234,6 +328,84 @@ clear_cache_files(const data::mapped_cache_files_t& _cache_files)
    }
 }

+void
+merge_perfetto_files(const std::vector<std::string>& perfetto_files,
+                     const std::string&              _filename)
+{
+    if(perfetto_files.empty())
+    {
+        ROCPROFSYS_VERBOSE(
+            0, "perfetto trace data is empty. File '%s' will not be written...\n",
+            _filename.c_str());
+        return;
+    }
+
+    std::vector<char> trace_data;
+    size_t            total_size = 0;
+
+    // Calculate total size for reservation
+    for(const auto& file : perfetto_files)
+    {
+        std::ifstream ifs(file, std::ios::binary | std::ios::ate);
+        if(ifs)
+        {
+            total_size += ifs.tellg();
+        }
+    }
+
+    trace_data.reserve(total_size);
+
+    // Read and concatenate all files
+    for(const auto& file : perfetto_files)
+    {
+        std::ifstream ifs(file, std::ios::binary);
+        if(!ifs)
+        {
+            ROCPROFSYS_VERBOSE(-1, "Error opening '%s'...\n", file.c_str());
+            continue;
+        }
+
+        ifs.seekg(0, std::ios::end);
+        size_t file_size = ifs.tellg();
+        ifs.seekg(0, std::ios::beg);
+
+        size_t current_size = trace_data.size();
+        trace_data.resize(current_size + file_size);
+
+        ifs.read(trace_data.data() + current_size, file_size);
+    }
+
+    if(!trace_data.empty())
+    {
+        operation::file_output_message<tim::project::rocprofsys> _fom{};
+        // Write the trace into a file.
+        if(config::get_verbose() >= 0)
+            _fom(_filename, std::string{ "perfetto" },
+                 " (%.2f KB / %.2f MB / %.2f GB)... ",
+                 static_cast<double>(trace_data.size()) / units::KB,
+                 static_cast<double>(trace_data.size()) / units::MB,
+                 static_cast<double>(trace_data.size()) / units::GB);
+        std::ofstream ofs{};
+        if(!filepath::open(ofs, _filename, std::ios::out | std::ios::binary))
+        {
+            _fom.append("Error opening '%s'...", _filename.c_str());
+        }
+        else
+        {
+            // Write the trace into a file.
+            ofs.write(trace_data.data(), trace_data.size());
+            if(config::get_verbose() >= 0) _fom.append("%s", "Done");  // NOLINT
+        }
+        ofs.close();
+    }
+    else
+    {
+        ROCPROFSYS_VERBOSE(
+            0, "perfetto trace data is empty. File '%s' will not be written...\n",
+            _filename.c_str());
+    }
+}
+
 }  // namespace filesystem_utils

 namespace processing_utils
@@ -244,13 +416,20 @@ configure_processors(const std::shared_ptr<sample_processor_t>&       _type_proc
                     const data::enabled_formats_t&                   _enabled_formats)
 {
    data::processor_storage_t processor_storage;
-    if(_enabled_formats.rocpd)
+    if(_enabled_formats.is_rocpd_enabled())
    {
        processor_storage.rocpd_processor = std::make_shared<rocpd_processor_t>(
            _processor_config->_metadata_registry, _processor_config->_agent_manager,
            _processor_config->_pid, _processor_config->_ppid);
        _type_processing->add_handler(*processor_storage.rocpd_processor);
    }
+    if(_enabled_formats.is_perfetto_enabled())
+    {
+        processor_storage.perfetto_processor = std::make_shared<perfetto_processor_t>(
+            _processor_config->_metadata_registry, _processor_config->_agent_manager,
+            _processor_config->_pid, _processor_config->_ppid);
+        _type_processing->add_handler(*processor_storage.perfetto_processor);
+    }
    return processor_storage;
 }

@@ -321,6 +500,37 @@ multithreaded_processing(
    }
 }

+void
+sequential_processing(
+    const std::vector<std::shared_ptr<data::processor_config_t>>& _processor_configs,
+    const data::enabled_formats_t&                                _enabled_formats)
+{
+    for(const auto& processor_config : _processor_configs)
+    {
+        process_buffered_storage(processor_config,
+                                 utility::get_buffered_storage_filename(
+                                     processor_config->_ppid, processor_config->_pid),
+                                 _enabled_formats);
+    }
+}
+
+void
+dispatch_processing(
+    const std::vector<std::shared_ptr<data::processor_config_t>>& _processor_configs,
+    const data::enabled_formats_t&                                _enabled_formats)
+{
+    if(_enabled_formats.has_sequential_formats())
+    {
+        auto sequential_formats = _enabled_formats.get_sequential_formats();
+        sequential_processing(_processor_configs, sequential_formats);
+    }
+    if(_enabled_formats.has_parallel_formats())
+    {
+        auto parallel_formats = _enabled_formats.get_parallel_formats();
+        multithreaded_processing(_processor_configs, parallel_formats);
+    }
+}
+
 }  // namespace processing_utils

 cache_manager&
@@ -361,7 +571,50 @@ cache_manager::post_process_bulk()
        getpid(), root_pid, m_metadata,
        std::make_shared<agent_manager>(get_agent_manager_instance().get_agents())));

-    processing_utils::multithreaded_processing(processor_configs, enabled_formats);
+    processing_utils::dispatch_processing(processor_configs, enabled_formats);
+
+    if(enabled_formats.is_perfetto_enabled())
+    {
+        std::vector<std::string> perfetto_files;
+
+        for(const auto& config : processor_configs)
+        {
+            // Check for both naming styles: default (current process) and PID-suffixed
+            auto filename_default = config::get_perfetto_output_filename();
+            auto filename_suffix  = config::get_perfetto_output_filename_with_suffix(
+                std::to_string(config->_pid));
+
+            if(static_cast<pid_t>(config->_pid) == getpid() &&
+               tim::filepath::exists(filename_default))
+            {
+                perfetto_files.push_back(filename_default);
+            }
+            else if(tim::filepath::exists(filename_suffix))
+            {
+                perfetto_files.push_back(filename_suffix);
+            }
+        }
+
+        if(config::get_perfetto_combined_traces() && perfetto_files.size() > 1)
+        {
+            // Use base filename without suffix for merged output
+            auto _filename = config::get_perfetto_output_filename();
+            filesystem_utils::merge_perfetto_files(perfetto_files, _filename);
+        }
+        else if(perfetto_files.size() > 1)
+        {
+            ROCPROFSYS_VERBOSE(
+                0,
+                "Generated %zu separate perfetto trace files. "
+                "Set ROCPROFSYS_PERFETTO_COMBINE_TRACES=ON to merge them.\n",
+                perfetto_files.size());
+
+            for(const auto& file : perfetto_files)
+            {
+                ROCPROFSYS_VERBOSE(1, "  - %s\n", file.c_str());
+            }
+        }
+    }

    filesystem_utils::clear_cache_files(cache_files);
 }
@@ -0,0 +1,88 @@
+// MIT License
+//
+// Copyright (c) 2025 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 "agent_manager.hpp"
+#include "config.hpp"
+#include "core/perfetto_fwd.hpp"
+#include "core/trace_cache/metadata_registry.hpp"
+#include "core/trace_cache/sample_processor.hpp"
+
+#include <functional>
+#include <memory>
+#include <perfetto.h>
+
+namespace rocprofsys
+{
+namespace trace_cache
+{
+using char_vec_t = std::vector<char>;
+
+struct pmc_track_info
+{
+    const char*                                                           default_units;
+    std::function<bool(uint64_t)>                                         exists_fn;
+    std::function<void(uint64_t, const std::string&, const std::string&)> emplace_fn;
+    std::function<void(uint64_t, uint64_t, uint64_t, double)>             trace_fn;
+};
+
+class perfetto_processor_t : public processor_t<perfetto_processor_t>
+{
+public:
+    perfetto_processor_t(const std::shared_ptr<metadata_registry>& metadata,
+                         const std::shared_ptr<agent_manager>& agent_mngr, int pid,
+                         int ppid);
+
+    void prepare_for_processing();
+    void finalize_processing();
+
+    void handle(const kernel_dispatch_sample& sample);
+    void handle(const memory_copy_sample& sample);
+    void handle(const memory_allocate_sample& sample);
+    void handle(const region_sample& sample);
+    void handle(const in_time_sample& sample);
+    void handle(const pmc_event_with_sample& sample);
+    void handle(const amd_smi_sample& sample);
+    void handle(const cpu_freq_sample& sample);
+    void handle(const backtrace_region_sample& sample);
+
+private:
+    void       initialize_perfetto();
+    void       setup_perfetto();
+    void       start_session();
+    void       stop_session();
+    void       flush(bool& perfetto_output_error);
+    char_vec_t get_session_data();
+
+    metadata_registry&                          m_metadata;
+    uint64_t                                    m_process_id;
+    uint64_t                                    m_parrent_pid;
+    agent_manager&                              m_agent_manager;
+    ::perfetto::TraceConfig                     m_session_config;
+    std::shared_ptr<tmp_file>                   m_tmp_file{ nullptr };
+    std::unique_ptr<::perfetto::TracingSession> m_tracing_session{ nullptr };
+    bool                                        m_use_annotations{ false };
+
+    std::unordered_map<size_t, pmc_track_info> m_pmc_track_map;
+};
+}  // namespace trace_cache
+}  // namespace rocprofsys
@@ -22,6 +22,7 @@

 #include "core/trace_cache/rocpd_processor.hpp"
 #include "core/agent_manager.hpp"
+#include "core/common_types.hpp"
 #include "core/config.hpp"
 #include "core/debug.hpp"
 #include "core/gpu_metrics.hpp"
@@ -234,48 +235,6 @@ void
 rocpd_processor_t::handle([[maybe_unused]] const region_sample& _rs)
 {
 #if ROCPROFSYS_USE_ROCM > 0
-    static auto parse_args = []([[maybe_unused]] const std::string& arg_str) {
-        rocprofiler_sdk::function_args_t args;
-        const std::string                delimiter = ";;";
-
-        auto split = [](const std::string& str, const std::string& _delimiter) {
-            std::vector<std::string> tokens;
-            size_t                   start = 0;
-            size_t                   end   = str.find(_delimiter);
-
-            while(end != std::string::npos)
-            {
-                tokens.push_back(str.substr(start, end - start));
-                start = end + _delimiter.length();
-                end   = str.find(_delimiter, start);
-            }
-
-            return tokens;
-        };
-
-        if(arg_str.empty())
-        {
-            return args;
-        }
-
-        auto tokens = split(arg_str, delimiter);
-
-        // Ensure the number of tokens is a multiple of 4
-        if(tokens.size() % 4 != 0)
-        {
-            throw std::invalid_argument("Malformed argument string.");
-        }
-
-        for(auto it = tokens.begin(); it != tokens.end(); it += 4)
-        {
-            rocprofiler_sdk::argument_info arg = { static_cast<uint32_t>(std::stoi(*it)),
-                                                   *(it + 1), *(it + 2), *(it + 3) };
-            args.push_back(arg);
-        }
-
-        return args;
-    };
-
    auto& n_info  = node_info::get_instance();
    auto  process = m_metadata->get_process_info();
    auto  thread_primary_key =
@@ -292,7 +251,7 @@ rocpd_processor_t::handle([[maybe_unused]] const region_sample& _rs)
        m_data_processor->insert_event(category_primary_key, stack_id, parent_stack_id,
                                       correlation_id, _rs.call_stack.c_str());

-    auto args = parse_args(_rs.args_str);
+    auto args = process_arguments_string(_rs.args_str);
    for(const auto& arg : args)
    {
        m_data_processor->insert_args(event_primary_key, arg.arg_number,
@@ -114,7 +114,7 @@ serialize(uint8_t* buffer, const kernel_dispatch_sample& item)
        item.correlation_id_internal, item.correlation_id_ancestor,
        item.private_segment_size, item.group_segment_size, item.workgroup_size_x,
        item.workgroup_size_y, item.workgroup_size_z, item.grid_size_x, item.grid_size_y,
-        item.grid_size_z, (uint64_t) item.stream_handle);
+        item.grid_size_z, static_cast<uint64_t>(item.stream_handle));
 }

 template <>
@@ -144,7 +144,7 @@ get_size(const kernel_dispatch_sample& item)
        item.correlation_id_internal, item.correlation_id_ancestor,
        item.private_segment_size, item.group_segment_size, item.workgroup_size_x,
        item.workgroup_size_y, item.workgroup_size_z, item.grid_size_x, item.grid_size_y,
-        item.grid_size_z, (uint64_t) item.stream_handle);
+        item.grid_size_z, static_cast<uint64_t>(item.stream_handle));
 }

 struct memory_copy_sample : cacheable_t
@@ -196,7 +196,8 @@ serialize(uint8_t* buffer, const memory_copy_sample& item)
                         item.dst_agent_id_handle, item.src_agent_id_handle, item.kind,
                         item.operation, item.bytes, item.correlation_id_internal,
                         item.correlation_id_ancestor, item.dst_address_value,
-                         item.src_address_value, (uint64_t) item.stream_handle);
+                         item.src_address_value,
+                         static_cast<uint64_t>(item.stream_handle));
 }

 template <>
@@ -218,11 +219,12 @@ template <>
 inline size_t
 get_size(const memory_copy_sample& item)
 {
-    return utility::get_size(
-        item.start_timestamp, item.end_timestamp, item.thread_id,
-        item.dst_agent_id_handle, item.src_agent_id_handle, item.kind, item.operation,
-        item.bytes, item.correlation_id_internal, item.correlation_id_ancestor,
-        item.dst_address_value, item.src_address_value, (uint64_t) item.stream_handle);
+    return utility::get_size(item.start_timestamp, item.end_timestamp, item.thread_id,
+                             item.dst_agent_id_handle, item.src_agent_id_handle,
+                             item.kind, item.operation, item.bytes,
+                             item.correlation_id_internal, item.correlation_id_ancestor,
+                             item.dst_address_value, item.src_address_value,
+                             static_cast<uint64_t>(item.stream_handle));
 }

 struct memory_allocate_sample : cacheable_t
@@ -270,7 +272,7 @@ serialize(uint8_t* buffer, const memory_allocate_sample& item)
                         item.agent_id_handle, item.kind, item.operation,
                         item.allocation_size, item.correlation_id_internal,
                         item.correlation_id_ancestor, item.address_value,
-                         (uint64_t) item.stream_handle);
+                         static_cast<uint64_t>(item.stream_handle));
 }

 template <>
@@ -291,10 +293,11 @@ template <>
 inline size_t
 get_size(const memory_allocate_sample& item)
 {
-    return utility::get_size(
-        item.start_timestamp, item.end_timestamp, item.thread_id, item.agent_id_handle,
-        item.kind, item.operation, item.allocation_size, item.correlation_id_internal,
-        item.correlation_id_ancestor, item.address_value, (uint64_t) item.stream_handle);
+    return utility::get_size(item.start_timestamp, item.end_timestamp, item.thread_id,
+                             item.agent_id_handle, item.kind, item.operation,
+                             item.allocation_size, item.correlation_id_internal,
+                             item.correlation_id_ancestor, item.address_value,
+                             static_cast<uint64_t>(item.stream_handle));
 }

 struct region_sample : cacheable_t
@@ -373,11 +376,12 @@ struct in_time_sample : cacheable_t
        type_identifier_t::in_time_sample;

    in_time_sample() = default;
-    in_time_sample(std::string _track_name, size_t _timestamp_ns,
-                   std::string _event_metadata, size_t _stack_id, size_t _parent_stack_id,
-                   size_t _correlation_id, std::string _call_stack,
-                   std::string _line_info)
-    : track_name(std::move(_track_name))
+    in_time_sample(size_t _category_enum_id, std::string _track_name,
+                   size_t _timestamp_ns, std::string _event_metadata, size_t _stack_id,
+                   size_t _parent_stack_id, size_t _correlation_id,
+                   std::string _call_stack, std::string _line_info)
+    : category_enum_id(_category_enum_id)
+    , track_name(std::move(_track_name))
    , timestamp_ns(_timestamp_ns)
    , event_metadata(std::move(_event_metadata))
    , stack_id(_stack_id)
@@ -387,6 +391,7 @@ struct in_time_sample : cacheable_t
    , line_info(std::move(_line_info))
    {}

+    size_t      category_enum_id;
    std::string track_name;
    size_t      timestamp_ns;
    std::string event_metadata;
@@ -402,10 +407,11 @@ inline void
 serialize(uint8_t* buffer, const in_time_sample& item)
 {
    utility::store_value(
-        buffer, std::string_view(item.track_name), (uint64_t) item.timestamp_ns,
-        std::string_view(item.event_metadata), (uint64_t) item.stack_id,
-        (uint64_t) item.parent_stack_id, (uint64_t) item.correlation_id,
-        std::string_view(item.call_stack), std::string_view(item.line_info));
+        buffer, std::string_view(item.track_name),
+        static_cast<uint64_t>(item.timestamp_ns), std::string_view(item.event_metadata),
+        static_cast<uint64_t>(item.stack_id), static_cast<uint64_t>(item.parent_stack_id),
+        static_cast<uint64_t>(item.correlation_id), std::string_view(item.call_stack),
+        std::string_view(item.line_info));
 }

 template <>
@@ -413,20 +419,22 @@ inline in_time_sample
 deserialize(uint8_t*& buffer)
 {
    in_time_sample   item;
+    size_t           category_enum_id;
    std::string_view track_name_view, event_metadata_view, call_stack_view,
        line_info_view;
    uint64_t timestamp_ns, stack_id, parent_stack_id, correlation_id;
-    utility::parse_value(buffer, track_name_view, timestamp_ns, event_metadata_view,
-                         stack_id, parent_stack_id, correlation_id, call_stack_view,
-                         line_info_view);
-    item.track_name      = std::string(track_name_view);
-    item.timestamp_ns    = timestamp_ns;
-    item.event_metadata  = std::string(event_metadata_view);
-    item.stack_id        = stack_id;
-    item.parent_stack_id = parent_stack_id;
-    item.correlation_id  = correlation_id;
-    item.call_stack      = std::string(call_stack_view);
-    item.line_info       = std::string(line_info_view);
+    utility::parse_value(buffer, category_enum_id, track_name_view, timestamp_ns,
+                         event_metadata_view, stack_id, parent_stack_id, correlation_id,
+                         call_stack_view, line_info_view);
+    item.category_enum_id = category_enum_id;
+    item.track_name       = std::string(track_name_view);
+    item.timestamp_ns     = timestamp_ns;
+    item.event_metadata   = std::string(event_metadata_view);
+    item.stack_id         = stack_id;
+    item.parent_stack_id  = parent_stack_id;
+    item.correlation_id   = correlation_id;
+    item.call_stack       = std::string(call_stack_view);
+    item.line_info        = std::string(line_info_view);
    return item;
 }

@@ -435,10 +443,11 @@ inline size_t
 get_size(const in_time_sample& item)
 {
    return utility::get_size(
-        std::string_view(item.track_name), (uint64_t) item.timestamp_ns,
-        std::string_view(item.event_metadata), (uint64_t) item.stack_id,
-        (uint64_t) item.parent_stack_id, (uint64_t) item.correlation_id,
-        std::string_view(item.call_stack), std::string_view(item.line_info));
+        item.category_enum_id, std::string_view(item.track_name),
+        static_cast<uint64_t>(item.timestamp_ns), std::string_view(item.event_metadata),
+        static_cast<uint64_t>(item.stack_id), static_cast<uint64_t>(item.parent_stack_id),
+        static_cast<uint64_t>(item.correlation_id), std::string_view(item.call_stack),
+        std::string_view(item.line_info));
 }

 struct pmc_event_with_sample : in_time_sample
@@ -447,15 +456,15 @@ struct pmc_event_with_sample : in_time_sample
        type_identifier_t::pmc_event_with_sample;

    pmc_event_with_sample() = default;
-    pmc_event_with_sample(std::string _track_name, size_t _timestamp_ns,
-                          std::string _event_metadata, size_t _stack_id,
-                          size_t _parent_stack_id, size_t _correlation_id,
-                          std::string _call_stack, std::string _line_info,
-                          uint32_t _device_id, uint8_t _device_type,
-                          std::string _pmc_info_name, double _value)
-    : in_time_sample(std::move(_track_name), _timestamp_ns, std::move(_event_metadata),
-                     _stack_id, _parent_stack_id, _correlation_id, std::move(_call_stack),
-                     std::move(_line_info))
+    pmc_event_with_sample(size_t _category_enum_id, std::string _track_name,
+                          size_t _timestamp_ns, std::string _event_metadata,
+                          size_t _stack_id, size_t _parent_stack_id,
+                          size_t _correlation_id, std::string _call_stack,
+                          std::string _line_info, uint32_t _device_id,
+                          uint8_t _device_type, std::string _pmc_info_name, double _value)
+    : in_time_sample(_category_enum_id, std::move(_track_name), _timestamp_ns,
+                     std::move(_event_metadata), _stack_id, _parent_stack_id,
+                     _correlation_id, std::move(_call_stack), std::move(_line_info))
    , device_id(_device_id)
    , device_type(_device_type)
    , pmc_info_name(std::move(_pmc_info_name))
@@ -473,12 +482,12 @@ inline void
 serialize(uint8_t* buffer, const pmc_event_with_sample& item)
 {
    utility::store_value(
-        buffer, std::string_view(item.track_name), (uint64_t) item.timestamp_ns,
-        std::string_view(item.event_metadata), (uint64_t) item.stack_id,
-        (uint64_t) item.parent_stack_id, (uint64_t) item.correlation_id,
-        std::string_view(item.call_stack), std::string_view(item.line_info),
-        item.device_id, item.device_type, std::string_view(item.pmc_info_name),
-        item.value);
+        buffer, item.category_enum_id, std::string_view(item.track_name),
+        static_cast<uint64_t>(item.timestamp_ns), std::string_view(item.event_metadata),
+        static_cast<uint64_t>(item.stack_id), static_cast<uint64_t>(item.parent_stack_id),
+        static_cast<uint64_t>(item.correlation_id), std::string_view(item.call_stack),
+        std::string_view(item.line_info), item.device_id, item.device_type,
+        std::string_view(item.pmc_info_name), item.value);
 }

 template <>
@@ -486,22 +495,24 @@ inline pmc_event_with_sample
 deserialize(uint8_t*& buffer)
 {
    pmc_event_with_sample item;
+    size_t                category_enum_id;
    std::string_view      track_name_view, event_metadata_view, call_stack_view,
        line_info_view, pmc_info_name_view;
    uint64_t timestamp_ns, stack_id, parent_stack_id, correlation_id;
-    utility::parse_value(buffer, track_name_view, timestamp_ns, event_metadata_view,
-                         stack_id, parent_stack_id, correlation_id, call_stack_view,
-                         line_info_view, item.device_id, item.device_type,
-                         pmc_info_name_view, item.value);
-    item.track_name      = std::string(track_name_view);
-    item.timestamp_ns    = timestamp_ns;
-    item.event_metadata  = std::string(event_metadata_view);
-    item.stack_id        = stack_id;
-    item.parent_stack_id = parent_stack_id;
-    item.correlation_id  = correlation_id;
-    item.call_stack      = std::string(call_stack_view);
-    item.line_info       = std::string(line_info_view);
-    item.pmc_info_name   = std::string(pmc_info_name_view);
+    utility::parse_value(buffer, category_enum_id, track_name_view, timestamp_ns,
+                         event_metadata_view, stack_id, parent_stack_id, correlation_id,
+                         call_stack_view, line_info_view, item.device_id,
+                         item.device_type, pmc_info_name_view, item.value);
+    item.category_enum_id = category_enum_id;
+    item.track_name       = std::string(track_name_view);
+    item.timestamp_ns     = timestamp_ns;
+    item.event_metadata   = std::string(event_metadata_view);
+    item.stack_id         = stack_id;
+    item.parent_stack_id  = parent_stack_id;
+    item.correlation_id   = correlation_id;
+    item.call_stack       = std::string(call_stack_view);
+    item.line_info        = std::string(line_info_view);
+    item.pmc_info_name    = std::string(pmc_info_name_view);
    return item;
 }

@@ -510,12 +521,12 @@ inline size_t
 get_size(const pmc_event_with_sample& item)
 {
    return utility::get_size(
-        std::string_view(item.track_name), (uint64_t) item.timestamp_ns,
-        std::string_view(item.event_metadata), (uint64_t) item.stack_id,
-        (uint64_t) item.parent_stack_id, (uint64_t) item.correlation_id,
-        std::string_view(item.call_stack), std::string_view(item.line_info),
-        item.device_id, item.device_type, std::string_view(item.pmc_info_name),
-        item.value);
+        item.category_enum_id, std::string_view(item.track_name),
+        static_cast<uint64_t>(item.timestamp_ns), std::string_view(item.event_metadata),
+        static_cast<uint64_t>(item.stack_id), static_cast<uint64_t>(item.parent_stack_id),
+        static_cast<uint64_t>(item.correlation_id), std::string_view(item.call_stack),
+        std::string_view(item.line_info), item.device_id, item.device_type,
+        std::string_view(item.pmc_info_name), item.value);
 }

 struct amd_smi_sample : cacheable_t
@@ -568,10 +579,10 @@ template <>
 inline void
 serialize(uint8_t* buffer, const amd_smi_sample& item)
 {
-    utility::store_value(buffer, item.settings, item.device_id, (uint64_t) item.timestamp,
-                         item.gfx_activity, item.umc_activity, item.mm_activity,
-                         item.power, item.temperature, (uint64_t) item.mem_usage,
-                         item.gpu_activity);
+    utility::store_value(
+        buffer, item.settings, item.device_id, static_cast<uint64_t>(item.timestamp),
+        item.gfx_activity, item.umc_activity, item.mm_activity, item.power,
+        item.temperature, static_cast<uint64_t>(item.mem_usage), item.gpu_activity);
 }

 template <>
@@ -592,10 +603,10 @@ template <>
 inline size_t
 get_size(const amd_smi_sample& item)
 {
-    return utility::get_size(item.settings, item.device_id, (uint64_t) item.timestamp,
-                             item.gfx_activity, item.umc_activity, item.mm_activity,
-                             item.power, item.temperature, (uint64_t) item.mem_usage,
-                             item.gpu_activity);
+    return utility::get_size(
+        item.settings, item.device_id, static_cast<uint64_t>(item.timestamp),
+        item.gfx_activity, item.umc_activity, item.mm_activity, item.power,
+        item.temperature, static_cast<uint64_t>(item.mem_usage), item.gpu_activity);
 }

 struct cpu_freq_sample : cacheable_t
@@ -634,7 +645,7 @@ template <>
 inline void
 serialize(uint8_t* buffer, const cpu_freq_sample& item)
 {
-    utility::store_value(buffer, (uint64_t) item.timestamp, item.page_rss,
+    utility::store_value(buffer, static_cast<uint64_t>(item.timestamp), item.page_rss,
                         item.virt_mem_usage, item.peak_rss, item.context_switch_count,
                         item.page_faults, item.user_mode_time, item.kernel_mode_time,
                         item.freqs);
@@ -657,7 +668,7 @@ template <>
 inline size_t
 get_size(const cpu_freq_sample& item)
 {
-    return utility::get_size((uint64_t) item.timestamp, item.page_rss,
+    return utility::get_size(static_cast<uint64_t>(item.timestamp), item.page_rss,
                             item.virt_mem_usage, item.peak_rss,
                             item.context_switch_count, item.page_faults,
                             item.user_mode_time, item.kernel_mode_time, item.freqs);
@@ -150,9 +150,6 @@ public:
        }

        ifs.close();
-        ROCPROFSYS_DEBUG("File parsing finished. Removing %s from file system.\n",
-                         m_filename.c_str());
-        std::remove(m_filename.c_str());
    }

 private:
@@ -262,9 +262,10 @@ cache_backtrace_metrics_events(const uint32_t device_id, uint64_t timestamp_ns,

    auto insert_event_and_sample = [&](const char* _track_name, double _value) {
        trace_cache::get_buffer_storage().store(trace_cache::pmc_event_with_sample{
-            _track_name, timestamp_ns, event_metadata, stack_id, parent_stack_id,
-            correlation_id, call_stack, line_info, device_id,
-            static_cast<uint8_t>(agent_type::CPU), _track_name, _value });
+            static_cast<size_t>(category_enum_id<Category>::value), _track_name,
+            timestamp_ns, event_metadata, stack_id, parent_stack_id, correlation_id,
+            call_stack, line_info, device_id, static_cast<uint8_t>(agent_type::CPU),
+            _track_name, _value });
    };

    if constexpr(std::is_same_v<Category, category::thread_hardware_counter>)
@@ -152,6 +152,7 @@ cache_comm_data_events(const uint32_t device_id, int bytes)
    const std::string line_info       = "{}";

    trace_cache::get_buffer_storage().store(trace_cache::pmc_event_with_sample{
+        static_cast<size_t>(category_enum_id<category::comm_data>::value),
        track_name.c_str(), timestamp_ns, event_metadata.c_str(), stack_id,
        parent_stack_id, correlation_id, call_stack.c_str(), line_info.c_str(), device_id,
        static_cast<uint8_t>(agent_type::CPU), track_name.c_str(),
@@ -192,6 +192,7 @@ cache_kokkos_event(const char* name, const char* event_type, const char* target,

    rocprofsys::trace_cache::get_buffer_storage().store(
        rocprofsys::trace_cache::in_time_sample{
+            static_cast<size_t>(rocprofsys::category_enum_id<category::kokkos>::value),
            rocprofsys::trait::name<category::kokkos>::value, timestamp_ns,
            event_metadata.dump().c_str(), stack_id, parent_stack_id, correlation_id,
            call_stack, line_info });
@@ -24,6 +24,7 @@
 #include "api.hpp"
 #include "common/synchronized.hpp"
 #include "core/common.hpp"
+#include "core/common_types.hpp"
 #include "core/config.hpp"
 #include "core/containers/stable_vector.hpp"
 #include "core/debug.hpp"
@@ -594,20 +595,6 @@ cache_memory_allocation(rocprofiler_buffer_tracing_memory_allocation_record_t* r
 }
 #endif

-std::string
-get_args_string(const function_args_t& args)
-{
-    std::string args_str;
-    std::for_each(args.begin(), args.end(), [&args_str](const argument_info& arg) {
-        const auto*       delimiter = ";;";
-        std::stringstream ss;
-        ss << arg.arg_number << delimiter << arg.arg_type << delimiter << arg.arg_name
-           << delimiter << arg.arg_value << delimiter;
-        args_str.append(ss.str());
-    });
-    return args_str;
-}
-
 template <typename CategoryT>
 void
 tool_tracing_callback_start(CategoryT, rocprofiler_callback_tracing_record_t record,
@@ -121,6 +121,8 @@ counter_event::operator()(const client_data* tool_data, ::perfetto::CounterTrack
        auto agent = get_agent_manager_instance().get_agent_by_handle(agent_handle);

        trace_cache::get_buffer_storage().store(trace_cache::pmc_event_with_sample{
+            static_cast<size_t>(
+                category_enum_id<category::rocm_counter_collection>::value),
            track_name.c_str(), _timing.start, event_metadata.c_str(), stack_id,
            parent_stack_id, correlation_id, call_stack.c_str(), line_info.c_str(),
            static_cast<uint32_t>(agent.device_id), static_cast<uint8_t>(agent.type),
@@ -174,8 +176,8 @@ counter_storage::operator()(const counter_event& _event, timing_interval _timing
 }

 void
-counter_storage::write(counter_storage_type* storage, std::string metric_name,
-                       std::string metric_description)
+counter_storage::write(counter_storage_type* storage, const std::string& metric_name,
+                       const std::string& metric_description)
 {
    if(!trait::runtime_enabled<counter_data_tracker>::get())
    {
@@ -109,8 +109,8 @@ struct counter_storage
    void operator()(const counter_event& _event, timing_interval _timing,
                    scope::config _scope = scope::get_default()) const;

-    static void write(counter_storage_type* storage, std::string metric_name,
-                      std::string metric_description);
+    static void write(counter_storage_type* storage, const std::string& metric_name,
+                      const std::string& metric_description);
 };
 }  // namespace rocprofiler_sdk
 }  // namespace rocprofsys
@@ -101,16 +101,6 @@ struct timing_interval
    rocprofiler_timestamp_t end   = 0;
 };

-struct argument_info
-{
-    uint32_t    arg_number = 0;
-    std::string arg_type   = {};
-    std::string arg_name   = {};
-    std::string arg_value  = {};
-};
-
-using function_args_t = std::vector<argument_info>;
-
 using agent_counter_info_map_t =
    std::unordered_map<rocprofiler_agent_id_t,
                       std::vector<rocprofiler_tool_counter_info_t>>;