CC BY
53
УДК 004.4'242, 004.4'422, 004.432.4
АВТОМАТИЧЕСКАЯ ГЕНЕРАЦИЯ ПРОГРАММ ДЛЯ ГРАФИЧЕСКИХ ПРОЦЕССОРОВ ПО НЕПРОЦЕДУРНЫМ СПЕЦИФИКАЦИЯМ1
А.Н. Андрианов, А.Б. Бугеря, Е.Н. Гладкова, К.Н. Ефимкин, П.И. Колударов
Рассмотрены методы и средства автоматического построения параллельных программ для графических процессоров по непроцедурной спецификации задачи, заданной на языке Норма. Приводятся алгоритмы автоматического построения программ для графических процессоров, основанные на анализе зависимостей и выявлении внутреннего параллелизма графа потока управления, полученного по программе на языке Норма. Приведены результаты применения описанных алгоритмов для создания расчетной программы с использованием технологии CUDA при решении задачи из области газодинамики и оценена производительность получающихся параллельных программ.
Ключевые слова: суперкомпьютеры, параллельное программирование, графические процессоры, CUDA, Норма, непроцедурные спецификации.
Введение
В современном научном сообществе задача разработки эффективных параллельных программ имеет важное стратегическое значение. Наверное, уже не осталось ни одной области науки или отрасли промышленности, так или иначе не связанной со сферой высокопроизводительных вычислений. Несмотря на то, что параллельное программирование появилось уже достаточно давно, успешно развивается и проводится масса исследований на эту тему, вопрос «как создать эффективную параллельную программу для решения такой-то задачи» до сих пор крайне актуален для программистов и математиков.
Достаточно быстрое развитие новых аппаратных возможностей для поддержки параллельных вычислений, наблюдаемое в последнее время, еще больше усложняет проблему. Например, появление массово доступных многоядерных процессоров поставило вопрос об эффективном программировании для них. Практически одновременно появились массово доступные графические ускорители (графические процессоры), и опять возник вопрос об эффективном программировании для них. Агрессивное продвижение своих решений фирмами-производителями вычислительных систем, обладающих этими возможностями, часто дезориентирует прикладных специалистов, разрабатывающих параллельные вычислительные программы, толкает их на изменение средств разработки программ, хотя ясные и достаточно убедительные аргументы в пользу таких изменений отсутствуют. Так, применение технологии CUDA для эффективного программирования для графических ускорителей в первых своих версиях являлось, фактически, программированием на уровне ассемблера, с учетом тонких особенностей аппаратуры.
1 Статья рекомендована к публикации программным комитетом Международной суперкомпью-терной конференции «Научный сервис в сети Интернет: все грани параллелизма 2013».
С помощью такого ручного низкоуровневого программирования за последние годы некоторые расчетные прикладные пакеты и математические библиотеки были портиро-ваны для использования на вычислительных системах с графическими процессорами [1]. Это, несомненно, существенно облегчает задачу прикладному специалисту, но только в том случае, если всего его потребности в высокопроизводительных вычислениях покрываются имеющимися распараллеленным пакетом и/или библиотеками. Если же с помощью таких готовых средств построить решение для своей задачи не удается, то иного пути, кроме как изучать CUDA или ATI Stream (или OpenCL - это уже более высокий уровень абстракции, но и эффективность его реализации пока далека от желаемого уровня) и начать реализовывать свой алгоритм на столь низком уровне, у прикладного специалиста нет.
Надежды на автоматическое распараллеливание уже написанных последовательных программ на графические процессоры пока совершенно не оправдываются, несмотря на то, что фирмы-производители графических процессоров, как NVIDIA, так и AMD, активно поддерживают данное направление исследований. Из уже реализованных подходов можно отметить те, которые базируются на вполне разумном симбиозе распараллеливающего компилятора и подсказок со стороны программиста, выполненных в виде специальных программных директив, например OpenACC [2]. Но, к сожалению, производительность получаемого таким путем программного кода далека от производительности программ, написанных ручным программированием для целевой платформы, и поэтому пока не может быть признанной удовлетворительной.
Работы по созданию и продвижению новых средств и языков программирования для графических процессоров, гибридных решений и различных нетрадиционных вычислительных архитектур (например, FPGA), ведутся весьма активно [3, 4], однако проблема простой разработки параллельных программ и утилизации новых возможностей вычислительной техники так и остается в настоящее время не решенной.
В данной статье рассмотрены методы и средства автоматического построения параллельных программ для графических процессоров по непроцедурной спецификации на примере решения этой задачи для спецификаций, заданных на языке Норма. В разделе 1 вкратце рассказывается о языке Норма и его основных свойствах. Во втором разделе сформулированы методы автоматического построения программы с использованием технологии CUDA по непроцедурным спецификациям, заданным на языке Норма. В разделе 3 приведен пример построения такой программы в соответствии с предложенными методами. В разделе 4 рассказано об опыте применения предложенных методов для создания расчетной программы с использованием технологии CUDA при решении задачи из области газодинамики и оценена производительность получающихся параллельных программ. В заключении говорится о продолжающихся работах коллектива авторов по данной тематике.
1. Декларативный подход. Язык Норма
Один из возможных подходов к решению задачи автоматизации параллельного программирования вычислительных задач, и, в частности, задачи автоматического построения эффективной программы для графических процессоров, является подход с использованием декларативных (непроцедурных) языков. При использовании этого подхода прикладной специалист программирует решение вычислительной задачи на непроцедур-
ном языке (понятия, связанные с архитектурой параллельного компьютера, моделями параллелизма и др. при этом не используются), а затем компилятор автоматически строит параллельную программу (учитывая архитектуру целевого параллельного компьютера, модели параллелизма и прочее.). С учетом отмеченных выше проблем привлекательность этого подхода в настоящее время только усиливается и интерес к идеям непроцедурного декларативного программирования и реализации этих идей в языках программирования неуклонно растет.
Идеи декларативного программирования были сформулированы еще в прошлом веке, теоретические исследования этого подхода для класса вычислительных задач проведены в пионерских работах И.Б. Задыхайло [5] еще в 1963 г. Непроцедурный язык Норма и система программирования Норма [6-8] разработаны в ИПМ им. М.В. Келдыша РАН также достаточно давно и предназначены для автоматизации решения вычислительных сеточных задач на параллельных компьютерах. Расчетные формулы записываются на языке Норма в математическом, привычном для прикладного специалиста виде. Язык Норма позволяет описывать решение широкого класса задач математической физики. Программа на языке Норма имеет очень высокий уровень абстракции и отражает метод решения, а не его реализацию при конкретных условиях. Такое описание не ориентировано на конкретную архитектуру компьютера, поэтому оно предоставляет большие возможности для выявления естественного параллелизма и организации вычислений.
В настоящее время в ИПМ им. М.В. Келдыша ведутся работы по созданию версии компилятора программ на языке Норма+, который на выходе создает исполняемую программу для графических процессоров фирмы М"УГО1А с использованием технологии CUDA.
2. Принципы построения CUDA программы по декларативным описаниям
При трансляции с языка Норма+ решается задача синтеза выходной параллельной программы, то есть выходная параллельная программа строится автоматически. В результате анализа зависимостей по данным между операторами программы на языке Норма+, в случае разрешимости этих зависимостей, представляется возможным построить так называемую «параллельную ярусную схему» выполнения программы. На каждом ярусе данной ярусной схемы располагаются операторы программы, которые не имеют зависимостей друг от друга и могут выполняться независимо и, соответственно, параллельно. В то же время каждый из этих операторов имеет зависимость от одного или более операторов, располагающихся на предыдущем ярусе ярусной схемы. Таким образом, группу операторов, располагающихся на одном уровне ярусной схемы, можно выполнять параллельно в результирующей программе, но только после того, как будут полностью выполнены все операторы предыдущего уровня ярусной схемы. Параллельная ярусная схема программы является, фактически, представлением идеального (естественного) параллелизма, определяемого соотношениями и зависимостями между расчетными переменными программы.
В результате ряда исследований по трансформации описанной параллельной ярусной схемы программы на языке Норма+ в программу с использованием технологии М"УГО1А CUDA для графических процессоров авторами предлагается использовать сле-
дующий подход для автоматического построения исполняемой программы для графических процессоров с использованием технологии CUDA.
Исполняемая программа стартует и завершается на центральном процессоре, на нем же выполняется ввод-вывод данных и итерационные циклы. Очевидно также, что вся логика по управлению вычислениями, выделением памяти на графическом процессоре, организация обменов данными между памятью графического и центрального процессора, тоже создается в программном коде, выполняющемся на центральном процессоре. Сами вычислительные операторы выполняются на графическом процессоре. Для этого они группируются в определенные наборы, каждый из которых может выполняться в пределах одного ядра программы с использованием технологии М"УГО1А CUDA. Программа, выполняющаяся на центральном процессоре, осуществляет контроль за общим выполнением программы, запускает полученные ядра в определенном порядке, ведет итерационные циклы и проверяет условие выхода из итерации. На рис.1 приведен пример общей схемы выполнения исполняемой программы.
Рис. 1. Пример общей схемы выполнения исполняемой программы
Один из ключевых моментов в этой схеме организации исполняемой программы, и он же алгоритмически самый тяжелый - это задача группировки вычислительных операторов в ядра, которые будут выполняться на графическом процессоре. С одной стороны, чем больше операторов будет содержать такое ядро, т.е. чем крупнее будут исполняемые ядра в программе, тем эффективнее она будет работать, т.к. будет меньше запусков ядер, меньше передачи параметров, в том числе и через память центрального процессора и др. Но, с другой стороны, все операторы, сгруппированные в одно ядро, должны работать с одним и тем же распределением рабочих областей на конфигурацию блоков и нитей, определяемую при запуске ядра. И может оказаться, что поиск такого распределения для большой группы операторов, объединенных в одно ядро, может дать худший результат, чем разбиение этой группы операторов на более мелкие отдельные ядра и поиск распределения для каждой группы в отдельности. Кроме того, при поиске
таких группировок необходимо соблюдать ряд обязательных условий, накладываемых ярусной схемой и сущностью группируемых операторов. Например, мы можем менять порядок выполнения операторов, но только в пределах одного яруса. А операторы ввода-вывода должны выполняться на центральном процессоре.
В итоге, после ряда исследований характеристик получающихся программ, предлагается использовать следующий набор эвристик и обязательных условий для решения задачи группировки операторов программы на языке Норма+ в один набор, который может выполняться в пределах одного ядра программы с использованием технологии ]^УГО1А CUDA:
1. Каждая нить графического процессора осуществляет вычисления в одной точке расчетной области Норма-программы (другими словами, одной точке индексного пространства вычисления).
2. Поскольку операторы, находящиеся на одном ярусе параллельной ярусной схемы, можно выполнять параллельно, они могут выполняться одним ядром CUDA программы, причем в любом порядке. Поэтому в очередном ярусе выбираются операторы, которые выполняются на одной и той же одномерной, двумерной или трехмерной подобласти. Вообще говоря, оператор может выполняться и на области большей размерности, но выбранная для распараллеливания подобласть должна присутствовать в области выполнения оператора. Эти выбранные операторы могут выполняться одним ядром, при условии соблюдения описанных ниже условий.
3. Выбранная для распараллеливания подобласть разбивается на блоки и нити для ядра графического процессора следующим образом:
• если подобласть одномерная, то за число нитей берется предопределенное число, кратное степени двойки (1024 по умолчанию) и равное максимальному количеству нитей в блоке на целевой архитектуре. А для блоков выбирается одномерный массив с учетом ограничений целевой архитектуры на размер массива блоков по одному направлению или двумерный массив таким образом, чтобы произведение количества блоков на количество нитей полностью покрывало выбранную подобласть;
Пусть D(N) - выбранная для распараллеливания подобласть из N точек. Тогда распределение выглядит так:
D(N), ((^ + 1023)/1024) <= 65535) => кете1<<< (^ + 1023) /1024), 1024 >>>()
D(N), ((^ + 1023)/1024) > 65535) =>
кете1<<< dim3((((N + 1023)/1024 + 65534)/65535), 65535), 1024 >>>()
• если подобласть двумерная, и одно из направлений меньше или равно максимальному количеству нитей в блоке на целевой архитектуре, то этому направлению будут соответствовать нити, количество которых определяется как ближайшая сверху степень двойки. А для второго направления подобласти выбирается одномерный массив блоков, равный размеру этого направления, с учетом ограничений целевой архитектуры на размер массивов блоков по одному направлению, или подбирается двумерный массив блоков таким образом, чтобы общее количество блоков полностью покрывало второе направления подобласти.
D(N1, N2), (N2 <= 1024), (N1 <= 65535) => кете1<<< N1, N2 >>>()
D(N1, N2), (N2 <= 1024), (N1 > 65535) => кете1<<< Шт3((Ш + 65534)/65535, 65535), N2 >>>()
• если подобласть двумерная, и ни одно из направлений не может быть представлено нитями в блоке на целевой архитектуре, то одному из направлений назначается двумерный массив, одно измерение которого - нити, количество которых определяется как предопределенное число, кратное степени двойки (1024 по умолчанию) и равное максимальному количеству нитей в блоке на целевой архитектуре. А второе измерение этого двумерного массива выбирается из одного из направлений в двумерном массиве блоков, таким образом, чтобы произведение количества блоков по данному направлению на количество нитей полностью покрывало рассматриваемое направление подобласти. Второму направлению подобласти ставится в соответствие второе направление в двумерном массиве блоков.
D(N1, N2), (N2 > 1024), (N1 <= 65535) => кете1<<< dim3(N1, (N1 + 1023)/1024), 1024 >>>()
• если подобласть трехмерная, и одно из направлений меньше или равно максимальному количеству нитей в блоке на целевой архитектуре (а остальные направления меньше или равны размеру блока), то этому направлению будут соответствовать нити количеством ближайшая сверху степень двойки. А для остальных направлений подобласти ставится в соответствие двумерный массив блоков соответствующей размерности.
D(N1, N2, N3), (N3 <= 1024), (N1, N2 <= 65535) => кете1<<< Шт3(Ш, N2), N3 >>>()
4. Если оператор является оператором ввода-вывода данных, то он должен выполняться на центральном процессоре, и он может выполняться параллельно с запуском ядра. При этом дальнейшее добавление в формируемое ядро операторов со следующего уровня (п.6) становится невозможным.
5. Если в операторе присутствует функция редукции, и область применения данной функции редукции полностью представлена нитями по какому-то индексу, а все остальные индексы из области применения функции редукции (если они есть) не входят в область, распараллеленную в текущем ядре, то этот оператор может быть выполнен ядром без исключений. В противном же случае, если область применения функции редукции имеет индексы, которым соответствует какое-либо направление массива блоков, данное ядро может выполнить функцию редукции только частично, и затем выполнение ядра должно быть прекращено и запущено следующее ядро (или даже последовательно несколько ядер) с другим распределением - так, чтобы в итоге все частичные результаты выполнения функции редукции были представлены исключительно нитями и было возможно получить окончательный результат. Такой оператор должен быть поставлен последним в ядре, и дальнейшее добавление в формируемое ядро операторов со следующего уровня (п.6) становится невозможным.
6. После завершения формирования последнего вычислительного ядра с операторами данного слоя ярусной схемы можно попробовать добавить в это ядро операторы со следующего слоя ярусной схемы при соблюдении двух описанных ниже условий. Это возможно, т.к. ядро выполняет операторы, содержащиеся в нем, последовательно в отдельно взятой точке распараллеленной области. Если зависимости между операторами разных уровней параллельной ярусной схемы простые, без смещения по какому-
либо индексу из области распараллеливания (например X = R = func(X)), то последовательно выполненные действия в каждой точке области дадут правильный результат. И не важно, скажем, что при вычислении R в точке k значение X в точке k+1 еще не было вычислено другой нитью - важно только, что значение X в точке k уже вычислено. Затем, после завершения очередного уровня параллельной ярусной схемы, можно начинать включать в текущее ядро операторы со следующего уровня и т.п., пока процесс добавления не закончен по одной из причин, описанной в п. 4-6. Добавление таких операторов возможно при соблюдении следующих условий:
• если на добавляемом слое есть операторы ввода-вывода, они должны быть выполнены на центральном процессоре после завершения выполнения создаваемого ядра, формирование которого должно быть закончено текущим слоем;
• в добавляемых операторах не должно быть использования переменных с индексами из подобласти распределения со смещением (например R = func(X[i+1])). Если такие операторы есть, то они не должны быть добавлены в создаваемое ядро, формирование которого должно быть закончено текущим слоем. С такими оставшимися операторами можно начать формировать следующее ядро.
3. Пример построения программы с использованием CUDA
В качестве примера рассмотрим применение описанной выше схемы для простого фрагмента программы на языке Норма+, см. рис. 2.
Oij:(Oj:(j=1..w);Oi:(i=1..v)).
VARIABLE Vij DEFINED ON Oij DOUBLE.
VARIABLE Vsum DEFINED ON Oi DOUBLE.
DOMAIN PARAMETERS v=2000,w=3000.
FOR Oij ASSUME Vij = j+(i-1)*w.
FOR Oi ASSUME Vsum = SUM((Oj)Vij).____________________________
Рис. 2. Фрагмент программы на языке Норма+
В приведенном фрагменте величине Vij, определенной на двумерной области Oij с индексами i и j, присваиваются начальные значения j+(i-1)*w, а затем производится суммирование по индексу j (по области Oj). В данном случае распараллеливается вся область Oij. По индексу j область распределяется на нити и на одну из размерностей блоков, а по индексу i - на вторую размерность блоков (З-й подпункт пункта З приведенной выше схемы). Начало оператора суммирования выполняется в первом же ядре, но, т.к. все направление индекса j не покрывается полностью нитями, необходимо создание второго ядра, завершающего суммирование.
Схема распределения области Oij на блоки и нити приведена на рис. 3.
1 -V 2 ■■■ V
jl j2 ■ j 1024 jl j2 ■ j 1024
j 1025 j 1026 ■ j2048 j 1025 j 1026 ■ j2048
j2049 j2050 j2049 j 2050 ■
С локи 2000 Ч, 3
Рис. 3. Отображение области Oij на блоки и нити
Жирными линиями выделены блоки, их конфигурация 2000 (размер по индексу i) на 3 (дополнительное направление по индексу j, которое в объединении с 1024 нитями в каждом блоке обеспечивает представление 3000 точек по индексу j). Нити в каждом блоке показаны пунктиром.
#define v 2000
#define w 3000
#define wthreadKl 1024 // Ближайшая сверху степень 2 от w,
// с учетом ограничения на кол-во нитей в блоке #define wBlockYKl (w+wthreadK1-1)/wthreadK1
#define wthreadKlRed 4 // Ближайшая сверху степень 2 от wBlockYKl
___global___ void SumKl() {
__shared___ double shared[wthreadK1];
// Получаем индекс j
int j = threadldx.x + blockIdx.y*wthreadK1;
if (j < w) { // Проверяем, что не вышли за границу области
int i = blockIdx.x; int gidx = i*w + j;
// Присваиваем начальные значения Vij dev[gidx] = j + 1 + i * w;
// Начинаем суммирование shared[threadIdx.x] = Vij dev[gidx]; for (int d = wthreadKl/2; d > 0; d /= 2) {
int from = threadldx.x + d;
___syncthreads();
if ((threadldx.x < d) &&
(from + blockldx.y*wthreadK1 < w)) shared[threadIdx.x] += shared[from];
}
// Сохраняем частичный результат if (threadldx.x == 0)
Vsum block[blockIdx.x*wBlockYK1+blockIdx.y] = shared[0];
}
J_____________________________________________________________________________
Рис. 4. Программа для первого ядра
На рис. 4 приведена получившаяся программа для первого ядра вместе с общими определениями констант. В ней производится инициализация переменной Vij (в ядре представлена переменной Vij_dev) и начинается ее суммирование по области Oj. Для этого каждый блок вычисляет сумму для всех входящих в него нитей. Но, т.к. точки области Oj распределены на 3 блока, то итоговый результат, который должен быть вычислен как сумма частичных результатов, полученных в каждом из этих 3-х блоков, в рамках данного ядра вычислить невозможно.
Поэтому для завершения суммирования организуется второе ядро. Программа для него приведена на рис. 5.
__global___ void SumK1Red() {
// Осуществляет суммирование Vsum block
__shared___ double shared[wBlockYK1];
if (threadIdx.x < wBlockYKl) { int i = blockIdx.x;
int gidx = i*wBlockYK1 + threadIdx.x;
// Копируем насчитанные предыдущим ядром частичные данные shared[threadIdx.x] = Vsum block[gidx];
// Завершаем суммирование
for (int d = wthreadK1Red/2; d > 0; d /= 2) {
int from = threadIdx.x + d;
__syncthreads();
if ((threadIdx.x < d) && (from < wBlockYKl)) shared[threadIdx.x] += shared[from];
}
// Сохраняем итоговый результат if (threadIdx.x == 0)
Vsum_dev[i] = shared[0];
}
J____________________________________________________________________________
Рис. 5. Программа для второго ядра
Фрагмент программы, выполняющейся на центральном процессоре, приведен на рис. 6. Он, как можно заметить, заключается в последовательном вызове ядер на нужной конфигурации блоков и нитей.
SumK1<<< dim3(v, wBlockYKl), wthreadKl >>>();
SumK1Red<<< v, wthreadKlRed >>>();
Рис. 6. Фрагмент программы для центрального процессора
4. Предварительные результаты
Приведенная выше схеме построения CUDA программы по декларативным описаниям была применена к программе на языке Норма+ из области газодинамики. Была взята программа на языке Си, полученная компиляцией указанной программы на языке Норма+ в последовательную программу, и ручным способом трансформирована в программу с использованием CUDA так, как это мог бы сделать компилятор по описанной выше схеме. Полученная параллельная программа была запущена на вычислительном кластере K-100 [9]. Производительность параллельной программы сравнивалась с последовательной программой и OpenMP версией той же самой программы. Времена выполнения версий программы приведены в табл.
Таблица
Время выполнения различных версий программы
Последовательная программа, 1 ядро Intel Xeon X5670 OpenMP программа, 11 ядер Intel Xeon X5670 CUDA программа, 1 графический процессор nVidia Fermi C2050
44,5 с 4,62 с 1,35 с
Полученный результат - ускорение в 33 раза по сравнению с последовательной программой - позволяет надеяться на успешное применение схемы для автоматического
построения эффективных программ для графических процессоров по декларативным описанием.
Заключение
В рамках данной работы были рассмотрены методы и средства автоматического построения параллельных программ для графических процессоров по непроцедурным спецификациям. Приведены алгоритмы автоматического построения программ для графических процессоров, основанные на анализе зависимостей и выявлении внутреннего параллелизма графа потока управления, полученного по непроцедурным спецификациям, заданным на языке Норма. Приведены результаты применения описанных алгоритмов для создания расчетной программы с использованием технологии CUDA при решении задачи из области газодинамики и оценена производительность получающихся параллельных программ
В настоящее время ведется разработка версии компилятора с языка Норма+ для графических процессоров с использованием технологии CUDA с применением описанной в данной работе схемы.
Работа выполнена при финансовой поддержке гранта РФФИ № 12-01-00527-а. Литература
1. Кривов, М.А. Опыт портирования среды для HDR-обработки изображений на GPU и АРи / М.А. Кривов, М.Н. Притула, С.Г. Елизаров. URL: http://pavt.susu.ru/2012/short/175.pdf (дата обращения: 16.08.2013).
2. ОрепАСС. URL: http://openacc.org (дата обращения: 16.08.2013).
3. Бахтин, В.А. Распараллеливание с помощью DVM-системы некоторых приложений гидродинамики для кластеров с графическими процессорами / В.А. Бахтин, И.Г. Бородич, Н.А. Катаев, М.С. Клинов, В.А. Крюков, Н.В. Поддерюгина, М.Н. Притула, Ю.Л. Сазанов // Научный сервис в сети Интернет: поиск новых решений: Труды Международной суперкомпьютерной конференции (17-22 сентября 2012 г., г. Новороссийск). — М.: Изд-во МГУ, 2012. — С. 444-450.
4. Описание языка программирования COLAMO. URL: http://colamo.parallel.ru (дата обращения: 16.08.2013).
5. Задыхайло, И.Б. Организация циклического процесса счета по параметрической записи специального вида / И.Б. Задыхайло // Журн. выч. мат. и мат. физ. — 1963. — Т. 3, № 2, — С. 337-357.
6. Андрианов, А.Н. Норма. Описание языка. Рабочий стандарт / А.Н. Андрианов,
А.Б. Бугеря, К.Н. Ефимкин, И.Б. Задыхайло — М.: Препринт ИПМ им.
М.В. Келдыша РАН, 1995. — № 120. — 52 с.
7. Андрианов, А.Н. Декларативный язык Норма и программирование для новых архитектур: многоядерные системы / А.Н. Андрианов, А.Б. Бугеря, К.Н. Ефимкин, П.И. Колударов // Научный сервис в сети Интернет: Суперкомпьютерные центры и задачи: Труды Международной суперкомпьютерной конференции (20-25 сентября 2010 г., г. Новороссийск). — М.: Изд-во МГУ, 2010. — С. 68-70.
8. Система Норма. URL: http://www.keldysh.ru/pages/norma (дата обращения:
16.08.2013).
9. Гибридный вычислительный кластер К-100. URL:
http://www.kiam.ru/MVS/resourses/k100.html (дата обращения: 16.08.2013).
Андрианов Александр Николаевич, д.ф.-м.н., в.н.с., Институт прикладной математики им. М.В. Келдыша РАН (Москва, Российская Федерация), and@a5.kiam.ru.
Бугеря Александр Борисович, к.ф.-м.н., с.н.с., Институт прикладной математики им. М.В. Келдыша РАН (Москва, Российская Федерация), bug@a5.kiam.ru.
Гладкова Екатерина Николаевна, электроник, Институт прикладной математики им. М.В. Келдыша РАН (Москва, Российская Федерация), gladkova@keldysh.ru.
Ефимкин Кирилл Николаевич, к.ф.-м.н., зав. отделом, Институт прикладной математики им. М.В. Келдыша РАН (Москва, Российская Федерация), bigcrocodile@yandex.ru.
Колударов Павел Иванович, стажер-исследователь, Институт прикладной математики им. М.В. Келдыша РАН (Москва, Российская Федерация), koludarov@yandex.ru.
Поступила в редакцию 18 сентября 2013 г.
Bulletin of the South Ural State University Series “Computational Mathematics and Software Engineering”
2014, vol. 3, no. 1, pp. 5-16
AUTOMATIC PROGRAMS GENERATING FROM NON-PROCEDURAL SPECIFICATIONS FOR GPUS
A.N. Andrianov, Keldysh Institute of Applied Mathematics RAS (Moscow, Russian Federation),
A.B. Bugerya, Keldysh Institute of Applied Mathematics RAS (Moscow, Russian Federation) ,
E.N. Gladkova, Keldysh Institute of Applied Mathematics RAS (Moscow, Russian Federation) ,
K.N. Efimkin, Keldysh Institute of Applied Mathematics RAS (Moscow, Russian Federation) ,
P.I. Koludarov, Keldysh Institute of Applied Mathematics RAS (Moscow, Russian Federation)
A new approach to automatic programs generating for GPUs from non-procedural task’s specification in Norma language is considered. Algorithms of automatic programs generating for GPUs based on dependencies analysis and discovering control flow internal parallelism obtained from Norma program are given. Described algorithms applied to gas-dynamics task solving are carried out in calculating CUDA program. The performance of generated CUDA program is estimated.
Keywords: HPC, parallel programming, GPU, CUDA, Norma language, non-procedural specifications.
References
1. Krivov M.A., Pritula M.N., Elizarov S.G. Opyt portirovaniya sredy dlya HDR- obrabot-ki izobrazheniy na GPU i APU [Porting HDR Image Processing System to GPU and APU Experience] URL: http://pavt.susu.ru/2012/short/175.pdf (accessed: 16.08.2013).
2. OpenACC. URL: http://openacc.org (accessed: 16.08.2013).
3. Bakhtin V.A., Borodich I.G., Kataev N.A., Klinov M.S., Krukov V.A., Podderugi-na N.V., Pritula M.N., Sazanov U.L. Rasparallelivanie s pomoschyu DVM-systemy nekotoryh prilozheniy gidrodinamiki dlya klasterov s graficheskimi processorami [Paral-lelization Some Hydrodynamics Applications Using DVM-System for Clusters with GPUs]. Nauchniy servis v seti Internet: poisk novyh resheniy: Trudy mezhdunarodnoi superkomputernoi konferentsii (17-22 sentyabrya 2012, Novorossiysk) [Scientific Service in Internet: Searching the New Solutions: Proceedings of the International Supercomputing Conference (Novorossiysk, Russia, September, 17-22, 2012)] — Moscow, Publishing of the Moscow State University, 2012. — P. 444-450.
4. Opisanie yazyka programmirovaniya COLAMO [COLAMO Programming Language Description]. URL: http://colamo.parallel.ru (accessed: 16.08.2013).
5. Zadykhailo I.B. Organizatsiya tsiklicheskogo protsessa scheta po parametricheskoy zapisi spetsialnogo vida [Organizing of Cycle Process Calculating by Special Parametrized Record]. Vychislitelnaya matematika i matematicheskaya fisika [Computing Mathematics and Mathematical Physics]. — 1963. — Vol. 3, No 2. — P. 337-357.
6. Andrianov A.N., Bugerya A.B., Efimkin K.N., Zadykhailo I.B. Norma. Opisanie yazyka. Rabochiy standart [Norma Programming Language. Draft Standard]. — Moscow, Preprint of the Keldysh Institute of Applied Mathematics RAS. — 1995. — No 120. — 52 p.
7. Andrianov A.N., Bugerya A.B., Efimkin K.N., Koludarov P.I. Deklarativniy yazyk Norma i programmirovanie dlya novyh arkhitektur: mnogoyadernie systemy [Declarative Norma Language and Programming for New Architectures: Multi-Core Systems]. Nauch-niy servis v seti Internet: Supercomputernie tsentry i zadachi: Trudy mezhdunarodnoi superkomputernoi konferentsii (20-25 sentyabrya 2010, Novorossiysk) [Scientific Service in Internet: Supercomputing Centers and Tasks: Proceedings of the International Supercomputing Conference (Novorossiysk, Russia, September, 20-25, 2010)] — Moscow, Publishing of the Moscow State University, 2010. — P. 68-70.
8. Norma. URL: http://www.keldysh.ru/pages/norma (accessed: 16.08.2013).
9. Hybrid Supercomputing K-100. URL: http://www.kiam.ru/MVS/resourses/k100.html
(accessed: 16.08.2013).
Received 18 September 2013
