[rocprofiler-sdk] On-demand GPU profile queue creation/destruction

Create and destroy GPU hardware profile queues per start_context/stop_context
cycle instead of permanently at SDK initialization. This eliminates ~70ms TTFT
overhead caused by 8 persistent GPU queues in the MES runlist.

Changes:
- agent_cache.hpp: Add destroy_device_counting_service_queue() declaration
- agent_cache.cpp: Defer queue creation from constructor to start_context;
  implement destroy_device_counting_service_queue() with hsa_queue_destroy
- device_counting.cpp: Call init_device_counting_service_queue() in
  start_agent_ctx before queue check; destroy signals, reset packet, and
  destroy queue in stop_agent_ctx after completion

Test results (10 rounds, Qwen3-VL-235B on MI300X x8 with old in-process RDC):
- Baseline (no RDC): Mean TTFT 889.67ms
- Old RDC (permanent queues): Mean TTFT 972.63ms (+83ms, +9.3%)
- On-demand queues: Mean TTFT 854.60ms (-35ms, -3.9% vs baseline)

Author: chun-wan

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/device_counting.cpp

@@ -432,6 +432,9 @@ start_agent_ctx(const context::context* ctx)
             break;
         }
 
+        // On-demand: create the profile queue now (destroyed in stop_agent_ctx)
+        agent->init_device_counting_service_queue(*hsa::get_core_table(), *hsa::get_amd_ext_table());
+
         // But if we have an agent cache, we need a profile queue.
         if(!agent->profile_queue())
         {
@@ -603,6 +606,19 @@ stop_agent_ctx(const context::context* ctx)
         {
             counters::counter_collection_device_unlock(agent->get_rocp_agent());
         }
+
+        // On-demand cleanup: destroy signals, reset packet, destroy queue
+        if(callback_data.completion.handle != 0) {
+            hsa::get_core_table()->hsa_signal_destroy_fn(callback_data.completion);
+            callback_data.completion.handle = 0;
+        }
+        if(callback_data.start_signal.handle != 0) {
+            hsa::get_core_table()->hsa_signal_destroy_fn(callback_data.start_signal);
+            callback_data.start_signal.handle = 0;
+        }
+        callback_data.packet.reset();
+        callback_data.queue = nullptr;
+        agent->destroy_device_counting_service_queue();
     }
 
     agent_ctx.status.exchange(rocprofiler::context::device_counting_service::state::DISABLED);

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/agent_cache.cpp

@@ -27,6 +27,7 @@
 #include <stdexcept>
 
 #include "lib/rocprofiler-sdk/context/context.hpp"
+#include "lib/rocprofiler-sdk/hsa/hsa.hpp"
 
 namespace
 {
@@ -171,7 +172,7 @@ AgentCache::AgentCache(const rocprofiler_agent_t* rocp_agent,
     {
         init_cpu_pool(ext_table, *this);
         init_gpu_pool(ext_table, *this);
-        init_device_counting_service_queue(api, ext_table);
+        // init_device_counting_service_queue(api, ext_table);  // Deferred to start_context for on-demand queue
     } catch(std::runtime_error& e)
     {
         ROCP_WARNING << fmt::format(
@@ -181,5 +182,20 @@ AgentCache::AgentCache(const rocprofiler_agent_t* rocp_agent,
     }
 }
 
+
+void
+AgentCache::destroy_device_counting_service_queue() const
+{
+    static std::mutex           m_mutex;
+    std::lock_guard<std::mutex> lock(m_mutex);
+
+    if(!m_profile_queue) return;
+
+    auto* api = hsa::get_core_table();
+    if(api && api->hsa_queue_destroy_fn)
+        api->hsa_queue_destroy_fn(m_profile_queue);
+    m_profile_queue = nullptr;
+}
+
 }  // namespace hsa
 }  // namespace rocprofiler

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/agent_cache.hpp

@@ -74,6 +74,7 @@ class AgentCache
     size_t                     index() const { return m_index; }
 
     void init_device_counting_service_queue(const CoreApiTable& api, const AmdExtTable& ext) const;
+    void destroy_device_counting_service_queue() const;
     bool operator==(const rocprofiler_agent_t*) const;
     bool operator==(hsa_agent_t) const;