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

Question

我正在研究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";

我运行的是Windows7，采用英特尔i5-3450芯片(SandyBridge架构)。相比之下，"直接"复制方式只需不到5毫秒。我不认为事件.getProfileInfo包括主机和设备之间的通信时间。想法？

编辑：

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

Answer 1

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

优化：使用较大的工作组调用内核。您甚至不必更改现有的内核。看问题：这个尺寸应该是多少？我用下面的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();

您还应该让您的内核做的不仅仅是复制单个值。我在这里回答了一个关于全球记忆的类似问题。

Answer 2

CPU与GPU有很大不同。在x86 CPU上运行此程序，获得良好性能的最佳方法是使用double16(最大的数据类型)，而不是char或float4(如其他人所建议的)。

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

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

Answer 3

我相信是cl:：NDRange(1)告诉运行时使用单项工作组。这是无效的。在C API中，您可以为此传递NULL，以便将工作组大小保留到运行时；在C++API中也应该有这样的方法(也许也只是NULL)。这在CPU上应该更快；它肯定会在GPU上。