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

Question

我使用两个内核，我们称它们为 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 的新手，我不明白到底发生了什么。我应该尝试使内核更高效还是它们已经高效？

Answer 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 来调试您的代码。