diff options
33 files changed, 349 insertions, 48 deletions
diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml index 39a73a3..a214b06 100644 --- a/.gitlab-ci.yml +++ b/.gitlab-ci.yml @@ -334,6 +334,16 @@ t:cuda11.8-minimal-ninja: variables: CMAKE_CI_NO_MR: "true" +t:hip5.5-nvidia: + extends: + - .hip5.5_nvidia + - .cmake_test_linux_release + - .linux_x86_64_tags_cuda + - .run_dependent + - .needs_centos7_x86_64 + variables: + CMAKE_CI_NO_MR: "true" + t:hip5.5-radeon: extends: - .hip5.5_radeon diff --git a/.gitlab/ci/configure_debian12_hip_radeon.cmake b/.gitlab/ci/configure_debian12_hip_radeon.cmake index 58036b0..c7d7004 100644 --- a/.gitlab/ci/configure_debian12_hip_radeon.cmake +++ b/.gitlab/ci/configure_debian12_hip_radeon.cmake @@ -1,3 +1,3 @@ -set(CMake_TEST_HIP "ON" CACHE BOOL "") +set(CMake_TEST_HIP "amd" CACHE BOOL "") include("${CMAKE_CURRENT_LIST_DIR}/configure_external_test.cmake") diff --git a/.gitlab/ci/configure_fedora38_hip_radeon.cmake b/.gitlab/ci/configure_fedora38_hip_radeon.cmake index 58036b0..c7d7004 100644 --- a/.gitlab/ci/configure_fedora38_hip_radeon.cmake +++ b/.gitlab/ci/configure_fedora38_hip_radeon.cmake @@ -1,3 +1,3 @@ -set(CMake_TEST_HIP "ON" CACHE BOOL "") +set(CMake_TEST_HIP "amd" CACHE BOOL "") include("${CMAKE_CURRENT_LIST_DIR}/configure_external_test.cmake") diff --git a/.gitlab/ci/configure_hip5.5_nvidia.cmake b/.gitlab/ci/configure_hip5.5_nvidia.cmake new file mode 100644 index 0000000..4b3511a --- /dev/null +++ b/.gitlab/ci/configure_hip5.5_nvidia.cmake @@ -0,0 +1,3 @@ +set(CMake_TEST_HIP "nvidia" CACHE BOOL "") + +include("${CMAKE_CURRENT_LIST_DIR}/configure_external_test.cmake") diff --git a/.gitlab/ci/configure_hip5.5_radeon.cmake b/.gitlab/ci/configure_hip5.5_radeon.cmake index 58036b0..c7d7004 100644 --- a/.gitlab/ci/configure_hip5.5_radeon.cmake +++ b/.gitlab/ci/configure_hip5.5_radeon.cmake @@ -1,3 +1,3 @@ -set(CMake_TEST_HIP "ON" CACHE BOOL "") +set(CMake_TEST_HIP "amd" CACHE BOOL "") include("${CMAKE_CURRENT_LIST_DIR}/configure_external_test.cmake") diff --git a/.gitlab/ci/docker/hip5.5/Dockerfile b/.gitlab/ci/docker/hip5.5/Dockerfile index 2deb3c6..3a4aa53 100644 --- a/.gitlab/ci/docker/hip5.5/Dockerfile +++ b/.gitlab/ci/docker/hip5.5/Dockerfile @@ -2,7 +2,13 @@ ARG BASE_IMAGE=rocm/dev-ubuntu-22.04:5.5 -FROM ${BASE_IMAGE} AS apt-cache +FROM ${BASE_IMAGE} AS cuda-keyring +ADD https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb /root/ +RUN --mount=type=tmpfs,target=/var/log \ + dpkg -i /root/cuda-keyring_1.1-1_all.deb \ + && rm /root/cuda-keyring_1.1-1_all.deb + +FROM cuda-keyring AS apt-cache # Populate APT cache w/ the fresh metadata and prefetch packages. # Use an empty `docker-clean` file to "hide" the image-provided # file to disallow removing packages after `apt-get` operations. @@ -12,9 +18,12 @@ RUN --mount=type=tmpfs,target=/var/log \ apt-get update \ && apt-get --download-only -y install $(grep -h '^[^#]\+$' /root/*.lst) -FROM ${BASE_IMAGE} +FROM cuda-keyring MAINTAINER Brad King <brad.king@kitware.com> +ENV NVIDIA_DRIVER_CAPABILITIES=compute,utility +ENV NVIDIA_REQUIRE_CUDA=cuda>=11.8 +ENV NVIDIA_VISIBLE_DEVICES=all ENV PATH="/opt/rocm/bin:$PATH" RUN --mount=type=bind,source=install_deps.sh,target=/root/install_deps.sh \ diff --git a/.gitlab/ci/docker/hip5.5/deps_packages.lst b/.gitlab/ci/docker/hip5.5/deps_packages.lst index 9847925..3276055 100644 --- a/.gitlab/ci/docker/hip5.5/deps_packages.lst +++ b/.gitlab/ci/docker/hip5.5/deps_packages.lst @@ -2,3 +2,20 @@ g++ curl git + +# NVIDIA CUDA Compiler +cuda-keyring +cuda-nvcc-11-8 +cuda-profiler-api-11-8 + +# NVIDIA CUDA Toolkit +# These are not needed for HIP, but having them in +# the environment allows us to run CUDA tests too. +cuda-nvrtc-dev-11-8 +cuda-nvtx-11-8 +libcublas-dev-11-8 +libcufft-dev-11-8 +libcurand-dev-11-8 +libcusolver-dev-11-8 +libcusparse-dev-11-8 +libnpp-dev-11-8 diff --git a/.gitlab/ci/env_hip5.5_nvidia.sh b/.gitlab/ci/env_hip5.5_nvidia.sh new file mode 100644 index 0000000..67d1ef2 --- /dev/null +++ b/.gitlab/ci/env_hip5.5_nvidia.sh @@ -0,0 +1,4 @@ +export HIP_PLATFORM=nvidia +export CUDA_PATH=/usr/local/cuda-11.8 +export PATH=/usr/local/cuda-11.8/bin:$PATH +export LD_LIBRARY_PATH=/usr/local/cuda-11.8/lib64 diff --git a/.gitlab/os-linux.yml b/.gitlab/os-linux.yml index 8398108..d42340c 100644 --- a/.gitlab/os-linux.yml +++ b/.gitlab/os-linux.yml @@ -345,7 +345,7 @@ ### HIP builds .hip5.5: - image: "kitware/cmake:ci-hip5.5-x86_64-2023-06-01" + image: "kitware/cmake:ci-hip5.5-x86_64-2023-09-18" variables: GIT_CLONE_PATH: "$CI_BUILDS_DIR/cmake ci" @@ -373,6 +373,13 @@ CMAKE_CONFIGURATION: fedora38_hip_radeon CTEST_LABELS: "HIP" +.hip5.5_nvidia: + extends: .hip5.5 + + variables: + CMAKE_CONFIGURATION: hip5.5_nvidia + CTEST_LABELS: "HIP" + ### C++ modules .gcc_cxx_modules_x86_64: diff --git a/Help/manual/cmake-variables.7.rst b/Help/manual/cmake-variables.7.rst index fa7a90f..536046a 100644 --- a/Help/manual/cmake-variables.7.rst +++ b/Help/manual/cmake-variables.7.rst @@ -590,6 +590,7 @@ Variables for Languages /variable/CMAKE_Fortran_MODOUT_FLAG /variable/CMAKE_HIP_ARCHITECTURES /variable/CMAKE_HIP_EXTENSIONS + /variable/CMAKE_HIP_PLATFORM /variable/CMAKE_HIP_STANDARD /variable/CMAKE_HIP_STANDARD_REQUIRED /variable/CMAKE_ISPC_HEADER_DIRECTORY diff --git a/Help/prop_tgt/HIP_ARCHITECTURES.rst b/Help/prop_tgt/HIP_ARCHITECTURES.rst index 06f956b..58a813d 100644 --- a/Help/prop_tgt/HIP_ARCHITECTURES.rst +++ b/Help/prop_tgt/HIP_ARCHITECTURES.rst @@ -3,7 +3,8 @@ HIP_ARCHITECTURES .. versionadded:: 3.21 -List of AMD GPU architectures to generate device code for. +List of GPU architectures to for which to generate device code. +Architecture names are interpreted based on :variable:`CMAKE_HIP_PLATFORM`. A non-empty false value (e.g. ``OFF``) disables adding architectures. This is intended to support packagers and rare cases where full control diff --git a/Help/release/dev/hip-nvidia.rst b/Help/release/dev/hip-nvidia.rst new file mode 100644 index 0000000..1d9814e --- /dev/null +++ b/Help/release/dev/hip-nvidia.rst @@ -0,0 +1,9 @@ +hip-nvidia +---------- + +* ``HIP`` language code may now be compiled for NVIDIA GPUs + using the NVIDIA CUDA Compiler (NVCC). + +* The :variable:`CMAKE_HIP_PLATFORM` variable was added to specify + the GPU platform for which HIP language sources are to be compiled + (``amd`` or ``nvidia``). diff --git a/Help/variable/CMAKE_HIP_ARCHITECTURES.rst b/Help/variable/CMAKE_HIP_ARCHITECTURES.rst index bcc6b35..3f17983 100644 --- a/Help/variable/CMAKE_HIP_ARCHITECTURES.rst +++ b/Help/variable/CMAKE_HIP_ARCHITECTURES.rst @@ -3,10 +3,14 @@ CMAKE_HIP_ARCHITECTURES .. versionadded:: 3.21 -Default value for :prop_tgt:`HIP_ARCHITECTURES` property of targets. +List of GPU architectures to for which to generate device code. +Architecture names are interpreted based on :variable:`CMAKE_HIP_PLATFORM`. -This is initialized to the architectures reported by ``rocm_agent_enumerator``, -if available, and otherwise to the default chosen by the compiler. +This is initialized based on the value of :variable:`CMAKE_HIP_PLATFORM`: + +``amd`` + Uses architectures reported by ``rocm_agent_enumerator``, if available, + and otherwise to a default chosen by the compiler. This variable is used to initialize the :prop_tgt:`HIP_ARCHITECTURES` property on all targets. See the target property for additional information. diff --git a/Help/variable/CMAKE_HIP_PLATFORM.rst b/Help/variable/CMAKE_HIP_PLATFORM.rst new file mode 100644 index 0000000..5e3a2b7 --- /dev/null +++ b/Help/variable/CMAKE_HIP_PLATFORM.rst @@ -0,0 +1,22 @@ +CMAKE_HIP_PLATFORM +------------------ + +.. versionadded:: 3.28 + +GPU platform for which HIP language sources are to be compiled. + +The value must be one of: + +``amd`` + AMD GPUs + +``nvidia`` + NVIDIA GPUs + +If not specified, a default is computed via ``hipconfig --platform``. + +:variable:`CMAKE_HIP_ARCHITECTURES` entries are interpreted with +as architectures of the GPU platform. + +:variable:`CMAKE_HIP_COMPILER <CMAKE_<LANG>_COMPILER>` must target +the same GPU platform. diff --git a/Modules/CMakeCompilerIdDetection.cmake b/Modules/CMakeCompilerIdDetection.cmake index 5fec06f..eef2bd5 100644 --- a/Modules/CMakeCompilerIdDetection.cmake +++ b/Modules/CMakeCompilerIdDetection.cmake @@ -103,7 +103,7 @@ function(compiler_id_detection outvar lang) endif() if("x${lang}" STREQUAL "xHIP") - set(ordered_compilers Clang) + set(ordered_compilers NVIDIA Clang) endif() if(CID_ID_DEFINE) diff --git a/Modules/CMakeDetermineHIPCompiler.cmake b/Modules/CMakeDetermineHIPCompiler.cmake index 9a40e82..21322f1 100644 --- a/Modules/CMakeDetermineHIPCompiler.cmake +++ b/Modules/CMakeDetermineHIPCompiler.cmake @@ -10,6 +10,24 @@ if( NOT ( ("${CMAKE_GENERATOR}" MATCHES "Make") OR message(FATAL_ERROR "HIP language not currently supported by \"${CMAKE_GENERATOR}\" generator") endif() +if(NOT CMAKE_HIP_PLATFORM) + execute_process(COMMAND hipconfig --platform + OUTPUT_VARIABLE _CMAKE_HIPCONFIG_PLATFORM OUTPUT_STRIP_TRAILING_WHITESPACE + RESULT_VARIABLE _CMAKE_HIPCONFIG_RESULT + ) + if(_CMAKE_HIPCONFIG_RESULT EQUAL 0 AND _CMAKE_HIPCONFIG_PLATFORM MATCHES "^(nvidia|nvcc)$") + set(CMAKE_HIP_PLATFORM "nvidia" CACHE STRING "HIP platform" FORCE) + else() + set(CMAKE_HIP_PLATFORM "amd" CACHE STRING "HIP platform" FORCE) + endif() +endif() +if(NOT CMAKE_HIP_PLATFORM MATCHES "^(amd|nvidia)$") + message(FATAL_ERROR + "The CMAKE_HIP_PLATFORM has unsupported value:\n" + " '${CMAKE_HIP_PLATFORM}'\n" + "It must be 'amd' or 'nvidia'." + ) +endif() if(NOT CMAKE_HIP_COMPILER) set(CMAKE_HIP_COMPILER_INIT NOTFOUND) @@ -34,15 +52,19 @@ if(NOT CMAKE_HIP_COMPILER) # finally list compilers to try if(NOT CMAKE_HIP_COMPILER_INIT) - set(CMAKE_HIP_COMPILER_LIST clang++) + if(CMAKE_HIP_PLATFORM STREQUAL "nvidia") + set(CMAKE_HIP_COMPILER_LIST nvcc) + elseif(CMAKE_HIP_PLATFORM STREQUAL "amd") + set(CMAKE_HIP_COMPILER_LIST clang++) - # Look for the Clang coming with ROCm to support HIP. - execute_process(COMMAND hipconfig --hipclangpath - OUTPUT_VARIABLE _CMAKE_HIPCONFIG_CLANGPATH - RESULT_VARIABLE _CMAKE_HIPCONFIG_RESULT - ) - if(_CMAKE_HIPCONFIG_RESULT EQUAL 0 AND EXISTS "${_CMAKE_HIPCONFIG_CLANGPATH}") - set(CMAKE_HIP_COMPILER_HINTS "${_CMAKE_HIPCONFIG_CLANGPATH}") + # Look for the Clang coming with ROCm to support HIP. + execute_process(COMMAND hipconfig --hipclangpath + OUTPUT_VARIABLE _CMAKE_HIPCONFIG_CLANGPATH + RESULT_VARIABLE _CMAKE_HIPCONFIG_RESULT + ) + if(_CMAKE_HIPCONFIG_RESULT EQUAL 0 AND EXISTS "${_CMAKE_HIPCONFIG_CLANGPATH}") + set(CMAKE_HIP_COMPILER_HINTS "${_CMAKE_HIPCONFIG_CLANGPATH}") + endif() endif() endif() @@ -63,17 +85,55 @@ mark_as_advanced(CMAKE_HIP_COMPILER) if(NOT CMAKE_HIP_COMPILER_ID_RUN) set(CMAKE_HIP_COMPILER_ID_RUN 1) - # Try to identify the compiler. + include(${CMAKE_ROOT}/Modules/CMakeDetermineCompilerId.cmake) + + # We determine the vendor to use the right flags for detection right away. + # The main compiler identification is still needed below to extract other information. + list(APPEND CMAKE_HIP_COMPILER_ID_VENDORS NVIDIA Clang) + set(CMAKE_HIP_COMPILER_ID_VENDOR_REGEX_NVIDIA "nvcc: NVIDIA \\(R\\) Cuda compiler driver") + set(CMAKE_HIP_COMPILER_ID_VENDOR_REGEX_Clang "(clang version)") + CMAKE_DETERMINE_COMPILER_ID_VENDOR(HIP "--version") + + if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA") + # Find the CUDA toolkit to get: + # - CMAKE_HIP_COMPILER_CUDA_TOOLKIT_VERSION + # - CMAKE_HIP_COMPILER_CUDA_TOOLKIT_ROOT + # - CMAKE_HIP_COMPILER_CUDA_LIBRARY_ROOT + # We save them in CMakeHIPCompiler.cmake. + # Match arguments with cmake_cuda_architectures_all call. + include(Internal/CMakeCUDAFindToolkit) + cmake_cuda_find_toolkit(HIP CMAKE_HIP_COMPILER_CUDA_) + + # If the user set CMAKE_HIP_ARCHITECTURES, validate its value. + include(Internal/CMakeCUDAArchitecturesValidate) + cmake_cuda_architectures_validate(HIP) + endif() + + if(CMAKE_HIP_COMPILER_ID STREQUAL "Clang") + list(APPEND CMAKE_HIP_COMPILER_ID_TEST_FLAGS_FIRST "-v") + elseif(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA") + # Tell nvcc to treat .hip files as CUDA sources. + list(APPEND CMAKE_HIP_COMPILER_ID_TEST_FLAGS_FIRST "-x cu -v") + endif() + + # We perform compiler identification for a second time to extract implicit linking info. + # We need to unset the compiler ID otherwise CMAKE_DETERMINE_COMPILER_ID() doesn't work. set(CMAKE_HIP_COMPILER_ID) set(CMAKE_HIP_PLATFORM_ID) file(READ ${CMAKE_ROOT}/Modules/CMakePlatformId.h.in CMAKE_HIP_COMPILER_ID_PLATFORM_CONTENT) - list(APPEND CMAKE_HIP_COMPILER_ID_TEST_FLAGS_FIRST "-v") - - include(${CMAKE_ROOT}/Modules/CMakeDetermineCompilerId.cmake) CMAKE_DETERMINE_COMPILER_ID(HIP HIPFLAGS CMakeHIPCompilerId.hip) + if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA") + include(Internal/CMakeCUDAArchitecturesAll) + # From CMAKE_HIP_COMPILER_CUDA_TOOLKIT_VERSION and CMAKE_HIP_COMPILER_{ID,VERSION}, get: + # - CMAKE_HIP_ARCHITECTURES_ALL + # - CMAKE_HIP_ARCHITECTURES_ALL_MAJOR + # Match arguments with cmake_cuda_find_toolkit call. + cmake_cuda_architectures_all(HIP CMAKE_HIP_COMPILER_CUDA_) + endif() + _cmake_find_compiler_sysroot(HIP) endif() @@ -104,7 +164,8 @@ if(NOT CMAKE_HIP_COMPILER_ROCM_ROOT) message(FATAL_ERROR "Failed to find ROCm root directory.") endif() -# Normally implicit link information is not detected until +# Normally implicit link information is not detected until ABI detection, +# but we need to populate CMAKE_HIP_LIBRARY_ARCHITECTURE to find hip-lang. cmake_parse_implicit_link_info("${CMAKE_HIP_COMPILER_PRODUCED_OUTPUT}" _CMAKE_HIP_COMPILER_ID_IMPLICIT_LIBS _CMAKE_HIP_COMPILER_ID_IMPLICIT_DIRS @@ -165,6 +226,26 @@ include(CMakeFindBinUtils) include(Compiler/${CMAKE_HIP_COMPILER_ID}-FindBinUtils OPTIONAL) unset(_CMAKE_PROCESSING_LANGUAGE) +if(CMAKE_HIP_COMPILER_ID STREQUAL "Clang") + set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED") +elseif(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA") + include(Internal/CMakeNVCCParseImplicitInfo) + # Parse CMAKE_HIP_COMPILER_PRODUCED_OUTPUT to get: + # - CMAKE_HIP_ARCHITECTURES_DEFAULT + # - CMAKE_HIP_HOST_IMPLICIT_LINK_DIRECTORIES + # - CMAKE_HIP_HOST_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES + # - CMAKE_HIP_HOST_IMPLICIT_LINK_LIBRARIES + # - CMAKE_HIP_HOST_LINK_LAUNCHER + # - CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT + # - CMAKE_HIP_CUDA_TOOLKIT_INCLUDE_DIRECTORIES + # Match arguments with cmake_nvcc_filter_implicit_info call in CMakeTestHIPCompiler. + cmake_nvcc_parse_implicit_info(HIP CMAKE_HIP_CUDA_) + + include(Internal/CMakeCUDAFilterImplicitLibs) + # Filter out implicit link libraries that should not be passed unconditionally. + cmake_cuda_filter_implicit_libs(CMAKE_HIP_HOST_IMPLICIT_LINK_LIBRARIES) +endif() + if(CMAKE_HIP_COMPILER_SYSROOT) string(CONCAT _SET_CMAKE_HIP_COMPILER_SYSROOT "set(CMAKE_HIP_COMPILER_SYSROOT \"${CMAKE_HIP_COMPILER_SYSROOT}\")\n" @@ -185,7 +266,20 @@ if(MSVC_HIP_ARCHITECTURE_ID) "set(MSVC_HIP_ARCHITECTURE_ID ${MSVC_HIP_ARCHITECTURE_ID})") endif() -if(NOT DEFINED CMAKE_HIP_ARCHITECTURES) +if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA") + if(NOT "$ENV{CUDAARCHS}" STREQUAL "") + set(CMAKE_HIP_ARCHITECTURES "$ENV{CUDAARCHS}" CACHE STRING "CUDA architectures") + endif() + + # If the user did not set CMAKE_HIP_ARCHITECTURES, use the compiler's default. + if("${CMAKE_HIP_ARCHITECTURES}" STREQUAL "") + set(CMAKE_HIP_ARCHITECTURES "${CMAKE_HIP_ARCHITECTURES_DEFAULT}" CACHE STRING "HIP architectures" FORCE) + if(NOT CMAKE_HIP_ARCHITECTURES) + message(FATAL_ERROR "Failed to detect a default HIP architecture.\n\nCompiler output:\n${CMAKE_HIP_COMPILER_PRODUCED_OUTPUT}") + endif() + endif() + unset(CMAKE_HIP_ARCHITECTURES_DEFAULT) +elseif(NOT DEFINED CMAKE_HIP_ARCHITECTURES) # Use 'rocm_agent_enumerator' to get the current GPU architecture. set(_CMAKE_HIP_ARCHITECTURES) find_program(_CMAKE_HIP_ROCM_AGENT_ENUMERATOR diff --git a/Modules/CMakeHIPCompiler.cmake.in b/Modules/CMakeHIPCompiler.cmake.in index 0fa5bf0..6d5e62a 100644 --- a/Modules/CMakeHIPCompiler.cmake.in +++ b/Modules/CMakeHIPCompiler.cmake.in @@ -1,4 +1,6 @@ set(CMAKE_HIP_COMPILER "@CMAKE_HIP_COMPILER@") +set(CMAKE_HIP_HOST_COMPILER "@CMAKE_HIP_HOST_COMPILER@") +set(CMAKE_HIP_HOST_LINK_LAUNCHER "@CMAKE_HIP_HOST_LINK_LAUNCHER@") set(CMAKE_HIP_COMPILER_ID "@CMAKE_HIP_COMPILER_ID@") set(CMAKE_HIP_COMPILER_VERSION "@CMAKE_HIP_COMPILER_VERSION@") set(CMAKE_HIP_STANDARD_COMPUTED_DEFAULT "@CMAKE_HIP_STANDARD_COMPUTED_DEFAULT@") @@ -45,14 +47,27 @@ if(CMAKE_HIP_LIBRARY_ARCHITECTURE) set(CMAKE_LIBRARY_ARCHITECTURE "@CMAKE_HIP_LIBRARY_ARCHITECTURE@") endif() -set(CMAKE_HIP_TOOLKIT_INCLUDE_DIRECTORIES "@CMAKE_HIP_TOOLKIT_INCLUDE_DIRECTORIES@") +set(CMAKE_HIP_COMPILER_CUDA_TOOLKIT_ROOT "@CMAKE_HIP_COMPILER_CUDA_TOOLKIT_ROOT@") +set(CMAKE_HIP_COMPILER_CUDA_TOOLKIT_LIBRARY_ROOT "@CMAKE_HIP_COMPILER_CUDA_TOOLKIT_LIBRARY_ROOT@") +set(CMAKE_HIP_COMPILER_CUDA_TOOLKIT_VERSION "@CMAKE_HIP_COMPILER_CUDA_TOOLKIT_VERSION@") +set(CMAKE_HIP_COMPILER_CUDA_LIBRARY_ROOT "@CMAKE_HIP_COMPILER_CUDA_LIBRARY_ROOT@") + +set(CMAKE_HIP_ARCHITECTURES_ALL "@CMAKE_HIP_ARCHITECTURES_ALL@") +set(CMAKE_HIP_ARCHITECTURES_ALL_MAJOR "@CMAKE_HIP_ARCHITECTURES_ALL_MAJOR@") +set(CMAKE_HIP_ARCHITECTURES_NATIVE "@CMAKE_HIP_ARCHITECTURES_NATIVE@") + +set(CMAKE_HIP_CUDA_TOOLKIT_INCLUDE_DIRECTORIES "@CMAKE_HIP_CUDA_TOOLKIT_INCLUDE_DIRECTORIES@") + +set(CMAKE_HIP_HOST_IMPLICIT_LINK_LIBRARIES "@CMAKE_HIP_HOST_IMPLICIT_LINK_LIBRARIES@") +set(CMAKE_HIP_HOST_IMPLICIT_LINK_DIRECTORIES "@CMAKE_HIP_HOST_IMPLICIT_LINK_DIRECTORIES@") +set(CMAKE_HIP_HOST_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES "@CMAKE_HIP_HOST_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES@") set(CMAKE_HIP_IMPLICIT_INCLUDE_DIRECTORIES "@CMAKE_HIP_IMPLICIT_INCLUDE_DIRECTORIES@") set(CMAKE_HIP_IMPLICIT_LINK_LIBRARIES "@CMAKE_HIP_IMPLICIT_LINK_LIBRARIES@") set(CMAKE_HIP_IMPLICIT_LINK_DIRECTORIES "@CMAKE_HIP_IMPLICIT_LINK_DIRECTORIES@") set(CMAKE_HIP_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES "@CMAKE_HIP_IMPLICIT_LINK_FRAMEWORK_DIRECTORIES@") -set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED") +set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "@CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT@") set(CMAKE_AR "@CMAKE_AR@") set(CMAKE_HIP_COMPILER_AR "@CMAKE_HIP_COMPILER_AR@") diff --git a/Modules/CMakeHIPCompilerABI.hip b/Modules/CMakeHIPCompilerABI.hip index 6c912bd..7d8b815 100644 --- a/Modules/CMakeHIPCompilerABI.hip +++ b/Modules/CMakeHIPCompilerABI.hip @@ -1,9 +1,13 @@ -#ifndef __HIP__ +#if !defined(__HIP__) && !defined(__NVCC__) # error "A C or C++ compiler has been selected for HIP" #endif #include "CMakeCompilerABI.h" +#if defined(__NVCC__) +# include "CMakeCompilerCUDAArch.h" +#endif + int main(int argc, char* argv[]) { int require = 0; @@ -11,6 +15,16 @@ int main(int argc, char* argv[]) #if defined(ABI_ID) require += info_abi[argc]; #endif - (void)argv; + static_cast<void>(argv); + +#if defined(__NVCC__) + if (!cmakeCompilerCUDAArch()) { + // Convince the compiler that the non-zero return value depends + // on the info strings so they are not optimized out. + return require ? -1 : 1; + } + return 0; +#else return require; +#endif } diff --git a/Modules/CMakeHIPCompilerId.hip.in b/Modules/CMakeHIPCompilerId.hip.in index 3c4a1d4..4ac0f30 100644 --- a/Modules/CMakeHIPCompilerId.hip.in +++ b/Modules/CMakeHIPCompilerId.hip.in @@ -1,4 +1,4 @@ -#ifndef __HIP__ +#if !defined(__HIP__) && !defined(__NVCC__) # error "A C or C++ compiler has been selected for HIP" #endif diff --git a/Modules/CMakeHIPInformation.cmake b/Modules/CMakeHIPInformation.cmake index 41a98db..3995c36 100644 --- a/Modules/CMakeHIPInformation.cmake +++ b/Modules/CMakeHIPInformation.cmake @@ -8,6 +8,19 @@ else() endif() set(CMAKE_INCLUDE_FLAG_HIP "-I") +# Set implicit links early so compiler-specific modules can use them. +set(__IMPLICIT_LINKS) +foreach(dir ${CMAKE_HIP_HOST_IMPLICIT_LINK_DIRECTORIES}) + string(APPEND __IMPLICIT_LINKS " -L\"${dir}\"") +endforeach() +foreach(lib ${CMAKE_HIP_HOST_IMPLICIT_LINK_LIBRARIES}) + if(${lib} MATCHES "/") + string(APPEND __IMPLICIT_LINKS " \"${lib}\"") + else() + string(APPEND __IMPLICIT_LINKS " -l${lib}") + endif() +endforeach() + # Load compiler-specific information. if(CMAKE_HIP_COMPILER_ID) include(Compiler/${CMAKE_HIP_COMPILER_ID}-HIP OPTIONAL) @@ -129,7 +142,7 @@ endif() # compile a HIP file into an object file if(NOT CMAKE_HIP_COMPILE_OBJECT) set(CMAKE_HIP_COMPILE_OBJECT - "<CMAKE_HIP_COMPILER> <DEFINES> <INCLUDES> <FLAGS> -o <OBJECT> ${_CMAKE_COMPILE_AS_HIP_FLAG} -c <SOURCE>") + "<CMAKE_HIP_COMPILER> ${_CMAKE_HIP_EXTRA_FLAGS} <DEFINES> <INCLUDES> <FLAGS> -o <OBJECT> ${_CMAKE_COMPILE_AS_HIP_FLAG} -c <SOURCE>") endif() # compile a cu file into an executable diff --git a/Modules/CMakeTestHIPCompiler.cmake b/Modules/CMakeTestHIPCompiler.cmake index 686f055..ec54d80 100644 --- a/Modules/CMakeTestHIPCompiler.cmake +++ b/Modules/CMakeTestHIPCompiler.cmake @@ -10,7 +10,10 @@ if(CMAKE_HIP_COMPILER_FORCED) endif() set(__CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS}") -string(APPEND CMAKE_HIP_FLAGS " --cuda-host-only") + +if(CMAKE_HIP_COMPILER_ID STREQUAL "Clang") + string(APPEND CMAKE_HIP_FLAGS " --cuda-host-only") +endif() include(CMakeTestCompilerCommon) @@ -31,6 +34,13 @@ if(CMAKE_HIP_ABI_COMPILED) # The compiler worked so skip dedicated test below. set(CMAKE_HIP_COMPILER_WORKS TRUE) message(STATUS "Check for working HIP compiler: ${CMAKE_HIP_COMPILER} - skipped") + + if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA") + include(Internal/CMakeCUDAArchitecturesNative) + # Run the test binary to get: + # - CMAKE_HIP_ARCHITECTURES_NATIVE + cmake_cuda_architectures_native(HIP) + endif() endif() # This file is used by EnableLanguage in cmGlobalGenerator to @@ -42,7 +52,7 @@ if(NOT CMAKE_HIP_COMPILER_WORKS) PrintTestCompilerStatus("HIP") __TestCompiler_setTryCompileTargetType() string(CONCAT __TestCompiler_testHIPCompilerSource - "#ifndef __HIP__\n" + "#if !defined(__HIP__) && !defined(__NVCC__)\n" "# error \"The CMAKE_HIP_COMPILER is set to a C/CXX compiler\"\n" "#endif\n" "int main(){return 0;}\n") @@ -76,6 +86,16 @@ unset(__CMAKE_HIP_FLAGS) include(${CMAKE_ROOT}/Modules/CMakeDetermineCompileFeatures.cmake) CMAKE_DETERMINE_COMPILE_FEATURES(HIP) +if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA") + include(Internal/CMakeNVCCFilterImplicitInfo) + # Match arguments with cmake_nvcc_parse_implicit_info call in CMakeDetermineHIPCompiler. + cmake_nvcc_filter_implicit_info(HIP CMAKE_HIP_CUDA_) + + include(Internal/CMakeCUDAFilterImplicitLibs) + # Filter out implicit link libraries that should not be passed unconditionally. + cmake_cuda_filter_implicit_libs(CMAKE_HIP_IMPLICIT_LINK_LIBRARIES) +endif() + # Re-configure to save learned information. configure_file( ${CMAKE_ROOT}/Modules/CMakeHIPCompiler.cmake.in diff --git a/Modules/Compiler/NVIDIA-HIP.cmake b/Modules/Compiler/NVIDIA-HIP.cmake new file mode 100644 index 0000000..e144ff8 --- /dev/null +++ b/Modules/Compiler/NVIDIA-HIP.cmake @@ -0,0 +1,16 @@ +include(Compiler/NVIDIA) +__compiler_nvidia_cxx_standards(HIP) +__compiler_nvidia_cuda_flags(HIP) + +# The ROCm hip-lang package does not work for nvcc, +# so provide a minimal one ourselves. +add_library(hip-lang::device INTERFACE IMPORTED) +set(_CMAKE_HIP_DEVICE_RUNTIME_TARGET hip-lang::device) + +set(CMAKE_HIP_STANDARD_INCLUDE_DIRECTORIES "${CMAKE_HIP_COMPILER_ROCM_ROOT}/include") + +set(CMAKE_HIP_LINK_EXECUTABLE + "<CMAKE_HIP_HOST_LINK_LAUNCHER> <LINK_FLAGS> <OBJECTS> -o <TARGET> <LINK_LIBRARIES>${__IMPLICIT_LINKS}") +set(CMAKE_HIP_CREATE_SHARED_LIBRARY + "<CMAKE_HIP_HOST_LINK_LAUNCHER> <CMAKE_SHARED_LIBRARY_HIP_FLAGS> <LINK_FLAGS> <CMAKE_SHARED_LIBRARY_CREATE_HIP_FLAGS> <SONAME_FLAG><TARGET_SONAME> -o <TARGET> <OBJECTS> <LINK_LIBRARIES>${__IMPLICIT_LINKS}") +set(CMAKE_HIP_CREATE_SHARED_MODULE "${CMAKE_HIP_CREATE_SHARED_LIBRARY}") diff --git a/Source/cmCoreTryCompile.cxx b/Source/cmCoreTryCompile.cxx index 2ec62d9..0674d5a 100644 --- a/Source/cmCoreTryCompile.cxx +++ b/Source/cmCoreTryCompile.cxx @@ -78,6 +78,7 @@ std::string const kCMAKE_EXECUTABLE_ENABLE_EXPORTS = std::string const kCMAKE_SHARED_LIBRARY_ENABLE_EXPORTS = "CMAKE_SHARED_LIBRARY_ENABLE_EXPORTS"; std::string const kCMAKE_HIP_ARCHITECTURES = "CMAKE_HIP_ARCHITECTURES"; +std::string const kCMAKE_HIP_PLATFORM = "CMAKE_HIP_PLATFORM"; std::string const kCMAKE_HIP_RUNTIME_LIBRARY = "CMAKE_HIP_RUNTIME_LIBRARY"; std::string const kCMAKE_ISPC_INSTRUCTION_SETS = "CMAKE_ISPC_INSTRUCTION_SETS"; std::string const kCMAKE_ISPC_HEADER_SUFFIX = "CMAKE_ISPC_HEADER_SUFFIX"; @@ -1081,6 +1082,7 @@ cm::optional<cmTryCompileResult> cmCoreTryCompile::TryCompileCode( vars.insert(kCMAKE_EXECUTABLE_ENABLE_EXPORTS); vars.insert(kCMAKE_SHARED_LIBRARY_ENABLE_EXPORTS); vars.insert(kCMAKE_HIP_ARCHITECTURES); + vars.insert(kCMAKE_HIP_PLATFORM); vars.insert(kCMAKE_HIP_RUNTIME_LIBRARY); vars.insert(kCMAKE_ISPC_INSTRUCTION_SETS); vars.insert(kCMAKE_ISPC_HEADER_SUFFIX); diff --git a/Source/cmGeneratorTarget.cxx b/Source/cmGeneratorTarget.cxx index 80e3392..ca38be6 100644 --- a/Source/cmGeneratorTarget.cxx +++ b/Source/cmGeneratorTarget.cxx @@ -3579,9 +3579,11 @@ void cmGeneratorTarget::AddISPCTargetFlags(std::string& flags) const } } -void cmGeneratorTarget::AddHIPArchitectureFlags(std::string& flags) const +void cmGeneratorTarget::AddHIPArchitectureFlags(cmBuildStep compileOrLink, + std::string const& config, + std::string& flags) const { - const std::string& arch = this->GetSafeProperty("HIP_ARCHITECTURES"); + std::string arch = this->GetSafeProperty("HIP_ARCHITECTURES"); if (arch.empty()) { this->Makefile->IssueMessage(MessageType::FATAL_ERROR, @@ -3594,6 +3596,11 @@ void cmGeneratorTarget::AddHIPArchitectureFlags(std::string& flags) const return; } + if (this->Makefile->GetSafeDefinition("CMAKE_HIP_PLATFORM") == "nvidia") { + return this->AddCUDAArchitectureFlagsImpl(compileOrLink, config, "HIP", + std::move(arch), flags); + } + cmList options(arch); for (std::string& option : options) { diff --git a/Source/cmGeneratorTarget.h b/Source/cmGeneratorTarget.h index 6a385ea..751f907 100644 --- a/Source/cmGeneratorTarget.h +++ b/Source/cmGeneratorTarget.h @@ -504,7 +504,9 @@ public: std::string& flags) const; void AddCUDAToolkitFlags(std::string& flags) const; - void AddHIPArchitectureFlags(std::string& flags) const; + void AddHIPArchitectureFlags(cmBuildStep compileOrLink, + std::string const& config, + std::string& flags) const; void AddISPCTargetFlags(std::string& flags) const; diff --git a/Source/cmLocalGenerator.cxx b/Source/cmLocalGenerator.cxx index f1fe7df..fe8d502 100644 --- a/Source/cmLocalGenerator.cxx +++ b/Source/cmLocalGenerator.cxx @@ -87,6 +87,8 @@ static auto ruleReplaceVars = { "CMAKE_${LANG}_COMPILER", "CMAKE_TAPI", "CMAKE_CUDA_HOST_COMPILER", "CMAKE_CUDA_HOST_LINK_LAUNCHER", + "CMAKE_HIP_HOST_COMPILER", + "CMAKE_HIP_HOST_LINK_LAUNCHER", "CMAKE_CL_SHOWINCLUDES_PREFIX" }; cmLocalGenerator::cmLocalGenerator(cmGlobalGenerator* gg, cmMakefile* makefile) @@ -2055,7 +2057,7 @@ void cmLocalGenerator::AddLanguageFlags(std::string& flags, this->Makefile->GetSafeDefinition("CMAKE_CXX_SIMULATE_ID"); } } else if (lang == "HIP") { - target->AddHIPArchitectureFlags(flags); + target->AddHIPArchitectureFlags(compileOrLink, config, flags); } // Add VFS Overlay for Clang compilers diff --git a/Tests/HIP/ArchitectureOff/CMakeLists.txt b/Tests/HIP/ArchitectureOff/CMakeLists.txt index bccb3b4..18f3a1e 100644 --- a/Tests/HIP/ArchitectureOff/CMakeLists.txt +++ b/Tests/HIP/ArchitectureOff/CMakeLists.txt @@ -2,7 +2,17 @@ cmake_minimum_required(VERSION 3.18) project(HIPArchitecture HIP) # Make sure CMake doesn't pass architectures if HIP_ARCHITECTURES is OFF. -string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908") +set(CMAKE_HIP_ARCHITECTURES OFF) + +# Pass our own architecture flags instead. +if(CMAKE_HIP_PLATFORM STREQUAL "amd") + string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908") +elseif(CMAKE_HIP_PLATFORM STREQUAL "nvidia") + string(APPEND CMAKE_HIP_FLAGS " -arch=sm_52") +endif() add_executable(HIPOnlyArchitectureOff main.hip) -set_property(TARGET HIPOnlyArchitectureOff PROPERTY HIP_ARCHITECTURES OFF) +get_property(hip_archs TARGET HIPOnlyArchitectureOff PROPERTY HIP_ARCHITECTURES) +if(NOT hip_archs STREQUAL "OFF") + message(FATAL_ERROR "CMAKE_HIP_ARCHITECTURES did not initialize HIP_ARCHITECTURES") +endif() diff --git a/Tests/HIP/CMakeLists.txt b/Tests/HIP/CMakeLists.txt index 9499be8..26d7459 100644 --- a/Tests/HIP/CMakeLists.txt +++ b/Tests/HIP/CMakeLists.txt @@ -9,7 +9,10 @@ add_hip_test_macro(HIP.CompileFlags HIPOnlyCompileFlags) add_hip_test_macro(HIP.EnableStandard HIPEnableStandard) add_hip_test_macro(HIP.InferHipLang1 HIPInferHipLang1) add_hip_test_macro(HIP.InferHipLang2 HIPInferHipLang2) -add_hip_test_macro(HIP.MathFunctions HIPOnlyMathFunctions) +if(CMake_TEST_HIP STREQUAL "amd") + # The NVIDIA CUDA compiler cannot handle device lambda markup. + add_hip_test_macro(HIP.MathFunctions HIPOnlyMathFunctions) +endif() add_hip_test_macro(HIP.MixedLanguage HIPMixedLanguage) add_hip_test_macro(HIP.TryCompile HIPOnlyTryCompile) add_hip_test_macro(HIP.WithDefs HIPOnlyWithDefs) diff --git a/Tests/HIP/CompileFlags/CMakeLists.txt b/Tests/HIP/CompileFlags/CMakeLists.txt index c808313..a3adb7b 100644 --- a/Tests/HIP/CompileFlags/CMakeLists.txt +++ b/Tests/HIP/CompileFlags/CMakeLists.txt @@ -3,6 +3,11 @@ project(CompileFlags HIP) add_executable(HIPOnlyCompileFlags main.hip) -set_property(TARGET HIPOnlyCompileFlags PROPERTY HIP_ARCHITECTURES gfx803) +if(CMAKE_HIP_PLATFORM STREQUAL "amd") + set(hip_archs gfx803) +elseif(CMAKE_HIP_PLATFORM STREQUAL "nvidia") + set(hip_archs 52) +endif() +set_property(TARGET HIPOnlyCompileFlags PROPERTY HIP_ARCHITECTURES ${hip_archs}) target_compile_options(HIPOnlyCompileFlags PRIVATE -DALWAYS_DEFINE) diff --git a/Tests/HIP/MathFunctions/CMakeLists.txt b/Tests/HIP/MathFunctions/CMakeLists.txt index 81e3ddb..7e768e8 100644 --- a/Tests/HIP/MathFunctions/CMakeLists.txt +++ b/Tests/HIP/MathFunctions/CMakeLists.txt @@ -14,5 +14,9 @@ project(MathFunctions HIP) # that hip needs that inject support for __half support # add_executable(HIPOnlyMathFunctions main.hip) -target_compile_options(HIPOnlyMathFunctions PRIVATE -Werror) +if(CMAKE_HIP_COMPILER_ID STREQUAL "NVIDIA") + target_compile_options(HIPOnlyMathFunctions PRIVATE "SHELL:-Werror all-warnings") +elseif(CMAKE_HIP_COMPILER_ID STREQUAL "Clang") + target_compile_options(HIPOnlyMathFunctions PRIVATE "-Werror") +endif() target_compile_features(HIPOnlyMathFunctions PRIVATE hip_std_14) diff --git a/Tests/HIP/TryCompile/CMakeLists.txt b/Tests/HIP/TryCompile/CMakeLists.txt index 92a834c..1022a58 100644 --- a/Tests/HIP/TryCompile/CMakeLists.txt +++ b/Tests/HIP/TryCompile/CMakeLists.txt @@ -4,7 +4,12 @@ project (TryCompile HIP) #Goal for this example: # Verify try_compile with HIP language works set(CMAKE_HIP_STANDARD 14) -set(CMAKE_HIP_ARCHITECTURES gfx803 gfx900) + +if(CMAKE_HIP_PLATFORM STREQUAL "amd") + set(CMAKE_HIP_ARCHITECTURES gfx803 gfx900) +elseif(CMAKE_HIP_PLATFORM STREQUAL "nvidia") + set(CMAKE_HIP_ARCHITECTURES 52) +endif() set(CMAKE_TRY_COMPILE_TARGET_TYPE STATIC_LIBRARY) try_compile(result "${CMAKE_CURRENT_BINARY_DIR}" diff --git a/Tests/HIP/WithDefs/CMakeLists.txt b/Tests/HIP/WithDefs/CMakeLists.txt index 270f957..5602111 100644 --- a/Tests/HIP/WithDefs/CMakeLists.txt +++ b/Tests/HIP/WithDefs/CMakeLists.txt @@ -2,12 +2,11 @@ cmake_minimum_required(VERSION 3.18) project (WithDefs HIP) -set(CMAKE_HIP_ARCHITECTURES OFF) set(release_compile_defs DEFREL) #Goal for this example: -#build a executable that needs to be passed a complex define through add_definitions -#this verifies we can pass C++ style attributes to hipcc +#Build an executable that needs to be passed a complex define through add_definitions. +#Verify we can pass C++ style attributes to the HIP compiler. add_definitions("-DPACKED_DEFINE=[[gnu::packed]]") add_executable(HIPOnlyWithDefs main.hip.cpp) @@ -17,9 +16,8 @@ target_compile_features(HIPOnlyWithDefs PRIVATE hip_std_17) target_compile_options(HIPOnlyWithDefs PRIVATE - --offload-arch=gfx900 -DFLAG_COMPILE_LANG_$<COMPILE_LANGUAGE> - $<$<HIP_COMPILER_ID:Clang>:-DFLAG_LANG_IS_HIP=$<COMPILE_LANGUAGE:HIP>> # Host-only defines are possible only on NVCC. + -DFLAG_LANG_IS_HIP=$<COMPILE_LANGUAGE:HIP> ) target_compile_definitions(HIPOnlyWithDefs diff --git a/Tests/HIP/WithDefs/main.hip.cpp b/Tests/HIP/WithDefs/main.hip.cpp index a8f2d18..b69fa02 100644 --- a/Tests/HIP/WithDefs/main.hip.cpp +++ b/Tests/HIP/WithDefs/main.hip.cpp @@ -51,6 +51,10 @@ static __global__ void DetermineIfValidHIPDevice() # undef PACKED_DEFINE # define PACKED_DEFINE #endif +#ifdef __NVCC__ +# undef PACKED_DEFINE +# define PACKED_DEFINE +#endif struct PACKED_DEFINE result_type { bool valid; |