c - 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 上，matrix_multiplications 在大小为 512*512 的矩阵上大约需要 1.2 毫秒，而在 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 数组访问仅使用 384 位内存访问中的 32 位，而在 330M 的情况下，每个 128 位访问中仅使用 32 位。所以b访问的有效内存带宽是14.8GB/s和6.4GB/s；所以现在总内存带宽只有 2 倍的差异，而不是 7 倍左右，而且速度更快的卡的大部分优势都被浪费了；此外，内存带宽必须除以 10 倍的内核数，因此每个内核访问和进行计算的延迟时间更长。我怀疑如果您使用更大的矩阵大小，您可以隐藏更多延迟并更接近最佳可能的 2 倍加速，而不是您看到的 2.5 倍减速。

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

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