Технологія паралельного програмування CUDA

-- українська Народна Казка

1. Введення

У цій статті ми поговоримо про одну з популярних технологій високопродуктивних обчислень, яка використовує GPU (graphics processing units). Спочатку цей клас пристроїв розроблявся для обробки графіки. Техніка вдосконалювалася, GPU нарощували продуктивність і в якийсь момент виявилося, що GPU можна успішно використовувати не тільки для задач комп'ютерної графіки, але і як математичний співпроцесор для CPU машини, отримуючи при цьому суттєвий приріст у продуктивності.

CUDA працює тільки з пристроями виробництва NVIDIA, але це не біда, крім CUDA існують інші аналогічні технології, наприклад OpenCL і AMD FireStream, але їх опис виходить за рамки цієї статті.

2. Апаратна частина

Архітектура GPU побудована інакше, ніж у універсальних CPU [2], і в ній спочатку закладено певну спеціалізацію. Завдання комп'ютерної графіки передбачають незалежну паралельну обробку даних, і GPU спочатку призначений для паралельних обчислень. Він спроектований так, щоб виконувати велику кількість тредів (елементарних паралельних процесів). GPU містить багато відносно простих арифметико-логічних пристроїв (Рис.2), які об'єднані у групи, та реалізує модель паралельного обчислювача над загальною пам'яттю, але з деякими особливостями.

GPU орієнтований виконання програм із великим обсягом обчислень (розпаралелювання за даними типу SIMD). Пам'ять GPU має ієрархічну структуру та оптимізована під максимальну пропускну здатність. Замість системи кешів CPU та складних арифметико-логічних схем (АЛУ), GPU має багато спрощених АЛУ, що мають спільну пам'ять. Це допомагає підвищити продуктивність у обчислювальних задачах.ускладнює програмування. Для досягнення найкращого прискорення необхідно продумувати стратегії доступу до пам'яті та враховувати апаратні особливості.

GPU є масив потокових процесорів (Streaming Processor Array), що складається з кластерів текстурних процесорів (Texture Processor Clusters, TPC). TPC складається з набору мультипроцесорів (SM – Streaming Multi-processor), кожен із яких містить кілька потокових процесорів (SP – Streaming Processors) чи ядер. Набір ядер кожного мультипроцесора працює за принципом SIMD (одинний потік команд, безліч потоків даних). На рис.3 представлена ​​загальна апаратна схема роботи мультипроцесорів (SM) GPU.

У порівнянні з універсальним CPU, конструкція GPU накладає ряд додаткових обмежень для програмування, вони залежать від конкретної моделі. Для NVIDIA Quadro FX1700 це такі особливості: GPU не підтримує рекурсію та обчислення з подвійною (double) точністю (можлива лише одинарна точність – float).

3. Програмна частина

Пакет інструментів для розробника та бібліотеки CUDA можна завантажити із сайту NVIDIA. На момент написання цієї статті доступні пакети для Windows, Linux і MacOSX. Ми будемо використовувати ОС Linux, але майже все, що викладається далі, справедливо і для інших ОС, можливо з невеликими виправленнями. Опис процедури інсталяції та налаштування середовища програмування залишимо за рамками статті та перейдемо до компілятора та програм.

Існують прив'язки CUDA для різних мов програмування, повний список можна переглянути на сайті NVIDIA. Для реалізації тестових прикладів ми будемо використовувати CUDA C -- адаптований для програмування GPU діалект мови C++.

Пакет CUDA-розробника для Linux містить компілятор nvccгенерує код для роботи з GPU.

Кожен пристрій CUDA має спеціальну властивість - рівень обчислювальних можливостей (compute capability), яке визначає набір доступних пристрою можливостей із усього функціоналу CUDA. Для NV >'-arch sm_11' .

4. Програмування на CUDA C

На лістингу 2 представлена ​​найпростіша програма CUDA C, вона не робить нічого. Ядро, яке описане функцією kern, виконується на GPU, ця функція має специфікатор __global__, який говорить компілятору, що функція викликається з CPU та виконується на GPU. Для функцій є інші специфікатори (див.табл.1).

специфікатор виконується

__host__	host	host
__global__	device	host
__device__	device	device

На функції, що виконуються на GPU накладається ряд обмежень: вони не можуть містити рекурсію, не можуть мати змінну кількість аргументів і не можуть мати static змінні в собі.

Ядро на GPU запускається рядком kern >>() , потрійні кутові дужки це специфічний CUDA C синтаксис, значення цих дужках визначають кількість копій ядра, що запускаються на GPU, тобто. кількість та конфігурацію паралельних процесів GPU. В даному випадку ми запускаємо один процес - один блок з одним тредом, всередині основного стриму і очікуємо на його завершення.

Запуск ядра породжує процеси (треди) на GPU відповідно до заданих параметрів (рис.4). Все безліч процесів, що породжуються запуском ядра, в термінології CUDA називається грид (gr & gt; блоків (block), блок з тредів (thread). Тред це елементарний паралельний процес. Треди в блоках і блоки в гриді можуть бути представлені у вигляді одно-, дво- або тривимірної решітки.

ДляFX1700 максимальний розмір гриду 65535 x 65535 блоків. Максимально можливі значення індексів номера треда в блоці 512 x 512 x 64, але кількість тредів в блоці не повинна перевищувати 512.

Ядра можна запускати асинхронно, тоді кілька гридів виконуватимуться на GPU паралельно.

При конфігуруванні топології процесів потрібно враховувати апаратні особливості. Тут слід запровадити поняття варпа. Варп (warp) - це група тредів, розмір варпа для FX1700 - 32 треди. Полуварп (half-warp) – половина тредів варпа. Усі треди одного варпа виконуються одночасно синхронно (SIMD) своєму мультипроцессоре. При доступі треду до основної пам'яті GPU її частини можуть кешуватися у локальній пам'яті даного мультипроцесора. Якщо всі дані, які потрібні тредам напівварпа будуть знаходитися в цьому кеші, то це може підвищити продуктивність.

GPU має складно організовану пам'ять (рис.3), крім основної (або глобальної) пам'яті кожен мультипроцесор має власну пам'ять. Програмна модель пам'яті CUDA представлена ​​на рис.5, далі ми розглянемо типи пам'яті CUDA та методи роботи з нею.

4.1. Основна (глобальна) пам'ять

4.2. Про компонування тредів та блоків

На лістингу 4 представлений приклад наступної організації тредів. Розмір блоку 8x8x8, що дорівнює 512 тредів на один блок, розмір грида 8х32 блоку, таким чином загальна кількість паралельних процесів 131072=8x8x8x8x32.

Для того, щоб задати три- або двомірну решітку тредів та/або блоків необхідно визначити змінну розмірів типу dim3 і передати її як параметр (у кутових дужках) при запуску ядра. Крім dim3 CUDA містить інші векторні типи char2, int3, float3 та ін.

Для визначення номера треду та розмірів блоку та гриду в CUDAІснують такі системні змінні (таблиця 2).

тип ім'я призначення

dim3	gridDim	розмір гриду
dim3	blockDim	розмір блоку
uint3	blockIdx	номер блоку у гриді
uint3	threadIdx	номер треду в блоці
int	warpSize	кількість тредів у варпі

4.3. Константна пам'ять

4.4. Пам'ять, що розділяється

На лістингу 6 представлений приклад роботи з пам'яттю, що розділяється, яка оголошується всередині ядра з модифікатором __shared__ . Кожен тред записує в "свою" комірку пам'яті, що розділяється дані, потім змінює дані в "чужій" комірці, після цього результат переписується в основну пам'ять GPU, а потім і в пам'ять CPU.

Оскільки треди конкурують між собою за загальний ресурс (пам'ять), виникає можливість конфліктів і неоднозначностей між тредами. З'являється необхідність у механізмі синхронізації тредів. Для цього можна використовувати функцію __syncthreads() . Цей метод працює для тредів одного блоку. Треди блоку, що виконали __syncthreads(), очікують інші треди цього блоку, поки всі виконають __syncthreads() . У цьому випадку якщо один з тредів за логікою програми ніколи не дійде до __syncthreads(), то програма зависає (дідлок).

4.5. Текстурна пам'ять

Розглянемо найпростіший приклад роботи з текстурною пам'яттю (листинг 7).

4.6. Атомарні операції

На лістингу 8 представлений простий приклад роботи з атомарними операціями. У пам'яті виділяється масив із цілих чисел, туди записуються нульові значення, далі кожен процес збільшує на 1 всі осередки цього масиву. Якщо замість атомарної складання atomicAdd() використовувати звичайний інкремент, то результатможе бути некоректним.

4.7. Паралельний запуск кількох ядер на GPU

За замовчуванням у CUDA-програмі є лише одна черга, що виконується або нуль-стрім. У представленій вище програмі ми за допомогою cudaStreamCreate() створюємо ще два стрими stream0, stream1 і запускаємо два ядра у новостворені стрими, додавши відповідний параметр при запуску ядра. Далі ми чекаємо поки обидва ядра завершать роботу, використовуючи cudaStreamSynchronize() , переписуємо результат у пам'ять CPU і закриваємо обидва відпрацьовані стрими за допомогою cudaStreamDestroy() .

Перед тим як асинхронно копіювати дані з пам'яті хоста на девайс та назад необхідно заблокувати цю область пам'яті на хості. Блокування області пам'яті (page-locked) означає, що ОС хоста забороняється переміщати дані із цієї області фізичної пам'яті, тобто. до цієї області пам'яті не можна застосовувати свопінг тощо. Цей метод потрібно використовувати обережно, тому що ми втрачаємо всі переваги віртуальної пам'яті, і повинні стежити за тим, щоб у нас залишалося достатньо фізичної пам'яті для роботи.

На закінчення треба відзначити ще один момент, для спільного використання двома стримами одного і того ж буфера блокованої пам'яті необхідно використовувати прапор cudaHostAllocPortable або блокованим цей буфер буде вважати тільки один стрим.

4.9. Використання пам'яті CPU

На початку ми перевіряємо чи підтримує пристрій відображення пам'яті. Далі вмикаємо цю функціональність за допомогою cudaSetDeviceFlags(cudaDeviceMapHost) . Виділяємо блоковану пам'ять CPU за допомогою cudaHostAlloc(), при цьому для вхідного буфера можна вказати прапор cudaHostAllocWriteCombined, який встановлює особливий режим кешування цієї області пам'яті для CPU, це прискорить читання цих данихдля GPU, але при цьому швидкість роботи з ними CPU значно падає. Відображаємо виділені буфери пам'яті CPU у пам'ять GPU за допомогою cudaHostGetDevicePointer() та запускаємо ядро. Перед тим як друкувати результат необхідно викликати cudaThreadSynchronize() , щоб буфери пам'яті GPU і CPU синхронізувалися. Наприкінці роботи звільняємо ресурси.

4.10. Операції з матрицями

На лістингу 12 представлений приклад використання cuBLAS, ця програма генерує дві матриці чисел однакового розміру, виконує з ними дії - додавання, транспонування та множення. Спочатку ми виділяємо пам'ять для даних звичайним способом, далі ініціалізуємо cuBLAS за допомогою cublasCreate() , копіюємо дані host на device за допомогою cublasSetMatrix() . Після цього виконуємо операції з матрицями: додавання - cublasSgeam() , множення - cublasSgemm() , домноження на скаляр - cublasSscal() . Кожна з цих функцій cuBLAS самостійно формує та запускає на device потрібну кількість процесів, при цьому кожна така функція може виконувати відразу кілька операцій, наприклад – транспонування та множення. Після завершення обчислень результат копіюємо назад на host за допомогою cublasGetMatrix(), друкуємо результат і звільняємо ресурси - cublasDestroy().

4.11. Оцінка витраченого на обчислення часу

На лістингу 13 наведено чат коду, яка замірює час виконання програми. На початку створюємо тимчасові мітки start, stop з допомогою cudaEventCreate() , далі фіксуємо час старту з допомогою cudaEventRecord() , виконуємо обчислення, фіксуємо час завершення. Оскільки деякі операції можуть виконуватися асинхронно Викликаємо cudaEventSynchronize() , щоб мітка записалася коректно. Далі обчислюємо різницю в часі за допомогою cudaEventElapsedTime() та друкуєморезультат.

---