同时使用2个GPU调用cudaMalloc时性能不佳

Question

我有一个应用程序,我在用户系统上的GPU之间分配处理负载.基本上,每个GPU都有CPU线程,当主应用程序线程定期触发时,它会启动GPU处理间隔.

考虑以下图像(使用NVIDIA的CUDA探查器工具生成)作为GPU处理间隔的示例- 此处应用程序使用单个GPU.

正如您所看到的,两个排序操作消耗了大部分GPU处理时间,而我正在使用Thrust库(thrust :: sort_by_key).此外,看起来push :: sort_by_key会在启动实际排序之前调用几个cudaMallocs.

现在考虑应用程序在两个GPU上分散处理负载的相同处理间隔:

在完美的世界中,您可以预期2 GPU处理间隔恰好是单GPU的一半(因为每个GPU的工作量只有一半).正如你所看到的,部分原因并非如此,因为cudaMallocs由于某种争用问题而被同时调用(有时长2-3倍)时似乎需要更长的时间.我不明白为什么会出现这种情况,因为2个GPU的内存分配空间是完全独立的,因此cudaMalloc上不应该有系统范围的锁定 - 每GPU锁定会更合理.

为了证明我的假设问题是同时使用cudaMalloc调用,我创建了一个非常简单的程序,它有两个CPU线程(每个GPU),每个线程多次调用cudaMalloc.我首先运行此程序,以便单独的线程不会同时调用cudaMalloc:

你看,每次分配需要大约175微秒.接下来,我用同时调用cudaMalloc的线程运行程序:

在这里,每个呼叫比前一个案例花了大约538微秒或3倍!毋庸置疑,这极大地减慢了我的应用程序,并且理所当然,只有2个以上的GPU才会使问题变得更糟.

我在Linux和Windows上注意到了这种行为.在Linux上,我使用的是Nvidia驱动程序版本319.60,而在Windows上我使用的是327.23版本.我正在使用CUDA工具包5.5.

可能的原因: 我在这些测试中使用的是GTX 690.这张卡基本上是2 680个GPU,安装在同一个单元中.这是我运行的唯一"多GPU"设置,所以cudaMalloc问题可能与690的2 GPU之间的硬件依赖性有关吗？

Answer 1

我将在免责声明的前言:我不知道NVIDIA驱动程序的内部,所以这有点推测.

您看到的速度减慢只是由多个线程同时调用设备malloc的竞争引起的驱动程序级别争用.设备内存分配需要许多OS系统调用,驱动程序级上下文切换也需要.两种操作都存在非常重要的延迟.当两个线程同时尝试并分配内存时,您看到的额外时间可能是由于在两个设备上分配内存所需的系统调用序列中从一个设备切换到另一个设备的额外驱动程序延迟.

我可以想到一些你应该能够缓解这种情况的方法:

• 您可以通过为在初始化期间分配的内存块工作的设备编写自己的自定义推力内存分配器,将推力内存分配的系统调用开销减少到零.这将消除每个系统中的所有系统调用开销sort_by_key,但编写自己的用户内存管理器的努力并非易事.另一方面,它保留了其余的推力代码.
• 您可以切换到备用排序库并自行恢复管理临时内存的分配.如果在初始化阶段进行所有分配,则一次性内存分配的成本可以在每个线程的生命周期内分摊到几乎为零.

在我编写的基于多GPU CUBLAS的线性代数代码中,我结合了这两个想法并编写了一个独立的用户空间设备内存管理器,它可以处理一次性分配的设备内存池.我发现去除中间设备内存分配的所有开销成本产生了有用的加速.您的用例可能会受益于类似的策略.

Answer 2

总结一下问题并给出可能的解决方案：

cudaMalloc 争用可能源于驱动程序级别争用（可能是由于需要按照 talonmies 的建议切换设备上下文），并且可以通过事先 cudaMalloc 和临时缓冲区来避免性能关键部分中的这种额外延迟。

看起来我可能需要重构我的代码，这样我就不会调用任何在幕后调用 cudaMalloc 的排序例程（在我的例子中推力::sort_by_key）。CUB图书馆 在这方面看起来很有前途。作为奖励，CUB 还向用户公开 CUDA 流参数，这也可以提高性能。

有关从推力移动到 CUB 的一些详细信息，请参阅相当于推力的 CUB (CUDA UnBound)::gather 。

更新：

我放弃了对 Thrust::sort_by_key 的调用，转而使用 cub::DeviceRadixSort::SortPairs。
这样做可以减少我的每个时间间隔的处理时间。此外，多 GPU 争用问题已自行解决 - 卸载到 2 个 GPU 几乎可以将处理时间缩短 50%，正如预期的那样。