Cuda profiler 说我的两个内核很贵，但是它们的执行时间似乎很短

Posted 2023-04-15

【中文标题】Cuda profiler 说我的两个内核很贵，但是它们的执行时间似乎很短

【英文标题】：Cuda profiler says that my two kernels are expensive, however their execution time seems to be small

【发布时间】：2013-05-10 00:18:07

【问题描述】：

我使用两个内核，我们称它们为 A 和 B。

我运行了 CUDA 分析器，这是它返回的内容：

第一个内核有 44% 的开销，而第二个内核有 20%。

但是，如果我决定按照这个逻辑找出实际的执行时间：

timeval tim;
gettimeofday(&tim, NULL);
double before = tim.tv_sec+(tim.tv_usec/1000000.0);

runKernel<<<...>>>(...)

gettimeofday(&tim, NULL);
double after=tim.tv_sec+(tim.tv_usec/1000000.0);
totalTime = totalTime + after - before;

totalTime 会非常小，大约为 0.0001 秒。

我是 CUDA 的新手，我不明白到底发生了什么。我应该尝试使内核更高效还是它们已经高效？

【参考方案1】：

从 CPU 的角度来看，内核调用是异步的（请参阅this answer）。如果您按照您在没有任何同步的情况下（即不调用 cudaDeviceSynchronize()）对内核进行计时，那么您的计时将没有任何意义，因为 GPU 上仍在进行计算。

在为内核计时 (nvprof / nvvp) 时，您可以信任 NVIDIA 的分析器。 NVIDIA Visual Profiler 还可以分析您的程序并就您的内核可能出现的问题提供一些建议：uncoalesced memory accesses、分配的线程/块数量效率低等。您还需要使用优化标志在发布模式下编译您的代码（例如-O3）来获取一些相关的时间。

关于内核优化，你需要找到你的瓶颈（例如你的 44% 内核），分析它，然后应用通常的optimization techniques：

使用设备的有效带宽来计算内核的性能上限 尽量减少主机和设备之间的内存传输 - 即使这意味着在设备上进行效率不高的计算 合并所有内存访问 首选共享内存访问而不是全局内存访问 避免在单个 warp 中执行代码执行分支，因为这会序列化线程

你也可以使用指令级并行（你应该阅读these slides）。

但是，很难知道您何时无法再优化内核。说你的内核的执行时间很小并不意味着什么：与什么相比很小？您是否正在尝试进行一些实时计算？可扩展性是一个问题吗？这些是您在尝试优化内核之前需要回答的一些问题。

另外，您还应该广泛使用error checking，并依赖cuda-memcheck/cuda-gdb 来调试您的代码。

【讨论】：

非常感谢，请问最后一件事与问题无关，如果内核结束后有一个cudaMemcpy 来自要托管的设备和cudaFree?更具体地说，在我上面描述的情况下，在内核声明之前和cudaFree 之后使用gettimeofday 函数会得到正确的结果吗？

cudaMemcpy() 是同步的，所以通常你的 CPU 定时器应该返回一个正确的值。请注意，还有一个异步版本：cudaMemcpyAsync()。您可以查看this page 上描述的示例。