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

我有一个应用程序，可以在用户系统中的GPU之间分配处理负载。基本上，每个GPU有一个CPU线程，当由主应用程序线程周期性地触发时，它启动一个GPU处理间隔。使用2个GPU同时调用cudaMalloc时的性能不佳

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

enter image description here

正如你所看到的GPU处理时间的很大一部分是由两个分拣作业消耗，我使用这个（推力:: sort_by_key）的推力库。另外，在启动实际排序之前，它看起来像thrust :: sort_by_key在引擎盖下调用一些cudaMallocs。

现在考虑同样的处理间隔，其中应用已经普及的处理负荷超过两个GPU：

enter image description here

在一个完美的世界里，你所期望的2 GPU处理间隔正好一半的单GPU（因为每个GPU都在做一半的工作）。正如您所看到的，这并非部分原因，因为由于某种争用问题，cudaMallocs在同时调用时（有时延长2-3倍）似乎需要更长的时间。我不明白为什么需要这样做，因为这两个GPU的内存分配空间是完全独立的，所以不应该有cudaMalloc上的全系统锁定 - 每个GPU锁定会更合理。

为了证明我的假设，即同时发生的cudaMalloc调用问题，我创建了一个带有两个CPU线程（每个GPU）的可笑简单程序，每个线程多次调用cudaMalloc。我第一次运行这个程序，使得单独的线程不会在同一时间拨打cudaMalloc：

enter image description here

你看，它需要每〜175分配微秒。接下来，我跑的程序与线程同时调用cudaMalloc：

enter image description here

在这里，每个呼叫拿了〜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之间的一些硬件依赖有关？

来源

2013-10-05 rmccabe3701

高性能代码的常见建议是使malloc操作脱离任何性能循环。我意识到这并不是一个微不足道的问题，因为你使用推力。有高性能的排序库可以替代推力sort_by_key，这将允许您提前进行分配并将其重新用于排序操作。 [CUB]（http://nvlabs.github.io/cub/），[b40c]（http://code.google.com/p/back40computing/）和[MGPU]（http：//nvlabs.github .io/moderngpu /）都是可能的。 –

是的，我已经看过CUB和B40C（B40C网站说该项目已被弃用）。在我做清除推力的工作之前，我想看看图书馆之间的一些比较图。你能指点我一些表演数字吗？你推荐哪个图书馆？ ......似乎推力并不是非常高的性能，例如，我已经用我自己的定制内核交换了一堆推力:: reduce和reduce_by_key调用 - 这样做将我的处理时间减半。不是开玩笑。 – rmccabe3701

推力实际上是基于b40c的一个特定变体（或曾经是）。对于等效的测试用例，我在b40c和MGPU之间的测试没有太大差异。在我运行的一个测试中，我只对32位值的22位进行排序。 MGPU有一个拨号盘，我可以转而只用22位分辨率，我观察到这样做的速度提高了40％。我没有使用CUB多。如果你浏览这些链接，你可能会发现一些性能数据。例如，某些MGPU性能数据[here]（http://nvlabs.github.io/moderngpu/performance.html#performance） –

总结问题和给出一个可能的解决方案：

的cudaMalloc争可能是由驱动程序级别的竞争（可能是由于需要切换设备上下文为talonmies suggestsed）茎和一个能避免这种额外的延迟性能关键部分预先通过cudaMalloc-ing和临时缓冲区。

它看起来像我可能需要重构我的代码，以便我不调用任何排序例程，调用引擎盖下的cudaMalloc（在我的情况下是thrust :: sort_by_key）。在这方面，CUB library 看起来很有希望。作为奖励，CUB还向用户公开了CUDA流参数，这也可以提高性能。

查看CUB (CUDA UnBound) equivalent of thrust::gather了解从推力转向CUB的一些细节。

UPDATE：

我退缩了有利于幼崽:: DeviceRadixSort :: SortPairs的推力:: sort_by_key呼叫。
这样做我的每间隔处理时间减少了毫秒。此外，多GPU争用问题已经解决 - 卸载至2个GPU几乎可以将处理时间缩短50％，如预期的那样。

来源

2013-10-06 16:00:21 rmccabe3701

如果你可以通过这个和你较旧的CUDA问题并接受一些你认为适合的答案，这将是一件好事。它将它们从未得到答复的列表中删除（我们主动尽量保持尽可能短的内容），并且通过搜索可以使其他人更容易找到您是否可以这样做。谢谢。 – talonmies

哎呀，对不起，我曾想过，当答案被投票时，它会被“接受”。我回去接受了一些对我的旧问题的回答。再一次，对不起，我对这个网站还是有些新鲜的。 – rmccabe3701

我将在前言中声明一个免责声明：我并不知道NVIDIA驱动程序的内部，所以这是有点投机的。

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

我能想到的几个方法，你应该能够减轻这一点：

你可以推的内存分配的系统调用的开销降低到零通过编写自己的自定义推力内存分配器的该设备在初始化过程中分配一块内存。这将消除每个sort_by_key内的所有系统调用开销，但编写自己的用户内存管理器的努力并非微不足道。另一方面，它会使您的推力代码的其余完好无损。
您可以切换到另一个排序库并收回自己管理临时内存的分配。如果在初始化阶段执行所有分配，则在每个线程的使用期限内，一次内存分配的成本可以分摊到几乎为零。

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

来源

2013-10-05 10:22:48 talonmies

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

回答

相关问题