SWDEV-294319 & SWDEV-294321:

Added Support for launch kernel functions to fill_api_db
Added support for hipMemcpyToSymbol in add_memcpy
Added support for hsa_amd_memory_pool_allocate to be counted as source of allocations

Change-Id: I456a242ca1bc0c1bd39ae687a455b02c588de466


[ROCm/rocprofiler commit: b6910a9bc6]

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)

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