SWDEV-318682: Using roctx to rename long kernel names to shorter ones.

To enable this feature use the --roctx-rename rocprof option. This implementation records all messages received in roctxPush calls and use them to replace corresponding kernel names. Tested with the following HIP program: \#include <hip/hip_runtime.h> \#include <roctracer/roctx.h> __global__ void ThisIsALongKernelName () { } int main (int argc, char* argv[]) { hipSetDevice (0); // Not in a roctx range. ThisIsALongKernelName<<<1, 1>>> (); roctxRangePush ("A"); // In a simple first level roctx range. ThisIsALongKernelName<<<1, 1>>> (); roctxRangePop (); roctxRangePush ("B"); roctxRangePush ("C"); // In a nested roctx range. ThisIsALongKernelName<<<1, 1>>> (); roctxRangePop (); roctxRangePop (); roctxRangePush ("D"); roctxRangePush ("E"); roctxRangePop (); // In a first level roctx range, but after a nested range. ThisIsALongKernelName<<<1, 1>>> (); roctxRangePop (); hipDeviceSynchronize (); return 0; } Change-Id: I629312234468daff8b017caa5cb0773707d98cce [ROCm/rocprofiler commit: 1078a088e9]
2022-04-04 22:38:35 -07:00
parent 6da8b33053
commit 26bb13ade0
4 changed files with 48 additions and 9 deletions
@@ -20,6 +20,7 @@
 # THE SOFTWARE.
 ################################################################################

+import os
 from sqlitedb import SQLiteDB

 def gen_message(outfile):
@@ -59,7 +60,7 @@ def gen_ext_json_trace(db, table, start_ns, outfile):
  gen_message(outfile)

 def gen_ops_json_trace(db, table, base_pid, start_ns, outfile):
-  db.execute('create view B as select "Index", Name as name, ("dev-id" + %d) as pid, __lane as tid, ((BeginNs - %d)/1000) as ts, (DurationNs/1000) as dur from %s;' % (base_pid, start_ns, table));
+  db.execute('create view B as select "Index", "%s" as name, ("dev-id" + %d) as pid, __lane as tid, ((BeginNs - %d)/1000) as ts, (DurationNs/1000) as dur from %s;' % ('roctx-range' if 'ROCP_RENAME_KERNEL' in os.environ else 'Name',base_pid, start_ns, table));
  db.dump_json('B', table, outfile)
  db.execute('DROP VIEW B')
  gen_message(outfile)
@@ -179,6 +179,8 @@ usage() {
  echo "        </parameters>"
  echo "      </trace>"
  echo ""
+  echo "  --roctx-rename - to rename kernels with their enclosing rocTX range's message."
+  echo ""
  echo "  --trace-start <on|off> - to enable tracing on start [on]"
  echo "  --trace-period <dealy:length:rate> - to enable trace with initial delay, with periodic sample length and rate"
  echo "    Supported time formats: <number(m|s|ms|us)>"
@@ -425,6 +427,9 @@ while [ 1 ] ; do
    export ROCP_TIMESTAMP_ON=1
    GEN_STATS=1
    HIP_TRACE=1
+  elif [ "$1" = "--roctx-rename" ] ; then
+    ARG_VAL=0
+    export ROCP_RENAME_KERNEL=1
  elif [ "$1" = "--trace-start" ] ; then
    if [ "$2" = "off" ] ; then
      export ROCP_CTRL_RATE="-1:0:0"
@@ -220,7 +220,7 @@ class SQLiteDB:
          value = data[value_index]
          if label[:3] == '"__': continue
          if name_ptrn.search(label): value = sub_ptrn.sub(r'', value)
-          if label != '"Index"': args_list.append('%s:"%s"' % (label, value))
+          if label != '"Index"' and label != '"roctx-range"': args_list.append('%s:"%s"' % (label, value))

        fd.write(',{"ph":"%s",%s,\n  "args":{\n    %s\n  }\n}\n' % ('X', ','.join(vals_list), ',\n    '.join(args_list)))

@@ -20,7 +20,7 @@
 # THE SOFTWARE.
 ################################################################################

-import os, sys, re, subprocess
+import os, sys, re, subprocess, bisect
 from sqlitedb import SQLiteDB
 from mem_manager import MemManager
 import dform
@@ -257,12 +257,15 @@ ext_table_descr = [
  {'BeginNs':'INTEGER', 'EndNs':'INTEGER', 'pid':'INTEGER', 'tid':'INTEGER', 'Name':'TEXT', 'Index':'INTEGER', '__section':'INTEGER', '__lane':'INTEGER'}
 ]
 def fill_ext_db(table_name, db, indir, trace_name, api_pid):
+  global range_data
+
  file_name = indir + '/' + trace_name + '_trace.txt'
  # tms pid:tid cid:rid:'.....'
-  ptrn_val = re.compile(r'(\d+) (\d+):(\d+) (\d+):(\d+):(.*)$')
+  ptrn_val = re.compile(r'(\d+) (\d+):(\d+) (\d+):(\d+):"(.*)"$')

  if not os.path.isfile(file_name): return 0

+  range_data = {}
  range_stack = {}
  range_map = {}

@@ -275,12 +278,13 @@ def fill_ext_db(table_name, db, indir, trace_name, api_pid):
      if m:
        tms = int(m.group(1))
        pid = m.group(2)
-        tid = m.group(3)
+        tid = int(m.group(3))
        cid = int(m.group(4))
        rid = int(m.group(5))
        msg = m.group(6)

        rec_vals = []
+        if not tid in range_data: range_data[tid] = {}

        if cid != 2:
          rec_vals.append(tms)
@@ -307,6 +311,12 @@ def fill_ext_db(table_name, db, indir, trace_name, api_pid):
          rec_stack = pid_stack[tid]
          rec_vals = rec_stack.pop()
          rec_vals[1] = tms
+          # record the range's start/stop timestamps, its parent (ranges can be nested), and its message.
+          range_start = rec_vals[0]
+          range_stop = tms
+          range_parent = rec_stack[-1][0] if len(rec_stack) != 0 else 0
+          range_msg = rec_vals[4]
+          range_data[tid][range_start] = (range_stop, range_parent, range_msg)

        # range start
        if cid == 3:
@@ -364,6 +374,7 @@ def fill_api_db(table_name, db, indir, api_name, api_pid, dep_pid, dep_list, dep
  global hsa_activity_found
  global memory_manager

+  range_start_times = {}
  copy_csv = ''
  copy_index = 0

@@ -520,7 +531,26 @@ def fill_api_db(table_name, db, indir, api_name, api_pid, dep_pid, dep_list, dep
          copy_index += 1

        if op_found:
-          ops_patch_data[(corr_id, proc_id)] = (thread_id, stream_id, kernel_str)
+          roctx_msg = ''
+
+          if not thread_id in range_start_times:
+            range_start_times[thread_id] = sorted(range_data[thread_id].keys()) if thread_id in range_data else []
+          start_times = range_start_times[thread_id]
+
+          index = bisect.bisect_right(start_times,int(rec_vals[0]))
+          if index > 0:
+            # We found the range that is closest to this operation. Iterate the
+            # range stack this range is part of until we find a range that entirely
+            # contains the operation.
+            range_start = start_times[index - 1]
+            while range_start != 0:
+              (range_end, range_start, msg) = range_data[thread_id][range_start]
+              if int(rec_vals[1]) < range_end:
+                # This range contains the operation.
+                roctx_msg = msg
+                break
+
+          ops_patch_data[(corr_id, proc_id)] = (thread_id, stream_id, kernel_str, roctx_msg)

        if op_found:
          op_found = 0
@@ -628,8 +658,8 @@ def fill_copy_db(table_name, db, indir):

 # fill HCC ops DB
 ops_table_descr = [
-  ['BeginNs', 'EndNs', 'dev-id', 'queue-id', 'Name', 'pid', 'tid', 'stream-id', 'Index', 'Data', '__section', '__lane'],
-  {'Index':'INTEGER', 'Name':'TEXT', 'args':'TEXT', 'BeginNs':'INTEGER', 'EndNs':'INTEGER', 'dev-id':'INTEGER', 'queue-id':'INTEGER', 'pid':'INTEGER', 'tid':'INTEGER', 'Data':'TEXT', 'stream-id':'INTEGER', '__section':'INTEGER', '__lane':'INTEGER'}
+  ['BeginNs', 'EndNs', 'dev-id', 'queue-id', 'Name', 'pid', 'tid', 'roctx-range', 'stream-id', 'Index', 'Data', '__section', '__lane'],
+  {'Index':'INTEGER', 'Name':'TEXT', 'args':'TEXT', 'BeginNs':'INTEGER', 'EndNs':'INTEGER', 'dev-id':'INTEGER', 'queue-id':'INTEGER', 'pid':'INTEGER', 'tid':'INTEGER', 'roctx-range':'TEXT', 'Data':'TEXT', 'stream-id':'INTEGER', '__section':'INTEGER', '__lane':'INTEGER'}
 ]
 def fill_ops_db(kernel_table_name, mcopy_table_name, db, indir):
  global max_gpu_id
@@ -689,9 +719,11 @@ def fill_ops_db(kernel_table_name, mcopy_table_name, db, indir):

        thread_id = 0
        stream_id = 0
+        roctx_range = ''
        if (corr_id, proc_id) in ops_patch_data:
-          (thread_id, stream_id, name_patch) = ops_patch_data[(corr_id, proc_id)]
+          (thread_id, stream_id, name_patch, roctx_range) = ops_patch_data[(corr_id, proc_id)]
          if name_patch != '': name = name_patch
+          if roctx_range == '': roctx_range = name
        else:
          if is_barrier: continue
          else: fatal("hcc ops data not found: '" + record + "', " + str(corr_id) + ", " + str(proc_id))
@@ -700,6 +732,7 @@ def fill_ops_db(kernel_table_name, mcopy_table_name, db, indir):
        rec_vals[4] = name                       # Name
        rec_vals.append(proc_id)                 # pid
        rec_vals.append(thread_id)               # tid
+        rec_vals.append(roctx_range)             # roctx-range
        rec_vals.append(stream_id)               # StreamId
        rec_vals.append(corr_id)                 # Index