adding kernel properties

Change-Id: I5d8672400357d5ef5cbfb5023d82f712fdaf7057


[ROCm/rocprofiler commit: 64e0325a0d]

projects/rocprofiler/inc/rocprofiler.h

@@ -43,6 +43,7 @@ THE SOFTWARE.
 #define INC_ROCPROFILER_H_
 
 #include <hsa.h>
+#include <amd_hsa_kernel_code.h>
 #include <hsa_ven_amd_aqlprofile.h>
 #include <stdint.h>
 
@@ -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;

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<hsa_packet_type_t>((*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<amd_kernel_code_t*>(dispatch_packet->kernel_object);
     }
-    return kernel_code->runtime_loader_kernel_symbol;
+    return kernel_code;
   }
 
   static const char* GetKernelName(const uint64_t kernel_symbol) {

projects/rocprofiler/test/tool/tool.cpp

@@ -69,6 +69,17 @@ struct callbacks_data_t {
   std::vector<uint32_t>* 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<callbacks_data_t*>(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;