From 2fa98c4d14f3c28d7ca86861291fbe2ea7ce0f72 Mon Sep 17 00:00:00 2001 From: "Welton, Benjamin" Date: Tue, 8 Jul 2025 23:28:02 -0700 Subject: [PATCH] [SDK] Added retry for signal wait (#494) Added retry for signal wait Co-authored-by: Benjamin Welton --- source/lib/rocprofiler-sdk/hsa/queue.cpp | 20 +++++++++++++------- 1 file changed, 13 insertions(+), 7 deletions(-) diff --git a/source/lib/rocprofiler-sdk/hsa/queue.cpp b/source/lib/rocprofiler-sdk/hsa/queue.cpp index c9d977205a..bf691443e3 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -569,15 +569,21 @@ Queue::Queue(const AgentCache& agent, counters::submitPacket(_intercept_queue, &pkt); constexpr auto timeout_hint = std::chrono::duration_cast(std::chrono::seconds{1}); - if(core_api.hsa_signal_wait_relaxed_fn(completion, - HSA_SIGNAL_CONDITION_EQ, - 0, - timeout_hint.count(), - HSA_WAIT_STATE_ACTIVE) != 0) + hsa_signal_value_t val; + for(int i = 0; i < 3; i++) { - ROCP_FATAL << "Could not set agent to be profiled"; + val = core_api.hsa_signal_wait_scacquire_fn(completion, + HSA_SIGNAL_CONDITION_EQ, + 0, + timeout_hint.count(), + HSA_WAIT_STATE_ACTIVE); + if(val == 0) + { + core_api.hsa_signal_destroy_fn(completion); + return; + } } - core_api.hsa_signal_destroy_fn(completion); + ROCP_FATAL << "Could not set agent to be profiled - Signal Value: " << val; }); }