opencl中速度慢的问题

Question

我第一次尝试使用opencl，目标是计算数组中每一行的argmin。由于每一行的操作都独立于其他行，因此我认为将其轻松放在图形卡上是很容易的。

与使用外部forloop在cpu上运行代码时相比，使用此代码获得的性能似乎更差，我们将不胜感激。

这是代码：

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

int argmin(global double *array, int end)
{
  double minimum = array[0];
  int index;
  for (int j = 0; j < end; j++)
  {
    if (array[j] < minimum)
    {
      minimum = array[j];
      index = j;
    }
  }
  return index;
}

kernel void execute(global double *dist, global long *res, global double *min_dist)
{
  int row_size = 0;
  int i = get_global_id(0);

  int row_index = i * row_size;
  res[i] = argmin(&dist[row_index], row_size);
  min_dist[i] = dist[res[i] + row_index];

}

Answer 1

评论者提出了一些有效的观点，但我会尝试更有建设性和更有组织性：

1. double您的数据似乎由精度浮点值组成。根据您的 GPU，这本身可能是个坏消息。消费级 GPU 通常未针对处理doubles 进行优化，与单精度运算相比，通常仅实现 1/32 或 1/16 的吞吐量float。不过，许多专业级 GPU（Quadro、Tesla、FirePro、一些 Radeon Pro 卡）都可以很好地使用它们，实现 1/2 或 1/4 的吞吐量，而float. 由于您只执行简单的算术运算（比较），并且您的运行时很可能由内存访问主导，因此它在消费类硬件上也可能没问题。
2. 我假设你的值row_size实际上不是 0，这将有助于了解真正的（典型）值是什么，以及它是固定的、按行变量还是每次运行变量但每行都相同。无论如何，除非row_size非常小，否则您在其上运行串行循环的事实for可能会阻碍您的代码。
3. 您的工作规模有多大？换句话说，数组中有多少行（如果有变化，请给出典型范围）？如果它非常小，您将看不到 GPU 并行带来的好处：GPU 具有大量处理器，并且可以为每个处理器调度几个线程。因此，您的工作项目需要数百甚至数千个才能实现良好的硬件利用率。
4. 您正在从（大概）系统内存中读取一个非常大的数组，并且没有对其执行任何密集操作。这意味着您的瓶颈通常出现在内存访问方面 - 对于独立 GPU，系统内存访问需要通过 PCIe，因此该链接的速度将为您的性能设置上限。此外，您的内存访问模式对于 GPU 来说远非理想 - 您通常希望工作项同时读取相邻的内存单元，因为内存单元通常一次获取 64 字节或更多字节。

改进建议：

• 分析。如果可能，请使用 GPU 供应商的分析工具来确定真正的瓶颈。否则我们只是猜测。
• 对于 (4) - 如果可能的话，尽量不要过多地移动大量数据。如果您可以在 GPU 上生成输入数组，请这样做，这样它们就永远不会离开 VRAM。
• 对于 (4) - 优化您的内存访问。AMD、NVidia 和 Intel 都有 OpenCL GPU 优化指南，解释了如何做到这一点。本质上，重新构建数据布局或内核，以便相邻的工作项读取相邻的内存片段。理想情况下，您希望工作项 0 读取数组项 0，工作项 1 读取数组项 1，等等。您可能需要使用本地内存在工作项之间进行协调。另一种选择是读取每个工作项的向量大小的数据块。（例如，每个工作项一次读取一个 double8）但在这种情况下请注意对齐。
• 对于 (2) 和 (3) - 除非row_size非常小（且固定），否则尝试将循环拆分为多个工作项，并使用本地内存（缩减算法）和全局内存中的原子操作进行协调。
• 对于 (1)：如果您已经优化了其他所有内容，并且分析告诉您double在消费类硬件上比较 s 太慢，请检查是否可以将数据生成为floats 而不会损失准确性（这也将使您的内存带宽问题减半） ），或者检查是否可以以其他方式做得更好，例如将 the 视为doublealong并使用整数运算手动解包和比较指数和尾数。