OpenCL障碍:如何为所有内核设置障碍

Question

我有一个OpenCL内核，它在全局配置[1024x1024]中运行，不与工作组一起工作（它是一个简单的颜色分割内核）。

我在那里放置了一个barrier(CLK_GLOBAL_MEM_FENCE)同步，以便等待[1024x1024]配置中的所有线程都到达它。但遗憾的是，它似乎不起作用。

你不能告诉我，我做错了什么吗？

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 %zun", n_pics, n_chans, pic_h, pic_w);
  for(int n = 0; n < 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 < 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("%fn", 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 < 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);
  }
}

Answer 1

CLK_GLOBAL_MEM_FENCE:

• 同步工作组内的全局读/写

CLK_LOCAL_MEM_FENCE:

• 同步工作组内的本地读/写

您真正想要的是在内核中同步全局读/写。这在OpenCL中是不可能的（在任何GPU编程语言中也是如此）。如果您的内核可以放在一个工作组中，那么它也可以，但我真的怀疑1M WI是否可以放在单个组中。

如果需要这种行为，最合适和可移植的方法是将内核拆分为两个内核