adding pid for kernel results to support multi-process profiling

Change-Id: I283228a4b4145599c5e637dd6faa771b9f4b6345

bin/rpl_run.sh

@@ -543,9 +543,9 @@ for name in $input_list; do
 done
 
 if [ -n "$csv_output" ] ; then
+  merge_output $OUTPUT_LIST
   if [ "$GEN_STATS" = "1" ] ; then
     db_output=$(echo $csv_output | sed "s/\.csv/.db/")
-    merge_output $OUTPUT_LIST
     $ROCP_PYTHON_VERSION $BIN_DIR/tblextr.py $db_output $OUTPUT_LIST
   else
     $ROCP_PYTHON_VERSION $BIN_DIR/tblextr.py $csv_output $OUTPUT_LIST

bin/tblextr.py

@@ -99,7 +99,7 @@ def parse_res(infile):
       if not dispatch_number in var_table: fatal("Error: dispatch number not found '" + str(dispatch_number) + "'")
       var = m.group(1)
       val = m.group(2)
-      var_table[dispatch_number][m.group(1)] = m.group(2)
+      var_table[dispatch_number][var] = val
       if not var in var_list: var_list.append(var)
 
     m = beg_pattern.match(record)

inc/rocprofiler.h

@@ -249,7 +249,7 @@ typedef struct {
   const char* kernel_name;                             // Kernel name
   uint64_t kernel_object;                              // Kernel object address
   const amd_kernel_code_t* kernel_code;                // Kernel code pointer
-  int64_t thread_id;                                   // Thread id
+  uint32_t thread_id;                                   // Thread id
   const rocprofiler_dispatch_record_t* record;         // Dispatch record
 } rocprofiler_callback_data_t;
 

src/core/intercept_queue.h

@@ -196,7 +196,7 @@ class InterceptQueue {
                                             kernel_name,
                                             kernel_object,
                                             kernel_code,
-                                            syscall(__NR_gettid),
+                                            (uint32_t)syscall(__NR_gettid),
                                             (tracker_entry) ? tracker_entry->record : NULL};
 
         // Calling dispatch callback

test/app/intercept_test.cpp

@@ -91,7 +91,7 @@ void dump_context_entry(context_entry_t* entry, rocprofiler_feature_t* features,
   const rocprofiler_dispatch_record_t* record = entry->data.record;
 
   fflush(stdout);
-  fprintf(stdout, "kernel symbol(0x%lx) name(\"%s\") tid(%ld) queue-id(%u) gpu-id(%u) ",
+  fprintf(stdout, "kernel symbol(0x%lx) name(\"%s\") tid(%u) queue-id(%u) gpu-id(%u) ",
     entry->data.kernel_object,
     kernel_name.c_str(),
     entry->data.thread_id,

test/tool/tool.cpp

@@ -149,6 +149,8 @@ bool is_spm_trace = false;
 static inline uint32_t GetPid() { return syscall(__NR_getpid); }
 static inline uint32_t GetTid() { return syscall(__NR_gettid); }
 
+uint32_t my_pid = GetPid();
+
 // Error handler
 void fatal(const std::string msg) {
   fflush(stdout);
@@ -483,11 +485,12 @@ bool dump_context_entry(context_entry_t* entry) {
   const std::string nik_name = (to_truncate_names == 0) ? entry->data.kernel_name : filtr_kernel_name(entry->data.kernel_name);
   const AgentInfo* agent_info = HsaRsrcFactory::Instance().GetAgentInfo(entry->agent);
 
-  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), sig(0x%lx), kernel-name(\"%s\")",
+  fprintf(file_handle, "dispatch[%u], gpu-id(%u), queue-id(%u), queue-index(%lu), pid(%u), tid(%u), grd(%u), wgr(%u), lds(%u), scr(%u), vgpr(%u), sgpr(%u), fbar(%u), sig(0x%lx), kernel-name(\"%s\")",
     index,
     agent_info->dev_index,
     entry->data.queue_id,
     entry->data.queue_index,
+    my_pid,
     entry->data.thread_id,
     entry->kernel_properties.grid_size,
     entry->kernel_properties.workgroup_size,
@@ -1004,7 +1007,7 @@ extern "C" PUBLIC_API void OnLoadToolProp(rocprofiler_settings_t* settings)
       abort();
     }
     std::ostringstream oss;
-    oss << result_prefix << "/results.txt";
+    oss << result_prefix << "/" << GetPid() << "_results.txt";
     result_file_handle = fopen(oss.str().c_str(), "w");
     if (result_file_handle == NULL) {
       std::ostringstream errmsg;