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.

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

Рассмотрим на простом примере как можно ускорить перемножение матриц:
  1. #include <openacc.h>
  2. #include <stdio.h>
  3. #include <stdlib.h>
  4. void main() {
  5.   int n = 100;
  6.   float a[n][n];
  7.   float b[n][n];
  8.   float c[n][n];
  9.   float elements [n];
  10.   for(int i = 0; i < n; i++)
  11.    for (int j=0; j<n; j++){
  12.     a[i][j] = i+j;
  13.     b[i][j] = 100 + 2 * i;
  14.   }
  15. #pragma acc kernels loop independent
  16.   for(int i = 0; i < n; i++)  
  17.     for (int j=0; j < n; j++){
  18.         for (int k=0; k<n; k++)
  19.                 c[i][j]=+a[i][k]*b[k][j];
  20.         }      
  21.   free(a); free(b); free(c);
  22. } // main
Эта программа отличается от простой версии, выполняемой на одном ядре 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

Tags: 

Comments

От редакции сайта :)

В надежде что автор прочтет комментарии, задам вопросы

1) Насколько ускорилось умножение матриц в данном примере. Тесты ведь были?
2) Насколько автоматическое ускорение оказалось медленнее, чем использование библиотечной процедуры DGEMM? Из, к примеру, Magma?

От автора:)

Спасибо за вопросы, пока не могу ответить - когда писал статью кончилась лицензия на сервере, а сейчас после восстановления перестали работать все примеры. Компиляция идет, а при запуске все просто висит. Думаю эту проблему мы в скором времени решим!

Умножение матриц - отличный

Умножение матриц - отличный пример для изучения, потому что "наивные" методы работают очень плохо.
Скажем, разница в быстродействии между блочной реализацией (Intel MKL) и "наивными" тремя вложенными циклами - у меня получалась больше двух порядков.

На GPU, так как относительная латентность памяти еще больше, разница тоже должна быть изрядной.

Как следствие, интересно было бы посмотреть, как выглядит блочный вариант с OpenACC.

код пахнет  free(a);

код пахнет

  1. ...
  2.   int n = 100;
  3.   float a[n][n];
  4.   float b[n][n];
  5.   float c[n][n];
  6. ...
  7.   free(a); free(b); free(c);

на нормальном компиляторе,

на нормальном компиляторе, как и ожидается - crash:
http://ideone.com/KzPvz

*** glibc detected *** ./prog: free(): invalid pointer: 0xbfa29710 ***
======= Backtrace: =========
/lib/libc.so.6[0xb76ddfd4]
/lib/libc.so.6(cfree+0x9c)[0xb76df87c]
./prog[0x8048c50]
/lib/libc.so.6(__libc_start_main+0xe5)[0xb7689725]
./prog[0x8048361]
======= Memory map: ========

Все хорошо только на словах,

Все хорошо только на словах, и возможно будет хорошо в будущем. Пока же PGI полон багов и не реализованных фич. По крайней мере версии для Fortran (http://www.pgroup.com/userforum//viewforum.php?f=12) с которой мне в основном и приходится работать .Не могу такого сказать про CRAY и CAPS, не работал.

Да и не любой код так легко переносится... практически всегда его приходится адаптировать к тому, чтобы его можно было распараллелить. Вы обратите внимание так же на то, что в сети присутствуют только примитивные примеры и (возможно пока) нет ни одной полноценной программы с применением OpenACC.

Ну это же рассчитано на то,

Ну это же рассчитано на то, чтобы очень малыми усилиями получить какое-то ускорение. В таком режиме - может и сработать, на мой взгляд.

Насчет багов - а где их нет?

Насчет багов - а где их нет? Только в gcc наверное. Со временем поправят, может и какой-нибудь open-source компилятор допилят, будет общедоступно.

На сайте NVIDIA есть раздел, посвященный этому: http://www.nvidia.ru/object/openacc-stories-ru.html . Плюс, там объявлен конкурс, можно получить бесплатную бессрочную лицензию PGI.