summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorRobert Maynard <robert.maynard@kitware.com>2020-09-02 18:53:41 (GMT)
committerZack Galbreath <zack.galbreath@kitware.com>2021-06-07 19:25:33 (GMT)
commit947dbed0aa502c70a9e165757450bd8409036980 (patch)
tree0a45b31986184368c7ad6df72d9f9186c7a24b4c
parentb50bfc89131e55685aa41146254b37d26049b4d5 (diff)
downloadCMake-947dbed0aa502c70a9e165757450bd8409036980.zip
CMake-947dbed0aa502c70a9e165757450bd8409036980.tar.gz
CMake-947dbed0aa502c70a9e165757450bd8409036980.tar.bz2
HIP: Automatically inject the `hip::device` runtime target
Any target that might need to link to hip code needs the `hip::device` target
-rw-r--r--Modules/CMakeHIPCompiler.cmake.in5
-rw-r--r--Modules/CMakeHIPInformation.cmake7
-rw-r--r--Modules/CMakeHIPRuntime.cmake.in99
-rw-r--r--Modules/CMakeTestHIPCompiler.cmake21
-rw-r--r--Modules/Compiler/Clang-HIP.cmake4
-rw-r--r--Modules/Compiler/ROCMClang-HIP.cmake4
-rw-r--r--Source/cmComputeLinkDepends.cxx14
-rw-r--r--Source/cmGeneratorTarget.cxx196
-rw-r--r--Source/cmGeneratorTarget.h6
-rw-r--r--Source/cmLinkItem.h5
-rw-r--r--Tests/HIP/CMakeLists.txt3
-rw-r--r--Tests/HIP/InferHipLang1/CMakeLists.txt12
-rw-r--r--Tests/HIP/InferHipLang1/interface.hip19
-rw-r--r--Tests/HIP/InferHipLang1/main.cxx19
-rw-r--r--Tests/HIP/InferHipLang2/CMakeLists.txt12
-rw-r--r--Tests/HIP/InferHipLang2/interface.hip19
-rw-r--r--Tests/HIP/InferHipLang2/main.cxx19
-rw-r--r--Tests/HIP/MixedLanguage/CMakeLists.txt19
-rw-r--r--Tests/HIP/MixedLanguage/main.cxx40
-rw-r--r--Tests/HIP/MixedLanguage/shared.c12
-rw-r--r--Tests/HIP/MixedLanguage/shared.cxx21
-rw-r--r--Tests/HIP/MixedLanguage/shared.hip26
-rw-r--r--Tests/HIP/MixedLanguage/static.c6
-rw-r--r--Tests/HIP/MixedLanguage/static.cxx7
-rw-r--r--Tests/HIP/MixedLanguage/static.hip21
25 files changed, 587 insertions, 29 deletions
diff --git a/Modules/CMakeHIPCompiler.cmake.in b/Modules/CMakeHIPCompiler.cmake.in
index eeb51b2..9a30a45 100644
--- a/Modules/CMakeHIPCompiler.cmake.in
+++ b/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@")
diff --git a/Modules/CMakeHIPInformation.cmake b/Modules/CMakeHIPInformation.cmake
index df1d98a..ec37e1c 100644
--- a/Modules/CMakeHIPInformation.cmake
+++ b/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()
diff --git a/Modules/CMakeHIPRuntime.cmake.in b/Modules/CMakeHIPRuntime.cmake.in
new file mode 100644
index 0000000..ade26bb
--- /dev/null
+++ b/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()
diff --git a/Modules/CMakeTestHIPCompiler.cmake b/Modules/CMakeTestHIPCompiler.cmake
index da70f8d..d9fcc9d 100644
--- a/Modules/CMakeTestHIPCompiler.cmake
+++ b/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()
diff --git a/Modules/Compiler/Clang-HIP.cmake b/Modules/Compiler/Clang-HIP.cmake
index 6ac39de..1030a43 100644
--- a/Modules/Compiler/Clang-HIP.cmake
+++ b/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 "")
diff --git a/Modules/Compiler/ROCMClang-HIP.cmake b/Modules/Compiler/ROCMClang-HIP.cmake
index 9e792b5..7af7699 100644
--- a/Modules/Compiler/ROCMClang-HIP.cmake
+++ b/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)
diff --git a/Source/cmComputeLinkDepends.cxx b/Source/cmComputeLinkDepends.cxx
index 4a6518fd..0b27e34 100644
--- a/Source/cmComputeLinkDepends.cxx
+++ b/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);
}
diff --git a/Source/cmGeneratorTarget.cxx b/Source/cmGeneratorTarget.cxx
index 812535b..cb10117 100644
--- a/Source/cmGeneratorTarget.cxx
+++ b/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;
}
diff --git a/Source/cmGeneratorTarget.h b/Source/cmGeneratorTarget.h
index ab9ccfd..ed66fb1 100644
--- a/Source/cmGeneratorTarget.h
+++ b/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;
diff --git a/Source/cmLinkItem.h b/Source/cmLinkItem.h
index db44938..0863edd 100644
--- a/Source/cmLinkItem.h
+++ b/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;
diff --git a/Tests/HIP/CMakeLists.txt b/Tests/HIP/CMakeLists.txt
index b3966c9..9499be8 100644
--- a/Tests/HIP/CMakeLists.txt
+++ b/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)
diff --git a/Tests/HIP/InferHipLang1/CMakeLists.txt b/Tests/HIP/InferHipLang1/CMakeLists.txt
new file mode 100644
index 0000000..63d77fd
--- /dev/null
+++ b/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)
diff --git a/Tests/HIP/InferHipLang1/interface.hip b/Tests/HIP/InferHipLang1/interface.hip
new file mode 100644
index 0000000..6ac8641
--- /dev/null
+++ b/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;
+}
diff --git a/Tests/HIP/InferHipLang1/main.cxx b/Tests/HIP/InferHipLang1/main.cxx
new file mode 100644
index 0000000..987b6c6
--- /dev/null
+++ b/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;
+}
diff --git a/Tests/HIP/InferHipLang2/CMakeLists.txt b/Tests/HIP/InferHipLang2/CMakeLists.txt
new file mode 100644
index 0000000..0e69de3
--- /dev/null
+++ b/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)
diff --git a/Tests/HIP/InferHipLang2/interface.hip b/Tests/HIP/InferHipLang2/interface.hip
new file mode 100644
index 0000000..6ac8641
--- /dev/null
+++ b/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;
+}
diff --git a/Tests/HIP/InferHipLang2/main.cxx b/Tests/HIP/InferHipLang2/main.cxx
new file mode 100644
index 0000000..987b6c6
--- /dev/null
+++ b/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;
+}
diff --git a/Tests/HIP/MixedLanguage/CMakeLists.txt b/Tests/HIP/MixedLanguage/CMakeLists.txt
new file mode 100644
index 0000000..4f6dd3b
--- /dev/null
+++ b/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)
diff --git a/Tests/HIP/MixedLanguage/main.cxx b/Tests/HIP/MixedLanguage/main.cxx
new file mode 100644
index 0000000..003d18a
--- /dev/null
+++ b/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;
+}
diff --git a/Tests/HIP/MixedLanguage/shared.c b/Tests/HIP/MixedLanguage/shared.c
new file mode 100644
index 0000000..347fba9
--- /dev/null
+++ b/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;
+}
diff --git a/Tests/HIP/MixedLanguage/shared.cxx b/Tests/HIP/MixedLanguage/shared.cxx
new file mode 100644
index 0000000..8e6c1d3
--- /dev/null
+++ b/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;
+}
diff --git a/Tests/HIP/MixedLanguage/shared.hip b/Tests/HIP/MixedLanguage/shared.hip
new file mode 100644
index 0000000..e6fea9f
--- /dev/null
+++ b/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;
+}
diff --git a/Tests/HIP/MixedLanguage/static.c b/Tests/HIP/MixedLanguage/static.c
new file mode 100644
index 0000000..06c33b4
--- /dev/null
+++ b/Tests/HIP/MixedLanguage/static.c
@@ -0,0 +1,6 @@
+
+
+int static_c_func(int x)
+{
+ return -x;
+}
diff --git a/Tests/HIP/MixedLanguage/static.cxx b/Tests/HIP/MixedLanguage/static.cxx
new file mode 100644
index 0000000..2c14fb1
--- /dev/null
+++ b/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;
+}
diff --git a/Tests/HIP/MixedLanguage/static.hip b/Tests/HIP/MixedLanguage/static.hip
new file mode 100644
index 0000000..359b9fa
--- /dev/null
+++ b/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;
+}