AnsweredAssumed Answered

CLK_LOCAL_MEM_FENCE wipes values at indices of >63 in local arrays

Question asked by animal_magic on Sep 8, 2014
Latest reply on Sep 12, 2014 by maxdz8

I've run into a weird bug in a kernel. I reproduced the bug with a simple kernel:

 

//data is initialized as {0, 0, 0}
//x has size 128 and is not initialized to anything
__kernel void testBarrier(__global int *data, __local uint* x) {
  x[0] = 1;
  x[63] = 1;
  x[64] = 1;

  barrier(CLK_LOCAL_MEM_FENCE);

  data[0] = x[0];
  data[1] = x[63];
  data[2] = x[64];
}

 

I then print off data on the host side and I see

   1, 1, 0


If I comment out "barrier(CLK_LOCAL_MEM_FENCE);", then I see

   1, 1, 1


The number of work items and work group size doesn't seem to matter. I even tried a single work item. I'm pretty sure I'm using the latest drivers and APP SDK, but I can double-check if it matters.

Any idea as to what's going on?


Outcomes