Con olle s: An abs ac ion o ease he use o ha dwa e accele a o s
Ana Mo e on-Fe nandez, Hec o O ega-A anz, and A u o Gonzalez-Esc ibano
Uni e sidad de Valladolid, Spain
Decembe 2016
Abs ac
Nowadays he use o ha dwa e accele a o s, such as he G aphics P ocessing Uni s (GPUs) o XeonPhi cop ocesso s, is key
o sol e compu a ionally cos ly p oblems ha equi e High Pe o mance Compu ing (HPC). Howe e , p og amming solu ions o
an e icien deploymen in his kind o de ices is a e y complex ask ha elies on he manual managemen o memo y ans e s
and con igu a ion pa ame e s. The p og amme has o ca y ou a deep s udy o he pa icula da a needed o be compu ed a each
momen , in di e en compu ing pla o ms, also conside ing a chi ec u al de ails.
We in oduce he Con olle concep as an abs ac en i y ha allows he p og amme o easily manage he communica ions
and ke nel launching de ails on ha dwa e accele a o s in a anspa en way. This model also p o ides he possibili y o de ining
and launching CPU ke nels in mul i-co e p ocesso s wi h he same abs ac ion and me hodology used o he accele a o s. I
in e nally combines di e en na i e p og amming models and echnologies o exploi he po en ial o each kind o de ice. Ad-
di ionally, he model also allows he p og amme o simpli y he p ope selec ion o alues o se e al con igu a ion pa ame e s
ha can be selec ed when a ke nel is launched. This is done h ough a quali a i e cha ac e iza ion p ocess o he ke nel code o
be execu ed.
Finally, we p esen he implemen a ion o he Con olle model in a p o o ype lib a y, oge he wi h i s applica ion in se e al
case s udies. I s use has led o educ ions in he de elopmen and po ing cos s, wi h signi ican ly low o e heads in he execu ion
imes when compa ed o manually p og ammed and op imized solu ions using di ec ly CUDA and OpenMP.
1 In oduc ion
Cu en ly, he sys ems used o High Pe o mance Compu ing (HPC) include accele a o de ices. This end is no iceable om
pe sonal compu e s o classical supe compu e s. Examples o hese HPC he e ogeneous en i onmen s can be seen in he i s
posi ions o he cu en op 500 lis o supe compu e s. When de eloping solu ions o be deployed in he e ogeneous sys ems,
we can: (1) Use a single p og amming model esponsible o managing he a chi ec u al and concep ual di e ences be ween he
di e en compu a ional de ices, e.g. OpenACC; o (2) Use a combina ion o di e en p og amming models speci ic o each
kind o compu a ional de ices, e.g. MPI oge he wi h CUDA, OpenCL, o OpenMP.
One o he main d awbacks o he i s app oach is he di icul y o ep esen non-comple ely egula p og ams, wi h non-
i ial communica ions o synch oniza ions. Besides, he inal gene a ed code esul ing om p og amming wi h his kind o
models is no usually as op imized as he o iginal code. Fo example, he cu en OpenACC compile implemen a ions do no
o e a comple e solu ion o he p oblem o au oma ically choose app op ia e alues o se e al ke nel-launching pa ame e s
wi hou any p og amme guidance, as equi ed by he s anda d. These pa ame e s include he h eadblock size and i s mul i-
dimensional geome y, o he con igu a ion o he sizes o he L1 cache s. he sha ed-memo y in mode n GPUs. On he o he
hand, implemen ing solu ions ollowing he second app oach equi es a deepe knowledge o he di e en pa allel p og amming
models in ol ed, using di e en synch oniza ion and memo y access s a egies o di e en de ices. Addi ionally, he p og am-
me is he inal esponsible o p ope ly managing he da a ans e among he di e en memo y spaces o he compu a ional
de ices, a he mos app op ia e imes, oge he wi h he choice o p ope alues o ke nel-launching con igu a ion pa ame e s.
Howe e , his manual adjus men gi es o he p og amme he possibili y o op imize he use o he pa icula esou ces o each
speci ic de ice. O he app oaches ha y o c ea e a concep ual b idge be ween hese wo app oaches a e domain speci ic, o a e
based in sophis ica ed compile echnology o he gene a ion o coo dina ion o ke nel codes.
In his pape we p esen he concep o Con olle as an abs ac en i y ha allows he p og amme o anspa en ly manage
he launching o se ies o asks on accele a o de ices, and/o mul i-co e CPU p ocesso s. I is a ede ini ion o he Communica o
en i y p esen ed in [1], wi h new unc ionali ies and a comple e new in e nal design ha leads o a e y e icien implemen a ion.
The Con olle uses he mos app op ia e p og amming models (CUDA, OpenMP, ...) o exploi he compu a ional esou ces
o he accele a o s and hos s machines. The model has se e al impo an ea u es: (1) A mechanism o de ine common ke nels
eusable ac oss di e en ypes o de ices, o specialized ke nels o speci ic de ice kinds, including he possibili y o conside
a subse o CPU co es as a single independen de ice; (2) An op imiza ion sys em o selec p ope alues o ke nel-launching
con igu a ion pa ame e s (such as he h eadblock geome y), guided by simple quali a i e code cha ac e iza ion p o ided by he
p og amme ; (3) A anspa en mechanism o memo y managemen , including op imized communica ions o he da a s uc u es
1
be ween he hos and he co esponding images in he accele a o s; (4) An abs ac ion o indexed da a s uc u es ha uni ies he
da a managemen in ke nels o di e en kinds o de ices (such as GPUs, o mul i- h eaded ec o CPU co es).
This pape also desc ibes a p o o ype, p esen ed as a C/C++ lib a y ex ension, ha implemen s he Con olle en i y. The
p o o ype is designed o exploi NVIDIA’s GPUs using he CUDA pa allel p og amming model, o subse s o CPU co es using
OpenMP. We also p esen an expe imen al s udy based on se e al case s udies, some o hem ob ained om he CUDA Toolki
Samples. The s udy shows how he use o he p o o ype implies a educ ion o he p og amming e o needed o implemen and
po hese applica ions, when i is compa ed wi h he o iginal e sions ha di ec ly use he speci ic na i e pa allel p og amming
models. Besides, we show ha ou implemen a ion does no in oduce signi ican pe o mance o e heads.
The s uc u e o he a icle is as ollows. In Sec . 2 we discuss some ela ed wo k. Sec ion 3 p esen s he Con olle model.
Sec ion 4 explains he p o o ype lib a y, i s usage o de elop a p og am and ele an implemen a ion de ails. Sec ion 5 desc ibes
he expe imen al s udy. Finally, Sec . 6 desc ibes he conclusions o his wo k and p oblems ha will be add essed in he u u e.
2 Rela ed wo k
In his sec ion we discuss some wo k ela ed o ou p oposal. We co e some examples o di e en app oaches ha ease he
in eg a ed p og amming o di e en compu a ional de ices, including accele a o s.
The e a e se e al wo ks in eg a ing, in he same ool, di e en pa allel languages o models, o conside ing di e en de ice
ypes. One o hese models, widely known, is OpenCL. I p o ides he Con ex abs ac ion, de ining a memo y model in which
associa ed da a is sha ed o mo ed be ween he hos and he de ice memo ies when needed. Al hough OpenCL is in e nally
exploi ing he endo d i e s and na i e p og amming ools, i s abs ac ions ha e been p o ed o p e en ob aining he same
e iciency as when using di ec ly he endo p og amming models, o se e al common si ua ions [11]. Mo eo e , he implemen-
a ions a e no easily eachable wi h espec o he de ini ion o he managemen policies o he in e nal queues, o he possibili y
o change hem. Ano he d awback o his model is he manual managemen o ke nel compila ion a un- ime, o di e en
a chi ec u es in di e en con ex s, ha is desi ed o be gene alized and simpli ied. The llCoMP ool [17] is a sou ce- o-sou ce
compile ha ansla es C anno a ed code o MPI+OpenMP o CUDA. Howe e , i does no suppo he join use o CUDA
wi h he o he pa allel models. SkelCL [18] enhances he OpenCL in e ace o allow he coo dina ion o di e en GPUs on he
same machine. The wo ks [3, 12] in oduce hyb id MPI+CUDA app oaches o coo dina e GPUs in he same o di e en hos
machines. Apa om hei speci ic limi a ions, hey do no ha e an abs ac suppo o easily manage homogeneously uni s o
di e en na u e, including CPU co es, as in ou app oach.
The e a e o he app oaches ha a e mo e domain speci ic, bu include small abs ac ions o ease he managemen be ween
he accele a o and he hos (e.g. MCUDA [19] o hiCUDA [8]). They do no conside guided op imiza ions, no allow he
p og amme o con ol he load dis ibu ion o he de ices coo dina ion. O he app oaches wi h simila limi a ions also conside
CUDA, MPI, and OpenMP combina ions (e.g. [21, 9]).
Mo e gene al app oaches p opose comple e in eg a ed amewo ks, such as OMPICUDA [13], S a PU [10], o he skele on
p og amming amewo k based on i , SkePU [4]. In gene al, hey hide he coo dina ion de ails o he p og amme , o he poin
o cons aining he po en ial op imiza ions ha could be achie ed manually. The selec ion o launching pa ame e s like he
h eadblock size is ackled in SkePU, bu using a ial-and-e o me hods, hus lea ing no possibili y o ex apola e he esul s o
o he ke nel codes o a chi ec u es. A highe -le el app oach is o ely on compile echnology o anspa en ly gene a e code o
di e en kinds o de ices (bo h o coo dina ion and o ke nels) om a single uni ied language. Fo example, C++14 is used in
PACXX [7], a ans o ma ion sys em in eg a ed in o he LLVM compile amewo k. I ans o m explici pa allel cons uc ions
ha use he concep o ke nel and launching is an abs ac an elegan o m. On he o he hand, some o he solu ions a e dependen
on ea u es o his language, and hey a e no easily po able o o he ones. The decisions abou launching pa ame e s, such as
he h eadblock choice, a e s ill he p og amme esponsibili y alone.
Ou solu ion sol es p e ious limi a ions. I in eg a es he coo dina ion o compu a ional de ices o e y di e en na u es
in he con ex o a compile lib a y, using a new app oach based on simple abs ac ions. I hides he di e ences o execu ion
models up o he poin o allowing he de elopmen o gene ic ke nels po able ac oss de ices. Bu a he same ime, i allows he
in eg a ion o na i e o endo p og amming models, applying speci ic op imiza ion echniques, and a oiding e iciency losses
associa ed wi h o he gene ic app oaches.
Ou app oach can be also used as a un- ime sys em o gene alize he po ing o p og ams ac oss di e en ype o accele a o
de ices. Se e al au oma ic code-gene a ion compile- ime ools [2, 5, 14] can de i e, om sequen ial code wi h p agma anno a-
ions: (1) The da a dependences among he ke nel launches; (2) The needed da a ans e s be ween he hos and he a ge de ices,
and; (3) The domains on which ke nels should be execu ed o accele a o s. Such ools ha e been used ypically o gene a e code
o only one kind o accele a o . Fo example, in [14], he OpenMP 4.0 #p agma o load ea u e is used o gene a e codes o
execu e on he XeonPhi accele a o . Using he same in o ma ion ex ac ed by his echnologies, oge he wi h he abs ac ions
and gene ic p og amming guidelines o ou p oposal, i is possible o gene a e a gene ic code alid o any kind o execu ion
de ice. This app oach will be s udied in a u u e wo k.
2
De ice
P ocesso
Memo y
Execu ion
policy
Queue
Hos
Memo y
In e nal
execu ion
Ke nel
launch
Ke nel
B
CC
A
B
Con olle
Code
a iables
Hi Tile
Bounded
In e nal
Bounded
Figu e 1: Diag am o he Con olle model a chi ec u e. The ke nel-launching eques s can be enqueued. The Con olle en i y
manages he execu ion o enqueued ke nels, and o bounded a iables, he da a ans e s be ween memo y spaces. In he igu e,
he hos a iable A is no bounded o he Con olle . Va iable B is a bounded a iable, wi h a duplica ed image o he da a in he
de ice memo y. Va iable C is an in e nal a iable o he Con olle , de ined in he hos , bu alloca ed only in he de ice.
3 Con olle Model
The Con olle model in oduces a simpli ied way o p og am applica ions ha can exploi he e ogeneous compu a ional pla -
o ms including accele a o s o /and mul i-co e CPUs. I s a chi ec u e is ep esen ed in Fig. 1. The de ice Con olle s coo dina e
he execu ion o se ies o ke nels. These ke nels a e decla ed as unc ions, ha a e managed by he Con olle en i ies. Con olle s
au oma ically manage he wo main concep s used in a p og am ha exploi s accele a o s:
•Ke nel managemen , including he ke nel launching and con igu a ion. The Con olle manages he deploymen /execu ion
o sequences o ke nel unc ions in he compu a ional de ice associa ed o he Con olle . The Con olle s can include
policies o exploi concu en ke nel execu ion echniques, in e lea e compu a ions wi h communica ions, o eo de he
sequence o ke nels. The ke nel con igu a ion is he selec ion o speci ic con igu a ion pa ame e s o he ke nel launching,
ha can be associa ed o a pa icula ke nel and compu a ional pla o m.
•Da a managemen , including he da a ans e s ca ied ou ac oss he memo y hie a chies o he hos and he accele a o s
and he abs ac ion used o access da a elemen s independen ly o he a ge de ice, he h eads indexes space, o he da a
layou .
3.1 Ke nel managemen
Ke nel launching is an ope a ion ha inse s in he Con olle queue an o de o execu e a ke nel, wi h gi en eal pa ame e s,
when he associa ed de ice (accele a o o se o CPU co es) is a ailable. This is done using a launching p imi i e (see Fig. 1).
The Con olle in e nals will ensu e ha he inpu da a ha e been ans e ed o he de ice memo y i needed, and ha p e ious
ke nels ha e ended, be o e doing he eal launch. The Con olle execu ion policy could eo de he ke nels in he queue o
maximize he execu ion e iciency as a as he inpu /ou pu dependencies ac oss ke nels a e no iola ed. The main ea u es and
con ibu ions o ou model in he ke nel managemen a e desc ibed below:
3.1.1 Ke nel de ini ion and launching: Suppo o bo h gene ic and specialized codes
The model p o ides a gene ic me hod o decla e ke nels. We de ine a ke nel implemen a ion as a uple: (< name >, <
de iceT ype >, < pa ame e s lis >, < code >), whe e de iceType is a symbol, o lis o symbols (e.g. GPU, CPU), in-
dica ing he speci ic kind o de ice/s whe e he ke nel can be deployed. Thus, we p opose o allow he decla a ion o se e al
implemen a ions o he same ke nel name, indica ing he a chi ec u e/s o which his pa icula ke nel is mo e sui able.
We also p opose he possibili y o de ining ke nel implemen a ions o a gene ic de ice ype. These ke nels will be po able
ac oss di e en ypes o de ices and a chi ec u es, and will be used by de aul when no speci ic ke nel implemen a ion is p o ided
in he code o a gi en a ge de ice. As hey a e simply conside ed ano he ke nel implemen a ion, wi h a di e en de ice ype
symbol, hey can be decla ed in he same p og am oge he wi h o he implemen a ions o he same ke nel name. This allows he
c ea ion o lib a ies o ke nel implemen a ions, de eloping i s a gene ic and po able one ha will wo k in any pla o m, and
in oducing g adually mo e speci ic and op imized ke nel implemen a ions o speci ic o new a ge de ices.
The code o he gene ic ke nels implemen a ion, in o de o be po able ac oss di e en de ice ypes, should comply in ou
model wi h a se o p ope ies (see also Sec . 4.6 o a discussion abou an example o a gene ic ke nel used on bo h mul i-co e
3
CPU and GPU de ices). P ope ies o a gene ic ke nel: (1) The code is pu e da a-pa allel, in he sense ha i can be execu ed by
many logical h eads, and he p og amme ensu es ha no da a dependencies o ace condi ions can a ise ac oss hem; (2) The
code ope a es on indexed da a s uc u es using abs ac 1, 2, o 3-dimensional h ead indexes p o ided by he un ime sys em
( x, y, z), ha a e in e nally adjus ed o each de ice o access da a in ow-majo o de keeping maximum coalescing and/o
ec o iza ion p ope ies in he speci ic de ice; (3) The code does no explici ly use esou ces, p imi i es, o synch oniza ion
mechanisms, ha a e speci ic o he p og amming model o a gi en de ice; (4) All accesses o da a s uc u es a e done h ough
an abs ac in e ace p o ided by he un ime sys em, ha is independen o he a ge de ice chosen (see Sec . 3.2.2 o ou
p oposal o choose such an abs ac in e ace).
A uni ied ke nel launching unc ion, a un- ime, ma ches he a chi ec u e o he de ice o he bes implemen a ion p o ided
a compile ime o ha a chi ec u e. In his way, i is possible o anspa en ly use ei he gene ic po able ke nels, op imized
ke nels p og ammed in he na i e models o a speci ic de ice, o e en w appe s o call specialized lib a ies o speci ic de ice
ypes. The Con olle un ime sys em can choose he mos app op ia e one.
Finally, we p opose o equi e in he pa ame e s lis o a ke nel implemen a ion, ha he p og amme epo s he inpu /ou pu
ole o each ke nel pa ame e . This will be used by he Con olle o de ec which da a should be ans e ed among di e en
memo y spaces, a each momen , depending on he sequence o ke nels launched.
3.1.2 Cha ac e isa ion o ke nels o execu ion
Ou model conside s he cha ac e iza ion o he ke nel code o au oma ically op imize launching pa ame e s, such as he h ead-
block geome ies. We p opose o in eg a e he model o quali a i e cha ac e is ics p esen ed in [15, 20] in ou Con olle . To use
his model, he p og amme should examine he ke nel code, and he should concep ually cha ac e ize i in classes acco ding o
h ee main c i e ia. We in oduce an ex ension o he ke nel implemen a ion uple wi h a new elemen o p o ide he classi ica ion
o he ke nel code. The Con olle in e nally uses an associa i e able o implemen he selec ion o launching pa ame e s
acco ding o he ules p oposed on he p e iously ci ed pape s. Mo e de ails abou he c i e ia and examples a e p esen ed in
Sec . 4.4.
3.2 Da a managemen
Accele a o s may ha e hei own memo y spaces, o cing o ans e he inpu da a o he ke nels, and he ob ained esul s,
be ween he memo y o he hos pla o m and he memo y o he de ice accele a o . The manual managemen o hese da a
mo emen s can be cumbe some and e o -p one. Mo eo e , i is di icul o p edic in ad ance when asynch onous da a ans e s
a e possible, o when da a should s ay in he accele a o de ice memo y, as his depends on he exac sequence o ke nels
launching. Ou model abs ac s om he p og amme all hese issues. The main ea u es/con ibu ions o ou model on da a
managemen a e he ollowing:
3.2.1 Da a ans e s be ween he hos and he accele a o s
In ou model, a Con olle is associa ed, in he momen o i s c ea ion, wi h a pa icula accele a o o subse o CPU co es, and
i anspa en ly manages in he accele a o memo y space he images o he hos memo y da a s uc u es. The Con olle can
decide when and how he ans e s should be ca ied ou , depending on he da a s uc u es used in he co esponding ke nels, and
he sequence o ke nels enqueued o launching. The p og amme can use he o iginal names o he hos a iables in he ke nel
anspa en ly. Depending on he ole o he a iables (named da a-s uc u es) used as eal pa ame e s in ke nel launches, we can
dis inguish wo ypes: Bounded a iables and in e nal a iables.
Bounded a iables
Bounded a iables a e hos a iables ha ha e an image in he memo y space o he accele a o (see Fig. 1). The model de ines
an ope a ion o bind a hos a iable o he Con olle . Once a a iable is bounded, i s da a should no be modi ied by he p og am
a he hos side un il an unbinding ope a ion is applied.
The i s ke nel equi ing he use o a bounded a iable as an inpu will o ce he Con olle o anspa en ly ensu e ha i s
da a ha e been al eady ans e ed. Applying an unbinding ope a ion o a bounded a iable will o ce he ans e o i s da a om
he accele a o o he hos , i i has been used as ou pu by any ke nel. The main p og am wai s un il he end o he ke nels using
ha a iable, and he end o he da a ansmission o he hos i needed.
In e nal a iables
In e nal a iables a e a iables whose scope is delimi ed o he ke nels execu ed in he accele a o . They a e only handled inside
he memo y space o he de ice, and hey will no ha e alloca ion in he hos memo y space. Thus, hey ne e imply a da a
ans e (see Fig. 1). In he pa icula case o a de ice ep esen ing a subse o CPU co es, he memo y should be anspa en ly
alloca ed and managed in he hos de ice.
4
The model de ines an ope a ion o c ea e an in e nal a iable in he Con olle . A da a s uc u e is decla ed wi hou alloca ing
memo y o i in he hos . The name o his da a s uc u e is used in he c ea ion ope a ion o clone he ype, size, and in e nal
s uc u e in he Con olle memo y space. Since ha momen , he name can be used as eal pa ame e in ke nel launching as a
e e ence o he in e nal a iable. To des oy an in e nal a iable i is needed o apply ano he ope a ion using he e e ence name
o he hos a iable.
3.2.2 Uni o m da a accesses
One o he key ea u es ha a p og amming model o he e ogeneous sys ems should p o ide is he abili y o manage uni o mly
he da a s uc u es on a p og am. As p e iously commen ed when discussing gene ic ke nel de ini ions, he da a accesses on he
body o he ke nel should be independen o he a ge de ice chosen. We p opose he use o abs ac h ead indexes and da a-
accessing me hods in he ke nel codes. They a e de ised in o de o design he codes o wo k e icien ly when accessing elemen s
in ow-majo o de , independen ly o he de ice. We p opose he use o he da a-handle abs ac ion o a ays in oduced by
Hi map [6], a lib a y o hie a chically dis ibu ed a ays. E icien implemen a ions ha e been p o ided o di e en de ices,
such as CPUs, o GPUs. See mo e de ails abou he Hi map unc ionali ies used in he implemen a ion o ou model in Sec . 4.1.
4 The Con olle s lib a y
We ha e de eloped an implemen a ion o he Con olle s model. I is designed as a lib a y w i en in C99 code. Thus, i can
be used o de elop C/C++ p og ams, independen ly o he chosen compile . The lib a y de ines unc ions, bu also elies on
p ep ocesso mac os o code ew i ing. Al hough his allows an e icien in e ace implemen a ion in a compile agnos ic way,
he p og amme should ake ca e o use he in e ace in he expec ed way o a oid p oblems de i ed om some common pi alls
o mac o subs i u ions (unexpec ed ype checking issues, misnes ing due o inco ec code injec ion in he mac o pa ame e s,
e c.).
The cu en implemen a ion suppo s NVIDIA’s GPU de ices using CUDA in e nally, and subse s o CPU codes using
OpenMP in e nally. In ou implemen a ion, a ke nel is a unc ion coded o any, o o a speci ic compu a ional de ice wi h
pa icula inpu /ou pu pa ame e s. The Con olle s lib a y in e ace de ines p imi i es o: C ea e Con olle s; Decla e and
cha ac e ize ke nels; Manage hos da a s uc u es ha can be bounded o c ea ed as in e nal a iables in he Con olle s, and
anspa en ly accessed inside he ke nels; and Launch he ke nels.
4.1 Da a s uc u es and Hi map
Rega ding he da a s uc u es ha he model handles, we ha e decided o in eg a e ou implemen a ion wi h Hi map [6], an
e icien lib a y o hie a chical iling and mapping o a ays. Hi map de ines he Hi Tile s uc u e, an abs ac en i y o n-
dimensional a ays and iles. A Hi Tile s uc u e is a handle o s o e a ay me a-da a, along wi h he poin e o he ac ual memo y
space o he da a. The ci cles o he a iables A, B, and C in Fig. 1 ep esen he handle s. The Hi Tile s uc u e should be
specialized o each a ay base ype a he beginning o he p og am.
The e a e only ou unc ions o Hi map needed o wo k wi h he Con olle s implemen a ion. Fi s , hi ileDomain and
hi ileDomainAlloc a e used o decla e he index domains o a ile a ay, also alloca ing he memo y o he da a in he second
case. The unc ion hi ileF ee is used o ee he da a memo y and clean he handle . The unc ion hi ileElem is used in hos
o ke nels code o access he elemen s o a ile. I ecei es a ile name, a numbe o dimensions, and he indexes alues o he
desi ed elemen . This unc ion p o ides an homogeneous in e ace o manage da a s uc u es in bo h, hos and accele a o de ice
ke nels. The da a a e s o ed and accessed in ow majo o de in all cases.
Hi map lib a y includes many o he unc ionali ies. They include managemen o hie a chical subselec ions o pa s o he
iles, and anspa en managemen o dis ibu ed-a ays, wi h abs ac pa i ion and communica ion unc ionali ies ha in e nally
use a message-passing pa adigm (exploi ing MPI). This will allow in he u u e he anspa en in eg a ion o he Con olle s in
dis ibu ed mul i-node clus e s.
Mos o he me a-da a in he Hi Tile s uc u e a e only needed o ad anced unc ionali ies, such as dis ibu ed-a ay pa i ions
and communica ions in he hos code. In ou new implemen a ion we de ine a new much smalle handle (Hi KTile) wi h he
minimum in o ma ion needed o da a accesses o mul idimensional a ays h ough a da a poin e . The ke nel launching in e ace
will anspa en ly ans o m he handle s o his new ype, subs i u ing he da a poin e wi h i s equi alen in he de ice memo y
space. The da a accesses inside he ke nels uses his in e nal poin e along wi h a minimum numbe o a i hme ic ope a ions,
exposing he exp essions o he na i e compile o open he possibili y o u he op imiza ions. The esul ob ains as good
pe o mance as he classical a ay accesses.
In Hi map, he implemen a ion o di e en kinds o iles hides o he p og amme he de ails o he da a access o di e en
in e nal da a layou s. The Hi map lib a y al eady in eg a es spa se domains and spa se da a s uc u es in o he Hi Tile abs ac-
ion. The u u e ans o ma ion o hese handle s uc u es, o implemen e icien and po able ke nels, should ollow a simila
app oach as he one used o dense a ays.
5

4.2 Con olle s and a iables managemen
Ini ializa ion and des uc ion
A Con olle is associa ed o a pa icula compu a ional pla o m (accele a o o CPU-co es se ) a he momen o i s c ea ion.
This is done h ough he C lC ea e unc ion. This p imi i e has wo main pa ame e s: The name o he Con olle a iable,
and he iden i ica ion o he associa ed compu ing de ice. The p og amme can ee he esou ces wi h he C lDes oy
unc ion.
The Con olle s associa ed o CPU co es in e nally use OpenMP. To allow he launching o CPU ke nels asynch onously, as
in he case o GPUs, ou implemen a ion uses OpenMP asks. A mas e h ead execu es he hos code, and one OpenMP ask is
ac i a ed o each co e associa ed o a CPU Con olle . The Con olle ini ializes on i s c ea ion an in e nal a ay o s uc u es
wi h one elemen pe assigned co e ask, o con ol hei ac i i ies and hei synch oniza ion.
Binding a iables
The unc ion C lA ach binds a ile de ined in he code o he hos wi h a Con olle . The unc ion C lDe ach unbinds
i . I he memo y space o he de ice is no he same as he hos , he a ach ope a ion alloca es memo y o he da a in he de ice
space. A e he binding, he hos should no manipula e he ile da a un il i is unbounded. Du ing he ime he a iable is
bounded, ope a ions o copy he da a o and om he associa ed de ice can happen a any ime, as an in e nal decision o he
Con olle . Fo he pa icula case o Con olle s associa ed o CPU co es, he e is no da a duplica ion. The ke nels may be
modi ying he hos a iable da a a any ime.
In he cu en implemen a ion we ha e in eg a ed pa o he a iables managemen in he Hi map ile handle s. The handle
s o es he poin e o he de ice memo y, a e e ence o he Con olle is bounded o, and lags o indica e he clean/di y s a e o
he da a. This a oids he duplica ion o bindings, a emp s o unbind a iables om he w ong Con olle , e c. The lags help he
Con olle o choose he p ope momen s o synch onize he da a be ween he hos and he de ice memo y image.
C ea ing in e nal a iables
The unc ion C lIn e nalC ea e c ea es an in e nal a iable on he de ice memo y space. On he o he hand, he unc ion
C lIn e nalDes oy is used o ee he memo y space in he assigned compu a ional de ice. Fo he case o Con olle s
associa ed o accele a o s, his kind o a iables does no need o ha e alloca ed memo y in he hos side. A ile ini ialized wi h a
domain in o ma ion is enough. E en i i would ha e alloca ed memo y, he e is no synch oniza ion o da a be ween he hos and
he de ice image c ea ed by his unc ionali y.
4.3 Decla a ion and con igu a ion o ke nels
A ke nel is decla ed by using he p imi i e KERNEL < ype>. Whe e ype may be emp y o indica e a gene ic ke nel, us-
able on any kind o de ice, o a speci ic alue o a specialized code o a gi en ype o de ice. This is use ul when di e -
en op imiza ions on he ke nel code a e equi ed o di e en de ices. Cu en ly, he lib a y suppo s he speci ic p imi i es
KERNEL GPU o CUDA code a ge ing NVIDIA’s GPUs, KERNEL CPU o hos machine code a ge ing se s o CPU co es, and
KERNEL GPU WRAPPER o hos machine code which includes calls o specialized GPU lib a ies, like o example cuBLAS
ou ines. The Con olle ensu es ha bo h kinds o GPU ke nels a e execu ed wi h exclusi e con ol o he associa ed de ice.
This allows o au oma ically coo dina e he launch o sequences o classical ke nels and calls o GPU lib a ies in he same de ice
wi h ou model.
The ke nel-de ini ion p imi i es decla e in b acke s he numbe o pa ame e s o he ke nel, wi h a uple o in o ma ion o
each pa ame e . The pa ame e in o ma ion includes i s ype, name, and inpu /ou pu ole:
•IN: o inpu Hi Tile pa ame e s, whose elemen s a e only ead.
•OUT: o ou pu Hi Tile pa ame e s, whose elemen s a e only w i en.
•IO: o inpu and ou pu Hi Tile pa ame e s, wi h elemen s can be bo h ead and w i en.
•INVAL: o inpu pa ame e s o any ype passed by alue.
Fo he case o accele a o de ices, wi h sepa a ed memo y spaces, his con igu a ion allows he Con olle s o de e mine i
i is necessa y o ca y ou da a ans e s wi h he main memo y o he hos when ke nels a e launched. The p imi i e is ollowed
by a s uc u ed block wi h he ke nel code. Thus, i esembles a C unc ion heade .
The Con olle CPU implemen a ion con ains he loops o execu e he ke nel code o each elemen o he h eads index
space, ha is in e nally assigned o a speci ic OpenMP h ead. Bo h CPU and GPU ke nels code use a p ede ined h ead s uc u e
wi h h ee in ege elemen s x,y,z indica ing he 3-dimensional indexes o he co esponding ine-g ain h ead. In o de o ease
he ke nel euse ac oss di e en de ice kinds, hese indexes a e adap ed o ha e he same ow-majo meaning in bo h ypes o
6
1/*CUDA ke nel launch */
2/*Pa ame e s:
3name: Ke nel name
4 h eads: C lTh eads, index space limi s
5a ch: GPU a chi ec u e
6pa ams: eal pa ame e s o he launch
7kcha : cha ac e iza ion o he ke nel
8*/
9dim3 block = CALBlockModel(name,kcha ,a ch);
10 dim3 g id = CALG idDi Up( h eads,block);
11 w appe _gpu_##name<<<g id, block>>>
12 ( h eads, pa ams);
13 ...
14 __global__ oid w appe _gpu_##name( ... ) {
15 CALTh ead h eadId = { h eads.dims,
16 h eadIdx.y + blockDim.y *blockIdx.y,
17 h eadIdx.x + blockDim.x *blockIdx.x,
18 h eadIdx.z + blockDim.z *blockIdx.z };
19 ...
20 ke nel_gpu_##name( h eadId, pa ams );
21 }
1/*OpenMP index space pa i ion and ke nel launch */
2/*Pa ame e s:
3name: Ke nel name
4n_ h: OpenMP h ead iden i ie
5numCo es: numbe o CPU-co es associa ed o
6 he cu en con olle
7 h eads: C lTh eads s uc u e wi h limi s o he
8 h eads index space
9pa ams: eal pa ame e s o he ke nel launch
10 */
11 in s ipSize =(in ) ceil( h eads->x/(numCo es) );
12 in begin= n_ h*(s ipSize);
13 in end= MIN( i s + (s ipSize) - 1,( h eads->x)-1);
14 o (i=begin; i<=end; i++)
15 o (j=0; j< h eads->y; j++){
16 h eadId.x = i;
17 h eadId.y = j;
18 h eadId.z = 0;
19 ke nel_cpu_##name( h eadId, pa ams);
20 }
Figu e 2: Exce p s o he Con olle lib a y code gene a ed o ke nel deploymen /launching on a CUDA capable GPU de ice
(le ), and a g oup o CPU-co es ( igh ). In he CUDA e sion he block geome y is selec ed using he ke nel cha ac e iza ion
p o ided by he p og amme . We also show he pa o he w appe unc ion launched, ha e e ses he h ead indexes be o e
calling he ac ual ke nel code. Fo he case o he CPU-co es, each OpenMP h ead (assigned o a co e) execu es his code. In
his example, he 2-dimensional index space is di ided in o blocks by ows wi hou balancing he emaining. The code execu es
he loops ha call he ke nel code o each index-elemen mapped o he h ead. I simula es he many- h ead app oach used o
GPU p og amming using coa se -g ain OpenMP asks.
ke nels. Also, he space o alid indexes ha can be de ined in he ke nel launching is independen o he h eadblock sizes in
GPUs. Ac ual h eads wi h iden i ie s ou side he chosen launching index space will be clipped anspa en ly by he Con olle
launching sys em. Fo specialized GPU ke nels, he ke nel code may use CUDA p imi i es mixed wi h he Hi map da a accesses
ha use he new po able index sys em.
4.4 Ke nel cha ac e iza ion
The ke nel cha ac e iza ion is a p og amme hin o help he sys em o au oma ically de e mine p ope ke nel launching pa am-
e e s in e ms o special code ea u es and pla o m a chi ec u e in o ma ion. The CPU h eads g anula i y in ou p o o ype is
de e mined by a simple egula blocking policy, ha does no equi e a speci ic ke nel cha ac e iza ion.
Fo GPU ke nels, ou cu en p o o ype lib a y in eg a es he model p esen ed in [15, 20]. This model allows o de e mine
con igu a ion pa ame e s (g id, h eadblock and L1 cache memo y sizes), o NVIDIA’s GPUs. The p imi i e KERNEL CHAR,
aken om [16], is used o p o ide o he Con olle he cha ac e iza ion o he ke nels. The p imi i e ecei es he ke nel name,
he numbe o dimensions o he h ead space (1, 2, o 3), and desc ip i e alues o he cha ac e iza ion model. These alues
a e a quali a i e desc ip ion o cha ac e is ics o he ke nel code p o ided by he p og amme . They a e ela ed o: (a) The
coalescing p ope y o he global memo y access pa e ns ( ull, medium, sca e ); (b) The a io o a i hme ic/logic ope a ions
pe global memo y access (high, medium, low); and (c) The a io o da a sha ing accesses in a block pe global memo y access
(high, medium, low). Fo he de aul case, when he p og amme canno p o ide a p ope cha ac e iza ion o all he pa ame e s,
ade keywo d can be used ins ead o one o he gi en alues, and he model p o ides ypical h eadblock alues o each CUDA
a chi ec u e, ha wo k well in a gene al case, maximizing occupancy i possible, e c. Fo a mo e de ailed desc ip ion o his
quali a i e desc ip o s wi h en a i e quan i a i e anges, see he wo ks [15, 20, 16].
Ou implemen a ion ex ends his cha ac e iza ion wi h he possibili y o speci y a ixed size o he h eadblock. This is use ul
o ke nels ha ely on speci ic block sizes o geome ies o manipula e sha ed memo y ( o example he ma ix mul iplica ion
code in he CUDA Toolki Samples).
4.5 Ke nel launching
The unc ion C lLaunch is used o launch a ke nel, wi h gi en eal pa ame e s, o he compu a ional de ice associa ed o a
Con olle . The launched ke nel will be enqueued, and e en ually execu ed wi h he co esponding con igu a ion de i ed om
he in o ma ion p o ided by he cha ac e iza ion p imi i es. Cu en ly, he p o o ype only suppo s a Fi s -Coming-Fi s -Se ice
policy o ke nels execu ion. The launching unc ion has he ollowing pa ame e s: (a) The Con olle ; (b) The name o he
7
1/*Hi Tile specializa ion
2*C ea es new ype Hi Tile_ loa */
3hi _ ileNewType( loa );
4
5/*Ke nel cha ac e iza ions */
6KERNEL_CHAR(copyCell,1,de ,de ,de )
7KERNEL_CHAR(upda eCell,1, ull,low,high)
8
9/*Gene ic ke nel codes o any de ice */
10 KERNEL(copyCell, 2,
11 OUT, Hi Tile_ loa , ds ,
12 IN, Hi Tile_ loa , s c) {
13
14 hi _ ileElem( ds , 2, h ead.x, h ead.y ) =
15 hi _ ileElem( s c, 2, h ead.x, h ead.y );
16 }
17
18 KERNEL(upda eCell, 2,
19 OUT, Hi Tile_ loa , ds ,
20 IN, Hi Tile_ loa , s c) {
21
22 in ow = h ead.x + 1;
23 in col = h ead.y + 1;
24 hi _ ileElem( ds , 2, ow, col ) =
25 ( hi _ ileElem( s c, 2, ow+1, col )
26 + hi _ ileElem( s c, 2, ow-1, col )
27 + hi _ ileElem( s c, 2, ow, col+1 )
28 + hi _ ileElem( s c, 2, ow, col-1 )
29 ) / 4;
30 }
1/*Tile decla a ion and alloca ion (hos ) */
2Hi Tile_ loa M, Mcopy;
3hi _ ileDomainAlloc(&M, loa , 2, ows, columns);
4hi _ ileDomain(&Mcopy, loa , 2, ows, columns);
5
6/*Tile ini ializa ion */
7 o (in i=0; i< ows; i++)
8 o (in j=0; j< ows; j++)
9hi _ ileElem( M, 2, i, j ) = ...
10
11 /*Con olle c ea ion */
12 Con olle c lGPU, c lCPU;
13 C lC ea e(&c lGPU, COMM_GPU, 2);
14 C lC ea e(&c lCPU, COMM_CPU, 4, 7);
15
16 /*Tile binding and c ea ion (de ice) */
17 C lA ach(&c lGPU, &M);
18 C lIn e nalC ea e(&c lGPU, &Mcopy);
19
20 /*Ke nel launching */
21 domain1 = C lTh eads( 2, ows, columns );
22 domain2 = C lTh eads( 2, ows-2, columns-2 );
23 o (i e =0; i e < num_i e a ions; i e ++) {
24 C lLaunch(&c lGPU, copyCell, domain1,
25 2, Mcopy, M);
26 C lLaunch(&c lGPU, upda eCell, domain2,
27 2, M, Mcopy);
28 }
29
30 /*Unbinding and eeing esou ces */
31 C lDe ach(&c lGPU, &M);
32 C lIn e nalDes oy(&c lGPU,&Mcopy);
33 C lDes oy(&c lGPU); C lDes oy(&c lCPU);
34 hi _ ileF ee(&M); hi _ ileF ee(MCopy);
Figu e 3: Example o he ke nel cha ac e iza ion and de ini ion (le ) and main code ( igh ), o a s encil p og am implemen ing an
i e a i e Jacobi PDE sol e o he Poisson’s hea di usion equa ion. The ke nels a e usable o bo h CPU and GPU Con olle s.
ke nel; (c) The index space o he h ead se ; (d) The numbe o pa ame e s equi ed by he ke nel; and (e) The eal pa ame e s
o he ke nel execu ion.
The dimensions o he h eads index space a e speci ied using a C lTh eads s uc u e ha s o es up o h ee in ege alues
ep esen ing he ca dinali y on each dimension. I is equi alen o he dim3 ype in CUDA, bu i is used o bo h, GPU o CPU
ke nels independen ly. The a iable pa ame e s should be ile a iables associa ed o he Con olle , in e nal o bounded, o any
alue o he p ope ype o he INVAL pa ame e s.
In e nally, he execu ion o a GPU ke nel implies: (1) he c ea ion o he small handle s using he o iginal ile handle s
in o ma ion; (2) The use o he cha ac e iza ion model o selec he g id and h eadblock geome ies; and (3) he execu ion o
a w appe unc ion. This w appe eo de s he h ead indexes o be used as in he CPU ke nels code o access da a elemen s
in ow-majo o de , and ensu es ha h eads ou side o he equi ed index space e u n immedia ely be o e execu ing he ke nel
code (see an exce p o his code on he le o Fig. 2). Fo CPUs, he ke nel execu ion in e nally implies he pa i ion o he index
space in coa se blocks. Figu e 2 ( igh ) shows an example o a simpli ied code ha pe o ms his pa i ion. The ke nel launching
(line 18) calls o an inlined unc ion ha is gene a ed when his ke nel implemen a ion o CPU-co es is decla ed. A inal global
synch oniza ion is needed o p ese e he launching seman ic be o e p oceeding o launch ano he ke nel.
4.6 P og amming example
In his sec ion we p esen a p ac ical applica ion o he lib a y concep s desc ibed in he p e ious sec ion by using an example.
Figu e 3 shows he implemen a ion o a S encil 2-D applica ion. On he le o he igu e, i s , we decla e a he beginning o he
p og am he specialized Hi map a ay ypes o be used in he p og am (see line 3 in Fig. 3 le ).
Ke nel cha ac e iza ion: Two ke nels a e cha ac e ized a lines 6 and 7 o Fig. 3 (le ). The i s ke nel cha ac e iza ion
uses he de aul con igu a ion, while he second one looks o a pa icula con igu a ion o he upda eCell ke nel. This las
cha ac e iza ion is based on he code p ope ies. Acco ding o he classi ica ions p esen ed in [15, 20], he global accesses o he
p og am sugges an almos ully coalesced memo y access pa e n. The a io o a i hme ic ope a ions pe da a elemen is low
because less a i hme ic ope a ions a e done compa ed o ead ope a ions. The a io o da a sha ed be ween h eads o he same
block is high, because mos neighbo alues a e eused ac oss hei compu a ions.
Ke nel de ini ion: The ke nel de ini ions s a in lines 10 and 18. The ke nel de ini ions include he name o he ke nel, he
8
1/*Recu ence equa ion ke nel */
2KERNEL_CHAR(kRecu ence, 2, ull, high, low)
3
4/*Black Scholes ke nel */
5KERNEL_CHAR(kBlackScholes, 1, ull, medium, low)
6
7/*S encil Jacobi ke nels */
8KERNEL_CHAR(kCopy, 2, ull, low, low)
9KERNEL_CHAR(kUpda e, 2, medium, medium, medium)
10
11 /*GPU ma ix mul iplica ion ke nel */
12 KERNEL_CHAR(kMa ixMul , 2, ixed-squa e-32)
1/*Ke nel w appe o cuBLAS ma ix mul . */
2KERNEL_GPU_WRAPPER( Hi Tile_ loa A,
3Hi Tile_ loa B,
4Hi Tile_ loa C ) {
5cons loa alpha = 1.0 ;
6cons loa be a = 0.0 ;
7cublasHandle_ handle;
8
9cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N,
10 B.dimx, A.dimx, A.dimy, &alpha,
11 hi _k ileRawDa a(B), B.dimx,
12 hi _k ileRawDa a(A), A.dimx, &be a,
13 hi _k ileRawDa a(C), B.dimx );
14 }
Figu e 4: Cha ac e iza ion o he gene ic o GPU specialized ke nels o he case s udies (le ). Example o ke nel w appe o
execu e a specialized GPU lib a y unc ion ( igh ).
numbe o pa ame e s, a speci ica ion o each pa ame e o i s inpu /ou pu mode and ype, and he pa ame e name. Bo h a e
gene ic ke nel codes alid o any kind o de ice as he KERNEL p imi i e poin s ou . The code in he ke nel bodies is execu ed
in pa allel on he de ice by as many h eads as we choose in he main p og am. Ano he ke nel de ini ion example can be seen on
Fig. 4 ( igh ). In his case, i is a speci ic ke nel code o GPUs ha calls a unc ion o a specialized lib a y o NVIDIA’s GPUs.
Da a managemen : All he accesses o he a iables in he gene ic ke nel body (ei he eads o w i es) mus be done h ough
he Hi map unc ions. Fo example, line 14 o Fig. 3 (le ) shows he copy o an elemen o he ma ix s c on he ma ix ds . I
uses he hi ileElem unc ion o access he da a elemen co esponding o he indexes assigned o he logical h ead ha execu es
he ke nel. We can see ano he example on line 24 o Fig. 3 (le ). The ds ma ix elemen co esponding o he h ead indexes is
upda ed using he alues o he neighbo cells in he ma ix s c.
We show on he igh o Fig. 3 he code o he main unc ion ha coo dina es he applica ion execu ion.
Da a s uc u es: Da a s uc u es a e c ea ed and ini ialized in he i s pa o he main p og am (see lines 1 o 9 in Fig. 3
igh ). The Mma ix is alloca ed on he hos , whe e i is ini ialized. Thus, we use he hi ileDomainAlloc unc ion o c ea e he
da a s uc u e. Howe e , he Mcopy ma ix is an in e nal a iable, and does no need o ha e alloca ion in he hos memo y space.
Only a i ual ep esen a ion is c ea ed, wi hou assigning ac ual memo y, by using he hi ileDomain unc ion.
Con olle en i y: Lines 13 and 14 o Fig. 3 ( igh ) show examples o he Con olle c ea ion. One associa ed o he hi d
GPU o he sys em, and ano he one associa ed o he subse o CPU co es wi h indexes in he ange 4 o 7. Once he Con olle is
c ea ed, hos da a s uc u es a e binded o he a ge de ices, and in e nal da a s uc u es a e c ea ed in he a ge de ices h ough
he Con olle (see lines 17 and 18 o Fig. 3 ( igh )). Ke nel launching: Wi h he da a s uc u es in he a ge de ice, we can
launch he ke nels. The ke nel would be execu ed by as many h eads as a C lTh ead objec speci ies. We see in lines 21 and 22
how wo C lTh ead objec s a e c ea ed. The i s one includes he domain o he whole da a s uc u e, and he second one he
whole da a s uc u e wi hou he bo de s. A e ha , in lines 24 and 26 he wo ke nels a e launched using hei di e en h ead-
index domains, by using he di e en C lTh ead objec s. A e he desi ed numbe o i e a ions o he ke nels, he con ol o he
esul da a s uc u e is ans e ed again o he hos by unbinding i , and he Con olle is des oyed (see lines 31 o 33) .
To summa ize, we obse e ha he inal p og am p esen s an o ganized sequence o p og amming phases, ha leads o a clea
s uc u e. Easy gene ic guidelines o p og am wi h his lib a y can be deduced om his simple example.
5 Expe imen al s udy
This sec ion desc ibes he expe imen s we ha e ca ied ou o check he unc ionali y, and o e alua e he po en ial pe o mance
issues, in oduced in he Con olle p o o ype implemen a ion. We also e alua e he de elopmen e o o using he Con olle s
p o o ype when compa ed o di ec ly using common na i e p og amming models (CUDA o OpenMP).
5.1 Case s udies
As case s udies we ha e selec ed he ollowing p og ams. All o hem wo k wi h loa ing poin numbe s. F om he CUDA Toolki
Samples, we ha e selec ed he Black-Scholes p og am, and wo Ma ix Mul iplica ion examples, one using GPU sha ed-memo y,
and ano he di ec ly calling he op imized CUBlas ou ine. We also use he s encil p og am ha is shown in Fig. 3. Finally, we
ha e selec ed a code o independen ly apply a i ial ecu ence equa ion o elemen s in a ma ix. A quick desc ip ion and he
mo i a ion o choose hese examples ollows. The cha ac e iza ions o he ke nels using he p oposed model, a e shown in he
le o Fig. 4. The pa ame e s whe e chosen ollowing he guidance p esen ed in [15, 20].
9