CUDA с нуля: пишем ядро для сложения векторов как первый шаг к пониманию Transformer

Зачем писать сложение векторов на CUDA в 2026 году?

Кажется глупым. Абсурдно. В PyTorch есть torch.add(), в NumPy - numpy.add(). Одна строчка кода. Зачем тратить время на низкоуровневый код?

Потому что все современные LLM - от GPT-4.5 до свежих open-source моделей 2026 года - работают на этих примитивах. Когда вы видите "CUDA out of memory" в трансформере на 70 миллиардов параметров, вы не просто увеличиваете batch size. Вы понимаете, почему память закончилась. Какие тензоры где лежат. Как данные движутся между CPU и GPU.

Сложение векторов - это "Hello, World!" мира GPU. Но не тот скучный "Hello, World!", который пишут и забывают. Это фундамент, на котором строится все: матричные умножения, attention механизмы, нормализация слоев.

Архитектура GPU: почему 10000 потоков работают быстрее одного?

CPU думает линейно. Одно ядро - одна задача. GPU думает параллельно. Тысячи ядер - тысячи задач одновременно.

Но не все так просто. GPU - не просто много CPU. Это другая архитектура с другой философией:

Почему такая сложная иерархия? Потому что потоки внутри одного блока могут общаться через быструю shared memory. Потоки из разных блоков - только через медленную global memory. Это ключевой момент для оптимизации.

Память CUDA: иерархия скоростей и ограничений

Самая частая ошибка новичков: думать, что память GPU - это просто быстрая RAM. На самом деле это сложная иерархия с разными скоростями доступа:

• Global memory - медленная, но много (до 48GB на RTX 5090). Сюда загружаются данные с CPU
• Shared memory - быстрая, но мало (до 256KB на блок). Общая для потоков внутри блока
• Registers - сверхбыстрая, индивидуальная для каждого потока
• Constant memory - кэшируемая read-only память для констант

Разница в скорости? Shared memory в 100 раз быстрее global memory. Registers еще быстрее. Но объемы... Объемы смешные по сравнению с global.

Вот почему оптимизация памяти важнее оптимизации вычислений. Можно написать идеальный алгоритм, но если он постоянно ходит в global memory - будет медленнее плохого алгоритма, который использует shared.

Подготовка: что нужно установить в 2026 году

CUDA Toolkit 12.5+ (на февраль 2026 это актуальная версия). NVIDIA не стоит на месте - каждый год новые фичи, новые оптимизации. Старые гайды 2023 года уже не работают с новыми картами.

# Проверяем установку CUDA
nvcc --version
# Должно быть что-то вроде:
# nvcc: NVIDIA (R) Cuda compiler driver
# Copyright (c) 2005-2025 NVIDIA Corporation
# Built on Wed_Dec__4_16:38:05_PST_2024
# Cuda compilation tools, release 12.5, V12.5.40

1 Пишем хост-код: подготовка данных на CPU

Сначала создаем массивы на CPU, копируем их в GPU, запускаем ядро, копируем результат обратно. Стандартный паттерн.

#include <iostream>
#include <vector>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

// Прототип нашего ядра (объявим позже)
__global__ void vectorAdd(const float* A, const float* B, float* C, int n);

int main() {
    const int N = 1000000;  // 1 миллион элементов
    size_t size = N * sizeof(float);
    
    // Создаем данные на хосте
    std::vector<float> h_A(N, 1.0f);  // Все элементы = 1.0
    std::vector<float> h_B(N, 2.0f);  // Все элементы = 2.0
    std::vector<float> h_C(N, 0.0f);  // Результат
    
    // Указатели на устройстве (GPU)
    float *d_A, *d_B, *d_C;
    
    // Выделяем память на GPU
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);
    
    // Копируем данные с CPU на GPU
    cudaMemcpy(d_A, h_A.data(), size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B.data(), size, cudaMemcpyHostToDevice);
    
    // Запускаем ядро
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // Копируем результат обратно на CPU
    cudaMemcpy(h_C.data(), d_C, size, cudaMemcpyDeviceToHost);
    
    // Проверяем результат
    bool correct = true;
    for (int i = 0; i < N; i++) {
        if (fabs(h_C[i] - 3.0f) > 1e-6) {  // 1.0 + 2.0 = 3.0
            correct = false;
            break;
        }
    }
    
    std::cout << "Result: " << (correct ? "CORRECT" : "WRONG") << std::endl;
    
    // Освобождаем память
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    
    return 0;
}

Пока ничего сложного. Просто выделили память, скопировали, запустили какое-то ядро vectorAdd. Но что внутри этого ядра? Вот где начинается магия.

2 Пишем само ядро: как 1000000 потоков работают вместе

Вот самая важная часть - само CUDA ядро. Всего 5 строк кода, но в них вся философия GPU программирования.

__global__ void vectorAdd(const float* A, const float* B, float* C, int n) {
    // Вычисляем глобальный индекс для этого потока
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    
    // Проверяем границы массива
    if (i < n) {
        // Собственно сложение
        C[i] = A[i] + B[i];
    }
}

Давайте разберем каждую строчку:

• __global__ - спецификатор CUDA. Говорит компилятору, что это ядро, которое запускается на GPU и вызывается с CPU
• blockIdx.x - индекс блока в grid
• blockDim.x - размер блока (сколько потоков в блоке)
• threadIdx.x - индекс потока внутри блока

Формула i = blockIdx.x * blockDim.x + threadIdx.x - это сердце CUDA. Она гарантирует, что каждый поток обрабатывает уникальный элемент массива.

Если у нас 1000000 элементов, 256 потоков в блоке, то нужно 3907 блоков (1000000 / 256, округлено вверх). Каждый блок получает свой blockIdx.x, каждый поток в блоке - свой threadIdx.x.

Компиляция и запуск: почему nvcc такой странный?

Компилируем наш код:

nvcc -o vector_add vector_add.cu
./vector_add

Если все сделано правильно - увидите "Result: CORRECT". Если нет... Добро пожаловать в мир CUDA debugging.

nvcc - это не просто компилятор C++. Это двухпроходной компилятор, который:

1. Компилирует host код (CPU часть) обычным C++ компилятором
2. Компилирует device код (GPU часть) в PTX (Parallel Thread Execution) ассемблер
3. Собирает все вместе в один исполняемый файл

PTX - это промежуточный байт-код, который JIT-компилируется в машинный код для конкретной GPU при запуске. Вот почему один бинарник может работать на разных GPU (но с разной производительностью).

Ошибки, которые совершают все (и вы тоже совершите)

Ошибка 1: Забыть проверку границ

// ПЛОХО: segmentation fault гарантирован
__global__ void badVectorAdd(const float* A, const float* B, float* C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    C[i] = A[i] + B[i];  // Без проверки i < n
}

Результат? Либо мусор в памяти, либо segfault, либо тихая порча данных. GPU не проверяет границы массивов как CPU. Вышли за пределы - пишите куда попало.

Ошибка 2: Неправильный расчет блоков

// ПЛОХО: часть элементов не обработается
int blocksPerGrid = N / threadsPerBlock;  // Деление без округления вверх
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

Если N=1000000, threadsPerBlock=256, то 1000000/256=3906.25. Целочисленное деление даст 3906. 3906*256=999936. Последние 64 элемента никогда не обработаются.

Правильно так:

int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

Ошибка 3: Игнорировать ошибки CUDA

CUDA функции возвращают cudaError_t. Игнорировать его - преступление.

// ПЛОХО: ошибка останется незамеченной
cudaMalloc(&d_A, size);

// ХОРОШО: проверяем ошибки
cudaError_t err = cudaMalloc(&d_A, size);
if (err != cudaSuccess) {
    std::cerr << "CUDA error: " << cudaGetErrorString(err) << std::endl;
    return 1;
}

Или используйте макросы-обертки. В продакшене без проверок ошибок CUDA жить нельзя.

А что насчет производительности? Это же GPU!

Запустите бенчмарк. Сравните с CPU. Удивитесь.

Для маленьких массивов (меньше 1000 элементов) CPU будет быстрее. Накладные расходы на копирование данных в GPU и обратно съедят всю выгоду.

Для больших массивов (миллионы элементов) GPU выигрывает в разы. Но только если правильно настроить размеры блоков.

Почему размер блока важен? Потоки внутри блока выполняются группами по 32 (warps). Если в блоке 256 потоков - это 8 warps. GPU переключается между warps, когда один ждет данных из памяти. Больше потоков в блоке - лучше скрытие задержек памяти.

Но! Большие блоки занимают больше регистров и shared memory. Может не хватить ресурсов.

Эмпирическое правило на 2026 год: 128, 256 или 512 потоков на блок. Для простых операций вроде сложения - 256 обычно оптимально.

Как это связано с Transformer?

Кажется, что сложение векторов - это слишком просто. Где же матричные умножения? Где attention? Где softmax?

А они все состоят из таких же примитивов. Только сложнее.

• Matrix multiplication (GEMM) - это множество скалярных произведений, каждое из которых - сумма произведений
• Attention - это матричные умножения плюс softmax
• Softmax - это экспоненты, суммы, деления
• Layer normalization - это средние, дисперсии, масштабирование

Каждый из этих операторов - это CUDA ядро. Только больше строк кода. Больше оптимизаций. Больше работы с памятью.

Когда вы пишете transformer_layer.forward() в PyTorch, под капотом запускаются десятки таких ядер. Тысячи потоков. Миллионы операций.

Наша следующая цель - написать матричное умножение на CUDA. Потом attention. Потом весь Transformer.

Но начинаем с малого. С сложения векторов. Потому что без этого фундамента не построить небоскреб.

Что дальше? План погружения в CUDA

1. Матричное умножение (GEMM) - основа всех нейросетей
2. Оптимизация через shared memory и регистры
3. Atomic операции для reduction (сумма, максимум)
4. Attention механизм с mask и softmax
5. Layer normalization и residual connections
6. Сборка всего вместе в Transformer блок

Каждый шаг - это новая статья. Новые сложности. Новые оптимизации.

Но если вы освоили сложение векторов - вы уже понимаете главное: как думает GPU. Как распределять работу между потоками. Как работать с памятью.

Остальное - технические детали.

P.S. Если кажется, что это слишком сложно - так и есть. CUDA не прощает ошибок. Но когда ваш код наконец заработает в 100 раз быстрее CPU версии... Эта магия стоит потраченного времени.

P.P.S. В следующей статье разберем, почему трансформеры на стероидах используют те же принципы, но с дополнительными оптимизациями. И как новые GPU RTX 5000 серии меняют правила игры.