[cfe-dev] Extent of CUDA support in clang as of 2.9-rc1

[Peter Collingbourne]

On Fri, Mar 18, 2011 at 11:54:06PM +0800, Tan Chee Eng wrote:
> Hi,
> 
> 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.

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.

Thanks,
-- 
Peter