Merge topic 'cuda_propagate_flags_when_dlink_on_executables'

8d1f9e5b CUDA: Now pass correct FLAGS when device link cuda executables.
dc5051f1 CUDA: Test that CUDA flags are used when device linking executables.

13 files changed, 75 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_COMPILER> <CMAKE_CUDA_LINK_FLAGS> ${CMAKE_CUDA_COMPILE_OPTIONS_PIC} -shared -dlink <OBJECTS> -o <TARGET> <LINK_LIBRARIES>")
+    "<CMAKE_CUDA_COMPILER> <FLAGS> <CMAKE_CUDA_LINK_FLAGS> ${CMAKE_CUDA_COMPILE_OPTIONS_PIC} -shared -dlink <OBJECTS> -o <TARGET> <LINK_LIBRARIES>")
 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
   "<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")
+  "<CMAKE_CUDA_COMPILER> <FLAGS> <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")
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: ",
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 <iostream>
+
+#include "file1.h"
+
+int main(int argc, char** argv)
+{
+  return 0;
+}