Мы живем в эпоху, когда ИИ стал доступен каждому. Но за магией PyTorch скрывается колоссальная инженерная работа и сложные вычислительные процессы, которые для большинства остаются черным ящиком.
Я хочу запустить большой цикл статей От MNIST к Transformer, цель которого пошагаво пройти путь от простого CUDA ядра до создания архитектуры Transformer - фундамента современных LLM моделей. Мы не будем использовать готовые высокоуровневые библиотеки. Мы будем разбирать, как все устроено под капотом, и пересобирать их ключевые механизмы своими руками на самом низком уровне. Только так можно по настоящему понять как работают LLM и что за этим стоит.
Приготовьтесь, будет много кода на C++ и CUDA, работы с памятью и погружения в архитектуру GPU. И конечно же математика что за этим стоит. Поехали!
Введение в цикл "От MNIST к Transformer"
Это будет долгий, сложный, но невероятно увлекательный путь. Мы начнем с самых азов и будем шаг за шагом усложнять задачу. Наша цель не только досконально понять как работают современные модели, но и научится программировать на GPU.
Самое начало (Мы здесь): Разберемся с основами работы GPU. Напишем наш первый код на C++ и CUDA для простейшей операции сложения векторов, hello world в мире GPGPU и CUDA, чтобы понять, как CPU и видеокарта общаются друг с другом.
Настройка окружения
Для разработки я использую WSL2 с Ubuntu и видеокарту NVIDIA RTX 4070 Ti. Эта карта базируется на архитектуре Ada Lovelace (Compute Capability = 8.9)
Compute Capability
Это номер поколения архитектуры видеокарты. Это не версия драйвера и не версия программного обеспечения CUDA Toolkit. Это идентификатор, который говорит компилятору и ОС, какие именно аппаратные возможности зашиты в чип
Установка драйверов и CUDA Toolkit
Первым делом убедимся, что система видит карту, а компилятор готов к работе. Для начала нужно установить актуальные драйвера на видиокарту и CUDA Toolkit с официального сайта NVIDIA. Инструкцию можно посмотреть например вот здесь для WSL. Не будем сильно погружаться в установку нужно просто следовать инструкции с официальных источников.
После установки нужно будет запустить в терминале для проверки установки.
nvcc --version # должна отобразится версия компилятора например: # NVIDIA (R) Cuda compiler driver # Copyright (c) 2005-2025 NVIDIA Corporation # Built on Wed_Jan_15_19:20:09_PST_2025 # Cuda compilation tools, release 12.8, V12.8.61 # Build cuda_12.8.r12.8/compiler.35404655_0
NVCC (NVIDIA CUDA Compiler) формально является полноценным компилятором C++. Его главная задача - это "распознать" в исходном коде (файлах с расширением .cu) разделение на логику для процессора (Host) и для видеокарты (Device). NVCC самостоятельно обрабатывает все специфические расширения CUDA и компилирует код кернелов в промежуточное представление (PTX) или бинарный код для GPU. Однако для сборки основной части приложения на C++ я буду использовать G++.
Немного про компиляцию
В процессе компиляции NVCC создает executable или library. И здесь есть два ключевых понятия cubin и fatbin.
cubin - это бинарный файл, содержащий исполняемый код для конкретной архитектуры GPU. Плюсы использования cubin - это быстрый запуск ядра, минусы невозможность запускать cubin на другой архитектуре.
fatbin — это своего рода контейнер, который решает проблему совместимости. Он может включать в себя одновременно и бинарный код GPU (cubin), и виртуальный код на языке архитектуры команд PTX. Код PTX может быть скомпилирован с помощью JIT-компиляции для работы на различных поколениях устройств.
Бинарный файл на выходе nvcc содержит как бинарный код для CPU, так и контейнер fatbin для GPU-кода.
GPU vs CPU
Для начала разберем ключевые отличия между CPU и GPU. В в литературе или статьях посвященных GPGPU центральный процессор (CPU) часто называют Host (хостом). Именно он выполняет основной код программы и берет на себя роль управляющего, он выстраивает пайплайны и координирует обмен данных между хостом и девайсом (GPU). GPU в этой связке выступает как мощный вычислительный узел, своего рода внешний сервер, которому мы делегируем вычеслительные задачи, чтобы затем получить готовый результат.
На схеме выше видна принципиальная разница между GPU и CPU. На CPU ядер меньше, но каждое из них это производительное и сложное устройство, способное быстро выполнять сложную логику и очень быстро переключаться между разными задачами. Так же на CPU довольна сложная и многоуровневая система система кэша (L1, L2, L3). И большой объем оперативной памяти (DRAM). На GPU же в свою очередь ядер намного больше, но GPU ядра (cores) спроектированы максимально простыми и компактными, чтобы их можно было уместить на кристалле в большом количестве. Они не обладают такой гибкостью и плохо справляются например с такими задачами как ветвление кода или переключение контекста. Так же на GPU более простой и меньший по объему кэш, его обычно два уровня L1 и L2. VRAM - это видеопамять она по сути выполняем роль оперативной памяти для CPU, но обычно ее меньше и она расположена на плате самой видеокарты. Если смотреть верхнеуровнево то GPU это такой компьютер в компьютере, со своими процессорами, материнской платой, оперативной памятью и тд. Самый же большой плюс GPU по сравнению с CPU это количество вычислительных ядер их может быть в тысячи раз больше чем на CPU, это идеально для проведения большого объема однотипных вычислений, например умножения огромных матриц, что и делает GPU очень важной частью работы современных ИИ моделей.
Архитектура GPU на уровне железа
Немного подробнее разберем архитектуру GPU на уровне железа. На GPU ядра организованы в группы, так называемые Streaming Multiprocessors или SMs. Каждый Streaming Multiprocessor включает в себя:
набор вычеслительных юнитов или тех самых ядер (core)
Локальный регистровый файл (Local Register File): Это самая быстрая память, расположенная физичес��и ближе всего к вычислительным ядрам. Она предназначена для хранения переменных и промежуточных данных с которыми ядра работают в данный момент
Унифицированный кэш данных (Unified Data Cache): Это универсальное хранилище которое объединяет в себе L1 cache и shared memory. Распределение физической памяти между L1 cache и shared memory может быть сделано в рантайме.
В GPU может быть много Streaming Multiprocessors, в самых новых архитектурах для которых Compute Capability = 9.0 и выше, Streaming Multiprocessors могут быть дополнительно объеденены в кластеры Graphics Processing Clusters (GPC), на моей домашней 4070 такого нет так как у нее Compute Capability = 8.9.
VRAM - это аналог оперативной памяти для CPU, это самый большой объем памяти доступный для GPU но и самый медленный, по сравнению с L1, L2 кэшем, Shared Memory и регистрами.
Логическая архитектура CUDA
Мы поговорили про архитектуру GPU на уровне железа, пришло время поговорить немного как устроена логическая архитектура на уровне софта.
В терминах CUDA одно физическое ядро соответствует одному вычислительному потоку (thread), потоки объеденены в логические блоки thread blocks, они в свою очередь обеднены в одну большую сеть (grid)
Один thread block должен выполняться и помещаться на одном Streaming Multiprocessor, далее буду называть просто SM. Один SM может обслуживать одновременно один или несколько блоков в зависимости от размера блока. Бывают ситуации когда состояние выполнения блока может быть сохранено в глобальную память vram и запущено потом на другом уже SM-e, это чем то похоже на переключение контекстов на CPU, но это довольно редкая техника работы с блоками и SM. Размер grid может быть огромный и содержать миллионы блоков, в свою очередь количество SM на видеокарте ограниченно и может быть меньше, в таком случае блоки распределяются на выполнения на доступные SM и когда какой либо SM освободиться от выполнения одного блока он может взять другой блок на выполнение. Порядок выполнения блоков не гарантирован, поэтому один блок не должен полагаться на результат выполнения другого блока. Я уже говорил что есть еще одна структура организации данных кластеры, но пока мы опустим ее.
Так же внутри блока потоки организованы в так называемые варпы или warps, один warp включает в себя 32 потока и орагнизованы по принципу SIMT (Single-Instruction Multiple-Threads) парадигмы. Что это такое и почему это важно поговорим в следующей статье.
Характеристики видеокарты
Характеристики видеокарты можно посмотреть вызвав вот такой код
void device_info() { int deviceCount; cudaGetDeviceCount(&deviceCount); for (int i = 0; i < deviceCount; i++) { cudaDeviceProp prop; cudaGetDeviceProperties(&prop, i); std::cout << "Device " << i << ": " << prop.name << std::endl; std::cout << " Number of SM (Streaming Multiprocessors): " << prop.multiProcessorCount << std::endl; std::cout << " Compute Capability: " << prop.major << "." << prop.minor << std::endl; std::cout << " Total global memory: " << prop.totalGlobalMem / (1024 * 1024) << " MB" << std::endl; std::cout << " Shared memory per block: " << prop.sharedMemPerBlock / 1024 << " KB" << std::endl; std::cout << " Registers per block: " << prop.regsPerBlock << std::endl; std::cout << " Warp size: " << prop.warpSize << std::endl; std::cout << " Maximum number of threads per block: " << prop.maxThreadsPerBlock << std::endl; std::cout << " Maximum number of threads per multiprocessor: " << prop.maxThreadsPerMultiProcessor << std::endl; } }
Device 0: NVIDIA GeForce RTX 4070 Ti SUPER Number of SM (Streaming Multiprocessors): 66 Compute Capability: 8.9 Total global memory: 16375 MB Shared memory per block: 48 KB Registers per block: 65536 Warp size: 32 Maximum number of threads per block: 1024 Maximum number of threads per multiprocessor: 1536 Maximum number of threads per block: 1024
Итого с текущими параметрами мы можем запустить 1 блок с размером 1024 ядра на одном SM или 6 блоков по 256 ядер в каждом на одном SM, что даст меньший размер блока но 100% нагрузку на SM.
Сложение векторов и написание первого ядра
Вектор и сложение векторов
Давайте немного отвлечемся от железа и кода и вернемся к математики. Что такое вектор? Вектор это набор чисел в математических формулах его можно записать как 
и он состоит из элементов 
. Сложение двух векторов и 
это простейшая операция где каждый элемент из вектора складывается с соответствующим элементом из вектора ![C = [a_1 + b_1, a_2 +b_2, ..., a_i+b_i]](https://habrastorage.org/getpro/habr/formulas/f/f2/f2b/f2bab39272816ce20d802a0be5bb375c.svg)
размер векторов и должен совпадать.
На CPU мы бы просто в цикле прошлись по элементам двух вектором и сложили их. Но так как элементов может быть много то и цикл будет очень большой, а сложение каждого элемента вектора это независимая простейшая операция, то это очень просто перекладывается на GPU.
Написание ядра
Теперь, когда окружение настроено, теория изучена можно перейти к написанию самого ядра (kernel) который будет выполняться непосредственно на видеокарте. В CUDA для этого используется спецификатор __global__
__global__ void vectorAdd(const float* A, const float* B, float* C, int numElements) { // Получаем индекс элемента в векторе int i = blockIdx.x * blockDim.x + threadIdx.x; // Делаем проверку границ что мы не вышли за размер вектора if (i < numElements) { // Складываем элементы C[i] = A[i] + B[i]; } }
Данная функция выполнится не зависимо на каждом ядре видеокарты. Например если у нас размер вектора 1000 элементов, мы сделали размер блока в 256 потоков итого у нас будет выполнено 4 блока что вполне влезает в один SM. Все эти блоки и потоки будут выполнены одновременно и нам уже не нужен ни какой цикл как на CPU.
СPU часть и взаимодействие с GPU
Для того что бы запустить код кернела на видеокарте нам нужно написать небольшую обвязку. Синтаксис вызова кернела выглядит так kernel_name<<<blocksPerGrid, threadsPerBlock>>>()
void launchVectorAdd(const float* d_A, const float* d_B, float* d_C, int numElements, int threadsPerBlock) { // скоролько блоков на сетку (grid) // например для 1000 элементов и размера блока 256 потоков // у нас получается что нам нужно запустить 4 блока int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; // Запускаем само ядро vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements); }
Теперь давайте разберем основную CPU часть
int numElements = 1000; size_t size = numElements * sizeof(float); // 1. Выделяем память на стороне хоста в оперативной памяти float *h_A = (float *)malloc(size); float *h_B = (float *)malloc(size); float *h_C = (float *)malloc(size); // Заполняем вектора A и B начальными значениями (можно добавить рандомные значения) std::cout << "Initializing data..." << std::endl; for (int i = 0; i < numElements; i++) { h_A[i] = 1.0f; h_B[i] = 2.0f; } // 2. Выделяем память для этих же трех векторов но уже на стороне GPU в VRAM float *d_A, *d_B, *d_C; cudaMalloc((void **)&d_A, size); cudaMalloc((void **)&d_B, size); cudaMalloc((void **)&d_C, size); // 3. Для того что бы работать с данными на GPU нам нужно явно скопировать данные // из оперативной памяти в VRAM std::cout << "Copying data to GPU..." << std::endl; cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); // 4. Тут запускаем сам kernel std::cout << "Running CUDA kernel..." << std::endl; int threadsPerBlock = 256; launchVectorAdd(d_A, d_B, d_C, numElements, threadsPerBlock); // Ждем выполнения кода на GPU cudaDeviceSynchronize(); // Проверяем на ошибки GPU auto err = cudaGetLastError(); if (err != cudaSuccess) { std::cerr << "Kernel launch error: " << cudaGetErrorString(err) << std::endl; return; } // 5. Так как результат работы сохранен в VRAM нам нужно перенести обратно // в память хоста (оперативную память компютера) std::cout << "Copying results back to CPU..." << std::endl; cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); // 6. Так как значения элементов у нас в данном примере не рандомные, то мы явно проверим // что сложение выполнено правильно. std::cout << "Verifying results..." << std::endl; bool success = true; for (int i = 0; i < numElements; i++) { if (h_C[i] != 3.0f) { std::cerr << "Error at index " << i << ": " << h_C[i] << " != 3.0" << std::endl; success = false; break; } } if (success) { std::cout << "Success! " << numElements << " vector additions completed correctly." << std::endl; } // 7. Память выделенную с помощью cudaMalloc нужно явно освободить с помощью вызова cudaFree cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); free(h_A); free(h_B); free(h_C);
Отдельно напишу про синхронизацию работы CPU и GPU. Вызов функции vectorAdd<<<blocksPerGrid, threadsPerBlock>>>() является асинхронным и возвращает управление основному коду программы сразу после вызова, мы можем относиться к этом как к запросу на сервер. Для того что бы подождать выполнения работы на GPU мы явно должны вызвать cudaDeviceSynchronize(); существуют разные техники синхронизации и вызов cudaDeviceSynchronize(); самый простой, но в данной статье остановимся пока на этом. В других статьях отдельно разберем другие способы синхронизации.
Сравнение с выполнением на СPU
Теперь давайте проведем небольшой бэнчмарк по сравнению запуском этого же алгоритма на CPU.
Для вектора размером 1000 элементов:
CPU time: 0.000889 ms GPU time (with transfers): 0.244128 ms GPU time (kernel only): 0.162944 ms Speedup (with transfers): 0.00364153x Speedup (kernel only): 0.00545586x
Ой, что мы видим, CPU намного быстрее. Все было зря? Выкидываем GPU? Или все таки нет?
А давайте вспомним что такое LLM, это же большая языковая модель и ключевое слово здесь большая. И количество параметров там не измеряется тысячами и даже миллионами, количество параметров там измеряется миллиардами. Так давайте же чуть чуть увеличим наш вектор и проведем сравнение в таком случае.
Давайте возьмем вектор например 50 млн элементов. И посмотрим на результат.
CPU time: 114.603 ms GPU time (with transfers): 108.736 ms GPU time (kernel only): 1.15523 ms Speedup (with transfers): 1.05395x Speedup (kernel only): 99.2035x
Так так так, тут уже картина совершенно другая, мы видимо что чисто вычислительная часть получилась в 99 раз быстрее чем на CPU. Очень много времени заняло как мы видим трансфер памяти от CPU к GPU и обратно, но даже не смотря на это весь процесс занял меньше времени чем цикл на CPU. Исходя из этого мы видим что при больших объемах данных вычисления на GPU происходят намного быстрее. Но при этом если размер вектора маленький нужно сто раз подумать нужно ли отправлять такой объем работы на GPU.
Заключение
Данная статья является начальной статьей цикла про программирование на CUDA и создание и разбора архитектуры transformer с нуля. Здесь мы разобрались как устроено железо видеокарты, как устроено логическое деление CUDA на потоки, блоки и grid. Так же написали наше первое ядро для сложения векторов и провели бенчмарк.
Подробный код для статьи можно посмотреть на github.
Так же если хотите следить и данная тема вам интересна можете подписаться на канал в телеграмме созданный специально для этого цикла статей, там будут анонсы статьей по циклу, обсуждения, ваши запросы и многое другое.
Продолжение следует ...
