Ask for CUDA supports in Clang

Hi there,

May I ask what’s the status of supporting CUDA on Clang?
Are there any incomplete patches to use for CUDA? I am doing
a source-to-source translation work for CUDA based on Clang.

Thanks,
Pengcheng

Hi there,

Hi Pengcheng,

CUDA is still only partially supported. This has been discussed a few times
this year:

http://lists.cs.uiuc.edu/pipermail/cfe-dev/2014-September/038979.html
http://lists.cs.uiuc.edu/pipermail/cfe-dev/2014-March/035782.html

Things haven't changed significantly since then. Patches welcome!
Eli

Thanks, Eli.

May I know what your initial support is? Can you please open it to me?
I don’t need too strong CUDA support, say generating CUDA code to LLVM
IR. What I only need to do is to parse CUDA code to Clang AST, and then
translate it back .

Thanks,
Pengcheng

Thanks, Eli.

May I know what your initial support is? Can you please open it to me?
I don’t need too strong CUDA support, say generating CUDA code to LLVM
IR. What I only need to do is to parse CUDA code to Clang AST, and then
translate it back .

So this part should already work, if you're careful. The main missing parts
for just parsing the CUDA source into the AST are:

1. Real support for the threadIdx, blockIdx, etc. special globals
2. Various header definitions for __global__ & friends, CUDA builtins, etc.

If all you need is to *parse*, then a simple header like
https://gist.github.com/eliben/b014ac17cbe5a452803f should do the trick.

Then you can do:

$ clang -cc1 -ast-dump -x cuda -include <path-to-header-file-in-gist>
cuda-input-file.cu

And this should dump the AST. If you need to support more CUDA builtins
(and/or runtime functions), you should add them to the header, *or* try to
get Clang to parse the NVIDIA CUDA headers.

Eli

Eli,

Lacking of real support for threadIdx, blockIdx, etc. may be fine for Pengcheng (?) because he’s doing CUDA->AST->CUDA transformation.

Jingyue

Eli,

Lacking of real support for threadIdx, blockIdx, etc. may be fine for
Pengcheng (?) because he's doing CUDA->AST->CUDA transformation.

Yep, I think this is what I wrote below suggests. The linked header creates
fake declarations of these globals, which is good enough for Clang's
parsing.

Eli

Thanks, Eli and Jingyue. I just thought to fake them that way. Plus, consider corresponding
semantics of CUDA syntax when processing AST.

Hope to see you guys at google next summer. :slight_smile:

Pengcheng

By the way, this simple header file https://gist.github.com/eliben/b014ac17cbe5a452803f was pushed into Github
3 hours ago. Haha~ :slight_smile:

By the way, this simple header file
https://gist.github.com/eliben/b014ac17cbe5a452803f
<https://urldefense.proofpoint.com/v1/url?u=https://gist.github.com/eliben/b014ac17cbe5a452803f&k=p4Ly7qpEBiYPBVenR9G2iQ%3D%3D &r=8jiWUOA7YNYuE%2FO4heswhF5KaS3dN4VEfeX8hooBpaw%3D &m=uCB6%2BIbeBWRMZqs%2B6BtY9ZPvLRQLB3nVYqkh0i2fkkA%3D &s=2d9719353c5102fec52c164e3cf7a04c32f063166ba55d31e0f51324118acbf7> was
pushed into Github
3 hours ago. Haha~ :slight_smile:

Yes, I pushed it specifically to provide you with a concrete header :slight_smile:

Eli