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.