c++ - 调用 cuda 内核时的性能损失

Question

我想知道在 C/C++ 中执行 cuda 内核调用的开销是多少，如下所示:

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

我问这个问题的原因是因为我正在构建的应用程序当前对多个内核进行重复调用(调用之间没有内存被重新读取/写入设备)，我想知道是否将这些内核调用包装到单个内核调用(某些内核 1-3 成为设备函数)会对性能产生任何有意义的差异。

Answer 1

在非 WDDM Windows 平台上，使用运行时 API 启动内核的主机端开销仅为大约 15-30 微秒。在 WDDM 平台(我不使用)上，我知道它可以高得多，而且驱动程序中有某种批处理机制，它试图通过在单个驱动程序端操作中执行多个操作来分摊成本。

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

您选择走哪条路应该真正以代码/算法的性质为指导。我不相信这个问题有一个可以适用于所有情况的单一“正确”答案。