From 26470eb98717a7f34e7ffe55d387040427f3c8c5 Mon Sep 17 00:00:00 2001 From: Brad King Date: Mon, 18 Sep 2023 14:18:41 -0400 Subject: ci: Put HIP GPU platform in CMake_TEST_HIP The `CMake_TEST_HIP` option is added explicitly to builds where we want the tests to run, so we can set it to a value indicating the HIP GPU platform. --- .gitlab/ci/configure_debian12_hip_radeon.cmake | 2 +- .gitlab/ci/configure_fedora38_hip_radeon.cmake | 2 +- .gitlab/ci/configure_hip5.5_radeon.cmake | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) 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_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") -- cgit v0.12 From cfec29196ef13b62fc391bcae6ebbde5e377fdb2 Mon Sep 17 00:00:00 2001 From: Brad King Date: Mon, 18 Sep 2023 18:34:05 -0400 Subject: ci: Add CUDA 11.8 to HIP 5.5 image --- .gitlab/ci/docker/hip5.5/Dockerfile | 13 +++++++++++-- .gitlab/ci/docker/hip5.5/deps_packages.lst | 17 +++++++++++++++++ .gitlab/os-linux.yml | 2 +- 3 files changed, 29 insertions(+), 3 deletions(-) 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 +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/os-linux.yml b/.gitlab/os-linux.yml index 8398108..408feee 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" -- cgit v0.12 From 9ebdf3281f431589af88e94470ef3274f2c6e753 Mon Sep 17 00:00:00 2001 From: Brad King Date: Mon, 18 Sep 2023 15:06:08 -0400 Subject: Tests/HIP/ArchitectureOff: Cover HIP_ARCHITECTURES initialization Verify that the property is initialized by `CMAKE_HIP_ARCHITECTURES`. --- Tests/HIP/ArchitectureOff/CMakeLists.txt | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/Tests/HIP/ArchitectureOff/CMakeLists.txt b/Tests/HIP/ArchitectureOff/CMakeLists.txt index bccb3b4..b40301a 100644 --- a/Tests/HIP/ArchitectureOff/CMakeLists.txt +++ b/Tests/HIP/ArchitectureOff/CMakeLists.txt @@ -2,7 +2,11 @@ cmake_minimum_required(VERSION 3.18) project(HIPArchitecture HIP) # Make sure CMake doesn't pass architectures if HIP_ARCHITECTURES is OFF. +set(CMAKE_HIP_ARCHITECTURES OFF) string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908") 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() -- cgit v0.12 From 90e23f40ee27c0990b30b3640731e89539cb3990 Mon Sep 17 00:00:00 2001 From: Brad King Date: Thu, 10 Aug 2023 16:05:25 -0400 Subject: Tests/HIP/WithDefs: Clean up test case Fix code copied from equivalent CUDA test. Drop `CMAKE_HIP_ARCHITECTURES=OFF`. That is already covered by a dedicated `HIP.ArchitectureOff` test. --- Tests/HIP/WithDefs/CMakeLists.txt | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) 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_$ - $<$:-DFLAG_LANG_IS_HIP=$> # Host-only defines are possible only on NVCC. + -DFLAG_LANG_IS_HIP=$ ) target_compile_definitions(HIPOnlyWithDefs -- cgit v0.12 From 127b6fa06bf53ad9f31d041a7d11434ca2856c8e Mon Sep 17 00:00:00 2001 From: Brad King Date: Mon, 18 Sep 2023 14:11:41 -0400 Subject: HIP: Add CMAKE_HIP_PLATFORM variable to specify GPU platform For now, require the value to be `amd`, since that is the only platform we currently support. --- Help/manual/cmake-variables.7.rst | 1 + Help/prop_tgt/HIP_ARCHITECTURES.rst | 3 ++- Help/release/dev/hip-nvidia.rst | 6 ++++++ Help/variable/CMAKE_HIP_ARCHITECTURES.rst | 10 +++++++--- Help/variable/CMAKE_HIP_PLATFORM.rst | 19 +++++++++++++++++++ Modules/CMakeDetermineHIPCompiler.cmake | 28 ++++++++++++++++++++-------- Source/cmCoreTryCompile.cxx | 2 ++ Tests/HIP/ArchitectureOff/CMakeLists.txt | 6 +++++- Tests/HIP/CompileFlags/CMakeLists.txt | 5 ++++- Tests/HIP/TryCompile/CMakeLists.txt | 5 ++++- 10 files changed, 70 insertions(+), 15 deletions(-) create mode 100644 Help/release/dev/hip-nvidia.rst create mode 100644 Help/variable/CMAKE_HIP_PLATFORM.rst 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..f4617e7 --- /dev/null +++ b/Help/release/dev/hip-nvidia.rst @@ -0,0 +1,6 @@ +hip-nvidia +---------- + +* The :variable:`CMAKE_HIP_PLATFORM` variable was added to specify + the GPU platform for which HIP language sources are to be compiled + (``amd``). 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..1715066 --- /dev/null +++ b/Help/variable/CMAKE_HIP_PLATFORM.rst @@ -0,0 +1,19 @@ +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 + +If not specified, the default is ``amd``. + +:variable:`CMAKE_HIP_ARCHITECTURES` entries are interpreted with +as architectures of the GPU platform. + +:variable:`CMAKE_HIP_COMPILER _COMPILER>` must target +the same GPU platform. diff --git a/Modules/CMakeDetermineHIPCompiler.cmake b/Modules/CMakeDetermineHIPCompiler.cmake index 9a40e82..e55648a 100644 --- a/Modules/CMakeDetermineHIPCompiler.cmake +++ b/Modules/CMakeDetermineHIPCompiler.cmake @@ -10,6 +10,16 @@ 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) + set(CMAKE_HIP_PLATFORM "amd" CACHE STRING "HIP platform" FORCE) +endif() +if(NOT CMAKE_HIP_PLATFORM MATCHES "^(amd)$") + message(FATAL_ERROR + "The CMAKE_HIP_PLATFORM has unsupported value:\n" + " '${CMAKE_HIP_PLATFORM}'\n" + "It must be 'amd'." + ) +endif() if(NOT CMAKE_HIP_COMPILER) set(CMAKE_HIP_COMPILER_INIT NOTFOUND) @@ -34,15 +44,17 @@ 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 "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() 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 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/Tests/HIP/ArchitectureOff/CMakeLists.txt b/Tests/HIP/ArchitectureOff/CMakeLists.txt index b40301a..9d0bf05 100644 --- a/Tests/HIP/ArchitectureOff/CMakeLists.txt +++ b/Tests/HIP/ArchitectureOff/CMakeLists.txt @@ -3,7 +3,11 @@ project(HIPArchitecture HIP) # Make sure CMake doesn't pass architectures if HIP_ARCHITECTURES is OFF. set(CMAKE_HIP_ARCHITECTURES OFF) -string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908") + +# Pass our own architecture flags instead. +if(CMAKE_HIP_PLATFORM STREQUAL "amd") + string(APPEND CMAKE_HIP_FLAGS " --offload-arch=gfx908") +endif() add_executable(HIPOnlyArchitectureOff main.hip) get_property(hip_archs TARGET HIPOnlyArchitectureOff PROPERTY HIP_ARCHITECTURES) diff --git a/Tests/HIP/CompileFlags/CMakeLists.txt b/Tests/HIP/CompileFlags/CMakeLists.txt index c808313..46a94a3 100644 --- a/Tests/HIP/CompileFlags/CMakeLists.txt +++ b/Tests/HIP/CompileFlags/CMakeLists.txt @@ -3,6 +3,9 @@ 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) +endif() +set_property(TARGET HIPOnlyCompileFlags PROPERTY HIP_ARCHITECTURES ${hip_archs}) target_compile_options(HIPOnlyCompileFlags PRIVATE -DALWAYS_DEFINE) diff --git a/Tests/HIP/TryCompile/CMakeLists.txt b/Tests/HIP/TryCompile/CMakeLists.txt index 92a834c..c98e59c 100644 --- a/Tests/HIP/TryCompile/CMakeLists.txt +++ b/Tests/HIP/TryCompile/CMakeLists.txt @@ -4,7 +4,10 @@ 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) +endif() set(CMAKE_TRY_COMPILE_TARGET_TYPE STATIC_LIBRARY) try_compile(result "${CMAKE_CURRENT_BINARY_DIR}" -- cgit v0.12 From 18158bf81ccb5269afcfa53cd671cf7f56dc70b5 Mon Sep 17 00:00:00 2001 From: Brad King Date: Mon, 18 Sep 2023 15:46:39 -0400 Subject: HIP: Add support for NVIDIA GPUs Add support for using the CUDA Toolkit's NVCC to compile HIP code. Fixes: #25143 --- Help/release/dev/hip-nvidia.rst | 5 +- Help/variable/CMAKE_HIP_PLATFORM.rst | 5 +- Modules/CMakeCompilerIdDetection.cmake | 2 +- Modules/CMakeDetermineHIPCompiler.cmake | 102 ++++++++++++++++++++++++++++--- Modules/CMakeHIPCompiler.cmake.in | 19 +++++- Modules/CMakeHIPCompilerABI.hip | 18 +++++- Modules/CMakeHIPCompilerId.hip.in | 2 +- Modules/CMakeHIPInformation.cmake | 15 ++++- Modules/CMakeTestHIPCompiler.cmake | 24 +++++++- Modules/Compiler/NVIDIA-HIP.cmake | 16 +++++ Source/cmGeneratorTarget.cxx | 11 +++- Source/cmGeneratorTarget.h | 4 +- Source/cmLocalGenerator.cxx | 4 +- Tests/HIP/ArchitectureOff/CMakeLists.txt | 2 + Tests/HIP/CMakeLists.txt | 5 +- Tests/HIP/CompileFlags/CMakeLists.txt | 2 + Tests/HIP/MathFunctions/CMakeLists.txt | 6 +- Tests/HIP/TryCompile/CMakeLists.txt | 2 + Tests/HIP/WithDefs/main.hip.cpp | 4 ++ 19 files changed, 221 insertions(+), 27 deletions(-) create mode 100644 Modules/Compiler/NVIDIA-HIP.cmake diff --git a/Help/release/dev/hip-nvidia.rst b/Help/release/dev/hip-nvidia.rst index f4617e7..1d9814e 100644 --- a/Help/release/dev/hip-nvidia.rst +++ b/Help/release/dev/hip-nvidia.rst @@ -1,6 +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``). + (``amd`` or ``nvidia``). diff --git a/Help/variable/CMAKE_HIP_PLATFORM.rst b/Help/variable/CMAKE_HIP_PLATFORM.rst index 1715066..5e3a2b7 100644 --- a/Help/variable/CMAKE_HIP_PLATFORM.rst +++ b/Help/variable/CMAKE_HIP_PLATFORM.rst @@ -10,7 +10,10 @@ The value must be one of: ``amd`` AMD GPUs -If not specified, the default is ``amd``. +``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. 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 e55648a..21322f1 100644 --- a/Modules/CMakeDetermineHIPCompiler.cmake +++ b/Modules/CMakeDetermineHIPCompiler.cmake @@ -11,13 +11,21 @@ if( NOT ( ("${CMAKE_GENERATOR}" MATCHES "Make") OR endif() if(NOT CMAKE_HIP_PLATFORM) - set(CMAKE_HIP_PLATFORM "amd" CACHE STRING "HIP platform" FORCE) + 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)$") +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'." + "It must be 'amd' or 'nvidia'." ) endif() @@ -44,7 +52,9 @@ if(NOT CMAKE_HIP_COMPILER) # finally list compilers to try if(NOT CMAKE_HIP_COMPILER_INIT) - if(CMAKE_HIP_PLATFORM STREQUAL "amd") + 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. @@ -75,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() @@ -116,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 @@ -177,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" @@ -197,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(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 - " -o ${_CMAKE_COMPILE_AS_HIP_FLAG} -c ") + " ${_CMAKE_HIP_EXTRA_FLAGS} -o ${_CMAKE_COMPILE_AS_HIP_FLAG} -c ") 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 + " -o ${__IMPLICIT_LINKS}") +set(CMAKE_HIP_CREATE_SHARED_LIBRARY + " -o ${__IMPLICIT_LINKS}") +set(CMAKE_HIP_CREATE_SHARED_MODULE "${CMAKE_HIP_CREATE_SHARED_LIBRARY}") 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 63c8aa8..2325579 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) @@ -2058,7 +2060,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 9d0bf05..18f3a1e 100644 --- a/Tests/HIP/ArchitectureOff/CMakeLists.txt +++ b/Tests/HIP/ArchitectureOff/CMakeLists.txt @@ -7,6 +7,8 @@ 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) 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 46a94a3..a3adb7b 100644 --- a/Tests/HIP/CompileFlags/CMakeLists.txt +++ b/Tests/HIP/CompileFlags/CMakeLists.txt @@ -5,6 +5,8 @@ add_executable(HIPOnlyCompileFlags main.hip) 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}) 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 c98e59c..1022a58 100644 --- a/Tests/HIP/TryCompile/CMakeLists.txt +++ b/Tests/HIP/TryCompile/CMakeLists.txt @@ -7,6 +7,8 @@ set(CMAKE_HIP_STANDARD 14) 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) 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; -- cgit v0.12 From 6546aa2a2a2f695237a2f14745d18fc8a52b1fcf Mon Sep 17 00:00:00 2001 From: Brad King Date: Mon, 18 Sep 2023 18:34:27 -0400 Subject: ci: Add HIP job using CUDA on NVIDIA GPUs --- .gitlab-ci.yml | 10 ++++++++++ .gitlab/ci/configure_hip5.5_nvidia.cmake | 3 +++ .gitlab/ci/env_hip5.5_nvidia.sh | 4 ++++ .gitlab/os-linux.yml | 7 +++++++ 4 files changed, 24 insertions(+) create mode 100644 .gitlab/ci/configure_hip5.5_nvidia.cmake create mode 100644 .gitlab/ci/env_hip5.5_nvidia.sh diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml index 7a66494..6d0c144 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_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/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 408feee..d42340c 100644 --- a/.gitlab/os-linux.yml +++ b/.gitlab/os-linux.yml @@ -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: -- cgit v0.12