Extent of CUDA support in clang as of 2.9-rc1


I’m working on a CUDA source to source compiler which is using clang as it’s frontend. I’ve read some of the messages on this mailing list which seems to indicate that CUDA support is on the way, and looking through the source for 2.9-rc1 shows signs that the AST should be able to handle at least some CUDA semantics.

Unfortunately, when I do pass some CUDA code through my program, it clang throws parse errors on tokens such as global and shared. In my initialization, I’ve set LangOptions.CUDA = 1. Did I miss something out, or is CUDA support in clang still incomplete? I’ve spent the better part of the last day looking through the docs and source trying to figure it out, and any help or pointers in the right direction would be very much appreciated.

Chee Eng

Hi Tan,

The symbols global and shared are not actual tokens but instead macros for the attributes global and shared. These are defined in headers included by cuda_runtime.h. If you include that, clang should be able to parse CUDA just fine.


Hi Chee,

The CUDA support in Clang is dependent on a set of runtime library
headers, and is currently designed to handle NVIDIA's headers, which
contains macro definitions resembling:

#define __global__ __attribute__((global))

You can see a partial emulation of these definitions in the test
suite header file test/SemaCUDA/cuda.h.

Furthermore, CUDA support is incomplete and consists of only parsing
and basic semantic checks. I have a patch series lined up covering
further semantic checks and code generation support, but this is
also incomplete. However, what exists at the moment may suffice for
a source to source compiler, if you know that the input program is
valid and that you do not need LLVM code generation.

If you are interested, I can send some test scripts I have been
using to test Clang against the NVIDIA headers, which may allow you
to parse code that uses these headers.

I have asked the release manager to disable CUDA support in 2.9 due
to its incomplete state, so if you would like to use CUDA I would
recommend that you follow SVN trunk.



Thank you, Peter and Gabriel. About the test scripts, yes, I would very thankful if you could pass me those scripts.

Also, I hope you don’t mind a further question (I’m relatively new to CUDA): The runtime library headers which you mentioned, are they included in the NVIDIA’s CUDA toolkit? Or are they from somewhere else? I attempted to look for them, but could not find them. Even Google wasn’t very forthcoming…

Chee Eng