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

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

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.

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

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

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

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)

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

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

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

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

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;