opencl速度慢的问题答案

【问题标题】：Troubles with slow speeds in openclopencl速度慢的问题
【发布时间】：2019-10-02 07:10:56
【问题描述】：

我第一次尝试使用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];

}

【问题讨论】：

“我的表现似乎变差了。”为什么不添加一些关于执行时间的信息（也许你的硬件）？并且“常规for循环”意味着cpu上的“相同”实现只是使用外部for循环而不是gpu上的线程？关于性能，我会说访问模式对 gpu 来说真的很糟糕。以列为主的顺序会更好。参见例如this question.
如果你的 opencl 设备是一个独立的 GPU，那么我并不奇怪你的性能会变差。您的内核是微不足道的，总运行时间主要是传输数据的开销。 CPU 和 iGPU 不需要传输任何数据，但独立 GPU 需要，而且通过 PCIe 的传输速度相对较慢。除非它是更大计算的一部分，否则在离散 GPU 上运行此代码确实没有意义。

标签： c gpu opencl

【解决方案1】：

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

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

改进建议：

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

【讨论】：