CUDA compilation fails due to unknown type name '__device__' separate compilation and linking solves the problem

[Quuxplusone]

Bugzilla Link	PR34379
Status	RESOLVED WONTFIX
Importance	P enhancement
Reported by	Lorenz Braun (lorenz.braun@ziti.uni-heidelberg.de)
Reported on	2017-08-30 06:45:14 -0700
Last modified on	2017-08-31 11:30:18 -0700
Version	5.0
Hardware	PC Linux
CC	hfinkel@anl.gov, llvm-bugs@lists.llvm.org, tra@google.com
Fixed by commit(s)
Attachments	streamcluster.tar.gz (21541 bytes, application/x-gzip)
Blocks
Blocked by
See also

Created attachment 19060
Streamcluster application with makefile to reproduce the error + log of the
compilation with error

The streamcluster application of the rodinia benchmark [0] cannot be compiled
with clang++. Compiling the source files in separate objects and linking them
afterwards works just fine.

Changing the order included headers might help.

Steps to reproduce:
1) extract tar
2) change path to cuda (i used cuda 8.0)
3) run make

Result: compilation fails

Expected Result: compilation is successful
(you can set NVCC to nvcc, then it works to)

Build-Date & Hardware:
2017-30-08
Linux octane 4.2.0-27-generic #32~14.04.1-Ubuntu SMP Fri Jan 22 15:32:26 UTC
2016 x86_64 x86_64 x86_64 GNU/Linux
CUDA 8.0
LLVM 5.0
not yet tested on trunk

[0] https://www.cs.virginia.edu/~skadron/wiki/rodinia/index.php/Main_Page

[Quuxplusone]

Attached streamcluster.tar.gz (21541 bytes, application/x-gzip): Streamcluster application with makefile to reproduce the error + log of the compilation with error

[Quuxplusone]

include of a CUDA file from inside a CPP file is not going to work as clang has no idea that the source will have CUDA in it.

If you want to do it this way, you'll need to explicitly tell clang to compile your top-level .cpp file as CUDA by passing '-x cuda' option.

[Quuxplusone]

(In reply to Artem Belevich from comment #1)
> #include of a CUDA file from inside a CPP file is not going to work as clang
> has no idea that the source will have CUDA in it.
>
> If you want to do it this way, you'll need to explicitly tell clang to
> compile your top-level .cpp file as CUDA by passing '-x cuda' option.

Should we add a better error message for this? Maybe something like:

#ifndef __CUDACC__
#error Say something useful.
#endif

[Quuxplusone]

I can't think of a good place for such a check. During C++ compilation we do not pre-include any CUDA headers, so we can't put the check there. We could always pre-include a special header with this check during C++ compilation, but I'm not sure whether it's worth it (all C++ compilations will have to carry the overhead, however small, just to catch a rare user error in non-C++ compilation).

We could add device and other key CUDA keywords in C++ mode and issue a warning or error if they are seen. Given that they all start with double-underscore, this should be OK to do. I think.

[Quuxplusone]

(In reply to Artem Belevich from comment #3)
> I can't think of a good place for such a check. During C++ compilation we do
> not pre-include any CUDA headers, so we can't put the check there. We could
> always pre-include a special header with this check during C++ compilation,
> but I'm not sure whether it's worth it (all C++ compilations will have to
> carry the overhead, however small, just to catch a rare user error in
> non-C++ compilation).

From the log attached to this bug report, it looks like the errors are
triggered when including: lib/clang/5.0.0/include/cuda_wrappers/new - I was
trying to suggest putting the ifdef checks in those headers.

[Quuxplusone]

I'm not sure how the cuda_wrapper/new got included in this case. My attempt to
reproduce the issue fails due to missing <cuda.h>. This is what I'd expect as
C++ compilation does not set up include paths to CUDA headers.

clang++ --cuda-path=/usr/local/cuda-8.0  streamcluster_cuda_cpu.cpp -c
tra@art2:~/tmp/Streamcluster
In file included from streamcluster_cuda_cpu.cpp:16:
./streamcluster_header.cu:29:10: fatal error: 'cuda.h' file not found
#include <cuda.h>
         ^~~~~~~~
1 error generated.

In any case CUDA sources often rely on the headers nvcc always pre-includes, so
they don't always #include <cuda.h>. E.g. your complete CUDA source may be just
this single line which I would expect to be compilable:
__device__ void foo() {}

That said, I'll add some sanity checks in cuda_wrappers so we at least catch
inclusion of CUDA wrappers during C++ compilation mode.