我想知道执行cuda内核调用的开销在C/C++中,如下所示：
somekernel1<<<blocks,threads>>>(args); somekernel2<<<blocks,threads>>>(args); somekernel3<<<blocks,threads>>>(args);
我问这个问题的原因是因为我正在构建的应用程序重复调用几个内核(没有内存被重新读取/写入到设备之间的调用),我想知道是否将这些内核调用包装到单个内核调用(使用somekernel1-3成为设备功能)将在性能上产生任何有意义的差异.
解决方法
在非WDDM Windows平台上,启动运行时API的内核的主机端开销只有15-30微秒.在WDDM平台(我不使用)上,我明白,它可以要高得多,加上在驱动程序中有一些批处理机制,通过在单个驱动程序端操作中执行多个操作来尝试平摊成本.
通常,“融合”多个数据操作的性能将会提高,否则这些数据操作将在单独的内核中进行单个内核,而算法允许它们. GPU具有比峰值存储器带宽更高的算术峰值性能,因此每个存储器事务(以及每个内核“设置代码”)可以执行的FLOP越多,内核的性能就越好.另一方面,试图写出一个“瑞士军刀”风格的内核,试图将完全不同的操作压缩成一个单独的代码,从来不是一个特别好的主意,因为它增加了寄存器压力并降低了像L1,不断的记忆和纹理缓存.
你选择哪一种方式应该是由代码/算法的本质来指导的.我不相信这个问题有一个单一的“正确的”答案可以在所有情况下应用.
