Better CUDA Support
Open, Needs TriagePublic

Description

Background

With commit 64d7d49e, our clang parser gains some basic ability to understand the kernel launch syntax in .cu files. However, nvcc and clang implicitly include many headers when compiling, which our parser doesn't, thus generating false alarms.

For now the user could manually include those headers in there file to get correct parsing result, but this of course is less than desirable.

Note: The include file list can be get by nvcc -E cuda.cu | egrep -o '/opt/cuda/.*\.h'| sort | uniq or clang++ --cuda-path=/opt/cuda -E cuda.cu | egrep -o '/opt/cuda/.*\.h'| sort | uniq. The file list is slightly different between nvcc and clang.

TODO

  1. nvcc knows where the CUDA installation is, but clang doesn't. Thus clang needs --cuda-path argument when compiling .cu files. Find a way to pass this argument to the clang parser. This information can be either provided directly by the user (similar to custom include path), or automatically found by looking for nvcc in PATH, or by looking at cmake generated information (not sure if the last one is possible, but cmake does has FindCuda)
  2. Not sure if adding --cuda-path will trigger the inclusion of additional header files by clang. If not, then we need to find a way to include these header files in our parser.
qi437103 created this task.Mar 23 2017, 1:19 AM
kfunk added a subscriber: kfunk.

One thing to keep in mind: you can make KDevelop print how clang is invoked by setting this env var: KDEV_CLANG_DISPLAY_ARGS=1 I'm not sure why libclang does not use the implicitly included headers automatically.

Please at least make sure the unit test for cuda.cu does not fail on our CI for now. We don't have a CUDA installation available on the CI anyway.

mwolff added a subscriber: mwolff.Mar 23 2017, 2:16 PM

We are explicitly excluding the standard includes for C and C++, so maybe something like that accidentally happens for cuda, too? Cf. nostdinc.

qi437103 added a comment.EditedMar 23 2017, 5:43 PM

Yeah there're nocudainc and nocudalib. But without cuda includes, some cuda structures like __device__ and __global__ will be undefined.

So are the standard includes never parsed, or only parsed once for all files? Since I still get correct highlight for vector and like. Is it possible to do similar thing to cuda includes?

Sorry I misread the last comment. No it's not affected by nostdinc or nostdinc++, as in current code those are not added to the command line.

After played with the command line directly, I find there might be more complex issues if we have to exclude standard includes.

  • clang --cuda-path=/opt/cuda -std=c++11 -nostdinc -nostdinc++ -xcuda -isystem/usr/include/c++/6.3.1 -isystem/usr/include/c++/6.3.1/x86_64-pc-linux-gnu -isystem/usr/include/c++/6.3.1/backward -isystem/usr/local/include -isystem/usr/lib/clang/3.9.1/include -isystem/usr/include -imacros /home/aetf/definestest_files-clang.J25682 /home/aetf/develop/vcs/kdevelop/languages/clang/tests/files/cuda.cu

gives various undeclared identifier errors for clang 3.9 and 4.0

  • clang --cuda-path=/opt/cuda -std=c++11 -nostdinc++ -xcuda -isystem/usr/include/c++/6.3.1 -isystem/usr/include/c++/6.3.1/x86_64-pc-linux-gnu -isystem/usr/include/c++/6.3.1/backward -isystem/usr/local/include -isystem/usr/include -imacros /home/aetf/definestest_files-clang.J25682 /home/aetf/develop/vcs/kdevelop/languages/clang/tests/files/cuda.cu

Removing -nostdinc and -isystem/usr/lib/clang/3.9.1/include makes it work well for clang 4.0 but generates the same error for clang 3.9

  • Whatever other arguemts passed in, nocudainc will make it generate errors about undefined cuda keywords.

I think it's OK to rely on clang only for cuda and not support GCC emulation. This doesn't even really work properly for C++ either, so don't go down that rabbit hole.

So from my side, feel free to remove the -nostdinc flags etc. from CUDA files. Or only support clang 4.0+. It's a new feature, so from my side it would be OK to only support such a new clang.

I mistakenly set wrong include paths for Clang 3.9 and 4.0. That's why I was getting so many errors. After fix that up, everything works as expected even with -nostdinc flags.

Remaining work:

cuda_runtime.h should already been implicitly included by Clang, but seems explicitly include that file gives more information. See limitations in D5210.