Una he amien a de benchma king pa a
compilado es de OpenACC
Daniel Ba ba,1A u o Gonz´alez-Esc ibano2y Diego R. Llanos3
Resumen— OpenACC exis e ya desde hace algunos
a˜nos y, du an e los mismos, han ido apa eciendo una
se ie de compilado es an o en el ´ambi o acad´emico
como en la indus ia. Debido a la no edad del
es ´anda OpenACC as´ı como al con inuo desa ollo
de los compilado es exis en es, una sui e de bench-
ma ks espec´ı icamen e c eada pa a analiza el com-
po amien o del c´odigo gene ado po es os compi-
lado es en dis in as m´aquinas adquie e una u ili-
dad impo an e. En es e a ´ıculo p esen amos TOR-
MENT OpenACC, una sui e de benchma ks p epa ada
pa a se compilada po di e en es compilado es y que
o ece un esumen de los esul ados ob enidos. As´ı
mismo, jun o a es a he amien a hemos desa ollado
una m´e ica adecuada pa a la pun uaci´on de los pa es
compilado -m´aquina y que hemos denominado Pun-
uaci´on TORMENT ACC.
Palab as cla e— OpenACC, compilado es, bench-
ma king.
I. In oducci´
on
OpenACC es un es ´anda abie o que de ine una
se ie de di ec i as de compilaci´on o p agmas pa a
ejecuci´on de agmen os de c´odigo en acele ado es
de ipo GPU y Xeon Phi. Su obje i o es acili a
la pa alelizaci´on de c´odigo secuencial en es e ipo
de acele ado es educiendo el iempo necesa io an o
en la p og amaci´on como en el ap endizaje [1]. La
especi icaci´on se encuen a en la e si´on 2.5 [2].
La esponsabilidad de la pa alelizaci´on au om´a ica
del c´odigo secuencial ecae sob e los di e en es com-
pilado es que sopo an OpenACC: PGI Compile [3],
de The Po land G oup, emp esa aho a pe enecien e
a N idia, es seg´un los es udios que hemos ealizado
el compilado con un g ado de madu ez m´as al o.
Se encuen a disponible como pa e del N idia Ope-
nACC Toolki , g a ui o du an e es meses y con la
posibilidad de adqui i licencias come ciales o de in-
es igaci´on. OpenUH [4], de la Uni e sidad de Hous-
on y accULL [5] de la Uni e sidad de La Laguna son
dos al e na i as de c´odigo abie o desa olladas en el
´ambi o acad´emico y que pueden desca ga se lib e-
men e pa a su uso.
Adem´as de los compilado es se˜nalados exis en
o os compilado es que sopo an OpenACC, pe o
debido a que en la echa de edacci´on de es e
a ´ıculo ue imposible ob ene licencias de p ueba o
acad´emicas, han sido dejados ue a de nues os es u-
dios. Es os compilado es es ´an siendo desa ollados
po CRAY Inc. yPa hscale Inc..
1Dp o. de In o m´a ica, Uni e sidad de Valladolid, e-mail:
[email p o ec ed].
2Dp o. de In o m´a ica, Uni e sidad de Valladolid, e-mail:
[email p o ec ed].
3Dp o. de In o m´a ica, Uni e sidad de Valladolid, e-mail:
[email p o ec ed].
Nues o abajo consis e en el desa ollo de una
he amien a de an´alisis de endimien o del c´odigo
gene ado po es os compilado es que hemos denomi-
nado TORMENT OpenACC (T asgo pe ORMance
and E alua ioN Tool o OpenACC). El mo i o que
nos ha mo ido a comenza es e abajo es la es-
casez de he amien as simila es pa a OpenACC y las
ca encias de las exis en es. Es as he amien as son,
a echa de edacci´on de es e a ´ıculo, EPCC Ope-
nACC Benchma k Sui e [6] y Rodinia [7]. El p ime o
ha sido desa ollado po el Edinbu gh Pa allel Com-
pu ing Cen e y consis e en una se ie de mic obench-
ma ks, des inados a medi el o e head de la imple-
men aci´on de los dis in os p agmas, y benchma ks
pa a medi el endimien o del c´odigo gene ado. Ro-
dinia es una aducci´on de los benchma ks o iginales
de la sui e de Rodinia [8] pa a OpenACC desa ol-
lado po Pa hscale.
Nues o abajo busca da espues a a la necesi-
dad de an´alisis del endimien o y compa aci´on del
c´odigo gene ado po los dis in os compilado es, in-
en ando man ene un equilib io en e las o alezas y
debilidades de los di e en es compilado es y a ando
de ob ene una he amien a que sea posible u iliza
con los compilado es disponibles pa a la comunidad
acad´emica. Debido a que los compilado es se en-
cuen an a´un en ase de desa ollo, es imp escindible
man ene el c´odigo lo m´as sencillo posible.
Adem´as del desa ollo de la he amien a como
al, hemos desa ollado ambi´en una m´e ica pa a
pode o ece al usua io de es a he amien a una
pun uaci´on ela i a que pe mi a la compa aci´on en-
e di e en es sis emas y compilado es analizados.
Es a m´e ica ha sido denominada Pun uaci´on TOR-
MENT ACC
El es o del a ´ıculo se o ganiza del siguien e modo:
La secci´on dos desc ibe con m´as de enimien o las he -
amien as de benchma king exis en es. La secci´on
es desc ibe los compilado es ac ualmen e sopo a-
dos po TORMENT OpenACC. En la secci´on cua-
o se desc ibe la he amien a, sus obje i os y ca ac-
e ´ıs icas a medi y enume a los benchma ks ac ual-
men e implemen ados, as´ı como la m´e ica u ilizada.
Finalmen e, la secci´on cinco concluye el a ´ıculo.
II. He amien as exis en es
Como se indicaba en la in oducci´on, a echa de
edacci´on de es e a ´ıculo exis en dos he amien as
de an´alisis de endimien o pa a OpenACC. A con in-
uaci´on desc ibimos m´as en p o undidad las mismas
y las ca encias de ec adas.
A. EPCC OpenACC Benchma k Sui e
La sui e de benchma ks desa ollada po el EPCC
ha sido dise˜nada espec´ı icamen e pa a OpenACC. Se
compone de es pa es di e enciadas y que denomina
Ni el 0,Ni el 1 yNi el de Aplicaciones.
En el Ni el 0 se pueden encon a una se ie de mi-
c obenchma ks cuyo obje i o es medi el o e head
gene ado po las implemen aciones de los p agmas.
Es os mic obenchma ks dan un esul ado que es la
di e encia de dos iempos de ejecuci´on en unci´on
del p agma que es ´an analizando, o bien el iempo
de ans e encia de da os en una di ec i a de ipo
#p agma acc da a. Es os esul ados son in e esan es
pa a el desa ollo de los compilado es, pe o al y
como es ´an dise˜nados no o ecen una idea cla a de
endimien o y los esul ados pueden se con usos.
En el Ni el 1 se encuen an un conjun o de bench-
ma ks ´ıpicos de ipo BLAS basados en Polybench y
Polybench/GPU [9]. Los esul ados o ecidos po es-
os benchma ks son iempos de ejecuci´on de los di e -
en es c´odigos po lo que son un buen indicado del
endimien o del c´odigo gene ado po los di e en es
compilado es.
En el Ni el de Aplicaciones hay 3 benchma ks de
mayo en idad que los an e io es. Es os benchma ks
ambi´en o ecen iempos de ejecuci´on, pe o al se
p oblemas menos sin ´e icos dan esul ados m´as in-
e esan es.
En gene al, la idea del EPCC OpenACC Bench-
ma k Sui e es ´a bien plan eada y dise˜nada. No ob-
s an e, los esul ados ob enidos en los mic obench-
ma ks no son adecuados pa a un an´alisis de
endimien o. Los dem´as benchma ks, si bien pod ´ıan
se ´u iles, siemp e han dado p oblemas en nues os
es udios. En unos casos no pod´ıan compila se con al-
gunos de los compilado es u ilizados, en o os casos
daban e o de ejecuci´on. Uno de los mayo es p oble-
mas ha sido la limi aci´on en los ama˜nos de los da os
a u iliza debido a p oblemas en la implemen aci´on
que lle aban a que se in en ase asigna mucha m´as
memo ia de la deseada, o iginando e o es en la eje-
cuci´on y limi ando a 10MB el ama˜no de los da os.
B. Rodinia pa a OpenACC
La emp esa Pa hscale Inc. ha desa ollado una
e si´on [7] de la sui e de benchma ks Rodinia [8],
[10] pa a su uso con compilado es de OpenACC. Con
la e si´on exis en e en Gi Hub a d´ıa 25 de Ab il
de 2014, en gene al los benchma ks disponibles no
compilan con ninguno de los compilado es exis en es,
sal o con adas excepciones. El compilado de PGI es
el ´unico que log a compila un n´ume o signi ica i o
de benchma ks.
La sui e se compone de benchma ks que dan el
iempo de ejecuci´on, po lo que se ´ıan un buen pun o
de pa ida pa a un an´alisis de endimien o. Po des-
g acia, la escasa madu ez del c´odigo y la poca com-
pa ibilidad con los compilado es disponibles hacen
imposible su uso pa a es ablece una compa a i a de
endimien os.
III. Compilado es Sopo ados
En la in oducci´on hemos enume ado b e emen e
los compilado es disponibles. Su elecci´on es ´a mo-
i ada po la exis encia de licencias g a ui as o
acad´emicas, o bien po se de c´odigo abie o. A con-
inuaci´on explica emos algunos de alles de los compi-
lado es sopo ados en la e si´on p elimina de TOR-
MENT OpenACC.
A. PGI Compile
El compilado de PGI [3], desa ollado po The
Po land G oup y N idia, es a echa de edacci´on de
es e a ´ıculo el compilado con un g ado de madu ez
m´as al o. Su uso es ´a muy ex endido en los di e -
en es alle es y con e encias que se ealizan sob e
OpenACC. Ac ualmen e es ´a disponible como pa e
del N idia OpenACC oolki que incluye una licencia
g a ui a de es meses y la posibilidad de adqui i
una licencia come cial o acad´emica. En nues os es-
udios, el compilado de PGI ha demos ado se el
m´as s´olido y el que m´as al o g ado de madu ez iene.
Se ajus a bien a la especi icaci´on de OpenACC. Su
ins alaci´on es simple y dispone de abundan e docu-
men aci´on.
B. OpenUH
El compilado OpenUH, desa ollado po la Uni-
e sidad de Hous on, es una al e na i a de c´odigo
abie o. Compa ado con el compilado de PGI, su
madu ez y solidez son meno es y ca ece de algunas
uncionalidades, como educciones sob e m´ınimos o
m´aximos, lo cual di icul a lige amen e la p og a-
maci´on de aplicaciones. Su ins alaci´on, a echa de
edacci´on de es e a ´ıculo, no es an sencilla como
cab ´ıa espe a . La e si´on p e-compilada disponible
en su web [4] ca ece algunas lib e ´ıas que deben se
compiladas a pa e u ob enidas de o a o ma.
C. accULL
El compilado accULL [5], desa ollado po la Uni-
e sidad de La Laguna, es al igual que OpenUH
un compilado de c´odigo abie o. De los es com-
pilado es sopo ados es el que iene una meno
solidez, eniendo p oblemas pa a la aducci´on uen e
a uen e de algunas ca ac e ´ıs icas de C, como pun-
e os a unci´on o de uncionalidades de OpenACC
como algunos ipos de educciones. La ins alaci´on
de la e si´on disponible en su web [11] es bas an e
simple siguiendo las ins ucciones que pueden encon-
a se adjun as al compilado en iche os de ex o.
IV. TORMENT OpenACC
Nues a p opues a en desa ollo se denomina
TORMENT OpenACC y es el ac ´onimo pa a T asgo
pe ORMance and E alua ioN Tool o OpenACC.
El p oyec o ha nacido pa a in en a da una soluci´on
a la necesidad de da espues a a los in en os de e-
aliza compa a i as de endimien o de c´odigo Ope-
nACC.
A. Mo i aci´on
Las he amien as ac uales han demos ado en nue-
s os es udios p e ios no se su icien es pa a da una
espues a sencilla a las compa a i as de c´odigo gen-
e ado po los di e en es compilado es de OpenACC.
En nues os es udios con dichas he amien as hemos
encon ado una se ie de p oblemas que es amos in-
en ando e i a en nues a p opues a.
El p ime o de los p oblemas es la escasa com-
pa ibilidad del c´odigo de los benchma ks en e los
di e en es compilado es. Es cie o que si el c´odigo
se ajus a al es ´anda debe ´ıa compila se y ejecu-
a se co ec amen e en cualquie implemen aci´on del
es ´anda OpenACC, pe o en la ac ual e apa en la
que se encuen an an o el es ´anda como los compi-
lado es es muy di ´ıcil. Debido a es o hemos in en-
ando desa olla TORMENT OpenACC eniendo
en cuen a an o la especi icaci´on del es ´anda como
las capacidades ac uales de los di e en es compi-
lado es.
O o p oblema de ec ado es que muchos de
los benchma ks no o ecen una idea cla a de
endimien o, como es el caso de los mic obenchma ks
del EPCC OpenACC Benchma k Sui e. Si bien es
cie o que es os mic obenchma ks son ´u iles pa a el
desa ollo de los compilado es, c eemos que el o e -
head de la implemen aci´on de los dis in os p agmas
no es ealmen e ele an e en una compa a i a de
endimien os. Po es e mo i o, en nues a p opues a
dejamos ue a ese ipo de p uebas.
Finalmen e, hemos obse ado ambi´en que, en mu-
chos casos, benchma ks que pa ecen compila co ec-
amen e luego gene an e o es en iempo de ejecuci´on
o esul ados inco ec os. Es a ´ul ima si uaci´on puede
supone que, si no se de ec a el esul ado e ´oneo, los
esul ados de endimien o ob enidos sean ambi´en
inco ec os. Pa a e i a es o, adem´as de inco po-
a comp obaci´on de esul ados a los benchma ks
ambi´en analizamos los esul ados du an e el desa -
ollo de nues a p opues a pa a que los benchma ks
que inco po amos en nues a he amien a compilen
y ejecu en co ec amen e con cualquie a de los com-
pilado es.
O o aspec o impo an e que hemos enido en
cuen a a la ho a de lle a a cabo el desa ollo de
TORMENT OpenACC es la o ece a la comu-
nidad una he amien a que adem´as de acili a el
an´alisis de endimien o del c´odigo gene ado po los
compilado es de OpenACC ejecu ados en dis in as
m´aquinas, o ezcan una medida del endimien o que
pe mi a una compa aci´on ´acil en e m´aquinas y
compilado es. Es a idea nos ha lle ado al desa ollo
de la m´e ica TORMENT ACC que desc ibi emos
pos e io men e.
B. Obje i os
El obje i o p incipal de TORMENT OpenACC es
la de pe mi i un an´alisis de endimien o de c´odigo
OpenACC gene ado po los dis in os compilado es
de o ma sencilla, gene ando un esumen de esul a-
dos ´acilmen e analizable y dando un alo , denom-
inado pun uaci´on TORMENT ACC, que pe mi e la
compa aci´on de los pa es m´aquina-compilado .
Nues a p opues a p e ende acili a a la comu-
nidad una he amien a p epa ada espec´ı icamen e
pa a OpenACC y el es ado de desa ollo emp ano
de los compilado es exis en es de o ma que se e i en
p oblemas en compilaci´on o ejecuci´on de o as he -
amien as de benchma king exis en es. TORMENT
OpenACC es a ´a p epa ado pa a compila se y ejecu-
a se con la meno in e enci´on posible del usua io.
Recopilando la in o maci´on del p oceso y o eciendo
inalmen e un in o me en o ma o HTML con los
da os ele an es.
Adem´as, TORMENT OpenACC u iliza ´a los com-
pilado es GCC y NVCC pa a ob ene da os de ejecu-
ciones de c´odigo secuencial y c´odigo CUDA espec i-
amen e. De es e modo, nues a p opues a o ece ´a
al usua io in o maci´on del speedup con espec o al
c´odigo secuencial y CUDA ejecu ado en la misma
m´aquina.
C. Es uc u a de la he amien a
TORMENT OpenACC se compone de una se ie de
sc ip s que se enca gan de odo el p oceso de com-
pilaci´on, ejecuci´on y ob enci´on de esul ados, elim-
inando es a ca ga al usua io. Es os sc ip s se di i-
den en es ca ego ´ıas. Los sc ip s de con igu aci´on
gu´ıan al usua io en la ob enci´on de las u as co -
ec as a lib e ´ıas CUDA (necesa ias po algunos de
los compilado es), u as de compilado es y coman-
dos de ejecuci´on (en caso de que el usua io u ilice,
po ejemplo, sis emas de colas ipo slu m). El sc ip
de ejecuci´on se enca ga de la compilaci´on y ejecuci´on
usando odos los compilado es que hayan sido halla-
dos en el sis ema del usua io. Finalmen e, el sc ip
de gene aci´on de esul ados p ocesa los esul ados
ob enidos y gene a un iche o HTML con el in o me
inal.
El p og ama p opiamen e dicho es ´a desa ollado
de modo que sea lo m´as sencillo posible pa a e i a
p oblemas de compilaci´on, e i ando ni eles de indi-
ecci´on que en o as si uaciones se ´ıan aconsejables.
Cada benchma k es ´a de inido en su p opia unidad
de compilaci´on y es ´a inco po ado en el mismo uen e
el c´odigo OpenACC y CUDA, u ilizando compilaci´on
condicional y e i ando ene de o ma simul anea
iche os .c y.cu con c´odigo duplicado. Los iche os .cu
son necesa ios pa a que el compilado NVCC gene e
el c´odigo CUDA co espondien e, pe o es o se ha e-
suel o median e el uso de enlaces simb´olicos a los
iche os .c co espondien es.
El p og ama p incipal se enca ga del lanzamien o
de los benchma ks. Cada uno se lanza con 10 epe i-
ciones con una epe ici´on ex a al comienzo cuyos
esul ados son desechados. Al inaliza es as diez
epe iciones se ob ienen los alo es denominados peak
ya e age y que co esponden a la mejo de las ejecu-
ciones y a la media a i m´e ica de odas ellas. Es os
alo es se i ´an pa a calcula la m´e ica de la que
habla emos pos e io men e.
CUDA_ATTR__ in ge Random(unsigned in * seed)
{
unsigned in nex = *seed;
in esul ;
nex *= 1103515245;
nex += 12345;
esul = (unsigned in ) (nex /65536) % 2048;
nex *= 1103515245;
nex += 12345;
esul <<= 10;
esul ^= (unsigned in ) (nex /65536) % 1024;
nex *= 1103515245;
nex += 12345;
esul <<= 10;
esul ^= (unsigned in ) (nex /65536) % 1024;
*seed = nex ;
e u n esul ;
}
Fig. 1
C´
odigo pa a la gene aci´
on de n´
ume os alea o ios.
D. Benchma ks implemen ados
La e si´on p elimina de TORMENT OpenACC
con iene ´unicamen e dos benchma ks. El desa ollo
de la he amien a sigue en p oceso y o os bench-
ma ks se i ´an a˜nadiendo p og esi amen e.
D.1 Mon eCa loPi
Es e benchma k consis e en una ap oximaci´on de
Pi po el m´e odo de Mon e Ca lo, que se basa en
la gene aci´on de pun os alea o ios en un cuad ado
de lado uni a io. Se comp ueba si es os pun os se
encuen an den o de un cua o de c´ı culo de adio
uni a io y se acumula el o al de pun os que cumplen
dicha condici´on. Finalmen e, se aplica la siguien e
´o mula:
π≈4∗P
T
Donde Pes el n´ume o de pun os den o del cua o
de c´ı culo y Tes el o al de pun os gene ados.
Mon eCa loPi es un benchma k que no iene
p ´ac icamen e ans e encias de memo ia y un
c´alculo compu acional muy simple, pe o puede se
op imizado en CUDA haciendo que cada hilo calcule
a ios pun os y u ilizando la sha ed memo y de los
bloques pa a e i a accesos a memo ia global. Un
buen esul ado de los compilado es en es e bench-
ma k depende ´a de es os ac o es.
Dado que no se puede hace uso de la unci´on s and
de C en el c´odigo ejecu ado en la GPU, y el uso de
la lib e ´ıa cu and se limi a a c´odigo CUDA, hemos
decidido eplica la unci´on s and pa a pe mi i su
ejecuci´on en la GPU, como se indica en la ig. 1.
El c´odigo u ilizado pa a las e siones de OpenACC
y CUDA se mues a en las igs. 2 y 3
#p agma acc pa allel loop
p i a e(i, d, x, y, seed)
educ ion(+:coun )
o (i = 0; i < COORD_NUM; ++i){
seed = 1987 ^ i*27;
x = (double)ge Random(&seed)/(double)RAND_MAX;
y = (double)ge Random(&seed)/(double)RAND_MAX;
d = sq (x*x + y*y);
i (d <= 1.0){
++coun ;
}
}
Fig. 2
C´
odigo de Mon eCa loPi pa a OpenACC.
__CUDA_GLOBAL__ oid piKe nel(
cons unsigned long in iesPe Th ead,
unsigned long in * hi s)
{
unsigned in seed;
in gid, id, bid;
in lhi s;
loa x, y;
ex e n __sha ed__ unsigned long in sda a[];
gid = (blockIdx.x*blockDim.x) + h eadIdx.x;
id = h eadIdx.x;
bid = blockIdx.x;
lhi s = 0;
seed = 1987 ^ gid*27;
o (in i = 0; i < iesPe Th ead; ++i){
x = ( loa )ge Random(&seed)/( loa )RAND_MAX;
y = ( loa )ge Random(&seed)/( loa )RAND_MAX;
loa d = sq (x*x + y*y);
i (d <= 1.0 ){
++lhi s;
}
}
sda a[ id] = lhi s;
__sync h eads();
i ( id == 0){
o (in i = 1; i < blockDim.x; ++i){
lhi s += sda a[i];
}
hi s[bid] = lhi s;
}
}
Fig. 3
C´
odigo de Mon eCa loPi pa a CUDA.
#p agma acc da a copyin(S[0:sizeS],B[0:sizeB])
copy(C[0:sizeB])
{
#p agma acc pa allel loop p i a e(s a B)
i s p i a e( esul ,sizeS,sizeB)
o (s a B = 0; s a B <= sizeB-sizeS; s a B++){
C[s a B] = 0;
i (s a B <= esul ) {
in ind;
o (ind = 0; ind < sizeS; ind++){
i (S[ind] != B[s a B+ind]) b eak;
}
i (ind == sizeS){
esul = s a B;
C[s a B] = 1;
}
}
}
}
esul = -1;
o (s a B = 0; s a B <= sizeB-sizeS; s a B++){
i (C[s a B] == 1){
esul = s a B;
b eak;
}
}
Fig. 4
C´
odigo de S ingMa ch pa a OpenACC.
D.2 S ingMa ch
El benchma k S ingMa ch es un p og ama de
alineamien o de cadenas de ca ac e es. El p og ama
consis e en la b´usqueda de la p ime a ocu encia
de una cadena peque˜na en una cadena g ande, u i-
lizando un algo i mo nai e. En es e benchma k, la
cadena g ande iene una longi ud de 10 millones de
ca ac e es, es deci 10MB. Las cadenas peque˜nas,
que son cua o, ienen una longi ud de 1000 ca ac-
e es o 1KB.
Es e algo i mo es in e esan e po que combina
ans e encia de da os con uso e icien e de la memo-
ia, especialmen e el uso de la sha ed memo y. Los
c´odigos de las e siones de OpenACC y CUDA
pueden e se en las igs. 4 y 5.
E. M´e ica u ilizada
Pa a la elecci´on de la m´e ica denominada
TORMENT ACC hemos decidido op a po la
me odolog´ıa SPEC [12]. El Sys em Pe o mance
E alua ion Coope a i e, com´unmen e conocido como
SPEC, es un e e en e ampliamen e conocido en
cuan o a benchma king y an´alisis de endimien o se
e ie e. Una de sus o alezas es el econoce que los
benchma ks en ejecen en unci´on del paso del iempo
y, en consecuencia, deben se ac ualizados.
SPEC u iliza la siguien e me odolog´ıa. En p ime
luga , cada p og ama de uel e su iempo de eje-
cuci´on y se calcula el SPEC a io, que consis e en el
a io ob enido de di idi un iempo de ejecuci´on de
e e encia suminis ado po SPEC en e el iempo de
ejecuci´on ob enido. Finalmen e, se ob iene la media
geom´e ica de odos los SPEC a ios del conjun o de
__global__ oid ke nel_busqueda(cons cha * b_ida a,
cons in sizeB, cons cha * s_ida a,
cons in sizeS, unsigned in * esul )
{
ex e n __sha ed__ unsigned cha sda a[];
in gid = (blockIdx.x*blockDim.x) + h eadIdx.x;
in id = h eadIdx.x;
in sS a = THREADS_PER_BLOCK + sizeS;
in o se ;
i (gid == 0){
* esul = -1;
}
__sync h eads();
i (gid < * esul ){
unsigned cha * bd = &sda a[ id];
b_ida a += gid;
unsigned cha * sd = &sda a[sS a + id];
s_ida a += id;
i ( id<THREADS_PER_BLOCK){
o (o se = 0; id+o se <sS a ;
o se += THREADS_PER_BLOCK)
{
i (o se ==0){
*bd = *b_ida a;
bd+=THREADS_PER_BLOCK;
b_ida a+=THREADS_PER_BLOCK;
con inue;
}
*bd = *b_ida a;
bd+=THREADS_PER_BLOCK;
b_ida a+=THREADS_PER_BLOCK;
*sd = *s_ida a;
sd+=THREADS_PER_BLOCK;
s_ida a+=THREADS_PER_BLOCK;
}
}
__sync h eads();
i (gid <= sizeB-sizeS){
unsigned in i;
unsigned in b = 1;
unsigned cha * sd = &sda a[sS a ];
unsigned cha * bd = &sda a[ id];
o (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++;
i (!b) b eak;
}
o (; b && i < sizeS ; i++){
b &= (*sd == *bd); sd++; bd++;
i (!b) b eak;
}
i ( b ) a omicMin( esul , gid);
}
}
}
Fig. 5
C´
odigo de S ingMa ch pa a CUDA.

benchma ks [13].
Nues a p opues a sigue una idea simila a la de
SPEC, pe o con algunas a iaciones. En p ime lu-
ga , la ejecuci´on de los benchma ks que componen
TORMENT OpenACC de uel en es alo es. Peak
Time es el mejo iempo de ejecuci´on ob enido, me-
dido en segundos. A e age Time es el iempo medio
de odas las ejecuciones del benchma k, ambi´en
en segundos. Finalmen e, S anda d De ia ion es la
des iaci´on es ´anda del conjun o de medidas e indica
la a iabilidad ob enida en las ejecuciones del bench-
ma k. Con los iempos peak ya e age se calcula un
a io con espec o a los iempos de e e encia sum-
inis ados con la he amien a, y que son los iem-
pos de ejecuci´on secuencial del benchma k en una
m´aquina de e e encia. Una ez ejecu ados odos los
benchma ks, se ob iene la media a m´onica de odos
los a ios, an o pa a peak ime como a e age ime.
Es os alo es que se ob ienen son las pun uaciones
TORMENT ACC peak ya e age.
La p incipal di e encia en e TORMENT Ope-
nACC y SPEC, a pa e de la me odolog´ıa de oma de
iempos de ejecuci´on, consis e en el uso de la media
a m´onica en luga de la media geom´e ica. Es a de-
cisi´on se undamen e en el hecho de que, pa a los ob-
je i os de TORMENT OpenACC la media a m´onica
iene m´as en ajas que la media geom´e ica. En
p ime luga , aunque la media geom´e ica siemp e
p oduce una o denaci´on consis en e, no necesa ia-
men e p oduce la o denaci´on co ec a [13] ya que es a
media no es in e samen e p opo cional al iempo de
ejecuci´on. En cambio, la media a m´onica si que es
in e samen e p opo cional al iempo de ejecuci´on lo
que hace que sea una media co ec a pa a exp esa
a ios. Es as a i maciones son compa idas po o os
a ´ıculos como [14].
V. Conclusiones y T abajo Fu u o
Como conclusi´on, TORMENT OpenACC es una
he amien a de an´alisis y compa aci´on de endimien-
os de c´odigo gene ado po compilado es de Ope-
nACC, eniendo en conside aci´on el ni el de madu ez
de an o el es ´anda OpenACC como de los di e en es
compilado es. TORMENT OpenACC desa olla una
sui e de benchma ks espec´ı icamen e dise˜nados pa a
OpenACC y man eniendo la m´axima po abilidad
en e compilado es, pe mi iendo su compilaci´on y
ejecuci´on en odos ellos.
Los esul ados o ecidos po TORMENT Ope-
nACC incluyen la denominada pun uaci´on TOR-
MENT ACC, dise˜nada pa a la compa aci´on de pa es
m´aquina-compilado . Adem´as, se o ece un esumen
de la ejecuci´on de los benchma ks con una abla de
iempos y a ios, an o m´ınimos como medios e in-
cluyendo la des iaci´on es ´anda pa a un an´alisis de
a iabilidad de los iempos de ejecuci´on del c´odigo
gene ado.
Jun o con los esul ados de los compilado es de
OpenACC se incluyen ambi´en esul ados de eje-
cuci´on de c´odigo gene ado po los compilado es GCC
y NVCC pa a c´odigo secuencial y CUDA. De es e
modo, se puede o ece al usua io una compa a i a a
ni el de m´aquina del endimien o del c´odigo gene -
ado po los compilado es de OpenACC con espec o
a e siones secuenciales y CUDA de los benchma ks.
El abajo u u o a desa olla consis e en la am-
pliaci´on de la sui e de benchma ks, pa a log a esul-
ados que sean e dade amen e ele an es. T a ando
de cub i los aspec os m´as in e esan es de la eje-
cuci´on de c´odigo en las GPUs. Adem´as, una pa e
impo an e del abajo es an e consis e en man ene
la compa ibilidad en e compilado es y asegu a el
co ec o uncionamien o de la he amien a en dis in-
as m´aquinas.
Ag adecimien os
En memo ia de Agus ´ın de Dios He n´andez.
Es a in es igaci´on ha sido pa cialmen e inanciada
po el MICINN y el p og ama ERDF de la Uni´on Eu-
opea: p oyec o HomP og-He Sys (TIN2014-58876-
P), la ed CAPAP-H5 (TIN2014-53522-REDT) y el
COST P og am Ac ion IC1305: Ne wo k o Sus-
ainable Ul ascale Compu ing (NESUS).
Re e encias
[1] OpenACC-s anda d.o g, “Abou OpenACC,” .
[2] OpenACC-s anda d.o g, “OpenACC 2.5 d a o public
commen ,” oc 2015.
[3] PGI, “Pgi accele a o compile s wi h OpenACC di ec-
i es,” h ps://www.pg oup.com/ esou ces/accel.h m,
no 2015.
[4] Uni e si y o Hous on, “Open-sou ce UH compile ,”
h p://web.cs.uh.edu/~openuh/download/, no 2015.
[5] Ruym´an Reyes, I ´an L´opez-Rod ´ıguez, Juan J Fume o,
and F ancisco de Sande, “accULL: an OpenACC imple-
men a ion wi h CUDA and OpenCL suppo ,” in Eu o-
Pa 2012 Pa allel P ocessing, pp. 871–882. Sp inge ,
2012.
[6] EPCC, “Epcc OpenACC benchma k sui e,” h ps:
//gi hub.com/EPCCed/epcc-openacc-benchma ks, sep
2013.
[7] Pa hscale, “Rodinia benchma k sui e 2.1 wi h Ope-
nACC po ,” h ps://gi hub.com/pa hscale/ odinia,
ap 2014.
[8] Shuai Che, Michael Boye , Jiayuan Meng, Da id Ta jan,
Je emy W Shea e , Sang-Ha Lee, and Ke in Skad on,
“Rodinia: A benchma k sui e o he e ogeneous com-
pu ing,” in Wo kload Cha ac e iza ion, 2009. (IISWC),
2009 IEEE In e na ional Symposium on. IEEE, 2009,
pp. 44–54.
[9] Louis-No¨el Pouche , “Polybench: The polyhed al
benchma k sui e,” URL: h p://www. cs. ucla. edu/˜
pouche /so wa e/polybench/[ci ed July,], 2012.
[10] Shuai Che, Je emy W Shea e , Michael Boye , Lukasz G
Sza a yn, Liang Wang, and Ke in Skad on, “A cha ac e -
iza ion o he Rodinia benchma k sui e wi h compa ison
o con empo a y CMP wo kloads,” in Wo kload Cha -
ac e iza ion (IISWC), 2010 IEEE In e na ional Sympo-
sium on. IEEE, 2010, pp. 1–11.
[11] Uni e sidad de La Laguna, “accULL,” h p://cap.pcg.
ull.es/es/accULL, no 2015.
[12] Kai alya M Dixi , “The spec benchma ks,” Pa allel
compu ing, ol. 17, no. 10, pp. 1195–1209, 1991.
[13] Da id J Lilja, Measu ing compu e pe o mance: a p ac-
i ione ’s guide, Camb idge Uni e si y P ess, 2005.
[14] John R Mashey, “Wa o he benchma k means: ime o
a uce,” ACM SIGARCH Compu e A chi ec u e News,
ol. 32, no. 4, pp. 1–14, 2004.
Benchma k e sion:0.91
Hos name:hyd a
Use name:daniel
So wa es ackin o
Ke nel:Linux3.10.0229.4.2.el7.x86_64x86_64
GCC:gcc(GCC)4.8.320140911(RedHa 4.8.39)
NVCC:Cudacompila ion ools, elease7.5,V7.5.17
PGICompile :pgcc15.7064bi  a ge onx8664Linux phaswell
OpenUHCompile :OpenUH3.1.0(basedonOpen64Compile Sui e:Ve sion5.0)
accULLCompile :Release0.4alpha
CPUin o
Model:In el(R)Xeon(R)CPUE52609 3@1.90GHz
A chi ec u e:x86_64
Numbe o Co es:6
MaxMHz:MHz
MinMHz:MHz
L1Cache:32K
L2Cache:256K
L3Cache:15360K
RAM:65687144kB
GPU(s)in o
GPU0:NVIDIACo po a ionGK110B[GeFo ceGTXTi anBlack]
( e a1)
GPU1:NVIDIACo po a ionGK110B[GeFo ceGTXTi anBlack]
( e a1)
GPU2:NVIDIACo po a ionGK110B[GeFo ceGTXTi anBlack]
( e a1)
GPU3:NVIDIACo po a ionGK110B[GeFo ceGTXTi anBlack]
( e a1)
TORMENT_ACC_0.91Peak:28.26 TORMENT_ACC_0.91A e age:25.71
Speedup sSequen ial(peak):57.92x Speedup sSequen ial(a e age):52.40x
Speedup sCUDA(peak):0.17x Speedup sCUDA(a e age):0.15x
TORMENT_ACC_0.91Peak:12.37 TORMENT_ACC_0.91A e age:12.30
Speedup sSequen ial(peak):25.35x Speedup sSequen ial(a e age):25.07x
Speedup sCUDA(peak):0.07x Speedup sCUDA(a e age):0.07x
TORMENT_ACC_0.91Peak:16.14 TORMENT_ACC_0.91A e age:16.07
Speedup sSequen ial(peak):33.08x Speedup sSequen ial(a e age):32.76x
Speedup sCUDA(peak):0.10x Speedup sCUDA(a e age):0.10x
Sys emIn o ma ion
GCC_Sequen ialCompile Resul s:
GCC_Sequen ial
Benchma k PeakTime(s) PeakRa io A g.Time(s) A g.Ra io S d.De ia ion
Mon eCa loPi 14.763701 0.43 14.764007 0.43 0.000336
S ingMa ch 20.118267 0.57 20.136980 0.57 0.021079
NVCC_CUDACompile Resul s:
NVCC_CUDA
Benchma k PeakTime(s) PeakRa io A g.Time(s) A g.Ra io S d.De ia ion
Mon eCa loPi 0.016409 383.96 0.017840 354.36 0.001254
S ingMa ch 0.105644 108.48 0.105889 109.21 0.000162
PGICompile Resul s:
PGI
Benchma k PeakTime(s) PeakRa io A g.Time(s) A g.Ra io S d.De ia ion
Mon eCa loPi 0.067226 93.72 0.077228 81.86 0.008222
S ingMa ch 0.688718 16.64 0.758494 15.25 0.068944
OpenUHCompile Resul s:
OpenUH
Benchma k PeakTime(s) PeakRa io A g.Time(s) A g.Ra io S d.De ia ion
Mon eCa loPi 0.153154 41.14 0.165321 38.24 0.006447
S ingMa ch 1.574083 7.28 1.578459 7.33 0.002537
accULLCompile Resul s:
accULL
Benchma k PeakTime(s) PeakRa io A g.Time(s) A g.Ra io S d.De ia ion
Mon eCa loPi 0.135701 46.43 0.139198 45.42 0.007257
S ingMa ch 1.173044 9.77 1.184555 9.76 0.004565
Peaksco euses hebes o 10 esul s.A e agesco euses hea i hme icmeano 10 esul s.Inbo hcaseshighe isbe e .
TORMENTOpenACCResul page
G upoT asgoUni e sidaddeValladolid
Benchma k e sion:0.91
Hos name:co ikLap op
Use name:dani
So wa es ackin o
Ke nel:Linux3.16.04amd64x86_64
GCC:gcc(Debian4.9.210)4.9.2
NVCC:Cudacompila ion ools, elease7.5,V7.5.17
PGICompile :pgcc15.7064bi  a ge onx8664Linux phaswell
OpenUHCompile :OpenUH3.1.0(basedonOpen64Compile Sui e:Ve sion5.0)
accULLCompile :Release0.4alpha
CPUin o
Model:In el(R)Co e(TM)i54200HCPU@2.80GHz
A chi ec u e:x86_64
Numbe o Co es: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)in o
GPU0:NVIDIACo po a ionGM107M[GeFo ceGTX850M]( e
a2)
TORMENT_ACC_0.91Peak:8.98 TORMENT_ACC_0.91A e age:8.22
Speedup sSequen ial(peak):8.84x Speedup sSequen ial(a e age):8.05x
Speedup sCUDA(peak):0.17x Speedup sCUDA(a e age):0.16x
TORMENT_ACC_0.91Peak:6.56 TORMENT_ACC_0.91A e age:6.60
Speedup sSequen ial(peak):6.46x Speedup sSequen ial(a e age):6.46x
Speedup sCUDA(peak):0.13x Speedup sCUDA(a e age):0.13x
TORMENT_ACC_0.91Peak:6.50 TORMENT_ACC_0.91A e age:6.54
Speedup sSequen ial(peak):6.40x Speedup sSequen ial(a e age):6.40x
Speedup sCUDA(peak):0.12x Speedup sCUDA(a e age):0.13x
Sys emIn o ma ion
GCC_Sequen ialCompile Resul s:
GCC_Sequen ial
Benchma k PeakTime(s) PeakRa io A g.Time(s) A g.Ra io S d.De ia ion
Mon eCa loPi 6.196389 1.02 6.196760 1.02 0.000310
S ingMa ch 11.292974 1.01 11.298524 1.02 0.006537
NVCC_CUDACompile Resul s:
NVCC_CUDA
Benchma k PeakTime(s) PeakRa io A g.Time(s) A g.Ra io S d.De ia ion
Mon eCa loPi 0.057900 108.81 0.062234 101.58 0.004161
S ingMa ch 0.335251 34.19 0.335390 34.48 0.000085
PGICompile Resul s:
PGI
Benchma k PeakTime(s) PeakRa io A g.Time(s) A g.Ra io S d.De ia ion
Mon eCa loPi 0.343979 18.32 0.378440 16.70 0.034409
S ingMa ch 1.925690 5.95 2.120944 5.45 0.192720
OpenUHCompile Resul s:
OpenUH
Benchma k PeakTime(s) PeakRa io A g.Time(s) A g.Ra io S d.De ia ion
Mon eCa loPi 0.697705 9.03 0.698031 9.06 0.000198
S ingMa ch 2.224097 5.15 2.226121 5.19 0.003758
accULLCompile Resul s:
accULL
Benchma k PeakTime(s) PeakRa io A g.Time(s) A g.Ra io S d.De ia ion
Mon eCa loPi 0.675184 9.33 0.676132 9.35 0.000704
S ingMa ch 2.298089 4.99 2.299703 5.03 0.003703
Peaksco euses hebes o 10 esul s.A e agesco euses hea i hme icmeano 10 esul s.Inbo hcaseshighe isbe e .
TORMENTOpenACCResul page
G upoT asgoUni e sidaddeValladolid