Пристрої перетворення персональних комп'ютерів на маленькі суперкомп'ютери відомі досить давно. Ще у 80-х роках минулого століття на ринку пропонувалися так звані трансп'ютери, які вставлялися у поширені тоді слоти розширення ISA. Перший час їхня продуктивність у відповідних завданнях вражала, але потім зростання швидкодії універсальних процесорів прискорилося, вони посилили свої позиції в паралельних обчисленнях, і сенсу в трансп'ютерах не залишилося. Хоча подібні пристрої існують і зараз – це різноманітні спеціалізовані прискорювачі. Але найчастіше сфера їх застосування вузького та особливого поширення такі прискорювачі не набули.

Але останнім часом естафета паралельних обчислень перейшла до масового ринку, так чи інакше пов'язаного із тривимірними іграми. Універсальні пристрої з багатоядерними процесорами для паралельних векторних обчислень, що використовуються в 3D-графіці, досягають високої пікової продуктивності, яка універсальним процесорам не під силу. Звичайно, максимальна швидкість досягається лише в ряді зручних завдань і має деякі обмеження, але такі пристрої вже почали досить широко застосовувати у сферах, для яких вони спочатку не призначалися. Відмінним прикладом такого паралельного процесора є процесор Cell, розроблений альянсом Sony-Toshiba-IBM і застосовується в ігровій приставці Sony PlayStation 3, а також всі сучасні відеокарти від лідерів ринку - компаній Nvidia і AMD.

Cell ми сьогодні чіпати не будемо, хоч він і з'явився раніше і є універсальним процесором з додатковими векторними можливостями, сьогодні не про нього. Для 3D-відеоприскорювачів ще кілька років тому з'явилися перші технології неграфічних розрахунків загального призначення GPGPU (General-Purpose computation on GPUs). Адже сучасні відеочіпи містять сотні математичних виконавчих блоків, і ця міць може використовуватися для значного прискорення безлічі обчислювально-інтенсивних додатків. І нинішні покоління GPU мають досить гнучку архітектуру, що разом з високорівневими мовами програмування і програмно-апаратними архітектурами, подібними до цієї статті, розкриває ці можливості і робить їх значно доступнішими.

На створення GPCPU розробників спонукала появу досить швидких та гнучких шейдерних програм, які здатні виконувати сучасні відеочіпи. Розробники задумали зробити так, щоб GPU розраховували не тільки зображення в 3D-додатках, але й застосовувалися в інших паралельних розрахунках. У GPGPU для цього використовувалися графічні API: OpenGL та Direct3D, коли дані до відеочіпа передавалися у вигляді текстур, а розрахункові програми завантажувалися у вигляді шейдерів. Недоліками такого методу є порівняно висока складність програмування, низька швидкість обміну даними між CPU та GPU та інші обмеження, про які ми поговоримо далі.

Обчислення на GPU розвивалися та розвиваються дуже швидко. І надалі, два основних виробника відеочіпів, Nvidia та AMD, розробили та анонсували відповідні платформи під назвою CUDA (Compute Unified Device Architecture) та CTM (Close To Metal або AMD Stream Computing), відповідно. На відміну від попередніх моделей програмування GPU, ці були виконані з урахуванням прямого доступу до апаратних можливостей відеокарт. Платформи не сумісні між собою, CUDA – це розширення мови програмування C, а CTM – віртуальна машина, що виконує асемблерний код. Натомість обидві платформи ліквідували деякі з важливих обмежень попередніх моделей GPGPU, які використовують традиційний графічний конвеєр та відповідні інтерфейси Direct3D чи OpenGL.

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

Саме тому компанія Nvidia випустила платформу CUDA C-подібну мову програмування зі своїм компілятором та бібліотеками для обчислень на GPU. Звичайно, написання оптимального коду для відеочіпів зовсім не таке просте і це завдання потребує тривалої ручної роботи, але CUDA якраз і розкриває всі можливості та дає програмісту більший контроль над апаратними можливостями GPU. Важливо, що підтримка Nvidia CUDA є у чіпів G8x, G9x і GT2xx, які застосовуються у відеокартах Geforce серій 8, 9 і 200, які дуже поширені. В даний час випущено фінальну версію CUDA 2.0, в якій з'явилися деякі нові можливості, наприклад, підтримка розрахунків з подвійною точністю. CUDA доступна на 32-бітових та 64-бітних операційних системах Linux, Windows та MacOS X.

Різниця між CPU та GPU у паралельних розрахунках

Зростання частот універсальних процесорів уперся у фізичні обмеження та високе енергоспоживання, і збільшення їхньої продуктивності все частіше відбувається за рахунок розміщення декількох ядер в одному чіпі. Процесори, що продаються зараз, містять лише до чотирьох ядер (подальше зростання не буде швидким) і вони призначені для звичайних додатків, використовують MIMD - множинний потік команд і даних. Кожне ядро ​​працює окремо від інших, виконуючи різні інструкції щодо різних процесів.

Спеціалізовані векторні можливості (SSE2 та SSE3) для чотирьохкомпонентних (одинарна точність обчислень з плаваючою точкою) та двокомпонентних (подвійна точність) векторів з'явилися в універсальних процесорах через збільшені вимоги графічних додатків, в першу чергу. Саме тому для певних завдань застосування GPU вигідніше, адже вони спочатку зроблено для них.

Наприклад, у відеочіпах Nvidia основний блок - це мультипроцесор з вісьма-десятьма ядрами і сотнями ALU в цілому, кількома тисячами регістрів і невеликою кількістю спільної пам'яті, що розділяється. Крім того, відеокарта містить швидку глобальну пам'ять із доступом до неї всіх мультипроцесорів, локальну пам'ять у кожному мультипроцесорі, а також спеціальну пам'ять для констант.

Найголовніше ці кілька ядер мультипроцесора в GPU є SIMD (одинаковий потік команд, безліч потоків даних) ядрами. І ці ядра виконують одні й самі інструкції одночасно, такий стиль програмування є звичайним для графічних алгоритмів та багатьох наукових завдань, але потребує специфічного програмування. Проте такий підхід дозволяє збільшити кількість виконавчих блоків за рахунок їх спрощення.

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

Розробники CPU намагаються домогтися виконання якомога більшої кількості інструкцій паралельно для збільшення продуктивності. Для цього починаючи з процесорів Intel Pentium з'явилося суперскалярне виконання, що забезпечує виконання двох інструкцій за такт, а Pentium Pro відзначився позачерговим виконанням інструкцій. Але в паралельного виконання послідовного потоку інструкцій є певні базові обмеження та збільшенням кількості виконавчих блоків кратного збільшення швидкості не досягти.

У відеочіпів робота проста і розпаралелена спочатку. Відеочіп приймає на вході групу полігонів, проводить усі необхідні операції і на виході видає пікселі. Обробка полігонів та пікселів незалежна, їх можна обробляти паралельно, окремо один від одного. Тому, через початково паралельну організацію роботи в GPU використовується велика кількість виконавчих блоків, які легко завантажити, на відміну від послідовного потоку інструкцій для CPU. Крім того, сучасні GPU можуть виконувати більше однієї інструкції за такт (dual issue). Так, архітектура Tesla в деяких умовах запускає виконання операції MAD+MUL або MAD+SFU одночасно.

GPU відрізняється від CPU ще й за принципами доступу до пам'яті. У GPU він пов'язаний і легко передбачуваний - якщо з пам'яті читається тексель текстури, через деякий час прийде час і для сусідніх текселів. Та й при записі те саме - піксель записується у фреймбуфер, і через кілька тактів записуватиметься розташований поруч із ним. Тому організація пам'яті відрізняється від тієї, що використовується CPU. І відеочіпу, на відміну від універсальних процесорів, просто не потрібна кеш-пам'ять великого розміру, а для текстур потрібно лише кілька (до 128-256 у нинішніх GPU) кілобайт.

Та й сама по собі робота з пам'яттю у GPU та CPU дещо відрізняється. Так, не всі центральні процесори мають вбудовані контролери пам'яті, а всі GPU зазвичай мають по кілька контролерів, аж до восьми 64-бітних каналів у чіпі Nvidia GT200. Крім того, на відеокартах застосовується швидша пам'ять, і в результаті відеочіп доступна в рази велика пропускна здатність пам'яті, що також дуже важливо для паралельних розрахунків, що оперують з величезними потоками даних.

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

Для відмінності в кешуванні. Універсальні центральні процесори використовують кеш-пам'ять для збільшення продуктивності за рахунок зниження затримок доступу до пам'яті, а GPU використовують кеш або загальну пам'ять для збільшення смуги пропускання. CPU знижують затримки доступу до пам'яті за допомогою кеш-пам'яті великого розміру та передбачення розгалужень коду. Ці апаратні частини займають більшу частину площі чіпа і споживають багато енергії. Відеочіпи обходять проблему затримок доступу до пам'яті за допомогою одночасного виконання тисяч потоків - в той час, коли один з потоків очікує даних із пам'яті, відеочіп може виконувати обчислення іншого потоку без очікування та затримок.

Є безліч відмінностей і підтримки багатопоточності. CPU виконує 1-2 потоки обчислень на одне процесорне ядро, а відеочіпи можуть підтримувати до 1024 потоків на кожен мультипроцесор, яких у чіпі кілька штук. І якщо перемикання з одного потоку на інший для CPU коштує сотні тактів, то GPU перемикає кілька потоків за один такт.

Крім того, центральні процесори використовують SIMD (одна інструкція виконується над численними даними) блоки векторних обчислень, а відеочіпи застосовують SIMT (одна інструкція і кілька потоків) для скалярної обробки потоків. SIMT не вимагає, щоб розробник перетворював дані на вектори, і допускає довільні розгалуження в потоках.

Коротко можна сказати, що на відміну від сучасних універсальних CPU відеочіпи призначені для паралельних обчислень з великою кількістю арифметичних операцій. І значно більше транзисторів GPU працює за прямим призначенням - обробці масивів даних, а чи не управляє виконанням (flow control) нечисленних послідовних обчислювальних потоків. Це схема того, скільки місця в CPU та GPU займає різноманітна логіка:

У результаті основою для ефективного використання потужності GPU в наукових та інших неграфічних розрахунках є розпаралелювання алгоритмів на сотні виконавчих блоків, що є у відеочіпах. Наприклад, безліч додатків з молекулярного моделювання добре пристосовано для розрахунків на відеочіпах, вони вимагають високих обчислювальних потужностей і тому зручні для паралельних обчислень. А використання кількох GPU дає ще більше обчислювальних потужностей для вирішення таких завдань.

Виконання розрахунків на GPU показує відмінні результати алгоритмах, що використовують паралельну обробку даних. Тобто, коли ту саму послідовність математичних операцій застосовують до великого обсягу даних. При цьому кращі результати досягаються, якщо відношення числа арифметичних інструкцій до звернень до пам'яті досить велике. Це пред'являє менші вимоги до управління виконанням (flow control), а висока щільність математики і великий обсяг даних скасовує необхідність великих кешах, як і CPU.

В результаті всіх описаних вище відмінностей теоретична продуктивність відеочіпів значно перевершує продуктивність CPU. Компанія Nvidia наводить такий графік зростання продуктивності CPU та GPU за останні кілька років:

Звичайно, ці дані не без частки лукавства. Адже на CPU набагато простіше на практиці досягти теоретичних цифр, та й цифри наведені для одинарної точності у випадку GPU, і для подвійної в разі CPU. У будь-якому випадку, для частини паралельних завдань одинарної точності вистачає, а різниця у швидкості між універсальними та графічними процесорами дуже велика, і тому шкурка коштує вичинки.

Перші спроби застосування розрахунків на GPU

Відеочіпи в паралельних математичних розрахунках намагалися використати досить давно. Найперші спроби такого застосування були вкрай примітивними та обмежувалися використанням деяких апаратних функцій, таких як растеризація та Z-буферизація. Але в нинішньому столітті з появою шейдерів почали прискорювати обчислення матриць. У 2003 році на SIGGRAPH окрема секція була виділена під обчислення на GPU і отримала назву GPGPU (General-Purpose computation on GPU) - універсальні обчислення на GPU).

Найбільш відомий BrookGPU - компілятор потокової мови програмування Brook, створений для виконання неграфічних обчислень на GPU. До появи розробники, які використовують можливості відеочипів для обчислень, вибирали одне із двох поширених API: Direct3D чи OpenGL. Це серйозно обмежувало застосування GPU, адже в 3D графіку використовуються шейдери та текстури, про які фахівці з паралельного програмування знати не зобов'язані, вони використовують потоки та ядра. Brook зміг допомогти полегшити їх завдання. Ці потокові розширення до мови C, розроблені в Стендфордському університеті, приховували від програмістів тривимірний API і представляли відеочіп у вигляді паралельного співпроцесора. Компілятор обробляв файл.br із кодом C++ та розширеннями, виробляючи код, прив'язаний до бібліотеки з підтримкою DirectX, OpenGL або x86.

Звичайно, у Brook було безліч недоліків, на яких ми зупинялися, і про які ще докладніше поговоримо далі. Але навіть його поява викликала значний приплив уваги тих же Nvidia і ATI до ініціативи обчислень на GPU, оскільки розвиток цих можливостей серйозно змінило ринок надалі, відкривши цілий новий його сектор - паралельні обчислювачі на основі відеочіпів.

В подальшому деякі дослідники з проекту Brook влилися в команду розробників Nvidia, щоб представити програмно-апаратну стратегію паралельних обчислень, відкривши нову частку ринку. І головною перевагою цієї ініціативи Nvidia стало те, що розробники чудово знають всі можливості своїх GPU до дрібниць, і у використанні графічного API немає необхідності, а працювати з апаратним забезпеченням можна за допомогою драйвера. Результатом зусиль цієї команди стала Nvidia CUDA (Compute Unified Device Architecture) – нова програмно-апаратна архітектура для паралельних обчислень на Nvidia GPU, якій присвячена ця стаття.

Області застосування паралельних розрахунків на GPU

Щоб зрозуміти, які переваги приносить перенесення розрахунків на відеочіпи, наведемо середні цифри, отримані дослідниками у всьому світі. У середньому, при перенесенні обчислень на GPU, у багатьох завданнях досягається прискорення у 5-30 разів, порівняно із швидкими універсальними процесорами. Найбільші цифри (близько 100-кратного прискорення і навіть більше!) досягаються на коді, який дуже добре підходить для розрахунків за допомогою блоків SSE, але цілком зручний для GPU.

Це лише деякі приклади прискорення синтетичного коду на GPU проти SSE-векторизованого коду на CPU (за даними Nvidia):

• Флуоресцентна мікроскопія: 12 х;
• Молекулярна динаміка (non-bonded force calc): 8-16x;
• Електростатика (пряме та багаторівневе підсумовування Кулону): 40-120x та 7x.

А це табличка, яку дуже любить Nvidia, показуючи її на всіх презентаціях, на якій ми докладніше зупинимося у другій частині статті, присвяченій конкретним прикладам практичних застосувань CUDA обчислень:

Як бачите, цифри дуже привабливі, особливо вражають 100-150-кратні прирости. У наступній статті, присвяченій CUDA, ми детально розберемо деякі з цих цифр. А зараз перерахуємо основні додатки, в яких зараз застосовуються обчислення на GPU: аналіз та обробка зображень та сигналів, симуляція фізики, обчислювальна математика, обчислювальна біологія, фінансові розрахунки, бази даних, динаміка газів та рідин, криптографія, адаптивна променева терапія, астрономія звуку, біоінформатика, біологічні симуляції, комп'ютерний зір, аналіз даних (data mining), цифрове кіно та телебачення, електромагнітні симуляції, геоінформаційні системи, військові застосування, гірське планування, молекулярна динаміка, магнітно-резонансна томографія (MRI), нейромережі, океанографічні дослідження, фізика частинок, симуляція згортання молекул білка, квантова хімія, трасування променів, візуалізація, радари, гідродинамічне моделювання (reservoir simulation), штучний інтелект, аналіз супутникових даних, сейсмічна розвідка, хірургія, ультразвук, відеоконференції.

Подробиці про багато застосування можна знайти на сайті компанії Nvidia в розділі . Як бачите, список досить великий, але це ще не все! Його можна продовжувати, і, напевно, можна припустити, що в майбутньому будуть знайдені й інші області застосування паралельних розрахунків на відеочіпах, про які ми поки не здогадуємося.

Можливості Nvidia CUDA

Технологія CUDA - це програмно-апаратна обчислювальна архітектура Nvidia, заснована на розширенні мови Сі, яка дає можливість організації доступу до набору інструкцій графічного прискорювача та управління його пам'яттю при організації паралельних обчислень. CUDA допомагає реалізовувати алгоритми, здійсненні на графічних процесорах відеоприскорювачів Geforce восьмого покоління та старші (серії Geforce 8, Geforce 9, Geforce 200), а також Quadro та Tesla.

Хоча трудомісткість програмування GPU за допомогою CUDA досить велика, вона нижча, ніж із ранніми GPGPU рішеннями. Такі програми вимагають розбиття програми між декількома мультипроцесорами подібно до MPI програмування, але без поділу даних, які зберігаються у спільній відеопам'яті. І так як CUDA програмування для кожного мультипроцесора подібно до OpenMP програмування, воно вимагає хорошого розуміння організації пам'яті. Але, звичайно ж, складність розробки та перенесення на CUDA сильно залежить від програми.

Набір розробників містить безліч прикладів коду і добре документований. Процес навчання вимагатиме близько двох-чотирьох тижнів для тих, хто вже знайомий з OpenMP та MPI. В основі API лежить розширена мова Сі, а для трансляції коду цієї мови до складу CUDA SDK входить компілятор командного рядка nvcc, створений на основі відкритого компілятора Open64.

Перерахуємо основні характеристики CUDA:

• уніфіковане програмно-апаратне рішення для паралельних обчислень на відеочіпах Nvidia;
• великий набір рішень, що підтримуються, від мобільних до мультичіпових.
• стандартна мова програмування Сі;
• стандартні бібліотеки чисельного аналізу FFT (швидке перетворення Фур'є) та BLAS (лінійна алгебра);
• оптимізований обмін даними між CPU та GPU;
• взаємодія з графічними API OpenGL та DirectX;
• підтримка 32- та 64-бітних операційних систем: Windows XP, Windows Vista, Linux та MacOS X;
• можливість розробки низький рівень.

Щодо підтримки операційних систем потрібно додати, що офіційно підтримуються всі основні дистрибутиви Linux (Red Hat Enterprise Linux 3.x/4.x/5.x, SUSE Linux 10.x), але, судячи з даних ентузіастів, CUDA чудово працює і на інших збірках: Fedora Core, Ubuntu, Gentoo та ін.

Середовище розробки CUDA (CUDA Toolkit) включає:

• компілятор nvcc;
• бібліотеки FFT та BLAS;
• профільувальник;
• налагоджувач gdb для GPU;
• CUDA runtime драйвер у комплекті стандартних драйверів Nvidia
• посібник з програмування;
• CUDA Developer SDK (початковий код, утиліти та документація).

У прикладах вихідного коду: паралельне бітонне сортування (bitonic sort), транспонування матриць, паралельне префіксне підсумовування великих масивів, згортка зображень, дискретне вейвлет-перетворення, приклад взаємодії з OpenGL і Direct3D, використання бібліотек CUBLAS і CUFFT, обчислення ціни опціону Шоулза, біноміальна модель, метод Монте-Карло), паралельний генератор випадкових чисел Mersenne Twister, обчислення гістограми великого масиву, шумозаглушення, фільтр Собеля (знаходження кордонів).

Переваги та обмеження CUDA

З погляду програміста графічний конвеєр є набором стадій обробки. Блок геометрії генерує трикутники, а блок растеризації - пікселі, що відображаються на моніторі. Традиційна модель програмування GPGPU виглядає так:

Щоб перенести обчислення на GPU у рамках такої моделі, потрібний спеціальний підхід. Навіть поелементне складання двох векторів вимагатиме відтворення фігури на екрані або у позаекранний буфер. Фігура розтеризується, колір кожного пікселя обчислюється за програмою (піксельному шейдеру). Програма зчитує вхідні дані з текстур кожного пікселя, складає їх і записує у вихідний буфер. І всі ці численні операції потрібні у тому, що у звичайній мові програмування записується одним оператором!

Тому застосування GPGPU для обчислень загального призначення має обмеження у вигляді занадто великої складності навчання розробників. Та й інших обмежень достатньо, адже піксельний шейдер - це лише формула залежності підсумкового кольору пікселя від його координати, а мова піксельних шейдерів - мова запису цих формул з Сі-подібним синтаксисом. Ранні методи GPGPU є хитрим трюком, що дозволяє використовувати потужність GPU, але без будь-якої зручності. Дані там представлені зображеннями (текстурами), а алгоритм – процесом растеризації. Потрібно особливо відзначити і дуже специфічну модель пам'яті та виконання.

Програмно-апаратна архітектура для обчислень на GPU компанії Nvidia відрізняється від попередніх моделей GPGPU тим, що дозволяє писати програми для GPU справжньою мовою Сі зі стандартним синтаксисом, покажчиками та необхідністю в мінімумі розширень для доступу до обчислювальних ресурсів відеочіпів. CUDA не залежить від графічних API, і має деякі особливості, призначені спеціально для обчислень загального призначення.

Переваги CUDA перед традиційним підходом до обчислень GPGPU:

• інтерфейс програмування додатків CUDA заснований на стандартній мові програмування Сі з розширеннями, що полегшує процес вивчення та впровадження архітектури CUDA;
• CUDA забезпечує доступ до розділяється між потоками пам'яті розміром 16 Кб на мультипроцесор, яка може бути використана для організації кешу з широкою смугою пропускання, в порівнянні з текстурними вибірками;
• більш ефективна передача даних між системною та відеопам'яттю
• відсутність необхідності в графічних API з надмірністю та накладними витратами;
• лінійна адресація пам'яті, і gather і scatter; можливість запису за довільними адресами;
• апаратна підтримка цілих і бітових операцій.

Основні обмеження CUDA:

• відсутність підтримки рекурсії для виконуваних функцій;
• мінімальна ширина блоку 32 потоку;
• закрита архітектура CUDA, що належить Nvidia.

Слабкими місцями програмування за допомогою попередніх методів GPGPU є те, що ці методи не використовують блоки виконання вершинних шейдерів у попередніх неуніфікованих архітектурах, дані зберігаються у текстурах, а виводяться у позаекранний буфер, а багатопрохідні алгоритми використовують піксельні шейдерні блоки. До обмежень GPGPU можна включити: недостатньо ефективне використання апаратних можливостей, обмеження смугою пропускання пам'яті, відсутність операції scatter (тільки gather), обов'язкове використання графічного API.

Основні переваги CUDA у порівнянні з попередніми методами GPGPU випливають з того, що ця архітектура спроектована для ефективного використання неграфічних обчислень на GPU та використовує мову програмування C, не вимагаючи перенесення алгоритмів у зручний для концепції графічного конвеєра вигляд. CUDA пропонує новий шлях обчислень на GPU, що не використовує графічні API, що пропонує довільний доступ до пам'яті (scatter або gather). Така архітектура позбавлена ​​недоліків GPGPU та використовує всі виконавчі блоки, а також розширює можливості за рахунок цілісної математики та операцій бітового зсуву.

Крім того, CUDA відкриває деякі апаратні можливості, недоступні з графічних API, такі як пам'ять, що розділяється. Це пам'ять невеликого об'єму (16 кілобайт на мультипроцесор), до якої мають доступ блоки потоків. Вона дозволяє кешувати дані, що найчастіше використовуються, і може забезпечити більш високу швидкість, порівняно з використанням текстурних вибірок для цього завдання. Що, своєю чергою, знижує чутливість до пропускної спроможності паралельних алгоритмів у багатьох додатках. Наприклад, це корисно для лінійної алгебри, швидкого перетворення Фур'є та фільтрів обробки зображень.

Зручніший у CUDA та доступ до пам'яті. Програмний код у графічних API виводить дані у вигляді 32-х значень з плаваючою точкою одинарної точності (RGBA значення одночасно у вісім render target) у заздалегідь визначені області, а CUDA підтримує scatter запис - необмежену кількість записів за будь-якою адресою. Такі переваги уможливлюють виконання на GPU деяких алгоритмів, які неможливо ефективно реалізувати за допомогою методів GPGPU, заснованих на графічних API.

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

Історія розвитку CUDA

Розробка CUDA була анонсована разом із чіпом G80 у листопаді 2006, а реліз публічної бета-версії CUDA SDK відбувся у лютому 2007 року. Версія 1.0 вийшла у червні 2007 року під запуск у продаж рішень Tesla, заснованих на чіпі G80, та призначених для ринку високопродуктивних обчислень. Потім, наприкінці року, вийшла бета-версія CUDA 1.1, яка, незважаючи на малозначне збільшення номера версії, ввела досить багато нового.

З CUDA 1.1, що з'явився, можна відзначити включення CUDA-функціональності в звичайні відеодрайвери Nvidia. Це означало, що у вимогах до будь-якої програми CUDA достатньо було вказати відеокарту серії Geforce 8 і вище, а також мінімальну версію драйверів 169.xx. Це дуже важливо для розробників, при дотриманні цих умов програми CUDA будуть працювати у будь-якого користувача. Також було додано асинхронне виконання разом із копіюванням даних (тільки для чіпів G84, G86, G92 і вище), асинхронне пересилання даних у відеопам'ять, атомарні операції доступу до пам'яті, підтримка 64-бітових версій Windows та можливість мультичіпової роботи CUDA у режимі SLI.

На даний момент актуальною є версія для рішень на основі GT200 CUDA 2.0, що вийшла разом з лінійкою Geforce GTX 200. Бета-версія була випущена ще навесні 2008 року. У другій версії з'явилися: підтримка обчислень подвійної точності (апаратна підтримка тільки у GT200), підтримується Windows Vista (32 і 64-бітні версії) і Mac OS X, додані засоби налагодження та профілювання, підтримуються 3D текстури, оптимізована пересилання даних.

Що стосується обчислень з подвійною точністю, то їх швидкість на поточному апаратному поколінні нижче одинарної точності в кілька разів. Причини розглянуті в нашій. Реалізація в GT200 цієї підтримки полягає в тому, що блоки FP32 не використовуються для отримання результату в чотири рази меншому темпі, для підтримки FP64 обчислень у Nvidia вирішили зробити виділені обчислювальні блоки. І в GT200 їх удесятеро менше, ніж блоків FP32 (по одному блоку подвійної точності на кожен мультипроцесор).

Реально продуктивність може бути навіть ще меншою, тому що архітектура оптимізована для 32-бітного читання з пам'яті та регістрів, крім того, подвійна точність не потрібна в графічних додатках, і в GT200 вона зроблена швидше, щоб просто була. Та й сучасні чотириядерні процесори показують не набагато меншу реальну продуктивність. Але навіть у 10 разів повільніше, ніж одинарна точність, така підтримка корисна для схем зі змішаною точністю. Одна з найпоширеніших технік - отримати спочатку наближені результати в одинарній точності, а потім їх уточнити в подвійній. Тепер це можна зробити безпосередньо на відеокарті, без пересилання проміжних даних до CPU.

Ще одна корисна особливість CUDA 2.0 не має відношення до GPU, як не дивно. Просто тепер можна компілювати код CUDA у високоефективний багатопотоковий SSE код для швидкого виконання на центральному процесорі. Тобто тепер ця можливість годиться не тільки для налагодження, але й реального використання на системах без відеокарти Nvidia. Адже використання CUDA у звичайному коді стримується тим, що відеокарти Nvidia хоч і найпопулярніші серед виділених відеорішень, але є не у всіх системах. І до версії 2.0 у таких випадках довелося б робити два різні коди: для CUDA та окремо для CPU. А тепер можна виконувати будь-яку програму CUDA на CPU з високою ефективністю, нехай і з меншою швидкістю, ніж на відеочіпах.

Рішення з підтримкою Nvidia CUDA

Всі відеокарти, що мають підтримку CUDA, можуть допомогти прискорити більшість вимогливих завдань, починаючи від аудіо- та відеообробки, і закінчуючи медициною та науковими дослідженнями. Єдине реальне обмеження полягає в тому, що багато CUDA програми вимагають мінімум 256 мегабайт відеопам'яті, і це одна з найважливіших технічних характеристик для CUDA-додатків.

Актуальний список продуктів, що підтримують CUDA, можна отримати на . На момент написання статті розрахунки CUDA підтримували всі продукти серій Geforce 200, Geforce 9 і Geforce 8, у тому числі і мобільні продукти, починаючи з Geforce 8400M, а також і чіпсети Geforce 8100, 8200 і 8300. Також підтримкою CUDA всі Tesla: S1070, C1060, C870, D870 та S870.

Особливо зазначимо, що разом з новими відеокартами Geforce GTX 260 та 280, були анонсовані та відповідні рішення для високопродуктивних обчислень: Tesla C1060 та S1070 (подано на фото вище), які будуть доступні для придбання восени цього року. GPU в них застосований той же - GT200, в C1060 він один, S1070 - чотири. Натомість, на відміну від ігрових рішень, у них використовується по чотири гігабайти пам'яті на кожен чіп. З мінусів хіба що менша частота пам'яті та ПСП, ніж у ігрових карток, що забезпечує по 102 гігабайт/с на чіп.

Склад Nvidia CUDA

CUDA включає два API: високого рівня (CUDA Runtime API) та низького (CUDA Driver API), хоча в одній програмі одночасне використання обох неможливе, потрібно використовувати один або інший. Високорівневий працює «зверху» низькорівневого, всі дзвінки runtime транслюються у прості інструкції, що обробляються низькорівневим Driver API. Але навіть «високорівневий» API передбачає знання про пристрій та роботу відеочіпів Nvidia, надто високого рівня абстракції там немає.

Є й ще один рівень, навіть вищий за дві бібліотеки:

CUBLAS CUDA варіант BLAS (Basic Linear Algebra Subprograms), призначений для обчислень завдань лінійної алгебри і використовує прямий доступ до ресурсів GPU;

CUFFT CUDA варіант бібліотеки Fast Fourier Transform для розрахунку швидкого перетворення Фур'є, що широко використовується при обробці сигналів. Підтримуються такі типи перетворень: complex-complex (C2C), real-complex (R2C) та complex-real (C2R).

Розглянемо ці бібліотеки докладніше. CUBLAS - це перекладені мовою CUDA стандартні алгоритми лінійної алгебри, на даний момент підтримується лише певний набір основних функцій CUBLAS. Бібліотеку дуже легко використовувати: потрібно створити матрицю та векторні об'єкти у пам'яті відеокарти, заповнити їх даними, викликати необхідні функції CUBLAS, та завантажити результати з відеопам'яті назад у системну. CUBLAS містить спеціальні функції для створення та знищення об'єктів у пам'яті GPU, а також для читання та запису даних у цю пам'ять. Функції BLAS, що підтримуються: рівні 1, 2 і 3 для дійсних чисел, рівень 1 CGEMM для комплексних. Рівень 1 – це векторно-векторні операції, рівень 2 – векторно-матричні операції, рівень 3 – матрично-матричні операції.

CUFFT CUDA варіант функції швидкого перетворення Фур'є широко використовуваної і дуже важливою при аналізі сигналів, фільтрації і т.п. CUFFT надає простий інтерфейс для ефективного обчислення FFT на відеочіпах виробництва Nvidia без необхідності розробки власного варіанту FFT для GPU. CUDA варіант FFT підтримує 1D, 2D, та 3D перетворення комплексних та дійсних даних, пакетне виконання для декількох 1D трансформацій у паралелі, розміри 2D та 3D трансформацій можуть бути в межах , для 1D підтримується розмір до 8 мільйонів елементів.

Основи створення програм на CUDA

Для розуміння подальшого тексту слід розумітися на базових архітектурних особливостях відеочіпів Nvidia. GPU складається з кількох кластерів текстурних блоків (Texture Processing Cluster). Кожен кластер складається з укрупненого блоку текстурних вибірок та двох-трьох потокових мультипроцесорів, кожен з яких складається з восьми обчислювальних пристроїв та двох суперфункціональних блоків. Всі інструкції виконуються за принципом SIMD, коли одна інструкція застосовується до всіх потоків у warp (термін з текстильної промисловості, в CUDA це група з 32 потоків - мінімальний обсяг даних, оброблюваних мультипроцесорами). Цей спосіб виконання назвали SIMT (single instruction multiple threads - одна інструкція і багато потоків).

Кожен із мультипроцесорів має певні ресурси. Так, є спеціальна пам'ять, що розділяється, обсягом 16 кілобайт на мультипроцесор. Але це не кеш, так як програміст може використовувати її для будь-яких потреб, подібно до Local Store в SPU процесорів Cell. Ця пам'ять, що розділяється, дозволяє обмінюватися інформацією між потоками одного блоку. Важливо, що всі потоки одного блоку завжди виконуються тим самим мультипроцесором. А потоки з різних блоків обмінюватись даними не можуть, і треба пам'ятати це обмеження. Пам'ять, що розділяється, часто буває корисною, крім тих випадків, коли кілька потоків звертаються до одного банку пам'яті. Мультипроцесори можуть звертатися і до відеопам'яті, але з великими затримками та гіршою пропускною здатністю. Для прискорення доступу та зниження частоти звернення до відеопам'яті, мультипроцесори мають по 8 кілобайт кешу на константи і текстурні дані.

Мультипроцесор використовує 8192-16384 (для G8x/G9x і GT2xx, відповідно) регістра, загальні всім потоків всіх блоків, виконуваних у ньому. Максимальне число блоків на один мультипроцесор для G8x/G9x дорівнює восьми, а число warp 24 (768 потоків на один мультипроцесор). Усього топові відеокарти серій Geforce 8 та 9 можуть обробляти до 12288 потоків одночасно. Geforce GTX 280 на основі GT200 пропонує до 1024 потоків на мультипроцесор, в ньому є 10 кластерів по три мультипроцесори, що обробляють до 30720 потоків. Знання цих обмежень дає змогу оптимізувати алгоритми під доступні ресурси.

Першим кроком при перенесенні існуючого додатка на CUDA є його профільування та визначення ділянок коду, що є «пляшковим шийкою», що гальмує роботу. Якщо серед таких ділянок є придатні для швидкого паралельного виконання, ці функції переносяться на Cі розширення CUDA для виконання на GPU. Програма компілюється за допомогою компілятора Nvidia, що генерує код і для CPU, і для GPU. При виконанні програми центральний процесор виконує свої порції коду, а GPU виконує CUDA код з найважчими паралельними обчисленнями. Ця частина, призначена для GPU називається ядром (kernel). У ядрі визначаються операції, які будуть виконані над даними.

Відеочіп отримує ядро ​​та створює копії для кожного елемента даних. Ці копії називають потоками (thread). Потік містить лічильник, регістри та стан. Для великих обсягів даних, таких як обробка зображень, запускаються мільйони потоків. Потоки виконуються групами по 32 штуки, званими warp"и. Warp"ам призначається виконання на певних потокових мультипроцесорах. Кожен мультипроцесор складається з восьми ядер - потокових процесорів, які виконують одну інструкцію MAD за один такт. Для виконання одного 32-потокового warp'а потрібно чотири такти роботи мультипроцесора (мова про частоту shader domain, яка дорівнює 1.5 ГГц і вище).

Мультипроцесор не є традиційним багатоядерним процесором, він відмінно пристосований для багатопоточності, підтримуючи до 32 warp"ів одночасно. Кожен такт апаратне забезпечення вибирає, який з warp"ів виконувати, і перемикається від одного до іншого без втрат у тактах. Якщо проводити аналогію з центральним процесором, це схоже на одночасне виконання 32 програм та перемикання між ними кожен такт без втрат на перемикання контексту. Реально ядра CPU підтримують одноразове виконання однієї програми та перемикаються на інші із затримкою в сотні тактів.

Модель програмування CUDA

Повторимося, що CUDA використовує паралельну модель обчислень, коли кожен із SIMD процесорів виконує ту ж інструкцію над різними елементами даних паралельно. GPU є обчислювальним пристроєм, співпроцесором (device) для центрального процесора (host), що має власну пам'ять і обробляє паралельно велику кількість потоків. Ядром (kernel) називається функція для GPU, що виконується потоками (аналогія з 3D графіки – шейдер).

Ми говорили вище, що відеочіп відрізняється від CPU тим, що може обробляти одночасно десятки тисяч потоків, що зазвичай для графіки, яка добре розпаралелюється. Кожен потік скалярний, не вимагає упаковки даних у 4-компонентні вектори, що зручніше для більшості завдань. Кількість логічних потоків і блоків потоків перевищує кількість фізичних виконавчих пристроїв, що дає хорошу масштабованість для всього модельного ряду рішень компанії.

Модель програмування в CUDA передбачає групування потоків. Потоки об'єднуються в блоки потоків (thread block) - одномірні або двовимірні сітки потоків, що взаємодіють між собою за допомогою пам'яті, що розділяється, і точок синхронізації. Програма (ядро, kernel) виконується над сіткою (grid) блоків потоків (thread blocks), див. малюнок нижче. Одночасно виконується одна сітка. Кожен блок може бути одно-, дво- або тривимірним формою, і може складатися з 512 потоків на поточному апаратному забезпеченні.

Блоки потоків виконуються у вигляді невеликих груп, званих варп (warp), розмір яких 32 потоку. Це мінімальний обсяг даних, які можуть бути оброблені в мультипроцесорах. І оскільки це не завжди зручно, CUDA дозволяє працювати і з блоками, що містять від 64 до 512 потоків.

Угруповання блоків у сітки дозволяє уникнути обмежень і застосувати ядро ​​до більшого числа потоків за один виклик. Це допомагає і за масштабування. Якщо GPU недостатньо ресурсів, він буде виконувати блоки послідовно. У протилежному випадку блоки можуть виконуватися паралельно, що важливо для оптимального розподілу роботи на відеочіпах різного рівня, починаючи від мобільних та інтегрованих.

Модель пам'яті CUDA

Модель пам'яті в CUDA відрізняється можливістю побайтної адресації, підтримкою як gather, і scatter. Доступно досить велику кількість регістрів на кожен потоковий процесор, до 1024 штук. Доступ до них дуже швидкий, зберігати в них можна 32-бітові цілі чи числа з плаваючою точкою.

Кожен потік має доступ до наступних типів пам'яті:

Глобальна пам'ятьНайбільший обсяг пам'яті, доступний для всіх мультипроцесорів на відеочіпі, розмір становить від 256 мегабайт до 1.5 гігабайт на поточних рішеннях (і до 4 Гбайт на Tesla). Має високу пропускну здатність, більше 100 гігабайт/с для топових рішень Nvidia, але дуже великими затримками в кілька сотень тактів. Не кешується, підтримує узагальнені вказівки load і store, і звичайні покажчики на згадку.

Локальна пам'ятьЦе невеликий обсяг пам'яті, до якого має доступ тільки один потоковий процесор. Вона відносно повільна така ж, як і глобальна.

Пам'ять, що розділяєтьсяЦе 16-кілобайтний (у відеочіпах нинішньої архітектури) блок пам'яті із загальним доступом для всіх потокових процесорів в мультипроцесорі. Ця пам'ять дуже швидка, така сама, як регістри. Вона забезпечує взаємодію потоків, управляється розробником безпосередньо та має низькі затримки. Переваги пам'яті, що розділяється: використання у вигляді керованого програмістом кеша першого рівня, зниження затримок при доступі виконавчих блоків (ALU) до даних, скорочення кількості звернень до глобальної пам'яті.

Пам'ять констант- Область пам'яті об'ємом 64 кілобайти (те ж - для нинішніх GPU), доступна тільки для читання всіма мультипроцесорами. Вона кешується по 8 кілобайт на кожен мультипроцесор. Досить повільна - затримка кілька сотень тактів за відсутності потрібних даних у кеші.

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

Природно, що глобальна, локальна, текстурна та пам'ять констант - це фізично та сама пам'ять, відома як локальна відеопам'ять відеокарти. Їх відмінності у різних алгоритмах кешування та моделях доступу. Центральний процесор може оновлювати та вимагати лише зовнішню пам'ять: глобальну, константну та текстурну.

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

Типовий, але не обов'язковий шаблон розв'язання задач:

• завдання розбивається на підзавдання;
• вхідні дані діляться на блоки, які вміщуються в пам'ять, що розділяється;
• кожний блок обробляється блоком потоків;
• підблок підвантажується в пам'ять, що розділяється, з глобальної;
• над даними в пам'яті, що розділяється, проводяться відповідні обчислення;
• результати копіюються з пам'яті, що розділяється, назад у глобальну.

Середовище програмування

До складу CUDA входять runtime бібліотеки:

• загальна частина, що надає вбудовані векторні типи та підмножини викликів RTL, що підтримуються на CPU та GPU;
• CPU-компонента, для керування одним або декількома GPU;
• GPU-компонента, що надає специфічні функції для GPU.

Основний процес CUDA працює на універсальному процесорі (host), він запускає кілька копій процесів kernel на відеокарті. Код для CPU робить наступне: ініціалізує GPU, розподіляє пам'ять на відеокарті та системі, копіює константи у пам'ять відеокарти, запускає кілька копій процесів kernel на відеокарті, копіює отриманий результат із відеопам'яті, звільняє пам'ять та завершує роботу.

Як приклад для розуміння наведемо CPU код для складання векторів, представлений у CUDA:

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

Програми на CUDA можуть взаємодіяти з графічними API: для рендерингу даних, згенерованих у програмі, зчитування результатів рендерингу та їх обробки засобами CUDA (наприклад, при реалізації фільтрів постобробки). Для цього ресурси графічних API можуть відображатися (з отриманням адреси ресурсу) в простір глобальної пам'яті CUDA. Підтримуються такі типи ресурсів графічних API: Buffer Objects (PBO/VBO) у OpenGL, вершинні буфери та текстури (2D, 3D та кубічні карти) Direct3D9.

Стадії компіляції CUDA-додатку:

Файли вихідного коду на CUDA C компілюються за допомогою програми NVCC, яка є оболонкою над іншими інструментами, і викликає їх: cudacc, g++, cl та ін. Сі, та об'єктний код PTX для відеочіпа. Виконувані файли з кодом на CUDA обов'язково вимагають наявності бібліотек CUDA runtime library (cudart) і CUDA core library (cuda).

Оптимізація програм на CUDA

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

Загалом, при оптимізації програми CUDA потрібно постаратися досягти оптимального балансу між розміром і кількістю блоків. Більша кількість потоків у блоці знизить вплив затримок пам'яті, але знизить доступне число регістрів. Крім того, блок з 512 потоків неефективний, сама Nvidia рекомендує використовувати блоки по 128 або 256 потоків як компромісне значення для досягнення оптимальних затримок і кількості регістрів.

Серед основних моментів оптимізації програм CUDA: якомога більш активне використання пам'яті, що розділяється, так як вона значно швидше глобальної відеопам'яті відеокарти; Операції з читання та запису з глобальної пам'яті повинні бути об'єднані (coalesced) по можливості. Для цього потрібно використовувати спеціальні типи даних для читання та запису відразу по 32/64/128 біта даних однією операцією. Якщо операції читання важко поєднати, можна спробувати використовувати текстурні вибірки.

Висновки

Представлена ​​компанією Nvidia програмно-апаратна архітектура для розрахунків на відеочіпах CUDA добре підходить для вирішення широкого кола завдань із високим паралелізмом. CUDA працює на великій кількості відеочіпів Nvidia, і покращує модель програмування GPU, значно спрощуючи її і додаючи велику кількість можливостей, таких як пам'ять, що розділяється, можливість синхронізації потоків, обчислення з подвійною точністю і цілочисленні операції.

CUDA - це доступна кожному розробнику ПЗ технологія, її може використовувати будь-який програміст, який знає мову Сі. Прийде лише звикнути до іншої парадигми програмування, властивої паралельним обчисленням. Але якщо алгоритм у принципі добре розпаралелюється, то вивчення та витрати часу на програмування на CUDA повернуться у багаторазовому розмірі.

Цілком імовірно, що через поширення відеокарт у світі, розвиток паралельних обчислень на GPU сильно вплине на індустрію високопродуктивних обчислень. Ці можливості вже викликали великий інтерес у наукових колах, та й не лише у них. Адже потенційні можливості прискорення алгоритмів, що добре піддаються розпаралелюванню (на доступному апаратному забезпеченні, що не менш важливо) відразу в десятки разів бувають не так часто.

Універсальні процесори розвиваються досить повільно, вони не мають таких стрибків продуктивності. По суті, нехай це і звучить занадто голосно, всі, хто потребує швидких обчислювачів, тепер можуть отримати недорогий персональний суперкомп'ютер на своєму столі, іноді навіть не вкладаючи додаткових коштів, оскільки відеокарти Nvidia широко поширені. Не говорячи вже про підвищення ефективності в термінах GFLOPS/$ і GFLOPS/Вт, які так подобаються виробникам GPU.

Майбутнє безлічі обчислень явно за паралельними алгоритмами, майже всі нові рішення та ініціативи направлені в цей бік. Поки що, втім, розвиток нових парадигм знаходиться на початковому етапі, доводиться вручну створювати потоки та планувати доступ до пам'яті, що ускладнює завдання порівняно зі звичним програмуванням. Але технологія CUDA зробила крок у правильному напрямку і в ній явно проглядається успішне рішення, особливо якщо Nvidia вдасться переконати якомога розробників у його користі та перспективах.

Але, звісно, ​​GPU не замінять CPU. У їхньому нинішньому вигляді вони й не призначені для цього. Зараз відеочіпи рухаються поступово в бік CPU, стаючи все більш універсальними (розрахунки з плаваючою точкою одинарної і подвійної точності, цілочисленні обчислення), так і CPU стають все більш «паралельними», обзаводячись великою кількістю ядер, технологіями багатопоточності, не кажучи про появу блоків SIMD та проектів гетерогенних процесорів. Швидше за все, GPU та CPU у майбутньому просто зіллються. Відомо, що багато компаній, у тому числі Intel та AMD, працюють над подібними проектами. І неважливо, чи будуть GPU поглинені CPU, чи навпаки.

У статті ми переважно говорили про переваги CUDA. Але є і ложечка дьогтю. Один з нечисленних недоліків CUDA – слабка переносимість. Ця архітектура працює тільки на відеочіпах цієї компанії, та ще й не на всіх, а починаючи з серії Geforce 8 та 9 та відповідних Quadro та Tesla. Так, таких рішень у світі дуже багато, Nvidia наводить цифру в 90 мільйонів CUDA-сумісних відеочіпів. Це просто чудово, але конкуренти пропонують свої рішення, відмінні від CUDA. Так, у AMD є Stream Computing, у Intel у майбутньому буде Ct.

Яка з технологій переможе, стане поширеною і проживе довше за інші - покаже лише час. Але у CUDA є непогані шанси, тому що в порівнянні з Stream Computing, наприклад, вона представляє більш розвинене та зручне для використання середовище програмування звичайною мовою Сі. Можливо, у визначенні допоможе третя сторона, випустивши спільне рішення. Наприклад, в наступному оновленні DirectX під версією 11, компанією Microsoft обіцяні обчислювальні шейдери, які можуть стати якимось усередненим рішенням, що влаштовує всіх, або багатьох.

Судячи з попередніх даних, цей новий тип шейдерів запозичує багато моделей CUDA. І програмуючи у цьому середовищі вже зараз, можна отримати переваги одразу та необхідні навички для майбутнього. З точки зору високопродуктивних обчислень, DirectX також має явний недолік у вигляді поганої переносимості, оскільки цей API обмежений платформою Windows. Втім, розробляється ще один стандарт - відкрита мультиплатформна ініціатива OpenCL, яка підтримується більшістю компаній, серед яких Nvidia, AMD, Intel, IBM та багато інших.

Не забувайте, що в наступній статті з CUDA на вас чекає дослідження конкретних практичних застосувань наукових та інших неграфічних обчислень, виконаних розробниками з різних куточків нашої планети за допомогою Nvidia CUDA.

Ядра CUDA - умовне позначення скалярних обчислювальних блоківу відео-чіпах NVidia, починаючи з G 80 (GeForce 8 xxx, Tesla C-D-S870, FX4/5600 , 360M). Самі чіпи є похідними архітектури. До речі, тому компанія NVidiaтак охоче взялася за розробку власних процесорів Tegra Series, заснованих також на RISCархітектури. Досвід роботи з цими архітектурами дуже великий.

CUDAядро містить у собі один один векторнийі один скалярнийюніт, які за один такт виконують по одній векторній і по одній скалярній операції, передаючи обчислення іншому мультипроцесору, або для подальшої обробки. Масив із сотень і тисяч таких ядер являє собою значну обчислювальну потужність і може виконувати різні завдання в залежності від вимог, за наявності певного софта підтримує . Застосуванняможе бути різноманітним: декодування відеопотоку, прискорення 2D/3D графіки, хмарні обчислення, спеціалізовані математичні аналізи тощо.

Досить часто, об'єднані професійні карти NVidia Teslaі NVidia Quadroє кістяком сучасних суперкомп'ютерів.

CUDA— ядра не зазнали якихось значних змін із часів G 80, але збільшується їх кількість (спільно з іншими блоками ROP, Texture Units& etc) та ефективність паралельних взаємодій один з одним (поліпшуються модулі Giga Thread).

Наприклад:

GeForce

GTX 460 - 336 CUDA ядер

GTX 580 - 512 CUDA ядер

8800GTX - 128 CUDA ядер

Від кількості потокових процесорів ( CUDA), практично пропорційно збільшується продуктивність у шейдерних обчисленнях (при рівномірному збільшенні кількості та інших елементів).

Починаючи з чіпа GK110(NVidia GeForce GTX 680) - CUDAядра тепер не мають подвоєну частоту, а загальну з рештою всіх блоків чіпа. Натомість було збільшено їх кількість приблизно в три разиу порівнянні з попереднім поколінням G110.

Нова технологія — як еволюційний вигляд, що знову виник. Дивне створення, несхоже на численних старожилів. Місцями незграбне, подекуди смішне. І спочатку його нові якості здаються ну ніяк не придатними для цього обжитого та стабільного світу.

Проте минає небагато часу, і виявляється, що новачок бігає швидше, стрибає вище та взагалі сильніше. І мух він лопає більше за його сусідів-ретроградів. І ось тоді ці сусіди починають розуміти, що сваритися з цим колишнім незграбним не варто. Краще дружити з ним, а ще краще організувати симбіоз. Дивишся, і мух перепаде більше.

Технологія GPGPU (General-Purpose Graphics Processing Units — графічний процесор загального призначення) тривалий час існувала лише в теоретичних викладках мозкових академіків. А як інакше? Запропонувати кардинально змінити обчислювальний процес, що склався за десятиліття, довіривши розрахунок його паралельних гілок відеокарті, — на це тільки теоретики і здатні.

Логотип технології CUDA нагадує про те, що вона виросла в надрах
3D графіки.

Але довго припадати пилом на сторінках університетських журналів технологія GPGPU не збиралася. Розпушивши пір'я своїх найкращих якостей, вона привернула до себе увагу виробників. Так з'явилася CUDA — реалізація GPGPU на графічних процесорах GeForce виробництва компанії nVidia.

Завдяки CUDA технології GPGPU стали мейнстрімом. І нині тільки недалекоглядний і покритий товстим шаром лінощів розробник систем програмування не заявляє про підтримку своїм продуктом CUDA. IT-видання вважали за честь викласти подробиці технології у численних пухких науково-популярних статтях, а конкуренти терміново посідали за лекала та кроскомпілятори, щоб розробити щось подібне.

Публічне визнання — це мрія не тільки старлеток-початківців, а й технологій, що знову зародилися. І CUDA пощастило. Вона на слуху, про неї говорять та пишуть.

Ось тільки пишуть так, ніби продовжують обговорювати GPGPU у товстих наукових журналах. Закидають читача купою термінів типу "grid", "SIMD", "warp", "хост", "текстурна та константна пам'ять". Занурюють його по саму маківку в схеми організації графічних процесорів nVidia, ведуть звивистими стежками паралельних алгоритмів і (найсильніший хід) показують довгі лістинги коду мовою Сі. В результаті виходить, що на вході статті ми маємо свіжого і палючого бажанням зрозуміти CUDA читача, а на виході - того ж читача, але з головою, що розпухла, заповненою кашею з фактів, схем, коду, алгоритмів і термінів.

А тим часом мета будь-якої технології зробити наше життя простіше. І CUDA чудово з цим справляється. Результати її роботи — саме це переконає будь-якого скептика краще за сотню схем та алгоритмів.

Далеко не скрізь

CUDA підтримується високопродуктивними суперкомп'ютерами
nVidia Tesla.

І все ж таки перш, ніж поглянути на результати праць CUDA на терені полегшення життя пересічного користувача, варто усвідомити всі її обмеження. Точно як із джинном: будь-яке бажання, але одне. У CUDA теж є свої ахіллесові п'яти. Одна з них – обмеження платформ, на яких вона може працювати.

Перелік відеокарт виробництва nVidia, що підтримують CUDA, представлений у спеціальному списку, що називається CUDA Enabled Products. Список дуже значний, але легко класифікується. У підтримці CUDA не відмовляють:

Моделі nVidia GeForce 8-й, 9-й, 100-й, 200-й та 400-й серій з мінімумом 256 мегабайт відеопам'яті на борту. Підтримка поширюється як на карти для настільних систем, так і на мобільні рішення.

Переважна більшість настільних та мобільних відеокарт nVidia Quadro.

Всі рішення нетбучного ряду nvidia ION.

Високопродуктивні HPC (High Performance Computing) і суперкомп'ютерні рішення nVidia Tesla, що використовуються як для персональних обчислень, так і для організації кластерних систем, що масштабуються.

Тому, перш ніж використовувати програмні продукти на базі CUDA, варто звіритися з цим переліком обраних.

Крім самої відеокарти для підтримки CUDA потрібен відповідний драйвер. Саме він є сполучною ланкою між центральним та графічним процесором, виконуючи роль своєрідного програмного інтерфейсу для доступу коду та даних програми до багатоядерної скарбниці GPU. Щоб напевно не помилитися, nVidia рекомендує відвідати сторінку драйверів та отримати найсвіжішу версію.

...але сам процес

Як працює CUDA? Як пояснити складний процес паралельних обчислень на особливій апаратній архітектурі GPU так, щоб не занурити читача в безодню специфічних термінів?

Можна спробувати це зробити, уявивши, як центральний процесор виконує програму в симбіозі з графічним процесором.

Архітектурно центральний процесор (CPU) та його графічний зібрат (GPU) влаштовані по-різному. Якщо проводити аналогію зі світом автопрому, то CPU — універсал, із тих, що називають «сарай». Виглядає легковим авто, але при цьому (з погляду розробників) «і швець, і жнець, і на дуді гравець». Виконує роль маленької вантажівки, автобуса та гіпертрофованого хечбека одночасно. Універсал, коротше. Циліндр-ядер у нього небагато, але вони «тягнуть» практично будь-які завдання, а велика кеш-пам'ять здатна розмістити купу даних.

А ось GPU – це спорткар. Функція одна: доставити пілота на фініш якнайшвидше. Тому ніякої великої пам'яті-багажника, жодних зайвих місць. Натомість циліндрів-ядер у сотні разів більше, ніж у CPU.

Завдяки CUDA розробникам програм GPGPU не потрібно вникати у складності програми.
рування під такі графічні движки, як DirectX і OpenGL

На відміну від центрального процесора, здатного вирішувати будь-яке завдання, у тому числі і графічну, але з усередненою продуктивністю, графічний процесор адаптований на високошвидкісне рішення одного завдання: перетворення куп полігонів на вході в купу пікселів на виході. Причому це завдання можна вирішувати паралельно на сотнях щодо простих обчислювальних ядер у складі GPU.

То який же може бути тандем з універсалу та спорткара? Робота CUDA відбувається приблизно так: програма виконується на CPU доти, доки в ній з'являється ділянка коду, яку можна виконати паралельно. Тоді замість того, щоб він повільно виконувався на двох (та нехай навіть і восьми) ядрах найкрутішого CPU, його передають на сотні ядер GPU. При цьому час виконання цієї ділянки скорочується в рази, а отже, скорочується час виконання всієї програми.

Технологічно для програміста нічого не змінюється. Код CUDA-програм пишеться мовою Сі. Точніше, на особливому діалекті «З with streams» (Сі з потоками). Розроблене в Стенфорді, це розширення мови Сі отримало назву Brook. Як інтерфейс, що передає Brook-код на GPU, виступає драйвер відеокарти, що підтримує CUDA. Він організує весь процес обробки цієї ділянки програми так, що для програміста GPU виглядає як співпроцесор CPU. Дуже схоже використання математичного співпроцесора на зорі персональних комп'ютерів. З появою Brook, відеокарт із підтримкою CUDA та драйверів для них будь-який програміст став здатний у своїх програмах звертатися до GPU. Адже раніше цим шаманством володів вузьке коло обраних, які роками відточують техніку програмування під графічні двигуни DirectX або OpenGL.

У бочку цього пафосного меду - дифірамбів CUDA - варто покласти ложку дьогтю, тобто обмежень. Не будь-яке завдання, яке потрібно запрограмувати, підходить для вирішення за допомогою CUDA. Домогтися прискорення вирішення рутинних офісних завдань не вдасться, а ось довірити CUDA обрахунок поведінки тисячі однотипних бійців у World of Warcraft – будь ласка. Але це завдання, висмоктане з пальця. Розглянемо приклади того, що CUDA вже дуже ефективно вирішує.

Праці праведні

CUDA – дуже прагматична технологія. Реалізувавши її підтримку у своїх відеокартах, компанія nVidia дуже справедливо розраховувала на те, що прапор CUDA буде підхоплений безліччю ентузіастів як в університетському середовищі, так і в комерції. Так і сталося. Проекти на базі CUDA живуть та приносять користь.

NVIDIA PhysX

Рекламуючи черговий ігровий шедевр, виробники часто напирають на його 3D-реалістичність. Але яким би реальним не був ігровий 3D-світ, якщо елементарні закони фізики, такі як тяжіння, тертя, гідродинаміка будуть реалізовані неправильно, фальш відчує моментально.

Одна з можливостей фізичного двигуна NVIDIA PhysX – реалістична робота з тканинами.

Реалізувати алгоритми комп'ютерної симуляції базових фізичних законів – справа дуже трудомістка. Найбільш відомими компаніями на цій ниві є ірландська компанія Havok з її міжплатформним фізичним Havok Physics і каліфорнійська Ageia - прабатько першого у світі фізичного процесора (PPU - Physics Processing Unit) та відповідного фізичного движка PhysX. Перша з них, хоч і придбана компанією Intel, активно працює зараз на терені оптимізації двигуна Havok для відеокарт ATI і процесорів AMD. А ось Ageia з її двигуном PhysX стала частиною nVidia. При цьому nVidia вирішила складне завдання адаптації PhysX під технологію CUDA.

Можливе це стало завдяки статистиці. Статистично було доведено, що який би складний рендеринг не виконував GPU, частина його ядер все одно простоює. Саме на цих ядрах і працює двигун PhysX.

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

CUDA-версія пакету стиснення текстур NVIDIA Texture Tools 2

Подобаються реалістичні об'єкти у сучасних іграх? Дякуємо розробникам текстур. Але що більше реальності у текстурі, то більше вписувалося її обсяг. Тим більше вона займає дорогоцінну пам'ять. Щоб цього уникнути, текстури попередньо стискають і динамічно розпаковують у міру потреби. А стиск та розпакування — це суцільні обчислення. Для роботи із текстурами nVidia випустила пакет NVIDIA Texture Tools. Він підтримує ефективне стискування та розпакування текстур стандарту DirectX (так званий ВЧЕ-формат). Друга версія цього пакета може похвалитися підтримкою алгоритмів стиснення BC4 і BC5, реалізованих у технології DirectX 11. Але головне те, що NVIDIA Texture Tools 2 реалізована підтримка CUDA. За оцінкою nVidia, це дає 12-кратний приріст продуктивності у завданнях стиснення та розпакування текстур. А це означає, що фрейми ігрового процесу вантажитимуться швидше і радуватимуть гравця своєю реалістичністю.

Пакет NVIDIA Texture Tools 2 заточено під роботу з CUDA. Приріст продуктивності при стисканні та розпакуванні текстур очевидна.

Використання CUDA дозволяє суттєво підвищити ефективність відеостеження.

Обробка відеопотоку в реальному часі

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

Повноводні річки відеоінформації стікаються до центрів її обробки і... натрапляють на вузьку ланку — людину. Саме він у більшості випадків остання інстанція, яка стежить за відеосвітом. Причому інстанція не найефективніша. Моргає, відволікається і намагається заснути.

Завдяки CUDA з'явилася можливість реалізації алгоритмів одночасного стеження безлічі об'єктів у відеопотоці. При цьому процес відбувається у реальному масштабі часу, а відео є повноцінним 30 fps. Порівняно з реалізацією такого алгоритму на сучасних багатоядерних CPU CUDA дає дво-, триразовий приріст продуктивності, а це, погодьтеся, чимало.

Конвертування відео, фільтрація аудіо

Відеоконвертер Badaboom - перша ластівка, що використовує CUDA для прискорення конвертування.

Приємно подивитися новинку відеопрокату у FullHD-якості та на великому екрані. Але великий екран не візьмеш із собою в дорогу, а відеокодек FullHD икатиме на малопотужному процесорі мобільного гаджета. На допомогу приходить конвертування. Але більшість тих, хто з ним стикався на практиці, нарікають на тривалий час конвертації. Воно і зрозуміло, процес рутинний, придатний до розпаралелювання, і його виконання на CPU не дуже оптимальне.

А ось CUDA з ним справляється на ура. Перша ластівка – конвертер Badaboom від компанії Elevental. Розробники Badaboom, обравши CUDA, не прорахувалися. Тести показують, що стандартний півторагодинний фільм на ньому конвертується у формат iPhone/iPod Touch менш як за двадцять хвилин. І це при тому, що при використанні CPU лише цей процес займає більше години.

Допомагає CUDA та професійним меломанам. Будь-який з них півцарства віддасть за ефективний FIR-кросовер - набір фільтрів, які розділяють звуковий спектр на кілька смуг. Процес цей дуже трудомісткий і за великого обсягу аудіоматеріалу змушує звукорежисера сходити на кілька годин «покурити». Реалізація FIR-кросоверу на базі CUDA прискорює його роботу в сотні разів.

CUDA Future

Зробивши технологію GPGPU реальністю, CUDA не має наміру спочивати на лаврах. Як це відбувається повсюдно, у CUDA працює принцип рефлексії: тепер не тільки архітектура відеопроцесорів nVidia впливає на розвиток версій CUDA SDK, а й сама технологія CUDA змушує nVidia переглядати архітектуру своїх чіпів. Приклад такої рефлексії – платформа nVidia ION. Її другу версію спеціально оптимізовано для вирішення CUDA-завдань. А це означає, що навіть у відносно недорогих апаратних рішеннях споживачі отримають усю міць та блискучі можливості CUDA.

Технологія CUDA

Володимир Фролов,[email protected]

Анотація

Стаття розповідає про технологію CUDA, що дозволяє програмісту використовувати відеокарти як потужні обчислювальні одиниці. Інструменти, надані Nvidia, дозволяють писати програми для графічного процесора (GPU) на підмножині мови С++. Це позбавляє програміста необхідності використання шейдерів і розуміння процесу роботи графічного конвеєра. У статті наведено приклади програмування з використанням CUDA та різні прийоми оптимізації.

1. Введення

Розвиток обчислювальних технологій останні десятки років йшов швидкими темпами. Настільки швидкими, що вже зараз розробники процесорів практично підійшли до так званого «кремнієвого глухого кута». Нестримне зростання тактової частоти стало неможливим через цілу низку серйозних технологічних причин.

Тому всі виробники сучасних обчислювальних систем йдуть у бік збільшення числа процесорів і ядер, а не збільшують частоту одного процесора. Кількість ядер центрального процесора (CPU) у передових системах зараз дорівнює 8.

Інша причина-відносно невисока швидкість роботи оперативної пам'яті. Як би швидко не працював процесор, вузькими місцями, як показує практика, є зовсім не арифметичні операції, а саме невдалі звернення до пам'яті-кеш-промахи.

Однак якщо подивитися у бік графічних процесорів GPU (Graphics Processing Unit), то там шляхом паралелізму пішли набагато раніше. У сьогоднішніх відеокартах, наприклад, у GF8800GTX, число процесорів може досягати 128. Продуктивність подібних систем при вмілому їх програмуванні може бути дуже значною (рис. 1).

Мал. 1. Кількість операцій з плаваючою точкою для CPU та GPU

Коли перші відеокарти тільки з'явилися у продажу, вони являли собою досить прості (порівняно з центральним процесором) вузькоспеціалізовані пристрої, призначені для того, щоб зняти з навантаження процесора з візуалізації двовимірних даних. З розвитком ігрової індустрії та появою таких тривимірних ігор як Doom (рис. 2) та Wolfenstein 3D (рис. 3) виникла потреба у 3D візуалізації.

Малюнки 2.3. Ігри Doom та Wolfenstein 3D

З часу створення компанією 3Dfx перших відеокарт Voodoo (1996 р.) і аж до 2001 року в GPU був реалізований тільки фіксований набір операцій над вхідними даними.

У програмістів був вибору в алгоритмі візуалізації, і підвищення гнучкості з'явилися шейдеры- невеликі програми, виконувані відеокартою кожної вершини чи кожного піксела. У їх завдання входили перетворення над вершинами і затінення-розрахунок освітлення в точці, наприклад, за моделлю Фонга.

Хоча зараз шейдери отримали дуже сильний розвиток, слід розуміти, що вони були розроблені для вузькоспеціалізованих завдань тривимірних перетворень та растеризації. У той час як GPU розвиваються у бік універсальних багатопроцесорних систем, мови шейдерів залишаються вузькоспеціалізованими.

Їх можна порівняти з мовою FORTRAN у тому сенсі, що вони, як і FORTRAN, були першими, але призначеними для вирішення лише одного типу завдань. Шейдери малопридатні для вирішення будь-яких інших завдань, крім тривимірних перетворень та растеризації, як і FORTRAN, не зручний для вирішення завдань, не пов'язаних з чисельними розрахунками.

Сьогодні з'явилася тенденція нетрадиційного використання відеокарт для вирішення завдань у галузях квантової механіки, штучного інтелекту, фізичних розрахунків, криптографії, фізично коректної візуалізації, реконструкції з фотографій, розпізнавання тощо. Ці завдання незручно вирішувати у межах графічних API (DirectX, OpenGL), оскільки ці API створювалися для інших застосувань.

Розвиток програмування загального призначення на GPU (General Programming on GPU, GPGPU) логічно призвело до виникнення технологій, націлених на ширше коло завдань, ніж розтеризація. В результаті компанією Nvidia була створена технологія Compute Unified Device Architecture (або скорочено CUDA), а компанією ATI, що конкурує, - технологія STREAM.

Слід зазначити, що на момент написання цієї статті, технологія STREAM сильно відставала у розвитку від CUDA, і тому вона розглядатися не буде. Ми зосередимося на CUDA – технології GPGPU, що дозволяє писати програми на підмножині мови C++.

2. Принципова різниця між CPU та GPU

Розглянемо коротко деякі істотні відмінності між областями та особливостями застосувань центрального процесора та відеокарти.

2.1. Можливості

CPU спочатку пристосований для вирішення завдань загального плану і працює з пам'яттю, що довільно адресується. Програми на CPU можуть звертатися безпосередньо до будь-яких осередків лінійної та однорідної пам'яті.

Для GPU це негаразд. Як ви дізнаєтеся, прочитавши цю статтю, CUDA має цілих 6 видів пам'яті. Читати можна з будь-якого осередку, доступного фізично, але записувати - не в усі осередки. Причина полягає в тому, що GPU в будь-якому випадку є специфічним пристроєм, призначеним для конкретних цілей. Це обмеження запроваджено задля збільшення швидкості роботи певних алгоритмів та зниження вартості обладнання.

2.2. Швидкодія пам'яті

Одвічна проблема більшості обчислювальних систем полягає в тому, що пам'ять працює повільніше процесора. Виробники CPU вирішують її шляхом запровадження кешів. Найчастіше використовувані ділянки пам'яті міститься в надоперативну або кеш-пам'ять, що працює на частоті процесора. Це дозволяє заощадити час при зверненні до даних, що найчастіше використовуються, і завантажити процесор власне обчисленнями.

Зауважимо, що кеші для програміста практично прозорі. Як при читанні, так і при записі дані не потрапляють одразу до оперативної пам'яті, а проходять через кеші. Це дозволяє, зокрема, швидко зчитувати деяке значення відразу після запису .

На GPU (тут мається на увазі відеокарти GF восьмої серії) кеші теж є, і вони теж важливі, але цей механізм не такий потужний, як на CPU. По-перше, кешується в повному обсязі типи пам'яті, а по-друге, кеші працюють лише з читання.

На GPU повільні звернення до пам'яті приховують за допомогою паралельних обчислень. Поки одні завдання чекають на дані, працюють інші, готові до обчислень. Це один із основних принципів CUDA, що дозволяють сильно підняти продуктивність системи в цілому.

3. Ядро CUDA

3.1. Поточна модель

Обчислювальна архітектура CUDA заснована на концепціїодна команда на безліч даних(Single Instruction Multiple Data, SIMD) та поняття мультипроцесора.

Концепція SIMD має на увазі, що одна інструкція дозволяє одночасно обробити безліч даних. Наприклад, команда addps у процесорі Pentium 3 і новіших моделях Pentium дозволяє складати одночасно 4 числа з плаваючою точкою одинарної точності.

Мультипроцесор - це багатоядерний SIMD процесор, що у кожен певний час виконувати на всіх ядрах лише одну інструкцію. Кожне ядро ​​мультипроцесора скалярне, тобто. воно не підтримує векторні операції у чистому вигляді.

Перед тим, як продовжити, введемо пару визначень. Зазначимо, що під пристроєм і хостом у цій статті розумітиметься зовсім не те, до чого звикла більшість програмістів. Ми будемо користуватися такими термінами, щоб уникнути розбіжностей із документацією CUDA.

Під пристроєм (device) у нашій статті ми розумітимемо відеоадаптер, що підтримує драйвер CUDA, або інший спеціалізований пристрій, призначений для виконання програм, що використовують CUDA (таке, наприклад, як NVIDIA Tesla). У нашій статті ми розглянемо GPU лише як логічний пристрій, уникаючи конкретних деталей реалізації.

Хостом (host ) ми називатимемо програму у звичайній оперативної пам'яті комп'ютера, використовує CPU і виконує функції з роботи з пристроєм.

Фактично, та частина вашої програми, яка працює на CPU – цехост, а ваша відеокарта -пристрій. Логічно пристрій можна подати як набір мультипроцесорів (рис. 4) плюс драйвер CUDA.

Мал. 4. Пристрій

Припустимо, що ми хочемо запустити на нашому пристрої деяку процедуру в N потоках (тобто хочемо розпаралелити її роботу). Відповідно до документації CUDA, назвемо цю процедуру ядром.

Особливістю архітектури CUDA є блочно-сіточна організація, незвичайна багатопоточних додатків (рис. 5). Драйвер CUDA самостійно розподіляє ресурси пристрою між потоками.

Мал. 5. Організація потоків

На рис. 5. ядро ​​позначене як Kernel. Всі потоки, що виконують це ядро, об'єднуються в блоки (Block), а блоки, своєю чергою, об'єднуються в сітку (Grid).

Як видно з рис 5, для ідентифікації потоків використовуються двомірні індекси. Розробники CUDA надали можливість працювати з тривимірними, двовимірними або простими (одномірними) індексами, залежно від того, як зручніше програмісту.

У загальному випадку індекси є тривимірними векторами. Для кожного потоку будуть відомі: індекс потоку всередині блоку threadIdx та індекс блоку всередині сітки blockIdx. При запуску всі потоки відрізнятимуться лише цими індексами. Фактично саме через ці індекси програміст здійснює управління, визначаючи, яка саме частина його даних обробляється в кожному потоці.

Відповідь на запитання, чому розробники обрали саме таку організацію, нетривіальна. Одна з причин полягає в тому, що один блок гарантовано виконується.на одному мультипроцесор пристрою, але один мультипроцесор може виконувати кілька різних блоків. Решта причин проясниться далі під час статті.

Блок завдань (потоків) виконується на мультипроцесорі частинами або пулами, званими warp. Розмір warp на даний момент у відеокартах з підтримкою CUDA дорівнює 32 потокам. Завдання всередині пулу warp виконуються у SIMD стилі, тобто. у всіх потоках всередині warp одночасно може виконуватися лише одна інструкція.

Тут слід зробити одне застереження. В архітектурах, сучасних на момент написання цієї статті, кількість процесорів всередині одного мультипроцесора дорівнює 8, а не 32. З цього випливає, що не весь warp виконується одночасно, він розбивається на 4 частини, які виконуються послідовно (т.к. скалярні процесори) .

Але, по-перше, розробники CUDA не регламентують жорстко розмір warp. У своїх роботах вони згадують параметр warp size, а не число 32. По-друге, з логічного погляду саме warp є тим мінімальним об'єднанням потоків, про який можна говорити, що всі потоки всередині нього виконуються одночасно - і при цьому жодних припущень щодо решти системи зроблено не буде.

3.1.1. Розгалуження

Відразу ж виникає питання: якщо в той самий момент часу всі потоки всередині warp виконують одну і ту ж інструкцію, то як бути з розгалуженнями? Адже якщо код програми зустрічається розгалуження, то інструкції будуть вже різні. Тут застосовується стандартне для SIMD програмування рішення (рис. 6).

Мал. 6. Організація розгалуження в SIMD

Нехай є наступний код:

if(cond)B;

У випадку SISD (Single Instruction Single Data) ми виконуємо оператор A, перевіряємо умову, потім виконуємо оператори B і D (якщо умова є істинною).

Нехай тепер ми маємо 10 потоків, що виконуються в стилі SIMD. У всіх 10 потоках ми виконуємо оператор A, потім перевіряємо умову cond і виявляється, що в 9 з 10 потоках воно є істинним, а в одному потоці - хибним.

Зрозуміло, що ми не можемо запустити 9 потоків для виконання оператора B, а один - для виконання оператора C, тому що одночасно у всіх потоках може виконуватися тільки одна інструкція. В цьому випадку потрібно вчинити так: спочатку «вбиваємо» потік, що відколовся, так, щоб він не псував нічиї дані, і виконуємо 9 потоків, що залишилися. Потім "вбиваємо" 9 потоків, що виконали оператор B, і проходимо один потік з оператором C. Після цього потоки знову об'єднуються і виконують оператор D все одночасно.

Виходить сумний результат: мало того що ресурси процесорів витрачаються на порожнє перемелювання бітів в потоках, що відкололися, так ще, що набагато гірше, ми будемо змушені в результаті виконати ОБІГІ гілки.

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

З усього вищесказаного випливає один дуже важливий висновок. Розгалуження є причиною падіння продуктивності власними силами. Шкідливі лише ті розгалуження, на яких потоки розходяться всередині одного пулу потоків warp. При цьому якщо потоки розійшлися всередині одного блоку, але в різних пулах warp, або всередині різних блоків, це не надає ніякого ефекту.

3.1.2. Взаємодія між потоками

На момент написання цієї статті будь-яка взаємодія між потоками (синхронізація та обмін даними) була можлива лише всередині блоку. Тобто, між потоками різних блоків не можна організувати взаємодію, користуючись лише документованими можливостями.

Щодо недокументованих можливостей, ними користуватися вкрай не рекомендується. Причина в тому, що вони спираються на конкретні апаратні особливості тієї чи іншої системи.

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

3.2. Пам'ять

У CUDA виділяють шість видів пам'яті (рис. 7). Це регістри, локальна, глобальна, розділяється, константна та текстурна пам'ять.

Така велика кількість обумовлена ​​специфікою відеокарти та первинним її призначенням, а також прагненням розробників зробити систему якомога дешевше, жертвуючи в різних випадках або універсальністю, або швидкістю.

Мал. 7. Види пам'яті у CUDA

3.2.0. Реєстри

По можливості компілятор намагається розміщувати всі локальні змінні функції у регістрах. Доступ до таких змінних здійснюється максимальною швидкістю. У поточній архітектурі на один мультипроцесор доступно 8192 32-розрядних регістру. Щоб визначити, скільки доступно регістрів одному потоку, треба розділити це число (8192) на розмір блоку (кількість потоків у ньому).

При нормальному поділі в 64 потоки на блок виходить всього 128 регістрів (існують деякі об'єктивні критерії, але 64 підходить в середньому для багатьох завдань). Реально, 128 регістрів nvcc ніколи не виділить. Зазвичай він не дає більше 40, а решта змінних попаде в локальну пам'ять. Так відбувається тому, що на одному мультипроцесорі може виконуватися кілька блоків. Компілятор намагається максимізувати кількість одночасно працюючих блоків. Для більшої більшої ефективності треба намагатися займати менше ніж 32 регістри. Тоді теоретично може бути запущено 4 блоки (8 warp-ів, якщо 64 треди в одному блоці) на одному мультипроцесорі. Однак тут ще слід враховувати об'єм пам'яті, що займається потоками, оскільки якщо один блок займає всю розділювану пам'ять, два таких блоки не можуть виконуватися на мультипроцесорі одночасно .

3.2.1. Локальна пам'ять

У випадках, коли локальні дані процедур займають занадто великий розмір, або компілятор не може обчислити для них постійний крок при зверненні, він може помістити їх у локальну пам'ять. Цьому може сприяти, наприклад, приведення покажчиків типів різних розмірів.

Фізично локальна пам'ять є аналогом глобальної пам'яті і працює з тією ж швидкістю. На момент написання статті не було жодних механізмів, що дозволяють явно заборонити компілятор використання локальної пам'яті для конкретних змінних. Так як проконтролювати локальну пам'ять досить важко, краще не використовувати її зовсім (див. розділ 4 «Рекомендації оптимізації»).

3.2.2. Глобальна пам'ять

У документації CUDA як одне з основних досягненьТехнологія наводить можливість довільної адресації глобальної пам'яті. Тобто можна читати з будь-якої комірки пам'яті, і писати можна теж у довільну комірку (на GPU це зазвичай не так).

Проте за універсальність у разі доводиться розплачуватися швидкістю. Глобальна пам'ять не кешується. Вона працює дуже повільно, кількість звернень до глобальної пам'яті слід у будь-якому разі мінімізувати.

Глобальна пам'ять необхідна переважно збереження результатів роботи програми перед відправкою їх у хост (у звичайну пам'ять DRAM). Причина цього в тому, що глобальна пам'ять є єдиним видом пам'яті, куди можна щось записувати.

Змінні, оголошені з кваліфікатором __global__, розміщуються у світовій пам'яті. Глобальну пам'ять можна також виділити динамічно, викликавши функцію cudaMalloc(void* mem, int size) на хості. З пристрою цю функцію не можна викликати. Звідси випливає, що розподілом пам'яті має займатися програма-хост, що працює на CPU. Дані з хоста можна надсилати у пристрій викликом функції cudaMemcpy:

cudaMemcpy(void* gpu_mem, void* cpu_mem, int size, cudaMemcpyHostToDevice);

Так само можна зробити і зворотну процедуру:

cudaMemcpy(void* cpu_mem, void* gpu_mem, int size, cudaMemcpyDeviceToHost);

Цей виклик також здійснюється з хоста.

При роботі з глобальною пам'яттю важливо дотримуватися правила коалесингу (coalescing). Основна ідея в тому, що треди повинні звертатися до послідовних осередків пам'яті, причому 4,8 або 16 байт. При цьому найперший тред повинен звертатися за адресою, вирівняною на кордон відповідно 4,8 або 16 байт. Адреси, що повертаються cudaMalloc, вирівняні як мінімум за кордоном 256 байт.

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

Пам'ять, що розділяється - це некешируемая, але швидка пам'ять. Її і рекомендується використовувати як керований кеш. На один мультипроцесор доступно всього 16KB пам'яті, що розділяється. Розділивши це число на кількість завдань у блоці, отримаємо максимальну кількість пам'яті, що розділяється, доступною на один потік (якщо планується використовувати її незалежно у всіх потоках).

Відмінною рисою пам'яті, що розділяється, є те, що вона адресується однаково для всіх завдань усередині блоку (рис. 7). Звідси випливає, що її можна використовувати для обміну даними між потоками лише одного блоку.

Гарантується, що під час виконання блоку на мультипроцесорі вміст пам'яті буде зберігатися. Однак після того, як на мультипроцесорі змінився блок, не гарантується, що вміст старого блоку зберігся. Тому не варто намагатися синхронізувати завдання між блоками, залишаючи в пам'яті якісь дані і сподіваючись на їх збереження.

Змінні, оголошені з кваліфікатором __shared__, розміщуються в пам'яті, що розділяється.

Shared__ float mem_shared;

Слід ще раз підкреслити, що пам'ять, що розділяється, для блоку одна. Тому якщо потрібно використовувати її просто як керований кеш, слід звертатися до різних елементів масиву, наприклад:

float x = mem_shared;

Де threadIdx.x – індекс x потоку всередині блоку.

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

Константна пам'ять кешується, як видно на рис. 4. Кеш існує в єдиному екземплярі одного мультипроцесора, отже, загальний всім завдань усередині блока. На хості в константну пам'ять можна щось записати, викликавши функцію cudaMemcpyToSymbol. З пристрою константна пам'ять доступна лише для читання.

Константна пам'ять дуже зручна у використанні. Можна розміщувати в ній дані будь-якого типу та читати їх за допомогою простого присвоєння.

#define N 100

Constant__ int gpu_buffer[N];

void host_function()

int cpu_buffer[N];

cudaMemcpyToSymbol(gpu_buffer, cpu_buffer, sizeof(int )*N);

//__global__ означає, що device_kernel - ядро, яке може бути запущене на GPU

Global__ void device_kernel()

int a = gpu_buffer;

int b = gpu_buffer + gpu_buffer;

// gpu_buffer = a; ПОМИЛКА! константна пам'ять доступна лише для читання

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

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

Текстурна пам'ять кешується (рис. 4). Для кожного мультипроцесора є лише один кеш, отже, цей кеш загальний всім завдань усередині блоку.

Назва текстурної пам'яті (і, на жаль, функціональність) успадковано від понять «текстура» та «текстурування». Текстурування - це процес накладання текстури (просто картинки) на полігон у процесі растеризації. Текстурна пам'ять оптимізована під вибірку 2D даних і має такі можливості:

швидка вибірка значень фіксованого розміру (байт, слово, подвійне або вчетверне слово) з одномірного або двомірного масиву;

нормалізована адресація числами типу float в інтервалі. Потім їх можна вибирати, використовуючи нормалізовану адресацію. Результуючим значенням буде слово типу float4, що відображається в інтервал ;

CudaMalloc((void**) &gpu_memory, N*sizeof (uint4 )); //виділимо пам'ять у GPU

// Налаштування параємтрів текстури texture

Texture.addressMode = cudaAddressModeWrap; //режим Wrap

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; //найближче значення

Texture.normalized = false; // не використовувати нормалізовану адресацію

CudaBindTexture (0, texture , gpu _ memory , N ) // відтепер ця пам'ять вважатиметься текстурною


CudaMemcpy (gpu _ memory , cpu _ buffer , N * sizeof (uint 4), cudaMemcpyHostToDevice ); // Копіюємо дані наGPU

//__global__ означає, що device_kernel - ядро, яке потрібно розпаралелити

Global__ void device_kernel()

uint4 a = tex1Dfetch(texture,0); // можна вибирати дані лише в такий спосіб!

uint4 b = tex1Dfetch(texture,1);

int c = a.x*b.y;

...

3.3. Простий приклад

Як простий приклад пропонується розглянути програму cppIntegration з CUDA SDK. Вона демонструє прийоми роботи з CUDA, а також використання nvcc (спеціальний компілятор підмножини С++ від Nvidia) у поєднанні з MS Visual Studio, що спрощує розробку програм на CUDA.

4.1. Правильно проводьте розбиття вашого завдання

Не всі завдання підходять для архітектури SIMD. Якщо ваше завдання для цього не придатне, можливо, не варто використовувати GPU. Але якщо ви твердо вирішили використовувати GPU, потрібно намагатися розбити алгоритм на такі частини, щоб вони могли ефективно виконуватись у стилі SIMD. Якщо потрібно – змініть алгоритм для вирішення вашого завдання, придумайте новий – той, який добре б лягав на SIMD. Як приклад відповідної області використання GPU можна навести реалізацію пірамідального складання елементів масиву.

4.2. Вибір типу пам'яті

Поміщайте свої дані в текстурну або константну пам'ять, якщо всі завдання одного блоку звертаються до однієї ділянки пам'яті або близько розташованих ділянок. Двовимірні дані можуть бути ефективно оброблені за допомогою функцій text2Dfetch та text2D. Текстурну пам'ять спеціально оптимізовано під двомірну вибірку.

Використовуйте глобальну пам'ять у поєднанні з пам'яттю, що розділяється, якщо всі завдання звертаються безсистемно до різних, далеко розташованих один від одного ділянок пам'яті (з дуже різними адресами або координатами, якщо це 2D/3D дані).

глобальна пам'ять => пам'ять, що розділяється

Syncthreads();

Обробити дані в пам'яті, що розділяється

Syncthreads();

глобальна пам'ять<= разделяемая память

4.3. Увімкніть лічильники пам'яті

Прапор компілятора --ptxas-options=-v дозволяє точно сказати, скільки і якої пам'яті (регістрів, локальної, константної, що розділяється) ви використовуєте. Якщо компілятор використовує локальну пам'ять, ви знаєте про це. Аналіз даних про кількість і типи пам'яті, що використовується, може сильно допомогти вам при оптимізації програми.

4.4. Намагайтеся мінімізувати використання регістрів і пам'яті, що розділяється

Чим більше ядро ​​використовує регістрів або пам'яті, що розділяється, тим менше потоків (вірніше warp-ов) одночасно можуть виконуватися на мультипроцесорі, т.к. ресурси мультипроцесора обмежені. Тому невелике збільшення зайнятості регістрів або пам'яті, що розділяється, може призводити в деяких випадках до падіння продуктивності в два рази - саме через те, що тепер рівно в два рази менше warp-ів одночасно виконуються на мультипроцесорі.

4.5. Пам'ять, що розділяється, замість локальної

Якщо компілятор Nvidia з якоїсь причини розташував дані в локальній пам'яті (зазвичай це помітно з дуже сильного падіння продуктивності в місцях, де нічого ресурсомісткого немає), з'ясуйте, які саме дані потрапили в локальну пам'ять, і помістіть їх у пам'ять, що розділяється (shared memory ).

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

Якщо ж ви дійсно рідко використовуєте змінну - краще помістити її в глобальну пам'ять.

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

4.6. Розгортання циклів

Розгортання циклів є стандартним прийомом підвищення продуктивності у багатьох системах. Суть його в тому, щоб на кожній ітерації виконувати більше дій, зменшивши у такий спосіб загальну кількість ітерацій, а отже, і кількість умовних переходів, які має виконати процесор.

Ось як можна розгорнути цикл знаходження суми масиву (наприклад, цілісного):

int a[N]; int summ;

for (int i=0;i

Очевидно, цикли можна розгорнути і вручну (як показано вище), але це малопродуктивна праця. Набагато краще використовувати шаблони С++ у поєднання з функціями, що вбудовуються.

template

class ArraySumm

Device__ static T exec (const T * arr) ( return arr + ArraySumm (arr+1); )

template

class ArraySumm<0,T>

Device__ static T exec(const T* arr) ( return 0; )

for (int i=0;i

summ+= ArraySumm<4,int>::exec(a);

Слід зазначити одну цікаву особливість компілятора nvcc. Компілятор завжди вбудуватиме функції типу __device__ за умовчанням (щоб це скасувати, існує спеціальна директива __noinline__) .

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

4.7. Вирівнювання даних та вибірка по 16 байт

Вирівнюйте структури даних по 16-байтовому кордону. У цьому випадку компілятор зможе використовувати для них спеціальні інструкції, що виконують завантаження даних одразу по 16 байт.

Якщо структура займає 8 б або менше, можна вирівнювати її по 8 б. Але в цьому випадку можна вибрати відразу дві змінні за один раз, об'єднавши дві 8-байтові змінні в структуру за допомогою union або приведення покажчиків. Приведенням слід користуватися обережно, оскільки компілятор може помістити дані локальну пам'ять, а чи не в регістри.

4.8. Конфлікти банків пам'яті, що розділяється

Пам'ять, що розділяється, організована у вигляді 16 (всього-то!) банків пам'яті з кроком в 4 байти. Під час виконання пулу потоків warp на мультипроцесорі, він ділиться на дві половинки (якщо warp-size = 32) по 16 потоків, які здійснюють доступ до пам'яті по черзі.

Завдання в різних половинах warp не конфліктують з пам'яті, що розділяється. Через завдання однієї половинки пулу warp будуть звертатися до однакових банків пам'яті, виникнуть колізії і, як наслідок, падіння продуктивності. Завдання в межах однієї половинки warp можуть звертатися до різних ділянок пам'яті, що розділяється, з певним кроком.

Оптимальні кроки - 4, 12, 28, ..., 2 n-4 байт (рис. 8).



Мал. 8. Оптимальні кроки.

Не оптимальні кроки - 1, 8, 16, 32, ..., 2 n байт (рис. 9).



Мал. 9. Неоптимальні кроки

4.9. Мінімізація переміщень даних Host<=>Device

Намагайтеся якнайрідше передавати проміжні результати на host для обробки за допомогою CPU. Реалізуйте якщо не весь алгоритм, то принаймні його основну частину на GPU, залишаючи CPU лише керуючі завдання.

5. CPU/GPU переносима математична бібліотека

Автором цієї статті написана бібліотека MGML_MATH, що переноситься, для роботи з простими просторовими об'єктами, код якої працездатний як на пристрої, так і на хості.

Бібліотека MGML_MATH може бути використана як каркас для написання CPU/GPU переносних (або гібридних) систем розрахунку фізичних, графічних чи інших просторових завдань. Основна її перевага в тому, що один і той же код може використовуватися як на CPU, так і на GPU, і при цьому на чільне місце вимог, що пред'являються до бібліотеки, ставиться швидкість.

6 . Література

Кріс Касперскі. Техніка оптимізації програм. Ефективне використання пам'яті. – Спб.: БХВ-Петербург, 2003. – 464 с.: іл.

CUDA Programming Guide 1.1 ( http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

CUDA Programming Guide 1.1. page 14-15

CUDA Programming Guide 1.1. page 48

Згідно з Дарвінською теорією еволюції, перша людиноподібна мавпа (якщо
бути точним – homo antecessor, людина-попередник) перетворилася згодом
у нас. Багатотонні обчислювальні центри з тисячею і більше радіоламп,
які займають цілі кімнати, змінилися напівкілограмовими ноутами, які, до речі,
не поступляться продуктивністю першим. Допотопні друкарські машинки перетворилися
у друкуючі що завгодно та на чому завгодно (навіть на тілі людини)
багатофункціональні пристрої. Процесорні гіганти раптом надумали замурувати
графічне ядро ​​в "камінь". А відеокарти стали не тільки показувати картинку з
прийнятним FPS та якістю графіки, але й проводити всілякі обчислення. Так
ще як робити! Про технологію багатопоточних обчислень засобами GPU, і йтиметься.

Чому GPU?

Цікаво, чому всю обчислювальну міць вирішили перекласти на графічний
адаптер? Як видно, процесори ще в моді, та й навряд чи поступляться своїм теплим
Містечко. Але у GPU є пара козирів у рукаві разом із джокером, та й рукавів
вистачає. Сучасний центральний процесор заточений під отримання максимальної
продуктивності при обробці цілих даних та даних з плаваючою
комою, особливо не переймаючись при цьому паралельною обробкою інформації. В теж
час архітектура відеокарти дозволяє швидко і без проблем «розпаралелити»
обробку даних. З одного боку, йде облік полігонів (за рахунок 3D-конвеєра),
з іншого – піксельна обробка текстур. Видно, що відбувається «злагоджена
розбивка навантаження в ядрі карти. Крім того, робота пам'яті та відеопроцесора
оптимальніша, ніж зв'язка «ОЗУ-кеш-процесор». У той момент, коли одиниця даних
у відеокарті починає оброблятися одним потоковим процесором GPU, інша
одиниця паралельно завантажується в інший, і, в принципі, легко можна досягти
завантаженості графічного процесора, порівнянної з пропускною здатністю шини,
однак для цього завантаження конвеєрів має здійснюватися однаково, без
всяких умовних переходів та розгалужень. Центральний процесор через свою
універсальності вимагає для своїх процесорних потреб кеш, заповнений
інформацією.

Вчені чоловіки задумалися щодо роботи GPU у паралельних обчисленнях та
математики і вивели теорію, що багато наукових розрахунків багато в чому схожі з
обробкою 3D-графіки. Багато експертів вважають, що основним фактором у
розвитку GPGPU (General Purpose computation on GPU – універсальні
розрахунки засобами відеокарти) стала поява у 2003 році проекту Brook GPU.

Творцям проекту зі Стендфордського університету треба було вирішити непросту
проблему: апаратно та програмно змусити графічний адаптер виробляти
різнопланові обчислення. І вони це вийшло. Використовуючи універсальну мову C,
американські вчені змусили працювати GPU як процесор, з поправкою на
паралельну обробку. Після Brook з'явилася ціла низка проектів за VGA-розрахунками,
таких як бібліотека Accelerator, бібліотека Brahma, система
метапрограмування GPU++ та інші.

CUDA!

Передчуття перспективності розробки змусило AMDі NVIDIA
вчепитися в Brook GPU, як пітбуль. Якщо опустити маркетингову політику, то,
реалізувавши все правильно, можна закріпитися не лише у графічному секторі
ринку, а й у обчислювальному (дивись на спеціальні обчислювальні карти та
сервери Teslaз сотнями мультипроцесорів), потіснивши звичні для всіх CPU.

Природно, «володарі FPS» розійшлися біля каменю спотикання кожен за своєю.
стежці, але основний принцип залишився незмінним - проводити обчислення
засобами GPU. І зараз ми докладніше розглянемо технологію «зелених». CUDA
(Compute Unified Device Architecture).

Робота нашої «героїні» полягає у забезпеченні API, причому одразу двох.
Перший – високорівневий, CUDA Runtime, є функціями, які
розбиваються більш прості рівні і передаються нижньому API – CUDA Driver. Так
що фраза «високорівневий» може бути застосована до процесу з натяжкою. Вся сіль знаходиться
саме у драйвері, і добути її допоможуть бібліотеки, люб'язно створені
розробниками NVIDIA: CUBLAS (засоби для математичних розрахунків) та
FFT (розрахунок за допомогою алгоритму Фур'є). Ну що ж, перейдемо до практичної
частини матеріалу.

Термінологія CUDA

NVIDIAоперує дуже своєрідними визначеннями для CUDA API. Вони
відрізняються від визначень, що застосовуються до роботи з центральним процесором.

Потік- Набір даних, який необхідно обробити (не
потребує великих ресурсів під час обробки).

Варп (warp)- Група з 32 потоків. Дані обробляються тільки
варпами, отже варп - це мінімальний обсяг даних.

Блок (блок)- Сукупність потоків (від 64 до 512) або сукупність
варпів (від 2 до 16).

Сітка (grid)- Це сукупність блоків. Такий поділ даних
застосовується виключно підвищення продуктивності. Так, якщо число
мультипроцесорів велике, то блоки виконуватимуться паралельно. Якщо ж з
карткою не пощастило (розробники рекомендують для складних розрахунків використовувати
адаптер не нижче рівня GeForce 8800 (GTS 320 Мб), то блоки даних опрацьовуються
послідовно.

Також NVIDIA вводить такі поняття, як ядро (kernel), хост (host)
і девайс (device).

Працюємо!

Для повноцінної роботи з CUDA потрібно:

1. Знати будову шейдерних ядер GPU, оскільки суть програмування
полягає в рівномірному розподілі навантаження між ними.
2. Вміти програмувати серед С, з урахуванням деяких аспектів.

Розробники NVIDIAрозкрили «начинки» відеокарти кілька
інакше, ніж ми звикли бачити. Так що мимоволі доведеться вивчати все
тонкощі архітектури. Розберемо будову «каменю» G80 легендарною GeForce 8800
GTX.

Шейдерне ядро ​​складається з восьми TPC (Texture Processor Cluster) – кластерів
текстурних процесорів (так, у GeForce GTX 280- 15 ядер, у 8800 GTS
їх шість, у 8600 - Чотири і т.д.). Ті, у свою чергу, складаються з двох
потокових мультипроцесорів (streaming multiprocessor – далі SM). SM (їх всього
16) складається з front end (розв'язує задачі читання та декодування інструкцій) та
back end (кінцевий висновок інструкцій) конвеєрів, а також восьми scalar SP (shader)
processor) та двома SFU (суперфункціональні блоки). За кожен такт (одиницю
часу) front end вибирає варп та обробляє його. Щоб усі потоки варпа
(Нагадаю, їх 32 штуки) обробилися, потрібно 32/8 = 4 такти в кінці конвеєра.

Кожен мультипроцесор має так звану загальну пам'ять (shared memory).
Її розмір складає 16 кілобайт та надає програмісту повну свободу
дій. Розподіляй як хочеш:). Shared memory забезпечує зв'язок потоків у
одному блоці та не призначена для роботи з піксельними шейдерами.

Також SM можуть звертатися до GDDR. Для цього їм «пришили» по 8 кілобайт
кеш-пам'яті, що зберігають все найголовніше для роботи (наприклад, обчислювальні
константи).

Мультипроцесор має 8192 регістри. Число активних блоків не може бути
більше восьми, а число варпів – не більше 768/32 = 24. З цього видно, що G80
може обробити максимум 32*16*24 = 12 288 потоків за одиницю часу. Не можна не
враховувати ці цифри при оптимізації програми надалі (на одній чаші ваг
- Розмір блоку, на інший - кількість потоків). Баланс параметрів може зіграти
важливу роль надалі, тому NVIDIAрекомендує використовувати блоки
зі 128 або 256 потоками. Блок з 512 потоків неефективний, оскільки має
підвищеними затримками. Враховуючи всі тонкощі будови GPU відеокарти плюс
непогані навички у програмуванні, можна створити дуже продуктивне
засіб для паралельних обчислень. До речі, про програмування...

Програмування

Для «творчості» разом із CUDA потрібно відеокарта GeForce не нижче
восьмий серії. З

офіційного сайту потрібно завантажити три програмні пакети: драйвер з
підтримкою CUDA (для кожної ОС – свій), безпосередньо пакет CUDA SDK (друга
бета-версія) та додаткові бібліотеки (CUDA toolkit). Технологія підтримує
операційні системи Windows (XP і Vista), Linux та Mac OS X. Для вивчення я
вибрав Vista Ultimate Edition x64 (забігаючи вперед, скажу, що система поводилася
просто чудово). У момент написання цих рядків актуальним для роботи був
драйвер для ForceWare 177.35. Як набор інструментів використовувався
програмний пакет Borland C++ 6 Builder (хоча підійде будь-яке середовище, що працює з
мовою C).

Людині, яка знає мову, буде легко освоїтися в новому середовищі. Потрібно лише
запам'ятати основні параметри. Ключове слово _global_ (ставиться перед функцією)
показує, що функція стосується kernel (ядра). Її викликатиме центральний
процесор, а вся робота відбудеться на GPU. Виклик _global_ вимагає більше
конкретних деталей, а саме розмір сітки, розмір блоку та яке ядро ​​буде
застосовано. Наприклад, рядок _global_ void saxpy_parallel<<>>, де X –
розмір сітки, а Y – розмір блоку, задає ці параметри.

Символ _device_ означає, що функцію викличе графічне ядро, воно ж
виконає усі інструкції. Ця функція знаходиться в пам'яті мультипроцесора,
отже, отримати її адресу неможливо. Префікс _host_ означає, що виклик
та обробка пройдуть лише за участю CPU. Треба враховувати, що _global_ і
_device_ не можуть викликати один одного і не можуть викликати себе.

Також мова для CUDA має низку функцій для роботи з відеопам'яттю: cudafree
(звільнення пам'яті між GDDR і RAM), cudamemcpy та cudamemcpy2D (копіювання
пам'яті між GDDR і RAM) і cudamalloc (виділення пам'яті).

Усі програмні коди проходять компіляцію із боку CUDA API. Спочатку береться
код, призначений виключно для центрального процесора, і піддається
стандартної компіляції, а інший код, призначений для графічного адаптера,
переписується в проміжну мову PTX (сильно нагадує асемблер) для
виявлення можливих помилок. Після всіх цих «танців» відбувається остаточний
переклад (трансляція) команд у зрозумілу для GPU/CPU мову.

Набір для вивчення

Практично всі аспекти програмування описані в документації, що йде
разом з драйвером та двома програмами, а також на сайті розробників. Розміру
статті не вистачить, щоб описати їх (зацікавлений читач має додати
малу дещицю старань і вивчити матеріал самостійно).

Спеціально для новачків розроблено CUDA SDK Browser. Будь-хто може
відчути силу паралельних обчислень на своїй шкурі (краща перевірка на
стабільність – робота прикладів без артефактів та вильотів). Додаток має
великий ряд показових міні-програм (61 «тест»). До кожного досвіду є
детальна документація програмного коду плюс PDF-файли. Відразу видно, що люди,
присутні зі своїми творіннями у браузері, займаються серйозною роботою.
Тут же можна порівняти швидкості роботи процесора та відеокарти під час обробки
даних. Наприклад, сканування багатовимірних масивів відеокартою GeForce 8800
GT 512 Мб з блоком з 256 потоками виробляє за 0.17109 мілісекунди.
Технологія не розпізнає SLI-тандеми, тому якщо у тебе дует або тріо,
відключай функцію «спарювання» перед роботою, інакше CUDA побачить лише один
девайс. Двоядерний AMD Athlon 64 X2(Частота ядра 3000 МГц) той же досвід
проходить за 2.761528 мілісекунди. Виходить, що G92 більш ніж у 16 ​​разів
швидше «каменю» AMD! Як бачиш, далеко не екстремальна система в
тандемі з нелюбою в масах операційною системою показує непогані
результати.

Крім браузера існує низка корисних суспільству програм. Adobe
адаптувала свої продукти до нової технології. Тепер Photoshop CS4 у повній
мірою використовує ресурси графічних адаптерів (необхідно завантажити спеціальний
плагін). Такими програмами, як Badaboom media converter та RapiHD, можна
зробити декодування відео у форматі MPEG-2. Для обробки звуку непогано
підійде безкоштовна утиліта Accelero. Кількість софту, заточеного під CUDA API,
безперечно, зростатиме.

А в цей час…

А поки ти читаєш цей матеріал, трудяги з процесорних концернів
розробляють свої технології з впровадження GPU у CPU. З боку AMDУсе
зрозуміло: у них є величезний досвід, набутий разом із ATI.

Творіння «мікродевайсерів», Fusion, складатиметься з кількох ядер під
кодовою назвою Bulldozer та відеочіпа RV710 (Kong). Їхній взаємозв'язок буде
здійснюватись за рахунок покращеної шини HyperTransport. Залежно від
кількості ядер та їх частотних характеристик AMD планує створити цілу цінову
ієрархію «каменів». Також планується виробляти процесори як для ноутбуків (Falcon),
і для мультимедійних гаджетів (Bobcat). Причому саме застосування технології
у портативних пристроях буде початковим завданням для канадців. З розвитком
паралельних обчислень застосування таких «камінців» має бути дуже популярним.

Intelтрішки відстає за часом зі своєю Larrabee. Продукти AMD,
якщо нічого не станеться, з'являться на прилавках магазинів наприкінці 2009 – на початку
2010 року. А рішення противника вийде на світ божий лише майже через два
року.

Larrabee налічуватиме велику кількість (читай – сотні) ядер. На початку
ж вийдуть товари, розраховані на 8 – 64 ядер. Вони дуже подібні до Pentium, але
досить сильно перероблені. Кожне ядро ​​має 256 кілобайт кешу другого рівня
(Згодом його розмір збільшиться). Взаємозв'язок здійснюватиметься за рахунок
1024-бітної двонаправленої кільцевої шини. Інтел каже, що їхня «дитя» буде
добре працювати з DirectX і Open GL API (для «яблучників»), тому ніяких
програмних втручань не потрібно.

А чого я все це тобі розповів? Очевидно, що Larrabee та Fusion не витіснять
звичайні, стаціонарні процесори з ринку, так само, як не витіснять з ринку
відеокарти. Для геймерів та екстремалів межею мрій, як і раніше, залишиться
багатоядерний CPU та тандем з декількох топових VGA. Але те, що навіть
процесорні компанії переходять на паралельні обчислення за принципами,
аналогічним GPGPU, говорить вже багато про що. Зокрема про те, що така
технологія, як CUDA, має право на існування і, мабуть, буде
дуже популярна.

Невелике резюме

Паралельні обчислення засобами відеокарти – лише хороший інструмент
у руках працьовитого програміста. Навряд чи процесорам на чолі із законом Мура
настане кінець. Компанії NVIDIAналежить пройти ще довгий шлях по
просування в маси свого API (те ж можна сказати і про дітище ATI/AMD).
Якою вона буде, покаже майбутнє. Отже, CUDA will be back:).

P.S. Початківцям програмістам і людям, що зацікавилися, рекомендую відвідати
наступні «віртуальні заклади»:

офіційний сайт NVIDIA та сайт
GPGPU.com. Вся
дана інформація - англійською мовою, але, дякую хоча б, що не на
китайською. Так що дерзай! Сподіваюся, що автор хоч трохи допоміг тобі у
захоплюючих починаннях пізнання CUDA!




Поділися статтею:
Схожі статті

17 квітня 2015
Як додати музику до ГТА Сан Андреас?

17 квітня 2015
Що робити, якщо GTA IV не запускається

17 квітня 2015
Інстаграм у цифрах - корисні дані, факти та фішки

17 квітня 2015
Зміна май модему щоб визначався як смартфон