调用CUDA内核时的性能惩罚

Question

我想知道在C/C++中执行cuda内核调用的开销，例如以下代码：

```c++ cuda_kernel<<>>(args); ```

somekernel1<<<blocks,threads>>>(args);
somekernel2<<<blocks,threads>>>(args);
somekernel3<<<blocks,threads>>>(args);

我之所以问这个问题，是因为我目前正在构建的应用程序会反复调用多个内核（在调用之间不会重新读取/写入设备内存），我想知道将这些内核调用包装成单个内核调用（其中一些内核变为设备函数）是否会对性能产生任何有意义的差异。

- NothingMore

Answer 1

使用运行时API进行内核启动的主机端开销在非WDDM Windows平台上只有大约15-30微秒。在我不使用的WDDM平台上，我了解到它可能会高得多，而且驱动程序中也有一些批处理机制，尝试通过在单个驱动程序端操作中执行多个操作来分摊成本。

通常情况下，将多个数据操作“融合”到一个内核中（如果算法允许）会提高性能。 GPU的算术峰值性能比峰值存储带宽高得多，因此每个内存事务（和每个内核“设置代码”）可以执行的FLOP越多，内核的性能就越好。另一方面，试图编写一个“瑞士军刀”样式的内核，试图将完全不同的操作塞进一段代码中，从来不是一个特别好的主意，因为它会增加寄存器压力并降低诸如L1、常量内存和纹理缓存等效率。

你选择哪种方式应该真正由代码/算法的性质来引导。我认为没有一个单一的“正确”答案可以适用于所有情况。

- talonmies

---

我正在尝试避免使用瑞士军刀式的方法，以便在项目之间保持内核共享。感谢您的回复，我只是想确保在进行多个CUDA调用时没有一些我不知道的疯狂性能问题。 - NothingMore

注意：如果您使用的是特斯拉GPU，则可以使用特斯拉计算集群（TCC）驱动程序将性能与非WDDM平台（如XP或Linux）保持一致。关于原始问题，我要强调的是：如果组合内核有助于减少所需的PCI Express传输，则可能值得这样做。如果没有，那么请确保至少重叠Kernel1的计算和向GPU传输Kernel2等数据的过程。 - harrism

talonmies，您所说的非WDDM Windows平台是指哪些？WinXP吗？我非常感兴趣，因为WDDM的惩罚是巨大的，而我无法切换到Linux。目前正在使用win7 x64，并且需要一个x64平台（RAM问题）。 - Dredok

Answer 2

如果您正在Windows上使用Visual Studio Pro，我建议您使用NVidia的Parallel NSight运行测试应用程序，它可以告诉您从方法调用到实际执行的时间戳，无论如何，惩罚是固有的，但如果您的内核持续足够长，则将是可以忽略不计的。

- Caian

---

我没有在Windows上运行（RHEL 6.0，Tesla C2075）。 - NothingMore