CC BY
40
УДК 519.685.1; 004.4'42 DOI: 10.17213/0321-2653-2017-3-13-21
ДИАЛОГОВЫЙ ГЕНЕРАТОР CUDA-КОДА
© 2017 г. А.Н. Аллазов1, С.А. Гуда2, Р.И. Морылев3
1IT Solutions LLC, Баку, Азербайджан, 2Южный федеральный университет, г. Ростов-на-Дону, Россия, 3ООО Синопсис СПб, г. Санкт-Петербург, Россия
INTERACTIVE CUDA CODE GENERATOR
A.N. Allazov1, S.A. Guda2, R.I. Morilev3
1IT Solutions LLC, Baku, Azerbaijan, 2Southern Federal University, Rostov-on-Don, Russia, 3LLC Synopsys SPb, Saint Petersburg, Russia
Аллазов Афар Назим оглы - программист, IT Solutions LLC Баку, Азербайджан, E-mail: [email protected]
Гуда Сергей Александрович - канд. физ.-мат. наук, доцент, Южный федеральный университет, г. Ростов-на-Дону, Россия. E-mail: [email protected]
Морылев Роман Игоревич - программист, ООО Синопсис СПб, г. Санкт-Петербург, Россия. E-mail: [email protected]
Allazov Afar Nazim ogli - programmer, IT Solutions LLC, Baku, Azerbaijan E-mail: [email protected]
Guda Sergey Alexandrovich - Candidate of Physical and Mathematical Sciences, Associate professor, Southern Federal University, Rostov-on-Don, Russia. E-mail: [email protected]
Morilev Roman Igorevich - programmer, LLC Synopsys SPb, Saint Petersburg, Russia. E-mail: [email protected]
Портируя программы на CUDA, программист сталкивается с множеством трудностей. Ему приходится анализировать зависимости в программе, искать распараллеливаемые циклы, трансформировать код так, чтобы достичь наилучшего отображения на архитектуру видеокарты. Избежать ошибок - невозможно. Описываемый в статье Диалоговый высокоуровневый оптимизирующий распаралле-ливатель программ (ДВОР) позволяет автоматизировать разработку CUDA-программ. Диалоговый подход имеет ряд преимуществ над полностью автоматическим распараллеливанием: пользователь может выбрать последовательность преобразований программы, попробовать несколько вариантов результирующего кода, задать параметры преобразований, сравнить производительность и выбрать лучшие значения. ДВОР может автоматически находить распараллеливаемые циклы, визуализировать зависимости по данным, выполнять множество преобразований кода (расщепление тела цикла, слияние, гнездование, раскрутка, векторизация циклов, преобразование рекуррентных циклов к распараллеливаемой форме и др.), генерировать CUDA-код, автоматически определять оптимальные параметры запуска задачи на видеокарте.
Ключевые слова: генерация CUDA-кода; диалоговая оптимизация; автоматическое распараллеливание; ДВОР.
When porting C programs to CUDA a programmer faces many difficulties. He has to manually analyze dependencies, to determine parallelizable loops, to perform program transformations and to aim at optimal code to GPU-architecture map. It is extremely difficult to avoid errors. The proposed Interactive High-level Optimizing Parallelizer (IHOP) allows of automating the steps of CUDA program development. Interactive approach has several advantages over a fully automatic parallelization: the user has the ability to choose transformation sequence, to try several variants of the result program, to set transformation parameters, to compare the performance of the result parallel programs and to choose the best one. IHOP has the following features: automatic detection of parallelizable loops, data dependencies visualization, a variety of automatic program transformations (fusion, nesting, unrolling, strip mining, recurrence elimination, etc), CUDA code generation, automatic optimal parameters determination for running loop nest on GPU (block sizes and loop order).
Keywords: CUDA code generation; interactive optimization; automatic parallelization; IHOP.
ISSN 0321-2653 IZVESTIYA VUZOV. SEVERO-KAVKAZSKIIREGION.
1. Введение
В последнее время большую популярность приобрели ускорители параллельных вычислений. На рынке активно конкурируют устройства NVIDIA, AMD и Intel. Разработано множество средств, облегчающих работу программиста: технологии программирования (CUDA, OpenCL), библиотеки программ, прагмы (OpenACC, OpenMP), расширения языков и т.д. В данной статье описано одно из таких средств, основанное на диалоговом подходе к оптимизации кода. Визуальный интерфейс пользователя позволяет удобно выделять фрагменты кода, применять преобразования, делать проверки, строить графы по программе.
Состоящий из более чем 180000 строк кода на C++ кросплатформенный Диалоговый высокоуровневый распараллеливатель программ (ДВОР) [1] обладает большой базой анализаторов и преобразователей кода, многофункциональным GUI с возможностью выделения участков кода. По своим параметрам он похож на ав-тораспараллеливающие системы: Rose Compiller [2], SUIF [3], Cetus [4], PPCG [5], Par4All [6], Parascope Editor [7]. Авторы статьи реализовали в данной распараллеливающей системе автоматический генератор CUDA-кода. ДВОР использует анализ информационных зависимостей в циклах, основанный на полиэдральной модели, с собственной реализацией параметрического метода Гомори [8]. ДВОР поддерживает создание проектов с множеством файлов, позволяет выполнять разнообразные преобразования циклов, производит проверку корректности применяемых преобразований. Так же, как и менеджер памяти времени выполнения [9], ДВОР находит оптимальное расположение операций копирования данных на ускоритель и обратно, предотвращая копирование не измененных данных, интенсивные перемещения данных внутри циклов и копирование ячеек памяти, используемых только на видеокарте. В отличие от менеджера [9], ДВОР определяет все это на этапе компиляции.
Модуль статической профилировки GPU-кода позволяет находить оптимальные размеры блока потоков, запускаемых на видеокарте, и оптимальное отображение циклов гнезда на измерения пространства потоков. ДВОР полагается на встроенный тайлинг циклов, который автоматически получается в результате разбиения видеокартой потоков на блоки, и на автоматическое кеширование, в отличие от программного
TECHNICAL SCIENCE. 2017. No 3
тайлинга и управления L1-кешем в компиляторе PPCG [5]. Это позволяет снизить накладные расходы и добиться лучшей производительности и компактности кода kernel-функций. Как показывают численные эксперименты, ДВОР позволяет получать преобразованную программу в компактной форме, производительность которой сравнима с результатами распараллеливающего без участия пользователя компилятора PPCG.
К функциям автоматической генерации GPU-кода в ДВОР организован доступ через веб-интерфейс [10]. Распараллеливаемые циклы в загружаемых на сайт программах должны быть помечены директивой «pragma target».
2. Диалоговая оптимизация программ
Портирование программ на новые архитектуры требует более деликатных инструментов разработки, чем традиционные компиляторы. Не ограниченное интерфейсом командной строки тесное взаимодействие с пользователем позволяет применять оптимизирующие преобразования к отдельным фрагментам программы. Разработанный инструмент ДВОР является системой диалогового преобразования и распараллеливания для портирования существующих программ на разнообразные архитектуры, в том числе видеокарты (рис. 1). Он включает в себя множество инструментов, которые можно разделить на три класса:
1. Автоматические трансформации программ, включающие в себя проверки применимости.
2. Анализаторы распараллеливаемых циклов.
3. Графовые визуализации программ (граф вызовов, потока управления, информационных связей).
Процесс портирования программы на GPU в ДВОР является итерационным. Одна итерация состоит из следующих шагов:
1. Определение фрагментов программы, которые нужно выполнять на GPU. Для этого можно воспользоваться автоматическим анализатором распараллеливаемых циклов или выделить фрагменты вручную.
2. Генерация kernel-функций выбранных фрагментов и результирующей программы.
3. Компиляция сгенерированной программы и ее запуск для замера производительности. Если пользователь удовлетворен скоростью работы программы, то процесс портирования окончен.
ISSN 0321-2653 IZVESTIYA VUZOV. SEVERO-KAVKAZSKIIREGION. TECHNICAL SCIENCE. 2017. No 3
& W <& IHOP File Edit Project Tools Window Help
nil ... »S IT; s,.,,.i.c.j.|6pjpu^ И «"К- Л II ff;
Reprise Explorer Y& IR;correlation.с
t _built in jva f init_array f print_array f та™ -СОГГ t _built injva V polybenchj V polybenchj V polybenchj V polybenchj V polybench_ V polybench f rtclock f polybench_ ff polybenchj f polybenchj f polybenchj ff polybench f xmalloc f posix_memi ff polybench_. kernel_correlation - \\<Щ\
} mean[j] = (meantj] / float n) >
tor j - 0; j < m; j - (j + Open Selector a[i, j] - mean[j]))]
stddev[j] = 0. ■for i = 0; i < n; i = <i ■[ stddev[j] = (stddev[j] XML Dump
Build DepGraph Build CallGraph
stddev[j] = (stddev[j] / stddev[j] = sqrt Cstddev[j if stddev[j] <= eps { uni49 = 1. } ■c uni49 = stddev[j] } stddev[j] = uni49 > for i = 0; i < n; i = (i + { Arithmetic Operator Expansion Exception Generator Full inline substitution Subroutine splitting Corruptor Loop Distribution Loop Full Unrolling Loop Header Removal Loop Nesting Loop Unrolling Recurrency Elimination Strip Mining
Category Description
► ® Loops Loop iterations can be ► • Loops Loop iterations can be ► • Loops Loop is sequental
Execute on GPU Parallel execute on GPU
► • Loops Loop iterations can be xecuted independently
GPGPU The loop can be executed in parallel. Do you want to parallelize it using GPGPU computations? ► * Loops Loop is sequental ► » Loops Loop iterations can be executed independently ► » Loops Loop iterations can be executed independently ► • Loops Loop iterations can be executed independents 1
Linear view
Parser Output Issues GPGPU
Рис. 1. Интерфейс пользователя в ДВОР / Fig. 1. IHOP screenshot
4. В противном случае необходимо выполнить дополнительные преобразования для улучшения соответствия программы архитектуре GPU, включая расщепление тела цикла, перестановку циклов, гнездование и т.п. Процесс генерации и проверки программы затем повторяется.
Однако пользователь не ограничен этой схемой и может применять любые инструменты в ДВОР в любом порядке, пробуя различные комбинации преобразований для достижения наилучшей производительности.
3. Архитектура генератора CUDA-кода
На вход генератору подаются выбранные пользователем фрагменты кода двух типов:
1) гнезда циклов размерности от 1 до 3, итерации которых следует запускать в разных потоках на видеокарте;
2) последовательно следующие друг за другом операторы, принадлежащие одному блоку, которые следует запускать в один поток на видеокарте.
Фрагменты второго типа приходится запускать на видеокарте, чтобы избежать дополнительных операций обмена данных с ускорителем.
Генератор GPU-кода работает в 4 этапа:
1) инициализация параметров участков кода, предварительные проверки и преобразования;
2) анализ участков кода и определение параметров отображения кода на ускоритель;
3) GPU-преобразования кода;
4) получение текстового представления
кода.
На первом этапе выполняются предварительные преобразования: канонизация циклов, обертка фрагментов в блоки. Происходит проверка сохранения корректности программы при запуске выделенного пользователем участка кода на ускорителе: распараллеливаемы ли циклы, являются ли алиасы фортрановскими. Также происходит проверка ограничений, накладываемых на входные данные текущей версией генератора кода. На данный момент запрещены структуры, фрагменты из разных функций, массивы указателей.
На втором этапе работы генератора определяются параметры запуска выделенных фрагментов на графическом ускорителе. Для общих с ускорителем переменных необходимо определить типы и размер создаваемых GPU-переменных, способ передачи в каждую kemel-функцию, оптимальное расположение операций выделения/освобождения памяти на ускорителе и функций копирования данных на и с ускорителя (см. ниже п. 4). Для каждого гнезда циклов размерности больше 1 определяется оптимальное отображение счетчиков циклов на координаты индекса потоков, оптимальные размеры блока потоков вдоль каждого измерения (см. ниже п. 5). Когда границы изменения счетчиков внутренних циклов распараллеливаемого гнезда зависят от счетчиков внешних циклов, приходится определять их наибольшие значения. В текущей версии генератора поддерживаются линейно зависящие от счетчиков выражения с числовыми коэффициентами. На втором этапе
также определяются функции, которые необходимо пометить как_device_, так как они вызываются из помеченных для запуска на GPU фрагментов.
Третий этап работы генератора занимается собственно преобразованиями кода. Вставляются операции выбора ускорителя и инициализации контекста на видеокарте, создаются GPU-переменные, вставляются в найденные на предыдущем этапе анализа позиции операции выделения/освобождения памяти на устройстве, обмена данными с ускорителем. Из выделенных пользователем фрагментов конструируются kernel-функции, а фрагменты заменяются на их вызовы. Так, например, трехмерное гнездо циклов
for (int io = 0; i0 < n0; io++) for (int ii = 0; ii < ni; ii++) for (int i2 = 0; i2 < n2; i2++) LoopBodyBlock
заменяется на вызов kernel-функции
if (n0 > 0 && n_max1 > 0 && n_max2 > 0) {
dim3 blockDim = dim3(b0, b1f b2); dim3 gridDim = dim3(g0, gif g2); kerneli<<<gridDim, blockDim>>>(...); gpuErrChk(cudaPeekAtLastError());
}
где kernel-функция kernel 1 имеет вид
_global_ void kerneli(...)
{
iV(0) = blockIdx.x*blockDim.x + threadldx.x iV(i) = blockIdx.y*blockDim.y + threadldx.y iy(2) = blockIdx.z*blockDim.z + threadldx.z if (i0<n0 && ii<ni(i0) && i2<n2(i0,ii)) LoopBodyBlock
}
Здесь bj - оптимальные размеры блока; gi = (n_maxV(!) + bj - 1) / bi, у - оптимальное отображение измерений пространства потоков на циклы гнезда (см. ниже п. 5), i = 0,1,2. Размеры b и отображение у определяются на этапе компиляции.
На четвертом этапе работы генератора преобразованный код во внутреннем представлении переводится в текстовое представление. Система старается оставить код по возможности похожим на его изначальную форму. Используемый парсер clang добавляет в программу объявления из всех подключаемых библиотек. На данном этапе происходит анализ используемых объявлений. Неиспользуемые удаляются, а вместо унаследованных из include-файлов вставляются соответствующие директивы include.
4. Анализ данных
Для выполнения выделенных пользователем фрагментов кода на ускорителе необходимо определить переменные, которые будут использованы на двух устройствах: CPU и GPU, выделить память на GPU и обеспечить синхронизацию данных. Для корректной работы анализа необходимо, чтобы код не содержал переменных, ссылающихся на общие или перекрывающиеся области памяти. Для каждого фрагмента кода общие для CPU и GPU ячейки памяти мы делим на три непересекающихся типа:
1) внешние (для фрагмента) скалярные переменные, значения которых не меняются;
2) внешние изменяемые скалярные переменные;
3) массивы.
Операторы выделения памяти на видеокарте cudaMalloc необходимы только для ячеек второго и третьего типа. Ячейки памяти первого типа передаются в kernel-функции по значению. Локальные переменные помеченных для выполнения на GPU фрагмента кода и счетчики распараллеливаемых циклов мы относим к локальным GPU-переменным и не считаем общими с CPU.
Для каждой общей с GPU ячейки памяти необходимо создать переменную, описывающую копию данной ячейки на GPU, определить расположение операций выделения, освобождения и копирования данных. Тип создаваемой переменной равен для скаляра указателю на тип скаляра, для массивов - указателю на тип ячейки массива и для указателей совпадает с их типом.
Операторы cudaMalloc и cudaFree ДВОР располагает в функциях, помеченных для выполнения на GPU фрагментах, в тех же местах, где происходит выделение и освобождение памяти на хосте. Это определяется расположением переменных в памяти хоста: на стеке или в куче. Для переменных на стеке оператор cudaMalloc вставляется в начало блока определения, cudaFree - в места выхода из блока определения переменной. Для динамически выделенных ячеек памяти операторы cudaMalloc и cudaFree размещаются непосредственно после операторов выделения/освобождения памяти на хосте, если это происходит в текущей функции; и в начале и в местах выхода из функции - в противном случае. Размер выделенной памяти должен быть легкодоступен из помеченного для выполнения на GPU фрагмента.
Правильная расстановка операций синхронизации cudaMemcpy экземпляров одного
ISSN 0321-2653 IZVESTIYA VUZOV. SEVERO-KAVKAZSKIIREGION.
массива в памятях разных устройств (хоста и GPU) имеет большое значение. Копирование не измененных или локальных для GPU данных, особенно внутри циклов, может драматически сказаться на производительности получившейся программы. Авторы разработали несколько инструментов, которые помогают исключить избыточное перемещение данных. Предлагаемый подход работает на этапе компиляции.
Для заданной программы, некоторые фрагменты которой помечены для выполнения на GPU, можно выделить несколько наивных схем синхронизации экземпляров ячеек памяти на хосте и GPU. Первая - ориентироваться на GPU-фрагменты кода, вставляя перед каждым фрагментом операции копирования данных с хоста на GPU и после - копирования в обратном направлении для изменившихся ячеек памяти. Вторая, симметричная первой, - ориентироваться на выполняемый на хосте код: перед началом каждого host-фрагмента вставлять операции копирования с GPU на хост для используемых на хосте ячеек памяти, а после host-фрагмента - с хоста на GPU для изменившихся ячеек.
В ДВОР используется третья, смешанная схема, ориентированная на изменения общих ячеек памяти. Синхронизации вставляются после GPU-фрагментов для измененных в данном фрагменте ячеек и после каждого изменения ячейки на хосте. Затем запускается алгоритм оптимизации операций синхронизации данных, который перемещает, выносит из циклов и удаляет лишние операции копирования. Важно при этом для любого обозначения dev1=host, dev2=GPU или dev1=GPU, dev2=host соблюсти следующее правило. К примеру, рассмотрим путь на графе потока управления, соединяющий оператор на devl, изменяющий некоторую общую для devl и dev2 переменную x, и оператор на dev2 с вхождением той же ячейки памяти. На каждом таком пути должна располагаться хотя бы одна операция копирования переменной x в направлении dev1^dev2, срабатывающая раньше любой операции копирования в противоположном dev2^dev1 направлении, если последняя присутствует на данном пути.
ДВОР применяет следующий алгоритм оптимизации размещения операций копирования, соблюдающий данное правило:
Для каждого оператора cudaMemcpy копирования x с devl на dev2 {
// переместить cudaMemcpy вниз по графу потока управления
while (true)
TECHNICAL SCIENCE. 2017. No 3 {
Если cudaMemcpy находится не в конце блока {
Если следующий оператор не содержит вхождение x на dev2 переместить cudaMemcpy за следующий оператор else break
}
else // cudaMemcpy стоит в конце блока {
// этот блок может является веткой оператора if, // телом while, for или просто оператором - блоком // Обозначим ParentStmt - оператор, к которому относится блок
Если ParentStmt содержит вхождение x на dev2, которое доступно по графу потока управления из текущего расположения cudaMemcpy тогда break в противном случае разместить операцию cudaMemcpy после ParentStmt
}
}}}
Несколько одинаковых операторов cudaMemcpy, оказавшихся в одном месте, ДВОР сливает в один. Данный алгоритм позволяет выносить операции cudaMemcpy из циклов, когда это возможно, что существенно влияет на скорость работы сгенерированной программы.
Рассмотрим пример - LU-разложение матрицы A
initMatrix(A); //initialization
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 не является распараллеливаемым. Два фрагмента внутри его тела помечены пользователем для выполнения на GPU. Анализ данных генератора CUDA-кода определяет, что циклы внутри фрагментов можно распараллелить. Общие для хоста и GPU переменные будут классифицированы следующим образом: k, N - не изменяющиеся внешние для фрагмента переменные (тип 1), A - массив (тип 3). В начале работы алгоритма операция cudaMemcpy копирования массива A с хоста на GPU будет размещена сразу за вызовом initMatrix(A). Операции копирования массива A в обратном направлении - после первого и
ISSN 0321-2653 IZVESTIYA VUZOV. SEVERO-KAVKAZSKIIREGION.
второго фрагментов. Алгоритм оптимизации размещений cudaMemcpy переместит оба последних оператора за цикл по k и сольет вместе в один оператор. В результате генерации 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));
где kernel-функции kernelO и kernell определяются следующим образом:
global void kernel0(intk,double*A gpu,int N)
{
int j = blockIdx.x*blockDim.x + threadldx.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 = blockIdx.y*blockDim.y + threadldx.y; int i = blockIdx.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];
}
5. Статическая профилировка и отображение циклов на измерения пространства потоков
Для двух- и трёхмерных гнёзд циклов встает вопрос о порядке назначения координат индекса потоков счетчикам циклов. Например, для трехмерного гнезда
for (int i0 = 0; i0 < n0; io++) for (int ii = 0; ii < ni; ii++) for (int i2 = 0; i2 < n2; i2++) LoopBodyBlock
в результирующем kernel-коде счетчикам можно назначить измерения пространства потоков шестью способами:
TECHNICAL SCIENCE. 2017. No 3 тождественным
10 = blockIdx.x*blockDim.x + threadldx.x
11 = blockIdx.y*blockDim.y + threadldx.y
12 = blockIdx.z*blockDim.z + threadldx.z
или другим, определяющимся некоторой перестановкой у чисел (0, 1, 2). От этого отображения зависит наличие оптимизаций доступа к памяти, в частности, слияние запросов к памяти (см. [11]). Например, операции доступа к ячейке массива A[i2][i1][i0] или x[i2*m+i1*n+i0] для тождественного отображения у будут объединяться (coalesce), а для перестановок с у(0)^0 - возможно нет, так как потоки с соседними номерами обращаются к ячейкам из разных кэш-линеек.
Чтобы найти оптимальное отображение, был реализован статический профилировщик, оценивающий время работы программы, исходя из паттернов обращения к памяти. Рассмотрим обращение x[e0][e1]...[ek] к элементу (k+1)-мерного массива x размера d0did2\..^dk, где e0, eb...,ek - аффинные функции счетчиков с коэффициентами - произвольными выражениями. ДВОР составляет выражение, описывающее номер ячейки от начала массива M = e0did2\..^dk + + e3d3\..-dk + ... +
+ ek-1dk + ek.
В нем группируются слагаемые с одними и теми же счетчиками циклов i0, i1, i2 и определяются коэффициенты при этих счетчиках. Обозначим их a0, a1, a2. Пусть CS (Cache Size) - количество элементов массива, помещающихся в одну кэш-линейку. Коэффициенты при счетчиках мы делим на три типа:
1) нулевой;
2) литерал c: |c| < CS;
3) литерал c: |c| > CS или выражение с неизвестным на этапе компиляции значением.
Составим таблицы времени tkjm обращения к памяти в расчете на один GPU-поток для тождественного отображения у: (i0^x, i1^y, i2^z) и всех возможных комбинаций типов коэффициентов: k, j, m=1, 2, 3, где k - тип коэффициента a0, j - тип a1, m - тип a2. За единицу примем время чтения одной кэш-линейки из глобальной памяти в Ll-кеш. Предположим для простоты, что потоки, принадлежащие разным блокам, не обращаются к общим кэш-линейкам.
Результаты расчета для двумерного случая: tn = 0, ti2 = min( (CS/ai)b0, BV )-1, ti3 = b0-1, t2i = min( (CS/a0)bi, BV )-1, t22 = min( CS/(a0b0a1)b0+CS%(a0b0a1)/a0, BV )-1,
t23 = min( CS/a0, b0 )-1,
t31 =b 1 , t32 = 1 t33 = 1,
ISSN 0321-2653 IZVESTIYA VUZOV. SEVERO-KAVKAZSKIIREGION.
для трехмерного:
= 0, ^21 = тш( ^/^^2, BV )-1, 1хэ1 = (Ь^)-1, 12„ = тш( (CS/ao)blb2, BV )-1, t22l = min( (CS/(aoboal)bo+CS%(aoboal)/ao)b2, Ъ0Ъ1>-1, t231 = min( (CS/ao)Ъ2, ЪоЪ2)-1, tзl1=(ЪlЪ2)-1, tз21 = Ъ2-1,
1-331 =Ъ2 ,
г112 = min( ^ЗДоЪь ВУ)-1,
1122 = min( (CS/(alЪla2)Ъl+CS%(alЪla2)/al)Ъo, ВУ )-1,
^32 = Ъ0 ,
1212 =min( (CS/(aoЪoa2)Ъo+CS%(aoЪoa2)/ao)Ъl, ВУ )-1,
1222 = min( CS/(aoboalbla2)bobl + CS%(aoЪoalЪla2)/ (аоЪоа1)Ъо + CS%(aoalbobla2)%(aoboal)/ao), ВУ )-1,
1232 = тт( CS/ao, Ъо)"1, 1з12=Ъ1-1, 1322=1, 1332=1, 1113= (ЪоЪ1)"1,
1123 = тт( (CS/al)Ъo, ЪоЪ1 )-1, 1133 = Ъо"1,
1213 = min( ^/ао^ьЪоЪО-1,
1223 = min( CS/(aoЪoal)Ъo +CS%(aoboal)/ao, ЪоЪ1)-1,
1233 = min(CS/ao, Ъо)-1, 1э13=Ъ1-1, 1323 = 1, 1333 = 1.
Здесь через Ь0, Ъ1, Ъ2 обозначены размеры блоков потоков, ВУ - число потоков в блоке (ВУ=Ъ0Ъ1 для двумерного и ВУ=Ь0Ъ1Ъ2 - для трехмерного случаев), символ "/" обозначает деление нацело, "%" - остаток от деления.
Покажем, как рассчитывалось время доступа к памяти для одного из значений двухмерного случая a0 = 0, a1 - маленький литерал. Паттерн доступа к памяти имеет вид
0^ + . (1) Для тождественного отображения у в первом блоке потоки нумеруются выражением ^ + Ъ0^. Оценим количество потоков, которые обращаются к ячейкам массива одной кэш" линейки. В силу (1) к первой кэш"линейке массива х (предполагается, что массив выравнен) обращаются потоки с индексами: ^ - любое, i1<CS/a1. Таким образом, количество потоков в первом блоке, обращающихся к одной кэш" линейке, составит (CS/a1)b0, если b1>CS/a1 и ВУ - в противном случае. В итоге время доступа к памяти (количество чтений кэш"линеек в расчете на один поток) составит min((CS/a1)Ъ0, ВУ)-1.
Если вхождение срабатывает в одном потоке несколько раз (например, когда оно находится в цикле), то его время доступа к памяти умножается на количество срабатываний (в случае неизвестного числа срабатываний статический профилировщик использует значения по умолчанию).
Пользуясь рассчитанными временами ^ или статический профилировщик перебирает все возможные перестановки у и все размеры блоков, являющиеся степенями двойки, для фик-
TECHNICAL SCIENCE. 2017. No 3
сированного объема блока BV, и находит те параметры, для которых суммарное время обращения к памяти меньше всего.
6. Результаты экспериментов
Для сравнения возможностей диалогового и автоматического подходов к распараллеливанию были выбраны бенчмарки Polybench 3.2. На рис. 2 представлены результаты ускорения программ с помощью генератора CUDA-кода ДВОР и автоматического генератора PPCG (см. [5]). Ускорение отсчитывалось от времени работы исходной программы на центральном процессоре Intel Core i7-3820 CPU 3.60GHz с памятью DDR3 1600MHz, откомпилированной gcc с флагом оптимизации -O2. Распараллеленные программы запускались на ускорителе NVIDIA Tesla C2075. Время работы бенчмарков замерялось для директивы LARGE_DATASET.
Рис. 2. Ускорение бенчмарков Polybench по сравнению с CPU после генерации GPU-кода при помощи ДВОР и PPCG / Fig. 2. Polybench tests speedup over CPU after parallelizing by IHOP and PPCG
Каждый бенчмарк проанализирован в ДВОР. Были найдены распараллеливаемые циклы, применены оптимизирующие преобразования, помечены фрагменты для запуска на ускорителе и запущен генератор CUDA-кода. Преобразованные версии бенчмарков с отмеченными для выполнения на GPU фрагментами кода доступны через веб-интерфейс [10].
К бенчмаркам были применены следующие преобразования:
- correlation - к последнему циклу внутри #pragma scop мы применили разрезание, отщепив оператор symmat[j 1][j1]=1;
- atax - ко второму циклу внутри #pragma scop мы применили разрезание на три цикла, в последнем из которых потом переставили операторы for;
- bicg - аналогично предыдущему;
- doitgen - нетесное гнездо циклов внутри #pragma scop разрезано на три трехмерных тесных гнезда;
- adi - переставлены циклы в четвертом и шестом гнездах;
- fdtd-apml - двумерное гнездо циклов по iz и iy разрезано на три части.
Бенчмарки: covariance, 2mm, 3mm, cholesky, gemm, mvt, gemver, gesummv, syrk, syr2k, gramschmidt, lu, ludcmp, reg\_detect, fdtd-2d, jacobi-ld-imper, jacobi-2d-imper были поданы на вход генератору CUDA-кода без предварительных преобразований. Программы symm, trisolv, trmm, durbin, dynprog, floyd-warshall, seidel-2d не были распараллелены.
Анализ рис. 2 показывает, что благодаря статической профилировке ДВОР генерирует гораздо более быстрый код для программ с многомерными гнездами циклов (doitgen). Во многих случаях (covariance, lu, ludcmp, 3mm) код, сгенерированный PPCG, выполняется медленнее из-за обилия синхронизаций при управлении кешированием и огромного количества арифметических операций в сгенерированных тайлин-гом циклах и индексных выражениях. Для этих программ встроенный тайлинг (путем разделения потоков на блоки) и автоматическое управление Ll-кешем оказываются эффективнее и приводят к коду kernel-функций гораздо меньшего объема. Небольшой код использует немного регистров и бо'льшее число блоков может быть запущено одновременно на одном мультипроцессоре. Например, самое горячее гнездо циклов в сгенерированном PPCG коде LU-разложения использует 35 регистров против 14 регистров - после генерации CUDA-кода с помощью ДВОР. Профилировщик показывает недостаточную загруженность мультипроцессора (процент времени, когда хотя бы один варп активен): 65 % у PPCG против 99 % у сгенерированного ДВОР кода. Несмотря на это, во многих бенчмарках программное управление кэшем и тайлинг PPCG оказались лучше подхода ДВОР (adi, gemm, reg_detect, syr2k, syrk, 2mm).
В программе 3mm PPCG не смог оптимизировать накопление суммы в локальной переменной. PPCG пытается распараллелить любую программу, зачастую это приводит к замедлению кода по сравнению с CPU: durbin, dynprog, floyd-marshall, gramschmidt, ludcmp, symm, trisolv, trmm. В некоторых случаях сгенерированный код оказался медленнее CPU-варианта для обоих генераторов: cholesky, fdtd-ampl, gesummv. Не
упомянутые тесты продемонстрировали одинаковую производительность.
Заключение
В данной работе предлагается новый инструмент ДВОР для диалоговой генерации CUDA-кода. Он помогает программисту автоматически анализировать зависимости в циклах, находить распараллеливаемые циклы, выполнять преобразования программ и генерировать CUDA-код для выделенных фрагментов. На этапе компиляции генератор CUDA-кода находит оптимальное расположение операций копирования данных на ускоритель и обратно, предотвращая избыточные перемещения не измененных данных, интенсивные копирования внутри циклов и копирование локальных для GPU массивов. Модуль статической профилировки GPU-кода автоматически находит оптимальные размеры блока потоков и отображение циклов на измерения пространства потоков. ДВОР полагается на встроенный тайлинг циклов и автоматическое кеширо-вание в отличие от программного в PPCG, уменьшая таким образом накладные расходы и добиваясь лучшей производительности и компактности кода kernel-функций.
Эксперименты, проведенные для бенчмар-ков Polybench, показали, что ДВОР генерирует программы в форме, больше соответствующей исходной, производительность которых сравнима с полностью автоматическим распараллеливающим компилятором PPCG. Благодаря модулю статической профилировки ДВОР лучше работает с многомерными гнездами циклов, так как определяет оптимальные размеры блока потоков и оптимальное отображение циклов на измерения пространства потоков. Благодаря компактности кода сгенерированные ДВОР программы потребляют меньше регистров и оказываются быстрее, чем сгенерированные PPCG. Несмотря на это, в некоторых случаях программное управление кэшем и тайлинг PPCG оказываются эффективнее.
Литература
1. ДВОР. URL: http://ops.rsu.ru/about.shtml (дата обращения 01.02.2017)
2. Lidman J., Quinlan D.J., Liao C., McKee S.A.: Rose: FTTransform - a source-to-source translation framework for exascale fault-tolerance research // Dependable systems and networks workshops. 2012. P. 1 - 6.
ISSN 0321-2653 IZVESTIYA VUZOV. SEVERO-KAVKAZSKIIREGION.
3. Liao S., Diwan A., Bosch R.P., Ghuloum, A., Lam M.S. : SUIF Explorer: an interactive and interprocedural parallelizer // SIGPLAN, 1999. № 34. P. 37 - 48.
4. Lee S.I., Johnson T., Eigenmann R.: Cetus - an extensible compiler infrastructure for source-to-source transformation // 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. Vol. 9, № 4. P. 54:1 - 54:23.
6. Amini M., Creusillet B., Even S., Keryell R., Goubier O.,
Guelton S., McMahon J.O., Pasquier FX., Pean G., Villalon P. Par4All: From convex array regions to heterogeneous computing // 2nd International Workshop on Polyhedral Compilation Techniques, Impact. 2012.
TECHNICAL SCIENCE. 2017. No 3
7. Balasundaram V., Kennedy K., Kremer U., McKinley K. Subhlok J. The parascope editor: an interactive parallel programming tool // Proceedings of the 1989 ACM/IEEE conference on Supercomputing. New York: ACM. P. 540 - 550.
8. Feautrier P.: Parametric Integer Programming // RAIRO Recherche Operationnelle. 1988. Vol. 22. P. 243 - 268.
9. Pai S., Govindarajan R., Thazhuthaveetil M.J. Fast and efficient automatic memory management for GPUs using compiler-assisted runtime coherence scheme // Proceedings of the 21st international conference on Parallel architectures and compilation techniques, 2012. New York:ACM. P. 33 - 42.
10. Веб-распараллеливатель. URL: http://ops.opsgroup.ru (дата обращения 01.02.2017)
11. CUDA C Programming Guide. URL: http://docs.nvidia.com/cuda/cuda-c-programming-guide (дата обращения 01.02.2017)
References
1. DVOR. Available at: http://ops.rsu.ru/about.shtml (accessed 01.02.2017)
2. Lidman J., Quinlan, D.J., Liao C., McKee S.A. Rose: FTTransform - a source-to-source translation framework for exascale fault-tolerance research // 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. 1999. № 34. Pp. 37 - 48.
4. Lee S.I., Johnson T., Eigenmann R. Cetus - an extensible compiler infrastructure for source-to-source transformation // 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. Vol. 9, № 4. Pp. 54:1 - 54:23.
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 // 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 // Proceedings of the 1989 ACM/IEEE conference on Supercomputing. New York: ACM. Pp. 540 - 550.
8. Feautrier P. Parametric Integer Programming // RAIRO Recherche Operationnelle. 1988. Vol. 22, Pp. 243 -268.
9. Pai S., Govindarajan R., Thazhuthaveetil M.J. Fast and efficient automatic memory management for GPUs using compiler-assisted runtime coherence scheme // Proceedings of the 21st international conference on Parallel architectures and compilation techniques, 2012. New York. ACM. Pp. 33 - 42.
10. Veb-rasparallelivatel'. Available at: http://ops.opsgroup.ru (accessed 01.02.2017)
11. CUDA C Programming Guide. Available at: http://docs.nvidia.com/cuda/cuda-c-programming-guide (accessed 01.02.2017)
Поступила в редакцию /Receive 03 мая 2017 г. /May 03, 2017
