If I pass restrict to the params as:
__global float2 *restrict a,
const __global float4 *restrict b
produces incorrect results with Cat.11.3/5750/Win7_64. If I remove the restrict keyword works ok.
__kernel __attribute__((reqd_work_group_size(WORK_WIDTH,WORK_HEIGHT,1))) /* work width=16, work height=8 */ void MyKernel ( __global float2 *a, const __global float4 *b ) { const size_t g1D = get_global_id(1U)*GRID_WIDTH + get_global_id(0U); /* grid width=256U */ if ( b[g1D].x>0.0f ) { a[g1D].x += 1.0f; } } /* note: I use float4 because I assign YZW components for other kernel */