Введение в CUDA: модель потоков, блоков и сеток (grid/block/thread) на простом примере

Введение

CUDA — платформа параллельных вычислений от NVIDIA, позволяющая выполнять код на тысячах ядер GPU одновременно. В отличие от CPU с несколькими быстрыми ядрами, GPU состоит из множества простых вычислительных блоков, эффективных при массово-параллельной обработке однотипных данных. Чтобы написать первую CUDA-программу, нужно понять иерархию выполнения: поток (thread) — блок (block) — сетка (grid). В статье разберём эту модель на примере простого покомпонентного сложения двух массивов.

Концепция

Каждый вызов CUDA-ядра (kernel) запускает множество параллельных потоков, организованных в иерархию. Потоки группируются в блоки (block), а блоки — в сетку (grid). Каждый поток получает уникальные индексы threadIdx (внутри блока) и blockIdx (блока внутри сетки), а размеры блока и сетки доступны через blockDim и gridDim. Комбинируя эти величины, каждый поток вычисляет свой глобальный индекс и обрабатывает свою часть данных — это и есть базовый паттерн SIMT (single instruction, multiple threads): все потоки выполняют один и тот же код ядра, но над разными элементами данных.

Размер блока обычно выбирают кратным 32 (размер warp — группы потоков, выполняемых физически синхронно на одном SM), типичные значения — 128, 256, 512 потоков на блок.

Пример кода

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

// CUDA-ядро: каждый поток складывает один элемент массивов
__global__ void vectorAdd(const float *a, const float *b, float *c, int n)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main()
{
    const int n = 1'000'000;
    const size_t bytes = n * sizeof(float);

    std::vector<float> hostA(n, 1.0f);
    std::vector<float> hostB(n, 2.0f);
    std::vector<float> hostC(n);

    float *deviceA, *deviceB, *deviceC;
    cudaMalloc(&deviceA, bytes);
    cudaMalloc(&deviceB, bytes);
    cudaMalloc(&deviceC, bytes);

    cudaMemcpy(deviceA, hostA.data(), bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(deviceB, hostB.data(), bytes, cudaMemcpyHostToDevice);

    const int threadsPerBlock = 256;
    const int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;

    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(deviceA, deviceB, deviceC, n);
    cudaDeviceSynchronize();

    cudaMemcpy(hostC.data(), deviceC, bytes, cudaMemcpyDeviceToHost);

    std::cout << "Результат: " << hostC[0] << ", " << hostC[n - 1] << std::endl;

    cudaFree(deviceA);
    cudaFree(deviceB);
    cudaFree(deviceC);
    return 0;
}

Пояснения к коду

Ядро vectorAdd помечено квалификатором __global__, указывающим, что функция выполняется на GPU, но вызывается с CPU (host-кода). Каждый поток вычисляет свой глобальный индекс idx через формулу blockIdx.x * blockDim.x + threadIdx.x — это стандартная идиома преобразования двухуровневой иерархии block/thread в линейный индекс элемента массива. Проверка if (idx < n) обязательна, поскольку число запущенных потоков (blocksPerGrid * threadsPerBlock) округляется вверх до кратного блоку числа и может превышать реальный размер данных n.

Вызов ядра в синтаксисе vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(...) — особый CUDA-синтаксис конфигурации запуска, задающий размер сетки и блока. cudaDeviceSynchronize() блокирует host до завершения всех запущенных на GPU операций — без этого вызова чтение результата может начаться до того, как GPU закончит вычисления.

Подводные камни

Отсутствие проверки границ массива (if (idx < n)) — частая причина выхода за пределы выделенной памяти, который на GPU не всегда вызывает явный краш (в отличие от CPU), а может привести к порче соседних данных или undefined behavior, которое проявляется непредсказуемо.
Игнорирование кодов ошибок CUDA API. Каждый вызов cudaMalloc, cudaMemcpy и другие возвращают cudaError_t — игнорирование этого значения (как сделано в упрощённом примере для краткости) на практике маскирует реальные проблемы (нехватка памяти GPU, неверный указатель) тихим отказом операции.
Запуск ядра с числом потоков, не кратным размеру warp (32). Хотя такая конфигурация технически работает, она может приводить к недозагрузке последнего warp в каждом блоке — для производительности стоит выбирать threadsPerBlock кратным 32.
Передача указателей на host-память напрямую в ядро. GPU не может напрямую разрешать указатели на CPU-память (за исключением Unified Memory, рассмотренной отдельно) — попытка передать host-указатель в vectorAdd приведёт к ошибке доступа к памяти во время выполнения ядра.