Створення методики проектування застосувань для програмно-апаратної платформи CUDA

Проаналізовано програмно-апаратні платформи, призначені для масивно-паралельних обчислень на відеоадаптерах. Описано принципи створення програм для платформи NVidia CUDA. Запропоновано методику роботи з динамічно розподіленою пам’яттю на основі шаблону «Стратегія» з використанням нотації UML. Відзна...

Повний опис

Збережено в:

Бібліографічні деталі
Дата:	2013
Автори:	Погорілий, С.Д., Верещинський, О.А.
Мова:	Ukrainian
Опубліковано:	Інститут програмних систем НАН України 2013
Назва видання:	Проблеми програмування
Теми:	Моделі та засоби паралельних і розподілених програм
Онлайн доступ:	https://nasplib.isofts.kiev.ua/handle/123456789/86677
Теги:	Додати тег Немає тегів, Будьте першим, хто поставить тег для цього запису!
Назва журналу:	Digital Library of Periodicals of National Academy of Sciences of Ukraine
Цитувати:	Створення методики проектування застосувань для програмно-апаратної платформи CUDA / С.Д. Погорілий, О.А. Верещинський // Проблеми програмування. — 2013. — № 3. — С. 47-60. — Бібліогр.: 10 назв. — укр.

Репозитарії

Digital Library of Periodicals of National Academy of Sciences of Ukraine

id	nasplib_isofts_kiev_ua-123456789-86677
record_format	dspace
spelling	nasplib_isofts_kiev_ua-123456789-866772025-02-09T14:23:38Z Створення методики проектування застосувань для програмно-апаратної платформи CUDA Погорілий, С.Д. Верещинський, О.А. Моделі та засоби паралельних і розподілених програм Проаналізовано програмно-апаратні платформи, призначені для масивно-паралельних обчислень на відеоадаптерах. Описано принципи створення програм для платформи NVidia CUDA. Запропоновано методику роботи з динамічно розподіленою пам’яттю на основі шаблону «Стратегія» з використанням нотації UML. Відзначено взаємозв’язок шаблонів паралельного програмування із математичним апаратом систем алгоритмічних алгебр. Описано підхід до проектування алгоритмів на основі шаблону «Команда». Реалізовано програмний інтерфейс, на основі якого створено паралельну версію алгоритму Данцига. 2013 Створення методики проектування застосувань для програмно-апаратної платформи CUDA / С.Д. Погорілий, О.А. Верещинський // Проблеми програмування. — 2013. — № 3. — С. 47-60. — Бібліогр.: 10 назв. — укр. 1727-4907 https://nasplib.isofts.kiev.ua/handle/123456789/86677 004.4 uk Проблеми програмування application/pdf Інститут програмних систем НАН України
institution	Digital Library of Periodicals of National Academy of Sciences of Ukraine
collection	DSpace DC
language	Ukrainian
topic	Моделі та засоби паралельних і розподілених програм Моделі та засоби паралельних і розподілених програм
spellingShingle	Моделі та засоби паралельних і розподілених програм Моделі та засоби паралельних і розподілених програм Погорілий, С.Д. Верещинський, О.А. Створення методики проектування застосувань для програмно-апаратної платформи CUDA Проблеми програмування
description	Проаналізовано програмно-апаратні платформи, призначені для масивно-паралельних обчислень на відеоадаптерах. Описано принципи створення програм для платформи NVidia CUDA. Запропоновано методику роботи з динамічно розподіленою пам’яттю на основі шаблону «Стратегія» з використанням нотації UML. Відзначено взаємозв’язок шаблонів паралельного програмування із математичним апаратом систем алгоритмічних алгебр. Описано підхід до проектування алгоритмів на основі шаблону «Команда». Реалізовано програмний інтерфейс, на основі якого створено паралельну версію алгоритму Данцига.
author	Погорілий, С.Д. Верещинський, О.А.
author_facet	Погорілий, С.Д. Верещинський, О.А.
author_sort	Погорілий, С.Д.
title	Створення методики проектування застосувань для програмно-апаратної платформи CUDA
title_short	Створення методики проектування застосувань для програмно-апаратної платформи CUDA
title_full	Створення методики проектування застосувань для програмно-апаратної платформи CUDA
title_fullStr	Створення методики проектування застосувань для програмно-апаратної платформи CUDA
title_full_unstemmed	Створення методики проектування застосувань для програмно-апаратної платформи CUDA
title_sort	створення методики проектування застосувань для програмно-апаратної платформи cuda
publisher	Інститут програмних систем НАН України
publishDate	2013
topic_facet	Моделі та засоби паралельних і розподілених програм
url	https://nasplib.isofts.kiev.ua/handle/123456789/86677
citation_txt	Створення методики проектування застосувань для програмно-апаратної платформи CUDA / С.Д. Погорілий, О.А. Верещинський // Проблеми програмування. — 2013. — № 3. — С. 47-60. — Бібліогр.: 10 назв. — укр.
series	Проблеми програмування
work_keys_str_mv	AT pogorílijsd stvorennâmetodikiproektuvannâzastosuvanʹdlâprogramnoaparatnoíplatformicuda AT vereŝinsʹkijoa stvorennâmetodikiproektuvannâzastosuvanʹdlâprogramnoaparatnoíplatformicuda
first_indexed	2025-11-26T20:12:03Z
last_indexed	2025-11-26T20:12:03Z
_version_	1849885119228149760
fulltext	Моделі та засоби паралельних і розподілених програм © С.Д. Погорілий, О.А. Верещинський, 2013 ISSN 1727-4907. Проблеми програмування. 2013. № 3 47 УДК 004.4 С.Д. Погорілий, О.А. Верещинський СТВОРЕННЯ МЕТОДИКИ ПРОЕКТУВАННЯ ЗАСТОСУВАНЬ ДЛЯ ПРОГРАМНО-АПАРАТНОЇ ПЛАТФОРМИ CUDA Проаналізовано програмно-апаратні платформи, призначені для масивно-паралельних обчислень на ві- деоадаптерах. Описано принципи створення програм для платформи NVidia CUDA. Запропоновано ме- тодику роботи з динамічно розподіленою пам’яттю на основі шаблону «Стратегія» з використанням нотації UML. Відзначено взаємозв’язок шаблонів паралельного програмування із математичним апара- том систем алгоритмічних алгебр. Описано підхід до проектування алгоритмів на основі шаблону «Ко- манда». Реалізовано програмний інтерфейс, на основі якого створено паралельну версію алгоритму Да- нцига. Вступ GPGPU – обчислення загального призначення з використанням графічних процесорів (GPU), які на сьогоднішній день є високопродуктивними багатоядер- ними процесорами, що характеризуються надвисокою обчислювальною потужністю та великою пропускною здатністю. Споча- тку – спеціально призначені для комп'юте- рної графіки і незручні для прикладного програмування, графічні процесори сьо- годні є універсальними паралельними пристроями з підтримкою цілком доступ- них програмних інтерфейсів та стандарт- них мов програмування, таких як С та С++ [1, 2]. У графічних API повністю відсутня підтримка взаємодії між пікселями, що об- робляються паралельно, які в принципі не дуже потрібно для графіки, але для обчис- лювальних задач є досить бажаним. Іншим суттєвим недоліком є відсутність підтрим- ки операції типу scatter (розподілити, рис. 1). Найпростішим прикладом операцій та- кого типу є побудова гістограм за вхідними даними, коли кожен новий елемент призво- дить до зміни наперед невідомого компо- нента (чи компонентів) гістограми. Рис. 1. Операції типу scatter Це пов’язано з тим, що в графічних API шейдер може здійснювати запис тільки в наперед визначене місце, оскільки для фрагментного шейдера наперед визначаєть- ся, який саме фрагмент він буде опрацьову- вати, і він може записати лише значення даного фрагмента. Ще одним «спадковим» недоліком є те, що розробка проходить одночасно двома мовами: для CPU та для GPU відпо- відно. Всі вищенаведені обставини усклад- нюють використання GPGPU і накладають ряд серйозних обмежень на алгоритми. То- му цілком природно виникла необхідність у створенні засобів розробки GPGPU про- грам, вільних від цих обмежень та орієнто- ваних на розв’язування складних задач. Як такі засоби виступають NVidia CUDA, OpenCL [3] та DX11 Compute Shaders [4]. OpenCL (Open Computing Language) – відкритий стандарт для паралельного програмування гетерогенних систем [3], який є першим відкритим безкоштовним стандартом для кросплатформенного пара- лельного програмування сучасних проце- сорів персональних комп'ютерах, серверах і портативних та вбудованих пристроях. OpenCL значно збільшує швидкість і гнуч- кість для широкого спектру застосувань (від ігор і розваг до наукових і медичних програм). OpenCL створюється групою бага- тьох передових компаній і установ під на- звою Khronos Group. Моделі та засоби паралельних і розподілених програм 48 У фреймворк OpenCL входять мова програмування, яка базується на стандарті C99, та інтерфейс програмування (API), що забезпечує паралельність на рівні інструк- цій та на рівні даних і є реалізацією техніки GPGPU. OpenCL – повністю відкритий стандарт, його використання доступне на базі вільних ліцензій. Мета OpenCL полягає у тому, щоб доповнити OpenGl і OpenAL, які є відкритими галузевими стандартами для тривимірної комп'ютерної графіки і звуку. DirectX Compute – технологія, створена компанією Microsoft на базі біблі- отеки Direct3D. Для розв’язування задач використовуються шейдери, але є можли- вість використання специфічних для Di- rectX API. NVidia CUDA. Запропонована ком- панією NVidia технологія CUDA помітно полегшує створення GPGPU-застосувань. Вона не використовує графічних АРІ і віль- на від обмежень, притаманних такого роду АРІ [5, 6]. При цьому слід чітко розуміти принципову різницю між потоками на CPU та на GPU: потоки на GPU мають дуже низьку «собівартість» створення, керування та ви- користання (контекст потоку мінімальний, всі регістри розподілені наперед); для ефективного використання GPU слід використовувати багато тисяч окремих потоків, тоді як для CPU достатньо 10 – 20 штук. Зазвичай типова програма з викори- станням CUDA має наступний вигляд: 1) виділення пам’яті на GPU; 2) копіювання даних між CPU та GPU; 3) запуск ядра (технологія Fermi дозволяє запуск кількох ядер одночасно); 4) копіювання результатів даних між GPU та CPU; 5) звільнення виділеної пам’яті на GPU. Для обробки масиву даних створю- ється багато потоків, і як наслідок, кожен потік обробляє один елемент вхідних да- них. Усі ці потоки виконуються паралель- но, потік може отримати інформацію про себе через набір змінних середовища. Важливим моментом є те, що хоча даний підхід і є подібним на роботу із SIMD-моделлю, але є принципові відмінно- сті (компанія NVidia використовує термін SIMT – Single Instruction Multiple Thread). Потоки розбиваються на групи по 32 шту- ки, які називаються варпами (warp). Тільки потоки в межах одного варпа виконуються фізично одночасно. Потоки з різних варпів можуть перебувати на різних стадіях вико- нання програми. При цьому керування вар- пами прозоро для програміста виконується самим GPU. Розподіл потоків насправді відображає ще один із основних прийомів використання CUDA–декомпозицію вихід- ної задачі на окремі підзадачі, які можуть бути розв’язані незалежно одна від одної. У зв’язку з усім вищезгаданим ме- тою роботи є створення методики проекту- вання застосувань для програмно-апаратної платформи CUDA на основі існуючого промислового інструментарію та застосу- вання UML і шаблонів паралельного про- грамування. Реалізація підходу до виділення пам’яті на базі шаблону «Стратегія» Робота з динамічно розподіленою пам’яттю. При використанні динамічно розподіленої пам’яті у програмах, створе- них мовами С та С++, зазвичай використо- вується один з методів виділення пам’яті (memory allocation) – функція malloc/free (С) чи new/delete (С++). Команда new забезпечує виклик конструк- торів за замовчуванням для кожного екзем- пляра класу, під який виділяється пам’ять. Після завершення роботи з даними виділе- ну пам’ять слід звільнити відповідно ко- мандою free або delete. Якщо цього не зробити, виникне класична проблема – про- тікання пам’яті (memory leak) – фактично пам’ять не звільнюється, а тому не може бути заново розподілена диспетчером [7, 8]. Технологія NVidia CUDA дозволяє напряму використовувати ресурси відеока- рти за допомогою набору API-функцій, які зокрема, забезпечують і доступ до відео- пам’яті. Щоб виділити необхідний обсяг слід скористатися функцією cudaMalloc (чи однією з подібних). Аналогічно, http://uk.wikipedia.org/wiki/%D0%9C%D0%BE%D0%B2%D0%B0_%D0%BF%D1%80%D0%BE%D0%B3%D1%80%D0%B0%D0%BC%D1%83%D0%B2%D0%B0%D0%BD%D0%BD%D1%8F http://uk.wikipedia.org/wiki/%D0%9C%D0%BE%D0%B2%D0%B0_%D0%BF%D1%80%D0%BE%D0%B3%D1%80%D0%B0%D0%BC%D1%83%D0%B2%D0%B0%D0%BD%D0%BD%D1%8F http://uk.wikipedia.org/w/index.php?title=%D0%A1%D1%82%D0%B0%D0%BD%D0%B4%D0%B0%D1%80%D1%82_C99&action=edit&redlink=1 http://uk.wikipedia.org/w/index.php?title=%D0%A1%D1%82%D0%B0%D0%BD%D0%B4%D0%B0%D1%80%D1%82_C99&action=edit&redlink=1 http://uk.wikipedia.org/wiki/%D0%86%D0%BD%D1%82%D0%B5%D1%80%D1%84%D0%B5%D0%B9%D1%81 http://uk.wikipedia.org/wiki/API http://uk.wikipedia.org/w/index.php?title=GPGPU&action=edit&redlink=1 http://uk.wikipedia.org/w/index.php?title=%D0%92%D1%96%D0%B4%D0%BA%D1%80%D0%B8%D1%82%D0%B8%D0%B9_%D1%81%D1%82%D0%B0%D0%BD%D0%B4%D0%B0%D1%80%D1%82&action=edit&redlink=1 http://uk.wikipedia.org/w/index.php?title=%D0%92%D1%96%D0%B4%D0%BA%D1%80%D0%B8%D1%82%D0%B8%D0%B9_%D1%81%D1%82%D0%B0%D0%BD%D0%B4%D0%B0%D1%80%D1%82&action=edit&redlink=1 http://uk.wikipedia.org/wiki/%D0%9B%D1%96%D1%86%D0%B5%D0%BD%D0%B7%D1%96%D1%8F http://uk.wikipedia.org/wiki/OpenGL http://uk.wikipedia.org/w/index.php?title=OpenAL&action=edit&redlink=1 Моделі та засоби паралельних і розподілених програм 49 пам’ять звільняється командою cudaFree. Доступ до даних, розміщених на відео карті може здійснювати лише GPU. Розглянемо копіювання динамічно розміщених даних. При простому копію- ванні вказівників отримуємо так зване «по- верхневе» копіювання, яке має сенс лише в межах одного типу пам’яті. В цьому випад- ку дані безпосередньо не копіюються. Це може призвести до проблем часу виконання при неправильному звільненні пам’яті. Для глибокого копіювання необхідно заново виділити потрібний обсяг пам’яті та скори- статися однією з багатьох функцій для ко- піювання (memcpy, std::copy, cudaMemcpy, тощо). Зокрема, копіювання між різними типами пам’яті можливе лише з використанням CUDA API. Розв’язати деякі з вказаних проблем дозволяють так звані «розумні» вказівники (smart pointers). Як приклад можна навести shared_ptr із стандартної бібліотеки STL чи із бібліотеки Boost. Вони дозволяють проводити вдосконалену політику керуван- ня пам’яттю (memory management). Але ці конструкції не розраховані на роботу в се- редовищі CUDA. Таким чином формується завдання про створення деякого уніфікованого інте- рфейсу розумного вказівника, який дозво- лив би однаково працювати з даними, неза- лежно від їх розміщення, приховуючи реа- лізацію низькорівневих операцій, які стали б прозорими для користувача (тут і надалі, користувач – розробник, що використовує можливості створеного вказівника). Можна сформувати низку вимог до створюваного інтерфейсу (надалі – UnifiedPtr). Автоматизація виділення та звіль- нення пам’яті, що дозволить наступні операції: // Вказівник для виділення пам’яті в «купі» для 10-ти цілих чисел HeapUnifiedPtr<int> h(10); // Вказівник для виділення пам’яті на відеокарті для 10-ти цілих чи- сел DeviceUnifiedPtr<int> d (10); // Власне, виділення пам’яті h.Allocate(); d.Allocate(); // Явне звільнення пам’яті в змін- ній h, h.Deallocate(); Реалізація прозорого копіювання да- них з використанням різних типів UnifiedPtr, що зробить можливим наступні присвоювання: HeapUnifiedPtr h1, h2; DeviceUnifiedPtr d1, d2; // Присвоювання працюють коректно h2 = h1; // Host => Host d1 = h2; // Host => Device d2 = d1; // Device => Device h1 = d2; // Device => Host Як описано вище, на UnifiedPtr на- кладається низка вимог, сформованих зі сторони функціоналу. Таким чином постає проблема настроювання компонента. Не- хай, наприклад, треба створити вказівник для виділення динамічної області з глибо- ким копіюванням для одно поточного сере- довища, а також вказівник для виділення пам’яті на відео карті з підрахунком поси- лань для багато поточного середовища. Виникає проблема створення досить гнучкого інтерфейсу, який міг би дати ко- ристувачу можливість настроювання окре- мих компонентів, і водночас мав розумні розміри. Реалізація всього можливого функ- ціоналу під оболонкою деякого універсаль- ного інтерфейсу не є вдалим вибором, оскі- льки призведе до створення громіздких компонентів (яскравим прикладом в даному випадку є клас std::string із стандартної бібліотеки С++, переповнений функціона- лом та дублюванням операцій). Програми, що будуть використовувати такі класи пра- цюють набагато повільніше, аніж аналогіч- ні програми, розроблені самотужки [2]. Але однією із найбільших проблем є втрата безпеки статичних типів. Однією із цілей будь якої архітектури є закладання деяких аксіом «за означенням». В ідеаль- ному випадку більшість обмежень мають накладатися ще на етапі компіляції коду. В межах громіздкого інтерфейсу дуже важко врахувати всі подібні обмеження. Це приз- водить до збільшення розриву між семан- тичною та синтаксичною правильністю програми. Тобто виникає небезпека напи- Моделі та засоби паралельних і розподілених програм 50 сання синтаксично правильного коду, який може зруйнувати основні зв’язки всередині бібліотеки. Один з найочевидніших підходів – фактично, груба сила, – полягає в реалізації різних проектних рішень у вигляді окремих класів більш скромного розміру. Кожен з таких класів міг би інкапсулювати набір взаємопов’язаних алгоритмів, які б забез- печували швидку і налагоджену поведінку. В такому випадку логічно очікувати появи численних реалізацій типу: HeapDeepCopySingleThreadPtr DeviceRefCountMultiThreadPtr Для такого розв’язку проблеми ха- рактерним є різкий ріст кількості різних варіантів проектних рішень. Тільки два ви- щезгаданих класи можуть легко призвести о виникнення багатьох нових комбінацій. Перетворення типів викличе ще більшу кількість варіацій, які, рано чи пізно, вий- дуть за рамки можливостей програміста та користувача бібліотеки. Як бачимо, подо- лати експоненційний ріст варіантів за до- помогою грубої сили неможливо. Отже проблеми такого підходу – на- яву: масштабованість даного рішення практично нульова. Щоб додати новий тип даних потрібно повністю його описувати; велика частина коду просто копі- юється (Copy-Paste); отримані структури даних ніяк не зв’язані між собою, хоча існують деякі ло- гічні операції (наприклад, присвоювання, копіювання), які передбачають існування зв’язку. Фактично, ми нехтуємо можливос- тями, які дає використання поліморфізму; бібліотеки такого типу не потре- бують значних інтелектуальних зусиль для розробки, але є вкрай негнучкими. Най- менша непередбачувана дрібниця може привести всі класи в плачевний стан. Щоб уникнути значної частки опи- саних проблем, можна використати спадку- вання. Це, в свою чергу призведе до ство- рення складних ієрархій, які доведеться описувати вручну. Знову ж таки при роз- ширенні архітектури програми в області алокації пам’яті кількість нових класів зро- статиме геометрично. Шаблон «Стратегія» (Strategy) Стратегія – шаблон поведінки об’єктів [5, 7]. Шаблон визначає сімейство алгори- тмів, інкапсулює їх та робить взаємозамін- ними. Стратегія дозволяє змінювати алго- ритми незалежно від клієнтів, що їх вико- ристовують. Учасники (рис. 2): Strategy – Стратегія. Оголошує спільний для всіх підтри- муваних алгоритмів інтерфейс. Concrete Strategy – конкретна страте- гія. Реалізує алгоритм, який використо- вує інтерфейс, оголошений стратегією. Context – контекст. Конфігурується об’єктом класу Concrete Strategy; зберігає вказівник на Strategy; може оголосити інтерфейс, який до- зволить стратегії отримати доступ до даних контексту. +ContextInterface() Context +AlgorithmInterface() Strategy +AlgorithmInterface() ConcreteStrategyC +AlgorithmInterface() ConcreteStrategyB +AlgorithmInterface() ConcreteStrategyA Рис. 2. Діаграма класів, що відображає шаблон «Стратегія» Моделі та засоби паралельних і розподілених програм 51 Результати застосування шаблону «Стратегія» (позитивні та негативні): спорідненість алгоритмів; альтернатива породженню підкласів; за допомогою стратегій можна уник- нути використання операторів розгалу- ження; вибір реалізації на етапі виконан- ня/компіляції; клієнти повинні «знати» про різні стратегії; обмін інформацією між клієнтом та стратегією; збільшення кількості об’єктів; реалізація стратегій, як шаблонних параметрів. В С++ для конфігурації стратегій можна використати шаблонні параметри. Клас конфігурується стратегією у момент інстанціювання. При такому використанні стратегій відпадає необхідність у абстрак- тному класі для визначення інтерфейсу. Крім того передача у вигляді шаблона до- зволяє статично зв’язати стратегію з кон- текстом, підвищуючи таким чином ефек- тивність програми. У роботі запропоновано викорис- тання класів стратегій для настроювання функціоналу на етапі компіляції програми (інформація про стратегії передається у вигляді шаблонних параметрів). Тип успа- дковує свої стратегії, успадковуючи таким чином необхідний функціонал [7]. Такий підхід призводить до автома- тичної генерації компілятором необхідних ієрархій класів. Перевагою є відсутність динамічного поліморфізму, а, як наслідок, – пов’язаних з цим затримок. Декомпозиція UnifiedPtr на стратегії Стратегія виділення пам’яті. Призначення цієї стратегії – динамічне виділення пам’яті. Таким чином вона пос- тачає набір функцій: void * AllocateMemory(size_t) void * AllocateMemory(size_t, size_t) void DeallocateMemory(void) для централізації інтерфейсу всі стратегії виділення пам’яті успадковуються від кла- су Allocator (фактично, це не обов’язко- во, достатньо реалізувати дві вищеописані функції. Додатково виділено дві групи роз- поділювачів пам’яті (рис. 3): лінійні роз- поділювачі (пам’ять виділяється одним нероздільним шматком в адресному прос- торі, всі дані розміщені послідовно) та ма- тричні розподілювачі (при виділенні на відеоадаптері пам’яті під матрицю кожен рядок вирівнюється за адресою, кратною 64 або 128 байт, що значно пришвидшує роботу з глобальною пам’яттю). +AllocateMemory(в i_size : unsigned long) : void +AllocateMemory(в i_width : unsigned long, в i_height : unsigned long) : void* +DeallocateMemory(в ip_data : void) +GetPitch() : long MemoryAllocator HeapAllocator HeapPitchAllocatorDeviceAllocator DevicePitchAllocator Рис. 3. Діаграма класів, що відображає стратегію виділення пам’яті Моделі та засоби паралельних і розподілених програм 52 Стратегія копіювання даних. Стратегія копіювання даних тісно пов’язана з методом виділення пам’яті та стратегією зберігання інформації. Для створення UnifiedPtr доцільно створити деякий універсальний клас, який реалізуватиме глибоке копіювання даних. Необхідну функцію компілятор може ви- бирати сам на основі інформації, закладе- ної в шаблонні параметри. Таким чином мовою С++ можна записати наступну сиг- натуру для функції копіювання: template < Class Allocator1, Class Allocator2 > void Copy(void ip_src, void * ip_dest, size_t i_size) Такий вигляд дозволяє уникнути операторів розгалуження у тексті підпрог- рами і полегшує копіювання даних між різними пристроями. Стратегії зберігання даних. Стра- тегія зберігання реалізує логіку зберігання та копіювання даних. Розглянемо існуючі стратегії та мо- жливість їх застосування до UnifiedPtr. Стратегія глибокого копіювання (DeepCopyStorage Policy). Кожен екземпляр класу містить вка- зівник на свою ділянку пам’яті (рис. 4). При копіюванні відбувається повне (гли- боке) копіювання даних. При знищенні об’єкта пам’ять звільняється. До плюсів такого підходу можна віднести відносну простоту реалізації. Суттєвим мінусом є значні накладні витрати на копіювання та зберігання інформації. Стратегію показано на рисунку. Heap p1 Data Heap p1 Data p2 Data p2=p1 Рис. 4. Схематичне зображення глибокого копіювання даних Для UnifiedPtr дана стратегія є ціл- ком застосовною, оскільки не накладає жодних обмежень, що легко реалізується через клас CopyFunctor. Стратегія володіння (Ownership Strategy). Ресурсом (у даному випадку – пам’яттю) володіє лише один об’єкт (рис. 5). При присвоювання відбувається пере- дача права власності. Пам’ять звільнюєть- ся лише один раз при знищенні екземпля- ра, який володіє ресурсом у поточний мо- мент часу. p1 Data Heap p1 Data Heap p2 p2=p1 Рис. 5. Схематичне зображення стратегії володіння даними Така стратегія значно зменшує за- трати, пов’язані з копіюванням інформації. Але, оскільки в нашому випадку, спосіб звільнення пам’яті жорстко зв’язаний із способом її виділення і задається ще на етапі компіляції, то копіювання можливе лише в межах одного пристрою. Стратегія підрахунку посилань (Reference Count Strategy) Окрім власне корисної інформації в динамічній пам’яті зберігається також лі- чильник посилань (рис. 6). Пам’ять звіль- няється тоді, коли лічильник стає рівний нулю. Таким чином копіювання призво- дить до збільшення лічильника на одини- цю. p1 Data Heap p1 Heap p2 p2=p1 1 Data 2 Рис. 6. Схематичне зображення стратегії підрахунку посилань Для використання з UnifiedPtr про- понується провести деяку модифікацію. Будемо зберігати два лічильники посилань (рис. 7). При копіюванні між CPU та відео- картою виконується глибоке копіювання, а після чого можна використовувати класи- чний підхід підрахунку посилань. Моделі та засоби паралельних і розподілених програм 53 Device Memory Data 01 h1 Heap Data 02 h1 Heap Device Memory h2 Data 12 Data h1 Heap Device Memory h2 d1 Data 22 Data h1 Heap Device Memory h2 d1 d2 h2=h1 d2=d1 d1=h2 Рис. 7. Стратегія підрахунку посилань, адаптована для CUDA Інкапсуляція операцій в об’єктах для декомпозиції задачі Шаблон «Команда». «Команда» (інші назви – оператор, функтор), як і «Стратегія» – шаблон поведінки об’єктів. Він дозволяє інкапсулювати запит як об’єкт, дозволяючи тим самим задавати параметри клієнта для обробки відповід- них запитів, ставити запити в чергу, про- токолювати і підтримувати операцію «Ві- дмінити» [5, 7]. Інколи необхідно посилати об’єктам запити, нічого не знаючи про те, виконання якої операції необхідно та хто є адресатом запиту. Шаблон проектування «Команда» дозволяє надсилати запити де- яким об’єктам програми, перетворюючи сам запит в об’єкт. Його можна передавати та зберігати як і будь-який інший об’єкт. В основі лежить абстрактний клас, який ого- лошує інтерфейс для використання опера- ції. У найпростішій формі цей інтерфейс містить єдину функцію Execute (виконати операцію). Конкретні підкласи містять інформацію про те, хто є адресатом, і що саме треба виконувати. У адресата є набір даних, необхідних для успішного вико- нання операції. Учасники (згідно з рис. 8Рис. 8): Command – команда. Оголошує інтерфейс для виконання операції. ConcreteCommand – конкретна ко- манда. Визначає зв’язок між об’єктом ад- ресатом та дією; реалізує операцію Execute шляхом виклику відповідних операцій. Client – клієнт. Створює об’єкт класу Concrete- Command та встановлює його одержувача. Invoker – ініціатор. +Execute() Command +Execute() -State ConcreteCommand Invoker 1 Client +Action() Receiver Рис. 8. Структура шаблону «Команда» Моделі та засоби паралельних і розподілених програм 54 Звертається до команди для вико- нання запиту. Receiver – одержувач запиту. Містить інформацію про способи виконання операції. У ролі одержувача може виступати будь-який клас. У результаті використання «Коман- ди» розривається зв’язок між об’єктом, що ініціює виконання операції та об’єктом, що містить операцію, необхідну для її вико- нання. Оскільки команди – це справжні об’єкти, то допускаються довільні маніпу- ляції з їх участю, розширення, тощо. Як наслідок із простих команд можна збирати складніші, об’єднуючи їх в ланцюжки опе- рацій. Для додавання нових операцій не потрібно змінювати вже існуючі класи. Зв’язок шаблону «Команда» та систем алгоритмічних алгебр (САА) В.М. Глушкова САА – потужний математичний апарат для формалізації алгоритмів, запро- понований В.М. Глушковим. Основною перевагою САА є спрямування на проек- тування схем алгоритмів і можливість представлення довільного алгоритму у вигляді алгебраїчної формули. Розвинені методи формальних перетворень САА до- зволяють проводити оптимізацію схеми, абстрагуючись від конкретної програмної реалізації. Крім того САА цілком відпові- дають основним принципам структурова- ного програмування [9]. Наступним кроком було розширен- ня математичного апарату на випадок ба- гатопроцесорних архітектур, за рахунок введення додаткових операцій, спрямова- них на паралельну обробку даних. Введен- ня нових операцій призвело до виникнення модифікованих САА (САА-М). Як відомо САА оперують поняттям операторів, пов’язуючи їх у цілісні алгори- тми з допомогою формульного запису. Використовуючи шаблон «Команда» для представлення таких «примітивних» алго- ритмів можна реалізовувати схеми, подані у вигляді САА, створюючи макрокоманди на рівні об’єктів (окремі оператори можна об’єднувати у послідовності). Таким чи- ном, САА виконують роль абстракції від конкретної платформи, зосереджуючись на особливостях роботи алгоритму. Перенесення схеми на конкретну робочу платформу покладається на окремі оператори, які виступають деяким промі- жним етапом в даному випадку. Як приклад, розглянемо дуже прос- ту схему, записану з використанням нота- ції САА: 1 ( 2 3)Scheme Step Step Step Cond (1) Наступна ієрархія класів дозволяє представити дану схему. На рис. 9 нижні два рівні представляють абстракції, приз- начені для уніфікації роботи з класами. На вищому рівні ієрархії (внизу рисунка) різ- ними кольорами показано по дві конкретні реалізації кожного оператора. Крім того ніщо не забороняє змішувати різнотипні реалізації, дозволяючи таким чином отри- мати вісім можливих програмних предста- влень наведеної схеми. +Operate() Step +Operate() Step1 +Operate() ConcreteStep1_1 +Operate() ConcreteStep1_2 +Operate() Step2 +Operate() ConcreteStep2_1 +Operate() ConcreteStep2_2 +Operate() Step3 +Operate() ConcreteStep3-1 +Operate() ConcreteStep3_2 Рис. 9. Ієрархія класів для реалізації простої САА схеми Моделі та засоби паралельних і розподілених програм 55 Тоді програмний код буде мати на- ступний вигляд: Step step1, step2, step3; // Ініціалізація даних та команд ... // Виконання схеми Step1->Operate(); If (Condition) Step2->Operate(); else Step3->Operate(); Крім того, умовний оператор також можна представити у вигляді команди, повністю «автоматизувавши» таким чи- ном виконання алгоритму. Можна доби- тися навіть зовнішньо настроювання ал- горитму з допомогою файлу конфігурації чи набору команд користувача. Тобто на вхід програми поступає САА схема, за якою збирається настроєний користува- чем алгоритм. Слід зауважити, що такий підхід потребує чіткої декомпозиції задачі на окремі підзадачі (не обов’язково при- мітивні). Таким чином, можна використати схеми, записані з допомогою САА для ре- алізації алгоритмів допомогою технології NVidia CUDA. Для цього слід використати механізм шаблонів С++: template<class Operation> __global__ void Kernel(Operation i_operation) { i_operation.Operate(); } Вищенаведена функція – узагальне- не ядро CUDA, яке абсолютно не залежить від реалізації оператора. Функція Operate() в свою чергу є тілом функтора і може бути як віртуальною, так і статич- ною. Програмна реалізація стратегій Стратегія виділення пам’яті. Для виділення пам’яті використовується клас Allocator. Він об’єднує дві функції – виді- лення та звільнення пам’яті. В роботі реа- лізовано чотири підкласи: HeapAllocator – стандартне ви- ділення динамічної пам’яті можливостями мови С++. Власне в програмі використову- ється метод виділення malloc/free, що до- зволяє уникнути необхідності в конструк- торі за замовчуванням, а також деяких на- кладних витрат, пов’язаних із початковою ініціалізацією даних; DeviceAllocator – виділення пам’яті на відео карті з використанням інтерфейсу CUDA; PitchDeviceAllocator – виді- лення пам’яті на відео карті з використан- ням інтерфейсу CUDA з використанням автоматичного вирівнювання даних в пам’яті, що призводить до покращення швидкодії при читанні даних; PitchHeapAllocator – відповід- ний розподілювач пам’яті для централь- ного процесора. Оскільки вирівнювання не є настільки важливим у даному випад- ку, то відбувається просто емуляція функ- ціоналу. У подальшому клас Allocator може бути легко розширено. Наприклад, можна реалізувати роботу з текстурною пам’яттю відеокарти, розміщення даних у стеку програми, тощо. Стратегія копіювання даних. Для забезпечення легкого і прозорого копію- вання інформації створено клас CopyFunctor, який має один статичний метод Copy. Для інстанціювання класу використовується інформація про два кла- си типу Allocator: template < class AllocationPolicy1, class AllocationPolicy2 > class CopyFunctor У залежності від вхідної комбінації типів компілятор автоматично вибирає необхідну функцію. На даному етапі пред- ставлено 4 перезавантажених методи для копіювання типу: host – host, host – device, Моделі та засоби паралельних і розподілених програм 56 device – device, device – host, які відповідають всім можливим комбіна- ціям нащадків класу Allocator. Стратегія зберігання даних. В ро- боті реалізовано вищеописані типи страте- гій зберігання даних. Для уніфікації коду та уникнення повторювань всі стратегії у свою чергу є нащадками неабстрактного базового класу StoragePolicy. Розгляне- мо його інтерфейсну частину. template < typename T, class AllocationPolicy > class Storage : public AllocationPolicy { public: void SetSize(size_t i_size); size_t GetSize() const; bool Empty() const; void Allocate(); void Deallocate(); T GetImpl(); }; Стратегія реалізує загальні функції для задання розміру необхідного буфера, виділення та звільнення пам’яті. Власне функція GetImpl() забезпечує прямий доступ до буфера пам’яті, покла- даючи таким чином відповідальність за роботу з буфером на користувача. У роботі реалізовано програмно три вищеописані стратегії зберігання даних. У нотації UML ми отримаємо наступну діаг- раму (рис. 10). Оскільки базова стратегія вже реа- лізує необхідні функції для роботи з пам’яттю, класи-нащадки мають реалізу- вати лише ту логіку, яка специфікується даним типом, делегуючи виконання більш загальних операцій базовому класу, який за необхідності може передати їх далі по ієрархії. Таким чином стратегії зберігання реалізовують лише набір необхідних ву- зькоспеціалізованих конструкторів та операторів присвоювання. Вони власне інкапсулюють основні відмінності між класами, які між іншим, також встанов- люються компілятором на основі інфор- мації про базові типи даних. +SetSize(в i_size : int) +GetSize() : int +Empty() : bool +Allocate() +Deallocate() +GetImpl() Storage typename T, class AllocationPolicy DeepCopyStorage typename T, class AllocationPolicy ReferenceCountStorage typename T, class AllocationPolicyDestroyCopyStorage typename T, class AllocationPolicy Рис. 10. Діаграма, що відображає ієрархію стратегій зберігання даних Моделі та засоби паралельних і розподілених програм 57 Опис інтерфейсу UnifiedPtr На основі описаних стратегій реалі- зовано клас UnifiedPtr, який дозволяє однаково працювати з даними розміщени- ми як на відео карті, так і в оперативній пам’яті комп’ютера. Далі наведено заголо- вок класу: template < typename T, class AllocationPolicy, template <class, class> class StoragePolicy > class UnifiedPtr : public StoragePolicy<T, AllocationPolicy> Як видно, клас повністю специфі- кується трьома шаблонними типами, ус- падковуючи функції стратегій. Таким чи- ном, для настроювання поведінки класу достатньо змінити будь-яку з стратегій, а для розширення слід просто описати від- повідну частину функціоналу. Наприклад, для додавання функцій, що забезпечать правильну роботу класу в багатопоточно- му середовищі процесора можна створити нову стратегію зберігання даних MultithreadingStoragePolicy і пере- дати її як шаблонний параметр. Таким чином значну роботу виконує компіля- тор, який при інстанціюванні шаблонно- го типу генерує відповідні ієрархії кла- сів. На діаграмі рис. 11 показано реаль- ну ієрархію класів, що створюється компі- лятором. Як бачимо, на основі вже реалі- зованих типів даних можна створити два- надцять нових типів, які легко зробити взаємозамінними, настроювати та підтри- мувати. Створений програмний інтерфейс успішно пройшов низку перевірок основ- ного функціоналу, для чого було викорис- тано бібліотеку UnitTest++. UnifiedPtr T, AllocationPolicy, StoragePolicy ConcreteStoragePolicy T, AllocationPolicy StoragePolicy T, AllocationPolicy AllocationPolicy T HeapAllocator DeviceAllocator DestroyCopyPolicy ReferenceCountPolicy DeepCopyPolicy PitchHeapAllocator PitchDeviceAllocator Рис. 11. Схема класу UnifiedPtr Реалізація базового класу-оператора Шаблон «Команда» дозволяє ство- рити узагальнене ядро CUDA. Для можли- вості вибору відповідного алгоритму мож- на застосувати принципи успадкування та поліморфізму. У роботі запропоновано підхід, що ґрунтується на можливостях статичного поліморфізму, а саме шаблонних парамет- рах С++. Прийом, що застосовується но- сить назву CRTP – рекурентний шаблон- ний паттерн, який дозволяє використову- вати клас нащадок при інстанціюванні ба- тьківського класу. Наприклад, template<class TDerived> class Base {}; class Derived : public Base<Derived> {}; Як не дивно, така конструкція не викликає зациклювання компілятора (до- зволено стандартом С++99) і водночас від- криває цікаві можливості. Зокрема, в робо- ті використано підхід, що дозволяє емулю- вати віртуальність функції, визначаючи необхідну адресу на етапі компіляції. Моделі та засоби паралельних і розподілених програм 58 На рис. 12 наведено відповідну діа- граму UML. +__device__ Operate() +GetData() : UnifiedPtr +Setdata(в ip_data : UnifiedPtr) Operation class concreteOperation static_cast<ConcreteOperation>(*this).DoOperate() +__device__ DoOperate()() ConcreteOperation Рис. 12. Схема класу-оператора з підтримкою статичного поліморфізму Як видно з рисунка, всередині класу застосовується складне приведення типів та виклик невіртуальної функції. Таким чином клас-нащадок зобов’язаний реалізу- вати функцію _DoOperate(). Як вже було сказано, перевагою та- кого підходу є, безумовно, краща швидко- дія. До мінусів можна віднести необхід- ність перекомпіляції ядра при використан- ні нового оператора. Втім цього недоліку можна уникнути створивши набір ядер, інстанційованих різними операторами. У роботі створено власне базовий каркас для оператора, який слід успадко- вувати при створенні нових алгоритмів. Як перспективу розвитку слід зауважити мож- ливість об’єднання команд в ланцюжки, створення об’єктів-умов, об’єктів-циклів і т. д. Як вже було сказано, це дає мож- ливість перенести алгоритми, записані за допомогою абстрактних засобів (САА) на конкретну програмну реалізацію. Приклад реалізації алгоритму Данцига для платформи CUDA Паралельна схема алгоритму в нотації САА. Одним із найефективніших методів пошуку всіх найкоротших шляхів у зваженому орієнтованому графі є алгоритм Данцига. Його складність становить O(n 3 ), де n – кількість вершин графа [10]. Вихідними даними для алгоритму є матриця вагових коефіцієнтів. Ідея полягає у послідовному обчисленні з допомогою рекурентної процедури підматриць найкоротших шляхів Dm зростаючої розмірності m x m. Кожна така матриця, фактично, є матрицею найкоротших шляхів підграфа з вершинами від 0 до m-1. Оскільки алгоритм має квадратичну складність, то він є перспективним з точки зору розпаралелювання. Розглянуто мож- ливість розпаралелювання алгоритму для систем із спільною пам’яттю. Паралельна версія алгоритму описується наступною САА-схемою: 2 2 1 2 2 1 2 2 1 { ( ( , 2 , ) ( )) ( ( , 2 , ) ( )) ( ( , 2 , ) ( )) ( )}. N N m N l N N l N N l Dancig PSub1 m N l B l PSub2 m N l B l PSub3 m N l B l m (2) Завдяки використанню шаблонів можна застосувати дану схему для реалізації програми на платформі CUDA. Результати тестування алгоритму На рис. 13 показано графік для по- рівняння часових характеристик роботи алгоритму з використанням технології CUDA та з використанням багатопо- точності у системі з чотирьохядерним процесором. Дані було отримано для ви- бірки графів з кількістю вершин від 100 до 1000. Як видно з графіка, для графів ни- зької розмірності (до 500 вершин) багато- поточна версія є виграшною за часом. При зростанні навантаження (збільшення кіль- кості вершин) версія для відеоадаптера дає виграш, який поступово зростає з ростом розмірності графа. Для пошуку найкорот- ших шляхів для 1000 вершин необхідно близько 900 мс на відеоадаптері, що у 2.5 рази швидше за програму для CPU. Такий хід графіків цілком пояс- нюється специфікою алгоритму Данци- га, який працює з ітеративним набором Моделі та засоби паралельних і розподілених програм 59 Рис. 13. Порівняння часових характеристик програм для GPU та CPU матриць зростаючої розмірності (1х1, 2х2, 3х3 і т. д.). Таким чином для графів малої розмірності завантаження відеоадаптера не є оптимальним, отже його ресурс не вико- ристовується повністю. При переході до великих розмірностей графів більшу час- тину роботи алгоритму GPU повністю за- вантажений обчисленнями, що й зумовлює виграш у часі. Як видно з графіка, час роботи ал- горитму зростає повільніше для GPU ніж для CPU, що дає змогу зробити висновок про доцільність його використання для графів високої розмірності. Висновки У роботі запропоновано підхід до реалізації уніфікованого інтерфейсу для роботи із динамічно розподіленою пам’яттю (діаграма класів на рис. 12), який базується на використанні шаблону проек- тування «Стратегія». Здійснено розклад класу на основні стратегії та розглянуто можливість їх застосування для програм, призначених для платформи CUDA. Виді- лено наступні класи стратегій: стратегія виділення пам’яті; стратегія копіювання даних; стратегія зберігання даних. Описано метод створення об’єктів- операторів на основі шаблону «Команда», який дозволяє реалізувати абстрактні схе- ми алгоритмів за рахунок використання можливостей статичного поліморфізму. Запропоновані підходи представле- но за допомогою UML-нотації, яка є уні- версальним засобом проектування як ієра- рхій класів, так і діаграм потоків, що відо- бражають роботу програми. Для використання описаних у робо- ті підходів створено бібліотеку мовою С++, яку використано для програмної реа- лізації алгоритму Данцига, представленого у вигляді абстрактної схеми за допомогою САА Глушкова. Алгоритм реалізований для відеоадаптера дає виграш у часових характеристиках приблизно у 2.5 рази по- рівняно з чотирьохядерним процесором Intel Core 2 Quad, що ефективну підтвер- джує застосовність запропонованого мето- ду проектування програм та перспектив- ність його розвитку. 1. Боресков А.В., Харламов А.А. Основы ра- боты с технологией CUDA. – М.: ДМК Пресс, 2010. – 232 с. 2. Sanders J., Kandrot E. CUDA by Example: An Introduction to General-Purpose GPU Programming. – Addison-Wesley Profes- sional, 2010. – 312 p. 0 500 1000 1500 2000 2500 0 200 400 600 800 1000 1200 GPU CPU (4 Cores) Кількість вершин Час, мс Моделі та засоби паралельних і розподілених програм 60 3. The Khronos Group: Open Standards, Royalty Free, Dynamic Media Techno-logies. http://www.khronos.org/opencl/ 4. Direct Compute. http://www.nvidia.com/object/cuda_directco mpute.html 5. Александреску А. Современное проектиро- вание на С++. Серия C++ in Depth.: Пер. с англ. – М.: Издательский дом «Вильямс», 2008. – 336 с. 6. Погорелый С.Д., Бойко Ю.В., Трибрат М.И., Грязнов Д.Б. Анализ методов повы- шения производительности компьютеров с использованием графических процессоров и апаратно-програмной платформы CUDA // Математичні машини і системи. – 2010. – № 1. – С. 40–54. 7. Гамма Э., Хелм Р., Джонсон Р., Влисси- дес Д. Приемы объекто-ориентированого проектирования // Паттерны проектирова- ния. – СПб: Питер, 2009. – 366 с. 8. Саттер Г. Решение сложных задач на С++. Серия C++ in Depth.: Пер. с англ. – М.: Издательский дом «Вильямс», 2008. – 400 с. 9. Ющенко Е.Л., Цейтлин Г.Е., Грицай В.П. и др. Многоуровневое структурное проекти- рование программ // Теоретические осно- вы, инструментарий.– М.: Финансы и ста- тистика, 1989. – 342 с. 10. Майника Э. Алгоритмы оптимизации на сетях и графах. – М.: Мир, 1981. – 324 с. Одержано 22.03.2013 Про авторів: Погорілий Сергій Демьянович, доктор технічних наук, професор, Верещинський Олег Андрійович, аспірант другого року навчання. Місце роботи авторів: Київський національний університет імені Тараса Шевченка, 01601, вул. Володимирська, 60, м. Київ, Україна. Тел. +38 (093) 080 7975. E-mail: sdp@univ.net.ua, oleg.vereshchynsky@gmail.com http://www.khronos.org/opencl/ http://www.nvidia.com/object/cuda_directcompute.html http://www.nvidia.com/object/cuda_directcompute.html mailto:sdp@univ.net.ua mailto:oleg.vereshchynsky@gmail.com

Створення методики проектування застосувань для програмно-апаратної платформи CUDA

Репозитарії

Схожі ресурси