SPIR Portability Discussion

Previous Topic Next Topic
 
classic Classic list List threaded Threaded
22 messages Options
12
Reply | Threaded
Open this post in threaded view
|

SPIR Portability Discussion

Ouriel, Boaz
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.


_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: SPIR Portability Discussion

Richard Smith-33
On Wed, Sep 12, 2012 at 12:27 PM, Ouriel, Boaz <[hidden email]> 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

_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: SPIR Portability Discussion

Nadav Rotem
In reply to this post by Ouriel, Boaz

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


_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: SPIR Portability Discussion

Nadav Rotem
In reply to this post by Richard Smith-33
>
> 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
> [hidden email]         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev

_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: SPIR Portability Discussion

Villmow, Micah
In reply to this post by Richard Smith-33

 

 

From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 1:55 PM
To: Ouriel, Boaz
Cc: [hidden email]; [hidden email]
Subject: Re: [LLVMdev] SPIR Portability Discussion

 

On Wed, Sep 12, 2012 at 12:27 PM, Ouriel, Boaz <[hidden email]> 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


_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: SPIR Portability Discussion

Villmow, Micah
In reply to this post by Nadav Rotem


> -----Original Message-----
> From: [hidden email] [mailto:[hidden email]]
> On Behalf Of Nadav Rotem
> Sent: Wednesday, September 12, 2012 2:05 PM
> To: Ouriel, Boaz
> Cc: [hidden email]; [hidden email]
> 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
> [hidden email]         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev



_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: SPIR Portability Discussion

Villmow, Micah
In reply to this post by Nadav Rotem


> -----Original Message-----
> From: [hidden email] [mailto:[hidden email]]
> On Behalf Of Nadav Rotem
> Sent: Wednesday, September 12, 2012 2:17 PM
> To: Richard Smith
> Cc: [hidden email]; [hidden email]
> Subject: Re: [LLVMdev] SPIR Portability Discussion
>
> >
> > 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.
[Villmow, Micah] Yep, this is the basic idea of the steps taken to make SPIR portable.
There is no restriction in the C99 or OpenCL specs that I know of that requires the frontend to make decisions on the size of device specific constructs. While it is the most logical and efficient choice, it cannot be determined if the target device is not known, so the decision is delayed via function calls or opaque types until this information is known, the results are still compile time constants.

>
> > -- Richard
> > _______________________________________________
> > LLVM Developers mailing list
> > [hidden email]         http://llvm.cs.uiuc.edu
> > http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
>
> _______________________________________________
> LLVM Developers mailing list
> [hidden email]         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev



_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: [cfe-dev] SPIR Portability Discussion

Eli Friedman-2
On Wed, Sep 12, 2012 at 2:30 PM, Villmow, Micah <[hidden email]> wrote:

>
>
>> -----Original Message-----
>> From: [hidden email] [mailto:[hidden email]]
>> On Behalf Of Nadav Rotem
>> Sent: Wednesday, September 12, 2012 2:17 PM
>> To: Richard Smith
>> Cc: [hidden email]; [hidden email]
>> Subject: Re: [LLVMdev] SPIR Portability Discussion
>>
>> >
>> > 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.
> [Villmow, Micah] Yep, this is the basic idea of the steps taken to make SPIR portable.
> There is no restriction in the C99 or OpenCL specs that I know of that requires the frontend to make decisions on the size of device specific constructs.

Well, if you ignore the preprocessor, it's true that there's no strict
requirement....

>While it is the most logical and efficient choice, it cannot be determined if the target device is not known, so the decision is delayed via function calls or opaque types until this information is known, the results are still compile time constants.

The issue is that it starts to get nasty really fast in the general case:

// We're required to diagnose this iff sizeof(size_t) != 4.
extern int x[20];
int x[sizeof(size_t) * 5];

// We're required to diagnose this iff sizeof(size_t) == 4.
void f(int x) {
  switch(x) {
  case 4:
  case sizeof(size_t):
    break;
  }
}

You basically have to delay lowering anything that can involve an
integer constant expression.

-Eli
_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: SPIR Portability Discussion

Richard Smith-33
In reply to this post by Villmow, Micah
On Wed, Sep 12, 2012 at 2:23 PM, Villmow, Micah <[hidden email]> wrote:

 

 

From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 1:55 PM
To: Ouriel, Boaz
Cc: [hidden email]; [hidden email]
Subject: Re: [LLVMdev] SPIR Portability Discussion

 

On Wed, Sep 12, 2012 at 12:27 PM, Ouriel, Boaz <[hidden email]> 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];


Oops, I meant sizeof(void*) - 5.
 

[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


My point here was that if sizeof(void*) is an integer constant expression, then this code has a constraint violation if sizeof(void*) is 4 but not if sizeof(void*) is 8 (and my example was intended to be of a global array, not a function-local one, so you can't fall back to treating it as a VLA). Another case which might depend on this:

enum E {
  a = sizeof(void*) // is this valid?
};

Based on the behavior you describe above, it looks like sizeof applied to a pointer (and to size_t etc.) isn't a constant expression in SPIR's model.


Another factor to consider, with size_t etc as defined in SPIR, is the usual arithmetic conversions. For instance (assuming a 64-bit long long), sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and would be unsigned if size_t is 64 bits wide. How is this handled?
 

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.


I see, that makes sense.

_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: SPIR Portability Discussion

Villmow, Micah

 

 

From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 2:51 PM
To: Villmow, Micah
Cc: Ouriel, Boaz; [hidden email]; [hidden email]
Subject: Re: [LLVMdev] SPIR Portability Discussion

 

On Wed, Sep 12, 2012 at 2:23 PM, Villmow, Micah <[hidden email]> wrote:

 

 

From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 1:55 PM
To: Ouriel, Boaz
Cc: [hidden email]; [hidden email]
Subject: Re: [LLVMdev] SPIR Portability Discussion

 

On Wed, Sep 12, 2012 at 12:27 PM, Ouriel, Boaz <[hidden email]> 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];

 

Oops, I meant sizeof(void*) - 5.

 

[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

 

My point here was that if sizeof(void*) is an integer constant expression, then this code has a constraint violation if sizeof(void*) is 4 but not if sizeof(void*) is 8 (and my example was intended to be of a global array, not a function-local one, so you can't fall back to treating it as a VLA).

[Villmow, Micah] OpenCL restricts this behavior, so it is illegal.

Another case which might depend on this:

 

enum E {

  a = sizeof(void*) // is this valid?

};

[Villmow, Micah] I will have to think on this one, good example.

 

Based on the behavior you describe above, it looks like sizeof applied to a pointer (and to size_t etc.) isn't a constant expression in SPIR's model.

 

 

Another factor to consider, with size_t etc as defined in SPIR, is the usual arithmetic conversions. For instance (assuming a 64-bit long long), sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and would be unsigned if size_t is 64 bits wide. How is this handled?

[Villmow, Micah] OpenCL C defines ‘int’ to be 32bits irrespective of the host/device bitness. So this would follow the normal integer promotion rules.

 

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.

 

I see, that makes sense.


_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: [cfe-dev] SPIR Portability Discussion

Eli Friedman-2
On Wed, Sep 12, 2012 at 2:58 PM, Villmow, Micah <[hidden email]> wrote:
> Another factor to consider, with size_t etc as defined in SPIR, is the usual
> arithmetic conversions. For instance (assuming a 64-bit long long),
> sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and would be
> unsigned if size_t is 64 bits wide. How is this handled?
>
> [Villmow, Micah] OpenCL C defines ‘int’ to be 32bits irrespective of the
> host/device bitness. So this would follow the normal integer promotion
> rules.

I think you're misunderstanding the issue: the point is, is
"sizeof(int) + -8LL < 0" true or false?

-Eli

_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: SPIR Portability Discussion

Richard Smith-33
In reply to this post by Villmow, Micah
On Wed, Sep 12, 2012 at 2:58 PM, Villmow, Micah <[hidden email]> wrote:

 

 

From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 2:51 PM
To: Villmow, Micah
Cc: Ouriel, Boaz; [hidden email]; [hidden email]


Subject: Re: [LLVMdev] SPIR Portability Discussion

 

On Wed, Sep 12, 2012 at 2:23 PM, Villmow, Micah <[hidden email]> wrote:

 

 

From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 1:55 PM
To: Ouriel, Boaz
Cc: [hidden email]; [hidden email]
Subject: Re: [LLVMdev] SPIR Portability Discussion

 

On Wed, Sep 12, 2012 at 12:27 PM, Ouriel, Boaz <[hidden email]> 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];

 

Oops, I meant sizeof(void*) - 5.

 

[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

 

My point here was that if sizeof(void*) is an integer constant expression, then this code has a constraint violation if sizeof(void*) is 4 but not if sizeof(void*) is 8 (and my example was intended to be of a global array, not a function-local one, so you can't fall back to treating it as a VLA).

[Villmow, Micah] OpenCL restricts this behavior, so it is illegal.

Another case which might depend on this:

 

enum E {

  a = sizeof(void*) // is this valid?

};

[Villmow, Micah] I will have to think on this one, good example.

 

Based on the behavior you describe above, it looks like sizeof applied to a pointer (and to size_t etc.) isn't a constant expression in SPIR's model.

 

 

Another factor to consider, with size_t etc as defined in SPIR, is the usual arithmetic conversions. For instance (assuming a 64-bit long long), sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and would be unsigned if size_t is 64 bits wide. How is this handled?

[Villmow, Micah] OpenCL C defines ‘int’ to be 32bits irrespective of the host/device bitness. So this would follow the normal integer promotion rules.


The value of sizeof(int) isn't what's relevant here. This code implicitly depends on sizeof(sizeof(int)): if size_t is a 32 bit unsigned type, then sizeof(int) + 1LL has a 64-bit signed type. If it's a 64 bit unsigned type, then sizeof(int) + 1LL has a 64-bit *unsigned* type. For instance, the value of "-1LL < sizeof(int) + 1LL" is 1 on 32-bit and 0 on 64-bit.

_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: [cfe-dev] SPIR Portability Discussion

Villmow, Micah
In reply to this post by Eli Friedman-2


> -----Original Message-----
> From: Eli Friedman [mailto:[hidden email]]
> Sent: Wednesday, September 12, 2012 3:22 PM
> To: Villmow, Micah
> Cc: Richard Smith; [hidden email]; [hidden email]
> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>
> On Wed, Sep 12, 2012 at 2:58 PM, Villmow, Micah <[hidden email]>
> wrote:
> > Another factor to consider, with size_t etc as defined in SPIR, is
> the usual
> > arithmetic conversions. For instance (assuming a 64-bit long long),
> > sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and
> would be
> > unsigned if size_t is 64 bits wide. How is this handled?
> >
> > [Villmow, Micah] OpenCL C defines 'int' to be 32bits irrespective of
> the
> > host/device bitness. So this would follow the normal integer
> promotion
> > rules.
>
> I think you're misunderstanding the issue: the point is, is
> "sizeof(int) + -8LL < 0" true or false?
[Villmow, Micah] Yep, I don't see why this is any different than "4 + -8LL < 0".  OpenCL C, and in turn SPIR, defines sizeof(int) == 4. While this might be a problem in C, this isn't an issue in OpenCL since there is no variance in the sizeof(int) across devices.
>
> -Eli



_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: [cfe-dev] SPIR Portability Discussion

Richard Smith-33
On Wed, Sep 12, 2012 at 3:26 PM, Villmow, Micah <[hidden email]> wrote:


> -----Original Message-----
> From: Eli Friedman [mailto:[hidden email]]
> Sent: Wednesday, September 12, 2012 3:22 PM
> To: Villmow, Micah
> Cc: Richard Smith; [hidden email]; [hidden email]
> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>
> On Wed, Sep 12, 2012 at 2:58 PM, Villmow, Micah <[hidden email]>
> wrote:
> > Another factor to consider, with size_t etc as defined in SPIR, is
> the usual
> > arithmetic conversions. For instance (assuming a 64-bit long long),
> > sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and
> would be
> > unsigned if size_t is 64 bits wide. How is this handled?
> >
> > [Villmow, Micah] OpenCL C defines 'int' to be 32bits irrespective of
> the
> > host/device bitness. So this would follow the normal integer
> promotion
> > rules.
>
> I think you're misunderstanding the issue: the point is, is
> "sizeof(int) + -8LL < 0" true or false?
[Villmow, Micah] Yep, I don't see why this is any different than "4 + -8LL < 0".  OpenCL C, and in turn SPIR, defines sizeof(int) == 4. While this might be a problem in C, this isn't an issue in OpenCL since there is no variance in the sizeof(int) across devices.

I think you're still misunderstanding. If size_t is 32 bits, sizeof(int) + -8LL is -4LL, so the comparison produces true. If it's 64 bits, the -8LL promotes to an unsigned long long, sizeof(int) + -8LL is 18446744073709551612ULL, the 0 promotes to 0ULL, and the comparison produces false.

_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: [cfe-dev] SPIR Portability Discussion

Villmow, Micah

 

 

From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 3:30 PM
To: Villmow, Micah
Cc: Eli Friedman; [hidden email]; [hidden email]
Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion

 

On Wed, Sep 12, 2012 at 3:26 PM, Villmow, Micah <[hidden email]> wrote:



> -----Original Message-----
> From: Eli Friedman [mailto:[hidden email]]
> Sent: Wednesday, September 12, 2012 3:22 PM
> To: Villmow, Micah

> Cc: Richard Smith; [hidden email]; [hidden email]
> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>
> On Wed, Sep 12, 2012 at 2:58 PM, Villmow, Micah <[hidden email]>
> wrote:
> > Another factor to consider, with size_t etc as defined in SPIR, is
> the usual
> > arithmetic conversions. For instance (assuming a 64-bit long long),
> > sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and
> would be
> > unsigned if size_t is 64 bits wide. How is this handled?
> >
> > [Villmow, Micah] OpenCL C defines 'int' to be 32bits irrespective of
> the
> > host/device bitness. So this would follow the normal integer
> promotion
> > rules.
>
> I think you're misunderstanding the issue: the point is, is
> "sizeof(int) + -8LL < 0" true or false?

[Villmow, Micah] Yep, I don't see why this is any different than "4 + -8LL < 0".  OpenCL C, and in turn SPIR, defines sizeof(int) == 4. While this might be a problem in C, this isn't an issue in OpenCL since there is no variance in the sizeof(int) across devices.

 

I think you're still misunderstanding. If size_t is 32 bits, sizeof(int) + -8LL is -4LL, so the comparison produces true. If it's 64 bits, the -8LL promotes to an unsigned long long, sizeof(int) + -8LL is 18446744073709551612ULL, the 0 promotes to 0ULL, and the comparison produces false.

[Villmow, Micah] I see now, I think you had a type-o in the previous email, “sizeof(sizeof(int))” should have been size_t(sizeof(int)), which was throwing me off. I view this case as being well defined in SPIR. It can be produced with something like the following:

%0 = call %spir.size_t @__spir_sizet_convert_size_t(i32 0)

%1 = call %spir.size_t @__spir_sizet_convert_size_t(i32 4)

%2 = call %spir.size_t @__spir_sizet_convert_size_t(i64 8)

%3 = call %spir.size_t @__spir_sizet_neg(%spir.size_t %2)

%4 = call %spir.size_t @__spir_sizet_add(%spir.size_t %1, %spir.size_t %3)

%5 = call %spir.size_t @__spir_sizet_cmp(%spir.size_t %4, %spir.size_t %0)

%6 = call i1 %spir.size_t @__spir_size_t_convert_i1(%spir.size_t %5)

 

While this is very verbose, it is possible to handle it correctly. Once you lower the SPIR to LLVMIR and run some basic optimizations, then resulting IR should be equivalent as if were generating LLVMIR directly.

 

Though I’m curious where it states we have to promote -8LL to unsigned long and not signed long, I would have thought it would be signed.\

 

Thanks, Micah

 

 


_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: [cfe-dev] SPIR Portability Discussion

Eli Friedman-2
On Wed, Sep 12, 2012 at 3:40 PM, Villmow, Micah <[hidden email]> wrote:

>
>
>
>
> From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard
> Smith
> Sent: Wednesday, September 12, 2012 3:30 PM
> To: Villmow, Micah
> Cc: Eli Friedman; [hidden email]; [hidden email]
>
>
> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>
>
>
> On Wed, Sep 12, 2012 at 3:26 PM, Villmow, Micah <[hidden email]>
> wrote:
>
>
>
>> -----Original Message-----
>> From: Eli Friedman [mailto:[hidden email]]
>> Sent: Wednesday, September 12, 2012 3:22 PM
>> To: Villmow, Micah
>
>> Cc: Richard Smith; [hidden email]; [hidden email]
>> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>>
>> On Wed, Sep 12, 2012 at 2:58 PM, Villmow, Micah <[hidden email]>
>> wrote:
>> > Another factor to consider, with size_t etc as defined in SPIR, is
>> the usual
>> > arithmetic conversions. For instance (assuming a 64-bit long long),
>> > sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and
>> would be
>> > unsigned if size_t is 64 bits wide. How is this handled?
>> >
>> > [Villmow, Micah] OpenCL C defines 'int' to be 32bits irrespective of
>> the
>> > host/device bitness. So this would follow the normal integer
>> promotion
>> > rules.
>>
>> I think you're misunderstanding the issue: the point is, is
>> "sizeof(int) + -8LL < 0" true or false?
>
> [Villmow, Micah] Yep, I don't see why this is any different than "4 + -8LL <
> 0".  OpenCL C, and in turn SPIR, defines sizeof(int) == 4. While this might
> be a problem in C, this isn't an issue in OpenCL since there is no variance
> in the sizeof(int) across devices.
>
>
>
> I think you're still misunderstanding. If size_t is 32 bits, sizeof(int) +
> -8LL is -4LL, so the comparison produces true. If it's 64 bits, the -8LL
> promotes to an unsigned long long, sizeof(int) + -8LL is
> 18446744073709551612ULL, the 0 promotes to 0ULL, and the comparison produces
> false.
>
> [Villmow, Micah] I see now, I think you had a type-o in the previous email,
> “sizeof(sizeof(int))” should have been size_t(sizeof(int)), which was
> throwing me off. I view this case as being well defined in SPIR. It can be
> produced with something like the following:
>
> %0 = call %spir.size_t @__spir_sizet_convert_size_t(i32 0)
>
> %1 = call %spir.size_t @__spir_sizet_convert_size_t(i32 4)
>
> %2 = call %spir.size_t @__spir_sizet_convert_size_t(i64 8)
>
> %3 = call %spir.size_t @__spir_sizet_neg(%spir.size_t %2)
>
> %4 = call %spir.size_t @__spir_sizet_add(%spir.size_t %1, %spir.size_t %3)
>
> %5 = call %spir.size_t @__spir_sizet_cmp(%spir.size_t %4, %spir.size_t %0)
>
> %6 = call i1 %spir.size_t @__spir_size_t_convert_i1(%spir.size_t %5)

This conversion simply isn't correct: the type of the comparison is
not size_t if size_t is 32 bits.

-Eli

_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: [cfe-dev] SPIR Portability Discussion

Richard Smith-33
In reply to this post by Villmow, Micah
On Wed, Sep 12, 2012 at 3:40 PM, Villmow, Micah <[hidden email]> wrote:

 

 

From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 3:30 PM
To: Villmow, Micah
Cc: Eli Friedman; [hidden email]; [hidden email]


Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion

 

On Wed, Sep 12, 2012 at 3:26 PM, Villmow, Micah <[hidden email]> wrote:



> -----Original Message-----
> From: Eli Friedman [mailto:[hidden email]]
> Sent: Wednesday, September 12, 2012 3:22 PM
> To: Villmow, Micah

> Cc: Richard Smith; [hidden email]; [hidden email]
> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>
> On Wed, Sep 12, 2012 at 2:58 PM, Villmow, Micah <[hidden email]>
> wrote:
> > Another factor to consider, with size_t etc as defined in SPIR, is
> the usual
> > arithmetic conversions. For instance (assuming a 64-bit long long),
> > sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and
> would be
> > unsigned if size_t is 64 bits wide. How is this handled?
> >
> > [Villmow, Micah] OpenCL C defines 'int' to be 32bits irrespective of
> the
> > host/device bitness. So this would follow the normal integer
> promotion
> > rules.
>
> I think you're misunderstanding the issue: the point is, is
> "sizeof(int) + -8LL < 0" true or false?

[Villmow, Micah] Yep, I don't see why this is any different than "4 + -8LL < 0".  OpenCL C, and in turn SPIR, defines sizeof(int) == 4. While this might be a problem in C, this isn't an issue in OpenCL since there is no variance in the sizeof(int) across devices.

 

I think you're still misunderstanding. If size_t is 32 bits, sizeof(int) + -8LL is -4LL, so the comparison produces true. If it's 64 bits, the -8LL promotes to an unsigned long long, sizeof(int) + -8LL is 18446744073709551612ULL, the 0 promotes to 0ULL, and the comparison produces false.

[Villmow, Micah] I see now, I think you had a type-o in the previous email, “sizeof(sizeof(int))” should have been size_t(sizeof(int)), which was throwing me off.


What I wrote was what I meant. The *value* of sizeof(int) is not relevant here, what matters is the precision of its type (or more specifically, its integer conversion rank).

Though I’m curious where it states we have to promote -8LL to unsigned long and not signed long, I would have thought it would be signed.\



_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: [cfe-dev] SPIR Portability Discussion

Villmow, Micah
In reply to this post by Eli Friedman-2


> -----Original Message-----
> From: Eli Friedman [mailto:[hidden email]]
> Sent: Wednesday, September 12, 2012 3:50 PM
> To: Villmow, Micah
> Cc: Richard Smith; [hidden email]; [hidden email]
> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>
> On Wed, Sep 12, 2012 at 3:40 PM, Villmow, Micah <[hidden email]>
> wrote:
> >
> >
> >
> >
> > From: [hidden email] [mailto:[hidden email]] On Behalf Of
> Richard
> > Smith
> > Sent: Wednesday, September 12, 2012 3:30 PM
> > To: Villmow, Micah
> > Cc: Eli Friedman; [hidden email]; [hidden email]
> >
> >
> > Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
> >
> >
> >
> > On Wed, Sep 12, 2012 at 3:26 PM, Villmow, Micah
> <[hidden email]>
> > wrote:
> >
> >
> >
> >> -----Original Message-----
> >> From: Eli Friedman [mailto:[hidden email]]
> >> Sent: Wednesday, September 12, 2012 3:22 PM
> >> To: Villmow, Micah
> >
> >> Cc: Richard Smith; [hidden email]; [hidden email]
> >> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
> >>
> >> On Wed, Sep 12, 2012 at 2:58 PM, Villmow, Micah
> <[hidden email]>
> >> wrote:
> >> > Another factor to consider, with size_t etc as defined in SPIR, is
> >> the usual
> >> > arithmetic conversions. For instance (assuming a 64-bit long
> long),
> >> > sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and
> >> would be
> >> > unsigned if size_t is 64 bits wide. How is this handled?
> >> >
> >> > [Villmow, Micah] OpenCL C defines 'int' to be 32bits irrespective
> of
> >> the
> >> > host/device bitness. So this would follow the normal integer
> >> promotion
> >> > rules.
> >>
> >> I think you're misunderstanding the issue: the point is, is
> >> "sizeof(int) + -8LL < 0" true or false?
> >
> > [Villmow, Micah] Yep, I don't see why this is any different than "4 +
> -8LL <
> > 0".  OpenCL C, and in turn SPIR, defines sizeof(int) == 4. While this
> might
> > be a problem in C, this isn't an issue in OpenCL since there is no
> variance
> > in the sizeof(int) across devices.
> >
> >
> >
> > I think you're still misunderstanding. If size_t is 32 bits,
> sizeof(int) +
> > -8LL is -4LL, so the comparison produces true. If it's 64 bits, the -
> 8LL
> > promotes to an unsigned long long, sizeof(int) + -8LL is
> > 18446744073709551612ULL, the 0 promotes to 0ULL, and the comparison
> produces
> > false.
> >
> > [Villmow, Micah] I see now, I think you had a type-o in the previous
> email,
> > "sizeof(sizeof(int))" should have been size_t(sizeof(int)), which was
> > throwing me off. I view this case as being well defined in SPIR. It
> can be
> > produced with something like the following:
> >
> > %0 = call %spir.size_t @__spir_sizet_convert_size_t(i32 0)
> >
> > %1 = call %spir.size_t @__spir_sizet_convert_size_t(i32 4)
> >
> > %2 = call %spir.size_t @__spir_sizet_convert_size_t(i64 8)
> >
> > %3 = call %spir.size_t @__spir_sizet_neg(%spir.size_t %2)
> >
> > %4 = call %spir.size_t @__spir_sizet_add(%spir.size_t %1,
> %spir.size_t %3)
> >
> > %5 = call %spir.size_t @__spir_sizet_cmp(%spir.size_t %4,
> %spir.size_t %0)
> >
> > %6 = call i1 %spir.size_t @__spir_size_t_convert_i1(%spir.size_t %5)
>
> This conversion simply isn't correct: the type of the comparison is
> not size_t if size_t is 32 bits.
[Villmow, Micah] Sorry, the result of sizet_cmp is i1, I should've looked at the spec for correctness. The point I was trying to make is that in both a 32bit system and a 64bit system, the sequence of instructions should still evaluate to the correct result. The difference is that evaluation does not occur in the frontend, but instead occurs once the SPIR binary is loaded and converted to the device binary. All of these function calls should collapse into a single true/false constant after constant propagation has occurred.
>
> -Eli



_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: [cfe-dev] SPIR Portability Discussion

Eli Friedman-2
On Wed, Sep 12, 2012 at 4:02 PM, Villmow, Micah <[hidden email]> wrote:

>
>
>> -----Original Message-----
>> From: Eli Friedman [mailto:[hidden email]]
>> Sent: Wednesday, September 12, 2012 3:50 PM
>> To: Villmow, Micah
>> Cc: Richard Smith; [hidden email]; [hidden email]
>> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>>
>> On Wed, Sep 12, 2012 at 3:40 PM, Villmow, Micah <[hidden email]>
>> wrote:
>> >
>> >
>> >
>> >
>> > From: [hidden email] [mailto:[hidden email]] On Behalf Of
>> Richard
>> > Smith
>> > Sent: Wednesday, September 12, 2012 3:30 PM
>> > To: Villmow, Micah
>> > Cc: Eli Friedman; [hidden email]; [hidden email]
>> >
>> >
>> > Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>> >
>> >
>> >
>> > On Wed, Sep 12, 2012 at 3:26 PM, Villmow, Micah
>> <[hidden email]>
>> > wrote:
>> >
>> >
>> >
>> >> -----Original Message-----
>> >> From: Eli Friedman [mailto:[hidden email]]
>> >> Sent: Wednesday, September 12, 2012 3:22 PM
>> >> To: Villmow, Micah
>> >
>> >> Cc: Richard Smith; [hidden email]; [hidden email]
>> >> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>> >>
>> >> On Wed, Sep 12, 2012 at 2:58 PM, Villmow, Micah
>> <[hidden email]>
>> >> wrote:
>> >> > Another factor to consider, with size_t etc as defined in SPIR, is
>> >> the usual
>> >> > arithmetic conversions. For instance (assuming a 64-bit long
>> long),
>> >> > sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and
>> >> would be
>> >> > unsigned if size_t is 64 bits wide. How is this handled?
>> >> >
>> >> > [Villmow, Micah] OpenCL C defines 'int' to be 32bits irrespective
>> of
>> >> the
>> >> > host/device bitness. So this would follow the normal integer
>> >> promotion
>> >> > rules.
>> >>
>> >> I think you're misunderstanding the issue: the point is, is
>> >> "sizeof(int) + -8LL < 0" true or false?
>> >
>> > [Villmow, Micah] Yep, I don't see why this is any different than "4 +
>> -8LL <
>> > 0".  OpenCL C, and in turn SPIR, defines sizeof(int) == 4. While this
>> might
>> > be a problem in C, this isn't an issue in OpenCL since there is no
>> variance
>> > in the sizeof(int) across devices.
>> >
>> >
>> >
>> > I think you're still misunderstanding. If size_t is 32 bits,
>> sizeof(int) +
>> > -8LL is -4LL, so the comparison produces true. If it's 64 bits, the -
>> 8LL
>> > promotes to an unsigned long long, sizeof(int) + -8LL is
>> > 18446744073709551612ULL, the 0 promotes to 0ULL, and the comparison
>> produces
>> > false.
>> >
>> > [Villmow, Micah] I see now, I think you had a type-o in the previous
>> email,
>> > "sizeof(sizeof(int))" should have been size_t(sizeof(int)), which was
>> > throwing me off. I view this case as being well defined in SPIR. It
>> can be
>> > produced with something like the following:
>> >
>> > %0 = call %spir.size_t @__spir_sizet_convert_size_t(i32 0)
>> >
>> > %1 = call %spir.size_t @__spir_sizet_convert_size_t(i32 4)
>> >
>> > %2 = call %spir.size_t @__spir_sizet_convert_size_t(i64 8)
>> >
>> > %3 = call %spir.size_t @__spir_sizet_neg(%spir.size_t %2)
>> >
>> > %4 = call %spir.size_t @__spir_sizet_add(%spir.size_t %1,
>> %spir.size_t %3)
>> >
>> > %5 = call %spir.size_t @__spir_sizet_cmp(%spir.size_t %4,
>> %spir.size_t %0)
>> >
>> > %6 = call i1 %spir.size_t @__spir_size_t_convert_i1(%spir.size_t %5)
>>
>> This conversion simply isn't correct: the type of the comparison is
>> not size_t if size_t is 32 bits.
> [Villmow, Micah] Sorry, the result of sizet_cmp is i1, I should've looked at the spec for correctness. The point I was trying to make is that in both a 32bit system and a 64bit system, the sequence of instructions should still evaluate to the correct result. The difference is that evaluation does not occur in the frontend, but instead occurs once the SPIR binary is loaded and converted to the device binary. All of these function calls should collapse into a single true/false constant after constant propagation has occurred.

The problem isn't the type of the result of __spir_sizet_cmp, it's the
type of the operands.  Assuming size_t is 32 bits, the type of the
operands in "sizeof(int) + -8LL < 0" is long long, which is not the
same as size_t.

-Eli

_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
Reply | Threaded
Open this post in threaded view
|

Re: [cfe-dev] SPIR Portability Discussion

Villmow, Micah
In reply to this post by Richard Smith-33

Ok, thanks for pointing out the location in the spec. There are some good points here, especially in the rank of ‘size_t’ that the SPIR WG will have to decide.

 

However, one thing to keep in mind, this code is by definition non-portable and so SPIR is not trying to make it more portable. This might be one of those cases where implementation defined behavior of C does not allow the program to be portable.


Thanks.

 

From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 4:03 PM
To: Villmow, Micah
Cc: Eli Friedman; [hidden email]; [hidden email]
Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion

 

On Wed, Sep 12, 2012 at 3:40 PM, Villmow, Micah <[hidden email]> wrote:

 

 

From: [hidden email] [mailto:[hidden email]] On Behalf Of Richard Smith
Sent: Wednesday, September 12, 2012 3:30 PM
To: Villmow, Micah
Cc: Eli Friedman; [hidden email]; [hidden email]


Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion

 

On Wed, Sep 12, 2012 at 3:26 PM, Villmow, Micah <[hidden email]> wrote:



> -----Original Message-----
> From: Eli Friedman [mailto:[hidden email]]
> Sent: Wednesday, September 12, 2012 3:22 PM
> To: Villmow, Micah

> Cc: Richard Smith; [hidden email]; [hidden email]
> Subject: Re: [cfe-dev] [LLVMdev] SPIR Portability Discussion
>
> On Wed, Sep 12, 2012 at 2:58 PM, Villmow, Micah <[hidden email]>
> wrote:
> > Another factor to consider, with size_t etc as defined in SPIR, is
> the usual
> > arithmetic conversions. For instance (assuming a 64-bit long long),
> > sizeof(int) + 1LL would be signed if size_t is 32 bits wide, and
> would be
> > unsigned if size_t is 64 bits wide. How is this handled?
> >
> > [Villmow, Micah] OpenCL C defines 'int' to be 32bits irrespective of
> the
> > host/device bitness. So this would follow the normal integer
> promotion
> > rules.
>
> I think you're misunderstanding the issue: the point is, is
> "sizeof(int) + -8LL < 0" true or false?

[Villmow, Micah] Yep, I don't see why this is any different than "4 + -8LL < 0".  OpenCL C, and in turn SPIR, defines sizeof(int) == 4. While this might be a problem in C, this isn't an issue in OpenCL since there is no variance in the sizeof(int) across devices.

 

I think you're still misunderstanding. If size_t is 32 bits, sizeof(int) + -8LL is -4LL, so the comparison produces true. If it's 64 bits, the -8LL promotes to an unsigned long long, sizeof(int) + -8LL is 18446744073709551612ULL, the 0 promotes to 0ULL, and the comparison produces false.

[Villmow, Micah] I see now, I think you had a type-o in the previous email, “sizeof(sizeof(int))” should have been size_t(sizeof(int)), which was throwing me off.

 

What I wrote was what I meant. The *value* of sizeof(int) is not relevant here, what matters is the precision of its type (or more specifically, its integer conversion rank).

 

Though I’m curious where it states we have to promote -8LL to unsigned long and not signed long, I would have thought it would be signed.\

 

C99 6.3.1.8/1.


_______________________________________________
LLVM Developers mailing list
[hidden email]         http://llvm.cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
12