2 Replies Latest reply on Mar 10, 2010 3:17 PM by northwind

    Compiler crash on named address space assignment by addition

    northwind
      Compiler crashes when compiling kernel that accepts named address space pointer and performs an assignment by addition

      Hey all,

      I've found an interesting compiler bug. Given a specific set of circumstances (which I will elaborate momentarily),  the compiler crashes when performing an assignment by addition (+=) on a float vector type. As well, given the same set of circumstances, an assignment by subtraction, multiplication or division will also crash the compiler.

      The requirements for the crash are not very strict but they must include the following:

      1) The lvalue being operated on must be stored in a named space and therefore must also be a pointer.

      2) The lvalue must be employing a component selector for assignment.

      3) On the right side of the assignment-by operation, a multiplication by a scalar must occur. How this scalar is obtained is irrelevant; it could be a scalar literal, it could be a single component from a larger vector or it could be a scalar variable.

      If all 3 of these conditions are met, the compiler will crash and throw up a message that looks something like the following:

       


      Stored value type does not match pointer operand type!
        store <4 x float> addrspace(1)* %arrayidx, <3 x float>** %tmp1 <3 x float>*Both operands to a binary operator are not of the same type!
        %tmp9 = fmul <4 x float> %tmp4, <3 x float> %conv8 ; <<4 x float>> [#uses=1]Stored value type does not match pointer operand type!
        store <4 x float> %tmp9, <3 x float>* %tmp2 <3 x float>Broken module found, compilation aborted!

      This application has requested the Runtime to terminate it in an unusual way.
      Please contact the application's support team for more information.


       

      This was generated from compiling the following code:

       


      __kernel void hello (__global float4* float_array)
      {
          float_array[0].xyz *= (float4)(0.0f).x;
      }


       

      I'm fairly confident that this has not been mentioned elsewhere but should I be wrong in this regard, feel free to close this thread.

      I've attached more code that will also cause this crash and its associated crash report from the compiler.

      I should note that I am performing compilation on an Intel Core2 Duo 1.6ghz intended to be used as an opencl cpu device. From my experiments thus far, the NVidia compiler, compiling for an 8800GTX, does not suffer from this issue. All the code that I have shown here compiles without trouble on the NVidia platform.

      I've listed what I think are useful crash cases; if you need me to produce other ones, please let me know. As well, note that while I have not shown use of the -= and /= operators, they too trigger a crash under the same expressions.

      Expression 1: ======= float_array[0].xyz += (float4)( 1.0f, 2.0f, 3.0f, 4.0f ).xyz * 5.0f; Crash String 1: ======== Stored value type does not match pointer operand type! store <4 x float> addrspace(1)* %arrayidx, <3 x float>** %tmp1 <3 x float>*Both operands to a binary operator are not of the same type! %tmp13 = fadd <4 x float> %tmp4, <3 x float> %tmp12 ; <<4 x float>> [#uses=1]Stored value type does not match pointer operand type! store <4 x float> %tmp13, <3 x float>* %tmp2 <3 x float>Broken module found, compilation aborted! This application has requested the Runtime to terminate it in an unusual way. Please contact the application's support team for more information. Expression 2: ======== float4 otherVar = { -76.21f, 1.0f, 5.8732f, -9.31f }; float_array[0].xyzw += otherVar.xyzw * 0.0f; Crash String 2: ========= Stored value type does not match pointer operand type! store <4 x float> addrspace(1)* %arrayidx, <4 x float>** %tmp <4 x float>*Broken module found, compilation aborted! This application has requested the Runtime to terminate it in an unusual way. Please contact the application's support team for more information. Expression 3: ======== DECLARATION >> __global float8* float_array float4 otherVar = (float4)(0.0f); float_array[0].s7264 += otherVar * 0; Crash String 3: ========= Stored value type does not match pointer operand type! store <8 x float> addrspace(1)* %arrayidx, <4 x float>** %tmp1 <4 x float>*Both operands to a binary operator are not of the same type! %tmp8 = fadd <8 x float> %tmp5, <4 x float> %tmp7 ; <<8 x float>> [#uses=1]Stored value type does not match pointer operand type! store <8 x float> %tmp8, <4 x float>* %tmp2 <4 x float>Broken module found, compilation aborted! This application has requested the Runtime to terminate it in an unusual way. Please contact the application's support team for more information. Expression 4: ======== DECLARATION >> __local float8* float_array float8 otherVar = (float8)(0.0f); float_array[0].s01234567 += otherVar * 0; Crash String 4: ========= Stored value type does not match pointer operand type! store <8 x float> addrspace(3)* %arrayidx, <8 x float>** %tmp1 <8 x float>*Broken module found, compilation aborted! This application has requested the Runtime to terminate it in an unusual way. Please contact the application's support team for more information.