search for: localbuff

...s meaning that accesses > across work-items invoke undefined behaviour, so in the example above, > the write to x via p would itself be undefined. > I was referring more to the front-end aspects here. Let's say we have: __kernel void foo() { float privateBuffer[8]; __local float localBuffer[8]; } What mechanisms, other than address spaces, can we use to tell the X86 back-end that privateBuffer is private to the thread, and localBuffer is shared among all threads in a work-group? > > Thanks, > -- > Peter > -- Thanks, Justin Holewinski -------------- next part...

On Thu, Oct 13, 2011 at 04:14:09PM -0400, Justin Holewinski wrote: > On Thu, Oct 13, 2011 at 11:57 AM, Peter Collingbourne <peter at pcc.me.uk>wrote: > > > Hi Justin, > > > > Thanks for bringing this up, I think it's important to discuss > > these issues here. > > > > On Thu, Oct 13, 2011 at 09:46:28AM -0400, Justin Holewinski wrote: > >

I am working on a custom LLVM backend for a 16 bit architecture. For my architecture, I need smaller array alignments for arrays created on the stack. For example, consider the following code at the start of a C function: char localBuff[20]; char localBuff2[6]; this gets converted by Clang into this: %localBuff = alloca [20 x i8], align 16 %localBuff2 = alloca [6 x i8], align 1 Note that char arrays smaller than 16 bytes are specified to be aligned to 1 byte, but bigger ones are aligned to 16 bytes. On the LLVM backend,...