HIP: Automatically inject the `hip::device` runtime target

Any target that might need to link to hip code needs the `hip::device`
target

Modules/CMakeHIPCompiler.cmake.in

@@ -22,7 +22,7 @@ set(CMAKE_HIP_COMPILER_ENV_VAR "HIPCXX")
 set(CMAKE_HIP_COMPILER_LOADED 1)
 set(CMAKE_HIP_COMPILER_ID_RUN 1)
 set(CMAKE_HIP_SOURCE_FILE_EXTENSIONS hip)
-set(CMAKE_HIP_LINKER_PREFERENCE 15)
+set(CMAKE_HIP_LINKER_PREFERENCE 90)
 set(CMAKE_HIP_LINKER_PREFERENCE_PROPAGATES 1)
 
 set(CMAKE_HIP_SIZEOF_DATA_PTR "@CMAKE_HIP_SIZEOF_DATA_PTR@")
@@ -48,8 +48,7 @@ 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_COMPILER_DEVICE_LIBRARY_ROOT_DIR "@CMAKE_HIP_COMPILER_DEVICE_LIBRARY_ROOT_DIR@")
-@_SET_CMAKE_HIP_DEVICE_LIBRARY_INCLUSION@
+set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED")
 
 set(CMAKE_AR "@CMAKE_AR@")
 set(CMAKE_HIP_COMPILER_AR "@CMAKE_HIP_COMPILER_AR@")

Modules/CMakeHIPInformation.cmake

@@ -130,3 +130,10 @@ if(NOT CMAKE_HIP_LINK_EXECUTABLE)
 endif()
 
 set(CMAKE_HIP_INFORMATION_LOADED 1)
+
+# Load the file and find the relevant HIP runtime.
+# This file will only exist after all compiler detection has finished
+include(${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPRuntime.cmake OPTIONAL)
+if(COMMAND _CMAKE_FIND_HIP_RUNTIME)
+  _CMAKE_FIND_HIP_RUNTIME()
+endif()

Modules/CMakeHIPRuntime.cmake.in

@@ -0,0 +1,99 @@
+# Distributed under the OSI-approved BSD 3-Clause License.  See accompanying
+# file Copyright.txt or https://cmake.org/licensing for details.
+
+
+function(_CMAKE_FIND_HIP_RUNTIME )
+  # Determined when hipcc is the HIP compiler
+  set(_CMAKE_HIP_COMPILER_ROCM_ROOT "@_CMAKE_HIP_COMPILER_ROCM_ROOT@")
+
+  # Forward facing value that can be provided by the user
+  set(CMAKE_HIP_COMPILER_TOOLKIT_ROOT @CMAKE_HIP_COMPILER_TOOLKIT_ROOT@)
+
+  if(NOT DEFINED _CMAKE_HIP_DEVICE_RUNTIME_TARGET)
+    set(message_on_found TRUE)
+  endif()
+
+  set(explicit_search_only FALSE)
+  set(rocm_root_dirs )
+  if(DEFINED CMAKE_HIP_COMPILER_TOOLKIT_ROOT)
+    set(rocm_root_dirs "${CMAKE_HIP_COMPILER_TOOLKIT_ROOT}")
+    set(explicit_search_only TRUE)
+    set(error_message_location "the variable CMAKE_HIP_COMPILER_TOOLKIT_ROOT [\"${CMAKE_HIP_COMPILER_TOOLKIT_ROOT}\"]")
+  elseif(DEFINED ENV{CMAKE_HIP_COMPILER_TOOLKIT_ROOT})
+    set(rocm_root_dirs "$ENV{CMAKE_HIP_COMPILER_TOOLKIT_ROOT}")
+    set(explicit_search_only TRUE)
+    set(error_message_location "CMAKE_HIP_COMPILER_TOOLKIT_ROOT")
+    set(error_message_location "the environment variable CMAKE_HIP_COMPILER_TOOLKIT_ROOT [\"$ENV{CMAKE_HIP_COMPILER_TOOLKIT_ROOT}\"]")
+  elseif(DEFINED _CMAKE_HIP_COMPILER_ROCM_ROOT)
+    set(rocm_root_dirs "${_CMAKE_HIP_COMPILER_ROCM_ROOT}")
+    set(explicit_search_only TRUE)
+    set(error_message_location "the associated hipconfig --rocmpath [\"${_CMAKE_HIP_COMPILER_ROCM_ROOT}\"]")
+  endif()
+
+  # Guess on where rocm is installed
+  if(NOT rocm_root_dirs AND (UNIX AND NOT APPLE))
+    set(platform_base "/opt/rocm-")
+
+    # Finad all default rocm installations
+    file(GLOB possible_paths "${platform_base}*")
+
+    set(versions)
+    foreach(p ${possible_paths})
+      # Extract version number from end of string
+      string(REGEX MATCH "[0-9]+\\.[0-9]+\\.[0-9]+$" p_version ${p})
+      if(IS_DIRECTORY ${p} AND p_version)
+        list(APPEND versions ${p_version})
+      endif()
+    endforeach()
+
+    # Sort numerically in descending order, so we try the newest versions first.
+    list(SORT versions COMPARE NATURAL ORDER DESCENDING)
+
+    # With a descending list of versions, populate possible paths to search.
+    set(rocm_root_dirs "/opt/rocm")
+    foreach(v IN LISTS versions)
+      list(APPEND rocm_root_dirs "${platform_base}${v}")
+    endforeach()
+  endif()
+
+  set(search_rel_path "/lib/cmake/hip-lang/")
+  list(TRANSFORM rocm_root_dirs APPEND "${search_rel_path}")
+
+  find_package(hip-lang
+    CONFIG
+    PATHS ${rocm_root_dirs}
+    QUIET
+    NO_DEFAULT_PATH
+  )
+  if(NOT DEFINED _CMAKE_HIP_DEVICE_RUNTIME_TARGET AND NOT explicit_search_only)
+    find_package(hip-lang CONFIG QUIET)
+  endif()
+
+  if(DEFINED _CMAKE_HIP_DEVICE_RUNTIME_TARGET)
+    set(CMAKE_HIP_RUNTIME_LIBRARIES_STATIC
+      ${CMAKE_HIP_RUNTIME_LIBRARIES_STATIC}
+      ${_CMAKE_HIP_DEVICE_RUNTIME_TARGET} PARENT_SCOPE)
+    set(CMAKE_HIP_RUNTIME_LIBRARIES_SHARED
+      ${CMAKE_HIP_RUNTIME_LIBRARIES_SHARED}
+      ${_CMAKE_HIP_DEVICE_RUNTIME_TARGET} PARENT_SCOPE)
+  endif()
+
+  if(DEFINED _CMAKE_HIP_DEVICE_RUNTIME_TARGET AND message_on_found)
+    message(STATUS "Found HIP runtime: ${hip-lang_DIR}")
+  elseif(NOT DEFINED _CMAKE_HIP_DEVICE_RUNTIME_TARGET)
+    if(explicit_search_only)
+      set(error_message "Failed to find the HIP runtime, Could not find hip-lang-config.cmake at the following location(s):\n")
+      foreach(p IN LISTS rocm_root_dirs)
+        string(APPEND error_message "\t${p}\n")
+      endforeach()
+      string(APPEND "which are computed from the location specified by ${error_message_location}. \
+        Please specify CMAKE_HIP_COMPILER_TOOLKIT_ROOT to the location of")
+      message(FATAL_ERROR "${error_message}")
+    else()
+      message(FATAL_ERROR
+        "Failed to find the HIP runtime, Could not find hip-lang-config.cmake.\
+        Try setting CMAKE_HIP_COMPILER_TOOLKIT_ROOT")
+    endif()
+  endif()
+
+endfunction()

Modules/CMakeTestHIPCompiler.cmake

@@ -78,6 +78,20 @@ unset(__CMAKE_HIP_FLAGS)
 include(${CMAKE_ROOT}/Modules/CMakeDetermineCompileFeatures.cmake)
 CMAKE_DETERMINE_COMPILE_FEATURES(HIP)
 
+
+# Setup the following:
+# Configure the new template file CMakeHipRuntime.cmake to
+# - ${CMAKE_PLATFORM_INFO_DIR}/
+# This file will do the actual find_package query. We than have
+# CMakeHIPInformation.cmake include `CMakeHipRuntime`
+# So it is included once system information has been finished
+#
+configure_file(
+ ${CMAKE_ROOT}/Modules/CMakeHIPRuntime.cmake.in
+ ${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPRuntime.cmake
+ @ONLY
+)
+
 # Re-configure to save learned information.
 configure_file(
   ${CMAKE_ROOT}/Modules/CMakeHIPCompiler.cmake.in
@@ -96,3 +110,10 @@ endif()
 set(CMAKE_TRY_COMPILE_TARGET_TYPE ${__CMAKE_SAVED_TRY_COMPILE_TARGET_TYPE})
 unset(__CMAKE_SAVED_TRY_COMPILE_TARGET_TYPE)
 unset(__CMAKE_HIP_COMPILER_OUTPUT)
+
+# Load the file and find the relevant HIP runtime.
+# This file will only exist after all compiler detection has finished
+include(${CMAKE_PLATFORM_INFO_DIR}/CMakeHIPRuntime.cmake)
+if(COMMAND _CMAKE_FIND_HIP_RUNTIME)
+  _CMAKE_FIND_HIP_RUNTIME()
+endif()

Modules/Compiler/Clang-HIP.cmake

@@ -14,3 +14,7 @@ endif()
 set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED")
 set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_STATIC  "")
 set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_SHARED  "")
+
+# Populated by CMakeHIPRuntime.cmake
+set(CMAKE_HIP_RUNTIME_LIBRARIES_STATIC "")
+set(CMAKE_HIP_RUNTIME_LIBRARIES_SHARED "")

Modules/Compiler/ROCMClang-HIP.cmake

@@ -42,4 +42,8 @@ set(CMAKE_HIP_RUNTIME_LIBRARY_DEFAULT "SHARED")
 set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_STATIC  "")
 set(CMAKE_HIP_RUNTIME_LIBRARY_LINK_OPTIONS_SHARED  "")
 
+# Populated by CMakeHIPRuntime.cmake
+set(CMAKE_HIP_RUNTIME_LIBRARIES_STATIC "")
+set(CMAKE_HIP_RUNTIME_LIBRARIES_SHARED "")
+
 __compiler_check_default_language_standard(HIP 3.5 11)

Source/cmComputeLinkDepends.cxx

@@ -7,6 +7,7 @@
 #include <cstdio>
 #include <iterator>
 #include <sstream>
+#include <unordered_map>
 #include <utility>
 
 #include <cm/memory>
@@ -371,6 +372,12 @@ void cmComputeLinkDepends::FollowLinkEntry(BFSEntry qe)
       // This target provides its own link interface information.
       this->AddLinkEntries(depender_index, iface->Libraries);
       this->AddLinkObjects(iface->Objects);
+      for (auto const& language : iface->Languages) {
+        auto runtimeEntries = iface->LanguageRuntimeLibraries.find(language);
+        if (runtimeEntries != iface->LanguageRuntimeLibraries.end()) {
+          this->AddLinkEntries(depender_index, runtimeEntries->second);
+        }
+      }
 
       if (isIface) {
         return;
@@ -516,6 +523,13 @@ void cmComputeLinkDepends::AddDirectLinkEntries()
     this->Target->GetLinkImplementation(this->Config);
   this->AddLinkEntries(-1, impl->Libraries);
   this->AddLinkObjects(impl->Objects);
+
+  for (auto const& language : impl->Languages) {
+    auto runtimeEntries = impl->LanguageRuntimeLibraries.find(language);
+    if (runtimeEntries != impl->LanguageRuntimeLibraries.end()) {
+      this->AddLinkEntries(-1, runtimeEntries->second);
+    }
+  }
   for (cmLinkItem const& wi : impl->WrongConfigLibraries) {
     this->CheckWrongConfigItem(wi);
   }

Source/cmGeneratorTarget.cxx

@@ -1235,6 +1235,20 @@ bool cmGeneratorTarget::IsSystemIncludeDirectory(
                               &dagChecker, result, excludeImported, language);
     }
 
+    cmLinkImplementation const* impl = this->GetLinkImplementation(config);
+    if (impl != nullptr) {
+      auto runtimeEntries = impl->LanguageRuntimeLibraries.find(language);
+      if (runtimeEntries != impl->LanguageRuntimeLibraries.end()) {
+        for (auto const& lib : runtimeEntries->second) {
+          if (lib.Target) {
+            handleSystemIncludesDep(this->LocalGenerator, lib.Target, config,
+                                    this, &dagChecker, result, excludeImported,
+                                    language);
+          }
+        }
+      }
+    }
+
     std::for_each(result.begin(), result.end(),
                   cmSystemTools::ConvertToUnixSlashes);
     std::sort(result.begin(), result.end());
@@ -1474,31 +1488,80 @@ void AddLangSpecificImplicitIncludeDirectories(
   }
 }
 
+void addInterfaceEntry(cmGeneratorTarget const* headTarget,
+                       std::string const& config, std::string const& prop,
+                       std::string const& lang,
+                       cmGeneratorExpressionDAGChecker* dagChecker,
+                       EvaluatedTargetPropertyEntries& entries,
+                       bool usage_requirements_only,
+                       std::vector<cmLinkImplItem> const& libraries)
+{
+  for (cmLinkImplItem const& lib : libraries) {
+    if (lib.Target) {
+      EvaluatedTargetPropertyEntry ee(lib, lib.Backtrace);
+      // Pretend $<TARGET_PROPERTY:lib.Target,prop> appeared in our
+      // caller's property and hand-evaluate it as if it were compiled.
+      // Create a context as cmCompiledGeneratorExpression::Evaluate does.
+      cmGeneratorExpressionContext context(
+        headTarget->GetLocalGenerator(), config, false, headTarget, headTarget,
+        true, lib.Backtrace, lang);
+      cmExpandList(lib.Target->EvaluateInterfaceProperty(
+                     prop, &context, dagChecker, usage_requirements_only),
+                   ee.Values);
+      ee.ContextDependent = context.HadContextSensitiveCondition;
+      entries.Entries.emplace_back(std::move(ee));
+    }
+  }
+}
+
+// IncludeRuntimeInterface is used to break the cycle in computing
+// the necessary transitive dependencies of targets that can occur
+// now that we have implicit language runtime targets.
+//
+// To determine the set of languages that a target has we need to iterate
+// all the sources which includes transitive INTERFACE sources.
+// Therefore we can't determine what language runtimes are needed
+// for a target until after all sources are computed.
+//
+// Therefore while computing the applicable INTERFACE_SOURCES we
+// must ignore anything in LanguageRuntimeLibraries or we would
+// create a cycle ( INTERFACE_SOURCES requires LanguageRuntimeLibraries,
+// LanguageRuntimeLibraries requires INTERFACE_SOURCES).
+//
+enum class IncludeRuntimeInterface
+{
+  Yes,
+  No
+};
 void AddInterfaceEntries(cmGeneratorTarget const* headTarget,
                          std::string const& config, std::string const& prop,
                          std::string const& lang,
                          cmGeneratorExpressionDAGChecker* dagChecker,
                          EvaluatedTargetPropertyEntries& entries,
+                         IncludeRuntimeInterface searchRuntime,
                          bool usage_requirements_only = true)
 {
-  if (cmLinkImplementationLibraries const* impl =
-        headTarget->GetLinkImplementationLibraries(config)) {
-    entries.HadContextSensitiveCondition = impl->HadContextSensitiveCondition;
-    for (cmLinkImplItem const& lib : impl->Libraries) {
-      if (lib.Target) {
-        EvaluatedTargetPropertyEntry ee(lib, lib.Backtrace);
-        // Pretend $<TARGET_PROPERTY:lib.Target,prop> appeared in our
-        // caller's property and hand-evaluate it as if it were compiled.
-        // Create a context as cmCompiledGeneratorExpression::Evaluate does.
-        cmGeneratorExpressionContext context(
-          headTarget->GetLocalGenerator(), config, false, headTarget,
-          headTarget, true, lib.Backtrace, lang);
-        cmExpandList(lib.Target->EvaluateInterfaceProperty(
-                       prop, &context, dagChecker, usage_requirements_only),
-                     ee.Values);
-        ee.ContextDependent = context.HadContextSensitiveCondition;
-        entries.Entries.emplace_back(std::move(ee));
+  if (searchRuntime == IncludeRuntimeInterface::Yes) {
+    if (cmLinkImplementation const* impl =
+          headTarget->GetLinkImplementation(config)) {
+      entries.HadContextSensitiveCondition =
+        impl->HadContextSensitiveCondition;
+
+      auto runtimeLibIt = impl->LanguageRuntimeLibraries.find(lang);
+      if (runtimeLibIt != impl->LanguageRuntimeLibraries.end()) {
+        addInterfaceEntry(headTarget, config, prop, lang, dagChecker, entries,
+                          usage_requirements_only, runtimeLibIt->second);
       }
+      addInterfaceEntry(headTarget, config, prop, lang, dagChecker, entries,
+                        usage_requirements_only, impl->Libraries);
+    }
+  } else {
+    if (cmLinkImplementationLibraries const* impl =
+          headTarget->GetLinkImplementationLibraries(config)) {
+      entries.HadContextSensitiveCondition =
+        impl->HadContextSensitiveCondition;
+      addInterfaceEntry(headTarget, config, prop, lang, dagChecker, entries,
+                        usage_requirements_only, impl->Libraries);
     }
   }
 }
@@ -1656,7 +1719,8 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetSourceFilePaths(
   // Collect INTERFACE_SOURCES of all direct link-dependencies.
   EvaluatedTargetPropertyEntries linkInterfaceSourcesEntries;
   AddInterfaceEntries(this, config, "INTERFACE_SOURCES", std::string(),
-                      &dagChecker, linkInterfaceSourcesEntries);
+                      &dagChecker, linkInterfaceSourcesEntries,
+                      IncludeRuntimeInterface::No, true);
   std::vector<std::string>::size_type numFilesBefore = files.size();
   bool contextDependentInterfaceSources = processSources(
     this, linkInterfaceSourcesEntries, files, uniqueSrcs, debugSources);
@@ -3587,7 +3651,7 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetIncludeDirectories(
   }
 
   AddInterfaceEntries(this, config, "INTERFACE_INCLUDE_DIRECTORIES", lang,
-                      &dagChecker, entries);
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes);
 
   if (this->Makefile->IsOn("APPLE")) {
     if (cmLinkImplementationLibraries const* impl =
@@ -3812,7 +3876,7 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetCompileOptions(
     this, config, language, &dagChecker, this->CompileOptionsEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_COMPILE_OPTIONS", language,
-                      &dagChecker, entries);
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes);
 
   processOptions(this, entries, result, uniqueOptions, debugOptions,
                  "compile options", OptionsParse::Shell);
@@ -3854,7 +3918,8 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetCompileFeatures(
     this, config, std::string(), &dagChecker, this->CompileFeaturesEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_COMPILE_FEATURES",
-                      std::string(), &dagChecker, entries);
+                      std::string(), &dagChecker, entries,
+                      IncludeRuntimeInterface::Yes);
 
   processOptions(this, entries, result, uniqueFeatures, debugFeatures,
                  "compile features", OptionsParse::None);
@@ -3898,7 +3963,7 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetCompileDefinitions(
     this, config, language, &dagChecker, this->CompileDefinitionsEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_COMPILE_DEFINITIONS", language,
-                      &dagChecker, entries);
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes);
 
   if (!config.empty()) {
     std::string configPropName =
@@ -3956,7 +4021,7 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetPrecompileHeaders(
     this, config, language, &dagChecker, this->PrecompileHeadersEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_PRECOMPILE_HEADERS", language,
-                      &dagChecker, entries);
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes);
 
   std::vector<BT<std::string>> list;
   processOptions(this, entries, list, uniqueOptions, debugDefines,
@@ -4341,7 +4406,7 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetLinkOptions(
     this, config, language, &dagChecker, this->LinkOptionsEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_LINK_OPTIONS", language,
-                      &dagChecker, entries,
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes,
                       this->GetPolicyStatusCMP0099() != cmPolicies::NEW);
 
   processOptions(this, entries, result, uniqueOptions, debugOptions,
@@ -4600,7 +4665,7 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetLinkDirectories(
     this, config, language, &dagChecker, this->LinkDirectoriesEntries);
 
   AddInterfaceEntries(this, config, "INTERFACE_LINK_DIRECTORIES", language,
-                      &dagChecker, entries,
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes,
                       this->GetPolicyStatusCMP0099() != cmPolicies::NEW);
 
   processLinkDirectories(this, entries, result, uniqueDirectories,
@@ -4639,7 +4704,7 @@ std::vector<BT<std::string>> cmGeneratorTarget::GetLinkDepends(
     }
   }
   AddInterfaceEntries(this, config, "INTERFACE_LINK_DEPENDS", language,
-                      &dagChecker, entries,
+                      &dagChecker, entries, IncludeRuntimeInterface::Yes,
                       this->GetPolicyStatusCMP0099() != cmPolicies::NEW);
 
   processOptions(this, entries, result, uniqueOptions, false, "link depends",
@@ -6382,6 +6447,7 @@ cmLinkInterface const* cmGeneratorTarget::GetLinkInterface(
     iface.AllDone = true;
     if (iface.Exists) {
       this->ComputeLinkInterface(config, iface, head, secondPass);
+      this->ComputeLinkInterfaceRuntimeLibraries(config, iface);
     }
   }
 
@@ -6905,6 +6971,83 @@ void cmGeneratorTarget::ComputeLinkInterfaceLibraries(
   }
 }
 
+namespace {
+
+template <typename ReturnType>
+ReturnType constructItem(cmGeneratorTarget* target,
+                         cmListFileBacktrace const& bt);
+
+template <>
+inline cmLinkImplItem constructItem(cmGeneratorTarget* target,
+                                    cmListFileBacktrace const& bt)
+{
+  return cmLinkImplItem(cmLinkItem(target, false, bt), false);
+}
+
+template <>
+inline cmLinkItem constructItem(cmGeneratorTarget* target,
+                                cmListFileBacktrace const& bt)
+{
+  return cmLinkItem(target, false, bt);
+}
+
+template <typename ValueType>
+std::vector<ValueType> computeImplicitLanguageTargets(
+  std::string const& lang, std::string const& config,
+  cmGeneratorTarget const* currentTarget)
+{
+  cmListFileBacktrace bt;
+  std::vector<ValueType> result;
+  cmLocalGenerator* lg = currentTarget->GetLocalGenerator();
+
+  std::string const& runtimeLibrary =
+    currentTarget->GetRuntimeLinkLibrary(lang, config);
+  if (cmProp runtimeLinkOptions = currentTarget->Makefile->GetDefinition(
+        "CMAKE_" + lang + "_RUNTIME_LIBRARIES_" + runtimeLibrary)) {
+    std::vector<std::string> libsVec = cmExpandedList(*runtimeLinkOptions);
+    result.reserve(libsVec.size());
+
+    for (std::string const& i : libsVec) {
+      cmGeneratorTarget::TargetOrString resolved =
+        currentTarget->ResolveTargetReference(i, lg);
+      if (resolved.Target) {
+        result.emplace_back(constructItem<ValueType>(resolved.Target, bt));
+      }
+    }
+  }
+
+  return result;
+}
+}
+
+void cmGeneratorTarget::ComputeLinkInterfaceRuntimeLibraries(
+  const std::string& config, cmOptionalLinkInterface& iface) const
+{
+  for (std::string const& lang : iface.Languages) {
+    if ((lang == "CUDA" || lang == "HIP") &&
+        iface.LanguageRuntimeLibraries.find(lang) ==
+          iface.LanguageRuntimeLibraries.end()) {
+      auto implicitTargets =
+        computeImplicitLanguageTargets<cmLinkItem>(lang, config, this);
+      iface.LanguageRuntimeLibraries[lang] = std::move(implicitTargets);
+    }
+  }
+}
+
+void cmGeneratorTarget::ComputeLinkImplementationRuntimeLibraries(
+  const std::string& config, cmOptionalLinkImplementation& impl) const
+{
+  for (std::string const& lang : impl.Languages) {
+    if ((lang == "CUDA" || lang == "HIP") &&
+        impl.LanguageRuntimeLibraries.find(lang) ==
+          impl.LanguageRuntimeLibraries.end()) {
+      auto implicitTargets =
+        computeImplicitLanguageTargets<cmLinkImplItem>(lang, config, this);
+      impl.LanguageRuntimeLibraries[lang] = std::move(implicitTargets);
+    }
+  }
+}
+
 const cmLinkInterface* cmGeneratorTarget::GetImportLinkInterface(
   const std::string& config, cmGeneratorTarget const* headTarget,
   bool usage_requirements_only, bool secondPass) const
@@ -7169,6 +7312,7 @@ const cmLinkImplementation* cmGeneratorTarget::GetLinkImplementation(
   if (!impl.LanguagesDone) {
     impl.LanguagesDone = true;
     this->ComputeLinkImplementationLanguages(config, impl);
+    this->ComputeLinkImplementationRuntimeLibraries(config, impl);
   }
   return &impl;
 }

Source/cmGeneratorTarget.h

@@ -1111,6 +1111,12 @@ private:
   cmProp GetPropertyWithPairedLanguageSupport(std::string const& lang,
                                               const char* suffix) const;
 
+  void ComputeLinkImplementationRuntimeLibraries(
+    const std::string& config, cmOptionalLinkImplementation& impl) const;
+
+  void ComputeLinkInterfaceRuntimeLibraries(
+    const std::string& config, cmOptionalLinkInterface& iface) const;
+
 public:
   const std::vector<const cmGeneratorTarget*>& GetLinkImplementationClosure(
     const std::string& config) const;

Source/cmLinkItem.h

@@ -7,6 +7,7 @@
 #include <map>
 #include <ostream>
 #include <string>
+#include <unordered_map>
 #include <vector>
 
 #include <cmext/algorithm>
@@ -80,6 +81,8 @@ struct cmLinkInterface : public cmLinkInterfaceLibraries
 {
   // Languages whose runtime libraries must be linked.
   std::vector<std::string> Languages;
+  std::unordered_map<std::string, std::vector<cmLinkItem>>
+    LanguageRuntimeLibraries;
 
   // Shared library dependencies needed for linking on some platforms.
   std::vector<cmLinkItem> SharedDeps;
@@ -115,6 +118,8 @@ struct cmLinkImplementation : public cmLinkImplementationLibraries
 {
   // Languages whose runtime libraries must be linked.
   std::vector<std::string> Languages;
+  std::unordered_map<std::string, std::vector<cmLinkImplItem>>
+    LanguageRuntimeLibraries;
 
   // Whether the list depends on a link language genex.
   bool HadLinkLanguageSensitiveCondition = false;

Tests/HIP/CMakeLists.txt

@@ -7,6 +7,9 @@ endmacro ()
 add_hip_test_macro(HIP.ArchitectureOff HIPOnlyArchitectureOff)
 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)
+add_hip_test_macro(HIP.MixedLanguage HIPMixedLanguage)
 add_hip_test_macro(HIP.TryCompile HIPOnlyTryCompile)
 add_hip_test_macro(HIP.WithDefs HIPOnlyWithDefs)

Tests/HIP/InferHipLang1/CMakeLists.txt

@@ -0,0 +1,12 @@
+cmake_minimum_required(VERSION 3.18)
+project(InferHipLang C CXX HIP)
+
+#Goal for this example:
+#make sure that we understand that HIP is the correct link language
+add_library(InterfaceWithHIP INTERFACE)
+target_sources(InterfaceWithHIP INTERFACE interface.hip main.cxx)
+target_compile_features(InterfaceWithHIP INTERFACE hip_std_14)
+target_compile_features(InterfaceWithHIP INTERFACE cxx_std_11)
+
+add_executable(HIPInferHipLang1 )
+target_link_libraries(HIPInferHipLang1 PRIVATE InterfaceWithHIP)

Tests/HIP/InferHipLang1/interface.hip

@@ -0,0 +1,19 @@
+#include <system_error>
+#include <type_traits>
+#include <hip/hip_runtime_api.h>
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ interface_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}

Tests/HIP/InferHipLang1/main.cxx

@@ -0,0 +1,19 @@
+
+#include <iostream>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+int interface_hip_func(int);
+
+int main(int argc, char** argv)
+{
+  interface_hip_func(int(42));
+
+  return 0;
+}

Tests/HIP/InferHipLang2/CMakeLists.txt

@@ -0,0 +1,12 @@
+cmake_minimum_required(VERSION 3.18)
+project(InferHipLang C CXX HIP)
+
+#Goal for this example:
+#make sure that we understand that HIP is the correct link language
+add_library(InterfaceWithHIP OBJECT)
+target_sources(InterfaceWithHIP PRIVATE interface.hip main.cxx)
+target_compile_features(InterfaceWithHIP INTERFACE hip_std_14)
+target_compile_features(InterfaceWithHIP INTERFACE cxx_std_11)
+
+add_executable(HIPInferHipLang2 )
+target_link_libraries(HIPInferHipLang2 PRIVATE InterfaceWithHIP)

Tests/HIP/InferHipLang2/interface.hip

@@ -0,0 +1,19 @@
+#include <system_error>
+#include <type_traits>
+#include <hip/hip_runtime_api.h>
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ interface_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}

Tests/HIP/InferHipLang2/main.cxx

@@ -0,0 +1,19 @@
+
+#include <iostream>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+int interface_hip_func(int);
+
+int main(int argc, char** argv)
+{
+  interface_hip_func(int(42));
+
+  return 0;
+}

Tests/HIP/MixedLanguage/CMakeLists.txt

@@ -0,0 +1,19 @@
+cmake_minimum_required(VERSION 3.18)
+project (MixedLanguage C CXX HIP)
+
+set(CMAKE_HIP_STANDARD 14)
+set(CMAKE_CXX_STANDARD 14)
+
+#Goal for this example:
+#make sure that we can build multiple languages into targets
+#and have the link language always be HIP
+add_library(MixedSharedLib SHARED shared.c)
+add_library(MixedObjectLib OBJECT shared.cxx shared.hip)
+set_target_properties(MixedObjectLib PROPERTIES POSITION_INDEPENDENT_CODE ON)
+target_link_libraries(MixedSharedLib PRIVATE MixedObjectLib)
+
+add_library(MixedStaticLib STATIC static.c static.cxx static.hip)
+set_target_properties(MixedStaticLib PROPERTIES POSITION_INDEPENDENT_CODE ON)
+
+add_executable(HIPMixedLanguage main.cxx)
+target_link_libraries(HIPMixedLanguage PRIVATE MixedStaticLib MixedSharedLib)

Tests/HIP/MixedLanguage/main.cxx

@@ -0,0 +1,40 @@
+
+#include <iostream>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#else
+#  define IMPORT
+#endif
+
+extern "C" {
+IMPORT int shared_c_func(int);
+int static_c_func(int);
+}
+
+IMPORT int shared_cxx_func(int);
+IMPORT int shared_hip_func(int);
+
+int static_cxx_func(int);
+int static_hip_func(int);
+
+int main(int argc, char** argv)
+{
+  static_c_func(int(42));
+  static_cxx_func(int(42));
+  static_hip_func(int(42));
+
+  shared_c_func(int(42));
+  shared_cxx_func(int(42));
+  shared_hip_func(int(42));
+
+  return 0;
+}

Tests/HIP/MixedLanguage/shared.c

@@ -0,0 +1,12 @@
+
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+EXPORT int shared_c_func(int x)
+{
+  return -x;
+}

Tests/HIP/MixedLanguage/shared.cxx

@@ -0,0 +1,21 @@
+
+#include <type_traits>
+
+#ifdef __HIP_PLATFORM_HCC__
+#  error "__HIP_PLATFORM_HCC__ propagated to C++ compilation!"
+#endif
+
+#ifdef __HIP_ROCclr__
+#  error "__HIP_ROCclr__ propagated to C++ compilation!"
+#endif
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+EXPORT int shared_cxx_func(int x)
+{
+  return x * x + std::integral_constant<int, 14>::value;
+}

Tests/HIP/MixedLanguage/shared.hip

@@ -0,0 +1,26 @@
+#include <system_error>
+#include <type_traits>
+#include <hip/hip_runtime_api.h>
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ shared_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}

Tests/HIP/MixedLanguage/static.c

@@ -0,0 +1,6 @@
+
+
+int static_c_func(int x)
+{
+  return -x;
+}

Tests/HIP/MixedLanguage/static.cxx

@@ -0,0 +1,7 @@
+
+#include <type_traits>
+
+int static_cxx_func(int x)
+{
+  return x * x + std::integral_constant<int, 14>::value;
+}

Tests/HIP/MixedLanguage/static.hip

@@ -0,0 +1,21 @@
+
+#include <type_traits>
+#include <system_error>
+#include <hip/hip_runtime_api.h>
+
+
+static __global__ void fake_hip_kernel()
+{
+}
+
+
+int __host__ static_hip_func(int x)
+{
+
+  fake_hip_kernel<<<1, 1>>>();
+  bool valid = (hipSuccess == hipGetLastError());
+  if (!valid) {
+    throw std::system_error(ENODEV, std::generic_category(), "no hip device");
+  }
+  return x * x + std::integral_constant<int, 17>::value;
+}