Controllers: an abstraction to ease the use of hardware accelerators

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