diff options
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) |