diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 17c34b0ad5..744885923f 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/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 diff --git a/rocclr/hip_device.cpp b/rocclr/hip_device.cpp index 8695ef43e5..b8a5e65d29 100644 --- a/rocclr/hip_device.cpp +++ b/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"); diff --git a/rocclr/hip_device_runtime.cpp b/rocclr/hip_device_runtime.cpp index 531f35c732..3823af5206 100644 --- a/rocclr/hip_device_runtime.cpp +++ b/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"); diff --git a/rocclr/hip_memory.cpp b/rocclr/hip_memory.cpp index 4bd75b8ffb..febf5cc700 100755 --- a/rocclr/hip_memory.cpp +++ b/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)); } diff --git a/rocclr/hip_module.cpp b/rocclr/hip_module.cpp index 7cda6864c6..3a51001063 100755 --- a/rocclr/hip_module.cpp +++ b/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); diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index 8759ef47f1..b67c6ac90a 100755 --- a/rocclr/hip_platform.cpp +++ b/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(stream); diff --git a/rocclr/hip_prof_gen.py b/rocclr/hip_prof_gen.py index c20df3c1aa..c508888140 100755 --- a/rocclr/hip_prof_gen.py +++ b/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("\(void\*\)", '', record); + record = re.sub("\s__dparm\([^\)]*\)", r'', 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: - api_valid = 0 - message("\t" + api_name) - else: - non_public_api = 1 - - if non_public_api == 1: + if api_valid == 1: message("\t" + api_name) + # Ignore if it is initialized as NONE + init_name = m.group(3) + if init_name != 'NONE': + # Check if init name matching API name + if init_name != api_name: + fatal("init name mismatch: '" + init_name + "' <> '" + api_name + "'") # Registering dummy API for non public API if the name in INIT is not NONE - init_name = m.group(3) - # 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 api_valid == 0: # 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]