Help with locks and global memory

Discussion created by tonyo_au on Jul 4, 2011
Latest reply on Jul 21, 2011 by tonyo_au
Global does not update correctly

Hi All, I am looking for some help with this code


__kernel void ReduceBoundsFromGBuffer(
  __global int   *p_Semaphors,
  __global PartitionBoundsC* p_PartionBounds
  uint2 groupId;            // SV_GroupID.xy
  uint2 groupThreadId;      // SV_GroupThreadID.xy
  uint  groupIndex ;        // SV_GroupIndex

  groupId.x = get_group_id(0);
  groupId.y = get_group_id(1);

  groupThreadId.x = get_local_id(0);
  groupThreadId.y = get_local_id(1);

  groupIndex = get_local_size(0)*groupThreadId.y+groupThreadId.x;

  // Now write out the result from this pass
  if (groupIndex == 0)
     //  p_Semaphors[0] += 1.0f;
      p_PartionBounds[1].c_MinCoord.y += 1.0;
      p_PartionBounds[2].c_MinCoord.y  += 1.0f;
      p_PartionBounds[3].c_MinCoord.y  += 1.0f;
      p_PartionBounds[0].c_MinCoord.y  += 1.0f;

If I run this on cpu's only it works correctly

p_PartionBounds[1].c_MinCoord.y has number of threads run.


If I run it on a GPU (5870)

p_PartionBounds[1].c_MinCoord.y has 1,

all other y values are correct


if I uncomment the line 

//  p_Semaphors[0] += 1.0f;

p_PartionBounds[1].c_MinCoord.y has the correct value


Am I doing something obviouly wrong?


PS the semaphore code seems to be ok