AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: I6fe9b3b4bde9cb49ff3b8119e66da23945ef7e2f
This commit is contained in:
Jenkins
2018-11-19 21:04:57 -06:00
parent 715134e2ff 32dd47c9a6
commit d581f58b1b
49 changed files with 4359 additions and 2581 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/CMakeLists.txt
+7 -11
View File
@@ -207,9 +207,6 @@ if(HIP_PLATFORM STREQUAL "hcc")
src/env.cpp
src/program_state.cpp)
set(SOURCE_FILES_DEVICE
src/device_util.cpp)
execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS)
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic")
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906")
@@ -222,19 +219,18 @@ if(HIP_PLATFORM STREQUAL "hcc")
target_link_libraries(hip_hcc PRIVATE hc_am)
target_link_libraries(hip_hcc_static PRIVATE hc_am)
endif()
add_library(hip_device STATIC ${SOURCE_FILES_DEVICE})
string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS})
foreach(TARGET hip_hcc hip_hcc_static hip_device)
foreach(TARGET hip_hcc hip_hcc_static)
target_include_directories(${TARGET} SYSTEM INTERFACE $<INSTALL_INTERFACE:$<INSTALL_PREFIX>/include>;${HSA_PATH}/include)
endforeach()
add_library(host INTERFACE)
target_link_libraries(host INTERFACE hip_hcc)
add_library(device INTERFACE)
if(HIP_COMPILER STREQUAL "hcc")
target_link_libraries(device INTERFACE host hip_device hcc::hccrt hcc::hc_am)
elseif(HIP_COMPILER STREQUAL "clang")
target_link_libraries(device INTERFACE host hip_device)
target_link_libraries(device INTERFACE host hcc::hccrt hcc::hc_am)
else()
target_link_libraries(device INTERFACE host)
endif()
# Generate .hipInfo
@@ -256,7 +252,7 @@ endif()
#############################
# Install hip_hcc if platform is hcc
if(HIP_PLATFORM STREQUAL "hcc")
install(TARGETS hip_hcc_static hip_hcc hip_device DESTINATION lib)
install(TARGETS hip_hcc_static hip_hcc DESTINATION lib)
# Install .hipInfo
install(FILES ${PROJECT_BINARY_DIR}/.hipInfo DESTINATION lib)
@@ -284,7 +280,7 @@ set(BIN_INSTALL_DIR ${CMAKE_INSTALL_PREFIX}/bin)
set(CONFIG_PACKAGE_INSTALL_DIR ${LIB_INSTALL_DIR}/cmake/hip)
if(HIP_PLATFORM STREQUAL "hcc")
install(TARGETS hip_hcc_static hip_hcc hip_device host device EXPORT hip-targets DESTINATION ${LIB_INSTALL_DIR})
install(TARGETS hip_hcc_static hip_hcc host device EXPORT hip-targets DESTINATION ${LIB_INSTALL_DIR})
install(EXPORT hip-targets DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR} NAMESPACE hip::)
include(CMakePackageConfigHelpers)
@@ -353,7 +349,7 @@ add_custom_target(pkg_hip_hcc COMMAND ${CMAKE_COMMAND} .
COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
WORKING_DIRECTORY ${BUILD_DIR}
DEPENDS hip_hcc hip_device hip_hcc_static)
DEPENDS hip_hcc hip_hcc_static)
# Package: hip_nvcc
set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_nvcc)
hipamd/bin/hipcc
+6 -2
View File
@@ -498,6 +498,10 @@ foreach $arg (@ARGV)
$obj = "$tmpdir/$obj";
my $fileType = `file $obj`;
my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/);
if ($fileType =~ m/ELF/) {
my $sections = `readelf -e -W $obj`;
$isObj = !($sections =~ m/__CLANG_OFFLOAD_BUNDLE__/);
}
$allIsObj = ($allIsObj and $isObj);
if ($isObj) {
$realObjs = ($realObjs . " " . $obj);
@@ -723,9 +727,9 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
if ($needHipHcc) {
if ($linkType eq 0) {
substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/libhip_hcc_static.a $HIP_PATH/lib/libhip_device.a " ;
substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/libhip_hcc_static.a " ;
} else {
substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so $HIP_PATH/lib/libhip_device.a ";
substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so ";
}
}
hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md
+65 -58
View File
@@ -22,12 +22,12 @@
| typedef |`CUDA_RESOURCE_VIEW_DESC_st` | |
| struct |`CUDA_TEXTURE_DESC` | |
| typedef |`CUDA_TEXTURE_DESC_st` | |
| struct |`CUdevprop` |`hipDeviceProp_t` |
| typedef |`CUdevprop_st` |`hipDeviceProp_t` |
| struct |`CUipcEventHandle` | |
| typedef |`CUipcEventHandle_st` | |
| struct |`CUipcMemHandle` | |
| typedef |`CUipcMemHandle_st` | |
| struct |`CUdevprop` | |
| typedef |`CUdevprop_st` | |
| struct |`CUipcEventHandle` |`ihipIpcEventHandle_t` |
| typedef |`CUipcEventHandle_st` |`ihipIpcEventHandle_t` |
| struct |`CUipcMemHandle` |`hipIpcMemHandle_t` |
| typedef |`CUipcMemHandle_st` |`hipIpcMemHandle_st` |
| union |`CUstreamBatchMemOpParams` | |
| typedef |`CUstreamBatchMemOpParams_union` | |
| enum |***`CUaddress_mode`*** | |
@@ -56,14 +56,14 @@
| 0x20 |*`CU_AD_FORMAT_FLOAT`* |*`HIP_AD_FORMAT_FLOAT`* |
| enum |***`CUctx_flags`*** | |
| typedef |***`CUctx_flags_enum`*** | |
| 0x00 |*`CU_CTX_SCHED_AUTO`* | |
| 0x01 |*`CU_CTX_SCHED_SPIN`* | |
| 0x02 |*`CU_CTX_SCHED_YIELD`* | |
| 0x04 |*`CU_CTX_SCHED_BLOCKING_SYNC`* | |
| 0x04 |*`CU_CTX_BLOCKING_SYNC`* | |
| 0x07 |*`CU_CTX_SCHED_MASK`* | |
| 0x08 |*`CU_CTX_MAP_HOST`* | |
| 0x10 |*`CU_CTX_LMEM_RESIZE_TO_MAX`* | |
| 0x00 |*`CU_CTX_SCHED_AUTO`* |`hipDeviceScheduleAuto` |
| 0x01 |*`CU_CTX_SCHED_SPIN`* |`hipDeviceScheduleSpin` |
| 0x02 |*`CU_CTX_SCHED_YIELD`* |`hipDeviceScheduleYield` |
| 0x04 |*`CU_CTX_SCHED_BLOCKING_SYNC`* |`hipDeviceScheduleBlockingSync` |
| 0x04 |*`CU_CTX_BLOCKING_SYNC`* |`hipDeviceScheduleBlockingSync` |
| 0x07 |*`CU_CTX_SCHED_MASK`* |`hipDeviceScheduleMask` |
| 0x08 |*`CU_CTX_MAP_HOST`* |`hipDeviceMapHost` |
| 0x10 |*`CU_CTX_LMEM_RESIZE_TO_MAX`* |`hipDeviceLmemResizeToMax` |
| 0x1f |*`CU_CTX_FLAGS_MASK`* | |
| enum |***`CUdevice_attribute`*** |***`hipDeviceAttribute_t`*** |
| typedef |***`CUdevice_attribute_enum`*** |***`hipDeviceAttribute_t`*** |
@@ -154,8 +154,8 @@
| 80 |*`CU_DEVICE_ATTRIBUTE_LOCAL_L1_CACHE_SUPPORTED`* | |
| 81 |*`CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR`* |*`hipDeviceAttributeMaxSharedMemoryPerMultiprocessor`* |
| 82 |*`CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_MULTIPROCESSOR`* | |
| 83 |*`CU_DEVICE_ATTRIBUTE_MANAGED_MEMORY`* |*`hipDeviceAttributeManagedMemory`* |
| 84 |*`CU_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD`* | |
| 83 |*`CU_DEVICE_ATTRIBUTE_MANAGED_MEMORY`* | |
| 84 |*`CU_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD`* |*`hipDeviceAttributeIsMultiGpuBoard`* |
| 85 |*`CU_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD_GROUP_ID`* | |
| 86 |*`CU_DEVICE_ATTRIBUTE_HOST_NATIVE_ATOMIC_SUPPORTED`* | |
| 87 |*`CU_DEVICE_ATTRIBUTE_SINGLE_TO_DOUBLE_PRECISION_PERF_RATIO`* | |
@@ -184,8 +184,8 @@
| typedef |***`CUfilter_mode_enum`*** |***`hipTextureFilterMode`*** |
| 0 |*`CU_TR_FILTER_MODE_POINT`* |*`hipFilterModePoint`* |
| 1 |*`CU_TR_FILTER_MODE_LINEAR`* |*`hipFilterModeLinear`* |
| enum |***`CUfunc_cache`*** |***`hipFuncCache`*** |
| typedef |***`CUfunc_cache_enum`*** |***`hipFuncCache`*** |
| enum |***`CUfunc_cache`*** |***`hipFuncCache_t`*** |
| typedef |***`CUfunc_cache_enum`*** |***`hipFuncCache_t`*** |
| 0x00 |*`CU_FUNC_CACHE_PREFER_NONE`* |*`hipFuncCachePreferNone`* |
| 0x01 |*`CU_FUNC_CACHE_PREFER_SHARED`* |*`hipFuncCachePreferShared`* |
| 0x02 |*`CU_FUNC_CACHE_PREFER_L1`* |*`hipFuncCachePreferL1`* |
@@ -301,12 +301,12 @@
| 0x1 |*`CU_MEM_ATTACH_GLOBAL`* | |
| 0x2 |*`CU_MEM_ATTACH_HOST`* | |
| 0x4 |*`CU_MEM_ATTACH_SINGLE`* | |
| enum |***`CUmemorytype`*** | |
| typedef |***`CUmemorytype_enum`*** | |
| 0x01 |*`CU_MEMORYTYPE_HOST`* | |
| 0x02 |*`CU_MEMORYTYPE_DEVICE`* | |
| 0x03 |*`CU_MEMORYTYPE_ARRAY`* | |
| 0x04 |*`CU_MEMORYTYPE_UNIFIED`* | |
| enum |***`CUmemorytype`*** |*`hipMemoryType`* |
| typedef |***`CUmemorytype_enum`*** |*`hipMemoryType`* |
| 0x01 |*`CU_MEMORYTYPE_HOST`* |*`hipMemoryTypeHost`* |
| 0x02 |*`CU_MEMORYTYPE_DEVICE`* |*`hipMemoryTypeDevice`* |
| 0x03 |*`CU_MEMORYTYPE_ARRAY`* |*`hipMemoryTypeArray`* |
| 0x04 |*`CU_MEMORYTYPE_UNIFIED`* |*`hipMemoryTypeUnified`* |
| enum |***`CUmem_range_attribute`*** | |
| typedef |***`CUmem_range_attribute_enum`*** | |
| 1 |*`CU_MEM_RANGE_ATTRIBUTE_READ_MOSTLY`* | |
@@ -513,28 +513,28 @@
| 0x01 |*`CU_D3D11_DEVICE_LIST_ALL`* | |
| 0x02 |*`CU_D3D11_DEVICE_LIST_CURRENT_FRAME`* | |
| 0x03 |*`CU_D3D11_DEVICE_LIST_NEXT_FRAME`* | |
| typedef |`CUarray` |`hipArray *` |
| struct |`CUarray_st` |`hipArray` |
| typedef |`CUarray` |`hipArray *` |
| struct |`CUctx` |`ihipCtx_t` |
| typedef |`CUcontext_st` |`hipCtx_t` |
| typedef |`CUdevice` |`hipDevice_t` |
| typedef |`CUdeviceptr` |`hipDeviceptr_t` |
| typedef |`CUeglStreamConnection` | |
| struct |`CUeglStreamConnection_st` | |
| typedef |`CUeglStreamConnection` | |
| typedef |`CUevent` |`hipEvent_t` |
| struct |`CUevent_st` |`ihipEvent_t` |
| typedef |`CUfunction` |`hipFunction_t` |
| struct |`CUfunc_st` |`ihipModuleSymbol_t` |
| typedef |`CUgraphicsResource` | |
| struct |`CUgraphicsResource_st` | |
| typedef |`CUmipmappedArray` | |
| struct |`CUmipmappedArray_st` | |
| typedef |`CUmipmappedArray` |`hipMipmappedArray_t` |
| struct |`CUmipmappedArray_st` |`hipMipmappedArray` |
| typedef |`CUmodule` |`hipModule_t` |
| struct |`CUmod_st` |`ihipModule_t` |
| typedef |`CUstream` |`hipStream_t` |
| struct |`CUstream_st` |`ihipStream_t` |
| typedef |`CUstreamCallback` |`hipStreamCallback_t` |
| typedef |`CUsurfObject` | |
| typedef |`CUsurfObject` |`hipSurfaceObject_t` |
| typedef |`CUsurfref` | |
| struct |`CUsurfref_st` | |
| typedef |`CUtexObject` |`hipTextureObject_t` |
@@ -544,12 +544,12 @@
| define |`CU_LAUNCH_PARAM_BUFFER_POINTER` |`HIP_LAUNCH_PARAM_BUFFER_POINTER` |
| define |`CU_LAUNCH_PARAM_BUFFER_SIZE` |`HIP_LAUNCH_PARAM_BUFFER_SIZE` |
| define |`CU_LAUNCH_PARAM_END` |`HIP_LAUNCH_PARAM_END` |
| define |`CU_MEMHOSTALLOC_DEVICEMAP` | |
| define |`CU_MEMHOSTALLOC_PORTABLE` | |
| define |`CU_MEMHOSTALLOC_WRITECOMBINED` | |
| define |`CU_MEMHOSTREGISTER_DEVICEMAP` | |
| define |`CU_MEMHOSTREGISTER_IOMEMORY` | |
| define |`CU_MEMHOSTREGISTER_PORTABLE` | |
| define |`CU_MEMHOSTALLOC_DEVICEMAP` |`hipHostMallocMapped` |
| define |`CU_MEMHOSTALLOC_PORTABLE` |`hipHostMallocPortable` |
| define |`CU_MEMHOSTALLOC_WRITECOMBINED` |`hipHostAllocWriteCombined` |
| define |`CU_MEMHOSTREGISTER_DEVICEMAP` |`hipHostRegisterMapped` |
| define |`CU_MEMHOSTREGISTER_IOMEMORY` |`hipHostRegisterIoMemory` |
| define |`CU_MEMHOSTREGISTER_PORTABLE` |`hipHostRegisterPortable` |
| define |`CU_PARAM_TR_DEFAULT` | |
| define |`CU_STREAM_LEGACY` | |
| define |`CU_STREAM_PER_THREAD` | |
@@ -558,11 +558,11 @@
| define |`CU_TRSF_READ_AS_INTEGER` | |
| define |`CU_TRSF_SRGB` | |
| define |`CUDA_ARRAY3D_2DARRAY` | |
| define |`CUDA_ARRAY3D_CUBEMAP` | |
| define |`CUDA_ARRAY3D_CUBEMAP` |`hipArrayCubemap` |
| define |`CUDA_ARRAY3D_DEPTH_TEXTURE` | |
| define |`CUDA_ARRAY3D_LAYERED` | |
| define |`CUDA_ARRAY3D_SURFACE_LDST` | |
| define |`CUDA_ARRAY3D_TEXTURE_GATHER` | |
| define |`CUDA_ARRAY3D_LAYERED` |`hipArrayLayered` |
| define |`CUDA_ARRAY3D_SURFACE_LDST` |`hipArraySurfaceLoadStore` |
| define |`CUDA_ARRAY3D_TEXTURE_GATHER` |`hipArrayTextureGather` |
| define |`CUDA_ARRAY3D_COLOR_ATTACHMENT` | |
| define |`CUDA_VERSION` | |
| typedef |`CUexternalMemory` | |
@@ -697,8 +697,8 @@
| 1 |*`CU_EGL_FRAME_TYPE_PITCH`* | |
| enum |***`CUeglResourceLocationFlags`*** | |
| typedef |***`CUeglResourceLocationFlags_enum`*** | |
| 0 |*`CU_EGL_RESOURCE_LOCATION_SYSMEM`* | |
| 1 |*`CU_EGL_RESOURCE_LOCATION_VIDMEM`* | |
| 0x00 |*`CU_EGL_RESOURCE_LOCATION_SYSMEM`* | |
| 0x01 |*`CU_EGL_RESOURCE_LOCATION_VIDMEM`* | |
| enum |***`CUexternalMemoryHandleType`*** | |
| typedef |***`CUexternalMemoryHandleType_enum`*** | |
| 1 |*`CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD`* | |
@@ -731,6 +731,8 @@
| define |`CUDA_CB` | |
| define |`CU_DEVICE_CPU` | |
| define |`CU_DEVICE_INVALID` | |
| struct |`CUuuid` | |
| typedef |`CUuuid_st` | |
## **2. Error Handling**
@@ -761,6 +763,7 @@
| `cuDeviceGetName` | `hipDeviceGetName` |
| `cuDeviceTotalMem` | `hipDeviceTotalMem` |
| `cuDeviceGetLuid` | |
| `cuDeviceGetUuid` | |
## **6. Device Management [DEPRECATED]**
@@ -790,9 +793,9 @@
| `cuCtxGetCurrent` | `hipCtxGetCurrent` |
| `cuCtxGetDevice` | `hipCtxGetDevice` |
| `cuCtxGetFlags` | `hipCtxGetFlags` |
| `cuCtxGetLimit` | |
| `cuCtxGetLimit` | `hipDeviceGetLimit` |
| `cuCtxGetSharedMemConfig` | `hipCtxGetSharedMemConfig` |
| `cuCtxGetStreamPriorityRange` | |
| `cuCtxGetStreamPriorityRange` | `hipDeviceGetStreamPriorityRange`|
| `cuCtxPopCurrent` | `hipCtxPopCurrent` |
| `cuCtxPushCurrent` | `hipCtxPushCurrent` |
| `cuCtxSetCacheConfig` | `hipCtxSetCacheConfig` |
@@ -833,16 +836,16 @@
|-----------------------------------------------------------|-------------------------------|
| `cuArray3DCreate` | `hipArray3DCreate` |
| `cuArray3DGetDescriptor` | |
| `cuArrayCreate` | |
| `cuArrayCreate` | `hipArrayCreate` |
| `cuArrayDestroy` | |
| `cuArrayGetDescriptor` | |
| `cuDeviceGetByPCIBusId` | `hipDeviceGetByPCIBusId` |
| `cuDeviceGetPCIBusId` | `hipDeviceGetPCIBusId` |
| `cuIpcCloseMemHandle` | |
| `cuIpcCloseMemHandle` | `hipIpcCloseMemHandle` |
| `cuIpcGetEventHandle` | |
| `cuIpcGetMemHandle` | |
| `cuIpcGetMemHandle` | `hipIpcGetMemHandle` |
| `cuIpcOpenEventHandle` | |
| `cuIpcOpenMemHandle` | |
| `cuIpcOpenMemHandle` | `hipIpcOpenMemHandle` |
| `cuMemAlloc` | `hipMalloc` |
| `cuMemAllocHost` | |
| `cuMemAllocManaged` | |
@@ -865,7 +868,7 @@
| `cuMemcpyDtoDAsync` | `hipMemcpyDtoDAsync` |
| `cuMemcpyDtoH` | `hipMemcpyDtoH` |
| `cuMemcpyDtoHAsync` | `hipMemcpyDtoHAsync` |
| `cuMemcpyHtoA` | |
| `cuMemcpyHtoA` | `hipMemcpyHtoA` |
| `cuMemcpyHtoAAsync` | |
| `cuMemcpyHtoD` | `hipMemcpyHtoD` |
| `cuMemcpyHtoDAsync` | `hipMemcpyHtoDAsync` |
@@ -873,11 +876,11 @@
| `cuMemcpyPeerAsync` | |
| `cuMemFree` | `hipFree` |
| `cuMemFreeHost` | `hipFreeHost` |
| `cuMemGetAddressRange` | |
| `cuMemGetAddressRange` | `hipMemGetAddressRange` |
| `cuMemGetInfo` | `hipMemGetInfo` |
| `cuMemHostAlloc` | `hipHostMalloc` |
| `cuMemHostGetDevicePointer` | |
| `cuMemHostGetFlags` | |
| `cuMemHostGetDevicePointer` | `hipHostGetDevicePointer` |
| `cuMemHostGetFlags` | `hipHostGetFlags` |
| `cuMemHostRegister` | `hipHostRegister` |
| `cuMemHostUnregister` | `hipHostUnregister` |
| `cuMemsetD16` | |
@@ -890,8 +893,8 @@
| `cuMemsetD2D8Async` | |
| `cuMemsetD32` | `hipMemset` |
| `cuMemsetD32Async` | `hipMemsetAsync` |
| `cuMemsetD2D8` | |
| `cuMemsetD2D8Async` | |
| `cuMemsetD8` | `hipMemsetD8` |
| `cuMemsetD8Async` | |
| `cuMipmappedArrayCreate` | |
| `cuMipmappedArrayDestroy` | |
| `cuMipmappedArrayGetLevel` | |
@@ -914,8 +917,8 @@
|-----------------------------------------------------------|-------------------------------|
| `cuStreamAddCallback` | `hipStreamAddCallback` |
| `cuStreamAttachMemAsync` | |
| `cuStreamCreate` | |
| `cuStreamCreateWithPriority` | |
| `cuStreamCreate` | `hipStreamCreateWithFlags` |
| `cuStreamCreateWithPriority` | `hipStreamCreateWithPriority` |
| `cuStreamDestroy` | `hipStreamDestroy` |
| `cuStreamGetFlags` | `hipStreamGetFlags` |
| `cuStreamGetPriority` | `hipStreamGetPriority` |
@@ -930,7 +933,7 @@
| **CUDA** | **HIP** |
|-----------------------------------------------------------|-------------------------------|
| `cuEventCreate` | `hipEventCreate` |
| `cuEventCreate` | `hipEventCreateWithFlags` |
| `cuEventDestroy` | `hipEventDestroy` |
| `cuEventElapsedTime` | `hipEventElapsedTime` |
| `cuEventQuery` | `hipEventQuery` |
@@ -965,10 +968,13 @@
| **CUDA** | **HIP** |
|-----------------------------------------------------------|-------------------------------|
| `cuFuncGetAttribute` | |
| `cuFuncSetAttribute` | |
| `cuFuncSetCacheConfig` | `hipFuncSetCacheConfig` |
| `cuFuncSetSharedMemConfig` | |
| `cuLaunchKernel` | `hipModuleLaunchKernel` |
| `cuLaunchHostFunc` | |
| `cuLaunchCooperativeKernel` | |
| `cuLaunchCooperativeKernelMultiDevice` | |
## **18. Execution Control [DEPRECATED]**
@@ -1045,8 +1051,8 @@
| `cuTexRefGetMipmapLevelBias` | |
| `cuTexRefGetMipmapLevelClamp` | |
| `cuTexRefGetMipmappedArray` | |
| `cuTexRefSetAddress` | |
| `cuTexRefSetAddress2D` | |
| `cuTexRefSetAddress` | `hipTexRefSetAddress` |
| `cuTexRefSetAddress2D` | `hipTexRefSetAddress2D` |
| `cuTexRefSetAddressMode` | `hipTexRefSetAddressMode` |
| `cuTexRefSetArray` | `hipTexRefSetArray` |
| `cuTexRefSetBorderColor` | |
@@ -1231,3 +1237,4 @@
| `cuEGLStreamProducerReturnFrame` | |
| `cuGraphicsEGLRegisterImage` | |
| `cuGraphicsResourceGetMappedEglFrame` | |
| `cuEventCreateFromEGLSync` | |
hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+667 -440
View File
File diff suppressed because it is too large Load Diff
hipamd/docs/markdown/hip-math-api.md
+64 -56
View File
@@ -1433,7 +1433,7 @@ __device__ float __expf(float x);
__device__ static float __fadd_rd(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fadd_rn
@@ -1441,7 +1441,7 @@ __device__ static float __fadd_rd(float x, float y);
__device__ static float __fadd_rn(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fadd_ru
@@ -1449,7 +1449,7 @@ __device__ static float __fadd_rn(float x, float y);
__device__ static float __fadd_ru(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fadd_rz
@@ -1457,7 +1457,7 @@ __device__ static float __fadd_ru(float x, float y);
__device__ static float __fadd_rz(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fdiv_rd
@@ -1465,7 +1465,7 @@ __device__ static float __fadd_rz(float x, float y);
__device__ static float __fdiv_rd(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fdiv_rn
@@ -1473,7 +1473,7 @@ __device__ static float __fdiv_rd(float x, float y);
__device__ static float __fdiv_rn(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fdiv_ru
@@ -1481,7 +1481,7 @@ __device__ static float __fdiv_rn(float x, float y);
__device__ static float __fdiv_ru(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fdiv_rz
@@ -1489,7 +1489,7 @@ __device__ static float __fdiv_ru(float x, float y);
__device__ static float __fdiv_rz(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fdividef
@@ -1505,7 +1505,7 @@ __device__ static float __fdividef(float x, float y);
__device__ float __fmaf_rd(float x, float y, float z);
```
**Description:** Supported
**Description:** Unsupported
### __fmaf_rn
@@ -1513,7 +1513,7 @@ __device__ float __fmaf_rd(float x, float y, float z);
__device__ float __fmaf_rn(float x, float y, float z);
```
**Description:** Supported
**Description:** Unsupported
### __fmaf_ru
@@ -1521,7 +1521,7 @@ __device__ float __fmaf_rn(float x, float y, float z);
__device__ float __fmaf_ru(float x, float y, float z);
```
**Description:** Supported
**Description:** Unsupported
### __fmaf_rz
@@ -1529,7 +1529,7 @@ __device__ float __fmaf_ru(float x, float y, float z);
__device__ float __fmaf_rz(float x, float y, float z);
```
**Description:** Supported
**Description:** Unsupported
### __fmul_rd
@@ -1537,7 +1537,7 @@ __device__ float __fmaf_rz(float x, float y, float z);
__device__ static float __fmul_rd(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fmul_rn
@@ -1545,7 +1545,7 @@ __device__ static float __fmul_rd(float x, float y);
__device__ static float __fmul_rn(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fmul_ru
@@ -1553,7 +1553,7 @@ __device__ static float __fmul_rn(float x, float y);
__device__ static float __fmul_ru(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fmul_rz
@@ -1561,7 +1561,7 @@ __device__ static float __fmul_ru(float x, float y);
__device__ static float __fmul_rz(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __frcp_rd
@@ -1569,7 +1569,7 @@ __device__ static float __fmul_rz(float x, float y);
__device__ float __frcp_rd(float x);
```
**Description:** Supported
**Description:** Unsupported
### __frcp_rn
@@ -1577,7 +1577,7 @@ __device__ float __frcp_rd(float x);
__device__ float __frcp_rn(float x);
```
**Description:** Supported
**Description:** Unsupported
### __frcp_ru
@@ -1585,7 +1585,7 @@ __device__ float __frcp_rn(float x);
__device__ float __frcp_ru(float x);
```
**Description:** Supported
**Description:** Unsupported
### __frcp_rz
@@ -1593,7 +1593,7 @@ __device__ float __frcp_ru(float x);
__device__ float __frcp_rz(float x);
```
**Description:** Supported
**Description:** Unsupported
### __frsqrt_rn
@@ -1601,7 +1601,7 @@ __device__ float __frcp_rz(float x);
__device__ float __frsqrt_rn(float x);
```
**Description:** Supported
**Description:** Unsupported
### __fsqrt_rd
@@ -1609,7 +1609,7 @@ __device__ float __frsqrt_rn(float x);
__device__ float __fsqrt_rd(float x);
```
**Description:** Supported
**Description:** Unsupported
### __fsqrt_rn
@@ -1617,7 +1617,7 @@ __device__ float __fsqrt_rd(float x);
__device__ float __fsqrt_rn(float x);
```
**Description:** Supported
**Description:** Unsupported
### __fsqrt_ru
@@ -1625,7 +1625,7 @@ __device__ float __fsqrt_rn(float x);
__device__ float __fsqrt_ru(float x);
```
**Description:** Supported
**Description:** Unsupported
### __fsqrt_rz
@@ -1633,7 +1633,7 @@ __device__ float __fsqrt_ru(float x);
__device__ float __fsqrt_rz(float x);
```
**Description:** Supported
**Description:** Unsupported
### __fsub_rd
@@ -1641,7 +1641,7 @@ __device__ float __fsqrt_rz(float x);
__device__ static float __fsub_rd(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fsub_rn
@@ -1649,7 +1649,7 @@ __device__ static float __fsub_rd(float x, float y);
__device__ static float __fsub_rn(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fsub_ru
@@ -1657,7 +1657,15 @@ __device__ static float __fsub_rn(float x, float y);
__device__ static float __fsub_ru(float x, float y);
```
**Description:** Supported
**Description:** Unsupported
### __fsub_rz
```cpp
__device__ static float __fsub_rz(float x, float y);
```
**Description:** Unsupported
### __log10f
@@ -1729,7 +1737,7 @@ __device__ float __tanf(float x);
__device__ static double __dadd_rd(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __dadd_rn
@@ -1737,7 +1745,7 @@ __device__ static double __dadd_rd(double x, double y);
__device__ static double __dadd_rn(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __dadd_ru
@@ -1745,7 +1753,7 @@ __device__ static double __dadd_rn(double x, double y);
__device__ static double __dadd_ru(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __dadd_rz
@@ -1753,7 +1761,7 @@ __device__ static double __dadd_ru(double x, double y);
__device__ static double __dadd_rz(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __ddiv_rd
@@ -1761,7 +1769,7 @@ __device__ static double __dadd_rz(double x, double y);
__device__ static double __ddiv_rd(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __ddiv_rn
@@ -1769,7 +1777,7 @@ __device__ static double __ddiv_rd(double x, double y);
__device__ static double __ddiv_rn(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __ddiv_ru
@@ -1777,7 +1785,7 @@ __device__ static double __ddiv_rn(double x, double y);
__device__ static double __ddiv_ru(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __ddiv_rz
@@ -1785,7 +1793,7 @@ __device__ static double __ddiv_ru(double x, double y);
__device__ static double __ddiv_rz(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __dmul_rd
@@ -1793,7 +1801,7 @@ __device__ static double __ddiv_rz(double x, double y);
__device__ static double __dmul_rd(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __dmul_rn
@@ -1801,7 +1809,7 @@ __device__ static double __dmul_rd(double x, double y);
__device__ static double __dmul_rn(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __dmul_ru
@@ -1809,7 +1817,7 @@ __device__ static double __dmul_rn(double x, double y);
__device__ static double __dmul_ru(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __dmul_rz
@@ -1817,7 +1825,7 @@ __device__ static double __dmul_ru(double x, double y);
__device__ static double __dmul_rz(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __drcp_rd
@@ -1825,7 +1833,7 @@ __device__ static double __dmul_rz(double x, double y);
__device__ double __drcp_rd(double x);
```
**Description:** Supported
**Description:** Unsupported
### __drcp_rn
@@ -1833,7 +1841,7 @@ __device__ double __drcp_rd(double x);
__device__ double __drcp_rn(double x);
```
**Description:** Supported
**Description:** Unsupported
### __drcp_ru
@@ -1841,7 +1849,7 @@ __device__ double __drcp_rn(double x);
__device__ double __drcp_ru(double x);
```
**Description:** Supported
**Description:** Unsupported
### __drcp_rz
@@ -1849,7 +1857,7 @@ __device__ double __drcp_ru(double x);
__device__ double __drcp_rz(double x);
```
**Description:** Supported
**Description:** Unsupported
### __dsqrt_rd
@@ -1857,7 +1865,7 @@ __device__ double __drcp_rz(double x);
__device__ double __dsqrt_rd(double x);
```
**Description:** Supported
**Description:** Unsupported
### __dsqrt_rn
@@ -1865,7 +1873,7 @@ __device__ double __dsqrt_rd(double x);
__device__ double __dsqrt_rn(double x);
```
**Description:** Supported
**Description:** Unsupported
### __dsqrt_ru
@@ -1873,7 +1881,7 @@ __device__ double __dsqrt_rn(double x);
__device__ double __dsqrt_ru(double x);
```
**Description:** Supported
**Description:** Unsupported
### __dsqrt_rz
@@ -1881,7 +1889,7 @@ __device__ double __dsqrt_ru(double x);
__device__ double __dsqrt_rz(double x);
```
**Description:** Supported
**Description:** Unsupported
### __dsub_rd
@@ -1889,7 +1897,7 @@ __device__ double __dsqrt_rz(double x);
__device__ static double __dsub_rd(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __dsub_rn
@@ -1897,7 +1905,7 @@ __device__ static double __dsub_rd(double x, double y);
__device__ static double __dsub_rn(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __dsub_ru
@@ -1905,7 +1913,7 @@ __device__ static double __dsub_rn(double x, double y);
__device__ static double __dsub_ru(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __dsub_rz
@@ -1913,7 +1921,7 @@ __device__ static double __dsub_ru(double x, double y);
__device__ static double __dsub_rz(double x, double y);
```
**Description:** Supported
**Description:** Unsupported
### __fma_rd
@@ -1921,7 +1929,7 @@ __device__ static double __dsub_rz(double x, double y);
__device__ double __fma_rd(double x, double y, double z);
```
**Description:** Supported
**Description:** Unsupported
### __fma_rn
@@ -1929,7 +1937,7 @@ __device__ double __fma_rd(double x, double y, double z);
__device__ double __fma_rn(double x, double y, double z);
```
**Description:** Supported
**Description:** Unsupported
### __fma_ru
@@ -1937,7 +1945,7 @@ __device__ double __fma_rn(double x, double y, double z);
__device__ double __fma_ru(double x, double y, double z);
```
**Description:** Supported
**Description:** Unsupported
### __fma_rz
@@ -1945,7 +1953,7 @@ __device__ double __fma_ru(double x, double y, double z);
__device__ double __fma_rz(double x, double y, double z);
```
**Description:** Supported
**Description:** Unsupported
### __brev
hipamd/hipify-clang/README.md
+157 -33
View File
@@ -11,6 +11,7 @@
- [Build and install](#build-and-install)
* [Building](#building)
* [Testing](#testing)
* [Linux](#linux)
* [Windows](#windows)
- [Running and using hipify-clang](#running-and-using-hipify-clang)
- [Disclaimer](#disclaimer)
@@ -80,7 +81,7 @@ Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested; `LLVM+
The binary can then be found at `./dist/bin/hipify-clang`.
### <a name="testing"></a> Test
### <a name="testing"></a> Testing
`hipify-clang` has unit tests using LLVM [`lit`](https://llvm.org/docs/CommandGuide/lit.html)/[`FileCheck`](https://llvm.org/docs/CommandGuide/FileCheck.html).
@@ -88,61 +89,183 @@ The binary can then be found at `./dist/bin/hipify-clang`.
To run it:
1. Download [`LLVM`](http://releases.llvm.org/6.0.1/llvm-6.0.1.src.tar.xz)+[`CLANG`](http://releases.llvm.org/6.0.1/cfe-6.0.1.src.tar.xz) sources.
2. Build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html).
For instance:
```shell
2. Build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html):
```shell
cd llvm
mkdir build dist
cd build
```
- **Linux**:
cmake \
-DCMAKE_INSTALL_PREFIX=../dist \
-DLLVM_SOURCE_DIR=../llvm \
-DCMAKE_BUILD_TYPE=Release \
-Thost=x64 \
../llvm
```shell
cmake \
-DCMAKE_INSTALL_PREFIX=../dist \
-DLLVM_SOURCE_DIR=../llvm \
-DCMAKE_BUILD_TYPE=Release \
../llvm
make -j install
```
- **Windows**:
```shell
cmake \
-G "Visual Studio 15 2017 Win64" \
-DCMAKE_INSTALL_PREFIX=../dist \
-DLLVM_SOURCE_DIR=../llvm \
-DCMAKE_BUILD_TYPE=Release \
-Thost=x64 \
../llvm
```
&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp; Run `Visual Studio 15 2017`, open the generated `LLVM.sln`, build all, build project `INSTALL`.
make -j install
```
On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"`; the generated `LLVM.sln` should be built by `Visual Studio 15 2017` instead of `make`.
3. Ensure [`CUDA`](https://developer.nvidia.com/cuda-toolkit-archive) of minimum version 7.5 is installed.
* Having multiple CUDA installations, in order to choose a particular version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified:
* Having multiple CUDA installations to choose a particular version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified:
`-DCUDA_TOOLKIT_ROOT_DIR="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"`
- Linux: `-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-8.0`
* On Windows `CUDA_SDK_ROOT_DIR` option should be specified as well:
- Windows: `-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"`
`-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0"`
`-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0"`
4. Ensure [`cuDNN`](https://developer.nvidia.com/rdp/cudnn-archive) of version corresponding to CUDA's version is installed.
* Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option:
`-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1`
- Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-8.0-v7.1`
- Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1`
5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed.
6. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM.
* installing `lit` into `python` might be required:
* Install `lit` into `python`:
`python f:/LLVM/6.0.1/llvm/utils/lit/setup.py install`,
- Linux: `python /srv/git/LLVM/6.0.1/llvm/utils/lit/setup.py install`
where `f:/LLVM/6.0.1/llvm` is LLVM sources root directory.
- Windows: `python f:/LLVM/6.0.1/llvm/utils/lit/setup.py install`
* Starting with LLVM 6.0.1 path to llvm-lit.py script should be specified by the `LLVM_EXTERNAL_LIT` option:
* Starting with LLVM 6.0.1 path to `llvm-lit` python script should be specified by the `LLVM_EXTERNAL_LIT` option:
`-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py`,
- Linux: `-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/6.0.1/build/bin/llvm-lit`
where `f:/LLVM/6.0.1/build/Release` is LLVM build directory.
- Windows: `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py`
7. Build with the `HIPIFY_CLANG_TESTS` option turned on: -DHIPIFY_CLANG_TESTS=1.
7. Set `HIPIFY_CLANG_TESTS` option turned on: `-DHIPIFY_CLANG_TESTS=1`.
8. `make test-hipify`
8. Run `cmake`:
* [Linux](#linux)
* [Windows](#windows)
On Windows after `cmake` the project `test-hipify` in the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make test-hipify`.
9. Run tests:
- Linux: `make test-hipify`.
- Windows: run `Visual Studio 15 2017`, open the generated `hipify-clang.sln`, build project `test-hipify`.
### <a name="linux"></a >Linux
On Linux (Ubuntu 14-18) the following configurations are tested:
LLVM 5.0.0 - 6.0.1, CUDA 8.0, cudnn-8.0
Build system for the above configurations:
Python 2.7 (min), cmake 3.5.2 (min), GNU C/C++ 5.4.0 (min).
Here is an example of building `hipify-clang` with testing support on `Ubuntu 16.04`:
```shell
cmake
-DHIPIFY_CLANG_TESTS=1 \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_INSTALL_PREFIX=../dist \
-DCMAKE_PREFIX_PATH=/srv/git/LLVM/6.0.1/dist \
-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-8.0 \
-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-8.0-v7.1 \
-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/6.0.1/build/bin/llvm-lit \
..
```
*A corresponding successful output:*
```shell
-- The C compiler identification is GNU 5.4.0
-- The CXX compiler identification is GNU 5.4.0
-- Check for working C compiler: /usr/bin/cc
-- Check for working C compiler: /usr/bin/cc -- works
-- Detecting C compiler ABI info
-- Detecting C compiler ABI info - done
-- Detecting C compile features
-- Detecting C compile features - done
-- Check for working CXX compiler: /usr/bin/c++
-- Check for working CXX compiler: /usr/bin/c++ -- works
-- Detecting CXX compiler ABI info
-- Detecting CXX compiler ABI info - done
-- Detecting CXX compile features
-- Detecting CXX compile features - done
-- Found LLVM 6.0.1:
-- - CMake module path: /srv/git/LLVM/6.0.1/dist/lib/cmake/llvm
-- - Include path : /srv/git/LLVM/6.0.1/dist/include
-- - Binary path : /srv/git/LLVM/6.0.1/dist/bin
-- Linker detection: GNU ld
-- Found PythonInterp: /usr/bin/python2.7 (found suitable version "2.7.12", minimum required is "2.7")
-- Found lit: /usr/local/bin/lit
-- Found FileCheck: /srv/git/LLVM/6.0.1/dist/bin/FileCheck
-- Looking for pthread.h
-- Looking for pthread.h - found
-- Looking for pthread_create
-- Looking for pthread_create - not found
-- Looking for pthread_create in pthreads
-- Looking for pthread_create in pthreads - not found
-- Looking for pthread_create in pthread
-- Looking for pthread_create in pthread - found
-- Found Threads: TRUE
-- Found CUDA: /usr/local/cuda-8.0 (found version "8.0")
-- Configuring done
-- Generating done
-- Build files have been written to: /srv/git/HIP/hipify-clang/build
```
```shell
make test-hipify
```
*A corresponding successful output:*
```shell
[100%] Running HIPify regression tests
-- Testing: 28 tests, 12 threads --
PASS: hipify :: allocators.cu (1 of 28)
PASS: hipify :: coalescing.cu (2 of 28)
PASS: hipify :: cuDNN/cudnn_softmax.cu (3 of 28)
PASS: hipify :: cuFFT/simple_cufft.cu (4 of 28)
PASS: hipify :: cuComplex/cuComplex_Julia.cu (5 of 28)
PASS: hipify :: cuBLAS/cublas_sgemm_matrix_multiplication.cu (6 of 28)
PASS: hipify :: cuBLAS/cublas_1_based_indexing.cu (7 of 28)
PASS: hipify :: cuBLAS/cublas_0_based_indexing.cu (8 of 28)
PASS: hipify :: axpy.cu (9 of 28)
PASS: hipify :: dynamic_shared_memory.cu (10 of 28)
PASS: hipify :: headers_test_01.cu (11 of 28)
PASS: hipify :: headers_test_02.cu (12 of 28)
PASS: hipify :: headers_test_03.cu (13 of 28)
PASS: hipify :: headers_test_05.cu (14 of 28)
PASS: hipify :: cuDNN/cudnn_convolution_forward.cu (15 of 28)
PASS: hipify :: cuRAND/poisson_api_example.cu (16 of 28)
PASS: hipify :: cudaRegister.cu (17 of 28)
PASS: hipify :: headers_test_06.cu (18 of 28)
PASS: hipify :: headers_test_04.cu (19 of 28)
PASS: hipify :: intro.cu (20 of 28)
PASS: hipify :: headers_test_07.cu (21 of 28)
PASS: hipify :: square.cu (22 of 28)
PASS: hipify :: static_shared_memory.cu (23 of 28)
PASS: hipify :: vec_add.cu (24 of 28)
PASS: hipify :: headers_test_08.cu (25 of 28)
PASS: hipify :: cuRAND/benchmark_curand_generate.cpp (26 of 28)
PASS: hipify :: cuRAND/benchmark_curand_kernel.cpp (27 of 28)
PASS: hipify :: headers_test_09.cu (28 of 28)
Testing Time: 1.71s
Expected Passes : 28
[100%] Built target test-hipify
```
### <a name="windows"></a >Windows
@@ -172,7 +295,7 @@ cmake
-Thost=x64
..
```
A corresponding successful output:
*A corresponding successful output:*
```shell
-- Found LLVM 6.0.1:
-- - CMake module path: F:/LLVM/6.0.1/dist/lib/cmake/llvm
@@ -194,12 +317,13 @@ To process a file, `hipify-clang` needs access to the same headers that would be
For example:
```shell
hipify-clang square.cu -- \
./hipify-clang \
square.cu \
-- \
-x cuda \
--cuda-path=/opt/cuda \
--cuda-gpu-arch=sm_30 \
-isystem /opt/cuda/samples/common/inc
-I /opt/cuda/cuDNN
--cuda-path=/usr/local/cuda-8.0 \
--cuda-gpu-arch=sm_50 \
-isystem /usr/local/cuda-8.0/samples/common/inc
```
`hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you
hipamd/hipify-clang/src/CUDA2HIP.cpp
+2
View File
@@ -51,6 +51,8 @@ const std::map<llvm::StringRef, hipCounter>& CUDA_RENAMES_MAP() {
ret.insert(CUDA_DRIVER_FUNCTION_MAP.begin(), CUDA_DRIVER_FUNCTION_MAP.end());
ret.insert(CUDA_RUNTIME_TYPE_NAME_MAP.begin(), CUDA_RUNTIME_TYPE_NAME_MAP.end());
ret.insert(CUDA_RUNTIME_FUNCTION_MAP.begin(), CUDA_RUNTIME_FUNCTION_MAP.end());
ret.insert(CUDA_COMPLEX_TYPE_NAME_MAP.begin(), CUDA_COMPLEX_TYPE_NAME_MAP.end());
ret.insert(CUDA_COMPLEX_FUNCTION_MAP.begin(), CUDA_COMPLEX_FUNCTION_MAP.end());
ret.insert(CUDA_BLAS_TYPE_NAME_MAP.begin(), CUDA_BLAS_TYPE_NAME_MAP.end());
ret.insert(CUDA_BLAS_FUNCTION_MAP.begin(), CUDA_BLAS_FUNCTION_MAP.end());
ret.insert(CUDA_RAND_TYPE_NAME_MAP.begin(), CUDA_RAND_TYPE_NAME_MAP.end());
hipamd/hipify-clang/src/CUDA2HIP.h
+4
View File
@@ -15,6 +15,10 @@ extern const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP;
extern const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_FUNCTION_MAP;
// Maps the names of CUDA RUNTIME API types to the corresponding HIP types
extern const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP;
// Maps the names of CUDA Complex API types to the corresponding HIP types
extern const std::map<llvm::StringRef, hipCounter> CUDA_COMPLEX_TYPE_NAME_MAP;
// Maps the names of CUDA Complex API functions to the corresponding HIP functions
extern const std::map<llvm::StringRef, hipCounter> CUDA_COMPLEX_FUNCTION_MAP;
// Maps the names of CUDA RUNTIME API functions to the corresponding HIP functions
extern const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP;
// Maps the names of CUDA BLAS API types to the corresponding HIP types
hipamd/hipify-clang/src/CUDA2HIP_BLAS_API_functions.cpp
+445 -445
View File
File diff suppressed because it is too large Load Diff
hipamd/hipify-clang/src/CUDA2HIP_Complex_API_functions.cpp
+28
View File
@@ -0,0 +1,28 @@
#include "CUDA2HIP.h"
// Maps the names of CUDA DRIVER API types to the corresponding HIP types
const std::map<llvm::StringRef, hipCounter> CUDA_COMPLEX_FUNCTION_MAP{
{"cuCrealf", {"hipCrealf", CONV_COMPLEX, API_COMPLEX}},
{"cuCimagf", {"hipCimagf", CONV_COMPLEX, API_COMPLEX}},
{"make_cuFloatComplex", {"make_hipFloatComplex", CONV_COMPLEX, API_COMPLEX}},
{"cuConjf", {"hipConjf", CONV_COMPLEX, API_COMPLEX}},
{"cuCaddf", {"hipCaddf", CONV_COMPLEX, API_COMPLEX}},
{"cuCsubf", {"hipCsubf", CONV_COMPLEX, API_COMPLEX}},
{"cuCmulf", {"hipCmulf", CONV_COMPLEX, API_COMPLEX}},
{"cuCdivf", {"hipCdivf", CONV_COMPLEX, API_COMPLEX}},
{"cuCabsf", {"hipCabsf", CONV_COMPLEX, API_COMPLEX}},
{"cuCreal", {"hipCreal", CONV_COMPLEX, API_COMPLEX}},
{"cuCimag", {"hipCimag", CONV_COMPLEX, API_COMPLEX}},
{"make_cuDoubleComplex", {"make_hipDoubleComplex", CONV_COMPLEX, API_COMPLEX}},
{"cuConj", {"hipConj", CONV_COMPLEX, API_COMPLEX}},
{"cuCadd", {"hipCadd", CONV_COMPLEX, API_COMPLEX}},
{"cuCsub", {"hipCsub", CONV_COMPLEX, API_COMPLEX}},
{"cuCmul", {"hipCmul", CONV_COMPLEX, API_COMPLEX}},
{"cuCdiv", {"hipCdiv", CONV_COMPLEX, API_COMPLEX}},
{"cuCabs", {"hipCabs", CONV_COMPLEX, API_COMPLEX}},
{"make_cuComplex", {"make_hipComplex", CONV_COMPLEX, API_COMPLEX}},
{"cuComplexFloatToDouble", {"hipComplexFloatToDouble", CONV_COMPLEX, API_COMPLEX}},
{"cuComplexDoubleToFloat", {"hipComplexDoubleToFloat", CONV_COMPLEX, API_COMPLEX}},
{"cuCfmaf", {"hipCfmaf", CONV_COMPLEX, API_COMPLEX}},
{"cuCfma", {"hipCfma", CONV_COMPLEX, API_COMPLEX}},
};
hipamd/hipify-clang/src/CUDA2HIP_Complex_API_types.cpp
+8
View File
@@ -0,0 +1,8 @@
#include "CUDA2HIP.h"
// Maps the names of CUDA DRIVER API types to the corresponding HIP types
const std::map<llvm::StringRef, hipCounter> CUDA_COMPLEX_TYPE_NAME_MAP{
{"cuFloatComplex", {"hipFloatComplex", CONV_TYPE, API_COMPLEX}},
{"cuDoubleComplex", {"hipDoubleComplex", CONV_TYPE, API_COMPLEX}},
{"cuComplex", {"hipComplex", CONV_TYPE, API_COMPLEX}},
};
hipamd/hipify-clang/src/CUDA2HIP_DNN_API_functions.cpp
+173 -173
View File
@@ -3,207 +3,207 @@
// Map of all functions
const std::map<llvm::StringRef, hipCounter> CUDA_DNN_FUNCTION_MAP{
{"cudnnGetVersion", {"hipdnnGetVersion", CONV_VERSION, API_DNN}},
{"cudnnGetCudartVersion", {"hipdnnGetCudartVersion", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}},
{"cudnnQueryRuntimeError", {"hipdnnQueryRuntimeError", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetProperty", {"hipdnnGetProperty", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetErrorString", {"hipdnnGetErrorString", CONV_ERROR, API_DNN}},
{"cudnnIm2Col", {"hipdnnIm2Col", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCreate", {"hipdnnCreate", CONV_MATH_FUNC, API_DNN}},
{"cudnnDestroy", {"hipdnnDestroy", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetStream", {"hipdnnSetStream", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetStream", {"hipdnnGetStream", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetCallback", {"hipdnnSetCallback", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetCallback", {"hipdnnGetCallback", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetVersion", {"hipdnnGetVersion", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetCudartVersion", {"hipdnnGetCudartVersion", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnQueryRuntimeError", {"hipdnnQueryRuntimeError", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetProperty", {"hipdnnGetProperty", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetErrorString", {"hipdnnGetErrorString", CONV_LIB_FUNC, API_DNN}},
{"cudnnIm2Col", {"hipdnnIm2Col", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCreate", {"hipdnnCreate", CONV_LIB_FUNC, API_DNN}},
{"cudnnDestroy", {"hipdnnDestroy", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetStream", {"hipdnnSetStream", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetStream", {"hipdnnGetStream", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetCallback", {"hipdnnSetCallback", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetCallback", {"hipdnnGetCallback", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
// cuDNN Tensor functions
{"cudnnCreateTensorDescriptor", {"hipdnnCreateTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetTensor4dDescriptor", {"hipdnnSetTensor4dDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetTensor4dDescriptorEx", {"hipdnnSetTensor4dDescriptorEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetTensor4dDescriptor", {"hipdnnGetTensor4dDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetTensorNdDescriptor", {"hipdnnSetTensorNdDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetTensorNdDescriptorEx", {"hipdnnSetTensorNdDescriptorEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetTensorNdDescriptor", {"hipdnnGetTensorNdDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetTensorSizeInBytes", {"hipdnnGetTensorSizeInBytes", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroyTensorDescriptor", {"hipdnnDestroyTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnTransformTensor", {"hipdnnTransformTensor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnAddTensor", {"hipdnnAddTensor", CONV_MATH_FUNC, API_DNN}},
{"cudnnCreateOpTensorDescriptor", {"hipdnnCreateOpTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetOpTensorDescriptor", {"hipdnnSetOpTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetOpTensorDescriptor", {"hipdnnGetOpTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnDestroyOpTensorDescriptor", {"hipdnnDestroyOpTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnOpTensor", {"hipdnnOpTensor", CONV_MATH_FUNC, API_DNN}},
{"cudnnCreateTensorDescriptor", {"hipdnnCreateTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetTensor4dDescriptor", {"hipdnnSetTensor4dDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetTensor4dDescriptorEx", {"hipdnnSetTensor4dDescriptorEx", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetTensor4dDescriptor", {"hipdnnGetTensor4dDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetTensorNdDescriptor", {"hipdnnSetTensorNdDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetTensorNdDescriptorEx", {"hipdnnSetTensorNdDescriptorEx", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetTensorNdDescriptor", {"hipdnnGetTensorNdDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetTensorSizeInBytes", {"hipdnnGetTensorSizeInBytes", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroyTensorDescriptor", {"hipdnnDestroyTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnTransformTensor", {"hipdnnTransformTensor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnAddTensor", {"hipdnnAddTensor", CONV_LIB_FUNC, API_DNN}},
{"cudnnCreateOpTensorDescriptor", {"hipdnnCreateOpTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetOpTensorDescriptor", {"hipdnnSetOpTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetOpTensorDescriptor", {"hipdnnGetOpTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnDestroyOpTensorDescriptor", {"hipdnnDestroyOpTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnOpTensor", {"hipdnnOpTensor", CONV_LIB_FUNC, API_DNN}},
// cuDNN Reduce Tensor functions
{"cudnnCreateReduceTensorDescriptor", {"hipdnnCreateReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetReduceTensorDescriptor", {"hipdnnSetReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetReduceTensorDescriptor", {"hipdnnGetReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnDestroyReduceTensorDescriptor", {"hipdnnDestroyReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetReductionIndicesSize", {"hipdnnGetReductionIndicesSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetReductionWorkspaceSize", {"hipdnnGetReductionWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
{"cudnnReduceTensor", {"hipdnnReduceTensor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetTensor", {"hipdnnSetTensor", CONV_MATH_FUNC, API_DNN}},
{"cudnnScaleTensor", {"hipdnnScaleTensor", CONV_MATH_FUNC, API_DNN}},
{"cudnnCreateReduceTensorDescriptor", {"hipdnnCreateReduceTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetReduceTensorDescriptor", {"hipdnnSetReduceTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetReduceTensorDescriptor", {"hipdnnGetReduceTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnDestroyReduceTensorDescriptor", {"hipdnnDestroyReduceTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetReductionIndicesSize", {"hipdnnGetReductionIndicesSize", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetReductionWorkspaceSize", {"hipdnnGetReductionWorkspaceSize", CONV_LIB_FUNC, API_DNN}},
{"cudnnReduceTensor", {"hipdnnReduceTensor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetTensor", {"hipdnnSetTensor", CONV_LIB_FUNC, API_DNN}},
{"cudnnScaleTensor", {"hipdnnScaleTensor", CONV_LIB_FUNC, API_DNN}},
// cuDNN Filter functions
{"cudnnCreateFilterDescriptor", {"hipdnnCreateFilterDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetFilter4dDescriptor", {"hipdnnSetFilter4dDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetFilter4dDescriptor", {"hipdnnGetFilter4dDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetFilterNdDescriptor", {"hipdnnSetFilterNdDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetFilterNdDescriptor", {"hipdnnGetFilterNdDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnDestroyFilterDescriptor", {"hipdnnDestroyFilterDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnCreateFilterDescriptor", {"hipdnnCreateFilterDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetFilter4dDescriptor", {"hipdnnSetFilter4dDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetFilter4dDescriptor", {"hipdnnGetFilter4dDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetFilterNdDescriptor", {"hipdnnSetFilterNdDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetFilterNdDescriptor", {"hipdnnGetFilterNdDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnDestroyFilterDescriptor", {"hipdnnDestroyFilterDescriptor", CONV_LIB_FUNC, API_DNN}},
// cuDNN Convolution functions
{"cudnnCreateConvolutionDescriptor", {"hipdnnCreateConvolutionDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetConvolutionMathType", {"hipdnnSetConvolutionMathType", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionMathType", {"hipdnnGetConvolutionMathType", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetConvolutionGroupCount", {"hipdnnSetConvolutionGroupCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetConvolutionGroupCount", {"hipdnnGetConvolutionGroupCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetConvolution2dDescriptor", {"hipdnnSetConvolution2dDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolution2dDescriptor", {"hipdnnGetConvolution2dDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolution2dForwardOutputDim", {"hipdnnGetConvolution2dForwardOutputDim", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetConvolutionNdDescriptor", {"hipdnnSetConvolutionNdDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionNdDescriptor", {"hipdnnGetConvolutionNdDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetConvolutionNdForwardOutputDim", {"hipdnnGetConvolutionNdForwardOutputDim", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroyConvolutionDescriptor", {"hipdnnDestroyConvolutionDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionForwardAlgorithmMaxCount", {"hipdnnGetConvolutionForwardAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindConvolutionForwardAlgorithm", {"hipdnnFindConvolutionForwardAlgorithm", CONV_MATH_FUNC, API_DNN}},
{"cudnnFindConvolutionForwardAlgorithmEx", {"hipdnnFindConvolutionForwardAlgorithmEx", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionForwardAlgorithm", {"hipdnnGetConvolutionForwardAlgorithm", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionForwardAlgorithm_v7", {"hipdnnGetConvolutionForwardAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetConvolutionForwardWorkspaceSize", {"hipdnnGetConvolutionForwardWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
{"cudnnConvolutionForward", {"hipdnnConvolutionForward", CONV_MATH_FUNC, API_DNN}},
{"cudnnConvolutionBiasActivationForward", {"hipdnnConvolutionBiasActivationForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnConvolutionBackwardBias", {"hipdnnConvolutionBackwardBias", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardFilterAlgorithmMaxCount", {"hipdnnGetConvolutionBackwardFilterAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindConvolutionBackwardFilterAlgorithm", {"hipdnnFindConvolutionBackwardFilterAlgorithm", CONV_MATH_FUNC, API_DNN}},
{"cudnnFindConvolutionBackwardFilterAlgorithmEx", {"hipdnnFindConvolutionBackwardFilterAlgorithmEx", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardFilterAlgorithm", {"hipdnnGetConvolutionBackwardFilterAlgorithm", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardFilterAlgorithm_v7", {"hipdnnGetConvolutionBackwardFilterAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetConvolutionBackwardFilterWorkspaceSize", {"hipdnnGetConvolutionBackwardFilterWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
{"cudnnConvolutionBackwardFilter", {"hipdnnConvolutionBackwardFilter", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardDataAlgorithmMaxCount", {"hipdnnGetConvolutionBackwardDataAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindConvolutionBackwardDataAlgorithm", {"hipdnnFindConvolutionBackwardDataAlgorithm", CONV_MATH_FUNC, API_DNN}},
{"cudnnFindConvolutionBackwardDataAlgorithmEx", {"hipdnnFindConvolutionBackwardDataAlgorithmEx", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardDataAlgorithm", {"hipdnnGetConvolutionBackwardDataAlgorithm", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardDataAlgorithm_v7", {"hipdnnGetConvolutionBackwardDataAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetConvolutionBackwardDataWorkspaceSize", {"hipdnnGetConvolutionBackwardDataWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
{"cudnnConvolutionBackwardData", {"hipdnnConvolutionBackwardData", CONV_MATH_FUNC, API_DNN}},
{"cudnnCreateConvolutionDescriptor", {"hipdnnCreateConvolutionDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetConvolutionMathType", {"hipdnnSetConvolutionMathType", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionMathType", {"hipdnnGetConvolutionMathType", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetConvolutionGroupCount", {"hipdnnSetConvolutionGroupCount", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetConvolutionGroupCount", {"hipdnnGetConvolutionGroupCount", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetConvolution2dDescriptor", {"hipdnnSetConvolution2dDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolution2dDescriptor", {"hipdnnGetConvolution2dDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolution2dForwardOutputDim", {"hipdnnGetConvolution2dForwardOutputDim", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetConvolutionNdDescriptor", {"hipdnnSetConvolutionNdDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionNdDescriptor", {"hipdnnGetConvolutionNdDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetConvolutionNdForwardOutputDim", {"hipdnnGetConvolutionNdForwardOutputDim", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroyConvolutionDescriptor", {"hipdnnDestroyConvolutionDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionForwardAlgorithmMaxCount", {"hipdnnGetConvolutionForwardAlgorithmMaxCount", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindConvolutionForwardAlgorithm", {"hipdnnFindConvolutionForwardAlgorithm", CONV_LIB_FUNC, API_DNN}},
{"cudnnFindConvolutionForwardAlgorithmEx", {"hipdnnFindConvolutionForwardAlgorithmEx", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionForwardAlgorithm", {"hipdnnGetConvolutionForwardAlgorithm", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionForwardAlgorithm_v7", {"hipdnnGetConvolutionForwardAlgorithm_v7", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetConvolutionForwardWorkspaceSize", {"hipdnnGetConvolutionForwardWorkspaceSize", CONV_LIB_FUNC, API_DNN}},
{"cudnnConvolutionForward", {"hipdnnConvolutionForward", CONV_LIB_FUNC, API_DNN}},
{"cudnnConvolutionBiasActivationForward", {"hipdnnConvolutionBiasActivationForward", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnConvolutionBackwardBias", {"hipdnnConvolutionBackwardBias", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardFilterAlgorithmMaxCount", {"hipdnnGetConvolutionBackwardFilterAlgorithmMaxCount", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindConvolutionBackwardFilterAlgorithm", {"hipdnnFindConvolutionBackwardFilterAlgorithm", CONV_LIB_FUNC, API_DNN}},
{"cudnnFindConvolutionBackwardFilterAlgorithmEx", {"hipdnnFindConvolutionBackwardFilterAlgorithmEx", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardFilterAlgorithm", {"hipdnnGetConvolutionBackwardFilterAlgorithm", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardFilterAlgorithm_v7", {"hipdnnGetConvolutionBackwardFilterAlgorithm_v7", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetConvolutionBackwardFilterWorkspaceSize", {"hipdnnGetConvolutionBackwardFilterWorkspaceSize", CONV_LIB_FUNC, API_DNN}},
{"cudnnConvolutionBackwardFilter", {"hipdnnConvolutionBackwardFilter", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardDataAlgorithmMaxCount", {"hipdnnGetConvolutionBackwardDataAlgorithmMaxCount", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindConvolutionBackwardDataAlgorithm", {"hipdnnFindConvolutionBackwardDataAlgorithm", CONV_LIB_FUNC, API_DNN}},
{"cudnnFindConvolutionBackwardDataAlgorithmEx", {"hipdnnFindConvolutionBackwardDataAlgorithmEx", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardDataAlgorithm", {"hipdnnGetConvolutionBackwardDataAlgorithm", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetConvolutionBackwardDataAlgorithm_v7", {"hipdnnGetConvolutionBackwardDataAlgorithm_v7", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetConvolutionBackwardDataWorkspaceSize", {"hipdnnGetConvolutionBackwardDataWorkspaceSize", CONV_LIB_FUNC, API_DNN}},
{"cudnnConvolutionBackwardData", {"hipdnnConvolutionBackwardData", CONV_LIB_FUNC, API_DNN}},
// cuDNN Sortmax functions
{"cudnnSoftmaxForward", {"hipdnnSoftmaxForward", CONV_MATH_FUNC, API_DNN}},
{"cudnnSoftmaxBackward", {"hipdnnSoftmaxBackward", CONV_MATH_FUNC, API_DNN}},
{"cudnnSoftmaxForward", {"hipdnnSoftmaxForward", CONV_LIB_FUNC, API_DNN}},
{"cudnnSoftmaxBackward", {"hipdnnSoftmaxBackward", CONV_LIB_FUNC, API_DNN}},
// cuDNN Pooling functions
{"cudnnCreatePoolingDescriptor", {"hipdnnCreatePoolingDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetPooling2dDescriptor", {"hipdnnSetPooling2dDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetPooling2dDescriptor", {"hipdnnGetPooling2dDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetPoolingNdDescriptor", {"hipdnnSetPoolingNdDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetPoolingNdDescriptor", {"hipdnnGetPoolingNdDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetPoolingNdForwardOutputDim", {"hipdnnGetPoolingNdForwardOutputDim", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetPooling2dForwardOutputDim", {"hipdnnGetPooling2dForwardOutputDim", CONV_MATH_FUNC, API_DNN}},
{"cudnnDestroyPoolingDescriptor", {"hipdnnDestroyPoolingDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnPoolingForward", {"hipdnnPoolingForward", CONV_MATH_FUNC, API_DNN}},
{"cudnnPoolingBackward", {"hipdnnPoolingBackward", CONV_MATH_FUNC, API_DNN}},
{"cudnnCreatePoolingDescriptor", {"hipdnnCreatePoolingDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetPooling2dDescriptor", {"hipdnnSetPooling2dDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetPooling2dDescriptor", {"hipdnnGetPooling2dDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetPoolingNdDescriptor", {"hipdnnSetPoolingNdDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetPoolingNdDescriptor", {"hipdnnGetPoolingNdDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetPoolingNdForwardOutputDim", {"hipdnnGetPoolingNdForwardOutputDim", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetPooling2dForwardOutputDim", {"hipdnnGetPooling2dForwardOutputDim", CONV_LIB_FUNC, API_DNN}},
{"cudnnDestroyPoolingDescriptor", {"hipdnnDestroyPoolingDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnPoolingForward", {"hipdnnPoolingForward", CONV_LIB_FUNC, API_DNN}},
{"cudnnPoolingBackward", {"hipdnnPoolingBackward", CONV_LIB_FUNC, API_DNN}},
// cuDNN Activation functions
{"cudnnCreateActivationDescriptor", {"hipdnnCreateActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetActivationDescriptor", {"hipdnnSetActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetActivationDescriptor", {"hipdnnGetActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnDestroyActivationDescriptor", {"hipdnnDestroyActivationDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnActivationForward", {"hipdnnActivationForward", CONV_MATH_FUNC, API_DNN}},
{"cudnnActivationBackward", {"hipdnnActivationBackward", CONV_MATH_FUNC, API_DNN}},
{"cudnnCreateActivationDescriptor", {"hipdnnCreateActivationDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetActivationDescriptor", {"hipdnnSetActivationDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetActivationDescriptor", {"hipdnnGetActivationDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnDestroyActivationDescriptor", {"hipdnnDestroyActivationDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnActivationForward", {"hipdnnActivationForward", CONV_LIB_FUNC, API_DNN}},
{"cudnnActivationBackward", {"hipdnnActivationBackward", CONV_LIB_FUNC, API_DNN}},
// cuDNN LRN functions
{"cudnnCreateLRNDescriptor", {"hipdnnCreateLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetLRNDescriptor", {"hipdnnSetLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetLRNDescriptor", {"hipdnnGetLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnDestroyLRNDescriptor", {"hipdnnDestroyLRNDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnLRNCrossChannelForward", {"hipdnnLRNCrossChannelForward", CONV_MATH_FUNC, API_DNN}},
{"cudnnLRNCrossChannelBackward", {"hipdnnLRNCrossChannelBackward", CONV_MATH_FUNC, API_DNN}},
{"cudnnCreateLRNDescriptor", {"hipdnnCreateLRNDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetLRNDescriptor", {"hipdnnSetLRNDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetLRNDescriptor", {"hipdnnGetLRNDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnDestroyLRNDescriptor", {"hipdnnDestroyLRNDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnLRNCrossChannelForward", {"hipdnnLRNCrossChannelForward", CONV_LIB_FUNC, API_DNN}},
{"cudnnLRNCrossChannelBackward", {"hipdnnLRNCrossChannelBackward", CONV_LIB_FUNC, API_DNN}},
// cuDNN Divisive Normalization functions
{"cudnnDivisiveNormalizationForward", {"hipdnnDivisiveNormalizationForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDivisiveNormalizationBackward", {"hipdnnDivisiveNormalizationBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDivisiveNormalizationForward", {"hipdnnDivisiveNormalizationForward", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDivisiveNormalizationBackward", {"hipdnnDivisiveNormalizationBackward", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
// cuDNN Batch Normalization functions
{"cudnnDeriveBNTensorDescriptor", {"hipdnnDeriveBNTensorDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnBatchNormalizationForwardTraining", {"hipdnnBatchNormalizationForwardTraining", CONV_MATH_FUNC, API_DNN}},
{"cudnnBatchNormalizationForwardInference", {"hipdnnBatchNormalizationForwardInference", CONV_MATH_FUNC, API_DNN}},
{"cudnnBatchNormalizationBackward", {"hipdnnBatchNormalizationBackward", CONV_MATH_FUNC, API_DNN}},
{"cudnnDeriveBNTensorDescriptor", {"hipdnnDeriveBNTensorDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnBatchNormalizationForwardTraining", {"hipdnnBatchNormalizationForwardTraining", CONV_LIB_FUNC, API_DNN}},
{"cudnnBatchNormalizationForwardInference", {"hipdnnBatchNormalizationForwardInference", CONV_LIB_FUNC, API_DNN}},
{"cudnnBatchNormalizationBackward", {"hipdnnBatchNormalizationBackward", CONV_LIB_FUNC, API_DNN}},
// cuDNN Spatial Transformer functions
{"cudnnCreateSpatialTransformerDescriptor", {"hipdnnCreateSpatialTransformerDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetSpatialTransformerNdDescriptor", {"hipdnnSetSpatialTransformerNdDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroySpatialTransformerDescriptor", {"hipdnnDestroySpatialTransformerDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSpatialTfGridGeneratorForward", {"hipdnnSpatialTfGridGeneratorForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSpatialTfGridGeneratorBackward", {"hipdnnSpatialTfGridGeneratorBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSpatialTfSamplerForward", {"hipdnnSpatialTfSamplerForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSpatialTfSamplerBackward", {"hipdnnSpatialTfSamplerBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCreateSpatialTransformerDescriptor", {"hipdnnCreateSpatialTransformerDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetSpatialTransformerNdDescriptor", {"hipdnnSetSpatialTransformerNdDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroySpatialTransformerDescriptor", {"hipdnnDestroySpatialTransformerDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSpatialTfGridGeneratorForward", {"hipdnnSpatialTfGridGeneratorForward", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSpatialTfGridGeneratorBackward", {"hipdnnSpatialTfGridGeneratorBackward", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSpatialTfSamplerForward", {"hipdnnSpatialTfSamplerForward", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSpatialTfSamplerBackward", {"hipdnnSpatialTfSamplerBackward", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
// cuDNN Dropout functions
{"cudnnCreateDropoutDescriptor", {"hipdnnCreateDropoutDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnDestroyDropoutDescriptor", {"hipdnnDestroyDropoutDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnDropoutGetStatesSize", {"hipdnnDropoutGetStatesSize", CONV_MATH_FUNC, API_DNN}},
{"cudnnDropoutGetReserveSpaceSize", {"hipdnnDropoutGetReserveSpaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetDropoutDescriptor", {"hipdnnSetDropoutDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetDropoutDescriptor", {"hipdnnGetDropoutDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnRestoreDropoutDescriptor", {"hipdnnRestoreDropoutDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDropoutForward", {"hipdnnDropoutForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDropoutBackward", {"hipdnnDropoutBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCreateDropoutDescriptor", {"hipdnnCreateDropoutDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnDestroyDropoutDescriptor", {"hipdnnDestroyDropoutDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnDropoutGetStatesSize", {"hipdnnDropoutGetStatesSize", CONV_LIB_FUNC, API_DNN}},
{"cudnnDropoutGetReserveSpaceSize", {"hipdnnDropoutGetReserveSpaceSize", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetDropoutDescriptor", {"hipdnnSetDropoutDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetDropoutDescriptor", {"hipdnnGetDropoutDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnRestoreDropoutDescriptor", {"hipdnnRestoreDropoutDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDropoutForward", {"hipdnnDropoutForward", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDropoutBackward", {"hipdnnDropoutBackward", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
// cuDNN RNN functions
{"cudnnCreateRNNDescriptor", {"hipdnnCreateRNNDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnDestroyRNNDescriptor", {"hipdnnDestroyRNNDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetRNNForwardInferenceAlgorithmMaxCount", {"hipdnnGetRNNForwardInferenceAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindRNNForwardInferenceAlgorithmEx", {"hipdnnFindRNNForwardInferenceAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNForwardTrainingAlgorithmMaxCount", {"hipdnnGetRNNForwardTrainingAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindRNNForwardTrainingAlgorithmEx", {"hipdnnFindRNNForwardTrainingAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNBackwardDataAlgorithmMaxCount", {"hipdnnGetRNNBackwardDataAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindRNNBackwardDataAlgorithmEx", {"hipdnnFindRNNBackwardDataAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNBackwardWeightsAlgorithmMaxCount", {"hipdnnGetRNNBackwardWeightsAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindRNNBackwardWeightsAlgorithmEx", {"hipdnnFindRNNBackwardWeightsAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCreatePersistentRNNPlan", {"hipdnnCreatePersistentRNNPlan", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetPersistentRNNPlan", {"hipdnnSetPersistentRNNPlan", CONV_MATH_FUNC, API_DNN}},
{"cudnnDestroyPersistentRNNPlan", {"hipdnnDestroyPersistentRNNPlan", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetRNNDescriptor", {"hipdnnSetRNNDescriptor", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetRNNDescriptor", {"hipdnnGetRNNDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetRNNProjectionLayers", {"hipdnnSetRNNProjectionLayers", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNProjectionLayers", {"hipdnnGetRNNProjectionLayers", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetRNNAlgorithmDescriptor", {"hipdnnSetRNNAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetRNNMatrixMathType", {"hipdnnSetRNNMatrixMathType", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNMatrixMathType", {"hipdnnGetRNNMatrixMathType", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNWorkspaceSize", {"hipdnnGetRNNWorkspaceSize", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetRNNTrainingReserveSize", {"hipdnnGetRNNTrainingReserveSize", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetRNNParamsSize", {"hipdnnGetRNNParamsSize", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetRNNLinLayerMatrixParams", {"hipdnnGetRNNLinLayerMatrixParams", CONV_MATH_FUNC, API_DNN}},
{"cudnnGetRNNLinLayerBiasParams", {"hipdnnGetRNNLinLayerBiasParams", CONV_MATH_FUNC, API_DNN}},
{"cudnnRNNForwardInference", {"hipdnnRNNForwardInference", CONV_MATH_FUNC, API_DNN}},
{"cudnnRNNForwardTraining", {"hipdnnRNNForwardTraining", CONV_MATH_FUNC, API_DNN}},
{"cudnnRNNBackwardData", {"hipdnnRNNBackwardData", CONV_MATH_FUNC, API_DNN}},
{"cudnnRNNBackwardWeights", {"hipdnnRNNBackwardWeights", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetRNNDescriptor_v5", {"hipdnnSetRNNDescriptor_v5", CONV_MATH_FUNC, API_DNN}},
{"cudnnSetRNNDescriptor_v6", {"hipdnnSetRNNDescriptor_v6", CONV_MATH_FUNC, API_DNN}},
{"cudnnCreateRNNDescriptor", {"hipdnnCreateRNNDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnDestroyRNNDescriptor", {"hipdnnDestroyRNNDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetRNNForwardInferenceAlgorithmMaxCount", {"hipdnnGetRNNForwardInferenceAlgorithmMaxCount", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindRNNForwardInferenceAlgorithmEx", {"hipdnnFindRNNForwardInferenceAlgorithmEx", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNForwardTrainingAlgorithmMaxCount", {"hipdnnGetRNNForwardTrainingAlgorithmMaxCount", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindRNNForwardTrainingAlgorithmEx", {"hipdnnFindRNNForwardTrainingAlgorithmEx", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNBackwardDataAlgorithmMaxCount", {"hipdnnGetRNNBackwardDataAlgorithmMaxCount", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindRNNBackwardDataAlgorithmEx", {"hipdnnFindRNNBackwardDataAlgorithmEx", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNBackwardWeightsAlgorithmMaxCount", {"hipdnnGetRNNBackwardWeightsAlgorithmMaxCount", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnFindRNNBackwardWeightsAlgorithmEx", {"hipdnnFindRNNBackwardWeightsAlgorithmEx", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCreatePersistentRNNPlan", {"hipdnnCreatePersistentRNNPlan", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetPersistentRNNPlan", {"hipdnnSetPersistentRNNPlan", CONV_LIB_FUNC, API_DNN}},
{"cudnnDestroyPersistentRNNPlan", {"hipdnnDestroyPersistentRNNPlan", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetRNNDescriptor", {"hipdnnSetRNNDescriptor", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetRNNDescriptor", {"hipdnnGetRNNDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetRNNProjectionLayers", {"hipdnnSetRNNProjectionLayers", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNProjectionLayers", {"hipdnnGetRNNProjectionLayers", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetRNNAlgorithmDescriptor", {"hipdnnSetRNNAlgorithmDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetRNNMatrixMathType", {"hipdnnSetRNNMatrixMathType", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNMatrixMathType", {"hipdnnGetRNNMatrixMathType", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetRNNWorkspaceSize", {"hipdnnGetRNNWorkspaceSize", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetRNNTrainingReserveSize", {"hipdnnGetRNNTrainingReserveSize", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetRNNParamsSize", {"hipdnnGetRNNParamsSize", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetRNNLinLayerMatrixParams", {"hipdnnGetRNNLinLayerMatrixParams", CONV_LIB_FUNC, API_DNN}},
{"cudnnGetRNNLinLayerBiasParams", {"hipdnnGetRNNLinLayerBiasParams", CONV_LIB_FUNC, API_DNN}},
{"cudnnRNNForwardInference", {"hipdnnRNNForwardInference", CONV_LIB_FUNC, API_DNN}},
{"cudnnRNNForwardTraining", {"hipdnnRNNForwardTraining", CONV_LIB_FUNC, API_DNN}},
{"cudnnRNNBackwardData", {"hipdnnRNNBackwardData", CONV_LIB_FUNC, API_DNN}},
{"cudnnRNNBackwardWeights", {"hipdnnRNNBackwardWeights", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetRNNDescriptor_v5", {"hipdnnSetRNNDescriptor_v5", CONV_LIB_FUNC, API_DNN}},
{"cudnnSetRNNDescriptor_v6", {"hipdnnSetRNNDescriptor_v6", CONV_LIB_FUNC, API_DNN}},
// cuDNN Connectionist Temporal Classification loss functions
{"cudnnCreateCTCLossDescriptor", {"hipdnnCreateCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetCTCLossDescriptor", {"hipdnnSetCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetCTCLossDescriptor", {"hipdnnGetCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroyCTCLossDescriptor", {"hipdnnDestroyCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCTCLoss", {"hipdnnCTCLoss", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetCTCLossWorkspaceSize", {"hipdnnGetCTCLossWorkspaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCreateCTCLossDescriptor", {"hipdnnCreateCTCLossDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetCTCLossDescriptor", {"hipdnnSetCTCLossDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetCTCLossDescriptor", {"hipdnnGetCTCLossDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroyCTCLossDescriptor", {"hipdnnDestroyCTCLossDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCTCLoss", {"hipdnnCTCLoss", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetCTCLossWorkspaceSize", {"hipdnnGetCTCLossWorkspaceSize", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
// cuDNN Algorithm functions
{"cudnnCreateAlgorithmDescriptor", {"hipdnnCreateAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetAlgorithmDescriptor", {"hipdnnSetAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetAlgorithmDescriptor", {"hipdnnGetAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCopyAlgorithmDescriptor", {"hipdnnCopyAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroyAlgorithmDescriptor", {"hipdnnDestroyAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCreateAlgorithmPerformance", {"hipdnnCreateAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetAlgorithmPerformance", {"hipdnnSetAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetAlgorithmPerformance", {"hipdnnGetAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroyAlgorithmPerformance", {"hipdnnDestroyAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetAlgorithmSpaceSize", {"hipdnnGetAlgorithmSpaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSaveAlgorithm", {"hipdnnSaveAlgorithm", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnRestoreAlgorithm", {"hipdnnRestoreAlgorithm", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCreateAlgorithmDescriptor", {"hipdnnCreateAlgorithmDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetAlgorithmDescriptor", {"hipdnnSetAlgorithmDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetAlgorithmDescriptor", {"hipdnnGetAlgorithmDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCopyAlgorithmDescriptor", {"hipdnnCopyAlgorithmDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroyAlgorithmDescriptor", {"hipdnnDestroyAlgorithmDescriptor", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnCreateAlgorithmPerformance", {"hipdnnCreateAlgorithmPerformance", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSetAlgorithmPerformance", {"hipdnnSetAlgorithmPerformance", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetAlgorithmPerformance", {"hipdnnGetAlgorithmPerformance", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnDestroyAlgorithmPerformance", {"hipdnnDestroyAlgorithmPerformance", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnGetAlgorithmSpaceSize", {"hipdnnGetAlgorithmSpaceSize", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnSaveAlgorithm", {"hipdnnSaveAlgorithm", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
{"cudnnRestoreAlgorithm", {"hipdnnRestoreAlgorithm", CONV_LIB_FUNC, API_DNN, HIP_UNSUPPORTED}},
};
hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp
+712 -387
View File
File diff suppressed because it is too large Load Diff
hipamd/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp
+165 -106
View File
@@ -2,7 +2,9 @@
// Maps the names of CUDA DRIVER API types to the corresponding HIP types
const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// 1. Structs
{"CUDA_ARRAY3D_DESCRIPTOR_st", {"HIP_ARRAY3D_DESCRIPTOR", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_ARRAY3D_DESCRIPTOR", {"HIP_ARRAY3D_DESCRIPTOR", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
@@ -41,35 +43,47 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUDA_KERNEL_NODE_PARAMS_st", {"hipKernelNodeParams", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_KERNEL_NODE_PARAMS", {"hipKernelNodeParams", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// no analogue
// NOTE: cudaLaunchParams struct differs
{"CUDA_LAUNCH_PARAMS_st", {"hipLaunchParams", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_LAUNCH_PARAMS", {"hipLaunchParams", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_MEMCPY2D_st", {"hip_Memcpy2D", CONV_TYPE, API_DRIVER}},
{"CUDA_MEMCPY2D", {"hip_Memcpy2D", CONV_TYPE, API_DRIVER}},
// no analogue
{"CUDA_MEMCPY3D_st", {"hip_Memcpy3D", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_MEMCPY3D", {"hip_Memcpy3D", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_MEMCPY3D_PEER_st", {"hip_Memcpy3D_Peer", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_MEMCPY3D_PEER", {"hip_Memcpy3D_Peer", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_MEMSET_NODE_PARAMS_st", {"hipMemsetNodeParams", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_MEMSET_NODE_PARAMS", {"hipMemsetNodeParams", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// cudaMemsetParams
{"CUDA_MEMSET_NODE_PARAMS_st", {"hipMemsetParams", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_MEMSET_NODE_PARAMS", {"hipMemsetParams", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_POINTER_ATTRIBUTE_P2P_TOKENS_st", {"HIP_POINTER_ATTRIBUTE_P2P_TOKENS", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_POINTER_ATTRIBUTE_P2P_TOKENS", {"HIP_POINTER_ATTRIBUTE_P2P_TOKENS", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// no analogue
// NOTE: cudaResourceDesc struct differs
{"CUDA_RESOURCE_DESC_st", {"HIP_RESOURCE_DESC", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_RESOURCE_DESC", {"HIP_RESOURCE_DESC", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// cudaResourceViewDesc
// NOTE: cudaResourceViewDesc hasn't reserved bytes in the end
{"CUDA_RESOURCE_VIEW_DESC_st", {"HIP_RESOURCE_VIEW_DESC", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_RESOURCE_VIEW_DESC", {"HIP_RESOURCE_VIEW_DESC", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// no analogue
// NOTE: cudaTextureDesc differs
{"CUDA_TEXTURE_DESC_st", {"HIP_TEXTURE_DESC", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUDA_TEXTURE_DESC", {"HIP_TEXTURE_DESC", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUdevprop_st", {"hipDeviceProp_t", CONV_TYPE, API_DRIVER}},
{"CUdevprop", {"hipDeviceProp_t", CONV_TYPE, API_DRIVER}},
// no analogue
// NOTE: cudaDeviceProp differs
{"CUdevprop_st", {"hipDeviceProp_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUdevprop", {"hipDeviceProp_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// cudaIpcEventHandle_st
{"CUipcEventHandle_st", {"ihipIpcEventHandle_t", CONV_TYPE, API_DRIVER}},
@@ -91,9 +105,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUctx_st", {"ihipCtx_t", CONV_TYPE, API_DRIVER}},
{"CUcontext", {"hipCtx_t", CONV_TYPE, API_DRIVER}},
// cudaEglStreamConnection
// CUeglStreamConnection_st
{"CUeglStreamConnection_st", {"hipEglStreamConnection", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUeglStreamConnection", {"hipEglStreamConnection", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// cudaEglStreamConnection
{"CUeglStreamConnection", {"hipEglStreamConnection *", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// the same - CUevent_st
{"CUevent_st", {"ihipEvent_t", CONV_TYPE, API_DRIVER}},
@@ -135,9 +150,9 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUgraphNode", {"hipGraphNode", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// cudaMipmappedArray
{"CUmipmappedArray_st", {"hipMipmappedArray_st", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUmipmappedArray_st", {"hipMipmappedArray_st", CONV_TYPE, API_DRIVER}},
// cudaMipmappedArray_t
{"CUmipmappedArray", {"hipMipmappedArray_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUmipmappedArray", {"hipMipmappedArray_t", CONV_TYPE, API_DRIVER}},
// no analogue
{"CUmod_st", {"ihipModule_t", CONV_TYPE, API_DRIVER}},
@@ -148,16 +163,21 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// cudaStream_t
{"CUstream", {"hipStream_t", CONV_TYPE, API_DRIVER}},
// surfaceReference
// NOTE: possibly surfaceReference is analogue
{"CUsurfref_st", {"ihipSurfaceReference_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// no analogue
{"CUsurfref", {"hipSurfaceReference_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// surfaceReference
// textureReference
{"CUtexref_st", {"textureReference", CONV_TYPE, API_DRIVER}},
{"CUtexref", {"hipTextureReference_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// CUuuid_st
// NOTE: the same struct and its name
{"CUuuid_st", {"hipUUID", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUuuid", {"hipUUID", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// 2. Unions
{"CUstreamBatchMemOpParams", {"hipStreamBatchMemOpParams", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUstreamBatchMemOpParams_union", {"hipStreamBatchMemOpParams", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
@@ -170,15 +190,21 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_TR_ADDRESS_MODE_MIRROR", {"HIP_TR_ADDRESS_MODE_MIRROR", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 2
{"CU_TR_ADDRESS_MODE_BORDER", {"HIP_TR_ADDRESS_MODE_BORDER", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 3
{"CUarray_cubemap_face", {"hipArray_cubemap_face", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUarray_cubemap_face_enum", {"hipArray_cubemap_face", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUarray_cubemap_face", {"hipGraphicsCubeFace", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUarray_cubemap_face_enum", {"hipGraphicsCubeFace", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// CUarray_cubemap_face enum values
{"CU_CUBEMAP_FACE_POSITIVE_X", {"HIP_CUBEMAP_FACE_POSITIVE_X", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x00
{"CU_CUBEMAP_FACE_NEGATIVE_X", {"HIP_CUBEMAP_FACE_NEGATIVE_X", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_CUBEMAP_FACE_POSITIVE_Y", {"HIP_CUBEMAP_FACE_POSITIVE_Y", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CU_CUBEMAP_FACE_NEGATIVE_Y", {"HIP_CUBEMAP_FACE_NEGATIVE_Y", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03
{"CU_CUBEMAP_FACE_POSITIVE_Z", {"HIP_CUBEMAP_FACE_POSITIVE_Z", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
{"CU_CUBEMAP_FACE_NEGATIVE_Z", {"HIP_CUBEMAP_FACE_NEGATIVE_Z", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x05
// cudaGraphicsCubeFacePositiveX
{"CU_CUBEMAP_FACE_POSITIVE_X", {"hipGraphicsCubeFacePositiveX", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x00
// cudaGraphicsCubeFaceNegativeX
{"CU_CUBEMAP_FACE_NEGATIVE_X", {"hipGraphicsCubeFaceNegativeX", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
// cudaGraphicsCubeFacePositiveY
{"CU_CUBEMAP_FACE_POSITIVE_Y", {"hipGraphicsCubeFacePositiveY", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
// cudaGraphicsCubeFaceNegativeY
{"CU_CUBEMAP_FACE_NEGATIVE_Y", {"hipGraphicsCubeFaceNegativeY", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03
// cudaGraphicsCubeFacePositiveZ
{"CU_CUBEMAP_FACE_POSITIVE_Z", {"hipGraphicsCubeFacePositiveZ", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
// cudaGraphicsCubeFaceNegativeZ
{"CU_CUBEMAP_FACE_NEGATIVE_Z", {"hipGraphicsCubeFaceNegativeZ", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x05
{"CUarray_format", {"hipArray_format", CONV_TYPE, API_DRIVER}},
{"CUarray_format_enum", {"hipArray_format", CONV_TYPE, API_DRIVER}},
@@ -198,8 +224,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// CUcomputemode enum values
// cudaComputeModeDefault
{"CU_COMPUTEMODE_DEFAULT", {"hipComputeModeDefault", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0
// Deprecated since CUDA 10.0
// cudaComputeModeExclusive
// NOTE: Deprecated since CUDA 10.0
{"CU_COMPUTEMODE_EXCLUSIVE", {"hipComputeModeExclusive", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 1
// cudaComputeModeProhibited
{"CU_COMPUTEMODE_PROHIBITED", {"hipComputeModeProhibited", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 2
@@ -209,15 +235,25 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUctx_flags", {"hipCctx_flags", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUctx_flags_enum", {"hipCctx_flags", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// CUctx_flags enum values
{"CU_CTX_SCHED_AUTO", {"HIP_CTX_SCHED_AUTO", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x00
{"CU_CTX_SCHED_SPIN", {"HIP_CTX_SCHED_SPIN", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_CTX_SCHED_YIELD", {"HIP_CTX_SCHED_YIELD", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CU_CTX_SCHED_BLOCKING_SYNC", {"HIP_CTX_SCHED_BLOCKING_SYNC", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
{"CU_CTX_BLOCKING_SYNC", {"HIP_CTX_BLOCKING_SYNC", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
{"CU_CTX_SCHED_MASK", {"HIP_CTX_SCHED_MASK", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x07
{"CU_CTX_MAP_HOST", {"HIP_CTX_MAP_HOST", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x08
{"CU_CTX_LMEM_RESIZE_TO_MAX", {"HIP_CTX_LMEM_RESIZE_TO_MAX", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x10
{"CU_CTX_FLAGS_MASK", {"HIP_CTX_FLAGS_MASK", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1f
// cudaDeviceScheduleAuto
{"CU_CTX_SCHED_AUTO", {"hipDeviceScheduleAuto", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x00
// cudaDeviceScheduleSpin
{"CU_CTX_SCHED_SPIN", {"hipDeviceScheduleSpin", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x01
// cudaDeviceScheduleYield
{"CU_CTX_SCHED_YIELD", {"hipDeviceScheduleYield", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x02
// cudaDeviceScheduleBlockingSync
{"CU_CTX_SCHED_BLOCKING_SYNC", {"hipDeviceScheduleBlockingSync", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x04
// cudaDeviceBlockingSync
// NOTE: Deprecated since CUDA 4.0 and replaced with CU_CTX_SCHED_BLOCKING_SYNC
{"CU_CTX_BLOCKING_SYNC", {"hipDeviceScheduleBlockingSync", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x04
// cudaDeviceScheduleMask
{"CU_CTX_SCHED_MASK", {"hipDeviceScheduleMask", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x07
// cudaDeviceMapHost
{"CU_CTX_MAP_HOST", {"hipDeviceMapHost", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x08
// cudaDeviceLmemResizeToMax
{"CU_CTX_LMEM_RESIZE_TO_MAX", {"hipDeviceLmemResizeToMax", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x10
// cudaDeviceMask
{"CU_CTX_FLAGS_MASK", {"hipDeviceMask", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1f
// cudaDeviceAttr
{"CUdevice_attribute", {"hipDeviceAttribute_t", CONV_TYPE, API_DRIVER}},
@@ -239,7 +275,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z", {"hipDeviceAttributeMaxGridDimZ", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 7
// cudaDevAttrMaxSharedMemoryPerBlock
{"CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK", {"hipDeviceAttributeMaxSharedMemoryPerBlock", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 8
// Deprecated, use CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK
// no analogue
// NOTE: Deprecated, use CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK
{"CU_DEVICE_ATTRIBUTE_SHARED_MEMORY_PER_BLOCK", {"hipDeviceAttributeMaxSharedMemoryPerBlock", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 8
// cudaDevAttrTotalConstantMemory
{"CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY", {"hipDeviceAttributeTotalConstantMemory", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 9
@@ -255,8 +292,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_DEVICE_ATTRIBUTE_CLOCK_RATE", {"hipDeviceAttributeClockRate", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 13
// cudaDevAttrTextureAlignment
{"CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT", {"hipDeviceAttributeTextureAlignment", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 14
// Deprecated, use instead CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT
// cudaDevAttrGpuOverlap
// NOTE: Deprecated, use instead CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT
{"CU_DEVICE_ATTRIBUTE_GPU_OVERLAP", {"hipDeviceAttributeAsyncEngineCount", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 15
// cudaDevAttrMultiProcessorCount
{"CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT", {"hipDeviceAttributeMultiprocessorCount", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 16
@@ -286,14 +323,14 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT", {"hipDeviceAttributeMaxTexture2DLayeredHeight", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 28
// cudaDevAttrMaxTexture2DLayeredLayers
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_LAYERS", {"hipDeviceAttributeMaxTexture2DLayeredLayers", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 29
// Deprecated, use CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH
// cudaDevAttrMaxTexture2DLayeredWidth
// NOTE: Deprecated, use CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_WIDTH", {"hipDeviceAttributeMaxTexture2DLayeredWidth", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 27
// Deprecated, use CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT
// cudaDevAttrMaxTexture2DLayeredHeight
// NOTE: Deprecated, use CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_HEIGHT", {"hipDeviceAttributeMaxTexture2DLayeredHeight", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 28
// Deprecated, use CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_LAYERS
// cudaDevAttrMaxTexture2DLayeredLayers
// NOTE: Deprecated, use CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_LAYERS
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_NUMSLICES", {"hipDeviceAttributeMaxTexture2DLayeredLayers", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 29
// cudaDevAttrSurfaceAlignment
{"CU_DEVICE_ATTRIBUTE_SURFACE_ALIGNMENT", {"hipDeviceAttributeSurfaceAlignment", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 30
@@ -323,8 +360,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_WIDTH", {"hipDeviceAttributeMaxTexture1DLayeredWidth", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 42
// cudaDevAttrMaxTexture1DLayeredLayers
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_LAYERS", {"hipDeviceAttributeMaxTexture1DLayeredLayers", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 43
// Deprecated, do not use
// no analogue
// NOTE: Deprecated, do not use
{"CU_DEVICE_ATTRIBUTE_CAN_TEX2D_GATHER", {"hipDeviceAttributeCanTex2DGather", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 44
// cudaDevAttrMaxTexture2DGatherWidth
{"CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_WIDTH", {"hipDeviceAttributeMaxTexture2DGatherWidth", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 45
@@ -454,9 +491,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// cudaDevP2PAttrNativeAtomicSupported = 3
{"CU_DEVICE_P2P_ATTRIBUTE_NATIVE_ATOMIC_SUPPORTED", {"hipDeviceP2PAttributeNativeAtomicSupported", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03
// cudaDevP2PAttrCudaArrayAccessSupported = 4
{"CU_DEVICE_P2P_ATTRIBUTE_ARRAY_ACCESS_ACCESS_SUPPORTED", {"hipDeviceP2PAttributeArrayAccessSupported", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
// NOTE" deprecated, use CU_DEVICE_P2P_ATTRIBUTE_CUDA_ARRAY_ACCESS_SUPPORTED instead
{"CU_DEVICE_P2P_ATTRIBUTE_ARRAY_ACCESS_ACCESS_SUPPORTED", {"hipDevP2PAttributeCudaArrayAccessSupported", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
// cudaDevP2PAttrCudaArrayAccessSupported = 4
{"CU_DEVICE_P2P_ATTRIBUTE_CUDA_ARRAY_ACCESS_SUPPORTED", {"hipDeviceP2PAttributeArrayAccessSupported", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
{"CU_DEVICE_P2P_ATTRIBUTE_CUDA_ARRAY_ACCESS_SUPPORTED", {"hipDevP2PAttributeCudaArrayAccessSupported", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
// cudaEGL.h - presented only on Linux in nvidia-cuda-dev package
// cudaEglColorFormat
@@ -657,8 +695,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_RESOURCE", {"hipExternalMemoryHandleTypeD3D12Resource", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 5
// cudaExternalSemaphoreHandleType
{"CUexternalSemaphoreHandleType_enum", {"hipExternalSemaphoreHandleType", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUexternalSemaphoreHandleType", {"hipExternalSemaphoreHandleType", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUexternalSemaphoreHandleType_enum", {"hipExternalSemaphoreHandleType", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// CUexternalSemaphoreHandleType enum values
// cudaExternalSemaphoreHandleTypeOpaqueFd
{"CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD", {"hipExternalSemaphoreHandleTypeOpaqueFD", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 1
@@ -679,8 +717,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_TR_FILTER_MODE_LINEAR", {"hipFilterModeLinear", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 1
// cudaFuncCache
{"CUfunc_cache", {"hipFuncCache", CONV_TYPE, API_DRIVER}},
{"CUfunc_cache_enum", {"hipFuncCache", CONV_TYPE, API_DRIVER}},
{"CUfunc_cache", {"hipFuncCache_t", CONV_TYPE, API_DRIVER}},
{"CUfunc_cache_enum", {"hipFuncCache_t", CONV_TYPE, API_DRIVER}},
// CUfunc_cache enum values
// cudaFilterModePoint = 0
{"CU_FUNC_CACHE_PREFER_NONE", {"hipFuncCachePreferNone", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x00
@@ -715,7 +753,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES", {"hipFuncAttributeMaxDynamicSharedMemorySize", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 8
// cudaFuncAttributePreferredSharedMemoryCarveout
{"CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT", {"hipFuncAttributePreferredSharedMemoryCarveout", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 9
// no analogue
// cudaFuncAttributeMax
{"CU_FUNC_ATTRIBUTE_MAX", {"hipFuncAttributeMax", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 10
//cudaGraphicsMapFlags
@@ -785,9 +823,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUjit_fallback", {"hipJitFallback", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUjit_fallback_enum", {"hipJitFallback", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// CUjit_fallback enum values
// no analogue
{"CU_PREFER_PTX", {"hipJitFallbackPreferPtx", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0
// no analogue
{"CU_PREFER_BINARY", {"hipJitFallbackPreferBinary", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}},
// no analogue
@@ -821,13 +857,13 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUjit_target", {"hipJitTarget", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUjit_target_enum", {"hipJitTarget", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// CUjit_target enum values
// Deprecated
// NOTE: Deprecated
{"CU_TARGET_COMPUTE_10", {"hipJitTargetCompute10", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 10
// Deprecated
// NOTE: Deprecated
{"CU_TARGET_COMPUTE_11", {"hipJitTargetCompute11", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 11
// Deprecated
// NOTE: Deprecated
{"CU_TARGET_COMPUTE_12", {"hipJitTargetCompute12", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 12
// Deprecated
// NOTE: Deprecated
{"CU_TARGET_COMPUTE_13", {"hipJitTargetCompute13", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 13
{"CU_TARGET_COMPUTE_20", {"hipJitTargetCompute20", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 20
{"CU_TARGET_COMPUTE_21", {"hipJitTargetCompute21", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 21
@@ -842,7 +878,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_TARGET_COMPUTE_61", {"hipJitTargetCompute61", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 61
{"CU_TARGET_COMPUTE_62", {"hipJitTargetCompute62", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 62
{"CU_TARGET_COMPUTE_70", {"hipJitTargetCompute70", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 70
// Deprecated
// NOTE: Deprecated
{"CU_TARGET_COMPUTE_73", {"hipJitTargetCompute73", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 73
{"CU_TARGET_COMPUTE_75", {"hipJitTargetCompute75", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 75
@@ -905,14 +941,14 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_MEM_ATTACH_SINGLE", {"hipMemAttachSingle", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x4
// no analogue
// NOTE: cudaMemoryType is not an analogue
{"CUmemorytype", {"hipMemType_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUmemorytype_enum", {"hipMemType_t", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// NOTE: cudaMemoryType is partial analogue
{"CUmemorytype", {"hipMemoryType", CONV_TYPE, API_DRIVER}},
{"CUmemorytype_enum", {"hipMemoryType", CONV_TYPE, API_DRIVER}},
// CUmemorytype enum values
{"CU_MEMORYTYPE_HOST", {"hipMemTypeHost", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_MEMORYTYPE_DEVICE", {"hipMemTypeDevice", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CU_MEMORYTYPE_ARRAY", {"hipMemTypeArray", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03
{"CU_MEMORYTYPE_UNIFIED", {"hipMemTypeUnified", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
{"CU_MEMORYTYPE_HOST", {"hipMemoryTypeHost", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x01
{"CU_MEMORYTYPE_DEVICE", {"hipMemoryTypeDevice", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x02
{"CU_MEMORYTYPE_ARRAY", {"hipMemoryTypeArray", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x03
{"CU_MEMORYTYPE_UNIFIED", {"hipMemoryTypeUnified", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x04
// cudaMemRangeAttribute
{"CUmem_range_attribute", {"hipMemRangeAttribute", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
@@ -936,9 +972,9 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// cudaOccupancyDisableCachingOverride
{"CU_OCCUPANCY_DISABLE_CACHING_OVERRIDE", {"hipOccupancyDisableCachingOverride", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
//no analogue
// TODO: Analogous enum is needed in HIP. Couldn't map enum to struct hipPointerAttribute_t.
// TODO: Do the same for Pointer Attributes as for Device Attributes.
//no analogue
{"CUpointer_attribute", {"hipPointerAttribute", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUpointer_attribute_enum", {"hipPointerAttribute", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// CUpointer_attribute enum values
@@ -953,17 +989,17 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL", {"hipPointerAttributeDeviceOrdinal", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 9
// cudaResourceType
{"CUresourcetype", {"hipResourceType", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUresourcetype_enum", {"hipResourceType", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUresourcetype", {"hipResourceType", CONV_TYPE, API_DRIVER}},
{"CUresourcetype_enum", {"hipResourceType", CONV_TYPE, API_DRIVER}},
// CUresourcetype enum values
// cudaResourceTypeArray
{"CU_RESOURCE_TYPE_ARRAY", {"hipResourceTypeArray", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x00
{"CU_RESOURCE_TYPE_ARRAY", {"hipResourceTypeArray", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x00
//cudaResourceTypeMipmappedArray
{"CU_RESOURCE_TYPE_MIPMAPPED_ARRAY", {"hipResourceTypeMipmappedArray", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_RESOURCE_TYPE_MIPMAPPED_ARRAY", {"hipResourceTypeMipmappedArray", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x01
//cudaResourceTypeLinear
{"CU_RESOURCE_TYPE_LINEAR", {"hipResourceTypeLinear", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CU_RESOURCE_TYPE_LINEAR", {"hipResourceTypeLinear", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x02
//cudaResourceTypePitch2D
{"CU_RESOURCE_TYPE_PITCH2D", {"hipResourceTypePitch2D", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03
{"CU_RESOURCE_TYPE_PITCH2D", {"hipResourceTypePitch2D", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0x03
// cudaResourceViewFormat
{"CUresourceViewFormat", {"hipResourceViewFormat", CONV_TYPE, API_DRIVER}},
@@ -1058,14 +1094,14 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUDA_ERROR_DEINITIALIZED", {"hipErrorDeinitialized", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 4
// cudaErrorProfilerDisabled = 55
{"CUDA_ERROR_PROFILER_DISABLED", {"hipErrorProfilerDisabled", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 5
// Deprecated since CUDA 5.0
// cudaErrorProfilerNotInitialized = 56
// NOTE: Deprecated since CUDA 5.0
{"CUDA_ERROR_PROFILER_NOT_INITIALIZED", {"hipErrorProfilerNotInitialized", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 6
// Deprecated since CUDA 5.0
// NOTE: Deprecated since CUDA 5.0
// cudaErrorProfilerAlreadyStarted = 57
{"CUDA_ERROR_PROFILER_ALREADY_STARTED", {"hipErrorProfilerAlreadyStarted", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 7
// Deprecated since CUDA 5.0
// cudaErrorProfilerAlreadyStopped = 58
// NOTE: Deprecated since CUDA 5.0
{"CUDA_ERROR_PROFILER_ALREADY_STOPPED", {"hipErrorProfilerAlreadyStopped", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 8
// cudaErrorNoDevice = 38
{"CUDA_ERROR_NO_DEVICE", {"hipErrorNoDevice", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 100
@@ -1076,8 +1112,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUDA_ERROR_INVALID_IMAGE", {"hipErrorInvalidImage", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 200
// no analogue
{"CUDA_ERROR_INVALID_CONTEXT", {"hipErrorInvalidContext", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 201
// Deprecated since CUDA 3.2
// no analogue
// NOTE: Deprecated since CUDA 3.2
{"CUDA_ERROR_CONTEXT_ALREADY_CURRENT", {"hipErrorContextAlreadyCurrent", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 202
// cudaErrorMapBufferObjectFailed = 14
// TODO: double check the matching
@@ -1268,22 +1304,22 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUGLDeviceList_enum", {"hipGLDeviceList", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// CUGLDeviceList enum values
// cudaGLDeviceListAll = 1
{"CU_GL_DEVICE_LIST_ALL", {"HIP_GL_DEVICE_LIST_ALL", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_GL_DEVICE_LIST_ALL", {"hipGLDeviceListAll", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
// cudaGLDeviceListCurrentFrame = 2
{"CU_GL_DEVICE_LIST_CURRENT_FRAME", {"HIP_GL_DEVICE_LIST_CURRENT_FRAME", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CU_GL_DEVICE_LIST_CURRENT_FRAME", {"hipGLDeviceListCurrentFrame", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
// cudaGLDeviceListNextFrame = 3
{"CU_GL_DEVICE_LIST_NEXT_FRAME", {"HIP_GL_DEVICE_LIST_NEXT_FRAME", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03
{"CU_GL_DEVICE_LIST_NEXT_FRAME", {"hipGLDeviceListNextFrame", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03
// cudaGLMapFlags
{"CUGLmap_flags", {"hipGLMapFlags", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUGLmap_flags_enum", {"hipGLMapFlags", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// CUGLmap_flags enum values
// cudaGLMapFlagsNone = 0
{"CU_GL_MAP_RESOURCE_FLAGS_NONE", {"HIP_GL_MAP_RESOURCE_FLAGS_NONE", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x00
{"CU_GL_MAP_RESOURCE_FLAGS_NONE", {"hipGLMapFlagsNone", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x00
// cudaGLMapFlagsReadOnly = 1
{"CU_GL_MAP_RESOURCE_FLAGS_READ_ONLY", {"HIP_GL_MAP_RESOURCE_FLAGS_READ_ONLY", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_GL_MAP_RESOURCE_FLAGS_READ_ONLY", {"hipGLMapFlagsReadOnly", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
// cudaGLMapFlagsWriteDiscard = 2
{"CU_GL_MAP_RESOURCE_FLAGS_WRITE_DISCARD", {"HIP_GL_MAP_RESOURCE_FLAGS_WRITE_DISCARD", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CU_GL_MAP_RESOURCE_FLAGS_WRITE_DISCARD", {"hipGLMapFlagsWriteDiscard", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
// cudaD3D9DeviceList
{"CUd3d9DeviceList", {"hipD3D9DeviceList", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
@@ -1296,8 +1332,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// cudaD3D9DeviceListNextFrame = 3
{"CU_D3D9_DEVICE_LIST_NEXT_FRAME", {"HIP_D3D9_DEVICE_LIST_NEXT_FRAME", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03
// Deprecated
// cudaD3D9MapFlags
// NOTE: Deprecated
{"CUd3d9map_flags", {"hipD3D9MapFlags", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
{"CUd3d9map_flags_enum", {"hipD3D9MapFlags", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}},
// CUd3d9map_flags enum values
@@ -1360,6 +1396,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CU_D3D11_DEVICE_LIST_NEXT_FRAME", {"HIP_D3D11_DEVICE_LIST_NEXT_FRAME", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03
// 4. Typedefs
// no analogue
{"CUdevice", {"hipDevice_t", CONV_TYPE, API_DRIVER}},
{"CUdeviceptr", {"hipDeviceptr_t", CONV_TYPE, API_DRIVER}},
@@ -1380,37 +1417,59 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
{"CUtexObject", {"hipTextureObject_t", CONV_TYPE, API_DRIVER}},
// 5. Defines
{"__CUDACC__", {"__HIPCC__", CONV_DEF, API_DRIVER}},
{"CUDA_CB", {"HIP_CB", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}},
{"CU_DEVICE_CPU", {"HIP_DEVICE_CPU", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // ((CUdevice)-1)
{"CU_DEVICE_INVALID", {"HIP_DEVICE_INVALID", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // ((CUdevice)-2)
{"CU_IPC_HANDLE_SIZE", {"HIP_LAUNCH_PARAM_END", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 64
{"CU_LAUNCH_PARAM_BUFFER_POINTER", {"HIP_LAUNCH_PARAM_BUFFER_POINTER", CONV_DEF, API_DRIVER}}, // ((void*)0x01)
{"CU_LAUNCH_PARAM_BUFFER_SIZE", {"HIP_LAUNCH_PARAM_BUFFER_SIZE", CONV_DEF, API_DRIVER}}, // ((void*)0x02)
{"CU_LAUNCH_PARAM_END", {"HIP_LAUNCH_PARAM_END", CONV_DEF, API_DRIVER}}, // ((void*)0x00)
{"CU_MEMHOSTALLOC_DEVICEMAP", {"HIP_MEMHOSTALLOC_DEVICEMAP", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CU_MEMHOSTALLOC_PORTABLE", {"HIP_MEMHOSTALLOC_PORTABLE", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_MEMHOSTALLOC_WRITECOMBINED", {"HIP_MEMHOSTALLOC_WRITECOMBINED", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
{"CU_MEMHOSTREGISTER_DEVICEMAP", {"HIP_MEMHOSTREGISTER_DEVICEMAP", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CU_MEMHOSTREGISTER_IOMEMORY", {"HIP_MEMHOSTREGISTER_IOMEMORY", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
{"CU_MEMHOSTREGISTER_PORTABLE", {"HIP_MEMHOSTREGISTER_PORTABLE", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_PARAM_TR_DEFAULT", {"HIP_PARAM_TR_DEFAULT", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // -1
{"CU_STREAM_LEGACY", {"HIP_STREAM_LEGACY", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // ((CUstream)0x1)
{"CU_STREAM_PER_THREAD", {"HIP_STREAM_PER_THREAD", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // ((CUstream)0x2)
{"CU_TRSA_OVERRIDE_FORMAT", {"HIP_TRSA_OVERRIDE_FORMAT", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_TRSF_NORMALIZED_COORDINATES", {"HIP_TRSF_NORMALIZED_COORDINATES", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CU_TRSF_READ_AS_INTEGER", {"HIP_TRSF_READ_AS_INTEGER", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_TRSF_SRGB", {"HIP_TRSF_SRGB", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x10
// Deprecated, use CUDA_ARRAY3D_LAYERED
{"CUDA_ARRAY3D_2DARRAY", {"HIP_ARRAY3D_2DARRAY", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CUDA_ARRAY3D_COLOR_ATTACHMENT", {"HIP_ARRAY3D_COLOR_ATTACHMENT", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x20
{"CUDA_ARRAY3D_CUBEMAP", {"HIP_ARRAY3D_CUBEMAP", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x04
{"CUDA_ARRAY3D_DEPTH_TEXTURE", {"HIP_ARRAY3D_DEPTH_TEXTURE", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x10
{"CUDA_ARRAY3D_LAYERED", {"HIP_ARRAY3D_LAYERED", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CUDA_ARRAY3D_SURFACE_LDST", {"HIP_ARRAY3D_SURFACE_LDST", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CUDA_ARRAY3D_TEXTURE_GATHER", {"HIP_ARRAY3D_TEXTURE_GATHER", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x08
{"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC", {"HIP_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC", {"HIP_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CUDA_EXTERNAL_MEMORY_DEDICATED", {"HIP_EXTERNAL_MEMORY_DEDICATED", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1
{"CUDA_VERSION", {"HIP_VERSION", CONV_DEF, API_DRIVER, HIP_UNSUPPORTED}}, // 10000
{"__CUDACC__", {"__HIPCC__", CONV_DEFINE, API_DRIVER}},
{"CUDA_CB", {"HIP_CB", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}},
// cudaCpuDeviceId ((int)-1)
{"CU_DEVICE_CPU", {"hipCpuDeviceId", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // ((CUdevice)-1)
// cudaInvalidDeviceId ((int)-1)
{"CU_DEVICE_INVALID", {"hipInvalidDeviceId", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // ((CUdevice)-2)
// CUDA_IPC_HANDLE_SIZE
{"CU_IPC_HANDLE_SIZE", {"HIP_IPC_HANDLE_SIZE", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 64
{"CU_LAUNCH_PARAM_BUFFER_POINTER", {"HIP_LAUNCH_PARAM_BUFFER_POINTER", CONV_DEFINE, API_DRIVER}}, // ((void*)0x01)
{"CU_LAUNCH_PARAM_BUFFER_SIZE", {"HIP_LAUNCH_PARAM_BUFFER_SIZE", CONV_DEFINE, API_DRIVER}}, // ((void*)0x02)
{"CU_LAUNCH_PARAM_END", {"HIP_LAUNCH_PARAM_END", CONV_DEFINE, API_DRIVER}}, // ((void*)0x00)
// cudaHostAllocPortable
{"CU_MEMHOSTALLOC_PORTABLE", {"hipHostMallocPortable", CONV_DEFINE, API_DRIVER}}, // 0x01
// cudaHostAllocMapped
{"CU_MEMHOSTALLOC_DEVICEMAP", {"hipHostMallocMapped", CONV_DEFINE, API_DRIVER}}, // 0x02
// cudaHostAllocWriteCombined
{"CU_MEMHOSTALLOC_WRITECOMBINED", {"hipHostAllocWriteCombined", CONV_DEFINE, API_DRIVER}}, // 0x04
// cudaHostRegisterPortable
{"CU_MEMHOSTREGISTER_PORTABLE", {"hipHostRegisterPortable", CONV_DEFINE, API_DRIVER}}, // 0x01
// cudaHostRegisterMapped
{"CU_MEMHOSTREGISTER_DEVICEMAP", {"hipHostRegisterMapped", CONV_DEFINE, API_DRIVER}}, // 0x02
// cudaHostRegisterIoMemory
{"CU_MEMHOSTREGISTER_IOMEMORY", {"hipHostRegisterIoMemory", CONV_DEFINE, API_DRIVER}}, // 0x04
{"CU_PARAM_TR_DEFAULT", {"HIP_PARAM_TR_DEFAULT", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // -1
// cudaStreamLegacy ((cudaStream_t)0x1)
{"CU_STREAM_LEGACY", {"hipStreamLegacy", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // ((CUstream)0x1)
// cudaStreamPerThread ((cudaStream_t)0x2)
{"CU_STREAM_PER_THREAD", {"hipStreamPerThread", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // ((CUstream)0x2)
{"CU_TRSA_OVERRIDE_FORMAT", {"HIP_TRSA_OVERRIDE_FORMAT", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_TRSF_NORMALIZED_COORDINATES", {"HIP_TRSF_NORMALIZED_COORDINATES", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
{"CU_TRSF_READ_AS_INTEGER", {"HIP_TRSF_READ_AS_INTEGER", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CU_TRSF_SRGB", {"HIP_TRSF_SRGB", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x10
// no analogue
// NOTE: Deprecated, use CUDA_ARRAY3D_LAYERED
{"CUDA_ARRAY3D_2DARRAY", {"HIP_ARRAY3D_2DARRAY", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
// cudaArrayLayered
{"CUDA_ARRAY3D_LAYERED", {"hipArrayLayered", CONV_DEFINE, API_DRIVER}}, // 0x01
// cudaArraySurfaceLoadStore
{"CUDA_ARRAY3D_SURFACE_LDST", {"hipArraySurfaceLoadStore", CONV_DEFINE, API_DRIVER}}, // 0x02
// cudaArrayCubemap
{"CUDA_ARRAY3D_CUBEMAP", {"hipArrayCubemap", CONV_DEFINE, API_DRIVER}}, // 0x04
// cudaArrayTextureGather
{"CUDA_ARRAY3D_TEXTURE_GATHER", {"hipArrayTextureGather", CONV_DEFINE, API_DRIVER}}, // 0x08
// no analogue
{"CUDA_ARRAY3D_DEPTH_TEXTURE", {"hipArrayDepthTexture", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x10
// cudaArrayColorAttachment
{"CUDA_ARRAY3D_COLOR_ATTACHMENT", {"hipArrayColorAttachment", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x20
// cudaCooperativeLaunchMultiDeviceNoPreSync
{"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_PRE_LAUNCH_SYNC", {"hipCooperativeLaunchMultiDeviceNoPreSync", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
// cudaCooperativeLaunchMultiDeviceNoPostSync
{"CUDA_COOPERATIVE_LAUNCH_MULTI_DEVICE_NO_POST_LAUNCH_SYNC", {"hipCooperativeLaunchMultiDeviceNoPostSync", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02
// cudaExternalMemoryDedicated
{"CUDA_EXTERNAL_MEMORY_DEDICATED", {"hipExternalMemoryDedicated", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01
{"CUDA_VERSION", {"HIP_VERSION", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 10000
};
hipamd/hipify-clang/src/CUDA2HIP_FFT_API_functions.cpp
+32 -32
View File
@@ -2,36 +2,36 @@
// Map of all functions
const std::map<llvm::StringRef, hipCounter> CUDA_FFT_FUNCTION_MAP{
{"cufftPlan1d", {"hipfftPlan1d", CONV_MATH_FUNC, API_FFT}},
{"cufftPlan2d", {"hipfftPlan2d", CONV_MATH_FUNC, API_FFT}},
{"cufftPlan3d", {"hipfftPlan3d", CONV_MATH_FUNC, API_FFT}},
{"cufftPlanMany", {"hipfftPlanMany", CONV_MATH_FUNC, API_FFT}},
{"cufftMakePlan1d", {"hipfftMakePlan1d", CONV_MATH_FUNC, API_FFT}},
{"cufftMakePlan2d", {"hipfftMakePlan2d", CONV_MATH_FUNC, API_FFT}},
{"cufftMakePlan3d", {"hipfftMakePlan3d", CONV_MATH_FUNC, API_FFT}},
{"cufftMakePlanMany", {"hipfftMakePlanMany", CONV_MATH_FUNC, API_FFT}},
{"cufftMakePlanMany64", {"hipfftMakePlanMany64", CONV_MATH_FUNC, API_FFT}},
{"cufftGetSizeMany64", {"hipfftGetSizeMany64", CONV_MATH_FUNC, API_FFT}},
{"cufftEstimate1d", {"hipfftEstimate1d", CONV_MATH_FUNC, API_FFT}},
{"cufftEstimate2d", {"hipfftEstimate2d", CONV_MATH_FUNC, API_FFT}},
{"cufftEstimate3d", {"hipfftEstimate3d", CONV_MATH_FUNC, API_FFT}},
{"cufftEstimateMany", {"hipfftEstimateMany", CONV_MATH_FUNC, API_FFT}},
{"cufftCreate", {"hipfftCreate", CONV_MATH_FUNC, API_FFT}},
{"cufftGetSize1d", {"hipfftGetSize1d", CONV_MATH_FUNC, API_FFT}},
{"cufftGetSize2d", {"hipfftGetSize2d", CONV_MATH_FUNC, API_FFT}},
{"cufftGetSize3d", {"hipfftGetSize3d", CONV_MATH_FUNC, API_FFT}},
{"cufftGetSizeMany", {"hipfftGetSizeMany", CONV_MATH_FUNC, API_FFT}},
{"cufftGetSize", {"hipfftGetSize", CONV_MATH_FUNC, API_FFT}},
{"cufftSetWorkArea", {"hipfftSetWorkArea", CONV_MATH_FUNC, API_FFT}},
{"cufftSetAutoAllocation", {"hipfftSetAutoAllocation", CONV_MATH_FUNC, API_FFT}},
{"cufftExecC2C", {"hipfftExecC2C", CONV_MATH_FUNC, API_FFT}},
{"cufftExecR2C", {"hipfftExecR2C", CONV_MATH_FUNC, API_FFT}},
{"cufftExecC2R", {"hipfftExecC2R", CONV_MATH_FUNC, API_FFT}},
{"cufftExecZ2Z", {"hipfftExecZ2Z", CONV_MATH_FUNC, API_FFT}},
{"cufftExecD2Z", {"hipfftExecD2Z", CONV_MATH_FUNC, API_FFT}},
{"cufftExecZ2D", {"hipfftExecZ2D", CONV_MATH_FUNC, API_FFT}},
{"cufftSetStream", {"hipfftSetStream", CONV_MATH_FUNC, API_FFT}},
{"cufftDestroy", {"hipfftDestroy", CONV_MATH_FUNC, API_FFT}},
{"cufftGetVersion", {"hipfftGetVersion", CONV_MATH_FUNC, API_FFT}},
{"cufftGetProperty", {"hipfftGetProperty", CONV_MATH_FUNC, API_FFT, HIP_UNSUPPORTED}},
{"cufftPlan1d", {"hipfftPlan1d", CONV_LIB_FUNC, API_FFT}},
{"cufftPlan2d", {"hipfftPlan2d", CONV_LIB_FUNC, API_FFT}},
{"cufftPlan3d", {"hipfftPlan3d", CONV_LIB_FUNC, API_FFT}},
{"cufftPlanMany", {"hipfftPlanMany", CONV_LIB_FUNC, API_FFT}},
{"cufftMakePlan1d", {"hipfftMakePlan1d", CONV_LIB_FUNC, API_FFT}},
{"cufftMakePlan2d", {"hipfftMakePlan2d", CONV_LIB_FUNC, API_FFT}},
{"cufftMakePlan3d", {"hipfftMakePlan3d", CONV_LIB_FUNC, API_FFT}},
{"cufftMakePlanMany", {"hipfftMakePlanMany", CONV_LIB_FUNC, API_FFT}},
{"cufftMakePlanMany64", {"hipfftMakePlanMany64", CONV_LIB_FUNC, API_FFT}},
{"cufftGetSizeMany64", {"hipfftGetSizeMany64", CONV_LIB_FUNC, API_FFT}},
{"cufftEstimate1d", {"hipfftEstimate1d", CONV_LIB_FUNC, API_FFT}},
{"cufftEstimate2d", {"hipfftEstimate2d", CONV_LIB_FUNC, API_FFT}},
{"cufftEstimate3d", {"hipfftEstimate3d", CONV_LIB_FUNC, API_FFT}},
{"cufftEstimateMany", {"hipfftEstimateMany", CONV_LIB_FUNC, API_FFT}},
{"cufftCreate", {"hipfftCreate", CONV_LIB_FUNC, API_FFT}},
{"cufftGetSize1d", {"hipfftGetSize1d", CONV_LIB_FUNC, API_FFT}},
{"cufftGetSize2d", {"hipfftGetSize2d", CONV_LIB_FUNC, API_FFT}},
{"cufftGetSize3d", {"hipfftGetSize3d", CONV_LIB_FUNC, API_FFT}},
{"cufftGetSizeMany", {"hipfftGetSizeMany", CONV_LIB_FUNC, API_FFT}},
{"cufftGetSize", {"hipfftGetSize", CONV_LIB_FUNC, API_FFT}},
{"cufftSetWorkArea", {"hipfftSetWorkArea", CONV_LIB_FUNC, API_FFT}},
{"cufftSetAutoAllocation", {"hipfftSetAutoAllocation", CONV_LIB_FUNC, API_FFT}},
{"cufftExecC2C", {"hipfftExecC2C", CONV_LIB_FUNC, API_FFT}},
{"cufftExecR2C", {"hipfftExecR2C", CONV_LIB_FUNC, API_FFT}},
{"cufftExecC2R", {"hipfftExecC2R", CONV_LIB_FUNC, API_FFT}},
{"cufftExecZ2Z", {"hipfftExecZ2Z", CONV_LIB_FUNC, API_FFT}},
{"cufftExecD2Z", {"hipfftExecD2Z", CONV_LIB_FUNC, API_FFT}},
{"cufftExecZ2D", {"hipfftExecZ2D", CONV_LIB_FUNC, API_FFT}},
{"cufftSetStream", {"hipfftSetStream", CONV_LIB_FUNC, API_FFT}},
{"cufftDestroy", {"hipfftDestroy", CONV_LIB_FUNC, API_FFT}},
{"cufftGetVersion", {"hipfftGetVersion", CONV_LIB_FUNC, API_FFT}},
{"cufftGetProperty", {"hipfftGetProperty", CONV_LIB_FUNC, API_FFT, HIP_UNSUPPORTED}},
};
hipamd/hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp
+53 -53
View File
@@ -3,60 +3,60 @@
// Map of all functions
const std::map<llvm::StringRef, hipCounter> CUDA_RAND_FUNCTION_MAP{
// RAND Host functions
{"curandCreateGenerator", {"hiprandCreateGenerator", CONV_MATH_FUNC, API_RAND}},
{"curandCreateGeneratorHost", {"hiprandCreateGeneratorHost", CONV_MATH_FUNC, API_RAND}},
{"curandCreatePoissonDistribution", {"hiprandCreatePoissonDistribution", CONV_MATH_FUNC, API_RAND}},
{"curandDestroyDistribution", {"hiprandDestroyDistribution", CONV_MATH_FUNC, API_RAND}},
{"curandDestroyGenerator", {"hiprandDestroyGenerator", CONV_MATH_FUNC, API_RAND}},
{"curandGenerate", {"hiprandGenerate", CONV_MATH_FUNC, API_RAND}},
{"curandGenerateLogNormal", {"hiprandGenerateLogNormal", CONV_MATH_FUNC, API_RAND}},
{"curandGenerateLogNormalDouble", {"hiprandGenerateLogNormalDouble", CONV_MATH_FUNC, API_RAND}},
{"curandGenerateLongLong", {"hiprandGenerateLongLong", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGenerateNormal", {"hiprandGenerateNormal", CONV_MATH_FUNC, API_RAND}},
{"curandGenerateNormalDouble", {"hiprandGenerateNormalDouble", CONV_MATH_FUNC, API_RAND}},
{"curandGeneratePoisson", {"hiprandGeneratePoisson", CONV_MATH_FUNC, API_RAND}},
{"curandGenerateSeeds", {"hiprandGenerateSeeds", CONV_MATH_FUNC, API_RAND}},
{"curandGenerateUniform", {"hiprandGenerateUniform", CONV_MATH_FUNC, API_RAND}},
{"curandGenerateUniformDouble", {"hiprandGenerateUniformDouble", CONV_MATH_FUNC, API_RAND}},
{"curandGetDirectionVectors32", {"hiprandGetDirectionVectors32", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGetDirectionVectors64", {"hiprandGetDirectionVectors64", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGetProperty", {"hiprandGetProperty", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGetScrambleConstants32", {"hiprandGetScrambleConstants32", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGetScrambleConstants64", {"hiprandGetScrambleConstants64", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGetVersion", {"hiprandGetVersion", CONV_MATH_FUNC, API_RAND}},
{"curandSetGeneratorOffset", {"hiprandSetGeneratorOffset", CONV_MATH_FUNC, API_RAND}},
{"curandSetGeneratorOrdering", {"hiprandSetGeneratorOrdering", CONV_MATH_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandSetPseudoRandomGeneratorSeed", {"hiprandSetPseudoRandomGeneratorSeed", CONV_MATH_FUNC, API_RAND}},
{"curandSetQuasiRandomGeneratorDimensions", {"hiprandSetQuasiRandomGeneratorDimensions", CONV_MATH_FUNC, API_RAND}},
{"curandSetStream", {"hiprandSetStream", CONV_MATH_FUNC, API_RAND}},
{"curandCreateGenerator", {"hiprandCreateGenerator", CONV_LIB_FUNC, API_RAND}},
{"curandCreateGeneratorHost", {"hiprandCreateGeneratorHost", CONV_LIB_FUNC, API_RAND}},
{"curandCreatePoissonDistribution", {"hiprandCreatePoissonDistribution", CONV_LIB_FUNC, API_RAND}},
{"curandDestroyDistribution", {"hiprandDestroyDistribution", CONV_LIB_FUNC, API_RAND}},
{"curandDestroyGenerator", {"hiprandDestroyGenerator", CONV_LIB_FUNC, API_RAND}},
{"curandGenerate", {"hiprandGenerate", CONV_LIB_FUNC, API_RAND}},
{"curandGenerateLogNormal", {"hiprandGenerateLogNormal", CONV_LIB_FUNC, API_RAND}},
{"curandGenerateLogNormalDouble", {"hiprandGenerateLogNormalDouble", CONV_LIB_FUNC, API_RAND}},
{"curandGenerateLongLong", {"hiprandGenerateLongLong", CONV_LIB_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGenerateNormal", {"hiprandGenerateNormal", CONV_LIB_FUNC, API_RAND}},
{"curandGenerateNormalDouble", {"hiprandGenerateNormalDouble", CONV_LIB_FUNC, API_RAND}},
{"curandGeneratePoisson", {"hiprandGeneratePoisson", CONV_LIB_FUNC, API_RAND}},
{"curandGenerateSeeds", {"hiprandGenerateSeeds", CONV_LIB_FUNC, API_RAND}},
{"curandGenerateUniform", {"hiprandGenerateUniform", CONV_LIB_FUNC, API_RAND}},
{"curandGenerateUniformDouble", {"hiprandGenerateUniformDouble", CONV_LIB_FUNC, API_RAND}},
{"curandGetDirectionVectors32", {"hiprandGetDirectionVectors32", CONV_LIB_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGetDirectionVectors64", {"hiprandGetDirectionVectors64", CONV_LIB_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGetProperty", {"hiprandGetProperty", CONV_LIB_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGetScrambleConstants32", {"hiprandGetScrambleConstants32", CONV_LIB_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGetScrambleConstants64", {"hiprandGetScrambleConstants64", CONV_LIB_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandGetVersion", {"hiprandGetVersion", CONV_LIB_FUNC, API_RAND}},
{"curandSetGeneratorOffset", {"hiprandSetGeneratorOffset", CONV_LIB_FUNC, API_RAND}},
{"curandSetGeneratorOrdering", {"hiprandSetGeneratorOrdering", CONV_LIB_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curandSetPseudoRandomGeneratorSeed", {"hiprandSetPseudoRandomGeneratorSeed", CONV_LIB_FUNC, API_RAND}},
{"curandSetQuasiRandomGeneratorDimensions", {"hiprandSetQuasiRandomGeneratorDimensions", CONV_LIB_FUNC, API_RAND}},
{"curandSetStream", {"hiprandSetStream", CONV_LIB_FUNC, API_RAND}},
// RAND Device functions
{"curand", {"hiprand", CONV_DEVICE_FUNC, API_RAND}},
{"curand_init", {"hiprand_init", CONV_DEVICE_FUNC, API_RAND}},
{"curand_log_normal", {"hiprand_log_normal", CONV_DEVICE_FUNC, API_RAND}},
{"curand_log_normal_double", {"hiprand_log_normal_double", CONV_DEVICE_FUNC, API_RAND}},
{"curand_log_normal2", {"hiprand_log_normal2", CONV_DEVICE_FUNC, API_RAND}},
{"curand_log_normal2_double", {"hiprand_log_normal2_double", CONV_DEVICE_FUNC, API_RAND}},
{"curand_log_normal4", {"hiprand_log_normal4", CONV_DEVICE_FUNC, API_RAND}},
{"curand_log_normal4_double", {"hiprand_log_normal4_double", CONV_DEVICE_FUNC, API_RAND}},
{"curand_mtgp32_single", {"hiprand_mtgp32_single", CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curand_mtgp32_single_specific", {"hiprand_mtgp32_single_specific", CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curand_mtgp32_specific", {"hiprand_mtgp32_specific", CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curand_normal", {"hiprand_normal", CONV_DEVICE_FUNC, API_RAND}},
{"curand_normal_double", {"hiprand_normal_double", CONV_DEVICE_FUNC, API_RAND}},
{"curand_normal2", {"hiprand_normal2", CONV_DEVICE_FUNC, API_RAND}},
{"curand_normal2_double", {"hiprand_normal2_double", CONV_DEVICE_FUNC, API_RAND}},
{"curand_normal4", {"hiprand_normal4", CONV_DEVICE_FUNC, API_RAND}},
{"curand_normal4_double", {"hiprand_normal4_double", CONV_DEVICE_FUNC, API_RAND}},
{"curand_uniform", {"hiprand_uniform", CONV_DEVICE_FUNC, API_RAND}},
{"curand_uniform_double", {"hiprand_uniform_double", CONV_DEVICE_FUNC, API_RAND}},
{"curand_uniform2_double", {"hiprand_uniform2_double", CONV_DEVICE_FUNC, API_RAND}},
{"curand_uniform4", {"hiprand_uniform4", CONV_DEVICE_FUNC, API_RAND}},
{"curand_uniform4_double", {"hiprand_uniform4_double", CONV_DEVICE_FUNC, API_RAND}},
{"curand_discrete", {"hiprand_discrete", CONV_DEVICE_FUNC, API_RAND}},
{"curand_discrete4", {"hiprand_discrete4", CONV_DEVICE_FUNC, API_RAND}},
{"curand_poisson", {"hiprand_poisson", CONV_DEVICE_FUNC, API_RAND}},
{"curand_poisson4", {"hiprand_poisson4", CONV_DEVICE_FUNC, API_RAND}},
{"curand_Philox4x32_10", {"hiprand_Philox4x32_10", CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curand", {"hiprand", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_init", {"hiprand_init", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_log_normal", {"hiprand_log_normal", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_log_normal_double", {"hiprand_log_normal_double", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_log_normal2", {"hiprand_log_normal2", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_log_normal2_double", {"hiprand_log_normal2_double", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_log_normal4", {"hiprand_log_normal4", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_log_normal4_double", {"hiprand_log_normal4_double", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_mtgp32_single", {"hiprand_mtgp32_single", CONV_LIB_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curand_mtgp32_single_specific", {"hiprand_mtgp32_single_specific", CONV_LIB_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curand_mtgp32_specific", {"hiprand_mtgp32_specific", CONV_LIB_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}},
{"curand_normal", {"hiprand_normal", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_normal_double", {"hiprand_normal_double", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_normal2", {"hiprand_normal2", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_normal2_double", {"hiprand_normal2_double", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_normal4", {"hiprand_normal4", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_normal4_double", {"hiprand_normal4_double", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_uniform", {"hiprand_uniform", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_uniform_double", {"hiprand_uniform_double", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_uniform2_double", {"hiprand_uniform2_double", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_uniform4", {"hiprand_uniform4", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_uniform4_double", {"hiprand_uniform4_double", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_discrete", {"hiprand_discrete", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_discrete4", {"hiprand_discrete4", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_poisson", {"hiprand_poisson", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_poisson4", {"hiprand_poisson4", CONV_LIB_DEVICE_FUNC, API_RAND}},
{"curand_Philox4x32_10", {"hiprand_Philox4x32_10", CONV_LIB_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}},
// unchanged function names: skipahead, skipahead_sequence, skipahead_subsequence
};
hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
+204 -147
View File
@@ -1,133 +1,185 @@
#include "CUDA2HIP.h"
// Map of all functions
// Map of all CUDA Runtime API functions
const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
/////////////////////////////// CUDA RT API ///////////////////////////////
// Error API
{"cudaGetLastError", {"hipGetLastError", CONV_ERROR, API_RUNTIME}},
{"cudaPeekAtLastError", {"hipPeekAtLastError", CONV_ERROR, API_RUNTIME}},
{"cudaGetErrorName", {"hipGetErrorName", CONV_ERROR, API_RUNTIME}},
{"cudaGetErrorString", {"hipGetErrorString", CONV_ERROR, API_RUNTIME}},
// memcpy
// memcpy structs
{"cudaMemcpy3DParms", {"hipMemcpy3DParms", CONV_MEM, API_RUNTIME}},
{"cudaMemcpy3DPeerParms", {"hipMemcpy3DPeerParms", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
// memcpy functions
{"cudaMemcpy", {"hipMemcpy", CONV_MEM, API_RUNTIME}},
{"cudaMemcpyToArray", {"hipMemcpyToArray", CONV_MEM, API_RUNTIME}},
{"cudaMemcpyToSymbol", {"hipMemcpyToSymbol", CONV_MEM, API_RUNTIME}},
{"cudaMemcpyToSymbolAsync", {"hipMemcpyToSymbolAsync", CONV_MEM, API_RUNTIME}},
{"cudaMemcpyAsync", {"hipMemcpyAsync", CONV_MEM, API_RUNTIME}},
{"cudaMemcpy2D", {"hipMemcpy2D", CONV_MEM, API_RUNTIME}},
{"cudaMemcpy2DAsync", {"hipMemcpy2DAsync", CONV_MEM, API_RUNTIME}},
{"cudaMemcpy2DToArray", {"hipMemcpy2DToArray", CONV_MEM, API_RUNTIME}},
{"cudaMemcpy2DArrayToArray", {"hipMemcpy2DArrayToArray", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpy2DFromArray", {"hipMemcpy2DFromArray", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpy2DFromArrayAsync", {"hipMemcpy2DFromArrayAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpy2DToArrayAsync", {"hipMemcpy2DToArrayAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpy3D", {"hipMemcpy3D", CONV_MEM, API_RUNTIME}},
{"cudaMemcpy3DAsync", {"hipMemcpy3DAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpy3DPeer", {"hipMemcpy3DPeer", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpy3DPeerAsync", {"hipMemcpy3DPeerAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpyArrayToArray", {"hipMemcpyArrayToArray", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpyFromArrayAsync", {"hipMemcpyFromArrayAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpyFromSymbol", {"hipMemcpyFromSymbol", CONV_MEM, API_RUNTIME}},
{"cudaMemcpyFromSymbolAsync", {"hipMemcpyFromSymbolAsync", CONV_MEM, API_RUNTIME}},
{"cudaMemAdvise", {"hipMemAdvise", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}}, //
{"cudaMemRangeGetAttribute", {"hipMemRangeGetAttribute", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}}, //
{"cudaMemRangeGetAttributes", {"hipMemRangeGetAttributes", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}}, //
// no analogue
// NOTE: Not equal to cuMemcpy due to different signatures
{"cudaMemcpy", {"hipMemcpy", CONV_MEMORY, API_RUNTIME}},
{"cudaMemcpyToArray", {"hipMemcpyToArray", CONV_MEMORY, API_RUNTIME}},
{"cudaMemcpyToSymbol", {"hipMemcpyToSymbol", CONV_MEMORY, API_RUNTIME}},
{"cudaMemcpyToSymbolAsync", {"hipMemcpyToSymbolAsync", CONV_MEMORY, API_RUNTIME}},
{"cudaMemcpyAsync", {"hipMemcpyAsync", CONV_MEMORY, API_RUNTIME}},
// no analogue
// NOTE: Not equal to cuMemcpy2D due to different signatures
{"cudaMemcpy2D", {"hipMemcpy2D", CONV_MEMORY, API_RUNTIME}},
// no analogue
// NOTE: Not equal to cuMemcpy2DAsync due to different signatures
{"cudaMemcpy2DAsync", {"hipMemcpy2DAsync", CONV_MEMORY, API_RUNTIME}},
{"cudaMemcpy2DToArray", {"hipMemcpy2DToArray", CONV_MEMORY, API_RUNTIME}},
{"cudaMemcpy2DArrayToArray", {"hipMemcpy2DArrayToArray", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpy2DFromArray", {"hipMemcpy2DFromArray", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpy2DFromArrayAsync", {"hipMemcpy2DFromArrayAsync", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpy2DToArrayAsync", {"hipMemcpy2DToArrayAsync", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// no analogue
// NOTE: Not equal to cuMemcpy3D due to different signatures
{"cudaMemcpy3D", {"hipMemcpy3D", CONV_MEMORY, API_RUNTIME}},
// no analogue
// NOTE: Not equal to cuMemcpy3DAsync due to different signatures
{"cudaMemcpy3DAsync", {"hipMemcpy3DAsync", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// no analogue
// NOTE: Not equal to cuMemcpy3DPeer due to different signatures
{"cudaMemcpy3DPeer", {"hipMemcpy3DPeer", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// no analogue
// NOTE: Not equal to cuMemcpy3DPeerAsync due to different signatures
{"cudaMemcpy3DPeerAsync", {"hipMemcpy3DPeerAsync", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// no analogue
// NOTE: Not equal to cuMemcpyAtoA due to different signatures
{"cudaMemcpyArrayToArray", {"hipMemcpyArrayToArray", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpyFromArrayAsync", {"hipMemcpyFromArrayAsync", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemcpyFromSymbol", {"hipMemcpyFromSymbol", CONV_MEMORY, API_RUNTIME}},
{"cudaMemcpyFromSymbolAsync", {"hipMemcpyFromSymbolAsync", CONV_MEMORY, API_RUNTIME}},
// cuMemAdvise
{"cudaMemAdvise", {"hipMemAdvise", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// cuMemRangeGetAttribute
{"cudaMemRangeGetAttribute", {"hipMemRangeGetAttribute", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// cuMemRangeGetAttributes
{"cudaMemRangeGetAttributes", {"hipMemRangeGetAttributes", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// memset
{"cudaMemset", {"hipMemset", CONV_MEM, API_RUNTIME}},
{"cudaMemsetAsync", {"hipMemsetAsync", CONV_MEM, API_RUNTIME}},
{"cudaMemset2D", {"hipMemset2D", CONV_MEM, API_RUNTIME}},
{"cudaMemset2DAsync", {"hipMemset2DAsync", CONV_MEM, API_RUNTIME}},
{"cudaMemset3D", {"hipMemset3D", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemset3DAsync", {"hipMemset3DAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemset", {"hipMemset", CONV_MEMORY, API_RUNTIME}},
{"cudaMemsetAsync", {"hipMemsetAsync", CONV_MEMORY, API_RUNTIME}},
{"cudaMemset2D", {"hipMemset2D", CONV_MEMORY, API_RUNTIME}},
{"cudaMemset2DAsync", {"hipMemset2DAsync", CONV_MEMORY, API_RUNTIME}},
{"cudaMemset3D", {"hipMemset3D", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemset3DAsync", {"hipMemset3DAsync", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// Memory management
{"cudaMemGetInfo", {"hipMemGetInfo", CONV_MEM, API_RUNTIME}},
{"cudaArrayGetInfo", {"hipArrayGetInfo", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaFreeMipmappedArray", {"hipFreeMipmappedArray", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGetMipmappedArrayLevel", {"hipGetMipmappedArrayLevel", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGetSymbolAddress", {"hipGetSymbolAddress", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGetSymbolSize", {"hipGetSymbolSize", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMemPrefetchAsync", {"hipMemPrefetchAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}}, // // API_Driver ANALOGUE (cuMemPrefetchAsync)
// cuMemGetInfo
{"cudaMemGetInfo", {"hipMemGetInfo", CONV_MEMORY, API_RUNTIME}},
{"cudaArrayGetInfo", {"hipArrayGetInfo", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// no analogue
// NOTE: Not equal to cuMipmappedArrayDestroy due to different signatures
{"cudaFreeMipmappedArray", {"hipFreeMipmappedArray", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGetMipmappedArrayLevel", {"hipGetMipmappedArrayLevel", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGetSymbolAddress", {"hipGetSymbolAddress", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGetSymbolSize", {"hipGetSymbolSize", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// TODO: double check cuMemPrefetchAsync
{"cudaMemPrefetchAsync", {"hipMemPrefetchAsync", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// malloc
{"cudaMalloc", {"hipMalloc", CONV_MEM, API_RUNTIME}},
{"cudaMallocHost", {"hipHostMalloc", CONV_MEM, API_RUNTIME}},
{"cudaMallocArray", {"hipMallocArray", CONV_MEM, API_RUNTIME}},
{"cudaMalloc3D", {"hipMalloc3D", CONV_MEM, API_RUNTIME}},
{"cudaMalloc3DArray", {"hipMalloc3DArray", CONV_MEM, API_RUNTIME}},
{"cudaMallocManaged", {"hipMallocManaged", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMallocMipmappedArray", {"hipMallocMipmappedArray", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMallocPitch", {"hipMallocPitch", CONV_MEM, API_RUNTIME}},
{"cudaMalloc", {"hipMalloc", CONV_MEMORY, API_RUNTIME}},
{"cudaMallocHost", {"hipHostMalloc", CONV_MEMORY, API_RUNTIME}},
{"cudaMallocArray", {"hipMallocArray", CONV_MEMORY, API_RUNTIME}},
{"cudaMalloc3D", {"hipMalloc3D", CONV_MEMORY, API_RUNTIME}},
{"cudaMalloc3DArray", {"hipMalloc3DArray", CONV_MEMORY, API_RUNTIME}},
{"cudaMallocManaged", {"hipMallocManaged", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
// no analogue
// NOTE: Not equal to cuMipmappedArrayCreate due to different signatures
{"cudaMallocMipmappedArray", {"hipMallocMipmappedArray", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaMallocPitch", {"hipMallocPitch", CONV_MEMORY, API_RUNTIME}},
{"cudaFree", {"hipFree", CONV_MEM, API_RUNTIME}},
{"cudaFreeHost", {"hipHostFree", CONV_MEM, API_RUNTIME}},
{"cudaFreeArray", {"hipFreeArray", CONV_MEM, API_RUNTIME}},
{"cudaHostRegister", {"hipHostRegister", CONV_MEM, API_RUNTIME}},
{"cudaHostUnregister", {"hipHostUnregister", CONV_MEM, API_RUNTIME}},
// hipHostAlloc deprecated - use hipHostMalloc instead
{"cudaHostAlloc", {"hipHostMalloc", CONV_MEM, API_RUNTIME}},
// cuMemFree
{"cudaFree", {"hipFree", CONV_MEMORY, API_RUNTIME}},
// cuMemFreeHost
{"cudaFreeHost", {"hipHostFree", CONV_MEMORY, API_RUNTIME}},
{"cudaFreeArray", {"hipFreeArray", CONV_MEMORY, API_RUNTIME}},
// cuMemHostRegister
{"cudaHostRegister", {"hipHostRegister", CONV_MEMORY, API_RUNTIME}},
// cuMemHostUnregister
{"cudaHostUnregister", {"hipHostUnregister", CONV_MEMORY, API_RUNTIME}},
// cuMemHostAlloc
// NOTE: hipHostAlloc deprecated - use hipHostMalloc instead
{"cudaHostAlloc", {"hipHostMalloc", CONV_MEMORY, API_RUNTIME}},
// make memory functions
{"make_cudaExtent", {"make_hipExtent", CONV_MEM, API_RUNTIME}},
{"make_cudaPitchedPtr", {"make_hipPitchedPtr", CONV_MEM, API_RUNTIME}},
{"make_cudaPos", {"make_hipPos", CONV_MEM, API_RUNTIME}},
// Host Malloc Flags (#defines)
{"cudaHostAllocDefault", {"hipHostMallocDefault", CONV_MEM, API_RUNTIME}},
{"cudaHostAllocPortable", {"hipHostMallocPortable", CONV_MEM, API_RUNTIME}},
{"cudaHostAllocMapped", {"hipHostMallocMapped", CONV_MEM, API_RUNTIME}},
{"cudaHostAllocWriteCombined", {"hipHostMallocWriteCombined", CONV_MEM, API_RUNTIME}},
{"make_cudaExtent", {"make_hipExtent", CONV_MEMORY, API_RUNTIME}},
{"make_cudaPitchedPtr", {"make_hipPitchedPtr", CONV_MEMORY, API_RUNTIME}},
{"make_cudaPos", {"make_hipPos", CONV_MEMORY, API_RUNTIME}},
// Host Register Flags
{"cudaHostGetFlags", {"hipHostGetFlags", CONV_MEM, API_RUNTIME}},
{"cudaHostRegisterDefault", {"hipHostRegisterDefault", CONV_MEM, API_RUNTIME}},
{"cudaHostRegisterPortable", {"hipHostRegisterPortable", CONV_MEM, API_RUNTIME}},
{"cudaHostRegisterMapped", {"hipHostRegisterMapped", CONV_MEM, API_RUNTIME}},
{"cudaHostRegisterIoMemory", {"hipHostRegisterIoMemory", CONV_MEM, API_RUNTIME}},
{"warpSize", {"hipWarpSize", CONV_SPECIAL_FUNC, API_RUNTIME}},
// cuMemHostGetFlags
{"cudaHostGetFlags", {"hipHostGetFlags", CONV_MEMORY, API_RUNTIME}},
// Events
{"cudaEventCreate", {"hipEventCreate", CONV_EVENT, API_RUNTIME}},
{"cudaEventCreateWithFlags", {"hipEventCreateWithFlags", CONV_EVENT, API_RUNTIME}},
{"cudaEventDestroy", {"hipEventDestroy", CONV_EVENT, API_RUNTIME}},
{"cudaEventRecord", {"hipEventRecord", CONV_EVENT, API_RUNTIME}},
{"cudaEventElapsedTime", {"hipEventElapsedTime", CONV_EVENT, API_RUNTIME}},
{"cudaEventSynchronize", {"hipEventSynchronize", CONV_EVENT, API_RUNTIME}},
{"cudaEventQuery", {"hipEventQuery", CONV_EVENT, API_RUNTIME}},
// Event Flags
{"cudaEventDefault", {"hipEventDefault", CONV_EVENT, API_RUNTIME}},
{"cudaEventBlockingSync", {"hipEventBlockingSync", CONV_EVENT, API_RUNTIME}},
{"cudaEventDisableTiming", {"hipEventDisableTiming", CONV_EVENT, API_RUNTIME}},
{"cudaEventInterprocess", {"hipEventInterprocess", CONV_EVENT, API_RUNTIME}},
// no analogue
// NOTE: Not equal to cuEventCreate due to different signatures
{"cudaEventCreate", {"hipEventCreate", CONV_EVENT, API_RUNTIME}},
// cuEventCreate
{"cudaEventCreateWithFlags", {"hipEventCreateWithFlags", CONV_EVENT, API_RUNTIME}},
// cuEventDestroy
{"cudaEventDestroy", {"hipEventDestroy", CONV_EVENT, API_RUNTIME}},
// cuEventRecord
{"cudaEventRecord", {"hipEventRecord", CONV_EVENT, API_RUNTIME}},
// cuEventElapsedTime
{"cudaEventElapsedTime", {"hipEventElapsedTime", CONV_EVENT, API_RUNTIME}},
// cuEventSynchronize
{"cudaEventSynchronize", {"hipEventSynchronize", CONV_EVENT, API_RUNTIME}},
// cuEventQuery
{"cudaEventQuery", {"hipEventQuery", CONV_EVENT, API_RUNTIME}},
// 5.6. External Resource Interoperability
// cuDestroyExternalMemory
{"cudaDestroyExternalMemory", {"hipDestroyExternalMemory", CONV_EXT_RES, API_RUNTIME, HIP_UNSUPPORTED}},
// cuDestroyExternalSemaphore
{"cudaDestroyExternalSemaphore", {"hipDestroyExternalSemaphore", CONV_EXT_RES, API_RUNTIME, HIP_UNSUPPORTED}},
// cuExternalMemoryGetMappedBuffer
{"cudaExternalMemoryGetMappedBuffer", {"hipExternalMemoryGetMappedBuffer", CONV_EXT_RES, API_RUNTIME, HIP_UNSUPPORTED}},
// cuExternalMemoryGetMappedMipmappedArray
{"cudaExternalMemoryGetMappedMipmappedArray", {"hipExternalMemoryGetMappedMipmappedArray", CONV_EXT_RES, API_RUNTIME, HIP_UNSUPPORTED}},
// cuImportExternalMemory
{"cudaImportExternalMemory", {"hipImportExternalMemory", CONV_EXT_RES, API_RUNTIME, HIP_UNSUPPORTED}},
// cuImportExternalSemaphore
{"cudaImportExternalSemaphore", {"hipImportExternalSemaphore", CONV_EXT_RES, API_RUNTIME, HIP_UNSUPPORTED}},
// cuSignalExternalSemaphoresAsync
{"cudaSignalExternalSemaphoresAsync", {"hipSignalExternalSemaphoresAsync", CONV_EXT_RES, API_RUNTIME, HIP_UNSUPPORTED}},
// cuWaitExternalSemaphoresAsync
{"cudaWaitExternalSemaphoresAsync", {"hipWaitExternalSemaphoresAsync", CONV_EXT_RES, API_RUNTIME, HIP_UNSUPPORTED}},
// Streams
// no analogue
// NOTE: Not equal to cuStreamCreate due to different signatures
{"cudaStreamCreate", {"hipStreamCreate", CONV_STREAM, API_RUNTIME}},
// cuStreamCreate
{"cudaStreamCreateWithFlags", {"hipStreamCreateWithFlags", CONV_STREAM, API_RUNTIME}},
{"cudaStreamCreateWithPriority", {"hipStreamCreateWithPriority", CONV_STREAM, API_RUNTIME, HIP_UNSUPPORTED}},
// cuStreamCreateWithPriority
{"cudaStreamCreateWithPriority", {"hipStreamCreateWithPriority", CONV_STREAM, API_RUNTIME}},
// cuStreamDestroy
{"cudaStreamDestroy", {"hipStreamDestroy", CONV_STREAM, API_RUNTIME}},
// cuStreamWaitEvent
{"cudaStreamWaitEvent", {"hipStreamWaitEvent", CONV_STREAM, API_RUNTIME}},
// cuStreamSynchronize
{"cudaStreamSynchronize", {"hipStreamSynchronize", CONV_STREAM, API_RUNTIME}},
// cuStreamGetFlags
{"cudaStreamGetFlags", {"hipStreamGetFlags", CONV_STREAM, API_RUNTIME}},
// cuStreamQuery
{"cudaStreamQuery", {"hipStreamQuery", CONV_STREAM, API_RUNTIME}},
// cuStreamAddCallback
{"cudaStreamAddCallback", {"hipStreamAddCallback", CONV_STREAM, API_RUNTIME}},
// cuStreamAttachMemAsync
{"cudaStreamAttachMemAsync", {"hipStreamAttachMemAsync", CONV_STREAM, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaStreamGetPriority", {"hipStreamGetPriority", CONV_STREAM, API_RUNTIME, HIP_UNSUPPORTED}},
// cuStreamBeginCapture
{"cudaStreamBeginCapture", {"hipStreamBeginCapture", CONV_STREAM, API_RUNTIME, HIP_UNSUPPORTED}},
// cuStreamEndCapture
{"cudaStreamEndCapture", {"hipStreamEndCapture", CONV_STREAM, API_RUNTIME, HIP_UNSUPPORTED}},
// cuStreamIsCapturing
{"cudaStreamIsCapturing", {"hipStreamIsCapturing", CONV_STREAM, API_RUNTIME, HIP_UNSUPPORTED}},
// cuStreamGetPriority
{"cudaStreamGetPriority", {"hipStreamGetPriority", CONV_STREAM, API_RUNTIME}},
// Other synchronization
{"cudaDeviceSynchronize", {"hipDeviceSynchronize", CONV_DEVICE, API_RUNTIME}},
{"cudaDeviceReset", {"hipDeviceReset", CONV_DEVICE, API_RUNTIME}},
{"cudaSetDevice", {"hipSetDevice", CONV_DEVICE, API_RUNTIME}},
{"cudaGetDevice", {"hipGetDevice", CONV_DEVICE, API_RUNTIME}},
// cuDeviceGetCount
{"cudaGetDeviceCount", {"hipGetDeviceCount", CONV_DEVICE, API_RUNTIME}},
{"cudaChooseDevice", {"hipChooseDevice", CONV_DEVICE, API_RUNTIME}},
@@ -143,39 +195,44 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
{"cudaDeviceGetAttribute", {"hipDeviceGetAttribute", CONV_DEVICE, API_RUNTIME}},
// Pointer Attributes
// struct cudaPointerAttributes
{"cudaPointerGetAttributes", {"hipPointerGetAttributes", CONV_MEM, API_RUNTIME}},
{"cudaHostGetDevicePointer", {"hipHostGetDevicePointer", CONV_MEM, API_RUNTIME}},
// no analogue
// NOTE: Not equal to cuPointerGetAttributes due to different signatures
{"cudaPointerGetAttributes", {"hipPointerGetAttributes", CONV_ADDRESSING, API_RUNTIME}},
// cuMemHostGetDevicePointer
{"cudaHostGetDevicePointer", {"hipHostGetDevicePointer", CONV_MEMORY, API_RUNTIME}},
// Device
{"cudaGetDeviceProperties", {"hipGetDeviceProperties", CONV_DEVICE, API_RUNTIME}},
// cuDeviceGetPCIBusId
{"cudaDeviceGetPCIBusId", {"hipDeviceGetPCIBusId", CONV_DEVICE, API_RUNTIME}},
// cuDeviceGetByPCIBusId
{"cudaDeviceGetByPCIBusId", {"hipDeviceGetByPCIBusId", CONV_DEVICE, API_RUNTIME}},
{"cudaDeviceGetStreamPriorityRange", {"hipDeviceGetStreamPriorityRange", CONV_DEVICE, API_RUNTIME, HIP_UNSUPPORTED}},
// cuCtxGetStreamPriorityRange
{"cudaDeviceGetStreamPriorityRange", {"hipDeviceGetStreamPriorityRange", CONV_DEVICE, API_RUNTIME}},
{"cudaSetValidDevices", {"hipSetValidDevices", CONV_DEVICE, API_RUNTIME, HIP_UNSUPPORTED}},
// Device Flags
{"cudaGetDeviceFlags", {"hipGetDeviceFlags", CONV_DEVICE, API_RUNTIME, HIP_UNSUPPORTED}},
// cuCtxGetFlags
{"cudaGetDeviceFlags", {"hipCtxGetFlags", CONV_DEVICE, API_RUNTIME}},
{"cudaSetDeviceFlags", {"hipSetDeviceFlags", CONV_DEVICE, API_RUNTIME}},
// Cache config
{"cudaDeviceSetCacheConfig", {"hipDeviceSetCacheConfig", CONV_CACHE, API_RUNTIME}},
{"cudaDeviceGetCacheConfig", {"hipDeviceGetCacheConfig", CONV_CACHE, API_RUNTIME}},
{"cudaFuncSetCacheConfig", {"hipFuncSetCacheConfig", CONV_CACHE, API_RUNTIME}},
{"cudaDeviceSetCacheConfig", {"hipDeviceSetCacheConfig", CONV_DEVICE, API_RUNTIME}},
{"cudaDeviceGetCacheConfig", {"hipDeviceGetCacheConfig", CONV_DEVICE, API_RUNTIME}},
{"cudaFuncSetCacheConfig", {"hipFuncSetCacheConfig", CONV_DEVICE, API_RUNTIME}},
// Execution control functions
{"cudaFuncGetAttributes", {"hipFuncGetAttributes", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaFuncSetSharedMemConfig", {"hipFuncSetSharedMemConfig", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGetParameterBuffer", {"hipGetParameterBuffer", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaSetDoubleForDevice", {"hipSetDoubleForDevice", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaSetDoubleForHost", {"hipSetDoubleForHost", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaFuncGetAttributes", {"hipFuncGetAttributes", CONV_EXECUTION, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaFuncSetSharedMemConfig", {"hipFuncSetSharedMemConfig", CONV_EXECUTION, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGetParameterBuffer", {"hipGetParameterBuffer", CONV_EXECUTION, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaSetDoubleForDevice", {"hipSetDoubleForDevice", CONV_EXECUTION, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaSetDoubleForHost", {"hipSetDoubleForHost", CONV_EXECUTION, API_RUNTIME, HIP_UNSUPPORTED}},
// Execution Control [deprecated since 7.0]
{"cudaConfigureCall", {"hipConfigureCall", CONV_EXEC, API_RUNTIME}},
{"cudaLaunch", {"hipLaunchByPtr", CONV_EXEC, API_RUNTIME}},
{"cudaSetupArgument", {"hipSetupArgument", CONV_EXEC, API_RUNTIME}},
{"cudaConfigureCall", {"hipConfigureCall", CONV_EXECUTION, API_RUNTIME}},
{"cudaLaunch", {"hipLaunchByPtr", CONV_EXECUTION, API_RUNTIME}},
{"cudaSetupArgument", {"hipSetupArgument", CONV_EXECUTION, API_RUNTIME}},
// Version Management
{"cudaDriverGetVersion", {"hipDriverGetVersion", CONV_VERSION, API_RUNTIME}},
@@ -194,8 +251,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
{"cudaDeviceDisablePeerAccess", {"hipDeviceDisablePeerAccess", CONV_PEER, API_RUNTIME}},
{"cudaDeviceEnablePeerAccess", {"hipDeviceEnablePeerAccess", CONV_PEER, API_RUNTIME}},
{"cudaMemcpyPeerAsync", {"hipMemcpyPeerAsync", CONV_MEM, API_RUNTIME}},
{"cudaMemcpyPeer", {"hipMemcpyPeer", CONV_MEM, API_RUNTIME}},
{"cudaMemcpyPeerAsync", {"hipMemcpyPeerAsync", CONV_MEMORY, API_RUNTIME}},
{"cudaMemcpyPeer", {"hipMemcpyPeer", CONV_MEMORY, API_RUNTIME}},
// Shared memory
{"cudaDeviceSetSharedMemConfig", {"hipDeviceSetSharedMemConfig", CONV_DEVICE, API_RUNTIME}},
@@ -204,39 +261,39 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
// {"cudaThreadGetSharedMemConfig", {"hipDeviceGetSharedMemConfig", CONV_DEVICE, API_RUNTIME}},
// {"cudaThreadSetSharedMemConfig", {"hipDeviceSetSharedMemConfig", CONV_DEVICE, API_RUNTIME}},
// cuCtxGetLimit
{"cudaDeviceGetLimit", {"hipDeviceGetLimit", CONV_DEVICE, API_RUNTIME}},
// Profiler
{"cudaProfilerInitialize", {"hipProfilerInitialize", CONV_OTHER, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuProfilerInitialize)
{"cudaProfilerStart", {"hipProfilerStart", CONV_OTHER, API_RUNTIME}}, // API_Driver ANALOGUE (cuProfilerStart)
{"cudaProfilerStop", {"hipProfilerStop", CONV_OTHER, API_RUNTIME}}, // API_Driver ANALOGUE (cuProfilerStop)
{"cudaProfilerInitialize", {"hipProfilerInitialize", CONV_PROFILER, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuProfilerInitialize)
{"cudaProfilerStart", {"hipProfilerStart", CONV_PROFILER, API_RUNTIME}}, // API_Driver ANALOGUE (cuProfilerStart)
{"cudaProfilerStop", {"hipProfilerStop", CONV_PROFILER, API_RUNTIME}}, // API_Driver ANALOGUE (cuProfilerStop)
{"cudaBindTexture", {"hipBindTexture", CONV_TEX, API_RUNTIME}},
{"cudaUnbindTexture", {"hipUnbindTexture", CONV_TEX, API_RUNTIME}},
{"cudaBindTexture2D", {"hipBindTexture2D", CONV_TEX, API_RUNTIME}},
{"cudaBindTextureToArray", {"hipBindTextureToArray", CONV_TEX, API_RUNTIME}},
{"cudaBindTextureToMipmappedArray", {"hipBindTextureToMipmappedArray", CONV_TEX, API_RUNTIME}}, // Unsupported yet on NVCC path
{"cudaGetTextureAlignmentOffset", {"hipGetTextureAlignmentOffset", CONV_TEX, API_RUNTIME}}, // Unsupported yet on NVCC path
{"cudaGetTextureReference", {"hipGetTextureReference", CONV_TEX, API_RUNTIME}}, // Unsupported yet on NVCC path
{"cudaBindTexture", {"hipBindTexture", CONV_TEXTURE, API_RUNTIME}},
{"cudaUnbindTexture", {"hipUnbindTexture", CONV_TEXTURE, API_RUNTIME}},
{"cudaBindTexture2D", {"hipBindTexture2D", CONV_TEXTURE, API_RUNTIME}},
{"cudaBindTextureToArray", {"hipBindTextureToArray", CONV_TEXTURE, API_RUNTIME}},
{"cudaBindTextureToMipmappedArray", {"hipBindTextureToMipmappedArray", CONV_TEXTURE, API_RUNTIME}}, // Unsupported yet on NVCC path
{"cudaGetTextureAlignmentOffset", {"hipGetTextureAlignmentOffset", CONV_TEXTURE, API_RUNTIME}}, // Unsupported yet on NVCC path
{"cudaGetTextureReference", {"hipGetTextureReference", CONV_TEXTURE, API_RUNTIME}}, // Unsupported yet on NVCC path
{"cudaCreateChannelDesc", {"hipCreateChannelDesc", CONV_TEX, API_RUNTIME}},
{"cudaGetChannelDesc", {"hipGetChannelDesc", CONV_TEX, API_RUNTIME}},
{"cudaCreateChannelDesc", {"hipCreateChannelDesc", CONV_TEXTURE, API_RUNTIME}},
{"cudaGetChannelDesc", {"hipGetChannelDesc", CONV_TEXTURE, API_RUNTIME}},
// Texture Object Management
{"cudaAddressModeWrap", {"hipAddressModeWrap", CONV_TEX, API_RUNTIME}},
{"cudaAddressModeClamp", {"hipAddressModeClamp", CONV_TEX, API_RUNTIME}},
{"cudaAddressModeMirror", {"hipAddressModeMirror", CONV_TEX, API_RUNTIME}},
{"cudaAddressModeBorder", {"hipAddressModeBorder", CONV_TEX, API_RUNTIME}},
{"cudaAddressModeWrap", {"hipAddressModeWrap", CONV_TEXTURE, API_RUNTIME}},
{"cudaAddressModeClamp", {"hipAddressModeClamp", CONV_TEXTURE, API_RUNTIME}},
{"cudaAddressModeMirror", {"hipAddressModeMirror", CONV_TEXTURE, API_RUNTIME}},
{"cudaAddressModeBorder", {"hipAddressModeBorder", CONV_TEXTURE, API_RUNTIME}},
// functions
{"cudaCreateTextureObject", {"hipCreateTextureObject", CONV_TEX, API_RUNTIME}},
{"cudaDestroyTextureObject", {"hipDestroyTextureObject", CONV_TEX, API_RUNTIME}},
{"cudaGetTextureObjectResourceDesc", {"hipGetTextureObjectResourceDesc", CONV_TEX, API_RUNTIME}},
{"cudaGetTextureObjectResourceViewDesc", {"hipGetTextureObjectResourceViewDesc", CONV_TEX, API_RUNTIME}},
{"cudaGetTextureObjectTextureDesc", {"hipGetTextureObjectTextureDesc", CONV_TEX, API_RUNTIME}},
{"cudaCreateTextureObject", {"hipCreateTextureObject", CONV_TEXTURE, API_RUNTIME}},
{"cudaDestroyTextureObject", {"hipDestroyTextureObject", CONV_TEXTURE, API_RUNTIME}},
{"cudaGetTextureObjectResourceDesc", {"hipGetTextureObjectResourceDesc", CONV_TEXTURE, API_RUNTIME}},
{"cudaGetTextureObjectResourceViewDesc", {"hipGetTextureObjectResourceViewDesc", CONV_TEXTURE, API_RUNTIME}},
{"cudaGetTextureObjectTextureDesc", {"hipGetTextureObjectTextureDesc", CONV_TEXTURE, API_RUNTIME}},
// Surface Reference Management
{"cudaBindSurfaceToArray", {"hipBindSurfaceToArray", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}},
@@ -256,10 +313,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
{"cudaIpcOpenMemHandle", {"hipIpcOpenMemHandle", CONV_DEVICE, API_RUNTIME}},
// OpenGL Interoperability
{"cudaGLGetDevices", {"hipGLGetDevices", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGraphicsGLRegisterBuffer", {"hipGraphicsGLRegisterBuffer", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGraphicsGLRegisterImage", {"hipGraphicsGLRegisterImage", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaWGLGetDevice", {"hipWGLGetDevice", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGLGetDevices", {"hipGLGetDevices", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGraphicsGLRegisterBuffer", {"hipGraphicsGLRegisterBuffer", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaGraphicsGLRegisterImage", {"hipGraphicsGLRegisterImage", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}},
{"cudaWGLGetDevice", {"hipWGLGetDevice", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}},
// Graphics Interoperability
{"cudaGraphicsMapResources", {"hipGraphicsMapResources", CONV_GRAPHICS, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGraphicsMapResources)
@@ -270,21 +327,21 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
{"cudaGraphicsUnmapResources", {"hipGraphicsUnmapResources", CONV_GRAPHICS, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGraphicsUnmapResources)
{"cudaGraphicsUnregisterResource", {"hipGraphicsUnregisterResource", CONV_GRAPHICS, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGraphicsUnregisterResource)
{"cudaGLGetDevices", {"hipGLGetDevices", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLGetDevices)
{"cudaGraphicsGLRegisterBuffer", {"hipGraphicsGLRegisterBuffer", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGraphicsGLRegisterBuffer)
{"cudaGraphicsGLRegisterImage", {"hipGraphicsGLRegisterImage", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGraphicsGLRegisterImage)
{"cudaWGLGetDevice", {"hipWGLGetDevice", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuWGLGetDevice)
{"cudaGLGetDevices", {"hipGLGetDevices", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLGetDevices)
{"cudaGraphicsGLRegisterBuffer", {"hipGraphicsGLRegisterBuffer", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGraphicsGLRegisterBuffer)
{"cudaGraphicsGLRegisterImage", {"hipGraphicsGLRegisterImage", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGraphicsGLRegisterImage)
{"cudaWGLGetDevice", {"hipWGLGetDevice", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuWGLGetDevice)
// OpenGL Interoperability [DEPRECATED]
{"cudaGLMapBufferObject", {"hipGLMapBufferObject__", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // Not equal to cuGLMapBufferObject due to different signatures
{"cudaGLMapBufferObjectAsync", {"hipGLMapBufferObjectAsync__", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // Not equal to cuGLMapBufferObjectAsync due to different signatures
{"cudaGLRegisterBufferObject", {"hipGLRegisterBufferObject", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLRegisterBufferObject)
{"cudaGLSetBufferObjectMapFlags", {"hipGLSetBufferObjectMapFlags", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLSetBufferObjectMapFlags)
{"cudaGLSetGLDevice", {"hipGLSetGLDevice", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // no API_Driver ANALOGUE
{"cudaGLUnmapBufferObject", {"hipGLUnmapBufferObject", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLUnmapBufferObject)
{"cudaGLUnmapBufferObjectAsync", {"hipGLUnmapBufferObjectAsync", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLUnmapBufferObjectAsync)
{"cudaGLUnregisterBufferObject", {"hipGLUnregisterBufferObject", CONV_GL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLUnregisterBufferObject)
{"cudaGLMapBufferObject", {"hipGLMapBufferObject__", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // Not equal to cuGLMapBufferObject due to different signatures
{"cudaGLMapBufferObjectAsync", {"hipGLMapBufferObjectAsync__", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // Not equal to cuGLMapBufferObjectAsync due to different signatures
{"cudaGLRegisterBufferObject", {"hipGLRegisterBufferObject", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLRegisterBufferObject)
{"cudaGLSetBufferObjectMapFlags", {"hipGLSetBufferObjectMapFlags", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLSetBufferObjectMapFlags)
{"cudaGLSetGLDevice", {"hipGLSetGLDevice", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // no API_Driver ANALOGUE
{"cudaGLUnmapBufferObject", {"hipGLUnmapBufferObject", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLUnmapBufferObject)
{"cudaGLUnmapBufferObjectAsync", {"hipGLUnmapBufferObjectAsync", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLUnmapBufferObjectAsync)
{"cudaGLUnregisterBufferObject", {"hipGLUnregisterBufferObject", CONV_OPENGL, API_RUNTIME, HIP_UNSUPPORTED}}, // API_Driver ANALOGUE (cuGLUnregisterBufferObject)
// Direct3D 9 Interoperability
hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp
+1024 -402
View File
File diff suppressed because it is too large Load Diff
hipamd/hipify-clang/src/HipifyAction.cpp
+10 -10
View File
@@ -24,7 +24,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
Statistics::current().incrementCounter(counter, name.str());
if (!counter.unsupported) {
clang::SourceLocation sl = start.getLocWithOffset(begin + 1);
ct::Replacement Rep(SM, sl, name.size(), repName);
ct::Replacement Rep(SM, sl, name.size(), repName.str());
clang::FullSourceLoc fullSL(sl, SM);
insertReplacement(Rep, fullSL);
}
@@ -70,7 +70,7 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
return;
}
StringRef repName = found->second.hipName;
ct::Replacement Rep(SM, sl, name.size(), repName);
ct::Replacement Rep(SM, sl, name.size(), repName.str());
clang::FullSourceLoc fullSL(sl, SM);
insertReplacement(Rep, fullSL);
}
@@ -219,7 +219,7 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
}
const char *B = SM.getCharacterData(sl);
const char *E = SM.getCharacterData(filename_range.getEnd());
ct::Replacement Rep(SM, sl, E - B, newInclude);
ct::Replacement Rep(SM, sl, E - B, newInclude.str());
insertReplacement(Rep, clang::FullSourceLoc{sl, SM});
}
@@ -270,21 +270,21 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
if (numArgs > 0) {
OS << ", ";
// Start of the first argument.
clang::SourceLocation argStart = launchKernel->getArg(0)->getLocStart();
clang::SourceLocation argStart = llcompat::getBeginLoc(launchKernel->getArg(0));
// End of the last argument.
clang::SourceLocation argEnd = launchKernel->getArg(numArgs - 1)->getLocEnd();
clang::SourceLocation argEnd = llcompat::getEndLoc(launchKernel->getArg(numArgs - 1));
OS << readSourceText(*SM, {argStart, argEnd});
}
OS << ")";
clang::SourceRange replacementRange = getWriteRange(*SM, {launchKernel->getLocStart(), launchKernel->getLocEnd()});
clang::SourceRange replacementRange = getWriteRange(*SM, {llcompat::getBeginLoc(launchKernel), llcompat::getEndLoc(launchKernel)});
clang::SourceLocation launchStart = replacementRange.getBegin();
clang::SourceLocation launchEnd = replacementRange.getEnd();
size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart);
ct::Replacement Rep(*SM, launchStart, length, OS.str());
clang::FullSourceLoc fullSL(launchStart, *SM);
insertReplacement(Rep, fullSL);
hipCounter counter = {"hipLaunchKernelGGL", ConvTypes::CONV_KERN, ApiTypes::API_RUNTIME};
hipCounter counter = {"hipLaunchKernelGGL", ConvTypes::CONV_EXECUTION, ApiTypes::API_RUNTIME};
Statistics::current().incrementCounter(counter, refName.str());
return true;
}
@@ -320,8 +320,8 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
}
if (!typeName.empty()) {
clang::SourceLocation slStart = sharedVar->getLocStart();
clang::SourceLocation slEnd = sharedVar->getLocEnd();
clang::SourceLocation slStart = llcompat::getBeginLoc(sharedVar->getTypeSourceInfo()->getTypeLoc());
clang::SourceLocation slEnd = llcompat::getEndLoc(sharedVar->getTypeSourceInfo()->getTypeLoc());
clang::SourceManager* SM = Result.SourceManager;
size_t repLength = SM->getCharacterData(slEnd) - SM->getCharacterData(slStart) + 1;
std::string varName = sharedVar->getNameAsString();
@@ -329,7 +329,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
ct::Replacement Rep(*SM, slStart, repLength, repName);
clang::FullSourceLoc fullSL(slStart, *SM);
insertReplacement(Rep, fullSL);
hipCounter counter = {"HIP_DYNAMIC_SHARED", ConvTypes::CONV_MEM, ApiTypes::API_RUNTIME};
hipCounter counter = {"HIP_DYNAMIC_SHARED", ConvTypes::CONV_MEMORY, ApiTypes::API_RUNTIME};
Statistics::current().incrementCounter(counter, refName.str());
}
return true;
hipamd/hipify-clang/src/HipifyAction.h
+1
View File
@@ -9,6 +9,7 @@
#include "Statistics.h"
namespace ct = clang::tooling;
using namespace llvm;
/**
* A FrontendAction that hipifies CUDA programs.
hipamd/hipify-clang/src/LLVMCompat.cpp
+34 -2
View File
@@ -8,11 +8,11 @@ void PrintStackTraceOnErrorSignal() {
#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8)
llvm::sys::PrintStackTraceOnErrorSignal();
#else
llvm::sys::PrintStackTraceOnErrorSignal(clang::StringRef());
llvm::sys::PrintStackTraceOnErrorSignal(StringRef());
#endif
}
ct::Replacements& getReplacements(ct::RefactoringTool& Tool, clang::StringRef file) {
ct::Replacements& getReplacements(ct::RefactoringTool& Tool, StringRef file) {
#if LLVM_VERSION_MAJOR > 3
// getReplacements() now returns a map from filename to Replacements - so create an entry
// for this source file and return a reference to it.
@@ -40,4 +40,36 @@ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *
#endif
}
clang::SourceLocation getBeginLoc(const clang::Stmt* stmt) {
#if LLVM_VERSION_MAJOR < 8
return stmt->getLocStart();
#else
return stmt->getBeginLoc();
#endif
}
clang::SourceLocation getBeginLoc(const clang::TypeLoc& typeLoc) {
#if LLVM_VERSION_MAJOR < 8
return typeLoc.getLocStart();
#else
return typeLoc.getBeginLoc();
#endif
}
clang::SourceLocation getEndLoc(const clang::Stmt* stmt) {
#if LLVM_VERSION_MAJOR < 8
return stmt->getLocEnd();
#else
return stmt->getEndLoc();
#endif
}
clang::SourceLocation getEndLoc(const clang::TypeLoc& typeLoc) {
#if LLVM_VERSION_MAJOR < 8
return typeLoc.getLocEnd();
#else
return typeLoc.getEndLoc();
#endif
}
} // namespace llcompat
hipamd/hipify-clang/src/LLVMCompat.h
+9 -1
View File
@@ -25,15 +25,23 @@ namespace llcompat {
#define LLVM_DEBUG(X) DEBUG(X)
#endif
clang::SourceLocation getBeginLoc(const clang::Stmt* stmt);
clang::SourceLocation getBeginLoc(const clang::TypeLoc& typeLoc);
clang::SourceLocation getEndLoc(const clang::Stmt* stmt);
clang::SourceLocation getEndLoc(const clang::TypeLoc& typeLoc);
void PrintStackTraceOnErrorSignal();
using namespace llvm;
/**
* Get the replacement map for a given filename in a RefactoringTool.
*
* Older LLVM versions don't actually support multiple filenames, so everything all gets
* smushed together. It is the caller's responsibility to cope with this.
*/
ct::Replacements& getReplacements(ct::RefactoringTool& Tool, clang::StringRef file);
ct::Replacements& getReplacements(ct::RefactoringTool& Tool, StringRef file);
/**
* Add a Replacement to a Replacements.
hipamd/hipify-clang/src/Statistics.cpp
+41 -11
View File
@@ -4,12 +4,42 @@
#include <iomanip>
const char *counterNames[NUM_CONV_TYPES] = {
"version", "init", "device", "mem", "kern", "coord_func", "math_func", "device_func",
"special_func", "stream", "event", "occupancy", "ctx", "peer", "module",
"cache", "exec", "external_resource_interop", "graph", "err", "def", "tex", "gl", "graphics",
"surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl", "complex",
"thread", "other", "include", "include_cuda_main_header", "type", "literal",
"numeric_literal"
"error", // CONV_ERROR
"init", // CONV_INIT
"version", // CONV_VERSION
"device", // CONV_DEVICE
"context", // CONV_CONTEXT
"module", // CONV_MODULE
"memory", // CONV_MEMORY
"addressing", // CONV_ADDRESSING
"stream", // CONV_STREAM
"event", // CONV_EVENT
"external_resource_interop", // CONV_EXT_RES
"stream_memory", // CONV_STREAM_MEMORY
"execution", // CONV_EXECUTION
"graph", // CONV_GRAPH
"occupancy", // CONV_OCCUPANCY
"texture", // CONV_TEXTURE
"surface", // CONV_SURFACE
"peer", // CONV_PEER
"graphics", // CONV_GRAPHICS
"profiler", // CONV_PROFILER
"openGL", // CONV_OPENGL
"D3D9", // CONV_D3D9
"D3D10", // CONV_D3D10
"D3D11", // CONV_D3D11
"VDPAU", // CONV_VDPAU
"EGL", // CONV_EGL
"thread", // CONV_THREAD
"complex", // CONV_COMPLEX
"library", // CONV_LIB_FUNC
"device_library", // CONV_LIB_DEVICE_FUNC
"include", // CONV_INCLUDE
"include_cuda_main_header", // CONV_INCLUDE_CUDA_MAIN_H
"type", // CONV_TYPE
"literal", // CONV_LITERAL
"numeric_literal", // CONV_NUMERIC_LITERAL
"define" // CONV_DEFINE
};
const char *apiNames[NUM_API_TYPES] = {
@@ -44,7 +74,7 @@ void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string
} // Anonymous namespace
void StatCounter::incrementCounter(const hipCounter& counter, std::string name) {
void StatCounter::incrementCounter(const hipCounter& counter, const std::string& name) {
counters[name]++;
apiCounters[(int) counter.apiType]++;
convTypeCounters[(int) counter.type]++;
@@ -70,7 +100,7 @@ int StatCounter::getConvSum() {
return acc;
}
void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix) {
void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& prefix) {
conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n", "[HIPIFY] info: " + prefix + " refs by type:\n");
for (int i = 0; i < NUM_CONV_TYPES; i++) {
if (convTypeCounters[i] > 0) {
@@ -87,7 +117,7 @@ void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, std::str
}
}
Statistics::Statistics(std::string name): fileName(name) {
Statistics::Statistics(const std::string& name): fileName(name) {
// Compute the total bytes/lines in the input file.
std::ifstream src_file(name, std::ios::binary | std::ios::ate);
src_file.clear();
@@ -99,7 +129,7 @@ Statistics::Statistics(std::string name): fileName(name) {
///////// Counter update routines //////////
void Statistics::incrementCounter(const hipCounter &counter, std::string name) {
void Statistics::incrementCounter(const hipCounter &counter, const std::string& name) {
if (counter.unsupported) {
unsupported.incrementCounter(counter, name);
} else {
@@ -188,7 +218,7 @@ Statistics& Statistics::current() {
return *Statistics::currentStatistics;
}
void Statistics::setActive(std::string name) {
void Statistics::setActive(const std::string& name) {
stats.emplace(std::make_pair(name, Statistics{name}));
Statistics::currentStatistics = &stats.at(name);
}
hipamd/hipify-clang/src/Statistics.h
+73 -27
View File
@@ -11,45 +11,91 @@
namespace chr = std::chrono;
enum ConvTypes {
CONV_VERSION = 0,
// Driver API: 5.2. Error Handling
// Runtime API: 5.3. Error Handling
CONV_ERROR = 0,
// Driver API : 5.3. Initialization
CONV_INIT,
// Driver API : 5.4. Version Management
// Runtime API: 5.28. Version Management
CONV_VERSION,
// Driver API : 5.5. Device Management, 5.6. Device Management [DEPRECATED]
// Runtime API: 5.1. Device Management
CONV_DEVICE,
CONV_MEM,
CONV_KERN,
CONV_COORD_FUNC,
CONV_MATH_FUNC,
CONV_DEVICE_FUNC,
CONV_SPECIAL_FUNC,
CONV_STREAM,
CONV_EVENT,
CONV_OCCUPANCY,
// Driver API : 5.7. Primary Context Management, 5.8.Context Management, 5.9. Context Management [DEPRECATED]
CONV_CONTEXT,
CONV_PEER,
// Driver API : 5.10. Module Management
CONV_MODULE,
CONV_CACHE,
CONV_EXEC,
CONV_EXTERNAL_RES,
// Driver API : 5.11. Memory Management
// Runtime API: 5.10. Memory Management
CONV_MEMORY,
// Driver API : 5.12. Unified Addressing
// Runtime API: 5.11. Unified Addressing
CONV_ADDRESSING,
// Driver API : 5.13. Stream Management
// Runtime API: 5.4. Stream Management
CONV_STREAM,
// Driver API : 5.14. Event Management
// Runtime API: 5.5. Event Management
CONV_EVENT,
// Driver API : 5.15. External Resource Interoperability
// Runtime API: 5.6.External Resource Interoperability
CONV_EXT_RES,
// Driver API : 5.16. Stream memory operations
CONV_STREAM_MEMORY,
// Driver API : 5.17. Execution Control, 5.18. Execution Control [DEPRECATED]
// Runtime API: 5.7.Execution Control, 5.9. Execution Control [DEPRECATED]
CONV_EXECUTION,
// Driver API : 5.19. Graph Management
// Runtime API: 5.29. Graph Management
CONV_GRAPH,
CONV_ERROR,
CONV_DEF,
CONV_TEX,
CONV_GL,
CONV_GRAPHICS,
// Driver API : 5.20. Occupancy
// Runtime API: 5.8. Occupancy
CONV_OCCUPANCY,
// Driver API : 5.21. Texture Reference Management, 5.22. Texture Reference Management [DEPRECATED], 5.24. Texture Object Management
// Runtime API: 5.24. Texture Reference Management, 5.26. Texture Object Management
CONV_TEXTURE,
// Driver API : 5.23. Surface Reference Management, 5.25. Surface Object Management
// Runtime API: 5.25. Surface Reference Management, 5.27. Surface Object Management
CONV_SURFACE,
CONV_JIT,
// Driver API : 5.26. Peer Context Memory Access
// Runtime API: 5.12. Peer Device Memory Access
CONV_PEER,
// Driver API : 5.27. Graphics Interoperability
// Runtime API: 5.23. Graphics Interoperability
CONV_GRAPHICS,
// Driver API : 5.28. Profiler Control
// Runtime API: 5.32. Profiler Control
CONV_PROFILER,
// Driver API : 5.29. OpenGL Interoperability
// Runtime API: 5.13. OpenGL Interoperability, 5.14. OpenGL Interoperability [DEPRECATED]
CONV_OPENGL,
// Driver API : 5.30. Direct3D 9 Interoperability
// Runtime API: 5.15. Direct3D 9 Interoperability, 5.16. Direct3D 9 Interoperability [DEPRECATED]
CONV_D3D9,
// Driver API : 5.31. Direct3D 10 Interoperability
// Runtime API: 5.17. Direct3D 10 Interoperability, 5.18. Direct3D 10 Interoperability [DEPRECATED]
CONV_D3D10,
// Driver API : 5.32. Direct3D 11 Interoperability
// Runtime API: 5.19. Direct3D 11 Interoperability, 5.20. Direct3D 11 Interoperability [DEPRECATED]
CONV_D3D11,
// Driver API : 5.33. VDPAU Interoperability
// Runtime API: 5.21. VDPAU Interoperability
CONV_VDPAU,
// Driver API : 5.34. EGL Interoperability
// Runtime API: 5.22. EGL Interoperability
CONV_EGL,
CONV_COMPLEX,
// Runtime API: 5.2. Thread Management [DEPRECATED]
CONV_THREAD,
CONV_OTHER,
CONV_COMPLEX,
CONV_LIB_FUNC,
CONV_LIB_DEVICE_FUNC,
CONV_INCLUDE,
CONV_INCLUDE_CUDA_MAIN_H,
CONV_TYPE,
CONV_LITERAL,
CONV_NUMERIC_LITERAL,
CONV_DEFINE,
CONV_LAST
};
constexpr int NUM_CONV_TYPES = (int) ConvTypes::CONV_LAST;
@@ -88,11 +134,11 @@ private:
int convTypeCounters[NUM_CONV_TYPES] = {};
public:
void incrementCounter(const hipCounter& counter, std::string name);
void incrementCounter(const hipCounter& counter, const std::string& name);
// Add the counters from `other` onto the counters of this object.
void add(const StatCounter& other);
int getConvSum();
void print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix);
void print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& prefix);
};
/**
@@ -110,8 +156,8 @@ class Statistics {
chr::steady_clock::time_point completionTime;
public:
Statistics(std::string name);
void incrementCounter(const hipCounter &counter, std::string name);
Statistics(const std::string& name);
void incrementCounter(const hipCounter &counter, const std::string& name);
// Add the counters from `other` onto the counters of this object.
void add(const Statistics &other);
void lineTouched(int lineNumber);
@@ -146,5 +192,5 @@ public:
* Set the active Statistics object to the named one, creating it if necessary, and write the completion
* timestamp into the currently active one.
*/
static void setActive(std::string name);
static void setActive(const std::string& name);
};
hipamd/hipify-clang/src/StringUtils.cpp
+1 -1
View File
@@ -7,7 +7,7 @@ llvm::StringRef unquoteStr(llvm::StringRef s) {
return s;
}
void removePrefixIfPresent(std::string &s, std::string prefix) {
void removePrefixIfPresent(std::string &s, const std::string& prefix) {
if (s.find(prefix) != 0) {
return;
}
hipamd/hipify-clang/src/StringUtils.h
+1 -1
View File
@@ -11,4 +11,4 @@ llvm::StringRef unquoteStr(llvm::StringRef s);
/**
* If `s` starts with `prefix`, remove it. Otherwise, does nothing.
*/
void removePrefixIfPresent(std::string &s, std::string prefix);
void removePrefixIfPresent(std::string &s, const std::string& prefix);
hipamd/include/hip/hcc_detail/code_object_bundle.hpp
+5
View File
@@ -84,6 +84,9 @@ class Bundled_code_header {
std::copy_n(f + y.header.offset, y.header.bundle_sz, std::back_inserter(y.blob));
it += y.header.triple_sz;
x.bundled_code_size = std::max(x.bundled_code_size,
y.header.offset + y.header.bundle_sz);
}
return true;
@@ -123,6 +126,8 @@ class Bundled_code_header {
// MANIPULATORS
Bundled_code_header& operator=(const Bundled_code_header&) = default;
Bundled_code_header& operator=(Bundled_code_header&&) = default;
size_t bundled_code_size = 0;
};
// CREATORS
hipamd/include/hip/hcc_detail/device_functions.h
+9
View File
@@ -876,6 +876,15 @@ static void __threadfence_system()
__atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_all_svm_devices);
}
// abort
__device__
inline
__attribute__((weak))
void abort() {
return __builtin_trap();
}
#endif // __HCC_OR_HIP_CLANG__
#ifdef __HCC__
hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
+34 -7
View File
@@ -33,6 +33,7 @@ THE SOFTWARE.
#include <cstddef>
#include <cstdint>
#include <cstring>
#include <functional>
#include <iostream>
#include <mutex>
@@ -56,7 +57,9 @@ template <
typename... Ts,
typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
inline std::vector<std::uint8_t> make_kernarg(
std::vector<std::uint8_t> kernarg, const std::tuple<Ts...>&) {
const std::tuple<Ts...>&,
const std::vector<std::pair<std::size_t, std::size_t>>&,
std::vector<std::uint8_t> kernarg) {
return kernarg;
}
@@ -65,7 +68,9 @@ template <
typename... Ts,
typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
inline std::vector<std::uint8_t> make_kernarg(
std::vector<std::uint8_t> kernarg, const std::tuple<Ts...>& formals) {
const std::tuple<Ts...>& formals,
const std::vector<std::pair<std::size_t, std::size_t>>& size_align,
std::vector<std::uint8_t> kernarg) {
using T = typename std::tuple_element<n, std::tuple<Ts...>>::type;
static_assert(
@@ -80,24 +85,44 @@ inline std::vector<std::uint8_t> make_kernarg(
#endif
kernarg.resize(round_up_to_next_multiple_nonnegative(
kernarg.size(), alignof(T)) + sizeof(T));
kernarg.size(), size_align[n].second) +
size_align[n].first);
new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::get<n>(formals)};
std::memcpy(
kernarg.data() + kernarg.size() - size_align[n].first,
&std::get<n>(formals),
size_align[n].first);
return make_kernarg<n + 1>(std::move(kernarg), formals);
return make_kernarg<n + 1>(formals, size_align, std::move(kernarg));
}
template <typename... Formals, typename... Actuals>
inline std::vector<std::uint8_t> make_kernarg(
void (*)(Formals...), std::tuple<Actuals...> actuals) {
void (*kernel)(Formals...), std::tuple<Actuals...> actuals) {
static_assert(sizeof...(Formals) == sizeof...(Actuals),
"The count of formal arguments must match the count of actuals.");
if (sizeof...(Formals) == 0) return {};
const auto it = function_names().find(
reinterpret_cast<std::uintptr_t>(kernel));
if (it == function_names().cend()) {
throw std::runtime_error{"Undefined __global__ function."};
}
const auto it1 = kernargs().find(it->second);
if (it1 == kernargs().end()) {
throw std::runtime_error{
"Missing metadata for __global__ function: " + it->second};
}
std::tuple<Formals...> to_formals{std::move(actuals)};
std::vector<std::uint8_t> kernarg;
kernarg.reserve(sizeof(to_formals));
return make_kernarg<0>(std::move(kernarg), to_formals);
return make_kernarg<0>(to_formals, it1->second, std::move(kernarg));
}
void hipLaunchKernelGGLImpl(std::uintptr_t function_address, const dim3& numBlocks,
@@ -120,6 +145,8 @@ inline void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimB
}
template <typename... Args, typename F = void (*)(hipLaunchParm, Args...)>
[[deprecated("hipLaunchKernel is deprecated and will be removed in the next "
"version of HIP; please upgrade to hipLaunchKernelGGL.")]]
inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t groupMemBytes, hipStream_t stream, Args... args) {
hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
hipamd/include/hip/hcc_detail/hip_memory.h
+7 -2
View File
@@ -27,6 +27,8 @@ THE SOFTWARE.
// HIP heap is implemented as a global array with fixed size. Users may define
// __HIP_SIZE_OF_PAGE and __HIP_NUM_PAGES to have a larger heap.
#if __HCC__ || __HIP__
// Size of page in bytes.
#ifndef __HIP_SIZE_OF_PAGE
#define __HIP_SIZE_OF_PAGE 64
@@ -39,11 +41,14 @@ THE SOFTWARE.
#define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE)
#if __HCC__ || __HIP__
#if __HIP__ && __HIP_DEVICE_COMPILE__
__attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
__attribute__((weak)) __device__
uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
#else
extern __device__ char __hip_device_heap[];
extern __device__ uint32_t __hip_device_page_flag[];
#endif
extern "C" inline __device__ void* __hip_malloc(size_t size) {
char* heap = (char*)__hip_device_heap;
hipamd/include/hip/hcc_detail/hip_runtime.h
+1 -2
View File
@@ -181,8 +181,7 @@ extern int HIP_TRACE_API;
#define __HCC_C__
#endif
// abort
__device__ void abort();
__host__ inline void* __get_dynamicgroupbaseptr() { return nullptr; }
#if __HIP_ARCH_GFX701__ == 0
hipamd/include/hip/hcc_detail/hip_runtime_api.h
+2 -2
View File
@@ -94,7 +94,7 @@ typedef struct ihipModule_t* hipModule_t;
typedef struct ihipModuleSymbol_t* hipFunction_t;
struct hipFuncAttributes {
typedef struct hipFuncAttributes {
int binaryVersion;
int cacheModeCA;
size_t constSizeBytes;
@@ -105,7 +105,7 @@ struct hipFuncAttributes {
int preferredShmemCarveout;
int ptxVersion;
size_t sharedSizeBytes;
};
} hipFuncAttributes;
typedef struct ihipEvent_t* hipEvent_t;
hipamd/include/hip/hcc_detail/math_functions.h
+38 -32
View File
@@ -514,38 +514,41 @@ float __exp10f(float x) { return __ocml_exp10_f32(x); }
__DEVICE__
inline
float __expf(float x) { return __ocml_exp_f32(x); }
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
inline
float __fadd_rd(float x, float y) { return __ocml_add_rtp_f32(x, y); }
float __fadd_rd(float x, float y) { return __ocml_add_rtn_f32(x, y); }
__DEVICE__
inline
float __fadd_rn(float x, float y) { return __ocml_add_rte_f32(x, y); }
__DEVICE__
inline
float __fadd_ru(float x, float y) { return __ocml_add_rtn_f32(x, y); }
float __fadd_ru(float x, float y) { return __ocml_add_rtp_f32(x, y); }
__DEVICE__
inline
float __fadd_rz(float x, float y) { return __ocml_add_rtz_f32(x, y); }
__DEVICE__
inline
float __fdiv_rd(float x, float y) { return x / y; }
float __fdiv_rd(float x, float y) { return __ocml_div_rtn_f32(x, y); }
__DEVICE__
inline
float __fdiv_rn(float x, float y) { return x / y; }
float __fdiv_rn(float x, float y) { return __ocml_div_rte_f32(x, y); }
__DEVICE__
inline
float __fdiv_ru(float x, float y) { return x / y; }
float __fdiv_ru(float x, float y) { return __ocml_div_rtp_f32(x, y); }
__DEVICE__
inline
float __fdiv_rz(float x, float y) { return x / y; }
float __fdiv_rz(float x, float y) { return __ocml_div_rtz_f32(x, y); }
#endif
__DEVICE__
inline
float __fdividef(float x, float y) { return x / y; }
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
inline
float __fmaf_rd(float x, float y, float z)
{
return __ocml_fma_rtp_f32(x, y, z);
return __ocml_fma_rtn_f32(x, y, z);
}
__DEVICE__
inline
@@ -557,7 +560,7 @@ __DEVICE__
inline
float __fmaf_ru(float x, float y, float z)
{
return __ocml_fma_rtn_f32(x, y, z);
return __ocml_fma_rtp_f32(x, y, z);
}
__DEVICE__
inline
@@ -567,13 +570,13 @@ float __fmaf_rz(float x, float y, float z)
}
__DEVICE__
inline
float __fmul_rd(float x, float y) { return __ocml_mul_rtp_f32(x, y); }
float __fmul_rd(float x, float y) { return __ocml_mul_rtn_f32(x, y); }
__DEVICE__
inline
float __fmul_rn(float x, float y) { return __ocml_mul_rte_f32(x, y); }
__DEVICE__
inline
float __fmul_ru(float x, float y) { return __ocml_mul_rtn_f32(x, y); }
float __fmul_ru(float x, float y) { return __ocml_mul_rtp_f32(x, y); }
__DEVICE__
inline
float __fmul_rz(float x, float y) { return __ocml_mul_rtz_f32(x, y); }
@@ -594,28 +597,29 @@ inline
float __frsqrt_rn(float x) { return __llvm_amdgcn_rsq_f32(x); }
__DEVICE__
inline
float __fsqrt_rd(float x) { return __ocml_sqrt_f32(x); }
float __fsqrt_rd(float x) { return __ocml_sqrt_rtn_f32(x); }
__DEVICE__
inline
float __fsqrt_rn(float x) { return __ocml_sqrt_f32(x); }
float __fsqrt_rn(float x) { return __ocml_sqrt_rte_f32(x); }
__DEVICE__
inline
float __fsqrt_ru(float x) { return __ocml_sqrt_f32(x); }
float __fsqrt_ru(float x) { return __ocml_sqrt_rtp_f32(x); }
__DEVICE__
inline
float __fsqrt_rz(float x) { return __ocml_sqrt_f32(x); }
float __fsqrt_rz(float x) { return __ocml_sqrt_rtz_f32(x); }
__DEVICE__
inline
float __fsub_rd(float x, float y) { return __ocml_sub_rtp_f32(x, y); }
float __fsub_rd(float x, float y) { return __ocml_sub_rtn_f32(x, y); }
__DEVICE__
inline
float __fsub_rn(float x, float y) { return __ocml_sub_rte_f32(x, y); }
__DEVICE__
inline
float __fsub_ru(float x, float y) { return __ocml_sub_rtn_f32(x, y); }
float __fsub_ru(float x, float y) { return __ocml_sub_rtp_f32(x, y); }
__DEVICE__
inline
float __fsub_rz(float x, float y) { return __ocml_sub_rtz_f32(x, y); }
#endif
__DEVICE__
inline
float __log10f(float x) { return __ocml_log10_f32(x); }
@@ -1034,39 +1038,40 @@ double yn(int n, double x)
}
// BEGIN INTRINSICS
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__DEVICE__
inline
double __dadd_rd(double x, double y) { return __ocml_add_rtp_f64(x, y); }
double __dadd_rd(double x, double y) { return __ocml_add_rtn_f64(x, y); }
__DEVICE__
inline
double __dadd_rn(double x, double y) { return __ocml_add_rte_f64(x, y); }
__DEVICE__
inline
double __dadd_ru(double x, double y) { return __ocml_add_rtn_f64(x, y); }
double __dadd_ru(double x, double y) { return __ocml_add_rtp_f64(x, y); }
__DEVICE__
inline
double __dadd_rz(double x, double y) { return __ocml_add_rtz_f64(x, y); }
__DEVICE__
inline
double __ddiv_rd(double x, double y) { return x / y; }
double __ddiv_rd(double x, double y) { return __ocml_div_rtn_f64(x, y); }
__DEVICE__
inline
double __ddiv_rn(double x, double y) { return x / y; }
double __ddiv_rn(double x, double y) { return __ocml_div_rte_f64(x, y); }
__DEVICE__
inline
double __ddiv_ru(double x, double y) { return x / y; }
double __ddiv_ru(double x, double y) { return __ocml_div_rtp_f64(x, y); }
__DEVICE__
inline
double __ddiv_rz(double x, double y) { return x / y; }
double __ddiv_rz(double x, double y) { return __ocml_div_rtz_f64(x, y); }
__DEVICE__
inline
double __dmul_rd(double x, double y) { return __ocml_mul_rtp_f64(x, y); }
double __dmul_rd(double x, double y) { return __ocml_mul_rtn_f64(x, y); }
__DEVICE__
inline
double __dmul_rn(double x, double y) { return __ocml_mul_rte_f64(x, y); }
__DEVICE__
inline
double __dmul_ru(double x, double y) { return __ocml_mul_rtn_f64(x, y); }
double __dmul_ru(double x, double y) { return __ocml_mul_rtp_f64(x, y); }
__DEVICE__
inline
double __dmul_rz(double x, double y) { return __ocml_mul_rtz_f64(x, y); }
@@ -1084,25 +1089,25 @@ inline
double __drcp_rz(double x) { return __llvm_amdgcn_rcp_f64(x); }
__DEVICE__
inline
double __dsqrt_rd(double x) { return __ocml_sqrt_f64(x); }
double __dsqrt_rd(double x) { return __ocml_sqrt_rtn_f64(x); }
__DEVICE__
inline
double __dsqrt_rn(double x) { return __ocml_sqrt_f64(x); }
double __dsqrt_rn(double x) { return __ocml_sqrt_rte_f64(x); }
__DEVICE__
inline
double __dsqrt_ru(double x) { return __ocml_sqrt_f64(x); }
double __dsqrt_ru(double x) { return __ocml_sqrt_rtp_f64(x); }
__DEVICE__
inline
double __dsqrt_rz(double x) { return __ocml_sqrt_f64(x); }
double __dsqrt_rz(double x) { return __ocml_sqrt_rtz_f64(x); }
__DEVICE__
inline
double __dsub_rd(double x, double y) { return __ocml_sub_rtp_f64(x, y); }
double __dsub_rd(double x, double y) { return __ocml_sub_rtn_f64(x, y); }
__DEVICE__
inline
double __dsub_rn(double x, double y) { return __ocml_sub_rte_f64(x, y); }
__DEVICE__
inline
double __dsub_ru(double x, double y) { return __ocml_sub_rtn_f64(x, y); }
double __dsub_ru(double x, double y) { return __ocml_sub_rtp_f64(x, y); }
__DEVICE__
inline
double __dsub_rz(double x, double y) { return __ocml_sub_rtz_f64(x, y); }
@@ -1110,7 +1115,7 @@ __DEVICE__
inline
double __fma_rd(double x, double y, double z)
{
return __ocml_fma_rtp_f64(x, y, z);
return __ocml_fma_rtn_f64(x, y, z);
}
__DEVICE__
inline
@@ -1122,7 +1127,7 @@ __DEVICE__
inline
double __fma_ru(double x, double y, double z)
{
return __ocml_fma_rtn_f64(x, y, z);
return __ocml_fma_rtp_f64(x, y, z);
}
__DEVICE__
inline
@@ -1130,6 +1135,7 @@ double __fma_rz(double x, double y, double z)
{
return __ocml_fma_rtz_f64(x, y, z);
}
#endif
// END INTRINSICS
// END DOUBLE
hipamd/include/hip/hcc_detail/math_fwd.h
+49 -1
View File
@@ -288,6 +288,30 @@ __attribute__((const))
float __ocml_mul_rtz_f32(float, float);
__device__
__attribute__((const))
float __ocml_div_rte_f32(float, float);
__device__
__attribute__((const))
float __ocml_div_rtn_f32(float, float);
__device__
__attribute__((const))
float __ocml_div_rtp_f32(float, float);
__device__
__attribute__((const))
float __ocml_div_rtz_f32(float, float);
__device__
__attribute__((const))
float __ocml_sqrt_rte_f32(float, float);
__device__
__attribute__((const))
float __ocml_sqrt_rtn_f32(float, float);
__device__
__attribute__((const))
float __ocml_sqrt_rtp_f32(float, float);
__device__
__attribute__((const))
float __ocml_sqrt_rtz_f32(float, float);
__device__
__attribute__((const))
float __ocml_fma_rte_f32(float, float, float);
__device__
__attribute__((const))
@@ -572,6 +596,30 @@ __attribute__((const))
double __ocml_mul_rtz_f64(double, double);
__device__
__attribute__((const))
double __ocml_div_rte_f64(double, double);
__device__
__attribute__((const))
double __ocml_div_rtn_f64(double, double);
__device__
__attribute__((const))
double __ocml_div_rtp_f64(double, double);
__device__
__attribute__((const))
double __ocml_div_rtz_f64(double, double);
__device__
__attribute__((const))
double __ocml_sqrt_rte_f64(double, double);
__device__
__attribute__((const))
double __ocml_sqrt_rtn_f64(double, double);
__device__
__attribute__((const))
double __ocml_sqrt_rtp_f64(double, double);
__device__
__attribute__((const))
double __ocml_sqrt_rtz_f64(double, double);
__device__
__attribute__((const))
double __ocml_fma_rte_f64(double, double, double);
__device__
__attribute__((const))
@@ -594,4 +642,4 @@ double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64");
#if defined(__cplusplus)
} // extern "C"
#endif
#endif
hipamd/include/hip/hcc_detail/program_state.hpp
+2
View File
@@ -99,6 +99,8 @@ const std::unordered_map<std::uintptr_t, std::vector<std::pair<hsa_agent_t, Kern
functions(bool rebuild = false);
const std::unordered_map<std::uintptr_t, std::string>& function_names(bool rebuild = false);
std::unordered_map<std::string, void*>& globals(bool rebuild = false);
std::unordered_map<
std::string, std::vector<std::pair<std::size_t, std::size_t>>>& kernargs();
hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable,
hsa_agent_t agent);
hipamd/packaging/hip-targets-release.cmake
-10
View File
@@ -41,15 +41,5 @@ endif()
list(APPEND _IMPORT_CHECK_TARGETS hip::hip_hcc )
list(APPEND _IMPORT_CHECK_FILES_FOR_hip::hip_hcc "/opt/rocm/hip/lib/libhip_hcc.so" )
# Import target "hip::hip_device" for configuration "Release"
set_property(TARGET hip::hip_device APPEND PROPERTY IMPORTED_CONFIGURATIONS RELEASE)
set_target_properties(hip::hip_device PROPERTIES
IMPORTED_LINK_INTERFACE_LANGUAGES_RELEASE "CXX"
IMPORTED_LOCATION_RELEASE "/opt/rocm/hip/lib/libhip_device.a"
)
list(APPEND _IMPORT_CHECK_TARGETS hip::hip_device )
list(APPEND _IMPORT_CHECK_FILES_FOR_hip::hip_device "/opt/rocm/hip/lib/libhip_device.a" )
# Commands beyond this point should not need to know the version.
set(CMAKE_IMPORT_FILE_VERSION)
hipamd/packaging/hip-targets.cmake
+4 -12
View File
@@ -16,7 +16,7 @@ set(CMAKE_IMPORT_FILE_VERSION 1)
set(_targetsDefined)
set(_targetsNotDefined)
set(_expectedTargets)
foreach(_expectedTarget hip::hip_hcc_static hip::hip_hcc hip::hip_device hip::host hip::device)
foreach(_expectedTarget hip::hip_hcc_static hip::hip_hcc hip::host hip::device)
list(APPEND _expectedTargets ${_expectedTarget})
if(NOT TARGET ${_expectedTarget})
list(APPEND _targetsNotDefined ${_expectedTarget})
@@ -57,14 +57,6 @@ set_target_properties(hip::hip_hcc PROPERTIES
INTERFACE_SYSTEM_INCLUDE_DIRECTORIES "${_IMPORT_PREFIX}/include;/opt/rocm/hsa/include"
)
# Create imported target hip::hip_device
add_library(hip::hip_device STATIC IMPORTED)
set_target_properties(hip::hip_device PROPERTIES
INTERFACE_INCLUDE_DIRECTORIES "${_IMPORT_PREFIX}/include;/opt/rocm/hsa/include"
INTERFACE_SYSTEM_INCLUDE_DIRECTORIES "${_IMPORT_PREFIX}/include;/opt/rocm/hsa/include"
)
# Create imported target hip::host
add_library(hip::host INTERFACE IMPORTED)
@@ -75,13 +67,13 @@ set_target_properties(hip::host PROPERTIES
# Create imported target hip::device
add_library(hip::device INTERFACE IMPORTED)
if(HIP_COMPILER STREQUAL "clang")
if(HIP_COMPILER STREQUAL "hcc")
set_target_properties(hip::device PROPERTIES
INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device"
INTERFACE_LINK_LIBRARIES "hip::host;hcc::hccrt;hcc::hc_am"
)
else()
set_target_properties(hip::device PROPERTIES
INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device;hcc::hccrt;hcc::hc_am"
INTERFACE_LINK_LIBRARIES "hip::host"
)
endif()
hipamd/packaging/hip_hcc.txt
-1
View File
@@ -3,7 +3,6 @@ project(hip_hcc)
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip)
install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/packaging/hip-targets-release.cmake DESTINATION lib/cmake/hip)
hipamd/samples/0_Intro/bit_extract/bit_extract.cpp
+1 -5
View File
@@ -23,10 +23,6 @@ THE SOFTWARE.
#include <stdio.h>
#include <iostream>
#include "hip/hip_runtime.h"
#ifdef __HIP_PLATFORM_HCC__
#include <hc.hpp>
#endif
#define CHECK(cmd) \
{ \
@@ -44,7 +40,7 @@ __global__ void bit_extract_kernel(uint32_t* C_d, const uint32_t* A_d, size_t N)
for (size_t i = offset; i < N; i += stride) {
#ifdef __HIP_PLATFORM_HCC__
C_d[i] = hc::__bitextract_u32(A_d[i], 8, 4);
C_d[i] = __bitextract_u32(A_d[i], 8, 4);
#else /* defined __HIP_PLATFORM_NVCC__ or other path */
C_d[i] = ((A_d[i] & 0xf00) >> 8);
#endif
hipamd/src/device_util.cpp
-33
View File
@@ -1,33 +0,0 @@
/*
Copyright (c) 2015-2017 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.
*/
#include <hc.hpp>
#include "device_util.h"
#include "hip/hcc_detail/device_functions.h"
#include "hip/hcc_detail/grid_launch.h"
#include "hip/hip_runtime.h"
#include <atomic>
// abort
__device__ void abort() { return hc::abort(); }
__host__ void* __get_dynamicgroupbaseptr() { return nullptr; }
hipamd/src/hip_clang.cpp
+57 -13
View File
@@ -22,6 +22,7 @@ THE SOFTWARE.
#include <unordered_map>
#include <string>
#include <fstream>
#include "hip/hip_runtime.h"
#include "hip_hcc_internal.h"
@@ -86,6 +87,7 @@ __hipRegisterFatBinary(const void* data)
std::string target{&desc->triple[sizeof(AMDGCN_AMDHSA_TRIPLE)],
desc->tripleSize - sizeof(AMDGCN_AMDHSA_TRIPLE)};
tprintf(DB_FB, "Found bundle for %s\n", target.c_str());
for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
hsa_agent_t agent = g_allAgents[deviceId + 1];
@@ -110,10 +112,35 @@ __hipRegisterFatBinary(const void* data)
if (module->executable.handle) {
modules->at(deviceId) = module;
tprintf(DB_FB, "Loaded code object for %s\n", name);
if (HIP_DUMP_CODE_OBJECT) {
char fname[30];
static std::atomic<int> index;
sprintf(fname, "__hip_dump_code_object%04d.o", index++);
tprintf(DB_FB, "Dump code object %s\n", fname);
std::ofstream ofs;
ofs.open(fname, std::ios::binary);
ofs << image;
ofs.close();
}
} else {
fprintf(stderr, "Failed to load code object for %s\n", name);
abort();
}
}
}
for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
hsa_agent_t agent = g_allAgents[deviceId + 1];
char name[64] = {};
hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, name);
if (!(*modules)[deviceId]) {
fprintf(stderr, "No device code bundle for %s\n", name);
abort();
}
}
tprintf(DB_FB, "__hipRegisterFatBinary succeeds and returns %p\n", modules);
return modules;
}
@@ -132,13 +159,20 @@ extern "C" void __hipRegisterFunction(
dim3* gridDim,
int* wSize)
{
HIP_INIT_API(modules, hostFunction, deviceFunction, deviceName);
std::vector<hipFunction_t> functions{g_deviceCnt};
assert(modules && modules->size() >= g_deviceCnt);
for (int deviceId = 0; deviceId < g_deviceCnt; ++deviceId) {
hipFunction_t function;
if (hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName)) {
if (hipSuccess == hipModuleGetFunction(&function, modules->at(deviceId), deviceName) &&
function != nullptr) {
functions[deviceId] = function;
}
else {
tprintf(DB_FB, "__hipRegisterFunction cannot find kernel %s for"
" device %d\n", deviceName, deviceId);
}
}
g_functions.insert(std::make_pair(hostFunction, std::move(functions)));
@@ -180,6 +214,7 @@ hipError_t hipSetupArgument(
size_t size,
size_t offset)
{
HIP_INIT_API(arg, size, offset);
auto ctx = ihipGetTlsDefaultCtx();
LockedAccessor_CtxCrit_t crit(ctx->criticalData());
auto& arguments = crit->_execStack.top()._arguments;
@@ -194,6 +229,7 @@ hipError_t hipSetupArgument(
hipError_t hipLaunchByPtr(const void *hostFunction)
{
HIP_INIT_API(hostFunction);
ihipExec_t exec;
{
auto ctx = ihipGetTlsDefaultCtx();
@@ -213,20 +249,28 @@ hipError_t hipLaunchByPtr(const void *hostFunction)
deviceId = 0;
}
hipError_t e = hipSuccess;
decltype(g_functions)::iterator it;
if ((it = g_functions.find(hostFunction)) == g_functions.end())
return hipErrorUnknown;
if ((it = g_functions.find(hostFunction)) == g_functions.end() ||
!it->second[deviceId]) {
e = hipErrorUnknown;
fprintf(stderr, "hipLaunchByPtr cannot find kernel with stub address %p"
" for device %d!\n", hostFunction, deviceId);
abort();
} else {
size_t size = exec._arguments.size();
void *extra[] = {
HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec._arguments[0],
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END
};
size_t size = exec._arguments.size();
void *extra[] = {
HIP_LAUNCH_PARAM_BUFFER_POINTER, &exec._arguments[0],
HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END
};
e = hipModuleLaunchKernel(it->second[deviceId],
exec._gridDim.x, exec._gridDim.y, exec._gridDim.z,
exec._blockDim.x, exec._blockDim.y, exec._blockDim.z,
exec._sharedMem, exec._hStream, nullptr, extra);
}
return hipModuleLaunchKernel(it->second[deviceId],
exec._gridDim.x, exec._gridDim.y, exec._gridDim.z,
exec._blockDim.x, exec._blockDim.y, exec._blockDim.z,
exec._sharedMem, exec._hStream, nullptr, extra);
return ihipLogStatus(e);
}
hipamd/src/hip_hcc.cpp
+6
View File
@@ -97,6 +97,8 @@ int HIP_INIT_ALLOC = -1;
int HIP_SYNC_STREAM_WAIT = 0;
int HIP_FORCE_NULL_STREAM = 0;
int HIP_DUMP_CODE_OBJECT = 0;
#if (__hcc_workweek__ >= 17300)
// Make sure we have required bug fix in HCC
@@ -1294,6 +1296,10 @@ void HipReadEnv() {
"overridden by specifying hipEventReleaseToSystem or hipEventReleaseToDevice flag "
"when creating the event.");
READ_ENV_I(release, HIP_DUMP_CODE_OBJECT, 0,
"If set, dump code object as __hip_dump_code_object[nnnn].o in the current directory,"
"where nnnn is the index number.");
// Some flags have both compile-time and runtime flags - generate a warning if user enables the
// runtime flag but the compile-time flag is disabled.
if (HIP_DB && !COMPILE_HIP_DB) {
hipamd/src/hip_hcc_internal.h
+1 -1
View File
@@ -83,11 +83,11 @@ extern int HIP_SYNC_NULL_STREAM;
extern int HIP_INIT_ALLOC;
extern int HIP_FORCE_NULL_STREAM;
extern int HIP_DUMP_CODE_OBJECT;
// TODO - remove when this is standard behavior.
extern int HCC_OPT_FLUSH;
// Class to assign a short TID to each new thread, for HIP debugging purposes.
class TidInfo {
public:
hipamd/src/hip_memory.cpp
+14 -45
View File
@@ -28,6 +28,9 @@ THE SOFTWARE.
#include "hip_hcc_internal.h"
#include "trace_helper.h"
__device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
__device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
// Internal HIP APIS:
namespace hip_internal {
@@ -574,24 +577,8 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si
size = size * height;
}
size_t allocSize = 0;
switch (desc->f) {
case hipChannelFormatKindSigned:
allocSize = size * sizeof(int);
break;
case hipChannelFormatKindUnsigned:
allocSize = size * sizeof(unsigned int);
break;
case hipChannelFormatKindFloat:
allocSize = size * sizeof(float);
break;
case hipChannelFormatKindNone:
allocSize = size * sizeof(size_t);
break;
default:
hip_status = hipErrorUnknown;
break;
}
const size_t allocSize = size * ((desc->x + desc->y + desc->z + desc->w) / 8);
hc::accelerator acc = ctx->getDevice()->_acc;
hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
@@ -800,24 +787,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
const unsigned am_flags = 0;
const size_t size = extent.width * extent.height * extent.depth;
size_t allocSize = 0;
switch (desc->f) {
case hipChannelFormatKindSigned:
allocSize = size * sizeof(int);
break;
case hipChannelFormatKindUnsigned:
allocSize = size * sizeof(unsigned int);
break;
case hipChannelFormatKindFloat:
allocSize = size * sizeof(float);
break;
case hipChannelFormatKindNone:
allocSize = size * sizeof(size_t);
break;
default:
hip_status = hipErrorUnknown;
break;
}
const size_t allocSize = size * ((desc->x + desc->y + desc->z + desc->w) / 8);
hc::accelerator acc = ctx->getDevice()->_acc;
hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
@@ -1015,10 +985,9 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t cou
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
if (kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost ||
if (kind == hipMemcpyHostToDevice || kind == hipMemcpyDefault ||
kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) {
stream->lockedSymbolCopySync(acc, dst, (void*)src, count, offset, kind);
// acc.memcpy_symbol(dst, (void*)src, count+offset);
stream->locked_copySync((char*)dst+offset, (void*)src, count, kind, false);
} else {
return ihipLogStatus(hipErrorInvalidValue);
}
@@ -1048,9 +1017,9 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count,
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
if (kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost ||
if (kind == hipMemcpyDefault || kind == hipMemcpyDeviceToHost ||
kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) {
stream->lockedSymbolCopySync(acc, dst, (void*)src, count, offset, kind);
stream->locked_copySync((void*)dst, (char*)src+offset, count, kind, false);
} else {
return ihipLogStatus(hipErrorInvalidValue);
}
@@ -1082,7 +1051,7 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_
if (stream) {
try {
stream->lockedSymbolCopyAsync(acc, dst, (void*)src, count, offset, kind);
hip_internal::memcpyAsync((char*)dst+offset, src, count, kind, stream);
} catch (ihipException& ex) {
e = ex._code;
}
@@ -1118,7 +1087,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co
stream = ihipSyncAndResolveStream(stream);
if (stream) {
try {
stream->lockedSymbolCopyAsync(acc, dst, src, count, offset, kind);
hip_internal::memcpyAsync(dst, (char*)src+offset, count, kind, stream);
} catch (ihipException& ex) {
e = ex._code;
}
@@ -1688,7 +1657,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
} else {
try {
if(!isLocked){
for (int i = 0; i < height; ++i)
for (int i = 0; i < height; ++i)
e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch,
(unsigned char*)src + i * spitch, width, kind, stream);
} else{
@@ -1738,7 +1707,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
stream->locked_wait();
} else {
e = hipErrorInvalidValue;
}
}
return ihipLogStatus(e);
}
hipamd/src/hip_module.cpp
+16 -7
View File
@@ -258,20 +258,29 @@ struct Agent_global {
uint32_t byte_cnt;
};
inline void track(const Agent_global& x) {
inline void track(const Agent_global& x, hsa_agent_t agent) {
tprintf(DB_MEM, " add variable '%s' with ptr=%p size=%u to tracker\n", x.name.c_str(),
x.address, x.byte_cnt);
auto device = ihipGetTlsDefaultCtx()->getWriteableDevice();
int deviceIndex =0;
for ( deviceIndex = 0; deviceIndex < g_deviceCnt; deviceIndex++) {
if(g_allAgents[deviceIndex] == agent)
break;
}
auto device = ihipGetDevice(deviceIndex - 1);
hc::AmPointerInfo ptr_info(nullptr, x.address, x.address, x.byte_cnt, device->_acc, true,
false);
hc::am_memtracker_add(x.address, ptr_info);
#if USE_APP_PTR_FOR_CTX
hc::am_memtracker_update(x.address, device->_deviceId, 0u, ihipGetTlsDefaultCtx());
#else
hc::am_memtracker_update(x.address, device->_deviceId, 0u);
#endif
}
template <typename Container = vector<Agent_global>>
inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t,
inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t agent,
hsa_executable_symbol_t x, void* out) {
assert(out);
@@ -281,7 +290,7 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t,
if (t == HSA_SYMBOL_KIND_VARIABLE) {
static_cast<Container*>(out)->push_back(Agent_global{name(x), address(x), size(x)});
track(static_cast<Container*>(out)->back());
track(static_cast<Container*>(out)->back(),agent);
}
return HSA_STATUS_SUCCESS;
@@ -342,7 +351,7 @@ hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes, hi
tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(), it0->second.cend(), name);
return dptr ? hipSuccess : hipErrorNotFound;
return *dptr ? hipSuccess : hipErrorNotFound;
}
hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, const char* name) {
@@ -367,7 +376,7 @@ hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, c
tie(*dptr, *bytes) = read_global_description(it->second.cbegin(), it->second.cend(), name);
return dptr ? hipSuccess : hipErrorNotFound;
return *dptr ? hipSuccess : hipErrorNotFound;
}
hsa_executable_symbol_t find_kernel_by_name(hsa_executable_t executable, const char* kname) {
hipamd/src/program_state.cpp
+115 -6
View File
@@ -209,10 +209,16 @@ const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs(bool reb
nullptr);
for (auto&& blob : blobs) {
Bundled_code_header tmp{blob};
if (valid(tmp)) {
for (auto&& bundle : bundles(tmp)) {
r[triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob);
for (auto sub_blob = blob.begin(); sub_blob != blob.end(); ) {
Bundled_code_header tmp(sub_blob, blob.end());
if (valid(tmp)) {
for (auto&& bundle : bundles(tmp)) {
r[triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob);
}
sub_blob+=tmp.bundled_code_size;
}
else {
break;
}
}
}
@@ -306,8 +312,8 @@ const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels(bool rebui
void load_code_object_and_freeze_executable(
const string& file, hsa_agent_t agent,
hsa_executable_t
executable) { // TODO: the following sequence is inefficient, should be refactored
hsa_executable_t executable) {
// TODO: the following sequence is inefficient, should be refactored
// into a single load of the file and subsequent ELFIO
// processing.
static const auto cor_deleter = [](hsa_code_object_reader_t* p) {
@@ -334,6 +340,90 @@ void load_code_object_and_freeze_executable(
code_readers.push_back(move(tmp));
}
}
size_t parse_args(
const string& metadata,
size_t f,
size_t l,
vector<pair<size_t, size_t>>& size_align) {
if (f == l) return f;
if (!size_align.empty()) return l;
do {
static constexpr size_t size_sz{5};
f = metadata.find("Size:", f) + size_sz;
if (l <= f) return f;
auto size = strtoul(&metadata[f], nullptr, 10);
static constexpr size_t align_sz{6};
f = metadata.find("Align:", f) + align_sz;
char* l{};
auto align = strtoul(&metadata[f], &l, 10);
f += (l - &metadata[f]) + 1;
size_align.emplace_back(size, align);
} while (true);
}
void read_kernarg_metadata(
elfio& reader,
unordered_map<string, vector<pair<size_t, size_t>>>& kernargs)
{ // TODO: this is inefficient.
auto it = find_section_if(
reader, [](const section* x) { return x->get_type() == SHT_NOTE; });
if (!it) return;
const note_section_accessor acc{reader, it};
for (decltype(acc.get_notes_num()) i = 0; i != acc.get_notes_num(); ++i) {
ELFIO::Elf_Word type{};
string name{};
void* desc{};
Elf_Word desc_size{};
acc.get_note(i, type, name, desc, desc_size);
if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA.
string tmp{
static_cast<char*>(desc), static_cast<char*>(desc) + desc_size};
auto dx = tmp.find("Kernels:");
if (dx == string::npos) continue;
static constexpr decltype(tmp.size()) kernels_sz{8};
dx += kernels_sz;
do {
dx = tmp.find("Name:", dx);
if (dx == string::npos) break;
static constexpr decltype(tmp.size()) name_sz{5};
dx = tmp.find_first_not_of(" '", dx + name_sz);
auto fn = tmp.substr(dx, tmp.find_first_of("'\n", dx) - dx);
dx += fn.size();
auto dx1 = tmp.find("CodeProps", dx);
dx = tmp.find("Args:", dx);
if (dx1 < dx) {
dx = dx1;
continue;
}
if (dx == string::npos) break;
static constexpr decltype(tmp.size()) args_sz{5};
dx = parse_args(tmp, dx + args_sz, dx1, kernargs[fn]);
} while (true);
}
}
} // namespace
namespace hip_impl {
@@ -495,6 +585,25 @@ unordered_map<string, void*>& globals(bool rebuild) {
return r;
}
unordered_map<string, vector<pair<size_t, size_t>>>& kernargs() {
static unordered_map<string, vector<pair<size_t, size_t>>> r;
static once_flag f;
call_once(f, []() {
for (auto&& blob : code_object_blobs()) {
stringstream tmp{std::string{
blob.second.front().cbegin(), blob.second.front().cend()}};
elfio reader;
if (!reader.load(tmp)) continue;
read_kernarg_metadata(reader, r);
}
});
return r;
}
hsa_executable_t load_executable(const string& file, hsa_executable_t executable,
hsa_agent_t agent) {
elfio reader;
hipamd/tests/src/deviceLib/hipDoublePrecisionIntrinsics.cpp
+2
View File
@@ -34,6 +34,7 @@ THE SOFTWARE.
#pragma clang diagnostic ignored "-Wunused-variable"
__device__ void double_precision_intrinsics() {
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__dadd_rd(0.0, 1.0);
__dadd_rn(0.0, 1.0);
__dadd_ru(0.0, 1.0);
@@ -62,6 +63,7 @@ __device__ void double_precision_intrinsics() {
__fma_rn(1.0, 2.0, 3.0);
__fma_ru(1.0, 2.0, 3.0);
__fma_rz(1.0, 2.0, 3.0);
#endif
}
__global__ void compileDoublePrecisionIntrinsics(int ignored) {
hipamd/tests/src/deviceLib/hipFloatMath.cpp
+6 -4
View File
@@ -38,11 +38,13 @@ __global__ void floatMath(float* In, float* Out) {
Out[tid] = __cosf(In[tid]);
Out[tid] = __exp10f(Out[tid]);
Out[tid] = __expf(Out[tid]);
#if defined OCML_BASIC_ROUNDED_OPERATIONS
Out[tid] = __frsqrt_rn(Out[tid]);
//Out[tid] = __fsqrt_rd(Out[tid]);
//Out[tid] = __fsqrt_rn(Out[tid]);
//Out[tid] = __fsqrt_ru(Out[tid]);
//Out[tid] = __fsqrt_rz(Out[tid]);
Out[tid] = __fsqrt_rd(Out[tid]);
Out[tid] = __fsqrt_rn(Out[tid]);
Out[tid] = __fsqrt_ru(Out[tid]);
Out[tid] = __fsqrt_rz(Out[tid]);
#endif
Out[tid] = __log10f(Out[tid]);
Out[tid] = __log2f(Out[tid]);
Out[tid] = __logf(Out[tid]);
hipamd/tests/src/deviceLib/hipSinglePrecisionIntrinsics.cpp
+4
View File
@@ -39,6 +39,7 @@ __device__ void single_precision_intrinsics() {
__cosf(0.0f);
__exp10f(0.0f);
__expf(0.0f);
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__fadd_rd(0.0f, 1.0f);
__fadd_rn(0.0f, 1.0f);
__fadd_ru(0.0f, 1.0f);
@@ -47,7 +48,9 @@ __device__ void single_precision_intrinsics() {
__fdiv_rn(4.0f, 2.0f);
__fdiv_ru(4.0f, 2.0f);
__fdiv_rz(4.0f, 2.0f);
#endif
__fdividef(4.0f, 2.0f);
#if defined OCML_BASIC_ROUNDED_OPERATIONS
__fmaf_rd(1.0f, 2.0f, 3.0f);
__fmaf_rn(1.0f, 2.0f, 3.0f);
__fmaf_ru(1.0f, 2.0f, 3.0f);
@@ -69,6 +72,7 @@ __device__ void single_precision_intrinsics() {
__fsub_rn(2.0f, 1.0f);
__fsub_ru(2.0f, 1.0f);
__fsub_rz(2.0f, 1.0f);
#endif
__log10f(1.0f);
__log2f(1.0f);
__logf(1.0f);