CUDA速度问题方面的疑问

本人在做一些算法的优化,单独测试每个在GPU上运行的函数速度都能达到120帧以上,但是把四五个这样速度的函数放在一起(都是在GPU上运行,没有CPU和GPU之间相互的数据传输),速度竟然降到了60帧,但是此时的GPU的资源利用率才达到70%左右,并没有满载。函数结合在一起速度怎么降低这么多那?迷惑中,求解答。

楼主您好,将N个小kernel合并成1个大kernel, 不一定是高效的。
也从来没有手册或者资料说合并后就会高效。

相反,很多情况下,合并后反而会降低效率,我给您举个例子,
您有2个kernel, 分别在缓冲区1和缓冲区2上操作。

(1)如果您分开连续运行2个kernel, 那么先操作第一个缓冲区的代码将会被集中执行,然后操作第二端缓冲区的代码也将会被集中执行。
(2)如果您合并在一起,每个线程都需要操作完缓冲区1里的东西继续干缓冲区2的。
此时,只要不是这2个缓冲区的总大小小于L2 cash的大小,那么就会引起L2中的内容在不停的刷出刷入,从而有效了降低了L2的性能,从而进一步降低了您的总体性能。

这是一个例子,但请您知悉这一点。

版主您好,谢谢您的详细解答。我的意思不是把几个kernel函数合成一个kernel。 我的意思是说:我有5个独立运行的函数(functions),运行完function1后运行function2,然后运行function3…单独测试每个function,都很快,但是这五个函数同时运行,就是连贯的运行速度就很慢了。但是此时GPU并没有满载。

如果连贯起来就很慢,您需要注意您的kernel是否过小,

例如一个过小的kernel, 可能执行只需要30us, 但是<<<>>>执行回来一次的时间就有100us的话,那么显然有70us的时间, GPU无工作可干的。

以及,如果能排除kernel过小的因素,您要考虑是否您执行kernel的方式不当,例如您不小心每次都等待(包括使用cudaMemcpy导致的等待)。我建议您这样:
kernel1<<<>>>();
kernel2<<<>>>();
kernel3<<<<>>();
kernel4<<<>>>();
kernel5<<<<>>();
cudaMemcpy(…);

请您考虑这2点可能的因素。

以及,补充下第一点:

如果过小的kernel无法避免,您可以考虑这2个解决方案:
(1)使用更快的CPU, 以便能更快的给GPU发布任务。
(2)同时进行多份任务链(您可以在2个stream这样搞)。

感谢来访。