AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Drop signal limiter (#382)

Browse Source 
* Drop signal limiter

* source formatting (clang-format v11) (#383)

Co-authored-by: bwelton <bwelton@users.noreply.github.com>

---------

Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: bwelton <bwelton@users.noreply.github.com>

[ROCm/rocprofiler-sdk commit: e0fd4c1716]
This commit is contained in:
Benjamin Welton
2024-01-16 17:43:24 -08:00
committed by GitHub
parent 3c47bc534a
commit 099db3bab7
1 changed files with 1 additions and 18 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp
+1 -18
View File
@@ -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<size_t>("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";