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.

22 files changed, 722 insertions, 3 deletions

diff --git a/Tests/CudaOnly/CMakeLists.txt b/Tests/CudaOnly/CMakeLists.txt
index a0575cd..cc1ee1a 100644
--- a/Tests/CudaOnly/CMakeLists.txt
+++ b/Tests/CudaOnly/CMakeLists.txt
@@ -5,10 +5,21 @@ ADD_TEST_MACRO(CudaOnly.ExportPTX CudaOnlyExportPTX)
 ADD_TEST_MACRO(CudaOnly.GPUDebugFlag CudaOnlyGPUDebugFlag)
 ADD_TEST_MACRO(CudaOnly.ResolveDeviceSymbols CudaOnlyResolveDeviceSymbols)
 ADD_TEST_MACRO(CudaOnly.SeparateCompilation CudaOnlySeparateCompilation)
+ADD_TEST_MACRO(CudaOnly.SharedRuntimePlusToolkit CudaOnlySharedRuntimePlusToolkit)
 ADD_TEST_MACRO(CudaOnly.Standard98 CudaOnlyStandard98)
 ADD_TEST_MACRO(CudaOnly.Toolkit CudaOnlyToolkit)
 ADD_TEST_MACRO(CudaOnly.WithDefs CudaOnlyWithDefs)
 
+# The CUDA only ships the shared version of the toolkit libraries
+# on windows
+if(NOT WIN32)
+  ADD_TEST_MACRO(Cuda.StaticRuntimePlusToolkit StaticRuntimePlusToolkit)
+endif()
+
+if(MSVC)
+  ADD_TEST_MACRO(CudaOnly.PDB CudaOnlyPDB)
+endif()
+
 add_test(NAME CudaOnly.DontResolveDeviceSymbols COMMAND
   ${CMAKE_CTEST_COMMAND} -C $<CONFIGURATION>
   --build-and-test
@@ -20,6 +31,14 @@ add_test(NAME CudaOnly.DontResolveDeviceSymbols COMMAND
   --test-command ${CMAKE_CTEST_COMMAND} -V -C $<CONFIGURATION>
   )
 
-if(MSVC)
-  ADD_TEST_MACRO(CudaOnly.PDB CudaOnlyPDB)
-endif()
+add_test(NAME CudaOnly.RuntimeControls COMMAND
+  ${CMAKE_CTEST_COMMAND} -C $<CONFIGURATION>
+  --build-and-test
+  "${CMAKE_CURRENT_SOURCE_DIR}/RuntimeControls/"
+  "${CMAKE_CURRENT_BINARY_DIR}/RuntimeControls/"
+  --build-two-config
+    ${build_generator_args}
+  --build-project RuntimeControls
+  --build-options ${build_options}
+  --test-command ${CMAKE_CTEST_COMMAND} -V -C $<CONFIGURATION>
+  )
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()
diff --git a/Tests/CudaOnly/SharedRuntimePlusToolkit/CMakeLists.txt b/Tests/CudaOnly/SharedRuntimePlusToolkit/CMakeLists.txt
new file mode 100644
index 0000000..03fba22
--- /dev/null
+++ b/Tests/CudaOnly/SharedRuntimePlusToolkit/CMakeLists.txt
@@ -0,0 +1,42 @@
+cmake_minimum_required(VERSION 3.15)
+project(SharedRuntimePlusToolkit CUDA)
+
+#Goal for this example:
+# Validate that with c++ we can use some components of the CUDA toolkit, and
+# specify the cuda runtime
+find_package(CUDAToolkit REQUIRED)
+
+add_library(Common OBJECT curand.cu nppif.cu)
+target_link_libraries(Common PRIVATE CUDA::toolkit)
+set_target_properties(Common PROPERTIES POSITION_INDEPENDENT_CODE ON)
+
+#shared runtime with shared toolkit libraries
+add_library(SharedToolkit SHARED shared.cu)
+target_link_libraries(SharedToolkit PRIVATE Common PUBLIC CUDA::curand CUDA::nppif)
+set_target_properties(SharedToolkit PROPERTIES CUDA_RUNTIME_LIBRARY none)
+target_link_libraries(SharedToolkit PUBLIC CUDA::cudart)
+
+# The CUDA only ships the shared version of the toolkit libraries
+# on windows
+if(NOT WIN32)
+  #shared runtime with static toolkit libraries
+  add_library(StaticToolkit SHARED static.cu)
+  target_link_libraries(StaticToolkit PRIVATE Common CUDA::curand_static CUDA::nppif_static)
+  set_target_properties(StaticToolkit PROPERTIES CUDA_RUNTIME_LIBRARY Shared)
+
+  #static runtime with mixed toolkit libraries
+  add_library(MixedToolkit SHARED mixed.cu)
+  target_link_libraries(MixedToolkit PRIVATE Common CUDA::curand_static CUDA::nppif)
+  set_target_properties(MixedToolkit PROPERTIES CUDA_RUNTIME_LIBRARY Shared)
+endif()
+
+add_executable(CudaOnlySharedRuntimePlusToolkit main.cu)
+target_link_libraries(CudaOnlySharedRuntimePlusToolkit PRIVATE SharedToolkit
+                      $<TARGET_NAME_IF_EXISTS:StaticToolkit>
+                      $<TARGET_NAME_IF_EXISTS:MixedToolkit>)
+
+if(UNIX)
+  # Help the shared cuda runtime find libcudart as it is not located
+  # in a default system searched location
+  set_property(TARGET CudaOnlySharedRuntimePlusToolkit PROPERTY BUILD_RPATH ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
+endif()
diff --git a/Tests/CudaOnly/SharedRuntimePlusToolkit/curand.cu b/Tests/CudaOnly/SharedRuntimePlusToolkit/curand.cu
new file mode 100644
index 0000000..fdd7b53
--- /dev/null
+++ b/Tests/CudaOnly/SharedRuntimePlusToolkit/curand.cu
@@ -0,0 +1,65 @@
+// Comes from:
+// https://docs.nvidia.com/cuda/curand/host-api-overview.html#host-api-example
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+/*
+ * This program uses the host CURAND API to generate 100
+ * pseudorandom floats.
+ */
+#include <cuda.h>
+#include <curand.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define CUDA_CALL(x)                                                          \
+  do {                                                                        \
+    if ((x) != cudaSuccess) {                                                 \
+      printf("Error at %s:%d\n", __FILE__, __LINE__);                         \
+      return EXIT_FAILURE;                                                    \
+    }                                                                         \
+  } while (0)
+#define CURAND_CALL(x)                                                        \
+  do {                                                                        \
+    if ((x) != CURAND_STATUS_SUCCESS) {                                       \
+      printf("Error at %s:%d\n", __FILE__, __LINE__);                         \
+      return EXIT_FAILURE;                                                    \
+    }                                                                         \
+  } while (0)
+
+EXPORT int curand_main()
+{
+  size_t n = 100;
+  size_t i;
+  curandGenerator_t gen;
+  float *devData, *hostData;
+
+  /* Allocate n floats on host */
+  hostData = (float*)calloc(n, sizeof(float));
+
+  /* Allocate n floats on device */
+  CUDA_CALL(cudaMalloc((void**)&devData, n * sizeof(float)));
+
+  /* Create pseudo-random number generator */
+  CURAND_CALL(curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT));
+
+  /* Set seed */
+  CURAND_CALL(curandSetPseudoRandomGeneratorSeed(gen, 1234ULL));
+
+  /* Generate n floats on device */
+  CURAND_CALL(curandGenerateUniform(gen, devData, n));
+
+  /* Copy device memory to host */
+  CUDA_CALL(
+    cudaMemcpy(hostData, devData, n * sizeof(float), cudaMemcpyDeviceToHost));
+
+  /* Cleanup */
+  CURAND_CALL(curandDestroyGenerator(gen));
+  CUDA_CALL(cudaFree(devData));
+  free(hostData);
+  return EXIT_SUCCESS;
+}
diff --git a/Tests/CudaOnly/SharedRuntimePlusToolkit/main.cu b/Tests/CudaOnly/SharedRuntimePlusToolkit/main.cu
new file mode 100644
index 0000000..2a4da22
--- /dev/null
+++ b/Tests/CudaOnly/SharedRuntimePlusToolkit/main.cu
@@ -0,0 +1,23 @@
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+IMPORT int shared_version();
+int static_version()
+{
+  return 0;
+}
+int mixed_version()
+{
+  return 0;
+}
+#else
+int shared_version();
+int static_version();
+int mixed_version();
+#endif
+
+int main()
+{
+  return mixed_version() == 0 && shared_version() == 0 &&
+    static_version() == 0;
+}
diff --git a/Tests/CudaOnly/SharedRuntimePlusToolkit/mixed.cu b/Tests/CudaOnly/SharedRuntimePlusToolkit/mixed.cu
new file mode 100644
index 0000000..6de6886
--- /dev/null
+++ b/Tests/CudaOnly/SharedRuntimePlusToolkit/mixed.cu
@@ -0,0 +1,16 @@
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#  define EXPORT __declspec(dllexport)
+#else
+#  define IMPORT
+#  define EXPORT
+#endif
+
+IMPORT int curand_main();
+IMPORT int nppif_main();
+
+EXPORT int mixed_version()
+{
+  return curand_main() == 0 && nppif_main() == 0;
+}
diff --git a/Tests/CudaOnly/SharedRuntimePlusToolkit/nppif.cu b/Tests/CudaOnly/SharedRuntimePlusToolkit/nppif.cu
new file mode 100644
index 0000000..ac5341c
--- /dev/null
+++ b/Tests/CudaOnly/SharedRuntimePlusToolkit/nppif.cu
@@ -0,0 +1,92 @@
+// Comes from
+// https://devtalk.nvidia.com/default/topic/1037482/gpu-accelerated-libraries/help-me-help-you-with-modern-cmake-and-cuda-mwe-for-npp/post/5271066/#5271066
+
+#ifdef _WIN32
+#  define EXPORT __declspec(dllexport)
+#else
+#  define EXPORT
+#endif
+
+#include <cstdio>
+#include <iostream>
+
+#include <assert.h>
+#include <cuda_runtime_api.h>
+#include <nppi_filtering_functions.h>
+
+EXPORT int nppif_main()
+{
+  /**
+   * 8-bit unsigned single-channel 1D row convolution.
+   */
+  const int simgrows = 32;
+  const int simgcols = 32;
+  Npp8u *d_pSrc, *d_pDst;
+  const int nMaskSize = 3;
+  NppiSize oROI;
+  oROI.width = simgcols - nMaskSize;
+  oROI.height = simgrows;
+  const int simgsize = simgrows * simgcols * sizeof(d_pSrc[0]);
+  const int dimgsize = oROI.width * oROI.height * sizeof(d_pSrc[0]);
+  const int simgpix = simgrows * simgcols;
+  const int dimgpix = oROI.width * oROI.height;
+  const int nSrcStep = simgcols * sizeof(d_pSrc[0]);
+  const int nDstStep = oROI.width * sizeof(d_pDst[0]);
+  const int pixval = 1;
+  const int nDivisor = 1;
+  const Npp32s h_pKernel[nMaskSize] = { pixval, pixval, pixval };
+  Npp32s* d_pKernel;
+  const Npp32s nAnchor = 2;
+  cudaError_t err = cudaMalloc((void**)&d_pSrc, simgsize);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  err = cudaMalloc((void**)&d_pDst, dimgsize);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  err = cudaMalloc((void**)&d_pKernel, nMaskSize * sizeof(d_pKernel[0]));
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  // set image to pixval initially
+  err = cudaMemset(d_pSrc, pixval, simgsize);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  err = cudaMemset(d_pDst, 0, dimgsize);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  err = cudaMemcpy(d_pKernel, h_pKernel, nMaskSize * sizeof(d_pKernel[0]),
+                   cudaMemcpyHostToDevice);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  // copy src to dst
+  NppStatus ret =
+    nppiFilterRow_8u_C1R(d_pSrc, nSrcStep, d_pDst, nDstStep, oROI, d_pKernel,
+                         nMaskSize, nAnchor, nDivisor);
+  assert(ret == NPP_NO_ERROR);
+  Npp8u* h_imgres = new Npp8u[dimgpix];
+  err = cudaMemcpy(h_imgres, d_pDst, dimgsize, cudaMemcpyDeviceToHost);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  // test for filtering
+  for (int i = 0; i < dimgpix; i++) {
+    if (h_imgres[i] != (pixval * pixval * nMaskSize)) {
+      fprintf(stderr, "h_imgres at index %d failed to match\n", i);
+      return 1;
+    }
+  }
+
+  return 0;
+}
diff --git a/Tests/CudaOnly/SharedRuntimePlusToolkit/shared.cu b/Tests/CudaOnly/SharedRuntimePlusToolkit/shared.cu
new file mode 100644
index 0000000..f3c3dbc
--- /dev/null
+++ b/Tests/CudaOnly/SharedRuntimePlusToolkit/shared.cu
@@ -0,0 +1,16 @@
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#  define EXPORT __declspec(dllexport)
+#else
+#  define IMPORT
+#  define EXPORT
+#endif
+
+int curand_main();
+int nppif_main();
+
+EXPORT int shared_version()
+{
+  return curand_main() == 0 && nppif_main() == 0;
+}
diff --git a/Tests/CudaOnly/SharedRuntimePlusToolkit/static.cu b/Tests/CudaOnly/SharedRuntimePlusToolkit/static.cu
new file mode 100644
index 0000000..6932fa3
--- /dev/null
+++ b/Tests/CudaOnly/SharedRuntimePlusToolkit/static.cu
@@ -0,0 +1,16 @@
+
+#ifdef _WIN32
+#  define IMPORT __declspec(dllimport)
+#  define EXPORT __declspec(dllexport)
+#else
+#  define IMPORT
+#  define EXPORT
+#endif
+
+IMPORT int curand_main();
+IMPORT int nppif_main();
+
+EXPORT int static_version()
+{
+  return curand_main() == 0 && nppif_main() == 0;
+}
diff --git a/Tests/CudaOnly/StaticRuntimePlusToolkit/CMakeLists.txt b/Tests/CudaOnly/StaticRuntimePlusToolkit/CMakeLists.txt
new file mode 100644
index 0000000..97ac229
--- /dev/null
+++ b/Tests/CudaOnly/StaticRuntimePlusToolkit/CMakeLists.txt
@@ -0,0 +1,29 @@
+cmake_minimum_required(VERSION 3.15)
+project(StaticRuntimePlusToolkit CUDA)
+
+#Goal for this example:
+# Validate that with cuda we can use some components of the CUDA toolkit, and
+# specify the cuda runtime
+find_package(CUDAToolkit REQUIRED)
+
+add_library(Common OBJECT curand.cu nppif.cu)
+target_link_libraries(Common PRIVATE CUDA::toolkit)
+set_target_properties(Common PROPERTIES POSITION_INDEPENDENT_CODE ON)
+
+#static runtime with shared toolkit libraries
+add_library(SharedToolkit SHARED shared.cu)
+target_link_libraries(SharedToolkit PRIVATE Common CUDA::curand CUDA::nppif )
+set_target_properties(SharedToolkit PROPERTIES CUDA_RUNTIME_LIBRARY none)
+target_link_libraries(SharedToolkit PUBLIC CUDA::cudart_static)
+
+#static runtime with static toolkit libraries
+add_library(StaticToolkit SHARED static.cu)
+target_link_libraries(StaticToolkit PRIVATE Common CUDA::curand_static CUDA::nppif_static)
+
+#static runtime with mixed toolkit libraries
+add_library(MixedToolkit SHARED mixed.cu)
+target_link_libraries(MixedToolkit PRIVATE Common CUDA::curand CUDA::nppif_static)
+set_target_properties(MixedToolkit PROPERTIES CUDA_RUNTIME_LIBRARY Static)
+
+add_executable(CudaOnlyStaticRuntimePlusToolkit main.cu)
+target_link_libraries(CudaOnlyStaticRuntimePlusToolkit PRIVATE SharedToolkit StaticToolkit MixedToolkit)
diff --git a/Tests/CudaOnly/StaticRuntimePlusToolkit/curand.cu b/Tests/CudaOnly/StaticRuntimePlusToolkit/curand.cu
new file mode 100644
index 0000000..95872f0
--- /dev/null
+++ b/Tests/CudaOnly/StaticRuntimePlusToolkit/curand.cu
@@ -0,0 +1,59 @@
+// Comes from:
+// https://docs.nvidia.com/cuda/curand/host-api-overview.html#host-api-example
+
+/*
+ * This program uses the host CURAND API to generate 100
+ * pseudorandom floats.
+ */
+#include <cuda.h>
+#include <curand.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define CUDA_CALL(x)                                                          \
+  do {                                                                        \
+    if ((x) != cudaSuccess) {                                                 \
+      printf("Error at %s:%d\n", __FILE__, __LINE__);                         \
+      return EXIT_FAILURE;                                                    \
+    }                                                                         \
+  } while (0)
+#define CURAND_CALL(x)                                                        \
+  do {                                                                        \
+    if ((x) != CURAND_STATUS_SUCCESS) {                                       \
+      printf("Error at %s:%d\n", __FILE__, __LINE__);                         \
+      return EXIT_FAILURE;                                                    \
+    }                                                                         \
+  } while (0)
+
+int curand_main()
+{
+  size_t n = 100;
+  size_t i;
+  curandGenerator_t gen;
+  float *devData, *hostData;
+
+  /* Allocate n floats on host */
+  hostData = (float*)calloc(n, sizeof(float));
+
+  /* Allocate n floats on device */
+  CUDA_CALL(cudaMalloc((void**)&devData, n * sizeof(float)));
+
+  /* Create pseudo-random number generator */
+  CURAND_CALL(curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT));
+
+  /* Set seed */
+  CURAND_CALL(curandSetPseudoRandomGeneratorSeed(gen, 1234ULL));
+
+  /* Generate n floats on device */
+  CURAND_CALL(curandGenerateUniform(gen, devData, n));
+
+  /* Copy device memory to host */
+  CUDA_CALL(
+    cudaMemcpy(hostData, devData, n * sizeof(float), cudaMemcpyDeviceToHost));
+
+  /* Cleanup */
+  CURAND_CALL(curandDestroyGenerator(gen));
+  CUDA_CALL(cudaFree(devData));
+  free(hostData);
+  return EXIT_SUCCESS;
+}
diff --git a/Tests/CudaOnly/StaticRuntimePlusToolkit/main.cu b/Tests/CudaOnly/StaticRuntimePlusToolkit/main.cu
new file mode 100644
index 0000000..5a09f8e
--- /dev/null
+++ b/Tests/CudaOnly/StaticRuntimePlusToolkit/main.cu
@@ -0,0 +1,11 @@
+
+
+int shared_version();
+int static_version();
+int mixed_version();
+
+int main()
+{
+  return mixed_version() == 0 && shared_version() == 0 &&
+    static_version() == 0;
+}
diff --git a/Tests/CudaOnly/StaticRuntimePlusToolkit/mixed.cu b/Tests/CudaOnly/StaticRuntimePlusToolkit/mixed.cu
new file mode 100644
index 0000000..a05140d
--- /dev/null
+++ b/Tests/CudaOnly/StaticRuntimePlusToolkit/mixed.cu
@@ -0,0 +1,8 @@
+
+int curand_main();
+int nppif_main();
+
+int mixed_version()
+{
+  return curand_main() == 0 && nppif_main() == 0;
+}
diff --git a/Tests/CudaOnly/StaticRuntimePlusToolkit/nppif.cu b/Tests/CudaOnly/StaticRuntimePlusToolkit/nppif.cu
new file mode 100644
index 0000000..2871090
--- /dev/null
+++ b/Tests/CudaOnly/StaticRuntimePlusToolkit/nppif.cu
@@ -0,0 +1,86 @@
+// Comes from
+// https://devtalk.nvidia.com/default/topic/1037482/gpu-accelerated-libraries/help-me-help-you-with-modern-cmake-and-cuda-mwe-for-npp/post/5271066/#5271066
+
+#include <cstdio>
+#include <iostream>
+
+#include <assert.h>
+#include <cuda_runtime_api.h>
+#include <nppi_filtering_functions.h>
+
+int nppif_main()
+{
+  /**
+   * 8-bit unsigned single-channel 1D row convolution.
+   */
+  const int simgrows = 32;
+  const int simgcols = 32;
+  Npp8u *d_pSrc, *d_pDst;
+  const int nMaskSize = 3;
+  NppiSize oROI;
+  oROI.width = simgcols - nMaskSize;
+  oROI.height = simgrows;
+  const int simgsize = simgrows * simgcols * sizeof(d_pSrc[0]);
+  const int dimgsize = oROI.width * oROI.height * sizeof(d_pSrc[0]);
+  const int simgpix = simgrows * simgcols;
+  const int dimgpix = oROI.width * oROI.height;
+  const int nSrcStep = simgcols * sizeof(d_pSrc[0]);
+  const int nDstStep = oROI.width * sizeof(d_pDst[0]);
+  const int pixval = 1;
+  const int nDivisor = 1;
+  const Npp32s h_pKernel[nMaskSize] = { pixval, pixval, pixval };
+  Npp32s* d_pKernel;
+  const Npp32s nAnchor = 2;
+  cudaError_t err = cudaMalloc((void**)&d_pSrc, simgsize);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  err = cudaMalloc((void**)&d_pDst, dimgsize);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  err = cudaMalloc((void**)&d_pKernel, nMaskSize * sizeof(d_pKernel[0]));
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  // set image to pixval initially
+  err = cudaMemset(d_pSrc, pixval, simgsize);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  err = cudaMemset(d_pDst, 0, dimgsize);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  err = cudaMemcpy(d_pKernel, h_pKernel, nMaskSize * sizeof(d_pKernel[0]),
+                   cudaMemcpyHostToDevice);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  // copy src to dst
+  NppStatus ret =
+    nppiFilterRow_8u_C1R(d_pSrc, nSrcStep, d_pDst, nDstStep, oROI, d_pKernel,
+                         nMaskSize, nAnchor, nDivisor);
+  assert(ret == NPP_NO_ERROR);
+  Npp8u* h_imgres = new Npp8u[dimgpix];
+  err = cudaMemcpy(h_imgres, d_pDst, dimgsize, cudaMemcpyDeviceToHost);
+  if (err != cudaSuccess) {
+    fprintf(stderr, "Cuda error %d\n", __LINE__);
+    return 1;
+  }
+  // test for filtering
+  for (int i = 0; i < dimgpix; i++) {
+    if (h_imgres[i] != (pixval * pixval * nMaskSize)) {
+      fprintf(stderr, "h_imgres at index %d failed to match\n", i);
+      return 1;
+    }
+  }
+
+  return 0;
+}
diff --git a/Tests/CudaOnly/StaticRuntimePlusToolkit/shared.cu b/Tests/CudaOnly/StaticRuntimePlusToolkit/shared.cu
new file mode 100644
index 0000000..9967b66
--- /dev/null
+++ b/Tests/CudaOnly/StaticRuntimePlusToolkit/shared.cu
@@ -0,0 +1,8 @@
+
+int curand_main();
+int nppif_main();
+
+int shared_version()
+{
+  return curand_main() == 0 && nppif_main() == 0;
+}
diff --git a/Tests/CudaOnly/StaticRuntimePlusToolkit/static.cu b/Tests/CudaOnly/StaticRuntimePlusToolkit/static.cu
new file mode 100644
index 0000000..ca7eb4c
--- /dev/null
+++ b/Tests/CudaOnly/StaticRuntimePlusToolkit/static.cu
@@ -0,0 +1,8 @@
+
+int curand_main();
+int nppif_main();
+
+int static_version()
+{
+  return curand_main() == 0 && nppif_main() == 0;
+}