[LLVMdev] SPIR Portability Discussion

[Nadav Rotem]

> 
>         ****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 ? 
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.  

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