diff --git a/projects/rocprofiler/bin/mem_manager.py b/projects/rocprofiler/bin/mem_manager.py index d466c37549..8b644a0014 100755 --- a/projects/rocprofiler/bin/mem_manager.py +++ b/projects/rocprofiler/bin/mem_manager.py @@ -24,7 +24,7 @@ import sys, os, re from sqlitedb import SQLiteDB pinned = ['hipMallocHost', 'hipHostMalloc', 'hipHostAlloc'] -ondevice = ['hipMalloc', 'hipMallocPitch', 'hipMallocArray', 'hipMalloc3DArray'] +ondevice = ['hipMalloc', 'hipMallocPitch', 'hipMallocArray', 'hipMalloc3DArray', 'hsa_amd_memory_pool_allocate'] mm_table_descr = [ ['BeginNs', 'EndNs', 'pid', 'tid', 'Name', 'Direction', 'SrcType', 'DstType', 'Size', 'BW', 'Async'], @@ -239,6 +239,9 @@ class MemManager: # hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, # size_t count, hipMemcpyKind kind); hip_memcpy_ptrn3 = re.compile(r'\(\s*dst\((.*)\) .* src\((.*)\) count\((\d+)\).*\)') + # hipMemcpyToSymbol(const void* symbolName, const void* src, size_t sizeBytes, + # size_t offset = 0, hipMemcpyKind kind) + hip_memcpy_ptrn4 = re.compile(r'\(\s*symbol\((.*)\) src\((.*)\) sizeBytes\((\d+)\).*\)') # memcopy with kind argument hip_memcpy_ptrn_kind = re.compile(r'.* kind\((\d+)\)\s*.*') #hsa_amd_memory_async_copy(void* dst, hsa_agent_t dst_agent, const void* src, @@ -257,6 +260,7 @@ class MemManager: # aysnc memcopy async_event_ptrn = re.compile(r'Async|async') m_basic_hip = hip_memcpy_ptrn.match(args) + m_basic_hsa3 = hip_memcpy_ptrn4.match(args) m_basic_hsa_prev = hsa_memcpy_ptrn_prev.match(args) m_basic_hsa = hsa_memcpy_ptrn.match(args) m_basic_hsa2 = hsa_memcpy_ptrn2.match(args) @@ -299,6 +303,7 @@ class MemManager: srcptr_type = self.get_ptr_type(src_agent_ptr) size = int(m_basic_hsa_prev.group(5)) condition_matched = True + if m_basic_hsa: dstptr = m_basic_hsa.group(1) dst_agent_ptr = m_basic_hsa.group(2) @@ -322,6 +327,14 @@ class MemManager: size = x*y*z condition_matched = True + if m_basic_hsa3: + dstptr = m_basic_hsa3.group(1) + dstptr_type = self.get_ptr_type(dstptr) + srcptr = m_basic_hsa3.group(2) + srcptr_type = self.get_ptr_type(srcptr) + size = int(m_basic_hsa3.group(3)) + condition_matched = True + if m_array: dstptr = m_array.group(1) dstptr_type = self.get_ptr_type(dstptr) diff --git a/projects/rocprofiler/bin/tblextr.py b/projects/rocprofiler/bin/tblextr.py index deafb1999f..2e52169003 100755 --- a/projects/rocprofiler/bin/tblextr.py +++ b/projects/rocprofiler/bin/tblextr.py @@ -354,6 +354,7 @@ def fill_api_db(table_name, db, indir, api_name, api_pid, dep_pid, dep_list, dep copy_index = 0 ptrn_val = re.compile(r'(\d+):(\d+) (\d+):(\d+) ([^\(]+)(\(.*)$') + hip_launch_kernel_ptrn = re.compile(r'hipExtLaunchKernel|hipLaunchKernel') hip_mcopy_ptrn = re.compile(r'hipMemcpy|hipMemset') hip_wait_event_ptrn = re.compile(r'WaitEvent') hip_sync_event_ptrn = re.compile(r'hipStreamSynchronize') @@ -488,6 +489,9 @@ def fill_api_db(table_name, db, indir, api_name, api_pid, dep_pid, dep_list, dep mcopy_found = 1 op_found = 1 + if hip_launch_kernel_ptrn.match(record_name): + op_found = 1 + # HIP WaitEvent API if wait_event_ptrn.search(record_name): op_found = 1