Retrieving address space from the AST using attributes

Clang Frontend
1

Hi there,
Part of a program I am writing traverses the AST for an OpenCL program. I am
attempting to differentiate between accesses to variables in global and
local memory. These are defined like this:

#define __global __attribute__((address_space(1)))
#define __local __attribute__((address_space(3)))

__kernel void test(__global float *in, __local float *out)
{
...
}

The problem is that I can't seem to retrieve the address space from the AST.
My current approaches are all in the function:
bool VisitDeclRefExpr(const DeclRefExpr *Expr)

I have tried:
Expr->getDecl()->getType().getAddressSpace() //returns 0
Expr->getDecl()->getType()->getAs<AttributedType>() //returns null pointer
Context->getDeclAttrs(Expr) //returns empty AttrVec

However, calling Expr->getNameInfo().getName().getAsString() returns
"__attribute__((address_space(1))) float *" for the "in" variable and
"__attribute__((address_space(3))) float *" for the "out" variable, as one
would expect.

Is there something I'm doing wrong or a correct way to do this? I am using
Clang v3.2.
Thanks

2

Hi Simon,
First of all, Clang 3.2 supports the OpenCL attributes, so you don't have to define them as preprocessor macros. The only thing that is needed in order to generate LLVM with the target specific address spaces is to add the address space mapping in Targets.cpp. You could also use one of the targets that already have a mapping, like SPIR, TCE or PTX, or use the fake address space map by using the -ffake-address-space-map command line option.

Except from that, the QualType has its address space in its qualifiers, so no other attributes needed. For a pointer, you should look at the address space of the pointee type - I guess that's the actual issue you see.

Thanks
    Guy

3

Hi Guy,
Thanks very much, that solved my issue.
Cheers,
Simon