cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

tonyo_au
Journeyman III

Help with locks and global memory

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)
  {
       GetSemaphor(&p_Semaphors[2]);
   {
     //  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;
    }
     ReleaseSemaphor(&p_Semaphors[2]);
  }
}

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

0 Likes
10 Replies