5 files changed, 151 insertions, 0 deletions
diff --git a/Tests/CudaOnly/CMakeLists.txt b/Tests/CudaOnly/CMakeLists.txt
index 9c4f86a..f1fd344 100644
--- a/Tests/CudaOnly/CMakeLists.txt
+++ b/Tests/CudaOnly/CMakeLists.txt
@@ -7,6 +7,17 @@ ADD_TEST_MACRO(CudaOnly.ResolveDeviceSymbols CudaOnlyResolveDeviceSymbols)
 ADD_TEST_MACRO(CudaOnly.SeparateCompilation CudaOnlySeparateCompilation)
 ADD_TEST_MACRO(CudaOnly.WithDefs CudaOnlyWithDefs)
 
+add_test(NAME CudaOnly.DontResolveDeviceSymbols COMMAND
+  ${CMAKE_CTEST_COMMAND} -C $<CONFIGURATION>
+  --build-and-test
+  "${CMAKE_CURRENT_SOURCE_DIR}/DontResolveDeviceSymbols/"
+  "${CMAKE_CURRENT_BINARY_DIR}/DontResolveDeviceSymbols/"
+  ${build_generator_args}
+  --build-project DontResolveDeviceSymbols
+  --build-options ${build_options}
+  --test-command ${CMAKE_CTEST_COMMAND} -V -C $<CONFIGURATION>
+  )
+
 if(MSVC)
   ADD_TEST_MACRO(CudaOnly.PDB CudaOnlyPDB)
 endif()
diff --git a/Tests/CudaOnly/DontResolveDeviceSymbols/CMakeLists.txt b/Tests/CudaOnly/DontResolveDeviceSymbols/CMakeLists.txt
new file mode 100644
index 0000000..6190089
--- /dev/null
+++ b/Tests/CudaOnly/DontResolveDeviceSymbols/CMakeLists.txt
@@ -0,0 +1,50 @@
+cmake_minimum_required(VERSION 3.13)
+project (DontResolveDeviceSymbols CUDA)
+
+# Find nm and dumpbin
+if(CMAKE_NM)
+  set(dump_command ${CMAKE_NM})
+  set(dump_args --defined-only)
+  set(symbol_name cudaRegisterLinkedBinary)
+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 /SYMBOLS)
+    set(symbol_name nv_fatb)
+  endif()
+endif()
+
+
+#Goal for this example:
+# Build a static library that defines multiple methods and kernels that
+# use each other.
+# Don't resolve the device symbols in the static library
+# Don't resolve the device symbols in the executable library
+# Verify that we can't use those device symbols from anything
+string(APPEND CMAKE_CUDA_FLAGS " -gencode arch=compute_30,code=[compute_30] -gencode arch=compute_50,code=\\\"compute_50\\\"")
+set(CMAKE_CXX_STANDARD 11)
+set(CMAKE_CUDA_STANDARD 11)
+
+add_library(CUDANoDeviceResolve SHARED file1.cu)
+set_target_properties(CUDANoDeviceResolve
+                      PROPERTIES
+                      CUDA_SEPARABLE_COMPILATION ON
+                      CUDA_RESOLVE_DEVICE_SYMBOLS OFF
+                      POSITION_INDEPENDENT_CODE ON)
+if(MSVC)
+  target_link_options(CUDANoDeviceResolve PRIVATE "/FORCE:UNRESOLVED")
+endif()
+
+if(dump_command)
+add_custom_command(TARGET CUDANoDeviceResolve POST_BUILD
+  COMMAND ${CMAKE_COMMAND}
+  -DDUMP_COMMAND=${dump_command}
+  -DDUMP_ARGS=${dump_args}
+  -DSYMBOL_NAME=${symbol_name}
+  -DTEST_LIBRARY_PATH=$<TARGET_FILE:CUDANoDeviceResolve>
+  -P ${CMAKE_CURRENT_SOURCE_DIR}/verify.cmake
+  )
+endif()
diff --git a/Tests/CudaOnly/DontResolveDeviceSymbols/file1.cu b/Tests/CudaOnly/DontResolveDeviceSymbols/file1.cu
new file mode 100644
index 0000000..3924f67
--- /dev/null
+++ b/Tests/CudaOnly/DontResolveDeviceSymbols/file1.cu
@@ -0,0 +1,69 @@
+
+#include <iostream>
+
+static __global__ void file1_kernel(int in, int* out)
+{
+  *out = in * in;
+}
+
+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 file1_launch_kernel()
+{
+  int ret = choose_cuda_device();
+  if (ret) {
+    return 0;
+  }
+
+  int input = 4;
+
+  int* output;
+  cudaError_t err = cudaMallocManaged(&output, sizeof(int));
+  cudaDeviceSynchronize();
+  if (err != cudaSuccess) {
+    return 1;
+  }
+
+  file1_kernel<<<1, 1>>>(input, output);
+  cudaDeviceSynchronize();
+  err = cudaGetLastError();
+  std::cout << err << " " << cudaGetErrorString(err) << std::endl;
+  if (err == cudaSuccess) {
+    // This kernel launch should failed as the device linking never occured
+    std::cerr << "file1_kernel: kernel launch should have failed" << std::endl;
+    return 1;
+  }
+  return 0;
+}
diff --git a/Tests/CudaOnly/DontResolveDeviceSymbols/main.cu b/Tests/CudaOnly/DontResolveDeviceSymbols/main.cu
new file mode 100644
index 0000000..84a7a19
--- /dev/null
+++ b/Tests/CudaOnly/DontResolveDeviceSymbols/main.cu
@@ -0,0 +1,7 @@
+
+#include <iostream>
+
+int main(int argc, char** argv)
+{
+  return 0;
+}
diff --git a/Tests/CudaOnly/DontResolveDeviceSymbols/verify.cmake b/Tests/CudaOnly/DontResolveDeviceSymbols/verify.cmake
new file mode 100644
index 0000000..9bb426d
--- /dev/null
+++ b/Tests/CudaOnly/DontResolveDeviceSymbols/verify.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 "${SYMBOL_NAME}")
+  message(FATAL_ERROR
+    "The '${SYMBOL_NAME}' symbol is defined; device linking occurred!")
+endif()