CUDA: Honor CUDA_RESOLVE_DEVICE_SYMBOLS for more target types

`CUDA_RESOLVE_DEVICE_SYMBOLS` can be used with shared, module, and
executable target types.  This relaxation is to allow for better
interoperability with linkers that automatically do CUDA device symbol
resolution and have no way to disable it.

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()