Ouriel, Boaz
2012-Sep-14 14:14 UTC
[LLVMdev] SPIR Review Status: after Introduction and 32bits vs. 64bits discussions
Hi All,
I have made an attempt to summarize the different comments that were raised so
far.
An answer has been provided inside. Please let me know if I missed any comments
or got them wrong.
*** Hal Finkel: Handling FP_CONTRACT ***
*****Comment: The current specification provides a mechanism for handling
FP_CONTRACT, but does so only at the module level. After much debate,
we have adopted and implemented a way
of handling FP_CONTRACT in clang/LLVM through the llvm.fmuladd intrinsic
(currently in trunk).
I suggest that the proposed
spir.disable.FP CONTRACT metadata be replaced with the current
llvm.fmuladd-based mechanism.
Answer: We were not aware of this, we will definitely adopt
this approach.
*** James Molloy: Why new calling conventions?***
*** Comment: What are their semantics? And what is their purpose? Why not use
metadata instead?
Answer: We still hold the opinion that this is the right way to go.
Our plan is to provide additional explanation
about the cc semantics and purpose and try to reach a final decision on whether
they are required or not.
*** James Molloy: Why disallow type conversion for vector types?***
***Comment: This might cause some LLVM optimizations to generate an invalid
SPIR module which
Answer: Type conversions in OpenCL between vector types is done
via builtin functions and not via implicit conversions, so there is no OpenCL
code that can generate these conversions
directly(OpenCL spec 6.2.1). In order to be portable, library functions cannot
be lowered to their IR equivalent until
after the device is known. This restriction is
not likely to change because of the importance that as many OpenCL
implementations support SPIR.
The implication is that a SPIR optimizer will
need to rule out such optimizations.
As a side note, we are running a check in Khronos
if this restriction can be removed.
*** Richard Smith, Eli Friedman & Nadav Rotem: Portability Issues ***
*****comment 1: int does_this_compile[sizeof(void*) - 3];
Answer: We are discussing this internally and will provide
an answer soon.
****comment 2: struct how_do_you_represent_this_in_IR {
int a : 1;
int b : sizeof(void*) * 4;
};
Answer: Bitfields are disallowed in OpenCL “C”
****comment 3: 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];
Answer: 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.
*****comment 4:
// 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;
}
}
Answer: We are discussing this an provide an answer soon.
*****comment 5: What about this case?
enum E {
a = sizeof(void*) // is this valid?
};
Answer: we are discussing this and will provide an answer
soon.
****comment 6: What is the rank of ‘size_t’?
example: is "sizeof(int) + -8LL < 0" true or
false?
Answer: we are discussing this and will provide an answer
soon.
****comment 7: Why can't we always make size_t 64 bits wide?
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.
Answer: we are discussing this and will provide an answer
soon.
I appreciate all of the good feedback,
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.
Benyei, Guy
2012-Sep-19 12:43 UTC
[LLVMdev] [cfe-dev] SPIR Review Status: after Introduction and 32bits vs. 64bits discussions
Ouriel, Boaz wrote:
...
*** Richard Smith, Eli Friedman & Nadav Rotem: Portability Issues ***
*****comment 1: int does_this_compile[sizeof(void*) - 3];
Answer: We are discussing this internally and will provide
an answer soon.
****comment 2: struct how_do_you_represent_this_in_IR {
int a : 1;
int b : sizeof(void*) * 4;
};
Answer: Bitfields are disallowed in OpenCL “C”
****comment 3: 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];
Answer: 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.
*****comment 4:
// 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;
}
}
Answer: We are discussing this an provide an answer soon.
[Guy Benyei] Some inherently non-portable code snippets won't be supported
in SPIR. IMO, these cases should be detected during compilation (and we
don't define anything about compilation in the SPIR spec). Especially, when
a given source code should raise compilation error in one architecture, and pass
in the other (32/64), the result must be a compilation error.
*****comment 5: What about this case?
enum E {
a = sizeof(void*) // is this valid?
};
Answer: we are discussing this and will provide an answer
soon.
[Guy Benyei] Same goes here - this source is not functionally portable.
****comment 6: What is the rank of ‘size_t’?
example: is "sizeof(int) + -8LL < 0" true or
false?
Answer: we are discussing this and will provide an answer
soon.
[Guy Benyei] We discussed this case a lot, and IMO there are two possible
solutions: we can either assign a rank to size_t, s.t. rank(long) <
rank(ptrdiff_t) < rank(size_t) < rank(ulong), or we can simply disallow
the case where the usual arithmetic conversions should decide between types
which would lead to different decision in 64bit and 32bit architectures.
Personally I don't really like the option of assigning a rank to size_t,
since it would change the behavior in some corner cases, and lead to results
that were not expected by the developer.
****comment 7: Why can't we always make size_t 64 bits wide?
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.
Answer: we are discussing this and will provide an answer
soon.
[Guy Benyei] First and most importantly, OpenCL embedded profile doesn't
require support for 64 bit integers. Making all size_t and ptrdiff_t 64 bit
would disallow the usage of SPIR in some embedded systems. Secondly, I'm not
sure target specific optimizations will be able to guess the bit width
correctly. Provided a kernel which uses a buffer and accesses elements from
get_global_id(0)*8 to get_global_id(0)*8+7, it won't be sufficient to assume
that get_global_id(0) returns a 32bit value, since this makes get_global_id(0)*8
is a 35bit value, and it would disable these optimizations. Assuming
get_global_id(0) returns less than 32 bits seems to me wrong too. Not being able
to optimize these size_t values in 32 bit architectures would cause huge
performance degradation, since emulation of 64 bit operations in 32 bit
architectures would be quite painful.
...
Thanks
Guy Benyei
---------------------------------------------------------------------
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.
Eli Friedman
2012-Sep-19 20:02 UTC
[LLVMdev] [cfe-dev] SPIR Review Status: after Introduction and 32bits vs. 64bits discussions
On Wed, Sep 19, 2012 at 5:43 AM, Benyei, Guy <guy.benyei at intel.com> wrote:> Ouriel, Boaz wrote: > > > ... > > > *** Richard Smith, Eli Friedman & Nadav Rotem: Portability Issues *** > > *****comment 1: int does_this_compile[sizeof(void*) - 3]; > Answer: We are discussing this internally and will provide an answer soon. > > ****comment 2: struct how_do_you_represent_this_in_IR { > int a : 1; > int b : sizeof(void*) * 4; > }; > Answer: Bitfields are disallowed in OpenCL “C” > > ****comment 3: 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]; > > Answer: 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. > > *****comment 4: > // 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; > } > } > Answer: We are discussing this an provide an answer soon. > > > [Guy Benyei] Some inherently non-portable code snippets won't be supported in SPIR. IMO, these cases should be detected during compilation (and we don't define anything about compilation in the SPIR spec). Especially, when a given source code should raise compilation error in one architecture, and pass in the other (32/64), the result must be a compilation error. > > > *****comment 5: What about this case? > enum E { > a = sizeof(void*) // is this valid? > }; > Answer: we are discussing this and will provide an answer soon. > > [Guy Benyei] Same goes here - this source is not functionally portable.Okay... how exactly do you plan to detect it? The only reasonable way I can think of is banning sizeof(void*) and friends in integer constant expressions, which will work, but might make existing code non-SPIR-compatible.> ****comment 6: What is the rank of ‘size_t’? > example: is "sizeof(int) + -8LL < 0" true or false? > Answer: we are discussing this and will provide an answer soon. > > [Guy Benyei] We discussed this case a lot, and IMO there are two possible solutions: we can either assign a rank to size_t, s.t. rank(long) < rank(ptrdiff_t) < rank(size_t) < rank(ulong), or we can simply disallow the case where the usual arithmetic conversions should decide between types which would lead to different decision in 64bit and 32bit architectures. Personally I don't really like the option of assigning a rank to size_t, since it would change the behavior in some corner cases, and lead to results that were not expected by the developer.Banning the conversions is reasonable, and should be straightforward to detect. -Eli
Possibly Parallel Threads
- [LLVMdev] SPIR Review Status: after Introduction and 32bits vs. 64bits discussions
- [LLVMdev] [cfe-dev] SPIR Review Status: after Introduction and 32bits vs. 64bits discussions
- [LLVMdev] SPIR Portability Discussion
- [LLVMdev] SPIR Portability Discussion
- [LLVMdev] SPIR: Answers to the issues raised so far