Вы когда-нибудь задумывались, как код, написанный на CUDA, превращается в молниеносные вычисления на GPU? Для большинства разработчиков запуск CUDA-ядра выглядит как магия: пишем функцию с приставкой __global__, вызываем её с помощью загадочного синтаксиса с тройными угловыми скобками <<<grid, block>>>, и вуаля — вычисления, которые на CPU занимали минуты, завершаются за доли секунды. Но что происходит на самом деле?
В этой статье мы детально разберем весь жизненный цикл CUDA-ядра: от компиляции и работы драйвера до планирования варпов на физических мультипроцессорах и оптимизации доступа к памяти.
1. Этап подготовки: Компиляция и создание контекста
Прежде чем видеокарта сможет выполнить хотя бы одну инструкцию, исходный код должен пройти сложный процесс трансформации. В отличие от стандартного C++ кода, компилятор NVCC (Nvidia CUDA Compiler) разделяет программу на две части: хост-код (выполняемый на CPU) и девайс-код (выполняемый на GPU).
От исходного кода к PTX и SASS
Когда вы компилируете CUDA-программу, девайс-код проходит через несколько стадий:
- PTX (Parallel Thread Execution): Это низкоуровневый виртуальный ассемблер, независимый от конкретной архитектуры GPU (например, Kepler, Pascal, Ampere или Hopper). PTX обеспечивает переносимость кода между поколениями видеокарт.
- SASS (Source Assembly): Это реальный машинный код, оптимизированный под конкретную микроархитектуру GPU. Перевод из PTX в SASS может происходить как на этапе компиляции (offline compilation), так и непосредственно во время выполнения программы с помощью JIT-компилятора, встроенного в CUDA Driver.
Инициализация CUDA Context
При первом вызове любой функции CUDA Runtime API (например, cudaMalloc) происходит неявная инициализация CUDA Context. Контекст — это аналог процесса в операционной системе CPU. Он изолирует ресурсы: выделенную память, стек вызовов, таблицы дескрипторов и очереди команд. Контекст привязывается к конкретному физическому устройству и управляет его жизненным циклом.
// Типичный пример подготовки данных на хостеfloat *h_data, *d_data;int size = 1024 * sizeof(float);h_data = (float*)malloc(size);// Инициализация контекста и выделение памяти на GPUcudaMalloc(&d_data, size); cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
На этом этапе драйвер выделяет виртуальное адресное пространство на GPU и связывает его с таблицей страниц памяти хоста, подготавливая физические каналы DMA (Direct Memory Access) для быстрой передачи данных по шине PCIe.
2. Анатомия вызова ядра: От хоста к командному буферу
Когда выполнение программы доходит до строки запуска ядра, например myKernel<<<blocks, threads>>>(d_data);, CPU не начинает выполнять эти вычисления сам. Вместо этого он формирует команду для отправки на GPU. Этот процесс асинхронен по своей природе.
Различие в философии вычислений
Когда мы проектируем высоконагруженные системы, мы всегда выбираем между оптимизацией под задержку (latency) и пропускную способность (throughput). Это фундаментальное различие можно увидеть даже в повседневных мобильных приложениях: если сравнить условные apple music vs whatsapp, то первое требует стабильного, широкого канала для непрерывного стриминга медиаданных (аналог GPU), тогда как второе ориентировано на мгновенную доставку коротких сообщений с минимальным пингом (аналог CPU). Архитектура GPU изначально создавалась для «тяжелого стриминга» вычислений, где задержка запуска отдельной нити не играет роли, если общая пропускная способность остается колоссальной.