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.

7 files changed, 221 insertions, 0 deletions

diff --git a/Tests/CudaOnly/RuntimeControls/CMakeLists.txt b/Tests/CudaOnly/RuntimeControls/CMakeLists.txt
new file mode 100644
index 0000000..8b58fec
--- /dev/null
+++ b/Tests/CudaOnly/RuntimeControls/CMakeLists.txt
@@ -0,0 +1,60 @@
+cmake_minimum_required(VERSION 3.7)
+project (RuntimeControls CUDA)
+
+# Find nm and dumpbin
+if(CMAKE_NM)
+  set(dump_command ${CMAKE_NM})
+  set(dump_args -g)
+else()
+  include(GetPrerequisites)
+  message(STATUS "calling list_prerequisites to find dumpbin")
+  list_prerequisites("${CMAKE_COMMAND}" 0 0 0)
+  if(gp_dumpbin)
+    set(dump_command ${gp_dumpbin})
+    set(dump_args /ARCHIVEMEMBERS)
+  endif()
+endif()
+
+string(APPEND CMAKE_CUDA_FLAGS " -gencode arch=compute_30,code=[compute_30]")
+
+set(CMAKE_CUDA_STANDARD 11)
+set(CMAKE_CUDA_RUNTIME_LIBRARY static)
+
+if(NOT "x${CMAKE_CUDA_SIMULATE_ID}" STREQUAL "xMSVC")
+  add_library(UsesNoCudaRT SHARED file1.cu)
+  set_target_properties(UsesNoCudaRT PROPERTIES CUDA_RUNTIME_LIBRARY none)
+endif()
+
+add_library(UsesStaticCudaRT SHARED file2.cu)
+
+add_executable(CudaOnlyRuntimeControls main.cu)
+set_target_properties(CudaOnlyRuntimeControls PROPERTIES CUDA_RUNTIME_LIBRARY shared)
+
+target_link_libraries(CudaOnlyRuntimeControls PRIVATE $<TARGET_NAME_IF_EXISTS:UsesNoCudaRT> UsesStaticCudaRT)
+
+
+if(dump_command)
+  if(TARGET UsesNoCudaRT)
+    add_custom_command(TARGET UsesNoCudaRT POST_BUILD
+    COMMAND ${CMAKE_COMMAND}
+    -DDUMP_COMMAND=${dump_command}
+    -DDUMP_ARGS=${dump_args}
+    -DTEST_LIBRARY_PATH=$<TARGET_FILE:UsesNoCudaRT>
+    -P ${CMAKE_CURRENT_SOURCE_DIR}/no_runtime.cmake
+    )
+  endif()
+  add_custom_command(TARGET UsesStaticCudaRT POST_BUILD
+    COMMAND ${CMAKE_COMMAND}
+    -DDUMP_COMMAND=${dump_command}
+    -DDUMP_ARGS=${dump_args}
+    -DTEST_LIBRARY_PATH=$<TARGET_FILE:UsesStaticCudaRT>
+    -P ${CMAKE_CURRENT_SOURCE_DIR}/uses_static_runtime.cmake
+  )
+  string(REPLACE ";" "|" dirs "${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}")
+  add_custom_command(TARGET CudaOnlyRuntimeControls POST_BUILD
+    COMMAND ${CMAKE_COMMAND}
+    -DEXEC_PATH=$<TARGET_FILE:CudaOnlyRuntimeControls>
+    -DEXTRA_LIB_DIRS="${dirs}"
+    -P ${CMAKE_CURRENT_SOURCE_DIR}/verify_runtime.cmake
+  )
+endif()
diff --git a/Tests/CudaOnly/RuntimeControls/file1.cu b/Tests/CudaOnly/RuntimeControls/file1.cu
new file mode 100644
index 0000000..28beb5e
--- /dev/null
+++ b/Tests/CudaOnly/RuntimeControls/file1.cu
@@ -0,0 +1,18 @@
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+void __global__ file1_kernel(int x, int& r)
+{
+  r = -x;
+}
+
+EXPORT int file1_launch_kernel(int x)
+{
+  int r = 0;
+  file1_kernel<<<1, 1>>>(x, r);
+  return r;
+}
diff --git a/Tests/CudaOnly/RuntimeControls/file2.cu b/Tests/CudaOnly/RuntimeControls/file2.cu
new file mode 100644
index 0000000..ff68a70
--- /dev/null
+++ b/Tests/CudaOnly/RuntimeControls/file2.cu
@@ -0,0 +1,18 @@
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+void __global__ file2_kernel(int x, int& r)
+{
+  r = -x;
+}
+
+EXPORT int file2_launch_kernel(int x)
+{
+  int r = 0;
+  file2_kernel<<<1, 1>>>(x, r);
+  return r;
+}
diff --git a/Tests/CudaOnly/RuntimeControls/main.cu b/Tests/CudaOnly/RuntimeControls/main.cu
new file mode 100644
index 0000000..0be22af
--- /dev/null
+++ b/Tests/CudaOnly/RuntimeControls/main.cu
@@ -0,0 +1,81 @@
+
+#include <iostream>
+
+#include "cuda.h"
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#else
+#  define IMPORT
+#endif
+
+#ifndef _WIN32
+IMPORT int file1_launch_kernel(int x);
+#endif
+
+IMPORT int file2_launch_kernel(int x);
+
+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;
+#ifndef _WIN32
+  file1_launch_kernel(1);
+  err = cudaGetLastError();
+  if (err != cudaSuccess) {
+    std::cerr << "file1_launch_kernel: kernel launch should have passed.\n "
+                 "Error message: "
+              << cudaGetErrorString(err) << std::endl;
+    return 1;
+  }
+#endif
+
+  file2_launch_kernel(1);
+  err = cudaGetLastError();
+  if (err != cudaSuccess) {
+    std::cerr << "file2_launch_kernel: kernel launch should have passed.\n "
+                 "Error message: "
+              << cudaGetErrorString(err) << std::endl;
+    return 1;
+  }
+
+  return 0;
+}
diff --git a/Tests/CudaOnly/RuntimeControls/no_runtime.cmake b/Tests/CudaOnly/RuntimeControls/no_runtime.cmake
new file mode 100644
index 0000000..55f28cc
--- /dev/null
+++ b/Tests/CudaOnly/RuntimeControls/no_runtime.cmake
@@ -0,0 +1,14 @@
+execute_process(COMMAND ${DUMP_COMMAND} ${DUMP_ARGS} ${TEST_LIBRARY_PATH}
+  RESULT_VARIABLE RESULT
+  OUTPUT_VARIABLE OUTPUT
+  ERROR_VARIABLE ERROR
+)
+
+if(NOT "${RESULT}" STREQUAL "0")
+  message(FATAL_ERROR "${DUMP_COMMAND} failed [${RESULT}] [${OUTPUT}] [${ERROR}]")
+endif()
+
+if(NOT "${OUTPUT}" MATCHES "(__cuda)")
+  message(FATAL_ERROR
+  "not missing cuda device symbols, static runtime linking was used.")
+endif()
diff --git a/Tests/CudaOnly/RuntimeControls/uses_static_runtime.cmake b/Tests/CudaOnly/RuntimeControls/uses_static_runtime.cmake
new file mode 100644
index 0000000..b372fea
--- /dev/null
+++ b/Tests/CudaOnly/RuntimeControls/uses_static_runtime.cmake
@@ -0,0 +1,14 @@
+execute_process(COMMAND ${DUMP_COMMAND} ${DUMP_ARGS} ${TEST_LIBRARY_PATH}
+  RESULT_VARIABLE RESULT
+  OUTPUT_VARIABLE OUTPUT
+  ERROR_VARIABLE ERROR
+)
+
+if(NOT "${RESULT}" STREQUAL "0")
+  message(FATAL_ERROR "${DUMP_COMMAND} failed [${RESULT}] [${OUTPUT}] [${ERROR}]")
+endif()
+
+if("${OUTPUT}" MATCHES "__cuda")
+  message(FATAL_ERROR
+    "missing cuda device symbols, static runtime linking was not used.")
+endif()
diff --git a/Tests/CudaOnly/RuntimeControls/verify_runtime.cmake b/Tests/CudaOnly/RuntimeControls/verify_runtime.cmake
new file mode 100644
index 0000000..b313dac
--- /dev/null
+++ b/Tests/CudaOnly/RuntimeControls/verify_runtime.cmake
@@ -0,0 +1,16 @@
+
+string(REPLACE "|" ";" dirs "${EXTRA_LIB_DIRS}")
+file(GET_RUNTIME_DEPENDENCIES
+  RESOLVED_DEPENDENCIES_VAR resolved_libs
+  UNRESOLVED_DEPENDENCIES_VAR unresolved_libs
+  DIRECTORIES ${dirs}
+  EXECUTABLES ${EXEC_PATH}
+  )
+
+list(FILTER resolved_libs INCLUDE REGEX ".*cudart.*")
+list(LENGTH resolved_libs has_cudart)
+
+if(has_cudart EQUAL 0)
+  message(FATAL_ERROR
+    "missing cudart shared library from runtime dependency output.")
+endif()