[rocprofiler-sdk] Add ROCPROFILER_ONDEMAND_QUEUE env flag for on-demand queue control

Gate the on-demand queue creation/destruction behavior behind the
ROCPROFILER_ONDEMAND_QUEUE environment variable. When set to 1 or true,
profile queues are created in start_context and destroyed in stop_context.
When unset (default), the original behavior is preserved with persistent
queues created during AgentCache initialization.

This allows controlled rollout and A/B testing of the on-demand queue
optimization that eliminates ~130ms TTFT overhead from persistent GPU
queues in the MES runlist.

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

Author: bwelton

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

@@ -433,7 +433,11 @@ start_agent_ctx(const context::context* ctx)
         }
 
         // 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());
+        if(hsa::use_ondemand_queue())
+        {
+            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())
@@ -608,17 +612,22 @@ stop_agent_ctx(const context::context* ctx)
         }
 
         // 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;
+        if(hsa::use_ondemand_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();
         }
-        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

@@ -21,6 +21,7 @@
 // THE SOFTWARE.
 
 #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp"
+#include "lib/common/environment.hpp"
 #include "lib/common/logging.hpp"
 
 #include <fmt/core.h>
@@ -120,6 +121,13 @@ namespace rocprofiler
 {
 namespace hsa
 {
+bool
+use_ondemand_queue()
+{
+    static bool value = rocprofiler::common::get_env("ROCPROFILER_ONDEMAND_QUEUE", false);
+    return value;
+}
+
 void
 AgentCache::init_device_counting_service_queue(const CoreApiTable& api,
                                                const AmdExtTable&  ext) const
@@ -172,7 +180,8 @@ 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);  // Deferred to start_context for on-demand queue
+        // When on-demand queue mode is enabled, defer queue creation to start_context
+        if(!use_ondemand_queue()) init_device_counting_service_queue(api, ext_table);
     } catch(std::runtime_error& e)
     {
         ROCP_WARNING << fmt::format(
@@ -182,7 +191,6 @@ AgentCache::AgentCache(const rocprofiler_agent_t* rocp_agent,
     }
 }
 
-
 void
 AgentCache::destroy_device_counting_service_queue() const
 {
@@ -192,8 +200,7 @@ AgentCache::destroy_device_counting_service_queue() const
     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);
+    if(api && api->hsa_queue_destroy_fn) api->hsa_queue_destroy_fn(m_profile_queue);
     m_profile_queue = nullptr;
 }
 

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

@@ -107,5 +107,11 @@ AgentCache::operator==(hsa_agent_t agent) const
 {
     return (agent.handle == m_hsa_agent.handle);
 }
+/// Returns true when ROCPROFILER_ONDEMAND_QUEUE=1|true is set.
+/// In on-demand mode, the profile queue is created in start_context
+/// and destroyed in stop_context instead of being persistent.
+bool
+use_ondemand_queue();
+
 }  // namespace hsa
 }  // namespace rocprofiler