11 files changed, 199 insertions, 0 deletions
diff --git a/Tests/Cuda/Complex/CMakeLists.txt b/Tests/Cuda/Complex/CMakeLists.txt
new file mode 100644
index 0000000..9a3703a
--- /dev/null
+++ b/Tests/Cuda/Complex/CMakeLists.txt
@@ -0,0 +1,40 @@
+
+cmake_minimum_required(VERSION 3.7)
+project (CudaComplex CXX CUDA)
+#Goal for this example:
+
+#build a cpp dynamic library base
+#build a cuda static library base that uses separable compilation
+
+#build a cuda dynamic library that uses the first dynamic library
+#build a mixed cpp & cuda dynamic library uses all 3 previous libraries
+
+#lastly build a cpp executable that uses this last cuda dynamic library
+
+#this tests that we can properly handle linking cuda and cpp together
+#and also bulding cpp targets that need cuda implicit libraries
+
+#verify that we can pass explicit cuda arch flags
+set(CMAKE_CUDA_FLAGS "-gencode arch=compute_30,code=compute_30")
+set(CMAKE_CUDA_STANDARD 11)
+set(CMAKE_CXX_STANDARD 11)
+set(CMAKE_CUDA_STANDARD_REQUIRED TRUE)
+set(CMAKE_CXX_STANDARD_REQUIRED TRUE)
+
+add_library(CudaComplexCppBase SHARED dynamic.cpp)
+add_library(CudaComplexSeperableLib STATIC file1.cu file2.cu file3.cu)
+set_target_properties(CudaComplexSeperableLib
+                       PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+set_target_properties( CudaComplexSeperableLib
+                       PROPERTIES POSITION_INDEPENDENT_CODE ON)
+
+add_library(CudaComplexSharedLib SHARED dynamic.cu)
+target_link_libraries(CudaComplexSharedLib PUBLIC CudaComplexCppBase)
+
+add_library(CudaComplexMixedLib SHARED mixed.cpp mixed.cu)
+target_link_libraries(CudaComplexMixedLib
+                      PUBLIC CudaComplexSharedLib
+                      PRIVATE CudaComplexSeperableLib)
+
+add_executable(CudaComplex main.cpp)
+target_link_libraries(CudaComplex PUBLIC CudaComplexMixedLib)
diff --git a/Tests/Cuda/Complex/dynamic.cpp b/Tests/Cuda/Complex/dynamic.cpp
new file mode 100644
index 0000000..d579f1e
--- /dev/null
+++ b/Tests/Cuda/Complex/dynamic.cpp
@@ -0,0 +1,5 @@
+
+int dynamic_base_func(int x)
+{
+  return x * x;
+}
diff --git a/Tests/Cuda/Complex/dynamic.cu b/Tests/Cuda/Complex/dynamic.cu
new file mode 100644
index 0000000..9540e86
--- /dev/null
+++ b/Tests/Cuda/Complex/dynamic.cu
@@ -0,0 +1,29 @@
+
+#include <string>
+#include <cuda.h>
+
+int dynamic_base_func(int);
+
+int __host__ cuda_dynamic_host_func(int x)
+{
+  return dynamic_base_func(x);
+}
+
+static
+__global__
+void DetermineIfValidCudaDevice()
+{
+}
+
+void cuda_dynamic_lib_func(std::string& contents )
+{
+  DetermineIfValidCudaDevice <<<1,1>>> ();
+  if(cudaSuccess == cudaGetLastError())
+    {
+    contents = "ran a cuda kernel";
+    }
+  else
+    {
+    contents = "cant run a cuda kernel";
+    }
+}
diff --git a/Tests/Cuda/Complex/file1.cu b/Tests/Cuda/Complex/file1.cu
new file mode 100644
index 0000000..a2e8bf3
--- /dev/null
+++ b/Tests/Cuda/Complex/file1.cu
@@ -0,0 +1,10 @@
+
+#include "file1.h"
+
+result_type __device__ file1_func(int x)
+{
+  result_type r;
+  r.input = x;
+  r.sum = x*x;
+  return r;
+}
diff --git a/Tests/Cuda/Complex/file1.h b/Tests/Cuda/Complex/file1.h
new file mode 100644
index 0000000..ff1945c
--- /dev/null
+++ b/Tests/Cuda/Complex/file1.h
@@ -0,0 +1,7 @@
+
+#pragma once
+struct result_type
+{
+  int input;
+  int sum;
+};
diff --git a/Tests/Cuda/Complex/file2.cu b/Tests/Cuda/Complex/file2.cu
new file mode 100644
index 0000000..6b8b06b
--- /dev/null
+++ b/Tests/Cuda/Complex/file2.cu
@@ -0,0 +1,20 @@
+
+#include "file2.h"
+
+result_type __device__ file1_func(int x);
+
+result_type_dynamic __device__ file2_func(int x)
+{
+  if(x!=42)
+    {
+    const result_type r = file1_func(x);
+    const result_type_dynamic rd { r.input, r.sum, true };
+    return rd;
+    }
+  else
+    {
+    const result_type_dynamic rd { x, x*x*x, false };
+    return rd;
+    }
+
+}
diff --git a/Tests/Cuda/Complex/file2.h b/Tests/Cuda/Complex/file2.h
new file mode 100644
index 0000000..d2dbaa4
--- /dev/null
+++ b/Tests/Cuda/Complex/file2.h
@@ -0,0 +1,10 @@
+
+#pragma once
+#include "file1.h"
+
+struct result_type_dynamic
+{
+  int input;
+  int sum;
+  bool from_static;
+};
diff --git a/Tests/Cuda/Complex/file3.cu b/Tests/Cuda/Complex/file3.cu
new file mode 100644
index 0000000..3c5e952
--- /dev/null
+++ b/Tests/Cuda/Complex/file3.cu
@@ -0,0 +1,25 @@
+
+#include <iostream>
+
+#include "file1.h"
+#include "file2.h"
+
+result_type __device__ file1_func(int x);
+result_type_dynamic __device__ file2_func(int x);
+
+static
+__global__
+void file3_kernel(result_type& r, int x)
+{
+  //call static_func which is a method that is defined in the
+  //static library that is always out of date
+  r = file1_func(x);
+  result_type_dynamic rd = file2_func(x);
+}
+
+int file3_launch_kernel(int x)
+{
+  result_type r;
+  file3_kernel <<<1,1>>> (r,x);
+  return r.sum;
+}
diff --git a/Tests/Cuda/Complex/main.cpp b/Tests/Cuda/Complex/main.cpp
new file mode 100644
index 0000000..a72ffd7
--- /dev/null
+++ b/Tests/Cuda/Complex/main.cpp
@@ -0,0 +1,14 @@
+#include <iostream>
+
+#include "file1.h"
+#include "file2.h"
+
+result_type call_cuda_seperable_code(int x);
+result_type mixed_launch_kernel(int x);
+
+int main(int argc, char** argv)
+{
+  call_cuda_seperable_code(42);
+  mixed_launch_kernel(42);
+  return 0;
+}
diff --git a/Tests/Cuda/Complex/mixed.cpp b/Tests/Cuda/Complex/mixed.cpp
new file mode 100644
index 0000000..205f091
--- /dev/null
+++ b/Tests/Cuda/Complex/mixed.cpp
@@ -0,0 +1,14 @@
+
+int dynamic_base_func(int);
+int cuda_dynamic_host_func(int);
+int file3_launch_kernel(int);
+
+int dynamic_final_func(int x)
+{
+  return cuda_dynamic_host_func(dynamic_base_func(x));
+}
+
+int call_cuda_seperable_code(int x)
+{
+  return file3_launch_kernel(x);
+}
diff --git a/Tests/Cuda/Complex/mixed.cu b/Tests/Cuda/Complex/mixed.cu
new file mode 100644
index 0000000..d2e8275
--- /dev/null
+++ b/Tests/Cuda/Complex/mixed.cu
@@ -0,0 +1,25 @@
+
+#include <iostream>
+
+#include "file1.h"
+#include "file2.h"
+
+result_type __device__ file1_func(int x);
+result_type_dynamic __device__ file2_func(int x);
+
+static
+__global__
+void mixed_kernel(result_type& r, int x)
+{
+  //call static_func which is a method that is defined in the
+  //static library that is always out of date
+  r = file1_func(x);
+  result_type_dynamic rd = file2_func(x);
+}
+
+int mixed_launch_kernel(int x)
+{
+  result_type r;
+  mixed_kernel <<<1,1>>> (r,x);
+  return r.sum;
+}