Как написать Transformer с нуля на CUDA: руководство по низкоуровневому программированию GPU

Зачем писать Transformer на CUDA в 2026 году?

PyTorch, TensorFlow, JAX - все эти фреймворки отлично работают. Зачем тогда лезть в дебри CUDA и писать все с нуля? Ответ прост: контроль и понимание. Когда вы видите ошибку "CUDA illegal instruction в llama.cpp на RTX 5060 Ti", вы не просто перезапускаете код - вы понимаете, почему новая архитектура GPU ломает старые оптимизации. Когда читаете про трансформеры на стероидах, вы не просто восхищаетесь - вы знаете, как они устроены изнутри.

Что мы будем писать?

Минимальный Transformer для классификации текста. Не гигантскую LLM на 70 миллиардов параметров (хотя принципы те же), а компактную модель, которую можно запустить на одной GPU. Архитектура: 4 слоя, 8 голов внимания, скрытый размер 512.

1 Подготовка окружения

Первое, что нужно понять: CUDA - это не библиотека, это платформа. Вы пишете код на C++ с расширениями, который компилируется в машинный код для GPU. Для работы потребуется:

• NVIDIA GPU с поддержкой Compute Capability 7.0+ (RTX 2000 серии и новее)
• CUDA Toolkit 12.5+ (на февраль 2026 это актуальная версия)
• Компилятор nvcc
• Базовое понимание C++ (указатели, шаблоны, RAII)

# Проверяем установку 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

2 Базовая структура памяти GPU

Самая частая ошибка новичков: думать, что память GPU работает как RAM. Это не так. Есть иерархия:

Вот как выглядит базовый класс для работы с памятью:

// gpu_memory.h
class GPUMemory {
private:
    void* d_ptr;
    size_t size;
    
public:
    GPUMemory(size_t bytes) : size(bytes) {
        cudaMalloc(&d_ptr, bytes);
        cudaMemset(d_ptr, 0, bytes);
    }
    
    ~GPUMemory() {
        if (d_ptr) cudaFree(d_ptr);
    }
    
    // Запрещаем копирование
    GPUMemory(const GPUMemory&) = delete;
    GPUMemory& operator=(const GPUMemory&) = delete;
    
    // Разрешаем перемещение
    GPUMemory(GPUMemory&& other) noexcept : d_ptr(other.d_ptr), size(other.size) {
        other.d_ptr = nullptr;
        other.size = 0;
    }
    
    void* get() const { return d_ptr; }
    size_t getSize() const { return size; }
};

3 Ядро матричного умножения (наша версия matmul)

Вся магия Transformer сводится к матричным умножениям. В PyTorch вы просто вызываете torch.matmul(). В CUDA вам нужно написать ядро, которое распределит работу между тысячами потоков.

Неправильный подход (так делают 90% новичков):

// ПЛОХО: наивное матричное умножение
__global__ void naive_matmul(float* A, float* B, float* C, 
                             int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < M && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < K; k++) {
            sum += A[row * K + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

Почему это плохо? Каждый поток читает из глобальной памяти K раз. При K=512 это 512 обращений к медленной памяти. Правильный подход - использовать shared memory:

// ХОРОШО: использование shared memory
template
__global__ void optimized_matmul(float* A, float* B, float* C,
                                 int M, int N, int K) {
    // Выделяем shared memory для блоков матриц
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    
    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    
    // Разбиваем умножение на плитки (tiling)
    for (int tile = 0; tile < (K + BLOCK_SIZE - 1) / BLOCK_SIZE; tile++) {
        // Загружаем плитку из A в shared memory
        int a_col = tile * BLOCK_SIZE + threadIdx.x;
        if (row < M && a_col < K) {
            As[threadIdx.y][threadIdx.x] = A[row * K + a_col];
        } else {
            As[threadIdx.y][threadIdx.x] = 0.0f;
        }
        
        // Загружаем плитку из B в shared memory
        int b_row = tile * BLOCK_SIZE + threadIdx.y;
        if (b_row < K && col < N) {
            Bs[threadIdx.y][threadIdx.x] = B[b_row * N + col];
        } else {
            Bs[threadIdx.y][threadIdx.x] = 0.0f;
        }
        
        __syncthreads();  // Ждем загрузки всех потоков
        
        // Умножаем плитки
        for (int k = 0; k < BLOCK_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        
        __syncthreads();  // Ждем перед следующей плиткой
    }
    
    if (row < M && col < N) {
        C[row * N + col] = sum;
    }
}

4 Attention - сердце Transformer

Multi-head attention - самый сложный компонент. Нужно вычислить QK^T, применить маску, softmax, умножить на V. И все это для 8 голов параллельно.

Ключевая оптимизация: fused attention kernel. Вместо 4 отдельных ядер (matmul, softmax, matmul) пишем одно:

// Fused attention kernel
// Вычисляет attention за один проход
__global__ void fused_attention(
    float* Q, float* K, float* V, float* output,
    int batch_size, int seq_len, int d_model, int num_heads,
    float scale, bool causal_mask) {
    
    extern __shared__ float shared_mem[];
    
    int head_idx = blockIdx.z;
    int batch_idx = blockIdx.y;
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row >= seq_len) return;
    
    int d_head = d_model / num_heads;
    int head_offset = head_idx * d_head;
    
    // Указатели на данные для текущей головы
    float* q_head = Q + batch_idx * seq_len * d_model + head_offset;
    float* k_head = K + batch_idx * seq_len * d_model + head_offset;
    float* v_head = V + batch_idx * seq_len * d_model + head_offset;
    
    // Вычисляем QK^T для текущей строки
    float max_val = -INFINITY;
    float* scores = shared_mem;  // Используем shared memory
    
    for (int col = 0; col < seq_len; col++) {
        if (causal_mask && col > row) {
            scores[col] = -INFINITY;
            continue;
        }
        
        float dot = 0.0f;
        for (int d = 0; d < d_head; d++) {
            dot += q_head[row * d_model + d] * 
                   k_head[col * d_model + d];
        }
        
        scores[col] = dot * scale;
        max_val = fmaxf(max_val, scores[col]);
    }
    
    // Softmax
    float sum_exp = 0.0f;
    for (int col = 0; col < seq_len; col++) {
        if (causal_mask && col > row) continue;
        scores[col] = expf(scores[col] - max_val);
        sum_exp += scores[col];
    }
    
    float inv_sum = 1.0f / sum_exp;
    for (int col = 0; col < seq_len; col++) {
        if (causal_mask && col > row) continue;
        scores[col] *= inv_sum;
    }
    
    // Умножение на V
    float* out_ptr = output + batch_idx * seq_len * d_model + head_offset;
    for (int d = 0; d < d_head; d++) {
        float val = 0.0f;
        for (int col = 0; col < seq_len; col++) {
            if (causal_mask && col > row) continue;
            val += scores[col] * v_head[col * d_model + d];
        }
        out_ptr[row * d_model + d] = val;
    }
}

5 LayerNorm и GeLU активация

LayerNorm стабилизирует обучение. GeLU (Gaussian Error Linear Unit) - современная активация. Обе операции можно объединить в одно ядро:

// Combined LayerNorm + GeLU
__global__ void layernorm_gelu(
    float* input, float* output, float* gamma, float* beta,
    int batch_size, int seq_len, int d_model, float eps = 1e-5f) {
    
    int batch = blockIdx.x;
    int pos = threadIdx.x;
    
    if (pos >= d_model) return;
    
    // Вычисляем mean и var для каждого токена
    __shared__ float shared_sum[1024];
    __shared__ float shared_sq_sum[1024];
    
    float sum = 0.0f;
    float sq_sum = 0.0f;
    
    for (int i = 0; i < seq_len; i++) {
        float val = input[batch * seq_len * d_model + i * d_model + pos];
        sum += val;
        sq_sum += val * val;
    }
    
    shared_sum[pos] = sum;
    shared_sq_sum[pos] = sq_sum;
    __syncthreads();
    
    // Reduce в shared memory
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (pos < stride) {
            shared_sum[pos] += shared_sum[pos + stride];
            shared_sq_sum[pos] += shared_sq_sum[pos + stride];
        }
        __syncthreads();
    }
    
    float mean = shared_sum[0] / (seq_len * d_model);
    float var = shared_sq_sum[0] / (seq_len * d_model) - mean * mean;
    
    // LayerNorm + GeLU
    for (int i = 0; i < seq_len; i++) {
        int idx = batch * seq_len * d_model + i * d_model + pos;
        float x = input[idx];
        float normalized = (x - mean) / sqrtf(var + eps);
        float scaled = normalized * gamma[pos] + beta[pos];
        
        // GeLU approximation
        output[idx] = 0.5f * scaled * 
                     (1.0f + tanhf(0.79788456f * 
                     (scaled + 0.044715f * scaled * scaled * scaled)));
    }
}

6 Собираем все вместе

Теперь у нас есть все компоненты. Собираем их в класс Transformer:

// transformer.h
class Transformer {
private:
    // Параметры модели
    int d_model;      // 512
    int num_heads;    // 8
    int num_layers;   // 4
    int vocab_size;   // 50000
    int max_seq_len;  // 1024
    
    // Веса на GPU
    std::unique_ptr embedding_weights;
    std::vector> layer_weights;
    
    // Временные буферы
    std::unique_ptr workspace;
    
public:
    Transformer(int d_model = 512, int num_heads = 8, 
                int num_layers = 4, int vocab_size = 50000);
    
    // Прямой проход
    void forward(const float* input_tokens, int batch_size, 
                 int seq_len, float* output);
    
    // Инициализация весов (Xavier/Glorot)
    void initialize_weights();
    
    // Сохранение/загрузка
    void save(const std::string& filename);
    void load(const std::string& filename);
};

Реализация forward pass:

// transformer.cu
void Transformer::forward(const float* input_tokens, 
                         int batch_size, int seq_len, 
                         float* output) {
    
    // 1. Embedding lookup
    embedding_kernel<<>>(input_tokens, 
                                      embedding_weights->get(),
                                      workspace->get(),
                                      batch_size, seq_len,
                                      d_model, vocab_size);
    
    // 2. Positional encoding
    add_positional_encoding<<>>(workspace->get(),
                                            batch_size, seq_len,
                                            d_model);
    
    float* current = workspace->get();
    
    // 3. Transformer layers
    for (int layer = 0; layer < num_layers; layer++) {
        // Self-attention
        fused_attention<<>>(
            current, current, current, attention_output,
            batch_size, seq_len, d_model, num_heads,
            1.0f / sqrtf(d_model / num_heads), true);
        
        // Add & Norm
        add_layernorm<<>>(current, attention_output,
                                      layer_norm_gamma[layer],
                                      layer_norm_beta[layer],
                                      batch_size, seq_len, d_model);
        
        // FFN
        ffn_kernel<<>>(current, ffn_output,
                                    ffn_weights1[layer],
                                    ffn_weights2[layer]);
        
        // Final Add & Norm
        add_layernorm<<>>(current, ffn_output,
                                      final_norm_gamma[layer],
                                      final_norm_beta[layer],
                                      batch_size, seq_len, d_model);
        
        current = layer_output[layer];
    }
    
    // 4. Final projection
    final_projection<<>>(current, output,
                                      projection_weights,
                                      batch_size, seq_len,
                                      d_model, vocab_size);
    
    cudaDeviceSynchronize();  // Ждем завершения всех ядер
}

Оптимизации, которые дают 10x ускорение

Наивная реализация будет работать в 10-50 раз медленнее PyTorch. Вот что нужно добавить:

1. Tensor Cores: Используйте wmma::mma_sync для матричных умножений на тензорных ядрах (доступно с Compute Capability 7.0)
2. FP8 вместо FP32: Новые карты RTX 5000 серии поддерживают FP8 natively. Экономит память и ускоряет вычисления в 2 раза. Почитайте про FP8 на RTX 3090 для вдохновения
3. Кэширование весов в shared memory: Если веса слоя помещаются в shared memory, загружайте их один раз
4. Pinned memory для host-GPU копирования: Используйте cudaMallocHost вместо malloc
5. Streams для overlap: Запускайте копирование данных и вычисления параллельно в разных стримах

Отладка - это боль

CUDA ошибки бывают трех видов:

• Runtime ошибки: cudaError_t. Легко отлавливаются cudaGetLastError()
• Memory errors: Out of bounds доступ. Используйте cuda-memcheck
• Race conditions: Самое страшное. Потоки пишут в одну память без синхронизации. Помогает только atomicAdd и тщательный дизайн

Мой workflow отладки:

# 1. Проверка памяти
cuda-memcheck ./transformer_test

# 2. Профилирование
nsys profile --stats=true ./transformer_test

# 3. Визуализация занятости SM
nvprof --metrics achieved_occupancy ./transformer_test

# 4. Отладчик CUDA-GDB (если совсем плохо)
cuda-gdb ./transformer_test

Сравнение с высокоуровневыми фреймворками

Написали 2000 строк кода на CUDA. Что получили взамен?

Выигрыш есть, но цена - тысячи строк кода и недели отладки. Стоило ли оно того? Если вы хотите понять, как на самом деле работают LLM - однозначно да. Если нужно просто запустить модель - используйте готовые решения на C++ и CUDA.

Что дальше?

Вы написали базовый Transformer. Теперь можно:

1. Добавить обучение (backpropagation на CUDA - это отдельная история)
2. Реализовать смешанную точность (FP16, BF16, FP8)
3. Добавить поддержку MoE (Mixture of Experts) как в cuda-nn
4. Оптимизировать для конкретной архитектуры GPU (RTX 4090 vs H100)
5. Добавить квантование весов для экономии памяти

Финальный совет

Пишите код на CUDA не для production, а для обучения. Каждая ошибка, каждый segfault, каждый race condition учит вас чему-то новому о том, как работает железо. После этого когда вы читаете про новые фичи в Transformers v5, вы понимаете не только что добавили, но и почему это сложно было реализовать.

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

И да, теперь когда вы видите статью про то, как Claude пишет CUDA-ядра за вас, вы не просто верите на слово - вы можете проверить, действительно ли сгенерированный код оптимален.