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.

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@")

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(

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

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);
+}

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

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)

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

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: ",

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)

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)

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;
+}

Tests/Cuda/ProperLinkFlags/file1.h

@@ -0,0 +1,7 @@
+
+#pragma once
+struct result_type
+{
+  int input;
+  int sum;
+};

Tests/Cuda/ProperLinkFlags/main.cxx

@@ -0,0 +1,9 @@
+
+#include <iostream>
+
+#include "file1.h"
+
+int main(int argc, char** argv)
+{
+  return 0;
+}