enable rocprofiler kernels trace; fixing signal value for proxy signal; fixing scope options in input txt format;
Change-Id: Ife63c4091e565a2158e292c7acf9718085c709be
This commit is contained in:
Evgeny
+15
-5
@@ -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)
|
||||
##############################################################################################
|
||||
|
||||
+7
-12
@@ -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
|
||||
|
||||
+5
-4
@@ -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)
|
||||
|
||||
+12
-8
@@ -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)
|
||||
|
||||
+15
-12
@@ -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 <<EOF
|
||||
if [ "$feature" == "pmc" ] ; then
|
||||
line=`echo "$line" | sed -e "s/ /,/g"`
|
||||
cat >> $output <<EOF
|
||||
<metric range="$range" kernel="$kernel" gpu_index="$gpu_index"></metric>
|
||||
<metric name=$line ></metric>
|
||||
EOF
|
||||
fi
|
||||
fi
|
||||
|
||||
if [ "$feature" == "sqtt" ] ; then
|
||||
cat >> $output <<EOF
|
||||
if [ "$feature" == "sqtt" ] ; then
|
||||
cat >> $output <<EOF
|
||||
<metric range="$range" kernel="$kernel" gpu_index="$gpu_index"></metric>
|
||||
<trace name="SQTT"><parameters $line ></parameters></trace>
|
||||
EOF
|
||||
fi
|
||||
fi
|
||||
|
||||
if [ "$feature" == "hsa" ] ; then
|
||||
cat >> $output <<EOF
|
||||
if [ "$feature" == "hsa" ] ; then
|
||||
cat >> $output <<EOF
|
||||
<trace name="HSA"><parameters api="$line"></parameters></trace>
|
||||
EOF
|
||||
fi
|
||||
fi
|
||||
fi
|
||||
|
||||
|
||||
@@ -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<amd_signal_t*>(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);
|
||||
}
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user