From d0682a1b2d6996f8ed97827df731c30bc979e0c9 Mon Sep 17 00:00:00 2001 From: Venkateshwar Reddy Kandula Date: Wed, 3 Jun 2026 16:06:16 +0000 Subject: [PATCH 1/4] fix(serializer): prevent barrier signal-handle reuse hang A serializer barrier could be retired (its HSA signal destroyed) while a queue still held an un-executed transition packet on that signal. HSA recycled the freed handle into the next barrier, so an armed barrier stranded the queue's stale packet -> deadlock (seen with roctx pause/resume). Defer retirement until every transition packet has executed: tag each packet with its dispatch id, clear the tag on drain/teardown, and gate retirement on safe_to_retire(). Destructor leaks a still-referenced signal as a backstop. Adds regression tests. --- .../lib/rocprofiler-sdk/hsa/hsa_barrier.cpp | 57 +++++-- .../lib/rocprofiler-sdk/hsa/hsa_barrier.hpp | 16 +- .../hsa/profile_serializer.cpp | 8 +- .../source/lib/rocprofiler-sdk/hsa/queue.cpp | 4 + .../source/lib/rocprofiler-sdk/hsa/queue.hpp | 15 ++ .../lib/rocprofiler-sdk/tests/hsa_barrier.cpp | 144 ++++++++++++++++++ 6 files changed, 228 insertions(+), 16 deletions(-) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp index a99d0acb0113..90a3cbdc719b 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp @@ -45,20 +45,24 @@ hsa_barrier::~hsa_barrier() return; } - // If this barrier is destroyed and has queued packets, a future signal creation may - // inadvertently block these packets if the same handle is re-used by HSA. - if(!complete()) + bool outstanding = false; + _barrier_enqueued.rlock( + [&](const auto& barrier_enqueued) { outstanding = !barrier_enqueued.empty(); }); + + // release the signal so any stale transition packet can pass + clear_barrier(); + + // Backstop: retirement is gated on safe_to_retire(), so this should not happen. If a packet + // still references this signal, leak it rather than let HSA recycle the handle into a barrier. + if(outstanding) { - _barrier_enqueued.rlock([&](auto& barrier_enqueued) { - if(!barrier_enqueued.empty()) - { - ROCP_WARNING << "An incomplete hsa_barrier which was enqueued is being destroyed."; - } - }); + ROCP_WARNING << "hsa_barrier (handle: " << _barrier_signal.handle + << ") destroyed with outstanding transition packets; leaking its signal to " + "avoid handle reuse."; + return; } - // Clear and destroy the barrier signal - clear_barrier(); + // Destroy the barrier signal _core_api.hsa_signal_destroy_fn(_barrier_signal); } @@ -93,7 +97,8 @@ hsa_barrier::enqueue_packet(const Queue* queue) if(barrier_enqueued.find(queue->get_id().handle) == barrier_enqueued.end()) { return_block = true; - barrier_enqueued.insert(queue->get_id().handle); + // stamp the queue with the dispatch id that carries this transition packet + barrier_enqueued.emplace(queue->get_id().handle, queue->serialized_dispatched()); } }); @@ -107,9 +112,27 @@ hsa_barrier::enqueue_packet(const Queue* queue) return barrier; } +void +hsa_barrier::notify_drain(const Queue* queue) +{ + _barrier_enqueued.wlock([&](auto& barrier_enqueued) { + auto it = barrier_enqueued.find(queue->get_id().handle); + if(it == barrier_enqueued.end()) return; + // packet executed once completions reach the carrying dispatch id (queue is in-order) + if(queue->serialized_completed() >= it->second) + { + barrier_enqueued.erase(it); + } + }); +} + void hsa_barrier::remove_queue(const Queue* queue) { + // a torn-down queue can't execute its packet, so release its pin too + _barrier_enqueued.wlock( + [&](auto& barrier_enqueued) { barrier_enqueued.erase(queue->get_id().handle); }); + _queue_waiting.wlock([&](auto& queue_waiting) { if(queue_waiting.find(queue->get_id().handle) == queue_waiting.end()) return; queue_waiting.erase(queue->get_id().handle); @@ -146,6 +169,16 @@ hsa_barrier::complete() const return _core_api.hsa_signal_load_scacquire_fn(_barrier_signal) == 0; } +bool +hsa_barrier::safe_to_retire() const +{ + if(!complete()) return false; + bool no_outstanding = false; + _barrier_enqueued.rlock( + [&](const auto& barrier_enqueued) { no_outstanding = barrier_enqueued.empty(); }); + return no_outstanding; +} + void hsa_barrier::clear_barrier() { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp index 9c78fe75b4bb..28755f8f84b9 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp @@ -35,7 +35,6 @@ #include #include #include -#include namespace rocprofiler { @@ -65,18 +64,29 @@ class hsa_barrier // barrier and marks it as complete. Returns true if the given queue had async packets waiting. bool register_completion(const Queue* queue); + // Drop the queue from the outstanding set once it has executed past the dispatch that + // carried this barrier's transition packet (called on each serialized completion). + void notify_drain(const Queue* queue); + // Checks if this barrier is complete bool complete() const; - // Removes a queue from the barrier dependency list. + // complete() AND no queue still references the signal -> safe to destroy (no handle reuse). + bool safe_to_retire() const; + + // Removes a queue from the barrier dependency list (waiting + outstanding transition packets). // If this is the last queue waiting, clears the barrier and marks it as complete. void remove_queue(const Queue* queue); + // Handle of the underlying barrier signal. + uint64_t signal_handle() const { return _barrier_signal.handle; } + private: std::function _barrier_finished = {}; CoreApiTable _core_api = {}; common::Synchronized> _queue_waiting = {}; - common::Synchronized> _barrier_enqueued = {}; + // queue id -> serialized-dispatch id carrying this barrier's transition packet + common::Synchronized> _barrier_enqueued = {}; void clear_barrier(); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp index 94d48897183c..9c6564f458bb 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp @@ -45,9 +45,11 @@ profiler_serializer_ready_signal_handler(hsa_signal_value_t /* signal_value */, void clear_complete_barriers(std::deque& barriers) { + // Retire only when safe_to_retire(): complete AND no queue still references the signal. + // Otherwise the barrier stays parked (signal at 0, stale packets pass) until its packets drain. while(!barriers.empty()) { - if(barriers.front().barrier->complete()) + if(barriers.front().barrier->safe_to_retire()) { barriers.pop_front(); } @@ -78,6 +80,10 @@ profiler_serializer::add_queue(hsa_queue_t** hsa_queues, const Queue& queue) void profiler_serializer::kernel_completion_signal(const Queue& completed) { + // let each barrier drop `completed` if its transition packet on that queue has now executed + for(auto& barrier : _barrier) + barrier.barrier->notify_drain(&completed); + // We do not want to track kernel compleiton signals before we have reached the barrier clear_complete_barriers(_barrier); 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 5ef8a86b47b5..425469243ab2 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -202,6 +202,8 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data) if(packet.is_serialized) { + // count this completion before notifying the serializer (drives barrier drain detection) + queue_info_session.queue.serialized_completed_inc(); CHECK_NOTNULL(hsa::get_queue_controller()) ->serializer(&queue_info_session.queue) .wlock([&](auto& serializer) { @@ -631,6 +633,8 @@ WriteInterceptor(const void* packets, if(_packet_data.is_serialized) { inserted_before = true; + // count this dispatch before enqueue_packet() stamps the barrier with its id + queue.serialized_dispatched_inc(); CHECK_NOTNULL(hsa::get_queue_controller()) ->serializer(&queue) .rlock([&](const auto& serializer) { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp index a27f51cbb94c..7dc1f2b69f83 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -170,6 +170,19 @@ class Queue { return _core_api.hsa_signal_load_scacquire_fn(_active_kernels); } + + // In-order dispatch ids for serialized kernels; hsa_barrier uses them to tell when a + // transition packet handed to this queue has executed (completed >= dispatched-at-enqueue). + void serialized_dispatched_inc() { _serialized_dispatched.fetch_add(1, std::memory_order_release); } + void serialized_completed_inc() { _serialized_completed.fetch_add(1, std::memory_order_release); } + uint64_t serialized_dispatched() const + { + return _serialized_dispatched.load(std::memory_order_acquire); + } + uint64_t serialized_completed() const + { + return _serialized_completed.load(std::memory_order_acquire); + } void sync() const; void register_callback(ClientID id, queue_callbacks_t callbacks); @@ -188,6 +201,8 @@ class Queue private: std::atomic _notifiers = {0}; std::atomic _active_async_packets = {0}; + std::atomic _serialized_dispatched = {0}; + std::atomic _serialized_completed = {0}; CoreApiTable _core_api = {}; AmdExtTable _ext_api = {}; const AgentCache& _agent; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp index 619f8526aad3..d2505216d937 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp @@ -388,3 +388,147 @@ TEST(hsa_barrier, block_multi) registration::set_init_status(1); registration::finalize(); } + +// Regression tests: a barrier must not be retired (its signal destroyed/recyclable) while a queue +// still holds an un-executed transition packet referencing it. + +// complete-but-referenced barrier is not safe_to_retire() until the carrying dispatch runs. +TEST(hsa_barrier, safe_to_retire_gated_on_transition_drain) +{ + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); + test_init(); + + registration::init_logging(); + registration::set_init_status(-1); + context::push_client(1); + + auto queues = create_queue_map(1); + ASSERT_EQ(queues.size(), 1u); + Queue* q = queues.begin()->second.get(); + + hsa_barrier::queue_map_ptr_t q_map; + for(const auto& [k, v] : queues) + q_map[k] = v.get(); + + q->serialized_dispatched_inc(); // pre-barrier dispatch (id 1), in flight + q->async_started(); + + hsa_barrier barrier([]() {}, get_api_table()); + barrier.set_barrier(q_map); + ASSERT_FALSE(barrier.complete()); + ASSERT_FALSE(barrier.safe_to_retire()); + + q->serialized_dispatched_inc(); // carrying dispatch (id 2) gets the transition packet + auto pkt = barrier.enqueue_packet(q); + ASSERT_TRUE(pkt.has_value()); + + // dispatch 1 completes -> barrier clears, but carrier (id 2) has not run -> not retirable + q->serialized_completed_inc(); // completion of dispatch id 1 + q->async_complete(); + barrier.register_completion(q); + barrier.notify_drain(q); // completed(1) < token(2): no-op + ASSERT_TRUE(barrier.complete()); + ASSERT_FALSE(barrier.safe_to_retire()); + + // carrier (id 2) completes -> retirable + q->serialized_completed_inc(); // completion of dispatch id 2 (the carrier) + barrier.notify_drain(q); + ASSERT_TRUE(barrier.safe_to_retire()); + + registration::set_init_status(1); + registration::finalize(); +} + +// teardown escape: removing an enqueued queue releases the barrier instead of pinning it forever. +TEST(hsa_barrier, safe_to_retire_after_enqueued_queue_removed) +{ + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); + test_init(); + + registration::init_logging(); + registration::set_init_status(-1); + context::push_client(1); + + auto queues = create_queue_map(1); + ASSERT_EQ(queues.size(), 1u); + Queue* q = queues.begin()->second.get(); + + hsa_barrier::queue_map_ptr_t q_map; + for(const auto& [k, v] : queues) + q_map[k] = v.get(); + + q->async_started(); + hsa_barrier barrier([]() {}, get_api_table()); + barrier.set_barrier(q_map); + + q->serialized_dispatched_inc(); + ASSERT_TRUE(barrier.enqueue_packet(q).has_value()); + + q->async_complete(); + barrier.register_completion(q); + ASSERT_TRUE(barrier.complete()); + ASSERT_FALSE(barrier.safe_to_retire()); // transition packet still outstanding + + barrier.remove_queue(q); + ASSERT_TRUE(barrier.safe_to_retire()); + + registration::set_init_status(1); + registration::finalize(); +} + +// a complete-but-referenced barrier's signal handle must not be recycled into a new barrier. +TEST(hsa_barrier, retired_barrier_signal_not_recycled_while_transition_outstanding) +{ + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); + test_init(); + + registration::init_logging(); + registration::set_init_status(-1); + context::push_client(1); + + auto queues = create_queue_map(1); + ASSERT_EQ(queues.size(), 1u); + Queue* q = queues.begin()->second.get(); + + hsa_barrier::queue_map_ptr_t q_map; + for(const auto& [k, v] : queues) + q_map[k] = v.get(); + + // arm the barrier on one in-flight kernel + q->async_started(); + auto* b_old = new hsa_barrier([]() {}, get_api_table()); + b_old->set_barrier(q_map); + ASSERT_FALSE(b_old->complete()); + const uint64_t old_handle = b_old->signal_handle(); + + // q gets a transition packet on the barrier signal (never executed below) + q->serialized_dispatched_inc(); + auto pkt = b_old->enqueue_packet(q); + ASSERT_TRUE(pkt.has_value()); + ASSERT_EQ(pkt->barrier_and.dep_signal[0].handle, old_handle); + + // barrier clears; transition packet still outstanding + q->async_complete(); + b_old->register_completion(q); + ASSERT_TRUE(b_old->complete()); + ASSERT_FALSE(b_old->safe_to_retire()); + + // destroy while still referenced; its signal handle must not become reusable + delete b_old; + + std::vector> fresh; + bool reused = false; + for(int i = 0; i < 8; ++i) + { + fresh.emplace_back(std::make_unique([]() {}, get_api_table())); + if(fresh.back()->signal_handle() == old_handle) reused = true; + } + + EXPECT_FALSE(reused) << "retired barrier signal " << old_handle + << " was recycled while a transition packet still referenced it"; + + fresh.clear(); + + registration::set_init_status(1); + registration::finalize(); +} From 8894cf1bf6b525d0ed3e90eac05c4634df5c1b29 Mon Sep 17 00:00:00 2001 From: Venkateshwar Reddy Kandula Date: Wed, 3 Jun 2026 20:18:32 +0000 Subject: [PATCH 2/4] format and misc. --- .../hsa/profile_serializer.cpp | 2 +- .../source/lib/rocprofiler-sdk/hsa/queue.cpp | 3 ++- .../source/lib/rocprofiler-sdk/hsa/queue.hpp | 24 ++++++++++++------- 3 files changed, 19 insertions(+), 10 deletions(-) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp index 9c6564f458bb..d69da49aa90e 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp @@ -84,7 +84,7 @@ profiler_serializer::kernel_completion_signal(const Queue& completed) for(auto& barrier : _barrier) barrier.barrier->notify_drain(&completed); - // We do not want to track kernel compleiton signals before we have reached the barrier + // We do not want to track kernel completion signals before we have reached the barrier clear_complete_barriers(_barrier); // Find the state of this barrier 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 425469243ab2..7b99e1cc101b 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -202,7 +202,8 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data) if(packet.is_serialized) { - // count this completion before notifying the serializer (drives barrier drain detection) + // count this completion before notifying the serializer (drives barrier drain + // detection) queue_info_session.queue.serialized_completed_inc(); CHECK_NOTNULL(hsa::get_queue_controller()) ->serializer(&queue_info_session.queue) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp index 7dc1f2b69f83..2d3aea0ca53d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -173,15 +173,23 @@ class Queue // In-order dispatch ids for serialized kernels; hsa_barrier uses them to tell when a // transition packet handed to this queue has executed (completed >= dispatched-at-enqueue). - void serialized_dispatched_inc() { _serialized_dispatched.fetch_add(1, std::memory_order_release); } - void serialized_completed_inc() { _serialized_completed.fetch_add(1, std::memory_order_release); } + // Monotonic ids; ordering/visibility comes from the serializer lock, so relaxed is sufficient + // (a stale read only delays barrier retirement, never makes it premature). + void serialized_dispatched_inc() + { + _serialized_dispatched.fetch_add(1, std::memory_order_relaxed); + } + void serialized_completed_inc() + { + _serialized_completed.fetch_add(1, std::memory_order_relaxed); + } uint64_t serialized_dispatched() const { - return _serialized_dispatched.load(std::memory_order_acquire); + return _serialized_dispatched.load(std::memory_order_relaxed); } uint64_t serialized_completed() const { - return _serialized_completed.load(std::memory_order_acquire); + return _serialized_completed.load(std::memory_order_relaxed); } void sync() const; @@ -199,12 +207,12 @@ class Queue hsa_signal_t ready_signal = {.handle = 0}; private: - std::atomic _notifiers = {0}; - std::atomic _active_async_packets = {0}; + std::atomic _notifiers = {0}; + std::atomic _active_async_packets = {0}; std::atomic _serialized_dispatched = {0}; std::atomic _serialized_completed = {0}; - CoreApiTable _core_api = {}; - AmdExtTable _ext_api = {}; + CoreApiTable _core_api = {}; + AmdExtTable _ext_api = {}; const AgentCache& _agent; common::Synchronized _callbacks = {}; hsa_queue_t* _intercept_queue = nullptr; From 855a89ff526fe5bc825a958abb28a90a694a822b Mon Sep 17 00:00:00 2001 From: Venkateshwar Reddy Kandula Date: Thu, 4 Jun 2026 19:35:35 +0000 Subject: [PATCH 3/4] address comments. better comments and function names. --- .../lib/rocprofiler-sdk/hsa/hsa_barrier.cpp | 24 +++--- .../lib/rocprofiler-sdk/hsa/hsa_barrier.hpp | 7 +- .../hsa/profile_serializer.cpp | 17 ++-- .../lib/rocprofiler-sdk/tests/hsa_barrier.cpp | 77 +++---------------- 4 files changed, 35 insertions(+), 90 deletions(-) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp index 90a3cbdc719b..ee9b1bbfdfe2 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp @@ -49,19 +49,14 @@ hsa_barrier::~hsa_barrier() _barrier_enqueued.rlock( [&](const auto& barrier_enqueued) { outstanding = !barrier_enqueued.empty(); }); + // retirement is gated on safe_to_destroy() and queue teardown clears entries via remove_queue, + // so a barrier should never be destroyed while a transition packet still references its signal. + ROCP_ERROR_IF(outstanding) << "hsa_barrier (handle: " << _barrier_signal.handle + << ") destroyed with outstanding transition packets"; + // release the signal so any stale transition packet can pass clear_barrier(); - // Backstop: retirement is gated on safe_to_retire(), so this should not happen. If a packet - // still references this signal, leak it rather than let HSA recycle the handle into a barrier. - if(outstanding) - { - ROCP_WARNING << "hsa_barrier (handle: " << _barrier_signal.handle - << ") destroyed with outstanding transition packets; leaking its signal to " - "avoid handle reuse."; - return; - } - // Destroy the barrier signal _core_api.hsa_signal_destroy_fn(_barrier_signal); } @@ -97,7 +92,8 @@ hsa_barrier::enqueue_packet(const Queue* queue) if(barrier_enqueued.find(queue->get_id().handle) == barrier_enqueued.end()) { return_block = true; - // stamp the queue with the dispatch id that carries this transition packet + // record the current serialized-dispatch id; the transition packet has executed + // once this queue's completed count reaches it barrier_enqueued.emplace(queue->get_id().handle, queue->serialized_dispatched()); } }); @@ -113,8 +109,10 @@ hsa_barrier::enqueue_packet(const Queue* queue) } void -hsa_barrier::notify_drain(const Queue* queue) +hsa_barrier::drain_queue(const Queue* queue) { + // transition packet can't have executed until the barrier clears; skip the lock while armed + if(!complete()) return; _barrier_enqueued.wlock([&](auto& barrier_enqueued) { auto it = barrier_enqueued.find(queue->get_id().handle); if(it == barrier_enqueued.end()) return; @@ -170,7 +168,7 @@ hsa_barrier::complete() const } bool -hsa_barrier::safe_to_retire() const +hsa_barrier::safe_to_destroy() const { if(!complete()) return false; bool no_outstanding = false; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp index 28755f8f84b9..f037366d421c 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp @@ -66,21 +66,18 @@ class hsa_barrier // Drop the queue from the outstanding set once it has executed past the dispatch that // carried this barrier's transition packet (called on each serialized completion). - void notify_drain(const Queue* queue); + void drain_queue(const Queue* queue); // Checks if this barrier is complete bool complete() const; // complete() AND no queue still references the signal -> safe to destroy (no handle reuse). - bool safe_to_retire() const; + bool safe_to_destroy() const; // Removes a queue from the barrier dependency list (waiting + outstanding transition packets). // If this is the last queue waiting, clears the barrier and marks it as complete. void remove_queue(const Queue* queue); - // Handle of the underlying barrier signal. - uint64_t signal_handle() const { return _barrier_signal.handle; } - private: std::function _barrier_finished = {}; CoreApiTable _core_api = {}; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp index d69da49aa90e..2eceebc93827 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp @@ -42,14 +42,23 @@ profiler_serializer_ready_signal_handler(hsa_signal_value_t /* signal_value */, return true; } +void +drain_barriers(std::deque& barriers, + const Queue& completed) +{ + // let each barrier drop `completed` if its transition packet on that queue has now executed + for(auto& barrier : barriers) + barrier.barrier->drain_queue(&completed); +} + void clear_complete_barriers(std::deque& barriers) { - // Retire only when safe_to_retire(): complete AND no queue still references the signal. + // Retire only when safe_to_destroy(): complete AND no queue still references the signal. // Otherwise the barrier stays parked (signal at 0, stale packets pass) until its packets drain. while(!barriers.empty()) { - if(barriers.front().barrier->safe_to_retire()) + if(barriers.front().barrier->safe_to_destroy()) { barriers.pop_front(); } @@ -80,9 +89,7 @@ profiler_serializer::add_queue(hsa_queue_t** hsa_queues, const Queue& queue) void profiler_serializer::kernel_completion_signal(const Queue& completed) { - // let each barrier drop `completed` if its transition packet on that queue has now executed - for(auto& barrier : _barrier) - barrier.barrier->notify_drain(&completed); + drain_barriers(_barrier, completed); // We do not want to track kernel completion signals before we have reached the barrier clear_complete_barriers(_barrier); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp index d2505216d937..d9057257977a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp @@ -392,8 +392,8 @@ TEST(hsa_barrier, block_multi) // Regression tests: a barrier must not be retired (its signal destroyed/recyclable) while a queue // still holds an un-executed transition packet referencing it. -// complete-but-referenced barrier is not safe_to_retire() until the carrying dispatch runs. -TEST(hsa_barrier, safe_to_retire_gated_on_transition_drain) +// complete-but-referenced barrier is not safe_to_destroy() until the carrying dispatch runs. +TEST(hsa_barrier, safe_to_destroy_gated_on_transition_drain) { ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); test_init(); @@ -416,7 +416,7 @@ TEST(hsa_barrier, safe_to_retire_gated_on_transition_drain) hsa_barrier barrier([]() {}, get_api_table()); barrier.set_barrier(q_map); ASSERT_FALSE(barrier.complete()); - ASSERT_FALSE(barrier.safe_to_retire()); + ASSERT_FALSE(barrier.safe_to_destroy()); q->serialized_dispatched_inc(); // carrying dispatch (id 2) gets the transition packet auto pkt = barrier.enqueue_packet(q); @@ -426,21 +426,21 @@ TEST(hsa_barrier, safe_to_retire_gated_on_transition_drain) q->serialized_completed_inc(); // completion of dispatch id 1 q->async_complete(); barrier.register_completion(q); - barrier.notify_drain(q); // completed(1) < token(2): no-op + barrier.drain_queue(q); // completed(1) < token(2): no-op ASSERT_TRUE(barrier.complete()); - ASSERT_FALSE(barrier.safe_to_retire()); + ASSERT_FALSE(barrier.safe_to_destroy()); // carrier (id 2) completes -> retirable q->serialized_completed_inc(); // completion of dispatch id 2 (the carrier) - barrier.notify_drain(q); - ASSERT_TRUE(barrier.safe_to_retire()); + barrier.drain_queue(q); + ASSERT_TRUE(barrier.safe_to_destroy()); registration::set_init_status(1); registration::finalize(); } // teardown escape: removing an enqueued queue releases the barrier instead of pinning it forever. -TEST(hsa_barrier, safe_to_retire_after_enqueued_queue_removed) +TEST(hsa_barrier, safe_to_destroy_after_enqueued_queue_removed) { ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); test_init(); @@ -467,67 +467,10 @@ TEST(hsa_barrier, safe_to_retire_after_enqueued_queue_removed) q->async_complete(); barrier.register_completion(q); ASSERT_TRUE(barrier.complete()); - ASSERT_FALSE(barrier.safe_to_retire()); // transition packet still outstanding + ASSERT_FALSE(barrier.safe_to_destroy()); // transition packet still outstanding barrier.remove_queue(q); - ASSERT_TRUE(barrier.safe_to_retire()); - - registration::set_init_status(1); - registration::finalize(); -} - -// a complete-but-referenced barrier's signal handle must not be recycled into a new barrier. -TEST(hsa_barrier, retired_barrier_signal_not_recycled_while_transition_outstanding) -{ - ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); - test_init(); - - registration::init_logging(); - registration::set_init_status(-1); - context::push_client(1); - - auto queues = create_queue_map(1); - ASSERT_EQ(queues.size(), 1u); - Queue* q = queues.begin()->second.get(); - - hsa_barrier::queue_map_ptr_t q_map; - for(const auto& [k, v] : queues) - q_map[k] = v.get(); - - // arm the barrier on one in-flight kernel - q->async_started(); - auto* b_old = new hsa_barrier([]() {}, get_api_table()); - b_old->set_barrier(q_map); - ASSERT_FALSE(b_old->complete()); - const uint64_t old_handle = b_old->signal_handle(); - - // q gets a transition packet on the barrier signal (never executed below) - q->serialized_dispatched_inc(); - auto pkt = b_old->enqueue_packet(q); - ASSERT_TRUE(pkt.has_value()); - ASSERT_EQ(pkt->barrier_and.dep_signal[0].handle, old_handle); - - // barrier clears; transition packet still outstanding - q->async_complete(); - b_old->register_completion(q); - ASSERT_TRUE(b_old->complete()); - ASSERT_FALSE(b_old->safe_to_retire()); - - // destroy while still referenced; its signal handle must not become reusable - delete b_old; - - std::vector> fresh; - bool reused = false; - for(int i = 0; i < 8; ++i) - { - fresh.emplace_back(std::make_unique([]() {}, get_api_table())); - if(fresh.back()->signal_handle() == old_handle) reused = true; - } - - EXPECT_FALSE(reused) << "retired barrier signal " << old_handle - << " was recycled while a transition packet still referenced it"; - - fresh.clear(); + ASSERT_TRUE(barrier.safe_to_destroy()); registration::set_init_status(1); registration::finalize(); From 8d409b651da2973ee36b33aa9fb09130c69305d4 Mon Sep 17 00:00:00 2001 From: Venkateshwar Reddy Kandula Date: Fri, 5 Jun 2026 18:34:07 +0000 Subject: [PATCH 4/4] move couting serial_counts dispatched and completed to serializer. --- .../lib/rocprofiler-sdk/hsa/hsa_barrier.cpp | 22 +++++++------- .../lib/rocprofiler-sdk/hsa/hsa_barrier.hpp | 15 +++++----- .../hsa/profile_serializer.cpp | 26 ++++++++++++---- .../hsa/profile_serializer.hpp | 25 ++++++++++++---- .../source/lib/rocprofiler-sdk/hsa/queue.cpp | 5 ---- .../source/lib/rocprofiler-sdk/hsa/queue.hpp | 30 +++---------------- .../lib/rocprofiler-sdk/tests/hsa_barrier.cpp | 22 +++++++------- 7 files changed, 75 insertions(+), 70 deletions(-) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp index ee9b1bbfdfe2..d86385c74112 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp @@ -49,8 +49,8 @@ hsa_barrier::~hsa_barrier() _barrier_enqueued.rlock( [&](const auto& barrier_enqueued) { outstanding = !barrier_enqueued.empty(); }); - // retirement is gated on safe_to_destroy() and queue teardown clears entries via remove_queue, - // so a barrier should never be destroyed while a transition packet still references its signal. + // retirement is gated on safe_to_destroy(), so a barrier should never be destroyed while a + // transition packet still references its signal; log an error (and still release) if it does. ROCP_ERROR_IF(outstanding) << "hsa_barrier (handle: " << _barrier_signal.handle << ") destroyed with outstanding transition packets"; @@ -84,17 +84,17 @@ hsa_barrier::set_barrier(const queue_map_ptr_t& q) } std::optional -hsa_barrier::enqueue_packet(const Queue* queue) +hsa_barrier::enqueue_packet(int64_t queue_id, uint64_t dispatch_id) { if(complete()) return std::nullopt; bool return_block = false; _barrier_enqueued.wlock([&](auto& barrier_enqueued) { - if(barrier_enqueued.find(queue->get_id().handle) == barrier_enqueued.end()) + if(barrier_enqueued.find(queue_id) == barrier_enqueued.end()) { return_block = true; - // record the current serialized-dispatch id; the transition packet has executed - // once this queue's completed count reaches it - barrier_enqueued.emplace(queue->get_id().handle, queue->serialized_dispatched()); + // record the dispatch id carrying this transition packet; it has executed once the + // queue's completed count reaches this id + barrier_enqueued.emplace(queue_id, dispatch_id); } }); @@ -104,20 +104,20 @@ hsa_barrier::enqueue_packet(const Queue* queue) barrier.barrier_and.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; barrier.barrier_and.dep_signal[0] = _barrier_signal; ROCP_INFO << "Barrier (handle: " << _barrier_signal.handle - << ") added to queue (handle: " << queue->get_id().handle << ")"; + << ") added to queue (handle: " << queue_id << ")"; return barrier; } void -hsa_barrier::drain_queue(const Queue* queue) +hsa_barrier::drain_queue(int64_t queue_id, uint64_t completed_id) { // transition packet can't have executed until the barrier clears; skip the lock while armed if(!complete()) return; _barrier_enqueued.wlock([&](auto& barrier_enqueued) { - auto it = barrier_enqueued.find(queue->get_id().handle); + auto it = barrier_enqueued.find(queue_id); if(it == barrier_enqueued.end()) return; // packet executed once completions reach the carrying dispatch id (queue is in-order) - if(queue->serialized_completed() >= it->second) + if(completed_id >= it->second) { barrier_enqueued.erase(it); } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp index f037366d421c..6145b948cb8e 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp @@ -54,19 +54,20 @@ class hsa_barrier // packets, clears the barrier and marks it as complete. void set_barrier(const queue_map_ptr_t& q); - // Returns the packets needed to place this barrier in a queue. - // If the barrier has already been enqueued in the queue given or the barrier is complete, no - // packets are returned. - std::optional enqueue_packet(const Queue* queue); + // Returns the transition packet to insert in the queue identified by queue_id, stamping it with + // dispatch_id (that queue's in-order serialized-dispatch id carrying the packet). Returns + // nothing if the barrier is already complete or the queue was already handed a packet. + std::optional enqueue_packet(int64_t queue_id, uint64_t dispatch_id); // To be called when 1 async packet is finished in the given queue to decrement the remaining // packets counter. If all packets in all queues given in set_barrier are completed, clears the // barrier and marks it as complete. Returns true if the given queue had async packets waiting. bool register_completion(const Queue* queue); - // Drop the queue from the outstanding set once it has executed past the dispatch that - // carried this barrier's transition packet (called on each serialized completion). - void drain_queue(const Queue* queue); + // Drop the queue from the outstanding set once its completed count (completed_id) reaches the + // dispatch id that carried this barrier's transition packet (called on each serialized + // completion). + void drain_queue(int64_t queue_id, uint64_t completed_id); // Checks if this barrier is complete bool complete() const; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp index 2eceebc93827..bab971b8aa83 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp @@ -44,11 +44,12 @@ profiler_serializer_ready_signal_handler(hsa_signal_value_t /* signal_value */, void drain_barriers(std::deque& barriers, - const Queue& completed) + int64_t queue_id, + uint64_t completed_id) { - // let each barrier drop `completed` if its transition packet on that queue has now executed + // let each barrier drop the queue if its transition packet on that queue has now executed for(auto& barrier : barriers) - barrier.barrier->drain_queue(&completed); + barrier.barrier->drain_queue(queue_id, completed_id); } void @@ -84,12 +85,21 @@ profiler_serializer::add_queue(hsa_queue_t** hsa_queues, const Queue& queue) profiler_serializer_ready_signal_handler, *hsa_queues); if(status != HSA_STATUS_SUCCESS) ROCP_FATAL << "hsa_amd_signal_async_handler failed"; + + // Track this queue's in-order serialized dispatch/completion ids. Created here (under wlock) so + // the dispatch path only ever increments an existing entry under the shared lock. + _serial.try_emplace(queue.get_id().handle); } void profiler_serializer::kernel_completion_signal(const Queue& completed) { - drain_barriers(_barrier, completed); + const auto qid = completed.get_id().handle; + uint64_t completed_id = 0; + if(auto it = _serial.find(qid); it != _serial.end()) + completed_id = it->second.completed.fetch_add(1, std::memory_order_relaxed) + 1; + + drain_barriers(_barrier, qid, completed_id); // We do not want to track kernel completion signals before we have reached the barrier clear_complete_barriers(_barrier); @@ -191,9 +201,15 @@ profiler_serializer::kernel_dispatch(const Queue& queue) const return barrier; }; + // advance this queue's in-order dispatch id (counts every serialized dispatch) + const auto qid = queue.get_id().handle; + uint64_t dispatch_id = 0; + if(auto it = _serial.find(qid); it != _serial.end()) + dispatch_id = it->second.dispatched.fetch_add(1, std::memory_order_relaxed) + 1; + if(!_barrier.empty()) { - if(auto maybe_barrier = _barrier.back().barrier->enqueue_packet(&queue)) + if(auto maybe_barrier = _barrier.back().barrier->enqueue_packet(qid, dispatch_id)) { ret.push_back(*maybe_barrier); } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp index d689e928c89b..fecde1feff9a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp @@ -29,6 +29,8 @@ #include #include +#include +#include #include #include @@ -82,13 +84,26 @@ class profiler_serializer void destroy_queue(hsa_queue_t* id, const Queue& queue); - static void add_queue(hsa_queue_t** hsa_queues, const Queue& queue); + void add_queue(hsa_queue_t** hsa_queues, const Queue& queue); private: - const Queue* _dispatch_queue{nullptr}; - std::deque _dispatch_ready; - std::atomic _serializer_status{Status::DISABLED}; - std::deque _barrier; + // Per-queue in-order serialized dispatch/completion ids. hsa_barrier uses them to tell when a + // transition packet handed to a queue has executed (completed >= dispatched-at-enqueue). + // Relaxed atomics: ordering/visibility comes from the serializer lock; dispatch only takes a + // shared lock, so the values must be atomic and the map structure must stay stable. Entries are + // created in add_queue under the write lock and only ever incremented (never inserted) on the + // shared dispatch path. + struct serial_counts + { + std::atomic dispatched{0}; + std::atomic completed{0}; + }; + + const Queue* _dispatch_queue{nullptr}; + std::deque _dispatch_ready; + std::atomic _serializer_status{Status::DISABLED}; + std::deque _barrier; + mutable std::unordered_map _serial; }; } // namespace hsa 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 7b99e1cc101b..5ef8a86b47b5 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -202,9 +202,6 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data) if(packet.is_serialized) { - // count this completion before notifying the serializer (drives barrier drain - // detection) - queue_info_session.queue.serialized_completed_inc(); CHECK_NOTNULL(hsa::get_queue_controller()) ->serializer(&queue_info_session.queue) .wlock([&](auto& serializer) { @@ -634,8 +631,6 @@ WriteInterceptor(const void* packets, if(_packet_data.is_serialized) { inserted_before = true; - // count this dispatch before enqueue_packet() stamps the barrier with its id - queue.serialized_dispatched_inc(); CHECK_NOTNULL(hsa::get_queue_controller()) ->serializer(&queue) .rlock([&](const auto& serializer) { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp index 2d3aea0ca53d..e3bfc2205e25 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -171,26 +171,6 @@ class Queue return _core_api.hsa_signal_load_scacquire_fn(_active_kernels); } - // In-order dispatch ids for serialized kernels; hsa_barrier uses them to tell when a - // transition packet handed to this queue has executed (completed >= dispatched-at-enqueue). - // Monotonic ids; ordering/visibility comes from the serializer lock, so relaxed is sufficient - // (a stale read only delays barrier retirement, never makes it premature). - void serialized_dispatched_inc() - { - _serialized_dispatched.fetch_add(1, std::memory_order_relaxed); - } - void serialized_completed_inc() - { - _serialized_completed.fetch_add(1, std::memory_order_relaxed); - } - uint64_t serialized_dispatched() const - { - return _serialized_dispatched.load(std::memory_order_relaxed); - } - uint64_t serialized_completed() const - { - return _serialized_completed.load(std::memory_order_relaxed); - } void sync() const; void register_callback(ClientID id, queue_callbacks_t callbacks); @@ -207,12 +187,10 @@ class Queue hsa_signal_t ready_signal = {.handle = 0}; private: - std::atomic _notifiers = {0}; - std::atomic _active_async_packets = {0}; - std::atomic _serialized_dispatched = {0}; - std::atomic _serialized_completed = {0}; - CoreApiTable _core_api = {}; - AmdExtTable _ext_api = {}; + std::atomic _notifiers = {0}; + std::atomic _active_async_packets = {0}; + CoreApiTable _core_api = {}; + AmdExtTable _ext_api = {}; const AgentCache& _agent; common::Synchronized _callbacks = {}; hsa_queue_t* _intercept_queue = nullptr; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp index d9057257977a..2349160ef2a8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp @@ -143,7 +143,7 @@ inject_barriers(hsa_barrier& barrier, QueueController::queue_map_t& queues) for(auto& [hsa_queue, fq] : queues) { auto complete = barrier.complete(); - auto _pkt = barrier.enqueue_packet(fq.get()); + auto _pkt = barrier.enqueue_packet(fq->get_id().handle, 0); // If barrier is complete, no packets should be generated ASSERT_NE(complete, _pkt.has_value()); @@ -410,29 +410,28 @@ TEST(hsa_barrier, safe_to_destroy_gated_on_transition_drain) for(const auto& [k, v] : queues) q_map[k] = v.get(); - q->serialized_dispatched_inc(); // pre-barrier dispatch (id 1), in flight - q->async_started(); + const int64_t qid = q->get_id().handle; + + q->async_started(); // pre-barrier dispatch (id 1) in flight hsa_barrier barrier([]() {}, get_api_table()); barrier.set_barrier(q_map); ASSERT_FALSE(barrier.complete()); ASSERT_FALSE(barrier.safe_to_destroy()); - q->serialized_dispatched_inc(); // carrying dispatch (id 2) gets the transition packet - auto pkt = barrier.enqueue_packet(q); + // carrying dispatch (id 2) gets the transition packet + auto pkt = barrier.enqueue_packet(qid, 2); ASSERT_TRUE(pkt.has_value()); // dispatch 1 completes -> barrier clears, but carrier (id 2) has not run -> not retirable - q->serialized_completed_inc(); // completion of dispatch id 1 q->async_complete(); barrier.register_completion(q); - barrier.drain_queue(q); // completed(1) < token(2): no-op + barrier.drain_queue(qid, 1); // completed(1) < token(2): no-op ASSERT_TRUE(barrier.complete()); ASSERT_FALSE(barrier.safe_to_destroy()); // carrier (id 2) completes -> retirable - q->serialized_completed_inc(); // completion of dispatch id 2 (the carrier) - barrier.drain_queue(q); + barrier.drain_queue(qid, 2); ASSERT_TRUE(barrier.safe_to_destroy()); registration::set_init_status(1); @@ -457,12 +456,13 @@ TEST(hsa_barrier, safe_to_destroy_after_enqueued_queue_removed) for(const auto& [k, v] : queues) q_map[k] = v.get(); + const int64_t qid = q->get_id().handle; + q->async_started(); hsa_barrier barrier([]() {}, get_api_table()); barrier.set_barrier(q_map); - q->serialized_dispatched_inc(); - ASSERT_TRUE(barrier.enqueue_packet(q).has_value()); + ASSERT_TRUE(barrier.enqueue_packet(qid, 1).has_value()); q->async_complete(); barrier.register_completion(q);