title: OpenCL barriers: how to make a barrier for all kernels?
body: <p>I have an OpenCL kernel, that runs in global configuration [1024x1024] and does not work with workgroups (it is a simple color-splitting kernel).</p>
  
  <p>I placed there a <code>barrier(CLK_GLOBAL_MEM_FENCE)</code> synchronization, in order to wait while all the threads from [1024x1024] configuration reach it. But alas, it seems like it does not work.</p>
  
  <p>Cannot you tell me, what I'm doing wrong?</p>
  
  <pre><code>typedef float dtype;
  
  // Splits color channels in-place
  __kernel void nhwc_to_nchw(__global dtype* pic_data, __global dtype* buffer,
                             const int n_pics, const int n_chans)
  {
    size_t pic_h = get_global_size(0);
    size_t pic_w = get_global_size(1);
  
    size_t pic_y = get_global_id(0);
    size_t pic_x = get_global_id(1);
  
    size_t nhwc_index;
    size_t nchw_index;
    size_t pic_index;
  
  //  printf("%i %i %zu %zu\n", n_pics, n_chans, pic_h, pic_w);
  
    for(int n = 0; n &lt; n_pics; n++) {
      // Writting the splitted channels from one pic to buffer
      pic_index = n * n_chans * pic_w * pic_h;
      for(int chan = 0; chan &lt; n_chans; chan++) {
        nhwc_index = pic_y * pic_w * n_chans + pic_x * n_chans + chan;
        nchw_index = chan * pic_h * pic_w + pic_y * pic_w + pic_x;
  //      printf("%f\n", pic_data[pic_index + nhwc_index]);
        buffer[nchw_index] = pic_data[pic_index + nhwc_index];
      }
  
      barrier(CLK_GLOBAL_MEM_FENCE);
  
      //Reading splitted channels from buffer and placing them to source pic
      for(int chan = 0; chan &lt; n_chans; chan++) {
        nchw_index = chan * pic_h * pic_w + pic_y * pic_w + pic_x;
        pic_data[pic_index + nchw_index] = buffer[nchw_index];
      }
  
      barrier(CLK_GLOBAL_MEM_FENCE);
    }
  
  }
  </code></pre>