title: Multiple read-write synchronization issues in opencl local and global memories
body: <p>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.</p>
  
  <p>eg: Input string:  "pirates of the carribean".</p>
  
  <p>Output String: 'r' (but it should be 's').</p>
  
  <p>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.</p>
  
  <p>The code is below:</p>
  
  <pre><code>__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&lt;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&lt;size; j+=globalSize )
  {
      if( data[j] &gt; privateMaxC )
      {
          privateMaxC = data[j];
      }
  }
  
  ///////////////////////////////////////////////////
  
  
  ///////////////////////////////////////////////////
  ////UPDATING SHARED MAX CHARACTER//////////////////
  ///////////////////////////////////////////////////
  
  temp = *localMaxC;
  read_mem_fence( CLK_LOCAL_MEM_FENCE );
  
  if(privateMaxC&gt;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) &amp;&amp; (temp &gt; temp1 ))
  {
              read_mem_fence( CLK_GLOBAL_MEM_FENCE );
      *result = temp;
      write_mem_fence( CLK_GLOBAL_MEM_FENCE );
  }
  
  
   }
  }
  </code></pre>