Технологія CUDA

Размещено на http://www.allbest.ru/

27

Размещено на http://www.allbest.ru/

Міністерство освіти і науки молоді та спорту України

Прикарпатський національний університет імені Василя Стефаника

Факультет математики та інформатики

Кафедра інформатики

КУРСОВА РОБОТА

на тему:

Технологія CUDA

Виконав:

Петрушка Тарас

Івано-Франківськ, 2012

Зміст

ВСТУП

РОЗДІЛ 1. Загальна теорія та поняття

1.1 Загальна термінологія CUDA

1.2 Структура NVIDIA CUDA

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

1.4 Різниця між GPU і CPU

1.5 Програмна модель та модель пам'яті CUDA

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

1.7 Проблеми CUDA

РОзділ 2. Модель програмування в СUDA

2.1 Основні поняття

2.2 Розширення мови С++

2.2.1 Специфікатори функцій й змінні

2.2.2 Додані змінні

2.2.3 Директива виклику ядра

2.2.4 Додані функції

РОЗДІЛ 3. ДЕМОНСТРАЦІЯ ТЕХНОЛОГІЇ CUDA НА ПРИКЛАДІ ПІДРАХУНКУ CRC32-КОДУ

ВИСНОВКИ

ЛІТЕРАТУРА

ДОДАТОК А. Диск з програмою та програмним забезпеченням

ДОДАТОК В. Мінімальні вимоги до програми

Вступ

Що таке CUDA ?

CUDA (англ. Compute Unified Device Architecture) - програмно-апаратна архітектура, яка дозволяє робити обчислення з використанням графічних процесорів NVIDIA, що підтримують технологію GPGPU (англ. General-Purpose computing on Graphics Processing Units) - обчислення загального призначення на графічних процесорах. Це архітектура паралельних обчислень від NVIDIA, що дозволяє істотно збільшити обчислювальну продуктивність завдяки використанню GPU (графічних процесорів).

Напрямок обчислень еволюціонує від «централізованої обробки даних» на центральному процесорі до «спільної обробки» на CPU і GPU. Для реалізації нової обчислювальної парадигми компанія NVIDIA винайшла архітектуру паралельних обчислень CUDA, на даний момент представлену в графічних процесорах GeForce, ION, Quadro, Tesla і забезпечує необхідну базу розробникам ПЗ.

З кожним днем зростає число завдань, пов'язаних з обробкою великої кількості інформації. Підвищена складність багатьох завдань, серед яких завдання планування і організації перевезеннями, вантажної роботи для транзитних і місцевих поїздів, стимулює до користування потужного апаратного забезпечення. Однак, враховуючи сьогоднішні економічні умови, необхідно грамотно підходити до використання великого числа електричної енергії, витрат на техобслуговування такого потужного апаратного забезпечення, а також витрат на утримання великих інформаційно-обчислювальних центрів. Найбільш практичним і економічно вигідним способом нарощування обчислювальних потужностей є організація паралельних обчислень, з використанням відповідних алгоритмів і функцій, зумовленими класом вирішуваних завдань. Для організації ефективного розпаралелювання обчислень необхідний перехід на нові технології та нові паралельні методи вирішення задач, які дозволять різко знизити вартість обчислень. Практично всі сегменти ринку напівпровідникової продукції, включаючи персональні комп'ютери, ігрові консолі, мобільні пристрої, сервери, суперкомп'ютери і мережеві пристрої переходять до використання паралельних платформ. Є дві основні причини. По-перше, паралельні процесори надають більш ефективне використання доступної площі кристала і бюджету енергоспоживання для багатьох вимогливих додатків. По-друге, велика кількість завдань, які традиційно вирішувалися з використанням спеціалізованих інтегральних схем тепер можуть бути реалізовані на паралельних процесорах, що дозволяє домогтися нового рівня функціональності і скоротити витрати на їх розробку. У 2007 році компанія NVIDIA заявила про створення нової технології CUDA (англ. Compute Unified Device Architecture), покликаної підвищити продуктивність звичайних персональних комп'ютерів за рахунок обчислювальних потужностей відеокарт цієї компанії. Початок масового використання паралельних обчислень поклало початок створення і розвиток перших технологій неграфічних розрахунків загального призначення GPGPU (General-Purpose computation on GPUs). Сучасні відеочіпи містять сотні математичних виконавчих блоків, і ця потужність може використовуватися для значного прискорення роботи обчислювально інтенсивних додатків. Нинішні покоління GPU (Graphics Processing Unit) мають досить гнучку архітектуру, що разом з високорівневими мовами програмування і програмно-апаратними архітектурами розкривають ці можливості і роблять їх значно більше доступними. На створення GPGPU розробників спонукала поява досить швидких і гнучких шейдерних програм, які здатні виконувати сучасні відеочіпи. Розробники задумали зробити так, щоб GPU розраховували не тільки зображення в 3D додатках, але і застосовувалися в інших паралельних розрахунках. В GPGPU для цього використовувалися Графічні API: OpenGL і Direct3D, коли дані до відеочіпа передавалися у вигляді текстур, а розрахункові програми завантажувалися у вигляді шейдерів. недостатками такого методу є порівняно висока складність програмування, низька швидкість обміну даними між CPU і GPU та інші подібні обмеження.

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

На сьогоднішній день продажу CUDA процесорів становлять близько 128 000 000. Область наукових досліджень з великим ентузіазмом зустріла технологію CUDA. Для прикладу, зараз CUDA прискорює AMBER, програму для моделювання молекулярної динаміки, використовувану більше 60000 дослідниками в академічному середовищі і фармацевтичними компаніями по всьому світу для скорочення термінів створення лікарських препаратів. На фінансовому ринку компанії Numerix і CompatibL анонсували підтримку CUDA в новому додатку аналізу ризику контрагентів і досягли прискорення роботи в 18 разів. Numerix використовується майже 400 фінансовими інститутами. Показником зростання застосування CUDA є також зростання використання графічних процесорів Tesla в GPU обчисленнях. На даний момент більш 700 GPU кластерів встановлені по всьому світу в компаніях в енергетичному секторі, а також у секторі банківських послуг.

У системах Microsoft Windows 7 і Apple Snow Leopard, обчислення на GPU зайняли свої позиції в секторі масових рішень. У цих операційних системах GPU постає не тільки графічним процесором, але також і універсальним процесором для паралельних обчислень, які працюють з будь-яким додатком.

Розділ 1. загальна теорія та поняття

1.1 Загальна термінологія CUDA

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

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

Ядро - це функція коду, який запускається на GPU. Це функція main для GPU коду. У системі може бути декілька ядер. Вони можуть бути одночасно запущені на одному або декількох пристроях і виконуватися послідовно, Але одночасно виконуватися на одному пристрої вони не можуть.

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

Сітка - це безліч потоків для одного ядра.

Блок - це набір потоків виконуваних на одному мультипроцесорі - це якийсь процесор, який є в складі GPU.

WARP - це набір потоків, фізично виконуваних паралельно. Тобто ця величина залежить безпосередньо від можливостей апаратних засобів. Кількість потоків може бути і 10 000, але фізично виконатися одночасно вони не можуть, а можуть лише 500. Так от ці 500 потоків і є WARP.

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

Потоки в блоці можуть взаємодіяти. Потоки з різних блоків взаємодіяти не можуть.

1.2 Структура NVIDIA CUDA

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

Підтримуються такі типи перетворень: complex-complex (C2C), real-complex (R2C) і complex-real (C2R).

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

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

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

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

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

1.4 Різниця між GPU і CPU

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

1.5 Програмна модель та модель пам'яті CUDA

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

Модель програмування в CUDA передбачає групування потоків. Потоки об'єднуються в блоки потоків (thread block) - одномірні або двомірні

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

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

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

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

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

· Глобальна пам'ять (global memory);

· Регістрова пам'ять (register);

· Локальна пам'ять (local memory);

· Поділювана пам'ять (shared memory);

· Константна пам'ять (constant memory);

· Текстурна пам'ять (texture memory);

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

Глобальна пам'ять (global memory)

Глобальна пам'ять (global memory) - це найповільніша із усіх доступних типів пам'яті на GPU. Вона служить для зберігання великих об'ємів даних. Саме в ній програміст може зберігати результати своїх обрахунків. Як при запису та і при читанні допускається довільна адресація, але затримки при її використанні можуть досягати сотні тактів. Через глобальну пам'ять проходить копіювання даних з host`a (CPU) на device (GPU). Для такого переміщення потрібно використовувати функцію cudaMemcpy. Глобальні змінні виділяються за допомогою специфікатора __global__, а також динамічно функціями типу cudMalloc. Для задач, які потребують швидкого виконання, потрібно зменшити кількість операцій з цією пам'яттю.

Регістрова пам'ять (register)

Регістрова пам'ять (register) - це найшвидша у усіх видів пам'яті. У ній компілятор зберігає локальні змінні. Для визначення кількість регістрів, потрібно використати функцію cudaGetDeviceProperties. Регістри ділять між всіма блоками потоку. Розрахувати кількість доступних регістрів для однієї нитки (thread) потрібно поділити загальне число регістрів на кількість ниток в блоці (block) і кількості блоків в сітці (grid). В середньому на один мультипроцесор доступно 8192 32-розрідних регістри.

Локальна пам'ять (local memory)

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

Поділювана пам'ять (shared memory)

Поділювана пам'ять (shared memory) - відноситься до швидкої пам'яті. Вона не кишується та її розмір становить 16 Кбайт на один мультипроцесор. ЇЇ рекомендовано використовувати для мінімізації запитів до глобальної пам'яті, а також зберігання локальних змінних функцій. Її використовують для синхронізації даних між потоками. Адресація поділюваної пам'яті між нитками (thread) однакова в межах одного блоку, що може бути використано для обміну даними між потоками в блоці. Для її використання потрібно скористатись специфікатором __shared__.

Константна пам'ять (constant memory)

Константна пам'ять (constant memory) - кешуємий, досить швидкий вид пам'яті. Його кеш не великий - всього 8 Кбайт на мультипроцесор, тому розміщати в ній можна тільки не великі обсяги часто використовуваних даних. Особливість константної пам'яті у тому, що до неї можна записувати данні з host`a, але при цьому у межах GPU можливе тільки зчитування з неї. Для розміщення даних користуються специфікатором __constant__. Якщо потрібно використовувати масив з константної пам'яті, то попередньо потрібно вказати його розмір, тому що динамічне виділення як у глобальній пам'яті відсутнє. Для запису даних з host`a використовують функцію cudaMemcpyToSymbol, а для копіювання з device на host - cudaMemcpyFromSymbol.

Текстурна пам'ять (texture memory)

Текстурна пам'ять (texture memory) - призначена для для роботи з текстурами. Вона має специфічні особливості в адресації, читання і запису даних. Текстурна пам'ять - це механізм кешування частин глобальної пам'яті. Текстура, як правило, представляє собою двохвимірний масив, тому ця пам'ять була оптимізована під швидку вибірку двохвимірних даних з фіксованим розміром. Після виділення глобальної пам'яті за допомогою спеціальних функцій cudaBindTexture вказується, що дана частина пам'яті повинна розглядатись як текстура. Вона доступна тільки на читання, для чого використовують спеціальні функції. Данні в текстурну пам'ять можна розміщати тільки фіксованого розміру (1, 2, 4, 8 і 16 байт)

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

Для ефективного використання можливостей CUDA потрібно забути про звичайні методи написання програм для CPU, і використовувати ті алгоритми, який добре розпаралелюється на тисячі потоків.

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

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

1.7 Проблеми CUDA

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

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

Про проблему сумісності можна і не говорити. Зараз багато навпаки погодилися б на втрату продуктивності, щоб тільки якусь то стару програму перенести без проблем з одної платформи на іншу. Тому, найбільш важлива проблема для ринку персональних ПК - це відсутність сумісних з CUDA технологій подібного рівня. Для комерційних розробників важлива підтримка всіх систем. Якби CUDA була реалізована, нехай навіть з меншою продуктивністю, в рішеннях AMD, то зараз було б набагато більше інвестицій в CUDA-ПО.

Але, AMD оголосила про підтримку остаточного варіанту архітектури OpenCL. Можна сказати, що це інша обгортка CUDA-архітектури, по суті, абсолютно те ж саме. Основні відмінності в синтаксисі. Так що, ця проблема сумісності частково вирішена. Ще одна сфера застосування - це професіональні рішення для підприємств і організацій, коли поставляється техніка і спеціальне ПЗ під конкретну задачу. Найчастіше такі рішення мають місце для конкретних замовників, вони не мають масового тиражу. Але на західному ринку співвідношення зарплат, тим більше, кваліфікованих програмістів і цін апаратного забезпечення, інше. Там частіше економічно вигідніше, при малих обсягах поставок, просто купити процесори тієї ж продуктивності, але за більш високу ціну, ніж залучати додаткових розробників, інженерів з якості ПО, системних архітекторів та інших співробітників в команду розробки. В цьому сенсі, нашими підприємствами, які ну дуже вже хочуть заощадити на розробці, технологія може бути затребувана набагато швидше. Інша проблема - це ще не повна сумісництвом

технології з деякими операційними системами. Наприклад, проблема, що деякі операційні системи Windows намагаються діяти як OS реального часу і норовлять перервати я CUDA-програми, якщо воно триває більше 5 секунд. Таким чином, OS Microsoft не може «достукатися» до GPU і намагається щось зробити. І розробникам доводиться обходити цю проблему з використанням специфічного системного пропрограмування для кожної операційної системи Microsoft.

І останнє - недостатня продуктивність обчислень з речовими числами, типу double з подвійною точністю. Звичайно, багатьом фізичним задачам не потрібно подвійна точність розрахунків. Але, double - це стандарт в індустрії, всі до нього звикли. Зрештою, рідко система збирається заради одного алгоритму. А якщо він зміниться, якщо потрібно вважати щось інше. Така вузька спеціалізація пристрою перешкоджає широкому застосування. Але відсутність повної підтримки double не являєся базової рисою архітектури CUDA. Можливо, так склалося тому, що в графіку тип double не дуже затребуваний. Однак, у новій версії CUDA архітектури - «Fermi» - уже цілеспрямовано призначеної для «неграфічних» застосований, double підтримується нарівні з float, як і в звичайних CPU. І є таким же базовим типом змінних.

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

Розділ 2. Модель програмування в СUDA

2.1 Основні поняття

Запропонована компанією Nvidia технолоrія CUDA (Cornpute Unified Device Architecture) помітно полегшує написання GРGРU додатків. Вона не використовує графічні АРI і вільна від обмежень, властивих цим API.

Дана технолоrія призначена для розробки додатків для масивно-паралельних обчислювальних пристроїв. На даний момент підтримуваними пристроями є всі GPU компанії Nvidia, починаючи з серії GeForce8, а також спеціалізовані для вирішення розрахункових завдань GPU ceмейства Tesla

Основними перевагами технолоrіі CUDA є її простота всі проrрамми пишуться на «розширеній» мові С++, наявність гарної документації, набір готових інструментів, що включають профайлер, набір готових бібліотек, кроссплатформенность (підтримуються Мiсrоsоft Windows, Linux і Мас OS Х).

СUDА є повністю безкоштовною, SDK, документацію і приклади можна викачати з сайту developer.nvidiacorn . На даний момент останньою версією є CUDA 2.3.

CUDA будується на концепції, що GPU(що називається пристроєм, device) виступає в ролі масивно-параллельноrо співпроцесора до CPU(званому host). Проrрамма на CUDA задіює як CPU, так і GPU. При цьому звичайний(послідовний, тобто непаралельний) код виконується на CPU, а для масивно-паралельних обчислень відповідний код виконується на GPU як набір ниток(потоків, threads), що одночасно виконуються.

Таким чином, GPU розглядається як спеціалізований обчислювальний пристрій, який :

- є співпроцесором до CPU;

- має власну пам'ять;

- має можливість паралельноrо виконання великої кількості окремих «ниток».

При цьому дуже важливо розуміти, що між нитками на CPU і нитками на GPU є принципові відмінності:

- нитки на GPU мають украй невелику вартість створення, управління і знищення (контекст нитки мінімальний, усі регістри розподілені заздалегідь);

- для ефективної заrрузки GPU необхідно використати багато тисяч окремих ниток, тоді як для CPU зазвичай достатні 10-20 ниток.

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

Специфікатор __global__ використовується для позначення тoro, що це ядро, тобто функція, яка працює на GPU і яка може бути викликана (точніше, запущена відразу на великій кількості ниток) тільки з CPU.

Ілюстрація використання CUDA :

- виділяємо пам'ять на GPU;

- копіюємо дані з пам'яті CPU у виділену пам'ять GPU;

- здійснюємо запуск ядра чи послідовно запускаємо декілька ядер;

- копіюємо результати обчислень назад в пам'ять CPU;

- звільняємо виділену пам'ять GPU.

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



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

двомірний масив блоків (block). Кожен блок - це одновимірний, двомірний або трьохвимірний мсив ниток (tread). При цьому всі блоки які складають сітку мають однакову розмірність і розмір.

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

Оскільки одно і те ж ядро виконується одночасно дуже великим числом ниток, то для того, щоб ядро могло однозначно визначити номер нитки, використовуються вкладені змінні treadIdx I blockIdx. Кожна із цих змінних являється трьохвимірним цілочисельним вектором. Зверніть увагу що вони доступні тільки для функцій, які виконуються на GPU, - для функцій на СPU, вони не мають смислу.

Також ядро може получити розміри сітки і блока через встроєні змінні gridDim I blockDim.

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

При цьому кожна під задача спільно розв'язується всіма нитками свого блоку. Розбиття ниток на warp'и проходять окремо для кожного блоку; таким чином, всі нитки одного warp'а завжди належать одному блоку. При цьому нитки можуть взаємодіяти між собою тільки в межах блоку. Нитки різних блоків взаємодіяти між собою не можуть.

Існують лише два механізми, за допомогою яких нитки всередині блока можуть взаємодіяти між собою:

- розподілювальна (shared) пам'ять;

- бар'єрна синхронізація.

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

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

2.2 Розширення мови С++

Проrрамми для CUDA (відповідні файли зазвичай мають розширення

*.сu) пишуться на «розширеному» С і компілюються за допомогою команди nvcc.exe

Введені в CUDA розширення мови С складаються з:

- специфікаторів функцій, що показують, де буде виконуватися функція і звідки вона може бути викликана;

- специфікаторів змінних, які задають тип пам'яті, що використовується для даних змінних;

- директиви, що служить для запуску ядра, що задає як дані, так і ієрархію ниток;

- вбудованих змінних, що містять інформацію про поточну нитки;

- ruпtime, що включає в себе додаткові типи даних.

2.2.1 Специфікатори функцій й змінні

У CUDA використовується наступні специфікатори функцій наведені в таблиці:

Специфікатор	Функція виконується на	Функція може викликатися із
__device__	devise (GPU)	devise (GPU)
__global__	devise (GPU)	host (CPU)
__ host__	host (CPU)	host (CPU)

Специфікатори __host__ і __device__ можуть бути використані разом (це означає, що відповідна функція може виконуватися як на GPU, так і на CPU - відповідний код для обох платформ буде автоматично згенерований компілятором). Специфікатори __global__ і __host__ не можуть бути використані разом. Специфікатор __global__ означає ядро, і відповідна функція повинна повертати значення типу void.

На функції, що виконуються на GPU(__device__ і __global__), накладаються наступні обмеження:

- не можна брати їх адресу(за виключенням - global - функцій);

- не підтримується рекурсія;

- не підтримуються stаtiс- змінні усередині функції;

- не підтримується змінне число вхідних aprументів.

Для завдання розміщення в пам'яті GPU змінних використовуються наступні специфікатори __device__, __constant__ і __shared__. На їх використання також накладається ряд обмежень:

- ці специфікатори не можуть бути застосовані до полів структури(struct чи union);

- запис в змінні типу - constant - може здійснюватися тільки CPU за допомогою спеціальних функцій;

- __shared__ змінні не можуть ініціалізуватися при оголошенні.

2.2.2 Додані змінні

В мову додані наступні спеціальні змінні:

- gridDim - розмір сітки (має тип dimз);

- blockDim - розмір блоку (має тип dim3);

- blockldx - індекс поточного блоку в сітці (має тип uint3);

- threadldx - індекс поточної нитки у блоці (має тип uint3);

- warpSize - розмір warp'a (має тип int).

2.2.3 Директива виклику ядра

Для запуску ядра на GPU використовується наступна конструкція:

// kernelName <<<Dg, Db, Ns, S>>> ( args );

Тут kemelName це ім'я (адреса) тієї, що відповідає __global__ функції. Через Dg позначена змінна (чи значення) типу dim3, яка задає розмірність і

розмір сітки (у блоках). Змінна (чи значення) Db - типу dim3, задає розмірність і розмір блоку (у нитках).

Необов'язкова змінна (чи значення) Ns типу size_t задає додатковий об'єм пам'яті у байтах, що розділяється, яка має бути динамічно виділена кожному блоку (до вже статично виділеної пам'яті, що розділяється) якщо не задано, то використовується значення О.

Змінна (чи значення) S типу cudaStream_t задає потік (CUDA stream)

у якому повинен статися виклик, за замовчуванням використовується потік О. Через args позначені арrументы виклику функції kernelName (їх може бути декілька ).

2.2.4 Добавлені функції

CUDA підтримує всі математичні функції із стандартної бібліотеки мови С++. Однак при цьому слід мати на увазі, що більшість стандартних математичних функцій використовують числа з подвійною точністю (double). Однак оскільки для сучасних GPU операції з double-числами виконуються повільніше, ніж операції з float-числами, то використовують там де це можливо float-аналоги стандартних функцій. Так, float-аналогом функції sin являється функція sinf.

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

Висновки

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

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

Основним завданням курсової роботи було порівняння швидкодії роботи GPU та CPU у підрахунку CRC32-коду.

Для вирішення завдання була написана програма на мові С з використанням технології СUDA.

Література

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

Эйнджел Э. Интерактивная компьютерная графика. Вводний курс на базе OpenGL / Э. Эйнджел. 2001. - 590 с.

Тихомиров Ю. Программирование трехмерной графики. Серия ''Мастер''. - СПб.: BHV-Санкт-Петербург. 2000. - 256 с.

Майер Р.В. Компьютерное моделирование физических явлений / Р.В. Майер. - Глазов: ГГПИ, 2009. - 111 с.

Мейсон Ву. OpenGL. Официальное руководство программиста / Ву Мейсон, Нейдер Д., Девис Т., Шрайнер Д. 2002. - 592 с.

Роджерс Д. Алгоритмические основы машинной графики / Д. Роджерс. М.: Мир. 1989. - 512 с.

Размещено на Allbest