Цей документ є другою частиною серії статей, присвячених CUDA. Він надає спрощене представлення деяких низькорівневих механізмів, що працюють у багатопотокових програмах._
Припустимо, ми працюємо з GPU від NVIDIA RTX A4500 , який працює на архітектурі Ampere і є дуже ефективним для високопродуктивних обчислень. Метою цієї частини є роз’яснення того, як організовано багатопоточне програмування в CUDA, поза межами користувацького інтерфейсу.
Багатопоточність в CUDA
Стрімінгові багатопроцесори і ядра CUDA
Для цього необхідно визначити фізичні одиниці всередині GPU, які є важливими для нашого завдання. На малюнку 1 показано, серед іншого, стрімінгові багатопроцесори (SM). Вони містять ядра CUDA, обчислювальні одиниці, які відповідають за виконання інструкцій потоку. Потоки, як нагадування, перебувають у блоках, які разом формують сітку. Крім того, є ще одна підсистема: варпи. Це групи з 32 потоків усередині блоків, які запускаються одночасно набором з 32 ядер CUDA.
Малюнок 1 — Структура GPU та запуск потоків ядрами CUDA
Стани варпів
RTX A4500 містить 56 SM, кожен з яких має 128 ядер CUDA. Це означає, що до 7168 потоків можуть виконувати інструкцію в межах одного такту. SM також оснащений чотирма планувальниками варпів, які під час кожного циклу вибирають варпи для виконання ядрами CUDA. Якщо всі потоки у варпі завершили виконання, один із планувальників вибирає інший варп, що чекає виконання. Отже, ключ до функціональності CUDA полягає не лише у багатопоточності, але й у швидкому перерасподілянні ресурсів, що дає змогу виконувати допустимі варпи.
Стан варпа можна визначити за п’ятьма статусами:
- Активний: керується ядрами CUDA;
- Очікуючий: активний, але не може виконати наступну інструкцію через залежності даних;
- Доступний: готовий до виконання своїх інструкцій;
- Завершений: усі інструкції оброблено. Апаратура звільнена.
Ми також можемо визначити п’ятий статус, який базується на апаратних або пам’ятевих обмеженнях, властивих SM. Ті, що складають RTX A4500, можуть містити максимум:
- 2048 потоків / 32 блоки;
- 164 КБ спільної пам’яті;
- 64K регістрів, що еквівалентно 64 × 2¹⁰ = 65,536 регістрів. Регістр — це невелика, надшвидка одиниця пам’яті розміром 32 біти, що виділяється для зберігання приватних змінних потоку. Кожен потік може використовувати кілька регістрів.
Таким чином, якщо варпи в блоках вимагатимуть надмірних апаратних або пам’ятевих ресурсів, багато з них потраплять у цей п’ятий стан:
- Не резидентний на SM: варп і його блок чекають на звільнення пам’яті або апаратних ресурсів, щоб залишитися в SM. Блоки або повністю знаходяться в одному з SM GPU, або зовсім поза їх межами. Це важливо для забезпечення узгодженості спільної пам’яті в межах блоку.
Давайте проілюструємо це спрощеним прикладом.
Практичний приклад: Обчислення кількості резидентних потоків
Припустимо, у нас є машина з RTX A4500. У наступному питанні ми не будемо враховувати оптимізації компілятора CUDA (NVCC).
Розглянемо наведений код, скільки потоків теоретично можуть одночасно перебувати в SM GPU?
// Запуск на GPU
__global__ void deviceFunction(float4* output, float3* additionalOutput, int numThreads)
{
int tidX = threadIdx.x + blockIdx.x * blockDim.x;
int tidY = threadIdx.y + blockIdx.y * blockDim.y;
int threadID = tidY * (gridDim.x * blockDim.x) + tidX;
if (threadID < numThreads)
{
float4 privateVar1[15];
float3 privateVar2;
for (int i = 0; i < 15; i++)
{
privateVar1[i].x = threadID * 0.1f;
privateVar1[i].y = threadID * 0.2f;
privateVar1[i].z = threadID * 0.3f;
privateVar1[i].w = threadID * 0.4f;
}
privateVar2.x = threadID * 0.5f;
privateVar2.y = threadID * 0.6f;
privateVar2.z = threadID * 0.7f;
output[threadID] = privateVar1[0];
additionalOutput[threadID] = privateVar2;
}
}
// Запуск на CPU
void hostFunction()
{
int gridSizeX = 128, gridSizeY = 128;
int blockSizeX = 32, blockSizeY = 16;
int numThreads = gridSizeX * gridSizeY * blockSizeX * blockSizeY; // = 8,388,608 потоків
float4* devOutput;
float3* devAdditionalOutput;
cudaMalloc(&devOutput, numThreads * sizeof(float4));
cudaMalloc(&devAdditionalOutput, numThreads * sizeof(float3));
dim3 gridDim(gridSizeX, gridSizeY);
dim3 blockDim(blockSizeX, blockSizeY);
deviceFunction<<>>(devOutput, devAdditionalOutput, numThreads);
cudaDeviceSynchronize();
cudaFree(devOutput);
cudaFree(devAdditionalOutput);
return;
}
При перевірці програми ми бачимо, що спільна пам'ять не використовується. Таким чином, факторами, що обмежують кількість потоків, які можуть одночасно перебувати в SM, є:
- Апаратна структура SM;
- Використання регістрів.
Обмеження апаратної структури
Функція host ініціалізує блоки по 32 × 16 = 512 потоків. Відповідно, оскільки GPU обмежує кількість потоків, що можуть бути резидентними на одному SM, до 2048, ми можемо розмістити максимум 2048 потоків/SM ÷ 512 потоків/блок = 4 блоки/SM.
Обмеження, що накладаються використанням регістрів
З функції пристрою можна зробити висновок про наступне використання регістрів.
Чотири змінні типу int (по 32 біти кожна): tidX, tidY, threadID і i. Це складає 4 регістри.
privateVar1: масив з 15 елементів типу float4, кожен з яких вимагає 4 × 32 біти. Це складає 15 × 4 = 60 регістрів.
privateVar2: тип float3. Однак CUDA обробляє дані та сутності кратно 2ⁿ × 32, де n ∈ N. Ця логіка пояснює, чому:
- Варп складається з 32 потоків;
- Блок завжди повинен містити 2ⁿ варпів.
Якби кількість потоків у блоці не була кратною 32, 1–31 надлишкових потоків все одно оброблялися б 32 ядрами CUDA. Наприклад, якби залишався 1 потік, 31 ядро CUDA було б виділено для порожніх слотів. Це було б марнотратством ресурсів.
Згідно з цією логікою, тип float3 займає 4 × 32 біти:
- 3 × 32 біти для зберігання даних
- 32 біти для доповнення
Це 4 регістри.
Отже, кожен потік використовує 68 регістрів для виконання своїх інструкцій, не враховуючи оптимізації компілятора. Оскільки кожен SM може вмістити максимум 65,536 регістрів, це встановлює межу на:
65,536 регістри/SM ÷ 68 регістри/потік ≈ 963 потоки/SM,
що призводить до 963 потоки/SM ÷ 512 потоки/блок ≈ 1.88 блоків/SM.
Апаратні критерії визначають, що кожен SM може вмістити максимум 4 блоки. Однак, через використання регістрів, це обмеження зменшується до 1.88. Оскільки блок повинен повністю розміщуватись в одному SM, один блок буде розміщений на кожному SM.
Висновок: Кількість потоків, що одночасно перебувають в SM GPU, становить:
56 SM × 512 потоків/SM = 28,672 потоки.
Другий сценарій
Тепер припустимо, що ми використовуємо той самий код, але з сіткою розміру 256×256 та блоками розміру 16×8.
Кількість потоків у сітці не змінюється, потоки виконують ті самі послідовності інструкцій, тільки змінюється поділ на блоки.
- Апаратне обмеження тепер дозволяє:
2048 потоків/SM ÷ 128 потоків/блок = 16 блоків/SM.
- Щодо обмежень пам'яті:
963 потоки/SM ÷ 128 потоків/блок ≈ 7 блоків/SM.
Висновок: Кількість потоків, що одночасно перебувають в SM GPU, становить:
56 SM × 128 потоків/блок × 7 блоків/SM = 50,176 потоки.
Просте змінення конфігурації сітки призвело до значного збільшення кількості потоків, що перебувають в SM.
Ця вправа, хоча й не враховує оптимізації NVCC, підкреслює важливість:
- Оптимізації використання регістрів і спільної пам'яті_;
- Правильного визначення розміру блоків для максимізації кількості потоків у SM та уникнення марнотратства ядер CUDA;
- Зберігання даних у структурах розміру 2ⁿ× 32 біти (наприклад, уникати структур типу
float3_ для запобігання витратам пам'яті)._
Посилання
- 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 II: From the Machine’s Perspective