Merge topic 'cuda-windows'

f9a810f7 CUDA: Enable CudaOnly.SeparateCompilation test runtime calls
5599d858 CUDA: Port test cases to Windows with MSVC host compiler
11551702 CUDA: Populate NVIDIA compiler information on Windows
5365421e CUDA: Detect implicit link information on Windows
522b913f CUDA: Find MSVC binutils on Windows
02582b91 CUDA: Populate compiler PDB placeholder during device linking
d470cb70 CUDA: Use `.obj` object file extension on Windows
a2e80cb0 CUDA: Detect MSVC architecture id
65c1e012 CUDA: Detect use of MSVC host compiler
945dd207 CUDA: Allow platform files to set device linking rules
95420cea CMakeParseImplicitLinkInfo: Add support for MSVC invoked by CUDA nvcc

29 files changed, 266 insertions, 46 deletions

diff --git a/Modules/CMakeCUDACompiler.cmake.in b/Modules/CMakeCUDACompiler.cmake.in
index 8a6c0bc..1ba42d9 100644
--- a/Modules/CMakeCUDACompiler.cmake.in
+++ b/Modules/CMakeCUDACompiler.cmake.in
@@ -4,6 +4,9 @@ set(CMAKE_CUDA_HOST_LINK_LAUNCHER "@CMAKE_CUDA_HOST_LINK_LAUNCHER@")
 set(CMAKE_CUDA_COMPILER_ID "@CMAKE_CUDA_COMPILER_ID@")
 set(CMAKE_CUDA_COMPILER_VERSION "@CMAKE_CUDA_COMPILER_VERSION@")
 set(CMAKE_CUDA_STANDARD_COMPUTED_DEFAULT "@CMAKE_CUDA_STANDARD_COMPUTED_DEFAULT@")
+set(CMAKE_CUDA_SIMULATE_ID "@CMAKE_CUDA_SIMULATE_ID@")
+set(CMAKE_CUDA_SIMULATE_VERSION "@CMAKE_CUDA_SIMULATE_VERSION@")
+@SET_MSVC_CUDA_ARCHITECTURE_ID@
 
 set(CMAKE_CUDA_COMPILER_ENV_VAR "CUDACXX")
 set(CMAKE_CUDA_HOST_COMPILER_ENV_VAR "CUDAHOSTCXX")
diff --git a/Modules/CMakeCUDACompilerId.cu.in b/Modules/CMakeCUDACompilerId.cu.in
index a20f1b1..018bab7 100644
--- a/Modules/CMakeCUDACompilerId.cu.in
+++ b/Modules/CMakeCUDACompilerId.cu.in
@@ -9,6 +9,9 @@
    because some compilers will just produce instructions to fill the
    array rather than assigning a pointer to a static array.  */
 char const* info_compiler = "INFO" ":" "compiler[" COMPILER_ID "]";
+#ifdef SIMULATE_ID
+char const* info_simulate = "INFO" ":" "simulate[" SIMULATE_ID "]";
+#endif
 
 @CMAKE_CUDA_COMPILER_ID_PLATFORM_CONTENT@
 @CMAKE_CUDA_COMPILER_ID_ERROR_FOR_TEST@
@@ -35,6 +38,12 @@ int main(int argc, char* argv[])
 #ifdef COMPILER_VERSION_MAJOR
   require += info_version[argc];
 #endif
+#ifdef SIMULATE_ID
+  require += info_simulate[argc];
+#endif
+#ifdef SIMULATE_VERSION_MAJOR
+  require += info_simulate_version[argc];
+#endif
   require += info_language_dialect_default[argc];
   (void)argv;
   return require;
diff --git a/Modules/CMakeCUDAInformation.cmake b/Modules/CMakeCUDAInformation.cmake
index 202a7a6..abc4b66 100644
--- a/Modules/CMakeCUDAInformation.cmake
+++ b/Modules/CMakeCUDAInformation.cmake
@@ -1,7 +1,11 @@
 # Distributed under the OSI-approved BSD 3-Clause License.  See accompanying
 # file Copyright.txt or https://cmake.org/licensing for details.
 
-set(CMAKE_CUDA_OUTPUT_EXTENSION .o)
+if(UNIX)
+  set(CMAKE_CUDA_OUTPUT_EXTENSION .o)
+else()
+  set(CMAKE_CUDA_OUTPUT_EXTENSION .obj)
+endif()
 set(CMAKE_INCLUDE_FLAG_CUDA "-I")
 
 # Load compiler-specific information.
@@ -177,11 +181,14 @@ if(NOT CMAKE_CUDA_LINK_EXECUTABLE)
 endif()
 
 #These are used when linking relocatable (dc) cuda code
-set(CMAKE_CUDA_DEVICE_LINK_LIBRARY
-      "<CMAKE_CUDA_COMPILER> <CMAKE_CUDA_LINK_FLAGS> <LANGUAGE_COMPILE_FLAGS> ${CMAKE_CUDA_COMPILE_OPTIONS_PIC} -shared -dlink <OBJECTS> -o <TARGET> <LINK_LIBRARIES>")
-set(CMAKE_CUDA_DEVICE_LINK_EXECUTABLE
-      "<CMAKE_CUDA_COMPILER> <CMAKE_CUDA_LINK_FLAGS> ${CMAKE_CUDA_COMPILE_OPTIONS_PIC} -shared -dlink <OBJECTS> -o <TARGET> <LINK_LIBRARIES>")
-
+if(NOT CMAKE_CUDA_DEVICE_LINK_LIBRARY)
+  set(CMAKE_CUDA_DEVICE_LINK_LIBRARY
+    "<CMAKE_CUDA_COMPILER> <CMAKE_CUDA_LINK_FLAGS> <LANGUAGE_COMPILE_FLAGS> ${CMAKE_CUDA_COMPILE_OPTIONS_PIC} -shared -dlink <OBJECTS> -o <TARGET> <LINK_LIBRARIES>")
+endif()
+if(NOT CMAKE_CUDA_DEVICE_LINK_EXECUTABLE)
+  set(CMAKE_CUDA_DEVICE_LINK_EXECUTABLE
+    "<CMAKE_CUDA_COMPILER> <CMAKE_CUDA_LINK_FLAGS> ${CMAKE_CUDA_COMPILE_OPTIONS_PIC} -shared -dlink <OBJECTS> -o <TARGET> <LINK_LIBRARIES>")
+endif()
 
 mark_as_advanced(
 CMAKE_CUDA_FLAGS
diff --git a/Modules/CMakeDetermineCUDACompiler.cmake b/Modules/CMakeDetermineCUDACompiler.cmake
index e03de7e..bef6d0e 100644
--- a/Modules/CMakeDetermineCUDACompiler.cmake
+++ b/Modules/CMakeDetermineCUDACompiler.cmake
@@ -60,7 +60,7 @@ if(NOT CMAKE_CUDA_COMPILER_ID_RUN)
   set(CMAKE_CXX_COMPILER_ID_TOOL_MATCH_REGEX "\nLd[^\n]*(\n[ \t]+[^\n]*)*\n[ \t]+([^ \t\r\n]+)[^\r\n]*-o[^\r\n]*CompilerIdCUDA/(\\./)?(CompilerIdCUDA.xctest/)?CompilerIdCUDA[ \t\n\\\"]")
   set(CMAKE_CXX_COMPILER_ID_TOOL_MATCH_INDEX 2)
 
-  set(CMAKE_CUDA_COMPILER_ID_FLAGS_ALWAYS "-v")
+  set(CMAKE_CUDA_COMPILER_ID_FLAGS_ALWAYS -v --keep --keep-dir tmp)
   if(CMAKE_CUDA_HOST_COMPILER)
       list(APPEND CMAKE_CUDA_COMPILER_ID_FLAGS_ALWAYS "-ccbin=${CMAKE_CUDA_HOST_COMPILER}")
   endif()
@@ -70,6 +70,10 @@ if(NOT CMAKE_CUDA_COMPILER_ID_RUN)
 endif()
 
 include(CMakeFindBinUtils)
+if(MSVC_CUDA_ARCHITECTURE_ID)
+  set(SET_MSVC_CUDA_ARCHITECTURE_ID
+    "set(MSVC_CUDA_ARCHITECTURE_ID ${MSVC_CUDA_ARCHITECTURE_ID})")
+endif()
 
 #if this compiler vendor is matches NVIDIA we can determine
 #what the host compiler is. This only needs to be done if the CMAKE_CUDA_HOST_COMPILER
@@ -102,6 +106,14 @@ if(CMAKE_CUDA_COMPILER_ID STREQUAL NVIDIA)
       if("${_nvcc_output_line}" MATCHES "^ *nvlink")
         string(APPEND _nvcc_log "    ignoring nvlink line\n")
       elseif(_nvcc_libraries)
+        if("${_nvcc_output_line}" MATCHES "(@\"?tmp/a\\.exe\\.res\"?)")
+          set(_nvcc_link_res_arg "${CMAKE_MATCH_1}")
+          set(_nvcc_link_res "${CMAKE_PLATFORM_INFO_DIR}/CompilerIdCUDA/tmp/a.exe.res")
+          if(EXISTS "${_nvcc_link_res}")
+            file(READ "${_nvcc_link_res}" _nvcc_link_res_content)
+            string(REPLACE "${_nvcc_link_res_arg}" "${_nvcc_link_res_content}" _nvcc_output_line "${_nvcc_output_line}")
+          endif()
+        endif()
         string(FIND "${_nvcc_output_line}" "${_nvcc_libraries}" _nvcc_libraries_pos)
         if(NOT _nvcc_libraries_pos EQUAL -1)
           set(_nvcc_link_line "${_nvcc_output_line}")
@@ -112,9 +124,13 @@ if(CMAKE_CUDA_COMPILER_ID STREQUAL NVIDIA)
   endif()
 
   if(_nvcc_link_line)
-    #extract the compiler that is being used for linking
-    separate_arguments(_nvcc_link_line_args UNIX_COMMAND "${_nvcc_link_line}")
-    list(GET _nvcc_link_line_args 0 CMAKE_CUDA_HOST_LINK_LAUNCHER)
+    if("x${CMAKE_CUDA_SIMULATE_ID}" STREQUAL "xMSVC")
+      set(CMAKE_CUDA_HOST_LINK_LAUNCHER "${CMAKE_LINKER}")
+    else()
+      #extract the compiler that is being used for linking
+      separate_arguments(_nvcc_link_line_args UNIX_COMMAND "${_nvcc_link_line}")
+      list(GET _nvcc_link_line_args 0 CMAKE_CUDA_HOST_LINK_LAUNCHER)
+    endif()
 
     #prefix the line with cuda-fake-ld so that implicit link info believes it is
     #a link line
diff --git a/Modules/CMakeDetermineCompilerId.cmake b/Modules/CMakeDetermineCompilerId.cmake
index eae139d..bb34de5 100644
--- a/Modules/CMakeDetermineCompilerId.cmake
+++ b/Modules/CMakeDetermineCompilerId.cmake
@@ -131,6 +131,7 @@ function(CMAKE_DETERMINE_COMPILER_ID_BUILD lang testflags src)
   # Create a clean working directory.
   file(REMOVE_RECURSE ${CMAKE_${lang}_COMPILER_ID_DIR})
   file(MAKE_DIRECTORY ${CMAKE_${lang}_COMPILER_ID_DIR})
+  file(MAKE_DIRECTORY ${CMAKE_${lang}_COMPILER_ID_DIR}/tmp)
   CMAKE_DETERMINE_COMPILER_ID_WRITE("${lang}" "${src}")
 
   # Construct a description of this test case.
diff --git a/Modules/CMakeFindBinUtils.cmake b/Modules/CMakeFindBinUtils.cmake
index 75a031e..4c0486e 100644
--- a/Modules/CMakeFindBinUtils.cmake
+++ b/Modules/CMakeFindBinUtils.cmake
@@ -25,6 +25,7 @@ if("x${CMAKE_C_SIMULATE_ID}" STREQUAL "xMSVC"
    OR "x${CMAKE_Fortran_SIMULATE_ID}" STREQUAL "xMSVC"
    OR "x${CMAKE_C_COMPILER_ID}" STREQUAL "xMSVC"
    OR "x${CMAKE_CXX_COMPILER_ID}" STREQUAL "xMSVC"
+   OR "x${CMAKE_CUDA_SIMULATE_ID}" STREQUAL "xMSVC"
    OR (CMAKE_GENERATOR MATCHES "Visual Studio"
        AND NOT CMAKE_VS_PLATFORM_NAME STREQUAL "Tegra-Android"))
 
diff --git a/Modules/CMakeParseImplicitLinkInfo.cmake b/Modules/CMakeParseImplicitLinkInfo.cmake
index 2031ba5..3469d34 100644
--- a/Modules/CMakeParseImplicitLinkInfo.cmake
+++ b/Modules/CMakeParseImplicitLinkInfo.cmake
@@ -45,6 +45,7 @@ function(CMAKE_PARSE_IMPLICIT_LINK_INFO text lib_var dir_var fwk_var log_var obj
       endif()
       list(GET args 0 cmd)
     endif()
+    set(is_msvc 0)
     if("${cmd}" MATCHES "${linker_regex}")
       string(APPEND log "  link line: [${line}]\n")
       string(REGEX REPLACE ";-([LYz]);" ";-\\1" args "${args}")
@@ -54,6 +55,17 @@ function(CMAKE_PARSE_IMPLICIT_LINK_INFO text lib_var dir_var fwk_var log_var obj
           string(REGEX REPLACE "^-L" "" dir "${arg}")
           list(APPEND implicit_dirs_tmp ${dir})
           string(APPEND log "    arg [${arg}] ==> dir [${dir}]\n")
+        elseif("${arg}" MATCHES "^[-/]LIBPATH:(.+)")
+          # MSVC search path.
+          set(dir "${CMAKE_MATCH_1}")
+          list(APPEND implicit_dirs_tmp ${dir})
+          string(APPEND log "    arg [${arg}] ==> dir [${dir}]\n")
+        elseif(is_msvc AND "${arg}" STREQUAL "-link")
+          string(APPEND log "    arg [${arg}] ==> ignore MSVC cl option\n")
+        elseif(is_msvc AND "${arg}" MATCHES "^(.*\\.[Ll][Ii][Bb])$")
+          set(lib "${CMAKE_MATCH_1}")
+          list(APPEND implicit_libs_tmp ${lib})
+          string(APPEND log "    arg [${arg}] ==> lib [${lib}]\n")
         elseif("${arg}" MATCHES "^-l([^:].*)$")
           # Unix library.
           set(lib "${CMAKE_MATCH_1}")
@@ -82,6 +94,9 @@ function(CMAKE_PARSE_IMPLICIT_LINK_INFO text lib_var dir_var fwk_var log_var obj
           # Link editor option.
           list(APPEND implicit_libs_tmp ${arg})
           string(APPEND log "    arg [${arg}] ==> opt [${arg}]\n")
+        elseif("${arg}" STREQUAL "cl.exe")
+          string(APPEND log "    arg [${arg}] ==> recognize MSVC cl\n")
+          set(is_msvc 1)
         else()
           string(APPEND log "    arg [${arg}] ==> ignore\n")
         endif()
diff --git a/Modules/CMakeTestCUDACompiler.cmake b/Modules/CMakeTestCUDACompiler.cmake
index 670b31d..80113cb 100644
--- a/Modules/CMakeTestCUDACompiler.cmake
+++ b/Modules/CMakeTestCUDACompiler.cmake
@@ -58,6 +58,11 @@ else()
   include(${CMAKE_ROOT}/Modules/CMakeDetermineCompilerABI.cmake)
   CMAKE_DETERMINE_COMPILER_ABI(CUDA ${CMAKE_ROOT}/Modules/CMakeCUDACompilerABI.cu)
 
+  if("x${CMAKE_CUDA_SIMULATE_ID}" STREQUAL "xMSVC")
+    set(CMAKE_CUDA_IMPLICIT_LINK_LIBRARIES "${CMAKE_CUDA_HOST_IMPLICIT_LINK_LIBRARIES}")
+    set(CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES "${CMAKE_CUDA_HOST_IMPLICIT_LINK_DIRECTORIES}")
+  endif()
+
   # Re-configure to save learned information.
   configure_file(
     ${CMAKE_ROOT}/Modules/CMakeCUDACompiler.cmake.in
diff --git a/Modules/Compiler/NVIDIA-CUDA.cmake b/Modules/Compiler/NVIDIA-CUDA.cmake
index 605d555..ae35132 100644
--- a/Modules/Compiler/NVIDIA-CUDA.cmake
+++ b/Modules/Compiler/NVIDIA-CUDA.cmake
@@ -1,14 +1,15 @@
 set(CMAKE_CUDA_VERBOSE_FLAG "-v")
 
-
-set(CMAKE_CUDA_COMPILE_OPTIONS_PIE -Xcompiler=-fPIE)
-set(CMAKE_CUDA_COMPILE_OPTIONS_PIC -Xcompiler=-fPIC)
-#CMAKE_SHARED_LIBRARY_CUDA_FLAGS is sent to the host linker so we don' need
-#to forward it through nvcc
-set(CMAKE_SHARED_LIBRARY_CUDA_FLAGS -fPIC)
+if(NOT "x${CMAKE_CUDA_SIMULATE_ID}" STREQUAL "xMSVC")
+  set(CMAKE_CUDA_COMPILE_OPTIONS_PIE -Xcompiler=-fPIE)
+  set(CMAKE_CUDA_COMPILE_OPTIONS_PIC -Xcompiler=-fPIC)
+  set(CMAKE_CUDA_COMPILE_OPTIONS_VISIBILITY -Xcompiler=-fvisibility=)
+  # CMAKE_SHARED_LIBRARY_CUDA_FLAGS is sent to the host linker so we
+  # don't need to forward it through nvcc.
+  set(CMAKE_SHARED_LIBRARY_CUDA_FLAGS -fPIC)
+endif()
 set(CMAKE_SHARED_LIBRARY_CREATE_CUDA_FLAGS -shared)
 set(CMAKE_INCLUDE_SYSTEM_FLAG_CUDA -isystem=)
-set(CMAKE_CUDA_COMPILE_OPTIONS_VISIBILITY -Xcompiler=-fvisibility=)
 
 string(APPEND CMAKE_CUDA_FLAGS_INIT " ")
 string(APPEND CMAKE_CUDA_FLAGS_DEBUG_INIT " -g")
@@ -16,9 +17,12 @@ string(APPEND CMAKE_CUDA_FLAGS_MINSIZEREL_INIT " -Os -DNDEBUG")
 string(APPEND CMAKE_CUDA_FLAGS_RELEASE_INIT " -O3 -DNDEBUG")
 string(APPEND CMAKE_CUDA_FLAGS_RELWITHDEBINFO_INIT " -O2 -g -DNDEBUG")
 
-set(CMAKE_CUDA98_STANDARD_COMPILE_OPTION "")
-set(CMAKE_CUDA98_EXTENSION_COMPILE_OPTION "")
-set(CMAKE_CUDA11_STANDARD_COMPILE_OPTION "-std=c++11")
-set(CMAKE_CUDA11_EXTENSION_COMPILE_OPTION "-std=c++11")
-
-set(CMAKE_CUDA_STANDARD_DEFAULT 98)
+if("x${CMAKE_CUDA_SIMULATE_ID}" STREQUAL "xMSVC")
+  set(CMAKE_CUDA_STANDARD_DEFAULT "")
+else()
+  set(CMAKE_CUDA_STANDARD_DEFAULT 98)
+  set(CMAKE_CUDA98_STANDARD_COMPILE_OPTION "")
+  set(CMAKE_CUDA98_EXTENSION_COMPILE_OPTION "")
+  set(CMAKE_CUDA11_STANDARD_COMPILE_OPTION "-std=c++11")
+  set(CMAKE_CUDA11_EXTENSION_COMPILE_OPTION "-std=c++11")
+endif()
diff --git a/Modules/Compiler/NVIDIA-DetermineCompiler.cmake b/Modules/Compiler/NVIDIA-DetermineCompiler.cmake
index 32ccf8a..cb0beaf 100644
--- a/Modules/Compiler/NVIDIA-DetermineCompiler.cmake
+++ b/Modules/Compiler/NVIDIA-DetermineCompiler.cmake
@@ -4,4 +4,14 @@ set(_compiler_id_pp_test "defined(__NVCC__)")
 set(_compiler_id_version_compute "
 # define @PREFIX@COMPILER_VERSION_MAJOR @MACRO_DEC@(__CUDACC_VER_MAJOR__)
 # define @PREFIX@COMPILER_VERSION_MINOR @MACRO_DEC@(__CUDACC_VER_MINOR__)
-# define @PREFIX@COMPILER_VERSION_PATCH @MACRO_DEC@(__CUDACC_VER_BUILD__)")
+# define @PREFIX@COMPILER_VERSION_PATCH @MACRO_DEC@(__CUDACC_VER_BUILD__)
+# if defined(_MSC_VER)
+   /* _MSC_VER = VVRR */
+#  define @PREFIX@SIMULATE_VERSION_MAJOR @MACRO_DEC@(_MSC_VER / 100)
+#  define @PREFIX@SIMULATE_VERSION_MINOR @MACRO_DEC@(_MSC_VER % 100)
+# endif")
+
+set(_compiler_id_simulate "
+# if defined(_MSC_VER)
+#  define @PREFIX@SIMULATE_ID \"MSVC\"
+# endif")
diff --git a/Modules/Platform/Windows-MSVC.cmake b/Modules/Platform/Windows-MSVC.cmake
index f506500..0e90d35 100644
--- a/Modules/Platform/Windows-MSVC.cmake
+++ b/Modules/Platform/Windows-MSVC.cmake
@@ -54,6 +54,8 @@ if(NOT MSVC_VERSION)
     set(_compiler_version ${CMAKE_CXX_SIMULATE_VERSION})
   elseif(CMAKE_Fortran_SIMULATE_VERSION)
     set(_compiler_version ${CMAKE_Fortran_SIMULATE_VERSION})
+  elseif(CMAKE_CUDA_SIMULATE_VERSION)
+    set(_compiler_version ${CMAKE_CUDA_SIMULATE_VERSION})
   elseif(CMAKE_C_COMPILER_VERSION)
     set(_compiler_version ${CMAKE_C_COMPILER_VERSION})
   else()
diff --git a/Modules/Platform/Windows-NVIDIA-CUDA.cmake b/Modules/Platform/Windows-NVIDIA-CUDA.cmake
new file mode 100644
index 0000000..80ecfbc
--- /dev/null
+++ b/Modules/Platform/Windows-NVIDIA-CUDA.cmake
@@ -0,0 +1,42 @@
+include(Platform/Windows-MSVC)
+
+set(CMAKE_CUDA_COMPILE_PTX_COMPILATION
+  "<CMAKE_CUDA_COMPILER> ${CMAKE_CUDA_HOST_FLAGS} <DEFINES> <INCLUDES> <FLAGS> -x cu -ptx <SOURCE> -o <OBJECT> -Xcompiler=-Fd<TARGET_COMPILE_PDB>,-FS")
+set(CMAKE_CUDA_COMPILE_SEPARABLE_COMPILATION
+  "<CMAKE_CUDA_COMPILER> ${CMAKE_CUDA_HOST_FLAGS} <DEFINES> <INCLUDES> <FLAGS> -x cu -dc <SOURCE> -o <OBJECT> -Xcompiler=-Fd<TARGET_COMPILE_PDB>,-FS")
+set(CMAKE_CUDA_COMPILE_WHOLE_COMPILATION
+  "<CMAKE_CUDA_COMPILER> ${CMAKE_CUDA_HOST_FLAGS} <DEFINES> <INCLUDES> <FLAGS> -x cu -c <SOURCE> -o <OBJECT> -Xcompiler=-Fd<TARGET_COMPILE_PDB>,-FS")
+
+set(__IMPLICT_LINKS )
+foreach(dir ${CMAKE_CUDA_HOST_IMPLICIT_LINK_DIRECTORIES})
+  string(APPEND __IMPLICT_LINKS " -LIBPATH:\"${dir}\"")
+endforeach()
+foreach(lib ${CMAKE_CUDA_HOST_IMPLICIT_LINK_LIBRARIES})
+  string(APPEND __IMPLICT_LINKS " \"${lib}\"")
+endforeach()
+set(CMAKE_CUDA_LINK_EXECUTABLE
+   "<CMAKE_CUDA_HOST_LINK_LAUNCHER> <CMAKE_CUDA_LINK_FLAGS> <LINK_FLAGS> <OBJECTS> /out:<TARGET> /implib:<TARGET_IMPLIB> /pdb:<TARGET_PDB> /version:<TARGET_VERSION_MAJOR>.<TARGET_VERSION_MINOR> <LINK_LIBRARIES>${__IMPLICT_LINKS}")
+
+set(_CMAKE_VS_LINK_DLL "<CMAKE_COMMAND> -E vs_link_dll --intdir=<OBJECT_DIR> --manifests <MANIFESTS> -- ")
+set(_CMAKE_VS_LINK_EXE "<CMAKE_COMMAND> -E vs_link_exe --intdir=<OBJECT_DIR> --manifests <MANIFESTS> -- ")
+set(CMAKE_CUDA_CREATE_SHARED_LIBRARY
+  "${_CMAKE_VS_LINK_DLL}<CMAKE_LINKER> ${CMAKE_CL_NOLOGO} <OBJECTS> ${CMAKE_START_TEMP_FILE} /out:<TARGET> /implib:<TARGET_IMPLIB> /pdb:<TARGET_PDB> /dll /version:<TARGET_VERSION_MAJOR>.<TARGET_VERSION_MINOR>${_PLATFORM_LINK_FLAGS} <LINK_FLAGS> <LINK_LIBRARIES>${__IMPLICT_LINKS} ${CMAKE_END_TEMP_FILE}")
+
+set(CMAKE_CUDA_CREATE_SHARED_MODULE ${CMAKE_CUDA_CREATE_SHARED_LIBRARY})
+set(CMAKE_CUDA_CREATE_STATIC_LIBRARY  "<CMAKE_LINKER> /lib ${CMAKE_CL_NOLOGO} <LINK_FLAGS> /out:<TARGET> <OBJECTS> ")
+set(CMAKE_CUDA_LINKER_SUPPORTS_PDB ON)
+set(CMAKE_CUDA_LINK_EXECUTABLE
+  "${_CMAKE_VS_LINK_EXE}<CMAKE_LINKER> ${CMAKE_CL_NOLOGO} <OBJECTS> ${CMAKE_START_TEMP_FILE} /out:<TARGET> /implib:<TARGET_IMPLIB> /pdb:<TARGET_PDB> /version:<TARGET_VERSION_MAJOR>.<TARGET_VERSION_MINOR>${_PLATFORM_LINK_FLAGS} <CMAKE_CUDA_LINK_FLAGS> <LINK_FLAGS> <LINK_LIBRARIES>${__IMPLICT_LINKS} ${CMAKE_END_TEMP_FILE}")
+unset(_CMAKE_VS_LINK_EXE)
+unset(_CMAKE_VS_LINK_EXE)
+
+set(CMAKE_CUDA_DEVICE_LINK_LIBRARY
+  "<CMAKE_CUDA_COMPILER> <CMAKE_CUDA_LINK_FLAGS> <LANGUAGE_COMPILE_FLAGS> -shared -dlink <OBJECTS> -o <TARGET> <LINK_LIBRARIES> -Xcompiler=-Fd<TARGET_COMPILE_PDB>,-FS")
+set(CMAKE_CUDA_DEVICE_LINK_EXECUTABLE
+  "<CMAKE_CUDA_COMPILER> <CMAKE_CUDA_LINK_FLAGS> -shared -dlink <OBJECTS> -o <TARGET> <LINK_LIBRARIES> -Xcompiler=-Fd<TARGET_COMPILE_PDB>,-FS")
+
+string(APPEND CMAKE_CUDA_FLAGS_INIT " -Xcompiler=-GR,-EHsc")
+string(APPEND CMAKE_CUDA_FLAGS_DEBUG_INIT " -Xcompiler=-MDd,-Zi,-RTC1")
+string(APPEND CMAKE_CUDA_FLAGS_RELEASE_INIT " -Xcompiler=-MD")
+string(APPEND CMAKE_CUDA_FLAGS_RELWITHDEBINFO_INIT " -Xcompiler=-MD")
+string(APPEND CMAKE_CUDA_FLAGS_MINSIZEREL_INIT " -Xcompiler=-MD")
diff --git a/Source/cmMakefileExecutableTargetGenerator.cxx b/Source/cmMakefileExecutableTargetGenerator.cxx
index 069011d..4bc706c 100644
--- a/Source/cmMakefileExecutableTargetGenerator.cxx
+++ b/Source/cmMakefileExecutableTargetGenerator.cxx
@@ -104,10 +104,12 @@ void cmMakefileExecutableTargetGenerator::WriteDeviceExecutableRule(
 
   // Get the language to use for linking this library.
   std::string linkLanguage = "CUDA";
+  std::string const objExt =
+    this->Makefile->GetSafeDefinition("CMAKE_CUDA_OUTPUT_EXTENSION");
 
   // Get the name of the device object to generate.
   std::string const targetOutputReal =
-    this->GeneratorTarget->ObjectDirectory + "cmake_device_link.o";
+    this->GeneratorTarget->ObjectDirectory + "cmake_device_link" + objExt;
   this->DeviceLinkObject = targetOutputReal;
 
   this->NumberOfProgressActions++;
@@ -228,6 +230,11 @@ void cmMakefileExecutableTargetGenerator::WriteDeviceExecutableRule(
         this->LocalGenerator->GetCurrentBinaryDirectory(), targetOutputReal),
       output);
 
+    std::string targetFullPathCompilePDB = this->ComputeTargetCompilePDB();
+    std::string targetOutPathCompilePDB =
+      this->LocalGenerator->ConvertToOutputFormat(targetFullPathCompilePDB,
+                                                  cmOutputConverter::SHELL);
+
     vars.Language = linkLanguage.c_str();
     vars.Objects = buildObjs.c_str();
     vars.ObjectDir = objectDir.c_str();
@@ -235,6 +242,7 @@ void cmMakefileExecutableTargetGenerator::WriteDeviceExecutableRule(
     vars.LinkLibraries = linkLibs.c_str();
     vars.Flags = flags.c_str();
     vars.LinkFlags = linkFlags.c_str();
+    vars.TargetCompilePDB = targetOutPathCompilePDB.c_str();
 
     std::string launcher;
 
diff --git a/Source/cmMakefileLibraryTargetGenerator.cxx b/Source/cmMakefileLibraryTargetGenerator.cxx
index 2b0e1b1..27b7c21 100644
--- a/Source/cmMakefileLibraryTargetGenerator.cxx
+++ b/Source/cmMakefileLibraryTargetGenerator.cxx
@@ -281,6 +281,8 @@ void cmMakefileLibraryTargetGenerator::WriteDeviceLibraryRules(
 
   // Get the language to use for linking this library.
   std::string linkLanguage = "CUDA";
+  std::string const objExt =
+    this->Makefile->GetSafeDefinition("CMAKE_CUDA_OUTPUT_EXTENSION");
 
   // Create set of linking flags.
   std::string linkFlags;
@@ -288,7 +290,7 @@ void cmMakefileLibraryTargetGenerator::WriteDeviceLibraryRules(
 
   // Get the name of the device object to generate.
   std::string const targetOutputReal =
-    this->GeneratorTarget->ObjectDirectory + "cmake_device_link.o";
+    this->GeneratorTarget->ObjectDirectory + "cmake_device_link" + objExt;
   this->DeviceLinkObject = targetOutputReal;
 
   this->NumberOfProgressActions++;
@@ -364,12 +366,18 @@ void cmMakefileLibraryTargetGenerator::WriteDeviceLibraryRules(
         this->LocalGenerator->GetCurrentBinaryDirectory(), targetOutputReal),
       output);
 
+    std::string targetFullPathCompilePDB = this->ComputeTargetCompilePDB();
+    std::string targetOutPathCompilePDB =
+      this->LocalGenerator->ConvertToOutputFormat(targetFullPathCompilePDB,
+                                                  cmOutputConverter::SHELL);
+
     vars.Objects = buildObjs.c_str();
     vars.ObjectDir = objectDir.c_str();
     vars.Target = target.c_str();
     vars.LinkLibraries = linkLibs.c_str();
     vars.ObjectsQuoted = buildObjs.c_str();
     vars.LinkFlags = linkFlags.c_str();
+    vars.TargetCompilePDB = targetOutPathCompilePDB.c_str();
 
     // Add language feature flags.
     std::string langFlags;
diff --git a/Source/cmNinjaNormalTargetGenerator.cxx b/Source/cmNinjaNormalTargetGenerator.cxx
index 0db5687..bc8d8ff 100644
--- a/Source/cmNinjaNormalTargetGenerator.cxx
+++ b/Source/cmNinjaNormalTargetGenerator.cxx
@@ -229,6 +229,7 @@ void cmNinjaNormalTargetGenerator::WriteDeviceLinkRule(bool useResponseFile)
     vars.SONameFlag = "$SONAME_FLAG";
     vars.TargetSOName = "$SONAME";
     vars.TargetPDB = "$TARGET_PDB";
+    vars.TargetCompilePDB = "$TARGET_COMPILE_PDB";
 
     vars.Flags = "$FLAGS";
     vars.LinkFlags = "$LINK_FLAGS";
@@ -602,10 +603,12 @@ void cmNinjaNormalTargetGenerator::WriteDeviceLinkStatement()
   // First and very important step is to make sure while inside this
   // step our link language is set to CUDA
   std::string cudaLinkLanguage = "CUDA";
+  std::string const objExt =
+    this->Makefile->GetSafeDefinition("CMAKE_CUDA_OUTPUT_EXTENSION");
 
   std::string const cfgName = this->GetConfigName();
-  std::string const targetOutputReal =
-    ConvertToNinjaPath(genTarget.ObjectDirectory + "cmake_device_link.o");
+  std::string const targetOutputReal = ConvertToNinjaPath(
+    genTarget.ObjectDirectory + "cmake_device_link" + objExt);
 
   std::string const targetOutputImplib =
     ConvertToNinjaPath(genTarget.GetFullPath(cfgName,
@@ -714,6 +717,8 @@ void cmNinjaNormalTargetGenerator::WriteDeviceLinkStatement()
     this->ConvertToNinjaPath(objPath), cmOutputConverter::SHELL);
   EnsureDirectoryExists(objPath);
 
+  this->SetMsvcTargetPdbVariable(vars);
+
   if (this->GetGlobalGenerator()->IsGCCOnWindows()) {
     // ar.exe can't handle backslashes in rsp files (implicitly used by gcc)
     std::string& linkLibraries = vars["LINK_LIBRARIES"];
diff --git a/Source/cmNinjaTargetGenerator.cxx b/Source/cmNinjaTargetGenerator.cxx
index 23caead..8ad2efe 100644
--- a/Source/cmNinjaTargetGenerator.cxx
+++ b/Source/cmNinjaTargetGenerator.cxx
@@ -344,7 +344,8 @@ bool cmNinjaTargetGenerator::SetMsvcTargetPdbVariable(cmNinjaVars& vars) const
 {
   cmMakefile* mf = this->GetMakefile();
   if (mf->GetDefinition("MSVC_C_ARCHITECTURE_ID") ||
-      mf->GetDefinition("MSVC_CXX_ARCHITECTURE_ID")) {
+      mf->GetDefinition("MSVC_CXX_ARCHITECTURE_ID") ||
+      mf->GetDefinition("MSVC_CUDA_ARCHITECTURE_ID")) {
     std::string pdbPath;
     std::string compilePdbPath = this->ComputeTargetCompilePDB();
     if (this->GeneratorTarget->GetType() == cmStateEnums::EXECUTABLE ||
diff --git a/Tests/CMakeTests/ImplicitLinkInfoTest.cmake.in b/Tests/CMakeTests/ImplicitLinkInfoTest.cmake.in
index 1313dbf..d6d2357 100644
--- a/Tests/CMakeTests/ImplicitLinkInfoTest.cmake.in
+++ b/Tests/CMakeTests/ImplicitLinkInfoTest.cmake.in
@@ -518,6 +518,13 @@ set(msys_g77_libs "frtbegin;g2c;mingw32;moldname;mingwex;msvcrt;user32;kernel32;
 set(msys_g77_dirs "C:/some-mingw/lib/gcc/mingw32/3.4.5;C:/some-mingw/lib/gcc;/some-mingw/lib;C:/some-mingw/lib")
 list(APPEND platforms msys_g77)
 
+#-----------------------------------------------------------------------------
+# MSVC from NVIDIA CUDA
+
+set(nvcc_msvc_text [[cuda-fake-ld cl.exe -nologo "tmp/a_dlink.obj" "tmp/CMakeCUDACompilerId.obj" -link -INCREMENTAL:NO   "/LIBPATH:C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/bin/../lib/x64" cudadevrt.lib  cudart_static.lib -Fe"a.exe"]])
+set(nvcc_msvc_libs "cudadevrt.lib;cudart_static.lib")
+set(nvcc_msvc_dirs "C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0/lib/x64")
+list(APPEND platforms nvcc_msvc)
 
 #-----------------------------------------------------------------------------
 # Test parsing for all above examples.
diff --git a/Tests/Cuda/Complex/dynamic.cpp b/Tests/Cuda/Complex/dynamic.cpp
index d579f1e..3848ce7 100644
--- a/Tests/Cuda/Complex/dynamic.cpp
+++ b/Tests/Cuda/Complex/dynamic.cpp
@@ -1,5 +1,11 @@
 
-int dynamic_base_func(int x)
+#ifdef _WIN32
+#define EXPORT __declspec(dllexport)
+#else
+#define EXPORT
+#endif
+
+EXPORT int dynamic_base_func(int x)
 {
   return x * x;
 }
diff --git a/Tests/Cuda/Complex/dynamic.cu b/Tests/Cuda/Complex/dynamic.cu
index ea52acb..a23dc25 100644
--- a/Tests/Cuda/Complex/dynamic.cu
+++ b/Tests/Cuda/Complex/dynamic.cu
@@ -3,9 +3,15 @@
 #include <cuda.h>
 #include <iostream>
 
+#ifdef _WIN32
+#define EXPORT __declspec(dllexport)
+#else
+#define EXPORT
+#endif
+
 int dynamic_base_func(int);
 
-int __host__ cuda_dynamic_host_func(int x)
+EXPORT int __host__ cuda_dynamic_host_func(int x)
 {
   return dynamic_base_func(x);
 }
@@ -16,7 +22,7 @@ void DetermineIfValidCudaDevice()
 {
 }
 
-void cuda_dynamic_lib_func()
+EXPORT void cuda_dynamic_lib_func()
 {
   DetermineIfValidCudaDevice <<<1,1>>> ();
   cudaError_t err = cudaGetLastError();
diff --git a/Tests/Cuda/Complex/main.cpp b/Tests/Cuda/Complex/main.cpp
index 32312d0..5a3f820 100644
--- a/Tests/Cuda/Complex/main.cpp
+++ b/Tests/Cuda/Complex/main.cpp
@@ -3,8 +3,14 @@
 #include "file1.h"
 #include "file2.h"
 
-int call_cuda_seperable_code(int x);
-int mixed_launch_kernel(int x);
+#ifdef _WIN32
+#define IMPORT __declspec(dllimport)
+#else
+#define IMPORT
+#endif
+
+IMPORT int call_cuda_seperable_code(int x);
+IMPORT int mixed_launch_kernel(int x);
 
 int main(int argc, char** argv)
 {
diff --git a/Tests/Cuda/Complex/mixed.cpp b/Tests/Cuda/Complex/mixed.cpp
index 205f091..bd32e51 100644
--- a/Tests/Cuda/Complex/mixed.cpp
+++ b/Tests/Cuda/Complex/mixed.cpp
@@ -1,6 +1,14 @@
 
+#ifdef _WIN32
+#define EXPORT __declspec(dllexport)
+#define IMPORT __declspec(dllimport)
+#else
+#define EXPORT
+#define IMPORT
+#endif
+
 int dynamic_base_func(int);
-int cuda_dynamic_host_func(int);
+IMPORT int cuda_dynamic_host_func(int);
 int file3_launch_kernel(int);
 
 int dynamic_final_func(int x)
@@ -8,7 +16,7 @@ int dynamic_final_func(int x)
   return cuda_dynamic_host_func(dynamic_base_func(x));
 }
 
-int call_cuda_seperable_code(int x)
+EXPORT int call_cuda_seperable_code(int x)
 {
   return file3_launch_kernel(x);
 }
diff --git a/Tests/Cuda/Complex/mixed.cu b/Tests/Cuda/Complex/mixed.cu
index 45b412f..7051de0 100644
--- a/Tests/Cuda/Complex/mixed.cu
+++ b/Tests/Cuda/Complex/mixed.cu
@@ -4,10 +4,18 @@
 #include "file1.h"
 #include "file2.h"
 
+#ifdef _WIN32
+#define EXPORT __declspec(dllexport)
+#define IMPORT __declspec(dllimport)
+#else
+#define EXPORT
+#define IMPORT
+#endif
+
 result_type __device__ file1_func(int x);
 result_type_dynamic __device__ file2_func(int x);
 
-void __host__ cuda_dynamic_lib_func();
+IMPORT void __host__ cuda_dynamic_lib_func();
 
 static
 __global__
@@ -17,7 +25,7 @@ void mixed_kernel(result_type& r, int x)
   result_type_dynamic rd = file2_func(x);
 }
 
-int mixed_launch_kernel(int x)
+EXPORT int mixed_launch_kernel(int x)
 {
   cuda_dynamic_lib_func();
 
diff --git a/Tests/CudaOnly/EnableStandard/main.cu b/Tests/CudaOnly/EnableStandard/main.cu
index 83e9dfd..f7144e6 100644
--- a/Tests/CudaOnly/EnableStandard/main.cu
+++ b/Tests/CudaOnly/EnableStandard/main.cu
@@ -1,8 +1,14 @@
 
 #include <iostream>
 
+#ifdef _WIN32
+#define IMPORT __declspec(dllimport)
+#else
+#define IMPORT
+#endif
+
 int static_cuda11_func(int);
-int shared_cuda11_func(int);
+IMPORT int shared_cuda11_func(int);
 
 void test_functions()
 {
diff --git a/Tests/CudaOnly/EnableStandard/shared.cu b/Tests/CudaOnly/EnableStandard/shared.cu
index 28555b3..ccdd0b2 100644
--- a/Tests/CudaOnly/EnableStandard/shared.cu
+++ b/Tests/CudaOnly/EnableStandard/shared.cu
@@ -1,9 +1,15 @@
 
 #include <type_traits>
 
+#ifdef _WIN32
+#define EXPORT __declspec(dllexport)
+#else
+#define EXPORT
+#endif
+
 using tt = std::true_type;
 using ft = std::false_type;
-int __host__ shared_cuda11_func(int x)
+EXPORT int __host__ shared_cuda11_func(int x)
 {
   return x * x + std::integral_constant<int, 17>::value;
 }
diff --git a/Tests/CudaOnly/SeparateCompilation/main.cu b/Tests/CudaOnly/SeparateCompilation/main.cu
index d4520ae..5c8e150 100644
--- a/Tests/CudaOnly/SeparateCompilation/main.cu
+++ b/Tests/CudaOnly/SeparateCompilation/main.cu
@@ -4,12 +4,12 @@
 #include "file1.h"
 #include "file2.h"
 
-// result_type file4_launch_kernel(int x);
-// result_type file5_launch_kernel(int x);
+int file4_launch_kernel(int x);
+int file5_launch_kernel(int x);
 
 int main(int argc, char **argv)
 {
-  // file4_launch_kernel(42);
-  // file5_launch_kernel(42);
+  file4_launch_kernel(42);
+  file5_launch_kernel(42);
   return 0;
 }
diff --git a/Tests/CudaOnly/WithDefs/CMakeLists.txt b/Tests/CudaOnly/WithDefs/CMakeLists.txt
index c4ca8b9..6c4011c 100644
--- a/Tests/CudaOnly/WithDefs/CMakeLists.txt
+++ b/Tests/CudaOnly/WithDefs/CMakeLists.txt
@@ -4,7 +4,12 @@ project (CudaOnlyWithDefs CUDA)
 
 #verify that we can pass explicit cuda arch flags
 set(CMAKE_CUDA_FLAGS "-gencode arch=compute_30,code=compute_30")
-set(debug_compile_flags --generate-code arch=compute_20,code=sm_20 -Xcompiler=-Werror)
+set(debug_compile_flags --generate-code arch=compute_20,code=sm_20)
+if(CMAKE_CUDA_SIMULATE_ID STREQUAL "MSVC")
+  list(APPEND debug_compile_flags -Xcompiler=-WX)
+else()
+  list(APPEND debug_compile_flags -Xcompiler=-Werror)
+endif()
 set(release_compile_defs DEFREL)
 
 #Goal for this example:
diff --git a/Tests/CudaOnly/WithDefs/main.notcu b/Tests/CudaOnly/WithDefs/main.notcu
index 33a49d2..67bf10c 100644
--- a/Tests/CudaOnly/WithDefs/main.notcu
+++ b/Tests/CudaOnly/WithDefs/main.notcu
@@ -2,12 +2,21 @@
 #include <cuda_runtime.h>
 #include <iostream>
 
+#ifndef PACKED_DEFINE
+#error "PACKED_DEFINE not defined!"
+#endif
+
 static
 __global__
 void DetermineIfValidCudaDevice()
 {
 }
 
+#ifdef _MSC_VER
+#pragma pack(push,1)
+#undef PACKED_DEFINE
+#define PACKED_DEFINE
+#endif
 struct PACKED_DEFINE result_type
 {
   bool valid;
@@ -16,6 +25,9 @@ struct PACKED_DEFINE result_type
 #error missing DEFREL flag
 #endif
 };
+#ifdef _MSC_VER
+#pragma pack(pop)
+#endif
 
 result_type can_launch_kernel()
 {
diff --git a/Tests/RunCMake/try_compile/CudaStandardNoDefault.cmake b/Tests/RunCMake/try_compile/CudaStandardNoDefault.cmake
new file mode 100644
index 0000000..ea22152
--- /dev/null
+++ b/Tests/RunCMake/try_compile/CudaStandardNoDefault.cmake
@@ -0,0 +1,9 @@
+enable_language(CUDA)
+try_compile(result ${CMAKE_CURRENT_BINARY_DIR}
+  SOURCES ${CMAKE_CURRENT_SOURCE_DIR}/src.cu
+  CUDA_STANDARD 3 # bogus, but not used
+  OUTPUT_VARIABLE out
+  )
+if(NOT result)
+  message(FATAL_ERROR "try_compile failed:\n${out}")
+endif()
diff --git a/Tests/RunCMake/try_compile/RunCMakeTest.cmake b/Tests/RunCMake/try_compile/RunCMakeTest.cmake
index 5452e6d..6a1bc64 100644
--- a/Tests/RunCMake/try_compile/RunCMakeTest.cmake
+++ b/Tests/RunCMake/try_compile/RunCMakeTest.cmake
@@ -36,7 +36,11 @@ elseif(DEFINED CMAKE_CXX_STANDARD_DEFAULT)
   run_cmake(CxxStandardNoDefault)
 endif()
 if(CMake_TEST_CUDA)
-  run_cmake(CudaStandard)
+  if(CMAKE_HOST_WIN32)
+    run_cmake(CudaStandardNoDefault)
+  else()
+    run_cmake(CudaStandard)
+  endif()
 endif()
 if(CMAKE_C_COMPILER_ID STREQUAL "GNU" AND NOT CMAKE_C_COMPILER_VERSION VERSION_LESS 4.4)
   run_cmake(CStandardGNU)