Путь задачи от строчки кода до вычислительных ядер видеокарты

MAATRIX

«Просто перенесите вычисления на GPU» — фраза, которая на практике разворачивается в цепочку из нескольких разных этапов: компилятора, шины передачи данных, драйвера и планировщика. Когда обучение модели или рендер тормозит, а видеокарта в nvidia-smi показывает низкую загрузку, причина почти всегда прячется в одном из этих промежуточных звеньев. Разберём весь путь по шагам — от строчки кода до реального выполнения на кристалле.

Код, который выполняется в двух мирах одновременно

Программа, использующая GPU, всегда состоит из двух частей на разном железе. Есть хост-код — обычный C++, Python или Rust, выполняющийся на CPU: он открывает файлы, читает аргументы, готовит данные, вызывает библиотеки. И есть код ядра (kernel) — функция, помеченная специальным образом, которая должна выполниться не на процессоре, а на графическом чипе.

В CUDA это выглядит так:

__global__ void add_vectors(float* a, float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

Ключевое слово __global__ — это не просто модификатор доступа. Оно говорит компилятору: эта функция вызывается с CPU, но выполняется на GPU, причём не один раз, а параллельно, множеством потоков одновременно. Вызов в хост-коде выглядит как обычная функция с дополнительным синтаксисом:

add_vectors<<<blocks, threads_per_block>>>(d_a, d_b, d_c, n);

Эта строчка — не вызов функции в привычном смысле, а постановка задачи в очередь на выполнение где-то в другом месте. Именно с этого начинается путь через компилятор, шину, драйвер и планировщик, прежде чем хоть один такт реально отработает на видеокарте. В экосистеме AMD аналогичную роль играет ROCm/HIP, у Intel — oneAPI, но принципиальная механика одна: разделение на хост- и device-код, отдельная компиляция, отдельная память, явная передача управления.

Важно сразу отделить это от классической многопоточности на CPU. Поток на процессоре — независимая единица выполнения с собственным стеком и полным набором регистров, их разумный предел — десятки или сотни. Поток на GPU (в терминологии CUDA — thread внутри warp) — гораздо более лёгкая сущность: тысячи таких потоков выполняются группами, разделяя логику исполнения. Именно поэтому нельзя просто «взять цикл for и отправить на GPU» — код ядра пишется с оглядкой на то, что каждая итерация станет отдельным потоком, а не последовательным шагом.

Компиляция ядра: от исходника до инструкций конкретного чипа

Компилятор для GPU-кода устроен не так, как для обычного CPU-кода. Например, nvcc в CUDA на самом деле — это драйвер компиляции, который разделяет исходный файл на две части и прогоняет их через разные бэкенды.

Хост-часть компилируется обычным способом — в машинный код x86-64 или ARM. Код ядра проходит два этапа: компиляция в PTX (Parallel Thread Execution) — промежуточное ассемблероподобное представление, независимое от конкретной модели видеокарты, архитектурный аналог байткода; затем компиляция PTX в SASS (Shader Assembly) — машинный код, специфичный для конкретной архитектуры GPU и реально выполняющийся на исполнительных блоках чипа.

Посмотреть промежуточный код можно напрямую:

nvcc -ptx kernel.cu -o kernel.ptx
nvcc -arch=sm_86 -cubin kernel.cu -o kernel.cubin
cuobjdump --dump-sass kernel.cubin | less

Здесь sm_86 — обозначение вычислительной возможности (compute capability), привязанное к поколению архитектуры. Собранный под одно поколение SASS-код может не запуститься или работать не оптимально на другом — отсюда и полезность PTX как промежуточного слоя: драйвер способен на лету дотранслировать PTX в SASS нужной архитектуры (JIT-компиляция), если готового бинарника под конкретный чип нет. Это одна из причин, по которой первый запуск программы после обновления драйвера иногда занимает заметно больше времени — идёт незаметная для пользователя JIT-компиляция ядер.

Итог этапа: то, что в коде выглядело как одна функция add_vectors, к моменту запуска уже превратилось в бинарный набор инструкций, привязанный к конкретной архитектуре исполнительных блоков видеокарты. Это первая точка, где абстракция «просто запусти на GPU» перестаёт быть простой.

Шина PCIe: данные и код должны физически попасть на карту

Видеокарта — это отдельное устройство со своей памятью (VRAM), физически не совпадающей с оперативной памятью хоста. Прежде чем ядро сможет что-то посчитать, нужные ему данные должны оказаться в памяти устройства, а сам скомпилированный код — быть загружен в GPU.

Типичная последовательность в CUDA:

cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

cudaMemcpy с флагом HostToDevice — не абстрактное копирование, а реальная передача байт через шину PCIe (на некоторых платформах есть более быстрый межсоединитель вроде NVLink между самими GPU, но связь CPU-GPU в подавляющем большинстве серверов идёт именно через PCIe). Эта шина устроена иначе, чем внутренняя память видеокарты: она рассчитана на совместимость и универсальность, а не на максимальную полосу пропускания, и для операций с интенсивным обменом данными часто оказывается самым узким местом всего конвейера — заметно медленнее, чем последующая работа GPU с уже загруженными данными в собственной VRAM.

Отсюда практическое следствие, которое легко упустить: если вычисление лёгкое, а объём передаваемых данных большой, время на копирование через PCIe может превысить время самого вычисления на порядок. Программа в этом случае не «ускоряется от GPU» — она просто тратит время на перекладывание данных туда и обратно, и именно эта передача данных, а не сам расчёт на исполнительных блоках, часто определяет итоговую скорость всего пайплайна.

Часть современных фреймворков и драйверов умеет скрывать эту задержку — например, через unified memory (cudaMallocManaged), где страницы переносятся между хостом и устройством по требованию, или через пины памяти хоста (cudaHostAlloc) для более быстрого DMA-переноса. Но это оптимизации поверх той же физической реальности: данные всё равно должны пересечь шину, просто более эффективным способом или в фоне, параллельно с вычислениями через отдельные CUDA-стримы.

Драйвер GPU: очередь команд вместо немедленного выполнения

Когда хост-код вызывает add_vectors<<<...>>>(...), ядро не начинает выполняться в этот же момент. Вызов асинхронный: он кладёт команду в очередь (command queue, в терминологии CUDA — stream) и немедленно возвращает управление CPU-коду.

Драйвер видеокарты — программный слой между пользовательским API (CUDA Runtime, Driver API, Vulkan, OpenCL) и физическим железом. Он отвечает за преобразование высокоуровневых вызовов в команды конкретной ревизии чипа, за управление контекстом выполнения (состоянием, которое GPU сохраняет между переключениями между процессами), за постановку команд в аппаратные очереди и их упорядочивание, за обработку прерываний и сигналов о завершении задач.

Проверить, что видит и как настроен драйвер, можно стандартными утилитами:

nvidia-smi -q | grep -A3 "Driver Version"
nvidia-smi --query-gpu=name,driver_version,memory.used,memory.total --format=csv

Асинхронность здесь не побочный эффект, а осознанная часть модели программирования: пока GPU выполняет одно ядро, CPU может готовить следующую партию данных или запускать другое ядро в параллельном стриме. Синхронизация — момент, когда хост-код явно ждёт завершения — вызывается отдельно:

cudaDeviceSynchronize();

Без такого вызова (или без синхронного cudaMemcpy без суффикса Async) программа на CPU может дойти до чтения результата раньше, чем GPU реально его посчитал — типичная причина трудноуловимых багов при портировании кода под GPU: логическая гонка между хостом и устройством, а не ошибка в самом вычислении.

Отдельный источник путаницы — то, что в один момент времени с GPU физически может работать несколько процессов (несколько контекстов), а драйвер должен разделять между ними и память, и время исполнительных блоков; это разделение не бесплатно, и именно из-за него один процесс на общей видеокарте может внезапно замедлить работу другого.

Планировщик GPU: как задача превращается в параллельную работу тысяч ядер

Здесь начинается то, ради чего всё затевалось. Когда команда дошла до аппаратного планировщика GPU, он должен раскидать миллионы логических потоков ядра по реальным физическим исполнительным блокам чипа. Архитектура современного GPU строится вокруг иерархии, а не однородного массива ядер:

• Потоковые мультипроцессоры (Streaming Multiprocessors, SM у NVIDIA; Compute Units у AMD) — крупные блоки чипа с собственными вычислительными ядрами, регистровым файлом, разделяемой памятью и планировщиком warp'ов. Именно SM, а не отдельное «ядро», — минимальная единица, которой GPU-планировщик реально назначает работу.
• Warp (у NVIDIA) или wavefront (у AMD) — группа потоков фиксированного размера, выполняющаяся в лок-шаге: все потоки warp'а исполняют одну и ту же инструкцию одновременно, но каждый — над своими данными (модель SIMT).
• Блоки потоков (thread blocks) — логическая единица из кода ядра, которую планировщик целиком назначает на один SM и не может произвольно разбить между разными SM.

Планировщик решает несколько задач сразу: на какой SM отправить очередной блок потоков; в каком порядке чередовать выполнение разных warp'ов на одном SM, чтобы скрыть задержки обращения к памяти (пока один warp ждёт данные, SM переключается на другой — это называется latency hiding, одна из причин, почему GPU эффективен именно при высокой степени параллелизма); как распределить ограниченные регистры и разделяемую память SM между одновременно резидентными блоками.

Именно здесь выясняется, почему число ядер в характеристиках видеокарты — не то же самое, что реальное ускорение задачи. Если ядро использует ветвления, из-за которых потоки одного warp'а расходятся по разным путям выполнения (divergent branching), часть исполнительных единиц простаивает, пока другие досчитывают свою ветку. Если ядру не хватает параллельной работы, чтобы заполнить все SM, часть чипа простаивает. Общая природа этого параллелизма разобрана в материале почему GPU быстрее CPU для нейросетей.

Возврат результата: копирование обратно и синхронизация

Когда все блоки потоков ядра завершили выполнение, результат лежит в памяти устройства — но хост-программа об этом ничего не знает, пока явно не спросит. Обратный путь данных зеркалит прямой:

cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);

Копирование DeviceToHost — снова операция через шину PCIe, с той же физической стоимостью, что и на входе. Если приложение гоняет данные туда-обратно на каждой мелкой операции (характерная ошибка при наивном портировании числодробительного кода), накладные расходы на передачу могут полностью съесть выигрыш от параллельных вычислений.

Есть и более тонкий момент: синхронный cudaMemcpy сам по себе выступает точкой синхронизации — вызывающий код не продолжит выполнение, пока копирование не завершится, а значит неявно дождётся завершения всех предыдущих операций в том же стриме. Каждая такая точка — гарантированный простой CPU в ожидании GPU или наоборот.

В реальных нагрузках — обучении моделей, инференсе, рендере — этот цикл (загрузка данных → запуск ядра → выгрузка результата) повторяется миллионы раз за сеанс, и поэтому фреймворки вроде PyTorch стараются держать данные в памяти устройства как можно дольше, минимизируя число пересечений границы CPU–GPU. Если после загрузки модели видеокарта в мониторинге показывает низкую активность, стоит проверить именно эту часть конвейера — материал почему модель загрузилась, а GPU простаивает разбирает типичные причины такого простоя.

Что может пойти не так на каждом из этапов

Каждый шаг пути — отдельная точка деградации, и стоит понимать, куда смотреть при диагностике.

Отдельно стоит упомянуть ситуацию, когда сам драйвер перестаёт корректно видеть карту после автоматического обновления — механика похожа на цепочку «драйвер → контекст → очередь», просто ломается на самом первом звене; разбор такого случая есть в материале как обновившийся драйвер перестаёт видеть видеокарту, хотя CUDA Toolkit не менялся.

Эта многоступенчатая цепочка одинаково присутствует что на локальной рабочей станции, что на арендованном сервере, но на выделенном GPU-сервере вы получаете предсказуемую полосу PCIe и полный контроль над версией драйвера, без соседей по железу, конкурирующих за очередь планировщика. Конфигурации под обучение моделей разобраны в материале про GPU-сервер в Великобритании для обучения моделей.

Частые вопросы