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

Київський національний університет імені Тараса Шевченка

Радіофізичний факультет

Аналітичний звіт на тему

«Програмно-апаратна платформа CUDA, архітектурний ряд Fermi, програмування за допомогою засобів програмної платформи CUDA SDK. Технологія GPGPU»

Київ 2012

Зміст

• Список використаних скорочень
• Вступ
• 1. Технологія GPGPU
• 1.1 Еволюція GPU та поява GPGPU
• 1.2 OpenCL та DirectX Compute
• 1.3 Nvidia CUDA
• 2. Програмно-апаратна платформа CUDA
• 2.1 Програмна архітектура CUDA
• 2.2 Архітектура пам'яті в CUDA
• 3. Програмування за допомогою CUDA SDK
• 3.1 Підготовка необхідного інструментарію
• 3.2 Приклад CUDA-застосування
• 3.3 Особливості програмування з CUDA SDK
• 4. Архітектурний ряд Fermi
• 4.1 Інновації в архітектурі Fermi
• 4.2 Огляд архітектури Fermi
• Висновки
• Перелік джерел інформації

Список використаних скорочень

ALU - Arithmetic Logic Unit

API - Application Programming Interface

CPU - Central Processing Unit

CUDA - Compute Unified Device Architecture

ECC - Error Correction Control

GPU - Graphics Processing Unit

GPGPU - General-Purpose computing on GPU

HPC - High Performance Cluster

SFU - Special Function Unit

SIMD - Single Instruction Multiple Data

SM - Streaming Multiprocessor

Вступ

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

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

На створення GPGPU розробників спонукало появу досить швидких і гнучких шейдерних програм, які здатні виконувати сучасні відеочіпи. Розробники задумали робити так, щоб 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 4.1. Версія 3.0 була останньою, в котрій ще підтримувалась можливість емуляції, тобто можливість виконання CUDA-програм і системах, в яких відсутні пристрої, що підтримують цю технологію.

1. Технологія GPGPU

GPGPU (General-Purpose computing on Graphics Processing Units) - техніка використання графічного процесора відеокарти для загальних обчислень, які зазвичай виконує центральний процесор. Це стало можливим завдяки додаванню програмованих шейдерних блоків і більш високої арифметичної точності растрових конвеєрів, що дозволяє розробникам ПО застосовувати потокові процесори для неграфічних даних.

1.1 Еволюція GPU та поява GPGPU

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

Зараз сучасні GPU є потужними масивно-паралельними обчислювальними пристроями з дуже великою швидкодією (більше одного терафлопа) і великим об'ємом власної пам'яті.

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

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

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

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

Після своєї появи прискорювач 3D графіки швидко еволюціонували при цьому окрім збільшення швидкодії також зростала і їх функціональність. Так, графічні прискорювачі наступного покоління (наприклад, Riva TNT) уже могли самостійно обробляти вершини, розвантажуючи при цьому центральний процесор і одночасно накладати кілька текстур.

Наступним кроком було збільшення гнучкості при обробці окремих фрагментів (пікселів) для реалізації цілого ряду ефектів, наприклад, піксельного освітлення. На цьому етапі розвитку створити повноцінний обробник фрагментів було неможливо, але доволі велику частину функціоналу можна було реалізувати за допомогою так-званих register combiner-ів. Які з'явилися в GeForce 256. Це були блоки, здатні реалізувати доволі прості операції (наприклад, обчислення скалярного добутку). При цьому ці блоки можна було доволі швидко і легко настроювати, з'єднуючи між собою відповідні входи та виходи. Виконавши необхідну конфігурацію, можна було реалізувати основні операції попіксельного освітлення.

Наступним кроком була поява вершинних програм (GeForce 2) - замість фіксованого набору кроків з обробки вершин можна було задати програму, написану на спеціальному асемблері. Дана програма виконувалася паралельно, а вся робота виконувалася над дійсними числами типу float (32 біти).

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

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

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

Сам модуль, який виконує перетворення вхідних потоків у вихідні прийнято називати ядром (kernel). Окремі ядра можуть сполучатися між собою, що призводить до утворення досить складних схем обробки потоків вхідних даних.

Рис.1.1.1. Приклад багатоядерної обробки даних на GPU

Додавання підтримки текстур із значеннями в форматах з плаваючою крапкою (16- та 32-бітові значення float замість використовуваних раніше 8-бітних без знакових цілих чисел) дозволило застосовувати фрагментні процесор для обробки великих масивів даних. Для виконання такого типу задач дані передавалися на GPU у вигляді текстур зі значеннями компонент типу float, а результат обробки зберігався в текстурах того ж типу.

Фактично було отримано потужний паралельний процесор, і на вхід якого можна передати великий масив даних і програму для їх обробки. А на виході отримати масив результатів. Поява низькорівневих мов для написання програм для GPU, таких як Cg, GLSL, HLSL, помітно полегшило процес створення таких програм.

В результаті з'явилася нова течія, яку було названо GPGPU - використання графічних процесорів для розв'язку нетипових (неграфічних) задач. За рахунок високого рівня паралелізму досить швидко вдалося одержати дуже високу швидкодію - пришвидшення у порівнянні із CPU для деяких задач сягало 10 разів.

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

В основному розробники використовували один із поширених графічних API (OpenGl або Direct3D) для доступу до ресурсів графічного процесора. Використовуючи такий API, готувалися текстури, що містили вхідні дані, а потім через операцію рендерінгу (зазвичай прямокутника) на графічному процесорі запускалася програма для оброби цих даних. Результат отримувався також у вигляді текстур, які потім зчитувалися у пам'ять центрального процесора.

Фактично програму доводилося писати одночасно двома мовами - традиційною мовою програмування, наприклад С++, і мовою написання шейдерів. Частина програми. Написана традиційною мовою, відповідала за підготовку вхідних даних та їх передачу, а також за запуск на GPU програм, написаних шейдерними мовами.

Однак GPGPU в класичному варіанті містить деякі недоліки, що ускладнюють його поширення. Ці обмеження пов'язані з тим, що використання можливостей GPU проходить через API, орієнтований власне на роботу із графікою (OpenGl, Direct3D). Як результат це впливає на реалізацію обчислювальних задач, де такі обмеження є явно надлишковими.

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

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

Рис. 1.1.2. Операції типу scatter (розподілення)

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

Ще одним «спадковим» недоліком є те, що розробка проходить одночасно двома мовами: для CPU та для GPU відповідно.

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

В якості таких засобів виступають NVidia CUDA, OpenCL[11] та DX11 Compute Shaders[12].

1.2 OpenCL та DirectX Compute

OpenCL - відкритий стандарт для паралельного програмування гетерогенних систем.

OpenCL є першим відкритим безплатним стандартом для кросплатформенного паралельного програмування сучасних процесорів персональних комп'ютерах, серверах і портативних та вбудованих пристроях. OpenCL (Open Computing Language) значно збільшує швидкість і гнучкість для широкого спектру застосувань (від ігор і розваг до наукових і медичних програм).

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

OpenCL створюється Khronos Group за участю багатьох передових компаній і установ, утому числі 3Dlabs, Activision Blizzard, AMD, Apple, ARM, Broadcom, Codeplay, Electronic Arts, Ericsson, Freescale, Fujitsu, GE, Graphic Remedy, HI, IBM, Intel, Imagination Technologies, Лос-Аламоської національної лабораторії, Motorola, Movidius, Nokia, NVIDIA, Petapath, QNX, Qualcomm, Rapid Mind, Samsung, S3, ST Microelectronics, Такумі, Texas Instruments, Toshibaі Vivante.

У фреймворк OpenCL входять мова програмування, яка базується на стандарті C99, та інтерфейс програмування (API). OpenCL забезпечує паралельність на рівні інструкцій та на рівні даних і є реалізацією техніки GPGPU. OpenCL--повністю відкритий стандарт, його використання доступне на базі вільних ліцензій. Мета OpenCL полягає в тому, щоб доповнити OpenGl і OpenAL, які є відкритими галузевими стандартами для тривимірної комп'ютерної графіки і звуку, користуючись можливостями GPU.

DirectX Compute - технологія, створена компанією Microsoft на базі бібліотеки Direct3D. Для розв'язування задач використовуються шейдери, але є можливість використання специфічних для DirectX API.

1.3 Nvidia CUDA

Запропонована компанією NVidia технологія CUDA помітно полегшує створення GPGPU-застосувань. Вона не використовує графічних АРІ і вільна від обмежень, притаманних такого роду АРІ[4].

Дана технологія призначена для розробки застосувань для масивно-паралельних обчислювальних приладів. На сьогоднішній день підтримуються всі GPU компанії NVidia, починаючи із GeForce 8, а також спеціалізовані для розв'язування розрахункових задач GPU сімейства Tesla.

Основними перевагами CUDA є її простота - всі програми пишуться на розширеній мові С (С++), наявність хорошої документації, набір готових інструментів, який включає профайлер, набір готових бібліотек, кросплатформенність (підтримуються Microsoft Windows, Linux, Mac OS X).

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

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

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

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

- може паралельно виконувати величезну кількість паралельних потоків.

При цьому слід чітко розуміти принципову різницю між потоками на CPU та на GPU:

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

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

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

Зазвичай типова програма з використанням CUDA має наступний вигляд:

1. Виділення пам'яті на GPU

2. Копіювання даних між CPU та GPU

3. Запуск ядра (технологія Fermi дозволяє запуск кількох ядер одночасно)

4. Копіювання результатів даних між GPU та CPU

5. Звільнення виділеної пам'яті на GPU

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

Важливим моментом є те, що хоча даний підхід і є подібним на роботу із SIMD-моделлю, але є принципові відмінності (компанія NVidia використовує термін SIMT - Single Instruction Multiple Thread). Потоки розбиваються на групи по 32 штуки, які називаються варпами (warp). Тільки потоки в межах одного варпа виконуються фізично одночасно. Потоки з різних варпів можуть перебувати на різни стадіях виконання програми. При цьому керування варпами прозоро для програміста виконується самим GPU.

Всі запущені на виконання потоки організовані в деяку ієрархію.

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

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

Оскільки одне ядро виконується в оди момент часу великою кількістю потоків, для того, щоб можна було визначити номер потоку використовуються вбудовані змінні threadIdx і blockIdx. Кожна з цих змінних є цілочисельним тривимірним вектором. Ці змінні мають сенс лише для функцій, що виконуються на GPU, для CPU вони не мають жодного сенсу. Також ядро може визначити розміри сітки та блока за допомогою змінних gridDim та blockDim.

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

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

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

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

- спільна пам'ять (shared memory);

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

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

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

2 Програмно-апаратна платформа CUDA

CUDA (англ. Compute Unified Device Architecture) -- технологія GPGPU (англ. General-purpose computing on Graphics Processing Units), що дозволяє програмістам реалізовувати мовою програмування С (а в останніх версіях CUDA і мовою C++) алгоритми, що виконуватимуться на графічних процесорах Geforce восьмого покоління і старше (Geforce 8 Series, Geforce 9 Series, Geforce 200 Series), NvidiaQuadro і Tesla компанії Nvidia. Технологія CUDA розроблена компанією Nvidia.

2.1 Програмна архітектура CUDA

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

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

· володіє власною пам'яттю;

· надає можливість параллельноrо виконання величезної кількості окремих потоків

CUDA використовує велику кількість окремих потоків для обчислень, часто кожному обчислювальному елементові відповідає одинпоток. Всі нитки групуються в ієрархію - сітка / блок / потік (див.рис.2.1.1). Верхній рівень - сітка - відповідає ядру і об'єднує всі потоки, що виконують дане ядро. Сітка (grid) є одновимірним або двомірним масивом блоків (block). Кожен блок zвляє себою одно- двох- або тривимірний масив потоків (threads).

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

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

При цьому потоки всередині блоку можуть взаємодіяти між собою через:

· загальну пам'ять (що розділяється пам'яті)

· функцію синхронізації всіх ниток блоку (__synchronize)

З апаратної точки зору всі потоки розбиваються на так звані warp-и - блоки підряд ідучихпотоків, які одночасно (фізично) виконуються і можуть взаємодіяти один з одним. Кожен блок пооків розбивається на декілька warp'ов, розмір warp'а для усіх існуючих зараз GPU дорівнює 32.

Рис.2.1.1. Ієрархія потоків в CUDA



Рис.2.1.2. Цикл виклику ядра

Типовий цикл CUDAпрограми протікає наступним чином:

1. Виділення пам'яті на GPU

2. Копіювання даних із пам'яті CPU на GPU.

3. Конфігурація та запуск ядра

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

5. Копіювання даних із GPUв пам'ять CPU.

6. Звільнення пам'яті GPU.

Рис.2.1.3. Послідовність виконання ядра

Проrрами для CUDA (відповідні файли зазвичай мають розширення.cu) пишуться на «розширеній» мові С (а тепер і C++) і компілюються компіляторомnvcc.

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

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

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

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

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

· runtime, що включає додаткові типи даних.

В CUDAвикористовуються наступні специфікатори функцій (табл.. 2.1.1):

Табл.2.1.1. Специфікатори функцій в CUDA

Специфікатор	Фукція виконується на	Функція може бути викликана із
device	device (GPU)	device (GPU)
global	device (GPU)	host (CPU)
host	host (CPU)	host (CPU)

CUDA надає ряд функцій, котрі можуть бути використані лише на CPU (так званий CUDA Host API).Ці функції відповідають за:

· керування GPU;

· роботу з контекстом;

· роботу з пам'яттю;

· роботу з модулями;

· керування виконанням коду;

· роботу з текстурами;

· взаємодію з OpenGL і Direct3D;

CUDAAPIдля CPUвиступає в двох формах:

· низькорівневий CUDA Driver API

· високорівневий CUDAruntimeAPI, реалізований через CUDADriverAPI

В одній програмі можна використовувати лише один із них.

Рис.2.1.4. Програмний стек CUDA

2.2 Архітектура пам'яті в CUDA

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

На рис.1.3.1 та 1.3.2 показані структура пам'яті GPUта схема доступу різних елементів ядра до різних видів пам'яті відповідно.



Рис.1.3.1. Структура пам'яті в CUDA

Рис.1.3.2. структура роботи ядра з пам'яттю GPU

В табл1.3.1 наведені доступні види пам'яті в CUDAі їх основні характеристики.

Табл.1.3.1. Види пам'яті в CUDA

Тип пам'яті	Доступ	Рівень виділення	Швидкість роботи
registers	R/W	per-thread	висока (on chip)
local	R/W	per-thread	низька (DRAM)
shared	R/W	per-block	висока (on-chip)
global	R/W	per-grid	низька(DRAM)
constant	R/O	per-grid	висока(on chip L1 cache)
texture	R/O	per-grid	висока(on chip L1 cache)

Чи не основними шляхами пришвидшення роботи програм CUDAє оптимізація звернень до глобальної то роздільної пам'яті.

Стосовно першої, слід намагатись здійснювати звернення потоків напівварпа до глобальної пам'яті таким чином, щоб здійснювалися зчитування чи запис 16-и послідовних 4-хбайтових слів відповідно до номерів потоків (рис.1.3.3). Це носить назву coalescingі дозволяє прискорити роботу з globalпам'яттю до 16 разів.

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

Рис.1.3.3. Coalescing присутній	Рис.1.3.4. Coalescing не спостерігається
Рис.1.3.5. Безконфліктний доступ до пам'яті виду shared	Рис.1.3.6. Конфлікт 8-го порядку

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

3 Програмування за допомогою CUDA SDK

3.1 Підготовка необхідного інструментарію

До початку створення CUDA-застосувань необхідно підготувати середовищу, в якому можна створювати програми мовами С/С++ (із застосуванням CUDA SDK) та виконувати їх.

Графічний процесор із підтримкою CUDA.

Технологія CUDA підтримується кожним відеоадаптером від компанії Nvidia, починаючи із GeForce 8800 GTX, випущеного в 2006 році. Повний перелік моделей можна переглянути на офіційному сайті компанії: www.nvidia.com/cuda .

Емуляція CUDA-процесорів для відлагодженя роботи CUDA-програм підтримується у всіх версіях CUDA SDK до CUDA 3.0 включно. Для від лагодження та виконання програм, створених з допомогою усіх наступних версій обов'язкова наявність GPU, створеного на основі архітектури CUDA.

Драйвер пристроїв Nvidia.

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

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

CUDA Development Toolkit.

А для можливості і розробляти CUDA-застосування, необхідна установка CUDA Toolkit, що містить спеціальний C-компілятор тих одиниць трансляції застосування, що містять серцевий код, призначений для виконання на GPU, тобто код, написаний на CUDA.

CUDA Toolkit можна завантажити на офіційному сайті компанії за адресою www.developer.nvidia.com/object/gpucomputing.html .

Стандартний C/C++ - компілятор.

Необхідний для компіляції частин програмного коду, призначеного для виконання на центральному процесорі. Зв'язування (англ. linking) об'єктних файлів, отриманих в результаті роботи обох компіляторів, відбувається наявним стандартним «компонувальником».

3.2 Приклад CUDA-застосування

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

//-----------------------------------------------------------------------

// main.cpp

//-----------------------------------------------------------------------

#include "resource.h"

#define N 10

int main( void ) {

int a[N], b[N], c[N];

int *dev_a, *dev_b, *dev_c;

// allocate the memory on the GPU

HANDLE_ERROR( cudaMalloc( (void**)&dev_a, N * sizeof(int) ) );

HANDLE_ERROR( cudaMalloc( (void**)&dev_b, N * sizeof(int) ) );

HANDLE_ERROR( cudaMalloc( (void**)&dev_c, N * sizeof(int) ) );

// fill the arrays 'a' and 'b' on the CPU

for (int i=0; i<N; i++) {

a[i] = -i;

b[i] = i * i;

}

// copy the arrays 'a' and 'b' to the GPU

HANDLE_ERROR( cudaMemcpy( dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice ));

HANDLE_ERROR( cudaMemcpy( dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice ) );

// call the kernel

add<<<N,1>>>( dev_a, dev_b, dev_c );

// copy the array 'c' back from the GPU to the CPU

HANDLE_ERROR( cudaMemcpy( c, dev_c, N *sizeof(int),cudaMemcpyDeviceToHost ) );

// display the results

for (int i=0; i<N; i++) {

printf( "%d + %d = %d\n", a[i], b[i], c[i] );

}

// free the memory allocated on the GPU

cudaFree( dev_a );

cudaFree( dev_b );

cudaFree( dev_c );

return 0;

}

//-----------------------------------------------------------------------

// kernel.cu

//-----------------------------------------------------------------------

__global__ void add( int *a, int *b, int *c ) {

int tid = blockIdx.x; // handle the data at this index

if (tid < N)

c[tid] = a[tid] + b[tid];

}

//-----------------------------------------------------------------------

3.3 Особливості програмування з CUDA SDK

Далі перелічено деякі основні особливості стосовно створення програмних застосувань з допомогою CUDA SDK. Вони ж відносяться до тих, що виділяють CUDA споміж інших технологій GPGPU.

Оголошення функцій.

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

Запуск «ядра».

В CUDA також розширено синтаксис C-функцій шляхом додавання параметри конфігурації виклику ядра, оточених символами << та >>. Ці параметри визначають розмірність сітки (англ. grid) ядра, що викликається, та розмірність кожного з її блоків (англ. block). Деталі описані в CUDA Programming Guide, що входить до складу CUDA SDK.

Наперед визначені змінні.

Ядро CUDA-застосування отримує доступ до деякої множини наперед визначених змінних, що, зокрема, дозволяють кожному потокові визначати своє положення серед інших потоків та, таким чином, визначити з якими саме даними в пам'яті йому слід працювати. Річ іде, зокрема, про змінні threadIdx, gridDim та blockDim.

Runtime API.

В CUDA наявний набір API-функцій, що забезпечують функціонування CUDA-застосувань. Наприклад, функції cudaMalloc() та cudaMemcpy(), що виділяють пам'ять на CUDA-пристрої та виконують перенесення даних між пам'яттю комп'ютера та пам'яттю адаптера відповідно.

Повний перелік функцій CUDA Runtime API також можна знайти в CUDA Programming Guide.

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

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

Синхронізація потоків.

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

Робота з пам'яттю.

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

І, нарешті, слід відзначити можливість застосування декількох GPU для роботи одного CUDA-застосування.

4 Архітектурний ряд Fermi

Архітектури під кодовою назвою «Fermi» - наступне архітектурне покоління відеоадаптерів із підтримкою CUDA від компанії NVIDIA.

4.1 Інновації в архітектурі Fermi

платформа cuda fermi програмування

Архітектура Fermi є найбільш значним кроком уперед в архітектурі GPU з часів оригінальної G80. G80 було початковим баченням того, як повинен виглядати універсальний графічний та масивно-паралельний процесор. GT200 розвинув продуктивність і функціональність G80. В Nvidia кажуть, що для створення Фермі, підсумовано всі аспеккти, про які вдалося дізнатися від двох попередніх процесорів і всіх застосувань, які були для них написані, і використано абсолютно новий підхід до дизайну, щоб створити перший у світі обчислювальний GPU. Коли почалося закладання фундаменту для Фермі, було зібрано велику кількість відгуків користувачів щодо обчислювальних GPU після введення G80 і GT200, і зосереджено сили на таких основних напрямках для покращення:

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

· Підтримка ECC (error correction control - контроль помилок пам'яті) - дозволяє користувачам GPGPU безпечно встановлювати велику кількість графічних процесорів в установках обробки даних, а також захистити «чутливі до оброблюваних даних» програми від помилок роботи пам'яті.

· «Правильна» ієрархія кеш-пам'яті - деякі паралельні алгоритми не надавали можливості використання розділюваної (shared) пам'яті GPU, так що користувачі зверталися з проханнями про «правильну» ієрархію, для вирішення подібних проблем.

· Більше розділюваної пам'яті - чимало програмістів CUDA також прохали про більшу ніж 16 Кб/мультипроцесор кількість розділюваної пам'яті для прискорення власних застосування.

· Більше швидке перемикання контексту - також поступали запити про більш швидке перемикання між програмними застосуваннями.

· Прискорення атомарних операцій - у відповідь на необхідність прискорення атомарних операцій «читання-модифікація-запис» під час роботи паралельних алгоритмів.

Маючи це на увазі, команда Fermi розробила процесор, який значно підвищ обчислювальну потужність комп'ютера, і, шляхом введення архітектурних інновацій, також пропонує різко зростаючу «програмованість» та обчислювальну ефективність. Основними архітектурними моменти Fermi є:

· Третє покоління потокових мультипроцесорів (SM - streaming multiprocessor)

o 32 CUDA ядра на один SM, 4x-приріст у пор. з GT200

o 8x-приріст піку обчислень подвійної точності з плаваючою точкою у пор. з GT200

o Dual Warp Scheduler - одночасно отримує і відправляє інструкції від двох незалежних warp-ів

o 64 КБ RAM (random access memory) з налаштованим розподіленням розділюваної пам'яті і кешу 1-го рівня (L1 cache)

· Паралельнео виконання потоків ISA другого покоління

o Єдиний адресний простір з повною підтримкою C + +

o Оптимізовано для OpenCL і DirectCompute

o Повна підтримка 32-бітної та 64-бітної точності згідно стандарту IEEE 754-2008

o Повний 32-бітовий цілочисельний шлях із 64-розрядними розширеннями

o Інструкції доступу до пам'яті для підтримки переходу на 64-розрядну адресацію

o Покращена продуктивність через предикації (predication)

· Покращена підсистема пам'яті

o ієрархія NVIDIA Parallel DataCache з налаштовуваним «однорідним» кешом рівнів L1 і L2

o перші GPU з підтримкою ECC для пам'яті

o Значно покращена продуктивність атомарних операцій з пам'яттю

· двигун NVIDIA GigaThread

o в10 разів швидше перемикання контексту застосувань

o паралельне виконання завдань ядра

o виконання блоку при «виході з ладу» одного з потоків

o двигуни передачі пам'яті типу «Dual overlapped»

4.2 Огляд архітектури Fermi

Перша відеокарта архітектури Fermi, реалізована на базі 3.0 млрд. транзисторів, містила 512 ядер CUDA. Таке ядро виконує інструкції над цілими числами та числами з плаваючою точкою для одного потоку й відповідно тактовому генератору. 512 ядер організовані в 16 потокових мультипроцесорів, по 32 на кожен. Розглядуваний GPU має 6 64-бітних відсіків пам'яті, 384-бітний інтерфейс пам'яті й підтримує до 6 Гігабайт глобальної пам'яті серії GDDR5. Підключається до комп'ютера через інтерфейс PCI-Express. Блоки CUDA-ядрарозподіляються між планувальниками потоків поткових мультипроцесорів з допомогою глобального диспетчера під назвою GigaThread global scheduler.



Рис. 4.2.1. 16 мультипроцесорів розташовані навколо спільної кеш-пам'яті рівня L2

Потоковий мультипроцесор третього покоління.

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

512 високоефективних ядер CUDA.

Отож, кожен потоковий мультипроцесор містить 32 CUDA-процесори - у 4 рази більше, ніж в архітектурах-попередниках. Кожне таке ядро володіє повністю конвеєризованим арифметико-логічним пристроєм (ALU) та модулем операцій з плаваючою очкою. Архітектура Fermi реалізовує новий стандарт щодо операцій із плаваючою точкою - IEEE 754-2008, на відміну від IEEE 754-1985 у попередників.



Рис. 4.2.2. Потоковий мультипроцесор архітектури Fermi.

Слід додати, що арифметико-логічний пристрій моделі GT200 був обмежений 24-бітною точністю при виконанні операцій множення, в той час, як відповідні ALU в Fermi підтримують 32-бітну точність для усіх операцій, та 64-бітну для деяких спеціальних.

16 модулів завантаження/збереження.

Кожен потоковий мультипроцесор містить 16 блоків «завантаження-збереження» (англ.. load/store unit), що дозволяє обрахунок «адреси джерела» та адреси призначення для 16-и потоків за один такт. Збереження даних відбувається в кеш-пам'ять або глобальну.

4 модулі спеціального призначення (SFU - special function unit).

Функція таких модулів - виконувати трансцендентні операції, такі, як sin, cos обернені їм та взяття квадратного кореню. Кожен киконують одну інструкцію за такт; warp виконується за 8 або більше тактів. Конвеєр SFU відділений від «диспетчера» інструкцій, що надає можливість останньому відправляти інструкції іншим модулям під час роботи пристроїв спеціального призначення.

Рис. 4.2.3. Попередня оцінка ефективності показує, що Fermi працює в 4.2 рази швидше, ніж GT200 при виконанні операцій із подвійною точністю.

Розроблено для пришвидшення виконання операцій із подвійною точністю.

Арифметика операцій із подвійною точністю лежить в серці HPC-застосувань, таких як обчислення із галузей лінійної алгебри, числового моделювання чи квантової хімії. Архітектура Fermi розроблена таким чином, щоб забезпечити безпрецедентну продуктивність таких операцій: один SM може виконати до 16 інструкцій типу «множення-додавання» за один такт - значне покращення в порівнянні з архітектурою GT200.

Dual Warp планувальник.

Мультипроцесор групує потоки в групи по 32 фізично-паралельні потоки, котрі називаються warp-ами. Кожен мультипроцесор містить два групувальники й два модулі «розсилки» інструкцій, що дозволяє паралельне виконання одразу двох warp-ів. Планувальник вибирає два warp-и і передає по одній інструкції від кожного групі із 16-и ядер, 16-и модулів «load/store». Оскільки warp-и виконуються незалежно, відпадає необхідність у перевірці залежностей всередині потоку інструкцій. Використовуючи таку елегантну двоварпову модель, карти Fermi досягають практично пікової апаратної продуктивності.

Рис. 4.2.4. Dual Warp Scheduler

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

64 КБ конфігурованих розділюваної пам'яті та кешпам'яті рівня L1.

Розділювана пам'ять є, без сумніву, однією із ключових архітектурних інновацій, що значно покращує програмованість та продуктивність GPU-застосувань. Відеоадаптери G80 та GT200 мають по 16 КБ на один мультипроцесор. В архітектурі Fermi кожен потоковий мультипроцесор включає 64 КБ пам'яті чіпу, що може бути сконфігурована, як 48 КБ розділюваної пам'яті та 16 КБ кеш--пам'яті рівня L1, або ж навпаки.

У табл. 3.2.5 наведена порівняльна характеристика архітектурного ряду Fermi та відеокарт G80 й GT200.

Рис. 4.2.5. Порівняльна таблиця

Висновки

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

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

CUDA SDK є чи не найзручнішим для програміста серед усіх існуючих інструментарієм для створення GPGPU-застосувань, але від цього не менш продуктивним.

Що ж стосується застосувань CUDA в сторонніх середовищах та мовах програмування, відмінних від C/C++, то, доводилось на практиці спостерігати застосування CUDA в середовищі Mathematica 8 для операцій редукції масивів великого розміру й результати роботи з точки зору продуктивності були доволі вражаючими, в хорошому сенсі.

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

Перелік джерел інформації

1. http://www.nvidia.com

2. http://en.wikipedia.org/wiki/CUDA

3. cudacsmsusu@googlegroups.com

4. Kirk D.B., Wen-mei W. Hwu «Programming massively parallel processors. A hands-on approach», 2010, 279 ст.

5. Боресков А.В., Харламов А.А. «Основы работы с технологией CUDA», Москва 2010, 234 с.

6. J. Sanders, E.Kandrot «CUDA by example. An introduction to general-purpose GPU programming», 2010, 313 ст.

7. http://wiki.tntu.edu.ua/GPGPU

8. The Khronos Group: Open Standards, Royalty Free, Dynamic Media Techno-logies. http://www.khronos.org/opencl/

9. http://www.nvidia.com/object/cuda_directcompute.html

10. Nvidia's next generation CUDA compute architecture: Fermi (www.nvidia.com)

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