CC BY
27
ИНФОРМАТИКА, ВЫЧИСЛИТЕЛЬНАЯ ТЕХНИКА И УПРАВЛЕНИЕ INFORMATICS, COMPUTER ENGINEERING AND CONTROL
УДК 004.4'242 DOI: 10.17213/0321-2653-2015-3-3-9
АВТОМАТИЧЕСКАЯ ГЕНЕРАЦИЯ GPU-КОДА В ДИАЛОГОВОМ ВЫСОКОУРОВНЕВОМ ОПТИМИЗИРУЮЩЕМ РАСПАРАЛЛЕЛИВАТЕЛЕ
AUTOMATIC GENERATION OF GPU CODE IN DVOR
© 2015 г. А.Н. Аллазов, С.А. Гуда, Р.И. Морылев
Аллазов Афар Назим оглы - студент, Южный Федеральный университет, г. Ростов-на-Дону, Россия. E-mail: afarallazov@ gmail.com
Гуда Сергей Александрович - канд. физ.-мат. наук, доцент, Южный Федеральный университет, г. Ростов-на-Дону, Россия. E-mail: gudasergey@gmail.com
Морылев Роман Игоревич - научный сотрудник, Южный Федеральный университет, г. Ростов-на-Дону, Россия. E-mail: frg10@yandex.ru
Allazov Afar Nazim - student, Southern Federal University, Rostov-on-Don, Russia. E-mail: afarallazov@gmail.com
Guda Sergey Alexandrovich - Candidate of Physical and Mathematical Sciences, assistant professor, Southern Federal University, Rostov-on-Don, Russia. E-mail: gudasergey@gmail.com
Morilev Roman Igorevich - researcher, Southern Federal University, Rostov-on-Don, Russia. E-mail: frg10@yandex.ru
Описываемый диалоговый высокоуровневый оптимизирующий распараллеливатель (ДВОР) позволяет автоматизировать этапы разработки CUDA-программ. ДВОР обладает следующими возможностями: автоматическое определение распараллеливаемых циклов, визуализация информационных зависимостей, автоматическая трансформация программ (расщепление тела циклов, слияние, гнездование, раскрутка, разбиение, оптимизация рекуррентности и т.д.), генерация CUDA-кода и автоматическое определение оптимальных параметров запуска задачи на видеокарте.
Ключевые слова: генерация CUDA-программ; оптимизация программ; ДВОР.
The proposed Dialogoviy Visokourovneviy Optimiziruyusshiy Rasparallelivatel (DVOR) allows of automating the steps of CUDA program development. DVOR has the following features: automatic detection ofparal-lelizable loops, data dependencies visualization, a variety of automatic program transformations (loop distribution, fusion, nesting, unrolling, strip mining, recurrence elimination, etc), CUDA code generation and automatic optimal parameters determination for running loop nest on GPU.
Keywords: CUDA code generation; program optimizations; DVOR.
Введение
В последнее время большую популярность приобрели ускорители параллельных вычислений. На рынке активно конкурируют устройства от NVIDIA, AMD и Intel. Разработано множество средств, облегчающих работу программиста: технологии программирования (CUDA, OpenCL), библиотеки программ, прагмы (OpenACC, OpenMP), расширения языков и т.д. В данной статье описано одно из таких средств, основанное на диалоговом подходе к оптимизации кода. Визуальный интерфейс пользователя позволяет удобно выделять фрагменты кода, применять преобразования, делать проверки, строить различные графы программы.
Состоящий из более чем 180 000 строк кода на C++, кроссплатформенный диалоговый высокоуровневый оптимизирующий распараллеливатель программ (ДВОР) [1] обладает большой базой анализаторов и преобразователей кода и многофункциональным GUI с возможностью выделения участков кода. По своим параметрам он похож на автораспараллели-вающие системы: Rose Compiller [2], SUIF [3], Cetus [4], PPCG [5], Par4All [6], Parascope Editor [7]. Авторы статьи реализовали в данной распараллеливающей системе графический интерфейс пользователя и автоматический генератор CUDA-кода. ДВОР использует анализ информационных зависимостей в циклах, основанный на полиэдральной модели, с собственной
реализацией параметрического метода Гомори [8]. Реализована возможность создания проектов с несколькими файлами, многочисленные преобразования циклов, проверки корректности применяемых преобразований. Так же, как и менеджер памяти времени выполнения [9], генератор CUDA-кода на этапе компиляции находит оптимальное расположение операций копирования данных на ускоритель и обратно, предотвращая копирование неизмененных данных, частые операции перемещения данных внутри циклов и лишние синхронизации данных между вызовами функций на видеокарте. Модуль статической профилировки GPU-кода позволяет находить оптимальные размеры блока и оптимальное соответствие циклов гнезда размерностям пространства параллельных потоков видеокарты. В отличие от PPCG [5] ДВОР полагается на аппаратный тайлинг циклов и на автоматическое кэширование, что позволяет снизить накладные расходы и добиться лучшей производительности и компактности кода kernel-функций. Как показывают численные эксперименты (параграф 5), ДВОР позволяет получать преобразованную программу в компактной форме, производительность которой сравнима с результатами компилятора PPCG [5], распараллеливающего без участия пользователя.
Статья организована следующим образом. В параграфе 2 описаны особенности GPU и диалога с пользователем в ДВОР. Параграф 3 посвящен структуре генератора GPU-кода. В параграфе 4 рассказано об анализе общих для хоста и GPU-ячеек памяти, выборе оптимального расположения операций копирования данных. В параграфе 5 представлены результаты экспериментальных исследований.
Диалоговая оптимизация программ
Перенос существующих программ на новые архитектуры, такие как GPU, требует более функциональные инструменты, чем традиционные компиляторы. Не ограниченное интерфейсом командной строки тесное взаимодействие с пользователем позволяет применять оптимизирующие преобразования для отдельных фрагментов программы. На мехмате Южного федерального университета разрабатывается ДВОР - диалоговый высокоуровневый оптимизирующий распараллеливатель - система для интерактивной трансформации, распараллеливания и переноса существующих программ на различные архитектуры, в том числе GPU (рис. 1).
Он включает в себя разнообразные инструменты, которые можно разделить на три класса:
- автоматические преобразования программ с проверкой;
- анализ распараллеливаемых циклов;
- построение графовых моделей программы (графа вызовов процедур, графа потока управления, графа информационных зависимостей).
Программа пользователя подается на вход парсе-ра (рис. 2). В ДВОР на данный момент работают два собственных парсера для языков Си и Фортран и пар-сер Clang из проекта LLVM. Возвращаемое объектное представление программы, специфическое для ее языка, транслируется специальным модулем в языко-независимое высокоуровневое внутреннее представление, с которым затем работает пользователь в графическом интерфейсе.
Flte Edit Profect Joots window Help t>- t-S Ш sen erste cods
- Analyse
i? G
Reprise Explorer
■n.c ©\
* correlation ▼ correlations
t _builtln_va. f init_array f print_array f kernel_corn f main * potybench.c
t ^builtin^va^ v polybenchj
V polybenchj
V polybenchj
V polybenchj
V polybenchj
V polybenchj f rtclock
f polybenchj f polybenchj f polybenchj f polybenchj f polybenchj f xmalloc f posixjriemî f polybench_
kernel j:orrelation
-a
meantj] = (meantjl / f\oat_n)
>
for j = 0; j < m; ] = (j + {
stddevCj] = 0, for i = o; i < n; i = (i f
stddevtj] = (stddevtj]
У
stddevtj] = (stddevtj] / \ stddevtj] = sqrttstddevlj~
if stddevtj] <= eps {
_uni49 = 1.
>
else {
_un149 = stddevtj]
}
stddevtj] = _uni49
>
for
{
i « o; i < n; i = (i
Category
Desmptton
Loops Loop iterations can be
• Loops Loop iterations can be
• Loops Loop is sequental
• Loops Loop is sequental
• Loops Loop iterations can be executed independently j is sequental_
Open Selector XML Dump Build DepGraph Suild callGraph
a[i, j] ■ meant]])))
Arithmetic Operator Expansion
Exception Generator
Full inline substitution
Subroutine splitting
Corrupt or
Loop Distribution
Loop Fragmentation
Loop Full Unrolling
Loop Header Removal
Loop Nesting
Loop Unrolling
Recurrency Elimination
Strip Mining
Execute on GPU
Parallel execute on GPU
Linear view Ready
GPGPU
► • Loops
* • Loops
► • Loops
* ♦ Loops Parser Output
The loop can be executed in parallel, Do you want to parallelize it using GPGPU computations? Loop is sequental
Loop iterations can be executed independently Loop iterations can be executed independently Loop iterations can be executed Independently Issues GPGPU
Рис. 1. Интерфейс пользователя ДВОР
Анализ граф потока управления анализ псевдонимов граф зависимостей решетчатый граф
граф вызовов граф вычислений статическое профилирование ^
Рис. 2. Схема работы пользователя с ДВОР
Процесс генерации CUDA-кода из внутреннего представления является итеративным. Каждая итерация состоит из следующих этапов:
1. Выбор фрагментов программы, которые должны выполняться на GPU. Для этого можно использовать автоматический анализатор, обнаруживающий распараллеливаемые без трансформаций циклы, или выбирать эти фрагменты вручную.
2. Генерация kernel-функций для выбранных фрагментов программы на CUDA.
3. Компиляция и профилирование сгенерированной программы. Если пользователь удовлетворен полученными результатами, то работа на этом заканчивается.
4. В противном случае необходимо сделать дополнительные трансформации, оптимизирующие программу под архитектуру GPU, такие как разрезание, перестановка, гнездование циклов и т.д. Затем процесс повторяется.
Тем не менее, пользователь не ограничен этой схемой и может использовать любой из инструментов, доступных в ДВОР, в любом порядке, комбинируя преобразования для достижения наилучшей производительности.
Структура генератора
На вход генератору GPU-кода подаются фрагменты двух типов:
- тесное гнездо циклов, размерности от 1 до 3, итерации которого должны быть запущены в параллельных потоках на ускорителе;
- последовательность операторов программы, принадлежащих одному блоку, которые должны быть запущены в одном потоке на ускорителе.
Фрагменты второго типа приходится запускать на ускорителе, чтобы избежать излишних пересылок данных между оперативной памятью и памятью ускорителя. Генератор GPU-кода работает в 4 этапа:
1. Инициализация параметров участков кода, предварительные проверки и преобразования.
2. Анализ участков кода и определение параметров отображения кода на ускоритель.
3. Преобразование внутреннего представления в форму, соответствующую технологии CUDA.
4. Получение текстового представления кода.
На первом этапе выполняются предварительные преобразования: канонизация циклов, объединение фрагментов в блоки. Происходит проверка сохранения корректности программы, при запуске выделенного пользователем участка кода на ускорителе: распараллеливаемы ли циклы, отсутствуют ли псевдонимы. Также происходит проверка ограничений, накладываемых на входные данные текущей версией генера-
тора кода. На данный момент запрещены структуры, фрагменты из разных функций, массивы указателей.
На втором этапе генератор анализирует вхождения переменных в предназначенных для запуска на графическом ускорителе фрагментах программы. Для общих с ускорителем переменных определяются типы и размер создаваемых GPU-переменных, способ передачи в каждую kernel-функцию, оптимальное расположение операций выделения/освобождения памяти на ускорителе и функций копирования данных на и с ускорителя. Для многомерных гнезд циклов генератор находит оптимальное отображение пространства итераций на пространство параллельных потоков и оптимальный размер блока потоков. Когда границы внутреннего цикла зависят от счетчика внешнего цикла, вычисляются их максимальные значения. В текущей версии генератора поддерживается линейно зависящие от счётчиков выражения с числовыми коэффициентами. На втором этапе также определяются функции, которые необходимо пометить ключевым словом _device_, так как они вызываются из предназначенных для запуска на видеокарте фрагментов.
На третьем этапе работы генератора происходят преобразования внутреннего представления программы: вставляются операции выбора ускорителя и инициализации контекста на видеокарте, создаются GPU-переменные, вставляются в найденные на предыдущем этапе анализа позиции операций выделения/освобождения памяти на устройстве, обмена данными с ускорителем. Из выделенных пользователем фрагментов конструируются kernel-функции, а фрагменты заменяются на их вызовы. Так, например, трехмерное гнездо циклов for (int i0 = 0; i0 < n0; i0++) for (int i1 = 0; i1 < n1(i0); i1++) for (int i2 = 0; i2 < n2(i0,i1); i1++) LoopBodyBlock
заменится на вызов kernel-функции:
if (n0 > 0 && n_max1 > 0 && n_max2 > 0) {
dim3 blockDim = dim3(b0, b1, b2); dim3 gridDim = dim3(g0, g1, g2); kernel1<<<gridDim,blockDim>>>(...);
gpuErrchk(cudaPeekAtLastError()); }
которая определена следующим образом:
_global_ void kernel1(...)
{
10 = blockIdx.x*blockDim.x+threadIdx.x
11 = blockIdx.y*blockDim.y+threadIdx.y
12 = blockIdx.z*blockDim.z+threadIdx.z if (i0<n0 && i1<n1(i0) && i2<n2(i0,i1))
LoopBodyBlock }
где b0, b1, b2 - оптимальные размеры блока; gi = = (n_maxi + bi - 1) / bi, i = 0, 1, 2; n_max1, n_max2 -максимальные значения границ вложенных циклов.
На четвёртом этапе работы генератора преобразованный код из внутреннего представления переводится в текстовое. Система старается оставить код по возможности похожим на его изначальную форму. Используемый парсер Clang добавляет в программу объявления из всех подключаемых библиотек. На
данном этапе происходит анализ используемых объявлений. Неиспользуемые объявления удаляются, а взятые из заголовочных файлов заменяются на соответствующие директивы include.
Анализ ячеек памяти
Для выполнения выделенных пользователем фрагментов кода на ускорителе необходимо определить переменные, которые будут использованы на двух устройствах: CPU и GPU (выделить память на видеокарте и обеспечить синхронизацию содержимого этих переменных). Для корректной работы анализатора необходимо, чтобы в коде отсутствовали псевдонимы. Для каждого фрагмента кода общие для CPU и GPU ячейки памяти мы делим на три непересекающихся типа:
1) скалярные переменные, значения которых не меняются в GPU-фрагментах программы;
2) изменяемые в GPU-фрагментах скалярные переменные;
3) массивы.
Операторы выделения памяти на видеокарте cudaMalloc необходимы только для ячеек второго и третьего типа. Ячейки памяти первого типа передаются в kernel-функции по значению. Локальные переменные GPU-фрагментов кода и счётчики распараллеливаемых циклов анализатор относит к локальным GPU-переменным и не считает общими с CPU.
Для каждой общей с GPU ячейки памяти необходимо создать переменную, описывающую копию данной ячейки на GPU, определить расположение операций выделения, освобождения и копирования данных. Тип создаваемой переменной равен для скаляра указателю на тип скаляра, для массивов - указателю на тип ячейки массива и для указателей - совпадает с их типом.
Операторы cudaMalloc и cudaFree размещаются анализатором в тех же местах, где происходит выделение и освобождение памяти на хосте. Для переменных на стеке оператор cudaMalloc вставляется в начало блока определения, cudaFree - в места выхода из блока определения переменной. Для динамически выделенных ячеек памяти операторы cudaMalloc и cudaFree размещаются непосредственно после операторов выделения/освобождения памяти на хосте, если это происходит в текущей функции; и в начале и в местах выхода из функции, в противном случае.
Правильная расстановка операций копирования cudaMemcpy для синхронизации экземпляров одного массива в памяти разных устройств имеет зачастую ключевое значение. Лишние синхронизации неизмяе-мых переменных или неиспользуемых на хосте между вызовами kernel-функций, особенно внутри циклов, могут существенно увеличить время работы программы. Предлагаемый подход основан на анализе программы во время компиляции.
Для программы с предназначенными для выполнения на GPU-фрагментами можно выделить несколько наивных схем синхронизации общих с GPU-переменных. Первая - ориентироваться на GPU-фраг-
менты кода, вставляя перед каждым фрагментом операции копирования данных с хоста на GPU и после -копирования в обратном направлении для изменившихся ячеек памяти. Вторая, симметричная первой, -ориентироваться на CPU-фрагменты кода: перед началом каждого такого фрагмента вставлять операции копирования c GPU на хост для используемых на хосте переменных, а после - копирование с хоста на GPU для изменившихся ячеек.
Мы выбрали третью смешанную схему, ориентирующуюся на изменения общих переменных. Синхронизации вставляются после GPU-фрагментов для измененных в данном фрагменте ячеек и после каждого изменения ячейки на хосте. Затем запускается алгоритм оптимизации операций синхронизации данных, который перемещает, выносит из циклов и удаляет лишние операции копирования. Важно при этом соблюсти следующее правило:
Пусть device1=CPU, device2=GPU или de-vice1=GPU, device2=CPU. Тогда на каждом пути графа потока управления, который соединяет генератор на devicel и вхождение той же переменной на device2, должна располагаться хотя бы одна операция копирования этих данных с devicel на de-vice2, срабатывающая раньше любой операции копирования в противоположном направлении, если последняя присутствует на данном пути.
В ДВОР применяется следующий алгоритм оптимизации размещения операций копирования некоторой переменной x:
для всех операторов cudaMemcpy, копирующих х из devicel в device2
двигаем cudaMemcpy вниз по графу потока
управления: {
while (true) {
если cudaMemcpy находится не в конце блока {
если следующий оператор не содержит вхождение текущей ячейки памяти на device2, переместить cudaMemcpy за следующий оператор,
в противном случае break }
в противном случае (cudaMemcpy находится в
конце блока) {
(этот блок может являться веткой оператора if, телом цикла или просто оператором-блоком. Обозначим ParentStmt - родительский оператор - цикл, оператор if или оператор-блок)
Если внутри ParentStmt есть вхождения текущей ячейки памяти на device2 и они доступны по графу потока управления из рассматриваемой cudaMemcpy, тогда break в противном случае поместить операцию cudaMemcpy непосредственно после Parent-
Stmt }
} }
Несколько одинаковых операторов cudaMemcpy, оказавшихся в одном месте, сливаются в один. Данный алгоритм позволяет выносить операции cuda-
Memcpy из циклов, когда это возможно, что существенно влияет на скорость работы сгенерированной программы.
Рассмотрим пример программы - LU-разложение матрицы A:
initMatrix(A);
for (int k = 0; k < N; k++) {
//fragment 1 begin
for (int j = k + 1; j < N; j++)
A[k*N+j] = A[k*N+j] / A[k*N+k]; //fragment 1 end
//fragment 2 begin for(int i = k + 1; i < N; i++) for (int j = k + 1; j < N; j++) A[i*N+j] = A[i*N+j] -
- A[i*N+k] * A[k*N+j];
//fragment 2 end }
Цикл со счетчиком k не является распараллеливаемым. Два фрагмента этого цикла помечены для выполнения на ускорителе. Анализ ячеек памяти определяет, что оба фрагмента распараллеливаемы. Общие переменные для CPU и GPU будут классифицированы следующим образом: k, N - внешние скалярные переменные, значения которых не меняется (тип 1), A - массив (тип 3). Функция копирования массива А с хоста в GPU будет расположена после "initMatrix(A)". Сначала копирования cudaMemcpy в обратном направлении будут размещены после первого и второго фрагментов. Затем алгоритм оптимизации расположения операций cudaMemcpy вынесет их из цикла и сольет вместе в одну операцию. В результате генерации CUDA-кода получим следующую программу:
gpuErrchk(cudaMalloc(&A_gpu, N*N*8))); gpuErrchk(cudaMemcpy(A_gpu, A, N*N*8, cudaMemcpyHostToDevice));
for (int k = 0; k < N; k = k + 1) {
if (N - (k + 1) > 0) {
int blockDim = 256; int gridDim =
= (N-k-1+blockDim-1) / blockDim; kernel0<<<gridDim,blockDim>>>(k, A_gpu, n);
gpuErrchk(cudaPeekAtLastError()); }
if (N - (k + 1) > 0 && N -
- (k + 1) > 0) {
dim3 blockDim = dim3(64,4,1); dim3 gridDim = dim3((N-k-1+ +blockDim.x-1)/blockDim.x, (N-k-1+blockDim.y-1)/blockDim.y, 1); kernel1<<<gridDim, blockDim>>>(k, A_gpu, N);
gpuErrchk(cudaPeekAtLastError()); }
}
gpuErrchk(cudaMemcpy(A, A_gpu, N*N*8, cudaMemcpyDeviceToHost)) ; gpuErrchk(cudaFree(A_gpu));
где kernelO и kernell определены следующим образом:
_global_ void kernel0(int k,
double *A_gpu, int N) {
int j = blockldx.x * blockDim.x + +threadIdx.x; if (j < N - (k + 1)) A_gpu[k*N+k+1+j]=
=A_gpu[k*N+k+1+j]/A_gpu[k*N+k]; }
_global_ void kernel1(int k,
double *A_gpu, int N) {
int j = blockldx.y * blockDim.y + + threadldx.y;
int i = blockldx.x * blockDim.x + + threadldx.x;
if (i < N - (k + 1) && j < N - (k + 1)) A_gpu[(k+1+i)*N+k+1+j] = = A_gpu[(k+1+i)*N+k+1+j] -
A_gpu[(k+1+i)*N+k] * A_gpu[k*N+k+1+j]; }
Результаты экспериментов
Для сравнения возможностей диалогового и автоматического подходов к распараллеливанию были выбраны бенчмарки Polybench 3.2. На рис. 3 представлены результаты ускорения программ с помощью генератора CUDA-кода ДВОР и полностью автоматического генератора PPCG. Ускорение отсчитывалось от времени работы исходной программы на центральном процессоре Intel Core i7-3820 CPU @ 3.60GHz с памятью DDR3 1600MHz, откомпилированной gcc с флагом оптимизации -O2. Распараллеленные программы запускались на ускорителе NVIDIA Tesla C2075. Компиляция бенчмарков осуществлялась с директивой LARGE_DATASET.
Каждый бенчмарк был проанализирован в ДВОР. Были найдены распараллеливаемые циклы, применены оптимизирующие преобразования, помечены фрагменты для запуска на ускорителе и запущен генератор CUDA-кода. Преобразованные версии бенчмарков с размеченными с помощью прагм GPU-фрагментами кода доступны через веб-интерфейс [10].
К бенчмаркам были применены следующие преобразования:
- correlation - к последнему циклу внутри #pragma scop region применено разрезание, отщепив оператор symmat[j1][j1] = 1.0;
- atax - ко второму циклу внутри #pragma scop region применено разрезание на три цикла, в последнем из которых потом переставлены операторы for;
- bieg - аналогично предыдущему;
- doitgen - нетесное гнездо разрезано на три трехмерных тесных гнезда;
- adi - переставлены циклы в четвертом и шестом подгнездах;
- fdtd-apml - двумерное гнездо циклов по iz и iy разрезано на три части, и в получившемся втором гнезде переставлены внутренние циклы по iy и ix.
10
1
0.1
Рис. 3. Ускорение бенчмарков Polybench по сравнению с CPU после генерации CUDA кода ДВОР и PPCG
Бенчмарки: covariance, 2mm, 3mm, cholesky, gemm, mvt, gemver, gesummv, syrk, syr2k, gramschmidt, lu, ludcmp, reg detect, fdtd-2d, jacobi-ld-imper, jacobi-2d-imper были переданы в генератор CUDA-кода без трансформаций. Программы: trisolv, trmm, durbin, dynprog, oyd-warshall, seidel-2d не были распараллелены.
Из рис. 3 видно, что ДВОР генерирует более быстрый код для программ с многомерными распараллеливаемыми циклами. Во многих тестах (covariance, lu, ludcmp, 3mm) сгенерированный PPCG код страдает от множества синхронизации из-за программного управления кэшем, большого количества арифметических операций в полученных после тайлинга циклах и индексных выражениях. В этих тестах аппаратный тайлинг циклов и управление кэшем показывают лучшую производительность и приводят к значительно меньшему размеру кода kernel-функции. Небольшой размер кода уменьшает количество используемых регистров, и большее количество блоков параллельных потоков объединяются видеокартой для выполнения на одном мультипроцессоре. Например, дольше всего работающая kernel-функция в сгенерированном с помощью PPCG LU-разложении требует 35 регистров против 13 у программы, сгенерированной с помощью ДВОР. Профилировщик NVIDIA показывает недостаточную загруженность мультипроцессора: 65 % у программы, сгенерированной PPCG, против 99 % - у сгенерированной ДВОР. Тем не менее во многих случаях программное управление кэшем и тайлинг циклов PPCG дают более высокую производительность (adi, gemm, reg detect, syr2k, syrk, 2mm).
В тесте 3mm PPCG не оптимизирует сумму, накапливая ее в локальной переменной. Генератор PPCG
пытается распараллелить каждую программу, и иногда это приводит к замедлению по сравнению с CPU. Замедление наблюдается в тестах durbin, dynprog, oyd-marshall, gramschmidt, ludcmp, symm, trisolv, trmm. В некоторых случаях сгенерированный код медленнее у обоих генераторов: cholesky, fdtd-ampl, gesummv.
Неупомянутые тесты работают практически одинаковое время.
Литература
1. Диалоговый высокоуровневый оптимизирующий распа-раллеливатель программ. URL: http://ops.rsu.ru/about.shtml (Дата обращения: 15.03.2015).
2. Lidman J., Quinlan D.J., Liao C., McKee S.A. Rose: fttrans-form - a source-to-source translation framework for exascale fault-tolerance research. In: Dependable systems and networks workshops. 2012. Р. 1 - 6.
3. Liao S., Diwan A., Bosch R.P., Ghuloum A., Lam M.S. SUIF Explorer: an interactive and interprocedural parallelizer. SIGPLAN Not. 1999. Vol. 34. P. 37 - 48.
4. Lee S.I., Johnson T., Eigenmann R. Cetus - an extensible compiler infrastructure for source-to-source transformation. In: Proc. Workshops on Languages and Compilers for Parallel Computing. 2003.
5. Verdoolaege S., Juega J.C., Cohen A., Gomez J.I., Tenllado Ch., Catthoor F. Polyhedral parallel code generation for CUDA. ACM Trans. Archit. Code Optim, 2013. № 9.
6. Amini M., Creusillet B., Even S., Keryell R., Goubier O., Guelton S., McMahon J.O., Pasquier F.X., Pean G.,
Villalon P. Par4All: From convex array regions to heterogeneous computing. In: 2nd International Workshop on Polyhedral Compilation Techniques, Impact. 2012.
7. Balasundaram V., Kennedy K., Kremer U., McKinley K. Subhlok J.: The parascope editor: an interactive parallel programming tool. In Proceedings of the 1989 ACM/IEEE conference on Supercomputing, ACM, New York. 1989. P. 540 - 550.
8. Feautrier P. Parametric Integer Programming. RAIRO Rech-
erche Opérationnelle. 1988. № 22, Р. 243 - 268.
9. Pai S., Govindarajan R., Thazhuthaveetil M.J. Fast and efficient automatic memory management for GPUs using compiler-assisted runtime coherence scheme. In Proceedings of the 21st international conference on Parallel architectures and compilation techniques, ACM. New York, 2012. Р. 33 - 42.
10. Веб-распараллеливатель. URL: http://ops.opsgroup.ru/ (Дата обращения: 12.03.2015).
References
1. Dialogovyj vysokourovnevyj optimiziruyuschij rasparallelivatel' programm [Interactive high-level optimizing Parallelizer programs.]. Available at: http://ops.rsu.ru/about.shtml (accessed 15.03.2015).
2. Lidman, J., Quinlan, D.J., Liao, C., McKee, S.A.: Rose::fttransform - a source-to-source translation framework for exascale fault-tolerance research. In Dependable systems and networks workshops, 2012, pp. 1-6.
3. Liao, S., Diwan, A., Bosch, R.P., Ghuloum, A., Lam M.S.: SUIF Explorer: an interactive and interprocedural parallelizer. SIGPLAN Not. 34, 37-48. 1999.
4. Lee, S.I., Johnson, T., Eigenmann, R.: Cetus - an extensible compiler infrastructure for source-to-source transformation. In Proc. Workshops on Languages and Compilers for Parallel Computing. 2003.
5. Verdoolaege, S., Juega, J.C., Cohen, A., Gomez, J.I., Tenllado, Ch., Catthoor, F.: Polyhedral parallel code generation for CUDA. ACM Trans. Archit. Code Optim, 9. 2013.
6. Amini, M., Creusillet, B., Even, S., Keryell, R., Goubier, O., Guelton, S., McMahon, J.O., Pasquier, F.X., Pean, G., Villalon, P.: Par4All: From convex array regions to heterogeneous computing. In: 2nd International Workshop on Polyhedral Compilation Techniques, Impact. 2012.
7. Balasundaram, V., Kennedy, K., Kremer, U., McKinley, K. Subhlok, J.: The parascope editor: an interactive parallel programming tool. In Proceedings of the 1989 ACM/IEEE conference on Supercomputing, C. 540-550 ACM, New York. 1989.
8. Feautrier, P.: Parametric Integer Programming. RAIRO Recherche Op'erationnelle. 22, 243-268. 1988.
9. Pai, S., Govindarajan, R., Thazhuthaveetil, M.J.: Fast and efficient automatic memory management for GPUs using compilerassisted runtime coherence scheme. In Proceedings of the 21st international conference on Parallel architectures and compilation techniques, С. 33-42. ACM, New York. 2012.
10. Veb-rasparallelivatel [Web Parallelizer]. Available at: http://ops.opsgroup.ru/ (accessed 12.03.2015).
Поступила в редакцию 8 апреля 2015 г.
