CUDA: Add a test to verify device linking can handle circular deps

Tests/CudaOnly/CMakeLists.txt

@@ -1,4 +1,5 @@
 
+ADD_TEST_MACRO(CudaOnly.CircularLinkLine CudaOnlyCircularLinkLine)
 ADD_TEST_MACRO(CudaOnly.EnableStandard CudaOnlyEnableStandard)
 ADD_TEST_MACRO(CudaOnly.ExportPTX CudaOnlyExportPTX)
 ADD_TEST_MACRO(CudaOnly.GPUDebugFlag CudaOnlyGPUDebugFlag)

Tests/CudaOnly/CircularLinkLine/CMakeLists.txt

@@ -0,0 +1,35 @@
+cmake_minimum_required(VERSION 3.7)
+project (CudaOnlyCircularLinkLine CUDA)
+
+#Goal for this example:
+# Verify that we de-duplicate the device link line
+# Verify that a de-duplicated link line still works with circular static libraries
+
+string(APPEND CMAKE_CUDA_FLAGS " -gencode arch=compute_30,code=[compute_30]")
+set(CMAKE_CXX_STANDARD 11)
+set(CMAKE_CUDA_STANDARD 11)
+
+add_library(CUDACircularDeviceLinking1 STATIC file1.cu)
+add_library(CUDACircularDeviceLinking2 STATIC file2.cu)
+add_library(CUDACircularDeviceLinking3 STATIC file3.cu)
+add_executable(CudaOnlyCircularLinkLine main.cu)
+
+target_link_libraries(CUDACircularDeviceLinking1 PUBLIC CUDACircularDeviceLinking2)
+target_link_libraries(CUDACircularDeviceLinking2 PUBLIC CUDACircularDeviceLinking3)
+#FIXME: complete the loop once supported
+#target_link_libraries(CUDACircularDeviceLinking3 PUBLIC CUDACircularDeviceLinking1)
+
+target_link_libraries(CudaOnlyCircularLinkLine PRIVATE CUDACircularDeviceLinking3)
+
+
+set_target_properties(CUDACircularDeviceLinking1
+                      PROPERTIES
+                      CUDA_SEPARABLE_COMPILATION ON)
+
+set_target_properties(CUDACircularDeviceLinking2
+                      PROPERTIES
+                      CUDA_SEPARABLE_COMPILATION ON)
+
+set_target_properties(CUDACircularDeviceLinking3
+                      PROPERTIES
+                      CUDA_SEPARABLE_COMPILATION ON)

Tests/CudaOnly/CircularLinkLine/file1.cu

@@ -0,0 +1,6 @@
+
+extern __device__ int file2_func(int);
+int __device__ file1_func(int x)
+{
+  return file2_func(x);
+}

Tests/CudaOnly/CircularLinkLine/file2.cu

@@ -0,0 +1,6 @@
+
+extern __device__ int file3_func(int);
+int __device__ file2_func(int x)
+{
+  return x + file3_func(x);
+}

Tests/CudaOnly/CircularLinkLine/file3.cu

@@ -0,0 +1,8 @@
+
+extern __device__ int file1_func(int);
+int __device__ file3_func(int x)
+{
+  if (x > 0)
+    return file1_func(-x);
+  return x;
+}

Tests/CudaOnly/CircularLinkLine/main.cu

@@ -0,0 +1,5 @@
+
+int main(int argc, char** argv)
+{
+  return 0;
+}