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

SWDEV-399631: Converting into singleton class

Browse Source 
Rocprofiler and HSAsupport classes have been implemented as
singletons which gets initialized lazily.

Change-Id: I98db4713c7282d88966aeb0ea9df83ba457b2ea3


[ROCm/rocprofiler commit: 4980409c5a]
This commit is contained in:
Sriraksha Nagaraj
2023-04-12 11:26:18 -05:00
parent 162eba6642
commit ca53c3f18d
45 changed files with 1995 additions and 1335 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler/src/core/hsa/hsa_common.cpp
-123
View File
@@ -1,123 +0,0 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE. */
#include "hsa_common.h"
#include "src/utils/exception.h"
namespace rocprofiler {
namespace hsa_support {
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()) {
return agent_info_map.at(handle);
} 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});
}
}
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()};
}
}
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{};
CoreApiTable& GetCoreApiTable() { return saved_core_api; }
void SetCoreApiTable(const CoreApiTable& table) { saved_core_api = table; }
AmdExtTable saved_amd_ext_api{};
AmdExtTable GetAmdExtTable() { return saved_amd_ext_api; }
void SetAmdExtTable(AmdExtTable* table) { saved_amd_ext_api = *table; }
hsa_ven_amd_loader_1_01_pfn_t hsa_loader_api{};
hsa_ven_amd_loader_1_01_pfn_t GetHSALoaderApi() { return hsa_loader_api; }
void SetHSALoaderApi() {
hsa_status_t status = saved_core_api.hsa_system_get_major_extension_table_fn(
HSA_EXTENSION_AMD_LOADER, 1, sizeof(hsa_ven_amd_loader_1_01_pfn_t), &hsa_loader_api);
if (status != HSA_STATUS_SUCCESS) fatal("hsa_system_get_major_extension_table failed");
}
void ResetMaps() {
if (hsa_status_t status = saved_amd_ext_api.hsa_amd_profiling_async_copy_enable_fn(false);
status != HSA_STATUS_SUCCESS)
assert(!"hsa_amd_profiling_async_copy_enable failed");
memset(&saved_core_api, '\0', sizeof(saved_core_api));
memset(&saved_amd_ext_api, '\0', sizeof(saved_amd_ext_api));
memset(&hsa_loader_api, '\0', sizeof(hsa_loader_api));
}
rocprofiler_timestamp_t GetCurrentTimestampNS() {
// If the HSA intercept is installed, then use the "original"
// 'hsa_system_get_info' function to avoid reporting calls for internal use
// of the HSA API by the tracer.
auto hsa_system_get_info_fn = saved_core_api.hsa_system_get_info_fn;
// If the HSA intercept is not installed, use the default
// 'hsa_system_get_info'.
if (hsa_system_get_info_fn == nullptr) hsa_system_get_info_fn = hsa_system_get_info;
uint64_t sysclock;
if (hsa_status_t status = hsa_system_get_info_fn(HSA_SYSTEM_INFO_TIMESTAMP, &sysclock);
status == HSA_STATUS_ERROR_NOT_INITIALIZED)
return rocprofiler_timestamp_t{0};
else if (status != HSA_STATUS_SUCCESS)
assert(!"hsa_system_get_info failed");
static uint64_t sysclock_period = [&]() {
uint64_t sysclock_hz = 0;
if (hsa_status_t status =
hsa_system_get_info_fn(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz);
status != HSA_STATUS_SUCCESS)
assert(!"hsa_system_get_info failed");
return (uint64_t)1000000000 / sysclock_hz;
}();
return rocprofiler_timestamp_t{sysclock * sysclock_period};
}
} // namespace hsa_support
} // namespace rocprofiler
projects/rocprofiler/src/core/hsa/hsa_common.h
-64
View File
@@ -1,64 +0,0 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE. */
#ifndef SRC_CORE_HSA_HSA_COMMON_H_
#define SRC_CORE_HSA_HSA_COMMON_H_
#include <hsa/hsa.h>
#include <hsa/hsa_api_trace.h>
#include <hsa/hsa_ext_amd.h>
#include <hsa/hsa_ven_amd_aqlprofile.h>
#include <hsa/hsa_ven_amd_loader.h>
#include <map>
#include <mutex>
#include "rocprofiler.h"
#include "src/core/hardware/hsa_info.h"
#define ASSERTM(exp, msg) assert(((void)msg, exp))
namespace rocprofiler {
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(uint64_t agent_index);
CoreApiTable& GetCoreApiTable();
void SetCoreApiTable(const CoreApiTable& table);
AmdExtTable GetAmdExtTable();
void SetAmdExtTable(AmdExtTable* table);
hsa_ven_amd_loader_1_01_pfn_t GetHSALoaderApi();
void SetHSALoaderApi();
void ResetMaps();
rocprofiler_timestamp_t GetCurrentTimestampNS();
} // namespace hsa_support
} // namespace rocprofiler
#endif // SRC_CORE_HSA_HSA_COMMON_H_
projects/rocprofiler/src/core/hsa/hsa_support.cpp
+397 -378
View File
File diff suppressed because it is too large Load Diff
projects/rocprofiler/src/core/hsa/hsa_support.h
+72 -12
View File
@@ -29,8 +29,10 @@
#include <atomic>
#include <string>
#include <unordered_map>
#include <memory>
#include "hsa_common.h"
#include "rocprofiler.h"
#include "src/core/hardware/hsa_info.h"
// HSA EVT data type
@@ -92,15 +94,78 @@ typedef struct {
namespace rocprofiler {
namespace hsa_support {
void Initialize(HsaApiTable* Table);
hsa_status_t hsa_iterate_agents_cb(hsa_agent_t agent, void* data);
void Finalize();
class HSAAgentInfo {
private:
hsa_agent_t agent_;
Agent::DeviceInfo device_info_;
hsa_agent_t near_cpu_agent_;
hsa_device_type_t type_;
bool IterateCounters(rocprofiler_counters_info_callback_t counters_info_callback);
} // namespace hsa_support
public:
HSAAgentInfo(hsa_agent_t agent, hsa_device_type_t type) : agent_(agent), type_(type){};
uint64_t getHandle() const;
const Agent::DeviceInfo& GetDeviceInfo() const;
void SetNearCpuAgent(hsa_agent_t near_cpu_agent);
void SetDeviceInfo(Agent::DeviceInfo device_info);
hsa_agent_t GetNearCpuAgent() const;
hsa_device_type_t GetType() const;
hsa_amd_memory_pool_t cpu_pool_;
hsa_amd_memory_pool_t kernarg_pool_;
hsa_amd_memory_pool_t gpu_pool_;
};
struct queues_deleter {
queues_deleter() {};
queues_deleter(queues_deleter&) { };
void operator() (void * queue) const;
};
class HSASupport_Singleton {
private:
HSASupport_Singleton() {};
~HSASupport_Singleton() = delete;
CoreApiTable saved_core_api;
AmdExtTable saved_amd_ext_api;
hsa_ven_amd_loader_1_01_pfn_t hsa_loader_api;
std::mutex info_map_mutex_;
std::unordered_map<uint64_t, HSAAgentInfo> HSAagent_info_map_;
std::atomic<bool> ksymbols_flag{true};
std::atomic<bool> kernel_names_flag{true};
std::mutex queues_mutex_;
std::unordered_map<hsa_queue_t*, std::unique_ptr<void, queues_deleter&>> queues;
void SetCoreApiTable(CoreApiTable& table);
void SetAmdExtTable(AmdExtTable& table);
void SetHSALoaderApi();
public:
std::vector<hsa_agent_t> gpu_agents;
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<std::string, std::vector<uint64_t>>* kernel_names;
std::mutex ksymbol_map_lock;
std::map<uint64_t, std::string>* ksymbols;
void SetHSAAgentInfo(hsa_agent_t agent, HSAAgentInfo hsa_agent_info);
static HSASupport_Singleton& GetInstance();
CoreApiTable& GetCoreApiTable();
AmdExtTable& GetAmdExtTable();
hsa_ven_amd_loader_1_01_pfn_t& GetHSALoaderApi();
void AddQueue(hsa_queue_t* queue, std::unique_ptr<void, queues_deleter&>);
void RemoveQueue(hsa_queue_t* queue);
void HSAInitialize(HsaApiTable* Table);
void HSAFinalize();
void InitKsymbols();
void FinitKsymbols();
HSASupport_Singleton(const HSASupport_Singleton&) = delete;
HSASupport_Singleton& operator=(const HSASupport_Singleton&) = delete;
};
bool hsa_support_IterateCounters(rocprofiler_counters_info_callback_t counters_info_callback);
} // namespace rocprofiler
#include "src/core/session/tracer/src/roctracer.h"
@@ -126,11 +191,6 @@ uint32_t GetApiCode(const char* str);
void RegisterTracerCallback(int (*function)(rocprofiler_tracer_activity_domain_t domain,
uint32_t operation_id, void* data));
rocprofiler_timestamp_t timestamp_ns();
void Initialize_roctracer(HsaApiTable* table);
} // namespace roctracer::hsa_support
#endif // SRC_CORE_HSA_HSA_SUPPORT_H_
projects/rocprofiler/src/core/hsa/packets/packets_generator.cpp
+60 -47
View File
@@ -40,7 +40,6 @@
#include "src/core/counters/basic/basic_counter.h"
#include "src/utils/exception.h"
#include "src/utils/logger.h"
#include "src/core/hsa/hsa_common.h"
#include "src/core/counters/metrics/metrics.h"
#include "src/core/hardware/hsa_info.h"
@@ -83,13 +82,14 @@ static hsa_status_t FindGlobalPool(hsa_amd_memory_pool_t pool, void* data, bool
if (nullptr == data) {
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
err = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
err = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment);
ASSERTM(err != HSA_STATUS_ERROR, "hsa_amd_memory_pool_get_info");
if (HSA_AMD_SEGMENT_GLOBAL != segment) {
return HSA_STATUS_SUCCESS;
}
err = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
err = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag);
ASSERTM(err != HSA_STATUS_ERROR, "hsa_amd_memory_pool_get_info");
uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT;
@@ -114,20 +114,22 @@ hsa_status_t FindKernArgPool(hsa_amd_memory_pool_t pool, void* data) {
return FindGlobalPool(pool, data, true);
}
void InitializePools(hsa_agent_t cpu_agent, Agent::AgentInfo* agent_info) {
void InitializePools(hsa_agent_t cpu_agent, rocprofiler::HSAAgentInfo* agent_info) {
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
hsa_status_t status =
rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
cpu_agent, FindStandardPool, &(agent_info->cpu_pool));
hsasupport_singleton.GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
cpu_agent, FindStandardPool, &(agent_info->cpu_pool_));
CHECK_HSA_STATUS("Error: Command Buffer Pool is not initialized", status);
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
cpu_agent, FindKernArgPool, &(agent_info->kernarg_pool));
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
cpu_agent, FindKernArgPool, &(agent_info->kernarg_pool_));
CHECK_HSA_STATUS("Error: Output Buffer Pool is not initialized", status);
}
void InitializeGPUPool(hsa_agent_t gpu_agent, Agent::AgentInfo* agent_info) {
void InitializeGPUPool(hsa_agent_t gpu_agent, rocprofiler::HSAAgentInfo* agent_info) {
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
hsa_status_t status =
hsa_amd_agent_iterate_memory_pools(gpu_agent, FindStandardPool, &(agent_info->gpu_pool));
hsasupport_singleton.GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(gpu_agent, FindStandardPool, &(agent_info->gpu_pool_));
CHECK_HSA_STATUS("hsa_amd_agent_iterate_memory_pools(gpu_pool)", status);
}
@@ -136,13 +138,18 @@ struct block_des_t {
uint32_t index;
};
std::map<uint32_t, rocprofiler::MetricsDict*> metricsDict;
static std::atomic<bool> counters_added{false};
void CheckPacketReqiurements(std::vector<hsa_agent_t>& gpu_agents) {
for (auto& gpu_agent : gpu_agents) {
std::map<uint32_t, rocprofiler::MetricsDict*> metricsDict;
void CheckPacketReqiurements() {
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
for (auto& gpu_agent : hsasupport_singleton.gpu_agents) {
// get the instance of MetricsDict
Agent::AgentInfo& agentInfo = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle);
rocprofiler::HSAAgentInfo& agentInfo = hsasupport_singleton.GetHSAAgentInfo(gpu_agent.handle);
metricsDict[gpu_agent.handle] = rocprofiler::MetricsDict::Create(&agentInfo);
}
}
@@ -155,18 +162,18 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
std::vector<std::string>& 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)) {
for (auto& name : counter_names) {
rocprofiler::GetROCProfilerSingleton()
->GetSession(session_id)
->GetProfiler()
->AddCounterName(name);
if (rocprofiler_singleton.HasActiveSession()) {
rocprofiler_singleton.GetSession(session_id)->GetProfiler()->AddCounterName(name);
}
}
counters_added.exchange(true, std::memory_order_release);
}
Agent::AgentInfo& agentInfo = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle);
rocprofiler::HSAAgentInfo& agentInfo = hsasupport_singleton.GetHSAAgentInfo(gpu_agent.handle);
std::map<std::string, rocprofiler::results_t*> results_map;
std::vector<rocprofiler::event_t> events_list;
std::vector<rocprofiler::results_t*> results_list;
@@ -330,8 +337,8 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
<< "Error: Command buffer given size is " << size << std::endl;
abort();
}
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.cpu_pool, size, 0, reinterpret_cast<void**>(&(profile->command_buffer.ptr)));
status =hsasupport_singleton.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(
@@ -344,11 +351,10 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
}
} else {
// Both the CPU and GPU can access the memory
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
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;
@@ -358,8 +364,8 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
<< "Error: Output buffer given size is " << size << std::endl;
abort();
}
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.kernarg_pool, size, 0, reinterpret_cast<void**>(&profile->output_buffer.ptr));
status =hsasupport_singleton.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(
@@ -372,7 +378,7 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
abort();
}
} else {
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
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);
@@ -420,29 +426,34 @@ hsa_ven_amd_aqlprofile_profile_t* InitializeDeviceProfilingAqlPackets(
// Preparing an Getting the size of the command and output buffers
status = hsa_ven_amd_aqlprofile_start(profile, NULL);
Agent::AgentInfo& agentInfo = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle);
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
rocprofiler::HSAAgentInfo& agentInfo = hsasupport_singleton.GetHSAAgentInfo(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
//FixMe: Command buffer and output buffers are allocated repetatively.
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 = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.cpu_pool, size, 0, reinterpret_cast<void**>(&(profile->command_buffer.ptr)));
status =hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
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) {
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
status =hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn(
ag_list_count, ag_list, NULL, profile->command_buffer.ptr);
CHECK_HSA_STATUS("Error: GPU Agent can't have command buffer access", status);
} else {
hsa_agent_t near_cpu_node = agentInfo.GetNearCpuAgent();
uint32_t near_cpu_node_id = 0;
hsasupport_singleton.GetCoreApiTable().hsa_agent_get_info_fn(near_cpu_node,
HSA_AGENT_INFO_NODE, &near_cpu_node_id);
profile->command_buffer.ptr = numa_alloc_onnode(
profile->command_buffer.size,
rocprofiler::hsa_support::GetAgentInfo(agentInfo.getNearCpuAgent().handle).getNumaNode());
near_cpu_node_id);
if (profile->command_buffer.ptr != nullptr) {
status = HSA_STATUS_SUCCESS;
} else {
@@ -455,12 +466,12 @@ hsa_ven_amd_aqlprofile_profile_t* InitializeDeviceProfilingAqlPackets(
size = profile->output_buffer.size;
profile->output_buffer.ptr = nullptr;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
agentInfo.gpu_pool, size, 0, reinterpret_cast<void**>(&(profile->output_buffer.ptr)));
status =hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
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) {
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
status =hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn(
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);
@@ -490,11 +501,12 @@ uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_po
hsa_status_t status = HSA_STATUS_ERROR;
uint8_t* buffer = NULL;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
status =hsasupport_singleton.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) {
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agents_allow_access_fn(
ag_list_count, ag_list, NULL, buffer);
}
uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
@@ -504,32 +516,33 @@ uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_po
// Allocate memory for use by a kernel of specified size
uint8_t* AllocateLocalMemory(size_t size, hsa_amd_memory_pool_t* gpu_pool) {
hsa_status_t status = HSA_STATUS_ERROR;
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
uint8_t* buffer = NULL;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = hsa_amd_memory_pool_allocate(*gpu_pool, size, 0, reinterpret_cast<void**>(&buffer));
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(*gpu_pool, size, 0, reinterpret_cast<void**>(&buffer));
uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
return ptr;
}
hsa_status_t Allocate(hsa_agent_t gpu_agent, hsa_ven_amd_aqlprofile_profile_t* profile,
size_t att_buffer_size) {
Agent::AgentInfo& agentInfo = rocprofiler::hsa_support::GetAgentInfo(gpu_agent.handle);
hsa_status_t Allocate(hsa_agent_t gpu_agent, hsa_ven_amd_aqlprofile_profile_t* profile, size_t att_buffer_size) {
rocprofiler::HSAAgentInfo& agentInfo = rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo(gpu_agent.handle);
profile->command_buffer.ptr =
AllocateSysMemory(gpu_agent, profile->command_buffer.size, &agentInfo.cpu_pool);
AllocateSysMemory(gpu_agent, profile->command_buffer.size, &agentInfo.cpu_pool_);
profile->output_buffer.size = att_buffer_size;
profile->output_buffer.ptr = (g_output_buffer_local)
? AllocateLocalMemory(profile->output_buffer.size, &agentInfo.gpu_pool)
: AllocateSysMemory(gpu_agent, profile->output_buffer.size, &agentInfo.cpu_pool);
? AllocateLocalMemory(profile->output_buffer.size, &agentInfo.gpu_pool_)
: AllocateSysMemory(gpu_agent, profile->output_buffer.size, &agentInfo.cpu_pool_);
return (profile->command_buffer.ptr && profile->output_buffer.ptr) ? HSA_STATUS_SUCCESS
: HSA_STATUS_ERROR;
}
bool AllocateMemoryPools(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
hsa_amd_memory_pool_t* cpu_pool, hsa_amd_memory_pool_t* gpu_pool) {
hsa_status_t status = hsa_amd_agent_iterate_memory_pools(cpu_agent, FindStandardPool, cpu_pool);
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
hsa_status_t status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(cpu_agent, FindStandardPool, cpu_pool);
CHECK_HSA_STATUS("hsa_amd_agent_iterate_memory_pools(cpu_pool)", status);
status = hsa_amd_agent_iterate_memory_pools(gpu_agent, FindStandardPool, gpu_pool);
status = hsasupport_singleton.GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(gpu_agent, FindStandardPool, gpu_pool);
CHECK_HSA_STATUS("hsa_amd_agent_iterate_memory_pools(gpu_pool)", status);
return true;
projects/rocprofiler/src/core/hsa/packets/packets_generator.h
+3 -4
View File
@@ -46,8 +46,8 @@ InitializeAqlPackets(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool);
void GetCommandBufferMap(std::map<size_t, uint8_t*>);
void GetOutputBufferMap(std::map<size_t, uint8_t*>);
void InitializePools(hsa_agent_t cpu_agent, Agent::AgentInfo* agent_info);
void InitializeGPUPool(hsa_agent_t gpu_agent, Agent::AgentInfo* agent_info);
void InitializePools(hsa_agent_t cpu_agent, rocprofiler::HSAAgentInfo* agent_info);
void InitializeGPUPool(hsa_agent_t gpu_agent, rocprofiler::HSAAgentInfo* agent_info);
hsa_ven_amd_aqlprofile_profile_t* InitializeDeviceProfilingAqlPackets(
hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, hsa_ven_amd_aqlprofile_event_t* events,
uint32_t event_count, packet_t* start_packet, packet_t* stop_packet, packet_t* read_packet);
@@ -65,8 +65,7 @@ uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_po
void get_command_buffer_map(std::map<size_t, uint8_t*>);
void get_outbuffer_map(std::map<size_t, uint8_t*>);
void initialize_pools(hsa_agent_t cpu_agent);
void CheckPacketReqiurements(std::vector<hsa_agent_t>& gpu_agents);
void CheckPacketReqiurements();
typedef struct {
hsa_amd_memory_pool_t cpu_mem_pool;
projects/rocprofiler/src/core/hsa/queues/queue.cpp
+96 -151
View File
@@ -35,6 +35,7 @@
#include "src/utils/helper.h"
#include "src/core/isa_capture/code_object_track.hpp"
#define CHECK_HSA_STATUS(msg, status) \
do { \
if ((status) != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) { \
@@ -74,37 +75,37 @@ static inline bool IsEventMatch(const hsa_ven_amd_aqlprofile_event_t& event1,
typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> att_trace_callback_data_t;
static std::mutex ksymbol_map_lock;
static std::map<uint64_t, std::string>* ksymbols;
static std::atomic<bool> ksymbols_flag{true};
void AddKernelName(uint64_t handle, std::string name) {
std::lock_guard<std::mutex> lock(ksymbol_map_lock);
ksymbols->emplace(handle, name);
HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance();
std::lock_guard<std::mutex> lock(hsasupport_singleton.ksymbol_map_lock);
hsasupport_singleton.ksymbols->emplace(handle, name);
}
void RemoveKernelName(uint64_t handle) {
std::lock_guard<std::mutex> lock(ksymbol_map_lock);
ksymbols->erase(handle);
HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance();
std::lock_guard<std::mutex> lock(hsasupport_singleton.ksymbol_map_lock);
hsasupport_singleton.ksymbols->erase(handle);
}
std::string GetKernelNameFromKsymbols(uint64_t handle) {
std::lock_guard<std::mutex> lock(ksymbol_map_lock);
if (ksymbols->find(handle) != ksymbols->end())
return ksymbols->at(handle);
HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance();
std::lock_guard<std::mutex> lock(hsasupport_singleton.ksymbol_map_lock);
if (hsasupport_singleton.ksymbols->find(handle) != hsasupport_singleton.ksymbols->end())
return hsasupport_singleton.ksymbols->at(handle);
else
return "Unknown Kernel!";
}
static std::mutex kernel_names_map_lock;
static std::map<std::string, std::vector<uint64_t>>* kernel_names;
static std::atomic<bool> kernel_names_flag{true};
void AddKernelNameWithDispatchID(std::string name, uint64_t id) {
std::lock_guard<std::mutex> lock(kernel_names_map_lock);
if (kernel_names->find(name) == kernel_names->end())
kernel_names->emplace(name, std::vector<uint64_t>());
kernel_names->at(name).push_back(id);
HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance();
std::lock_guard<std::mutex> 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<uint64_t>());
hsasupport_singleton.kernel_names->at(name).push_back(id);
}
std::string GetKernelNameUsingDispatchID(uint64_t given_id) {
std::lock_guard<std::mutex> lock(kernel_names_map_lock);
for (auto kernel_name : (*kernel_names)) {
HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance();
std::lock_guard<std::mutex> 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;
}
@@ -112,34 +113,6 @@ std::string GetKernelNameUsingDispatchID(uint64_t given_id) {
return "Unknown Kernel!";
}
void InitKsymbols() {
if (ksymbols_flag.load(std::memory_order_relaxed)) {
{
std::lock_guard<std::mutex> lock(ksymbol_map_lock);
ksymbols = new std::map<uint64_t, std::string>();
ksymbols_flag.exchange(false, std::memory_order_release);
}
{
std::lock_guard<std::mutex> lock(kernel_names_map_lock);
kernel_names = new std::map<std::string, std::vector<uint64_t>>();
kernel_names_flag.exchange(false, std::memory_order_release);
}
}
}
void FinitKsymbols() {
if (!ksymbols_flag.load(std::memory_order_relaxed)) {
std::lock_guard<std::mutex> lock(ksymbol_map_lock);
ksymbols->clear();
delete ksymbols;
ksymbols_flag.exchange(true, std::memory_order_release);
}
if (!kernel_names_flag.load(std::memory_order_relaxed)) {
std::lock_guard<std::mutex> lock(kernel_names_map_lock);
kernel_names->clear();
delete kernel_names;
kernel_names_flag.exchange(true, std::memory_order_release);
}
}
struct kernel_descriptor_t {
@@ -185,7 +158,8 @@ enum amd_kernel_code_property_t {
static const kernel_descriptor_t* GetKernelCode(uint64_t kernel_object) {
const kernel_descriptor_t* kernel_code = NULL;
hsa_status_t status = hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_query_host_address(
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
hsa_status_t status = hsasupport_singleton.GetHSALoaderApi().hsa_ven_amd_loader_query_host_address(
reinterpret_cast<const void*>(kernel_object), reinterpret_cast<const void**>(&kernel_code));
if (HSA_STATUS_SUCCESS != status) {
kernel_code = reinterpret_cast<kernel_descriptor_t*>(kernel_object);
@@ -193,8 +167,8 @@ static const kernel_descriptor_t* GetKernelCode(uint64_t kernel_object) {
return kernel_code;
}
static uint32_t arch_vgpr_count(Agent::AgentInfo& info, const kernel_descriptor_t& kernel_code) {
const std::string_view& name = info.getName();
static uint32_t arch_vgpr_count(const std::string_view& name, const kernel_descriptor_t& kernel_code) {
std::string info_name(name.data(), name.size());
if (strcmp(name.data(), "gfx90a") == 0 || strncmp(name.data(), "gfx94", 5) == 0)
return (AMD_HSA_BITS_GET(kernel_code.compute_pgm_rsrc3,
@@ -210,23 +184,23 @@ static uint32_t arch_vgpr_count(Agent::AgentInfo& info, const kernel_descriptor_
? 8
: 4);
}
static uint32_t accum_vgpr_count(Agent::AgentInfo& info, const kernel_descriptor_t& kernel_code) {
const std::string_view& name = info.getName();
static uint32_t accum_vgpr_count(const std::string_view& name, const kernel_descriptor_t& kernel_code) {
std::string info_name(name.data(), name.size());
if (strcmp(info_name.c_str(), "gfx908") == 0) return arch_vgpr_count(info, kernel_code);
if (strcmp(info_name.c_str(), "gfx908") == 0) return arch_vgpr_count(name, kernel_code);
if (strcmp(info_name.c_str(), "gfx90a") == 0 || strncmp(info_name.c_str(), "gfx94", 5) == 0)
return (AMD_HSA_BITS_GET(kernel_code.compute_pgm_rsrc1,
AMD_COMPUTE_PGM_RSRC_ONE_GRANULATED_WORKITEM_VGPR_COUNT) +
1) *
8 -
arch_vgpr_count(info, kernel_code);
arch_vgpr_count(name, kernel_code);
return 0;
}
static uint32_t sgpr_count(Agent::AgentInfo& info, const kernel_descriptor_t& kernel_code) {
static uint32_t sgpr_count(const std::string_view& name, const kernel_descriptor_t& kernel_code) {
// GFX10 and later always allocate 128 sgprs.
const std::string_view name = info.getName();
// TODO(srnagara): Recheck the extraction of gfxip from gpu name
const char* name_data = name.data();
const size_t gfxip_label_len = std::min(name.size() - 2, size_t{63});
@@ -259,10 +233,10 @@ rocprofiler_kernel_properties_t set_kernel_properties(hsa_kernel_dispatch_packet
kernel_properties_ptr.workgroup_size = (uint32_t)workgroup_size;
kernel_properties_ptr.lds_size = packet.group_segment_size;
kernel_properties_ptr.scratch_size = packet.private_segment_size;
Agent::AgentInfo agent_info = hsa_support::GetAgentInfo(agent.handle);
kernel_properties_ptr.arch_vgpr_count = arch_vgpr_count(agent_info, *kernel_code);
kernel_properties_ptr.accum_vgpr_count = accum_vgpr_count(agent_info, *kernel_code);
kernel_properties_ptr.sgpr_count = sgpr_count(agent_info, *kernel_code);
HSAAgentInfo agent_info = HSASupport_Singleton::GetInstance().GetHSAAgentInfo(agent.handle);
kernel_properties_ptr.arch_vgpr_count = arch_vgpr_count(agent_info.GetDeviceInfo().getName(), *kernel_code);
kernel_properties_ptr.accum_vgpr_count = accum_vgpr_count(agent_info.GetDeviceInfo().getName(), *kernel_code);
kernel_properties_ptr.sgpr_count = sgpr_count(agent_info.GetDeviceInfo().getName(), *kernel_code);
kernel_properties_ptr.wave_size =
AMD_HSA_BITS_GET(kernel_code->kernel_code_properties,
AMD_KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32)
@@ -275,9 +249,7 @@ rocprofiler_kernel_properties_t set_kernel_properties(hsa_kernel_dispatch_packet
namespace queue {
using rocprofiler::GetROCProfilerSingleton;
hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
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;
pmc_callback_data_t* passed_data = reinterpret_cast<pmc_callback_data_t*>(data);
@@ -330,7 +302,7 @@ void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_sign
rocprofiler_record_counter_value_t{value}});
}
record->counters = counters;
rocprofiler::Session* session = GetROCProfilerSingleton()->GetSession(pending->session_id);
rocprofiler::Session* session = rocprofiler::ROCProfiler_Singleton::GetInstance().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);
@@ -347,7 +319,8 @@ void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_sign
void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent,
att_pending_signal_t& pending) {
Agent::AgentInfo agent_info = hsa_support::GetAgentInfo(gpu_agent.handle);
HSASupport_Singleton& hsasupport_singleton = HSASupport_Singleton::GetInstance();
HSAAgentInfo agent_info = hsasupport_singleton.GetHSAAgentInfo(gpu_agent.handle);
att_trace_callback_data_t data;
hsa_status_t status =
hsa_ven_amd_aqlprofile_iterate_data(pending.profile, attTraceDataCallback, &data);
@@ -373,11 +346,11 @@ void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent
void* buffer = NULL;
if (data_size != 0) {
// Allocate buffer on CPU to copy out trace data
buffer = Packet::AllocateSysMemory(gpu_agent, data_size, &agent_info.cpu_pool);
buffer = Packet::AllocateSysMemory(gpu_agent, data_size, &agent_info.cpu_pool_);
if (buffer == NULL) fatal("Trace data buffer allocation failed");
auto status = rocprofiler::hsa_support::GetCoreApiTable().hsa_memory_copy_fn(buffer, data_ptr,
data_size);
auto status =
hsasupport_singleton.GetCoreApiTable().hsa_memory_copy_fn(buffer, data_ptr, data_size);
if (status != HSA_STATUS_SUCCESS) fatal("Trace data memcopy to host failed");
record->shader_engine_data[se_index].buffer_ptr = buffer;
@@ -392,12 +365,13 @@ void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent
bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
auto queue_info_session = static_cast<queue_info_session_t*>(data);
if (!queue_info_session || !GetROCProfilerSingleton() ||
!GetROCProfilerSingleton()->GetSession(queue_info_session->session_id) ||
!GetROCProfilerSingleton()->GetSession(queue_info_session->session_id)->GetProfiler())
rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton = rocprofiler::ROCProfiler_Singleton::GetInstance();
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
if (!queue_info_session ||
!rocprofiler_singleton.GetSession(queue_info_session->session_id) ||
!rocprofiler_singleton.GetSession(queue_info_session->session_id)->GetProfiler())
return true;
rocprofiler::Session* session =
GetROCProfilerSingleton()->GetSession(queue_info_session->session_id);
rocprofiler::Session* session = rocprofiler_singleton.GetSession(queue_info_session->session_id);
std::lock_guard<std::mutex> lock(session->GetSessionLock());
rocprofiler::profiler::Profiler* profiler = session->GetProfiler();
std::vector<pending_signal_t*> pending_signals = const_cast<std::vector<pending_signal_t*>&>(
@@ -407,10 +381,9 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
for (auto it = pending_signals.begin(); it != pending_signals.end();
it = pending_signals.erase(it)) {
auto& pending = *it;
if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->new_signal))
return true;
if (hsasupport_singleton.GetCoreApiTable().hsa_signal_load_relaxed_fn(pending->new_signal)) return true;
hsa_amd_profiling_dispatch_time_t time;
hsa_support::GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn(
hsasupport_singleton.GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn(
queue_info_session->agent, pending->original_signal, &time);
uint32_t record_count = 1;
bool is_individual_xcc_mode = false;
@@ -440,7 +413,7 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
record.correlation_id = rocprofiler_correlation_id_t{pending->correlation_id};
if (pending->session_id.handle == 0) {
pending->session_id = GetROCProfilerSingleton()->GetCurrentSessionId();
pending->session_id = rocprofiler_singleton.GetCurrentSessionId();
}
if (pending->counters_count > 0) {
if (xcc_id == 0 && pending->context && pending->context->metrics_list.size() > 0 &&
@@ -456,7 +429,7 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
pending->context->metrics_list,
time.end - time.start);
AddRecordCounters(&record, pending);
} else {
}else {
if (session->FindBuffer(pending->buffer_id)) {
Memory::GenericBuffer* buffer = session->GetBuffer(pending->buffer_id);
buffer->AddRecord(record);
@@ -467,13 +440,12 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
// 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 =
rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
(pending->profile->output_buffer.ptr));
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 = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
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;
@@ -483,9 +455,9 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
delete pending->context;
}
if (pending->new_signal.handle)
hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(pending->new_signal);
hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn(pending->new_signal);
if (queue_info_session->interrupt_signal.handle)
hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(queue_info_session->interrupt_signal);
hsasupport_singleton.GetCoreApiTable().hsa_signal_destroy_fn(queue_info_session->interrupt_signal);
}
}
delete queue_info_session;
@@ -496,15 +468,13 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) {
bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
auto queue_info_session = static_cast<queue_info_session_t*>(data);
if (!queue_info_session || !GetROCProfilerSingleton())
rocprofiler::ROCProfiler_Singleton& rocprofiler_singleton = rocprofiler::ROCProfiler_Singleton::GetInstance();
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
if (!queue_info_session ||
!rocprofiler_singleton.GetSession(queue_info_session->session_id) ||
!rocprofiler_singleton.GetSession(queue_info_session->session_id)->GetAttTracer())
return true;
rocprofiler::Session* session =
GetROCProfilerSingleton()->GetSession(queue_info_session->session_id);
if (!session) return true;
std::lock_guard<std::mutex> lock(session->GetSessionLock());
rocprofiler::Session* session = rocprofiler_singleton.GetSession(queue_info_session->session_id);
rocprofiler::att::AttTracer* att_tracer = session->GetAttTracer();
if (!session->GetAttTracer()) return true;
@@ -516,9 +486,8 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
for (auto it = pending_signals.begin(); it != pending_signals.end();
it = pending_signals.erase(it)) {
auto& pending = *it;
if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.new_signal))
return true;
std::lock_guard<std::mutex> lock(session->GetSessionLock());
if (hsasupport_singleton.GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.new_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)queue_info_session->gpu_index};
@@ -540,7 +509,7 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
std::atomic_thread_fence(std::memory_order_release);
if (pending.session_id.handle == 0) {
pending.session_id = GetROCProfilerSingleton()->GetCurrentSessionId();
pending.session_id = rocprofiler_singleton.GetCurrentSessionId();
}
if (session->FindBuffer(pending.buffer_id)) {
Memory::GenericBuffer* buffer = session->GetBuffer(pending.buffer_id);
@@ -549,10 +518,10 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) {
}
codeobj_record::free_capture(record.header.id);
hsa_status_t status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
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);
status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(
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;
@@ -580,20 +549,20 @@ 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(
hsa_status_t status = HSASupport_Singleton::GetInstance().GetAmdExtTable().hsa_amd_signal_async_handler_fn(
signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalHandler, data);
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(
hsa_status_t status = HSASupport_Singleton::GetInstance().GetAmdExtTable().hsa_amd_signal_async_handler_fn(
signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalHandlerATT, data);
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);
HSASupport_Singleton::GetInstance().GetAmdExtTable().hsa_amd_signal_create_fn(1, 0, nullptr, attribute, signal);
CHECK_HSA_STATUS("Error: hsa_amd_signal_create failed", status);
}
@@ -635,17 +604,16 @@ void ResetSessionID(rocprofiler_session_id_t id) { session_id = id; }
void CheckNeededProfileConfigs() {
rocprofiler_session_id_t internal_session_id;
if (GetROCProfilerSingleton())
// Getting Session ID
internal_session_id = GetROCProfilerSingleton()->GetCurrentSessionId();
else
internal_session_id = {0};
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) {
session_id = internal_session_id;
// Getting Counters count from the Session
if (session_id.handle > 0 && GetROCProfilerSingleton()) {
session = GetROCProfilerSingleton()->GetSession(session_id);
if (session_id.handle > 0 ) {
session = rocprofiler_singleton.GetSession(session_id);
if (session && session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION)) {
rocprofiler_filter_id_t filter_id =
session->GetFilterIdWithKind(ROCPROFILER_COUNTERS_COLLECTION);
@@ -690,9 +658,9 @@ std::pair<std::vector<bool>, bool> GetAllowedProfilesList(const void* packets, i
std::vector<bool> can_profile_packet;
bool b_can_profile_anypacket = false;
can_profile_packet.reserve(pkt_count);
std::lock_guard<std::mutex> lock(ksymbol_map_lock);
assert(ksymbols);
rocprofiler::HSASupport_Singleton& hsasupport_singleton = rocprofiler::HSASupport_Singleton::GetInstance();
std::lock_guard<std::mutex> lock(hsasupport_singleton.ksymbol_map_lock);
assert(hsasupport_singleton.ksymbols);
uint32_t current_writer_id = WRITER_ID.load(std::memory_order_relaxed);
@@ -710,7 +678,7 @@ std::pair<std::vector<bool>, bool> GetAllowedProfilesList(const void* packets, i
for (auto id : kernel_profile_dispatch_ids) b_profile_this_object |= id == current_writer_id;
try {
// Can throw
const std::string& kernel_name = ksymbols->at(kdispatch.kernel_object);
const std::string& kernel_name = hsasupport_singleton.ksymbols->at(kdispatch.kernel_object);
// If no filters specified, auto profile this kernel
if (kernel_profile_names.size() == 0 && kernel_profile_dispatch_ids.size() == 0 &&
@@ -739,7 +707,7 @@ ProcessATTParams(
Packet::packet_t& start_packet,
Packet::packet_t& stop_packet,
Queue& queue_info,
Agent::AgentInfo& agentInfo
rocprofiler::HSAAgentInfo& agentInfo
) {
std::vector<hsa_ven_amd_aqlprofile_parameter_t> att_params;
int num_att_counters = 0;
@@ -805,7 +773,7 @@ ProcessATTParams(
* pointer to the packet. This packet is written into the queue by this
* interceptor by invoking the writer function.
*/
void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index, void* data,
void Queue::WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index, void* data,
hsa_amd_queue_intercept_packet_writer writer) {
static const char* env_MAX_ATT_PROFILES = getenv("ROCPROFILER_MAX_ATT_PROFILES");
@@ -882,7 +850,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
rocprofiler_kernel_properties_t kernel_properties =
set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent());
if (session) {
uint64_t record_id = GetROCProfilerSingleton()->GetUniqueRecordId();
uint64_t record_id = rocprofiler::ROCProfiler_Singleton::GetInstance().GetUniqueRecordId();
AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object),
record_id);
if (session_data_count > 0 && profile.second) {
@@ -936,19 +904,21 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
(reinterpret_cast<Packet::packet_t*>(&barrier));
transformed_packets.emplace_back(*pkt);
}
Agent::AgentInfo& agentInfo =
rocprofiler::hsa_support::GetAgentInfo(queue_info.GetGPUAgent().handle);
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.getIndex(), agentInfo.getXccCount()});
new queue_info_session_t{queue_info.GetGPUAgent(), session_id_snapshot, queue_info.GetQueueID(),
writer_id, interrupt_signal, agentInfo.GetDeviceInfo().getGPUId(),
agentInfo.GetDeviceInfo().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());
} else if (session_id_snapshot.handle > 0 && pkt_count > 0 && is_att_collection_mode && session &&
@@ -957,7 +927,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);
rocprofiler::HSAAgentInfo& agentInfo = rocprofiler::HSASupport_Singleton::GetInstance().GetHSAAgentInfo(queue_info.GetGPUAgent().handle);
bool can_profile_anypacket = false;
std::vector<bool> can_profile_packet;
@@ -1021,7 +991,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
// list to be processed by the signal interrupt
rocprofiler_kernel_properties_t kernel_properties =
set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent());
uint64_t record_id = GetROCProfilerSingleton()->GetUniqueRecordId();
uint64_t record_id = rocprofiler::ROCProfiler_Singleton::GetInstance().GetUniqueRecordId();
AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object),
record_id);
@@ -1080,31 +1050,13 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt
/* Write the original packets to the hardware queue if no profiling session
* is active */
writer(packets, pkt_count);
}
}
Queue::Queue(const hsa_agent_t& cpu_agent, const hsa_agent_t& gpu_agent, uint32_t size,
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)
: cpu_agent_(cpu_agent), gpu_agent_(gpu_agent) {
[[maybe_unused]] hsa_status_t status =
hsa_support::GetAmdExtTable().hsa_amd_queue_intercept_create_fn(
gpu_agent, size, type, callback, data, private_segment_size, group_segment_size,
&intercept_queue_);
assert(status == HSA_STATUS_SUCCESS);
status = hsa_support::GetAmdExtTable().hsa_amd_profiling_set_profiler_enabled_fn(intercept_queue_,
true);
assert(status == HSA_STATUS_SUCCESS);
hsa_support::GetAmdExtTable().hsa_amd_queue_intercept_register_fn(intercept_queue_,
WriteInterceptor, this);
assert(status == HSA_STATUS_SUCCESS);
*queue = intercept_queue_;
}
Queue::Queue(const hsa_agent_t cpu_agent, const hsa_agent_t gpu_agent, hsa_queue_t* queue)
: cpu_agent_(cpu_agent), gpu_agent_(gpu_agent), intercept_queue_(queue) { }
Queue::~Queue() {
while (ACTIVE_INTERRUPT_SIGNAL_COUNT.load(std::memory_order_acquire) > 0) {
@@ -1119,15 +1071,8 @@ hsa_agent_t Queue::GetCPUAgent() { return cpu_agent_; }
uint64_t Queue::GetQueueID() { return intercept_queue_->id; }
void InitializePools(hsa_agent_t cpu_agent, Agent::AgentInfo* agent_info) {
Packet::InitializePools(cpu_agent, agent_info);
}
void InitializeGPUPool(hsa_agent_t gpu_agent, Agent::AgentInfo* agent_info) {
Packet::InitializeGPUPool(gpu_agent, agent_info);
}
void CheckPacketReqiurements(std::vector<hsa_agent_t>& gpu_agents) {
Packet::CheckPacketReqiurements(gpu_agents);
}
void CheckPacketReqiurements() {
Packet::CheckPacketReqiurements();}
} // namespace queue
} // namespace rocprofiler
projects/rocprofiler/src/core/hsa/queues/queue.h
+7 -11
View File
@@ -39,8 +39,7 @@
namespace rocprofiler {
void InitKsymbols();
void FinitKsymbols();
void AddKernelName(uint64_t handle, std::string kernel_name);
void RemoveKernelName(uint64_t handle);
void AddKernelNameWithDispatchID(std::string name, uint64_t id);
@@ -52,12 +51,12 @@ namespace queue {
class Queue {
public:
Queue(const hsa_agent_t& cpu_agent, const hsa_agent_t& gpu_agent, uint32_t size,
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(const hsa_agent_t cpu_agent, const hsa_agent_t gpu_agent,
hsa_queue_t* queue);
~Queue();
static void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index,
void* data, hsa_amd_queue_intercept_packet_writer writer);
hsa_queue_t* GetCurrentInterceptQueue();
hsa_agent_t GetGPUAgent();
hsa_agent_t GetCPUAgent();
@@ -69,7 +68,6 @@ class Queue {
std::mutex mutex_;
hsa_agent_t cpu_agent_;
hsa_agent_t gpu_agent_;
hsa_queue_t* original_queue_;
hsa_queue_t* intercept_queue_;
hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
@@ -88,12 +86,10 @@ struct queue_info_session_t {
void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t& pending);
void InitializePools(hsa_agent_t cpu_agent, Agent::AgentInfo* agent_info);
void InitializeGPUPool(hsa_agent_t gpu_agent, Agent::AgentInfo* agent_info);
void CheckPacketReqiurements(std::vector<hsa_agent_t>& gpu_agents);
void ResetSessionID(rocprofiler_session_id_t id = rocprofiler_session_id_t{0});
void CheckPacketReqiurements();
} // namespace queue
} // namespace rocprofiler