Merge topic 'hip-nvidia'

6546aa2a2a ci: Add HIP job using CUDA on NVIDIA GPUs
18158bf81c HIP: Add support for NVIDIA GPUs
127b6fa06b HIP: Add CMAKE_HIP_PLATFORM variable to specify GPU platform
90e23f40ee Tests/HIP/WithDefs: Clean up test case
9ebdf3281f Tests/HIP/ArchitectureOff: Cover HIP_ARCHITECTURES initialization
cfec29196e ci: Add CUDA 11.8 to HIP 5.5 image
26470eb987 ci: Put HIP GPU platform in CMake_TEST_HIP

Acked-by: Kitware Robot <[email protected]>
Merge-request: !8817

.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

.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")

.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")

.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")

.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")

.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 <[email protected]>
 
+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 \

.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

.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

.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:

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

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

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``).

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.

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.

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)

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

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@")

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
 }

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
 

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

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

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}")

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);

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) {

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;
 

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

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()

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)

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)

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)

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}"

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

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;