Improves the handling of CUDA layouts where we have multiple include
and library directories listed in the output of `nvcc -v`. This
updates both when the CUDA language is enabled or not.
Fixes: #24915
Due to an oversight in cmLinkLineDeviceComputer object
files did not get propagate to the device linking phase
when given via the `$<TARGET_OBJECTS>` generator expression.
Without `CudaLink` fields, MSBuild tries to perform a device link step
for object libraries with CUDA_SEPARABLE_COMPILATION on them even when
not needed. This commonly was encountered when generating PTX or
OPTIX-IR output.
Since commit fa45594407 (file(GET_RUNTIME_DEPENDENCIES): Preserve casing
for Windows PE binaries, 2023-03-06) we may find the cudart runtime
library with its on-disk casing instead of always lower-case.
Since commit 2def6a874b (CUDA: Add support for CUBIN, FATBIN, and
OPTIXIR compilation, 2023-01-27) added this test, it has failed when
executed with a CUDA toolkit that does not support the true native
architecture of the GPU. Detect this case and skip the test.
We need to revert this change as it can disable error messages
when compiling invalid CUDA code.
This reverts commit ea659b155d (CUDA: Always mark cuda toolkit as system
include, 2022-06-27, v3.25.0-rc1~269^2).
CUDA 11.6 added the `nvcc -arch=native` flag to automatically compile
for the host GPUs' architectures. Add support for specifying this
special `native` value in `CMAKE_CUDA_ARCHITECTURES` and
`CUDA_ARCHITECTURES`. During the compiler ABI detection step,
detect the native architectures so we can pass them explicitly
when using Clang or older versions of nvcc.
Fixes: #22375
The purpose of the test is to cover special values of
`CMAKE_CUDA_ARCHITECTURES`. Prepare to add more than
just `all` and `all-major`.
While at it, fix the `project()` name in the test to
match what we pass to `ctest` with `--build-project`.
Fix the test added by commit 61b9764b03 (CUDA: Allow both
CUDA_SEPARABLE_COMPILATION and CUDA_PTX_COMPILATION, 2021-09-17,
v3.23.0-rc1~516^2) to have a `project()` name matching what
is passed to CTest's `--build-project` option. Otherwise, on
VS builds, MSBuild warns:
MSBUILD : error MSB1009: Project file does not exist.
Other CUDA tests use a matching project name already.
Commit 14d8a276 (CUDA: Support nvcc 11.5 new -arch=all|all-major flags,
2021-08-17) added all and all-major options to CUDA_ARCHITECTURES. These are
fairly generic and likely to see real-world use by distributors. Thus it's
desirable to support these also for Clang and older NVCC versions.
The supported architectures are dependent on the toolkit version. We determine
the toolkit version prior to compiler detection. For NVCC we get the version
from the vendor identification output, but for Clang we need to invoke NVCC
separately.
The architecture information is mostly based on the Wikipedia list with the
earliest supported version being CUDA 7.0. This could be documented and
expanded in the future to allow projects to query CUDA toolkit version and
architecture information.
For Clang we additionally constrain based on its support.
Additionally the architecture mismatch detection logic is fixed, improved and
updated for generic support:
* Commit 01428c55 (CUDA: Fail fast if CMAKE_CUDA_ARCHITECTURES doesn't work
during detection, 2020-08-29) enabled CMAKE_CUDA_COMPILER_ID_REQUIRE_SUCCESS
if CMAKE_CUDA_ARCHITECTURES is specified. This results in
CMakeDetermineCompilerID.cmake printing the compiler error and our code for
presenting the mismatch in a user-friendly way being useless. The custom
logic seems preferable so go back to not enabling it.
* Commit 14d8a276 (CUDA: Support nvcc 11.5 new -arch=all|all-major flags,
2021-08-17) tried to support CMP0054 but forgot to add x to the interpolated
result. Thus the conditions would always evaluate to false. This is fixed as
a byproduct of removing NVIDIA specific checks, improving the error message
and replacing architectures_mode with a simpler architectures_explicit.
Visual Studio support omits testing the flags during detection due to
complexities in determining the toolkit version when using it.
A long-term proper implementation would be #23161.
Implements #22860.
From CUDA 9.0 to CUDA 11.4 the CUDA Visual Studio integration
defines omitted user defines from PTX generation.
With CUDA 11.5 this has been resolved, so we backport the fix
to allow for consistent behavior when using CMake
The target properties `CUDA_SEPARABLE_COMPILATION` and `CUDA_PTX_COMPILATION`
now aren't mutually exclusive and can now be used together on the same
target.
Seems the relative paths were wrong basically all around such that only
compiling files in the top-level directory would work. I've modified
CudaOnly.SeparateCompilation to cover this.
Fixes#22482.
With CUDAARCHS there's the possibility of tests being run with multiple
architectures or having CMAKE_CUDA_ARCHITECTURES set with NVCC despite CMP0104
being OLD.
Fix a few tests to work properly in such cases.
For NVCC the compiler takes care of device linking when passed the "-dlink"
flag.
Clang doesn't support such magic and requires the buildsystem to do the work
that NVCC does behind the scenes.
The implementation is based on Bazel's device linking documentation:
7cabcdf073/third_party/nccl/build_defs.bzl.tpl (L259)Closes: #20726
Architecture 30 was removed with CUDA 11, so most of the CUDA tests fail with
it.
Remove setting the architecture and bump the minimum version to 3.18, so
CMP0104 takes effect and we can rely on the default architecture, which is
guaranteed to be compilable.
Use of __ldg() in ProperLinkFlags was removed as it only affects performance
and is available only on sm_35 and above.
Testing the functionality of CUDA_ARCHITECTURES is already covered by
CudaOnly.Architecture and CudaOnly.CompileFlags.
The ability to disable adding architectures completely for packaging purposes
and cases requiring passing the architectures flags explicitly has been
requested.
Support a false value for CUDA_ARCHITECTURES and CMAKE_CUDA_ARCHITECTURES
for this purpose.
Implements #20821.
Since commit 0d0145138f (CUDA: Add abstraction for cuda runtime
selection, 2019-11-29, v3.17.0-rc1~83^2) we add CUDA runtime library
selection flags by default.
To maintain backwards compatibility the default CUDA runtime
library needs to be computed based on what libraries are found
on the initial compiler invocation. For example a toolchain
could establish initial flags that have all CUDA compilations
using the runtime version, and if we don't detect this we will
try to link to both the static and shared runtime.
Co-Author: Brad King <brad.king@kitware.com>
Fixes: #20708
In commit a653ca9504 (Tests: Update CUDA tests to work with Clang,
2020-03-27) some tests were conditioned using `CMAKE_CUDA_COMPILER_ID`.
That is not defined when configuring CMake itself, so it accidentally
turned off NVCC-specific CUDA tests altogether. Convert the conditions
to check `CMake_TEST_CUDA` for `Clang` instead. That option is added
explicitly to builds where we want the tests to run, so we can set it to
a value indicating the CUDA compiler vendor.
In commit a653ca9504 (Tests: Update CUDA tests to work with Clang,
2020-03-27) the NVCC-specific `CudaOnly.GPUDebugFlag` test was
accidentally broken by removing a space when appending `-G` to the CUDA
flags. This was covered by the test not running. Restore the space.
Fixes: #20727
Simplifies CUDA target architecture handling.
Required for Clang support as Clang doesn't automatically select a supported architecture.
We detect a supported architecture during compiler identification and set CMAKE_CUDA_ARCHITECTURES to it.
Introduces CMP0104 for backwards compatibility with manually setting code generation flags with NVCC.
Implements #17963.
Fixes#17559
Replace our hard-coded default of cudart=static with a first-class abstraction to select the runtime library from an enumeration of logical names.
