Is it required that an entire .cu file must use syntax below C++17?

[Eric_Song]

Hi, guys,
It is said that CUDA-C++ compiling only supports the standard syntax up to C++17. Does this mean that the entire .cu file must use the syntax below C++17, or just __global__ and __device__ functions should be implemented in the syntax below C++17?

Your answer and guidance will be appreciated!

[Robert_Crovella]

It applies to host code as well in that file, for a few reasons:

1. The processing of a .cu file involves preprocessing steps that may not understand newer syntax.
2. The configured host compiler may or may not support newer syntax.
3. The nvcc command has no way to explicitly articulate to the host compiler that it should process a newer syntax, if that is required.

None of these reasons are absolutes. For example, it may be that the particular newer syntax you have does not trip up the preprocessing steps. It may be that the configured host compiler will support the newer syntax. It may be that you can use for example -Xcompiler with whatever switch your host compiler may need to accept the newer syntax successfully. So whatever you try may work.

But as far as what NVIDIA tests and supports for nvcc, support for anything newer than C++17 is not yet advertised for the latest CUDA 11.8. And it applies to host code in a .cu file.

The usual suggestions in these cases is to partition your “newer” C++ code into a separate .cpp file, and use wrapper functions as needed to connect the two. In your build process, do not use nvcc to process those .cpp files. Use the host compiler directly.

[Eric_Song]

Hi, @Robert_Crovella , thanks sincerely for your guidance.
I tried using a separate .cppfile to create a function, which is then referenced in a .cu file, but it finally failed. The failure reason seems to be that nvcc cannot recognize the C++20 syntax.
As you said, if the host compiler is used directly, then the .cu source should not be able to use any function from .cpp file of C++20.
Is that correct?

[Robert_Crovella]

The wrapper functions I referred to cannot use C++20 in their prototypes either, that is correct. If you do that, then there is no need for nvcc to recognize the C++20 syntax. It won’t be visible in the prototype.

And you can’t use a host compiler to compile a .cu source file.

[Eric_Song]

I think “using wrapper functions as needed to connect the two” seems a really hard task for me as a beginner in CUDA programming. I think I should tend to use C+±17 instead.
Appreciate your help very much!

[Robert_Crovella]

A wrapper function isn’t really unique or specific to CUDA. Here is an example. Using the [[likely]] and [[unlikely]] attributes is a C++20 feature.

If we wanted to isolate usage of that in a .cpp file, it could look like this:

f.cpp:

int f(bool b){
    if(b) [[likely]] {
        return 1;
    }
    else{
        return 8675309;
    }
}

test.cu:

int f(bool);

int main(){
  int a = f(true);
}

compile with:

g++ -c f.cpp
nvcc -o test test.cu f.o

The prototype of f does not contain any C++20 syntax. The implementation of f uses C++20 features.

(This isn’t a really great example, because this particular functionality can be handled by nvcc directly in some cases, even though it is a C++20 feature, for the reasons I previously indicated.
But my purpose is to indicate how wrapper functions may be used.)

[Eric_Song]

I think this code sample is awesome!
I have learned that the main idea is using different source files to isolate different syntax usages.
By the way, currently, I am using CMake to build my project. So, I guess I might be able to translate it as:

# `g++ -c f.cpp` to some CMake command I don't know yet,
add_executable(test test.cu f.o)

Am I right?

[Robert_Crovella]

I don’t have any advice for you with CMake. Perhaps others will know.

[Robert_Crovella]

FWIW, the recently released CUDA 12.0 advertises support for C++20.

[Eric_Song]

@Robert_Crovella

FWIW

Thanks sincerely for the latest information.