Retrieving address space from the AST using attributes

classic Classic list List threaded Threaded
3 messages Options
Reply | Threaded
Open this post in threaded view
|

Retrieving address space from the AST using attributes

Simon Brand
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
Reply | Threaded
Open this post in threaded view
|

Re: Retrieving address space from the AST using attributes

gbenyei
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


-----Original Message-----
From: [hidden email] [mailto:[hidden email]] On Behalf Of Simon Brand
Sent: Monday, January 14, 2013 14:14
To: [hidden email]
Subject: [cfe-dev] Retrieving address space from the AST using attributes

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
Expr->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




--
View this message in context: http://clang-developers.42468.n3.nabble.com/Retrieving-address-space-from-the-AST-using-attributes-tp4029688.html
Sent from the Clang Developers mailing list archive at Nabble.com.
_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
---------------------------------------------------------------------
Intel Israel (74) Limited

This e-mail and any attachments may contain confidential material for
the sole use of the intended recipient(s). Any review or distribution
by others is strictly prohibited. If you are not the intended
recipient, please contact the sender and delete all copies.


_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: Retrieving address space from the AST using attributes

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