Adding rocprofilerv2
Change-Id: Ic0cc280ba207d2b8f6ccae1cd4ac3184152fc1ad
[ROCm/rocprofiler commit: 8032adb64f]
This commit is contained in:
Ammar ELWazir committed
1 parent
b23a2f3029
commit
de4abd0d0f
263 files changed
+607729
-307
No files matched your search
@@ -0,0 +1,112 @@
|
||||
/* 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 rocmtools {
|
||||
|
||||
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
|
||||
throw(std::string("Error: Can't find Agent with handle(") + std::to_string(handle) +
|
||||
") in this system");
|
||||
}
|
||||
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);
|
||||
}
|
||||
|
||||
hsa_agent_t GetAgentByIndex(int 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");
|
||||
}
|
||||
|
||||
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 rocmtools
|
||||
@@ -0,0 +1,61 @@
|
||||
/* 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 rocmtools {
|
||||
namespace hsa_support {
|
||||
|
||||
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);
|
||||
|
||||
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 rocmtools
|
||||
|
||||
#endif // SRC_CORE_HSA_HSA_COMMON_H_
|
||||
@@ -0,0 +1,895 @@
|
||||
/* 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_support.h"
|
||||
|
||||
#include <hsa/amd_hsa_signal.h>
|
||||
#include <hsa/hsa.h>
|
||||
#include <hsa/hsa_ven_amd_loader.h>
|
||||
|
||||
#include <mutex>
|
||||
#include <optional>
|
||||
#include <unordered_map>
|
||||
#include <string>
|
||||
#include <atomic>
|
||||
#include <cstdint>
|
||||
#include <cstring>
|
||||
#include <iostream>
|
||||
#include <map>
|
||||
#include <memory>
|
||||
#include <mutex>
|
||||
#include <string>
|
||||
#include <unordered_map>
|
||||
#include <utility>
|
||||
#include <vector>
|
||||
|
||||
#include "src/core/session/tracer/src/correlation_id.h"
|
||||
#include "src/core/session/tracer/src/exception.h"
|
||||
#include "src/core/session/tracer/src/roctracer.h"
|
||||
#include "src/utils/helper.h"
|
||||
|
||||
#include "src/core/hsa/queues/queue.h"
|
||||
#include "src/api/rocmtool.h"
|
||||
|
||||
namespace {
|
||||
|
||||
hsa_status_t hsa_executable_iteration_callback(hsa_executable_t executable, hsa_agent_t agent,
|
||||
hsa_executable_symbol_t symbol, void* args) {
|
||||
hsa_symbol_kind_t type;
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
|
||||
symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &type);
|
||||
if (type == HSA_SYMBOL_KIND_KERNEL) {
|
||||
uint32_t name_length;
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
|
||||
symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &name_length);
|
||||
// TODO(aelwazir): to be removed if the HSA fixed the issue of corrupted
|
||||
// names overflowing the length given
|
||||
if (name_length > 1) {
|
||||
char name[name_length + 1];
|
||||
uint64_t kernel_object;
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
|
||||
symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_symbol_get_info_fn(
|
||||
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel_object);
|
||||
std::string kernel_name = std::string(name).substr(0, name_length);
|
||||
rocmtools::AddKernelName(kernel_object, kernel_name);
|
||||
}
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
std::atomic<int (*)(rocprofiler_tracer_activity_domain_t domain, uint32_t operation_id, void* data)>
|
||||
report_activity;
|
||||
|
||||
bool IsEnabled(rocprofiler_tracer_activity_domain_t domain, uint32_t operation_id) {
|
||||
auto report = report_activity.load(std::memory_order_relaxed);
|
||||
return report && report(domain, operation_id, nullptr) == 0;
|
||||
}
|
||||
|
||||
void ReportActivity(rocprofiler_tracer_activity_domain_t domain, uint32_t operation_id, void* data) {
|
||||
if (auto report = report_activity.load(std::memory_order_relaxed))
|
||||
report(domain, operation_id, data);
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
#include "hsa_prof_str.inline.h"
|
||||
|
||||
namespace roctracer::hsa_support {
|
||||
|
||||
namespace {
|
||||
|
||||
CoreApiTable saved_core_api{};
|
||||
AmdExtTable saved_amd_ext_api{};
|
||||
hsa_ven_amd_loader_1_01_pfn_t hsa_loader_api{};
|
||||
|
||||
struct AgentInfo {
|
||||
uint32_t id;
|
||||
hsa_device_type_t type;
|
||||
};
|
||||
std::unordered_map<decltype(hsa_agent_t::handle), AgentInfo> agent_info_map;
|
||||
|
||||
class Tracker {
|
||||
public:
|
||||
enum { ENTRY_INV = 0, ENTRY_INIT = 1, ENTRY_COMPL = 2 };
|
||||
|
||||
enum entry_type_t {
|
||||
DFLT_ENTRY_TYPE = 0,
|
||||
API_ENTRY_TYPE = 1,
|
||||
COPY_ENTRY_TYPE = 2,
|
||||
KERNEL_ENTRY_TYPE = 3,
|
||||
NUM_ENTRY_TYPE = 4
|
||||
};
|
||||
|
||||
struct entry_t {
|
||||
std::atomic<uint32_t> valid;
|
||||
entry_type_t type;
|
||||
uint64_t correlation_id;
|
||||
roctracer_timestamp_t begin; // begin timestamp, ns
|
||||
roctracer_timestamp_t end; // end timestamp, ns
|
||||
hsa_agent_t agent;
|
||||
uint32_t dev_index;
|
||||
hsa_signal_t orig;
|
||||
hsa_signal_t signal;
|
||||
void (*handler)(const entry_t*);
|
||||
union {
|
||||
struct {
|
||||
} copy;
|
||||
struct {
|
||||
const char* name;
|
||||
hsa_agent_t agent;
|
||||
uint32_t tid;
|
||||
} kernel;
|
||||
};
|
||||
};
|
||||
|
||||
// Add tracker entry
|
||||
inline static void Enable(entry_type_t type, const hsa_agent_t& agent, const hsa_signal_t& signal,
|
||||
entry_t* entry) {
|
||||
hsa_status_t status = HSA_STATUS_ERROR;
|
||||
|
||||
// Creating a new tracker entry
|
||||
entry->type = type;
|
||||
entry->agent = agent;
|
||||
entry->dev_index = 0; // hsa_rsrc->GetAgentInfo(agent)->dev_index;
|
||||
entry->orig = signal;
|
||||
entry->valid.store(ENTRY_INIT, std::memory_order_release);
|
||||
|
||||
// Creating a proxy signal
|
||||
status = rocmtools::hsa_support::GetCoreApiTable().hsa_signal_create_fn(1, 0, NULL,
|
||||
&(entry->signal));
|
||||
if (status != HSA_STATUS_SUCCESS) rocmtools::fatal("hsa_signal_create failed");
|
||||
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_signal_async_handler_fn(
|
||||
entry->signal, HSA_SIGNAL_CONDITION_LT, 1, Handler, entry);
|
||||
if (status != HSA_STATUS_SUCCESS) rocmtools::fatal("hsa_amd_signal_async_handler failed");
|
||||
}
|
||||
|
||||
// Delete tracker entry
|
||||
inline static void Disable(entry_t* entry) {
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(entry->signal);
|
||||
entry->valid.store(ENTRY_INV, std::memory_order_release);
|
||||
}
|
||||
|
||||
private:
|
||||
// Entry completion
|
||||
inline static void Complete(hsa_signal_value_t signal_value, entry_t* entry) {
|
||||
static roctracer_timestamp_t sysclock_period = []() {
|
||||
uint64_t sysclock_hz = 0;
|
||||
hsa_status_t status = rocmtools::hsa_support::GetCoreApiTable().hsa_system_get_info_fn(
|
||||
HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz);
|
||||
if (status != HSA_STATUS_SUCCESS) rocmtools::fatal("hsa_system_get_info failed");
|
||||
return (uint64_t)1000000000 / sysclock_hz;
|
||||
}();
|
||||
|
||||
if (entry->type == COPY_ENTRY_TYPE) {
|
||||
hsa_amd_profiling_async_copy_time_t async_copy_time{};
|
||||
hsa_status_t status =
|
||||
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_profiling_get_async_copy_time_fn(
|
||||
entry->signal, &async_copy_time);
|
||||
if (status != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_amd_profiling_get_async_copy_time failed");
|
||||
entry->begin = async_copy_time.start * sysclock_period;
|
||||
entry->end = async_copy_time.end * sysclock_period;
|
||||
} else {
|
||||
assert(false && "should not reach here");
|
||||
}
|
||||
|
||||
hsa_signal_t orig = entry->orig;
|
||||
hsa_signal_t signal = entry->signal;
|
||||
|
||||
// Releasing completed entry
|
||||
entry->valid.store(ENTRY_COMPL, std::memory_order_release);
|
||||
|
||||
assert(entry->handler != nullptr);
|
||||
entry->handler(entry);
|
||||
|
||||
// Original intercepted signal completion
|
||||
if (orig.handle) {
|
||||
amd_signal_t* orig_signal_ptr = reinterpret_cast<amd_signal_t*>(orig.handle);
|
||||
amd_signal_t* prof_signal_ptr = reinterpret_cast<amd_signal_t*>(signal.handle);
|
||||
orig_signal_ptr->start_ts = prof_signal_ptr->start_ts;
|
||||
orig_signal_ptr->end_ts = prof_signal_ptr->end_ts;
|
||||
|
||||
[[maybe_unused]] const hsa_signal_value_t new_value =
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(orig) - 1;
|
||||
assert(signal_value == new_value && "Tracker::Complete bad signal value");
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_signal_store_screlease_fn(orig, signal_value);
|
||||
}
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_signal_destroy_fn(signal);
|
||||
delete entry;
|
||||
}
|
||||
|
||||
// Handler for packet completion
|
||||
static bool Handler(hsa_signal_value_t signal_value, void* arg) {
|
||||
// Acquire entry
|
||||
entry_t* entry = reinterpret_cast<entry_t*>(arg);
|
||||
while (entry->valid.load(std::memory_order_acquire) != ENTRY_INIT) sched_yield();
|
||||
|
||||
// Complete entry
|
||||
Tracker::Complete(signal_value, entry);
|
||||
return false;
|
||||
}
|
||||
};
|
||||
|
||||
hsa_status_t HSA_API MemoryAllocateIntercept(hsa_region_t region, size_t size, void** ptr) {
|
||||
hsa_status_t status =
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_memory_allocate_fn(region, size, ptr);
|
||||
if (status != HSA_STATUS_SUCCESS) return status;
|
||||
|
||||
if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_ALLOCATE)) {
|
||||
hsa_evt_data_t data{};
|
||||
data.allocate.ptr = *ptr;
|
||||
data.allocate.size = size;
|
||||
if (rocmtools::hsa_support::GetCoreApiTable().hsa_region_get_info_fn(
|
||||
region, HSA_REGION_INFO_SEGMENT, &data.allocate.segment) != HSA_STATUS_SUCCESS ||
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_region_get_info_fn(
|
||||
region, HSA_REGION_INFO_GLOBAL_FLAGS, &data.allocate.global_flag) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_region_get_info failed");
|
||||
|
||||
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_ALLOCATE, &data);
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t MemoryAssignAgentIntercept(void* ptr, hsa_agent_t agent,
|
||||
hsa_access_permission_t access) {
|
||||
hsa_status_t status =
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_memory_assign_agent_fn(ptr, agent, access);
|
||||
if (status != HSA_STATUS_SUCCESS) return status;
|
||||
|
||||
if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_DEVICE)) {
|
||||
hsa_evt_data_t data{};
|
||||
data.device.ptr = ptr;
|
||||
if (rocmtools::hsa_support::GetCoreApiTable().hsa_agent_get_info_fn(
|
||||
agent, HSA_AGENT_INFO_DEVICE, &data.device.type) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_agent_get_info failed");
|
||||
|
||||
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_DEVICE, &data);
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t MemoryCopyIntercept(void* dst, const void* src, size_t size) {
|
||||
hsa_status_t status =
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_memory_copy_fn(dst, src, size);
|
||||
if (status != HSA_STATUS_SUCCESS) return status;
|
||||
|
||||
if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_MEMCOPY)) {
|
||||
hsa_evt_data_t data{};
|
||||
data.memcopy.dst = dst;
|
||||
data.memcopy.src = src;
|
||||
data.memcopy.size = size;
|
||||
|
||||
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_MEMCOPY, &data);
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t MemoryPoolAllocateIntercept(hsa_amd_memory_pool_t pool, size_t size, uint32_t flags,
|
||||
void** ptr) {
|
||||
hsa_status_t status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
|
||||
pool, size, flags, ptr);
|
||||
if (size == 0 || status != HSA_STATUS_SUCCESS) return status;
|
||||
|
||||
if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_ALLOCATE)) {
|
||||
hsa_evt_data_t data{};
|
||||
data.allocate.ptr = *ptr;
|
||||
data.allocate.size = size;
|
||||
|
||||
if (rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
|
||||
pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &data.allocate.segment) != HSA_STATUS_SUCCESS ||
|
||||
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_get_info_fn(
|
||||
pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &data.allocate.global_flag) !=
|
||||
HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_region_get_info failed");
|
||||
|
||||
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_ALLOCATE, &data);
|
||||
}
|
||||
|
||||
if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_DEVICE)) {
|
||||
auto callback_data = std::make_pair(pool, ptr);
|
||||
auto agent_callback = [](hsa_agent_t agent, void* iterate_agent_callback_data) {
|
||||
auto [pool, ptr] = *reinterpret_cast<decltype(callback_data)*>(iterate_agent_callback_data);
|
||||
|
||||
if (hsa_amd_memory_pool_access_t value;
|
||||
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agent_memory_pool_get_info_fn(
|
||||
agent, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &value) != HSA_STATUS_SUCCESS ||
|
||||
value != HSA_AMD_MEMORY_POOL_ACCESS_ALLOWED_BY_DEFAULT)
|
||||
return HSA_STATUS_SUCCESS;
|
||||
|
||||
auto it = agent_info_map.find(agent.handle);
|
||||
if (it == agent_info_map.end()) rocmtools::fatal("agent was not found in the agent_info map");
|
||||
|
||||
hsa_evt_data_t data{};
|
||||
data.device.type = it->second.type;
|
||||
data.device.id = it->second.id;
|
||||
data.device.agent = agent;
|
||||
data.device.ptr = ptr;
|
||||
|
||||
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_DEVICE, &data);
|
||||
return HSA_STATUS_SUCCESS;
|
||||
};
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(agent_callback, &callback_data);
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t MemoryPoolFreeIntercept(void* ptr) {
|
||||
if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_ALLOCATE)) {
|
||||
hsa_evt_data_t data{};
|
||||
data.allocate.ptr = ptr;
|
||||
data.allocate.size = 0;
|
||||
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_ALLOCATE, &data);
|
||||
}
|
||||
|
||||
if (ptr)
|
||||
return rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn(ptr);
|
||||
else
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
// Agent allow access callback 'hsa_amd_agents_allow_access'
|
||||
hsa_status_t AgentsAllowAccessIntercept(uint32_t num_agents, const hsa_agent_t* agents,
|
||||
const uint32_t* flags, const void* ptr) {
|
||||
hsa_status_t status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
|
||||
num_agents, agents, flags, ptr);
|
||||
if (status != HSA_STATUS_SUCCESS) return status;
|
||||
|
||||
if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_DEVICE)) {
|
||||
while (num_agents--) {
|
||||
hsa_agent_t agent = *agents++;
|
||||
auto it = agent_info_map.find(agent.handle);
|
||||
if (it == agent_info_map.end()) rocmtools::fatal("agent was not found in the agent_info map");
|
||||
|
||||
hsa_evt_data_t data{};
|
||||
data.device.type = it->second.type;
|
||||
data.device.id = it->second.id;
|
||||
data.device.agent = agent;
|
||||
data.device.ptr = ptr;
|
||||
|
||||
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_DEVICE, &data);
|
||||
}
|
||||
}
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
struct CodeObjectCallbackArg {
|
||||
activity_rtapi_callback_t callback_fun;
|
||||
void* callback_arg;
|
||||
bool unload;
|
||||
};
|
||||
|
||||
hsa_status_t CodeObjectCallback(hsa_executable_t executable,
|
||||
hsa_loaded_code_object_t loaded_code_object, void* arg) {
|
||||
hsa_evt_data_t data{};
|
||||
|
||||
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
|
||||
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE,
|
||||
&data.codeobj.storage_type) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
|
||||
|
||||
if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_FILE) {
|
||||
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
|
||||
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_FILE,
|
||||
&data.codeobj.storage_file) != HSA_STATUS_SUCCESS ||
|
||||
data.codeobj.storage_file == -1)
|
||||
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
|
||||
data.codeobj.memory_base = data.codeobj.memory_size = 0;
|
||||
} else if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY) {
|
||||
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
|
||||
loaded_code_object,
|
||||
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_BASE,
|
||||
&data.codeobj.memory_base) != HSA_STATUS_SUCCESS ||
|
||||
rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
|
||||
loaded_code_object,
|
||||
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_SIZE,
|
||||
&data.codeobj.memory_size) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
|
||||
data.codeobj.storage_file = -1;
|
||||
} else if (data.codeobj.storage_type == HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_NONE) {
|
||||
return HSA_STATUS_SUCCESS; // FIXME: do we really not care about these
|
||||
// code objects?
|
||||
} else {
|
||||
rocmtools::fatal("unknown code object storage type: %d", data.codeobj.storage_type);
|
||||
}
|
||||
|
||||
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
|
||||
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE,
|
||||
&data.codeobj.load_base) != HSA_STATUS_SUCCESS ||
|
||||
rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
|
||||
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE,
|
||||
&data.codeobj.load_size) != HSA_STATUS_SUCCESS ||
|
||||
rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
|
||||
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA,
|
||||
&data.codeobj.load_delta) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
|
||||
|
||||
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
|
||||
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH,
|
||||
&data.codeobj.uri_length) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
|
||||
|
||||
std::string uri_str(data.codeobj.uri_length, '\0');
|
||||
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
|
||||
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI, uri_str.data()) !=
|
||||
HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
|
||||
|
||||
if (rocmtools::hsa_support::GetHSALoaderApi().hsa_ven_amd_loader_loaded_code_object_get_info(
|
||||
loaded_code_object, HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_AGENT,
|
||||
&data.codeobj.agent) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_ven_amd_loader_loaded_code_object_get_info failed");
|
||||
|
||||
data.codeobj.uri = uri_str.c_str();
|
||||
data.codeobj.unload = *static_cast<bool*>(arg) ? 1 : 0;
|
||||
ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ, &data);
|
||||
|
||||
hsa_executable_iterate_agent_symbols(executable, data.codeobj.agent,
|
||||
hsa_executable_iteration_callback, nullptr);
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t ExecutableFreezeIntercept(hsa_executable_t executable, const char* options) {
|
||||
hsa_status_t status =
|
||||
rocmtools::hsa_support::GetCoreApiTable().hsa_executable_freeze_fn(executable, options);
|
||||
if (status != HSA_STATUS_SUCCESS) return status;
|
||||
|
||||
// if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ)) {
|
||||
bool unload = false;
|
||||
rocmtools::hsa_support::GetHSALoaderApi()
|
||||
.hsa_ven_amd_loader_executable_iterate_loaded_code_objects(executable, CodeObjectCallback,
|
||||
&unload);
|
||||
// }
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
hsa_status_t ExecutableDestroyIntercept(hsa_executable_t executable) {
|
||||
// if (IsEnabled(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ)) {
|
||||
bool unload = true;
|
||||
rocmtools::hsa_support::GetHSALoaderApi()
|
||||
.hsa_ven_amd_loader_executable_iterate_loaded_code_objects(executable, CodeObjectCallback,
|
||||
&unload);
|
||||
// }
|
||||
|
||||
return rocmtools::hsa_support::GetCoreApiTable().hsa_executable_destroy_fn(executable);
|
||||
}
|
||||
|
||||
bool profiling_async_copy_enable = false;
|
||||
|
||||
hsa_status_t ProfilingAsyncCopyEnableIntercept(bool enable) {
|
||||
hsa_status_t status =
|
||||
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_profiling_async_copy_enable_fn(enable);
|
||||
if (status == HSA_STATUS_SUCCESS) profiling_async_copy_enable = enable;
|
||||
return status;
|
||||
}
|
||||
|
||||
void MemoryASyncCopyHandler(const Tracker::entry_t* entry) {
|
||||
activity_record_t record{};
|
||||
record.domain = ACTIVITY_DOMAIN_HSA_OPS;
|
||||
record.op = HSA_OP_ID_COPY;
|
||||
record.begin_ns = entry->begin;
|
||||
record.end_ns = entry->end;
|
||||
record.device_id = 0;
|
||||
record.correlation_id = entry->correlation_id;
|
||||
ReportActivity(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY, &record);
|
||||
}
|
||||
|
||||
hsa_status_t MemoryASyncCopyIntercept(void* dst, hsa_agent_t dst_agent, const void* src,
|
||||
hsa_agent_t src_agent, size_t size, uint32_t num_dep_signals,
|
||||
const hsa_signal_t* dep_signals,
|
||||
hsa_signal_t completion_signal) {
|
||||
bool is_enabled = IsEnabled(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY);
|
||||
|
||||
// FIXME: what happens if the state changes before returning?
|
||||
[[maybe_unused]] hsa_status_t status =
|
||||
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_profiling_async_copy_enable_fn(
|
||||
profiling_async_copy_enable | is_enabled);
|
||||
assert(status == HSA_STATUS_SUCCESS && "hsa_amd_profiling_async_copy_enable failed");
|
||||
|
||||
if (!is_enabled) {
|
||||
return rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_fn(
|
||||
dst, dst_agent, src, src_agent, size, num_dep_signals, dep_signals, completion_signal);
|
||||
}
|
||||
|
||||
Tracker::entry_t* entry = new Tracker::entry_t();
|
||||
entry->handler = MemoryASyncCopyHandler;
|
||||
entry->correlation_id = CorrelationId();
|
||||
Tracker::Enable(Tracker::COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
|
||||
|
||||
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_fn(
|
||||
dst, dst_agent, src, src_agent, size, num_dep_signals, dep_signals, entry->signal);
|
||||
if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry);
|
||||
|
||||
return status;
|
||||
}
|
||||
|
||||
hsa_status_t MemoryASyncCopyRectIntercept(const hsa_pitched_ptr_t* dst,
|
||||
const hsa_dim3_t* dst_offset,
|
||||
const hsa_pitched_ptr_t* src,
|
||||
const hsa_dim3_t* src_offset, const hsa_dim3_t* range,
|
||||
hsa_agent_t copy_agent, hsa_amd_copy_direction_t dir,
|
||||
uint32_t num_dep_signals, const hsa_signal_t* dep_signals,
|
||||
hsa_signal_t completion_signal) {
|
||||
bool is_enabled = IsEnabled(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY);
|
||||
|
||||
// FIXME: what happens if the state changes before returning?
|
||||
[[maybe_unused]] hsa_status_t status =
|
||||
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_profiling_async_copy_enable_fn(
|
||||
profiling_async_copy_enable | is_enabled);
|
||||
assert(status == HSA_STATUS_SUCCESS && "hsa_amd_profiling_async_copy_enable failed");
|
||||
|
||||
if (!is_enabled) {
|
||||
return rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_rect_fn(
|
||||
dst, dst_offset, src, src_offset, range, copy_agent, dir, num_dep_signals, dep_signals,
|
||||
completion_signal);
|
||||
}
|
||||
|
||||
Tracker::entry_t* entry = new Tracker::entry_t();
|
||||
entry->handler = MemoryASyncCopyHandler;
|
||||
entry->correlation_id = CorrelationId();
|
||||
Tracker::Enable(Tracker::COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
|
||||
|
||||
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_async_copy_rect_fn(
|
||||
dst, dst_offset, src, src_offset, range, copy_agent, dir, num_dep_signals, dep_signals,
|
||||
entry->signal);
|
||||
if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry);
|
||||
|
||||
return status;
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
rocprofiler_timestamp_t timestamp_ns() {
|
||||
// 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 = rocmtools::hsa_support::GetCoreApiTable().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)
|
||||
rocmtools::fatal("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)
|
||||
rocmtools::fatal("hsa_system_get_info failed");
|
||||
|
||||
return (uint64_t)1000000000 / sysclock_hz;
|
||||
}();
|
||||
|
||||
return rocprofiler_timestamp_t{sysclock * sysclock_period};
|
||||
}
|
||||
|
||||
void Initialize_roctracer(HsaApiTable* table) {
|
||||
// Save the HSA core api and amd_ext api.
|
||||
saved_core_api = rocmtools::hsa_support::GetCoreApiTable();
|
||||
saved_amd_ext_api = rocmtools::hsa_support::GetAmdExtTable();
|
||||
hsa_loader_api = rocmtools::hsa_support::GetHSALoaderApi();
|
||||
|
||||
// Enumerate the agents.
|
||||
if (rocmtools::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(
|
||||
[](hsa_agent_t agent, void* data) {
|
||||
hsa_support::AgentInfo agent_info;
|
||||
if (rocmtools::hsa_support::GetCoreApiTable().hsa_agent_get_info_fn(
|
||||
agent, HSA_AGENT_INFO_DEVICE, &agent_info.type) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_agent_get_info failed");
|
||||
switch (agent_info.type) {
|
||||
case HSA_DEVICE_TYPE_CPU:
|
||||
static int cpu_agent_count = 0;
|
||||
agent_info.id = cpu_agent_count++;
|
||||
break;
|
||||
case HSA_DEVICE_TYPE_GPU: {
|
||||
uint32_t driver_node_id;
|
||||
if (rocmtools::hsa_support::GetCoreApiTable().hsa_agent_get_info_fn(
|
||||
agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_DRIVER_NODE_ID),
|
||||
&driver_node_id) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_agent_get_info failed");
|
||||
|
||||
agent_info.id = driver_node_id;
|
||||
} break;
|
||||
default:
|
||||
static int other_agent_count = 0;
|
||||
agent_info.id = other_agent_count++;
|
||||
break;
|
||||
}
|
||||
hsa_support::agent_info_map.emplace(agent.handle, agent_info);
|
||||
return HSA_STATUS_SUCCESS;
|
||||
},
|
||||
nullptr) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_iterate_agents failed");
|
||||
}
|
||||
|
||||
const char* GetApiName(uint32_t id) { return detail::GetApiName(id); }
|
||||
|
||||
const char* GetEvtName(uint32_t id) {
|
||||
switch (id) {
|
||||
case HSA_EVT_ID_ALLOCATE:
|
||||
return "ALLOCATE";
|
||||
case HSA_EVT_ID_DEVICE:
|
||||
return "DEVICE";
|
||||
case HSA_EVT_ID_MEMCOPY:
|
||||
return "MEMCOPY";
|
||||
case HSA_EVT_ID_SUBMIT:
|
||||
return "SUBMIT";
|
||||
case HSA_EVT_ID_KSYMBOL:
|
||||
return "KSYMBOL";
|
||||
case HSA_EVT_ID_CODEOBJ:
|
||||
return "CODEOBJ";
|
||||
case HSA_EVT_ID_NUMBER:
|
||||
break;
|
||||
}
|
||||
throw ApiError(ROCTRACER_STATUS_ERROR_INVALID_ARGUMENT, "invalid HSA EVT callback id");
|
||||
}
|
||||
|
||||
const char* GetOpsName(uint32_t id) {
|
||||
switch (id) {
|
||||
case HSA_OP_ID_DISPATCH:
|
||||
return "DISPATCH";
|
||||
case HSA_OP_ID_COPY:
|
||||
return "COPY";
|
||||
case HSA_OP_ID_BARRIER:
|
||||
return "BARRIER";
|
||||
case HSA_OP_ID_RESERVED1:
|
||||
return "PCSAMPLE";
|
||||
}
|
||||
throw ApiError(ROCTRACER_STATUS_ERROR_INVALID_ARGUMENT, "invalid HSA OPS callback id");
|
||||
}
|
||||
|
||||
uint32_t GetApiCode(const char* str) { return detail::GetApiCode(str); }
|
||||
|
||||
void RegisterTracerCallback(int (*function)(activity_domain_t domain, uint32_t operation_id,
|
||||
void* data)) {
|
||||
report_activity.store(function, std::memory_order_relaxed);
|
||||
}
|
||||
|
||||
} // namespace roctracer::hsa_support
|
||||
|
||||
namespace rocmtools {
|
||||
|
||||
namespace hsa_support {
|
||||
|
||||
hsa_agent_t cpu_agent;
|
||||
std::map<uint32_t, std::unique_ptr<queue::Queue>> queues;
|
||||
std::atomic<uint32_t> active_queues{0};
|
||||
|
||||
/**
|
||||
* @brief This function is a queue create interceptor. It intercepts the queue
|
||||
* creation, registers the profiler, and registers a packet write interceptor.
|
||||
* It also creates a Queue Interceptor object to store the
|
||||
* newly created queue information.
|
||||
**/
|
||||
hsa_status_t QueueCreateInterceptor(hsa_agent_t 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) {
|
||||
// TODO(aelwazir): Queue ID
|
||||
static std::mutex qc_mutex;
|
||||
std::lock_guard<std::mutex> lk(qc_mutex);
|
||||
queues.emplace(active_queues.fetch_add(1, std::memory_order_release),
|
||||
std::make_unique<queue::Queue>(cpu_agent, agent, size, type, callback, data,
|
||||
private_segment_size, group_segment_size, queue));
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief This function is a queue destroy interceptor. It intercepts the queue
|
||||
* destroy. It deletes the queue entry from the template storage and calls the
|
||||
* hsa_queue_destroy_fn.
|
||||
**/
|
||||
|
||||
hsa_status_t QueueDestroyInterceptor(hsa_queue_t* hsa_queue) {
|
||||
static std::mutex qd_mutex;
|
||||
std::lock_guard<std::mutex> lk(qd_mutex);
|
||||
ASSERTM(GetCoreApiTable().hsa_queue_destroy_fn(hsa_queue) == HSA_STATUS_SUCCESS,
|
||||
"Queue couldn't be destroyed!");
|
||||
queues.erase(active_queues.fetch_sub(1, std::memory_order_release));
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
void Initialize(HsaApiTable* table) {
|
||||
InitKsymbols();
|
||||
// Save the HSA core api and amd_ext api.
|
||||
SetCoreApiTable(*table->core_);
|
||||
SetAmdExtTable(table->amd_ext_);
|
||||
// Enumerate the agents.
|
||||
if (GetCoreApiTable().hsa_iterate_agents_fn(
|
||||
[](hsa_agent_t agent, void* data) {
|
||||
Agent::AgentInfo agent_info{agent, &GetCoreApiTable()};
|
||||
static int cpu_agent_count = 0;
|
||||
static int gpu_agent_count = 0;
|
||||
static int other_agent_count = 0;
|
||||
switch (agent_info.getType()) {
|
||||
case HSA_DEVICE_TYPE_CPU:
|
||||
agent_info.setIndex(cpu_agent_count++);
|
||||
cpu_agent = agent;
|
||||
rocmtools::queue::InitializePools(cpu_agent);
|
||||
break;
|
||||
case HSA_DEVICE_TYPE_GPU:
|
||||
// XXX FIXME: When multiple ranks are used, each rank's first
|
||||
// logical device always has GPU ID 0, regardless of which
|
||||
// physical device is selected with CUDA_VISIBLE_DEVICES.
|
||||
// Because of this, when merging traces from multiple ranks,
|
||||
// GPU IDs from different processes may overlap.
|
||||
//
|
||||
// The long term solution is to use KFD's gpu_id, which is
|
||||
// stable across APIs and processes, but it isn't currently
|
||||
// exposed by ROCr. We could use the agent's
|
||||
// HSA_AMD_AGENT_INFO_DRIVER_NODE_ID in the meantime, as even
|
||||
// that would be an improvement--it's what legacy roctracer
|
||||
// is currently doing as well as the roctracer compatibility
|
||||
// code earlier in this file.
|
||||
agent_info.setIndex(gpu_agent_count++);
|
||||
break;
|
||||
default:
|
||||
agent_info.setIndex(other_agent_count++);
|
||||
break;
|
||||
}
|
||||
SetAgentInfo(agent.handle, agent_info);
|
||||
return HSA_STATUS_SUCCESS;
|
||||
},
|
||||
nullptr) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_iterate_agents failed");
|
||||
|
||||
SetHSALoaderApi();
|
||||
|
||||
roctracer::hsa_support::Initialize_roctracer(table);
|
||||
|
||||
// Install the Queue intercept
|
||||
table->core_->hsa_queue_create_fn = QueueCreateInterceptor;
|
||||
table->core_->hsa_queue_destroy_fn = QueueDestroyInterceptor;
|
||||
|
||||
// Install the HSA_OPS intercept
|
||||
table->amd_ext_->hsa_amd_memory_async_copy_fn = roctracer::hsa_support::MemoryASyncCopyIntercept;
|
||||
table->amd_ext_->hsa_amd_memory_async_copy_rect_fn =
|
||||
roctracer::hsa_support::MemoryASyncCopyRectIntercept;
|
||||
table->amd_ext_->hsa_amd_profiling_async_copy_enable_fn =
|
||||
roctracer::hsa_support::ProfilingAsyncCopyEnableIntercept;
|
||||
|
||||
// Install the HSA_EVT intercept
|
||||
table->core_->hsa_memory_allocate_fn = roctracer::hsa_support::MemoryAllocateIntercept;
|
||||
table->core_->hsa_memory_assign_agent_fn = roctracer::hsa_support::MemoryAssignAgentIntercept;
|
||||
table->core_->hsa_memory_copy_fn = roctracer::hsa_support::MemoryCopyIntercept;
|
||||
table->amd_ext_->hsa_amd_memory_pool_allocate_fn =
|
||||
roctracer::hsa_support::MemoryPoolAllocateIntercept;
|
||||
table->amd_ext_->hsa_amd_memory_pool_free_fn = roctracer::hsa_support::MemoryPoolFreeIntercept;
|
||||
table->amd_ext_->hsa_amd_agents_allow_access_fn =
|
||||
roctracer::hsa_support::AgentsAllowAccessIntercept;
|
||||
table->core_->hsa_executable_freeze_fn = roctracer::hsa_support::ExecutableFreezeIntercept;
|
||||
table->core_->hsa_executable_destroy_fn = roctracer::hsa_support::ExecutableDestroyIntercept;
|
||||
|
||||
// Install the HSA_API wrappers
|
||||
roctracer::hsa_support::detail::InstallCoreApiWrappers(table->core_);
|
||||
roctracer::hsa_support::detail::InstallAmdExtWrappers(table->amd_ext_);
|
||||
roctracer::hsa_support::detail::InstallImageExtWrappers(table->image_ext_);
|
||||
}
|
||||
|
||||
void Finalize() {
|
||||
while (active_queues.load(std::memory_order_relaxed) != 0) {
|
||||
}
|
||||
|
||||
// FinitKsymbols();
|
||||
ResetMaps();
|
||||
}
|
||||
|
||||
static std::map<uint64_t, rocmtools::MetricsDict*> metricsDicts;
|
||||
|
||||
bool IterateCounters(rocprofiler_counters_info_callback_t counters_info_callback) {
|
||||
if (GetCoreApiTable().hsa_iterate_agents_fn(
|
||||
[](hsa_agent_t agent, void* data) {
|
||||
Agent::AgentInfo agent_info{agent, &GetCoreApiTable()};
|
||||
if (agent_info.getType() == HSA_DEVICE_TYPE_GPU) {
|
||||
metricsDicts.emplace(agent.handle, rocmtools::MetricsDict::Create(&agent_info));
|
||||
}
|
||||
return HSA_STATUS_SUCCESS;
|
||||
},
|
||||
nullptr) != HSA_STATUS_SUCCESS)
|
||||
rocmtools::fatal("hsa_iterate_agents failed");
|
||||
uint32_t gpu_counter = 0;
|
||||
for (auto metricsDictAgent : metricsDicts) {
|
||||
rocmtools::MetricsDict* metricsDict = metricsDictAgent.second;
|
||||
std::string gpu_name = metricsDict->GetAgentName();
|
||||
auto nodes_vec = metricsDict->GetNodes();
|
||||
|
||||
for (auto* node : nodes_vec) {
|
||||
const std::string& name = node->opts["name"];
|
||||
const std::string& descr = node->opts["descr"];
|
||||
const std::string& expr = node->opts["expr"];
|
||||
// Getting the block name
|
||||
const std::string block_name = node->opts["block"];
|
||||
uint32_t block_counters = 0;
|
||||
|
||||
hsa_ven_amd_aqlprofile_id_query_t query;
|
||||
|
||||
if (expr.empty()) {
|
||||
// Querying profile
|
||||
hsa_ven_amd_aqlprofile_profile_t profile = {};
|
||||
profile.agent = hsa_agent_t{metricsDictAgent.first};
|
||||
profile.type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC;
|
||||
|
||||
// Query block id info
|
||||
query = {block_name.c_str(), 0, 0};
|
||||
hsa_status_t status =
|
||||
hsa_ven_amd_aqlprofile_get_info(&profile, HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_ID, &query);
|
||||
if (status != HSA_STATUS_SUCCESS)
|
||||
AQL_EXC_RAISING(HSA_STATUS_ERROR, "get block id info: '" << block_name << "'");
|
||||
|
||||
// Metric object
|
||||
const std::string metric_name = (query.instance_count > 1) ? name + "[0]" : name;
|
||||
const Metric* metric = metricsDict->Get(metric_name);
|
||||
if (metric == NULL) EXC_RAISING(HSA_STATUS_ERROR, "metric '" << name << "' is not found");
|
||||
|
||||
// Process metrics counters
|
||||
const counters_vec_t& counters_vec = metric->GetCounters();
|
||||
if (counters_vec.size() != 1)
|
||||
EXC_RAISING(HSA_STATUS_ERROR, "error: '" << metric->GetName() << "' is not basic");
|
||||
|
||||
// Query block counters number
|
||||
profile.events = &(counters_vec[0]->event);
|
||||
status = hsa_ven_amd_aqlprofile_get_info(
|
||||
&profile, HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_COUNTERS, &block_counters);
|
||||
if (status != HSA_STATUS_SUCCESS) continue;
|
||||
}
|
||||
|
||||
const rocprofiler_counter_info_t counter_info =
|
||||
rocprofiler_counter_info_t{strdup(name.c_str()),
|
||||
strdup(descr.c_str()),
|
||||
expr.empty() ? nullptr : strdup(expr.c_str()),
|
||||
query.instance_count,
|
||||
block_name.c_str(),
|
||||
block_counters};
|
||||
counters_info_callback(counter_info, gpu_name.c_str(), gpu_counter);
|
||||
}
|
||||
gpu_counter++;
|
||||
|
||||
// auto start = metricsDict->Begin();
|
||||
// while (start != metricsDict->End()) {
|
||||
// const xml::Expr* expr = start->second->GetExpr();
|
||||
// std::string expr_str;
|
||||
// if (expr) expr_str = expr->GetStr().c_str();
|
||||
// const rocprofiler_counter_info_t counter_info =
|
||||
// rocprofiler_counter_info_t{start->first.c_str(), "", expr ? expr_str.c_str() : nullptr};
|
||||
// counters_info_callback(counter_info, gpu_name.c_str(), gpu_counter);
|
||||
// start++;
|
||||
// }
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
} // namespace hsa_support
|
||||
} // namespace rocmtools
|
||||
@@ -0,0 +1,136 @@
|
||||
/* 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_SUPPORT_H_
|
||||
#define SRC_CORE_HSA_HSA_SUPPORT_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 <atomic>
|
||||
#include <string>
|
||||
|
||||
#include "hsa_common.h"
|
||||
#include "src/core/hardware/hsa_info.h"
|
||||
|
||||
// HSA EVT data type
|
||||
typedef struct {
|
||||
union {
|
||||
struct {
|
||||
const void* ptr; // allocated area ptr
|
||||
size_t size; // allocated area size, zero size means 'free' callback
|
||||
hsa_amd_segment_t segment; // allocated area's memory segment type
|
||||
hsa_amd_memory_pool_global_flag_t global_flag; // allocated area's memory global flag
|
||||
int is_code; // equal to 1 if code is allocated
|
||||
} allocate;
|
||||
|
||||
struct {
|
||||
hsa_device_type_t type; // type of assigned device
|
||||
uint32_t id; // id of assigned device
|
||||
hsa_agent_t agent; // device HSA agent handle
|
||||
const void* ptr; // ptr the device is assigned to
|
||||
} device;
|
||||
|
||||
struct {
|
||||
const void* dst; // memcopy dst ptr
|
||||
const void* src; // memcopy src ptr
|
||||
size_t size; // memcopy size bytes
|
||||
} memcopy;
|
||||
|
||||
struct {
|
||||
const void* packet; // submitted to GPU packet
|
||||
const char* kernel_name; // kernel name, NULL if not a kernel dispatch packet
|
||||
hsa_queue_t* queue; // HSA queue the packet was submitted to
|
||||
uint32_t device_type; // type of device the packet is submitted to
|
||||
uint32_t device_id; // id of device the packet is submitted to
|
||||
} submit;
|
||||
|
||||
struct {
|
||||
uint64_t object; // kernel symbol object
|
||||
const char* name; // kernel symbol name
|
||||
uint32_t name_length; // kernel symbol name length
|
||||
int unload; // symbol executable destroy
|
||||
} ksymbol;
|
||||
|
||||
struct {
|
||||
uint32_t storage_type; // code object storage type
|
||||
int storage_file; // origin file descriptor
|
||||
uint64_t memory_base; // origin memory base
|
||||
uint64_t memory_size; // origin memory size
|
||||
uint64_t load_base; // code object load base
|
||||
uint64_t load_size; // code object load size
|
||||
uint64_t load_delta; // code object load size
|
||||
uint32_t uri_length; // URI string length (not including the terminating
|
||||
// NUL character)
|
||||
const char* uri; // URI string
|
||||
hsa_agent_t agent; // device HSA agent handle
|
||||
int unload; // unload flag
|
||||
} codeobj;
|
||||
};
|
||||
|
||||
} hsa_evt_data_t;
|
||||
|
||||
namespace rocmtools {
|
||||
|
||||
namespace hsa_support {
|
||||
|
||||
void Initialize(HsaApiTable* Table);
|
||||
hsa_status_t hsa_iterate_agents_cb(hsa_agent_t agent, void *data);
|
||||
void Finalize();
|
||||
|
||||
bool IterateCounters(rocprofiler_counters_info_callback_t counters_info_callback);
|
||||
|
||||
} // namespace hsa_support
|
||||
} // namespace rocmtools
|
||||
|
||||
#include "src/core/session/tracer/src/roctracer.h"
|
||||
|
||||
namespace roctracer::hsa_support {
|
||||
|
||||
struct hsa_trace_data_t {
|
||||
hsa_api_data_t api_data;
|
||||
uint64_t phase_enter_timestamp;
|
||||
uint64_t phase_data;
|
||||
|
||||
void (*phase_enter)(hsa_api_id_t operation_id, hsa_trace_data_t* data);
|
||||
void (*phase_exit)(hsa_api_id_t operation_id, hsa_trace_data_t* data);
|
||||
};
|
||||
|
||||
void Initialize(HsaApiTable* table);
|
||||
void Finalize();
|
||||
|
||||
const char* GetApiName(uint32_t id);
|
||||
const char* GetEvtName(uint32_t id);
|
||||
const char* GetOpsName(uint32_t id);
|
||||
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_
|
||||
@@ -0,0 +1,489 @@
|
||||
/* 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 "packets_generator.h"
|
||||
#include "src/api/rocmtool.h"
|
||||
|
||||
#include <hsa/hsa_ext_amd.h>
|
||||
#include <hsa/hsa_ven_amd_aqlprofile.h>
|
||||
#include <stddef.h>
|
||||
#include <stdint.h>
|
||||
|
||||
#include <algorithm>
|
||||
#include <iostream>
|
||||
#include <map>
|
||||
#include <string>
|
||||
#include <set>
|
||||
#include <vector>
|
||||
#include <utility>
|
||||
|
||||
#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"
|
||||
|
||||
|
||||
#define ASSERTM(exp, msg) assert(((void)msg, exp))
|
||||
|
||||
#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>"); \
|
||||
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,
|
||||
// the function adds an additional requirement that the pool have the
|
||||
// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT property. If kern_arg is false,
|
||||
// pools must NOT have this property.
|
||||
// Upon finding a pool that meets these conditions, HSA_STATUS_INFO_BREAK is
|
||||
// returned. HSA_STATUS_SUCCESS is returned if no errors were encountered, but
|
||||
// no pool was found meeting the requirements. If an error is encountered, we
|
||||
// return that error.
|
||||
static hsa_status_t FindGlobalPool(hsa_amd_memory_pool_t pool, void* data, bool kern_arg) {
|
||||
[[maybe_unused]] hsa_status_t err;
|
||||
hsa_amd_segment_t segment;
|
||||
uint32_t flag;
|
||||
if (nullptr == data) {
|
||||
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
|
||||
}
|
||||
err = rocmtools::hsa_support::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 = rocmtools::hsa_support::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;
|
||||
if ((karg_st == 0 && kern_arg) || (karg_st != 0 && !kern_arg)) {
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
*(reinterpret_cast<hsa_amd_memory_pool_t*>(data)) = pool;
|
||||
return HSA_STATUS_INFO_BREAK;
|
||||
}
|
||||
|
||||
// This is the call-back function for hsa_amd_agent_iterate_memory_pools() that
|
||||
// finds a pool with the properties of HSA_AMD_SEGMENT_GLOBAL and that is NOT
|
||||
// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT
|
||||
hsa_status_t FindStandardPool(hsa_amd_memory_pool_t pool, void* data) {
|
||||
return FindGlobalPool(pool, data, false);
|
||||
}
|
||||
|
||||
// This is the call-back function for hsa_amd_agent_iterate_memory_pools() that
|
||||
// finds a pool with the properties of HSA_AMD_SEGMENT_GLOBAL and that IS
|
||||
// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT
|
||||
hsa_status_t FindKernArgPool(hsa_amd_memory_pool_t pool, void* data) {
|
||||
return FindGlobalPool(pool, data, true);
|
||||
}
|
||||
|
||||
void InitializePools(hsa_agent_t cpu_agent) {
|
||||
hsa_status_t status =
|
||||
rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
|
||||
cpu_agent, FindStandardPool, &command_pool);
|
||||
if ((status != HSA_STATUS_INFO_BREAK)) printf("Error: Command Buffer Pool is not initialized\n");
|
||||
|
||||
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
|
||||
cpu_agent, FindKernArgPool, &output_pool);
|
||||
if ((status != HSA_STATUS_INFO_BREAK)) printf("Error: Output Buffer Pool is not initialized\n");
|
||||
}
|
||||
|
||||
struct block_des_t {
|
||||
uint32_t id;
|
||||
uint32_t index;
|
||||
};
|
||||
|
||||
// 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*>>*
|
||||
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;
|
||||
|
||||
Agent::AgentInfo& agentInfo = rocmtools::hsa_support::GetAgentInfo(gpu_agent.handle);
|
||||
// get the instance of MetricsDict
|
||||
rocmtools::MetricsDict* metricsDict = rocmtools::MetricsDict::Create(&agentInfo);
|
||||
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++;
|
||||
}
|
||||
|
||||
rocmtools::metrics::ExtractMetricEvents(counter_names, gpu_agent, metricsDict, 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");
|
||||
}
|
||||
|
||||
// std::cout << "Max Block Counters: " << max_block_counters << std::endl;
|
||||
|
||||
std::vector<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>>*
|
||||
profiles = new 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))]++;
|
||||
// std::cout << "Block Name: " << event->block_name << " Block Index: " <<
|
||||
// event->block_index
|
||||
// << " Current Count: "
|
||||
// << 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))]
|
||||
// << std::endl;
|
||||
// std::cout << "Counter Taken: " << event->block_index << ", " << event->counter_id << " "
|
||||
// << 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))]
|
||||
// << std::endl;
|
||||
results_list.erase(result);
|
||||
events_list.erase(event);
|
||||
} else {
|
||||
// std::cout << "Counter Left: " << event->block_index << ", " << event->counter_id
|
||||
// << std::endl;
|
||||
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->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);
|
||||
// std::cout << "Metric to be checked from map: " << result.first << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
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()) {
|
||||
// std::cout << metric_counter_name << " for " << metric_name << " is not found!"
|
||||
// << std::endl;
|
||||
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->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;
|
||||
|
||||
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;
|
||||
} 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(
|
||||
command_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, profile->command_buffer.ptr);
|
||||
|
||||
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;
|
||||
}
|
||||
} 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 (!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(
|
||||
output_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));
|
||||
}
|
||||
}
|
||||
} while (events_list.size() > 0);
|
||||
return profiles;
|
||||
}
|
||||
|
||||
// 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
|
||||
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) {
|
||||
hsa_status_t status = HSA_STATUS_SUCCESS;
|
||||
|
||||
// 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");
|
||||
}
|
||||
|
||||
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"
|
||||
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
|
||||
|
||||
// Preparing an Getting the size of the command and output buffers
|
||||
status = hsa_ven_amd_aqlprofile_start(profile, NULL);
|
||||
|
||||
// Preparing and Initializing bool of buffers for command and output buffers
|
||||
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
|
||||
cpu_agent, FindStandardPool, &command_pool);
|
||||
if ((status != HSA_STATUS_INFO_BREAK)) printf("Error: Command Buffer Pool is not initialized\n");
|
||||
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agent_iterate_memory_pools_fn(
|
||||
cpu_agent, FindKernArgPool, &output_pool);
|
||||
if ((status != HSA_STATUS_INFO_BREAK)) printf("Error: Output Buffer Pool is not initialized\n");
|
||||
|
||||
// Allocating Command Buffer
|
||||
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(
|
||||
command_pool, size, 0, reinterpret_cast<void**>(&command_buffer));
|
||||
// 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);
|
||||
}
|
||||
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;
|
||||
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
|
||||
status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
|
||||
output_pool, size, 0, reinterpret_cast<void**>(&output_buffer));
|
||||
// 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;
|
||||
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);
|
||||
status = hsa_ven_amd_aqlprofile_stop(profile, stop_packet);
|
||||
status = hsa_ven_amd_aqlprofile_read(profile, read_packet);
|
||||
|
||||
if (status == HSA_STATUS_ERROR) return nullptr;
|
||||
return profile;
|
||||
}
|
||||
|
||||
} // namespace Packet
|
||||
@@ -0,0 +1,59 @@
|
||||
/* 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_PACKETS_PACKETS_GENERATOR_H_
|
||||
#define SRC_CORE_HSA_PACKETS_PACKETS_GENERATOR_H_
|
||||
|
||||
#include "inc/rocprofiler.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 <string>
|
||||
#include <vector>
|
||||
#include <mutex>
|
||||
|
||||
#include "src/core/counters/metrics/eval_metrics.h"
|
||||
|
||||
namespace Packet {
|
||||
|
||||
static std::mutex pool_lock;
|
||||
|
||||
typedef hsa_ext_amd_aql_pm4_packet_t packet_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);
|
||||
void GetCommandBufferMap(std::map<size_t, uint8_t*>);
|
||||
void GetOutputBufferMap(std::map<size_t, uint8_t*>);
|
||||
void InitializePools(hsa_agent_t cpu_agent);
|
||||
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);
|
||||
hsa_amd_memory_pool_t& GetCommandPool();
|
||||
hsa_amd_memory_pool_t& GetOutputPool();
|
||||
|
||||
|
||||
} // namespace Packet
|
||||
#endif // SRC_CORE_HSA_PACKETS_PACKETS_GENERATOR_H_
|
||||
@@ -0,0 +1,705 @@
|
||||
/* 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 "queue.h"
|
||||
|
||||
#include <atomic>
|
||||
#include <map>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <utility>
|
||||
#include <algorithm>
|
||||
|
||||
#include "rocprofiler.h"
|
||||
#include "src/api/rocmtool.h"
|
||||
#include "src/core/hsa/packets/packets_generator.h"
|
||||
#include "src/core/hsa/hsa_support.h"
|
||||
#include "src/utils/helper.h"
|
||||
|
||||
#define __NR_gettid 186
|
||||
std::mutex sessions_pending_signal_lock;
|
||||
|
||||
namespace rocmtools {
|
||||
|
||||
std::atomic<uint32_t> ACTIVE_INTERRUPT_SIGNAL_COUNT{0};
|
||||
|
||||
uint32_t GetCurrentActiveInterruptSignalsCount() {
|
||||
return ACTIVE_INTERRUPT_SIGNAL_COUNT.load(std::memory_order_relaxed);
|
||||
}
|
||||
|
||||
typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> pmc_callback_data_t;
|
||||
|
||||
static inline bool IsEventMatch(const hsa_ven_amd_aqlprofile_event_t& event1,
|
||||
const hsa_ven_amd_aqlprofile_event_t& event2) {
|
||||
return (event1.block_name == event2.block_name) && (event1.block_index == event2.block_index) &&
|
||||
(event1.counter_id == event2.counter_id);
|
||||
}
|
||||
|
||||
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);
|
||||
}
|
||||
std::string GetKernelNameFromKsymbols(uint64_t handle) {
|
||||
std::lock_guard<std::mutex> lock(ksymbol_map_lock);
|
||||
return ksymbols->at(handle);
|
||||
}
|
||||
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);
|
||||
}
|
||||
}
|
||||
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);
|
||||
}
|
||||
}
|
||||
|
||||
struct kernel_descriptor_t {
|
||||
uint8_t reserved0[16];
|
||||
int64_t kernel_code_entry_byte_offset;
|
||||
uint8_t reserved1[20];
|
||||
uint32_t compute_pgm_rsrc3;
|
||||
uint32_t compute_pgm_rsrc1;
|
||||
uint32_t compute_pgm_rsrc2;
|
||||
uint16_t kernel_code_properties;
|
||||
uint8_t reserved2[6];
|
||||
};
|
||||
// AMD Compute Program Resource Register Three.
|
||||
typedef uint32_t amd_compute_pgm_rsrc_three32_t;
|
||||
enum amd_compute_gfx9_pgm_rsrc_three_t {
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_COMPUTE_PGM_RSRC_THREE_ACCUM_OFFSET, 0, 5),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_COMPUTE_PGM_RSRC_THREE_TG_SPLIT, 16, 1)
|
||||
};
|
||||
enum amd_compute_gfx10_gfx11_pgm_rsrc_three_t {
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_COMPUTE_PGM_RSRC_THREE_SHARED_VGPR_COUNT, 0, 4),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_COMPUTE_PGM_RSRC_THREE_INST_PREF_SIZE, 4, 6),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_COMPUTE_PGM_RSRC_THREE_TRAP_ON_START, 10, 1),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_COMPUTE_PGM_RSRC_THREE_TRAP_ON_END, 11, 1),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_COMPUTE_PGM_RSRC_THREE_IMAGE_OP, 31, 1)
|
||||
};
|
||||
|
||||
// Kernel code properties.
|
||||
enum amd_kernel_code_property_t {
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER, 0,
|
||||
1),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR, 1, 1),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR, 2, 1),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR, 3, 1),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID, 4, 1),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT, 5, 1),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE, 6, 1),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_RESERVED0, 7, 3),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32, 10,
|
||||
1), // GFX10+
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_USES_DYNAMIC_STACK, 11, 1),
|
||||
AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_RESERVED1, 12, 4),
|
||||
};
|
||||
|
||||
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(
|
||||
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);
|
||||
}
|
||||
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();
|
||||
std::string info_name(name.data(), name.size());
|
||||
if (strcmp(name.data(), "gfx90a") == 0 || strcmp(name.data(), "gfx940") == 0)
|
||||
return (AMD_HSA_BITS_GET(kernel_code.compute_pgm_rsrc3,
|
||||
AMD_COMPUTE_PGM_RSRC_THREE_ACCUM_OFFSET) +
|
||||
1) *
|
||||
4;
|
||||
|
||||
return (AMD_HSA_BITS_GET(kernel_code.compute_pgm_rsrc1,
|
||||
AMD_COMPUTE_PGM_RSRC_ONE_GRANULATED_WORKITEM_VGPR_COUNT) +
|
||||
1) *
|
||||
(AMD_HSA_BITS_GET(kernel_code.kernel_code_properties,
|
||||
AMD_KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32)
|
||||
? 8
|
||||
: 4);
|
||||
}
|
||||
static uint32_t accum_vgpr_count(Agent::AgentInfo& info, const kernel_descriptor_t& kernel_code) {
|
||||
const std::string_view& name = info.getName();
|
||||
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(), "gfx90a") == 0 || strcmp(info_name.c_str(), "gfx940") == 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);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
static uint32_t sgpr_count(Agent::AgentInfo& info, 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});
|
||||
if (gfxip_label_len > 0 && strlen(name_data) >= gfxip_label_len) {
|
||||
char gfxip[gfxip_label_len];
|
||||
memcpy(gfxip, name_data, gfxip_label_len);
|
||||
// TODO(srnagara): Check if it is hardcoded
|
||||
if (std::atoi(&gfxip[3]) >= 10) return 128;
|
||||
return (AMD_HSA_BITS_GET(kernel_code.compute_pgm_rsrc1,
|
||||
AMD_COMPUTE_PGM_RSRC_ONE_GRANULATED_WAVEFRONT_SGPR_COUNT) /
|
||||
2 +
|
||||
1) *
|
||||
16;
|
||||
} else {
|
||||
return 0;
|
||||
}
|
||||
}
|
||||
|
||||
rocprofiler_kernel_properties_t set_kernel_properties(hsa_kernel_dispatch_packet_t packet,
|
||||
hsa_agent_t agent) {
|
||||
const uint64_t kernel_object = packet.kernel_object;
|
||||
rocprofiler_kernel_properties_t kernel_properties_ptr = {};
|
||||
const kernel_descriptor_t* kernel_code = GetKernelCode(kernel_object);
|
||||
uint64_t grid_size = packet.grid_size_x * packet.grid_size_y * packet.grid_size_z;
|
||||
if (grid_size > UINT32_MAX) abort();
|
||||
kernel_properties_ptr.grid_size = grid_size;
|
||||
uint64_t workgroup_size =
|
||||
packet.workgroup_size_x * packet.workgroup_size_y * packet.workgroup_size_z;
|
||||
if (workgroup_size > UINT32_MAX) abort();
|
||||
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);
|
||||
kernel_properties_ptr.wave_size =
|
||||
AMD_HSA_BITS_GET(kernel_code->kernel_code_properties,
|
||||
AMD_KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32)
|
||||
? 32
|
||||
: 64;
|
||||
kernel_properties_ptr.signal_handle = packet.completion_signal.handle;
|
||||
|
||||
return kernel_properties_ptr;
|
||||
}
|
||||
|
||||
namespace queue {
|
||||
|
||||
using rocmtools::GetROCMToolObj;
|
||||
|
||||
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);
|
||||
|
||||
pmc_callback_data_t::iterator data_it;
|
||||
if (info_data->sample_id == 0) {
|
||||
passed_data->emplace_back(*info_data);
|
||||
} else {
|
||||
for (data_it = passed_data->begin(); data_it != passed_data->end(); ++data_it) {
|
||||
if (info_type == HSA_VEN_AMD_AQLPROFILE_INFO_PMC_DATA) {
|
||||
if (IsEventMatch(info_data->pmc_data.event, data_it->pmc_data.event)) {
|
||||
data_it->pmc_data.result += info_data->pmc_data.result;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
return status;
|
||||
}
|
||||
|
||||
void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t& pending) {
|
||||
rocmtools::metrics::GetCounterData(pending.profile, pending.context->results_list);
|
||||
rocmtools::metrics::GetMetricsData(pending.context->results_map, pending.context->metrics_list);
|
||||
|
||||
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];
|
||||
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()) {
|
||||
value = it->second->val_double;
|
||||
}
|
||||
counters_vec.emplace_back(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()};
|
||||
}
|
||||
|
||||
// static const size_t MEM_PAGE_BYTES = 0x1000;
|
||||
// static const size_t MEM_PAGE_MASK = MEM_PAGE_BYTES - 1;
|
||||
// static std::mutex begin_signal_lock;
|
||||
|
||||
// bool BeginSignalHandler(hsa_signal_value_t signal_value, void* data) {
|
||||
// std::lock_guard<std::mutex> lock(begin_signal_lock);
|
||||
// auto profiling_context =
|
||||
// static_cast<std::pair<rocmtools::profiling_context_t*, hsa_ven_amd_aqlprofile_profile_t*>*>(
|
||||
// data);
|
||||
// if (!profiling_context->first->begin_completed.load(std::memory_order_relaxed)) {
|
||||
// std::cout << "BeginSignalHandler is called" << std::endl;
|
||||
// hsa_status_t status = HSA_STATUS_ERROR;
|
||||
// size_t size = profiling_context->second->command_buffer.size;
|
||||
// size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
|
||||
// status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
|
||||
// Packet::GetCommandPool(), size, 0,
|
||||
// reinterpret_cast<void**>(&(profiling_context->second->command_buffer.ptr)));
|
||||
|
||||
// // Both the CPU and GPU can access the memory
|
||||
// if (status == HSA_STATUS_SUCCESS) {
|
||||
// hsa_agent_t ag_list[1] = {profiling_context->first->gpu_agent};
|
||||
// status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
|
||||
// 1, ag_list, NULL, profiling_context->second->command_buffer.ptr);
|
||||
|
||||
// if (status != HSA_STATUS_SUCCESS) {
|
||||
// printf("Error: Can't allow access for both agents to Command Buffer\n");
|
||||
// }
|
||||
// } else if (status == HSA_STATUS_ERROR_OUT_OF_RESOURCES) {
|
||||
// printf("Error: Ran out of GPU memory to allocate Command Buffer\n");
|
||||
// } 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);
|
||||
// }
|
||||
|
||||
// status = HSA_STATUS_ERROR;
|
||||
// size = profiling_context->second->output_buffer.size;
|
||||
// size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
|
||||
// status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_allocate_fn(
|
||||
// Packet::GetOutputPool(), size, 0,
|
||||
// reinterpret_cast<void**>(&profiling_context->second->output_buffer.ptr));
|
||||
|
||||
// if (status == HSA_STATUS_ERROR_OUT_OF_RESOURCES) {
|
||||
// printf("Error: Ran out of GPU memory to allocate Output Buffer\n");
|
||||
// }
|
||||
|
||||
// if (status == HSA_STATUS_SUCCESS) {
|
||||
// hsa_agent_t ag_list[1] = {profiling_context->first->gpu_agent};
|
||||
// status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_agents_allow_access_fn(
|
||||
// 1, ag_list, NULL, profiling_context->second->output_buffer.ptr);
|
||||
|
||||
// if (status == HSA_STATUS_SUCCESS) {
|
||||
// memset(profiling_context->second->output_buffer.ptr, 0x0,
|
||||
// profiling_context->second->output_buffer.size);
|
||||
|
||||
// // Creating the start/stop/read packets
|
||||
// status = hsa_ven_amd_aqlprofile_start(profiling_context->second,
|
||||
// profiling_context->first->start_packet);
|
||||
// status = hsa_ven_amd_aqlprofile_stop(profiling_context->second,
|
||||
// profiling_context->first->stop_packet);
|
||||
// status = hsa_ven_amd_aqlprofile_read(profiling_context->second,
|
||||
// profiling_context->first->read_packet);
|
||||
// } else {
|
||||
// printf("Error: Can't allow access for both agents to output Buffer\n");
|
||||
// }
|
||||
// } 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);
|
||||
// }
|
||||
|
||||
// profiling_context->first->begin_completed.exchange(true, std::memory_order_relaxed);
|
||||
// }
|
||||
// return true;
|
||||
// }
|
||||
|
||||
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 || !GetROCMToolObj() ||
|
||||
!GetROCMToolObj()->GetSession(queue_info_session->session_id) ||
|
||||
!GetROCMToolObj()->GetSession(queue_info_session->session_id)->GetProfiler())
|
||||
return true;
|
||||
rocmtools::Session* session = GetROCMToolObj()->GetSession(queue_info_session->session_id);
|
||||
rocmtools::profiler::Profiler* profiler = session->GetProfiler();
|
||||
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;
|
||||
hsa_amd_profiling_dispatch_time_t time;
|
||||
hsa_support::GetAmdExtTable().hsa_amd_profiling_get_dispatch_time_fn(
|
||||
queue_info_session->agent, pending.signal, &time);
|
||||
rocprofiler_record_profiler_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.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) {
|
||||
AddRecordCounters(&record, pending);
|
||||
}
|
||||
record.header = {ROCPROFILER_PROFILER_RECORD,
|
||||
rocprofiler_record_id_t{GetROCMToolObj()->GetUniqueRecordId()}};
|
||||
|
||||
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);
|
||||
record.header.id = rocprofiler_record_id_t{GetROCMToolObj()->GetUniqueRecordId()};
|
||||
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 {
|
||||
buffer->AddRecord(record);
|
||||
}
|
||||
}
|
||||
if (pending.counters_count > 0 && pending.profile && pending.profile->events) {
|
||||
// TODO(aelwazir): we need a better way of distributing events and free them
|
||||
// free(const_cast<hsa_ven_amd_aqlprofile_event_t*>(pending.profile->events));
|
||||
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");
|
||||
}
|
||||
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");
|
||||
}
|
||||
delete pending.profile;
|
||||
}
|
||||
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);
|
||||
}
|
||||
}
|
||||
delete queue_info_session;
|
||||
ACTIVE_INTERRUPT_SIGNAL_COUNT.fetch_sub(1, std::memory_order_relaxed);
|
||||
return false;
|
||||
}
|
||||
|
||||
void CreateBarrierPacket(const hsa_signal_t& packet_completion_signal,
|
||||
std::vector<Packet::packet_t>* transformed_packets) {
|
||||
hsa_barrier_and_packet_t barrier{0};
|
||||
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
|
||||
barrier.dep_signal[0] = packet_completion_signal;
|
||||
void* barrier_ptr = &barrier;
|
||||
transformed_packets->emplace_back(*reinterpret_cast<Packet::packet_t*>(barrier_ptr));
|
||||
}
|
||||
|
||||
void AddVendorSpecificPacket(const Packet::packet_t* packet,
|
||||
std::vector<Packet::packet_t>* transformed_packets,
|
||||
const hsa_signal_t& packet_completion_signal) {
|
||||
transformed_packets->emplace_back(*packet).completion_signal = packet_completion_signal;
|
||||
}
|
||||
|
||||
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");
|
||||
}
|
||||
|
||||
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");
|
||||
}
|
||||
|
||||
template <typename Integral = uint64_t> constexpr Integral bit_mask(int first, int last) {
|
||||
assert(last >= first && "Error: hsa_support::bit_mask -> invalid argument");
|
||||
size_t num_bits = last - first + 1;
|
||||
return ((num_bits >= sizeof(Integral) * 8) ? ~Integral{0}
|
||||
/* num_bits exceed the size of Integral */
|
||||
: ((Integral{1} << num_bits) - 1))
|
||||
<< first;
|
||||
}
|
||||
|
||||
/* Extract bits [last:first] from t. */
|
||||
template <typename Integral> constexpr Integral bit_extract(Integral x, int first, int last) {
|
||||
return (x >> first) & bit_mask<Integral>(0, last - first);
|
||||
}
|
||||
|
||||
std::atomic<uint32_t> WRITER_ID{0};
|
||||
/**
|
||||
* @brief This function is a queue write interceptor. It intercepts the
|
||||
* packet write function. Creates an instance of packet class with the raw
|
||||
* pointer. invoke the populate function of the packet class which returns a
|
||||
* 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,
|
||||
hsa_amd_queue_intercept_packet_writer writer) {
|
||||
const Packet::packet_t* packets_arr = reinterpret_cast<const Packet::packet_t*>(packets);
|
||||
std::vector<Packet::packet_t> transformed_packets;
|
||||
rocprofiler_session_id_t session_id;
|
||||
if (GetROCMToolObj())
|
||||
// Getting Session ID
|
||||
session_id = GetROCMToolObj()->GetCurrentSessionId();
|
||||
else
|
||||
session_id = {0};
|
||||
|
||||
// Counter Names declaration
|
||||
std::vector<std::string> session_data;
|
||||
|
||||
rocprofiler_buffer_id_t buffer_id;
|
||||
|
||||
uint64_t session_data_count = 0;
|
||||
|
||||
bool is_counter_collection_mode = false;
|
||||
bool is_timestamp_collection_mode = false;
|
||||
bool is_pc_sampling_collection_mode = false;
|
||||
|
||||
uint32_t replay_mode_count = 0;
|
||||
|
||||
rocmtools::Session* session = nullptr;
|
||||
|
||||
// Getting Counters count from the Session
|
||||
if (session_id.handle > 0 && GetROCMToolObj()) {
|
||||
session = GetROCMToolObj()->GetSession(session_id);
|
||||
if (session && session->FindFilterWithKind(ROCPROFILER_COUNTERS_COLLECTION)) {
|
||||
rocprofiler_filter_id_t filter_id = session->GetFilterIdWithKind(ROCPROFILER_COUNTERS_COLLECTION);
|
||||
rocmtools::Filter* filter = session->GetFilter(filter_id);
|
||||
session_data = filter->GetCounterData();
|
||||
is_counter_collection_mode = true;
|
||||
session_data_count = session_data.size();
|
||||
buffer_id = filter->GetBufferId();
|
||||
} else if (session && session->FindFilterWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION)) {
|
||||
is_timestamp_collection_mode = true;
|
||||
rocprofiler_filter_id_t filter_id =
|
||||
session->GetFilterIdWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION);
|
||||
rocmtools::Filter* filter = session->GetFilter(filter_id);
|
||||
buffer_id = filter->GetBufferId();
|
||||
} else if (session && session->FindFilterWithKind(ROCPROFILER_PC_SAMPLING_COLLECTION)) {
|
||||
is_pc_sampling_collection_mode = true;
|
||||
}
|
||||
}
|
||||
|
||||
if (session_id.handle > 0 && pkt_count > 0 &&
|
||||
(is_counter_collection_mode || is_timestamp_collection_mode ||
|
||||
is_pc_sampling_collection_mode) &&
|
||||
session) {
|
||||
// Getting Queue Data and Information
|
||||
auto& queue_info = *static_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;
|
||||
|
||||
// Searching accross all the packets given during this write
|
||||
for (size_t i = 0; i < pkt_count; ++i) {
|
||||
auto& original_packet = static_cast<const hsa_barrier_and_packet_t*>(packets)[i];
|
||||
|
||||
// Skip packets other than kernel dispatch packets.
|
||||
if (bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE,
|
||||
HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) !=
|
||||
HSA_PACKET_TYPE_KERNEL_DISPATCH) {
|
||||
transformed_packets.emplace_back(packets_arr[i]);
|
||||
continue;
|
||||
}
|
||||
|
||||
// If counters found in the session
|
||||
if (session_data_count > 0 && is_counter_collection_mode) {
|
||||
// Get the PM4 Packets using packets_generator
|
||||
profiles = Packet::InitializeAqlPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(),
|
||||
session_data);
|
||||
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);
|
||||
|
||||
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) {
|
||||
// hsa_signal_t begin_signal{};
|
||||
// CreateSignal(0, &begin_signal);
|
||||
// hsa_barrier_and_packet_t barrier{0};
|
||||
// barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
|
||||
// CreateSignal(0, &barrier.completion_signal);
|
||||
// barrier.dep_signal[0] = hsa_signal_t{};
|
||||
// Packet::packet_t* __attribute__((__may_alias__)) pkt =
|
||||
// (reinterpret_cast<Packet::packet_t*>(&barrier));
|
||||
// transformed_packets.emplace_back(*pkt);
|
||||
// hsa_status_t status = hsa_support::GetAmdExtTable().hsa_amd_signal_async_handler_fn(
|
||||
// barrier.completion_signal, HSA_SIGNAL_CONDITION_GTE, 1, BeginSignalHandler,
|
||||
// &profiles->at(profile_id));
|
||||
// if (status != HSA_STATUS_SUCCESS)
|
||||
// fatal("hsa_amd_signal_async_handler failed for begin signal");
|
||||
|
||||
// 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);
|
||||
}
|
||||
|
||||
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();
|
||||
|
||||
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) {
|
||||
if (profiles && replay_mode_count > 0) {
|
||||
session->GetProfiler()->AddPendingSignals(
|
||||
writer_id, dispatch_packet.kernel_object, 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, dispatch_packet.kernel_object, 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);
|
||||
} 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);
|
||||
}
|
||||
// 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
|
||||
}
|
||||
/* Write the transformed packets to the hardware queue. */
|
||||
writer(&transformed_packets[0], transformed_packets.size());
|
||||
} else {
|
||||
/* 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() { std::lock_guard<std::mutex> lk(mutex_); }
|
||||
|
||||
hsa_queue_t* Queue::GetCurrentInterceptQueue() { return intercept_queue_; }
|
||||
|
||||
hsa_agent_t Queue::GetGPUAgent() { return gpu_agent_; }
|
||||
|
||||
hsa_agent_t Queue::GetCPUAgent() { return cpu_agent_; }
|
||||
|
||||
uint64_t Queue::GetQueueID() { return intercept_queue_->id; }
|
||||
|
||||
void InitializePools(hsa_agent_t cpu_agent) { Packet::InitializePools(cpu_agent); }
|
||||
|
||||
} // namespace queue
|
||||
} // namespace rocmtools
|
||||
@@ -0,0 +1,91 @@
|
||||
/* 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_QUEUES_QUEUE_H_
|
||||
#define SRC_CORE_HSA_QUEUES_QUEUE_H_
|
||||
|
||||
#include "inc/rocprofiler.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 <hsa/amd_hsa_kernel_code.h>
|
||||
|
||||
#include <map>
|
||||
#include <mutex>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "src/core/session/profiler/profiler.h"
|
||||
|
||||
namespace rocmtools {
|
||||
|
||||
void InitKsymbols();
|
||||
void FinitKsymbols();
|
||||
void AddKernelName(uint64_t handle, std::string kernel_name);
|
||||
std::string GetKernelNameFromKsymbols(uint64_t handle);
|
||||
uint32_t GetCurrentActiveInterruptSignalsCount();
|
||||
|
||||
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() {}
|
||||
|
||||
hsa_queue_t* GetCurrentInterceptQueue();
|
||||
hsa_agent_t GetGPUAgent();
|
||||
hsa_agent_t GetCPUAgent();
|
||||
uint64_t GetQueueID();
|
||||
static void PrintCounters();
|
||||
std::mutex qw_mutex;
|
||||
|
||||
private:
|
||||
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,
|
||||
hsa_ven_amd_aqlprofile_info_data_t* info_data, void* data);
|
||||
};
|
||||
|
||||
struct queue_info_session_t {
|
||||
hsa_agent_t agent;
|
||||
rocprofiler_session_id_t session_id;
|
||||
uint64_t queue_id;
|
||||
uint32_t writer_id;
|
||||
hsa_signal_t interrupt_signal;
|
||||
};
|
||||
|
||||
void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t& pending);
|
||||
|
||||
void InitializePools(hsa_agent_t cpu_agent);
|
||||
|
||||
} // namespace queue
|
||||
} // namespace rocmtools
|
||||
|
||||
#endif // SRC_CORE_HSA_QUEUES_QUEUE_H_
|
||||
Reference in new issue
Block a user