diff --git a/bin/dform.py b/bin/dform.py index f797d63703..9319460828 100644 --- a/bin/dform.py +++ b/bin/dform.py @@ -1,13 +1,18 @@ #!/usr/bin/python from sqlitedb import SQLiteDB -def post_process_data(db, table_name, outfile = ''): +def gen_message(outfile): + if outfile != '': + print("File '" + outfile + "' is generating") + +def post_process_data(db, table_name, outfile = ''): # db.add_data_column('A', 'DispDurNs', 'INTEGER', 'BeginNs - DispatchNs') # db.add_data_column('A', 'ComplDurNs', 'INTEGER', 'CompleteNs - EndNs') # db.add_data_column('A', 'TotalDurNs', 'INTEGER', 'CompleteNs - DispatchNs') # db.add_data_column(table_name, 'TimeNs', 'INTEGER', 'BeginNs - %d' % start_ns) db.add_data_column(table_name, 'DurationNs', 'INTEGER', 'EndNs - BeginNs') if outfile != '': db.dump_csv(table_name, outfile) + gen_message(outfile) def gen_data_bins(db, outfile): db.execute('create view C as select Name, Calls, TotalDurationNs, TotalDurationNs/Calls as AverageNs, TotalDurationNs*100.0/(select sum(TotalDurationNs) from %s) as Percentage from %s order by TotalDurationNs desc;' % ('B', 'B')); @@ -18,24 +23,29 @@ def gen_table_bins(db, table, outfile, name_var, dur_ns_var): db.execute('create view B as select (%s) as Name, count(%s) as Calls, sum(%s) as TotalDurationNs from %s group by %s' % (name_var, name_var, dur_ns_var, table, name_var)) gen_data_bins(db, outfile) db.execute('DROP VIEW B') + gen_message(outfile) def gen_api_json_trace(db, table, start_us, outfile): - db.execute('create view B as select "Index", Name as name, pid, tid, (BeginNs/1000 - %d) as ts, (DurationNs/1000) as dur from %s order by ts asc;' % (start_us, table)); + db.execute('create view B as select "Index", Name as name, pid, tid, (BeginNs/1000 - %d) as ts, (DurationNs/1000) as dur from %s;' % (start_us, table)); db.dump_json('B', table, outfile) db.execute('DROP VIEW B') + gen_message(outfile) def gen_ext_json_trace(db, table, start_us, outfile): - db.execute('create view B as select Name as name, pid, tid, (BeginNs/1000 - %d) as ts, ((EndNs - BeginNs)/1000) as dur from %s order by ts asc;' % (start_us, table)); + db.execute('create view B as select Name as name, pid, tid, (BeginNs/1000 - %d) as ts, ((EndNs - BeginNs)/1000) as dur from %s;' % (start_us, table)); db.dump_json('B', table, outfile) db.execute('DROP VIEW B') + gen_message(outfile) def gen_ops_json_trace(db, table, base_pid, start_us, outfile): - db.execute('create view B as select "Index", Name as name, ("dev-id" + %d) as pid, tid, (BeginNs/1000 - %d) as ts, (DurationNs/1000) as dur from %s order by ts asc;' % (base_pid, start_us, table)); + db.execute('create view B as select "Index", Name as name, ("dev-id" + %d) as pid, tid, (BeginNs/1000 - %d) as ts, (DurationNs/1000) as dur from %s;' % (base_pid, start_us, table)); db.dump_json('B', table, outfile) db.execute('DROP VIEW B') + gen_message(outfile) def gen_kernel_json_trace(db, table, base_pid, start_us, outfile): - db.execute('create view B as select "Index", KernelName as name, ("gpu-id" + %d) as pid, (0) as tid, (BeginNs/1000 - %d) as ts, (DurationNs/1000) as dur from %s order by ts asc;' % (base_pid, start_us, table)); + db.execute('create view B as select "Index", KernelName as name, ("gpu-id" + %d) as pid, (0) as tid, (BeginNs/1000 - %d) as ts, (DurationNs/1000) as dur from %s;' % (base_pid, start_us, table)); db.dump_json('B', table, outfile) db.execute('DROP VIEW B') + gen_message(outfile) ############################################################################################## diff --git a/bin/rpl_run.sh b/bin/rpl_run.sh index 8a0ef86bc3..e87a6130d3 100755 --- a/bin/rpl_run.sh +++ b/bin/rpl_run.sh @@ -115,7 +115,7 @@ usage() { echo " --list-derived - to print the list of derived metrics with formulas" echo "" echo " -i <.txt|.xml file> - input file" - echo " Input file .txt format, automatically rerun application for every pmc line:" + echo " Input file .txt format, automatically rerun application for every profiling features line:" echo "" echo " # Perf counters group 1" echo " pmc : Wavefronts VALUInsts SALUInsts SFetchInsts FlatVMemInsts LDSInsts FlatLDSInsts GDSInsts VALUUtilization FetchSize" @@ -221,7 +221,6 @@ run() { fi fi mkdir -p "$ROCP_OUTPUT_DIR" - OUTPUT_LIST="$OUTPUT_LIST $ROCP_OUTPUT_DIR/results.txt" fi @@ -229,9 +228,6 @@ run() { if [ "$ROCTX_TRACE" = 1 ] ; then API_TRACE=${API_TRACE}":roctx" fi - if [ "$HSA_TRACE" = 1 ] ; then - API_TRACE=${API_TRACE}":hsa" - fi if [ "$HIP_TRACE" = 1 ] ; then API_TRACE=${API_TRACE}":hip" fi @@ -239,11 +235,12 @@ run() { API_TRACE=${API_TRACE}":sys" fi - if [ -n "$API_TRACE" ] ; then + if [ "$HSA_TRACE" = 1 ] ; then + export ROCTRACER_DOMAIN=$API_TRACE":hsa" + export HSA_TOOLS_LIB="$HSA_TOOLS_LIB $TTLIB_PATH/libtracer_tool.so" + elif [ -n "$API_TRACE" ] ; then export ROCTRACER_DOMAIN=$API_TRACE - if [ "$HSA_TRACE" = 0 ] ; then - OUTPUT_LIST="$ROCP_OUTPUT_DIR/" - fi + OUTPUT_LIST="$ROCP_OUTPUT_DIR/" export HSA_TOOLS_LIB="$TTLIB_PATH/libtracer_tool.so" fi @@ -492,9 +489,7 @@ if [ -n "$csv_output" ] ; then else python $BIN_DIR/tblextr.py $csv_output $OUTPUT_LIST fi - if [ "$?" -eq 0 ] ; then - echo "RPL: '$csv_output' is generated" - else + if [ "$?" -ne 0 ] ; then echo "Data extracting error: $OUTPUT_LIST'" fi fi diff --git a/bin/sqlitedb.py b/bin/sqlitedb.py index e02d413625..02994ac624 100644 --- a/bin/sqlitedb.py +++ b/bin/sqlitedb.py @@ -6,7 +6,7 @@ class SQLiteDB: def __init__(self, file_name): self.connection = sqlite3.connect(file_name) self.tables = {} - self.json_arg_list_enabled = 0 + self.section_index = 0 def __del__(self): self.connection.close() @@ -111,7 +111,8 @@ class SQLiteDB: if not re.search(r'\.json$', file_name): raise Exception('wrong output file type: "' + file_name + '"' ) with open(file_name, mode='a') as fd: - fd.write(',{"args":{"name":"%s"},"ph":"M","pid":%s,"name":"process_name"}\n' %(label, pid)); + fd.write(',{"args":{"name":"%s %s"},"ph":"M","pid":%s,"name":"process_name"}\n' %(self.section_index, label, pid)); + self.section_index += 1 def flow_json(self, base_id, from_pid, from_tid, from_us_list, to_pid, to_us_dict, corr_id_list, start_us, file_name): if not re.search(r'\.json$', file_name): @@ -137,9 +138,9 @@ class SQLiteDB: name_ptrn = re.compile(r'(name|Name)') table_fields = self._get_fields(table_name) - table_raws = self._get_raws_indexed(table_name) + table_raws = self._get_raws(table_name) data_fields = self._get_fields(data_name) - data_raws = self._get_raws_indexed(data_name) + data_raws = self._get_raws(data_name) with open(file_name, mode='a') as fd: table_raws_len = len(table_raws) diff --git a/bin/tblextr.py b/bin/tblextr.py index 895f41b112..8ff735810c 100755 --- a/bin/tblextr.py +++ b/bin/tblextr.py @@ -32,11 +32,11 @@ import dform # SQ_WAVES (4096) # SQ_INSTS_VMEM_RD (36864) -COPY_PID = 0 -OPS_PID = 1 -HSA_PID = 2 -HIP_PID = 3 -EXT_PID = 4 +EXT_PID = 0 +COPY_PID = 1 +HIP_PID = 2 +HSA_PID = 3 +OPS_PID = 4 GPU_BASE_PID = 5 max_gpu_id = -1 START_US = 0 @@ -166,6 +166,8 @@ def dump_csv(file_name): if ind != dispatch_number: fatal("Dispatch #" + ind + " index mismatch (" + dispatch_number + ")\n") val_list = [entry[var] for var in var_list] fd.write(','.join(val_list) + '\n'); + + print("File '" + file_name + "' is generating") ############################################################# # fill kernels DB @@ -337,6 +339,8 @@ def fill_copy_db(table_name, db, indir): else: fatal("async-copy bad record: '" + record + "'") dep_dict[COPY_PID]['to'] = dep_to_us_dict + + return 1 ############################################################# # fill HCC ops DB @@ -434,14 +438,14 @@ else: if ext_trace_found: db.label_json(EXT_PID, "Markers and Ranges", jsonfile) + if hip_trace_found: + db.label_json(HIP_PID, "CPU HIP API", jsonfile) + if hsa_trace_found: db.label_json(HSA_PID, "CPU HSA API", jsonfile) if hsa_activity_found: db.label_json(COPY_PID, "COPY", jsonfile) - if hip_trace_found: - db.label_json(HIP_PID, "CPU HIP API", jsonfile) - if any_trace_found and max_gpu_id >= 0: for ind in range(0, int(max_gpu_id) + 1): db.label_json(int(ind) + int(GPU_BASE_PID), "GPU" + str(ind), jsonfile) diff --git a/bin/txt2xml.sh b/bin/txt2xml.sh index 27bbe8c474..126337ed03 100755 --- a/bin/txt2xml.sh +++ b/bin/txt2xml.sh @@ -64,29 +64,32 @@ parse() { gpu_index=$line fi else - output=$outdir/input${index}.xml - header="# $timestamp '$output' generated with '$0 $*'" - echo $header > $output + found=$(echo $feature | sed -n "/^\(pmc\|sqtt\|hsa\)$/ p") + if [ -n "$found" ] ; then + output=$outdir/input${index}.xml + header="# $timestamp '$output' generated with '$0 $*'" + echo $header > $output - if [ "$feature" == "pmc" ] ; then - line=`echo "$line" | sed -e "s/ /,/g"` - cat >> $output <> $output < EOF - fi + fi - if [ "$feature" == "sqtt" ] ; then - cat >> $output <> $output < EOF - fi + fi - if [ "$feature" == "hsa" ] ; then - cat >> $output <> $output < EOF + fi fi fi diff --git a/src/core/tracker.h b/src/core/tracker.h index c4d619c9f3..2f7cf7abe3 100644 --- a/src/core/tracker.h +++ b/src/core/tracker.h @@ -105,9 +105,10 @@ class Tracker { entry->record = record; // Creating a proxy signal - status = hsa_api_.hsa_signal_create(1, 0, NULL, &(entry->signal)); + const hsa_signal_value_t signal_value = hsa_api_.hsa_signal_load_relaxed(orig); + status = hsa_api_.hsa_signal_create(signal_value, 0, NULL, &(entry->signal)); if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_signal_create"); - status = hsa_api_.hsa_amd_signal_async_handler(entry->signal, HSA_SIGNAL_CONDITION_LT, 1, Handler, entry); + status = hsa_api_.hsa_amd_signal_async_handler(entry->signal, HSA_SIGNAL_CONDITION_LT, signal_value, Handler, entry); if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_signal_async_handler"); // Adding antry to the list @@ -210,9 +211,6 @@ class Tracker { amd_signal_t* prof_signal_ptr = reinterpret_cast(entry->signal.handle); orig_signal_ptr->start_ts = prof_signal_ptr->start_ts; orig_signal_ptr->end_ts = prof_signal_ptr->end_ts; - - const hsa_signal_value_t new_value = hsa_api_.hsa_signal_load_relaxed(orig) - 1; - if (signal_value != new_value) EXC_ABORT(HSA_STATUS_ERROR, "Tracker::Complete bad signal value"); hsa_api_.hsa_signal_store_screlease(orig, signal_value); } }