OpenCL 代码在 MBP 上比在 NVIDIA GTX 480 上运行得更快

Question

我遇到了一个奇怪的问题。我在 OpenCL 中实现了一些线性代数，到目前为止只有矩阵乘法，并且一直在我的笔记本电脑上进行测试。代码很简单：

__kernel void matrix_mult(__global float* a, 
              __global float* b, 
              __global float* c,
              const int N) 
{
  int row = get_global_id(1);
  int col = get_global_id(0);
  float sum = 0.0f;
  for (int i = 0; i < N; i++) {
    sum += a[row*N+i] * b[i*N+col];
  }
  c[row*N+col] = sum;
}

我通过像这样运行代码 100 次来测试硬件：

  clock_t begin=clock(); 

  const unsigned int repeats = 100;
  for(int  i = 0; i != repeats; i++){
    runCL(a, b, results,N, N*N);
  }

  clock_t end=clock();

在我的 MBP 矩阵乘法上大约需要 1.2 毫秒，在大小为 512*512 的矩阵上，而在 GTX 480 Linux 机器上运行相同的代码大约需要 3 毫秒。这让我很困扰，因为我不希望昂贵的 GTX 卡比笔记本电脑快一点。

据我所知，要么我的代码“错误”，要么我以某种错误的方式计时。

我尝试在 OpenCL 规范中使用基于事件的计时系统，这给出了一些更真实的结果。

cl_event event = {0}; 
err = clEnqueueNDRangeKernel(cmd_queue, kernel[0], 2, NULL, global_work_size, NULL, 0, NULL, &event);
assert(err == CL_SUCCESS);


cl_int err =  clWaitForEvents (1,&event);
cl_ulong start, end; 
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END,   sizeof(cl_ulong), &end,   NULL); 
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL); 
double executionTimeInMilliseconds = (end - start) * 1.0e-6f;
std::cout << "execution time in milis : " << executionTimeInMilliseconds << std::endl;

现在 GT330M 将在 46 毫秒内完成操作，而 GTX480 在 2.5 毫秒内完成。这就引出了另一个非常有趣的问题，开启 PROFILING 后，GT 330M 的速度会慢 30 倍左右，这是有道理的，但 GTX480 保持了相同的性能。谁能解释这是为什么？

Answer 1

在为原始问题计时，您在这里看到的是，使用这种幼稚的代码，更好的 GTX480 规格实际上正在伤害您。

代码示例，矩阵乘法的第一遍，完全由内存带宽控制；每个线程正在访问 B 的不同元素，由于步幅而无法合并。

GTX480 的内存总线比 GT330M（128 位，800 MHz）大 3 倍（384 位）和快 2 倍（1840 MHz）内存总线。名义上，这提供了 177.4GB/s 与 25.6GB/s 的峰值带宽优势，并且由于这是内存带宽占主导地位，您可能认为这会赢。然而，由于非合并读取和更宽的内存总线，b-array 访问仅使用 384 位内存访问中的 32 位，而在 330M 的情况下，每个 128 位访问中仅使用 32 位。所以b访问的有效内存带宽分别为14.8GB/s和6.4GB/s；所以现在总内存带宽的差异只有 2 倍，而不是 7 倍左右，而且更快的卡的大部分优势都被浪费了；此外，内存带宽必须除以 10 倍的内核数，因此每个内核访问和进行计算的延迟时间更长。我怀疑如果你使用更大的矩阵大小，你可以隐藏更多的延迟并接近最好的 2 倍加速，而不是你看到的 2.5 倍减速。

这里的最终解决方案是使用内存更友好的矩阵乘法算法作为基准。

不过，我不知道您看到的分析结果。或许 330M 对性能分析没有那么好的硬件支持，所以必须用软件来实现？由于 GTX 编号几乎相同，所以我现在只使用更简单的计时方法，因为您没有使用异步内核或传输，所以应该没问题。

Answer 2

我认为您正在突破 Nvidia 计时器分辨率的限制。尝试 clGetDeviceInfo() 和 CL_DEVICE_PROFILING_TIMER_RESOLUTION 来检查它。在那些微小的时间里，我不会真正得出任何结论。

Answer 3

每个代码路径的初始化例程之间可能会有几毫秒的差异，尤其是当两个测试系统具有不同的硬件时。 我建议先在笔记本电脑和 nVidia 卡上测试一个更大的集合，这至少需要几秒钟。