Kernel Serialization Support (#379)

* Serialization-rebased with main branch * Removing client_id from queue completion callbacks * removing debugging code * source formatting (clang-format v11) (#449) Co-authored-by: SrirakshaNag <SrirakshaNag@users.noreply.github.com> * moving ready signal handler to anonymous namespace * source formatting (clang-format v11) (#450) Co-authored-by: SrirakshaNag <SrirakshaNag@users.noreply.github.com> * Handling deque search better in queue destructor * source formatting (clang-format v11) (#451) Co-authored-by: SrirakshaNag <SrirakshaNag@users.noreply.github.com> * disabling test_total_runtime test in code coverage --------- Co-authored-by: Benjamin Welton <bewelton@amd.com> Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Co-authored-by: SrirakshaNag <SrirakshaNag@users.noreply.github.com> [ROCm/rocprofiler-sdk commit: f6198f226a]
2024-02-05 12:42:59 -06:00
@@ -160,6 +160,10 @@ AQLPacketConstruct::construct_packet(const AmdExtTable& ext) const
                    "could not generate stop packet");
    throw_if_failed(hsa_ven_amd_aqlprofile_read(&profile, &pkt.read),
                    "could not generate read packet");
+    pkt.start.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
+    pkt.stop.header  = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
+    pkt.read.header  = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
+
    return pkt_ptr;
 }

@@ -37,6 +37,9 @@ namespace rocprofiler
 {
 namespace counters
 {
+using ClientID   = int64_t;
+using inst_pkt_t = common::container::
+    small_vector<std::pair<std::unique_ptr<rocprofiler::hsa::AQLPacket>, ClientID>, 4>;
 class CounterController
 {
 public:
@@ -136,9 +139,8 @@ destroy_counter_profile(uint64_t id)
 * We return an AQLPacket containing the start/stop/read packets for injection.
 */
 std::unique_ptr<rocprofiler::hsa::AQLPacket>
-queue_cb(const std::shared_ptr<counter_callback_info>& info,
-         const hsa::Queue&                             queue,
-         hsa::ClientID,
+queue_cb(const std::shared_ptr<counter_callback_info>&                   info,
+         const hsa::Queue&                                               queue,
         const hsa::rocprofiler_packet&                                  pkt,
         uint64_t                                                        kernel_id,
         const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids,
@@ -245,9 +247,38 @@ queue_cb(const std::shared_ptr<counter_callback_info>& info,
        ret_pkt = prof_config->pkt_generator->construct_packet(
            hsa::get_queue_controller().get_ext_table());
    }
+    ret_pkt->before_krn_pkt.clear();
+    ret_pkt->after_krn_pkt.clear();
+    if(ret_pkt->empty) return ret_pkt;

    info->packet_return_map.wlock([&](auto& data) { data.emplace(ret_pkt.get(), prof_config); });

+    auto&& CreateBarrierPacket =
+        [](hsa_signal_t*                                                     dependency_signal,
+           hsa_signal_t*                                                     completion_signal,
+           common::container::small_vector<hsa_ext_amd_aql_pm4_packet_t, 3>& _packets) {
+            hsa::rocprofiler_packet barrier{};
+            barrier.barrier_and.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
+            if(dependency_signal != nullptr) barrier.barrier_and.dep_signal[0] = *dependency_signal;
+            if(completion_signal != nullptr)
+                barrier.barrier_and.completion_signal = *completion_signal;
+            _packets.emplace_back(barrier.ext_amd_aql_pm4);
+        };
+
+    hsa_signal_t ready_signal = queue.ready_signal;
+    hsa_signal_t block_signal = queue.block_signal;
+    CreateBarrierPacket(nullptr, &ready_signal, ret_pkt->before_krn_pkt);
+    CreateBarrierPacket(&block_signal, &block_signal, ret_pkt->before_krn_pkt);
+
+    ret_pkt->before_krn_pkt.push_back(ret_pkt->start);
+    ret_pkt->before_krn_pkt.end()->completion_signal.handle = 0;
+    ret_pkt->after_krn_pkt.push_back(ret_pkt->stop);
+    ret_pkt->after_krn_pkt.push_back(ret_pkt->read);
+    for(auto& aql_pkt : ret_pkt->after_krn_pkt)
+    {
+        aql_pkt.completion_signal.handle = 0;
+    }
+
    return ret_pkt;
 }

@@ -257,20 +288,32 @@ queue_cb(const std::shared_ptr<counter_callback_info>& info,
 void
 completed_cb(const std::shared_ptr<counter_callback_info>& info,
             const hsa::Queue&,
-             hsa::ClientID,
             hsa::rocprofiler_packet,
-             const hsa::Queue::queue_info_session_t&      session,
-             std::unique_ptr<rocprofiler::hsa::AQLPacket> pkt)
+             const hsa::Queue::queue_info_session_t& session,
+             inst_pkt_t&                             pkts)
 {
-    if(!info || !pkt) return;
+    if(!info || pkts.empty()) return;

    std::shared_ptr<profile_config> prof_config;
    // Get the Profile Config
+    std::unique_ptr<rocprofiler::hsa::AQLPacket> pkt = nullptr;
    info->packet_return_map.wlock([&](auto& data) {
-        prof_config = data.at(pkt.get());
-        data.erase(pkt.get());
+        for(auto& [aql_pkt, _] : pkts)
+        {
+            const auto* profile = rocprofiler::common::get_val(data, aql_pkt.get());
+            if(profile)
+            {
+                prof_config = *profile;
+                data.erase(aql_pkt.get());
+                pkt = std::move(aql_pkt);
+                return;
+            }
+        }
    });

+    if(!pkt) return;
+    hsa::profiler_serializer_kernel_completion_signal(session.queue.block_signal);
+
    auto decoded_pkt = EvaluateAST::read_pkt(prof_config->pkt_generator.get(), *pkt);
    EvaluateAST::read_special_counters(
        *prof_config->agent, prof_config->required_special_counters, decoded_pkt);
@@ -333,21 +376,17 @@ start_context(const context::context* ctx)
            cb->queue_id = controller.add_callback(
                std::nullopt,
                [=](const hsa::Queue&                                               q,
-                    hsa::ClientID                                                   c,
                    const hsa::rocprofiler_packet&                                  kern_pkt,
                    uint64_t                                                        kernel_id,
                    const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids,
                    const context::correlation_id* correlation_id) {
-                    return queue_cb(cb, q, c, kern_pkt, kernel_id, extern_corr_ids, correlation_id);
+                    return queue_cb(cb, q, kern_pkt, kernel_id, extern_corr_ids, correlation_id);
                },
                // Completion CB
                [=](const hsa::Queue&                       q,
-                    hsa::ClientID                           c,
                    hsa::rocprofiler_packet                 kern_pkt,
                    const hsa::Queue::queue_info_session_t& session,
-                    std::unique_ptr<hsa::AQLPacket>         aql) {
-                    completed_cb(cb, q, c, kern_pkt, session, std::move(aql));
-                });
+                    inst_pkt_t& aql) { completed_cb(cb, q, kern_pkt, session, aql); });
        }
        enabled = true;
    });
@@ -22,8 +22,11 @@

 #pragma once

+#include "lib/common/container/small_vector.hpp"
+
 #include <hsa/hsa_ext_amd.h>
 #include <hsa/hsa_ven_amd_aqlprofile.h>
+#include <deque>

 namespace rocprofiler
 {
@@ -53,14 +56,16 @@ struct AQLPacket
    AQLPacket(const AQLPacket&) = delete;
    AQLPacket& operator=(const AQLPacket&) = delete;

+    bool                             command_buf_mallocd    = false;
+    bool                             output_buffer_malloced = false;
+    bool                             empty                  = {true};
    hsa_ven_amd_aqlprofile_profile_t profile                = {};
    hsa_ext_amd_aql_pm4_packet_t     start                  = null_amd_aql_pm4_packet;
    hsa_ext_amd_aql_pm4_packet_t     stop                   = null_amd_aql_pm4_packet;
    hsa_ext_amd_aql_pm4_packet_t     read                   = null_amd_aql_pm4_packet;
-    bool                             command_buf_mallocd    = false;
-    bool                             output_buffer_malloced = false;
    memory_pool_free_func_t          free_func              = nullptr;
-    bool                             empty                  = {true};
+    common::container::small_vector<hsa_ext_amd_aql_pm4_packet_t, 3> before_krn_pkt = {};
+    common::container::small_vector<hsa_ext_amd_aql_pm4_packet_t, 2> after_krn_pkt  = {};
 };

 inline AQLPacket::AQLPacket(memory_pool_free_func_t func)
@@ -156,30 +156,15 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data)
            }
        }
    }
-
    // Calls our internal callbacks to callers who need to be notified post
    // kernel execution.
    queue_info_session.queue.signal_callback([&](const auto& map) {
        for(const auto& [client_id, cb_pair] : map)
        {
-            // If this is the client that gave us the AQLPacket,
-            // return it to that client otherwise notify.
-            if(queue_info_session.inst_pkt_id == client_id)
-            {
-                cb_pair.second(queue_info_session.queue,
-                               client_id,
-                               queue_info_session.kernel_pkt,
-                               queue_info_session,
-                               std::move(queue_info_session.inst_pkt));
-            }
-            else
-            {
-                cb_pair.second(queue_info_session.queue,
-                               client_id,
-                               queue_info_session.kernel_pkt,
-                               queue_info_session,
-                               nullptr);
-            }
+            cb_pair.second(queue_info_session.queue,
+                           queue_info_session.kernel_pkt,
+                           queue_info_session,
+                           queue_info_session.inst_pkt);
        }
    });

@@ -244,17 +229,13 @@ WriteInterceptor(const void* packets,
 {
    using context_array_t = Queue::context_array_t;

-    auto&& AddVendorSpecificPacket = [](hsa_ext_amd_aql_pm4_packet_t     _packet,
-                                        hsa_signal_t                     _signal,
-                                        std::vector<rocprofiler_packet>& _packets) {
-        _packets.emplace_back(_packet).ext_amd_aql_pm4.completion_signal = _signal;
-    };
-
-    auto&& CreateBarrierPacket = [](hsa_signal_t                     _signal,
+    auto&& CreateBarrierPacket = [](hsa_signal_t*                    dependency_signal,
+                                    hsa_signal_t*                    completion_signal,
                                    std::vector<rocprofiler_packet>& _packets) {
        hsa_barrier_and_packet_t barrier{};
-        barrier.header        = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
-        barrier.dep_signal[0] = _signal;
+        barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
+        if(dependency_signal != nullptr) barrier.dep_signal[0] = *dependency_signal;
+        if(completion_signal != nullptr) barrier.completion_signal = *completion_signal;
        _packets.emplace_back(barrier);
    };

@@ -319,32 +300,37 @@ WriteInterceptor(const void* packets,
        // Stores the instrumentation pkt (i.e. AQL packets for counter collection)
        // along with an ID of the client we got the packet from (this will be returned via
        // completed_cb_t)
-        ClientID                   inst_pkt_id = -1;
-        std::unique_ptr<AQLPacket> inst_pkt;
+
+        inst_pkt_t inst_pkt;

        // Signal callbacks that a kernel_pkt is being enqueued
        queue.signal_callback([&](const auto& map) {
            for(const auto& [client_id, cb_pair] : map)
            {
-                if(auto maybe_pkt = cb_pair.first(
-                       queue, client_id, kernel_pkt, kernel_id, extern_corr_ids, corr_id))
+                if(auto maybe_pkt =
+                       cb_pair.first(queue, kernel_pkt, kernel_id, extern_corr_ids, corr_id))
                {
-                    LOG_IF(FATAL, inst_pkt)
-                        << "We do not support two injections into the HSA queue";
-                    inst_pkt    = std::move(maybe_pkt);
-                    inst_pkt_id = client_id;
+                    inst_pkt.push_back(std::make_pair(std::move(maybe_pkt), client_id));
                }
            }
        });

-        constexpr auto dummy_signal = hsa_signal_t{.handle = 0};
-
-        // Write instrumentation start packet (if one exists)
-        if(inst_pkt && !inst_pkt->empty)
+        bool inserted_before = false;
+        for(const auto& pkt_injection : inst_pkt)
        {
-            inst_pkt->start.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
-            AddVendorSpecificPacket(inst_pkt->start, dummy_signal, transformed_packets);
-            CreateBarrierPacket(inst_pkt->start.completion_signal, transformed_packets);
+            for(const auto& pkt : pkt_injection.first->before_krn_pkt)
+            {
+                inserted_before = true;
+                transformed_packets.emplace_back(pkt);
+            }
+        }
+
+        // Barrier packet is last packet inserted into queue
+        if(inserted_before)
+        {
+            CreateBarrierPacket(&transformed_packets.back().ext_amd_aql_pm4.completion_signal,
+                                nullptr,
+                                transformed_packets);
        }

        transformed_packets.emplace_back(kernel_pkt);
@@ -363,15 +349,20 @@ WriteInterceptor(const void* packets,
        // Adding a barrier packet with the original packet's completion signal.
        queue.create_signal(0, &interrupt_signal);

-        if(inst_pkt && !inst_pkt->empty)
+        bool injected_end_pkt = false;
+        for(const auto& pkt_injection : inst_pkt)
        {
-            inst_pkt->stop.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
-            AddVendorSpecificPacket(inst_pkt->stop, dummy_signal, transformed_packets);
-            inst_pkt->read.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
-            AddVendorSpecificPacket(inst_pkt->read, interrupt_signal, transformed_packets);
+            for(const auto& pkt : pkt_injection.first->after_krn_pkt)
+            {
+                transformed_packets.emplace_back(pkt);
+                injected_end_pkt = true;
+            }
+        }

-            // Added Interrupt Signal with barrier and provided handler for it
-            CreateBarrierPacket(interrupt_signal, transformed_packets);
+        if(injected_end_pkt)
+        {
+            transformed_packets.back().ext_amd_aql_pm4.completion_signal = interrupt_signal;
+            CreateBarrierPacket(&interrupt_signal, nullptr, transformed_packets);
        }
        else
        {
@@ -391,7 +382,6 @@ WriteInterceptor(const void* packets,
            interrupt_signal,
            new Queue::queue_info_session_t{.queue            = queue,
                                            .inst_pkt         = std::move(inst_pkt),
-                                            .inst_pkt_id      = inst_pkt_id,
                                            .interrupt_signal = interrupt_signal,
                                            .tid              = thr_id,
                                            .kernel_id        = kernel_id,
@@ -462,6 +452,8 @@ Queue::Queue(const AgentCache&  agent,
           _ext_api.hsa_amd_queue_intercept_register_fn(_intercept_queue, WriteInterceptor, this))
        << "Could not register interceptor";

+    create_signal(0, &ready_signal);
+    create_signal(0, &block_signal);
    *queue = _intercept_queue;
 }

@@ -504,5 +496,17 @@ Queue::remove_callback(ClientID id)
        if(map.erase(id) == 1) _notifiers--;
    });
 }
+
+queue_state
+Queue::get_state() const
+{
+    return _state;
+}
+
+void
+Queue::set_state(queue_state state)
+{
+    _state = state;
+}
 }  // namespace hsa
 }  // namespace rocprofiler
@@ -57,6 +57,8 @@ struct correlation_id;
 namespace hsa
 {
 using ClientID = int64_t;
+using inst_pkt_t =
+    common::container::small_vector<std::pair<std::unique_ptr<AQLPacket>, ClientID>, 4>;

 union rocprofiler_packet
 {
@@ -92,6 +94,12 @@ union rocprofiler_packet
    rocprofiler_packet& operator=(const rocprofiler_packet&) = default;
    rocprofiler_packet& operator=(rocprofiler_packet&&) noexcept = default;
 };
+enum class queue_state
+{
+    normal       = 0,
+    to_destroy   = 1,
+    done_destroy = 2
+};

 // Interceptor for a single specific queue
 class Queue
@@ -109,8 +117,7 @@ public:
            std::unordered_map<const context_t*, rocprofiler_user_data_t>;

        Queue&                     queue;
-        std::unique_ptr<AQLPacket> inst_pkt         = {};
-        ClientID                   inst_pkt_id      = 0;
+        inst_pkt_t                 inst_pkt         = {};
        hsa_signal_t               interrupt_signal = {};
        rocprofiler_thread_id_t    tid              = common::get_tid();
        rocprofiler_kernel_id_t    kernel_id        = 0;
@@ -128,19 +135,16 @@ public:
    // the queue.
    using queue_cb_t = std::function<std::unique_ptr<AQLPacket>(
        const Queue&,
-        ClientID,
        const rocprofiler_packet&,
        uint64_t,
        const queue_info_session_t::external_corr_id_map_t&,
        const context::correlation_id*)>;
    // Signals the completion of the kernel packet.
    using completed_cb_t = std::function<void(const Queue&,
-                                              ClientID,
                                              const rocprofiler_packet&,
                                              const Queue::queue_info_session_t&,
-                                              std::unique_ptr<AQLPacket>)>;
+                                              inst_pkt_t&)>;
    using callback_map_t = std::unordered_map<ClientID, std::pair<queue_cb_t, completed_cb_t>>;
-
    Queue(const AgentCache&  agent,
          uint32_t           size,
          hsa_queue_type32_t type,
@@ -178,8 +182,14 @@ public:
    void register_callback(ClientID id, queue_cb_t enqueue_cb, completed_cb_t complete_cb);
    void remove_callback(ClientID id);

-    const CoreApiTable& core_api() const { return _core_api; }
-    const AmdExtTable&  ext_api() const { return _ext_api; }
+    const CoreApiTable&             core_api() const { return _core_api; }
+    const AmdExtTable&              ext_api() const { return _ext_api; }
+    mutable std::mutex              cv_mutex;
+    mutable std::condition_variable cv_ready_signal;
+    hsa_signal_t                    block_signal;
+    hsa_signal_t                    ready_signal;
+    queue_state                     get_state() const;
+    void                            set_state(queue_state state);

 private:
    std::atomic<int>                                  _notifiers            = {0};
@@ -189,6 +199,7 @@ private:
    const AgentCache&                                 _agent;
    rocprofiler::common::Synchronized<callback_map_t> _callbacks       = {};
    hsa_queue_t*                                      _intercept_queue = nullptr;
+    queue_state                                       _state           = queue_state::normal;
 };

 inline rocprofiler_queue_id_t
@@ -205,4 +216,4 @@ Queue::signal_callback(FuncT&& func) const
 }

 }  // namespace hsa
-}  // namespace rocprofiler
+}  // namespace rocprofiler
@@ -61,7 +61,11 @@ create_queue(hsa_agent_t        agent,
                                                     get_queue_controller().get_core_table(),
                                                     get_queue_controller().get_ext_table(),
                                                     queue);
+
+            get_queue_controller().profiler_serializer_register_ready_signal_handler(
+                new_queue->ready_signal, *queue);
            get_queue_controller().add_queue(*queue, std::move(new_queue));
+
            return HSA_STATUS_SUCCESS;
        }
    }
@@ -104,6 +108,48 @@ QueueController::add_queue(hsa_queue_t* id, std::unique_ptr<Queue> queue)
 void
 QueueController::destory_queue(hsa_queue_t* id)
 {
+    const auto*                  queue = get_queue_controller().get_queue(*id);
+    std::unique_lock<std::mutex> cvlock(queue->cv_mutex);
+    profiler_serializer([&](auto& data) {
+        /*Deletes the queue to be destructed from the dispatch ready.*/
+        data.dispatch_ready.erase(
+            std::remove_if(
+                data.dispatch_ready.begin(),
+                data.dispatch_ready.end(),
+                [&](auto& it) {
+                    /*Deletes the queue to be destructed from the dispatch ready.*/
+                    if(it->get_id().handle == queue->get_id().handle)
+                    {
+                        if(data.dispatch_queue &&
+                           data.dispatch_queue->get_id().handle == queue->get_id().handle)
+                        {
+                            // insert fatal condition here
+                            // ToDO [srnagara]: Need to find a solution rather than abort.
+                            LOG(FATAL)
+                                << "Queue is being destroyed while kernel launch is still active";
+                        }
+                        return true;
+                    }
+                    return false;
+                }),
+            data.dispatch_ready.end());
+        set_queue_state(queue_state::to_destroy, id);
+        /*
+          This lambda triggers the async ready handler.
+          The async ready handler then unregisters itself
+          and sets the queue state to done_destroy for which
+          the condition variable here is waiting for.
+        */
+        auto trigger_ready_async_handler = [queue]() {
+            get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(
+                queue->ready_signal, 0);
+        };
+        trigger_ready_async_handler();
+    });
+    queue->cv_ready_signal.wait(
+        cvlock, [queue] { return queue->get_state() == queue_state::done_destroy; });
+    if(queue->block_signal.handle != 0)
+        get_queue_controller().get_core_table().hsa_signal_destroy_fn(queue->block_signal);
    _queues.wlock([&](auto& map) { map.erase(id); });
 }

@@ -216,6 +262,103 @@ QueueController::get_queue(const hsa_queue_t& _hsa_queue) const
        _hsa_queue);
 }

+template <typename FuncT>
+void
+QueueController::profiler_serializer(FuncT&& lambda)
+{
+    _profiler_serializer.wlock(std::forward<FuncT>(lambda));
+}
+
+namespace
+{
+/*
+    Function name:  AsyncSignalReadyHandler
+    Argument:    hsa signal value for which the async handler was called
+                 and pointer to the data.
+    Description: This async handler is invoked when the queue is ready
+                 to launch a kernel. It first, resets the queue's ready signal to 1.
+                 It then checks if there is any queue which has a kernel currently dispatched.
+                 If yes, it pushes the queue to the dispatch ready else
+                 it enables the dispatch for the given queue.
+    Return :     It returns true since we need this handler to be invoked
+                 each time the queue's ready signal (used for entire queue) is set to 0.
+                 If we had a separate signal for every dispatch in the queue then we don't
+                 need this to be invoked more than once in which case we would return false.
+*/
+bool
+profiler_serializer_ready_signal_handler(hsa_signal_value_t /* signal_value */, void* data)
+{
+    auto*       hsa_queue = static_cast<hsa_queue_t*>(data);
+    const auto* queue     = get_queue_controller().get_queue(*hsa_queue);
+    get_queue_controller().profiler_serializer([&](auto& serializer) {
+        {
+            std::lock_guard<std::mutex> cv_lock(queue->cv_mutex);
+            if(queue->get_state() == queue_state::to_destroy)
+            {
+                get_queue_controller().set_queue_state(queue_state::done_destroy, hsa_queue);
+                get_queue_controller().get_core_table().hsa_signal_destroy_fn(queue->ready_signal);
+                queue->cv_ready_signal.notify_one();
+                return;
+            }
+        }
+        get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(queue->ready_signal,
+                                                                              1);
+        if(serializer.dispatch_queue == nullptr)
+        {
+            get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(
+                queue->block_signal, 0);
+            serializer.dispatch_queue = queue;
+        }
+        else
+        {
+            serializer.dispatch_ready.push_back(queue);
+        }
+    });
+    return true;
+}
+}  // namespace
+
+void
+profiler_serializer_kernel_completion_signal(hsa_signal_t queue_block_signal)
+{
+    get_queue_controller().profiler_serializer([queue_block_signal](auto& serializer) {
+        assert(serializer.dispatch_queue != nullptr);
+        serializer.dispatch_queue = nullptr;
+        get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(queue_block_signal,
+                                                                              1);
+        if(!serializer.dispatch_ready.empty())
+        {
+            auto queue = serializer.dispatch_ready.front();
+            serializer.dispatch_ready.erase(serializer.dispatch_ready.begin());
+            get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(
+                queue->block_signal, 0);
+            serializer.dispatch_queue = queue;
+        }
+    });
+}
+
+void
+QueueController::set_queue_state(enum queue_state state, hsa_queue_t* hsa_queue)
+{
+    _queues.wlock([&](auto& map) { map[hsa_queue]->set_state(state); });
+}
+
+/*
+    Function name:  SignalAsyncReadyHandler.
+    Argument :      The signal value and pointer to the data to
+                    pass to the handler.
+    Description :   Registers a asynchronous callback function
+                    for the ready signal to be invoked when it goes to zero.
+*/
+void
+QueueController::profiler_serializer_register_ready_signal_handler(const hsa_signal_t& signal,
+                                                                   void*               data) const
+{
+    hsa_status_t status = get_ext_table().hsa_amd_signal_async_handler_fn(
+        signal, HSA_SIGNAL_CONDITION_EQ, 0, profiler_serializer_ready_signal_handler, data);
+    if(status != HSA_STATUS_SUCCESS) LOG(FATAL) << "hsa_amd_signal_async_handler failed";
+}
+
 void
 QueueController::iterate_queues(const queue_iterator_cb_t& cb) const
 {
@@ -36,6 +36,26 @@ namespace rocprofiler
 {
 namespace hsa
 {
+/*This is a profiler serializer. It should be instantiated
+only once for the profiler. The following is the
+description of each field.
+1. dispatch_queue - The queue to which the currently dispatched kernel
+        belongs to.
+        At any given time, in serialization only one kernel
+        can be executing.
+2. dispatch_ready- It is a software data structure which holds
+        the queues which have a kernel ready to be dispatched.
+        This stores the queues in FIFO order.
+3. serializer_mutex - The mutex is used for thread synchronization
+        while accessing the singleton instance of this structure.
+Currently, in case of profiling kernels are serialized by default.
+*/
+struct profiler_serializer_t
+{
+    const Queue*             dispatch_queue{nullptr};
+    std::deque<const Queue*> dispatch_ready;
+};
+
 // Tracks and manages HSA queues
 class QueueController
 {
@@ -69,6 +89,12 @@ public:
    const Queue* get_queue(const hsa_queue_t&) const;

    void iterate_queues(const queue_iterator_cb_t&) const;
+    void set_queue_state(queue_state state, hsa_queue_t* hsa_queue);
+    void profiler_serializer_register_ready_signal_handler(const hsa_signal_t& signal,
+                                                           void*               data) const;
+    void add_dispatch_ready(const Queue* queue);
+    template <typename FuncT>
+    void profiler_serializer(FuncT&& lambda);

 private:
    using agent_callback_tuple_t =
@@ -77,11 +103,12 @@ private:
    using client_id_map_t   = std::unordered_map<ClientID, agent_callback_tuple_t>;
    using agent_cache_map_t = std::unordered_map<uint32_t, AgentCache>;

-    CoreApiTable                          _core_table       = {};
-    AmdExtTable                           _ext_table        = {};
-    common::Synchronized<queue_map_t>     _queues           = {};
-    common::Synchronized<client_id_map_t> _callback_cache   = {};
-    agent_cache_map_t                     _supported_agents = {};
+    CoreApiTable                                _core_table       = {};
+    AmdExtTable                                 _ext_table        = {};
+    common::Synchronized<queue_map_t>           _queues           = {};
+    common::Synchronized<client_id_map_t>       _callback_cache   = {};
+    agent_cache_map_t                           _supported_agents = {};
+    common::Synchronized<profiler_serializer_t> _profiler_serializer;
 };

 QueueController&
@@ -92,5 +119,9 @@ queue_controller_init(HsaApiTable* table);

 void
 queue_controller_fini();
+
+void
+profiler_serializer_kernel_completion_signal(hsa_signal_t queue_block_signal);
+
 }  // namespace hsa
 }  // namespace rocprofiler
@@ -49,3 +49,4 @@ add_subdirectory(async-copy-tracing)

 # rocprofv3 validation tests
 add_subdirectory(rocprofv3)
+add_subdirectory(counter-collection)
@@ -9,6 +9,7 @@ set(CMAKE_BUILD_RPATH "\$ORIGIN:\$ORIGIN/../lib")

 # applications used by integration tests which DO NOT link to rocprofiler-sdk-roctx
 add_subdirectory(simple-transpose)
+add_subdirectory(multistream)

 set(CMAKE_BUILD_RPATH
    "\$ORIGIN:\$ORIGIN/../lib:$<TARGET_FILE_DIR:rocprofiler-sdk-roctx::rocprofiler-sdk-roctx-shared-library>"
@@ -0,0 +1,46 @@
+#
+#
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+if(NOT CMAKE_HIP_COMPILER)
+    find_program(
+        amdclangpp_EXECUTABLE
+        NAMES amdclang++
+        HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
+        PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
+        PATH_SUFFIXES bin llvm/bin NO_CACHE)
+    mark_as_advanced(amdclangpp_EXECUTABLE)
+
+    if(amdclangpp_EXECUTABLE)
+        set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}")
+    endif()
+endif()
+
+project(rocprofiler-test-app-multistream LANGUAGES CXX HIP)
+
+foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO)
+    if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "")
+        set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}")
+    endif()
+endforeach()
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_EXTENSIONS OFF)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_HIP_STANDARD 17)
+set(CMAKE_HIP_EXTENSIONS OFF)
+set(CMAKE_HIP_STANDARD_REQUIRED ON)
+
+set_source_files_properties(multistream_app.cpp PROPERTIES LANGUAGE HIP)
+add_executable(multistream)
+target_sources(multistream PRIVATE multistream_app.cpp)
+target_compile_options(multistream PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow -Werror)
+
+find_package(Threads REQUIRED)
+target_link_libraries(multistream PRIVATE Threads::Threads)
+
+install(
+    TARGETS multistream
+    DESTINATION bin
+    COMPONENT tests)
@@ -0,0 +1,116 @@
+// 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 <hip/hip_runtime.h>
+#include <vector>
+#define HIP_ASSERT(call)                                                                           \
+    do                                                                                             \
+    {                                                                                              \
+        hipError_t err = call;                                                                     \
+        if(err != hipSuccess)                                                                      \
+        {                                                                                          \
+            fprintf(stderr, "%s\n", hipGetErrorString(err));                                       \
+            abort();                                                                               \
+        }                                                                                          \
+    } while(0)
+
+__device__ int counter = 0;
+__global__ void
+add(int n, float* x, float* y)
+{
+    if(__hip_atomic_load(&counter, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_AGENT) != 0)
+    {
+        abort();
+    }
+    __hip_atomic_fetch_add(&counter, 1, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
+
+    int index  = blockIdx.x * blockDim.x + threadIdx.x;
+    int stride = blockDim.x * gridDim.x;
+    for(int i = index; i < n; i += stride)
+        y[i] = x[i] + y[i];
+    __hip_atomic_fetch_add(&counter, -1, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+void
+LaunchMultiStreamKernels()
+{
+    int    N = 1 << 4;
+    float* x = new float[N];
+    float* y = new float[N];
+    float* d_x;
+    float* d_y;
+    //   Allocate Unified Memory -- accessible from CPU or GPU
+    HIP_ASSERT(hipMallocManaged(&d_x, N * sizeof(float)));
+    HIP_ASSERT(hipMallocManaged(&d_y, N * sizeof(float)));
+
+    //   initialize x and y arrays on the host
+    for(int i = 0; i < N; i++)
+    {
+        x[i] = 1.0f;
+        y[i] = 2.0f;
+    }
+    std::vector<hipStream_t> hip_streams;
+    for(int i = 0; i < 100; i++)
+    {
+        hipStream_t stream;
+        HIP_ASSERT(hipStreamCreate(&stream));
+        hip_streams.push_back(stream);
+    }
+    HIP_ASSERT(hipMemcpy(d_x, x, N * sizeof(float), hipMemcpyHostToDevice));
+    HIP_ASSERT(hipMemcpy(d_y, y, N * sizeof(float), hipMemcpyHostToDevice));
+
+    // Launch kernel on 1M elements on the GPU
+    int blockSize = 64;
+    // This Kernel will always be launched with one wave
+    int numBlocks = 1;
+    for(int i = 0; i < 100; i++)
+    {
+        for(size_t j = 0; j < hip_streams.size(); j++)
+        {
+            hipLaunchKernelGGL(add, numBlocks, blockSize, 0, hip_streams[j], N, d_x, d_y);
+        }
+    }
+
+    // Wait for GPU to finish before accessing on host
+    HIP_ASSERT(hipDeviceSynchronize());
+
+    HIP_ASSERT(hipMemcpy(x, d_x, N * sizeof(float), hipMemcpyDeviceToHost));
+    HIP_ASSERT(hipMemcpy(y, d_y, N * sizeof(float), hipMemcpyDeviceToHost));
+
+    //   Free memory
+    HIP_ASSERT(hipFree(d_x));
+    HIP_ASSERT(hipFree(d_y));
+
+    delete[] x;
+    delete[] y;
+
+    for(size_t i = 0; i < hip_streams.size(); i++)
+    {
+        HIP_ASSERT(hipStreamDestroy(hip_streams[i]));
+    }
+}
+
+int
+main()
+{
+    LaunchMultiStreamKernels();
+}
@@ -231,6 +231,15 @@ save(ArchiveT& ar, rocprofiler_buffer_tracing_hsa_api_record_t data)
    save_buffer_tracing_api_record(ar, data);
 }

+template <typename ArchiveT>
+void
+save(ArchiveT& ar, rocprofiler_record_counter_t data)
+{
+    SAVE_DATA_FIELD(id);
+    SAVE_DATA_FIELD(counter_value);
+    SAVE_DATA_FIELD(corr_id);
+}
+
 template <typename ArchiveT>
 void
 save(ArchiveT& ar, rocprofiler_buffer_tracing_hip_api_record_t data)
@@ -0,0 +1,47 @@
+#
+#
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+project(
+    rocprofiler-tests-counter-collection
+    LANGUAGES CXX
+    VERSION 0.0.0)
+
+find_package(rocprofiler-sdk REQUIRED)
+
+if(ROCPROFILER_MEMCHECK_PRELOAD_ENV)
+    set(PRELOAD_ENV
+        "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}:$<TARGET_FILE:rocprofiler-sdk-json-tool>")
+else()
+    set(PRELOAD_ENV "LD_PRELOAD=$<TARGET_FILE:rocprofiler-sdk-json-tool>")
+endif()
+
+add_test(NAME test-counter-collection-execute COMMAND $<TARGET_FILE:multistream>)
+
+set_tests_properties(
+    test-counter-collection-execute
+    PROPERTIES
+        TIMEOUT
+        45
+        LABELS
+        "integration-tests"
+        ENVIRONMENT
+        "${PRELOAD_ENV};HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>;ROCPROFILER_TOOL_OUTPUT_FILE=counter-collection-test.json;ROCPROFILER_TOOL_CONTEXTS=COUNTER_COLLECTION;ROCPROF_COUNTERS=SQ_WAVES_sum"
+        FAIL_REGULAR_EXPRESSION
+        "threw an exception")
+
+foreach(FILENAME validate.py pytest.ini conftest.py)
+    configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME}
+                   ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY)
+endforeach()
+
+add_test(NAME test-counter-collection-validate
+         COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input
+                 ${CMAKE_CURRENT_BINARY_DIR}/counter-collection-test.json)
+
+set_tests_properties(
+    test-counter-collection-validate
+    PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS
+               test-counter-collection-execute FAIL_REGULAR_EXPRESSION
+               "threw an exception")
@@ -0,0 +1,20 @@
+#!/usr/bin/env python3
+
+import json
+import pytest
+
+
+def pytest_addoption(parser):
+    parser.addoption(
+        "--input",
+        action="store",
+        default="counter-collection-test.json",
+        help="Input JSON",
+    )
+
+
+@pytest.fixture
+def input_data(request):
+    filename = request.config.getoption("--input")
+    with open(filename, "r") as inp:
+        return json.load(inp)
@@ -0,0 +1,4 @@
+
+[pytest]
+addopts = --durations=20 -ras -vv
+testpaths = validate.py
@@ -0,0 +1,34 @@
+#!/usr/bin/env python3
+
+import sys
+import pytest
+
+
+# helper function
+def node_exists(name, data, min_len=1):
+    assert name in data
+    assert data[name] is not None
+    assert len(data[name]) >= min_len
+
+
+def test_data_structure(input_data):
+    """verify minimum amount of expected data is present"""
+    node_exists("rocprofiler-sdk-json-tool", input_data)
+    rocp_data = input_data
+    node_exists("names", rocp_data["rocprofiler-sdk-json-tool"]["buffer_records"])
+    node_exists(
+        "counter_collection", rocp_data["rocprofiler-sdk-json-tool"]["buffer_records"]
+    )
+
+
+def test_counter_values(input_data):
+    data = input_data
+    counter_info = {}
+    for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["counter_collection"]:
+        value = itr["counter_value"]
+        assert value == 1
+
+
+if __name__ == "__main__":
+    exit_code = pytest.main(["-x", __file__] + sys.argv[1:])
+    sys.exit(exit_code)
@@ -11,7 +11,7 @@ project(
 find_package(rocprofiler-sdk REQUIRED)

 set(PYTEST_ARGS)
-if(ROCPROFILER_MEMCHECK MATCHES "(Address|Thread)Sanitizer")
+if(ROCPROFILER_MEMCHECK MATCHES "(Address|Thread)Sanitizer" OR ROCPROFILER_BUILD_CODECOV)
    set(PYTEST_ARGS -k "not test_total_runtime")
 endif()

@@ -43,6 +43,8 @@
 #include <rocprofiler-sdk/registration.h>
 #include <rocprofiler-sdk/rocprofiler.h>

+#include <glog/logging.h>
+
 #include <unistd.h>
 #include <atomic>
 #include <cassert>
@@ -57,6 +59,7 @@
 #include <iostream>
 #include <map>
 #include <mutex>
+#include <shared_mutex>
 #include <string>
 #include <string_view>
 #include <thread>
@@ -361,15 +364,123 @@ struct marker_api_callback_record_t
    }
 };

-auto code_object_records   = std::deque<code_object_callback_record_t>{};
-auto kernel_symbol_records = std::deque<kernel_symbol_callback_record_t>{};
-auto hsa_api_cb_records    = std::deque<hsa_api_callback_record_t>{};
-auto marker_api_cb_records = std::deque<marker_api_callback_record_t>{};
-auto hip_api_cb_records    = std::deque<hip_api_callback_record_t>{};
+auto code_object_records           = std::deque<code_object_callback_record_t>{};
+auto kernel_symbol_records         = std::deque<kernel_symbol_callback_record_t>{};
+auto hsa_api_cb_records            = std::deque<hsa_api_callback_record_t>{};
+auto marker_api_cb_records         = std::deque<marker_api_callback_record_t>{};
+auto counter_collection_bf_records = std::deque<rocprofiler_record_counter_t>{};
+auto hip_api_cb_records            = std::deque<hip_api_callback_record_t>{};

 rocprofiler_thread_id_t
 push_external_correlation();

+void
+counter_collection_buffered(rocprofiler_context_id_t, /*context*/
+                            rocprofiler_buffer_id_t,  /*buffer_id*/
+                            rocprofiler_record_header_t** headers,
+                            size_t                        num_headers,
+                            void*, /*user_data*/
+                            uint64_t /*drop_count*/)
+{
+    if(num_headers == 0)
+        throw std::runtime_error{"rocprofiler invoked a buffer callback with no headers "
+                                 "this should never happen"};
+
+    else if(headers == nullptr)
+        throw std::runtime_error{"rocprofiler invoked a buffer callback with a null pointer to the "
+                                 "array of headers. this should never happen"};
+    for(size_t i = 0; i < num_headers; ++i)
+    {
+        auto* header = headers[i];
+        if(header->category == ROCPROFILER_BUFFER_CATEGORY_COUNTERS && header->kind == 0)
+        {
+            auto* profiler_record = static_cast<rocprofiler_record_counter_t*>(header->payload);
+            counter_collection_bf_records.emplace_back(*profiler_record);
+        }
+    }
+}
+
+void
+dispatch_callback(rocprofiler_queue_id_t, /*queue_id*/
+                  const rocprofiler_agent_t* agent,
+                  rocprofiler_correlation_id_t,        /*correlation_id*/
+                  const hsa_kernel_dispatch_packet_t*, /*dispatch_packet*/
+                  uint64_t,                            /*kernel_id*/
+                  void* /*callback_data_args*/,
+                  rocprofiler_profile_config_id_t* config)
+{
+    static std::shared_mutex                                             m_mutex       = {};
+    static std::unordered_map<uint64_t, rocprofiler_profile_config_id_t> profile_cache = {};
+
+    auto search_cache = [&]() {
+        if(auto pos = profile_cache.find(agent->id.handle); pos != profile_cache.end())
+        {
+            *config = pos->second;
+            return true;
+        }
+        return false;
+    };
+
+    {
+        auto rlock = std::shared_lock{m_mutex};
+        if(search_cache()) return;
+    }
+
+    auto wlock = std::unique_lock{m_mutex};
+    if(search_cache()) return;
+
+    // Counters we want to collect (here its SQ_WAVES_sum)
+    auto* counters_env = getenv("ROCPROF_COUNTERS");
+    if(std::string(counters_env) != "SQ_WAVES_sum")
+        LOG(FATAL) << "Counter not supported in the test tool";
+
+    std::set<std::string> counters_to_collect = {"SQ_WAVES_sum"};
+    // GPU Counter IDs
+    std::vector<rocprofiler_counter_id_t> gpu_counters;
+
+    // Iterate through the agents and get the counters available on that agent
+    ROCPROFILER_CALL(rocprofiler_iterate_agent_supported_counters(
+                         agent->id,
+                         []([[maybe_unused]] rocprofiler_agent_id_t id,
+                            rocprofiler_counter_id_t*               counters,
+                            size_t                                  num_counters,
+                            void*                                   user_data) {
+                             std::vector<rocprofiler_counter_id_t>* vec =
+                                 static_cast<std::vector<rocprofiler_counter_id_t>*>(user_data);
+                             for(size_t i = 0; i < num_counters; i++)
+                             {
+                                 vec->push_back(counters[i]);
+                             }
+                             return ROCPROFILER_STATUS_SUCCESS;
+                         },
+                         static_cast<void*>(&gpu_counters)),
+                     "Could not fetch supported counters");
+
+    std::vector<rocprofiler_counter_id_t> collect_counters;
+    // Look for the counters contained in counters_to_collect in gpu_counters
+    for(auto& counter : gpu_counters)
+    {
+        const char* name;
+        size_t      size;
+        ROCPROFILER_CALL(rocprofiler_query_counter_name(counter, &name, &size),
+                         "Could not query name");
+        if(counters_to_collect.count(std::string(name)) > 0)
+        {
+            collect_counters.push_back(counter);
+        }
+    }
+
+    // Create a colleciton profile for the counters
+    rocprofiler_profile_config_id_t profile;
+    ROCPROFILER_CALL(rocprofiler_create_profile_config(
+                         agent->id, collect_counters.data(), collect_counters.size(), &profile),
+                     "Could not construct profile cfg");
+
+    profile_cache.emplace(agent->id.handle, profile);
+    // Return the profile to collect those counters for this dispatch
+    *config = profile;
+}
+
 void
 tool_tracing_callback(rocprofiler_callback_tracing_record_t record,
                      rocprofiler_user_data_t* /*user_data*/,
@@ -584,12 +695,14 @@ rocprofiler_context_id_t hip_api_buffered_ctx    = {};
 rocprofiler_context_id_t marker_api_buffered_ctx = {};
 rocprofiler_context_id_t kernel_dispatch_ctx     = {};
 rocprofiler_context_id_t memory_copy_ctx         = {};
+rocprofiler_context_id_t counter_collection_ctx  = {};
 // buffers
 rocprofiler_buffer_id_t hsa_api_buffered_buffer    = {};
 rocprofiler_buffer_id_t hip_api_buffered_buffer    = {};
 rocprofiler_buffer_id_t marker_api_buffered_buffer = {};
 rocprofiler_buffer_id_t kernel_dispatch_buffer     = {};
 rocprofiler_buffer_id_t memory_copy_buffer         = {};
+rocprofiler_buffer_id_t counter_collection_buffer  = {};

 auto contexts = std::unordered_map<std::string_view, rocprofiler_context_id_t*>{
    {"HSA_API_CALLBACK", &hsa_api_callback_ctx},
@@ -600,13 +713,15 @@ auto contexts = std::unordered_map<std::string_view, rocprofiler_context_id_t*>{
    {"HIP_API_BUFFERED", &hip_api_buffered_ctx},
    {"MARKER_API_BUFFERED", &marker_api_buffered_ctx},
    {"KERNEL_DISPATCH", &kernel_dispatch_ctx},
-    {"MEMORY_COPY", &memory_copy_ctx}};
+    {"MEMORY_COPY", &memory_copy_ctx},
+    {"COUNTER_COLLECTION", &counter_collection_ctx}};

-auto buffers = std::array<rocprofiler_buffer_id_t*, 5>{&hsa_api_buffered_buffer,
+auto buffers = std::array<rocprofiler_buffer_id_t*, 6>{&hsa_api_buffered_buffer,
                                                       &hip_api_buffered_buffer,
                                                       &marker_api_buffered_buffer,
                                                       &kernel_dispatch_buffer,
-                                                       &memory_copy_buffer};
+                                                       &memory_copy_buffer,
+                                                       &counter_collection_buffer};

 auto agents = std::vector<rocprofiler_agent_t>{};

@@ -785,6 +900,20 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
                                                     marker_api_buffered_buffer),
        "buffer tracing service configure");

+    ROCPROFILER_CALL(rocprofiler_create_buffer(counter_collection_ctx,
+                                               4096,
+                                               2048,
+                                               ROCPROFILER_BUFFER_POLICY_LOSSLESS,
+                                               counter_collection_buffered,
+                                               nullptr,
+                                               &counter_collection_buffer),
+                     "buffer creation");
+
+    ROCPROFILER_CALL(
+        rocprofiler_configure_buffered_dispatch_profile_counting_service(
+            counter_collection_ctx, counter_collection_buffer, dispatch_callback, nullptr),
+        "setup buffered service");
+
    ROCPROFILER_CALL(
        rocprofiler_configure_buffer_tracing_service(kernel_dispatch_ctx,
                                                     ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH,
@@ -885,7 +1014,8 @@ tool_fini(void* tool_data)
              << ", memory_copy_records=" << memory_copy_records.size()
              << ", hsa_api_bf_records=" << hsa_api_bf_records.size()
              << ", hip_api_bf_records=" << hip_api_bf_records.size()
-              << ", marker_api_bf_records=" << marker_api_bf_records.size() << " ...\n"
+              << ", marker_api_bf_records=" << marker_api_bf_records.size()
+              << ", counter_collection_records" << counter_collection_bf_records.size() << "...\n"
              << std::flush;

    auto* _call_stack = static_cast<call_stack_t*>(tool_data);
@@ -966,6 +1096,7 @@ tool_fini(void* tool_data)
            json_ar(cereal::make_nvp("hsa_api_traces", hsa_api_bf_records));
            json_ar(cereal::make_nvp("hip_api_traces", hip_api_bf_records));
            json_ar(cereal::make_nvp("marker_api_traces", marker_api_bf_records));
+            json_ar(cereal::make_nvp("counter_collection", counter_collection_bf_records));
        } catch(std::exception& e)
        {
            std::cerr << "[" << getpid() << "][" << __FUNCTION__
@@ -994,6 +1125,11 @@ start()
    {
        if(itr.second && !is_active(*itr.second))
        {
+            if(itr.first == "COUNTER_COLLECTION")
+            {
+                auto* counters = getenv("ROCPROF_COUNTERS");
+                if(!counters) continue;
+            }
            ROCPROFILER_CALL(rocprofiler_start_context(*itr.second), "context start");
        }
    }