diff options
25 files changed, 587 insertions, 29 deletions
diff --git a/Modules/CMakeHIPCompiler.cmake.in b/Modules/CMakeHIPCompiler.cmake.in index eeb51b281a..9a30a45b3a 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 df1d98ae41..ec37e1c91a 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 0000000000..ade26bb9b1 --- /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 da70f8d271..d9fcc9d8fe 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 6ac39def09..1030a43cda 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 9e792b5059..7af769987f 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 4a6518fde4..0b27e34c2b 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 812535b099..cb10117704 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 ab9ccfd58d..ed66fb1a81 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 db44938a37..0863edd558 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 b3966c9ebf..9499be8f6c 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 0000000000..63d77fdaa4 --- /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 0000000000..6ac8641cca --- /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 0000000000..987b6c6516 --- /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 0000000000..0e69de38fe --- /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 0000000000..6ac8641cca --- /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 0000000000..987b6c6516 --- /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 0000000000..4f6dd3b0fb --- /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 0000000000..003d18ac4c --- /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 0000000000..347fba9255 --- /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 0000000000..8e6c1d35f8 --- /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 0000000000..e6fea9fe21 --- /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 0000000000..06c33b4cc7 --- /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 0000000000..2c14fb1018 --- /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 0000000000..359b9facda --- /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; +} |