AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

adding code object freeze handling, kernel code allocate/device callbacks

Explorar el Código 
Change-Id: I6ec32a0c4e945260cd11f88c1d2a0a7f22d6e10c

documenting comments

Change-Id: I1b7af70d51042c48aef0474d0b6b54275b42c456


[ROCm/rocprofiler commit: 086c29fff1]
Este commit está contenido en:
Evgeny
2018-08-20 20:45:13 -05:00
padre 3b84973737
commit 59bc8da846
Se han modificado 5 ficheros con 161 adiciones y 54 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
projects/rocprofiler/inc/rocprofiler.h
+16 -16
Ver fichero
@@ -465,34 +465,34 @@ hsa_status_t rocprofiler_pool_flush(
// HSA callbacks ID enumeration
enum rocprofiler_hsa_cb_id_t {
ROCPROFILER_HSA_CB_ID_ALLOCATE = 0,
ROCPROFILER_HSA_CB_ID_DEVICE = 1,
ROCPROFILER_HSA_CB_ID_MEMCOPY = 2,
ROCPROFILER_HSA_CB_ID_SUBMIT = 3
ROCPROFILER_HSA_CB_ID_ALLOCATE = 0, // Memory allocate callback
ROCPROFILER_HSA_CB_ID_DEVICE = 1, // Device assign callback
ROCPROFILER_HSA_CB_ID_MEMCOPY = 2, // Memcopy callback
ROCPROFILER_HSA_CB_ID_SUBMIT = 3 // Packet submit callback
};
// HSA callback data type
struct rocprofiler_hsa_callback_data_t {
union {
struct {
const void* addr;
size_t size;
hsa_amd_segment_t segment;
hsa_amd_memory_pool_global_flag_t global_flag;
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
} allocate;
struct {
hsa_device_type_t type;
uint32_t id;
const void* mem;
hsa_device_type_t type; // type of assigned device
uint32_t id; // id of assigned device
const void* ptr; // ptr the device is assigned to
} device;
struct {
const void* dst;
const void* src;
size_t size;
const void* dst; // memcopy dst ptr
const void* src; // memcopy src ptr
size_t size; // memcopy size bytes
} memcopy;
struct {
const void* packet;
const amd_kernel_code_t* kernel_code;
const void* packet; // submitted to GPU packet
const char* kernel_name; // kernel name, not NULL if dispatch
} submit;
};
};
projects/rocprofiler/src/core/hsa_interceptor.h
+113 -28
Ver fichero
@@ -27,6 +27,7 @@ SOFTWARE.
#include <hsa.h>
#include <hsa_ext_amd.h>
#include <hsa_ven_amd_loader.h>
#include <atomic>
#include <mutex>
@@ -55,15 +56,17 @@ SOFTWARE.
do { __callback(__id, &data, __arg); } while (0)
#define ISSUE_HSA_CALLBACK(ID) \
IS_HSA_CALLBACK(ID) { DO_HSA_CALLBACK; }
do { IS_HSA_CALLBACK(ID) { DO_HSA_CALLBACK; } } while(0)
namespace rocprofiler {
extern decltype(hsa_memory_allocate)* hsa_memory_allocate_fn;
extern decltype(hsa_memory_assign_agent)* hsa_memory_assign_agent_fn;
extern decltype(hsa_memory_copy)* hsa_memory_copy_fn;
extern decltype(hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate_fn;
extern decltype(hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_fn;
extern decltype(hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn;
extern decltype(hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn;
extern decltype(hsa_executable_freeze)* hsa_executable_freeze_fn;
class HsaInterceptor {
public:
@@ -73,25 +76,34 @@ class HsaInterceptor {
static void Enable(const bool& enable) { enable_ = enable; }
static void HsaIntercept(HsaApiTable* table) {
fprintf(stderr, "HsaInterceptor ...\n");
if (enable_) {
fprintf(stderr, "HsaInterceptor enabled\n");
// saving original API functions
// Fetching AMD Loader HSA extension API
HSA_RT(hsa_system_get_major_extension_table(
HSA_EXTENSION_AMD_LOADER,
1,
sizeof(hsa_ven_amd_loader_1_01_pfn_t),
&LoaderApiTable));
// Saving original API functions
hsa_memory_allocate_fn = table->core_->hsa_memory_allocate_fn;
hsa_memory_assign_agent_fn = table->core_->hsa_memory_assign_agent_fn;
hsa_memory_copy_fn = table->core_->hsa_memory_copy_fn;
hsa_amd_memory_pool_allocate_fn = table->amd_ext_->hsa_amd_memory_pool_allocate_fn;
hsa_amd_memory_pool_free_fn = table->amd_ext_->hsa_amd_memory_pool_free_fn;
hsa_amd_agents_allow_access_fn = table->amd_ext_->hsa_amd_agents_allow_access_fn;
hsa_amd_memory_async_copy_fn = table->amd_ext_->hsa_amd_memory_async_copy_fn;
// intercepting API
hsa_executable_freeze_fn = table->core_->hsa_executable_freeze_fn;
// Intercepting HSA API
table->core_->hsa_memory_allocate_fn = MemoryAllocate;
table->core_->hsa_memory_assign_agent_fn = MemoryAssignAgent;
table->core_->hsa_memory_copy_fn = MemoryCopy;
table->amd_ext_->hsa_amd_memory_pool_allocate_fn = MemoryPoolAllocate;
table->amd_ext_->hsa_amd_memory_pool_free_fn = MemoryPoolFree;
table->amd_ext_->hsa_amd_agents_allow_access_fn = AgentsAllowAccess;
table->amd_ext_->hsa_amd_memory_async_copy_fn = MemoryAsyncCopy;
table->core_->hsa_executable_freeze_fn = ExecutableFreeze;
}
fprintf(stderr, "HsaInterceptor done\n");
}
static void SetCallbacks(rocprofiler_hsa_callbacks_t callbacks, void* arg) {
@@ -110,7 +122,7 @@ class HsaInterceptor {
HSA_RT(hsa_memory_allocate_fn(region, size, ptr));
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE) {
rocprofiler_hsa_callback_data_t data{};
data.allocate.addr = *ptr;
data.allocate.ptr = *ptr;
data.allocate.size = size;
HSA_RT(hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &data.allocate.segment));
@@ -131,7 +143,7 @@ class HsaInterceptor {
HSA_RT(hsa_memory_assign_agent_fn(ptr, agent, access));
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_DEVICE) {
rocprofiler_hsa_callback_data_t data{};
data.device.mem = ptr;
data.device.ptr = ptr;
HSA_RT(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &data.device.type));
@@ -150,7 +162,7 @@ class HsaInterceptor {
hsa_agent_t agent = *agent_p;
rocprofiler_hsa_callback_data_t data{};
data.device.id = util::HsaRsrcFactory::Instance().GetAgentInfo(agent)->dev_index;
data.device.mem = ptr;
data.device.ptr = ptr;
HSA_RT(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &data.device.type));
@@ -176,7 +188,7 @@ class HsaInterceptor {
// Callback function to get available in the system agents
struct agent_callback_data_t {
hsa_amd_memory_pool_t pool;
void* addr;
void* ptr;
};
static hsa_status_t AgentCallback(hsa_agent_t agent, void* data) {
agent_callback_data_t* callback_data = reinterpret_cast<agent_callback_data_t*>(data);
@@ -184,7 +196,7 @@ class HsaInterceptor {
hsa_amd_memory_pool_access_t value;
HSA_RT(hsa_amd_agent_memory_pool_get_info(agent, callback_data->pool, attribute, &value));
if (value == HSA_AMD_MEMORY_POOL_ACCESS_ALLOWED_BY_DEFAULT) {
DeviceCallback(1, &agent, callback_data->addr);
DeviceCallback(1, &agent, callback_data->ptr);
}
return HSA_STATUS_SUCCESS;
}
@@ -197,27 +209,39 @@ class HsaInterceptor {
{
hsa_status_t status = HSA_STATUS_SUCCESS;
HSA_RT(hsa_amd_memory_pool_allocate_fn(pool, size, flags, ptr));
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE) {
rocprofiler_hsa_callback_data_t data{};
data.allocate.addr = *ptr;
data.allocate.size = size;
HSA_RT(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &data.allocate.segment));
HSA_RT(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &data.allocate.global_flag));
#if 0
HSA_RT(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_ACCESSIBLE_BY_ALL, &data.allocate.global_mem));
#endif
DO_HSA_CALLBACK;
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_DEVICE) {
// Scan the pool assigned devices
agent_callback_data_t callback_data{pool, *ptr};
hsa_iterate_agents(AgentCallback, &callback_data);
if (size != 0) {
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE) {
rocprofiler_hsa_callback_data_t data{};
data.allocate.ptr = *ptr;
data.allocate.size = size;
HSA_RT(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &data.allocate.segment));
HSA_RT(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &data.allocate.global_flag));
DO_HSA_CALLBACK;
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_DEVICE) {
// Scan the pool assigned devices
agent_callback_data_t callback_data{pool, *ptr};
hsa_iterate_agents(AgentCallback, &callback_data);
}
}
}
return status;
}
static hsa_status_t MemoryPoolFree(
void* ptr)
{
hsa_status_t status = HSA_STATUS_SUCCESS;
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE) {
rocprofiler_hsa_callback_data_t data{};
data.allocate.ptr = ptr;
data.allocate.size = 0;
DO_HSA_CALLBACK;
}
HSA_RT(hsa_amd_memory_pool_free_fn(ptr));
return status;
}
static hsa_status_t MemoryCopy(
void *dst,
@@ -257,7 +281,68 @@ class HsaInterceptor {
return status;
}
static hsa_status_t CodeObjectCallback(
hsa_executable_t executable,
hsa_loaded_code_object_t loaded_code_object,
void *)
{
rocprofiler_hsa_callback_data_t data{};
HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE,
&data.allocate.ptr));
HSA_RT(LoaderApiTable.hsa_ven_amd_loader_loaded_code_object_get_info(
loaded_code_object,
HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE,
&data.allocate.size));
// Local GPU memory
// GLOBAL; FLAGS: COARSE GRAINED
data.allocate.segment = HSA_AMD_SEGMENT_GLOBAL;
data.allocate.global_flag = HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED;
ISSUE_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE);
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_DEVICE) {
hsa_amd_pointer_info_t pointer_info{};
uint32_t num_agents = 0;
hsa_agent_t* agents = NULL;
pointer_info.size = sizeof(hsa_amd_pointer_info_t);
HSA_RT(hsa_amd_pointer_info(
const_cast<void*>(data.allocate.ptr),
&pointer_info,
malloc,
&num_agents,
&agents));
DeviceCallback(num_agents, agents, data.allocate.ptr);
}
return HSA_STATUS_SUCCESS;
}
static hsa_status_t ExecutableFreeze(
hsa_executable_t executable,
const char *options)
{
hsa_status_t status = HSA_STATUS_SUCCESS;
HSA_RT(hsa_executable_freeze_fn(executable, options));
IS_HSA_CALLBACK(ROCPROFILER_HSA_CB_ID_ALLOCATE) {
LoaderApiTable.hsa_ven_amd_loader_executable_iterate_loaded_code_objects(
executable,
CodeObjectCallback,
NULL);
}
return status;
}
static bool enable_;
static hsa_ven_amd_loader_1_01_pfn_t LoaderApiTable;
static rocprofiler_hsa_callbacks_t callbacks_;
static arg_t arg_;
static mutex_t mutex_;
projects/rocprofiler/src/core/intercept_queue.h
+2 -2
Ver fichero
@@ -136,8 +136,8 @@ class InterceptQueue {
reinterpret_cast<const hsa_kernel_dispatch_packet_t*>(packet);
rocprofiler_hsa_callback_data_t data{};
data.submit.packet = (void*)packet;
data.submit.kernel_code =
(GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH) ? GetKernelCode(dispatch_packet) : NULL;
data.submit.kernel_name =
(GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH) ? GetKernelName(dispatch_packet) : NULL;
submit_callback_fun_(ROCPROFILER_HSA_CB_ID_SUBMIT, &data, submit_callback_arg_);
}
}
projects/rocprofiler/src/core/rocprofiler.cpp
+3
Ver fichero
@@ -89,9 +89,11 @@ decltype(hsa_memory_allocate)* hsa_memory_allocate_fn;
decltype(hsa_memory_assign_agent)* hsa_memory_assign_agent_fn;
decltype(hsa_memory_copy)* hsa_memory_copy_fn;
decltype(hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate_fn;
decltype(hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_fn;
decltype(hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn;
decltype(hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn;
decltype(hsa_amd_memory_async_copy_rect)* hsa_amd_memory_async_copy_rect_fn;
decltype(hsa_executable_freeze)* hsa_executable_freeze_fn;
::HsaApiTable* kHsaApiTable;
@@ -884,6 +886,7 @@ hsa_status_t rocprofiler_get_time(
bool rocprofiler::HsaInterceptor::enable_ = false;
rocprofiler_hsa_callbacks_t rocprofiler::HsaInterceptor::callbacks_{};
rocprofiler::HsaInterceptor::arg_t rocprofiler::HsaInterceptor::arg_{};
hsa_ven_amd_loader_1_01_pfn_t rocprofiler::HsaInterceptor::LoaderApiTable{};
rocprofiler::HsaInterceptor::mutex_t rocprofiler::HsaInterceptor::mutex_;
extern "C" {
projects/rocprofiler/test/tool/tool.cpp
+27 -8
Ver fichero
@@ -834,15 +834,34 @@ hsa_status_t hsa_unified_callback(
const rocprofiler_hsa_callback_data_t* data,
void* arg)
{
printf("hsa_unified_callback(%d, %p, %p\n", (int)id, data, arg);
printf("hsa_unified_callback(%d, %p, %p):\n", (int)id, data, arg);
if (id == ROCPROFILER_HSA_CB_ID_SUBMIT) {
if (data->submit.kernel_code != NULL) {
printf(" submit:\n");
printf(" kernel_code_entry_byte_offset %lx\n", data->submit.kernel_code->kernel_code_entry_byte_offset);
printf(" kernel_code_prefetch_byte_offset %lx\n", data->submit.kernel_code->kernel_code_prefetch_byte_offset);
printf(" kernel_code_prefetch_byte_size %lx\n", data->submit.kernel_code->kernel_code_prefetch_byte_size);
}
switch (id) {
case ROCPROFILER_HSA_CB_ID_ALLOCATE:
printf(" alloc ptr = %p\n", data->allocate.ptr);
printf(" alloc size = %zu\n", data->allocate.size);
printf(" segment type = 0x%x\n", data->allocate.segment);
printf(" global flag = 0x%x\n", data->allocate.global_flag);
break;
case ROCPROFILER_HSA_CB_ID_DEVICE:
printf(" device type = 0x%x\n", data->device.type);
printf(" device id = %u\n", data->device.id);
printf(" assigned ptr = %p\n", data->device.ptr);
break;
case ROCPROFILER_HSA_CB_ID_MEMCOPY:
printf(" memcopy dst = %p\n", data->memcopy.dst);
printf(" memcopy src = %p\n", data->memcopy.src);
printf(" memcopy size = %zu\n", data->memcopy.size);
break;
case ROCPROFILER_HSA_CB_ID_SUBMIT:
printf(" packet %p\n", data->submit.packet);
if (data->submit.kernel_name != NULL) {
printf(" submit kernel \"%s\"\n", data->submit.kernel_name);
}
break;
default:
printf("Unknown callback id(%u)\n", id);
abort();
}
fflush(stdout);