<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=us-ascii">
<meta name="Generator" content="Microsoft Word 14 (filtered medium)">
<style><!--
/* Font Definitions */
@font-face
        {font-family:Calibri;
        panose-1:2 15 5 2 2 2 4 3 2 4;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
        {margin:0in;
        margin-bottom:.0001pt;
        text-align:right;
        direction:rtl;
        unicode-bidi:embed;
        font-size:11.0pt;
        font-family:"Calibri","sans-serif";}
a:link, span.MsoHyperlink
        {mso-style-priority:99;
        color:blue;
        text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
        {mso-style-priority:99;
        color:purple;
        text-decoration:underline;}
p.MsoPlainText, li.MsoPlainText, div.MsoPlainText
        {mso-style-priority:99;
        mso-style-link:"Plain Text Char";
        margin:0in;
        margin-bottom:.0001pt;
        text-align:right;
        direction:rtl;
        unicode-bidi:embed;
        font-size:11.0pt;
        font-family:"Calibri","sans-serif";}
span.EmailStyle17
        {mso-style-type:personal-compose;
        font-family:"Calibri","sans-serif";
        color:windowtext;}
span.PlainTextChar
        {mso-style-name:"Plain Text Char";
        mso-style-priority:99;
        mso-style-link:"Plain Text";
        font-family:"Calibri","sans-serif";}
.MsoChpDefault
        {mso-style-type:export-only;
        font-family:"Calibri","sans-serif";}
@page WordSection1
        {size:8.5in 11.0in;
        margin:1.0in 1.25in 1.0in 1.25in;}
div.WordSection1
        {page:WordSection1;}
--></style><!--[if gte mso 9]><xml>
<o:shapedefaults v:ext="edit" spidmax="1026" />
</xml><![endif]--><!--[if gte mso 9]><xml>
<o:shapelayout v:ext="edit">
<o:idmap v:ext="edit" data="1" />
</o:shapelayout></xml><![endif]-->
</head>
<body lang="EN-US" link="blue" vlink="purple">
<div class="WordSection1">
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
Hi All,<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<o:p> </o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
Here are answers to the questions / comments that were raised so far during the SPIR discussions.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<o:p> </o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
**** A general clarification regarding sizeof ****<o:p></o:p></p>
<p class="MsoNormal" style="text-align:left;direction:ltr;unicode-bidi:embed">In SPIR, sizeof in an integer constant expression is illegal.
<o:p></o:p></p>
<p class="MsoNormal" style="text-align:left;direction:ltr;unicode-bidi:embed">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.  <o:p></o:p></p>
<p class="MsoNormal" style="text-align:left;direction:ltr;unicode-bidi:embed">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)."<o:p></o:p></p>
<p class="MsoNormal" style="text-align:left;direction:ltr;unicode-bidi:embed">SPIR defines size_t as a unsized unsigned integer and as such cannot be known until the SPIR binary is lowered to the device.
<o:p></o:p></p>
<p class="MsoNormal" style="text-align:left;direction:ltr;unicode-bidi:embed">At that time, SPIR is known and as such can be lowered correctly to its type.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
*** Hal Finkel: Handling FP_CONTRACT ***<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
     *****Comment: The current specification provides a mechanism for handling FP_CONTRACT, but does so only at the module level. After much debate,<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                         we have adopted and implemented a way of handling FP_CONTRACT in clang/LLVM through the llvm.fmuladd intrinsic (currently in trunk).
<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                                         I suggest that the proposed spir.disable.FP CONTRACT metadata be replaced with the current llvm.fmuladd-based mechanism.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span lang="HE" dir="RTL" style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                    Answer: We were not aware of this, we will definitely adopt this approach.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
*** James Molloy: Why new calling conventions?***<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span lang="HE" dir="RTL" style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>   *** Comment: What are their semantics? And what is their purpose? Why not use metadata instead?<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
            Answer: The semantics of the new calling conventions are as follows:<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>           spirfunc calling convention -<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                  All arguments are passed as is to the function with the exception of structures, no lowering or expansion is allowed.<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                  Structures are passed as a pointer to struct with the byval attribute set.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                  Functions marked with spirfunc calling convention can only be called by another function marked with either spirfunc or spirkrnl calling conventions.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                  Functions marked with spirfunc calling convention can only call functions marked with spirfunc or spirkrnl calling conventions.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                  Functions marked with spirfunc calling convention can only have zero or 1 return values.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                  Functions marked with spirfunc calling convention are only visible to the device.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                  Variable arguments are not allowed, except for printf.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                  This calling convention does not specify the how it is lowered to registers or how the return value is specified.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
            spirkrnl calling convention -<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                 Inherits the definition of spirfunc calling convention.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                Functions marked with spirkrnl calling convention cannot have any return values.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                Functions marked with spirkrnl calling convention cannot have variable arguments.<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                Functions marked with spirkrnl calling convention can be called by the host.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                Functions marked with spirkrnl calling convention are externally visible.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<o:p> </o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
*** James Molloy: Why disallow type conversion for vector types?***<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
     ***Comment: This might cause some LLVM optimizations to generate an invalid SPIR module which<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span lang="HE" dir="RTL" style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>              Answer: Type conversions in OpenCL between vector types is done via builtin functions and not via implicit conversions, so there is no OpenCL<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                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<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                               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). <o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                               OpenCL requires explicit vector conversion functions, and as such SPIR inherits those rules.
<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                               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.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                               The implication is that a SPIR optimizer will need to rule out such optimizations.
<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
*** Richard Smith, Eli Friedman & Nadav Rotem: Portability Issues ***<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
         *****comment 1: int does_this_compile[sizeof(void*) - 3];<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                      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<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
        ****comment 2: struct how_do_you_represent_this_in_IR {<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                                  int a : 1;<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                  int b : sizeof(void*) * 4;<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="RTL"></span><span lang="HE" dir="RTL" style="font-size:10.5pt;font-family:"Arial","sans-serif""><span dir="RTL"></span>                        };<o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                         Answer: Bitfields are disallowed in OpenCL “C”<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
              ****comment 3:  How do you perform record layout if the size of a pointer is unknown? For instance:<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span lang="HE" dir="RTL" style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                struct A {<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                    int *p;<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                    int n;<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                } a;<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>               int arr[offsetof(A, n) - 3]; // or, int arr[(char*)&a.n - (char*)&a.p - 3];<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                       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.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
        *****comment 4: <o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                  // We're required to diagnose this iff sizeof(size_t) != 4. <span lang="HE" dir="RTL">
<o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                 extern int x[20]; <o:p>
</o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                 int x[sizeof(size_t) * 5;<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>               // We're required to diagnose this iff sizeof(size_t) == 4.
<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                void f(int x) {<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                       switch(x) { <o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                          case 4: <o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                          case sizeof(size_t): <o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                break; <o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="RTL"></span><span lang="HE" dir="RTL" style="font-size:10.5pt;font-family:"Arial","sans-serif""><span dir="RTL"></span>                    }
</span><span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="RTL"></span><span lang="HE" dir="RTL" style="font-size:10.5pt;font-family:"Arial","sans-serif""><span dir="RTL"></span>               }</span><span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
              Answer:  In this case the expression is no longer a front-end constant expression. This will be a SPIR frontend compile time error.
<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
       *****comment 5: What about this case?<span lang="AR-SA" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                   enum E {<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                         a = sizeof(void*) // is this valid?<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="RTL"></span><span lang="HE" dir="RTL" style="font-size:10.5pt;font-family:"Arial","sans-serif""><span dir="RTL"></span>                   };<o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                   Answer: sizeof(void*) is not a frontend compile time constant, so you get a SPIR frontend compile error.<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
       ****comment 6: What is  the rank of ‘size_t’? <o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                  example: is "sizeof(int) + -8LL < 0" true or false?<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                   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)<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                     This means that the only ambiguity is when comparing it to long. Whenever this ambiguity is encountered a frontend compilation error should occur.<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p> </o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
       ****comment 7: Why can't we always make size_t 64 bits wide? <o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                          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. 
<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                           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). 
<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                           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.<span lang="HE" dir="RTL"><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<span dir="LTR"></span><span dir="LTR"></span>                  Answer:  First and most importantly, OpenCL embedded profile doesn't require support for 64 bit integers (it is optional).<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                     Making all size_t and ptrdiff_t 64 bit would disallow the usage of SPIR in some embedded systems.
<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                     Secondly, I'm not sure target specific optimizations will be able to guess the bit width correctly.
<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                     Provided a kernel which uses a buffer and accesses elements from get_global_id(0)*8 to get_global_id(0)*8+7,
<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                     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.
<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                     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,
<o:p></o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
                                    since emulation of 64 bit operations in 32 bit architectures would be quite painful<span style="font-size:10.5pt;font-family:"Arial","sans-serif""><o:p></o:p></span></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
<o:p> </o:p></p>
<p class="MsoPlainText" style="text-align:left;direction:ltr;unicode-bidi:embed">
I appreciate all of the good feedback,<o:p></o:p></p>
<p class="MsoNormal" style="text-align:left;direction:ltr;unicode-bidi:embed">Boaz<o:p></o:p></p>
</div>
<p>---------------------------------------------------------------------<br>
Intel Israel (74) Limited</p>

<p>This e-mail and any attachments may contain confidential material for<br>
the sole use of the intended recipient(s). Any review or distribution<br>
by others is strictly prohibited. If you are not the intended<br>
recipient, please contact the sender and delete all copies.</p></body>
</html>