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

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

Технологія 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.

- Набір низькорівневих програмних інтерфейсів ( API) для створення ігор та інших високопродуктивних мультимедіа-додатків. Включає підтримку високопродуктивної 2D- І 3D-графіки, звуку та пристроїв введення.

Direct3D (D3D) - інтерфейс виведення тривимірних примітивів(Геометричних тіл). Входить в .

OpenGL(Від англ. Open Graphics Library, Дослівно - відкрита графічна бібліотека) - специфікація, що визначає незалежний від мови програмування крос-платформний програмний інтерфейс для написання додатків, що використовують двовимірну та тривимірну комп'ютерну графіку. Включає понад 250 функцій для малювання складних тривимірних сцен із простих примітивів. Використовується для створення відеоігор, віртуальної реальності, візуалізації в наукових дослідженнях. На платформі Windowsконкурує з .

OpenCL(Від англ. Open Computing Language, дослівно – відкрита мова обчислень) – фреймворк(Каркас програмної системи) для написання комп'ютерних програм, пов'язаних з паралельними обчисленнями на різних графічних ( GPU) та ( ). У фреймворк OpenCLвходять мову програмування та інтерфейс програмування додатків ( API). OpenCLзабезпечує паралелізм на рівні інструкцій та на рівні даних та є реалізацією техніки GPGPU.

GPGPU(скор. від англ. General-P urpose G raphics P rokussing U nits, дослівно – GPUзагального призначення) - техніка використання графічного процесора відеокарти для загальних обчислень, які зазвичай проводить .

Шейдер(Англ. shader) – програма побудови тіней на синтезованих зображеннях, використовують у тривимірної графіці визначення остаточних параметрів об'єкта чи зображення. Як правило, включає довільної складності опис поглинання та розсіювання світла, накладання текстури, відображення та заломлення, затінювання, зміщення поверхні та ефекти пост-обробки. Складні поверхні можуть бути візуалізовані за допомогою простих геометричних форм.

Рендеринг(Англ. rendering) – візуалізація, у комп'ютерній графіці процес отримання зображення за моделлю за допомогою програмного .

SDK(скор. від англ. Software Development Kit) - Набір інструментальних засобів розробки програмного забезпечення.

CPU(скор. від англ. Central Processing Unit, дослівно - центральний/основний/головний обчислювальний пристрій) - центральний (мікро); пристрій, що виконує машинні інструкції; частина апаратного забезпечення, що відповідає за виконання обчислювальних операцій (заданих операційною системою та прикладним програмним) і координує роботу всіх пристроїв.

GPU(скор. від англ. Graphic Processing Unit, дослівно - графічний обчислювальний пристрій) - графічний процесор; окремий пристрій або ігрової приставки, що виконує графічний рендеринг (візуалізацію). Сучасні графічні процесори дуже ефективно обробляють і реалістично відображають комп'ютерну графіку. Графічний процесор у сучасних відеоадаптерах застосовується як прискорювач тривимірної графіки, проте його можна використовувати в деяких випадках і для обчислень ( GPGPU).

Проблеми CPU

Довгий час підвищення продуктивності традиційних переважно відбувалося рахунок послідовного збільшення тактової частоти (близько 80% продуктивності визначала саме тактова частота) з одночасним збільшенням кількості транзисторів однією кристалі. Однак подальше підвищення тактової частоти (при тактовій частоті більше 3,8 ГГц чіпи просто перегріваються!) впирається в ряд фундаментальних фізичних бар'єрів (оскільки технологічний процес майже впритул наблизився до розмірів атома: , А розміри атома кремнію – приблизно 0,543 нм):

По-перше, зі зменшенням розмірів кристала і підвищенням тактової частоти зростає струм витоку транзисторів. Це веде до підвищення споживаної потужності та збільшення викиду тепла;

По-друге, переваги вищої тактової частоти частково зводяться нанівець через затримки при зверненні до пам'яті, так як час доступу до пам'яті не відповідає зростаючим тактовим частотам;

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

Розвиток GPU

Паралельно з йшло (і йде!) розвиток GPU:

…

Листопад 2008 р. – Intelпредставила лінійку 4-ядерних Intel Core i7, в основу яких покладено мікроархітектуру нового покоління Nehalem. Процесори працюють на тактовій частоті 26-32 ГГц. Виконані за 45-нм техпроцесом.

Грудень 2008 р. – розпочалися поставки 4-ядерного AMD Phenom II 940(кодова назва – Deneb). Працює на частоті 3 ГГц, випускається за техпроцесом 45-нм.

…

Травень 2009 р. – компанія AMDпредставила версію графічного процесора ATI Radeon HD 4890із тактовою частотою ядра, збільшеною з 850 МГц до 1 ГГц. Це перший графічнийпроцесор, працюючий на частоті 1 ГГц. Обчислювальна потужність чіпа завдяки збільшенню частоти зросла з 1,36 до 1,6 терафлоп. Процесор містить 800 (!) обчислювальних ядер, підтримує відеопам'ять GDDR5, DirectX 10.1, ATI CrossFireXта всі інші технології, властиві сучасним моделям відеокарт. Чіп виготовлений на базі 55-нм технології.

Основні відмінності GPU

Відмінними рисами GPU(порівняно з ) є:

- архітектура, максимально націлена на збільшення швидкості розрахунку текстур та складних графічних об'єктів;

- пікова потужність типового GPUнабагато вище, ніж у ;

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

«Криза жанру»

«Криза жанру» для назріло до 2005 р., – саме тоді з'явилися . Але, незважаючи на розвиток технології, зростання продуктивності звичайних помітно знизився. Водночас продуктивність GPUпродовжує зростати. Так, до 2003 р. і кристалізувалась ця революційна ідея – використовувати для потреб обчислювальну міць графічного. Графічні процесори стали активно використовуватися для «неграфічних» обчислень (симуляція фізики, обробка сигналів, обчислювальна математика/геометрія, операції з базами даних, обчислювальна біологія, обчислювальна економіка, комп'ютерний зір тощо).

Головна проблема в тому, що не було жодного стандартного інтерфейсу для програмування GPU. Розробники використовували OpenGLабо Direct3Dале це було дуже зручно. Корпорація NVIDIA(один з найбільших виробників графічних, медіа- та комунікаційних процесорів, а також бездротових медіа-процесорів; заснована у 1993 р.) зайнялася розробкою якогось єдиного та зручного стандарту, – і представила технологію CUDA.

Як це починалося

2006 р. – NVIDIAдемонструє CUDA™; початок революції у обчисленнях на GPU.

2007 р. – NVIDIAвипускає архітектуру CUDA(Початкова версія CUDA SDKбула представлена ​​15 лютого 2007 р.); номінація « Найкраща новинка» від журналу Popular Scienceта «Вибір читачів» від видання HPCWire.

2008 р. – технологія NVIDIA CUDAперемогла у номінації «Технічна перевага» від PC Magazine.

Що таке CUDA

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

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

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

Уральський, провідний спеціаліст з технологій NVIDIA, порівнюючи GPUі , каже так: « - Це позашляховик. Він їздить завжди та скрізь, але не дуже швидко. А GPU- Це спорткар. На поганій дорозі він просто нікуди не поїде, але дайте хороше покриття – і він покаже всю свою швидкість, яка позашляховику і не снилася!..».

Можливості технології CUDA

Пристрої перетворення персональних комп'ютерів на маленькі суперкомп'ютери відомі досить давно. Ще у 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):

• Флуоресцентна мікроскопія: 12x;
• Молекулярна динаміка (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. Також підтримкою CU всі 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.

Дозвольте звернутися до історії - повернутися в 2003 рік, коли Intel та AMD брали участь у спільних перегонах за найпотужніший процесор. Усього за кілька років в результаті цих перегонів тактові частоти суттєво зросли, особливо після виходу Intel Pentium 4.

Але гонка швидко наближалася до краю. Після хвилі величезного приросту тактових частот (між 2001 і 2003 роками тактова частота Pentium 4 подвоїлася з 1,5 до 3 ГГц), користувачам довелося задовольнятися десятими частками гігагерц, які змогли вичавити виробники (з 2003 до 2005 8 ГГц).

Навіть архітектури, оптимізовані під високі тактові частоти, та сама Prescott, стали відчувати труднощі, причому цього разу не тільки виробничі. Виробники чіпів просто вперлися у закони фізики. Деякі аналітики навіть пророкували, що закон Мура перестане діяти. Але цього не сталося. Оригінальний сенс закону часто спотворюють, проте стосується кількості транзисторів на поверхні кремнієвого ядра. Довгий час підвищення числа транзисторів у CPU супроводжувалося відповідним зростанням продуктивності - що призвело до спотворення сенсу. Але потім ситуація ускладнилася. Розробники архітектури CPU підійшли до закону скорочення приросту: число транзисторів, яке потрібно додати для необхідного збільшення продуктивності, ставало все більшим, заводячи в глухий кут.

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

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

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

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

Хай живе GeForce FX!

Два світи тривалий час залишалися розділеними. Ми використовували CPU (або навіть кілька CPU) для офісних завдань та інтернет-застосунків, а GPU добре підходили лише для прискорення візуалізації. Але одна особливість змінила все: зокрема, поява програмованих GPU. Спочатку центральним процесорам не було чого боятися. Перші так звані програмовані GPU (NV20 та R200) ​​навряд чи становили загрозу. Число інструкцій у програмі залишалося обмеженим близько 10, вони працювали над екзотичними типами даних, такими як 9 або 12-бітними числами з фіксованою комою.

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

• Підтримка обчислень з плаваючою комою одинарної точності (нехай це навіть не відповідало стандарту IEEE754);
• підтримка числа інструкцій понад тисячу.

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

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

У 2003 році, з появою шейдерів, що еволюціонували, була досягнута нова планка - цього разу виконання матричних обчислень. Це був рік, коли цілу секцію SIGGRAPH ("Computations on GPUs/Обчислення на GPU") було виділено під нову область ІТ. Ця рання ініціатива дістала назву GPGPU (General-Purpose computation on GPU, універсальні обчислення на GPU). І раннім поворотним моментом стала поява.

Щоб зрозуміти роль BrookGPU, необхідно розібратися, як все відбувалося до появи. Єдиним способом отримати ресурси GPU у 2003 році було використання одного з двох графічних API – Direct3D або OpenGL. Отже, розробникам, які хотіли отримати можливості GPU для своїх обчислень, доводилося спиратися на два згадані API. Проблема в тому, що вони не завжди були експертами у програмуванні відеокарт, а це серйозно ускладнювало доступ до технологій. Якщо 3D-програмісти оперують шейдерами, текстурами та фрагментами, то фахівці в галузі паралельного програмування спираються на потоки, ядра, розкиди тощо. Тому спочатку треба було навести аналогії між двома світами.

• Потік (stream)є потік елементів одного типу, в GPU він може бути представлений текстурою. У принципі, у класичному програмуванні є такий аналог як масив.
• Ядро (Kernel)- функція, яка застосовуватиметься незалежно до кожного елемента потоку; є еквівалентом піксельного шейдера. У класичному програмуванні можна навести аналогію циклу - він застосовується до великої кількості елементів.
• Щоб зчитувати результати застосування ядра до потоку, має бути створена текстура. На CPU еквівалента немає, бо там є повний доступдо пам'яті.
• Управління місцезнаходженням у пам'яті, куди робитиметься запис (в операціях розкиду/scatter), здійснюється через вершинний шейдер, оскільки піксельний шейдер не може змінювати координати пікселя, що обробляється.

Як можна бачити, навіть з урахуванням наведених аналогій, завдання не виглядає простим. І на допомогу прийшов Brook. Під цією назвою маються на увазі розширення до мови C ("C with streams", "C з потоками"), як назвали їх розробники у Стенфорді. По суті, завдання Brook зводилася до приховування від програміста всіх складових 3D API, що дозволяло уявити GPU як співпроцесор для паралельних обчислень. Для цього компілятор Brook обробляв файл.br із кодом C++ та розширеннями, після чого генерував код C++, який прив'язувався до бібліотеки з підтримкою різних виходів (DirectX, OpenGL ARB, OpenGL NV3x, x86).

Brook має кілька заслуг, перша з яких полягає у виведенні GPGPU з тіні, щоб з цією технологією могли знайомитися і широкі маси. Хоча після оголошення про проект ряд ІТ-сайтів дуже оптимістично повідомив про те, що вихід Brook ставить під сумнів існування CPU, які незабаром будуть замінені на більш потужні GPU. Але, як бачимо, і за п'ять років цього не сталося. Чесно кажучи, ми не думаємо, що це взагалі колись станеться. З іншого боку, дивлячись на успішну еволюцію CPU, які все більше орієнтуються у бік паралелізму (більше ядер, технологія багатопоточності SMT, розширення блоків SIMD), а також і на GPU, які, навпаки, стають все більш універсальними (підтримка розрахунків із плаваючою комою) одинарної точності, цілочисленні обчислення, підтримка розрахунків з подвійною точністю), схоже, що GPU та CPU незабаром просто зіллються. Що ж тоді станеться? Чи будуть GPU поглинені CPU, як свого часу сталося з математичними співпроцесорами? Цілком можливо. Intel та AMD сьогодні працюють над подібними проектами. Але ще дуже багато може змінитися.

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

Одна з проблем пов'язана з різними рівнями абстракції, а також, зокрема, з надмірним додатковим навантаженням, яке створюється 3D API, яке може бути дуже відчутним. Але серйознішою можна вважати проблему сумісності, з якою розробники Brook нічого не могли зробити. Між виробниками GPU існує жорстка конкуренція, тому вони часто оптимізують свої драйвери. Якщо такі оптимізації, здебільшого, хороші для геймерів, вони можуть відразу покінчити з сумісністю Brook. Тому складно уявити використання цього API у промисловому коді, який десь працюватиме. І довгий час Brook залишався долею дослідників-аматорів та програмістів.

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

Дослідники, які спочатку залучені до проекту Brook, швидко приєдналися до команд розробників у Санта-Кларі, щоб представити глобальну стратегію для розвитку нового ринку. Ідея полягала у створенні комбінації апаратного та програмного забезпечення, що підходить для завдань GPGPU. Оскільки розробники nVidia знають усі секрети своїх GPU, то графічне API можна було й не спиратися, а зв'язуватися з графічним процесором через драйвер. Хоча, звісно, ​​у своїй виникають свої проблеми. Команда розробників CUDA (Compute Unified Device Architecture) створила набір програмних рівнів для роботи з GPU.

Як можна побачити на діаграмі, CUDA забезпечує два API.

• Високорівневий API: CUDA Runtime API;
• низькорівневий API: CUDA Driver API.

Оскільки високорівневий API реалізований над низькорівневим, кожен виклик функції рівня Runtime розбивається на простіші вказівки, які обробляє Driver API. Зверніть увагу, що два API взаємно виключають один одного: програміст може використовувати один або інший API, але змішувати дзвінки функцій двох API не вдасться. Взагалі термін "високорівневий API" відносний. Навіть Runtime API такий, що багато хто вважає його низькорівневим; втім, він все ж таки надає функції, дуже зручні для ініціалізації або управління контекстом. Але не чекайте особливо високого рівня абстракції - вам все одно потрібно мати гарний набір знань про nVidia GPU і про те, як вони працюють.

З Driver API працювати ще складніше; для запуску обробки на GPU вам знадобиться більше зусиль. З іншого боку, низькорівневий API більш гнучкий, надаючи програмісту додатковий контроль, якщо потрібно. Два API здатні працювати з ресурсами OpenGL або Direct3D (тільки дев'ята версія на сьогодні). Користь від такої можливості очевидна - CUDA може використовуватися для створення ресурсів (геометрія, процедурні текстури тощо), які можна передати на графічне API або, навпаки, можна зробити так, що 3D API надсилатиме результати рендерингу програмі CUDA, яка, у свою чергу, виконуватиме пост-обробку. Є багато прикладів таких взаємодій і перевага полягає в тому, що ресурси продовжують зберігатися в пам'яті GPU, їх не потрібно передавати через шину PCI Express, яка, як і раніше, залишається "вузьким місцем".

Втім, слід зазначити, що спільне використання ресурсів у відеопам'яті не завжди проходить ідеально і може призвести до деяких головних болів. Наприклад, при зміні роздільної здатності або глибини кольору, графічні дані є пріоритетними. Тому якщо потрібно збільшити ресурси в кадровому буфері, то драйвер без проблем зробить це рахунок ресурсів додатків CUDA, які просто "вилетять" з помилкою. Звичайно, не дуже елегантно, але така ситуація не повинна траплятися дуже часто. І якщо вже ми почали говорити про недоліки: якщо ви хочете використовувати кілька GPU для програм CUDA, то вам потрібно спочатку відключити режим SLI, інакше програми CUDA зможуть "бачити" лише один GPU.

Нарешті, третій програмний рівень відданий бібліотекам - двом, якщо точним.

• CUBLAS де є необхідні блоки для обчислень лінійної алгебри на GPU;
• CUFFT, яка підтримує розрахунок перетворень Фур'є - алгоритм, який широко використовується в області обробки сигналів.

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

Другий термін, який часто зустрічається в документації CUDA - варп (warp). Тут плутанини немає, оскільки в російській аналога не існує (хіба що ви не є фанатом Start Trek або ігри Warhammer). Насправді термін взятий із текстильної промисловості, де через основну пряжу (warp yarn), яка розтягнута на верстаті, простягається уточна пряжа (weft yarn). Варп в CUDA є групою з 32 потоків і є мінімальним обсягом даних, оброблюваних SIMD-способом в мультипроцесорах CUDA.

Але подібна "зернистість" який завжди зручна для програміста. Тому в CUDA, замість роботи з варпами безпосередньо, можна працювати з блоками/block, що містять від 64 до 512 потоків

Нарешті, ці блоки збираються разом у сітки/grid. Перевага подібного угруповання полягає в тому, що число блоків, що одночасно обробляються GPU, тісно пов'язане з апаратними ресурсами, як ми побачимо нижче. Угруповання блоків у сітки дозволяє повністю абстрагуватися від цього обмеження та застосувати ядро/kernel до більшого числа потоків за один виклик, не думаючи про фіксовані ресурси. За це відповідають бібліотеки CUDA. Крім того, подібна модель добре масштабується. Якщо GPU має мало ресурсів, він буде виконувати блоки послідовно. Якщо число обчислювальних процесорів велике, блоки можуть виконуватися паралельно. Тобто, той самий код може працювати на GPU як початкового рівня, так і на топових і навіть майбутніх моделях.

Є ще пара термінів у CUDA API, які позначають CPU ( хост/host) та GPU ( пристрій/device). Якщо це невелике введення вас не злякало, настав час ближче познайомитися з CUDA.

Якщо ви регулярно читаєте Tom"s Hardware Guide, то архітектура останніх GPU від nVidia вам знайома. Якщо ні, ми рекомендуємо ознайомитись зі статтею " nVidia GeForce GTX 260 та 280: нове покоління відеокартЩо стосується CUDA, то nVidia представляє архітектуру дещо по-іншому, демонструючи деякі деталі, що раніше залишалися прихованими.

Як можна бачити по ілюстрації вище, ядро ​​шейдерів nVidia складається з кількох кластерів текстурних процесорів (Texture Processor Cluster, TPC). Відеокарта 8800 GTX, наприклад, використала вісім кластерів, 8800 GTS – шість і т.д. Кожен кластер, по суті, складається з текстурного блоку та двох потокових мультипроцесорів (streaming multiprocessor). Останні включають початок конвеєра (front end), що виконує читання та декодування інструкцій, а також відсилання їх на виконання, і кінець конвеєра (back end), що складається з восьми обчислювальних пристроїв та двох суперфункціональних пристроїв SFU (Super Function Unit)де інструкції виконуються за принципом SIMD, тобто одна інструкція застосовується до всіх потоків у варпі. nVidia називає такий спосіб виконання SIMT(Single instruction multiple threads, одна інструкція, багато потоків). Важливо відзначити, що кінець конвеєра працює на частоті вдвічі перевершує його початок. Насправді це означає, що ця частина виглядає вдвічі "ширше", ніж вона є насправді (тобто як 16-канальний блок SIMD замість восьмиканального). Потокові мультипроцесори працюють так: кожен такт початок конвеєра вибирає варп, готовий до виконання, і запускає виконання інструкції. Щоб інструкція застосувалася до всіх 32 потоків у варпі, кінцю конвеєра знадобиться чотири такти, але оскільки він працює на подвоєній частоті в порівнянні з початком, знадобиться лише два такти (з точки зору початку конвеєра). Тому, щоб початок конвеєра не простоювало такт, а апаратне забезпечення було максимально завантажено, в ідеальному випадку можна чергувати інструкції кожен такт – класична інструкція в один такт та інструкція для SFU – в інший.

Кожен мультипроцесор має певний набір ресурсів, у яких варто розібратися. Є невелика область пам'яті під назвою "Спільна пам'ять/Shared Memory", по 16 кбайт на мультипроцесор Це аж ніяк не кеш-пам'ять: програміст може використовувати її на власний розсуд. Тобто перед нами щось близьке до Local Store у SPU на процесорах Cell. Ця деталь дуже цікава, оскільки вона підкреслює, що CUDA - це комбінація програмних та апаратних технологій. Ця область пам'яті не використовується для піксельних шейдерів, що nVidia дотепно підкреслює "нам не подобається, коли пікселі розмовляють один з одним".

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

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

Мультипроцесор має 8 192 регістри, які загальні всім потоків всіх блоків, активних на мультипроцессоре. Число активних блоків на мультипроцесор не може перевищувати восьми, а кількість активних варпів обмежена 24 (768 потоків). Тому 8800 GTX може обробляти до 12288 потоків в один момент часу. Всі ці обмеження варто згадати, оскільки вони дозволяють оптимізувати алгоритм залежно від доступних ресурсів.

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

З програмної точки зору CUDA складається із набору розширень до мови C, що нагадує BrookGPU, а також кількох специфічних викликів API. Серед розширень присутні специфікатори типу, що належать до функцій та змінних. Важливо запам'ятати ключове слово __global__, яке, будучи наведеним перед функцією, показує, що остання відноситься до ядра/kernel - цю функцію викликатиме CPU, а виконуватиметься вона на GPU. Префікс __device__вказує, що функція буде виконуватися на GPU (який, до речі, CUDA і називає "пристрій/device"), але вона може бути викликана тільки з GPU (іншими словами, з іншої функції __device__ або з функції __global__). Зрештою, префікс __host__опціональний, він позначає функцію, яка викликається CPU і виконується CPU - тобто звичайну функцію.

Є ряд обмежень, пов'язаних з функціями __device__ і __global__: вони не можуть бути рекурсивними (тобто викликати себе), і не можуть мати змінну кількість аргументів. Нарешті, оскільки функції __device__ розташовуються у просторі пам'яті GPU, цілком логічно, що отримати їхню адресу не вдасться. Змінні теж мають низку кваліфікаторів, які вказують на область пам'яті, де вони зберігатимуться. Змінна з префіксом __shared__означає, що вона зберігатиметься у спільній пам'яті потокового мультипроцесора. Виклик функції __global__ трохи відрізняється. Справа в тому, при виклику потрібно задати конфігурацію виконання - більш конкретно розмір сітки/grid, до якої буде застосовано ядро/kernel, а також розмір кожного блоку. Візьмемо, наприклад, ядро ​​з наступним підписом.

Воно буде викликатись у вигляді

де Dg є розміром сітки, а Db – розміром блоку. Дві ці змінні відносяться до нового типу вектора, що з'явився з CUDA.

API CUDA містить функції для роботи з пам'яттю у VRAM: cudaMalloc для виділення пам'яті, cudaFree для звільнення та cudaMemcpy для копіювання пам'яті між RAM та VRAM та навпаки.

Ми закінчимо цей огляд дуже цікавим способом, яким компілюється програма CUDA: компіляція виконується у кілька етапів. Спочатку виймається код, що відноситься до CPU, який передається стандартному компілятору. Код, призначений для GPU, спочатку перетворюється на проміжну мову PTX. Він подібний до асемблеру і дозволяє вивчати код у пошуках потенційних неефективних ділянок. Нарешті, остання фаза полягає у трансляції проміжної мови в специфічні команди GPU та створення двійкового файлу.

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

У подібний проект поринути досить легко. Сьогодні для завантаження доступна велика кількість безкоштовних, але якісних інструментів. Для нашого тесту ми використовували Visual C++ Express 2005 де є все необхідне. Найскладніше полягати в тому, щоб знайти програму, портування якої на GPU не зайняло б кілька тижнів, і водночас вона була б досить цікавою, щоб наші зусилля не пропали даремно. Зрештою, ми вибрали відрізок коду, який бере карту висот та розраховує відповідну карту нормалей. Ми не детально заглиблюватимемося в цю функцію, оскільки в даній статті це навряд чи цікаво. Якщо бути коротким, то програма займається викривленням ділянок: для кожного пікселя початкового зображення ми накладаємо матрицю, що визначає колір результуючого пікселя в зображенні прилеглих пікселів, використовуючи більш-менш складну формулу. Перевага цієї функції в тому, що її дуже легко розпаралелити, тому даний тестчудово показує можливості CUDA.

Ще одна перевага полягає в тому, що у нас вже є реалізація на CPU, тому ми можемо порівнювати її результат із версією CUDA – і не винаходити колесо наново.

Ще раз повторимо, що метою тесту було знайомство з утилітами CUDA SDK, а не порівняльне тестування версій під CPU та GPU. Оскільки це була перша наша спроба створення програми CUDA, ми не дуже сподівалися отримати високу продуктивність. Так як ця частина коду не є критичною, то версія під CPU була не оптимізована, тому пряме порівняння результатів навряд чи цікаве.

Продуктивність

Однак ми заміряли час виконання, щоб подивитися, чи є перевага у використанні CUDA навіть з найбрутальнішою реалізацією, чи нам буде потрібно тривала та стомлююча практика, щоб отримати якийсь виграш при використанні GPU. Тестова машина була взята з нашої лабораторії розробки – ноутбук з процесором Core 2 Duo T5450 та відеокартою GeForce 8600M GT, що працює під Vista. Це далеко не суперкомп'ютер, але результати дуже цікаві, оскільки тест не заточений під GPU. Завжди приємно бачити, коли nVidia демонструє величезний приріст на системах з монстроподібними GPU та чималою пропускною спроможністю, але на практиці багато хто з 70 мільйонів GPU з підтримкою CUDA на сучасному ринку ПК далеко не такі потужні, тому і наш тест має право на життя.

Для зображення 2 048 x 2 048 пікселів ми отримали такі результати.

• CPU 1 потік: 1419 мс;
• CPU 2 потоку: 749 мс;
• CPU 4 потоки: 593 мс
• GPU (8600M GT) блоки по 256 потоків: 109 мс;
• GPU (8600M GT) блоки по 128 потоків: 94 мс;
• GPU (8800 GTX) блоки по 128 потоків/256 потоків: 31 мс.

За наслідками можна зробити кілька висновків. Почнемо з того, що, незважаючи на розмови про очевидну лінь програмістів, ми модифікували початкову версію CPU під кілька потоків. Як ми вже згадували, код ідеальний для цієї ситуації - все, що потрібно, це розбити початкове зображення на стільки зон, скільки потоків. Зверніть увагу, що від переходу від одного потоку на два на нашому двоядерному CPU прискорення вийшло майже лінійним, що також вказує на паралельну природу тестової програми. Дуже несподівано, але версія із чотирма потоками теж виявилася швидше, хоча на нашому процесорі це дуже дивно - можна було, навпаки, очікувати падіння ефективності через накладні витрати на управління додатковими потоками. Як можна пояснити такий результат? Важко сказати, але, можливо, винний планувальник потоків під Windows; у будь-якому разі, результат повторюємо. З текстурами меншого розміру (512x512) приріст від поділу на потоки був не такий виражений (приблизно 35% проти 100%), і поведінка версії з чотирма потоками була логічніше, без приросту порівняно з версією на два потоки. GPU працював все ще швидше, але вже не так виражено (8600M GT була втричі швидше ніж версія з двома потоками).

Друге значуще спостереження - навіть найповільніша реалізація GPU виявилася майже вшестеро швидше, ніж найпродуктивніша версія CPU. Для першої програми та неоптимізованої версії алгоритму результат дуже навіть підбадьорливий. Зверніть увагу, що ми отримали відчутно кращий результатна невеликих блоках, хоча інтуїція може підказувати про інше. Пояснення просте - наша програма використовує 14 регістрів на потік, і з 256-потоковими блоками потрібно 3584 регістру на блок, а для повного навантаження процесора потрібно 768 потоків, як ми показували. У нашому випадку це становить три блоки або 10572 регістру. Але мультипроцесор має всього 8192 регістру, тому він може підтримувати активними тільки два блоки. Навпаки, з блоками по 128 потоків нам потрібно 1792 регістра на блок; якщо 8192 поділити на 1792 і округлити до найближчого цілого, то ми отримаємо чотири блоки. На практиці число потоків буде таким самим (512 на мультипроцесор, хоча для повного навантаження теоретично потрібно 768), але збільшення числа блоків дає GPU перевагу гнучкості по доступу до пам'яті - коли йде операція з великими затримками, то можна запустити виконання інструкцій іншого блоку, очікуючи надходження результатів. Чотири блоки явно знижують затримки, особливо з огляду на те, що наша програма використовує кілька доступів на згадку.

Аналіз

Нарешті, незважаючи на те, що ми сказали вище, ми не змогли встояти перед спокусою і запустили програму на 8800 GTX, яка виявилася втричі швидше за 8600, незалежно від розміру блоків. Можна подумати, що на практиці на відповідних архітектурах результат буде в чотири або більше разів вищий: 128 АЛУ/шейдерних процесорів проти 32 і вища тактова частота (1,35 ГГц проти 950 МГц), але так не вийшло. Швидше за все, обмежуючим фактором виявився доступ до пам'яті. Якщо бути точнішим, доступ до початкового зображення здійснюється як до багатовимірного масиву CUDA - дуже складний термін для того, що є не більш ніж текстурою. Але є кілька переваг.

• доступи виграють від кешу текстури;
• ми використовуємо wrapping mode, в якому не потрібно обробляти межі зображення на відміну від версії CPU.

Крім того, ми можемо отримати перевагу від "безкоштовної" фільтрації з нормалізованою адресацією між замість і, але в нашому випадку це навряд чи корисно. Як ви знаєте, 8600 оснащений 16 текстурними блоками порівняно з 32 у 8800 GTX. Тому між двома архітектурами співвідношення лише два до одного. Додайте до цього різницю в частотах, і ми отримаємо співвідношення (32 x 0,575)/(16 x 0,475) = 2,4 - близько до "трьома до одного", що ми отримали насправді. Ця теорія також пояснює, чому розмір блоків багато на G80 не змінює, оскільки АЛУ все одно упирається в текстурні блоки.

Окрім перспективних результатів, наше перше знайомство з CUDA пройшло дуже добре, враховуючи не найсприятливіші вибрані умови. Розробка на ноутбуці під Vista має на увазі, що доведеться використовувати CUDA SDK 2.0, що все ще знаходиться в стані бета-версії, з драйвером 174.55, який теж бета-версія. Незважаючи на це ми не можемо повідомити про якісь неприємні сюрпризи - тільки початкові помилки під час першої налагодження, коли наша програма, все ще дуже "глючна", спробувала адресувати пам'ять за межами виділеного простору.

Монітор почав дико мерехтіти, потім екран почорнів... поки Vista не запустила службу відновлення драйвера, і все стало гаразд. Але все ж таки трохи дивно це спостерігати, якщо ви звикли бачити типову помилку Segmentation Fault на стандартних програмах, подібно до нашої. Нарешті, невелика критика у бік nVidia: у всій документації, доступної для CUDA, немає невеликого керівництва, яке крок за кроком розповідало про те, як налаштувати оточення розробки під Visual Studio. Власне, проблема невелика, оскільки у SDK є повний набірприкладів, які можна вивчити для розуміння каркасу для програм CUDA, але керівництво для новачків не завадило б.

nVidia представила CUDA з випуском GeForce 8800. І в той час обіцянки здавались дуже спокусливими, але ми притримали свій ентузіазм до реальної перевірки. Справді, на той час це здавалося більше розміткою території, щоб залишатися на хвилі GPGPU. Без доступного SDK складно сказати, що перед нами не чергова маркетингова пустушка, з якої нічого не вийде. Вже не вперше хороша ініціатива була оголошена дуже рано і на той час не вийшла на світ через брак підтримки - особливо в такому конкурентному секторі. Тепер, через півтора роки після оголошення, ми з упевненістю можемо сказати, що nVidia дотрималася слова.

SDK досить швидко з'явився в бета-версії на початку 2007 року, з того часу він швидко оновлювався, що доводить важливість цього проекту для nVidia. Сьогодні CUDA дуже приємно розвивається: SDK доступний вже у бета-версії 2.0 для основних операційних систем (Windows XP і Vista, Linux, а також 1.1 для Mac OS X), а для розробників nVidia виділила цілий розділ сайту.

На більш професійному рівні враження від перших кроків з CUDA виявилося дуже позитивним. Якщо ви навіть знайомі з архітектурою GPU, ви легко розберетеся. Коли API виглядає зрозумілим з першого погляду, то відразу починаєш вважати, що отримаєш переконливі результати. Але чи не втрачатиметься обчислювальний час від численних передач із CPU на GPU? І як використовувати ці тисячі потоків майже без примітива синхронізації? Ми розпочинали наші експерименти з усіма цими побоюваннями в умі. Але вони швидко розвіялися, коли перша версія нашого алгоритму, нехай і досить тривіального, виявилася значно швидшою, ніж на CPU.

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

Що ж потрібно CUDA, щоб стати відомим API? Якщо казати одним словом: переносимість. Ми знаємо, що майбутнє ІТ криється в паралельних обчисленнях – сьогодні вже кожен готується до подібних змін, і всі ініціативи, як програмні, так і апаратні, спрямовані у цьому напрямі. Однак на даний момент, якщо дивитися на розвиток парадигм, ми ще на початковому етапі: ми створюємо потоки вручну і намагаємося спланувати доступ до загальних ресурсів; з усім цим ще якось можна впоратися, якщо кількість ядер можна перерахувати на пальцях однієї руки. Але через кілька років, коли кількість процесорів обчислюватиметься сотнями, такої можливості вже не буде. З випуском CUDA nVidia зробила перший крок у вирішенні цієї проблеми – але, звичайно, дане рішенняпідходить тільки для GPU від цієї компанії, та й то не всім. Тільки GF8 та 9 (і їх похідні Quadro/Tesla) сьогодні можуть працювати з програмами CUDA. І нова лінійка 260/280, звісно.

nVidia може хвалитися тим, що продала 70 мільйонів CUDA-сумісних GPU у всьому світі, але цього все одно мало, щоб стати стандартом де-факто. З урахуванням того, що конкуренти не сидять, склавши руки. AMD пропонує свій SDK (Stream Computing), та й Intel оголосила про рішення (Ct), хоча воно ще не доступне. Настає війна стандартів, і на ринку явно не буде місця для трьох конкурентів, поки інший гравець, наприклад Microsoft, не вийде з пропозицією загального API, що, звичайно, полегшить життя розробникам.

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

Технологія 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 Кбайт (на весь пристрій). З цього випливає, що в контекстній пам'яті має сенс зберігати лише невелику кількість даних, що часто використовуються.

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