本地内存的opencl大小对速度有影响吗？答案

【问题标题】：Opencl size of local memory has impact on speed?本地内存的opencl大小对速度有影响吗？
【发布时间】：2014-05-09 20:37:24
【问题描述】：

我是 OpenCL 的新手，我正在尝试计算灰度图像的直方图。我正在 GPU nvidia GT 330M 上执行此计算。

代码是

__kernel void histogram(__global struct gray * input, __global int * global_hist, __local volatile int * histogram){
    int local_offset = get_local_id(0) * 256;
    int histogram_global_offset = get_global_id(0) * 256;
    int offset = get_global_id(0) * 1920;
    int value;
    for(unsigned int i = 0; i < 256; i++){
        histogram[local_offset + i] = 0;
    }
    barrier(CLK_LOCAL_MEM_FENCE);

    for(unsigned int i = 0; i < 1920; i++){
        value = input[offset + i].i;
        histogram[local_offset + value]++;
    }

    barrier(CLK_LOCAL_MEM_FENCE);
    for(unsigned int i = 0; i < 256; i++){
        global_hist[histogram_global_offset + i] = histogram[local_offset + i];
    }
}

此计算是在图像 1920*1080 上执行的。

我正在用

发射内核

queue.enqueueNDRangeKernel(kernel_histogram, cl::NullRange, cl::NDRange(1080), cl::NDRange(1));

当直方图的本地大小设置为 256 * sizeof(cl_int) 时，此计算的速度为（通过 nvidia nsight 性能分析）11 675 微秒。

因为本地工作组大小设置为一。我尝试将本地工作组大小增加到 8。但是当我将直方图的本地大小增加到 256 * 8 * sizeof(cl_int) 并使用本地 wg 大小 1 进行计算时。我得到 85 177 微秒。

因此，当我使用每个工作组 8 个内核启动它时，我不会从 11 毫秒获得加速，而是从 85 毫秒获得加速。因此，每个工作组 8 个内核的最终速度为 13 714 微秒。

但是当我创建计算错误时，将 local_offset 设置为零，并且局部直方图的大小为 256 * sizeof(cl_int)，并且每个工作组使用 8 个内核，我得到了更好的时间 - 3 854 微秒。

有没有人有一些想法来加快这个计算？

谢谢！

【问题讨论】：

您发布的示例代码似乎正在为每个工作项处理一行图像。当您启动每个工作组的多个工作项时，您是否仍在这样做，或者您是否让每个工作组合作处理一行？您能向我们展示您在执行此操作时使用的内核和主机代码吗？
感谢您的回复。是的，我使用的是相同的方法，一个线程计算一行。
这会导致内存访问性能不佳。理想情况下，您希望合并您的内存访问 - 也就是说，相邻的工作项应该尽可能访问相邻的内存地址。
当我在每个工作组的本地内存中使用一个直方图（8 个本地大小）并使用 atom_inc(&histogram[value]);这个计算的速度是 5.5 毫秒，所以它更好，但是，我仍然不明白为什么随着本地内存大小的增加，我的时间会变慢......
好的，谢谢，我试试这个

标签： opencl histogram

【解决方案1】：

此答案假定您最终希望将直方图一直减少到 256 个 int 值。您调用内核的工作组与设备上的计算单元一样多，并且组大小应该（一如既往）是设备上 CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE 的倍数。

__kernel void histogram(__global struct gray * input, __global int * global_hist){
  int group_id = get_group_id(0);
  int num_groups = get_num_groups(0);
  int local_id = get_local_id(0);
  int local_size = get_local_size(0);

  volatile __local int histogram[256];

  int i;
  for(i=local_id; i<256; i+=local_size){
    histogram[i] = 0;
  }

  int rowNum, colNum, value, global_hist_offset

  for(rowNum = group_id; rowNum < 1080; rowNum+=num_groups){
    for(colNum = local_id; colNum < 1920; colNum += local_size){
      value = input[rowNum*1920 + colNum].i;
      atomic_inc(histogram[input]);
    }
  }

  barrier(CLK_LOCAL_MEM_FENCE);
  global_hist_offset = group_id * 256;
  for(i=local_id; i<256; i+=local_size){
    global_hist[global_hist_offset + i] = histogram[i];
  }

}

每个工作组一次合作处理一排图像。然后该组移动到另一行，使用 num_groups 值计算。无论您有多少组，这都会很好地工作。例如，如果您有 7 个计算单元，则第 3 组（第四组）将从映像中的第 3 行开始，然后每隔 7 行开始。第 3 组将总共计算 153 行，其最后一行将是第 1074 行。一些工作组可能会再计算 1 行——在本例中为组 0 和 1。

在查看图像的列时，在工作组内执行相同的隔行扫描。在 colNum 循环中，第 N 个工作项从第 N 列开始，并向前跳过 local_size 列。此循环的其余部分不应经常出现，因为 CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE 可能是 1920 的一个因子。尝试从 (1..X) * CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE 到设备的最大工作组大小的所有工作组大小。

关于此内核的最后一点：结果与您的原始内核不同。您的 global_hist 数组是 1080 * 256 个整数。我需要的是 num_groups * 256 个整数。如果您想要完全减少，这会有所帮助，因为内核执行后要添加的内容要少得多。

【讨论】：