diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp index 7a8d93df0f..9e5bf71bb8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -66,16 +66,6 @@ namespace hsa { namespace { -common::active_capacity_gate& -signal_limiter() -{ - // Limit the maximun number of HSA signals created. - // There is a hard limit to the maximum that can exist. - static auto _gate = - common::active_capacity_gate{common::get_env("ROCPROFILER_GATE_CAPACITY", 4)}; - return _gate; -} - bool context_filter(const context::context* ctx) { @@ -192,24 +182,18 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data) } }); - size_t signals_to_remove = 0; // Delete signals and packets, signal we have completed. if(queue_info_session.interrupt_signal.handle != 0u) { - signals_to_remove++; queue_info_session.queue.core_api().hsa_signal_destroy_fn( queue_info_session.interrupt_signal); } if(queue_info_session.kernel_pkt.ext_amd_aql_pm4.completion_signal.handle != 0u) { - signals_to_remove++; queue_info_session.queue.core_api().hsa_signal_destroy_fn( queue_info_session.kernel_pkt.ext_amd_aql_pm4.completion_signal); } - if(signals_to_remove > 0) - { - signal_limiter().remove_active(signals_to_remove); - } + queue_info_session.queue.async_complete(); if(_corr_id) @@ -443,7 +427,6 @@ Queue::signal_async_handler(const hsa_signal_t& signal, Queue::queue_info_sessio void Queue::create_signal(uint32_t attribute, hsa_signal_t* signal) const { - signal_limiter().add_active(1); hsa_status_t status = _ext_api.hsa_amd_signal_create_fn(1, 0, nullptr, attribute, signal); LOG_IF(FATAL, status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) << "Error: hsa_amd_signal_create failed";