为什么这个OpenCL算法中的本地内存如此之慢？

Question

我正在写一些OpenCL代码.我的内核应该根据输入图像创建一个特殊的"累加器"输出.我尝试了两个概念,两者同样很慢,尽管第二个使用本地内存.你能帮我解一下为什么本地内存版本这么慢吗？内核的目标GPU是AMD Radeon Pro 450.

// version one
__kernel void find_points(__global const unsigned char* input, __global unsigned int* output) {
  const unsigned int x = get_global_id(0);
  const unsigned int y = get_global_id(1);

  int ind;

  for(k = SOME_BEGINNING; k <= SOME_END; k++) {
    // some pretty wild calculation
    // ind is not linear and accesses different areas of the output
    ind = ...
    if(input[y * WIDTH + x] == 255) {
      atomic_inc(&output[ind]);
    }
  }

}

// variant two
__kernel void find_points(__global const unsigned char* input, __global unsigned int* output) {
  const unsigned int x = get_global_id(0);
  const unsigned int y = get_global_id(1);

  __local int buf[7072];
  if(y < 221 && x < 32) {
    buf[y * 32 + x] = 0;
  }
  barrier(CLK_LOCAL_MEM_FENCE);

  int ind;
  int k;

  for(k = SOME_BEGINNING; k <= SOME_END; k++) {
    // some pretty wild calculation
    // ind is not linear and access different areas of the output
    ind = ...
    if(input[y * WIDTH + x] == 255) {
      atomic_inc(&buf[ind]);
    }
  }

  barrier(CLK_LOCAL_MEM_FENCE);
  if(get_local_id(0) == get_local_size(0) - 1)
    for(k = 0; k < 7072; k++)
      output[k] = buf[k];
  }

}

我希望第二种变体比第一种更快,但事实并非如此.有时甚至更慢.

Answer 1

本地缓冲区大小__local int buf[7072](28288字节)太大.我不知道AMD Radeon Pro 450的共享内存有多大,但每台计算单元可能只有32kB或64kB.

32768/28288 = 1,65536/28288 = 2意味着只有1个或最多2个波前(64个工作项)只能同时运行,因此计算单元的占用率非常低,因此性能较差.

您的目标应该是尽可能减少本地缓冲区,以便可以同时处理更多的波前.

使用CodeXL分析您的内核-有工具向您展示这一切.或者,CUDA occupancy calculator如果您不想运行探查器以更好地了解它是什么,您可以查看Excel电子表格.