P oceedings o he 17 h In e na ional Con e ence
on Compu a ional and Ma hema ical Me hods
in Science and Enginee ing, CMMSE 2017
4–8 July, 2017.
Analysis o OpenACC Pe o mance
Using Di e en Block Geome ies
Daniel Ba ba1, A u o Gonzalez-Esc ibano1and Diego R. Llanos1
1Depa amen o de In o ma ica, Uni e sidad de Valladolid, Spain
emails: [email p o ec ed],[email p o ec ed],[email p o ec ed]
Abs ac
OpenACC is a pa allel p og amming model o au oma ic pa alleliza ion o sequen-
ial code using compile di ec i es o p agmas. OpenACC is in ended o be used wi h
accele a o s such as GPUs and Xeon Phi. The di e en implemen a ions o he s an-
da d, al hough s ill in ea ly de elopmen , a e p ima ily ocused on GPU execu ion. In
his s udy, we analyze how he di e en OpenACC compile s a ailable unde ce ain
p emises beha e when he clauses a ec ing he unde lying block geome y implemen-
a ion a e modi ied. These clauses a e he Gang numbe , Wo ke numbe , and Vec o
Size de ined by he s anda d.
Key wo ds: OpenACC, GPU, block geome y, h ead geome y
1 In oduc ion
OpenACC is an open s anda d in ended o au oma ically pa allelize sequen ial code and
manage i s execu ion in accele a o s like GPUs o Xeon Phi cop ocesso s. I de ines a
numbe o compile di ec i es, also called p agmas. The main goal o OpenACC is o
educe bo h lea ning and coding ime in a po able way [1]. The e sion o he OpenACC
s anda d a he ime o w i ing is he 2.5 [2].
The OpenACC s anda d was ounded by N idia, CRAY, CAPS and PGI. The numbe
o membe s now is la ge , including bo h academic ins i u ions and companies like he Oak
Ridge Na ional Labo a o y, he Uni e si y o Hous on, AMD, and he Edinbu gh Pa allel
Compu ing Cen e (EPCC), among o he s.
The e a e se e al compile s ha implemen he OpenACC s anda d. The PGI com-
pile , de eloped by he Po land G oup (subsidia y o N idia) is being dis ibu ed as pa o
c
CMMSE ISBN: 978-84-617-8694-7
Analysis o OpenACC Pe o mance Using Di e en Block Geome ies
he N idia OpenACC Toolki unde a ee 90-day license. C ay Inc. has i s own OpenACC
compile , only a ailable o use wi h hei supe compu e s. Pa hscale Inc., a so wa e de-
elope o compile s and mul ico e so wa e, also has an OpenACC implemen a ion, he
ENZO compile .
Among he many academic o open-sou ce al e na i es he e a e he OpenUH compile
[3] by he Uni e si y o Hous on and accULL [4] om Uni e sidad de La Laguna (Spain).
This wo k p esen s a s udy on he impac o di e en alues o he clauses ha a ec
he unde lying block geome y o OpenACC-gene a ed code. GPUs a e e y sensi i e o
he geome y o he h ead-block chosen [5], and OpenACC makes use o he e ms “gang”,
“wo ke ” and “ ec o ” in o de o de ine di e en le els o pa allelism. Acco ding o [6], he
speci ica ion is ambiguous and his unc ionali y depends di ec ly on how each compile is
implemen ed. In his wo k, we measu e he impac o he choice o an app op ia e h ead-
block geome y when unning a ep esen a i e benchma k. By de aul , he geome y is
decided by he compile unless he speci ic clause is used inside he OpenACC di ec i e.
Ou aim is o compa e he esul ing beha iou among di e en compile s and op ions. Fo
h ead-block geome y es ing, we will modi y he mos ep esen a i e o he benchma ks,
es ing se e al combina ions o alues o he clauses o speci y gang, wo ke s and ec o s,
and analyzing he di e ences in execu ion ime o each compile . This will o e some
insigh on he implemen a ion o hese clauses on each compile .
Ou con ibu ion shows ha he decisions made by each compile is no always op imal,
bu manual uning o he di e en alues is no always possible o e e y compile .
The es o his pape is o ganized as ollows. Sec ion 2 desc ibes he selec ed compile s.
Sec ion 3 shows ou selec ed mic obenchma k o es ing he beha iou o he gene a ed code
when modi ying he block geome y. Sec ion 4 con ains he esul o ou analysis abou he
impac on pe o mance when changing he unde lying block geome y. Finally, Sec ion 5
concludes ou pape .
2 A ailable Compile s
We men ioned se e al compile s in he In oduc ion. In his sec ion we desc ibe wi h mo e
de ail he compile s we we e able o use o his s udy.
2.1 PGI Compile
The PGI Compile [7] is being de eloped by The Po land G oup, being owned by N idia.
This compile is equen ly p esen ed in webina s, wo kshops, and con e ences.
A he ime o w i ing his pape , he PGI compile is a ailable o download as pa o
he OpenACC Toolki om N idia. This oolki includes a 90-day ee ial, he possibili y
o acqui ing an academic license o a whole yea , o buying a comme cial license.
c
CMMSE ISBN: 978-84-617-8694-7
Daniel Ba ba, A u o Gonzalez-Esc ibano, Diego R. Llanos
2.2 accULL
The accULL [4] compile de eloped by Uni e sidad o La Laguna (Spain) is an open-sou ce
ini ia i e. accULL consis s on a s uc u e o wo laye s con aining YaCF [8] (Ye ano he
Compile F amewo k) and F angollo [9], a un ime lib a y. YaCF ac s as a sou ce- o-sou ce
ansla o i, while F angollo wo ks as an in e ace ha p o ides he mos common ope a ions
ound in accele a o s.
2.3 OpenUH
The OpenUH [3] compile , de eloped by he Uni e si y o Hous on (USA) is ano he open-
sou ce ini ia i e. I makes use o Open64, a discon inued open-sou ce op imizing compile .
3 Mic obenchma k Desc ip ion
In OpenACC, block size is de ined by gangs, wo ke s and ec o s.Thei choices a ec he
pe o mance on memo y-bound applica ions. To s udy his issue, we a e going o use a
e y simple ma ix addi ion implemen ed bo h in CUDA and OpenACC. Ou decision is
made by he ac ha he p oblem is emba assingly pa allel, memo y acceses a e pe ec ly
coalesced, and he compu a ional load pe global memo y access is low (memo y-bound
applica ion).
The s anda d only de ines he di e en le els o pa allelism, bu i is up o each imple-
men a ion o decide how a e hese le els exploi ed in he ac ual a chi ec u e. In o de o
ob ain compa able esul s, some de ails should be aken in o accoun . The CUDA e sion
needs o be implemen ed using elas ic ke nels, using a ixed numbe o blocks, which equals
he gang numbe in he OpenACC code. Also, since he OpenACC s anda d es ablishes
ha g id dimensions depend on he use o he collapse clause wi h nes ed loops, we ha e
decided o make a one-dimensional g id. We e alua e a numbe o blocks in he g id anging
om one o 2048 (using only powe s o wo). The sequen ial code can be seen in Fig. 1 and
he CUDA ke nel in Fig. 2.
In OpenACC he X dimension o he CUDA block ansla es o ec o leng h, whe eas
he Y dimension equals he wo ke numbe . We ha e also decided o use 512 h eads pe
block, ying each possible combina ion o X and Y dimension alues using powe s o wo.
4 E alua ion
In his sec ion we analyze he impac o di e en choices o he geome y o he unde lying
h ead-blocks in he OpenACC gene a ed code.
c
CMMSE ISBN: 978-84-617-8694-7
Analysis o OpenACC Pe o mance Using Di e en Block Geome ies
#p agma acc da a copyin(p_A[0:Size*Size],p_B[0:Size*Size]), copyou (p_C[0:Size*Size])
{
in j;
#p agma acc ke nels
#p agma acc loop independen gang(GANG), wo ke (WORKER), ec o (VECTOR)
o (j = 0; j < Size*Size; ++j)
{
p_C[j] = ALPHA*p_A[j] + p_B[j];
}
} //End da a egion
Figu e 1: Sequen ial Code o he Ma ix Addi ion
__global__ oid ma ixKe nel( loa * p_A, loa * p_B, loa * p_C)
{
in i e a ions = ((SIZE/blockDim.x)*(SIZE/blockDim.y)) / GANG;
in i e ;
o (i e = 0; i e < i e a ions; ++i e )
{
in x = (blockIdx.x + i e *GANG)*blockDim.x + h eadIdx.x;
in y = h eadIdx.y;
in i = ( x/SIZE)*blockDim.y + y;
in j = ( x%SIZE);
in o se = i*SIZE + j;
i (o se < SIZE*SIZE) p_C[o se ] = ALPHA*p_A[o se ] + p_B[o se ];
}
}
Figu e 2: CUDA Ke nel o he Ma ix Addi ion
c
CMMSE ISBN: 978-84-617-8694-7
Daniel Ba ba, A u o Gonzalez-Esc ibano, Diego R. Llanos
4.1 Expe imen al Se up
We used a N idia GTX Ti an Black o un he expe imen s. This GPU con ains 2 880
CUDA co es wi h a clock a e o 980MHz and 15 SMs. I has 6GB o RAM, and Compu e
Capabili y 3.5. The hos is a Xeon E5-2690 3 wi h 12 co es a a clock a e o 1.9GHz, and
64GB in ou 12GB modules.
The PGI compile is he one con ained in he N idia OpenACC Toolki , e sion 15.7-
0, published in Jul 13, 2015. We used OpenUH e sion 3.1.0 (published in No embe
4, 2015), based on Open64 e sion 5.0 and using GCC 4.2.0, p ebuil , downloaded om
he High Pe o mance Compu ing Tools g oup websi e [10]. accULL is e sion 0.4alpha
(published in No embe 28, 2013), downloaded om Uni e sidad de La Laguna’s esea ch
g oup “Compu aci´on de Al as P es aciones” [11].
4.2 Block Geome y Sensibili y o Gene a ed Code
Obse ing he esul s ob ained using he CUDA code in Fig. 3, we can see ha he bes
esul s a e ob ained when he block numbe is high enough o make he GPU each p ope
le els o occupa ion. The X dimension plays a huge ole, bu we expec ed o see an im-
p o emen in pe o mance when he X dimension was a leas 16. We suspec his is due o
se e al ac o s, being he mos impo an he beha iou o he cache when elas ic ke nels
a e used.
The esul s o he OpenACC code gene a ed by PGI (Fig. 4) show ha he only ac o
a ec ing he pe o mance is he gangs numbe , whe eas he a ia ion o wo ke s numbe
and ec o leng h does no play a signi ican ole excep when a ec o leng h o one is
used. In his case, pe o mance is se e ely a ec ed, which indica es a poo use o he cache.
O e all, PGI’s beha iou is he closes one o CUDA o his example.
When using OpenUH as he compile o he OpenACC code (Fig. 5), all h ee pa-
ame e s a ec he pe o mance o he gene a ed code, which means ha he pa ame e s
a e ac ually being used o map he compu a ion o he a chi ec u e esou ces. The e is an
excep ion when any o he h ee pa ame e s is se o one. In his case, OpenUH seems o
assume di ec con ol and choose wha i conside s an adequa e se o pa ame e s o gangs
numbe , wo ke s numbe and ec o leng h.
Finally, he esul s ob ained by using accULL (Fig. 6) as ou OpenACC compile show
ha none o he pa ame e s ha e any e ec on he pe o mance o he gene a ed code, and
i is he compile i sel who decides he alue o be used.
5 Conclusions
Du ing his wo k, we ha e ealized ha he OpenACC s anda d is e y unspeci ic abou
how he di e en compile s should implemen he h ee le els o pa allelism. This allows
c
CMMSE ISBN: 978-84-617-8694-7

Analysis o OpenACC Pe o mance Using Di e en Block Geome ies
y1x512
y2x256
y4x128
y8x64
y16x32
y32x16
y64x8
y128x4
y256x2
y512x1
1 4 16 64 256 1024
Block.y Block.x
Numbe o blocks
Th ead-Block sensibili y in CUDA
Time (ms)
4
4.5
5
5.5
6
6.5
7
7.5
8
Figu e 3: E ec s o he Gang Numbe , Wo ke and Vec o Leng h in he measu ed execu ion
ime using CUDA. Execu ion ime is in milliseconds (lowe is be e ). Da ke is lowe .
B igh e is highe .
w1 512
w2 256
w4 128
w8 64
w16 32
w32 16
w64 8
w128 4
w256 2
w512 1
1 4 16 64 256 1024
Wo ke s, Vec o Leng h
Gang Numbe
Th ead-Block sensibili y in PGI
Time (ms)
4
4.5
5
5.5
6
6.5
7
7.5
8
Figu e 4: E ec s o he Gang Numbe , Wo ke and Vec o Leng h in he measu ed execu ion
ime using PGI compile . Execu ion ime is in milliseconds (lowe is be e ). Da ke is lowe .
B igh e is highe .
c
CMMSE ISBN: 978-84-617-8694-7
Daniel Ba ba, A u o Gonzalez-Esc ibano, Diego R. Llanos
w1 512
w2 256
w4 128
w8 64
w16 32
w32 16
w64 8
w128 4
w256 2
w512 1
1 4 16 64 256 1024
Wo ke s, Vec o Leng h
Gang Numbe
Th ead-Block sensibili y in OpenUH
Time (ms)
4
4.5
5
5.5
6
6.5
7
7.5
8
Figu e 5: E ec s o he Gang Numbe , Wo ke and Vec o Leng h in he measu ed execu ion
ime using OpenUH compile . Execu ion ime is in milliseconds (lowe is be e ). Da ke
is lowe . B igh e is highe .
w1 512
w2 256
w4 128
w8 64
w16 32
w32 16
w64 8
w128 4
w256 2
w512 1
1 4 16 64 256 1024
Wo ke s, Vec o Leng h
Gang Numbe
Th ead-Block sensibili y in accULL
Time (ms)
4
4.5
5
5.5
6
6.5
7
7.5
8
Figu e 6: E ec s o he Gang Numbe , Wo ke and Vec o Leng h in he measu ed execu ion
ime using accULL compile . Execu ion ime is in milliseconds (lowe is be e ). Da ke is
lowe . B igh e is highe .
c
CMMSE ISBN: 978-84-617-8694-7
Analysis o OpenACC Pe o mance Using Di e en Block Geome ies
e y di e en beha iou s while uni ying he basic concep s o au oma ic pa alleliza ion o
bo h GPUs and Xeon Phi cop ocesso s.
Due o he s anda d lea ing eedom o he implemen a ion o he di e en le els
o pa allelism o he di e en compile s, he ma u i y o he la e di ec ly a ec s he
pe o mance o he OpenACC-gene a ed code. Thus we ind ha he PGI compile , being
he mo e ma u e o he analyzed compile s, gene a es code ha esembles an op imized
CUDA implemen a ion. OpenUH shows an implemen a ion ha akes he de ined le els
o pa allelism in o conside a ion, bu wi h oom o a pe o mance imp o emen . On he
o he hand, accULL seems o a oid hand-made changes o he le els o pa allelism, choosing
always he same con igu a ion.
Al hough compile implemen a ions a e no e y ma u e ye , he simplici y o ou mi-
c obenchma k allows us o see he e ec s o he a ia ions in he di e en clauses: Gang
numbe , Wo ke numbe , Vec o leng h. Ou esul s ema k ha he pe o mance boos
ob ained by he es ed OpenACC compile s in GPUs is dependan on he implemen a ion
o hese clauses. Howe e , since he mission o hese clauses is o uni y di e en concep s
among GPUs and Xeon Phi accele a o s, we a gue his is complex ask o he compile
implemen a ions.
Acknowledgemen s
This esea ch has been pa ially suppo ed by MICINN (Spain) and ERDF p og am o
he Eu opean Union: HomP og-He Sys p ojec (TIN2014-58876-P), CAPAP-H5 ne wo k
(TIN2014-53522-REDT), and COST P og am Ac ion IC1305: Ne wo k o Sus ainable Ul-
ascale Compu ing (NESUS).
Re e ences
[1] OpenACC-s anda d.o g, “Abou OpenACC.”
[2] OpenACC-S anda d.o g, “The OpenACC applica ion p og amming in e ace e sion
2.5,” oc 2015.
[3] X. Tian, R. Xu, Y. Yan, Z. Yun, S. Chand aseka an, and B. Chapman, “Compiling a
high-le el di ec i e-based p og amming model o GPGPUs,” in Languages and Com-
pile s o Pa allel Compu ing, pp. 105–120, Sp inge , 2014.
[4] R. Reyes, I. L´opez-Rod ´ıguez, J. J. Fume o, and F. de Sande, “accULL: an OpenACC
implemen a ion wi h CUDA and OpenCL suppo ,” in Eu o-Pa 2012 Pa allel P o-
cessing, pp. 871–882, Sp inge , 2012.
c
CMMSE ISBN: 978-84-617-8694-7
Daniel Ba ba, A u o Gonzalez-Esc ibano, Diego R. Llanos
[5] H. O ega-A anz, Y. To es, A. Gonzalez-Esc ibano, and D. R. Llanos, “Op imizing
an apsp implemen a ion o n idia gpus using ke nel cha ac e iza ion c i e ia,” The
Jou nal o Supe compu ing, ol. 70, no. 2, pp. 786–798, 2014.
[6] C. Wang, R. Xu, S. Chand aseka an, B. Chapman, and O. He nandez, “A alida ion
es sui e o OpenACC 1.0,” in Pa allel Dis ibu ed P ocessing Symposium Wo kshops
(IPDPSW), 2014 IEEE In e na ional, pp. 1407–1416, May 2014.
[7] PGI, “Pgi accele a o compile s wi h OpenACC di ec i es.” h ps://www.pg oup.
com/ esou ces/accel.h m, no 2015.
[8] U. de La Laguna, “YaCF.” h ps://bi bucke .o g/ uyman/llcomp, no 2015.
[9] U. de La Laguna, “F angollo.” h ps://bi bucke .o g/ uyman/ angollo, no
2015.
[10] U. o Hous on, “Open-sou ce UH compile .” h p://web.cs.uh.edu/~openuh/
download/, no 2015.
[11] U. de La Laguna, “accULL.” h p://cap.pcg.ull.es/es/accULL, no 2015.
c
CMMSE ISBN: 978-84-617-8694-7