From d89316545d7ded10b5c94016d4ff2c42b0f3796b Mon Sep 17 00:00:00 2001 From: Emily Vo Date: Tue, 18 Sep 2018 20:33:45 -0400 Subject: [PATCH] Finish implementation and README --- README.md | 81 ++++++++++++- img/runtime_vs_size.PNG | Bin 0 -> 78327 bytes src/main.cpp | 192 +++++++++++++++++++++++++------ src/testing_helpers.hpp | 17 +-- stream_compaction/CMakeLists.txt | 2 +- stream_compaction/cpu.cu | 66 ++++++++--- stream_compaction/efficient.cu | 173 ++++++++++++++++++++++++++-- stream_compaction/naive.cu | 68 +++++++++-- stream_compaction/thrust.cu | 24 ++-- 9 files changed, 539 insertions(+), 84 deletions(-) create mode 100644 img/runtime_vs_size.PNG diff --git a/README.md b/README.md index 0e38ddb..d9dec4a 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,81 @@ CUDA Stream Compaction **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 2** -* (TODO) YOUR NAME HERE - * (TODO) [LinkedIn](), [personal website](), [twitter](), etc. -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +* Emily Vo + * [LinkedIn](linkedin.com/in/emilyvo), [personal website](emilyhvo.com) +* Tested on: Windows 10, i7-7700HQ @ 2.8GHz 16GB, GTX 1060 6GB (Personal Computer) +Updated the CMakeLists.txt to sm_61. -### (TODO: Your README) +### PERFORMANCE ANALYSIS -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +* Compare all of these GPU Scan implementations (Naive, Work-Efficient, and Thrust) to the serial CPU version of Scan. Plot a graph of the comparison (with array size on the independent axis). +![](img/runtime_vs_size.PNG) +* To guess at what might be happening inside the Thrust implementation (e.g. allocation, memory copy), take a look at the Nsight timeline for its execution. Your analysis here doesn't have to be detailed, since you aren't even looking at the code for the implementation. Write a brief explanation of the phenomena you see here. + +The first thrust run (power of 2) is much slower than the second (non power of 2), and much slower than all implementation at all array sizes except those greater than 2^20. This might be due to the first instance of invoking thrust requiring a lot of extra time to set up the library and any utility classes. + +The CPU implementation is surprisingly fast compared to the parallel implementations. This is likely due to the lack of overhead in kernel invocations, and speed of the cache, as the CPU processes the array sequentially. + +The work efficient implementation is slower than the naive solution. This is surprising, as the naive implementation does less work, but it appears that once again, like the CPU implementation, the fewer kernal invocations, as well as the lack of copying memory between the host and device, makes the naive solution faster. As the arrays get larger, I would expect this amount of overhead to stay constant, and so I would expect the work efficient implementation to eventually be faster. + +When we look at the curves, we see that the CPU implementation grows linearly as the array size grows (the line looks exponential as the x axis increases exponentially). This is expected, as the number of operations is linear. The other curves grow much slower. The Thrust implementation stays almost constant, indicating that the majority of the work it does is not actually related to computing the scan. + +* Paste the output of the test program into a triple-backtick block in your README. +``` +**************** +** SCAN TESTS ** +**************** + [ 0 1 2 3 4 5 6 0 ] +==== cpu scan, power-of-two ==== + elapsed time: 0.000365ms (std::chrono Measured) + [ 0 0 1 3 6 10 15 21 ] +==== cpu scan, non-power-of-two ==== + elapsed time: 0ms (std::chrono Measured) + [ 0 0 1 3 6 ] + passed +==== naive scan, power-of-two ==== + elapsed time: 0.014464ms (CUDA Measured) + passed +==== naive scan, non-power-of-two ==== + elapsed time: 0.012608ms (CUDA Measured) + passed +==== work-efficient scan, power-of-two ==== + elapsed time: 0.116896ms (CUDA Measured) + [ 0 0 1 3 6 10 15 21 ] + passed +==== work-efficient scan, non-power-of-two ==== + elapsed time: 0.08192ms (CUDA Measured) + passed +==== thrust scan, power-of-two ==== + elapsed time: 4.4583ms (CUDA Measured) + passed +==== thrust scan, non-power-of-two ==== + elapsed time: 0.014304ms (CUDA Measured) + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 0 3 2 1 3 2 0 0 ] +==== cpu compact without scan, power-of-two ==== + elapsed time: 0ms (std::chrono Measured) + [ 3 2 1 3 2 ] + passed +==== cpu compact without scan, non-power-of-two ==== + elapsed time: 0.000365ms (std::chrono Measured) + [ 3 2 1 3 ] + passed +==== cpu compact with scan ==== + elapsed time: 2.09687ms (std::chrono Measured) + [ 3 2 1 3 2 ] + passed +==== work-efficient compact, power-of-two ==== + elapsed time: 0.088352ms (CUDA Measured) + [ 3 2 1 3 2 ] + passed +==== work-efficient compact, non-power-of-two ==== + elapsed time: 0.081344ms (CUDA Measured) + [ 3 2 1 3 ] + passed +``` diff --git a/img/runtime_vs_size.PNG b/img/runtime_vs_size.PNG new file mode 100644 index 0000000000000000000000000000000000000000..525b037592f2cf7333c53aa36386b03ea098960e GIT binary patch literal 78327 zcmeFZXH=BS)-~GXoMTJQNQM@X)Z~l;N)*XRqa-B>k~2s~P=e4zB?%2Es32KDKpPN| zAUR6T8T<omkH_agUQN{_Vno9x>dTVrKzKX9((Vjrlw|)^=l>*3}Rq@ zc6aac!C){}*6e};6&h%MS(%WC$ogP~jnGpq8bPL(mX?tkm-whCC0yyuzf>S%58*!77ySl5UC zetlI{ReSq|^>q|qpOS{gM~CwF17l-ln|5I6(NR%N{@-8Sx$*F!)UKnK*OeKJN+}VM zY(7B=FD-$4Q6iZO3-;RWjd2{m7OPgYCvvF5Cc?~&lUjAz^u~?6M^A!jreKj-6iN+S zpQZWAV^2>{Nf<6KkG*HWEqFwH-PkyC;b%u8#{hxHmoHylT+4j8^VKNFtxJqkV*N=A z1DT6l9f3jb_gnS8-Ydhe6B84`=exY*qkcU|8QkZ7xU9CJXx4R=VTbQgg!oUjG&MQ- z%li8IDqY6mt(jf~IleYFdC#1}?($ZkO52>TRHj_bxSaENe}6wP5TaZL9&*F|>9+Si zdvVON7pSR2tG$E-1wDw`_0-ZWyItPz7@C=l${@U%uH1uU-fHkSn?*!yv@y6mlM7(p z`Y8BLl1+(QD-+DpG=iPA7*u5-wx|+l$1de-ln*g)s}0?W`Q->+%!x~P7RA1pCem*RA+z1nvT3G zV)&Qd-XLRg3bsC??KC3`KN>?;gaj84j}oNxM%2TH4>?Pf+0y&Jed~JPAmfKy!79QpPNrduBK zwtH51D8E~vs`5J7`L+=QwW3dTv9o(lqtRk6(tl86YG4p%=w;~5pd0t*&6`*T39vsj zqxQ+F`OCPrFyFp?tD>T!D^P>cP4=A_R{Zj%!GTE+Z5FQEV}65W_?PkVc#==g?q4BF zw^7{{?9`K9RuG{3F276a89O4~KGhsXH24-GYdo!n?qjeMG!G)klJTZ=A#-4oN#rma zq_B6eU?=jHm|BTo!V3yrk3?JIh*$=`)z{b0mc$LTn$GlG?7tX8I6D|siRnd?fkOV| z_OXTPT;M$00{3ydCSa~BuNOKo{cugNn+CW>BMm98u=hL_kZ0;L$M%VG{Jlhg=S14YOzjag7479hpb}g<`4!=71;Ig>5h9}st zmo8B!aX9eeMT2EIYmi2V8-G##<8dTuPK4jl?&v!YtrU3VDB^pU* zO$6WOZq_mgWn^S(QCDahS6f&?uqudSjID5%Wc6uXi6r?IBfa%w4v8^qVTjMHuCE2G zgWw_D_`zI6Gsj%DlIocUer-#J%hzWNSEfEhM>>y^23qR#wrcUYC@t2iv}9pDZtB9o zFqmeiM&YKaoVT>0>1m)KoW&vzN?#9q4&sMvR?S&P8Uv_lBp?Q7&QjtLTj1b$$JO zF!z|lHCi=ef$L_dXKN-xp|c9>z+KvUh~> z1|gVyJ}OpYelMQn!af084q>Zi|$p?>9m+kIENKaGR5t*YR!MnWT2;CSXvr+=TYK^ zXEI;)kvR=9PH_g25^$`yePOKw_y{UW%F;%V65hJ{w!zqu5jeQ@Ss{2?U$F(K_1Q#4 zDXeWQBqrY0Ie)M7b^`$#e~3$U@8LW|d6GOMCWaV$%F4=$=zzknb!um4C+j&IkAu@7 zO2HQGwll-WO8jx9`%*cP7}>(m(ngxoNVUn)=OyL`O1Lo!53L)N?*`Qxd}gWo8CM+B zR3Bm+&HMJ=B<60UDe_l_oDmgS1nKZgUmyAI*RNkSqQ1TPA{cxtlG1CS;;Dcwf8)#a z@ABQJ<%b7}NKtS5i6u!d%kuC-x+NR4Wm1$gm?M|Ejc|`F=7ibIr9)_-YuyDJniM3` zWz-Q86B}!_+QaE!uAUwDg3}Afrcu4rl*qVa(Ili*%5p~lsc?Ir-f>`fhjz*gf1A!L z5Z}+bFuK8CRj4OECMG!~g4VRg`9sIJ_g>pT)*98sThp%2&gCr}bF$`{nVD?1U2fGS zPY&^4V04q7eh1n!I?BYz7;{Bu)d@$;&gIl0b1Q_GQM5qzY(VNvM>|8ihUR&u=qq1e z4!<<50az*KTIrE9hP}MAbJ;hfYTDAN9T}4h%IL9wwad7(SxOu|)9HlgPthOkOEg=w z_(K)djb*Z~C+|J9a3RQ`>YU*Z57yiQ1&8<2nYy&EA>mFEvI1~x_oasqYSGyjDiiRQ zO>~YF3C6Gr9e)O0>wplxA9K07>Vk_?X`Zr2d}y(M@~OvmZ*|nX+>%1PGSmqKpS$$( zAn|Bppk1Rc*X%W@3rlTNEB-yi)Lm$3H69uBE4)Zo$_wo3t&s!q6@&Z&$^?9Td_CA! zLbgiH?mqnGy(>z7xcy(hvK?YCtYkhI)6mr1+}i^onJ3q&$WrZ;W$oYy_G_S3xdAcx zg@uK6b&Olhx}*LqxHd{T->NzUX0mSX%3H{rnrs~(9|M?aHBe|{HXwH8%A5Pssz;JW z&cwvT9SNu^RM!NhF_n8{SS4wLC1D{e+-+ykz>&yOqKMzduEJkHU8?N>c$7hu*Vsfu1+Z|Y({aE`V?hC()Tl$Q3% ze#K}LVgRf;6KxMbgp-`)}`C7815#%B{-#+)oSvqRZ~ z@U&aiQOBP1_C;E@w}S)!OjG*P6MroHPaN#*)s90`JJooYv)#YdGHRCl8KK?+y->BHp5^m)QY~keQ#>K{_?=kxow7z2b+5IRgDIp>zcCS78 zc3o6T>YdyC=hRgCrBy^LPe*t6_SV)&z2DKT2PUk)ZY3TAd?GQr0{{a9e~lkMehh&m z@r|MMR}dtx9M1m#M#U()0t0)wq`W*nHufw{^1TaaC@#MHsTN>TViJ-fx?hO8fy%*WcUezmZ*Ba&&x7_jCK3`7KAmS_caTle0-zKGCF3bKf11kezdq(I@b29^ zun5IKxwBc-gLw97^O*~-PsNcosFw^3Qab;t%E_L;c)`!lkA;PGa&iLxP z)5w#Aguob^l6>E?{a^8yZb*^4*Qbql!8Ta=A*(_Cz4vQdlzb&k!q{PO?c0Cg3 zhix3fGMRUj| zq;zi{5syjH37gk>u3R{^n`rVPQv5#4^(!W{lRvz7{FsenOOCRWTi=Q7(W6H*y14R< zg6)XXLBD(V$nPjmQkH9p>mWaV{HSsD>g)EKS)=Ur5uB`wN>Pl3P26*9hVo2_RQZ!U z6ciMLzhyMl8A^%&+u@um@>ouLuMu>188m#U_*9y zJyY0*uRK`&zQ6yblbmFGd)qh%B})4PhL0?5!P#O#g2_)IlAC+fvB6iL#LmI7@pWxT zP=4FO%q%vn@^X;#%d)w}fh5BG74YOvUrxhVg@l9v3OwOh*HTvx)XOIKSl+=AC9gS| zrM)C{yiY9u6IE1ND$w7q3`*?V(Z(%+<$#9kKX$EbHlcurJTP{05RrttGAy~02y{@kXx`W`?^;A^_04|t3cFW2kl$xC+qRCK1-H{Ya)$m=j z7fQwVv*nV9hX>;GGf+W5T8t-lU{q98qN1V;S1TjE2g=O&CGv_vM7a6cw58E@kev-s+ zj3$@YpL=jo&f~YfxB(o_dlE`HMo0|6y0t<&J3ABF=2LkY z`@F)xebmKR#&6>)iLkO1xqi)1tZNz;Pw3YM*@lkxAiCv|y(E5Z#Nfq#F4{}XXxSMz z^b`!)^HU(&cs_%|n)r}duLt3-gcQZ{Q{ZgpYB~FT&wUkv1Ux;{3<;{=T>(BbHw{?>{FF@VHa-3ovDc zf8XoBZakd4T^`9s`0u~@?;8lcf_v_Lv$cx$AR>eR{Y?J#fQZ@Rg5XNj|D2N^1pa+6 zH@}VT|G^Lv4gggMcGCLR7XQVIXBn1^gd{8?;(@1Uv5X1y{JTnH@y&1?sdjRA*Plh? zpi8ySnT1&4!|v^EcMu>gEwdXM0`A^j92}&_bZGrfKtRAd1b9g!Bf4VoOP7|1s?^ex zlar0_SlzpK?`U_$-`^i3+xF_Dq<^z^AeRu$?7nxrM=*ifpO|fQ=Nj$pcv7x_yul;3F&U)IF-+1sfmb) z;4p9Z;o#uFf(PPD1FG5Teb+3XkdW#EF+Vo%jk6vv2C& z&Q23y54KYCAmm|qby!$f4naQAV(>fwQ~-`E24%B@$TE0|@*WQ_o>Y^UE{hLiON(pN zj)~U_bha?&^Ju;9ayvq#JU5q}laxs+)Fh8Qr9y>M0Z;*ydT)eJv}$?K{zkp)tL}){Y=I=jc!+XnvSak%(S{NlHo<6&0N& zlC_tn9T07-Ah_m?JjrFzfSOC6?}!s24hYu8lKUIZz6!Gt}y8m1Elm!4h< z#~W%C0f~dn?0O^yR?Ov#7YQ=+Bp2m_$UkN?nx>_uMgjt3#ab#=Y!QOC@>zir=x2ec zmaL8wY(Cjt88On+3yN)1k(akaDzUrDw5P9+l9qOEe!j-;Q+h%|Li;YX76Q?tT&L{gRk-PkyECgV z^rQ7clI;o%X8@7_@%Hw{x2deGOixeuTp14QER3?!ioh(^FtX=jP|0B&-5iGTb?GLe z9-?&c*r}O*wU)P_~y7_WT zNnJq7UBBW(e%r2lK~hMFls$rQfsrZ#9TOM#&dCTxhPE!t(lhZ@?Zl}?cL=bG(DPNN zf;lg})dDs+&llzNb_DTx!SG2JLaN~?_JCz;KjuX7>G$v7>*{2%2X>bboTqL7%A^z; zYv12~VQx~=CpR@RG|Yt;N1?quqrD(({%&qnw@xd(x9rN{SQjrqp%o2!vP;XJO#QQqvmJ8wuA4iHHEtgc{x0*iecf^?2}L zBmyC7gf3kE0W&U(jEn@h^`oJ_J~l3H=TqtG=$BpI4yhN*k)j0WEQA~LeIl@$Z;O(7c_8;}q|&g0_ZDqc7}$;n|8FhU?+bW*W| zo|F35eX9)c>shOvMlMT96`vKijF5949SBr+?CsUn)d%1E)F6=-%F3;cjSxaSJUmd0 z#3fBkSo%V8u`v$gn55UAWoDXz7?=deh*z&(Js9)GyVI|56-cIl&MGRpm~8Cs;Sm-d zUS?8-i&oFG_4Snj`6~rf&0Sqxy=3y3Ce(mI0TCl-7|5hxV@Ws-hblVuzcj`0==<_z z^fJloMt9xd&Eesj?x4T~nLboPj*pXv$H>4S=cbMhnO!I^PDYEW9?M7)h4s~JUXX!R ze!?=m=|B>ux`9Ps_diK^lbumfvc9~c*cuK}hOz2)&)^{Mw2p=bl$kkeC9o9#4>{QD zq$*F%-|w`&^@EXt6B&6T9u5w`1e5PE9BUnx5XhdsjrHgR(K4m4}}Gf+AW(UuJJ_&*l3Y3PS9} z!jh5}Ur;#$O2f^~O;n#()7^a=tfo17N=kNTo5|OR38WO5c*%pQl9ir? zM?|nWULWrxK1w$RFymJ`MwvM>nTD1vkJJFDfm!ktWDu3M)RzTHR{HwXSzKCAn9d0> z)5J)H+=2ow0fBBb+T~NaI3P%>r1E@ms~=}&#TiD<&Dz2)CAyVh*Gw>v-_VT&EMMn= z-9pn=bJE7feeAe?6BQ1zbBo7l2urkEK$SI(D25MFJ)%6x=`c!aDoN3 zcZ&Q&mk2BH?WbNN^q*YhV3VETI=aNB2nv!x=bJY#)K}i#nW0bgJ2p287vMN&)quVs?#E$xBc{MxHIJfWQ_4L}{iW0Iejii;^uAS5}IF{Mu{e>svRnk(Qo*LHXE= zfXgs49{X}=E2dwY6r2hOzf7~KvXNr2%Z?PLI+~h_VSt7JK7g12clZ96^_CV?(QhBB z-rj&-q(kXGd8U2!YK{axarh%aZfvsGU`&=PYWIoj-%z}fkdRPPy!Px0$hC7OnFnEL)=DxWqQ}uz2~I99YD~ujdHz6W zPlF-sq%(2=(C>O82;hYmK`?e600O8oWK+qYu9lV-5ci(F3l|no?graSL`+N!fS8dz zV!)Nlg|JyIxtYxdYglo*&guZq2Oy$BX{{Hx4BQThs(#Ahfa>ttU4*&qEPq9S^#eF& zfWA4S32;Gz-V?75J4?{y%;s`c)}Si;t%B>Cf(tDd^`3GN#Ym%b;;ouO%g~<8&l#4NK40foT~zTm z&Oy?L`*(Jk_q(@>2)3b$&oF+^V2As1Y~@XeU46zpP-YXK^LJWb_ENm{dcJa($1j{5 zaSp`f-u(Qv4or^Ry8}?`IG+{GxWDgB`>BFb{r#BndcY{0TWw!G?hS|RcTYZl_=h(z zeZRmI)P1}dtuqd7Sh7Uke|rm5c)$Xo_>A=Tfk|~pn&Kc2ycROlr3Vop$j~w~kBy9! zY}ytzArY3Amg?#qJv|9$gm_Ml6`;D0F>eCJ=Qoz2DPHVB3j%W&!TxFMQ^zz2MHEz2 zH%v`SHrEu9CaULsf~9OL0!?jUuRKiWhfCq6sMrj)*U9lQJ3BkrN4i&lV&!+_nT(Bx zH;qNk1qf&S(vAPuz(JEkW^b>aHXau}B{4=Opc~pge7J3AmwT@zoTNk34TELP7{fgb zqtvX{Fg$egj7Nq>qk#OgY(jn0tz8X(bLy-_RgG)cNGZi`^tTVZyqFJ2bsQY4#KhP4 zmY`UT!`4xd^g>NQ1|=t}TLCEK<_6o1Y}LoJQ@}L2wEW{o)3~@H9FPb%0Ew{k&TvNW z_t0VYu|g;dihQQ*Y2rKg;nCoWF829=I+M(@yFP)s3>XhgO&P`*)~OaSa;=nSfH?p> zdQeP7zO}-Cg3k5_9O|NR)1`j^h=;obm{Z*a$zyCN>BpBi1QRiFB{qr;9~U=V5kKY)NqCUAmFI4xg*#?= zAh$7I2UhyT5{*;A$0#bE2})h+Zjew1fx}Fm`4@TK2(h{XnEX_+>a3?vBeFIq2Rc`l zS{%?~xZ8Pr?};>Z7&|`o7;nijwY!Z3v5WS+eGcp%9C7b+k_^wX0?qA)cYkpe!LOK$It&5 zi#ofk?DHbYgaP(6l9Fa@7(JA6^D$9-@o!%|e{QXsstzzIn3PgnS_~W2n!+fE7(YMJsN&7_>(ZSux3#zLt&BuS^|e_0;>zX)II8e<#K`zP2dvFjQg3H}>1nTWLIK2Z zUOoFvl>`h!+rtN8`|73AMJiPH-h#Lg>ReQEZ&0#sSZpLc9i`z&{Wd*QWMJzp*j! z1$t1fZ{=hp9&M~9S(?YIo+qKMVrWQf-aJrMRizi~>f=Y`;N@AjEAx$UCGSH%DiEOP zRMK=#PEO0~ysag>O6D@~D*fO2#;RA_#AHLlSr{ORyGXE@b|`(_g;2nnoi(D-6VVKs zJ>S0z(0MtI)eU_@Ob`R}$G{H*plua~$XWHPq5S8%(}q6VEv&7LSKihFGPQXyD;p#7 zzi|nU;=S$v04A8TsVhFXXsO6)gfQ}cga2F(97|B)F);L;*$?5WQ{E{=d>kl~TF;la zCPS<^1bH4>2?3rR?5@K2wb?wicy=oiblkVBmm4LTYr!ro`U$<(UdC>vWCHj{_(TpX zn5FdPNE6rJL$t|Fv0|MT#3`R|NT={)v;bvm4}L9|AUp&cz8xGKjERXU27|raY9+rH zjc(*EmCn$h?e)HQ4~0as1iK8zB@->&1<})C4iNU3hoe;z#!4|+>FJM*^gewt2+Vlo zHK1_8O7@{!ZNm^SJ-|hNo_xv|gO@h9w9Fnd%BZQT-rNdH&8hc$ZVvFvt|LnW6Ltwu zzjR9{C#`JR%mxZH)v0M{G%f^u1Mns3_u{s@Z}a38=YBT>o~yx_gkM_(VC;g`q8tU- z+}>iF%%1mbeMJ`L*VAl1(dlQw^QfM7y>=O|2eCfILTUll%74x_WOLoJe z1?B(a^z_BGwFW?LZf?5nRN|$jrQzV1_um&1$!C)GEUngA2E;EQ%LJZ}tnyyG=-tnx zEO5@oka3dmVT1jgoSl~ep8tQ@Bk24=jbi8I#VB~*VE8WLdX0pZEbBc_}$O1W`Ld9UFr>p zU-DNZfZ|zRR(?D;HwRqP`v(VBoZwo#3d*Djsp;u_yuA3}i;%DxA*R_H>A5x*dkFSg zq7GqZ)fZ0wRlY(6|9AJpQyvKv>u!s zdbkTI^ahzE$kRH2DG#|+&AYR;wIymIKzHbSv{D1y2g1=DUb`_yq7-RYa7!0Q8{*A4 z!UkJG+Ne;{izFKLoz4^UIXGY|XO!J3$&G3gYdF$w(&_|a>b zQd|yU)e8zvcWBF{al%sL%qFn~G6gD2FDn;3e~x=CQUI{}n4c;zAxzda*7U}Vcj`OO z0g)RWA7BPdWQ>t0V`ij=&+cjb_*Q~ieM9jwMt#xaq@?xRu^nEZaHFB2+4s@b)>gJ8 zK_~tk$6ES%ND0+q>$AskQTV>Rt8)V4u=1|!3p3m&x`tv;#AC(ym112Ey=d8xT}gctRJCwi>4DS=K>;Gg zaUiohiUP*u@dtBN(-=uHvFp3kjU{57|5d2sju&l1v3d}+8d@&#w`8ViYfXWq%la%9 zv?uI&oRKHGV9iprPD9<4d1C^Ni%kZs$Tkv52?-Uqrlzd>>C8-EKJScf+<5FknH<=; zbrfS{CC?_$v|S{XjuW<<$jX0*8dPSLfFYnS_F#wd_R zLowdJuC4}-C1A_Ws}~oSgk$xf-n|>1bv0qD|916_)gH#t1{7U@1;CoI)upCel$4Z~ z)V#_}K7JTCP`+&dfvnh|Y*=77t%MP`Z+~P)fByW$|NAaAVZpQi%Ee7a7te(EG?mTR ze5Um@K`wKlT=FNM5J^a{4jX4A23VItK|z~NNf{Zn&ZAP2NBDY`-q>zO0S?izR(NRQ+Y_0OPpBS$IXU2X0A zPF#4IDCH5)Wn~;L05}(D-Kft#6rwOUJIi)U!FU`@`dyxD&j#zftBjCA)K=-c4IDAw z{ZC>BlfM-xUAWO`C_>~&`dH;}lLT?@V_^CM)^tEGnfs^$)4@f4{*Zgt-Rd?hFF_Hh zP2))GNTJg9G@+xwg;JxfJ)ANB1h15A60 zs`j%dPdwKr6bLeL`qlVfg$blQeLBtME*+0%l|`R?$`&P$4>SR036zS?kUVCunX)oA zyZ_NwG0Nw1+oyJ37eU-VLHrVca#QBHmv4}Zxb73<#+0GVUMc2$E5H&VUs|N&d}O!Q z3ejx<32x`*TlCq8vbqqgj=+%!`kH{I+1Osn`z*O-$uVUVvfST^pAZ_eqSG zN8aZ<2aJW6b2*2gm7g#hdb|7y+l24E-^AM2HZzJiyXv_pYOdnWVm)qJP+#FO;~y>~ zb9p5GX(OVZlCEIW?GM0jYtUC5egRHp_I$NWnhx+v=eaiGV{gfPxp%@=27=Ui=yA(S z-iFq7gBZY_^UNLIGPzxkCSVq^@rDO)K`c1`J*T3VDdANoUQwNBl0Tp<8xdJqK|1s$ z4~rWrZhuh9w&?pkJOtMNh;+B;m6Msn(Y}KBGdj^kes&+55W;f)A1lwP)@|0wL@CzO zQa7)sAfJAng8~2KnV~W*EJ(uY7w4o7XxF=UiyIn#4!?E)!p_$&W*}1nV?<#$W@Ccq z@?ewBa=VQq!18QOXmol044qLlxX9m5k>lduZ7OMMw;w({l&d>t^qsRtdl+02{x8_3 zvLFM8$L-Bfh7)e5eFzk;1#bM)_7Ot5Nd z+9_Cr>B-J!1z%tbBR}#|a14WPi$yhMW#I45e*RT3_g+6m4~>+vc(=;1^<5CtK+vRT z#*NB(_;T)TRN5*tHb;DPbQiEt8PV5iV-sZ{k@3COR0RbE^fH7Mh9E*3lk|6k7_CTD zm!6d+R1#agS&Wdodi5%39pF`yl9YT4SfQ_=l?zoJ`9^o4T9ck^mrixzx%XgiDqVG$ zV6#XCYiemJbM@^ZBFYZUgP|(Kymt*NFmF;^xbWz8^9WV|5p2Z|`a7}9oFd2-ww~s~ zxuSazQ6T`Q@}Jy9XzG0Wsb2Rz=dKR~kN0n-2|t*kVQQMxik}kFBq4y~l7WRs%6hEw zCZ>?C8X~7awE^vr_x1G!f`DTHVZU-DG9yD;R5a+Gq?FV{TU+I1+?Tov!|YZkM-Bg0 z3WT1Pym>>P9(2vl!>$cM{>esv0b`ngh^Xay9v-oA_|qp(a19ciW#iGdwziCtuGpgk z-@e&!Mrsd_j8y2^Lx!geUVW$*Q78i4ezB2^P47Tv3`&fq%w|jP;QxKj0of4%AuNt|%ud?dOZKD$iM_$V29-vao` zZD0J7NMfRtRBFbVuY0H7wLcAG zQyP-(T6J$ZS4d%03f6I&u>7v!q^ac==dC0P0m5n)W+7Of{!=s_-Eki)G(QsfLe+qC zGclz^Y#C2mh=L1HtN!xt6uMl!8i$5lDlIGXe0z&NCiKW$VlteISiNgN^udumottl) zgB)*Ebaabq1skMR=%+xBf0QP0Q|R65jZ7f zC_#-Eik3-Z%-bZg1j^tB=$be)#MEdwJFfx>Y(|?i6lGqHtZ077A6xFzTK-QM2xZAP zWORA`y7@M6w}QSFwV1f-$PBJdVF3Ycb#?h{kO2Ha=j-aK^DoG||E5iS897K!ldKP{ zKBnDzzHR;|#JFkTx!l9~hvBFB(n^3@7ChtiSBusICf2IP@G`=MR0OVO^L?d`Mtjc@Y z7(hC*w$LK-Rqc)~B|nFIx$^;Y3*e1{a{y=akXJqgz6B_x6bt%4 zU@)S2nCW;hniDBxbPxuvaUKqn1C^1k$rdG&-nu~6iA?YngOUG^pbGN zaqu4?nF%Wc=~@KZmq*-umX7!jkPP!czn^<)13ewUe3Oz99a)-^Q2j z9Ej0r5ndjFw$(=(9;mg>`{G`pqNjgjN89>6A)&ds8Bps7M@I+*0`$X7<_WX0O|I08 zn>YB&@|iZsW zgwN8pdLFPe;Y9)Rx-=$+#U(#&*Vg15vcSKMhbLg?!wZ(9DJV(+euJt1;S7tze845; zPX~Lmp1yWK@&l&{fbMW!M;8}5K>P@=G3p?L(pWF4`6aiuw`a06t#ZH*X^Ah6;)MhH$Yed4zy|Sg85Y~EvoT##+2@@4x^p@I1Ybx~)(H(C5@m+ALY%`xTWb zeE|?|h5&V+SKF-~nqO@>H=j_SSVjXcl`)cv=bHRG5yc-t0-*T%r#HPD-K#av1`8lx zez6(<+4Vqs_ot-jPn$_NM`8_47)N4qGH3b2mpxy;P+hn%JUD1C=Q#BP<~6NAL&_{; zANFH`4Ui^>gi4Z&*;MMm9ZF_qneJ0>d(&{qf))SOrE`yy7jKHZNEbACJ$n{&3P@a5 zvs8*_JfPcgx?PK$M^vaN9%F)QJL+y~y}LtD>9;JF#k2ODNV%%YO4J8nKh?erRCJcW zQ~^ibQyD#cp1H{G|!H2ZIAf?0$^bTLZlhdfpm?SKMau&)iS--rHM$seJe?f2n-Hoel&bYp9;J$CPC8uX(NGWa%H+ zfF5nIEb$V;olrZ#BdtwX&J7g1U|wZYz%qHXO6dF3)}#~3!LdjguRt<@VNB7ln_pVm z^m9>5u_li?@EQ@GQyO{$+2pObRPH0@%-b7jrt@n8GoZ{60cF|(3u)l17m}=2dr<80 zJ9qBj*NzMi1MN#U2N14J64rcmuh*0IiL-#2fA*EAf>6vIC#Nv9zL8N2FR(b3E+N+M zy1Tn;X%RJtgSKUM*G5l3-Op`8-GiklPl$M5S;%3n^Tl**@zLoAwphCHs`1d%RBk>9 z+F*@!CH?`tbfDepO#WIKDnVPOgVe|#1fjZB>nTW${{!sye7)dHhm$)nev;oK++5Cr zB9`qO*|&p|V;fLJRjFGt0e?5J+V5xk6rw?uDDOWrX9ENTB^-M!PQ3eXJpy27T?i&1 z%uRVIZ}pL9{{&J&Y1-bY-2!lxvv3urWo8P>dc}46Yir1}`Qo~Y{nG4k1!}0V{?g*6 zKYse8;s;z!frr5L{4ERg6kYM&V)jD(g`g8I{_QMntyMaW=OW{_h&V*7`<|#3byk$5+0e#sDt`K5)Hr zr|87@G9a1Gku}qzrKcaz1^kgzNsDf|x|5J5H%y+r8UH6~>(>5CS0MVBbIBfG=0G*q1JMc9C zU+KAUpl(ysS;}|r#SlzRr9u6VZ{o?ooP^O!}uQo$7Qrtfw`%x$!kz@$E$b8#+*brKP14 z$EpxIN=hNg!r@6tl|8jI8X-Gh=f1@1Zk>;b^gV5*Dy)n9wPF0l5WjZ3iTYgI&-^dBwk+!!v-e z$;^}$7Y~cC2fe_5x5XAZ5|DKQND_3Zc@^l0?!nwZH3Za6AqP`#@X6zIE)va7pB7FZrk5DCPD2x_|rnj?P#d88z`Qqal`@^z*>@)C-z+tt~AfVO}+$4Hpj^ zJNqY)m(s3*9-)UsKz;BwF-bFHu&o?qTr4-M9k|f9u)zDUZ9o(SjLc8F!hkVVe}c=` z6P&3r3GfH#DPj#$X}q2dZxeuhSO@;>YG77!k1Gu`MGrpv_O)v5Iq>_~+~#~EB?8Eq zTz_DEet6cB@3H{n;=1kS_JP7kOZFWcJJ6b zldrptEBoDmou;te+uhAYy?@Wd;)kD1aj8ohkLJ5C?uHcR!+tdlU?x>Q$ZAJ`;0U;a z;*t_JtjLZWk_B?}pu2y2men20{vu=c&09)9$fYdxiPJoM6_O~g2l{*S*&eF0g|m@t zm=QcJEGkO0Qp(jOF>R_wsr8j~+-QE6Wm5FiG`fhvx-JaV1gRz)R2 z{sdGt-@o6|hCZTXVv-dT3wGV+{@cgQ5_4-BcMwmJUo8Er=V*SwT?Rc-xfd^*|NJ@N2*9#xel2pf2ulOfk;~2g z5h!7{pb@mk4fu`Voj@s`ymp<1$$DTg)X$3+s1carK;k)*l_z4=m5fBxw*!YXM$U}C zs@TizaUC|lGqy2IZ@uP>EPNmOx?ZOdOa~|<@(7>|ilzpDGFtGU4%{sczr>)T%k)Q19!F(Ym1WHAURchyCrD z4j_yx-GMC1|0SXMm|D6hGN03ev^?|dWia~KdoY8qfw~kI3C>BVGD>h20kp3XIGrS+ zTG37CmXHxpw{P8IY)0+xd(}+CfOQA38^)^1#+hDmcY)6+yJtA}lJf1N4de$#B!Jh| zajwQ{dh5hR9i5%QAy#f@9e=Ba7JoM)eQ~8EI~HOo%KhD^tA!&hxSY{$?YJp_*5( zSRDu7dGNF2IZAMGsW3LW=`Ub=EF<3J=<;D{g(j@mojHv{B&~%~CxDY+ov(UohmVff z1Du`K!L620xzV=3#YKjFRh$yrKr5wV-=<#tD7Wb?i9G;-{@moUd{LbGzYhk2LXn1&D5NI(QC!A@hYY2H0UwM2tbB zb?q5_(E_@F2``<27INfGdm*ES2_O@@-rIF{gsfsftI%EWGf&>x0GhN0oE2#uoXvPb zl=buy^Yvl@g~9iG(y8dVzd`&zIOW)?2<59W`*HF2Iu#n-oeo=1({VI;v8ja(x{1Ab z0GZs`8^Z+}YnFcxwKf|q+h*%&--s#={GNC2mD7crWB`lnPrxfkM3ij*7^+ggeMWX!o<9^!;( zg#SuATKnCT(+6LEoMD(R@DQE(!^6YDLC_hrSDZVv&vE67v48WxlCV(kZ!{Q&51&pm z;`Fj>UZ?QIGr}X}4icO{@wvMj90xUCg@{ik21kaes3^FCuBSu}HDF_r6`wtIq5r+l ze4^Hcrx7tgGgow+^hHg*p&S|aEyvGR!ibj;Qu*mM)*%IGVmw! zK5SpxyW|Pvm$k{*|E#?UFccB3EWoF%K5b%NnT@6o{6_=zFVbp*4Y+d(_iUMc+YpP3 zi{#JA{hrDC&>kt?j>JXP`7$toBQvTiD)hk1#igZ#YnLHoANgM297?UFgKy0DK@YYS+^@bs(((#FV?T;Hm3 z+I|fB)|-16p4YZ}L7_c9UM{AL1}=8|-DPfTQkGZEY$Jg}(uHMt#mnKgEQ z^+uY(#~Zbt&4^$^SAGAYAB&5*XaTx-CiyQ2TCB0H|L``+aPNjqgHi-L%jG#kmWy5Z zmdv;|OlG{BSY~_~tNcANQgUb@cT+1?ILqa-tgIO*{lMv40AHo26Y?v`Hvotd8j79l zIQ)v>K6NphaZRoOTG!*R_O6{?9dB>RLuD&KS=-weB9Y_Zbc6~q-o!{2igapoFsK}n zzCqgRXu<_>Y1u1R=*@ue-}UJk;9br-%zOe(2Ug1t01fFejB`qHO z+S{i1PKm24D=L!aDUBC0S#{gy1rLK1!SZ4CSE7zECOrYQu8l}!;L*_$n1|9{!bo@9+sDT)}9ZES_gQrWX-nJiIc9~2?kE6WIDDP&Ey?E4bKknHG$GNVqbB&qL`hLHb=kxJOcMUo5=6M1$It>6Z6{1hpOo4h_ ze!lRzySj2uGB^tZcXWs;X-2YoiCNs@UZtYXK0a?zxX(YULi(52JfX>R#T^!eM9fX1o>#Hv(~Sr7x~B=6?FHx*7Jo_w`*mS2)pUd=%CjV)=qh!h#ef_&3kM|nM zjxw+sJZtfc44d(GkDJi9HgCi~cD-pDRf^VlGeRyo)72E3<9$VT91C>5K;7JH3o-~z zj<@xGpu5&+jouXO&25lG8&*Y|RNj@a^p>+W*CPBAdr@%;n6Y%GpcswYlyHmAaKeb5 z=b&}yt$K%(*!`uIZz1MokJfC?T3+vb0zL80?0?U*Dt+)jvZtcyp}gX|jaJrz+MG=- zE%GvNA|gN=E#JC_pP84}l>WXVOaFw6_5JRQ`?43j?SH&mCNTpMsu@rX0;*3Y{kL}3 zop;(-BxfF}aP&4Blhv*as&5`-=~l4;S=?uGd|ieWsJ1#P*?495`_yYsMH4>IoqwMF zLBC<54-zIw5Tx%`WxQajz~Zt*l98p^IxieW@p~knX^(V06p%u+g*`2HKxgn({=p@F ze~N{BDH}#!FY`c-NA7iR_O$c(1R&oI?+ASfLF@S6&8%f{sGn6LI(L`r_}NBSJL~ww z&!1LNE#_2|`C1n3-S#6%<|y&bKyeo$DE6fTo=`p zl){!to(YAb$(anTt*y<>CRbLjUYjW#b!hFVfX1s6X_^1UoJqdOB<=RRDuEBh1y+7t zF;(*%f!?0#|H!K_`tHa20A%JFS66%SfKuc?Yv0?@WHghp$S!KC!mtVD<$@IpXcEcQ zhhp|(Oo3qnl-H)%@L~g3*H?EJwat}+LY7&I+ z;j<6NPF~E-M#?&TB~BuL)(f=n`#}26oCqm=Xn+@7Jtb~rf`iA-T>tc=keFHeV^Y&! zfn*DH&#=Y1rR#nLaFn>0<>>E4lyJ{R5O(9mrzrJzk^Wxx)vbEc|xjT&`K%gsSQvA&r_m(>IT|i z_k;+azB}uAGahIHZ}8zWKB_R85beD>Y@!#1nb_S7@achs<_`vY!Zi< zmd`X$-0BD(inEfl94inr-pTC_A@SC}DpfjWDDNIA--DqgxeF0=S3vXcmWJO7<$c}} z>ogiYc@Jwr;ME5P{1|NzdWV30W1ZGi5?=9aIiT13Cc>pHDaZbr0*WB#Wy_E|{aODKR5{zBZC zY7rhb+W%Q;*C_!0z5&zGaput@efgOWkO3JNebDfyu&>S5M+&lP(>D-2ZwhN+h3{Oj zz8iSs|DXLS@F8ZrVlI7ii|MG~(+^0o8Qrtr7rT%iS@P0}gXUr4l2;I3uLAD6#{Ju8 zFapQ*#vtJz{WT4Fj%>JTrJ~E0BRHyHZH~pb-eT$_WAaQnq}YfDl@QU#d#vHuy9Zz``j{>rbLd^Y?(0Z1cCOdOZgrA2M!}ot+6RH zv$5Z^4HF=j-!I6xa42Rs2QpJLbi5QV<|@8SGcf_MTt4|-v`7pP{=~(Fy3W4>g*qUY zUuhX%eQpQ}@OU-3b6V2wmj7V^kHbCm@*lW|%xlO1em;z4m)+vlAs_E$F!x^qcDKoNOyVRn|BiVA4b0r-_qi}d`-RnWt`J?2Zn z2q1ml!iBzzNrWXTpdv<1t)LBryr8&$k&TUQu);$kJ&_IBaPb&0L9t%5k>NCyBN5hz z1P)yUC5urW7j!Z(s^eikr#9Eu6R5rEg~M;@Qqf5zdjj#tFRQPvZl6AV^~uB(8>I`5 z+bbq>JV4>KpP;|ORt7e5Ji+gBsq?ae$sr-3O4*sl+obN76Ok{|uWmnDAd-5b@xtvp zcg(2-RZtY}^di8|^s&8W{)_u21m~9YIU#{yBiBuc#Fg9R$(=j&YDihWcU4v( z&jVnb4xWm5&|#&`N_q=gbY}@7RB7+G6w=ai4-E_i+L!U(yCe>tvuY4$%DOxN^u3GB zxyvAzYgYrHKT;nVkk4>%{pj!K<$lSK5Hpi@_V}w0t3f-!QG25)#Q7%B%>fbj$+V!= zgczWU`e{mcRT5tGkp9uEBbN@BC{ninHp*t#p#42103<+x#6>n7A3$-t`7q@d6eO~c z2H{~wOnw)Su_TF!hdcHH^r~A`QwKvgt%#uPr8;lZc_yFrvo8D+C!}@}UQ%O{cd4Iw z3!2ui6d9L?ucX_M8^4SesVYloPP^o@TH>}j0bx?R*MSO_!t;tEWI1_F1C?Bc6!Xqp z{;T`pQ)%+w^wdNXix*N^1q{KN&+)FpUTK0Lr7aN@r z#J3=Y zMS(;^UVb-?ppbfA=A!pv65q>(SGTU}ihX|Gd{ME+>Qz8O#SOI9IebiFlz;|LvrG}F zA8HN{{3ptEV8Uy)CYT5y5x!sW5St5@6lFgP7t*HtXim&LK>wmqd!Hoqm zYek@cX-cJ@mNv6iGkP?zx11p7xPlC)J$;^Bbmm0PqHL;fIkk8=liPpd>;J#F^TFGj z{W0%0(lx>Y5(Z!Y2#=5%yscz@tPj9OfVy)2hS0hnbV`!~E%u<33`ESUY<_H>;g|lB zZ4p}{2zO$o-dYFSGy#ax(B`Kun)GAnjT;`6dg7<f7Z;}Y4tgua(~pce(4UTJIANHTJ3TXP1RG&xn~)y2k}k%D+B_LN=@34B$7yo~ z&epd7i$JV+RF_s zd7GY2SeYs@hc1w*Z2zcnqT6pdIrg3~r@!*@U3$Dt;87sSDdejDlN^V2eZHV4cz7a@ z4I^m_+1wsmD3DmJ3#|qlH0ANF_}EWL{B7tG!BgZ>gxgxV*;lZr7;&xxhH80 z105Yu#&Y)kifNC#0|bddn;xbl)I#sbm11LCSp_CxHobo!kY08vvJrNgEXfQ{O@z^& zO2R{*YTEJ8Cmu}zePXbppa5hMz3RgJraQkpf;DpQxT8L4S-Q0 zfI9{ZKWF(8{PC=(K7!#rK%_dENn~@U#w~-BOCBub{SI8wXD=%)?fZI8>QKdSZ#TT| z%ZUsRJM%aT(<;fx<{OW{7B)31H_LV?Q|2}fEWWqr(LdLgiW>UQHK6Q^KCBf*5lR8{AFKu-c zoU4QV()c|C8b4N05k>iL29n{F!)v>Ixf{mWNX!&D0j_QF&3&U&J?us3Q^8D+XhhOyU&VVQw2H@~Q{qO$;*afExbaTgWP z(i(;s@z_u0wPQQ-yhFAS8~&TMuy_v%*{R^aqEt4oo%6bXASIuaq2WSS-0%w-M!Deh2zV!bd&iLecSm6EJ>j?A7aCiei zVz~UFnr+)eI6n}Ov+!IF*4FS81Lb{0e0+9hCeTc*Z*2id*cn#Ve^e~-%y}cABxg=i5R^v$dADC0i?e`lT+WEKpY38Y~0tbX`K+z|N7ay<@AT#D>jS_ z?T>+w9!xtaX-1<#MO1EVbaWKZ0Pj&vZ3`9v6|Obho8SayWocgjd2-L&^iMVUmcO)&K}f9>*9;Q)L%(za4Aqm z?dtA64EL(UKhTrEv7zvO@eC;6n^&!F-Xz(Pa)lSF_SN^1h8R6Oa~6SRi~ zqrswhaNzmsuC;YW5beQptN3K^nX3CMhY{b)OS0!c4IP*+!(_~Z$hrU<64(U+pS_gf zy-aOiyv6e#&}%cPN;}QL5qhkp7pMV(ZUCP}B;7{tsaYeps#7at*Sw0%>VID;NMXCE zksE$hu3jB{^`SliXk2t746M#TS#;L2Ov;#Xh56;QygVdzQI<4J;i7aX`tLDW z|MLW+VY{oeOccDRuH)Ycj=xhPG&X)h@QN%(p6Mv^5IEkE$SZ$fTjoFvL<}dU-t(gkC^s8fVMy(#npKH#tP?r9IT!adU72P2WOvN zj?C2gy*V{%j0CKZy)H(i!e~Vh=l~h5GndMdzbd?f%ze1(3gp;9wmCkkz)bUkeflS{ z<#XcO=Z#7Uz$Km-Khe^tQ>7kkh$ zJS)dufjZa<(Y;-K9Sl`n;z`Pol$@;h;K6i>XM@$F9-9AF!j$O#UJ1MXJWHns%_i;1 z_S(Bz{_5q4FES24ZXjtXQBT3doap4TmyZL`I#e@^q2S3!5q5AFMu50bYqk^o!tR+3 zp!5Lu_sPz>B_$=rEAtyP>&$;@{HZq)m6S#_$UZ%X2UAf2AVIt7eSH{nU_o9t_2Cj2 zj6vRqXI`9SV~dTCxBc|$^UCgfsUI*9q8wOk6Guw;;5j$)%%KDOt#=6xfcgvcR8e(x z#9s*AlvGtC-&2jYascF`;+v?;$XpgKp;$%9_O2I1WB?{&eFaqF0D5Z7q~k8xK=7+q zVL4!;g1{fIQ|eBT<0=Xek~H&`5BKaxtb(Fkc3{tDldtm&Gxz_!K&85YcoPT&nm9OpMg3gjAb_G5Ww4DE1u6aMq5}c z0a$vZ!hoy)kg@z4=wm_i7SJjIjS_$(fg7sqGI~k(MSJj|b3tEkU1$Xx@ zE%fa@E`@wB3hhVeA!S-dC1U_^IDgO?s{p6~(=FhT1l`bS-x>fazMC<>EiZ z%F5@<$gkh~#hl`M?}=&e6VWyx_`j^==?V3FDHlYad{g9IOV#s%8|`;Vr@IQwZhfS; z(Y_F3;G<$(!Q^0Rd8yD5z#D#qfBso@S?O#b)ZNWZvWM)>>qE8qwV!b;6^v=IY_%v@@*~?3MFR`GYYs#L1k+Igd{aa8gYKM}`?q3iMYOtz1 zy{G?!x&nH&k@#6+qS`{h7Cfd_b-DdC#+Nmr7UV%6+}+*pg=n9Dr^vG+{w&0kwxS2f z)L-gEaw~=;AVII^28On#_yQMZgLoj~e7)0{!I!qQN&wbd_e)IDv$?ZLfjTDdoz^5A z9$~(!&vE%u$zv<4lU{i%3(UTTJyg0y`@Y$CIrv}V=?yjW4{y^zflK$FFIem(1iMkN z?0WD`dWbD;jn4-x^3zchEhD%rEXU$J1OO6ORN?^2RKWUW2ao{CkA_s@8Pav7D{R7KrRESsVnC*`r$EOvsQN-Vy%PsP6DsaVeE+qxDYq?x8-9&$Q1G zoF-d(QJqj6q)Yw7SlmIS?q@sIZuH^Eik0&h{ZWh>nwq!#?FT8Xea#!Wj9phxYIM%; zB(Q8|dR!v-V(7l9FZw*Qpz`yT4|I}%5MUAr+ECGq!#pl`Tntcs)NoeNvKyuGgC>Kr zHqd|K=LS^Yf@19P=)M>a2fw)%Z%{Q$S^@KJ;6|wcOighf)K^yd;^|Hg)H=6_ucR|X zai0N`8K>*(>(d0rmSg%l%R%Lqkceo?=KQgU#ns%d(5}3mR~*34m+62c{*RsSbd=sc zn24t3VI16S?f>Smjr+GUlm9)FFd)m+d3>ep^oiMK7Dx>DP*;~ZN1o=f;|nYosbu^= zhXZv>Qv3Xl4uB@G`KCeC&{{P=Wm%GlMZy*Adb2?MJ4<|s)~fWd+PqNl2dED`0A6fg zPD9muT~w#ENkPK~6s5@(=+1NVM&Q132g)!D04&0{n=HoJ6A;j-1jIjZuWY?)+CPA3 z@rO26ccD7FD|CwWmyNGKG?)A@7vIz3cSIk+JC;=l#P!yC-JB#jXKw$hAc{eKhCde{ zHZy%M8FNLq>C@+mKdXgW%x?#IRCM1vl?ekajdd7O*9GCIM+{dBjk5CB0h z80&(d*H8ev(x3R>IIeD$+`ioeV7=G#0M%Q$asCEy6O4M&vYtwcxPc=q1{~xS~v6t7>!Y7jb`#m#mAt4S7sdZrryvHAiBQt4?~oai7j&i z#Rh7M{zy(u;TA>)BJgQbJCaX(uTAha$5HolfTUfKvsxc)tPNeIg}awmA{b`%EbH)w z1OTkNLGOAz0s5eCnb*%2uj`yC2F0Iez4xjvgDQ)yg|GMEzeg0ebei^C zB_-Yo%MhMVTX6D%z-5U6rceWrkY;Hgc*~&!Wat2)heoQZF>XId;Q3+=>@1HE)2C2r z&7#tM{TN83Ls4w|`UNQLiA%MN)wfJ=Z+0 zIp^xp)s`xlXLJPxl-@;VKYsM~k*!i6BKS1U!+{E(vVQ;VzkShm2^t{mvL*vzfYg0N zC~LHZSpTWMf~XAXO<3cjIkIR|u>=_zcuR~tE*_nM{|xGl58BXk?~VB%<{I$H%B=>z z6~x39MorJD+?6uLqfRX-o8}_xbn1o5b}W;H6tU^hx&OY0nTef zK(hLOhlb$@5gLHW@h$|juhy3Leq(dJASE~lJXV0-0#JLvbhLs90s~(01HL?e0o1J? z3{g{2C58}K%f$wV12c4FRMegPp#hJumbrtb@hr45OLEQvB+ zes}sDSTW`0I2!^wl31F}jjJAWJ-SGQVd-{PCHC#5yIAeleua|ORG}7hgKKK!G6}j9 zGc!@yR4xp0#`4k0Y-?|{hjsv96fEt%#V*G=%0Sr`Z+PgPQ~J6n@~(I5`j^x36tj*o zn{nRe)!s^GF2<+Nh27kbk7v(WE;9cECJ$YNpey3QTcRDEy2keg1UQlGJ%^*g$VCJH zf_V{|yZGK{05GqDLztHv*JTCbP!NEBVwmH4d6SrpgTqK$TXXAc&aBsExUL z)pahu(+Rx5@Z96Ec>vg!%Qk&3TX9i>sb+H5004};j=iyMqofCPp;k(sD-A?7)l{fY zwsVx!OOW~072-=qlx33}`xp9tTdykef!24$0J$3Qs{t!F{LfIbgD0Mje!x_n?|GqO zkrHTPB)Wp<=jR=zE3nf5=s^N#1OW){4dn>v9 zL>vjaokM5ig$VC|qpS>j;P}#|X_H6ypFPVT!1XM?$mt4Y zC{8j1vFM{m(hXd_irydyNn{YGP-<*!)O!r(rgax))>0V$3PauyfL3I=iNmkKFK@sb z^3ZYCZ^mgz%WgWYd4e6xDpNe!JE5zq%fx;y3Rn=Q47XzNP8LR5%OClXO@vZ95+rer z_4Q+q{i>61$;lC3IIp+9spX?!0qX)oX0ZvjfPl6gsZ-3orN}Cs8G1TgVTJgnQ z8Bi<+1~d5k$?55Ay&~1N>vDB-8#~N=B?aHNOB8QV}LUoM#I4a5LW8&Z9+l%*gtQWM}w}JCrn}!G0$4)bw1(N;8~Ys zqW$p9=etc{;)|_&LiK#ZiD)B(+YfLRs{wA)0n{#%4^=&B@Q7IHXYEDfcTeG`sx+= z%lgey(<^%i7V*Cl42?(<{B$<~^5QWaFo#n~EvY8BM@Zbq4jTYnI^NHf;1o$)lx_IjAYsWaslQ;80G19csX5c z@i-CE@yNm?cm?^lHsmZ0YFp4(uvjtiqXW=D=-VgbOz>$6i5d-0^CPujKs+)19xvpV z3#3WkYd&%|n~wlqFQz^uh~a+D$8luj2c3bmmPy3`MFb+^mg|;UFQp~+Teo9;N=!??=f*2sO!hD1VM%j6Ccg8zmIZkO@_~}1KbfjNzFwc{=z7IK-#x-*ECz{vs zpYUR+FzN5HZC^P0pd#k*m2)bD>jujryDfzfh7%E`-VP+Ebp`;Mbho zJ@Q%B5rdizC5R=hf9AC(KJuZ_dId@Vs?Ag}(r3zVF^G{IeX;8fFLv%U7Tm)OuGFKS zgY(-#jy@LWG{Sw3KIR4dv0xh!pVFIjk3L!_WH_rZE@vZ0RarHi3wGo)Y2xakVB5ZD zXSk#4U>C}3&vxW#MRve4S%vMS4U=BS0!#i8*LiB~@Wxk~2J+fpJo-CYV4ksv69{au z2zrF&%)Oh`M{9e(YzO7pwJ$Va+qA&dYx_gZJnhj?B8%ZGo=a~kjxF!o9R9!)ld}Rr z(DMAlUx{~cVW1P+til7>wK3XlWoyOKW_mgJ9{@V0O+;2>te1~P;9Duivk;sqddszf z`G_CB;6X~wE&3<~$VpzawenCH`k-INeMJ>=ks!<~E#j6jt*D*jlzy4BF6xg`85I8~dtKX)G$)_wcXoQk zZk&XtjNbnEPFdg01I&~v5$9Kmq#f`4eJ__qqMN8|=eXr`<6qc6=E1a& z;n%9ra)xuFmkp_uzh01zK9Jv5?kga@#LS1g3bz=qK?9A{5opoKLHvq>TTMk9ly~L4 z);cDr;=f1+yO5JAZ=n!4EGi|Yd=%-*#L^?BxM*giZ2io(a$6rSINYxij?2zll4QWX z2XW_s^lZ;|jQF7k=X+nU9rbAWoj+Xg@!F!B89uXqgiWX@8h@ot!+FvsRB!G`skD@c z=d$OGYGeGM9L3+?Q1Isnq**!D|Mn-XW65tP?4P8h*iHL5yCvC8u2lM^&1u%du3{e_^l-=36(gV8|Tm=&s zcZK`?xnTI6U&XIttf8mUlDeu_>SmYA_AkRzo*Fld>pRCNP#=6(8q4{r@F0BOb;qDd z`APTLBk#Y-EjWHn)gEI_(Qc)R@uez2gduqDaG!JES~m)1gQIUuV3d%yao#AnpiTbA zCsk!RlaLBY3ryAi1M@{aE`RWLAhi!;h; zjopm&HblI7_dVS>+wZm9&rF9eV7b(-^haG+uX@f%r%G;A7>s_Ob1V0fHZeYQBsi+#UUG_^%^A;_LA13>d)R%;-T9f*$h{pDMqlw)Yv%^BtT} zFgw<98)c@(|AliG#&+TA?ouBgB&r>LkR=gUR>{m6^2e?ax`=;wy!IzR(3Vr;LnGQi z>IK1m!{yJv?@`zZkGqcJ1`lLFD7VLI2?aH#ua#SSHx+l@?bMu9Bbm$oQ5#L@3_koQ zdUq;4sjIGqhkp3|dYzi;xwDHbmS-;QF>D7}Z!fIx?~dGwE~cu=o}f1&gN1xAd*c+f zMRfL0s@&Kr{-FqhMHs2*v(DkvN-MDMzo`u~vEJCU?E?GG#7w+GNrKZ!YK%ZfazqrY zO?X#sRtXnpd1MF?dbm`K=wS9*C5kjVtj4ix(OVa!#+=2O=RO@wgdxIfake{48w1X% zm26N4@4oPzTVkCHReiU997f<;FVUbqm*pA-hzyCGTM|>9zxgm$((rpm@#MIxW_}E~ z$y*+!6{|?~_!z0{5GkiU`4xowm1>tuJ59uF(;t77b|WmlAFlN}bWnAnz-AYp-A_!D}|z$qJ)x(}Be#whi%# z^5VKHF}E9StmR(W<#E9;aC0XEEQWcREK-o*@58cdxv{90KKHYb3=C;&g6r(gu9L?^ zk(0F4RZp0ky-WVA#|q+`^(U7(aJ22wD>cW#P}zJ8t+n_U@ZDpg{#Z+rpswBrKdp*+E)C3q(m87Pjdu zKS#FmmTO_dO{!oa4ytir2t#b8!b_JPn(q2eaQjxCK+ZjCvMqUr|iVY9@1 zgPomVZaopOUfgvChb5CX)YFp+{s;LnlOQm${+_mW58iC}Zdl^e&OY2E z!My;i9rCrwv>i!kEvEI^U;m-1e*GKuISkY4RF;=+*G_4*!!K{)IkN~;~iol+~x_ll62 z+c_ibV$u#ehbfB=-OPJ{K>Ibq)B+HGfte`%2kGyc6fjx9(PtVt_2E-emp@x9c1Q|6 z6jvhJSHLO+d9YqF_ie0y^U@>KjJ#q6i-?0-p8Sl(lI1i@rS`@uvURZYvXEL#Ctlgh zH@|}@{vAeO;)%r0El|ZI3e&P`=I?C7O+YL$gO%S(RZt8)xCXcQ0&)88i2D|2+8)hu zu2oVG7E?e{L529Vt3T?|gUV$&IzeGSOfTL?RRGUl@@p*)^B$J75DsPB*h_2sBr$R! z2v{bzmepYK6=cEK)FuQs#Ar~67ZyXMUGNUq{g$B;vku(3V6q)PSEz-J z^EPXAq)&nvsv35Z|CHKDIn;i(XB48o-$6n5z@}g8{;KR+4W!(6o{0%YcH z3QsRzYa2}I6YpH;&UI@-eZw6rxb&!4S|6FsRd*f+L4Yr|oNuUcvw?a3{VZ>Z>$OSy z)t;XV)DRr^GFmMP>bew!`Z$|LbAgv_Uaxq%SN>Vzl<$>hZe~?;J`Uw$?Y5IF2D%#y<@$nsaU}lOq7| z<1C%#e9wd-{Jf~cGg<%T5b*)pSa6O?2|KQ-f$*uySKIhA6<;&v>q6ge;--kMeZkSE z@)&uda<|ido319J<%;?&;n*Cld9#QY&g}CUtQrz0qHdZO^7lV#F?U z=t;5IhTE0|QBh8iikj4ujyp~Rx0)#5HB7}YUD+pqPaEoO&(`G+wo}r~E9V{5ia=m7 z`v(s6QbixKjy}|d!(?wu8f)wBg?x8eY->rmnu8@tqB-8OX+KZhsrC9cA{}xN^#D@h zAETJUn0}Q8mg$_7W0iKLG!?qn9Dvlnz7|?g?B({o!oYPBJ~=ri(5+F{tFfh}H$S|T zQ^v3+%~gr)aoQQ4Dzcz%d*?ZNuwK4cPk$9u(u5}n_sua`DxkUL(LG|0Ku^srXf zqt%Ru8~nG_mZ^^)ALhH4<$ca^s+~&8iElml+{y)oC(R<;lFBf~BTbI-MVE;>v*#oX zoZVFityGqwrJu6){AVbnjp8tL169KwlLvj@YLT}Da1ALAtgo%_yT|X$1*0=dLCwf9 zXJ^0;K7OT1c|uhF{5->v;ZY?1Ev#hWxH1)J((mY%<~q4HyaX;1wP)eaWAfXx-@c^I zvnFLf8{qfTJ^pn`g~o`zT&v$pu$XUjZiuxqECN%s z8gIF|-$A_$i#dMeV^s%ps&F_ZgpsiJvfq6=(gk6x{XuP{ypZmaXo)kDeA0E|3xmM5 zvGW=_+x1sl>aznCM?g%;v$kj89OJjOKQO=Q)gG?xeEoyzkT$<3^br#@Z>)v!>Zd_z zUeB4L3L@mCZJ(m>R@^+~qgt1C`LI3bft7YgYxNy&^$Dw#BM)g5qSBLn02XOXi-4r# zbaS%DGi6FOL+NYnw2mZ+V7m2MSfC02a$GT!Lf0nA1^u^R0q2CLe3%~`>hzMbUy+*F z7%3KY#05P5d5s`vsayj3WO=%8thw@pH%bEzH|>x{4sGR<@e_YxCxaR|f>rLdHfy%i zcd!`|{$3+bgX-=2&+p-Bf_7|8nl~{*_KUWSFNR|rvu+xFa;#*U&d6h|J+&F^^uCnI z!ZN^8oDzXGtF~4zCj(o#KW=FK#>mo|I~RT*5x*`3w!x)IkRZ+BHuioj<{~UQw-ntQ z-1*1`zPPtNwv=9-NZ~fhZpFSeip^FyDo-E#6n(J?1ys8BD$-ACA4>QzF6CQ1L%K~p ztRB&NAXp;V9G`6$v^2|K>9`!SxhpY0z4W7eWC#OhN|YWSL=59jBCxwx9aD6B zLVUsy&C>hde$2|DKj8=EGW{dV+m1vlw+;_wxNa4sM7O2%W-yM!IF1`I(wFjIa5Na~ zay%Y7X&rSxJ5kkvm%T@CxMKK?Q@Kcasfg+HQ&+pKDU|JXdA-dA?$lJNErs0*)>zc?8M7AJtF75@4y+Y-b0Rk7?__hZe)WVW zu!cI5mbBPyXbk=odzvvbnZ?ROW*5s+_S9|V#_C7h_IHOcL{269y$uUYxIj-o32ErK zmzcYg>3cOJN{rIW?nRI+7d+~`(b;ks*q5d(dr3Gc!&A&QmGe5f=Rp`F47d22uKVE1 zpZ$T><}E~ccIlIVlGewHm9A|nq*~kDOX$}Joz7lhY01kWv*>FMl^SUO+QrFOD|xb( zKDiRJnYV)2v1S0V>j~%ti)5VMM_Zy%q#z~=DkY}5`;DLpHPz#ECxhhARao~0l$*?i z|Gbd|cFmUvX&4?`(?o_G!Wpt9eEoY;hn%fNVWFfX}{)wd|&gq(FSy*+%P6K%}Khp8Tz|GT{Jq zurvGlGyMG-GPl&Fo$m$p>&!!$v`FS+ry_8JP_of>UpZd(;TTg}kkTWk+&3D&$(99c z`b4yrfBIB;Vgy_XALGigQpgrdPEz?x}MF~|^ z2Z%uwV#Hg#c5g}A_7vR8@cn!L%V-rb){+b!sx$6|)vu|yJQjX@(&l$My{Uq&QF#kQ zq`Lgbdt@f}v*Y&WZVHf@9L(9YmUM4=h^~-oQNjIzG%1J27POdPIK~{R;Wu(S%Ep>3 zX48l-X}J>|-e_h06JH1CPa<5!9Kp?Ze=FjZn*!1mtC7#zF_I_$w*KmlK|9r-X4pCm zRHL=}OzXbb^$g_5Jnz-0=!Wn2KDWvP28Y>(59lC$j%D5w-L1$dK9fzanl0)YRDFyG za~Q%1H@C%-jpHgaoPt9BR&9AA2ug?#_TGcg9T^e_QdkGOPmtFB5oI6LkG(0DSHbyi z=l*tU@(4XFKrnEnbf)eOnJBADF+ZyH@tXFzMHq}DXnVx3aC4qotPJKh`%*!{l z@0XlO_22R>vFCz2_-KQ<=VeBv?X;6ws$sFqXeUkY25NZ_*Vz{%syBeKr7RMg*{plY zU=E7E8_B67mgQ-$Ep(@#6TfOXaGZ{?GAR_&$BZuk~o8XM% zKT%s1Ked0-UvNL^4)BbE3aCP9V-9<BSw3mSTL?KEc6KDY{?W3K7B#Z6g8$|Cdf=K;TGW zS5W5vFNim<3x|fUZ3(x zn!%PWGXKDxkMml3oBbFclt|15UHU@m{>r!8k4!qWtwyk(rsgsZ9@`PqKdj^Ycii~= z{RSGBhCE~GE%U#)Aa#AoOv>y_$~TN~+ssWUJpghE4+bzB0i?1En0_oR&Mx+!lxPGM z2kKF&qQQePQk|Jyc8q2zo9f557*~;}{uV0%$`z$Zs0@CKLt~o+&SY?(H63YGJ+mX{ z*jF^Kw!bjm7yOuUJg0EEw8OyDdZi|=RFntB!R#)HWM0fY^AelZC>uR9B1N8OSmIkyCc|fq zk6Iip*IfRV+%-g_m4_S+q_5pAJy@$oORgZ7ohwFq*EToCHW7Iuk+F%9;g|u56VgVH zA1_sjM zl-$IZKo@!ynoz&$%#|S5XMLXYILX!Tu4j3Cbz}6m2;E~YLUTS@Q8C&;fd@X*p@|o?O-H(&jVDTSQt<6fNmk)Ii8UWaqeu#sFkc$miTO`RLjNJ4S zICl>Mv+0RMJ5}YjVhRAy(Jvm_1sUOX>m-jvh{q`?{68@r6l=6|kJ4;< zFv@!wJ@$eEn8yBF*A|C2h=>j8i4p6`RP6-zuGaE|cw)vFx8_Opa%d_E5xJyJvlLx- zV@R*JZO{N0#<&}dc(&pb8}GW^C3|u(}!-& zEv!vT5Quioj}cqda?hO^#-R|hHvQuecUp!03#hen2JKT@DJjr1<;R6E?sjGN*3WK4 zxS2Zcpd1>D%0=e_wD0uDStx|=?PMkDA&0MXLpXOESAS-iZiFE~BxzALpYp>CZma*5 zaVIw;?hNl8@>TO+kyh);9n^M~yY`wpUEcIQR$$!X1Kk_LzLDKDa4k}BgNqI(i)6SZ z#_A=I^Tmz~3%We5FkZ?PT!1_YX!G&V%qjXBci&e`j;7?b9duYQ?EK{L;~(N-9t01EToZ_NT>4z1uKVTpDBrnf-hw z>;cdBK@SJtP=ZL`!P2tb!6YVW1Z!f>~nI-5(7$VqX5%J@te+1bLHkt+c<= zRO8HW@4kxHsfbA&;=B6exKQkX$jdMW65#l0J0msXH5CcUxsvf{AezD2AGWbr@3J6w>QDo0lhj9pm!>fpcD;j(?7d-P@nJ;C6aTN<_*-dMWUm zL?sOZ>u;JWf$x~$n=nhx5$jP>S#tUNVDnp$pn4^u_dR<{i+Z?RERi?-CPoPg8H9G1 z?XrBRfBk;-`F9PFw3ed=>-S5FpIUwXS#Y{+K2t0bJ~3EtU=S!)k(&>7n6HVuQ(DsO zD^6RIdc}pwmV7Exc z?ZDiZKr6i9=_Kg%)dkB&Q^KberscHlRM+;Y72qTwb^6lDu-J}gY^gqn=D8I+yMHl9 zJMK>P8*=R8#@BOUtX87>wym>E87w6OYPGJ)r^g|ps?^oSfVMs z{48oj9AlNYZ=V3lRwmDqr9v-X3*170eD86eRH~O}$Ku)>u=z`XO;!D=wY`ld$2&0B ziWkf?!Y$rG%UO6+aEvuBXh)!btQ+{D_Edj%dXOD-HB}q&N;}_d6CSuHi3wK6sj;4g z1dmWgG>4++49xBJhhTTqJhRs|TC2|f)#=MrSb7#xt+wf*zhB9G+u;)SL1!NYK%^j; z;>&a;;9?mw3n_&YXpz%CXRl%MKDqbSHDkyy1@a%y6pR@ z^3j)~l}jy_cDlZ1C2w0hZbrlwfAzZKZZ~yu&%RO(A~N=qcTRdA6WFb5D8?c7qKISI zc*2J!VCD1@(^CO7$pa z3%WVMer8kfIjv1mbm!ck2K7Ktx3p=QL(Hp)RHBvDRJMAkzz$chC%UEPwdxbP3-vsa z?saLZA7$_Xd*>2H4yk(~r9FC#dyl8U z;%=DpLKn3k)3|fN4NX~C=23p(Yrz77)Y@UwpGkE{uldQ+kW<1~kW=W0Rf)sIx^d$R z!^{Tk;u#QQ80KJyvFDwsZ2E{z9Jt?(MnSnU7Zg)Jr!z=feK}1V zy2cS-`K7XF@wA?yheev}2apTZ2dZ>iMy_4-NN4w=^AM3wXV(Q@W8bS&_TQE*JVy(n z?1qc8ZgcnoR7_*GsEmatZZVaVdapT{8;7~0XJIz2&dD9+ZqW0PM%F?Vmovjd#`tt& zbkb#OQ^0faE~lAA3zm%V_^+pWLI;$cp&;6b4Nb$zFnVPz-S*6xACASj8upJMndjra zKDFw8GHi3R=~3m2h>hzjpVFNxq)o3s7L(d{tmWHm$mz`9o$%QLxk7>oNa;-TL3yRk zw?Mn1TOUM&$s^abw~MWcF1ofg`ojd4LQCE5VC*I(n99_$Cx@z}EXICj&S{ ztDPnb?IO`ZL(wv=keq9!``^tww9(@RJs$8X!J+^NrH2KlWVr-zqO)4Z+bycYM5TB$ z0L4akg#~V**TxYHbPe-%Mv*=cI{$`c!{U3Z5tB8zH$Gm?5MJZkpT@Z^rJHFeu^K#j@E9dt~hdpJoctJ{Bz)%Mt6(d$E?e zgWXm&L?u`sDV1h#HJxzVAzeJpel3*|6i}XkIacL{0>V8LAiRH#o4R~%F@NM^9_v8N zLvcMw3CIik^u&aJNPh49zQt52EkFJ6q01Mkt%lWQsV8^l>7|SBvftAy$nFy_z0ck& z;xggL{!?$;gQQ8n9malp^aWcJs?uYwuRkNS@@_^oEJR<50)cHG7-5rKX}Rfozy-m5 zHecUNn_u{UbbV!5RBIQuihxR~NFyNvisTHP($XN^Aj8ltT_Pn7l0%J1w{(L@!@$s; z(%muSw>j_oz1R89@$;9LKiqqsy`Hu1b+3D^Wug~|V83D4Q}txrA=kQTxH5}`!Qr`y z_GPyr2lun!BKhMSjyBzjYl4(hZw2_&VB_B0e*ZQ&)xZ*<33a7w)Q;VH25eFsEI|%~7s8<2L=p@!;s+g*!?#wSgreIkWW1BL2U-Ns|tb7ZuVsi7`y6FK2 zHmN9|$fEsP-uM>RzJ`l^(S1D&B=s1i)^jy(&eWLk6ziuZoxeyZ>z#SZ$gbBL$F6^~ zL#_dOd`cba6%hdp?T4YHcyw`LC?H*Cu+$m$@dvgz{iPafh2D56TwgTSi_c(C6`$*w zh*3aR5I3lS;kCt^XCti>)0E3Eg`r+5ULQQVdwKyn@HSlz#V0nW$bDp4@{3z+wvc2d z!c8^66}}J2*tB*=6^7K+GU!_fo`*S%hHgdp!gD`I4hL3M?bg{XW^8|x-ZP?55e@`0 zEnkk5{U(*F8XWmA^VRuhQ2pb6#h$gN*R!}lwbsZJpUFmF=wPP%a9=(7%0Zw9SP)1t z&)RhGyJ`UpFcf)fc6Rm1XF9olh)d*fF%IQr=GfCWV$jqevbvdJ&3*kQ95_jV>3dbm zl-x??yY5vY&_;DNHAB2ux()SGk-XS#g(ab<$ zXrt%eEym-uC8@qefgKNSYJVB zy4QcMZ9*`5#_1_AgCh54xZ>vKX8A|b3-IPiSe4EA|0M+@_)hCu2Fw7oqW^|@og!t| zBI-6_+XK$Ek;C$;G0U7Q3?C>D@nd1*+byoHbphP^A)#~%sF1z~Hi7?@c(BYd=yPi^r z0}9Gjo0BRdlExdeQzV=dytod%i}8qIE|;5QGqyi&C6oYB z*tNs*eD)B@3RAF2-iKB7NaPD%mjU(Xd{57Sywf#u#KZ#_>OVdN4w_XOa8vR7s+w(pMU`>6dn#=J6awOgl#diljI)5+T`B0B z>)zlX+)eUtt3tPAYjkcWoWW2mcZC|+>SKfu8!CXX#L3fiH)Ibsqv?IQTvfX~kHDXj z(7yMc3e4eM13jNLO0`q8W8W{!q!)5N66pf zQeRzM#!ne&q-aYADCfB=6o>X{FM09=P*S7^Aw{7K*PB(6ea?Aq5%1)-qmYz^pO7)C zgn8cs|0-;s{XmYBG&vUMXEf~;R(_@5uopb|H1!ix{(9+GW{d?Ipxf+{12*>pHN673 z5X*IZ?MXdZ(L2|I)6g zk>KfQDB>&|I|a1}h&xlX?Nu8jGiow7LTLYjWCmR{(*$9+mTdb4ISwO-ZSc6}+i7!h7`OAJ12=9TfLKEj?HDr3ug-xjw9=~CZ=(e3Q1@eXq!K`qa=hq)_B+hb# zwwr)HoR6Y#U}>`iFnwk+U;X+@p%Dw_O3S57VD{5h5PwDyC=}F<8<--9T6X49WJ%Vk z=Pi@~Bj!Y|5=-IVyuY6v@k|6dYAhiaG5bCf>p%J%alys*l)e&x+sT_V^%2ofqlx(3 zl*_Uk&l^CJ08oaeR?pU9@Du$;ERX+LQIhHoumzTD*%{pQl!!y1 zaGz$%P!P<^a}6BP#SN{;?k<ee~ZSkPa~ZZ!qOpu_ZO-9@Wxm35l`){mxZxIiO7v4X8O zeQW)EIwFrRvVbKa zX)mKlSnXkeN{}Pk4R=^i&s%5_t(ZdVreZD`Jx|$cXJ$fbYW(yLhH#`6`Z>aF!QMq# z?$GM4ulS>0y9;Ucf`HKKhD`B;-a;?m=;K2jTv~>2q6U0sL^*|5G=j{GN@lQvBpQb* zvNiA}h9}hGEqsA%P7eMA=LqQW@tJO@ZDu-`&^PdaeigiVGe<_L{mi%h)mivYh2x}L zDzN|VtTDFDIJVX1hA^~`o*vIt*E>j)>m|(fN}n7jaBXcX?d#+F6J6hs^)JhwhomgD z(B>6Nb#w!-h5aEu$^BSWBE`=A)ppQTR>o{&AZPox)Osdi;1$-u+Q1I3CKT>@sG&V$ zC*t&_&2lA|vec!dfyauWCKb=%RzX`QZCx8U&$Gs676VT5uq#u#An9A5btG&kmm~!@ z+Z7Q{exR(<8UzM?lT|ttEpmbv$5P#_AU8*mi0CuR0b#+42RoQWLrmdTarfo=NLmp@o znEQsb!~u?fg>Mb`+g0UhX&tejGR>Rnpa>+8%T_ws_zMhHNHLfB6X*tKv809@!DB z5E!#L=i%&;W&I;!noF6zKU#WHz1CK%+P)R1XK_NR_33#6R&mmc;MuFG{LEcp&c_@JoLTo~3X?gPY zr`XtD@3m1o?82|bC0Z=e|064uSVfX|o@UeQGn@!YCdM~GC1KT|WWNNlP z!bvm|&N>d>f#pt34b>A*V78I4->0pbLu6-?Lrv8m1+$4E?JeS;i5TQFts|Qk{PD+r z3q!>+K&IuA?u}`G_#qGKk*pew_H(j|q>srM9D63#1s#~!ff{b4l!Ay}n|NDzvi{s1 zIIO!Mt#qvUeu^C}eB8YKm$vV!OXN7G;i9Z8^TQe0`>>i9iM&mF8lJa{LH>DI4=Rzi z)zjc2PAagF>eF4AN0)dxr~mV9a7DV&(z~+OmsiSJhYP#8pAhBUweKf34dS|LbyWcC zA~5aaIKztA2YuiX5N&dyCxKvi{W4-W*6ufXEN1W)F8u0mCVhX(q|+08BeNoJJZI6^ z4*d5W&&`2rJLS(u`;jd5!5!FZ)Z4EO)uXRc(NlV#yeUM7DPcPr*y@!t#k!~o7FCj*4lcucOLB(j{0F`$*6WbWw) zIBbRvjliFG%$U@{+wr*vl(#9W+v^D)QS zcJERntTsqVWpEQ3c#c35`Y?%*9PO$>XNCqgyoE{&MKAV`xlo7qMJ)>BPoy_to^mMyD%R3*U91&zSLb+21goOZXJRK zkw{pBh*IyivdYGgK{glTy4bBP8~o$$KS)4g%PJIZC3UUTBZ7Z=eCR+STKP{CCK%5U zlse`2e!u@c7u$WYT*2Hx#c&}mc8mi!1zF_!ut%WDF?vM9f(3nh9!dNsgogufUhVaS z#Sq&BW0Oe0`v+p)>UgnVs5<2_{$US5fYtqeC8bURS)rln9jgGmz_^`Y`s3>RpuXeIu{eu9C;Oc zbboqXJj)nF4eq7}vt^D8{S}O{MSOs$+HX<`bU}I1unM*qlsIXV0@VmTQ6~?taOoM& zEnxnnYz96}uk5+^s%x)2JuZaqbc~CT)U9iJRT(7c_%i>@7HSx@pw1sldDuP)7>dEV z3J&OSWdFLc+if9i<_7#)O%Pl&f7{5o4-i#+LVyKP4hzED`?M ztAaxc+{?Pyv5Ju-SxWrQ9|_<8@kY0pWVBqqg*LKyUR3NRp%|TlkxNxSPFQ*nN#Sw3 zXUWHc1{k1dR1HIXTb49`QXk5Hh$oW@oOZ(H^8O}5o#GvEt@o;(B1yiCFEKNi;gB=) zRR$-sw%GTq4CMIu|M6PPgX2h~fVM_PJDQTl0O?7utp`BeXPz8H4OxkB@)A0#-^$6p z3yQu<(sp1%k{?z4dTh1sp-0jgS{8VZ8f=%R(k}6?Z;RR(qaZ+93oWeqW8X{TfamJs zjHG)^e|jR{jLG-Bo*Sv(7(j=N7h!BU)5NtP7}O$QD9>nm;+$X3Ro_RyRGZo0iszB_ zmOh+e-;!=E)WV;D3f%5P=ZC0H)vzMnd?`@1n{V`l3QTb~Ol5G zWi}1~5Zyfhas1tg%F;dQsv1yp8GgdhVWXCujO}jDQ&X;Wo#PgHqWLP1bPDluooZBmb zdxSL<7cCrbVwO%4-qu4x^F)c>0gBB;xL#vYYL}>N16Rpxf@>&wjoEMjKLYQi{F@<1w&6K$=|g&Mi!) zfGMg**Ev=|`pX>ch0=hbXtT<2slkt*)Upu&>R7QS$R~XZZGFFJW`EPAe8w{S4nMZa znlzEViRU@`Ef6}0DWW1wT2`=%?|Q_rQUb4$WcX7|2mD_w5T>$fC2<{(L2;Gj{p%9HUbjpUR2+DOG`#B?_zRrQQo0XX^I+haGg zhEM*5Q^boSTQn`e)0(71bg_{nv5HPRScw4S`nS{P>oW;3F%e`R1l&zSz*grLc5LnT zM1Y|3bjoF8I$9?}FMZ<2@2yYwRs@2Ax8|nr!OVWhe}5!E7Lw%@341evZp6^^a`aAT zMLK9y&UtRD`H+nYJh)L~d4dlb;_4a>11^c-v+V6Q_eXyM8-RxHAio~U<2Sx3W(msE~fRa(A01g}U<=b|$l18p7uh9Xf7Fos$6z+uD@_=&0|520faj$BewK>osEznrC~w7O zu04{@(BM8~1>$6vzdq)N;HFrzgYy(F=Z1zdJ-UTJPolatz6&vfb(r}Y6BAh%VLe7a za*l|6{&e@QJOi3PtiKeC+63JMmVZ%_CQm9W89j6|NIF9t4Ym&S%LqBR|Cmb9NI_*C zzPPNXn~FKUn4$u2Gt8>gm_&1cytvEOOx5fP7hOf5*w+Z7g8UnZ|2oMzPt?mu7}-;^ zL#&9guBXw)$Y;(KP5FiIoO-rOhj+;-GeJ(IQ_Ut?9hl_%8NOp38HWcuR9o1m>e*giKj zrOVW@prOy**H?!$QOn7kmCQPR2rDO0}`HD{TJsC}!> zG_rfo{{;RftOddBV9)7+&+M={i_R?05hW77CjjYCfha04CbHw>Y47{oy>KK8+xFlCn9{@L7-ywi7GV3}zh8)RAC|Sy?44j5 zH3RY-&xqj^t|gHE57Gnv z5Xa&S42VR2Y24wSvmZjkv)RaIUW8kxD9-L-C!(=z;1*mmbRvE}Qx=MgB-wstjpU9! zfAl9GQtt#e{Gp+>|Kv86r8|U!s(J^XI{mgsus}jQ=`#9Gv-QM~2+7857EPjPpuGlR zAdiUirp5y)J4gfca5IXPo4o6QHZ0hOvwqXMleK*mnCzrjPD=r&Y_zpU2%i4G0#fFsDwat(w@N838cp(pk+DFI4>F> z8h#EJw4Rjf4i&}*lF{E5w2`n4BRrmRqR?fd7Q4J8Af!oKn9ZsKhmT^jDG2MVB+P z*>19u5JFGmCk3lnc0T?#z+S^#_M!i&k-3*RbR3K2Ea`j6L;y}96}T!m5Bl`uLs--m zMREQY;1!CzaDa-$^N3drWLaN{8RJZt0W_Spg^*ydi~iHib!8I7*HSu=iN%CeA>=6I zbKAnKSex)C?!ng+HZFg{iH%Miim^_ohu<27!<4Y@AQTwYc2v<}xx6|1A~2X?>C0{) zRb*aO4~+?U3PMu?kYjPg&o62s&~WjL)%W4eDOv`J^iM$PB+X_NFYvjZ{0UbuoxrZw zyGDB0Oe(Ao;*X{Xy)I%`N zBVq7CZN|>6f&McOvVnp@vL7cFeAM39Yif0Z#Wx)Z zxXzZICHH>Ud-fc%;@9>8d8grgTr?6ETpX)0#B!FcC?w!a7eCH^=w;SBtcb;ChyiIo zzF+SOkadx%Q>V*pBG6xplU>i849RMA>3W!ws>-`!f2JlqA62&z+9MIuSEr8+{&&(e zK?=_$+tStyN1=)0oW*vZ_5GxW8`^3(e_iT)2o`A+>5#)17LETUoj?T+3?(w>CFaOJ zdJFAtiZqe0|EB)Z9dfCm#s&lwNBorDg>?WBgidtu(}gjV;M#{^H1Sg01ufza%3k&@ zl#0j%u|cJ&x6U?jTnF4V7@0O|ZJ)_zE$+oP;Hv#8vud${rYX zOz0dq>y+wf<^z$T16ikjvv;CVC99cJvo9oNSH5?{XrEi!<9sUMm6&k_&aPAeLXr+1 zc1RU|>nxOj?WbwP=UfXnIj5!u<9T-^k)=r&_?G4k6e~Oi8SN^`TsC$Z@IoGg(n;HM zCzo*ixepy^vT7k!vc;X^ZYE@Xe~3~G`N*Y;s`^4&1*abJQt&q1s~oM%IZW;pmN*a~ zbpEvL#(U9Get2Xp4y({TbchqdV=mn(`&O{wI{0;`?0Vq)qqJ33zBO8I>;%LD;(LFp zu%0%j!ZCdeB9!s8q#$DBeBVZ1(g#62rS`LUdf_%=vKdMU#Z6Up?OuE^CYfLw@>&ad zs6fvglB4m5dq5@#sDyui{sJd60|}|eJz!?Z0yXAa79qtrByT)i>ky}&@-sg|D&@1l z8ud{mJ0zRpktEUMneEzfRAAa&8OGxh)DOxYP@kHQoLb2-FtSFIr7L>6m7uP7jGPLCfsrsi>+)Jm zYbvm~cWu+S;5)(1T)?L=FUR{LrWx``N!M3kmnp0ysNF;J7r9S2)Bj69k^q5tipdXBjkkz}UFw6Je?l7RU1bC~ zL;F4`xpqR&WOt|c0`IE-gm`T@Dh+N58&Lb|Xx#PU57oE_t1+$8D4Lf9dS`c-*aeD> z^}!qtyn!0=C}GWQ4-+;HLy3#i`m}79Uk=3NdfawQ6c7oW*GW8V$Z}b_+w7w1(A9UE zv-`cM)oH9U;O?tKaN_rmj2j33zzK4;nQUNO2(k`S)frzKejj1KtWWjXr7KeoHhnx`5a&jz@e$7BIo?qNuqN5z9k*sq&9u>NU17rUp+R_8%2LaU-Tl%L-G zS7>h$j$FcBI=Xm2YdRn)i|x@dCo?u>v57ZJZB3ahU+~xA^+e2mznWhat3q(3W&eM^ z@L{l^s1Yb8{d3oH!j0rkk*5>sDPIlwA(krvyoBMHx~k_g5c%*h;Unpy5>AJD$y+Ec z3Cl|hdSAq|dRmY=svb$Z+o#7$@*sXxzGx1|{FiJDe}t&)FHNB`hfn}}3q!5nWiSy` zDo+`^tsD?1jlRY_xrbXe&2)&WPN%KLiQUKkV<>MkDb*tLE~b_Nyo+-=apJ@{(dFE0 zy*?8S@{OumTq;ZF)X zmvxiuxN!x7c+*j>0ZD#2bY?NO5L#YXCk#FCcDJ}ChTT68(4KGqW{8Agxf)xc8u=%7 zmGGg2bU9Bq1DXv&-WdBL_^Y2cL4MTh-sUSV0o`64)r>2Cb)py4-?g(Yx?gL_(Z-et z{^X13&APi;qPha170;Iks z2ZZp($#IFaDvr}&otL%7xE`EPqhFeC1z>Z5KVjyAA(L_Pa80dW-rds}(I~FhhTDGc zFqk{)pNy`9!6DPpUX?;xsjOukyBfRmlRh4dxC*0iAo-1J6KmhyS-M$i7wZOJuyxaC zD?k8(wliLpX?%-XR^nvN#n?;bL}K2y8t}GM;MCI>|B3i((iA56P=u{!!9y04Jnfai z{jmSfMJ%I$=kDyHk{U@Tm=0@T@>6T9u|5S!zy?8KzHK8sbl{3DMB>YyD?8Z@gk%9O zpDRV)P4dO=*jUFRuuJ>3neCD-VMz?k+U%jsAr8QvdbMrle=Fs>7vD1kVn(fKxCFQm zQ-c}Ok`I^}c1ONBS02wchydoLMQTr|@3r;fpC%8moU9rJe&^JRaQ z1W6)EdXTP%EH%BEC01=+(F<~et=|#1d>)V!?LV@e*pd=D4X4ou60wk2W0m%veu|Qs)kMdt({mqAc|%8g>5aF1szo+);5ZT-Vaw7fZ zV#UfcAgEjMAXC}Bf1Z71Hu9rfC@hQC!#J(#H^XD^5#@m|I0_bdNNcOl5Aqo~6cl2x zj|J8ts79Yjyk&89B0c<@{8uM<_=U=SYKpUrOw<|Qt8>ac4cvm=1iTJZfF)z_x|MCQ z)F~EKiq;WTC6dmb14g)14%Bn|v9NVKs$@5)Xx5Zr(T<^R95b#e(QeiV3j|(@mN(|_ z!;U`wRzL_@`6w?Hwt1`$V-@}VeEMildi?tfYB1xC2!%yU{(}Pl9@ry5I)1>BAUo!t z9+D}pa{lI@sbawPBAA@H>RY<%YlgE3Ugfh3z(Dpzom@ zD=N=RzWP1zDm!bkNzrvr;VG9IZZ#bc-#4w8Py(X?T%J1Yi@#ajI^5$Leji9J5_-HYFoXg2sph%eiqvZS10 z^-lw1L-bxDTX7_8qbmrc5=qkCe||!}IwT{Acz}c8S7VWo_N<+MC}rmi8rj2FOw%!i z>cos{n`^dO$_?(qobP&YC|76VFq+09(KWSP+vw3^7STk*=<7am@-`V)AU}pj6ft&O zn&>Vcr1ruN!|qskWHm-GLAb6>oTc7F*^Lw&G!*9Aj6K5L`vS(zC9c)zI%SKEpT+KT z_m&-+I0QMuc2}0d^REITNy;Kz%S2U9+NBh>8K2Taph`^1Lzf2~Y4i43BA2}3%+5W@ zeDy&ypf~^Yo(qmB#Q0^hw(zK(WXMeleVi?Lr3cO3k_{D-(vU)mD|cStiKnLG6vK7W;pe)qZ>j zH+0}`e{Hwac3l_q!|Od_RAkI-|i zZx4v0cGvU28aI}Jel91V?yiGdz&dOSldEw&O&Y1JhnTxVs$rEXLG&6E<}1M;wA*cL zs`6a`i^QJ4?hyZpSoG)dGAW*(`;ZO#+lhEfn<10hc;l6(5mfrPKvP! z+`;e%8{CaaTXCANy3+?ycm8RX02>0*CN*wBn;neNHaiz3F9(?LJH55q{WQOLQ}b@+qSDzBZCRu`rUzYSfY&dNeHwRX`#OZpWkeV@ zu%%AHciLRyoe}H0)xH`J7lR{#4l}Z#ni0*etHm$|*w3>IV#iRQwSYtn9Jw=-F8Pr}{h&GWH?h7zbAjM% z_mYpmzMPiGkMLUI;W+q%Qez z+^G1n$+ssZ=;0>!HO>SR7=FlaHO2T(-1B4omNB)8PjTz-gR~h^%_>I6;ox;zx}I;m z|178KZec$W&@>=9svw7_J?vZA<1()s0g#$PC-o56t5zc+8`_8F`6O4GHxZ5itDG7G zdyN5b9ypx{Sm?A{N)tY4(KD#yyYd9fXW;Yub3uU5@5xtg383MZjkvAx4Ev9Mv`d^D z?+9$fP2^dH{Ok2J1(hep3R51V(xTdfTL8cqrXW{=6K%AKd66qJCzVWCcdrKlQ zrQ&1yI%A9x;|b+xm3Fni>MkvPz}&I*^%1KOleU}A8BcAQ>o-^p(eiQ5mopq6S_(nY zXwq29NgH_=?v>Eht+7sQs}dv2@{Br~#-i>zZX_eq-2%L6Jae>-rIf%<)=wi^w2A z(suS!aWTi2haAUzM39f^dtSk}^a=-$kLzYr>NE=IfY8k4A>60wdGYt@vLeo#*N3LM zi@j_|Wx)~>?~~5662=d@(BsG2wOihg@u6Az{T06qemMNva9I9M7{}q9 zS~GZqIOM*#dsd}d!7#JI;uC@wEO`ipQABV?jU#?DD zw)r-s1>rqU6bzp} z74o#0w0ZZ;ddev8;712aK1bM`R|DXmG`qEi2+}dE_UvQ1Nbe_I{#68=!#kkO5KmSi z_Vi=hh-wW1#78;^-8IFj@*1IX;YLsfqG?*P3?0FX&sK8*S3oYUrat*T(W%zR}MM5eF3)~v5T)=I33>&m$XPHGz` zR~P&=&J2y;pU2db251?*oAid6=E%BI_=>WjoeSu;uPw@k2_Ll5Px+?nw|yAt?%R;D55?z%6pPM`gOP zsatcVDTuGH?#T7i=y>kzG}8udU-|j4BfJ&INAPgN*0nyze#OoK@cmdxp>r#zPOFErJFj@y?2BxBPYB$+#qKSZ`c$0Va$mE=v8w8G!;yvVK&SVb?A} zR^IEKRhq{$D`Qib+pvp$%L*vwk>8ua`+s*BQe4O%uH2*5M%^-EW|5&q8xCNi;Bfcpu$-t_D%(m>~UP6-4laeQU z(fz?slM0@hzL^7jf~5jB^kRZ{Yu|fm=UQ0%i5M=q6es^rcBbayj62g0XQ|$kIzKKR zGl80h7#xr=;q-fONe7`-2yRsCQV7#$@<@Ji?T-y=Ww2|v!m4>;?opAZw+qq`% zPKw{Fp3s-oT?=vVtPA`v$Y*a!5M*s9?DD1X__`Z$}wIHuSBAIa!hII-VPB=zd3nUBGI z0oE|Lp)80^Wn&Xio86Mvuy}~Eh>W7^6ZWBs-!wbuE{|$7c1wr7)wQF@U!GvaVgd+v zr-GAr7fMgDg`)?Fm+Kt?D2;Wvmz8e-yn9otm}DcZARAm%)Xf1WMqq5b!-e{0?Mb5@ zY3}=^wR5RqK$1U}-dFy-QVPTZWowM;SEdTa`p&1VXIoFh+)Bo}3D0EbhQ*|uCB;-> zR%&_XvTKeutDc+^RDL(K4kmIVg+cH)GM6o4n|{%>!%bfALro0j&6S_Ox{7|Ht-Y?V zHJ^}V{go82ssWP2;=(k@8rIbomI=ZreL&=Ft!q3ju4sGXb{K6c-0-< z!;tuc9)K)u6UWZIPE25J(YaV3Qyhe6!q}-596M#-_#amj*=BLBPWjOrB4zQ+UVohd zp*{LtQ1XPqi0psdeTl&7ZK|J7xkZmz2Ji=+0hk4V`BK&p;X$>Z_d{*~P5xL<&zvy6 zrr-cX3m_ zJNur^Mgg|}sjyv$oY@l4M2J7NxgW>}R$1^F(?K~Z14HJ+U zy71`r5-aN5axbl)^<&uQRyc+o5T4)lb>;rDicUFtWCR~yjvk^flX(lxpqXAJ@amKS zXMy4jK|JEk9GM47GCibdY#9^+snebAAEe6VAZf26HsW|{Oy~r^0y4$_WqXzia93&| z)bxMov$6@#V|Hfg#8_1ddDuWS?>ahyL0Y;?meY_%cjdVr=1BAjb&1kGe4@um*|Mau zT8{wy#GuDF*!9zuGB5buj#oS8m~(eF6~i8)1B&*${_ciNm>M|MILVwYxZ}9ocnNrC zNGQ(o_bA!_t}3!iu;%ZcQO7PfY+K~`Q1+yLa;JlI~TBW?vg zVR)8-{D{pOzNXM3ENk5kqKD=V8xIC{;Ign7XKmWFzDLcC;6v#GRvilL`HV%S!Zazg z;!4N4#f%RzPo}qi_+S8~7ZwK_ZO6aWhDt;|p;rZ*lo7v@<=Ef=0;TF*jVT);J@lz` zQPoG9A8Kyj;*ewCt6Vw9RDYiS_kqy!hx9#wUZVN#Sbb~ar?_*#@uepVXC^a~^v888M*bojWWnCIz`?g?pM3ac z5ITEctZsTxjntmE3qyv=u_XcEh!)2fV+d-Ti(HH2F&{&j^;@x$e3>$ogl6aEh#;XANJhreV*ahb#V_5% z{TP(J4G_EPT1O@WP$2lxS4M?>9fdt&ELn`~mpQfX1;1Wz_m>KXt?if7eF!rVY=Oba z*hcl%;|#<1s)*JI{wCsaeK?vBF;UyQ+GM;7Ffe97TnFkFF2SBoS`bH9Gv^K`ZQ2j? zpZ}8bj*7YxU3_L5U@Ak}Ub`jSAmi(dy>>l0E8m58cy=@1&};6K&_&baP2dLB=1sHB z9GlXO^M}yts-b}#Kur15^A38Yv#s+wsp9g*dDLv*f8&r7=)lkLdI zi5MiDdXgUs^71B15@=z@05bmI4M*Eo)V<9Mm6Gn2q09p|%kiHPI?~c>gBR|~E*??P z-&<9pJy|g=ZVe%o2BswG#hAm!isFgG#Ch7#AMozk+5b6!TOH_F-3`JSSjh(nLl4P8 z8=4|&K=|`YN8evdm~+DiYAE}>yD6UVj)`$AsWWU}!!)h?l6-mK0ky*%lPoLi1raCe zPDDy)B){tjFUt145fuiz5gj-Y#akXKlg$!G7(M|zo#wCm&h}V$6;~k;UD4m^-k*e# zYha)SAUvO3YzrHpGhDYm#$Zn{IZ(}O>>4~$HKqB4Jj=;w9>Co|#u0|S&*j(aIvWl` zd@fU}@#?;dW*^i3ShNQ*7^j7*4kdQ@X2Z+yQBf{yalA;9)L|{e3+AruGka7SXT`UB zqxturPxjxY6~JEaYVETTAn;d!pHIDr9+t6nb4Xij?6u?A%(+b%fAtlffA@VeCicpF zLYG+tBW+&YYLT{Efyh$*3uBOSx1pzC+cZz(^kI)Ko464!p#T?Bz&-(3o2Ab$XkH`) zPiod2zk9=>fmruPt#kzyr1|{Zk?$FmXsgZJTyDEKEW7ptiqJsM*Z-ps2E@9FX=X(k zcqKA)tyU51d_jo-uqV2v*Q%w-5!v(=Xkt&%o--kvvK<1gC(k@dk>bvhe*<=Rg@V2);E z6H*sqOXVEj9UC#69Hdg$cC9HeRqz_Et08IfRy#UAR>XY!7(x$l+kfx)`G$(h0F%Q-JWDdc=}{&XRmNWD3{S5TBK4_!Z3{!RmHy8D*rvlPXl}>~PePs7 zeW)hK@Pt@bLOs}BzajI>W^!uu^0I!5m@SOC!fvW0RMaGePxQM(PDfOAEzDHbcFKuK zA0(3_brNhy$t~Nif!yV`_WCzHG%(8en?(#Dq1%XyI!IFQWWCEr8ujY5P!`l29HJ)Tw)HQJ1K%NKcoffUvi$`YiK@zvy2&nNEE9 zV7e_zw$qxf+Nbw?WshorAf5R_Og(g5?CWQw19(o0h9N(lb@eLvYs>LS)nuZod#z;vrf3$OJrR+XW&zJezrv@(KUCZs z7X@1&htp-mm|*~CXIQfKe*99ZpMw=f>8hs|%u^;tmeK8rITF!lTLjRk|yQDXS_T-QjKar6?ic zhp=N!N9~7j4~}nCf&{-i)+Uz}=y8)Iw6_vd8SZc1IlXfWb;$10YZ|?d5MR-5bG^x7 z;G0l%FK$>i)Q=civpqtnJPaFcSuU`k96bJKMF`Nucpc6h!D#7_oc<-=6QX?@EEuXs zCmVN-R#+gO!{dmmVJ~_@sC`aYH@L+FXIx1m@Q&^9;z=*b(%!k{raH90qGKQ zWM;}JoMfm|<#0QhT7bM^%OoGQii|O#aHC~gRfI;p&*1a?mdYD}Bl58=MyT-3N_Vc( zn{@xrmblxrr_=X!(^zo&kQ}mdD-mI>I42iQ&RXRh`J?J{d9QeM5)QAmlr#xkVK(tg z-g>&s!y5Fn19?+(o1C0ah4^$E%h`V}m*x`93;oHd&k4g{?|#rPO|Lb^{;<|(JPS|M zD#N4KYLPypKk+yL7pU& z*AxBZ>zYJL2k%+soAx}(vtPk2Z$Ih_2$);dD@OJvkYzDB_4cXAhns|A`=6fHIyxo* zUx#yarZ~4+QP1NC_bw$;|COcVr=n-8-dalb)s+YKMky}ZWm?YJZ3azqtuB+7Is(yb z!YQMi|4nEAA1&4j-*501k5;9ASXZ0vDGTE3EFnEQ#*2$FE;Jzz(^pF@xr4U9Y~3na zk}HzqD%b4#nTl1UG-y~uf~g%344-Ta&=E>Gx0K?U@PCcNH{%(@|4vK8&w<9o*f891 zRq_2hq6}1DHYf^15Fv>Tv5)0iT~CQ;te##K=>6WUc(ki-(Wq{AsTT``oG!_{N8zyq z+f3b8d#J-Lb-;(iWga!=Fbe3YJ2vF%l>*mFA^Z3({B-Tz)-Z&k8mPnHN$Q|!UDvm1b z^ONPZk$Yts#Dk1vr1{UxEQqrgo)kaDH+XEiG0#jnPxaN{wTfwZozjJkXkJsASyPLH zV0Vb>LM-M%z=mtGWSOLSUJnr=2H{Flc6M;P1bzU_k2UU7pK%VXO_9H*~^>z>_~;W{&9+6{kp?W75GoL3TWCT zXLkM&G^;bA_B$$>X+5DvlJ*qGoD`cR*Ctpi`C}tXlDl-1GkoDP_~#Z^a+4)@LGesv zZYntU(v$~O`~F?FrD>6fMzWegX^E!q`u@XZxLw8#jQ;qB!vFHrbp^ind1uew($#8$ ze-hgjB-_5UCA++4tD=z2ks#q%(~m`XSNr*T6U#fVs+Y%$d7l??IGnXe33#WlFKSZ) zFg~|;yoYAV5D`rhhwc|9n=~o^NijSjEC{X@D~6}`5>sGpP}rJogK}&}WM2=KJ*s$G zaWI#kyq~pa?&I^|T1;xusgt9yn}&|&%{?@`pSC-ddhlp9xiq(xQ2W;4mZPqKXEK41 z*VXY1uW{bpEVDWKL|)UUo+R;CDc9eZ({tpuDUa*sOI=Er(B>23L{p+wR8%)#PaI_EOXiUI#b zkA|2i?D(EkxCQar;l{Sp|LN{6o8k(WXyF6`1PKy?TX1)GcS~>^2=4A4g1h_R7Tn$4 z-95Ow`@55K?)miofLoUus;1_ds(Gxt*KS#jxqkyd4NNRJ0-WFsOzG5=kxY!r!OsHs=fBNCcj1+ z_W#2BQ7+zL*e!G6>m#A!rF$f7*yh`LN!zx6$i-z^172fIa$jNaN~hO)1Y2uB3inPg z%w^S@dMP#adnrYAx0*`1lHyeZldmRjML}&W3;QO_=l2kX$i7rl(irH+vq8Z{VM7wz^~Py{-g_=pM@2 zE$vm;YK{+!OGDH*tVF;;0GHo#wCJRC33SrOn>OdUDhKJAE)3~IER3R`1B<$KpcwnW zJZ*So?-&YSDqW)u?HSjbvSG77^HAzLzt7wXUsDx)aPNZxB;(}-ur-D}4nLq_2q~U% zhz=b*nGJKQt+8Bi_y&;5qSdaw?O15e0KN8DWe}K|7sxePK4VcK?oQXS@teIbLJPMqZTglnp(2*9q9C6>8=3E6RjV_dbNWiP#*eRV zR$#npW*v&Ly{XDATju{7n%(B1O_Flbv_##$tJWcAx3kH1AZWI$?%fODCXR4sHA)gf znEK)GE(6Wv2EkxV=@ScP0CKb%@vZ=%@u=9^sqVnJ=AgFTq=toxWmcMPX3d|?R<_q4 zz+Gshfx~d^l!I?26nDPQ$NW1+ZMBaFso2)L#xGd_&fI*K&*aqB_Q#d11>(>rqrz$-Y$5hm3g@ZL`AoqoNrciW=;QdaSD&Y_^%> zM6{DOJG=0jlHl}IW6j)CIX@_dN4MW@f@}c(|7Sqjm}ML;&4_|^SIWp;QrcZf80Zmd zfd=1q7hQQS{$(v=5cqigxvcnPPtz=${%X*3t9fg!=xUH$Mk?pXF z6fJc_5SH2i0=;3MYUR|SdgEwG45E`PeY0uQH0ux+cvEJ?elHRK6Z4i`63ra0FJoKH zOnbuwAOAc9w<062G6!cmCtnH&0AE;;Stj;Y=eJg-4OeG%A)q1fA=n_~AdDcqAQB+T zAo?KIAZ`_suGkrp`{gf5iRrH_&&#U1R&Ab_ii*P>^-5)j(kug7S8sh^+H$%lZDOjY zcK}I3cpT`NvKT)0INcXsv8Y=YJ z%gsfS?X`>aIR^IQ?t(jW>*U=1j3kTAh=eI~9|cXe(*(SAhd*w(i#qdxz0WhwC;OG2 z6V;#EBm7RzB>9)MT(XMpDYTYSjz}SK(dZiMt`(%CUBvVCjk5*1i${H)OQDihj9So#Y+F1nCVN+z@ubXFlyg0At+0$>qB-JjY zFE9V))>!H`vOk>hVyTtXU#b>CB zZwnFXX##Zkm~*>4>}lVpUkH^!31pnH?!JtT=o9HYmt6CbKfPS(bj% zh@B8bs;}xjLtjNHZy{W@k<5At(|EHjon(zax3iV?9+7FL6wIkmB`wotVmY+> zeMd{ao_Yz}<(`_mo-+3mDR)G$=hB9P@#y5X`L<)s2Ravc9=dER9OYb&v#X}Pwkp_K zrMod-Rh6$Jc+?9oE)nUk1OZQ*NEDJ0DM#{;P2*(!v?>XVPwtH+7!7d$9BBWF3dO#u zrR>!1omW}9ywNW3bL{Q30 zgXMO1xgvMpgVj9l^yc2qsK!1$izT51Ekd);@s_cHPl)p>fKQ9CIKq~lyb0= zJEWX0Ihj4&ag#^%e(A1GcvjEtsXs7Z_zfX;WNM zPQ9Go1(~Nz%7-%>p*okiZI>Kw(-dvXNbkr41Um08cV4-)wgA7t?N^luyR|M@FjCPZ*-v+AQ{4V}{ z_nj1mj5*i?&Le#4UJKXty4PmZFEX4JewuF zh3DoNuJ+n(HV$#t?|u9N;==&;;Dn!hmK?k1iPd3a5J+bBId|{@gH-TryZ?=ne(L}g z1jqEIm0sCQcHyFbzK9-JU3y5qRK~}~>MhIDUcE#5AJ&h5CDg}i`EIT^7cB;$^k(Er zl6~yXpI^y}()aSPE>dEEw7!|@^ik+Jc1XOU@Lke_^Of=Q>$iqm$H8|6myJR#T;OdH z@|zyx4((iemOFVPC8P%<2OR%8369zHIqDZRuD6Smui9i@o~_%S!A0FE;aw_thkP%$ zwmH#I%%ZZ=tHfGlp>7YQex#AH>nCGJVw_JnsIscOwET`~CQz$Nk}uEh(DCdD zsdQLQX>7P$dIK1OZgRMDX?%s89&&kmAs)p+GVw_c(O{*b0Hs`?u*ke$t<^<7i4js! z(OS8|LfOH3358PmKDy}^2a{W zxhdo?4Vg79Ia`L0KXBkmAbLHTFa|BmIBR`gw-0x3?>b3^E?o_;=dvJ-sEn+;3AK&FXUzYi35a&e# zvR9~oX!A@I9LRJZ5>a@N(Vdo(*<%e0*z+y^S;-Qjrb;kR_aI<5ijt1q>)oc)sPjoK z2Q*kc&*er!f`M4Oz&4<+_ierwNP+~iS}u~Kv8<-Q7G%cb>GWO>8!wSC*AVYBc}dUX zVnimFS_biB1%~&JtFyr`@9^0~%-Pu0)=t8*wPh50RjI!c(z?#&|gwo8ycD+Y!Me5DR zH$Seqsir<8tUbn};+IxZiAPPG&@*H&31c41_ni`Msz#zBfA(3a@?ur($=|KHvo6}j zJ4FgxycY@wnu{T-iZk>*hj9Mbcz5+2+q(7hlRM_p<>O%xd{c-|GL#3!?xf9ftOWOV z=lcZYW3NuHk8ez(Z>iKT_GKd(3spK}k~)~N3AhQ9Skbk~ad+92eDWFzUxcA|gv>~V zna=_P)?ja1aJC)R$K{r~#b!I#CT6;|sb|FbfY`s+(8Ey6i=nH|mz|b&qXQ_Ssr8Yt z7xRGieo@4XxqYuMhR>bpKZ*J8n(n#FtIL)mVP)nUyI*36gqNtcB2=a~QYRKKKmVhw z`VZ)F9Qab5&t#6C@~Fm4-~KDM0|>VC z<_I&mO3UDfcY^#f?J@Ryz&D&9v7Hn8(zLhb=xqwn8)D&Cf?bLF_8Rn&t5+t~>GpdV zdRUxIQn(pwx4`MPa(9>o@Op96oI?kgJ2jCuO97mv-QWa5q7)d; zAH?C1A}rUa%)*CH3yz-0OWJzap4#fe9rvHV(4Fa#&sIRdCnK67v{E#I-==6WmLx+R z9V@gZ&0{9U6&OP1Snj5(fxJkUeLZTlAFjN8+Xii?{6hx(WvH(eqWivOF9bCw4@J7` z(M2zN!?)Ui-HLguH90NF<|GI*dV8=MTP89;rS1$};a{P8mdQOFI4LuChQ<$yJfa%m zG-dcD({M2crYPrZ?_lhY60AOdzj}^xbB<-G2erM|6DK5*>L5|cF**8~1pR;%buTOZ z)KCTUk2-4FM6xQ+!VDKHo8Xqs?^zVJobKzo-QJL?#<1;!fj^SgiQmojqoVMWpYZDG zYBM-joUW(Qv*c6gy;oe0a&|)r>hZE2@=ovf?vC3Yvi1yvy4S>0EVjvBg2vD_9677o z6M62?@NK1N@spn5UEt$n6^Q^-2~^*uC> z*<>n8{?S0}#skD|?~%au$h@GJv0>->13MGdSJ4kmV7k336|~SNSDxF8F&1bz_DCkz zC?dPy7nXV-ao`eG#1%QT7#X4&8LSqj!~##~+aMiqkr=*JoRir#4%MtZQGL43*(qFvX+dsBGHwDZk`%ZZ|@uD}ir z;wt_Pow`rymEK?4hmoZe9HdJ^fNnsH_mNk20)7-Zk!X=ht{>TgGMmq)Lu+=|HgppWp z>8dgT=r03dGZJCu(?GFc(GPC8XEHf*MQ{$$oKwgGN+(%8QGg-_x8O3W8Qw}J22mb{j11$9y6~h z{oP8-BC%b)xKfq4R++d|p7^b9WVT>rk!WPTWn`UnWSeTF?T<*;ZS!F5&;h>G1+RX; z_w{jvXVTs?(4=O zuiTo)?zk^VE?(>|>MAXDH29VQMWMz5#uw-ZJyoKpS3P&@kcvHbD|k6QiHj3ar3Z$p zPL46B`w66fDOVlY6$1&O?sh-U0F z(!Xi5KZzOV*tzfE;!KJROY7!~w~R)(Boo}S@E_~ie(YMk{OXrF`(E|qvVMAms_!bv zX6tM<#K^GFDt2~vGb&Tx^kMKr4;J#Ll98$ClA4g?i2)0YR9sYlEz&Hlyy*w20=PIi zbc%EeB_8cuml8GA+=W+N$tYD*DKT|ZO?C3-mQb}q*s{E$RM4uh2bRuye*nA%K1o|k z;L!;-F07Va@@mHe42IGQp_c;}jh#T=F`tGo$Amp)ymjpQ6?3M_H~pRkI^c zqb^60(btBS_4?d-*7knB=4Gm!ewNyPhC+FEFWGCJur$xIny6J8|IIAP%?*?(#;l^j z&Va=YIYZkG*#8H6A196zDGpw|!w#jk5IVV$dY z8cee)l6oDx!WtN<{z*Qav@0^IBmKwa_c!q}ZG(AQf2*Ff=rl-m1&un8E#$XJ;X!~5GT3OENqx2&P9Gs+Iny4dy(~f$VR%Q;sx|^Ey|ED z?aw`dEN(d$PaJF3^^rai_pjE;HCuw6)tfyG4<6(FmlKUnb&}jXtTrE#E9VQu_%bl> zUxn8J2aXHE!M6_6L{^?jrwmkFI7ImV!1*bDtPka{=N7qr516yk7dAyUV>Yi-U7R-G z_3lDM&Jx#9%KhNmDV4ot+jgoHoG5uacWh6I)H+bmR*UHOVfuNG>&Ki;SF!OOR#R?xGeM4lzUV7Jy`p$RTBl9$Tq}zDp%CL zw@3O^LPa~#l4LGV`Gh-*)lJRWUga0jR?rv;C*)I=2hNJM-zFHs2+rBg^$}g@(;4UQp_KAs|7-qS=HY>e=;P> zCW(cQyP}>cN<}!-EKTKJS#pPHe$syooDS>{=_5bNH(uUDzo#JA!jBvzYg!QK?;M6H zXyn}j*P&=ApULR<8;D(*c&vUCQT7%vEjQMUwAYTbRH=vBca|#m(#w#0{;4uj;Wm?G z))QxU54PK%LLIsY5d?B55=rYU2V*|UM+HYHNrp`fhgFS*+6_fs4#+yl4?BgFRNc@_ z*U~6YP@%;aH8@r8juh?^ADt)_@>aEg$?-(co#|yj99CNnipt=!9J9T_e=G@lXD9ENl&itCTXt}%R_h7nQ#$PhZf5I zXQS%;$5~id_74xKsHwR*IArAI6$^`#(o;aw-@09xGB4boHx?+v2?=Xo?T2_V(73bE z+-2)vv8%p{(&N<|9FGop4TB%3F<$6K~$mh-#m@uChId`VDzqRqcemPgfPs~&eq z_!!&h*=KJ~W$E5W*m|&qB+asQgukPEr8folo@9(dQncZF*k}ZH_9*%vM_fFF;7Y{PjWD~iZPGM zF)@B)BB7KdA!jFH7A1WtOZt+Pgj<;OnKlEPIu@NiKH>f=u_z6(G7IyImC^1=Nr|R9 z=X|wYxzP@2Y^*&AWJA|{E~Bz>P?8<1E*i+h&MCJWr!`8iOo?DmGlMkToDzLg-%l1~ z+?J|JqqLN{K9Rx0z~5$S*YyKsI+$!P7)ywE3!e2w(LYV!bX3^Z@r4krm5+3l?qStw z@yFNIzQI;E*cT-IK$dLr86t9}F5|fFbbf=QJbev90};Ahak5-xu3R~$8X2M*0Unk; zL6-mnks=4B7z-J&heeH$N3McLl9W(|nO>$wcDq}4hDuC^TvfZE*-ulqC_}Lz&$xaj_rWF?vk$H*B(B_O4>0V3 zBuM>A49wVSd3Ze$*)OrU$XrD#iR!uXaX+= zUIf}_{7*5jX;T~BRg7Dgwp$^vTSJoiMT7c9lsQy|#8jr;R9M1P$-Y$TFGGu-Qj4u2 zoV+1eCm{#4QiGqQf@_Ct3gFz*v zLLyVb4=X1kDtr84G1`e4wB$Qk8)+rFNv73?)ZPk7(EOsZ?&yU_o`yZHF_&H);9huH zW#6lOcZtJiNFHFJ%n2{1S%fpHK%$$ zqVhc~Cc-Y(Pci$PYObhxShaL`wTvUhg+twxYVNL?`na?N&D9jh^(5XT-VbDgKkbsi zm`wY&G@|G^%{^mU>;77xceD)=PbCv$pc}p1x954}sl5N7ADvXu7`Iyy2!piGPE?eGIH#QkqI{l2K9F& zY))b`H4llkO&F})yKI_)eIt3-EY7PZWDn(4vMbr`uT#5_~vXzv;4oYz{4h@> zkRKQ1_dVB{y{|c+>h2Z+YS*`wBM3^8^3O|!dNZ0WsQ+_2w|He7?C;yy+7iyIo}8V* z)W7d0tX4@%PxH^;D+I^byBey(QCXoYuhABw)0f6mmaCSSeJi(`Dzjw!ZOv+KQEzqX zU~yt+es*zm2!DKvbcPppKp=;WD@lMqjF(1%ok5cUE=x75^(zhF^RC( zC}E&yL4VCuN2q12Xkd}h{H4P5VFsxtNjwp<$Q&&9K?$#>fG{6R=Jy~L*<@D$Hw+cJ zo=>l~L1k7tVJ58Ho=xUbA^RUqr_)CxExOcULt35b&#A2q6)Y_k91Qgn%-rjoyz^|F z%RG$J416GNE_GHWbv_<#9u^f2hLmOgm|1S{8h6|jbJ7AY^7<^n^diCRB+l9#&eRg> z!qmz7%Fgo4!R*9__WG&9`l-^qkPir+}%bsoO93-Oc-(Ufm$TP?26jQA!GJ~0usNx-jU z96B=f?kI$$*;K__g!p@f6oOKsRhzs7BAkA-8T|;>_?Bn)tyICLSkI3KJh`eecVl$W@k^k({f(3 zpEW!)Ol-vf-K9IPI*(bdW7WSDtlnlg@Y}4sbbP<_6<owZ%F@ge5X|C6#=>4_qwxCkLx~I5wU8B`W7zRr9JG|ihkif zx^rz$&tO$9xWd=k;MrvRRHJ{#8@5`2)vQBo)UGk&HM{>ZGzk-zjYJ_v z5EaSy6~|M6?C_`3%=^LLkVr&GLE`wxp*$4Ww&&VCo!lre*5bIIBRCwWHz&f+h3C7G zK%clhNRQ9+D!MU*k#s)tdeB`1jeWX>vI6A={c!nE?6%d;qlr7k5n%{8zc^g~n3#MK z6u*PQd*M1?*s^s{{WG~bd?EX1l2=gtmaaegPj|l%A^{So1JU+GVU%!zFe0o!E)Ncz z?H`YQANVX7x&ZtJ+@?3Pg;#UkCu_g6cC5BDm3zL8^Hw4hcWd$ER@BNn?H>HeAl^IT zyCBh1l*s2@!EN@K=;pj^r(uzkgzzO=j9N`TvqYg8y59Zr`0RF-h*u5LNAubJ@>%`1 zsh|fb=3VO2q;mW|k3zhsJh@|xrb?atUau^Ir$)hJ;{e+XoL*08UEDiLKgpE>qiCH} z(PO5_1w)iJoGz+koZ<@Jxh!`}q1%Ok<7)UdbBvaCK9hFg39tSR$Os>pG)SbT54d(0 z{KwE6&NHiLL4jwC^6UEAe8}V;9uZG28P^wl8+)Ep<41ArdxZky(?0mdzj($zxPw1A z@#7pi(S-$FZCGItfe7iu8j(@9Jg|u79?Zbg@h%0$uo3fLKjW2=b85-gpvu)WrkQxs zPhNUu`-{Xh8pIUJ!~|=_R0>r^D^&SNRkgWQrHh4>sf0v071dc4C3vMI=r-kl4+xOV zZ&J@saZN}utSfSi2$3zTvP`e?Y>N^0iBZhTwvW!14*q3c_)D|#m$-k6cwUQfE!a66 z?3@N3m+p`9ILB^w2~LVZ!ixK3nh0B%h)5R=H!^@W zAFEXJS%s}N)5=9g(YQXMPaImxn<4teKW|~&;`C)kaN9f0y38~PUJWR>eoh_MzDPey zMtT|Z*!BMFw0IFTjtBtr;<=p)-TRxPVEE_Itw%PJVL2=MoQD7$ReAwfF-GcDb(~Ui z4~N>KT(5Znx@hQ=vrfT|bXc-$jkJP!X_LyllHrhxmfgLrTeHN|ma$i!q1UvlPlnuo zJI?$oa9O|tS`E!*l<9_5LcRYqf#9h4F&j(DkCR1=d6uRssA5FHU!@Rjk^y9uN*8EJ z=f~s+!biCkq2Ltob!;3xY*?`%g~Towc|yy^6zW97t}@U>&z9+Im082XuHCs55tyzM zT0_;SM^?VnHV}}m*GM-xo*Q})hzg%~OMGEca&R*#t5F^Pr$X`&N> z*!WY92%>o#qB$`~obIJZBO}6r<$kD zeyh&G4QEX!uQzEs4M(pxyhuTyYL~?~t;{cKWZ@rzQ&$A5* z;eMBagJ?~K@s9_yeaWo(i>&C4j&lz;PX6I?y|QI_qxfj0Ys!|dvv`7zX}X_QzsvGg z(t5&LzEI-%9;^y;wh}|V==b_=?R2Bg<|y85O^+&8<1t_H+QUP_*PviyLb;>xMfLah zH#avAA{qO6|L!6^$EPTF#7<=A6W_B%sMTsDk}kS$GjswoW8R!V zGhLzk7(!^>En=fBTB|K6LfsQms}oAyt%nZmE=Zv%xGp@Qp&{9!DcL@_F6Ft;U7}W; z+q7wDG~lmaM$^z}!ErRW@wB+gBe=?=qq&YJak6+bBX~2k$E-BRthC3qG=#_a8jGfy ziwA0pLi-wvhw96Rn#-6pfmyY+(Y2raDr>`oYikeelGHUd7t}P>G&PshG}VczF|n`! zT>Gv4tA{WMwAux6%I~LxiHYg4S9d6e{NB)bB;&mY$_GK@4}xrw{{B!!(LMl!?C#AI z!a`_Xo}r=Pj{O|FFfjV>c8)~ z(STURf4}EMwY+L_lDGgH&nmM-IS=r0ve z6aW`!7C1yc9i0eyXs@S7#K*VR)sSh=t-D)TR<^XPEJ8|1JE1^@$wP~SgM)&a`snbGjfEvD zctCo3ecjs1N^Z{B(2#U7`)7#1jx99(uVa(C=JD;_T`o>eS~|LjhzJc04JSuOfKaTM zk(qgGb2BC`?n7nS=DEGT z*3v&72uMmvfzT)H$19oH<8tm0Mm9*TUiHk+&hA=VOg_%2tgk1)#?EGn8vSweCjAbZ zVxo(Tpc{tQVRWw|sy<^CB=2AHrVV+6{R71nO)-PqoV2vz-5b|;cgssln;-pEVa!+o?Y%Eqtk&k{#L$JcA)=B;Mr0}yXV*!6T02%X)zyQ6o?`w9~9imo`u zwzeo^74`KP_`c&3l9JKr?|*IY?HQYx+}nBb(SK?w3kVPs6nua4#001s8$Fv1n?YS& z-KM6dZFMLZ7;;k5Xv2R8yGZf(_;^rFT$`Jk`uh5cii%?D-T{}2lA7Ar2Qa$jUNX1* zA7NqfxEuzirat3Beh~$MJWWku)8kW8%&e@&jH(BH0iiZAGn+DDxtt#w8VUpHI_bhbr?qjoW#EP`8-g>0#l^$;rufaC9r3O{6`L)-s!;@^MI1JcmIUOM`U04Tj;_a8<^MkJ)9fGX$&+6x&V z)8hvSy20dugS)#spv#MTm^Ox^B2}Ldd^=`Ca6jcITK82(z`j75}8I1obZrNr%M@N>9g+9R$`*s(Bdoar8RgDsVP3 zYe9+}aslcdX|Wa3@-8|W<(y~&nkn(a(nXZ};SA-u&*P(8m2LZyu>cjI03jeD35%<$ zVt7y`0P7UkZ?Fdq%LC8j=Yhai?e-eh)Wju$A`uFlY7mxspEnjqN>&!37B)bxf{cj) z{+#O@EGsWhO+^*6W6FZvf{W_}EahE!Gn1W#uCAc6a&2MZ!2=)U9#RbSruZwI{|3nh z#GuYUgtmzM?}+g1zMwaLephqW36q(J7^%vBq{^Q^eSnXTPcENM0{V44{$BGP*1*8P zS9kt_nBw=Qg=J;vXlSIGZ^fSZnVABTlEKFvJ~`2Wr0Trm_-ZdDvoFllC6azQ(FLE7 z-dS6prKDgR8yj<=5)%_sAe7*rLw8*G#D{`(cfNIm^zrl?3zTi|>;NnHy{^g|`iS>e z88aV>EooaoQ^{C>e=g|uVW@c11$5}j-kh$aUIy*awh={%-Y?qS&$P7?pR;H#K zc8!?=6PbZ}QF4z6p}kB1e{x8O#=#U89$u^W8=rxol)b&ZjLgX4TDmb%OYq)fYw&?N zijSQg2E)+I3=XPey-d6HaAQL-ft{V*+1WYp(#8$Ir{{m~aK$U|Rm9Vp$7~6k*MA11 z1@7M=Lz(moZ@J$fc$q{1J1vCc=~(S`qfhY;(dJKw4-68%nwr{Ao~VP~-d<|^b=99b z)YR0XqM}*PgR$g;$#=j^M^ZMl|3LrvNc}Y1G0y7#jnCM#8ZBLG9pUo^CAI!xlFaVm z7%Bt%IW7e~2}vBFYIjt1j?Lyutf498Z?Nv;Tfk@o5H?*6ZCHP;RSV1CS)BudQ4w@_ zXm;rQ`i?2{?&Ou38*v3d3cwA}{>$4fcrOa5_J^xI5D27BgC5_*4(PYKU4OR3%1TCI zVPU%|lz+DGKTin1-X#u)M@E*Hmgae_!*+1O01zIFJMf=z3HW V49l*w?RUVRq^O)onb414{~tYJe}w=5 literal 0 HcmV?d00001 diff --git a/src/main.cpp b/src/main.cpp index 1850161..37729d3 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,83 +13,208 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 8; // feel free to change the size of array -const int NPOT = SIZE - 3; // Non-Power-Of-Two -int *a = new int[SIZE]; -int *b = new int[SIZE]; -int *c = new int[SIZE]; int main(int argc, char* argv[]) { - // Scan tests + //for (int i = 12; i <= 20; i++) { + int SIZE = 1 << 8; // feel free to change the size of array + int NPOT = SIZE - 3; // Non-Power-Of-Two + int *a = new int[SIZE]; + int *b = new int[SIZE]; + int *c = new int[SIZE]; + + genArray(SIZE - 1, a, 50); // Leave a 0 at the end to test that edge case + a[SIZE - 1] = 0; + //printArray(SIZE,x` a, true); + + // initialize b using StreamCompaction::CPU::scan you implement + // We use b for further comparison. Make sure your StreamCompaction::CPU::scan is correct. + // At first all cases passed because b && c are all zeroes. + zeroArray(SIZE, b); + //printDesc("cpu scan, power-of-two"); + StreamCompaction::CPU::scan(SIZE, b, a); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + //printArray(SIZE, b, true); + //printCmpResult(NPOT, b, c); + + zeroArray(SIZE, c); + //printDesc("cpu scan, non-power-of-two"); + StreamCompaction::CPU::scan(NPOT, c, a); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + //printArray(NPOT, b, true); + //printCmpResult(NPOT, b, c); + + zeroArray(SIZE, c); + //printDesc("naive scan, power-of-two"); + StreamCompaction::Naive::scan(SIZE, c, a); + printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(SIZE, c, true); + //printCmpResult(SIZE, b, c); + + + zeroArray(SIZE, c); + //printDesc("naive scan, non-power-of-two"); + StreamCompaction::Naive::scan(NPOT, c, a); + printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(SIZE, c, true); + //printCmpResult(NPOT, b, c); + + zeroArray(SIZE, c); + //printDesc("work-efficient scan, power-of-two"); + StreamCompaction::Efficient::scan(SIZE, c, a); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(SIZE, c, true); + //printCmpResult(SIZE, b, c); + + zeroArray(SIZE, c); + //printDesc("work-efficient scan, non-power-of-two"); + StreamCompaction::Efficient::scan(NPOT, c, a); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(NPOT, c, true); + //printCmpResult(NPOT, b, c); + + zeroArray(SIZE, c); + //printDesc("thrust scan, power-of-two"); + StreamCompaction::Thrust::scan(SIZE, c, a); + printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(SIZE, c, true); + //printCmpResult(SIZE, b, c); + + zeroArray(SIZE, c); + //printDesc("thrust scan, non-power-of-two"); + StreamCompaction::Thrust::scan(NPOT, c, a); + printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(NPOT, c, true); + //printCmpResult(NPOT, b, c); + + + // Compaction tests + + genArray(SIZE - 1, a, 4); // Leave a 0 at the end to test that edge case + a[SIZE - 1] = 0; + //printArray(SIZE, a, true); + int count, expectedCount, expectedNPOT; + + // initialize b using StreamCompaction::CPU::compactWithoutScan you implement + // We use b for further comparison. Make sure your StreamCompaction::CPU::compactWithoutScan is correct. + zeroArray(SIZE, b); + //printDesc("cpu compact without scan, power-of-two"); + count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + expectedCount = count; + //printArray(count, b, true); + //printCmpLenResult(count, expectedCount, b, b); + + zeroArray(SIZE, c); + //printDesc("cpu compact without scan, non-power-of-two"); + count = StreamCompaction::CPU::compactWithoutScan(NPOT, c, a); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + expectedNPOT = count; + //printArray(count, c, true); + //printCmpLenResult(count, expectedNPOT, b, c); + + zeroArray(SIZE, c); + //printDesc("cpu compact with scan"); + count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + //printArray(count, c, true); + //printCmpLenResult(count, expectedCount, b, c); + + zeroArray(SIZE, c); + //printDesc("work-efficient compact, power-of-two"); + count = StreamCompaction::Efficient::compact(SIZE, c, a); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(count, c, true); + //printCmpLenResult(count, expectedCount, b, c); + + zeroArray(SIZE, c); + //printDesc("work-efficient compact, non-power-of-two"); + count = StreamCompaction::Efficient::compact(NPOT, c, a); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(count, c, true); + //printCmpLenResult(count, expectedNPOT, b, c); + std::cout << "" << std::endl; + delete[] a; + delete[] b; + delete[] c; + + //} + + // Scan tests +//int a[SIZE] = { 0, 1, 2, 3, 4, 5, 6, 7 }; + + + + system("pause"); // stop Win32 console from closing on exit + + + /* + // Scan tests + //int a[SIZE] = { 0, 1, 2, 3, 4, 5, 6, 7 }; printf("\n"); printf("****************\n"); printf("** SCAN TESTS **\n"); printf("****************\n"); - genArray(SIZE - 1, a, 50); // Leave a 0 at the end to test that edge case - a[SIZE - 1] = 0; - printArray(SIZE, a, true); + //genArray(SIZE - 1, a, 50); // Leave a 0 at the end to test that edge case + a[SIZE - 1] = 0; + //printArray(SIZE, a, true); // initialize b using StreamCompaction::CPU::scan you implement // We use b for further comparison. Make sure your StreamCompaction::CPU::scan is correct. // At first all cases passed because b && c are all zeroes. zeroArray(SIZE, b); - printDesc("cpu scan, power-of-two"); + //printDesc("cpu scan, power-of-two"); StreamCompaction::CPU::scan(SIZE, b, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(SIZE, b, true); + //printArray(SIZE, b, true); + //printCmpResult(NPOT, b, c); zeroArray(SIZE, c); - printDesc("cpu scan, non-power-of-two"); + //printDesc("cpu scan, non-power-of-two"); StreamCompaction::CPU::scan(NPOT, c, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(NPOT, b, true); + //printArray(NPOT, b, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); - printDesc("naive scan, power-of-two"); + //printDesc("naive scan, power-of-two"); StreamCompaction::Naive::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); //printArray(SIZE, c, true); printCmpResult(SIZE, b, c); - /* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan - onesArray(SIZE, c); - printDesc("1s array for finding bugs"); - StreamCompaction::Naive::scan(SIZE, c, a); - printArray(SIZE, c, true); */ zeroArray(SIZE, c); - printDesc("naive scan, non-power-of-two"); + //printDesc("naive scan, non-power-of-two"); StreamCompaction::Naive::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); //printArray(SIZE, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); - printDesc("work-efficient scan, power-of-two"); + //printDesc("work-efficient scan, power-of-two"); StreamCompaction::Efficient::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); +//printArray(SIZE, c, true); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); - printDesc("work-efficient scan, non-power-of-two"); + //printDesc("work-efficient scan, non-power-of-two"); StreamCompaction::Efficient::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); //printArray(NPOT, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); - printDesc("thrust scan, power-of-two"); + //printDesc("thrust scan, power-of-two"); StreamCompaction::Thrust::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); //printArray(SIZE, c, true); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); - printDesc("thrust scan, non-power-of-two"); + //printDesc("thrust scan, non-power-of-two"); StreamCompaction::Thrust::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); //printArray(NPOT, c, true); @@ -111,37 +236,37 @@ int main(int argc, char* argv[]) { // initialize b using StreamCompaction::CPU::compactWithoutScan you implement // We use b for further comparison. Make sure your StreamCompaction::CPU::compactWithoutScan is correct. zeroArray(SIZE, b); - printDesc("cpu compact without scan, power-of-two"); + //printDesc("cpu compact without scan, power-of-two"); count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); expectedCount = count; - printArray(count, b, true); + //printArray(count, b, true); printCmpLenResult(count, expectedCount, b, b); zeroArray(SIZE, c); - printDesc("cpu compact without scan, non-power-of-two"); + //printDesc("cpu compact without scan, non-power-of-two"); count = StreamCompaction::CPU::compactWithoutScan(NPOT, c, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); expectedNPOT = count; - printArray(count, c, true); + //printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); zeroArray(SIZE, c); - printDesc("cpu compact with scan"); + //printDesc("cpu compact with scan"); count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(count, c, true); + //printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); zeroArray(SIZE, c); - printDesc("work-efficient compact, power-of-two"); + //printDesc("work-efficient compact, power-of-two"); count = StreamCompaction::Efficient::compact(SIZE, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); //printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); zeroArray(SIZE, c); - printDesc("work-efficient compact, non-power-of-two"); + //printDesc("work-efficient compact, non-power-of-two"); count = StreamCompaction::Efficient::compact(NPOT, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); //printArray(count, c, true); @@ -151,4 +276,5 @@ int main(int argc, char* argv[]) { delete[] a; delete[] b; delete[] c; + */ } diff --git a/src/testing_helpers.hpp b/src/testing_helpers.hpp index 46337ab..6cbcd68 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -1,8 +1,8 @@ #pragma once -#include -#include -#include +#include +#include +#include #include #include @@ -18,7 +18,7 @@ int cmpArrays(int n, T *a, T *b) { } void printDesc(const char *desc) { - printf("==== %s ====\n", desc); + printf("%s\n", desc); } template @@ -69,8 +69,9 @@ void printArray(int n, int *a, bool abridged = false) { printf("]\n"); } -template -void printElapsedTime(T time, std::string note = "") -{ - std::cout << " elapsed time: " << time << "ms " << note << std::endl; +template +void printElapsedTime(T time, std::string note = "") +{ + //std::cout << " elapsed time: " << time << "ms " << note << std::endl; + std::cout << time << std::endl; } \ No newline at end of file diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index cdbef77..4bb0dc2 100644 --- a/stream_compaction/CMakeLists.txt +++ b/stream_compaction/CMakeLists.txt @@ -13,5 +13,5 @@ set(SOURCE_FILES cuda_add_library(stream_compaction ${SOURCE_FILES} - OPTIONS -arch=sm_20 + OPTIONS -arch=sm_61 ) diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 05ce667..b3f5567 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -1,26 +1,45 @@ #include #include "cpu.h" -#include "common.h" +#include "common.h" namespace StreamCompaction { namespace CPU { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } + void printArr(int n, int *arr) { + printf("["); + for (int i = 0; i < n; i++) { + printf("%d ", arr[i]); + } + printf("]"); + } + /** * CPU scan (prefix sum). * For performance analysis, this is supposed to be a simple for loop. * (Optional) For better understanding before starting moving to GPU, you can simulate your GPU scan in this function first. */ void scan(int n, int *odata, const int *idata) { - timer().startCpuTimer(); - // TODO - timer().endCpuTimer(); + bool end = true; + try { + timer().startCpuTimer(); + } + catch (std::exception) { + end = false; + } + odata[0] = 0; + for (int i = 1; i < n; i++) { + odata[i] = idata[i - 1] + odata[i - 1]; + } + if (end) { + timer().endCpuTimer(); + } } /** @@ -30,9 +49,15 @@ namespace StreamCompaction { */ int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + int odataIdx = 0; + for (int i = 0; i < n; i++) { + if (idata[i] != 0) { + odata[odataIdx] = idata[i]; + odataIdx++; + } + } timer().endCpuTimer(); - return -1; + return odataIdx; } /** @@ -42,9 +67,24 @@ namespace StreamCompaction { */ int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + int *tmpArr = new int[n]; + int *scanRes = new int[n]; + for (int i = 0; i < n; i++) { + tmpArr[i] = idata[i] != 0 ? 1 : 0; + } + + scan(n, scanRes, tmpArr); + //printf("scanRes: "); + //printArr(n, scanRes); + + for (int i = 0; i < n; i++) { + if (tmpArr[i] == 1) { + odata[scanRes[i]] = idata[i]; + } + } + timer().endCpuTimer(); - return -1; + return scanRes[n - 1]; } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 36c5ef2..4b17aad 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -12,13 +12,108 @@ namespace StreamCompaction { return timer; } + __global__ void kernUpward(int n, int d, int *data) { + int k = (blockIdx.x * blockDim.x) + threadIdx.x; + if (k >= n) { + return; + } + + int power_d = 1 << d; + int power_d1 = 1 << (d + 1); + + if (k % power_d1 == 0) { + int idx = k + power_d - 1; + int idx_1 = k + power_d1 - 1; + data[idx_1] += data[idx]; + } + + } + + __global__ void kernDownward(int n, int d, int *data) { + int k = (blockIdx.x * blockDim.x) + threadIdx.x; + if (k >= n) { + return; + } + + int power_d = 1 << d; + int power_d1 = 1 << (d + 1); + + if (k % power_d1 == 0) { + int idx = k + power_d - 1; + int idx_1 = k + power_d1 - 1; + int t = data[idx]; + data[idx] = data[idx_1]; + data[idx_1] += t; + } + } + + void scanEfficient(int n, int *dev_data) { + int logn = ilog2ceil(n); + int length = 1 << ilog2ceil(n); + int blockSize = 128; + dim3 blocksPerGrid((length + blockSize - 1) / blockSize); + dim3 threadsPerBlock(blockSize); + + for (int i = 0; i < logn; i++) { + kernUpward << > > (length, i, dev_data); + checkCUDAError("kernUpwardSweep failed", __LINE__); + } + int zero = 0; + cudaMemcpy(dev_data + length - 1, &zero, sizeof(int), cudaMemcpyHostToDevice); + for (int i = logn - 1; i >= 0; i--) { + kernDownward << > > (length, i, dev_data); + checkCUDAError("kernDownwardSweep failed", __LINE__); + } + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO - timer().endGpuTimer(); + int length = 1 << ilog2ceil(n); + int size = length * sizeof(int); + int *idata_padded = new int[length]; + + // pad the array if the length is not a power of 2 + for (int i = 0; i < n; i++) { + idata_padded[i] = idata[i]; + } + for (int i = n; i < length; i++) { + idata_padded[i] = 0; + } + + // copy the padded data to the device + int *dev_data; + cudaMalloc((void**)&dev_data, size); + checkCUDAError("ERROR: cudaMalloc dev_dat", __LINE__); + cudaMemcpy(dev_data, idata_padded, size, cudaMemcpyHostToDevice); + checkCUDAError("ERROR: cudaMemcpy to device", __LINE__); + + // perform the scan + bool end = true; + try { + timer().startGpuTimer(); + } + catch (std::exception) { + end = false; + } + scanEfficient(n, dev_data); + if (end) { + timer().endGpuTimer(); + } + + // copy the results back to host and free data + cudaMemcpy(idata_padded, dev_data, size, cudaMemcpyDeviceToHost); + checkCUDAError("ERROR: cudaMemcpy to host", __LINE__); + cudaFree(dev_data); + + // copy to output data and free padded array + for (int i = 0; i < n; i++) { + odata[i] = idata_padded[i]; + } + + free(idata_padded); + } /** @@ -30,11 +125,75 @@ namespace StreamCompaction { * @param idata The array of elements to compact. * @returns The number of elements remaining after compaction. */ + __global__ void kernCreateMask(int n, int *dev_odata, int *dev_idata) { + int k = (blockIdx.x * blockDim.x) + threadIdx.x; + if (k >= n) { + return; + } + + dev_odata[k] = dev_idata[k] == 0 ? 0 : 1; + } + + __global__ void kernCompact(int n, int *tmp, int* scanRes, int *dev_odata, int *dev_idata) { + int k = (blockIdx.x * blockDim.x) + threadIdx.x; + if (k >= n) { + return; + } + if (tmp[k] == 1) { + dev_odata[scanRes[k]] = dev_idata[k]; + } + } + int compact(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO - timer().endGpuTimer(); - return -1; + int blockSize = 128; + dim3 blocksPerGrid((n + blockSize - 1) / blockSize); + dim3 threadsPerBlock(blockSize); + + int length = 1 << ilog2ceil(n); + int size = length * sizeof(int); + int *idata_padded = new int[length]; + + // pad the array if the length is not a power of 2 + for (int i = 0; i < n; i++) { + idata_padded[i] = idata[i]; + } + for (int i = n; i < length; i++) { + idata_padded[i] = 0; + } + + int *dev_odata, *dev_idata, *dev_tmp, *dev_scanRes; + int *scanRes = new int[length]; + + // allocate the buffers and copy the data + cudaMalloc((void**)&dev_odata, size); + checkCUDAError("ERROR: cudaMalloc of dev_odata"); + cudaMalloc((void**)&dev_idata, size); + checkCUDAError("ERROR: cudaMalloc of dev_idata"); + cudaMalloc((void**)&dev_tmp, size); + checkCUDAError("ERROR: cudaMalloc of tmp"); + cudaMalloc((void**)&dev_scanRes, size); + checkCUDAError("ERROR: cudaMalloc of scanRes"); + cudaMemcpy(dev_idata, idata_padded, size, cudaMemcpyHostToDevice); + checkCUDAError("ERROR: cudaMemcpy idata failed"); + + timer().startGpuTimer(); + kernCreateMask << > > (n, dev_tmp, dev_idata); + cudaMemcpy(dev_scanRes, dev_tmp, size, cudaMemcpyDeviceToDevice); + scanEfficient(n, dev_scanRes); + kernCompact << > > (n, dev_tmp, dev_scanRes, dev_odata, dev_idata); + timer().endGpuTimer(); + cudaMemcpy(odata, dev_odata, size, cudaMemcpyDeviceToHost); + cudaMemcpy(scanRes, dev_scanRes, size, cudaMemcpyDeviceToHost); + int result = scanRes[length - 1]; + checkCUDAError("ERROR: cudaMemcpy of output data"); + + cudaFree(dev_odata); + cudaFree(dev_idata); + cudaFree(dev_tmp); + cudaFree(dev_scanRes); + free(scanRes); + + return result; } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 9218f8e..549225e 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -5,21 +5,75 @@ namespace StreamCompaction { namespace Naive { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } // TODO: __global__ + __global__ void kernScanNaive(int n, int d, int *odata, int *idata) { + int k = (blockIdx.x * blockDim.x) + threadIdx.x; + if (k >= n) { + return; + } + + int power = 1 << (d - 1); + if (k >= power) { + odata[k] = idata[k - power] + idata[k]; + } + else { + odata[k] = idata[k]; + } + } + + __global__ void kernShiftRight(int n, int *odata, int *idata) { + int k = (blockIdx.x * blockDim.x) + threadIdx.x; + if (k >= n) { + return; + } + + odata[k] = (k == 0) ? 0 : idata[k - 1]; + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { + int blockSize = 128; + dim3 blocksPerGrid((n + blockSize - 1) / blockSize); + dim3 threadsPerBlock(blockSize); + + int *dev_odata, *dev_idata; + int size = n * sizeof(int); + + // allocate the buffers and copy the data + cudaMalloc((void**)&dev_odata, size); + checkCUDAError("ERROR: cudaMalloc of dev_odata", __LINE__); + cudaMalloc((void**)&dev_idata, size); + checkCUDAError("ERROR: cudaMalloc of dev_idata", __LINE__); + cudaMemcpy(dev_idata, idata, size, cudaMemcpyHostToDevice); + checkCUDAError("ERROR: cudaMemcpy idata failed", __LINE__); + timer().startGpuTimer(); - // TODO - timer().endGpuTimer(); + int its = ilog2ceil(n); + for (int i = 1; i <= its; i++) { + kernScanNaive << > > (n, i, dev_odata, dev_idata); + checkCUDAError("ERROR: naive scan", __LINE__); + + std::swap(dev_odata, dev_idata); + } + + // convert from inclusive to exclusive + kernShiftRight << > > (n, dev_odata, dev_idata); + timer().endGpuTimer(); + + cudaMemcpy(odata, dev_odata, size, cudaMemcpyDeviceToHost); + checkCUDAError("ERROR: cudaMemcpy of output data", __LINE__); + + cudaFree(dev_odata); + cudaFree(dev_idata); } } } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 36b732d..405d710 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -8,21 +8,27 @@ namespace StreamCompaction { namespace Thrust { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO use `thrust::exclusive_scan` - // example: for device_vectors dv_in and dv_out: - // thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + + thrust::host_vector hv(n); + thrust::copy(idata, idata + n, hv.begin()); + thrust::device_vector i_dv = hv; + thrust::device_vector o_dv(n); + + timer().startGpuTimer(); + thrust::exclusive_scan(i_dv.begin(), i_dv.end(), o_dv.begin()); timer().endGpuTimer(); + + thrust::copy(o_dv.begin(), o_dv.end(), odata); } } }