[cfe-dev] SPIR Review Status: after Introduction and 32bits vs. 64bits discussions

Benyei, Guy guy.benyei at intel.com
Wed Sep 19 05:43:22 PDT 2012

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): 
              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.


    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.

More information about the cfe-dev mailing list