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