Параллельные вычисления общего назначения на графических процессорах [Кирилл Александрович Некрасов] (pdf) читать онлайн

К. А. НЕКРАСОВ
С. И. ПОТАШНИКОВ
А. С. БОЯРЧЕНКОВ
А. Я. КУПРЯЖКИН

ПАРАЛЛЕЛЬНЫЕ ВЫЧИСЛЕНИЯ
ОБЩЕГО НАЗНАЧЕНИЯ
НА ГРАФИЧЕСКИХ ПРОЦЕССОРАХ
Учебное пособие

Министерство образования и науки Российской Федерации
Уральский федеральный университет
имени первого Президента России Б. Н. Ельцина

К. А. Некрасов, С. И. Поташников,
А. С. Боярченков, А. Я. Купряжкин

Параллельные вычисления
общего назначения
на графических процессорах
Учебное пособие
Рекомендовано методическим советом УрФУ
для студентов, обучающихся по направлениям подготовки
14.04.02 — Ядерная физика и технологии;
09.04.02 — Информационные системы и технологии;
14.04.01 — Ядерные реакторы и материалы

Екатеринбург
Издательство Уральского университета
2016

УДК 004.032.24:004.383.5(075.8)
ББК 32.97я73
П18

Авторы:
Некрасов К. А., Поташников С. И., Боярченков А. С., Купряжкин А. Я.

Рецензенты:
Институт теплофизики УрО РАН (д‑р физ.-мат. наук, проф. В. Г. Байда‑
ков); гл. науч. сотр. лаборатории математического моделирования Ин‑
ститута промышленной экологии УрО РАН д‑р физ.-мат. наук, проф.
А. Н. Вадаксин
    Параллельные вычисления общего назначения на графических процессо­
П18 рах : учебное пособие / К. А. Некрасов, С. И. Поташников, А. С. Боярчен‑
ков, А. Я. Купряжкин. — Екатеринбург : Изд-во Урал. ун-та, 2016. — 104 с.
ISBN 978-5-7996-1722-6
В учебном пособии изложены основные принципы организации высокоско‑
ростных параллельных вычислений на графических процессорах. Рассмотрены
подходы к программированию графических процессоров с использованием шей‑
дерной модели и NVIDIA CUDA. Проанализированы примеры.
Пособие предназначено для проведения практических занятий по програм‑
мированию графических процессоров для магистрантов.
Библиогр.: 22 назв. Рис. 24. Табл. 3. Прил. 1.

ISBN 978-5-7996-1722-6

УДК 004.032.24:004.383.5(075.8)
ББК 32.97я73

© Уральский федеральный
университет, 2016

Введение

Д

ля предсказания характеристик и поведения больших си‑
стем широко используется вычислительное моделирова‑
ние, наиболее принципиальным методом повышения про‑
изводительности которого является распараллеливание вычислений.
До последнего времени наиболее доступными системами для парал‑
лельных расчетов были кластеры персональных компьютеров или
близких к ним по архитектуре машин, в которых вычисления про‑
изводились на центральных процессорах общего назначения (CPU).
Однако такие кластеры достаточно дороги и сложны в эксплуатации.
К тому же архитектура CPU персональных компьютеров не оптими‑
зирована для интенсивных математических вычислений, поскольку
основной задачей таких процессоров является исполнение последо‑
вательных программ со сложным ветвлением.
Во второй половине 1990‑х годов началось быстрое развитие графических процессоров (GPU) — дополнительных вычислительных
устройств для ускоренного исполнения алгоритмов визуализации трех‑
мерных сцен [1–3]. Поскольку трехмерная визуализация допускает
эффективное распараллеливание расчетов, графические процессоры
разрабатывались как поточно-параллельные системы с большим ко‑
личеством вычислительных блоков, конвейерной обработкой данных
и памятью с максимальной пропускной способностью.
Современные графические процессоры выполняют не только стан‑
дартные алгоритмы визуализации, но и сложные пользовательские
программы, что позволяет решать на них задачи общего назначения,
включая физико-математическое моделирование [4–6]. При парал‑
лельных расчетах GPU может обеспечить производительность класте‑
ра из сотен обычных персональных компьютеров. По соотношению
производительности и цены графические процессоры имеют большое
преимущество перед другими вычислительными системами, в том чис‑
ле перед специализированными суперкомпьютерами.
3

1. Структура и возможности
вычислительной системы
с графическим процессором

1.1. Задача компьютерной визуализации трехмерных сцен

В

о многих областях использования компьютеров существует
задача визуализации (представления на экране) трехмерных
изображений (в дальнейшем — сцен). Трехмерная (3D) визу‑
ализация необходима, например, в компьютерных играх, при созда‑
нии анимации и спецэффектов, а также для инженерного проектиро‑
вания, наглядного представления физико-математических моделей.
Поскольку графические процессоры создавались именно для реше‑
ния задачи визуализации, рассмотрим ее основные составляющие.
Пусть, для определенности, трехмерная сцена представляет со‑
бой совокупность поверхностей, разбитых для дискретной компью‑
терной обработки на плоские треугольники. В некоторых массивах,
которые обычно называют текстурами, хранится информация о цве‑
тах треугольников. Кроме того, известны положения и характеристи‑
ки источников света. Задача визуализации состоит в том, чтобы сфор‑
мировать изображение этой сцены на плоскости экрана, положение
которой определяется точкой зрения наблюдателя.
Процесс формирования изображения включает в себя следующие
основные этапы:
· проектирование треугольников, представляющих сцену, на пло‑
скость экрана;
· разбиение полученных проекций на отдельные пиксели, для ко‑
торых будут определяться цвета (стадия растеризации треуголь‑
ников);
4

1.2. Архитектура графического процессора (GPU)

· определение видимого цвета элементов поверхности треуголь‑
ников с учетом исходного цвета и отражающих свойств самой
поверхности, прозрачности других поверхностей, освещенно‑
сти и теней.
Обрабатываемые сцены обычно состоят из очень многих треуголь‑
ников, которые разбиваются на еще большее количество пикселей, так
что их визуализация требует больших вычислительных ресурсов. Вме‑
сте с тем и вершины, и пиксели можно обрабатывать почти или со‑
всем независимо друг от друга. Соответственно задача визуализации
допускает очень эффективное распараллеливание.
Именно в целях распараллеливания графических вычислений цен‑
тральные процессоры персональных компьютеров стали с у п е р с к а л я р н ы ми — получили возможность одновременного (в екторного)
исполнения некоторых операций сразу над несколькими числами (рас‑
ширения 3DNow! и SSE).
Большинство центральных процессоров ПК состоит из несколь‑
ких ядер, что дополнительно увеличивает их потенциал как систем
для параллельных вычислений. Тем не менее исторически сложилось
так, что наиболее эффективными устройствами для распараллелива‑
ния вычислений на ПК стали и до сих пор остаются специализиро‑
ванные графические процессоры (GPU). Эти процессоры, разраба‑
тывавшиеся именно для обработки графики, практически полностью
ориентированы на параллельную обработку данных.

1.2. Архитектура графического процессора (GPU)
1.2.1. Распараллеливание вычислений по данным
Конструктивно графический процессор представляет собой вычис‑
лительное устройство, работающее отдельно от центрального процес‑
сора, параллельно с ним. Обычно графические процессоры размещают
на отдельных печатных платах с собственной системой охлаждения,
которые называют графическими ускорителями (или видеокартами).
Вместе с GPU на плате графического ускорителя расположена видео­
память — специализированная оперативная память, в которой хра‑
нятся обрабатываемые графическим процессором массивы данных.
5

1. Структура и возможности вычислительной системы с графическим процессором

На рис. 1.1 в качестве примера показана архитектура графическо‑
го процессора G80 — одного из процессоров компании NVIDIA [4].
Главной (и общей для всех графических процессоров) характеристи‑
кой этой архитектуры является то, что GPU представляет собой си‑
стему из параллельных вычислительных устройств, каждое из которых
применяет заданную, единую для всех устройств, программу (в ычисл и т е ль н о е я д р о , англ. kernel) к различным элементам входных мас‑
сивов данных, расположенных в общей памяти.
Отдельный «вычислитель» (процессор),
осуществляющий конвейерную обработку
данных

Центральный процессор (Host)
Устройство управления вычислительными процессами
(Tread Execution Manager)

SM

Параллельный кеш данных (Parallel Data
Cache, Shared Memory) – разделяемая память, одновременно доступная всем «вычислителям» одного мультипроцессора

16 мультипроцессоров
SM

SM

SM

SM

ј

SM

SM

SM

SM

Кеш графического процессора, доступный всем мультипроцессорам.
Центральному процессору доступен частично и только для записи
Общая память для размещения входных данных и результатов
(видеопамять, Global Memory), доступная как всем мультипроцессорам
GPU, так и центральному процессору. Размер − до нескольких гигабайтов
Рис. 1.1. Архитектура графического процессора NVIDIA G80 [4]

Параллельная архитектура графических процессоров ориентиро‑
вана на исполнение алгоритмов, в которых элементы больших вход‑
ных массивов обрабатываются одинаковым образом независимо или
почти независимо друг от друга, то есть использующих распарал л еливание вычислений по данным. Множества элементов, подверга‑
емых однотипной независимой обработке, называют потоками (дан‑
ных либо результатов), так что графические процессоры осуществляют
поточно-параллельную обработку данных.
6

1.2. Архитектура графического процессора (GPU)

Концепция программирования, заключающаяся в потоко‑
вой обработке данных, известна под аббревиатурой SIMD (от англ.
Single Instruction — Multiple Data — одна инструкция для множества
данных, рис. 1.2). Процессор, работающий по принципу SIMD, пре‑
образует поток данных в поток результатов, используя программу как
функцию преобразования.
Инструкции
(вычислительное ядро)

GPU
Данные

Результаты

Рис. 1.2. Поточно-параллельная обработка данных SIMD

Выбор концепции SIMD для графических процессоров обусловлен
тем, что она обеспечивает параллельное использование большого ко‑
личества «вычислителей» без явного управления ими: распределения
задач, синхронизации вычислений и коммуникации между параллель‑
ными расчетами. Разработчикам GPU это позволяет за счет упроще‑
ния архитектуры добиваться большей производительности, а при про‑
граммировании сокращает работу.
C 2001 по 2006 годы графические процессоры включали в себя «вы‑
числители» двух типов: вер ш инные и пиксел ьные конв ейеры (или
шейдеры, см. пояснения ниже) [7–8]. Первые были предназначены для
проектирования на плоскость экрана вершин, задающих отображае‑
мые поверхности, а вторые — для расчета цветов пикселей на экране.
Так выпущенные в 2006 году процессоры ATI Radeon X1900–1950 име‑
ли 8 вершинных и 48 пиксельных конвейеров, которые к тому же были
суперскалярными (одновременно обрабатывали по 4 компоненты век‑
тора).
В 2007 году производители GPU перешли от различающихся вер‑
шинных и пиксельных конвейеров к унифицированным потоковым
процессорам, заменяющим как вершинные, так и пиксельные конвей‑
еры. Графические процессоры 2014–2015 годов выпуска ATI Radeon
Fiji XT и NVIDIA GM200–400 включают 4096 и 3072 потоковых процес‑
7

1. Структура и возможности вычислительной системы с графическим процессором

сора (для 32‑битовых вычислений «одинарной» точности) при несколь‑
ко различной внутренней архитектуре. При этом производительность
графических процессоров продолжает быстро увеличиваться.
1.2.2. Взаимодействие графического и центрального процессоров
Графический процессор не имеет средств прямого взаимодействия
с устройствами ввода-вывода (кроме монитора), а также доступа к опе‑
ративной памяти компьютера. Поэтому управление графическим про‑
цессором осуществляется только через центральный процессор. Схема
взаимодействия центрального и графического процессоров приведе‑
на на рис. 1.3.
Графический ускоритель (видеокарта)
Регистры и разделяемая память вычислительных блоков

GPU

• • • Языки для GPU: HLSL, CUDA

Кеш для
констант

Вычислительные
блоки GPU

Программа
для GPU

Общая память GPU (видеопамять)
Шина
данных(обычно
PCI-Express
либо AGP
Шина
данных
− PCI-Express)
• • Драйвер графического процессора
• Языки программирования высокого
уровня для CPU

Оперативная память
компьютера

CPU

Пользователь

Внешние устройства

Рис. 1.3. Схема взаимодействия CPU и GPU с памятью и между собой

Графические ускорители подключаются к системной плате персо‑
нального компьютера через высокоскоростную шину данных (в на‑
стоящее время PCI-Express). Посредством этой шины центральный
процессор получает доступ к видеопамяти, а также к некоторым раз‑
делам кеш-памяти, расположенной на самом графическом процессо‑
ре. Через эту же шину CPU загружает в графический процессор про‑
грамму и запускает ее.
8

1.2. Архитектура графического процессора (GPU)

Перед запуском программы, исполняемой на GPU, центральный
процессор передает графическому процессору данные двух видов:
· значения констант, используемых в программе;
· один или несколько больших массивов данных для потоковой
обработки.
К константам необходим постоянный и быстрый доступ, поэто‑
му они записываются в кеш-память (или регистры), расположенную
на кристалле графического процессора. Массивы данных часто быва‑
ют настолько велики, что целиком в кеш-память не помещаются. Од‑
нако при простейшей потоковой обработке каждый из элементов мас‑
сивов данных используется только один раз. Так что для хранения этих
массивов предназначена видеопамять (общая память), представляю‑
щая собой отдельные микросхемы на плате графического ускорителя.
Она работает медленнее кеш-памяти и регистров, зато имеет бόльший
объем, до нескольких гигабайтов.
Результаты своей работы графический процессор может сразу за‑
писывать в раздел видеопамяти, называемый буфером кадра, откуда
они передаются на монитор. Но существует также возможность вооб‑
ще не отображать расчет на экране, а копировать результаты из видео‑
памяти в оперативную память компьютера, где они становятся доступ‑
ными для дальнейшей обработки центральным процессором. На этом
и основано использование графических процессоров для вычислений
общего назначения, не связанных с обработкой графики.
На схеме (рис. 1.3) указаны также типы программ, которые исполня‑
ются центральным и графическим процессорами на различных этапах
обработки данных. Такие программы будут далее подробно рассмотрены.
1.2.3. Иерархия памяти, доступной центральному
и графическому процессорам
Как показано на схемах (рис. 1.1, рис. 1.3), в программах для гра‑
фических процессоров используется память нескольких разновидно‑
стей с различными характеристиками и назначением. Это так, посколь‑
ку время, затрачиваемое центральным и графическим процессорами
на операции чтения данных из памяти и записи в память, является
одним из важнейших факторов, определяющих быстродействие про‑
грамм для графических процессоров.
9

1. Структура и возможности вычислительной системы с графическим процессором

Использование памяти различных типов обусловлено необходимо‑
стью баланса между объемом памяти и скоростью доступа к данным.
Разновидности памяти, имеющие наибольшую емкость, обычно ха‑
рактеризуются бόльшим временем доступа к данным, и наоборот. Бы‑
стродействие памяти, в свою очередь, определяется двумя характери‑
стиками — латентностью и пропускной способностью.
Латентность — это время доступа к памяти, точнее, время ожида‑
ния процессором данных после запроса. Латентностью определяется
производительность вычислений при решении задач, требующих ча‑
стого обращения к различным по расположению неупорядоченным
ячейкам памяти (пр о изво ль ный д о с ту п к памяти). Такой обмен
с памятью характерен для интерактивных приложений (программы,
интенсивно взаимодействующие с другими приложениями и с поль‑
зователями), а также для приложений, управляющих сложными про‑
цессами. Подобные алгоритмы обычно исполняются центральным
процессором.
Как можно меньшая латентность памяти необходима современ‑
ным центральным процессорам, работающим на очень высоких ча‑
стотах, еще и потому, что таким частотам должна соответствовать вы‑
сокая скорость доступа к данным.
Пропускная способность памяти характеризует объем данных, ко‑
торые могут быть переданы к процессору или от процессора за еди‑
ницу времени. Высокая пропускная способность оказывается эф‑
фективнее низкой латентности в задачах, позволяющих организовать
по с ле до в ат е ль ный д о с туп к памяти — считывание (или запись)
данных из ячеек памяти, расположенных друг за другом, непрерыв‑
ным потоком.
При поточно-параллельной обработке данных предпочтителен
именно последовательный доступ к памяти, поэтому видеопамять,
предназначенная для обмена данными с графическим процессо‑
ром, должна обладать максимальной пропускной способностью даже
в ущерб латентности.
Иерархия памяти, используемой центральным и графическим про‑
цессорами, уже отчасти отражена на рис. 1.1 и рис. 1.3. На рис. 1.4 при‑
ведена еще одна схема, иллюстрирующая основные принципы исполь‑
зования памяти различных типов. Особенности и назначение каждого
вида памяти кратко пояснены рядом и будут подробнее прокоммен‑
тированы в примерах.
10

1.2. Архитектура графического процессора (GPU)

1. Регистры и кеш CPU – память небольшого (до нескольких мегабайтов) объема, расположенная на кристалле центрального процессора.
Характеризуется минимальной латентностью. Доступна только центральному процессору. Обычно этой памятью автоматически управляРегистры вычислиет процессор, записывая в нее данные, используемые наиболее часто
тельных блоков GPU
2. Оперативная память компьютера – память большого (до десятков гигабайтов) объема в форме микросхем, расположенных на системной
плате. Имеет среднюю латентность и пропускную способность. ДоРазделяемая память
ступна только центральному процессору. Используется для хранения
вычислительных бловсех обрабатываемых CPU данных, не помещающихся в кеше
ков GPU
3. Видеопамять – память, в которую CPU копирует данные, предназнаПамять для констант на ченные для обработки графическим процессором. Сюда же GPU записывает результаты расчета. Видеопамять имеет большой объем (до некристалле GPU
скольких гигабайтов) и максимально возможную пропускную способность для потокового доступа к данным. Расположена на плате графиВидеопамять на плате
ческого ускорителя, но не на самом кристалле GPU. Доступна как граграфического ускорифическому, так и центральному процессорам для чтения и записи. Лателя (Общая память)
тентность со стороны GPU средняя, а со стороны CPU очень большая
(медленная передача данных)
Оперативная память
4. Память для констант на кристалле GPU. Имеет низкую латентность со
стороны GPU и предназначена для быстрого произвольного доступа
компьютера
GPU к константам. Доступна центральному процессору для записи
констант, а графическому процессору – для их чтения
Регистры и кеш CPU
5-6. Разделяемая память и регистры вычислительных блоков GPU. Расположены на кристалле графического процессора и доступны только
ему. Имеют очень низкую латентность (высокую скорость), но небольЦентральный процесшой объем (десятки килобайтов). Предназначены для произвольного
сор (CPU)
доступа к данным со стороны параллельных вычислительных процессов

Графический процессор
(GPU)

6
5
4
3
2
1

Риc. 1.4. Использование памяти графическим
и центральным процессорами

Отметим основные принципы разделения памяти между централь‑
ным и графическим процессорами. В персональном компьютере для
хранения большей части данных и программ предназначена так на‑
зываемая о п е р ат ивная памя т ь . Микросхемы этой памяти распо‑
ложены на системной плате (не на самом процессоре), они доступны
центральному процессору для произвольных чтения и записи данных.
Графическому процессору эта память не доступна.
На плате графического ускорителя расположена видеопамять (об‑
щая память), технологически и по объему подобная оперативной па‑
мяти компьютера. Она доступна графическому процессору для чтения
и записи в последовательном и в произвольном режиме, но оптимизи‑
рована для последовательного (потокового) доступа, поскольку пред‑
назначена для хранения массивов данных, обрабатываемых GPU в по‑
точно-параллельном режиме.
Видеопамять является основным средством обмена данными между
графическим и центральным процессорами. Центральный процессор
11

1. Структура и возможности вычислительной системы с графическим процессором

тоже имеет к этой памяти доступ для чтения и записи. Перед началом
вычислений на GPU он копирует исходные данные из оперативной
памяти в видеопамять, а после окончания расчета копирует результа‑
ты обратно в оперативную память. Обмен данными между централь‑
ным процессором и видеопамятью происходит сравнительно медлен‑
но, так что проводится обычно только в начале и конце расчета.
Поскольку для размещения исходных потоков данных для GPU
и результатов его работы видеопамять должна иметь сравнитель‑
но большой объем (на современных графических ускорителях
до нескольких гигабайтов), она выпускается в виде отдельных ми‑
кросхем и расположена не на самом кристалле графического процес‑
сора, а на плате графического ускорителя. Поэтому работает видео‑
память медленно по сравнению с регистрами и разделяемой памятью
GPU (ниже).
Для повышения производительности GPU оснащены небольшой
быстрой памятью (регистры процессора, память для констант, раз‑
деляемая память вычислительных блоков), расположенной прямо
на кристалле процессора. Доступ центрального процессора к этой
памяти, наоборот, медленен и ограничен (см. следующий подпара‑
граф). Быстрая память используется для хранения данных, к кото‑
рым графическому процессору нужен произвольный, а не потоковый
доступ. Это могут быть, например, константы, задаваемые перед на‑
чалом расчета.
Улучшение микросхем памяти связано с уменьшением латентно‑
сти и увеличением пропускной способности. В последние годы про‑
пускная способность памяти увеличивается почти в 2 раза быстрее,
чем уменьшается латентность. Быстрое возрастание пропускной спо‑
собности памяти означает перспективность разработки и реализации
алгоритмов, ориентированных на потоковую обработку данных, эф‑
фективность которой как раз определяется пропускной способно‑
стью памяти.
1.2.4. Конвейерная обработка данных
Уже в первых графических процессорах была реализована конвей‑
ерная обработка потоков данных. Заключается она в том, что каждый
из параллельных «вычислителей» может одновременно обрабатывать
12

1.2. Архитектура графического процессора (GPU)

несколько элементов потока, применяя к ним различные операции,
например: загружать из памяти одно число, модифицировать другое
и сохранять третье. Необходимо, чтобы операции, исполняемые кон‑
вейером одновременно, не зависели друг от друга, что как раз харак‑
терно для поточно-параллельных расчетов.
Совершенствование вершинных и пиксельных шейдеров создало
возможность программирования GPU для решения неграфических
задач. Изначально при этом вместо координат вершин требовалось
передавать исходные числовые данные, алгоритм обработки которых
программировали в форме вершинного шейдера. В настоящее вре‑
мя графические процессоры позволяют формулировать задачи об‑
щего назначения «естественным» для них образом, однако конвейер‑
ную обработку данных сохраняют. В качестве иллюстрации принципа
конвейерной обработки данных на рис. 1.5 показаны основные эта‑
пы гр аф и ч е с ко го ко нве йе р а, используемого графическим процес‑
сором визуализации трехмерных сцен [9–12].
Поток
Поток преобразотреугольников
ванных вершин
Отсечение
Вершинный
Сборка
невидимых
шейдер
треугольников
треугольников

Поток
вершин

Текстуры – массивы с данными о цветах,
неровностях и других характеристиках
поверхностей сцены, которые используются
для их отображения

Поток
необработанных
фрагментов

Пиксельный
шейдер

Композиция
Поток
фрагментов

Разбиение
проекций на
фрагменты
(растеризация)

Буфер кадра
(изображение в
памяти)

Поток
пикселей

Рис. 1.5. Этапы графического конвейера, исполняемые
на графическом процессоре [9]

13

1. Структура и возможности вычислительной системы с графическим процессором

1.3. Уровни управления графическим процессором
и основные системы программирования GPU
1.3.1. Уровни управления графическим процессором
Как уже отмечено выше, графический процессор управляется поль‑
зователем не напрямую, а через центральный процессор компьютера.
При этом CPU в большинстве случаев исполняет как минимум три вза‑
имодействующие программы, которые различаются по происхожде‑
нию, назначению и возможностям. Можно сказать, что эти програм‑
мы работают на трех уровнях управления графическим процессором
(см. рис. 1.7). На двух уровнях, «ближайших» к графическому процес‑
сору, исполняются стандартные программы, избавляющие програм‑
миста от необходимости написания рутинных процедур, таких, напри‑
мер, как выделение видеопамяти для данных или представление этих
данных в формате, необходимом графическому процессору.
1.3.2. Драйвер графического процессора
На самом низком уровне, «ближайшем» к графическому процес‑
сору, исполняется драйвер графического процессора — программа,
которая непосредственно управляет как самим GPU, так и видеопа‑
мятью. Фактически драйвер работает как часть операционной систе‑
мы. Он получает от других приложений запросы — задачи для графи‑
ческого процессора, после чего передает их на GPU в оптимальном
порядке и необходимом формате. В частности, функциями драйвера
являются такие:
· распределение ресурсов графического процессора между
несколькими приложениями;
· загрузка данных в видеопамять и регистры графического про‑
цессора, копирование данных из видеопамяти в оперативную
память компьютера;
· загрузка в графический процессор и запуск на исполнение поль‑
зовательских программ;
· автоматическое распределение расчетов между параллельными
вычислительными блоками графического процессора;
14

1.3. Уровни управления графическим процессором и основные системы программирования GPU

· исполнение стандартных алгоритмов обработки графики;
· регулирование частоты графического процессора для поддер‑
жания рабочей температуры.
Драйверы обычно разрабатываются самими производителями гра‑
фических процессоров, поскольку это требует детального знания их
технического устройства и особенностей. Они обновляются для под‑
держки процессоров новых моделей.
Благодаря существованию драйверов, пользовательские программы
могут выполняться на различных GPU без перекомпиляции, если GPU
и драйвер поддерживают функции, задействованные в программе.
1.3.3. Интерфейсы программирования приложений
К драйверу GPU можно обращаться из пользовательских прило‑
жений напрямую, однако это требует явного включения в текст про‑
граммы большого количества стандартных операций управления гра‑
фическим процессором. Для того чтобы автоматизировать выполнение
подобных операций и унифицировать взаимодействие пользователь‑
ских приложений с различными устройствами различных произво‑
дителей, существуют библиотеки специализированных процедур, из‑
вестные как п р о г р а м м н ы й и н т е р ф е й с п р и л о ж е н и й (application
programming interface, или API).
В качестве разработчиков API для графических процессоров вы‑
ступают как сами производители GPU, так и компании, специализи‑
рующиеся на программном обеспечении. Под управлением Windows
наиболее широко используется пакет API DirectX от Microsoft, а для
других операционных систем такие библиотеки разрабатываются в ос‑
новном по спецификациям OpenСL и OpenGL.
Пакет мультимедийных библиотек DirectX и библиотека DirectCompute
Для операционной системы Windows корпорацией Microsoft уже
давно выпускается собственный пакет мультимедийных библиотек
API, который в целом называется DirectX. В этот пакет входит библио‑
тека DirectCompute, предназначенная для реализации вычислений об‑
щего назначения на GPU (см., например, [13]). Начиная с Windows 98
пакет DirectX поставляется вместе с операционной системой, так что
практически входит в ее состав. Интерфейс программирования при‑
15

1. Структура и возможности вычислительной системы с графическим процессором

ложений DirectCompute поддерживается операционными системами
Windows, начиная с Windows Vista.
В состав DirectX входит компилятор языка высокого уровня для
пользовательского программирования графических процессоров HLSL
(High Level Shading Language). Этот язык разработан Microsoft специаль‑
но для DirectX, он совершенствуется вместе с этим пакетом библиотек.
Отметим, что программы, непосредственно предназначенные для
выполнения на графических процессорах, традиционно называют
ш е й д е р а м и , от англ. shader. Это название относится к программам
построения теней при отображении трехмерных сцен в компьютер‑
ной графике.
OpenGL и OpenCL
OpenGL (Open Graphics Library — открытая графическая библио‑
тека) представляет собой с пе циф икацию , определяющую независи‑
мый от языка программирования кросс-платформенный программный
интерфейс (API) для написания приложений, использующих двумер‑
ную и трехмерную компьютерную графику. Спецификация не явля‑
ется конкретной библиотекой, она только описывает набор функций
и точное их поведение.
На основе этой спецификации OpenGL разработчики создают би‑
блиотеки функций, называемые р е али зациями. Реализации по воз‑
можности используют аппаратные средства GPU, а другие требуемые
функции эмулируют программно. Чтобы называться реализациями
OpenGL, библиотеки должны удовлетворять определенным тестам
на соответствие (conformance tests). Существует возможность написа‑
ния на OpenGL вычис лит е ль ных шейдеров (compute shaders), реа‑
лизующих вычисления общего назначения на GPU.
Поддержкой и модернизацией OpenGL с 2006 года занимается ин‑
дустриальное объединение Khronos Group. В 2008 году этой группой
(при значительном участии Nvidia) была выпущена спецификация
OpenGL 3.0, которая предоставляет для разработки графических при‑
ложений примерно те же возможности, что и DirectX 10.
Консорциумом Khronos Group разработана новая спецификация
OpenCL (Open Computing Language — открытый язык для вычисле‑
ний), включающая в себя интерфейсы программирования приложе‑
ний и язык программирования для гетерогенных систем, имеющих
в своем составе центральные и графические процессоры, как и дру‑
16

1.3. Уровни управления графическим процессором и основные системы программирования GPU

гие процессоры для высокопроизводительных параллельных вычис‑
лений. Существуют многочисленные реализации OpenCL, в том чис‑
ле от Nvidia, AMD, IBM, Apple и Intel.
Главным преимуществом OpenGL и OpenCL перед DirectX являет‑
ся существование их реализаций не только для Windows, но и для боль‑
шинства других операционных систем, в том числе Unix-подобных си‑
стем (Linux, Mac OS), систем игровых приставок. Такие реализации
разрабатываются производителями центральных и графических про‑
цессоров, поэтому реализации максимально используют аппаратное
ускорение вычислений.
1.3.4. Пользовательское приложение
На «верхнем» уровне управления графическим процессором нахо‑
дятся пользовательские приложения — программы, разрабатываемые
для решения конечных задач, таких как создание трехмерных сцен,
проведение физико-математических расчетов. Структура пользова‑
тельского приложения, использующего API типа DirectCompute или
OpenCL, показана на рис. 1.6.
При использовании DirectCompute или OpenCL пользовательское
приложение состоит по крайней мере из двух частей, для написания
которых используются различные языки программирования. Одна
из этих частей — это программа для графического процессора, кото‑
рую обычно называют в ы ч и с л и т е л ь н ы м я д р о м, а другая часть —
це н т р аль н ы й пр о цес с о р .
Для написания вычислительных ядер используются специаль‑
ные языки программирования. При работе с DirectX это может быть
либо язык высоко уровня HLSL [11], либо низкоуровневый ассемблер
(assembly shader language). Вычислительное ядро хранится в отдельном
текстовом файле, оно компилируется и загружается в память GPU
средствами DirectX или OpenCL непосредственно перед запуском. Эти
API используются также для загрузки в графический процессор и ви‑
деопамять данных и констант, необходимых для расчета.
Часть пользовательского приложения, которая исполняется на цен‑
тральном процессоре, обеспечивает:
· взаимодействие с пользователем (пользовательский интерфе й с );
17

1. Структура и возможности вычислительной системы с графическим процессором

· получение исходных данных от пользователя или из внешних
источников, сохранение результатов расчета;
· подготовку и запуск расчетов на графическом процессоре;
· исполнение участков расчетного алгоритма, предназначенных
для центрального процессора.
Пользовательский интерфейс. Установка и
отображение параметров моделирования
Формирование модельной физической
системы, подготовка исходных данных для
графического процессора

CPU,
программирование на
C#, С++ или другом
языке общего назначения с использованием
стандартных библиотек
DirectX или OpenCL

Проведение поточно-параллельных
расчетов на участках, критичных к
производительности моделирования

GPU,
программирование на
HLSL или другом
специальном языке

Исполнение участков алгоритма, плохо
распараллеливаемых или требующих
повышенной точности вычислений

CPU,
программирование на
C#, С++ или другом
языке общего назначения с использованием
DirectX или OpenСL

Обработка, сохранение и отображение
результатов моделирования

Рис. 1.6. Структура пользовательского приложения,
решающего задачу общего назначения (физического моделирования)
на графическом процессоре

Программы для центрального процессора обычно разрабатываются
на традиционных языках высокого уровня, в частности, на C++ или
C#, которые достаточно легко обращаются к процедурам из OpenCL
или DirectX.

18

1.3. Уровни управления графическим процессором и основные системы программирования GPU

1.3.5. Программно-аппаратная платформа NVIDIA CUDA
Средства программирования GPU, описываемые выше, примени‑
мы для программирования графических процессоров всех произво‑
дителей, в частности, компаний NVIDIA и AMD. Однако внутренние
архитектуры GPU разных производителей и поколений существенно
различаются между собой, а возможности использования особенно‑
стей каждой из архитектур универсальными системами программи‑
рования ограничены. Преимущества конкретных GPU доступны при
использовании специализированных инструментов программирова‑
ния, одним из которых является NVIDIA CUDA.
Платформа для параллельных вычислений NVIDIA CUDA
(Compute Unified Device Architecture — архитектура единого вы‑
числительного устройства) фактически представляет собой про‑
граммно-аппаратный комплекс, в котором физическое устройство
графического процессора и его программная модель соответствуют
подходам к организации параллельных вычислений, разрабатыва‑
емым NVIDIA.
CUDA позволяет программировать только совместимые GPU (про‑
изводства NVIDIA), однако позволяет управлять ими на более низком
уровне, чем DirectX и OpenCL. В частности, преимуществами CUDA
являются:
· доступ к программируемой разделяемой памяти;
· произвольная адресация при записи в память;
· значительно ускоренное взаимодействие CPU и GPU, некото‑
рые операции выполняются асинхронно.
Интерфейс программирования приложений реализован в форме
расширения языка C++. Ниже CUDA будет рассмотрена более под‑
робно.
Общая схема, иллюстрирующая взаимодействие приложений в со‑
ставе программного комплекса, для расчетов с использованием GPU
приведена на рис. 1.7.

19

1. Структура и возможности вычислительной системы с графическим процессором

Программа для центрального процессора: пользовательский интерфейс и плохо распараллеливаемые расчеты.
Языки высокого уровня общего назначения (C#, С++)
Программа для GPU с использованием специализированных
инструментов (HLSL, CUDA). Инициализируется
и запускается центральным процессором
Библиотеки API (DirectX, OpenCL, CUDA). Обеспечивают
и унифицируют взаимодействие пользовательских
приложений с драйвером GPU

Драйвер графического процессора. Управляет GPU
с учетом архитектуры и технологических особенностей
конкретного устройства
Графический процессор
Рис. 1.7. Иерархия приложений, составляющих программный комплекс
для расчетов на графических процессорах

1.3.6. Выбор платформы программирования GPU
На сегодня возможными платформами программирования GPU для
вычислений общего назначения являются NVIDIA CUDA, Microsoft
DirectX и OpenCL. Все платформы достаточно функциональны и ди‑
намично развиваются.
Преимуществами CUDA являются лучшая управляемость парал‑
лельных вычислений, а также использование обычного языка C++
с дополнительными библиотеками.
Достоинство платформы Microsoft DirectX (включая язык HLSL
и вычислительные шейдеры) — совместимость приложений с графи‑
ческими процессорами всех ведущих производителей (в частности,
AMD и NVIDIA), что актуально, поскольку по теоретической произ‑
водительности GPU AMD и NVIDIA примерно равны.
В ближайшие годы будет актуальной разработка приложений фи‑
зико-математического моделирования как для DirectX/OpenCL, так
и для CUDA.

20

1.4. Области применения графических процессоров

1.4. Области применения графических процессоров
Современные графические процессоры могут эффективно решать
очень многие задачи. Основное требование к таким задачам — возмож‑
ность распараллеливания по данным или по задачам, а также алго‑
ритм, в котором вычисления превалируют над условными переходами
и обменом данными между параллельными процессами. Вычисления
на графических процессорах уже сегодня применяются в следующий
областях:
· матричные преобразования;
· молекулярная динамика;
· астрофизические расчеты;
· глобальная оптимизация;
· дискретное преобразование Фурье;
· кодирование видео;
· визуализация (томография, нанотехнологии);
· нейронные сети;
· биофизические расчеты.
Наряду с симуляцией реальных процессов, GPU используются
и для визуализации (отображение результатов томографии, представ‑
ление сложных молекул и больших систем).

1.5. Необходимое аппаратное и программное обеспечение
Требования к архитектуре и производительности компьютера, на ко‑
торый устанавливается графический процессор, состоят в следующем.
Системная плата должна быть оборудована шиной данных PCI Express
либо AGP (устаревший и не лучший вариант) для подключения гра‑
фического ускорителя. Шина PCI Express обычна как для настольных
компьютеров, так и для ноутбуков.
Наиболее производительные графические ускорители потребля‑
ют сравнительно большую мощность, так что вместо блоков питания,
обычно поставляемых с корпусами настольных ПК, может потребовать‑
ся блок мощностью до 1000 Вт (как показала практика, такой блок необ‑
ходим, например, для видеокарты с GPU NVIDIA GeForce GTX 780 Ti).
21

1. Структура и возможности вычислительной системы с графическим процессором

Скорость вычислений на GPU мало зависит от производительно‑
сти центрального процессора, так что практически любой CPU мо‑
жет быть использован.
Объемы оперативной памяти персонального компьютера (RAM)
и видеопамяти графического ускорителя должны быть достаточными
для хранения обрабатываемых данных. Для моделирования больших
систем может потребоваться несколько гигабайтов памяти этих ти‑
пов. Связь между производительностью вычислений и объемами па‑
мяти практически полностью определяется пользовательским алго‑
ритмом моделирования.
Для расчетов общего назначения требуется графический ускори­
тель с процессором, поддерживающим шейдерную модель 4.0 или
старше. Для использования CUDA нужен графический ускоритель
NVIDIA с процессором G80 или новее. Драйвер графического процес‑
сора должен поддерживать выбранную платформу программирования
GPU. Для обеспечения поддержки наиболее новых версий платформы
может потребоваться обновление драйвера (даже когда графический
ускоритель нормально работает в других приложениях).
Для работы с платформой CUDA достаточно операционной си­
стемы Microsoft Windows XP Service Pack 2 (SP2). Для использования
DirectCompute и OpenCL (реализующих шейдерную модель 4.0) необ‑
ходимы операционные системы Microsoft Windows Vista и старше. Су‑
ществуют реализации OpenCL для Unix-подобных операционных си‑
стем, включая Linux и MacOS.
Специализированные библиотеки (API) включают в себя следую‑
щие пакеты.
Пакет DirectX 10, необходимый для программирования в рамках
шейдерной модели 4.0, входит в дистрибутив Windows Vista и после‑
дующих версий.
В примерах мы будем использовать процедуры из библиотеки
Microsoft XNA Game Studio 2.0 (XNA 2.0), которая является надстрой‑
кой над DirectX 9.0c. Эта библиотека была разработана для упроще‑
ния работы пользователей с DirectX и для упрощения доступа к DirectX
в рамках.Net Framework. Для ее использования требуется установка
дистрибутивов XNA 2.0 или 3.0, доступных бесплатно.
В соответствии со стратегией, которую предлагает и осуществляет
корпорация Microsoft, большинство приложений для Windows в насто‑
ящее время разрабатываются на базе программной оболочки Microsoft.
22

1.5. Необходимое аппаратное и программное обеспечение

Net Framework. Эта оболочка включает в себя библиотеку базовых клас‑
сов (Base Class Library), которая содержит реализации стандартных за‑
дач, выполняемых приложениями [15], а также компонент Common
Language Runtime [16], который исполняет пользовательские прило‑
жения. Приложения для .Net Framework компилируются не в машин‑
ные коды, а в промежуточный байт-код (Common IntermediateLanguage,
CIL), который окончательно компилируется для конкретной версии
операционной системы и конкретного процессора уже перед самым
исполнением с помощью JIT-компилятора (Jast In Time compiler), тоже
входящего в состав .Net Framework.
Среда .Net Framework необходима для работы с Microsoft XNA
Game Studio. Компиляторы, встроенные в среду программирования
Microsoft Visual Studio (см. пункт «Средства программирования» ниже),
тоже по умолчанию создают именно CIL-код, а не машинные коды.
Оболочка .Net Framework автоматически устанавливается вместе
с Windows, при необходимости ее новые версии (доступные бесплат‑
но) можно устанавливать вручную. Для использования XNA 2.0 до‑
статочно версии .Net Framework 2.0, поставляемой, например, вместе
с Microsoft Visual Studio 2005.
Для выполнения программ, написанных на CUDA, требуется
установка библиотек, формирующих эту платформу: CUDA Toolkit
и CUDA SDK. Дистрибутивы доступны бесплатно (на сайте NVIDIA).
Версии CUDA регулярно обновляются.
Средства программирования
Центральный процессор
При использовании любой платформы программирования GPU
часть кода исполняется на центральном процессоре. Средой програм‑
мирования CPU для Windows может быть Microsoft Visual Studio. Эта
среда хорошо подходит для написания приложений с вычисления‑
ми на графических процессорах, при использовании DirectX/XNA
и CUDA. Библиотеки XNA 2.0 и 3.0 автоматически встраиваются в сре‑
ду Visual Studio при своей установке, как и компилятор CUDA.
Для обращения к XNA можно использовать такие языки про‑
граммирования, как C++, C# и Visual Basic. Процедуры, написан‑
ные на этих языках, могут входить и в состав приложений, исполь‑
23

1. Структура и возможности вычислительной системы с графическим процессором

зующих CUDA. Собственно CUDA представляет собой расширение
C++, для компилирования процедур на CUDА существует специаль‑
ный компилятор.
Вместо полных версий Visual Studio вполне возможно использова‑
ние Microsoft Visual C# и С++ Express.
DirectX/XNA/HLSL
Обращение к процедурам DirectX можно осуществлять посредством
Microsoft XNA. При использовании технологии .Net Framework библи‑
отека XNA стандартным образом подключается к программам на C#
или C++. В результате процедуры XNA становятся доступными ана‑
логично процедурам из любых других классов.
Шейдеры на языке HLSL хранятся в отдельных текстовых файлах.
Для их написания можно использовать любой текстовый редактор. Су‑
ществуют специализированные среды программирования для напи‑
сания шейдеров: ATI Render Monkey (разрабатывается AMD/ATI), FX
Composer (разрабатывается NVIDIA), Shader Maker (OpenGL, OpenСL).
NVIDIA CUDA
Программу на CUDA можно редактировать в любой среде, напри‑
мер, в Microsoft Visual Studio. Для компилирования программы проще
всего использовать компилятор CUDA (nvcc.exe [17]), который входит
в состав CUDA Toolkit. Если при компиляции создать не исполняе‑
мый файл (.exe), а динамическую библиотеку (.dll), то она затем мо‑
жет быть подключена к проекту, разрабатываемому на любом из рас‑
пространенных языков высокого уровня для CPU.

24

2. Поточно-параллельное
программирование GPU

2.1. Распараллеливание расчетов

П

усть имеется вычислительная задача,
предполагающая обработку некоторо‑
го массива из N однотипных элемен‑
тов (исходных данных, моделируемых объектов,
случайных событий), по одному и тому же алго‑
ритму, включающему m операций. Может требо‑
ваться исполнение нескольких таких задач, тогда
блок-схема последовательных вычислений, без
распараллеливания, будет иметь общий вид, по‑
казанный на рис. 2.1. Из этой блок-схемы вид‑
но, что возможно три следующих подхода к рас‑
параллеливанию расчетов [18].
2.1.1. Распараллеливание по задачам
Распараллеливание по задачам возможно,
если задачи (1 — k на рис. 2.1) независимы друг
от друга. Такое распараллеливание актуально
для сетевых серверов и других вычислительных
систем, выполняющих одновременно несколько
функций либо обслуживающих многих пользо‑
вателей. Оно может быть востребовано и в фи‑
зическом моделировании: например, для ис‑
следования одной и той же системы при разных

Задача 1
i=0
Операция 1
над элементом i
ј
Операция m
над элементом i
i=i+1
Да

i= 1)
{
__syncthreads ();
for (int a = threadIdx.x; a < stride; a += blockDim.x)
accumulator [a] += accumulator [stride + a];
}
if (threadIdx.x == 0) C [v] = accumulator [0];
}
}

Компиляция программ на CUDA
Опции компилятора nvcc
Nvcc — инструмент для управления всеми этапами компиляции, на‑
зываемый также драйвером компиляции (CUDA compiler driver). Ком‑
пиляция при помощи nvcc происходит в несколько этапов (phases).
Выбор этих фаз производится на основе опций, передаваемых через
командную строку, и расширений входных файлов, причем на вы‑
полнение этих фаз можно влиять при помощи других опций команд‑
ной строки.
В процессе своей работы nvcc вызывает следующие инструменты:
· CUDAFE (CUDA Language Front End);
· NVOPENCC (CUDA Open64 Compiler);
· PTXAS (PTX Optimizing Assembler);
· хостовые компилятор и линковщик (Microsoft Visual C++ для
ОС Windows или gcc для ОС Linux).
88

Компиляция программ на CUDA

Компилятор поддерживает следующие расширения файлов:
· .cu — исходные файлы CUDA, содержащие код для хоста и для
устройств;
· .cup — исходные файлы CUDA после предварительной обра‑
ботки (preprocessed);
· .cc,.cxx,.cpp — исходные файлы C++;
· .gpu — промежуточный файл для GPU;
· .ptx — файл с промежуточным ассемблерным PTX-кодом;
· .o,.obj — объектный файл;
· .a,.lib — библиотечный файл;
· .res — файл с ресурсами;
· .so — разделяемый объектный файл.
Nvcc не различает объектные, библиотечные и ресурсные файлы,
а сразу пропускает их к линковщику на соответствующей фазе. В от‑
личие от gcc, nvcc при получении на входе файла с незнакомым рас‑
ширением (то есть не перечисленном выше) генерирует ошибку.
Nvcc различает 3 варианта опций командной строки: булевы (фла‑
ги) — без аргументов, опции с одним аргументом и опции с нескольки‑
ми аргументами (списки). Аргументы отделяются от названия опции
символом равенства либо одним или более пробелами. В некоторых
случаях совместимости с компилятором gcc (опции –I, — l, — L) аргу‑
мент может не отделяться от названия опции. Несколько аргументов
одной опции могут разделяться запятыми, допускается повтор таких
опций, возможна также комбинация этих двух вариантов.
Каждая опция имеет длинное и сокращенное имя, которые от‑
личаются количеством дефисов: два для длинного имени и один для
короткого. Длинные имена, как правило, используются в сценариях
автоматической компиляции, где длина менее важна, чем ценность
самоописания.
Список опций, определяющих фазу компиляции
Имя
Описание
длинное
короткое
—cuda
-cuda
Компиляция.cu →.cu.c/.cu.cpp
Компиляция.cu/.gpu/.ptx →.cubin.
—cubin
-cubin
Хостовый код отбрасывается
Компиляция.cu/.gpu →.ptx.
—ptc
-ptx
Хостовый код отбрасывается

89

Приложение

Имя
длинное
короткое
—gpu

-gpu

—preprocess
—generatedependencies
—compile

-E
-M

—link

-link

—lib

-lib

—run

-run

-c

Описание
Компиляция.cu →.gpu
Хостовый код отбрасывается
Предварительная обработка.c/.cc/.cpp/.cxx/.cu.
Генерация для одного.c/.cc/.cpp/.cxx/.cu файла зави‑
симостей, который может быть включен в make-файл
Компиляция.c/.cc/.cpp/.cxx/.cu →.obj
Поведение по умолчанию: компиляция и линковка
всех файлов
Компиляция всех входных файлов в объектные и объ‑
единение результатов в один библиотечный файл
Компиляция, линковка и запуск. При этом путь
к необходимым динамическим библиотекам CUDA
берется из файла nvcc.profile

Список опций, указывающих файлы и пути к файлам
Имя
длинное
—output-file

—pre-incude
, …
—library , …
—define-macro
, …
—undefine-macro
, …
—include-path
, …
—system-include
, …
—library-path
, …

короткое
-o

Выходной файл. Более одного аргумента допуска‑
ется только в режиме линковки/архивации

-include Заголовочные файлы
-l

Библиотечные файлы. Путь для их поиска указыва‑
ется опцией ‘-L’

-D

Определение макросов

-U

Отмена уже определенных макросов

-I

Пути поиска заголовочных файлов

-isystem Пути поиска системных заголовочных файлов
-L

—output-directory


-odir

—compiler-bindir


-ccbin

90

Описание

Пути поиска библиотечных файлов
Путь для сохранения выходного файла. Полезна
для генерации зависимостей через ‘—generate-de‑
pendencies’
Путь к хостовому компилятору:
Microsoft Visual Studio cl или gcc

Компиляция программ на CUDA

Список опций для изменения поведения компилятора/линковщика
длинное

Имя

короткое

Описание

—profile

-pg

—debug
—device-debug
—optimize


-g
-G

Инструментирование кода для профилировщика
gprof (только в ОС Linux)
Генерация отлаживаемого кода
Генерация отлаживаемого кода устройства

-O

Генерация оптимизированного кода

—shared

-shared

—machine

-m

Генерация разделяемой библиотеки во время лин‑
ковки. Если требуются другие опции линковщика,
используйте опцию ‘-Xlinker’
Выбор 32‑битной или 64‑битной архитектуры.
В данный момент она нужна только на платформе
linux64 при компиляции с опцией ‘-cubin’

Список опций, управляющих вызываемыми nvcc инструментами
Имя
Описание
длинное
короткое
—compiler-options
-Xcompiler Передача опций компилятору/препроцессору
,..
—linker-options
-Xlinker
Передача опций линковщику
, …
—cudafe-options
-Xcudafe Передача опций cudafe
, …
Передача опций nvopencc, обычно для управ‑
—opencc-options
ления оптимизацией. К примеру, ‘‑LIST:
-Xopencc
, …
source=on’ приводит к включению в ptx-код
комментариев с исходным кодом
Передача опций оптимизирующему PTX—ptxas-options
-Xptxas
ассемблеру. К примеру, ‘-v’ приводит к выводу
, …
статистики генерации кода
Список опций, управляющих поведением nvcc
длинное

Имя

короткое

—dryrun

-dryrun

—verbose

-v

Описание
Перечислить команды компиляции, генерируе‑
мые nvcc, вместо их выполнения
Перечислить команды компиляции, генерируе‑
мые nvcc, вместе с выполнением

91

Приложение

Имя
Описание
длинное
короткое
—keep
-keep
Сохранить все промежуточные файлы
—save-temps
-save-temps То же, что и ‘— keep’
—dont-use-profile
-noprof
Не использовать файл настроек nvcc.profile
Вместо каждой фазы компиляции удалить фай‑
—clean-targets
-clean
лы, которые бы остались после нее
—run-args ,
Используется вместе с опцией ‘-R’ для передачи
-run-args
…
аргументов исполняемому файлу
В ОС Windows: сконвертировать все аргументы,
—input-driveвключающие имена файлов, в родной формат Win‑
-idp
prefix
dows. Это относится к представлению абсолютных
путей к файлам в «текущей» среде разработки
В ОС Windows при генерации файлов с зависи‑
—dependencyмостями (опция ‘-M’) все имена файлов конвер‑
drive-prefix
-ddp
тируются в такие, которые поддерживаются ис‑

пользуемой утилитой make. К примеру, некоторые
варианты make не поддерживают символ двоеточия
Задает префикс логического диска для обеих оп‑
—drive-prefix
-dp
ций: input-drive-prefix и dependency-drive-prefix

Список опций, управляющих компиляцией CUDA
длинное

Имя

короткое

Описание

Сгенерировать код для библиотеки эмуляции
GPGPU. В релизе CUDA 3.0 режим эмуляции
стал считаться устаревшим
Использование библиотеки быстрой арифме‑
—use_fast_math -use_fast_math
тики
Управление денормализацией чисел с одинар‑
—ftz
-ftz
ной точностью (false означает поддержку де‑
нормализованных чисел, true — их обнуление)
Управление точностью 32‑битного деле‑
—prec-div
-prec-div
ния (true — для соответствия стандарту IEEE,

false — использование аппроксимации)
Управление точностью 32‑битного квадратного
—prec-sqrt
-prec-sqrt
корня (true — для соответствия стандарту IEEE,

false — использование аппроксимации)
В случае компиляции.ptx/.gpu →.cubin: опреде‑
—entries ,
лить функции, являющиеся глобальными точ‑
-e
…
ками входа, для которых будет сгенерирован
код. По умолчанию код генерируется для всех
—deviceemulation

92

-deviceemu

Компиляция программ на CUDA

Список опций, управляющих генерацией GPU-кода
Имя
длинное

корот‑
кое

Описание

Задает архитектуру GPU — реальную или вир‑
туальную (PTX) — для всех этапов компи‑
ляции вплоть до генерации PTX-кода. Под‑
— gpu-architecture
-arch
держиваются виртуальные архитектуры:

compute_10, compute_11, compute_12, com‑
pute_13; и реальные: sm_10, sm_11, sm_12,
sm_13, — которые их реализуют
Влияет на генерацию бинарного кода, кото‑
рый задействуется только в случае его совме‑
стимости с используемым GPU, иначе проис‑
ходит компиляция встроенного PTX-кода.
— gpu-code ,
Указанная архитектура должна быть совме‑
-code
стима с указанной в опции ‘-arch’. Если опция
…
‘-arch’ указана, то здесь должна быть выбрана
виртуальная PTX-архитектура.
Если опция ‘-code’ не указана, то берется зна‑
чение из опции ‘-arch’
Позволяет многократно вызывать компилятор
— generate-code’‑arch=
nvopencc, генерируя код для нескольких вир‑
‑code=, -gencode
туальных архитектур. Аргументы ‘arch’ и ‘code’
…’
соответствуют двум предыдущим опциям
Указывает директорию или zip-файл, в ко‑
торые будут скопированы все сгенерирован‑
ные образы кода. Этот файл будет играть роль
репозитория кода, инспектируемого драйве‑
ром CUDA во время выполнения программы.
В обоих случаях nvcc сохранит структуру ди‑
— export-dir
-dir
ректорий для помощи драйверу CUDA при
поиске.
Если эта опция не указана, то все внешние
образы будут отброшены. Если указано имя
несуществующего файла, то будет создана ди‑
ректория с таким именем
Определяет, какие образы кода будут скопи‑
рованы в директорию, указанную в опции ‘ex‑
— extern-mode
-ext
port-dir’.

Если эта опция и ‘intern-mode’ не указаны, то все
образы кода будут считаться внутренними

93

Приложение

Имя
длинное
— intern-mode


— maxrregcount


корот‑
кое

Описание

Определяет, какие образы кода будут встроены
в выходной (например, исполняемый) файл.
-int
Если эта опция и ‘extern-mode’ не указаны,
то все образы кода будут считаться внутренними
Определяет максимальное число регистров,
которые может использовать один поток GPU.
Из-за ограниченных ресурсов мультипроцес‑
сора GPU слишком большое значение приве‑
дет к ограничению числа потоков в связке. Од‑
нако при недостаточно большом для какого-то
вычислительного ядра значении его произво‑
дительность будет ниже в том числе за счет ис‑
-maxr‑
regcount пользования вместо регистров медленной ло‑
кальной памяти GPU.
Указанное значение округляется в большую
сторону до кратного четырем, вплоть до мак‑
симума в 128 регистров.
Если опция не указана, то ограничение
не применяется (только физические ограниче‑
ния)

Профайлер
В инструментарий CUDA Toolkit входит профайлер, который
во время выполнения вычислительного ядра и операций копирова‑
ния собирает информацию об их длительности. Это позволяет нахо‑
дить узкие места в программах и численно оценить выгоду от оптими‑
зации отдельного ядра. Собирание этой информации контролируется
следующими переменными среды:
· CUDA_PROFILE — значения 1 или 0 включают и отключают
сбор информации;
· CUDA_PROFILE_LOG — указание имени файла и пути к нему,
в который будет произведена запись собранной информации.
Если путь не указан, то запись производится в файл “cuda_pro‑
file.log” в текущей директории. При запуске на нескольких GPU
можно записывать данные для каждого GPU в отдельный файл,
94

Профайлер

если использовать при указании имени файла переменную %d,
обозначающую номер устройства;
· CUDA_PROFILE_CSV — включение и отключение формата за‑
писанной информации с разделением столбцов запятыми (Com‑
ma-Separated Values);
· CUDA_PROFILE_CONFIG — использование файла с конфи‑
гурацией счетчиков производительности GPU.
В GPU встроены следующие счетчики производительности:
timestamp — моменты времени запуска ядер и операций копирования;
· gpustarttimestamp — время начала выполнения ядра на GPU;
· gpuendtimestamp — время завершения выполнения ядра на GPU;
· gridsize — число связок потоков в сетке по всем измерениям X и Y;
· threadblocksize — число потоков в связке по всем измерениям X,
Y и Z;
· dynsmemperblock — размер динамически выделяемой разделя‑
емой памяти в байтах у каждой связки потоков;
· statsmemperblock — размер статически выделяемой разделяемой
памяти в байтах у каждой связки потоков;
· regperthread — число регистров на один поток;
· memtransferdir — направление копирования памяти — 0 — для
операций хост→устройство, 1 — для операций устройство→хост;
· memtransfersize — объем копируемого блока данных;
· memtransferhostmemtype — тип памяти хоста, участвующий
в операции копирования: со страничной организацией (page‑
able) или с ее блокировкой (page-locked);
· streamid — идентификатор потока при запуске ядра.
Для детализации процесса выполнения ядра поддерживаются сле‑
дующие счетчики:
· local_load — число инструкций чтения из локальной памяти
на 1 варп мультипроцессора;
· local_store — число инструкций записи в локальную память
на 1 варп мультипроцессора;
· gld_request — число инструкций чтения из глобальной памяти
на 1 варп мультипроцессора;
· gst_request — число инструкций чтения из глобальной памяти
на 1 варп мультипроцессора;
· divergent_branch — число уникальных ветвлений, которые про‑
являются;
95

Приложение

· branch — число уникальных инструкций ветвления в програм‑
ме;
· sm_cta_launched — число связок потоков, выполняемых на муль‑
типроцессоре.
На устройствах с вычислительной способностью 1.x имеются сле‑
дующие счетчики для процесса выполнения ядра:
· gld_incoherent — число необъединенных (некогерентных) опе‑
раций чтения из глобальной памяти;
· gld_coherent — число объединенных операций чтения из гло‑
бальной памяти;
· gld_32b — число 32‑байтных транзакций чтения из глобальной
памяти;
· gld_64b — то же, 64‑байтных;
· gld_128b — то же, 128‑байтных;
· gst_incoherent — число необъединенных (некогерентных) опе‑
раций записи в глобальную память;
· gst_coherent — число объединенных операций записи в глобаль‑
ную память;
· gst_32b — число 32‑байтных транзакций записи в глобальную
память;
· gst_64b — то же, 64‑байтных;
· gst_128b — то же, 128‑байтных;
· instructions — число выполненных инструкций;
· warp_serialize — число потоков варпа, которым приходится се‑
риализовать операции доступа к разделяемой или константной
памяти;
· cta_launched — число выполненных связок потоков;
· prof_trigger_00, …, prof_trigger_07 — триггеры профайлера
(до 8 штук), которые можно включать в пользовательском коде
ядра через вызов встроенной функции __prof_trigger (int N), где
N — индекс триггера;
· tex_cache_hit — число попаданий текстурного кеша;
· tex_cache_miss — число промахов текстурного кеша.
На устройствах с вычислительной способностью 2.0 и более высо‑
кой возможен одновременный сбор информации по большему чис‑
лу счетчиков производительности в зависимости от типа выбранных
счетчиков, кроме того, добавлены следующие счетчики:
96

Профайлер

· shared_load — число инструкций чтения из разделяемой памя‑
ти на 1 варп мультипроцессора;
· shared_store — число инструкций записи в разделяемую память
на 1 варп мультипроцессора;
· inst_issued — число выполненных инструкций, включая реплеи
(replays);
· inst_executed — число выполненных инструкций, исключая ре‑
плеи;
· warps_launched — число запущенных варпов на мультипроцес‑
соре;
· threads_launched — число запущенных потоков на мультипро‑
цессоре;
· l1_global_load_hit — число попаданий кеша L1 при чтении;
· l1_global_load_miss — число промахов кеша L1 при чтении.

97

Библиографический список

1. ATI Technologies [Электронный ресурс]. — Режим досту‑
па: http://en.wikipedia.org/wiki/ATI_Technologies. — Заглавие
с экрана.
2. Nvidia [Электронный ресурс]. — Режим доступа:
http://en.wikipedia.org/wiki/Nvidia.– Заглавие с экрана.
3. S3 Graphics [Электронный ресурс]. — Режим доступа:
http://en.wikipedia.org/wiki/S3_Graphics. — Заглавие с экрана.
4. NVIDIA Corp. NVIDIA CUDA Compute Unified Device Ar‑
chitecture. Programming Guide Version 1.0 [Электронный ре‑
сурс]/NVIDIA Corp. 2701 San Tomas Expressway, Santa Clara,
CA 95050. — Режим доступа: http://www.nvidia.com. — Заглавие
с экрана.
5. NVIDIA Corp. CUDA Technical Training [Электронный ре‑
сурс]/NVIDIA Corp. 2701 San Tomas Expressway, Santa Clara,
CA 95050. — Режим доступа: http://www.nvidia.com. — Заглавие
с экрана.
6. CUDA C Programming Guide [Электронный ресурс]. — Ре‑
жим доступа: http://docs.nvidia.com/cuda/cuda-c‑programmingguide/#axzz3k3HvOyd7. — Заглавие с экрана.
7. Comparison of ATI graphics processing units [Электронный ре‑
сурс]. — Режим доступа: http://en.wikipedia.org/wiki/Comparison_
of_ATI_graphics_processing_units. — Заглавие с экрана.
8. Comparison of Nvidia graphics processing units [Электронный ре‑
сурс]. — Режим доступа: http://en.wikipedia.org/wiki/Comparison_
of_Nvidia_graphics_processing_units. — Заглавие с экрана.
9. Графический конвейер [Электронный ресурс]. — Режим досту‑
па:/http://ru.wikipedia.org/wiki. — Заглавие с экрана.
10. Shaders With Ogre [Электронный ресурс]. — Режим до‑
ступа: http://www.ogre3d.ru/wik/pmwiki.php?n=Main.
ShadersWithOgre. — Заглавие с экрана.

98

11. Reference for HLSL [Электронный ресурс]. — Режим досту‑
па: https://msdn.microsoft.com/ru-ru/library/windows/desk‑
top/bb509638 (v=vs.85).aspx. — Заглавие с экрана.
12. Медведев, A. Современная терминология 3D графики [Элек‑
тронный ресурс]/А. Медведев. Режим доступа: http://www.ixbt.
com/video2/terms2k5.shtml.
13. DirectCompute [Электронный ресурс]. — Режим доступа:
https://en.wikipedia.org/wiki/DirectCompute. — Заглавие с экрана.
14. OpenGL 3 против DirectX 11: война закончена [Электронный
ресурс]. — Режим доступа: http://www.thg.ru/graphic/open_gl_3_
directx_11/index.html. — Заглавие с экрана.
15. Net Framework [Электронный ресурс]. — Режим доступа:
http://ru.wikipedia.org/wiki/.NET_Framework. — Заглавие с экрана.
16. Common Language Runtime [Электронный ресурс]. — Режим
доступа: http://ru.wikipedia.org/wiki/Common_Language_Run‑
time. — Заглавие с экрана.
17. NVIDIA CUDA Compiler Driver NVCC [Электронный ресурс]. —
Режим доступа: http://docs.nvidia.com/cuda/cuda-compilerdriver-nvcc/#axzz3k3HvOyd7. — Заглавие с экрана.
18. Немнюгин, С. А. Параллельное программирование для мно‑
гопроцессорных вычислительных систем/С. А. Немнюгин,
О. Л. Стесик. — СПб. : БХВ‑Петербург, 2002. — 400 с.
19. Обзор видеокарты NVIDIA GeForce GTX 280 — очень бы‑
стро и жутко горячо [Электронный ресурс]. — Режим доступа:
http://tfile.ru/forum/viewtopic.php?t=191087. — Заглавие с экрана.
20. GPU-accelerated libraries [Электронный ресурс]. — Режим до‑
ступа: https://developer.nvidia.com/gpu-accelerated-libraries. —
Заглавие с экрана.
21. Arora, N. Direct N‑body Kernels for Multicore Platforms/N. Arora,
A. Shringarpure, R. W. Vuduc//Proceedings of the International
Conference on Parallel Processing. — 2009. — P. 379–387.
22. Satish, N. Designing efficient sorting algorithms for manycore
GPUs/N. Satish, M. Harris, M. Garland//Proceeding IPDPS
‘09 Proceedings of the 2009 IEEE International Symposium on
Parallel&Distributed Processing. — 2009. — P. 1–10.

99

Оглавление

Введение........................................................................................................... 3
1. Структура и возможности вычислительной системы
с графическим процессором.............................................................................. 4
1.1. Задача компьютерной визуализации трехмерных сцен....................... 4
1.2. Архитектура графического процессора (GPU).................................... 5
1.2.1. Распараллеливание вычислений по данным.............................. 5
1.2.2. Взаимодействие графического и центрального процессоров.... 8
1.2.3. Иерархия памяти, доступной центральному
и графическому процессорам...................................................... 9
1.2.4. Конвейерная обработка данных................................................ 12
1.3. Уровни управления графическим процессором и основные
системы программирования GPU..................................................... 14
1.3.1. Уровни управления графическим процессором....................... 14
1.3.2. Драйвер графического процессора............................................ 14
1.3.3. Интерфейсы программирования приложений......................... 15
1.3.4. Пользовательское приложение.................................................. 17
1.3.5. Программно-аппаратная платформа NVIDIA CUDA.............. 19
1.3.6. Выбор платформы программирования GPU............................ 20
1.4. Области применения графических процессоров............................... 21
1.5. Необходимое аппаратное и программное обеспечение.................... 21
Средства программирования.............................................................. 23
2. Поточно-параллельное программирование GPU......................................... 25
2.1. Распараллеливание расчетов.............................................................. 25
2.1.1. Распараллеливание по задачам.................................................. 25
2.1.2. Распараллеливание по инструкциям......................................... 26
2.1.3. Распараллеливание по данным.................................................. 27
2.2. Преимущества графических процессоров при параллельных
расчетах............................................................................................... 28
2.3. Принцип программирования SIMD на примере пиксельного
шейдера............................................................................................... 30
2.4. Пример сложения матриц................................................................... 32
2.4.1. Распараллеливание независимых вычислений......................... 32
2.4.2. Сложение матриц в рамках шейдерной модели 3.0.................. 34
2.4.3. Структура программы для центрального процессора............... 36
2.4.4. Реализация программы для центрального процессора на C#.... 37

100

Содержание

2.4.5. Программа для графического процессора................................ 43
2.4.6. Вычислительные шейдеры модели 5.0...................................... 49
3. Программирование графических процессоров на CUDA............................. 50
3.1. Модель программирования графических процессоров
как универсальных вычислительных систем..................................... 50
3.1.1. Взаимодействие параллельных вычислительных процессов... 50
3.1.2. Концепция универсального вычислительного
устройства CUDA...................................................................... 51
3.1.3. Иерархия вычислительных процессов и памяти CUDA.......... 53
3.1.4. Возможности и ограничения процессоров
архитектуры CUDA.................................................................... 55
3.1.5. Конвейерная обработка данных в архитектуре CUDA............. 56
3.2. Особенности программирования на CUDA...................................... 56
3.2.1. Идентификация вычислительного потока................................ 56
3.2.2. Совместимость с шейдерными моделями................................. 58
3.2.3. Язык программирования CUDA............................................... 58
3.2.4. Структура программы на CUDA................................................ 59
3.3. Анализ алгоритма параллельного перемножения матриц................. 62
3.3.1. Алгоритм перемножения матриц.............................................. 62
3.3.2. Процедура перемножения матриц на CUDA............................ 64
3.3.3. Оптимизация доступа к памяти при умножении матриц......... 69
3.4. Динамика N тел на CUDA. Пример ускорения программы
за счет скорости GPU......................................................................... 73
3.5. Распараллеливание алгоритмов сортировки. Пример ускорения
программы за счет скорости GPU...................................................... 76
Заключение.................................................................................................... 79
Приложение................................................................................................... 80
Перемножение матриц на CUDA. Программа, исполняемая
центральным процессором..................................................................... 80
Вычисление скалярного произведения векторов на CUDA.................. 84
Компиляция программ на CUDA........................................................... 88
Профайлер............................................................................................... 94
Библиографический список.......................................................................... 98

101

Для заметок

102

Для заметок

103

Учебное издание

Некрасов Кирилл Александрович,
Поташников Святослав Игоревич,
Боярченков Антон Сергеевич,
Купряжкин Анатолий Яковлевич

ПАРАЛЛЕЛЬНЫЕ ВЫЧИСЛЕНИЯ
ОБЩЕГО НАЗНАЧЕНИЯ
НА ГРАФИЧЕСКИХ ПРОЦЕССОРАХ

Редактор И. В. Меркурьева
Верстка О. П. Игнатьевой
Подписано в печать 06.05.2016. Формат 70×100/16.
Бумага писчая. Печать цифровая. Гарнитура Newton.
Уч.-изд. л. 5,2. Усл. печ. л. 8,4. Тираж 50 экз.
Заказ 123
Издательство Уральского университета
Редакционно-издательский отдел ИПЦ УрФУ
620049, Екатеринбург, ул. С. Ковалевской, 5
Тел.: 8(343)375-48-25, 375-46-85, 374-19-41
E-mail: rio@urfu.ru
Отпечатано в Издательско-полиграфическом центре УрФУ
620075, Екатеринбург, ул. Тургенева, 4
Тел.: 8(343) 350-56-64, 350-90-13
Факс: 8(343) 358-93-06
E-mail: press-urfu@mail.ru

I SBN 579961722 - 3

9 785799 617226