如何测量 OpenCL 读出时间？答案

【问题标题】：How to measure OpenCL readout time?如何测量 OpenCL 读出时间？
【发布时间】：2014-06-09 20:44:35
【问题描述】：

我想确保clEnqueueReadBuffer() 的测量时间正确，如何测量将数据从 GPU 复制到主内存所需的时间。我有点怀疑clFinish() 也有贡献。

我运行 1000000 个浮点数的 100 倍点积并以此测量时间

  // RUN TIME BLOCK
  println( " Running OpenCL program ... " );
  t1 = System.nanoTime(); 
  for (int reps = 0; reps < 100; reps++) {
    clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, global_work_size, local_work_size, 0, null, null);    
  }
  clFinish(commandQueue); // VERSION 1
  t2 = System.nanoTime(); 
  println( " OpenCL Run Time : "+ ((t2-t1)*1e-9)+" [s] " );

  // READ OUT TIME BLOCK
  t1 = System.nanoTime(); 
  //clFinish(commandQueue); // VERSION 2
  clEnqueueReadBuffer(commandQueue, memObjects[2], CL_TRUE, 0, n * Sizeof.cl_float, dst, 0, null, null);
  t2 = System.nanoTime(); 
  println( " Read-out Time: "+ ((t2-t1)*1e-9)+" [s] " );

得到结果

 OpenCL Run Time : 2.5124469 [s] 
 Read-out Time: 0.002145424 [s]

这对我来说似乎有点太好了，只有 2 毫秒和

当我将 clFinish(commandQueue); 放入 Read-out Time 块中时，我得到了这些结果

OpenCL Run Time : 1.0892084 [s] 
Read-out Time: 1.4300439 [s]

另一方面，这似乎太糟糕了……在 GPU 上进行 100 次乘法比通过 PCI-express 复制它更快？ ....也许

只是为了完整性：

我为processing 使用了openclp5 库，它在带有Quadro FX 580 GPU 的Ubuntu 12.04 64 位和Java jdk 1.7 上使用jocl 和

我的内核很简单（没有优化或任何东西）

  String programSource =     
   "__kernel void sampleKernel(                                                 "+
   "__global const float *a, __global const float *b, __global float *c) {      "+
   "  int gid = get_global_id(0);                                               "+
   "  c[gid] = a[gid] * b[gid];                                                 "+
   "}";

【问题讨论】：

任何“clEnqueue*”命令都是这样做的：enqueueing。您只是在测量向队列添加内容的开销。如果您想知道复制所需的时间，您需要添加 clFinish。或者使用 OpenCL 事件和命令队列分析。
嗨，我很清楚，为了测量 OpenCL 程序的运行时间，有必要计算 clFinish()，但我不清楚，如果测量 clEnqueueReadBuffer() 确实如此包括将数据从 GPU 复制到 RAM 所需的所有时间，或者如果此操作的一部分已在 clFinish() 中完成
clEnqueueReadBuffer 排队工作。工作在不久之后异步开始。 clFinish 本身并不做工作，它只是等待所有工作完成。如果工作没有完成，它会阻塞并且在它完成之前不会返回。如果工作已经完成（可能是因为你开始计算 pi 到一百万位），那么clFinish 就会返回。因此，如果您从调用 clEnqueueReadBuffer 之前到调用 clFinish 之后进行计时，那么您就是在计算复制所需的时间，以及一些较小的 API 开销。
aha 好吧，我第一次在clEnqueue* 中没有得到那个*。现在我知道我必须在clEnqueueReadBuffer 之后再打电话给clFinish()
您不需要在clEnqueueReadBuffer 之后调用clFinish，因为您设置了阻止标志。这是一个隐含的clFinish。

标签： java performance profiling opencl

【解决方案1】：

您测量将数据读回主机的时间的第一种方法是正确的。第二种方法将包括一些计算时间。

您正在读取的数据量是1000000 * sizeof(float) = 4MB。如果这需要 2 毫秒，那么这意味着您正在实现4MB/0.002s = 2 GB/s 的带宽。为什么你认为这好得令人难以置信？您的卡支持 PCIe x16，其理论峰值带宽为 8 GB/s（单向）。

【讨论】：

好的，谢谢，有道理。我只是太认真地对待像GPU-RAM overhead is so big that it pays-off only if you do a lot of operations with each number 这样的声明，据我了解，每个浮点数复制到/从 GPU 至少有 10-50 个 add/mul ops。如果我的处理器可以执行 3 Gflop/s，那么如果我每个浮点数执行 6 次操作（带宽为 2GB/s = 0.5 Gfloats/s），那么将其传递给 GPU 将会得到回报。
测量不是 100% 正确的。操作可能需要更少的设备时间，其余的是驱动程序开销，因此传输可能要快得多。但作为一个经验法则，这些措施是正确的。一般来说，您需要在 GPU 上投入大量工作才能使麻烦得到回报。不仅足以“隐藏”传输瓶颈。否则，CPU 中的驱动程序使用和开销将使其不值得。顺便说一句，您正在运行 100 次操作与 1 次读取，通常 1write+1run+1read 将是 4ms IO 和 25ms 操作，这是很多 IO %。