7

I would like NVCC to treat the warning below as an error:

warning : calling a __host__ function("foo") from a __host__ __device__ function("bar")

NVCC documentation "NVIDIA CUDA Compiler Driver NVCC" doesn't even contain the word "warning".

talonmies
  • 70,661
  • 34
  • 192
  • 269
  • 1
    Why is this not an error by design? I just experienced nvcc compiling this successfully (with only the warning you mentioned): `__host__ int c() { return 0; } __host__ __device__ void b(){int a = c();} __global__ void a() {b();} /*...*/ a <<<1, 1 >>>();` and the line `a = c();` is turned into a read from 0: `mov.u32 %r1, 0; ld.volatile.u32 %r2, [%r1]; ` which can *never* work and was certainly not what I intended. Why should the compiler proceed with this? – masterxilo Jul 22 '16 at 20:38

1 Answers1

6

Quoting the CUDA COMPILER DRIVER NVCC reference guide, Section 3.2.8. "Generic Tool Options":

--Werror kind Make warnings of the specified kinds into errors. The following is the list of warning kinds accepted by this option:

cross-execution-space-call Be more strict about unsupported cross execution space calls. The compiler will generate an error instead of a warning for a call from a __host__ __device__ to a __host__ function.

Therefore, do the following:

Project -> Properties -> Configuration Properties -> CUDA C/C++ -> Command Line -> Additional Optics -> add --Werror cross-execution-space-call

This test program

#include <cuda.h>
#include <cuda_runtime.h>

void foo() { int a = 2;}

__host__ __device__ void test() {
    int tId = 1;
    foo();
}

int main(int argc, char **argv) { }

returns the following warning

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

without the above mentioned additional compilation option and returns the following error

Error   3   error : calling a __host__ function("foo") from a __host__ __device__ function("test") is not allowed

with the above mentioned additional compilation option.

Community
  • 1
  • 1
Vitality
  • 20,705
  • 4
  • 108
  • 146
  • 1
    I just hunted for a crash caused by nvcc compiling a `__host__ __device__` function calling a `__host__` function into nonsensical device code (namely an explicit read from 0: `mov.u32 %r31, 0; ld.u32 %r32, [%r31];`). Is there ever any reason not to use `--Werror cross-execution-space-call`? – masterxilo Jul 22 '16 at 20:35
  • Is there a case where "a call from a __host__ __device__ to a __host__ function" is not an error? It seems that any such call will crash. – alfC Apr 18 '21 at 08:56
  • Yes, a call from a `__host__ __device__` function to a `__host__` function is not illegal if the calling function is executing on the host. The compiler does not know at the point of compilation how it will be called, or even if it is ever called, in the objectionable configuration. Flagging some of these as definite errors and in other cases only warnings when it can't be ascertained can lead to a false sense of security "I didn't get an error, so...". If you'd like to see a change in CUDA behavior you can [file a bug](https://forums.developer.nvidia.com/t/how-to-report-a-bug/67911). – Robert Crovella Feb 15 '22 at 21:30