Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил
Использование графических ускорителей позволяет достичь высокой производительности, однако требует от разработчика низкоуровневого программирования и знания деталей аппаратной и программной платформы. В работе предложен подход к автоматизации разработки приложений для графических ускорителей, основа...
Збережено в:
| Дата: | 2009 |
|---|---|
| Автори: | , |
| Формат: | Стаття |
| Мова: | Russian |
| Опубліковано: |
Інститут програмних систем НАН України
2009
|
| Теми: | |
| Онлайн доступ: | https://nasplib.isofts.kiev.ua/handle/123456789/4419 |
| Теги: |
Додати тег
Немає тегів, Будьте першим, хто поставить тег для цього запису!
|
| Назва журналу: | Digital Library of Periodicals of National Academy of Sciences of Ukraine |
| Цитувати: | Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил / А.Е. Дорошенко, К.А. Жереб // Пробл. програмув. — 2009. — № 3. — С. 3-18 — Бібліогр.: 22 назв. — рос. |
Репозитарії
Digital Library of Periodicals of National Academy of Sciences of Ukraine| id |
nasplib_isofts_kiev_ua-123456789-4419 |
|---|---|
| record_format |
dspace |
| spelling |
nasplib_isofts_kiev_ua-123456789-44192025-02-09T14:42:18Z Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил Розробка високопаралельних застосувань для графічних прискорювачів з використанням переписувальних правил Development of highly-parallel applications for graphical processing units using rewriting rules Дорошенко, А.Е. Жереб, К.А. Моделі та засоби паралельних і розподілених програм Использование графических ускорителей позволяет достичь высокой производительности, однако требует от разработчика низкоуровневого программирования и знания деталей аппаратной и программной платформы. В работе предложен подход к автоматизации разработки приложений для графических ускорителей, основанный на использовании парадигмы переписывающих правил. Використання графічних прискорювачів до-зволяє досягти високої продуктивності, проте потребує від розробника низькорівневого програмування та знання деталей апаратної та програмної платформ. В роботі запропоновано підхід до автоматизації розробки застосувань для графічних прискорювачів, заснований на використанні парадигми переписувальних правил. Using graphical processing units (GPUs) allows increasing performance significantly, but requires low-level programming and detailed understanding of underlying hardware and software platform. The paper presents a technique for automating GPU application development, based on rewriting rules paradigm. 2009 Article Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил / А.Е. Дорошенко, К.А. Жереб // Пробл. програмув. — 2009. — № 3. — С. 3-18 — Бібліогр.: 22 назв. — рос. 1727-4907 https://nasplib.isofts.kiev.ua/handle/123456789/4419 681.3 ru application/pdf Інститут програмних систем НАН України |
| institution |
Digital Library of Periodicals of National Academy of Sciences of Ukraine |
| collection |
DSpace DC |
| language |
Russian |
| topic |
Моделі та засоби паралельних і розподілених програм Моделі та засоби паралельних і розподілених програм |
| spellingShingle |
Моделі та засоби паралельних і розподілених програм Моделі та засоби паралельних і розподілених програм Дорошенко, А.Е. Жереб, К.А. Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил |
| description |
Использование графических ускорителей позволяет достичь высокой производительности, однако требует от разработчика низкоуровневого программирования и знания деталей аппаратной и программной платформы. В работе предложен подход к автоматизации разработки приложений для графических ускорителей, основанный на использовании парадигмы переписывающих правил. |
| format |
Article |
| author |
Дорошенко, А.Е. Жереб, К.А. |
| author_facet |
Дорошенко, А.Е. Жереб, К.А. |
| author_sort |
Дорошенко, А.Е. |
| title |
Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил |
| title_short |
Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил |
| title_full |
Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил |
| title_fullStr |
Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил |
| title_full_unstemmed |
Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил |
| title_sort |
разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил |
| publisher |
Інститут програмних систем НАН України |
| publishDate |
2009 |
| topic_facet |
Моделі та засоби паралельних і розподілених програм |
| url |
https://nasplib.isofts.kiev.ua/handle/123456789/4419 |
| citation_txt |
Разработка високопараллельних приложений для графических ускорителей с использованием переписывающих правил / А.Е. Дорошенко, К.А. Жереб // Пробл. програмув. — 2009. — № 3. — С. 3-18 — Бібліогр.: 22 назв. — рос. |
| work_keys_str_mv |
AT dorošenkoae razrabotkavisokoparallelʹnihpriloženijdlâgrafičeskihuskoritelejsispolʹzovaniemperepisyvaûŝihpravil AT žerebka razrabotkavisokoparallelʹnihpriloženijdlâgrafičeskihuskoritelejsispolʹzovaniemperepisyvaûŝihpravil AT dorošenkoae rozrobkavisokoparalelʹnihzastosuvanʹdlâgrafíčnihpriskorûvačívzvikoristannâmperepisuvalʹnihpravil AT žerebka rozrobkavisokoparalelʹnihzastosuvanʹdlâgrafíčnihpriskorûvačívzvikoristannâmperepisuvalʹnihpravil AT dorošenkoae developmentofhighlyparallelapplicationsforgraphicalprocessingunitsusingrewritingrules AT žerebka developmentofhighlyparallelapplicationsforgraphicalprocessingunitsusingrewritingrules |
| first_indexed |
2025-11-27T00:05:00Z |
| last_indexed |
2025-11-27T00:05:00Z |
| _version_ |
1849899776635568128 |
| fulltext |
Моделі та засоби паралельних і розподілених програм
3
УДК 681.3
А.Е. Дорошенко, К.А. Жереб
РАЗРАБОТКА ВЫСОКОПАРАЛЛЕЛЬНЫХ ПРИЛОЖЕНИЙ
ДЛЯ ГРАФИЧЕСКИХ УСКОРИТЕЛЕЙ С
ИСПОЛЬЗОВАНИЕМ ПЕРЕПИСЫВАЮЩИХ ПРАВИЛ
Использование графических ускорителей позволяет достичь высокой производительности, однако тре-
бует от разработчика низкоуровневого программирования и знания деталей аппаратной и программной
платформы. В работе предложен подход к автоматизации разработки приложений для графических ус-
корителей, основанный на использовании парадигмы переписывающих правил.
Введение
В настоящее время параллельное
программирование становится все более
актуальным при разработке программ для
стандартных, широкодоступных компью-
теров, а не только для специализирован-
ных высокопроизводительных систем.
Этому способствует бурное развитие мно-
гоядерных процессоров, которые сейчас
устанавливаются в большинстве настоль-
ных и мобильных компьютеров [1]. Но
есть еще одно направление параллельного
программирования, которое получило осо-
бенное развитие в последнее время. Речь
идет о программировании общецелевых
задач для графических ускорителей [2].
Рыночные требования привели к
бурному развитию графических ускорите-
лей, в результате чего их вычислительная
мощность на данный момент значительно
превышает возможности обычных процес-
соров. Поэтому возник интерес к исполь-
зованию графических ускорителей для ре-
шения задач, не связанных напрямую с об-
работкой графики [3–4]. Первоначально
для этого использовались средства про-
граммирования графических задач, не при-
способленные для произвольных вычисле-
ний. В последнее время интерес к исполь-
зованию графических ускорителей в каче-
стве средств высокопроизводительных вы-
числений поддерживается усилиями веду-
щих разработчиков аппаратуры. Так, ком-
пания NVidia представляет платформу для
вычислений на графическом ускорителе
CUDA [5]. Аналогично, компания AMD
выступила с инициативой Stream [6].
Такие платформы облегчают реали-
зацию различных задач на графических
ускорителях, поскольку предоставляют
модель программирования, которая более
приспособлена к разработке произвольных
вычислений, по сравнению со средствами
программирования графики. Тем не менее,
разработка приложений для графических
ускорителей остается достаточно сложной
задачей. Разработчик должен быть знаком
с устройством графического ускорителя,
понимать особенности работы его компо-
нент. Модель программирования, предос-
тавляемая выбранной платформой, являет-
ся новой по сравнению как с последова-
тельным программированием, так и с мно-
гопоточным программированием для мно-
гоядерных процессоров. Многие задачи
требуют достаточно низкоуровневой реа-
лизации и оптимизации. Все это приводит
к необходимости автоматизации процесса
разработки.
В данной работе предложен подход
к автоматизации разработки приложений
для графических ускорителей, на примере
платформы CUDA. Для этого используется
система переписывающих правил Term-
ware [7–9]. Предложенный подход позво-
ляет как переходить от последовательных
программ к параллельным, выполняющим-
ся на графическом ускорителе, так и опти-
мизировать программы для CUDA. При
этом возможно распараллеливание про-
грамм, изначально написанных на языке
© А.Е. Дорошенко, К.А. Жереб, 2009
ISSN 1727-4907. Проблеми програмування. 2009. № 3
Моделі та засоби паралельних і розподілених програм
4
C# для платформы Microsoft .NET, с со-
хранением возможности использования
всех средств платформы .NET.
Данная работа продолжает исследо-
вания, начатые в работах [10–11], по авто-
матизации процесса проектирования и раз-
работки эффективных параллельных про-
грамм. Особенность данной работы – ис-
пользование новой аппаратной платформы
для параллельных вычислений, графиче-
ских ускорителей, что позволяет добиться
значительного повышения производитель-
ности по сравнению с использованием
многоядерных процессоров общего
назначения.
Область исследований, связанная с
автоматизацией разработки приложений
для графических процессоров, в настоящее
время активно развивается. При этом рас-
сматриваются как задачи перехода от по-
следовательных к параллельным програм-
мам, так и задачи оптимизации сущест-
вующих параллельных программ с исполь-
зованием возможностей графических ус-
корителей. Так, в работе [12] рассмотрен
автоматический переход от многопоточной
программы, реализованной с использова-
нием технологии OpenMP [13], к реализа-
ции данной программы на платформе
CUDA. Работа [14] описывает платформу
для оптимизации циклов в программах для
графических ускорителей. Разработаны
системы для автоматического распаралле-
ливания и оптимизации программ из кон-
кретной предметной области, например,
data mining [15] или обработка изображе-
ний [16]. Авторы работы [17] описывают
библиотеку высокоуровневых структур
данных для графических ускорителей.
Также разрабатываются платформы про-
граммирования графических ускорителей,
предоставляющие более высокоуровневые
средства по сравнению с CUDA, такие как
hiCUDA [18] и BSGP [19].
В отличие от существующих работ
по данной тематике, в данной работе рас-
смотрена автоматизация перехода на
платформу CUDA с высокоуровневого
языка C#. Это позволяет использовать
возможности платформы Microsoft .NET,
которая в настоящее время широко ис-
пользуется для разработки приложений в
различных областях. Кроме того, исполь-
зование переписывающих правил для опи-
сания распараллеливающих и оптимизи-
рующих преобразований позволяет легко
добавлять новые преобразования.
Материал данной работы организо-
ван следующим образом. В разделе 1 опи-
саны особенности платформы CUDA,
влияющие на разработку приложений для
этой платформы. В разделе 2 описан авто-
матизированный процесс распараллелива-
ния программы с целью исполнения на
графическом ускорителе. Раздел 3 содер-
жит некоторые примеры оптимизирующих
преобразований, автоматизированные с
помощью платформы Termware. Работу
завершают выводы и направления даль-
нейшей работы.
1. Особенности платформы
CUDA
1.1. Аппаратные составляющие.
Современные графические ускорители
поддерживают высокую степень паралле-
лизма, специально приспособленную для
выполнения графических задач. Использо-
вание возможностей графических ускори-
телей для общецелевых вычислений тре-
бует от разработчика понимания особен-
ностей аппаратной платформы и модели
программирования CUDA [5].
Общая схема устройства графиче-
ского ускорителя показана на рис. 1. Гра-
фическое устройство (device) содержит не-
сколько мультипроцессоров, а также об-
щую для них графическую память. Каж-
дый мультипроцессор содержит несколько
вычислительных ядер (скалярных процес-
соров), а также один управляющий блок,
поддерживающий многопоточное испол-
нение. В результате количество вычисли-
тельных ядер (а значит, и степень возмож-
ного параллелизма) оказывается сущест-
венно выше, чем у общецелевых много-
ядерных процессоров.
Графические ускорители поддержи-
вают несколько разных видов памяти, от-
личающихся по объему, скорости доступа
и особенностям реализации. Самая быст-
рая память – регистры вычислительных
ядер; однако их количество в каждом ядре
ограничено.
Моделі та засоби паралельних і розподілених програм
5
Рис. 1. Основные аппаратные компоненты графического ускорителя
Кеш данных, или разделяемая па-
мять (shared memory) поддерживает произ-
вольный доступ из любого вычислитель-
ного блока, однако имеет ограниченный
размер и требует явной синхронизации
доступа. Есть также два специфических
вида памяти – кеш констант и кеш текстур.
Они поддерживают только чтение, но
имеют больший объем по сравнению с
shared памятью. Разница между текстур-
ной и константной памятью заключается в
том, что текстурная память поддерживает
специальные режимы доступа, полезные
для графических задач. Всем мультипро-
цессорам доступна также графическая па-
мять устройства. Она является самой
большой по объему и самой медленной
памятью, доступной графическому уско-
рителю.
1.2. Структура параллельной
программы. Параллельная программа для
графических ускорителей состоит из
большого количества потоков. Особенно-
сти аппаратного обеспечение, а именно
большое количество вычислительных
ядер, позволяют использовать очень мел-
козернистый параллелизм, вплоть до вы-
деления отдельного потока для каждого
элемента данных.
Для исполнения на графическом ус-
корителе потоки объединяются в блоки.
Каждый блок выполняется на одном муль-
типроцессоре. Различные блоки, соответ-
ствующие одной или нескольким про-
граммам, по возможности распределяются
равномерно по доступным мультипро-
цессорам.
При исполнении на мультипроцес-
соре, потоки объединяются в так называе-
мые warp’ы по 32 потока. На каждом шаге
многопоточный планировщик выбирает
один из доступных warp’ов. Далее очеред-
ная инструкция для потоков этого warp’а
выполняется одновременно на вычисли-
тельных ядрах. Следует отметить, что на-
личие ветвлений в потоках понижает про-
изводительность, поскольку каждый из ва-
риантов исполнения в таких случаях вы-
полняется последовательно, а остальные
потоки warp’а ожидают выполнения.
1.3. Программная модель CUDA.
Программа для графического ускорителя,
Моделі та засоби паралельних і розподілених програм
6
написанная на CUDA, разделяется на две
части: код для исполнения на графическом
устройстве (device code) и код для испол-
нения на обычном процессоре (host code).
Код для исполнения на устройстве пред-
ставлен в виде ядер (kernel). Ядра могут
быть написаны как на специальном низко-
уровневом языке (PTX), так и на расшире-
нии языка C (C for CUDA). В последнем
случае ядро имеет вид функции языка C,
описывающей поведение одного потока.
Платформа CUDA предлагает два
разных способа вызова ядер. Более про-
стой способ, C for CUDA, является расши-
рением языка C. Вызов ядра при этом по-
хож на вызов функции C, но с передачей
дополнительных параметров. Эти пара-
метры определяют размер блока, а также
количество одновременно исполняемых
блоков в сетке (grid). Каждый поток имеет
доступ к своему номеру в блоке и номеру
блока в сетке, что позволяет различным
потокам работать с разными данными.
Более сложным способом вызова
ядер является driver API. В этом случае яд-
ра необходимо предварительно скомпили-
ровать в бинарный формат, после чего ис-
пользуются функции для загрузки и вы-
полнения ядер. Driver API является более
сложным в разработке и может привести к
появлению ошибок, так как требует ручно-
го выполнения многих действий, которые
автоматизированы в C for CUDA. Пример
таких действий – передача параметров яд-
ра. С другой стороны, driver API дает раз-
работчику больший контроль над выпол-
нением программы.
Оба способа вызова CUDA предос-
тавляют функции для работы с графиче-
ской памятью. Обычно работа программы
состоит из таких шагов: инициализация
структур данных в графической памяти,
копирование из обычной памяти в графи-
ческую входных данных, вызов ядра (или
нескольких ядер, или многократный вызов
одного ядра), копирование результатов из
графической памяти, освобождение гра-
фической памяти.
1.4. Программирование графиче-
ского ускорителя из .NET. Платформа
CUDA на данный момент поддерживает
только язык C, поэтому для использования
возможностей графического ускорителя из
.NET программ необходимы дополнитель-
ные компоненты. К подобным компонен-
там относится библиотека CUDA .NET
[20]. Эта библиотека позволяет .NET коду
выполнять ядра CUDA, используя воз-
можности, близкие к driver API.
Таким образом, для использования
CUDA .NET необходимо сначала реализо-
вать ядро CUDA, что делается средствами
C for CUDA. После этого необходимо на-
писать C# код для вызова ядра. Так как
CUDA .NET использует более сложный
driver API, от разработчика требуется кор-
ректная реализация многих низкоуровне-
вых деталей. В частности, при передаче
параметров необходимо иметь представле-
ние об их размещении в памяти и коррект-
но передать размеры и относительные
смещения параметров. Необходимость вы-
полнения этих, довольно сложных, опера-
ций вручную затрудняет разработку и по-
вышает вероятность ошибки.
Кроме CUDA .NET, для разработки
приложений для .NET можно использовать
также проект Accelerator от Microsoft Re-
search [21]. Этот проект предоставляет на-
бор функций, которые вызываются из C# и
выполняются на графическом ускорителе.
Он не требует низкоуровневой разработки
на языках, не принадлежащих к семейству
.NET. Однако его возможности сущест-
венно ограничены, так как любой алгоритм
необходимо представить в виде комбина-
ции доступных функций. Это затрудняет
использование Accelerator во многих прак-
тических классах задач. Кроме того, про-
ект более не развивается и поэтому не
поддерживает новейшие достижения раз-
работчиков графических ускорителей.
2. Переход к вычислениям на
графическом ускорителе
2.1. Общая схема преобразования.
В данной работе используется метод авто-
матизированного распараллеливания кода
для выполнения на графическом процессо-
ре, основанный на использовании перепи-
сывающих правил. В качестве входных
данных для такого преобразования исполь-
зуется последовательный код на языке C#.
Этот код преобразуется в представление в
Моделі та засоби паралельних і розподілених програм
7
виде термов с использованием анализатора
(парсера) языка C#. В качестве входных
данных может использоваться также фор-
мальное представление алгоритма, напри-
мер, из системы ИПС [10], однако этот
случай в работе не рассматривается.
В исходной (последовательной)
программе выделяются фрагменты, в ко-
торых происходят интенсивные вычисле-
ния. Такие фрагменты целесообразно ис-
полнять на графическом процессоре, что
позволит существенно повысить произво-
дительность программы в целом. При этом
обычно такой код составляет лишь незна-
чительную часть всего кода программы.
После выделения перспективных
участков кода к ним применяются разра-
ботанные правила Termware, которые пе-
реводят последовательный код в парал-
лельный код для исполнения на графиче-
ском ускорителе. Преобразованный код
состоит из двух частей. Код, описываю-
щий алгоритм, переводится в представле-
ние, соответствующее языку С for CUDA.
Этот код представлен в виде ядра CUDA,
сохраняется в файл с расширением .cu и
компилируется в бинарный формат (.cubin)
с использованием компилятора NVCC от
NVidia. После генерации кода для CUDA
возможно также применение оптимизи-
рующих преобразований.
Кроме алгоритмического кода, ге-
нерируется служебный код на C#, который
вызывает сгенерированное ядро CUDA с
использованием средств CUDA .NET. Ре-
зультатом преобразования является про-
грамма, преобладающая часть кода кото-
рой по-прежнему реализована на C# и мо-
жет использовать все возможности плат-
формы .NET. При этом критические по
производительности участки программы
исполняются на графическом ускорителе с
повышенной производительностью.
2.2. Система Termware. Для осу-
ществления всех преобразований исполь-
зуется платформа переписывающих пра-
вил Termware, подробно описанная в [7–
11]. Termware предназначена для описания
преобразования над термами, т. е. выраже-
ниями вида f(t1, …, tn). Для задания пре-
образований используются правила Term-
ware, т. е. конструкции вида source
[condition] -> destination [action].
Здесь source – исходный терм
(образец для поиска), condition – усло-
вие применения правила, destination –
преобразованный терм, action – допол-
нительное действие при срабатывании
правила. Каждый из 4 компонентов прави-
ла может содержать переменные (которые
записываются в виде $var), что обеспечи-
вает общность правил. Компоненты con-
dition и action являются необязатель-
ными. Они могут исполнять произвольный
процедурный код, в частности использо-
вать дополнительные данные о программе.
Применение правила происходит
следующим образом: сначала находится
подтерм входного терма (дерева програм-
мы), который подходит под source. Да-
лее проверяется условие применения (если
оно присутствует). Если условие выполня-
ется, происходит замена source на des-
tination. При этом переменные в desti-
nation заменяются соответствующими зна-
чениями из source. Также выполняется
действие action (если оно при-
сутствовало).
Каждое преобразование задается
системой правил, т.е. набором правил, ко-
торые последовательно применяются к
данному терму (дереву программы). Поря-
док применения правил определяется стра-
тегией. В систему Termware встроены не-
сколько основных стратегий, таких как
TopDown, BottomUp, FirstTop. Кроме того,
возможно создание дополнительных
стратегий.
2.3. Пример распараллеливания.
В качестве примера использования опи-
санного подхода рассмотрим задачу распа-
раллеливания алгоритма битонической
сортировки (bitonic sort, [22]). Последова-
тельная реализация алгоритма имеет сле-
дующий вид:
for (k = 2; k <= ARR_SIZE; k <<= 1)
{
for (j = k >> 1; j > 0; j = j >> 1)
{
for (int i = 0; i < ARR_SIZE;
i++)
{
int ixj = i ^ j;
if ((ixj) > i)
{
if ((i & k) == 0 &&
Моделі та засоби паралельних і розподілених програм
8
toSort[i] > toSort[ixj])
{
swap(toSort,i,ixj);
}
if ((i & k) != 0 &&
toSort[i] < toSort[ixj])
{
swap(toSort,i,ixj);
}
}
}
}
}
Для обработки этого кода средства-
ми Termware он сначала преобразуется в
терм с помощью анализатора языка C#.
При этом создается следующий терм:
For(Assignment(k,2),k<=ARR_SIZE,
Assignment(k,k<<1),
[For(Assignment(j,k>>1),
j>0,Assignment(j,j>>1),
[For(Assignment(i,0),i<=ARR_SIZE,
Increment(i),
[DeclarationAssign-
ment(ixj,int,bit_xor(i,j)),
If(ixj>i,[
If(And(Equals(i & k,0),
ArrayElement(toSort,i) >
ArrayElement(toSort,ixj)),
[MethodCall(swap,toSort,i,ixj)]),
If(And(NotEquals(i & k,0),
ArrayElement(toSort,i) < ArrayE-
lement(toSort,ixj)), [MethodCall
(swap,toSort,i,ixj)])
])])])])
2.4. Выбор участка кода для рас-
параллеливания. Первым шагом перехо-
да от последовательной программы К па-
раллельной программе с использованием
графического ускорителя является выбор
критичных по производительности участ-
ков кода, которые будут распараллелены.
Во многих случаях такие участки опреде-
ляются циклической конструкцией (цикл
типа for или while) или вызовом метода.
Выбор кода для распараллеливания
может осуществляться по-разному в зави-
симости от специфики задачи. Возможен
выбор необходимого кода вручную, на ос-
новании знаний и опыта разработчика.
Возможно также использование внешнего
инструментария, такого как профилиров-
щик. В этих случаях информация о месте
применения распараллеливающих преоб-
разований задается разработчиком через
пользовательский интерфейс. Возможно
указание терма, к которому применяется
преобразование, или непосредственно уча-
стка кода (в этом случае система определя-
ет необходимый терм автоматически, ис-
пользуя информацию о связях между тер-
мами и элементами кода). Возможна также
автоматизированная оценка сложности
различных участков алгоритма, описанная
в работе [11].
Независимо от способа выбора кода
для распараллеливания, этот код помеча-
ется специальной меткой, для использова-
ния в дальнейших правилах. Так, для цик-
ла for это сводится к добавлению подтер-
ма-метки _MARK_CUDA к терму, обозна-
чающего циклическую конструкцию. Эта
метка указывает место применения даль-
нейших правил, а также задает имя, иден-
тифицирующее данный участок кода (на
случай, если преобразование можно при-
менить в нескольких местах). После до-
бавления метки терм приобретает вид
For($init, $condition, $step, $body,
_MARK_CUDA(bitonic_kernel))
В примере алгоритма битонической
сортировки кандидатом для распараллели-
вания является внутренний цикл (по i).
Внешние циклы имеют более сложную
структуру, что препятствует их распарал-
леливанию средствами CUDA.
2.5. Структура преобразования.
После добавления метки, выделяющей
участки кода для преобразования, стано-
вится возможным применение правил
Termware для перехода к параллельному
коду. Первое правило, которое применяет-
ся в данном случае, задает общую струк-
туру преобразования: перевод алгоритми-
ческого кода в ядро CUDA, а также гене-
рацию служебного кода для вызова этого
ядра средствами CUDA .NET. Это правило
имеет следующий вид:
For(Assignment($idx,0),$idx<$size,
Increment($idx),
$body,_MARK_CUDA($name)) ->
[_InitCuda,_InitCudaKernel($name),
_CallCudaKernel($name,$size,$params)]
[GetExtraParams($body,$params,$idx);
GenerateKernel1D($name, $params,
$body, $idx)]
Оно срабатывает для циклов типа
for, причем только специального вида –
фактически цикл со счетчиком. (Внешние
циклы в алгоритме битонической сорти-
ровки не имеют такого вида, поэтому к
ним данное преобразование неприменимо).
Моделі та засоби паралельних і розподілених програм
9
Кроме того, дополнительное условие сра-
батывания правила – наличие метки
_MARK_CUDA.
Срабатывание правила приводит к
замене цикла на 3 терма: _InitCuda – ини-
циализация платформы CUDA;
_InitCudaKernel – инициализация данно-
го ядра; и _CallCudaKernel – собственно
вызов ядра. Эти термы являются заготов-
ками для служебного кода, вызывающего
ядро CUDA. Кроме того, используются
средства Termware для вызова двух проце-
дурных методов. Метод GetExtraParams
вычисляет список параметров ядра, т.е. пе-
ременных, которые используются в теле
цикла, но не определяются в нем. (Этот
метод можно было бы реализовать систе-
мой правил, однако в данном случае про-
цедурная реализация оказывается более
простой и наглядной). Метод Generate-
Kernel1D создает из тела цикла ядро
CUDA; этот метод использует отдельную
систему правил для перехода от кода на C#
к коду на C for CUDA.
После срабатывания данного пра-
вила происходит промежуточный этап –
перемещение созданных термов-заготовок
для служебного кода. Терм
_CallCudaKernel никуда не перемещает-
ся, поскольку он должен находиться в том
же месте в коде, где был исходный цикл.
Однако термы _InitCuda и
_InitCudaKernel должны использоваться
однократно, в начале программы. Поэтому
используется система правил, которая пе-
реводит эти термы в начало программы
(т.е. в начало метода Main или специаль-
ный метод, отведенный для инициализа-
ции). Кроме того, эта система устраняет
повторения: если распараллеливание про-
исходит в нескольких местах кода, дубли-
каты термов _InitCuda и _InitCudaKernel
устраняются. В результате остается один
терм _InitCuda, за которым следуют тер-
мы _InitCudaKernel (по одному для каж-
дого уникального ядра).
2.6. Преобразование алгоритми-
ческого кода в код для CUDA. Как уже
упоминалось, для преобразования алго-
ритмического кода в ядро CUDA исполь-
зуется метод GenerateKernel1D
($name,$params,$body,$idx). Этот метод
реализован на C# и вызывается средствами
Termware как действие (action) при сраба-
тывании правила. Метод сводится к созда-
нию терма, соответствующего новому
файлу и вызову системы правил для созда-
ния содержимого этого файла.
Имя метода содержит суффикс 1D,
что указывает на использование одномер-
ных структур данных. Параметрами мето-
да являются: $name – имя ядра; $params –
список параметров ядра; $body – исходное
тело цикла, которое преобразуется в код
ядра и $idx – исходный индекс цикла.
В ходе работы метода создается но-
вый исходный файл в проекте, с именем,
совпадающим с именем ядра и расширени-
ем .cu (исходный файл C for CUDA). В
этот файл добавляются необходимые для
CUDA файлы заголовков, а затем – код яд-
ра, который получается преобразованием
исходного терма, _CudaKernel1D($name,
$params, $body, $idx) с использованием
разработанной системы правил. В самом
общем виде эта система имеет следующий
вид:
1. _CudaKernel1D($name, $params,
$body, $idx) ->
_CudaFunctionKernel($name,
$params,
[_PrepareIdx($idx,x),_СsToCuda(
$body)])
2. _CudaFunctionKernel($name,
$params, $body) ->
Function([ExternC,__global__],
void,$name,$params,$body)
3. _PrepareIdx($idx,$coor) -> De-
clarationAssignment(
$idx,int,_CudaIdx($coor))
4. _CudaIdx($coor) ->
Dot(blockIdx,$coor) *
Dot(blockDim,$coor) +
Dot(threadIdx,$coor)
Первое правило задает общую
структуру ядра – специальное описание
функции, с телом, состоящим из определе-
ния индекса, за которым следует тело ис-
ходного цикла. Правило 2 реализует опи-
сание функции, которая является ядром
CUDA. Эта функция имеет модификатор
__global__ , не возвращает значения
(void), а ее имя и список параметров опре-
деляется данными, переданными в метод
GenerateKernel1D. Кроме того, добавляет-
ся модификатор extern "C" , необходи-
мый для вызова ядра из CUDA .NET. (Вы-
Моделі та засоби паралельних і розподілених програм
10
зов осуществляется по имени ядра, и ис-
пользование такого модификатора требует
от компилятора сохранить имя в простом
виде, без добавления информации, харак-
терной для C++). Правила 3 и 4 определя-
ют индекс исходного массива из парамет-
ров запуска ядра (индекс блока в сетке и
потока в блоке). Поскольку в данном слу-
чае ядро работает с одномерным массивом,
единственный индекс определяется соот-
ношением
i = blockIdx.x * blockDim.x +
threadIdx.x;
Вышеприведенная общая система
правил дополняется специфическими пра-
вилами, которые переводят код из C# в C
for CUDA. Эти правила работают с термом
_СsToCuda( $body). Такие правила могут
быть двух видов: технические, которые
переводят определенные конструкции C# в
соответствующие конструкции C, и опти-
мизационные, которые вносят изменения в
структуру кода с целью повышения произ-
водительности. Правила первого типа дос-
таточно простые и одинаковы для различ-
ных программ. Особый интерес представ-
ляют именно оптимизационные преобра-
зования, которые являются специфичными
для каждой задачи. Следует отметить, что
оптимизационные преобразования имеет
смысл выделить в отдельные системы пра-
вил, которые применяются отдельно от
процесса распараллеливания. Некоторые
примеры оптимизационных преобразова-
ний рассмотрены в разделе 3.
Для алгоритма битонической сор-
тировки, в силу его простоты, практически
не требуются специальные преобразования
(как технические, так и оптимизационные).
Можно упомянуть лишь одно необходимое
преобразование, связанное с вызовом ме-
тода swap в исходном коде. В данном слу-
чае этот метод преобразуется в функцию с
модификатором __device__ , которая так-
же выполняется на графическом ускорите-
ле (на самом деле функции такого рода
встраиваются в код ядра на этапе
компиляции).
В результате работы метода Gener-
ateKernel1D получается следующий код
ядра:
extern "C" __global__ void
bitonic_kernel(int * toSort, int j,
int k)
{
int i = threadIdx.x + blockDim.x
* blockIdx.x;
int ixj = i ^ j;
if ((ixj) > i)
{
if ((i & k) == 0 && toSort[i]
> toSort[ixj])
{
swap(toSort,i,ixj);
}
if ((i & k) != 0 && toSort[i]
< toSort[ixj])
{
swap(toSort,i,ixj);
}
}
}
Далее этот код компилируется в би-
нарный формат (.cubin) с использованием
компилятора NVCC. В таком виде он ста-
новится доступным для использования
средствами CUDA .NET.
2.7. Генерация служебного кода.
Служебный код для вызова полученного
ядра генерируется из трех термов
_InitCuda, _InitCudaKernel и
_CallCudaKernel. Первые два терма, от-
вечающие за инициализацию платформы и
данного ядра, генерируют достаточно про-
стой и стандартный код:
// Init CUDA, select 1st device.
CUDA cuda = new CUDA(0, true);
// load module, select kernel
string path = Path.Combine (
Environment.CurrentDirectory,
"bitonic_kernel.cubin");
cuda.LoadModule(path);
CUfunction bitonic_kernel =
cuda.GetModuleFunction
("bitonic_kernel");
Этот код инициализирует платфор-
му CUDA, выбирает для вычислений пер-
вое из доступных графических устройств.
Затем загружается сгенерированный би-
нарный модуль, и из него выбирается (по
имени) ядро.
Правила для вызова ядра являются
более сложными. Сам вызов ядра состоит
из нескольких шагов:
1) копирование входных данных в
графическую память;
2) задание параметров ядра;
3) собственно вызов ядра;
Моделі та засоби паралельних і розподілених програм
11
4) копирование результатов работы
из графической памяти;
5) освобождение структур данных
CUDA.
При этом различные шаги выпол-
няются в разных местах программы и раз-
личное количество раз. Шаги 2 и 3 выпол-
няются для каждого вызова: в случае зада-
чи битонической сортировки они оказы-
ваются внутри двух циклов (по k и по j) и
поэтому выполняются многократно. Шаг 1
выполняется после того, как входные дан-
ные проинициализированы в C# коде, и до
вызова ядра. При этом рассматриваются
только массивы данных (toSort в алго-
ритме сортировки), а не единичные пере-
менные (j и k). Аналогично, шаг 4 выпол-
няется после того как произведены вы-
числения на графическом ускорителе, до
того как их результаты используются в
коде на C#.
В общем случае, возможно много-
кратное выполнение шагов 1 и 4 – либо
для разных массивов, либо для одного
массива, если он используется в коде меж-
ду вызовами ядра (или различных ядер).
Для задачи битонической сортировки реа-
лизуется более типичный сценарий, когда
массив копируется в графическую память
и из нее однократно.
Шаг 5 – освобождение структур
данных CUDA – происходит после по-
следнего использования этих структур в
шагах 1–4. В примере с сортировкой это
происходит сразу после копирования дан-
ных из графической памяти (шаг 4).
Таким образом, система правил для
генерации вызова ядра устроена следую-
щим образом. Сначала исходный терм
_CallCudaKernel($name,$size,$params)
разбивается на 5 термов:
_CopyHostToDevice($params),
_SetParameters($name,$params),
_ExecuteKernel($name,$size),
_CopyDeviceToHost ($params) и
_FreeCuda($params). Далее терм
_CopyHostToDevice поднимается «вверх»,
т.е. он переставляется местами с термами,
в которых не происходит запись в массив
данных. В результате он оказывается сразу
после записи входных данных в массив.
Аналогично, термы CopyDeviceToHost и
_FreeCuda перемещаются вниз.
Далее каждый из термов расшифро-
вывается. Для _CopyHostToDevice получа-
ется достаточно простой код, поскольку
CUDA .NET предоставляет удобные мето-
ды для работы с массивами. Полученный
код имеет вид
CUdeviceptr dev_toSort =
cuda.CopyHostToDevice<int>(toSort);
Аналогично термы CopyDeviceTo-
Host и _FreeCuda преобразуются в сле-
дующий код:
cuda.CopyDeviceToHost<int>(
dev_toSort, toSort);
cuda.Free(dev_toSort);
Терм _SetParameters преобразует-
ся в код, который задает значение каждого
параметра и его размер в памяти. Размеры
параметров вычисляются автоматически,
что позволяет избежать ошибок при руч-
ном определении размеров. Сгенерирован-
ный код имеет вид:
cuda.SetParameter(bitonic_kernel, 0,
(uint) dvalues.Pointer);
cuda.SetParameter(bitonic_kernel,
IntPtr.Size, (uint)j);
cuda.SetParameter(bitonic_kernel,
IntPtr.Size + 4, (uint)k);
cuda.SetParameterSize(bitonic_kernel,
(uint) (IntPtr.Size + 8));
В результате сам вызов ядра (терм
_ExecuteKernel) генерирует достаточно
простой код:
cuda.SetFunctionBlockShape
(bitonic_kernel, MAX_THREAD, 1, 1);
cuda.Launch(bitonic_kernel, ARR_SIZE
/ MAX_THREAD, 1);
В этом коде появляется дополни-
тельный параметр MAX_THREAD, опреде-
ляющий количество потоков в блоке.
Таким образом, использование сис-
тем переписывающих правил позволяет
полностью автоматизировать процесс соз-
дания служебного кода для вызова ядер
CUDA средствами CUDA .NET. Тем са-
мым разработчик может концентрировать
внимание на особенностях алгоритма, а не
на технических средствах его реализации.
Кроме того, исключаются ошибки, воз-
можные при создании этого служебного
кода вручную.
2.8. Сравнение производительно-
сти. Для оценки эффекта от применения
распараллеливающих преобразований бы-
ли проведены измерения скорости работы
алгоритма битонической сортировки в
Моделі та засоби паралельних і розподілених програм
12
двух версиях: исходная реализация (CPU)
и распараллеленная с использованием
CUDA (GPU). Результаты измерений при-
ведены в табл. 1.
Таблица 1. Сравнение производи-
тельности
Размер GPU, c CPU, c Ускорение
1048576 0.12 3.3 27.5
2097152 0.26 8.98 34.53846
4194304 0.5 16 32
8388608 0.98 32.34 33
16777216 2.42 75 30.99174
Из результатов видно, что распа-
раллеливание дает существенный эффект:
производительность повышается в 30 раз.
Отметим, что этот результат достигается
при автоматизированном преобразовании,
и для его реализации от разработчика не
требуется знаний платформы CUDA или
CUDA .NET.
3. Оптимизация кода для
CUDA
3.1. Особенности оптимизации
CUDA. Для программ, выполняющихся на
графическом ускорителе и написанных на
CUDA, оптимизация имеет особое значе-
ние. Достаточно типичной является ситуа-
ция, когда производительность неоптими-
зированной программы в десятки раз от-
личается от производительности ее опти-
мизированного аналога. Так как именно
повышение производительности является
причиной перехода на вычисления на гра-
фическом ускорителе, то оптимизация –
важная часть жизненного цикла программ
на CUDA.
При этом оптимизация приложений,
написанных на CUDA – достаточно слож-
ный процесс и требует понимания архи-
тектуры CUDA, взаимосвязи различных
аппаратных компонент графического ус-
корителя. Сам алгоритм должен поддер-
живать очень широкий параллелизм, что-
бы задействовать возможности современ-
ных графических процессоров. Кроме то-
го, его реализация должна учитывать мно-
гие мелкие детали, такие как размещение
данных в памяти, порядок доступа к ним,
использование оптимальных типов данных
и операций и т. д.
Учитывая все вышеперечисленные
особенности, особую важность приобрета-
ет задача автоматизации процесса оптими-
зации программ для CUDA. В данном раз-
деле описаны примеры автоматизирован-
ных преобразований в виде правил Term-
ware, направленных на повышение произ-
водительности приложения.
3.2. Пример: иерархическое па-
раллельное суммирование. Рассмотрим
оптимизирующие преобразования для за-
дачи нахождения суммы большого массива
данных. Используется иерархический ал-
горитм суммирования: на каждом этапе
элементы массива разбиваются на пары и
суммируются независимо друг от друга.
Это позволяет получить высокую степень
параллелизма.
Простейшая реализация данного ал-
горитма на CUDA имеет следующий вид:
extern "C" __global__ void
parallelAdd ( int * inData, int *
outData )
{
int tid = threadIdx.x;
int i = blockIdx.x * blockDim.x
+ threadIdx.x;
outData[i]=inData[i];
__syncthreads ();
for ( int s = 1; s <
blockDim.x; s *= 2 ) {
if ( tid % (2*s) == 0 )
outData [i] += outData [i + s];
__syncthreads ();
}
}
3.3. Использование разделяемой
(shared) памяти. Часто производитель-
ность программ для CUDA зависит в ос-
новном от правильной реализации доступа
к памяти, а не от эффективности самих
вычислений. В частности, в примере с па-
раллельным суммированием все операции
производятся с глобальной памятью (кото-
рая реализована на видеопамяти). Ее пре-
имущества заключаются в том, что она
доступна всем процессорам, а также с ней
можно работать непосредственно из обыч-
ного кода. Однако скорость доступа к ней
существенно меньше, чем к более быст-
рым видам памяти.
Более быстрой альтернативой гло-
бальной памяти является так называемая
разделяемая (shared) память. Она располо-
жена на каждом процессоре, поэтому ско-
Моделі та засоби паралельних і розподілених програм
13
рость доступа к ней максимальна. Ее не-
достатками являются малый размер (по-
рядка 16К в современных ускорителях), а
также то, что доступ к ней возможен лишь
в пределах одного блока, при этом требует
синхронизации. Тем не менее, если алго-
ритм позволяет использовать небольшое
локальное хранилище данных, целесооб-
разно использовать shared память. (Алго-
ритм битонической сортировки, рассмот-
ренный в разделе 2, такими свойствами не
обладает).
Обычно работа с shared памятью
выполняется следующим образом: данные
из обрабатываемого участка глобальной
памяти копируются в shared память (этот
процесс также является параллельным –
каждый поток копирует свою часть дан-
ных). После этого выполняется синхрони-
зация потоков командой __syncthreads
(). Далее вся работа идет уже в shared па-
мяти. Когда необходимо записать резуль-
таты в глобальную память, также прово-
дится синхронизация и результаты вычис-
лений по частям копируются из shared в
глобальную память.
Для перехода к использованию
shared памяти в задаче параллельного сум-
мирования можно использовать следую-
щую систему правил:
1. Function($x, $y,
parallelAdd,$z,[$b1:$b2]) ->
Function($x, $y,
parallelAdd1,
$z,[SharedDecl:$b1:CopyBack($b2
)])
2. SharedDecl -> ArrayDeclaration
(localData,int, BLOCK_SIZE,
__shared__)
3. Assignment (ArrayElement
(outData,i), ArrayEle-
ment(inData,i)) -> Assignment
(ArrayElement (localData
,tid),ArrayElement(inData,i))
4. PlusAssignment(ArrayElement
(out-
Data,i),ArrayElement(outData,i+
s)) -> PlusAssignment
(ArrayElement (localData,
tid),ArrayElement (localData,
tid+s))
5. CopyBack([$x:$y]) ->
[$x:CopyBack($y)]
6. CopyBack(NIL) -> If
(Equals(tid,0),Assignment(
ArrayElement(outData,Dot
(blockIdx,x)),ArrayElement
(localData,0)))
Правила 1 и 2 добавляют объявле-
ние локального массива shared памяти.
Правило 3 заменяет копирование напря-
мую в выходной массив копированием в
локальный массив. Правило 4 аналогично
заменяет суммирование в глобальной па-
мяти на суммирование в shared памяти.
Правила 5 и 6 добавляют в конце ядра ко-
пирование в выходной массив результата
вычисления.
В результате действия системы пра-
вил получается следующее модифициро-
ванное ядро:
extern "C" __global__ void
parallelAdd1 ( int * inData, int *
outData )
{
__shared__ int localData
[BLOCK_SIZE];
int tid = threadIdx.x;
int i = blockIdx.x * blockDim.x
+ threadIdx.x;
localData [tid] = inData [i];
__syncthreads ();
for ( int s = 1; s <blockDim.x;
s *= 2 ) {
if ( tid % (2*s) == 0 )
localData [tid] += localData
[tid + s];
__syncthreads ();
}
if ( tid == 0 )
outData[blockIdx.x] = localData
[0];
}
3.4. Уменьшение количества
ветвлений. Еще одним недостатком ис-
ходной реализации алгоритма параллель-
ного суммирования является большое ко-
личество ветвлений. В коде используется
условие цикла:
if ( tid % (2*s) == 0 )
Это условие часто меняется для со-
седних потоков (в частности, для s=1 все
соседние потоки идут по разным ветвям).
Поскольку в архитектуре графического ус-
корителя предусмотрено большее количе-
ство вычислительных блоков по сравне-
нию с управляющими, такое ветвление
вынуждает запускать потоки последова-
тельно, а не одновременно.
Возможно уменьшение степени
ветвления, за счет перераспределения вы-
числений между потоками. В исходной
реализации, вычисления выполнялись на
потоках с номерами, кратными степеням
Моделі та засоби паралельних і розподілених програм
14
двойки (на первом шаге – четные номера,
на втором – кратные 4, и т.д.). Рассмотрим
преобразование, которое переносит вы-
числения на последовательные потоки (на
первом шаге – первая половина потоков,
на втором – первая четверть и т.д.). Преоб-
разование состоит из единственного
правила:
If(Equals(tid % (2*s),0), PlusAs-
signment( ArrayElement (localData,
tid) ,ArrayElement(localData,tid+s)))
-> [DeclarationAssignment(index,int,2
* s * tid),If(index<Dot(blockDim,x),
PlusAssignment( ArrayElement( local-
Data,index), ArrayElement( localData,
index+s)))]
Это правило заменяет предыдущее
условия на такой код:
int index = 2 * s * tid;
if ( index < blockDim.x )
data [index] += data [index + s];
3.5. Конфликты доступа к shared
памяти. После применения предыдущего
правила количество ветвлений существен-
но понижается. Однако теперь недостат-
ком являются конфликты доступа к shared
памяти. Особенностью этой памяти явля-
ется то, что она разделена на 16 банков,
при этом доступ соседних потоков к одно-
му банку осуществляется последовательно,
а к различным банкам – одновременно. В
варианте реализации, полученном в пре-
дыдущем пункте, все данные сохраняются
в первых элементах массива, которые ле-
жат в одном банке. Поэтому все обраще-
ния к элементам массива оказываются по-
следовательными.
Чтобы избежать такой ситуации,
можно провести преобразование, заклю-
чающееся в изменении порядка суммиро-
вания. Ранее на первом этапе суммирова-
лись соседние элементы, потом удаленные
на 2 и т.д. Перейдем к другому порядку
суммирования: на первом этапе будем
суммировать максимально удаленные эле-
менты и далее будем уменьшать расстоя-
ние между элементами. Для этого исполь-
зуется следующая система правил:
1. For(DeclarationAssignment(s,int
,1),
s<Dot(blockDim,x),MulAssignment
(s,2), $body) ->
For(DeclarationAssignment(s,int
, Dot(blockDim,x)/2),s>0,
RShiftAssignment(s,1),$body)
2. If(Equals(tid % (2*s),0),$body)
-> If(tid < s,$body)
Эта система правил применяется к
программе, использующей shared память,
из пункта 3.3. Первое правило заменяет
параметры цикла: расстояние между сум-
мируемыми элементами уменьшается, а не
увеличивается. Второе правило заменяет
условие суммирования. Преобразованный
код имеет вид:
for ( int s = blockDim.x / 2; s > 0;
s >>= 1 )
{
if ( tid < s )
localData [tid] += localData
[tid + s];
__syncthreads ();
}
Заметим, что при использовании
данного преобразования вводятся две оп-
тимизации: уменьшение конфликтов при
доступе к банкам памяти и снижения ко-
личества ветвлений. В этом смысле преоб-
разование является более эффективным,
чем рассмотренное в п. 3.4. Однако данное
преобразование в большей мере использу-
ет знания о характере задачи (т.е. о воз-
можности проводить суммирование в дру-
гом порядке).
3.6. Изменение параметров вызо-
ва ядра. Преобразование из п. 3.5 позво-
лило избавиться от ветвлений и конфлик-
тов по доступу к памяти, но при этом по-
низилась эффективность использования
потоков. Половина из них вообще не про-
изводит вычислений (поскольку для них не
выполняется условие tid < s).
Поэтому можно модифицировать
ядро, исключив потоки, которые не вы-
полняют вычислений. В новой реализации,
первое суммирование производится еще на
стадии копирования данных в shared па-
мять. Поэтому в нем участвуют все пото-
ки. Правила для выполнения этого преоб-
разования имеют вид:
1. DeclarationAssignment(
i,int,Dot(blockIdx,x) *
Dot(blockDim,x) +
Dot(threadIdx,x)) ->
DeclarationAssign-
ment(i,int,2*Dot(blockIdx,x) *
Dot(blockDim,x) +
Dot(threadIdx,x))
2. Assignment (ArrayElement (
localData,tid), ArrayElement
(inData,i)) -> Assignment (
Моделі та засоби паралельних і розподілених програм
15
ArrayElement(localData,tid),
ArrayElement(inData,i)
+ArrayElement(inData,i+
Dot(blockDim,x)))
3. CallKernel(
Shape(ARR_SIZE/BLOCK_SIZE,
BLOCK_SIZE),parallelSum,[inData
,outData]) -> CallKernel(
Shape(ARR_SIZE/(2*
BLOCK_SIZE),BLOCK_SIZE),
parallelSum,[inData,outData])
Правило 2 определяет, что при ко-
пировании в shared память происходит
также суммирование. При этом получает-
ся, что один блок может обработать вдвое
больший объем информации. Поэтому
правило 1 задает изменение индекса во
входном массиве, а правило 3 – изменение
параметров вызова ядра (поскольку каж-
дый блок обрабатывает вдвое больший
объем данных, размер сетки уменьшается
вдвое). В результате код ядра меняется
следующим образом:
int i = 2*blockIdx.x * blockDim.x +
threadIdx.x;
localData [tid] = inData [i] + inData
[i+blockDim.x];
3.7. Развертывание цикла и уст-
ранение лишних синхронизаций. Одним
из недостатков исходной реализации, ко-
торый так и не был устранен после всех
преобразований, является достаточно
большое количество синхронизаций (на
каждом шаге суммирования). Для первых
итераций эти синхронизации являются не-
обходимыми, так как разные потоки могут
исполняться на разных warp’ах. Но на по-
следних итерациях, когда количество вы-
числяющих потоков не превосходит 32
(размер warp’а), синхронизации происхо-
дят автоматически. Поэтому можно изба-
виться от явного вызова __syncthreads
(). Кроме того, в этом случае можно раз-
вернуть цикл и убрать ветвление по при-
знаку tid < s.
Для этого используется следующая
система правил:
1. For(DeclarationAssignment(s,int
,Dot(blockDim,x)/2),s>0,
RShiftAssignment(s,1),$body) ->
[For(DeclarationAssignment(s,
int,Dot(blockDim,x)/2),s>32,
RShiftAssignment(
s,1),$body),_Unroll($body)]
2. _Unroll(If(tid<s,$body)) ->
If(tid<32,_Unroll($body))
3. _Unroll[$x:$y] ->
[_Unroll($x):_Unroll($y)]
4. _Unroll(PlusAssignment($x,$y))
->
_Unroll(PlusAssignment($x,$y),3
2)
5. _Unroll(PlusAssignment($x,$y),$
n) -> [PlusAssignment(
$x,_Replace($y,$n)),_Unroll(
PlusAssignment( $x,$y),$n/2)]
[$n>1]
6. _Unroll(PlusAssignment($x,$y),1
) -> PlusAssignment(
$x,_Replace($y,1))
7. _Replace(ArrayElement(localData
,tid+s),$n) -> ArrayElement(
localData,tid+$n)
8. _Unroll(SyncThreads) -> NIL
В этой системе правило 1 заменяет
предел цикла с 0 до 32, а также добавляет
за циклом заготовку для развернутого ва-
рианта. Правило 2 удаляет проверку tid <
s и добавляет проверку на срабатывание
развернутого цикла tid < 32. Правила 3–7
определяют замену тела цикла (суммиро-
вание) на развернутый вариант, который
повторяется для значений 32, 16, 8, 4, 2, 1.
В результате правило 8 удаляет явные вы-
зовы синхронизационной функции.
После применения правил код ядра
принимает следующий вид:
extern "C" __global__ void
parallelAdd5 ( int * inData, int *
outData )
{
__shared__ int localData
[BLOCK_SIZE];
int tid = threadIdx.x;
int i = 2*blockIdx.x *
blockDim.x + threadIdx.x;
localData [tid] = inData [i] +
inData [i+blockDim.x];
__syncthreads ();
for ( int s = blockDim.x / 2; s
> 32; s >>= 1 )
{
if ( tid < s )
localData [tid] +=
localData [tid + s];
__syncthreads ();
}
if ( tid < 32
{
localData [tid] +=
localData [tid + 32];
localData [tid] +=
localData [tid + 16];
localData [tid] +=
localData [tid + 8];
localData [tid] +=
localData [tid + 4];
localData [tid] +=
Моделі та засоби паралельних і розподілених програм
16
localData [tid + 2];
localData [tid] +=
localData [tid + 1];
}
if ( tid == 0 )
outData [blockIdx.x] =
localData [0];
}
Заметим, что для tid > 0 некото-
рые вычисления в развернутом цикле яв-
ляются избыточными, поскольку исполь-
зуется в конечном счете только значение
localData[0]. Это произошло из-за того,
что было исключено условие tid < s . Для
обычных платформ программирования та-
кие избыточные вычисления привели бы к
понижению производительности. Однако
для графического ускорителя это не про-
исходит: из-за большого количества вы-
числительных блоков можно выполнять
даже ненужные вычисления, если это уп-
рощает работу управляющего блока как и
происходит в данном случае.
3.8. Сравнение производительно-
сти. Для оценки эффективности различных
преобразований приведем данные сравне-
ния производительности различных вари-
антов реализации. Измерялось время
100-кратного исполнения алгоритма. Так-
же для сравнения приведено время испол-
нения последовательной версии алгоритма.
Результаты приведены в табл. 2.
Таблица 2. Сравнение различных
способов оптимизации
Вариант Время исполнения,
мс
Исходный 625
Shared память 300
Без ветвлений 186
Без конфликтов 77
Без простаивания 44
Развертывание 30
Последовательный 578
Из результатов видно, что исходная
реализация была весьма неэффективной (и
даже менее производительной, чем после-
довательная версия). Каждое из преобра-
зований давало вполне ощутимый эффект
по повышению производительности. В
итоге после применения всех преобразова-
ний получили программу, которая в 20 раз
эффективнее исходной.
Выводы
В работе рассмотрен подход к ав-
томатизации процесса разработки прило-
жений для графических ускорителей, ос-
нованный на использовании системы пе-
реписывающих правил Termware. Рас-
смотрены преобразования, которые приво-
дят как к распараллеливанию программ с
использованием платформы CUDA, так и к
оптимизации параллельных программ для
графических ускорителей. Применение
разработанных преобразований дает суще-
ственный эффект: производительность
приложений повышается в десятки раз.
При этом автоматизированное осуществ-
ление преобразований позволяет упро-
стить разработку и понизить вероятность
ошибки разработчика.
Дальнейшие исследования в данном
направлении предполагают разработку до-
полнительных преобразований для распа-
раллеливания и оптимизации кода для
графических ускорителей, а также оценку
их эффективности на различных примерах.
Также предполагается поддержка других
технологий программирования для графи-
ческих ускорителей, таких как AMD
Stream.
Авторы выражают благодарность
Н. Котюку за помощь в проведении
вычислительных экспериментов.
1. Эндрюс Г. Основы многопоточного,
параллельного и распределенного про-
граммирования. – М.: ИД ”Вильямс”,
2003. – 512 с.
2. Ryoo S., Rodrigues C.I., Baghsorkhi S.S.,
Stone S.S., Kirk D.B., and Hwu W.W. Op-
timization principles and application per-
formance evaluation of a multithreaded
GPU using CUDA. In Proceedings of the
13th ACM SIGPLAN Symposium on
Principles and Practice of Parallel Pro-
gramming (Salt Lake City, UT, USA,
February 20–23, 2008). PPoPP '08. ACM,
New York, NY. – P. 73–82.
3. Fatahalian K., Sugerman J., and Hanra-
han P. Understanding the efficiency of
GPU algorithms for matrix-matrix multi-
plication. In ACM SIGGRAPH/EURO-
GRAPHICS Conf. on Graphics Hardware,
2004. – P. 133–137.
Моделі та засоби паралельних і розподілених програм
17
4. General-Purpose Computation Using
Graphics Hardware. http://www.gpgpu.
org .
5. NVidia CUDA technology. http://www.
nvidia.com/cuda .
6. AMD (ATI) Stream technology. http://
www.amd.com/stream .
7. Doroshenko A., Shevchenko R. A Rewrit-
ing Framework for Rule-Based Program-
ming Dynamic Applications, Fundamenta
Informaticae. – 2006. – Vol. 72, N 1–3. –
P. 95–108.
8. TermWare. – http://www.gradsoft.com.ua/
products/termware_rus.html .
9. Дорошенко А.Е., Шевченко Р.С. Систе-
ма символьных вычислений для
программирования динамических при-
ложений // Проблемы программирова-
ния. – 2005. – № 4. – С. 718–727.
10. Дорошенко А.Е., Жереб К.А.,
Яценко Е.А. Формализованное проек-
тирование эффективных многопоточ-
ных программ // Проблемы программи-
рования. – 2007. – № 1. – С. 17–30.
11. Дорошенко А.Е., Жереб К.А.,
Яценко Е.А. Об оценке сложности и ко-
ординации вычислений в многопоточ-
ных программах // Проблемы програм-
мирования. – 2007. – № 2. – С. 41–55.
12. Lee S., Min S., and Eigenmann R.
OpenMP to GPGPU: a compiler
framework for automatic translation and
optimization. In Proceedings of the 14th
ACM SIGPLAN Symposium on
Principles and Practice of Parallel
Programming (Raleigh, NC, USA,
February 14–18, 2009). PPoPP '09. ACM,
New York, NY. – P. 101–110.
13. OpenMP specification. http://openmp.org/
wp/ .
14. Baskaran M., Bondhugula U., Krishna-
moorthy S., Ramanujam J., Rountev A.,
and Sadayappan P. A compiler frame-
work for optimization of affine loop nests
for gpgpus. In Proceedings of the 22nd
Annual international Conf. on Supercom-
uting (Island of Kos, Greece, June 07–12,
2008). ICS '08. ACM, New York, NY. –
P. 225–234.
15. Ma W. and Agrawal G. A compiler and
runtime system for enabling data mining
applications on gpus. In Proceedings of
the 14th ACM SIGPLAN Symposium on
Principles and Practice of Parallel
Programming (Raleigh, NC, USA,
February 14 – 18, 2009). PPoPP '09.
ACM, New York, NY. – P. 287–288.
16. Allusse Y., Horain P., Agarwal A., and
Saipriyadarshan C. GpuCV: an open
source GPU-accelerated framework for
image processing and computer vision. In
Proceeding of the 16th ACM International
Conf. on Multimedia (Vancouver, British
Columbia, Canada, October 26–31, 2008).
MM '08. ACM, New York, NY. –
P. 1089–1092.
17. Lefohn A.E., Sengupta S., Kniss J.,
Strzodka R., and Owens J.D. Glift:
Generic, efficient, random-access GPU
data structures. ACM Trans. Graph. 25,
1 Jan. 2006. – P. 60–99.
18. Han T.D. and Abdelrahman T.S.
hiCUDA: a high-level directive-based
language for GPU programming. In
Proceedings of 2nd Workshop on General
Purpose Processing on Graphics
Processing Units (Washington, D.C.,
March 08 – 08, 2009). GPGPU-2, Vol.
383. ACM, New York, NY. – P. 52–61.
19. Hou, Q., Zhou, K., and Guo, B. BSGP:
bulk-synchronous GPU programming. In
ACM SIGGRAPH 2008 Papers (Los
Angeles, California, August 11 – 15,
2008). SIGGRAPH '08. ACM, New York,
NY. – P. 1–12.
20. CUDA .NET http://www.gass-ltd.co.il/en/
products/cuda.net/ .
21. Microsoft Research Accelerator Project
http://research.microsoft.com/en-
us/downloads/648909e1-cb85-46c4-9a94-
3cca55971b1d/ .
22. Bitonic Sort Algorithm. http://www.itl.
nist.gov/div897/sqg/dads/HTML/bitonicS
ort.html .
Получено 12.06.2009
Моделі та засоби паралельних і розподілених програм
18
Об авторах:
Дорошенко Анатолий Ефимович,
доктор физико-математических наук,
профессор, заведующий отделом теории
компьютерных вычислений Института
программных систем НАН Украины,
Жереб Константин Анатольевич,
аспирант Физико-технического учебно-
научного центра НАН Украины.
Место работы авторов:
Институт программных систем
НАН Украины,
03680, Киев-187,
Проспект Академика Глушкова, 40.
Тел.: (044) 526 1538.
e-mail: dor@isofts.kiev.ua
Физико-технический учебно-научный
центр НАН Украины,
03142, Киев-142,
Бульвар Вернадского, 36.
Тел.: (044) 424 3025.
e-mail: zhereb@gmail.com
|