scieee Open visual document viewer

Implementing P Systems Parallelism by Means of GPUs

Cecilia, José M.; García, José M.; Guerrero, Ginés D.; Martínez del Amor, Miguel Ángel; Pérez Hurtado de Mendoza, Ignacio; Pérez Jiménez, Mario de Jesús

Abstract

Software development for Membrane Computing is growing up yielding new applications. Nowadays, the efficiency of P systems simulators have become a critical point when working with instances of large size. The newest generation of GPUs (Graphics Processing Units) provide a massively parallel framework to compute general purpose computations. We present GPUs as an alternative to obtain better performance in the simulation of P systems and we illustrate it by giving a solution to the N-Queens problem as an example.

Full text

Implemen ing P Sys ems Pa allelism by Means o GPUs Jose M. Cecilia2,Jos´eM.Ga c´ıa2,Gin´es D. Gue e o2, Miguel A. Ma ´ınez–del–Amo 1, Ignacio P´e ez–Hu ado1, and Ma io J. P´e ez–Jim´enez1 1Resea ch G oup on Na u al Compu ing Depa men o Compu e Science and A ificial In elligence Uni e si y o Se illa A da. Reina Me cedes s/n, 41012 Se illa, Spain {mdelamo ,pe ezh,ma pe }@us.es 2G upo de A qui ec u a y Compu aci´on Pa alela Dp o. Ingenie ´ıa y Tecnolog´ıa de Compu ado es Uni e sidad de Mu cia Campus de Espina do, 30100 Mu cia, Spain {chema,jmga cia,gines.gue e o}@di ec.um.es Abs ac . So wa e de elopmen o Memb ane Compu ing is g owing up yielding new applica ions. Nowadays, he efficiency o P sys ems sim- ula o s ha e become a c i ical poin when wo king wi h ins ances o la ge size. The newes gene a ion o GPUs (G aphics P ocessing Uni s) p o- ide a massi ely pa allel amewo k o compu e gene al pu pose compu- a ions. We p esen GPUs as an al e na i e o ob ain be e pe o mance in he simula ion o P sys ems and we illus a e i by gi ing a solu ion o he N-Queens p oblem as an example. 1 In oduc ion Memb ane Compu ing is an eme ging b anch wi hin Na u al Compu ing ha was in oduced by Gh. P˘aun [24]. The main idea is o conside biochemical p ocesses aking place inside li ing cells om a compu a ional poin o iew, in a way ha gi es us a new nonde e minis ic model o compu a ion by using cellula machines. Up o now, i has no been possible o ha e implemen a ions nei he in i o no in i o o P sys ems, so he compu a ion and analysis o hese de ices a e pe o med by simula o s. The e o e, P sys ems simula o s a e ools ha help he esea che s o ex ac esul s om models. Since P sys ems was p esen ed, many so wa e applica ions ha e been p oduced [11]. These simula o s ha e o be as much efficien as possible when handling la ge p oblem sizes. Thus, he massi ely pa allel na u e o P sys ems compu a ions poin s ou o look o a massi ely pa allel echnology whe e he simula o can un efficien ly. Pa allel compu a ion on clus e s is he adi ional en i onmen o speed- up pa allel applica ions. Pa icula ly, many simula o s o P sys ems ha e been designed o clus e s o compu e s [4]. Howe e , his compu a ion is ela i ely expensi e and i is a ailable o o ganiza ions ha ha e enough esou ces o buy and main ain hose clus e s. Nowadays, he e a e o he cheape solu ions in he compu e ma ke ha also p o ides pa allel en i onmen s. Among hese solu- ions, he newes gene a ion o g aphics p ocesso uni s (GPUs) a e massi ely pa allel p ocesso s which allow o de elop a wide ange o pa allel applica ions. We also ecall ha o he pa allel compu ing pla o ms o P sys ems simula o s a e being in es iga ed, such as special ha dwa e ci cui s [6] and FPGAs [20]. GPUs can suppo se e al housand o concu en h eads p o iding a mas- si ely pa allel en i onmen whe e pa allel applica ions can ob ain huge pe o - mance [14][17][29]. Cu en N idia’s GPUs, o example, con ain up o 240 scala p ocessing elemen s pe chip [16], hey a e p og ammed using C and CUDA [32][21], and hey ha e low cos compa ed wi h a clus e o compu e s. In his pape , we use CUDA as pa allel p og amming en i onmen o P sys- ems simula o in o de o speedup he simula ion. The inpu o he simula o is a P sys em which is defined by using he P-Lingua [5] p og amming language, and he ou pu is a de ailed lis o in o ma ion o e e y configu a ion o he compu a ion. The simula ion is di ided in wo main s ages: selec ion s age and execu ion s age. A his s age o de elopmen , he simula o simula es ecognize P sys ems wi h ac i e memb anes, he selec ion s age is execu ed on he GPU and he execu ion s age is execu ed on he CPU. The es o he pape is s uc u ed as ollows. In Sec ion 2 se e al defini ions and concep s a e gi en o a co ec unde s anding o he pape . Sec ion 3 in- oduces he Compu e Unified De ice A chi ec u e (CUDA) and some concep s o p og amming on GPUs a e specified. In Sec ion 4 we explain he design o he simula o . In Sec ion 5 we implemen a solu ion o he N-Queens p oblem using he simula o and P-Lingua. Finally, in Sec ion 6 we show some esul s and compa e hem wi h he sequen ial e sion o he simula o . The pape ends wi h some conclusions and ideas o u u e wo k in Sec ion 7. 2 P elimina ies Polynomial ime solu ions o NP-comple e p oblems in Memb ane Compu ing a e achie ed by ading ime o space. This is inspi ed by he capabili y o cells o p oduce an exponen ial numbe o new memb anes in polynomial ime. The e a e many ways a li ing cell can p oduce new memb anes: mi osis (cell di ision), au opoiesis (memb ane c ea ion), gemma ion, e c. Following hese in- spi a ions a numbe o diffe en models o P sys ems has a isen, and many o hem p o ed o be compu a ional comple eness ( hey a e equi alen in powe o Tu ing machines). In his pape we ocus on he model o P sys ems wi h ac i e memb anes.I is one o he mos s udied models in Memb ane Compu ing and one o he fi s models p esen ed by Gh. P˘aun [25]. P sys ems wi h ac i e memb anes is o med by a memb ane s uc u e, whe e a label and a pola iza ion is associa ed o each memb ane. In his model, e e y elemen a y memb ane is able o di ide i sel by ep oducing i s con en in o a new memb ane. He e we p o ide a sho ecall o i s ea u es (see [25] o de ails). The model o P sys em wi h ac i e memb anes is a cons uc o he o m Π=(O, H, μ, ω1,...,ω m,R), whe e m≥1 is he ini ial deg ee o he sys em; O is he alphabe o objec s,His a fini e se o labels o memb anes; μis a mem- b ane s uc u e (a oo ed ee), consis ing o mmemb anes injec i ely labelled wi h elemen s o H,ω1,...,ω ma e s ings o e O, desc ibing he mul ise s o objec s placed in he m egions o μ;andRis a fini e se o ules,whe eeach ule is o one o he ollowing o ms: (a) [a→ ]α hwhe e h∈H,α∈{+,−,0}(elec ical cha ges), a∈Oand is a s ing o e Odesc ibing a mul ise o objec s associa ed wi h memb anes and depending on he label and he cha ge o he memb anes (e olu ion ules). (b) a[] α h→[b]β hwhe e h∈H,α, β ∈{+,−,0},a, b ∈O(send-in communi- ca ion ules). An objec is in oduced in he memb ane, possibly modified, and he ini ial cha ge αis changed o β. (c) [a]α h→[] β hbwhe e h∈H,α, β ∈{+,−,0},a, b ∈O(send-ou communica- ion ules). An objec is sen ou o he memb ane, possibly modified, and he ini ial cha ge αis changed o β. (d) [a]α h→bwhe e h∈H,α∈{+,−,0},a, b ∈O(dissolu ion ules). A memb ane wi h a specific cha ge is dissol ed in eac ion wi h a (possibly modified) objec . (e) [a]α h→[b]β h[c]γ hwhe e h∈H,α, β, γ ∈{+,−,0},a, b, c ∈O(di ision ules). A memb ane is di ided in o wo memb anes. The objec s inside he mem- b ane a e eplica ed, excep o a, ha may be modified in each memb ane. Rules a e applied acco ding o he ollowing p inciples: –All he elemen s which a e no in ol ed in any o he ope a ions o be applied emain unchanged. –Rules associa ed wi h label ha e used o all memb anes wi h his label, no ma e whe he he memb ane is an ini ial one o whe he i was gene a ed by di ision du ing he compu a ion. –Rules om (a) o (e) a e used as usual in he amewo k o memb ane com- pu ing, i.e., in a maximal pa allel way. In one s ep, each objec in a memb ane can only be used by a mos one ule (non-de e minis ically chosen), bu any objec which can e ol e by a ule mus do i (wi h he es ic ions indica ed below). –Rules (b) o (e) canno be applied simul aneously in a memb ane in one compu a ion s ep. –An objec ain a memb ane labelled wi h handwi hcha geαcan igge a di ision, yielding wo memb anes wi h label h, one o hem ha ing cha ge β and he o he one ha ing cha ge γ.No e ha all hecon en sp esen be o e he di ision, excep o objec a, can be he subjec o ules in pa allel wi h he di ision. In his case we conside ha in a single s ep wo p ocesses ake place: “fi s ” he con en s a e affec ed by he ules applied o hem, and “a e ha ” he esul s a e eplica ed in o he wo new memb anes. –I a memb ane is dissol ed, i s con en (mul ise and in e io memb anes) becomes pa o he immedia ely ex e nal one. The skin is ne e dissol ed nei he di ided. No e ha P sys ems can be seen as de ices wi h wo le els o pa allelism: among memb anes (e e y memb ane wo ks independen ly, wi h he excep ion o when he e a e communica ion ac oss hem) and among objec s inside a mem- b ane ( he ules a e applied o he exis ing mul ise o objec s in a maximal pa allel way). Recognize P sys ems we e in oduced in [26], and cons i u e he na u al amewo k o s udy he sol abili y o decision p oblems. The da a ep esen ing an ins ance o he p oblem has o be p o ided o he P sys em o compu e he app op ia e answe . This is done by codi ying each ins ance as a mul ise placed in an inpu memb ane. The ou pu o he compu a ion, yes o no,issen o he en i onmen in e e y hal ing configu a ion. Fu he mo e, he ac o simula ing some hing gene ally en ails ep esen ing ce ain key cha ac e is ics o beha iou s o some physical, o abs ac , sys em. Howe e , an emula ion ool duplica es he unc ions o one sys em by using a diffe en sys em, so ha he second sys em beha es like (and appea s o be) he fi s sys em. Wi h he cu en echnology, we can no emula e he unc ionali y o a cellula machine by using a con en ional compu e o sol e NP-comple e p oblems in polynomial ime, bu we can simula e hese cellula machines, no necessa ily in polynomial ime, in o de o aid esea che s. Howe e , depending on he unde lying echnology whe e he simula o is execu ed, he simula ions can ake oo much ime. The echnology used o his wo k is called CUDA (Compu e Unified De- ice A chi ec u e). CUDA is a co-designed ha dwa e and so wa e solu ion o make easie de eloping gene al-pu pose applica ions on he G aphics P ocesso Uni (GPU) [34]. GPUs, ha a e one o he main componen s o adi ional compu e s, o iginally we e specialized o ma h-in ensi e, highly pa allel com- pu a ion which is he na u e o g aphics applica ions. These cha ac e is ics o he GPU we e e y a ac i e o accele a e scien ific applica ions which ha e massi ely pa allel compu a ions. Howe e , he p oblem was he way o p o- g am gene al pu pose applica ions on he GPU. This way in ol ed o deal wi h GPUs designed o ideo games, so hey ha e had o une hei applica ions using p og amming idioms ied o compu e g aphics, p og aming en i onmen igh ly cons ained, e c [17] [14]. The CUDA ex ensions de eloped by N idia p o ides an easie en i onmen o p og am gene al-pu pose applica ions on o he GPU, because i is based on ANSI C, suppo ed by se e al keywo ds and cons uc s. ANSI C is he s anda d published by he Ame ican Na ional S an- da ds Ins i u e (ANSI) o he C p og amming language, which is one o he mos used. P sys ems de ices a e massi ely pa allel, wha fi s, in a simila way, in o mas- si ely pa allel na u e o he GPUs wi h housands o h eads unning in pa allel. These h eads a e uni s o execu ion which execu e he same code concu en ly on diffe en pieces o da a. 3 G aphics P ocessing Uni D i en by he ideo games ma ke , p og ammable GPUs (G aphics P ocessing Uni s) ha e e ol ed in o a highly pa allel, mul i h eaded, manyco e p ocesso . They we e designed o accele a e g aphics applica ions, which ans o m h ee- dimensional da a (coo dina es o iangle e ices) in o pixels ha a e displayed on a sc een, using o his ask p og amming in e aces such as OpenGL and Di ec X. The massi ely pa allel na u e o g aphics applica ions and i s a i h- me ic in ensi y leads he esea ches o explo e mo e gene al non-g aphics ap- plica ions on o he GPU, c ea ing a new p og amming field called GPGPU (Gene al-Pu pose on GPUs). GPUs ha e become an inexpensi e and eadily a ailable single-chip massi ely pa allel sys em. Howe e , GPGPU p og amme s had o deal wi h he limi a ions and difficul ies o cons ained g aphics p imi i es o compu e hei non-g aphics compu a ions. The eme gence o Compu e Unified De ice A chi ec u e (CUDA) [34] p og amming model, p oposed by N idia Co po a ion in 2007, has helped o de elop highly-pa allel applica ions on o he GPU easie han i was be o e. CUDA allows GPGPU p og amme s o de elop hei applica ions in a mo e a- milia en i onmen by using C/C++ p og amming language, wi h some ex en- sions o manipula e special aspec s o he GPU. Mo eo e , N idia consolida ed his end launching a line o GPUs op imized o gene al pu pose compu a ions called TESLA [16]. In his wo k we use a Tesla C1060 g aphics p ocesso uni (GPU) om N idia as ha dwa e a ge o i s s udy. This sec ion in oduces he Tesla C1060 compu ing a chi ec u e. In addi ion, i analyses he h eading model o Tesla a chi ec u es, and also he mos impo an issues in he CUDA p og amming en i onmen . 3.1 Tesla C1060 Base Mic oa chi ec u e The Tesla C1060 [16] is based on a scalable p ocesso a ay which has 240 s eaming-p ocesso (SP) co es o ganised as 30 s eaming mul ip ocesso (SM). The applica ions s a a he hos side ( he CPU) which communica es wi h he de ice side ( he GPU) h ough a PCI-Exp ess x16 bus (see he op o figu e 1). The SM is he p ocessing uni , and i is unified g aphics and compu ing mul- ip ocesso . E e y SM con ains eigh SPs a i hme ic co es, one double p ecision uni , 16-Kby e ead/w i e sha ed memo y, a se o 16384 egis e s, and access o he off-chip memo y (global/local memo y). The access o sha ed memo y is e y cheap, howe e , he access o he off-chip memo y has low pe o mance because i is ou o he chip, as i is shown on figu e 1. In addi ion, able 1 shows all memo ies a ailable on he GPU and also he cos o access hem. Fig. 1. Tesla C1060 GPU wi h 240 SPs: S eamming P ocesso s, o ganised in 30 SMs: S eamming Mul ip ocesso s Table 1. Memo y Sys em on he Tesla C1060 Memo y Loca ion Size La ency Access Regis e s On-Chip 16384 32-bi s Regis e s pe SM ≃0cycles R/W Sha ed Memo y On-Chip 16 KB pe SM ≃ egis e s R/W Cons an On-Chip 64 KB ≃ egis e s R Tex u e On-Chip Up o Global >100 cycles R Local Off-Chip 4GB 400-600 cycles R/W Global Off-Chip 4GB 400-600 cycles R/W 3.2 Pa allel Compu ing wi h CUDA The GPU is seen as a coop ocesso ha execu es da a-pa allel ke nel unc ions. The use c ea es a p og am encompassing CPU code (Hos code) and GPU code (Ke nel code). They a e sepa a ed and compiled by n cc (N idia’s compile o CUDA code) as shown in figu e 2. Fi s ly, he hos code is esponsible o ans e ing da a om he main memo y (RAM o hos memo y) o he GPU memo y (de ice memo y), using CUDA in- s uc ions, such as cudamemcpy. Mo eo e , he hos code has o s a e he numbe o h eads execu ing he ke nel unc ion and he o ganiza ion o hem. Th eads execu e he ke nel code, and hey a e o ganized in o a h ee-le el hie a chy as i is shown in figu e 3. A he highes le el, each ke nel c ea es a single g id ha consis s o many h ead blocks. Each h ead block can con ain up o 512 C/C++ wi h CUDA Ex ensions NVCC CPU Code PTX Code PTX o Ta ge Compile G80 PTX Code T10 Fig. 2. N cc compila ion p ocess h eads, which can sha e da a h ough Sha ed Memo y and can pe o m ba - ie synch oniza ion by in oking he –sync h eads p imi i e [31]. Besides, h ead blocks can no pe o m synch oniza ion. The synch oniza ion ac oss blocks can only be ob ained by e mina ing he ke nel. Fu he mo e, he hos code calls he ke nel unc ion like a C unc ion by passing pa ame e s i i is needed, and also by speci ying he numbe o h eads pe block and he numbe o blocks making up he g id. Each block wi hin he g id has hei own iden ifie [22]. This iden ifie can be one, wo o h ee Fig. 3. Th ead o ganiza ion in CUDA p og amming model dimensions depending on how he p og amme has decla ed he g id, accessed ia .x,.y,and.z index fields. Each h ead wi hin he block ha e hei own iden ifie which can be one, wo o h ee dimensions as well. Combining h ead and block iden ifie s, he h eads can access o diffe en da a add ess, and also selec he wo k ha hey ha e o do. The ke nel code is specified h ough he key wo d global and he syn ax is: global ke nelName <<< dimG id, dimBlock >>> (...pa ame e lis ...) whe e dimG id and dimBlock a e h ee-elemen s ec o s ha speci y he dimensions o he g id in blocks and he dimensions o he blocks in h eads, espec i ely [21]. 3.3 Th eading Model A SM is a ha dwa e de ice specifically designed wi h mul i h eaded capabil- i ies. Each SM manages and execu es up o 1024 h eads in ha dwa e wi h ze o scheduling o e head. Each h ead has i s own h ead execu ion s a e and can execu e an independen code pa h. The SMs execu e h eads in a Single- Ins uc ion Mul iple-Th ead (SIMT) ashion [16]. Basically, in he SIMT model all he h eads execu e he same ins uc ion on diffe en piece o da a. The SMs c ea e, manage, schedule and execu e h eads in g oups o 32 h eads. This se o 32 h eads is called Wa p. Each SM can handle up o 32 Wa ps (1024 h eads in o al, see able 2). Indi idual h eads o he same Wa p mus be o he same ype and s a oge he a he same p og am add ess, bu hey a e ee o b anch and execu e independen ly. Table 2. Majo Ha dwa e and So wa e Limi a ions p og aming on CUDA Con igu a ion Pa ame e s Limi a ion Th eads/SM 1024 Th ead Blocks/SM 8 32-bi Regis e s/SM 16384 Sha ed Memo y/SM 16KB Th eads/Block 512 Th eads/Wa p 32 Wa ps/SM 32 The execu ion flow begins wi h a se o Wa ps eady o be selec ed. The ins uc- ion uni selec s one o hem, which is eady o issue and execu ing ins uc ions. The SM maps all he h eads in an ac i e Wa p pe SP co e, and each h ead ex- ecu es independen ly wi h i s own ins uc ions and egis e s a e. Some h eads o he ac i e Wa p can be inac i e due o b anching o p edica ion, and i is also ano he c i ical poin in he op imisa ion p ocess. The maximum pe o mance is achie ed when all he h eads in an ac i e Wa p akes he same pa h ( he same execu ion flow). I he h eads o a Wa p di e ge, he Wa p se ially execu es each b anch pa h aken, disabling h eads ha a e no on ha pa h, and when all he pa hs comple e, he h eads econ e ge o he o iginal execu ion pa h. 4 Design o he Simula o o Recognize P Sys ems In his sec ion we b iefly desc ibe he simula o o ecognize P sys ems wi h ac- i e memb anes, elemen a y di ision and pola iza ion. Fi s ly, we explain he p e- ious wo k ha we ha e done in o de o p epa e he de elopmen o he pa allel simula o on he GPU. Then, we in oduce he algo i hm design in he CUDA p og amming language, and finally, we finish wi h ou simula o ’s design. 4.1 Design o he Baseline Simula o As p e iously men ioned, CUDA p og amming model is based on C/C++ lan- guage. The e o e, he fi s ecommended s ep when de eloping applica ions in CUDA is o s a om a baseline algo i hm w i en in C++, whe e some pa s can be suscep ible o be pa allelized on he GPU. In his wo k, we ha e based on he simula o o P sys ems wi h ac i e mem- b anes de eloped in PLinguaCo e [5]. This sequen ial (o single- h eaded) simu- la o is p og ammed in JAVA, so he fi s s ep was o ansla e he code o C++. The simula o is execu ed in o wo main s ages: selec ion s age and execu ion s age.Theselec ion s age consis s o he sea ch o he ules o be execu ed in each memb ane. Once he ules ha e been selec ed, he execu ion s age consis s o he execu ion o hese ules. The inpu da a o he selec ion s age consis s o he desc ip ion o he mem- b anes wi h hei mul ise s (s ings o e he wo king alphabe O, labels associ- a ed wi h he memb ane in H,e c.),and hese o ulesR o be selec ed. The ou pu da a o his s age is he se o selec ed ules. Only he execu ion s age changes he in o ma ion o he configu a ion. I is he eason because execu ion s age needs synch oniza ion when accessing o he memb ane s uc u e and he mul ise s. A his poin o implemen a ion, we ha e pa allelized he selec ion s age on he GPU, and he execu ion s age is s ill execu ed on he CPU because o he synch oniza ion p oblem. We also ha e de eloped an adap ed sequen ial simula o o he CPU (called as sequen ial simula o ), which has he same cons ain s as he CUDA simu- la o explained in he nex subsec ions o make a ai compa ison among hem. This simula o achie es much be e pe o mance han he o iginal sequen ial simula o . 4.2 Algo i hm Design in CUDA Whene e we design algo i hms in he CUDA p og amming model, ou main effo is di iding he equi ed wo k in o p ocessing pieces, which ha e o be p ocessed by TB h ead blocks o T h eads each. Using a h ead block size o T=256, i is empi ically de e mined o ob ain he o e all bes pe o mance on he Tesla C1060 [28]. Each h ead block access o one diffe en se o inpu da a, and assigns a single o small cons an numbe o inpu elemen s o each h ead. Each h ead block can be conside ed independen o he o he , and i is a his le el a which in e nal communica ion (among h eads) is cheap using explici