diff --git a/bin/mem_manager.py b/bin/mem_manager.py new file mode 100755 index 0000000000..ed9aa67a30 --- /dev/null +++ b/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) + diff --git a/bin/tblextr.py b/bin/tblextr.py index baff8eb24d..a8ed23fb96 100755 --- a/bin/tblextr.py +++ b/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)