Parsing CUDA file to AST

Clang Frontend
1

Hi,
I ma building rewriter based on clang (based on some tutorial available on
the internet).
I am loading CUDA file and trying to traverse AST. I got a problem. It
doesn't recognize CUDA directives like __global__ etc.
I set:
TheCompInst.getLangOpts().CUDA = 1;
and source search patch to:
/usr/include/linux

Sould I do someting more?

I get an error:

:~/Documents/rewriter$ ./rewritersample bigfunc2.c
bigfunc2.c:17:1: error: unknown type name '__global__'
__global__ void foo22(int *a, int *b, int *c) {
^
bigfunc2.c:17:12: error: expected identifier or '('
__global__ void foo22(int *a, int *b, int *c) {
           ^
#include <cuda.h>

#define N_WAVE 1024
#define LOG2_N_WAVE 10

2

These are usually defined as preprocessor macros in the CUDA headers,
something like:

#define __device__ __attribute__((device))

You can take a peek at the public CUDA headers for more information.

Eli

3

Actually I have same problem as this guy:
http://lists.cs.uiuc.edu/pipermail/cfe-dev/2011-June/015596.html

I set search patches and now my program sees __global__ etc. but not
blockDim.x or threadIdx.x.
I set flag -D__CUDACC__ in makefile but it doesnt work. setting:
TheCompInst.getPreprocessorOpts().addMacroDef("__CUDACC__"); gives even
more errors like:

/usr/local/cuda-5.5/include/cuda_runtime_api.h:1029:17: warning: unknown
attribute 'cudart_builtin' ignored
extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI
cudaGetLastError(void);
                ^
/usr/local/cuda-5.5/include/host_defines.h:199:22: note: expanded from
macro '__cudart_builtin__'
        __location__(cudart_builtin)
                     ^
/usr/local/cuda-5.5/include/host_defines.h:83:22: note: expanded from macro
'__location__'
        __annotate__(a)

4

Actually I have same problem as this guy:
http://lists.cs.uiuc.edu/pipermail/cfe-dev/2011-June/015596.html

I set search patches and now my program sees __global__ etc. but not
blockDim.x or threadIdx.x.

Yeah, these don't work in Clang yet. Patches welcome!

I have some very initial patches to implement this, but very far from
something I could actually send upstream. It may take time until I get to
that.

Eli

5

Yep, but there are lots of publications about CU2CL which is CUDA to
Open-Cl source-to-source compiler. I'm curious how they made it if this
approach is not supported yet. Unfortunately there is no code of this
project in the Internet :confused: What should I do?

6

Is it possible to do it by hand. Define it somewhere to don;t get error and
clang will recognize it as statements?

7

one quick approach would be to generate a “dummy” C code out of CUDA device code.

By “dummy”, I mean replace CUDA specific variables with something else and take special care of them after parsing.

However, this is a hacky solution and may not work in some cases but if your needs are not huge then I think this should work.

I am also badly waiting for clang to compile CUDA code to LLVM IR.