调用cuda内核时的性能损失

Question

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

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

我之所以这么说是因为我正在构建的应用程序重复调用几个内核(没有内存在调用之间重新读/写到设备)而且我想知道是否将这些内核调用包装到单个内核中调用(somekernel1-3成为设备功能)会在性能上产生任何有意义的差异.

Answer 1

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

通常,"融合"多个数据操作会有性能提升,否则这些操作将在单独的内核中完成到单个内核中,算法允许这样做.GPU具有比峰值内存带宽高得多的算术峰值性能,因此每个内存事务(以及每个内核"设置代码")可以执行的FLOP越多,内核的性能就越好.另一方面,尝试编写一个"瑞士军刀"风格的内核试图将完全不同的操作塞进一段代码中,这绝不是一个特别好的主意,因为它增加了寄存压力并降低了L1之类的效率,常量内存和纹理缓存.

您选择哪种方式应该真正受到代码/算法性质的指导.我不相信这个问题有一个"正确"的答案可以适用于所有情况.