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 `$<TARGET_OBJECTS>` generator expression.

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<std::string> 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(

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.

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 $<TARGET_OBJECTS:foo> 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()

Tests/CudaOnly/SeparateCompilationTargetObjects/bar.cu

@@ -0,0 +1,18 @@
+
+#include <iostream>
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT __attribute__((__visibility__("default")))
+#endif
+
+__global__ void b1()
+{
+}
+
+EXPORT int bar()
+{
+  b1<<<1, 1>>>();
+  return 0;
+}

Tests/CudaOnly/SeparateCompilationTargetObjects/foo.cu

@@ -0,0 +1,18 @@
+
+#include <iostream>
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT __attribute__((__visibility__("default")))
+#endif
+
+__global__ void k1()
+{
+}
+
+EXPORT int foo()
+{
+  k1<<<1, 1>>>();
+  return 0;
+}

Tests/CudaOnly/SeparateCompilationTargetObjects/main.cu

@@ -0,0 +1,16 @@
+// main.cu
+#include <iostream>
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#else
+#  define IMPORT
+#endif
+
+IMPORT int foo();
+IMPORT int bar();
+
+int main(int argc, char**)
+{
+  return foo() && bar();
+}