Цей документ є першою частиною серії статей, присвячених CUDA. Він надає введення в мультипоточність.
Припустимо, ми маємо чорно-біле зображення розміру 4096×2048 для обробки. Метою є використання масивної паралелізації, яку надає CUDA, для виконання операцій на графічному процесорі, звичайно відомому як GPU. Кожному обчислювальному елементу, який зазвичай називається потоком (thread), призначено піксель на зображенні.
CUDA дозволяє паралелізувати будь-яку програму за допомогою сітки, яка може бути:
- Одновимірною: проста спискова структура, де елементи ідентифікуються через координату x;
- Двовимірною: таблиця, де кожен елемент розташований у точці (x, y);
- Тривимірною: тензор третього порядку, компоненти якого ідентифікуються за їх координатами (x, y, z).
Крім того, ця сітка поділяється на кілька блоків однакового розміру (див. Рисунок 1). Ці блоки містять різні потоки, які працюють разом над зображенням. Перевага цієї проміжної структури полягає в концепції спільної пам'яті. Усі потоки в блоці мають спільний простір пам'яті, що дозволяє швидший доступ, але є значно більш обмеженим порівняно з глобальною пам'яттю, яка доступна всім потокам по всій сітці. Це може бути корисно для дуже специфічних оптимізацій.
Рисунок 1 — Організація потоків у двовимірній сітці
Функціональний код CUDA визначений двома окремими частинами: одна виконується на процесорі (CPU), інша — на графічному процесорі (GPU). Ось приклад псевдокоду. Він містить функцію hostFunction(). Це функція, яка виконується на процесорі і, з одного боку, встановлює розміри сітки і блоків, а з іншого боку, викликає функцію deviceFunction(), яка має працювати на GPU. Хоча це не показано тут, функція host також відповідає за виділення необхідної пам'яті для передачі даних від процесора до графічного процесора.
// Запуск на GPU
__global__ void deviceFunction(deviceArg1, deviceArg2, ...)
{
int tidx = blockIdx.x * blockDim.x + threadIdx.x;
int tidy = blockIdx.y * blockDim.y + threadIdx.y;
int threadID = tidy * imgWidth + tidx;
...
}
// Запуск на CPU
void hostFunction(hostArg1, hostArg2, ...)
{
...
int blockWidth = 32, blockHeight = 16;
int imgWidth = 4096, imgHeight = 2048;
dim3 gridDim((imgWidth - 1 + blockWidth) / blockWidth,
(imgHeight - 1 + blockHeight) / blockHeight);
// = (128, 128, 1)
dim3 blockDim(blockWidth, blockHeight); // = (32, 16, 1)
deviceFunction<<>>(deviceArg1, deviceArg2, ...);
...
}
Крім того, в функції hostFunction(), оскільки кожен потік асоціюється з пікселем на зображенні, кількість потоків повинна дорівнювати кількості пікселів у сітці. Таким чином, зображення розміру 4096×2048 ділиться на сітку розміру 128×128, де кожен блок має розміри 32×16. Оскільки 4096 = 128×32 і 2048 = 128×16, обчислення вірне. Формула в псевдокоді є загальною і дозволяє правильно встановити розміри сітки незалежно від розмірів блоків.
У функції device кожен потік ідентифікується унікальною парою (tidx, tidy), таким чином:
tidx = blockIdx.x * blockDim.x + threadIdx.x;tidy = blockIdx.y * blockDim.y + threadIdx.y;
де:
blockIdxпредставляє 2D координати блоку, що містить активний потік;threadIdxпредставляє індекс активного потоку в межах блоку;blockDimпредставляє розміри блоків, які складають сітку.
Нарешті, CUDA спирається на програмну філософію, звану SIMT (Single Instruction, Multiple Threads — одна інструкція, багато потоків). Цей підхід ілюструється наступною операцією: threadID = tidy * imgWidth + tidx.
Це означає, що, незважаючи на те, що кожен потік виконує одну й ту саму інструкцію, унікальність пари індексів (tidx, tidy) забезпечує те, що кожен потік обробляє різні дані.
Висновок: Цей документ став можливістю для першого знайомства з мультипоточністю за допомогою CUDA, зокрема з організацією потоків усередині блоків, які в свою чергу структуровані в сітку. Однак слід зазначити, що все, що було охоплено тут, стосується інтерфейсу користувача. У майбутньому документі буде детальніше розглянуто, що відбувається на стороні машини, зокрема концепції, такі як **warp-и*, *CUDA ядра** та **стрімінгові мультипроцесори_**.
Джерела
- Ampere Tuning Guide, NVIDIA Doc, 20 листопада 2024
- CUDA C++ Programming Guide, NVIDIA Doc, 20 листопада 2024
- Issue Efficiency, NVIDIA Doc, 2015
Перекладено з: How to optimise your CUDA code — GPU programming, Part I: First approach