OpenACC

Amb l'objectiu de garantir una àmplia disponibilitat a totes les arquitectures informàtiques disponibles al moment de la seva creació però també assegurant-se les d'un futur, OpenACC defineix un model abstracte per a computació accelerada.

Aquest model exposa múltiples nivells de paral·lelisme que poden aparèixer en un processador, així com una jerarquia de memòries amb diferents graus de velocitat i adreçabilitat. Al seu nucli, OpenACC suporta la descàrrega tant de dades com de computació des d'un dispositius amfitrió a un dispositiu accelerador. En aquest moment és quan l'arquitectura dels dispositius pot ser la mateixa o completament diferent, com el cas d'un host CPU i un accelerador GPU, a més de tenir un únic espai de memòria o separada. En aquest segon cas el compilador d'OpenACC analitzarà el codi en temps d'execució i manejarà qualsevol gestió de memòria i transferència de dades entre el host i la memoria del dispositiu.

La següent figura mostra un diagrama d'alt nivell de l'accelerador abstracte OpenACC, encara que hem de tenir present que els dispositius i les memòries poden ser físicament iguals en algunes arquitectures.^[2]

Els compiladors comercials que donen suport a OpenACC son PGI (a partir de la versió 12.6) i Cray (només per a maquinari Cray).

OpenUH és un compilador Open64 de codi obert que dona suport a C i FORTRAN, desenvolupat pel grup HPCTools de la Universitat d'Houston.

OpenARC^[3] és un compilador de codi obert C desenvolupat a Oak Ridge National Laboratory per donar suport a totes les característiques de OpenACC 1.0. Un compilador de codi obert experimental, accULL, està sent desenvolupat per la Universitat de La Laguna (només en el llenguatge C).

IPMACC^[4] és un compilador de codi obert C desenvolupat per la Universitat de Victoria que tradueix de OpenACC a CUDA, OpenCL i ISPC. Actualment, només s'admeten les següents directives: data, kernels, loop and cache.

El suport de GCC a OpenACC va ser lent en arribar. Una implementació d'orientació GPU des de Samsung es va anunciar el setembre de 2013; això va traduir el codi obert OpenACC 1.1 a OpenCL. L'anunci d'una implementació "real" va arribar dos mesos més tard, aquesta vegada des de NVIDIA i basat en OpenACC 2.0. Això va generar certa controvèrsia, ja que la implementació només s'orientaria al propi llenguatge de muntatge PTX de NVIDIA, per a la qual no hi havia cap assemblador de codi obert o el temps d'execució disponible. El suport experimental per OpenACC / PTX va acabar en GCC a partir de la versió 5.1. Les sèries de llançament de GCC6 i GCC7 inclouen una implementació molt millorada de les especificacions OpenACC 2.0a.

D'una manera similar a OpenMP 3.x en un sistema homogeni o en l'anterior OpenHMPP, la principal forma de programar en OpenACC és mitjançant directives.^[5] Les especificacions també inclouen una biblioteca en temps d'execució que defineix diverses funcions de suport. Per fer-les servir, l'usuari ha d'incloure "openacc.h" a C o "openacc_lib.h" a Fortran;^[6] i després cridar a la funció acc_init ().

Aquestes directrius del compilador especifiquen bucles i regions de codi en C, C ++ i Fortran estàndard per ser executades en paral·lel, ja sigui en una CPU multinucli, o bé per ser descarregades i executades en paral·lel en un dispositiu accelerador que es trobi connectat, proporcionant així portabilitat tant entre sistemes operatius com CPUs i acceleradors.

La gran majoria de les directrius s'apliquen al bloc estructurat o bucle que es troba immediatament seguit. Quan parlem d'un bloc estructurat en C i C++ ens referim a una sola declaració o una declaració composta, i en el cas de Fortran a una seqüència de sentències, amb un sol punt d'entrada a la part superior i una sola sortida a la part inferior.^[7]

OpenACC defineix una extensa llista de directives,^[1] on la seva sintaxi general és la següent:

En C/C++:

#pragma acc directive [clause [[,] clause]...] new-line

En Fortran:

!$acc directive [clause [[,] clause]...]

En OpenACC una directiva construct s'aplica de manera immediata al bucle o bloc estructurat següent. Aquestes directives definides com "Compute Construct" són parallel, kernels i serial.

La directiva parallel és una construcció paral·lela que llança una sèrie de gangs que s'executen en paral·lel, on cada gang pot donar suport a múltiples treballadors, cadascun amb operacions vectorials o SIMD, permetent així definir els nuclis de computació paral·lela en l'accelerador.

En C/C++:

#pragma acc parallel
{ structured block }

En Fortran:

!$acc parallel
structured block
!$acc end parallel

La directiva kernels defineix la regió del programa que s'executarà al dispositiu accelerador, normalment compilada com una seqüència d'operacions del kernel.

En C/C++:

#pragma acc kernels [clause [[,] clause]…] new-line
{ structured block }

En Fortran:

!$acc kernels [clause [[,] clause]…]
structured block
!$acc end kernels

La directiva serial envolta els bucles o el codi que s'ha d'executar en sèrie al dispositiu.

En C/C++:

#pragma acc serial [clause [[,] clause]…] new-line
{ structured block }

En Fortran:

!$acc serial [clause [[,] clause]…]
structured block
!$acc end serial

Les directives de Compute Constructs definides, és a dir, parallel, kernels i serial, suporten una sèrie de clàusules de còmput, i altres de dades les quals modifiquen el recompte de referències estructurades per a les dades associades. Les tres Compute Constructs comparteixen les clàusules de dades que suporten, a diferència de les de còmput, que no les tres suporten les mateixes.^[7] Alguns examples de les clàusules de càlcul serien les següents:

Quan la condició és nonzero o TRUE la regió de còmput s'executarà en el dispositiu, en cas contrari, serà executada pel thread local.

if(condition)

Quan la condició no està present o quan la condició és nonzero o TRUE, la regió de còmput s'executarà al dispositiu local, en cas contrari, serà executada al dispositiu actual.

self [(condition)]

Prevé que el compilador determini implícitament atributs de dades per a qualsevol variable utilitzada o assignada a la construcció.

default(none)

Assumeix implícitament que no hi ha cap dada no escalable que no hagi estat especificada en una clàusula de dades.

default(present)

La regió de computació s'executa de manera asíncrona amb el thread local a la cua asíncrona corresponent. Sense argument, la regió de comput s'executarà a la cua asíncrona predeterminada.

async [(expression)]

La regió de computació no començarà a executar-se fins que totes les accions de la cua o cues asíncrones corresponents no es completin. Sense cap argument, la regió de computació esperarà a totes les cues asíncrones.

wait [(expression-list)]

Aquesta directiva de dispositiu defineix una regió del programa dins de la qual el dispositiu accelerador pot accedir a les dades. És útil per a la copia de dades a l'accelerador, i viceversa.

En C/C++:

#pragma acc data [clause[[,] clause]…] new-line
{ structured block

En Fortran:

!$acc data [clause[[,] clause]…]
structured block
!$acc end data

Algunes de les clàusules de dades mencionades prèviament serien les següents:

Copy: Només estructurat. Assigna memòria al dispositiu i copia les dades de l'host al dispositiu en entrar a la regió i copia les dades a l'amfitrió quan surti de la regió.

copy(list)

Dades d'entrada: Defineix l'inici d'una vida de dades no estructurada.

Copyin: Assigna memòria al dispositiu i copia dades de l'host al dispositiu en entrar a la regió.

copyin([readonly:] list)

Create: Assigna memòria al dispositiu però no la copia.

create(list)

Dades de sortida: Defineix el final d'una vida de dades no estructurada.

Copyout: Assigna la memòria al dispositiu i copia les dades a l'amfitrió quan surti de la regió.

copyout(list)

Delete: Es distribueix la memòria del dispositiu sense copiar-la.

delete(list)

També trobem algunes funcions API en temps d'execució definides: acc_get_num_devices(), acc_set_device_type(), acc_get_device_type(), acc_set_device_num(), acc_get_device_num(), acc_async_test(), acc_async_test_all(), acc_async_wait(), acc_async_wait_all(), acc_init(), acc_shutdown(), acc_on_device(), acc_malloc(), acc_free().

OpenACC generalment s'encarrega d'organitzar del treball per al dispositiu objectiu, però això pot ser sobreescrit mitjançant gangs i workers (bandes i treballadors). Un gang es compon de workers i opera sobre diversos elements de processament (com amb un grup de treball en OpenCL).

El següent codi mostra com es pot accelerar mitjançant OpenACC la multiplicació de matrius:

#include <iostream>
#include <stdio.h>
#include <stdlib.h>

#define size 1000

float a[size][size];
float b[size][size];
float c[size][size];

int main()
{
 int i, j, k;

 // Initialize matrices
 for(i=0; i<size; i++){
 for(j=0; j<size; j++){
 a[i][j] = (float)i + j;
 b[i][j] = (float)i - j;
 c[i][j] = 0.0;
 }
 }

 // Matrix multiplication
 #pragma acc kernels copyin(a,b) copy(c)
 for(i=0; i<size; i++){
 for(j=0; j<size; j++){
 for(k=0; k<size; k++){
 c[i][j] += a[i][k] * b[k][j];
 }
 }
 }

 return 0;
}

• OpenCL
• OpenMP

• http://www.openacc.org/
• Usage example from NVIDIA: part1, part2