From 8ba7120b630626228be165f536d0538af81efc2e Mon Sep 17 00:00:00 2001 From: itrowbri Date: Tue, 16 Sep 2025 10:35:16 -0500 Subject: [PATCH] [rocprofiler-sdk] Verify there are callbacks for every kernel dispatch when syncing (#321) * Added check in Queue::sync to verify that there is a callback for every dispatch * Removed new atomic, using get_balanced_signal_slots() atomic with initial value of NUM_SIGNALS to verify dispatches complete --- .../source/lib/rocprofiler-sdk/hsa/queue.cpp | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) 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 bf691443e3..7e431fbf2f 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -68,11 +68,11 @@ namespace hsa { namespace { +constexpr int64_t NUM_SIGNALS = 16; std::atomic& get_balanced_signal_slots() { - constexpr int64_t NUM_SIGNALS = 16; - static auto*& atomic = common::static_object>::construct(NUM_SIGNALS); + static auto*& atomic = common::static_object>::construct(NUM_SIGNALS); return *atomic; } @@ -638,6 +638,11 @@ Queue::sync() const _core_api.hsa_signal_wait_relaxed_fn( _active_kernels, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); } + // get_balanced_signal_slots() increments upon kernel dispatch completion and decrements in + // WriteInterceptor with a starting value of NUM_SIGNALS, so the get_balanced_signal_slots() + // should be equivalent to NUM_SIGNALS if all kernel dispatches are completed + ROCP_CI_LOG_IF(WARNING, get_balanced_signal_slots().load() != NUM_SIGNALS) << fmt::format( + "There are {} incomplete dispatches", NUM_SIGNALS - get_balanced_signal_slots().load()); } void