[cfe-dev] Determine CUDA qualifiers in AST

[Felix Schmitt]

Peter Collingbourne <peter at ...> 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