1 files changed, 85 insertions, 0 deletions
diff --git a/Tests/CudaOnly/ResolveDeviceSymbols/main.cu b/Tests/CudaOnly/ResolveDeviceSymbols/main.cu
new file mode 100644
index 0000000..b4b5b9e
--- /dev/null
+++ b/Tests/CudaOnly/ResolveDeviceSymbols/main.cu
@@ -0,0 +1,85 @@
+
+#include <iostream>
+
+#include "file1.h"
+#include "file2.h"
+
+int file2_launch_kernel(int x);
+
+result_type_dynamic __device__ file2_func(int x);
+static __global__ void main_kernel(result_type_dynamic& r, int x)
+{
+  // call function that was not device linked to us, this will cause
+  // a runtime failure of "invalid device function"
+  r = file2_func(x);
+}
+
+int main_launch_kernel(int x)
+{
+  result_type_dynamic r;
+  main_kernel<<<1, 1>>>(r, x);
+  return r.sum;
+}
+
+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;
+    if (prop.major >= 3) {
+      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 supporting compute >=3.0"
+            << std::endl;
+
+  return 1;
+}
+
+int main(int argc, char** argv)
+{
+  int ret = choose_cuda_device();
+  if (ret) {
+    return 0;
+  }
+
+  cudaError_t err;
+  file2_launch_kernel(42);
+  err = cudaGetLastError();
+  if (err != cudaSuccess) {
+    std::cerr << "file2_launch_kernel: kernel launch failed: "
+              << cudaGetErrorString(err) << std::endl;
+    return 1;
+  }
+
+  main_launch_kernel(1);
+  err = cudaGetLastError();
+  if (err == cudaSuccess) {
+    // This kernel launch should fail as the file2_func was device linked
+    // into the static library and is not usable by the executable
+    std::cerr << "main_launch_kernel: kernel launch should have failed"
+              << std::endl;
+    return 1;
+  }
+
+  return 0;
+}