Диагностика предзагрузки управляемой памяти GPU на Linux с Nsight Systems и cudaMemPrefetchAsync

Управляемая память CUDA избавляет разработчика от ручного копирования буферов между оперативной памятью и видеокартой: один указатель, полученный вызовом cudaMallocManaged, одинаково работает и на процессоре, и на GPU. Плата за это удобство проявляется внезапно: первый запуск ядра тянется в разы дольше расчетного времени, хотя сам код безупречен. Виноваты страничные ошибки, из-за которых драйвер посреди вычислений перетаскивает страницы данных по шине PCIe и каждый раз останавливает занятые мультипроцессоры. На Linux всю эту скрытую работу можно увидеть поименно: трассировщик Nsight Systems показывает на таймлайне каждую ошибку страницы и каждую миграцию, а вызов cudaMemPrefetchAsync заранее переносит буфер в память конкретного GPU до запуска ядра. В статье разобрано, откуда берутся ошибки страниц, как посчитать их количество и цену в микросекундах, как оформить предзагрузку и как убедиться числами, что она действительно сработала.

Устройство управляемой памяти и происхождение страничных ошибок на GPU

Управляемая память живет в едином виртуальном адресном пространстве: процессор и видеокарта видят один и тот же указатель, а физическое место каждой страницы определяет драйвер. Сразу после выделения вызовом cudaMallocManaged страницы размещаются в оперативной памяти системы, поэтому первый доступ ядра к ним оборачивается страничной ошибкой на стороне GPU. Драйвер останавливает варп, переносит страницу по шине в память видеокарты, обновляет таблицы страниц и только потом разрешает вычисления продолжаться. Полноценная обработка ошибок страниц на GPU появилась с архитектуры Pascal, то есть с вычислительной способности 6.0, а на более старых картах система мигрирует весь буфер целиком при запуске ядра. Степень поддержки конкретной платы отражает атрибут cudaDevAttrConcurrentManagedAccess, о его проверке сказано в конце статьи.

Размер страницы управляемой памяти типично равен двум мегабайтам, что делает цену одной ошибки наглядной. Перенос 2 МБ по PCIe 3.0 при практической пропускной способности около 16 ГБ/с занимает порядка 130 микросекунд, по PCIe 4.0 с ее 32 ГБ/с вдвое меньше, а сверху добавляются десятки микросекунд на обновление таблиц и сброс TLB. Буфер размером 1 ГиБ разбивается на 512 страниц, и если ядро трогает его без подготовки, его время разрастается на сотни остановок. Современные драйверы NVIDIA на Linux опираются на механизм ядра HMM, Heterogeneous Memory Management: он зеркалирует таблицы страниц процессора и позволяет ядру ОС переносить страницы между памятью хоста и памятью GPU по требованию. Механизм вошел в основное дерево ядра Linux в середине 2010-х годов, и с ним поведение управляемой памяти на Linux приблизилось к привычной страничной подкачке. Для разработчика вывод простой: каждая ошибка страницы реальна, стоит реальное время и оставляет след, который видит трассировщик.

Трассировка Nsight Systems с событиями Unified Memory на Linux

Инструмент nsys входит в состав CUDA Toolkit, отдельно ставить его не требуется. Ключевая опция -t перечисляет собираемые виды событий, и среди них um отвечает именно за управляемую память: ошибки страниц, миграции и предзагрузки попадают в отчет отдельными событиями. Компиляция и первый прогон под трассировкой выглядят так:

nvcc -O2 -o um_demo um_demo.cu
nsys profile -t cuda,um,nvtx -o um_report --force-overwrite true ./um_demo

Опция nvtx добавляет пользовательские диапазоны, которыми код помечает фазы вроде заполнения буфера и итераций ядра; в коде их открывают вызовом nvtxRangePushA и закрывают nvtxRangePop. На выходе получается файл um_report.nsys-rep, его открывают в графическом интерфейсе Nsight Systems или разбирают консольной командой nsys stats. Трассировщик фиксирует события и на стороне GPU, и на стороне драйвера, поэтому после прогона видна вся работа по переносу страниц, которая обычно остается невидимой. Для повторяемости стоит убрать постороннюю нагрузку: закрыть браузер и фоновые вычисления, проверить частоты и температуру карты командой nvidia-smi, иначе миграции чужих процессов исказят картину.

Чтение таймлайна и подсчет ошибок в отчете

В графическом интерфейсе под каждым ядром лежит ряд Unified Memory: каждая ошибка страницы рисуется отдельной меткой, каждая миграция показывается стрелкой между хостом и устройством. Первый прогон без предзагрузки выглядит характерно: ядро тянется длинной полосой, усыпанной метками ошибок, а после каждой миграции виден зазор, в течение которого мультипроцессоры простаивают. Разбор отчета укладывается в короткую схему:

1. Найти ряд Unified Memory под первым запуском ядра и посчитать метки ошибок страниц;
2. Сверить число меток с количеством страниц буфера, для 1 ГиБ при страницах по 2 МБ это ровно 512;
3. Сложить зазоры простоя возле миграций и сравнить сумму с расчетным временем ядра;
4. Повторить прогон после правок и убедиться, что все три числа пошли вниз.

Консольный вариант быстрее: команда nsys stats печатает сводку по ядрам и переносам памяти без запуска графики. Фрагмент сводки для примера из следующего раздела, прогон без предзагрузки, выглядит примерно так:

kernel vec_add      96.4 ms
UM page faults      512
UM bytes migrated   1024 MiB

Отдельного внимания заслуживает коварство второго прогона: после первого переноса страницы уже лежат в памяти GPU, ошибки исчезают, и ядро летает. Без холодного запуска и трассировки проблему легко списать на прогрев и забыть о ней до реального запуска, где приложение стартует с чистого листа. Поэтому измеряют именно первый прогон после выделения и заполнения буфера, а диапазоны NVTX помогают отличить итерации друг от друга и не спутать холодный запуск с разогретым.

Предзагрузка буфера в конкретный GPU вызовом cudaMemPrefetchAsync

Функция cudaMemPrefetchAsync переносит диапазон управляемой памяти в память заданного устройства и, как любая асинхронная операция CUDA, ставится в поток. Третьим аргументом служит номер устройства, а константа cudaCpuDeviceId возвращает данные на хост перед обработкой результатов процессором. Правильное место вызова: сразу после заполнения буфера на процессоре и до запуска ядра.

float *data;
size_t bytes = 1024U * 1024U * 1024U;   // 1 GiB
cudaMallocManaged(&data, bytes);
size_t n = bytes / sizeof(float);
for (size_t i = 0; i < n; ++i) data[i] = 1.0f;

cudaMemPrefetchAsync(data, bytes, 0);   // на GPU 0 до запуска ядра
vec_add<<<256, 256>>>(data, n);
cudaDeviceSynchronize();

cudaMemPrefetchAsync(data, bytes, cudaCpuDeviceId); // обратно на хост

Поскольку вызов стоит в том же потоке перед ядром, порядок выполнения CUDA гарантирует, что миграция завершится до начала вычислений, дополнительная синхронизация внутри потока не нужна. Вызов cudaDeviceSynchronize в примере нужен ради замера времени, а не ради корректности. Ценность предзагрузки не только в порядке операций: перенос может перекрываться с независимой работой, пока страницы едут по шине, процессор заполняет следующий буфер, и пауза в конвейере почти исчезает. Обратная предзагрузка на cudaCpuDeviceId перед чтением результатов на хосте экономит тот же такт, но в обратную сторону.

На системе с двумя видеокартами буфер делят на куски и переносят каждый к своему потребителю:

cudaMemPrefetchAsync(part0, half, 0);
cudaMemPrefetchAsync(part1, half, 1);

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

Числовой пример расчета выигрыша от предзагрузки страниц

Возьмем буфер 1 ГиБ из 268 435 456 чисел float и ядро сложения вектора, которое на реальной карте считает свои 8 миллисекунд. Чистый перенос 1073741824 байт по PCIe 4.0 при практических 25 ГБ/с занимает 1073741824 / 25000000000, то есть около 43 миллисекунд. Без предзагрузки эти 512 страниц мигрируют по одной по факту ошибок: к 43 миллисекундам переноса добавляются служебные остановки по 150-200 микросекунд на страницу, и суммарно ядро набирает порядка 80-100 лишних миллисекунд. С предзагрузкой перенос уходит в отдельную операцию потока и стоит свои 45-50 миллисекунд, а ядро возвращает расчетные 8-9 миллисекунд, и часть переноса прячется в перекрытии с работой процессора.

Грубую оценку считают по формуле: время без предзагрузки равно времени чистого ядра плюс число страниц, умноженное на цену одной ошибки. Для примера выходит 8 мс + 512 × 0,17 мс ≈ 95 мс, что совпадает с наблюдаемой картиной на таймлайне.

Для точного замера годится пара событий CUDA:

cudaEvent_t t0, t1;
cudaEventCreate(&t0); cudaEventCreate(&t1);
cudaEventRecord(t0);
vec_add<<<256, 256>>>(data, n);
cudaEventRecord(t1);
cudaEventSynchronize(t1);
float ms = 0.0f;
cudaEventElapsedTime(&ms, t0, t1);

Типичный прогон на карте уровня GeForce RTX с линией PCIe 4.0 дает 96 миллисекунд без предзагрузки против 9 миллисекунд после нее, плюс отдельные 46 миллисекунд на саму предзагрузку, значительная часть которой перекрывается с заполнением следующего буфера. Числа условные, но порядок именно такой, и он проверяется трассировкой за пару минут. Если выигрыш вышел меньше ожиданий, первым делом смотрят долю ошибок на чтение и запись: чтение лечится советами из следующего раздела, запись требует переделки схемы доступа.

Советы cudaMemAdvise и защита от перетаскивания страниц между GPU

Худший сценарий для управляемой памяти: два ядра на разных GPU пишут в один буфер по очереди, и каждая итерация тянет страницы через шину туда и обратно. На таймлайне это видно сразу: частые ошибки на обоих устройствах без пауз между ними, шина загружена почти постоянно, оба ядра тормозят поочередно. Драйвер называет такое поведение перетаскиванием, и лечится оно тремя способами: разделением буфера на куски по GPU, советами cudaMemAdvise или изменением порядка доступа.

Функция cudaMemAdvise подсказывает драйверу намерения приложения:

cudaMemAdvise(data, bytes, cudaMemAdviseSetReadMostly, 0);
cudaMemAdvise(data, bytes, cudaMemAdviseSetPreferredLocation, 0);
cudaMemAdvise(data, bytes, cudaMemAdviseSetAccessedBy, 1);

Совет ReadMostly разрешает держать у каждого читающего устройства собственные копии страниц, и чтение на нескольких GPU перестает провоцировать миграции вообще. PreferredLocation указывает дом, куда страницы стремятся вернуться, AccessedBy связывает диапазон с устройством, которое работает с ним удаленно. С предзагрузкой советы сочетаются в строгом порядке: сначала cudaMemAdvise, затем cudaMemPrefetchAsync для нужного устройства, и только потом ядро. В отчете эффект заметен мгновенно: ошибки на чтение исчезают, записи остаются у владельца страниц, а зазоры простоя сжимаются до незаметности. Бюджет памяти при этом считают заранее, ведь копии для чтения занимают место у каждого читающего устройства, и для гигабайтных буферов эта цена ощутима.

Ограничения архитектур и проверка окружения на Linux

Полноценные страничные ошибки и одновременный доступ к управляемой памяти работают на картах с вычислительной способностью 6.0 и выше, начиная с архитектуры Pascal. На старом железе предзагрузка остается полезной, но миграция идет укрупненными кусками, и тонкая настройка советами может игнорироваться. Поддержку конкретной платы проверяют из кода:

int cm = 0;
cudaDeviceGetAttribute(&cm, cudaDevAttrConcurrentManagedAccess, 0);
printf("concurrent managed access: %d\n", cm);

Версии инструментов и параметры шины на Linux проверяют двумя командами:

nvcc --version | tail -n 1
nvidia-smi --query-gpu=name,pcie.link.gen.current,pcie.link.width.current --format=csv

Вывод второй команды сразу показывает поколение и ширину линии PCIe, а от них напрямую зависит цена каждой миграции:

name, pcie.link.gen.current, pcie.link.width.current
NVIDIA GeForce RTX 4090, 4, 16

Поведение управляемой памяти заметно менялось между выпусками драйвера и Toolkit, поэтому их держат в одной свежести: свежий nsys умеет события um, свежее ядро Linux с HMM ближе к честной постраничной работе, свежий драйвер точнее исполняет советы. Итог укладывается в три строчки таймлайна: трассировка показывает правду о страницах, предзагрузка двигает их вовремя, а числа из отчета превращают догадку о тормозах в измеренный факт.