Multiple read-write synchronization issues in opencl local and global memories

Question

I have an opencl kernel that finds the maximum ASCII character in a string. The problem is I cannot synchronize the multiple read-writes to global and local memories. I am trying to update a local_maximum character in shared memory, and at the end of the workgroup (last thread), the global_maximum character, by comparing it with the local_maximum. The threads are writing one over another, I guess.

eg: Input string: "pirates of the carribean".

Output String: 'r' (but it should be 's').

Please have a look at the code and give a solution as to what I can do to get everything synchronized. I am sure people having sound knowledge can understand the code. Optimization tips are welcome.

The code is below:

__kernel void find_highest_ascii( __global const char* data, __global char* result, unsigned int size,  __local char* localMaxC )
{
//creating variables and initialising..
unsigned int i, localSize, globalSize, j;
char privateMaxC,temp,temp1;

i = get_global_id(0);
localSize = get_local_size(0);
globalSize = get_global_size(0);

privateMaxC = '\0';

if(i<size){
if(i == 0)
read_mem_fence( CLK_LOCAL_MEM_FENCE );
*localMaxC = '\0';
mem_fence( CLK_LOCAL_MEM_FENCE);

////////////////////////////////////////////////////
/////UPDATING PRIVATE MAX CHARACTER/////////////////
////////////////////////////////////////////////////

for( j = i; j<size; j+=globalSize )
{
    if( data[j] > privateMaxC )
    {
        privateMaxC = data[j];
    }
}

///////////////////////////////////////////////////


///////////////////////////////////////////////////
////UPDATING SHARED MAX CHARACTER//////////////////
///////////////////////////////////////////////////

temp = *localMaxC;
read_mem_fence( CLK_LOCAL_MEM_FENCE );

if(privateMaxC>temp)
{
    *localMaxC = privateMaxC;
    write_mem_fence( CLK_LOCAL_MEM_FENCE );
    temp = privateMaxC;
}

//////////////////////////////////////////////////


//UPDATING GLOBAL MAX CHARACTER.

temp1 = *result;

if(( (i+1)%localSize == 0 || i==size-1) && (temp > temp1 ))
{
            read_mem_fence( CLK_GLOBAL_MEM_FENCE );
    *result = temp;
    write_mem_fence( CLK_GLOBAL_MEM_FENCE );
}


 }
}

Answer 1

You are correct that threads will be overwriting each other's values, since your code is riddled with race conditions. In OpenCL, there is no way to synchronise between work-items that are in different work-groups. Instead of trying to achieve this kind of synchronisation with explicit fences, you can make your code much simpler by using the built-in atomic functions instead. In particular, there is an atomic_max built-in which solves your problem perfectly.

So, instead of the code you currently have to update both your local and global memory maximum values, just do something like this:

kernel void ascii_max(global int *input, global int *output, int size,
                      local int *localMax)
{
  int i = get_global_id(0);
  int l = get_local_id(0);

  // Private reduction                                                          
  int privateMax = '\0';
  for (int idx = i; idx < size; idx+=get_global_size(0))
  {
    privateMax = max(privateMax, input[idx]);
  }

  // Local reduction                                                            
  atomic_max(localMax, privateMax);
  barrier(CLK_LOCAL_MEM_FENCE);

  // Global reduction                                                           
  if (l == 0)
  {
    atomic_max(output, *localMax);
  }
}

This will require you to update your local memory scratch space and final result to use 32-bit integer values, but on the whole is a significantly cleaner approach to solving this problem (not to mention it actually works).

---

NON-ATOMIC SOLUTION

If you really don't want to use atomics, then you can implement a bog-standard reduction using local memory and work-group barriers. Here's an example:

kernel void ascii_max(global int *input, global int *output, int size,
                      local int *localMax)
{
  int i = get_global_id(0);
  int l = get_local_id(0);

  // Private reduction                                                          
  int privateMax = '\0';
  for (int idx = i; idx < size; idx+=get_global_size(0))
  {
    privateMax = max(privateMax, input[idx]);
  }

  // Local reduction                                                            
  localMax[l] = privateMax;
  for (int offset = get_local_size(0)/2; offset > 1; offset>>=1)
  {
    barrier(CLK_LOCAL_MEM_FENCE);
    if (l < offset)
    {
      localMax[l] = max(localMax[l], localMax[l+offset]);
    }
  }

  // Store work-group result in global memory                                   
  if (l == 0)
  {
    output[get_group_id(0)] = max(localMax[0], localMax[1]);
  }
}

This compares pairs of elements at a time using local memory as a scratch space. Each work-group will produce a single result, which is stored in global memory. If your data-set is small, you could run this with a single work-group (i.e. make global and local sizes the same), and this will work just fine. If it is larger, you could run a two-stage reduction by running this kernel twice, e.g.:

size_t N = ...; // something big

size_t local  = 128;
size_t global = local*local; // Must result in at most 'local' number of work-groups

// First pass - run many work-groups using temporary buffer as output
clSetKernelArg(kernel, 1, sizeof(cl_mem), d_temp);
clEnqueueNDRangeKernel(..., &global, &local, ...);

// Second pass - run one work-group with temporary buffer as input
global = local;
clSetKernelArg(kernel, 0, sizeof(cl_mem), d_temp);
clSetKernelArg(kernel, 1, sizeof(cl_mem), d_output);
clEnqueueNDRangeKernel(..., &global, &local, ...);

I'll leave it to you to run them and decide which approach would be best for your own data-set.