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

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

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

...

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

...

[Guy Benyei] I see no usable way to make sizeof(void*) or sizeof(size_t) compile time constants while creating a portable intermediate representation. You're right, some existing code might be non-SPIR-compatible; however I think that's a reasonable price to pay when using SPIR. Clang has very good capabilities for providing clear and understandable error messages, so in this case the developer would be informed so he could choose a more portable solution.

Thanks
   Guy Benyei