3

我遇到了一个奇怪的问题。我在 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 保持了相同的性能。谁能解释这是为什么？

4

3 回答 3

4

在计时原始问题时，您在这里看到的是，使用这个幼稚的代码，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 编号几乎相同，所以我现在只使用更简单的计时方法，因为您没有使用异步内核或传输，所以应该没问题。

于 2011-05-25T18:29:02.237 回答
2

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

于 2011-05-25T18:43:03.773 回答
1

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

于 2011-05-25T14:16:30.097 回答