diff --git a/bin/rpl_run.sh b/bin/rpl_run.sh index eb432a9685..ec4a2b0418 100755 --- a/bin/rpl_run.sh +++ b/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 diff --git a/bin/tblextr.py b/bin/tblextr.py index 9faa5522b2..b5672c54e6 100755 --- a/bin/tblextr.py +++ b/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) diff --git a/inc/rocprofiler.h b/inc/rocprofiler.h index fac482eaf9..b356b132bc 100644 --- a/inc/rocprofiler.h +++ b/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; diff --git a/src/core/intercept_queue.h b/src/core/intercept_queue.h index 0501bd6fcf..a52d8c1d6b 100644 --- a/src/core/intercept_queue.h +++ b/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 diff --git a/test/app/intercept_test.cpp b/test/app/intercept_test.cpp index c2905d1eea..e62bf6cee9 100644 --- a/test/app/intercept_test.cpp +++ b/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, diff --git a/test/tool/tool.cpp b/test/tool/tool.cpp index f0a94063aa..77c7a97a2a 100644 --- a/test/tool/tool.cpp +++ b/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;