Hi All,
Here are answers to the questions / comments that were raised so far during the
SPIR discussions.
**** A general clarification regarding sizeof ****
In SPIR, sizeof in an integer constant expression is illegal.
The reason behind it is that once the device is no longer known, the width of
the unsigned integer that represents the size_t type is no longer known.
The C spec states in section 6.5.3.4.4 that "The value of the result is
implementation-defined, and its type(an unsigned integer type) is size_t,
defined in <stddef.h> (and other headers)."
SPIR defines size_t as a unsized unsigned integer and as such cannot be known
until the SPIR binary is lowered to the device.
At that time, SPIR is known and as such can be lowered correctly to its type.
*** 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: The semantics of the new calling conventions are as follows:
spirfunc calling convention -
All arguments are passed as is to the function with the
exception of structures, no lowering or expansion is allowed.
Structures are passed as a pointer to struct with the byval
attribute set.
Functions marked with spirfunc calling convention can only be
called by another function marked with either spirfunc or spirkrnl calling
conventions.
Functions marked with spirfunc calling convention can only
call functions marked with spirfunc or spirkrnl calling conventions.
Functions marked with spirfunc calling convention can only
have zero or 1 return values.
Functions marked with spirfunc calling convention are only
visible to the device.
Variable arguments are not allowed, except for printf.
This calling convention does not specify the how it is lowered
to registers or how the return value is specified.
spirkrnl calling convention -
Inherits the definition of spirfunc calling convention.
Functions marked with spirkrnl calling convention cannot have
any return values.
Functions marked with spirkrnl calling convention cannot have
variable arguments.
Functions marked with spirkrnl calling convention can be called
by the host.
Functions marked with spirkrnl calling convention are externally
visible.
*** 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. When OpenCL C was
developed, there were many different implementations that supported vector
casts, but they were not conversion casts(e.g. Altivec/Cell C extensions).
OpenCL requires explicit vector conversion
functions, and as such SPIR inherits those rules.
Since many compilers already had the vector casts
that weren't conversions, the function calls were required so that code
could not be written expecting the prior behavior of other compilers in OpenCL.
The implication is that a SPIR optimizer will
need to rule out such optimizations.
*** Richard Smith, Eli Friedman & Nadav Rotem: Portability Issues ***
*****comment 1: int does_this_compile[sizeof(void*) - 3];
Answer: In this scenario the sizeof(void*) is no longer a
frontend compile time constant. This case is not supported, and a compile time
error should be raised by a SPIR frontend
****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: In this case the expression is no longer a front-end
constant expression. This will be a SPIR frontend compile time error.
*****comment 5: What about this case?
enum E {
a = sizeof(void*) // is this valid?
};
Answer: sizeof(void*) is not a frontend compile time
constant, so you get a SPIR frontend compile error.
****comment 6: What is the rank of 'size_t'?
example: is "sizeof(int) + -8LL < 0" true or
false?
Answer: The rank of size_t is either uint for 32bit devices
and ulong for 64bit devices (int < uint ( == size_t for 32bit) < long (
== size_t for 64bit) < ulong)
This means that the only ambiguity is when
comparing it to long. Whenever this ambiguity is encountered a frontend
compilation error should occur.
****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: First and most importantly, OpenCL embedded profile
doesn't require support for 64 bit integers (it is optional).
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
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.
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20120927/f3e574e7/attachment.html>
