OpenACC: Вычисления на GPU с помощью простых директив

В ноябре 2011 года был анонсирован стандарт OpenACC ? совместное детище суперкомпьютерных гигантов CRAY, CAPS и PGI и лидера рынка графических процессоров NVIDIA. Сам стандарт призван значительно упростить работу программиста и создать высокоуровневую прослойку над уже известными CUDA и OpenCL.

Стоит отметить, что до недавнего времени стандарт не поддерживался в полной мере ни одним компилятором, но даже то, что уже есть, впечатляет своей простотой и результативностью. Теперь написание программы, выполняемой параллельно на тысячах ядер современных GPU не требует почти никаких усилий и практически полностью перекладывается на компилятор. Все что нужно сделать ? расставить директивы по коду на манер OpenMP. Набор директив достаточно велик (полную спецификацию можно посмотреть по ссылке) и за один день его весь не освоить, но простейшую программу можно сделать за 5 минут, особенно если есть однопоточная реализация. Отсюда и вытекает основная идея ? спрятать от разработчика почти все детали архитектуры, освободить его от тонкостей (а ведь лет шесть назад до появления CUDA использовать GPU могли только знатоки шейдеров) и оставить время на работу над научным или пользовательским проектом.

Как и его прародители (PGI accelerator и CAPS HMPP) OpenACC поддерживает языки С и Fortran. Итак, все директивы в С-версии стандарта начинаются как обычно с #pragma, далее ставится спецификатор ?acc? и одна из основных директив, дополненная одним, или несколькими условиями. Чаще всего используются 3 директивы: parallel, kernels и data.

љ

Как использовать:

Рассмотрим на простом примере как можно ускорить перемножение матриц:

Эта программа отличается от простой версии, выполняемой на одном ядре CPU только строкой 15, где мы видим директиву kernels, говорящую компилятору создать потоки, сгруппированные в несколько блоков, количество которых он выбирает на свое усмотрение. Кроме того, здесь же добавлена директива loop, после которой обязан начинаться цикл, loop служит для того, чтобы указать, как выполнять итерации цикла: independent ? независимо, seq ? последовательно.

Попробуем скомпилировать программу с помощью компилятора PGI:

pgcc -Minfo=accel -acc -ta=nvidia -o e:\1.exe e:\2.c
main:
     16, Generating copyout(c[0:100][0:100])
         Generating copyin(a[0:100][0:100])
         Generating copyin(b[0:100][0:100])
         Generating compute capability 1.0 binary
         Generating compute capability 2.0 binary
     17, Loop is parallelizable
     18, Loop is parallelizable
     20, Loop carried reuse of 'c' prevents parallelization
         Inner sequential loop scheduled on accelerator
         Accelerator kernel generated
         17, #pragma acc loop gang /* blockIdx.y */
         18, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
         20, CC 1.0 : 17 registers; 68 shared, 4 constant, 0 local memory bytes
             CC 2.0 : 19 registers; 0 shared, 84 constant, 0 local memory bytes

Здесь опция ?acc означает использование OpenACC, а ?Minfo=accel выводит на экран лог компиляции. Рассмотрим его более детально: в нашей функции main компилятор определил четыре строки, которые необходимо обработать:

• 16, где он генерирует код копирования массивов b и a на устройство и результирующего массива обратно в память CPU;
• 17 и 18 где генерируются kernel?ы (именно из них и состоят программы на CUDA) и определяется размерность сетки;
• 20, где компилятор сообщает о невозможности одновременно всем нитям использовать элемент массива ?с? для суммирования.

Лог может быть очень полезен при оптимизации программы (не все же компилятору за нас делать) и при поиске ошибок, так, например, в последних строках описано количество используемых регистров, переменных в разделяемой и константной памяти устройства и некое occupancy, введенное в обиход NVIDIA. Оно означает соотношение возможной вычислительной мощности к полученной, или попросту эффективность использования.

Немного о директивах

Рассмотрим краткое описание некоторых директив и условий к ним:

• Директива parallel указывает на необходимость распараллеливания. Компилятор, проводя анализ кода, определяет необходимость исполнения различных его частей на GPU, или на хосте.
• Директива kernels ? аналог parallel, указывает на то, что для каждого нового цикла необходимо создать отдельную __device__ функцию.
• Директива loop предшествует оператору цикла и используется для спецификации его свойств. Современные компиляторы не требуют ее явного указания.

Несмотря на всю мощь компилятора, иногда нужно подсказывать, какие данные необходимо передать с хоста на устройство и обратно, а поскольку зачастую копирование выполняется дольше расчетов, нужно заранее продумать, где и как оптимизировать доступ к данным. Все условия передачи данных требуют входные данные, выглядящие следующим образом: a[start:length], где a ? массив, или указатель на него, start ? номер стартового элемента для копирования, а length ?длина региона данных, копируемого на GPU, или с него; start и length указываются в элементах массива (для Fortran есть существенное отличие ? вместо length указывается end ? конечный элемент). Эти условия можно использовать только с директивами kernels, parallel и data region. Ниже представлены те из них, которые используются наиболее часто:

• copy ? говорит компилятору скопировать данные на устройство перед выполнением ядра и назад после его завершения.
• copyin ? указывает, что данные на GPU используются только для чтения, и нет необходимости копировать их обратно на хост.
• copyout ? данные появятся только в результате выполнения ядра на GPU и никак не зависят от предыдущих значений по этому адресу, их нужно скопировать на хост после выполнения кернела.
• create ? выделяет в памяти устройства место для данных, не требующих какого-либо копирования, например массив для хранения промежуточных результатов.
• present - подсказывает компилятору, что эти данные уже были переданы на устройство ранее. Вызывает ошибку, если данных на GPU нет.

Плюсы и минусы

Вот, как прост и неприхотлив в использовании OpenACC, как вы могли заметить, он очень сильно напоминает OpenMP (это точно неспроста) ? он задумывался как ответвление и создатели на своем официальном сайте http://openacc.org говорят о скорой его интеграции в последующие релизы OpenMP. Значит, скоро можно будет легко распараллеливать свои задачи на огромных гетерогенных кластерах, почти не имея представления об их архитектуре. К плюсам можно отнести также высокую степень абстракции и кроссплатформенность ? сразу после выхода новых архитектур необязательно переписывать весь код, большую часть компилятор сделает за нас. К примеру, CAPS HMPP уже объявил о поддержке ускорителей не только NVIDIA, но и Intel MIC и даже AMD FirePro.

Плюсов и правда много, но не может же быть все так хорошо. Давайте обратимся к минусам: самое первое, что бросается в глаза ? все компиляторы с поддержкой OpenACC стоят денег. Может для научных лабораторий лицензия и не такое уж дорогое удовольствие, но студенты вряд-ли соберутся потратиться на это. Второй минус ? производительность: ни один компилятор не сможет оптимизировать код лучше, чем это можно сделать вручную, или с использованием библиотек от NVIDIA.

В заключение можно отметить, что OpenACC и правда дает возможность по-быстрому переписать свои проекты под использование GPU и практически не требует навыков их программирования. С его помощью уже ускорены десятки проектов в областях изучения и прогнозирования поведения атмосферы, газо- и гидродинамики и финансовых потоков. Пять лет назад началась революция массивно-параллельных вычислений и на сегодня OpenACC ? лучший способ остаться на плаву, не потеряв позиции и не потратив сотни часов на изучение всех тонкостей CUDA или OpenCL.

Автор: Ивахненко Алексей, аспирант ФГБОУ ВПО ?Юго-Западный государственный университет?, преподаватель ?APPLIED PARALLEL COMPUTING? E&R Center