How to disable a specific nvcc compiler warnings

Question

I want to disable a specific compiler warning with nvcc, specifically

warning: NULL reference is not allowed

The code I am working on uses NULL references are part of SFINAE, so they can't be avoided.

An ideal solution would be a #pragma in just the source file where we want to disable the warnings, but a compiler flag would also be fine, if one exists to turn off only the warning in question.

Answer 1

It is actually possible to disable specific warnings on the device with NVCC. It took me ages to figure out how to do it.

You need to use the -Xcudafe flag combined with a token listed on this page. For example, to disable the "controlling expression is constant" warning, pass the following to NVCC:

-Xcudafe "--diag_suppress=boolean_controlling_expr_is_constant"

For other warnings, see the above link.

Answer 2

Just to add to the previous answer about -xcudafe (not enough reputation yet to leave a comment)

MAJOR EDIT:

cudaFE is apparently Nvidia's custom version of C++ Front End from the Edison Design Group. You can find the docs for it here: http://www.edg.com/docs/edg_cpp.pdf. I am currently referring to page numbers from the July 2019 v5.1 manual.

@einpoklum notes that just doing push/pop as I had originally stated in the initial post doesn't work and specifically that #pragma push generates a warning that it is ignored. I could not replicate the warning, but indeed in the test program below neither CUDA 10.1 nor CUDA 9.2 did the push/pop actually do anything (note that lines 20 and 22 generate no warning).

However, on page 75 of this manual they go through how to do localized diagnostic severity control without push/pop:

The following example suppresses the “pointless friend declaration” warning on the declaration of class A:

#pragma diag_suppress 522
class A { friend class A; };
#pragma diag_default 522
class B { friend class B; };

The #pragma diag_default returns the warning to default state. Another example would be:

#pragma diag_suppress = code_is_unreachable
...
#pragma diag_default = code_is_unreachable

The equal symbol is optional. Testing shows that this works and truly localized the severity control. Further, testing reveals that adding diagnostic suppressions in this way adds to previous diagnostic suppressions - it doesn’t replace. Also of note, in CUDA 10.1 unreachable code did not generate a warning while it did in CUDA 9.2. Finally, Page 77 of the manual mentions a new push/pop syntax:

#pragma push_macro(“identifier”)
#pragma pop_macro(“identifier”)

But I couldn't get it to work in the program below.

All of the above is tested in the program below, compiled with nvcc -std=c++14 test_warning_suppression.cu -o test_warning_suppression:

#include <cuda_runtime.h>

__host__ __device__ int return1(){
    int j = -1; //warning given for both CUDA 9.2 and 10.1
    return 1;
    if(false){ return 0; } //warning given here for CUDA 9.2
}

#pragma push
#pragma diag_suppress = code_is_unreachable 
#pragma diag_suppress = declared_but_not_referenced
__host__ __device__ int return2(){
    int j = -1;
    return 2;
    if(false){ return 0; }
}
#pragma pop

__host__ __device__ int return3(){
    int j = -1; //no warning given here
    return 3;
    if(false){ return 0; } //no warning here even in CUDA 9.2 
}

//push/pop did not localize warning suppression, so reset to default
#pragma diag_default = declared_but_not_referenced
#pragma diag_default = code_is_unreachable

//warning suppression localized to lines above by diag_default!
__host__ __device__ int return4(){
    int j = -1; //warning given for both CUDA 9.2 and 10.1
    return 4;
    if(false){ return 0; } //warning given here for CUDA 9.2
}

/* below does not work as of CUDA 10.1
#pragma push_macro(“identifier”)
#pragma diag_suppress = code_is_unreachable 
__device__ int return5(){
    return 5;
    if(false){ return 0; }
}
#pragma pop_macro(“identifier”)

__device__ int return6(){
    return 6;
    if(false){ return 0; }
} */

int main(){ return 0; }

Answer 3

To augment user2333829's answer: if you know the warning name you can disable it like this:

-Xcudafe "--diag_suppress=boolean_controlling_expr_is_constant"

If you don't know the name, get warning numbers by compiling with:

-Xcudafe --display_error_number

And then with:

-Xcudafe --diag_suppress=<warning_number>

(Note: both options at the same time apparently don't work.)

Answer 4

You can use w flag to suppress the warnings nvcc -w

Answer 5

I struggled to find a matching -Xcudafe for my warning. So here is another way.

You can pass a compiler flag to CL.exe that will disable a specific warning. For example, to disable the warnings about unchecked iterators you can pass /wd4996.

warning C4996: 'std::_Copy_impl': Function call with parameters that may be
unsafe - this call relies on the caller to check that the passed values are 
correct. To disable this warning, use -D_SCL_SECURE_NO_WARNINGS. See 
documentation on how to use Visual C++ 'Checked Iterators'

The tricky thing here is that by default the arguments from the host compiler settings are not passed to nvcc, so you need to add it via the CUDA C/C++ dialog.