Una herramienta de benchmarking para compiladores de OpenACC

Abstract

Resumen— OpenACC es un modelo de programación paralela para aceleradores de tipo GPU y Xeon PHI que lleva en desarrollo algunos años. Durante este tiempo han aparecido distintos compiladores, tanto comerciales como de código abierto, que se encuentran aún en un estado temprano de desarrollo. Dado que tanto el estándar como sus implementaciones son relativamente recientes, disponer de una suite de benchmarks diseñada para soportar varios compiladores puede facilitar enormemente el análisis comparativo de rendimiento de los distintos códigos generados. En este artículo presentamos TORMENT OpenACC, una suite de benchmarks preparada para comparar tanto arquitecturas como el código generado por diferentes compiladores. Junto a esta herramienta hemos desarrollado unas métricas adecuadas para la comparación del rendimiento de los pares compilador-máquina y que hemos denominado TORMENT ACC Score. Aquí se presenta la version 0.91 de la suite, desarrollada como prueba de concepto, y que incluye dos benchmarks. La version 1.0, considerada estable, incluirá seis benchmarks y estará disponible próximamente.

Full text

Una herramienta de benchmarking para
compiladores de OpenACC
Daniel Barba1, Arturo Gonz´alez-Escribano2y Diego R. Llanos3
Resumen— OpenACC es un modelo de programa-
ci´on paralela para aceleradores de tipo GPU y Xeon
PHI que lleva en desarrollo algunos a˜nos. Durante es-
te tiempo han aparecido distintos compiladores, tanto
comerciales como de c´odigo abierto, que se encuen-
tran a´un en un estado temprano de desarrollo. Da-
do que tanto el est´andar como sus implementaciones
son relativamente recientes, disponer de una suite de
benchmarks dise˜nada para soportar varios compilado-
res puede facilitar enormemente el an´alisis comparati-
vo de rendimiento de los distintos c´odigos generados.
En este art´ıculo presentamos TORMENT OpenACC,
una suite de benchmarks preparada para comparar
tanto arquitecturas como el c´odigo generado p or dife-
rentes compiladores. Junto a esta herramienta hemos
desarrollado unas m´etricas adecuadas para la com-
paraci´on del rendimiento de los pares compilador-
m´aquina y que hemos denominado TORMENT ACC
Score. Aqu´ı se presenta la version 0.91 de la suite,
desarrollada como prueba de concepto, y que incluye
dos benchmarks. La version 1.0, considerada estable,
incluir´a seis benchmarks y estar´a disponible pr´oxima-
mente.
Palabras clave— OpenACC, compiladores, bench-
marking.
I. Introducci´
on
OpenACC es un est´andar abierto que define una
serie de directivas de compilaci´on o pragmas para
ejecuci´on de fragmentos de c´odigo en aceleradores
de tipo GPU y Xeon Phi. Su objetivo es facilitar
la paralelizaci´on de c´odigo secuencial en este tipo
de aceleradores reduciendo el tiempo necesario tanto
en la programaci´on como en el aprendizaje [1]. La
especificaci´on se encuentra en su versi´on 2.5 [2].
La responsabilidad de la paralelizaci´on autom´atica
del c´odigo secuencial recae sobre los diferentes com-
piladores que soportan OpenACC: PGI Compiler [3],
de The Portland Group, empresa ahora pertenecien-
te a Nvidia, es (seg´un los estudios que hemos reali-
zado) el compilador con un grado de madurez m´as
alto. Se encuentra disponible como parte del Nvi-
dia OpenACC Toolkit, gratuito durante tres meses y
con la posibilidad de adquirir licencias comerciales o
de investigaci´on. OpenUH [4], de la Universidad de
Houston y accULL [5] de la Universidad de La Lagu-
na son dos alternativas de c´odigo abierto desarrolla-
das en el ´ambito acad´emico y que pueden descargarse
libremente para su uso.
Adem´as de los compiladores se˜nalados, existen
otros compiladores que soportan OpenACC, pero de-
bido a que en la fecha de redacci´on de este art´ıculo
1Dpto. de Inform´atica, Universidad de Valladolid, Espa˜na.
e-mail: daniel@infor.uva.es.
2Dpto. de Inform´atica, Universidad de Valladolid, Espa˜na.
e-mail: arturo@infor.uva.es.
3Dpto. de Inform´atica, Universidad de Valladolid, Espa˜na.
e-mail: diego@infor.uva.es.
fue imposible obtener licencias de prueba o acad´emi-
cas, han sido dejados fuera de nuestros estudios.
Estos compiladores est´an siendo desarrollados por
CRAY Inc. yPathscale Inc.
Nuestro trabajo ha consistido en el desarrollo de
una herramienta de an´alisis de rendimiento del c´odi-
go generado por estos compiladores, que hemos de-
nominado TORMENT OpenACC (Trasgo perfOR-
Mance and EvaluatioN Tool for OpenACC). El mo-
tivo que nos ha movido a comenzar este trabajo es la
escasez de herramientas similares para OpenACC y
las carencias que presentan las existentes. Estas he-
rramientas son, a fecha de redacci´on de este art´ıcu-
lo, EPCC OpenACC Benchmark Suite [6] y Rodi-
nia [7]. El primero ha sido desarrollado por el Edin-
burgh Parallel Computing Centre y consiste en una
serie de microbenchmarks destinados a medir el over-
head de la implementaci´on de los distintos pragmas,
y benchmarks para medir el rendimiento del c´odi-
go generado. Rodinia [7] es una traducci´on de los
benchmarks originales de la suite del mismo nombre
[8] para OpenACC desarrollado por Pathscale.
Nuestro trabajo busca dar respuesta a la necesi-
dad de an´alisis del rendimiento y comparaci´on del
c´odigo generado por los distintos compiladores, in-
tentando mantener un equilibrio entre las fortalezas
y debilidades de los diferentes compiladores y tra-
tando de obtener una herramienta que sea posible
utilizar con los compiladores disponibles para la co-
munidad acad´emica. Debido a que los compiladores
se encuentran a´un en fase de desarrollo, es imprescin-
dible mantener el c´odigo de los benchmarks lo m´as
sencillo posible.
Adem´as del desarrollo de la herramienta como tal,
hemos desarrollado tambi´en una m´etrica para po-
der ofrecer al usuario de esta herramienta una pun-
tuaci´on relativa que permita la comparaci´on entre
diferentes sistemas y compiladores analizados. Esta
m´etrica ha sido denominada TORMENT ACC Sco-
re.
El resto del art´ıculo se organiza del siguiente mo-
do: La secci´on II describe con m´as detenimiento las
herramientas de benchmarking existentes. La secci´on
III describe los compiladores actualmente soportados
por TORMENT OpenACC. En la secci´on IV se des-
cribe la herramienta, sus objetivos y caracter´ısticas a
medir y enumera los benchmarks actualmente imple-
mentados, as´ı como la m´etrica utilizada. Finalmente,
la secci´on V concluye el art´ıculo.
II. Herramientas existentes
Como se indicaba en la Introducci´on, a fecha de
redacci´on de este art´ıculo existen dos herramientas

de an´alisis de rendimiento para OpenACC. A conti-
nuaci´on describimos m´as en profundidad las mismas
y las carencias detectadas.
A. EPCC OpenACC Benchmark Suite
La suite de benchmarks desarrollada por el EPCC
ha sido dise˜nada espec´ıficamente para OpenACC. Se
compone de tres partes diferenciadas y que se deno-
minan Nivel 0,Nivel 1 yNivel de Aplicaciones.
En el Nivel 0 se pueden encontrar una serie de mi-
crobenchmarks cuyo objetivo es medir el overhead
generado por las implementaciones de los pragmas.
Estos microbenchmarks dan un resultado que es la
diferencia de dos tiempos de ejecuci´on en funci´on
del pragma que est´an analizando, o bien el tiempo
de transferencia de datos en una directiva de tipo
#pragma acc data. Estos resultados son interesan-
tes para el desarrollo de los compiladores, pero tal y
como est´an dise˜nados no ofrecen una idea clara de
rendimiento y los resultados pueden ser dif´ıciles de
interpretar.
En el Nivel 1 se encuentran un conjunto de bench-
marks t´ıpicos de tipo BLAS basados en Polybench y
Polybench/GPU [9]. Los resultados ofrecidos por es-
tos benchmarks son tiempos de ejecuci´on de los dife-
rentes c´odigos, por lo que son un buen indicador del
rendimiento del c´odigo generado por los diferentes
compiladores.
En el Nivel de Aplicaciones hay tres benchmarks
de mayor entidad que los anteriores. Estos bench-
marks tambi´en ofrecen tiempos de ejecuci´on, pero al
no ser problemas sint´eticos dan resultados m´as in-
teresantes.
En general, la idea del EPCC OpenACC Bench-
mark Suite est´a bien planteada y dise˜nada. No obs-
tante, los resultados obtenidos en los microbench-
marks no son adecuados para un an´alisis de rendi-
miento. Los dem´as benchmarks, si bien podr´ıan ser
´utiles, siempre han dado problemas en nuestros es-
tudios. En unos casos no pod´ıan compilarse con al-
gunos de los compiladores utilizados, en otros casos
daban error de ejecuci´on. Uno de los mayores pro-
blemas ha sido la limitaci´on en los tama˜nos de los
datos a utilizar debido a problemas en la implemen-
taci´on que llevaban a que se intentase asignar mucha
m´as memoria de la deseada, originando errores en la
ejecuci´on y limitando a 10MB el tama˜no de los datos.
B. Rodinia para OpenACC
La empresa Pathscale Inc. ha desarrollado una ver-
si´on [7] de la suite de benchmarks Rodinia [8], [10]
para su uso con compiladores de OpenACC. Con la
versi´on existente en GitHub a d´ıa 25 de Abril de 2014
(versi´on m´as reciente durante el desarrollo de nues-
tro trabajo), en general los benchmarks disponibles
no compilan con ninguno de los compiladores exis-
tentes, salvo contadas excepciones. El compilador de
PGI es el ´unico que logra compilar un n´umero signi-
ficativo de benchmarks.
La suite se compone de benchmarks que devuelven
el tiempo de ejecuci´on, por lo que ser´ıan un buen
punto de partida para un an´alisis de rendimiento.
Por desgracia, la escasa madurez del c´odigo y la po-
ca compatibilidad con los compiladores disponibles
hacen imposible su uso para establecer una compa-
rativa de rendimientos.
III. Compiladores Soportados
En la Introducci´on hemos enumerado brevemen-
te los compiladores disponibles. Su elecci´on est´a
motivada por la existencia de licencias gratuitas o
acad´emicas, o bien por ser de c´odigo abierto. A
continuaci´on explicaremos algunos detalles de los
compiladores soportados en la versi´on preliminar de
TORMENT OpenACC.
A. PGI Compiler
El compilador de PGI [3], desarrollado por The
Portland Group y Nvidia, es a fecha de redacci´on de
este art´ıculo el compilador con un grado de madu-
rez m´as alto. Su uso est´a muy extendido en los di-
ferentes talleres y conferencias que se realizan sobre
OpenACC. Actualmente est´a disponible como parte
del Nvidia OpenACC toolkit, que incluye una licen-
cia gratuita de tres meses y la posibilidad de adquirir
una licencia comercial o acad´emica. En nuestros es-
tudios, el compilador de PGI ha demostrado ser el
m´as s´olido y el que m´as alto grado de madurez tiene.
Se ajusta bien a la especificaci´on de OpenACC. Su
instalaci´on es simple y dispone de abundante docu-
mentaci´on.
B. OpenUH
El compilador OpenUH, desarrollado por la Uni-
versidad de Houston, es una alternativa de c´odigo
abierto. Comparado con el compilador de PGI, su
madurez y solidez son menores y carece de algunas
funcionalidades, como reducciones sobre m´ınimos o
m´aximos, lo cual dificulta ligeramente la programa-
ci´on de aplicaciones. Su instalaci´on, a fecha de redac-
ci´on de este art´ıculo, no es tan sencilla como cabr´ıa
esperar. La versi´on pre-compilada disponible en su
web [4] carece de algunas librer´ıas que deben ser com-
piladas aparte u obtenidas de otra forma.
C. accULL
El compilador accULL [5], desarrollado por la Uni-
versidad de La Laguna, es (al igual que OpenUH) un
compilador de c´odigo abierto. De los tres compilado-
res soportados, es el que tiene una menor robustez,
teniendo problemas para la traducci´on fuente a fuen-
te de algunas caracter´ısticas de C, como punteros a
funci´on, o de funcionalidades de OpenACC, como al-
gunos tipos de reducciones. La instalaci´on de la ver-
si´on disponible en su web [11] es bastante simple,
siguiendo las instrucciones que pueden encontrarse
adjuntas al compilador en ficheros de texto.
IV. TORMENT OpenACC
Nuestra propuesta en desarrollo se denomina
TORMENT OpenACC y es el acr´onimo para Trasgo
perfORMance and EvaluatioN Tool for OpenACC.

El proyecto ha nacido para intentar dar una solu-
ci´on a la necesidad de dar respuesta a los intentos
de realizar comparativas de rendimiento de c´odigo
OpenACC.
A. Motivaci´on
Las herramientas actuales han demostrado en
nuestros estudios previos no ser suficientes para dar
una respuesta sencilla a la evaluaci´on del c´odigo ge-
nerado por los diferentes compiladores de OpenACC.
El primero de los problemas es la escasa compatibi-
lidad del c´odigo de los benchmarks entre los diferen-
tes compiladores. Es cierto que si el c´odigo se ajusta
al est´andar deber´ıa compilarse y ejecutarse correc-
tamente en cualquier implementaci´on del est´andar
OpenACC, pero en la actual etapa en la que se en-
cuentran tanto el est´andar como los compiladores es-
to no se cumple. Debido a esto, hemos desarrollado
TORMENT OpenACC teniendo en cuenta tanto la
especificaci´on del est´andar como las capacidades ac-
tuales de los diferentes compiladores.
Otro problema detectado es que muchos de los
benchmarks no ofrecen una imagen clara de rendi-
miento, como es el caso de los microbenchmarks del
EPCC OpenACC Benchmark Suite. Si bien es cierto
que estos microbenchmarks son ´utiles para el desa-
rrollo de los compiladores, creemos que el overhead
de la implementaci´on de los distintos pragmas no
es realmente relevante en una comparativa de ren-
dimientos. Por este motivo, en nuestra propuesta de-
jamos fuera ese tipo de pruebas.
Finalmente, hemos observado tambi´en que, en mu-
chos casos, benchmarks que parecen compilar correc-
tamente luego generan errores en tiempo de ejecuci´on
o resultados incorrectos. Esta ´ultima situaci´on pue-
de suponer que, si no se detecta el resultado err´oneo,
los resultados de rendimiento obtenidos sean tambi´en
incorrectos. Para evitar esto, adem´as de incorporar
comprobaci´on de resultados a los benchmarks, tam-
bi´en hemos analizado los resultados durante el desa-
rrollo de nuestra propuesta, para que los benchmarks
que incorporamos en nuestra herramienta compilen
y ejecuten correctamente con cualquiera de los com-
piladores.
Otro aspecto importante que hemos tenido en
cuenta a la hora de llevar a cabo el desarrollo de
TORMENT OpenACC es la ofrecer a la comunidad
una herramienta que, adem´as de facilitar el an´alisis
de rendimiento del c´odigo generado por los compi-
ladores de OpenACC ejecutados en distintas m´aqui-
nas, ofrezcan una medida del rendimiento que per-
mita una comparaci´on f´acil entre m´aquinas y compi-
ladores. Esta idea nos ha llevado al desarrollo de las
m´etricas TORMENT ACC Score que describiremos
posteriormente.
B. Objetivos
El objetivo principal de TORMENT OpenACC es
la de permitir un an´alisis de rendimiento de c´odigo
OpenACC generado por los distintos compiladores
de forma sencilla, generando un resumen de resulta-
dos f´acilmente analizable y ofreciendo unas m´etricas,
denominadas TORMENT ACC Score, que permiten
la comparaci´on de los pares m´aquina-compilador.
Nuestra propuesta pretende facilitar a la comuni-
dad una herramienta preparada espec´ıficamente para
OpenACC que sea consciente del estado de desarro-
llo temprano de los compiladores existentes, de forma
que se eviten problemas en compilaci´on o ejecuci´on
como sucede con otras herramientas de benchmar-
king existentes. TORMENT OpenACC estar´a pre-
parado para compilarse y ejecutarse con la menor in-
tervenci´on posible del usuario. Los scripts que acom-
pa˜nan a la suite recopilan la informaci´on del proceso
y ofrecen finalmente un informe en formato HTML
con los datos relevantes.
Adem´as, TORMENT OpenACC utiliza los compi-
ladores GCC y NVCC para obtener datos de ejecu-
ciones de c´odigo secuencial y c´odigo CUDA respec-
tivamente. De este modo, nuestra propuesta ofrece
al usuario informaci´on del speedup con respecto al
c´odigo secuencial y CUDA ejecutado en la misma
m´aquina.
C. Estructura de la herramienta
TORMENT OpenACC se compone de una serie de
scripts que se encargan de todo el proceso de com-
pilaci´on, ejecuci´on y obtenci´on de resultados, elimi-
nando esta carga al usuario. Estos scripts se dividen
en tres categor´ıas. Los scripts de configuraci´on gu´ıan
al usuario en la obtenci´on de las rutas correctas a
librer´ıas CUDA (necesarias por algunos de los com-
piladores), rutas de compiladores y comandos de eje-
cuci´on (en caso de que el usuario utilice, por ejemplo,
sistemas de colas tipo slurm). El script de ejecuci´on
se encarga de la compilaci´on y ejecuci´on usando to-
dos los compiladores que hayan sido hallados en el
sistema del usuario. Finalmente, el script de genera-
ci´on de resultados procesa los resultados obtenidos y
genera un fichero HTML con el informe final.
Cada benchmark est´a desarrollado de modo que
sea lo m´as sencillo posible, para evitar problemas
de compilaci´on. Por este motivo se procura evitar
el uso de niveles de indirecci´on que en otras situacio-
nes ser´ıan aconsejables. Cada benchmark est´a defi-
nido en su propia unidad de compilaci´on y el c´odigo
OpenACC y CUDA est´a incorporado en el mismo
fuente, utilizando compilaci´on condicional y evitan-
do tener de forma simultanea ficheros .c y.cu con
c´odigo duplicado. Los ficheros .cu son necesarios pa-
ra que el compilador NVCC genere el c´odigo CUDA
correspondiente, pero esto se ha resuelto mediante el
uso de enlaces simb´olicos a los ficheros .c correspon-
dientes.
El programa principal se encarga del lanzamiento
de los benchmarks. Cada uno se lanza con diez repe-
ticiones, con una repetici´on extra al comienzo cuyos
resultados son desechados. Al finalizar estas diez re-
peticiones se obtienen los valores denominados peak
yaverage y que corresponden a la mejor de las ejecu-
ciones y a la media aritm´etica de todas ellas. Estos
valores servir´an para calcular la m´etrica de la que

CUDA_ATTR__ int getRandom(unsigned int* seed)
{
unsigned int next = *seed;
int result;
next *= 1103515245;
next += 12345;
result = (unsigned int) (next/65536) % 2048;
next *= 1103515245;
next += 12345;
result <<= 10;
result ^= (unsigned int) (next/65536) % 1024;
next *= 1103515245;
next += 12345;
result <<= 10;
result ^= (unsigned int) (next/65536) % 1024;
*seed = next;
return result;
}
Fig. 1
C´
odigo para la generaci´
on de n´
umeros aleatorios.
hablaremos posteriormente.
D. Benchmarks implementados
La versi´on preliminar de TORMENT OpenACC
contiene ´unicamente dos benchmarks. El desarrollo
de la herramienta sigue en proceso y otros bench-
marks se ir´an a˜nadiendo progresivamente.
D.1 MonteCarloPi
Este benchmark consiste en una aproximaci´on de
Pi por el m´etodo de Monte Carlo, que se basa en
la generaci´on de puntos aleatorios en un cuadrado
de lado unitario. Se comprueba si estos puntos se
encuentran dentro de un cuarto de c´ırculo de radio
unitario y se acumula el total de puntos que cumplen
dicha condici´on. Finalmente, se aplica la siguiente
f´ormula:
π≈4∗P
T
Donde Pes el n´umero de puntos dentro del cuarto
de c´ırculo y Tes el total de puntos generados.
MonteCarloPi es un benchmark que no tiene
pr´acticamente transferencias de memoria y que reali-
za un c´alculo computacional muy simple, pero puede
ser optimizado en CUDA haciendo que cada hilo cal-
cule varios puntos y utilizando la shared memory de
los bloques para evitar accesos a memoria global. Un
buen resultado de los compiladores en este bench-
mark depender´a de estos factores.
Dado que no se puede hacer uso de la funci´on srand
de C en el c´odigo ejecutado en la GPU, y el uso de
la librer´ıa curand se limita a c´odigo CUDA, hemos
decidido replicar la funci´on srand para permitir su
ejecuci´on en la GPU, como se indica en la Fig. 1.
El c´odigo utilizado para las versiones de OpenACC
y CUDA se muestra en las Figs. 2 y 3.
#pragma acc parallel loop \
private(i, d, x, y, seed) \
reduction(+:count)
for(i = 0; i < COORD_NUM; ++i){
seed = 1987 ^ i*27;
x = (double)getRandom(&seed)/(double)RAND_MAX;
y = (double)getRandom(&seed)/(double)RAND_MAX;
d = sqrt(x*x + y*y);
if(d <= 1.0){
++count;
}
}
Fig. 2
C´
odigo de MonteCarloPi para OpenACC.
__CUDA_GLOBAL__ void piKernel(
const unsigned long int triesPerThread,
unsigned long int* hits)
{
unsigned int seed;
int gid, tid, bid;
int lhits;
float x, y;
extern __shared__ unsigned long int sdata[];
gid = (blockIdx.x*blockDim.x) + threadIdx.x;
tid = threadIdx.x;
bid = blockIdx.x;
lhits = 0;
seed = 1987 ^ gid*27;
for (int i = 0; i < triesPerThread; ++i){
x = (float)getRandom(&seed)/(float)RAND_MAX;
y = (float)getRandom(&seed)/(float)RAND_MAX;
float d = sqrt(x*x + y*y);
if (d <= 1.0f){
++lhits;
}
}
sdata[tid] = lhits;
__syncthreads();
if (tid == 0){
for (int i = 1; i < blockDim.x; ++i){
lhits += sdata[i];
}
hits[bid] = lhits;
}
}
Fig. 3
C´
odigo de MonteCarloPi para CUDA.

#pragma acc data copyin(S[0:sizeS],B[0:sizeB]) \
copy(C[0:sizeB])
{
#pragma acc parallel loop private(startB) \
firstprivate(result,sizeS,sizeB)
for (startB = 0; startB <= sizeB-sizeS; startB++){
C[startB] = 0;
if (startB <= result) {
int ind;
for(ind = 0; ind < sizeS; ind++){
if (S[ind] != B[startB+ind]) break;
}
if (ind == sizeS){
result = startB;
C[startB] = 1;
}
}
}
}
result = -1;
for (startB = 0; startB <= sizeB-sizeS; startB++){
if (C[startB] == 1){
result = startB;
break;
}
}
Fig. 4
C´
odigo de StringMatch para OpenACC.
D.2 StringMatch
El benchmark StringMatch es un programa de ali-
neamiento de cadenas de caracteres. El programa
consiste en la b´usqueda de la primera ocurrencia de
una cadena peque˜na en una cadena grande, utilizan-
do un algoritmo na¨ıve. En este benchmark, la cadena
grande tiene una longitud de 10 millones de carac-
teres, es decir, aproximadamente 10MB. Las cade-
nas peque˜nas, que son cuatro, tienen una longitud
de 1000 caracteres.
Este algoritmo es interesante porque combina
transferencia de datos con la necesidad de un uso
eficiente de la memoria, especialmente de la shared
memory. Los c´odigos de las versiones de OpenACC
y CUDA pueden verse en las Figs. 4 y 5.
E. M´etrica utilizada
Para la elecci´on de la m´etrica denominada
TORMENT ACC Score hemos decidido optar por la
metodolog´ıa utilizada por SPEC [12]. El System Per-
formance Evaluation Council, com´unmente conocido
como SPEC, es un referente ampliamente conocido
en cuanto a benchmarking y an´alisis de rendimiento
se refiere. Una de sus fortalezas es el reconocer que los
benchmarks envejecen en funci´on del paso del tiempo
y, en consecuencia, deben ser actualizados.
SPEC utiliza la siguiente metodolog´ıa. En primer
lugar, cada programa devuelve su tiempo de ejecu-
ci´on y se calcula el SPECratio, que consiste en el
ratio obtenido de dividir un tiempo de ejecuci´on de
referencia suministrado por SPEC entre el tiempo de
ejecuci´on obtenido. Finalmente, se obtiene la media
geom´etrica de todos los SPECratios del conjunto de
__global__ void kernel_busqueda(const char* b_idata,
const int sizeB, const char* s_idata,
const int sizeS, unsigned int* result)
{
extern __shared__ unsigned char sdata[];
int gid = (blockIdx.x*blockDim.x) + threadIdx.x;
int tid = threadIdx.x;
int sStart = THREADS_PER_BLOCK + sizeS;
int offset;
if (gid == 0){
*result = -1;
}
__syncthreads();
if (gid < *result){
unsigned char* bd = &sdata[tid];
b_idata += gid;
unsigned char* sd = &sdata[sStart+tid];
s_idata += tid;
if(tid<THREADS_PER_BLOCK){
for(offset = 0; tid+offset<sStart;
offset += THREADS_PER_BLOCK)
{
if (offset==0){
*bd = *b_idata;
bd+=THREADS_PER_BLOCK;
b_idata+=THREADS_PER_BLOCK;
continue;
}
*bd = *b_idata;
bd+=THREADS_PER_BLOCK;
b_idata+=THREADS_PER_BLOCK;
*sd = *s_idata;
sd+=THREADS_PER_BLOCK;
s_idata+=THREADS_PER_BLOCK;
}
}
__syncthreads();
if (gid <= sizeB-sizeS){
unsigned int i;
unsigned int b = 1;
unsigned char* sd = &sdata[sStart];
unsigned char* bd = &sdata[tid];
for(i = 0; b && i+15 < sizeS; i+=16){
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
b &= (*sd == *bd); sd++; bd++;
if (!b) break;
}
for(; b && i < sizeS ; i++){
b &= (*sd == *bd); sd++; bd++;
if (!b) break;
}
if ( b ) atomicMin(result, gid);
}
}
}
Fig. 5
C´
odigo de StringMatch para CUDA.

benchmarks [13].
Nuestra propuesta sigue una idea similar a la de
SPEC, pero con algunas variaciones. En primer lu-
gar, la ejecuci´on de los benchmarks que componen
TORMENT OpenACC devuelven tres valores. Peak
Time es el mejor tiempo de ejecuci´on obtenido, me-
dido en segundos. Average Time es el tiempo medio
de todas las ejecuciones del benchmark, tambi´en en
segundos. Finalmente, Standard Deviation es la des-
viaci´on est´andar del conjunto de medidas e indica la
variabilidad obtenida en las ejecuciones del bench-
mark. Con los tiempos peak yaverage se calcula un
ratio con respecto a los tiempos de referencia sumi-
nistrados con la herramienta, y que son los tiempos
de ejecuci´on secuencial del benchmark en una m´aqui-
na de referencia. Una vez ejecutados todos los bench-
marks, se obtiene la media arm´onica de todos los ra-
tios, tanto para peak time como average time. Estos
valores que se obtienen son TORMENT ACC Peak
Score yTORMENT ACC Average Score.
La principal diferencia entre TORMENT
OpenACC y SPEC, aparte de la metodolog´ıa
de toma de tiempos de ejecuci´on, consiste en el
uso de la media arm´onica en lugar de la media
geom´etrica. Esta decisi´on se fundamenta en el
hecho de que, para los objetivos de TORMENT
OpenACC, la media arm´onica es m´as adecuada
que la media geom´etrica. En primer lugar, aunque
la media geom´etrica siempre produce una orde-
naci´on consistente, no necesariamente produce la
ordenaci´on correcta [13], ya que esta media no es
inversamente proporcional al tiempo de ejecuci´on.
En cambio, la media arm´onica si que es inversamente
proporcional al tiempo de ejecuci´on, lo que hace que
sea una media correcta para expresar ratios. Estas
afirmaciones son compartidas por otros trabajos,
como por ejemplo [14].
Como ejemplo de uso, las figuras 6 y 7 muestran
dos informes generados para sistemas con diferente
configuraci´on. En el primero puede verse que el com-
pilador de PGI obtiene el TORMENT ACC Score
m´as alto, tanto Peak como Average, con un valor
de 28.56 y 25.71 respectivamente. Los resultados con
respecto a la implementaci´on en CUDA de los bench-
marks, muestran que los tres compiladores generan
c´odigo cuyo rendimiento queda muy lejos del conse-
guido usando CUDA. En el segundo puede observarse
que a pesar de utilizar una GPU mucho m´as modesta
en comparaci´on, los resultados son consistentes con
los del primer informe. En ambos resultados se apre-
cia como el compilador de PGI obtiene los valores
m´as grandes de la desviaci´on t´ıpica. Esto parece in-
dicar que al generar el c´odigo se est´a modificando
la sem´antica del programa y la primera ejecuci´on de
cada benchmark que deber´ıa ser desechada se est´a
incluyendo en las medidas.
V. Conclusiones y Trabajo Futuro
TORMENT OpenACC es una herramienta de
an´alisis y comparaci´on de rendimientos de c´odigo
generado por compiladores de OpenACC, teniendo
en consideraci´on el nivel de madurez de tanto el
est´andar OpenACC como de los diferentes compi-
ladores. TORMENT OpenACC desarrolla una sui-
te de benchmarks espec´ıficamente dise˜nados para
OpenACC y manteniendo la m´axima portabilidad
entre compiladores, permitiendo su compilaci´on y
ejecuci´on en todos ellos.
Los resultados ofrecidos por TORMENT
OpenACC incluyen la denominada
TORMENT ACC Score, dise˜nada para la com-
paraci´on de pares m´aquina-compilador. Adem´as, se
ofrece un resumen de la ejecuci´on de los benchmarks
con una tabla de tiempos y ratios, tanto m´ınimos
como medios e incluyendo la desviaci´on est´andar
para un an´alisis de variabilidad de los tiempos de
ejecuci´on del c´odigo generado.
Junto con los resultados de los compiladores de
OpenACC se incluyen tambi´en resultados de ejecu-
ci´on de c´odigo generado por los compiladores GCC
y NVCC para c´odigo secuencial y CUDA. De este
modo, se puede ofrecer al usuario una comparativa a
nivel de m´aquina del rendimiento del c´odigo genera-
do por los compiladores de OpenACC con respecto
a versiones secuenciales y CUDA de los benchmarks.
El trabajo futuro a desarrollar consiste en la am-
pliaci´on de la suite de benchmarks, para lograr re-
sultados que sean verdaderamente relevantes. Aqu´ı
se presenta la versi´on 0.91 de la suite, desarrollada
como prueba de concepto, y que incorpora dos bench-
marks. La versi´on 1.0, considerada estable, incluir´a
seis benchmarks y estar´a disponible pr´oximamente,
Tratando de cubrir los aspectos m´as interesantes de
la ejecuci´on de c´odigo en las GPUs. Adem´as, una par-
te importante del trabajo restante consiste en mante-
ner la compatibilidad entre compiladores y asegurar
el correcto funcionamiento de la herramienta en dis-
tintas m´aquinas.
Agradecimientos
En memoria de Agust´ın de Dios Hern´andez.
Esta investigaci´on ha sido parcialmente financiada
por el MICINN y el programa ERDF de la Uni´on Eu-
ropea: proyecto HomProg-HetSys (TIN2014-58876-
P), la red CAPAP-H5 (TIN2014-53522-REDT) y el
COST Program Action IC1305: Network for Sustai-
nable Ultrascale Computing (NESUS).
Referencias
[1] OpenACC-standard.org, “About OpenACC,” .
[2] OpenACC-standard.org, “OpenACC 2.5 draft for public
comment,” oct 2015.
[3] PGI, “Pgi accelerator compilers with OpenACC direc-
tives,” https://www.pgroup.com/resources/accel.htm,
nov 2015.
[4] University of Houston, “Open-source UH compiler,”
http://web.cs.uh.edu/~openuh/download/, nov 2015.
[5] Ruym´an Reyes, Iv´an L´opez-Rodr´ıguez, Juan J Fumero,
and Francisco de Sande, “accULL: an OpenACC im-
plementation with CUDA and OpenCL support,” in
Euro-Par 2012 Parallel Processing, pp. 871–882. Sprin-
ger, 2012.
[6] EPCC, “Epcc OpenACC benchmark suite,” https:
//github.com/EPCCed/epcc-openacc- benchmarks, sep
2013.
[7] Pathscale, “Rodinia benchmark suite 2.1 with OpenACC

port,” https://github.com/pathscale/rodinia, apr
2014.
[8] Shuai Che, Michael Boyer, Jiayuan Meng, David Tarjan,
Jeremy W Sheaffer, Sang-Ha Lee, and Kevin Skadron,
“Rodinia: A benchmark suite for heterogeneous compu-
ting,” in Workload Characterization, 2009. (IISWC),
2009 IEEE International Symposium on. IEEE, 2009,
pp. 44–54.
[9] Louis-No¨el Pouchet, “Polybench: The polyhedral bench-
mark suite,” URL: http://www. cs. ucla. edu/˜ pou-
chet/software/polybench/[cited July,], 2012.
[10] Shuai Che, Jeremy W Sheaffer, Michael Boyer, Lukasz G
Szafaryn, Liang Wang, and Kevin Skadron, “A characte-
rization of the Rodinia benchmark suite with comparison
to contemporary CMP workloads,” in Workload Cha-
racterization (IISWC), 2010 IEEE International Sym-
posium on. IEEE, 2010, pp. 1–11.
[11] Universidad de La Laguna, “accULL,” http://cap.pcg.
ull.es/es/accULL, nov 2015.
[12] Kaivalya M Dixit, “The spec benchmarks,” Parallel com-
puting, vol. 17, no. 10, pp. 1195–1209, 1991.
[13] David J Lilja, Measuring computer performance: a prac-
titioner’s guide, Cambridge University Press, 2005.
[14] John R Mashey, “War of the benchmark means: time for
a truce,” ACM SIGARCH Computer Architecture News,
vol. 32, no. 4, pp. 1–14, 2004.

Benchmarkversion:0.91
Hostname:hydra
Username:daniel
Softwaresta ckinfo
Kernel:Linux3.10.0229.4.2.el7.x86_64x86_64
GCC:gcc(GCC)4.8.320140911(RedHat4.8.39)
NVCC:Cudacompilationtools,release7.5,V7.5.17
PGICompiler:pgcc15.7064bittargetonx8664Linuxtphaswell
OpenUHCompiler:OpenUH3.1.0(bas edonOpen64CompilerSuite:Version5.0)
accULLCompiler:Release0.4alpha
CPUinfo
Model:Intel(R)Xeon(R)CPU E52609v3@1.90GHz
Architecture:x86_64
NumberofCores:6
MaxMHz:MHz
MinMHz:MHz
L1Cache:32K
L2Cache:256K
L3Cache:15360K
RAM:65687144kB
GPU(s)info
GPU0:NVIDIACorporationGK110B[GeForceGTX Tit anBlack]
(reva1)
GPU1:NVIDIACorporationGK110B[GeForceGTX Tit anBlack]
(reva1)
GPU2:NVIDIACorporationGK110B[GeForceGTX Tit anBlack]
(reva1)
GPU3:NVIDIACorporationGK110B[GeForceGTX Tit anBlack]
(reva1)
TORMENT_ACC_0.91PeakScore:28.26 TORMENT_ACC_0.91AverageScore:25.71
SpeedupvsSequential(pea k):57.92x SpeedupvsSequential(a verage):52.40x
SpeedupvsCUDA(peak):0.17x SpeedupvsCUDA(average ):0.15x
TORMENT_ACC_0.91PeakScore:12.37 TORMENT_ACC_0.91AverageScore:12.30
SpeedupvsSequential(pea k):25.35x SpeedupvsSequential(a verage):25.07x
SpeedupvsCUDA(peak):0.07x SpeedupvsCUDA(average ):0.07x
TORMENT_ACC_0.91PeakScore:16.14 TORMENT_ACC_0.91AverageScore:16.07
SpeedupvsSequential(pea k):33.08x SpeedupvsSequential(a verage):32.76x
SpeedupvsCUDA(peak):0.10x SpeedupvsCUDA(average ):0.10x
SystemInformation
GCC_SequentialCompilerResults:
GCC_Sequential
Benchmark PeakTime (s) PeakRatio Avg.Time (s) Avg.Ratio Std.Deviation
MonteCarloPi 14.763701 0.43 14.764007 0.43 0.000336
StringMatch 20.118267 0.57 20.136980 0.57 0.021079
NVCC_CUDACompilerResults:
NVCC_CUDA
Benchmark PeakTime (s) PeakRatio Avg.Time (s) Avg.Ratio Std.Deviation
MonteCarloPi 0.016409 383.96 0.017840 354.36 0.001254
StringMatch 0.105644 108.48 0.105889 109.21 0.000162
PGICompilerResults:
PGI
Benchmark PeakTime (s) PeakRatio Avg.Time (s) Avg.Ratio Std.Deviation
MonteCarloPi 0.067226 93.72 0.077228 81.86 0.008222
StringMatch 0.688718 16.64 0.758494 15.25 0.068944
OpenUHCompilerResults:
OpenUH
Benchmark PeakTime (s) PeakRatio Avg.Time (s) Avg.Ratio Std.Deviation
MonteCarloPi 0.153154 41.14 0.165321 38.24 0.006447
StringMatch 1.574083 7.28 1.578459 7.33 0.002537
accULLCompilerResults:
accULL
Benchmark PeakTime (s) PeakRatio Avg.Time (s) Avg.Ratio Std.Deviation
MonteCarloPi 0.135701 46.43 0.139198 45.42 0.007257
StringMatch 1.173044 9.77 1.184555 9.76 0.004565
Peaksc oreusesthebestof10results.Averagescoreusesthearithmeticmean of10results.Inbothcaseshigherisbetter.
TORMENTOpenACCResultpage
GrupoTrasgoUniversidaddeValladolid
Fig. 6
P´
agina de Resultados TORMENT OpenACC

Benchmarkversion:0.91
Hostname:corikLaptop
Username:dani
Softwaresta ckinfo
Kernel:Linux3.16.04amd64x 86_64
GCC:gcc(Debian4.9.210) 4.9.2
NVCC:Cudacompilationtools,release7.5,V7.5.17
PGICompiler:pgcc15.7064bittargetonx8664Linuxtphaswell
OpenUHCompiler:OpenUH3.1.0(bas edonOpen64CompilerSuite:Version5.0)
accULLCompiler:Release0.4alpha
CPUinfo
Model:Intel(R)Core(TM)i54200HCPU@2.80GHz
Architecture:x86_64
NumberofCores:4
MaxMHz:3400.0000MHz
MinMHz:800.0000MHz
L1Cache:32K
L2Cache:256K
L3Cache:3072K
RAM:3939364kB
GPU(s)info
GPU0:NVIDIACorporationGM107M[GeForceGTX850M](rev
a2)
TORMENT_ACC_0.91PeakScore:8.98 TORMENT_ACC_0.91Ave rageScore:8.22
SpeedupvsSequential(pea k):8.84x Speedupvs Sequential (average):8.05x
SpeedupvsCUDA(peak):0.17x SpeedupvsCUDA(average ):0.16x
TORMENT_ACC_0.91PeakScore:6.56 TORMENT_ACC_0.91Ave rageScore:6.60
SpeedupvsSequential(pea k):6.46x Speedupvs Sequential (average):6.46x
SpeedupvsCUDA(peak):0.13x SpeedupvsCUDA(average ):0.13x
TORMENT_ACC_0.91PeakScore:6.50 TORMENT_ACC_0.91Ave rageScore:6.54
SpeedupvsSequential(pea k):6.40x Speedupvs Sequential (average):6.40x
SpeedupvsCUDA(peak):0.12x SpeedupvsCUDA(average ):0.13x
SystemInformation
GCC_SequentialCompilerResults:
GCC_Sequential
Benchmark PeakTime (s) PeakRatio Avg.Time (s) Avg.Ratio Std.Deviation
MonteCarloPi 6.196389 1.02 6.196760 1.02 0.000310
StringMatch 11.292974 1.01 11.298524 1.02 0.006537
NVCC_CUDACompilerResults:
NVCC_CUDA
Benchmark PeakTime (s) PeakRatio Avg.Time (s) Avg.Ratio Std.Deviation
MonteCarloPi 0.057900 108.81 0.062234 101.58 0.004161
StringMatch 0.335251 34.19 0.335390 34.48 0.000085
PGICompilerResults:
PGI
Benchmark PeakTime (s) PeakRatio Avg.Time (s) Avg.Ratio Std.Deviation
MonteCarloPi 0.343979 18.32 0.378440 16.70 0.034409
StringMatch 1.925690 5.95 2.120944 5.45 0.192720
OpenUHCompilerResults:
OpenUH
Benchmark PeakTime (s) PeakRatio Avg.Time (s) Avg.Ratio Std.Deviation
MonteCarloPi 0.697705 9.03 0.698031 9.06 0.000198
StringMatch 2.224097 5.15 2.226121 5.19 0.003758
accULLCompilerResults:
accULL
Benchmark PeakTime (s) PeakRatio Avg.Time (s) Avg.Ratio Std.Deviation
MonteCarloPi 0.675184 9.33 0.676132 9.35 0.000704
StringMatch 2.298089 4.99 2.299703 5.03 0.003703
Peaksc oreusesthebestof10results.Averagescoreusesthearithmeticmean of10results.Inbothcaseshigherisbetter.
TORMENTOpenACCResultpage
GrupoTrasgoUniversidaddeValladolid
Fig. 7
P´
agina de Resultados TORMENT OpenACC