As far as using nvcc
, one needs to use the corresponding gcc
(currently max. 5.4 I believe) in conjunction. This of course somewhat prevents one from using C++17 on the host side.
Since C++17 can be compiled using clang 5
and upwards (see here), and one can compile cuda code as well (see here), is it possible to use both C++17 and CUDA at the same time (or can there be problems, e.g. with the CUDA runtime)?
Clang currently supports CUDA 7.0 through 11.5. If clang detects a newer CUDA version, it will issue a warning and will attempt to use detected CUDA SDK it as if it were CUDA 11.5. Before you build CUDA code, you'll need to have installed the CUDA SDK.
All non-CUDA compilation steps are forwarded to a C++ host compiler that is supported by nvcc , and nvcc translates its options to appropriate host compiler command line options.
Yes, as you already guessed the CUDA clang frontend is indeed ahead in C++ feature support, even in device code. It was already in the past, introducing C++14 features before NVCC which was mostly unnoticed by the community.
Take this C++17, unnecessarily modified if constexpr
, snippet: Fibo
#include <cuda_runtime.h>
#include <cstdio>
constexpr unsigned
fibonacci(const unsigned x) {
if constexpr (false)
{
return 0u;
}
if( x <= 1 )
return 1;
return fibonacci(x - 1) + fibonacci(x - 2);
}
__global__
void k()
{
constexpr unsigned arg = fibonacci(5);
printf("%u", arg);
}
int main()
{
k<<<1,1>>>();
return 0;
}
It already runs with clang++ -std=c++17 -x cuda
: https://cuda.godbolt.org/z/GcIqeW
Nevertheless, for this specific example, C++17 extended lambdas and C++14 relaxed constexpr are that important in modern C++, that even in C++11 and C++14 mode of NVCC 8.0+ flags were added to enable those features already: https://devblogs.nvidia.com/new-compiler-features-cuda-8/
That means the above example compiles for example with NVCC 9.2 even without __device__
qualifiers when removing the demonstrating C++17 if constexpr
construct and adding -std=c++14 --expt-relaxed-constexpr
flags.
Here is a list about C++ standard support on the device side for nvcc
and clang -x cuda
: https://gist.github.com/ax3l/9489132#device-side-c-standard-support (NVCC 11.0 supports device-side C++17 now.)
Currently up to C++14 is supported in device code (Introduced in CUDA 9)
--std {c++03|c++11|c++14}
Options for Specifying Behavior of Compiler/Linker
However, if your host is only using C++17, it should be possible to use separate compilation and link them with library. Separate Compilation and Linking of CUDA C++ Device Code
Update: formatting and more info
If you love us? You can donate to us via Paypal or buy me a coffee so we can maintain and grow! Thank you!
Donate Us With