From 2dff8e83f4163e9ae094fdd2a439f82f9fe93baf Mon Sep 17 00:00:00 2001 From: Evgeny Date: Fri, 21 Dec 2018 18:02:26 -0600 Subject: [PATCH] adding kernel properties Change-Id: I5d8672400357d5ef5cbfb5023d82f712fdaf7057 [ROCm/rocprofiler commit: 64e0325a0dc0d97bdd87081aa03545de7a2918ab] --- projects/rocprofiler/inc/rocprofiler.h | 2 + .../rocprofiler/src/core/intercept_queue.h | 8 ++-- projects/rocprofiler/test/tool/tool.cpp | 38 ++++++++++++++++++- 3 files changed, 44 insertions(+), 4 deletions(-) diff --git a/projects/rocprofiler/inc/rocprofiler.h b/projects/rocprofiler/inc/rocprofiler.h index 4a25919e6a..a26071d392 100644 --- a/projects/rocprofiler/inc/rocprofiler.h +++ b/projects/rocprofiler/inc/rocprofiler.h @@ -43,6 +43,7 @@ THE SOFTWARE. #define INC_ROCPROFILER_H_ #include +#include #include #include @@ -221,6 +222,7 @@ typedef struct { const hsa_kernel_dispatch_packet_t* packet; // HSA dispatch packet const char* kernel_name; // Kernel name uint64_t kernel_object; // Kernel object pointer + const amd_kernel_code_t* kernel_code; // Kernel code pointer int64_t thread_id; // Thread id const rocprofiler_dispatch_record_t* record; // Dispatch record } rocprofiler_callback_data_t; diff --git a/projects/rocprofiler/src/core/intercept_queue.h b/projects/rocprofiler/src/core/intercept_queue.h index c36f623cbb..bc63a3bcc3 100644 --- a/projects/rocprofiler/src/core/intercept_queue.h +++ b/projects/rocprofiler/src/core/intercept_queue.h @@ -147,7 +147,8 @@ class InterceptQueue { } // Prepareing dispatch callback data - uint64_t kernel_symbol = GetKernelSymbol(dispatch_packet); + const amd_kernel_code_t* kernel_code = GetKernelCode(dispatch_packet); + const uint64_t kernel_symbol = kernel_code->runtime_loader_kernel_symbol; const char* kernel_name = GetKernelName(kernel_symbol); rocprofiler_callback_data_t data = {obj->agent_info_->dev_id, obj->agent_info_->dev_index, @@ -157,6 +158,7 @@ class InterceptQueue { dispatch_packet, kernel_name, kernel_symbol, + kernel_code, syscall(__NR_gettid), (tracker_entry) ? tracker_entry->record : NULL}; @@ -239,7 +241,7 @@ class InterceptQueue { return static_cast((*header >> HSA_PACKET_HEADER_TYPE) & header_type_mask); } - static uint64_t GetKernelSymbol(const hsa_kernel_dispatch_packet_t* dispatch_packet) { + static const amd_kernel_code_t* GetKernelCode(const hsa_kernel_dispatch_packet_t* dispatch_packet) { const amd_kernel_code_t* kernel_code = NULL; hsa_status_t status = util::HsaRsrcFactory::Instance().LoaderApi()->hsa_ven_amd_loader_query_host_address( @@ -248,7 +250,7 @@ class InterceptQueue { if (HSA_STATUS_SUCCESS != status) { kernel_code = reinterpret_cast(dispatch_packet->kernel_object); } - return kernel_code->runtime_loader_kernel_symbol; + return kernel_code; } static const char* GetKernelName(const uint64_t kernel_symbol) { diff --git a/projects/rocprofiler/test/tool/tool.cpp b/projects/rocprofiler/test/tool/tool.cpp index 0db9e290f7..2ed5d32aed 100644 --- a/projects/rocprofiler/test/tool/tool.cpp +++ b/projects/rocprofiler/test/tool/tool.cpp @@ -69,6 +69,17 @@ struct callbacks_data_t { std::vector* range; }; +// kernel properties structure +struct kernel_properties_t { + uint32_t grid_size; + uint32_t workgroup_size; + uint32_t lds_size; + uint32_t scratch_size; + uint32_t vgpr_count; + uint32_t sgpr_count; + uint32_t fbarrier_count; +}; + // Context stored entry type struct context_entry_t { bool valid; @@ -79,6 +90,7 @@ struct context_entry_t { rocprofiler_feature_t* features; unsigned feature_count; rocprofiler_callback_data_t data; + kernel_properties_t kernel_properties; FILE* file_handle; }; @@ -413,11 +425,19 @@ bool dump_context_entry(context_entry_t* entry) { FILE* file_handle = entry->file_handle; const std::string nik_name = (to_truncate_names == 0) ? entry->data.kernel_name : filtr_kernel_name(entry->data.kernel_name); - fprintf(file_handle, "dispatch[%u], gpu-id(%u), queue-id(%u), queue-index(%lu), kernel-name(\"%s\")", + fprintf(file_handle, "dispatch[%u], gpu-id(%u), queue-id(%u), queue-index(%lu), tid(%lu), grd(%u), wgr(%u), lds(%u), scr(%u), vgpr(%u), sgpr(%u), fbar(%u), kernel-name(\"%s\")", index, HsaRsrcFactory::Instance().GetAgentInfo(entry->agent)->dev_index, entry->data.queue_id, entry->data.queue_index, + entry->data.thread_id, + entry->kernel_properties.grid_size, + entry->kernel_properties.workgroup_size, + entry->kernel_properties.lds_size, + entry->kernel_properties.scratch_size, + entry->kernel_properties.vgpr_count, + entry->kernel_properties.sgpr_count, + entry->kernel_properties.fbarrier_count, nik_name.c_str()); if (record) fprintf(file_handle, ", time(%lu,%lu,%lu,%lu)", record->dispatch, @@ -563,6 +583,8 @@ bool check_filter(const rocprofiler_callback_data_t* callback_data, const callba hsa_status_t dispatch_callback(const rocprofiler_callback_data_t* callback_data, void* user_data, rocprofiler_group_t* group) { // Passed tool data + const hsa_kernel_dispatch_packet_t* packet = callback_data->packet; + const amd_kernel_code_t* kernel_code = callback_data->kernel_code; callbacks_data_t* tool_data = reinterpret_cast(user_data); // HSA status hsa_status_t status = HSA_STATUS_ERROR; @@ -578,6 +600,20 @@ hsa_status_t dispatch_callback(const rocprofiler_callback_data_t* callback_data, rocprofiler_t* context = NULL; // Context entry context_entry_t* entry = alloc_context_entry(); + // kernel properties + kernel_properties_t* kernel_properties_ptr = &(entry->kernel_properties); + 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 = (uint32_t)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; + kernel_properties_ptr->vgpr_count = kernel_code->reserved_vgpr_count; + kernel_properties_ptr->sgpr_count = kernel_code->reserved_sgpr_count; + kernel_properties_ptr->fbarrier_count = kernel_code->workgroup_fbarrier_count; + // context properties rocprofiler_properties_t properties{}; properties.handler = (result_prefix != NULL) ? context_handler : NULL;