Блок потоків () — програмна абстракція в CUDA, яка представляє групу потоків, які можуть виконуватися послідовно або паралельно. Для зручного співставлення процесів і даних потоки групуються в блоки потоків. Кількість потоків у блоці потоків раніше обмежувалася архітектурою до 512 потоків на блок, але станом на березень 2010 року, з обчислювальними можливостями 2.x і вище, блоки можуть містити до 1024 потоків. Потоки в одному блоці потоків виконуються на одному потоковому мультипроцесорі. Потоки в одному блоці можуть взаємодіяти один з одним через спільну пам'ять, бар'єрну синхронізацію або інші примітиви синхронізації, такі як атомарні операції.

Кілька блоків формують сітку (grid). Усі блоки в одній сітці містять однакову кількість потоків. Кількість потоків у блоці обмежена, але сітки можна використовувати для обчислень, які потребують паралельної роботи великої кількості блоків потоків та використати всі доступні мультипроцесори.

CUDA — це платформа паралельних обчислень та модель програмування, яку мови вищого рівня можуть використовувати для реалізації паралелізму. У CUDA ядра виконуються в потоках. Потік — це абстрактна сутність, яка представляє виконання ядра. Ядро — це функція, яка компілюється для запуску на CUDA пристрої. В багатопотоковій програмі потоки працюють одночасно для організації паралельних обчислень. Кожен потік має індекс, який використовується для визначення адрес пам'яті, а також для прийняття рішень.

Розміри CUDA працює на гетерогенній моделі програмування, яка використовується для запуску прикладних програм на хост-пристроях. Вона має модель виконання, подібну до OpenCL. У цій моделі ми починаємо виконувати програму на хост-пристрої, яким зазвичай є ядро процесора. Пристрій є пропускно-орієнтованим пристроєм, тобто ядром графічного процесора, яке виконує паралельні обчислення. Для цих паралельних виконань використовуються функції ядра. Після виконання цих функцій ядра керування передається назад на хост-пристрій, який відновлює послідовне виконання.

Оскільки багато паралельних програм використовують багатовимірні дані, зручно організовувати блоки потоків в 1D, 2D або 3D масиви потоків. Блоки в сітці повинні мати можливість виконуватися незалежно, оскільки зв'язок або співпраця між блоками в сітці неможливі. "Коли ядро запускається, кількість потоків на блок потоку, а також кількість блоків потоків, вказуються, що, у свою чергу, визначає загальну кількість запущених потоків CUDA. " Максимальні розміри x, y та z блоку становлять 1024, 1024 та 64, і він повинен бути розподілений таким чином, щоб x × y × z ≤ 1024, що є максимальною кількістю потоків на блок. Блоки можуть бути організовані в одно-, дво- або тривимірні сітки до , 65 535 та 65 535 блоків у вимірах x, y та z відповідно. На відміну від максимальної кількості потоків на блок, немає обмеження на кількість блоків на сітку, відмінного від максимальних розмірів сітки.

Індексування

1D-індексування Кожен потік у CUDA пов'язаний з певним індексом, щоб він міг обчислювати та отримувати доступ до комірок пам'яті в масиві.

Розглянемо приклад, у якому є масив із 512 елементів. Одна з можливих структур організації обчислень — це сітка з одним блоком, який має 512 потоків. Розглянемо задачу знаходження масиву C із 512 елементів, який складається з поелементного множення двох масивів A та B, кожен з яких містить по 512 елементів. Кожен потік має індекс i, і він виконує множення i-го елемента A та B, а потім зберігає результат в i-му елементі C. Індекс i обчислюється за допомогою blockIdx (який у цьому випадку дорівнює 0, оскільки є лише один блок), blockDim (512 у цьому випадку, оскільки блок має 512 елементів) та threadIdx, який змінюється від 0 до 511 для кожного потоку.

Індекс потоку i розраховується за такою формулою :

blockIdx.x — це х індекс блоку.

blockDim.x — це розмір блоку вздовж вісі x

threadIdx.x — це x індекс потоку

Якщо ми хочемо розглянути обчислення для масиву, більшого за 1024, ми можемо мати кілька блоків з 1024 потоками кожен. Розглянемо приклад з 2048 елементами масиву. У цьому випадку у нас є 2 блоки потоків з 1024 потоками кожен. Таким чином, значення ідентифікаторів потоків змінюватимуться від 0 до 1023, ідентифікатор блоку змінюватиметься від 0 до 1, а розмірність блоку становитиме 1024. Таким чином, перший блок отримуватиме значення індексу від 0 до 1023, а останній матиме значення індексу від 1024 до 2047.

Таким чином, кожен потік спочатку обчислить індекс пам'яті, до якої він має доступ, а потім виконає потрібне обчислення. Розглянемо приклад, у якому елементи масивів A та B додаються паралельно в потоках, а результати зберігаються в масиві C. Відповідний код, що виконується в кожному потоці показано нижче.

2D-індексування Так само в ідентифікатор блоку (blockId), а також ідентифікатор потоку (threadId) потрібно обчислити для кожного потоку залежно від геометрії сітки. Розглянемо двовимірну сітку з двовимірними блоками. Ідентифікатор потоку (threadId) та ідентифікатор блоку (blockId) будуть обчислюватися за такими формулами:

Апаратна перспектива Хоча ми вже визначили ієрархію потоків, слід зазначити, що потоки, блоки потоків та сітка — це, по суті, перспектива програміста. Щоб отримати повне уявлення про блок потоків, важливо знати їх з точки зору апаратного забезпечення. Апаратне забезпечення групує потоки, які виконують одну й ту саму інструкцію, у варпи. Кілька варпінгів утворюють блок потоку. Кілька блоків потоків призначаються потоковому мультипроцесору (SM). Кілька SM утворюють цілий блок графічного процесора (який виконує всю сітку ядра).

Потокові мультипроцесори Кожна архітектура графічного процесора (наприклад, Kepler або Fermi) складається з кількох потокових мультипроцесорів (SM) або потокових процесорів (Streaming Multiprocessors). Це процесори загального призначення з низькою цільовою тактовою частотою та невеликим кешем. SM здатний виконувати кілька блоків потоків паралельно. Як тільки один з його блоків потоків завершує виконання, він послідовно переходить до наступного блоку потоку. Загалом, SM підтримують паралелізм на рівні інструкцій, але не прогнозування розгалужень.

Примітки