Suppressing warnings in CUDA

2020-12-08

Consider the following program:

#include <cmath>

__host__ constexpr double foo(double x){
  return std::sin(x);
}

__host__ __device__ double boo(double x, double y){

  return foo(x)+y;
}

__global__ void kern(){
  boo(3,5);
}

int main(){
  kern<<<1,20>>>();

  return 0;
}

Compiling this with CUDA 10.1

nvcc test.cu

gives some warnings

test.cu(9): warning: calling a constexpr __host__ function("foo") from a __host__ __device__ function("boo") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

test.cu(9): warning: calling a constexpr __host__ function from a __host__ __device__ function is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

test.cu(9): warning: calling a constexpr __host__ function("foo") from a __host__ __device__ function("boo") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

We could follow the compiler's advice for this simple program, but that isn't always possible and not every warning has an accompanying flag. So how do we suppress warnings in a more general way?

We can find their diagnostic error numbers.

To do so with CUDA 10.1 we compile as follows:

nvcc -Xcudafe --display_error_number test.cu

This gives:

test.cu(9): warning #2930-D: calling a constexpr __host__ function("foo") from a __host__ __device__ function("boo") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

test.cu(9): warning #2932-D: calling a constexpr __host__ function from a __host__ __device__ function is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

test.cu(9): warning #2930-D: calling a constexpr __host__ function("foo") from a __host__ __device__ function("boo") is not allowed. The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

We can now rewrite the program to suppress these warnings:

#include <cmath>

#pragma diag_suppress 2930,2932
__host__ constexpr double foo(double x){
  return std::sin(x);
}

__host__ __device__ double boo(double x, double y){
  return foo(x)+y;
}

__global__ void kern(){
  boo(3,5);
}

int main(){
  kern<<<1,20>>>();

  return 0;
}

In theory you can return the warning levels to what they should be using

#pragma diag_default 2930,2932 //Reset the warnings

But I couldn't find a place to put the reset that would still leave all the warnings suppressed.

Richbits