6

我有一个应用程序，我在用户系统上的 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 之间的某些硬件依赖有关？

4

2 回答 2

6

我将以免责声明开头：我不了解 NVIDIA 驱动程序的内部结构，所以这有点推测。

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

我可以想出一些你应该能够减轻这种情况的方法：

  • 您可以通过为设备编写自己的自定义推力内存分配器来将推力内存分配的系统调用开销减少到零，该设备在初始化期间分配的内存板工作。这将消除每个内部的所有系统调用开销sort_by_key，但编写自己的用户内存管理器的工作并非微不足道。另一方面，它使您的其余推力代码保持不变。
  • 您可以切换到另一个排序库并自己收回管理临时内存的分配。如果您在初始化阶段完成所有分配，那么在每个线程的生命周期内，一次性内存分配的成本可以摊销到几乎为零。

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

于 2013-10-05T10:22:48.447 回答
4

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

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

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

有关从推力移动到 CUB 的一些详细信息，请参见推力::聚集的 CUB (CUDA UnBound) 等效项。

更新：

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

于 2013-10-06T16:00:21.577 回答