Técnicas de implementación de Stencils en multi-GPU distribuidas

T´ecnicas de implemen aci´on de S encils en mul i-GPU dis ibuidas Senmao Ji Ye 1, A u o Gonz´alez Esc ibano2, Diego R. Llanos3 Resumen— En el pa ´on de compu aci´on denomi- nado s encil cada elemen o de una es uc u a de da os de ipo a ay se ac ualiza i e a i amen e en unci´on de los alo es de sus ecinos. En e o as aplica- ciones, es e pa ´on pe mi e esol e num´e icamen e sis emas de ecuaciones en de i adas pa ciales, po lo que es de g an in e ´es en el compu o cien ´ı ico, c e- ciendo incesan emen e los eque imien os de ama˜no de da os y ca ga compu acional en p oblemas eales. La es uc u a de es e pa ´on pe mi e u iliza es a e- gias sencillas de pa alelismo de da os, po lo que su pa alelizaci´on, an o en CPUs como en acele - ado es es de g an in e ´es. Sin emba go, la necesi- dad de sinc onizaci´on y comunicaci´on en e elemen- os de p oceso de i a en p oblemas elacionados con la capacidad de dis ibui la ca ga y explo a mul i- ples disposi i os simul aneamen e. En es e abajo p esen amos un epaso y ac ualizaci´on de ´ecnicas de p og amaci´on e icien es basadas en MPI y CUDA pa a explo a es e pa ´on de compu aci´on en sis emas mul i-GPU dis ibuidos. Nues os esul ados mues- an c´omo las ´ecnicas u ilizadas pueden ali ia los p oblemas de comunicaci´on en e hos y GPUs, ob e- niendo endimien os y escalabilidad en unci´on de las capacidades del sis ema de in e conexi´on en e nodos. Palab as cla e— S encil, Mul i-GPU, compu aci´on dis ibuida I. In oducci´ on PARA la esoluci´on de nume osos p oblemas cien ´ı icos ac uales, en los que se aplican ecua- ciones de i adas pa ciales, ales como el c´alculo de la p opagaci´on del sonido o del calo , se u iliza una ´ecnica de compu aci´on llamada S encil. Una ´ecnica S ecil consis e en una i e aci´on empo al en la que en cada i e aci´on se ac ualizan odas y cada una de las celdas de una ma iz mul idimensional, a pa i de los alo es an e io es de un conjun o de celdas eci- nas. Cada p oblema de ine cu´an as celdas ecinas se u ilizan y qu´e pesos se aplican a cada una an es de aplica la unci´on que calcula el nue o alo de la celda que es ´a siendo p ocesada. La u ilidad de es e pa ´on y su es uc u a que pe - mi e la aplicaci´on de ´ecnicas sencillas de pa alelismo de da os, hace que haya sido ampliamen e es udiado en la li e a u a. Desde implemen aciones sencillas ilus a i as de ´ecnicas b´asicas de p og amaci´on dis- ibuida (p.e. [1]), o e siones op imizadas en acele - ado es ha dwa e (p.e. [2]). Se han p opues o ´ecnicas pa a su explo aci´on e icien e en sis emas mul i-GPU (p.e. [3]), o en cl´us e s he e og´eneos (p.e. [4]). En es e abajo p esen amos un epaso ac ual- 1Dp o. de In o m´a ica, Uni . Valladolid, Espa˜na. e-mail: [email p o ec ed]. 2Dp o. de In o m´a ica, Uni . Valladolid, Espa˜na. e-mail: [email p o ec ed]. 3Dp o. de In o m´a ica, Uni . Valladolid, Espa˜na. e-mail: [email p o ec ed]. izado de ´ecnicas de implemen aci´on de s encils pa a sis emas dis ibuidos de acele ado es GPU, o ien- adas ocul a las la encias de las comunicaciones necesa ias en cada i e aci´on empo al. U ilizando como caso de es udio un s encil b´asico con una ca ga de c´ompu o m´ınima po elemen o (el m´e odo de Jacobi), mos amos con di e en es ap oximaciones ´ecnicas e icien es y poco complejas pa a solapa compu aci´on y comunicaci´on. Es udiamos el im- pac o de las copias as´ınc onas en e GPU y hos , y en e nodos de la ed, explo amos uncionalidades del API de CUDA no u ilizadas en es udios p e ios, y ex- plo amos la escalabilidad en p esencia de ecnolog´ıas de ed asequibles y de ´acil acceso. Nues os esul- ados mues an c´omo las ´ecnicas u ilizadas pueden ali ia los p oblemas de comunicaci´on en e hos y disposi i os acele ado es, as´ı como en e nodos, pa a p oblemas de ama˜nos adecuados pa a los espacios de memo ia de los disposi i os acele ado es ac uales. El es o del a ´ıculo se o ganiza de la siguien e o ma. En la secci´on 2 se mues a una desc ipci´on m´as de allada de los pa ones s encil y del caso de es udio. En la secci´on 3 se desc iben las ´ecnicas u i- lizadas. En la secci´on 4 se p esen an un es udio ex- pe imen al pa a de e mina la u ilidad de las ´ecnicas dicu idas. La secci´on 5 p esen a las conclusiones y desc ibe el abajo u u o. II. Compu aci´ on S encil Los p og amas basados en compu aci´on de ipo s encil epi en i e a i amen e la ac ualizaci´on de odas las celdas de un a ay mul idimensional en unci´on de los alo es an e io es de celdas ecinas has a una condici´on de e minaci´on que puede de- pende de los alo es calculados. Las celdas ecinas se de inen en unci´on de su posici´on ela i a espec o al alo de los ´ındices de cada dimensio´n de la celda que se compu a. Cada p oblema de ine un pa ´on o s encil, que indica qu´e celdas se conside an ecinas, apo ando su alo a la unci´on que calcula el nue o alo de una celda, y opcionalmen e pesos a aplica de- pendiendo de la posici´on ela i a de la cada celda ecina espec o a la que se es ´a compu ando. Uno de los m´e odos s encil m´as simples y es udiados es el m´e odo de Jacobi ( e Fig. 1). En es e c´odigo se conside an como celdas ecinas las que ienen´ındices mul idimensionales que a ´ıan con espec o a los de la celda ac ual en s´olo una unidad, en un ´unico´ındice. La unci´on de ac ualizaci´on es la media a i m´e ica de los alo es ecinos, odos con el mismo peso. El pa ´on o s encil se puede complica con cel- das ecinas a mayo dis ancia ( adio del s encil), de o ma no equilib ada en las di e en es dimensiones o ejes [5], has a llega a m´e odos complejos con pa- ones que se al e an de o ma din´amica mien as a anzan las i e aciones. A. Jacobi 2D o (i=1; i< ows-1; i++) { o (j=1; j<columns-1; j++) { ma ix[i][j] = ( ma ixCopy[i-1][j] + ma ixCopy[i+1][j] + ma ixCopy[i][j-1] + ma ixCopy[i][j+1] ) / 4; } } Fig. 1. Codigo de Jacobi en 2D El m´e odo Jacobi aplicado a una ma iz de dos di- mensiones ( e Fig. 1) calcula el nue o alo de cada celda c´omo la media a i m´e ica de los alo es an e- io es de los cua o ecinos adyacen es. U ilizando una es uc u a de apoyo pa a almacena los alo es de odos los elemen os de la ma iz en la i e aci´on an e io , es e p oceso es ideal pa a se ealizado en pa alelo, ya que odos los alo es de salida son inde- pendien es en e s´ı. La dependencia del c´alculo de una celda, con los alo es de las ecinas calculados en la i e aci´on an e- io , implica que en ejecuciones con mul iples p oce- sos, un p oceso equie a pa a el c´ompu o de la pa e que se le asigna, de alo es que se encuen an en un p oceso dis in o. Pa a pode calcula los alo es de las celdas si uadas en los bo des de las subma ices de cada p oceso (bounda y), una soluci´on ´ıpica es amplia las subma ices locales con una ila/columna de celdas en cada di ecci´on, pa a almacena los da os adyacen e a las bounda ies, que han sido calcula- dos en el p oceso ecino en la i e aci´on an e io ( e Fig. 3). Es as celdas adicionales eciben el nomb e de halo. T as calcula odo los alo es de la subma- iz, cada p oceso comunica los da os ac ualizados de sus bo des (bounda ies), en i´andolos hacia los halos ecinos. que se encuen an en o os p ocesos (E), ecibiendo a su ez los da os de los bo des (bound- a ies) de los p ocesos ecinos, ecinas (A)(B)(C) y (D), que se almacena ´an en los halos locales. En ejecuciones con un ´unico disposi i o, es a de- pendencia no nos p oduce muchos incon enien es, ya que disponemos de oda la ma iz en la misma memo ia global, po lo que se puede accede desde cualquie pun o independien emen e de po qui´en ha sido ac ualizado en la i e aci´on an e io . El ´unico equisi o es una ba e a de sinc onizaci´on de los el- emen os de p oceso implicados en e una i e aci´on y o a. O as ´ecnicas asociadas a la explo aci´on de la je a qu´ıa de memo ia son posibles. Po ejemplo, en uno de los p og amas de ejemplo que acompa˜nan al Toolki de CUDA suminis ado po NVIDIA, que implemen a una con oluci´on de imagen basada en el m´e odo de Jacobi, cada bloque de hilos ealiza la copia local de la subma iz co espondien e en memo ia compa ida, ademas del halo de dicha sub- ma iz [6]. Los accesos a memo ia compa ida son mucho m´as ´apidos, con lo que se acele an los accesos a los elemen os cuando se eu ilizan, en el momen o en que se calcula cada elemen o ecino. __sha ed__ loa sDa a [THREADS_BLOCK_Y + 2 * RADIUS] [THREADS_BLOCK_X + 2 * RADIUS]; unsigned in index = (i)* Nj + (j) ; i ( li<RADIUS ) // copy op and bo om halo { //Copy Top Halo Elemen // Bounda y check i (blockIdx.y > 0) sDa a[li][e_lj] = inpu [index - RADIUS * Nj]; //Copy Bo om Halo Elemen // Bounda y check i (blockIdx.y < (g idDim.y-1)) sDa a[e_li+THREADS_BLOCK_Y][e_lj] = inpu [index + THREADS_BLOCK_Y * Nj]; } i ( lj<RADIUS ) // copy le and igh halo { // Bounda y check i ( blockIdx.x > 0) sDa a[e_li][lj] = inpu [index - RADIUS]; // Bounda y check i (blockIdx.x < (g idDim.x-1)) sDa a[e_li][e_lj+THREADS_BLOCK_X] = inpu [index + THREADS_BLOCK_X]; } Fig. 2. Ex ac o del co´ıgo de Jacobi 2D implemen ado en un ejemplo de con oluci´on de imagen suminis ado con el Toolki de CUDA Pa a la implemen aci´on en un cl´us e mul i-GPU dis ibuido, se a necesa io u iliza una API pa a el paso de mensajes en e p ocesos, como po ejem- plo MPI (Message Passing In e ace). Tambi´en se ´a necesa io u iliza una API pa a la implemen aci´on del pa alelismo en el disposi i o o disposi i os lo- cales. En es e abajo, u iliza emos MPI como la API empleada pa a el paso de mensajes, y CUDA como la API pa a explo a los disposi i os GPU. El m´e odo de comunicaci´on en e p ocesos po cada i e aci´on que e aluamos en es e abajo es el siguien e: (E) (E) (E) (E) (A) (B) (D) (C) Fig. 3. Comunicaci´on en e p ocesos en Jacobi 2D. Se mues- an los halos o ex ensiones de la subma iz asignada al p oceso, y los mo ien os de da os en e bo des y halos. 1. Se p og aman cinco ke nels pa a compu a la subma iz del p oceso asignada a un disposi- i o. Un ke nel se enca ga a de la pa e cen- al, sin inclui los bo des. El es o de ke nels compu a an cada uno un bo de (a iba, abajo, de echa, izquie da), en caso de que exis a. Al epa i la ma iz en e los disposi i os, apa e- cen subma ices que ocan alguno de los bo - des. Dichas subma ices no ienen halo en esa di ecci´on, ni necesidad de comunica o p oce- sa esos bo des de o ma especial, ya que no hay o o p oceso emo o que necesi e esos da os. Cada uno de los cinco ke nels se puede ejecu a de o ma as´ınc ona. En CUDA se pueden u i- liza s eams o colas de ins ucciones di e en es pa a lanza cada ke nel pa a consegui dicha as- inc on´ıa, ejecu ando odos los ke nels de o ma concu en e. 2. En el s eam (cola de ins ucciones) donde se lanza la ejecuci´on de cada ke nel que compu a un bo de, encolamos inmedia amen e despu´es una ope aci´on de copia de los da os almacenados pa a ese bo de, en la memo ia del disposi i o, hacia la memo ia del hos . De es a o ma, en el momen o en que acabe la compu aci´on de un bo de, comienza la copia as´ınc ona de los nue os da os hacia el hos . U ilizamos la unci´on cud- aMemcpyAsync. 3. T ans e encias de da os de los bo des en e dis- posi i o y hos . Es amos abajando en C, con ma ices almacenadas en lo que se conoce como ow majo o de . Los elemen os se almacenan po ilas, de o ma que los elemen os de colum- nas consecu i as den o de la misma ila, es ´an consecu i os en memo ia. Po an o, el en io de los bo des ho izon ales (pa e de una ila) se ealizan di ec amen e con una llamada a la unci´on de copia as´ınc ona de CUDA, pasando el pun e o de comienzo del bo de y su ama˜no. En cambio, pa a los bo des la e ales ( e icales), no podemos ealiza el paso de memo ia de es a mane a ya que los elemen os del bo de no es ´an consecu i os en la memo ia. Mo e da o a da o con di e en es ope aciones de ans e encia de memo ia es muy ine icien e. En la siguien e secci´on se discu en dos mecanismos di e en es pa a mejo a es as ans e encias. Finalmen e se ealiza una sinc onizaci´on o es- pe a pa a asegu a que los da os de los cua o bo des se han ans e ido a la memo ia del hos . Se puede ealiza con llamadas a la unci´on cu- daS eamSynch onize con cada uno de los cua- o s eams asociados a las ope aciones sob e los bo des. 4. Comunicaci´on de da os de bo des en e p oce- sos MPI. Las ope aciones de comunicaci´on en MPI pueden se muy cos osas, especialmen e en- e p ocesos asignados a nodos de ed di e en es. Es undamen al que es as ope aciones se eali- cen de o ma as´ınc ona, pudiendo solapa se con la compu aci´on p incipal de la subma iz cen al en el disposi i o. Dado que el ke nel que ejecu a la pa e cen al en el disposi i o ya ha sido lanzado, cualquie ipo de ope aciones de comunicaci´on en e eci- nos es iable. En es e abajo u ilizamos un- ciones MPI Isend, MPI I ec pa a el en ´ıo y e- cepci´on de los bo des hacia los halos emo os. Se inician las ope aciones de en ´ıo y ecepci´on y se ealiza una sinc onizaci´on con MPI Wai all pa a espe a a que se eciban los bo des emo- os en los halos locales. La u ilizaci´on de op- e aciones as´ınc onas pe mi e que los mensajes se p ocesen en cuan o es ´an disponibles, inde- pendien emen e del o den en que se llama a las unciones de en ´ıo o ecepci´on. 5. Se p ocede a ans e i los bo des que acaban de llega a los halos a las co espondien es ima- genes en la memo ia del disposi i o. Se u ilizan de nue o las ´ecnicas comen adas pa a la ans- e encia in e sa. 6. Se ealiza una sinc onizaci´on global del dispos- i i o pa a asegu a que el ke nel que ejecu a la compu aci´on de la pa e p incipal en el disposi- i o ha e minado. U ilizamos una llamada a la unci´on cudaDe iceSynch onize. 7. El ejemplo escogido ejecu a un n´ume o de e - minado de i e aciones. Si se desea in oduci una condici´on de e minaci´on basada en el alo de esiduos, debe ´ıa se implemen ada en es e pun o. 8. Finalmen e, se p ocede a in e cambia los pun- e os de las es uc u as de da os de las dos ma- ices en el disposi i o, pa a e i a ene que ealiza una copia de la ma iz con los nue os da os a la ma iz donde se man ienen las copias de la i e aci´on an e io . 9. T as la inalizaci´on del bucle de i e acions, se se comp ueba el n´ume o de i e aciones ealizadas, pa a conoce donde se encuen a la ma iz esul- ado as los cambios de pun e os, e ec uando la co eci´on de pun e os si es necesa io. La ma iz esul ado inal puede se ans e ida a la memo- ia del hos pa a su esc i u a en iche o, o su u ilizaci´on en pos e io es ases de la aplicaci´on. Es e m´e odo as´ınc ono nos pe mi e ealiza las op aciones de o ma ag egada, con la mayo g an- ula idad posible, man eniendo la mayo concu en- cia posible, y solapando compu aci´on y comunicaci´on pa a educiendo no ablemen e el iempo compa ado con una implem aci´on s´ınc ono m´as di ec a y simple. III. Implemen aci´ on y op imizaciones Pa a mejo a las ans e encias de memo ia y e- duci el uido en las mediciones de los iempos, se han aplicado dis in as ´ecnicas pa a op imiza la im- plemen aci´on de la soluci´on comen ada en la secci´on an e io . A. T ans e encia de bo des e icales Discu imos dos posibles ´ecinas pa a e i a que la ans e encia de los bo des e icales en e hos y disposi i o, cuyos elemen os no son consecu i os en memo ia, u ilizando di e en es llamadas a unciones de ans e encia pa a cada elemen o. La p ime a es ´a basada en u iliza bu e s ex a pa a copia los da os de cada bo de en una memo ia con igua an es de las ans e encias. Es o nos obliga a ealiza manualmen e la p og amaci´on de es as op- e aciones de ma shalling/unma shalling en ke nels y en el hos . Rese a manualmen e espacio adicional en la memo ia de hos y disposi i o pa a dos ma i- ces unidimensionales (bu e s), uno de en ´ıo y o o de ecpci´on, po cada bo de e ical. El ke nel que compu a un bo de e ical, copia los esul ados di- ec amen e en la co espondien e posic´on del bu e con iguo co espondien e (ma shalling). Es necesa io p og ama un ke nel pa a ealiza la ope aci´on in- e sa (unma shalling) con los da os de los halos e i- cales ecibidos desde p ocesos emo os, colocando los da os en sus posiciones co espondien es en la ma iz de da os. La segunda se basa basada en u iliza las unciones de copia de ma ices bidimensionales inclu´ıdas en el API de CUDA. En conc e o, la unci´on cudaMem- cpy2D pe mi e a a ´es de sus pa ´ame os indica el sal o que hay en e los elemen os discon iguos, ans- i iendo una subma iz con menos columnas de las que es ´an ese adas en la ma iz comple a. En el es- udio expe imen al compa amos ambas opciones en ´e minos de endimien o. B. Memo ia pinned La as e encia de da os del hos al disposi i o y ice e sa, implica un cos e de iempo pa a ealiza la comunicaci´on a a ´es de los buses PCI en los que se alojan las a je as GPU. Puede a ia seg´un c´omo se ese e la memo ia en el hos . Po de ec o, la memo- ia que ese amos en el hos con las u inas alloc de C, es memo ia paginada iden i icada po una Um- mmm... es as segu o de odo es o? No es s´olo que iene que hace los accesos esol iendo las di ecciones de memo ia i ual a ´ısica? . Es o p o oca que pa a ealiza una ansmisi´on de da os desde el disposi- i o al hos sea necesa ia la ese a de un bloque de memo ia no paginada, seguido de una copia en el hos desde la memo ia ´ısica a la eci´en ese ada, sum´andole el cos e de la ansmisi´on, espe a y lib- e acion de memo ia de odo es e p oceso. En cam- bio, si ese amos memo ia pinned (memo ia no pag- inada), la comunicacion es mas ´apida debido a que la GPU no necesi a hace la esoluci´on de la di eccion de memo ia, ya que la memo ia pinned u iliza una di- ecci´on de memo ia ´ısica (RAM), y se puede ealiza di ec amen e la ansmisi´on de da os sin necesidad del sob ecos e en iempo mencionado an e io men e. C. A inidad en e CPU y GPU O o ac o impo an e pa a op imiza las ans- e encias de memo ia es de ini la a inidad del p o- ceso/ h ead que es ´a ejecu ando el c´odigo del hos , de o ma que se ejecu e en los co es que es ´an en el mismo nodo NUMA que es ´a conec ado al bus PCI donde es ´a conec ada Es o pe mi e e i a el sob e- cos e de la comunicaci´on en e dis in os nodos de la CPU. No con ola es a a inidad nos puede p oduci esul ados es oc´as icamen e peo es en los iempos de ans e encia du an e una se ie de p uebas del mismo p og ama. El co e de la CPU que se enca ga de p ocesa el p og ama po de ec o lo escoge el sis ema ope a i o en unci´on de pa ´ame os no ienen en cuen a el uso de los disposi i os ex e nos. Adem´as, si no se de ine la a inidad, el sis ema ope a i o ambi´en puede decidi mig a el p oceso a o os co es. En el ejemplo de la igu a 4 hay una buena a inidad en e la CPU0 y las GPU0 y GPU1, ya que se encuen an en un mismo nodo NUMA y la memo ia ges ionada po sus co es es a ´a no malmen e ese ada en los ban- cos de memo ia m´as p ´oximos a ella. Po an o, las comunicaciones en e esa memo ia y el bus PCIe de esa CPU es m´as e icien e que si iene que asmi i se a a ´es de la conexi´on en e la CPU0 y la CPU1 pa a di igi se a las GPUs que se encuen an en el o o nodo NUMA. Fig. 4. A inidad en e CPU y GPU Hyd a Chime a Phoenix Thunde bi d CPU Xeon E5-2690 3 Xeon E5-2620 2 Co e2 Quad Q6600 Co e i5330 Co es 12 24 4 4 Clock 1.9 GHz 2.1 GHz 2.4 GHz 3 GHz Memo y 64 GB 32 GB 6 GB 8 GB Num.GPUs 4 1 1 1 Model GTX Ti an Black GTX Ti an Black Tesla K40c Tesla K40c Co es 2880 2880 2880 2880 Clock 980 MHz 980 MHz 745 MHz 745 MHz Memo y 6 GB 6GB 12 Gb 12 Gb TABLA I Cl´ us e T asgo: M´ aquinas y ca ac e ´ ıs icas IV. Expe imen aci´ on En es a secci´on p esen amos los esul ados de un es udio expe imen al esalizado pa a comp oba el impac o eala i o de las di e en es opciones y op i- mizaciones conside adas, y pa a e i ica el ni el de escalabilidad conseguido po la soluci´on. Es e es ´a en g an medida asociado a la capacidad de solapa la compu aci´on y comunicaci´on de o ma sos enida a lo la go de las i e aciones del p og ama. Hemos escogido una implemen aci´on de Jacobi2D, con una pa ici´on de da os cl´asica y sencilla: blo- ques bidimensionales homogeneos. Es e p oblema es uno de los s encils m´as sencillos en dos dimensiones, con una ca ga compu acional muy baja po cada ele- men o conside ado, con un esquema de comunicaci´on en e ecinos con has a cua o ecinos po p oceso. Es a disposici´on, con baja ca ga compu acional po a ea, hace muy isible el impac o de las comuni- caciones. Adem´as, disponemos de implemen aciones sencillas y e icien es de e e encia. Po si al inal me is e el que has pues o en la igu a. Bo a la ase que no sea cie a. // Hemos escogido como ke nel de base la implemen aci´on con enida en el ejemplo dis ibuido con el Toolki de CUDA en su e si´on 8.0. // Hemos escogido como ke nel de base una a- ducci´on di ec a de la unci´on de ac ualizaci´on de un elemen o como se mues a en la igu a ??. (Y pones el c´odigo del ke nel b´asico, que son pocas l´ıneas). U ilizamos como e si´on de e e encia una im- plemen aci´on de la pa e de comunicaci´on en MPI basada en el ejemplo p esen ado en [1]. En nues- a e si´on de e e encia la pa e compu acional se ejecu a en la GPU, lanzado un ´unico ke nel pa a cal- cula oda la subma iz asignada al p oceso en cada i e aci´on. Se ealizan las ans e encias de memo ia pa a mo e los da os de los bo des al hos , o los ha- los ecibidos al disposi i o, de o ma s´ınc ona con la unci´on cudaMemcpy. Compa amos los esul ados de endimien o con la e si´on modi icada seg´un las es- peci icaciones comen adas en la secci´on II, p obando las di e nen es opciones y op imizaciones de alladas en la secci´on III. En es e abajo nos cen amos en la escalabilidad de es e p oblema en cl´us e s de GPUs dis ibuidas. u ilizan p ´ac icamen e oda la memo ia global de los disposi i os disponibles, maximizando el ama˜no global del g id compu ado. El ipo base de los a ays es loa , pa a pode ap o echa el m´aximo n´ume o de unidades uncionales en las GPUs. En la ex- pe imen aci´on u ilizamos ama˜nos de p oblema que de i an en una ocupaci´on de la memo ia global de los disposi i os que a desde ap oximadamen e de un 5% a un 75% pa a las GPUs de e e encia Ti an Black. A. Pla a o mas Los expe imen os se han ealizado en dos cl´us e s mul i-GPU di e en es, ambos ges ionados con un sis- ema de colas Slu m. El p ime o es ´a ges ionado di- ec amen e po el g upo de in es igaci´on T asgo de la Uni e sidad de Valladolid. Se a a de un cl´us e he e ogeneos con nodos de di e en es capacidades. El la abla I se esumen las m´aquinas que compo- nen el cl´us e y sus ca ac e ´ıs icas, incluyendo las de las GPUs ins aladas. La ecnolog´ıa de ed de in- e conexi´on es E he ne 1Gb en e los dos p ime os nodos (Hyd a y Chime a), y de 1Mb con los o os dos (Phoenix y Thunde bi d). La segunda pla a o ma es el cl´us e mul i-GPUs del Cen o Ex eme˜no de Tecnolog´ıas A anzadas (CETA-Ciema ). Cada nodo iene dos CPUs In el Xeon E5-2620, con seis co es a 2.0 GHz, wi h 32 GB de RAM, y dos GPUs ...... modelo y ca ac e - is icas...... La ecnolog´ıa de ed de in e conexi´on es In iniband QDR (40Gb/s) y FDR (56Gb/s). B. Comunicaci´on s´ınc ona y as´ınc ona Una de las op imizaciones com m´as impac o en compu os i e a i os que necesi an as e i da os en- e hos y disposi i os en cada i e aci´on, es el so- lapamien o del compu o con la comunicaci´on. Es o es ac ible en el caso de nues a compu aci´on s en- cil, que no p esen a dependencias en e i e aciones si se man iene una copia con los da os de la i e aci´on an e io . Po an o es posible ealiza las ans e en- cias de da os y comunicaciones en e nodos de o ma as´ınc ona. En la abla II p esen amos el desglose de iem- pos de ejecuci´on dedicados a compu aci´on y a ans- e encias de memo ia y/o comunicaci´on en e nodos pa a di e en es ama˜nos de p oblema en la e si´on de e e encia. Los iempos es ´an medidos en En que m´aquina o combinaci´on de m´aquinas? Es o es eje- cu ando en co es o en GPUs??. Obse amos que los iempos de compu o y comunicaci´on son muy simila es, c eciendo ambos p opo cionalmen e con el TABLA II Tabla con los iempos de compu aci´ on y comunicaci´ on, con 50 i e aciones (en segundos) Tama˜no ma iz C´ompu o Comumicaci´on 6000x6000 0.3104 0.3263 20000x20000 3.0809 3.1852 25000x25000 5.0184 5.2029 TABLA III Tabla con modelo s´ ınc ono y a´ sinc ono, con 50 i e aciones En que m´ aquina o combinaci´ on de m´ aquinas?? (en segundos) ama˜no ma iz s´ınc ono As´ınc ono 6000x6000 0.6367 0.3302 20000x20000 6.2661 3.2516 25000x25000 10.2213 5.3099 ama˜no de p oblema. En la abla III mos amos los iempos de ejecuci´on o ales de la e si´on de e e encia y de la e si´on con ans e encias y comunicaciones as´ınc onas. En la e si´on s´ınc ona de e e encia, los iempos de com- pu o y comunicaci´on se acumulan en el o al, ya que las ope aciones se secuencializan. En la e si´on as´ınc ona, la mayo pa e de los iempos de comuni- caci´on se solapan con la compu aci´on, obse andose iempos o ales que ienen apenas un inc emen o de en e un 1% y un 2% sob e los iempos de comuni- caci´on de la e si´on de e e encia. C. Bu e s s. cudaMemcpy2D En es e abajo compa amos dos ap oximaciones pa a ans e i los da os de los bo des e icales en- e el hos y el disposi i o y ice e sa. Una con p o- g amaci´on manual, ejecu ando ope aciones de ma - shalling y unma shalling en bu e s con da os con- iguos. La o a u ilizando di ec amen e la unci´on cudaMemcpy2D con los pa ´ame os adecuados. En la abla IV mos amos una compa aci´on de iempos de ejecuci´on o ales en e la cada una de las dos e - siones. Obse amos que en e las dos ´ecnicas apli- cadas, el uso de cudaMemcpy2D nos da un lige a mejo a sob e el uso de bu e s, ya que es e m´e odo iene di e en es op imicaciones a ni el del d i e de CUDA. Adem´as es mucho m´as sencillo que p og a- TABLA IV Tabla con iempo de comunicaci´ on de bu e s s. cudaMempcy2D (en segundos) ama˜no ma iz bu e cudaMemcpy2D 6000x6000 0.3268 0.3255 20000x20000 3.1919 3.1484 25000x25000 5.1980 5.1371 ma manualmen e la ope aciones sob e los bu e s, y al se pa e del API de CUDA es m´as po able con a qui ec u as u u as. D. Escalabilidad de comunicaciones MPI En es a pa e del abajo expe imen al es udiamos la escalabilidad de las comunicaciones as´ınc onas en MPI, ejecu ando la pa e compu acional en los co es de la CPU. Ejecu amos dos ipos de expe imen os. El p ime o pa a medi la escabilidad d´ebil (weak scal- ing), aumen ando el ama˜no del p oblema p opo - cionalmen e al n´ume o de disposi i os in oluc ados. As´ı podemos obse a el endimien o de la aplicacion cuando se amplia el nume o de nodos dis ibuidos, con una ca ga cons an e pa a cada nodo. El segundo ipo de expe imen os es ´a o ien ado a medi la escal- abilidad pa a un ama˜no de p oblema ijo (s ong scaling). As´ı podemos obse a el compo amien o de la apliaci´on con una ca ga o al cons an e, mien- as inc emen amos el n´ume o de nodos usados. Los iempos de compu o se educen, ya que al habe m´as p ocesos in oluc ados le oca una pa e m´as peque˜na de las es uc u as de da os a cada uno. El impac o y e ec os de los iempos de comunicaci´on se hacen m´as e iden es, Tama˜nos de p oblema pa a es e caso??? En el caso del cl´us e T asgo, empezamos usando co es de una ´unica m´aquina, inc emen ando el nu- me o de co es de cua o en cua o, has a comple a su capacidad. Se con inuan in oduciendo co es de cua o en cua o de las siguien es m´aquinas, siemp e comple ´ando su capacidad an es de a˜nadi la sigu- ien e. El o den de in oducci´on de las m´aquinas se ha escogido po su can idad de co es, quedando es- ablecido el o den siguien e: chime a, hyd a, phoenix y hunde bi d. En el caso inal se es ´an u ilizando los 24 co es de chime a, los 12 de hyd a, los 4 de phoenix y de hunde bi d. En el caso del cl´us e CETA p o- cedemos de mane a simila ellenando m´aquinas con bloques de 4 p ocesos. Al se m´aquinas homogeneas el o den en que se a˜naden los nodos no es ele an e. Tan o en la igu a 5, como en la igu a 6, emos que la aplicaci´on inc emen a el iempo cuando in- c emen amos los p ocesos en un mismo nodo, de- bido a que con pocos p ocesos, nos encon amos en una si uacion de al a a inidad, donde los co es que se es an ejecu ando se encuen an en un mismo nodo NUMA (misma CPU) y los mo imien os de memo- ia y las comunicaciones en e los p ocesos son muy ´apidas. Al inc emen a el n´ume o de p ocesos, se in oducen p ocesos en o as pa es de la je a qu´ıa, inc emen ´andose los cos es de comunicaci´on. En el momen o en que se empiezan a in oduci p ocesos en o os nodos de la ed, los cos es de comunicaci´on son ya simila es. Po lo que la escalabilidad d´ebil se man iene es able, con un lige ´ısimo inc emen o hacia el inal en el caso del cl´us e T asgo al in oduci se encnolog´ıa de ed m´as limi ada (E he ne 1Mb). En la segunda pa e del es udio comp obamos la escalabilidad pa a un ama˜no de p oblema ijo (s ong scaling). En la igu a 7 se mues an los esul- 0 5 10 15 20 4 8 12 16 20 24 24+4 24+8 24+12 (36)+4 (40)+4 Tiempo (Segundos) p ocesos o ales Fig. 5. Cl´us e T asgo: Weak scaling con n´ume o de p oce- sado es 0 10 20 30 40 50 1,6 1,12 2,18 2,24 3,30 3,36 4,42 4,48 5,54 5,60 6,66 6,72 Tiempo (Segundos) (nodos, p ocesos) Fig. 6. Cl´us e CETA: Weak scaling con n´ume o de p oce- sado es ados pa a el cl´us e T asgo. Podemos e la mejo a cla a en los iempos del p og ama a medida que in- c emen amos el n´ume o de nodos que usamos, ya que la ca ga pe nodo que ealizamos es cada ez meno . Incluso con una ecnolog´ıa de ed basad en bus, como es E he ne , la can idad y ecuencia de las comunicaciones de es e ipo de aplicaciones no llega a se su icien e Las i egula ides de la cu a en los p ime os pun os se debe de nue o al cambio b usco en los cos es de comunicaci´on den o de la je a qu´ıa NUMA del mismo nodo. E. Mul i-GPU dis ibuidas En es a pa e del es udio nos cen amos en es u- dia la escalabilidad, an o d´ebil como ue e (weak, s ong scaling) cuando el c´ompu o se ealiza con los disposi i os GPU, educiendose los iempos de c´ompu o e inc men ´andose los de comunicaci´on de- bido a la in oducci´on de las ans e encias en e la je a qu´ıa de memo ia del disposi i o y el hos . Tama˜nos de p oblema pa a es e caso??? En la igu a 8 mos amos los esul ados de escala- bilidad d´ebil pa a el cl´us e CETA. Podemos obse - a que apa ecen m´as i egula idades debidas a los e ec os es oc´as icos de la ed. Sin emba go, la es- calabilidad se man iene g acias a la capacidad de la ecnolog´ıa de ed de es e cl´us e (In iniband). 5 10 15 20 25 30 35 40 45 50 4 8 12 16 20 24 24+4 24+8 24+12 (36)+4 (40)+4 Tiempo (Segundos) p ocesos o ales Fig. 7. Cl´us e T asgo: S ong scaling con n´ume o de p oce- sado es 0 5 10 15 20 1,1 1,2 2,3 2,4 3,5 3,6 4,7 4,8 5,9 5,10 Tiempo (Segundos) (nodos, p ocesos) Fig. 8. Cl´us e CETA: Weak scaling con GPUs en nodos dis ibuidos En la igu a 9 mos amos los esul ados de escal- abilidad ue e en el clus e T asgo. Las p ime as GPUs que se in oducen son las de hyd a (4 en el mismo nodo). Luego se an a˜nadiendo GPUs de los o os nodos, in oduciendo al inal las de los no- dos conec ados con la ecnolog´ıa de ed m´as limi- ada. En es e caso el iempo de c´ompu o se e- duce d ´as icamen e g acias al uso de las GPUs, so- lap´andose po comple o con el iempo de comuni- caci´on que domina el iempo o al. Los iempos en los ´ul imos pun os de la g ´a ica se en a ec ados po la capacidad de la ecnolog´ıa de ed m´as limi ada de los ´ul imos nodos in oducidos (E he ne 1Mb). Conclusiones En es e abajo se han epasado y e isado ´ecnicas de implemen aci´on e icien e de compu aciones de ipo s encil pa a la ocul aci´on de las la encias de comunicaci´on en e hos y GPU, y en e nodos, en un cl´us e de GPUs dis ibuido. Se han e isado di e en es op imizaciones y usos del API de CUDA. Los esul ados mues an que la aplicaci´on de ´ecnicas sencillas de p og amaci´on puede consegui minimiza el impac o de las comunicaciones incluso en el caso de u iliza ecnolog´ıas de ed de medio compa ido de bajo cos e pa a g ados de escalabilidad aco ados. Los esul ados son p ome edo es en cuan o a la posi- 2 4 6 8 10 12 14 1 2 3 4 4,1 4,1,1 4,1,1,1 4,1,1,1,1 Tiempo (Segundos) p ocesos po nodo Fig. 9. Cl´us e T asgo: S ong scaling con GPUs en nodos dis ibuidos bilidad de esol e e icien emen e p oblemas basados en s encils en g andes cl´us e s con nodos do ados con acele ado es ha dwa e y ecnolg´ıa de ed escalable. El abajo u u o incluye el es udio de aplicaciones s encil m´as complejas, o as ´ecnicas no con em- pladas en el es udio, ajus a la pa ici´on de da os a la he e ogenidad de los disposi i os y nodos, y la in- eg aci´on de las ´ecnicas es udiadas en abs acciones de p og amaci´on de sis emas he e og´eneos de m´as al o ni el. Re e encias [1] Einwg Lusk William G opp, Using MPI. Po able Pa allel P og amming wi h he Message-Passing In e ace, MIT P ess, 3 d edi ion, 2015. [2] B ian Hamil on, C aig J. Webb, Alan G ay, and S e an Bilbao, “La ge s encil ope a ions o GPU-based 3-D acous ics simula ions,” in 18 h In . Con e ence on Dig- i al Audio E ec s (DAFx-15), 2015. [3] Paulius Micike icius, “3d ini e di e ence compu a ion on gpus using cuda,” in GPGPU2. Ma ch 2009, ACM P ess. [4] Mohammed Sou ou i, Johannes Langgu h, Filippo Spiga, Sco B. Baden, and Xing Cai, “CPU+GPU p og am- ming o s encil compu a ions o esou ce-e icien use o GPU clus e s,” in IEEE 18 h In e na ional Con e ence on Compu a ional Science and Enginee ing. 2015, IEEE. [5] Kaushik Da a, Samuel Williams, Vasily Volko 2, Jona han Ca e , Leonid Olike , John Shal , and Ka he - ine Yelick, “Au o- uning he 27-poin s encil o mul i- co e,” in In e na ional Wo kshop on Au oma ic Pe o - mance Tuning (iWAPT), 2009. [6] Da id Ki k and Wen mei Hwu, P og amming Massi ely Pa allel P ocesso s, Mo gan Kau man, 2nd edi ion, 2013.