Плохая производительность при одновременном вызове cudaMalloc с двумя графическими процессорами

Question

У меня есть приложение, в котором я разделяю нагрузку обработки между графическими процессорами в пользовательской системе. В принципе, на каждый графический процессор имеется поток ЦП, который инициирует цикл обработки GPU, когда он запускается периодически по основному потоку приложения.

Рассмотрите следующее изображение (сгенерированное с использованием инструмента профилирования NVIDIA CUDA) для примера интервала обработки графического процессора - здесь приложение использует один графический процессор.

Как вы можете видеть, большая часть времени обработки GPU потребляется двух сортировочных операций, и я использую библиотеку Thrust для этого (тяги :: sort_by_key). Кроме того, похоже, что thrust :: sort_by_key вызывает несколько cudaMallocs под капотом, прежде чем он начнет фактический сортировку.

Теперь рассмотрим тот же интервал обработки, когда приложение распределения нагрузки обработки в течение двух графических процессоров:

В идеальном мире можно было бы ожидать интервал обработки 2 GPU, чтобы быть точно половину от один GPU (потому что каждый GPU выполняет половину работы). Как вы можете видеть, это не так, частично потому, что cudaMallocs, кажется, занимает больше времени, когда их называют одновременно (иногда в 2-3 раза больше) из-за какой-то вопрос раздора. Я не понимаю, почему это должно быть так, потому что пространство для выделения памяти для 2 графических процессоров полностью независимо, поэтому не должно быть общесистемной блокировки cudaMalloc - блокировка за один GPU была бы более разумной.

Чтобы доказать свою гипотезу, что проблема связана с одновременными вызовами cudaMalloc, я создал смехотворно простую программу с двумя потоками ЦП (для каждого графического процессора), каждый раз вызывающий cudaMalloc несколько раз. Я первый запустил эту программу так, что отдельные нити не называют cudaMalloc в то же время:

Вы видите это занимает ~ 175 микросекунд на распределение. Далее, я запустил программу с нити вызова cudaMalloc одновременно:

Здесь каждый вызов принял 538 ~ микросекунд или в 3 раза дольше, чем в предыдущем случае! Излишне говорить, что это значительно замедляет мое приложение, и разумно, что проблема будет только ухудшаться с более чем 2 графическими процессорами.

Я заметил это поведение в Linux и Windows. В Linux я использую драйвер Nvidia версии 319.60, а в Windows я использую версию 327.23. Я использую CUDA toolkit 5.5.

Возможная причина: Я использую GTX 690 в этих тестах. Эта карта в основном состоит из 6 680-подобных графических процессоров, размещенных в одном устройстве. Это единственная настройка «multi-GPU», которую я запускал, поэтому, возможно, проблема cudaMalloc связана с некоторой аппаратной зависимостью между 6 GPU 690?

Answer 1

Резюмируя проблему и дать возможное решение:

cudaMalloc утверждение, вероятно, связано с утверждением на уровне драйвера (возможно, из-за необходимости переключения контекстов устройств, как talonmies suggestsed) и один может избежать этой дополнительной задержки в высокопроизводительные критические разделы с помощью cudaMalloc-ing и временных буферов заранее.

Похоже, мне, вероятно, нужно реорганизовать мой код, чтобы я не вызывал никакой процедуры сортировки, которая вызывает cudaMalloc под капотом (в моем случае thrust :: sort_by_key). CUB library выглядит перспективным в этом отношении. В качестве бонуса CUB также предоставляет пользователю параметр потока CUDA, что также может повысить производительность.

Для получения подробной информации о переходе от тяги к CUB см. CUB (CUDA UnBound) equivalent of thrust::gather.

UPDATE:

Я попятился вызовы упорным :: sort_by_key в пользу куб :: DeviceRadixSort :: SortPairs.
Выполнение этой бритой миллисекунды с моего времени обработки в течение интервала. Также проблема разногласий с несколькими GPU разрешилась сама собой - разгрузка до 2 графических процессоров почти снижает время обработки на 50%, как и ожидалось.

Answer 2

Я предопределю это с отказом от ответственности: я не привязан к внутренним устройствам драйвера NVIDIA, поэтому это несколько умозрительно.

Замедление, которое вы видите, - это просто конфликт на уровне драйвера, вызванный конкуренцией со стороны нескольких потоков, вызывающих устройство malloc одновременно. Для распределения памяти устройства требуется ряд системных вызовов ОС, равно как и переключение контекста на уровне драйвера. В обеих операциях существует нетривиальная сумма задержек. Вероятно, дополнительное время, которое вы видите, когда два потока пытаются и выделяют память одновременно, вызвано дополнительной задержкой драйвера от переключения с одного устройства на другое в течение последовательности системных вызовов, необходимых для распределения памяти на обоих устройствах.

Я могу придумать несколько способов, вы должны быть в состоянии смягчить это:

• Вы можете уменьшить системный вызов накладных расходы упорного распределения памяти до нуля, написав свой собственный засунул распределитель памяти для устройство, которое работает с плиты памяти, выделенной во время инициализации . Это избавит вас от всей накладной системного вызова в каждом sort_by_key, но усилие написания собственного пользователя диспетчер памяти является нетривиальным. С другой стороны, он оставляет остатки вашего кода тяги.
• Вы можете переключиться на альтернативную библиотеку сортировки и вернуть самостоятельно управлять распределением временной памяти. Если вы выполняете все выделение на этапе инициализации, стоимость одноразового выделения может быть амортизирована почти до нуля в течение срока службы каждого потока.

В коде на основе линейной алгебры с несколькими GPU на основе CUBLAS, которые я написал, я объединил обе идеи и написал автономный диспетчер памяти пространства пользователя, который работает с одноразовым распределенным пулом памяти устройства. Я обнаружил, что удаление всех накладных расходов на распределение памяти промежуточных устройств дало полезную скорость. Ваш вариант использования может воспользоваться аналогичной стратегией.