CUDA: Add abstraction for cuda runtime selection

Fixes #17559
Replace our hard-coded default of cudart=static with a first-class abstraction to select the runtime library from an enumeration of logical names.

4 files changed, 33 insertions, 14 deletions

diff --git a/Tests/Cuda/Complex/CMakeLists.txt b/Tests/Cuda/Complex/CMakeLists.txt
index d3d4b7c..08d1e16 100644
--- a/Tests/Cuda/Complex/CMakeLists.txt
+++ b/Tests/Cuda/Complex/CMakeLists.txt
@@ -22,18 +22,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(CudaComplexSeperableLib STATIC file1.cu file2.cu file3.cu)
 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)
@@ -41,7 +34,27 @@ target_link_libraries(CudaComplexMixedLib
 add_executable(CudaComplex main.cpp)
 target_link_libraries(CudaComplex PUBLIC CudaComplexMixedLib)
 
+
+set_target_properties(CudaComplexMixedLib
+                      CudaComplexSeperableLib
+                      PROPERTIES
+                        POSITION_INDEPENDENT_CODE  ON
+                        CUDA_SEPARABLE_COMPILATION ON
+                      )
+set_target_properties(CudaComplexMixedLib
+                      CudaComplexSharedLib
+                      PROPERTIES
+                        CUDA_RUNTIME_LIBRARY shared
+                      )
+
+
 if(APPLE)
   # Help the static cuda runtime find the driver (libcuda.dyllib) at runtime.
   set_property(TARGET CudaComplex PROPERTY BUILD_RPATH ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
 endif()
+
+if(UNIX)
+  # Help the shared cuda runtime find libcudart as it is not located
+  # in a default system searched location
+  set_property(TARGET CudaComplexMixedLib PROPERTY BUILD_RPATH ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
+endif()
diff --git a/Tests/Cuda/Complex/dynamic.cu b/Tests/Cuda/Complex/dynamic.cu
index 9da8853..7f2f2b5 100644
--- a/Tests/Cuda/Complex/dynamic.cu
+++ b/Tests/Cuda/Complex/dynamic.cu
@@ -54,17 +54,20 @@ EXPORT int choose_cuda_device()
   return 1;
 }
 
-EXPORT void cuda_dynamic_lib_func()
+EXPORT bool cuda_dynamic_lib_func()
 {
-  DetermineIfValidCudaDevice<<<1, 1>>>();
   cudaError_t err = cudaGetLastError();
   if (err != cudaSuccess) {
-    std::cerr << "DetermineIfValidCudaDevice [SYNC] failed: "
+    std::cerr << "DetermineIfValidCudaDevice [Per Launch] failed: "
               << cudaGetErrorString(err) << std::endl;
+    return false;
   }
+  DetermineIfValidCudaDevice<<<1, 1>>>();
   err = cudaDeviceSynchronize();
   if (err != cudaSuccess) {
-    std::cerr << "DetermineIfValidCudaDevice [ASYNC] failed: "
+    std::cerr << "DetermineIfValidCudaDevice [SYNC] failed: "
               << cudaGetErrorString(cudaGetLastError()) << std::endl;
+    return false;
   }
+  return true;
 }
diff --git a/Tests/Cuda/Complex/main.cpp b/Tests/Cuda/Complex/main.cpp
index 6ca5952..da09b44 100644
--- a/Tests/Cuda/Complex/main.cpp
+++ b/Tests/Cuda/Complex/main.cpp
@@ -22,5 +22,6 @@ int main(int argc, char** argv)
 
   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/Complex/mixed.cu b/Tests/Cuda/Complex/mixed.cu
index 5b85aec..76119ad 100644
--- a/Tests/Cuda/Complex/mixed.cu
+++ b/Tests/Cuda/Complex/mixed.cu
@@ -15,7 +15,7 @@
 result_type __device__ file1_func(int x);
 result_type_dynamic __device__ file2_func(int x);
 
-IMPORT void __host__ cuda_dynamic_lib_func();
+IMPORT bool __host__ cuda_dynamic_lib_func();
 
 static __global__ void mixed_kernel(result_type* r, int x)
 {
@@ -25,7 +25,9 @@ static __global__ void mixed_kernel(result_type* r, int x)
 
 EXPORT int mixed_launch_kernel(int x)
 {
-  cuda_dynamic_lib_func();
+  if (!cuda_dynamic_lib_func()) {
+    return x;
+  }
 
   result_type* r;
   cudaError_t err = cudaMallocManaged(&r, sizeof(result_type));