diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp index f9f46e66220e..b8ae91ec2833 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp @@ -23,6 +23,7 @@ #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/common/container/small_vector.hpp" #include "lib/common/container/stable_vector.hpp" +#include "lib/common/scope_destructor.hpp" #include "lib/common/static_object.hpp" #include "lib/common/synchronized.hpp" #include "lib/common/utility.hpp" @@ -38,23 +39,38 @@ #include #include +#include #include #include #include #include #include -#include #include +#include +#include +#include namespace rocprofiler { namespace context { +// Owns the contexts it names. A reader holding a snapshot keeps them alive for the duration of +// the read, so deregister_client_contexts() dropping the registry's reference cannot destroy a +// context out from under an in-flight completion. The parallel raw-pointer array is what the +// snapshot iterates, so walking one costs no reference counting. +struct registered_context_storage +{ + std::vector> owned = {}; + std::vector view = {}; +}; + namespace { -using reserve_size_t = common::container::reserve_size; -using stable_context_vec_t = common::container::stable_vector, 8>; -using active_context_vec_t = common::container::stable_vector, 8>; +using reserve_size_t = common::container::reserve_size; +using context_ptr_t = std::shared_ptr; +using stable_context_vec_t = common::container::stable_vector; +using active_context_vec_t = common::container::stable_vector, 8>; +using context_snapshot_ptr_t = std::shared_ptr; constexpr auto invalid_client_idx = std::numeric_limits::max(); @@ -77,14 +93,6 @@ get_client_index() return _v; } -stable_context_vec_t*& -get_registered_contexts_impl() -{ - static auto*& _v = common::static_object::construct( - reserve_size_t{stable_context_vec_t::chunk_size}); - return _v; -} - auto& get_num_active_contexts() { @@ -98,8 +106,150 @@ get_active_contexts_impl() static auto* _v = new active_context_vec_t{reserve_size_t{active_context_vec_t::chunk_size}}; return *_v; } + +// C++20 replaces these with std::atomic>; the free functions are the C++17 +// spelling and are deprecated there, so select on the feature macro rather than leaving a +// deprecation warning to trip -Werror on a future standard bump. +#if defined(__cpp_lib_atomic_shared_ptr) && __cpp_lib_atomic_shared_ptr >= 201711L +using published_snapshot_t = std::atomic; + +context_snapshot_ptr_t +load_snapshot(published_snapshot_t& _v) +{ + return _v.load(std::memory_order_acquire); +} + +void +store_snapshot(published_snapshot_t& _v, context_snapshot_ptr_t&& _new) +{ + _v.store(std::move(_new), std::memory_order_release); +} +#else +using published_snapshot_t = context_snapshot_ptr_t; + +context_snapshot_ptr_t +load_snapshot(published_snapshot_t& _v) +{ + return std::atomic_load_explicit(&_v, std::memory_order_acquire); +} + +void +store_snapshot(published_snapshot_t& _v, context_snapshot_ptr_t&& _new) +{ + std::atomic_store_explicit(&_v, std::move(_new), std::memory_order_release); +} +#endif + +// The registry and the snapshot published from it live in one static object rather than two. +// common::destroy_static_objects() tears static objects down in reverse construction order, so two +// objects would be ordered by whichever happened to be touched first -- and if the snapshot +// outlived the registry it would be holding the last reference to every context, destroying them +// after the HSA tables that a context destructor calls into are already gone. As members, the +// snapshot is destroyed before the registry by the ordinary reverse-member rule. +struct registry_state +{ + stable_context_vec_t registry = + stable_context_vec_t{reserve_size_t{stable_context_vec_t::chunk_size}}; + published_snapshot_t published = {}; +}; + +registry_state* +get_registry_state() +{ + static auto*& _v = common::static_object::construct(); + return _v; +} + +stable_context_vec_t* +get_registered_contexts_impl() +{ + auto* _state = get_registry_state(); + return (_state) ? &_state->registry : nullptr; +} + +published_snapshot_t* +get_published_snapshot() +{ + auto* _state = get_registry_state(); + return (_state) ? &_state->published : nullptr; +} + +context_snapshot_ptr_t +current_registered_contexts() +{ + auto* _pub = get_published_snapshot(); + if(!_pub) return {}; + return load_snapshot(*_pub); +} + +// Rebuilds the published snapshot from the registry. Must be called with get_contexts_mutex() +// held, i.e. by whichever function just mutated the registry. +void +publish_registered_contexts() +{ + auto* _pub = get_published_snapshot(); + if(!_pub) return; + + auto _data = std::make_shared(); + + if(auto* _impl = get_registered_contexts_impl()) + { + _data->owned.reserve(_impl->size()); + _data->view.reserve(_impl->size()); + for(const auto& itr : *_impl) + { + if(!itr) continue; + _data->owned.emplace_back(itr); + _data->view.emplace_back(itr.get()); + } + } + + store_snapshot(*_pub, std::move(_data)); +} + +// The registry is index-addressed: allocate_context() derives context_idx from the slot position, +// so a context id maps straight back to its slot. +context_ptr_t +get_registered_context_slot(uint64_t handle) +{ + auto* _impl = get_registered_contexts_impl(); + if(!_impl || handle < get_contexts_offset()) return {}; + + auto _idx = handle - get_contexts_offset(); + if(_idx >= _impl->size()) return {}; + + return _impl->at(_idx); +} + +// Context ids whose stop is in progress. stop_context() releases get_contexts_mutex() across the +// GPU drain, so for that window the context is still in the active array with no lock held. This +// set is what keeps start_context() and a second stop_context() from acting on that half-stopped +// state. Guarded by get_contexts_mutex(). +std::unordered_set& +get_stopping_contexts() +{ + static auto _v = std::unordered_set{}; + return _v; +} + +std::condition_variable& +get_contexts_cv() +{ + static auto _v = std::condition_variable{}; + return _v; +} + } // namespace +// Waits until no context is mid-stop, which restores the property the old single-lock +// stop_context() had: no other lifecycle operation ever observes a context whose services have +// been torn down but whose active slot is still populated. +void +wait_for_stopping_contexts(std::unique_lock& _lk) +{ + get_contexts_cv().wait(_lk, []() { return get_stopping_contexts().empty(); }); +} + std::mutex& get_contexts_mutex() { @@ -121,22 +271,33 @@ dispatch_counter_collection_service::intersects( return false; } +registered_contexts_snapshot::registered_contexts_snapshot(context_snapshot_ptr_t&& data) +: m_data{std::move(data)} +{ + if(m_data) + { + m_begin = m_data->view.data(); + m_end = m_begin + m_data->view.size(); + } +} + +registered_contexts_snapshot +get_registered_contexts_snapshot() +{ + return registered_contexts_snapshot{current_registered_contexts()}; +} + context_array_t& get_registered_contexts(context_array_t& data, context_filter_t filter) { data.clear(); - if(!get_registered_contexts_impl()) return data; - - auto num_ctx = get_registered_contexts_impl()->size(); - if(num_ctx <= 0) return data; + auto snapshot = get_registered_contexts_snapshot(); + if(snapshot.empty()) return data; - data.reserve(num_ctx); - for(auto& itr : *get_registered_contexts_impl()) + data.reserve(snapshot.size()); + for(const auto* ctx : snapshot) { - if(!itr.has_value()) continue; - - const auto* ctx = &itr.value(); if(!filter || (filter && filter(ctx))) data.emplace_back(ctx); } return data; @@ -232,13 +393,9 @@ allocate_context() // initial context identifier number auto _idx = get_registered_contexts_impl()->size() + get_contexts_offset(); - // make space in registered - get_registered_contexts_impl()->emplace_back(); - // create an entry in the registered - auto& _cfg_v = get_registered_contexts_impl()->back(); - _cfg_v.emplace(); - auto* _cfg = &_cfg_v.value(); + auto& _cfg_v = get_registered_contexts_impl()->emplace_back(std::make_shared()); + auto* _cfg = _cfg_v.get(); // ... if(!_cfg) return std::nullopt; @@ -251,20 +408,15 @@ allocate_context() << " rocprofiler internal error: a context was allocated without an associated tool client " "identifier"; + publish_registered_contexts(); + return rocprofiler_context_id_t{_idx}; } context* get_mutable_registered_context(rocprofiler_context_id_t id) { - if(id.handle < get_contexts_offset()) return nullptr; - if(!get_registered_contexts_impl()) return nullptr; - auto _idx = id.handle - get_contexts_offset(); - if(_idx >= get_registered_contexts_impl()->size()) - return nullptr; - else if(get_registered_contexts_impl()->at(_idx).has_value()) - return &get_registered_contexts_impl()->at(_idx).value(); - return nullptr; + return get_registered_context_slot(id.handle).get(); } const context* @@ -284,43 +436,52 @@ start_context(rocprofiler_context_id_t context_id) { if(context_id.handle < get_contexts_offset()) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; - if(!get_registered_contexts_impl()) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; - - if((context_id.handle - get_contexts_offset()) >= get_registered_contexts_impl()->size()) - return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; - - const auto* cfg = get_registered_context(context_id); + // Held for the rest of the function: the service starts below run without the contexts mutex, + // and deregister_client_contexts() may drop the registry's reference while they do. + auto _owner = context_ptr_t{}; + const context* cfg = nullptr; + { + auto _lk = std::unique_lock{get_contexts_mutex()}; + _owner = get_registered_context_slot(context_id.handle); + cfg = _owner.get(); + } if(!cfg) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; if(validate_context(cfg) != ROCPROFILER_STATUS_SUCCESS) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; - auto current_contexts = context_array_t{}; - for(const auto* itr : get_active_contexts(current_contexts)) - { - if(cfg->context_idx == itr->context_idx) - { - return ROCPROFILER_STATUS_SUCCESS; - } - else if(cfg->dispatch_counter_collection && itr->dispatch_counter_collection && - cfg->dispatch_counter_collection->intersects(*itr->dispatch_counter_collection)) - { - // Conflicting context. Two counter-collection contexts can run concurrently as - // long as they target disjoint sets of GPU agents -- the hardware counters they - // program are per-agent, so contexts that never touch the same agent cannot - // contend. A context with no agent restriction claims every agent and therefore - // still conflicts with any other counter-collection context. - return ROCPROFILER_STATUS_ERROR_CONTEXT_CONFLICT; - } - } - uint64_t rocp_tot_contexts = get_registered_contexts_impl()->size(); auto idx = rocp_tot_contexts; auto& active_contexts = get_active_contexts_impl(); { // hold a lock here to prevent multiple threads from finding the same nullptr slot auto _lk = std::unique_lock{get_contexts_mutex()}; + + // A context that is mid-stop is still in the active array while its GPU drain runs, so + // scanning against that state would either report a conflict against a context that is on + // its way out, or hand back a slot that the stop is about to clear. + wait_for_stopping_contexts(_lk); + + auto current_contexts = context_array_t{}; + for(const auto* itr : get_active_contexts(current_contexts)) + { + if(cfg->context_idx == itr->context_idx) + { + return ROCPROFILER_STATUS_SUCCESS; + } + else if(cfg->dispatch_counter_collection && itr->dispatch_counter_collection && + cfg->dispatch_counter_collection->intersects(*itr->dispatch_counter_collection)) + { + // Conflicting context. Two counter-collection contexts can run concurrently as + // long as they target disjoint sets of GPU agents -- the hardware counters they + // program are per-agent, so contexts that never touch the same agent cannot + // contend. A context with no agent restriction claims every agent and therefore + // still conflicts with any other counter-collection context. + return ROCPROFILER_STATUS_ERROR_CONTEXT_CONFLICT; + } + } + // try to find a nullptr slot first for(size_t i = 0; i < active_contexts.size(); ++i) { @@ -335,6 +496,7 @@ start_context(rocprofiler_context_id_t context_id) return ROCPROFILER_STATUS_SUCCESS; } } + // if no nullptr slot was found, then create one while lock is held if(idx == rocp_tot_contexts) { @@ -349,8 +511,7 @@ start_context(rocprofiler_context_id_t context_id) // atomic swap the pointer into the "active" array used internally const context* _expected = nullptr; - bool success = active_contexts.at(idx).compare_exchange_strong( - _expected, get_registered_context(context_id)); + bool success = active_contexts.at(idx).compare_exchange_strong(_expected, cfg); if(!success) { @@ -384,72 +545,117 @@ start_context(rocprofiler_context_id_t context_id) rocprofiler_status_t stop_context(rocprofiler_context_id_t idx) { - // hold a lock here to prevent other thread from changing the active contexts array - auto _lk = std::unique_lock{get_contexts_mutex()}; - - // atomically assign the context pointer to NULL so that it is skipped in future - // callbacks - for(auto& itr : get_active_contexts_impl()) + std::atomic* slot = nullptr; + const context* _expected = nullptr; + // Held for the rest of the function. The teardown below runs without the contexts mutex, so + // without this the registry's reference is the only thing keeping the context alive and + // deregister_client_contexts() could destroy it out from under the drain. + auto _owner = context_ptr_t{}; + + // Phase one, locked: claim the context. Only the search of the active array needs the lock; + // the teardown below must not hold it. { - const context* _expected = itr.load(std::memory_order_acquire); - if(_expected && _expected->context_idx == idx.handle) + auto _lk = std::unique_lock{get_contexts_mutex()}; + + // another thread may already be tearing this context down + wait_for_stopping_contexts(_lk); + + for(auto& itr : get_active_contexts_impl()) { - // Stop queue-interposed services before clearing the active slot so - // disable_serialization() always runs before dispatches stop being instrumented. - // Clearing the slot first opens the opposite window, in which - // kernel_dispatch_phase_enter_hook sees no active context and a dispatch is submitted - // without serializer packets while the serializer is still enabled. - if(_expected->dispatch_counter_collection) + const context* _ctx = itr.load(std::memory_order_acquire); + if(_ctx && _ctx->context_idx == idx.handle) { - rocprofiler::counters::stop_context(const_cast(_expected)); + slot = &itr; + _expected = _ctx; + break; } + } - if(_expected->dispatch_spm) - rocprofiler::spm::stop_context(const_cast(_expected)); - - if(_expected->device_thread_trace) _expected->device_thread_trace->stop_context(); - if(_expected->dispatch_thread_trace) _expected->dispatch_thread_trace->stop_context(); + if(!slot) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; - bool success = itr.compare_exchange_strong(_expected, nullptr); + _owner = get_registered_context_slot(idx.handle); + get_stopping_contexts().emplace(idx.handle); + } - if(success) - { - auto nactive = get_num_active_contexts().load(std::memory_order_acquire); - if(nactive > 0) get_num_active_contexts().fetch_sub(1, std::memory_order_release); + auto _unclaim = common::scope_destructor{[&idx]() { + { + auto _lk = std::unique_lock{get_contexts_mutex()}; + get_stopping_contexts().erase(idx.handle); + } + get_contexts_cv().notify_all(); + }}; + + // Phase two, unlocked: the service teardowns below call hsa::queue_controller_sync(), an + // unbounded wait on in-flight GPU work. Holding get_contexts_mutex() across it stalls every + // context lifecycle operation in the process behind one context's dispatches, and it puts the + // mutex on the far side of a wait that the completion path has to get through -- so any future + // completion-path read that took the mutex would deadlock rather than merely block. + // + // The active slot stays populated for the whole phase, which is deliberate: + // kernel_dispatch_phase_enter_hook has to keep seeing the context so the serialized -> + // unserialized transition stays coordinated until disable_serialization() has run. Clearing + // the slot first opens the opposite window, in which the enter hook sees no active context and + // a dispatch is submitted without serializer packets while the serializer is still enabled. + // get_stopping_contexts() is what keeps other lifecycle callers from reading that state as a + // running context. + if(_expected->dispatch_counter_collection) + { + rocprofiler::counters::stop_context(const_cast(_expected)); + } - rocprofiler::hsa::queue_interposition:: - notify_queue_interposition_consumer_context_stopped(_expected); + if(_expected->dispatch_spm) rocprofiler::spm::stop_context(const_cast(_expected)); - if(_expected->device_counter_collection) - { - rocprofiler::counters::stop_agent_ctx(const_cast(_expected)); - } + if(_expected->device_thread_trace) _expected->device_thread_trace->stop_context(); + if(_expected->dispatch_thread_trace) _expected->dispatch_thread_trace->stop_context(); -#if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0 - if(_expected->pc_sampler) - { - rocprofiler::pc_sampling::stop_service(_expected); - } -#endif + // Phase three, relocked: retire the slot. The element address is stable across phase two -- + // stable_vector never moves an element once constructed -- so slot is still the entry the + // search found. + { + auto _lk = std::unique_lock{get_contexts_mutex()}; - return ROCPROFILER_STATUS_SUCCESS; - } + if(!slot->compare_exchange_strong(_expected, nullptr)) + { + // Not reachable as written: deactivate_client_contexts() is the only other writer of + // this slot and the stopping marker holds it off until phase four returns. Kept + // because the pre-phasing code bailed out the same way on a lost exchange. + return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; } + + auto nactive = get_num_active_contexts().load(std::memory_order_acquire); + if(nactive > 0) get_num_active_contexts().fetch_sub(1, std::memory_order_release); } - return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; // compare exchange failed + // Phase four, unlocked: the remaining services are not queue-interposed, so they only need + // the slot to be cleared first, not the lock to be held. They wait on the GPU too -- + // stop_agent_ctx() waits for its stop packet and pc_sampling::stop_service() drains the + // device buffers -- and the stopping marker above, not the mutex, is what keeps a concurrent + // start_context() from running before they finish. + rocprofiler::hsa::queue_interposition::notify_queue_interposition_consumer_context_stopped( + _expected); + + if(_expected->device_counter_collection) + { + rocprofiler::counters::stop_agent_ctx(const_cast(_expected)); + } + +#if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0 + if(_expected->pc_sampler) + { + rocprofiler::pc_sampling::stop_service(_expected); + } +#endif + + return ROCPROFILER_STATUS_SUCCESS; } context_id_array_t get_client_contexts(rocprofiler_client_id_t id) { auto _data = context_id_array_t{}; - if(!get_registered_contexts_impl()) return _data; - for(auto& itr : *get_registered_contexts_impl()) + for(const auto* itr : get_registered_contexts_snapshot()) { - if(!itr.has_value()) continue; - if(itr->client_idx == id.handle) { _data.emplace_back(rocprofiler_context_id_t{.handle = itr->context_idx}); @@ -464,10 +670,12 @@ stop_client_contexts(rocprofiler_client_id_t client_id) if(!get_registered_contexts_impl()) return ROCPROFILER_STATUS_ERROR; auto ret = ROCPROFILER_STATUS_SUCCESS; - for(auto& itr : *get_registered_contexts_impl()) - { - if(!itr.has_value()) continue; + // Held across the loop: stop_context() may block on a GPU drain, and a snapshot is what keeps + // the contexts it is iterating alive for that long. + auto snapshot = get_registered_contexts_snapshot(); + for(const auto* itr : snapshot) + { if(itr->client_idx == client_id.handle) { auto status = stop_context(rocprofiler_context_id_t{itr->context_idx}); @@ -483,6 +691,11 @@ void deactivate_client_contexts(rocprofiler_client_id_t client_id) { auto _lk = std::unique_lock{get_contexts_mutex()}; + + // a context mid-stop is still in the active array; let its teardown finish rather than + // clearing the slot underneath it + wait_for_stopping_contexts(_lk); + for(auto& itr : get_active_contexts_impl()) { const context* itr_v = itr.load(std::memory_order_acquire); @@ -500,18 +713,38 @@ deactivate_client_contexts(rocprofiler_client_id_t client_id) void deregister_client_contexts(rocprofiler_client_id_t client_id) { - for(auto& itr : *get_registered_contexts_impl()) + if(!get_registered_contexts_impl()) return; + + // Outlives the lock below so the contexts are destroyed without it held: a context destructor + // calls back into HSA, which is not something to do underneath the registry mutex. + auto _retired = std::vector{}; + { - if(!itr.has_value()) continue; + // Mutates the registry, so it needs the same lock allocate_context() takes. + auto _lk = std::unique_lock{get_contexts_mutex()}; + + // a context mid-stop is still being torn down through a raw pointer; retiring its registry + // entry now would pull the object out from under that teardown + wait_for_stopping_contexts(_lk); - if(itr->client_idx == client_id.handle && buffer::get_buffers()) + for(auto& itr : *get_registered_contexts_impl()) { - for(auto& bitr : *buffer::get_buffers()) + if(!itr) continue; + + if(itr->client_idx == client_id.handle && buffer::get_buffers()) { - if(bitr && bitr->context_id == itr->context_idx) bitr.reset(); + for(auto& bitr : *buffer::get_buffers()) + { + if(bitr && bitr->context_id == itr->context_idx) bitr.reset(); + } + // Moved rather than reset: dropping the registry's reference here would destroy + // the context under the lock, and a reader holding an older snapshot still needs + // it to stay alive until that snapshot is released. + _retired.emplace_back(std::move(itr)); } - itr.reset(); } + + publish_registered_contexts(); } } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp index 8d52e2f8108d..c1c4c376c2ad 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp @@ -217,6 +217,40 @@ stop_context(rocprofiler_context_id_t id); using context_array_t = common::container::small_vector; using context_id_array_t = common::container::small_vector; +// defined in context.cpp +struct registered_context_storage; + +/// \brief An immutable view of the registered contexts, taken at the point of construction. +/// +/// The registry itself may only be walked with get_contexts_mutex() held: allocate_context() +/// appends to it while dispatch and completion paths read it. The completion path cannot take +/// that mutex, because stop_context() waits on the GPU drain and the drain waits on completions, +/// so it reads a snapshot instead. A snapshot also owns the contexts it names, which is what makes +/// it safe against deregister_client_contexts() destroying a context mid-iteration. +class registered_contexts_snapshot +{ +public: + registered_contexts_snapshot() = default; + + const context* const* begin() const { return m_begin; } + const context* const* end() const { return m_end; } + size_t size() const { return m_end - m_begin; } + bool empty() const { return m_begin == m_end; } + +private: + friend registered_contexts_snapshot get_registered_contexts_snapshot(); + + explicit registered_contexts_snapshot(std::shared_ptr&&); + + std::shared_ptr m_data = {}; + const context* const* m_begin = nullptr; + const context* const* m_end = nullptr; +}; + +/// \brief Returns the currently published snapshot of the registered contexts. +registered_contexts_snapshot +get_registered_contexts_snapshot(); + context* get_mutable_registered_context(rocprofiler_context_id_t id); @@ -228,6 +262,10 @@ using context_filter_t = bool (*)(const context*); inline bool default_context_filter(const context* val); +// These read the published snapshot and then drop it, so the pointers they return outlive the +// guarantee that the contexts are alive. That is fine for callers that cannot overlap with +// deregister_client_contexts(); anything on a dispatch or completion path should hold a +// registered_contexts_snapshot for as long as it uses the pointers. context_array_t& get_registered_contexts(context_array_t& data, context_filter_t filter = default_context_filter); @@ -268,6 +306,15 @@ deregister_client_contexts(rocprofiler_client_id_t id); std::mutex& get_contexts_mutex(); +// Blocks until no context stop is in progress. stop_context() releases get_contexts_mutex() +// across the GPU drain, so holding the mutex alone is not enough to exclude a stop: for that +// window the context is still in the active array with nothing held. Anything that reads +// activation state to decide whether a context is running has to wait here first. +// +// _lk must own get_contexts_mutex(); it is released while waiting and re-acquired on return. +void +wait_for_stopping_contexts(std::unique_lock& _lk); + inline bool default_context_filter(const context* val) { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp index 848ba3eb1be0..576752a2ab46 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp @@ -226,6 +226,12 @@ set_dispatch_agents(rocprofiler_context_id_t context_id, // set for serialization) and then race with the unordered-set assignment below while // dispatch hooks read it. auto _lk = std::unique_lock{context::get_contexts_mutex()}; + + // A context that is mid-stop is still in the active array while its GPU drain runs, and that + // drain happens with the mutex released, so the lock alone does not exclude it. Without this + // the scan below would read a context that is on its way out as a running one. + context::wait_for_stopping_contexts(_lk); + for(const auto* itr : context::get_active_contexts()) { if(itr && itr->context_idx == ctx_p->context_idx) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/device_counting.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/device_counting.cpp index 22ed73bbaae0..7c8755c1a23c 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/device_counting.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/device_counting.cpp @@ -698,7 +698,17 @@ device_counting_service_hsa_registration() agent_callback_data::~agent_callback_data() { - if(completion.handle != 0) hsa::get_core_table()->hsa_signal_destroy_fn(completion); + if(completion.handle == 0) return; + + // Guarded the same way as hsa::active_signals::destroy(). The owning context is not + // guaranteed to be destroyed while HSA is still up -- teardown can reach here after the + // runtime has been shut down, and the table pointer is also null under unit test -- so + // reaching through the table unconditionally faults instead of leaking a signal. + const auto* _core_table = hsa::get_core_table(); + if(hsa::get_hsa_ref_count() > 0 && _core_table && _core_table->hsa_signal_destroy_fn) + { + _core_table->hsa_signal_destroy_fn(completion); + } } } // namespace counters } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/queue_hooks.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/queue_hooks.cpp index a40fea7f31bf..e7464db13e3d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/queue_hooks.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/queue_hooks.cpp @@ -84,7 +84,7 @@ kernel_dispatch_phase_enter_hook( } void -kernel_dispatch_phase_exit_hook(const hsa::Queue& /*queue*/, +kernel_dispatch_phase_exit_hook(const hsa::Queue* /*queue*/, const hsa::rocprofiler_packet& /*kernel_packet*/, std::shared_ptr& session, hsa::packet_data_t& packet, @@ -107,9 +107,17 @@ kernel_dispatch_phase_exit_hook(const hsa::Queue& /*queue*/, // context from the active list. This is what guarantees a completion that arrives is delivered; // the drain in stop_context is what bounds when completions arrive. Kernel replay needs the // same property, since each pass completes separately. - auto contexts = context::get_registered_contexts(counter_contexts_filter()); + // + // Read through a snapshot rather than the registry itself. This runs on the HSA async signal + // handler, so it cannot take get_contexts_mutex() -- stop_context() waits on the drain that + // this completion is part of -- and the registry is not safe to walk unlocked while + // allocate_context() appends to it. Holding the snapshot also keeps every context it names + // alive for the duration of the loop. + auto contexts = context::get_registered_contexts_snapshot(); for(const auto* ctx : contexts) { + if(!counter_contexts_filter()(ctx)) continue; + for(auto& cb : ctx->dispatch_counter_collection->callbacks) { completed_cb(ctx, cb, session, packet, inst_pkt, dispatch_time); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/queue_hooks.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/queue_hooks.hpp index 6fc716d59258..1a76e7d66b3e 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/queue_hooks.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/queue_hooks.hpp @@ -55,8 +55,10 @@ kernel_dispatch_phase_enter_hook( // registered dispatch_counter_collection contexts (not only active ones) and calls // each callback's completed_cb; completed_cb self-filters via packet_return_map so // in-flight dispatches still complete after stop_context. +// queue is nullable: the hook never dereferences it, and tests exercise the +// early-return path without an HSA runtime to build a queue from. void -kernel_dispatch_phase_exit_hook(const hsa::Queue& queue, +kernel_dispatch_phase_exit_hook(const hsa::Queue* queue, const hsa::rocprofiler_packet& kernel_packet, std::shared_ptr& session, hsa::packet_data_t& packet, diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/queue_hooks_test.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/queue_hooks_test.cpp index 165e9e80b996..9404ac7ef9e4 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/queue_hooks_test.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/queue_hooks_test.cpp @@ -176,25 +176,19 @@ TEST(counters_queue_hooks, is_any_active_false_when_no_context_active) TEST(counters_queue_hooks, exit_hook_skips_when_inst_pkt_has_no_counter_client_id) { - ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); - test_init(); - - auto agents = hsa::get_queue_controller()->get_supported_agents(); - ASSERT_FALSE(agents.empty()); - const auto& [_, agent] = *agents.begin(); - hsa::QueueHooksFakeQueue fq(agent, {.handle = 1}); - hsa::inst_pkt_t inst_pkt; inst_pkt.emplace_back(std::make_pair(std::make_unique(), rocprofiler::hsa::queue_hooks::THREAD_TRACE_CLIENT_ID)); - auto sess = std::make_shared(hsa::queue_info_session_t{.queue = fq}); + // The hook returns before it dereferences the session, so a null one keeps this test free of + // any HSA runtime: queue_info_session_t holds a Queue& and cannot be default constructed. + auto sess = std::shared_ptr{}; auto packet = rocprofiler::hsa::packet_data_t{}; auto fq_pkt = rocprofiler::hsa::rocprofiler_packet{}; // Must return without touching registered counter contexts. rocprofiler::counters::kernel_dispatch_phase_exit_hook( - fq, fq_pkt, sess, packet, inst_pkt, rocprofiler::kernel_dispatch::profiling_time{}); + nullptr, fq_pkt, sess, packet, inst_pkt, rocprofiler::kernel_dispatch::profiling_time{}); SUCCEED(); } @@ -353,7 +347,7 @@ TEST(counters_queue_hooks, stop_context_in_flight_completion_routes_via_hook_pat inst_pkt.emplace_back(std::move(in_flight[i]), hsa::queue_hooks::COUNTERS_CLIENT_ID); rocprofiler::counters::kernel_dispatch_phase_exit_hook( - fq, pkt, sess, packet_data, inst_pkt, rocprofiler::kernel_dispatch::profiling_time{}); + &fq, pkt, sess, packet_data, inst_pkt, rocprofiler::kernel_dispatch::profiling_time{}); size_t remaining = num_dispatches; cb_info->packet_return_map.rlock([&](const auto& data) { remaining = data.size(); }); 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 de36f78bd07d..6cce070b81b8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -297,7 +297,7 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data) // Counter collection completion is migrated off the callback registry (see // WriteInterceptor); invoke it explicitly here. - counters::kernel_dispatch_phase_exit_hook(queue_info_session.queue, + counters::kernel_dispatch_phase_exit_hook(&queue_info_session.queue, packet.kernel_packet, _session, packet, diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp index ea45c17d2530..27625b0532ce 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp @@ -769,15 +769,19 @@ QueueController::update_serialization(const agent_handle_set_t& agents, bool ena { auto itr = was_enabled.find(agent_id); if(itr == was_enabled.end()) continue; - if(itr->second == state.enabled(agent_id)) continue; + const bool now_enabled = state.enabled(agent_id); + if(itr->second == now_enabled) continue; auto queues = hsa_barrier::queue_map_ptr_t{}; if(auto it = pd_map.find(agent_id); it != pd_map.end()) queues = it->second; auto ser_itr = serializers.find(agent_id); if(ser_itr == serializers.end() || !ser_itr->second) continue; + // Derived from the post-update refcount rather than the caller's `enable` + // flag, so the applied state cannot disagree with the count that produced + // the transition. ser_itr->second->wlock([&](auto& serializer) { - if(enable) + if(now_enabled) serializer.enable(queues); else serializer.disable(queues); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_hooks/client_ids.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_hooks/client_ids.hpp index 815ae5879ee3..b7a45802fef9 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_hooks/client_ids.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_hooks/client_ids.hpp @@ -30,30 +30,35 @@ namespace hsa { namespace queue_hooks { -// Tags identifying the producer subsystem of each inst_pkt_t entry. -// Only distinctness matters; values are fixed for test-order stability. -// Services are migrated off the per-queue callback registry one at a time; the -// unused ids are reserved so the tag values stay stable across those PRs. +// Tags identifying the producer subsystem of each inst_pkt_t entry, stored in the +// inst_pkt_t pair as an hsa::ClientID. // -// The underlying type is fixed to int64_t to match hsa::ClientID, which is how the -// tag is stored in inst_pkt_t. The enum is deliberately unscoped rather than an -// enum class: unscoped keeps the implicit conversion to ClientID, so the emplace -// sites need no cast, while still grouping the values under one type. +// The underlying type is fixed to int64_t to match hsa::ClientID. The enum is +// deliberately unscoped rather than an enum class: unscoped keeps the implicit +// conversion to ClientID, so the emplace sites need no cast, while still grouping +// the values under one type. // -// These values numerically overlap the per-queue registry's auto-incrementing -// ClientID, which also starts at 1 and is still in use by the services that have -// not been migrated yet. That is inert only because no consumer routes on the tag: -// counters' completed_cb and thread trace's post_kernel_call both identify their -// own packets by pointer lookup or dynamic_cast. Before adding a consumer that -// dispatches on these values, either finish migrating the remaining services off -// the registry or move this enum to a range the registry cannot produce. +// The values are negative because they share the inst_pkt_t tag space with the +// per-queue callback registry, whose QueueController::add_callback hands out +// strictly positive ids counting up from 1. Migrated services (counters, thread +// trace, SPM) select their own packets by comparing against these tags, so a tag +// drawn from the registry's range could match a packet the service does not own. +// Services that have not migrated yet still receive registry ids, so the two +// ranges must stay disjoint until the registry is gone. Any tag added here must +// likewise be negative; -1 is left unused because QueueController::add_callback +// uses it as a local "unassigned" initializer. enum client_id : int64_t { - COUNTERS_CLIENT_ID = 1, - THREAD_TRACE_CLIENT_ID = 2, - PC_SAMPLING_CLIENT_ID = 3, - SPM_CLIENT_ID = 4, + COUNTERS_CLIENT_ID = -1001, + THREAD_TRACE_CLIENT_ID = -1002, + PC_SAMPLING_CLIENT_ID = -1003, + SPM_CLIENT_ID = -1004, }; + +static_assert(COUNTERS_CLIENT_ID < 0 && THREAD_TRACE_CLIENT_ID < 0 && PC_SAMPLING_CLIENT_ID < 0 && + SPM_CLIENT_ID < 0, + "queue hook tags must stay disjoint from the positive ids handed out by " + "QueueController::add_callback"); } // namespace queue_hooks } // namespace hsa } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/tests/counter-collection-per-agent/per_agent_scoping_app.cpp b/projects/rocprofiler-sdk/tests/counter-collection-per-agent/per_agent_scoping_app.cpp index f9b5cd723d51..57fa6cabbc19 100644 --- a/projects/rocprofiler-sdk/tests/counter-collection-per-agent/per_agent_scoping_app.cpp +++ b/projects/rocprofiler-sdk/tests/counter-collection-per-agent/per_agent_scoping_app.cpp @@ -91,6 +91,10 @@ rocprofiler_context_id_t ctx_all_second = {.handle = 0}; // reorder HIP ordinals. std::vector gpu_agents = {}; +// HIP ordinal that could not be matched to a rocprofiler agent, or -1 when every device +// mapped. Reported by main() so a skip is distinguishable from a machine with one GPU. +int unmapped_device = -1; + // Intermediate storage used to build the HIP-to-rocprofiler mapping: keeps the // location_id (BDF) and PCI domain alongside the agent id until the mapping is done. struct agent_info_t @@ -254,6 +258,10 @@ tool_init(rocprofiler_client_finalize_t, void*) int device_count = 0; if(hipGetDeviceCount(&device_count) != hipSuccess || device_count < 2) return 0; + // A device with no matching agent must not be skipped silently: dropping it would shift + // every later ordinal, so gpu_agents[i] would stop being HIP device i and the contexts + // would be scoped to a different physical GPU than the workload runs on. Bail instead and + // let main() report the skip. for(int dev = 0; dev < device_count; ++dev) { hipDeviceProp_t prop{}; @@ -261,15 +269,25 @@ tool_init(rocprofiler_client_finalize_t, void*) const auto bus = static_cast(prop.pciBusID); const auto did = static_cast(prop.pciDeviceID); const auto dom = static_cast(prop.pciDomainID); + + bool matched = false; for(const auto& info : all_gpu_agent_info) { if(info.domain == dom && (info.location_id >> 8) == bus && ((info.location_id >> 3) & 0x1Fu) == did) { gpu_agents.push_back(info.id); + matched = true; break; } } + + // Only the first two ordinals are used, so an unmapped device beyond them is harmless. + if(!matched) + { + unmapped_device = dev; + break; + } } if(gpu_agents.size() < 2) return 0; @@ -371,7 +389,12 @@ main() if(device_count < 2 || gpu_agents.size() < 2) { - printf("SKIP: per-agent scoping requires at least 2 GPUs (found %d)\n", device_count); + if(unmapped_device >= 0 && unmapped_device < 2) + printf("SKIP: HIP device %d has no rocprofiler agent with matching PCI " + "coordinates, so per-agent scoping cannot be checked on this machine\n", + unmapped_device); + else + printf("SKIP: per-agent scoping requires at least 2 GPUs (found %d)\n", device_count); return 0; } @@ -406,6 +429,17 @@ main() scoped_ms, kNumStreams, kKernelsPer); + // This test only runs on a multi-GPU machine, so print the mapping and the raw counts the + // assertions below read: a control failure is otherwise indistinguishable from a scoping bug. + printf("hip device 0 -> agent %lu, hip device 1 -> agent %lu\n", + static_cast(agent_0.handle), + static_cast(agent_1.handle)); + printf("unrestricted: dispatches(gpu0)=%zu | scoped: records(gpu1)=%zu dispatches(gpu0)=%zu " + "records(gpu0)=%zu\n", + unrestricted_dispatches_gpu0, + scoped_records_gpu1, + scoped_dispatches_gpu0, + scoped_records_gpu0); // Exact assertions. These hold regardless of how fast the machine is. check(unrestricted_dispatches_gpu0 > 0,