cuda - 同时使用 2 个 GPU 调用 cudaMalloc 时性能较差

Question

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

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

如您所见，两个排序操作消耗了很大一部分 GPU 处理时间，我为此使用 Thrust 库 (thrust::sort_by_key)。另外，看起来 Thrust::sort_by_key 在开始实际排序之前会在后台调用一些 cudaMalloc。

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

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

为了证明我的假设(问题在于同时调用 cudaMalloc)，我创建了一个极其简单的程序，其中有两个 CPU 线程(针对每个 GPU)，每个线程调用 cudaMalloc 多次。我首先运行这个程序，以便单独的线程不会同时调用 cudaMalloc:

您会发现每次分配大约需要 175 微秒。接下来，我使用同时调用 cudaMalloc 的线程运行该程序:

这里，每个调用花费了约 538 微秒，或者说是前一个案例的 3 倍!不用说，这极大地减慢了我的应用程序的速度，而且按理说，如果 GPU 数量超过 2 个，问题只会变得更糟。

我注意到 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 的多个线程的竞争而引起的驱动程序级别争用。设备内存分配需要许多操作系统系统调用，驱动程序级别上下文切换也是如此。这两个操作都存在相当大的延迟。当两个线程尝试同时分配内存时，您看到的额外时间很可能是由在两个设备上分配内存所需的系统调用序列中从一个设备切换到另一个设备的额外驱动程序延迟引起的。

我可以想到一些可以减轻这种情况的方法:

• 您可以减少推力内存分配的系统调用开销通过编写您自己的自定义推力内存分配器来为零使用期间分配的一 block 内存工作的设备初始化。这将消除所有系统调用开销在每个 sort_by_key 中，但需要编写自己的用户内存管理器并不简单。另一方面它留下了其余的您的推力代码完好无损。
• 您可以切换到替代排序库并收回自己管理临时内存的分配。如果你做了所有的在初始化阶段分配，一次性的成本内存分配在整个生命周期内可以摊销到几乎为零每个线程。

在我编写的基于多 GPU CUBLAS 的线性代数代码中，我结合了这两种想法并编写了一个独立的用户空间设备内存管理器，该管理器在一次性分配的设备内存池上工作。我发现消除中间设备内存分配的所有开销成本可以带来有用的加速。您的用例可能受益于类似的策略。