search for: __constant

...gt; Cc: cfe-dev at lists.llvm.org; llvm-dev <llvm-dev at lists.llvm.org> Subject: Re: builtins name mangling in SPIR 2.0 Hi Sam and others, I saw that in [1], printf is mangled to _Z6printfPrU3AS2cz, while in clang's opencl-c.h[2], printf does not have the overload attribute: int printf(__constant const char* st, ...); (and it is different from the standard, which is printf(restrict __constant char *, ...)) I try the following code: #include <opencl-c.h> __kernel void vadd(__global const int* a, __global const int* b, __global int* c) { printf("aaaaa"); } and get a printf...

Thanks a lot. On Mon, Sep 12, 2016 at 1:42 PM, Liu, Yaxun (Sam) <Yaxun.Liu at amd.com> wrote: > If you use the default header file under clang/lib/Headers/opencl-c.h, > get_global_id will be mangled. > > > > If you want to declare get_global_id in your own header, add > __attribute__((overloadable)), then it will be mangled. > > > > Sam > > > >

...ev <llvm-dev at lists.llvm.org<mailto:llvm-dev at lists.llvm.org>> Subject: Re: builtins name mangling in SPIR 2.0 Hi Sam and others, I saw that in [1], printf is mangled to _Z6printfPrU3AS2cz, while in clang's opencl-c.h[2], printf does not have the overload attribute: int printf(__constant const char* st, ...); (and it is different from the standard, which is printf(restrict __constant char *, ...)) I try the following code: #include <opencl-c.h> __kernel void vadd(__global const int* a, __global const int* b, __global int* c) { printf("aaaaa"); } and get a printf...

...d. > > > > Which mechanisms could be used to differentiate between thread-private and > __local data? In OpenCL C, it is illegal to declare a variable with static storage duration in the __private address space (section 6.5: "All program scope variables must be declared in the __constant address space."; section 6.8g: "The extern, static, auto and register storage-class specifiers are not supported."). This implies that there is no way for pointers to the __private address space to be usefully shared between work-items without invoking undefined behaviour, so the qu...

...mechanisms could be used to differentiate between thread-private > and > > __local data? > > In OpenCL C, it is illegal to declare a variable with static storage > duration in the __private address space (section 6.5: "All program > scope variables must be declared in the __constant address space."; > section 6.8g: "The extern, static, auto and register storage-class > specifiers are not supported."). This implies that there is no way > for pointers to the __private address space to be usefully shared > between work-items without invoking undefined...

...ess spaces. > (approx. July 8 - July 22) > > Operations on __global, __local, and __private memory will be > dealt with using the resource access opcodes, __kernel function > parameters will be accessed through a special resource meant for > parameter passing, __constant memory will be mapped to constant > buffers. > > * Get function calls working reliably. > (approx. July 22 - August 5) > > This will involve fixing the passing of aggregate types and > anything that doesn't fit in a 32-bit register, fixing stack >...

Francisco Jerez <currojerez at riseup.net> writes: >[...] > Pity not to see any interest in this since I brought up the idea two > weeks ago. I've uploaded a first attempt at writing a TGSI back-end > here [2]. It's able to generate code -- though only in assembly form > and with many loose ends still. > > Also, would it be possible for Tom Stellard

...(approx. July 8 - July 22) >> >> Operations on __global, __local, and __private memory will be >> dealt with using the resource access opcodes, __kernel function >> parameters will be accessed through a special resource meant for >> parameter passing, __constant memory will be mapped to constant >> buffers. >> >> * Get function calls working reliably. >> (approx. July 22 - August 5) >> >> This will involve fixing the passing of aggregate types and >> anything that doesn't fit in a 32-bit reg...

On Mon, Feb 28, 2011 at 4:41 PM, Peter Collingbourne <peter at pcc.me.uk> wrote: > > The more I think about it, the more I become uncomfortable with the > concept of language-specific address spaces in LLVM.  These are the > main issues I see with language-specific address spaces: ... > Instead of language-specific address spaces, each target should > concentrate on

On Tue, Mar 1, 2011 at 4:06 PM, David Neto wrote: > On Mon, Feb 28, 2011 at 4:41 PM, Peter Collingbourne wrote: >> >> The more I think about it, the more I become uncomfortable with the >> concept of language-specific address spaces in LLVM.  These are the >> main issues I see with language-specific address spaces: > > ... > >> Instead of language-specific

On Fri, Feb 25, 2011 at 02:55:33PM -0500, Ken Dyck wrote: > The address space mechanism is used by some code generators to > differentiate between physical memory spaces. The PIC16 code generator > uses address spaces 0 and 1 to select between its RAM and ROM spaces. > And X86 uses address space 256 for GS and 257 for FS. In the back end > for a dual-harvard DSP that I've been

...re status, using two target platforms X86 and R600. The part related to the description of source language address spaces features through TBAA like metadata is omitted because I think we agree on that part. /// test.cl /// __kernel void convolve(const __global int *input, __constant int *mask, __global int *output) { unsigned x = get_global_id(0); output[x] = input[x] + mask[get_local_id(0)]; } The IR for R600 now is: /// test.r600.ll /// target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-v16:16:16-v24:32:...

...2) > >> > >> Operations on __global, __local, and __private memory will be > >> dealt with using the resource access opcodes, __kernel function > >> parameters will be accessed through a special resource meant for > >> parameter passing, __constant memory will be mapped to constant > >> buffers. > >> > >> * Get function calls working reliably. > >> (approx. July 22 - August 5) > >> > >> This will involve fixing the passing of aggregate types and > >> anything t...

> -----Original Message----- > From: Michele Scandale [mailto:michele.scandale at gmail.com] > Sent: Saturday, August 10, 2013 6:29 AM > To: Micah Villmow > Cc: LLVM Developers Mailing List > Subject: Re: [LLVMdev] Address space extension > > On 08/10/2013 02:47 PM, Micah Villmow wrote: > > Michele, > > The information you are trying to gather is fundamentally

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: > > It is becoming increasingly clear to me that LLVM address spaces are not > the > > general solution

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: > It is becoming increasingly clear to me that LLVM address spaces are not the > general solution to OpenCL/CUDA memory spaces. They are a convenient hack to > get things working in the short term, but I think a more