Когда обучающая ферма гонит градиенты между узлами через сотни гигабит, каждая лишняя копия данных стоит дороже, чем кажется в презентациях. Классический путь выглядел так: буфер видеокарты копировался в оперативную память, оттуда сетевая карта читала данные по RDMA, а на приёме всё повторялось в обратном порядке, и на каждом узле накапливались две лишние пересадки через PCIe и две работы процессора. GPUDirect RDMA убирает промежуточные остановки: сетевой контроллер читает и пишет память видеокарты напрямую через PCIe, минуя оперативную память и процессор, и путь данных становится одним отрезком вместо четырёх. Linux поддерживает механизм двумя путями, через модуль совместимости от вендора и через стандартный интерфейс dma-buf, и разница между ними хорошо видна в коде регистрации памяти. Заслуживают упоминания и родственные механизмы с похожими названиями: GPUDirect Storage строит путь между накопителями и видеокартой со своим стеком, а библиотека gdrcopy решает задачу быстрого доступа процессора к памяти видеокарты, и все три живут рядом, но решают разные уравнения.
Разберём механику пути данных, оба способа регистрации, влияние топологии PCIe, проверку и типичные ловушки.
Как данные попадают из видеокарты в сеть
Секрет механизма в отображении памяти видеокарты для сторонних устройств. Драйвер видеокарты умеет выставить окно своей памяти в конфигурационное пространство, и сетевой контроллер получает возможность адресовать эту память как обычный регион DMA, то есть читать и писать её напрямую по шине. Для приложения меняется только регистрация: буфер, выделенный вызовом cudaMalloc, регистрируется в защите домена RDMA так же, как обычная память, и дальше RDMA-операции пишут и читают его напрямую:
void *gpu_buf;
cudaMalloc(&gpu_buf, bytes);
struct ibv_mr *mr = ibv_reg_mr(pd, gpu_buf, bytes,
IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE);Ключевой момент в том, что вызов регистрации выглядит обычным, а происходит необычное: драйвер сетевой карты видит, что адрес принадлежит видеокарте, и вместо отображения страниц оперативной памяти строит отображение окна видеокарты. Физически сетевой контроллер адресует область через окно BAR1 видеокарты, поэтому объём и раскладка этого окна становятся частью плана: регистрируемые под сеть буферы должны попадать в выставленное окно, и драйвер видеокарты управляет им динамически, подстраиваясь под запросы. Пакет при этом движется так: контроллер получает дескриптор с адресом в видеокарте, делает запись DMA прямо в неё, а процессор видит только завершение, и ни байта данных через него не проходит. Выгод двойная: исчезают копии через процессор, экономя его время и полосу PCIe от видеокарты к памяти, и падает задержка, потому что путь пакета сокращается на две пересадки. Для сообщений в мегабайты выигрыш в основном в полосе, для мелких в задержке, и оба эффекта на фермах обучения складываются в заметную долю общей производительности.
Два пути регистрации GPU-памяти
Первый путь, исторический, это модуль совместимости от вендора. Модуль nvidia_peermem подключается к API одноранговой памяти в стеке RDMA от Mellanox и перехватывает регистрации адресов видеокарты, отдавая сетевому драйверу корректные физические окна:
modprobe nvidia_peermem
lsmod | grep -E 'peermem|ib_uverbs'
dmesg | grep -i peermemМодуль появился в драйверах серии 450 и требует согласованной пары версий с драйвером видеокарты и стеком RDMA, поэтому сборка фермы начинается со сверки матрицы совместимости, и несогласованные версии всплывают как пустой список устройств при регистрации. Сравнение путей прагматичное: модуль совместимости проверен годами на многих фермах, но живёт вне мейнстрима, а dma-buf путь моложе, зато опирается на стандартные интерфейсы и переживает обновления без плясок с проприетарными пакетами.
Второй путь, современный, стандартный dma-buf. Ядро начиная с 5.12 умеет регистрировать память через дескриптор dma-buf в вызове ibv_reg_dmabuf_mr, а драйвер видеокарты свежих версий умеет такой дескриптор отдавать:
int dmabuf_fd = cuda_get_dmabuf(gpu_buf, bytes);
struct ibv_mr *mr = ibv_reg_dmabuf_mr(pd, 0, bytes,
(uint64_t)gpu_buf, dmabuf_fd,
IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE);Здесь буфер экспортируется подсистеме dma-buf, и сетевой драйвер строит отображение из него, без сторонних модулей совместимости. Поддержка на стороне стека RDMA появилась в ядре 5.12 и в rdma-core начиная с версии 36, где и объявлена функция регистрации через дескриптор dma-buf, а на стороне видеокарты требуется драйвер, умеющий экспортировать память CUDA в эту подсистему, что пришло в свежих сериях драйвера. Путь через dma-buf чище архитектурно, лучше живёт в мейнстриме и не требует согласованных проприетарных модулей, но требует свежести по всей линии: ядро, rdma-core, драйвер видеокарты и CUDA должны поддерживать экспорт. Для новых фер разумный выбор dma-buf, для проверенных конфигураций на старых стеках модуль совместимости остаётся рабочим решением.
Топология PCIe и её влияние
Прямой путь чувствителен к топологии сильнее обычного RDMA, потому что трафик теперь идёт не к процессору, а между устройствами. Лучшая схема, когда сетевой контроллер и видеокарта сидят за одним коммутатором PCIe или на одном корневом порту, худшая, когда между ними стоит процессор другого NUMA-узла. Раскладку показывает nvidia-smi:
nvidia-smi topo -mЛегенда вывода расшифровывает соседство: PIX означает один коммутатор, PXB общий корневой порт, PHB общий процессор, SYS пересечение сокетов, и для GPUDirect RDMA целевой класс это PIX или PXB, где трафик между устройствами не заходит в когерентную фабрику процессора. Ввод-вывод через IOMMU тоже влияет: трансляция адресов на пути между устройствами добавляет задержку и иногда просто ломает одноранговые передачи, поэтому на узлах под прямой RDMA трансляцию либо выключают, либо ставят в проходной режим с параметром iommu=pt, сверяя решение с политикой безопасности стенда. Типовая ферма под обучение ставит раскладку рельсами: каждая видеокарта получает собственный контроллер на том же коммутаторе PCIe, восемь карт дают восемь рельс, и кольца обмена распределяются по ним равномерно, что делает прямой путь естественным, а не надстройкой. Линии PCIe важны по-прежнему: шестнадцать линий пятого поколения дают паре порядка шестидесяти гигабайт в секунду, и конфигурация с урезанными линиями прячет выгоду прямого пути.
Проверка механизма по слоям и тестирование
Проверка работоспособности делается по слоям. Первый слой, факт загрузки модулей и сообщений в журнале, о котором речь выше. Второй слой, микротест на регистрирующей стороне: небольшая программа выделяет буфер cudaMalloc, регистрирует его в домене и гоняет ibv_post_send с чтением или записью на соседний узел, и успех регистрации уже доказывает, что связка драйверов работает. Стандартные измерители полосы в такой конфигурации исторически не работали с памятью видеокарты, поэтому мерить приходилось собственными средствами, и только свежие сборки перфестов начали догонять, так что собственный микротест в наборе инструментов фермы не лишний. Базовую проверку сети без видеокарты при этом никто не отменял:
ib_write_bw -d mlx5_0 -Rс серверной и клиентской стороной на паре узлов. Третий слой, полоса и задержка на реальном размере сообщений, где стандартом де-факто служат тесты NCCL:
NCCL_DEBUG=INFO all_reduce_perf -b 8M -e 512M -g 8С включённой отладкой NCCL печатает, какие сетевые пути выбрал, и среди строк видны упоминания прямого пути, когда механизм задействован. Сравнение с выключенным механизмом делается принудительным ограничением уровня, и разница в полосе на больших сообщениях обычно составляет десятки процентов, а в задержке мелких сообщений единицы микросекунд.
NCCL и обучающие фермы
На практике большинство фер получают прямой путь через NCCL, не написав ни строчки RDMA-кода. Библиотека сама регистрирует GPU-буферы при обмене между узлами и сама проверяет топологию, а настройка сводится к переменным окружения. Уровень допуска задаёт NCCL_NET_GDR_LEVEL, ограничивая, насколько близко по топологии должны сидеть карта и контроллер, чтобы путь считался допустимым, список контроллеров управляется NCCL_IB_HCA, а отладочный уровень NCCL_DEBUG выводит принятые решения. Для фер с несколькими сетевыми интерфейсами на узел, так называемыми рельсами, NCCL раскладывает кольца по рельсам автоматически, и прямой путь делает каждую рельс полноценным участником, а не узким горлышком с копированиями. Отдельные переменные управляют деталями: NCCL_NET_GDR_READ ограничивает, на каких направлениях данные читаются сетевым контроллером напрямую из памяти видеокарты, поскольку чтение дороже записи и на некоторых топологиях выгоднее отдать его другой стороне, а версия библиотеки определяет, насколько умно раскладываются кольца без ручных подсказок. На сетях с аппаратной поддержкой коллективных операций прямой путь сочетается с выгрузкой коллективов в коммутаторы, и выигрыши складываются.
Экономика на уровне фермы прямая: копирования через процессор съедали не только полосу, но и ядра, которые иначе считали батчи. Прямой путь возвращает процессоры расчётам, и на узлах с восемью видеокартами это заметный процент узла целиком, что объясняет, почему механизм давно стал обязательной частью спецификаций кластеров под обучение.
Типичные ошибки и ограничения
Первая ошибка, неверная топология: контроллер за процессором чужого узла превращает прямой путь в обходной, и полоса падает до уровня обычного пути с копированием, поэтому раскладку PCIe сверяют до монтажа, а не после недельного поиска потерь. Вторая, IOMMU: включённая трансляция без проходного режима блокирует одноранговые передачи, и симптом выглядит как отказ регистрации либо как сообщения драйвера о невозможности маппинга. Третья, рассинхрон версий модуля совместимости с драйвером видеокарты, что лечится либо согласованием версий, либо переходом на dma-buf путь.
Отдельные ограничения стоит знать заранее. Регистрация памяти видеокарты удерживает страницы резидентными, и большие объёмы под прямой путь уменьшают гибкость памяти, поэтому размер буферов под сеть планируют отдельно от рабочих наборов. Мелкие сообщения через прямой путь не всегда выгоднее обычных, и библиотеки сами переключают порог, так что принудительное включение на всё подряд не рекомендуется. И наконец, не путать с родственным механизмом GPUDirect Storage: тот решает задачу пути между накопителями и видеокартой, со своим стеком и своими инструментами, и требования к топологии у него свои.
Практический порядок внедрения в кластере
Ввод механизма в строй укладывается в шесть шагов:
- сверить матрицу версий драйвера видеокарты, стека RDMA и модуля совместимости либо подтвердить готовность всей линии к dma-buf;
- проверить топологию через nvidia-smi topo -m и убедиться, что пары карта-контроллер сидят в целевом классе соседства;
- выставить режим трансляции адресов согласованно с политикой безопасности и проверить регистрации тестовой программой;
- прогнать тесты NCCL с отладкой и зафиксировать полосу и задержку на целевых размерах сообщений;
- сравнить с отключённым прямым путём на том же стенде и записать выигрыш в базу знаний фермы;
- закрепить конфигурацию в провижининге, включая переменные окружения и проверку модулей в самодиагностике узла.
Итог у механизма рабочий и проверенный годами: прямой путь превращает сетевой обмен видеокарт из трёхостановочного маршрута в один отрезок шины, возвращает процессорам время на расчёты и делает масштабирование колец обмена предсказуемым. Появление стандартного dma-buf пути снимает исторический минус в виде проприетарных модулей совместимости, и на свежих стеках настройка сводится к правильной топологии и паре переменных окружения, что делает механизм доступным без глубокого погружения в RDMA-код.