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

[Ouriel, Boaz]

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.