From 6947ec68f40d9e39922cc9fe64255f460c03ad34 Mon Sep 17 00:00:00 2001 From: Alexander Chan Date: Tue, 18 Sep 2018 21:38:37 -0400 Subject: [PATCH] Submission --- README.md | 85 ++++++++++- img/scan.png | Bin 0 -> 33307 bytes img/stream-compact.png | Bin 0 -> 29217 bytes src/main.cpp | 252 +++++++++++++++---------------- src/testing_helpers.hpp | 6 +- stream_compaction/CMakeLists.txt | 2 +- stream_compaction/cpu.cu | 56 +++++-- stream_compaction/efficient.cu | 152 +++++++++++++++++-- stream_compaction/naive.cu | 60 ++++++-- stream_compaction/thrust.cu | 24 ++- 10 files changed, 458 insertions(+), 179 deletions(-) create mode 100644 img/scan.png create mode 100644 img/stream-compact.png diff --git a/README.md b/README.md index 0e38ddb..a75554b 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,85 @@ 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) +* Alexander Chan +* Tested on: Windows 10 Version 1803, i7-5820k @ 3.70 GHz 16GB, GTX 1080 @ 1620 MHz 8GB (Personal Computer) -### (TODO: Your README) +### README -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +* Implemented CPU scan +* Implemented naive and work efficient scan +* Implemented stream compaction using work efficient scan +## Performance Analysis +Here are metrics for scan, including CPU, naive, work efficient, and thrust. + +![](img/scan.png) + +The horizontal axis indicates the array size, doubling with every tick. Thus, to visualize a better relationship, the vertical axis is a log scale. + +As we can see, the CPU scan is almost a perfect linear relationship between array size and time. This makes sense as there are O(N) operations in scan. For smaller array sizes, +the CPU scan is faster than all GPU implementations. This is likely due to constant overhead of kernel invocations, and the fact that the GPU implementation only uses slow global memory, while the CPU was able to take advantage of its cache, which excelled in the sequential lookups and writes of small elements in the scan algorithm. The work efficient implementations are slower than the naive implementations, probably because there are twice as many kernel invocations, in addition to using more global memory. Thrust's implementation is slower, but constant. This probably means that Thrust is doing other work in addition to performing the scan. + +Here are metrics for stream compaction. Once again, the horizontal axis indicates the array size, doubling with every tick, and the vertical axis is a log scale. + +![](img/stream-compact.png) + +## Output + +``` +**************** +** SCAN TESTS ** +**************** + [ 38 4 37 27 41 44 46 5 2 8 23 9 12 ... 13 0 ] +cpu scan, power-of-two +0.032968 + [ 0 38 42 79 106 147 191 237 242 244 252 275 284 ... 399406 399419 ] +cpu scan, non-power-of-two +0.03359 + [ 0 38 42 79 106 147 191 237 242 244 252 275 284 ... 399313 399349 ] + passed +naive scan, power-of-two +0.041984 + passed +naive scan, non-power-of-two +0.04096 + passed +work-efficient scan, power-of-two +0.136192 + passed +work-efficient scan, non-power-of-two +0.136192 + a[8384] = 203041, b[8384] = 602399 + FAIL VALUE +thrust scan, power-of-two +4.56704 + passed +thrust scan, non-power-of-two +0.048096 + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 2 0 1 0 0 1 0 2 1 2 3 1 1 ... 1 0 ] +cpu compact without scan, power-of-two +0.039811 + [ 2 1 1 2 1 2 3 1 1 1 2 1 3 ... 2 1 ] + passed +cpu compact without scan, non-power-of-two +0.038878 + [ 2 1 1 2 1 2 3 1 1 1 2 1 3 ... 1 2 ] + passed +cpu compact with scan +4.0355 + [ 2 1 1 2 1 2 3 1 1 1 2 1 3 ... 2 1 ] + passed +work-efficient compact, power-of-two +3.71302 + [ 2 1 1 2 1 2 3 1 1 1 2 1 3 ... 2 1 ] + passed +work-efficient compact, non-power-of-two +4.68685 + [ 2 1 1 2 1 2 3 1 1 1 2 1 3 ... 1 2 ] + passed +``` \ No newline at end of file diff --git a/img/scan.png b/img/scan.png new file mode 100644 index 0000000000000000000000000000000000000000..46daa0390d69612ebd6e5e68be9e3b9e55292078 GIT binary patch literal 33307 zcmd42RaBPW*DeeQ(%l^*ARW>j3X+0KNH<6~NOyNPh=7E2Bi&un-JR087x??fyZ8PM zzO(ItgTZ)s?t85@*PK_(8TejS92tQC0RjR7`JIHA0t5t983Y8R4IC8sjf}XfDg=ZF z#5*zJkB&2Y4X~DKlZh`2n+)!+8dQr24!@Z2r=37Mca<-cI+zaq=p+eUAO9SQ^(qqU zjbx0`*f)%qwDQoA&)_j20=prgC<9X1grK3s1_Qf2jVLKLryzJZuLTDX`33Bo%sm2- zmw3+*&W^4doDZFr(`^2HAb|1|B141o@plX|hkW`Gx(o*T>6hI&XmC$IsXl%J{ABw7 z+6PM9-qT>&V!}PWvM_FKZLNWU0WvZ&ufMgMn*hz^`1r^eHQM7ko@aif&5R@=AyHIR zoS2wMOG|Tda^k8{*#`eSzxw+6MlxtIb8>c#soZ_a&dzph9~c-o9mIswLxh9Y?$%>v zVq${(UQl2nEiG+oN?!r}P*-0sMTt=cGu!jHOJK`dF}nNveR!xSDe39yOI}QA!)l{< z<3uoktpi`e3&&|b=fQ)5f^ve|ia#|o^I28ZMngFB{al0d#Y#tTTwGjOSlIFLv5k$* z+|}yp>d45*=1_W9F&8yE`zP91>!!DF-jE}tL`Ow=nV|2*(kac%&c+)fQczGF9S#qp zwOUzP3SrHyV`F0<92{`1?Cqt*#dRLabx{!GEzXgO!!`Poxs8BQy>Ujw6KW>1knM;Y)0D?{X|@bw>3%=rc)aX(VLiH^LLGtvc4$ z)>l;!-+ZT{qoeQc?v9x%&{|fQ++-vqJV;ix6cnPuzE982zF4n8Fcz@cKpGTZ@*)WQ zEwW&_Ht|yxoM;HlD1CH~Vsyw$Z+17EFJHbaaYvjW;xcN%%K514-=-1PNv5+p$4nNf z_jPv*Nun{)|88z>W+f$xg5>mwA;WiMvD~0iWJP~k9lMZ)1rzI6y^Dytx`GuOb}K&-9=|nq1a*4-YqG@dyY;(Z9bmgJfKm z6?$K|P`$6^Dcg7fNuR)RRIXXnW7^+|6m9SGT!6t+B} z&%3we^~8iDLgjs)H8fRVkI-9*5=bFBFb(BjZnd=$}AvIv^Gn;3G;uJBV)RBXV<@_3~Kn_8EiEOZU$=g!sVT6Gf6>ma*h=@RiWta|pi9%L~!uDrP<##?GOy#Yt zto-@&XSKy5hm2}OQc{pN8Z&XI9#LvaO4(U+sj@I@9OhoHX-m}O1IciLOWpr$$fPER zYg*kN%-8oxt2#J12nl)k`oioQqvPY}eYg`Md%XWuNZ%EVKYYM!5>67Ph|GWsP$SWo z%vQ)7gVlgFrGM5m@5%lnIfKE&Xqn6;YSV-=sCLT1-k!s{STQ1G9lD-XwPcN7a}9zE$tK%VBzP{cz zqFmkD-0VWG_eTy23EA7*Q>3=D&F$6hab|RkMbYO-p%f5EX$Et+cJ=GaBvEd(_T%l_r`qTpI^tDEIxdWPENufk1%CTEG+!_pkZLl>KKf^d}%d~;DOtWDkoYI&@(jj z7T^-ufqM3g8?|)hZ{`>TM+xeKrE}&=->*$CrR@)jc^|pDM>n>%0@}BRtd020Qd3j! z@9&p<1D3E_Vme@OI$HEWF{ZrX`s;bjaPjKL=^grT&|A#<93ES}O?C#Fa z7kzzp0f9l{UoKbHQc{Q@4jyrgJ&um9E-064)>o~qkHG= ztVI3zK@%&PO-dhDQCUf@M+A?A6eF3<KZ(~60dsfz$ds|x}rp%1FNH)Vkq0x|8D2c|{SnM@4s^<44lG#xv3F6Z{e0(YI z1i-hx^cmz<`-*lyUZ6BBT={u_X^EMSj~-;fE2CtI^#zkY4GnRR)v zzTSZv=gxWRtv-r_7dZLeLl03Jro>b%t{dY6Mg2f`Q6~ro+UYBS6Ff~ zq3vWx2h0Q{JP9%J+3BgCeDbje#8rz(?ey{9o@t8~w`y}Le9z_u12a}}rfn2MhMHeD zhKi1)-emCJmus91yaDW8`4oLuM6L#(?$Xh#rKP2-Kby$YP$0_Q)4WEdTeg_1iHwiW zB{h|k3(v}0ZEl8;^;0hD$FbZNz62m>y;l_qRty)@SBsv;ebLG0I zD<&ppZIufs;tLB4+uOz|S-Oap+z7n@j{zWg$J1ZiOyvzFUi z7Y7GU&8D-lwrhlsN=mC*1_o3JyM+TJEGuh4K?r;CF)?HVg|=bLpa!AAG5UNo)Ypff zff8*&n1LcCP5<_dBA_AC5gA6J%EKC;LP{_2RV2H78Aw*cx**dBaIRr~H8wQF#m6@` zF+sfA{CY2nYjxY|iwLsrbtoJ+H}|l*iODWFs);rsYUOTn#){HXD+dQT2?;HR1Fx!# ztgN=azF-SP=&`ZMNq={AfwjclN{BOGzP5TTcO9M9)>g0GgXTIcp%Sa{=4Ou!aA1X? zE>C$|xVbeRJHkApR23p6`6rH%%luBQ+umzv3?|`8Q}Xk|>LLE!(75iT6o!ECaB=u# zQ^^19bKGNM{nuYdF@O}Cfg1id4giJd#u@e(*9)X)d07S!wI1SD+lw->xC9IpkbTQw zPM6VDD5RtV!GZnPp9Ru;nzr!h>giGP^3L}46*M%Y@=lG9uOA$k>*}staJv=Ggq zB(pK9geW4n+#LaR)Jh^#0f<3)aRw^k8LbQ2WB7yjFTpfx0mvJK`}BQxcR)h}|38lz z6W`64VUWQ|vr5+s6spo{>cBeG5%Xij4HTnk1qtKk<|epfB{P#87dJj30Z&zRmQJZCBrNO;2`MBBvUk)r zw1S4l0VvH<-@bkONE_=`!sVqvB+?n#+$@y$G3Z4=-}E#$1H-_=0tq!L3d#=k&@eX* z?AOPl67wuKCnpX`CL&&o3(!Rv7?>&H6<4ZF>~ zo(c{LLHZ86Zn2sPe}aJo7Dz=y1J|Ueq@?=!^9sh^s)V|Fsx%d^^Eth}rJJQCGoyB7 zEB0|$BmDEe6k%=cOcddws%D8_eo)}9KE|`EP`Djkrpd>BC-LW`%}6cFbblhVg92-P zD___EXS^fSUfg7+?*P|QXe0W zE2w<0K8|YtPF0IiLlp9F1pxnd)xl&an4BcVYWhed5;hxp~gz-ax@2_WBC*Ybg#A|vxK z@%H~hMMY(8b_7(GhK2?Jwz$N*aF(pGEWslU4GoV?%*~73N+=@XIy6(*1O&pE6Hnj~ z06qittlCPodFB$po=70?NMn_$EJCeT2|ey-2)b<1j~~r$@7}@Jl#g|7&c zih~8x94%<~-<=T>AU2+BArfKghc@Ihoa*uN@@8fL-3ra3o5Xf-aFCIeb#!6^=Va(? z%0F%S(iagIK*&==;sK}8*48fc_mnM(U{Nyyh2Y>IEh3`%b1#hVucrdoSHm#O zw+9obN9NK^yIWcEodInvn-rx$)`PF@nfAZxAVbl};VfviD-kQs=u&wgip8R;aEG6r#k}YJ`)mi14pvkJ?r#q>* z%jxGf&BpE~LYi_jG&dP&BhRB#jnfZ@NV*)X#Pu{O;?lsL`*NNoA-UU(7C5|=jU~+?=fQb1BtDJD|bHK!Xf@r6|9T+ zyzM=FCWD5R*HSnK$Y%W%WD#X!7rScpV&r$sV|(??>s|S1bXIb4g|f}XT(E{aAE{46 z3L47Sit7b;Wbq=HNU@oQAJ<#*^Sz342kZN_G@9v>o%eCZ1c%Q%K8vcz(X3T<+ecoZ ztjS2f4&WrSd7Y>B+$Hz(<`ln=P3upg&Y)SiE6+R~jMl}7KV?uf)ysN{%opFijT91JBWzN%>QB`qEG2=b+4L zRZLqpxA?A=F3jv?JkuZg*Eqydlv$nTF!j!}d&C<`)Nm8N&D_hmyyvKDKX2q*ruwS> z7+xQAhM+3#G&1bXPA)3)#YF+4XX@Fr&RX^gh80Dzd}d`*$+huWEnA;geE$~}Hrf*` zBTDP3qW0yM)}EdXm7Wrp0=>NM!HB#}tbiD{Ssx^I28G7Yt-|eP}`Jb zZHTQ}*vhtL50**rn+-%vk2yVVyK-h^uWUoN+C)vAjrdsk#$^PWBYd??=xr4Nm6ZPc zcW4Asf^*k5DCdv;;(5eb??pN?a_I`bl)1`h<;F1)!b1QT81SJ!4zfPgm3fz7kk(!`-G~o8z zb16$Ts^4#1R!|GBld$%yH$tConM54G>-LgyL0-i1=ovq#mCel&Nzv%($_<>W^hxtp z*B%0qy7o75{k2N2;fqAP3OuKRTxkL3pT^EIRC8rnANks)>Hm(7m($Qq-&j}ufs=Yh z^VTl-DxOjbVE>rxB4lL;+p5x1!@??+TcI+{{G`V}Ue793(kG-_N98glT}Wx^IaxRm zzF{*RUo8@?CoHOqrU&lB`?yDx%`#6`)r_sRW@bbfgU@f3dUCc*c*=sN|3 zo2LfSh{M4gH_%(jkN!XOqDPzqXj1aqsv=c=>^q|X&krX#$ShhBRj{uPokf8{L!t8!VO@iUgVEXfO&@#$tnyD9E!>`~~Wwi#{#+ z7hy)tN$LUuS%nI2yv5&am0ptoFasKvo1G18jI7_hTwL8@fRokM)F`W}!ajcv)L}(c zz@Gh5kB*Lj17rUga4IV+7JGpE`S|!u&CFg%FlfBl7mY>g#96#xOIluC8QcWR@jtZEb~xA?*PZ z0(c3i1f}T!tm5M4raM9hL;nx107=@M>W4YiGB>BTdb|S*82UfwhS1X+n*=b6c|eOW z(>|KF%C-Rk)PA?<1)}MicGtn?RCgP)V_8|9u1%gC{$h-)u%9U;bm0^7r=+A5+DP;h zE~u-kqr?a-F1Aoq#971v0?rbdJwVr~sq?3&6;EK_R9&4Ii=T^&>%|LiMA9Y^bMmR0 zx!KwJ!a^iou2H~3AZ2A`--dXNz*9*J9uLm|e@!jG22B7(hg25o0y9upl5C2m>yfU& z7#|AR*;Qs|zhrC?2qPP~4m6Z+8z7@$WNgs_Qo3}Mo0;c78pJ;P+PBUiFGCb&WqtMPAMMmusU>FTpSKIwy&lYgTyv<8Yw+&i&lKazh#`ivD^2O_w`R=oo1QOE)+rv z^a^<#<%lh#FHox^7>vx!&W?_>Zh+$lz#~m?8-kCo@39#q8U}{fiTRf=v<2i{JqCPR zM*06k{aV$N$dhF~eGDVfKv*Ek1CfAiv|pKdNklD-n3I!pt0*`avCum@dJG5zxw*Mo zEB~A36-d}|sH#PyA-WB?*c0Muw>2)?Z~KzFCN7Fp_$}LRK3_&I(|&E4bYbl+iKHq4 z=Xy_3nT2QL_4(;37Y`2#64H;4v_Ik-89b;xh(wNnwmIDd=s>IShde9<-wo}Tx1ofb zCnqO09B_(~Z=F6f*HKr(ps6dWdK?w2*LV6_SXeABE{5N&t*wD;fJd;mSZ9a$kn{aJ z9F;~%6%ZLaO%Ox=gUvz83yu8gmcc`OY!uMa5hFvQ)6A^-sgu-+OG)F`oAknr>u)aM z)(n*HE-oPYlHU<5;&iWDcjGCQ3JMAyE-k%RRi%sW$*?c9fe^K^VMFMJ`vQy{GBPqw z&iKl_RW7KW?(SycmMv&ZzW`t$Uj-VvA}&40z{CWsVnQjSu8shfjg4(tT{)Rwr@aSw zxl>H9UX@l<=w;j7wIlxDm9oE4n)m+Iho8rC=&rUVrQ*xG2Y&!flRw^II(vm#$w%OH zc_3hAu;V3JY4^AYk#s;9Kq{M?pU1?&P{!n1#DQpOX&H7r(43$}9q?8Xs!b<-18BYN z<72TvyB;T65SYZ_>8I4vK1@e12)pJ)l(SLL_)e>a+$vYzqFVVCQ+jf$* z@csF@NqD4Zb(cnK+8Ku@fBy=2i&;K0RUW30u8!bI=I%K^y?2V+zN;Y%yJUYW%l|s| zA)D91}gkI1U&==%%Q?H-JOS58-`3e*BK-E1{duacFHcW@l>dtQ?4v1sFG8+3)^ zTEqAKy>CgLp!UB4TxBcWrJ`)40C&?i=hd88*~(r>*IN>fiIM-VkQxaIsx?usHjz9X z;)ujKQXjumNufEEZ|`ae?mnAZ>OaM$ftgPIXh0(+FJeG+STJwFwk&@;ujfa^fx`uh z=+RPYg9-bIrAleI!EdFU|J3|F6dl?~lSx*`c~>ijXS2$AnP|poodti!;a4u^-2V<8bm2ldBI(`|e8Ov9Y45`6bUD`~<5qbr~p(rG!c~9iep#F8a4|)jE!rq85G1Rg@38rsywJ24)e<$ajF# z^iRr|O)+k9T_4#Ew!w6@m8zdvrivQ3r>18jDY=cRg*0_dC&e{1G*xzP(P91WdIb51 zIGRcKno<7^EgW#h%Fe~haoh0qIB7W}3uPTz!{&!VYcwMQKXY8?(9>WOGycI}9IqnK zbHTmg@@gxW)K?o4lGMlkua$7-*v-@4>-i=kOU8a3>qfAp4nfxZ;}jgPFJt=|`TI;_ zq+d2`dKD}z0Gyi_!ePX}LEN6b){J1Dy>yIKIqt5sVoju; zOeQb+5IPOEy2hWx#j&qmr-w8%Lj@nBMoR-fjED9k@AWcTBUT)gT=hKr^N>Vf>-Ar6 z#4#p+4CQ{Vb4Y(hP~`E;rHpDcH3D^9Eb{x(W?_`jSBrGhBDze6BV}V`HfUJZgX^1= zgxAh>z32YE__#4&#^W6QyE;21czgbZFtMB4=b)q^I8`!Llce3xSCGj}Xy@UtNY;-P zy-Qo_?}_;fpFvjX_OQTf_-`~_Kd=(Hr6{gF6D|3)undeXo6W&jv69Hz#QzP@8dZ!z z5Io{nb$Em9a2Z7-z4H~jmnJ;(@J?Yud3(p}pOHT&&ogmUS`)*`c3j8s@N-^_YeIA1 z*@*Qp=^HIJ(?IBb8qSH&7+~CL^ja|f`61^ka?f*4jTHzF_ zfRcvx?zX>yw&aN@zH z0=rThRy-*wi7ch3rzb$3nWx2vf%+H@dM5lPu$}BB+UJ<2P=7zo%wu49*5EWzr24hy z<>kR6AhZ^bkBCXC{Dr%2pME zMfV2Sp@kLsKqSU7#|i_W4KyzGBB=Z4+u7QxHGGZ5hpB1{YSqxtP`iw90WMZIwURG{ zzJ3vuM+^npVNfgJInryKtm>tpTw^KOg->%`H@>d+0kVD?xA#7J(!={7g!Y` zh}{lO^5-4tUDNmbhyukbF|epoITJTMMX~=oOcecZn5bj*2ovT1!oooA!x3;*LnKKmfGCo~bL8L*j zKKD#XNzr?_zPf_<3rE93Fw;Z`6tiCK4E^--2~Bp7y`=>NiJJa1%iKjr)Q)0RzNmW*u0HDKx`*7=ljtTBpgHp^!V@dRQv@85e&?=t##)+*jakNx;=Q8&}^JG|`mbw~!Al-eo_`ikHnJjgOA;{_>+?VltFL+T9GbX9Z26 zZovPfg^2z9=S@J&8H?JVK*Y0_I`k5iz#|8W_HaTSErjq&D@jw$#7a6ttNzSS-^B-N z4S#Fis+8S@yF2g4va>+zA;yIR5gV)WQN&i3dTh-J$tcUe?o$;#g0`rUU1Z4DJlrBM zSHs+Q5B!8tk>IKk9c@N?vAodWG1|N8bhT(t)_s9g;;40}yFK|Z{ct;UAAh*Z4BZsW z@n=X*txe(y5B{j~Ko}OAs7Uy_5Lx;O&AD^3TGUgP6UtaBE0AwOC7TtX1f$~>B43O7 zcc^Cq0>JwS+NsY0`un@<@W0Dfj#U@D=HWPwOJ#epN9OP5qPnK~vlx}s92Ngu zeGp|7ygBW3;~*Zs#(!*SDuV%;=2c*!G0T97Vp(lyX*tO<(~pmA+dXrp&@bW#Ux|xz zT$JYH;7E2+neHCi&=HZmlR~MxA1S#)eXYqizIj$4=zAZC@1@z{rf(F~ynl?M+2~T6 zo&5&zlagoL(F#Bc%MUNJAD0^ZEQh9cgKGURnu(| ze-~*LznFNx;ha>Qrobck3G+- z?QreIRnJi~Hoaqwq8?#xyE((Y#hMZ`=h0img>_W}hI7F&m;0ERjaP|7AFf_KRPkP! z45D4W&X)QVl==t72K1>YB6Cy_Jsvlc==S#2k~b=lz6b~juMT0)XJ=vNiB2?#6%FF5J-r72> zz#G8bbDZAEut4YNB8C(am>O|wK}~y!*I2DSSH0*#1I@0MD1op5EK+sbrpECpTXA}^ zQx<$y{(+ZnYh>j5JD!<;d~vhgV;z+}U9fWL9F_Iq!GZBAFKBjgk??|g5RGfHU8Zl- zVS~3Hh#*GeTA~F=;1a+IWxl7RrZyodBE`nX$45sekM5CllZhMzQV6Fct@$K_!uUkN5ix#m-BH~k`hsM^MZiY|lk``2E1 zLM+iVU9KSRr3kv;pKgx<%05MYOjaf~@j9{SDhk$QG*pXS? zT0?34K(F}o=TE)EVVW}>(cdRQC$?g8wVUK}3T&|pzHXQ?bw#d6FAOp53%Rs5p>ATz zH1vSkmHJ_Im7Q}m)9M&nxM7Cpm9utu&$4sTkLDt{_2Z4xQHfsAg7))(hRRA{05^7S z<+phKedyB9>V=)Xz5c;L0lOWnd(cWW*W@nvucJba=ZpCLu}GWKi<*M>=ljkxB$tmG zU7t`>g6gdJ*G$Ru{$j?$ zS8VEcQLu=#9am>Ny*SHG=AnKP?z?B2(OX{EQVmJ5ah`XV(d4XM zFw`kLoMe(a2A|Q*iTDBOM!SC6OHT0acS{R7pg`WR*GG@Q9VVv$To~c(q=CMbYi#?6 z<$>wexiczdb{2hA_w)jiOIF6EuvrZ=&DOESaWVIABPuhF!hkmt_g;G1{bu44wWOQ{ zBrbVD8%lf3w56c9kv3A|sqt#!cS4b4&oev#(`a2F*r3(@?d4vI{WRf(3Y3hU<7}7t z_=~!rnXUO-+XnT_jwjua39VO!D?F+H7kX!|LtMiww|h=V*s2P0KwziLtC~kg%kuMs z1o(i{ArrLS8^`AX6lKc~=#N@OVF9_;J`4~JaH(l!rNSCd((FHp?luwBIZ3n46uAd} zEpq*C%Cv|^)aY!yO8`oj&%Oa7phFsNSL>((mrIx=Pp@o`nOLON^X#n9k(Dok4Z4Ap z6x+~{{>OxszIN)l*N0MpH;h9ju()dOpFR-g%F~_oFw*205XfqTT;UtSX-L`75SZqG z7-1$?wLIODS5le;pi1(euP#pz*p7nD4?GS;YZ`QAVkGdA3m=HZDkOzzIQP2kIEUJ8 z8l9_J@*~)ZQhy&*U!)dx(;)|S!9d{kF?nD6pBR~{cVIeWc+5m%2?8Tgu*}f9ll;@x z&rM@IzZZSI!)Th)INDe15Lbw<*oQHAkg?GX&mj*a{PIK2-kfV+a?~b;-0|hKDeq(T ze%u033;+`RyY{g}M)#g)9YTxU+2CV%6@o9!jMbX-5UX*g?oz4dh1tdpeF|Ob79Ivk z2GMTC_W~=@XutZyLq}J^STx<8BQpIu+E{&R$D0l{tV?-`3W8PJ(3#@{`8|qyiAs|b z9L_AFB#si#g3X1tv*)~A<@ht1yzB92F7z~p#V?!Af|nYC{m8E6jOW(@)`3%P#;HK` zd^&H=gzYFPl5Q~u?{h=tGY|WmR;f)OZtlhmf)0qe(XJW9!Mo_=9Z&IC{F80D)h2f6 z$ip1_^~W{=Q~ZA#29pPWUq-=3)$QC{$kG#d0rg7#Laiw;xN;w6+%CKlSItEj+fss8 z{+M+aZ{ubm)D78h3D(oXm(z!C{Ar#4bP(37<-CpKOQ3tTb#u$b2Az7Nkpg$BPA!^Y z#Z^;N%3uK8vCo}_x_gPrVZE@De{40o;sy4wsOe{K-4B%!f;XtoZSL7G?Hs)fMHctc z4sskEZ>GJbrl!7sr>s?LEGcPT)NG^N7g8y}$QT(>c<1Nmfi4aZ7GP}j zQvTSMwzeSw0X4SU7>0NXfcjW4^uTbY+jwh4+%fAIEHC8UJ02@nkL=Mi2(qt74jj-3NGf#?AAg>)oYh+aw z$NT2LorS##JK~uZA(o%O*cDXM-+p)ABKO6RMMGEnIhhUdTvGZ6*`j-XH8xdCFz3R= zR9s{MX|Cl(2<`WTIg2nq+(*F%B_bzPJ8GE^=2Gr9&;JtNhk?|T+f>_NviSvtRXul^ z+|0QQuMxzPSm)i7@_Q;sIdL(I=LV5SkWtemjbkd zAM%gQ=HKt&znL2wtEd*(p>ioGDZ!TdQ>?raqGDrZCF7$B7bztZLHm5$YYK@XqGpFP z7~7+F3GcZy1I9Vd>0TF>ltiR9@7v6T6aRP%zRi{ebO*nF9dq+e@O50sQ>S}mM1)5) z^7vDmqB0zZ$e;6bMmoB&a5x3zpT2IMF@NpIXqOfmTT{r_B%H{&D?~m+(oy%WY}V<> zrMjF~k;J$Nikn4AK@5(61~?OEbrKR2ThNs=!7x#TBM{c=Nu+CEqS9L+g}_*}|LXWO z|Fg;gle7}IeV!($O7Z8nuO6ZXcd?|SHe0=>DTv|7Y ziHYRoNwG`a;@05G)0->ZhQ4RA~k}A2KO=7zX zh>?6b$sEqgZ|=KsOoicw#q-!pDEufzB{4WIhRhkR?p;2Y6;#Hogb@!T(zbhn1Z*5G zR&MUA1?R)arxkBVPe;Mt?0?M?v{5g13dSDwgk^uT-owczBA5H7;n?iVqpE>yM@M{N z;FBV%rNoaCXCBbT-(vWRbPH8%E+{-U(PO!#X0Qq|)b3CT>%l(*{zu;2?98o<=FmW3 zDN$a9MWvpSHX;PDAK^lP{YW{`(*qA=_1IVG<9+LM>^=$~nvV1x!I(`C@il5P(tlbj7p~DuX&zfEPsJ~R zG{6%z6y6%a0P)uWg+5Gm4jW3$C$oUV1v|6F4;q?#=(AdKeaQv3Jl^X8#9=?t9YxOpkME_O^zw)j^e&01(=( zv7edpCPks8ZgJ16RQahDt-B9U(UtkNy!I0$uE@0zI@1SytcGygTI}fU*{gi zcKv4pxpR9>Qql&x zNIvBo2S;pIM32ALXv}4)`|bq8ZE{ppYwyID;%oZz>#G!k*~E>BO$l#q`< z`*DQ!RLne=9&b))`KnSg+;fwlN|-`|{EUpH^Zi*sW(%odN~-Rp-aRU^`#)>b&+m;# zqzaO&l{pHw^F*etIC{#c_WA;Z*nraaBJUlKz|zlt>Dk{U0sRp{y)oO3&sQp0p7!v<<0>#&Esc3jxQ+omjfrZ+A0q(~du57&aCg=1WnB;Jr29K!bvX^#{EN&`|ukKhJ=eX2L73vu5oC=FWtKATFZa zjQn2Dc6Gt}TKu3&Cq~%%J5B+|_Bc*K0mdm-cpt|pF3!&ldKvIw6h`Og2PQ2hsv)Xr zSXd&@7|?!MygewlvJV$fd)AWlgF$G8gwOH4ipuoFgn^g2M@&?=x!tKAZY;@7!Y9;b z99!-*^@Cp>i%AeNE5?C~_irzhFn*zM`=;>55oN7DL(yALG|Zi&uT`B<Xvy(Uf%ldZV6RYRmu>;z>KGI_hRo74QT?)tOX>GKuqfM!^MYXf6{~5 z>chK>MMHMNztwKSvTuFA^wlT-1#Q2K`*B&w#agWMhxi?7d> zfe<)x4C!ngB!|6m^?RWYpXeZUF{5%Mwzu|T1@*yoIQg)XFa1+R0oP$VS0F>r!{SA9 zBLTkTTF_;mQz7d#Sv%eYBreJ+jA+^$h7FqDHt3I753sFjVHO^Bt7}QU4Ir#rd77P? z-)<#Cn3lYN`&H1UZ_A~xIQL5jv!qRM>Syp%BW;F7NEPHnKTh)ifC}X;8JC23c$ec! zntu~D!H-OC0b%6gLmj`2>q_-kxFXj#Ii$MRHSohz4-FfCa!HgQE)TNjCvRm_LlCEI zIMWZG*{{c>(G=CVG~Eq0OLXm}d}*UQ_Y*%!v3Tedyfbvkd}B;V8|+oz*j_ZQ8!LHZ z$(u_R2*8*9;Wr?R$eq(BPa6(Q!r*~Ig zJGmVoq?YTr{uP;+_2K1_ywiNg#<|aeaHOMRk)j13aklmIBm{iM!G}%p>_2l4@U_Hy zsNyP%SR{0(Boo>Obq3bIASSf+Xyo}mSp2Z^T+nW{??{^V0_@7ej#?Y^ee+qZKnKFM zd5G;eH9b?cb#Ncq;tDdti2YI@^-{G&vYSUN6lcPWiGTKNr|GoeUJkh{&7%G&?PNnZ zw2z5Bx2Z4mH6ibxMXA&@#Jc;)Pcf68^>qSEM(-+hvuKJ8GYqW1wrOy<_7B;e+G`dU z3-lKUg%4bIA@&UYu3l6&D1Bn6QZrrve!JeBr}8JjZ+2wEfzyaRf$EJ}|7o-OarDI?<%OMokKHk^g4q}u!cx1?p{if@I_U1^i zK=6jmsA(rykdWg+_-@+0r}k3vH1?wCDzz-M(S?HjK+B-&Ai}d%(%)|)0QL@hHF^wT zok{>EfM0SJIDh?HTf+SrI|A)>rMzN+AFZVB(*JztC6CGy_+tey}Z0nL)RA58h^SlheNi& zh3>P^(D$@L1#UKW`MF7(KW<}67rmMjhasIg)5{LunHVmyW&O}rJ`=+-?_G(k`Yx{3 z|H5;)B2G`%1VZm6>%^hu=OFZb!vRRgYgrjHN4X_>ZKxxtHKjx!M{sPpRDE?kYPROG zX5)$ zmws9bMULOl;tsJFXbW%4@{e_FoZG$cZ?(NOv{$VYYgvBYU+s^jTn}a?Dh|2GtzL&v z9hu2&ILX+>*!!_h7rqiyNi?<8q0ilH83!ohjJlAa|B4iVi1)YqhiB<&`D;8`dE;ey zjkEi)q>3B2J0n*rf_2f?#lnFW)9r}TnBQ++W3nyCc0c5QKYOd{AaY37uL(LFB#4_g zXtPCFSorWJM;)UCk5CtW)GkAZ?Mm||D7$;hn7Z&SM<~2x%R`NO)cITMro4>XrnpH- zdM#;2tt#s?)m{6ThKmc7lrMDFKc?p#C2h+gxZ&1TC|pzsXoM>-)_mh zcG5`7`8aTNC*x9v(c@hjY0F{7ui7f>K5>1nSfA-hb1dEPnd2vcF==D^q5{q*yNm02 z%E_$+M9#J$ZeQ-ypfm0yjBxQIS^E^TFkk=|{n{h7S)Lno>*T+A@rQ5m19q%o-+?*; z@;dcydVjJ`$I_i2l-No%I(7=Yg$>Ygf9i>+aXI`^T|o{m_+r5s&0`^P~{bZ zU#gan1bMq~j&NVNjlPM4*Fz-n_V*Xp%_>ct0_AB^55_3^?nnzx_Tz7x$~zMixDaRG z&R55&p6Cx@BwcUOM~~jR`DV6m70s1py&Ny(xwe(T9CR=-Zlc-|`I1$CbBGeAQ;?q0 zGQrD)AJ-dqFw8tt)MGv3cz+%_)dujGlm&m43(7fZjIsaka(mm{bV0@V=BD#oxk`5A zz+Co@Mmlp3<&&N=Sua!UdVm0~VjqG9$R6C46iH}{uX&2+US$tczh?W*#y(8WePpY! zH1b>Q&eQ#RHga!{>CYcxx?YRdlR2;b*B5zJVjhaacL)#ef6QHw?hf47t~VEpE|k|> zW|tPJf_*B_%`QmT6otcq{*y0--V`jGBDlVsWDAzJ_2<*SjgZA_UWB? z4-Q`uN>*07>-jX8e4(YIJN|ZrpdB?jszCH``)3IBrBp5Y=YCf9i21fOi0|dI&?1A6 zKG^kLEprw7d^3(QPtT9z5{WBNH=ErOV~UZZx_-XksELq-Kqqy{ZVQyWD@LY?^amN& zE!?(P`8KxbzEnKJ70;ZPqN-hJ&I&enpUl@94$jw*(XQnbvy+k&i#$C7KldMg+N1rQ zTt-EA6E8pQ+)yr1+WYI$K(^P~yl;fpnEG&zW2kOFE`CKbj#ynOZEWnB_{%A8u|Jk! zt(pA=>}gF5h&TRS^b#Xdh4Cv08{Mj?$6DCb5>Z+m@P!ZsT%JnGgYWOz5;aQ#L+^t6 ztBVv4q*af;2h&G~MKKIBc{1o37`PpG;4(Gco)B96vRdno1Z+g>0ce`Ny}dJTQZHFV zAxiZqJc&OrLwEL244^TXG9#mv+%8+a7b+92zt$p~`#e+1vJXqF4EX{ZRm$7~Td^z& z-jte-P`TuGATKYp3F3vzxj3c~=>Kp2ut9RYzk>DNDh-XE@7lJKYw#)_Unl1@-_ z{oZ{3ds1)0fx}C-FC?K8+< zr1|#6cY^l3=eGZ=wXY6~x{KBY1W|@kU=Wa$h9MP@l#(uKMgeJ%k`5&W6r@DDLAq<` zQo51uZjkPnyYc((Irlv0-t*Vx4}6fA-|W5DTJK(Kz3)nERQ1_i`3_GdX?Dgk2+ZZ) zn#v7e@J}kIkQQP|`if;PQWdWtqgHJix|Or#v0ur|EHb#8!N>JIdncpHdy$iODo|QL z!j#Nnb|jhNk(BM@#>K}8)mQqF79%CgBAiNL%C1gBB3h*wN*0i(N>v?P8;lcZzMzH< z(#kG!Ztm_A52L2GfBBw>%iyg70iHmr!>nAg*!%biNpB4in)ZQgCZ17o*X;b0QPI~+F$r_7#>h_k z_LHSr-Z`aU;Qjk~rbl4oY>XW=Y}2dOb*dk0Or=x?w8U!=Vg}Y$76`&$z92(7b;{bO z5>vbTCXYRWov!~>DYp@3-Z{Sb;_Bl3X9BVrimCI;Uvp3ELApz(hZz+Xok!weWP*Ts zzB^JhYieA&Ce_$88bNw6QMnlUli-SKPyc;)a6#dBfl;X7AWip=biiF$3!odaf-ZqkybzIYY)j zLa=OIbZ3m=@yzn`JU*|c?bV%nm&)XbQ=x}niaR)f5r3kYWb9A1vDeAo3jstQ@u9nO zP%~T4qQ<^gKf!&&yw>I!XVeouqIo-tJlSWwLh4__;PH(QULRnF{K&ZxuflIBrI*7~ z>Ne30{O!bXB{P5HZuNHR1(Wl4*gX5oD~W$BXIz_-0vq*cyji_V-UW9>_kn(k-u zWa-yGF;)&1L5kKkk0PMX^z_R|A?)UsX}d#(d(UMz)1dY@2F>}NbNZfqjh@N#hJNKr zA;d?=)0WnJsbylEm5&FnGmG>IGwB1P=k!bDvAa)p4`7!ONId0c`9? z`vkDkTc(q$aGgvwO?Fu{H#!LS7WC1a%XZ~&Sgtm3CqquxJej9sL$+S69mrXoy7T!9 z7CDwe&63m0KfBM5F^HK> zG9Y=WqcbZx+@xVJZjfMb+SP6Sl{Phhw*Dr;U(C~W&<0)L{!b(9BaEWcZz2Ut1&!E3 z_muXEnv3v4qO8dNKJN=HLpW8qG#VP3@WpyM@7xChr2nT+<&~8-gLruUO<6mryu2vb zqSrsX(}>jSU`lXI{Od1>DZ$3OBtf846?4T(&ne&lo zjaWjM8j-#q!KI zc(U2fQUN|iiPb;in{10lrkbTe%910#n|d*Dvoelx)lvz$hZLRowscAPq_(H?E&>f) zkxJVRMpte4uBWPqTufUDs69Y-e%Xys=R^3_tU4#Q_lKGCnBJ$;6HeamsJyt%vR#R# zKn`PVjh|;l^t*#8(s7>J#ph#qt*L3P4AWm_0GgBS32xpjb!8LC>e`w!xK#NG3qC@# z&Jn+Iv?MggL6xYa6VGd=-Zh6@UzI*%7%A*$RU99Ss^n2XxrgqNMlC~Ta-}&*>LKxQ z)&1_y&zxN|I7?w;uEJe4R+AE`QD!u_EMwWre}W*9`wkO|#dZTIAH zHrk}olQ&*ZIG;9i4BglxsmgBR+vTyZaKx0QE7vb`U-dp4z43mQ5Q>{+X%v+k_h>BS z`7()*Qje_SU70{h=18}QxdjDL<=#aD%q#*uxN@Jh%|9XxNsj<61{3!&pX&jwr8iJI zIC9R7r~$I<1bwU^y2;c=yv)5zRxz^+DY`#} z0cvBoe1yhi?*_-QUYy4i3~OqfeVu>Uw)H&&6w==GbBg0%_VIF{BT9Zp_QLCgT7K0o zV#SY-bMAz_gw4zv{S}AE#|K4raESr}aUiQ`P%>`x{8U0~WX@&RA-_M3x+;Fx{;(U= z@GGp)6-D&6Degti@F(tmN)5B{->o$IWEsk#5m%SZm(dPHZmmmuc(^^ga;dBx`3PSo6YvQDwzhzV!jEJ0ZIepP(%tLhHi#r<6irpO6U40UrY#f0vLCJa9(g5J& zeR}HgC{PB_`1oOsZ1w(PkRsr18UgT&ZQ$T=1hkgIELbQ;pS1p0ZHku?66*=`;L;cw zxw5@IJvvJN$j4c6)w}NZfW}}OtD|&v4>RBn=wEkq@k?Cyg<^|#xH~bhBU&|?K?4{0 zjn*3B?^A2f+{AIdd3~o0jQU>^pgp|lk?{4ADoEIYS_?MCs3h&h1!xH5V>bXTWiD|6 z0Rhx-Wslw2-w;eyHbXQHu62E(4#1StPFzq@HTyG?G?uTH<9uP*yH#+-`fFRNc@o!O z1Z!GPZA*HW<5s6b3k~?8TnU;9FG5hjvSo4aTCoG&63xc3U&jG|4T1g@k6#O3%B3P# z>qo+k4r3QtbKm~0jr3O7RxzUqtuQFQkw>g3iQ|^d`i&>ycvNKUK>vAieRC}x9SwDL zs~3S}q@;EA^%D)Raz4TE%b$RhV{?m&dZhOE)9GIuR~{-V6g3_XMp`8_kQ<>1)K&Xg za=5|@m4|Uo2fd-}^8ArU=OO`ht<~Oo=sP8WANQ?CiH{>Uj_e9~0jqJU1^V>av-+!p zZpsgP%?yH;@>brjvV%}Ypq6h+yGb01-Ae5%MsSip{Dc|V?zA`}r=HONPPbL1cz5b> z+A4Nw`U~&9^?a6SQL^HPg7VS#qLWhptc(lSV<9&nnpXh?^1cNeOYi>|36T1Eb0Y|X zal#mvUkhhqzwYWn`=^A>N|1% zOlsE9e}48|?NyIYDKou7vDCs{!BO)KBjw2U*_bDmnPNN*^Df(4l`K1ihclqOs|^~3 zM$ottA|lE-zX0=i>0mkQGoiiH;|q)wFt;p|Et@rHAp`@n7L`|`m`V~(q8S#U3e~55 zPsH&p33Nu7C%i+|bMd7#-0R7aeW!PzwYd6gwmoDIjCsCQV}e(hP-p=qODqqXm_nG# zJQ~L2GQY>xEe1tm?AaLBsBlP zpQ0NVQA!_4lO`zt$#2h%TIHEJd%BTGi9s=6gIh;fjnwfzA5eqbzO<^P=>I{XoZl^a z`wi?zDGWb>#|xCpYajOn#%uDdH{F{!*Vgvdbm>ic6Z8!I64%6)A!zNpz%Np zY)0K7fJGDv&+(m_==$^B>%3;RH+oNCqNSORnB+Lkx;R@NB5JSxXeXiTXj#N3I3>@b(1@^6*B&)%qDIWtZ%?- z+Ukpot`Y`!TQhJuei!7QayDP*4E@=w5O4 zBU~HzQqXT_SwZ|{YNlDG`Q1Ap4d&xhd`AU%BA_7BDjWjLPjox=t@I-!t|w3KJPi3J zw2V^&No_6afzDI4sGHisW99vm@$VJl;NqGziVXq*oQeC>!dchj-gdmto{1iV+6AcL zS0^8d8+I}>F^O(GrKUz@Aj^ftBqhBp1_khYy?+Us@**@^G}11bn#9uQ7Z)bR#+=l| zU&Y#?>&b@84J3G00>jXGr?SS2TU9L5pUjx#>nZj27l>e`>pXT=9+JmymCQUgS%MlB zQ%K zFAFr5%7JXB$1l{W{+=G}`}YUUA%K1`Unn^^(cS${2JW>^)tQjDzP>*1@f#161kX^U zi>To!ot2=ec*gtwd_tI_glvW#s1KHWRZ;eWr@GdC`9$d#0vix9?YuN$EZ@h0^j!&l7gPhAXo!@03bm zIf3x8SkYaM2_J@1jJSZH9?{CzuU{+IINI5XHBSPjZaejqJQWFObpeGAM@w%%zt-M( z-aBY$;}1Hh^V<45J2{uU-xy5aDw%DtdbJ~E7@m4?g)#pC#gPF)ORl~D5HrIY2>mR% z0Gh$AtwZ)enZ$uMnbMCM2$Pj9e_ZNYwkD3TCV=ag24j5hj-i^UAa&pnY79#lc#M>i z8x4L2^{$94Zk-LCN?MgIPMfYcd;Dm;ErAZHy;(}@=LT%70Lq~FV zFNCf(t^wl(XwU)e1+rJ;jX(hLMnOjAt?|*jvBLKmy6*0mH#Z*Sa4%nQux#)0tIh+l zUQkZnkT9BtQt=n!hpwM}gdB;(8(_*gYWZXd5;WECmtN}XCV5DJyE@Ryy>sUdA4!f* z22AUqfdtDce%#07&g~gGD53YooeUUh7N*MyP~zYnIQd%KNf9&^tgC_pyh-`LFC%z8qPwE8%bGK=X!*?7u26UjS#= zLGI8h$OE`cHRj%)9tV4SL#t)};CpcaWf1*<;{ro{?ylj{Q5yUX+II!K))1`-!G}+x zyFtqWxB*oGQ{nzmDl9C_cq;z&d>iaEAhZJcnFl7f@-q@N;T@ouz4kRU6tG^n18Dy# z09s8$SzYH4)i9SCwU#;a9@9F?k>3p^=jyDd-6T&&v8-xpufmy3`?1P9Bq`r@*nB5 zcmBIuTjJ+aU0qmq-8yp^R;Vb~wXTQsjRB}z=qM51|5=3_kh`tIp&b6L!gYy#1)`}S z@|~jtt;Bg|TR}h=RT=tnuQMblsO(<>wln0E06}(kp=1B4&dfW&s`$XHXpPHocJ&3 z!_i=?v1lX^#+NxhX@Sh^c)X{}KQ`6w2i^@t+9a;3$aUw+u%?nG%8 z2O7>N37^gtA%~6eep6`EF8J?~7-AhbIG#cy(hth+cf)T02qR*Ne(LvWGt$?7zpK<5 z#VS3y&PC~E+c&RIm6mCC#I2OKf z=3UNgASb+x(6_q?t;^B$mdvv*`0Sj|@;_+X-Mp>IL0n=F&M^CHs#E^VYW8>>gZA?4nk0EE)h6llKYtQLW~*;L z*lu@+a-iN1$hY?~Y5C`&ykJeiyaBC>AAfr*ig?s&n35HiP9=9JaAsc~rT>ALo~vRf zRcO__Eyj8wY!ya$uQ9qIHVqG%JeoN#^u|6R&x2XI{p7tyF2_Oyg$TTEcL@c&&Dm)J z0Un;oIrrL{pY*4rdCT`@{}c>bxGU%c1T0}_++4R8{G#pFkUzoC*Pf&#dn?#FZyt30 zUgAYfCf^P${7kZPotMWCxiKde{V3jRO^m0Jz|#!Asy+9+#ncdOj~Bb-5Y^+SvGo%> z>Yvd0w)hAG%nQiYyM$dMWR%bwGmqb>Wb9hb`K_i>RMD;TH=~ZeIlS@;oH5!ZR}Z`= zdz>E@)qmhxeU|qb10qH&GizJ^r?%0>P8D0n%1PKRu$kb?O7Oj> zbJ|G>M>UP>>RTOi$mf;|$m+Vbar8QpigzD#qqk-l-!CPW&Sp-WqlsSOh;R3RlH=TSe8TDEU}8(GN0TND!80zO;|-{spXal&$G$I_{`^97Ve;s^ z2ztv6fxCJW=sDx!K`rvPfnV3hC(=I6KhaQe_La?Ro8xHdQ-Z%JK~zF6|EMJIW&zsyZ@6pF1WBug*Q zL?^aN%$rZkDJNl?=z<-%d*-kwytbyho+8?RXF`9k4j*A`{;l51qr+J}r~06)`tm^k zb_IcK=%0)_!WleDVR>o|j%UTOggvVSl=kYuB-`CvBP;(NiPVxuj|nu34x2-W>=baatU{eyiB zxCimg3gYWfL&lfRUT4p!`)IdE8B&Zv=%8)mt#2{^euycimqB=yLp}d|)@R>5)8nN> z0PF?q3E}Ru8KQ5}fv4#^gXu702DkaVhFG&{VpfegS5wBN)>O)y%FQ-UeYuZ|bA+!i zr@o)9r9IUBTECC{?coA!f;a-{TEFhAr>kv6FNl28kvT?_uZL6O{=8|aUsrFp_;PsE zgcIgeNcP51DWQttfHg2(!`)B%4=RNisySGo`tM5bS-^VF#UdL z7#U!=f`9L2-lgD8^qKguXo#!Y3gk>Z0+-D*{yu!nSVfP7Y*MqG#nwBr>!*=Co9-+D z_8#2dFGQ-tl;paPk*~j<%?&LgscBC#xYP6mULv7=%(dD1Rk#XWZvprTc5}rus3zWS7~Yt3t($^q~L3^c+d~88^X=@T;qzoiShJ zKW;q{aar3u;QcuoXTE~ux1?*m`TSka-g@5TqVd!04NcXVfJK+;aM9*FPHyHR67tZR z^fwzq>s6W`oF};Y@U!!}-W#V98mxuE@W0Oeb<_ErD0rBc6L-#^v?ePZoo9*kS6o60 zuXmp9v!h8groBWWK39s=T{^81NgV`S(+Yb);TMcz2 z-Z#nzAu)A+jNM2%E_!BU@UD0&2#-Li1 z%V5a4UsxZ)ccG?QqWPxElJ62{q!;-r`HXjB-;mu{NiT z+9Pzr;M(6Y%xA)yjnV%VL(;*7HADjG zuP}rc@-gy;mJq4!$G>6bjJNB@W@mQj! zudSPEN|Saqv`In>Hn~&C5)z6ACGTY9bknb9pUPcc7sx7GTy76Vz4)2$(&*dQh+iRw zCzB@{YB|bSCHRoPfQ_xK?Il1dWX8GxJ1pm;O!&OoxENjC?Z6TIqLtycZJFC<-l$b- z*=mINf@R}+8Mv9i$BRz}ju-pyPRkJ|+LapCCjRp` zAu#4ntG}l98KW4ZYTpDhd|_jCQ>a#xBqFaIbE+$zDCMcwFpvGK9iOGXsO*$=)r6vp zC;zd0JtrAGX^O1mk59Mga`iG;syRta7{db#>GKPnn-NiM zwY4Sr*O}a&p82***LnZ@2c@&iji191jQkM}{3AVq>bdP=g^uVp$41Z zV_iQp6NNuCK=p#fgz_fMb2QIf;ejQ&bL}_e_F#dqb@OP6_wb_OdUGmJayOhh8(x*N z(o!O{RQ-i+EEzwg!zR=@{ZlFsWz5Ljmw?6FzQ=?gpi?QD;ueeR_G3Ptyb<5e_dc)f{&;%7!HTe*uUtxpwbe1TG z*Gh2&+;}P(G_S3f9Im_5BgF5@A?tbg9B_LDC9Yn8of-Fe4ibCY zK-BdqX)lmyvuyrQxBrw>(sOs`WK>zZ0Pb3%-n=8SPmWT#ymnQXZ+3WncQ=bueowmO zAgW5DU3G#h259pHBM(%w6e|#jB)<2c3YJFRDu`eUmo}qhV8~iuH%nA(XOHk? z^_?-3v$BGbe?Dn+*!Vp^FR=qu%GXv_B;jIWV#5~C1W-XExo;Vl*Ga0)7YQ2XtL=$0 zEI%JB8w-nrUjeYyjv-Pmctwdq6p7 z{&e)pYwg_{WnY!nWPr@72Z8V0=+*mR$R}9L2rq@7ePn57?Usq$Q04G6JM-&4Pqi?S zr?yI^#iczpDysnfah#=mM=~#(_Jr9jPc94@%hA7j91g3p?{L(b^}F5nci{%R{9~;* ztBCWA4Yx5hHwdHu%F2p)#U#+2&C@KhBV$Sfa=?&DAJBsxm`?abompL9Ums8@^;srd zcy&@?KwXN^3&^~0*b;IvJq0D6-z9y?BFRxvCY$+Rtx#Jh7Z(jKXm-R`LX`yjp2;g3 z(7Z3yN&NOrEK6eN3}{sY8S8On%vSe{JF|B+~#=C5q0rdB_WJR^SRq`w^}^ts9gv zgzMLlNy-LZdlpo^G~${7v`StMC9AE^pYN%f8e2lDK4xeC3cyqSoN>^fDA3NcdF%#g zZa~7Ixmi>vwMxzpU?C|fsiCZ+HwTPBvU>Vv_`~VpA;Dwy2dpRIh^C(R3Pktddxz65 zgndles%AxbPAMGy{UR-T9go8#snV%zf~baGjJ}2D4sEoT70!P9P!mYvv+d}aqbODU z379<0aafLjjDn{{ zCdfd?^RxU1O3$R~{fgn+Agz~01fT)y+R^`Wpi0dF;z+4lT?*eb#0T6N~K{UDq+dN#o4y50A_%*IP->A zy(`Z2O1^0cJhy~-+XZ_pHRJ2*O=}1KTs20&aaWoXnIu?D)UD`5dnhGuF$RC4V%|uI z$o88?+XXQblT09)KO3yG_Sy!u;#-c;%=Gm3pFd!NfTDnE@drWp3(@FfsPZ`~D(dTV zHa0ep3NKHJiD7*B1t{;EU4toDM{|Itr2{n2?&IR}Eb%I58_jl1RG< z!zj!ljA?*04$|M)*x0WdKEzJ=balyEpq?Hd1LdK~4hrt&uH_JeIx5UP3&1ARhmAnWudD=S z65>AXE%fRIYJ9`U?5wp#Tc&=A55fy65c9lH-~0CFnkj0ojwFh=$%8g@)00Xk3O$<< z`ga?WKo7_`{vA_@$!0L*kOV{v85v*tmi=}@snZir?&05oI(LCC9JP}!H8d#44+n%s zjv|}73&oEE`|+;>(-+BWF}u4}x|^UcSpUY^Mr%PXbnlyU^JHyKab<*um68e6f@B=+ zHm(AYFcv&EnuR=h<$86O3`A5-HWgakc!R(v|6(2#kV&6*F}m0H5}`Y@*h5dM2D{7_ zS_T)M$Ekl8`tKjP3LeWNg~)E-Kl>f3@bos?kIZTE4%mbnsAPySwfuyljQcPRj*q*5 zB~6zKy^Dz{;1VYG@=1*P**p%(UFnf0y3o>3P|r@%|we_ zQ~|iRENsm@Db*R~!Vl`kDnV(2UeQtqh^2yie9}c0NQKyUKsm(sg2n$7M8?NbXU{bl zuUiXSkbG`6{<}EzIk~xu4;3t)l7cqVC__W$? zixPK+xr#Av?-)JAaAggRHoG_YAPy6r0k@QOmg5TNAPXaeo_`@j)T;a0I2rk;m`q@< zG%#d5>=Z(2`hfBFeiKwaz0&I^l29bEw;qH4+QnjG9oMt8(0v?h((1oZD2@<$&|J=wrX;p4&1g@qmo5yqe`sbAo^78To;g6F*pC>B|6VW7QT}N6y+l2qLPRt_BU;LyTbylxp)-PE3*o;@O>3*XrCyDLxvgzl zQhr_@m{61k1B610inv5C5LlYEN%Sp}!CVGCJw0X^A752tBQiW=Fw4OqTv~#`^<-Py zPwZ`OLQG5ys5a>UU0T2uBq1eDJMQV}0kj+)b8G99{ryz2b?0=Hek9I$05@CI$y*Wf z)2W}(o)Mi)fajSh7RY@3E95Sh(vRzinNwu_v_SO9qDEvMC_7z<#Od*`yjR_TrFa3+ zw5$;e5KxPB>$|Led=>^1fy^Lx@7`s+Q9)PFQ3GXGP*d`$$7rh^#InlC&Q58Hjfs&C zNXM7sEG+@KEXSLxbD-V!7ax2*O%-E28N!o0p#k1s@yq1gs$9AfQyUP(9dh1KJ@I+m zY!?_9SopeHEEx#&f+EcYR>4pdApU`|Ie?I;lX?vbq8lt%b72Be`XxOifLeFO>iMY{ z1E6|E`rr@tVPoQZ1NeeV)OY?mqGAjs>;DF2h@aKT_K!m!?Y<0;1e=tl%pB1P#+g;4 zI$8utNl3)-v)tkwyvY9gg_lKRwKX+PPEK*`!APdW-wVG2m`os1+67Y~ouQNzR%wGt0Q z6O2SE4;Owo9r2TUtH#FCa=G0GWSPJ?nr6E;O8T|*7houpj{MsipXK3{)0m3{u6~v z=PmQWQ-(IQG~yN(7Qlo@&<71r;nGaDdM|()_OX+5pbii6n?QuCoZGLBp&Uca15oia z8vxN74v#M!X6T6(s-y!bYy!zSaPR8~SdQzeykb@#Dup8(<4sI1LGdnkM)i)JixubBK!hKJQGrprtrJyU2WfO${%7j{r-;i^j9B8>qzU0sT6uKenL-4b!X0b-^T z5IMk5CDm47PO)U6*v%vGTmn5Qd;b)VL)K^L|3>6ZWg|NWhn}u3w;!-ap5xrRfZn>W z@PL@O9ft|V%Z=2Dahx$mMJ@2D*gTM@(23*-f=|>%p@ev83*w$>A5{U7{`ZG z4@ruYrw8>vZW~A&Wr*8=%&T)F)p1owXgS6YHfH9o&dx6l9la79Yypl8Km#g3Vhd3I z!Ks}#B-jo|4;FKWH_607a3N%~XYB2CUeFf8IR)}&vN?Tlg8|zcuOENdW51f2oeV@(5SYcb`#)>3nlRbvgXBiiV6YSrL)cJm zO3h24&W;uM-+%1(wrIXcqT&Dfy*s?KiHkQZJR+n%r8Bk33J7r! zaKXhZs2Cljf(C;%)y&QAskOe6BMM?BlBI^;rVa2L_GRU2C3u7rs2OLg-IbXu)JpzY z?@<&Z0rX;wh~Y8DmP;i~*oSBmANR{f5>_Tr*qz%$R1Bh@zQ+szC54n{gZPBopcy## z?mus+kGJHCwqrE{gUl55ZZSrb?NWutv$?Cbcy9D>LQu1N%C~StY}@b%8Y2U>WQK5n z1jZEYALpWLm*y7u7hzpnK+i5bAca8xebg*tNDnQ9Clcckur{j(XMQ>V`B9H(Ywe)c zpvdl(!?mtH%95&l{L3TD=Bf>&=FD%9J0L@ff&#^ugT~1I`x@YDeU-fks`&r*A2LwZ znA^V%egXyLa&JGt`vV_du-nflmZ-O%^&zaGK=}dG8o+eC3(x|auEu-}sEZB`f}S@Y zY;1}!Sd4*`Ci!!_mT$8{2C|ltl73)PAMk3eNF3g$PZOTo``0HX2|cwYxi=dRv}b5f z)(6vRPnvx(qhO;n?-;-i2ko7lD9Xtl9vy*5W(WuKs({_w+n2pefAB>jn(a-gTE6UDZ%ilaMtuFc^bi#TaF{;O1W_n_;2>EOHEs+R!>jmIO@TqoMJc9Edv2e08VfNN*|W zVGyAl%mmDI`Llbz+!+y-ssbRNm6Z^n7*1s?-xAMZTKUFfir_qGV9LF8bK|cNVb$6I zGk(=`2qjr(Z`sy?-+{LSis$7s2fE!|#Bw<6Ap21=>jUCziI%17DXX&$Xb*f(;wD`L z@VrJIIEPVh1W?C)CafWhc8^6oroXir2S>iD_*t3-uY;xJ$)Nh`=ny^!`4vZy4T89H z=QT%W9E00(#Ij*~NceeKF-vzNZ+BN0m^3+Y0;J@>fu0o090OSG-*-LN zG#h;9;pI)f-I}%OMJv&F@A3iyFvnsLU_zb#FOV+)n$yNyB(5m{69IRv_3f3f`%qP- zSC4l_V8!+%0E|7l@DzyERBDD6B74vq`L)!@uRDKpyo#0a}wI0z_f+oa@cD zQxkW0fr|JEaFGXqF{t&mYI9(R&)=Pw%Tr5Klq$#rsKfO)`_CA>(8s@x>w&u0Nba)7 z`V$Bg2;na+FFPP_KHA*@;tj8T!z;J}?d1!B~<;1pp(Apno6gH*-Zz&14Er zFCA$BEM~|6i1asH{t*BicP}p#3k@B_I6&|~;sX;C^C3%gH%vS|$Gr@6iZ7R)Lbc=> z!cw%l3lsn2;NKp9kpH{=_-j`H#Q*)`f3TIm-~WI9%Kv0#|MPDG!@T+m%uHyOmrp?c`q*hD^nr<_s|e!Z8^}o1s&l-f%5clAatv)J_2ak z8$($@R+09{8w1i!en zbg^hi3MgEq56JZ%fhYzBF$_XM_7fzDPFewFYVVGxqvLTP5la>F)%no|XiN>D^pHjd z2L)Y*VM>0Sly3HD468rNo3I@*>?LR&wNPx|E^!3dG z>AC@NjRGA>44_5}rp*BWVt=XZpjTy^>aC)meWn zwPAo8C;H06`x*pwr@+>4PfAz+Ql*EZ@*R zEMfZ?cs*RTyS{!anrF&-M&C3ymyEj)IOa3e_U+@)un{U2<=3x!u9}mC+{aGiFqQ<7 z{3a2$)-#28p8alG$}-SFmZN6zat5dtl9Q9;AZ!ik{_aFm=(+Km5LKJacp)bMs4@}W^ZwYtx)?7&a{Fh>6_|Gbw9v|d z!VOyEYc?R4WC|wIn?kY^h*&g$vDjV$OY%6GFa@gt?jwMIy*+_Z*!Xtgr{8un1{Ylb zE;*M-men{&QxV+Us&SW!AiG6uRi*%1#uKT_qgJ#m_LVuv=HN?_)-i) z(6Rp^dqffD=;?_BBkS$-$^jj^zn@oF0%*%tdrJmE3*eQ1fah?lPb$Lz;pr+QFeN2r zKOmHFd$QskI7A?w2KM|a4Kx6c58Yoj5x!z-3!oorQVb2LvYJA~^V)=ne1Li@+&Ven z6_c=*iE{Ju#)|YrmWqTpsK5-wU%!@QWcfUA1ZTm|0RJ)1^zgxh5{nV$=0YEmQ_9Pe zoh3Ye`5*)WF;nYICU_105O82U>R{9YOwSDbVEBKiwgo{5pV U_AOJ+oVQ=gNGeJcis}3OFMaSw?f?J) literal 0 HcmV?d00001 diff --git a/img/stream-compact.png b/img/stream-compact.png new file mode 100644 index 0000000000000000000000000000000000000000..3c21ae3681b5fe164841c5cba8bdc2ea04397c31 GIT binary patch literal 29217 zcmd43Wl)@J+bxI_JOp=wGz19*4+M8AkqBk(Wzr+ZJv_V1H z>`96UDmyJ4xg)84o&8-$u&`q_$zibo8cx4HJq9OV*0q4G@|W0tUvY?qjraUPossEL zVM9JHW8v>?bZ4UquOn^o?mUfrOH*C@S>!`C6 zhgo^cjLXe}tNSV8_5IEsumIp?*Z8>s{-3`IVEyw31^)Mq4f>z=F}VLf{pY5SV&NaK zKsPiLu)HCMocw&N`KsK8hPf>DoV>huA<@zAl!>8$olL^UNr_rimY0{8mX@ZbrlzOs zXlM|^BX6pyt1G|_ZcxYcfBe=U&|OC)Bp@K*;83ShoC^K*hZ;!&lTO|NJrh%U4lxlC zQNNvkV4yJ8(lF7dD;z|qF}QYo@`S;sr>DaA*VotKz$+*ylOyDL`L}JpGF|OA@YwQl zi+(RUh0r@1K0ZDQiq7fjxMyCwy-2a_Bxi_j6*QPx5&{L4fds`Wq^S6=w$p(94xNpp zLSOZL$fm8f5&}&ex50ShnFxnH`mc2*h2bv zzS&ziD~nk`@3}KXkR;QSo{P>)y|}U>yAOPpQi_)s)PTS2_Rh}9((f3dFJHclr16@4 zH&ih(`9S^iXzo+XSy~z?^3C7B?}`FLLs6w>W@h56vX=^-9M$n1j4YQL1mUgi>~dm# zeS8iM55+yP5`8{Oyacx|p) z7DGeBsaR}WT*Zq<`$H8@PEPI&0lh^L3taY-Cdbp(pWDc0riGuf&R`MIQ+OOt#xjM6 z5C*q@I^YChC=vmeXt1TIsHg@vLH$SzZ?^qmzQy$hL`CHn5HQg;)v4%knK07R+p9z` z=S?G*#wHdVS6WpyaW=Ky71)nnIsGoyL`@C98x<8*Mn)#5xTLf+NY>2ka2Y#b`($fN zq~5vMF3PmwO5a zor)>RoL1r5STeV%pA7WwZ3uNfiOI}d4u4XC&nvbcpbHNcRP>S(>dpuZ{C*I2_v#C} z*kCYtDnq_n2y7}Y-n-ZPMSaokY)w!PBb1N_IY%3vOsE%I2#xF^ft&N>NwNs`_ zcwxC;T01Lvgl}Tng1Jqxg>nhE!MYAUgm!!|_NE{FLj5Wa##k9C74`k8D5dUwMH`&X zDg3EuXd?dE(767mjZL=J@1coD#+KoFIH78#N=h~c&-<%IMRXqRLH$0ht$b{bz?PU< z3dZ$H5e`Gsxa|kziSZ}l7#SHm+W`YTs6hg=*Y?>Rx4J)sZjY3}JWUg!?NG|&mh>_- z2itds+MBktwCpn&T0k1o)5-SV^X|91M4Ah@y7}uliaFsdzjX+_%2NnkhNI)Rzw+Cr zffdNjjTJZ15X>cSxAmeQjdiWx&PJWMb=0M%le3JKP~5#geo zoS9KkP+A{L(22C*eAG)UwCrYZ$-JRq7*> zmFDNu0A4VZR$9tP=AGWmY`)#p)bzm*BC2;1FpkO&&5zB1Tva79+1o45lp@txE=;VW zhkA-i)>l6TvV@^ew2pUS>*5faWFa&S@TZC~k+8Ug=n9>LMs^#M=G$QFv-OAX?EEr? znH;^IQvPOd&%V8}wr0KY2g#*ST)F2pX5z!wCQjw*rc(k&9U+M&uNon=w7jf@+Rhnb zpi%I935X>LIv(oE%63opS4#P^;091eIC<*v!2wH0Oy+<7O9P^c{RglY=D{X95s408}A|0^&sFo<~G2)*++ z631_zN{APs;L6b6dk&6xh;*D(t}@++a6Q%2iCk%HSx=gozm*h~{XZr-y$AB@A?_s* zUACzqajZ9||MP(4utIOmng=G} zjO|I&T61pSm$b1<@av_Uu;@|BexL$@{<2O|%1TE_49l>2ucZDDKe%P<_%+Q9k0VaO z$IIJNQ85-7zGuYsdRa@}YiqiuOhA0{lP)Wm6bnmeYT_Z=+uJLct$jBQ*WcfdjsNal zrzZwDAdK4w@XsTi>mV&#Th>*drr*DjZZIKl!2*nwwWU)hC`951o@Bf(%5!4X-H$dl zHf||-d7!RS^d6(TISD3V8=Fu5XVP5F)@=tSTwZ$K)$Fb=QRq*$9PRScJf$BQBj!q_NlniCU$J|90h<+WfG=Q&Y1mew1w7zG zT9F#x^J8O=Aj{eGSt4d-`sIQgN zKd=AKYrr7`JA35P$oxSU48%f*uuyNkC1R-bydz`d~%YNn|q<7<2^Zf zMs;<}1{)_Q5Q(dA4KDirCdv8<5fTzUK0eY?QBmd8*Cz;RXlgDmEp2RU5R;I!G&j=< zsi`G~g$tyommq!p`W0JNQo$`r1_YrbF((i*HKo%j+U$+cX?BW2M?gR*DJh|P zheVQ}o7>&fgN}}_XmGFU^d6^ue*W%gzFJmRHr7c~lSsNZGGq@8lM(hk4NY%%x4(p_ zsA%D|>G+~yUlj7@`nsMPzVU0@psUCRog`NE6|fUze~S0-*Z0P=4TAgJaB*-Xb0= zE>u>8R;$YT7P|bli{G%BWFFjuorkB@T_ zz9wJ+YJ(!t*ugz^Z%gAm-P4s82M33!s3+5HSgrH~@s}%QtVRbc0%>Lw>5*hiHgv5QT50jhEdJXEy_5O1IlL?H%#~#aH zdISY?Qn`wpY>1VfFV8hKY+aopKm7o2L{KD$S7Z&2%4<{BZvJ_8^prx#<;q_OJzPrh ze6;(E$Mn#xQJp zz`DoaaB~qXaiU;!VPSI+BnXTVime<=TnV-L>YvhSmWDD5fD~$AWo}{7&lT1GUTJG@ z51)Wwe{D@5FI?wQ?nxEd(bhI3EQ~#6aYkB%vxXrmOW{qDB!@a5Ay_w&oj`YBd}YOB zcO>=v{QPmMRWID2?)d0vVsg?So>EvA;>Ey8a&vj9rm0!&?4Ai)V`7_O=iK^!6!0s4 zz0)6MWMl-Jl$3Pj=Pdm{OP`vWx_x*M#>zl~{SfH`sz5Dhnl$))5is5U8xeb(sAj_D z**vi3=3?jxBe3OZd2IW|;Z~TH;bqw~+bT$16Zz-QpL0VnY|s$qNSF`CFc&JYoFjcz zyDaSFk}eVomFo0XZOeLJ>7U9#kC%INim*9IM;D!ww$&L(8Av^T4Gj%DJ3s!{!Tl&f z;FgtQP1~9F0^GOw2uDn!VD!m-2tg6{()F-e^|HY{B7x6zoq05FjR-`-}m}4S` zTI{_g9%ui|JxoIOT$u|igZ2fhjk=}Q?p8$3UGyg$L<&ex5U3o$&>2#q#Kd;eRblNi z%->8Y!FJ}t^7CJp`_3vV6V}RRPvCSPTr*6$^X7p+Y5bW0J?XTupALfY=zkmNOoA#wp zp0-AV@%6F)J~-(E<*;sTmN}4?f-Z2YErPDdzj2?1oM9-9a+&zz;-jM`>ai2p%9-zx z%=xYAaT}YP1K~ASgo1*iuTiP$HCTLk$_nA1JF z4+R7Yuevzy>V#;e5!>jte_{8|bp5Vy+75a|Pl$~@J1^WkIXz`+e26rzqBMoUx+9c+ z>je_XMg^h?*t&kK2w_X~^i4fAM;-#5B}dEO+(oOKkLUJJUvKiP!hY(u53zfn(BJ&Y zTJg5;c|te(@xv0+>m4>o`bV;VeooFW(4ZL%A86in-33Pfj*y-PSfm)i>mt`I0}z zlFXJFgCJm;VJv{9uyx#KmRuO2AmtLs0l1L7$y8vquvO(J||Ik4{De?Les zNMz&aNG52^-ylCGCPtLqyV3yoV{E^o(k+bojR01L0>Z0H<73neXnS0FWto?*X#44e zy)OukZmA~r?7=|o(@G=kZTE6Eg;5yo1NI$uM>6AH8|qcUt~I@|U^Wn7g75!kbv z6KTytf&mgmtRQj@;JB#va(PtSHaiUxcCZv$a_)|o`J9}d$&%3B_YaFcIs`ZWd}srO z9|&rU))z2g2eAP{x$DBr%L!Hula`fsuQFN{0X=@)ON^TVfdc9jD;^2&DVgs%CsKN( zp;Wt`8Ev9A{gSCy7>R@hbAOkId7jT!2=WJ7Rg|%S`g6?!@k1m;J+y|do862Wx&L$7xbV!b;QHEccwRDGUJ9H6OobyznTMT-w!kSN18oM$7aAiy8PK z!wY!VI{v%j?BaI^whCat?Y!zqas*5fSSCEQ)0he_+BUx_l>8WUlt8mHcUxm` zYin`rDA$~c_#Yn~8CAqC6@5$2sbjrv1ZjO-q}hUn=#y2gdovxAxpKID{jT)>e{5Mm zEXB2}xOm>pyV?5llk3Ab{nm*)BlAt~*8`R8S?Qm10&EsIAzc$7qUNQ)rfQdS&Y}6r zBAfIzKPlT4t?CoLx|B0Ay;GUlv)3j741B%j^_Y=gB2dXp%7d{FZOcESVgM!l7~B{fun2PY z523O4&s`^ItZ6f-Di`a)3q#t6WYL89-#puTPn?}Oi{!yxZ?&|v8j$|;1P}sX1zB#s zzy|X;HH(i@UBom*6Bw+{D_pz&qSe>a&lC`TpDmnSmiEc8JG^nAcI}^%`r^t;`v25b zIG%oXkumfF0VfC;2{Xu4=AtP_Uj$7OorhXo74tRZ1@~FdwzN{o4*VS!lt9A2hcB|@DLIfet3B3|LEf4 zA|@`*!^H&$2WRKtpqK(Ac1OpX_2h2~g7Kqv6xb7EV{5y+;_rc1OUn|oqM2E7aqybH8c>aDlI z0!*rl_{2ns^I8o?6Q5Gedc}dCzf(}#04ys>ta0mXy(?cf^}EU#vl+{-!H2_x16u3V zR(AlR$@(08J&uBcf>6{(XNP|bN1!>%NV{Yo2#&sO!S_X11>#i#BFJ2?B>UZ*mnX*MA^YB-yB`-a z0~Zh^|K{ph{g+sF=7^)XIHCy*pP|&B*;&Vi`=9@BHfH}MNft^)h__i)qSYX$wJ+R9 zZO^hJv)p!c6DuM2GW(JBSxWwu=1q+LI{x_5wg^1L=;2k%)tf4+hYZUK>0(0i;&M}@ z{EvM5P>Qfwblx(2CT|qq4Fm|-J3>^(zgkl#d?awNN@R`DDH&t@?aX2Pz7T7mVB)#g zzE-`u6mp`7W$@#Nn25+5;me&V5)zW8#zvXf#*~GHB~eeG{KE6;-aZqQ;dv+W_u>MB z^oucrS$r#CxN&hf>gQ)?*5kzXVmp7`-Q4>8TwPqEo5_) zT))+Q2Kul9=lx=O$?{`{^;w3^Vw1%B-U9r}-|KMlNfL&5%Wuzyg zR=%T^ik_G@w!4FVp=2@}O5DQ}Mr2p%+~{mbPcT58wT?=q>ei)saf%v2njl(zf`uQTZ=sc!>S#8_#`2Ex~j1V>ctHbmvjCDYR>~4B`LV8pp znsfsJLy=VV#sD1yqwn4qqf}R#fV+-wO2L6}nr%ccAhZAP-@iWrUk4nyE-8JA#?=h~ z$2kh+#?8&mArQ#@{XK1X+Utg0x(r5MpDYtEi}P1Phz&?Cea~sHospW|Um9L3?Hv5a4bD{T+V1H}cy4zT`F; zK7K|b&;kee<89Mw;0K8Sk@vqT6(aeeY;m>;+`MFcsZN`sG1D*zJ4~ogrmFuhfVp>- z^g}6y1GE@Gm4!w`?B-vi*-}*l&{;xSnn$lMO-vMsd=T8r%gYO(y5`zC1waoTWJOMYuUx&@hLy(aW3&|(XY)zZhwg&-Sl_i>9jeUkQJG9 z+T5(}X?~xH_W^PTM>Q1K8WEr#qoSjQeik;~U#sK+FcwvngsH!u-%_JJ%0H#3pMY&~ z!FAuwwD93Ay0`~hhpF&_YPx#`g|Xo1%qcP=28)Hq*6{^iAY5{b;jWYm;ji(`;2iRZ z7lT2hrln`$S)~Y|3{JH86%t}*X&C?-3yg@6`}wmfFi`yK^cm?6VOURhcWy<6JSe97 z&ma7WWwL4=XqNi-wfRffZ2|kAE{glFnOF58k32cSrZIYJFR?@Ui#WO7m`o1j@#yU zDQwF=KXR7B4cLgr3dj1dqYu*SRjPWgpYef^Mpu@3`~ePF0^-!NiMZp7n=wy@DL`f;8zP=tn@m_vsjIfa42w};1fs8FI zwjJXdZ8ig9Vz6Gj-@>>c6Me-V&LY6+%+bmffH3$%35s@j#(tY*dIzw z*VL9g9nH6{=}Y*}=!*6~j?;|0kePC|en(M)Ww&CA*)mGT_@IyOjMf61?67P$G+e32 z4SK|z6x7$JWIu|P5D|e>ca)j|0N!u_OZJ;507=^O8g~aPN9<2wEN5^)>%H04ZP=aLU!%V_}@l9-b7Y5+}b_SxJm}s zz`$S+bL&4{uQ5z}G=zEN{qXi~Vj#X@9P5H%WTUHMlB&6Ws)v+MM(=fCW5-A3$%gUy z^5xT{*y>vxjxX|U%aYU1U8s^_D?Ak@vn^Sc;i8kD*4r9#a{MCX+`o8wdX8&_io}p8 zU?uj=s`&uC={B35MqKln{m3Vo3E^n0C7l&l2H;VIWQD;CZ^$BE$OXw?HVzx>6oE%Y zK@?pr?`UMuV%uKAOKRrK|9OtF3`?%Lb!%DK$k=eDF*o<-a|kOdYyBGmBxKg@3n@uS zOmO@xoSRJ5P3IL{&tG;GSzvG9A?oV{85d!Y|Z z03?Hsg!Cz%0YkqZLQ2CNRbB}Pun~jCod8iP9qP9gxpu%pICl4|qP z{VwRw{Cv{8Oq70E8XB4qQyBgJU{jdPi|+vbA^jEtRpAAaQQsXceNPT=?Sa*{TRij3 zXPgsP}dovS#g!{3~bv9o!#~h(Y%WMaCw?{wusSxYOe)0 zq5w&Rmm$Q3fMcgk?Pi(9I`aVJ@`FN+}hbB0<>|}{?LPoFT zyn-a>f+y8op&0hUUCoURX!UZZ>U{v7(bkWKW0&H02Uj>Qaxsjiad>6q@{>(xsl0ep zj6ICh7v7oLs-**HU0bU&Y0IHUNrB#3W_+=Xpn-$o&BvC~YS&mSxXkiZ%pgFXTP&#?9JV~2*WL#WBxDv~3)?bYIz>amkUb^^2P;<ZE|i&uLBR`Lp&)$D>Z!I(lWNYXI=kMP79nNfqpwqBoh#)67VA(qc$LghQsB_2^(dyXaVcc%=FgkQC40}-s5_Jb(?9|GjQMKK#I7BS{ z#YOw~92V-eioc(r0#Nb+%(GDeyythnxb>-Z{Z(p&5rQu_?S7{Ys9@dw9PXT!wy@iT zgv0XB#|?MMjU%SYF@-lK8Es|}I#+I%N}X8=J5~=2f2|v3N;|3%0mwn=vF_pob%ehN z#Uls{&LQF_3X(Y7!$+IRSt=Xbt!{cfKEdGL48?1iqk9auwC%J=`G$$=WF7$2+%aFFJcS7>I1Jyovcy$Os<)BTBVn4nCfTsJPYq4 z|KkWQ9VQREFnAIiOTD}bCeH;R{m8Ge9PW)jZ+W_AXG?z4hK(o`vHyFr+$#aqjeyJuaQR$7(Md^aL@kK+`4Y>D@#2Xc}nnmu6LZ^2efKs7`m+z9spT z)53j7X8q9ueJJY3GbVV>9JKAo?wuA%o1bnI5%Wa3)9ih2_cf-f_P70^hDVjI|1x{A zgt}`w?!SJwxL9@~V{n2x)lY{0z{!$$37}Nh@u?30F9CuFz`yDEql&<;mAuXBuscVQ zmYN7nO-uJx-A>}K(p->=c}cpg0}5Uu^fWo!h{+zSJ&j;_u#oZeL2_@WDXz`pTO3A+ zRXxXvy{Dv9A@aXZ++FloOcgxf4_qclslpZ8i@aXg^^bX@YwBen2q29z!(kTDQo8pU zHW?S!@Pu(APn^ItH3d5G%=C@<4Ra_4f&FvvbS@r}e=Lt7^oUj#fZiM6%?$YAj&qMLI{r?L#M8u1L46vP4|&QbVb1c~}xeuiijQve?umsy1A zyJnlG)Mp;2ba*nWeho(vj`egxA*{sDjg5`g0O2%rVlg`a0rr^lvEg{t*cb=@>O%W& zB1{r`01;7Fs=ZW44uRA!dC@_PXm3vzP^j~R;q={ zohU~6y++0`tS*aIO{@)PM$2{S7{~nla(S_&g)&b{M`9l1Q0n(QJMe!5ehM1sb(}c{ z-AGXG(w3=bBPFA~#sN9fr@i&D`rzX>z_~2jDq_(>X__FhBo-Fe$#ZAz34?rYz3U6u za=;1_0j*LCkTMfT>QilvtoI2lZ#Ga~GHvgu*;fZrM~{Zk@<~aRVB3Q4?X2d1^Z6UZ zCIvGw=&n+8qa8ail!{|<#7Ni=-bg1z^t(haLmVDFQ^+fx%gXXe zr*pm*yXuF3WlEDaT|4N+BjyWPA^sFZgdwwCoqr;$Pez|~+h4T{k;nwpg^P=eBO@Q# z*k1KkM_Kv#Ms{{JzP`eMPSuEJHt~N+75MspLH-Y^LIP0^z+Tz|EM-WDt)1QF-@jRJ z^z`&NI5+?)_8He?3VVgrLb zG%%n@7pDvr&kzAf83$rxk+q}^2S-LY7y^t8@P+cCMg@zpud#s*+t}4`YXYZ3{&;`K$cWvdU)t4Hz$;a47} zJ7!DO99rE{0(iX_6$}!5G|GqJJ9fj?=wIS?hXEY*eaLCs(>YO6JdO{{G-97abky5E zGg5th|JB%NqS_I_K1)DJA01@!8vi%ISC;%=fG^NY@c#h#5=eG;b`%pd9Cycf(`%)2 z4^B>gcu(!~I~*k%hi5^&xKCo)`!O4X8~3Mv=L-ycw|)aKRBA%J5Q9?y?v+%Fgkcu@ zR#TG~;}1%dZN_|Ar_^(-DFVoAa48*o>r@U9*+R3;_`>V86k0P+FL}i`(`6Q)QB0e7 zIQY(H-CyRSNmO_+ANqeWFuZk3pdCy?1nZ*b>`4{_kYJgowdgx_y9u95Vm=3h+)o#B z174TVyahQp5})K5V_ep7k>we8Nhi9k#_pdqnyy;cOUw4~p2iu<^EygGrma5enISoP z?pfJi&eX`^CT4>AhK53a{dya)H(bcW&Fx@spY(b#K$FC5+b6K|a({Iw1QvMt8??2% zOZLTSOl^k29RmXcu)L~`_4W0v&j+C;G-v1M?>}z@N1DC*q0p z^W{nNB>Gf8pLUG8O|=Qj%5@oXZ|b<*6UEUq?_zpRoo3a#nYdE(?`!uB@~VGqW*P5B zVq>SENinYyDLEA)kf#zYa`7>oCi@vk^RFSTM(ZWMI0rj0x2Q%{HGMej{*T4ju{zgd zT=0t3GVj7M|30$Y@H{X zL&Kj}m{A9!gDs#u&3pB;$o6&tzCZo_g}J$gG7zuP#>&cZpzkgFT0Bj~o*2hIG@b~O zkdXKl9?ti3>kSDRm$lb3ulxPwEgLkBK;c%MUzAkt!_BFjv^49h_hcds;XyP-gAcDK zYG|C`mt!uoOQfw34Nui#>Oyo5n9MT@ZN*Ai+4o9)c^&WSlr~$h9(%~hws6s$8f-`$ zukdtKmGi*}N>TpsndP>t)Gi}YKfaVs7BSQyRW~^NmV2L>#@lA8eEV8JxHPu*aSLDq z++JYQbN$0_t;QlYEiv+XbX+v~M{W$7e5kreFd{V=R#Rst9hu%DTC@p%<<%}lZ9KR*p8zv|RJ_^CT~5&v!-TWOgR zr&(lRAnML!17uPAT=euIYVtEK5U;ken)AGzFTA{kRIsqHvHGT_MS7n^q*1`SNE^Cq zkgA2a5=tg?svwMkkd=Lta7n8V85|JvB0Zy%-QMS`8E2-0V`>f@$jSU+%P~N?6#(tZ z9CEzu-pOvj6A^e}<|hRT3U{$R(eLu<+}@o62nJAUaC~?dWp{BG=QZ`SQl!^%n%oGl3%5?K_C}E1pySM;yk)(RZq6^?)fJ_R|(HYd;b-zp<=2~-d*3ZqU zbrIOKh}a1azWYGvAK7iFbd)^~7oZ-Km_ZC^1szY82*}7rnUt7^qNsiWj_2ZXPakD?k6?H@nP5$_*mg0)nO~H3w`)vsRJrWy;C#k8PiTx81gfcT ze9VPYBDlCIRXYeY#5)p-q-4FX4N*2p#IeaOl|B4pk zX{;CK7g8-`{V?wksW2sIBfsI_=R}gKI-d^!N;N=TLodzlepj8DC&M}+k;~`}{Z*%w z+ODEWH!TJpR{=bE8}@v(L<$9IZ|?3G3xMuwv&E24E;dD)Eefc2%RJEnm=vHE_hFR{ z80<(3S`~^XZ=eVYJ@-!;;z_|FaUT6}k71urDo@>NI@C69AmFLzU6TM78pIewM0lMQ-wAPVc7cda#hxjnhb z`A8d0NUJS*Z(lbZOG{}l=a?JDp3ETOan(Ws-Ti}4ota3R77C)P5)A9_LHSXM9;M0D zmp+=Wf9^r%5qCvBN@eAvq;tw*3zzfB8I{eSgF=JMwhFKN65%vk9B5nYZT%R~7+iqv zO7R^d9B#lzezL1dmP@oJc<3@Yn`$VkzJu!D`z3u8GGAj!P6&Z8Q5rcSzfg7mC2>skTS|>n;U*x z!Q;Ede{e*KdCv)v+GAqK2;bket%HnevxC z7%T7ANx0khjh23TF9|nRTJdC_aP`aSOGc9_eybww z(kfRU&cgy=z{Oj{;Exl~*_eozb!J~-%|iv_Z~aF67n3D+ZeG7Z*=lfjbIJeqjrss1i7|akh`3ZP zfM}hjbQF4Xzu34DvG}qJsybMnAC_0v`mLLNVZyr_$uhzKgfR?a#5I1l((SGHUqXrS z8y=U!ExSc7k8C!{fgTFm!5}D+i1Yly7Zm_#)J81e za09C}Jyd~4;0E$tQw3QEK2JI_C@(#oq`uwZiL z3E_0~G}|3MbmE0=6E#FP54pPK2v*_MW_1@0PRz1uK#~MVX)M?y&)Il|HW$DD1Wm^l zUGxZuna+~77jQ9-`f@Pm)ji;=D{D*b#ZdQG6VY8@^ImuO zl>MSm9s)6pFkp17vA`#taZUMAOO_r$2Dc4cF`WJ;TX;&uI6%nxDl-RL|0>?30$ z3F#dKSUxTF>NWMuH?+?C{xMdiS5K;AyU7zAM~9(y~6$iWYyR^V|^P$Yx3F4OyTl)S{cU3;^7+bfa3eXDIe?s6A2~LOEg;6cFu5 z$URyA`FxZ#>1U$c7#*|pHmX7+)x&}H^2%LY^_lLLy>7l;bGL;!sU9C zgQB_dKHMV6wc046a$F)8P?D@9>RZTa9!|DYdLxKXX@{De8A^sdl#VByg)jvm zUa`>ozO7q-=s%;KApK2fp_V~rSo*a0ma<8?kmH;k$YH=V#hOL>w=)zJ<$Q~N8B^k& zK570m6R?`*hh4S4dqf>HR5$SPLs#pW>*mJuX_5ZR(_E%rJ$4(aPA$Eos+?znT-E|LfOdM{^w}O)~IBI(B8?tzqOHD@+ zhcdNkn(>iPJmi0*U<|hwyLh;2I;a!#KUI`Sf2Ie8mr%3o{e`weubN z!O9}jv_DGf?Oa()Ret*XQ&?oWJ4i!_tC+OSJ5>3p!%K9O)4j3EEMiqaqpC{(meqU2 z!{U~Dd%#^$8@ahjTc68Z6V?XVf!>BKM2(Z&r(2l&b7nN05k7%>Z8cB8M4Ia*X>JEK%)X2#eeAUVNKIlwLf3#kdZu97hro;{w9iiRgT+bI7($sXcZ~*^!DSO;BO_MG|=>2m`nyjC}X~0WDcYPJJpP!|mnR;!%^x{M}0xmEu>ErF<+M z3$B7%*PI`#I>(*QE`iv;G|Oq=bRpn3=45Rh zi9SaSh`9jmbF7xJv9YIT+gYr(X>4s^VDEgjQEE!cUUVIYWmfP#v2;A#Kpz4#4tNMI{|Hl!|BQ}7ESp0__%U_mzNiybzk4f{~tOsHWlVqZpQXl z-?O$NW9%>jIHs?zCBI`RVAGHKElAF%E#Ic{c`ji+-lPXC$HRDrJdYKVXI;;rJ|%wV zmasp@&?UnJ+@oDc!_pKzkA7v@*7bKg1a@MN*wCoWVWw3cHwOv%KO!uH@qRd592=c} z#n)jLb(g9JWeFDh*Lj7uLOHp-D059&S@cvxh!;%%Uj;b5DnREYt3Y9Q42YxW)me54 z2nZNpJco!pxXEobRaC@4{Qr*NJbvh!Qu|a@RTV2`=;roClw*d3Q${5g?;o|;JNJQ* zf5c+*_Kkj3JoX3JlW;jnv@lw-pAq2UKfM>d1~|WNm?O1s4- zWrXXJb-jKyd9H%ZqL3fMc0Ez_{PW(_$#2g$&xorqt@xE}dsX6x?P&PPNVssWkBQ{? zeQS+8+l-LfKuc-wdKjwrg~12tXxg+W6L3=@_6`o9Y$sqk?!ZA@)JC7*z^lI^>RX2P z-!rp(U}7~q9DUslnnzYjiu>o*uN7sc zHD^@+rVj|^)D^UIko=GU8BNvZ($Z2z1@m$$7IiE-(wk^(y11gIret_+U?l1!(>NSp ze?>X;%=Nt6dQayZg8)_51;w;dIgft1t&^KmxbCt$e41SOo7)F~b*K!mS{Zk4icw*! z&j&~}gb%!ek%~p%=@GQAcQ^>?3RE}!t3<#8uXs!L9)P!qjQ~^l)zy;#;?mjwgO##Z z;KyG8u>soR`$=Gc3=xCW;Jty_Jo%kJ3SQlt_K%-Rkyuz+KggGdjHXXWgVj*!aI8Af zz~y>XSL?gaa*+q-9od3m0@>==+7evB#n2Lyi@b}5& zA59i6FId1EnC~)xo|n~wba!_HQY;q&0J4fDS#8NX@;gLARU&EubWJL;0Due~M8hFF zHaE8i;J33y)23NnY~FgW7n1vWAsxMZo_ed)bW-)QG=<@2OAg~Y7MwXt)}FB!ZEekm zr}eo4mrQ0%k7*YxIZ?xP5D1H+kRUldV?13A`X`n;zcaf`WTq)p@~_(3vcS;X(l_pr zfq|;Z%Jqmn{7~L*zgSg}!z9ED ze4(U=jV${8JB9l9=U6?{9=48^7T3aL+4w}Jpx3Jr%SM`WfoVMW6NF{XD`;q#+tl>% zUw~wvu8SBLyZMNQhW08f#)Yk`ARj%@8$ms7)S0_Wqfl%aPQfzja1hSbY-@X!igRy4 zeo1;C_*o4AN=e7up&bYWMiD3|Y2R=^)S97Nqr7>pRAQNrfiKAW088+>$^pP@{(~Bn z-kPQd0ixQfa9~8lU}Emi?*b5cqzxK6I&K-AF4jv9Gc9~o10Wz1van!OcN_>mqMesr z|IjiOY64Sr3Cw)$@9zuwxbz4AB@6=rzda!|G}O-8I-~li33daRH88WXY6FIGZ*QGE zl6>^S-8xK7P3I1fU*ly+2T0i|5`9D$*{jPq4?NK@Dg2oz{7N%-hm1A?VDP&SndI_p z>{M|D+Lvp%Z`DKxKn88l%2;e2M^kXX=rz*vNBSCee@RQci4!sx_m3Lt75Jr`mMa0Z0{^PhaZ0dsQtzKX?o!FioYnUe=Yb z+hSOc_*>y*=H*4lHs&{-`*KRY9b5IDL0{bLqDLOZDLxbO#)bluhd({s)YsG`{u*0a zlCc4rPr^Xy(OS!`R%vtIODlb9Y9%T9GLaoJm!_Wm7eMqrc<9xNro+E~hNZg7uj&8d zu6lG@CF|A(Anf(_?Ma`bFZd%3N585xg`X}8)eJfcOVPerIB4QSjb~Z?zl!_HpsJ&< zT}tJk7?cW#h%dbwK^U?}vBp z%$?PfKf`T{He=Iz^cBO_I%rKIM2qFOSv^tlI@`Zn{21)0k&`&LXGjSYo{ zJKSrQMN9blif2{7rg``=d4yQDOc zq;tb+qMmUMHIeJP2c2aJ7rWWgDL3Mq&qS^x1>13wbL>~gRg26(AwbNZ1pmB;aPo`r z|3Cp1h=;cjq6gjft(n!o_r5+Gj(4Ls8w1dbdK|CWVI~&7c%qmC#tFv*ERykIDb@u0kSK9 z&$0bj-`s?!l;>T$N3RwR5Z_|=+bY8{`HtxOw!Yat#?PoK8d~*etu}9xKImyC#&{io zerm4q9QBtbuQ89T@e`Ks4EYj%R#U|MYjOo0+3!13>U3xG?><(X;RM*2H4~qrd22PB zv8T7UcO5|cpMFk-y&_Xs)xp+)@HPS9IB+Uj8aKD`zG2T$=N^Re?bvG7UO%B3OBsbly%5^{_w^o28PPsv8KNwBF5#q5!h zz9D?w9d*}apABRe9!enF<>?d@6?=e|l!@r?@6W=R5S~9&Dj9lH%Pe81mq)<1I8z`KdoK!J^x8Us&!W~T7AfBnR!hrB( zBfEZA?TN2ZB#C5o83xHFGUM)&Lf>a*ZHoGL?EF%3WQph&8TR1By8cef&tEUGnkF~} z3^~Ac=t=lp#zBykw^)hY_+&^2%)0;994!VF;5TnL@*&Mzb5=Vu5F|`$G-_gZ7LWNW zckX?HLlnZZW*)1>_CMdpgmY@eJj`xHI$u3|BQhfudC`%dBA8QOGm;&XA7^)Ww>T4s z$ADPctbg#9G;XAI^FgzIAk3EyBTt0IHb7_2C1zxp9PaO?D`g)lV+`s7zgE}mqy4?n zLt1KFlfbBDKVd#6^p3*F6C8K%zyOE~@k;;sN5zk&TW5kuQzMTuYV zuuqP>WRA{vg8Zk_$26lyf%3f+f=FBr^PM#~0(~SP{)G65X1{x6XD`Fn?YU|PnAzbS zDG3gj$W z%M}LG1Z7IiU%ac3YhlRFe<0H8jUKLJ$^>M4n*q)ss% zFukW~L65h(fL7f(CX^xE@6$0JJwG+wH!~R~kV9WsOZ)<0JAmLih8+(HqvTZ9J02l% z1uQDe;{V&Ht0L(#gSh8EeP~-$jV`iZ4UOq#wEJg?E{vuyW?&r8FD||~CL@uG`=8BX zbPMPbn)GiwUA3HR;}#>9O;ekH1|#$n!qdwOHs|w|o^X;23sVvja)l$43Y&S77r>W>wJt&kb0NiN|Oj<0%trH zDv=)3ax#jijn=c(QHK;{6ynOda?1Vu5=MA!AFe*ja1H$*E(ppD&YeF`%By?sgcHV6 zt#oIL8M&I<*&n@AqJBmKg5+bf z3}J0;Z8V>QY`jcR#6^m>)1AEjZ1(_uPaW&LHkY1#IInT^wA6a`lS2B_r%#W)M3<$6^cigNZ_uM_Abw|Ap2xnf!csJra%SQJv4Xri$he)t>7$0`y( z?um$a#m2_MLzcKW?p}GcwY$H+gs7-DFt)N0A7#@3H*)Wu2g@WdH8eCd0R8XxxG?t6 z>?)SZgR#}oQOavM5+_tK+Oqcc_QOM0^p2Fq_`-$26OvZPZ3r|(FaM9u_5&IUVq+mh zbz$Fi61Yj}e>QjF)CL=*cMsYcZiFp0`QVeCbl#it$Y}eO!WO_zr(251)*AKnVV*?Y zl9W+awx4Og!9(dUh0IA!^?gJ8vA8&#fnP<9?p8!;ZLMhDNLLDVr?{!>zoR;i%>SRG z`osb$wen|!OI1GZLvO-Rj)+JMa3n*w;UeOzniNq*t&5Sp>6i67r$D3bITckHTqNi& zG4}v~_w{(L_P5l5YsAE29y;!6x2a2nDs)IB(TY+3H_qQ`b2D>G(KPYa;zB_mWzyt_ zhH8u*nNa28+cIq5`jQZe*{GJvRfr*-;8ECNzS_Py~h5V>wdm>3vX zbMz8@sCf7T2M4WU;Rh-fI=?_O3LHA2M2Pd{T6n^Gy5nq}`7{&B z16mtXsZpjQ*Dg+ngoK2Lhl6Lso4yo>0Osxe!12=s>u>1OBu4aVu)7Be)Ak&npa-99 za&oc@PBtyw_=}v{+S)b0ZA|{BW3jh@S;Iv@Uggg^DeEedd!!(HyL$yLh0_wg*v&gK>tBDL$K z*f}|aZ9W$i+zzfswgRKw%%F#hi%Yrqle^@8@Ygt^hVeJ42we1I{SYm0UZM`~i<9L{V%?bOnMnOX=EuWKn zgl0w*>lG332T{+$JtE`UxK=RY;j>F`Wj{H#r8AD38gr<~YBp9ol)!MVMsd#igepM{i}(stbG>h&5(|95t17O8r`tHYcxA@=8X7e|XZja7 zX?go)U>0Chd!o(fASNsO17r)Q+G39y>9aGGbDyn_*IAcq)%bt_&yEm?TM!7x|9GJx z%d|STtpcc7BjqxaFG5f6pnHJPj3fEHELV<(l#C3P-*8vgM|nC6fAps4qEP@-nrH!) zOkO`7GX#AT`%#}>+ZG|glkeCh;nfnGR=Zd}?T|&ev^ib?4uhZKrSZvzRgsKLLg0kG zkm1U>F0AVwQ{^fbk5wyLsATK#=h)ADj5UAcuTkk?$|sXQ+V>&H#qKEW25dAieJ*ZR zL6kn`h$lzLJNVa63BCOGO{=Etd=RPfhI3g|MMC9i_EY!u#2$6c=Ct~T>2!oU^CV<> z<9D(y#`LvQ-&u}S36Igr6FzEkW{VFV&vks^gSTIZ|%tY=!AcFXDd72oogtzqrH!gBPKwZu>R>CrX#Xl>&)Ym<&EZ} z#Po$r>r6$e7+9YR?j9GI8Z|WfFyKSdtYZ7A61~zyfk$^7Xxkk|;k=M=-ET3|Etq$wQ_WB%)gb^>% zQWu_oIzYmVdb};q`MEpStU;wtg1G_*9GA&f&|F4_^l+ZG{dT;ATiVVy&(Ml{gq4%9 zv;0u;y$7d!!u?Xkt#ww1=?37l9XG$@>UOJ1KlPnhrALqa4D!m6quTFG39_iK1S=UB zcv_6tToBYbPYs|dDXDrcOu#q`JxW)!Q!frGn!D+$Nx2#wl7>yP^iOxTp&U{tnZd~&sco`EN{8(1kKN?+037^f41oQPmJt%aAPjeTAIEo>iO z!=6P#0`{N1m@t|&v=3cv{U>_ncH8c-2OT!9)x5dXjqjC8%M$r8Qi$_in9dUz(MMl; z2VJo42TQD8VU02+BJvI3hH$*6#5-A9X>I6d zI6Ntig|T2-Tzn|1<~m6a!-u^|$L(!XzoR9NoxqrPHCG!kBS?-r23I?H%awasMMZXS z6Zrw77zj(v*Pp&fuqu+?+&|JYid#D1)E=VJu$ZM}&w2h$QzkD#Izdd*iyTn9E9Rnsb?V8!~-0$f3+vkjqtq)BNADf?d8vA-` zP?1fG#2a3l+0PccyqptM%jKjwxXyd~n4RQOHn!O73ipTg|wGE$m+5Q-BG)so@H2OHI=L@(_nwzNpG8U#hnFO-SjiI(V-H=PrHb&z-8X!aHRcu3#SDq{1c;h6GB3* zV&rFEQ{jzLd&w8a1-T=7aKdj25$Y9=h7tBKsa`DMjKRNP@+Qb;%lO2iEVPkvb-3ty zBrih|>71FW^CMrkZoh7Ql5pI7O#=2WSU0QH?-Oz`USy&!A9Z@|V!ET1(^Wf+{l{vW z&)6`RkdZ9wm{rj^TVBKicKFQH=N>7B+Zy`*=@0s?w2fT$>!kOKR#a1Y$)vaWDRI2H z?voP2@gy~MbzgunP5LdIxW|mop5IBy9mAaF+NG%=J!uErR*R~u4?&OzFdj%7gSc#x z-Bx`787(a>5QL|4O72AF(7He89^1s9S9D?zCbjQVwailfV~B_h1RH4pxC=lP0*2lTBn>Ss0uPUr173~{vC`yu6eZhxdTdYS zD68>X=FGjnERHW(i;u7E`f-F|wud%(I&Pp9XtJm$Ct8?YI8Mm%^_9KGWG;1(H0jvtW)$~_WP zdT24ii6ey)JBrK7rY9%A>*S%;LQ`o-h6e}FNwRpM`aQq)_NWCoLejTpjFIl(DEw;f14Q6E+TZauQurgffxM? zw61nvBRNTS?NK`Gr}kv;KmXehcVzW|fR#8Rj4(7~WG7j5VCJ92j%U+FS%91P6outY zdS5!Ks^DEUntQszb?8p5Mtq07s<|r0j zTNYVsEMZJ*EWI*1VQz-dv1mC`uYorcRG-%lUT%W04 zlb_ghOmArwzvQQdHG^Y~)$JWw?jmA#fz5TY(ynb$4B=A?M32*F{xqf?(M8b6!qoZ(LphPafVWgcHMOz`9?TdavG`=VreCU8*_jS3j|9yaxXl`}BX>Vv zAo)*YyJQMib5n9({@SW;Gq&9m6%=?rt6k1?%Xmnxal4;zK>AgqdQfurJ#T{S%{hn~USu1xYPs*lP(2YJZkJ$JExv`1k zGwr_(km)7Rn=>#+UA|RxrgF%6P4<1Cp8mfdI|-(8`ug*$!UnMILidvOdt7{&`Kra= z`pRI$xc41ui3Av4?>+63S>Cxiq(=zmC%A-Xjg@MGfJ0Hs@7!+fcmwPUd=+jy+oM*l zW0>bT{?3Do`~gF{m{DNaZqo;4u~)93EhZC3D5Re6VyTmGYu~ZJHJh6lmkMWj|88bR}MpnQ3-oUVc4{6T^f4+eMxz@%L58k)FV;WH~IIK zpWfc*4%rX;UTa+*?^Y&)Zxa6dsZ?*z>$;o52OgP1i$JN)rqaiM`fglUG-8^CJIF9l z%^%AiXr_?QMCd2#jTM727rA(|bnHhAYaV`=6H{hm3C~^{_P+w>+qMi7lNgph*B|Uk z5c#`FBuZE8cVYJEo*1UOPJi!UopZjNZ*oiMp+Rl>2cGJ`brvs}nA2ufO;kx0Vp`mD zZXme-uPxigdlJGoLsHHU@;U|Fel2rsf0!O-glXNHgOC~8=pJAo_qp}&!A9M z-$a31X`Vl<(tfV`O{&KdW+)4~4Hz*nI2E&7HD#5?x@4)igM+R`5&xygUPH$wj-|Jz zqhprKX9N!${}pLmDE;Pfh-fmsxrY1)`>AC<8+V%sbW3J%o>%sn#pEjJk?e0T`rCec zGvlN36;=k?`N7DHxG`m;W~r|;>@XUnZ{4GpxDr2u;|(4FEa_rrobHJ#cj>@hRwz)S zM}Ky_DA9e(NKYRsmn`%s%?=Gs(b#D!Vt$6WiJn322sNWZX}iKA@5La~(T;b_+4ra4 z2t?_@uQBfE4vOWgfA(1NeV&BBzY=I3HBhhouiOrk*xBCAYs1U~?)5HOiZeyM%dh(v z7__`kTxhCZp+5&pIC;8_2IWBY?B;is7r!kMy}&*vzPzB>w(m3#|KY&t%I{0Yv-!L> zG{-}G>?KASf8KxKQzf~qRwz@6@dVXA7GQC$cRrJ&jPT_%b>v@9aHvRz%QyJ;aCugG zM-y}YPQ17%O*>-ciGG=jktG|mQX%H5#*^yG>V}@2^<<8UAqSCM^xQ-Hi4~H3^+Jy< zVt-w-V)C5jc@cd2&VVA)3}`>NL>%xktOM6bn({sQlJEyn*Qom4Ro~q>80nPLdl~HW zuL+}!dSwRK+Jcv@_n75|iw#tFhQRl#JLoOC4?R4o3_a1!p)8JpodT*P{4N~xpSed4 zWGnoJ42F3;h9{kdg?>_0aFO{&#`;^+I0to^t}jP~rqzt!gVarSO|CO&qWcYt{!2>l z&`?ZdBq=fR`=}^hR@PNy5fCauLqgE^*XiA^0{JddWbJ0>qyp(lOoarL+xd0u5;U&; zT{zArwq^WGOj(tc!zLV;iHX(A?N=quBp*JcWnu!*(A><-&SSl@F|#@detG!TXZF;j zcEcWwdmCV#ChUeh9pjY7cTl8DjI# zJJVHF6;kQS3MdDu>anVjx9!?<8**Q4$^!1yX6w?*RHzkpsWW&FIOU3oh8w-AX zFx_Tt??&4%^s@J;Cpqc<2PbeojUl=Dy#vOWnk47N;|R-}1U$I5s>U*7jbf_>*tvs< zZWf)fb`oiu5>)$OL}aS@-J+&sgn=ZTet+(H)`|b{49u*m5)=?Xp(x17U%z_gu&yfj zkeWqaiU0fxfoWXRygzv2QFBQMZPmA}M(&n1?a$@>adxALu~q|%u=fQ~9liXLoSYnZ zl0{{A&Z4a5sb_I)gAG}@6(N9|4CisDNfdg-$=y^~#%Ggkmb3Z z#Y!6R;0^iPFpbL1U~4Qn5%soy;i`X$U$HLew&+ z%D%IOL{Jr-dE~;0bH9+z)E(9$M1vR%Ypenmy2wWsgfY|*BH)g@$OI=Ztt8#mK4kj~Li z#o-)n`Bh+^k@%$R=XE9XA_d-i&F*A{AoK-o))G zMfK)T?_nCg+K$c4!&-q7{R+bI#u(qHTYCRg?R?v{J+<%wRu@I~-+XouoK zijVx(xad*o5`P!lQI9ts9dJCP;$#D{%Xz{W$@fY|Mci&ypq3qvw*L+Y28Iwd=$Tt% z{crxL-eZMpW?5OwwOcy?QD$aKJ3GalFCy)htnKXlWguA&Rab4Tx3Rb7mI9VFRx`j! zAd?UmEp1w2q6jLncV7G1Ge;*UMiMx3jsf0PoigOtb1nc1*8Nzpj?o+zO}Gg9@W7My zeQOyo!p|xMbmPqI?8~oQf6u0rm-9_OVp|@}9asb)6IDo`2@@S%sr_ogS#L`Zbamyj zPoGBSG5He~05zL#nsZ#f`jqfPxYxYQq;`?IN2(VwC3#E0+2Lkwn#e1w8Bi9dt9tn@ z@VZ?V!+Ptr7^$%EaGGp#Qc@o0og64L;<(I?F*!O~+DrNzS(h3O*!+@I$ZgaT39p+5S(Hp`>V zv|jmCRK&^7t`O~0`k1GI{?zCKz%DX=Dp7TrZg-nByw zQ(tVU?h`3#ro-gHJS-&)Z_%^B$xKhTijV;i!47xtNq?`M>k4{mY@-U@Sg6MrR{ zjML0FNo9k~BKYjNPLNXrc0CCEixC5n*Sd82N-M9Os5T*WCphM39|35577tI2GNdVk z0Aez8&m@SYai?^V2KY#=Tt$*zzkYPMw+vGhgv@IZsmAg}fSA?(_C`!p)K|u_Vz?x) z+-PZIW5YLmZ4(h~#;ZRp!<7HXT?*F93}kQkzxBueyzqwM;iuAW?`wQ~JT${@uOIT- z3=g9Yr?KM*F|zZGd$Rzec)_AAkb)l%A78!H+7QGEo;-P5tY>1<3{sE~hE10n$DVt? zcmV;{uL%&TfV%~x!2occQQ1nQUOV^4M;xluQaVlhKN~taMvk2ZS1`rN!jdz0K{0*v zV0iu4FG!N450qIh0y`URbE-lTlJVyM!E7lNnS%))|N?z0#76Iml3gTOpGtxO)_ zMIxe*sHnY9PT~?0^fWY@wV3S)_ua|qX-K!U+?Sv7>C+=r=g`pN^mJHIP!8Nc^@Y_* z!h6+tF}+lKVnV{}`-#O!DynfHf3BSc7sI-j3V9Od78f&}X_hnwc?1M14LYb+$0nP? zF1Jvf-yMHs@o^+i*WZ9(oHw#%iaY8m$Ut6%EEA+zs0zAlxX2|E3@oNJdRbLfC6z*W zgU_pd`1C(;z+r!59^bnVy@OW^EqtgpS z14VKl@k+w6N^-L3%Z8>V!dNvFsuEIOS!h{X%UhqaSjrmA=i-2eet$U!T_+70 z9*=~C1S(3(>YbMg>58v-*{+;pWn~3FX2}dhKXimH-s1Hm$rxkBpS`X`aq2uAfa(a0i#T zfjybW#j?G<-P8mbBP^*k$am@t55PN{4;IIN_Upr^ym?c^*&oOlBjD~Pr6?VJW-zOc z4$BzM@Y*f}^+MXNm)_nqu|0qka%#Hw!>9nR1haYS{SAzJGA`L=bj0p#BjpZq zyDQBRcN6%VmvF{=Hs^ckH2@HuwXoi{`s+I(WGa{x@%lnpscm7gYW#zpjm?+zhU;i+ ztJLJ9)kT!XP4pUPC-`Op$c>J3A$QGzQIR4YTltBQ&v{3NT(^#TV|&|wD~?Zg^f^~; zJSQuw1xU@XvomT`qSac-H#Z$!05c9N@&}^!z~0HJsGxv0R*;*!1aPxphr|wLNM042 z9k;*+7llpK^uG7=An8Z%(J7CeXnZmb-**;}{wapn&M#KxFdBtIW&hR|!+?8)3J3#r zC3_3T35khyZ+8?RGRVPZy>_>h>p@#X11JT{Rj!PF@!DPcku^U7vErbh6}zP|&>Srnz=S)Z zm3TD(A9R0GST`dm-(JgwaqkOcFMkh?&2*r!#i1TD;UF+l*!cMNTeH73?_GoLfq&(J htcU+^{?K1N;*=rVwPMT!x+Vt3Loq4QEMZO0{{_|vaIpXY literal 0 HcmV?d00001 diff --git a/src/main.cpp b/src/main.cpp index 1850161..3129c96 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,141 +13,135 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 8; // feel free to change the size of array +const int SIZE = 1 << 14; // 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 - - 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); - - // 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); - - 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); - - /* 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"); - 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); - - printf("\n"); - printf("*****************************\n"); - printf("** STREAM COMPACTION TESTS **\n"); - printf("*****************************\n"); - - // 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); - - system("pause"); // stop Win32 console from closing on exit + + + 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); + + // 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); + + 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); + + printf("\n"); + printf("*****************************\n"); + printf("** STREAM COMPACTION TESTS **\n"); + printf("*****************************\n"); + + // 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); + + system("pause"); // stop Win32 console from closing on exit delete[] a; delete[] b; delete[] c; diff --git a/src/testing_helpers.hpp b/src/testing_helpers.hpp index 46337ab..37a3072 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -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 @@ -72,5 +72,7 @@ void printArray(int n, int *a, bool abridged = false) { template void printElapsedTime(T time, std::string note = "") { - std::cout << " elapsed time: " << time << "ms " << note << std::endl; + + std::cout << time << std::endl; + //std::cout << " elapsed time: " << time << "ms " << note << 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..62dca77 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -1,15 +1,15 @@ #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; } /** @@ -18,9 +18,19 @@ namespace StreamCompaction { * (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 isTiming = true; + try { + timer().startCpuTimer(); + } catch (std::exception &) { + isTiming = false; + } + odata[0] = 0; + for (int i = 1; i < n; ++i) { + odata[i] = odata[i - 1] + idata[i - 1]; + } + if (isTiming) { + timer().endCpuTimer(); + } } /** @@ -30,9 +40,15 @@ namespace StreamCompaction { */ int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + int count = 0; + for (int i = 0; i < n; ++i) { + if (idata[i] != 0) { + odata[count] = idata[i]; + count++; + } + } timer().endCpuTimer(); - return -1; + return count; } /** @@ -42,9 +58,23 @@ namespace StreamCompaction { */ int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + int *tmpdata = new int[n]; + int *scandata = new int[n]; + for (int i = 0; i < n; ++i) { + tmpdata[i] = idata[i] == 0 ? 0 : 1; + } + scan(n, scandata, tmpdata); + int count = 0; + for (int i = 0; i < n; ++i) { + if (tmpdata[i] != 0) { + odata[scandata[i]] = idata[i]; + count++; + } + } + delete tmpdata; + delete scandata; timer().endCpuTimer(); - return -1; + return count; } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 36c5ef2..c0581a8 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -5,22 +5,104 @@ namespace StreamCompaction { namespace Efficient { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } + + __global__ + void kernUpSweep(int n, int d, int *data, int offset_1, int offset_2) { + int k = threadIdx.x + (blockIdx.x * blockDim.x); + if (k % offset_1 != 0) { return; } + if (k > n) { return; } + data[k + offset_1 - 1] += data[k + offset_2 - 1]; + if (k == n - 1) { data[k] = 0; } + } + + __global__ + void kernDownSweep(int n, int d, int *data, int offset_1, int offset_2) { + int k = threadIdx.x + (blockIdx.x * blockDim.x); + if (k % offset_1 != 0) { return; } + if (k > n) { return; } + int t = data[k + offset_2 - 1]; + data[k + offset_2 - 1] = data[k + offset_1 - 1]; + data[k + offset_1 - 1] += t; + } + + __global__ + void kernZeroCorrect(int n, int *data) { + int k = threadIdx.x + (blockIdx.x * blockDim.x); + if (k > n) { return; } + data[k] -= data[0]; + } /** * 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 paddedSize = 1 << ilog2ceil(n); + + int *idataPadded = new int[paddedSize]; + for (int i = 0; i < paddedSize; ++i) { + idataPadded[i] = i < n ? idata[i] : 0; + } + + int blockSize = 128; + dim3 blocksPerGrid((paddedSize + blockSize - 1) / blockSize); + dim3 threadsPerBlock(blockSize); + + int *dv_data; + cudaMalloc((void **) &dv_data, paddedSize * sizeof(int)); + checkCUDAError("cudaMalloc dv_data failed!"); + + cudaMemcpy(dv_data, idataPadded, paddedSize * sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy to dv_data failed!"); + + bool end = true; + try { + timer().startGpuTimer(); + } catch (std::exception &) { + end = false; + } + + for (int i = 0; i < ilog2ceil(n); ++i) { + kernUpSweep << > > (paddedSize, i, dv_data, 1 << (i + 1), 1 << i); + } + + // set root to 0 + int z = 0; + cudaMemcpy(dv_data + n - 1, &z, sizeof(int), cudaMemcpyHostToDevice); + + for (int i = ilog2ceil(n) - 1; i >= 0; i--) { + kernDownSweep << > > (paddedSize, i, dv_data, 1 << (i + 1), 1 << i); + } + + if (end) { timer().endGpuTimer(); } + kernZeroCorrect << > > (paddedSize, dv_data); + cudaMemcpy(odata, dv_data, n * sizeof(int), cudaMemcpyDeviceToHost); + + delete idataPadded; + cudaFree(dv_data); } + __global__ + void kernMapToBoolean(int n, int *odata, int *idata) { + int idx = threadIdx.x + (blockIdx.x * blockDim.x); + if (idx > n) { return; } + odata[idx] = idata[idx] == 0 ? 0 : 1; + } + + __global__ + void kernScatter(int n, int *odata, int *bdata, int *scandata, int *idata) { + int idx = threadIdx.x + (blockIdx.x * blockDim.x); + if (idx > n) { return; } + if (bdata[idx] == 1) { + odata[scandata[idx]] = idata[idx]; + } + } + /** * Performs stream compaction on idata, storing the result into odata. * All zeroes are discarded. @@ -31,10 +113,60 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { + + int blockSize = 128; + dim3 blocksPerGrid((n + blockSize - 1) / blockSize); + dim3 threadsPerBlock(blockSize); + + int *dv_bdata, *dv_scandata, *dv_idata, *dv_data; + cudaMalloc((void **) &dv_bdata, n * sizeof(int)); + checkCUDAError("cudaMalloc dv_bdata failed!"); + + cudaMalloc((void **) &dv_scandata, n * sizeof(int)); + checkCUDAError("cudaMalloc dv_scandata failed!"); + + cudaMalloc((void **) &dv_data, n * sizeof(int)); + checkCUDAError("cudaMalloc dv_data failed!"); + + cudaMalloc((void **) &dv_idata, n * sizeof(int)); + checkCUDAError("cudaMalloc dv_idata failed!"); + + cudaMemcpy(dv_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy to dv_idata failed!"); + + int *cpu_bdata, *cpu_scandata; + cpu_bdata = new int[n]; + cpu_scandata = new int[n]; + timer().startGpuTimer(); - // TODO + + kernMapToBoolean << > > (n, dv_bdata, dv_idata); + + cudaMemcpy(cpu_bdata, dv_bdata, n * sizeof(int), cudaMemcpyDeviceToHost); + + int count = 0; + for (int i = 0; i < n; ++i) { + if (cpu_bdata[i] == 1) { count++; } + } + + scan(n, cpu_scandata, cpu_bdata); + + cudaMemcpy(dv_scandata, cpu_scandata, n * sizeof(int), cudaMemcpyHostToDevice); + + kernScatter<<>>(n, dv_data, dv_bdata, dv_scandata, dv_idata); + timer().endGpuTimer(); - return -1; + + cudaMemcpy(odata, dv_data, count * sizeof(int), cudaMemcpyDeviceToHost); + + delete(cpu_bdata); + delete(cpu_scandata); + cudaFree(dv_bdata); + cudaFree(dv_scandata); + cudaFree(dv_idata); + cudaFree(dv_data); + + return count; } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 9218f8e..0a37415 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -4,22 +4,60 @@ #include "naive.h" namespace StreamCompaction { - namespace Naive { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; - } - // TODO: __global__ + namespace Naive { + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() { + static PerformanceTimer timer; + return timer; + } + + __global__ + void kernNaiveScanIteration(int n, int d, int *o, const int *i) { + int k = threadIdx.x + (blockIdx.x * blockDim.x); + if (k >= n) { return; } + int offset = 1 << (d - 1); + o[k] = k >= offset ? i[k - offset] + i[k] : i[k]; + } + + __global__ + void kernShiftRight(int n, int *o, int *i) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { return; } + o[index] = index == 0 ? 0 : i[index - 1]; + } /** * 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 blockSize = 128; + dim3 blocksPerGrid((n + blockSize - 1) / blockSize); + dim3 threadsPerBlock(blockSize); + + int *dv_idata, *dv_odata; + cudaMalloc((void **) &dv_idata, n * sizeof(int)); + checkCUDAError("cudaMalloc dv_idata failed!"); + + cudaMalloc((void **) &dv_odata, n * sizeof(int)); + checkCUDAError("cudaMalloc dv_odata failed!"); + + cudaMemcpy(dv_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy to dv_idata failed!"); + + timer().startGpuTimer(); + + for (int d = 1; d <= ilog2ceil(n); ++d) { + kernNaiveScanIteration << > > (n, d, dv_odata, dv_idata); + std::swap(dv_idata, dv_odata); + } + kernShiftRight << > > (n, dv_odata, dv_idata); + + timer().endGpuTimer(); + + cudaMemcpy(odata, dv_odata, n * sizeof(int), cudaMemcpyDeviceToHost); + + cudaFree(dv_idata); + cudaFree(dv_odata); } } } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 36b732d..536e7cb 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -8,21 +8,31 @@ 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(); + thrust::device_vector dv_in(n); + thrust::copy(idata, idata + n, dv_in.begin()); + + thrust::device_vector dv_out(n); + + timer().startGpuTimer(); + + thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + + timer().endGpuTimer(); + + thrust::copy(dv_out.begin(), dv_out.end(), odata); // 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()); - timer().endGpuTimer(); } } }