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

Posted 2023-02-19

【中文标题】同时使用 2 个 GPU 调用 cudaMalloc 时性能不佳

【英文标题】：Poor performance when calling cudaMalloc with 2 GPUs simultaneously

【发布时间】：2013-10-12 02:48:57

【问题描述】：

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

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

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

现在考虑应用程序将处理负载分散到两个 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 之间的某些硬件依赖有关？

【问题讨论】：

对高性能代码的通常建议是让 malloc 操作脱离任何性能循环。我意识到这不是一件小事，因为您使用的是推力。有高性能的排序库可以替换推力 sort_by_key，这将允许您提前进行分配并将它们重用于排序操作。 CUB、b40c、MGPU都是可能的。

是的，我研究过 CUB 和 b40c（b40c 网站说该项目已被弃用）。在我做去除推力的工作之前，我想看看这些库之间的一些比较图。你能指出一些性能数据吗？你推荐哪个图书馆？ ... 似乎推力不是很高的性能，例如，我已经用我自己的自定义内核切换了一堆推力::reduce 和 reduce_by_key 调用——这样做将我的处理时间缩短了一半。不开玩笑。

Thrust 实际上是基于 b40c 的特定变体（或曾经是）。对于等效的测试用例，我的测试在 b40c 和 MGPU 之间没有太大区别。在我运行的一项测试中，我只对 32 位值的大约 22 位进行排序。 MGPU 有一个表盘，我只能在 22 位上进行排序，我观察到这样做比推力快了 40%。我没怎么用过CUB。如果您浏览这些链接，您可能会发现一些性能数据。例如一些 MGPU 性能数据here

如果不清楚，我建议这些备用库的意思并不是说它们具有比推力更高的排序性能（尽管它们可能，但我不确定您的测试用例中的结果是什么) 但它们允许选择拆分推力正在执行的临时数据分配，以便您可以预先处理这些。

【参考方案1】：

我先声明一下：我不了解 NVIDIA 驱动程序的内部结构，所以这有点推测。

您看到的减速只是由同时调用设备 malloc 的多个线程的竞争引起的驱动程序级争用。设备内存分配需要许多操作系统系统调用，驱动程序级别的上下文切换也是如此。两种操作都有不小的延迟。当两个线程同时尝试分配内存时，您看到的额外时间很可能是由在两个设备上分配内存所需的系统调用序列中从一个设备切换到另一个设备的额外驱动程序延迟引起的。

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

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

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

【参考方案2】：

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

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

看起来我可能需要重构我的代码，这样我就不会调用任何在后台调用 cudaMalloc 的排序例程（在我的例子中是thrust::sort_by_key）。 CUB library 在这方面看起来很有希望。作为奖励，CUB 还向用户公开了一个 CUDA 流参数，这也可以提高性能。

有关从推力转移到 CUB 的一些详细信息，请参阅 CUB (CUDA UnBound) equivalent of thrust::gather。

更新：

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

【讨论】：

如果你能通过这个和你以前的 CUDA 问题并在你认为合适的地方接受一些答案，那就太好了。它将它们从未答复的列表中删除（我们积极尝试使其尽可能短），并且如果您这样做，其他人可以更容易地通过搜索找到它们。谢谢。

哎呀，对不起，我原以为当一个答案被投票时，它会被“接受”。我回去接受了一堆旧问题的答案。再次抱歉，我对这个网站还是有点陌生​​。