memory - 如何设置OpenCL的本地工作空间大小？

Question

我正在使用 OpenCL 进行一些图像处理。

例如，我使用了100*200大小的图像。在 .cl 代码中，我将图像像素值减半:

{
  int width=get_group_id(0);
  int height=get_group_id(1);

  // col(width)
  int x= get_global_id(0);
  // row(height)
  int y= get_global_id(1);
  (unsigned char) data_output[x*width+y]= 
     (unsigned char)data_input[x*width+y]/2;
}

内核参数设置后，我通过以下方式运行内核:

clEnqueueNDRangeKernel( queue,kernel_DIP,2,NULL,global_work_size,local_work_size, 0,NULL,NULL);

我使用的global_work_size是图像大小:

   size_t global_work_size[2] = {100,200};

我发现甚至 .cl 代码也不包含诸如“get_local_id(0);”之类的代码

local_work_size对性能也有很大影响。

两者“size_t local_work_size[2]= {1,1};”(小型本地工作尺寸)和“size_t local_work_size[2]= {50,50};” (工作量大)速度慢。

像下面这样的一些合适的尺寸会快得多:

size_t local_work_size[2]= {10,10};

这是我的问题:

1. 为什么没有 get_local_id() 的代码也会受到本地内存的影响？

2. 如何设置最佳的本地大小，使其以最高的速度运行？

3. 我还在其他平台上测试了运行速度，例如freescale的IMX.6，似乎改变大小的本地work-size在那里根本不起作用!那么为什么呢？

如果有人知道答案，请帮忙。太感谢了!

Answer 1

DarkZeros 已经提到，您可以将本地工作大小设置为 null，以便 OpenCL 在给定全局工作大小及其执行设备的情况下选择它认为“合适”的大小。

但是，对于某些全局工作大小，OpenCL 可能无法选择“合适的”本地工作大小。特别是当全局工作大小是大于最大本地工作大小的质数时。那么它可能会被迫使用本地工作大小 1。您可以考虑填充您的输入数据，以便它可以很好地分布在多个工作组之间。(我最近在 https://stackoverflow.com/a/22969485 中写了几句话)

对于复杂的内核，您可以考虑查询CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE并以此为基础进行计算，但对于这个简单的内核，这应该不是必需的。

此外，您可能想看看“AMD APP KernelAnalyzer”或“NVIDIA Occupancy Calculator” - 这些工具可能会给出有关某些目标平台的适当配置的一些提示(尽管最好应该编写代码尽可能通用，只要它不会对性能产生太严重的影响)