memcpy info added to csv, sqlitedb and json

Change-Id: I56f2ec667c7f3741b519f3335a3eec26b34c0ab4

bin/mem_manager.py

@@ -0,0 +1,207 @@
+################################################################################
+# Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in
+# all copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+# THE SOFTWARE.
+################################################################################
+
+import os, re
+from sqlitedb import SQLiteDB
+
+pinned = ['hipMallocHost', 'hipHostMalloc', 'hipHostAlloc']
+ondevice = ['hipMallocPitch', 'hipMallocArray', 'hipMalloc3DArray']
+
+mm_table_descr = [
+  ['BeginNs', 'EndNs', 'pid', 'tid', 'Name', 'Direction', 'SrcType', 'DstType', 'Size', 'BW', 'Async'],
+  {'BeginNs':'INTEGER', 'EndNs':'INTEGER', 'pid':'INTEGER', 'tid':'INTEGER', 'Name':'TEXT', 'Direction':'TEXT', 'SrcType':'TEXT', 'DstType':'TEXT', 'Size':'INTEGER', 'BW':'TEXT', 'Async':'TEXT'}
+]
+
+def fatal(msg):
+  sys.stderr.write(sys.argv[0] + ": " + msg + "\n");
+  sys.exit(1)
+
+# Mem copy manager class
+class MemManager:
+
+  def __init__(self, db):
+    self.db = db
+    self.allocations = {}
+    self.memcopies = {}
+    self.filename = ''
+    self.fd = 0
+
+  def __del__(self):
+    if self.fd != 0: self.fd.close()
+
+  # register allo and memcpy API calls
+  def register_api(self, rec_vals):
+    res = ''
+    malloc_ptrn = re.compile(r'hip.*Malloc')
+    record_name = rec_vals[4]
+    record_args = rec_vals[5]
+    mallocm = malloc_ptrn.match(record_name)
+    if mallocm:
+      self.add_allocation(record_name, record_args)
+    else:
+      hip_ptrn_mcopy = re.compile(r'hipMemcpy')
+      memcpym = hip_ptrn_mcopy.match(record_name)
+      if memcpym:
+        res = self.add_memcpy(rec_vals)
+
+    return res
+
+  # add allocation to map
+  def add_allocation(self, event, args):
+    choice = 0
+    if event == "hipMallocPitch":
+      hipMalloc_ptrn = re.compile(r'\(ptr\((.*)\) width\((.*)\) height\((.*)\)\)')
+      choice = 1
+    elif event == "hipMallocArray":
+      hipMalloc_ptrn = re.compile(r'\(array\((.*)\) width\((.*)\) height\((.*)\)\)')
+      choice = 1
+    elif event == "hipMalloc3DArray":
+      hipMalloc_ptrn = re.compile(r'\(array\((.*)\) width\((.*)\) height\((.*)\) depth\((.*)\)\)')
+      choice = 2
+    else:
+      hipMalloc_ptrn = re.compile(r'\(ptr\((.*)\) size\((.*)\)\)')
+      choice = 3
+    mhipMalloc = hipMalloc_ptrn.match(args)
+    if mhipMalloc:
+      ptr = mhipMalloc.group(1)
+      if choice == 3:
+        size = mhipMalloc.group(2)
+      elif choice == 1:
+        size = mhipMalloc.group(2) * mhipMalloc.group(3)
+      else:
+        size = mhipMalloc.group(2) * mhipMalloc.group(3) * mhipMalloc.group(4)
+      self.allocations[ptr]=(size, event)
+    return
+
+  #get type of ptr
+  def get_ptr_type(self, ptr):
+    ptr_type = ''
+    found = 0
+    for pt in self.allocations.keys():
+      (size, event) = self.allocations[pt]
+      if int(ptr,16) >= int(pt, 16) and int(ptr,16) <= int(pt, 16) + int(size, 16):
+        found = 1
+        break
+    if not found:
+      ptr_type = 'pageable'
+    elif event in pinned:
+      ptr_type = 'pinned'
+    elif event in ondevice:
+      ptr_type = 'device'
+    else:
+      fatal('internal error: ptr(' + ptr + ') cannot be identified')
+    return ptr_type
+
+  # add memcpy to map
+  def add_memcpy(self, recvals):
+    recordid = recvals[6]  #same as corrid
+    event = recvals[4]
+    start_time = recvals[0] # sync time stamp
+    end_time = recvals[1] # sync time stamp
+    args = recvals[5]
+    procid = recvals[2] # used to query async entries
+    pid = recvals[2]
+    tid = recvals[3]
+    separator = ','
+
+    select_expr = '"Index" = ' + str(recordid) + ' AND "proc-id" = ' + str(procid)
+    async_copy_records = self.db.table_get_record('COPY', select_expr)  #List of async copy record fields
+    async_copy_start_time = async_copy_records[0]
+    async_copy_end_time = async_copy_records[1]
+
+    #hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind)
+    hipMemcpy_ptrn = re.compile(r'\(dst\((.*)\) src\((.*)\) sizeBytes\((\d+)\).*\)')
+    #hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
+    #                   size_t height, hipMemcpyKind kind);
+    hipMemcpy_ptrn2 = re.compile(r'\(dst\((.*)\) .* src\((.*)\) .* width\((\d+)\) height\((\d+)\).*\)')
+    #hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
+    #                        size_t count, hipMemcpyKind kind);
+    hipMemcpy_ptrn3 = re.compile(r'\(dst\((.*)\) .* src\((.*)\) count\((\d+)\).*\)')
+    mhip = hipMemcpy_ptrn.match(args)
+    mhip2 = hipMemcpy_ptrn2.match(args)
+    mhip3 = hipMemcpy_ptrn3.match(args)
+
+    if mhip or mhip2 or mhip3:
+      size = 0
+      dstptr_type = 'unknown'
+      srcptr_type = 'unknown'
+      if mhip:
+        dstptr = mhip.group(1)
+        dstptr_type = self.get_ptr_type(dstptr)
+        srcptr = mhip.group(2)
+        srcptr_type = self.get_ptr_type(srcptr)
+        size = int(mhip.group(3), 16)
+      if mhip3:
+        dstptr = mhip3.group(1)
+        dstptr_type = self.get_ptr_type(dstptr)
+        srcptr = mhip3.group(2)
+        srcptr_type = self.get_ptr_type(srcptr)
+        size = int(mhip3.group(3), 16)
+      if mhip2:
+        dstptr = mhip2.group(1)
+        dstptr_type = self.get_ptr_type(dstptr)
+        srcptr = mhip2.group(2)
+        srcptr_type = self.get_ptr_type(srcptr)
+        size = int(mhip2.group(3)*mhip2.group(4), 16)
+      bandwidth = float(size) / (int(end_time) - int(start_time))
+      copy_line = str(start_time) + ',' + str(end_time) + ',' + pid + ',' + tid + ',' + event
+      hipMemcpy_ptrn_kind = re.compile(r'.* kind\((\d+)\)\s*.*')
+      mkind = hipMemcpy_ptrn_kind.match(args)
+      if mkind:
+        if mkind.group(1) == '0': copy_line += separator + "Direction=hipMemcpyHostToHost"
+        if mkind.group(1) == '1': copy_line += separator + "Direction=hipMemcpyHostToDevice"
+        if mkind.group(1) == '2': copy_line += separator + "Direction=hipMemcpyDeviceToHost"
+        if mkind.group(1) == '3': copy_line += separator + "Direction=hipMemcpyDeviceToDevice"
+        if mkind.group(1) == '4': copy_line += separator + "Direction=runtime copy-kind"
+      else:
+        copy_line += ", -"
+      copy_line += separator + 'SrcType=' + srcptr_type + separator + 'DstType=' + dstptr_type + separator + "Size=" + str(round(size*1e-6, 2)) + separator + "BW=" + str(round(bandwidth,2))
+      async_event_ptrn = re.compile(r'hipMemcpy.*Async')
+      masync = async_event_ptrn.match(event)
+      if masync:
+        copy_line += separator + 'Async=yes'
+      else:
+        copy_line += separator + 'Async=no'
+
+    self.memcopies[recordid] = copy_line
+    return copy_line;
+
+  def dump_data(self):
+    # To create “MM” table in DB on the finish:
+    table_name = "MM"
+    file_name = os.environ['PWD'] + '/results_memcopy_info.csv'
+    print("File '" + file_name + "' is generating")
+    table_handle = self.db.add_table(table_name, mm_table_descr)
+
+    fld_ptrn = re.compile(r'(.*)=(.*)')
+    for (key, record) in self.memcopies.items():
+      rec_vals_array = []
+      for rec in record.split(','):
+        fld_ptrnm = fld_ptrn.match(rec)
+        if fld_ptrnm:
+          rec_vals_array.append(fld_ptrnm.group(2))
+        else:
+          rec_vals_array.append(rec)
+      self.db.insert_entry(table_handle, rec_vals_array)
+    # To dump the MM table as CSV by:
+    self.db.dump_csv(table_name, file_name)
+

bin/tblextr.py

@@ -22,6 +22,7 @@
 
 import os, sys, re, subprocess
 from sqlitedb import SQLiteDB
+from mem_manager import MemManager
 import dform
 
 EXT_PID = 0
@@ -298,8 +299,8 @@ def set_field(args, field, val):
 
 # Fill API DB
 api_table_descr = [
-  ['BeginNs', 'EndNs', 'pid', 'tid', 'Name', 'args', 'Index'],
-  {'BeginNs':'INTEGER', 'EndNs':'INTEGER', 'pid':'INTEGER', 'tid':'INTEGER', 'Name':'TEXT', 'args':'TEXT', 'Index':'INTEGER'}
+  ['BeginNs', 'EndNs', 'pid', 'tid', 'Name', 'args', 'Index', 'Data'],
+  {'BeginNs':'INTEGER', 'EndNs':'INTEGER', 'pid':'INTEGER', 'tid':'INTEGER', 'Name':'TEXT', 'args':'TEXT', 'Index':'INTEGER', 'Data':'TEXT'}
 ]
 # Filling API records DB table
 # table_name - created DB table name
@@ -312,6 +313,7 @@ api_table_descr = [
 # dep_filtr - registered dependencies by record ID
 def fill_api_db(table_name, db, indir, api_name, api_pid, dep_pid, dep_list, dep_filtr, expl_id):
   global hsa_activity_found
+  global memory_manager
   copy_raws = []
   if (hsa_activity_found): copy_raws = db.table_get_raws('COPY')
   copy_csv = ''
@@ -353,11 +355,10 @@ def fill_api_db(table_name, db, indir, api_name, api_pid, dep_pid, dep_list, dep
       m = ptrn_val.match(record)
       if m:
         rec_vals = []
-        rec_len = len(api_table_descr[0])
-        for ind in range(1,rec_len):
+        rec_len = len(api_table_descr[0]) - 1
+        for ind in range(1, rec_len):
           rec_vals.append(m.group(ind))
         proc_id = rec_vals[2]
-        rec_vals[2] = api_pid
         record_name = rec_vals[4]
         record_args = rec_vals[5]
         rec_vals.append(record_id)
@@ -409,17 +410,22 @@ def fill_api_db(table_name, db, indir, api_name, api_pid, dep_pid, dep_list, dep
                 dev_id = m.group(2)
                 select_expr += ' AND "dev-id" = ' + dev_id
                 activity_record_patching(db, ops_table_name, 1, kernel_name, stream_found, stream_id, select_expr)
-              fatal('Bad multi-kernel format: "' + kernel_item + '" in "' + kernel_str + '"')
+              else:
+                fatal('Bad multi-kernel format: "' + kernel_item + '" in "' + kernel_str + '"')
           else:
             activity_record_patching(db, ops_table_name, kernel_found, kernel_str, stream_found, stream_id, select_expr)
 
+        rec_vals.append(memory_manager.register_api(rec_vals))
+
+        rec_vals[2] = api_pid
+
         db.insert_entry(table_handle, rec_vals)
         record_id += 1
       else: fatal(api_name + " bad record: '" + record + "'")
 
   # inserting of dispatch events correlated to the dependent dispatches
   for (tid, from_ns) in dep_list:
-    db.insert_entry(table_handle, [from_ns, from_ns, api_pid, tid, 'hsa_dispatch', '', record_id])
+    db.insert_entry(table_handle, [from_ns, from_ns, api_pid, tid, 'hsa_dispatch', '', record_id, ''])
     record_id += 1
 
   # registering dependencies informatino
@@ -479,8 +485,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', 'Index', 'proc-id'],
-  {'Index':'INTEGER', 'proc-id':'INTEGER', 'Name':'TEXT', 'args':'TEXT', 'BeginNs':'INTEGER', 'EndNs':'INTEGER', 'dev-id':'INTEGER', 'queue-id':'INTEGER', 'pid':'INTEGER', 'tid':'INTEGER'}
+  ['BeginNs', 'EndNs', 'dev-id', 'queue-id', 'Name', 'pid', 'tid', 'Index', 'proc-id', 'Data'],
+  {'Index':'INTEGER', 'proc-id':'INTEGER', 'Name':'TEXT', 'args':'TEXT', 'BeginNs':'INTEGER', 'EndNs':'INTEGER', 'dev-id':'INTEGER', 'queue-id':'INTEGER', 'pid':'INTEGER', 'tid':'INTEGER', 'Data':'TEXT'}
 ]
 def fill_ops_db(kernel_table_name, mcopy_table_name, db, indir):
   global max_gpu_id
@@ -534,6 +540,7 @@ def fill_ops_db(kernel_table_name, mcopy_table_name, db, indir):
         rec_vals.append(0)                       # tid
         rec_vals.append(corr_id)                 # Index
         rec_vals.append(proc_id)                 # proc-id
+        rec_vals.append('')                      # Data
         db.insert_entry(table_handle, rec_vals)
 
         # registering a dependency filtr
@@ -596,6 +603,7 @@ else:
 
   with open(dbfile, mode='w') as fd: fd.truncate()
   db = SQLiteDB(dbfile)
+  memory_manager = MemManager(db)
 
   ext_trace_found = fill_ext_db('rocTX', db, indir, 'roctx', EXT_PID)
 
@@ -695,6 +703,9 @@ else:
   if any_trace_found:
     db.metadata_json(jsonfile, sysinfo_file)
     db.close_json(jsonfile);
+
+  memory_manager.dump_data()
+
   db.close()
 
 sys.exit(0)