From fa1b9e67abf01a450e9e9808d2e6ae10d6296723 Mon Sep 17 00:00:00 2001 From: Giovanni Lenzi Baraldi Date: Thu, 15 Aug 2024 13:57:13 -0300 Subject: [PATCH] ATT Agent fixes and improvements (#1011) * Tidying ATT dispatch API. ATT Agent to be initialized with rest of profiler. Removing read_index-based wait. * Formatting * Adding some input validation * Add perf test for agent * Removing async --- .../amd_detail/thread_trace_dispatch.h | 2 - .../lib/rocprofiler-sdk/context/context.cpp | 7 +- .../lib/rocprofiler-sdk/context/context.hpp | 3 +- source/lib/rocprofiler-sdk/hsa/aql_packet.cpp | 19 +- .../rocprofiler-sdk/hsa/queue_controller.cpp | 12 +- source/lib/rocprofiler-sdk/registration.cpp | 3 +- .../rocprofiler-sdk/thread_trace/att_core.cpp | 209 +++++++++++++----- .../rocprofiler-sdk/thread_trace/att_core.hpp | 62 +++--- .../thread_trace/att_service.cpp | 23 +- .../thread_trace/tests/att_packet_test.cpp | 60 ++++- tests/thread-trace/multi_dispatch.cpp | 25 +-- 11 files changed, 286 insertions(+), 139 deletions(-) diff --git a/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h b/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h index 686a77d1dd..7e156cdfd0 100644 --- a/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h +++ b/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h @@ -40,8 +40,6 @@ ROCPROFILER_EXTERN_C_INIT typedef enum { ROCPROFILER_ATT_CONTROL_NONE = 0, - ROCPROFILER_ATT_CONTROL_START = 1, - ROCPROFILER_ATT_CONTROL_STOP = 2, ROCPROFILER_ATT_CONTROL_START_AND_STOP = 3 } rocprofiler_att_control_flags_t; diff --git a/source/lib/rocprofiler-sdk/context/context.cpp b/source/lib/rocprofiler-sdk/context/context.cpp index c7839d6035..976a7e7c97 100644 --- a/source/lib/rocprofiler-sdk/context/context.cpp +++ b/source/lib/rocprofiler-sdk/context/context.cpp @@ -321,7 +321,8 @@ start_context(rocprofiler_context_id_t context_id) auto status = ROCPROFILER_STATUS_SUCCESS; if(cfg->counter_collection) rocprofiler::counters::start_context(cfg); - if(cfg->thread_trace) cfg->thread_trace->start_context(); + if(cfg->agent_thread_trace) cfg->agent_thread_trace->start_context(); + if(cfg->dispatch_thread_trace) cfg->dispatch_thread_trace->start_context(); if(cfg->agent_counter_collection) status = rocprofiler::counters::start_agent_ctx(cfg); #if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0 if(cfg->pc_sampler) status = rocprofiler::pc_sampling::start_service(cfg); @@ -355,7 +356,9 @@ stop_context(rocprofiler_context_id_t idx) rocprofiler::counters::stop_context(const_cast(_expected)); } - if(_expected->thread_trace) _expected->thread_trace->stop_context(); + if(_expected->agent_thread_trace) _expected->agent_thread_trace->stop_context(); + if(_expected->dispatch_thread_trace) + _expected->dispatch_thread_trace->stop_context(); if(_expected->agent_counter_collection) { diff --git a/source/lib/rocprofiler-sdk/context/context.hpp b/source/lib/rocprofiler-sdk/context/context.hpp index 895b6be4bc..449cb4db38 100644 --- a/source/lib/rocprofiler-sdk/context/context.hpp +++ b/source/lib/rocprofiler-sdk/context/context.hpp @@ -128,7 +128,8 @@ struct context std::unique_ptr agent_counter_collection = {}; std::unique_ptr pc_sampler = {}; - std::unique_ptr thread_trace = {}; + std::unique_ptr dispatch_thread_trace = {}; + std::unique_ptr agent_thread_trace = {}; }; // set the client index needs to be called before allocate_context() diff --git a/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp b/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp index eeb10c4cd9..410db946b1 100644 --- a/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp +++ b/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp @@ -36,6 +36,9 @@ namespace rocprofiler { namespace hsa { +constexpr uint16_t VENDOR_BIT = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; +constexpr uint16_t BARRIER_BIT = 1 << HSA_PACKET_HEADER_BARRIER; + hsa_status_t CounterAQLPacket::CounterMemoryPool::Alloc(void** ptr, size_t size, desc_t flags, void* data) { @@ -114,13 +117,10 @@ CounterAQLPacket::CounterAQLPacket(aqlprofile_agent_handle_t ag reinterpret_cast(&pool)); if(status != HSA_STATUS_SUCCESS) ROCP_FATAL << "Could not create PMC packets!"; - packets.start_packet.header = 1 << HSA_PACKET_HEADER_BARRIER; - // Read and stop packets require the barrier bit set to wait for the dispatch packet - auto header = (HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE) | - (1 << HSA_PACKET_HEADER_BARRIER); - packets.stop_packet.header = header; - packets.read_packet.header = header; - empty = false; + packets.start_packet.header = VENDOR_BIT; + packets.stop_packet.header = VENDOR_BIT | BARRIER_BIT; + packets.read_packet.header = VENDOR_BIT | BARRIER_BIT; + empty = false; } hsa_status_t @@ -181,9 +181,8 @@ TraceControlAQLPacket::TraceControlAQLPacket(const TraceMemoryPool& _tr tracepool.get()); CHECK_HSA(status, "failed to create ATT packet"); - packets.start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - packets.stop_packet.header = (HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE) | - (1 << HSA_PACKET_HEADER_BARRIER); + packets.start_packet.header = VENDOR_BIT | BARRIER_BIT; + packets.stop_packet.header = VENDOR_BIT | BARRIER_BIT; packets.start_packet.completion_signal = hsa_signal_t{.handle = 0}; packets.stop_packet.completion_signal = hsa_signal_t{.handle = 0}; this->empty = false; diff --git a/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp b/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp index 3398b97a3d..49518d9e97 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp @@ -147,8 +147,8 @@ QueueController::add_queue(hsa_queue_t* id, std::unique_ptr queue) { for(const auto& itr : context::get_registered_contexts()) { - if(itr->thread_trace) - itr->thread_trace->resource_init(queue->get_agent(), get_core_table(), get_ext_table()); + if(auto* trace = itr->dispatch_thread_trace.get()) + trace->resource_init(queue->get_agent(), get_core_table(), get_ext_table()); } CHECK(queue); @@ -175,11 +175,11 @@ QueueController::destroy_queue(hsa_queue_t* id) for(const auto& itr : context::get_registered_contexts()) { - if(!itr->thread_trace) continue; + if(!itr->dispatch_thread_trace) continue; _queues.wlock([&](auto& map) { if(map.find(id) != map.end()) - itr->thread_trace->resource_deinit(map.at(id)->get_agent()); + itr->dispatch_thread_trace->resource_deinit(map.at(id)->get_agent()); }); } @@ -264,7 +264,7 @@ QueueController::init(CoreApiTable& core_table, AmdExtTable& ext_table) auto enable_intercepter = false; for(const auto& itr : context::get_registered_contexts()) { - constexpr auto expected_context_size = 208UL; + constexpr auto expected_context_size = 216UL; static_assert( sizeof(context::context) == expected_context_size, "If you added a new field to context struct, make sure there is a check here if it " @@ -277,7 +277,7 @@ QueueController::init(CoreApiTable& core_table, AmdExtTable& ext_table) itr->buffered_tracer->domains(ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH)); if(itr->counter_collection || itr->pc_sampler || has_kernel_tracing || - itr->agent_counter_collection || itr->thread_trace) + itr->agent_counter_collection || itr->agent_thread_trace || itr->dispatch_thread_trace) { enable_intercepter = true; break; diff --git a/source/lib/rocprofiler-sdk/registration.cpp b/source/lib/rocprofiler-sdk/registration.cpp index ba5d0b3960..3b03ec4f00 100644 --- a/source/lib/rocprofiler-sdk/registration.cpp +++ b/source/lib/rocprofiler-sdk/registration.cpp @@ -631,6 +631,7 @@ finalize() set_fini_status(-1); hsa::async_copy_fini(); hsa::queue_controller_fini(); + thread_trace::finalize(); page_migration::finalize(); #if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0 // WARNING: this must precede `code_object::finalize()` @@ -788,7 +789,7 @@ rocprofiler_set_api_table(const char* name, rocprofiler::hsa::async_copy_init(hsa_api_table, lib_instance); rocprofiler::code_object::initialize(hsa_api_table); - rocprofiler::thread_trace::code_object::initialize(hsa_api_table); + rocprofiler::thread_trace::initialize(hsa_api_table); #if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0 rocprofiler::pc_sampling::code_object::initialize(hsa_api_table); #endif diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp b/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp index 5be0eed9c7..2f3c30ae45 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp @@ -51,12 +51,14 @@ } \ } -constexpr size_t ROCPROFILER_QUEUE_SIZE = 64; - namespace rocprofiler { namespace thread_trace { +constexpr size_t QUEUE_SIZE = 128; +constexpr uint64_t MIN_BUFFER_SIZE = 1 << 18; // 2 pages per SE +constexpr uint64_t MAX_BUFFER_SIZE = std::numeric_limits::max(); // aqlprofile limit + struct cbdata_t { rocprofiler_att_shader_data_callback_t cb_fn; @@ -65,10 +67,78 @@ struct cbdata_t common::Synchronized> client; -bool -ThreadTracerQueue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet) +CoreApiTable& +get_core() { - const uint64_t write_idx = add_write_index_relaxed_fn(queue, 1); + static CoreApiTable api{}; + return api; +} + +AmdExtTable& +get_ext() +{ + static AmdExtTable api{}; + return api; +} + +bool +thread_trace_parameter_pack::are_params_valid() const +{ + if(shader_cb_fn == nullptr) + { + ROCP_WARNING << "Callback cannot be null!"; + return false; + } + + if(shader_engine_mask == 0) return false; + + if(buffer_size > MAX_BUFFER_SIZE || buffer_size < MIN_BUFFER_SIZE) + { + ROCP_WARNING << "Invalid buffer size: " << buffer_size; + return false; + } + + if(target_cu > 0xF) return false; + if(simd_select > 0xF) return false; // Only 16 CUs and 4 SIMDs + + return true; +} + +class Signal +{ +public: + Signal(hsa_ext_amd_aql_pm4_packet_t* packet) + { + get_ext().hsa_amd_signal_create_fn(0, 0, nullptr, 0, &signal); + packet->completion_signal = signal; + get_core().hsa_signal_store_screlease_fn(signal, 1); + }; + ~Signal() + { + WaitOn(); + get_core().hsa_signal_destroy_fn(signal); + } + Signal(Signal& other) = delete; + Signal(const Signal& other) = delete; + Signal& operator=(Signal& other) = delete; + Signal& operator=(const Signal& other) = delete; + + void WaitOn() const + { + auto* wait_fn = get_core().hsa_signal_wait_scacquire_fn; + while(wait_fn(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED)) + ; + } + + hsa_signal_t signal; + std::atomic released{false}; +}; + +std::unique_ptr +ThreadTracerQueue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet, bool bWait) +{ + std::unique_ptr signal{}; + const uint64_t write_idx = add_write_index_relaxed_fn(queue, 1); size_t index = (write_idx % queue->size) * sizeof(hsa_ext_amd_aql_pm4_packet_t); auto* queue_slot = reinterpret_cast(size_t(queue->base_address) + index); // NOLINT @@ -76,23 +146,15 @@ ThreadTracerQueue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet) const auto* slot_data = reinterpret_cast(packet); memcpy(&queue_slot[1], &slot_data[1], sizeof(hsa_ext_amd_aql_pm4_packet_t) - sizeof(uint32_t)); + if(bWait) + signal = + std::make_unique(reinterpret_cast(queue_slot)); auto* header = reinterpret_cast*>(queue_slot); header->store(slot_data[0], std::memory_order_release); signal_store_screlease_fn(queue->doorbell_signal, write_idx); - int loops = 0; - while(load_read_index_relaxed_fn(queue) <= write_idx) - { - loops++; - usleep(1); - if(loops > 10000) // Add loop limit to prevent hang. TODO: Remove once stability proven - { - ROCP_ERROR << "Codeobj packet submission failed!"; - return false; - } - } - return true; + return signal; } ThreadTracerQueue::ThreadTracerQueue(thread_trace_parameter_pack _params, @@ -106,7 +168,7 @@ ThreadTracerQueue::ThreadTracerQueue(thread_trace_parameter_pack _params, control_packet = factory->construct_control_packet(); auto status = coreapi.hsa_queue_create_fn(cache.get_hsa_agent(), - ROCPROFILER_QUEUE_SIZE, + QUEUE_SIZE, HSA_QUEUE_TYPE_SINGLE, nullptr, nullptr, @@ -147,8 +209,10 @@ ThreadTracerQueue::~ThreadTracerQueue() control_packet->clear(); control_packet->populate_after(); + std::vector> wait_idx{}; + for(auto& after_packet : control_packet->after_krn_pkt) - Submit(&after_packet); + wait_idx.emplace_back(Submit(&after_packet, true)); } /** @@ -200,11 +264,8 @@ ThreadTracerQueue::load_codeobj(code_object_id_t id, uint64_t addr, uint64_t siz if(!queue || active_traces.load() < 1) return; - auto packet = factory->construct_load_marker_packet(id, addr, size); - bool bSuccess = Submit(&packet->packet); - - if(!bSuccess) // If something went wrong, don't delete packet to avoid CP memory access fault - packet.release(); + auto packet = factory->construct_load_marker_packet(id, addr, size); + Submit(&packet->packet, true)->WaitOn(); } void @@ -215,11 +276,8 @@ ThreadTracerQueue::unload_codeobj(code_object_id_t id) if(!control_packet->remove_codeobj(id)) return; if(!queue || active_traces.load() < 1) return; - auto packet = factory->construct_unload_marker_packet(id); - bool bSuccess = Submit(&packet->packet); - - if(!bSuccess) // If something went wrong, don't delete packet to avoid CP memory access fault - packet.release(); + auto packet = factory->construct_unload_marker_packet(id); + Submit(&packet->packet, true)->WaitOn(); } void @@ -301,15 +359,13 @@ DispatchThreadTracer::pre_kernel_call(const hsa::Queue& queue, auto it = agents.find(queue.get_agent().get_hsa_agent()); assert(it != agents.end() && it->second != nullptr); - auto packet = it->second->get_control(bool(control_flags & ROCPROFILER_ATT_CONTROL_START)); + auto packet = it->second->get_control(true); post_move_data.fetch_add(1); maybe_add_serialization(packet); - if((control_flags & ROCPROFILER_ATT_CONTROL_START) != 0) packet->populate_before(); - - if((control_flags & ROCPROFILER_ATT_CONTROL_STOP) != 0) packet->populate_after(); - + packet->populate_before(); + packet->populate_after(); return packet; } @@ -383,7 +439,7 @@ DispatchThreadTracer::start_context() } void -DispatchThreadTracer::stop_context() +DispatchThreadTracer::stop_context() // NOLINT { client.wlock([&](auto& client_id) { if(!client_id) return; @@ -398,34 +454,33 @@ DispatchThreadTracer::stop_context() } void -AgentThreadTracer::resource_init(const hsa::AgentCache& cache, - const CoreApiTable& coreapi, - const AmdExtTable& ext) +AgentThreadTracer::resource_init(const CoreApiTable& coreapi, const AmdExtTable& ext) { - auto id = cache.get_rocp_agent()->id; + auto rocp_agents = rocprofiler::agent::get_agents(); + std::unique_lock lk(agent_mut); - if(params.find(id) == params.end()) return; - - if(tracers.find(id) != tracers.end()) + for(const auto* rocp_agent : rocp_agents) { - tracers.at(id)->active_queues.fetch_add(1); - return; + auto id = rocp_agent->id; + const auto* cache = rocprofiler::agent::get_agent_cache(rocp_agent); + + if(tracers.find(id) != tracers.end()) + ROCP_WARNING << "Agent configured twice: " << id.handle; + else if(params.find(id) == params.end()) + ROCP_INFO << "Skipping agent " << id.handle; + else if(cache == nullptr) + ROCP_WARNING << "Invalid agent id: " << id.handle; + else + tracers[id] = std::make_unique(params.at(id), *cache, coreapi, ext); } - tracers.emplace(id, std::make_unique(params.at(id), cache, coreapi, ext)); } void -AgentThreadTracer::resource_deinit(const hsa::AgentCache& cache) +AgentThreadTracer::resource_deinit() { - auto id = cache.get_rocp_agent()->id; std::unique_lock lk(agent_mut); - - if(params.find(id) == params.end()) return; - if(tracers.find(id) == tracers.end()) return; - - auto& tracer = *tracers.at(id); - if(tracer.active_queues.fetch_sub(1) == 1) tracers.erase(id); + tracers.clear(); } void @@ -439,13 +494,15 @@ AgentThreadTracer::start_context() return; } + std::vector> wait_list{}; + for(auto& [_, tracer] : tracers) { auto packet = tracer->get_control(true); packet->populate_before(); - for(auto& start : packet->before_krn_pkt) - tracer->Submit(&start); + auto sig = tracer->SubmitAndSignalLast(packet->before_krn_pkt); + if(sig) wait_list.emplace_back(std::move(sig)); } } @@ -460,16 +517,50 @@ AgentThreadTracer::stop_context() return; } - for(auto& [_, tracer] : tracers) + std::vector>> + wait_list{}; + + for(auto& [id, tracer] : tracers) { auto packet = tracer->get_control(false); packet->populate_after(); - for(auto& stop : packet->after_krn_pkt) - tracer->Submit(&stop); + std::optional write_index{}; + auto signal = tracer->SubmitAndSignalLast(packet->after_krn_pkt); + if(signal) + wait_list.push_back({tracer.get(), packet->GetHandle(), std::move(signal)}); // NOLINT + } + for(auto& [tracer, handle, signal] : wait_list) + { + signal->WaitOn(); rocprofiler_user_data_t userdata{.ptr = tracer->params.callback_userdata}; - tracer->iterate_data(packet->GetHandle(), userdata); + tracer->iterate_data(handle, userdata); + } +} + +void +initialize(HsaApiTable* table) +{ + assert(table->core_ && table->amd_ext_); + get_core() = *table->core_; + get_ext() = *table->amd_ext_; + + code_object::initialize(table); + + for(auto& ctx : context::get_registered_contexts()) + { + if(ctx->agent_thread_trace) + ctx->agent_thread_trace->resource_init(*table->core_, *table->amd_ext_); + } +} + +void +finalize() +{ + for(auto& ctx : context::get_registered_contexts()) + { + if(ctx->agent_thread_trace) ctx->agent_thread_trace->resource_deinit(); } } diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp b/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp index dcaab786fc..6011cac3d4 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp +++ b/source/lib/rocprofiler-sdk/thread_trace/att_core.hpp @@ -71,6 +71,8 @@ struct thread_trace_parameter_pack static constexpr size_t DEFAULT_SE_MASK = 0x21; static constexpr size_t DEFAULT_BUFFER_SIZE = 0x8000000; static constexpr size_t PERFCOUNTER_SIMD_MASK_SHIFT = 28; + + bool are_params_valid() const; }; class ThreadTracerQueue @@ -99,7 +101,19 @@ public: std::unique_ptr control_packet; std::unique_ptr factory; - bool Submit(hsa_ext_amd_aql_pm4_packet_t* packet); + [[nodiscard]] std::unique_ptr Submit(hsa_ext_amd_aql_pm4_packet_t* packet, + bool bWait); + + template + [[nodiscard]] std::unique_ptr SubmitAndSignalLast(VecType vec) + { + for(size_t i = 0; i < vec.size(); i++) + { + auto sig = Submit(&vec.at(i), i == vec.size() - 1); + if(sig) return sig; + } + return nullptr; + } private: std::unique_ptr codeobj_reg{nullptr}; @@ -112,19 +126,7 @@ private: decltype(hsa_queue_destroy)* queue_destroy_fn{nullptr}; }; -class ThreadTracerInterface -{ -public: - ThreadTracerInterface() = default; - virtual ~ThreadTracerInterface() = default; - - virtual void start_context() = 0; - virtual void stop_context() = 0; - virtual void resource_init(const hsa::AgentCache&, const CoreApiTable&, const AmdExtTable&) = 0; - virtual void resource_deinit(const hsa::AgentCache&) = 0; -}; - -class DispatchThreadTracer : public ThreadTracerInterface +class DispatchThreadTracer { using code_object_id_t = uint64_t; using AQLPacketPtr = std::unique_ptr; @@ -134,12 +136,12 @@ public: DispatchThreadTracer(thread_trace_parameter_pack _params) : params(std::move(_params)) {} - ~DispatchThreadTracer() override = default; + ~DispatchThreadTracer() = default; - void start_context() override; - void stop_context() override; - void resource_init(const hsa::AgentCache&, const CoreApiTable&, const AmdExtTable&) override; - void resource_deinit(const hsa::AgentCache&) override; + void start_context(); + void stop_context(); + void resource_init(const hsa::AgentCache&, const CoreApiTable&, const AmdExtTable&); + void resource_deinit(const hsa::AgentCache&); std::unique_ptr pre_kernel_call(const hsa::Queue& queue, uint64_t kernel_id, @@ -157,16 +159,16 @@ public: thread_trace_parameter_pack params; }; -class AgentThreadTracer : public ThreadTracerInterface +class AgentThreadTracer { public: - AgentThreadTracer() = default; - ~AgentThreadTracer() override = default; + AgentThreadTracer() = default; + ~AgentThreadTracer() = default; - void start_context() override; - void stop_context() override; - void resource_init(const hsa::AgentCache&, const CoreApiTable&, const AmdExtTable&) override; - void resource_deinit(const hsa::AgentCache&) override; + void start_context(); + void stop_context(); + void resource_init(const CoreApiTable&, const AmdExtTable&); + void resource_deinit(); void add_agent(rocprofiler_agent_id_t id, thread_trace_parameter_pack _params) { @@ -180,11 +182,17 @@ public: } std::map> tracers{}; - std::map params; + std::map params{}; std::mutex agent_mut; }; +void +initialize(HsaApiTable* table); + +void +finalize(); + }; // namespace thread_trace } // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp b/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp index 277029cec1..64471d51ef 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/att_service.cpp @@ -30,6 +30,9 @@ #include "lib/rocprofiler-sdk/registration.hpp" #include "rocprofiler-sdk/amd_detail/thread_trace.h" +using DispatchThreadTracer = rocprofiler::thread_trace::DispatchThreadTracer; +using AgentThreadTracer = rocprofiler::thread_trace::AgentThreadTracer; + extern "C" { rocprofiler_status_t ROCPROFILER_API rocprofiler_configure_dispatch_thread_trace_service( @@ -45,7 +48,8 @@ rocprofiler_configure_dispatch_thread_trace_service( auto* ctx = rocprofiler::context::get_mutable_registered_context(context_id); if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_STARTED; - if(ctx->thread_trace) return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED; + if(ctx->dispatch_thread_trace) return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED; + if(ctx->agent_thread_trace) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; auto pack = rocprofiler::thread_trace::thread_trace_parameter_pack{}; @@ -54,6 +58,8 @@ rocprofiler_configure_dispatch_thread_trace_service( pack.shader_cb_fn = shader_callback; pack.callback_userdata = callback_userdata; + if(pack.dispatch_cb_fn == nullptr) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; + auto id_map = rocprofiler::counters::getPerfCountersIdMap(); for(size_t p = 0; p < num_parameters; p++) { @@ -83,7 +89,9 @@ rocprofiler_configure_dispatch_thread_trace_service( } } - ctx->thread_trace = std::make_unique(pack); + if(!pack.are_params_valid()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; + + ctx->dispatch_thread_trace = std::make_unique(pack); return ROCPROFILER_STATUS_SUCCESS; } @@ -96,14 +104,14 @@ rocprofiler_configure_agent_thread_trace_service( rocprofiler_att_shader_data_callback_t shader_callback, void* callback_userdata) { - using AgentThreadTracer = rocprofiler::thread_trace::AgentThreadTracer; if(rocprofiler::registration::get_init_status() > -1) return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED; auto* ctx = rocprofiler::context::get_mutable_registered_context(context_id); if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_STARTED; + if(ctx->dispatch_thread_trace) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; - if(!ctx->thread_trace) ctx->thread_trace = std::make_unique(); + if(!ctx->agent_thread_trace) ctx->agent_thread_trace = std::make_unique(); auto pack = rocprofiler::thread_trace::thread_trace_parameter_pack{}; @@ -140,11 +148,10 @@ rocprofiler_configure_agent_thread_trace_service( } } - auto* agent_tracer = dynamic_cast(ctx->thread_trace.get()); - if(agent_tracer == nullptr || agent_tracer->has_agent(agent)) - return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED; + if(!pack.are_params_valid()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - agent_tracer->add_agent(agent, pack); + assert(ctx->agent_thread_trace); + ctx->agent_thread_trace->add_agent(agent, pack); return ROCPROFILER_STATUS_SUCCESS; } } diff --git a/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp b/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp index 5ac8e46544..a8929ef1ad 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp @@ -209,7 +209,7 @@ TEST(thread_trace, perfcounters_configure_test) nullptr); auto* context = rocprofiler::context::get_mutable_registered_context(ctx); - auto* tracer = dynamic_cast(context->thread_trace.get()); + auto* tracer = context->dispatch_thread_trace.get(); ASSERT_NE(tracer, nullptr); ASSERT_EQ(tracer->params.perfcounter_ctrl, 1); @@ -250,3 +250,61 @@ TEST(thread_trace, perfcounters_aql_options_test) context::pop_client(1); hsa_shut_down(); } + +rocprofiler_status_t +query_available_agents(rocprofiler_agent_version_t /* version */, + const void** agents, + size_t num_agents, + void* ctx_ptr) +{ + for(size_t idx = 0; idx < num_agents; idx++) + { + const auto* agent = static_cast(agents[idx]); + if(agent->type != ROCPROFILER_AGENT_TYPE_GPU) continue; + + std::vector params; + params.push_back({ROCPROFILER_ATT_PARAMETER_TARGET_CU, {1}}); + params.push_back({ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, {0xF}}); + params.push_back({ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, {0x1000000}}); + params.push_back({ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, {0xF}}); + params.push_back({ROCPROFILER_ATT_PARAMETER_PERFCOUNTERS_CTRL, {1}}); + + { + auto metrics = rocprofiler::counters::getMetricsForAgent("gfx90a"); + + rocprofiler_att_parameter_t att_param; + att_param.type = ROCPROFILER_ATT_PARAMETER_PERFCOUNTER; + att_param.simd_mask = 0xF; + for(auto& metric : metrics) + if(metric.name() == "SQ_WAVES") rocprofiler_counter_id_t{.handle = metric.id()}; + + params.push_back(att_param); + } + + rocprofiler_configure_agent_thread_trace_service( + *reinterpret_cast(ctx_ptr), + params.data(), + params.size(), + agent->id, + [](int64_t, void*, size_t, rocprofiler_user_data_t) {}, + nullptr); + } + return ROCPROFILER_STATUS_SUCCESS; +} + +TEST(thread_trace, agent_configure_test) +{ + test_init(); + + registration::init_logging(); + registration::set_init_status(-1); + context::push_client(1); + rocprofiler_context_id_t ctx; + ROCPROFILER_CALL(rocprofiler_create_context(&ctx), "context creation failed"); + + ROCPROFILER_CALL(rocprofiler_query_available_agents(ROCPROFILER_AGENT_INFO_VERSION_0, + &query_available_agents, + sizeof(rocprofiler_agent_t), + &ctx), + "Failed to find GPU agents"); +} diff --git a/tests/thread-trace/multi_dispatch.cpp b/tests/thread-trace/multi_dispatch.cpp index 6a5dc4577f..287d1766bb 100644 --- a/tests/thread-trace/multi_dispatch.cpp +++ b/tests/thread-trace/multi_dispatch.cpp @@ -54,37 +54,18 @@ rocprofiler_att_control_flags_t dispatch_callback(rocprofiler_queue_id_t /* queue_id */, const rocprofiler_agent_t* /* agent */, rocprofiler_correlation_id_t /* correlation_id */, - rocprofiler_kernel_id_t kernel_id, + rocprofiler_kernel_id_t /* kernel_id */, rocprofiler_dispatch_id_t /* dispatch_id */, rocprofiler_user_data_t* dispatch_userdata, void* userdata) { C_API_BEGIN + assert(userdata && "Dispatch callback passed null!"); - ToolData& tool = *reinterpret_cast(userdata); dispatch_userdata->ptr = userdata; - static std::atomic call_id{0}; - static std::string_view desired_func_name = "branching_kernel"; - - try - { - auto& kernel_name = tool.kernel_id_to_kernel_name.at(kernel_id); - if(kernel_name.find(desired_func_name) == std::string::npos) - return ROCPROFILER_ATT_CONTROL_NONE; - - int id = call_id.fetch_add(1); - if(id == 0) - return ROCPROFILER_ATT_CONTROL_START; - else if(id == 1) - return ROCPROFILER_ATT_CONTROL_STOP; - } catch(...) - { - std::cerr << "Could not find kernel id: " << kernel_id << std::endl; - } - C_API_END - return ROCPROFILER_ATT_CONTROL_NONE; + return ROCPROFILER_ATT_CONTROL_START_AND_STOP; } int