Determine CUDA qualifiers in AST

Hi everbody,

I am currently working on CUDA source-to-source transformations using clang
version 3.0 (trunk 132310).

I understand from the sources that there is support for CUDA qualifiers such as
__global__, __shared__ etc. However, I don't know how to find out for example
which functions have the __global__ qualifiers attached.

I am currently using an ASTConsumer to traverse the functions and statements:

...
for (DeclContext::decl_iterator
            D = C.getTranslationUnitDecl()->decls_begin(),
                DEnd = C.getTranslationUnitDecl()->decls_end();
                D != DEnd;
                ++D)
{
            FunctionDecl *FD;
            Stmt *Body = FD->getBody();

            ...
for (Stmt::child_iterator stmt = Body->child_begin(); stmt != Body->child_end();
stmt++) {...}
}
...

Could anyone please tell me which methods to use to extract this information
(or is there a better way than mine to achieve this?) ?

Thanks a lot
Felix

Hi Felix,

The __global__, __shared__ etc. qualifiers are represented as
declaration attributes. You can test for the presence of attributes
using the Decl::hasAttr function. For example:

if (FD->hasAttr<CUDAGlobalAttr>()) {
  ...
}

Thanks,

Peter Collingbourne <peter@...> writes:

Hi Felix,

The __global__, __shared__ etc. qualifiers are represented as
declaration attributes. You can test for the presence of attributes
using the Decl::hasAttr function. For example:

if (FD->hasAttr<CUDAGlobalAttr>()) {
  ...
}

Hi Peter,

I expected it to work this way.
Unfortunately, it didn't work out for me.
I used the following code:
void HandleTranslationUnit(ASTContext &C) {
...

for (DeclContext::decl_iterator
                D = C.getTranslationUnitDecl()->decls_begin(),
                    DEnd = C.getTranslationUnitDecl()->decls_end();
                    D != DEnd;
                    ++D)
            {
                FunctionDecl *FD;
                Stmt *Body;
                if ((FD = dyn_cast<FunctionDecl > (*D)) != 0/)
                {
                    if (FD->hasAttr<CUDAGlobalAttr>())
                        std::cout << "Found __global__ qualifier";
}
}
}

on the file my_kernel.cu:

#include <cuda_runtime.h>

__global__
void kernel(int *a)
{
  int tx = blockIdx.x * blockDim.x + threadIdx.x;
}

However, no message is written.
My diagnostic printer prints:

Diagnostics:
my_kernel.cu:6:11: error: use of undeclared identifier 'blockIdx'
my_kernel.cu:6:24: error: use of undeclared identifier 'blockDim'
my_kernel.cu:6:37: error: use of undeclared identifier 'threadIdx'

Can this have anything to do with it?

Thanks,
Felix

Assuming you are using the NVIDIA header files, you should add
-D__CUDACC__ to your frontend command line. This will cause the
headers to both define the __global__ (etc.) macros correctly and
declare the blockIdx (etc.) variables.

Thanks,

Peter Collingbourne <peter@...> writes:

Assuming you are using the NVIDIA header files, you should add
-D__CUDACC__ to your frontend command line. This will cause the
headers to both define the __global__ (etc.) macros correctly and
declare the blockIdx (etc.) variables.

Thanks a lot,
using addMacroDef("__CUDACC__") on the PreprocessorOpts did the trick.

Felix