Программная модель CUDA: грид, блоки, треды

1. Почему видеокарта считает быстрее — и когда нет

Центральный процессор спроектирован, чтобы выполнить одну цепочку инструкций как можно быстрее: большие кэши, предсказание ветвлений, высокая тактовая частота. Ядер у него немного — единицы или десятки. Видеокарта устроена наоборот: тысячи простых ядер без хитрой логики. Каждое из них медленнее ядра CPU, но их очень много.

Отсюда правило, которое стоит запомнить до всякого кода: CPU выигрывает по латентности (время одной операции), GPU — по пропускной способности (операций в секунду на всём устройстве). Если задача — миллион одинаковых независимых вычислений, GPU разложит их по своим ядрам и обгонит процессор на порядок. Если задача — длинная цепочка, где каждый шаг ждёт предыдущий, тысячи ядер будут простаивать, и CPU окажется быстрее. Какие алгоритмы «раскладываются», а какие нет — отдельная большая тема (C04 и конспект P02 соседней линии); пока договоримся о словах.

Модель выполнения GPU называется SIMT (Single Instruction, Multiple Threads): одна инструкция, много тредов. Группа тредов шагает по программе строем — все выполняют одну и ту же строку кода, но каждый над своими данными.

2. Словарь CUDA: одна аналогия на восемь терминов

Представьте завод, получивший заказ. Заказ — это ядро: одна инструкция «что делать». Под заказ собирают всех исполнителей — это грид. Исполнители организованы в бригады-блоки: у бригады общий склад инструментов (общая память) и планёрки (синхронизация); бригада целиком работает в одном цехе-мультипроцессоре. Внутри бригады рабочие-треды ходят шеренгами по 32 человека — варпами — и шагают строго в ногу: одна команда на шеренгу. Теперь строго:

Почему архитекторы CUDA не остановились на паре «грид — треды», зачем прослойка из блоков? Потому что железо не резиновое: общая быстрая память и синхронизация возможны только в пределах одного SM. Блок — это обещание программиста «эти треды будут общаться между собой», и планировщик кладёт их на один SM. Треды разных блоков друг о друге ничего не знают — зато блоков можно запустить сколько угодно, и грид масштабируется на любую карту: у дешёвой SM четыре, у топовой — за сотню, а программа одна и та же.

3. Спецификаторы: кто вызывает и где выполняется

В .cu-файле живут функции трёх сортов, и различаются они префиксом-спецификатором:

Полезная комбинация __host__ __device__ помечает функцию, которая компилируется дважды — и для CPU, и для GPU: удобно для маленькой математики вроде clamp или преобразования координат, чтобы не писать её два раза.

// вспомогательная функция устройства: вызывается из ядра
__device__ float squared(float x)
{
    return x * x;
}

// ядро: вызов с хоста, исполнение на устройстве, возвращает всегда void
__global__ void kernel(float* data)
{
    data[threadIdx.x] = squared(data[threadIdx.x]);
}

4. Запуск ядра и встроенные переменные

Тройные угловые скобки из C01 — это конфигурация грида:

kernel<<< 3, 64 >>>(аргументы);   // 3 блока × 64 треда = 192 треда

Первое число — сколько блоков в гриде, второе — сколько тредов в каждом блоке. (У конструкции есть ещё два необязательных параметра — размер динамической разделяемой памяти и поток выполнения; они понадобятся в C05 и позже.) Внутри ядра каждый тред ориентируется по четырём встроенным переменным:

У всех четырёх есть поля .x, .y, .z: грид и блоки могут быть одномерными, двумерными и трёхмерными. Для картинок и матриц удобна двумерная конфигурация — она задаётся типом dim3 (незаполненные компоненты становятся единицами):

dim3 block(16, 16);                // блок 16×16 = 256 тредов
dim3 grid (32, 32);                // грид 32×32 блока — хватит на картинку 512×512
kernel<<< grid, block >>>(…);     // внутри ядра работают threadIdx.y, blockIdx.y

Соберём всё в работающий пример — «перекличку» тредов:

#include <cstdio>

__global__ void rollcall()
{
    printf("блок %d, тред %d\n", blockIdx.x, threadIdx.x);
}

int main()
{
    rollcall<<<2, 3>>>();        // 2 блока по 3 треда
    cudaDeviceSynchronize();      // дождаться видеокарту (C01)
    return 0;
}

> nvcc rollcall.cu -o rollcall && ./rollcall
блок 1, тред 0
блок 1, тред 1
блок 1, тред 2
блок 0, тред 0
блок 0, тред 1
блок 0, тред 2

Обратите внимание: блок 1 отчитался раньше блока 0 — и при следующем запуске порядок может стать другим. Треды внутри одного варпа шагают в ногу, но блоки и варпы планировщик запускает в том порядке, в каком ему удобно.

5. Тренажёр: грид на ладони

Контрольные вопросы

• Чем GPU принципиально отличается от CPU и в каких задачах проигрывает?
  CPU оптимизирован под латентность — быстро выполнить одну цепочку инструкций; GPU — под пропускную способность: тысячи простых ядер выполняют массу одинаковых независимых операций. GPU проигрывает там, где шаги зависят друг от друга и параллелить нечего: длинные последовательные цепочки, ветвистая логика.
• Зачем в иерархии CUDA нужны блоки — почему не просто «грид из тредов»?
  Общая быстрая память и синхронизация физически возможны только в пределах одного SM. Блок — это группа тредов, которой обещаны эти возможности, поэтому он целиком выполняется на одном SM. А граница между блоками даёт масштабируемость: блоки независимы, и один и тот же грид раскладывается и на 4 SM дешёвой карты, и на 100 SM топовой.
• Что такое варп и почему его размер важен при выборе blockDim?
  Варп — 32 подряд идущих треда блока, физически выполняющихся одновременно одной инструкцией (SIMT). Блок нарезается на варпы целиком: при blockDim = 40 получится варп из 32 тредов и варп, где занято только 8 мест из 32, — остальные 24 «места в шеренге» пропадают. Поэтому размер блока выбирают кратным 32.
• Куда __global__-функция девает результат, если возвращать значение ей запрещено?
  Записывает в память устройства через переданные указатели-аргументы, а хост потом копирует данные себе. Возврат значения невозможен из-за асинхронности запуска: хост не ждёт ядро в точке вызова.
• Запуск kernel<<<4, 256>>>(): сколько тредов, чему равны blockDim.x и gridDim.x, каков диапазон threadIdx.x?
  Всего 4 × 256 = 1024 треда; gridDim.x = 4, blockDim.x = 256; threadIdx.x пробегает 0…255 в каждом блоке, blockIdx.x — 0…3.
• Гарантирует ли CUDA порядок выполнения блоков и что из этого следует?
  Нет: планировщик запускает блоки в произвольном порядке, и от запуска к запуску он меняется. Следствие: нельзя строить логику на том, что один блок «успеет раньше» другого; треды пишут каждый в свои ячейки, а координация — только явными механизмами синхронизации.

Источники

1. CUDA C++ Programming Guide. Programming Model // NVIDIA Docs : [сайт]. — URL: https://docs.nvidia.com/cuda/cuda-c-programming-guide/ (дата обращения: 09.07.2026).
2. An Even Easier Introduction to CUDA // NVIDIA Developer Blog : [сайт]. — URL: https://developer.nvidia.com/blog/even-easier-introduction-cuda/ (дата обращения: 09.07.2026).
3. CUDA Refresher: The CUDA Programming Model // NVIDIA Developer Blog : [сайт]. — URL: https://developer.nvidia.com/blog/cuda-refresher-cuda-programming-model/ (дата обращения: 09.07.2026).