0

Please look at this code:

void bar() {}

__host__ __device__ void foo()
{
  bar();
}

__global__ void kernel()
{
  foo();
}

int main()
{
  kernel<<<1, 1>>>();
  gpuErrchk(cudaPeekAtLastError());

  gpuErrchk(cudaDeviceSynchronize());

  return 0;
}

I spent hours trying to solve the an illegal memory access was encountered runtime error. As it turned out, the reason is the bar() function - it's not declared as __device__. But! But the code compiles! It produces a warning, but compiles! The warning says:

warning: calling a __host__ function("bar") from a __host__ __device__
function("Test::foo") is not allowed

Since the compilation for my project produces a lot of output, I simply didn't see that warning. But if I remove the __device__ attribute from the foo() function, I get the expected error:

error: identifier "foo" is undefined in device code

The question is why the compiler prints only a warning and how to turn it into an error?

John Kugelman
  • 349,597
  • 67
  • 533
  • 578
nikitablack
  • 4,359
  • 2
  • 35
  • 68

2 Answers2

4

The question is why the compiler prints only a warning and how to turn it into an error?

The compiler prints only a warning because it doesn't know (at the point of compilation of the calling function) if the function will actually be called at runtime, in the objectionable configuration (i.e. on or from device code).

and how to turn it into an error?

From the nvcc manual you can add either:

-Werror all-warnings

to flag all warnings as errors

or

-Werror cross-execution-space-call

to only flag this type of warning as an error.

Also see here. To those who will ask why I didn't flag as a dupe, that other question doesn't include a question (or in the answer itself) about why the compiler behaves this way.

Robert Crovella
  • 143,785
  • 11
  • 213
  • 257
  • Thank you. `-Werror` flag works. But how is it that _the compiler doesn't know if the function will actually be called at runtime_? If it prints a **compilation** warning, then it knows. – nikitablack Feb 16 '22 at 09:02
  • 1
    The undecorated function, `bar`, is being called from `foo`. `foo` is decorated as `__host__ __device__`. That means `foo` can be called from host code or device code. When the compiler is compiling `foo`, it does not know if it will be called from host code, or device code, or both. It is acceptable for `foo` to be called from host code, because then `bar` runs as host code. All of that will work. The compiler does not know your intention at that point, so it issues a warning. Furthermore, when compiler is compiling `foo`, it does not know if `foo` will be called **at all**. – Robert Crovella Feb 16 '22 at 14:13
-3

I spent hours trying to solve the... error. ... But the code compiles! It produces a warning, but compiles!

You need to revisit your debugging methodology right there :-(

Any warning which you have not positively proven to yourself is immaterial - is where you need to look for your errors. And it is far easier and more rewarding to resolve warnings than to prove them invalid. (And by resolve, I mean address the underlying condition, not suppressing the warning, or const_cast'ing etc.)

So, don't turn warnings into errors with the compiler, turn them into essentially-errors in your mind. Clean, warning-free code = happy life.

einpoklum
  • 118,144
  • 57
  • 340
  • 684
  • I agree with you and this is what I'm trying to do with _my_ code. But there're some cases where it's impossible to achieve. For example, with a third-party lib. Or, as in my case, we're a hundred devs company working in a single repo, we have hundreds of thousands of lines of legacy code. – nikitablack Feb 16 '22 at 08:31
  • @nikitablack: Actually, if you are including the header of a third-party library, you can suppress the warnings it emits, with compiler-specific preprocessor directives before the inclusion; then drop the suppression after the inclusion. – einpoklum Feb 16 '22 at 09:45
  • Yes, if there are no macros and inline functions it will work. Also, it's not easy if you need to do that in every file for a large project or there are no specific compiler requirements. – nikitablack Feb 16 '22 at 10:12
  • 1
    @nikitablack: Fair enough. – einpoklum Feb 16 '22 at 11:02