[LLVMdev] SPIR Portability Discussion

[Villmow, Micah]

> -----Original Message-----
> From: llvmdev-bounces at cs.uiuc.edu [mailto:llvmdev-bounces at cs.uiuc.edu]
> On Behalf Of Nadav Rotem
> Sent: Wednesday, September 12, 2012 2:05 PM
> To: Ouriel, Boaz
> Cc: cfe-dev at cs.uiuc.edu; llvmdev at cs.uiuc.edu
> Subject: Re: [LLVMdev] SPIR Portability Discussion
> 
> 
> >
> >         ****Pointers****
> > During SPIR generation, the size, and the alignment of pointers is
> unknown (32 vs. 64 bits).
> > The SPIR representation shouldn't assume anything about the size and
> the alignment of pointers,
> > but it might use pointers in the usual way (except from using GEP
> when the pointed type has unknown size - this one is illegal in SPIR
> and will fail the SPIR verification pass which was written by Khronos
> members)
> 
> I don't understand the GEP restriction. Can I use GEP on strucst with
> pointers or size_t ?
[Villmow, Micah] Yes, pointers inside of structures are well defined.
> This is important if OpenCL 2.0 allows structs with pointers (for
> implementing linked lists, etc).
> 
> Also, future OpenCL versions may introduce C++ features to the
> language.  You need to be prepared to supports these features in SPIR.
> For example, c++ references may require SPIR to handle GEPs to structs
> that contain size_t members.
[Villmow, Micah] SPIR 1.0 targets OpenCL 1.2, so features outside of OpenCL 1.2 are also outside of the scope of SPIR 1.0.
> 
> >         ****size_t*****
> > SPIR tries to deal with size_t , ptrdiff_t, uintptr_t, intptr_t.
> Since these types have device specific size and alignment, their
> behavior is uncertain during compilation time.
> > SPIR represents these types as opaque types, and defines "builtin"
> functions to handle them.
> 
> If we ignore the issue of size_t inside structs,  I don't see the
> problem with deciding that size_t is 64bits, even on 32bit systems.
> The only place that I saw that size_t was used, in user code, is in the
> get_global_id() family of functions (and other APIs which require
> offsets).  A target-specific compiler optimization can reduce the bit
> width of the get_global_id (and friends) back to 32bits and propagate
> this, if needed.
> 
> 
> _______________________________________________
> LLVM Developers mailing list
> LLVMdev at cs.uiuc.edu         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev