From dc5051f1c1b7604b9c05e04bdccdff222b69efa0 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 5 Jan 2017 14:21:23 -0500 Subject: CUDA: Test that CUDA flags are used when device linking executables. --- Tests/Cuda/CMakeLists.txt | 1 + Tests/Cuda/ProperLinkFlags/CMakeLists.txt | 20 ++++++++++++++++++++ Tests/Cuda/ProperLinkFlags/file1.cu | 11 +++++++++++ Tests/Cuda/ProperLinkFlags/file1.h | 7 +++++++ Tests/Cuda/ProperLinkFlags/main.cxx | 9 +++++++++ 5 files changed, 48 insertions(+) create mode 100644 Tests/Cuda/ProperLinkFlags/CMakeLists.txt create mode 100644 Tests/Cuda/ProperLinkFlags/file1.cu create mode 100644 Tests/Cuda/ProperLinkFlags/file1.h create mode 100644 Tests/Cuda/ProperLinkFlags/main.cxx diff --git a/Tests/Cuda/CMakeLists.txt b/Tests/Cuda/CMakeLists.txt index 5772fcf..40c9675 100644 --- a/Tests/Cuda/CMakeLists.txt +++ b/Tests/Cuda/CMakeLists.txt @@ -2,3 +2,4 @@ ADD_TEST_MACRO(Cuda.Complex CudaComplex) ADD_TEST_MACRO(Cuda.ConsumeCompileFeatures CudaConsumeCompileFeatures) ADD_TEST_MACRO(Cuda.ObjectLibrary CudaObjectLibrary) +ADD_TEST_MACRO(Cuda.ProperLinkFlags ProperLinkFlags) diff --git a/Tests/Cuda/ProperLinkFlags/CMakeLists.txt b/Tests/Cuda/ProperLinkFlags/CMakeLists.txt new file mode 100644 index 0000000..a9eb8bd --- /dev/null +++ b/Tests/Cuda/ProperLinkFlags/CMakeLists.txt @@ -0,0 +1,20 @@ + +cmake_minimum_required(VERSION 3.7) +project (ProperLinkFlags CUDA CXX) + +#Goal for this example: +#Verify that when we have CXX and CUDA enabled and we link an executable that +#has CUDA and CXX we use the CUDA link flags when doing the device link +#step + +#Specify a set of valid CUDA flags and an invalid set of CXX flags ( for CUDA ) +#to make sure we don't use the CXX flags when linking CUDA executables +set(CMAKE_CUDA_FLAGS "-arch=sm_35 --use_fast_math") +set(CMAKE_CXX_FLAGS "-Wall") + +set(CMAKE_CXX_STANDARD 11) +set(CMAKE_CUDA_STANDARD 11) +add_executable(ProperLinkFlags file1.cu main.cxx) + +set_target_properties( ProperLinkFlags + PROPERTIES CUDA_SEPARABLE_COMPILATION ON) diff --git a/Tests/Cuda/ProperLinkFlags/file1.cu b/Tests/Cuda/ProperLinkFlags/file1.cu new file mode 100644 index 0000000..d93dc9f --- /dev/null +++ b/Tests/Cuda/ProperLinkFlags/file1.cu @@ -0,0 +1,11 @@ + +#include "file1.h" + +result_type __device__ file1_func(int x) +{ + __ldg(&x); + result_type r; + r.input = x; + r.sum = x*x; + return r; +} diff --git a/Tests/Cuda/ProperLinkFlags/file1.h b/Tests/Cuda/ProperLinkFlags/file1.h new file mode 100644 index 0000000..ff1945c --- /dev/null +++ b/Tests/Cuda/ProperLinkFlags/file1.h @@ -0,0 +1,7 @@ + +#pragma once +struct result_type +{ + int input; + int sum; +}; diff --git a/Tests/Cuda/ProperLinkFlags/main.cxx b/Tests/Cuda/ProperLinkFlags/main.cxx new file mode 100644 index 0000000..7c0ee9e --- /dev/null +++ b/Tests/Cuda/ProperLinkFlags/main.cxx @@ -0,0 +1,9 @@ + +#include + +#include "file1.h" + +int main(int argc, char** argv) +{ + return 0; +} -- cgit v0.12 From 8d1f9e5b850e02d304e4c209e720f06e25837470 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 5 Jan 2017 16:31:36 -0500 Subject: CUDA: Now pass correct FLAGS when device link cuda executables. Previously we had a two issues when building cuda executables that required separable compilation. The first was that we didn't propagate FLAGS causing any -arch / -gencode flags to be dropped, and secondly generators such as ninja would use the CXX language flags instead of CUDA when the executable was mixed language. --- Modules/CMakeCUDACompiler.cmake.in | 2 ++ Modules/CMakeCUDAInformation.cmake | 2 +- Modules/Platform/Windows-NVIDIA-CUDA.cmake | 2 +- Source/cmLinkLineComputer.cxx | 6 ++++++ Source/cmLinkLineComputer.h | 4 ++++ Source/cmLinkLineDeviceComputer.cxx | 6 ++++++ Source/cmLinkLineDeviceComputer.h | 3 +++ Source/cmLocalGenerator.cxx | 5 ++++- 8 files changed, 27 insertions(+), 3 deletions(-) diff --git a/Modules/CMakeCUDACompiler.cmake.in b/Modules/CMakeCUDACompiler.cmake.in index 1ba42d9..7e8efa7 100644 --- a/Modules/CMakeCUDACompiler.cmake.in +++ b/Modules/CMakeCUDACompiler.cmake.in @@ -13,6 +13,8 @@ set(CMAKE_CUDA_HOST_COMPILER_ENV_VAR "CUDAHOSTCXX") set(CMAKE_CUDA_COMPILER_ID_RUN 1) set(CMAKE_CUDA_SOURCE_FILE_EXTENSIONS cu) +set(CMAKE_CUDA_LINKER_PREFERENCE 10) +set(CMAKE_CUDA_LINKER_PREFERENCE_PROPAGATES 1) set(CMAKE_CUDA_HOST_IMPLICIT_LINK_LIBRARIES "@CMAKE_CUDA_HOST_IMPLICIT_LINK_LIBRARIES@") set(CMAKE_CUDA_HOST_IMPLICIT_LINK_DIRECTORIES "@CMAKE_CUDA_HOST_IMPLICIT_LINK_DIRECTORIES@") diff --git a/Modules/CMakeCUDAInformation.cmake b/Modules/CMakeCUDAInformation.cmake index abc4b66..13b1789 100644 --- a/Modules/CMakeCUDAInformation.cmake +++ b/Modules/CMakeCUDAInformation.cmake @@ -187,7 +187,7 @@ if(NOT CMAKE_CUDA_DEVICE_LINK_LIBRARY) endif() if(NOT CMAKE_CUDA_DEVICE_LINK_EXECUTABLE) set(CMAKE_CUDA_DEVICE_LINK_EXECUTABLE - " ${CMAKE_CUDA_COMPILE_OPTIONS_PIC} -shared -dlink -o ") + " ${CMAKE_CUDA_COMPILE_OPTIONS_PIC} -shared -dlink -o ") endif() mark_as_advanced( diff --git a/Modules/Platform/Windows-NVIDIA-CUDA.cmake b/Modules/Platform/Windows-NVIDIA-CUDA.cmake index 80ecfbc..809ee06 100644 --- a/Modules/Platform/Windows-NVIDIA-CUDA.cmake +++ b/Modules/Platform/Windows-NVIDIA-CUDA.cmake @@ -33,7 +33,7 @@ unset(_CMAKE_VS_LINK_EXE) set(CMAKE_CUDA_DEVICE_LINK_LIBRARY " -shared -dlink -o -Xcompiler=-Fd,-FS") set(CMAKE_CUDA_DEVICE_LINK_EXECUTABLE - " -shared -dlink -o -Xcompiler=-Fd,-FS") + " -shared -dlink -o -Xcompiler=-Fd,-FS") string(APPEND CMAKE_CUDA_FLAGS_INIT " -Xcompiler=-GR,-EHsc") string(APPEND CMAKE_CUDA_FLAGS_DEBUG_INIT " -Xcompiler=-MDd,-Zi,-RTC1") diff --git a/Source/cmLinkLineComputer.cxx b/Source/cmLinkLineComputer.cxx index cf0cf88..e728632 100644 --- a/Source/cmLinkLineComputer.cxx +++ b/Source/cmLinkLineComputer.cxx @@ -184,3 +184,9 @@ std::string cmLinkLineComputer::ComputeLinkLibraries( return fout.str(); } + +std::string cmLinkLineComputer::GetLinkerLanguage(cmGeneratorTarget* target, + std::string const& config) +{ + return target->GetLinkerLanguage(config); +} diff --git a/Source/cmLinkLineComputer.h b/Source/cmLinkLineComputer.h index bb13717..57a70bc 100644 --- a/Source/cmLinkLineComputer.h +++ b/Source/cmLinkLineComputer.h @@ -11,6 +11,7 @@ #include "cmStateDirectory.h" class cmComputeLinkInformation; +class cmGeneratorTarget; class cmOutputConverter; class cmLinkLineComputer @@ -36,6 +37,9 @@ public: virtual std::string ComputeLinkLibraries(cmComputeLinkInformation& cli, std::string const& stdLibString); + virtual std::string GetLinkerLanguage(cmGeneratorTarget* target, + std::string const& config); + protected: std::string ComputeLinkLibs(cmComputeLinkInformation& cli); std::string ComputeRPath(cmComputeLinkInformation& cli); diff --git a/Source/cmLinkLineDeviceComputer.cxx b/Source/cmLinkLineDeviceComputer.cxx index 75e5ef5..6a700ff 100644 --- a/Source/cmLinkLineDeviceComputer.cxx +++ b/Source/cmLinkLineDeviceComputer.cxx @@ -59,6 +59,12 @@ std::string cmLinkLineDeviceComputer::ComputeLinkLibraries( return fout.str(); } +std::string cmLinkLineDeviceComputer::GetLinkerLanguage(cmGeneratorTarget*, + std::string const&) +{ + return "CUDA"; +} + cmNinjaLinkLineDeviceComputer::cmNinjaLinkLineDeviceComputer( cmOutputConverter* outputConverter, cmStateDirectory stateDir, cmGlobalNinjaGenerator const* gg) diff --git a/Source/cmLinkLineDeviceComputer.h b/Source/cmLinkLineDeviceComputer.h index d1079d7..f4bb3eb 100644 --- a/Source/cmLinkLineDeviceComputer.h +++ b/Source/cmLinkLineDeviceComputer.h @@ -17,6 +17,9 @@ public: std::string ComputeLinkLibraries(cmComputeLinkInformation& cli, std::string const& stdLibString) CM_OVERRIDE; + + std::string GetLinkerLanguage(cmGeneratorTarget* target, + std::string const& config) CM_OVERRIDE; }; class cmNinjaLinkLineDeviceComputer : public cmLinkLineDeviceComputer diff --git a/Source/cmLocalGenerator.cxx b/Source/cmLocalGenerator.cxx index ead1e72..44c390c 100644 --- a/Source/cmLocalGenerator.cxx +++ b/Source/cmLocalGenerator.cxx @@ -13,6 +13,7 @@ #include "cmInstallScriptGenerator.h" #include "cmInstallTargetGenerator.h" #include "cmLinkLineComputer.h" +#include "cmLinkLineDeviceComputer.h" #include "cmMakefile.h" #include "cmRulePlaceholderExpander.h" #include "cmSourceFile.h" @@ -979,7 +980,9 @@ void cmLocalGenerator::GetTargetFlags( linkFlags += this->Makefile->GetSafeDefinition(build); linkFlags += " "; } - std::string linkLanguage = target->GetLinkerLanguage(buildType); + + const std::string linkLanguage = + linkLineComputer->GetLinkerLanguage(target, buildType); if (linkLanguage.empty()) { cmSystemTools::Error( "CMake can not determine linker language for target: ", -- cgit v0.12