Archives Discussions

[toastedcrumpets]

This kernel fails to compile on OpenCL 1.0 ATI-Stream-v2.0.0. The error is arising from the += operator on the local memory float4. Is this an issue or is it not allowed by the standard?

The build error is

Stored value type does not match pointer operand type!
  store <4 x float> addrspace(3)* %arrayidx, float** %tmp float*Both operands to a binary operator are not of the same type!
  %tmp5 = fadd <4 x float> %tmp3, float %tmp4     ; <<4 x float>> [#uses=1]Stored value type does not match pointer operand type!
  store <4 x float> %tmp5, float* %tmp1 floatBroken module found, compilation aborted!
Stack dump:
0.    Running pass 'Module Verifier' on function '@__OpenCL_testKernel_kernel'
Aborted

Thanks,

Marcus

__kernel void testKernel(__local float4* localVec) { float X = 12; localVec[get_local_id(0)].x += X; }

[genaganna]

Originally posted by: toastedcrumpets This kernel fails to compile on OpenCL 1.0 ATI-Stream-v2.0.0. The error is arising from the += operator on the local memory float4. Is this an issue or is it not allowed by the standard?

The build error is

Stored value type does not match pointer operand type!   store <4 x float> addrspace(3)* %arrayidx, float** %tmp float*Both operands to a binary operator are not of the same type!   %tmp5 = fadd <4 x float> %tmp3, float %tmp4     ; <<4 x float>> [#uses=1]Stored value type does not match pointer operand type!   store <4 x float> %tmp5, float* %tmp1 floatBroken module found, compilation aborted! Stack dump: 0.    Running pass 'Module Verifier' on function '@__OpenCL_testKernel_kernel' Aborted

Toastedcrumpets,

             Thanks for reporting this.  I hope you have a workaround already like following

localVec[get_local_id(0)].x = localVec[get_local_id(0)].x + X;

[toastedcrumpets]

Already done,

I'd also like to say the speed of your OpenCL platform (GPU and CPU) is pleasantly suprising in our MD benchmarks. So, thanks!