diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index b26ee22d62..2f65d75ae1 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -421,9 +421,9 @@ endif() ############################# # Package: hip_base set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_base) -configure_file(packaging/hip_base.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) -configure_file(packaging/hip_base.postinst ${BUILD_DIR}/postinst @ONLY) -configure_file(packaging/hip_base.prerm ${BUILD_DIR}/prerm @ONLY) +configure_file(packaging/hip-base.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip-base.postinst ${BUILD_DIR}/postinst @ONLY) +configure_file(packaging/hip-base.prerm ${BUILD_DIR}/prerm @ONLY) add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz @@ -441,9 +441,9 @@ endif() # Package: hip_hcc set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_hcc) -configure_file(packaging/hip_hcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) -configure_file(packaging/hip_hcc.postinst ${BUILD_DIR}/postinst @ONLY) -configure_file(packaging/hip_hcc.prerm ${BUILD_DIR}/prerm @ONLY) +configure_file(packaging/hip-hcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip-hcc.postinst ${BUILD_DIR}/postinst @ONLY) +configure_file(packaging/hip-hcc.prerm ${BUILD_DIR}/prerm @ONLY) add_custom_target(pkg_hip_hcc COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -455,7 +455,7 @@ add_custom_target(pkg_hip_hcc COMMAND ${CMAKE_COMMAND} . # Package: hip_nvcc set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_nvcc) -configure_file(packaging/hip_nvcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip-nvcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) add_custom_target(pkg_hip_nvcc COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -466,7 +466,7 @@ add_custom_target(pkg_hip_nvcc COMMAND ${CMAKE_COMMAND} . # Package: hip_doc set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_doc) -configure_file(packaging/hip_doc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip-doc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) add_custom_target(pkg_hip_doc COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -477,7 +477,7 @@ add_custom_target(pkg_hip_doc COMMAND ${CMAKE_COMMAND} . # Package: hip_samples set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_samples) -configure_file(packaging/hip_samples.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip-samples.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) add_custom_target(pkg_hip_samples COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package diff --git a/projects/hip/INSTALL.md b/projects/hip/INSTALL.md index 273c77a23c..2dc12f81a8 100644 --- a/projects/hip/INSTALL.md +++ b/projects/hip/INSTALL.md @@ -24,9 +24,9 @@ HIP code can be developed either on AMD ROCm platform using hcc or clang compile ## HIP-hcc * Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories). -* Install the "hip_hcc" package. This will install HCC and the HIP porting layer. +* Install the "hip-hcc" package. This will install HCC and the HIP porting layer. ``` -apt-get install hip_hcc +apt-get install hip-hcc ``` * Default paths and environment variables: @@ -67,9 +67,9 @@ apt-get install hip_hcc ## HIP-nvcc * Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories). -* Install the "hip_nvcc" package. This will install CUDA SDK and the HIP porting layer. +* Install the "hip-nvcc" package. This will install CUDA SDK and the HIP porting layer. ``` -apt-get install hip_nvcc +apt-get install hip-nvcc ``` * Default paths and environment variables: diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index b804d8d62d..d838678fdf 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -73,7 +73,7 @@ $no_output = 1 if $examine; push(@whitelist, split(',', $whitelist)); -@statNames = ("error", "init", "version", "device", "context", "module", "memory", "addressing", "stream", "event", "external_resource_interop", "stream_memory", "execution", "graph", "occupancy", "texture", "surface", "peer", "graphics", "profiler", "openGL", "D3D9", "D3D10", "D3D11", "VDPAU", "EGL", "thread", "complex", "library", "device_library", "device_function", "include", "include_cuda_main_header", "type", "literal", "numeric_literal", "define", "extern_shared", "kernel_launch"); +@statNames = ("error", "init", "version", "device", "context", "module", "memory", "virtual_memory", "addressing", "stream", "event", "external_resource_interop", "stream_memory", "execution", "graph", "occupancy", "texture", "surface", "peer", "graphics", "profiler", "openGL", "D3D9", "D3D10", "D3D11", "VDPAU", "EGL", "thread", "complex", "library", "device_library", "device_function", "include", "include_cuda_main_header", "type", "literal", "numeric_literal", "define", "extern_shared", "kernel_launch"); sub totalStats { my %count = %{ shift() }; @@ -1480,6 +1480,7 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bcudaErrorAssert\b/hipErrorAssert/g; $ft{'numeric_literal'} += s/\bcudaErrorCudartUnloading\b/hipErrorDeinitialized/g; $ft{'numeric_literal'} += s/\bcudaErrorDeviceAlreadyInUse\b/hipErrorContextAlreadyInUse/g; + $ft{'numeric_literal'} += s/\bcudaErrorDeviceUninitialized\b/hipErrorInvalidContext/g; $ft{'numeric_literal'} += s/\bcudaErrorDeviceUninitilialized\b/hipErrorInvalidContext/g; $ft{'numeric_literal'} += s/\bcudaErrorECCUncorrectable\b/hipErrorECCNotCorrectable/g; $ft{'numeric_literal'} += s/\bcudaErrorFileNotFound\b/hipErrorFileNotFound/g; @@ -2215,6 +2216,8 @@ sub warnUnsupportedDeviceFunctions { "__finite", "__finitef", "__finitel", + "__habs", + "__habs2", "__isinf", "__isinff", "__isinfl", @@ -2434,7 +2437,7 @@ while (@ARGV) { transformHostFunctions(); # TODO: would like to move this code outside loop but it uses $_ which contains the whole file unless ($no_output) { - my $apiCalls = $ft{'error'} + $ft{'init'} + $ft{'version'} + $ft{'device'} + $ft{'context'} + $ft{'module'} + $ft{'memory'} + $ft{'addressing'} + $ft{'stream'} + $ft{'event'} + $ft{'external_resource_interop'} + $ft{'stream_memory'} + $ft{'execution'} + $ft{'graph'} + $ft{'occupancy'} + $ft{'texture'} + $ft{'surface'} + $ft{'peer'} + $ft{'graphics'} + $ft{'profiler'} + $ft{'openGL'} + $ft{'D3D9'} + $ft{'D3D10'} + $ft{'D3D11'} + $ft{'VDPAU'} + $ft{'EGL'} + $ft{'thread'} + $ft{'complex'} + $ft{'library'} + $ft{'device_library'} + $ft{'include'} + $ft{'include_cuda_main_header'} + $ft{'type'} + $ft{'literal'} + $ft{'numeric_literal'} + $ft{'define'}; + my $apiCalls = $ft{'error'} + $ft{'init'} + $ft{'version'} + $ft{'device'} + $ft{'context'} + $ft{'module'} + $ft{'memory'} + $ft{'virtual_memory'} + $ft{'addressing'} + $ft{'stream'} + $ft{'event'} + $ft{'external_resource_interop'} + $ft{'stream_memory'} + $ft{'execution'} + $ft{'graph'} + $ft{'occupancy'} + $ft{'texture'} + $ft{'surface'} + $ft{'peer'} + $ft{'graphics'} + $ft{'profiler'} + $ft{'openGL'} + $ft{'D3D9'} + $ft{'D3D10'} + $ft{'D3D11'} + $ft{'VDPAU'} + $ft{'EGL'} + $ft{'thread'} + $ft{'complex'} + $ft{'library'} + $ft{'device_library'} + $ft{'include'} + $ft{'include_cuda_main_header'} + $ft{'type'} + $ft{'literal'} + $ft{'numeric_literal'} + $ft{'define'}; my $kernStuff = $hasDeviceCode + $ft{'kernel_launch'} + $ft{'device_function'}; my $totalCalls = $apiCalls + $kernStuff; $is_dos = m/\r\n$/; diff --git a/projects/hip/docker/dockerfile-hip-ubuntu-16.04 b/projects/hip/docker/dockerfile-hip-ubuntu-16.04 index 0852ae38c5..93dd17b948 100644 --- a/projects/hip/docker/dockerfile-hip-ubuntu-16.04 +++ b/projects/hip/docker/dockerfile-hip-ubuntu-16.04 @@ -10,10 +10,10 @@ COPY *.deb /tmp/ # Install the debian package RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install --no-install-recommends -y curl \ && apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install --no-install-recommends --allow-unauthenticated -y \ - /tmp/hip_base-*.deb \ - /tmp/hip_hcc-*.deb \ - /tmp/hip_doc-*.deb \ - /tmp/hip_samples-* \ + /tmp/hip-base-*.deb \ + /tmp/hip-hcc-*.deb \ + /tmp/hip-doc-*.deb \ + /tmp/hip-samples-* \ && rm -f /tmp/*.deb \ && apt-get clean \ && rm -rf /var/lib/apt/lists/* diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index c4fbd1db6d..ca61ddbb72 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -173,7 +173,11 @@ | 99 |*`CU_DEVICE_ATTRIBUTE_HOST_REGISTER_SUPPORTED`* | | 9.2 | | 100 |*`CU_DEVICE_ATTRIBUTE_PAGEABLE_MEMORY_ACCESS_USES_HOST_PAGE_TABLES`*| | 9.2 | | 101 |*`CU_DEVICE_ATTRIBUTE_DIRECT_MANAGED_MEM_ACCESS_FROM_HOST`* | | 9.2 | -| 102 |*`CU_DEVICE_ATTRIBUTE_MAX`* | | +| 102 |*`CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED`* | | 10.2 | +| 103 |*`CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR_SUPPORTED`* | | 10.2 | +| 104 |*`CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_WIN32_HANDLE_SUPPORTED`* | | 10.2 | +| 105 |*`CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_WIN32_KMT_HANDLE_SUPPORTED`* | | 10.2 | +| 106 |*`CU_DEVICE_ATTRIBUTE_MAX`* | | | enum |***`CUevent_flags`*** | | | typedef |***`CUevent_flags_enum`*** | | | 0x00 |*`CU_EVENT_DEFAULT`* |*`hipEventDefault`* | @@ -335,6 +339,11 @@ | 7 |*`CU_POINTER_ATTRIBUTE_BUFFER_ID`* | | | 8 |*`CU_POINTER_ATTRIBUTE_IS_MANAGED`* | | | 9 |*`CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL`* | | 9.2 | +| 10 |*`CU_POINTER_ATTRIBUTE_IS_LEGACY_CUDA_IPC_CAPABLE`* | | 10.2 | +| 11 |*`CU_POINTER_ATTRIBUTE_RANGE_START_ADDR`* | | 10.2 | +| 12 |*`CU_POINTER_ATTRIBUTE_RANGE_SIZE`* | | 10.2 | +| 13 |*`CU_POINTER_ATTRIBUTE_MAPPED`* | | 10.2 | +| 14 |*`CU_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES`* | | 10.2 | | enum |***`CUresourcetype`*** | | | typedef |***`CUresourcetype_enum`*** | | | 0x00 |*`CU_RESOURCE_TYPE_ARRAY`* | | @@ -452,6 +461,8 @@ | 906 |*`CUDA_ERROR_STREAM_CAPTURE_IMPLICIT`* | | 10.0 | | 907 |*`CUDA_ERROR_CAPTURED_EVENT`* | | 10.0 | | 908 |*`CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD`* | | 10.1 | +| 909 |*`CUDA_ERROR_TIMEOUT`* | | 10.2 | +| 910 |*`CUDA_ERROR_GRAPH_EXEC_UPDATE_FAILURE`* | | 10.2 | | 999 |*`CUDA_ERROR_UNKNOWN`* |*`hipErrorUnknown`* | | enum |***`CUsharedconfig`*** |***`hipSharedMemConfig`*** | | typedef |***`CUsharedconfig_enum`*** |***`hipSharedMemConfig`*** | @@ -716,7 +727,14 @@ | 3 |*`CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT`* | | 10.0 | | 4 |*`CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_HEAP`* | | 10.0 | | 5 |*`CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_RESOURCE`* | | 10.0 | +| 6 |*`CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE`* | | 10.2 | +| 7 |*`CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE_KMT`* | | 10.2 | +| 8 |*`CU_EXTERNAL_MEMORY_HANDLE_TYPE_NVSCIBUF`* | | 10.2 | | define |`CUDA_EXTERNAL_MEMORY_DEDICATED` | | 10.0 | +| define |`CUDA_EXTERNAL_SEMAPHORE_SIGNAL_SKIP_NVSCIBUF_MEMSYNC` | | 10.2 | +| define |`CUDA_EXTERNAL_SEMAPHORE_WAIT_SKIP_NVSCIBUF_MEMSYNC` | | 10.2 | +| define |`CUDA_NVSCISYNC_ATTR_SIGNAL` | | 10.2 | +| define |`CUDA_NVSCISYNC_ATTR_WAIT` | | 10.2 | | struct |`CUDA_EXTERNAL_MEMORY_HANDLE_DESC` | | 10.0 | | typedef |`CUDA_EXTERNAL_MEMORY_HANDLE_DESC_st` | | 10.0 | | struct |`CUDA_EXTERNAL_MEMORY_BUFFER_DESC` | | 10.0 | @@ -729,6 +747,10 @@ | 2 |*`CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32`* | | 10.0 | | 3 |*`CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT`* | | 10.0 | | 4 |*`CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D12_FENCE`* | | 10.0 | +| 5 |*`CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_FENCE`* | | 10.2 | +| 6 |*`CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC`* | | 10.2 | +| 7 |*`CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX`* | | 10.2 | +| 8 |*`CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX_KMT`* | | 10.2 | | struct |`CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC` | | 10.0 | | typedef |`CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC_st` | | 10.0 | | struct |`CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS` | | 10.0 | @@ -743,6 +765,42 @@ | define |`CU_DEVICE_INVALID` | | 8.0 | | struct |`CUuuid` | | | typedef |`CUuuid_st` | | +| enum |***`CUmemAllocationHandleType`*** | | 10.2 | +| typedef |***`CUmemAllocationHandleType_enum`*** | | 10.2 | +| 0x1 |*`CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR`* | | 10.2 | +| 0x2 |*`CU_MEM_HANDLE_TYPE_WIN32`* | | 10.2 | +| 0x4 |*`CU_MEM_HANDLE_TYPE_WIN32_KMT`* | | 10.2 | +| 0xFFFFFFFF |*`CU_MEM_HANDLE_TYPE_MAX`* | | 10.2 | +| enum |***`CUmemAccess_flags`*** | | 10.2 | +| typedef |***`CUmemAccess_flags_enum`*** | | 10.2 | +| 0x1 |*`CU_MEM_ACCESS_FLAGS_PROT_NONE`* | | 10.2 | +| 0x2 |*`CU_MEM_ACCESS_FLAGS_PROT_READ`* | | 10.2 | +| 0x3 |*`CU_MEM_ACCESS_FLAGS_PROT_READWRITE`* | | 10.2 | +| 0xFFFFFFFF |*`CU_MEM_ACCESS_FLAGS_PROT_MAX`* | | 10.2 | +| enum |***`CUmemLocationType`*** | | 10.2 | +| typedef |***`CUmemLocationType_enum`*** | | 10.2 | +| 0x0 |*`CU_MEM_LOCATION_TYPE_INVALID`* | | 10.2 | +| 0x1 |*`CU_MEM_LOCATION_TYPE_DEVICE`* | | 10.2 | +| 0xFFFFFFFF |*`CU_MEM_LOCATION_TYPE_MAX`* | | 10.2 | +| enum |***`CUmemAllocationGranularity_flags`*** | | 10.2 | +| typedef |***`CUmemAllocationGranularity_flags_enum`*** | | 10.2 | +| 0x0 |*`CU_MEM_ALLOC_GRANULARITY_MINIMUM`* | | 10.2 | +| 0x1 |*`CU_MEM_ALLOC_GRANULARITY_RECOMMENDED`* | | 10.2 | +| struct |`CUmemLocation` | | 10.2 | +| typedef |`CUmemLocation_st` | | 10.2 | +| struct |`CUmemAllocationProp` | | 10.2 | +| typedef |`CUmemAllocationProp_st` | | 10.2 | +| struct |`CUmemAccessDesc` | | 10.2 | +| typedef |`CUmemAccessDesc_st` | | 10.2 | +| enum |***`CUgraphExecUpdateResult`*** | | 10.2 | +| typedef |***`CUgraphExecUpdateResult_enum`*** | | 10.2 | +| 0x0 |*`CU_GRAPH_EXEC_UPDATE_SUCCESS`* | | 10.2 | +| 0x1 |*`CU_GRAPH_EXEC_UPDATE_ERROR`* | | 10.2 | +| 0x2 |*`CU_GRAPH_EXEC_UPDATE_ERROR_TOPOLOGY_CHANGED`* | | 10.2 | +| 0x3 |*`CU_GRAPH_EXEC_UPDATE_ERROR_NODE_TYPE_CHANGED`* | | 10.2 | +| 0x4 |*`CU_GRAPH_EXEC_UPDATE_ERROR_FUNCTION_CHANGED`* | | 10.2 | +| 0x5 |*`CU_GRAPH_EXEC_UPDATE_ERROR_PARAMETERS_CHANGED`* | | 10.2 | +| 0x6 |*`CU_GRAPH_EXEC_UPDATE_ERROR_NOT_SUPPORTED`* | | 10.2 | ## **2. Error Handling** @@ -771,6 +829,7 @@ | `cuDeviceGetAttribute` | `hipDeviceGetAttribute` | | `cuDeviceGetCount` | `hipGetDeviceCount` | | `cuDeviceGetName` | `hipDeviceGetName` | +| `cuDeviceGetNvSciSyncAttributes` | | 10.2 | | `cuDeviceTotalMem` | `hipDeviceTotalMem` | | `cuDeviceGetLuid` | | 10.0 | | `cuDeviceGetUuid` | | 9.2 | @@ -909,7 +968,25 @@ | `cuMipmappedArrayDestroy` | | | `cuMipmappedArrayGetLevel` | | -## **12. Unified Addressing** +## **12. Virtual Memory Management** + +| **CUDA** | **HIP** |**CUDA version\***| +|-----------------------------------------------------------|-------------------------------|:----------------:| +| `cuMemAddressFree` | | 10.2 | +| `cuMemAddressReserve` | | 10.2 | +| `cuMemCreate` | | 10.2 | +| `cuMemExportToShareableHandle` | | 10.2 | + +| `cuMemGetAccess` | | 10.2 | +| `cuMemGetAllocationGranularity` | | 10.2 | +| `cuMemGetAllocationPropertiesFromHandle` | | 10.2 | +| `cuMemImportFromShareableHandle` | | 10.2 | +| `cuMemMap` | | 10.2 | +| `cuMemRelease` | | 10.2 | +| `cuMemSetAccess` | | 10.2 | +| `cuMemUnmap` | | 10.2 | + +## **13. Unified Addressing** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| @@ -921,7 +998,7 @@ | `cuPointerGetAttributes` | | | `cuPointerSetAttribute` | | -## **13. Stream Management** +## **14. Stream Management** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| @@ -942,7 +1019,7 @@ | `cuStreamIsCapturing` | | 10.0 | | `cuThreadExchangeStreamCaptureMode` | | 10.1 | -## **14. Event Management** +## **15. Event Management** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|------------------| @@ -953,7 +1030,7 @@ | `cuEventRecord` | `hipEventRecord` | | `cuEventSynchronize` | `hipEventSynchronize` | -## **15. External Resource Interoperability** +## **16. External Resource Interoperability** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| @@ -966,7 +1043,7 @@ | `cuImportExternalSemaphore` | | 10.0 | | `cuDestroyExternalSemaphore` | | 10.0 | -## **16. Stream Memory Operations** +## **17. Stream Memory Operations** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| @@ -976,7 +1053,7 @@ | `cuStreamWriteValue32` | | 8.0 | | `cuStreamWriteValue64` | | 9.0 | -## **17. Execution Control** +## **18. Execution Control** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| @@ -989,7 +1066,7 @@ | `cuLaunchCooperativeKernel` | | 9.0 | | `cuLaunchCooperativeKernelMultiDevice` | | 9.0 | -## **18. Execution Control [DEPRECATED]** +## **19. Execution Control [DEPRECATED]** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|------------------| @@ -1003,7 +1080,7 @@ | `cuParamSetTexRef` | | | `cuParamSetv` | | -## **19. Graph Management** +## **20. Graph Management** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| @@ -1038,9 +1115,13 @@ | `cuGraphInstantiate` | | 10.0 | | `cuGraphExecDestroy` | | 10.0 | | `cuGraphExecKernelNodeSetParams` | | 10.1 | +| `cuGraphExecMemcpyNodeSetParams` | | 10.2 | +| `cuGraphExecMemsetNodeSetParams` | | 10.2 | +| `cuGraphExecHostNodeSetParams` | | 10.2 | +| `cuGraphExecUpdate` | | 10.2 | | `cuGraphDestroy` | | 10.0 | -## **20. Occupancy** +## **21. Occupancy** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|---------------------------------------------------------|------------------| @@ -1049,7 +1130,7 @@ | `cuOccupancyMaxPotentialBlockSize` |`hipOccupancyMaxPotentialBlockSize` | | `cuOccupancyMaxPotentialBlockSizeWithFlags` | | -## **21. Texture Reference Management** +## **22. Texture Reference Management [DEPRECATED]** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| @@ -1078,15 +1159,10 @@ | `cuTexRefSetMipmapLevelBias` | | | `cuTexRefSetMipmapLevelClamp` | | | `cuTexRefSetMipmappedArray` | | - -## **22. Texture Reference Management [DEPRECATED]** - -| **CUDA** | **HIP** |**CUDA version\***| -|-----------------------------------------------------------|-------------------------------|------------------| | `cuTexRefCreate` | | | `cuTexRefDestroy` | | -## **23. Surface Reference Management** +## **23. Surface Reference Management [DEPRECATED]** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|------------------| diff --git a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index 4fa033e21c..49d092c218 100644 --- a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -9,6 +9,7 @@ | `cudaDeviceGetByPCIBusId` | `hipDeviceGetByPCIBusId` | | `cudaDeviceGetCacheConfig` | `hipDeviceGetCacheConfig` | | `cudaDeviceGetLimit` | `hipDeviceGetLimit` | +| `cudaDeviceGetNvSciSyncAttributes` | | 10.2 | | `cudaDeviceGetPCIBusId` | `hipDeviceGetPCIBusId` | | `cudaDeviceGetSharedMemConfig` | `hipDeviceGetSharedMemConfig` | | `cudaDeviceGetStreamPriorityRange` | `hipDeviceGetStreamPriorityRange` | @@ -121,15 +122,15 @@ | `cudaOccupancyMaxActiveBlocksPerMultiprocessor` |`hipOccupancyMaxActiveBlocksPerMultiprocessor` | | `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` |`hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags`| -## **9. Execution Control [DEPRECATED since 7.0]** +## **Former 9. Execution Control [DEPRECATED since 7.0, REMOVED since 10.1]** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| | `cudaConfigureCall` | `hipConfigureCall` | | `cudaLaunch` | `hipLaunchByPtr` | -| `cudaSetupArgument` | `hipSetupArgument` | 7.0 - 10.0 | +| `cudaSetupArgument` | `hipSetupArgument` | -## **10. Memory Management** +## **9. Memory Management** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| @@ -168,16 +169,11 @@ | `cudaMemcpy3DAsync` | | | `cudaMemcpy3DPeer` | | | `cudaMemcpy3DPeerAsync` | | -| `cudaMemcpyArrayToArray` | | | `cudaMemcpyAsync` | `hipMemcpyAsync` | -| `cudaMemcpyFromArray` | `hipMemcpyFromArray` | -| `cudaMemcpyFromArrayAsync` | | | `cudaMemcpyFromSymbol` | `hipMemcpyFromSymbol` | | `cudaMemcpyFromSymbolAsync` | `hipMemcpyFromSymbolAsync` | | `cudaMemcpyPeer` | `hipMemcpyPeer` | | `cudaMemcpyPeerAsync` | `hipMemcpyPeerAsync` | -| `cudaMemcpyToArray` | `hipMemcpyToArray` | -| `cudaMemcpyToArrayAsync` | | | `cudaMemcpyToSymbol` | `hipMemcpyToSymbol` | | `cudaMemcpyToSymbolAsync` | `hipMemcpyToSymbolAsync` | | `cudaMemset` | `hipMemset` | @@ -190,6 +186,16 @@ | `make_cudaPitchedPtr` | `make_hipPitchedPtr` | | `make_cudaPos` | `make_hipPos` | +## **10. Memory Management [DEPRECATED since 10.1]** + +| **CUDA** | **HIP** |**CUDA version\***| +|-----------------------------------------------------------|-------------------------------|:----------------:| +| `cudaMemcpyArrayToArray` | | +| `cudaMemcpyFromArray` | `hipMemcpyFromArray` | +| `cudaMemcpyFromArrayAsync` | | +| `cudaMemcpyToArray` | `hipMemcpyToArray` | +| `cudaMemcpyToArrayAsync` | | + ## **11. Unified Addressing** | **CUDA** | **HIP** |**CUDA version\***| @@ -329,7 +335,7 @@ | `cudaGraphicsUnmapResources` | | | `cudaGraphicsUnregisterResource` | | -## **24. Texture Reference Management** +## **24. Texture Reference Management [DEPRECATED]** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|----------------------------------|:----------------:| @@ -343,7 +349,7 @@ | `cudaGetTextureReference` | `hipGetTextureReference` | | `cudaUnbindTexture` | `hipUnbindTexture` | -## **25. Surface Reference Management** +## **25. Surface Reference Management [DEPRECATED]** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------|:----------------:| @@ -399,6 +405,10 @@ | `cudaGraphHostNodeSetParams` | | 10.0 | | `cudaGraphInstantiate` | | 10.0 | | `cudaGraphExecKernelNodeSetParams` | | 10.1 | +| `cudaGraphExecMemcpyNodeSetParams` | | 10.2 | +| `cudaGraphExecMemsetNodeSetParams` | | 10.2 | +| `cudaGraphExecHostNodeSetParams` | | 10.2 | +| `cudaGraphExecUpdate` | | 10.2 | | `cudaGraphKernelNodeGetParams` | | 10.0 | | `cudaGraphKernelNodeSetParams` | | 10.0 | | `cudaGraphLaunch` | | 10.0 | @@ -412,8 +422,7 @@ | `cudaGraphNodeGetType` | | 10.0 | | `cudaGraphRemoveDependencies` | | 10.0 | -## **30. C++ API Routines** -*(7.0 contains, 7.5 doesn't)* +## **30. C++ API Routines [DEPRECATED since 7.5]** | **CUDA** | **HIP** |**CUDA version\***| |-----------------------------------------------------------|-------------------------------------------------------|:----------------:| @@ -685,11 +694,18 @@ | 3 |*`cudaExternalMemoryHandleTypeOpaqueWin32Kmt`* | 10.0 | | | 4 |*`cudaExternalMemoryHandleTypeD3D12Heap`* | 10.0 | | | 5 |*`cudaExternalMemoryHandleTypeD3D12Resource`* | 10.0 | | +| 6 |*`cudaExternalMemoryHandleTypeD3D11Resource`* | 10.2 | | +| 7 |*`cudaExternalMemoryHandleTypeD3D11ResourceKmt`* | 10.2 | | +| 8 |*`cudaExternalMemoryHandleTypeNvSciBuf`* | 10.2 | | | enum |***`cudaExternalSemaphoreHandleType`*** | 10.0 | | | 1 |*`cudaExternalSemaphoreHandleTypeOpaqueFd`* | 10.0 | | | 2 |*`cudaExternalSemaphoreHandleTypeOpaqueWin32`* | 10.0 | | | 3 |*`cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt`* | 10.0 | | | 4 |*`cudaExternalSemaphoreHandleTypeD3D12Fence`* | 10.0 | | +| 5 |*`cudaExternalSemaphoreHandleTypeD3D11Fence`* | 10.2 | | +| 6 |*`cudaExternalSemaphoreHandleTypeNvSciSync`* | 10.2 | | +| 7 |*`cudaExternalSemaphoreHandleTypeKeyedMutex`* | 10.2 | | +| 8 |*`cudaExternalSemaphoreHandleTypeKeyedMutexKmt`* | 10.2 | | | enum |***`cudaFuncAttribute`*** | 9.0 | | | 8 |*`cudaFuncAttributeMaxDynamicSharedMemorySize`* | 9.0 | | | 9 |*`cudaFuncAttributePreferredSharedMemoryCarveout`* | 9.0 | | @@ -799,6 +815,8 @@ | 906 |*`cudaErrorStreamCaptureImplicit`* | 10.0 | | | 907 |*`cudaErrorCapturedEvent`* | 10.0 | | | 908 |*`cudaErrorStreamCaptureWrongThread`* | 10.1 | | +| 909 |*`cudaErrorTimeout`* | 10.2 | | +| 910 |*`cudaErrorGraphExecUpdateFailure`* | 10.2 | | | 999 |*`cudaErrorUnknown`* | |*`hipErrorUnknown`* | 1030 | | 10000 |*`cudaErrorApiFailureBase`* | | | | enum |***`cudaFuncCache`*** | |***`hipFuncCache_t`*** | @@ -1069,6 +1087,10 @@ | struct |`CUeglStreamConnection_st` | 9.1 | | | typedef |`cudaEglStreamConnection` | 9.1 | | | define |`cudaExternalMemoryDedicated` | 10.0 | | +| define |`cudaExternalSemaphoreSignalSkipNvSciBufMemSync` | 10.2 | | +| define |`cudaExternalSemaphoreWaitSkipNvSciBufMemSync` | 10.2 | | +| define |`cudaNvSciSyncAttrSignal` | 10.2 | | +| define |`cudaNvSciSyncAttrWait` | 10.2 | | | typedef |`cudaExternalMemory_t` | 10.0 | | | struct |`CUexternalMemory_st` | 10.0 | | | typedef |`cudaExternalSemaphore_t` | 10.0 | | @@ -1087,5 +1109,13 @@ | 0 |*`MAJOR_VERSION`* | 8.0 | | | 1 |*`MINOR_VERSION`* | 8.0 | | | 2 |*`PATCH_LEVEL`* | 8.0 | | +| enum |***`cudaGraphExecUpdateResult`*** | 10.2 | | +| 0x0 |*`cudaGraphExecUpdateSuccess`* | 10.2 | | +| 0x1 |*`cudaGraphExecUpdateError`* | 10.2 | | +| 0x2 |*`cudaGraphExecUpdateErrorTopologyChanged`* | 10.2 | | +| 0x3 |*`cudaGraphExecUpdateErrorNodeTypeChanged`* | 10.2 | | +| 0x4 |*`cudaGraphExecUpdateErrorFunctionChanged`* | 10.2 | | +| 0x5 |*`cudaGraphExecUpdateErrorParametersChanged`* | 10.2 | | +| 0x6 |*`cudaGraphExecUpdateErrorNotSupported`* | 10.2 | | \* CUDA version, in which API has appeared and (optional) last version before abandoning it; no value in case of earlier versions < 7.5. diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp index 5b4f560e15..87fe2ac86c 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp @@ -548,6 +548,8 @@ const std::map CUDA_DEVICE_FUNC_MAP{ {"__hfma", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hfma_sat", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hneg", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"__habs2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, + {"__habs", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, {"__hbeq2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hbne2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"__hble2", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index 3c00660cae..0445d5cb31 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp @@ -52,6 +52,8 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ {"cuDeviceGetLuid", {"hipDeviceGetLuid", "", CONV_DEVICE, API_DRIVER, HIP_UNSUPPORTED}}, // no analogue {"cuDeviceGetName", {"hipDeviceGetName", "", CONV_DEVICE, API_DRIVER}}, + // cudaDeviceGetNvSciSyncAttributes + {"cuDeviceGetNvSciSyncAttributes", {"hipDeviceGetNvSciSyncAttributes", "", CONV_DEVICE, API_DRIVER, HIP_UNSUPPORTED}}, // no analogue {"cuDeviceGetUuid", {"hipDeviceGetUuid", "", CONV_DEVICE, API_DRIVER, HIP_UNSUPPORTED}}, // no analogue @@ -310,7 +312,22 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // NOTE: Not equal to cudaGetMipmappedArrayLevel due to different signatures {"cuMipmappedArrayGetLevel", {"hipMipmappedArrayGetLevel", "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, - // 5.12. Unified Addressing + // 5.12. Virtual Memory Management + // no analogue + {"cuMemAddressFree", {"hipMemAddressFree", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemAddressReserve", {"hipMemAddressReserve", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemCreate", {"hipMemCreate", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemExportToShareableHandle", {"hipMemExportToShareableHandle", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemGetAccess", {"hipMemGetAccess", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemGetAllocationGranularity", {"hipMemGetAllocationGranularity", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemGetAllocationPropertiesFromHandle", {"hipMemGetAllocationPropertiesFromHandle", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemImportFromShareableHandle", {"hipMemImportFromShareableHandle", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemMap", {"hipMemMap", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemRelease", {"hipMemRelease", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemSetAccess", {"hipMemSetAccess", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + {"cuMemUnmap", {"hipMemUnmap", "", CONV_VIRTUAL_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, + + // 5.13. Unified Addressing // cudaMemAdvise {"cuMemAdvise", {"hipMemAdvise", "", CONV_ADDRESSING, API_DRIVER, HIP_UNSUPPORTED}}, // TODO: double check cudaMemPrefetchAsync @@ -327,7 +344,7 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // no analogue {"cuPointerSetAttribute", {"hipPointerSetAttribute", "", CONV_ADDRESSING, API_DRIVER, HIP_UNSUPPORTED}}, - // 5.13. Stream Management + // 5.14. Stream Management // cudaStreamAddCallback {"cuStreamAddCallback", {"hipStreamAddCallback", "", CONV_STREAM, API_DRIVER}}, // cudaStreamAttachMemAsync @@ -364,7 +381,7 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // cudaThreadExchangeStreamCaptureMode {"cuThreadExchangeStreamCaptureMode", {"hipThreadExchangeStreamCaptureMode", "", CONV_STREAM, API_DRIVER, HIP_UNSUPPORTED}}, - // 5.14. Event Management + // 5.15. Event Management // cudaEventCreateWithFlags {"cuEventCreate", {"hipEventCreateWithFlags", "", CONV_EVENT, API_DRIVER}}, // cudaEventDestroy @@ -379,7 +396,7 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // cudaEventSynchronize {"cuEventSynchronize", {"hipEventSynchronize", "", CONV_EVENT, API_DRIVER}}, - // 5.15. External Resource Interoperability + // 5.16. External Resource Interoperability // cudaDestroyExternalMemory {"cuDestroyExternalMemory", {"hipDestroyExternalMemory", "", CONV_EXT_RES, API_DRIVER, HIP_UNSUPPORTED}}, // cudaDestroyExternalSemaphore @@ -397,7 +414,7 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // cudaWaitExternalSemaphoresAsync {"cuWaitExternalSemaphoresAsync", {"hipWaitExternalSemaphoresAsync", "", CONV_EXT_RES, API_DRIVER, HIP_UNSUPPORTED}}, - // 5.16. Stream Memory Operations + // 5.17. Stream Memory Operations // no analogues {"cuStreamBatchMemOp", {"hipStreamBatchMemOp", "", CONV_STREAM_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, {"cuStreamWaitValue32", {"hipStreamWaitValue32", "", CONV_STREAM_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, @@ -405,7 +422,7 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ {"cuStreamWriteValue32", {"hipStreamWriteValue32", "", CONV_STREAM_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, {"cuStreamWriteValue64", {"hipStreamWriteValue64", "", CONV_STREAM_MEMORY, API_DRIVER, HIP_UNSUPPORTED}}, - // 5.17.Execution Control + // 5.18.Execution Control // no analogue {"cuFuncGetAttribute", {"hipFuncGetAttribute", "", CONV_EXECUTION, API_DRIVER}}, // no analogue @@ -429,7 +446,7 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // NOTE: Not equal to cudaLaunchKernel due to different signatures {"cuLaunchKernel", {"hipModuleLaunchKernel", "", CONV_EXECUTION, API_DRIVER}}, - // 5.18.Execution Control [DEPRECATED] + // 5.19.Execution Control [DEPRECATED] // no analogue {"cuFuncSetBlockShape", {"hipFuncSetBlockShape", "", CONV_EXECUTION, API_DRIVER, HIP_UNSUPPORTED}}, // no analogue @@ -452,7 +469,7 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // no analogue {"cuParamSetv", {"hipParamSetv", "", CONV_EXECUTION, API_DRIVER, HIP_UNSUPPORTED}}, - // 5.19. Graph Management + // 5.20. Graph Management // cudaGraphAddChildGraphNode {"cuGraphAddChildGraphNode", {"hipGraphAddChildGraphNode", "", CONV_GRAPH, API_DRIVER, HIP_UNSUPPORTED}}, // cudaGraphAddDependencies @@ -517,8 +534,16 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ {"cuGraphNodeGetType", {"hipGraphNodeGetType", "", CONV_GRAPH, API_DRIVER, HIP_UNSUPPORTED}}, // cudaGraphRemoveDependencies {"cuGraphRemoveDependencies", {"hipGraphRemoveDependencies", "", CONV_GRAPH, API_DRIVER, HIP_UNSUPPORTED}}, + // cudaGraphExecMemcpyNodeSetParams + {"cuGraphExecMemcpyNodeSetParams", {"hipGraphExecMemcpyNodeSetParams", "", CONV_GRAPH, API_DRIVER, HIP_UNSUPPORTED}}, + // cudaGraphExecMemsetNodeSetParams + {"cuGraphExecMemsetNodeSetParams", {"hipGraphExecMemsetNodeSetParams", "", CONV_GRAPH, API_DRIVER, HIP_UNSUPPORTED}}, + // cudaGraphExecHostNodeSetParams + {"cuGraphExecHostNodeSetParams", {"hipGraphExecHostNodeSetParams", "", CONV_GRAPH, API_DRIVER, HIP_UNSUPPORTED}}, + // cudaGraphExecUpdate + {"cuGraphExecUpdate", {"hipGraphExecUpdate", "", CONV_GRAPH, API_DRIVER, HIP_UNSUPPORTED}}, - // 5.20. Occupancy + // 5.21. Occupancy // cudaOccupancyMaxActiveBlocksPerMultiprocessor {"cuOccupancyMaxActiveBlocksPerMultiprocessor", {"hipOccupancyMaxActiveBlocksPerMultiprocessor", "", CONV_OCCUPANCY, API_DRIVER}}, // cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags @@ -528,7 +553,7 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ // cudaOccupancyMaxPotentialBlockSizeWithFlags {"cuOccupancyMaxPotentialBlockSizeWithFlags", {"hipOccupancyMaxPotentialBlockSizeWithFlags", "", CONV_OCCUPANCY, API_DRIVER, HIP_UNSUPPORTED}}, - // 5.21. Texture Reference Management + // 5.22. Texture Reference Management [DEPRECATED] // no analogues {"cuTexRefGetAddress", {"hipTexRefGetAddress", "", CONV_TEXTURE, API_DRIVER}}, {"cuTexRefGetAddress_v2", {"hipTexRefGetAddress", "", CONV_TEXTURE, API_DRIVER}}, @@ -559,13 +584,10 @@ const std::map CUDA_DRIVER_FUNCTION_MAP{ {"cuTexRefSetMipmapLevelBias", {"hipTexRefSetMipmapLevelBias", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, {"cuTexRefSetMipmapLevelClamp", {"hipTexRefSetMipmapLevelClamp", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, {"cuTexRefSetMipmappedArray", {"hipTexRefSetMipmappedArray", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, - - // 5.22. Texture Reference Management [DEPRECATED] - // no analogues {"cuTexRefCreate", {"hipTexRefCreate", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, {"cuTexRefDestroy", {"hipTexRefDestroy", "", CONV_TEXTURE, API_DRIVER, HIP_UNSUPPORTED}}, - // 5.23. Surface Reference Management + // 5.23. Surface Reference Management [DEPRECATED] // no analogues {"cuSurfRefGetArray", {"hipSurfRefGetArray", "", CONV_SURFACE, API_DRIVER, HIP_UNSUPPORTED}}, {"cuSurfRefSetArray", {"hipSurfRefSetArray", "", CONV_SURFACE, API_DRIVER, HIP_UNSUPPORTED}}, diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp index 4603af69d0..f43f0c4e4c 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -198,6 +198,18 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CUuuid_st", {"hipUUID", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, {"CUuuid", {"hipUUID", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + // no analogue + {"CUmemLocation_st", {"hipMemoryLocation", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUmemLocation", {"hipMemoryLocation", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + + // no analogue + {"CUmemAllocationProp_st", {"hipMemoryAllocationProperties", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUmemAllocationProp", {"hipMemoryAllocationProperties", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + + // no analogue + {"CUmemAccessDesc_st", {"hipMemoryAccessDescription", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUmemAccessDesc", {"hipMemoryAccessDescription", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + // 2. Unions {"CUstreamBatchMemOpParams", {"hipStreamBatchMemOpParams", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, @@ -500,7 +512,15 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // cudaDevAttrDirectManagedMemAccessFromHost {"CU_DEVICE_ATTRIBUTE_DIRECT_MANAGED_MEM_ACCESS_FROM_HOST", {"hipDeviceAttributeDirectManagedMemAccessFromHost", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 101 // no analogue - {"CU_DEVICE_ATTRIBUTE_MAX", {"hipDeviceAttributeMax", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 102 + {"CU_DEVICE_ATTRIBUTE_VIRTUAL_ADDRESS_MANAGEMENT_SUPPORTED", {"hipDeviceAttributeVirtualAddressManagementSupported", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 102 + // no analogue + {"CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR_SUPPORTED", {"hipDeviceAttributeHandleTypePosixFileDescriptorSupported", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 103 + // no analogue + {"CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_WIN32_HANDLE_SUPPORTED", {"hipDeviceAttributeHandleTypeWin32HandleSupported", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 104 + // no analogue + {"CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_WIN32_KMT_HANDLE_SUPPORTED", {"hipDeviceAttributeHandleTypeWin32KmtHandleSupported", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 105 + // no analogue + {"CU_DEVICE_ATTRIBUTE_MAX", {"hipDeviceAttributeMax", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 106 // cudaDeviceP2PAttr {"CUdevice_P2PAttribute", {"hipDeviceP2PAttribute", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, @@ -717,6 +737,12 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_HEAP", {"hipExternalMemoryHandleTypeD3D12Heap", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 4 // cudaExternalMemoryHandleTypeD3D12Resource {"CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_RESOURCE", {"hipExternalMemoryHandleTypeD3D12Resource", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 5 + // cudaExternalMemoryHandleTypeD3D11Resource + {"CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE", {"hipExternalMemoryHandleTypeD3D11Resource", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 6 + // cudaExternalMemoryHandleTypeD3D11ResourceKmt + {"CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE_KMT", {"hipExternalMemoryHandleTypeD3D11ResourceKmt", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 7 + // cudaExternalMemoryHandleTypeNvSciBuf + {"CU_EXTERNAL_MEMORY_HANDLE_TYPE_NVSCIBUF", {"hipExternalMemoryHandleTypeNvSciBuf", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 8 // cudaExternalSemaphoreHandleType {"CUexternalSemaphoreHandleType", {"hipExternalSemaphoreHandleType", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, @@ -731,6 +757,15 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // cudaExternalSemaphoreHandleTypeD3D12Fence {"CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D12_FENCE", {"hipExternalSemaphoreHandleTypeD3D12Fence", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 4 + // cudaExternalSemaphoreHandleTypeD3D11Fence + {"CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_FENCE", {"hipExternalSemaphoreHandleTypeD3D11Fence", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 5 + // cudaExternalSemaphoreHandleTypeNvSciSync + {"CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC", {"hipExternalSemaphoreHandleTypeNvSciSync", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 6 + // cudaExternalSemaphoreHandleTypeKeyedMutex + {"CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX", {"hipExternalSemaphoreHandleTypeKeyedMutex", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 7 + // cudaExternalSemaphoreHandleTypeKeyedMutexKmt + {"CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX_KMT", {"hipExternalSemaphoreHandleTypeKeyedMutexKmt", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 8 + // cudaTextureFilterMode {"CUfilter_mode", {"hipTextureFilterMode", "", CONV_TYPE, API_DRIVER}}, {"CUfilter_mode_enum", {"hipTextureFilterMode", "", CONV_TYPE, API_DRIVER}}, @@ -825,6 +860,25 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // cudaGraphNodeTypeCount {"CU_GRAPH_NODE_TYPE_COUNT", {"hipGraphNodeTypeCount", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 6 + // cudaGraphExecUpdateResult + {"CUgraphExecUpdateResult", {"hipGraphExecUpdateResult", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUgraphExecUpdateResult_enum", {"hipGraphExecUpdateResult", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + // CUgraphExecUpdateResult enum values + // cudaGraphExecUpdateSuccess + {"CU_GRAPH_EXEC_UPDATE_SUCCESS", {"hipGraphExecUpdateSuccess", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x0 + // cudaGraphExecUpdateError + {"CU_GRAPH_EXEC_UPDATE_ERROR", {"hipGraphExecUpdateError", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 + // cudaGraphExecUpdateErrorTopologyChanged + {"CU_GRAPH_EXEC_UPDATE_ERROR_TOPOLOGY_CHANGED", {"hipGraphExecUpdateErrorTopologyChanged", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x2 + // cudaGraphExecUpdateErrorNodeTypeChanged + {"CU_GRAPH_EXEC_UPDATE_ERROR_NODE_TYPE_CHANGED", {"hipGraphExecUpdateErrorNodeTypeChanged", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x3 + // cudaGraphExecUpdateErrorFunctionChanged + {"CU_GRAPH_EXEC_UPDATE_ERROR_FUNCTION_CHANGED", {"hipGraphExecUpdateErrorFunctionChanged", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x4 + // cudaGraphExecUpdateErrorParametersChanged + {"CU_GRAPH_EXEC_UPDATE_ERROR_PARAMETERS_CHANGED", {"hipGraphExecUpdateErrorParametersChanged", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x5 + // cudaGraphExecUpdateErrorNotSupported + {"CU_GRAPH_EXEC_UPDATE_ERROR_NOT_SUPPORTED", {"hipGraphExecUpdateErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x6 + // no analogue {"CUipcMem_flags", {"hipIpcMemFlags", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, {"CUipcMem_flags_enum", {"hipIpcMemFlags", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, @@ -1012,6 +1066,11 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CU_POINTER_ATTRIBUTE_BUFFER_ID", {"hipPointerAttributeBufferId", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 7 {"CU_POINTER_ATTRIBUTE_IS_MANAGED", {"hipPointerAttributeIsManaged", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 8 {"CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL", {"hipPointerAttributeDeviceOrdinal", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 9 + {"CU_POINTER_ATTRIBUTE_IS_LEGACY_CUDA_IPC_CAPABLE", {"hipPointerAttributeIsLegacyCudaIpcCapable", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 10 + {"CU_POINTER_ATTRIBUTE_RANGE_START_ADDR", {"hipPointerAttributeRangeStartAddress", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 11 + {"CU_POINTER_ATTRIBUTE_RANGE_SIZE", {"hipPointerAttributeRangeSize", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 12 + {"CU_POINTER_ATTRIBUTE_MAPPED", {"hipPointerAttributeMapped", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 13 + {"CU_POINTER_ATTRIBUTE_ALLOWED_HANDLE_TYPES", {"hipPointerAttributeAllowedHandleTypes", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 14 // cudaResourceType {"CUresourcetype", {"hipResourceType", "", CONV_TYPE, API_DRIVER}}, @@ -1291,6 +1350,10 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CUDA_ERROR_CAPTURED_EVENT", {"hipErrorCapturedEvent", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 907 // cudaErrorStreamCaptureWrongThread {"CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD", {"hipErrorStreamCaptureWrongThread", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 908 + // cudaErrorTimeout + {"CUDA_ERROR_TIMEOUT", {"hipErrorTimeout", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 909 + // cudaErrorGraphExecUpdateFailure + {"CUDA_ERROR_GRAPH_EXEC_UPDATE_FAILURE", {"hipErrorGraphExecUpdateFailure", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 910 // cudaErrorUnknown = 999, CUDA_ERROR_UNKNOWN = 999, hipErrorUnknown = 1030 // TODO [HIP]: make hipErrorUnknown = 999 // cudaErrorUnknown @@ -1471,6 +1534,39 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // cudaD3D11DeviceListNextFrame = 3 {"CU_D3D11_DEVICE_LIST_NEXT_FRAME", {"HIP_D3D11_DEVICE_LIST_NEXT_FRAME", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x03 + // no analogue + {"CUmemAllocationHandleType", {"hipMemoryAllocationHandleType", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUmemAllocationHandleType_enum", {"hipMemoryAllocationHandleType", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + // CUmemAllocationHandleType enum values + {"CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR", {"HIP_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 + {"CU_MEM_HANDLE_TYPE_WIN32", {"HIP_MEM_HANDLE_TYPE_WIN32", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x2 + {"CU_MEM_HANDLE_TYPE_WIN32_KMT", {"HIP_MEM_HANDLE_TYPE_WIN32_KMT", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x4 + {"CU_MEM_HANDLE_TYPE_MAX", {"HIP_MEM_HANDLE_TYPE_MAX", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0xFFFFFFFF + + // no analogue + {"CUmemAccess_flags", {"hipMemoryAccessFlags", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUmemAccess_flags_enum", {"hipMemoryAccessFlags", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + // CUmemAccess_flags enum values + {"CU_MEM_ACCESS_FLAGS_PROT_NONE", {"HIP_MEM_ACCESS_FLAGS_PROT_NONE", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 + {"CU_MEM_ACCESS_FLAGS_PROT_READ", {"HIP_MEM_ACCESS_FLAGS_PROT_READ", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x2 + {"CU_MEM_ACCESS_FLAGS_PROT_READWRITE", {"HIP_MEM_ACCESS_FLAGS_PROT_READWRITE", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x3 + {"CU_MEM_ACCESS_FLAGS_PROT_MAX", {"HIP_MEM_ACCESS_FLAGS_PROT_MAX", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0xFFFFFFFF + + // no analogue + {"CUmemLocationType", {"hipMemoryLocationType", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUmemLocationType_enum", {"hipMemoryLocationType", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + // CUmemLocationType enum values + {"CU_MEM_LOCATION_TYPE_INVALID", {"HIP_MEM_LOCATION_TYPE_INVALID", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x0 + {"CU_MEM_LOCATION_TYPE_DEVICE", {"HIP_MEM_LOCATION_TYPE_DEVICE", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 + {"CU_MEM_LOCATION_TYPE_MAX", {"HIP_MEM_LOCATION_TYPE_MAX", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0xFFFFFFFF + + // no analogue + {"CUmemAllocationGranularity_flags", {"hipMemoryAllocationGranularityFlags", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUmemAllocationGranularity_flags_enum", {"hipMemoryLocationType", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + // CUmemAllocationGranularity_flags enum values + {"CU_MEM_ALLOC_GRANULARITY_MINIMUM", {"HIP_MEM_ALLOC_GRANULARITY_MINIMUM", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x0 + {"CU_MEM_ALLOC_GRANULARITY_RECOMMENDED", {"HIP_MEM_ALLOC_GRANULARITY_RECOMMENDED", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 + // 4. Typedefs // no analogue @@ -1546,6 +1642,14 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // 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_EXTERNAL_MEMORY_DEDICATED", {"hipExternalMemoryDedicated", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 + // cudaExternalSemaphoreSignalSkipNvSciBufMemSync + {"CUDA_EXTERNAL_SEMAPHORE_SIGNAL_SKIP_NVSCIBUF_MEMSYNC", {"hipExternalSemaphoreSignalSkipNvSciBufMemSync", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01 + // cudaExternalSemaphoreWaitSkipNvSciBufMemSync + {"CUDA_EXTERNAL_SEMAPHORE_WAIT_SKIP_NVSCIBUF_MEMSYNC", {"hipExternalSemaphoreWaitSkipNvSciBufMemSync", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02 + // cudaNvSciSyncAttrSignal + {"CUDA_NVSCISYNC_ATTR_SIGNAL", {"hipNvSciSyncAttrSignal", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 + // cudaNvSciSyncAttrWait + {"CUDA_NVSCISYNC_ATTR_WAIT", {"hipNvSciSyncAttrWait", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x1 {"CUDA_VERSION", {"HIP_VERSION", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 10000 }; diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp index 831f32267d..fbb5f22ae4 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp @@ -35,6 +35,8 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ {"cudaDeviceGetCacheConfig", {"hipDeviceGetCacheConfig", "", CONV_DEVICE, API_RUNTIME}}, // cuCtxGetLimit {"cudaDeviceGetLimit", {"hipDeviceGetLimit", "", CONV_DEVICE, API_RUNTIME}}, + // cuDeviceGetNvSciSyncAttributes + {"cudaDeviceGetNvSciSyncAttributes", {"hipDeviceGetNvSciSyncAttributes", "", CONV_DEVICE, API_RUNTIME, HIP_UNSUPPORTED}}, // cuDeviceGetP2PAttribute {"cudaDeviceGetP2PAttribute", {"hipDeviceGetP2PAttribute", "", CONV_DEVICE, API_RUNTIME, HIP_UNSUPPORTED}}, // cuDeviceGetPCIBusId @@ -209,7 +211,7 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ {"cudaSetDoubleForHost", {"hipSetDoubleForHost", "", CONV_EXECUTION, API_RUNTIME, HIP_UNSUPPORTED}}, // 5.8. Occupancy - // + // cuOccupancyMaxActiveBlocksPerMultiprocessor {"cudaOccupancyMaxActiveBlocksPerMultiprocessor", {"hipOccupancyMaxActiveBlocksPerMultiprocessor", "", CONV_OCCUPANCY, API_RUNTIME}}, // cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags {"cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags", {"hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags", "", CONV_OCCUPANCY, API_RUNTIME}}, @@ -222,7 +224,8 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ // no analogue {"cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags", {"hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags", "", CONV_OCCUPANCY, API_RUNTIME, HIP_UNSUPPORTED}}, - // 5.9. Execution Control [DEPRECATED] + // Former 5.9. Execution Control [DEPRECATED] + // NOTE: Removed in CUDA 10.1 // no analogue {"cudaConfigureCall", {"hipConfigureCall", "", CONV_EXECUTION, API_RUNTIME}}, // no analogue @@ -231,7 +234,7 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ // no analogue {"cudaSetupArgument", {"hipSetupArgument", "", CONV_EXECUTION, API_RUNTIME}}, - // 5.10. Memory Management + // 5.9. Memory Management // no analogue {"cudaArrayGetInfo", {"hipArrayGetInfo", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, // cuMemFree @@ -313,16 +316,9 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ // 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}}, - // no analogue // NOTE: Not equal to cuMemcpyAsync due to different signatures {"cudaMemcpyAsync", {"hipMemcpyAsync", "", CONV_MEMORY, API_RUNTIME}}, // no analogue - {"cudaMemcpyFromArray", {"hipMemcpyFromArray", "", CONV_MEMORY, API_RUNTIME}}, - // no analogue - {"cudaMemcpyFromArrayAsync", {"hipMemcpyFromArrayAsync", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, - // no analogue {"cudaMemcpyFromSymbol", {"hipMemcpyFromSymbol", "", CONV_MEMORY, API_RUNTIME}}, // no analogue {"cudaMemcpyFromSymbolAsync", {"hipMemcpyFromSymbolAsync", "", CONV_MEMORY, API_RUNTIME}}, @@ -333,10 +329,6 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ // NOTE: Not equal to cuMemcpyPeerAsync due to different signatures {"cudaMemcpyPeerAsync", {"hipMemcpyPeerAsync", "", CONV_MEMORY, API_RUNTIME}}, // no analogue - {"cudaMemcpyToArray", {"hipMemcpyToArray", "", CONV_MEMORY, API_RUNTIME}}, - // no analogue - {"cudaMemcpyToArrayAsync", {"hipMemcpyToArrayAsync", "", CONV_MEMORY, API_RUNTIME}}, - // no analogue {"cudaMemcpyToSymbol", {"hipMemcpyToSymbol", "", CONV_MEMORY, API_RUNTIME}}, // no analogue {"cudaMemcpyToSymbolAsync", {"hipMemcpyToSymbolAsync", "", CONV_MEMORY, API_RUNTIME}}, @@ -367,6 +359,19 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ // no analogue {"make_cudaPos", {"make_hipPos", "", CONV_MEMORY, API_RUNTIME}}, + // 5.10. Memory Management [DEPRECATED] + // no analogue + // NOTE: Not equal to cuMemcpyAtoA due to different signatures + {"cudaMemcpyArrayToArray", {"hipMemcpyArrayToArray", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, + // no analogue + {"cudaMemcpyFromArray", {"hipMemcpyFromArray", "", CONV_MEMORY, API_RUNTIME}}, + // no analogue + {"cudaMemcpyFromArrayAsync", {"hipMemcpyFromArrayAsync", "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, + // no analogue + {"cudaMemcpyToArray", {"hipMemcpyToArray", "", CONV_MEMORY, API_RUNTIME}}, + // no analogue + {"cudaMemcpyToArrayAsync", {"hipMemcpyToArrayAsync", "", CONV_MEMORY, API_RUNTIME}}, + // 5.11.Unified Addressing // no analogue // NOTE: Not equal to cuPointerGetAttributes due to different signatures @@ -546,7 +551,7 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ // cuGraphicsUnregisterResource {"cudaGraphicsUnregisterResource", {"hipGraphicsUnregisterResource", "", CONV_GRAPHICS, API_RUNTIME, HIP_UNSUPPORTED}}, - // 5.24. Texture Reference Management + // 5.24. Texture Reference Management [DEPRECATED] // no analogue {"cudaBindTexture", {"hipBindTexture", "", CONV_TEXTURE, API_RUNTIME}}, // no analogue @@ -568,7 +573,7 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ // no analogue {"cudaUnbindTexture", {"hipUnbindTexture", "", CONV_TEXTURE, API_RUNTIME}}, - // 5.25. Surface Reference Management + // 5.25. Surface Reference Management [DEPRECATED] // no analogue {"cudaBindSurfaceToArray", {"hipBindSurfaceToArray", "", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, // TODO: double check cuModuleGetSurfRef @@ -646,6 +651,14 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ {"cudaGraphInstantiate", {"hipGraphInstantiate", "", CONV_GRAPH, API_RUNTIME, HIP_UNSUPPORTED}}, // cuGraphExecKernelNodeSetParams {"cudaGraphExecKernelNodeSetParams", {"hipGraphExecKernelNodeSetParams", "", CONV_GRAPH, API_RUNTIME, HIP_UNSUPPORTED}}, + // cuGraphExecMemcpyNodeSetParams + {"cudaGraphExecMemcpyNodeSetParams", {"hipGraphExecMemcpyNodeSetParams", "", CONV_GRAPH, API_RUNTIME, HIP_UNSUPPORTED}}, + // cuGraphExecMemsetNodeSetParams + {"cudaGraphExecMemsetNodeSetParams", {"hipGraphExecMemsetNodeSetParams", "", CONV_GRAPH, API_RUNTIME, HIP_UNSUPPORTED}}, + // cuGraphExecHostNodeSetParams + {"cudaGraphExecHostNodeSetParams", {"hipGraphExecHostNodeSetParams", "", CONV_GRAPH, API_RUNTIME, HIP_UNSUPPORTED}}, + // cuGraphExecUpdate + {"cudaGraphExecUpdate", {"hipGraphExecUpdate", "", CONV_GRAPH, API_RUNTIME, HIP_UNSUPPORTED}}, // cuGraphKernelNodeGetParams {"cudaGraphKernelNodeGetParams", {"hipGraphKernelNodeGetParams", "", CONV_GRAPH, API_RUNTIME, HIP_UNSUPPORTED}}, // cuGraphKernelNodeSetParams diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp index d8acb9882c..5ad4659d06 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp @@ -720,6 +720,9 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaErrorStartupFailure", {"hipErrorStartupFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 127 // CUDA_ERROR_INVALID_IMAGE {"cudaErrorInvalidKernelImage", {"hipErrorInvalidImage", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 200 + // Typo fixed in 10.2 + // CUDA_ERROR_INVALID_CONTEXT + {"cudaErrorDeviceUninitialized", {"hipErrorInvalidContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 201 // CUDA_ERROR_INVALID_CONTEXT {"cudaErrorDeviceUninitilialized", {"hipErrorInvalidContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 201 // cudaErrorMapBufferObjectFailed = 205, CUDA_ERROR_MAP_FAILED = 205, hipErrorMapFailed = 205, hipErrorMapBufferObjectFailed = 1071 @@ -868,6 +871,10 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaErrorCapturedEvent", {"hipErrorCapturedEvent", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 907 // CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD {"cudaErrorStreamCaptureWrongThread", {"hipErrorStreamCaptureWrongThread", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 908 + // CUDA_ERROR_TIMEOUT + {"cudaErrorTimeout", {"hipErrorTimeout", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 909 + // CUDA_ERROR_GRAPH_EXEC_UPDATE_FAILURE + {"cudaErrorGraphExecUpdateFailure", {"hipErrorGraphExecUpdateFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 910 // cudaErrorUnknown = 999, CUDA_ERROR_UNKNOWN = 999, hipErrorUnknown = 1030 // TODO [HIP]: make hipErrorUnknown = 999 // CUDA_ERROR_UNKNOWN @@ -888,6 +895,12 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaExternalMemoryHandleTypeD3D12Heap", {"hipExternalMemoryHandleTypeD3D12Heap", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 4 // CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D12_RESOURCE {"cudaExternalMemoryHandleTypeD3D12Resource", {"hipExternalMemoryHandleTypeD3D12Resource", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 5 + // CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE + {"cudaExternalMemoryHandleTypeD3D11Resource", {"hipExternalMemoryHandleTypeD3D11Resource", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 6 + // CU_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_RESOURCE_KMT + {"cudaExternalMemoryHandleTypeD3D11ResourceKmt", {"hipExternalMemoryHandleTypeD3D11ResourceKmt", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 7 + // CU_EXTERNAL_MEMORY_HANDLE_TYPE_NVSCIBUF + {"cudaExternalMemoryHandleTypeNvSciBuf", {"hipExternalMemoryHandleTypeNvSciBuf", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 8 // CUexternalSemaphoreHandleType {"cudaExternalSemaphoreHandleType", {"hipExternalSemaphoreHandleType", "", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, @@ -900,6 +913,14 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt", {"hipExternalSemaphoreHandleTypeOpaqueWin32KMT", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 3 // CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D12_FENCE {"cudaExternalSemaphoreHandleTypeD3D12Fence", {"hipExternalSemaphoreHandleTypeD3D12Fence", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 4 + // CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_FENCE + {"cudaExternalSemaphoreHandleTypeD3D11Fence", {"hipExternalSemaphoreHandleTypeD3D11Fence", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 5 + // CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_NVSCISYNC + {"cudaExternalSemaphoreHandleTypeNvSciSync", {"hipExternalSemaphoreHandleTypeNvSciSync", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 6 + // CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX + {"cudaExternalSemaphoreHandleTypeKeyedMutex", {"hipExternalSemaphoreHandleTypeKeyedMutex", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 7 + // CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_D3D11_KEYED_MUTEX_KMT + {"cudaExternalSemaphoreHandleTypeKeyedMutexKmt", {"hipExternalSemaphoreHandleTypeKeyedMutexKmt", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 8 // CUfunction_attribute // NOTE: only last, starting from 8, values are presented and are equal to Driver's ones @@ -982,6 +1003,24 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { // CU_GRAPH_NODE_TYPE_COUNT {"cudaGraphNodeTypeCount", {"hipGraphNodeTypeCount", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, + // CUgraphExecUpdateResult + {"cudaGraphExecUpdateResult", {"hipGraphExecUpdateResult", "", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}}, + // cudaGraphExecUpdateResult enum values + // CU_GRAPH_EXEC_UPDATE_SUCCESS + {"cudaGraphExecUpdateSuccess", {"hipGraphExecUpdateSuccess", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x0 + // CU_GRAPH_EXEC_UPDATE_ERROR + {"cudaGraphExecUpdateError", {"hipGraphExecUpdateError", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x1 + // CU_GRAPH_EXEC_UPDATE_ERROR_TOPOLOGY_CHANGED + {"cudaGraphExecUpdateErrorTopologyChanged", {"hipGraphExecUpdateErrorTopologyChanged", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x2 + // CU_GRAPH_EXEC_UPDATE_ERROR_NODE_TYPE_CHANGED + {"cudaGraphExecUpdateErrorNodeTypeChanged", {"hipGraphExecUpdateErrorNodeTypeChanged", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x3 + // CU_GRAPH_EXEC_UPDATE_ERROR_FUNCTION_CHANGED + {"cudaGraphExecUpdateErrorFunctionChanged", {"hipGraphExecUpdateErrorFunctionChanged", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x4 + // CU_GRAPH_EXEC_UPDATE_ERROR_PARAMETERS_CHANGED + {"cudaGraphExecUpdateErrorParametersChanged", {"hipGraphExecUpdateErrorParametersChanged", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x5 + // CU_GRAPH_EXEC_UPDATE_ERROR_NOT_SUPPORTED + {"cudaGraphExecUpdateErrorNotSupported", {"hipGraphExecUpdateErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x6 + // CUlimit {"cudaLimit", {"hipLimit_t", "", CONV_TYPE, API_RUNTIME}}, // cudaLimit enum values @@ -1381,7 +1420,15 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { // CU_EVENT_INTERPROCESS {"cudaEventInterprocess", {"hipEventInterprocess", "", CONV_DEFINE, API_RUNTIME}}, // 0x04 // CUDA_EXTERNAL_MEMORY_DEDICATED - {"cudaExternalMemoryDedicated", {"hipExternalMemoryDedicated", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x01 + {"cudaExternalMemoryDedicated", {"hipExternalMemoryDedicated", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x1 + // CUDA_EXTERNAL_SEMAPHORE_SIGNAL_SKIP_NVSCIBUF_MEMSYNC + {"cudaExternalSemaphoreSignalSkipNvSciBufMemSync", {"hipExternalSemaphoreSignalSkipNvSciBufMemSync", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x01 + // CUDA_EXTERNAL_SEMAPHORE_WAIT_SKIP_NVSCIBUF_MEMSYNC + {"cudaExternalSemaphoreWaitSkipNvSciBufMemSync", {"hipExternalSemaphoreWaitSkipNvSciBufMemSync", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x02 + // CUDA_NVSCISYNC_ATTR_SIGNAL + {"cudaNvSciSyncAttrSignal", {"hipNvSciSyncAttrSignal", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x1 + // CUDA_NVSCISYNC_ATTR_WAIT + {"cudaNvSciSyncAttrWait", {"hipNvSciSyncAttrWait", "", CONV_DEFINE, API_RUNTIME, HIP_UNSUPPORTED}}, // 0x2 // no analogue {"cudaHostAllocDefault", {"hipHostMallocDefault", "", CONV_DEFINE, API_RUNTIME}}, // 0x00 // CU_MEMHOSTALLOC_PORTABLE diff --git a/projects/hip/hipify-clang/src/Statistics.cpp b/projects/hip/hipify-clang/src/Statistics.cpp index 686b51da41..1f7713cd88 100644 --- a/projects/hip/hipify-clang/src/Statistics.cpp +++ b/projects/hip/hipify-clang/src/Statistics.cpp @@ -34,6 +34,7 @@ const char *counterNames[NUM_CONV_TYPES] = { "context", // CONV_CONTEXT "module", // CONV_MODULE "memory", // CONV_MEMORY + "virtual_memory", // CONV_VIRTUAL_MEMORY "addressing", // CONV_ADDRESSING "stream", // CONV_STREAM "event", // CONV_EVENT @@ -76,6 +77,7 @@ const char *counterTypes[NUM_CONV_TYPES] = { "CONV_CONTEXT", "CONV_MODULE", "CONV_MEMORY", + "CONV_VIRTUAL_MEMORY", "CONV_ADDRESSING", "CONV_STREAM", "CONV_EVENT", diff --git a/projects/hip/hipify-clang/src/Statistics.h b/projects/hip/hipify-clang/src/Statistics.h index 030ea6b2d6..6cff9cd9d6 100644 --- a/projects/hip/hipify-clang/src/Statistics.h +++ b/projects/hip/hipify-clang/src/Statistics.h @@ -49,36 +49,38 @@ enum ConvTypes { // Driver API : 5.10. Module Management CONV_MODULE, // Driver API : 5.11. Memory Management - // Runtime API: 5.10. Memory Management + // Runtime API: 5.9. Memory Management, 5.10. Memory Management [DEPRECATED] CONV_MEMORY, - // Driver API : 5.12. Unified Addressing + // Driver API : 5.12. Virtual Memory Management + CONV_VIRTUAL_MEMORY, + // Driver API : 5.13. Unified Addressing // Runtime API: 5.11. Unified Addressing CONV_ADDRESSING, - // Driver API : 5.13. Stream Management + // Driver API : 5.14. Stream Management // Runtime API: 5.4. Stream Management CONV_STREAM, - // Driver API : 5.14. Event Management + // Driver API : 5.15. Event Management // Runtime API: 5.5. Event Management CONV_EVENT, - // Driver API : 5.15. External Resource Interoperability + // Driver API : 5.16. External Resource Interoperability // Runtime API: 5.6.External Resource Interoperability CONV_EXT_RES, - // Driver API : 5.16. Stream memory operations + // Driver API : 5.17. 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] + // Driver API : 5.18. Execution Control, 5.19. Execution Control [DEPRECATED] + // Runtime API: 5.7.Execution Control, Former 5.9. Execution Control [DEPRECATED] CONV_EXECUTION, - // Driver API : 5.19. Graph Management + // Driver API : 5.20. Graph Management // Runtime API: 5.29. Graph Management CONV_GRAPH, - // Driver API : 5.20. Occupancy + // Driver API : 5.21. 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 + // Driver API : 5.22. Texture Reference Management [DEPRECATED], 5.24. Texture Object Management + // Runtime API: 5.24. Texture Reference Management [DEPRECATED], 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 + // Driver API : 5.23. Surface Reference Management [DEPRECATED], 5.25. Surface Object Management + // Runtime API: 5.25. Surface Reference Management [DEPRECATED], 5.27. Surface Object Management CONV_SURFACE, // Driver API : 5.26. Peer Context Memory Access // Runtime API: 5.12. Peer Device Memory Access diff --git a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp index c9953c27d4..e7d8cc623f 100644 --- a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp @@ -34,6 +34,9 @@ THE SOFTWARE. #include #include +hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, + unsigned int flags, hip_impl::program_state& ps); + #pragma GCC visibility push(hidden) namespace hip_impl { @@ -176,4 +179,15 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, numBlocks, dimBlocks, sharedMemBytes, stream, &config[0]); } + +inline +__attribute__((visibility("hidden"))) +hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, + int numDevices, unsigned int flags) { + hip_impl::hip_init(); + auto& ps = hip_impl::get_program_state(); + return ihipExtLaunchMultiKernelMultiDevice(launchParamsList, numDevices, flags, ps); + +} + #pragma GCC visibility pop diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 48cec81da4..4397e7fe50 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -1277,7 +1277,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height * @param[in] height Requested pitched allocation height * * If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned. - * The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array. + * The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array. * Given the row and column of an array element of type T, the address is computed as: * T* pElement = (T*)((char*)BaseAddress + Row * Pitch) + Column; * @@ -1361,10 +1361,6 @@ hipError_t hipHostFree(void* ptr); */ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind); -// TODO: Add description -hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes, - hipMemcpyKind kind, hipStream_t stream); - /** * @brief Copy data from Host to Device * @@ -1502,7 +1498,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, #else hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*); -#ifdef __cplusplus //Start : Not supported in gcc +#ifdef __cplusplus //Start : Not supported in gcc namespace hip_impl { inline __attribute__((visibility("hidden"))) @@ -2960,6 +2956,7 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags); +#if defined(__clang__) && defined(__HIP__) /** * @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched * on respective streams before enqueuing any other work on the specified streams from any other threads @@ -2974,6 +2971,7 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags); +#endif // doxygen end Version Management /** @@ -3395,12 +3393,14 @@ inline hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchP return hipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags); } +#if defined(__clang__) && defined(__HIP__) template inline hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, unsigned int numDevices, unsigned int flags = 0) { return hipExtLaunchMultiKernelMultiDevice(launchParamsList, numDevices, flags); } +#endif /* * @brief Unbinds the textuer bound to @p tex diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index f3ce66c94e..9cb59f14ea 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -610,16 +610,6 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, cudaMemcpy(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind))); } -inline hipError_t hipMemcpyWithStream(void* dst, const void* src, - size_t sizeBytes, hipMemcpyKind copyKind, - hipStream_t stream) { - cudaError_t error = cudaMemcpyAsync(dst, src, sizeBytes, - hipMemcpyKindToCudaMemcpyKind(copyKind), - stream); - if (error != cudaSuccess) return hipCUDAErrorTohipError(error); - - return hipCUDAErrorTohipError(cudaStreamSynchronize(stream)); -} inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind, hipStream_t stream __dparm(0)) { diff --git a/projects/hip/install.sh b/projects/hip/install.sh index c7dc48e1a8..012fea543f 100755 --- a/projects/hip/install.sh +++ b/projects/hip/install.sh @@ -31,12 +31,11 @@ function setupENV() function buildHIP() { pushd $BUILD_ROOT - cmake $SRC_ROOT -DCMAKE_BUILD_TYPE=Release -DCOMPILE_HIP_ATP_MARKER=1 + cmake $SRC_ROOT -DCMAKE_BUILD_TYPE=Release make $DASH_JAY make package - rename -v 's/([a-z0-9_.\-]).deb/$1-amd64.deb/' *.deb;rename -v 's/([a-z0-9_.\-]).rpm/$1.x86_64.rpm/' *.rpm - cp hip_*.deb $WORKING_DIR - sudo dpkg -i hip_base*.deb hip_hcc*.deb hip_sample*.deb hip_doc*.deb + cp hip-*.deb $WORKING_DIR + sudo dpkg -i -B hip-base*.deb hip-hcc*.deb hip-sample*.deb hip-doc*.deb popd rm -rf $BUILD_ROOT } diff --git a/projects/hip/packaging/hip_base.postinst b/projects/hip/packaging/hip-base.postinst similarity index 100% rename from projects/hip/packaging/hip_base.postinst rename to projects/hip/packaging/hip-base.postinst diff --git a/projects/hip/packaging/hip_base.prerm b/projects/hip/packaging/hip-base.prerm similarity index 100% rename from projects/hip/packaging/hip_base.prerm rename to projects/hip/packaging/hip-base.prerm diff --git a/projects/hip/packaging/hip_base.txt b/projects/hip/packaging/hip-base.txt similarity index 87% rename from projects/hip/packaging/hip_base.txt rename to projects/hip/packaging/hip-base.txt index a44f0acb74..927812f910 100644 --- a/projects/hip/packaging/hip_base.txt +++ b/projects/hip/packaging/hip-base.txt @@ -16,7 +16,7 @@ install(DIRECTORY @hip_SOURCE_DIR@/cmake DESTINATION .) ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip_base") +set(CPACK_PACKAGE_NAME "hip-base") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [BASE]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -29,12 +29,18 @@ set(CPACK_GENERATOR "TGZ;DEB;RPM") set(CPACK_BINARY_DEB "ON") set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm") set(CPACK_DEBIAN_PACKAGE_DEPENDS "perl (>= 5.0)") +set(CPACK_DEBIAN_PACKAGE_PROVIDES "hip_base") +set(CPACK_DEBIAN_PACKAGE_REPLACES "hip_base") +set(CPACK_DEBIAN_PACKAGE_CONFLICTS "hip_base") set(CPACK_BINARY_RPM "ON") set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}") set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst") set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm") set(CPACK_RPM_PACKAGE_AUTOREQPROV " no") set(CPACK_RPM_PACKAGE_REQUIRES "perl >= 5.0") +set(CPACK_RPM_PACKAGE_OBSOLETES "hip_base") +set(CPACK_RPM_PACKAGE_CONFLICTS "hip_base") +set(CPACK_BINARY_RPM "ON") set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/packaging/hip_doc.txt b/projects/hip/packaging/hip-doc.txt similarity index 82% rename from projects/hip/packaging/hip_doc.txt rename to projects/hip/packaging/hip-doc.txt index 2c1619f221..9d39822340 100644 --- a/projects/hip/packaging/hip_doc.txt +++ b/projects/hip/packaging/hip-doc.txt @@ -20,7 +20,7 @@ endif() ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip_doc") +set(CPACK_PACKAGE_NAME "hip-doc") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [DOCUMENTATION]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -31,12 +31,17 @@ set(CPACK_PACKAGE_VERSION_PATCH @HIP_VERSION_PATCH@) set(CPACK_PACKAGE_FILE_NAME ${CPACK_PACKAGE_NAME}-${CPACK_PACKAGE_VERSION_MAJOR}.${CPACK_PACKAGE_VERSION_MINOR}.${CPACK_PACKAGE_VERSION_PATCH}) set(CPACK_GENERATOR "TGZ;DEB;RPM") set(CPACK_BINARY_DEB "ON") -set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION})") +set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-base (= ${CPACK_PACKAGE_VERSION})") +set(CPACK_DEBIAN_PACKAGE_PROVIDES "hip_doc") +set(CPACK_DEBIAN_PACKAGE_REPLACES "hip_doc") +set(CPACK_DEBIAN_PACKAGE_CONFLICTS "hip_doc") set(CPACK_BINARY_RPM "ON") set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}") set(CPACK_RPM_PACKAGE_AUTOREQPROV " no") string(REPLACE "-" "_" HIP_BASE_VERSION ${CPACK_PACKAGE_VERSION}) -set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${HIP_BASE_VERSION}") +set(CPACK_RPM_PACKAGE_REQUIRES "hip-base = ${HIP_BASE_VERSION}") +set(CPACK_RPM_PACKAGE_OBSOLETES "hip_doc") +set(CPACK_RPM_PACKAGE_CONFLICTS "hip_doc") set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/packaging/hip_hcc.postinst b/projects/hip/packaging/hip-hcc.postinst similarity index 100% rename from projects/hip/packaging/hip_hcc.postinst rename to projects/hip/packaging/hip-hcc.postinst diff --git a/projects/hip/packaging/hip_hcc.prerm b/projects/hip/packaging/hip-hcc.prerm similarity index 100% rename from projects/hip/packaging/hip_hcc.prerm rename to projects/hip/packaging/hip-hcc.prerm diff --git a/projects/hip/packaging/hip_hcc.txt b/projects/hip/packaging/hip-hcc.txt similarity index 83% rename from projects/hip/packaging/hip_hcc.txt rename to projects/hip/packaging/hip-hcc.txt index c688433cc0..919f7dabf5 100644 --- a/projects/hip/packaging/hip_hcc.txt +++ b/projects/hip/packaging/hip-hcc.txt @@ -15,7 +15,7 @@ install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/pack ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip_hcc") +set(CPACK_PACKAGE_NAME "hip-hcc") if(@HCC_VERSION_MAJOR@ EQUAL 0) set(HCC_PACKAGE_NAME "hcc_lc") else() @@ -33,10 +33,13 @@ set(CPACK_GENERATOR "TGZ;DEB;RPM") set(CPACK_BINARY_DEB "ON") set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm") if(@COMPILE_HIP_ATP_MARKER@) - set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), rocm-profiler, comgr (>= 1.1)") + set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), rocm-profiler, comgr (>= 1.1)") else() - set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), comgr (>= 1.1)") + set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), comgr (>= 1.1)") endif() +set(CPACK_DEBIAN_PACKAGE_PROVIDES "hip_hcc") +set(CPACK_DEBIAN_PACKAGE_REPLACES "hip_hcc") +set(CPACK_DEBIAN_PACKAGE_CONFLICTS "hip_hcc") set(CPACK_BINARY_RPM "ON") set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}") set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst") @@ -44,10 +47,12 @@ set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm") set(CPACK_RPM_PACKAGE_AUTOREQPROV " no") string(REPLACE "-" "_" HIP_BASE_VERSION ${CPACK_PACKAGE_VERSION}) if(@COMPILE_HIP_ATP_MARKER@) - set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, rocm-profiler, comgr >= 1.1") + set(CPACK_RPM_PACKAGE_REQUIRES "hip-base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, rocm-profiler, comgr >= 1.1") else() - set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, comgr >= 1.1") + set(CPACK_RPM_PACKAGE_REQUIRES "hip-base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, comgr >= 1.1") endif() +set(CPACK_RPM_PACKAGE_OBSOLETES "hip_hcc") +set(CPACK_RPM_PACKAGE_CONFLICTS "hip_hcc") set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/packaging/hip_nvcc.txt b/projects/hip/packaging/hip-nvcc.txt similarity index 79% rename from projects/hip/packaging/hip_nvcc.txt rename to projects/hip/packaging/hip-nvcc.txt index 975176f934..010d770458 100644 --- a/projects/hip/packaging/hip_nvcc.txt +++ b/projects/hip/packaging/hip-nvcc.txt @@ -6,7 +6,7 @@ project(hip_nvcc) ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip_nvcc") +set(CPACK_PACKAGE_NAME "hip-nvcc") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [NVCC]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -18,14 +18,19 @@ set(CPACK_PACKAGE_FILE_NAME ${CPACK_PACKAGE_NAME}-${CPACK_PACKAGE_VERSION_MAJOR} set(CPACK_GENERATOR "TGZ;DEB;RPM") set(CPACK_BINARY_DEB "ON") #set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm") -set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), cuda (>= 7.5)") +set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-base (= ${CPACK_PACKAGE_VERSION}), cuda (>= 7.5)") +set(CPACK_DEBIAN_PACKAGE_PROVIDES "hip_nvcc") +set(CPACK_DEBIAN_PACKAGE_REPLACES "hip_nvcc") +set(CPACK_DEBIAN_PACKAGE_CONFLICTS "hip_nvcc") set(CPACK_BINARY_RPM "ON") set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}") #set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst") #set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm") set(CPACK_RPM_PACKAGE_AUTOREQPROV " no") string(REPLACE "-" "_" HIP_BASE_VERSION ${CPACK_PACKAGE_VERSION}) -set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${HIP_BASE_VERSION}, cuda >= 7.5") +set(CPACK_RPM_PACKAGE_REQUIRES "hip-base = ${HIP_BASE_VERSION}, cuda >= 7.5") +set(CPACK_RPM_PACKAGE_OBSOLETES "hip_nvcc") +set(CPACK_RPM_PACKAGE_CONFLICTS "hip_nvcc") set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/packaging/hip_samples.txt b/projects/hip/packaging/hip-samples.txt similarity index 75% rename from projects/hip/packaging/hip_samples.txt rename to projects/hip/packaging/hip-samples.txt index 4c1e04cb23..497ffe6ffc 100644 --- a/projects/hip/packaging/hip_samples.txt +++ b/projects/hip/packaging/hip-samples.txt @@ -8,7 +8,7 @@ install(DIRECTORY @hip_SOURCE_DIR@/samples DESTINATION .) ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip_samples") +set(CPACK_PACKAGE_NAME "hip-samples") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [SAMPLES]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -19,12 +19,17 @@ set(CPACK_PACKAGE_VERSION_PATCH @HIP_VERSION_PATCH@) set(CPACK_PACKAGE_FILE_NAME ${CPACK_PACKAGE_NAME}-${CPACK_PACKAGE_VERSION_MAJOR}.${CPACK_PACKAGE_VERSION_MINOR}.${CPACK_PACKAGE_VERSION_PATCH}) set(CPACK_GENERATOR "TGZ;DEB;RPM") set(CPACK_BINARY_DEB "ON") -set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION})") +set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-base (= ${CPACK_PACKAGE_VERSION})") +set(CPACK_DEBIAN_PACKAGE_PROVIDES "hip_samples") +set(CPACK_DEBIAN_PACKAGE_REPLACES "hip_samples") +set(CPACK_DEBIAN_PACKAGE_CONFLICTS "hip_samples") set(CPACK_BINARY_RPM "ON") set(CPACK_RPM_PACKAGE_ARCHITECTURE "${CMAKE_SYSTEM_PROCESSOR}") set(CPACK_RPM_PACKAGE_AUTOREQPROV " no") string(REPLACE "-" "_" HIP_BASE_VERSION ${CPACK_PACKAGE_VERSION}) -set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${HIP_BASE_VERSION}") +set(CPACK_RPM_PACKAGE_REQUIRES "hip-base = ${HIP_BASE_VERSION}") +set(CPACK_RPM_PACKAGE_OBSOLETES "hip_samples") +set(CPACK_RPM_PACKAGE_CONFLICTS "hip_samples") set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 00a7af59f8..266f9b51d6 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -27,7 +27,6 @@ THE SOFTWARE. #include "hip_hcc_internal.h" #include "trace_helper.h" -#include #include __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP]; @@ -36,369 +35,23 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; // Internal HIP APIS: namespace hip_internal { -namespace { - inline - const char* hsa_to_string(hsa_status_t err) noexcept - { - const char* r{}; +hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, + hipStream_t stream) { + hipError_t e = hipSuccess; - if (hsa_status_string(err, &r) != HSA_STATUS_SUCCESS) return r; - - return "Unknown."; - } - - template - inline - void throwing_result_check(hsa_status_t res, const char (&file)[m], - const char (&function)[n], int line) { - if (res == HSA_STATUS_SUCCESS) return; - if (res == HSA_STATUS_INFO_BREAK) return; - - throw std::runtime_error{"Failed in file " + (file + - (", in function \"" + (function + - ("\", on line " + std::to_string(line))))) + - ", with error: " + hsa_to_string(res)}; - } - - inline - hsa_agent_t cpu_agent() noexcept { - static hsa_agent_t cpu{[]() { - hsa_agent_t r{}; - throwing_result_check( - hsa_iterate_agents([](hsa_agent_t x, void* pr) { - hsa_device_type_t t{}; - hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &t); - - if (t != HSA_DEVICE_TYPE_CPU) return HSA_STATUS_SUCCESS; - - *static_cast(pr) = x; - - return HSA_STATUS_INFO_BREAK; - }, &r), __FILE__, __func__, __LINE__); - - return r; - }()}; - - return cpu; - } - - inline - hsa_device_type_t type(hsa_agent_t x) noexcept - { - hsa_device_type_t r{}; - throwing_result_check(hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &r), - __FILE__, __func__, __LINE__); - - return r; - } - - const auto is_large_BAR{[](){ - std::unique_ptr hsa{ - (throwing_result_check(hsa_init(), __FILE__, __func__, __LINE__), - nullptr), - [](void*) { hsa_shut_down(); }}; - bool r{true}; - - throwing_result_check(hsa_iterate_agents([](hsa_agent_t x, void* pr) { - if (x.handle == cpu_agent().handle) return HSA_STATUS_SUCCESS; - - throwing_result_check( - hsa_agent_iterate_regions(x, [](hsa_region_t y, void* p) { - hsa_region_segment_t seg{}; - throwing_result_check( - hsa_region_get_info(y, HSA_REGION_INFO_SEGMENT, &seg), - __FILE__, __func__, __LINE__); - - if (seg != HSA_REGION_SEGMENT_GLOBAL) { - return HSA_STATUS_SUCCESS; - } - - uint32_t flags{}; - throwing_result_check(hsa_region_get_info( - y, HSA_REGION_INFO_GLOBAL_FLAGS, &flags), - __FILE__, __func__, __LINE__); - - if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) { - hsa_amd_memory_pool_access_t tmp{}; - throwing_result_check( - hsa_amd_agent_memory_pool_get_info( - cpu_agent(), - hsa_amd_memory_pool_t{y.handle}, - HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, - &tmp), - __FILE__, __func__, __LINE__); - - *static_cast(p) &= - tmp != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED; - } - - return HSA_STATUS_SUCCESS; - }, pr), __FILE__, __func__, __LINE__); - - return HSA_STATUS_SUCCESS; - }, &r), __FILE__, __func__, __LINE__); - - return r; - }()}; - - inline - hsa_amd_pointer_info_t info(const void* p) noexcept - { - hsa_amd_pointer_info_t r{sizeof(hsa_amd_pointer_info_t)}; - throwing_result_check( - hsa_amd_pointer_info( - const_cast(p), &r, nullptr, nullptr, nullptr), - __FILE__, __func__, __LINE__); - - r.size = is_large_BAR || (type(r.agentOwner) == HSA_DEVICE_TYPE_CPU) ? - UINT32_MAX : sizeof(hsa_amd_pointer_info_t); - - return r; - } - - constexpr size_t staging_sz{4 * 1024 * 1024}; // 2 Pages. - - thread_local const std::unique_ptr staging_buffer{ - []() { - hsa_region_t r{}; - throwing_result_check(hsa_agent_iterate_regions( - cpu_agent(), [](hsa_region_t x, void *p) { - hsa_region_segment_t seg{}; - throwing_result_check( - hsa_region_get_info(x, HSA_REGION_INFO_SEGMENT, &seg), - __FILE__, __func__, __LINE__); - - if (seg != HSA_REGION_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS; - - uint32_t flags{}; - throwing_result_check(hsa_region_get_info( - x, HSA_REGION_INFO_GLOBAL_FLAGS, &flags), - __FILE__, __func__, __LINE__); - - if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) { - *static_cast(p) = x; - - return HSA_STATUS_INFO_BREAK; - } - - return HSA_STATUS_SUCCESS; - }, &r), __FILE__, __func__, __LINE__); - - void *tp{}; - throwing_result_check(hsa_memory_allocate(r, staging_sz, &tp), - __FILE__, __func__, __LINE__); - - return tp; - }(), - [](void *ptr) { hsa_memory_free(ptr); }}; - - thread_local hsa_signal_t copy_signal{[]() { - hsa_agent_t cpu{cpu_agent()}; - hsa_signal_t sgn{}; - throwing_result_check(hsa_signal_create(1, 1, &cpu, &sgn), - __FILE__, __func__, __LINE__); - - return sgn; - }()}; -} // Unnamed namespace. - -inline -void do_copy(void* __restrict dst, const void* __restrict src, std::size_t n, - hsa_agent_t da, hsa_agent_t sa) { - hsa_signal_silent_store_relaxed(copy_signal, 1); - throwing_result_check( - hsa_amd_memory_async_copy(dst, da, src, sa, n, 0, nullptr, copy_signal), - __FILE__, __func__, __LINE__); - - while (hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_EQ, 0, - UINT64_MAX, HSA_WAIT_STATE_ACTIVE)); -} - -inline -void do_std_memcpy( - void* __restrict dst, const void* __restrict src, std::size_t n) { - std::memcpy(dst, src, n); - - return std::atomic_thread_fence(std::memory_order_seq_cst); -} - -inline -void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n, - hsa_amd_pointer_info_t si) { - if (si.size == UINT32_MAX) return do_std_memcpy(dst, src, n); - - const auto di{info(dst)}; - - if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) { - dst = static_cast(di.agentBaseAddress) + - (static_cast(dst) - - static_cast(di.hostBaseAddress)); - do_copy(dst, src, n, si.agentOwner, si.agentOwner); - } - else if (n <= staging_sz) { - do_copy(staging_buffer.get(), src, n, si.agentOwner, si.agentOwner); - std::memcpy(dst, staging_buffer.get(), n); - } - else { - std::unique_ptr lck{ - dst, [](void* p) { hsa_amd_memory_unlock(p); }}; - - throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1, - const_cast(&dst)), - __FILE__, __func__, __LINE__); - - do_copy(dst, src, n, si.agentOwner, si.agentOwner); - } -} - -inline -void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n, - hsa_amd_pointer_info_t di) { - if (di.size == UINT32_MAX) return do_std_memcpy(dst, src, n); - - const auto si{info(const_cast(src))}; - - if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) { - src = static_cast(si.agentBaseAddress) + - (static_cast(src) - - static_cast(si.hostBaseAddress)); - do_copy(dst, src, n, di.agentOwner, di.agentOwner); - } - else if (n <= staging_sz) { - std::memcpy(staging_buffer.get(), src, n); - do_copy(dst, staging_buffer.get(), n, di.agentOwner, di.agentOwner); - } - else { - std::unique_ptr lck{ - const_cast(src), [](void* p) { hsa_amd_memory_unlock(p); }}; - - throwing_result_check(hsa_amd_memory_lock(const_cast(src), n, - &di.agentOwner, 1, - const_cast(&src)), - __FILE__, __func__, __LINE__); - - do_copy(dst, src, n, di.agentOwner, di.agentOwner); - } -} - -inline -void generic_copy(void* __restrict dst, const void* __restrict src, size_t n, - hsa_amd_pointer_info_t di, hsa_amd_pointer_info_t si) { - if (di.size == UINT32_MAX && si.size == UINT32_MAX) { - return do_std_memcpy(dst, src, n); - } - - std::unique_ptr lck0{ - nullptr, [](void* p) { hsa_amd_memory_unlock(p); }}; - std::unique_ptr lck1{nullptr, lck0.get_deleter()}; - - switch (si.type) { - case HSA_EXT_POINTER_TYPE_HSA: - if (di.type == HSA_EXT_POINTER_TYPE_HSA) { - hsa_memory_copy(dst, src, n); - return; // TODO: do_copy(dst, src, n, di.agentOwner, si.agentOwner); - } - - if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN || - di.type == HSA_EXT_POINTER_TYPE_LOCKED) { - return d2h_copy(dst, src, n, si); - } - break; - case HSA_EXT_POINTER_TYPE_LOCKED: - if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) { - std::memcpy(dst, si.hostBaseAddress, n); - - return; - } - if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) { - std::memcpy(di.hostBaseAddress, si.hostBaseAddress, n); - - return; - } - src = si.agentBaseAddress; - si.agentOwner = di.agentOwner; - break; - case HSA_EXT_POINTER_TYPE_UNKNOWN: - if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) { - std::memcpy(dst, src, n); - - return; - } - if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) { - std::memcpy(di.hostBaseAddress, src, n); - - return; - } - return h2d_copy(dst, src, n, di); - default: do_copy(dst, src, n, di.agentOwner, si.agentOwner); break; - } -} - -inline -void memcpy_impl(void* __restrict dst, const void* __restrict src, size_t n, - hipMemcpyKind k) noexcept { - switch (k) { - case hipMemcpyHostToHost: std::memcpy(dst, src, n); break; - case hipMemcpyHostToDevice: - return is_large_BAR ? do_std_memcpy(dst, src, n) - : h2d_copy(dst, src, n, info(dst)); - case hipMemcpyDeviceToHost: - return is_large_BAR ? do_std_memcpy(dst, src, n) - : d2h_copy(dst, src, n, info(src)); - case hipMemcpyDeviceToDevice: - return do_copy(dst, src, n, info(dst).agentOwner, - info(const_cast(src)).agentOwner); - default: return generic_copy(dst, src, n, info(dst), info(src)); - } -} - -hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, - hipMemcpyKind kind, hipStream_t stream) { - if (sizeBytes == 0) return hipSuccess; + // Return success if number of bytes to copy is 0 + if (sizeBytes == 0) return e; if (!dst || !src) return hipErrorInvalidValue; + if (!(stream = ihipSyncAndResolveStream(stream))) { + return hipErrorInvalidValue; + } + try { - stream = ihipSyncAndResolveStream(stream); - - if (!stream) return hipErrorInvalidValue; - stream->locked_copyAsync(dst, src, sizeBytes, kind); } - catch (const ihipException& ex) { - return ex._code; - } - catch (const std::exception& ex) { - std::cerr << ex.what() << std::endl; - throw; - } - catch (...) { - return hipErrorUnknown; - } - - return hipSuccess; -} - -hipError_t memcpySync(void* dst, const void* src, size_t sizeBytes, - hipMemcpyKind kind, hipStream_t stream) { - if (sizeBytes == 0) return hipSuccess; - if (!dst || !src) return hipErrorInvalidValue; - - try { - stream = ihipSyncAndResolveStream(stream); - - if (!stream) return hipErrorInvalidValue; - - LockedAccessor_StreamCrit_t cs{stream->criticalData()}; - cs->_av.wait(); - - memcpy_impl(dst, src, sizeBytes, kind); - } - catch (const ihipException& ex) { - return ex._code; - } - catch (const std::exception& ex) { - std::cerr << ex.what() << std::endl; - throw; + catch (ihipException& ex) { + e = ex._code; } catch (...) { return hipErrorUnknown; @@ -1243,8 +896,20 @@ hipError_t hipMemcpyToSymbol(void* dst, const void* src, size_t count, tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst); - return ihipLogStatus( - hipMemcpy(static_cast(dst) + offset, src, count, kind)); + if (dst == nullptr) { + return ihipLogStatus(hipErrorInvalidSymbol); + } + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + if (kind == hipMemcpyHostToDevice || kind == hipMemcpyDefault || + kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) { + stream->locked_copySync((char*)dst+offset, (void*)src, count, kind, false); + } else { + return ihipLogStatus(hipErrorInvalidValue); + } + + return ihipLogStatus(hipSuccess); } hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count, @@ -1255,8 +920,20 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count, tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst); - return ihipLogStatus( - hipMemcpy(dst, static_cast(src) + offset, count, kind)); + if (dst == nullptr) { + return ihipLogStatus(hipErrorInvalidSymbol); + } + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + if (kind == hipMemcpyDefault || kind == hipMemcpyDeviceToHost || + kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) { + stream->locked_copySync((void*)dst, (char*)src+offset, count, kind, false); + } else { + return ihipLogStatus(hipErrorInvalidValue); + } + + return ihipLogStatus(hipSuccess); } @@ -1318,49 +995,120 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* src, size_t count, hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { HIP_INIT_SPECIAL_API(hipMemcpy, (TRACE_MCMD), dst, src, sizeBytes, kind); - return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind, - hipStreamNull)); + hipError_t e = hipSuccess; + + // Return success if number of bytes to copy is 0 + if (sizeBytes == 0) return ihipLogStatus(e); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + if(dst==NULL || src==NULL) + { + e=hipErrorInvalidValue; + return ihipLogStatus(e); + } + try { + stream->locked_copySync(dst, src, sizeBytes, kind); + } catch (ihipException& ex) { + e = ex._code; + } + + return ihipLogStatus(e); } + hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes); - return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, - hipMemcpyHostToDevice, - hipStreamNull)); + hipError_t e = hipSuccess; + if (sizeBytes == 0) return ihipLogStatus(e); + + if(dst==NULL || src==NULL){ + return ihipLogStatus(hipErrorInvalidValue); + } + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + try { + stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToDevice, false); + } catch (ihipException& ex) { + e = ex._code; + } + + return ihipLogStatus(e); } + hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes); - return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, - hipMemcpyDeviceToHost, - hipStreamNull)); + hipError_t e = hipSuccess; + if (sizeBytes == 0) return ihipLogStatus(e); + + if(dst==NULL || src==NULL){ + return ihipLogStatus(hipErrorInvalidValue); + } + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + try { + stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToHost, false); + } catch (ihipException& ex) { + e = ex._code; + } + + return ihipLogStatus(e); } hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes); - return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, - hipMemcpyDeviceToDevice, - hipStreamNull)); + hipError_t e = hipSuccess; + if (sizeBytes == 0) return ihipLogStatus(e); + + if(dst==NULL || src==NULL){ + return ihipLogStatus(hipErrorInvalidValue); + } + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + try { + stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToDevice, false); + } catch (ihipException& ex) { + e = ex._code; + } + + return ihipLogStatus(e); } hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemcpyHtoH, (TRACE_MCMD), dst, src, sizeBytes); - return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, - hipMemcpyHostToHost, - hipStreamNull)); -} + hipError_t e = hipSuccess; + if (sizeBytes == 0) return ihipLogStatus(e); -hipError_t hipMemcpyWithStream(void* dst, void* src, size_t sizeBytes, - hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_SPECIAL_API(hipMemcpyWithStream, (TRACE_MCMD), dst, src, sizeBytes, - kind, stream); + if(dst==NULL || src==NULL){ + return ihipLogStatus(hipErrorInvalidValue); + } - return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind, - stream)); + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + try { + stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToHost, false); + } catch (ihipException& ex) { + e = ex._code; + } + + return ihipLogStatus(e); } hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 498ee5bb03..552ccfccac 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -311,18 +311,18 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0)); } -hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, - int numDevices, unsigned int flags) { - HIP_INIT_API(hipExtLaunchMultiKernelMultiDevice, launchParamsList, numDevices, flags); +__attribute__((visibility("default"))) +hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, + int numDevices, unsigned int flags, hip_impl::program_state& ps) { hipError_t result; if ((numDevices > g_deviceCnt) || (launchParamsList == nullptr)) { - return ihipLogStatus(hipErrorInvalidValue); + return hipErrorInvalidValue; } hipFunction_t* kds = reinterpret_cast(malloc(sizeof(hipFunction_t) * numDevices)); if (kds == nullptr) { - return ihipLogStatus(hipErrorNotInitialized); + return hipErrorNotInitialized; } // prepare all kernel descriptors for each device as all streams will be locked in the next loop @@ -330,15 +330,15 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, const hipLaunchParams& lp = launchParamsList[i]; if (lp.stream == nullptr) { free(kds); - return ihipLogStatus(hipErrorNotInitialized); + return hipErrorNotInitialized; } - kds[i] = hip_impl::get_program_state().kernel_descriptor(reinterpret_cast(lp.func), + kds[i] = ps.kernel_descriptor(reinterpret_cast(lp.func), hip_impl::target_agent(lp.stream)); if (kds[i] == nullptr) { free(kds); - return ihipLogStatus(hipErrorInvalidValue); + return hipErrorInvalidValue; } - hip_impl::kernargs_size_align kargs = hip_impl::get_program_state().get_kernargs_size_align( + hip_impl::kernargs_size_align kargs = ps.get_kernargs_size_align( reinterpret_cast(lp.func)); kds[i]->_kernarg_layout = *reinterpret_cast>*>( kargs.getHandle()); @@ -352,6 +352,7 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, #endif } + GET_TLS(); // launch kernels for each device for (int i = 0; i < numDevices; ++i) { const hipLaunchParams& lp = launchParamsList[i]; @@ -377,7 +378,7 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, free(kds); - return ihipLogStatus(result); + return result; } namespace {