diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp index 3196cf1f72..a7075d58d9 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp @@ -160,6 +160,10 @@ AQLPacketConstruct::construct_packet(const AmdExtTable& ext) const "could not generate stop packet"); throw_if_failed(hsa_ven_amd_aqlprofile_read(&profile, &pkt.read), "could not generate read packet"); + pkt.start.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + pkt.stop.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + pkt.read.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + return pkt_ptr; } 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 04c7aed9b6..32a506b165 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp @@ -37,6 +37,9 @@ namespace rocprofiler { namespace counters { +using ClientID = int64_t; +using inst_pkt_t = common::container:: + small_vector, ClientID>, 4>; class CounterController { public: @@ -136,9 +139,8 @@ destroy_counter_profile(uint64_t id) * We return an AQLPacket containing the start/stop/read packets for injection. */ std::unique_ptr -queue_cb(const std::shared_ptr& info, - const hsa::Queue& queue, - hsa::ClientID, +queue_cb(const std::shared_ptr& info, + const hsa::Queue& queue, const hsa::rocprofiler_packet& pkt, uint64_t kernel_id, const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids, @@ -245,9 +247,38 @@ queue_cb(const std::shared_ptr& info, ret_pkt = prof_config->pkt_generator->construct_packet( hsa::get_queue_controller().get_ext_table()); } + ret_pkt->before_krn_pkt.clear(); + ret_pkt->after_krn_pkt.clear(); + if(ret_pkt->empty) return ret_pkt; info->packet_return_map.wlock([&](auto& data) { data.emplace(ret_pkt.get(), prof_config); }); + auto&& CreateBarrierPacket = + [](hsa_signal_t* dependency_signal, + hsa_signal_t* completion_signal, + common::container::small_vector& _packets) { + hsa::rocprofiler_packet barrier{}; + barrier.barrier_and.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; + if(dependency_signal != nullptr) barrier.barrier_and.dep_signal[0] = *dependency_signal; + if(completion_signal != nullptr) + barrier.barrier_and.completion_signal = *completion_signal; + _packets.emplace_back(barrier.ext_amd_aql_pm4); + }; + + hsa_signal_t ready_signal = queue.ready_signal; + hsa_signal_t block_signal = queue.block_signal; + CreateBarrierPacket(nullptr, &ready_signal, ret_pkt->before_krn_pkt); + CreateBarrierPacket(&block_signal, &block_signal, ret_pkt->before_krn_pkt); + + ret_pkt->before_krn_pkt.push_back(ret_pkt->start); + ret_pkt->before_krn_pkt.end()->completion_signal.handle = 0; + ret_pkt->after_krn_pkt.push_back(ret_pkt->stop); + ret_pkt->after_krn_pkt.push_back(ret_pkt->read); + for(auto& aql_pkt : ret_pkt->after_krn_pkt) + { + aql_pkt.completion_signal.handle = 0; + } + return ret_pkt; } @@ -257,20 +288,32 @@ queue_cb(const std::shared_ptr& info, void completed_cb(const std::shared_ptr& info, const hsa::Queue&, - hsa::ClientID, hsa::rocprofiler_packet, - const hsa::Queue::queue_info_session_t& session, - std::unique_ptr pkt) + const hsa::Queue::queue_info_session_t& session, + inst_pkt_t& pkts) { - if(!info || !pkt) return; + if(!info || pkts.empty()) return; std::shared_ptr prof_config; // Get the Profile Config + std::unique_ptr pkt = nullptr; info->packet_return_map.wlock([&](auto& data) { - prof_config = data.at(pkt.get()); - data.erase(pkt.get()); + for(auto& [aql_pkt, _] : pkts) + { + const auto* profile = rocprofiler::common::get_val(data, aql_pkt.get()); + if(profile) + { + prof_config = *profile; + data.erase(aql_pkt.get()); + pkt = std::move(aql_pkt); + return; + } + } }); + if(!pkt) return; + hsa::profiler_serializer_kernel_completion_signal(session.queue.block_signal); + auto decoded_pkt = EvaluateAST::read_pkt(prof_config->pkt_generator.get(), *pkt); EvaluateAST::read_special_counters( *prof_config->agent, prof_config->required_special_counters, decoded_pkt); @@ -333,21 +376,17 @@ start_context(const context::context* ctx) cb->queue_id = controller.add_callback( std::nullopt, [=](const hsa::Queue& q, - hsa::ClientID c, const hsa::rocprofiler_packet& kern_pkt, uint64_t kernel_id, const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids, const context::correlation_id* correlation_id) { - return queue_cb(cb, q, c, kern_pkt, kernel_id, extern_corr_ids, correlation_id); + return queue_cb(cb, q, kern_pkt, kernel_id, extern_corr_ids, correlation_id); }, // Completion CB [=](const hsa::Queue& q, - hsa::ClientID c, hsa::rocprofiler_packet kern_pkt, const hsa::Queue::queue_info_session_t& session, - std::unique_ptr aql) { - completed_cb(cb, q, c, kern_pkt, session, std::move(aql)); - }); + inst_pkt_t& aql) { completed_cb(cb, q, kern_pkt, session, aql); }); } enabled = true; }); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp index 238e2fd88f..93e41d4916 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp @@ -22,8 +22,11 @@ #pragma once +#include "lib/common/container/small_vector.hpp" + #include #include +#include namespace rocprofiler { @@ -53,14 +56,16 @@ struct AQLPacket AQLPacket(const AQLPacket&) = delete; AQLPacket& operator=(const AQLPacket&) = delete; + bool command_buf_mallocd = false; + bool output_buffer_malloced = false; + bool empty = {true}; hsa_ven_amd_aqlprofile_profile_t profile = {}; hsa_ext_amd_aql_pm4_packet_t start = null_amd_aql_pm4_packet; hsa_ext_amd_aql_pm4_packet_t stop = null_amd_aql_pm4_packet; hsa_ext_amd_aql_pm4_packet_t read = null_amd_aql_pm4_packet; - bool command_buf_mallocd = false; - bool output_buffer_malloced = false; memory_pool_free_func_t free_func = nullptr; - bool empty = {true}; + common::container::small_vector before_krn_pkt = {}; + common::container::small_vector after_krn_pkt = {}; }; inline AQLPacket::AQLPacket(memory_pool_free_func_t func) 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 a8a436fe1b..9efa8fd192 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -156,30 +156,15 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data) } } } - // Calls our internal callbacks to callers who need to be notified post // kernel execution. queue_info_session.queue.signal_callback([&](const auto& map) { for(const auto& [client_id, cb_pair] : map) { - // If this is the client that gave us the AQLPacket, - // return it to that client otherwise notify. - if(queue_info_session.inst_pkt_id == client_id) - { - cb_pair.second(queue_info_session.queue, - client_id, - queue_info_session.kernel_pkt, - queue_info_session, - std::move(queue_info_session.inst_pkt)); - } - else - { - cb_pair.second(queue_info_session.queue, - client_id, - queue_info_session.kernel_pkt, - queue_info_session, - nullptr); - } + cb_pair.second(queue_info_session.queue, + queue_info_session.kernel_pkt, + queue_info_session, + queue_info_session.inst_pkt); } }); @@ -244,17 +229,13 @@ WriteInterceptor(const void* packets, { using context_array_t = Queue::context_array_t; - auto&& AddVendorSpecificPacket = [](hsa_ext_amd_aql_pm4_packet_t _packet, - hsa_signal_t _signal, - std::vector& _packets) { - _packets.emplace_back(_packet).ext_amd_aql_pm4.completion_signal = _signal; - }; - - auto&& CreateBarrierPacket = [](hsa_signal_t _signal, + auto&& CreateBarrierPacket = [](hsa_signal_t* dependency_signal, + hsa_signal_t* completion_signal, std::vector& _packets) { hsa_barrier_and_packet_t barrier{}; - barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; - barrier.dep_signal[0] = _signal; + barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; + if(dependency_signal != nullptr) barrier.dep_signal[0] = *dependency_signal; + if(completion_signal != nullptr) barrier.completion_signal = *completion_signal; _packets.emplace_back(barrier); }; @@ -319,32 +300,37 @@ WriteInterceptor(const void* packets, // Stores the instrumentation pkt (i.e. AQL packets for counter collection) // along with an ID of the client we got the packet from (this will be returned via // completed_cb_t) - ClientID inst_pkt_id = -1; - std::unique_ptr inst_pkt; + + inst_pkt_t inst_pkt; // Signal callbacks that a kernel_pkt is being enqueued queue.signal_callback([&](const auto& map) { for(const auto& [client_id, cb_pair] : map) { - if(auto maybe_pkt = cb_pair.first( - queue, client_id, kernel_pkt, kernel_id, extern_corr_ids, corr_id)) + if(auto maybe_pkt = + cb_pair.first(queue, kernel_pkt, kernel_id, extern_corr_ids, corr_id)) { - LOG_IF(FATAL, inst_pkt) - << "We do not support two injections into the HSA queue"; - inst_pkt = std::move(maybe_pkt); - inst_pkt_id = client_id; + inst_pkt.push_back(std::make_pair(std::move(maybe_pkt), client_id)); } } }); - constexpr auto dummy_signal = hsa_signal_t{.handle = 0}; - - // Write instrumentation start packet (if one exists) - if(inst_pkt && !inst_pkt->empty) + bool inserted_before = false; + for(const auto& pkt_injection : inst_pkt) { - inst_pkt->start.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - AddVendorSpecificPacket(inst_pkt->start, dummy_signal, transformed_packets); - CreateBarrierPacket(inst_pkt->start.completion_signal, transformed_packets); + for(const auto& pkt : pkt_injection.first->before_krn_pkt) + { + inserted_before = true; + transformed_packets.emplace_back(pkt); + } + } + + // Barrier packet is last packet inserted into queue + if(inserted_before) + { + CreateBarrierPacket(&transformed_packets.back().ext_amd_aql_pm4.completion_signal, + nullptr, + transformed_packets); } transformed_packets.emplace_back(kernel_pkt); @@ -363,15 +349,20 @@ WriteInterceptor(const void* packets, // Adding a barrier packet with the original packet's completion signal. queue.create_signal(0, &interrupt_signal); - if(inst_pkt && !inst_pkt->empty) + bool injected_end_pkt = false; + for(const auto& pkt_injection : inst_pkt) { - inst_pkt->stop.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - AddVendorSpecificPacket(inst_pkt->stop, dummy_signal, transformed_packets); - inst_pkt->read.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - AddVendorSpecificPacket(inst_pkt->read, interrupt_signal, transformed_packets); + for(const auto& pkt : pkt_injection.first->after_krn_pkt) + { + transformed_packets.emplace_back(pkt); + injected_end_pkt = true; + } + } - // Added Interrupt Signal with barrier and provided handler for it - CreateBarrierPacket(interrupt_signal, transformed_packets); + if(injected_end_pkt) + { + transformed_packets.back().ext_amd_aql_pm4.completion_signal = interrupt_signal; + CreateBarrierPacket(&interrupt_signal, nullptr, transformed_packets); } else { @@ -391,7 +382,6 @@ WriteInterceptor(const void* packets, interrupt_signal, new Queue::queue_info_session_t{.queue = queue, .inst_pkt = std::move(inst_pkt), - .inst_pkt_id = inst_pkt_id, .interrupt_signal = interrupt_signal, .tid = thr_id, .kernel_id = kernel_id, @@ -462,6 +452,8 @@ Queue::Queue(const AgentCache& agent, _ext_api.hsa_amd_queue_intercept_register_fn(_intercept_queue, WriteInterceptor, this)) << "Could not register interceptor"; + create_signal(0, &ready_signal); + create_signal(0, &block_signal); *queue = _intercept_queue; } @@ -504,5 +496,17 @@ Queue::remove_callback(ClientID id) if(map.erase(id) == 1) _notifiers--; }); } + +queue_state +Queue::get_state() const +{ + return _state; +} + +void +Queue::set_state(queue_state state) +{ + _state = state; +} } // namespace hsa } // namespace rocprofiler 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 43c1b683dc..d2b05b9337 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -57,6 +57,8 @@ struct correlation_id; namespace hsa { using ClientID = int64_t; +using inst_pkt_t = + common::container::small_vector, ClientID>, 4>; union rocprofiler_packet { @@ -92,6 +94,12 @@ union rocprofiler_packet rocprofiler_packet& operator=(const rocprofiler_packet&) = default; rocprofiler_packet& operator=(rocprofiler_packet&&) noexcept = default; }; +enum class queue_state +{ + normal = 0, + to_destroy = 1, + done_destroy = 2 +}; // Interceptor for a single specific queue class Queue @@ -109,8 +117,7 @@ public: std::unordered_map; Queue& queue; - std::unique_ptr inst_pkt = {}; - ClientID inst_pkt_id = 0; + inst_pkt_t inst_pkt = {}; hsa_signal_t interrupt_signal = {}; rocprofiler_thread_id_t tid = common::get_tid(); rocprofiler_kernel_id_t kernel_id = 0; @@ -128,19 +135,16 @@ public: // the queue. using queue_cb_t = std::function( const Queue&, - ClientID, const rocprofiler_packet&, uint64_t, const queue_info_session_t::external_corr_id_map_t&, const context::correlation_id*)>; // Signals the completion of the kernel packet. using completed_cb_t = std::function)>; + inst_pkt_t&)>; using callback_map_t = std::unordered_map>; - Queue(const AgentCache& agent, uint32_t size, hsa_queue_type32_t type, @@ -178,8 +182,14 @@ public: void register_callback(ClientID id, queue_cb_t enqueue_cb, completed_cb_t complete_cb); void remove_callback(ClientID id); - const CoreApiTable& core_api() const { return _core_api; } - const AmdExtTable& ext_api() const { return _ext_api; } + const CoreApiTable& core_api() const { return _core_api; } + const AmdExtTable& ext_api() const { return _ext_api; } + mutable std::mutex cv_mutex; + mutable std::condition_variable cv_ready_signal; + hsa_signal_t block_signal; + hsa_signal_t ready_signal; + queue_state get_state() const; + void set_state(queue_state state); private: std::atomic _notifiers = {0}; @@ -189,6 +199,7 @@ private: const AgentCache& _agent; rocprofiler::common::Synchronized _callbacks = {}; hsa_queue_t* _intercept_queue = nullptr; + queue_state _state = queue_state::normal; }; inline rocprofiler_queue_id_t @@ -205,4 +216,4 @@ Queue::signal_callback(FuncT&& func) const } } // namespace hsa -} // namespace rocprofiler +} // namespace rocprofiler \ No newline at end of file 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 54785186df..106bdacd37 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 @@ -61,7 +61,11 @@ create_queue(hsa_agent_t agent, get_queue_controller().get_core_table(), get_queue_controller().get_ext_table(), queue); + + get_queue_controller().profiler_serializer_register_ready_signal_handler( + new_queue->ready_signal, *queue); get_queue_controller().add_queue(*queue, std::move(new_queue)); + return HSA_STATUS_SUCCESS; } } @@ -104,6 +108,48 @@ QueueController::add_queue(hsa_queue_t* id, std::unique_ptr queue) void QueueController::destory_queue(hsa_queue_t* id) { + const auto* queue = get_queue_controller().get_queue(*id); + std::unique_lock cvlock(queue->cv_mutex); + profiler_serializer([&](auto& data) { + /*Deletes the queue to be destructed from the dispatch ready.*/ + data.dispatch_ready.erase( + std::remove_if( + data.dispatch_ready.begin(), + data.dispatch_ready.end(), + [&](auto& it) { + /*Deletes the queue to be destructed from the dispatch ready.*/ + if(it->get_id().handle == queue->get_id().handle) + { + if(data.dispatch_queue && + data.dispatch_queue->get_id().handle == queue->get_id().handle) + { + // insert fatal condition here + // ToDO [srnagara]: Need to find a solution rather than abort. + LOG(FATAL) + << "Queue is being destroyed while kernel launch is still active"; + } + return true; + } + return false; + }), + data.dispatch_ready.end()); + set_queue_state(queue_state::to_destroy, id); + /* + This lambda triggers the async ready handler. + The async ready handler then unregisters itself + and sets the queue state to done_destroy for which + the condition variable here is waiting for. + */ + auto trigger_ready_async_handler = [queue]() { + get_queue_controller().get_core_table().hsa_signal_store_screlease_fn( + queue->ready_signal, 0); + }; + trigger_ready_async_handler(); + }); + queue->cv_ready_signal.wait( + cvlock, [queue] { return queue->get_state() == queue_state::done_destroy; }); + if(queue->block_signal.handle != 0) + get_queue_controller().get_core_table().hsa_signal_destroy_fn(queue->block_signal); _queues.wlock([&](auto& map) { map.erase(id); }); } @@ -216,6 +262,103 @@ QueueController::get_queue(const hsa_queue_t& _hsa_queue) const _hsa_queue); } +template +void +QueueController::profiler_serializer(FuncT&& lambda) +{ + _profiler_serializer.wlock(std::forward(lambda)); +} + +namespace +{ +/* + Function name: AsyncSignalReadyHandler + Argument: hsa signal value for which the async handler was called + and pointer to the data. + Description: This async handler is invoked when the queue is ready + to launch a kernel. It first, resets the queue's ready signal to 1. + It then checks if there is any queue which has a kernel currently dispatched. + If yes, it pushes the queue to the dispatch ready else + it enables the dispatch for the given queue. + Return : It returns true since we need this handler to be invoked + each time the queue's ready signal (used for entire queue) is set to 0. + If we had a separate signal for every dispatch in the queue then we don't + need this to be invoked more than once in which case we would return false. +*/ +bool +profiler_serializer_ready_signal_handler(hsa_signal_value_t /* signal_value */, void* data) +{ + auto* hsa_queue = static_cast(data); + const auto* queue = get_queue_controller().get_queue(*hsa_queue); + get_queue_controller().profiler_serializer([&](auto& serializer) { + { + std::lock_guard cv_lock(queue->cv_mutex); + if(queue->get_state() == queue_state::to_destroy) + { + get_queue_controller().set_queue_state(queue_state::done_destroy, hsa_queue); + get_queue_controller().get_core_table().hsa_signal_destroy_fn(queue->ready_signal); + queue->cv_ready_signal.notify_one(); + return; + } + } + get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(queue->ready_signal, + 1); + if(serializer.dispatch_queue == nullptr) + { + get_queue_controller().get_core_table().hsa_signal_store_screlease_fn( + queue->block_signal, 0); + serializer.dispatch_queue = queue; + } + else + { + serializer.dispatch_ready.push_back(queue); + } + }); + return true; +} +} // namespace + +void +profiler_serializer_kernel_completion_signal(hsa_signal_t queue_block_signal) +{ + get_queue_controller().profiler_serializer([queue_block_signal](auto& serializer) { + assert(serializer.dispatch_queue != nullptr); + serializer.dispatch_queue = nullptr; + get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(queue_block_signal, + 1); + if(!serializer.dispatch_ready.empty()) + { + auto queue = serializer.dispatch_ready.front(); + serializer.dispatch_ready.erase(serializer.dispatch_ready.begin()); + get_queue_controller().get_core_table().hsa_signal_store_screlease_fn( + queue->block_signal, 0); + serializer.dispatch_queue = queue; + } + }); +} + +void +QueueController::set_queue_state(enum queue_state state, hsa_queue_t* hsa_queue) +{ + _queues.wlock([&](auto& map) { map[hsa_queue]->set_state(state); }); +} + +/* + Function name: SignalAsyncReadyHandler. + Argument : The signal value and pointer to the data to + pass to the handler. + Description : Registers a asynchronous callback function + for the ready signal to be invoked when it goes to zero. +*/ +void +QueueController::profiler_serializer_register_ready_signal_handler(const hsa_signal_t& signal, + void* data) const +{ + hsa_status_t status = get_ext_table().hsa_amd_signal_async_handler_fn( + signal, HSA_SIGNAL_CONDITION_EQ, 0, profiler_serializer_ready_signal_handler, data); + if(status != HSA_STATUS_SUCCESS) LOG(FATAL) << "hsa_amd_signal_async_handler failed"; +} + void QueueController::iterate_queues(const queue_iterator_cb_t& cb) const { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp index dd008da3f2..9d4832447e 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp @@ -36,6 +36,26 @@ namespace rocprofiler { namespace hsa { +/*This is a profiler serializer. It should be instantiated +only once for the profiler. The following is the +description of each field. +1. dispatch_queue - The queue to which the currently dispatched kernel + belongs to. + At any given time, in serialization only one kernel + can be executing. +2. dispatch_ready- It is a software data structure which holds + the queues which have a kernel ready to be dispatched. + This stores the queues in FIFO order. +3. serializer_mutex - The mutex is used for thread synchronization + while accessing the singleton instance of this structure. +Currently, in case of profiling kernels are serialized by default. +*/ +struct profiler_serializer_t +{ + const Queue* dispatch_queue{nullptr}; + std::deque dispatch_ready; +}; + // Tracks and manages HSA queues class QueueController { @@ -69,6 +89,12 @@ public: const Queue* get_queue(const hsa_queue_t&) const; void iterate_queues(const queue_iterator_cb_t&) const; + void set_queue_state(queue_state state, hsa_queue_t* hsa_queue); + void profiler_serializer_register_ready_signal_handler(const hsa_signal_t& signal, + void* data) const; + void add_dispatch_ready(const Queue* queue); + template + void profiler_serializer(FuncT&& lambda); private: using agent_callback_tuple_t = @@ -77,11 +103,12 @@ private: using client_id_map_t = std::unordered_map; using agent_cache_map_t = std::unordered_map; - CoreApiTable _core_table = {}; - AmdExtTable _ext_table = {}; - common::Synchronized _queues = {}; - common::Synchronized _callback_cache = {}; - agent_cache_map_t _supported_agents = {}; + CoreApiTable _core_table = {}; + AmdExtTable _ext_table = {}; + common::Synchronized _queues = {}; + common::Synchronized _callback_cache = {}; + agent_cache_map_t _supported_agents = {}; + common::Synchronized _profiler_serializer; }; QueueController& @@ -92,5 +119,9 @@ queue_controller_init(HsaApiTable* table); void queue_controller_fini(); + +void +profiler_serializer_kernel_completion_signal(hsa_signal_t queue_block_signal); + } // namespace hsa } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/tests/CMakeLists.txt b/projects/rocprofiler-sdk/tests/CMakeLists.txt index 83749a972f..c92cbb99b0 100644 --- a/projects/rocprofiler-sdk/tests/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/CMakeLists.txt @@ -49,3 +49,4 @@ add_subdirectory(async-copy-tracing) # rocprofv3 validation tests add_subdirectory(rocprofv3) +add_subdirectory(counter-collection) diff --git a/projects/rocprofiler-sdk/tests/apps/CMakeLists.txt b/projects/rocprofiler-sdk/tests/apps/CMakeLists.txt index 86d833ceff..9b16f9d734 100644 --- a/projects/rocprofiler-sdk/tests/apps/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/apps/CMakeLists.txt @@ -9,6 +9,7 @@ set(CMAKE_BUILD_RPATH "\$ORIGIN:\$ORIGIN/../lib") # applications used by integration tests which DO NOT link to rocprofiler-sdk-roctx add_subdirectory(simple-transpose) +add_subdirectory(multistream) set(CMAKE_BUILD_RPATH "\$ORIGIN:\$ORIGIN/../lib:$" diff --git a/projects/rocprofiler-sdk/tests/apps/multistream/CMakeLists.txt b/projects/rocprofiler-sdk/tests/apps/multistream/CMakeLists.txt new file mode 100644 index 0000000000..67144e8353 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/apps/multistream/CMakeLists.txt @@ -0,0 +1,46 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +if(NOT CMAKE_HIP_COMPILER) + find_program( + amdclangpp_EXECUTABLE + NAMES amdclang++ + HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm + PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm + PATH_SUFFIXES bin llvm/bin NO_CACHE) + mark_as_advanced(amdclangpp_EXECUTABLE) + + if(amdclangpp_EXECUTABLE) + set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}") + endif() +endif() + +project(rocprofiler-test-app-multistream LANGUAGES CXX HIP) + +foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO) + if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "") + set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}") + endif() +endforeach() + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_HIP_STANDARD 17) +set(CMAKE_HIP_EXTENSIONS OFF) +set(CMAKE_HIP_STANDARD_REQUIRED ON) + +set_source_files_properties(multistream_app.cpp PROPERTIES LANGUAGE HIP) +add_executable(multistream) +target_sources(multistream PRIVATE multistream_app.cpp) +target_compile_options(multistream PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow -Werror) + +find_package(Threads REQUIRED) +target_link_libraries(multistream PRIVATE Threads::Threads) + +install( + TARGETS multistream + DESTINATION bin + COMPONENT tests) diff --git a/projects/rocprofiler-sdk/tests/apps/multistream/multistream_app.cpp b/projects/rocprofiler-sdk/tests/apps/multistream/multistream_app.cpp new file mode 100644 index 0000000000..35138a97b0 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/apps/multistream/multistream_app.cpp @@ -0,0 +1,116 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include +#define HIP_ASSERT(call) \ + do \ + { \ + hipError_t err = call; \ + if(err != hipSuccess) \ + { \ + fprintf(stderr, "%s\n", hipGetErrorString(err)); \ + abort(); \ + } \ + } while(0) + +__device__ int counter = 0; +__global__ void +add(int n, float* x, float* y) +{ + if(__hip_atomic_load(&counter, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_AGENT) != 0) + { + abort(); + } + __hip_atomic_fetch_add(&counter, 1, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM); + + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for(int i = index; i < n; i += stride) + y[i] = x[i] + y[i]; + __hip_atomic_fetch_add(&counter, -1, __ATOMIC_RELEASE, __HIP_MEMORY_SCOPE_SYSTEM); +} + +void +LaunchMultiStreamKernels() +{ + int N = 1 << 4; + float* x = new float[N]; + float* y = new float[N]; + float* d_x; + float* d_y; + // Allocate Unified Memory -- accessible from CPU or GPU + HIP_ASSERT(hipMallocManaged(&d_x, N * sizeof(float))); + HIP_ASSERT(hipMallocManaged(&d_y, N * sizeof(float))); + + // initialize x and y arrays on the host + for(int i = 0; i < N; i++) + { + x[i] = 1.0f; + y[i] = 2.0f; + } + std::vector hip_streams; + for(int i = 0; i < 100; i++) + { + hipStream_t stream; + HIP_ASSERT(hipStreamCreate(&stream)); + hip_streams.push_back(stream); + } + HIP_ASSERT(hipMemcpy(d_x, x, N * sizeof(float), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(d_y, y, N * sizeof(float), hipMemcpyHostToDevice)); + + // Launch kernel on 1M elements on the GPU + int blockSize = 64; + // This Kernel will always be launched with one wave + int numBlocks = 1; + for(int i = 0; i < 100; i++) + { + for(size_t j = 0; j < hip_streams.size(); j++) + { + hipLaunchKernelGGL(add, numBlocks, blockSize, 0, hip_streams[j], N, d_x, d_y); + } + } + + // Wait for GPU to finish before accessing on host + HIP_ASSERT(hipDeviceSynchronize()); + + HIP_ASSERT(hipMemcpy(x, d_x, N * sizeof(float), hipMemcpyDeviceToHost)); + HIP_ASSERT(hipMemcpy(y, d_y, N * sizeof(float), hipMemcpyDeviceToHost)); + + // Free memory + HIP_ASSERT(hipFree(d_x)); + HIP_ASSERT(hipFree(d_y)); + + delete[] x; + delete[] y; + + for(size_t i = 0; i < hip_streams.size(); i++) + { + HIP_ASSERT(hipStreamDestroy(hip_streams[i])); + } +} + +int +main() +{ + LaunchMultiStreamKernels(); +} diff --git a/projects/rocprofiler-sdk/tests/common/serialization.hpp b/projects/rocprofiler-sdk/tests/common/serialization.hpp index 9e638498f1..389be0899d 100644 --- a/projects/rocprofiler-sdk/tests/common/serialization.hpp +++ b/projects/rocprofiler-sdk/tests/common/serialization.hpp @@ -231,6 +231,15 @@ save(ArchiveT& ar, rocprofiler_buffer_tracing_hsa_api_record_t data) save_buffer_tracing_api_record(ar, data); } +template +void +save(ArchiveT& ar, rocprofiler_record_counter_t data) +{ + SAVE_DATA_FIELD(id); + SAVE_DATA_FIELD(counter_value); + SAVE_DATA_FIELD(corr_id); +} + template void save(ArchiveT& ar, rocprofiler_buffer_tracing_hip_api_record_t data) diff --git a/projects/rocprofiler-sdk/tests/counter-collection/CMakeLists.txt b/projects/rocprofiler-sdk/tests/counter-collection/CMakeLists.txt new file mode 100644 index 0000000000..e4409a5ad5 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/counter-collection/CMakeLists.txt @@ -0,0 +1,47 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-tests-counter-collection + LANGUAGES CXX + VERSION 0.0.0) + +find_package(rocprofiler-sdk REQUIRED) + +if(ROCPROFILER_MEMCHECK_PRELOAD_ENV) + set(PRELOAD_ENV + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}:$") +else() + set(PRELOAD_ENV "LD_PRELOAD=$") +endif() + +add_test(NAME test-counter-collection-execute COMMAND $) + +set_tests_properties( + test-counter-collection-execute + PROPERTIES + TIMEOUT + 45 + LABELS + "integration-tests" + ENVIRONMENT + "${PRELOAD_ENV};HSA_TOOLS_LIB=$;ROCPROFILER_TOOL_OUTPUT_FILE=counter-collection-test.json;ROCPROFILER_TOOL_CONTEXTS=COUNTER_COLLECTION;ROCPROF_COUNTERS=SQ_WAVES_sum" + FAIL_REGULAR_EXPRESSION + "threw an exception") + +foreach(FILENAME validate.py pytest.ini conftest.py) + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} + ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) +endforeach() + +add_test(NAME test-counter-collection-validate + COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input + ${CMAKE_CURRENT_BINARY_DIR}/counter-collection-test.json) + +set_tests_properties( + test-counter-collection-validate + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS + test-counter-collection-execute FAIL_REGULAR_EXPRESSION + "threw an exception") diff --git a/projects/rocprofiler-sdk/tests/counter-collection/conftest.py b/projects/rocprofiler-sdk/tests/counter-collection/conftest.py new file mode 100644 index 0000000000..4f87dc38cf --- /dev/null +++ b/projects/rocprofiler-sdk/tests/counter-collection/conftest.py @@ -0,0 +1,20 @@ +#!/usr/bin/env python3 + +import json +import pytest + + +def pytest_addoption(parser): + parser.addoption( + "--input", + action="store", + default="counter-collection-test.json", + help="Input JSON", + ) + + +@pytest.fixture +def input_data(request): + filename = request.config.getoption("--input") + with open(filename, "r") as inp: + return json.load(inp) diff --git a/projects/rocprofiler-sdk/tests/counter-collection/pytest.ini b/projects/rocprofiler-sdk/tests/counter-collection/pytest.ini new file mode 100644 index 0000000000..899bc38408 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/counter-collection/pytest.ini @@ -0,0 +1,4 @@ + +[pytest] +addopts = --durations=20 -ras -vv +testpaths = validate.py diff --git a/projects/rocprofiler-sdk/tests/counter-collection/validate.py b/projects/rocprofiler-sdk/tests/counter-collection/validate.py new file mode 100644 index 0000000000..42e9d6a9f8 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/counter-collection/validate.py @@ -0,0 +1,34 @@ +#!/usr/bin/env python3 + +import sys +import pytest + + +# helper function +def node_exists(name, data, min_len=1): + assert name in data + assert data[name] is not None + assert len(data[name]) >= min_len + + +def test_data_structure(input_data): + """verify minimum amount of expected data is present""" + node_exists("rocprofiler-sdk-json-tool", input_data) + rocp_data = input_data + node_exists("names", rocp_data["rocprofiler-sdk-json-tool"]["buffer_records"]) + node_exists( + "counter_collection", rocp_data["rocprofiler-sdk-json-tool"]["buffer_records"] + ) + + +def test_counter_values(input_data): + data = input_data + counter_info = {} + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["counter_collection"]: + value = itr["counter_value"] + assert value == 1 + + +if __name__ == "__main__": + exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) + sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/kernel-tracing/CMakeLists.txt b/projects/rocprofiler-sdk/tests/kernel-tracing/CMakeLists.txt index 6cfc0c1876..68af6c6e76 100644 --- a/projects/rocprofiler-sdk/tests/kernel-tracing/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/kernel-tracing/CMakeLists.txt @@ -11,7 +11,7 @@ project( find_package(rocprofiler-sdk REQUIRED) set(PYTEST_ARGS) -if(ROCPROFILER_MEMCHECK MATCHES "(Address|Thread)Sanitizer") +if(ROCPROFILER_MEMCHECK MATCHES "(Address|Thread)Sanitizer" OR ROCPROFILER_BUILD_CODECOV) set(PYTEST_ARGS -k "not test_total_runtime") endif() diff --git a/projects/rocprofiler-sdk/tests/tools/json-tool.cpp b/projects/rocprofiler-sdk/tests/tools/json-tool.cpp index cfe1fc95e1..883677aac9 100644 --- a/projects/rocprofiler-sdk/tests/tools/json-tool.cpp +++ b/projects/rocprofiler-sdk/tests/tools/json-tool.cpp @@ -43,6 +43,8 @@ #include #include +#include + #include #include #include @@ -57,6 +59,7 @@ #include #include #include +#include #include #include #include @@ -361,15 +364,123 @@ struct marker_api_callback_record_t } }; -auto code_object_records = std::deque{}; -auto kernel_symbol_records = std::deque{}; -auto hsa_api_cb_records = std::deque{}; -auto marker_api_cb_records = std::deque{}; -auto hip_api_cb_records = std::deque{}; +auto code_object_records = std::deque{}; +auto kernel_symbol_records = std::deque{}; +auto hsa_api_cb_records = std::deque{}; +auto marker_api_cb_records = std::deque{}; +auto counter_collection_bf_records = std::deque{}; +auto hip_api_cb_records = std::deque{}; rocprofiler_thread_id_t push_external_correlation(); +void +counter_collection_buffered(rocprofiler_context_id_t, /*context*/ + rocprofiler_buffer_id_t, /*buffer_id*/ + rocprofiler_record_header_t** headers, + size_t num_headers, + void*, /*user_data*/ + uint64_t /*drop_count*/) +{ + if(num_headers == 0) + throw std::runtime_error{"rocprofiler invoked a buffer callback with no headers " + "this should never happen"}; + + else if(headers == nullptr) + throw std::runtime_error{"rocprofiler invoked a buffer callback with a null pointer to the " + "array of headers. this should never happen"}; + for(size_t i = 0; i < num_headers; ++i) + { + auto* header = headers[i]; + if(header->category == ROCPROFILER_BUFFER_CATEGORY_COUNTERS && header->kind == 0) + { + auto* profiler_record = static_cast(header->payload); + counter_collection_bf_records.emplace_back(*profiler_record); + } + } +} + +void +dispatch_callback(rocprofiler_queue_id_t, /*queue_id*/ + const rocprofiler_agent_t* agent, + rocprofiler_correlation_id_t, /*correlation_id*/ + const hsa_kernel_dispatch_packet_t*, /*dispatch_packet*/ + uint64_t, /*kernel_id*/ + void* /*callback_data_args*/, + rocprofiler_profile_config_id_t* config) +{ + static std::shared_mutex m_mutex = {}; + static std::unordered_map profile_cache = {}; + + auto search_cache = [&]() { + if(auto pos = profile_cache.find(agent->id.handle); pos != profile_cache.end()) + { + *config = pos->second; + return true; + } + return false; + }; + + { + auto rlock = std::shared_lock{m_mutex}; + if(search_cache()) return; + } + + auto wlock = std::unique_lock{m_mutex}; + if(search_cache()) return; + + // Counters we want to collect (here its SQ_WAVES_sum) + auto* counters_env = getenv("ROCPROF_COUNTERS"); + if(std::string(counters_env) != "SQ_WAVES_sum") + LOG(FATAL) << "Counter not supported in the test tool"; + + std::set counters_to_collect = {"SQ_WAVES_sum"}; + // GPU Counter IDs + std::vector gpu_counters; + + // Iterate through the agents and get the counters available on that agent + ROCPROFILER_CALL(rocprofiler_iterate_agent_supported_counters( + agent->id, + []([[maybe_unused]] rocprofiler_agent_id_t id, + rocprofiler_counter_id_t* counters, + size_t num_counters, + void* user_data) { + std::vector* vec = + static_cast*>(user_data); + for(size_t i = 0; i < num_counters; i++) + { + vec->push_back(counters[i]); + } + return ROCPROFILER_STATUS_SUCCESS; + }, + static_cast(&gpu_counters)), + "Could not fetch supported counters"); + + std::vector collect_counters; + // Look for the counters contained in counters_to_collect in gpu_counters + for(auto& counter : gpu_counters) + { + const char* name; + size_t size; + ROCPROFILER_CALL(rocprofiler_query_counter_name(counter, &name, &size), + "Could not query name"); + if(counters_to_collect.count(std::string(name)) > 0) + { + collect_counters.push_back(counter); + } + } + + // Create a colleciton profile for the counters + rocprofiler_profile_config_id_t profile; + ROCPROFILER_CALL(rocprofiler_create_profile_config( + agent->id, collect_counters.data(), collect_counters.size(), &profile), + "Could not construct profile cfg"); + + profile_cache.emplace(agent->id.handle, profile); + // Return the profile to collect those counters for this dispatch + *config = profile; +} + void tool_tracing_callback(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t* /*user_data*/, @@ -584,12 +695,14 @@ rocprofiler_context_id_t hip_api_buffered_ctx = {}; rocprofiler_context_id_t marker_api_buffered_ctx = {}; rocprofiler_context_id_t kernel_dispatch_ctx = {}; rocprofiler_context_id_t memory_copy_ctx = {}; +rocprofiler_context_id_t counter_collection_ctx = {}; // buffers rocprofiler_buffer_id_t hsa_api_buffered_buffer = {}; rocprofiler_buffer_id_t hip_api_buffered_buffer = {}; rocprofiler_buffer_id_t marker_api_buffered_buffer = {}; rocprofiler_buffer_id_t kernel_dispatch_buffer = {}; rocprofiler_buffer_id_t memory_copy_buffer = {}; +rocprofiler_buffer_id_t counter_collection_buffer = {}; auto contexts = std::unordered_map{ {"HSA_API_CALLBACK", &hsa_api_callback_ctx}, @@ -600,13 +713,15 @@ auto contexts = std::unordered_map{ {"HIP_API_BUFFERED", &hip_api_buffered_ctx}, {"MARKER_API_BUFFERED", &marker_api_buffered_ctx}, {"KERNEL_DISPATCH", &kernel_dispatch_ctx}, - {"MEMORY_COPY", &memory_copy_ctx}}; + {"MEMORY_COPY", &memory_copy_ctx}, + {"COUNTER_COLLECTION", &counter_collection_ctx}}; -auto buffers = std::array{&hsa_api_buffered_buffer, +auto buffers = std::array{&hsa_api_buffered_buffer, &hip_api_buffered_buffer, &marker_api_buffered_buffer, &kernel_dispatch_buffer, - &memory_copy_buffer}; + &memory_copy_buffer, + &counter_collection_buffer}; auto agents = std::vector{}; @@ -785,6 +900,20 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) marker_api_buffered_buffer), "buffer tracing service configure"); + ROCPROFILER_CALL(rocprofiler_create_buffer(counter_collection_ctx, + 4096, + 2048, + ROCPROFILER_BUFFER_POLICY_LOSSLESS, + counter_collection_buffered, + nullptr, + &counter_collection_buffer), + "buffer creation"); + + ROCPROFILER_CALL( + rocprofiler_configure_buffered_dispatch_profile_counting_service( + counter_collection_ctx, counter_collection_buffer, dispatch_callback, nullptr), + "setup buffered service"); + ROCPROFILER_CALL( rocprofiler_configure_buffer_tracing_service(kernel_dispatch_ctx, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, @@ -885,7 +1014,8 @@ tool_fini(void* tool_data) << ", memory_copy_records=" << memory_copy_records.size() << ", hsa_api_bf_records=" << hsa_api_bf_records.size() << ", hip_api_bf_records=" << hip_api_bf_records.size() - << ", marker_api_bf_records=" << marker_api_bf_records.size() << " ...\n" + << ", marker_api_bf_records=" << marker_api_bf_records.size() + << ", counter_collection_records" << counter_collection_bf_records.size() << "...\n" << std::flush; auto* _call_stack = static_cast(tool_data); @@ -966,6 +1096,7 @@ tool_fini(void* tool_data) json_ar(cereal::make_nvp("hsa_api_traces", hsa_api_bf_records)); json_ar(cereal::make_nvp("hip_api_traces", hip_api_bf_records)); json_ar(cereal::make_nvp("marker_api_traces", marker_api_bf_records)); + json_ar(cereal::make_nvp("counter_collection", counter_collection_bf_records)); } catch(std::exception& e) { std::cerr << "[" << getpid() << "][" << __FUNCTION__ @@ -994,6 +1125,11 @@ start() { if(itr.second && !is_active(*itr.second)) { + if(itr.first == "COUNTER_COLLECTION") + { + auto* counters = getenv("ROCPROF_COUNTERS"); + if(!counters) continue; + } ROCPROFILER_CALL(rocprofiler_start_context(*itr.second), "context start"); } }