Одна видеокарта в сервере обычно достаётся нескольким задачам сразу, и порядок исполнения становится отдельной дисциплиной. Ядра и копирования толпятся в одной очереди, а утилизация SM висит на непонятных процентах. Недобор от дорогого ускорителя чаще сидит в планировании, а не в нехватке мощности. CUDA даёт инструменты, чтобы расписать порядок руками: потоки с приоритетами внутри приложения, служба MPS для дележа карты между процессами, nvidia-smi для контроля. Ниже по нарастающей: очереди и приоритеты, события, MPS с лимитами и партициями, графы и мониторинг.
Потоки CUDA как очереди задач внутри одного приложения
Поток в CUDA, он же stream, это очередь операций, куда программа кладёт запуски ядер и копирования памяти. Внутри очереди всё исполняется в порядке постановки, а потоков может быть сколько угодно: среда распределяет работу между ними по мере освобождения GPU. Очередь создаётся вызовом cudaStreamCreate(&computeStream), а ядро отправляется в неё четвёртым параметром тройных угловых скобок: kernel<<<grid, block, 0, computeStream>>>(d_data).
Главная ловушка новичка сидит в default stream, потоке с номером 0: все операции без явного потока попадают в него, а legacy default stream синхронизируется со всеми блокирующими потоками. Если kernel1 идёт в stream1, kernel2 в default stream, а kernel3 в stream2, то kernel2 ждёт kernel1, затем kernel3 ждёт kernel2, хотя все три могли бы идти параллельно. Лечится двумя способами: потоками с флагом cudaStreamNonBlocking через cudaStreamCreateWithFlags либо переходом с CUDA 7 на per-thread default stream через ключ --default-stream per-thread или макрос CUDA_API_PER_THREAD_DEFAULT_STREAM.
Копирования выпускаются в поток функцией cudaMemcpyAsync, но настоящую асинхронность она показывает только с закреплённой памятью от cudaMallocHost, со страничной вызов деградирует до синхронного. Цифры: буфер в 512 МБ по PCIe 4.0 при реальной пропускной способности около 25 ГБ/с копируется порядка 20 мс. Если копирование идёт параллельно с ядром на 30 мс, конвейер укладывается в 30 мс вместо 50 мс, экономия 40 процентов. Отсюда два правила документации: независимые операции раньше зависимых, а синхронизацию откладывать как можно дольше.
Приоритеты потоков через cudaStreamCreateWithPriority и диапазон уровней
Приоритет задаётся потоку при создании функцией cudaStreamCreateWithPriority: чем меньше число, тем выше место при разборе очередей. Диапазон уровней конкретной карты узнаётся запросом cudaDeviceGetStreamPriorityRange, по умолчанию поток получает приоритет 0.
int minPriority, maxPriority;
cudaDeviceGetStreamPriorityRange(&minPriority, &maxPriority);
cudaStream_t hi, lo;
cudaStreamCreateWithPriority(&hi, cudaStreamNonBlocking, maxPriority);
cudaStreamCreateWithPriority(&lo, cudaStreamNonBlocking, minPriority);Диапазон обычно невелик, часто пара уровней вроде -1 и 0, но зависит от поколения GPU и драйвера, проверять стоит запросом.
Приоритет это подсказка, а не приказ: он не вытесняет запущенные ядра, не гарантирует порядок и касается прежде всего запусков ядер, копирования могут его игнорировать. Смысл у подсказки солидный: когда ядро доедает блоки и на SM освобождаются слоты, планировщик первым разбирает очередь с более высоким приоритетом. Классическая раскладка: латентно-чувствительные запросы в высокоприоритетный поток, объёмный фон в низкоприоритетный, и фон не задерживает интерактив.
События, синхронизация и точный замер времени ядер
Склеивают потоки в конвейер события. cudaEventRecord ставит маркер в поток, cudaStreamWaitEvent заставляет другой поток ждать этот маркер, на таких связках собирается граф зависимостей. Проверка без блокировки идёт через cudaStreamQuery и cudaEventQuery: они возвращают cudaSuccess, когда всё завершено, и cudaErrorNotReady, когда работа идёт. cudaStreamSynchronize синхронизирует один поток, cudaDeviceSynchronize дожидается всех сразу и дороже всего.
События дают и простейший профилировщик. Создав маркеры через cudaEventCreate(&start) и cudaEventCreate(&stop), окружите ими ядро:
cudaEventRecord(start, computeStream);
kernel<<<grid, block, 0, computeStream>>>(d_data);
cudaEventRecord(stop, computeStream);
cudaStreamSynchronize(computeStream);
float ms;
cudaEventElapsedTime(&ms, start, stop);Если замер показал 0.42 мс на ядро, а таких ядер в кадре двадцать, то 8.4 мс уходит на вычисления, а остальное съедают копирования и лишние синхронизации. Для уведомлений на стороне хоста есть cudaLaunchHostFunc, ставящая колбэк в очередь потока, а устаревшая cudaStreamAddCallback готовится к удалению.
Разделение одной видеокарты между процессами через MPS
Дело потоками не заканчивается: на карте живут несколько процессов. В режиме DEFAULT они делят карту тайм-слайсингом, и каждое переключение контекста стоит ресурсов. Режим EXCLUSIVE_PROCESS запрещает второму процессу открывать контекст, и очереди приходится строить на уровне планировщика задач. Оба варианта проигрывают службе MPS, Multi-Process Service.
MPS устроена как пара процессов: контрольная служба nvidia-cuda-mps-control принимает клиентов и порождает сервер nvidia-cuda-mps-server, который владеет контекстом на GPU и исполняет операции клиентов. Множество контекстов сводится в один, переключений меньше, утилизация выше. Побочный эффект: в nvidia-smi владельцем карты числится nvidia-cuda-mps-server с типом процесса M, и свои PID в списке вычислительных процессов не найти.
Запуск на многопользовательском сервере:
nvidia-smi -i 0 -c EXCLUSIVE_PROCESS # режим EXCLUSIVE_PROCESS для GPU 0
export CUDA_VISIBLE_DEVICES=0 # карта, которую обслуживает служба
nvidia-cuda-mps-control -d # контрольная служба в фоновом процессеПорядок внедрения на рабочем сервере такой:
- Проверить видимость карт и UUID командой nvidia-smi -L;
- Перевести GPU в режим EXCLUSIVE_PROCESS через nvidia-smi -i 0 -c EXCLUSIVE_PROCESS;
- Задать export CUDA_VISIBLE_DEVICES=0 в окружении контрольной службы;
- Запустить службу фоновым процессом через nvidia-cuda-mps-control -d;
- Убедиться, что у клиентов CUDA_VISIBLE_DEVICES не установлена, иначе сервер откажет в подключении.
Остановка выполняется через echo quit | nvidia-cuda-mps-control в legacy-интерфейсе MPS v2 либо через nvidia-cuda-mps-control -q в MPS v3, которая включается ключом -p 3 и поддерживает server create и server list. На однопользовательской машине служба стартует под тем же UID, что и клиенты, пути задают переменные CUDA_MPS_PIPE_DIRECTORY=/tmp/nvidia-mps и CUDA_MPS_LOG_DIRECTORY=/tmp/nvidia-log, значения клиенту нужны те же. Логи лежат в /var/log/nvidia-mps: control.log у контрольной службы, server.log у сервера.
Лимиты клиентов MPS и статическое разбиение SM
Одного включения MPS мало, есть лимиты на нескольких уровнях. Переменная CUDA_MPS_ACTIVE_THREAD_PERCENTAGE ограничивает долю потоков SM клиента: строка CUDA_MPS_ACTIVE_THREAD_PERCENTAGE=50 оставляет программе половину ресурсов. Заданная в окружении контрольной службы, она становится дефолтом для всех серверов, у клиента только сужает установленный лимит. Неравномерное разбиение между контекстами одного процесса включается переменной CUDA_MPS_ENABLE_PER_CTX_DEVICE_MULTIPROCESSOR_PARTITIONING, после чего cudaDevAttrMultiProcessorCount отражает долю конкретного контекста.
Память ограничивается похоже: строка CUDA_MPS_PINNED_DEVICE_MEM_LIMIT='0=1G,1=512MB' закрепляет за клиентом гигабайт на устройстве 0 и 512 МБ на устройстве 1. Контрольная служба задаёт дефолт командой set_default_device_pinned_mem_limit 0 3G, отдельный сервер ограничивает set_device_pinned_mem_limit с PID сервера, а клиентская переменная только ужесточает порог.
Приоритеты клиентов задаются переменной CUDA_MPS_CLIENT_PRIORITY, где 0 это уровень Normal, а 1 уровень Below Normal, и снова подсказка, а не гарантия. За число аппаратных очередей между хостом и устройством отвечает CUDA_DEVICE_MAX_CONNECTIONS: под MPS все клиенты одного сервера делят общий пул соединений, для клиентов на Volta значение по умолчанию равно 2, и завышенный запрос уменьшит число доступных клиентов.
Самое жёсткое средство дележа, статическое разбиение SM, доступно с Ampere. В legacy-интерфейсе служба стартует с ключом -S, партицию создаёт команда sm_partition add с UUID карты и числом частей, конфигурацию показывает lspart, клиенту выдаётся переменная CUDA_MPS_SM_PARTITION. В MPS v3 то же делается нагляднее:
nvidia-cuda-mps-control -d -p 3
nvidia-cuda-mps-control server create training
nvidia-cuda-mps-control sm-partition create large --server=training --device=0 --chunks=5
nvidia-cuda-mps-control sm-partition list --server=training --allВ примере из документации пять частей дали sm-count=40, по 8 SM на часть при карте в 80 SM и десяти частях. Каждая SM закреплена за партицией намертво: соседи не займут простаивающие SM, зато авария одного клиента не тянет остальных. Начиная с драйвера ветки r610, работает частичная изоляция ошибок: сбой SM атрибутируется партиции, аварийный клиент останавливается и получает CUDA_ERROR_LAUNCH_FAILED.
CUDA Graphs против накладных расходов на многократные запуски
Когда цепочка запусков повторяется тысячи раз, узким местом становится сама постановка в очередь: каждый запуск ядра это работа на стороне CPU. CUDA Graphs позволяют один раз описать граф операций, собрать его через захват потока и дальше исполнять одним вызовом.
cudaGraph_t graph;
cudaGraphExec_t graphExec;
cudaStreamBeginCapture(computeStream, cudaStreamCaptureModeGlobal);
kernelA<<<grid, block, 0, computeStream>>>(d_data);
kernelB<<<grid, block, 0, computeStream>>>(d_data);
cudaStreamEndCapture(computeStream, &graph);
cudaGraphInstantiate(&graphExec, graph, 0);Дальше в цикле достаточно cudaGraphLaunch(graphExec, computeStream) и одной синхронизации на весь пакет. Разветвлённые схемы захватываются теми же cudaStreamWaitEvent, граф зависимостей переносится без переделки логики. В свежих версиях CUDA функция cudaGraphInstantiate принимает три аргумента, в старых набор параметров был шире, при переносе сигнатуру стоит сверить.
Переменная CUDA_LAUNCH_BLOCKING=1 полезна при отладке: среда синхронизируется после каждого запуска ядра, и ошибка привязывается к конкретной операции. Производительность проседает сильно, поэтому переменную включают только на время отладки.
Мониторинг нагрузки и учёт вычислительных процессов через консольные утилиты
Вся настройка держится на цифрах от nvidia-smi: список карт выводит nvidia-smi -L, текущая загрузка снимается одной строкой:
nvidia-smi --query-gpu=index,name,utilization.gpu,utilization.memory,memory.used,power.draw --format=csv -l 1Ключ -l 1 обновляет показания раз в секунду, формат csv пригоден для скриптов. Кто держит карту сейчас, отвечает запрос --query-compute-apps:
nvidia-smi --query-compute-apps=pid,process_name,used_memory --format=csvПод MPS в списке фигурирует nvidia-cuda-mps-server, это нормально.
Для наблюдения в терминале есть dmon: nvidia-smi dmon -s pucmet -d 1 -o T печатает по строке на замер, где p это мощность и температура, u загрузка SM и памяти, c частоты, m память кадра и BAR1, e ошибки ECC, t трафик PCIe. Строка выглядит примерно так: gpu 0, 247 Вт, 62 градуса, SM 98 процентов, память 45 процентов. По процессам смотрит режим pmon: nvidia-smi pmon -d 1 -s um показывает загрузку SM и память по каждому PID.
Учёт за всё время работы включается режимом accounting: nvidia-smi --accounting-mode=1 -i 0, после чего nvidia-smi --query-accounted-apps=pid,process_name,gpu_time --format=csv отдаёт накопленное время GPU по PID, а nvidia-smi --clear-accounted-apps обнуляет счётчики.
nvidia-smi -pm 1 включает режим постоянства драйвера, не переживающий перезагрузку. Частоты на Volta и новее фиксирует nvidia-smi -lgc 1500,1500, сбрасывает -rgc, предел мощности ставит nvidia-smi -pl 250, и замеры до и после сопоставимы.
Вся связка выстраивается по нарастающей: потоки и приоритеты, замер через события, MPS с лимитами, статические партиции для изоляции, мониторинг поверх всего. Каждый шаг проверяется числами: если утилизация SM ползёт вверх, а хвосты латентности укорачиваются, планирование вычислительного конвейера настроено по-настоящему.