diff --git a/projects/rocprofiler/src/core/counters/metrics/derived_counters.xml b/projects/rocprofiler/src/core/counters/metrics/derived_counters.xml index 5ab53a0c28..d9217bc50d 100755 --- a/projects/rocprofiler/src/core/counters/metrics/derived_counters.xml +++ b/projects/rocprofiler/src/core/counters/metrics/derived_counters.xml @@ -353,8 +353,8 @@ - - + + @@ -393,43 +393,43 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/projects/rocprofiler/src/core/counters/metrics/eval_metrics.cpp b/projects/rocprofiler/src/core/counters/metrics/eval_metrics.cpp index c1f5b427ff..ee2e84b01a 100644 --- a/projects/rocprofiler/src/core/counters/metrics/eval_metrics.cpp +++ b/projects/rocprofiler/src/core/counters/metrics/eval_metrics.cpp @@ -27,40 +27,43 @@ struct block_status_t { uint32_t group_index; }; +struct aqlprofile_event_t : public hsa_ven_amd_aqlprofile_event_t +{ + bool operator==(const aqlprofile_event_t& other) const { + return this->block_name == other.block_name && + this->block_index == other.block_index && + this->counter_id ==other.counter_id; + } +}; + +template <> +struct std::hash +{ + std::size_t operator()(const aqlprofile_event_t& k) const { + return (int(k.block_name)<<20) ^ (int(k.counter_id)<<10) ^ int(k.block_index); + } +}; + typedef struct { - std::vector* results; + std::unordered_map results; size_t index; } callback_data_t; -static inline bool IsEventMatch(const hsa_ven_amd_aqlprofile_event_t& event1, - const hsa_ven_amd_aqlprofile_event_t& event2) { - return (event1.block_name == event2.block_name) && (event1.block_index == event2.block_index) && - (event1.counter_id == event2.counter_id); -} - hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type, hsa_ven_amd_aqlprofile_info_data_t* info_data, void* data) { - hsa_status_t status = HSA_STATUS_SUCCESS; + if (info_type != HSA_VEN_AMD_AQLPROFILE_INFO_PMC_DATA) return HSA_STATUS_SUCCESS; callback_data_t* passed_data = reinterpret_cast(data); - try { - for (auto data_it = passed_data->results->begin(); data_it != passed_data->results->end(); - ++data_it) { - if (info_type != HSA_VEN_AMD_AQLPROFILE_INFO_PMC_DATA) continue; - if (!IsEventMatch(info_data->pmc_data.event, (*data_it)->event)) continue; - - // stores event result from each event separately - (*data_it)->xcc_vals.push_back(info_data->pmc_data.result); - // stores accumulated event result from all xccs - (*data_it)->val_double += info_data->pmc_data.result; - } - } catch (std::exception& e) { - std::cout << "caught an exception in eval_metrics.cpp:pmcCallback(): " << e.what() << std::endl; - } + auto it = passed_data->results.find(aqlprofile_event_t{info_data->pmc_data.event}); + if (it == passed_data->results.end()) + return HSA_STATUS_ERROR; + auto* res = it->second; + res->xcc_vals.push_back(info_data->pmc_data.result); + res->val_double += info_data->pmc_data.result; passed_data->index += 1; - return status; + return HSA_STATUS_SUCCESS; } @@ -187,10 +190,13 @@ bool metrics::ExtractMetricEvents( bool metrics::GetCounterData(hsa_ven_amd_aqlprofile_profile_t* profile, hsa_agent_t gpu_agent, std::vector& results_list) { size_t gpu_xcc_count = HSASupport_Singleton::GetInstance().GetHSAAgentInfo(gpu_agent.handle).GetDeviceInfo().getXccCount(); - callback_data_t callback_data{&results_list, 0}; + + callback_data_t callback_data{}; + for (auto* res : results_list) + callback_data.results[aqlprofile_event_t{res->event}] = res; hsa_status_t status = hsa_ven_amd_aqlprofile_iterate_data(profile, pmcCallback, &callback_data); - for (auto& data : *(callback_data.results)) + for (auto& data : results_list) { size_t xcc_count = (data->event.block_name != HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_UMC) ? gpu_xcc_count : 1; std::vector xcc_results = std::move(data->xcc_vals); diff --git a/projects/rocprofiler/src/core/hsa/hsa_support.cpp b/projects/rocprofiler/src/core/hsa/hsa_support.cpp index 7daab9ffbc..1356f1916a 100644 --- a/projects/rocprofiler/src/core/hsa/hsa_support.cpp +++ b/projects/rocprofiler/src/core/hsa/hsa_support.cpp @@ -844,8 +844,8 @@ void HSASupport_Singleton::InitKsymbols() { ksymbols_flag.exchange(false, std::memory_order_release); } { - std::lock_guard lock(kernel_names_map_lock); - kernel_names = new std::map>(); + std::unique_lock lock(kernel_names_map_lock); + kernel_names = new std::unordered_map(); kernel_names_flag.exchange(false, std::memory_order_release); } } @@ -858,7 +858,7 @@ void HSASupport_Singleton::FinitKsymbols() { ksymbols_flag.exchange(true, std::memory_order_release); } if (!kernel_names_flag.load(std::memory_order_relaxed)) { - std::lock_guard lock(kernel_names_map_lock); + std::unique_lock lock(kernel_names_map_lock); kernel_names->clear(); delete kernel_names; kernel_names_flag.exchange(true, std::memory_order_release); diff --git a/projects/rocprofiler/src/core/hsa/hsa_support.h b/projects/rocprofiler/src/core/hsa/hsa_support.h index 4cf2302a8c..5aeb86edbb 100644 --- a/projects/rocprofiler/src/core/hsa/hsa_support.h +++ b/projects/rocprofiler/src/core/hsa/hsa_support.h @@ -33,7 +33,7 @@ #include #include #include - +#include #include "rocprofiler.h" #include "src/core/hardware/hsa_info.h" @@ -152,8 +152,8 @@ class HSASupport_Singleton { HSAAgentInfo& GetHSAAgentInfo(uint64_t agent_handle); HSAAgentInfo& GetHSAAgentInfo(Agent::DeviceInfo device_info); Agent::DeviceInfo& GetDeviceInfo(HSAAgentInfo* agent_info); - std::mutex kernel_names_map_lock; - std::map>* kernel_names; + std::shared_mutex kernel_names_map_lock; + std::unordered_map* kernel_names; std::mutex ksymbol_map_lock; std::map* ksymbols; std::mutex signals_timestamps_map_lock; diff --git a/projects/rocprofiler/src/core/hsa/packets/packets_generator.cpp b/projects/rocprofiler/src/core/hsa/packets/packets_generator.cpp index 0394ef574d..d991232a65 100644 --- a/projects/rocprofiler/src/core/hsa/packets/packets_generator.cpp +++ b/projects/rocprofiler/src/core/hsa/packets/packets_generator.cpp @@ -43,6 +43,7 @@ #include "src/core/counters/metrics/metrics.h" #include "src/core/hardware/hsa_info.h" +#include "src/core/hsa/packets/packets_generator.h" #define ASSERTM(exp, msg) assert(((void)msg, exp)) @@ -160,16 +161,20 @@ void CheckPacketReqiurements() { // Initialize the PM4 commands with having the CPU&GPU agents, the counters, // counters count to output three packets which are start, stop and read // packets -std::vector> -InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, - std::vector& counter_names, rocprofiler_session_id_t session_id, - bool is_spm) { +std::unique_ptr InitializeAqlPackets( + hsa_agent_t cpu_agent, + hsa_agent_t gpu_agent, + std::vector& counter_names, + rocprofiler_session_id_t session_id, + bool is_spm +) { hsa_status_t status = HSA_STATUS_SUCCESS; rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton = rocprofiler::ROCProfiler_Singleton::GetInstance(); rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance(); - if (!counters_added.load(std::memory_order_acquire)) { + if (!counters_added.load(std::memory_order_acquire)) + { for (auto& name : counter_names) { if (rocprofiler_singleton.HasActiveSession()) { rocprofiler_singleton.GetSession(session_id)->GetProfiler()->AddCounterName(name); @@ -205,7 +210,10 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, } // do { - rocprofiler::profiling_context_t* context = new rocprofiler::profiling_context_t(); + auto prof_context = std::make_unique(hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn); + auto* context = prof_context->context.get(); + auto* profile = prof_context->profile.get(); + context->gpu_agent = gpu_agent; auto result = results_list.begin(); std::map, uint32_t> block_max_events_count; @@ -292,120 +300,93 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, context->results_map = results_map; context->metrics_dict = metricsDict[gpu_agent.handle]; - hsa_ven_amd_aqlprofile_parameter_t* params = {}; - packet_t* start_packet = new packet_t(); packet_t* stop_packet = new packet_t(); packet_t* read_packet = new packet_t(); - std::vector> - profiles = std::vector< - std::pair>(); - - // if (context->events_list.size() <= 0) { // std::cerr << "Error: No events to profile" << std::endl; // abort(); // } // Preparing the profile structure to get the packets - hsa_ven_amd_aqlprofile_event_type_t profile_type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC; - if (is_spm) profile_type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_TRACE; - hsa_ven_amd_aqlprofile_profile_t* profile = - new hsa_ven_amd_aqlprofile_profile_t{gpu_agent, - profile_type, - &(context->events_list[0]), - static_cast(context->events_list.size()), - params, - 0, - 0, - 0}; + profile->agent = gpu_agent; + profile->type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC; + profile->events = &(context->events_list[0]); + profile->event_count = static_cast(context->events_list.size()); size_t ag_list_count = 1; // rocprofiler::hsa_support::GetCPUAgentList().size(); hsa_agent_t ag_list[ag_list_count]; ag_list[0] = gpu_agent; - if (context->events_list.size() > 0) { - // Preparing an Getting the size of the command and output buffers - status = hsa_ven_amd_aqlprofile_start(profile, NULL); - // CHECK_HSA_STATUS("Error: Getting Buffers Size", status); + if (context->events_list.size() == 0) + return prof_context; - if (profile->command_buffer.size > 0 && profile->output_buffer.size > 0) { - status = HSA_STATUS_ERROR; - size_t size = profile->command_buffer.size; - size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK; - if (size <= 0) { - std::cerr << __FILE__ << ":" << __LINE__ << " " - << "Error: Command buffer given size is " << size << std::endl; - abort(); - } - status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn( - agentInfo.cpu_pool_, size, 0, reinterpret_cast(&(profile->command_buffer.ptr))); - if (status != HSA_STATUS_SUCCESS) { - profile->command_buffer.ptr = malloc(size); - /*numa_alloc_onnode( - size, - rocprofiler::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle).getNumaNode());*/ - if (profile->command_buffer.ptr == NULL) { - std::cerr << __FILE__ << ":" << __LINE__ << " " - << "Error: allocating memory for command buffer using NUMA" << std::endl; - abort(); - } - } else { - // Both the CPU and GPU can access the memory - status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn( - ag_list_count, ag_list, NULL, profile->command_buffer.ptr); - CHECK_HSA_STATUS("Error: Allowing access to Command Buffer", status); - } - if (!is_spm) { - status = HSA_STATUS_ERROR; - size_t size = profile->output_buffer.size; - size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK; - if (size <= 0) { - std::cerr << __FILE__ << ":" << __LINE__ << " " - << "Error: Output buffer given size is " << size << std::endl; - abort(); - } - status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn( - agentInfo.kernarg_pool_, size, 0, - reinterpret_cast(&profile->output_buffer.ptr)); - if (status != HSA_STATUS_SUCCESS) { - profile->output_buffer.ptr = malloc(size); - /*numa_alloc_onnode( - size, - rocprofiler::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle) - .getNumaNode());*/ - if (profile->output_buffer.ptr == NULL) { - std::cerr << __FILE__ << ":" << __LINE__ << " " - << "Error: allocating memory for output buffer using NUMA" << std::endl; - abort(); - } - } else { - status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn( - ag_list_count, ag_list, NULL, profile->output_buffer.ptr); - CHECK_HSA_STATUS("Error: GPU Agent can't have output buffer access", status); - memset(profile->output_buffer.ptr, 0x0, profile->output_buffer.size); - } - } else { - profile->output_buffer.size = 0; - } - status = hsa_ven_amd_aqlprofile_start(profile, start_packet); - // CHECK_HSA_STATUS("Error: Creating Start Packet\n", status); - status = hsa_ven_amd_aqlprofile_stop(profile, stop_packet); - // CHECK_HSA_STATUS("Error: Creating Stop Packet\n", status); - status = hsa_ven_amd_aqlprofile_read(profile, read_packet); - // CHECK_HSA_STATUS("Error: Creating Read Packet\n", status); + // Preparing an Getting the size of the command and output buffers + status = hsa_ven_amd_aqlprofile_start(profile, NULL); + CHECK_HSA_STATUS("Error: Getting Buffers Size", status); - context->start_packet = start_packet; - context->stop_packet = stop_packet; - context->read_packet = read_packet; - } + if (profile->command_buffer.size == 0 || profile->output_buffer.size == 0) + { + std::cerr << __FILE__ << ":" << __LINE__ << "Error: Did not return buffer size" << std::endl; + abort(); } - // add profiles - profiles.emplace_back(std::make_pair(context, profile)); - // } while (events_list.size() > 0); - return profiles; + status = HSA_STATUS_ERROR; + size_t size = (profile->command_buffer.size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK; + status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn( + agentInfo.cpu_pool_, size, 0, reinterpret_cast(&(profile->command_buffer.ptr))); + if (status != HSA_STATUS_SUCCESS) { + profile->command_buffer.ptr = malloc(size); + /*numa_alloc_onnode( + size, + rocprofiler::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle).getNumaNode());*/ + if (profile->command_buffer.ptr == NULL) { + std::cerr << __FILE__ << ":" << __LINE__ << " " + << "Error: allocating memory for command buffer using NUMA" << std::endl; + abort(); + } + } else { + // Both the CPU and GPU can access the memory + status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn( + ag_list_count, ag_list, NULL, profile->command_buffer.ptr); + CHECK_HSA_STATUS("Error: Allowing access to Command Buffer", status); + } + if (!is_spm) { + status = HSA_STATUS_ERROR; + size_t size = (profile->output_buffer.size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK; + status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn( + agentInfo.kernarg_pool_, size, 0, &profile->output_buffer.ptr + ); + if (status != HSA_STATUS_SUCCESS) { + profile->output_buffer.ptr = malloc(size); + if (profile->output_buffer.ptr == NULL) { + std::cerr << __FILE__ << ":" << __LINE__ << " " + << "Error: allocating memory for output buffer using NUMA" << std::endl; + abort(); + } + } else { + status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn( + ag_list_count, ag_list, NULL, profile->output_buffer.ptr); + CHECK_HSA_STATUS("Error: GPU Agent can't have output buffer access", status); + hsa_amd_memory_fill(profile->output_buffer.ptr, 0, profile->output_buffer.size/sizeof(uint32_t)); + } + } else { + profile->output_buffer.size = 0; + } + + status = hsa_ven_amd_aqlprofile_start(profile, start_packet); + CHECK_HSA_STATUS("Error: Creating Start Packet\n", status); + status = hsa_ven_amd_aqlprofile_stop(profile, stop_packet); + CHECK_HSA_STATUS("Error: Creating Stop Packet\n", status); + status = hsa_ven_amd_aqlprofile_read(profile, read_packet); + CHECK_HSA_STATUS("Error: Creating Read Packet\n", status); + + context->start_packet = start_packet; + context->stop_packet = stop_packet; + context->read_packet = read_packet; + + return prof_context; } // Initialize the PM4 commands with having the CPU&GPU agents, the counters, @@ -665,10 +646,7 @@ void CreateBarrierPacket(std::vector* transformed_packets, const hsa_signal_t* packet_completion_signal ) { hsa_barrier_and_packet_t barrier{0}; - barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE | - (1 << HSA_PACKET_HEADER_BARRIER) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); + barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE | (1 << HSA_PACKET_HEADER_BARRIER); if (packet_completion_signal != nullptr) barrier.completion_signal = *packet_completion_signal; if (packet_dependency_signal != nullptr) barrier.dep_signal[0] = *packet_dependency_signal; void* barrier_ptr = &barrier; @@ -714,4 +692,84 @@ std::vector ExtractDispatchPackets( return ret; } +std::atomic AQLPacketProfile::valid_profiles{0}; +std::condition_variable_any AQLPacketProfile::delete_cv; +std::shared_mutex AQLPacketProfile::deleter_mutex; +bool AQLPacketProfile::IsDeletingBegin = false; + +std::unordered_map>> _cache; +std::mutex cache_mutex; + +void AQLPacketProfile::WaitForProfileDeletion() +{ + { + std::unique_lock lk(cache_mutex); + _cache.clear(); + } + + std::unique_lock lk(deleter_mutex); + + if (valid_profiles.load() == 0) + { + IsDeletingBegin = true; + return; + } + + delete_cv.wait_for(lk, std::chrono::seconds(2), [] () { + return ::Packet::AQLPacketProfile::valid_profiles.load() == 0; + }); + IsDeletingBegin = true; +} + +std::unique_ptr AQLPacketProfile::MoveFromCache(hsa_agent_t gpu_agent) +{ + std::lock_guard lk(cache_mutex); + + auto agent_it = _cache.find(gpu_agent.handle); + if (agent_it == _cache.end()) return nullptr; + + auto& profile_set = agent_it->second; + if (!profile_set.size()) return nullptr; + + auto moved = std::move(profile_set.back()); + profile_set.resize(profile_set.size()-1); + return moved; +} + +void AQLPacketProfile::MoveToCache(hsa_agent_t gpu_agent, std::unique_ptr&& packet) +{ + if (!packet.get()) return; + + auto& output_buffer = packet->profile->output_buffer; + + std::lock_guard lk(cache_mutex); + + auto agent_it = _cache.find(gpu_agent.handle); + if (agent_it == _cache.end()) + agent_it = _cache.emplace(gpu_agent.handle, std::vector>{}).first; + + auto& profile_set = agent_it->second; + profile_set.emplace_back(std::move(packet)); +} + +AQLPacketProfile::~AQLPacketProfile() +{ + std::shared_lock deleter_lock(deleter_mutex); + if (IsDeletingBegin) return; + + int old_valid = valid_profiles.fetch_sub(1); + + if (profile->output_buffer.ptr) + free_fn(profile->output_buffer.ptr); + profile->output_buffer.ptr = nullptr; + profile->output_buffer.size = 0; + if (profile->command_buffer.ptr) + free_fn(profile->command_buffer.ptr); + profile->command_buffer.ptr = nullptr; + profile->command_buffer.size = 0; + + if (old_valid <= 1) + delete_cv.notify_all(); +} + } // namespace Packet diff --git a/projects/rocprofiler/src/core/hsa/packets/packets_generator.h b/projects/rocprofiler/src/core/hsa/packets/packets_generator.h index 42914bdfd6..d68a4ea48c 100644 --- a/projects/rocprofiler/src/core/hsa/packets/packets_generator.h +++ b/projects/rocprofiler/src/core/hsa/packets/packets_generator.h @@ -32,6 +32,7 @@ #include #include #include +#include #include "src/core/counters/metrics/eval_metrics.h" @@ -39,10 +40,43 @@ namespace Packet { typedef hsa_ext_amd_aql_pm4_packet_t packet_t; -std::vector> -InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, - std::vector& counter_names, rocprofiler_session_id_t session_id, - bool is_spm = false); + +class AQLPacketProfile +{ +public: + AQLPacketProfile(decltype(hsa_amd_memory_pool_free) _free_fn) + { + profile = std::make_unique(); + context = std::make_unique(); + this->free_fn = _free_fn; + valid_profiles.fetch_add(1); + }; + + ~AQLPacketProfile(); + + std::unique_ptr profile; + std::unique_ptr context; + + static std::unique_ptr MoveFromCache(hsa_agent_t gpu_agent); + static void MoveToCache(hsa_agent_t gpu_agent, std::unique_ptr&& packet); + + static void WaitForProfileDeletion(); + + static std::atomic valid_profiles; + static std::condition_variable_any delete_cv; + static std::shared_mutex deleter_mutex; + static bool IsDeletingBegin; + decltype(hsa_amd_memory_pool_free)* free_fn; +}; + +std::unique_ptr InitializeAqlPackets( + hsa_agent_t cpu_agent, + hsa_agent_t gpu_agent, + std::vector& counter_names, + rocprofiler_session_id_t session_id, + bool is_spm = false +); + uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool); void GetCommandBufferMap(std::map); void GetOutputBufferMap(std::map); diff --git a/projects/rocprofiler/src/core/hsa/queues/queue.cpp b/projects/rocprofiler/src/core/hsa/queues/queue.cpp index 6ae2387817..9d6949413d 100644 --- a/projects/rocprofiler/src/core/hsa/queues/queue.cpp +++ b/projects/rocprofiler/src/core/hsa/queues/queue.cpp @@ -60,12 +60,6 @@ std::mutex sessions_pending_signal_lock; namespace rocprofiler { -// std::atomic ACTIVE_INTERRUPT_SIGNAL_COUNT{0}; - -// uint32_t GetCurrentActiveInterruptSignalsCount() { -// return ACTIVE_INTERRUPT_SIGNAL_COUNT.load(std::memory_order_relaxed); -// } - typedef std::vector pmc_callback_data_t; static inline bool IsEventMatch(const hsa_ven_amd_aqlprofile_event_t& event1, @@ -98,19 +92,16 @@ std::string GetKernelNameFromKsymbols(uint64_t handle) { void AddKernelNameWithDispatchID(std::string name, uint64_t id) { HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance(); - std::lock_guard lock(hsasupport_singleton.kernel_names_map_lock); - if (hsasupport_singleton.kernel_names->find(name) == hsasupport_singleton.kernel_names->end()) - hsasupport_singleton.kernel_names->emplace(name, std::vector()); - hsasupport_singleton.kernel_names->at(name).push_back(id); + std::unique_lock lock(hsasupport_singleton.kernel_names_map_lock); + (*hsasupport_singleton.kernel_names)[id] = name; } std::string GetKernelNameUsingDispatchID(uint64_t given_id) { HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance(); - std::lock_guard lock(hsasupport_singleton.kernel_names_map_lock); - for (auto kernel_name : (*hsasupport_singleton.kernel_names)) { - for (auto dispatch_id : kernel_name.second) { - if (dispatch_id == given_id) return kernel_name.first; - } - } + std::shared_lock lock(hsasupport_singleton.kernel_names_map_lock); + + auto it = hsasupport_singleton.kernel_names->find(given_id); + if (it != hsasupport_singleton.kernel_names->end()) + return it->second; return "Unknown Kernel!"; } @@ -273,21 +264,23 @@ hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type, return status; } -void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t* pending) { +void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t* pending) +{ + auto* context = pending->profile->context.get(); record->counters_count = - rocprofiler_record_counters_instances_count_t{pending->context->metrics_list.size()}; + rocprofiler_record_counters_instances_count_t{context->metrics_list.size()}; size_t counters_list_size = record->counters_count.value * sizeof(rocprofiler_record_counter_instance_t); rocprofiler_record_counter_instance_t* counters = static_cast(malloc(counters_list_size)); - for (size_t i = 0; i < pending->context->metrics_list.size(); i++) { - const rocprofiler::Metric* metric = pending->context->metrics_list[i]; + for (size_t i = 0; i < context->metrics_list.size(); i++) { + const rocprofiler::Metric* metric = context->metrics_list[i]; double value = 0; std::string metric_name = metric->GetName(); - auto it = pending->context->results_map.find(metric_name); - if (it != pending->context->results_map.end()) { + auto it = context->results_map.find(metric_name); + if (it != context->results_map.end()) value = it->second->val_double; - } + counters[i] = (rocprofiler_record_counter_instance_t{ // TODO(aelwazir): Moving to span once C++20 is adopted, strdup can be // removed after that @@ -309,6 +302,12 @@ void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_sign static_cast(data); }); } + + // Reset counters + for (auto& [key, value] : context->results_map) + value->val_double = 0; + for (auto* res : context->results_list) + res->val_double = 0; } /* @@ -384,14 +383,15 @@ void SignalAsyncReadyHandler(const hsa_signal_t& signal, void* data) { signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalReadyHandler, data); if (status != HSA_STATUS_SUCCESS) fatal("hsa_amd_signal_async_handler failed"); } -bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) { +bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) +{ auto queue_info_session = static_cast(data); if (!queue_info_session) return true; rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton = - rocprofiler::ROCProfiler_Singleton::GetInstance(); + rocprofiler::ROCProfiler_Singleton::GetInstance(); rocprofiler::HSASupport_Singleton& hsasupport_singleton = - rocprofiler::HSASupport_Singleton::GetInstance(); + rocprofiler::HSASupport_Singleton::GetInstance(); rocprofiler::Session* session = rocprofiler_singleton.GetSession(queue_info_session->session_id); if (!session) return true; @@ -402,112 +402,98 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) { auto pending_signals = profiler->MovePendingSignals(queue_info_session->writer_id); - for (auto& pending : pending_signals) { - if (hsasupport_singleton.GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->new_signal)) - return true; - hsa_amd_profiling_dispatch_time_t time; - hsasupport_singleton.GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn( - queue_info_session->agent, pending->new_signal, &time); - { - std::lock_guard lock(hsasupport_singleton.signals_timestamps_map_lock); - hsasupport_singleton.signals_timestamps[pending->original_signal.handle].time = - std::make_optional(time); - } - uint32_t record_count = 1; - bool is_individual_xcc_mode = false; - uint32_t xcc_count = queue_info_session->xcc_count; - if (xcc_count > 1 && pending->counters_count > 0) { // for MI300 - const char* str = getenv("ROCPROFILER_INDIVIDUAL_XCC_MODE"); - if (str != NULL) is_individual_xcc_mode = (atol(str) > 0); - // for individual xcc mode, there will be xcc_count records for each dispatch - // for accumulation mode, there will be only one record for a dispatch - if (is_individual_xcc_mode) record_count = xcc_count; - } - for (uint32_t xcc_id = 0; xcc_id < record_count; xcc_id++) { - rocprofiler_record_profiler_t record{}; - // TODO: (sauverma) gpu-id will need to support xcc like so- 1.1, 1.2, 1.3 ... 1.5 for - // different xcc - record.gpu_id = rocprofiler_agent_id_t{(uint64_t)queue_info_session->gpu_index}; - record.kernel_properties = pending->kernel_properties; - record.thread_id = rocprofiler_thread_id_t{pending->thread_id}; - record.queue_idx = rocprofiler_queue_index_t{pending->queue_index}; - record.timestamps = rocprofiler_record_header_timestamp_t{time.start, time.end}; - record.queue_id = rocprofiler_queue_id_t{queue_info_session->queue_id}; - record.xcc_index = xcc_id; - // Kernel Descriptor is the right record id generated in the WriteInterceptor function and - // will be used to handle the kernel name of that dispatch - record.header = rocprofiler_record_header_t{ - ROCPROFILER_PROFILER_RECORD, rocprofiler_record_id_t{pending->kernel_descriptor}}; - record.kernel_id = rocprofiler_kernel_id_t{pending->kernel_descriptor}; - record.correlation_id = rocprofiler_correlation_id_t{pending->correlation_id}; - - if (pending->session_id.handle == 0) { - pending->session_id = rocprofiler_singleton.GetCurrentSessionId(); + for (auto& pending : pending_signals) + { + if (hsasupport_singleton.GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->new_signal)) + return true; + hsa_amd_profiling_dispatch_time_t time; + hsasupport_singleton.GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn( + queue_info_session->agent, pending->new_signal, &time); + { + std::lock_guard lock(hsasupport_singleton.signals_timestamps_map_lock); + hsasupport_singleton.signals_timestamps[pending->original_signal.handle].time = + std::make_optional(time); } - if (pending->counters_count > 0) { - if (xcc_id == 0 && pending->context && pending->context->metrics_list.size() > 0 && - pending->profile) // call to GetCounterData() is required only once for a dispatch - rocprofiler::metrics::GetCounterData(pending->profile, queue_info_session->agent, - pending->context->results_list); - if (is_individual_xcc_mode) - rocprofiler::metrics::GetCountersAndMetricResultsByXcc( - xcc_id, pending->context->results_list, pending->context->results_map, - pending->context->metrics_list, time.end - time.start); - else - rocprofiler::metrics::GetMetricsData( - pending->context->results_map, pending->context->metrics_list, time.end - time.start); - AddRecordCounters(&record, pending.get()); - } else { - if (session->FindBuffer(pending->buffer_id)) { - Memory::GenericBuffer* buffer = session->GetBuffer(pending->buffer_id); - buffer->AddRecord(record); + //hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn(pending->new_signal); + uint32_t record_count = 1; + uint32_t xcc_count = queue_info_session->xcc_count; + static thread_local bool is_individual_xcc_mode = [xcc_count]() { + if (xcc_count < 2) return false; + const char* str = getenv("ROCPROFILER_INDIVIDUAL_XCC_MODE"); + if (str != NULL) return (atol(str) > 0); + return false; + }(); + if (is_individual_xcc_mode) record_count = xcc_count; + for (uint32_t xcc_id = 0; xcc_id < record_count; xcc_id++) { + rocprofiler_record_profiler_t record{}; + // TODO: (sauverma) gpu-id will need to support xcc like so- 1.1, 1.2, 1.3 ... 1.5 for + // different xcc + record.gpu_id = rocprofiler_agent_id_t{(uint64_t)queue_info_session->gpu_index}; + record.kernel_properties = pending->kernel_properties; + record.thread_id = rocprofiler_thread_id_t{pending->thread_id}; + record.queue_idx = rocprofiler_queue_index_t{pending->queue_index}; + record.timestamps = rocprofiler_record_header_timestamp_t{time.start, time.end}; + record.queue_id = rocprofiler_queue_id_t{queue_info_session->queue_id}; + record.xcc_index = xcc_id; + // Kernel Descriptor is the right record id generated in the WriteInterceptor function and + // will be used to handle the kernel name of that dispatch + record.header = rocprofiler_record_header_t{ + ROCPROFILER_PROFILER_RECORD, rocprofiler_record_id_t{pending->kernel_descriptor}}; + record.kernel_id = rocprofiler_kernel_id_t{pending->kernel_descriptor}; + record.correlation_id = rocprofiler_correlation_id_t{pending->correlation_id}; + + if (pending->session_id.handle == 0) { + pending->session_id = rocprofiler_singleton.GetCurrentSessionId(); + } + if (pending->counters_count > 0) + { + auto* context = pending->profile->context.get(); + auto* profile = pending->profile->profile.get(); + if (xcc_id == 0 && context && context->metrics_list.size() > 0 && profile) + rocprofiler::metrics::GetCounterData(profile, queue_info_session->agent, + context->results_list); + if (is_individual_xcc_mode) + rocprofiler::metrics::GetCountersAndMetricResultsByXcc( + xcc_id, context->results_list, context->results_map, + context->metrics_list, time.end - time.start); + else + rocprofiler::metrics::GetMetricsData(context->results_map, + context->metrics_list, + time.end - time.start); + AddRecordCounters(&record, pending.get()); + } else { + if (session->FindBuffer(pending->buffer_id)) { + Memory::GenericBuffer* buffer = session->GetBuffer(pending->buffer_id); + buffer->AddRecord(record); + } } } - } - if (pending->counters_count > 0 && pending->profile && pending->profile->events) { - // TODO(aelwazir): we need a better way of distributing events and free them - // if (pending->profile->output_buffer.ptr) - // numa_free(pending->profile->output_buffer.ptr, pending->profile->output_buffer.size); - hsa_status_t status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn( - (pending->profile->output_buffer.ptr)); - CHECK_HSA_STATUS("Error: Couldn't free output buffer memory", status); - // if (pending->profile->command_buffer.ptr) - // numa_free(pending->profile->command_buffer.ptr, pending->profile->command_buffer.size); - status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_free_fn( - (pending->profile->command_buffer.ptr)); - CHECK_HSA_STATUS("Error: Couldn't free command buffer memory", status); - delete pending->profile; - for (auto& it : pending->context->results_map) { - delete it.second; - } - delete pending->context; - /* - Check if the dispatch ready is empty, If so, there is no more - dispatches to be launched and we return. Else, dispatch the - kernel of the queue in the front of the dispatch_ready. - */ + auto* profile = pending->profile ? pending->profile->profile.get() : nullptr; + if (pending->counters_count > 0 && profile && profile->events) + { + Packet::AQLPacketProfile::MoveToCache(queue_info_session->agent, std::move(pending->profile)); - profiler_serializer_t& serializer = + profiler_serializer_t& serializer = rocprofiler::ROCProfiler_Singleton::GetInstance().GetSerializer(); - std::lock_guard serializer_lock(serializer.serializer_mutex); - assert(serializer.dispatch_queue != nullptr); - hsasupport_singleton.GetCoreApiTable().hsa_signal_store_screlease_fn( - queue_info_session->block_signal, 1); - serializer.dispatch_queue = nullptr; - if (serializer.dispatch_ready.empty()) return false; - Queue* queue = serializer.dispatch_ready.front(); - serializer.dispatch_ready.erase(serializer.dispatch_ready.begin()); - enable_dispatch(queue); - } + std::lock_guard serializer_lock(serializer.serializer_mutex); + assert(serializer.dispatch_queue != nullptr); + hsasupport_singleton.GetCoreApiTable().hsa_signal_store_screlease_fn( + queue_info_session->block_signal, 1); + serializer.dispatch_queue = nullptr; + if (serializer.dispatch_ready.empty()) return false; + Queue* queue = serializer.dispatch_ready.front(); + serializer.dispatch_ready.erase(serializer.dispatch_ready.begin()); + enable_dispatch(queue); - if (pending->new_signal.handle) - hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn(pending->new_signal); - if (queue_info_session->interrupt_signal.handle) - hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn( - queue_info_session->interrupt_signal); + } + + if (pending->new_signal.handle) + hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn(pending->new_signal); + if (queue_info_session->interrupt_signal.handle) + hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn( + queue_info_session->interrupt_signal); } delete queue_info_session; - // ACTIVE_INTERRUPT_SIGNAL_COUNT.fetch_sub(1, std::memory_order_relaxed); return false; } @@ -541,23 +527,27 @@ uint32_t replay_mode_count = 0; rocprofiler::Session* session = nullptr; -void Queue::ResetSessionID(rocprofiler_session_id_t id) { +void Queue::ResetSessionID(rocprofiler_session_id_t id) +{ std::unique_lock session_id_lock(session_id_mutex); session_id = id; + if (session_id.handle != 0) + session = rocprofiler::ROCProfiler_Singleton::GetInstance().GetSession(session_id); + else + session = nullptr; } -void Queue::CheckNeededProfileConfigs() { - rocprofiler_session_id_t internal_session_id; +bool Queue::CheckNeededProfileConfigs() +{ + std::unique_lock session_id_lock(session_id_mutex); + // Getting Session ID rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton = rocprofiler::ROCProfiler_Singleton::GetInstance(); - internal_session_id = rocprofiler_singleton.GetCurrentSessionId(); - - if (session_id.handle > 0 && internal_session_id.handle == session_id.handle) return; - if (internal_session_id.handle == 0) return; - session_id = internal_session_id; + session_id = rocprofiler_singleton.GetCurrentSessionId(); // Getting Counters count from the Session + if (session_id.handle == 0) return false; session = rocprofiler_singleton.GetSession(session_id); if (session && session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION)) { @@ -568,7 +558,8 @@ void Queue::CheckNeededProfileConfigs() { is_counter_collection_mode = true; session_data_count = session_data.size(); buffer_id = filter->GetBufferId(); - } else if (session && session->FindFilterWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION)) { + } else if (session && + session->FindFilterWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION)) { is_timestamp_collection_mode = true; rocprofiler_filter_id_t filter_id = session->GetFilterIdWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION); @@ -583,15 +574,18 @@ void Queue::CheckNeededProfileConfigs() { att_tracer->SetParameters(filter->GetAttParametersData()); is_att_collection_mode = true; buffer_id = session->GetFilter(session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) - ->GetBufferId(); + ->GetBufferId(); att_tracer->SetCountersNames(filter->GetCounterData()); - att_tracer->SetKernelsNames( - std::get>(filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES))); - att_tracer->SetDispatchIds(std::get>>( - filter->GetProperty(ROCPROFILER_FILTER_DISPATCH_IDS))); + att_tracer->SetKernelsNames(std::get>( + filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES) + )); + att_tracer->SetDispatchIds(std::get>>( + filter->GetProperty(ROCPROFILER_FILTER_DISPATCH_IDS) + )); } else if (session && session->FindFilterWithKind(ROCPROFILER_PC_SAMPLING_COLLECTION)) { is_pc_sampling_collection_mode = true; } + return true; } std::atomic WRITER_ID{0}; @@ -604,79 +598,85 @@ std::atomic WRITER_ID{0}; * interceptor by invoking the writer function. */ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index, - void* data, hsa_amd_queue_intercept_packet_writer writer) { - std::shared_lock session_id_lock(session_id_mutex); + void* data, hsa_amd_queue_intercept_packet_writer writer) +{ const Packet::packet_t* packets_arr = reinterpret_cast(packets); std::vector transformed_packets; - CheckNeededProfileConfigs(); - rocprofiler_session_id_t session_id_snapshot = session_id; + std::shared_lock session_id_lock(session_id_mutex); + + if (session_id.handle == 0 || session_id.handle != rocprofiler::ROCProfiler_Singleton::GetInstance().GetCurrentSessionId().handle) + { + session_id_lock.unlock(); + CheckNeededProfileConfigs(); + session_id_lock.lock(); + } auto& queue_info = *reinterpret_cast(data); std::lock_guard lk(queue_info.qw_mutex); - if (session_id_snapshot.handle > 0 && pkt_count > 0 && + if (session_id.handle > 0 && pkt_count > 0 && (is_counter_collection_mode || is_timestamp_collection_mode || is_pc_sampling_collection_mode) && session) { + // hsa_ven_amd_aqlprofile_profile_t* profile; - std::vector> - profiles; + // Searching accross all the packets given during this write for (size_t i = 0; i < pkt_count; ++i) { auto& original_packet = static_cast(packets)[i]; // +Skip kernel dispatch IDs not wanted // Skip packets other than kernel dispatch packets. - if (session_id_snapshot.handle == 0 || !Packet::IsDispatchPacket(original_packet)) { + if (session_id.handle == 0 || !Packet::IsDispatchPacket(original_packet)) { transformed_packets.emplace_back(packets_arr[i]); continue; } + std::unique_ptr profile_packet; // If counters found in the session if (session_data_count > 0 && is_counter_collection_mode) { // Get the PM4 Packets using packets_generator - profiles = Packet::InitializeAqlPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(), - session_data, session_id_snapshot); - replay_mode_count = profiles.size(); + profile_packet = Packet::AQLPacketProfile::MoveFromCache(queue_info.GetGPUAgent()); + if (!profile_packet) + profile_packet = Packet::InitializeAqlPackets( + queue_info.GetCPUAgent(), queue_info.GetGPUAgent(), session_data, session_id); } - uint32_t profile_id = 0; - std::pair profile; - if (session_data_count > 0 && is_counter_collection_mode) { - if (profiles.size() > 0 && replay_mode_count > 0) { - profile = profiles.at(profile_id); - hsa_signal_t ready_signal = queue_info.GetReadySignal(); - hsa_signal_t block_signal = queue_info.GetBlockSignal(); - /* - Creates a barrier packet with its completion signal as the - queue's ready signal. - */ - Packet::CreateBarrierPacket(&transformed_packets, nullptr, &ready_signal); - /* - Creates a barrier packet with queue's blocksignal as its input and - completion signal.This will ensure it is no longer 0 so a later barrier - packet waiting on it to be 0 will be blocked - */ - Packet::CreateBarrierPacket(&transformed_packets, &block_signal, &block_signal); - } + if (profile_packet.get()) + { + hsa_signal_t ready_signal = queue_info.GetReadySignal(); + hsa_signal_t block_signal = queue_info.GetBlockSignal(); + + /* + Creates a barrier packet with its completion signal as the + queue's ready signal. + */ + Packet::CreateBarrierPacket(&transformed_packets, nullptr, &ready_signal); + /* + Creates a barrier packet with queue's blocksignal as its input and + completion signal.This will ensure it is no longer 0 so a later barrier + packet waiting on it to be 0 will be blocked + */ + Packet::CreateBarrierPacket(&transformed_packets, &block_signal, &block_signal); } uint32_t writer_id = WRITER_ID.fetch_add(1, std::memory_order_release); - - if (session_data_count > 0 && is_counter_collection_mode && profiles.size() > 0 && - replay_mode_count > 0 && profile.first && profile.first->start_packet) { + if (session_data_count > 0 && is_counter_collection_mode && profile_packet.get()) + { + auto* start_packet = profile_packet->context->start_packet; // Adding start packet and its barrier with a dummy signal hsa_signal_t dummy_signal{}; dummy_signal.handle = 0; - profile.first->start_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC - << HSA_PACKET_HEADER_TYPE; - Packet::AddVendorSpecificPacket(profile.first->start_packet, &transformed_packets, - dummy_signal); + start_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + Packet::AddVendorSpecificPacket(start_packet, &transformed_packets, dummy_signal); - Packet::CreateBarrierPacket(&transformed_packets, - &profile.first->start_packet->completion_signal, nullptr); + Packet::CreateBarrierPacket( + &transformed_packets, + &start_packet->completion_signal, + nullptr + ); } auto& packet = transformed_packets.emplace_back(packets_arr[i]); @@ -689,19 +689,24 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u // list to be processed by the signal interrupt rocprofiler_kernel_properties_t kernel_properties = set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent()); + + auto* context_backup = profile_packet.get() ? profile_packet->context.get() : nullptr; if (session) { uint64_t record_id = rocprofiler::ROCProfiler_Singleton::GetInstance().GetUniqueRecordId(); AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object), record_id); - if (session_data_count > 0 && profile.second) { + if (session_data_count > 0 && profile_packet.get()) + { session->GetProfiler()->AddPendingSignals( writer_id, record_id, original_packet.completion_signal, packet.completion_signal, - session_id, buffer_id, profile.first, session_data_count, profile.second, + session_id, buffer_id, session_data_count, std::move(profile_packet), kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index, correlation_id); - } else { + } + else + { session->GetProfiler()->AddPendingSignals( writer_id, record_id, original_packet.completion_signal, packet.completion_signal, - session_id, buffer_id, nullptr, session_data_count, nullptr, kernel_properties, + session_id, buffer_id, session_data_count, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index, correlation_id); } } @@ -710,7 +715,7 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u // packet and create a new signal for it to get timestamps if (original_packet.completion_signal.handle) { hsa_barrier_and_packet_t barrier{}; - barrier.header = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE); + barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; Packet::packet_t* __attribute__((__may_alias__)) pkt = (reinterpret_cast(&barrier)); transformed_packets.emplace_back(*pkt).completion_signal = @@ -730,56 +735,47 @@ void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t u CreateSignal(0, &interrupt_signal); // Adding Stop and Read PM4 Packets - if (session_data_count > 0 && is_counter_collection_mode && profiles.size() > 0 && - profile.first && profile.first->read_packet) { + if (session_data_count > 0 && is_counter_collection_mode && context_backup) + { hsa_signal_t dummy_signal{}; - profile.first->read_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC - << HSA_PACKET_HEADER_TYPE; - Packet::AddVendorSpecificPacket(profile.first->read_packet, &transformed_packets, - dummy_signal); - profile.first->stop_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC - << HSA_PACKET_HEADER_TYPE; - Packet::AddVendorSpecificPacket(profile.first->stop_packet, &transformed_packets, - interrupt_signal); + context_backup->read_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + Packet::AddVendorSpecificPacket(context_backup->read_packet, &transformed_packets, dummy_signal); + context_backup->stop_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + Packet::AddVendorSpecificPacket(context_backup->stop_packet, &transformed_packets, interrupt_signal); // Added Interrupt Signal with barrier and provided handler for it - Packet::CreateBarrierPacket(&transformed_packets, &interrupt_signal, nullptr); - rocprofiler::HSAAgentInfo& agentInfo = - rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo( - queue_info.GetGPUAgent().handle); - // Creating Async Handler to be called every time the interrupt signal is - // marked complete - SignalAsyncHandler( - interrupt_signal, - new queue_info_session_t{ - queue_info.GetGPUAgent(), session_id_snapshot, queue_info.GetQueueID(), writer_id, - interrupt_signal, agentInfo.GetDeviceInfo().getNumaNode(), - agentInfo.GetDeviceInfo().getXccCount(), queue_info.GetBlockSignal()}); - - } else { - rocprofiler::HSAAgentInfo& agentInfo = - rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo( - queue_info.GetGPUAgent().handle); - Packet::CreateBarrierPacket(&transformed_packets, nullptr, &interrupt_signal); - // Creating Async Handler to be called every time the interrupt signal is - // marked complete - SignalAsyncHandler( - interrupt_signal, - new queue_info_session_t{ - queue_info.GetGPUAgent(), session_id_snapshot, queue_info.GetQueueID(), writer_id, - interrupt_signal, agentInfo.GetDeviceInfo().getNumaNode(), - agentInfo.GetDeviceInfo().getXccCount(), queue_info.GetBlockSignal()}); + Packet::CreateBarrierPacket( &transformed_packets, &interrupt_signal, nullptr); } + else + Packet::CreateBarrierPacket(&transformed_packets, nullptr, &interrupt_signal); - - // ACTIVE_INTERRUPT_SIGNAL_COUNT.fetch_add(1, std::memory_order_relaxed); + rocprofiler::HSAAgentInfo& agentInfo = + rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo( + queue_info.GetGPUAgent().handle); + // Creating Async Handler to be called every time the interrupt signal is + // marked complete + SignalAsyncHandler( + interrupt_signal, + new queue_info_session_t{ + queue_info.GetGPUAgent(), session_id, queue_info.GetQueueID(), writer_id, + interrupt_signal, agentInfo.GetDeviceInfo().getNumaNode(), + agentInfo.GetDeviceInfo().getXccCount(), queue_info.GetBlockSignal()}); } /* Write the transformed packets to the hardware queue. */ writer(&transformed_packets[0], transformed_packets.size()); - } else if (!is_att_collection_mode || - !session->GetAttTracer()->ATTWriteInterceptor(packets, pkt_count, user_pkt_index, - *static_cast(data), writer, - buffer_id)) { + } else if ( + !is_att_collection_mode|| + !session || + !session->GetAttTracer()|| + !session->GetAttTracer()->ATTWriteInterceptor( + packets, + pkt_count, + user_pkt_index, + *static_cast(data), + writer, + buffer_id + ) + ) { /* Write the original packets to the hardware queue if no profiling session is active */ writer(packets, pkt_count); } @@ -830,11 +826,16 @@ hsa_agent_t Queue::GetCPUAgent() { return cpu_agent_; } uint64_t Queue::GetQueueID() { return intercept_queue_->id; } -void CheckPacketReqiurements() { Packet::CheckPacketReqiurements(); } +void CheckPacketReqiurements() { + Packet::CheckPacketReqiurements(); +} hsa_signal_t Queue::GetReadySignal() { return ready_signal_; } hsa_signal_t Queue::GetBlockSignal() { return block_signal_; } + + + } // namespace queue } // namespace rocprofiler diff --git a/projects/rocprofiler/src/core/hsa/queues/queue.h b/projects/rocprofiler/src/core/hsa/queues/queue.h index 3a8b72bf15..38aefc06c0 100644 --- a/projects/rocprofiler/src/core/hsa/queues/queue.h +++ b/projects/rocprofiler/src/core/hsa/queues/queue.h @@ -89,7 +89,7 @@ class Queue { hsa_signal_t GetBlockSignal(); static void ResetSessionID(rocprofiler_session_id_t id = rocprofiler_session_id_t{0}); - static void CheckNeededProfileConfigs(); + static bool CheckNeededProfileConfigs(); private: static std::shared_mutex session_id_mutex; static rocprofiler_session_id_t session_id; diff --git a/projects/rocprofiler/src/core/session/profiler/profiler.cpp b/projects/rocprofiler/src/core/session/profiler/profiler.cpp index 2fa0f838c5..ffb589f127 100644 --- a/projects/rocprofiler/src/core/session/profiler/profiler.cpp +++ b/projects/rocprofiler/src/core/session/profiler/profiler.cpp @@ -119,8 +119,8 @@ bool Profiler::HasActivePass() { void Profiler::AddPendingSignals( uint32_t writer_id, uint64_t kernel_object, const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id, - rocprofiler_buffer_id_t buffer_id, rocprofiler::profiling_context_t* context, - uint64_t session_data_count, hsa_ven_amd_aqlprofile_profile_t* profile, + rocprofiler_buffer_id_t buffer_id, + uint64_t session_data_count, std::unique_ptr&& profile, rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id, uint64_t queue_index, uint64_t correlation_id) { @@ -134,7 +134,7 @@ void Profiler::AddPendingSignals( sessions_pending_signals_.at(writer_id).emplace_back( new pending_signal_t{ kernel_object, original_completion_signal, new_completion_signal, - session_id_, buffer_id, context, session_data_count, profile, + session_id_, buffer_id, session_data_count, std::move(profile), kernel_properties, thread_id, queue_index, correlation_id } ); diff --git a/projects/rocprofiler/src/core/session/profiler/profiler.h b/projects/rocprofiler/src/core/session/profiler/profiler.h index aa5c389296..1416e20086 100644 --- a/projects/rocprofiler/src/core/session/profiler/profiler.h +++ b/projects/rocprofiler/src/core/session/profiler/profiler.h @@ -30,7 +30,7 @@ #include #include #include - +#include "src/core/hsa/packets/packets_generator.h" #include "rocprofiler.h" #include "src/core/counters/basic/basic_counter.h" #include "src/core/counters/metrics/eval_metrics.h" @@ -48,9 +48,8 @@ typedef struct { hsa_signal_t new_signal; rocprofiler_session_id_t session_id; rocprofiler_buffer_id_t buffer_id; - rocprofiler::profiling_context_t* context; uint64_t counters_count; - hsa_ven_amd_aqlprofile_profile_t* profile; + std::unique_ptr<::Packet::AQLPacketProfile> profile; rocprofiler_kernel_properties_t kernel_properties; uint32_t thread_id; uint64_t queue_index; @@ -73,8 +72,8 @@ class Profiler { const hsa_signal_t& original_completion_signal, const hsa_signal_t& new_completion_signal, rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id, - rocprofiler::profiling_context_t* context, uint64_t session_data_count, - hsa_ven_amd_aqlprofile_profile_t* profile, + uint64_t session_data_count, + std::unique_ptr&& profile, rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id, uint64_t queue_index, uint64_t correlation_id); diff --git a/projects/rocprofiler/src/core/session/session.cpp b/projects/rocprofiler/src/core/session/session.cpp index 4fea3b4489..d72b0e59ab 100644 --- a/projects/rocprofiler/src/core/session/session.cpp +++ b/projects/rocprofiler/src/core/session/session.cpp @@ -216,7 +216,7 @@ void Session::Terminate() GetProfiler()->WaitForPendingAndDestroy(); if (GetAttTracer()) GetAttTracer()->WaitForPendingAndDestroy(); - + Packet::AQLPacketProfile::WaitForProfileDeletion(); std::lock_guard lock(session_lock_); if (FindFilterWithKind(ROCPROFILER_SPM_COLLECTION)) {