GPU и data-parallel: модель SIMT (CUDA)

Урок про массовый параллелизм видеокарты и про то, какие задачи ей подходят.

SIMT (Single Instruction, Multiple Threads) — модель GPU: тысячи лёгких потоков выполняют одну и ту же программу (kernel) над разными элементами данных.

Чем GPU отличается от CPU

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

Иерархия потоков

В CUDA потоки организованы в иерархию: thread (один поток) → warp (32 потока, выполняющиеся синхронно) → block (группа потоков с общей быстрой памятью) → grid (все блоки запуска). Программист пишет ядро так, будто описывает работу одного потока, а GPU запускает тысячи копий, каждая со своим индексом. Ниже — CUDA-ядро сложения векторов; оно исполняется только на GPU, поэтому блок нечитаемый для браузерного раннера.

__global__ void add(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];   // каждый поток считает один элемент
    }
}

// запуск: тысячи потоков, по одному на элемент
// add<<<blocks, threadsPerBlock>>>(a, b, c, n);

Каждый поток вычисляет свой индекс i и обрабатывает один элемент. Запуск add<<<blocks, threads>>> порождает тысячи потоков сразу. Это чистый data-parallel: операция одна, данные разные.

Дивергенция ветвлений

Главная ловушка SIMT — warp из 32 потоков выполняет инструкции синхронно. Если внутри warp потоки идут разными ветвями if (одни в then, другие в else), GPU вынужден выполнить обе ветви последовательно, маскируя ненужные потоки. Это дивергенция, и она режет производительность вдвое и сильнее. Поэтому GPU-код пишут так, чтобы соседние потоки шли одной ветвью.

warp без дивергенции: все 32 потока -> одна ветвь (быстро)
warp с дивергенцией:  16 в then, 16 в else
   -> GPU гонит обе ветви по очереди, маскируя половину (медленно)

Когда GPU выигрывает, а когда нет

Стоимость передачи данных

Данные живут в памяти CPU, а считать их надо на GPU — значит, их нужно переслать через шину PCIe, что медленно. Если работы на GPU мало, передача туда-обратно съест весь выигрыш. Правило: GPU окупается, когда вычислений много относительно объёма пересылаемых данных. Эмуляция логики (без GPU) — обычный поэлементный map:

# логика GPU-ядра add: каждый "поток" считает один элемент
a = [1.0, 2.0, 3.0, 4.0]
b = [10.0, 20.0, 30.0, 40.0]
c = [a[i] + b[i] for i in range(len(a))]   # на GPU — тысячи потоков сразу
print("c =", c)

Вывод:

c = [11.0, 22.0, 33.0, 44.0]

Как работает под капотом

GPU прячет задержку памяти не кэшами, а переключением warp'ов: пока один warp ждёт данные из памяти, планировщик мгновенно запускает другой, готовый к счёту. Поэтому GPU нужно много потоков — чтобы всегда было кого запустить вместо ожидающих. Это противоположно CPU, который прячет задержку кэшами и предсказанием. Отсюда правило: загружай GPU тысячами потоков, иначе его счётные блоки простаивают в ожидании памяти.

Понять GPU — значит принять, что он спроектирован по противоположной CPU философии. CPU — спринтер: несколько мощных ядер, огромные кэши, хитрое предсказание ветвлений, всё ради того, чтобы один поток домчался до результата как можно быстрее. GPU — толпа: тысячи слабых ядер, крошечные кэши, никакого хитроумия на поток, зато способность держать в полёте десятки тысяч потоков сразу. На задаче «один сложный путь» побеждает CPU; на задаче «миллион одинаковых простых путей» побеждает GPU, и с огромным отрывом. Выбор ускорителя — это, по сути, диагноз формы вашей задачи: она длинная и ветвистая или широкая и однородная?

Частые ошибки

• Гнать на GPU задачу с малым параллелизмом или сильными ветвлениями — он проиграет CPU.
• Игнорировать стоимость передачи CPU↔GPU — для мелкой работы она доминирует.
• Допускать дивергенцию warp'ов ветвлениями — производительность падает кратно.

Итоги

• GPU (SIMT) запускает тысячи лёгких потоков, выполняющих одно ядро над разными данными.
• Дивергенция ветвлений в warp'е сериализует ветви и режет производительность.
• GPU выигрывает на регулярных data-parallel задачах с большим объёмом вычислений на элемент.
• Задержку памяти GPU прячет переключением warp'ов — поэтому нужно много потоков.