Una herramienta de benchmarking para compiladores de OpenACC

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