[LLVMdev] SPIR Portability Discussion
nrotem at apple.com
Wed Sep 12 14:04:34 PDT 2012
> 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.
> 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.
More information about the llvm-dev