9 files changed, 171 insertions, 11 deletions
diff --git a/Tests/CudaOnly/CMakeLists.txt b/Tests/CudaOnly/CMakeLists.txt
index aa4755d30..39634ac87 100644
--- a/Tests/CudaOnly/CMakeLists.txt
+++ b/Tests/CudaOnly/CMakeLists.txt
@@ -7,13 +7,13 @@ endmacro ()
 add_cuda_test_macro(CudaOnly.Architecture Architecture)
 add_cuda_test_macro(CudaOnly.ArchSpecial CudaOnlyArchSpecial)
 add_cuda_test_macro(CudaOnly.CompileFlags CudaOnlyCompileFlags)
-
 add_cuda_test_macro(CudaOnly.EnableStandard CudaOnlyEnableStandard)
 add_cuda_test_macro(CudaOnly.ExportPTX CudaOnlyExportPTX)
 add_cuda_test_macro(CudaOnly.SharedRuntimePlusToolkit CudaOnlySharedRuntimePlusToolkit)
 add_cuda_test_macro(CudaOnly.Standard98 CudaOnlyStandard98)
 add_cuda_test_macro(CudaOnly.Toolkit CudaOnlyToolkit)
 add_cuda_test_macro(CudaOnly.ToolkitBeforeLang CudaOnlyToolkitBeforeLang)
+add_cuda_test_macro(CudaOnly.ToolkitMultipleDirs CudaOnlyToolkitMultipleDirs)
 add_cuda_test_macro(CudaOnly.WithDefs CudaOnlyWithDefs)
 add_cuda_test_macro(CudaOnly.CircularLinkLine CudaOnlyCircularLinkLine)
 add_cuda_test_macro(CudaOnly.ResolveDeviceSymbols CudaOnlyResolveDeviceSymbols)
@@ -28,6 +28,19 @@ if(CMake_TEST_CUDA AND NOT CMake_TEST_CUDA STREQUAL "Clang")
   add_cuda_test_macro(CudaOnly.GPUDebugFlag CudaOnlyGPUDebugFlag)
 endif()
 
+# The CUDA only ships the shared version of the toolkit libraries
+# on windows
+if(NOT WIN32)
+  add_cuda_test_macro(CudaOnly.StaticRuntimePlusToolkit CudaOnlyStaticRuntimePlusToolkit)
+endif()
+
+add_cuda_test_macro(CudaOnly.DeviceLTO CudaOnlyDeviceLTO)
+
+if(MSVC)
+  # Tests for features that only work with MSVC
+  add_cuda_test_macro(CudaOnly.PDB CudaOnlyPDB)
+endif()
+
 add_test(NAME CudaOnly.DontResolveDeviceSymbols COMMAND
   ${CMAKE_CTEST_COMMAND} -C $<CONFIGURATION>
   --build-and-test
@@ -41,16 +54,6 @@ add_test(NAME CudaOnly.DontResolveDeviceSymbols COMMAND
 set_property(TEST "CudaOnly.DontResolveDeviceSymbols" APPEND
   PROPERTY LABELS "CUDA")
 
-# The CUDA only ships the shared version of the toolkit libraries
-# on windows
-if(NOT WIN32)
-  add_cuda_test_macro(CudaOnly.StaticRuntimePlusToolkit CudaOnlyStaticRuntimePlusToolkit)
-endif()
-
-if(MSVC)
-  add_cuda_test_macro(CudaOnly.PDB CudaOnlyPDB)
-endif()
-
 add_test(NAME CudaOnly.RuntimeControls COMMAND
   ${CMAKE_CTEST_COMMAND} -C $<CONFIGURATION>
   --build-and-test
diff --git a/Tests/CudaOnly/DeviceLTO/CMakeLists.txt b/Tests/CudaOnly/DeviceLTO/CMakeLists.txt
new file mode 100644
index 000000000..653b35db7
--- /dev/null
+++ b/Tests/CudaOnly/DeviceLTO/CMakeLists.txt
@@ -0,0 +1,37 @@
+cmake_minimum_required(VERSION 3.18)
+project(DeviceLTO CUDA)
+
+# Goal:
+# Verify that we correctly compile with device LTO
+# Verify that device LTO requirements are propagated to
+# the final device link line
+
+add_library(CUDA_dlto STATIC file1.cu file2.cu file3.cu)
+add_executable(CudaOnlyDeviceLTO main.cu)
+
+set_target_properties(CUDA_dlto
+                      PROPERTIES
+                      CUDA_ARCHITECTURES "${CMAKE_CUDA_ARCHITECTURES_ALL}"
+                      CUDA_SEPARABLE_COMPILATION ON
+                      POSITION_INDEPENDENT_CODE ON)
+
+set_target_properties(CudaOnlyDeviceLTO
+                      PROPERTIES
+                      CUDA_SEPARABLE_COMPILATION ON
+                      CUDA_ARCHITECTURES "${CMAKE_CUDA_ARCHITECTURES_ALL}"
+                      )
+
+target_link_libraries(CudaOnlyDeviceLTO PRIVATE CUDA_dlto)
+
+include(CheckIPOSupported)
+check_ipo_supported(LANGUAGES CUDA RESULT ipo_supported)
+if(ipo_supported)
+  set_target_properties(CUDA_dlto
+                        PROPERTIES
+                        INTERPROCEDURAL_OPTIMIZATION ON)
+
+  # When non-LTO variants (i.e. virtual) are built together with LTO ones the
+  # linker warns about missing device LTO for the virtual architectures.
+  # Ignore these warnings.
+  target_link_options(CudaOnlyDeviceLTO PRIVATE "$<DEVICE_LINK:-w>")
+endif()
diff --git a/Tests/CudaOnly/DeviceLTO/file1.cu b/Tests/CudaOnly/DeviceLTO/file1.cu
new file mode 100644
index 000000000..703927cfd
--- /dev/null
+++ b/Tests/CudaOnly/DeviceLTO/file1.cu
@@ -0,0 +1,17 @@
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+extern __device__ int file2_func(int);
+void __global__ kernel(int x)
+{
+  file2_func(x);
+}
+
+EXPORT int launch_kernel(int x)
+{
+  kernel<<<1, 1>>>(x);
+  return x;
+}
diff --git a/Tests/CudaOnly/DeviceLTO/file2.cu b/Tests/CudaOnly/DeviceLTO/file2.cu
new file mode 100644
index 000000000..73d646819
--- /dev/null
+++ b/Tests/CudaOnly/DeviceLTO/file2.cu
@@ -0,0 +1,5 @@
+extern __device__ int file3_func(int);
+int __device__ file2_func(int x)
+{
+  return x + file3_func(x);
+}
diff --git a/Tests/CudaOnly/DeviceLTO/file3.cu b/Tests/CudaOnly/DeviceLTO/file3.cu
new file mode 100644
index 000000000..235ac0665
--- /dev/null
+++ b/Tests/CudaOnly/DeviceLTO/file3.cu
@@ -0,0 +1,4 @@
+int __device__ file3_func(int x)
+{
+  return x * x * x;
+}
diff --git a/Tests/CudaOnly/DeviceLTO/main.cu b/Tests/CudaOnly/DeviceLTO/main.cu
new file mode 100644
index 000000000..8ef487338
--- /dev/null
+++ b/Tests/CudaOnly/DeviceLTO/main.cu
@@ -0,0 +1,62 @@
+#include <iostream>
+
+#include "cuda.h"
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#else
+#  define IMPORT
+#endif
+
+IMPORT int launch_kernel(int x);
+
+int choose_cuda_device()
+{
+  int nDevices = 0;
+  cudaError_t err = cudaGetDeviceCount(&nDevices);
+  if (err != cudaSuccess) {
+    std::cerr << "Failed to retrieve the number of CUDA enabled devices"
+              << std::endl;
+    return 1;
+  }
+  for (int i = 0; i < nDevices; ++i) {
+    cudaDeviceProp prop;
+    cudaError_t err = cudaGetDeviceProperties(&prop, i);
+    if (err != cudaSuccess) {
+      std::cerr << "Could not retrieve properties from CUDA device " << i
+                << std::endl;
+      return 1;
+    }
+    std::cout << "prop.major: " << prop.major << std::endl;
+    err = cudaSetDevice(i);
+    if (err != cudaSuccess) {
+      std::cout << "Could not select CUDA device " << i << std::endl;
+    } else {
+      return 0;
+    }
+  }
+
+  std::cout << "Could not find a CUDA enabled card" << std::endl;
+
+  return 1;
+}
+
+int main()
+{
+  int ret = choose_cuda_device();
+  if (ret) {
+    return 0;
+  }
+
+  cudaError_t err;
+  launch_kernel(1);
+  err = cudaGetLastError();
+  if (err != cudaSuccess) {
+    std::cerr << "launch_kernel: kernel launch should have passed.\n "
+                 "Error message: "
+              << cudaGetErrorString(err) << std::endl;
+    return 1;
+  }
+
+  return 0;
+}
diff --git a/Tests/CudaOnly/ToolkitMultipleDirs/CMakeLists.txt b/Tests/CudaOnly/ToolkitMultipleDirs/CMakeLists.txt
new file mode 100644
index 000000000..92551499d
--- /dev/null
+++ b/Tests/CudaOnly/ToolkitMultipleDirs/CMakeLists.txt
@@ -0,0 +1,22 @@
+cmake_minimum_required(VERSION 3.15)
+project(CudaOnlyToolkitMultipleDirs CUDA)
+
+find_package(CUDAToolkit REQUIRED)
+
+add_subdirectory(subdir)
+
+if(NOT DEFINED CUDAToolkit_VERSION)
+  message(FATAL_ERROR "expected CUDAToolkit variable CUDAToolkit_VERSION not found")
+endif()
+
+message(STATUS "CUDAToolkit_VERSION: ${CUDAToolkit_VERSION}")
+message(STATUS "CUDAToolkit_VERSION_MAJOR: ${CUDAToolkit_VERSION_MAJOR}")
+message(STATUS "CUDAToolkit_VERSION_MINOR: ${CUDAToolkit_VERSION_MINOR}")
+message(STATUS "CUDAToolkit_VERSION_PATCH: ${CUDAToolkit_VERSION_PATCH}")
+message(STATUS "CUDAToolkit_BIN_DIR: ${CUDAToolkit_BIN_DIR}")
+message(STATUS "CUDAToolkit_INCLUDE_DIRS: ${CUDAToolkit_INCLUDE_DIRS}")
+message(STATUS "CUDAToolkit_LIBRARY_DIR: ${CUDAToolkit_LIBRARY_DIR}")
+message(STATUS "CUDAToolkit_NVCC_EXECUTABLE ${CUDAToolkit_NVCC_EXECUTABLE}")
+
+add_executable(CudaOnlyToolkitMultipleDirs main.cu)
+target_link_libraries(CudaOnlyToolkitMultipleDirs PRIVATE CUDA::toolkit)
diff --git a/Tests/CudaOnly/ToolkitMultipleDirs/main.cu b/Tests/CudaOnly/ToolkitMultipleDirs/main.cu
new file mode 100644
index 000000000..0f3ccdcc8
--- /dev/null
+++ b/Tests/CudaOnly/ToolkitMultipleDirs/main.cu
@@ -0,0 +1,8 @@
+// Only thing we care about is that these headers are found
+#include <cuda.h>
+#include <cuda_runtime_api.h>
+
+int main(int argc, char** argv)
+{
+  return 0;
+}
diff --git a/Tests/CudaOnly/ToolkitMultipleDirs/subdir/CMakeLists.txt b/Tests/CudaOnly/ToolkitMultipleDirs/subdir/CMakeLists.txt
new file mode 100644
index 000000000..f1f4b8925
--- /dev/null
+++ b/Tests/CudaOnly/ToolkitMultipleDirs/subdir/CMakeLists.txt
@@ -0,0 +1,2 @@
+
+find_package(CUDAToolkit REQUIRED)