similar to: [LLVMdev] RFC: Representation of OpenCL Memory Spaces

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

Hi, Tanya and I also prefer the extended TBAA solution as it naturally fits with LLVM. From my understanding of TBAA, it seems to provide the power to describe the relationship between address spaces for alias analysis, i.e., it can describe if two address spaces are disjoint or one may nest within another. For OpenCL, it is most useful to indicate that address spaces are disjoint from the

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

Justin, Out of these options, I would take the metadata approach for AA support. This doesn't solve the problem of different frontend/backends choosing different address space representations for the same language, but is the correct approach for providing extra information to the optimizations. The issue about memory spaces in general is a little different. For example, based on the

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

Hi, I'm working on OpenCL and I'm using clang as compiler (based on clang 3.7.0). I have a issue, I'm generating a bitcode file (that I can print before before the generation). But when I'm trying to read it again with clang, I have this issue: "error: Invalid record" How can I managed to know where it comes from? Thank you, Romaric Here is what is print before the

Hi, After a few distractions I'm back to work on the llvm tgsi backend. I've added clang integration and I can now compile a simple opencl program to something which sort of looks like tgsi. You can find my latest work on this here: http://cgit.freedesktop.org/~jwrdegoede/llvm http://cgit.freedesktop.org/~jwrdegoede/clang (the latter may still need to sync) I've a little test

On Aug 7, 2013, at 7:23 PM, Michele Scandale <michele.scandale at gmail.com> wrote: > On 08/08/2013 03:52 AM, Pete Cooper wrote: >>> Why a backend should be responsible (meaning have knowledge) for a >>> mapping between high level address spaces and low level address spaces? >> Thats true. I’m thinking entirely from the persecutive of the backend >> doing

On Aug 7, 2013, at 2:07 PM, Matt Arsenault <Matthew.Arsenault at amd.com> wrote: > On 08/07/2013 01:52 PM, Michele Scandale wrote: >> >> IMHO this information should be a plus that could be *safely* ignored when not necessary and used where it can provide an improvement in optimizations. This does not necessary mean the the middle-end (and the back-ends) must be aware of the

It's not clear what is actually wrong from your original message, I think you need to give some more information as to what you are doing: Example source, what target GPU, compiler error messages or other evidence of "it's wrong" (llvm IR, disassembly, etc) ... -- Mats On 8 April 2016 at 09:55, Liu Xin via llvm-dev <llvm-dev at lists.llvm.org> wrote: > I built it

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

Hi all, According to the SPIR 2.0 spec[1], the name of OpenCL builtins are mangled. However, when I compile OpenCl code with Clang 3.9 with the "spir64-unknown-unknown" target, Clang generates IR without mangling the builtins, e.g. for: __kernel void input_zip_int(__global int *in0) { *in0 = get_global_id(0); } clang generates: define spir_kernel void @input_zip_int(i32

On Aug 7, 2013, at 2:54 PM, Michele Scandale <michele.scandale at gmail.com> wrote: >> I don’t know if CUDA has aliasing address spaces, but that would also be >> useful to consider. Something simple like this might work. Note i’m >> using the examples from the clang discussion, that is "1 = opencl/cuda >> global, 2 = opencl_local/cuda_shared, 3 = opencl/cuda

+ Alexey Anastasia According to SPIR spec v1.2 s2.10.3 2.10.3 The printf function The printf function is supported, and is mangled according to its prototype as follows: int printf(constant char * restrict fmt, ... ) Note that the ellipsis formal argument (...) is mangled to argument type specifier z It seems printf should be mangled. Alexey/Anastasia, What do you think? Thanks. Sam From:

On Thu, Oct 13, 2011 at 06:59:47PM +0000, Villmow, Micah wrote: > Justin, > Out of these options, I would take the metadata approach for AA support. > > This doesn't solve the problem of different frontend/backends choosing different > address space representations for the same language, but is the correct > approach for providing extra information to the optimizations.

Hi Tom, Thanks for taking the time to answer this. On 11-01-16 18:10, Tom Stellard wrote: > On Mon, Jan 11, 2016 at 12:07:14PM +0100, Hans de Goede wrote: >> Hi, >> >> After a few distractions I'm back to work on the llvm tgsi backend. I've >> added clang integration and I can now compile a simple opencl program >> to something which sort of looks like

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

On 08/07/2013 01:52 PM, Michele Scandale wrote: > > IMHO this information should be a plus that could be *safely* ignored > when not necessary and used where it can provide an improvement in > optimizations. This does not necessary mean the the middle-end (and > the back-ends) must be aware of the semantic of these logical address > spaces, it would be enough just to

On 8 Aug 2013, at 04:23, Pete Cooper <peter_cooper at apple.com> wrote: > > On Aug 7, 2013, at 7:23 PM, Michele Scandale <michele.scandale at gmail.com> wrote: > >> On 08/08/2013 03:52 AM, Pete Cooper wrote: >> >> From here I understand that in the IR there are addrspace(N) where N=0,1,2,3,... according to the target independent mapping done by the

> I don’t know if CUDA has aliasing address spaces, but that would also be > useful to consider. Something simple like this might work. Note i’m > using the examples from the clang discussion, that is "1 = opencl/cuda > global, 2 = opencl_local/cuda_shared, 3 = opencl/cuda constant" You are assuming that the target device has different physical address spaces (like, PTX