测量 CUDA 内核时时序不同答案

【问题标题】：Timings differ while measuring a CUDA kernel测量 CUDA 内核时时序不同
【发布时间】：2019-06-09 10:20:31
【问题描述】：

我尝试测量 CUDA 内核函数所花费的时间。我测量了它的 CPU 和 GPU 时间。但我发现两者之间存在巨大差异。

当我使用 NVIDIA 分析器对其进行分析时，内核大约需要 6 毫秒，这正是我想要的。但是当我在内核调用周围使用 gettimeofday() 来获取 CPU 计时时，测量结果是 15 毫秒。我那里也没有任何 memcpy 调用。内核在单独的流中运行。类似的内核也在并发流中运行。

示例代码：

gettimeofday(start);
cudaEventRecord(startGPU);

Kernel <<<abc, xyz,stream>>>();
cudaDeviceSynchronize();

cudaEventRecord(stopGPU);
printf("Elapsed GPU time  = ");

gettimeofday(stop);
printf("Elapsed CPU time  = ");

上述代码得到的结果：

经过的 GPU 时间 = 6 毫秒 CPU 运行时间 = 15 毫秒

这很奇怪，因为只有内核执行行存在。然而，内核参数是指针。内存副本是否占用了额外的时间？但我也没有在个人资料的任何地方找到内存副本。任何线索将不胜感激。

【问题讨论】：

什么平台？如果 Windows 带有 WDDM 驱动程序，CPU 时间将另外包括一个批处理超时。如果 GPU 还驱动显示器，则在 GPU 更新屏幕时可能会涉及额外的延迟。
是的，我知道这一点。但我使用的是 NVIDIA Jetson - TX2。

标签： c parallel-processing cuda gpu nvidia

【解决方案1】：

基本上，您测量的 CPU 时间是它所花费的时间

记录第一个事件，
使用相应的参数设置内核启动，
向 GPU 发送必要的命令，
在 GPU 上启动内核，
在 GPU 上执行内核，
等待 GPU 执行完成的通知返回 CPU，然后
记录第二个事件。

另外，请注意，您测量 CPU 时间的方法不仅仅测量您的进程/线程所花费的处理时间，而是所用的总系统时间（这可能包括其他进程/线程在您的进程/线程甚至不一定在运行）。我不得不承认，即使考虑到所有这些，与 GPU 时间相比，您报告的 CPU 时间仍然比我通常预期的要大得多。但我不确定那里真的是你的整个代码。事实上，我对此表示怀疑，例如，printf()s 并没有真正打印任何东西。因此，可能还有一些我们不知道的其他因素仍然需要考虑才能充分解释您的时间安排。

无论如何，您进行的两次测量很可能都没有真正测量您真正想要测量的内容。如果您对内核运行所需的时间感兴趣，请使用 CUDA 事件。但是，如果您先同步然后才记录结束事件，则开始和结束事件之间的时间将是内核执行开始、CPU 等待内核执行完成之间的时间，以及到那时可能需要的时间记录第二个事件并让那个事件到达 GPU，这样您就可以询问 GPU 是什么时候得到它的。考虑事件，例如标记发送到 GPU 的命令流中的特定点的标记。最有可能的是，您实际上想这样写：

cudaEventRecord(startGPU, stream);       // mark start of kernel execution
Kernel<<<abc, xyz, stream>>>();
cudaEventRecord(stopGPU, stream);        // mark end of kernel execution
cudaEventSynchronize(stopGPU);   // wait for results to be available

然后使用cudaEventElapsedTime() 获取两个事件之间的时间。

另外，请注意gettimeofday() 是not necessarily 获得高分辨率时序的可靠方法。在 C++ 中，您可以使用 std::steady_clock 或 std::high_resolution_clock 等（我只会在无法避免的情况下使用后者，因为它不能保证稳定；并确保时钟周期实际上足够对于您要测量的内容）。

【讨论】：

要测量内核在特定流上执行的经过时间，必须通过将流句柄作为cudaEventRecord 的第二个参数传递来在该流上记录事件。像这样：cudaEventRecord(startGPU, stream);.
@sgarizvi 是的，我不知何故忽略了他显然正在使用单独的流来启动内核。感谢您指出了这一点！固定。
感谢@MichaelKenzel 的回复。正如我所说，这只是示例代码，功能非常相似。我尝试了 std::steady_clock 并在 cudaEventRecord() 中使用了流，根据您的建议，它给了我与 +/- 0.05 ms 相同的时间，我猜这可以忽略不计。从 cuda profiler 中可以看出，内核需要 5-6 毫秒。但是为什么我的 CPU 时间显示为 15 毫秒？挠头

【解决方案2】：

在调试同一问题后，我发现 cuda 通常在第一次内核启动之前需要一些时间，如论坛中所述：https://devtalk.nvidia.com/default/topic/1042733/extremely-slow-cuda-api-calls-/?offset=3。

内核之前的 cuda 运行时 API 有 6ms 的 cudaMalloc 和 14ms 的 cudaLaunch，这是额外延迟的原因。但是，随后的内核可以正常工作。 cudaLaunch 需要的时间通常以微秒为单位，因此如果超出此范围，肯定需要进行一些修复。

注意：如果您在 while(1) 循环中运行任何 cuda 内核（仅一次），则分配必须在循环外完成。否则你最终会像这样延迟。

【讨论】：