Paralelización mediante procesadores gráficos (GPU) del cálculo de la dinámica de electrones en plasma confinado magnéticamente

Más e Uni e si a io en Compu ación Pa alela y Dis ibuida Depa amen o de Sis emas In o má icos y Compu ación Pa alelización median e p ocesado es g á icos (GPU) del cálculo de la dinámica de elec ones en plasma con inado magné icamen e TRABAJO FINAL DE MÁSTER Au o : Tomás Na a o Cosme Di ec o : José En ique Román Mol ó Sep iemb e de 2015 . Abs ac CUDA pa alleliza ion o a p og am, o iginally w i en in Fo an wi h MPI, o he calcula ion o he dynamic o elec ons, aking pa icula a en ion o he o ma ion o unaway elec ons in a mangen i- cally con ined plasma, unde he e ec o an elec ic ield, using he Lange in’s app oxima e calcula ion me hod. The inal implemen a ion will e ain he mos basic pa sou ce code w i en in Fo an, a lan- guage wi h a la ge base o use in he ield o Physics, o show i s compa ibili y wi h he newes NVIDIA’s G aphical P ocesso s Uni s ia CUDA. The implemen a ion is p og essi e, showing successi e echnical imp o emen s ha can be applied o such codes, analyzing he ob ained speedups. Resumen Pa alelización en CUDA de un p og ama, o iginalmen e esc i o en Fo an usando MPI, pa a el cálculo de la dinámica de los elec ones, con especial a ención a la apa ición de elec ones ugi i os ( unaway), den o de un plasma con inado magné icamen e y que se mue en po e ec o de un campo eléc ico, u ilizando el mé odo de cálculo ap oximado de Lange in. En la pa alelización se man end á la pa e básica de Fo an, lenguaje con una g an base de uso en el campo de la Física, pa a mos a su compa ibilidad de uso con los acele ado es g á icos de NVIDIA usando CUDA. La implemen ación es p og esi a, mos ando las sucesi as écnicas y mejo as que pueden aplica se a es e ipo de códigos, analizando las acele aciones de cálculo ob enidas. Palab as cla e: CUDA, Fo an, GPU, GPGPU, compu ación numé ica, código cien í ico, HPC, plasma, Lange in, elec ones unaway i ÍNDICE ABREVIADO Índice ab e iado · iii Índice gene al · Índice de igu as · ix Índice de ablas · ix Índice de lis ados de código · x 1 In oducción · 1 2 Conocimien os p e ios · 3 3 Ecuaciones de Lange in pa a desc ibi un plasma de pa ículas y la gene ación de elec ones unaway · 13 4 Implemen ación de la pa alelización con CUDA · 17 5 Resul ados · 43 6 T abajo u u o · 71 7 Conclusiones · 73 Glosa io · 75 Siglas · 77 Bibliog a ía · 79 iii ÍNDICE GENERAL Índice ab e iado iii Índice gene al Índice de igu as ix Índice de ablas ix Índice de lis ados de código x 1 In oducción 1 2 Conocimien os p e ios 3 2.1 La simulación compu acional en la Física . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 4 2.2 Ca ac e ís icas de los códigos cien í icos . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 5 2.3 GPGPU .................................................. 7 2.4 CUDA ................................................... 9 2.4.1 Lenguaje CUDA C . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 9 2.4.2 Tipos de código en CUDA C . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 9 2.4.3 Compilación de aplicaciones CUDA . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 10 3 Ecuaciones de Lange in pa a desc ibi un plasma de pa ículas y la gene ación de elec- ones unaway 13 3.1 El plasma y su aplicación p ác ica: eacción con olada de usión . . . . . . . . . . . . . . . . 13 3.2 Con inamien o a i icial del plasma . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 14 3.3 Física de pa ículas: la ecuación de Lange in . . . . . . . . . . . . . . . . . . . . . . . . . . . . 15 4 Implemen ación de la pa alelización con CUDA 17 4.1 Selección del lenguaje auxilia de p og amación . . . . . . . . . . . . . . . . . . . . . . . . . . 17 4.2 U ilización de w appe s en Cpa a se in ocados po Fo an . . . . . . . . . . . . . . . . . . 18 4.2.1 Respe o de la nomencla u a (name mangling) de los símbolos según compilado . 19 4.2.2 Esquema gene al de llamadas a CUDA desde Fo an a a és de w appe s C. . . . 21 4.3 Di e encias p og amá icas impo an es en e Fo an yCyCUDA-C . . . . . . . . . . . . . . 23 4.4 Implemen ación en CUDA del código de Lange in o iginalmen e en Fo an . . . . . . . . 24 4.4.1 Aplanamien o de sub u inas po inclusión en un ke nel aglu inado . . . . . . . . . 24 4.4.2 Minimización de la ans e encia de da os en e CPU yGPU . . . . . . . . . . . . . . 25 4.4.3 U ilización adecuada de los gene ado es de núme os alea o ios de CUDA . . . . . . 26 4.4.4 Mejo as en el código: u ilización de unciones es ánda en e a mé odos numé i- cos pesados . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 27 4.4.5 Algo i mos de educción en la GPU . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 28 4.4.5.1 Reducción global i e a i a . . . . . . . . . . . . . . . . . . . . . . . . . . . . 29 4.4.5.2 Reducción median e el uso de ope aciones a ómicas . . . . . . . . . . . . 29 4.4.6 Mejo a en el algo i mo de clasi icación de elocidades . . . . . . . . . . . . . . . . . . 30 4.4.7 De ección de e o es (bugs) en el código . . . . . . . . . . . . . . . . . . . . . . . . . . 31 4.5 Desacople de compu ación y comunicación usando OpenMP . . . . . . . . . . . . . . . . . . 33 4.5.1 Desacople usando OpenMP con hilo en espe a ac i a . . . . . . . . . . . . . . . . . . 34 4.5.2 Implemen ación inal con ce ojos de OpenMP (omp_locks) . . . . . . . . . . . . . 36 4.5.2.1 Conside aciones con la adición de OpenMP en el código Fo an . . . . . 40 4.6 Valo ación cuali a i a del p oceso de pa alelización con CUDA . . . . . . . . . . . . . . . . . 40 5 Resul ados 43 5.1 P uebas ealizadas . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 43 5.1.1 P ime a e apa: op imización del ke nel p incipal según los pa áme os p opios de su in ocación . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 44 5.1.2 Segunda e apa: medición de p es aciones de las e siones inales en un en o no HPC eal . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 45 5.2 En o no de ejecución . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 46 5.2.1 P ime a e apa: op imización de los pa áme os de in ocación del ke nel p incipal . 46 5.2.2 Segunda e apa: en o no HPC eal, MinoTau o . . . . . . . . . . . . . . . . . . . . . . . 47 5.3 Ejecución de las p uebas . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 47 5.3.1 Pa áme os óp imos de in ocación al ke nel p incipal ad ance . . . . . . . . . . . . 47 5.3.1.1 B e e desc ipción de la p ueba . . . . . . . . . . . . . . . . . . . . . . . . . . 47 5.3.1.2 Resul ados ob enidos . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 49 5.3.1.3 Fenómenos obse ados . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 49 5.3.1.4 Conclusión de la p ueba . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 51 5.3.2 Acele ación en un en o no de ejecución de al as p es aciones . . . . . . . . . . . . . 51 5.3.2.1 Ca ac e ización del compo amien o del p og ama o iginal . . . . . . . . 52 5.3.2.2 Ca ac e ización básica de la implemen ación consis en e en la pa aleli- zación con CUDA . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 54 5.3.2.2.1 P ueba de co ección . . . . . . . . . . . . . . . . . . . . . . . . . . 55 5.3.2.2.2 Acele ación y e iciencia espec o del código en CPU . . . . . . . 55 5.3.2.3 Ca ac e ización básica de la implemen ación con CUDA yOpenMP . . . . 57 5.3.2.4 P uebas masi as de escalabilidad de las implemen aciones con CUDA . . 58 5.3.2.4.1 Resul ados ob enidos . . . . . . . . . . . . . . . . . . . . . . . . . . 58 5.3.2.4.2 Escalabilidad de las implemen aciones . . . . . . . . . . . . . . . 61 5.3.2.4.3 Compa ación di ec a de las implemen aciones . . . . . . . . . . 64 5.3.2.4.4 Cálculo de la acele ación en una ca ga g ande . . . . . . . . . . . 66 5.4 Valo ación . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 67 5.4.1 Valo ación numé ica cuan i a i a . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 67 5.4.2 Valo ación cuali a i a de los esul ados . . . . . . . . . . . . . . . . . . . . . . . . . . . 67 5.4.2.1 Facilidad de p og amación . . . . . . . . . . . . . . . . . . . . . . . . . . . . 68 5.4.2.2 Necesidad de la e o mulación de los algo i mos . . . . . . . . . . . . . . . 68 5.4.2.3 Depu ación semán ica de las implemen aciones his ó icas . . . . . . . . . 69 5.4.2.4 Selección de códigos a pa aleliza . . . . . . . . . . . . . . . . . . . . . . . . 69 6 T abajo u u o 71 i 7 Conclusiones 73 Glosa io 75 Siglas 77 Bibliog a ía 79 ii Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e Finalmen e se expond án o as ecnologías de pa alelización, algunas de las cuales ambién han sido u ilizadas e incluso inco po adas. Así po ejemplo enemos MPI, una biblio eca de p og amación basada en el pa adigma de memo ia dis ibuida, median e la u ilización de paso de mensajes; y ambién OpenMP, que es á especialmen e pensado pa a memo ia compa ida. Úl imamen e han su gido o as posibilidades que in en an acili a la labo de la pa alelización abs ayéndose un poco más del ha dwa e exis en e po debajo, de en e los cuales des aca emos el caso de OpenACC. Po o a pa e, la pa e ísica del p oblema, po su especial ele ancia y en idad, se á a ada en el si- guien e Capí ulo 3 “Ecuaciones de Lange in pa a desc ibi un plasma de pa ículas y la gene ación de elec- ones unaway”, aunque sin en a en excesi o de alle cien í ico, sino más desc ip i a de su signi icado, así como de las p incipales ca ac e ís icas a ene en cuen a a la ho a de pa aleliza el algo i mo de simulación. 2.1 La simulación compu acional en la Física En el ámbi o de la ísica de pa iculas es muy habi ual el uso de la simulación compu acional. Es a u i- lización no se ealiza pa a su es udio di ec o, sino como mé odo de comp obación de las de hipó esis y nue os desa ollos eó icos. En es e campo de la ísica de pa ículas elemen ales, el mé odo de obse a- ción es á limi ado po la máxima p ecisión alcanzable, de e minada po el p incipio de inde e minación de Heisenbe g, y ese lími e ísico aco a ambién la ex acción de conocimien o a pa i de la expe imen- ación. A es o se une la complejidad de las ac uales eo ías, que deben compagina an o enómenos bien conocidos y eó icamen e bien desa ollados, como la g a i ación, la dinámica y el elec omagne ismo, con los e ec os cuán icos y la aplicación de la eo ía de la ela i idad, ya que es habi ual abaja con pa ículas que se mue en a al as elocidades (es o es, son de al a ene gía). Po odo ello, cuando se abaja a ni el de pa ículas ales como elec ones, neu ones, p o ones y pa ículas al a, las ecuaciones eó icas no pueden esol e se analí icamen e. El p ocedimien o habi ual es adap a las al caso pa icula en es udio, ealizando las opo unas simpli icaciones e ndo que no a ec en al esul ado a ob ene , y inalmen e, con el conjun o de ó mulas simpli icadas, y po lo an o ap oximadas, ealiza in ensos cálculos de simulación pa a comp o- ba si se pueden ep oduci los esul ados ob enidos expe imen al o empí icamen e. En caso de que así sea, ese conjun o de ó mulas ap oximadas puede some e se a un nue o p oceso de simulación que se ex ienda a casos más complejos que di ícil o muy one ósamen e pueden con as a se expe imen almen e. De es a o ma se ob iene in o mación aliosa que pe mi e acele a el p oceso ecnológico y de ingenie ía pa a la cons ucción de nue as ins alaciones de expe imen ación. Es e es el caso de la ísica de plasma, en donde pa ículas elemen ales de los á omos más lige os se hallan con inadas po el e ec o de ue es campos elec omagné icos a al as empe a u as, inonizadas, lib es en e sí, y p oduciéndose choques ocasionales de al a ene gía, que, en condiciones de iempo, empe a u a y densidad su icien e, pueden p o oca el inicio del p oceso de la usión nuclea . Desde el pun o de is a in o má ico, los p ocesos de simulación se ca ac e izan po se compu acional- men e in ensos, dado que se debe a a un núme o lo más g ande posible de pa ículas, y po que gene- almen e in e esa la dinámica del sis ema, es o es, cómo e oluciona desde un de e minado es ado inicial. Es o implica gene almen e que se debe calcula el es ado inicial, después supone un inc emen o de iem- po pequeño, pa a in en a linealiza lo más posible la dinámica del sis ema, es o es, que los e o es po las ap oximaciones ealizadas no se p opaguen en exceso, pe o no an pequeño como pa a supone una inmensa can idad de inc emen os “di e enciales” del iempo, y con ello, una g an can idad de nue os es a- dos a calcula has a alcanza el es ado inal buscado (gene almen e, e si alcanza un es ado es aciona io, o analiza y comp ende cómo se compo a si no. Po lo an o, las ecuaciones son complejas, con muchos cálculos po cada es ado, y muchos es ados a compu a . Es o puede complica se en unción de la compleji- dad del desa ollo eó ico u ilizado. De es a o ma, in e esa ob ene unas ecuaciones que calculen el es ado 4 Capí ulo 2. Conocimien os p e ios 2.2. Ca ac e ís icas de los códigos cien í icos de cada pa ícula de o ma independien e al compo amien o indi idual de cada una del es o, pues o que si exis e in e dependencia en e las pa ículas, el o den del núme o de ope aciones a compu a po cada es- ado pasa de se de O(N), siendo Nel núme o de pa ículas, a o den O(N·m), siendo mel núme o de pa iculas ecinas con las que in e ac úa, o incluso O(N2) si odas in e ac úan con odas en cada es ado. Es a al a necesidad de pode de cálculo implica que, si se u iliza un único p ocesado , el iempo de cálculo sea excesi amen e ele ado, lo que conlle a p oblemas po : inc emen o del e aso en e di e en es simulaciones, con la consiguien e pa alización de la alidación de los desa ollos eó icos ap oxima i os ealizados; mayo iesgo de inalización ab up a de la simulación, an o po enómenos ex e nos (co es de luz) como in e nos (sob ecalen amien o del nodo, e o de so wa e,...). La solución que se ha enido aplicando en los úl imos 20 años ha sido la pa alelización de los p ocesos, bien sea median e OpenMP (en sis emas mul ip ocesado es de memo ia compa ida), o median e MPI (paso de mensajes en e nodos con memo ia dis ibuida). Todo ello es posible u ilizando el lenguaje Fo an, que iene su p incipal nicho de uso den o del e eno de cómpu o cien í ico, pues o que: es á pensado pa a ansc ibi las ó mulas ma- emá icas; pe mi e la p og amación es uc u ada, gene almen e idónea pa a implemen a los habi uales algo i mos de cálculo cien í ico; pe mi e un uso muy simpli icado de las ope aciones ec o iales y ma i- ciales, muy habi uales en ciencia; es á o ien ado al cómpu o, con compilado es muy op imizados al e ec o; dispone de lib e ías especí icas de cálculo numé ico, algeb aico y cien í ico muy op imizadas, obus as y con as adas; iene una muy buena implemen ación y sopo e en los dos his ó icamen e p incipales pa a- digmas de p og amación pa alela, OpenMP yMPI, dado que sus es ánda es se desa ollan e implemen an eniendo muy en cuen a las pa icula idades y en ajas de Fo an. Es e úl imo aspec o es undamen al pa a con inua la he encia de muchos p og amas y biblio ecas de cálculo ma emá ico y cien í ico desde los años 60 has a la ac ualidad, lo cual iene aspec os muy posi i os (p og amas con as ados, e icien es y obus os), pe o que su ilmen e esconde algunos incon enien es que con iene ene en cuen a y analiza , como se e eló duan e la ase de ansc ipción y adap ación al nue o pa adigma SIMD de CUDA, al como se e á en es e abajo. 2.2 Ca ac e ís icas de los códigos cien í icos Po código cien í ico en ende emos en es e abajo aquel p og ama o biblio eca que implemen e di ec a- men e un de e minado algo i mo de compu ación cien í ica. Es e algo i mo de compu ación cien í ica se á gene almen e un conjun o de ecuaciones in e dependien es que pe mi en esol e un de e minado p o- blema ísico. Gene almen e, es e conjun o de ó mulas p o iene del desa ollo eó ico en una de e minada á ea cien í ica, y más especí icamen e, nos cen a emos en aquellas que si en pa a la simulación de sis e- mas ísicos. Es os p oblemas ienen la ca ac e ís ica común de que son lo su icien emen e complejos como pa a que el conjun o de ó mulas de i adas es ic amen e de la eo ía sean demasiado complejas pa a su esolu- ción analí ica. Es a imposibilidad de esolución analí ica impide que se pueda analiza el compo amien o de es os sis emas en di e en es si uaciones p ác icas. Al obje o de pode es udia los, se hace necesa io eali- za de e minadas suposiciones que pe mi en alige a la complejidad del conjun o de ecuaciones, median- e la inclusión de ap oximaciones azonables pa a el caso conc e o en es udio. T as es as ap oximaciones, es muy ecuen e que se ob enga o o conjun o de ecuaciones más simple, pe o aún así i esolubles o de di ícil in e p e ación. Pa a esol e las se ecu en a mé odos numé icos median e la implemen ación de p og amas in o má icos especí icos pa a ese p oblema. Es os códigos cien í icos se ejecu an en múl iples ins ancias, pues o que es ecuen e que es as unciones engan é minos es ocás icos, es o es, que pa a cap u a la a iabilidad de las condiciones ex e nas o de la complejidad in e na inna a del sis ema, y an e la imposibilidad o g an di icul ad de exp esión eó ica es ici a, en el cue po eó ico se in oducen é minos 5 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e de a iabilidad alea o ia (denominados é minos es ocás icos), que pe mi en do a al sis ema de un cie o compo amien o y adap ación pseudoalea o ia necesa io. Es as múl iples ejecuciones se denominan simu- laciones, y pe mi en ex ae nue o conocimien o del sis ema ísico en es udio del análisis de sus esul ados. His ó icamen e, es os códigos cien í icos han sido implemen ados po los p opios cien í icos eó icos enca gados de desa olla las ó mulas básicas y sus de i aciones ap oximadas, y p esen an las siguien es ca ac e ís icas: • Lenguaje de p og amación: p e e encia po lenguajes como Fo an o en o nos como Ma Lab, sob e odo en compa ación con C, con sin axis más sin é icas en las ins ucciones de cálculo ma emá ico, especialmen e al pe mi i abaja con ec o es y ma ices sin necesidad de gene a bucles que los eco an. • Es uc u ación: pensada pa a la implemen ación ápida de algo i mos ma emá icos. Los módulos pe mi en de ini biblio ecas de unciones, y la e i ación de bucles explíci os pa a ope aciones ma i- ciales pe mi e una mayo cla idad del código, pe mi iendo al p og amado a cen a se en el algo i mo y no en el código. • Es uc u a: Los códigos cien í icos son p incipalmen e de compu ación numé ica. Po lo an o suelen es a cla amen e es uc u ados en en ada o lec u a de da os, p epa ación del cálculo, ejecución del cálculo (pa e compu acionalmen e más in ensa) y ecopilación de esul ados pa a mos a o gua - da . Po es e hecho u ilizan biblio ecas pa a la en ada y salida de da os ( al como silo, con o ma os especiales como HDF5, y p ecisan de un lenguaje y biblio ecas muy comple as y e icien es compu- acionalmen e. Es á ue emen e basado en bibli ecas que esuel en los p oblemas compu acionales más gene ales. • Independiencia en e cálculo y isualización: En elación con lo an e io , el código cien í ico es á pen- sado pa a se e isado y eesc i o, y po ello se hace independien e de p ocesos pos e io es como el de la isualización. Los p og amas gene an iche os en o ma os es anda izados pa a su pos e io a- amien o po o os códigos, o po paque es de isualización de da os, como VisI , o el p opio Ma Lab • Legibilidad: Pensados pa a i ac ualizándolos con o me e olucione la eo ía que gene a el conjun o de ecuaciones a esol e . Po lo an o, in e esa que el lenguaje de p og amación e i e a e ac os p o- pios del lenguaje que enmasca e o di icul e la exp esión y comp ensión del algo i mo implemen ado. • Fue e in luencia de los algo i mos: lo cual puede se un p oblema. Muchas eces, la ansc ipción di ec a de la ó mula a código, hace que se incu a en la u ilización de algo i mos compu acional- men e ine icien es cuando son implemen ados po un cien í ico no expe o y que no conside e el cos e compu acional de su implemen ación. • Necesidad de ipos especí icos de a iables: ejemplo cla o es la necesidad de abaja con núme os complejos, y que la biblio ecas ma emá icas y compu acionales ambién lo hagan. También el e- conocimien o de núme os especiales (i, e) y del co ec o a amien o de alo es ex emos ( alo es mínimos y máximos p ede inidos, gene ación y abajo con a iables que hayan excedido el in e a- lo de inido (NaN, No a Numbe ), pa a lo cual es muy in e esan e que se sigan el es ánda IEEE 754 (Ins i u e o Elec ical and Elec onics Enginee s S anda d o Floa ing-Poin A i hme ic). 6 Capí ulo 2. Conocimien os p e ios 2.3. GPGPU 2.3 GPGPU GPGPU [5], Gene al-Pu pose compu ing on G aphics P ocessing Uni s, consis e en la ealización de cómpu- os de p opósi o gene al u ilizando pa a ello p ocesado es g á icos, GPU, en luga de los habi uales p oce- sado es gené icos, CPU. Debido a la economía de escala y al exi oso me cado de a je as g á icas impulsado po el sec o de los juegos en el PC y consolas, los p ocesado es g á icos que se u ilizan en esas a je as g á icas han e oluciona- do desde se unos cop ocesado es que simplemen e ep esen aban ca ac e es o g á icos de mapa de bi s, a se complejos p ocesado es capaces de ealiza no an simples cálculos ma emá icos y ec o iales sob e ca- da píxel pa a pode do a al conjun o de la imagen de un mucho mayo ealismo, a la ez que p opo ciona un al o e esco en la imagen pa a sua iza las ansiciones en e cada ins an ánea [11]. De es a o ma, el pode compu acional de las GPU se ha mul iplicado mucho más en es os úl imos años que el de los p ocesado es de p opósi o gene al, CPU. Una a je a g á ica de al a gama puede ene pe - ec amen e más de 2800 p ocesado es indi iduales odos ellos abajando de o ma (casi) independien e en pa alelo, y con capacidad de comunica se a a és de ápidas memo ias DDR5, incluso de a ios GB de capacidad, ubicadas en la misma a je a [9]. No obs an e, su conjun o de ins ucciones es mucho más especí ico, o ien ado a la ealización de cálculos g á icos, y se esien e de la al a de ins ucciones op imi- zadas muy comunes en la p og amación es ánda no g á ica (sal os y compa ación, con inua necesidad de as asa g an can idad de in o mación de la CPU a la GPU y ice e sa, po ejemplo). Tampoco ienen la uni e sal lexibilidad de las CPU pues iene di e en es zonas de memo ia, cada una con usos especí icos y no in e cambiables, o la coo dinación en e p ocesado es den o de un hilo de ejecución. O a ca ac e ís- ica que penalizaba mucho las an e io es a je as g á icas e a la inexis encia de ins ucciones especí icas pa a el lujo de con ol, oda ez que desde el pun o de is a me amen e g á ico, los p og amas siemp e han sido p e e en emen e secuenciales. En la ac ualidad, los p ocesado es g á icos ya inco po an ins ucciones pa a el con ol de lujo, pe o las ami icaciones (i - hen-else) son muy cos osas, sob e odo si en cada hilo el esul ado es di e en e, lo que p o oca g a es penalizaciones en iempo, pa a consegui una adecuada sinc onización en la ejecución del p og ama, po lo que con iene ene lo muy en cuen a pa a ob ene bue- nos esul ados [6]. Es as es icciones se han sua izado con la llegada de las nue as a qui ec u as Keple yMaxwell, pe o aún así son aspec os de uncionamien o que con iene segui eniendo en conside ación pa a la op imización de la ejecución del código en las GPU. Aún así, a mediados de los 2000 empeza on los expe imen os de ealiza cómpu os ma emá icos pesa- dos usando ó denes pu amen e g á icas, p incipalmen e basadas en el es ánda g á ico OpenGL. A pesa de la di icul ad de p og amación o ecida po es a ía, sí que se comp obó empí icamen e la g an po encia- lidad que enían las GPU pa a el cómpu o numé ico, y u o de ello, NVIDIA lanzó en 2007 su lenguaje de p og amación p opie a io denominado CUDA C. Desde en onces p ác icamen e NVIDIA ha ido p opo cio- nando la capacidad de ejecución de p og amas CUDA a la g an mayo ía de sus a je as g á icas. Es a ca ac- e ís ica es á especi icada median e la denominada “capacidad de compu ación CUDA” (en inglés, CUDA capabili ies), que consis e en un núme o que pe mi e conoce qué ca ac e ís icas y qué compo amien os iene la a je a g á ica pa a la ejecución de código CUDA, sob e odo eniendo en cuen a que el lenguaje y las p es aciones ha e olucionado mucho desde 2007 has a la echa. Además, la nume ación ambién iene cie a co elación con las di e en es a qui ec u as ha dwa e que han apa ecido: Tesla,Fe mi,Keple y la más ecien e Maxwell. Pos e io men e, Apple, lide ando un conjun o de o as compañías (AMD,IBM,In el yNVIDIA), impulsó el lanzamien o de o o lenguaje pa a GPGPU, denominado OpenCL [13], que al con a io que CUDA es de especi icaciones abie as. Ac ualmen e el enca gado de la especi icación es el G upo Kh onos [8], quien ambién es el esponsable de la especi icación de OpenGL [14]. 7 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e Po ello, mien as CUDA únicamen e es á implemen ado en los p ocesado es g á icos de NVIDIA po se un lenguaje p opie a io y o almen e ce ado, OpenCL puede u iliza se no sólo en los p ocesado es g á icos de di e en es ab ican es, como ATI [2], In el [7] y la p opia NVIDIA [12], sino que incluso hay CPUs que son capaces de ejecu a código de OpenCL (IBM Powe 77X eIBM BladeCen e s,In el Ion y la e ce a gene ación de In el Co es [7] y los ecien es p ocesado es de AMD, incluidos los AMD Fusion [1]). Samsung ambién dispone de una implemen ación de OpenCL que puede ejecu a se sob e p ocesado es Cell BE, ARM y sob e DSP [15]. A pesa del mayo núme o de ab ican es que p opo cionan ha dwa e compa ible con OpenCL,CUDA es un lenguaje con una mayo e olución, y su compilado es á mucho más op imizado pa a el ha dwa e al que a des inado, dado que es á especí icamen e diseñado pa a u iliza odo el po encial de los p oce- sado es g á icos de NVIDIA.OpenCL es mucho más gené ico, e incluso, como ya se ha comen ado, puede se ejecu ado po CPU, po lo que iene un mayo g ado de abs acción sob e el ha dwa e y po an o, no pe mi e emplea ins ucciones o ecu sos especí icos de cie os p ocesado es g á icos. En la Tabla 2.1 se compa a el desa ollo de cada uno de es os lenguajes. CUDA OpenCL Ve sión del Toolki Fecha Especi icación 1.0 julio 2007 1.1 no iemb e 2007 junio 2008 1.0 2.0 agos o 2008 2.1 diciemb e 2008 2.2 mayo 2009 2.3 julio 2009 3.0 ma zo 2010 3.1 mayo 2010 junio 2010 1.1 3.2 no iemb e 2010 4.0 mayo 2011 no iemb e 2011 1.2 4.1 ene o 2012 4.2 ab il 2012 5.0 oc ub e 2012 5.5 julio 2013 no iemb e 2013 2.0 6.0 ab il 2014 6.5 agos o 2014 ene o 2015 2.1 p o isional 7.0 ma zo 2015 7.5 sep iemb e 2015 Tabla 2.1: Compa ación de la e olución c onológica de CUDA yOpenCL No obs an e, es necesa io ecalca que la compu ación a a és de GPU puede en algún momen o so- b e alo a se. Exis en campos en los que las ac uales CPU son an capaces como las más po en es GPU, y es po ello que siemp e es necesa io alo a p ime o si el p oblema a abo da , o cada una de las pa es en que se pueda di idi , debe se lo a a és de compu ación gené ica o de compu ación basada en GPU. Así, 8 Capí ulo 2. Conocimien os p e ios 2.4. CUDA en campos ales como las ma ices dispe sas, los cálculos en que en cada p ocesado pueden di e gi en cuan o a su lujo de con ol, o en los que exis e una g an dependencia en e da os en di e en es ubicaciones espaciales, la compu ación gené ica puede se más e icien e, al como se señala en [16]. 2.4 CUDA CUDA son las siglas de Compu e Uni ied De ice A chi ec u e. Ha sido p opues a, implemen ada y desa o- llada exclusi amen e y de o ma ce ada po NVIDIA. Se a a de una a qui ec u a ha dwa e pa a p ocesa- do es g á icos que pe mi e que és os puedan ealiza cómpu os de ca ác e gene al (no g á icos) de una o ma sencilla y di ec a, sin necesidad de in oca di ec amen e las p imi i as g á icas o iginales basadas en OpenGL oDi ec X. Todo ello sin disminui su p opia capacidad in ínseca de ejecución de a eas g á icas basadas en cualquie a de es os es ánda es g á icos. 2.4.1 Lenguaje CUDA C Es e concep o a qui ec ónico es ap o echado en la p ác ica median e la de inición de un nue o lenguaje que es capaz de abs ae la complejidad del mul ipa alelismo de los ac uales p ocesado es g á icos man e- niendo una es uc u a y nomencla u a muy simila al lenguaje C, uni e salmen e u ilizado pa a la p og a- mación gené ica usando CPUs. Es e nue o lenguaje se denomina CUDA C, y en lo sucesi o, en es e abajo, se ab e ia á simplemen e po CUDA, el mismo nomb e que la a qui ec u a. Se a a de un lenguaje de p og amación de especi icaciones ce adas, que desa olla en exclusi a NVI- DIA, que pe mi e la u ilización de sus p ocesado es g á icos pa a compu ación gené ica, no exclusi amen e g á ica. Los p og amas ealizados con CUDA solamen e puede ejecu a se en aquellos p ocesado es g á icos que engan «capacidad CUDA» (CUDA capabili y). Es impo an e di e encia en e la e sión del Toolki , que ep esen a el API de p og amación en CU- DA con una sin axis basada en el lenguaje C, y la capacidad de compu ación CUDA. Es a úl ima se a a de una ca ac e ís ica del ha dwa e, es o es, de la GPU, y ep esen a qué subconjun o de ins ucciones es capaz de ejecu a di ec amen e cada p ocesado g á ico. Exis e una e iden e in e elación, po cuan o en cada e sión del Toolki se menciona exp esamen e qué capacidad CUDA es necesa ia pa a cada una de las unciones de lib e ía que se p opo ciona, al obje o de u iliza op imamen e el lenguaje en unción del ha dwa e que se aya a u iliza . En [3] se lis an odas las GPU que ienen capacidad de compu ación CUDA indicando exp esamen e su ni el de capacidad, que ac ualmen e llega has a la 5.2 pa a las GPU GeFo ce más a anzadas (se ie GTX 970 y 980). 2.4.2 Tipos de código en CUDA C La p og amación en CUDA se ca ac e iza po con ene dos ipos bien di e enciados de bloques de código: el código del an i ión (hos code) y el código del disposi i o (de ice code). • An i ión (hos ): Po an i ión se en iende el en o no ha dwa e de ejecución gené ico p opo cionado po las CPU con encionales. Su código es pu o C/C++ y no p ecisa nada de los p ocesado es g á icos pa a su ejecución, pe o ob iamen e, no se ap o echa el po encial de ellos. • Disposi i o (de ice): Po disposi i o se en iende el en o no ha dwa e de ejecución p opio de los p o- cesado es g á icos, GPU. Su código, como ya se ha comen ado, iene una es uc u a y sin axis muy simila al C, pe o ex endido pa a pe mi i p og ama ácilmen e las nue as capacidades que posee la a qui ec u a CUDA. 9 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e CUDA p opo ciona un compilado p opio, denominado n cc, que es á especializado en la compilación de código de disposi i o, pe o que ambién compila la pa e del an i ión y e ec úa el enlace en e ambos códigos en un único p og ama ejecu able. El código del disposi i o se ca ac e iza po es a basado en ke nels, que ienen la apa iencia de uncio- nes es ánda C, pe o con una nomencla u a de llamada especial, y que ins uyen al p ocesado gené ico de la CPU a en ia el código y los da os necesa ios a la GPU, donde dicho ke nel e ec i amen e se ejecu a. Los esul ados de esos ke nels pueden de ol e se di ec amen e a la CPU, pa a lo cual es necesa io hace una nue a ans e encia de da os, en es e caso en e la GPU y la CPU, o incluso man ene se den o de la memo ia de la GPU pa a eu iliza en pos e io es ejecuciones de nue os ke nels. Consiguien emen e, la CPU siemp e es la que lle a el con ol de ejecución del p og ama global, pe o a delegando en la GPU aquellas pa es del p og ama que pe mi an ealmen e ob ene una g an en aja de la pa alelización masi a que o ecen las GPU. 2.4.3 Compilación de aplicaciones CUDA La compilación de las aplicaciones CUDA se ealiza undamen almen e con el d i e compilado n cc.n cc puede ealiza muchas unciones, siguiendo caminos de compilación di e en es según las necesidades, al como se mues a en la Figu a 2.1. Una de las p incipales consis e en sepa a el código del an i ión (hos , que es á en C) del disposi i o (de ice, pa a el que se usan ex ensiones especí icas CUDA C). Una ez se- pa ados los códigos, lanza los compilado es especí icos de cada ipo de código po sepa ado pa a después ensambla los en un único ejec uable, gene almen e. La pa e del an i ión se delega en un compilado de uso gené ico, como pueden se gcc oicc. La compilación del código de la a je a g á ica es mucho más c í ica, y sólo puede ealiza se con las he amien as especí icas p opo cionadas po NVIDIA, dado el ca ác e ce ado de la especi icación ha d- wa e de sus a je as, a pesa de se pública la especi icación del código in e medio, o ma o denominado PTX (Pa allel Th ead eXecu ion), pa a el cual NVIDIA p o ee el compilado p xas.p xas puede se in oca- do di ec amen e po n cc, pe o ambién puede hace se en iempo eal po la p opia aplicación, ealizando una compilación JIT (Jus In Time) del código in e medio a código máquina especí ico de la a qui ec u a y capacidades CUDA de la a je a conc e a en la que se a a ejec u a el ke nel. La compilación del código puede ealiza se de es o mas básicas, hacia es o ma os di e en es, a sabe : •.cubin: El o ma o cubin es mic ocódigo na i o especí ico de una a je a conc e a NVIDIA con capa- cidad CUDA. Po lo an o, es di ec amen e ejecu able po ella. Los códigos bina ios cubin en gene al no pueden in e cambia se en e di e en es a qui ec u as o di e en es capacidades CUDA (CUDA ca- pabili ies. •.p x: Es la ep esen ación in e media del código a ejecu a en la GPU. És e puede se compilado a su ez y con e i se en código bina io cubin, a iba desc i o. El o ma o p x es neu o, en el sen ido de que es independien e de la a qui ec u a y de las capacidades compu acionales de la a je a NVIDIA en la que inalmen e se ejecu a á. Es á diseñado pa a se esis en e a cambios u u os (« u u e-p oo ») po inno aciones ecnológicas ha dwa e y so wa e. Puede se compilado o line po el compilado d i e n cc, o compilado au omá icamen e online JIT según se necesi e. •. a bin: Es un o ma o especial que con iene en su seno di e en es compilaciones o p ecompila- ciones, bien sea po con ene an o código cubin como código p x, bien po ene código di igido a di e en es a qui ec u as o con di e en es capacidades CUDA, así como odo ello a la ez. 10 Capí ulo 2. Conocimien os p e ios 2.4. CUDA SOFTWARE ARCHITECTURE 58 n cc and PTX PTX (“Pa allel Th ead eXecu ion”) is he in e media e ep esen a ion o com- piled GPU code ha can be compiled in o na i e GPU mic ocode. I is he mecha- nism ha enables CUDA applica ions o be “ u u e-p oo ” agains ins uc ion se inno a ions by NVIDIA—as long as he PTX o a gi en CUDA ke nel is a ailable, he CUDA d i e can ansla e i in o mic ocode o whiche e GPU he applica- ion happens o be unning on (e en i he GPU was no a ailable when he code was w i en). PTX can be compiled in o GPU mic ocode bo h “o line” and “online.” O line compila ion e e s o building so wa e ha will be execu ed by some compu e .cu ile (Mixed CPU/GPU) n cc Hos -only Code GPU Code .p x, . a bin Hos Execu able (wi h embedded GPU code) CUDA Run ime (e.g., libcuda .so.4) Hos Compile .cu ile (GPU only) n cc GPU Code .p x, .cubin CUDA D i e Hos code (wi h embedded GPU code) CUDA D i e (e.g., libcuda.so) CPU Sou ce Code Hos Execu able O line componen s o applica ion build p ocess CUDA Run ime D i e API Figu e 3.2 n cc wo k lows. Figu a 2.1: Diag ama de los lujos de compilación más habi uales de n cc Es e iden e que los cubin son más ápidos en ca ga se en la GPU, pues o que se en ían di ec amen e en cuan o debe ejecu a se dicho ke nel. Po su pa e, los p x son más po ables, pues pueden compila se a cualquie a qui ec u a o capacidad igual o supe io a su o ma o, pe o incu en en la conside able sob eca - ga de la compilación en línea jus o en el momen o en que debe lanza se a la GPU. Los a bin son iche os más oluminosos, pe o con mayo compa ibilidad, según lo que se incluya, sob e odo si se incluyen p x de di e en es a qui ec u as y capacidades. Nó ese (y es o es c ucial pa a el p esen e abajo) que la compilación especí ica del p x se ealiza jus o en el momen o en que se a a lanza el ke nel, no du an e la ca ga en memo ia de la aplicación al inicio de su ejecución. De igual o ma que la selección del cubin conc e o, en el caso de que ya es én incluidos a ios y no sea necesa io compila . Es in e esan e indica que los ke nels CUDA pueden es a inc us ados en el p opio iche o ejecu able ELF (o EXE, en Windows), o ambién pueden ca ga se du an e la ejecución desde o os iche os, en cual- quie a de los es o ma os a iba indicados. De hecho, cuando el ke nel es á incluido en el p og ama p in- cipal, es á en o ma o cadena (s ing li e als), y lo que se hace es lee los y pasa los a la GPU median e unas ins ucciones especí icas p e ias, que se enca gan ambién de lanza su ejecución den o de la a je a g á- ica. 11 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e Po úl imo, la co ec a selección del ipo de a je a g á ica pa a el cual se quie e compila (di ec amen e a código bina io .cubin u ob ene código in e medio PTX) es esencial. Es a selección se e á de i al impo - ancia a la ho a de la mig ación, dado que en odo caso, la aplicación que se lanza debe á, como equisi o básico, pode ejecu a se an o en la GPU inicial como en la GPU a la que se mig a. 12 CA P Í T U L O 3 ECUACIONES DE LANGEVIN PARA DESCRIBIR UN PLASMA DE PARTÍCULAS Y LA GENERACIÓN DE ELECTRONES unaway En es e capí ulo esumen b e emen e la base ísica nuleó ica que sopo a el conjun o de ecuaciones que son esuel as numé ica a a és del código o iginal de Lange in ealizado en Fo an, y que pe mi en la simulación del compo amien o de los elec ones de un plasma a al a empe a u a con inado magné ica- men e, y e en ualmen e, bajo la acción de un campo eléc ico pe pendicula . P ime amen e se in oduci á el concep o de plasma, y su impo ancia económica po el impac o que pod ía llega a supone la ob ención de una eacción de usión nuclea con balance posi i o de ene gía, que esul a ía en una uen e ela i amen e limpia de ene gía a pa i de elemen os muy comunes en nues o en o no, ales como el Helio, el Deu e io, el T i io y el Li io. Pos e io men e, se epesa á muy po encima, la base eó ica de la ísica de pa ículas que explica el compo amien o de los plasmas, haciendo especial mención a una o ma de abo da las denominada ecua- ciones de Lange in. Es e mé odo de Lange in pe mi e, median e unas ap oximaciones azonables, ob- ene unas ecuaciones ela i amen e sencillas, sob e las que es posible aplica mé odos numé icos y con ello consegui la simulación del compo amien o del plasma. Es e abajo se cen a á especí icamen e en la simulación de un ipo de pa ículas muy i e esan es, los elec ones ugi i os ( unaway). 3.1 El plasma y su aplicación p ác ica: eacción con olada de usión El plasma es un es ado de la ma e ia en el cual és a es á some ida a al as empe a u as (gene almen e mi- llones de g ados Kel in), a esul as de lo cual, sus pa ículas es án o almen e ionizadas. En es e es ado su compo amien o es simila al de un gas, po la g an libe ad de mo imien o de sus pa ículas. Esa g an em- pe a u a y g an libe ad, implica una al a ene gía ciné ica, es o es, unas al as elocidades den o del medio. Las condiciones necesa ias pa a la gene ación y man enimien o de la ma e ia en es ado de plasma se dan comúnmen e en la na u aleza, pe o no en nues o inmedia o en o no, sino en el in e io de las es ellas, po ejemplo. 13 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e Pa a su u ilización en C(básicamen e, den o de las unciones w appe ), se c ean unos iche os de cabe- ce a (.h), que median e di ec i as de compilado #de ine, se ealiza una aducción del nomb e complejo en que apa ecen es as a iables globales en los iche os obje o compilados con Fo an. Pa a se gené ico, con iene inclui di ec i as compilado es de compilado que sean capaces de de ec a los posibles compila- do es que se an a usa , y según eso, u iliza la aducción adecuada. A con inuación se expone un ejemplo de cómo di e ncia en e compilado es pa a di ecciona co ec amen e el símbolo de una a iable global expo ado desde código Fo an: Lis ado 4.1: Compilación condicional median e de ección del compilado 1#i de ined(__ICC) || de ined(__INTEL_COMPILER) 2/*In el ICC/ICPC. ------------------------------------------ */ 3#de ine L_TWO_NOISES __lange in_pa ams_mp_l_ wo_noises 4#de ine L_DE_LOR __lange in_pa ams_mp_l_de_lo 5... 6#eli (de ined(__GNUC__) || de ined(__GNUG__)) && !(de ined(__clang__) || de ined( __INTEL_COMPILER)) 7/*GNU GCC/G++. --------------------------------------------- */ 8#de ine L_TWO_NOISES __lange in_pa ams_MOD_l_ wo_noises 9#de ine L_DE_LOR __lange in_pa ams_MOD_l_de_lo 10 ... 11 #endi En h p://nadeauso wa e.com/a icles/2012/10/c_c_ ip_how_de ec _compile _name_ and_ e sion_using_compile _p ede ined_mac os se mues a cómo de ec a más amilias de com- pilado , y con ello pode amplia es e ejemplo a una mayo a iedad de compilado es de C. Di e encias en los p o o ipos de unción en e Fo an yC A pesa de que Fo an no es á o almen e es anda izado, especialmen e en lo e e en e al name man- gling, los compilado es Fo an g o an de GNU ei o de In el compa en las siguien es ca ac e ís icas pa icula es en la c eación de los símbolos de las sub u inas y unciones: • Las unciones y sub u inas en el iche o obje o es án siemp e con minúsculas • Además, se les añade a su nomb e el ca ác e de sub ayado ( ambién conocido po “guión bajo”) 0_0. Resumen compa a i o de las di e en es nomencla u as según lenguaje y amilia de compilado Como e e encia y esumen inal, en la Tabla 4.1 se ejempli ican algunos casos: Tipo de símbolo Tipo Nomb e decla ado Símbolo c eado po Símbolo c eado po nm en Fo an g o an (GNU)i o (In el) Función p opia T ADVANCE ad ance_ad ance_ Función ex e na (inde inida) U RHS_COEFFS_ hs_coe s_ hs_coe s_ Va iable p opia inicializada D MYDT __ad ance_MOD_myd __ad ance_mp_myd Va iable p opia no inicializada B MY_ITER __ad ance_MOD_my_i e __ad ance_mp_my_i e Va iable ex e na (inde inida) U DT __o b_pa ams_MOD_d __o b_pa ams_mp_d Tabla 4.1: Ejemplos del name mangling de los compilado es Fo an usados 20 Capí ulo 4. Implemen ación de la pa alelización con CUDA Sección 4.2 En la abla an e io , se mues an los símbolos que se c ean en el iche o obje o p oduc o de la compila- ción de un código uen e de un supues o iche o ADVANCE.F que de ine la sub u ina ADVANCE(). Se supone que es a sub u ina ealiza una llamada a la unción RHS_COEFFS que es á decla ada y de inida en un iche- o ex e no. En la sub u ina ADVANCE se de inen dos a iables, MYDT yMY_ITER, las cuales son globales pa a odo el p og ama Fo an, una de ellas inicializada y o a no, y u iliza la a iable ex e na DT decla ada en el iche o ORB_PARAMS. Más in o mación sob e el signi icado y ipos decla ados po la u ilidad GNU nm se puede ob ene en h p://www. eebsd.o g/cgi/man.cgi?que y=nm&ap opos=0&sek ion=0&manpa h=F eeBSD+9. 0-RELEASE&a ch=de aul & o ma =h ml 4.2.2 Esquema gene al de llamadas a CUDA desde Fo an a a és de w appe s C El mé odo elegido pa a adap a el p og ama en Fo an 90 a CUDA consis e en la c eación de una unción w appe en C, in ocada de o ma no mal desde Fo an, que ecoge los pa áme os y accede a las a iables globales decla adas en Fo an necesa ias pa a pasa al ke nel. A pa i de ese momen o, odo el código puede hace se en Co en C++, dado que se compila á con el compilado n cc p opo cionado po CUDA. Se debe án ene en cuen a las cues iones de nomencla u a y de paso de pa áme os ya expues as en el an e io apa ado. En gene al el p oceso se á: • La unción w appe C, debe á de ini se siemp e con los cali icado es ex e n "C" oid, dado que una sub u ina en Fo an nunca de uel e alo . El ex e n "C" pe mi e la compa ibilidad con la nomecla u a de unciones de Cdesde C++, pues o que n ccen ealidad es un en ol o io so is icado del compilado C++ p esen e. • En Cla unción debe á e mina siemp e con el ca ác e de sub ayado adicional ( ambién conocido como guión bajo) ’_’, y con odos sus ca ac e es en minúsculas. Todos sus pa áme os se án pun e- os, y es ecomendable que aquellos que no deban se modi icados po esa unción o sus de i adas (pa áme o sólo de en ada), lle en el cali icado cons . • Desde Fo an se in oca á a la unción po el nomb e de inido en C, con cualquie combinación de mayúsculas y minúsculas, pe o obliga o iamen e sin el úl imo ca ác e de sub ayado. Además, debe in oca se como se hace a una sub u ina de Fo an (con CALL <nomb e_sub u ina(lis a_pa áme os)), y no como se hace a una unción de C. • En la unción w appe Cse ges iona án los pa áme os de en ada, y sob e odo, se ecomienda allí u iliza mac os pa a enomb a las a iables globales de Fo an, po se emendamen e a agosas de esc ibi al cual desde C. A pa i de allí, con esos sinónimos en Cse ope a á no malmen e, siemp e eniendo en cuen a que se a a de pun e os. Ex ema la p ecaución con los pa áme os que sean ma ices mul idimensionales po la can idad de indi ecciones que se án necesa ias pa a su co ec a e e encia al alo , o a la di ección (pun e o donde se encuen e el alo eal). • En el w appe Cse ges iona án las ans e encias de da os que sean necesa ias en e la CPU y la GPU. Se ecomienda, po e iciencia, e i a en odo lo posible la c eación de nue as zonas de memo ia pa a cambia el o den de acceso de las ma ices. • Desde el w appe Cse in oca á ya de o ma no mal a los ke nels, siguiendo las pa icula idades del lenguaje CUDA-C, en especial, de la in ocación a dichos ke nels. 21 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e • Den o de los ke nels, donde se u iliza el lenguaje CUDA-C, se deben ene en cuen a, apa e de odo lo ca ac e ís ico de ese lenguaje ex ensión de C, la p ocedencia de las zonas de memo ia mul idimen- sionales (a ays). Si son en Cdeben accede se de la o ma habi ual (o den po ilas), pe o si p o ienen ( ue on de inidas) en Fo an, en onces el acceso debe ealiza se eniendo en cuen a que el almace- namien o es po columnas, es o es, el segundo elemen o en la zona de memo ia no es el elemen o A(1,2), sino el elemen o A(2,1). Es po ello que se ecomienda el uso de mac os en Cpa a di eccio- na siemp e con ellas las zonas de memo ia de inidas en Fo an. A con inuación se mues an las mac os u ilizadas en es e p oyec o pa a ealiza al acceso, en unción de si se p e ende accede al alo , o a la di ección de memo ia donde es á almacenado dicho alo . Es as mac os debe án de ini se pa a cada ipo de a ay mul idimensional. Las que se p esen an son pa a ma ices bidimensionales de núme os lo an es de doble p ecisión (**double odouble[][]). Lis ado 4.2: Mac os en C pa a el acceso a ma ices 2D de inidas en Fo an 1// Con e linea ec o ep esen a ion o 2D ma ix in C-like ile/column o de 2// Use: MATRIX2D_DBL_PTR(poin e _ o_ma ix_s a , ile, column, ile_wid h) 3#de ine MATRIX2D_DBL_PTR(m,i,j,N) (((m))+((j)*(N)+(i))) // A poin e o alue 4#de ine MATRIX2D_DBL_VAL(m,i,j,N) (*(((m))+((j)*(N)+(i)))) // The alue i se Siguiendo es as indicaciones, supongamos una unción denominada ADVANCE, que es á decla ada en el código o iginal en Fo an, y que al ealiza g an compu ación numé ica la amos a sus i ui po un ke nel. Sea PARAM1 sea un pa áme o de en ada, y PARAM2 o o de salida. Su de inición en Fo an se á: Lis ado 4.3: De inición de una sub u ina común en Fo an 90 1SUBROUTINE ADVANCE(PARAM1, PARAM2) 2INTEGER,INTENT(IN) :: PARAM1 3REAL*8, INTENT(OUT):: PARAM2 4... 5END SUBROUTINE ADVANCE Pa a sus i ui esa unción po un w appe en C, su decla ación y con enido básico se á la siguien e: Lis ado 4.4: De inición de la unción w appe en C 1ex e n "C" oid ke nel_ad ance_w appe _(cons in *pa am1, double *pa am2) { 2... 3// Copy da a om CPU o GPU, i apply 4ke nel_ad ance<<<G idSize,BlockSize(pa am1, ...); 5// Copy da a om GPU o CPU, i apply 6... 7 e u n;// I ’s oid. So we ha e o e u n no hing! 8} Finalmen e, es a unción w appe en Cse in oca desde Fo an de la siguien e o ma: Lis ado 4.5: Ejemplo de llamada a un w appe Cdesde Fo an 1CALL KERNEL_ADVANCE_WRAPPER(PARAM1, PARAM2) 22 Capí ulo 4. Implemen ación de la pa alelización con CUDA Sección 4.3 4.3 Di e encias p og amá icas impo an es en e Fo an yCyCUDA-C En las secciones an e io es ya se han comen ado aspec os di e enciado es en e CyFo an, pe o que e- nían más que e con la sin axis. En es a sección se ha á especial hincapié en las p incipales di e encias semán icas y p og amá icas que deben conside a se en la pa alelización con CUDA-C de p og amas Fo - an, que son las siguien es: Paso de pa áme os: Po alo en e a po e e encia En Cel paso de pa áme os es po alo siemp e (pa a pasa po e e encia, en ealidad se hace un pase po alo , pe o del pun e o a la zona de memo ia e e enciada), en Fo an el paso es siemp e po e e encia. En un apa ado pos e io se analiza án las consecuencias de es e hecho. De es a o ma, en Fo an no se puede inclui una cons an e como pa áme o, y po o a, desde el w appe en C, al ecibi un pa áme o siemp e se á un pun e o, y como al, es a á expues o a se modi icado o eliminado desde C. Es po ello que se ecomienda que sea decla ado como pun e o cons an e (cons ) en el p o o ipo en C. Va iables globales Fo an den o de una sub u ina llamada po di e en es hilos OpenMP Las sub u inas Fo an, cuando son llamadas desde di e en es hilos de ejecución, no compa en las a iables globales de inidas en Fo an. Es án disponibles, pe o cada hilo iene su p opia e sión lo- cal, sin que se copie el alo global cuando és a enía al se llamada. Pa a soluciona es e p oblema, se p opone el uso de pa áme os adicionales en el p o o ipo de la sub u ina Fo an. Así, po ejemplo, si se p ecisa ene acceso a la a iable global h ead del p o- g ama Fo an, se debe añadi un nue o pa áme o, con el modi icado INTENT adecuado ("IN", o "INOUT"), pa a que en la sub u ina se u ilice es a nue a a iable local con el alo que enía su co- espondien e a iable global en el p og ama p incipal en el momen o en que ue llamada. De lo con a io se inicializa a ce o. Es e p oblema a ec a ambién a las ma ices y a iables de inidas como ALLOCATABLE. Di e en e o den de almacenamien o en memo ia de ma ices mul idimensionales Ya ha sido comen ado en el apa ado an e io , pe o es an impo an e que me ece un nue o eco da- o io. En Cel acceso a las zonas de memo ia mul idimensionales se de ine y ealiza po ilas p ime o, po columnas después (si es mul idimensional, de izquie da a de echa en el o den de las dimensiones al cual se esc iben). En Fo an, pa a ma ices 2D, es po columnas p ime o y después po ilas. Si son mul idimensionales, de de echa a izquie da en el o den que apa ecen las dimensiones o índices. Di e en e inicio del índice de ec o es y ma ices En C, el índice de inicio siemp e es 0, an o en ilas como en columnas. No es lo mismo en Fo an, en donde el índice del p ime elemen o no malmen e empieza po uno ( éase la excepción en la siguien e ca ac e ís ica). Eso es muy impo an e a la ho a de eu iliza el código en Fo an como base pa a su aducción a C: los bucles e índices deben e isa se con enien emen e pa a que accedan co ec amen e al elemen o co espondien e en cada momen o. Ma ices mul idimensionales con índices que no empiezan en ce o En Fo an es posible de ini ma ices en las que una (o a ias) dimensiones no empiecen en 1 como no malmen e se hace. Puede de ini se que el índice de una de e minada dimensión enga un in e - alo cualquie a, como po ejemplo, en e -3 y +6, en luga de 1 a 10. En Clos índices siemp e se án, en ese ejemplo, en e 0 y 9, po lo que a la ho a de mig a el código de Fo an aCdebe án con em- 23 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e pla se es as e en ualidades pa a que el acceso sea siemp e al elemen o co ec o (apa e de que se pod ía, además, es a accediendo a zona de memo ia ue a de la ma iz, con iesgo de co upción de la misma). Compa ibilidad en e los di e en es ipos numé icos Es muy impo an e ene en cuen a ealmen e cuál es el espacio jus o de almacenamien o de las a iables en Csegún su decla ación en Fo an, donde se pueden de ini , po ejemplo, núme os eales de X decimales e Y ci as en e as. Es o da luga a ipos de a iables que deben analiza se pa a asegu a que en Cse le asigna un ipo compa ible en signo, en capacidad (4, 8, 16 By es) y en p ecisión. Especial a ención a los ipos inexis en es en C, como los LOGICAL de Fo an. O as ca ac e ís icas meno es No me ece la pena en a en de alle de la mul i ud de pequeños de alles de di e gencia en e es os lenguajes. Como me o ejemplo, en es e p oyec o en algún momen o se pensó en la con eniencia de u iliza un enume a e de C, pe o se obse ó que en Fo an no exis en las denominadas “cons an- es nomb adas”. Po lo an o, se iene que abaja con a iables no males di ec amen e, con alo es en e os. Po ello se p ecisa un expe o conocedo de ambos lenguajes (Fo an,CyCUDA-C) pa a ealiza una co ec a y óp ima mig ación o adap ación del código his ó ico en Fo an aCUDA. Como co ola io, es emendamen e impo an e ene p esen e en odo momen o de dónde p oceden las ma ices, pues o que según hayan sido de inidas en Co en Fo an, su acceso y ipología son di e en es. Además, es muy impo an e eco da que según en qué lenguaje es emos p og amando, el acceso a las ma ices y su de inición, son di e en es. 4.4 Implemen ación en CUDA del código de Lange in o iginalmen e en Fo an En es a sección se an a comen a b e emen e las p incipales ca ac e ís icas de la implemen ación ealiza- da esul an e de la mig ación pa cial del código de Lange in en Fo an aCUDA. En una sección pos e io se comen a á la modi icación que a és a p ime a implemen ación en CUDA debe ealiza se pa a ob ene una mejo a de endimien o po el desacoplamien o en e compu ación e in e cambio de mensajes y esc i- u a en disco. 4.4.1 Aplanamien o de sub u inas po inclusión en un ke nel aglu inado En CUDA la in ocación a los ke nels, y los cambios de con ex os in e nos en e los di e en es hilos de la GPU son menos cos osos que en los lenguajes adicionales de CyFo an. Po o o lado, es muy impo an e el núme o de egis os y o os ecu sos, como las unidades lógicas de coma lo an e (simple y doble, pues son independien es) no se sa u en, po que eso disminuye la can idad de hilos que pueden es a en ejecución simul áneamen e en el mismo núcleo in e no (bien sea en es ado ac i o, con ac ualización cons an e del con ado de p og ama, bien en es ado pasi o, po es a a la espe a de la inalización de alguna ope ación la ga, como las a i mé icas de doble p ecisión). Consiguien emen e, es muy ac ible en CUDA ob ene buenos esul ados incluso aunque se u ilicen muchos ke nels simples, siemp e que és os p ecisen pocos ecu sos, pe o el núme o o al de ecu sos que- den bien sa u ados, de o ma lo más homogénea posible, po la in ocación y ejecución en cada núcleo del 24 Capí ulo 4. Implemen ación de la pa alelización con CUDA Sección 4.4 núme o de bloques óp imo (o un poco más, nunca menos), como se demos a á en el siguien e Capí ulo 5 “Valo ación cuali a i a del p oceso de pa alelización con CUDA”. A pesa de ello, y como demues a es e abajo, es posible ambién con a con un buen endimien o ocupacional aunque el ke nel sea más bien complejo, si la u ilización de los ecu sos es homogéneo. Tal es el caso que se p oduce en el ke nel_ad ance, que es el p incipal esponsable del iempo de ejecución global del p og ama. Es e ke nel se enca ga de, pa a cada pa ícula, calcula la nue a elocidad ec o ial en unción de la elocidad an e io , del alo de los campos magné ico y eléc ico, y de un componen e es ocás ico ( undamen al al se unas ecuaciones del ipo Lange in) que simula el e ec o colisional del es o de pa ículas (elec ones e iones) a su al ededo . En el código o iginal, al como se obse a en la Figu a 5.1, la sub u ina ADVANCE es esponsable del 5,45% del iempo de compu ación, pe o és a llama exclusi amen e a o as sub u inas de ayuda. Es as sub u inas de ayuda si en pa a calcula algunos pa áme os conc e os de la ó mula de Lange in, que es la que se compu a en la sub u ina pad e ADVANCE, ales como el e ec o de la colisión con iones (RHS_ION) y el ángulo de la elocidad esul an e as impac a con iones o elec ones, (PHI_AND_PSI, llamada a a és de la o a sub u ina RHS_COEFFS). También se u ilizan sub u inas auxilia es, como GET_GAUSSIAN_NOISE, enca gada de ealimen a el plasma con un elec ón es adís icamen e simila a la condición inicial po cada elec ón ugi i o ( unaway) que se escapa del con inamien o magné ico. Es a sub u ina, GET_GAUSSIAN_NOISE, a su ez, llama a una sub u ina de biblio eca, RAN2, que es la que gene a los núme os alea o ios necesa ios pa a el compo amien o es ocás ico de la ecuación de Lange in. Po lo an o, en ealidad, la sub u ina ADVANCE, y las dependien es de ella, es esponsable del 98,5% del iempo de ejecución del código o iginal. Es a cadena de llamadas de sub u inas no puede ealiza se de o ma equi alen e con ke nels. No en el sen ido de que un ke nel llame a o o ke nel. Aunque eso es posible a pa i de las a je as Keple , esas in ocaciones de un ke nel po o o, es á pe mi ido sólo con un ni el de anidamien o de uno, e implica una nue a c eación de muchos más bloques con sus espec i os hilos. Es á pensado pa a pe mi i , has a cie o pun o, la p og amación dinámica, un pa adigma de p og amación que no es aplicable a nues o p og ama. Lo que se hace, es in oca a un único ke nel, el cual puede llama a o as unciones en CUDA-C que lo que hacen es simpli ica la expansión del código comple o en una única unción ke nel. Pe o en ealidad, el ke nel es único, lo que pasa es que és e llama a unciones, compiladas en el código p opio de la GPU (PTX). Po lo an o, en el código p esen ado, el ke nel_ad ance_w appe _incluye o as unciones CUDA-C, hs_coe s y hs_ion, que son la mig ación de sus co espondien es sub u inas de Fo an. Se obse a que “desapa ecen” las sub u inas GET_GAUSSIAN_NOISE yRAN2. Eso es debido a que las a je as NVIDIA, a a és de CUDA ienen sus p opios (y a iados) gene ado es de núme os alea o ios, muy p obablemen e más e icaces y segu os que cualquie o o mé odo que se quisie a implemen a di ec amen e con CUDA-C en la GPU. 4.4.2 Minimización de la ans e encia de da os en e CPU yGPU Las implemen aciones p esen adas han minimizado las ans e encias de da os en ambos sen idos en e la CPU y la GPU. Únicamen e se ansmi en los siguien es, que son ine i ables: • Vec o es de elocidades iniciales cuando és os son leídos de un iche o de einicio ( iche os con el nomb e lange in. es a .????): E iden emen e, pues la GPU no puede accede di ec amen e al disco du o. Nó ese que si el p og ama debe gene a las elocidades iniciales, según los pa áme os indicados en el iche o de con igu ación (dis ibución uni o me o gaussiana), és as se gene an ya 25 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e di ec amen e en la GPU, sin in e ención de la CPU, quien desconce á has a la inalización de la eje- cución el alo conc e o de la elocidad de cada pa ícula del plasma. • Vec o es de elocidades inales: Es necesa io pasa los de la GPU a la CPU al inal de la compu ación, po la misma azón: la GPU no puede gua da esos da os en iche os lange in. es a .???? en disco du o. • Valo es es adís icos in e medios: La p opia GPU es la enca gada de ca aloga la elocidad de cada pa - ícula en su co espondien e casille o pa a inalmen e en ia a la CPU sólo los da os es adís icos ya elabo ados que pe iódicamen e, si así se ha indicado, deben g aba se en disco o mos a se po conso- la. Es as es adís icas pueden se de di e en es alo es y ca ac e ís icas, pe o siemp e son compu adas po la GPU y ansmi idos los esul ados es adís icos elabo ados a la CPU. • La con abilización del núme o de elec ones ugi i os gene ados: Que en ealidad es edundan e con lo an e io , pues o que es o a es adís ica más. Con ello se consigue elimina la necesidad de in e cambia se la in o mación sob e las elocidades de las pa ículas en e la GPU y la CPU. La e sión que de ini i amen e eliminó esa necesidad p opo cionó una acele ación in e na en e e siones de p ác icamen e el 100%. Es o es, el in e cambio en cada i e ación empo al de los ec o es de elocidad suponía un cos e semejan e al cos e compu acional de los ke nels. 4.4.3 U ilización adecuada de los gene ado es de núme os alea o ios de CUDA CUDA, median e la biblio eca es ánda CURAND, p opo ciona un muy a iado y e icien e conjun o de ge- ne ado es de núme os alea o ios. Tal es su a iedad, que se pueden selecciona di e en es mé odos de ob- ención de núme os, que se basan en di e en es mé odos, con en ajas y des en ajas espec o de la ag u- pación, segu idad, dis ibución y elocidad de gene ación de los núme os alea o ios. En es e abajo no se ha en ado a alo a las di e en es écnicas base, y se han seleccionado los que pa ecían más usualmen e usados en la bibliog a ía, en u o iales y o os p og amas. Además, las unciones que apo a, pe mi en la gene ación di ec amen e de núme os con dis ibución uni o me (la más usualmen e implemen ada de o ma es ánda en los demás lenguajes), pe o ambién de núme os con dis ibución no mal ( ambién conocida como dis ibución gaussiana). Es más, incluso es po- sible gene a los núme os alea o ios po pa es, con una mejo a sob e un 25% espec o a gene a dos nú- me os alea o ios de o ma indi idual. Ambas ca ac e ís icas han sido u ilizadas en es as implemen aciones, pe mi iendo que ue a la p opia GPU la enca gada de sin e iza los é minos es ocás icos di ec amen e den- o del ke nel. O o aspec o de endimien o muy impo an e a ene en cuen a es que los núme os alea o ios deben se inicializados. Es a inicialización no debe ealiza se en cada in ocación del ke nel, sino que si las a iables de los gene ado es alea o ios son gua dados en la zona de almacenamien o global de la GPU, se asegu a su pe manencia du an e oda la ejecución de la aplicación, aunque el ke nel conc e o que los haya inicializado o usado po úl ima ez inalice. De es a o ma, es posible ene los gene ado es ya inicializados pa a cada pa ícula, e in oca la pe ición de nue os núme os alea o ios ap o echando las semillas exis en es de una in ocación a o a del ke nel. Cuando es e enónemo se descub ió, se ob u o una acele ación ela i a del 35% en el iempo global de ejecución del p og ama global, en e las dos e siones implicadas. Es o signi ica que aunque los núme os alea o ios pueden se ápida y e icien emen e gene ados en la GPU, no es así su p oceso de inicialización. Es po ello que es con enien e inicializa los una única ez du an e oda la ejecución de la aplicación, y eusa los siemp e que sea posible. 26 Capí ulo 4. Implemen ación de la pa alelización con CUDA Sección 4.4 O a mejo a ue la de asegu a se que es os gene ado es es án en memo ia compa ida en el momen o de la ejecución del ke nel. Es o es, en gene al es án en memo ia global, pe o cuando son necesa ios se in ocan a ias eces. Pa a asegu a nos de que únicamen e se bajan y suben de la memo ia global a la más ce cana a los núcleos CUDA una ez (al inicio y al inal de cada in ocación al ke nel) y no depende de que es én o no oda ía p esen es en la cache, se hace una copia de ellas en memo ia sha ed y después, al e mina el ke nel, se sube el es ado del gene ado alea o io de nue o a la memo ia global. Es o es posible po que enemos un gene ado po cada una de las pa ículas. No obs an e, hay que señala un p oblema que se p odujo con los núme os alea o ios: en de e minadas simulaciones, es necesa io con a con núme os alea o ios dis ibuidos uni o memen e pa a unas cosas, pe o o os dis ibuidos de o ma no mal (gaussiana) pa a o as. En esos casos, no puede u iliza se la mis- ma a iable gene ado a, siendo necesa io c ea gene ado es di e en es pa a dis ibuciones di e en es. Es o implica un sob ecos e, pe o se ealiza una única ez du an e el p og ama. Como idea de op imización se deja la posibilidad de p oba di e en es gene ado es, pa a e si su ini- cialización y cons an e gene ación es más ápida en unos que en o os. Nó ese que nues o p og ama, en p incipio, no equie e de especiales condiciones de ex ema segu idad en cuan o a la calidad de los núme- os alea o ios gene ados. O a idea de op imización consis e en segui los consejos que en di e sos o os apa ecen, consis en es en que se inicializa sólo una a iable gene ado a de núme os alea o ios, o un núme o educido de ellos, po ejemplo, una po cada bloque. Pos e io men e, cuando se solici an núme os alea o ios, en luga de pedi el siguien e, se pide el segundo, e ce o o cua o, e c. según el núme o de hilo. Dicen que si se hace a anza la cuen a, es posible que no exis an colisiones de núme os gene ados. No obs an e, hay c í icas a es os mé odos po cuan o que su u ilización implica la pé dida de la segu idad eó ica de los algo i mos base de la implemen ación de es os núme os alea o ios. Po lo an o, los c í icos indican que los núme os ob enidos pueden es a sesgados, o ene cie as ag upaciones locales indeseadas. Es un ema abie o a la expe imen ación, an o pa a e si acele a ealmen e las ac uales implemenaciones, y ambién sob e si los esul ados inales se mues an es adís icamen e semejan es o di e en es a los ob enidos con la ac ual implemen ación, más conse ado a y posiblemen e más len a, pe o segu a. 4.4.4 Mejo as en el código: u ilización de unciones es ánda en e a mé odos numé icos pesados En el código o iginal, la sub u ina ad ance debe implemen a una de e minada in eg al. En el código o igi- nal, la implemen ación es al cual, es o es, se ealiza una in eg ación numé ica usando el mé odo apezoi- dal con unos 200 pasos. Es o compu acionalmen e es cos oso. Si se analiza la in eg al, y como en algunos a íclos pu amen e cien í icos donde se desa ollan es as ecuaciones de Lange in pa a plasma, esa in eg al ha dado luga a la de inición como al de una unción ampliamen e conocida: la unción e o . La so p esa es á en que es a unción es an ecuen e, que p ác icamen e odos los lenguajes compu- acionales (es o es, que mínimamen e si an o se usen pa a cálculos ma emá icos), ienen p ác icamen e de se ie, alguna unción p opia o en una biblio eca numé ica es ánda , implemen ada es a unción. Tal co- mo incluso se ha comp obado expe imen almen e en es e p oyec o, la llamada a es a unción es mucho más p ecisa, exac a y ápida que el mé odo de in eg ación implemen ado en Fo an den o del código o i- ginal. Es más, po ejemplo, en C, el e o de la unción es á asegu ada en se como máximo dos unidades en la úl ima posición ep esen a i a del núme o lo an e de doble p ecisión esul an e, siendo de una como máximo en la mayo pa e del in e alo azonable y no mal de uso. Es e iden e que doscien as sumas y mul iplicaciones pa a un cálculo numé ico ap oximado a a ene una exac i ud de a ios ó denes mucho peo . 27 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e La unción e o en Cy en CUDA-C iene el mismo nomb e, y se a a de e (), la cual se ha u ilizado en el ke nel_ad ance, con mejo a de p es aciones. Es a mejo a ha sido del 1–2% en el momen o de ealiza el cambio. 4.4.5 Algo i mos de educción en la GPU El sec e o e iden e de la buena acele ación que se ob iene en los esul ados de las implemen aciones p e- sen adas en es e abajo, es la absolu a independencia de los cálculos en e cada una de las pa ículas simu- ladas. Po ello, no exis e dependencia de da os en la GPU y p ác icamen e se asegu a que odos los accesos a memo ia son coalescen es, p opiedad que mejo a las p es aciones compu acionales de la a je a g á ica, al mejo a el asiego de in o mación in e no en e la memo ia y los núcleos de la GPU. No obs an e, exis e un pa de si ios en el algo i mo, en donde la GPU se e o zada a ealiza una educ- ción. Es o es, debe ob ene un esul ado a pa i de los alo es calculados po el es o de hilos y bloques, es o es, de odas las pa ículas: • Cálculos es adís icos pa a ob ene es adís icas sob e el alo de las componen es de la elocidad de los elec ones (pe pendicula y pa alela al campo magné ico): ca ac e izado po que se deben siemp e ene en conside ación el alo de odas las pa ículas (po lo an o, de odos los hilos pa icipan es en el ke nel). • Recuen o del núme o de elec ones que en cada i e ación pasan a se elec ones ugi i os ( unaway), po excede la elocidad c í ica de con inamien o: se ca ac e iza po que, en condiciones no males del plasma, se a a de una baja p opo ción de elec ones (hilos) los que deben compu a se. Es as educciones se pueden abo da al menos desde dos pe spec i as di e en es: • Median e algo i mos p opios de educción, gene almen e con cos e compu acional de o den O(log2(N)): Son los indicados pa a cuando el po cen aje de hilos (pa ículas) a pa icipa en la educción es ele- ada espec o al o al. • Median e ope aciones a ómicas que cada pa ícula ealiza sob e una a iable global común: Son los indicados cuando el po cen aje de hilos (pa ículas) a pa icipa es desp eciable en e al conjun o o al de pa ículas. La elección del mé odo es muy impo an e po su impac o en el iempo de compu ación. Es a decisión es un cla o ejemplo de las di icul ades que suelen ene los p og amado es que no son expe os en inge- nie ía in o má ica, dado que ecuen emen e se ob ia la ealización de un es udio p e io (o aunque sea pos e io , expe imen al) de cos es del algo i mo a implemen a . La elección pa a cada uno de los dos casos es di e en e, y se basa en un es udio de cos es: Es adís ica de elocidades Cada hilo se ubica en su con enedo según sus componen es de elocidad pe - pendicula y pa alela, pe o hay que e cuán as pa ículas hay en cada con enedo . Po lo an o es una educción en las que siemp e pa icipan odas las pa ículas. Es po ello que se equie e un algo i mo de educción basado en educción bina ia en á bol, que iene un cos e de o den O(log2(N))- Con eo de elec ones unaway En condiciones no males de plasma, el núme o de elec ones ugi i os es muy pequeño. Si ue a g ande, el okamak es a ía descon olado y la eacción nuclea de usión se pa a ía en pocos milisegundos. Po lo an o, no iene sen ido es a simulando en esas condiciones. Consiguien emen e, in e esa el algo i mo basado en ope aciones a ómicas sob e una misma a iable, 28 Capí ulo 4. Implemen ación de la pa alelización con CUDA Sección 4.4 pues en el caso ex emo de que coincidan odas las ope aciones de adición en el mismo momen o, és as se se ializa ían, y el cos e se ía de o denlineal, O(log2(n)), con nsiendo el núme o de colisiones. Analicemos cada una de es e ipo de educción po sepa ado. 4.4.5.1 Reducción global i e a i a Es a educción iene sus iquiñuelas cuando se iene que ealiza en GPU, sob e odo cuando la capacidad compu acional CUDA de la a je a usada es más baja. Es o es así, po que és as manejan mucho peo los accesos no coalescen es a la memo ia. Pa a esol e es a e en ualidad, es impo an e plani ica es os acce- sos pa a que sean lo más coalescen es posibles, y siemp e bajo el p isma de que se á necesa io ealiza una en ada ecu si a (en ealidad, i e a i a) de o den O(log2(N)), siendo Nel núme o de pa ículas, pa a su inal educción. En el SDK de CUDA apa ece un mé odo e iblemen e e icien e que u iliza unas mac os de compilado pa a ob ene la máxima po encia de la a je a g á ica, independien emen e de su capacidad compu acio- nal, al op imiza mucho el acceso coalescen e a la memo ia. No obs an e, es un mé odo muy ebuscado como pa a pode eu iliza lo. Es po ello la implemen ación se ha ealizado u ilizando la écnica del á bol bina io, con cos e O(log2(m)), siendo mel núme o de hilos en cada bloque (32, o sea, 6 i e aciones), pa a pos e io men e, hace o a pasada y ol e a o aliza esas sumas pa ciales ( o al de bloques di idido po 32), una po cada bloque lanzado. Como ejemplo, se puede acudi al código uen e de la implemen ación, al iche o ke nel_calcula e_s a is ics.cu, conc e amen e a los ke nels ke nel_calcula e_s a is ics_ a ios yke nel_calcula e_s a is ics_ a ios2D, que son demasiado ex ensos pa a su inclusión en es a memo ia. 4.4.5.2 Reducción median e el uso de ope aciones a ómicas CUDA p o ee de ope aciones a ómicas sob e a iables globales. El único p oblema con las ope aciones a ómicas en la GPU es que has a el momen o (has a CUDA e sión 7.0, capacidades compu acionales has a 5.2), no es á pe mi ido ealiza una educción sob e a iables de longi ud de 64 bi s. Es o es, no se puede ha- ce un a omicAdd sob e un en e o la go (long in ) o sob e un núme o en coma lo an e de doble p ecisión (double). Como en nues o caso, el con eo se ealiza con en e os no males, de 4 oc e os, es pe ec amen e iable, y ecomendable desde el pun o de is a de eo ía de cos e compu acional. Simplemen e consis e en selecciona sólo aquellos hilos (pa ículas) que deban con abiliza se, y con ello, in oca di ec amen e a la unción a omicAdd. Veamos como ejemplo, la o ma en que se u iliza en el ke nel que ecuen a los elec ones con elocidad mayo que la c í ica, conside ados po ello en la simula- ción como elec ones unaway. El Código 4.6 es un ejemplo de su u ilización. 29 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e El p oblema con odo es e esquema es que las de enciones en espe a del semá o o, es que se a a de espe as ac i as, en un bucle acío, simplemen e comp obando cons an emen e la condición del cambio de es ado. Es o es ine icien e ene gé icamen e, pues obliga a que la CPU es é siemp e al 100% de u iliza- ción. Además, según el plani icado , puede que le es é obando ciclos p eciosos al o o hilo que sí que es á haciendo a eas p o echosas (compu ando, o ansmi iendo o esc ibiendo en disco). Se p opusie on soluciones al e na i as que ealmen e no pa ecían llega a soluciona o almen e el p o- blema: • Uso de señales: La idea es c ea los se icios de la a iable condición u ilizando los se icios del sis- ema ope a i o. Se desconoce cómo usa , lanza y ges iona señales del sis ema ope a i o en Fo an. Po ello, se desca a inmedia amen e. • Do mi el hilo en espe a, pa a que no es é en espe a ac i a, y despe a lo pasado cie o iempo. El p oblema es á de nue o en que Fo an no p o ee de unciones de de ención in e io es a un segundo. Se llegó a c ea un w appe pa a in oca a una nue a unción en Cque pe mi ía ealiza espe as desde un mic osegundos a a ios milisegundos. En las p uebas ealizadas, no obs an e, se comp obó que la g anula idad e ec i a mínima e a de unos 250 µs (0,25 ms), y que es e alo e a a iable según el en o no de ejecución. Una espe a de 0,5 ms im- plica 2.000 espe as así cada segundo. Los iempos de ejecución indicaban que en ese iempo, incluso el hilo 0 compu acional hab ía ealizado algunas i e aciones. Es o implica que si el hilo compu acio- nal, el eó icamen e más len o, in oca a a es a llamada, se pe de ía un mínimo de 5 i e aciones, lo cual pa ecía poco azonable, pues o que a lo mejo el hilo de comunicaciones ya había e minado an es. Po es as p ime as p uebas, ambién se desca ó es a opción. Desca adas es as opciones, se pensó en u iliza los ce ojos de OpenMP (omp_locks). 4.5.2 Implemen ación inal con ce ojos de OpenMP (omp_locks) La e sión 27 inal ecoge las ideas y la expe iencia de la an e io 26, pe o la aplica a pa i de la 24. Añade a es a e sión dos ce ojos, con los cuales se ges iona el acceso a las i e aciones en que se ealiza el cómpu o es adís ico, y las inmedia amen e siguien es a ella. Se a a pues de una ges ión de g ado más g ueso, pues o que ya no se p e ende ges iona el caso de cada una de las es adís icas po sepa ado, sino que simplemen e se ac úa sob e oda la i e ación, lo cual incluso iene cie o sen ido compu acional po la exis encia de dependencia de da os en e las di e en es es adís icas de una misma i e ación. Los equisi os son los mismos: se c ean pe manen emen e zonas de memo ia auxilia (bu e s) pa a al- macena los esul ados es adís icos de una i e ación, de al o ma que en la siguien e los da os o iginales puedan se sob eesc i os sin p oblema. De paso, oda la ges ión de memo ia dinámica se pasa a es á ica pa- a p ác icamen e odas las a iables, dado que se conside a ambién muy cos oso el es a cons an emen e ese ando, inicializando y libe ando memo ia en cada i e ación es adís ica. Los ce ojos en OpenMP pe mi en bloquea en una espe a no ac i a al bloque que in en a oma el ce ojo pe o és e ya es á omado po o o hilo. Sólo cuando el o o hilo lo libe a, el sis ema ope a i o a isa el hilo bloqueado pa a que és e pueda coge el ce ojo y segui su camino. Se a a pues de una ges ión p opia de OpenMP, y po lo an o op imizada pa a la a qui ec u a subyacen e y eó icamen e más e icien e que el in en o ealizado con la e sión 26 an e io men e explicada. 36 Capí ulo 4. Implemen ación de la pa alelización con CUDA Sección 4.5 Pa a el caso que nos ocupa, la solución plan eada ha sido la u ilización de dos ce ojos di e en es, así como dos a iables auxilia es de seguimien o de quién es su p opie a io (es impo an e, dado que es á p ohibido libe a un ce ojo que no se posea, ni ol e a coge un ce ojo que ya se posee, en ambos casos, la ejecución e mina ab up amen e con e o ). La idea básica es que inicialmen e el hilo 0 posee ambos ce ojos, y siemp e end á al menos uno de ellos. Nunca se le es a á pe mi ido libe a los dos ce ojos. Po su pa e, el hilo 1 pod á ene o ninguno, o sólo uno de los dos ce ojos. Pod án habe ce ojos lib es, sin ningún p opie a io, que siguiendo las eglas an e io es, y o as po de ini , pod án se omados po cualquie a de los dos hilos, si se le pe mi e. Depen- diendo de qué ce ojo o ce ojos poseea cada hilo, pod á a anza en su a ea, o debe á queda se bloqueado en espe a de un cambio de si uación, es o es, coge uno de los o os ce ojos. Cada ez que se coge el ce o- jo, el nue o p opie a io ano a en la a iable auxilia co espondien e que él es el p opie a io de ese ce ojo. Cada ez que se libe a, esa a iable se pone a -1, indicando que no iene p opie a io. Así, cualquie hilo puede en cualquie momen o sabe quién es el p opie a io del ce ojo, o si és e es á lib e. Los hilos 0 y 1 ienen asignadas las mismas unciones que las desc i as en el caso an e io de 26. El esquema de allado de uncionamien o es el siguien e: 1. La ejecución empieza con el hilo 0 poseyendo ambos ce ojos, que denomina emos ab e iadamen e WORK yWAIT. 2. Al inicio de la i e ación de gene ación de es adís icas, el hilo 0 comp ueba po las a iables auxilia- es si es el p opie a io de WORK. Si no lo es, in en a coge lo, quedando bloqueado si no lo consigue (llamada a omp_se _lock()). 3. Si ya e a p opie a io de WORK, en onces in en a coge WAIT, de al o ma que si alla, ambién se queda bloqueado has a ob ene lo. 4. Sólo cuando el hilo 0 posee ambos ce ojos, con inua la ejecución de la i e ación de es adís icas. Es o es, p ecisa ene los dos ce ojos pa a elabo a es adís icas. 5. No obs an e lo an e io , cuando ha ob enido los dos ce ojos, y a a con inua la ejecución, lo p ime o que hace es libe a el ce ojo de WAIT. 6. E en ualmen e el hilo 0 e mina el ciclo de gene ación de las es adís icas, que las gua da en las a ia- bles auxilia es al e ec o, y llega a la siguien e i e ación, la de es adís icas+1. 7. Cuando el hilo 0 llega al inicio de la i e ación es adís icas+1 (o e mina el bucle, que es equi alen e), el hilo 0 in en a coge el WAIT. Si no puede coge lo, queda bloqueado has a consegui lo. Si lo consigue, o cuando, as es a bloqueado, lo consigue, es po que el hilo 1 ya lo ha omado y lo ha de uel o, y po lo an o, si e pa a que el hilo 0 ob enga el econocimien o del hilo 1 de que és e sabe que el hilo 0 es á p epa ando los nue os da os. 8. El hilo 0, as ob ene de nue o el WAIT, libe a el WORK, con lo cual ha espe ado en odo momen o ene al menos uno de los dos ce ojos en su posesión. 9. Pasemos al hilo 1. Cuando és e llega a la i e ación de gene ación de es adís icas, ope a de o ma di e- en e al hilo 0. P ime o mi a si él es el poseedo de alguno de los dos ce ojos, WORK oWAIT, si es así, se ha p oducido un e o inespe ado en el p o ocolo, y di ec amen e abo a la ejecución, e minando anómalamen e el p og ama. 37 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e 10. Si sigue i o (no enía ninguno de los dos ce ojos), el hilo 1 in en a ob ene el WAIT, o se queda blo- queado en espe a de él. 11. Al consegui el WAIT se en e a de que el hilo 1 ya ha llegado ambién al hilo de gene ación de es adís- icas, así que lo de uel e inmedia amen e como econocimen o de habe se en e ado. 12. Inmedia amen e, in en a coge el ce ojo de WORK, dado que el hilo 0 debe es a en esos momen os gene ando las es adís icas, y no lo libe a á has a que el hilo 0 no e mine esa i e ación de abajo es adís ico, y al inicio de la i e ación de es adís icas+1, ecoja el WAIT pa a sabe que el hilo 1 ya es á p epa ado, y de uel a el WORK como señal al hilo 1 de que los da os ya es án gene ados y en si io segu o pa a pode se ges ionados po el hilo 1. 13. El hilo 1, cuando oma el WORK, se desbloquea, y con inua con su i e ación, que se á la de es adís icas acabas de gene a po el hilo 0. 14. Cuando el hilo 1 e mina su i e ación de in e cambio y esumen de es adís icas, y almacenamien o en disco de ellas, si p ocedía, llega al inicio de la i e ación de es adís icas+1. 15. Al inicio de la i e ación es adís icas+1, el hilo 1 de uel e el ce ojo de WORK como señal de que ya ha e minado su abajo. 16. Consecuen emen e, el hilo 1 sólo ha enido como máximo un único ce ojo en su posesión, el WAIT pa a econoce la si uación de abajo del hilo 0, o el WORK cuando es á ges ionando los da os es a- dís icos gene ados po el hilo 0. Cuando es á i e ando sin hace nada has a la siguien e i e ación de es adís icas, o bloqueado al inicio de és a, no posee ningún ce ojo. Con es e p o ocolo se ha implemen ado la e sión 27, que como se e á en el Capí ulo 5 “Resul ados”, es más ápida que la e sión de pa ida 24 en ap oximadamen e un 3–4% en el peo de los casos, excep o en un caso pun ual que se explica á en su momen o, y que pudie a ene elación con el en o no de ejecución. A con inuación se incluye el código al inicio de la i e ación p incipal del bucle, donde se ges ionan es os ce ojos. El Código 4.9 es á ampliamen e comen ado pa a segui el p o ocolo an e io . Lis ado 4.9: Ges ión de ce ojos pa a pe mi i el co ec o solapamien o de cálculo y comunicación 1DO jk = 1, ni e - 1 ! Main loop o empo al i e a ions 2IF(l_p oduc ion) THEN ! Check i he s a is ics a e ac i e 3i ( h ead==0) hen ! Th ead 0 does he compu a ional asks 4i (np od*in (jk/np od) == jk .o . jk==ni e -1 .o . 5& ncon ol*in (jk/ncon ol) == jk) hen 6! I e a ion when he special s a is ical calcula ions mus be done 7! Mas e ( h ead==0) canno s a i h ead==1 has no inished 8! messaging and s o ing he p e ious s a s 9! The h ead==0 (mas e ) has o ha e bo h locks o con inue 10 ! I always ha e one a leas , don’ need o check i ha e none 11 ! Then, when he has go bo h, he inmedia ly elease lock_wai 12 i (lock_wo k_owne .ne. h ead) hen 13 call omp_se _lock(lock_wo k) 38 Capí ulo 4. Implemen ación de la pa alelización con CUDA Sección 4.5 14 lock_wo k_owne = h ead 15 i e _wo k = jk 16 elsei (lock_wai _owne .ne. h ead) hen 17 call omp_se _lock(lock_wai ) 18 lock_wai _owne = h ead 19 i e _wai = jk 20 endi 21 call omp_unse _lock(lock_wai ) 22 lock_wai _owne = -1 23 elsei ((jk .ne. 1) .and. (np od*in (jk/np od)==jk-1 .o . 24 & ncon ol*in (jk/ncon ol)==jk-1) ) hen 25 ! I e a ion a e he special s a s ha e been done 26 ! The mas e ( h ead==0) has only he lwo k and wai s o lwai 27 ! When i ge s lwa , hen i eleases lwo k so i can be aken 28 ! by h ead==1 who will message and/o s o e ha in o 29 call omp_se _lock(lock_wai ) 30 lock_wai _owne = h ead 31 i e _wai = jk 32 call omp_unse _lock(lock_wo k) 33 lock_wo k_owne = -1 34 endi ! end o checking i e a ion numbe o h ead==0 35 else ! I ’s h ead 1, which does he MPI and I/O pa 36 i (np od*in (jk/np od) == jk .o . 37 & ncon ol*in (jk/ncon ol) == jk) hen 38 ! I e a ion when he special s a is ical calcula ions mus be done 39 ! Th ead==1 canno s a i mas e ( h ead==0) has no ye inished 40 ! his e y same i e a ion, because needs o be su e he da a is OK. 41 ! Th ead==1 ha e no lock, we check ha , and elease hem i so . 42 ! Then h ead==1 ies o ge lock_wai be o e con inue he mes &IO wo k. 43 ! When i ge s i , eleases i inmedia ly and hen ies o ge 44 ! he lock_wo k 45 i (lock_wo k_owne == h ead .o . 46 & lock_wai _owne == h ead) hen 47 p in *,’ h ead ’, h ead,’ha e one lock a i =’,jk, 48 &’Tha is no allowed. STOPING!!!’ 49 STOP ! No allowed, some hing wen w ong. Abno mal e mina ion 50 endi 51 call omp_se _lock(lock_wai ) 52 lock_wai _owne = h ead 53 i e _wai = jk 54 call omp_unse _lock(lock_wai ) 55 lock_wai _owne = -1 39 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e 56 call omp_se _lock(lock_wo k) 57 lock_wo k_owne = h ead 58 i e _wo k = jk 59 elsei ((jk .ne. 1) .and. (np od*in (jk/np od)==jk-1 .o . 60 & (ncon ol*in (jk/ncon ol)==jk-1 .o . jk==ni e -1))) 61 & hen 62 ! I e a ion a e he special s a s ha e been done 63 ! The h ead==1 has he lock_wo k, and eleases i 64 ! o le know he mas e i has al eady messaged and 65 ! s o ed p e ious s a s 66 call omp_unse _lock(lock_wo k) 67 lock_wo k_owne = -1 68 endi ! end o checking i e a ion numbe o h ead==1 69 endi ! i ( h ead==0) hen 70 ENDIF ! IF(l_p oduc ion) THEN 71 ! End o lock managemen , ne e make his egion OMP CRITICAL o will deadlock 4.5.2.1 Conside aciones con la adición de OpenMP en el código Fo an Al u iliza hilos de OpenMP se pueden de ini qué a iables an a se locales y cuáles an a se compa idas en e ambos hilos pa icipan es de la pa alelización de la egión dada. No obs an e, hay un caso especial y es cuando desde Fo an se in oca a una sub u ina. En ese momen o, las a iables que has a aho a e an globales, desapa ecen como ales, y pasan ealmen e a se locales. Es más, se einicializan a su alo po de ec o (gene almen e ce o), sin ni siquie a conse a el alo que la a iable global enía en el momen o de in oca se la sub u ina. Ese compo amien o no es así si la c eación de los hilos es den o de la sub u ina. Pa a sol en a ese p oblema se hace necesa io, pues, pasa como pa áme os en la in ocación a la su- b u ina de odas aquellas a iables que el hilo p ecise conoce pa a la ejecución co ec a de la sub u ina in ocada. Es o hace que sea necesa ia la modi icación del p o o ipo de casi odas las sub u inas in ocadas desde el p og ama p incipal, pues es allí donde se c ean los hilos. 4.6 Valo ación cuali a i a del p oceso de pa alelización con CUDA La ealización de la mig ación paso a paso ha pe mi ido e alua en cada momen o cuáles han sido los p incipales sal os cuali a i os y cuan i a i os ob enidos du an e el p oceso, así como ene que en en a se a decisiones pa a sol en a aquellos p oblemas que han ido apa eciendo. T as es e p oces, se p esen an de o ma esumida las p incipales ca ac e ís icas que se conside an in e esan es: • La mig ación a CUDA ha esul ado se , en lo que se e ie e a las unciones indi iduales de cómpu o, ela i amen e sencilla. • Dicha mig ación sí que ha p esen ado mayo es p oblemas a la ho a de de ini la es a egia de mig a- ción. • También ha esul ado se p oblemá ico analiza el algo i mo en el código implemen ado, pues o que en ocasiones és e no e a el más e iden e ni el más óp imo. 40 Capí ulo 4. Implemen ación de la pa alelización con CUDA Sección 4.6 • Se han enido que esol e p oblemas de p og amación in e esan es en CUDA-C, ales como la u ili- zación e icien e de los gene ado es de núme os alea o ios, la u ilización y p og amación de di e en es ipos de educción, y la op imización en la in ocación a los ke nels. • El mé odo de los w appe s, jun o con los iche os de cabece a pa a accede más ácilmen e a las a- iables globales decla adas en Fo an ha esul ado e icien e. No obs an e, apo a un g ado de com- plejidad que puede ep esen a un escalón impo an e pa a p og amado es no amilia izados en C. • El análisis del código, jun o con la comp ensión de la pa e ísica de las ecuaciones de Lange in con ellas implemen ada, ha pe mi ido la de ección de dos e o es de p og amación que se han co egido en la implemen ación p opues a, así como de la p opues a e implemen ación de ella de algo i mos y uso de unciones más e icien es que los mé odos numé icos u ilizados en el código o iginal. • La inclusión de OpenMP en el código ya pa alelizado, pa a ob ene una acele ación aún mayo , ha esul ado se mucho más complejo que los esul ados con él ob enidos. E iden emen e es o depen- de del caso conc e o, pe o el g ado de complejidad, no solo en p og amación, sino ambién en la necesidad de idea algo i mos co ec os lib es de bloqueo, hace que es a úl ima pa e no sea eco- mendable in en a la si no se iene a una pe sona muy expe a en compu ación pa alela y dis ibuida en el equipo. Se incluye la dis ibuida, po que el uncionamien o asínc ono de los hilos plan ea p o- blemas alejados de la compu ación pa alela, y mucho más p opios de la dis ibuida, con su necesidad de es udio de co ección y i eza de los algo i mos que se p opongan implemen a . • Desde el pun o de is a cuan i a i o, la mig ación es un o al éxi o, dado que es co ec a en sus esul a- dos, y apo a una acele ación en o no a 700 espec o al código o iginal ejecu ándose en un núcleo de CPU. Económicamen e, si las simulaciones son muchas, implica un impo an e aho o en el iempo de uso del clus e de p oducción u ilizado, y con ello, ambién económico. 41 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e (Página in encionadamen e en blanco) 42 CA P Í T U L O 5 RESULTADOS En es e capí ulo se p esen an los esul ados ob enidos en di e en es en o nos con la ejecución del p og ama implemen ado, o de algunas de sus e siones p e ias, si ya e a su icien e pa a saca conclusiones álidas. Las p uebas se han ealizado en a ios en o nos de ejecución di e en es, po lo que se á in e esan e p esen a una b e e eseña de los mismos con an elación a la p esen ación de dichos esul ados. Finalmen e, se e ec ua á un análisis de los esul ados, que pe mi i á alo a la implemen ación, en sus dos aspec os cuan i a i os más impo an es a alo a : • su co ección, es o es, si p oduce los esul ados que de él se espe a, • su e iciencia, y po lo an o, qué endimien o se ex ae de su u ilización, de o ma compa ada con o as implemen aciones exis en es. No obs an e, y eco dando que la pa alelización del código en conc e o u ilizado es una caso de p ueba y demos ación, se debe alo a cuali a i amen e el p oceso de pa alelización u ilizado, en compa ación con o as posibles al e na i as. Todo ello en is as a pode in o ma a la pa e cien í ica de los equisi os que debe con a el pe sonal que deba acome e en el u u o es a labo , pe sonal que ac ualmen e, en su mayo pa e, es á o mada po la p opia comunidad cien í ico eó ica del campo en cues ión, la ísica de pa ículas. 5.1 P uebas ealizadas Las p uebas se han ealizado en dos e apas di e en es: • inicio de la implemen ación, pa a la op imización del amaño de la malla y es ima ya desde las p i- me as e apas del abajo la acele ación mínima ob enible, • con las implemen aciones inales e minadas de las dos e siones del código, la pu amen e MPI con CUDA, y la que se añade OpenMP pa a desliga las ope aciones compu acionales de los ke nels de las de comunicación y esc i u a en disco. 43 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e A con inuación se de allan ambas e apas, y las azones y obje i os de las mismas: 5.1.1 P ime a e apa: op imización del ke nel p incipal según los pa áme os p opios de su in ocación En la p ime a e apa el abajo se ocalizó en comp oba si el p oyec o e a iable, o mejo dicho, si el código inicialmen e escogido e a pa alelizable con CUDA y suscep ible de ene una buena acele ación. Pa a ello se iden i icó el mínimo ozo de código que e a esponsable de una g an p opo ción del iempo o al de ejecución. En la Figu a 5.1 se ep esen a g á icamen e el diag ama de llamadas ob enido con el p o iling del p og ama o iginal. En él se de ec ó que la sub u ina ad ance, jun o con aquellas a las cuáles és a llama- ba ( hs_coe s,phi_and_psi, hs_ion yge _gaussian_noise) e a la consumido a del 85% del iempo de ejecución, al cual hay que añadi p ác icamen e o o 13% po la pa e co espondien e de las sub u inas es á icas de biblio eca hpso y sob e odo de an2, que son p incipalmen e u ilizadas po las sub u inas dependien es de ad ance. 84.54% 999000× 50.91% 1998000× 18.17% 1998000× 10.00% 1998000× 85.45% 1× 39.09% 1998000× MAIN__ 85.45% (0.91%) 1× ad ance_ 84.54% (5.45%) 999000× hs_coe s_ 50.91% (11.82%) 1998000× ge _gaussian_noise_ 18.18% (18.18%) 1999000× hs_ion_ 10.00% (10.00%) 1998000× main 85.45% (0.00%) phi_and_psi_ 39.09% (39.09%) 1998000× an2_ 13.64% (13.64%) hpso 2_ 0.91% (0.91%) Figu a 5.1: Diag ama de llamadas del p og ama o iginal ob enido del pe ilado Es po ello que la p ime a e apa de implemen ación consis ió en pa aleliza median e ke nels de CUDA la sub u ina ad ance y odas aquellas que son llamadas po és a. La ejecución del código p elimina así ob enido mos ó que e ec i amen e es a pa e del código e a la que más a ec aba al iempo de ejecución, al ob ene se ya un speedup en o no a 30–50, según el en o no de ejecución. T as op imiza dicho código, se lanzó una p ueba de op imización de los pa áme os de llamada a los 44 Capí ulo 5. Resul ados 5.1. P uebas ealizadas ke nels. Adicionalmen e a los ípicos pa áme os de en ada/salida de oda unción C/C++, los ke nels de CUDA se lanzan indicando en e dos y cua o pa áme os más, que se incluyen en la llamada con una no- ación especial que p ecisamen e ca ac e iza su código de llamada: Ke nel_Name<<< G idSize, BlockSize, SMEMSize, S eam >>> (a gumen s,....); Los alo es ele an es pa a nues o caso son BlockSize yG idSize, que de inen espec i amen e la can idad de bloques y cómo és os se es uc u an al lanza la ejecución del ke nel en la GPU. Aunque es os pa áme os son de un ipo especial de ec o de es dimensiones (x,y,z) de inido po el p opio CUDA, en nues o código no es necesa ia al complejidad y se u iliza á únicamen e las dimensiones xeyen el del amaño del bloque y xen el amaño de la malla. La o ma en que se in ocan los ke nels en la GPU es impo an e po las siguien e azones: • el núme o de bloques jun o con el amaño de cada bloque de e mina á el núme o de ejecuciones indi iduales que se ealiza á del código. Debe se el núme o su icien e pa a que acoja, en nues o caso, odas y cada una de las pa ículas a simula , pues o que cada hilo de ejecución se enca ga á de una de ellas en exclusi a, • el núme o de bloques (G idSize) iene que se su icien emen e ele ado como pa a pe mi i la máxi- ma u ilización de los mul ip ocesado es de la a je a g á ica, • el amaño del bloque (BlockSize) de e mina á un ac o muy impo an e pa a la e iciencia de la eje- cución pa alela en la a je a g á ica, que es el g ado de ocupación de los egis os, y con ello, el g ado de pa alelismo solapado exis en e en la ejecución. Téngase en cuen a que hay mucho cálculo de doble p ecisión, po lo que las unidades de coma lo an e es a án cons an emen e en u ilización, y con iene u iliza el iempo en que ese wa p es á inac i o en espe a del esul ado, en o os cálculos u ilizando egis os y o as unidades lógicas del co e. La p ueba consis e en a ia es os pa áme os, eniendo siemp e en cuen a que exis e un lími e máximo según la a qui ec u a de núme o de hilos que pueden ejecu a se concu en emen e en un bloque. Pa a ello, se a ían las dimensiones xeyde BlockSize, cambiándose el alo de G idSize pa a que se cumpla que se lancen los mínimos bloques necesa ios pa a que cada i e ación (pa ícula) sea ejecu ada una y al menos una ez po un único hilo. Es as p ebas se ealizan en di e sos en o nos educidos a una única máquina, dado que en es e caso, el esul ado no depende de la masi idad del cómpu o, sino de las ca ac e ís icas ha dwa e de la a je a g á ica u ilizada, es udiando pa a ello es di e en es gene aciones de a qui ec u a CUDA:Fe mi,Keple y Maxwell (p ime a gene ación). 5.1.2 Segunda e apa: medición de p es aciones de las e siones inales en un en o no HPC eal Una ez ob enidas las dimensiones óp imas pa a el lanzamien o del p incipal ke nel, el abajo se cen a en pa aleliza el es o de unciones que es án den o del bucle p incipal del p og ama. De es a o ma se ob iene un p og ama inal que ealiza p ác icamen e odos los cálculos sob e cada pa ícula u ilizando CUDA. Es o pe mi e e i a aspasos con inuos de memo ia en e la CPU y la GPU en cada i e ación empo al. Sólo se ealizan aquellas ans e encias de memo ia necesa ias, siemp e de GPU aCPU, pa a en ia los da os mínimos necesa ios pa a que la CPU pueda gua da los da os es adís icos y de p og eso de la simulación en disco (o p esen a los po consola). 45 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e También se ob end án conclusiones in e esan es sob e aspec os inicialmen e no plani icados, como po ejemplo, in luencia del lanzamien o de un núme o no di isible po 2 del amaño de los bloques, de ec- ción de posibles p oblemas con la u ilización de OpenMPI con In iniBand. Las p uebas se ealizan en el en o no de ejecución de MinoTau o. Se u iliza un núme o a iable de nodos, así como un núme o a iable de GPU. También se a ía el núme o de GPU po nodo que se u iliza en cada p ueba, pa a e si la mayo localidad de los p ocesos en menos nodos inc emen a la e iciencia de la ejecución de la simulación. Se seleccionan di e en es amaños de ca ga, a iando an o el núme o de i e aciones empo ales, como el núme o de pa ículas que son conside adas en cada i e ación. También se cambia el ipo de simulación (con gene ación o no de esul ados es adís icos in e medios, con ue e ca ga de in e cambio de mensajes y esc i u a en disco). Y inalmen e se u ilizan es p og amas: el código o iginal en Fo an+MPI; la nue a implemen ación sus i uyendo la pa e compu acional con CUDA, y algo de C pa a acili a la in ocación a la API; y la que mejo a és a úl ima con la adición de OpenMP pa a desacopla el cómpu o de la comunicación, y así solapa lo más posibles es as dos ac i idades. Todas es as p uebas se hacen de una o ma o ogonal (o ogonalidad no comple a, pe o sí su icien e pa a ex ae conclusiones adecuadas). De es a o ma se puede es udia la escalabilidad del p og ama bajo di e en es pun os de is a: espacial, empo al y localidad, así como e alua si la inclusión de CUDA p ime o y de OpenMP después compensan espec o de su espec i o iempo de desa ollo. A con inuación se p ocede a desc ibi , p esen a esul ados, analiza los esul ados, y ob ene conclu- siones sob e cada una de las p uebas ealizadas. 5.3.2.1 Ca ac e ización del compo amien o del p og ama o iginal Pa a pode e alua la nue a implemen ación, se necesi a hace lo po compa ación con el código o iginal de pa ida. En es a p ueba se i á pa a ca ac e iza el compo amien o del p og ama o iginal, ya que la nue a implemen ación debe á mejo a los aspec os más débiles del o iginal, sin que los aspec os ue es de és e su an me ma. La ca ac e ización consis i á en analiza su escalabilidad espacial y empo al. B e e desc ipción de las p uebas Consis e en la ejecución del código o iginal en el en o no de MinoTau o median e las siguien es p ue- bas: 1. Con una ca ga pequeña cons an e (10.000 i e aciones empo ales de 20.000 pa ículas), a ia el nú- me o de p ocesos que se ejecu an en un único nodo. Con es o se comp oba á la escalabilidad en un en o no pu o de memo ia compa ida. 2. Con la misma ca ga pequeña cons an e an e io , ejecu a siemp e con 12 p ocesos, pe o con un nú- me o a iable de nodos, a iando po lo an o el núme o de p ocesos po nodo. Así se pod á compa a la elocidad de in e comunicación en memo ia compa ida con la exis en e en e di e en es nodos (memo ia dis ibuida). 3. Con una ca ga g ande cons an e (500.000 i e aciones empo ales de 200.000 pa ículas), a ia el nú- me o de nodos, y po lo an o de p ocesos, haciendo siemp e que cada nodo abaje al 100% de su capacidad, es o es, con 2x6 =12 p ocesos po nodo. La idea es comp oba la escalabilidad del p o- g ama en un en o no mix o de memo ia compa ida y dis ibuida a g an escala. Los amaños son escogidos con dos ideas: 52 Capí ulo 5. Resul ados 5.3. Ejecución de las p uebas • La ca ga pequeña es pa a que las ejecuciones no sean excesi amen e la gas, sob e odo cuando se ejecu an en un único nodo. • La ca ga mayo es á elacionada con la memo ia máxima que una única GPU del clus e puede so- po a , y el núme o de i e aciones es uno ya su icien emen e ele ado como pa a que el enómeno del pla eau (mese a) en la o mación de los elec ones ugi i os sea cla o, se ea cla amen e es abilizado, en es ado es aciona io. Resul ados ob enidos Las p uebas 1 y 2 se ealizan en un único nodo, pa a e la escalabilidad en memo ia compa ida. En la Tabla 5.4 se mues an los esul ados ob enidos. P ocesos Nodos P ocesos Tiempo de Acele ación E iciencia po nodo ejecución 1 1 1 2045,806 s 1,000 1,000 2 1 2 1021,929 s 2,002 1,001 3 1 3 683,256 s 2,994 0,998 4 1 4 511,999 s 3,996 0,999 6 1 6 343,132 s 5,962 0,994 9 1 9 228,449 s 8,955 0,995 12 1 12 170,748 s 11,981 0,998 12 2 6 171,560 s 11,925 0,994 12 3 4 173,361 s 11,801 0,983 12 4 3 171,696 s 11,915 0,993 12 6 2 172,575 s 11,855 0,988 12 12 1 170,895 s 11,971 0,998 Tabla 5.4: Acele ación y e iciencia del código o iginal de CPU (memo ia compa ida y dis ibuida) En la G á ica 5.2, en ep esen ación doble loga í mica, se ep esen an los da os ob enidos de la P ueba 3, con g an can idad de p ocesos y una ca ga g ande. Se obse a la disminución del iempo de ejecución con o me se inc emen a el núme o de nodos u ilizados en la compu acion de o ma muy p opo cional al núme o de p ocesos pa icipan es. Fenómenos obse ados En la p uebas 1, sob e un único nodo, es o es, sólo con memo ia compa ida, al como se e en la Ta- bla 5.4 en las ilas somb eadas de ama illo, la escalabilidad es p ác icamen e lineal en odo el in e alo. Lo mismo es cie o pa a cuando se cambia p og esi amen e a memo ia dis ibuida ( ilas somb eadas de azul), donde igualmen e la e iciencia es p ác icamen e la unidad en odos los casos, nunca bajando del 98%. En la p ueba 3, mezcla de memo ia compa ida y dis ibuida, Figu as 5.2 y 5.3, se obse a que la línea es casi ec a, con una pendien e de p ác icamen e la unidad, lo que adelan a que la acele ación se á bas an e lineal. Ese hecho se comp ueba en la G á ica 5.3 que isualiza que an o la acele ación como la e iciencia se ap oximan a la idealidad (e iciencia muy ce cana a 1 incluso con 64 nodos, que son 768 p ocesos, y la línea muy coinciden e con la ideal). Aún así, con 64 nodos (768 p ocesos MPI) la e iciencia sigue siendo muy al a, p ác icamen e del 95%. 53 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e 1; 66988 4; 16910 16; 4382 32; 2142 64; 1105 1 10 100 1000 10000 100000 1 10 100 Núme o de nodos (cada nodo ejecu a 12 p ocesos en sendos núcleos) Tiempo (s) Tiempo ejecución o al Figu a 5.2: Escalabilidad empo al del código o iginal de CPU (500.000 pa ículas y 200.000 i e aciones) 16; 15.287 32; 31.276 64; 60.615 4; 3.962 1; 1.000 0.977 0.947 0.990 1.000 0.955 0 10 20 30 40 50 60 70 0 8 16 24 32 40 48 56 64 72 Núme o de nodos (cada nodo ejecu a 12 p ocesos en sendos núcleos) Acele ación (-) 0.880 0.900 0.920 0.940 0.960 0.980 1.000 1.020 E iciencia (-) Acele ación Escalabilidad lineal pe ec a E iciencia Figu a 5.3: Acele ación y e iciencia del código o iginal de CPU (500.000 pa ículas y 200.000 i e aciones) Conclusiones de las p uebas A aíz de es as es p uebas, se ob ienen las siguien es conclusiones espec o al p og ama o iginal en Fo an pa alelizado con MPI ejecu ándose sólo en CPU y ambién sob e el en o no de ejecución de Mino- Tau o. • El p og ama o iginal escala muy bien en memo ia compa ida, pe o al se an len o, no es su icien e pa a simulaciones que pe mi an ex ae conclusiones ísicas in e esan es. • El p og ama o iginal escala muy bien en un clus e HPC, lo que has a aho a ha pe mi ido ealiza simulaciones su icien e pa a ex ae conclusiones. • La ed de comunicaciones del clus e MinoTau o es muy ápida. No hay di e encia sus ancial en e el uso de memo ia compa ida y la memo ia dis ibuida. No obs an e, se obse a que pa a pode alcanza amaños de simulación in e esan es, se necesi ang an- des ecu sos compu acionales du an e un iempo nada desp eciable. Como se e á, la pa alelización con CUDA pe mi i á educi d ás icamen e es a elación en e iempo necesa io y ecu sos ha dwa e u ilizados, lo que implica á o un impo an e aho o de explo ación, y la posibilidad de lanza simulaciones mucho más g andes que pe mi an ob ene esul ados más conclusi os, o disminui la can idad de ap oximaciones ea- lizadas pa a alcanza un conjun o de ecuaciones más complicadas de compu a , pe o más ieles espec o del compo amien o espe ado del plasma de pa ículas simulado. 5.3.2.2 Ca ac e ización básica de la implemen ación consis en e en la pa alelización con CUDA A con inuación se a a p ocede a es udia el compo amien o de la e sión inal del p og ama en el que, basándose en el código o iginal en Fo an+MPI ejecu able en CPU, se ealiza una pa alelización de oda la pa e de compu ación numé ica y del cálculo de las es adís icas in e medias median e CUDA. De es a o ma, Fo an queda sólo como sopo e básico pa a la lec u a y esc i u a de iche os y la in ocación a la lib e ía MPI pa a el in e cambio de da os en e los di e en es p ocesos, así como ambién del lujo de con ol básico del p og ama, incluyendo la ges ión de las i e aciones empo ales, y de cuándo co esponde ealiza las es adís icas. La implemen ación ha sido al que an sólo es necesa ia la ans e encia de los da os de elocidades de las pa ículas al p incipio y al inal del p og ama. Du an e odas las i e aciones empo ales, la ans e encia de da os se ealiza únicamen e de la GPU a la CPU solamen e en aquellas i e aciones en las que hay que 54 Capí ulo 5. Resul ados 5.3. Ejecución de las p uebas ecopila y gua da en disco las es adís icas in e medias solici adas, y sólo de los mínimos da os necesa ios. De esa o ma se ha conseguido disminui mucho la ans e encia de memo ia en e la GPU y la CPU, lo que pe mi e inc emen a el po cen aje de código que el p og ama es á ealizando cómpu os, y po lo an o, po la Ley de Amdahl, pe mi i á un mucho mejo ap o echamien o del pode de la pa alelización, como se obse a á a a és de los esul ados que a con inuación se p esen an y discu en. Al igual que en la sección an e io , p ime o se de alla án las p uebas ealizadas, después se p esen a án los da os cuan i a i os esumidos ob enidos de dichas p uebas, que pe mi i án obse a las p incipales ca ac e ís icas de la ejecución, y con ello, ca ac e iza la implemen ación, a pa i de lo cual se ex ae án conclusiones sob e ella. 5.3.2.2.1 P ueba de co ección Du an e odo el p oceso de implemen ación se ha ido comp obando que los esul ados de cada pa e que se implemen aba no cambiaban los esul ados in e medios es adís i- cos que gene a el p og ama. En ealidad, los mismos esul ados no pueden gene a se, dado que se cambió el gene ado de núme os alea o ios. Sin emba go, o zando a que unos pocos núme os sean los mismos se comp obó que las elocidades y su clasi icación en con enedo es ue an iguales a los del código o iginal. De hecho, du an e es a e isión se localiza on a ios posibles bugs en el código o iginal que ue on pues os en conocimien o de sus esponsables. A pa i de ese momen o, los núme os gene ados alea o iamen e se ija on pa a la GPU has a la e sión inal, y aunque al p ocesa con di e en e núme o de GPU esos alo es uel en a cambia , sí que se obse a que la dis ibución de los alo es en cada columna en los iche os gene ada es es adís icamen e simila a los del código o iginal. Además, isualizando los iche os .silo con VisI , se obse a que los mapas de elocidades y su dis i- bución son simila es a los que se mues an en [4]. Es po odo ello que se conside a que las implemen aciones p esen adas son co ec as en cuan o a la idelidad de la esolución de las ecuaciones ísicas que se usan pa a simula el plasma con inado magné i- camen e. 5.3.2.2.2 Acele ación y e iciencia espec o del código en CPU B e e desc ipción de las p ueba La p ueba consis e en ejecu a la misma ca ga que se u ilizó en la p ueba de la ca ac e ización de la CPU, cuyos da os se han mos ado g á icamen e en la Figu as 5.2 y 5.3, es o es, 200.000 i e aciones empo ales sob e 500.000 pa ículas. Se ealizan a ias simulaciones a iando dos pa áme os: • El núme o de p ocesos, con cada p oceso u ilizando una GPU. Se a ía desde 1 nodo con 1 GPU (caso base) has a 128 p ocesos, u ilizando 128 GPU. • El núme o de GPU que se u ilizan en cada nodo. Las mismas p uebas se ealiza án u ilizando una GPU po nodo, o dos. Es o úl imo pe mi i á es ablece si hay alguna di e encia impo an e en e la ejecución den o de un mismo nodo o con un mayo núme o de nodos, y po lo an o, con una mayo u ilización de la ed de in e conexión. Resul ados ob enidos 55 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e En la Tabla 5.5 se mues an los da os de es a p ueba. Los expe imen os se han lanzado po duplicado, siendo muy coinciden es en odos los casos pa ejos, po lo que en la abla se mues a el alo p omedio medido. P ocesos Nodos GPU nodo Tiempo de ejecución Acele ación E iciencia Acele ación Acele ación To al GPU I/O MPI O os GPU/CPU ela i izada 1 1 1 1181,871 1177,111 4,559 0,057 0,144 1,000 1,000 680 680 2 2 1 600,646 597,715 2,596 0,386 −0,051 1,968 0,984 1338 669 4 4 1 317,959 315,880 1,631 0,525 −0,077 3,717 0,929 2528 632 8 8 1 167,233 165,699 1,141 0,515 −0,122 7,067 0,883 4807 601 16 16 1 93,997 92,619 0,901 0,583 −0,106 12,573 0,786 8552 534 32 32 1 55,610 54,437 0,792 0,565 −0,184 21,253 0,664 14455 452 64 64 1 44,437 38,298 0,716 5,525 −0,101 26,597 0,416 18090 283 2 1 2 601,306 598,022 2,780 0,512 −0,007 1,966 0,983 1337 668 4 2 2 317,888 315,954 1,625 0,370 −0,060 3,718 0,929 2529 632 8 4 2 167,683 166,181 1,151 0,431 −0,079 7,048 0,881 4794 599 16 8 2 93,997 92,619 0,901 0,583 −0,106 12,573 0,786 8552 534 32 16 2 55,610 54,437 0,792 0,565 −0,184 21,253 0,664 14455 452 64 32 2 44,437 38,298 0,716 5,525 −0,101 26,597 0,416 18090 283 128 64 2 26,550 25,318 0,656 1,086 −0,509 44,515 0,348 30277 237 Tabla 5.5: Acele ación y e iciencia de la implemen ación Fo an+MPI+CUDA con misma ca ga que Tabla 5.4 Pa a el cálculo de la acele ación sob e la CPU, y dado que el código o iginal escala linealmen e de o ma pe ec a, se ha omado el iempo o al de ejecución en un nodo comple o, co igiéndose el hecho de que ese iempo se co esponde con 12 p ocesos. Es os da os se isualizan en las G á icas 5.4. 0 5000 10000 15000 20000 25000 30000 35000 40000 0 16 32 48 64 80 96 112 128 Núme o de p ocesos (gpu) Acele ación espec o 1 CPU (-) 0 100 200 300 400 500 600 700 800 Acele ación ela i izada (-) Acele ación 2GPU/nodo espec o CPU Acele ación 1GPU/nodo espec o CPU Acele ación 2GPU/nodo ela i izada Acele ación 1GPU/nodo ela i izada (a) Compa ación con el código o iginal en CPU 0 20 40 60 80 100 120 140 0 16 32 48 64 80 96 112 128 Núme o de p ocesos (gpu) Acele ación en e GPU (-) 0.0 0.2 0.4 0.6 0.8 1.0 1.2 1.4 E iciencia (-) Acel. 2GPU/n CPU Acel. 1GPU/n sob e CPU Escalabilidad lineal pe ec a Acel. 2GPU/n ela . Acel. 1GPU/n ela . E iciencia unidad (b) Compe ación en e las p opias e siones en GPU Figu a 5.4: Acele ación y e iciencia de implemen ación Fo an+MPI+CUDA con misma ca ga que Tabla 5.4 Fenómenos obse ados En la Tabla 5.5 se obse an los siguien es enómenos: • La acele ación de una única GPU en e a un único núcleo de CPU alcanza el ele ado alo de 680. Es o quie e deci que una única GPU ealiza el mismo abajo p ác icamen e que 32 nodos bip ocesado como los exis en es en MinoTau o. 56 Capí ulo 5. Resul ados 5.3. Ejecución de las p uebas • Con o me se u ilizan más GPU la acele ación inc emen a, sin no a se que llegue un momen o de desacele ación. • Con o me se u ilizan más GPU la acele ación ela i a dec emen a. Es o es, añadi el doble de nodos con GPU al mismo abajo no disminuye a la mi ad el iempo global de ejecución. Es o es, la e iciencia se obse a que baja, has a llega a un 35% al usa 128 nodos. Es e enómeno es el espe able, dado que el man ene se cons an e abajo o al, al inc emen a se el núme o de GPU, disminuye su ca ga indi idual, po lo que en gene al, al disminui la ca ga compu acional po nodo, es espe able cie o dec emen o de p es aciones. En odo caso, en la p ác ica, es una si uación que no suele da se, pues o que el lanzamien o se suele ya ealiza sob e los mínimos p ocesos necesa ios pa a e i a el cos e de comunicaciones y ap o echa la localidad de la ca ga. • Al igual que en p uebas an e io es, la misma ca ga epa ida en e el doble de nodos que abajan con la mi ad de sus ecu sos (sólo 1 p ocesado CPU y sólo 1 a je a g á ica) iene p ác icamen e el mismo iempo de ejecución que si se ejecu an de o ma más concen ada en la mi ad de nodos a ope de capacidad (2 p ocesado es CPU con 2 GPU). • So p enden emen e, cuando se u ilizan 64 nodos en MinoTau o, en odas las ejecuciones ealizadas en es a p ueba, el iempo pa cial empleado po la aplicación en comunicaciones MPI se inc emen a de mane a súbi a de 0,5 segundos a 5,5 segundos. Cuando se u ilizan 128 nodos, uel e a baja a un alo jus i icable en e 0,5 y 1,0 segundos. • Los iempos pa ciales de esc i u a (I/O) dec emen an con o me aumen a el núme o de nodos. Es o es así po las ca ac e ís icas del sis ema de iche os dis ibuido GPFS que usa MinoTau o. Conclusiones de las p uebas De las obse aciones an e io es se concluye que: • La implemen ación en GPU del mé odo ap oxima i o de Lange in es muy e icien e compa ada con la o iginal en CPU+MPI, con una acele ación de 680. • Es a p ime a implemen ación con CUDA no escala bien cuando el núme o de GPU inc emen a más allá de 16, aunque sigue acele ándose el cómpu o. Es e hecho es el que impulsa la siguien e e sión añadiendo OpenMP pa a desacopla más el cálculo de la comunicación y la esc i u a en disco. • Se con i ma que la ed de in e conexión MinoTau o es excelen e. No obs an e, hay algunas combina- ciones de núme o de nodos, quizá ligada a la opología esul an e y p obablemen e ambién a algún p oblema en e OpenMPI eIn iniBand, que hace que el iempo de comunicación po MPI se dispa e, con la consiguien e pé dida de endimien o global. 5.3.2.3 Ca ac e ización básica de la implemen ación con CUDA yOpenMP Consis e en las mismas p uebas y me odología que en el an e io apa ado 5.3.2.2. Los esul ados son muy simila es a la e sión an e io , sólo que un 3–5% más ápidos, según la ca ga. La acele ación máxima ob enida en es e caso es de 695 sob e el código o iginal en CPU. Po mo de la b e edad, se omi e la abla y las g á icas de esul ados, al no apo a mayo in o mación, y se edundan es con las de la p ime a implemen ación po su semejanza. 57 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e 5.3.2.4 P uebas masi as de escalabilidad de las implemen aciones con CUDA T as obse a el uncionamien o básico de la nue a implemen ación, y siendo és a co ec a y ob eniendo esul ados de acele ación muy impo an es espec o al código o iginal en CPU, se p ocede a ealiza a ias p uebas masi as a ni el de p oducción de los dos códigos p esen ados, al que sólo se le ha inco po ado CUDA ( e e enciada como p og ama o e sión 24, con Fo an+MPI+CUDA) y al que añade a la an e io OpenMP pa a ges iona dos hilos en cada p oceso, uno enca gado de la compu ación numé ica en la GPU p incipalmen e y o o esponsable de las comunicaciones y la esc i u a en disco de los esul ados ( e e en- ciada como 27). B e e desc ipción de las p uebas Las p uebas se ealizan en el en o no de ejecución de MinoTau o. Se ealizan di e en es lanzamien os a iando la ca ga, pa a lo cual básicamen e se man iene cons an e el núme o de i e aciones empo ales, ijándolas en 200.000 i e aciones (dado que llega un momen o en que se alcanza el es ado es aciona io y no iene sen ido ísico alguno con inua la simulación), y a ia el núme o de pa ículas, desde muy pocas a muchas. Es impo an e indica que exis e un lími e pa a el núme o de pa ículas que se pueden lanza en una única GPU. La a je a g á ica NVIDIA M2090, que dispone de 6GB de memo ia RAM, y con la ac ual imple- men ación, sólo pe mi e un máximo de unas 825.000 pa ículas. Pa a abaja con un poco de ma gen, se escoge lanza simulaciones con 500.000 u 800.000 pa ículas po cada GPU. P esen an las siguien es ca ac e ís icas, lanzadas de o ma o ogonal en e odas ellas: • Ca ga de en e 1 millón y 50 millones de pa ículas, siemp e con 200.000 i e aciones empo ales. • Se a ía el núme o de GPU, desde el mínimo núme o posible po la capacidad en memo ia de las a je as g á icas, que coincide con unas 825.000 pa ículas/GPU, has a el máximo disponible en el clus e po polí ica de u ilización, que es de un máximo de 128 GPU. • Se p ueban las dos e siones 24 (con CUDA) y 27 (con CUDA más OpenMP). Es as p uebas se componen de o ma o ogonal, ba iendo odas las posibilidades combina i as, pe- o po mo de la b e edad, sólo se mos a án las más ele an es, sob e odo eniendo en cuen a que se comp ueba que alguna de las a iables inicialmen e p e is a no iene en la p ác ica, en es e en o no de ejecución, ninguna in luencia es adís ica. 5.3.2.4.1 Resul ados ob enidos Una impo an e obse ación as las p uebas, y conco dan e con lo is o con an e io idad, es que da p ác icamen e lo mismo lanza Xp ocesos en Xnodos (u ilizando sólo 1 GPU po nodo), que lanza esos mismos Xp ocesos en X 2nodos (u ilizando 2 GPU po nodo). Es po ello que en adelan e sólo se mos a án los da os co espondien es a la u ilización de 2 GPU po nodo, excep o en el caso de las p uebas con 128 nodos, que po polí ica del clus e , sólo podían se lanzadas en 64 nodos. En las Tablas 5.6 y 5.7 se mues an los esul ados de los iempos de ejecución de las di e en es p uebas. Las p uebas se ealizan po duplicado, epi iéndose si se obse a algún alo ex año, que gene almen e se p esen a como un alo ex añamen e al o en el iempo pa cial de las comunicaciones. En las ablas se mues an el alo p omedio de las ejecuciones ealizadas, eliminando los alo es ex emos ex años, que suelen se en o no al 10–15% de las ejecuciones, especialmen e las que se ejecu an con mayo núme o de GPU() o nodos. 58 Capí ulo 5. Resul ados 5.3. Ejecución de las p uebas P omedio de Tiempo Pa ículas P ocesos Pa e 1,000,000 2,000,000 4,000,000 8,000,000 16,000,000 32,000,000 50,000,000 2 GPU 1177.366 I/O 4.628 MPI 0.570 O os 0.024 To al 1182.587 4 GPU 597.950 1177.426 I/O 2.663 4.643 MPI 0.364 0.397 O os -0.065 -0.072 To al 600.912 1182.394 8 GPU 315.947 597.928 1177.343 I/O 1.621 2.784 4.596 MPI 0.550 0.488 0.802 O os -0.091 -0.087 -0.006 To al 318.027 601.112 1182.736 16 GPU 165.991 316.038 597.802 1177.352 I/O 1.126 1.621 2.618 4.636 MPI 0.447 0.503 0.460 0.470 O os -0.141 -0.089 -0.105 -0.076 To al 167.422 318.072 600.775 1182.382 32 GPU 92.590 166.030 315.896 597.833 1177.577 I/O 0.848 1.133 1.634 2.621 4.598 MPI 0.751 0.536 0.557 0.629 0.570 O os -0.195 -0.149 -0.165 -0.149 -0.125 To al 93.994 167.550 317.922 600.929 1182.620 64 GPU 54.031 92.574 166.073 315.856 597.692 1177.630 1861.650 I/O 0.763 0.906 1.139 1.633 2.619 4.569 6.751 MPI 0.773 0.889 0.719 1.238 9.418 0.716 8.898 O os -0.299 -0.467 -0.223 -0.291 -0.253 -0.243 -0.169 To al 55.267 93.902 167.707 318.436 609.475 1182.672 1877.130 128 GPU 54.140 92.571 166.113 315.868 597.822 954.424 I/O 0.788 0.932 1.209 1.660 2.627 3.681 MPI 11.897 23.777 49.865 101.112 204.925 355.542 O os -0.588 -0.535 -0.618 -0.691 -0.734 -0.515 To al 66.237 116.745 216.569 417.949 804.639 1313.133 Tabla 5.6: Tiempos de ejecución de las p uebas de p oducción pa a la e sión 24 (CUDA+MPI) Pa a pode obse a mejo el compo amien o, pa a cada una de las implemen aciones y ipo de simu- lación, se p esen an ambién las siguien es g á icas: Figu a 5.5: Tiempos de ejecución en unción del núme o de GPU, pa a cada una de las ca gas (núme o de pa ículas), en escala doble loga í mica. Pe mi i á isualiza cómo escala el p og ama, pues o que lo ideal se ía ob ene una ec a de pendien e la unidad. Figu a 5.6: Tiempos de ejecución en unción del núme o de pa ículas de la ca ga, según el núme o de GPU u ilizado. Figu a 5.7: Tiempos pa ciales de la pa e de comunicaciones MPI en unción del núme o de GPU y del núme o de pa ículas simuladas. Pe mi e comp oba si la pé dida de escalabilidad se puede achaca p incipalmen e a es a pa e del p og ama, po lo que es in e esan e obse a cómo a ía ambién en e las di e en es implemen aciones y ipos de simulación ealizadas. 59 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e P omedio de Tiempo Pa ículas P ocesos Pa e 1,000,000 2,000,000 4,000,000 8,000,000 16,000,000 32,000,000 50,000,000 2 GPU 1127.802 I/O 3.932 MPI 1.416 O os 5.009 To al 1138.158 4 GPU 570.129 1127.719 I/O 2.413 3.817 MPI 0.573 0.884 O os 5.479 5.422 To al 578.594 1137.842 8 GPU 299.392 569.961 1127.621 I/O 1.759 2.446 3.700 MPI 0.576 0.626 2.696 O os 5.536 5.528 4.884 To al 307.263 578.560 1138.901 16 GPU 155.517 299.533 570.070 1127.730 I/O 1.243 1.671 2.334 3.776 MPI 0.471 0.396 0.499 0.552 O os 5.533 5.711 5.700 5.987 To al 162.763 307.310 578.602 1137.905 32 GPU 85.486 155.836 299.634 570.313 1127.866 I/O 0.945 1.263 1.675 2.397 3.748 MPI 0.750 0.546 0.702 0.540 3.493 O os 5.460 5.475 5.831 5.578 5.938 To al 92.641 163.119 307.842 578.828 1141.045 64 GPU 50.852 86.174 156.630 300.415 571.260 1128.100 I/O 0.939 0.977 1.281 1.796 2.399 3.830 MPI 0.733 0.628 0.619 0.588 3.734 9.993 O os 4.704 5.360 5.354 5.510 5.554 5.661 To al 57.228 93.138 163.884 308.309 582.947 1147.586 128 GPU 52.316 87.130 157.008 302.132 657.749 1141.044 I/O 1.017 1.146 1.478 2.077 2.578 3.297 MPI 5.944 19.578 45.918 111.002 365.214 801.091 O os 6.205 7.089 7.253 -8.189 -241.537 -670.985 To al 65.481 114.941 211.656 407.021 784.003 1274.447 Tabla 5.7: Tiempos de ejecución de las p uebas de p oducción pa a la e sión 27 (CUDA+MPI) Figu a 5.8 a) y b): E iciencia empo al ela i a del p og ama en unción del núme o de GPU, pa a cada ca ga. Se ha denominado ela i a po que no se u iliza como e e encia el código o iginal Fo an en CPU, sino que se compa a con la ejecución que u iliza menos ecu sos. En es e caso, se a a se compa a con la ejecución en un único nodo con dos GPU, dado que la ca ga meno no puede se ejecu ada po una única GPU po al a de memo ia. Figu a 5.8 c) y d): E iciencia espacial ela i a del p og ama, es o es, cómo se compo an los p og amas con el inc emen o de la ca ga, ep esen ada po el núme o de pa ículas, pues el núme o de i e aciones se man iene cons an e en 200,000 en odas es as p uebas. Figu a 5.9 Compa ación di ec a en e los iempos de ejecución de las dos e siones con GPU p esen adas ( 24 y 27). Es impo an e explica po qué apa ece el alo nega i o en la pa e de compu ación “O os”, que es 60 Capí ulo 5. Resul ados 5.3. Ejecución de las p uebas aquella que no es á con abilizada ni como de cálculo en la GPU y su ges ión, ni de ansmisiones MPI, ni de esc i u a en disco du o (“I/O”). Apa ece sob e odo en la e sión con OpenMP, y especialmen e cuando más p ocesos es án ejecu ando, y en las ca gas mayo es. Es o es debido a que hay dos hilos den o de cada p oceso, y el hilo que se enca ga de la ansmisión MPI con abiliza los iempos de comunicación y ambién la de esc i u a en disco. Po su pa e, el hilo que ejecu a los ke nels c onome a el iempo de GPU u ilizado. El hecho de que la suma de la compu ación numé ica más la de comunicaciones y esc i u a en disco es mayo al iempo o al de ejecución eal de la aplicación demues a que exis e asinc onía en e ambos conjun os de ope aciones. Así se jus i ica ambién que el iempo de ejecución global sea meno que en la e sión 24, sin OpenMP, dado que allí es án se alizadas la compu ación, la compu ación y la esc i u a en disco, de al o ma que cuando aumen a el amaño del p oblema, o el núme o de nodos pa icipando de la compu ación, aumen a la can idad de da os a ansmi i y a esc ibi en disco. 5.3.2.4.2 Escalabilidad de las implemen aciones Pa a es udia la escalabilidad de las implemen aciones p esen adas nos se i emos de las Figu as 5.5 y 5.6, donde se mues an los esul ados ob enidos en ocando la ep esen ación a un análisis espacial o empo al, espec i amen e. Núme o de pa ículas en la simulación 10 100 1000 10000 1 10 100 1000 Núme o de p ocesos (gpu) Tiempo de ejecución (s) 1,000,000 2,000,000 4,000,000 8,000,000 16,000,000 32,000,000 50,000,000 (a) Código 24 (CUDA+MPI) Núme o de pa ículas en la simulación 10 100 1000 10000 1 10 100 1000 Núme o de p ocesos (gpu) Tiempo de ejecución (s) 1,000,000 2,000,000 4,000,000 8,000,000 16,000,000 32,000,000 50,000,000 (b) Código 27 (CUDA+MPI+OpenMP) Figu a 5.5: Escalabilidad espacial de las implemen aciones ealizadas De las Tablas 5.6 y 5.7 ambién se obse a un súbi o inc emen o del iempo empleado po las comunica- ciones con MPI cuando se llega a usa g an pa e de nodos del clus e MinoTau o. Las G á icas 5.7 mues an mejo ese enómeno. De odos es os da os se obse a que: • El código con OpenMP se ejecu a en unos iempos in e io es a la e sión con sólo CUDA yMPI. • Pa ece exis i algún p oblema con la e sión 27 con OpenMP cuando se abaja con 64 nodos (128 p ocesos) y con más de 30.000.000 de pa ículas. No se ha analizado la causa, pues el enómeno de pé dida de escalabilidad en esas condiciones no se ep oduce en la e sión 24. Una hipó esis po- d ía se que cuando en ese clus e se ejecu an simulaciones con la mi ad de los nodos globales de 61 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e de la pa alelización de o os códigos cien í icos his ó icos, se p ocede a alo a cuali a i amen e la expe- iencia del p oceso de implemen ación, las decisiones que se han omado en dicha implemen ación y o as al e na i as que se desca a on al inicio del abajo. Como oda alo ación cuali a i a, és a es subje i a, po lo que a con inuación se incluye se exp esa como opinión o mada a a és de odo el p oceso de ealización del p esen e abajo. Sin emba go, se conside a muy in e esan e su exp esión explíci a en es a memo ia po que es la pa e quizá más in e esan e pa a el abajo u u o, más que la me a pe ección de las implemen aciones conc e as p esen adas. A con inuación se p esen an di e en es ca ac e ís icas que se conside an de in e és. 5.4.2.1 Facilidad de p og amación CUDA es á basado en CyC++, y és os ya son menos in ui i os de usa que Fo an pa a los cien í icos sin base especí ica de p og amación. CUDA ep esen a o o paso más allá en la p og amación, pues se cambia el pa adigma de p og amación. Po ello, la pa alelización de código his ó ico en Fo an equie e: • que una pa e del g upo de cien í icos se especialice se iamen e en p og amación pa alela, • que la a ea de p og amación sea enca gada a un expe o ingenie o in o má ico. Es a segunda opción se ía la p e e ible, pe o siemp e bajo el aseso amien o de algún cien í ico, pa a po- de cla i ica las ó mulas, algo i mos, y en de ini i a, da a en ende al p og amado qué se es á haciendo. Po la azón an e io , CUDA es un lenguaje complicado de usa , aunque exp ime a ondo las a je as g á icas de NVIDIA, muy po en es compu acionalmen e. Es in e esan e es a a en o a o as al e na i as, conc e amen e con el a ance de OpenMP, e sión 4, y sob e odo, de OpenACC. Ambos pe mi en el o load (en ío de a eas de compu ación) a los acele ado es (p incipalmen e, pe o no exclusi amen e, GPU, pe o ambién a o os como In el Xeon Phi o a je as g á icas de AMD) de una o ma más in ui i a que CUDA. En es e úl imo imes e de año se espe a el lanzamien o de e siones ya inales de compilado es lib es, como gcc, con sopo e pleno a es as nue as ecnologías. Se p opone es a a en os a es e hecho, y compa a la pa alelización usando las GPU a a és de OpenMP 4.0 y OpenACC. 5.4.2.2 Necesidad de la e o mulación de los algo i mos El g an lei mo i de Fo an es su p opio o igina io ac ónimo: FORmula TRANsla ion sys em, hecho po el cual es an ampliamen e u ilizado en el mundo cien í ico. Pe o ambién es uno de sus p incipales p oble- mas. Al se an sencillo ansc ibi las ó mulas, los mé odos ma emá icos y los algo i mos en ellas ex e - namen e isualizados, es muy ácil cae en la p og amación di ec a de la ó mula, sin pensa en aspec os impo an es desde el pun o de is a de ingenie ía del so wa e: • ene en cuen a los cos es espaciales y empo ales a la ho a de implemen a algo i mos (po ejemplo, clasi icación en canas as, que puede hace se con cos e 1 po pa ícula, o al O(N), siendo innecesa io hace p e iamen e una o denación o al de los alo es, con cos e O(N·log(N)), y después eco e o o bucle con cos e O(N)), • implemen a e icien emen e según la a qui ec u a ha dwa e y so wa e disponible en el momen o, (po ejemplo, es a ac ualizado en cuan o al a ance de la ecnología de la pa alelización, ales como OpenMP,OpenACC yCUDA), 68 Capí ulo 5. Resul ados 5.4. Valo ación • amplio conocimien o de las unciones y biblo ecas disponibles en el en o no (ejemplo, u ilización de la unción e o , e en Fo an,CyCUDA), en luga de calcula la cada ez median e mé odos numé icos de in eg ación ( apecio o Simpson). • ac ualización del p og ama base a las nue as e siones de Fo an (po ejemplo, ac ualiza los p o- g amas desde Fo an 77 o Fo an 90 a Fo an 95 (bindings con C) o incluso Fo an 2003 (in e - ope abilidad con Cy po lo an o, más ácil abaja con CUDA, y obje os, incluso aunque sean como me as es uc u as de da os, pa a ocul ación de a iables ac ualmen e decla adas globalmen e) 5.4.2.3 Depu ación semán ica de las implemen aciones his ó icas Du an e el ancu so del abajo se de ec a on lo que podían se dos e o es de p og amación que pod ían in lui de mane a muy impo an e en los esul ados que és e gene a. A la echa de hoy, el ísico esponsable de los códigos no ha con i mado ni negado esos in o mes de e o en iados. Los códigos his ó icos u ilizando ecnologías de pa alelización basadas en clus e s de CPU ienen cie a limi ación en cuan o al amaño y, po lo an o, alidez, de las simulaciones que con ellos se pueden ealiza . Ésa es la azón que undamen a es e abajo, de hecho, el es udio de iabilidad de su con e sión pa a la u ilización de a je as g á icas de compu ación. Dado ese limi ado pode de simulación, es posible que en los códigos his ó icos, aunque es én en cons an e e olución y e isión, pe sis an e o es de base, di íciles de de ec a en código y ambién en simulaciones a pequeña escala, pe o que en simulaciones mayo es, p o oquen esul ados mucho más di e gen es con la ealidad que lo azonablemen e espe ado. Es o lle a ía a alsas conclusiones y gene ación de dudas sob e el p opio cue po eó ico cien í ico en que se basan. Es po ello que an es de cada pa alelización, se e ise comple amen e el código o iginal, o incluso en mucho casos, que la implemen ación se ealice desde ce o pa a e i a de aíz es e p oblema. Es o iene la en aja de que p opo ciona ía un código mucho más adap ado a las ac uales ecnologías de p og amación pa alela. 5.4.2.4 Selección de códigos a pa aleliza Finalmen e, y es una ob iedad, dadas odas las di icul ades a iba des acadas, es siemp e con enien e ea- liza un adecuado es udio de selección de qué códigos his ó icos deben se pa alelizados con GPU. La p o- g amación de GPU pa a GPGPU es delicada, labo iosa y di ícil de depu a , pe o su ecnología (ha dwa e y so wa e), po encial y iabilidad económica, no pa an de c ece . Aún así, no odos los p oblemas ap o e- chan es a nue a ecnología. Es una labo c ucial ene a algún expe o que aseso e sob e la iabilidad de una implemen ación con CUDA u o as ecnologías eme gen es (OpenACC,OpenMP 4), y de es a o ma, di igi los es ue zos en pa- aleliza p ime o aquellas aplicaciones más in e esan es pe o que sean iables, an o desde de el pun o de is a ecnológico ( endimien o ob enido), como empo al (cos e económico de la implemen ación). 69 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e (Página in encionadamen e en blanco) 70 CA P Í T U L O 6 TRABAJO FUTURO La pa alelización ealizada de CUDA, con y sin OpenMP, sob e un código cien í ico basado en Fo an con MPI, ha pe mi ido comp oba que la aplicación del modelo SIMT a cie o ipo de códigos pe mi e ob ene unas p es aciones ex ao dina ias compa a i amen e con el código uncionando sob e CPU. También ha se ido pa a de ec a di icul ades que debe án se supe adas pa a la aplicación de es e p ocedimien o de pa alelización a o os códigos cien í icos. Las líneas de u u o a pa i de es e abajo se di iden en es ca ego ías: 1. Ampliación del abajo sob e las implemen aciones en CUDA p esen adas, con obje o de mejo a su es abilidad, endimien o y aplicabilidad: • Depu ación del p og ama, median e de ección de posibles bugs, limpieza de código y o as a- eas básicas pa a que el código quede en un es ado de mejo man enibilidad. • Mejo a la ubicación de los ce ojos de OMP, que pe mi an, po ejemplo, disminui la g anula i- dad, y con ello, in en a ex ae oda ía más p es aciones basadas en un mayo solapamien o de la compu ación con las comunicaciones y la esc i u a en disco o consola. • Mejo a p og amá icamen e lasimplemen aciones ac uales, po ejemplo, e i ando la copia, den- o de la GPU de los ec o es de elocidad nue os sob e los iejos, median e el sis ema de o a el pun e o pa a ealiza dichos cambios con copia ce o, o in es igando sob e la mejo a de p es- aciones en el uso de los gene ado es de núme os alea o ios. • In es iga po qué azón no se pe mi en ejecuciones de simulaciones de más de 825.000 pa í- culas po cada GPU, cuando és as pa ecen dispone de su icien e capacidad de almacenamien o (6GB). Posiblemen e a a és de una mejo ges ión de los di e en es ipos de memo ia p esen es en la a je a g á ica. • Reimplemen a el código u ilizando o as ecnologías de pa alelización eme gen es que am- bién se ap o echan de los acele ado es g á icos, pe o pe mi en ambién hace lo sob e o os ace- le ado es, como sob e las a je as g á icas de AMD, o sob e odo, del acele ado In el Xeon Phi. Ejemplos de es as ecnologías eme gen es son OpenMP 4 y OpenACC, de los cuales se espe a una implemen ación inal lib e en es e úl imo imes e g acias a GNU. 71 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e 2. Aplicación de la pa alelización a o os códigos cien í icos • Mig a o os códigos simila es a CUDA, u o as ecnologías eme gen es, pa a ap o echa el ex- emo pode compu acional de las GPU. • Modi ica las ecuaciones base de Lange in ob enidas inicialmen e, a a és de la p og esi a eliminación de algunas de las ap oximaciones e ec uadas pa a la deducción del conjun o de ecuaciones implemen adas en el p esen e código. Es o debe ía pe mi e ob ene simulaciones más ce canas a la ealidad. • A pa i de es a expe iencia, y iendo que es posible, añadi complejidad al cue po eó ico inicial, po ejemplo, median e la conside ación de la ela i idad den o del conjun o de ecuaciones, pa a de es a o ma ace ca se más si cabe al compo amien o que end á el plasma cuando es é ealmen e en eacción de usión con olada de o ma con inuada, bien sea en un eac o de ipo okamak o del ipo s ella a o . 3. U ilización a ondo de la nue a capacidad de simulación en es e campo, y es udia nue as condicio- nes, a iando los pa áme os de en ada, ealizando las opo unas mínimas a iaciones en el código si es necesa io, de al o ma que se con ie a en una he amien a más pa a la comp ensión del com- po amien o eal del plasma de pa ículas con inado magné icamen e. 72 CA P Í T U L O 7 CONCLUSIONES En el p esen e T abajo Fin de Más e se ha es udiado la iabilidad y me odología pa a pode mig a códi- go his ó ico cien í ico, gene almen e en lenguaje Fo an, al nue o pa adigma de p og amación median e acele ado es ex e nos, conc e amen e, median e a je as g á icas de NVIDIA a a és de su ecnología p o- pia p i a i a CUDA. Pa a ello se ha u ilizado como caso de es udio un código en Fo an ya pa alelizado median e MPI, que pe mi e la simulación de un plasma de pa ículas a al a empe a u a con inado magné- icamen e, que es la ísica básica en la que se basan los ac uales in en os de usión en calien e en eac o es de ipo okamak. Los esul ados cuan i a i os son plenamen e sa is ac o ios: se han ob enido acele aciones de has a 700 en compa ación con el código ejecu ándose en un único núcleo de CPU, odo ello, ob eniendo esul ados co ec os en la simulación y gene ando co ec amen e los iche os in e medios es adís icos pe inen es. Es- o ep esen a g andes en ajas e iden es: pe mi e la ejecución de simulaciones más g andes en un iempo azonable, y disminuye el cos e de explo ación (núme o de ho as) de u ilización del clus e de compu ación de al as p es aciones. Asímismo, se ha edac ado una memo ia que puede se i de guía sob e los p oblemas, di icul ades, éc- nicas, plan eamien os y decisiones que pod ían su gi du an e la ealización a g an escala de la mig ación de código cien í ico a las nue as ecnologías de GPGPU. A pesa de ello, la iabilidad de ex ende es e caso de es udio a o os p oyec os más ambiciosos no debe sólo se in o mado po es a alo ación muy posi i a desde el pun o de is a cuan i a i o. Como ya se ha indicado al inal de los Capí ulos 4 y 5, hay o os aspec os cuali a i os que sean an o o más impo an es a la ho a de oma la decisión: • La mig ación es écnicamen e iable y se han demos ado cómo ealiza la. No obs an e, no es i- ial. El g ado de complejidad del código esul ado es ele ado, sob e odo po la mezcla de di e en es lenguajes y ecnologías. Es po ello que deben conclui se se p ecisa una pe sona cona la su icien e o mación en compu ación pa alela y en ingenie ía del so wa e pa a pode en en a se con ga an ías y exi osamen e al e o. • A pesa de lo an e io , es muy impo an e des aca la impo ancia que iene el signi icado ísico de las ó mulas y algo i mos que implemen an, así como el conocimien o del enómeno ísico es udiado, y 73 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e con ello, los esul ados que a p io i pueden ob ene se de la simulación. Es po ello que en el p oceso de mig ación ambién debe es a p esen e un expe o en la ma e ia cien í ica de que se a e, al obje o de guia , comp oba y e isa que la implemen ación se ajus e a la eo ía, es o es, e ec úe una e isión semán ica de los códigos. • Con ela i amen e poco es ue zo, se pueden consegui esul ados no ables. Mejo a los es ambién posible, pe o muchas eces equie e de una in e sión en iempo y ecnología más ele ado. Pe o odo depende del caso conc e o. • Hay un aspec o que no se ha podido analiza , y es la exis encia de compilado es especí icos p i a i os de pago que pe mi en el uso de CUDA di ec amen e sob e un dialec o o a iación del lenguaje Fo - an. Cabe la posibilidad de que su uso sea más sencillo que la solución p opues a, y que los equisi os del pe sonal a ca go del p oyec o sean ambién di e en es. • No odo código cien í ico his ó ico puede se ácilmen e pa alelizable con CUDA, o incluso puede que sea con ap oducen e hace lo. Se equie e un buen conocimien o de CUDA y del p og ama a mig a , pa a pode de e mina si el nue o pa adigma de pa alelización masi a de la GPGPU es adecuado al p oyec o en cues ión. Como ejemplos, los mé odos ecu si os deben pasa se a i e a i os. Tampoco se ole a la p og amación dinámica, aunque es poco común en el mundo cien í ico. Los algo i mos con g an dependencia en e da os pueden p esen a p oblemas de acceso a memo ia que disminuyan mucho el endimien o de las GPU. • Es po ello que es necesa ia una adecuada selección del código o algo i mo que a a se mig ado o implemen ado con CUDA, pa a ex ae de es a ecnología odo lo que puede apo a . 74 GLOSARIO bu e Espacio de memo ia in e media, que puede se so wa e o ha dwa e. bug En español, simplemen e, e o , aunque a eces se aduce li e almen e como bicho, o gusano. E o , gene almen e su il, en el código que p o oca el compo amien o anómalo del p og ama, incluso p o- ocando su inalización ab up a. Especialmen e pelig osos cuando sus mani es aciones son e á i- cas, dependiendo de alguna combinación de e minada de condiciones pa a e ela se. Muy di íciles de depu a en compu ación pa alela y dis ibuida. cache Memo ia ápida in e media de pequeña capacidad pa a acele a el acceso a una memo ia p incipal de g an capacidad más len a. clus e Conjun o de o denado es in e conec ados po una ed de al as p es aciones pa a la ejecución con- jun a de aplicaciones pa alelas. compu ación g id En o no de ejecución HPC o mado po a ios clus e s localizados geog á icamen e en ubicaciones muy dispe sas y adminis ados po di e en es o ganizaciones que, no obs an e, compa - en sus ecu sos en e sí pa a log a una mayo ocupación y u ilización de sus ins alaciones, o ecien- do una in e az de lanzamien o común. Las ejecuciones deben se plani icadas y coo dinadas, pa a lo cual exis en ges o es de colas especializados que pueden unciona en base a c édi os. CUDA ke nel Función cuyo código que se ejecu a en una GPU, y que po lo an o de mane a masi amen e pa alelizada. daemon Apicación se ido a que se inicia usualmen e en el a anque del sis ema ope a i o y que se enca - ga de ges iona un de e minado se icio pa a el sis ema ope a i o o las aplicaciones. Aunque no mal- men e se suele aduci al español po «demonio», en es e abajo se op a po el é mino «duende», é mino que pa ece más co ec o, pues demonio iene conno aciones malignas que no cabe espe a de un se icio legí mo a ancado po el sis ema ope a i o, mien as que el é mino «duende» e oca más al e ec o «mágico» de aquello que se ocupa de hace algo po noso os sin casi pe cibi esa ayuda. dis upción Fenómeno b usco en un plasma en con inamien o magné ico po el cual pie de sus condi- ciones de es abilidad y que p o oca que en muy pocos milisegundos desapa ezcan las condiciones necesa ias pa a la usión nuclea , apagándose el eac o como consecuencia. HPC High Pe o mance Compu ing, compu ación de al o endimien o. Consis e en la ejecución de p o- g amas eales en en o nos masi os de p oducción do ados de la úl ima ecnología especialmen e diseñados y ges ionados con el in de ob ene las máximas p es aciones compu acionales.. mapea Tecnicismo u ilizado pa a desc ibi la acción de e e encia un obje o o zona de memo ia desde o a ubicación, gene almen e con el obje o de i ualiza o acili a el acceso di ec o al obje o o iginal. 75 Pa alelización median e GPU de la dinámica de elec ones en plasma con inado magné icamen e middlewa e Pieza de so wa e que o ece se icios a o as aplicaciones más allá de los que o ece el sis ema ope a i o, ales como pe mi i la mú ua in e acción o el in e cambio de in o mación. Es comúnmen e u ilizado en sis emas dis ibuidos, en donde un único sis ema ope a i o no puede cen aliza y ges- iona oda la in e comunicación. O o uso es acili a al p og amado el acceso a di e en es ecu sos de en ada/salida, ya sean locales o emo os, y de ipo so wa e o ha dwa e. name mangling Nomencla u a pa icula u ilizada po cada compilado pa a cada lenguaje pa a la deno- minación de sus símbolos en los iche os obje o .o. plasma Es ado de la ma e ia, simila al gas, en el que po el e ec o de las al ísimas empe a u as, la ma e ia se haya comple amen e ionizada, dada la al a elocidad é mica de las pa ículas. p o iling In es igación del compo amien o de un p og ama a a és de la in o mación ecopilada desde el análisis dinámico del mismo. Exis en p og amas especí icos pa a ales a eas, siendo el más común gp o , den o del so wa e lib e en Linux. NVIDIA pone a disposición den o de su Toolki , la u ilidad n p o . Aunque el é mino inglés es á ampliamen e di undido, puede aduci se po “pe ilaje” o po “análisis del endimien o”.. un ime Biblio eca de unciones que dan sopo e pa a la ejecución de aplicaciones que u ilicen un de e - minado lenguaje, en nues o caso CUDA C. speedup Acele ación. Es una mé ica de la mejo a ela i a en p es aciones del endimien o pa a la ejecu- ción de una misma a ea. En el ámbi o de la compu ación pa alela, hace e e encia al cocien e en e el iempo de ejecución de la a ea secuencial y el iempo de ejecución del p og ama pa alelizado. s ella a o Disposi i o donde se con ina magné icamen e un plasma de pa ículas a al a empe a u a con la inalidad de ob ene la usión nuclea con olada. Tiene la pa icula idad de que el campo magné- ico es al que las ó bi as de las pa ículas con enidas iene una o ma helicolidal. okamak Disposi i o simila al s ella a o , pe o en el que el campo magné ico con ina al plasma en una conjun o de ó bi as de o ma o oidal. wa p En español, u dimb e. Rep esen a la ag upación de 32 ó 64 hilos de ejecución (según la a qui ec u a CUDA) que ienen un ámbi o de ejecución común en una a je a g á ica de NVIDIA. Compa en e- cu sos esenciales y escasos como egis os y memo ia cache, pe o sob e odo, ienen un mismo lujo de ejecución, de al o ma que si és e es di e gen e pa a cada hilo, los hilos que no es én en esa bi u - cación pe manecen pa ados, y luego al e és, de al o ma que se acumulan los iempos. Es o es así po que cada u dimb e sólo iene ac i a una ins ucción, la cual se á o no ejecu ada po el hilo que la componga. La u dimb e es, en la indus ia ex il, el hilo longi udinal que se man ienen en ensión en un ma co o ela, y que si e de guía y es uc u a sob e la que se en e ejen los hilos. w appe Función en ol o io, que o ece una in e az conocida pa a accede a o a unción o secuencia de unciones. Su uso puede es a mo i ado po di e en es causas, ales como acili a el acceso a o as unciones, simpli icando la o ma de in oca las, o de esconde el uncionamien o in e no de un so - wa e po emas de segu idad o de p o ección de la p opiedad in elec ual. 76 SIGLAS API Applica ion P og amming In e ace, In e az de p og amación de aplicaciones. ARPACK ARnoldi PACKage, Paque e de (mé odo de) A noldi. BLAS Basic Linea Algeb a Subp og ams, Subp og amas de álgeb a lineal básica. CPU Cen al P ocessing Uni , Unidad cen al de p ocesamien o. CUDA Compu e Uni ied De ice A chi ec u e, A qui ec u a uni icada de disposi i os de compu ación. FORTRAN An e io men e FORmula TRAsla ion sys em, Sis ema de aducción de ó mulas. En la ac uali- dad ya no se in e p e a de es a mane a, se conside a una palab a p opia. GPFS Gene al Pa allel File Sys em, Sis ema de iche os pa alelo gene al (desa ollado po IBM. GPGPU Gene al-Pu pose compu ing on G aphics P ocessing Uni s, Compu ación de p opósi o gene al en unidades de p ocesamien o g á ico. GPU G aphics P ocessing Uni , Unidad de p ocesamien o g á ico. ITER In e na ional The monuclea Expe imen al Reac o , Reac o expe imen al e monuclea in e nacio- nal. LAPACK Linea Algeb a PACKage, Paque e de álgeb a lineal. Lis Lib a y o I e a i e Sol e s, Biblio eca de esol edo es i e a i os. MPI Message Passing In e ace, In e az de in e cambio de mensajes. OpenCL Open Compu ing Language, Lenguaje abie o de compu ación. OpenMP Open Mul i-P ocessing, Mul ip ocesamien o abie o. PETSc Po able, Ex ensible Toolki o Scien i ic Compu a ion, Conjun o po able y ex ensible de he a- mien as pa a la compu ación cien í ica. ScaLAPACK Scalable Linea Algeb a PACKage, Paque e de álge a lineal escalable. SIMT Single Ins uc ion Mul iple Th eads, Ins ucción única, múl iples hilos. Modelo de ejecución en pa- alelo, u ilizada po las GPU de NVIDIA con capacidad CUDA, en e o as, en la que múl iples hilos ejecu an concu en emen e la misma ins ucción sob e di e en es da os. SLEPc Scalable Lib a y o Eigen alue P oblem Compu a ions, Biblio eca escalable pa a compu aciones sob e el p oblema de los alo es p opios o ca ac e ís icos. 77