AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-403050: Multiple fixes for Memory Leaks in profiler

Browse Source 
Change-Id: Ib720a81105af13898ff745ce0cbc2a48c1f4a980
This commit is contained in:
Ammar ELWazir
2023-05-27 09:34:42 +00:00
committed by Ammar ELWazir
parent 00ecca25c7
commit 08fc21ac31
26 changed files with 673 additions and 592 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/core/hsa/hsa_common.cpp
+20 -7
View File
@@ -30,26 +30,39 @@ std::mutex agents_map_lock;
std::map<decltype(hsa_agent_t::handle), Agent::AgentInfo> agent_info_map;
Agent::AgentInfo& GetAgentInfo(decltype(hsa_agent_t::handle) handle) {
std::lock_guard<std::mutex> lock(agents_map_lock);
if (agent_info_map.find(handle) != agent_info_map.end())
if (agent_info_map.find(handle) != agent_info_map.end()) {
return agent_info_map.at(handle);
else
throw(std::string("Error: Can't find Agent with handle(") + std::to_string(handle) +
") in this system");
} else {
std::cerr << std::string("Error: Can't find Agent with handle(") << std::to_string(handle) <<
") in this system" << std::endl;
abort();
}
}
std::vector<hsa_agent_t> cpu_agents_list;
void SetAgentInfo(decltype(hsa_agent_t::handle) handle, const Agent::AgentInfo& agent_info) {
std::lock_guard<std::mutex> lock(agents_map_lock);
agent_info_map.emplace(handle, agent_info);
if (agent_info.getType() == HSA_DEVICE_TYPE_GPU) {
cpu_agents_list.emplace_back(hsa_agent_t{handle});
}
}
hsa_agent_t GetAgentByIndex(int agent_index) {
std::vector<hsa_agent_t>& GetCPUAgentList() {
return cpu_agents_list;
}
hsa_agent_t GetAgentByIndex(uint64_t agent_index) {
std::lock_guard<std::mutex> lock(agents_map_lock);
for (auto& agent_info : agent_info_map) {
if (agent_info.second.getIndex() == agent_index) {
return hsa_agent_t{agent_info.second.getHandle()};
}
}
throw(std::string("Error: Can't find Agent with Index(") + std::to_string(agent_index) +
") in this system");
std::cerr << std::string("Error: Can't find Agent with Index(") << std::to_string(agent_index) <<
") in this system" << std::endl;
abort();
}
CoreApiTable saved_core_api{};
src/core/hsa/hsa_common.h
+4 -1
View File
@@ -38,9 +38,12 @@
namespace rocmtools {
namespace hsa_support {
std::vector<hsa_agent_t>& GetCPUAgentList();
Agent::AgentInfo& GetAgentInfo(decltype(hsa_agent_t::handle) handle);
void SetAgentInfo(decltype(hsa_agent_t::handle) handle, const Agent::AgentInfo& agent_info);
hsa_agent_t GetAgentByIndex(int agent_index);
hsa_agent_t GetAgentByIndex(uint64_t agent_index);
CoreApiTable& GetCoreApiTable();
void SetCoreApiTable(const CoreApiTable& table);
src/core/hsa/packets/packets_generator.cpp
+263 -280
View File
@@ -25,8 +25,11 @@
#include <hsa/hsa_ven_amd_aqlprofile.h>
#include <stddef.h>
#include <stdint.h>
#include <numa.h>
#include <algorithm>
#include <atomic>
#include <exception>
#include <iostream>
#include <map>
#include <string>
@@ -47,23 +50,22 @@
#define CHECK_HSA_STATUS(msg, status) \
do { \
if ((status) != HSA_STATUS_INFO_BREAK) { \
const char* emsg = 0; \
hsa_status_string(status, &emsg); \
printf("%s: %s\n", msg, emsg ? emsg : "<unknown error>"); \
if ((status) != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) { \
try { \
const char* emsg = nullptr; \
hsa_status_string(status, &emsg); \
if (!emsg) emsg = "<Unknown HSA Error>"; \
std::cerr << msg << std::endl; \
std::cerr << emsg << std::endl; \
} catch (std::exception & e) { \
} \
abort(); \
} \
} while (0)
namespace Packet {
static const size_t MEM_PAGE_BYTES = 0x1000;
static const size_t MEM_PAGE_MASK = MEM_PAGE_BYTES - 1;
// hsa_amd_memory_pool_t command_pool;
// hsa_amd_memory_pool_t output_pool;
// hsa_amd_memory_pool_t& GetCommandPool() { return command_pool; }
// hsa_amd_memory_pool_t& GetOutputPool() { return output_pool; }
// This function checks to see if the provided
// pool has the HSA_AMD_SEGMENT_GLOBAL property. If the kern_arg flag is true,
@@ -116,11 +118,11 @@ void InitializePools(hsa_agent_t cpu_agent, Agent::AgentInfo* agent_info) {
hsa_status_t status =
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
cpu_agent, FindStandardPool, &(agent_info->cpu_pool));
if ((status != HSA_STATUS_INFO_BREAK)) printf("Error: Command Buffer Pool is not initialized\n");
CHECK_HSA_STATUS("Error: Command Buffer Pool is not initialized", status);
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
cpu_agent, FindKernArgPool, &(agent_info->kernarg_pool));
if ((status != HSA_STATUS_INFO_BREAK)) printf("Error: Output Buffer Pool is not initialized\n");
CHECK_HSA_STATUS("Error: Output Buffer Pool is not initialized", status);
}
void InitializeGPUPool(hsa_agent_t gpu_agent, Agent::AgentInfo* agent_info) {
@@ -135,6 +137,7 @@ struct block_des_t {
};
std::map<uint32_t, rocmtools::MetricsDict*> metricsDict;
static std::atomic<bool> counters_added{false};
void CheckPacketReqiurements(std::vector<hsa_agent_t>& gpu_agents) {
for (auto& gpu_agent : gpu_agents) {
@@ -147,253 +150,240 @@ void CheckPacketReqiurements(std::vector<hsa_agent_t>& gpu_agents) {
// 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<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>*
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
std::vector<std::string>& counter_names, bool is_spm) {
hsa_status_t status = HSA_STATUS_SUCCESS;
if (!counters_added.load(std::memory_order_acquire)) {
for (auto& name : counter_names) {
if (rocmtools::GetROCMToolObj()->HasActiveSession()) {
rocmtools::GetROCMToolObj()
->GetSession(rocmtools::GetROCMToolObj()->GetCurrentSessionId())
->GetProfiler()
->AddCounterName(name);
}
}
counters_added.exchange(true, std::memory_order_release);
}
Agent::AgentInfo& agentInfo = rocmtools::hsa_support::GetAgentInfo(gpu_agent.handle);
std::map<std::string, rocmtools::results_t*> results_map;
std::vector<rocmtools::event_t> events_list;
std::vector<rocmtools::results_t*> results_list;
std::map<std::pair<uint32_t, uint32_t>, uint64_t> event_to_max_block_count;
std::map<std::string, std::set<std::string>> metrics_counters;
uint32_t counters_count = 0;
for (auto& name : counter_names) {
// std::cout << "Counter from Counter Names: " << name << std::endl;
if (rocmtools::GetROCMToolObj()->HasActiveSession()) {
rocmtools::GetROCMToolObj()
->GetSession(rocmtools::GetROCMToolObj()->GetCurrentSessionId())
->GetProfiler()
->AddCounterName(name);
}
counters_count++;
if (!rocmtools::metrics::ExtractMetricEvents(
counter_names, gpu_agent, metricsDict[gpu_agent.handle], results_map, events_list,
results_list, event_to_max_block_count, metrics_counters)) {
std::cerr << "Error: Failed to extract metric events" << std::endl;
abort();
}
rocmtools::metrics::ExtractMetricEvents(counter_names, gpu_agent, metricsDict[gpu_agent.handle],
results_map, events_list, results_list,
event_to_max_block_count, metrics_counters);
// TODO: validate needs to be called on each events_list[i]
// Validating the events array for the specified gpu agent
bool result;
hsa_ven_amd_aqlprofile_validate_event(gpu_agent, &events_list[0], &result);
if (!result) {
printf("Error: Events are not valid for the current gpu agent\n");
throw("Error: Events are not valid for the current gpu agent");
bool validate_event_result;
status =
hsa_ven_amd_aqlprofile_validate_event(gpu_agent, &events_list[0], &validate_event_result);
CHECK_HSA_STATUS("Error: Validating Counters", status);
if (!validate_event_result) {
std::cerr << "Error: Events are not valid for the current gpu agent" << std::endl;
abort();
}
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>*
profiles = new std::vector<
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
profiles = std::vector<
std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>();
do {
rocmtools::profiling_context_t* context = new rocmtools::profiling_context_t();
context->gpu_agent = gpu_agent;
uint64_t i = 0;
uint32_t counter_val_iteration = 0;
auto result = results_list.begin();
std::map<std::pair<uint32_t, uint32_t>, uint32_t> block_max_events_count;
std::set<hsa_ven_amd_aqlprofile_block_name_t> block_names_taken;
for (auto event = events_list.begin(); event != events_list.end();) {
if (block_max_events_count[std::make_pair<uint32_t, uint32_t>(
static_cast<uint32_t>(event->block_name),
static_cast<uint32_t>(event->block_index))] <
event_to_max_block_count[std::make_pair<uint32_t, uint32_t>(
static_cast<uint32_t>(event->block_name),
static_cast<uint32_t>(event->block_index))]) {
context->events_list.push_back(*event);
context->results_list.emplace_back(*result);
counter_val_iteration++;
block_max_events_count[std::make_pair<uint32_t, uint32_t>(
static_cast<uint32_t>(event->block_name), static_cast<uint32_t>(event->block_index))]++;
results_list.erase(result);
events_list.erase(event);
} else {
event++;
result++;
}
i++;
}
std::set<std::string> counters_taken;
std::set<std::string> metrics_counters_taken;
for (auto result : context->results_list) {
rocmtools::Metric* metric;
if (std::find(counter_names.begin(), counter_names.end(), result->name) !=
counter_names.end()) {
// std::cout << "Counter from Result List: " << result->name << std::endl;
counters_taken.insert(result->name);
metric = const_cast<rocmtools::Metric*>(metricsDict[gpu_agent.handle]->Get(result->name));
if (metric == nullptr) std::cout << result->name << " not found in metricsDict\n";
context->metrics_list.push_back(metric);
} else {
metrics_counters_taken.insert(result->name);
// std::cout << "Counter Added: " << result->name << std::endl;
}
}
std::set<std::string> metrics_taken;
for (auto result : results_map) {
if (counters_taken.find(result.first) == counters_taken.end() &&
std::find(counter_names.begin(), counter_names.end(), result.first) !=
counter_names.end()) {
bool flag = true;
for (auto result_basic : results_list) {
if (result_basic->name.compare(result.first)) {
flag = false;
break;
}
}
if (flag) metrics_taken.insert(result.first);
}
}
for (auto metric_name : metrics_taken) {
bool flag = true;
if (metrics_counters.find(metric_name) == metrics_counters.end()) continue;
for (auto metric_counter_name : metrics_counters.at(metric_name)) {
if (metrics_counters_taken.find(metric_counter_name) == metrics_counters_taken.end() &&
counters_taken.find(metric_counter_name) == counters_taken.end()) {
flag = false;
continue;
}
}
if (flag) {
// std::cout << "Counter from Result Map: " << metric_name << std::endl;
counters_taken.insert(metric_name);
rocmtools::Metric* metric =
const_cast<rocmtools::Metric*>(metricsDict[gpu_agent.handle]->Get(metric_name));
if (metric == nullptr) std::cout << metric_name << " not found in metricsDict\n";
context->metrics_list.push_back(metric);
}
}
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();
if (context->events_list.size() <= 0) {
continue;
}
// Preparing the profile structure to get the packets
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wconversion-null"
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<uint32_t>(context->events_list.size()),
params,
0,
NULL,
NULL};
#pragma GCC diagnostic pop
// Preparing an Getting the size of the command and output buffers
status = hsa_ven_amd_aqlprofile_start(profile, NULL);
if (status != HSA_STATUS_SUCCESS) {
const char* hsa_err_str = nullptr;
if (hsa_status_string(status, &hsa_err_str) != HSA_STATUS_SUCCESS) hsa_err_str = "Unknown";
printf("Error: %s\n", hsa_err_str);
continue;
// do {
rocmtools::profiling_context_t* context = new rocmtools::profiling_context_t();
context->gpu_agent = gpu_agent;
auto result = results_list.begin();
std::map<std::pair<uint32_t, uint32_t>, uint32_t> block_max_events_count;
std::set<hsa_ven_amd_aqlprofile_block_name_t> block_names_taken;
for (auto event = events_list.begin(); event != events_list.end();) {
if (block_max_events_count[std::make_pair<uint32_t, uint32_t>(
static_cast<uint32_t>(event->block_name), static_cast<uint32_t>(event->block_index))] <
event_to_max_block_count[std::make_pair<uint32_t, uint32_t>(
static_cast<uint32_t>(event->block_name), static_cast<uint32_t>(event->block_index))]) {
context->events_list.push_back(*event);
context->results_list.emplace_back(*result);
block_max_events_count[std::make_pair<uint32_t, uint32_t>(
static_cast<uint32_t>(event->block_name), static_cast<uint32_t>(event->block_index))]++;
results_list.erase(result);
events_list.erase(event);
} else {
status = HSA_STATUS_ERROR;
size_t size = profile->command_buffer.size;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.cpu_pool, size, 0, reinterpret_cast<void**>(&(profile->command_buffer.ptr)));
event++;
result++;
}
}
// Both the CPU and GPU can access the memory
if (status == HSA_STATUS_SUCCESS) {
hsa_agent_t ag_list[1] = {gpu_agent};
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
1, ag_list, NULL, profile->command_buffer.ptr);
std::set<std::string> counters_taken;
if (status != HSA_STATUS_SUCCESS) {
printf("Error: Can't allow access for both agents to Command Buffer\n");
continue;
} else if (status == HSA_STATUS_ERROR_OUT_OF_RESOURCES) {
printf("Error: Ran out of GPU memory to allocate Command Buffer\n");
continue;
std::set<std::string> metrics_counters_taken;
for (auto result : context->results_list) {
rocmtools::Metric* metric;
if (std::find(counter_names.begin(), counter_names.end(), result->name) !=
counter_names.end()) {
// std::cout << "Counter from Result List: " << result->name << std::endl;
counters_taken.insert(result->name);
metric = const_cast<rocmtools::Metric*>(metricsDict[gpu_agent.handle]->Get(result->name));
if (metric == nullptr) std::cout << result->name << " not found in metricsDict\n";
context->metrics_list.push_back(metric);
} else {
metrics_counters_taken.insert(result->name);
// std::cout << "Counter Added: " << result->name << std::endl;
}
}
std::set<std::string> metrics_taken;
for (auto result : results_map) {
if (counters_taken.find(result.first) == counters_taken.end() &&
std::find(counter_names.begin(), counter_names.end(), result.first) !=
counter_names.end()) {
bool flag = true;
for (auto result_basic : results_list) {
if (result_basic->name.compare(result.first)) {
flag = false;
break;
}
} else {
const char* hsa_err_str = NULL;
if (hsa_status_string(status, &hsa_err_str) != HSA_STATUS_SUCCESS) hsa_err_str = "Unknown";
printf("Error: Allocating command Buffer (Size=%lu) (%s)\n", size, hsa_err_str);
}
if (flag) metrics_taken.insert(result.first);
}
}
if (!is_spm) {
status = HSA_STATUS_ERROR;
size = profile->output_buffer.size;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.kernarg_pool, size, 0, reinterpret_cast<void**>(&profile->output_buffer.ptr));
if (status == HSA_STATUS_ERROR_OUT_OF_RESOURCES) {
printf("Error: Ran out of GPU memory to allocate Output Buffer\n");
continue;
}
if (status == HSA_STATUS_SUCCESS) {
hsa_agent_t ag_list[1] = {gpu_agent};
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
1, ag_list, NULL, profile->output_buffer.ptr);
if (status == HSA_STATUS_SUCCESS) {
memset(profile->output_buffer.ptr, 0x0, profile->output_buffer.size);
// Creating the start/stop/read packets
status = hsa_ven_amd_aqlprofile_start(profile, start_packet);
status = hsa_ven_amd_aqlprofile_stop(profile, stop_packet);
status = hsa_ven_amd_aqlprofile_read(profile, read_packet);
context->start_packet = start_packet;
context->stop_packet = stop_packet;
context->read_packet = read_packet;
// add profiles
profiles->emplace_back(std::make_pair(context, profile));
} else {
printf("Error: Can't allow access for both agents to output Buffer\n");
continue;
}
} else {
const char* hsa_err_str = NULL;
if (hsa_status_string(status, &hsa_err_str) != HSA_STATUS_SUCCESS)
hsa_err_str = "Unknown";
printf("Error: Allocating output Buffer (%s)\n", hsa_err_str);
continue;
}
} else {
profile->output_buffer.size = 0;
status = hsa_ven_amd_aqlprofile_start(profile, start_packet);
status = hsa_ven_amd_aqlprofile_stop(profile, stop_packet);
status = hsa_ven_amd_aqlprofile_read(profile, read_packet);
context->start_packet = start_packet;
context->stop_packet = stop_packet;
context->read_packet = read_packet;
// add profiles
profiles->emplace_back(std::make_pair(context, profile));
for (auto metric_name : metrics_taken) {
bool flag = true;
if (metrics_counters.find(metric_name) == metrics_counters.end()) continue;
for (auto metric_counter_name : metrics_counters.at(metric_name)) {
if (metrics_counters_taken.find(metric_counter_name) == metrics_counters_taken.end() &&
counters_taken.find(metric_counter_name) == counters_taken.end()) {
flag = false;
continue;
}
}
} while (events_list.size() > 0);
if (flag) {
// std::cout << "Counter from Result Map: " << metric_name << std::endl;
counters_taken.insert(metric_name);
rocmtools::Metric* metric =
const_cast<rocmtools::Metric*>(metricsDict[gpu_agent.handle]->Get(metric_name));
if (metric == nullptr) std::cout << metric_name << " not found in metricsDict\n";
context->metrics_list.push_back(metric);
}
}
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();
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<uint32_t>(context->events_list.size()),
params,
0,
0,
0};
size_t ag_list_count = 1; // rocmtools::hsa_support::GetCPUAgentList().size();
hsa_agent_t ag_list[ag_list_count];
ag_list[0] = gpu_agent;
// 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 (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 = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.cpu_pool, size, 0, reinterpret_cast<void**>(&(profile->command_buffer.ptr)));
if (status != HSA_STATUS_SUCCESS) {
profile->command_buffer.ptr = malloc(size);
/*numa_alloc_onnode(
size,
rocmtools::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 = rocmtools::hsa_support::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 = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.kernarg_pool, size, 0, reinterpret_cast<void**>(&profile->output_buffer.ptr));
if (status != HSA_STATUS_SUCCESS) {
profile->output_buffer.ptr = malloc(size);
/*numa_alloc_onnode(
size,
rocmtools::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 = rocmtools::hsa_support::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);
context->start_packet = start_packet;
context->stop_packet = stop_packet;
context->read_packet = read_packet;
// add profiles
profiles.emplace_back(std::make_pair(context, profile));
}
// } while (events_list.size() > 0);
return profiles;
}
@@ -407,67 +397,71 @@ hsa_ven_amd_aqlprofile_profile_t* InitializeDeviceProfilingAqlPackets(
// Validating the events array for the specified gpu agent
bool result;
hsa_ven_amd_aqlprofile_validate_event(gpu_agent, events, &result);
if (!result) {
printf("Error: Events are not valid for the current gpu agent\n");
throw("Error: Events are not valid for the current gpu agent");
}
status = hsa_ven_amd_aqlprofile_validate_event(gpu_agent, events, &result);
CHECK_HSA_STATUS("Error: Events are not valid for the current gpu agent\n", status);
hsa_ven_amd_aqlprofile_parameter_t* params = {};
uint8_t* command_buffer = nullptr;
uint8_t* output_buffer = nullptr;
// Preparing the profile structure to get the packets
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wconversion-null"
// Preparing the profile structure to get the packets
hsa_ven_amd_aqlprofile_profile_t* profile = new hsa_ven_amd_aqlprofile_profile_t{
gpu_agent, HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC, events, event_count, params, 0, NULL, NULL};
#pragma GCC diagnostic pop
gpu_agent, HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC, events, event_count, params, 0, 0, 0};
// Preparing an Getting the size of the command and output buffers
status = hsa_ven_amd_aqlprofile_start(profile, NULL);
Agent::AgentInfo& agentInfo = rocmtools::hsa_support::GetAgentInfo(gpu_agent.handle);
size_t ag_list_count = 1;
hsa_agent_t ag_list[ag_list_count];
ag_list[0] = gpu_agent;
// Allocating Command Buffer
status = HSA_STATUS_ERROR;
size_t size = profile->command_buffer.size;
profile->command_buffer.ptr = nullptr;
if (size <= 0) return nullptr;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.cpu_pool, size, 0, reinterpret_cast<void**>(&command_buffer));
agentInfo.cpu_pool, size, 0, reinterpret_cast<void**>(&(profile->command_buffer.ptr)));
// Both the CPU and GPU can access the memory
if (status == HSA_STATUS_SUCCESS) {
hsa_agent_t ag_list[1] = {gpu_agent};
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
1, ag_list, NULL, command_buffer);
ag_list_count, ag_list, NULL, profile->command_buffer.ptr);
CHECK_HSA_STATUS("Error: GPU Agent can't have command buffer access", status);
} else {
profile->command_buffer.ptr = numa_alloc_onnode(
profile->command_buffer.size,
rocmtools::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle).getNumaNode());
if (profile->command_buffer.ptr != nullptr) {
status = HSA_STATUS_SUCCESS;
} else {
CHECK_HSA_STATUS("Error: Allocating Command Buffer", status);
}
}
profile->command_buffer.ptr = (status == HSA_STATUS_SUCCESS) ? command_buffer : nullptr;
if (status != HSA_STATUS_SUCCESS) printf("Error: Allocating Command Buffer\n");
// Allocating Output Buffer
status = HSA_STATUS_ERROR;
size = profile->output_buffer.size;
profile->output_buffer.ptr = nullptr;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.kernarg_pool, size, 0, reinterpret_cast<void**>(&output_buffer));
agentInfo.gpu_pool, size, 0, reinterpret_cast<void**>(&(profile->output_buffer.ptr)));
CHECK_HSA_STATUS("Error: Can't Allocate Output Buffer", status);
// Both the CPU and GPU can access the kernel arguments
if (status == HSA_STATUS_SUCCESS) {
hsa_agent_t ag_list[1] = {gpu_agent};
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
1, ag_list, NULL, output_buffer);
}
if (status == HSA_STATUS_SUCCESS) {
profile->output_buffer.ptr = output_buffer;
ag_list_count, ag_list, NULL, profile->output_buffer.ptr);
CHECK_HSA_STATUS("Error: Can't allow access on the Output Buffer for the GPU", status);
memset(profile->output_buffer.ptr, 0x0, profile->output_buffer.size);
} else {
profile->output_buffer.ptr = nullptr;
}
// Creating the start/stop/read packets
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);
if (status == HSA_STATUS_ERROR) return nullptr;
return profile;
@@ -479,17 +473,19 @@ bool g_output_buffer_local = true;
// Allocate system memory accessible by both CPU and GPU
uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool) {
size_t ag_list_count = 1; // rocmtools::hsa_support::GetCPUAgentList().size();
hsa_agent_t ag_list[ag_list_count];
ag_list[0] = gpu_agent;
hsa_status_t status = HSA_STATUS_ERROR;
uint8_t* buffer = NULL;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
// if (!cpu_agents_.empty()) {
status = hsa_amd_memory_pool_allocate(*cpu_pool, size, 0, reinterpret_cast<void**>(&buffer));
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
*cpu_pool, size, 0, reinterpret_cast<void**>(&buffer));
// Both the CPU and GPU can access the memory
if (status == HSA_STATUS_SUCCESS) {
hsa_agent_t ag_list[1] = {gpu_agent};
status = hsa_amd_agents_allow_access(1, ag_list, NULL, buffer);
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
ag_list_count, ag_list, NULL, buffer);
}
// }
uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
return ptr;
}
@@ -542,18 +538,6 @@ att_mem_pools_map_t* GetAttMemPoolsMap() {
return agent_att_mem_pools_map;
}
att_memory_pools_t* GetAttMemPools(hsa_agent_t gpu_agent) {
auto it = GetAttMemPoolsMap()->find(gpu_agent.handle);
if (it != GetAttMemPoolsMap()->end()) {
return it->second;
}
printf("Error: att_memory_pools_t instance not found for given gpu agent handle: %lu\n",
gpu_agent.handle);
return nullptr;
}
// Generate start and stop packets for collecting ATT traces
// Also generate and return the profile object which has the PM4
// command buffer and the output buffer for retrieving the traces
@@ -573,7 +557,7 @@ hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets(
// Check the profile buffer sizes
hsa_status_t status = hsa_ven_amd_aqlprofile_start(profile, NULL);
if (status != HSA_STATUS_SUCCESS) printf("Error: aqlprofile_start(NULL)");
CHECK_HSA_STATUS("Error: Getting PM4 Start Packet", status);
// TODO: create a separate class for memory allocations
// Maintain pools per device
// handle allocation and resource cleanup
@@ -582,14 +566,13 @@ hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets(
// command buffer -> from CPU memory pool
// output buffer -> from GPU memory pool
status = Allocate(gpu_agent, profile);
if (status != HSA_STATUS_SUCCESS) printf("Error: Allocate()");
CHECK_HSA_STATUS("Error: Att Buffers Allocation", status);
// Generate start/stop/read profiling packets
status = hsa_ven_amd_aqlprofile_start(profile, start_packet);
if (status != HSA_STATUS_SUCCESS) printf("Error: aqlprofile_start");
CHECK_HSA_STATUS("Error: Creating Start PM4 Packet", status);
status = hsa_ven_amd_aqlprofile_stop(profile, stop_packet);
if (status != HSA_STATUS_SUCCESS) printf("Error: aqlprofile_stop");
if (status == HSA_STATUS_ERROR) return nullptr;
CHECK_HSA_STATUS("Error: Creating Stop PM4 Packet", status);
return profile;
}
src/core/hsa/packets/packets_generator.h
+1 -1
View File
@@ -39,7 +39,7 @@ namespace Packet {
typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>*
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
std::vector<std::string>& counter_names, bool is_spm = false);
uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool);
src/core/hsa/queues/queue.cpp
+196 -173
View File
@@ -26,6 +26,7 @@
#include <vector>
#include <utility>
#include <algorithm>
#include <numa.h>
#include "rocprofiler.h"
#include "src/api/rocmtool.h"
@@ -33,6 +34,21 @@
#include "src/core/hsa/hsa_support.h"
#include "src/utils/helper.h"
#define CHECK_HSA_STATUS(msg, status) \
do { \
if ((status) != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) { \
try { \
const char* emsg = nullptr; \
hsa_status_string(status, &emsg); \
if (!emsg) emsg = "<Unknown HSA Error>"; \
std::cerr << msg << std::endl; \
std::cerr << emsg << std::endl; \
} catch (std::exception & e) { \
} \
abort(); \
} \
} while (0)
#define __NR_gettid 186
#define MAX_ATT_PROFILES 16
@@ -290,27 +306,41 @@ hsa_status_t attTraceDataCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
return status;
}
void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t& pending) {
std::vector<rocprofiler_record_counter_instance_t> counters_vec;
for (size_t i = 0; i < pending.context->metrics_list.size(); i++) {
const rocmtools::Metric* metric = pending.context->metrics_list[i];
void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t* pending) {
record->counters_count =
rocprofiler_record_counters_instances_count_t{pending->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<rocprofiler_record_counter_instance_t*>(malloc(counters_list_size));
for (size_t i = 0; i < pending->context->metrics_list.size(); i++) {
const rocmtools::Metric* metric = pending->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 = pending->context->results_map.find(metric_name);
if (it != pending->context->results_map.end()) {
value = it->second->val_double;
}
counters_vec.emplace_back(rocprofiler_record_counter_instance_t{
counters[i] = (rocprofiler_record_counter_instance_t{
// TODO(aelwazir): Moving to span once C++20 is adopted, strdup can be
// removed after that
rocprofiler_counter_id_t{rocmtools::profiler::GetCounterID(metric_name)},
rocprofiler_record_counter_value_t{value}});
}
record->counters = static_cast<rocprofiler_record_counter_instance_t*>(
malloc(counters_vec.size() * sizeof(rocprofiler_record_counter_instance_t)));
::memcpy(record->counters, &(counters_vec)[0],
counters_vec.size() * sizeof(rocprofiler_record_counter_instance_t));
record->counters_count = rocprofiler_record_counters_instances_count_t{counters_vec.size()};
record->counters = counters;
rocmtools::Session* session = GetROCMToolObj()->GetSession(pending->session_id);
void* initial_handle = const_cast<rocprofiler_record_counter_instance_t*>(record->counters);
if (session->FindBuffer(pending->buffer_id)) {
Memory::GenericBuffer* buffer = session->GetBuffer(pending->buffer_id);
buffer->AddRecord(*record, record->counters, counters_list_size,
[initial_handle](auto& record, const void* data) {
if (record.counters == initial_handle && data != initial_handle) {
free(initial_handle);
}
record.counters =
static_cast<const rocprofiler_record_counter_instance_t*>(data);
});
}
}
void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent,
@@ -330,7 +360,6 @@ void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent
for (trace_data_it = data.begin(); trace_data_it != data.end(); trace_data_it++) {
const void* data_ptr = trace_data_it->trace_data.ptr;
const uint32_t data_size = trace_data_it->trace_data.size;
// fprintf(arg->file, " SE(%u) size(%u)\n", data.sample_id, data_size);
void* buffer = NULL;
if (data_size != 0) {
@@ -359,23 +388,22 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
!GetROCMToolObj()->GetSession(queue_info_session->session_id)->GetProfiler())
return true;
rocmtools::Session* session = GetROCMToolObj()->GetSession(queue_info_session->session_id);
std::lock_guard<std::mutex> lock(session->GetSessionLock());
rocmtools::profiler::Profiler* profiler = session->GetProfiler();
std::vector<pending_signal_t>& pending_signals = const_cast<std::vector<pending_signal_t>&>(
std::vector<pending_signal_t*> pending_signals = const_cast<std::vector<pending_signal_t*>&>(
profiler->GetPendingSignals(queue_info_session->writer_id));
if (!pending_signals.empty()) {
for (auto it = pending_signals.begin(); it != pending_signals.end();
it = pending_signals.erase(it)) {
auto& pending = *it;
std::lock_guard<std::mutex> lock(session->GetSessionLock());
if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.signal)) return true;
if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->signal)) return true;
hsa_amd_profiling_dispatch_time_t time;
hsa_support::GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn(
queue_info_session->agent, pending.signal, &time);
queue_info_session->agent, pending->signal, &time);
uint32_t record_count = 1;
bool is_individual_xcc_mode = false;
uint32_t xcc_count =
hsa_support::GetAgentInfo(queue_info_session->agent.handle).getXccCount();
uint32_t xcc_count = queue_info_session->xcc_count;
if (xcc_count > 1) { // for MI300
const char* str = getenv("ROCPROFILER_INDIVIDUAL_XCC_MODE");
if (str != NULL) is_individual_xcc_mode = (atol(str) > 0);
@@ -387,70 +415,62 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
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)hsa_support::GetAgentInfo(queue_info_session->agent.handle).getIndex()};
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.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};
if (pending.counters_count > 0 && pending.context->metrics_list.size() > 0 &&
pending.profile) {
if (xcc_id == 0) // call to GetCounterData() is required only once for a dispatch
rocmtools::metrics::GetCounterData(pending.profile, queue_info_session->agent,
pending.context->results_list);
if (is_individual_xcc_mode)
rocmtools::metrics::GetCountersAndMetricResultsByXcc(
xcc_id, pending.context->results_list, pending.context->results_map,
pending.context->metrics_list);
else
rocmtools::metrics::GetMetricsData(pending.context->results_map,
pending.context->metrics_list);
AddRecordCounters(&record, pending);
}
// 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_PROFILER_RECORD,
rocprofiler_record_id_t{pending.kernel_descriptor}};
record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor};
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 = GetROCMToolObj()->GetCurrentSessionId();
if (pending->session_id.handle == 0) {
pending->session_id = GetROCMToolObj()->GetCurrentSessionId();
}
if (session->FindBuffer(pending.buffer_id)) {
Memory::GenericBuffer* buffer = session->GetBuffer(pending.buffer_id);
if (pending.profile && pending.counters_count > 0) {
rocprofiler_record_counter_instance_t* record_counters = record.counters;
buffer->AddRecord(
record, record.counters,
(record.counters_count.value * (sizeof(rocprofiler_record_counter_instance_t) + 1)),
[](auto& record, const void* data) {
record.counters = const_cast<rocprofiler_record_counter_instance_t*>(
static_cast<const rocprofiler_record_counter_instance_t*>(data));
});
free(record_counters);
} else {
if (pending->counters_count > 0 && pending->context->metrics_list.size() > 0 &&
pending->profile) {
if (xcc_id == 0) // call to GetCounterData() is required only once for a dispatch
rocmtools::metrics::GetCounterData(pending->profile, queue_info_session->agent,
pending->context->results_list);
if (is_individual_xcc_mode)
rocmtools::metrics::GetCountersAndMetricResultsByXcc(
xcc_id, pending->context->results_list, pending->context->results_map,
pending->context->metrics_list);
else
rocmtools::metrics::GetMetricsData(pending->context->results_map,
pending->context->metrics_list);
AddRecordCounters(&record, pending);
} 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) {
if (pending->counters_count > 0 && pending->profile && pending->profile->events) {
// TODO(aelwazir): we need a better way of distributing events and free them
// free(const_cast<hsa_ven_amd_aqlprofile_event_t*>(pending.profile->events));
// if (pending->profile->output_buffer.ptr)
// numa_free(pending->profile->output_buffer.ptr, pending->profile->output_buffer.size);
hsa_status_t status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
(pending.profile->output_buffer.ptr));
if (status != HSA_STATUS_SUCCESS) {
printf("Error: Couldn't free output buffer memory\n");
}
(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 = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
(pending.profile->command_buffer.ptr));
if (status != HSA_STATUS_SUCCESS) {
printf("Error: Couldn't free command buffer memory\n");
(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.profile;
delete pending->context;
}
if (pending.signal.handle)
hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(pending.signal);
if (pending->signal.handle)
hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(pending->signal);
if (queue_info_session->interrupt_signal.handle)
hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(queue_info_session->interrupt_signal);
}
@@ -483,8 +503,7 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.signal)) return true;
rocprofiler_record_att_tracer_t record{};
record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor};
record.gpu_id = rocprofiler_agent_id_t{
(uint64_t)hsa_support::GetAgentInfo(queue_info_session->agent.handle).getIndex()};
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};
@@ -506,14 +525,10 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
}
hsa_status_t status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
(pending.profile->output_buffer.ptr));
if (status != HSA_STATUS_SUCCESS) {
printf("Error: Couldn't free output buffer memory\n");
}
CHECK_HSA_STATUS("Error: Couldn't free output buffer memory", status);
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
(pending.profile->command_buffer.ptr));
if (status != HSA_STATUS_SUCCESS) {
printf("Error: Couldn't free command buffer memory\n");
}
CHECK_HSA_STATUS("Error: Couldn't free command buffer memory", status);
delete pending.profile;
}
}
@@ -540,19 +555,19 @@ void AddVendorSpecificPacket(const Packet::packet_t* packet,
void SignalAsyncHandler(const hsa_signal_t& signal, void* data) {
hsa_status_t status = hsa_support::GetAmdExtTable().hsa_amd_signal_async_handler_fn(
signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalHandler, data);
if (status != HSA_STATUS_SUCCESS) fatal("hsa_amd_signal_async_handler failed");
CHECK_HSA_STATUS("Error: hsa_amd_signal_async_handler failed", status);
}
void signalAsyncHandlerATT(const hsa_signal_t& signal, void* data) {
hsa_status_t status = hsa_support::GetAmdExtTable().hsa_amd_signal_async_handler_fn(
signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalHandlerATT, data);
if (status != HSA_STATUS_SUCCESS) fatal("hsa_amd_signal_async_handler failed");
CHECK_HSA_STATUS("Error: hsa_amd_signal_async_handler for ATT failed", status);
}
void CreateSignal(uint32_t attribute, hsa_signal_t* signal) {
hsa_status_t status =
hsa_support::GetAmdExtTable().hsa_amd_signal_create_fn(1, 0, nullptr, attribute, signal);
if (status != HSA_STATUS_SUCCESS) fatal("hsa_amd_signal_create failed");
CHECK_HSA_STATUS("Error: hsa_amd_signal_create failed", status);
}
template <typename Integral = uint64_t> constexpr Integral bit_mask(int first, int last) {
@@ -659,13 +674,13 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
is_pc_sampling_collection_mode) &&
session) {
// Getting Queue Data and Information
auto& queue_info = *static_cast<Queue*>(data);
Queue& queue_info = *reinterpret_cast<Queue*>(data);
std::lock_guard<std::mutex> lk(queue_info.qw_mutex);
// hsa_ven_amd_aqlprofile_profile_t* profile;
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>*
profiles = nullptr;
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>
profiles;
// Searching accross all the packets given during this write
@@ -686,106 +701,109 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
// Get the PM4 Packets using packets_generator
profiles = Packet::InitializeAqlPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(),
session_data);
replay_mode_count = profiles->size();
replay_mode_count = profiles.size();
}
uint32_t profile_id = 0;
hsa_signal_t interrupt_signal;
do {
std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*> profile;
if (profiles && replay_mode_count > 0) profile = profiles->at(profile_id);
// do {
std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*> profile;
if (profiles.size() > 0 && replay_mode_count > 0) profile = profiles.at(profile_id);
uint32_t writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
uint32_t writer_id = WRITER_ID.fetch_add(1, std::memory_order_release);
if (session_data_count > 0 && is_counter_collection_mode && profiles &&
replay_mode_count > 0) {
// 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;
AddVendorSpecificPacket(profile.first->start_packet, &transformed_packets, dummy_signal);
if (session_data_count > 0 && is_counter_collection_mode && profiles.size() > 0 &&
replay_mode_count > 0) {
// 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;
AddVendorSpecificPacket(profile.first->start_packet, &transformed_packets, dummy_signal);
CreateBarrierPacket(profile.first->start_packet->completion_signal, &transformed_packets);
}
CreateBarrierPacket(profile.first->start_packet->completion_signal, &transformed_packets);
}
auto& packet = transformed_packets.emplace_back(packets_arr[i]);
auto& dispatch_packet = reinterpret_cast<hsa_kernel_dispatch_packet_t&>(packet);
auto& packet = transformed_packets.emplace_back(packets_arr[i]);
auto& dispatch_packet = reinterpret_cast<hsa_kernel_dispatch_packet_t&>(packet);
/*
* Only PC sampling relies on this right now, so it would be better to
* only generate an ID if PC sampling is active to conserve IDs, but it's
* unlikely 64 bits' worth of identifiers will be exhausted during the
* lifetime of the ROCMToolObj.
*/
dispatch_packet.reserved2 = GetROCMToolObj()->GetUniqueKernelDispatchId();
/*
* Only PC sampling relies on this right now, so it would be better to
* only generate an ID if PC sampling is active to conserve IDs, but it's
* unlikely 64 bits' worth of identifiers will be exhausted during the
* lifetime of the ROCMToolObj.
*/
uint64_t correlation_id = dispatch_packet.reserved2;
// dispatch_packet.reserved2 = GetROCMToolObj()->GetUniqueKernelDispatchId();
CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal);
// Adding the dispatch packet newly created signal to the pending signals
// list to be processed by the signal interrupt
rocprofiler_kernel_properties_t kernel_properties =
set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent());
if (session) {
uint64_t record_id = GetROCMToolObj()->GetUniqueRecordId();
AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object),
record_id);
if (profiles && replay_mode_count > 0) {
session->GetProfiler()->AddPendingSignals(
writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id,
profile.first, profile.first->metrics_list.size(), profile.second,
kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index);
} else {
session->GetProfiler()->AddPendingSignals(
writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id,
nullptr, 0, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid),
user_pkt_index);
}
}
// Make a copy of the original packet, adding its signal to a barrier
// packet and create a new signal for it to get timestamps
if (original_packet.completion_signal.handle) {
hsa_barrier_and_packet_t barrier{0};
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
Packet::packet_t* __attribute__((__may_alias__)) pkt =
(reinterpret_cast<Packet::packet_t*>(&barrier));
transformed_packets.emplace_back(*pkt).completion_signal =
original_packet.completion_signal;
}
// Adding a barrier packet with the original packet's completion signal.
CreateSignal(0, &interrupt_signal);
// Adding Stop and Read PM4 Packets
if (session_data_count > 0 && is_counter_collection_mode && profiles &&
replay_mode_count > 0) {
hsa_signal_t dummy_signal{};
profile.first->stop_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
AddVendorSpecificPacket(profile.first->stop_packet, &transformed_packets, dummy_signal);
profile.first->read_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
AddVendorSpecificPacket(profile.first->read_packet, &transformed_packets,
interrupt_signal);
// Added Interrupt Signal with barrier and provided handler for it
CreateBarrierPacket(interrupt_signal, &transformed_packets);
CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal);
// Adding the dispatch packet newly created signal to the pending signals
// list to be processed by the signal interrupt
rocprofiler_kernel_properties_t kernel_properties =
set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent());
if (session) {
uint64_t record_id = GetROCMToolObj()->GetUniqueRecordId();
AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object),
record_id);
if (profiles.size() > 0 && replay_mode_count > 0) {
session->GetProfiler()->AddPendingSignals(
writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id,
profile.first, profile.first->metrics_list.size(), profile.second, kernel_properties,
(uint32_t)syscall(__NR_gettid), user_pkt_index, correlation_id);
} else {
hsa_barrier_and_packet_t barrier{0};
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
barrier.completion_signal = interrupt_signal;
Packet::packet_t* __attribute__((__may_alias__)) pkt =
(reinterpret_cast<Packet::packet_t*>(&barrier));
transformed_packets.emplace_back(*pkt);
session->GetProfiler()->AddPendingSignals(
writer_id, record_id, dispatch_packet.completion_signal, session_id, buffer_id,
nullptr, 0, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid),
user_pkt_index, correlation_id);
}
// 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});
ACTIVE_INTERRUPT_SIGNAL_COUNT.fetch_add(1, std::memory_order_relaxed);
profile_id++;
} while (replay_mode_count > 0 && profile_id < replay_mode_count); // Profiles loop end
}
// Make a copy of the original packet, adding its signal to a barrier
// packet and create a new signal for it to get timestamps
if (original_packet.completion_signal.handle) {
hsa_barrier_and_packet_t barrier{0};
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
Packet::packet_t* __attribute__((__may_alias__)) pkt =
(reinterpret_cast<Packet::packet_t*>(&barrier));
transformed_packets.emplace_back(*pkt).completion_signal =
original_packet.completion_signal;
}
hsa_signal_t interrupt_signal{};
// Adding a barrier packet with the original packet's completion signal.
CreateSignal(0, &interrupt_signal);
// Adding Stop and Read PM4 Packets
if (session_data_count > 0 && is_counter_collection_mode) {
hsa_signal_t dummy_signal{};
profile.first->stop_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
AddVendorSpecificPacket(profile.first->stop_packet, &transformed_packets, dummy_signal);
profile.first->read_packet->header = HSA_PACKET_TYPE_VENDOR_SPECIFIC
<< HSA_PACKET_HEADER_TYPE;
AddVendorSpecificPacket(profile.first->read_packet, &transformed_packets, interrupt_signal);
// Added Interrupt Signal with barrier and provided handler for it
CreateBarrierPacket(interrupt_signal, &transformed_packets);
} else {
hsa_barrier_and_packet_t barrier{0};
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
barrier.completion_signal = interrupt_signal;
Packet::packet_t* __attribute__((__may_alias__)) pkt =
(reinterpret_cast<Packet::packet_t*>(&barrier));
transformed_packets.emplace_back(*pkt);
}
Agent::AgentInfo& agentInfo =
rocmtools::hsa_support::GetAgentInfo(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.getIndex(),
agentInfo.getXccCount()});
ACTIVE_INTERRUPT_SIGNAL_COUNT.fetch_add(1, std::memory_order_relaxed);
// profile_id++;
// } while (replay_mode_count > 0 && profile_id < replay_mode_count); // Profiles loop end
}
/* Write the transformed packets to the hardware queue. */
writer(&transformed_packets[0], transformed_packets.size());
@@ -795,7 +813,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
// Getting Queue Data and Information
auto& queue_info = *static_cast<Queue*>(data);
std::lock_guard<std::mutex> lk(queue_info.qw_mutex);
Agent::AgentInfo* agentInfo = &(hsa_support::GetAgentInfo(queue_info.GetGPUAgent().handle));
Agent::AgentInfo agentInfo = hsa_support::GetAgentInfo(queue_info.GetGPUAgent().handle);
bool can_profile_anypacket = false;
std::vector<bool> can_profile_packet;
@@ -858,7 +876,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
}
if (att_counters_names.size() > 0) {
MetricsDict* metrics_dict_ = MetricsDict::Create(agentInfo);
MetricsDict* metrics_dict_ = MetricsDict::Create(&agentInfo);
for (const std::string& counter_name : att_counters_names) {
const Metric* metric = metrics_dict_->Get(counter_name);
@@ -1007,6 +1025,11 @@ Queue::Queue(const hsa_agent_t& cpu_agent, const hsa_agent_t& gpu_agent, uint32_
*queue = intercept_queue_;
}
Queue::~Queue() {
while (ACTIVE_INTERRUPT_SIGNAL_COUNT.load(std::memory_order_acquire) > 0) {
}
}
hsa_queue_t* Queue::GetCurrentInterceptQueue() { return intercept_queue_; }
hsa_agent_t Queue::GetGPUAgent() { return gpu_agent_; }
src/core/hsa/queues/queue.h
+3 -1
View File
@@ -56,7 +56,7 @@ class Queue {
hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data,
uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue);
~Queue() {}
~Queue();
hsa_queue_t* GetCurrentInterceptQueue();
hsa_agent_t GetGPUAgent();
@@ -82,6 +82,8 @@ struct queue_info_session_t {
uint64_t queue_id;
uint32_t writer_id;
hsa_signal_t interrupt_signal;
uint64_t gpu_index;
uint32_t xcc_count;
};
void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t& pending);