From aa8facefe8a451918f9be143c2adc1fffe162c4a Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Fri, 22 Sep 2023 11:13:25 -0400 Subject: CUDA: Visual Studio propagate objects to device linking When given objects via `target_link_libraries(consumer PRIVATE producer)` the VisualStudio solution adds the objects under as `` entries in the solution. This works for host side linking but isn't handled by the cuda msbuild extensions. So to work around this we manually add the objects as additional link items. --- Source/cmVisualStudio10TargetGenerator.cxx | 47 +++++++++++++++++++++--------- 1 file changed, 33 insertions(+), 14 deletions(-) diff --git a/Source/cmVisualStudio10TargetGenerator.cxx b/Source/cmVisualStudio10TargetGenerator.cxx index ea4bd06..2a54a55 100644 --- a/Source/cmVisualStudio10TargetGenerator.cxx +++ b/Source/cmVisualStudio10TargetGenerator.cxx @@ -3864,22 +3864,41 @@ bool cmVisualStudio10TargetGenerator::ComputeCudaLinkOptions( } cudaLinkOptions.AppendFlagString("AdditionalOptions", linkFlags); - // For static libraries that have device linking enabled compute - // the libraries - if (this->GeneratorTarget->GetType() == cmStateEnums::STATIC_LIBRARY && - doDeviceLinking) { - cmComputeLinkInformation& cli = *pcli; - cmLinkLineDeviceComputer computer( - this->LocalGenerator, - this->LocalGenerator->GetStateSnapshot().GetDirectory()); - std::vector> btLibVec; - computer.ComputeLinkLibraries(cli, std::string{}, btLibVec); + if (doDeviceLinking) { std::vector libVec; - for (auto const& item : btLibVec) { - libVec.emplace_back(item.Value); + auto const& kinded = this->GeneratorTarget->GetKindedSources(configName); + // CMake conversion uses full paths when possible to allow deeper trees. + // However, CUDA 8.0 msbuild rules fail on absolute paths so for CUDA + // we must use relative paths. + const bool forceRelative = true; + for (cmGeneratorTarget::SourceAndKind const& si : kinded.Sources) { + switch (si.Kind) { + case cmGeneratorTarget::SourceKindExternalObject: { + std::string path = + this->ConvertPath(si.Source.Value->GetFullPath(), forceRelative); + ConvertToWindowsSlash(path); + libVec.emplace_back(std::move(path)); + } break; + default: + break; + } + } + // For static libraries that have device linking enabled compute + // the libraries + if (this->GeneratorTarget->GetType() == cmStateEnums::STATIC_LIBRARY) { + cmComputeLinkInformation& cli = *pcli; + cmLinkLineDeviceComputer computer( + this->LocalGenerator, + this->LocalGenerator->GetStateSnapshot().GetDirectory()); + std::vector> btLibVec; + computer.ComputeLinkLibraries(cli, std::string{}, btLibVec); + for (auto const& item : btLibVec) { + libVec.emplace_back(item.Value); + } + } + if (!libVec.empty()) { + cudaLinkOptions.AddFlag("AdditionalDependencies", libVec); } - - cudaLinkOptions.AddFlag("AdditionalDependencies", libVec); } this->CudaLinkOptions[configName] = std::move(pOptions); -- cgit v0.12 From cd984261e1c2f27c2c716d43f3502c829990685d Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Fri, 22 Sep 2023 11:13:13 -0400 Subject: CUDA: Device linking now uses TARGET_OBJECTS content Due to an oversight in cmLinkLineDeviceComputer object files did not get propagate to the device linking phase when given via the `$` generator expression. --- Source/cmLinkLineDeviceComputer.cxx | 16 ++++++++-------- Tests/CudaOnly/CMakeLists.txt | 1 + .../SeparateCompilationTargetObjects/CMakeLists.txt | 18 ++++++++++++++++++ Tests/CudaOnly/SeparateCompilationTargetObjects/bar.cu | 18 ++++++++++++++++++ Tests/CudaOnly/SeparateCompilationTargetObjects/foo.cu | 18 ++++++++++++++++++ .../CudaOnly/SeparateCompilationTargetObjects/main.cu | 16 ++++++++++++++++ 6 files changed, 79 insertions(+), 8 deletions(-) create mode 100644 Tests/CudaOnly/SeparateCompilationTargetObjects/CMakeLists.txt create mode 100644 Tests/CudaOnly/SeparateCompilationTargetObjects/bar.cu create mode 100644 Tests/CudaOnly/SeparateCompilationTargetObjects/foo.cu create mode 100644 Tests/CudaOnly/SeparateCompilationTargetObjects/main.cu diff --git a/Source/cmLinkLineDeviceComputer.cxx b/Source/cmLinkLineDeviceComputer.cxx index ded6466..28aa5d9 100644 --- a/Source/cmLinkLineDeviceComputer.cxx +++ b/Source/cmLinkLineDeviceComputer.cxx @@ -101,9 +101,7 @@ void cmLinkLineDeviceComputer::ComputeLinkLibraries( ItemVector const& items = cli.GetItems(); std::string config = cli.GetConfig(); bool skipItemAfterFramework = false; - // Note: - // Any modification of this algorithm should be reflected also in - // cmVisualStudio10TargetGenerator::ComputeCudaLinkOptions + for (auto const& item : items) { if (skipItemAfterFramework) { skipItemAfterFramework = false; @@ -132,11 +130,13 @@ void cmLinkLineDeviceComputer::ComputeLinkLibraries( BT linkLib; if (item.IsPath == cmComputeLinkInformation::ItemIsPath::Yes) { - // nvcc understands absolute paths to libraries ending in '.a' or '.lib'. - // These should be passed to nvlink. Other extensions need to be left - // out because nvlink may not understand or need them. Even though it - // can tolerate '.so' or '.dylib' it cannot tolerate '.so.1'. - if (cmHasLiteralSuffix(item.Value.Value, ".a") || + // nvcc understands absolute paths to libraries ending in '.o', .a', or + // '.lib'. These should be passed to nvlink. Other extensions need to be + // left out because nvlink may not understand or need them. Even though + // it can tolerate '.so' or '.dylib' it cannot tolerate '.so.1'. + if (cmHasLiteralSuffix(item.Value.Value, ".o") || + cmHasLiteralSuffix(item.Value.Value, ".obj") || + cmHasLiteralSuffix(item.Value.Value, ".a") || cmHasLiteralSuffix(item.Value.Value, ".lib")) { linkLib.Value = item .GetFormattedItem(this->ConvertToOutputFormat( diff --git a/Tests/CudaOnly/CMakeLists.txt b/Tests/CudaOnly/CMakeLists.txt index aa25c4c..b7ce5a1 100644 --- a/Tests/CudaOnly/CMakeLists.txt +++ b/Tests/CudaOnly/CMakeLists.txt @@ -20,6 +20,7 @@ add_cuda_test_macro(CudaOnly.CircularLinkLine CudaOnlyCircularLinkLine) add_cuda_test_macro(CudaOnly.ResolveDeviceSymbols CudaOnlyResolveDeviceSymbols) add_cuda_test_macro(CudaOnly.SeparateCompilation main/CudaOnlySeparateCompilation) add_cuda_test_macro(CudaOnly.SeparateCompilationPTX CudaOnlySeparateCompilationPTX) +add_cuda_test_macro(CudaOnly.SeparateCompilationTargetObjects CudaOnlySeparateCompilationTargetObjects) if(CMake_TEST_CUDA AND NOT CMake_TEST_CUDA STREQUAL "Clang") # Clang doesn't have flags for selecting the runtime. diff --git a/Tests/CudaOnly/SeparateCompilationTargetObjects/CMakeLists.txt b/Tests/CudaOnly/SeparateCompilationTargetObjects/CMakeLists.txt new file mode 100644 index 0000000..7dbc0d5 --- /dev/null +++ b/Tests/CudaOnly/SeparateCompilationTargetObjects/CMakeLists.txt @@ -0,0 +1,18 @@ +cmake_minimum_required(VERSION 3.25.5) + +project(SeparateCompilationObjects LANGUAGES CUDA) + +add_library(foo OBJECT foo.cu) +set_target_properties(foo PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + +add_library(bar OBJECT bar.cu) +set_target_properties(bar PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + +add_executable(CudaOnlySeparateCompilationTargetObjects main.cu) +set_target_properties(CudaOnlySeparateCompilationTargetObjects PROPERTIES CUDA_SEPARABLE_COMPILATION ON) +target_link_libraries(CudaOnlySeparateCompilationTargetObjects PRIVATE $ bar) + +if(APPLE) + # Help the static cuda runtime find the driver (libcuda.dyllib) at runtime. + set_property(TARGET CudaOnlySeparateCompilationTargetObjects PROPERTY BUILD_RPATH ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}) +endif() diff --git a/Tests/CudaOnly/SeparateCompilationTargetObjects/bar.cu b/Tests/CudaOnly/SeparateCompilationTargetObjects/bar.cu new file mode 100644 index 0000000..234586f --- /dev/null +++ b/Tests/CudaOnly/SeparateCompilationTargetObjects/bar.cu @@ -0,0 +1,18 @@ + +#include + +#ifdef _WIN32 +# define EXPORT __declspec(dllexport) +#else +# define EXPORT __attribute__((__visibility__("default"))) +#endif + +__global__ void b1() +{ +} + +EXPORT int bar() +{ + b1<<<1, 1>>>(); + return 0; +} diff --git a/Tests/CudaOnly/SeparateCompilationTargetObjects/foo.cu b/Tests/CudaOnly/SeparateCompilationTargetObjects/foo.cu new file mode 100644 index 0000000..75c04af --- /dev/null +++ b/Tests/CudaOnly/SeparateCompilationTargetObjects/foo.cu @@ -0,0 +1,18 @@ + +#include + +#ifdef _WIN32 +# define EXPORT __declspec(dllexport) +#else +# define EXPORT __attribute__((__visibility__("default"))) +#endif + +__global__ void k1() +{ +} + +EXPORT int foo() +{ + k1<<<1, 1>>>(); + return 0; +} diff --git a/Tests/CudaOnly/SeparateCompilationTargetObjects/main.cu b/Tests/CudaOnly/SeparateCompilationTargetObjects/main.cu new file mode 100644 index 0000000..78b10b1 --- /dev/null +++ b/Tests/CudaOnly/SeparateCompilationTargetObjects/main.cu @@ -0,0 +1,16 @@ +// main.cu +#include + +#ifdef _WIN32 +# define IMPORT __declspec(dllimport) +#else +# define IMPORT +#endif + +IMPORT int foo(); +IMPORT int bar(); + +int main(int argc, char**) +{ + return foo() && bar(); +} -- cgit v0.12