4 Replies Latest reply on Sep 12, 2014 2:44 AM by maxdz8

    CLK_LOCAL_MEM_FENCE wipes values at indices of >63 in local arrays

    animal_magic

      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?