为什么这个简单的 OpenCL 内核运行如此缓慢？答案

【问题标题】：Why is this simple OpenCL kernel running so slowly?为什么这个简单的 OpenCL 内核运行如此缓慢？
【发布时间】：2012-10-02 13:05:34
【问题描述】：

我正在研究 OpenCL，与我预期的运行方式相比，我有点困惑为什么这个内核运行得这么慢。这是内核：

__kernel void copy(
  const __global char* pSrc, 
  __global __write_only char* pDst, 
  int length)
{
  const int tid = get_global_id(0);

  if(tid < length) {
    pDst[tid] = pSrc[tid];
  }
}

我通过以下方式创建了缓冲区：

char* out = new char[2048*2048];
cl::Buffer(
  context,
  CL_MEM_USE_HOST_PTR | CL_MEM_WRITE_ONLY,
  length,
  out);

输入缓冲区同上，除了我已将 in 指针初始化为随机值。最后，我这样运行内核：

cl::Event event;
queue.enqueueNDRangeKernel(
  kernel, 
  cl::NullRange,
  cl::NDRange(length),
  cl::NDRange(1), 
  NULL, 
  &event);

event.wait();

平均而言，时间约为 75 毫秒，计算公式为：

cl_ulong startTime = event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
cl_ulong endTime = event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
std::cout << (endTime - startTime) * SECONDS_PER_NANO / SECONDS_PER_MILLI << "\n";

我运行的是带有 Intel i5-3450 芯片（Sandy Bridge 架构）的 Windows 7。相比之下，进行复制的“直接”方式只需不到 5 毫秒。我不认为 event.getProfilingInfo 包括主机和设备之间的通信时间。想法？

编辑：

在 ananthonline 的建议下，我将内核更改为使用 float4s 而不是 chars，这将平均运行时间降低到大约 50 毫秒。仍然没有我希望的那么快，但有所改善。谢谢 ananthonline！

【问题讨论】：

这是什么实现？英特尔 OpenCL 实施？您是否尝试过使用 4 浮点数组的相同内核？这可能是一种更好的内存访问模式。
是的，英特尔 OpenCL 实施。我还没有尝试过 4-floats，这是一个好主意。我会调查的。
出于好奇，clEnqueueCopyBuffer 的表现如何？这就是你所说的“直接”吗？
@willglynn 好多了。分析信息将时间报告为大约 2 毫秒。在 QueryPerformanceCounters 中包装对 event.wait() 的调用会报告时间稍高一些，大约 8 毫秒，但当然这包括等待调用本身的开销。
@willglynn 但我所说的“直接”是一个 memcopy 调用。我知道，这有点像苹果到橘子，但它至少提供了某种基线。

标签： performance opencl gpgpu

【解决方案1】：

我认为您的主要问题是您使用的 2048*2048 工作组。如果您有这么多单项工作组，系统上的 opencl 驱动程序必须管理更多开销。如果您使用 gpu 执行此程序，这将特别糟糕，因为您会得到非常低的硬件饱和度。

优化：用更大的工作组调用你的内核。您甚至不必更改现有内核。 see question: What should this size be?下面我以64为例。在大多数硬件上，64 恰好是一个不错的数字。

cl::size_t myOptimalGroupSize = 64;
cl::Event event;
queue.enqueueNDRangeKernel(
  kernel, 
  cl::NullRange,
  cl::NDRange(length),
  cl::NDRange(myOptimalGroupSize), 
  NULL, 
  &event);

event.wait();

您还应该让您的内核做更多的事情，而不是复制单个值。我已经回答了一个关于全局内存over here.

的类似问题

【讨论】：

【解决方案2】：

CPU 与 GPU 非常不同。在 x86 CPU 上运行它，获得良好性能的最佳方法是使用 double16（最大的数据类型）而不是 char 或 float4（如其他人所建议的那样）。

根据我在 CPU 上使用 OpenCL 的少量经验，我从未达到使用 OpenMP 并行化所能达到的性能水平。与CPU并行进行复制的最佳方法是将要复制的块划分为少量的大子块，并让每个线程复制一个子块。 GPU 方法是正交的：每个线程都参与同一块的副本。这是因为在 GPU 上，不同的线程可以有效地访问连续的内存区域（合并）。

要使用 OpenCL 在 CPU 上进行高效复制，请在内核中使用循环来复制连续数据。然后使用不大于可用内核数的工作组大小。

【讨论】：

【解决方案3】：

我相信是 cl::NDRange(1) 告诉运行时使用单项工作组。这效率不高。在 C API 中，您可以为此传递 NULL 以将工作组大小留给运行时；在 C++ API 中也应该有一种方法可以做到这一点（也许也只是 NULL）。这在 CPU 上应该更快；它肯定会在 GPU 上。

【讨论】：