diff options
20 files changed, 323 insertions, 50 deletions
diff --git a/Help/prop_tgt/CUDA_RESOLVE_DEVICE_SYMBOLS.rst b/Help/prop_tgt/CUDA_RESOLVE_DEVICE_SYMBOLS.rst index 127d79f..ef74ae2 100644 --- a/Help/prop_tgt/CUDA_RESOLVE_DEVICE_SYMBOLS.rst +++ b/Help/prop_tgt/CUDA_RESOLVE_DEVICE_SYMBOLS.rst @@ -1,12 +1,18 @@ CUDA_RESOLVE_DEVICE_SYMBOLS --------------------------- -CUDA only: Enables device linking for the specific static library target +CUDA only: Enables device linking for the specific library target -If set this will enable device linking on this static library target. Normally +If set this will enable device linking on the library target. Normally device linking is deferred until a shared library or executable is generated, allowing for multiple static libraries to resolve device symbols at the same -time. +time when they are used by a shared library or executable. + +By default static library targets have this property is disabled, +while shared, module, and executable targets have this property enabled. + +Note that device linking is not supported for :ref:`Object Libraries`. + For instance: diff --git a/Help/release/dev/CUDA_RESOLVE_DEVICE_SYMBOLS.rst b/Help/release/dev/CUDA_RESOLVE_DEVICE_SYMBOLS.rst new file mode 100644 index 0000000..32db233 --- /dev/null +++ b/Help/release/dev/CUDA_RESOLVE_DEVICE_SYMBOLS.rst @@ -0,0 +1,6 @@ +CUDA_RESOLVE_DEVICE_SYMBOLS +--------------------------- + +* The :prop_tgt:`CUDA_RESOLVE_DEVICE_SYMBOLS` target property is now supported + on shared library, module library, and executable targets. Previously it was + only honored on static libraries. diff --git a/Source/CMakeVersion.cmake b/Source/CMakeVersion.cmake index 39a5b3d..324d825 100644 --- a/Source/CMakeVersion.cmake +++ b/Source/CMakeVersion.cmake @@ -1,5 +1,5 @@ # CMake version number components. set(CMake_VERSION_MAJOR 3) set(CMake_VERSION_MINOR 13) -set(CMake_VERSION_PATCH 20190205) +set(CMake_VERSION_PATCH 20190206) #set(CMake_VERSION_RC 1) diff --git a/Source/cmMakefileExecutableTargetGenerator.cxx b/Source/cmMakefileExecutableTargetGenerator.cxx index e576b5f..e8ae5ae 100644 --- a/Source/cmMakefileExecutableTargetGenerator.cxx +++ b/Source/cmMakefileExecutableTargetGenerator.cxx @@ -95,7 +95,13 @@ void cmMakefileExecutableTargetGenerator::WriteDeviceExecutableRule( const bool hasCUDA = (std::find(closure->Languages.begin(), closure->Languages.end(), cuda_lang) != closure->Languages.end()); - if (!hasCUDA) { + + bool doDeviceLinking = true; + if (const char* resolveDeviceSymbols = + this->GeneratorTarget->GetProperty("CUDA_RESOLVE_DEVICE_SYMBOLS")) { + doDeviceLinking = cmSystemTools::IsOn(resolveDeviceSymbols); + } + if (!hasCUDA || !doDeviceLinking) { return; } diff --git a/Source/cmMakefileLibraryTargetGenerator.cxx b/Source/cmMakefileLibraryTargetGenerator.cxx index 72181ab..5a1ef4e 100644 --- a/Source/cmMakefileLibraryTargetGenerator.cxx +++ b/Source/cmMakefileLibraryTargetGenerator.cxx @@ -133,9 +133,12 @@ void cmMakefileLibraryTargetGenerator::WriteStaticLibraryRules() (std::find(closure->Languages.begin(), closure->Languages.end(), cuda_lang) != closure->Languages.end()); - const bool resolveDeviceSymbols = - this->GeneratorTarget->GetPropertyAsBool("CUDA_RESOLVE_DEVICE_SYMBOLS"); - if (hasCUDA && resolveDeviceSymbols) { + bool doDeviceLinking = false; + if (const char* resolveDeviceSymbols = + this->GeneratorTarget->GetProperty("CUDA_RESOLVE_DEVICE_SYMBOLS")) { + doDeviceLinking = cmSystemTools::IsOn(resolveDeviceSymbols); + } + if (hasCUDA && doDeviceLinking) { std::string linkRuleVar = "CMAKE_CUDA_DEVICE_LINK_LIBRARY"; this->WriteDeviceLibraryRules(linkRuleVar, false); } @@ -168,7 +171,12 @@ void cmMakefileLibraryTargetGenerator::WriteSharedLibraryRules(bool relink) const bool hasCUDA = (std::find(closure->Languages.begin(), closure->Languages.end(), cuda_lang) != closure->Languages.end()); - if (hasCUDA) { + bool doDeviceLinking = true; + if (const char* resolveDeviceSymbols = + this->GeneratorTarget->GetProperty("CUDA_RESOLVE_DEVICE_SYMBOLS")) { + doDeviceLinking = cmSystemTools::IsOn(resolveDeviceSymbols); + } + if (hasCUDA && doDeviceLinking) { std::string linkRuleVar = "CMAKE_CUDA_DEVICE_LINK_LIBRARY"; this->WriteDeviceLibraryRules(linkRuleVar, relink); } @@ -209,7 +217,12 @@ void cmMakefileLibraryTargetGenerator::WriteModuleLibraryRules(bool relink) const bool hasCUDA = (std::find(closure->Languages.begin(), closure->Languages.end(), cuda_lang) != closure->Languages.end()); - if (hasCUDA) { + bool doDeviceLinking = true; + if (const char* resolveDeviceSymbols = + this->GeneratorTarget->GetProperty("CUDA_RESOLVE_DEVICE_SYMBOLS")) { + doDeviceLinking = cmSystemTools::IsOn(resolveDeviceSymbols); + } + if (hasCUDA && doDeviceLinking) { std::string linkRuleVar = "CMAKE_CUDA_DEVICE_LINK_LIBRARY"; this->WriteDeviceLibraryRules(linkRuleVar, relink); } diff --git a/Source/cmNinjaNormalTargetGenerator.cxx b/Source/cmNinjaNormalTargetGenerator.cxx index 0d05782..cbc0103 100644 --- a/Source/cmNinjaNormalTargetGenerator.cxx +++ b/Source/cmNinjaNormalTargetGenerator.cxx @@ -566,22 +566,23 @@ void cmNinjaNormalTargetGenerator::WriteDeviceLinkStatement() (std::find(closure->Languages.begin(), closure->Languages.end(), cuda_lang) != closure->Languages.end()); - bool shouldHaveDeviceLinking = false; - switch (genTarget.GetType()) { - case cmStateEnums::SHARED_LIBRARY: - case cmStateEnums::MODULE_LIBRARY: - case cmStateEnums::EXECUTABLE: - shouldHaveDeviceLinking = true; - break; - case cmStateEnums::STATIC_LIBRARY: - shouldHaveDeviceLinking = - genTarget.GetPropertyAsBool("CUDA_RESOLVE_DEVICE_SYMBOLS"); - break; - default: - break; + bool doDeviceLinking = false; + if (const char* resolveDeviceSymbols = + genTarget.GetProperty("CUDA_RESOLVE_DEVICE_SYMBOLS")) { + doDeviceLinking = cmSystemTools::IsOn(resolveDeviceSymbols); + } else { + switch (genTarget.GetType()) { + case cmStateEnums::SHARED_LIBRARY: + case cmStateEnums::MODULE_LIBRARY: + case cmStateEnums::EXECUTABLE: + doDeviceLinking = true; + break; + default: + break; + } } - if (!(shouldHaveDeviceLinking && hasCUDA)) { + if (!(doDeviceLinking && hasCUDA)) { return; } diff --git a/Source/cmQtAutoGeneratorMocUic.cxx b/Source/cmQtAutoGeneratorMocUic.cxx index 0ba5224..ddff4cf 100644 --- a/Source/cmQtAutoGeneratorMocUic.cxx +++ b/Source/cmQtAutoGeneratorMocUic.cxx @@ -678,19 +678,21 @@ void cmQtAutoGeneratorMocUic::JobMocT::Process(WorkerT& wrk) BuildFile += '/'; BuildFile += IncludeString; } else { - std::string rel = wrk.FileSys().GetFilePathChecksum(SourceFile); - rel += "/moc_"; - rel += wrk.FileSys().GetFilenameWithoutLastExtension(SourceFile); - rel += ".cpp"; - // Register relative file path - wrk.Gen().ParallelMocAutoRegister(rel); + // Relative build path + std::string relPath = wrk.FileSys().GetFilePathChecksum(SourceFile); + relPath += "/moc_"; + relPath += wrk.FileSys().GetFilenameWithoutLastExtension(SourceFile); + + // Register relative file path with duplication check + relPath = wrk.Gen().ParallelMocAutoRegister(relPath); + // Absolute build path if (wrk.Base().MultiConfig) { BuildFile = wrk.Base().AutogenIncludeDir; BuildFile += '/'; - BuildFile += rel; + BuildFile += relPath; } else { - BuildFile = wrk.Base().AbsoluteBuildPath(rel); + BuildFile = wrk.Base().AbsoluteBuildPath(relPath); } } @@ -1953,11 +1955,31 @@ bool cmQtAutoGeneratorMocUic::ParallelMocIncluded( return (MocIncludedFiles_.find(sourceFile) != MocIncludedFiles_.end()); } -void cmQtAutoGeneratorMocUic::ParallelMocAutoRegister( - std::string const& mocFile) +std::string cmQtAutoGeneratorMocUic::ParallelMocAutoRegister( + std::string const& baseName) { - std::lock_guard<std::mutex> mocLock(JobsMutex_); - MocAutoFiles_.emplace(mocFile); + std::string res; + { + std::lock_guard<std::mutex> mocLock(JobsMutex_); + res = baseName; + res += ".cpp"; + if (MocAutoFiles_.find(res) == MocAutoFiles_.end()) { + MocAutoFiles_.emplace(res); + } else { + // Append number suffix to the file name + for (unsigned int ii = 2; ii != 1024; ++ii) { + res = baseName; + res += '_'; + res += std::to_string(ii); + res += ".cpp"; + if (MocAutoFiles_.find(res) == MocAutoFiles_.end()) { + MocAutoFiles_.emplace(res); + break; + } + } + } + } + return res; } void cmQtAutoGeneratorMocUic::ParallelMocAutoUpdated() diff --git a/Source/cmQtAutoGeneratorMocUic.h b/Source/cmQtAutoGeneratorMocUic.h index 32a6006..c22df29 100644 --- a/Source/cmQtAutoGeneratorMocUic.h +++ b/Source/cmQtAutoGeneratorMocUic.h @@ -389,7 +389,7 @@ public: bool ParallelJobPushMoc(JobHandleT& jobHandle); bool ParallelJobPushUic(JobHandleT& jobHandle); bool ParallelMocIncluded(std::string const& sourceFile); - void ParallelMocAutoRegister(std::string const& mocFile); + std::string ParallelMocAutoRegister(std::string const& baseName); void ParallelMocAutoUpdated(); private: diff --git a/Source/cmVisualStudio10TargetGenerator.cxx b/Source/cmVisualStudio10TargetGenerator.cxx index 178e717..8e08417 100644 --- a/Source/cmVisualStudio10TargetGenerator.cxx +++ b/Source/cmVisualStudio10TargetGenerator.cxx @@ -2998,18 +2998,19 @@ bool cmVisualStudio10TargetGenerator::ComputeCudaLinkOptions( // Determine if we need to do a device link bool doDeviceLinking = false; - switch (this->GeneratorTarget->GetType()) { - case cmStateEnums::SHARED_LIBRARY: - case cmStateEnums::MODULE_LIBRARY: - case cmStateEnums::EXECUTABLE: - doDeviceLinking = true; - break; - case cmStateEnums::STATIC_LIBRARY: - doDeviceLinking = this->GeneratorTarget->GetPropertyAsBool( - "CUDA_RESOLVE_DEVICE_SYMBOLS"); - break; - default: - break; + if (const char* resolveDeviceSymbols = + this->GeneratorTarget->GetProperty("CUDA_RESOLVE_DEVICE_SYMBOLS")) { + doDeviceLinking = cmSystemTools::IsOn(resolveDeviceSymbols); + } else { + switch (this->GeneratorTarget->GetType()) { + case cmStateEnums::SHARED_LIBRARY: + case cmStateEnums::MODULE_LIBRARY: + case cmStateEnums::EXECUTABLE: + doDeviceLinking = true; + break; + default: + break; + } } cudaLinkOptions.AddFlag("PerformDeviceLink", diff --git a/Tests/CudaOnly/CMakeLists.txt b/Tests/CudaOnly/CMakeLists.txt index 9c4f86a..f1fd344 100644 --- a/Tests/CudaOnly/CMakeLists.txt +++ b/Tests/CudaOnly/CMakeLists.txt @@ -7,6 +7,17 @@ ADD_TEST_MACRO(CudaOnly.ResolveDeviceSymbols CudaOnlyResolveDeviceSymbols) ADD_TEST_MACRO(CudaOnly.SeparateCompilation CudaOnlySeparateCompilation) ADD_TEST_MACRO(CudaOnly.WithDefs CudaOnlyWithDefs) +add_test(NAME CudaOnly.DontResolveDeviceSymbols COMMAND + ${CMAKE_CTEST_COMMAND} -C $<CONFIGURATION> + --build-and-test + "${CMAKE_CURRENT_SOURCE_DIR}/DontResolveDeviceSymbols/" + "${CMAKE_CURRENT_BINARY_DIR}/DontResolveDeviceSymbols/" + ${build_generator_args} + --build-project DontResolveDeviceSymbols + --build-options ${build_options} + --test-command ${CMAKE_CTEST_COMMAND} -V -C $<CONFIGURATION> + ) + if(MSVC) ADD_TEST_MACRO(CudaOnly.PDB CudaOnlyPDB) endif() diff --git a/Tests/CudaOnly/DontResolveDeviceSymbols/CMakeLists.txt b/Tests/CudaOnly/DontResolveDeviceSymbols/CMakeLists.txt new file mode 100644 index 0000000..6190089 --- /dev/null +++ b/Tests/CudaOnly/DontResolveDeviceSymbols/CMakeLists.txt @@ -0,0 +1,50 @@ +cmake_minimum_required(VERSION 3.13) +project (DontResolveDeviceSymbols CUDA) + +# Find nm and dumpbin +if(CMAKE_NM) + set(dump_command ${CMAKE_NM}) + set(dump_args --defined-only) + set(symbol_name cudaRegisterLinkedBinary) +else() + include(GetPrerequisites) + message(STATUS "calling list_prerequisites to find dumpbin") + list_prerequisites("${CMAKE_COMMAND}" 0 0 0) + if(gp_dumpbin) + set(dump_command ${gp_dumpbin}) + set(dump_args /SYMBOLS) + set(symbol_name nv_fatb) + endif() +endif() + + +#Goal for this example: +# Build a static library that defines multiple methods and kernels that +# use each other. +# Don't resolve the device symbols in the static library +# Don't resolve the device symbols in the executable library +# Verify that we can't use those device symbols from anything +string(APPEND CMAKE_CUDA_FLAGS " -gencode arch=compute_30,code=[compute_30] -gencode arch=compute_50,code=\\\"compute_50\\\"") +set(CMAKE_CXX_STANDARD 11) +set(CMAKE_CUDA_STANDARD 11) + +add_library(CUDANoDeviceResolve SHARED file1.cu) +set_target_properties(CUDANoDeviceResolve + PROPERTIES + CUDA_SEPARABLE_COMPILATION ON + CUDA_RESOLVE_DEVICE_SYMBOLS OFF + POSITION_INDEPENDENT_CODE ON) +if(MSVC) + target_link_options(CUDANoDeviceResolve PRIVATE "/FORCE:UNRESOLVED") +endif() + +if(dump_command) +add_custom_command(TARGET CUDANoDeviceResolve POST_BUILD + COMMAND ${CMAKE_COMMAND} + -DDUMP_COMMAND=${dump_command} + -DDUMP_ARGS=${dump_args} + -DSYMBOL_NAME=${symbol_name} + -DTEST_LIBRARY_PATH=$<TARGET_FILE:CUDANoDeviceResolve> + -P ${CMAKE_CURRENT_SOURCE_DIR}/verify.cmake + ) +endif() diff --git a/Tests/CudaOnly/DontResolveDeviceSymbols/file1.cu b/Tests/CudaOnly/DontResolveDeviceSymbols/file1.cu new file mode 100644 index 0000000..3924f67 --- /dev/null +++ b/Tests/CudaOnly/DontResolveDeviceSymbols/file1.cu @@ -0,0 +1,69 @@ + +#include <iostream> + +static __global__ void file1_kernel(int in, int* out) +{ + *out = in * in; +} + +int choose_cuda_device() +{ + int nDevices = 0; + cudaError_t err = cudaGetDeviceCount(&nDevices); + if (err != cudaSuccess) { + std::cerr << "Failed to retrieve the number of CUDA enabled devices" + << std::endl; + return 1; + } + for (int i = 0; i < nDevices; ++i) { + cudaDeviceProp prop; + cudaError_t err = cudaGetDeviceProperties(&prop, i); + if (err != cudaSuccess) { + std::cerr << "Could not retrieve properties from CUDA device " << i + << std::endl; + return 1; + } + std::cout << "prop.major: " << prop.major << std::endl; + if (prop.major >= 3) { + err = cudaSetDevice(i); + if (err != cudaSuccess) { + std::cout << "Could not select CUDA device " << i << std::endl; + } else { + return 0; + } + } + } + + std::cout << "Could not find a CUDA enabled card supporting compute >=3.0" + << std::endl; + + return 1; +} + +int file1_launch_kernel() +{ + int ret = choose_cuda_device(); + if (ret) { + return 0; + } + + int input = 4; + + int* output; + cudaError_t err = cudaMallocManaged(&output, sizeof(int)); + cudaDeviceSynchronize(); + if (err != cudaSuccess) { + return 1; + } + + file1_kernel<<<1, 1>>>(input, output); + cudaDeviceSynchronize(); + err = cudaGetLastError(); + std::cout << err << " " << cudaGetErrorString(err) << std::endl; + if (err == cudaSuccess) { + // This kernel launch should failed as the device linking never occured + std::cerr << "file1_kernel: kernel launch should have failed" << std::endl; + return 1; + } + return 0; +} diff --git a/Tests/CudaOnly/DontResolveDeviceSymbols/main.cu b/Tests/CudaOnly/DontResolveDeviceSymbols/main.cu new file mode 100644 index 0000000..84a7a19 --- /dev/null +++ b/Tests/CudaOnly/DontResolveDeviceSymbols/main.cu @@ -0,0 +1,7 @@ + +#include <iostream> + +int main(int argc, char** argv) +{ + return 0; +} diff --git a/Tests/CudaOnly/DontResolveDeviceSymbols/verify.cmake b/Tests/CudaOnly/DontResolveDeviceSymbols/verify.cmake new file mode 100644 index 0000000..9bb426d --- /dev/null +++ b/Tests/CudaOnly/DontResolveDeviceSymbols/verify.cmake @@ -0,0 +1,14 @@ +execute_process(COMMAND ${DUMP_COMMAND} ${DUMP_ARGS} ${TEST_LIBRARY_PATH} + RESULT_VARIABLE RESULT + OUTPUT_VARIABLE OUTPUT + ERROR_VARIABLE ERROR +) + +if(NOT "${RESULT}" STREQUAL "0") + message(FATAL_ERROR "${DUMP_COMMAND} failed [${RESULT}] [${OUTPUT}] [${ERROR}]") +endif() + +if("${OUTPUT}" MATCHES "${SYMBOL_NAME}") + message(FATAL_ERROR + "The '${SYMBOL_NAME}' symbol is defined; device linking occurred!") +endif() diff --git a/Tests/QtAutogen/SameName/CMakeLists.txt b/Tests/QtAutogen/SameName/CMakeLists.txt index 6d42499..8d4f71f 100644 --- a/Tests/QtAutogen/SameName/CMakeLists.txt +++ b/Tests/QtAutogen/SameName/CMakeLists.txt @@ -17,6 +17,10 @@ add_executable(sameName ccc/item.cpp ccc/data.qrc item.cpp + object.h + object.h++ + object.hpp + object.hxx data.qrc main.cpp ) diff --git a/Tests/QtAutogen/SameName/main.cpp b/Tests/QtAutogen/SameName/main.cpp index a4ffcb3..92f15cd 100644 --- a/Tests/QtAutogen/SameName/main.cpp +++ b/Tests/QtAutogen/SameName/main.cpp @@ -3,14 +3,25 @@ #include "bbb/aaa/item.hpp" #include "bbb/item.hpp" #include "ccc/item.hpp" +#include "item.hpp" +#include "object.h" +#include "object.h++" +#include "object.hpp" +#include "object.hxx" int main(int argv, char** args) { - // Object instances + // Item instances + ::Item item; ::aaa::Item aaa_item; ::aaa::bbb::Item aaa_bbb_item; ::bbb::Item bbb_item; ::bbb::aaa::Item bbb_aaa_item; ::ccc::Item ccc_item; + // Object instances + ::Object_h obj_h; + ::Object_hplpl obj_hplpl; + ::Object_hpp obj_hpp; + ::Object_hxx obj_hxx; return 0; } diff --git a/Tests/QtAutogen/SameName/object.h b/Tests/QtAutogen/SameName/object.h new file mode 100644 index 0000000..8662094 --- /dev/null +++ b/Tests/QtAutogen/SameName/object.h @@ -0,0 +1,13 @@ +#ifndef OBJECT_H +#define OBJECT_H + +#include <QObject> + +class Object_h : public QObject +{ + Q_OBJECT + Q_SLOT + void go(){}; +}; + +#endif diff --git a/Tests/QtAutogen/SameName/object.h++ b/Tests/QtAutogen/SameName/object.h++ new file mode 100644 index 0000000..64222b7 --- /dev/null +++ b/Tests/QtAutogen/SameName/object.h++ @@ -0,0 +1,13 @@ +#ifndef OBJECT_HPLPL +#define OBJECT_HPLPL + +#include <QObject> + +class Object_hplpl : public QObject +{ + Q_OBJECT + Q_SLOT + void go(){}; +}; + +#endif diff --git a/Tests/QtAutogen/SameName/object.hpp b/Tests/QtAutogen/SameName/object.hpp new file mode 100644 index 0000000..035050e --- /dev/null +++ b/Tests/QtAutogen/SameName/object.hpp @@ -0,0 +1,13 @@ +#ifndef OBJECT_HPP +#define OBJECT_HPP + +#include <QObject> + +class Object_hpp : public QObject +{ + Q_OBJECT + Q_SLOT + void go(){}; +}; + +#endif diff --git a/Tests/QtAutogen/SameName/object.hxx b/Tests/QtAutogen/SameName/object.hxx new file mode 100644 index 0000000..c3c050f --- /dev/null +++ b/Tests/QtAutogen/SameName/object.hxx @@ -0,0 +1,13 @@ +#ifndef OBJECT_HXX +#define OBJECT_HXX + +#include <QObject> + +class Object_hxx : public QObject +{ + Q_OBJECT + Q_SLOT + void go(){}; +}; + +#endif |