llvm dev - Sep 2012 - [LLVMdev] SPIR Portability Discussion

Hey All,

This is a very big topic in SPIR and probably a very controversial one as well.
It includes dealing with 32 vs. 64 bit architectures and OpenCL "C"
endianness.
We have written down some of the aspects, but of course did not cover everything
- let's start a discussion on the portability and see where it takes us.
I suggest we start with the 32 vs. 64 bits discussion and then move to the
Endianness part.

****Introduction****
As a reminder, portability is one of SPIR's goals. 
However, SPIR does not attempt to solve inherent portability issues, which exist
in OpenCL "C" or in C99.
It is clear that OpenCL programs could be written in a way which make them non
portable and very device specific.
Such programs will never be portable. In addition, some corner case
scenario's which have been identified by Khronos members have been
disallowed in SPIR.
So, SPIR aims at being portable but not for every scenario. 

1) ****Portability between Devices with different address width (32 vs. 64
bits)****
During the design stages, Khronos members needed to decide on its philosophy
when it comes to dealing with the address width of devices (32 vs. 64bits).
During internal discussions, two alternatives came up. The first alternative was
to split SPIR into two sub-cases: SPIR 32bits and SPIR 64bits.
The second alternative was to try and abstract this information at SPIR level. 

Splitting SPIR into 32bit and 64bit is simpler to implement. However, it is less
portable.
This will require OpenCL developers to pre-compile two versions of their code
one for 32bit and another for 64bit devices and make their application aware at
runtime to the underlying device architecture.
OpenCL applications would need to load the proper SPIR binary based on the
device architecture.
An option that was raised during the discussions was to have a fat binary that
contains both 32bit and 64bit versions of SPIR binaries.
However, this option was controversial inside Khronos and eventually was not
accepted.
The decision was to pursue the second alternative. Khronos members understand
that this is a more complex alternative and does not guarantee 100% percent
coverage to all cases.
However, as stated before, SPIR attempts to solve the important cases. Those
particular cases which SPIR will not be able to address are explicitly
documented in the specification.

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

         *****Sizeof******
Most valid built-in and user specific types in OpenCL have known non
device-specific size.
However, for some types (pointers, size_t, ptrdiff_t) the size is unknown during
compilation.
To overcome this issue, SPIR provides functions to substitute the constant
values of the sizeof operator.
These functions should be resolved by the device backend compiler when producing
the final machine code of the OpenCL program.

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

          ****structures****
Structures are also a major issue in OpenCL in general, since their layout and
size are compiler specific. To handle this issue, SPIR defines a standard layout
for structures.

2) ****Host and Device Endianness*****
Before diving into the details of how Endianness is dealt in SPIR, an
introduction to Endianness in OpenCL is required.
In a nutshell, OpenCL standard facilitates the means to mark the endianness type
of variables, which reside in global or constant address space memory.
Since such variables reside in global memory they might have conflicting
endianness between the host and the device.
Hence, OpenCL standard facilitates two types of endianness - a
"device" and "host" types.
The "host" type indicates that the variable uses the endianness of the
host processor.
The "device" type indicates that the variable uses the endianness of
the device on which the program will be executed.
The default type is the "device" type. When the user writes down
programs which rely on the endianness of a particular device -
his code becomes incompatible with devices whose endianness differ, and by
definition is non-portable at OpenCL level.

SPIR specification attempts to facilitate the same mechanism that OpenCL does.
Since "device" type is the default, the only type which requires
special handling is "host".
Initially, Khronos members considered the usage of metadata as the preferred
method for achieving this goal.
Every variable that needs to be marked with "host" endianness type
would be associated with a metadata that indicates it.
This approach could work but is not guaranteed to be enforced by the different
LLVM optimization passes since it is a metadata and as such could be disregarded
by optimization passes.
After a few discussions, Khronos members decided that usage of address space
qualifier could achieve the same effect with better support from the different
optimization passes.
For example, a function that accepts an argument with "host" type can
pass this variable as an argument to another function where the argument is not
marked as well with this type.
Finally, this approach was chosen and is now a part of the specification
(described in section 2.8.2.2 of the specification)

3)****Materialization of a SPIR program****
Since device information is abstracted during SPIR generation, the build phase
of SPIR binaries to device binaries includes an additional phase which is called
"materialization" phase.
This phase resolves the abstracted information and "materializes" a
SPIR binary it to a specific device.

Thanks,
Boaz
---------------------------------------------------------------------
Intel Israel (74) Limited

This e-mail and any attachments may contain confidential material for
the sole use of the intended recipient(s). Any review or distribution
by others is strictly prohibited. If you are not the intended
recipient, please contact the sender and delete all copies.

On Wed, Sep 12, 2012 at 12:27 PM, Ouriel, Boaz <boaz.ouriel at
intel.com>wrote:
> Hey All,
>
> This is a very big topic in SPIR and probably a very controversial one as
> well. It includes dealing with 32 vs. 64 bit architectures and OpenCL
"C"
> endianness.
> We have written down some of the aspects, but of course did not cover
> everything - let's start a discussion on the portability and see where
it
> takes us.
> I suggest we start with the 32 vs. 64 bits discussion and then move to the
> Endianness part.
>
> ****Introduction****
> As a reminder, portability is one of SPIR's goals.
> However, SPIR does not attempt to solve inherent portability issues, which
> exist in OpenCL "C" or in C99.
> It is clear that OpenCL programs could be written in a way which make them
> non portable and very device specific.
> Such programs will never be portable. In addition, some corner case
> scenario's which have been identified by Khronos members have been
> disallowed in SPIR.
> So, SPIR aims at being portable but not for every scenario.
>
> 1) ****Portability between Devices with different address width (32 vs. 64
> bits)****
> During the design stages, Khronos members needed to decide on its
> philosophy when it comes to dealing with the address width of devices (32
> vs. 64bits).
> During internal discussions, two alternatives came up. The first
> alternative was to split SPIR into two sub-cases: SPIR 32bits and SPIR
> 64bits.
> The second alternative was to try and abstract this information at SPIR
> level.
>
> Splitting SPIR into 32bit and 64bit is simpler to implement. However, it
> is less portable.
> This will require OpenCL developers to pre-compile two versions of their
> code one for 32bit and another for 64bit devices and make their application
> aware at runtime to the underlying device architecture.
> OpenCL applications would need to load the proper SPIR binary based on the
> device architecture.
> An option that was raised during the discussions was to have a fat binary
> that contains both 32bit and 64bit versions of SPIR binaries.
> However, this option was controversial inside Khronos and eventually was
> not accepted.
> The decision was to pursue the second alternative. Khronos members
> understand that this is a more complex alternative and does not guarantee
> 100% percent coverage to all cases.
> However, as stated before, SPIR attempts to solve the important cases.
> Those particular cases which SPIR will not be able to address are
> explicitly documented in the specification.
>
>          ****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)
>
>          *****Sizeof******
> Most valid built-in and user specific types in OpenCL have known non
> device-specific size.
> However, for some types (pointers, size_t, ptrdiff_t) the size is unknown
> during compilation.
> To overcome this issue, SPIR provides functions to substitute the constant
> values of the sizeof operator.
> These functions should be resolved by the device backend compiler when
> producing the final machine code of the OpenCL program.
>
OpenCL 1.2 (6.3)/k says the result of sizeof is an ICE. So these are valid:

int does_this_compile[sizeof(void*) - 3];

struct how_do_you_represent_this_in_IR {
  int a : 1;
  int b : sizeof(void*) * 4;
};

Is OpenCL going to be changed to reject these cases?


How do you perform record layout if the size of a pointer is unknown? For
instance:

struct A {
  int *p;
  int n;
} a;
int arr[offsetof(A, n) - 3]; // or, int arr[(char*)&a.n - (char*)&a.p -
3];

-- Richard
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20120912/f25bf10b/attachment.html>

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

> 
> OpenCL 1.2 (6.3)/k says the result of sizeof is an ICE. So these are valid:
> 
> int does_this_compile[sizeof(void*) - 3];
> 
> struct how_do_you_represent_this_in_IR {
>   int a : 1;
>   int b : sizeof(void*) * 4;
> };
> 
> Is OpenCL going to be changed to reject these cases?
> 
I don't think that they plan to allow it. I am not sure how much value
dynamic sized bitfields bring to OpenCL users.
In theory they could use opaque types and a number of external functions which
can be lowered to legal LLVM-IR once the value of sizeof is known.
> How do you perform record layout if the size of a pointer is unknown? For
instance:
> 
> struct A {
>   int *p;
>   int n;
> } a;
> int arr[offsetof(A, n) - 3]; // or, int arr[(char*)&a.n -
(char*)&a.p - 3];
> 
They can replace LLVM's alloca with a fake function which can be lowered to
a regular alloca once the size is known.
> -- Richard
> _______________________________________________
> LLVM Developers mailing list
> LLVMdev at cs.uiuc.edu         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev

From: llvmdev-bounces at cs.uiuc.edu [mailto:llvmdev-bounces at cs.uiuc.edu] On
Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 1:55 PM
To: Ouriel, Boaz
Cc: cfe-dev at cs.uiuc.edu; llvmdev at cs.uiuc.edu
Subject: Re: [LLVMdev] SPIR Portability Discussion

On Wed, Sep 12, 2012 at 12:27 PM, Ouriel, Boaz <boaz.ouriel at
intel.com<mailto:boaz.ouriel at intel.com>> wrote:
Hey All,

This is a very big topic in SPIR and probably a very controversial one as well.
It includes dealing with 32 vs. 64 bit architectures and OpenCL "C"
endianness.
We have written down some of the aspects, but of course did not cover everything
- let's start a discussion on the portability and see where it takes us.
I suggest we start with the 32 vs. 64 bits discussion and then move to the
Endianness part.

****Introduction****
As a reminder, portability is one of SPIR's goals.
However, SPIR does not attempt to solve inherent portability issues, which exist
in OpenCL "C" or in C99.
It is clear that OpenCL programs could be written in a way which make them non
portable and very device specific.
Such programs will never be portable. In addition, some corner case
scenario's which have been identified by Khronos members have been
disallowed in SPIR.
So, SPIR aims at being portable but not for every scenario.

1) ****Portability between Devices with different address width (32 vs. 64
bits)****
During the design stages, Khronos members needed to decide on its philosophy
when it comes to dealing with the address width of devices (32 vs. 64bits).
During internal discussions, two alternatives came up. The first alternative was
to split SPIR into two sub-cases: SPIR 32bits and SPIR 64bits.
The second alternative was to try and abstract this information at SPIR level.

Splitting SPIR into 32bit and 64bit is simpler to implement. However, it is less
portable.
This will require OpenCL developers to pre-compile two versions of their code
one for 32bit and another for 64bit devices and make their application aware at
runtime to the underlying device architecture.
OpenCL applications would need to load the proper SPIR binary based on the
device architecture.
An option that was raised during the discussions was to have a fat binary that
contains both 32bit and 64bit versions of SPIR binaries.
However, this option was controversial inside Khronos and eventually was not
accepted.
The decision was to pursue the second alternative. Khronos members understand
that this is a more complex alternative and does not guarantee 100% percent
coverage to all cases.
However, as stated before, SPIR attempts to solve the important cases. Those
particular cases which SPIR will not be able to address are explicitly
documented in the specification.

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

         *****Sizeof******
Most valid built-in and user specific types in OpenCL have known non
device-specific size.
However, for some types (pointers, size_t, ptrdiff_t) the size is unknown during
compilation.
To overcome this issue, SPIR provides functions to substitute the constant
values of the sizeof operator.
These functions should be resolved by the device backend compiler when producing
the final machine code of the OpenCL program.

OpenCL 1.2 (6.3)/k says the result of sizeof is an ICE. So these are valid:

int does_this_compile[sizeof(void*) - 3];
[Villmow, Micah] 'ICE'? Integer compile time expression? While not
pretty, this can be represented in SPIR with the following sequence on
instructions
%1 = call %spir.size_t @__spir_sizet_convert_size_t(i32 3)
%2 = call %spir.size_t @__spir_size_of_sizet()
%3 = call %spir.size_t @__spir_sizet_sub(%spir.size_t %1, %spir.size_t %2)
%4 = call %spir.size_t @__spir_sizet_convert_i32(%spir.size_t %3)
%5 = alloca i32, i32 %4

struct how_do_you_represent_this_in_IR {
  int a : 1;
  int b : sizeof(void*) * 4;
};
[Villmow, Micah] Bitfields are illegal in OpenCL (See 6.9.c);

Is OpenCL going to be changed to reject these cases?


How do you perform record layout if the size of a pointer is unknown? For
instance:

struct A {
  int *p;
  int n;
} a;
int arr[offsetof(A, n) - 3]; // or, int arr[(char*)&a.n - (char*)&a.p -
3];
[Villmow, Micah] Since in the current implementation of SPIR, a pointer is
defined as 64bits when in a structure(SPIR spec 2.1.5), the offsets themselves
are well defined.

-- Richard
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20120912/bcd8f4be/attachment.html>

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