Всім привіт! Досить багато часу пройшло від того моменту, коли я написав статтю «Знайомство із технологією CUDA», в якій було розглянуто основні принципи роботи з технологією, архітектуру графічних процесорів та інші цікаві речі.
Я займався написанням дипломного проекту, що також пов'язаний з CUDA та нейронними мережами, тому не мав змоги швидше написати цю статтю. Якщо вам буде цікаво, то я напишу окрему статтю про мій дипломний проект (після захисту, звичайно 😅).
Отож, сьогодні розглянемо як правильно встановити CUDA Toolkit, основні принципи написання CUDA-коду, та напишемо просту програму з використанням технології. Стаття буде складніша за попередню, тому я буду старатися доступно пояснити усі моменти. Поїхали!
Можливості CUDA
Розглянемо основні можливості технології.
- стандартна мова програмування C для паралельної розробки застосунків з використанням GPU;
- стандартні бібліотеки для швидкого перетворення Фур'є та базові пакети програм лінійної алгебри;
- спеціальний драйвер CUDA зі швидкою передачею даних між GPU та CPU;
- драйвер CUDA взаємодіє з OpenGL і DirectX;
- підтримка операційних систем Linux, Windows і Mac.
Набір інструментів CUDA
Інструмент CUDA Toolkit – середовище розробки для GPU з підтримкою CUDA, засноване на мові C. Це середовище включає:
- C-компілятор nvcc;
- бібліотеки FFT і BLAS для GPU;
- налагоджувач для GPU;
- драйвер CUDA Runtime;
- гід з програмування.
Встановлення CUDA
Для того, щоб встановити програмне забезпечення CUDA (далі буду писати просто ПЗ) переходимо на сайт Nvidia, обираємо платформу та операційну систему. Після цього розпочнеться завантаження останньої версії CUDA Toolkit (на час написання статті остання доступна версія – 9.0).
Версію ПЗ потрібно обирати в залежності від того, яка у вас відеокарта. Тобто яку версію CUDA підтримує ваш графічний процесор і чи підтримує взагалі. Перейшовши за цим посиланням ви можете побачити усі відеокарти, що підтримують технологію CUDA. Мені пощастило взяти в тимчасове користування GTX 950 Strix від Asus (для виконання дипломного проекту 😁), для якої я завантажив та встановив CUDA 8.0. Старіші версії ПЗ можна завантажити тут.
Під час встановлення ПЗ є момент, коли потрібно обрати або перезапис графічного драйвера новим, або ж залишити драйвер, що вже встановлений в системі. Я б рекомендував обрати перезапис, якщо у вас досить «старий» драйвер. В іншому разі, ви навіть не зможете завершити встановлення (так було в моєму випадку).
У цьому відео досить добре продемонстровано усі кроки встановлення ПЗ та запуск тестового прикладу:
Для перевірки правильності встановлення CUDA, ви можете відкрити приклад проекту Visual Studio deviceQuery, що знаходиться за шляхом C:\\ProgramData\\NVIDIA Corporation\\CUDA Samples\\версія_CUDA\\1_Utilities\\deviceQuery. Скомпілювавши та запустивши проект, ви побачите схоже вікно:
Як бачите, у вікні виводиться вся основна інформація про відеокарту, її обчислювальні можливості та інші цікаві показники. Усі інші приклади програм розміщені за шляхом C:\\ProgramData\\NVIDIA Corporation\\CUDA Samples\\версія_CUDA.
Програмна модель
Програму на CUDA можна логічно розділити на дві частини, перша частина (керуюча) виконується на CPU, друга частина (обчислювальна) виконується на GPU. CPU в термінології CUDA називається host, GPU – device. Про це я писав у попередній статті. Частина коду, яка повинна виконуватися на GPU, називається ядром (kernel), вона описується у вигляді функції, пізніше ми розглянемо приклади програм.
Розширення мови С, що входять у CUDA складаються з:
- специфікаторів функцій, які описують, де буде виконуватися функція і звідки вона може бути викликана;
- специфікаторів змінних, що задають тип пам'яті, який використовується для даної змінної;
- директив, що використовуються для запуску ядра і задає як дані, так і ієрархію потоків;
- вбудованих змінних, що містять інформацію про поточний потік;
- runtime, що містить в собі додаткові типи даних.
Програми зберігаються у файлі з розширенням .cu.
Всі програмні коди компілюються використовуючи CUDA API. Спочатку компілюється код, що призначений виключно для центрального процесора, а інший код, призначений для графічного процесора, компілюється в проміжну мову PTX (щось схоже до байт-коду в Java) для виявлення можливих помилок. Після чого, компілюється в «зрозумілу» для CPU/GPU мову.
Специфікатори (модифікатори)
Перед функціями в .cu файлі можуть стояти наступні модифікатори:
-
__device__– означає, що функція виконується тільки на відеокарті. З програми, що виконується на звичайному процесорі (хості), її викликати не можна; -
__global__– функція – початок вашого обчислювального ядра. Виконується на відеокарті, але запускається з хоста; -
__host__– виконується і запускається тільки з хоста (тобто звичайна функція C).
При цьому модифікатори __host__ і __device__ можуть бути використані разом (це означає, що відповідна функція може виконуватися як на GPU, так і на CPU – відповідний код для обох платформ буде автоматично згенерований компілятором). Модифікатори __global__ і __host__ не можуть бути використані разом. Детальніше можна почитати тут.
На функції, що виконуються на GPU, накладено певні обмеження: вони не можуть містити рекурсії, не можуть мати змінне число вхідних аргументів, не можуть містити статичні змінні, а також не можна взяти адресу такої функції.
Додані змінні
В мову додані наступні спеціальні змінні:
- gridDim – розмір grid (тип dim3);
- blockDim – розмір блоку (тип dim3);
- blockIdx – індекс поточного блоку в grid (тип uint3);
- threadIdx – індекс поточного потоку 😕 в блоці (тип uint3);
- warpSize – розмір warp (тип int).
Також додаються 1 / 2 / 3 / 4 – мірні вектори з базових типів: char1, char2, char3, char4, uchar1, uchar2, uchar3, uchar4, short1, short2, short3, short4, ushort1, ushort2, ushort3, ushort4, int1, int2, int3, int4, uint1, uint2, uint3, uint4 і так далі.
Директива виклику ядра
Для запуску ядра на GPU використовується наступна конструкція:
kernelName <<<Dg,Db,Ns,S>>> (args)
Тут kernelName – ім'я (адреса) відповідної __global__ функції, Dg – змінна (або значення) типу dim3, що задає розмірність grid (в блоках), Db – змінна (або значення) типу dim3, що задає розмірність блоку (в потоках), Ns – змінна (або значення) типу size_t, що задає додатковий обсяг спільної пам'яті, яка повинна бути динамічно виділена (до вже статично виділеної shared-пам'яті; параметр не є обов'язковим), S – змінна (або значення) типу cudaStream_t, що задає потік (потік CUDA), в якому має викликатися ядро, за замовчуванням використовується потік 0. Через args позначено аргументи виклику функції kernelName.
Як ви вже зрозуміли, множину потоків у блоці та блоків у grid можна задавати у вигляді 1 / 2 / 3 -мірних векторів. Розміри сітки та максимально можлива кількість потоків напряму залежать від відеокарти.
Також в мову С додана функція __syncthreads (детальніше тут), яка здійснює синхронізацію всіх потоків блоку. Управління з неї буде повернуто тільки тоді, коли всі потоки даного блоку викличуть цю функцію. Тобто, коли весь код, що йде перед цим викликом, вже виконано. Ця функція дуже зручна для організації безконфліктної роботи зі спільною пам'яттю.
Як потік знає над якими даними йому працювати?
Припустимо, що потрібно зробити деякі операції над зображенням (зберігається у змінній fox) розміром 400х400 пікселів. Зображення можна розділити на ділянки (блоки) 10х10 пікселів і для опрацювання кожної ділянки запустити окремий потік.
{@gist:https://gist.github.com/liashchynskyi/43d8db14e42944f753556faba5bf6140}
Оскільки параметри, що передаються у ядро однакові для всіх потоків, то кожен потік повинен сам «отримати дані для себе» 🤔. Щоб це зробити, потоку «потрібно розрахувати», в якому місці зображення він знаходиться.
{@gist:https://gist.github.com/liashchynskyi/32b21f39b404ec7e0ae1ac4c9e3d82bd}
ix та iy – координати, за допомогою яких можна отримати вихідні дані з масиву зображення. Як бачите, для цього якраз і застосовуються описані вище змінні.
Пишемо програму з використанням CUDA
Для прикладу, потрібно обчислити суму двох векторів розмірністю N. Для цього опишемо наступне ядро:
{@gist:https://gist.github.com/liashchynskyi/edadd1eae9ae86c74b3c0ca7e710abc4}
Таким чином, розпаралелювання буде виконано автоматично під час запуску ядра. У цій функції також використовується вбудована змінна threadIdx (дивися вище) та її поле x, що дозволяє отримати координату x потоку в блоці. Після чого, проводимо розрахунок кожного елемента вектора в окремому потоці.
{@gist:https://gist.github.com/liashchynskyi/a91d7cd48ee6e8e7cb37d234c21d9e6d}
Для виділення пам'яті на відеокарті використовується функція cudaMalloc, що має наступний прототип:
cudaError_t cudaMalloc( void** devPtr, size_t count ), де:
- devPtr – вказівник, який містить адресу виділеної пам'яті;
- count – розмір пам'яті, що виділяється.
Для копіювання даних в пам'ять відеокарти використовується функція cudaMemcpy, яка має наступний прототип:
cudaError_t cudaMemcpy(void* dst, const void* src,size_t count, enum cudaMemcpyKind kind), де:
- dst – вказівник, що містить адресу місця призначення копіювання (тобто destination, призначення);
- src – вказівник, що містить адресу джерела копіювання;
- count – розмір ресурсу, який необхідно скопіювати (в байтах);
- cudaMemcpyKind – enum, що вказує напрямок копіювання (може бути cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyHostToHost, cudaMemcpyDeviceToDevice).
Переходимо безпосередньо до виклику ядра.
{@gist:https://gist.github.com/liashchynskyi/273fbea272f1e87e3e9efcb80fdd4052}
Нам потрібно скопіювати результат розрахунку з відеопам'яті в пам'ять хоста. Але тут є одна особливість – асинхронне виконання, тобто, якщо після виклику ядра почав працювати наступний блок коду (мається на увазі код хоста), то це ще не означає, що GPU виконав розрахунки 🙄 Для завершення роботи заданої функції ядра необхідно використовувати засоби синхронізації, наприклад, events. Тому перед копіюванням результатів на хост виконується синхронізація потоків GPU. Якщо цього не робити, то ви можете отримати неправильні результати після виконання ядра 😮
{@gist:https://gist.github.com/liashchynskyi/1998b47805b806829cde30499d1b3d12}
Event створюється за допомогою функції cudaEventCreate, прототип якої має такий вигляд:
cudaError_t cudaEventCreate( cudaEvent_t* event ), де:
- event – вказівник для запису дескриптора події.
Запис event виконується за допомогою функції cudaEventRecord, прототип якої виглядає так:
cudaError_t cudaEventRecord( cudaEvent_t event, CUstream stream ), де:
- event – дескриптор event,
- stream – номер потоку, в якому працюємо (по-дефолту 0).
Синхронізація event виконується функцією cudaEventSynchronize. Функція очікує завершення роботи всіх потоків GPU і виклику заданого event, і тільки тоді передає управління керуючій програмі. Прототип функції виглядає так: cudaError_t cudaEventSynchronize( cudaEvent_t event ), де:
- event – дескриптор event, виклик якого очікується.
Тепер залишається вивести результат на екран і вивільнити ресурси.
{@gist:https://gist.github.com/liashchynskyi/a5deca2b9160ea3c4a6f9621403e02ed}
Компілюємо:
$ nvcc vector.cu
І отримуємо помилку:
Ми забули під'єднати бібліотеку iostream 😏. Додаємо її у початок файлу та компілюємо знову. Після чого запускаємо:
Весь код виглядає наступним чином:
{@gist:https://gist.github.com/liashchynskyi/62c1e9e4c8ed5baecd899bdb3b7523af}
Оцінка витраченого часу на виконання обчислень
Для того, щоб порахувати скільки часу зайняло виконання обрахунків на GPU, можна використати наступний фрагмент коду:
{@gist:https://gist.github.com/liashchynskyi/c4874211b5f55942897216cb644be7e7}
Для CPU:
{@gist:https://gist.github.com/liashchynskyi/8a5ca813853504c617835c344f6df093}
Запускаємо програму і бачимо результат (я збільшив кількість знаків після коми, щоб краще було видно результат).
Щось не дуже, правда? 🤔 До речі, у мене старенький AMD Athlon 64 x2 Dual Core Processor 4800+. Тут якраз добре видно те, про що я писав у попередній статті:
Не варто виконувати на GPU надто легкі завдання... Але для демонстрації роботи – можна 😆
Збільшуємо розмір вектора до 1_000_000_00 (мільярд) і запускаємо знову.
Як бачите, ми досягнули досить хорошого результату шляхом збільшення кількості елементів у векторі 🤗
Оптимізація коду
Трохи розповім про те, як не зробити програму повільнішою на GPU (як я вже писав у попередній статті: «Сповільнити програму на GPU набагато простіше, ніж прискорити»).
- старайтеся використовувати глобальну пам'ять не дуже часто, оскільки це найповільніший тип пам'яті;
- під час роботи з shared-пам'яттю уникайте конфліктів синхронізації;
- не використовуйте код, що містить багато розгалужень;
- старайтеся використовувати якнайменше пам'яті, що значно зменшить затримки.
Детальніше про типи пам'яті можна почитати тут і тут.
Висновки
На сьогодні ідея застосування графічних процесорів та технології CUDA є перспективною та популярною. CUDA широко застосовується для вирішення задач обробки зображень, машинного навчання та інженерних розрахунків, часто дозволяючи порівняно недорого та без громіздкого обладнання забезпечити задовільну продуктивність.
Звичайно, існують й інші технології, які також дозволяють виконувати обчислення на графічних процесорах, зокрема, OpenCL. Проте, як показують експерименти та дослідження, функція-ядро OpenCL виконується повільніше на 13-63%, а під час end-to-end тестування на 16-67% повільніше.
У статті я старався розповісти про основні аспекти написання програм з використанням CUDA. Оскільки тема є доволі широкою та містить багато матеріалу, то неможливо розповісти всі деталі у двох статтях 😉. Тому я підготував деякі матеріали для самостійного вивчення.
На цій ноті я хочу завершити серію статей по технології CUDA. Надіюся, що вам було цікаво читати 😊. Швидкого вам коду та до зустрічі!
Вихідний код на GitHub.
Ще немає коментарів