Merge "API tracing instrumentation" into amd-master-next

include/hip/hcc_detail/hip_runtime_api.h

@@ -513,7 +513,7 @@ hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig* pConfig);
  *
  * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue
  */
-hipError_t hipGetDeviceFlags(unsigned *flags);
+hipError_t hipGetDeviceFlags(unsigned int* flags);
 
 /**
  * @brief The bank width of shared memory on current device is set

rocclr/hip_device.cpp

@@ -99,7 +99,7 @@ hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device
 }
 
 hipError_t hipDeviceGetCount(int* count) {
-  HIP_INIT_API(NONE, count);
+  HIP_INIT_API(hipDeviceGetCount, count);
 
   HIP_RETURN(ihipDeviceGetCount(count));
 }
@@ -238,7 +238,7 @@ hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device )
 }
 
 hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) {
-  HIP_INIT_API(NONE, deviceId, acc);
+  HIP_INIT_API(hipHccGetAccelerator, deviceId, acc);
 
   assert(0 && "Unimplemented");
 
@@ -246,7 +246,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) {
 }
 
 hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av) {
-  HIP_INIT_API(NONE, stream, av);
+  HIP_INIT_API(hipHccGetAcceleratorView, stream, av);
 
   assert(0 && "Unimplemented");
 

rocclr/hip_device_runtime.cpp

@@ -480,7 +480,7 @@ hipError_t hipGetDeviceFlags ( unsigned int* flags ) {
 }
 
 hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event ) {
-  HIP_INIT_API(NONE, handle, event);
+  HIP_INIT_API(hipIpcGetEventHandle, handle, event);
 
   assert(0 && "Unimplemented");
 
@@ -488,7 +488,7 @@ hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event
 }
 
 hipError_t hipIpcOpenEventHandle ( hipEvent_t* event, hipIpcEventHandle_t handle ) {
-  HIP_INIT_API(NONE, event, handle);
+  HIP_INIT_API(hipIpcOpenEventHandle, event, handle);
 
   assert(0 && "Unimplemented");
 
@@ -542,7 +542,7 @@ hipError_t hipSetDeviceFlags ( unsigned int  flags ) {
 }
 
 hipError_t hipSetValidDevices ( int* device_arr, int  len ) {
-  HIP_INIT_API(NONE, device_arr, len);
+  HIP_INIT_API(hipSetValidDevices, device_arr, len);
 
   assert(0 && "Unimplemented");
 

rocclr/hip_memory.cpp

@@ -2090,7 +2090,7 @@ hipError_t hipArrayGetDescriptor(HIP_ARRAY_DESCRIPTOR* pArrayDescriptor,
 
 hipError_t hipMemcpyParam2DAsync(const hip_Memcpy2D* pCopy,
                                  hipStream_t stream) {
-  HIP_INIT_API(hipMemcpyParam2D, pCopy);
+  HIP_INIT_API(hipMemcpyParam2DAsync, pCopy);
 
   HIP_RETURN(ihipMemcpyParam2D(pCopy, stream, true));
 }

rocclr/hip_module.cpp

@@ -119,7 +119,7 @@ hipError_t hipModuleLoadDataEx(hipModule_t *module, const void *image,
                                void** optionsValues)
 {
   /* TODO: Pass options to Program */
-  HIP_INIT_API(hipModuleLoadData, module, image);
+  HIP_INIT_API(hipModuleLoadDataEx, module, image);
 
   HIP_RETURN(ihipModuleLoadData(module, image, 0));
 }
@@ -369,7 +369,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f,
                                  hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags = 0,
                                  uint32_t params = 0, uint32_t gridId = 0, uint32_t numGrids = 0,
                                  uint64_t prevGridSum = 0, uint64_t allGridSum = 0, uint32_t firstDevice = 0) {
-  HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
+  HIP_INIT_API(ihipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
     sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags, params);
 
   hip::Function* function = hip::Function::asFunction(f);
@@ -493,7 +493,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                     hipStream_t hStream, void** kernelParams, void** extra,
                                     hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags)
 {
-  HIP_INIT_API(NONE, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
+  HIP_INIT_API(hipExtModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
                localWorkSizeX, localWorkSizeY, localWorkSizeZ,
                sharedMemBytes, hStream,
                kernelParams, extra, startEvent, stopEvent, flags);
@@ -512,7 +512,7 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX,
                                     hipEvent_t startEvent,
                                     hipEvent_t stopEvent)
 {
-  HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ,
+  HIP_INIT_API(hipHccModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ,
                blockDimX, blockDimY, blockDimZ,
                sharedMemBytes, hStream,
                kernelParams, extra, startEvent, stopEvent);
@@ -529,7 +529,7 @@ hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t gridDimX,
                                     hipEvent_t startEvent,
                                     hipEvent_t stopEvent)
 {
-  HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ,
+  HIP_INIT_API(hipModuleLaunchKernelExt, f, gridDimX, gridDimY, gridDimZ,
                blockDimX, blockDimY, blockDimZ,
                sharedMemBytes, hStream,
                kernelParams, extra, startEvent, stopEvent);

rocclr/hip_platform.cpp

@@ -706,7 +706,7 @@ extern "C" hipError_t hipConfigureCall(
   size_t sharedMem,
   hipStream_t stream)
 {
-  HIP_INIT_API(NONE, gridDim, blockDim, sharedMem, stream);
+  HIP_INIT_API(hipConfigureCall, gridDim, blockDim, sharedMem, stream);
 
   PlatformState::instance().configureCall(gridDim, blockDim, sharedMem, stream);
 
@@ -719,7 +719,7 @@ extern "C" hipError_t __hipPushCallConfiguration(
   size_t sharedMem,
   hipStream_t stream)
 {
-  HIP_INIT_API(NONE, gridDim, blockDim, sharedMem, stream);
+  HIP_INIT_API(__hipPushCallConfiguration, gridDim, blockDim, sharedMem, stream);
 
   PlatformState::instance().configureCall(gridDim, blockDim, sharedMem, stream);
 
@@ -730,7 +730,7 @@ extern "C" hipError_t __hipPopCallConfiguration(dim3 *gridDim,
                                                 dim3 *blockDim,
                                                 size_t *sharedMem,
                                                 hipStream_t *stream) {
-  HIP_INIT_API(NONE, gridDim, blockDim, sharedMem, stream);
+  HIP_INIT_API(__hipPopCallConfiguration, gridDim, blockDim, sharedMem, stream);
 
   ihipExec_t exec;
   PlatformState::instance().popExec(exec);
@@ -747,7 +747,7 @@ extern "C" hipError_t hipSetupArgument(
   size_t size,
   size_t offset)
 {
-  HIP_INIT_API(NONE, arg, size, offset);
+  HIP_INIT_API(hipSetupArgument, arg, size, offset);
 
   PlatformState::instance().setupArgument(arg, size, offset);
 
@@ -756,7 +756,7 @@ extern "C" hipError_t hipSetupArgument(
 
 extern "C" hipError_t hipLaunchByPtr(const void *hostFunction)
 {
-  HIP_INIT_API(NONE, hostFunction);
+  HIP_INIT_API(hipLaunchByPtr, hostFunction);
 
   ihipExec_t exec;
   PlatformState::instance().popExec(exec);
@@ -1301,7 +1301,7 @@ extern "C" hipError_t hipLaunchKernel(const void *hostFunction,
                                       size_t sharedMemBytes,
                                       hipStream_t stream)
 {
-  HIP_INIT_API(NONE, hostFunction, gridDim, blockDim, args, sharedMemBytes,
+  HIP_INIT_API(hipLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes,
                stream);
 
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);

rocclr/hip_prof_gen.py

@@ -68,8 +68,8 @@ def filtr_api_name(name):
   return name
 
 def filtr_api_decl(record):
-  record = re.sub("\s__dparm\([^\)]*\)", '', record);
+  record = re.sub("\s__dparm\([^\)]*\)", r'', record);
-  record = re.sub("\(void\*\)", '', record);
+  record = re.sub("\(void\*\)", r'', record);
   return record
 
 # Normalizing API arguments
@@ -216,6 +216,8 @@ def parse_content(inp_file_p, api_map, out):
   api_name = ""
   # Valid public API found flag
   api_valid = 0
+  # API overload (parameters mismatch)
+  api_overload = 0
 
   # Input file patched content
   content = ''
@@ -240,6 +242,7 @@ def parse_content(inp_file_p, api_map, out):
 
     # Looking for API begin
     if found == 0:
+      record = re.sub(r'\s*extern\s+"C"\s+', r'', record);
       if beg_pattern.match(record):
         found = 1
         record = filtr_api_decl(record)
@@ -255,6 +258,8 @@ def parse_content(inp_file_p, api_map, out):
       # Checking if complete API matched
       if m:
         found = 2
+        api_valid = 0
+        api_overload = 0
         api_name = filtr_api_name(m.group(2))
         # Checking if API name is in the API map
         if (private_check_mode == 0) or (api_name in api_map):
@@ -280,48 +285,34 @@ def parse_content(inp_file_p, api_map, out):
             out[api_name] = filtr_api_opts(api_args)
             # Register missmatched API methods
           else:
+            api_overload = 1
             # Warning about mismatched API, possible non public overloaded version
             api_diff = '\t\t' + inp_file + " line(" + str(line_num) + ")\n\t\tapi: " + api_types + "\n\t\teta: " + eta_types
             message("\t" + api_name + ' args mismatch:\n' + api_diff + '\n')
 
-    if hip_patch_mode != 0:
-      # Looking for INIT macro
-      m = init_pattern.match(line)
-      if m:
-        if api_valid == 0: api_name = 'NONE'
-
-        if api_name == m.group(3):
-          if hip_patch_mode == 1: hip_patch_mode = 0
-          else: fatal("patching failed")
-        else:
-          hip_patch_mode = 2
-          init_args =  m.group(2)
-          if init_args != '': init_args = ', ' + init_args
-          line = m.group(1) + '(' + api_name + init_args + m.group(5) + '\n'
-          non_public_api = 0
-
     # API found action
     if found == 2:
-      # Looking for INIT macro
+      if hip_patch_mode != 0:
+        # Looking for INIT macro
+        m = init_pattern.match(line)
+        if m:
+          init_name = api_name
+          if api_overload == 1: init_name = 'NONE'
+          init_args = m.group(4)
+          line = m.group(1) + '(' + init_name + init_args + m.group(5) + '\n'
+
       m = init_pattern.match(line)
       if m:
         found = 0
-        non_public_api = 0
+        if api_valid == 1: message("\t" + api_name)
-
+        # Ignore if it is initialized as NONE
-        if api_valid == 1:
+        init_name = m.group(3)
-          api_valid = 0
+        if init_name != 'NONE':
-          message("\t" + api_name)
+          # Check if init name matching API name
-        else:
+          if init_name != api_name:
-          non_public_api = 1
+            fatal("init name mismatch: '" + init_name +  "' <> '" + api_name + "'")
-
-        if non_public_api == 1:
           # Registering dummy API for non public API if the name in INIT is not NONE
-          init_name = m.group(3)
+          if api_valid == 0:
-          # Ignore if it is initialized as NONE
-          if init_name != 'NONE':
-            # Check if init name matching API name
-            if init_name != api_name:
-              fatal("init name mismatch: '" + init_name +  "' <> '" + api_name + "'")
             # If init name is not in public API map then it is private API
             # else it was not identified and will be checked on finish
             if not init_name in api_map:
@@ -333,7 +324,6 @@ def parse_content(inp_file_p, api_map, out):
         # Expect INIT macro for valid public API
         # Removing and registering non-conformant APIs with missing HIP_INIT macro
         if api_valid == 1:
-          api_valid = 0
           if api_name in out:
             del out[api_name]
             del api_map[api_name]