OpenACC (Open Accelerators)

Материал из Национальной библиотеки им. Н. Э. Баумана
Последнее изменение этой страницы: 17:02, 24 декабря 2016.
OpenACC
Openacc 1.jpg
Предыдущий выпуск: 2.5 / November 2015; 4 years ago (2015-11)
Написана на: C++, C, Fortran
Веб-сайт www.openacc.org

OpenACC — стандарт, описывающий набор директив для написания гетерогенных программ, задействующих как центральный, так и графический процессор. Используется для распараллеливания программ на языках C, C++ и Fortran. Стандарт был создан группой, в которую вошли CAPS, Cray, NVIDIA и PGI.

OpenACC предоставляет программисту набор директив компилятора, с помощью которых можно отметить части кода, которые следует выполнять на графическом процессоре, обозначить какие из переменных являются общими, а какие индивидуальными для потока и т. п. По синтаксису схож с OpenMP. Позволяет программисту абстрагироваться от таких особенностей, как инициализация графического процессора, передача данных на графический процессор и обратно и т. д.

История

Создатели OpenACC также участвуют в работе над стандартами OpenMP и планируют расширить будущие версии OpenMP для поддержки вычислительных ускорителей. В ноябре 2012 года был опубликован технический отчет для обсуждения и добавления поддержки акселераторов, произведенных не Nvidia. На конференции ISC’12 продемонстрирована работа OpenACC на ускорителях производства NVIDIA, AMD и Intel без публикации данных о производительности.

Планируется объединить спецификации OpenACC и OpenMP, включив в последний поддержку работы с ускорителями, в том числе GPU. Черновик второй версии стандарта, OpenACC 2.0 был представлен в ноябре 2012 года на конференции SC12. В стандарт были добавлены директивы управления пересылкой данных, поддержка явных вызовов функций и раздельная компиляция.

Поддержка в компиляторах

Реализация OpenACC доступна в компиляторах от PGI (с версии 12.6), Cray и CAPS.

Группа HPCTools из Университета Хьюстона добавила поддержку OpenACC в открытый компилятор OpenUH, основанный на кодах Open64.

В национальной лаборатории ORNL был разработан компилятор с открытыми исходными текстами OpenARC для языка Си, поддерживающий OpenACC версии 1.0.

Бесплатный компилятор GNU GCC поддерживает OpenACC начиная с версии 5. GCC 5 включена в Ubuntu 15.10, в ОС Fedora 22, в DragonFly BSD 4.2. В версии GCC 5.1 (22 апреля 2015 года) была добавлена библиотека поддержки openacc.h.

Использование

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

Директивы

В OpenACC описаны различные директивы компилятора (прагмы), в том числе:

# pragma acc parallel 
# pragma acc kernels 

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

Основная директива для определения и копирования данных:

# pragma acc data 

Директива, определяющая тип параллелизма в регионах parallel и kernels:

# pragma acc loop 

Директива 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

  • GPU
    • выполняет большую часть программы
    • выделяет память на ускорителе
    • инициирует копирование данных из памяти хоста в память ускорителя
    • отправляет код ядра на ускоритель
    • устанавливает ядра в очередь для исполнения на ускорителе
    • ожидает выполнения ядра
    • инициирует копирование данных из памяти ускорителя в память хоста
    • освобождает память ускорителя
  • Ускоритель
    • исполняет ядра одно за другим
    • одновременно может передавать данные между хостом и ускорителем
  • Модель исполнения OpenACC имеет три уровня: gang, worker, vector
  • На архитектуру эта модель отображается, как набор обрабатываемых элементов (PEs)
  • Каждый PE содержит много рабочих и каждый рабочий может выполнять векторные инструкции
  • Для GPU в большинстве случаев отображение происходит так: gang=block, worker=warp, vector=threads-in-warp
  • Зависит от того, как компилятор посчитает нужным

Пример

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

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 – последовательно.

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

Вот, как прост и неприхотлив в использовании OpenACC, он очень сильно напоминает OpenMP. Скоро можно будет легко распараллеливать свои задачи на огромных гетерогенных кластерах, почти не имея представления об их архитектуре. К плюсам можно отнести также высокую степень абстракции и кроссплатформенность – сразу после выхода новых архитектур необязательно переписывать весь код, большую часть компилятор сделает за нас. К примеру, CAPS HMPP уже объявил о поддержке ускорителей не только NVIDIA, но и Intel MIC и даже AMD FirePro.

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

Источники