From 8016fa32df714b478f2d9431d85282a71c61c11c Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 14 Nov 2019 11:40:34 +0530 Subject: [PATCH 01/27] [packaging] Conform to debian package naming rules (#1647) * [packaging] Conform to debian package naming rules * Add hip_hcc metapackage for backward compatibility [ROCm/hip commit: 2ca1a82036ee775039da06220b4cc824153b2ff9] --- projects/hip/CMakeLists.txt | 31 +++++++---- projects/hip/INSTALL.md | 8 +-- .../hip/docker/dockerfile-hip-ubuntu-16.04 | 8 +-- projects/hip/install.sh | 4 +- .../{hip_base.postinst => hip-base.postinst} | 0 .../{hip_base.prerm => hip-base.prerm} | 0 .../packaging/{hip_base.txt => hip-base.txt} | 3 +- .../packaging/{hip_doc.txt => hip-doc.txt} | 7 +-- .../{hip_hcc.postinst => hip-hcc.postinst} | 0 .../{hip_hcc.prerm => hip-hcc.prerm} | 0 projects/hip/packaging/hip-hcc.txt | 54 +++++++++++++++++++ .../packaging/{hip_nvcc.txt => hip-nvcc.txt} | 7 +-- .../{hip_samples.txt => hip-samples.txt} | 7 +-- projects/hip/packaging/hip_hcc.txt | 31 ++--------- 14 files changed, 102 insertions(+), 58 deletions(-) rename projects/hip/packaging/{hip_base.postinst => hip-base.postinst} (100%) rename projects/hip/packaging/{hip_base.prerm => hip-base.prerm} (100%) rename projects/hip/packaging/{hip_base.txt => hip-base.txt} (96%) rename projects/hip/packaging/{hip_doc.txt => hip-doc.txt} (89%) rename projects/hip/packaging/{hip_hcc.postinst => hip-hcc.postinst} (100%) rename projects/hip/packaging/{hip_hcc.prerm => hip-hcc.prerm} (100%) create mode 100644 projects/hip/packaging/hip-hcc.txt rename projects/hip/packaging/{hip_nvcc.txt => hip-nvcc.txt} (87%) rename projects/hip/packaging/{hip_samples.txt => hip-samples.txt} (85%) diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index b26ee22d62..5f388f080b 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -421,9 +421,9 @@ endif() ############################# # Package: hip_base set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_base) -configure_file(packaging/hip_base.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) -configure_file(packaging/hip_base.postinst ${BUILD_DIR}/postinst @ONLY) -configure_file(packaging/hip_base.prerm ${BUILD_DIR}/prerm @ONLY) +configure_file(packaging/hip-base.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip-base.postinst ${BUILD_DIR}/postinst @ONLY) +configure_file(packaging/hip-base.prerm ${BUILD_DIR}/prerm @ONLY) add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz @@ -441,9 +441,9 @@ endif() # Package: hip_hcc set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_hcc) -configure_file(packaging/hip_hcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) -configure_file(packaging/hip_hcc.postinst ${BUILD_DIR}/postinst @ONLY) -configure_file(packaging/hip_hcc.prerm ${BUILD_DIR}/prerm @ONLY) +configure_file(packaging/hip-hcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip-hcc.postinst ${BUILD_DIR}/postinst @ONLY) +configure_file(packaging/hip-hcc.prerm ${BUILD_DIR}/prerm @ONLY) add_custom_target(pkg_hip_hcc COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -455,7 +455,7 @@ add_custom_target(pkg_hip_hcc COMMAND ${CMAKE_COMMAND} . # Package: hip_nvcc set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_nvcc) -configure_file(packaging/hip_nvcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip-nvcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) add_custom_target(pkg_hip_nvcc COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -466,7 +466,7 @@ add_custom_target(pkg_hip_nvcc COMMAND ${CMAKE_COMMAND} . # Package: hip_doc set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_doc) -configure_file(packaging/hip_doc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip-doc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) add_custom_target(pkg_hip_doc COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -477,7 +477,7 @@ add_custom_target(pkg_hip_doc COMMAND ${CMAKE_COMMAND} . # Package: hip_samples set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_samples) -configure_file(packaging/hip_samples.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip-samples.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) add_custom_target(pkg_hip_samples COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -486,6 +486,17 @@ add_custom_target(pkg_hip_samples COMMAND ${CMAKE_COMMAND} . COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR} WORKING_DIRECTORY ${BUILD_DIR}) +# Package: hip_hcc_legacy +set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_hcc_legacy) +configure_file(packaging/hip_hcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +add_custom_target(pkg_hip_hcc_legacy COMMAND ${CMAKE_COMMAND} . + COMMAND rm -rf *.deb *.rpm *.tar.gz + COMMAND make package + COMMAND cp *.deb ${PROJECT_BINARY_DIR} + COMMAND cp *.rpm ${PROJECT_BINARY_DIR} + COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR} + WORKING_DIRECTORY ${BUILD_DIR}) + # Package: all if(POLICY CMP0037) cmake_policy(PUSH) @@ -498,7 +509,7 @@ for i in *.rpm ; do mv \$i \${i/.rpm/.x86_64.rpm} ; done add_custom_target(package COMMAND bash ${PROJECT_BINARY_DIR}/fixnames WORKING_DIRECTORY ${PROJECT_BINARY_DIR} - DEPENDS pkg_hip_base pkg_hip_hcc pkg_hip_nvcc pkg_hip_doc pkg_hip_samples) + DEPENDS pkg_hip_base pkg_hip_hcc pkg_hip_nvcc pkg_hip_doc pkg_hip_samples pkg_hip_hcc_legacy) if(POLICY CMP0037) cmake_policy(POP) endif() diff --git a/projects/hip/INSTALL.md b/projects/hip/INSTALL.md index 273c77a23c..2dc12f81a8 100644 --- a/projects/hip/INSTALL.md +++ b/projects/hip/INSTALL.md @@ -24,9 +24,9 @@ HIP code can be developed either on AMD ROCm platform using hcc or clang compile ## HIP-hcc * Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories). -* Install the "hip_hcc" package. This will install HCC and the HIP porting layer. +* Install the "hip-hcc" package. This will install HCC and the HIP porting layer. ``` -apt-get install hip_hcc +apt-get install hip-hcc ``` * Default paths and environment variables: @@ -67,9 +67,9 @@ apt-get install hip_hcc ## HIP-nvcc * Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories). -* Install the "hip_nvcc" package. This will install CUDA SDK and the HIP porting layer. +* Install the "hip-nvcc" package. This will install CUDA SDK and the HIP porting layer. ``` -apt-get install hip_nvcc +apt-get install hip-nvcc ``` * Default paths and environment variables: diff --git a/projects/hip/docker/dockerfile-hip-ubuntu-16.04 b/projects/hip/docker/dockerfile-hip-ubuntu-16.04 index 0852ae38c5..93dd17b948 100644 --- a/projects/hip/docker/dockerfile-hip-ubuntu-16.04 +++ b/projects/hip/docker/dockerfile-hip-ubuntu-16.04 @@ -10,10 +10,10 @@ COPY *.deb /tmp/ # Install the debian package RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install --no-install-recommends -y curl \ && apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install --no-install-recommends --allow-unauthenticated -y \ - /tmp/hip_base-*.deb \ - /tmp/hip_hcc-*.deb \ - /tmp/hip_doc-*.deb \ - /tmp/hip_samples-* \ + /tmp/hip-base-*.deb \ + /tmp/hip-hcc-*.deb \ + /tmp/hip-doc-*.deb \ + /tmp/hip-samples-* \ && rm -f /tmp/*.deb \ && apt-get clean \ && rm -rf /var/lib/apt/lists/* diff --git a/projects/hip/install.sh b/projects/hip/install.sh index c7dc48e1a8..51a303e812 100755 --- a/projects/hip/install.sh +++ b/projects/hip/install.sh @@ -35,8 +35,8 @@ function buildHIP() 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 hip-base*.deb hip-hcc*.deb hip-sample*.deb hip-doc*.deb popd rm -rf $BUILD_ROOT } diff --git a/projects/hip/packaging/hip_base.postinst b/projects/hip/packaging/hip-base.postinst similarity index 100% rename from projects/hip/packaging/hip_base.postinst rename to projects/hip/packaging/hip-base.postinst diff --git a/projects/hip/packaging/hip_base.prerm b/projects/hip/packaging/hip-base.prerm similarity index 100% rename from projects/hip/packaging/hip_base.prerm rename to projects/hip/packaging/hip-base.prerm diff --git a/projects/hip/packaging/hip_base.txt b/projects/hip/packaging/hip-base.txt similarity index 96% rename from projects/hip/packaging/hip_base.txt rename to projects/hip/packaging/hip-base.txt index a44f0acb74..381c8b66d1 100644 --- a/projects/hip/packaging/hip_base.txt +++ b/projects/hip/packaging/hip-base.txt @@ -16,7 +16,7 @@ install(DIRECTORY @hip_SOURCE_DIR@/cmake DESTINATION .) ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip_base") +set(CPACK_PACKAGE_NAME "hip-base") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [BASE]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -29,6 +29,7 @@ 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_REPLACES "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") diff --git a/projects/hip/packaging/hip_doc.txt b/projects/hip/packaging/hip-doc.txt similarity index 89% rename from projects/hip/packaging/hip_doc.txt rename to projects/hip/packaging/hip-doc.txt index 2c1619f221..5f07785ba5 100644 --- a/projects/hip/packaging/hip_doc.txt +++ b/projects/hip/packaging/hip-doc.txt @@ -20,7 +20,7 @@ endif() ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip_doc") +set(CPACK_PACKAGE_NAME "hip-doc") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [DOCUMENTATION]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -31,12 +31,13 @@ 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_REPLACES "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_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/packaging/hip_hcc.postinst b/projects/hip/packaging/hip-hcc.postinst similarity index 100% rename from projects/hip/packaging/hip_hcc.postinst rename to projects/hip/packaging/hip-hcc.postinst diff --git a/projects/hip/packaging/hip_hcc.prerm b/projects/hip/packaging/hip-hcc.prerm similarity index 100% rename from projects/hip/packaging/hip_hcc.prerm rename to projects/hip/packaging/hip-hcc.prerm diff --git a/projects/hip/packaging/hip-hcc.txt b/projects/hip/packaging/hip-hcc.txt new file mode 100644 index 0000000000..0ccf75af50 --- /dev/null +++ b/projects/hip/packaging/hip-hcc.txt @@ -0,0 +1,54 @@ +cmake_minimum_required(VERSION 2.8.3) +project(hip_hcc) + +install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib) +install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib) +if(NOT @HIP_COMPILER@ STREQUAL "clang") + install(FILES @PROJECT_BINARY_DIR@/libhiprtc.so DESTINATION lib) +endif() +install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib) +install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip) +install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/packaging/hip-targets-release.cmake DESTINATION lib/cmake/hip) + +############################# +# Packaging steps +############################# +set(CPACK_SET_DESTDIR TRUE) +set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") +set(CPACK_PACKAGE_NAME "hip-hcc") +if(@HCC_VERSION_MAJOR@ EQUAL 0) + set(HCC_PACKAGE_NAME "hcc_lc") +else() + set(HCC_PACKAGE_NAME "hcc") +endif() +set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [HCC]") +set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") +set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") +set(CPACK_PACKAGE_VERSION @HIP_VERSION_MAJOR@.@HIP_VERSION_MINOR@.@HIP_VERSION_PATCH@) +set(CPACK_PACKAGE_VERSION_MAJOR @HIP_VERSION_MAJOR@) +set(CPACK_PACKAGE_VERSION_MINOR @HIP_VERSION_MINOR@) +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_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)") +else() + set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), comgr (>= 1.1)") +endif() +set(CPACK_DEBIAN_PACKAGE_REPLACES "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") +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") +else() + set(CPACK_RPM_PACKAGE_REQUIRES "hip-base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, comgr >= 1.1") +endif() +set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") +set(CPACK_SOURCE_GENERATOR "TGZ") +include(CPack) diff --git a/projects/hip/packaging/hip_nvcc.txt b/projects/hip/packaging/hip-nvcc.txt similarity index 87% rename from projects/hip/packaging/hip_nvcc.txt rename to projects/hip/packaging/hip-nvcc.txt index 975176f934..11c1c290a7 100644 --- a/projects/hip/packaging/hip_nvcc.txt +++ b/projects/hip/packaging/hip-nvcc.txt @@ -6,7 +6,7 @@ project(hip_nvcc) ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip_nvcc") +set(CPACK_PACKAGE_NAME "hip-nvcc") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [NVCC]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -18,14 +18,15 @@ 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_REPLACES "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_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/packaging/hip_samples.txt b/projects/hip/packaging/hip-samples.txt similarity index 85% rename from projects/hip/packaging/hip_samples.txt rename to projects/hip/packaging/hip-samples.txt index 4c1e04cb23..8a8bd05ceb 100644 --- a/projects/hip/packaging/hip_samples.txt +++ b/projects/hip/packaging/hip-samples.txt @@ -8,7 +8,7 @@ install(DIRECTORY @hip_SOURCE_DIR@/samples DESTINATION .) ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip_samples") +set(CPACK_PACKAGE_NAME "hip-samples") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [SAMPLES]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -19,12 +19,13 @@ 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_REPLACES "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_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/packaging/hip_hcc.txt b/projects/hip/packaging/hip_hcc.txt index c688433cc0..083288d5bd 100644 --- a/projects/hip/packaging/hip_hcc.txt +++ b/projects/hip/packaging/hip_hcc.txt @@ -1,14 +1,5 @@ cmake_minimum_required(VERSION 2.8.3) -project(hip_hcc) - -install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib) -install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib) -if(NOT @HIP_COMPILER@ STREQUAL "clang") - install(FILES @PROJECT_BINARY_DIR@/libhiprtc.so DESTINATION lib) -endif() -install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib) -install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip) -install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/packaging/hip-targets-release.cmake DESTINATION lib/cmake/hip) +project(hip_hcc_legacy) ############################# # Packaging steps @@ -16,11 +7,6 @@ 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") -if(@HCC_VERSION_MAJOR@ EQUAL 0) - set(HCC_PACKAGE_NAME "hcc_lc") -else() - set(HCC_PACKAGE_NAME "hcc") -endif() set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [HCC]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -31,23 +17,12 @@ 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_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)") -else() - set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), comgr (>= 1.1)") -endif() +set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-hcc (= ${CPACK_PACKAGE_VERSION})") 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}) -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") -else() - set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, comgr >= 1.1") -endif() +set(CPACK_RPM_PACKAGE_REQUIRES "hip-hcc = ${HIP_BASE_VERSION}") set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) From 5536940e125372a37527d58847149d450f7ff6bf Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 15 Nov 2019 07:58:40 +0530 Subject: [PATCH 02/27] [ci] Fix failures on nvcc path (#1654) Change-Id: I97f6b51314e14965618ab510bf57a2b034183e78 [ROCm/hip commit: f9fec8fdcaa256624853407f2e9ecaba83aceac4] --- projects/hip/tests/src/hipEnvVar.cpp | 2 +- projects/hip/tests/src/hipHcc.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/tests/src/hipEnvVar.cpp b/projects/hip/tests/src/hipEnvVar.cpp index a0b9503b41..7cf38fc77e 100644 --- a/projects/hip/tests/src/hipEnvVar.cpp +++ b/projects/hip/tests/src/hipEnvVar.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s test_common.cpp + * BUILD: %t %s test_common.cpp NVCC_OPTIONS -std=c++11 * HIT_END */ diff --git a/projects/hip/tests/src/hipHcc.cpp b/projects/hip/tests/src/hipHcc.cpp index 2d4523ecd6..9bcc999b34 100644 --- a/projects/hip/tests/src/hipHcc.cpp +++ b/projects/hip/tests/src/hipHcc.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. // Test the HCC-specific API extensions for HIP: /* HIT_START - * BUILD: %t %s + * BUILD: %t %s EXCLUDE_HIP_PLATFORM nvcc * TEST: %t EXCLUDE_HIP_PLATFORM all * HIT_END */ From 3ea404a8ba74fd626c40bb2ad59f99e04e326dd9 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Sun, 17 Nov 2019 07:24:07 +0100 Subject: [PATCH 03/27] Add support for extended launch syntax. (#1530) * Add support for extended launch syntax. * Add unit test. * Fix typo * hipExtLaunchKernelGGL lives in hip_ext.h Change-Id: Ice32dab0d43475fda65c6a910c11416871a8f2ff * [dtest] remove redundant include from hipModuleGetGlobal dtest [ROCm/hip commit: 69e74c3e96d63e1dd0ddd43cc9f70262a5b8a677] --- .../hip/hcc_detail/functional_grid_launch.hpp | 1 - .../include/hip/hcc_detail/hip_runtime_api.h | 1 + projects/hip/include/hip/hip_ext.h | 60 ++++++++++++++++-- .../src/kernel/hipExtLaunchKernelGGL.cpp | 62 +++++++++++++++++++ .../runtimeApi/module/hipModuleGetGlobal.cpp | 1 - 5 files changed, 119 insertions(+), 6 deletions(-) create mode 100644 projects/hip/tests/src/kernel/hipExtLaunchKernelGGL.cpp diff --git a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp index c493eec933..c9953c27d4 100644 --- a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp @@ -176,5 +176,4 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, numBlocks, dimBlocks, sharedMemBytes, stream, &config[0]); } - #pragma GCC visibility pop diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index ba7f2f4ade..4ae92f1c03 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -2885,6 +2885,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigne unsigned int sharedMemBytes, hipStream_t stream, void** kernelParams, void** extra); + /** * @brief launches kernel f with launch parameters and shared memory on stream with arguments passed * to kernelparams or extra, where thread blocks can cooperate and synchronize as they execute diff --git a/projects/hip/include/hip/hip_ext.h b/projects/hip/include/hip/hip_ext.h index 3df911ee7a..a3b1d3e878 100644 --- a/projects/hip/include/hip/hip_ext.h +++ b/projects/hip/include/hip/hip_ext.h @@ -60,12 +60,11 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a #endif // #ifdef __HCC__ - /** * @brief launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelparams or extra * - * @param [in[ f Kernel to launch. + * @param [in[ f Kernel to launch. * @param [in] gridDimX X grid dimension specified in work-items * @param [in] gridDimY Y grid dimension specified in work-items * @param [in] gridDimZ Z grid dimension specified in work-items @@ -88,7 +87,6 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a * * @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please refer to hip_porting_driver_api.md for sample usage. - * HIP/ROCm actually updates the start event when the associated kernel completes. */ HIP_PUBLIC_API @@ -111,8 +109,62 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, hipEvent_t stopEvent = nullptr) __attribute__((deprecated("use hipExtModuleLaunchKernel instead"))); +#ifdef __cplusplus + +namespace hip_impl { +inline +__attribute__((visibility("hidden"))) +void hipExtLaunchKernelGGLImpl( + std::uintptr_t function_address, + const dim3& numBlocks, + const dim3& dimBlocks, + std::uint32_t sharedMemBytes, + hipStream_t stream, + hipEvent_t startEvent, + hipEvent_t stopEvent, + std::uint32_t flags, + void** kernarg) { + + const auto& kd = hip_impl::get_program_state() + .kernel_descriptor(function_address, target_agent(stream)); + + hipExtModuleLaunchKernel(kd, numBlocks.x * dimBlocks.x, + numBlocks.y * dimBlocks.y, + numBlocks.z * dimBlocks.z, + dimBlocks.x, dimBlocks.y, dimBlocks.z, + sharedMemBytes, stream, nullptr, kernarg, + startEvent, stopEvent, flags); +} +} // namespace hip_impl + +template +inline +void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks, + const dim3& dimBlocks, std::uint32_t sharedMemBytes, + hipStream_t stream, hipEvent_t startEvent, + hipEvent_t stopEvent, std::uint32_t flags, + Args... args) { + hip_impl::hip_init(); + auto kernarg = + hip_impl::make_kernarg(kernel, std::tuple{std::move(args)...}); + std::size_t kernarg_size = kernarg.size(); + + void* config[]{ + HIP_LAUNCH_PARAM_BUFFER_POINTER, + kernarg.data(), + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &kernarg_size, + HIP_LAUNCH_PARAM_END}; + + hip_impl::hipExtLaunchKernelGGLImpl(reinterpret_cast(kernel), + numBlocks, dimBlocks, sharedMemBytes, + stream, startEvent, stopEvent, flags, + &config[0]); +} +#endif + // doxygen end AMD-specific features /** * @} */ -#endif // #ifdef HIP_INCLUDE_HIP_HIP_EXT_H +#endif // #iidef HIP_INCLUDE_HIP_HIP_EXT_H diff --git a/projects/hip/tests/src/kernel/hipExtLaunchKernelGGL.cpp b/projects/hip/tests/src/kernel/hipExtLaunchKernelGGL.cpp new file mode 100644 index 0000000000..4164a87f1e --- /dev/null +++ b/projects/hip/tests/src/kernel/hipExtLaunchKernelGGL.cpp @@ -0,0 +1,62 @@ +/* +Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +// Test the Grid_Launch syntax. + +/* HIT_START + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "hip/hip_ext.h" +#include "test_common.h" + +void test(size_t N) { + size_t Nbytes = N * sizeof(int); + + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + hipExtLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), + dim3(threadsPerBlock), 0, 0, nullptr, nullptr, 0, + static_cast(A_d), static_cast(B_d), C_d, N); + + HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + HIPCHECK(hipDeviceSynchronize()); + + HipTest::checkVectorADD(A_h, B_h, C_h, N); +} + +int main(int argc, char* argv[]) { + HipTest::parseStandardArguments(argc, argv, true); + + test(N); + + passed(); +} diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp index f98d5e4ec4..5896794e90 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp @@ -32,7 +32,6 @@ THE SOFTWARE. #include #include #include -#include #define LEN 64 #define SIZE LEN * sizeof(float) From fac2e14dab018766d00e0419e4dab6c5ec17033f Mon Sep 17 00:00:00 2001 From: kjayapra-amd <54370791+kjayapra-amd@users.noreply.github.com> Date: Mon, 18 Nov 2019 01:48:11 -0500 Subject: [PATCH 04/27] [dtests] Fixing compilation error in saxpy.cpp (#1608) [ROCm/hip commit: 0a68be8b5b6edeae737ca0b9851ace634ab9119b] --- projects/hip/tests/src/hiprtc/saxpy.cpp | 1 + 1 file changed, 1 insertion(+) mode change 100644 => 100755 projects/hip/tests/src/hiprtc/saxpy.cpp diff --git a/projects/hip/tests/src/hiprtc/saxpy.cpp b/projects/hip/tests/src/hiprtc/saxpy.cpp old mode 100644 new mode 100755 index 191e48b846..156a44afe4 --- a/projects/hip/tests/src/hiprtc/saxpy.cpp +++ b/projects/hip/tests/src/hiprtc/saxpy.cpp @@ -35,6 +35,7 @@ THE SOFTWARE. #include #include #include +#include static constexpr auto NUM_THREADS{128}; static constexpr auto NUM_BLOCKS{32}; From cd3bf4e205926d50848b5324a88772e472785b7b Mon Sep 17 00:00:00 2001 From: jiabaxie <57198431+jiabaxie@users.noreply.github.com> Date: Mon, 18 Nov 2019 01:48:29 -0500 Subject: [PATCH 05/27] Updated hipEnvVarDriver to work with Windows (#1614) * Updated hipEnvVarDriver to work with Windows * Cleaned up a bit of code * Fixed a part where putenv was used for both win and linux * Defines moved to test_common.h and cleaned up code * Cleaned up some macro defines and used const char instead * Got rid of some excess commenting * directory paths are unconditional * Cleaned some duplicate code, and variables are now declared and defined together [ROCm/hip commit: e73927caeeb7b40dfd0be42d7a3db0e7762fc6dd] --- projects/hip/tests/src/hipEnvVarDriver.cpp | 51 +++++++++++----------- projects/hip/tests/src/test_common.cpp | 10 ++++- projects/hip/tests/src/test_common.h | 4 ++ 3 files changed, 39 insertions(+), 26 deletions(-) diff --git a/projects/hip/tests/src/hipEnvVarDriver.cpp b/projects/hip/tests/src/hipEnvVarDriver.cpp index 4817d157a4..07379f0878 100644 --- a/projects/hip/tests/src/hipEnvVarDriver.cpp +++ b/projects/hip/tests/src/hipEnvVarDriver.cpp @@ -34,41 +34,43 @@ OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWA using namespace std; +const string directed_dir = "directed_tests" + string(PATH_SEPERATOR_STR) + "hipEnvVar"; +const string dir = "." + string(PATH_SEPERATOR_STR) + "hipEnvVar"; + int getDeviceNumber() { - FILE* in; char buff[512]; - string str; std::this_thread::sleep_for(std::chrono::milliseconds(10)); - if (!(in = popen("./directed_tests/hipEnvVar -c", "r"))) { - // Check at same level - if (!(in = popen("./hipEnvVar -c", "r"))) { + FILE* in = popen((directed_dir + " -c").c_str(), "r"); + if(fgets(buff, 512, in) == NULL){ + pclose(in); + //Check at same level + in = popen((dir + " -c").c_str(), "r"); + if(fgets(buff, 512, in) == NULL){ + pclose(in); return 1; } } - while (fgets(buff, 512, in) != NULL) { - cout << buff; - } + cout << buff; pclose(in); return atoi(buff); } // Query the current device ID remotely to hipEnvVar -void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) { - FILE* in; - string str = "./directed_tests/hipEnvVar -d "; - str += std::to_string(deviceID); +void getDevicePCIBusNumRemote(int deviceID, char* pciBusID) { std::this_thread::sleep_for(std::chrono::milliseconds(10)); - if (!(in = popen(str.c_str(), "r"))) { - // Check at same level - if (!(in = popen("./hipEnvVar -d ", "r"))) { - exit(1); + FILE* in = popen((directed_dir + " -d " + std::to_string(deviceID)).c_str(), "r"); + if(fgets(pciBusID, 100, in) == NULL){ + pclose(in); + //Check at same level + in = popen((dir + " -d").c_str(), "r"); + if(fgets(pciBusID, 100, in) == NULL){ + pclose(in); + return; } - - } - while (fgets(pciBusID, 100, in) != NULL) { - cout << pciBusID; } + cout << pciBusID; pclose(in); + return; } // Query the current device ID locally on AMD path @@ -81,9 +83,8 @@ void getDevicePCIBusNum(int deviceID, char* pciBusID) { } int main() { - unsetenv("HIP_VISIBLE_DEVICES"); - unsetenv("CUDA_VISIBLE_DEVICES"); - + unsetenv(HIP_VISIBLE_DEVICES_STR); + unsetenv(CUDA_VISIBLE_DEVICES_STR); std::vector devPCINum; char pciBusID[100]; // collect the device pci bus ID for all devices @@ -126,8 +127,8 @@ int main() { setenv("CUDA_VISIBLE_DEVICES", "0,1,2", 1); assert(getDeviceNumber() == 3); // test if CUDA_VISIBLE_DEVICES will be accepted by the runtime - unsetenv("HIP_VISIBLE_DEVICES"); - unsetenv("CUDA_VISIBLE_DEVICES"); + unsetenv(HIP_VISIBLE_DEVICES_STR); + unsetenv(CUDA_VISIBLE_DEVICES_STR); setenv("CUDA_VISIBLE_DEVICES", "0,1,2", 1); assert(getDeviceNumber() == 3); } diff --git a/projects/hip/tests/src/test_common.cpp b/projects/hip/tests/src/test_common.cpp index 0205310351..e7a2622662 100644 --- a/projects/hip/tests/src/test_common.cpp +++ b/projects/hip/tests/src/test_common.cpp @@ -33,7 +33,15 @@ unsigned threadsPerBlock = 256; int p_gpuDevice = 0; unsigned p_verbose = 0; int p_tests = -1; /*which tests to run. Interpretation is left to each test. default:all*/ - +#ifdef _WIN64 +const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES="; +const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES="; +const char* PATH_SEPERATOR_STR = "\\"; +#else +const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES"; +const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES"; +const char* PATH_SEPERATOR_STR = "/"; +#endif namespace HipTest { diff --git a/projects/hip/tests/src/test_common.h b/projects/hip/tests/src/test_common.h index 92528e38e3..426ea846b1 100644 --- a/projects/hip/tests/src/test_common.h +++ b/projects/hip/tests/src/test_common.h @@ -104,6 +104,7 @@ THE SOFTWARE. #define popen(x,y) _popen(x,y) #define pclose(x) _pclose(x) #define setenv(x,y,z) _putenv_s(x,y) +#define unsetenv _putenv #else #define aligned_free(x) free(x) #endif @@ -120,6 +121,9 @@ extern unsigned threadsPerBlock; extern int p_gpuDevice; extern unsigned p_verbose; extern int p_tests; +extern const char* HIP_VISIBLE_DEVICES_STR; +extern const char* CUDA_VISIBLE_DEVICES_STR; +extern const char* PATH_SEPERATOR_STR; // ********************* CPP section ********************* #ifdef __cplusplus From a4f2c7717f9faa8340a2a7c41b0541219db207a0 Mon Sep 17 00:00:00 2001 From: Sarbojit2019 <52527887+SarbojitAMD@users.noreply.github.com> Date: Mon, 18 Nov 2019 12:18:55 +0530 Subject: [PATCH 06/27] Added null check in hipEventSynchronize (#1627) * Added missing null check in hipEventSynchronize * Minor correction in the Event API description [ROCm/hip commit: b865a50e44443b9f306b2b090d09cd6f3605c5d8] --- projects/hip/include/hip/hcc_detail/hip_runtime_api.h | 4 ++-- projects/hip/src/hip_event.cpp | 10 +++++----- 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 4ae92f1c03..afb99da6a3 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -883,7 +883,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback * #hipEventDefault : Default flag. The event will use active synchronization and will support timing. Blocking synchronization provides lowest possible latency at the expense of dedicating a - CPU to poll on the eevent. + CPU to poll on the event. * #hipEventBlockingSync : The event will use blocking synchronization : if hipEventSynchronize is called on this event, the thread will block until the event completes. This can increase latency for the synchroniation but can result in lower power and more resources for other CPU threads. @@ -935,7 +935,7 @@ hipError_t hipEventCreate(hipEvent_t* event); * If hipEventRecord() has been previously called on this event, then this call will overwrite any * existing state in event. * - * If this function is called on a an event that is currently being recorded, results are undefined + * If this function is called on an event that is currently being recorded, results are undefined * - either outstanding recording may save state into the event, and the order is not guaranteed. * * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventSynchronize, diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp index dcab059774..3d1a366463 100644 --- a/projects/hip/src/hip_event.cpp +++ b/projects/hip/src/hip_event.cpp @@ -143,14 +143,14 @@ hipError_t hipEventDestroy(hipEvent_t event) { hipError_t hipEventSynchronize(hipEvent_t event) { HIP_INIT_SPECIAL_API(hipEventSynchronize, TRACE_SYNC, event); - if (!(event->_flags & hipEventReleaseToSystem)) { - tprintf(DB_WARN, + if (event){ + if (!(event->_flags & hipEventReleaseToSystem)) { + tprintf(DB_WARN, "hipEventSynchronize on event without system-scope fence ; consider creating with " "hipEventReleaseToSystem\n"); - } - auto ecd = event->locked_copyCrit(); + } + auto ecd = event->locked_copyCrit(); - if (event) { if (ecd._state == hipEventStatusUnitialized) { return ihipLogStatus(hipErrorInvalidResourceHandle); } else if (ecd._state == hipEventStatusCreated) { From 05ef9bd386e03951997f59f2f43181d966e055d5 Mon Sep 17 00:00:00 2001 From: Nick Curtis Date: Mon, 18 Nov 2019 00:49:12 -0600 Subject: [PATCH 07/27] fix complex conjugate for double-complex (#1659) The sign in the y component returned from hipConj incorrect for double-complex. Fix to match as in hipConjf above. [ROCm/hip commit: 3f2316086f652deb14bf8e17b53c87fd587da397] --- projects/hip/include/hip/hcc_detail/hip_complex.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_complex.h b/projects/hip/include/hip/hcc_detail/hip_complex.h index 4acee9e235..11648ce123 100644 --- a/projects/hip/include/hip/hcc_detail/hip_complex.h +++ b/projects/hip/include/hip/hcc_detail/hip_complex.h @@ -186,7 +186,7 @@ __device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double __device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) { hipDoubleComplex ret; ret.x = z.x; - ret.y = z.y; + ret.y = -z.y; return ret; } From ecdba45d291e31a265a3c19e4b9323671577e093 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sun, 17 Nov 2019 22:49:22 -0800 Subject: [PATCH 08/27] Fix gcc build on NVCC path (#1661) * Fix gcc build on NVCC path * Fix CI build errors * [dtest] Fix texture and surface obj2D tests [ROCm/hip commit: e39d7497eca7c8ab50a772416a0e60adc760299a] --- .../include/hip/nvcc_detail/hip_runtime_api.h | 176 +++++++++--------- projects/hip/tests/src/gcc/hipMalloc.c | 6 +- .../hip/tests/src/surface/hipSurfaceObj2D.cpp | 4 +- .../hip/tests/src/texture/hipTextureObj2D.cpp | 4 +- 4 files changed, 95 insertions(+), 95 deletions(-) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index d2f81c8b7b..9cb59f14ea 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -66,24 +66,24 @@ typedef enum hipMemcpyKind { #define HIP_LIBRARY_PATCH_LEVEL PATCH_LEVEL // hipTextureAddressMode -#define hipTextureAddressMode cudaTextureAddressMode +typedef enum cudaTextureAddressMode hipTextureAddressMode; #define hipAddressModeWrap cudaAddressModeWrap #define hipAddressModeClamp cudaAddressModeClamp #define hipAddressModeMirror cudaAddressModeMirror #define hipAddressModeBorder cudaAddressModeBorder // hipTextureFilterMode -#define hipTextureFilterMode cudaTextureFilterMode +typedef enum cudaTextureFilterMode hipTextureFilterMode; #define hipFilterModePoint cudaFilterModePoint #define hipFilterModeLinear cudaFilterModeLinear // hipTextureReadMode -#define hipTextureReadMode cudaTextureReadMode +typedef enum cudaTextureReadMode hipTextureReadMode; #define hipReadModeElementType cudaReadModeElementType #define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat // hipChannelFormatKind -#define hipChannelFormatKind cudaChannelFormatKind +typedef enum cudaChannelFormatKind hipChannelFormatKind; #define hipChannelFormatKindSigned cudaChannelFormatKindSigned #define hipChannelFormatKindUnsigned cudaChannelFormatKindUnsigned #define hipChannelFormatKindFloat cudaChannelFormatKindFloat @@ -167,7 +167,7 @@ typedef cudaIpcMemHandle_t hipIpcMemHandle_t; typedef enum cudaLimit hipLimit_t; typedef enum cudaFuncCache hipFuncCache_t; typedef CUcontext hipCtx_t; -typedef cudaSharedMemConfig hipSharedMemConfig; +typedef enum cudaSharedMemConfig hipSharedMemConfig; typedef CUfunc_cache hipFuncCache; typedef CUjit_option hipJitOption; typedef CUdevice hipDevice_t; @@ -177,8 +177,8 @@ typedef CUdeviceptr hipDeviceptr_t; typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_t; typedef struct cudaArray* hipArray_const_t; -typedef cudaFuncAttributes hipFuncAttributes; -typedef CUfunction_attribute hipFunction_attribute; +typedef struct cudaFuncAttributes hipFuncAttributes; +#define hipFunction_attribute CUfunction_attribute #define hip_Memcpy2D CUDA_MEMCPY2D #define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault @@ -196,8 +196,8 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t; #define hipTextureType3D cudaTextureType3D #define hipDeviceMapHost cudaDeviceMapHost -#define hipExtent cudaExtent -#define hipPitchedPtr cudaPitchedPtr +typedef struct cudaExtent hipExtent; +typedef struct cudaPitchedPtr hipPitchedPtr; #define make_hipExtent make_cudaExtent #define make_hipPos make_cudaPos #define make_hipPitchedPtr make_cudaPitchedPtr @@ -205,10 +205,10 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t; #define hipStreamDefault cudaStreamDefault #define hipStreamNonBlocking cudaStreamNonBlocking -#define hipChannelFormatDesc cudaChannelFormatDesc -#define hipResourceDesc cudaResourceDesc -#define hipTextureDesc cudaTextureDesc -#define hipResourceViewDesc cudaResourceViewDesc +typedef struct cudaChannelFormatDesc hipChannelFormatDesc; +typedef struct cudaResourceDesc hipResourceDesc; +typedef struct cudaTextureDesc hipTextureDesc; +typedef struct cudaResourceViewDesc hipResourceViewDesc; // adding code for hipmemSharedConfig #define hipSharedMemBankSizeDefault cudaSharedMemBankSizeDefault #define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte @@ -394,7 +394,7 @@ inline static enum cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind ki } } -inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode( +inline static enum cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode( hipTextureAddressMode kind) { switch (kind) { case hipAddressModeWrap: @@ -410,7 +410,7 @@ inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMo } } -inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode( +inline static enum cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode( hipTextureFilterMode kind) { switch (kind) { case hipFilterModePoint: @@ -422,7 +422,7 @@ inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode( } } -inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) { +inline static enum cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) { switch (kind) { case hipReadModeElementType: return cudaReadModeElementType; @@ -433,7 +433,7 @@ inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTex } } -inline static cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind( +inline static enum cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind( hipChannelFormatKind kind) { switch (kind) { case hipChannelFormatKindSigned: @@ -510,14 +510,14 @@ inline static hipError_t hipMallocManaged(void** ptr, size_t size, unsigned int return hipCUDAErrorTohipError(cudaMallocManaged(ptr, size, flags)); } -inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, +inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags __dparm(hipArrayDefault)) { return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } -inline static hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc, - struct hipExtent extent, unsigned int flags) { +inline static hipError_t hipMalloc3DArray(hipArray** array, const hipChannelFormatDesc* desc, + hipExtent extent, unsigned int flags) { return hipCUDAErrorTohipError(cudaMalloc3DArray(array, desc, extent, flags)); } @@ -1048,7 +1048,6 @@ inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attribut return err; } - inline static hipError_t hipMemGetInfo(size_t* free, size_t* total) { return hipCUDAErrorTohipError(cudaMemGetInfo(free, total)); } @@ -1073,7 +1072,6 @@ inline static hipError_t hipEventDestroy(hipEvent_t event) { return hipCUDAErrorTohipError(cudaEventDestroy(event)); } - inline static hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) { return hipCUDAErrorTohipError(cudaStreamCreateWithFlags(stream, flags)); } @@ -1360,79 +1358,15 @@ inline static hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gri kernelParams, extra)); } - inline static hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t cacheConfig) { return hipCUDAErrorTohipError(cudaFuncSetCacheConfig(func, cacheConfig)); } -#ifdef __cplusplus -} -#endif - -#ifdef __CUDACC__ - -template -inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int* blockSize, T func, - size_t dynamicSMemSize = 0, - int blockSizeLimit = 0) { - cudaError_t cerror; - cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit); - return hipCUDAErrorTohipError(cerror); -} - -template -inline static hipError_t hipBindTexture(size_t* offset, const struct texture& tex, - const void* devPtr, size_t size = UINT_MAX) { - return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size)); -} - -template -inline static hipError_t hipBindTexture(size_t* offset, struct texture& tex, - const void* devPtr, const struct hipChannelFormatDesc& desc, - size_t size = UINT_MAX) { +inline static hipError_t hipBindTexture(size_t* offset, struct textureReference* tex, const void* devPtr, + const hipChannelFormatDesc* desc, size_t size __dparm(UINT_MAX)){ return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); } -template -inline static hipError_t hipUnbindTexture(struct texture* tex) { - return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); -} - -template -inline static hipError_t hipUnbindTexture(struct texture &tex) { - return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); -} - -inline static hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr, - const hipChannelFormatDesc* desc, size_t size = UINT_MAX){ - return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); -} - -template -inline static hipError_t hipBindTextureToArray(struct texture& tex, - hipArray_const_t array, - const struct hipChannelFormatDesc& desc) { - return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); -} - -template -inline static hipError_t hipBindTextureToArray(struct texture *tex, - hipArray_const_t array, - const struct hipChannelFormatDesc* desc) { - return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); -} - -template -inline static hipError_t hipBindTextureToArray(struct texture& tex, - hipArray_const_t array) { - return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array)); -} - -template -inline static hipChannelFormatDesc hipCreateChannelDesc() { - return cudaCreateChannelDesc(); -} - inline static hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { return cudaCreateChannelDesc(x, y, z, w, hipChannelFormatKindToCudaChannelFormatKind(f)); @@ -1464,7 +1398,7 @@ inline static hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDe return hipCUDAErrorTohipError(cudaGetTextureObjectResourceDesc( pResDesc, textureObject)); } -inline static hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* texref) +inline static hipError_t hipGetTextureAlignmentOffset(size_t* offset, const struct textureReference* texref) { return hipCUDAErrorTohipError(cudaGetTextureAlignmentOffset(offset,texref)); } @@ -1473,6 +1407,70 @@ inline static hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_ { return hipCUDAErrorTohipError(cudaGetChannelDesc(desc,array)); } + + +#ifdef __cplusplus +} +#endif + +#ifdef __CUDACC__ + +template +inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int* blockSize, T func, + size_t dynamicSMemSize = 0, + int blockSizeLimit = 0) { + cudaError_t cerror; + cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit); + return hipCUDAErrorTohipError(cerror); +} + +template +inline static hipError_t hipBindTexture(size_t* offset, const struct texture& tex, + const void* devPtr, size_t size = UINT_MAX) { + return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size)); +} + +template +inline static hipError_t hipBindTexture(size_t* offset, struct texture& tex, + const void* devPtr, const hipChannelFormatDesc& desc, + size_t size = UINT_MAX) { + return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); +} + +template +inline static hipError_t hipUnbindTexture(struct texture* tex) { + return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); +} + +template +inline static hipError_t hipUnbindTexture(struct texture &tex) { + return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); +} + +template +inline static hipError_t hipBindTextureToArray(struct texture& tex, + hipArray_const_t array, + const hipChannelFormatDesc& desc) { + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); +} + +template +inline static hipError_t hipBindTextureToArray(struct texture *tex, + hipArray_const_t array, + const hipChannelFormatDesc* desc) { + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); +} + +template +inline static hipError_t hipBindTextureToArray(struct texture& tex, + hipArray_const_t array) { + return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array)); +} + +template +inline static hipChannelFormatDesc hipCreateChannelDesc() { + return cudaCreateChannelDesc(); +} #endif //__CUDACC__ #endif // HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H diff --git a/projects/hip/tests/src/gcc/hipMalloc.c b/projects/hip/tests/src/gcc/hipMalloc.c index e3c7a66919..19e90d5222 100644 --- a/projects/hip/tests/src/gcc/hipMalloc.c +++ b/projects/hip/tests/src/gcc/hipMalloc.c @@ -18,8 +18,10 @@ * */ /* HIT_START - * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t EXCLUDE_HIP_PLATFORM nvcc - * TEST: %t EXCLUDE_HIP_PLATFORM nvcc + * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_NVCC__ -I%hip-path/include -I/usr/local/cuda/include %S/%s -o %T/hipMalloc_nv -L/usr/local/cuda/lib64 -lcudart EXCLUDE_HIP_PLATFORM hcc + * BUILD_CMD: hipMalloc %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc + * TEST: hipMalloc_nv EXCLUDE_HIP_PLATFORM hcc + * TEST: hipMalloc_hcc EXCLUDE_HIP_PLATFORM nvcc * HIT_END */ diff --git a/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp b/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp index b6c016ecc0..4580220d1d 100644 --- a/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp +++ b/projects/hip/tests/src/surface/hipSurfaceObj2D.cpp @@ -53,7 +53,7 @@ int runTest(int argc, char** argv) { hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); - struct hipResourceDesc resDesc; + hipResourceDesc resDesc; memset(&resDesc, 0, sizeof(resDesc)); resDesc.resType = hipResourceTypeArray; resDesc.res.array.array = hipArray; @@ -62,7 +62,7 @@ int runTest(int argc, char** argv) { hipCreateSurfaceObject(&surfaceObject, &resDesc); hipMallocArray(&hipOutArray, &channelDesc, width, height); - struct hipResourceDesc resOutDesc; + hipResourceDesc resOutDesc; memset(&resOutDesc, 0, sizeof(resOutDesc)); resOutDesc.resType = hipResourceTypeArray; resOutDesc.res.array.array = hipOutArray; diff --git a/projects/hip/tests/src/texture/hipTextureObj2D.cpp b/projects/hip/tests/src/texture/hipTextureObj2D.cpp index 148c966e80..649c748fd2 100644 --- a/projects/hip/tests/src/texture/hipTextureObj2D.cpp +++ b/projects/hip/tests/src/texture/hipTextureObj2D.cpp @@ -53,13 +53,13 @@ int runTest(int argc, char** argv) { hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice); - struct hipResourceDesc resDesc; + hipResourceDesc resDesc; memset(&resDesc, 0, sizeof(resDesc)); resDesc.resType = hipResourceTypeArray; resDesc.res.array.array = hipArray; // Specify texture object parameters - struct hipTextureDesc texDesc; + hipTextureDesc texDesc; memset(&texDesc, 0, sizeof(texDesc)); texDesc.addressMode[0] = hipAddressModeWrap; texDesc.addressMode[1] = hipAddressModeWrap; From 97646b52f1a5a917eb6516192ea2ae4c5a68c2d0 Mon Sep 17 00:00:00 2001 From: mhbliao <47895780+mhbliao@users.noreply.github.com> Date: Mon, 18 Nov 2019 11:18:20 -0500 Subject: [PATCH 09/27] Fix mathlib and app builds with hip-clang. (#1665) [ROCm/hip commit: ebe0c56f4f700863a85984ccfbaf4c4d53e2805c] --- projects/hip/include/hip/hip_ext.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hip_ext.h b/projects/hip/include/hip/hip_ext.h index a3b1d3e878..5f6ba31d1e 100644 --- a/projects/hip/include/hip/hip_ext.h +++ b/projects/hip/include/hip/hip_ext.h @@ -109,7 +109,7 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, hipEvent_t stopEvent = nullptr) __attribute__((deprecated("use hipExtModuleLaunchKernel instead"))); -#ifdef __cplusplus +#if !__HIP_VDI__ && defined(__cplusplus) namespace hip_impl { inline @@ -161,7 +161,7 @@ void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks, stream, startEvent, stopEvent, flags, &config[0]); } -#endif +#endif // !__HIP_VDI__ && defined(__cplusplus) // doxygen end AMD-specific features /** From cfaf04a467ef8b2a8781b516930118023feb24ee Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <41661541+ssahasra@users.noreply.github.com> Date: Tue, 19 Nov 2019 09:22:40 +0530 Subject: [PATCH 10/27] hostcall: add tests that are only used by the VDI build (#1664) [ROCm/hip commit: 363caf29c6c74dcb2dfd5d58e8e14b26eb67cac3] --- .../src/hostcall/hipHostcallFuncCall.cpp | 106 +++++++++++++++ .../src/hostcall/hipHostcallPrintThings.cpp | 125 ++++++++++++++++++ 2 files changed, 231 insertions(+) create mode 100644 projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp create mode 100644 projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp diff --git a/projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp b/projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp new file mode 100644 index 0000000000..4b73352bb7 --- /dev/null +++ b/projects/hip/tests/src/hostcall/hipHostcallFuncCall.cpp @@ -0,0 +1,106 @@ +/* +Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM all + * HIT_END + */ + +#include + +extern "C" __device__ HIP_vector_base::Native_vec_ __ockl_call_host_function( + ulong fptr, ulong arg0, ulong arg1, ulong arg2, ulong arg3, ulong arg4, ulong arg5, ulong arg6); + +static void callee(ulong* output, ulong* input) { + output[0] = input[0] + 1; + output[1] = input[1] + input[2]; +} + +__global__ void kernel(ulong fptr, ulong* retval0, ulong* retval1) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + ulong arg0 = (ulong)fptr; + ulong arg1 = tid; + ulong arg2 = 42; + ulong arg3 = tid % 23; + ulong arg4 = 0; + ulong arg5 = 0; + ulong arg6 = 0; + ulong arg7 = 0; + + long2 result = {0, 0}; + if (tid % 71 != 1) { + result.data = __ockl_call_host_function(arg0, arg1, arg2, arg3, arg4, arg5, arg6, arg7); + retval0[tid] = result.x; + retval1[tid] = result.y; + } +} + +static bool test() { + uint num_blocks = 5; + uint threads_per_block = 1000; + uint num_threads = num_blocks * threads_per_block; + + void* retval0_void; + HIPCHECK(hipHostMalloc(&retval0_void, 8 * num_threads)); + auto retval0 = reinterpret_cast(retval0_void); + for (uint i = 0; i != num_threads; ++i) { + retval0[i] = 0x23232323; + } + + void* retval1_void; + HIPCHECK(hipHostMalloc(&retval1_void, 8 * num_threads)); + auto retval1 = reinterpret_cast(retval1_void); + for (uint i = 0; i != num_threads; ++i) { + retval1[i] = 0x23232323; + } + + hipLaunchKernelGGL(kernel, dim3(num_blocks), dim3(threads_per_block), 0, 0, (ulong)callee, + retval0, retval1); + + hipStreamSynchronize(0); + + for (uint i = 0; i != num_threads; ++i) { + ulong value = retval0[i]; + if (i % 71 == 1) { + if (value != 0x23232323) { + printf("failed\n"); + return false; + } + } else { + if (value != i + 1) { + printf("failed\n"); + return false; + } + } + } + + return true; +} + +int main(int argc, char** argv) { + if (!test()) { + return 1; + } + + printf("passed\n"); + return 0; +} diff --git a/projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp b/projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp new file mode 100644 index 0000000000..742798c0b9 --- /dev/null +++ b/projects/hip/tests/src/hostcall/hipHostcallPrintThings.cpp @@ -0,0 +1,125 @@ +/* +Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM all + * HIT_END + */ + +#include + +// This is NOT a real printf test. It is a test for calling a host function +// which happens to be a wrapper around system printf. + +extern "C" __device__ HIP_vector_base::Native_vec_ __ockl_call_host_function( + ulong fptr, ulong arg0, ulong arg1, ulong arg2, ulong arg3, ulong arg4, ulong arg5, ulong arg6); + +// FuncCall service function that expects three arguments bundled in the +// request: the format string, and two uint64_t arguments. +void print_things_0(ulong* output, ulong* input) { + auto fmt = reinterpret_cast(input); + auto arg0 = input[2]; + auto arg1 = input[3]; + output[0] = fprintf(stdout, fmt, arg0, arg1); +} + +__global__ void kernel0(ulong fptr, ulong* retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + ulong arg0 = fptr; + + const char* str = "(%lu -> %lu)\n"; + ulong arg1 = 0; + for (int ii = 0; ii != 8; ++ii) { + arg1 |= (ulong)str[ii] << (8 * ii); + } + ulong arg2 = 0; + for (int ii = 0; ii != 7; ++ii) { + arg2 |= (ulong)str[ii + 8] << (8 * ii); + } + + ulong arg3 = 42; + ulong arg4 = tid; + ulong arg5 = 0; + ulong arg6 = 0; + ulong arg7 = 0; + + long2 result = {0, 0}; + result.data = __ockl_call_host_function(arg0, arg1, arg2, arg3, arg4, arg5, arg6, arg7); + *retval = result.x; +} + +// FuncCall service function that expects two arguments bundled in the request: +// a kernel "name" and a uint64_t thread ID. The format string is built into the +// service function itself. +void print_things_1(ulong* output, const ulong* input) { + auto name = reinterpret_cast(input[0]); + auto tid = input[1]; + output[0] = fprintf(stdout, "kernel: %s; tid: %lu\n", name, tid); +} + +__global__ void kernel1(ulong fptr, ulong name, ulong* retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + ulong arg0 = fptr; + ulong arg1 = name; + ulong arg2 = tid; + ulong arg3 = 0; + ulong arg4 = 0; + ulong arg5 = 0; + ulong arg6 = 0; + ulong arg7 = 0; + + long2 result = {0, 0}; + result.data = __ockl_call_host_function(arg0, arg1, arg2, arg3, arg4, arg5, arg6, arg7); + *retval = result.x; +} + +static bool test() { + void* retval_void; + HIPCHECK(hipHostMalloc(&retval_void, 8)); + auto retval = reinterpret_cast(retval_void); + *retval = 0x23232323; + + hipLaunchKernelGGL(kernel0, dim3(1), dim3(1), 0, 0, (ulong)print_things_0, retval); + hipStreamSynchronize(0); + if (*retval != strlen("(42 -> 0)\n")) { + return false; + } + + *retval = 0x23232323; + const char* name = "kernel1"; + hipLaunchKernelGGL(kernel1, dim3(1), dim3(1), 0, 0, (ulong)print_things_1, (ulong)name, retval); + hipStreamSynchronize(0); + if (*retval != strlen("kernel: kernel1; tid: 0\n")) { + return false; + } + + return true; +} + +int main(int argc, char** argv) { + if (!test()) { + printf("failed\n"); + return 1; + } + printf("passed\n"); + return 0; +} From f55b1cd6a4b553984d2783617893607c1c172935 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 19 Nov 2019 09:41:13 -0800 Subject: [PATCH 11/27] [dtest] Rename hipExtLaunchMultiKernelMultiDevice dtest (#1671) [ROCm/hip commit: 188752b2cb9326ac6cdba78f2d2efe2ba82071ef] --- ...rnelMultiDevice.cpp => hipExtLaunchMultiKernelMultiDevice.cpp} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename projects/hip/tests/src/runtimeApi/module/{hipLaunchMultiKernelMultiDevice.cpp => hipExtLaunchMultiKernelMultiDevice.cpp} (100%) diff --git a/projects/hip/tests/src/runtimeApi/module/hipLaunchMultiKernelMultiDevice.cpp b/projects/hip/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp similarity index 100% rename from projects/hip/tests/src/runtimeApi/module/hipLaunchMultiKernelMultiDevice.cpp rename to projects/hip/tests/src/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice.cpp From 492248ef12a6f5ef1b6e993baa25eda04afd330b Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Tue, 19 Nov 2019 09:42:43 -0800 Subject: [PATCH 12/27] hipStreamSynchronize can skip marker if stream is empty (#1667) [ROCm/hip commit: 3a7eb694f508517b3389ed8a9295bf41f8d16685] --- projects/hip/src/hip_hcc.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 377ce0aad4..cc83ca0657 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -310,6 +310,8 @@ void ihipStream_t::locked_wait() { hc::completion_future marker; { LockedAccessor_StreamCrit_t crit(_criticalData); + // skipping marker since stream is empty + if (crit->_av.get_is_empty()) return; marker = crit->_av.create_marker(hc::no_scope); } From 6679cd39984eebe82dc1f3aee39d2bbd993460be Mon Sep 17 00:00:00 2001 From: Sarbojit2019 <52527887+SarbojitAMD@users.noreply.github.com> Date: Wed, 20 Nov 2019 11:55:46 +0530 Subject: [PATCH 13/27] Revert [HIP] Fixed hipStreamAddCallback (#1674) This reverts commit fa1e44aa0e201c34744da84348d949e3f85bf052. Addresses SWDEV#212675. [ROCm/hip commit: 153a9592805cf4072bc6560c72e3bbbd90f52448] --- projects/hip/src/hip_hcc.cpp | 15 +- projects/hip/src/hip_hcc_internal.h | 4 +- projects/hip/src/hip_stream.cpp | 127 +------------ .../runtimeApi/stream/StreamAddCallback.cpp | 145 -------------- .../hipStreamAddCallbackCrossStream.cpp | 133 ------------- .../hipStreamAddCallbackMultiThread.cpp | 178 ------------------ 6 files changed, 9 insertions(+), 593 deletions(-) delete mode 100644 projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp delete mode 100644 projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp delete mode 100644 projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index cc83ca0657..e93490a08e 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -1499,21 +1499,18 @@ hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream) { return e; } -bool ihipStreamCallbackHandler(hsa_signal_value_t value, void* cbArgs) { +void ihipStreamCallbackHandler(ihipStreamCallback_t* cb) { hipError_t e = hipSuccess; - ihipStreamCallback_t* cb = static_cast (cbArgs); - - if(cb->comFuture.valid()) - cb->comFuture.wait(); + // Synchronize stream + tprintf(DB_SYNC, "ihipStreamCallbackHandler wait on stream %s\n", + ToString(cb->_stream).c_str()); + GET_TLS(); + e = ihipStreamSynchronize(tls, cb->_stream); // Call registered callback function cb->_callback(cb->_stream, e, cb->_userData); - - hsa_signal_store_screlease(cb->_signal,0); - delete cb; - return false; } //--- diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h index 6337edd370..8ee47eba4a 100644 --- a/projects/hip/src/hip_hcc_internal.h +++ b/projects/hip/src/hip_hcc_internal.h @@ -650,9 +650,7 @@ class ihipStreamCallback_t { : _stream(stream), _callback(callback), _userData(userData) { }; hipStream_t _stream; - hsa_signal_t _signal; hipStreamCallback_t _callback; - hc::completion_future comFuture; void* _userData; }; @@ -970,7 +968,7 @@ hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod, hipStream_t ihipSyncAndResolveStream(hipStream_t, bool lockAcquired = 0); hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream); -bool ihipStreamCallbackHandler(hsa_signal_value_t value, void* cbArgs); +void ihipStreamCallbackHandler(ihipStreamCallback_t* cb); // Stream printf functions: inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) { diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index bf3289bb49..e3e4975b7e 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -254,137 +254,14 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) { //--- -void setCallbackPacket(hsa_queue_t* queue, - uint64_t& index, uint64_t& nextIndex, - hsa_barrier_and_packet_t** barrier1, - hsa_barrier_and_packet_t** barrier2){ - - uint64_t tempIndex = 0; - uint32_t mask = queue->size - 1; - hsa_barrier_and_packet_t* tempBarrier; - - // Check for empty packets - do{ - tempIndex = hsa_queue_load_write_index_scacquire(queue); - tempBarrier = &(((hsa_barrier_and_packet_t*)(queue->base_address))[tempIndex & mask]); - }while(!(tempBarrier->header & HSA_PACKET_TYPE_INVALID)); - - // Reserve two packets for two barriers - index = hsa_queue_add_write_index_scacquire(queue, 2); - - if(index > mask) - { - index = 0; - nextIndex = 1; - } - else if(index == mask) - nextIndex = 0; - else - nextIndex = index + 1; - - tempBarrier = new hsa_barrier_and_packet_t; - memset(tempBarrier, 0, sizeof(hsa_barrier_and_packet_t)); - tempBarrier->header = HSA_PACKET_TYPE_INVALID; - - // Barrier 1 - *barrier1 = &(((hsa_barrier_and_packet_t*)(queue->base_address))[index & mask]); - memcpy(*barrier1,tempBarrier,sizeof(hsa_barrier_and_packet_t)); - - // Barrier 2 - *barrier2 = &(((hsa_barrier_and_packet_t*)(queue->base_address))[nextIndex & mask]); - memcpy(*barrier2,tempBarrier,sizeof(hsa_barrier_and_packet_t)); - - delete tempBarrier; -} - hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, unsigned int flags) { - HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags); hipError_t e = hipSuccess; - if(stream == hipStreamNull) - { - ihipCtx_t* device = ihipGetTlsDefaultCtx(); - stream = device->_defaultStream; - } - - stream = ihipSyncAndResolveStream(stream); - - // Lock the stream - LockedAccessor_StreamCrit_t crit(stream->criticalData()); - - // Device synchronization - hc::completion_future marker = crit->_av.create_marker(hc::system_scope); - - // 1. Lock the queue - hsa_queue_t* lockedQ = static_cast (crit->_av.acquire_locked_hsa_queue()); - - if(lockedQ == nullptr) - { - // No queue attached to stream hence exiting early - return ihipLogStatus(hipErrorMissingConfiguration); - } - - // 2. Allocate a singals - hsa_signal_t signal; - hsa_status_t status = hsa_signal_create(1, 0, NULL, &signal); - - if(status != HSA_STATUS_SUCCESS) - { - crit->_av.release_locked_hsa_queue(); - return ihipLogStatus(hipErrorInvalidValue); - } - - hsa_signal_t depSignal; - status = hsa_signal_create(1, 0, NULL, &depSignal); - - if(status != HSA_STATUS_SUCCESS) - { - crit->_av.release_locked_hsa_queue(); - return ihipLogStatus(hipErrorInvalidValue); - } - - // 3. Store callback details, will destroy allocation in callback handler + // Create a thread in detached mode to handle callback ihipStreamCallback_t* cb = new ihipStreamCallback_t(stream, callback, userData); - if(cb == nullptr) - { - crit->_av.release_locked_hsa_queue(); - return ihipLogStatus(hipErrorMemoryAllocation); - } - cb->_signal = depSignal; - cb->comFuture = marker ; - - // 4. Create barrier packets - uint64_t index ; - uint64_t nextIndex; - - hsa_barrier_and_packet_t* barrier; - hsa_barrier_and_packet_t* depBarrier; - - setCallbackPacket(lockedQ, index, nextIndex, &barrier, &depBarrier); - - barrier->completion_signal = signal; - - depBarrier->dep_signal[0] = depSignal; - - uint16_t header = (HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE)| 1 << HSA_PACKET_HEADER_BARRIER; - - // 5. Update packet header, - // Intentionally updated second barrier header before first in order to avoid race - depBarrier->header = header; - barrier->header = header; - - // 6. Trigger the doorbell - nextIndex = nextIndex + 1; - hsa_queue_store_write_index_screlease(lockedQ, nextIndex); - hsa_signal_store_relaxed(lockedQ->doorbell_signal, index+1); - - // 7. Release queue - crit->_av.release_locked_hsa_queue(); - - // 8. Register signal callback - hsa_amd_signal_async_handler(signal, HSA_SIGNAL_CONDITION_EQ, 0, ihipStreamCallbackHandler, cb); + std::thread(ihipStreamCallbackHandler, cb).detach(); return ihipLogStatus(e); } diff --git a/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp b/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp deleted file mode 100644 index e6492c7ce2..0000000000 --- a/projects/hip/tests/src/runtimeApi/stream/StreamAddCallback.cpp +++ /dev/null @@ -1,145 +0,0 @@ -#include -#include -#include -#include "test_common.h" -#include - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 - * TEST: %t - * HIT_END - */ - -enum class ExecState -{ - EXEC_NOT_STARTED, - EXEC_STARTED, - EXEC_CB_STARTED, - EXEC_CB_FINISHED, - EXEC_FINISHED -}; - -struct UserData -{ - size_t size; - int* ptr; -}; - -// Global variable to check exection order -std::atomic gData(ExecState::EXEC_NOT_STARTED); - - -void myCallback(hipStream_t stream, hipError_t status, void* user_data) -{ - if(gData.load() != ExecState::EXEC_STARTED) - return; // Error hence return early - - gData.store(ExecState::EXEC_CB_STARTED); - - UserData* data = reinterpret_cast(user_data); - printf("Callback started\n"); - - sleep(1); - - printf("Callback ending.\n"); - gData.store(ExecState::EXEC_CB_FINISHED); -} - -bool test(int count) -{ - printf("\n============ Test iteration %d =============\n",count); - // Stream - hipStream_t stream; - bool result = true; - - gData.store(ExecState::EXEC_STARTED); - - HIPCHECK(hipStreamCreate(&stream)); - - // Array size - size_t size = 10000; - - // Device array - int *data = NULL; - HIPCHECK(hipMalloc((void**)&data, sizeof(int) * size)); - - // Initialize device array to -1 - HIPCHECK(hipMemset(data, -1, sizeof(int) * size)); - - // Host array - int *host = NULL; - HIPCHECK(hipHostMalloc((void**)&host, sizeof(int) * size)); - - // Print host ptr address - printf("In main thread\n"); - - // Initialize user_data for callback - UserData arg; - arg.size = size; - arg.ptr = host; - - // Synchronize device - HIPCHECK(hipDeviceSynchronize()); - - // Asynchronous copy from device to host - HIPCHECK(hipMemcpyAsync(host, data, sizeof(int) * size, hipMemcpyDeviceToHost, stream)); - - // Asynchronous memset on device - HIPCHECK(hipMemsetAsync(data, 0, sizeof(int) * size, stream)); - - // Add callback - should happen after hipMemsetAsync() - HIPCHECK(hipStreamAddCallback(stream, myCallback, &arg, 0)); - - printf("Will wait in main thread until callback completes\n"); - - //This should synchronize the stream (including the callback) - HIPCHECK(hipStreamSynchronize(stream)); - - if(gData.load() != ExecState::EXEC_CB_FINISHED) - { - std::cout<<"Callback is not finished\n"; - return false; - } - printf("Callback completed will resume main thread execution\n"); - - if(host[size/2] != -1) - { - // Print some host data that just got copied - printf("Pseudo host data printing (should be -1): %d\n", host[size/2]); - result = false; - } - - HIPCHECK(hipMemcpy(host, data, sizeof(int)*size, hipMemcpyDeviceToHost)); - - if(host[size-1] != 0) - { - printf("Pseudo host data printing (should be 0): %d\n", host[size-1]); - result = false; - } - - HIPCHECK(hipFree(data)); - HIPCHECK(hipHostFree(host)); - HIPCHECK(hipStreamDestroy(stream)); - - gData.store(ExecState::EXEC_FINISHED); - return result; -} - -int main() -{ - // Test involves multithreading hence running multiple times - // to make sure consitency in the behavior - bool status = true; - - for(int i=0; i < 10; i++){ - status = test(i+1); - if(status == false) - { - failed("Test Failed!\n"); - break; - } - } - - if(status == true) passed(); - return 0; -} diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp deleted file mode 100644 index b2fb168fbc..0000000000 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackCrossStream.cpp +++ /dev/null @@ -1,133 +0,0 @@ -/* - Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved. - Permission is hereby granted, free of charge, to any person obtaining a copy - of this software and associated documentation files (the "Software"), to deal - in the Software without restriction, including without limitation the rights - to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - copies of the Software, and to permit persons to whom the Software is - furnished to do so, subject to the following conditions: - The above copyright notice and this permission notice shall be included in - all copies or substantial portions of the Software. - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - THE SOFTWARE. -*/ - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 - * TEST: %t - * HIT_END -*/ -#include "test_common.h" -#include -#include -#include - - -//Globals -const int workloadCount = 1000000; -std::mutex gMutx; -bool callbackCompleted = false; -std::condition_variable condVar; - -// Device function -__global__ void increment(int *data,int N) -{ - int i = blockIdx.x*blockDim.x + threadIdx.x; - - if(i < N) - data[i] = 1 + data[i]; -} - -struct USER_DATA -{ - int* result; // Data received from device - int* copyOfOriginalData; // Copy of initial data which will be used for validation -}; - -// Callback -void callback(hipStream_t event, hipError_t status, void *userData) -{ - USER_DATA *data = (USER_DATA *)userData; - - if(!(data == nullptr || data->result == nullptr || data->copyOfOriginalData == nullptr)) - { - for(int i=0;iresult[i] != data->copyOfOriginalData[i]+1) - { - std::cout<<"Error value : "<result[i]<<"| Expected value :"<copyOfOriginalData[i]+1<result = hResultData; - inputParam->copyOfOriginalData = hData; - - HIPCHECK(hipStreamCreate(&stream)); - HIPCHECK(hipStreamAddCallback(stream,callback,inputParam,0)); - - // Wait for stream add callback to complete - std::unique_lock l(gMutx); - - while(!callbackCompleted) - condVar.wait(l); - - // Will destroy device memory hence explicite hipFree is not needed - HIPCHECK(hipDeviceReset()); - - free(hData); - free(hResultData); - passed(); -} diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp deleted file mode 100644 index 3babf5869e..0000000000 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamAddCallbackMultiThread.cpp +++ /dev/null @@ -1,178 +0,0 @@ -/* - Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved. - Permission is hereby granted, free of charge, to any person obtaining a copy - of this software and associated documentation files (the "Software"), to deal - in the Software without restriction, including without limitation the rights - to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - copies of the Software, and to permit persons to whom the Software is - furnished to do so, subject to the following conditions: - The above copyright notice and this permission notice shall be included in - all copies or substantial portions of the Software. - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - THE SOFTWARE. -*/ - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 - * TEST: %t - * HIT_END -*/ - -#include -#include -#include -#include -#include "test_common.h" - -// Will indicate completion of callback added as part of hipStreamAddCallback -struct signal -{ - int completedThreads; - std::mutex mu; - std::condition_variable cv; -}; - -struct workload -{ - int _workloadId; - int _deviceID; - - int *copyOf_hData; // copy of host data which will be used to validation - int *hData; // will contain host data - int *dData; // device data will be stored - hipStream_t _stream;// stream on which data will be processed - - bool success; // start will be stored -}; - -// Global data -int numWorkloads = 8; -const int perWorkloadSize = 1000000; -signal completionSignal; - -// Kernel run on device -__global__ void increment(int *data, int N) -{ - int i = blockIdx.x * blockDim.x + threadIdx.x; - - if (i < N) - data[i] = data[i]+1; -} - -/* - * Method validates processed data array with saved copy and notifies via conditional variable to all waiting threads -*/ -void Analyze(hipStream_t event, hipError_t status, void *data) -{ - HIPCHECK(status); - - workload *W = (workload *) data; - - if(W != NULL) - { - W->success = true; - - for (int i=0; i< perWorkloadSize; ++i) - { - W->success &= (W->copyOf_hData[i] == (W->hData[i]+1)); - if(!W->success) - { - std::cout<<"\nExpected Data :"<<(W->hData[i]+1)<<" Current Data :"<copyOf_hData[i]< guard(completionSignal.mu); - completionSignal.completedThreads += 1; - - completionSignal.cv.notify_all(); -} - -/* - * Thread routine to launch workloads into separate stream - */ -void LaunchWorkload(void *args) -{ - workload *W = (workload *) args; - - if (W == nullptr) return; - - std::srand(std::time(nullptr)); - - HIPCHECK(hipSetDevice(W->_deviceID)); - - // Allocate memory - HIPCHECK(hipMalloc(&W->dData,perWorkloadSize*sizeof(int))); - - size_t s = perWorkloadSize*sizeof(int); - HIPCHECK(hipMemset(W->dData,0,s)) - - W->hData = (int *) malloc(perWorkloadSize*sizeof(int)); - W->copyOf_hData = (int *) malloc(perWorkloadSize*sizeof(int)); - - // Initialize host array - for(int i =0;ihData[i] = W->copyOf_hData[i] = std::rand() % perWorkloadSize; - } - - HIPCHECK(hipStreamCreate(&W->_stream)); - - dim3 block(256); - dim3 grid((perWorkloadSize + block.x-1) / block.x); - - HIPCHECK(hipMemcpyAsync(W->dData,W->hData,perWorkloadSize*sizeof(int),hipMemcpyHostToDevice,W->_stream)); - - hipLaunchKernelGGL(increment, grid, block,0, W->_stream,W->dData, perWorkloadSize); - - HIPCHECK(hipMemcpyAsync(W->copyOf_hData,W->dData,perWorkloadSize*sizeof(int),hipMemcpyDeviceToHost,W->_stream)); - - HIPCHECK(hipStreamAddCallback(W->_stream, Analyze, W,0)) -} - -int main(int argc, char* argv[]) { - - int numDevice = 0; - - HIPCHECK(hipGetDeviceCount(&numDevice)); - - std::thread workerThread[numWorkloads]; - - workload *workloads; - workloads = (workload *) malloc(numWorkloads * sizeof(workload)); - - for(int i =0;i l(completionSignal.mu); - while(completionSignal.completedThreads != numWorkloads) - { - completionSignal.cv.wait(l); - } - - //clean-up - hipDeviceReset(); - free(workloads); - - passed(); - return 0; -} From 4ec70a7e03be747fecc6e47f70f7047410710784 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 20 Nov 2019 16:02:32 +0000 Subject: [PATCH 14/27] Extend vector type capabilities and add tests to reflect it. (#1656) [ROCm/hip commit: b5b3d1bbaae45df1b87e596c0dd87b4c9d6ddb99] --- .../include/hip/hcc_detail/hip_vector_types.h | 37 +++++++++++++++++++ .../tests/src/deviceLib/hipVectorTypes.cpp | 19 ++++++++++ .../src/deviceLib/hipVectorTypesDevice.cpp | 14 +++++++ 3 files changed, 70 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h index 3a10a06372..706b9aa3a7 100644 --- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/hip/include/hip/hcc_detail/hip_vector_types.h @@ -40,6 +40,7 @@ THE SOFTWARE. #endif #if defined(__cplusplus) && defined(__clang__) + #include #include namespace hip_impl { @@ -68,6 +69,23 @@ THE SOFTWARE. } }; + friend + inline + std::ostream& operator<<(std::ostream& os, + const Scalar_accessor& x) noexcept { + return os << x.data[idx]; + } + friend + inline + std::istream& operator>>(std::istream& is, + Scalar_accessor& x) noexcept { + T tmp; + is >> tmp; + x.data[idx] = tmp; + + return is; + } + // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor Vector data; @@ -76,6 +94,17 @@ THE SOFTWARE. __host__ __device__ operator T() const volatile noexcept { return data[idx]; } + __host__ __device__ + operator T&() noexcept { + return reinterpret_cast< + T (&)[sizeof(Vector) / sizeof(T)]>(data)[idx]; + } + __host__ __device__ + operator volatile T&() volatile noexcept { + return reinterpret_cast< + volatile T (&)[sizeof(Vector) / sizeof(T)]>(data)[idx]; + } + __host__ __device__ Address operator&() const noexcept { return Address{this}; } @@ -198,6 +227,8 @@ THE SOFTWARE. Native_vec_ data; hip_impl::Scalar_accessor x; }; + + using value_type = T; }; template @@ -209,6 +240,8 @@ THE SOFTWARE. hip_impl::Scalar_accessor x; hip_impl::Scalar_accessor y; }; + + using value_type = T; }; template @@ -367,6 +400,8 @@ THE SOFTWARE. T z; }; }; + + using value_type = T; }; template @@ -380,6 +415,8 @@ THE SOFTWARE. hip_impl::Scalar_accessor z; hip_impl::Scalar_accessor w; }; + + using value_type = T; }; template diff --git a/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp b/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp index 14479881ff..70c8320073 100644 --- a/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp +++ b/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp @@ -34,6 +34,7 @@ THE SOFTWARE. #include #include #include +#include #include using namespace std; @@ -157,6 +158,24 @@ bool TestVectorType() { if (f1 == f2) return false; if (!(f1 != f2)) return false; + using T = typename V::value_type; + + const T& x = f1.x; + T& y = f2.x; + const volatile T& z = f3.x; + volatile T& w = f2.x; + + if (x != T{3}) return false; + if (y != T{4}) return false; + if (z != T{3}) return false; + if (w != T{4}) return false; + + stringstream str; + str << f1.x; + str >> f2.x; + + if (f1.x != f2.x) return false; + return true; } diff --git a/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp b/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp index ba23931ee5..4bf5d2c87d 100644 --- a/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp +++ b/projects/hip/tests/src/deviceLib/hipVectorTypesDevice.cpp @@ -149,6 +149,20 @@ bool TestVectorType() { if (f1 == f2) return false; if (!(f1 != f2)) return false; + #if 0 // TODO: investigate on GFX8 + using T = typename V::value_type; + + const T& x = f1.x; + T& y = f2.x; + const volatile T& z = f3.x; + volatile T& w = f2.x; + + if (x != T{3}) return false; + if (y != T{4}) return false; + if (z != T{3}) return false; + if (w != T{4}) return false; + #endif + return true; } From 27112bbdf0317eadff2f6d2826a944b578285707 Mon Sep 17 00:00:00 2001 From: Paul Fultz II Date: Wed, 20 Nov 2019 10:03:42 -0600 Subject: [PATCH 15/27] Fix helper header when using c++17 (#1666) This will fix issue #1621. It also adds tests for is_callable with c++11, c++14, and c++17. The fallback implementation was completely broken so I rewrote it so it pass the tests as well. This should be used instead of PR #1631. [ROCm/hip commit: 8519a1411cfeed69fd574c9cf3bdc26d86d27351] --- .../hip/include/hip/hcc_detail/helpers.hpp | 92 ++++---- .../hip/tests/src/cppstd/hipInvocable11.cpp | 31 +++ .../src/cppstd/hipInvocable11Fallback.cpp | 34 +++ .../hip/tests/src/cppstd/hipInvocable14.cpp | 31 +++ .../hip/tests/src/cppstd/hipInvocable17.cpp | 31 +++ .../hip/tests/src/cppstd/is_callable_test.hpp | 209 ++++++++++++++++++ 6 files changed, 391 insertions(+), 37 deletions(-) create mode 100644 projects/hip/tests/src/cppstd/hipInvocable11.cpp create mode 100644 projects/hip/tests/src/cppstd/hipInvocable11Fallback.cpp create mode 100644 projects/hip/tests/src/cppstd/hipInvocable14.cpp create mode 100644 projects/hip/tests/src/cppstd/hipInvocable17.cpp create mode 100644 projects/hip/tests/src/cppstd/is_callable_test.hpp diff --git a/projects/hip/include/hip/hcc_detail/helpers.hpp b/projects/hip/include/hip/hcc_detail/helpers.hpp index cdae0e0fc4..b94b126994 100644 --- a/projects/hip/include/hip/hcc_detail/helpers.hpp +++ b/projects/hip/include/hip/hcc_detail/helpers.hpp @@ -27,6 +27,26 @@ THE SOFTWARE. // std::false_type, std result_of and std::true_type. #include // For std::declval. +#ifdef __has_include // Check if __has_include is present +# if __has_include() // Check for version header +# include +# if defined(__cpp_lib_is_invocable) && !defined(HIP_HAS_INVOCABLE) +# define HIP_HAS_INVOCABLE __cpp_lib_is_invocable +# endif +# if defined(__cpp_lib_result_of_sfinae) && !defined(HIP_HAS_RESULT_OF_SFINAE) +# define HIP_HAS_RESULT_OF_SFINAE __cpp_lib_result_of_sfinae +# endif +# endif +#endif + +#ifndef HIP_HAS_INVOCABLE +#define HIP_HAS_INVOCABLE 0 +#endif + +#ifndef HIP_HAS_RESULT_OF_SFINAE +#define HIP_HAS_RESULT_OF_SFINAE 0 +#endif + namespace std { // TODO: these should be removed as soon as possible. #if (__cplusplus < 201406L) #if (__cplusplus < 201402L) @@ -48,58 +68,56 @@ namespace hip_impl { template using void_t_ = void; -#if (__cplusplus < 201402L) -template -struct is_callable_impl : is_callable_impl {}; +#if HIP_HAS_INVOCABLE +template +struct is_callable_impl; -// Pointer to member function, call through non-pointer. -template -struct is_callable_impl< - F(C, Ts...), 0u, - void_t_().*std::declval())(std::declval()...))> > - : std::true_type {}; - -// Pointer to member function, call through pointer. -template -struct is_callable_impl< - F(C, Ts...), 1u, - void_t_()).*std::declval())(std::declval()...))> > - : std::true_type {}; - -// Pointer to member data, call through non-pointer, no args. -template -struct is_callable_impl().*std::declval())> > - : std::true_type {}; - -// Pointer to member data, call through pointer, no args. -template -struct is_callable_impl().*std::declval())> > - : std::true_type {}; - -// General call, n args. template -struct is_callable_impl()(std::declval()...))> > - : std::true_type {}; - -// Not callable. -template -struct is_callable_impl : std::false_type {}; -#elif (__cplusplus < 201703L) +struct is_callable_impl : std::is_invocable {}; +#elif HIP_HAS_RESULT_OF_SFINAE template struct is_callable_impl : std::false_type {}; template struct is_callable_impl::type > > : std::true_type {}; #else +template +auto simple_invoke(T Base::*pmd, Derived&& ref) +-> decltype(static_cast(ref).*pmd); + +template +auto simple_invoke(PMD&& pmd, Pointer&& ptr) +-> decltype((*static_cast(ptr)).*static_cast(pmd)); -// C++17 +template +auto simple_invoke(T Base::*pmd, const std::reference_wrapper& ref) +-> decltype(ref.get().*pmd); + +template +auto simple_invoke(T Base::*pmf, Derived&& ref, Args&&... args) +-> decltype((static_cast(ref).*pmf)(static_cast(args)...)); + +template +auto simple_invoke(PMF&& pmf, Pointer&& ptr, Args&&... args) +-> decltype(((*static_cast(ptr)).*static_cast(pmf))(static_cast(args)...)); + +template +auto simple_invoke(T Base::*pmf, const std::reference_wrapper& ref, Args&&... args) +-> decltype((ref.get().*pmf)(static_cast(args)...)); + +template +auto simple_invoke(F&& f, Ts&&... xs) +-> decltype(f(static_cast(xs)...)); template struct is_callable_impl : std::false_type {}; template -struct is_callable_impl > > : std::true_type {}; +struct is_callable_impl(), std::declval()...))> > + : std::true_type {}; + #endif + template struct is_callable : is_callable_impl {}; diff --git a/projects/hip/tests/src/cppstd/hipInvocable11.cpp b/projects/hip/tests/src/cppstd/hipInvocable11.cpp new file mode 100644 index 0000000000..6fa7e53622 --- /dev/null +++ b/projects/hip/tests/src/cppstd/hipInvocable11.cpp @@ -0,0 +1,31 @@ +/* +Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "is_callable_test.hpp" +#include + +int main() { passed(); } diff --git a/projects/hip/tests/src/cppstd/hipInvocable11Fallback.cpp b/projects/hip/tests/src/cppstd/hipInvocable11Fallback.cpp new file mode 100644 index 0000000000..071dce8997 --- /dev/null +++ b/projects/hip/tests/src/cppstd/hipInvocable11Fallback.cpp @@ -0,0 +1,34 @@ +/* +Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +// Ensure fallback path is chosen +#define HIP_HAS_RESULT_OF_SFINAE 0 +#define HIP_HAS_INVOCABLE 0 +#include "is_callable_test.hpp" +#include + +int main() { passed(); } diff --git a/projects/hip/tests/src/cppstd/hipInvocable14.cpp b/projects/hip/tests/src/cppstd/hipInvocable14.cpp new file mode 100644 index 0000000000..41b58bd7dd --- /dev/null +++ b/projects/hip/tests/src/cppstd/hipInvocable14.cpp @@ -0,0 +1,31 @@ +/* +Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++14 EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "is_callable_test.hpp" +#include + +int main() { passed(); } diff --git a/projects/hip/tests/src/cppstd/hipInvocable17.cpp b/projects/hip/tests/src/cppstd/hipInvocable17.cpp new file mode 100644 index 0000000000..aa37aac5fb --- /dev/null +++ b/projects/hip/tests/src/cppstd/hipInvocable17.cpp @@ -0,0 +1,31 @@ +/* +Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +/* HIT_START + * BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -std=c++17 EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "is_callable_test.hpp" +#include + +int main() { passed(); } diff --git a/projects/hip/tests/src/cppstd/is_callable_test.hpp b/projects/hip/tests/src/cppstd/is_callable_test.hpp new file mode 100644 index 0000000000..87e4e35b8b --- /dev/null +++ b/projects/hip/tests/src/cppstd/is_callable_test.hpp @@ -0,0 +1,209 @@ +/* +Copyright (c) 2017 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +using hip_impl::is_callable; + +template +struct callable_rank : callable_rank +{}; + +template<> +struct callable_rank<0> +{}; + +struct test1 +{ + struct is_callable_class + { + void operator()(int) const + { + } + }; + struct callable_test_param {}; + + void is_callable_function(int) + { + } + + struct is_callable_rank_class + { + void operator()(int, callable_rank<3>) const + { + } + + void operator()(int, callable_rank<4>) const + { + } + }; + + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); + + typedef void (*is_callable_function_pointer)(int); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(is_callable::value, "Not callable"); + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); + + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + static_assert(is_callable)>::value, "Not callable"); + + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable, callable_test_param)>::value, "callable failed"); + static_assert(not is_callable, callable_test_param)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable)>::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); + static_assert(not is_callable::value, "callable failed"); +}; + +struct test2 +{ + typedef int(callable_rank<0>::*fn)(int); + + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(not is_callable&)>::value, "Failed"); + static_assert(not is_callable const&, int)>::value, "Failed"); +}; + +struct test3 +{ + typedef int(callable_rank<0>::*fn)(int); + + typedef callable_rank<0>* T; + typedef callable_rank<1>* DT; + typedef const callable_rank<0>* CT; + typedef std::unique_ptr> ST; + + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(not is_callable::value, "Failed"); + +}; + +struct test4 +{ + typedef int(callable_rank<0>::*fn); + + static_assert(not is_callable::value, "Failed"); +}; + +struct test5 +{ + typedef int(callable_rank<0>::*fn); + + static_assert(is_callable&)>::value, "Failed"); + static_assert(is_callable&&)>::value, "Failed"); + static_assert(is_callable&)>::value, "Failed"); + static_assert(is_callable&)>::value, "Failed"); +}; + +struct test6 +{ + typedef int(callable_rank<0>::*fn); + + typedef callable_rank<0>* T; + typedef callable_rank<1>* DT; + typedef const callable_rank<0>* CT; + typedef std::unique_ptr> ST; + + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + static_assert(is_callable::value, "Failed"); + +}; + +struct test7 +{ + typedef void(*fp)(callable_rank<0>&, int); + + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(not is_callable&, int)>::value, "Failed"); + static_assert(not is_callable::value, "Failed"); + static_assert(not is_callable&)>::value, "Failed"); +}; + +struct test8 +{ + typedef void(&fp)(callable_rank<0>&, int); + + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(is_callable&, int)>::value, "Failed"); + static_assert(not is_callable&, int)>::value, "Failed"); + static_assert(not is_callable::value, "Failed"); + static_assert(not is_callable&)>::value, "Failed"); +}; From 7e75cf12df6731a59bd37b71885a5f8942aaaa3d Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 20 Nov 2019 08:05:27 -0800 Subject: [PATCH 16/27] Update error codes for hipGetDevice for doxygen and move up null check (#1668) * [docs] Update error codes for hipGetDevice * Move up out ptr check [ROCm/hip commit: b3161e9fa0323c2a355113eb9ca181aaf005b04f] --- .../hip/include/hip/hcc_detail/hip_runtime_api.h | 2 +- projects/hip/src/hip_device.cpp | 14 ++++++-------- 2 files changed, 7 insertions(+), 9 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index afb99da6a3..ffb03d23d7 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -391,7 +391,7 @@ hipError_t hipSetDevice(int deviceId); * This device is used implicitly for HIP runtime APIs called by this thread. * hipGetDevice returns in * @p device the default device for the calling host thread. * - * @returns #hipSuccess + * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue * * @see hipSetDevice, hipGetDevicesizeBytes */ diff --git a/projects/hip/src/hip_device.cpp b/projects/hip/src/hip_device.cpp index 6406b48c2d..403194483a 100644 --- a/projects/hip/src/hip_device.cpp +++ b/projects/hip/src/hip_device.cpp @@ -33,18 +33,16 @@ hipError_t hipGetDevice(int* deviceId) { HIP_INIT_API(hipGetDevice, deviceId); hipError_t e = hipSuccess; + if (deviceId == nullptr) + return ihipLogStatus(hipErrorInvalidValue); auto ctx = ihipGetTlsDefaultCtx(); - if (deviceId != nullptr) { - if (ctx == nullptr) { - e = hipErrorInvalidDevice; // TODO, check error code. - *deviceId = -1; - } else { - *deviceId = ctx->getDevice()->_deviceId; - } + if (ctx == nullptr) { + e = hipErrorInvalidDevice; // TODO, check error code. + *deviceId = -1; } else { - e = hipErrorInvalidValue; + *deviceId = ctx->getDevice()->_deviceId; } return ihipLogStatus(e); From 1baa84eae40e1e4a98bff0c5e8083c6e70c9b7be Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 20 Nov 2019 16:06:37 +0000 Subject: [PATCH 17/27] General sync memcpy improvements. Add `hipMemcpyWithStream` (#1673) * General sync memcpy improvements. Add `hipMemcpyWithStream` * Update hip_memory.cpp [ROCm/hip commit: 5a1f8237399368d2273bd9045fa10d4346697e5a] --- .../include/hip/hcc_detail/hip_runtime_api.h | 8 +- .../include/hip/nvcc_detail/hip_runtime_api.h | 10 + projects/hip/src/hip_memory.cpp | 512 +++++++++++++----- 3 files changed, 398 insertions(+), 132 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index ffb03d23d7..48cec81da4 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -1277,7 +1277,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height * @param[in] height Requested pitched allocation height * * If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned. - * The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array. + * The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array. * Given the row and column of an array element of type T, the address is computed as: * T* pElement = (T*)((char*)BaseAddress + Row * Pitch) + Column; * @@ -1361,6 +1361,10 @@ 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 * @@ -1498,7 +1502,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"))) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 9cb59f14ea..f3ce66c94e 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -610,6 +610,16 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, cudaMemcpy(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind))); } +inline hipError_t hipMemcpyWithStream(void* dst, const void* src, + size_t sizeBytes, hipMemcpyKind copyKind, + hipStream_t stream) { + cudaError_t error = cudaMemcpyAsync(dst, src, sizeBytes, + hipMemcpyKindToCudaMemcpyKind(copyKind), + stream); + if (error != cudaSuccess) return hipCUDAErrorTohipError(error); + + return hipCUDAErrorTohipError(cudaStreamSynchronize(stream)); +} inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind, hipStream_t stream __dparm(0)) { diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 266f9b51d6..be360428c4 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -27,6 +27,7 @@ THE SOFTWARE. #include "hip_hcc_internal.h" #include "trace_helper.h" +#include #include __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP]; @@ -35,23 +36,369 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; // Internal HIP APIS: namespace hip_internal { -hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, - hipStream_t stream) { - hipError_t e = hipSuccess; +namespace { + inline + const char* hsa_to_string(hsa_status_t err) noexcept + { + const char* r{}; - // Return success if number of bytes to copy is 0 - if (sizeBytes == 0) return e; - if (!dst || !src) return hipErrorInvalidValue; + if (hsa_status_string(err, &r) != HSA_STATUS_SUCCESS) return r; - if (!(stream = ihipSyncAndResolveStream(stream))) { - return hipErrorInvalidValue; + return "Unknown."; } + template + inline + void throwing_result_check(hsa_status_t res, const char (&file)[m], + const char (&function)[n], int line) { + if (res == HSA_STATUS_SUCCESS) return; + if (res == HSA_STATUS_INFO_BREAK) return; + + throw std::runtime_error{"Failed in file " + (file + + (", in function \"" + (function + + ("\", on line " + std::to_string(line))))) + + ", with error: " + hsa_to_string(res)}; + } + + inline + hsa_agent_t cpu_agent() noexcept { + static hsa_agent_t cpu{[]() { + hsa_agent_t r{}; + throwing_result_check( + hsa_iterate_agents([](hsa_agent_t x, void* pr) { + hsa_device_type_t t{}; + hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &t); + + if (t != HSA_DEVICE_TYPE_CPU) return HSA_STATUS_SUCCESS; + + *static_cast(pr) = x; + + return HSA_STATUS_INFO_BREAK; + }, &r), __FILE__, __func__, __LINE__); + + return r; + }()}; + + return cpu; + } + + inline + hsa_device_type_t type(hsa_agent_t x) noexcept + { + hsa_device_type_t r{}; + throwing_result_check(hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &r), + __FILE__, __func__, __LINE__); + + return r; + } + + const auto is_large_BAR{[](){ + std::unique_ptr hsa{ + (throwing_result_check(hsa_init(), __FILE__, __func__, __LINE__), + nullptr), + [](void*) { hsa_shut_down(); }}; + bool r{true}; + + throwing_result_check(hsa_iterate_agents([](hsa_agent_t x, void* pr) { + if (x.handle == cpu_agent().handle) return HSA_STATUS_SUCCESS; + + throwing_result_check( + hsa_agent_iterate_regions(x, [](hsa_region_t y, void* p) { + hsa_region_segment_t seg{}; + throwing_result_check( + hsa_region_get_info(y, HSA_REGION_INFO_SEGMENT, &seg), + __FILE__, __func__, __LINE__); + + if (seg != HSA_REGION_SEGMENT_GLOBAL) { + return HSA_STATUS_SUCCESS; + } + + uint32_t flags{}; + throwing_result_check(hsa_region_get_info( + y, HSA_REGION_INFO_GLOBAL_FLAGS, &flags), + __FILE__, __func__, __LINE__); + + if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) { + hsa_amd_memory_pool_access_t tmp{}; + throwing_result_check( + hsa_amd_agent_memory_pool_get_info( + cpu_agent(), + hsa_amd_memory_pool_t{y.handle}, + HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, + &tmp), + __FILE__, __func__, __LINE__); + + *static_cast(p) &= + tmp != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED; + } + + return HSA_STATUS_SUCCESS; + }, pr), __FILE__, __func__, __LINE__); + + return HSA_STATUS_SUCCESS; + }, &r), __FILE__, __func__, __LINE__); + + return r; + }()}; + + inline + hsa_amd_pointer_info_t info(const void* p) noexcept + { + hsa_amd_pointer_info_t r{sizeof(hsa_amd_pointer_info_t)}; + throwing_result_check( + hsa_amd_pointer_info( + const_cast(p), &r, nullptr, nullptr, nullptr), + __FILE__, __func__, __LINE__); + + r.size = is_large_BAR || (type(r.agentOwner) == HSA_DEVICE_TYPE_CPU) ? + UINT32_MAX : sizeof(hsa_amd_pointer_info_t); + + return r; + } + + constexpr size_t staging_sz{4 * 1024 * 1024}; // 2 Pages. + + thread_local const std::unique_ptr staging_buffer{ + []() { + hsa_region_t r{}; + throwing_result_check(hsa_agent_iterate_regions( + cpu_agent(), [](hsa_region_t x, void *p) { + hsa_region_segment_t seg{}; + throwing_result_check( + hsa_region_get_info(x, HSA_REGION_INFO_SEGMENT, &seg), + __FILE__, __func__, __LINE__); + + if (seg != HSA_REGION_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS; + + uint32_t flags{}; + throwing_result_check(hsa_region_get_info( + x, HSA_REGION_INFO_GLOBAL_FLAGS, &flags), + __FILE__, __func__, __LINE__); + + if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) { + *static_cast(p) = x; + + return HSA_STATUS_INFO_BREAK; + } + + return HSA_STATUS_SUCCESS; + }, &r), __FILE__, __func__, __LINE__); + + void *tp{}; + throwing_result_check(hsa_memory_allocate(r, staging_sz, &tp), + __FILE__, __func__, __LINE__); + + return tp; + }(), + [](void *ptr) { hsa_memory_free(ptr); }}; + + thread_local hsa_signal_t copy_signal{[]() { + hsa_agent_t cpu{cpu_agent()}; + hsa_signal_t sgn{}; + throwing_result_check(hsa_signal_create(1, 1, &cpu, &sgn), + __FILE__, __func__, __LINE__); + + return sgn; + }()}; +} // Unnamed namespace. + +inline +void do_copy(void* __restrict dst, const void* __restrict src, std::size_t n, + hsa_agent_t da, hsa_agent_t sa) { + hsa_signal_silent_store_relaxed(copy_signal, 1); + throwing_result_check( + hsa_amd_memory_async_copy(dst, da, src, sa, n, 0, nullptr, copy_signal), + __FILE__, __func__, __LINE__); + + while (hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_EQ, 0, + UINT64_MAX, HSA_WAIT_STATE_ACTIVE)); +} + +inline +void do_std_memcpy( + void* __restrict dst, const void* __restrict src, std::size_t n) { + std::memcpy(dst, src, n); + + return std::atomic_thread_fence(std::memory_order_seq_cst); +} + +inline +void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n, + hsa_amd_pointer_info_t si) { + if (si.size == UINT32_MAX) return do_std_memcpy(dst, src, n); + + const auto di{info(dst)}; + + if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) { + dst = static_cast(di.agentBaseAddress) + + (static_cast(dst) - + static_cast(di.hostBaseAddress)); + do_copy(dst, src, n, si.agentOwner, si.agentOwner); + } + else if (n <= staging_sz) { + do_copy(staging_buffer.get(), src, n, si.agentOwner, si.agentOwner); + std::memcpy(dst, staging_buffer.get(), n); + } + else { + std::unique_ptr lck{ + dst, [](void* p) { hsa_amd_memory_unlock(p); }}; + + throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1, + const_cast(&dst)), + __FILE__, __func__, __LINE__); + + do_copy(dst, src, n, si.agentOwner, si.agentOwner); + } +} + +inline +void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n, + hsa_amd_pointer_info_t di) { + if (di.size == UINT32_MAX) return do_std_memcpy(dst, src, n); + + const auto si{info(const_cast(src))}; + + if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) { + src = static_cast(si.agentBaseAddress) + + (static_cast(src) - + static_cast(si.hostBaseAddress)); + do_copy(dst, src, n, di.agentOwner, di.agentOwner); + } + else if (n <= staging_sz) { + std::memcpy(staging_buffer.get(), src, n); + do_copy(dst, staging_buffer.get(), n, di.agentOwner, di.agentOwner); + } + else { + std::unique_ptr lck{ + const_cast(src), [](void* p) { hsa_amd_memory_unlock(p); }}; + + throwing_result_check(hsa_amd_memory_lock(const_cast(src), n, + &di.agentOwner, 1, + const_cast(&src)), + __FILE__, __func__, __LINE__); + + do_copy(dst, src, n, di.agentOwner, di.agentOwner); + } +} + +inline +void generic_copy(void* __restrict dst, const void* __restrict src, size_t n, + hsa_amd_pointer_info_t di, hsa_amd_pointer_info_t si) { + if (di.size == UINT32_MAX && si.size == UINT32_MAX) { + return do_std_memcpy(dst, src, n); + } + + std::unique_ptr lck0{ + nullptr, [](void* p) { hsa_amd_memory_unlock(p); }}; + std::unique_ptr lck1{nullptr, lck0.get_deleter()}; + + switch (si.type) { + case HSA_EXT_POINTER_TYPE_HSA: + if (di.type == HSA_EXT_POINTER_TYPE_HSA) { + hsa_memory_copy(dst, src, n); + return; // TODO: do_copy(dst, src, n, di.agentOwner, si.agentOwner); + } + + if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN || + di.type == HSA_EXT_POINTER_TYPE_LOCKED) { + return d2h_copy(dst, src, n, si); + } + break; + case HSA_EXT_POINTER_TYPE_LOCKED: + if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) { + std::memcpy(dst, si.hostBaseAddress, n); + + return; + } + if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) { + std::memcpy(di.hostBaseAddress, si.hostBaseAddress, n); + + return; + } + src = si.agentBaseAddress; + si.agentOwner = di.agentOwner; + break; + case HSA_EXT_POINTER_TYPE_UNKNOWN: + if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) { + std::memcpy(dst, src, n); + + return; + } + if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) { + std::memcpy(di.hostBaseAddress, src, n); + + return; + } + return h2d_copy(dst, src, n, di); + default: do_copy(dst, src, n, di.agentOwner, si.agentOwner); break; + } +} + +inline +void memcpy_impl(void* __restrict dst, const void* __restrict src, size_t n, + hipMemcpyKind k) noexcept { + switch (k) { + case hipMemcpyHostToHost: std::memcpy(dst, src, n); break; + case hipMemcpyHostToDevice: + return is_large_BAR ? do_std_memcpy(dst, src, n) + : h2d_copy(dst, src, n, info(dst)); + case hipMemcpyDeviceToHost: + return is_large_BAR ? do_std_memcpy(dst, src, n) + : d2h_copy(dst, src, n, info(src)); + case hipMemcpyDeviceToDevice: + return do_copy(dst, src, n, info(dst).agentOwner, + info(const_cast(src)).agentOwner); + default: return generic_copy(dst, src, n, info(dst), info(src)); + } +} + +hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, + hipMemcpyKind kind, hipStream_t stream) { + if (!dst || !src) return hipErrorInvalidValue; + if (sizeBytes == 0) return hipSuccess; + try { + stream = ihipSyncAndResolveStream(stream); + + if (!stream) return hipErrorInvalidValue; + stream->locked_copyAsync(dst, src, sizeBytes, kind); } - catch (ihipException& ex) { - e = ex._code; + 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 (!dst || !src) return hipErrorInvalidValue; + if (sizeBytes == 0) return hipSuccess; + + 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 (...) { return hipErrorUnknown; @@ -896,20 +1243,8 @@ hipError_t hipMemcpyToSymbol(void* dst, const void* src, size_t count, tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst); - 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); + return ihipLogStatus( + hipMemcpy(static_cast(dst) + offset, src, count, kind)); } hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count, @@ -920,20 +1255,8 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count, tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst); - 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); + return ihipLogStatus( + hipMemcpy(dst, static_cast(src) + offset, count, kind)); } @@ -995,120 +1318,49 @@ 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); - 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); + return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind, + hipStreamNull)); } - hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes); - 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); + return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, + hipMemcpyHostToDevice, + hipStreamNull)); } - hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes); - 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); + return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, + hipMemcpyDeviceToHost, + hipStreamNull)); } hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes); - 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); + return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, + hipMemcpyDeviceToDevice, + hipStreamNull)); } hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { HIP_INIT_SPECIAL_API(hipMemcpyHtoH, (TRACE_MCMD), dst, src, sizeBytes); - hipError_t e = hipSuccess; - if (sizeBytes == 0) return ihipLogStatus(e); + return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, + hipMemcpyHostToHost, + hipStreamNull)); +} - if(dst==NULL || src==NULL){ - return ihipLogStatus(hipErrorInvalidValue); - } +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); - 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); + return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind, + stream)); } hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, From 3848cd97ea4ebf84317b7eff3fc4496247587d4c Mon Sep 17 00:00:00 2001 From: satyanveshd <53337087+satyanveshd@users.noreply.github.com> Date: Wed, 20 Nov 2019 21:37:52 +0530 Subject: [PATCH 18/27] fixed directed tests fail when hcc bumped to 3.0 (#1678) Handled the HCC version check appropriately as few of the directed tests (SWDEV-212161) were failing when hcc was bumped to 3.0. [ROCm/hip commit: 6b06911ef177d1a55c2da9a399efa013ced2d3ac] --- projects/hip/include/hip/hcc_detail/device_functions.h | 2 +- projects/hip/include/hip/hcc_detail/device_library_decls.h | 2 +- projects/hip/src/hip_stream.cpp | 6 +++--- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index a6f9cc0826..68e3277270 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -794,7 +794,7 @@ void *__amdgcn_get_dynamicgroupbaseptr() { return __get_dynamicgroupbaseptr(); } -#if defined(__HCC__) && (__hcc_minor__ < 3) +#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3) // hip.amdgcn.bc - sync threads #define __CLK_LOCAL_MEM_FENCE 0x01 typedef unsigned __cl_mem_fence_flags; diff --git a/projects/hip/include/hip/hcc_detail/device_library_decls.h b/projects/hip/include/hip/hcc_detail/device_library_decls.h index 8cfb020caa..ac35823cd2 100644 --- a/projects/hip/include/hip/hcc_detail/device_library_decls.h +++ b/projects/hip/include/hip/hcc_detail/device_library_decls.h @@ -80,7 +80,7 @@ extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(vo __device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; } #endif //__HIP_DEVICE_COMPILE__ -#if defined(__HCC__) && (__hcc_minor__ < 3) +#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3) // __llvm_fence* functions from device-libs/irif/src/fence.ll extern "C" __device__ void __llvm_fence_acq_sg(void); extern "C" __device__ void __llvm_fence_acq_wg(void); diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index e3e4975b7e..a9f30197d5 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -31,7 +31,7 @@ THE SOFTWARE. //------------------------------------------------------------------------------------------------- // Stream // -#if defined(__HCC__) && (__hcc_minor__ < 3) +#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3) enum queue_priority { priority_high = 0, @@ -73,7 +73,7 @@ hipError_t ihipStreamCreate(TlsData *tls, hipStream_t* stream, unsigned int flag // Obtain mutex access to the device critical data, release by destructor LockedAccessor_CtxCrit_t ctxCrit(ctx->criticalData()); -#if defined(__HCC__) && (__hcc_minor__ < 3) +#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3) auto istream = new ihipStream_t(ctx, acc.create_view(), flags); #else auto istream = new ihipStream_t(ctx, acc.create_view(Kalmar::execute_any_order, Kalmar::queuing_mode_automatic, (Kalmar::queue_priority)priority), flags); @@ -242,7 +242,7 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) { } else if (stream == hipStreamNull) { return ihipLogStatus(hipErrorInvalidResourceHandle); } else { -#if defined(__HCC__) && (__hcc_minor__ < 3) +#if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3) *priority = 0; #else LockedAccessor_StreamCrit_t crit(stream->criticalData()); From fd08d068c770a631f57a60519630c18cd14866f2 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 22 Nov 2019 02:23:48 +0000 Subject: [PATCH 19/27] Use native support for atomic FADD when address is in LDS (#1591) [ROCm/hip commit: d597e7ca20ac636a5992718715f26abbc4b85e66] --- .../hip/include/hip/hcc_detail/hip_atomic.h | 33 ++++++++++++++- .../src/deviceLib/hipSimpleAtomicsTest.cpp | 42 ++++++++++++++++--- 2 files changed, 68 insertions(+), 7 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_atomic.h b/projects/hip/include/hip/hcc_detail/hip_atomic.h index 263f639e96..c8dcf02322 100644 --- a/projects/hip/include/hip/hcc_detail/hip_atomic.h +++ b/projects/hip/include/hip/hcc_detail/hip_atomic.h @@ -55,7 +55,7 @@ unsigned long long atomicAdd( } __device__ inline -float atomicAdd(float* address, float val) +float atomicAdd_impl(float* address, float val) { unsigned int* uaddr{reinterpret_cast(address)}; unsigned int r{__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; @@ -73,6 +73,37 @@ float atomicAdd(float* address, float val) return __uint_as_float(r); } +#if !__has_builtin(__builtin_amdgcn_is_shared) + __device__ + inline + bool __builtin_amdgcn_is_shared( + const __attribute__((address_space(0))) void* ptr) noexcept + { + #if defined(__HIP_DEVICE_COMPILE__) + const unsigned int gp = reinterpret_cast(ptr); + + return gp == + (__builtin_amdgcn_s_getreg((15 << 11) | (16 << 6) | 15) << 16); + #else + return false; + #endif + } +#endif +__device__ +inline +float atomicAdd(float* address, float val) +{ + using GP = const __attribute__((address_space(0))) void*; + using LP = __attribute__((address_space(3))) float*; + + #if __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ || __HIP_ARCH_GFX908__ + if (__builtin_amdgcn_is_shared((GP) address)) { + return __builtin_amdgcn_ds_faddf((LP) address, val, 0, 0, false); + } + #endif + + return atomicAdd_impl(address, val); +} __device__ inline double atomicAdd(double* address, double val) diff --git a/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp index 760d65e555..5a09799c0e 100644 --- a/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp +++ b/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp @@ -29,6 +29,7 @@ THE SOFTWARE. // includes, system #include +#include #include #include #include @@ -212,7 +213,7 @@ __device__ void testKernelSub(...) {} template< - typename T, + typename T, typename enable_if< is_same{} || is_same{}>::type* = nullptr> __device__ @@ -259,13 +260,37 @@ void testKernelIntegral(T* g_odata) { testKernelSub(g_odata); } +namespace { + constexpr unsigned int numData = 11; +} + template -__global__ void testKernel(T* g_odata) { +__global__ void testKernel(T* g_odata, T* g_sdata) { // Atomic addition atomicAdd(&g_odata[0], 10); testKernelIntegral(g_odata); testKernelExch(g_odata); + + #if !defined(HIP_PLATFORM_NVCC) + // Shared Atomic addition. + __shared__ T s_odata[numData]; + + if (threadIdx.x == 0) { s_odata[8] = s_odata[10] = 0xff; } + + __syncthreads(); + + atomicAdd(&s_odata[0], 10); + + testKernelIntegral(s_odata); + testKernelExch(s_odata); + + __syncthreads(); + + if (threadIdx.x == 0) { + __builtin_memcpy(g_sdata, s_odata, sizeof(T) * numData); + } + #endif } template @@ -286,33 +311,38 @@ void runTest() { unsigned int numThreads = 256; unsigned int numBlocks = 64; - unsigned int numData = 11; unsigned int memSize = sizeof(T) * numData; // allocate mem for the result on host side T* hOData = (T*)malloc(memSize); + T* hSData = (T*)malloc(memSize); // initialize the memory - for (unsigned int i = 0; i < numData; i++) hOData[i] = 0; + std::memset(hOData, 0, memSize); + std::memset(hSData, 0, memSize); // To make the AND and XOR tests generate something other than 0... hOData[8] = hOData[10] = 0xff; // allocate device memory for result T* dOData; + T* dSData; hipMalloc((void**)&dOData, memSize); + hipMalloc((void**)&dSData, memSize); // copy host memory to device to initialize to zero hipMemcpy(dOData, hOData, memSize, hipMemcpyHostToDevice); // execute the kernel hipLaunchKernelGGL( - testKernel, dim3(numBlocks), dim3(numThreads), 0, 0, dOData); + testKernel, dim3(numBlocks), dim3(numThreads), 0, 0, dOData, dSData); // Copy result from device to host hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost); + hipMemcpy(hSData, dSData, memSize, hipMemcpyDeviceToHost); // Compute reference solution - testResult = computeGold(hOData, numThreads * numBlocks); + testResult = std::equal(hOData, hOData + numData, hSData) && + computeGold(hOData, numThreads * numBlocks); // Cleanup memory free(hOData); From 92837d4b4de36f5b31de9a0aa30d267bdd68ed9e Mon Sep 17 00:00:00 2001 From: ansurya <50609411+ansurya@users.noreply.github.com> Date: Fri, 22 Nov 2019 07:54:17 +0530 Subject: [PATCH 20/27] Fix rocBLAS compilation failure (#1677) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit SWDEV-212749: o Recent changes to “add support for extended launch” require hip_runtime.h to be include in hip_ext.h o Order in which external applications include hip_hcc.h/hip_runtime.h causes compilation failure [ROCm/hip commit: e60dec51da64df8fb5784412c5b28b0aed1fd402] --- projects/hip/include/hip/hip_ext.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hip_ext.h b/projects/hip/include/hip/hip_ext.h index 5f6ba31d1e..a618462d6e 100644 --- a/projects/hip/include/hip/hip_ext.h +++ b/projects/hip/include/hip/hip_ext.h @@ -22,7 +22,7 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HIP_EXT_H #define HIP_INCLUDE_HIP_HIP_EXT_H -#include "hip/hip_runtime_api.h" +#include "hip/hip_runtime.h" #ifdef __HCC__ // Forward declarations: From c8c265891afb7568cd36fcce853f2cedd4fea77a Mon Sep 17 00:00:00 2001 From: Shyam Saladi Date: Thu, 21 Nov 2019 18:24:37 -0800 Subject: [PATCH 21/27] [HIPIFY] fix typo (#1681) [ROCm/hip commit: 70f457d7ea6de1310fba6f3c8601e3d179544f5c] --- projects/hip/hipify-clang/src/ArgParse.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/hip/hipify-clang/src/ArgParse.cpp b/projects/hip/hipify-clang/src/ArgParse.cpp index 4f648c996f..cbf2f941e8 100644 --- a/projects/hip/hipify-clang/src/ArgParse.cpp +++ b/projects/hip/hipify-clang/src/ArgParse.cpp @@ -30,7 +30,7 @@ cl::opt OutputFilename("o", cl::cat(ToolTemplateCategory)); cl::opt OutputDir("o-dir", - cl::desc("Output direcory"), + cl::desc("Output directory"), cl::value_desc("directory"), cl::cat(ToolTemplateCategory)); @@ -45,17 +45,17 @@ cl::opt GeneratePython("python", cl::cat(ToolTemplateCategory)); cl::opt OutputHipifyPerlDir("o-hipify-perl-dir", - cl::desc("Output direcory for hipify-perl script"), + cl::desc("Output directory for hipify-perl script"), cl::value_desc("directory"), cl::cat(ToolTemplateCategory)); cl::opt OutputPythonMapDir("o-python-map-dir", - cl::desc("Output direcory for Python map"), + cl::desc("Output directory for Python map"), cl::value_desc("directory"), cl::cat(ToolTemplateCategory)); cl::opt TemporaryDir("temp-dir", - cl::desc("Temporary direcory"), + cl::desc("Temporary directory"), cl::value_desc("directory"), cl::cat(ToolTemplateCategory)); From f220a7592e4db6b0984d3e15e27f30c416caf023 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 22 Nov 2019 02:24:53 +0000 Subject: [PATCH 22/27] Unary operators were too restrictive in the type of their argument. (#1683) [ROCm/hip commit: aaf31b6b96a71bf25d35c420e05c05271d22a403] --- .../include/hip/hcc_detail/hip_vector_types.h | 54 +++++++++++++------ 1 file changed, 39 insertions(+), 15 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h index 706b9aa3a7..582a359fbd 100644 --- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/hip/include/hip/hcc_detail/hip_vector_types.h @@ -144,22 +144,40 @@ THE SOFTWARE. return *this; } + // TODO: convertibility is too restrictive, constraint should be on + // the operator being invocable with a value of type U. + template< + typename U, + typename std::enable_if< + std::is_convertible{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator+=(T x) noexcept { + Scalar_accessor& operator+=(U x) noexcept { data[idx] += x; return *this; } + template< + typename U, + typename std::enable_if< + std::is_convertible{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator-=(T x) noexcept { + Scalar_accessor& operator-=(U x) noexcept { data[idx] -= x; return *this; } + template< + typename U, + typename std::enable_if< + std::is_convertible{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator*=(T x) noexcept { + Scalar_accessor& operator*=(U x) noexcept { data[idx] *= x; return *this; } + template< + typename U, + typename std::enable_if< + std::is_convertible{}>::type* = nullptr> __host__ __device__ Scalar_accessor& operator/=(T x) noexcept { data[idx] /= x; @@ -167,50 +185,56 @@ THE SOFTWARE. } template< typename U = T, - typename std::enable_if{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator%=(T x) noexcept { + Scalar_accessor& operator%=(U x) noexcept { data[idx] %= x; return *this; } template< typename U = T, - typename std::enable_if{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator>>=(T x) noexcept { + Scalar_accessor& operator>>=(U x) noexcept { data[idx] >>= x; return *this; } template< typename U = T, - typename std::enable_if{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator<<=(T x) noexcept { + Scalar_accessor& operator<<=(U x) noexcept { data[idx] <<= x; return *this; } template< typename U = T, - typename std::enable_if{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator&=(T x) noexcept { + Scalar_accessor& operator&=(U x) noexcept { data[idx] &= x; return *this; } template< typename U = T, - typename std::enable_if{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator|=(T x) noexcept { + Scalar_accessor& operator|=(U x) noexcept { data[idx] |= x; return *this; } template< typename U = T, - typename std::enable_if{}>::type* = nullptr> + typename std::enable_if{} && + std::is_integral{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator^=(T x) noexcept { + Scalar_accessor& operator^=(U x) noexcept { data[idx] ^= x; return *this; } From f1544b621dd7d78362937180ede87fe2bf115972 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 25 Nov 2019 19:03:03 -0800 Subject: [PATCH 23/27] fix hipMempcy precondition, test for zero size first (#1697) [ROCm/hip commit: 299bbe232506782d368f5fb3c8b9518dbfc691b0] --- projects/hip/src/hip_memory.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index be360428c4..00a7af59f8 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -354,8 +354,8 @@ void memcpy_impl(void* __restrict dst, const void* __restrict src, size_t n, hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) { - if (!dst || !src) return hipErrorInvalidValue; if (sizeBytes == 0) return hipSuccess; + if (!dst || !src) return hipErrorInvalidValue; try { stream = ihipSyncAndResolveStream(stream); @@ -380,8 +380,8 @@ hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipError_t memcpySync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) { - if (!dst || !src) return hipErrorInvalidValue; if (sizeBytes == 0) return hipSuccess; + if (!dst || !src) return hipErrorInvalidValue; try { stream = ihipSyncAndResolveStream(stream); From dc9bc8ee862f59d17bc878bb4f3fd24187e81c4a Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 28 Nov 2019 10:09:05 +0200 Subject: [PATCH 24/27] Uniform `is_shared` query. (#1695) [ROCm/hip commit: b6514fffb9b1dadd84856eb0947bc17a405ab54a] --- .../hip/include/hip/hcc_detail/hip_atomic.h | 29 +++++++++---------- 1 file changed, 13 insertions(+), 16 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_atomic.h b/projects/hip/include/hip/hcc_detail/hip_atomic.h index c8dcf02322..699dbe7816 100644 --- a/projects/hip/include/hip/hcc_detail/hip_atomic.h +++ b/projects/hip/include/hip/hcc_detail/hip_atomic.h @@ -73,22 +73,19 @@ float atomicAdd_impl(float* address, float val) return __uint_as_float(r); } -#if !__has_builtin(__builtin_amdgcn_is_shared) - __device__ - inline - bool __builtin_amdgcn_is_shared( - const __attribute__((address_space(0))) void* ptr) noexcept - { - #if defined(__HIP_DEVICE_COMPILE__) - const unsigned int gp = reinterpret_cast(ptr); +__device__ +inline +bool __hip_is_shared(const __attribute__((address_space(0))) void* ptr) noexcept +{ // TODO: this is ersatz for __builtin_amdgcn_is_shared. + #if defined(__HIP_DEVICE_COMPILE__) + const unsigned int gp = reinterpret_cast(ptr); - return gp == - (__builtin_amdgcn_s_getreg((15 << 11) | (16 << 6) | 15) << 16); - #else - return false; - #endif - } -#endif + return + gp ==(__builtin_amdgcn_s_getreg((15 << 11) | (16 << 6) | 15) << 16); + #else + return false; + #endif +} __device__ inline float atomicAdd(float* address, float val) @@ -97,7 +94,7 @@ float atomicAdd(float* address, float val) using LP = __attribute__((address_space(3))) float*; #if __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ || __HIP_ARCH_GFX908__ - if (__builtin_amdgcn_is_shared((GP) address)) { + if (__hip_is_shared((GP) address)) { return __builtin_amdgcn_ds_faddf((LP) address, val, 0, 0, false); } #endif From 4e564f783e48aae354820f02157d962c6f7ee307 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 29 Nov 2019 11:58:12 +0530 Subject: [PATCH 25/27] Revert changes for atomic FADD support when address is in LDS (#1701) This reverts PR #1591 and follow-on PR #1695 [ROCm/hip commit: 32442c65064f835f4acaa716ca246e100d69500b] --- .../hip/include/hip/hcc_detail/hip_atomic.h | 30 +------------ .../src/deviceLib/hipSimpleAtomicsTest.cpp | 42 +++---------------- 2 files changed, 7 insertions(+), 65 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_atomic.h b/projects/hip/include/hip/hcc_detail/hip_atomic.h index 699dbe7816..263f639e96 100644 --- a/projects/hip/include/hip/hcc_detail/hip_atomic.h +++ b/projects/hip/include/hip/hcc_detail/hip_atomic.h @@ -55,7 +55,7 @@ unsigned long long atomicAdd( } __device__ inline -float atomicAdd_impl(float* address, float val) +float atomicAdd(float* address, float val) { unsigned int* uaddr{reinterpret_cast(address)}; unsigned int r{__atomic_load_n(uaddr, __ATOMIC_RELAXED)}; @@ -75,34 +75,6 @@ float atomicAdd_impl(float* address, float val) } __device__ inline -bool __hip_is_shared(const __attribute__((address_space(0))) void* ptr) noexcept -{ // TODO: this is ersatz for __builtin_amdgcn_is_shared. - #if defined(__HIP_DEVICE_COMPILE__) - const unsigned int gp = reinterpret_cast(ptr); - - return - gp ==(__builtin_amdgcn_s_getreg((15 << 11) | (16 << 6) | 15) << 16); - #else - return false; - #endif -} -__device__ -inline -float atomicAdd(float* address, float val) -{ - using GP = const __attribute__((address_space(0))) void*; - using LP = __attribute__((address_space(3))) float*; - - #if __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ || __HIP_ARCH_GFX908__ - if (__hip_is_shared((GP) address)) { - return __builtin_amdgcn_ds_faddf((LP) address, val, 0, 0, false); - } - #endif - - return atomicAdd_impl(address, val); -} -__device__ -inline double atomicAdd(double* address, double val) { unsigned long long* uaddr{reinterpret_cast(address)}; diff --git a/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp b/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp index 5a09799c0e..760d65e555 100644 --- a/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp +++ b/projects/hip/tests/src/deviceLib/hipSimpleAtomicsTest.cpp @@ -29,7 +29,6 @@ THE SOFTWARE. // includes, system #include -#include #include #include #include @@ -213,7 +212,7 @@ __device__ void testKernelSub(...) {} template< - typename T, + typename T, typename enable_if< is_same{} || is_same{}>::type* = nullptr> __device__ @@ -260,37 +259,13 @@ void testKernelIntegral(T* g_odata) { testKernelSub(g_odata); } -namespace { - constexpr unsigned int numData = 11; -} - template -__global__ void testKernel(T* g_odata, T* g_sdata) { +__global__ void testKernel(T* g_odata) { // Atomic addition atomicAdd(&g_odata[0], 10); testKernelIntegral(g_odata); testKernelExch(g_odata); - - #if !defined(HIP_PLATFORM_NVCC) - // Shared Atomic addition. - __shared__ T s_odata[numData]; - - if (threadIdx.x == 0) { s_odata[8] = s_odata[10] = 0xff; } - - __syncthreads(); - - atomicAdd(&s_odata[0], 10); - - testKernelIntegral(s_odata); - testKernelExch(s_odata); - - __syncthreads(); - - if (threadIdx.x == 0) { - __builtin_memcpy(g_sdata, s_odata, sizeof(T) * numData); - } - #endif } template @@ -311,38 +286,33 @@ void runTest() { unsigned int numThreads = 256; unsigned int numBlocks = 64; + unsigned int numData = 11; unsigned int memSize = sizeof(T) * numData; // allocate mem for the result on host side T* hOData = (T*)malloc(memSize); - T* hSData = (T*)malloc(memSize); // initialize the memory - std::memset(hOData, 0, memSize); - std::memset(hSData, 0, memSize); + for (unsigned int i = 0; i < numData; i++) hOData[i] = 0; // To make the AND and XOR tests generate something other than 0... hOData[8] = hOData[10] = 0xff; // allocate device memory for result T* dOData; - T* dSData; hipMalloc((void**)&dOData, memSize); - hipMalloc((void**)&dSData, memSize); // copy host memory to device to initialize to zero hipMemcpy(dOData, hOData, memSize, hipMemcpyHostToDevice); // execute the kernel hipLaunchKernelGGL( - testKernel, dim3(numBlocks), dim3(numThreads), 0, 0, dOData, dSData); + testKernel, dim3(numBlocks), dim3(numThreads), 0, 0, dOData); // Copy result from device to host hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost); - hipMemcpy(hSData, dSData, memSize, hipMemcpyDeviceToHost); // Compute reference solution - testResult = std::equal(hOData, hOData + numData, hSData) && - computeGold(hOData, numThreads * numBlocks); + testResult = computeGold(hOData, numThreads * numBlocks); // Cleanup memory free(hOData); From 389263471e7b9b54f2ebd5e0ff9fa4dc4a63a4b0 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 29 Nov 2019 19:29:11 +0300 Subject: [PATCH 26/27] [HIPIFY] Sync cudaError/CUresult between CUDA Driver API/CUDA RT API/HIP + Update *.md docs and hipify-perl accordingly [Reason] Starting with CUDA 10.1 all error codes are merged between Driver and RT APIs [ToDo] Do the same merge in HIP API as there is no need in distinguishing return codes by API [ROCm/hip commit: d33dd71ed3b5ef73fb1a11a6ba6dfd5a2793bec2] --- projects/hip/bin/hipify-perl | 22 +- ...A_Driver_API_functions_supported_by_HIP.md | 10 +- ..._Runtime_API_functions_supported_by_HIP.md | 108 +++--- .../src/CUDA2HIP_Driver_API_types.cpp | 198 ++++++----- .../src/CUDA2HIP_Runtime_API_types.cpp | 312 +++++++++++------- 5 files changed, 392 insertions(+), 258 deletions(-) diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 7b286d638d..b804d8d62d 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -1027,7 +1027,7 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_CONTEXT\b/hipErrorInvalidContext/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_DEVICE\b/hipErrorInvalidDevice/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_GRAPHICS_CONTEXT\b/hipErrorInvalidGraphicsContext/g; - $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_HANDLE\b/hipErrorInvalidResourceHandle/g; + $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_HANDLE\b/hipErrorInvalidHandle/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_IMAGE\b/hipErrorInvalidImage/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_PTX\b/hipErrorInvalidKernelFile/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_INVALID_SOURCE\b/hipErrorInvalidSource/g; @@ -1042,13 +1042,15 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bCUDA_ERROR_NOT_MAPPED_AS_ARRAY\b/hipErrorNotMappedAsArray/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_NOT_MAPPED_AS_POINTER\b/hipErrorNotMappedAsPointer/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_NOT_READY\b/hipErrorNotReady/g; + $ft{'numeric_literal'} += s/\bCUDA_ERROR_NOT_SUPPORTED\b/hipErrorNotSupported/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_NO_BINARY_FOR_GPU\b/hipErrorNoBinaryForGpu/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_NO_DEVICE\b/hipErrorNoDevice/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_OPERATING_SYSTEM\b/hipErrorOperatingSystem/g; - $ft{'numeric_literal'} += s/\bCUDA_ERROR_OUT_OF_MEMORY\b/hipErrorMemoryAllocation/g; + $ft{'numeric_literal'} += s/\bCUDA_ERROR_OUT_OF_MEMORY\b/hipErrorOutOfMemory/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED\b/hipErrorPeerAccessAlreadyEnabled/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_PEER_ACCESS_NOT_ENABLED\b/hipErrorPeerAccessNotEnabled/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_PEER_ACCESS_UNSUPPORTED\b/hipErrorPeerAccessUnsupported/g; + $ft{'numeric_literal'} += s/\bCUDA_ERROR_PRIMARY_CONTEXT_ACTIVE\b/hipErrorSetOnActiveProcess/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_PROFILER_ALREADY_STARTED\b/hipErrorProfilerAlreadyStarted/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_PROFILER_ALREADY_STOPPED\b/hipErrorProfilerAlreadyStopped/g; $ft{'numeric_literal'} += s/\bCUDA_ERROR_PROFILER_DISABLED\b/hipErrorProfilerDisabled/g; @@ -1472,8 +1474,15 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bcudaDevAttrTextureAlignment\b/hipDeviceAttributeTextureAlignment/g; $ft{'numeric_literal'} += s/\bcudaDevAttrTotalConstantMemory\b/hipDeviceAttributeTotalConstantMemory/g; $ft{'numeric_literal'} += s/\bcudaDevAttrWarpSize\b/hipDeviceAttributeWarpSize/g; + $ft{'numeric_literal'} += s/\bcudaErrorAlreadyAcquired\b/hipErrorAlreadyAcquired/g; + $ft{'numeric_literal'} += s/\bcudaErrorAlreadyMapped\b/hipErrorAlreadyMapped/g; + $ft{'numeric_literal'} += s/\bcudaErrorArrayIsMapped\b/hipErrorArrayIsMapped/g; $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/\bcudaErrorDeviceUninitilialized\b/hipErrorInvalidContext/g; $ft{'numeric_literal'} += s/\bcudaErrorECCUncorrectable\b/hipErrorECCNotCorrectable/g; + $ft{'numeric_literal'} += s/\bcudaErrorFileNotFound\b/hipErrorFileNotFound/g; $ft{'numeric_literal'} += s/\bcudaErrorHostMemoryAlreadyRegistered\b/hipErrorHostMemoryAlreadyRegistered/g; $ft{'numeric_literal'} += s/\bcudaErrorHostMemoryNotRegistered\b/hipErrorHostMemoryNotRegistered/g; $ft{'numeric_literal'} += s/\bcudaErrorIllegalAddress\b/hipErrorIllegalAddress/g; @@ -1488,17 +1497,22 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bcudaErrorInvalidMemcpyDirection\b/hipErrorInvalidMemcpyDirection/g; $ft{'numeric_literal'} += s/\bcudaErrorInvalidPtx\b/hipErrorInvalidKernelFile/g; $ft{'numeric_literal'} += s/\bcudaErrorInvalidResourceHandle\b/hipErrorInvalidResourceHandle/g; + $ft{'numeric_literal'} += s/\bcudaErrorInvalidSource\b/hipErrorInvalidSource/g; $ft{'numeric_literal'} += s/\bcudaErrorInvalidSymbol\b/hipErrorInvalidSymbol/g; $ft{'numeric_literal'} += s/\bcudaErrorInvalidValue\b/hipErrorInvalidValue/g; $ft{'numeric_literal'} += s/\bcudaErrorLaunchFailure\b/hipErrorLaunchFailure/g; $ft{'numeric_literal'} += s/\bcudaErrorLaunchOutOfResources\b/hipErrorLaunchOutOfResources/g; $ft{'numeric_literal'} += s/\bcudaErrorLaunchTimeout\b/hipErrorLaunchTimeOut/g; - $ft{'numeric_literal'} += s/\bcudaErrorMapBufferObjectFailed\b/hipErrorMapFailed/g; + $ft{'numeric_literal'} += s/\bcudaErrorMapBufferObjectFailed\b/hipErrorMapBufferObjectFailed/g; $ft{'numeric_literal'} += s/\bcudaErrorMemoryAllocation\b/hipErrorMemoryAllocation/g; $ft{'numeric_literal'} += s/\bcudaErrorMissingConfiguration\b/hipErrorMissingConfiguration/g; $ft{'numeric_literal'} += s/\bcudaErrorNoDevice\b/hipErrorNoDevice/g; $ft{'numeric_literal'} += s/\bcudaErrorNoKernelImageForDevice\b/hipErrorNoBinaryForGpu/g; + $ft{'numeric_literal'} += s/\bcudaErrorNotMapped\b/hipErrorNotMapped/g; + $ft{'numeric_literal'} += s/\bcudaErrorNotMappedAsArray\b/hipErrorNotMappedAsArray/g; + $ft{'numeric_literal'} += s/\bcudaErrorNotMappedAsPointer\b/hipErrorNotMappedAsPointer/g; $ft{'numeric_literal'} += s/\bcudaErrorNotReady\b/hipErrorNotReady/g; + $ft{'numeric_literal'} += s/\bcudaErrorNotSupported\b/hipErrorNotSupported/g; $ft{'numeric_literal'} += s/\bcudaErrorOperatingSystem\b/hipErrorOperatingSystem/g; $ft{'numeric_literal'} += s/\bcudaErrorPeerAccessAlreadyEnabled\b/hipErrorPeerAccessAlreadyEnabled/g; $ft{'numeric_literal'} += s/\bcudaErrorPeerAccessNotEnabled\b/hipErrorPeerAccessNotEnabled/g; @@ -1508,8 +1522,10 @@ sub simpleSubstitutions { $ft{'numeric_literal'} += s/\bcudaErrorProfilerAlreadyStopped\b/hipErrorProfilerAlreadyStopped/g; $ft{'numeric_literal'} += s/\bcudaErrorProfilerDisabled\b/hipErrorProfilerDisabled/g; $ft{'numeric_literal'} += s/\bcudaErrorProfilerNotInitialized\b/hipErrorProfilerNotInitialized/g; + $ft{'numeric_literal'} += s/\bcudaErrorSetOnActiveProcess\b/hipErrorSetOnActiveProcess/g; $ft{'numeric_literal'} += s/\bcudaErrorSharedObjectInitFailed\b/hipErrorSharedObjectInitFailed/g; $ft{'numeric_literal'} += s/\bcudaErrorSharedObjectSymbolNotFound\b/hipErrorSharedObjectSymbolNotFound/g; + $ft{'numeric_literal'} += s/\bcudaErrorSymbolNotFound\b/hipErrorNotFound/g; $ft{'numeric_literal'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g; $ft{'numeric_literal'} += s/\bcudaErrorUnmapBufferObjectFailed\b/hipErrorUnmapFailed/g; $ft{'numeric_literal'} += s/\bcudaErrorUnsupportedLimit\b/hipErrorUnsupportedLimit/g; diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 21401b0626..c4fbd1db6d 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -382,7 +382,7 @@ | typedef |`cudaError_enum` |***`hipError_t`*** | | 0 |*`CUDA_SUCCESS`* |*`hipSuccess`* | | 1 |*`CUDA_ERROR_INVALID_VALUE`* |*`hipErrorInvalidValue`* | -| 2 |*`CUDA_ERROR_OUT_OF_MEMORY`* |*`hipErrorMemoryAllocation`* | +| 2 |*`CUDA_ERROR_OUT_OF_MEMORY`* |*`hipErrorOutOfMemory`* | | 3 |*`CUDA_ERROR_NOT_INITIALIZED`* |*`hipErrorNotInitialized`* | | 4 |*`CUDA_ERROR_DEINITIALIZED`* |*`hipErrorDeinitialized`* | | 5 |*`CUDA_ERROR_PROFILER_DISABLED`* |*`hipErrorProfilerDisabled`* | @@ -416,7 +416,7 @@ | 302 |*`CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND`* |*`hipErrorSharedObjectSymbolNotFound`* | | 303 |*`CUDA_ERROR_SHARED_OBJECT_INIT_FAILED`* |*`hipErrorSharedObjectInitFailed`* | | 304 |*`CUDA_ERROR_OPERATING_SYSTEM`* |*`hipErrorOperatingSystem`* | -| 400 |*`CUDA_ERROR_INVALID_HANDLE`* |*`hipErrorInvalidResourceHandle`* | +| 400 |*`CUDA_ERROR_INVALID_HANDLE`* |*`hipErrorInvalidHandle`* | | 401 |*`CUDA_ERROR_ILLEGAL_STATE`* | | 10.0 | | 500 |*`CUDA_ERROR_NOT_FOUND`* |*`hipErrorNotFound`* | | 600 |*`CUDA_ERROR_NOT_READY`* |*`hipErrorNotReady`* | @@ -426,7 +426,7 @@ | 703 |*`CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING`* | | | 704 |*`CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED`* |*`hipErrorPeerAccessAlreadyEnabled`* | | 705 |*`CUDA_ERROR_PEER_ACCESS_NOT_ENABLED`* |*`hipErrorPeerAccessNotEnabled`* | -| 708 |*`CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE`* | | +| 708 |*`CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE`* |*`hipErrorSetOnActiveProcess`* | | 709 |*`CUDA_ERROR_CONTEXT_IS_DESTROYED`* | | | 710 |*`CUDA_ERROR_ASSERT`* |*`hipErrorAssert`* | | 711 |*`CUDA_ERROR_TOO_MANY_PEERS`* | | @@ -439,7 +439,7 @@ | 718 |*`CUDA_ERROR_INVALID_PC`* | | | 719 |*`CUDA_ERROR_LAUNCH_FAILED`* |*`hipErrorLaunchFailure`* | | 800 |*`CUDA_ERROR_NOT_PERMITTED`* | | -| 801 |*`CUDA_ERROR_NOT_SUPPORTED`* | | +| 801 |*`CUDA_ERROR_NOT_SUPPORTED`* |*`hipErrorNotSupported`* | | 802 |*`CUDA_ERROR_SYSTEM_NOT_READY`* | | 10.0 | | 803 |*`CUDA_ERROR_SYSTEM_DRIVER_MISMATCH`* | | 10.1 | | 804 |*`CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE`* | | 10.1 | @@ -452,7 +452,7 @@ | 906 |*`CUDA_ERROR_STREAM_CAPTURE_IMPLICIT`* | | 10.0 | | 907 |*`CUDA_ERROR_CAPTURED_EVENT`* | | 10.0 | | 908 |*`CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD`* | | 10.1 | -| 999 |*`CUDA_ERROR_UNKNOWN`* | | +| 999 |*`CUDA_ERROR_UNKNOWN`* |*`hipErrorUnknown`* | | enum |***`CUsharedconfig`*** |***`hipSharedMemConfig`*** | | typedef |***`CUsharedconfig_enum`*** |***`hipSharedMemConfig`*** | | 0x00 |*`CU_SHARED_MEM_CONFIG_DEFAULT_BANK_SIZE`* |*`hipSharedMemBankSizeDefault`* | diff --git a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index cbe4afb009..4fa033e21c 100644 --- a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -700,21 +700,17 @@ | enum |***`cudaError`*** | |***`hipError_t`*** | | typedef |***`cudaError_t`*** | |***`hipError_t`*** | | 0 |*`cudaSuccess`* | |*`hipSuccess`* | -| 1 |*`cudaErrorMissingConfiguration`* | |*`hipErrorMissingConfiguration`* | 1001 | +| 1 |*`cudaErrorInvalidValue`* | |*`hipErrorInvalidValue`* | 1011 | | 2 |*`cudaErrorMemoryAllocation`* | |*`hipErrorMemoryAllocation`* | 1002 | | 3 |*`cudaErrorInitializationError`* | |*`hipErrorInitializationError`* | 1003 | -| 4 |*`cudaErrorLaunchFailure`* | |*`hipErrorLaunchFailure`* | 1004 | -| 5 |*`cudaErrorPriorLaunchFailure`* | |*`hipErrorPriorLaunchFailure`* | 1005 | -| 6 |*`cudaErrorLaunchTimeout`* | |*`hipErrorLaunchTimeOut`* | 1006 | -| 7 |*`cudaErrorLaunchOutOfResources`* | |*`hipErrorLaunchOutOfResources`* | 1007 | -| 8 |*`cudaErrorInvalidDeviceFunction`* | |*`hipErrorInvalidDeviceFunction`* | 1008 | +| 4 |*`cudaErrorCudartUnloading`* | |*`hipErrorDeinitialized`* | +| 5 |*`cudaErrorProfilerDisabled`* | |*`hipErrorProfilerDisabled`* | +| 6 |*`cudaErrorProfilerNotInitialized`* | |*`hipErrorProfilerNotInitialized`* | +| 7 |*`cudaErrorProfilerAlreadyStarted`* | |*`hipErrorProfilerAlreadyStarted`* | +| 8 |*`cudaErrorProfilerAlreadyStopped`* | |*`hipErrorProfilerAlreadyStopped`* | | 9 |*`cudaErrorInvalidConfiguration`* | |*`hipErrorInvalidConfiguration`* | 1009 | -| 10 |*`cudaErrorInvalidDevice`* | |*`hipErrorInvalidDevice`* | 1010 | -| 11 |*`cudaErrorInvalidValue`* | |*`hipErrorInvalidValue`* | 1011 | | 12 |*`cudaErrorInvalidPitchValue`* | | | | 13 |*`cudaErrorInvalidSymbol`* | |*`hipErrorInvalidSymbol`* | 701 | -| 14 |*`cudaErrorMapBufferObjectFailed`* | |*`hipErrorMapFailed`* | 205 | -| 15 |*`cudaErrorUnmapBufferObjectFailed`* | |*`hipErrorUnmapFailed`* | 206 | | 16 |*`cudaErrorInvalidHostPointer`* | | | | 17 |*`cudaErrorInvalidDevicePointer`* | |*`hipErrorInvalidDevicePointer`* | 1017 | | 18 |*`cudaErrorInvalidTexture`* | | | @@ -728,63 +724,70 @@ | 26 |*`cudaErrorInvalidFilterSetting`* | | | | 27 |*`cudaErrorInvalidNormSetting`* | | | | 28 |*`cudaErrorMixedDeviceExecution`* | | | -| 29 |*`cudaErrorCudartUnloading`* | |*`hipErrorDeinitialized`* | 4 | -| 30 |*`cudaErrorUnknown`* | |*`hipErrorUnknown`* | 1030 | | 31 |*`cudaErrorNotYetImplemented`* | | | | 32 |*`cudaErrorMemoryValueTooLarge`* | | | -| 33 |*`cudaErrorInvalidResourceHandle`* | |*`hipErrorInvalidResourceHandle`* | 1033 | -| 34 |*`cudaErrorNotReady`* | |*`hipErrorNotReady`* | 1034 | | 35 |*`cudaErrorInsufficientDriver`* | |*`hipErrorInsufficientDriver`* | -| 36 |*`cudaErrorSetOnActiveProcess`* | | | | 37 |*`cudaErrorInvalidSurface`* | | | -| 38 |*`cudaErrorNoDevice`* | |*`hipErrorNoDevice`* | 1038 | -| 39 |*`cudaErrorECCUncorrectable`* | |*`hipErrorECCNotCorrectable`* | 214 | -| 40 |*`cudaErrorSharedObjectSymbolNotFound`* | |*`hipErrorSharedObjectSymbolNotFound`* | 302 | -| 41 |*`cudaErrorSharedObjectInitFailed`* | |*`hipErrorSharedObjectInitFailed`* | 303 | -| 42 |*`cudaErrorUnsupportedLimit`* | |*`hipErrorUnsupportedLimit`* | 215 | | 43 |*`cudaErrorDuplicateVariableName`* | | | | 44 |*`cudaErrorDuplicateTextureName`* | | | | 45 |*`cudaErrorDuplicateSurfaceName`* | | | | 46 |*`cudaErrorDevicesUnavailable`* | | | -| 47 |*`cudaErrorInvalidKernelImage`* | |*`hipErrorInvalidImage`* | 200 | -| 48 |*`cudaErrorNoKernelImageForDevice`* | |*`hipErrorNoBinaryForGpu`* | 209 | | 49 |*`cudaErrorIncompatibleDriverContext`* | | | -| 50 |*`cudaErrorPeerAccessAlreadyEnabled`* | |*`hipErrorPeerAccessAlreadyEnabled`* | 1050 | -| 51 |*`cudaErrorPeerAccessNotEnabled`* | |*`hipErrorPeerAccessNotEnabled`* | 1051 | -| 52 | | | | -| 53 | | | | -| 54 |*`cudaErrorDeviceAlreadyInUse`* | | | -| 55 |*`cudaErrorProfilerDisabled`* | |*`hipErrorProfilerDisabled`* | 5 | -| 56 |*`cudaErrorProfilerNotInitialized`* | |*`hipErrorProfilerNotInitialized`* | 6 | -| 57 |*`cudaErrorProfilerAlreadyStarted`* | |*`hipErrorProfilerAlreadyStarted`* | 7 | -| 58 |*`cudaErrorProfilerAlreadyStopped`* | |*`hipErrorProfilerAlreadyStopped`* | 8 | -| 59 |*`cudaErrorAssert`* | |*`hipErrorAssert`* | 1081 | -| 60 |*`cudaErrorTooManyPeers`* | | | -| 61 |*`cudaErrorHostMemoryAlreadyRegistered`* | |*`hipErrorHostMemoryAlreadyRegistered`* | 1061 | -| 62 |*`cudaErrorHostMemoryNotRegistered`* | |*`hipErrorHostMemoryNotRegistered`* | 1062 | -| 63 |*`cudaErrorOperatingSystem`* | |*`hipErrorOperatingSystem`* | 304 | -| 64 |*`cudaErrorPeerAccessUnsupported`* | |*`hipErrorPeerAccessUnsupported`* | 217 | +| 52 |*`cudaErrorMissingConfiguration`* | |*`hipErrorMissingConfiguration`* | 1001 | +| 53 |*`cudaErrorPriorLaunchFailure`* | |*`hipErrorPriorLaunchFailure`* | 1005 | | 65 |*`cudaErrorLaunchMaxDepthExceeded`* | | | | 66 |*`cudaErrorLaunchFileScopedTex`* | | | | 67 |*`cudaErrorLaunchFileScopedSurf`* | | | | 68 |*`cudaErrorSyncDepthExceeded`* | | | | 69 |*`cudaErrorLaunchPendingCountExceeded`* | | | -| 70 |*`cudaErrorNotPermitted`* | | | -| 71 |*`cudaErrorNotSupported`* | | | -| 72 |*`cudaErrorHardwareStackError`* | | | -| 73 |*`cudaErrorIllegalInstruction`* | | | -| 74 |*`cudaErrorMisalignedAddress`* | | | -| 75 |*`cudaErrorInvalidAddressSpace`* | | | -| 76 |*`cudaErrorInvalidPc`* | | | -| 77 |*`cudaErrorIllegalAddress`* | |*`hipErrorIllegalAddress`* | 700 | -| 78 |*`cudaErrorInvalidPtx`* | |*`hipErrorInvalidKernelFile`* | 218 | -| 79 |*`cudaErrorInvalidGraphicsContext`* | |*`hipErrorInvalidGraphicsContext`* | 219 | -| 80 |*`cudaErrorNvlinkUncorrectable`* | 8.0 | | -| 81 |*`cudaErrorJitCompilerNotFound`* | 9.0 | | -| 82 |*`cudaErrorCooperativeLaunchTooLarge`* | 9.0 | | -| 83 |*`cudaErrorSystemNotReady`* | 10.0 | | -| 84 |*`cudaErrorIllegalState`* | 10.0 | | +| 98 |*`cudaErrorInvalidDeviceFunction`* | |*`hipErrorInvalidDeviceFunction`* | 1008 | +| 100 |*`cudaErrorNoDevice`* | |*`hipErrorNoDevice`* | 1038 | +| 101 |*`cudaErrorInvalidDevice`* | |*`hipErrorInvalidDevice`* | 1010 | | 127 |*`cudaErrorStartupFailure`* | 10.0 | | +| 200 |*`cudaErrorInvalidKernelImage`* | |*`hipErrorInvalidImage`* | +| 201 |*`cudaErrorDeviceUninitilialized`* | |*`hipErrorInvalidContext`* | +| 205 |*`cudaErrorMapBufferObjectFailed`* | |*`hipErrorMapBufferObjectFailed`* | 1071 | +| 206 |*`cudaErrorUnmapBufferObjectFailed`* | |*`hipErrorUnmapFailed`* | +| 209 |*`cudaErrorNoKernelImageForDevice`* | |*`hipErrorNoBinaryForGpu`* | +| 214 |*`cudaErrorECCUncorrectable`* | |*`hipErrorECCNotCorrectable`* | +| 215 |*`cudaErrorUnsupportedLimit`* | |*`hipErrorUnsupportedLimit`* | +| 216 |*`cudaErrorDeviceAlreadyInUse`* | | | +| 217 |*`cudaErrorPeerAccessUnsupported`* | |*`hipErrorPeerAccessUnsupported`* | +| 218 |*`cudaErrorInvalidPtx`* | |*`hipErrorInvalidKernelFile`* | +| 219 |*`cudaErrorInvalidGraphicsContext`* | |*`hipErrorInvalidGraphicsContext`* | +| 220 |*`cudaErrorNvlinkUncorrectable`* | 8.0 | | +| 221 |*`cudaErrorJitCompilerNotFound`* | 9.0 | | +| 300 |*`cudaErrorInvalidSource`* | 10.1 |*`hipErrorInvalidSource`* | +| 301 |*`cudaErrorFileNotFound`* | 10.1 |*`hipErrorFileNotFound`* | +| 302 |*`cudaErrorSharedObjectSymbolNotFound`* | |*`hipErrorSharedObjectSymbolNotFound`* | +| 303 |*`cudaErrorSharedObjectInitFailed`* | |*`hipErrorSharedObjectInitFailed`* | +| 304 |*`cudaErrorOperatingSystem`* | |*`hipErrorOperatingSystem`* | +| 400 |*`cudaErrorInvalidResourceHandle`* | |*`hipErrorInvalidResourceHandle`* | 1033 | +| 401 |*`cudaErrorIllegalState`* | 10.0 | | +| 500 |*`cudaErrorSymbolNotFound`* | 10.1 |*`hipErrorNotFound`* | +| 600 |*`cudaErrorNotReady`* | |*`hipErrorNotReady`* | 1034 | +| 700 |*`cudaErrorIllegalAddress`* | |*`hipErrorIllegalAddress`* | +| 701 |*`cudaErrorLaunchOutOfResources`* | |*`hipErrorLaunchOutOfResources`* | 1007 | +| 702 |*`cudaErrorLaunchTimeout`* | |*`hipErrorLaunchTimeOut`* | 1006 | +| 703 |*`cudaErrorLaunchIncompatibleTexturing`* | | | +| 704 |*`cudaErrorPeerAccessAlreadyEnabled`* | |*`hipErrorPeerAccessAlreadyEnabled`* | 1050 | +| 705 |*`cudaErrorPeerAccessNotEnabled`* | |*`hipErrorPeerAccessNotEnabled`* | 1051 | +| 708 |*`cudaErrorSetOnActiveProcess`* | |*`hipErrorSetOnActiveProcess`* | 305 | +| 709 |*`cudaErrorContextIsDestroyed`* | | | +| 710 |*`cudaErrorAssert`* | |*`hipErrorAssert`* | 1081 | +| 711 |*`cudaErrorTooManyPeers`* | | | +| 712 |*`cudaErrorHostMemoryAlreadyRegistered`* | |*`hipErrorHostMemoryAlreadyRegistered`* | 1061 | +| 713 |*`cudaErrorHostMemoryNotRegistered`* | |*`hipErrorHostMemoryNotRegistered`* | 1062 | +| 714 |*`cudaErrorHardwareStackError`* | | | +| 715 |*`cudaErrorIllegalInstruction`* | | | +| 716 |*`cudaErrorMisalignedAddress`* | | | +| 717 |*`cudaErrorInvalidAddressSpace`* | | | +| 718 |*`cudaErrorInvalidPc`* | | | +| 719 |*`cudaErrorLaunchFailure`* | |*`hipErrorLaunchFailure`* | 1004 | +| 720 |*`cudaErrorCooperativeLaunchTooLarge`* | 9.0 | | +| 800 |*`cudaErrorNotPermitted`* | | | +| 801 |*`cudaErrorNotSupported`* | |*`hipErrorNotSupported`* | 1082 | +| 802 |*`cudaErrorSystemNotReady`* | 10.0 | | | 803 |*`cudaErrorSystemDriverMismatch`* | 10.0 | | | 804 |*`cudaErrorCompatNotSupportedOnDevice`* | 10.0 | | | 900 |*`cudaErrorStreamCaptureUnsupported`* | 10.0 | | @@ -796,6 +799,7 @@ | 906 |*`cudaErrorStreamCaptureImplicit`* | 10.0 | | | 907 |*`cudaErrorCapturedEvent`* | 10.0 | | | 908 |*`cudaErrorStreamCaptureWrongThread`* | 10.1 | | +| 999 |*`cudaErrorUnknown`* | |*`hipErrorUnknown`* | 1030 | | 10000 |*`cudaErrorApiFailureBase`* | | | | enum |***`cudaFuncCache`*** | |***`hipFuncCache_t`*** | | 0 |*`cudaFuncCachePreferNone`* | |*`hipFuncCachePreferNone`* | diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp index f88086f2cb..4603af69d0 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -1105,161 +1105,195 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CUresult", {"hipError_t", "", CONV_TYPE, API_DRIVER}}, {"cudaError_enum", {"hipError_t", "", CONV_TYPE, API_DRIVER}}, // CUresult enum values - // cudaSuccess = 0 + // cudaSuccess {"CUDA_SUCCESS", {"hipSuccess", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0 - // cudaErrorInvalidValue = 11 + // cudaErrorInvalidValue = 1, CUDA_ERROR_INVALID_VALUE = 1, hipErrorInvalidValue = 1011 + // TODO [HIP]: make hipErrorInvalidValue = 1 + // cudaErrorInvalidValue {"CUDA_ERROR_INVALID_VALUE", {"hipErrorInvalidValue", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 1 - // cudaErrorMemoryAllocation = 2 - {"CUDA_ERROR_OUT_OF_MEMORY", {"hipErrorMemoryAllocation", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 2 - // cudaErrorInitializationError = 3; hipErrorInitializationError = 1003 - // TODO: the same error code should be merged into 1 + // cudaErrorMemoryAllocation = 2, CUDA_ERROR_OUT_OF_MEMORY = 2, hipErrorOutOfMemory = 2, hipErrorMemoryAllocation = 1002 + // TODO [HIP]: remove hipErrorMemoryAllocation + // cudaErrorMemoryAllocation + {"CUDA_ERROR_OUT_OF_MEMORY", {"hipErrorOutOfMemory", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 2 + // cudaErrorInitializationError = 3, CUDA_ERROR_NOT_INITIALIZED = 3, hipErrorNotInitialized = 3, hipErrorInitializationError = 1003 + // TODO [HIP]: remove hipErrorInitializationError + // cudaErrorInitializationError {"CUDA_ERROR_NOT_INITIALIZED", {"hipErrorNotInitialized", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 3 - // cudaErrorCudartUnloading = 29 - // TODO: double check, that these errors match + // cudaErrorCudartUnloading {"CUDA_ERROR_DEINITIALIZED", {"hipErrorDeinitialized", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 4 - // cudaErrorProfilerDisabled = 55 + // cudaErrorProfilerDisabled {"CUDA_ERROR_PROFILER_DISABLED", {"hipErrorProfilerDisabled", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 5 - // cudaErrorProfilerNotInitialized = 56 + // cudaErrorProfilerNotInitialized // NOTE: Deprecated since CUDA 5.0 {"CUDA_ERROR_PROFILER_NOT_INITIALIZED", {"hipErrorProfilerNotInitialized", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 6 // NOTE: Deprecated since CUDA 5.0 - // cudaErrorProfilerAlreadyStarted = 57 + // cudaErrorProfilerAlreadyStarted {"CUDA_ERROR_PROFILER_ALREADY_STARTED", {"hipErrorProfilerAlreadyStarted", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 7 - // cudaErrorProfilerAlreadyStopped = 58 + // cudaErrorProfilerAlreadyStopped // NOTE: Deprecated since CUDA 5.0 {"CUDA_ERROR_PROFILER_ALREADY_STOPPED", {"hipErrorProfilerAlreadyStopped", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 8 - // cudaErrorNoDevice = 38 + // cudaErrorNoDevice = 100, CUDA_ERROR_NO_DEVICE = 100, hipErrorNoDevice = 1038 + // TODO [HIP]: make hipErrorNoDevice = 100 + // cudaErrorNoDevice {"CUDA_ERROR_NO_DEVICE", {"hipErrorNoDevice", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 100 - // cudaErrorInvalidDevice = 10 + // cudaErrorInvalidDevice = 101, CUDA_ERROR_INVALID_DEVICE = 101, hipErrorInvalidDevice = 1010 + // TODO [HIP]: make hipErrorInvalidDevice = 101 + // cudaErrorInvalidDevice {"CUDA_ERROR_INVALID_DEVICE", {"hipErrorInvalidDevice", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 101 - // cudaErrorInvalidKernelImage = 47 - // TODO: double check the matching + // cudaErrorInvalidKernelImage {"CUDA_ERROR_INVALID_IMAGE", {"hipErrorInvalidImage", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 200 - // no analogue + // cudaErrorDeviceUninitilialized {"CUDA_ERROR_INVALID_CONTEXT", {"hipErrorInvalidContext", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 201 // no analogue // NOTE: Deprecated since CUDA 3.2 {"CUDA_ERROR_CONTEXT_ALREADY_CURRENT", {"hipErrorContextAlreadyCurrent", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 202 - // cudaErrorMapBufferObjectFailed = 14 - // TODO: double check the matching + // cudaErrorMapBufferObjectFailed = 205, CUDA_ERROR_MAP_FAILED = 205, hipErrorMapFailed = 205, hipErrorMapBufferObjectFailed = 1071 + // TODO [HIP]: remove hipErrorMapBufferObjectFailed + // TODO [HIPIFY]: rename hipErrorMapBufferObjectFailed to hipErrorMapFailed + // cudaErrorMapBufferObjectFailed {"CUDA_ERROR_MAP_FAILED", {"hipErrorMapFailed", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 205 - // cudaErrorUnmapBufferObjectFailed = 15 - // TODO: double check the matching + // cudaErrorUnmapBufferObjectFailed {"CUDA_ERROR_UNMAP_FAILED", {"hipErrorUnmapFailed", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 206 - // no analogue + // cudaErrorArrayIsMapped {"CUDA_ERROR_ARRAY_IS_MAPPED", {"hipErrorArrayIsMapped", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 207 - // no analogue + // cudaErrorAlreadyMapped {"CUDA_ERROR_ALREADY_MAPPED", {"hipErrorAlreadyMapped", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 208 - // cudaErrorNoKernelImageForDevice = 48 - // TODO: double check the matching + // cudaErrorNoKernelImageForDevice {"CUDA_ERROR_NO_BINARY_FOR_GPU", {"hipErrorNoBinaryForGpu", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 209 - // no analogue + // cudaErrorAlreadyAcquired {"CUDA_ERROR_ALREADY_ACQUIRED", {"hipErrorAlreadyAcquired", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 210 - // no analogue + // cudaErrorNotMapped {"CUDA_ERROR_NOT_MAPPED", {"hipErrorNotMapped", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 211 - // no analogue + // cudaErrorNotMappedAsArray {"CUDA_ERROR_NOT_MAPPED_AS_ARRAY", {"hipErrorNotMappedAsArray", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 212 - // no analogue + // cudaErrorNotMappedAsPointer {"CUDA_ERROR_NOT_MAPPED_AS_POINTER", {"hipErrorNotMappedAsPointer", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 213 - // cudaErrorECCUncorrectable = 39 + // cudaErrorECCUncorrectable {"CUDA_ERROR_ECC_UNCORRECTABLE", {"hipErrorECCNotCorrectable", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 214 - // cudaErrorUnsupportedLimit = 42 + // cudaErrorUnsupportedLimit {"CUDA_ERROR_UNSUPPORTED_LIMIT", {"hipErrorUnsupportedLimit", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 215 - // no analogue + // cudaErrorDeviceAlreadyInUse {"CUDA_ERROR_CONTEXT_ALREADY_IN_USE", {"hipErrorContextAlreadyInUse", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 216 - // cudaErrorPeerAccessUnsupported = 64 + // cudaErrorPeerAccessUnsupported {"CUDA_ERROR_PEER_ACCESS_UNSUPPORTED", {"hipErrorPeerAccessUnsupported", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 217 - // cudaErrorInvalidPtx = 78 + // cudaErrorInvalidPtx {"CUDA_ERROR_INVALID_PTX", {"hipErrorInvalidKernelFile", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 218 - // cudaErrorInvalidGraphicsContext = 79 + // cudaErrorInvalidGraphicsContext {"CUDA_ERROR_INVALID_GRAPHICS_CONTEXT", {"hipErrorInvalidGraphicsContext", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 219 - // cudaErrorNvlinkUncorrectable = 80 + // cudaErrorNvlinkUncorrectable {"CUDA_ERROR_NVLINK_UNCORRECTABLE", {"hipErrorNvlinkUncorrectable", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 220 - // cudaErrorJitCompilerNotFound = 81 + // cudaErrorJitCompilerNotFound {"CUDA_ERROR_JIT_COMPILER_NOT_FOUND", {"hipErrorJitCompilerNotFound", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 221 - // no analogue + // cudaErrorInvalidSource {"CUDA_ERROR_INVALID_SOURCE", {"hipErrorInvalidSource", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 300 - // no analogue + // cudaErrorFileNotFound {"CUDA_ERROR_FILE_NOT_FOUND", {"hipErrorFileNotFound", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 301 - // cudaErrorSharedObjectSymbolNotFound = 40 + // cudaErrorSharedObjectSymbolNotFound {"CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND", {"hipErrorSharedObjectSymbolNotFound", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 302 - // cudaErrorSharedObjectInitFailed = 41 + // cudaErrorSharedObjectInitFailed {"CUDA_ERROR_SHARED_OBJECT_INIT_FAILED", {"hipErrorSharedObjectInitFailed", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 303 - // cudaErrorOperatingSystem = 63 + // cudaErrorOperatingSystem {"CUDA_ERROR_OPERATING_SYSTEM", {"hipErrorOperatingSystem", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 304 - // cudaErrorInvalidResourceHandle = 33 - {"CUDA_ERROR_INVALID_HANDLE", {"hipErrorInvalidResourceHandle", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 400 - // cudaErrorIllegalState = 84 + // cudaErrorInvalidResourceHandle = 400, CUDA_ERROR_INVALID_HANDLE = 400, hipErrorInvalidHandle = 400, hipErrorInvalidResourceHandle = 1033 + // TODO [HIP]: remove hipErrorInvalidResourceHandle + // TODO [HIPIFY]: rename hipErrorInvalidResourceHandle to hipErrorInvalidHandle + // cudaErrorInvalidResourceHandle + {"CUDA_ERROR_INVALID_HANDLE", {"hipErrorInvalidHandle", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 400 + // cudaErrorIllegalState {"CUDA_ERROR_ILLEGAL_STATE", {"hipErrorIllegalState", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 401 - // no analogue + // cudaErrorSymbolNotFound {"CUDA_ERROR_NOT_FOUND", {"hipErrorNotFound", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 500 - // cudaErrorNotReady = 34 + // cudaErrorNotReady = 600, CUDA_ERROR_NOT_READY = 600, hipErrorNotReady = 1034 + // TODO [HIP]: make hipErrorNotReady = 600 + // cudaErrorNotReady {"CUDA_ERROR_NOT_READY", {"hipErrorNotReady", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 600 - // cudaErrorIllegalAddress = 77 + // cudaErrorIllegalAddress {"CUDA_ERROR_ILLEGAL_ADDRESS", {"hipErrorIllegalAddress", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 700 - // cudaErrorLaunchOutOfResources = 7 + // cudaErrorLaunchOutOfResources = 701, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701, hipErrorLaunchOutOfResources = 1007 + // TODO [HIP]: make hipErrorLaunchOutOfResources = 701 + // cudaErrorLaunchOutOfResources {"CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES", {"hipErrorLaunchOutOfResources", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 701 - // cudaErrorLaunchTimeout = 6 + // cudaErrorLaunchTimeout = 702, CUDA_ERROR_LAUNCH_TIMEOUT = 702, hipErrorLaunchTimeOut = 1006 + // TODO [HIP]: make hipErrorLaunchTimeOut = 702 + // cudaErrorLaunchTimeout {"CUDA_ERROR_LAUNCH_TIMEOUT", {"hipErrorLaunchTimeOut", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 702 - // no analogue + // cudaErrorLaunchIncompatibleTexturing {"CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING", {"hipErrorLaunchIncompatibleTexturing", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 703 - // cudaErrorPeerAccessAlreadyEnabled = 50 + // cudaErrorPeerAccessAlreadyEnabled = 704, CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED = 704, hipErrorPeerAccessAlreadyEnabled = 1050 + // TODO [HIP]: make hipErrorPeerAccessAlreadyEnabled = 704 + // cudaErrorPeerAccessAlreadyEnabled {"CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED", {"hipErrorPeerAccessAlreadyEnabled", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 704 - // cudaErrorPeerAccessNotEnabled = 51 + // cudaErrorPeerAccessNotEnabled = 705, CUDA_ERROR_PEER_ACCESS_NOT_ENABLED = 705, hipErrorPeerAccessNotEnabled = 1051 + // TODO [HIP]: make hipErrorPeerAccessNotEnabled = 705 + // cudaErrorPeerAccessNotEnabled {"CUDA_ERROR_PEER_ACCESS_NOT_ENABLED", {"hipErrorPeerAccessNotEnabled", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 705 - // no analogue - {"CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE", {"hipErrorPrimaryContextActive", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 708 - // no analogue + // cudaErrorSetOnActiveProcess = 708, CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE = 708, hipErrorSetOnActiveProcess = 305 + // TODO [HIP]: make hipErrorSetOnActiveProcess = 708 + // cudaErrorSetOnActiveProcess + {"CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE", {"hipErrorSetOnActiveProcess", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 708 + // cudaErrorContextIsDestroyed {"CUDA_ERROR_CONTEXT_IS_DESTROYED", {"hipErrorContextIsDestroyed", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 709 - // cudaErrorAssert = 59 + // cudaErrorAssert = 710, CUDA_ERROR_ASSERT = 710, hipErrorAssert = 1081 + // TODO [HIP]: make hipErrorAssert = 710 + // cudaErrorAssert {"CUDA_ERROR_ASSERT", {"hipErrorAssert", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 710 - // cudaErrorTooManyPeers = 60 + // cudaErrorTooManyPeers {"CUDA_ERROR_TOO_MANY_PEERS", {"hipErrorTooManyPeers", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 711 - // = cudaErrorHostMemoryAlreadyRegistered = 61 + // cudaErrorHostMemoryAlreadyRegistered = 712, CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED = 712, hipErrorHostMemoryAlreadyRegistered = 1061 + // TODO [HIP]: make hipErrorHostMemoryAlreadyRegistered = 712 + // cudaErrorHostMemoryAlreadyRegistered {"CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED", {"hipErrorHostMemoryAlreadyRegistered", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 712 - // cudaErrorHostMemoryNotRegistered = 62 + // cudaErrorHostMemoryNotRegistered = 713, CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED = 713, hipErrorHostMemoryNotRegistered = 1062 + // TODO [HIP]: make hipErrorHostMemoryNotRegistered = 713 + // cudaErrorHostMemoryNotRegistered {"CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED", {"hipErrorHostMemoryNotRegistered", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 713 - // cudaErrorHardwareStackError= 72 + // cudaErrorHardwareStackError {"CUDA_ERROR_HARDWARE_STACK_ERROR", {"hipErrorHardwareStackError", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 714 - // cudaErrorIllegalInstruction = 73 + // cudaErrorIllegalInstruction {"CUDA_ERROR_ILLEGAL_INSTRUCTION", {"hipErrorIllegalInstruction", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 715 - // cudaErrorMisalignedAddress= 74 + // cudaErrorMisalignedAddress {"CUDA_ERROR_MISALIGNED_ADDRESS", {"hipErrorMisalignedAddress", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 716 - // cudaErrorInvalidAddressSpace= 75 + // cudaErrorInvalidAddressSpace {"CUDA_ERROR_INVALID_ADDRESS_SPACE", {"hipErrorInvalidAddressSpace", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 717 - // cudaErrorInvalidPc = 76 + // cudaErrorInvalidPc {"CUDA_ERROR_INVALID_PC", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 718 - // cudaErrorLaunchFailure = 4 + // cudaErrorLaunchFailure = 719, CUDA_ERROR_LAUNCH_FAILED = 719, hipErrorLaunchFailure = 1004 + // TODO [HIP]: make hipErrorSetOnActiveProcess = 719 + // cudaErrorLaunchFailure {"CUDA_ERROR_LAUNCH_FAILED", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 719 - // cudaErrorNotPermitted = 70 + // cudaErrorNotPermitted {"CUDA_ERROR_NOT_PERMITTED", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 800 - // cudaErrorNotSupported = 71 - {"CUDA_ERROR_NOT_SUPPORTED", {"hipErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 801 - // cudaErrorSystemNotReady = 83 + // cudaErrorNotSupported = 801, CUDA_ERROR_NOT_SUPPORTED = 801, hipErrorNotSupported = 1082 + // TODO [HIP]: make hipErrorNotSupported = 801 + // cudaErrorNotSupported + {"CUDA_ERROR_NOT_SUPPORTED", {"hipErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 801 + // cudaErrorSystemNotReady {"CUDA_ERROR_SYSTEM_NOT_READY", {"hipErrorSystemNotReady", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 802 - // cudaErrorSystemDriverMismatch = 803 + // cudaErrorSystemDriverMismatch {"CUDA_ERROR_SYSTEM_DRIVER_MISMATCH", {"hipErrorSystemDriverMismatch", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 803 - // cudaErrorCompatNotSupportedOnDevice = 804 + // cudaErrorCompatNotSupportedOnDevice {"CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE", {"hipErrorCompatNotSupportedOnDevice", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 804 - // cudaErrorStreamCaptureUnsupported = 900 + // cudaErrorStreamCaptureUnsupported {"CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED", {"hipErrorStreamCaptureUnsupported", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 900 - // cudaErrorStreamCaptureInvalidated = 901 + // cudaErrorStreamCaptureInvalidated {"CUDA_ERROR_STREAM_CAPTURE_INVALIDATED", {"hipErrorStreamCaptureInvalidated", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 901 - // cudaErrorStreamCaptureMerge = 902 + // cudaErrorStreamCaptureMerge {"CUDA_ERROR_STREAM_CAPTURE_MERGE", {"hipErrorStreamCaptureMerge", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 902 - // cudaErrorStreamCaptureUnmatched = 903 + // cudaErrorStreamCaptureUnmatched {"CUDA_ERROR_STREAM_CAPTURE_UNMATCHED", {"hipErrorStreamCaptureUnmatched", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 903 - // cudaErrorStreamCaptureUnjoined = 904 + // cudaErrorStreamCaptureUnjoined {"CUDA_ERROR_STREAM_CAPTURE_UNJOINED", {"hipErrorStreamCaptureUnjoined", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 904 - // cudaErrorStreamCaptureIsolation = 905 + // cudaErrorStreamCaptureIsolation {"CUDA_ERROR_STREAM_CAPTURE_ISOLATION", {"hipErrorStreamCaptureIsolation", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 905 - // cudaErrorStreamCaptureImplicit = 906 + // cudaErrorStreamCaptureImplicit {"CUDA_ERROR_STREAM_CAPTURE_IMPLICIT", {"hipErrorStreamCaptureImplicit", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 906 - // cudaErrorCapturedEvent = 907 + // cudaErrorCapturedEvent {"CUDA_ERROR_CAPTURED_EVENT", {"hipErrorCapturedEvent", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 907 - // cudaErrorStreamCaptureWrongThread = 908 + // cudaErrorStreamCaptureWrongThread {"CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD", {"hipErrorStreamCaptureWrongThread", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 908 - // cudaErrorUnknown = 30 + // cudaErrorUnknown = 999, CUDA_ERROR_UNKNOWN = 999, hipErrorUnknown = 1030 + // TODO [HIP]: make hipErrorUnknown = 999 + // cudaErrorUnknown {"CUDA_ERROR_UNKNOWN", {"hipErrorUnknown", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 999 // cudaSharedMemConfig diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp index 7b34d97ab4..d8acb9882c 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp @@ -595,42 +595,51 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaError", {"hipError_t", "", CONV_TYPE, API_RUNTIME}}, {"cudaError_t", {"hipError_t", "", CONV_TYPE, API_RUNTIME}}, // cudaError enum values - // CUDA_SUCCESS = 0 + // CUDA_SUCCESS {"cudaSuccess", {"hipSuccess", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 0 - // no analogue - {"cudaErrorMissingConfiguration", {"hipErrorMissingConfiguration", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 1 - // CUDA_ERROR_OUT_OF_MEMORY = 2 + // cudaErrorInvalidValue = 1, CUDA_ERROR_INVALID_VALUE = 1, hipErrorInvalidValue = 1011 + // TODO [HIP]: make hipErrorInvalidValue = 1 + // CUDA_ERROR_INVALID_VALUE + {"cudaErrorInvalidValue", {"hipErrorInvalidValue", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 1 + // cudaErrorMemoryAllocation = 2, CUDA_ERROR_OUT_OF_MEMORY = 2, hipErrorOutOfMemory = 2, hipErrorMemoryAllocation = 1002 + // TODO [HIP]: remove hipErrorMemoryAllocation + // TODO [HIPIFY]: rename hipErrorMemoryAllocation to hipErrorOutOfMemory + // CUDA_ERROR_OUT_OF_MEMORY {"cudaErrorMemoryAllocation", {"hipErrorMemoryAllocation", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 2 - // CUDA_ERROR_NOT_INITIALIZED = 3 + // cudaErrorInitializationError = 3, CUDA_ERROR_NOT_INITIALIZED = 3, hipErrorNotInitialized = 3, hipErrorInitializationError = 1003 + // TODO [HIP]: remove hipErrorInitializationError + // TODO [HIPIFY]: rename hipErrorInitializationError to hipErrorNotInitialized + // CUDA_ERROR_NOT_INITIALIZED {"cudaErrorInitializationError", {"hipErrorInitializationError", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 3 - // CUDA_ERROR_LAUNCH_FAILED = 719 - {"cudaErrorLaunchFailure", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 4 - // no analogue - {"cudaErrorPriorLaunchFailure", {"hipErrorPriorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 5 - // CUDA_ERROR_LAUNCH_TIMEOUT = 702 - {"cudaErrorLaunchTimeout", {"hipErrorLaunchTimeOut", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 6 - // CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701 - {"cudaErrorLaunchOutOfResources", {"hipErrorLaunchOutOfResources", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 7 - // no analogue - {"cudaErrorInvalidDeviceFunction", {"hipErrorInvalidDeviceFunction", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 8 + // CUDA_ERROR_DEINITIALIZED + {"cudaErrorCudartUnloading", {"hipErrorDeinitialized", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 4 + // CUDA_ERROR_PROFILER_DISABLED + {"cudaErrorProfilerDisabled", {"hipErrorProfilerDisabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 5 + // Deprecated since CUDA 5.0 + // CUDA_ERROR_PROFILER_NOT_INITIALIZED + {"cudaErrorProfilerNotInitialized", {"hipErrorProfilerNotInitialized", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 6 + // Deprecated since CUDA 5.0 + // CUDA_ERROR_PROFILER_ALREADY_STARTED + {"cudaErrorProfilerAlreadyStarted", {"hipErrorProfilerAlreadyStarted", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 7 + // Deprecated since CUDA 5.0 + // CUDA_ERROR_PROFILER_ALREADY_STOPPED + {"cudaErrorProfilerAlreadyStopped", {"hipErrorProfilerAlreadyStopped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 8 + // cudaErrorInvalidConfiguration = 9, hipErrorInvalidConfiguration = 1009 + // TODO [HIP]: make hipErrorInvalidConfiguration = 9 // no analogue {"cudaErrorInvalidConfiguration", {"hipErrorInvalidConfiguration", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 9 - // CUDA_ERROR_INVALID_DEVICE = 101 - {"cudaErrorInvalidDevice", {"hipErrorInvalidDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 10 - // CUDA_ERROR_INVALID_VALUE = 1 - {"cudaErrorInvalidValue", {"hipErrorInvalidValue", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 11 // no analogue {"cudaErrorInvalidPitchValue", {"hipErrorInvalidPitchValue", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 12 + // cudaErrorInvalidSymbol = 13, hipErrorInvalidSymbol = 701, but also cudaErrorLaunchOutOfResources = CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701 + // TODO [HIP]: make hipErrorInvalidSymbol = 13 // no analogue {"cudaErrorInvalidSymbol", {"hipErrorInvalidSymbol", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 13 - // CUDA_ERROR_MAP_FAILED = 205 - // TODO: double check the matching - {"cudaErrorMapBufferObjectFailed", {"hipErrorMapFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 14 - // CUDA_ERROR_UNMAP_FAILED = 206 - // TODO: double check the matching - {"cudaErrorUnmapBufferObjectFailed", {"hipErrorUnmapFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 15 + // Deprecated since CUDA 10.1 // no analogue {"cudaErrorInvalidHostPointer", {"hipErrorInvalidHostPointer", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 16 + // cudaErrorInvalidDevicePointer = 17, hipErrorInvalidDevicePointer = 1017 + // TODO [HIP]: make hipErrorInvalidDevicePointer = 17 + // Deprecated since CUDA 10.1 // no analogue {"cudaErrorInvalidDevicePointer", {"hipErrorInvalidDevicePointer", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 17 // no analogue @@ -639,6 +648,8 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaErrorInvalidTextureBinding", {"hipErrorInvalidTextureBinding", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 19 // no analogue {"cudaErrorInvalidChannelDescriptor", {"hipErrorInvalidChannelDescriptor", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 20 + // cudaErrorInvalidMemcpyDirection = 21, hipErrorInvalidMemcpyDirection = 1021 + // TODO [HIP]: make hipErrorInvalidMemcpyDirection = 21 // no analogue {"cudaErrorInvalidMemcpyDirection", {"hipErrorInvalidMemcpyDirection", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 21 // no analogue @@ -655,36 +666,16 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaErrorInvalidNormSetting", {"hipErrorInvalidNormSetting", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 27 // no analogue {"cudaErrorMixedDeviceExecution", {"hipErrorMixedDeviceExecution", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 28 - // CUDA_ERROR_DEINITIALIZED = 4 - {"cudaErrorCudartUnloading", {"hipErrorDeinitialized", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 29 - // CUDA_ERROR_UNKNOWN = 999 - {"cudaErrorUnknown", {"hipErrorUnknown", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 30 // Deprecated since CUDA 4.1 // no analogue {"cudaErrorNotYetImplemented", {"hipErrorNotYetImplemented", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 31 // Deprecated since CUDA 3.1 // no analogue {"cudaErrorMemoryValueTooLarge", {"hipErrorMemoryValueTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 32 - // CUDA_ERROR_INVALID_HANDLE = 400 - {"cudaErrorInvalidResourceHandle", {"hipErrorInvalidResourceHandle", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 33 - // CUDA_ERROR_NOT_READY = 600 - {"cudaErrorNotReady", {"hipErrorNotReady", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 34 // no analogue {"cudaErrorInsufficientDriver", {"hipErrorInsufficientDriver", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 35 // no analogue - {"cudaErrorSetOnActiveProcess", {"hipErrorSetOnActiveProcess", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 36 - // no analogue {"cudaErrorInvalidSurface", {"hipErrorInvalidSurface", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 37 - // CUDA_ERROR_NO_DEVICE = 100 - {"cudaErrorNoDevice", {"hipErrorNoDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 38 - // CUDA_ERROR_ECC_UNCORRECTABLE = 214 - {"cudaErrorECCUncorrectable", {"hipErrorECCNotCorrectable", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 39 - // CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND = 302 - {"cudaErrorSharedObjectSymbolNotFound", {"hipErrorSharedObjectSymbolNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 40 - // CUDA_ERROR_SHARED_OBJECT_INIT_FAILED = 303 - {"cudaErrorSharedObjectInitFailed", {"hipErrorSharedObjectInitFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 41 - // CUDA_ERROR_UNSUPPORTED_LIMIT = 215 - {"cudaErrorUnsupportedLimit", {"hipErrorUnsupportedLimit", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 42 // no analogue {"cudaErrorDuplicateVariableName", {"hipErrorDuplicateVariableName", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 43 // no analogue @@ -693,41 +684,16 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaErrorDuplicateSurfaceName", {"hipErrorDuplicateSurfaceName", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 45 // no analogue {"cudaErrorDevicesUnavailable", {"hipErrorDevicesUnavailable", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 46 - // CUDA_ERROR_INVALID_IMAGE = 200 - {"cudaErrorInvalidKernelImage", {"hipErrorInvalidImage", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 47 - // CUDA_ERROR_NO_BINARY_FOR_GPU = 209 - {"cudaErrorNoKernelImageForDevice", {"hipErrorNoBinaryForGpu", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 48 // no analogue {"cudaErrorIncompatibleDriverContext", {"hipErrorIncompatibleDriverContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 49 - // CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED = 704 - {"cudaErrorPeerAccessAlreadyEnabled", {"hipErrorPeerAccessAlreadyEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 50 - // CUDA_ERROR_PEER_ACCESS_NOT_ENABLED = 705 - {"cudaErrorPeerAccessNotEnabled", {"hipErrorPeerAccessNotEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 51 + // cudaErrorMissingConfiguration = 52, hipErrorMissingConfiguration = 1001 + // TODO [HIP]: make hipErrorMissingConfiguration = 52 // no analogue - {"cudaErrorDeviceAlreadyInUse", {"hipErrorDeviceAlreadyInUse", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 54 - // CUDA_ERROR_PROFILER_DISABLED = 5 - {"cudaErrorProfilerDisabled", {"hipErrorProfilerDisabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 55 - // Deprecated since CUDA 5.0 - // CUDA_ERROR_PROFILER_NOT_INITIALIZED = 6 - {"cudaErrorProfilerNotInitialized", {"hipErrorProfilerNotInitialized", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 56 - // Deprecated since CUDA 5.0 - // CUDA_ERROR_PROFILER_ALREADY_STARTED = 7 - {"cudaErrorProfilerAlreadyStarted", {"hipErrorProfilerAlreadyStarted", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 57 - // Deprecated since CUDA 5.0 - // CUDA_ERROR_PROFILER_ALREADY_STOPPED = 8 - {"cudaErrorProfilerAlreadyStopped", {"hipErrorProfilerAlreadyStopped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 58 - // CUDA_ERROR_ASSERT = 710 - {"cudaErrorAssert", {"hipErrorAssert", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 59 - // CUDA_ERROR_TOO_MANY_PEERS = 711 - {"cudaErrorTooManyPeers", {"hipErrorTooManyPeers", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 60 - // CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED = 712 - {"cudaErrorHostMemoryAlreadyRegistered", {"hipErrorHostMemoryAlreadyRegistered", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 61 - // CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED = 713 - {"cudaErrorHostMemoryNotRegistered", {"hipErrorHostMemoryNotRegistered", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 62 - // CUDA_ERROR_OPERATING_SYSTEM = 304 - {"cudaErrorOperatingSystem", {"hipErrorOperatingSystem", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 63 - // CUDA_ERROR_PEER_ACCESS_UNSUPPORTED = 217 - {"cudaErrorPeerAccessUnsupported", {"hipErrorPeerAccessUnsupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 64 + {"cudaErrorMissingConfiguration", {"hipErrorMissingConfiguration", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 52 + // cudaErrorPriorLaunchFailure = 53, hipErrorPriorLaunchFailure = 1005 + // TODO [HIP]: make hipErrorPriorLaunchFailure = 53 + // no analogue + {"cudaErrorPriorLaunchFailure", {"hipErrorPriorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 53 // no analogue {"cudaErrorLaunchMaxDepthExceeded", {"hipErrorLaunchMaxDepthExceeded", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 65 // no analogue @@ -738,60 +704,174 @@ const std::map CUDA_RUNTIME_TYPE_NAME_MAP { {"cudaErrorSyncDepthExceeded", {"hipErrorSyncDepthExceeded", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 68 // no analogue {"cudaErrorLaunchPendingCountExceeded", {"hipErrorLaunchPendingCountExceeded", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 69 - // CUDA_ERROR_NOT_PERMITTED = 800 - {"cudaErrorNotPermitted", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 70 - // CUDA_ERROR_NOT_SUPPORTED = 801 - {"cudaErrorNotSupported", {"hipErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 71 - // CUDA_ERROR_HARDWARE_STACK_ERROR = 714 - {"cudaErrorHardwareStackError", {"hipErrorHardwareStackError", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 72 - // CUDA_ERROR_ILLEGAL_INSTRUCTION = 715 - {"cudaErrorIllegalInstruction", {"hipErrorIllegalInstruction", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 73 - // CUDA_ERROR_MISALIGNED_ADDRESS = 716 - {"cudaErrorMisalignedAddress", {"hipErrorMisalignedAddress", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 74 - // CUDA_ERROR_INVALID_ADDRESS_SPACE = 717 - {"cudaErrorInvalidAddressSpace", {"hipErrorInvalidAddressSpace", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 75 - // CUDA_ERROR_INVALID_PC = 718 - {"cudaErrorInvalidPc", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 76 - // CUDA_ERROR_ILLEGAL_ADDRESS = 700 - {"cudaErrorIllegalAddress", {"hipErrorIllegalAddress", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 77 - // CUDA_ERROR_INVALID_PTX = 218 - {"cudaErrorInvalidPtx", {"hipErrorInvalidKernelFile", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 78 - // CUDA_ERROR_INVALID_GRAPHICS_CONTEXT = 219 - {"cudaErrorInvalidGraphicsContext", {"hipErrorInvalidGraphicsContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 79 - // CUDA_ERROR_NVLINK_UNCORRECTABLE = 220 - {"cudaErrorNvlinkUncorrectable", {"hipErrorNvlinkUncorrectable", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 80 + // cudaErrorInvalidDeviceFunction = 98, hipErrorInvalidDeviceFunction = 1008 + // TODO [HIP]: make hipErrorInvalidDeviceFunction = 98 // no analogue - {"cudaErrorJitCompilerNotFound", {"hipErrorJitCompilerNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 81 + {"cudaErrorInvalidDeviceFunction", {"hipErrorInvalidDeviceFunction", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 98 + // cudaErrorNoDevice = 100, CUDA_ERROR_NO_DEVICE = 100, hipErrorNoDevice = 1038 + // TODO [HIP]: make hipErrorNoDevice = 100 + // CUDA_ERROR_NO_DEVICE + {"cudaErrorNoDevice", {"hipErrorNoDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 100 + // cudaErrorInvalidDevice = 101, CUDA_ERROR_INVALID_DEVICE = 101, hipErrorInvalidDevice = 1010 + // TODO [HIP]: make hipErrorInvalidDevice = 101 + // CUDA_ERROR_INVALID_DEVICE + {"cudaErrorInvalidDevice", {"hipErrorInvalidDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 101 // no analogue - {"cudaErrorCooperativeLaunchTooLarge", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 82 - // CUDA_ERROR_SYSTEM_NOT_READY = 802 - {"cudaErrorSystemNotReady", {"hipErrorSystemNotReady", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 83 - // CUDA_ERROR_ILLEGAL_STATE = 401 - {"cudaErrorIllegalState", {"hipErrorIllegalState", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 84 + {"cudaErrorStartupFailure", {"hipErrorStartupFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 127 + // CUDA_ERROR_INVALID_IMAGE + {"cudaErrorInvalidKernelImage", {"hipErrorInvalidImage", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 200 + // CUDA_ERROR_INVALID_CONTEXT + {"cudaErrorDeviceUninitilialized", {"hipErrorInvalidContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 201 + // cudaErrorMapBufferObjectFailed = 205, CUDA_ERROR_MAP_FAILED = 205, hipErrorMapFailed = 205, hipErrorMapBufferObjectFailed = 1071 + // TODO [HIP]: remove hipErrorMapBufferObjectFailed + // TODO [HIPIFY]: rename hipErrorMapBufferObjectFailed to hipErrorMapFailed + // CUDA_ERROR_MAP_FAILED + {"cudaErrorMapBufferObjectFailed", {"hipErrorMapBufferObjectFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 205 + // CUDA_ERROR_UNMAP_FAILED + {"cudaErrorUnmapBufferObjectFailed", {"hipErrorUnmapFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 206 + // CUDA_ERROR_ARRAY_IS_MAPPED + {"cudaErrorArrayIsMapped", {"hipErrorArrayIsMapped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 207 + // CUDA_ERROR_ALREADY_MAPPED + {"cudaErrorAlreadyMapped", {"hipErrorAlreadyMapped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 208 + // CUDA_ERROR_NO_BINARY_FOR_GPU + {"cudaErrorNoKernelImageForDevice", {"hipErrorNoBinaryForGpu", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 209 + // CUDA_ERROR_ALREADY_ACQUIRED + {"cudaErrorAlreadyAcquired", {"hipErrorAlreadyAcquired", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 210 + // CUDA_ERROR_NOT_MAPPED + {"cudaErrorNotMapped", {"hipErrorNotMapped", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 211 + // CUDA_ERROR_NOT_MAPPED_AS_ARRAY + {"cudaErrorNotMappedAsArray", {"hipErrorNotMappedAsArray", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 212 + // CUDA_ERROR_NOT_MAPPED_AS_POINTER + {"cudaErrorNotMappedAsPointer", {"hipErrorNotMappedAsPointer", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 213 + // CUDA_ERROR_ECC_UNCORRECTABLE + {"cudaErrorECCUncorrectable", {"hipErrorECCNotCorrectable", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 214 + // CUDA_ERROR_UNSUPPORTED_LIMIT + {"cudaErrorUnsupportedLimit", {"hipErrorUnsupportedLimit", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 215 + // CUDA_ERROR_CONTEXT_ALREADY_IN_USE + {"cudaErrorDeviceAlreadyInUse", {"hipErrorContextAlreadyInUse", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 216 + // CUDA_ERROR_PEER_ACCESS_UNSUPPORTED + {"cudaErrorPeerAccessUnsupported", {"hipErrorPeerAccessUnsupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 217 + // CUDA_ERROR_INVALID_PTX + {"cudaErrorInvalidPtx", {"hipErrorInvalidKernelFile", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 218 + // CUDA_ERROR_INVALID_GRAPHICS_CONTEXT + {"cudaErrorInvalidGraphicsContext", {"hipErrorInvalidGraphicsContext", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 219 + // CUDA_ERROR_NVLINK_UNCORRECTABLE + {"cudaErrorNvlinkUncorrectable", {"hipErrorNvlinkUncorrectable", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 220 + // CUDA_ERROR_JIT_COMPILER_NOT_FOUND + {"cudaErrorJitCompilerNotFound", {"hipErrorJitCompilerNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 221 + // CUDA_ERROR_INVALID_SOURCE + {"cudaErrorInvalidSource", {"hipErrorInvalidSource", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 300 + // CUDA_ERROR_FILE_NOT_FOUND + {"cudaErrorFileNotFound", {"hipErrorFileNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 301 + // CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND + {"cudaErrorSharedObjectSymbolNotFound", {"hipErrorSharedObjectSymbolNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 302 + // CUDA_ERROR_SHARED_OBJECT_INIT_FAILED + {"cudaErrorSharedObjectInitFailed", {"hipErrorSharedObjectInitFailed", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 303 + // CUDA_ERROR_OPERATING_SYSTEM + {"cudaErrorOperatingSystem", {"hipErrorOperatingSystem", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 304 + // cudaErrorInvalidResourceHandle = 400, CUDA_ERROR_INVALID_HANDLE = 400, hipErrorInvalidHandle = 400, hipErrorInvalidResourceHandle = 1033 + // TODO [HIP]: remove hipErrorInvalidResourceHandle + // TODO [HIPIFY]: rename hipErrorInvalidResourceHandle to hipErrorInvalidHandle + // CUDA_ERROR_INVALID_HANDLE + {"cudaErrorInvalidResourceHandle", {"hipErrorInvalidResourceHandle", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 400 + // CUDA_ERROR_ILLEGAL_STATE + {"cudaErrorIllegalState", {"hipErrorIllegalState", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 401 + // CUDA_ERROR_NOT_FOUND + {"cudaErrorSymbolNotFound", {"hipErrorNotFound", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 500 + // cudaErrorNotReady = 600, CUDA_ERROR_NOT_READY = 600, hipErrorNotReady = 1034 + // TODO [HIP]: make hipErrorNotReady = 600 + // CUDA_ERROR_NOT_READY + {"cudaErrorNotReady", {"hipErrorNotReady", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 600 + // CUDA_ERROR_ILLEGAL_ADDRESS + {"cudaErrorIllegalAddress", {"hipErrorIllegalAddress", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 700 + // cudaErrorLaunchOutOfResources = 701, CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 701, hipErrorLaunchOutOfResources = 1007 + // TODO [HIP]: make hipErrorLaunchOutOfResources = 701 + // CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES + {"cudaErrorLaunchOutOfResources", {"hipErrorLaunchOutOfResources", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 701 + // cudaErrorLaunchTimeout = 702, CUDA_ERROR_LAUNCH_TIMEOUT = 702, hipErrorLaunchTimeOut = 1006 + // TODO [HIP]: make hipErrorLaunchTimeOut = 702 + // CUDA_ERROR_LAUNCH_TIMEOUT + {"cudaErrorLaunchTimeout", {"hipErrorLaunchTimeOut", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 702 + // CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING + {"cudaErrorLaunchIncompatibleTexturing", {"hipErrorLaunchIncompatibleTexturing", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 703 + // cudaErrorPeerAccessAlreadyEnabled = 704, CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED = 704, hipErrorPeerAccessAlreadyEnabled = 1050 + // TODO [HIP]: make hipErrorPeerAccessAlreadyEnabled = 704 + // CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED + {"cudaErrorPeerAccessAlreadyEnabled", {"hipErrorPeerAccessAlreadyEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 704 + // cudaErrorPeerAccessNotEnabled = 705, CUDA_ERROR_PEER_ACCESS_NOT_ENABLED = 705, hipErrorPeerAccessNotEnabled = 1051 + // TODO [HIP]: make hipErrorPeerAccessNotEnabled = 705 + // CUDA_ERROR_PEER_ACCESS_NOT_ENABLED + {"cudaErrorPeerAccessNotEnabled", {"hipErrorPeerAccessNotEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 705 + // cudaErrorSetOnActiveProcess = 708, CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE = 708, hipErrorSetOnActiveProcess = 305 + // TODO [HIP]: make hipErrorSetOnActiveProcess = 708 + // CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE + {"cudaErrorSetOnActiveProcess", {"hipErrorSetOnActiveProcess", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 708 + // CUDA_ERROR_CONTEXT_IS_DESTROYED + {"cudaErrorContextIsDestroyed", {"hipErrorContextIsDestroyed", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 709 + // cudaErrorAssert = 710, CUDA_ERROR_ASSERT = 710, hipErrorAssert = 1081 + // TODO [HIP]: make hipErrorAssert = 710 + // CUDA_ERROR_ASSERT + {"cudaErrorAssert", {"hipErrorAssert", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 710 + // CUDA_ERROR_TOO_MANY_PEERS + {"cudaErrorTooManyPeers", {"hipErrorTooManyPeers", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 711 + // cudaErrorHostMemoryAlreadyRegistered = 712, CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED = 712, hipErrorHostMemoryAlreadyRegistered = 1061 + // TODO [HIP]: make hipErrorHostMemoryAlreadyRegistered = 712 + // CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED + {"cudaErrorHostMemoryAlreadyRegistered", {"hipErrorHostMemoryAlreadyRegistered", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 712 + // cudaErrorHostMemoryNotRegistered = 713, CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED = 713, hipErrorHostMemoryNotRegistered = 1062 + // TODO [HIP]: make hipErrorHostMemoryNotRegistered = 713 + // CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED + {"cudaErrorHostMemoryNotRegistered", {"hipErrorHostMemoryNotRegistered", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 713 + // CUDA_ERROR_HARDWARE_STACK_ERROR + {"cudaErrorHardwareStackError", {"hipErrorHardwareStackError", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 714 + // CUDA_ERROR_ILLEGAL_INSTRUCTION + {"cudaErrorIllegalInstruction", {"hipErrorIllegalInstruction", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 715 + // CUDA_ERROR_MISALIGNED_ADDRESS + {"cudaErrorMisalignedAddress", {"hipErrorMisalignedAddress", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 716 + // CUDA_ERROR_INVALID_ADDRESS_SPACE + {"cudaErrorInvalidAddressSpace", {"hipErrorInvalidAddressSpace", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 717 + // CUDA_ERROR_INVALID_PC + {"cudaErrorInvalidPc", {"hipErrorInvalidPc", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 718 + // cudaErrorLaunchFailure = 719, CUDA_ERROR_LAUNCH_FAILED = 719, hipErrorLaunchFailure = 1004 + // TODO [HIP]: make hipErrorSetOnActiveProcess = 719 + // CUDA_ERROR_LAUNCH_FAILED + {"cudaErrorLaunchFailure", {"hipErrorLaunchFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 719 + // no analogue + {"cudaErrorCooperativeLaunchTooLarge", {"hipErrorCooperativeLaunchTooLarge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 720 + // CUDA_ERROR_NOT_PERMITTED + {"cudaErrorNotPermitted", {"hipErrorNotPermitted", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 800 + // cudaErrorNotSupported = 801, CUDA_ERROR_NOT_SUPPORTED = 801, hipErrorNotSupported = 1082 + // TODO [HIP]: make hipErrorNotSupported = 801 + // CUDA_ERROR_NOT_SUPPORTED + {"cudaErrorNotSupported", {"hipErrorNotSupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 801 + // CUDA_ERROR_SYSTEM_NOT_READY + {"cudaErrorSystemNotReady", {"hipErrorSystemNotReady", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 802 // CUDA_ERROR_SYSTEM_DRIVER_MISMATCH {"cudaErrorSystemDriverMismatch", {"hipErrorSystemDriverMismatch", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 803 // CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE {"cudaErrorCompatNotSupportedOnDevice", {"hipErrorCompatNotSupportedOnDevice", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 804 - // no analogue - {"cudaErrorStartupFailure", {"hipErrorStartupFailure", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 127 - // CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED = 900 + // CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED {"cudaErrorStreamCaptureUnsupported", {"hipErrorStreamCaptureUnsupported", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 900 - // CUDA_ERROR_STREAM_CAPTURE_INVALIDATED = 901 + // CUDA_ERROR_STREAM_CAPTURE_INVALIDATED {"cudaErrorStreamCaptureInvalidated", {"hipErrorStreamCaptureInvalidated", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 901 - // CUDA_ERROR_STREAM_CAPTURE_MERGE = 902 + // CUDA_ERROR_STREAM_CAPTURE_MERGE {"cudaErrorStreamCaptureMerge", {"hipErrorStreamCaptureMerge", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 902 - // CUDA_ERROR_STREAM_CAPTURE_UNMATCHED = 903 + // CUDA_ERROR_STREAM_CAPTURE_UNMATCHED {"cudaErrorStreamCaptureUnmatched", {"hipErrorStreamCaptureUnmatched", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 903 - // CUDA_ERROR_STREAM_CAPTURE_UNJOINED = 904 + // CUDA_ERROR_STREAM_CAPTURE_UNJOINED {"cudaErrorStreamCaptureUnjoined", {"hipErrorStreamCaptureUnjoined", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 904 - // CUDA_ERROR_STREAM_CAPTURE_ISOLATION = 905 + // CUDA_ERROR_STREAM_CAPTURE_ISOLATION {"cudaErrorStreamCaptureIsolation", {"hipErrorStreamCaptureIsolation", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 905 - // CUDA_ERROR_STREAM_CAPTURE_IMPLICIT = 906 + // CUDA_ERROR_STREAM_CAPTURE_IMPLICIT {"cudaErrorStreamCaptureImplicit", {"hipErrorStreamCaptureImplicit", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 906 - // CUDA_ERROR_CAPTURED_EVENT = 907 + // CUDA_ERROR_CAPTURED_EVENT {"cudaErrorCapturedEvent", {"hipErrorCapturedEvent", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 907 - // CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD = 908 + // CUDA_ERROR_STREAM_CAPTURE_WRONG_THREAD {"cudaErrorStreamCaptureWrongThread", {"hipErrorStreamCaptureWrongThread", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 908 + // cudaErrorUnknown = 999, CUDA_ERROR_UNKNOWN = 999, hipErrorUnknown = 1030 + // TODO [HIP]: make hipErrorUnknown = 999 + // CUDA_ERROR_UNKNOWN + {"cudaErrorUnknown", {"hipErrorUnknown", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 999 // Deprecated since CUDA 4.1 {"cudaErrorApiFailureBase", {"hipErrorApiFailureBase", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 10000 From ee44fffd35af7e30b5e8b5ac9978b04ae2532ff1 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sun, 1 Dec 2019 14:47:43 +0530 Subject: [PATCH 27/27] Revert "[packaging] Conform to debian package naming rules (#1647)" (#1688) This reverts commit 8016fa32df714b478f2d9431d85282a71c61c11c. [ROCm/hip commit: 4585b7f6f3df21526b3fabe73d2b40b0d8068778] --- projects/hip/CMakeLists.txt | 31 ++++------- projects/hip/INSTALL.md | 8 +-- .../hip/docker/dockerfile-hip-ubuntu-16.04 | 8 +-- projects/hip/install.sh | 4 +- projects/hip/packaging/hip-hcc.txt | 54 ------------------- .../{hip-base.postinst => hip_base.postinst} | 0 .../{hip-base.prerm => hip_base.prerm} | 0 .../packaging/{hip-base.txt => hip_base.txt} | 3 +- .../packaging/{hip-doc.txt => hip_doc.txt} | 7 ++- .../{hip-hcc.postinst => hip_hcc.postinst} | 0 .../{hip-hcc.prerm => hip_hcc.prerm} | 0 projects/hip/packaging/hip_hcc.txt | 31 +++++++++-- .../packaging/{hip-nvcc.txt => hip_nvcc.txt} | 7 ++- .../{hip-samples.txt => hip_samples.txt} | 7 ++- 14 files changed, 58 insertions(+), 102 deletions(-) delete mode 100644 projects/hip/packaging/hip-hcc.txt rename projects/hip/packaging/{hip-base.postinst => hip_base.postinst} (100%) rename projects/hip/packaging/{hip-base.prerm => hip_base.prerm} (100%) rename projects/hip/packaging/{hip-base.txt => hip_base.txt} (96%) rename projects/hip/packaging/{hip-doc.txt => hip_doc.txt} (89%) rename projects/hip/packaging/{hip-hcc.postinst => hip_hcc.postinst} (100%) rename projects/hip/packaging/{hip-hcc.prerm => hip_hcc.prerm} (100%) rename projects/hip/packaging/{hip-nvcc.txt => hip_nvcc.txt} (87%) rename projects/hip/packaging/{hip-samples.txt => hip_samples.txt} (85%) diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index 5f388f080b..b26ee22d62 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -421,9 +421,9 @@ endif() ############################# # Package: hip_base set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_base) -configure_file(packaging/hip-base.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) -configure_file(packaging/hip-base.postinst ${BUILD_DIR}/postinst @ONLY) -configure_file(packaging/hip-base.prerm ${BUILD_DIR}/prerm @ONLY) +configure_file(packaging/hip_base.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip_base.postinst ${BUILD_DIR}/postinst @ONLY) +configure_file(packaging/hip_base.prerm ${BUILD_DIR}/prerm @ONLY) add_custom_target(pkg_hip_base COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz @@ -441,9 +441,9 @@ endif() # Package: hip_hcc set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_hcc) -configure_file(packaging/hip-hcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) -configure_file(packaging/hip-hcc.postinst ${BUILD_DIR}/postinst @ONLY) -configure_file(packaging/hip-hcc.prerm ${BUILD_DIR}/prerm @ONLY) +configure_file(packaging/hip_hcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip_hcc.postinst ${BUILD_DIR}/postinst @ONLY) +configure_file(packaging/hip_hcc.prerm ${BUILD_DIR}/prerm @ONLY) add_custom_target(pkg_hip_hcc COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -455,7 +455,7 @@ add_custom_target(pkg_hip_hcc COMMAND ${CMAKE_COMMAND} . # Package: hip_nvcc set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_nvcc) -configure_file(packaging/hip-nvcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip_nvcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) add_custom_target(pkg_hip_nvcc COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -466,7 +466,7 @@ add_custom_target(pkg_hip_nvcc COMMAND ${CMAKE_COMMAND} . # Package: hip_doc set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_doc) -configure_file(packaging/hip-doc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip_doc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) add_custom_target(pkg_hip_doc COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -477,7 +477,7 @@ add_custom_target(pkg_hip_doc COMMAND ${CMAKE_COMMAND} . # Package: hip_samples set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_samples) -configure_file(packaging/hip-samples.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +configure_file(packaging/hip_samples.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) add_custom_target(pkg_hip_samples COMMAND ${CMAKE_COMMAND} . COMMAND rm -rf *.deb *.rpm *.tar.gz COMMAND make package @@ -486,17 +486,6 @@ add_custom_target(pkg_hip_samples COMMAND ${CMAKE_COMMAND} . COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR} WORKING_DIRECTORY ${BUILD_DIR}) -# Package: hip_hcc_legacy -set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_hcc_legacy) -configure_file(packaging/hip_hcc.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) -add_custom_target(pkg_hip_hcc_legacy COMMAND ${CMAKE_COMMAND} . - COMMAND rm -rf *.deb *.rpm *.tar.gz - COMMAND make package - COMMAND cp *.deb ${PROJECT_BINARY_DIR} - COMMAND cp *.rpm ${PROJECT_BINARY_DIR} - COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR} - WORKING_DIRECTORY ${BUILD_DIR}) - # Package: all if(POLICY CMP0037) cmake_policy(PUSH) @@ -509,7 +498,7 @@ for i in *.rpm ; do mv \$i \${i/.rpm/.x86_64.rpm} ; done add_custom_target(package COMMAND bash ${PROJECT_BINARY_DIR}/fixnames WORKING_DIRECTORY ${PROJECT_BINARY_DIR} - DEPENDS pkg_hip_base pkg_hip_hcc pkg_hip_nvcc pkg_hip_doc pkg_hip_samples pkg_hip_hcc_legacy) + DEPENDS pkg_hip_base pkg_hip_hcc pkg_hip_nvcc pkg_hip_doc pkg_hip_samples) if(POLICY CMP0037) cmake_policy(POP) endif() diff --git a/projects/hip/INSTALL.md b/projects/hip/INSTALL.md index 2dc12f81a8..273c77a23c 100644 --- a/projects/hip/INSTALL.md +++ b/projects/hip/INSTALL.md @@ -24,9 +24,9 @@ HIP code can be developed either on AMD ROCm platform using hcc or clang compile ## HIP-hcc * Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories). -* Install the "hip-hcc" package. This will install HCC and the HIP porting layer. +* Install the "hip_hcc" package. This will install HCC and the HIP porting layer. ``` -apt-get install hip-hcc +apt-get install hip_hcc ``` * Default paths and environment variables: @@ -67,9 +67,9 @@ apt-get install hip-hcc ## HIP-nvcc * Add the ROCm package server to your system as per the OS-specific guide available [here](https://rocm.github.io/ROCmInstall.html#installing-from-amd-rocm-repositories). -* Install the "hip-nvcc" package. This will install CUDA SDK and the HIP porting layer. +* Install the "hip_nvcc" package. This will install CUDA SDK and the HIP porting layer. ``` -apt-get install hip-nvcc +apt-get install hip_nvcc ``` * Default paths and environment variables: diff --git a/projects/hip/docker/dockerfile-hip-ubuntu-16.04 b/projects/hip/docker/dockerfile-hip-ubuntu-16.04 index 93dd17b948..0852ae38c5 100644 --- a/projects/hip/docker/dockerfile-hip-ubuntu-16.04 +++ b/projects/hip/docker/dockerfile-hip-ubuntu-16.04 @@ -10,10 +10,10 @@ COPY *.deb /tmp/ # Install the debian package RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install --no-install-recommends -y curl \ && apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install --no-install-recommends --allow-unauthenticated -y \ - /tmp/hip-base-*.deb \ - /tmp/hip-hcc-*.deb \ - /tmp/hip-doc-*.deb \ - /tmp/hip-samples-* \ + /tmp/hip_base-*.deb \ + /tmp/hip_hcc-*.deb \ + /tmp/hip_doc-*.deb \ + /tmp/hip_samples-* \ && rm -f /tmp/*.deb \ && apt-get clean \ && rm -rf /var/lib/apt/lists/* diff --git a/projects/hip/install.sh b/projects/hip/install.sh index 51a303e812..c7dc48e1a8 100755 --- a/projects/hip/install.sh +++ b/projects/hip/install.sh @@ -35,8 +35,8 @@ function buildHIP() 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 hip_base*.deb hip_hcc*.deb hip_sample*.deb hip_doc*.deb popd rm -rf $BUILD_ROOT } diff --git a/projects/hip/packaging/hip-hcc.txt b/projects/hip/packaging/hip-hcc.txt deleted file mode 100644 index 0ccf75af50..0000000000 --- a/projects/hip/packaging/hip-hcc.txt +++ /dev/null @@ -1,54 +0,0 @@ -cmake_minimum_required(VERSION 2.8.3) -project(hip_hcc) - -install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib) -install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib) -if(NOT @HIP_COMPILER@ STREQUAL "clang") - install(FILES @PROJECT_BINARY_DIR@/libhiprtc.so DESTINATION lib) -endif() -install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib) -install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip) -install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/packaging/hip-targets-release.cmake DESTINATION lib/cmake/hip) - -############################# -# Packaging steps -############################# -set(CPACK_SET_DESTDIR TRUE) -set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip-hcc") -if(@HCC_VERSION_MAJOR@ EQUAL 0) - set(HCC_PACKAGE_NAME "hcc_lc") -else() - set(HCC_PACKAGE_NAME "hcc") -endif() -set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [HCC]") -set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") -set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") -set(CPACK_PACKAGE_VERSION @HIP_VERSION_MAJOR@.@HIP_VERSION_MINOR@.@HIP_VERSION_PATCH@) -set(CPACK_PACKAGE_VERSION_MAJOR @HIP_VERSION_MAJOR@) -set(CPACK_PACKAGE_VERSION_MINOR @HIP_VERSION_MINOR@) -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_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)") -else() - set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), comgr (>= 1.1)") -endif() -set(CPACK_DEBIAN_PACKAGE_REPLACES "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") -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") -else() - set(CPACK_RPM_PACKAGE_REQUIRES "hip-base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, comgr >= 1.1") -endif() -set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") -set(CPACK_SOURCE_GENERATOR "TGZ") -include(CPack) diff --git a/projects/hip/packaging/hip-base.postinst b/projects/hip/packaging/hip_base.postinst similarity index 100% rename from projects/hip/packaging/hip-base.postinst rename to projects/hip/packaging/hip_base.postinst diff --git a/projects/hip/packaging/hip-base.prerm b/projects/hip/packaging/hip_base.prerm similarity index 100% rename from projects/hip/packaging/hip-base.prerm rename to projects/hip/packaging/hip_base.prerm diff --git a/projects/hip/packaging/hip-base.txt b/projects/hip/packaging/hip_base.txt similarity index 96% rename from projects/hip/packaging/hip-base.txt rename to projects/hip/packaging/hip_base.txt index 381c8b66d1..a44f0acb74 100644 --- a/projects/hip/packaging/hip-base.txt +++ b/projects/hip/packaging/hip_base.txt @@ -16,7 +16,7 @@ install(DIRECTORY @hip_SOURCE_DIR@/cmake DESTINATION .) ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip-base") +set(CPACK_PACKAGE_NAME "hip_base") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [BASE]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -29,7 +29,6 @@ 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_REPLACES "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") diff --git a/projects/hip/packaging/hip-doc.txt b/projects/hip/packaging/hip_doc.txt similarity index 89% rename from projects/hip/packaging/hip-doc.txt rename to projects/hip/packaging/hip_doc.txt index 5f07785ba5..2c1619f221 100644 --- a/projects/hip/packaging/hip-doc.txt +++ b/projects/hip/packaging/hip_doc.txt @@ -20,7 +20,7 @@ endif() ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip-doc") +set(CPACK_PACKAGE_NAME "hip_doc") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [DOCUMENTATION]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -31,13 +31,12 @@ 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_REPLACES "hip_doc") +set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION})") 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_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/packaging/hip-hcc.postinst b/projects/hip/packaging/hip_hcc.postinst similarity index 100% rename from projects/hip/packaging/hip-hcc.postinst rename to projects/hip/packaging/hip_hcc.postinst diff --git a/projects/hip/packaging/hip-hcc.prerm b/projects/hip/packaging/hip_hcc.prerm similarity index 100% rename from projects/hip/packaging/hip-hcc.prerm rename to projects/hip/packaging/hip_hcc.prerm diff --git a/projects/hip/packaging/hip_hcc.txt b/projects/hip/packaging/hip_hcc.txt index 083288d5bd..c688433cc0 100644 --- a/projects/hip/packaging/hip_hcc.txt +++ b/projects/hip/packaging/hip_hcc.txt @@ -1,5 +1,14 @@ cmake_minimum_required(VERSION 2.8.3) -project(hip_hcc_legacy) +project(hip_hcc) + +install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib) +install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib) +if(NOT @HIP_COMPILER@ STREQUAL "clang") + install(FILES @PROJECT_BINARY_DIR@/libhiprtc.so DESTINATION lib) +endif() +install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib) +install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip) +install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/packaging/hip-targets-release.cmake DESTINATION lib/cmake/hip) ############################# # Packaging steps @@ -7,6 +16,11 @@ project(hip_hcc_legacy) set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") set(CPACK_PACKAGE_NAME "hip_hcc") +if(@HCC_VERSION_MAJOR@ EQUAL 0) + set(HCC_PACKAGE_NAME "hcc_lc") +else() + set(HCC_PACKAGE_NAME "hcc") +endif() set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [HCC]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -17,12 +31,23 @@ 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-hcc (= ${CPACK_PACKAGE_VERSION})") +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)") +else() + set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), ${HCC_PACKAGE_NAME} (= @HCC_PACKAGE_VERSION@), comgr (>= 1.1)") +endif() 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-hcc = ${HIP_BASE_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") +else() + set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${HIP_BASE_VERSION}, ${HCC_PACKAGE_NAME} = @HCC_PACKAGE_VERSION@, comgr >= 1.1") +endif() set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/packaging/hip-nvcc.txt b/projects/hip/packaging/hip_nvcc.txt similarity index 87% rename from projects/hip/packaging/hip-nvcc.txt rename to projects/hip/packaging/hip_nvcc.txt index 11c1c290a7..975176f934 100644 --- a/projects/hip/packaging/hip-nvcc.txt +++ b/projects/hip/packaging/hip_nvcc.txt @@ -6,7 +6,7 @@ project(hip_nvcc) ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip-nvcc") +set(CPACK_PACKAGE_NAME "hip_nvcc") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [NVCC]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -18,15 +18,14 @@ 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_REPLACES "hip_nvcc") +set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), cuda (>= 7.5)") 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_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack) diff --git a/projects/hip/packaging/hip-samples.txt b/projects/hip/packaging/hip_samples.txt similarity index 85% rename from projects/hip/packaging/hip-samples.txt rename to projects/hip/packaging/hip_samples.txt index 8a8bd05ceb..4c1e04cb23 100644 --- a/projects/hip/packaging/hip-samples.txt +++ b/projects/hip/packaging/hip_samples.txt @@ -8,7 +8,7 @@ install(DIRECTORY @hip_SOURCE_DIR@/samples DESTINATION .) ############################# set(CPACK_SET_DESTDIR TRUE) set(CPACK_INSTALL_PREFIX "/opt/rocm/hip") -set(CPACK_PACKAGE_NAME "hip-samples") +set(CPACK_PACKAGE_NAME "hip_samples") set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [SAMPLES]") set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") set(CPACK_PACKAGE_CONTACT "Maneesh Gupta ") @@ -19,13 +19,12 @@ 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_REPLACES "hip_samples") +set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION})") 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_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") set(CPACK_SOURCE_GENERATOR "TGZ") include(CPack)