From 84f3c87b030037646c1110b1e4237f597ad3dd88 Mon Sep 17 00:00:00 2001
From: Pierre Moreau <dev@pmoreau.org>
Date: Thu, 2 Feb 2017 13:40:42 +0100
Subject: Tests/Cuda: Print error message if an error occurred

Fixes c59811a2 "CUDA: Tests now state why they are failing when no CUDA card is found."
---
 Tests/Cuda/Complex/dynamic.cu | 2 +-
 Tests/Cuda/Complex/file3.cu   | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/Tests/Cuda/Complex/dynamic.cu b/Tests/Cuda/Complex/dynamic.cu
index a23dc25..b0a27d4 100644
--- a/Tests/Cuda/Complex/dynamic.cu
+++ b/Tests/Cuda/Complex/dynamic.cu
@@ -26,7 +26,7 @@ EXPORT void cuda_dynamic_lib_func()
 {
   DetermineIfValidCudaDevice <<<1,1>>> ();
   cudaError_t err = cudaGetLastError();
-  if(err == cudaSuccess)
+  if(err != cudaSuccess)
     {
     std::cerr << cudaGetErrorString(err) << std::endl;
     }
diff --git a/Tests/Cuda/Complex/file3.cu b/Tests/Cuda/Complex/file3.cu
index 47e64c5..0c8a09b 100644
--- a/Tests/Cuda/Complex/file3.cu
+++ b/Tests/Cuda/Complex/file3.cu
@@ -20,7 +20,7 @@ int file3_launch_kernel(int x)
   result_type r;
   file3_kernel <<<1,1>>> (r,x);
   cudaError_t err = cudaGetLastError();
-  if(err == cudaSuccess)
+  if(err != cudaSuccess)
     {
     std::cerr << cudaGetErrorString(err) << std::endl;
     return x;
-- 
cgit v0.12


From eebb2be8b0db17bb5f760e7fa04c020406e6f6eb Mon Sep 17 00:00:00 2001
From: Pierre Moreau <dev@pmoreau.org>
Date: Thu, 2 Feb 2017 23:22:16 +0100
Subject: Tests/Cuda: Add identifiers to error messages

---
 Tests/Cuda/Complex/dynamic.cu | 3 ++-
 Tests/Cuda/Complex/file3.cu   | 3 ++-
 2 files changed, 4 insertions(+), 2 deletions(-)

diff --git a/Tests/Cuda/Complex/dynamic.cu b/Tests/Cuda/Complex/dynamic.cu
index b0a27d4..fc22c8b 100644
--- a/Tests/Cuda/Complex/dynamic.cu
+++ b/Tests/Cuda/Complex/dynamic.cu
@@ -28,6 +28,7 @@ EXPORT void cuda_dynamic_lib_func()
   cudaError_t err = cudaGetLastError();
   if(err != cudaSuccess)
     {
-    std::cerr << cudaGetErrorString(err) << std::endl;
+    std::cerr << "DetermineIfValidCudaDevice [SYNC] failed: "
+              << cudaGetErrorString(err) << std::endl;
     }
 }
diff --git a/Tests/Cuda/Complex/file3.cu b/Tests/Cuda/Complex/file3.cu
index 0c8a09b..1daf47e 100644
--- a/Tests/Cuda/Complex/file3.cu
+++ b/Tests/Cuda/Complex/file3.cu
@@ -22,7 +22,8 @@ int file3_launch_kernel(int x)
   cudaError_t err = cudaGetLastError();
   if(err != cudaSuccess)
     {
-    std::cerr << cudaGetErrorString(err) << std::endl;
+    std::cerr << "file3_kernel [SYNC] failed: "
+              << cudaGetErrorString(err) << std::endl;
     return x;
     }
   return r.sum;
-- 
cgit v0.12


From 21a125cdbff24efa746249a3b8ccc6296c5f0aff Mon Sep 17 00:00:00 2001
From: Pierre Moreau <dev@pmoreau.org>
Date: Thu, 2 Feb 2017 13:44:49 +0100
Subject: Tests/Cuda: Print error message if mixed_kernel failed

---
 Tests/Cuda/Complex/mixed.cu | 7 +++++++
 1 file changed, 7 insertions(+)

diff --git a/Tests/Cuda/Complex/mixed.cu b/Tests/Cuda/Complex/mixed.cu
index 7051de0..429f1f3 100644
--- a/Tests/Cuda/Complex/mixed.cu
+++ b/Tests/Cuda/Complex/mixed.cu
@@ -31,5 +31,12 @@ EXPORT int mixed_launch_kernel(int x)
 
   result_type r;
   mixed_kernel <<<1,1>>> (r,x);
+  cudaError_t err = cudaGetLastError();
+  if(err != cudaSuccess)
+    {
+    std::cerr << "mixed_kernel [SYNC] failed: "
+              << cudaGetErrorString(err) << std::endl;
+    return x;
+    }
   return r.sum;
 }
-- 
cgit v0.12


From c0d7bb8368c3d157d1a2758b620fc726355e554d Mon Sep 17 00:00:00 2001
From: Pierre Moreau <dev@pmoreau.org>
Date: Thu, 2 Feb 2017 13:58:09 +0100
Subject: Tests/Cuda: Print asynchronous error messages, if any

As kernel launches are asynchronous, a `cudaGetLastError()` right after
the kernel launch might be executed while the kernel is still running.
Synchronizing the device will ensure that all the work is completed
before progressing further on, and allows to catch errors that were
previously missed.
The `cudaGetLastError()` after the `cudaDeviceSynchronize()` is there
to reset the error variable to `cudaSuccess`.
---
 Tests/Cuda/Complex/dynamic.cu | 6 ++++++
 Tests/Cuda/Complex/file3.cu   | 7 +++++++
 Tests/Cuda/Complex/mixed.cu   | 7 +++++++
 3 files changed, 20 insertions(+)

diff --git a/Tests/Cuda/Complex/dynamic.cu b/Tests/Cuda/Complex/dynamic.cu
index fc22c8b..0206bb4 100644
--- a/Tests/Cuda/Complex/dynamic.cu
+++ b/Tests/Cuda/Complex/dynamic.cu
@@ -31,4 +31,10 @@ EXPORT void cuda_dynamic_lib_func()
     std::cerr << "DetermineIfValidCudaDevice [SYNC] failed: "
               << cudaGetErrorString(err) << std::endl;
     }
+  err = cudaDeviceSynchronize();
+  if(err != cudaSuccess)
+    {
+    std::cerr << "DetermineIfValidCudaDevice [ASYNC] failed: "
+              << cudaGetErrorString(cudaGetLastError()) << std::endl;
+    }
 }
diff --git a/Tests/Cuda/Complex/file3.cu b/Tests/Cuda/Complex/file3.cu
index 1daf47e..912105a 100644
--- a/Tests/Cuda/Complex/file3.cu
+++ b/Tests/Cuda/Complex/file3.cu
@@ -26,5 +26,12 @@ int file3_launch_kernel(int x)
               << cudaGetErrorString(err) << std::endl;
     return x;
     }
+  err = cudaDeviceSynchronize();
+  if(err != cudaSuccess)
+    {
+    std::cerr << "file3_kernel [ASYNC] failed: "
+              << cudaGetErrorString(cudaGetLastError()) << std::endl;
+    return x;
+    }
   return r.sum;
 }
diff --git a/Tests/Cuda/Complex/mixed.cu b/Tests/Cuda/Complex/mixed.cu
index 429f1f3..a7bcd4e 100644
--- a/Tests/Cuda/Complex/mixed.cu
+++ b/Tests/Cuda/Complex/mixed.cu
@@ -38,5 +38,12 @@ EXPORT int mixed_launch_kernel(int x)
               << cudaGetErrorString(err) << std::endl;
     return x;
     }
+  err = cudaDeviceSynchronize();
+  if(err != cudaSuccess)
+    {
+    std::cerr << "mixed_kernel [ASYNC] failed: "
+              << cudaGetErrorString(cudaGetLastError()) << std::endl;
+    return x;
+    }
   return r.sum;
 }
-- 
cgit v0.12


From 008ed80dcf1d03640879b4168f4fba956aa03196 Mon Sep 17 00:00:00 2001
From: Pierre Moreau <dev@pmoreau.org>
Date: Thu, 2 Feb 2017 22:52:53 +0100
Subject: Tests/Cuda: Output error messages to std::cerr instead of std::cout

---
 Tests/Cuda/ObjectLibrary/static.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/Tests/Cuda/ObjectLibrary/static.cu b/Tests/Cuda/ObjectLibrary/static.cu
index 2374c23..a801153 100644
--- a/Tests/Cuda/ObjectLibrary/static.cu
+++ b/Tests/Cuda/ObjectLibrary/static.cu
@@ -10,8 +10,8 @@ int __host__ file1_sq_func(int x)
   err = cudaGetDeviceCount(&nDevices);
   if(err != cudaSuccess)
   {
-    std::cout << "nDevices: " << nDevices << std::endl;
-    std::cout << "err: " << err << std::endl;
+    std::cerr << "nDevices: " << nDevices << std::endl;
+    std::cerr << "err: " << err << std::endl;
     return 1;
   }
   std::cout << "this library uses cuda code" << std::endl;
-- 
cgit v0.12


From ce19607fed3990b8e828330e77f09061c99aa113 Mon Sep 17 00:00:00 2001
From: Pierre Moreau <dev@pmoreau.org>
Date: Thu, 2 Feb 2017 22:38:32 +0100
Subject: Tests/Cuda: Fix missing CUDA static library at runtime on macOS

Suggested-by: Robert Maynard
---
 Tests/Cuda/Complex/CMakeLists.txt | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/Tests/Cuda/Complex/CMakeLists.txt b/Tests/Cuda/Complex/CMakeLists.txt
index 9a3703a..14454cf 100644
--- a/Tests/Cuda/Complex/CMakeLists.txt
+++ b/Tests/Cuda/Complex/CMakeLists.txt
@@ -38,3 +38,9 @@ target_link_libraries(CudaComplexMixedLib
 
 add_executable(CudaComplex main.cpp)
 target_link_libraries(CudaComplex PUBLIC CudaComplexMixedLib)
+
+if(APPLE)
+  # We need to add the default path to the driver (libcuda.dylib) as an rpath, so that
+  # the static cuda runtime can find it at runtime.
+  target_link_libraries(CudaComplex PRIVATE -Wl,-rpath,/usr/local/cuda/lib)
+endif()
-- 
cgit v0.12


From 0ae5386aa953d1670074c2f1bfc9a04ddb382684 Mon Sep 17 00:00:00 2001
From: Pierre Moreau <dev@pmoreau.org>
Date: Thu, 2 Feb 2017 22:41:14 +0100
Subject: Tests/Cuda: Add missing separable compilation property

This resulted in `mixed_kernel()` returning an "invalid device function"
at runtime for `file1_func()`.

Suggested-by: Robert Maynard
---
 Tests/Cuda/Complex/CMakeLists.txt | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/Tests/Cuda/Complex/CMakeLists.txt b/Tests/Cuda/Complex/CMakeLists.txt
index 14454cf..2dd6f73 100644
--- a/Tests/Cuda/Complex/CMakeLists.txt
+++ b/Tests/Cuda/Complex/CMakeLists.txt
@@ -32,6 +32,8 @@ add_library(CudaComplexSharedLib SHARED dynamic.cu)
 target_link_libraries(CudaComplexSharedLib PUBLIC CudaComplexCppBase)
 
 add_library(CudaComplexMixedLib SHARED mixed.cpp mixed.cu)
+set_target_properties(CudaComplexMixedLib
+                       PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
 target_link_libraries(CudaComplexMixedLib
                       PUBLIC CudaComplexSharedLib
                       PRIVATE CudaComplexSeperableLib)
-- 
cgit v0.12


From 8731701cb636317df2691359361562f32adfe759 Mon Sep 17 00:00:00 2001
From: Pierre Moreau <dev@pmoreau.org>
Date: Thu, 2 Feb 2017 22:54:24 +0100
Subject: Tests/Cuda: Use memory allocated on the GPU in the kernels

---
 Tests/Cuda/Complex/file3.cu | 27 ++++++++++++++++++++++-----
 Tests/Cuda/Complex/mixed.cu | 28 +++++++++++++++++++++++-----
 2 files changed, 45 insertions(+), 10 deletions(-)

diff --git a/Tests/Cuda/Complex/file3.cu b/Tests/Cuda/Complex/file3.cu
index 912105a..bd8198d 100644
--- a/Tests/Cuda/Complex/file3.cu
+++ b/Tests/Cuda/Complex/file3.cu
@@ -9,17 +9,25 @@ result_type_dynamic __device__ file2_func(int x);
 
 static
 __global__
-void file3_kernel(result_type& r, int x)
+void file3_kernel(result_type* r, int x)
 {
-  r = file1_func(x);
+  *r = file1_func(x);
   result_type_dynamic rd = file2_func(x);
 }
 
 int file3_launch_kernel(int x)
 {
-  result_type r;
+  result_type* r;
+  cudaError_t err = cudaMallocManaged(&r, sizeof(result_type));
+  if(err != cudaSuccess)
+    {
+    std::cerr << "file3_launch_kernel: cudaMallocManaged failed: "
+              << cudaGetErrorString(err) << std::endl;
+    return x;
+    }
+
   file3_kernel <<<1,1>>> (r,x);
-  cudaError_t err = cudaGetLastError();
+  err = cudaGetLastError();
   if(err != cudaSuccess)
     {
     std::cerr << "file3_kernel [SYNC] failed: "
@@ -33,5 +41,14 @@ int file3_launch_kernel(int x)
               << cudaGetErrorString(cudaGetLastError()) << std::endl;
     return x;
     }
-  return r.sum;
+  int result = r->sum;
+  err = cudaFree(r);
+  if(err != cudaSuccess)
+    {
+    std::cerr << "file3_launch_kernel: cudaFree failed: "
+              << cudaGetErrorString(err) << std::endl;
+    return x;
+    }
+
+  return result;
 }
diff --git a/Tests/Cuda/Complex/mixed.cu b/Tests/Cuda/Complex/mixed.cu
index a7bcd4e..d96cc7c 100644
--- a/Tests/Cuda/Complex/mixed.cu
+++ b/Tests/Cuda/Complex/mixed.cu
@@ -19,9 +19,9 @@ IMPORT void __host__ cuda_dynamic_lib_func();
 
 static
 __global__
-void mixed_kernel(result_type& r, int x)
+void mixed_kernel(result_type* r, int x)
 {
-  r = file1_func(x);
+  *r = file1_func(x);
   result_type_dynamic rd = file2_func(x);
 }
 
@@ -29,9 +29,17 @@ EXPORT int mixed_launch_kernel(int x)
 {
   cuda_dynamic_lib_func();
 
-  result_type r;
+  result_type* r;
+  cudaError_t err = cudaMallocManaged(&r, sizeof(result_type));
+  if(err != cudaSuccess)
+    {
+    std::cerr << "mixed_launch_kernel: cudaMallocManaged failed: "
+              << cudaGetErrorString(err) << std::endl;
+    return x;
+    }
+
   mixed_kernel <<<1,1>>> (r,x);
-  cudaError_t err = cudaGetLastError();
+  err = cudaGetLastError();
   if(err != cudaSuccess)
     {
     std::cerr << "mixed_kernel [SYNC] failed: "
@@ -45,5 +53,15 @@ EXPORT int mixed_launch_kernel(int x)
               << cudaGetErrorString(cudaGetLastError()) << std::endl;
     return x;
     }
-  return r.sum;
+
+  int result = r->sum;
+  err = cudaFree(r);
+  if(err != cudaSuccess)
+    {
+    std::cerr << "mixed_launch_kernel: cudaFree failed: "
+              << cudaGetErrorString(err) << std::endl;
+    return x;
+    }
+
+  return result;
 }
-- 
cgit v0.12


From cbe4d5957b526f9a9bb6954d0c6944ad7642bb90 Mon Sep 17 00:00:00 2001
From: Pierre Moreau <dev@pmoreau.org>
Date: Thu, 2 Feb 2017 23:12:15 +0100
Subject: Tests/Cuda: Return a non-zero code if errors occurred

---
 Tests/Cuda/Complex/main.cpp       | 6 +++---
 Tests/Cuda/ObjectLibrary/main.cpp | 8 +++++---
 2 files changed, 8 insertions(+), 6 deletions(-)

diff --git a/Tests/Cuda/Complex/main.cpp b/Tests/Cuda/Complex/main.cpp
index 5a3f820..92d1fb0 100644
--- a/Tests/Cuda/Complex/main.cpp
+++ b/Tests/Cuda/Complex/main.cpp
@@ -14,7 +14,7 @@ IMPORT int mixed_launch_kernel(int x);
 
 int main(int argc, char** argv)
 {
-  call_cuda_seperable_code(42);
-  mixed_launch_kernel(42);
-  return 0;
+  int r1 = call_cuda_seperable_code(42);
+  int r2 = mixed_launch_kernel(42);
+  return (r1 == 42 || r2 == 42) ? 1 : 0;
 }
diff --git a/Tests/Cuda/ObjectLibrary/main.cpp b/Tests/Cuda/ObjectLibrary/main.cpp
index 1a70a99..4d2f890 100644
--- a/Tests/Cuda/ObjectLibrary/main.cpp
+++ b/Tests/Cuda/ObjectLibrary/main.cpp
@@ -4,14 +4,16 @@
 int static_func(int);
 int file1_sq_func(int);
 
-void test_functions()
+int test_functions()
 {
-  file1_sq_func(static_func(42));
+  return file1_sq_func(static_func(42));
 }
 
 int main(int argc, char** argv)
 {
-  test_functions();
+  if (test_functions() == 1) {
+    return 1;
+  }
   std::cout
     << "this executable doesn't use cuda code, just call methods defined"
     << std::endl;
-- 
cgit v0.12


From 166b3393875e4f6752907b9b06c2518dce0432e9 Mon Sep 17 00:00:00 2001
From: Pierre Moreau <dev@pmoreau.org>
Date: Wed, 8 Feb 2017 22:50:58 +0100
Subject: Tests/Cuda: Select a CUDA device supporting compute 3.0

---
 Tests/Cuda/Complex/dynamic.cu | 40 ++++++++++++++++++++++++++++++++++++++++
 Tests/Cuda/Complex/main.cpp   |  6 ++++++
 2 files changed, 46 insertions(+)

diff --git a/Tests/Cuda/Complex/dynamic.cu b/Tests/Cuda/Complex/dynamic.cu
index 0206bb4..2b04ac9 100644
--- a/Tests/Cuda/Complex/dynamic.cu
+++ b/Tests/Cuda/Complex/dynamic.cu
@@ -22,6 +22,46 @@ void DetermineIfValidCudaDevice()
 {
 }
 
+EXPORT 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;
+      }
+    if (prop.major >= 4)
+      {
+        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;
+}
+
 EXPORT void cuda_dynamic_lib_func()
 {
   DetermineIfValidCudaDevice <<<1,1>>> ();
diff --git a/Tests/Cuda/Complex/main.cpp b/Tests/Cuda/Complex/main.cpp
index 92d1fb0..2498235 100644
--- a/Tests/Cuda/Complex/main.cpp
+++ b/Tests/Cuda/Complex/main.cpp
@@ -9,11 +9,17 @@
 #define IMPORT
 #endif
 
+IMPORT int choose_cuda_device();
 IMPORT int call_cuda_seperable_code(int x);
 IMPORT int mixed_launch_kernel(int x);
 
 int main(int argc, char** argv)
 {
+  int ret = choose_cuda_device();
+  if (ret) {
+    return 0;
+  }
+
   int r1 = call_cuda_seperable_code(42);
   int r2 = mixed_launch_kernel(42);
   return (r1 == 42 || r2 == 42) ? 1 : 0;
-- 
cgit v0.12