AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge 'master' into 'amd-master'

Просмотр исходного кода 
Change-Id: I11a48b1198f902fa74e9c0cf7e234c81fb704968
Этот коммит содержится в:
Jenkins
2019-12-06 10:51:37 +00:00
родитель 355b722246 e53fc316f1
Коммит 655f16bdb1
28 изменённых файлов: 602 добавлений и 523 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
CMakeLists.txt
+9 -9
Просмотреть файл
@@ -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
INSTALL.md
+4 -4
Просмотреть файл
@@ -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:
bin/hipify-perl
+5 -2
Просмотреть файл
@@ -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$/;
docker/dockerfile-hip-ubuntu-16.04
+4 -4
Просмотреть файл
@@ -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/*
docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md
+93 -17
Просмотреть файл
@@ -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\***|
|-----------------------------------------------------------|-------------------------------|------------------|
docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+42 -12
Просмотреть файл
@@ -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.
hipify-clang/src/CUDA2HIP_Device_functions.cpp
+2
Просмотреть файл
@@ -548,6 +548,8 @@ const std::map<llvm::StringRef, hipCounter> 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}},
hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp
+36 -14
Просмотреть файл
@@ -52,6 +52,8 @@ const std::map<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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}},
hipify-clang/src/CUDA2HIP_Driver_API_types.cpp
+106 -2
Просмотреть файл
@@ -198,6 +198,18 @@ const std::map<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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
};
hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp
+29 -16
Просмотреть файл
@@ -35,6 +35,8 @@ const std::map<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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
hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp
+48 -1
Просмотреть файл
@@ -720,6 +720,9 @@ const std::map<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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
hipify-clang/src/Statistics.cpp
+2
Просмотреть файл
@@ -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",
hipify-clang/src/Statistics.h
+16 -14
Просмотреть файл
@@ -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
include/hip/hcc_detail/functional_grid_launch.hpp
+14
Просмотреть файл
@@ -34,6 +34,9 @@ THE SOFTWARE.
#include <type_traits>
#include <utility>
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
include/hip/hcc_detail/hip_runtime_api.h
+6 -6
Просмотреть файл
@@ -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 <class T>
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
include/hip/nvcc_detail/hip_runtime_api.h
-10
Просмотреть файл
@@ -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)) {
install.sh
+3 -4
Просмотреть файл
@@ -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
}
packaging/hip_base.postinst → packaging/hip-base.postinst
Просмотреть файл
packaging/hip_base.prerm → packaging/hip-base.prerm
Просмотреть файл
packaging/hip_base.txt → packaging/hip-base.txt
+7 -1
Просмотреть файл
@@ -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 <maneesh.gupta@amd.com>")
@@ -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)
packaging/hip_doc.txt → packaging/hip-doc.txt
+8 -3
Просмотреть файл
@@ -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 <maneesh.gupta@amd.com>")
@@ -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)
packaging/hip_hcc.postinst → packaging/hip-hcc.postinst
Просмотреть файл
packaging/hip_hcc.prerm → packaging/hip-hcc.prerm
Просмотреть файл
packaging/hip_hcc.txt → packaging/hip-hcc.txt
+10 -5
Просмотреть файл
@@ -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)
packaging/hip_nvcc.txt → packaging/hip-nvcc.txt
+8 -3
Просмотреть файл
@@ -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 <maneesh.gupta@amd.com>")
@@ -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)
packaging/hip_samples.txt → packaging/hip-samples.txt
+8 -3
Просмотреть файл
@@ -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 <maneesh.gupta@amd.com>")
@@ -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)
src/hip_memory.cpp
+131 -383
Просмотреть файл
@@ -27,7 +27,6 @@ THE SOFTWARE.
#include "hip_hcc_internal.h"
#include "trace_helper.h"
#include <functional>
#include <fstream>
__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<std::size_t m, std::size_t n>
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<hsa_agent_t *>(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<void, void (*)(void*)> 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<bool*>(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<void*>(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<void, void (*)(void *)> 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<hsa_region_t *>(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<char*>(di.agentBaseAddress) +
(static_cast<char*>(dst) -
static_cast<char*>(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<void, void (*)(void*)> lck{
dst, [](void* p) { hsa_amd_memory_unlock(p); }};
throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1,
const_cast<void**>(&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<void*>(src))};
if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) {
src = static_cast<char*>(si.agentBaseAddress) +
(static_cast<const char*>(src) -
static_cast<char*>(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<void, void (*)(void*)> lck{
const_cast<void*>(src), [](void* p) { hsa_amd_memory_unlock(p); }};
throwing_result_check(hsa_amd_memory_lock(const_cast<void*>(src), n,
&di.agentOwner, 1,
const_cast<void**>(&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<void, void (*)(void*)> lck0{
nullptr, [](void* p) { hsa_amd_memory_unlock(p); }};
std::unique_ptr<void, void (*)(void*)> 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<void*>(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<char*>(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<const char*>(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,
src/hip_module.cpp
+11 -10
Просмотреть файл
@@ -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<hipFunction_t*>(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<std::uintptr_t>(lp.func),
kds[i] = ps.kernel_descriptor(reinterpret_cast<std::uintptr_t>(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<std::uintptr_t>(lp.func));
kds[i]->_kernarg_layout = *reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(
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 {