CUDA: Add tests to verify CUDA compiler works properly.

9 files changed, 170 insertions, 0 deletions

diff --git a/Tests/CudaOnly/SeparateCompilation/CMakeLists.txt b/Tests/CudaOnly/SeparateCompilation/CMakeLists.txt
new file mode 100644
index 0000000..7055eef
--- /dev/null
+++ b/Tests/CudaOnly/SeparateCompilation/CMakeLists.txt
@@ -0,0 +1,33 @@
+
+cmake_minimum_required(VERSION 3.7)
+project (CudaOnlySeparateCompilation CUDA)
+
+#Goal for this example:
+#Build a static library that defines multiple methods and kernels that
+#use each other.
+#After that confirm that we can call those methods from dynamic libraries
+#and executables.
+#We complicate the matter by also testing that multiple static libraries
+#all containing cuda separable compilation code links properly
+set(CMAKE_CUDA_FLAGS "-gencode arch=compute_30,code=compute_30")
+set(CMAKE_CXX_STANDARD 11)
+set(CMAKE_CUDA_STANDARD 11)
+add_library(CUDASerarateLibA STATIC file1.cu file2.cu file3.cu)
+
+#Having file4/file5 in a shared library causes serious problems
+#with the nvcc linker and it will generate bad entries that will
+#cause a segv when trying to run the executable
+#
+add_library(CUDASerarateLibB STATIC file4.cu file5.cu)
+target_link_libraries(CUDASerarateLibB PRIVATE CUDASerarateLibA)
+
+add_executable(CudaOnlySeparateCompilation main.cu)
+target_link_libraries(CudaOnlySeparateCompilation PRIVATE CUDASerarateLibB)
+
+set_target_properties( CUDASerarateLibA
+                       CUDASerarateLibB
+                       PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+
+set_target_properties( CUDASerarateLibA
+                       CUDASerarateLibB
+                       PROPERTIES POSITION_INDEPENDENT_CODE ON)
diff --git a/Tests/CudaOnly/SeparateCompilation/file1.cu b/Tests/CudaOnly/SeparateCompilation/file1.cu
new file mode 100644
index 0000000..a2e8bf3
--- /dev/null
+++ b/Tests/CudaOnly/SeparateCompilation/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/CudaOnly/SeparateCompilation/file1.h b/Tests/CudaOnly/SeparateCompilation/file1.h
new file mode 100644
index 0000000..ff1945c
--- /dev/null
+++ b/Tests/CudaOnly/SeparateCompilation/file1.h
@@ -0,0 +1,7 @@
+
+#pragma once
+struct result_type
+{
+  int input;
+  int sum;
+};
diff --git a/Tests/CudaOnly/SeparateCompilation/file2.cu b/Tests/CudaOnly/SeparateCompilation/file2.cu
new file mode 100644
index 0000000..6b8b06b
--- /dev/null
+++ b/Tests/CudaOnly/SeparateCompilation/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/CudaOnly/SeparateCompilation/file2.h b/Tests/CudaOnly/SeparateCompilation/file2.h
new file mode 100644
index 0000000..d2dbaa4
--- /dev/null
+++ b/Tests/CudaOnly/SeparateCompilation/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/CudaOnly/SeparateCompilation/file3.cu b/Tests/CudaOnly/SeparateCompilation/file3.cu
new file mode 100644
index 0000000..670a18b
--- /dev/null
+++ b/Tests/CudaOnly/SeparateCompilation/file3.cu
@@ -0,0 +1,25 @@
+
+
+#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);
+}
+
+result_type file3_launch_kernel(int x)
+{
+  result_type r;
+  file3_kernel <<<1,1>>> (r,x);
+  return r;
+}
diff --git a/Tests/CudaOnly/SeparateCompilation/file4.cu b/Tests/CudaOnly/SeparateCompilation/file4.cu
new file mode 100644
index 0000000..86ef623
--- /dev/null
+++ b/Tests/CudaOnly/SeparateCompilation/file4.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 file4_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 file4_launch_kernel(int x)
+{
+  result_type r;
+  file4_kernel <<<1,1>>> (r,x);
+  return r.sum;
+}
diff --git a/Tests/CudaOnly/SeparateCompilation/file5.cu b/Tests/CudaOnly/SeparateCompilation/file5.cu
new file mode 100644
index 0000000..6fdb32a
--- /dev/null
+++ b/Tests/CudaOnly/SeparateCompilation/file5.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 file5_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 file5_launch_kernel(int x)
+{
+  result_type r;
+  file5_kernel <<<1,1>>> (r,x);
+  return r.sum;
+}
diff --git a/Tests/CudaOnly/SeparateCompilation/main.cu b/Tests/CudaOnly/SeparateCompilation/main.cu
new file mode 100644
index 0000000..d4520ae
--- /dev/null
+++ b/Tests/CudaOnly/SeparateCompilation/main.cu
@@ -0,0 +1,15 @@
+
+#include <iostream>
+
+#include "file1.h"
+#include "file2.h"
+
+// result_type file4_launch_kernel(int x);
+// result_type file5_launch_kernel(int x);
+
+int main(int argc, char **argv)
+{
+  // file4_launch_kernel(42);
+  // file5_launch_kernel(42);
+  return 0;
+}