From 5d5f7cca75115b1a47255120e4ca1236f01a2828 Mon Sep 17 00:00:00 2001 From: Andrzej Janik Date: Sun, 14 Apr 2024 02:39:34 +0200 Subject: [PATCH] Rewrite surface implementation to more accurately support unofficial CUDA semantics (#203) This fixes black screen in some CompuBench tests (TV-L1 Optical Flow) and other apps that use CUDA surfaces incorrectly --- README.md | 4 - ptx/lib/zluda_ptx_impl.bc | Bin 144764 -> 232076 bytes ptx/lib/zluda_ptx_impl.cpp | 662 +++++++++++++++++++++++++++---------- ptx/src/translate.rs | 4 +- zluda/src/cuda.rs | 2 +- zluda/src/impl/surface.rs | 136 +++----- zluda/tests/kernel_suld.rs | 4 - zluda/tests/kernel_sust.rs | 12 +- 8 files changed, 547 insertions(+), 277 deletions(-) diff --git a/README.md b/README.md index 52927d0b..5be6a8ad 100644 --- a/README.md +++ b/README.md @@ -215,10 +215,6 @@ Performance is currently much lower than the native HIP backend, see the discuss This is a ROCm/HIP bug. Currently, CompuBench tests have to be run one at a time. -- Some tests output black screen. - - This is due to a bug (or an unintended hardware feature) in CompuBench that just happens to work on NVIDIA GPUs. - #### V-Ray Benchmark - Currently, ZLUDA crashes when running V-Ray benchmark. Nonetheless, certain "lucky" older combinations of ZLUDA and ROCm/HIP are known to run V-Ray Benchmark successfully. diff --git a/ptx/lib/zluda_ptx_impl.bc b/ptx/lib/zluda_ptx_impl.bc index 48ea22b8902095dd36ce79adfedb42c3358515e1..1edcbd5cc43abb641e28b19d37651f905d1e0986 100644 GIT binary patch literal 232076 zcmeEP4O~-I|G#&4+)Wr`DCowU7w`p1bbulV>IV3dSr3{PnlBq~BB4%1Mbq3EZyM?} zE%ebEsimY(SXSC&wE-#;7Mhlt+AK9BEiyIV{r}D#gWWSA?bYZ1X`hfu~TXbvZ6qnK3Q6!Z4iwI)VHiB^B`4c?2?hOA6 zbb%|C)%-_kNm$e>GEdaqNhtH-6ae`JVg;iSU{08lQfL*f23PU zkbighlJykjZ!)M}OPR=Y%gAL;o)=0Mkt&HlL%VBL%L$Q3u2z|so8Zi&-Ni|K6TgFj zH=A#i7ZDz5BeXiYn(QQ3NF>DTDlOST%4mDYT_};nL{K6*`jRL-i+Qz#^bAiFr1vKR zCzg}5RD4kc<5sL834X4-oY2P0D`-N}VK(33-bb!`Le8UmBsdQ{X5vpU3%f`ZAySg! z)jI`dDu}(RJQt}^R;(qdv3mCim3D@lM|4RL50jtaoiGde6BUD`I-N_sAdt|KPHKK9 zxzWE^LC|CK+@&V%a%k+G&~cbtO--l}@-r2WLu1$a4gn?+KVC^Z8d{{_r!iV*-j=0G zev+IbpOz>RMS+q$u|)gtenVHwZwM}H1az zex;C9(;+6IzDlm|EXq@+hywW~vYkYwXkevif?P#!4lP!w(h{j+d68@QHnT!{wm|j; zROeQTo`vcIP(3lt2-RD};dM}*T_~%E>Y7SXHdJ2A!KO#Q3UV?~sU0npk<0KQRfH))UpF|Hd)$9Fa@r*F8&i&DDxrT0O zM!jxuQOEE)vVcihBCA&^Bo$Sn3EEPpX-Ra5v@SO?RNq;(P5DZR(uMwlNvJKSr4tN6 zX^DC=q{t(D8@bObWvT27mC7}CuP7UC+nOXC7_cP8H3=@|MyK7V&qTc{l=2mAQ@$;Y zb&(4hZF*>4XHX~Nc#<&0#1K@)UdaSG$ad2U{$CJ!f>+cfZkricYv@~$NQroIMkU%o zQ+$=8LYgZ^1d2lWzKDPSSi%WN=kj@70kTaIw+CEqzOT| zBmBbiK)CUr$#aPkv!A@e_*|{ZjZ!BJt`z3VO4ME9hUaU`p~NsGtsZU=GT~R?hHiCo z9$BJe;D+2y~RtocECI0?!Lw>CRN+fY z=#+xF(t3UGusEseic-4PUqs(zmO;Iz zwH|yHfh75LU1L1(o=2l639gi4T)w(~11b3s0$kRrH*+E-TUfjk*<(O>8a%$aa3uA>(s zf2=U5A0ofQZG1XQ-mSt)zU!t^Sz5Bf*71>^NN0OIE< z#mCc&=AtGQO-(Ktmb2Dc!T4F591`D1FOs7syu(&1YJ;Ib2=-XP1X!E6)3Ry8&rp-t zrY5hfTvY&!fF?bxP0q-orU{>8u{1TgW>{T-jIDwRwl-0_jh|Muh{e*>nkg_umZTrjIlO3BU_t+1OQl? znp`uKu>!csJY{X7cH5nS1OQl?n%pvMV+C-NnP6?Q!|iwm5&&RnYO=`iA65W2naS2B zhXgk=kN^P77gj9qtSn~*aFa>0HVO5XJ&ObYSelyfSH8^(;3ku1ZK9M!J&ObYSelx2 zSh#^*hVt6;cWPgD%aMcJov*KX6nf8`c7bLQmD_wUC>$XxD+s!*^*@%Gg`cV-) zAX+m_Wl@irFCm#xD-~4l@J9K&F-{sj&(ANuF(1|-opaTEp>&v{RwX6Kr7M;4KR733pa%~@jP*xWXMJ3 zIw?=<79q(adV60tUPrZLBBcm8dB*pI*6*ypkA?I*Z1(-M)b)hFPd({NEHD%<2oh*Z z1Pg?Pvl0s>$DXbp8Ahf<-|B!cFVCMRgkyCWP8v z(65PX6lnPgims9$8pYEUZkHwEM#AlaMf_bIeMus|q7@T^s^bQ&rWCc*puEb+ZE@Zv zeas$e(CW&_y2|M4IPapkK_&XA)%uv7=9o?T!K*2)z(hY8c^CRvOz>5%#FNBr!az9J zQx(3aOI_=geifvj#ozbj8J`nI-wK^y9cdKI-%^-G6wNmlYNUm8(y1In;XGqumQ=7H z2f$Oawg_g+3$tKEC572_(6W&@rf{xC!*DZs5w8>1fqNcxg1qQGao%=R6gO~J zWz_;MZnO@iBbmG~UvHeK#^ z!3ZraWVvnxkN8RvL^8B@krPqw%BP5=a;3*Q)N5BFiWo&UQiPaTqA8r0PJtjaT0zbT z!Mw=Ai1b2@QIJi*NE=Fhh$O#yf1iD7-xGx2*)zWNr9S&~ey5dwClo#nW}n(ppR=W| z=ly-JK+6~+LYw*opu z3lsv-v{Jtsk~y#Qt5ExVQR-?@`kmJK)u?^y&-en`s#3pNvri2SLFIQ?*|ml=Q7S#@ zWa5qW9_4=%^#Dq1@*su~cmD*rYe%3Sh7cVUU@uCK*U?p1c0`Sof_GEsTl7R+P4k4(#`jWniyF)YTAHnF)cZD?32NuFB}0Bp6_%i2~z_ULWVZGcK}@8c1AW z+^!~wiBPvI67jd?;^~Cj_YsikVIdBtnlVG(0d^Md>Y6 zPs*K4ZrUM4w?EYLsB#c(&_Ga$)jy)&0`PJjjM~A+6j~k8Xr{Z1tlsYg!Rn-n3ZO{? z;i3?D8$u;oexf#mQqp{rz`HoC*~y>%+d%qt;~zw4B%s$fJ_dBzfnv{P<>5SY)O{d%$xfa(>xifQ(Ix{~&RHYJC zt8yE=cgp+MT4v!V{4>ct_^EP#xxPb%%xKPITzd#p%O984C|n`y37B8gqzkel3ukEx zBZP%>8SA$oyQENaOw#2X_-LbGu0pUtCCG-~67~mM0=k}M6wGD>Fu7$xm7M*-mOxdu z{ap#1s=0P|T>!FHYAUx`D!~GYU~V0KjX{M#dAiE2nf^6}Ajj5nKHD-&iyF)_2mnS_ zLeq_Npo$T|H&o4$pepnYG$jHJOLi6^$Wox%EJBd2MU6E;xsLt;Mq#BdcR^iCoafZh zZ%72wE#eDC!93#?t;AoJQaxg&yCN`8y*aa#wk&@A`{5;1f6D*vxIZ~wzHHJMf`H(S z;LFP6YO07xf+6?{xE6+~3U>-p6Go*B0FLvaL`h#Xm~^7$TQ3&y2rU#a)mq<ah)^98Ni~GrlcnC&<7@w3O7TqgJ&tS8k^i|9sgOm%lDgsMK@11Jp8Vu20K>MA8n+4jIPE(Xg%I;6cEPcKTs zfqnH(6%OpuFH3M>iSlxXW=Yb84%8Avy+g51kA~2%;@E2ZsK@Xj>|JFZj*oQ0Czb#T zW5|-U#+5j{itt_4_$aQ+7N+2X^>*zOhQV9&Z@Q}tV_n%e{R=p}q5mPJ>!IHc_VOAL zLFmH?Z7*Ol`dFoZoK|UIAl>an7I2CXeM%6og95NESp_)jtyOwqT`$d}9&yzY9N})R zB&D11Azz6sKa3B#+{cli^DB-7QxCU7uy&v$!G*qBvjf8S_0Qr6`+rSJx8VrWVs?Fj zBh2t~BslV@BSFs}S|J!a*pXmlSFO^UFvSDM(8p?$KT`5^iYO66@94w&iZjm>^mLfN z?5#5RwwuZ5N~i)jHm>JM_rO^A_JVFqAoK$W%*^{Zs~SPj99LE`oD`h~L312qWG}7K z2M5{LNnQjXIll4C>zci!grEt&7I0f#=;TW{XFRrU^r*uSTLT@N8=!exTrJ$lEd(ZX{6`}oZ*W?v}Rv|>!{1# z&vz3MsxkJm-urc|6RjY6gQ@a>73}o4bk-RP#^PbGJm%dAI*^MC!Pi1Ft6WA-x^4pI zII(v&x%9Jr#;ORu3A`>tzz0o~Wed9uhUcM=a$LlxE}l|?EVmb5*`@q{Il#M_+Z~73 zaq*c>IJ|s|Y6%W+hUR(|4sUK`ViE=~_$i;kp%}cQ1H~`l@PfM)uE60ry;M0AhxfrJ zX996}$Iok4;_yz)zFv*P>pgH{3I?w%)OT@<`}cYA>~O-xKsy zFnNj~kuDJ^~-|hd-${ z<3k=fzZHT}ZVm)9UTuZor7MmEzga-&`{M{R-M+*TI)5mTC*lZ8uWPpB2%lf*Nbu;Z zjs$DIXocYMTaE-LJWuE&a4Y`GO;rbOlM`!Y+#y9La949_h@)>WbLt?^fw6FW`@4`x zor$}er@EOe;M+Oz+tKBG*0*z9&4g8s4tCr#9UUC((taHrJ@(3;vm8A3wHF-;9$oH8 zaO4xM5IhjFY+QGVDz?|C_uukerCL_t=+bee(&ZgqOYT zUyV|q1N9aYd@WeM)nlh^xPBRC3YtK^zR6=>ej;%`!PkO|UKalKZ3MyJob+r91(u*+bx!L=io8^Rl zFm5k>JFDs(rsn-T(M{Z5294J$yJBHRSzj0Vb=+RArjhhwEEX!e+uLLTkepcP=td{$ zZ9Epbx$}A~iLr6dgt4UrmyKmPAs;1RZ1k+{JT3@lW9AEeJuoQ`{a~o4GbZJlQ{T8a zVNyObb96orlk$XJgnl&E?>O%+U*P>_tag!0ae~WctC!&fFYj)J;1iAn8@9JX@WZ)| z1Z)0F=*M6Y-MA%PR0+;lY!m!5Pc#x(%?ZFZPu41X;jzt07kMO%g%eD!SxC}v;8K3O zugL=3=6J<%H=U&Gc~}5e^YQgDxRkr4kKKw(*{9z}V=-=f`gx5D#-;4Jtgj~~<-N;? zt-+=2{OGrj;X=`E9=#bCO2c77KMw16+;p!m@P0D}xylo9g4?&Kx8npWjizC5? zr&=L+>P<(26JQ6^Q@E7B@>11-l(}=?BuX>^m-3}&waQ0uDd)S&b8soYyMv_P!lnFF zgvkQj=1BSIr?hkrF69X)kW1_#w0DV_(hLXv4c9qxT z{&0vIP|qbrTl z_6~RoJSSJmzV58&X2ip6LN!;shHaS|NC_hqRF_FQ;W4` z2dBN-eSiT#_OSPL2gARi--aj=d@SplSRY$H?4+VIp;~9}Jg=T8@~#Guod`bLkoAf; zy=PJ53GFm{uXuD_;0(rvkUwtk6>H{HxuyU}j#uotVwbKrVd}^6ia)dha@mkpKuSlo z0)ASQCkZYdkF+BkXIq={n*9;uwm9;ySAsEQuQ7=Yp!8y4*<1qa$&dEiO zpD^N_%zy7u?;|)TGX}H*QrOcG4_6W7fjqv z-#+SX!9m`PnMO^-Ir%})=~|qVAKcjF3LEgb;@9G%PzW_5D3y-qrh-ivgyhqL`mJ+q9}$#{-VKJCdGPLRFfY~RM}WG5$k z!`U;A)ydcF4QI|rOIe-tus58H#8yBCJ>m#*>C{$04%yxc$amMUI>`Zf;eA#oxgf)z zW_6MSa_ILXS)Js7jGfHtB!`o0zIue!Ne(AZEZU3=hs(*JU93)WIVpbG5#*9@TLJm< zL`RT&Uw+IH%;0K2c*~0%}6J?AR9kmb&?Bm>dUO*a5%a32UaJ!oV+!O z)k!WV_dUkyB!`np^IHKqq^~2$3E#H@a@U8gfUJ5sR3acsg6(y(^hO1JmLOiV*U8`s zQ}$VW2vr$JCx3MD-b4|jI@{~yzUO7sESPcZ4QI`TzL^R!As=sVIIHfZiarJF<`_=@ z>T=~k!jxifIKdukLsdA)m4THsq{+D;OH%g14ihfO{x6HXt8tLezbMl$z(E#;Woq0o zkXcWri7w+Hchr=d2V)@1rmPKx0GZ24T}UO(kQm4>QunF4VIZ%5B=TO($2d8+L`E&f zL3Rz#BsyXs$4pHV4dY>)3^`b?d=vv2{MOpg$PO6DpCT&hUN}gPsryvjF_7<`7kTf* zLC)A9)4z;^{3!h8Q3NTsPc8+Ymqi$%>)!U7c6CrV9YGRLax|@R>M^1xp?ce1)Ar^b zl-5Jnp7xq{x?D}YLTC@zYueF15y_TLgnXR6rd{4&;A-L#JzVTvbjI;w-4hVO+G|=_ zfA0tj4)Wa8aC$Hf(r@}Pc*?>BnO}BL3VZapAisE5O}&kSJQ$Iz>WqOr{8fQ#9S-vR zSH-#s7|7uKf4>l-C8WR~6hpC(%boVQdhdd+IhL5XTSo1|L~}HHKqk`#)1MO;3Pkm| z{wzMVQ#T3IpQRsqhnZY3{W0f!PA|apr)==dYHc4(e>{J9-1`^~a^gq+)E*4vnv(;j z8*z{~?-Yox;~)d;cj~5KAf3g&VT3CN^3ZFa)34zm4cXP&ei+DcFNu3}!ZwgJJobj` zC<@b`h6@HV0~1Zn&CfJ`#WE9(>5t`$Y1DB{ zf5!bdFatAp7bEpVJ>uo;hh0&w zS{x)Lc`C@)HE-n__vw&KCJ=LI!FLKp>3nQLyl`))ZW-n)YH0s36O5V@|6K4tLJ#C| zXYI1td9~W%m_i<1Rp)(Oh$-atjsqz<$&G)G=EQ249vH}Vqw7W0ILHOgX5Cs0 z;fhWeNOSKibOsLcVSdFln!(VEyi_?(>R0{8B^0m796D9HQziR zpWpve`a=vhE1f$b>E07F!;UkHTwx~X&c%JiD|H(%g&g^5X@m(z&GDso2FBA1aD`l1 zctZOxOd*|qG03jNs5uIm*>wSxKym|yaUF(DWr8uQNvafzM>t_1_eeMBwqhVNx&=lm zJTZ_5Q&Z?N9OTL3)7o(u$o$t1$#ywoAW!$pr4liaN8cJY)rdzmVaJ8y96Vr{>%CI9 z4FfrHUSMdXgX^1lb_4GEvLnk>e9R1I`n*eS#}snu%NxVWVbmO}$b`qfHrNI!97mH3!})D+BW#;qH!A^KoeQ*%}5TIGb9wtv^Xya?Qw*vxvSJ$P^u(LE^cToWAg@;5{(26D|0hh+%p_qz++*{E}R*i$S%$3UKY z>4;1VkpRbVGF}|0--Ux*JVv7!fN?VENoVmz9ORyn8_3fb$c&V&;dMC3tsz(Go*2lo zlbcU#XJR1pD^|-^;~-C#gi#;kAVZ$im>@{wq-z4EIg1qnOx)RzZ_q8pLB78=T;hv~ zdr72|z#*p;^@c^8je+d`!x6_wa?C3Osn2ncC!cf-{JvSV)HMMI8I-t8_W};m_07-( z9OUk#qjUt$$wk{vXy;%c^I!6m)#02>d?$eV5(l~Psp&>MI&C;q<_h^NuHlqTSV=C& zK-NrpB+P(=Jdk;m-brEB{qgn_rnwl%y&rkXEMiP2Rl5SHQ#eT9r)OGlkY3+4xaQy> z!zQgIS70EW-+Cmx5eIo=u9F}M<79A`w=Bd$4CJ}*j>@!noL2MUK}1f}!{uEuoxD)$B!EZf96uR6XQxH62m`tK#V=&#ZWt$H zKOac_f`eQzVWtrW`N8P#T`f4s*E80V>oAZ@CxwI?aF7SKISFRqoILcdMf*GkvirvX z66a*rv4PYr9AwZlj&XkOR2MOvDX()ayi+Qk^i z)P^slCLHA9Q(n|P9ONg@OxO0rTu{w(Z@Mak7|8Q8TUk3e^$nBwW8&U@<|w@o16ej{ z&k1cF1~PwD7g;?H^4R47suPd9tnV7O2U-5%X`%oF>Ga>19Tzwke+i>xI7qiC8U-wHIL63NOb&ZoTFt*n)u!zVnfVD8xW+KL4c*_Hc2h#iTm}DPJ7qM^iMWI1FU$Ef?`Z z9OSOW8%Qa}$%cDxg-Zf4aj*5iMsLGFmQAQStu4kt4qa6u+k%5U(lMM0$3bQ$X^2NK zki5|%F-!$q!_h3=pnDVJWW!TC!V_?iHy)t{F<2B7{Qh2xb{Pip+<%YB91`ForP0&? z9He*hbYeILvLscIRk7X-KUyPF@w}*z6;~;lOoTqd6n4i46@1$Y{26FTHL|LOd zChp0;k5EH!kRK+`AVy*!J3d$I8s7oqWci9UWG@V4+1(vs@HP@R=)IsiPjA6Nj;%eZ zSc8E~T`^u(-Vp;iH|!DWX&hvFnud4`1DX4oNSuO$RIS@UhF~BY#=a9?j)UC&Y-hnP z3}o4>2ftCgjDd9eugt3v2iZ6vj+%mllun&RJb{5+dak#)3I`dmp_GimK#trV-cQ~W z1Nqu5k@zeQ^6aY{ z$Uzv$45xR)%N;gT{H^^M{! z4CLOAyuBK6km)&b)W6p#@to zkilKQwkY1kK%Q#=PIksP*{~>@`Vt5E8Gw8S2YDw%oPvYA`*A5b9S13j>R0XnGN6lK z7Y4Fy@!4+_?_(gHUiR^7#6fCajHABBK^A5C=8&dr`v9zLc+|)}TEc`LFR&kYA=>ke z^-@U4jgk@Q*{bb5i#3%V4AGA*fG_mY_x4BT(dV6rCu#OMR?~YHYtL0{_YkTYcqWTJ z?33Bl_a}Ol-~^w4Igq-{)ApjYY(sYY@loOIX{yPDJk|cmbyD&SQMC*4y2$>?^;@;2 zy7{DOwf$HcUWs!1kWihl?|R(lKCc5oPT-gE?59b6vCP}Dh~e}%z-{|CNy~mdANC86 zn=1D#n*HJelH@$r&)Bzn1`kV1uAe>L343re_PfaI7`1V`XS*-NsMVZ~dg3EY#NB`P z5lq2Gk@Z}k+`ibeygP&T9{@gco`BDN^%LM8H#WQSS=3n!Ua`mDCw?AnFcw&3u zMjxML0lgB~=~dZ(FK)ZT&EV3Vn*OyP|vGE>#Ym#S@0(fBehUu!~Gn_-Q zqYTzefBz1M1Ku!Y5wed~74Lt07nGUvhUuQK9`~pDTF^ZPHDn+2{c`Dpdpzwl`}dVR zUFMA|B1sizOG$CG7)HUP?8kd`)&M2;$SL>Sr+U2E?%CfhlVWC`_`#dFndi^xLU*UR zGkEZ!A2xNOFrA-1Pi;uTw6yHUCnf+}I1jrwqWxrZYcuL4M-`XUEHj^IYeXq=ZK7V?S`@MZk?u193B7vS`0aDj|}9$qZ_TSRMTS zk1;^-fIUIa@8^k%Ns@u%j#3UpnImHFo+;ak_};5ys)_ z^Up0%09RLXxY}^C6Xgx(a&Vx|`Q9%-8PoO&^ZJT1K~fx=o}WC|X*uR>8@?M9S_dqx z<*>B$$I+2A-(;^OWhZ|5$eq$&;B-Co-syfx3c_fWCK40q=M#-@hGBl-C(F~Rn7de7 z`BsgVkb`<4t+CFMOQ)Q}tUII2$2Hg?DxNXVuMzOMzGBtwvH=9o1PpB660qip=Z|_5 zqzZ(B95Aat4Iv*lz#$KuPKU0WG7etk;kdGzgKw2l*bsJq_w-SCk%x1L!LTcD{s7b| zf~SBk2v!>-Z;qQsVDKXDlttt4wzLB8DMxsF|1;$(jlpw1?g;PrMMrpFTzj(thiAOy z2=7EI@NPbT)EDR8;*Wdo#khB_qU-7;4Bq84ZX@xDVE0eIYy!8#;oi;b&rJ|v+}n0r z8A5Zp7d-jyE1fYNI681~=M|VH2XFfLll@xmyqy1e*Y4myRUjPe2fSj^_dj-ou;H++ zh(PBBZM-w60gnfc-wn8iMKwor7O&1BNdw2te(>XuWH_cRYp?&Jm0{z}xc%l~e1A*E z-SJ77p#>-0>m7#8ENiweerYW>-leBLexVi{Z~x0*r>0@9;P8#N4uLq1aK=0J-uQHy zukdOvgvg=n$03g$)Dj~oTOkT@58nL}MJ#Gt8G+`D#(Qr^07BR1!Z1aWUZH*!Ucl)s z84IOhesEg9`zS+khaSWC@#CBor;Po+T+=WxVcLHv_fM`}*>!T=8wXMoat8PP&fDso|I_F!H!_rSKVt$_b910EX9rcJ<=qJ)lZ;ipDCc*dH2Ms3a=tx4E0mH! z{z*oUQ!qZlk@RNW6bx|(z?M$VKOyx!DypK<}`9*jB(3; zs^$t1`E&cBh#hwTh#gXbFv8Hog<7AgGd^h^JeMvQkY1Pt-`YYA96YHK%ncIEz3v2L=V=S4z)_&m!mN_Q`6asEaH!`G z7V%Y+_@bOjsz!%RT92IcIjQ!k@%O3H`PQiEFAQ-})p1cJ?Dww9$klO?MO5_a%D&L) z?cC@iRCIZzX#ccAO=4k|MxUK7$Waw$OAE7&g|bhdbo&909!wCQHK3yh&oJVPFbXEf z_oUgcLg{zd-{+*x@1(^CPQGj){nDhPBg^9kuGdG`CUk2Wh^la2jv%LG^r$NaH^c$I z4ECf+f8Wz)-{(mcNB#|;98+^&c@8y8&!mzBkPonfSf^8gQ= zUKu3FD$z~+RwBM=5}!*D|11wQA6CM-n_Z5b@jVRyM#!Se^@Dcmz2PX+;<%{V%9vf| zsOm}>V@&OcowI}jKsAd!OI7yOKJi7B+a(J+EA*V)4GssrmLQ%1r!C6GaQU-^hC@=% zCWx=g1HZYb_NmbM9y`;k|! zqHvz3P_}=7+hw`P63IGb zg^geWb;T=J9*Q?gf>L5a!(x{k{5M{g_oljs#;izHuY29|v2$aUkH<5~nd%u+nYpQ1 zv!~6R!GsRzBkv<;g3`0IHM0ly?K^c^cKV#8KFKqu_Z_27Nu4?4Uz*ezaWkiD=47Y# z&CHxTJw%f=Gb1%Qo5`A)9MZqffIi_&Sa^737?T#B8akk#Iy5CLG9t8pWU@RWwO?{_ zWY~a+@Q8?{(3E~j$@1i=(0;w9&B&fPea1A+%0l?Oxz=nFk@qt zO4LoC%$YMaeWuTuoelj>%}Sj;d&-2dPliThr{-kO$x0oQnwFZCIwLuiEsac{rWtFi zP^UaTGleY=led%)o0gfKn#HM!OIOdBk(xPdX4Z6dc6{pWWJHU-hg0*|v>7RBnKRYd z>KXGH%}hX_n#K0oAE+y{vefgZSV<&hsb?goPnnlGZEAY96>fjTxoMf|Y+yC&yC4jN z1|NB@Y4s@~FL`dDu=e;eMxH((*R%-=Hll(PP>=})JE7oPRPYrPXrQ193IgXr!AU5{ zhJqtdpoW5~Q&5lt1!oZr)Z+9qMzU}~uHp)!feOBc7Di~1cN%CcfP#7`*aQWup&%I* zT!4ZtP_PLKHll(CDA)!CJE7oPRPa3%ltV!kqVXIQT!ez1P;dm%Kn0heU>6jeMKn;s z4^XfN3a%g;sNhE^us}iH*FfWWD7X#tZ`-KQ{stzv(H-&%Gbw5X~IQ&OfGez z3IQE1dO})c7MZ2Bj6o@@m6Vu8l~Gq>2UKdr(3n|Idet(UbY8WRZDtYAd#72H>%9kX z`qVO2I-goeEtL8mF^hD*waiJ6B5Rs46GCOnOYY(q?t)vP{d)Go~NG)?DH>6f_67K7Bwn9`{ z36SGPRq?X}$s(Qfulz5uvzJllotUnr(+M~9%>>S{&`B2B$4di0pv1?^NEV_?_O~Y8 zR7%2c_%5Wa-_Sr;%!n%aNSN_iK(b}=cj$J%4@o;s`$@*a( zvXCWhh56VY{&jKA3W9k~upm)W0=~5pW|XM9%G~ol`>cmcn~@8yXXlU3{k~O+ze|8= z$?YtfoZP-k5ML?R>ICz_)xxX^e%3l?Bx=@IM(!k|H|gPAYu>Xk0sFx5)F&CAvugBR z&z^)_SL#!3b@b#RSn>cILQu<|$_tYvoHDJCfr-;lNot5wa4>cyOmi{d{7rG*Fo~~@ zgBd5rM}9^lK4%eMsiS|ih~bD+nEM6$qIbazuY}1>nT5{o{az!6qkOLtZa`nnSeJ;M z_)5CVNt(TOg7Qv$bCGT(?Jw-CGAxTPXB`b{i`Ok_r!K$kI@7W#ZPGO5w^Mg#a@+S^ zwX5T?V^Gr4*2iSgwT;bkC$7ytEH_%KGqZdm==Zy1LbqHV+UnkIr|zA0>UOkQ$Ix%J zQ@6dHx;NXY+uCNGf_|f&y3Os>ZEUBmw584_vo6Q1dibE%&AMsjxq*JJg|lG;6^)j*IVjreYC*`jo_D- zwpbq;!L@ekuC`KlsLB7xTX@t=(^l7~ox0xb)CIR$$IwCT)b(hmu3I~G0d3YPXr`UI zyPX|%*E)6ueY>5yn=N%VnRPjW`Jel|#&~X^ue8v1sJq-wU1K|SKeSm_PG4-NuA!Z} z`gZEhx77V!{@Cu9Hu1BWCpI5*r4`+_)&1N~-Q_mxB!WFH__6*Zf)CoM+tp6pdo6Xg zG5uDZvEAPWPZ9(_wK=u~L1R00Kep7_`tWX@?;61mEp0I#X#|(rsk_)x$7ka$bCL24 z_#%3~Hp6VSA*@{0M!BMma>iC}8$%B`j#JIF#HT28;@-I?a=moOIwG!bM4eYA2$Bk`e;KAbxF{Nt!-<4*y@_)@DuAh6t?kU zuLnDNTH`;p|b1`Hly?=40(2^jd`7f4|pUK9mx+abe|_$+-Tq{46!KQEqI# ze51CFa#b7UiZ;p_TRA3stM4#na{RcBmkw=_zi4XfP>1|QQ=N0u{A0^mApc}CWlB0y z*15uEsZywY5ia1PW?JSCn+zQ8M?M6y zt^Nm>$xAIV!DLvM1XvvUvuoc4bf*(PE=C0ZqP`!=TYDXiryKkw+77uo=U8H$YL~yX z+LAX!d!az(mN!d(mdRB}_$uk+!rViAvLsc)Q?JmXydFaF49(Q%hr|UOK*|9hHABED zYrmxrm4yD&`_TRceHj0OJ`~cwKpxV+KpuZVAI86+4+ZlV$b;ILsYACsx|!9>x=QUQp-cGt}CbbAzg-vHmW- z;wu7aflnTjs8uc~ieplYduK=_@LoII`wZ}Z|Q#nq113G znAF0hh{8YP2YI>QpBy;y&jjJOCVTY8ciXbJ?xCTejmz>}X&$jF3VyO5HA4`;OCOLj zz^64#=-ng^rq=brH5MqeU^8*BJs^+B72D&Ows?g0d~opCI7_Z5jg#B$Gb$i7HV|s& zH}l7|{}|XyJNyyjcyRuh_8$g&Vflmj>rz1E*yFP15v_KN_!qAjFQK|Ca%^Oe(zGfYbw1E%Wz?r=^Jjqtoe?nw|>e-uzBoWDm?xGru|~x2cOM0J1fV zRv*}&9wzAxU~6dmn9I7Mj~Xh2ieQ)Hdc;5UQ8NT_pL}h?mVz{0q}-h^l9z|_haDiO zem`IcQQ3&C_=dGnw{DI(DN(nallU$i^kjHm>M# zxK~u}@8PSVLzWk>ExxvE--aa`uiPY-g@}LXePEkjqUSPJKbgSKGpCz`#P(2udhxZX zi+>T^d4Vi-3T*P3xmi}9S-gs*BML0N!GX??uh>A4$I;-U3TU!0s=#76FXhE~g(@rz z-^>%$*7}M9sueGsV%X0LnR4got9OtOOcw65`cq~& zEj=d9%?0P0+9^Lx(qSHtPj(3_%KJ|G(xxuZJ%3%YP@HaajPiHp$cvB%htaDUg6Isi zY-2|ZL1U)_p<5Yej7Gys~j82^zEcMgBpQ3cwPWv1P4l+uh4H^}eo?t3+RrccdWpsfJV!Ax3>Wd+Y-ozZNBBt-ol#UQ zMJ$h$!Pu0&veBfu+U9N{Fg7zbciSpY?bkzy=iRDB;BR?EluE^A=_q-56kW9|KdL~X zXqH%*w?J=oy@+RM54+8-+r!-w%13rVGyV811bS5q=$aPLt?5Bqn8(|6K#Qi-cmP~L z?LdGdTL5>m0S2rH^FO@)>v(;?`f2^72vgEm3ql=1si2RhJZgs-a^L~*U%fuKw{@U zpS(NAU%hjxE4bF7;KP+Wn$`s)`bZPuKQVISOcrwEf#4B!q0gf!swApFrD}Fv;m;Iq zWPL3=LAEvy>w7}h*bnvJydD|7`0CS(uMHQiRv^W+V50&H~=j zl5if*V3aklK}5dG*)@pO9(;bo7_GK31C6dLnxy70fYIfwhS4RojP9|P(X}=CLeFe6Z=%Z!`;V;8TZPr1A3W`C1-t zpTF|cjv~r4ucI(FKvexWIaz)|6Q>fZ9;D1ZN$er9Ljy}{|cgRipggmG)w`tf? z%99@}+*@QO1?>5Ie$V%FlGX0fizxP>p zj;&<2*h(g{l@3;08P;qo^+P5(+RDjJ7N-?pD`%Q)CDdvwBjwVzZDnCUr(_{HFh#PD z-Rsbnt@N|Qk9;K8kEIh=LtkS5>cIaeU$%ly=$hPtApaM9S$pz{|Jf=Zv}fhvkCVs6L3RT3ym@adzBxcpy7=b)dGiKzCi(kFkR|zzQ%~FTnMP>^Uu=sZ z1b)piguoAC2%eKQhEN`|#oc4%@kEs_ z1FD%~^UJYaqYEUKIQrH_CXS)i^6G?QrBoCqR7%liD;Rfsv3Z;WkDBG(Zv`zOhsr6% zgj!{z@7;4#?v4iY%v0VfrVX8EOBT+XcjM*7vnB8aKa<@RDzvsv7Jk~Nn+GR0nbQ)R z%t5h9M{8_S5O0g)Hs@X-&V! zR97gvi|AOXQLCLb!tf4%pTbEei49HCn3Td?C+Vi)y&hN6{eEQP2>KMGm(Y`(7hROb zCb-EizM~hB(uI1_Q1lhs$cFiC#`*R(@KDE+23_a>DixykqZn(eN= zHE~pnEk*qqYhoX&%ce@(!k7rmh_^0`buS@{X)D%mPfqu-F8({@#5&O5DW~?<$A5#I zP*B;voa|$k2lQ)&;$~Ya6#p4(J_R)7yq$IP-(W{55!k*R zwI{3lTs!|Avid`OUOQ{(zd=^N?(^D{)qSp`{|;HPi|OBKSM9B*{{~s1MfGpDtNUD2 z{~fZrZy#u%qqh2Z`+9Ok5~bX&`K6=c zy3w8~il4<{{A7aGi+D1vl3hz0tZPX#{CP1ZPDv-ps~LV$`LC?OhC1-6S?+(-8e7#8 z)2RNGHTHfvO~$muXB~$(B_Z;3kc24GS(6aQtVswsXbqAOaP$ZyA;_jAL|bu>{qxR; zKPGKjaDzPCTiZS0vzxy@9&K89whI{7) zo!P0oU4E>x`59ZGbF=*xIwSk#Tk~UNbAD{3T+&K@Y&Swo$${-}TG9#lP%Fc`mic_8=ba zNfy?GbYtKBXntWU@tB z&0mvkl${QvXLs$}EqZnYdndcK*lX;@W_uO1*sB29Yi`p%k~7WE#D~d=R_t{&IavYr zsI5@1NwCn>Rl^OtyPJvGKp`qBz|+B?9kR!TS3?O*QCzhf{D zh8;ZyGc4mhQss~ZiIvmr;-bB{Y@{6zJy;&L6PIx}VO;8Z?S*ihSgYlPX<%Kn1)opQ{uzd#w?wBB0C#&HEw(P@V` zj-bsk5UWllXj_6vw_a;Y5G6Ibh6$4!H?vH;X=Pdw+JNDaEI~|@uuN-ft!=;V{2z_k zo0B5#>4p6>;{UC9-L^}i?O5G@ErLJ0JF;Iy9E_n9>S8pjpAh1S0+GkYI` z?*~G5hCXONtnGbJcsIsoSne%0<~{^$tlVm2-}P>`F^Om)YhzSb(4(^ zwc6MSv&F~3#-1(*8#BM$oX(Jk#iBh7yPDG(b~$558EKE;pD1j+qylIOY|gV+F{B|aH^N@ zmuQ+(7epX~QWs{o)@`=ADw`rlxDTaqmZpOWTMEqgsY)dITB=HAICH(4K z_G@xOG2?{dp#;C;;j_2%|bF(Ue0fM zx4;&+l(xhzr6_Ju=-9QvywSF}h4vJE%Em3vaN?E;+Air3x9n<)Tk?mqYXgIsa7+oP zBz2WOaay~SfJx{AVKj0-a`YshF_29ufm_zqZLtomZc`Yvx@D{>fp%k?VZb9>Okm}g zy&gyoZY1Ls?zQl83A|MB#^PJ=FGka+x^D1Bl6^Vj)Tz6dzV@lQd#UT!r%omEB~qkI zvZ2;Pl&!Bl@4wR=5ADAL@lbBl_U$Ie=DT zCVM;i--v6Qm;CL@-F{#0LmSiD?}cbZHusUwhj)bDSiEq!$b+;k^g%!2Y_UUjdm(K0 z>^``8t?xbIDX}gRCd){JOhYVxwS!32ZuJtDnZnH%? zX*MkQWz7qI-TJx}O$&bC_tApCC0FCP7O|h3&{>{#YmuhC5b(~=js5eg6z&9TR(V+fibA6t(hX!Y^2N*N9I_;tj7Ns;|Nh}(&+ z?1T8da_;i9(nL6$3B~Ka|yQj(t6~9 zzmpz6@}M>}(Ts9)lfXW;=Kd`Y8*U}S-#};-%kr$RezlpOdqBV2oK*N<_Ph4a0<8P) zn&Tu~o^`dhBm?O1!(^il&)HhB(Fc;PUE<-Plx{IB`sCqN%< zYjo>#`<8RZ;;nnWTjC~eg5&Nli|^qHjyqpNg5&nrEeQ_znScM>cRU)K&I^zg!#f_M>16bdM>6*v563l=eNM>Azc%M? zZo=cs#dp8M6CNAx-T2xk@7@i&gvXUkC!hn$fbVU|&$`aCo{!rD+EpoT+-v%*_B#`^K7_x~Aig9Kf0rQs*d+c@B0k&b z(kK^yPtYlZAQQeN>^CD&DNwVY%f$-_!F36`yig}zAQ5aL=&Kg-kBr+TV{K*MlW~!o zs6oa0=sHSzEH0Y;;bl@n*spTSg#YAjnHINOb>ee%=$0FGZtoI!E7CCprSCO@~6${0zPVnAQr>F zo%r$WnT+G%H)pIDVonC}XEl#i?slb7d>IB!xP1$0MFp)f(eOQs{mvrjvvuM#a<}h| z0*%&f7F@!OM!_tFvns*u2e`3LykbzfZcw>7rd&7VVqEm2zjy=FYuy?)jf{|M%L?L}PM%WvR+4NTl?MTbTn39<~-g>~q;B z2LD{NiEk7Kozn!K0~clvx~>+EBk?QV->%CGadW%a%X0YK_OZ`xM2~dRo?QxBNj-S~ zv=s`OF~<0I`(7kDSQ-fSEm4;DzAw@=-k>Y9QDCbqI&sEAoivHalS=Yqj7ecVs7nUj zkMaUb7-wGHJ4(BJydg=Dfnjh=PBg)1E{G^9U7@>jRJX74*|Wt{EuTIC{ulGnkr?&> zIBFvt#Poqka11g3!PDEOv|65`U!F9ktwvE<8mm?XJ?dW)xi-0D!T*fkN_Qy3qo|}rYiYmokSS7|#}cx`cR9-gDrjXb zT~Q+WMz_x{t0|tEPt33Cl{z^dUL5%7NDO-g$Tq@MrcKkqvFkRC59;fe6y|8t3PBL9H)=p1v9h51a! zeggyC37L}D;+fH^@z<({gVDNLX^asnv*BoL zMKkZ+3mg0v>Q>n0F-o{KllEXyhCEqL=af?vkUyL>4DTs_Vo8bQtgg(?jOYIQh`}qn zsP2awH^E#X$1@B2*GWhJGQ<8g!Vo_*c*Eo2A4seA1Wq*lG6N28=A$DqthxKF!=C!A z5smvS3fwjFSrfpqlg|RgXKOEJzbPwALfw65LM42Ec>k_SaN{!)ynQ1Ug0cPAyb)Vz z^pxbZ_xJ1aO5EHp_I4QbTT%lv-&oB0tzt7C9Q*2clUlZ0Moxy^JVph3K6!AJMLx~f z#7##oZ-_F*7^A{~U)wD+&9P#K{7tWADv5cVx#i$mYRoiWaRA)Z=U?Zywl%rm0;8zK zZ`JdT-uf*`eiJ&bzt?Y>m*L&vXx#}$tIZb9n`Lh43TK&_XqNdVm}TmurxXGrL#ivx z^Y8JAwkl5CsW^{wbALqp%crieSHkUD|J8o0?QYLvtvFNXN+!?(?cpxg()0Y*H@+(g##}ZSbLBY4Tw|HOa0kR(;7Nf?ZNiu2;D9XqASNG@(ZB(^Mb_(`QOr=SA^iu*!-OHU?A8 zw>BPwYxN_{4l~W^urvknTMN|nN z69WsofUW$b5I(8uy+|ew-fLR}2CO~=4cH<$U<2V~(}S4nT^Ms612LEQ+-<~MzKLGq zgv!4RCp`Tnyt{^R!hziZPACVy-33m#i-!|7;=3d%m%NU-DqP`1ZekF1Gl;o<&0D3S59-j zi*oi|p5cW9y^W4t_6M)6!wWSXdH(CkzJ?Nsn4w{Bswv{RMSrjFYKMH6Vvl8MIO~_Rhd05gCfEd>ff6rQ zL>QxKIAE>sf&p8E2JDuuaKI)*gx{ms>)p{ESpQ|J=fBuIC7b8!0=~-$M{5RoE_iBFE8@;rZ% zMIV$7^kIY_Kp(UaeW>TN+~&B(fWCxsT=#)v*YDS6kmCZ^c5czhFLUo@=gn6$%+H=& z&E&Ttvnj`<*=>GHL&$|%oT``KIww}GY&w1GX(TKjj`P3wb=Cl8yxlqfn8Pb^82*bv^(e}>`B!Y*=&hteR!Le(!$^;AnKDsa}*QN-65KOc{~m@H>AEw#RblxbZv90NaD@-%xBcFrby8e%&7* z&!fQksD{EtE~*93Ybu}tJcovJQC!$B-}Y#Tfd}TzGilBGxrIq4wH5&bz3)}#bESXv z<01vb@`_;|D%%;$(=h#j%|p45i<$wJ2d=Ye;-Z1)*pbt?AL8S8{Fq1!uGO_4YQgWC zYFITHeut)kW47$pN_gPih7d!NxL46SFhC2@0NoM=2PlKwd60K1g1FsLSnGGzIq%e9 zO!TX__319%zgP9?p!+xIn=V7#ZraxjXvdF<_QU@GA03Hd&?h1PGyohsVUw*ieofSm znnS(`%WDfU(1k1p>d)5QnH{mpq0ck8SfzSP7pGX|0f|D<{RvcPI)k9uFr~h4NV@VB zm*=VH8$63snn6Ym^pB?T4VNKi#`WAr{lm9>|L_#kpIOW-2^_oj55na{6GgzyV9mwf z96w-n@xvF65z^Z^epm$J2lXN@{bD&-WFtAq%A>k~_qe)%+j>9WDR~viSGT;M`Rj&S zFe}HZlqMRnn9FrTEtplkH9}413(iAdz>ShdeSy>3pcc%^dHr7tX7$%#Rt?4lcL-)B zZGEfX;p0{=vG2Em@=+(pypr{`v*+HAu%d;yWe#PWAwBF=4 z?l%023F{lnu{&W`Pcr?u1-s($g)|vQFE;9EoqbZe+Y@gGIxz(8`mX0>rGC zoWjq}IfZ`AGj_vy6YT9;76<1I`L}W2*uzR&+X=7+ovWH2HWGy8fcVDL#qP8oljT9V z@HsEXKr#31rtV6S_~*HA<~GiKV?Ok!xo`Nt;x+e8jQ6+e;J)Esi09llVRW0bJF^92 z00GVrjsQnk&LhB4ZF3^P88;6uP4;66aCBik@>v2L8zR6-+vXs^Vf=)+GvL2ThWt1B zD?b6gmS-;)Pc8cN8Sv2OKzSFa22#9R9C(%O#Mi z9^v)vIyi9n7xGVW;DnX7FD}g%=l~JU2#yHHsLbIJ;f&hPC&D3*ybFkMNWoJq5l;0< zVgyTsa~|c-AN;~WgyYta6+#{y4dlVu^d88KfB9(8M6t5=JP!I6{OCvwYwmu`Ef0=M zKL%+$8uw#-Ow85#7NhrTGhM#B+^x+V=ItBO!9w$=-;fL+PMP=j>+(w6+^-8hoUqb1 zHNcmZnDKB#inE4Cieo?gI-eA0?0!IsBOJz(;)v&3Nw$XGfJkw)N7gw=asG9_?CjXhyYb0!#;pf4P`> zOO9j7^H;LV)2wKD+CwAj$Z`IRFN+60t&Onx#$;}}aa{Q_2g3l`YnI`qbc)9e%>A~%` zPfX^PC#R_Z5%4_z)>`EQG5JX~fu#Y)!Q!&RGp7anB{OlDW8fC^DF?=$g9EoWY#V~lS3a_alD2@qq1 zR3ENmvy(iGkp_2<>X=>lG>P!HKhK-PkDA=XD(Yqso4Kx6XkI?;1qXRf&Ne6VoN=^7 zu@Q1{U|e!pm@{5kdg+60I< zLOLo>W3>LD-Qh<^V%RDm5$CoR9J|IFHC_X74h|iD^yChJTl>%BjqYj~nrg^3iC-%l zg#21+uWH16l@lAyUne^5V~JuvpxK!sQO+_4Lv)ueocJd z*{|`rw{PLmX$HA7FcfZGk7Ey`TOa4=&AJ1hR!61-T^P4_fK!Y{0X_0(Em$KZX{xbw zEBE#-SYt0fq0X&YGes0fs6*!YN3e*6%7IwSLWDZ?e3;vu89z3=5A+Gw@7iV%n}KWl z8`lPbnPEE%KVJq(;Oa(E12uRqVzXD6vDFx_&fOE6$rxUpe<{Y;%ZXg))@)g7FOFPC z{InJOf?>-yAaWh;k*$ba#{<9RHdodl|Jxnqw^-^BcRuS`p3fQ&j_a)@y_;~wCPOuD z<=5frJl=QUhpi(Jt%qux6UolFOo^fkL$b5eksY3gNOr*UtLLNK;)i^W?G_(ZS0M#` z6e8QX9qe!g&p%BB$F6<;UF4s>`14#le9-Ej;@SzL+i!ri9FF`~YpJA)i;+ex_LJ1@ zXG9*kcVKD^sr^xIBs(}TkiQBaY|tMy1AGv4`Nrd@2_TMQOKjPFAh^0iJ|@Vo*J!~6 z9s5AkacMVIPz9c+2YD{8pXbEp^9gvyJp%}MRzLj;K+3ZmBcYCfr=E{Wx&p938_aR# zIAejHtzl{gSRiPtCXAyd03XGSqQ?9aAOGXWPg-!juKiOB{?}B)tBLVHl+rrVqnDgl zD?e@v9H`Rh$IJ>@5kmeJ`$;nx0OBX9fP%o_3~K_`l~2!^H~la9{$weujApV zkNb^X5%SdY6}l_uxx7(5U*TCi(+o0wpuaSsuNVvC8AQs{h`)G)?=PNa z`n3n^nwltTZU$>Xz;pZiI6b*IqZKdC*oxwehh{m)8J$6#QHf}KKnVgOPV?Y>z2O(G z{@}LW*LM`&#nO)O-`D(gO)YqrgCn7dMy@Mvo))~z?LOX3>mPz7*a zE@w^zSqt6;c9i?;>fsuUAMPZ)i=!Oj`y)QyMf2i^7QCz3<6TWs1Ls*Rb64VC4g#OX zd=)>IX{mj4_EjylZ`@Yie{=0ygFfs|!oXbQZ1@-$^jj?$Sb1{)O9D?YFh@q@dp@Dh zxCb1xB@bqj4nm)L_*W9dznEM}mN&VQyAuEMBv;~ueR<4B-9!=eU59<45fbYwHECDa zs+4<=Lu*XX(Tvys(+deb~7Hc zG#1mVAMkd;frBKV;;$h%$&*8_u)RFKdx){aTID3Memu9aw|QF^SuYIxL0)KY-$w#= z^5(BAO}=mr1|#63BQY!&Nc8Sb2gh#i|3R-81~`aFN<1JUng83LB_iSfinoimQ`O_W zzFh|q3I9Uet%CsV{XY^u&yj{;8I~zI0QWBl6STbqawc>5oIk#h2SFUXP{QrfP8}r` zfR5diBu-w)J@ew&$jIHYO>;2uX>;H?i_bhB zDU?@qO{WV02_D4hMUiNa(4w8(##@`_)NQ=wc8^fF@zylFu-+b_HL?V5?=6Va!KzrV-wA{{Mf|8v@x4!d8F|+)*SJ=tjfVRrsVOBxht!h^rthl z{9o}ZLz|)&wbmT<`gR?uV*Z7=yIt64R{Qxxv3*9|XWrF}^1Z9s{n|d5SKRo1ZIs?7 z)=$J{BGAoAGT?KS8Z;73Pd*q!4CS#lK5{h$)}#1EQ4J0L$e2^>55!B!gK+Fe=Lpv0q?ka*C0 zHaM5>IX7%@l}EgNn>sk>{`A{aIF2NVsovkO%PVnnzo5N8_Cl53FRz%M`3r4Q7@c(CSK5D%K2- z;x6*#T4g)D?e3%|18UAK!aL=4<;92gO8iX2W+b(OX>9LB#T>anYyxK!4_uAFPD9s1~fy za*c~#hhGFfIud*zkSHE>;R_{ozVJWZ$4B#A;a+ayqvj9~KF9NgC&C$DC~IE6@X$}0 zKEe4z5lnilvqBDGQj_5o2TU48j{O-J8KS_rF`*zbeUX7+!BOzs;Xs&J7Zd8Qs54l9 zSm^8zZxTN{tvY&W`kG{Si+*=dlqm>XJg`>{Xm~a5NM)=W@7U z6Tq+;;G>84)cM2{CtlKO3?*s8rW}r-A2oqbOmT4YrS`4{<5to=G)4u}fz+{HaBwe4 zGG%;$2KTmTa%{jKRbbs=sW{kxbC%e#VvCyFyw-!*a8poc)+dU8arTJ`piiRM@IG+t z6dyIKTp@-%$?S{0fXx%`?at(t+MH(dgt{?~%mwg=O1xL`qkUnP5W2s{;=&eev?lO^ zIq15&`C+4T8IGgIvxKdC>w1&qZkc$#Rm-ue|A)(~W5?dbbklsrV8a~H7wY*zw|PQ- zTo?t;Yb0)I26;m00~J5z-iLC87Mp@i_C(JyJ0|mPbA(zVSU3eY*UJ(9B2#5bLgN|K zx_uFwBTPcsS@X6?{zz+#`^Y&**n*2brwy73{<*7gQNaLA#8@idwe#@4`k9 zW{RnQv6buDDbfA&iX2m-vg@sC+n=6ndA{U^@3s3by`dbgu4pYw&=4vGx<7SnO*gQasfpXmb7p`a*p5KT-$w*Hwu*}6qPM!(jfVR> z>L7+$33GasGlt3MTld7T{qq=RlgDnAAa?Vpu9J^j@?*A44}4k+ZrNlh^`^uv#la@& zt5YKjP`&0Or3Xkbh?k&2d=3uc`A28dTy!5rQ>XS~uwKliY= zIXstB46im*@X*f{13%ZRa25nB`7v9r5DYNz(UBN74M-HT4Ft#a{G4B1-mZ!ITQkVp zfoAYzKgM2%G61~(HWy1_$3nPR>eL(*O93j{&4&>23<#s-1t^TFj}_{#T;pm;@}2$7 z?Z#36d=+ey`yDtx%$qa!hl0ut@{(GOlj;#f&vxzU-#5anA&S?VAaXG3c8r@)MxnIZDg4Jy; zRapgz7CdY%e%R--O$`32|D>AS$1wRQc0@JmW1QBjwP1Dba#1Zz~KhPwNVR%WV)jD3j z2sI(F$754Yp;L;$yA@nWXb6~Dy@)Y9cxaAH_b;X8mFM;7>_CUA!>@E2L&=^FnU zW#G@4cnan-6a1Aie@P#$8k$`!4tyH(sddo3$b8n;_^&4YH_QA!#KgbZ;ga%m-Z8Qq z$O2c10ivQI-B-=c68HHzMlwx}rMH6nW}Mv_mRu7#RthG59-&`8yXZ ztEvnJ}j1I|q_-usF8;Fks)Pt5X z=GT%t^tl*dDv-<75`%upXUSMN7>UVki%Y|jF>xzEQpIJNhm8az?-&(liptM}WQeZT zB-hAlil?%5z;Ez#N^dc3&HAS)jkR@Z(93=$;f}Vh5F~D=6F{TeN>v*Kax)%Q44OEJ z`Ku43@VKqXs*2M1YCO#DV2Mi#Gw(@p9GHo~9XL>!j2pe*e8UxT*ql zuYNv@2qN`rM)jeSYbYwFB1KqwqE+n>wsKDdEBEZr?~KM{hO2ATz6!A+R5%we=NIq& zCQIETRP9lUeo^j`6_nh6TwP7b9fTcv=?{>GgI$_TKLD6Hz?x(3y`LJpUcOs)z}A5o z@Zx8!WWADf?U4`A2ECv^G5*t}P=PVu%38Vt_x%RsjbUH35c8{gG5)joakw69Nb89v zW$iEmoC7~P62sI$qV-sJg8w{x+y0Z`u4>eOE@Ik&t?h%clPmuTD6EW^vfq&%PeP+* zXF?@>ll=wnLN`^VU{5M9U%igPlE3cN+e)KnHl+8q1sc92H!T33MW$qUk5z9%ADZOP z`OxGx*MJXI#sVLjES`UOCWzURA=m_t8=vrocw)oJ%^Xi~1%x60Gn?A7^{8N%pFWy6& z+Iu^T@Y0DknGc=sC=Q7c_A1PiO7b~GX4E;qg>|ted1OO|JXuvo`)C;9O#4Vb_%$5x zkmlph_)my=-%0)xvcSNP$MGlNyg~%X0&pJ3>EIN5@_px{z{a)j6yMdp6EKZAHfA~Y z{PkncXSP!7!obd!P3%lFC4oKD4?|x%G7!b1ZSz62qfG+bIVp9%)LMZ&;{z`ekV@a- zNb_al;LX*+fHCJ$Zua`%Gln;T1!T;r=R5T<{zCKh=Yj|ju=RLWP=^7lwp@e zdlG|gf-&fqAO;nm@8%1-A@k9Zp2eVX-eR~8N}9h8!#({aoH2OlL$e@;%kX8nj6qil znXh7E@d9v&FkJV3R3iM%Yp3C0409P2x{1J=RSfzY&yPO&B{3m?rC~pCUhWn0JST0LD*W|u* z0>pCVRUTuxL)jt9jM3Zjr5dlH%u)#L0uZ>`CAjJi%4|+puJZP~XFg|KdxlnMekcYC4gR5vK$HBFU$HC<`4+=TB z;66P{6SdMg54tP4CJf|31#LF~WS}OAHww3<@<-wA#i1=2Zje?RbWYiV;W{wf&Vn|5 z8p?uRkfOofu}3*K*gK-ZzO6eP?7aQ!yv%{}_bd4K`Ai-^mkWPsV%~jh-qT&(5TNP0 ztFharM}Od)sDZj6p(E#6zJNmU<>6IM@#RK^af75PPOOMTn|nNT&UZ-$SF-9cPRr+K zZt;zkA*1&8ZM?v95br2}T1eJx2mvTC`vU3Ro zy5G|-Q}O_)4H6FzTg$^?lX~YZWN=t}Wl7_;LAaO*6=vf$er?btSqiKTB1PvxZBV`* z)CNs)_{1;byu)x;R7ZPc%?t^zjy52~P>=6Rdt{Xj>;lWMHP_2s>LA`) z33Hc}Gv3PQYHrjgHYdEb>Ep#L5HGr2&%yVHe2kS}FFPGvbA$1C3&zUjE1$gEFjkbR ze7N+077TI+z(QW(0(4uUXv8M=Ns-VeHN-2%^GQzY3|nwSPxEUnIARNq$nFZyz9T)TKY=(R zpL2}!fqWdXrGE1}d*CHXoNX3sI)cIzSG6a17>?K^^_#a9ueIQb4f2ukeIp-JY^mGy z&^NZ!ZK8zwk*2KM{8!_OF7j`DToLBnT5!dN{U$q4@d#Jsa$njb^6C8C*Q{UzLKHWy zlsIvhnH+pEI*2JUIj`(Lk@IT76g|m#IpK*c_d2`2y#-HfjOqIy!4uKwnH54S_)Fua zOHe^#>_RCRSPYWrK@I-~h$AvNrnAmDrh&}+kk4u6yZk2i8GUBV;oczg(hBSV>?OID zgUf5KD=sgu$hA%vh`}{yO7goGC_0zN5xn=n!K}3h!ZIt9>n~ zwBHun4VCt@c`(>d1f9UXY$vab%h)G-uG{-lZT(k3KM6lN62strv#74WJ2<{IM+3>t zf{tN>5&=(OTYRLqPqg6<v$HSF=5_#!}6y@>oWzQRLh_1G;nOyn2;im7+ z`@n{!TZF=(<+3THElJy+nuN8N^~%s<3W+%WiJ+4CBdK}Pak8L!Qhg^e_!B&?G+?;A}9qff~0^8kjhD_tIG55^NAMJF?Xs#C6Cz~{pC}g z@^#_}zvgRTjzazIa1?^Aul^)Qq0681mZMOTuO9Dp z3pn_dTwap9mAzoR764y+|GuO}aBfwGe61=8`r16;Yts&9C;w{36>x~D!4PAq&e%nC z)X}wkI?Z#%s0?xZGnyx{p{N=ElAP^b@#Tms|V-;nQ3!Q)_@3nHf}hq99u%j$gZ*7=+# zrH&Ty3p;1diU3;`&t3>vloTA}3wCS`WKrU+KXCZk5*XKFFs?0hj%z0}{bvxvd4H$H zwTOT4AM~}_C3?RfU{HHIG3`4{o3kY{8@|lIj(tjih3jfyjCj# z;01gu2ODS2@`u>4JVt3Jg(bG5_Yag$(Fy`Wr$~e|+0}l6^ioKp2pLJawR{cx+RZSo z{Tjrz;$Jd-!L2eM9qDOYd(>MHWFDBOflVYVK7Snu^7?*w@wbH^&xiOiAaQaN*KXtC z$4kL?qqw#Kf19hGDi?h$YT`YN)w^*`+uh1nH?-iDe z_zM3ipF0U+$>JwG#*&A_2LnDj62orS@BYVoD7E4Dnj0v!Sz*ae`+J`71+T_}L709{ z(`Sa@yS&N$ZnnekjsyH^oKJn^YyQfQ;@HKq7Xl#dqIMskT~vg!VKH`WGo)RFRHjVc zwh_O3kn_8ydVW_8`YhrCj{?Uo@zds%w*}V-dB6)4sN8Q+S%M<;S(rpl{wu`C4$fV4cWcvnN%l-0*`_uC^*cSP65r zDM>I_D+Y)|zMh|}eHHlEHl*O0Ct#-5H7A?9Tam@OqO3hX&A}wuyKuBm@i2G8SSv)Q9SE5Ps`t9-D3sdNmVMGy{RcGYkZ zr@Fz0Couhl#fDpaD6ihh10PC1$2!p4!G+iG{O1&K+*0eb7C-C`7jE*|E{t!MgF2^i z4_wW#mFxKUuWEiVQ|BZe9p0nyI;RZMnxt-z@#~zX1?ibOrw1@t=hOz&IVA;aN&|Kc zM!DD2-lSu&Z|epJyQ9Qui$w7`MZK0OBc65CTogNY)9@q)b%*_W<|~^K_6x?$ZN;GT zK@7@7pU&|oAMfSYF}2{mZsN~77w_ewQTw)`p#~n<`yzmNbuU_yOlr!3^hQF~CkE-Q z7lW>bIB(6fAh&9E#(5jx7umUHxH;jxU?}}#G3caMjapCYg!$NR5x8CydwP;nZNYZA zoa%|Y7~AF2sW~Gb;IG0d-ZoPr4n9hL3`#3bI!Y@xL25^KgDMYvYB9uhzk+qb=biJZ zp83>fR8)SuK6NIuPP_d+i_10~$Q?T~6yYTQE$O zRgg&O6StPa0}CG3f?;y&M_MpUPddfz`bvmlN?YIJEBP2^OUvy;I z-%t!OOePO@j?IH9Vlb==xJUnQR)<&V!VfFLD^G@1tPU&Jg_SA7Y$w}4ySjajuKj#P z`|OkLX0C3RrE51u(N25vp7hoCr0MP%r?^LRvh9e~ZBun^Qxt90Cxfn8f@-xvm&l+R zd*E41V6`^zG#O~O2OPHqRA~balL3`>X@x~vu9cRNQk(r^{Om7!o!#8^u@@q$=Y&if z)B3d1FGEIHeuNDVJ~|S^Isw^6IG9-z4bxU-n}5Y{At129TWJK}3$SsR%rumu%sa3f z#w($NrBbj6Q@y=Bvx**X$SAX617*^(^$C*xV0|V9rhvZ=%Zn)N5{c{%l-AI#1Ol+c zO!nS?R0jXR*SbXJSZJ9zMavbV5a2f_iphLmAmK+xV%TUPQBDBrov-NuwMecGT8hbV z(U=>6-OpzbFu^*Bg))QIQsmPeGwNEC zFqdh;O4TB=+#(ZXT2AdMvycT}cOBGOqW|vR?bsI+iBCp<5vuC1sO)42pRFeDFPxqK z0!9Z3*M$cOpK~QvSQ`=+ZMH@@eR*nKm>z zx8XQm8~Sec*qyplZ@HO;;buiFhnoZ4}v zb30xwOh?8uIu?O;v~cZMX;CumSOMDc=N_;fFLvFohV7`dUlK1JExDAh%?mZ#rRzdp zh_sPXwe};0?3K=veDpZ)iJ`L{4}C#ziKewcZjI1H`*am9kz6w1J(P4P*Y0Z;P=P6`_&(L~%Yy=I}v=1y~@*`F>aUbY^&ic1-%tf2aC!$z^d& zWj8Q0k{fyz&L%CGK#A$knAF-A;7Uq%;yHG$)Q08?>$_Btwn(YgzEHg^Y@JLMQRy$X zTBNhNcAs&tA+9h}Nn@RVkQ;VM8)(;A%v@YnA{i3tQC zxz&J)=&wo)6M_HV&A75Q_a0*8wVi&e@BObXSKDR>*1Hux=gGIs3x!_^Pl#+X?RstdXo^`FHuXM`E1r@xOfrLOO2j?r7VQTqoi46ni z=POc25hD%y?U?iwaZXv)OIwEN1u`N1j!4kOf39K=rI*KK+c9B@Wp0TY%Tsqaoy7d3 zYY*nZdOzO?(;5LO?C6UbPnxIzI8%!>gvrbO50O10=~A22@rAxVw9u~{aq6`pj{_2eEOvk zpOvk^{VR3*EJJ(0KNUQ{@DHR}>fSPBc-RL%A4?EuXocJnk#syGcm~aPL;MFwOrgX8l0IX4wwAUV22^#rp?=iYsP`iV z7z$G)B_HZkW<@{DQ=L)1(ibB+xjl zjhcs;s{#NqK7pLgn^J^!VU6NiSSkM(5_{F;e$rdgPB5?T3y(>@ z0bfY@FkhG%#(m*emaxlqGEMqT-4$kcvgMsLksX;Ul>Udi!YzU)s`k;^7+ap>mi?BJ zoxv?HV{iEscgw#7{h%A&*0780?6}C`iNs6bA|E1{i&Xit7ddELm5*FLTa$)^!6#9m zhI~3`rD`!~+KOK1%M3=>skQq_ds-9*6e?Ps;I^^G_Q6w4et#>05Y`)xtmNTIYn*@ScLeNrnDl z{FBD*rQt~z;YlCa`6rF&l6V&06FkZ4ujHTf;8%j4;H2N-N!!U6dH0mwrGI56xF>j$ z-T!|6Ne}&$nFvn03QyXh<)1XO%YftHq)K>_EI`FS>ES;F{lQ5$;Yo)q{F9#Tl5!fH zBx?muQU@gPPf8zVgeSGZ!AVE#{F5rW3_1%=(!i6l12p`TCQUe<0#0fVPpTqc0{8T& zZ@%7dpgc9jN)Db6*GLG!dki?T}1 zZQ}4*@+?DpOF-$UvqI?aI*CPLfQr88;b*PMOZ39{;WFvCR6;2X#}slwYfMrakQ73d zM1z#Wg~rVMrrU?+#MK*JevY?=UUZUfM( zwB*9PS(bqpj89jIyUK?Y?q$NO>I3NGoF66OQk6;f8g8_?NcS{Gnm9ruEErW zU9#daqLa8SChJA2O$+kXdv^@J`O5VfweOwTvHNDoJ<`JQhG`awx~)OcE`$9iIFjfr z>L`CPZA3=Qf?b{4!GocCEck~ABI;(?^?8>Mo*eM@&9IyEZv2u|LXIsVpDw{(>5@Oq zcTFHgzJ!Y~qf}znMMRh{^gOg9>t@7_MSt|Xup?DOXIcCydQU*uvaFOft$b5Elgp?S zqBVOU{z!iRQKvM+-+P&8lOmRuMJKKB6MT-#s>yP4;ae273>U-4iNLrKKubPf5`1l)RiV=A|Lwxh45}0h~}BSyh3jN#-irv;Xvq8Ey$!E<39%7VHdc zy)0mPYPGVOSmX0Thbe+C(tgsjR53ke$`;|*YLPDMf>JteyF4u`BL;NOd~r~z8Z^}< zWkoOPab zi>KP)9B_FO6*y6RKUnj>-bO#UO(K!CX&0p@m*KY;XycE@1)IGPAo zlvbk*M^BL9IBBxW=uUyFW=s6QJ`7s$K4s6b>=>JTVw$f>?`Mz*A%`k#>-u?*&OnEr z1?;@ed5^y_7ekCQw6SVWI&um1@-hGDbybm&)|17Qn#*+=h8wJ1Y)qAetLMnphzzO2 z6JSm)MSbO-8W3Fya1D_HCJBNr*c2w{7lOWJhydX*{HqEynU7nBBXJa4Hq3@v=Pq3=ZoeaS-k%z#~lV#I7eN2{y8un6HVHf3C5Cl5T z>SJJjF4K}L^Xo0!iYuxk8hAD#OkZYs%LS+}U#fva4vUk?I*A8`fGLbI?Ga(Hj2&cz zP-+S_3&DKEkPM^6PD}$c?Q2pn`9V0}gWu&Zvqk17NowKRMEKE>7zXm|3COYbfm+FnDyGWT>5cC zF1%j&|6tAyy^96lcz%h*CdVcH!6BHxNx=LqjVf8NaTh6M{Yrg%BWrIUZ{T(I3*U0s z2l2iiUf)(^FENmk3pc?zLqG$kQF!5%on(PLY&w|v;bN^gP+#Zg_`eC?iJ*&u_whhd z+)3;MCz$5;6J)!TC$z$yET+<^{mbK_y$0Fy{Tgf+?(xI>2?t+qt4afNa?qGM62ls^ zO}oVwJ}dAok)Ot7XbywP1pdkqL$=-mCchnw00aUb_`Lx@qUQkb@ltjG_X2Fyi+6Ie#y^ zo$nSaVqjnMUt*9Ov6?t=$b57pYTx>8rhnyg1DMHYmw;x|GtCBVJ!><|v=3xpl%#7w zv)Pzt^XzwU95mZ)JRAP+XkpVa=emA2b@u~)g0b`L+0^|8>^Sf~Sf7{d{0~(Q4-1CuZM$@V9g%TTjwPSa)5ZKg;&*76@h7RJHK1 z`buSKClJfxIM>IQ@9E;h^>NFz?J^kwablUWaH_! zc!#t44yQ{Bx@WcPk`(wRDOOyqqbGv9yQYP0SZ8+P8o z{~l9o^o|<{eyfdeME!RV*c9;l2mHQ&SSlI|P%7a3VJrjt7Rn8*h4V^6B;7Q|PKbVi zf8Gu?S-`+<*Q;=<%wRW)9`!YawgYdPaS*ijivXaHspMX;%`pK|Q}USCC}q9~VA>Kt z@S>1`YB&olh#FjQpJlw^lsPpo)bKETxs`&Kn;(pa%AU=5rHg9Rg85cOz6hL>YDT9> zqt}f$oGzSL36Su1AXT9ob2#0R`bd%=qkFCIqq_JQ zegBp|Y4c{#eLn7!-)SZt|5=}M!7RG}u|AU1v*>}}_em|AO%J}ls&B3EX?o1teO2E-O+UD;Z`MoC&<}soH~*ez>G8Yz7QOT= zJ+ZQH`G20HCwBabqan|Txy8r4f-uyADLiZ2hfZ=^rLPsP`UQk&=!{CvDBO#L3c3;!cN zCk(UFpW&u{)-aou-jBb~uR6^9A$<_H^pk{Ve?)(dZ|WBruKAe$0{^gIYB;u?K7#M* zmlIyOo&FO4qMtR~{0V&m|F&OkxaL#(B>rcpv>6UXqa2QBy(xj(?Vr-?6fS zzJhtkvbS`l{AtzE< zMK2&0Bw8bDs_4bUONq6S#b46Ph}DUaQPfxTi^PV+)F|cG^h?A?iCIzF9+_||D3M2YWf}Gh5kid&1dNK#B2S_ zyQ+VpHxldnS9i_+nSPJ>xWA-Z_F4LU;*0)~-E8OSEky2s)Nb}G^fuyu1G2hd*XjQe zX9nbVBPGTUiC5J{-PHcZkBR;2@^0Ed<0r&Tb#*r?*tmmuKS>g82{V=u-z7yx+d3Lc ziI)bZM%%j@cM*FAW<_H?jC+VH1M{QFIOAu;`sAW$b%JppaU!`qT06kFpIDkw9Zd~0 z9wbUqB;73|jh_=2QzEYtmGMX7=>P;BCi|s1oqMS zD(y|<0)a_Wq*D6jE)rNZ<*G_)?qWferdp*5%3Ug0I7Sjz6P&wD@a~w%IBi(&i-NDm zq{h{B$bCsbjm?R(MdcO>-XB{Sr-{yeMNm1`8mH`;yGpS1fvULd_}te8)(2|iiW74! zg5MvI$JeCht`S(qsp82IxoZU{#-+t;M(3^*EP60M-aaAs9l>W07R6Ui%3Uw0ey}`V znU%Xyu;QWWcx*=QdxFx3B>n8Oa^DyH`cP!Q;+)(qg4Z8T?MFS6yG^j~;jDg||K|Qr zaQ)%Je%T9hKN7r~X6>gg$lWeDnpV|MyEOMx!P4=y{VHF{wF!2OmnUdm&)q4wFGyM68?1F8D1&-k*v$oe~_&%<6AXHT@_knUde18f`iwSUzw8my&w>}GmG`&I zG@TP%pH|&p^Ni^iLG`1O0oX#*uYxbKA_o)~nl1`~voCL*@cgqm!-~JkD-xc5u5eh*nY?#}NB?Ubro583LAd|BnJK3}68@ZrTZd~G%Bk(bq6K)>@X9yj)ThFgi}2dv>QCg9O?a#Tmyam^ zN>1$*mM_6oBb2prs!V7y<7p#`1A?jD!rEnc&IqzwFjX$x@*-Y1LYo>)?G>(m3Ac{0 zObVtdghyB4RU^z#2U7=xI|}jI5!p+FsY>CMmAHIl@rGdPkno3Ban;D0&w{DLLh3a< zZKU!`Zlf_MJ)=UP76PJ2TvP?y%R$HAl$ef&lzR;G=!=a zp4)&Ij#3^Ep?(tDH{sS%wksjjS)rvEuNtKZ3Z>2qmv6>vN6ET}QZ>S(TX6Yk+ptjT zf^hFPTs2yg6-r$a=KT*(8*O_ol)5aewcZEf-y@#M8!W!o#T*qRWTyobj4|;Z&jMlf!u7 zc+L25>J^doC~h53JswW266ve(s_~Y^;neG*>tEuv zxN3s=+i+^F=-M}U+63kGaB7|C7dxIa!4}eiDiUozg%?hsR2`^yMem%(trN7PI#3%# z2foLvCTJe(Ky4De{UcsGLH*wjRI%vCYFs|i{%Qwmv*@dzaMi@@4?9p>MV7O8+QjUG z9jFgPubjtoCT5@OKv_lGf58hUlD-|Ok3`FU#jO*S;T@^%qHizYRTC}!I#QpCPF})m zCn_g&q->(%%eXwfIHx1EQ}li6VQhsof&OH9RL>UD1&$7ah8e z7p7PK)REdNI&u@YrfYm7sS42=JyDfz>ljHL5Urt!+H_g}NUBnF#z@E~k?E1tA<=gx zLN$qeI+8jp+LTYEP0|)dQb$D_7ZN#>Y?~sfD$&tJMB${Gy^+*cqOFSw>m=p3k<@X~ z)g?sLB=z-3>KoB-W}a^(Rmx#1Slry8K zA4I!fCUPFJDeiT@tNdOQcOUr*)++i?*&KawcmZ?@HB*j73D@ zWZ9Cg)E}aS?-JI@#T&X(*F;Cw6IGLC`?^v$L`OCfwUf1Hx>A(*t@j9d7AEON8N~0t zPpGnrJ9MLr;$xeMv@B)6Zj?#<@m3-y%Qm_jl`l4ZKon+?v%66X#U?9Z&B`w5Miq#w zJ|wEL)a$!ZOT^!POw?v!d%IC)@uE)%`E>0!-KZDDFYX{xr&CwEQ7gm?Y(&oVnt*7k zP+VL>6i(MhMpLhdkL@I^(<|ensa4_|r9{5)YGtw0Z&6+;s0;q`yU*&~VU&v5ogB72Z#z~J>GiR?RoKpcKQlE}U{&_jD9 zk^KXlJ(9>ii?c@(+2?b5B$2*|(<6!W<(wW#q^}10j^KKbM0y|4FW~hf3H7jj!hjyy zBZ=($arQ_edktrgB(i^;vquuy8##L? zDrD`EMD`It5APRAWS_*@BZ=%E=IoI~_D^v3NFw_J&K^l*zaHq@gWp9G+3)1+kwo?< zIeR3LeJy8?B(iT005tghNFw_lKo9R1Nn}5gvquuyPXT%o*dvMTjhr4yq+iSFkwp43 zPLCwgpXT&PBE1kK9N_as66#_93;}w0JxC(^c+MV4WUt}ukwo^{oIR4r-o)7>iR{;L z_DCZ8?VLT5$o@Fc!}doK*PEsy^7N#iS%il9!aFn;q*u%eIchu z66vj+9!aFH0(#iaNJ2gASHE!fNFw_n;54Bhkgo4sE5z16)5V0_DCXoCD22AB$53f&K^l*KZUbL z64^h+*&~VU3pslvk^M%_9!X?h0rc=ah9t5-#n~f??BRxi&>l%--yZ1U{UV9%`*8M1 zBKr}XJ(9?N24{~Xve$F=NFsX+(8KXOASZKg!u7iR{mD_DCXoDM+fr??)2Z zcLI8Nzepncft)>($bLM~lfWKHWd9VWM-u5*a(X0@{zFcWB+`G$>5)YGD?ksQ7m`p9 z`)3eHiNWhZ6500zdT5U%vLD3RBZ=&%a`s3f`)4?NB$55AoIR4r{(a6KNo4;y(8Kme z64_UC_DCXo2`I3E_DCZ8jzCWWJ(9?N0B4URvVVxPM-tgT!Pz5;>=$$PNFw`{vA$_B+`Gz z>5)YGGe8gDS4bkg1gyz};~hz;htIbI&_jD9k$oa(k0i2xfU`#u+0W(dkwo_SoIR4r z{%y`4No4;C(39YLkVN(;I6aa`f1T4KiS*%Z+3!IT=@U6Ul1M*^(<6!W&vAMrk^VJK zk0jE6!s(Gj`V*WUNu5)YGa4>m*?SLfGCjvcuUPvN+I?%)U29i(@pYI&b9!X?x z=IoI~_HS|aNFw_(&K^l*|0QRSB(ksN?2$zF0iYAY`$7`gcL#ddKaoWC$(%is$Uc*^ zM-tgT#n~f?>{oL3NFw|9IC~_K{XxziNo4;$XOASZ7lD@>ye}k?z6+;E66ptVdL)s4 zDyK&h>HlBZ*#OyAm1nwx!H|dyGXfDr0xF{XG~6HXO2mlrBMDp}TxA5z@!amy_cok9 zee6Dc@1<2N6q3<_3`50a1}70?QDdPY#;6$*Es2I?ijoXv!VCsStPoAi7^5lFNyfx^ zzrDWa{I2!w-BXx%?z7(Y?RT&B?Y;Z-&pEd{N8$H5N8t}UN8w)+xBfxlKdNuKY0KtN zhe#x*x@VtNT&>_qrd&|4H$al?TQDxVW}v4#od}+>hda zj=o80eiZ*H;#MEUf1&$P{MWi4#edNKDE>R$kK+F==P3M9=P3N|oTKpP>YJaIABCUg z9ED%u9EGnqN8xvhTR))ikBY0c=1}9-Kc9C$ivKU&kK+GF_oMiquWtfdeH8!M;+6-+ zzuWyN{*L&`@}u}~caFmEbB@9vc83$S{&;2O=+uV=h|B(Ap{11xT z{-XH5ByRnK;{Ue$QT#{ihYRLM@t-Dc^-=tnx*x^=3+_kp-{O80|NGpJ;{T*`6#lq# z6#iZ3DEwG$mgPs`=Q>B>Z*Y#nH=LvJ_lfJjITZd0LA&3g#;t!IaX*Ux8}3K(|IqzI z)qwpR=Omq!)kpE~6xWnFl=|1WAH}~cuK(sx{O=L8en9cxE3W_MQ2Y-Gnjgjgzr^+5 z9E$(@g62o@A9YIdqxk2=&5z>0*!?K}j&l@#yK@wNpK}!cuyYjtHRmY&@E4{0DEw6C zDEwmad7T#uUw4keKj<8V{}1OV{GXkp@H1YVj*G(IDEw{CQTT_QqwvR_qwvFCnvRRY z&vuT&Z*Y#n-{TyGf66%uf7&?;Kl-$k2ZdkY9EIQL9EJaya}@qZ&QbVxoTKoQUY7Ep z@Sk&z!Vft|;rEGKzoGCybB@A)zt$Tw>n4RA8?Ms9~HOj3x)raa}<8s8Oe{r_c%x4w>n4R4>(8RPdZ28_B#`n6NR7a z9EI<7j>7+ga}@q*=P3Le&QbVr3S)Ut`0Jdb@V;{t{_DYT_wGmGFW8>)pzv3V+y0{P8{Ln>?{q&3f6)CX{LAh~;YaL9c~JNn;+6-6 zU+aDpzA0{BPol={>+}EUbx{1D@j58{tL{hPNA65HQTW;7mJ@|v?|u}1i~CXd{q9HM zkGda)|D*d+_=)GHJShD2;?{>K{3iFK@LzL33jclgqwuf0ABF$;d1+rzV6dd zjQdgeOMWKpI|{#4-1Z$cZu`Dp{DtBu{C08McNG4g#ckhF7*`-1Z%X z|E0L?J8Ina{ixTaeMI5g#cdx^c+34L{N3(H;U5#Xbw=U;Q{2`Wg&%%VS~nDarns#e z3ctbqDE!^-N8!IKZh27nm&7d(3P0lFlm~^MEpB;Gc+34L{1@Gi!apVc64gQBeIp(}%7enUi(4KPzE|Axpz!yK zTOJht2kuAVPrDz5A9GpScND%$-1Z$cZu@?nxa~U%KO}yvI10a0-1Z$cZu|cG;DEuwrr-`HRUlO-QT=JvvUE=0P z;mhKd2Zi4yZh27nLtY<+|Bcs2;m2Q|@}Tew#VrpCUlX_cFbe+-=P3LU@fXXF!vD@W z3P1XawC^bV0_Q0FMsc;p918zs=P3N+UI&FgoC5Jg)Ff|2GQ1SKRsvg+JyTh0k4;`V)np<{X9Z z7PtB+{1)dZ`~k0n!XI~z!hhg(Q23cwr{kjV8^rDULgDXrj>7L3KU;nj{w3!q{D_}V z>yN_coulwQ;^(M73ct-c3jc(-O-0lx3 z{L9W!_z`>3{Q-ra;T(nU7PspQg}==?3jeUU-5*f+7o4N;Z+jgS{=zq;4THN{vg?F5z@H@N?3jefo6#h-GgTkNpZ_{y6`1#_NABErO9EIQGbx`;pI!EDu zBW~-D!hiBjiKFmK#BKdic-J`!zgyh$qwvo;N8$e8oulx#i(7x9@DDjh;g5J76#h@nQTXYtR3C+3?Hq+461V!9!loTKoscpVge_}+9}6n>_-U0*2tP0msHFN)jsi^A`Bj>7*`d|rMO{sZSI z{PczN`Voa+=^TX*#H~IG|4rv8{4uYC!oTYrg`cpP>Z9=2J4fLM#O?Jm3V*+I6#kI7 z^$!aFwsRDIawqL03cu7j3h#+;Q+*Wv0p}?EF>!l+jl%!0a}<8^QmT)_uXc{Y-z9GK zQTYANQTP|V4hnzPzI0p^zEj-lqwpJ?qww3rZT(UBC!C}3zY;%9eiZ(k{fVRS?c%on zD15;=3jY=Hm&uR9KkXca|CQH4;m0ngR@|zbbC`2NeE!=P3N!;xCaO zg&%Wc;wXHVxLv;}e91Wq{}u655AUYqqVUtjtv^xtb} z;dh81D?e)7_UnV|7Ffm_+{eH)p1ed zmh+A7NAVwYKZ^f0aa$i0|8I&PDUKSq^LoJjDE=?FAI1NKxZUSb{67?*7aytytbfir zl>8|EJ>up^;qMZ+=K>1;+_=5XU`1{1qk{^YC zPJE|03jdz?HgObw+NO>vj>22w=11XoxgUl9ckvzaqwxPHe!e&gKlg3vK7+!q6+cFP z)VTG-P2zT4p!n|)x9btb{{gR$;{QXhkK%t^{HIh0#s7@gM~z$m|Jd8p^@~#HMCT}c zn{yO?iE|Wwy>k@acaFkucaFmEagM@&+c^sVBj+gmOU_aFx16K!XT2k>4+?+2a}<7- za}@qs=P3MI=P11E9EIQN9EJa?a}@p&=P3Lk=P3Mf=P3Mdoulv{ivKJ1+fX%NeSXv} z`rBM_6#q%iQTV*L)j{DGI7i`^dmR*hgZokVP2$I?J_^58{3pdx_lZa{*ULxU zkHSAMZuwE_d{NxKZbacvdwrBT-}gEw{%5~4)kopSire!Sg`X~N=ZnH$;~a&*(K!n5 zirf02@OL{$;rEE!a~3sj>-Jy0J_`Sga}@q(UI&GLRs49>N8v&hM~&MH4)faAhp6$N z_)-1dwwQf?1FigO*tdwOeTNFaP3%Pd52cRK+?(~_Kl0@2y?OT@@oYV7S_`Ya^}&3r z)m~rktPNUA?d9(L(#oLK+U#yDwp;zdq1N(BzuRiP?#e5!y6o}`_gr@6?$#AoUUYHm zy6Z0A_KH?(|8oD*&fu!$L!Iu0i;HW^%gYzP@{-lHmG2uJNuSsW>S*C-p~lq+OiT;_AlUZl~s?%Cn-gD$j~e zt2`?@FY6=yuedGYxc5|XzV}aY+3jY0;Sl)1p%grbXwb^IF$M6!uJU+>l zZ#6aC+v2$IZE@W9wm9y4TO9YjEsp!%7RP;Wi{rkx#c{tti{rjq#c|)3;&{(qSlSkh zeNSa!e|ybdT3%Fs%oHbW7ALKgkJKwp>X#?=i<5LtIw3o$QcuO>N*c~rsBm1z*F_f| zKOEQb!*Lxy9M|!+-{JAYaUDM#*YU&g)KiL>>8Yjio~RGAb2dwJQb%Ry^h$G5LuKdm zOLM%Rs!78yFxU65>f4XHN=mafQ#c}V`;<)!|aoqc~IPQH~9Piod z!1RnBZtF&ahje*X^pq~oiXPMDS(_hx>{c;L#eOgaUDOL zuj7Z~I==QiJbpN?uqemI``N)g9io(}t}{8+QC&QA)SYUN46L#;e1c&1f7Q{@}J zGTmtKu#Toh=VcvSI9haO!L;bybfoh0a5MFrJo$e1el3oBzZS>6UyI{9zV%yi-21gS z?)_REPyMPx=sgMCUeuu!`^68dDt7ZX4ctuu+~iGMrPKz zk(qV2O{tAR=0;{RH!_pCk(tcZ)$Pxh#oTW>6`wCn%f9<~Qhy0G@JKkD5qc6Eu@ zo+Y`PxmoUJZr0m2r+MyXZgMwsle?Ll+|^#D`(Al-Snft)SAiMwDRx<(xme6{!weDpKS+g&17z5 zW}O?ES!dgn+8AVRWF~VXGnwnmgU+GWw#C+g^z&o;ansC8G_x%C7WI;4VbH`oDwq8Z zL7`mHo61?Trs+cOs9f=gQzvWXir!SthMT4fy`yr)BTl6=U(xI3^qRhLy2v{!S3Kf$ zI{S8l5@wrf`xg4Hfp=sn9&wnifj3#gUYaTwdPn7&rn7QIZz^ZQ^*h_ZJF*myICZk- zD|)>gy;S|-%JNEkU#GR+)}Pqu*V?vUl<$e@qflb{Z*=S~GCE2zNTyuTyJHLO*uXoo z6puJe*T9=BVT(;kMJMO4!$Egyzb{tuX zL8f-dej-%#w$t6#zH?f?rcKOap5j4=0h^fJ6ZYA3!O%QbuyH~wSTuV<8*jXZp?R!e z@u1VEY{sIQf&`JxG0aTzus_r7w_$%aFtLOyhnO_#0n zTiKmkTY`GK7p*&`iL9zx=l8O8r!-x*&aZ2;*@d({x{#&@s|(3kT~m!8z2G%V-P#OK zx{295b*nN_u+FT44X%vwk`3%$bSr1x#vN7s`Yj+x}_y(a3Oiw zx>K4CTj$qngk~4gj_5*~8mulPV}*-j7;^mR1+Q5euVoXndFobWB4C|a1shx$<0Tu| zz35iXyvc>+W$R99I&7Wa%GO^Q6IJW&UUW-K(BMMyvUR649k$MIFK3!vN!81z;EOxG zxEdSvukJFhxTe3nn@CW&nkrA@6>a3MOm(X>T+K#qGc~wICW(94Q%6S16ww2&^6h#fqy)0U>i>x z*k;pK4~E89{(;aw%o@6eQA1ZYY0x#Z>#FjBe>`ho8xI@UvT1{^$)*kb<7oric-p`= znbyAeYI@s>y;a@VQYDNnv+06`pIk^Q9GHjs0?#mD;F(PrEcR@|z&uPCc!miB&uqfr zrZAf@Fb@+3o?*hkGm|iWgw14(?WBxd(9 z4Q%6S16ww2&^6h#fqy)0U>i>x*k;p)kFeRWp?#P&bPc11u58kvYqC)T|9IBGHXb&x zWzz;-lT91=$I}M3@w9<$GOfNZ+xY1jd#k#!rAioEX43@=Ke>=rI4}?M1)gERz%!dL zSnSz^fq9rP@C*|Mp4o)KO<^`+U>+t6Ji~;6XC`6%2%E_m+esO_$Qiq`NrUdlMh*Pq zSp(a6*ua)e8+1xGZQvhI8`#Fv2DWV4plh;e1OIs1z&4&Xu+63oA7Qg$L;Emm=o&^1 zUD>2T*JPsx{_(7VZ9HsX%cc#wCYv_!kEac6<7oriWLo>lOXJ69jk&5^X15+ktaEHjG4BVrXfoYU7FwLfnpJ}r> zWB)K|Y#U~cZTYl8tK{Pb6DINozKOtrFP}JQn|$J6!bIZ0H<38- zCK3m}OyckvmkAvDM|nfrC~j!Wrwv*tA2*mVkvH&71P*-p#6b(?69*F}5(mDC#DQ-r zar{l!skE^sGE2BuNSz?4ZDta>J8;2xz6Orw;6DU&j|TVzrO z?orCXG)ftmW>dz`y4jquf0#734YS6!eA=K@@^OO+6L|yQMBu=ePaL#OK5;N%B5~lG zNF4a`iGvo(Ck`e|Bo2HNi349Iarmsu1P=Y9yrFFrH?-x`1}&728%&tU8~7#y2flpb zpoQ{@g9#If1K&jAvTv=kuhnh$I?IdIXAZVk7WXam=2!ME^ak_Y+^;0z39!ot%5Hu)o#r-P~HgS?@P!zn3(#_6Pg} zBhzm=`#;7%Hquf4QzIt_^B)?yveDINV^;6soett)SvRAZI4k<=In2|)*Xi@@%|#~) zRba-P3T4OkE_JH6NKI9S|C%mQ4dg=FK=lWi)o-@izm{`LTiNpSm(ynW{Pc#bGi$MT zNw(bl_QGKKKxd`Bep7Xgt#x#+x>}dl2c5NfEo+YE0!^qI=7!znUPqrWkr`QE-c+G#bw+mcO>a?k^j=jwJES)qnkc_- zZB=)ua;0g1Z)2rZ&B=UPUUlpRoytufb>ZfPuHFP^t$O?0eALRG+fwe0cA<{nIM>T} zP5VTmb$y~`^)>-|@c1>+SzK=Sc9h>vx|*VHsQ!$$TH+O^@)ob@dyD?UG;@qJ#wq9pnTXSpN-EO0z*cuhpTY+|02h}|zKKhukQQ@&ASncKd zGFpb(8Jy#e?TxL&)O9n)&a}S#n7Q!*hD(z^W3~Pu!wG(f#`9RFEl=3ufyP#$>7mA^ zG&-1WV|o%T>e5Hqb{d>8|RxNnrY%swtR0TxtrYPmHQ=nWLKNW2UBH`fR8v zn5O$z8DmG+Sxvz-eG2+?yQHU&ZXc#H27US@)O3LulfXOwd3gD3mtupKKf^x zJ{hzEDcWdKFoX8;wqBDavVX4phR9T62x>kNZ{ir&#g z(KNwbvU{=gW?kz$1E@skC6iLK!cs%8bGyYDg>ask&+}Q| z&*%9(%f!YRzv>@A5FP|Uq(~H9#!US%W9cskUTpI$J&p&@!XyN-?L~rU!w)0)@OL}- zU!?22MpY*qO-dsNuH_X-x_XG^p~96ZbdM4I*tKn9KV-j+6v>XMaD@4F`Jt|%q~G7@MVS;uRuP)eo5E;Y^blu>jyMH_)cY9s^_1^OXX2A1#}2-$hQq=!C? z=sc~8H&-o`L{q*cS{@#qU%~wW53$#1N{W!I)r)drD%dLrih+1SM<%nwgcYiWat}d%Vi+PmtVN25cU)Ht#N(vNwLqK__fH`Q8C&$B@_0%d*XTF8 zOTM8?M!8-$tk^%Qkyl8iEtj{bmD1{3NitpGF*AkiEo;n=iO{!G?^C@|rfNfeMkO{> zk+Nh%kBlTeuXk}k)IQ$PptKe8&(vz~*dvl$_}JbQ@t{u2)4WsQPer-Md#KMOy*fhB zL9$Qvwk)=dLQK(_5e4nQorp6j;@(Dz5L6$LCM&?UJ648$M(7ECar?M^CQz-RZ())^ z!dFmg$$nBGR4c1x`96q1afGk~A|Ou{4?+ZDk4keH0wp~UAOd14V_b!YKnQOcnCiv! zArS{n%mkt`wenc5NW$VgX390`(0!K6D@^=vVdnmtY>Vn{7Y z^S5E@kg#*P1{BA+NTg(|AcUv5zQ3N(qsYQ*{SGEGPLz$*<^@We1YqXj-7 zrlM}Z2f8#W_`EV51s}+7P)eZ(RgJhnUKSPxA6V94fG5(p3`zw45f-)2Z*)iD8hM?j z+;?rA#j)o9fZpI^JCU@>(uo-fV)$6ks4lhevC^=Vl6R6IGarbImmz1Fn=XFY2v)ydFTKTBeSO$1hCr{ls|RDwu_M2(0wliC>OPjr^#Bt*t_k$Uo8lBO2|xGUF7JU*S;UrZ9gsC3!iaUtAYr zr0f^zOoqU%uumBRH&lBV0xkMU3;~sNKSMxI@q+#z`~Gf!@dOqf~BY9>VN~qozq3@7y|o!BW5rJd~;4P1Ud@0AOgPm34<5{LcjS80sl__ zVhD7JKg|$m%s;>osPLGJ2sGx0Tn7U2+l~YY-vAMmZE_NZfCz={zLzcn5qQ1NhW&&@ zNG?#dARQbK>rWvaWNLi(0|8lMb!g5%K?fB_h6%SK0xOp$3}Xl^l`drn+@XG92uR~^ zGX$vO8w`Op(jOTDox-j&1eUcEmNOjmPEbzlLXZGaal%c|C9-s6V}K{EB6#0R^LGJP zd~s$`K)(@uwN$B;)t2}4SA+jmQ=;*}sZgr+Rb{Ra!e~ysoDgd!bu}))OiJDHZ40TT z(n2$-7JWZ6DMD3SVJ0Q+mu*j1ZP61!6oc2x_-Bzj`w@HTv zkM7=}U5K~Dpl!Y=wsrs&34|`>K&2yLynp+o*i)fQf;W>gs`qgwIoeQ&5&;?2*V{y( z_*HMpN2F-N|I+}JKLNds0^UfqaonI!#lo`DGlO&lkvYCIkUcoFKqxLBuu5z42-qFg z=TLKgo6M&>1G%P|VkY%^|5I^rwsO%4KjAKfE#lL!-R@WO#>@W7S2pSBp`MiTcvW}a z%WKO=LTjf}5}~xJyF7~=)G3UX_7`?o6Q9rPnN{c=eVUf_uMI+f4$La<-)VWgH2;+J zrFj33g1YSW1EM|p^93)j-KiH7niawz<2v2Yo}c!lJ4>s+qw@O-_2T_fNoS=|!A~Le zVwI|!R6r2V$_U~$NPvS4&Hsg$h?dY|L$snA{o$(&CSD=#@+jWYDxHi^`$kI(h+e_B z%WtDrUXnoB>D>7aXX(%jVIj@D(37SPCn~(phK027IuMHuMT>ig=rYk_ananQqG(#Q zC{MJwN73B!qJ@T{Xlc=6g=k(6(V|TKq6E>>9wNe*_THnuiwF~9x+nB&_PI#;&QkbZE{AtE z^Qv?SeBujb4<1DZ7kdy@-a-M9Ql$#mgnI2r3?#+Hz4Y0g_&+2q98JTf#z)YC4=HLNA=}&{7d{R2TY{D)g)}(37f;^*p0Mt><|d z`ICajh223tfYTZSh~dOq)Pk8T2i_S<1dls!=ih7{ z`e8??-vIuFqX4+NK3_iyyO2vE+moE~<%oHn7q&5)j;?3k^rrU-vEuSG5B zMDh(po5WevQJT1|)~iVL#^JA2w0yctTVdyoR7ky$ z_iQ(DVWcW%y(jMU0GQTiSz{hjSMe`^T3~D*hq+otwoCS`ytweHgaGCbXonXqkxkg?af#qj@Q3-8W;BHNix!J%9>G_lM|?_9IO*{D@k&Ah zoB$p8DOc<`>H8L+uZE5gcQjE)hXt;pRV2|CxL8`(WhzC!av#6)(Fl{far zP{65};esLcv#$fG>lHYurVyP2si~DXsfxC;JVN2e<|?+Q;u@h22fLsbMMSC-mQ7JN z^Nf=SCG$v0fJ(g%8jS*aFEEc@W_2xIswRj}In1|x(+Hdk3BHnm7o+(hw{$#WoQSQi zTl@F|3V2}2XgVt*2}K7sV~gMQlF+~dI1JEY7#oq(QAQYZ*^Fgf4Of75cx*dXD5@WCQb+V%zQ(fg`{Ysp@5r5^$HbcLjYgpFSGT}kZLUm(v@wpPG1s|69lmf9DsQz;{X?eX;l!`^@;+)Q{FVeG49S9 zo@@s$@{O3PleoxLA&vyS-f$$Ce$olShCz-5mvo>_4jjI#e+lQX|JOX(KAgjhA%{N0 zIm`-mBsiwYkzn8tP6);hb0iqkkyZs0#(2;edYMD=-vULZfG86~@94!~#ranVau$qV z?5GUC>uW6cCe)ocF|Oy!4udR0wxO>j68ZrI7Uq4t)Yrh!94MKLVgnl4FF9&x;FUDyn|DTsyA;JLg%2b}@5kjocE)la3{;zr@nxq#;hceE@l--T7-w>L8?_dgdmig+M&Rtd|MB_G zID2O5`=YqWtm8kfQ$UWkLsPc z$YT~cAvn<2fne5FCj`I0;YjeC#e}{;&SAFi=QxL6ABYr5IENLtwQu1ZzOuxTVAEDd zg7u#{A$aB&M}o<(5c+65ioftxHv(*Oa;>aid6Kg@)ObVTem-RC+W_N~_(2{x^EBsk`2 zCj>tVcO-bS)CpqO`#U1`nCJP9h`o2LBVvay^>jq+oeG{KV)yIjh}fm|LPx|txzPz? zk7?_O*kfMeJ0kYp@s5Zc{=xFTzJxl99jeQ|DJT^wVYb9V?CW#C`50Uxh#mVk2Wl>Y zd;1VV8Y16}*cn@HUx$%`Bv^=j{pqAdgpdY^pa^&|Q=4_&KO6!BE*r6%&a|BpO7NOF zi2ZL*#cAj_l+*7p9{Dn)ogB>ht-DbTf-GhW(xms0U4>E4myJ1hstElsJYM>?Q#U#c z&ByqXJ9xYdol2`ZVrfQY-!_Wdc)Z-q;E}IlxlrYAy^PJkBqtYY+U_B{i|0aj4&09A zVQM^_JgI`r8IzmtR8e}?y)wXTgqh7(*hPqPXqc>O&m1fO;!c;zi81V32dNU;7Ngnj~+ z(M?|7UY+QL6z@g-wG)tm%u=L@u|C!X7!Yomw(SvbjL=@K5f8@KY?eT~haZ4N3< zzT+X=%*PV2`v1E<0k?9O%t?E3D~I&^a1y5Nz|f${;kcCpSM?3VtbAnk$Wq+OUZcMq ziyK9^^VyxaQLdaM^pmlEC(jD{4DUBQhb#?r|h|Wr`DmU%u%` zFd24vOu?=EMUc7?tjrzzrVAv=xRt+uiB>&@TX~tcA`iFn!Tmh)Z@87GL>rqy+Z-!5 zogih0aVsaEMUdD{%!b75SJx_x@aj*nWhNhAw~sC?C~-N$7_{hQsc6 z@<+kyD!{|h?D=*0%L2(Q5X6%$$j&*mDU>jF0S?iNnI%=;^;Wdtcv#xZBbVYjj2vPt z$MJCT3X$vvj)(O>68hmdXpat6U&BE={9}RSHypHgE}~5ipgrU-2ogDJ9I>qkK_W+u zCqMZXK_W+u7cXlNBy!aF!cBzOTr2;uI0+D&W99I*LlI(gtlV@ubP6CR*UEEz8RX>l z`@>v>oSc4}|HB|BhhU(j|j;$0$v`I!yuO^zyM5kV`JU% zsBg%(Axi{|WuS?{*s77|lBS1^^JXR=YTX;bIf zR2xDuo{ft2^J~4+fJqK2cK^#ET`$7ekAsRoIAL>9a(8%-kN|jt*|dnY=S(reGSb z$@ja@qH#^W|ML!S*rd!g_l~Uxcztk9{^o?qpUa&vIeC&Jlgs`x#*xYT35-uJ=lEoH zIOCJgvVGF)*(1m&yLquk>N%B+PsVe6^0`1JaC)!<=dFE=Pqy`72hQQij8AT32hRKt zS1>*qzz&?OBqvPvc*>E<6*HVLIs7drOdc#{e3HZD<#!pMdRgWb))Z z#wWQ5NZuHF`Vfk#?7&(3-*m|dh;B~c^slQ@4I+$b?7#^R z*chS4nOxJkhJ-RXm&vlUqp-7t%VhrzlHfX=$yZ*J>lfop7Dr}leK96;CT2*k<4o?a zuQCn8n5=wpV+16~TuthF*N_wsWAd}~qv|dglQ%z<1g{ffnp{vO7rcry*)b}c@W+^( zFe5`Ul8b2;Rn-JRY5@Zig{>@{2<6Mx4n@UzF&QF($*8 zy}q(HO~^n$D2HN}%RP>I2Oom2Igyz3o?LJUGfmUr0ohc0%zw^aE|j$3{`2aW2Xxah z|5@=taHO#f=0B#qPszoY|5Of}Q%Co~{3q~-@xiBYCZ~NECOC{SS$b~3ta6;mJHHl6 zZsSaLZaJWP5o6NJr$Z#+jWPMjwol1zIFp9lI=Uam3C->6q_MV0_Is2-CmyIzwX8o3E1J2~H@5GS@a3=RJIBSf?nDqFvx1J6z7GX}l z;A^@5I_^LJ9jwuYVn)1U+CA78>-xv@aXUp)LO=*<;0ETEP@+-)sL|iVWCkOUSu-*v%{Y?^?`5V% z`1tIDku%EitS0h|*e4H97#0Mt z(e1;S9J8=sb) zdyL7;uLOu7)8Wh-Dg)P@P)^5~>{{d>Qq=}yvi8|n!9tA5(lKM^5`8fy&yMx5);h56RSQu8O?~I1ji5oSS-+1>jwV%v{oIB zG1=+c%M#el!e#Q!Atv2fjLEF6-J{gO7?aCx-XLe=OrCz@1Th<9a{kwCL+WrQPb?WE zcpqc3^v!3q>H!#&*Z(t90y`DCOlpQtUM%1y4B`LSdYYbt`A^s1+RAT3*BmDwXUGyX zV*b+EZCFVcj(w)upe2htt z-KXS-d@v?wm*xvj<4k@uA%hx>F`4!J74J-($-`r}=)T98%qr*^W5k&}9xWi6WZU@ty4#rl)PK-3lIn>0&qpy$@m=ciLbZsSZ&`C^bD5o5A+ z$@5xj49?`dHr@%In3Gq|KBDW1F`4ys?|#az7?Z0$2ozP~OrCvr&=fv>nOzl}eE0cv zPw{v(syQ$(IOkF+>>EJ?!AN#Na?AN;J(WbJ^=;_%(4CrpNz#Br3A=cBDQ^w07tgqk z9SdW+{;i(@CwTng2(k*?jAZH^mTBx4hfgZk6Y6?`VmdNc@V0#}Gx`(b}vO^4}hTM^qF z|M@U735(^x6)*bq;_)~Y?E0gBlWZo`->|KA^U`zOo+cD>rXE4SPjHynn;*~ZN)Zxv zjqA*v(lu>~aUpEue)a1&N#IvLY+F3_>eTOqLR!RF2Y#TfN>%^2FPG0Z3aroI{94nO zFQynFl%})$t^akb1QFy^xqo?eIqxZ+F_-Nvm7<7QTL|?XwjlG1%1BJ!@OJwvx6p(F z-e#*ar>{PJ>S-m^o>{w9>qlij-<>FH$Bv^(`7m(Tx~koXkN7obPx-Ky7WD;I#g1iQ$yG$jP25W+20%=;KUuXr{)A1gSF*i|ON z-{wEL;bmZt6Xv?7rr=>-`CKo_HxL>e?M1H~{S)r<2aZ&izkzva>HB55&tP79c6u+# zbv)MYu3Uz2gKLhW{p5O@5Zerl4i)Z=Xoqm)Q})V8<>aG5Nqjy9x0YF6=mg)#Wg)JotEDQ+oggz(1*>Hw=+l2 z8MhMZAOYLT%VHLgzX6jnLdm?-d|2st#DHs{@T1T#^OWiR@2hw65_raVjs@>5l=vvU zh;ET!YxFYX+ZU-XH9tcrR>AFL%u~jBayz~zc$H7jM69A0gPBM2&jdNLI=JD2Bda@8 zKO-p*qOF|GYJIn;K}1iYtd7HK<;Ekr;XGqM4y$MXDe=7v!39T4nEuVIZqAtN$ZB2x z8Yis2b0!4NNsmk8X!Y~sH;5guMVX`3o!V|>Cl6vGht;fAQ7d*}-hT7gR9P|5OlQ*! z*PiIM76aR+Av03ooRDH8N2OVVwkP1|J#)*_e}h3eU3(pT=Q7}clH;G>FPM^!;XwYs z_tj%?G5MRRr97S*bS?(1G43(u*gN&u+1RGnXOxPBG#q53;9X`AnDEbVUkoeWdt>q( z0(;xSe^w44_#CWgs+-XiOmI8(`tF@dAB?xATT>^)Hck$d`Tx4FQh-7Cm#<9u61H)2 zj~^UV#&c9HgDRGTDK>38UU}VDH!MQ<`x0 zGAB5)H*>Eed;fQ;axkvF?5t;o;wBr?f6p5jdtJ8%O#;N@YVXIN3^q=!nBKp>lqQVeI_4LkqQ~23FOE?X zaLCm1LMuD7r@Wkv=K~pUjyQxNeC0oydOQa(>uvokIw*bc{j)U_FnLc-|6bD*lXvIe zrZ}u9-F@zzX?URZ+c#C=AvA)wQv|%w>PC3~&w6dd2Qo>77>(c*wNzT*>h zd{7|;4UwU6Q>EWHLy2d%QAdSwUdk`a`+Z(|Wl-YGzY6?5zIEf!xlOx2N>9ui*7xJf zm){K*_Bs9xH819PpVy{)g{ynPA$_$tpAZWvDPQL!Tgj&>S-wgtRLO)YAWo4Z39kyu zYrE2GSC!YkMz6nK^HP*k0Qj{s336B>G%1psQET zQ$ANg3aUqkbEeR0Rp`mEkaN1wbIl=e^UoDtXol?Bn5wuzoArYm61!LgQ5P-D6XlgX zJMMA$O~XT!X7Nr%DuAnan`Gofgq7PMF7a}n;FVz~Xc!Ut}ufg5bV<%mTTK~%xNIiU$uT(Ly% znG=~nb=H-vS@TJJxwJ>xkci0G)rPR`EsEdxT_cA4C0(;=SK!#mJk|JkDm7a(dj_>2 zJ!js`IkTyV0euvG6jYDQ++6LvL4Etqn3%zj3jK09vCEbaW< z^uF2I3ug7!=FG`TPtB!r=A`!S-)BIdC@L~4DkhT3h)Rzb&`%SQ78w&A(LW|t5uM&I zH8m!3Ky*}ebV@{8zm!x(>cEJ8J!j6&oil6pOzoW9^x3&oYNjS<;>=mv>~t#bsi&w( zu__hnrcd^qIodw6=FiK8{$}K)&ztvR@}!9o(Yfh)x$|?vf`P))tuBl8Lf-4YkqKucit;GNK=i@FQ+Ufhmqm zl*g;7ExbT-vMecGTfr+NP5!-dD3{2T+00(f^3v z(C_WY0q;MX??dbb3=IR_wDHR60oWv3oTM#-c&Gss7}!{oe<|ds`9d@k3Y!)t=G*oA zR_*g$B0!1n1q3L*-zECosG@bEMLig#%PWE?MRA^_-CPrMfH!!D9`0h}zXYHe0#~Y? zqe3od(9ag;_O8Z?kUDd4@;-r?9I!)(8kn1;0N~(eH2n|&J3|dmOMD5JxzzwP4S|3) z#svfLu8RYd7!sm5ul2du>~o`${CBetT+ah|Ejl{*5TLRO07;dD?wk6T)(5Vix=Hu~ zeGO$!?%N7gWcza@^S1{)Z!0tv>(-EA;&y7ome?v7<;?HJ+cvssD|Ovw+&k4JO=E?d zwvBGu*09?6j0`GL=3J)cHLiJUUgf53xwE$YjGhQ`Si5WoJz3j^xM>^YrY**08%0LD zY3t{vEz(V!!eyJ1?Bk}bmz%b5H*G!aZC0DLKW+BI2mRaTn`WIG$WC_J4sDd1w%-C> z+0j7Wb=g)$-f`1*+fCaoH*GiUZ4assSMQfu^lQ_JInJ+<6;emJt!?FQ+FrG{S^Erj z-iK6l*)9*(zf|;{o3@MgHftZh*H5gUccF(w(JOX&Fdh>{OWm|BvA0?KaG?jtu&mGFWYH5v@LSew$M%6e3xximmrP}$&`jd(@ZrX;sX&Y*9v&!^)ea8BE7dlCV z`w3lPa-t~LP1`(so3)Q%=X%$Qa_sNLbfgu{b*9H$ z^XJ4+1^Iy&)=+a@!a6zxz@$=MC)@*Hp4S7 z+4X1Dk=4H)-kW8AuS45RH*L^|Re#n#T-d?z4D?}rud_a^Z8m+MXRxDc46ax6bAyZL zS{Kg~t-|ydC{Zwd2In_@^qKoGW7tf{6pKrf5(kcDz zjV_*7xp;1{KF4ft#tvgP2S=^Cba*eq3(I>Q+7Mn?+R$aiteOSWhtZfV^%t1OiZ&}$ zVj>X!fR~1{j~`YWID8&q2>$tY_Rl-KAK?o2ewy`ut_}BIi*u^uy^SuOSGjm@aQ^&F zt&8W0E}kn~JeOLZW4bqE&HdSs!2zqC9Oxjdac)DV4_v22d%F#+aqYr}X8gfz!umXB zL&Gz`i$i}hdw)UTJ8gw=z0f8L)OWnnvjqMjepHzMeo7d_ygOd1$neb9dBqZ&)Q7@) z)HN6IcLnFi$_6C(cM`*bPT?m7lx)0|o=S#EZ;6X}sE6lCO04f4*XaYM(q{J2hteOZ zdA$9JA9H_RkGa2s$85jC$1Gp*WA1OoWA0BF`*`DlOnJQhi63)+UXQuIC-{km97018 z#6uY$p5P}^=fl(I6a0L*`h0?)4_BX0@bls7^9g=FTzx*l&xfneC-~`9pJQ0iJE5i6BU+uNpiu3jk5S4xFy*?4jOCqiCXx|FY3MKg2$ z=L_A>-+y<1YAN|x`=cLof8~$4Klpmq$Er`Y>@nN#G51&gnEO*gSU%SHqaL&U9&>-? zkGa1m`1x?hqbK()W;}a>pAXr7j~UON;O9fO-($`PpWx?1w%-%{L<-eV-1_|Q zpZ`5(JfrWwzpfysbv^ud%Xc6685Thhp3qmDdbOmMk{FwR8Z`UU$g6LMXRYnjGB+sW zS^o2rOPs;o_~>r%#r4+dF15? z(qz!J`v-J_-iN;H`=0%t^OTZL=@as5iA8#;u#G~yH%i9uBK={FwWqNA9QYu&5u)GI8AjyhnVero#ycM zVEQRP6)beEA0$7UA6snnI6|5%`ayEEDI|)mOlYXzshh-}F*3eu-Jk=hY+i8POse!Q zEVt;!!jpl`lX*?mj2lRI@ye-EL04ryf|g%bR#X*Q@nh6_9vNL|2$8GRQ{$_*5T?_S z#RcEV+6g3=Wcl0V*0@wxdanpWaY<36Gh4F#1GU#GJq=V8L8OY~N>K-iQs`iLe1#nK zp)xXktlj>4s<&{yhpQKo7*PhPtO-&7}4KsZkQc0I{ly5R`@M!<3PYDFa|P zl5_k)_zQmQ^=u}&M%!H}-6838-cNABJ677uNbIUri&sbvoQ6O_{XfxRE7jp?a6+)7 zGe2aHmw&_$)!>B42iWwSjK8jYFd8|aGycq&=g$AUV^csjw9m55rV*nJH9d-9$7d+8 z!@M-qefy=Ew?hyg$D{#$1d?ot?%7DA@seykd_A^|Utx>ww3OLvJO2)yj<$w#In&4L zUZ{^S71PJs3W)6jTWlLV4~<^^`*pf8>Gd$BOP!i?8s|Dyc5BQKukdDU393V}&*W0b zf6SPeC@e5#8pXt05u#6KEd6Z8E99>$c^{Qn@zHAof-hqafUi{XiM2d({irT5&Pf$@ zl>Zy4nTE)+RI$&@)nSSqRv1z4x?rU85sdItocaJ8R}de7l>t1BB1pkkzSN+)Cg1ww z)UC+ZQpF?#&NKq&RIwN8e0;zr73#mB59D2|HKS|QtHYG#0}Gq!kya4gxLBYUBM<}< zjNI+uU#r~6kC5DCXXsYW84K+G@|`T;g3OBiX;&m-EAl&!Ig3x*_LCNv#6(0QNpSq* zJDJaRJIxgqnux_Gwz2r+6Wc)CprC=5GwcL+wfXU*ikoPQKL)b;`ki)R5Y?F+?702+%{`46@YNmE-S*(y-!`5X||Lz`Pg*<^DHZkI8TFkBeQ&e(eFF83W<+~s%(zEH1CBI&zR`*01uLi05H516S?i&~wg|>kal3qnIN_t_I0KZRk4EnbE#s5^tYS(oiS3YsyehkSY z@`?5!tJSx_tr5mKi@W>qVKcMN{Fj|CJ-Gj>^#%wuE>0625wGZ_p?*gkPBad z7-kqe*8^Qfy-2qm2)ln+b9{0;uQPGrXh{6ln^U*`EC=`;0T``1W6AN~Q>)>>fe5B` z*8rKKs*RUqGCdl8^cx`u&#=OiVef#L6>hMa?VG4?+q}XwW69AOOGir9DG{6#09f$f zdsZ|=1(#2mp`kPMnIDJxw2cpjp3&w$K|&2h5Fzk+tN-7peCAkahF*SRdKra!8O7~o zuca54wK{e_`n9N`u!% z&KDmz&MbHr{ODqgk;>VagZJqQJ02jum_ibF9HKen(C>&lyI%rgoirIEs{xVAbO=#& z9IRBYTnhys*ZIKx+K?mfJd-g8bqjs23qyUXI4GuWvxX@$A)LhZ$<#dDj{|nueqwrX zf|Kf%Yb+zKGaLR+j09<%U?kmJd8vneo`dF2ZH1$l8Vfpm4DtlT%B2F2ph^52mnFoq z2tJHyVV|tF&F`SJJ|HC#tqYnKHLq0Ie*IW+j_RjYd@L24{Jf9*OGYcG58$f#I;xze z=O!555!NU@Dunq4u^?w6-G1DlbJBf0Swmz%y!Ob5cZ4H{tBVgt_vk^XL;NrF(nTby zsDi{2+ajVaa*_eQfU!=Reb&4NiTn?(0m7N!wgsm3VQs460lYL6+}#NOa%Boz$BRfT zS9MixJwYVOIzxM^c+~pQsOB}qD!{3+Hd#)s1X+5SWs$eB%3?ZI&B!u_Bg?aN`*@J0 z+#<^evn<;F{awoPY~)0(1+myEw+c2`+%QlfL+gNHMwzz>NtR@0QR$I)Rw#UqA6f+h-^RxNc`BSw{V#eZU zJGg$9chI$;X(A^n3Z+IBsj;lvyNj7VEP0SW#w)EcgJR`K=_pRDG}vRM0mVueORRWW zV?}rBLz>S&+EaDHUm^%B@E6B+lGF|5y>P8X8>c4W-HB3DVB6#VvR`RkQ)9UmB91CO zbkf)e+qQ#ab>m=5K_e+I*dHvcMS67=E9jPLYphuESDQUgYvYQ5$$+V_zsv>VxRlY)zID&E^b@L9h9pX+NJstI+V}- z88B`%OBR3lot;Y~jPq0Q6NOR4jw2z$J zvq?g!jQJ)fs3`>rdI+A`kl!PLTysqpJD==K9tnm^ayE~u6Wf24)OhE6knc!z(h$uj zy-Wt7;A9(F+FRayFR`phIB={&GD$jy5-(R!$2{_QyKB689cWXpIF;}@n#T;qs?V>w z9EwLZB&x{rAE-Db4BvIsI-YDtfTi& z2;E%dBL>-y0Q2s{0JQs18~U4niOl+IC&>VD0L(aqv04(c^f-S>LSJ=W|Ndf_bZmx6 zhe?sB>Y<~$RMVT=HG6FG86ccXfHpn+rk($!m+li^+3Z<9UKUlVL6zdLfX?JGFhdfp zT70M>AdR@<6Ddq-_Kp;%Ht%GVv&5{N0@HxH#A5hC3BQJv-&GF7js>JUw*SDNtt6*J zn>khRoloXkw5vH-FKD;Mb*suJRpxE1u57@@*$1Bwm)LXrl1DHfZirZ-3b9R!6|sFM znyLe;!mRT+~&bpct)Al@{egt#kh6t@L*e;KMU&|yv^9&=;@3GOn=f#iT zMoa?s$86+gZv6Od#K?0G8?kE_4~At0_82O71oPvDh+GrP=Q{PB43p$q+a%dt4E22= zOuIeCR!_RU;^EJckCl65E36Y`W8FUzVWQk@pD4SEo&GEte&0O#@y8F7O#ZR>+3I?( z{P^RCk>?)q&yxb2o>6{I>I-xqaM>dTJ#UwJX-XRM=_5+Mw7Zw zd#Za9Qg;Sk)}eW`I73SIys^-q{Yrcp4D9w9F_QF>{y2_?hs@{$oOeLcdt z^kw^8`tl=~OV5*}OV62#O+T8ytfljvSPwbE@25-0#(j=(pN6cc;(iNR6nIgVqJEJde+zt%()EP8uelOoZ>K=Qw*`J zyZKtz-BeBWPS&U=)9ok08g-R*?QkTs-frkW09&JOh)C0gNFs$A!z$~_dyan}OOe66 z^8V25YMarvM!PbVyoYU9YiXkO*p(T9+}m{^%XiTk$9u&qyanP^bABx=PJQfgT4j&Z zDio)_<~Xg5v&Lys|0u^eod$7Q3UNBi5~pj-ahmbHH&&0GCz+~%IMvD0>~UJWGO#2} zxWJ&2p|YjBIA!PZXf5&{aSC?)X#H5uy7&Qmd>%mYY1t1Wx9x}FH63*lpO*bFd${{y z=>GVAn1+aHJQSZ__qR^`>puG1By{LrWL+VDXz^*AF}KF}yg%HH0UKIz*F4tP@?z(> zYu_f&dSaAGnBB+d1COb&!h%*lX(m=%+n9QP*!zICWtU7^k8SS*e;?GguxPflq$|Pf zB5Qs5Qu&cbkS|3Mr;a7@VWWSm$(LBW6CNHO+gHO|pIsjj9^1Dyv>w~u2OcM~#vQaD z`Cd7bdBxkZ-c1MXElu8wuXoeQmGJbqYF+YPU2>&my&L=>Qs2|QFMz+xwlARQueN$d z(O;pS5lfo)1wiIx*%y%I?c=yF;6EODCfFBXvhE8gG4Bhg`20G)FW{JhRCR<@BwGP} znis)!o2y_-#K3nts$b;Gi&b7rugYSTUh=E&=p`z7iHVBSkg2eJK$u!}?=|Fc4tjJu zeluGvEE_Cpx2_7>y9VE2QMx>D6=N$B z?a|kB_UPZsR=QScD^yqgqqdUnX;rpj#vS*zV&etAw+F=)dZGQ3mh+o%ZpZ3>B701U z-mv@MX1o6-&SLzpyEFgOas7`>N_FCYM)1FDEdSdx29+)5S;`iy`oCjLvCq38ykyla zAf^anIVYAz*<)!Gilta{ES+0xi={E%j4c1*Y#IEIFwe9dQUkdqmd2T5Y06j!S>gf@ zfGo!8_E_4?$fC8%;wqM+fmiE%9c_GM=S^nZ3t}cWNJdm+`hN}kw)4#)~-;NQLMH8a9qJZN3^%H03Q$16iSYoVZ3Y7zq_SDfxXAw5&5f-2+dfMXIbF-Lm$;GN|w<@7I;O zaJ~HuK(Puflat>?gq!TYR%`WY9c<{Qgd3*z8>{GIqHpjEA%RCH@#b?*;?3o~V7#o2 zOB7j7;;EI)is0#kgB?}`i&a&kc@kZafweS$5Yl}Fp(rNy!2FIUk4 zffPku;yN9g!|G(&Fd+b)*ynMM`O7E=ow5Bux@S8t-Y8c3H82aM72UG7M`xjdQqwpL zTHroVSU#;-fS%5i!VmZrmQmGkp~b)v=vq))4!6fPhY;<7yoYZ-DzSNltt6tJ#T=E` z&B@_*d*tbM)p+YdEO}3!d0sSayF=b1kO}h?$E2He_9?xh(^J+dJ=*NDDyW<$I?D~?c{S4(~ zR(ZRuTBV5zo%tMMc+TyWjBgA9AxlS?df^8l$lAA9 zJY?$2j`7&80p=kDc`hVQL7uS|E2!>^Qq08)(=Jr30PJ~4wzW?6TcvH$$*qTOTYCwh zHQ5%8O!u@cdtZaOcIY;4!KM{I+V?bA$CP7sA3nzN;mvj*PWqnl;Tg_+SjY8YGFk7$ zha16%rO1bsY#&~n;sLt_L(`;3QIYf><14#H<*H6t&(+xC zbiyh|#sM4|yTiB%wKtQs_K{?c)1;RiWYh=Nf{e{G=SftuY1X2FZu7b2r~oqGQZ#TC zqeED^gUkJ7Z2YyC5L#uNJ|YaZmKIu5oT5p>J>&EN?aMAXv`YIvA{@4sB3hGu?+b?? z(7x;vMXR)LGE9IT`Z#MZU9=|qqN2>hj$d|3qgC4Xe({-|U%1+PW*c)+j%+V=v?d$h z7e?RLeh2#)%#=WE1(00}vUWQbK;DF>(W1FoySrt*X)b`kI1h8CSk`_A>zv82yKTEo4a11+vmZq=w{2V%Vn}g=9K@E@xBhbN>ik_Aw{Vo zGhps?$Yh`2=fs9jv7F;z${iV$dC6Ya(|k-JT>(q>%h0|Dy|e^7$Kg1y8pq0I0^xVo zI*x6>qRnTRvWTrHG9_zQhGS7=7kWBR3O_1#S&AYvIW>`@*X?t#*U=ozToaj=VXKK` z9A|1GpL23M%k>c|>oyU~^%0uSe{h&@5~L}#Awrm`BW-sEx4aqkqm^UY=>a-^yK`jJ&fc1&dsAn_SNKl3V1m z?j^U#<+zu;BrI3U$VGePu;ns*G#Sao*viSBgimbs=dl%g{Rqd!w4)MqACdWvodURa zXwx3F#&FqQ<46@hV?Vd?i~ZcjFHX;Gs1Le1w=tW0ZX?qhKVKy$Gazp^TDCvP8cQ}N z#g#MPHE74fIEQ;a2LUcv*ILl`nz6>q40KaaXShMreA!a+o@GI{FTyE7I|KKwfvcsI z8dxPj_Vkf;;oci`*S`PnF53TRIZrkcKTpQJH|W04ll5feZJq^7D20}zR?4W%6h9qV z3;ei>J@z@WR*92}!fP{oN~LgGPAM9Q{*7%;uR6NgyulF!d^?O0(49@$W4U?nqHQgx zTR{BJzEU2#Y!|E=Z@xmyY(M;*+yy&!UO?CxV>u`1V;vKmpOc$Jw@+l9lk;$OPVT~t z>KCdsq_Ea|bB0pc5Rq-_069*)ER7{>*4tf?u=it2e2?a-W9GZU_>B9;!S?xjQ%q|yWL*Q<Lo~_rloWU~lV~g73XWxoAY9FdbnIcfSQ1L2R*j#~zD!P%QeGvp6 zSq#D@cX7*J)5PZ)0poNChEYWd7WdXM=ZsRt$2UB~IR|;wj=xtC{wgi_E4AXUuIYr6 zbC6T$_DO)hhpaPlGyazMkHK~&q{*aJ?Q`>Diw5N8#qaK(x43rqJjvH+pMoh7%`*x3 zq2iM)XP0bVGM-5&T;VZ0m%wwSb&ckOfx=YJYG%U?+(`Mq%a!yJP1Q{K4L`jkwQ&~n zFCa~k2+#%y`Z}AbZ@Kys@ zm=|HSu!+aEaLMb8g$c9rEce1Y&co8)wI@iL37+?g&0Z{gLbBR(UWBpa`UFlKHrnH`5yhdyJ>3e6FgMOT*bj@t9e=m;J}j2u-Ie~?xxj& z9}uvBZ^K+3SW6VNUs-J!Zdqv~_`bk_=_DEo@UeCoy5`Hf|lU=|SrQ zkJ_%z>I28rq{w=cTls+KZRWw#l|9*U+WPHjtr5UoYzpgLY~N*|3NG5?^J1&16&!0OeR|l^-(j6dW+ZZ-{|;s8;`h}I z%Ip|jhA{enRx_9(>L659pg>nbx@-gWcaJiPsbE+V3~J>-0|ZKU_~Z z5uweuUe8`Wpw|Se46f|K_Cx&q1*+HjDS{My^o=3#O(F1|D{ynlGW1ml?$%Q9yCfg- zT1u)oCDB~JeL*XN66HL}^fl;i$)H?I?oU6nJkx}7e|Pccm{-_xC%Bv0wn`@4qYK)z z3~o9}lnE^Rx}jQYt8Kw0l2)HQTyM{d{T$ZGxV5eOT_-CETav>Rp^YB7s0e+mDsTZ@ zda}-X@Pv1}xVW8!0oNz?<$j0d|6lCO&KtZTF4=hMq(0#X%fDzzAFbe9-5O%~JWtTg zc9u$(A8)${N%WTO9wgCQ=w4@``7)F++XlB6#yk0bq#`(E1c&}q)Rb6d0QwiwPpG0fe+07!;#tOP9v z+i}c%ri!yyW4}G8-0yVG=X9@zbuHMsSHrRv>}sz@BOKQ_qXyhk$? zv6@Ef!fr(&K~D1$)`rBjKc(`4f8nxi75UQD5Pgx^+rP33KG7<;E87oa*&dH`r~4f& z?hY58l_p!<9dau#JuA+uC)whBls(Qzp*XL$#Ce1*&Mk*~x!;4f0WJt)AMVX{akzKH zdR51WIw{RV8w8FS{a%JqY!7dMWiR0E%?ZpXcE2{qVtuKzb#xZa=`6!H@iTKSi?i-#?Rqd_f5dX2bSB;;g~*Le8Ajx4xjtq zjo(LPeEB&QY~%`o!~2UAg<3MjvS}hsQCLo1>|1P5h;rB;Mw=HLILF@nL17tltz8n_l?0~#sC;VNfvoCDc?!WJYvw12+B93>FrPY~`DzH-X1?F% zn$71jrth`q<0gD-3))Aw+P0_FdafS7!_so*Jfzlk=G^HWmXv;`qEmFzzEV)g~@s};RjoogJ=e5}9JIrgaw)kYe9}Bg+zqcsga*yS9 zn_N!su@u2{lzoq7qIo*%c#oy5`j#1slcZ+3rbd8VI+@lgm$y|eCgz;nXZgfte?FVJ z!%0pJr5I`lXW_e)!7+bbY5Li|m)sh6qqf@41CH7X=K-rNcRk#&`Dd~>Q!ZJ^IS=?8 zOrzjV7L)Zpi!`%;Ca=WHC5246M0+MAF8UlCs8o|*TbEe`9dVp=j5qckD{S7(u9v`u zdGnWRBOgD*xyu5_mDq=SH|~hV=om}+^m&rcl@NC(nYEJ#SleFPx2qIJ4HE&Fj|oIfboG z!Q7gR;?dciPa^vmDMDkU#aAENoli2ptkGPOXn)!ytCr=Nr`-h1v@@=m&w4<3hQT)n zq@iLoj516y+MsLe1*!huz0+tw%6dV{`9>8fA~yQ7B&d z5BW|gS!{btey)JhgW6~oTQ{IkK~6{2nwtNwy)yxCs!aQSPL`5Ffq;l<3m6czY~9!j z2&AlKFM?7wo6?0AONq2#WgI6>(ggvn3RqFJ$e=}Ktg=|QnIel-Tu>RUh+0vc;#w7F zaR2W6oaf)BZBBByuJ?WC{pS0)fPedY{?BvnXFY39EVuZ!hTH&UMW667eB}Mj4UjJy zKy97&y?5O9*(^WLZ6yEOj+^Mxf;D%sd}+a2utZjiQy%{xJjuuR5u2LpJn8Z!ZO%8_ z&u`yjjW+AGpX&;I-gZ}O698* zaJgXpzCPEry@uGn|AM_Y^WN6Sye;p0FZFrfd#}5jkLhO(mEYgiq17?Hc;CB(_qQqU zdkqrKG;~Yn^(jiXDAxx39Pb92d##?^P_AZmGfYvAj|{=znc2i| zYPA`sjcb1Yi{E!@9RcgYQ4Jy6#w=E$zhf>Hoifsh4!RG@DNM z|4RS1zc$^nuC}`7puBBCJh$Q9bMN*&7}j)r)wjFPa~pV{(~~XlzZd`XU z$k$!$As_Xc@XAju z94E@5*YZ$erb>bREDqq8xRN6=t zr<Xe%`m~n0K8@$PUOcK| zYJOB@=}_g7v{uJsE9iZ?bd6QQBWc;nqbf`8ZhJgdB0s89ut@Issh?8E(+m2!0gA-y zzNHJUX+62))%rcq+TQy5)&I7BkE^qXJfK|2>ih?)b7c{C`VB2>`xz~3`~P))&gX~D z6!~i9$Oh{0Wcp$OzKYk};lG5)9ZvfxwLYIlX}B)Z_L^n$b&=-hH+EmQa(C8B%0ly- zp2=!|Oun;ae(&^okNv{t_t=-vi3HE@%vSSTyvN>x`K`Ri{^NH(Ycs#&XoGY~JCxP&-?M+e#Y5e-L( zqal-b7Ga0vXZsp@|B0qfbx}Aq$m*u%oyBK|G#nGQUpJf2?fhIBw7eZTi0rzzhH(R>e5d@n@%dysCv&ru#Ke+@<= ze+@>X+H1L|F3$kLB6JRc!2V? z7w_}eUeu%Oc114au2WtAuaMXOsY5KJeE-qA=Ubg*w1;kyx+NQ~(4FNqAVX&z=5bT{5UGsy~w@%NUd(7C2Z-2pu$4M7*aEA=Nl!6>6zraT4KtH?wKiK zI@hJbv8~b?7wfVZ4GSOLSQ0LNe8q58P34=5-0luz=y^EfmDW6a|JaJm*xMZ<=k%CA zXNwNX$F>G*m94JQVo0xC{@{^H+on$jP3dng`r^vxb#E`xkBdEVzU20~ z=JREa(#UdM=Z7O7>R;PnC^}bS&qkxBc*2fh3!yz3pUN$$Q`)ekbsJvm(}w(9X7Dz&)7=k?VE}nWTOKiu_S=@vI<~d6<(}w>EYS=D zyWa29md|c{32k}G5LMIAx-FwxwdIDGsvJ>BXpcv!Eh|M^o~VqK+wxUv%kQJPEq@;L zyq()JuHk&i9casju!ke>imJ7R*&WaJwQTQi##r3_Zh~PLdUv~NxH-AI!CKQVd4XY& z@pkihw-g=giA0YVS9bZ#`noNlcyx(lw5PLvIW>jfw*AiXLeva$k;CK1_uNt0uYO5t}*3Jb!?f0Luc2Ya~T{{1}hNi!E z{m&Vf@ArGY_N|%vandK9YVE7d|160;G1%QZ&#I5B$a9tH3|);6N9v2P|69B{&f*cf z?cED)9$nwAA-9`*yk_r6`}P+cu1-f9u>W7#w@lgd_-nuaym2W{ZEDBFPdXiu+c8dQ zM|XTho~=@6h}H4Yk>f6#4GpyPaM&w2K$_n3bz5yzn4{q~dtB#-EeZWjbPB1d49h#B ztlFWM>Bm&0$9iged}ehY4I61LqqkQ2Yu8+PzE78rHFe^Ksj_oe`}om4E<<(q|TOLNk8H$^w80!`D9k8yXdwOT}h?EW(*6}c|R_x)O!}u z_j}%Tm;L*4*S>PsFo!MQVVi3{qi>8Yi+3E+-E6QJJsbYM{fg?W(2}B#p<&*%MxFW4 zWW6q}rmw|7M^&~G+qlr-p&{StSp+~v)vMQ<&ybN>B620v(mxQKNvi^ zMAF?%XV^s)8qo3Twv^NwhkfRxkr?BdZi&Y))PjR`H(8(==ij; z<2&*pZ~28<>+M}!mP$#kL_?Qa$2MuSyO*`D)7Q3Wr{{1(=jPa_f6#rH`E$n+myROn zanmPs3ZUaPc1cNpr)TKlelHx_CUrXPI9@s3S)ECJ*ZE*rw(aPlZGF4wN7cu_TGXFA zi2kJ6l+RjTV#Jq#9Go0-D8vu186HwZwIP2I_5UTtDRMs#)-@j!hg+l(vHE2D+4(N{ zWFk3sKyv()`5~F}&BgOO?w%iZbbc4pJe$dy?9!W(lDnJAl6#oU^Lm=h*H}&F1ra8* zHMy6`9F^S1WKK$Uxy%Euak(s0=s#N~{bw(x|6I~;{$GET|2LlH|IH?|**uSau<;KI z=m(d?KSfay5YjtCB`R;;Y0uFa)bVxRr zM7k%~Jm*T1b;S+(qiU4|k^;+k1Csr9IKH zr)O~ujXC+WaQB_&TWMhGI>gqwTnW2+_MkALlm2;dN&WZkqqUbsVrX_=RESo}T=c0=JvbP=V!2f$6^YDEb9=at+WR*L85*tM zMnNvt%R$;hl^~}qsPAC4<|U^{VOinkzG%otJ8duAKoj~(-abx-Lu2yX187_K>pbT| zDk?WL=$B6uXVVv!UzM@Erf)-NY(-vn=>6u~&Fip?4ZpSDm6WNPGi`0y&`{<;rlj;C zS(M4vNzPR5n`y-B4{@gbl&R!N+eV*EDd|J=D3gOTRfg{O%{21EepgYZ8qRc|!#7iQ zdUPRWs^m;3LTh|8jrw{yXL_E)d!W)cQ(=0{dae^^Y6^YcH`CR>o#9L;I8$|lZ>GBR z;V*KXI#VV~rxU)J#?46|LYYo-rl)M1ed<)7KH?Rw6K9I;bkaA|`0GF6O!b^;zr#0E zWBSN{a-BF+N~e0?OcU}vLn+f~%2aZ&(l^tE^iij|PMoQ*(`nyKlW#RdQ>G@)RMX&_ zDJkRXPq|KADN{wKCf`g`UiWaO^PK5f+Z{f2O34`aCD(~F)pR=Vn`!EYh8W6pfipet z@XeH+G5)P|suO2w=ybt1)3mSC$5EzVIMd%MeKQqiOgK%M8aR`s^Dn-cX8mRuPnj$d zWqPZ@H&b24vVnTb8fuv==5U4_ha(j zj{owhK`Za-*LY(@*1YbY#&%q8(Kj6HIr)yaV>jMqE1hTYSaierQAc@S@H6Qp#%1!p zAe??aLYKT~_5ZyCeL04`BBHC^uIn)|aZC4Z$#v3z-pQX!XL~0%Nv6o;W75%H^xe(7 zdnMl{o$Zx;yELFzvRlgRliVO3?af!EPV&3LC)*P{^jJqfg-oW#(GB2l_?FjjzBbrz zT3z|Pj*|^0+8^a-uU=D-yE<=$98c}*SDI^g?6ovH=x}9CY=_8@1f6HaVMm8`O}W-H zy08a28~QXH;&u>MP3zxUb~15Wk#WpSOW1-WeXOy!u5Xl~yN>QJ%1@%vfyTWd-Lt!g zeQL_JTvkY zPW;ZhzCL7-by`*FUW3hCU^nQ}Dh%e~W8b%Lwn;9!-4g$5xaCl{l_gc-yRyo+&s$lt zC;XAD%EfoCEPnLT$F8eh{N~Ey$1Z*1y4uCvZYX~I(tXBb_HK?YrZJXAd!?bMi&6O@ zZl*JAf$k@}`I!Ca_zilgEnEr<-RkNhvFq|7KR2%X|FMv0(YU z4O0g6+?v8a3^3_4DmI$`oSZ-TweS(8C-1)|f7WZG4|I5^-=d9qowpAt%L;v{i_3Pe zWYT%V%;jnQ`jt-{+WkQG$$r;toG~>1K+-!86-Dd1^cm$$dZ_P)lK!^4TuJ(Vp=HVs z5sk~gJo99ym!cz{%Px@KG9_JRTjo-Js9dDasWj<4M{;_VmDM`Jv(wJ^dvoK|7+b%$ zHtv>eOAk2@>pIbU(<^5ub~hebyIOad+4YxLZ*6E%mwmd>hi^E!>;*dSD094J?ZDW?j9UAv-81&YT6*h>y12t0O6h2}hR0g=85fNY zo7}-(=dpR_-J=hw80UR&JdU4K;m+PmyCp0bA? z>S7pc`_k6!J^Kc47f-Zpz3p7$$mAER57%}vk4>67a%Oes1(uR4yB_9VES}o)u)Wxu z-4q+PAhv93*TX{&*LJ=l6e&AOm(_v%SW-pS(c{)r*R+7NoCjOWIZl74&~*!Q(V z+jjPCbR2h>@~Sgy`PgquVnfK1kl|)JIE?L36qy@7t1meXgN^H|Gs}x@VF$x%=sVr0 z2zg!LW2H_dx&%JX<3DLO<+z=X7mL27W9*RUdt5THTK4-%>wgJ-ot#b|=`Y`>-PQA* z@9uFH=^|Ti?~uH`>LcAf7CzOG%jE0YmlzY}^@oF?WkNSP*0Ru8X1Gy*8NGtraO31k z|K)7)Er-`Xl*4tNY?tIUZj{+M(Lx=5lkukFZlt4(&otMS4!J3ha?M}K^JCX7b35x6 zhSAHv!EBYVK-Xq&}hE zLXhuwMCuK6Jg2AQxu{R+c1Oeg$!Ct@;P;&~vH0;rT__!6Q)~&&XZuFcLg8BEPCHUr zPV26xV9kbz`vzF)IJ=+4%a@b6t=rz~7$C>}Z@K@XyQjez?I`3ut2i%ib04Pvzc)7d zgxz18{Tt(IxXX5XD1FtW)n8n`&&%;?syqfK+kf0@m*~i1f_xaTN!LS)OV;bWJjUE6 z8e{H$);QaTG{$s;HqiL#)N@?gp-w%II$SQgj}`yrtj2vGzh1^??Vf&}>5kf|oBhG7 z&2y!suY6ySkMp>!PL|rJvPAQ5{G1B^Nwev6yw=eAM;{I;OO)5AuL`=-Z_R5^`}75^ z(s`}meS##f(aBe5*l9tbKWVm>>kCNZ1m+@7d6A-Q>tx9vBn#Xs&qv%jlP`Fy{6Km2C5&HUZ(lxG?> zaEeQsx#5tFKL3?NyK^aI$dr(xmVKS?HkEzPkF+LFwBgegwcW>Bj>dlAFutfiabUniau(2-16iTiDtJnJLg4{oA#0htEfwuPkh|2w2hY7I!n#vFLg>(exii( zd96?@KabMVY&soZ1VY_|l^y9-zJILK?QW}Ar{j;p$bZsoI$hg#P;0}HyF1eaz05*K z1Qu>a&f7)T5`)}Vs%_|)zm!f6oJCwsxfy->n&w)ob}X8m$p?TY4>hHcn^N}r#;IOk z$NVL{9qz6oxjE&${`;$WJKUtQPg{RAZ-YjIl26(HOp=c^cztPetRb!yH;=79V_c3j zg@pFy%XT_Cu_>{eaS(A@KN>%NYNO)`{p7GLx=eT5AKkjXHHJgyN_5?5EsriQsOeky zqVGNb5M_NtT*|S85*Lpt9^(%mGZ-zpOuO-9t=azLi3hHrGhI3q!nmwDI?yb}rDOhb zOR}C%2)ol&U$Yn2=0c;pXB-jtLAz|BoAXI@bKY*Ps=q4K;vU=m{yej#uR-6SpHUjA zkCLyl^VnD-)^5c1Z)~%Q+W*aM=I*H4X8ow~E;0Ta##!_6P3!TMeo zu05t-a_yqaw&g^=d-LU1Z;^GA6uzY>{C+9o&qWcBOOgAFBA=If9VzPds?_^rQSZ}I zpD&C0d|EVnXm-y5iCuQtLi*}nq&a-og?2uq-6(>$C=aNX&Bj?~%xoN?I5q;wv&O?gnqj3R9&SM z(-9**MCG|j-yzh{A%qTY^>lFS$!f9@W+FS=~k7w$2KqU$_gy2rl}ZO-`0 zJ@JF+q{&~qr+yJ#c>Zg5()ZC71>d-5{2X0#{u}qKl9>9F7u<8U#9Z+H=$`jrOhV>Q z?rZnPWMBHTd*OkYvaLV67rhWub?GneC3P`%Km6ie_D)RW?qA)hAIF$WesiaN6O)kj zyF2sOn8K-MPtL~Jnkz#+tL}}x5Yo|e<0G*N|LEvhdmy&(rB0rk564#B6z2KE@z}b5 zg?aw?cI<^mI(v$wk7E<8t}dQp>6_R>tD~!DgY;``jn!`TY?e00HCmgjo;#$w&o1eUxTFZ%<(}Qrk8#-%CocCqB6;J>BC4Z2d!)PKYa(0&J$t1G z;x9zl2YH^99*vKRY#8L(Cp{IP6zLf3IUqe7pB>pS*z>gXVtiSoV~FP&>6Q4Z$c7=F z=cKpd>mnUPJugV7;x9z%qCJPC&*Bq$)kb@cO5ep7_HxB|UY34~ujyq^@VqA7HmtFi zWwhrFsbZM9chhLko6^I>qI#bg<9SOliJ6{%Nr#4)U7kG4^P}|g@T$vC%<}v! zy*a$@^5of`U#0hkH(q{Xw&!>0WBMY@JR)ge^%SOVHN;uxIj-Mzb={Dr_ z?_b?G#QBEj4gKBY%tM<_dEV3?9~U*$`Jv}+y=Q#V(57>qzw4hHpFPz1mFHdk_v6ck zHeK+%r+;ih)llazp7-^iOsE^$^tMQPFD6z+I}^NL=yy%3i*6e0{YwA-q{e9H z1n)QcyC<7tnr3*v)gPZ672}-a{a){xk`&Xl$oqr-xhdH(&SlA#;+7SojN{ZW5# zYE?{SuJ>pC$*FZQm3iJ@_5YmO7*kp3{awFhnmM6zy|>74a9UJC`7K_Lp?+FY!iihG z#fIOeWhYoRcuNfrCY2@FH+nZ1jwDqj=x+CJH2jcMmrz;Z-E7!8y)nUZpZ5;KvFYZ~ z)jPa*8a|sIHQH9?z1y&DM$+iYN4;AO&&?OSA22k{yfE7KvbWOkms!>^j#s@88Qz_hFsA$s@1G4@W~Yp)KIMJbaA4KL4W925Js_X)!fbIfC#E_k0bJTy0IZ0*n9 zeTLWPCXIFd<~?BeX>Rt|*wEsq4UZ*Pj7<(Je%A1Aa?My>*W%|5?s@fNEfK|kH9R}- z!dO@D;#$Ku^Q>1lL>0eic;uRdtCNQmA2FQ1CgtktnBrrGyXO~Pts7fhXLxOX)z!7* zi;o+;3+k@6%qV`%P_w`~?!=tpw+s(nn=sD4toWp%WBB~K!g0Fw#r1}_ z7FLXNl@*^dJeyK8PIpJ~X~RQ{>c_cu6n|*gxcI_2$KK)w!+lGv;~SnYK4bW0N%nZ# z@#23N&Mqw*U;SS3Im7GAs>a7Q7Jp&be|_Ef>aUBxGCYvlI6nE;;%^LP4)cVjP9@(O zzH>xPD32)l-ta+M(uDFsB|jLNmS<0}k16?=;beNv1k3c2UktBg)K74xl>BC>&Ac$7 zI<>@Qd@9SD)Rb4^Htt!GkmR_v#A~d~PDwhkxunFnEvGOkc1KB>ar4TGq^7+kw;A`| zP?O~RYsn_#j$HF}`)eh48p~EiO>cU?yW z?@_OLrJIt!bQIKD6{f<3su88P0K~JB?)pDKnC1l|F3z zy`XSL?UK?e<3HC{%&^>0`lzvSUCj)~`qIaZ4L8-#XxLEt7h`?lg&B@}ORJ41Znn;B zs49KR`11OMnbpsf?l+$ML(0t9S4s~W-@B!3W>bCXGsYu-teRQ*QR#EWCvL5q>HNC% z1>>$iHO{R4we)YsZ7#_?t0~OtI%F&@lA>nmdRbkEjlUF0NwZE2wz`fQe{xINvm9fs zu9uA8d8D#g_8C^!%f`37Qq8Q&6szl1?;Cb-iOOy-g~dUHQ1xRd3v|QL32jc;4zdW%O>6 zYG%8RTV1D(jhm##*^X0I*GI;N&60Ue!#S(#V`Kg8Qq&yhk5<q?ai>%^ry(lbbI*9*LmZfd!&@PvA2f1{%PD^Ar;O|-W=|_VBCDK zR57>w{&3e%#-i;~&D;}@g}Z(+p5HFj&viW)?)uI6uluA6bCZvUyIiIw`gYHhQZ!3 z5w6Xq7ao!tQ%=N0xb84LyGt@JicN}e-D!I1&r;N)>XZoA-KOKarKCmn>jl$h z`XunuhFOuW7fqk`XnSR;b7_nLRbqb;{&9DAg=GksaweZMyHERKLu= zGSc;->EJ=>!m{MtNLPdD?*}F8^}4)B*BR3{2c?AT%U4IbJ~fp-Eu~z4Vtu6RGt?RiErr|RyL&mH&}Haaa`O;+Zz3_ihX4gXM2v%kG7qH-0jG)ep?^^DpNR4GV<>nJs(_6+lzBHQ zXSm={@CAZH!3zY3f^QKV3cjECrL+Yn3jb>1hr;g_eklCAgdYljjqpR^f0uYqiW3U|SHcg4pTE?D{ZRP(5$F0s;ZGEP zDEteB9}53E;fKP%N%*1gKSumA+Ab9S7lj`R|Hr})h5v%^L*c)SE--NUQ23*XbA6%k z&lY|t{29b;+(_3#xWAyxd3;6?=Xjv-j}d+-{K>)( zh2JUsQ20xP9}54&!ViW2dEtk`f0{VAKNSA&g&zw4C3K;O{ZRM^5$F0s;h!P=Q25h? z9}0hw@I&FR5Pm59HG)II>jj5`Ul1G$Zlw$Q96uC1L2xK|ir`T2LcyWn6@o*-YXpaa z*AwSEA8Nv^R|7PKb!e1f$Q26%>KNS8~iSs%Jh5wB3L*bX`c#!>2 z_%9>Q^@YMeQuv|p&k}wp{5J?c6#gRNhr(Y;oX0;D{-=c>3jaI84~73L;fKN>MmGz& zd?@@^5a;?r;h!k{Q1}-Sw~-$T|ILC!!S59u3jU1XQ1JHzhl2k|ocjyPoX2NZy5Y+4 zK;e%e&VDHTQ-mK1f2Qz5;s1m1L*c(q_@VGWDg03Q-z3iM4~4%`_@VHd=^+O8L*c)i zxQ#dz{t3blh5tI?hr+*3_@VG`7JexFdj*Grza}^o{42qs;Pjxc96uC%sNhiW*@8pC zR|yUU-z+#3d@pfsPbhO9KSzWg3je3V4~73{;fKN>NhgS0J{10V;@n?Q_^%OuDEv9X z4~2h&;85^AfL2xMe9>JmDFAEL@ zKS!L`D=4`65@o(anREa4ATxb6#n;w9}52k;fKQCnRYr{FDU#& ziSzyl3jZYGhr*vO{80FB7JexF_XzEPLypv-yvgcD~!6#fL^hr&Nw_@VIU z3O^M7BH@R^Un%@h_zw`bQ9Mxi|4y8`MhAue3*m>t--#Y9XFnAFzQnnFDEx`S4~2iB z@I&FhN%*1gZx((i{Cfq5g1;s>6#Of}q2QhA;wQ%s1s^In6nwVeQ1DfPL%}x_=kWjq zf1EfCEgh6OkDtTB4~74X@I&GMQTUCX9^AlFBTjMzFTl8_))>3;Ae@;>w{2m z`Y4q=U!csne&lDUAUMx5ie7E3G@S}o5!Osc~1vm9m;)j9{5F83V zg?Kd83kseuI23%R;85^41c!qEDmWB;@D<8-q2P-Ihk}<14h26TI28O7!J*(hcDcS# z@CkxL!B-IKQA~Gyzc;|d?@%F!J*)P5F83#B{&rP?}9_Y^;as} zg@VTj4h3H-I23$?;85^`ff)@%71>a4a#~T#*;85^O z2PyGD!6yk01;0^nDENbdL&0Ao&hrZj{*&NP@T&$Zekk|?!J*(qfQ7e;VYJq=SOb6C4VDtKd-Z#{`Fh*9#5>H^tDs0@^MVJc>9qoem1VK=`5HHwix! ze3$S;!H)?)6#P5khk{=ktHc8ZPbALqL%}nI9|~SB{7~?H!Vd-iNcf@P`Zy&XD0nn+ zjt2@pPxztWe-wTwc%|?|!Cw`ADEK$R4+W2mSK@(!k0;Lk3k6>({7~@C!Vd+1M);xN zr-dI1ZXTw@0|g&Uoa2FlFBE<#_y7btV?mu19xJVW8XO_T!# z|BLWL!QUs&^9u_88*w@X){59c+f`2RgQ1CwE=oum^2MRu!I2}Uhpy2t$=~P7r1>a7bcJVqW_$$H>1^-s~ zq2PVTEAc?VrxE9Ppy2C>b39P+9mF{vDEMo_4+TFj{7~@z6O?|3f=?sP{SIZ${hm#{ zH*qL&s9DEPg^x!5$E|11urMwm*Wx2ocsMr`knh73jR8A?sq8o8RFdUQ0Cn4T_!302nCNP z&ix1lPZfSBxJUS*;MK&rouT0G5a)J=f}18Q?FI!OLY&(T3cghMq2M0jhl1}T&hbFO z-y_cPK*7yZlz5=vF~m6@D0r&yL%~ag9}0eucz-Gf3Vw=sKjKj4-0!~-??W649x+w9 zuL=c^AWzPM+l=uMRQ1Bw+bjqoNf4XDRVO!4rvdJW%j8#IGO@1>Z`X`w_~V`~6YkLx@A+e}_0F>!9F2 z5$Aq~f?qXT@k7C95@$aYypZ^n#G&915a)Fg3jP{#Zf7X?dE(s8Q1Je96h9PvCUN#d z!T&&<cQB&k*N+go1}8EByimA4Z)0Q1FGq4+ZxU=XQgFKT14?I28On;@po=aNRtm9iZT` z#Muu8Unu-g@KWL&4-|YaagGNHep-|d1vgxy!~+E%Mx5h;f-e?+DEMu}xgVk6dx>*D zLYZ?vzDS(=5eom;#JOLf;N9md{Rjo0K%D(h@O0saf|nELc7uW+B+l&y1wTui+W`tr zFRzuyGZg%4;_Qclrwcz6{BGhL4;1_v;v5eYyit@71@C&T5)Tx70&$K93Z5nWQ1ETU zxjmua&l2xR9Lk*g@dWXq#G&y2N}TtPQ1E`&Df>GpcrtPJL&4VyKNNfyajq8>{1xJ{ z#G&9n5$F9Q6ujR;W&Z*NpHH0qQ1CwpKNS3N;v5eY{B7bK4;0*xqLdE>A51)s{7~?v z!Vd-a3O^M50CDa|DENECdlH8-=YIT#IQJtI{$7ieeu07~5|1W76nwStL&3KQKNS2$ z;@oaf@N>ku-Jsy%i^Xw*Q1C?ho!boxzDD?<;M;{C3Vxh8#{&ibjyT5y1@E^+i3bWk zgE+?n1^@3SLB<$2k=I z9^zatDEM>4!-+$|PZGa`IFvchmwyoNO&khtyk02>3f`M|FY-f~^LfN*;@yZt;h!P= zQ24V1hl0CBIZ%{yx9~&Ze}XvI7YbfSybEzC%6Uilq40krI24>-Rv?!HML9i*^LC-| z#|RDuPZH%oQO4K7WLwoc9EWg8wWy6uh^C zp0%WMpx~nfhcf4JHIMkEn4b2~u6tBG?tP?U3+IFDy2 z{0+o89w_`j5a;rt@OMj7%7?<=pLiegL*bt&I27DLoVN>Q&i%EPcwgdB__q=7P8`ad z+p|jeq42*f{80Gc7k()GzX%Qm@4Z}kzbceD$2o|2Uy26`|4iYB!k7c;HDLd9||5JI21faa47g>!J*)b z1c!q2m#}c0P;ignQ1A-Dq2PN3hl0N-I28O{!J*(^3JwLAvXwZY;5`M0f)5rP3T_u1 z3VyBNQ1F$4L&5(fI23$~;85^Af9Z0SOrJX^b@trJQ&JZ%o)tGLH8nfOIW4mw zIj1l)H!&k)O-@eEl+jaHuUVB=Fzea{0Za*V*JNa_VNbhFO`o5-aQd9da~ICf&0V)D zx?oM3Q~n`0Co8icXH{l&enHxrf@tT6oP^ZW)#*3nre-s#s*-A^Ky{+=A#W)Gb;?-1c6n}QkV>iN3RElg zT!E^ko-0tjcpT9`)ouxxF2D zEtfKTRdPAia#6jwH{=NY3>dX+RIvcr0@W-)wm?-2kS$Q%s9yOri2{xp)pRjtRMW+n zQB4Zok;16(%lo{p2@Ko?0HI@BHR*s^wAoQ_j@|N1q9kL|@^@t7s8b z-QP|0h@anshWxEL>1ynsg0??=@eXdm`>Xdm`>ZL z{tmePfa$dTfa$dTfa%JZq7eOzsVwzAp>gOp&w8~y%82rtCr>SpGNAnCajNAJ<4Mlc ze1Mssy_zmR(kxNSwO$z|{&NMY6OS$bxdPQnJy)Q5(RO%Ss`DgZx)`Ub`NcR@O&8-- zHJw_J$C+xn7^kY~Vw|d`E8~ET)pRjVRnx^d zRZY+1^FU=E-Q2BfHSDC-a|PNEmGMfO@Hbx7cW*q3{AUWVcT>+4VCSZuDZsu> z9(Zzme#UEnTJZ={$rh+y0kQ?ES%7ST>PCGTXumD5k5q|(<5e|Xj91lkFE#duXs z7vohmU5r1hS4SEa7W%uOq#Clk^O(o$Ds=4GzQNl(r3cN0~1r={PxHfK$yMmLq6 znU=5BO=WX5+RK*XUg2DO5pipCBW`VO#I4PZ_RhaK|03en=0@Dw+=yG78*$6yBzN`l zv|O#jB$yddYc08DkZ35mWsC$fV`$6W7|e{AHJQ=Sn#^ctxmSXPftWR!5wj*UV%B6v z%-n=GuoLuC1@CkS(6zt zYceBdO=iR_&u+1Q$?*4-lj{CO-HoUM2 z?aiH|J#KAo#I4PZxV5r-$50B5#vsSx5c8m#pdWVcoj>)vDBuxfZmqkv_{9psd|;-I5~LFYVfvL zRJYisvg}j!isCSqwiKr|x`4w`DUOQ@J{&c8TP#|mW1p&56o;|o zXE+WEHXPd&;cqy$SX2XP(?fhmNY$IBthRjRGQ;7G$k#TZq5W>Elrird{L#$EULV zY8;;$%xNti4H0TEivj}N$-d}=U@2)RJb5e77)zyRYj zKKS?)k)sg5s{7*+O+tB2@9TDleFs3#7i+USKdY z6}T?Gr4=>AsKG7@4OSW7nscea{G!-kRcc#oFh3VN{NkD*xLp()tV*q7rBZ|WMX|xE z)VA1Qewv%wCQ^K$iPR=ho=D7s4JzNQt9aQGtZkYrt;NiqV6D=sz+h%7P{YjdEmDJB z6nZf=*D{ervB9d;HrQZ(E;jhgXjN)(yD0Rcnn1%u62%6qQrlpI`RN=%+e8``Xd<;q zlqV81mC4Z@r0>>Mylm0jmRij03DzpD0t{xR0yWGG-y${GMWGi{b1f4|6dSBcZG#Qw z=VF7;j8>%vw~InAstGhqBvEXzDzyzZn4eE)w9O>>^eMo>9i3eHq>7_EKWAKR5ARw< zpfa0ekLDKDeiVCg4WVUraYK+V zSO>QYI6#K{0`FGY8%2 zw2P4k87}9jS3AfU_gkul$_E+Yev8%81OW$l`%;73xn0Q>tV+IprNu9b4OXSL#Rl{9 z3W`4U*^MZ3Azs5=@GU@PP_RXFb7(QMCs@Ur8$yejJQ@bKZ;=}8qR@+}xt6&jioKZn zYneM-Z1BO|stJPIMWGkf1RCa!DE6Y7fQk*~r&A^MxiCz_T<|SGWiGHqb8~1hvnN=^ zntM%)nLHW>w{MXe?4r<%skxT9Ad0=1`fHgBTx{^c-Kq(K+eM)l)dU*mf++T)nn1%` z;8Pp*$>8^vX#U#PyUSKZs7wdhqq#{mxn+~)=FsG3lZHXwJRmgsVY{gMYnc`(_@Y`t z%fyfi4nEXdH$-qhioLjo&@eer@WnNRmgzyeAoc0t_l|74>ET<1+Vo(L_Eyp2=F>Ip zI}A;3`3x}lxk1aQpb#`nClq^8_182##P+q13XK`$h=ULJ)(sKdk76&bA+$^nu9t?< z-Y#JEmRpm#A~iQHFEb}YevfwAs*DxsdC{wur{@(!=dSj1YMq^xo4VF156S%8)i?R( zx+!g4X7uv3HEVL{DS-ATU@e*|9 z%Dwo-_f_Zu`F|67>qP%ALSMBum)`d-zn4bYhyXbw{d|ad6+ICGbK8|vS4wIAt!)u}%Ua^49a@!u=-E z2WeUrzhcd5TA|cSYU|HiyDC-A~^0w2dDmREw({E1Cr4M$jkzeWHKUCT7uS>Bu z@zS z_q7E1*;%RdDfwJ_8l4BV9DG)Oc5bxaIl=BV?cFqN{cSwMIXq3Csr+c9Qopqss|)c#3}O?S z^P!n8I>SY7$fp&wZptT%rj5}uMth#Tw)-c_`ktt(hlK(7&{RaGpLfXcg9+d$=o zscoS07F2o3Y$;QDb!r=^yzjIPRI-bs44g22y$uxETZg1HojCYInId~@rxJV{rxJV{s07~v`TKx; zdc0uGTG~yt^ecTU3Ju4W-_tYc7`^3JC3kZy!rlD4;#Pj8Zy!<~pyu27*>tCJx!Aq6 z{%BX+WZyBWoO)&F__7Q-KyYTRS+$mSOv+r#Sd}VH66C!g{hE;}1K(_X$jIk*VjyI# zqHjKG8Gu=<(hBiwt4LJ>mG895Tk{QH{$N^rI*hKqx2J6%EmA(*MnzEQ^leV!7KQ>h;SpWb4 diff --git a/ptx/lib/zluda_ptx_impl.cpp b/ptx/lib/zluda_ptx_impl.cpp index 420ce651..ecbe691c 100644 --- a/ptx/lib/zluda_ptx_impl.cpp +++ b/ptx/lib/zluda_ptx_impl.cpp @@ -11,6 +11,7 @@ // https://llvm.org/docs/AMDGPUUsage.html #include +#include #include #define HIP_NO_HALF #include @@ -155,6 +156,399 @@ static __device__ float4::Native_vec_ __pack_to_float4(const T &t) return result; } +typedef uint32_t uint8 __attribute__((ext_vector_type(8))); +typedef uint32_t zluda_uint3 __attribute__((ext_vector_type(3))); +typedef uint8 CONSTANT_SPACE *surface_ptr; + +template +static __device__ To transmute(From f) +{ + if constexpr (sizeof(To) == sizeof(From)) + { + return std::bit_cast(f); + } + else if constexpr (sizeof(To) > sizeof(From)) + { + union + { + To t; + From f; + } u = {To{0}}; + u.f = f; + return u.t; + } + else if constexpr (sizeof(To) < sizeof(From)) + { + union + { + From f; + To t; + } u = {From{f}}; + return u.t; + } + else + { + static_assert(sizeof(To) == 0); + } +} + +enum class ImageGeometry +{ + _1D, + _2D, + _3D, + A1D, + A2D +}; + +// clang-format off +template struct Coordinates; +template <> struct Coordinates { using type = uint1::Native_vec_; }; +template <> struct Coordinates { using type = uint2::Native_vec_; }; +template <> struct Coordinates { using type = uint4::Native_vec_; }; +template <> struct Coordinates +{ + using type = uint2::Native_vec_; using arg_type = uint1::Native_vec_; + static __device__ type pack_layer(uint32_t layer, arg_type coord) + { + return type { coord.x, layer }; + } +}; +template <> struct Coordinates +{ + using type = zluda_uint3; using arg_type = uint2::Native_vec_; + static __device__ type pack_layer(uint32_t layer, arg_type coord) + { + return type { coord.x, coord.y, layer }; + } +}; +// clang-format on + +template +static __device__ void image_store_pck(T value, typename Coordinates::type coord, surface_ptr surface) +{ + if constexpr (sizeof(T) <= sizeof(uint)) + { + uint value_dword = transmute(value); + if constexpr (geo == ImageGeometry::_1D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0x1 dim:1D unorm" : : "v"(value_dword), "v"(coord.x), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_2D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0x1 dim:2D unorm" : : "v"(value_dword), "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_3D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0x1 dim:3D unorm" : : "v"(value_dword), "v"(transmute(coord)), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A1D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0x1 dim:1D_ARRAY unorm" : : "v"(value_dword), "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A2D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0x1 dim:2D_ARRAY unorm" : : "v"(value_dword), "v"(coord), "s"(*surface) : "memory"); + } + else + { + static_assert(sizeof(T) == 0, "Invalid geometry"); + } + } + else if constexpr (sizeof(T) == sizeof(uint2::Native_vec_)) + { + uint2::Native_vec_ value_dword2 = transmute(value); + if constexpr (geo == ImageGeometry::_1D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0x3 dim:1D unorm" : : "v"(value_dword2), "v"(coord.x), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_2D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0x3 dim:2D unorm" : : "v"(value_dword2), "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_3D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0x3 dim:3D unorm" : : "v"(value_dword2), "v"(transmute(coord)), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A1D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0x3 dim:1D_ARRAY unorm" : : "v"(value_dword2), "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A2D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0x3 dim:2D_ARRAY unorm" : : "v"(value_dword2), "v"(coord), "s"(*surface) : "memory"); + } + else + { + static_assert(sizeof(T) == 0, "Invalid geometry"); + } + } + else if constexpr (sizeof(T) == sizeof(uint4::Native_vec_)) + { + uint4::Native_vec_ value_dword4 = transmute(value); + if constexpr (geo == ImageGeometry::_1D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0xf dim:1D unorm" : : "v"(value_dword4), "v"(coord.x), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_2D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0xf dim:2D unorm" : : "v"(value_dword4), "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_3D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0xf dim:3D unorm" : : "v"(value_dword4), "v"(transmute(coord)), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A1D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0xf dim:1D_ARRAY unorm" : : "v"(value_dword4), "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A2D) + { + asm volatile("image_store_pck %0, %1, %2 dmask:0xf dim:2D_ARRAY unorm" : : "v"(value_dword4), "v"(coord), "s"(*surface) : "memory"); + } + else + { + static_assert(sizeof(T) == 0, "Invalid geometry"); + } + } + else + { + static_assert(sizeof(T) == 0, "Invalid vector size"); + } +} + +template +static __device__ T image_load_pck(typename Coordinates::type coord, surface_ptr surface) +{ + if constexpr (sizeof(T) <= sizeof(uint)) + { + uint data; + if constexpr (geo == ImageGeometry::_1D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0x1 dim:1D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord.x), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_2D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0x1 dim:2D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_3D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0x1 dim:3D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(transmute(coord)), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A1D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0x1 dim:1D_ARRAY unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A2D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0x1 dim:2D_ARRAY unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else + { + static_assert(sizeof(ImageGeometry) == 0, "Invalid geometry"); + } + return transmute(data); + } + else if constexpr (sizeof(T) == sizeof(uint2::Native_vec_)) + { + uint2::Native_vec_ data; + if constexpr (geo == ImageGeometry::_1D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0x3 dim:1D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord.x), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_2D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0x3 dim:2D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_3D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0x3 dim:3D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(transmute(coord)), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A1D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0x3 dim:1D_ARRAY unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A2D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0x3 dim:2D_ARRAY unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else + { + static_assert(sizeof(ImageGeometry) == 0, "Invalid geometry"); + } + return transmute(data); + } + else if constexpr (sizeof(T) == sizeof(uint4::Native_vec_)) + { + uint4::Native_vec_ data; + if constexpr (geo == ImageGeometry::_1D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0xf dim:1D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord.x), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_2D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0xf dim:2D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_3D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0xf dim:3D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(transmute(coord)), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A1D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0xf dim:1D_ARRAY unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A2D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0xf dim:2D_ARRAY unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else + { + static_assert(sizeof(ImageGeometry) == 0, "Invalid geometry"); + } + return transmute(data); + } + else + { + static_assert(sizeof(T) == 0, "Invalid vector size"); + } +} + +template +static __device__ uint4::Native_vec_ image_load_pck_full(typename Coordinates::type coord, surface_ptr surface) +{ + uint4::Native_vec_ data; + if constexpr (geo == ImageGeometry::_1D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0xf dim:1D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord.x), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_2D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0xf dim:2D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::_3D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0xf dim:3D unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(transmute(coord)), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A1D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0xf dim:1D_ARRAY unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else if constexpr (geo == ImageGeometry::A2D) + { + asm volatile("image_load_pck %0, %1, %2 dmask:0xf dim:2D_ARRAY unorm\ns_waitcnt vmcnt(0)" : "=v"(data) : "v"(coord), "s"(*surface) : "memory"); + } + else + { + static_assert(sizeof(ImageGeometry) == 0, "Invalid geometry"); + } + return data; +} + +template +static __device__ void image_store_pck_full_with(uint4::Native_vec_ data, T value, typename Coordinates::type coord, surface_ptr surface) +{ + // We avoid unions for types smaller than sizeof(uint32_t), + // because in those cases we get this garbage: + // ds_write_b128 v2, v[5:8] + // ds_write_b16 v2, v9 + // ds_read_b128 v[5:8], v2 + // tested with ROCm 5.7.1 on gfx1030 + if constexpr (sizeof(T) == sizeof(uint8_t)) + { + uint32_t x = uint32_t(std::bit_cast(value)); + uint32_t data_0 = ((data[0]) >> 8) << 8; + data[0] = data_0 | x; + } + else if constexpr (sizeof(T) == sizeof(uint16_t)) + { + uint32_t x = uint32_t(std::bit_cast(value)); + uint32_t data_0 = ((data[0]) >> 16) << 16; + data[0] = data_0 | x; + } + else + { + union + { + uint4::Native_vec_ full_vec; + T value; + } u = {0}; + u.full_vec = data; + u.value = value; + data = u.full_vec; + } + image_store_pck(data, coord, surface); +} + +constexpr auto IMAGE_RESERVED_TOP_BITS = 3; + +static __device__ surface_ptr get_surface_pointer(uint64_t s) +{ + return (surface_ptr)((s << IMAGE_RESERVED_TOP_BITS) >> IMAGE_RESERVED_TOP_BITS); +} + +static __device__ surface_ptr get_surface_pointer(struct textureReference GLOBAL_SPACE *surf_ref) +{ + return (surface_ptr)(surf_ref->textureObject); +} + +static __device__ uint32_t x_coordinate_shift(uint64_t s) +{ + return uint32_t(s >> (64 - IMAGE_RESERVED_TOP_BITS)); +} + +static __device__ uint32_t x_coordinate_shift(struct textureReference GLOBAL_SPACE *ptr) +{ + uint32_t channels = uint32_t(ptr->numChannels); + uint32_t format_width = 0; + hipArray_Format format = ptr->format; + switch (format) + { + case hipArray_Format::HIP_AD_FORMAT_UNSIGNED_INT8: + case hipArray_Format::HIP_AD_FORMAT_SIGNED_INT8: + format_width = 1; + break; + case hipArray_Format::HIP_AD_FORMAT_UNSIGNED_INT16: + case hipArray_Format::HIP_AD_FORMAT_SIGNED_INT16: + case hipArray_Format::HIP_AD_FORMAT_HALF: + format_width = 2; + break; + case hipArray_Format::HIP_AD_FORMAT_UNSIGNED_INT32: + case hipArray_Format::HIP_AD_FORMAT_SIGNED_INT32: + case hipArray_Format::HIP_AD_FORMAT_FLOAT: + format_width = 4; + break; + default: + __builtin_unreachable(); + } + return uint32_t(__builtin_ctz(format_width * channels)); +} + +template +static __device__ T suld_b_zero(Surface surf_arg, typename Coordinates::type coord) +{ + surface_ptr surface = get_surface_pointer(surf_arg); + uint32_t shift_x = x_coordinate_shift(surf_arg); + coord.x = coord.x >> shift_x; + return image_load_pck(coord, surface); +} + +template +static __device__ void sust_b_zero(Surface surf_arg, typename Coordinates::type coord, T data) +{ + surface_ptr surface = get_surface_pointer(surf_arg); + uint32_t shift_x = x_coordinate_shift(surf_arg); + coord.x = coord.x >> shift_x; + if (shift_x <= __builtin_ctz(sizeof(T))) [[likely]] + { + image_store_pck(data, coord, surface); + } + else + { + uint4::Native_vec_ pixel = image_load_pck_full(coord, surface); + image_store_pck_full_with(pixel, data, coord, surface); + } +} + extern "C" { #define atomic_inc(NAME, SUCCESS, FAILURE, SCOPE, SPACE) \ @@ -620,179 +1014,101 @@ extern "C" suld_b_a2d_vec(_v4, b32, uint4); // suld_b_a2d_vec(_v4, b64, ulong4); -#define sust_b_1d_vec(VEC, TYPE, HIP_TYPE) \ - void FUNC(sust_b_1d##VEC##_##TYPE##_trap)(struct textureReference GLOBAL_SPACE * ptr, int1::Native_vec_ coord, HIP_TYPE::Native_vec_ data) \ - { \ - hipTextureObject_t textureObject = ptr->textureObject; \ - TEXTURE_OBJECT_PARAMETERS_INIT; \ - (void)s; \ - int byte_coord = __hipGetPixelAddr(coord.x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i)); \ - HIP_TYPE hip_data; \ - hip_data.data = data; \ - auto tmp = __pack_to_float4(hip_data); \ - __ockl_image_store_1D(i, byte_coord, tmp); \ - } \ - void FUNC(sust_b_indirect_1d##VEC##_##TYPE##_trap)(uint64_t serf_arg, int1::Native_vec_ coord, HIP_TYPE::Native_vec_ data) \ - { \ - hipSurfaceObject_t surfObj = (hipSurfaceObject_t)serf_arg; \ - HIP_TYPE hip_data; \ - hip_data.data = data; \ - surf1Dwrite(hip_data, surfObj, coord.x); \ - } - - sust_b_1d_vec(, b8, uchar1); - sust_b_1d_vec(, b16, ushort1); - sust_b_1d_vec(, b32, uint1); - // sust_b_1d_vec(, b64, ulong1); - sust_b_1d_vec(_v2, b8, uchar2); - sust_b_1d_vec(_v2, b16, ushort2); - sust_b_1d_vec(_v2, b32, uint2); - // sust_b_1d_vec(_v2, b64, ulong2); - sust_b_1d_vec(_v4, b8, uchar4); - sust_b_1d_vec(_v4, b16, ushort4); - sust_b_1d_vec(_v4, b32, uint4); - // sust_b_1d_vec(_v4, b64, ulong4); - -#define sust_b_2d_vec(VEC, TYPE, HIP_TYPE) \ - void FUNC(sust_b_2d##VEC##_##TYPE##_trap)(struct textureReference GLOBAL_SPACE * ptr, int2::Native_vec_ coord, HIP_TYPE::Native_vec_ data) \ - { \ - hipTextureObject_t textureObject = ptr->textureObject; \ - TEXTURE_OBJECT_PARAMETERS_INIT; \ - (void)s; \ - int byte_coord = __hipGetPixelAddr(coord.x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i)); \ - HIP_TYPE hip_data; \ - hip_data.data = data; \ - auto tmp = __pack_to_float4(hip_data); \ - __ockl_image_store_2D(i, int2(byte_coord, coord.y).data, tmp); \ - } \ - void FUNC(sust_b_indirect_2d##VEC##_##TYPE##_trap)(uint64_t serf_arg, int2::Native_vec_ coord, HIP_TYPE::Native_vec_ data) \ - { \ - hipSurfaceObject_t surfObj = (hipSurfaceObject_t)serf_arg; \ - HIP_TYPE hip_data; \ - hip_data.data = data; \ - surf2Dwrite(hip_data, surfObj, coord.x, coord.y); \ - } - - sust_b_2d_vec(, b8, uchar1); - sust_b_2d_vec(, b16, ushort1); - sust_b_2d_vec(, b32, uint1); - // sust_b_2d_vec(, b64, ulong1); - sust_b_2d_vec(_v2, b8, uchar2); - sust_b_2d_vec(_v2, b16, ushort2); - sust_b_2d_vec(_v2, b32, uint2); - // sust_b_2d_vec(_v2, b64, ulong2); - sust_b_2d_vec(_v4, b8, uchar4); - sust_b_2d_vec(_v4, b16, ushort4); - sust_b_2d_vec(_v4, b32, uint4); - // sust_b_2d_vec(_v4, b64, ulong4); - -#define sust_b_3d_vec(VEC, TYPE, HIP_TYPE) \ - void FUNC(sust_b_3d##VEC##_##TYPE##_trap)(struct textureReference GLOBAL_SPACE * ptr, int4::Native_vec_ coord, HIP_TYPE::Native_vec_ data) \ - { \ - hipTextureObject_t textureObject = ptr->textureObject; \ - TEXTURE_OBJECT_PARAMETERS_INIT; \ - (void)s; \ - int byte_coord = __hipGetPixelAddr(coord.x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i)); \ - HIP_TYPE hip_data; \ - hip_data.data = data; \ - auto tmp = __pack_to_float4(hip_data); \ - __ockl_image_store_3D(i, int4(byte_coord, coord.y, coord.z, 0).data, tmp); \ - } \ - void FUNC(sust_b_indirect_3d##VEC##_##TYPE##_trap)(uint64_t serf_arg, int4::Native_vec_ coord, HIP_TYPE::Native_vec_ data) \ - { \ - hipSurfaceObject_t surfObj = (hipSurfaceObject_t)serf_arg; \ - __HIP_SURFACE_OBJECT_PARAMETERS_INIT; \ - int byte_coord = __hipGetPixelAddr(coord.x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i)); \ - HIP_TYPE hip_data; \ - hip_data.data = data; \ - auto tmp = __pack_to_float4(hip_data); \ - __ockl_image_store_3D(i, int4(byte_coord, coord.y, coord.z, 0).data, tmp); \ - } - - sust_b_3d_vec(, b8, uchar1); - sust_b_3d_vec(, b16, ushort1); - sust_b_3d_vec(, b32, uint1); - // sust_b_3d_vec(, b64, ulong1); - sust_b_3d_vec(_v2, b8, uchar2); - sust_b_3d_vec(_v2, b16, ushort2); - sust_b_3d_vec(_v2, b32, uint2); - // sust_b_3d_vec(_v2, b64, ulong2); - sust_b_3d_vec(_v4, b8, uchar4); - sust_b_3d_vec(_v4, b16, ushort4); - sust_b_3d_vec(_v4, b32, uint4); - // sust_b_3d_vec(_v4, b64, ulong4); - -#define sust_b_a1d_vec(VEC, TYPE, HIP_TYPE) \ - void FUNC(sust_b_a1d##VEC##_##TYPE##_trap)(struct textureReference GLOBAL_SPACE * ptr, uint layer, int x, HIP_TYPE::Native_vec_ data) \ - { \ - hipTextureObject_t textureObject = ptr->textureObject; \ - TEXTURE_OBJECT_PARAMETERS_INIT; \ - (void)s; \ - int byte_coord = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1Da(i), __ockl_image_channel_order_1Da(i)); \ - HIP_TYPE hip_data; \ - hip_data.data = data; \ - auto tmp = __pack_to_float4(hip_data); \ - __ockl_image_store_1Da(i, int2(byte_coord, int(layer)).data, tmp); \ - } \ - void FUNC(sust_b_indirect_a1d##VEC##_##TYPE##_trap)(uint64_t serf_arg, uint layer, int x, HIP_TYPE::Native_vec_ data) \ - { \ - hipSurfaceObject_t surfObj = (hipSurfaceObject_t)serf_arg; \ - __HIP_SURFACE_OBJECT_PARAMETERS_INIT; \ - int byte_coord = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1Da(i), __ockl_image_channel_order_1Da(i)); \ - HIP_TYPE hip_data; \ - hip_data.data = data; \ - auto tmp = __pack_to_float4(hip_data); \ - __ockl_image_store_1Da(i, int2(byte_coord, int(layer)).data, tmp); \ - } - - sust_b_a1d_vec(, b8, uchar1); - sust_b_a1d_vec(, b16, ushort1); - sust_b_a1d_vec(, b32, uint1); - // sust_b_a1d_vec(, b64, ulong1); - sust_b_a1d_vec(_v2, b8, uchar2); - sust_b_a1d_vec(_v2, b16, ushort2); - sust_b_a1d_vec(_v2, b32, uint2); - // sust_b_a1d_vec(_v2, b64, ulong2); - sust_b_a1d_vec(_v4, b8, uchar4); - sust_b_a1d_vec(_v4, b16, ushort4); - sust_b_a1d_vec(_v4, b32, uint4); - // sust_b_a1d_vec(_v4, b64, ulong4); - -#define sust_b_a2d_vec(VEC, TYPE, HIP_TYPE) \ - void FUNC(sust_b_a2d##VEC##_##TYPE##_trap)(struct textureReference GLOBAL_SPACE * ptr, uint layer, int x, int y, HIP_TYPE::Native_vec_ data) \ - { \ - hipTextureObject_t textureObject = ptr->textureObject; \ - TEXTURE_OBJECT_PARAMETERS_INIT; \ - (void)s; \ - int byte_coord = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2Da(i), __ockl_image_channel_order_2Da(i)); \ - HIP_TYPE hip_data; \ - hip_data.data = data; \ - auto tmp = __pack_to_float4(hip_data); \ - __ockl_image_store_2Da(i, int4(byte_coord, y, int(layer), 0).data, tmp); \ - } \ - void FUNC(sust_b_indirect_a2d##VEC##_##TYPE##_trap)(uint64_t serf_arg, uint layer, int x, int y, HIP_TYPE::Native_vec_ data) \ - { \ - hipSurfaceObject_t surfObj = (hipSurfaceObject_t)serf_arg; \ - __HIP_SURFACE_OBJECT_PARAMETERS_INIT; \ - int byte_coord = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2Da(i), __ockl_image_channel_order_2Da(i)); \ - HIP_TYPE hip_data; \ - hip_data.data = data; \ - auto tmp = __pack_to_float4(hip_data); \ - __ockl_image_store_2Da(i, int4(byte_coord, y, int(layer), 0).data, tmp); \ - } - - sust_b_a2d_vec(, b8, uchar1); - sust_b_a2d_vec(, b16, ushort1); - sust_b_a2d_vec(, b32, uint1); - // sust_b_a2d_vec(, b64, ulong1); - sust_b_a2d_vec(_v2, b8, uchar2); - sust_b_a2d_vec(_v2, b16, ushort2); - sust_b_a2d_vec(_v2, b32, uint2); - // sust_b_a2d_vec(_v2, b64, ulong2); - sust_b_a2d_vec(_v4, b8, uchar4); - sust_b_a2d_vec(_v4, b16, ushort4); - sust_b_a2d_vec(_v4, b32, uint4); - // sust_b_a2d_vec(_v4, b64, ulong4); +#define SUST_B_ZERO(TYPE, GEOMETRY, HIP_TYPE) \ + HIP_TYPE::Native_vec_ FUNC(suld_b_indirect_##TYPE##_zero)(uint64_t surf_arg, typename Coordinates::type coord) \ + { \ + return suld_b_zero(surf_arg, coord); \ + } \ + void FUNC(sust_b_indirect_##TYPE##_zero)(uint64_t surf_arg, typename Coordinates::type coord, HIP_TYPE::Native_vec_ data) \ + { \ + sust_b_zero(surf_arg, coord, data); \ + } \ + HIP_TYPE::Native_vec_ FUNC(suld_b_##TYPE##_zero)(struct textureReference GLOBAL_SPACE * ptr, typename Coordinates::type coord) \ + { \ + return suld_b_zero(ptr, coord); \ + } \ + void FUNC(sust_b_##TYPE##_zero)(struct textureReference GLOBAL_SPACE * ptr, typename Coordinates::type coord, HIP_TYPE::Native_vec_ data) \ + { \ + sust_b_zero(ptr, coord, data); \ + } + +#define SUST_B_ZERO_ARRAY(TYPE, GEOMETRY, HIP_TYPE) \ + HIP_TYPE::Native_vec_ FUNC(suld_b_indirect_##TYPE##_zero)(uint64_t surf_arg, uint32_t layer, typename Coordinates::arg_type coord) \ + { \ + auto coord_array = Coordinates::pack_layer(layer, coord); \ + return suld_b_zero(surf_arg, coord_array); \ + } \ + void FUNC(sust_b_indirect_##TYPE##_zero)(uint64_t surf_arg, uint32_t layer, typename Coordinates::arg_type coord, HIP_TYPE::Native_vec_ data) \ + { \ + auto coord_array = Coordinates::pack_layer(layer, coord); \ + sust_b_zero(surf_arg, coord_array, data); \ + } \ + HIP_TYPE::Native_vec_ FUNC(suld_b_##TYPE##_zero)(struct textureReference GLOBAL_SPACE * ptr, uint32_t layer, typename Coordinates::arg_type coord) \ + { \ + auto coord_array = Coordinates::pack_layer(layer, coord); \ + return suld_b_zero(ptr, coord_array); \ + } \ + void FUNC(sust_b_##TYPE##_zero)(struct textureReference GLOBAL_SPACE * ptr, uint32_t layer, typename Coordinates::arg_type coord, HIP_TYPE::Native_vec_ data) \ + { \ + auto coord_array = Coordinates::pack_layer(layer, coord); \ + sust_b_zero(ptr, coord_array, data); \ + } + + SUST_B_ZERO(1d_b8, ImageGeometry::_1D, uchar1); + SUST_B_ZERO(1d_b16, ImageGeometry::_1D, ushort1); + SUST_B_ZERO(1d_b32, ImageGeometry::_1D, uint1); + SUST_B_ZERO(1d_b64, ImageGeometry::_1D, ulong1); + SUST_B_ZERO(1d_v2_b8, ImageGeometry::_1D, uchar2); + SUST_B_ZERO(1d_v2_b16, ImageGeometry::_1D, ushort2); + SUST_B_ZERO(1d_v2_b32, ImageGeometry::_1D, uint2); + SUST_B_ZERO(1d_v2_b64, ImageGeometry::_1D, ulong2); + SUST_B_ZERO(1d_v4_b8, ImageGeometry::_1D, uchar4); + SUST_B_ZERO(1d_v4_b16, ImageGeometry::_1D, ushort4); + SUST_B_ZERO(1d_v4_b32, ImageGeometry::_1D, uint4); + SUST_B_ZERO(2d_b8, ImageGeometry::_2D, uchar1); + SUST_B_ZERO(2d_b16, ImageGeometry::_2D, ushort1); + SUST_B_ZERO(2d_b32, ImageGeometry::_2D, uint1); + SUST_B_ZERO(2d_b64, ImageGeometry::_2D, ulong1); + SUST_B_ZERO(2d_v2_b8, ImageGeometry::_2D, uchar2); + SUST_B_ZERO(2d_v2_b16, ImageGeometry::_2D, ushort2); + SUST_B_ZERO(2d_v2_b32, ImageGeometry::_2D, uint2); + SUST_B_ZERO(2d_v2_b64, ImageGeometry::_2D, ulong2); + SUST_B_ZERO(2d_v4_b8, ImageGeometry::_2D, uchar4); + SUST_B_ZERO(2d_v4_b16, ImageGeometry::_2D, ushort4); + SUST_B_ZERO(2d_v4_b32, ImageGeometry::_2D, uint4); + SUST_B_ZERO(3d_b8, ImageGeometry::_3D, uchar1); + SUST_B_ZERO(3d_b16, ImageGeometry::_3D, ushort1); + SUST_B_ZERO(3d_b32, ImageGeometry::_3D, uint1); + SUST_B_ZERO(3d_b64, ImageGeometry::_3D, ulong1); + SUST_B_ZERO(3d_v2_b8, ImageGeometry::_3D, uchar2); + SUST_B_ZERO(3d_v2_b16, ImageGeometry::_3D, ushort2); + SUST_B_ZERO(3d_v2_b32, ImageGeometry::_3D, uint2); + SUST_B_ZERO(3d_v2_b64, ImageGeometry::_3D, ulong2); + SUST_B_ZERO(3d_v4_b8, ImageGeometry::_3D, uchar4); + SUST_B_ZERO(3d_v4_b16, ImageGeometry::_3D, ushort4); + SUST_B_ZERO(3d_v4_b32, ImageGeometry::_3D, uint4); + SUST_B_ZERO_ARRAY(a1d_b8, ImageGeometry::A1D, uchar1); + SUST_B_ZERO_ARRAY(a1d_b16, ImageGeometry::A1D, ushort1); + SUST_B_ZERO_ARRAY(a1d_b32, ImageGeometry::A1D, uint1); + SUST_B_ZERO_ARRAY(a1d_b64, ImageGeometry::A1D, ulong1); + SUST_B_ZERO_ARRAY(a1d_v2_b8, ImageGeometry::A1D, uchar2); + SUST_B_ZERO_ARRAY(a1d_v2_b16, ImageGeometry::A1D, ushort2); + SUST_B_ZERO_ARRAY(a1d_v2_b32, ImageGeometry::A1D, uint2); + SUST_B_ZERO_ARRAY(a1d_v2_b64, ImageGeometry::A1D, ulong2); + SUST_B_ZERO_ARRAY(a1d_v4_b8, ImageGeometry::A1D, uchar4); + SUST_B_ZERO_ARRAY(a1d_v4_b16, ImageGeometry::A1D, ushort4); + SUST_B_ZERO_ARRAY(a1d_v4_b32, ImageGeometry::A1D, uint4); + SUST_B_ZERO_ARRAY(a2d_b8, ImageGeometry::A2D, uchar1); + SUST_B_ZERO_ARRAY(a2d_b16, ImageGeometry::A2D, ushort1); + SUST_B_ZERO_ARRAY(a2d_b32, ImageGeometry::A2D, uint1); + SUST_B_ZERO_ARRAY(a2d_b64, ImageGeometry::A2D, ulong1); + SUST_B_ZERO_ARRAY(a2d_v2_b8, ImageGeometry::A2D, uchar2); + SUST_B_ZERO_ARRAY(a2d_v2_b16, ImageGeometry::A2D, ushort2); + SUST_B_ZERO_ARRAY(a2d_v2_b32, ImageGeometry::A2D, uint2); + SUST_B_ZERO_ARRAY(a2d_v2_b64, ImageGeometry::A2D, ulong2); + SUST_B_ZERO_ARRAY(a2d_v4_b8, ImageGeometry::A2D, uchar4); + SUST_B_ZERO_ARRAY(a2d_v4_b16, ImageGeometry::A2D, ushort4); + SUST_B_ZERO_ARRAY(a2d_v4_b32, ImageGeometry::A2D, uint4); __device__ static inline bool is_upper_warp() { diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs index 61a74c95..10852583 100644 --- a/ptx/src/translate.rs +++ b/ptx/src/translate.rs @@ -2934,7 +2934,7 @@ fn replace_instructions_with_builtins_impl<'input>( vector, "_", suld.type_.to_ptx_name(), - "_trap", + "_zero", ] .concat(); statements.push(instruction_to_fn_call( @@ -2955,7 +2955,7 @@ fn replace_instructions_with_builtins_impl<'input>( vector, "_", sust.type_.to_ptx_name(), - "_trap", + "_zero", ] .concat(); statements.push(instruction_to_fn_call( diff --git a/zluda/src/cuda.rs b/zluda/src/cuda.rs index 1d054c38..1f37dbfe 100644 --- a/zluda/src/cuda.rs +++ b/zluda/src/cuda.rs @@ -1245,7 +1245,7 @@ mod definitions { pub(crate) unsafe fn cuSurfObjectDestroy( surfObject: hipSurfaceObject_t, ) -> hipError_t { - hipDestroySurfaceObject(surfObject) + surface::destroy(surfObject) } pub(crate) unsafe fn cuTexObjectCreate( diff --git a/zluda/src/impl/surface.rs b/zluda/src/impl/surface.rs index fcf9a52c..0f24fa38 100644 --- a/zluda/src/impl/surface.rs +++ b/zluda/src/impl/surface.rs @@ -1,23 +1,65 @@ +use super::hipfix; +use crate::hip_call_cuda; use cuda_types::*; use hip_runtime_sys::*; use std::{mem, ptr}; -use crate::hip_call_cuda; - -use super::{hipfix, FromCuda}; +// Same as in zluda_ptx_impl.cpp +const IMAGE_RESERVED_TOP_BITS: u32 = 3; pub(crate) unsafe fn create( - p_surf_object: *mut hipSurfaceObject_t, + result: *mut hipSurfaceObject_t, p_res_desc: *const CUDA_RESOURCE_DESC, ) -> Result<(), CUresult> { if p_res_desc == ptr::null() { return Err(CUresult::CUDA_ERROR_INVALID_VALUE); } let desc = to_surface_desc(*p_res_desc)?; - hip_call_cuda!(hipCreateSurfaceObject(p_surf_object, &desc)); + // We need to check array format and channel count to set top bits of the surface object. + // HIP does not support non-Array sources anyway + if desc.resType != hipResourceType::hipResourceTypeArray { + return Err(CUresult::CUDA_ERROR_NOT_SUPPORTED); + } + let mut surf_obj = mem::zeroed(); + hip_call_cuda!(hipCreateSurfaceObject(&mut surf_obj, &desc)); + let top_reserved_bits = surf_obj as usize >> (usize::BITS - IMAGE_RESERVED_TOP_BITS); + if top_reserved_bits != 0 { + #[allow(unused_must_use)] + { + hipDestroySurfaceObject(surf_obj); + } + return Err(CUresult::CUDA_ERROR_UNKNOWN); + } + let format_size = format_size((&*desc.res.array.array).Format)?; + let channels = (&*desc.res.array.array).NumChannels; + let pixel_size = format_size * channels as usize; + let shift_amount = + (pixel_size.trailing_zeros() as usize) << (usize::BITS - IMAGE_RESERVED_TOP_BITS); + surf_obj = (surf_obj as usize | shift_amount) as _; + *result = surf_obj; Ok(()) } +pub(crate) unsafe fn destroy(surf_object: hipSurfaceObject_t) -> hipError_t { + hipDestroySurfaceObject( + (((surf_object as usize) << IMAGE_RESERVED_TOP_BITS) >> IMAGE_RESERVED_TOP_BITS) as _, + ) +} + +pub(crate) fn format_size(f: hipArray_Format) -> Result { + Ok(match f { + hipArray_Format::HIP_AD_FORMAT_UNSIGNED_INT8 + | hipArray_Format::HIP_AD_FORMAT_SIGNED_INT8 => 1, + hipArray_Format::HIP_AD_FORMAT_UNSIGNED_INT16 + | hipArray_Format::HIP_AD_FORMAT_SIGNED_INT16 + | hipArray_Format::HIP_AD_FORMAT_HALF => 2, + hipArray_Format::HIP_AD_FORMAT_UNSIGNED_INT32 + | hipArray_Format::HIP_AD_FORMAT_SIGNED_INT32 + | hipArray_Format::HIP_AD_FORMAT_FLOAT => 4, + _ => return Err(CUresult::CUDA_ERROR_NOT_SUPPORTED), + }) +} + unsafe fn to_surface_desc(res_desc: CUDA_RESOURCE_DESC) -> Result { let res_type = mem::transmute(res_desc.resType); let res: hipResourceDesc__bindgen_ty_1 = match res_desc.resType { @@ -26,92 +68,10 @@ unsafe fn to_surface_desc(res_desc: CUDA_RESOURCE_DESC) -> Result hipResourceDesc__bindgen_ty_1 { - mipmap: hipResourceDesc__bindgen_ty_1__bindgen_ty_2 { - mipmap: mem::transmute(res_desc.res.mipmap.hMipmappedArray), - }, - }, - CUresourcetype::CU_RESOURCE_TYPE_LINEAR => hipResourceDesc__bindgen_ty_1 { - linear: hipResourceDesc__bindgen_ty_1__bindgen_ty_3 { - devPtr: res_desc.res.linear.devPtr.0, - desc: channel_format_desc( - FromCuda::from_cuda(res_desc.res.linear.format), - res_desc.res.linear.numChannels, - )?, - sizeInBytes: res_desc.res.linear.sizeInBytes, - }, - }, - CUresourcetype::CU_RESOURCE_TYPE_PITCH2D => hipResourceDesc__bindgen_ty_1 { - pitch2D: hipResourceDesc__bindgen_ty_1__bindgen_ty_4 { - devPtr: res_desc.res.pitch2D.devPtr.0, - desc: channel_format_desc( - FromCuda::from_cuda(res_desc.res.pitch2D.format), - res_desc.res.pitch2D.numChannels, - )?, - width: res_desc.res.pitch2D.width, - height: res_desc.res.pitch2D.height, - pitchInBytes: res_desc.res.pitch2D.pitchInBytes, - }, - }, - _ => todo!(), + _ => return Err(CUresult::CUDA_ERROR_NOT_SUPPORTED), }; Ok(hipResourceDesc { resType: res_type, res, }) } - -fn channel_format_desc( - format: hipArray_Format, - num_channels: u32, -) -> Result { - let mut bits = match num_channels { - 1 => (1, 0, 0, 0), - 2 => (1, 1, 0, 0), - 3 => (1, 1, 1, 0), - 4 => (1, 1, 1, 1), - _ => return Err(CUresult::CUDA_ERROR_INVALID_VALUE), - }; - let (kind, bit_width) = match format { - hipArray_Format::HIP_AD_FORMAT_UNSIGNED_INT8 => { - (hipChannelFormatKind::hipChannelFormatKindUnsigned, u8::BITS) - } - hipArray_Format::HIP_AD_FORMAT_UNSIGNED_INT16 => ( - hipChannelFormatKind::hipChannelFormatKindUnsigned, - u16::BITS, - ), - hipArray_Format::HIP_AD_FORMAT_UNSIGNED_INT32 => ( - hipChannelFormatKind::hipChannelFormatKindUnsigned, - u32::BITS, - ), - hipArray_Format::HIP_AD_FORMAT_SIGNED_INT8 => { - (hipChannelFormatKind::hipChannelFormatKindSigned, i8::BITS) - } - hipArray_Format::HIP_AD_FORMAT_SIGNED_INT16 => { - (hipChannelFormatKind::hipChannelFormatKindSigned, i16::BITS) - } - hipArray_Format::HIP_AD_FORMAT_SIGNED_INT32 => { - (hipChannelFormatKind::hipChannelFormatKindSigned, i32::BITS) - } - hipArray_Format::HIP_AD_FORMAT_HALF => ( - hipChannelFormatKind::hipChannelFormatKindFloat, - mem::size_of::() as u32 * u8::BITS, - ), - hipArray_Format::HIP_AD_FORMAT_FLOAT => ( - hipChannelFormatKind::hipChannelFormatKindFloat, - mem::size_of::() as u32 * u8::BITS, - ), - _ => return Err(CUresult::CUDA_ERROR_INVALID_VALUE), - }; - bits.0 *= bit_width; - bits.1 *= bit_width; - bits.2 *= bit_width; - bits.3 *= bit_width; - Ok(hipChannelFormatDesc { - x: bits.0 as i32, - y: bits.0 as i32, - z: bits.0 as i32, - w: bits.0 as i32, - f: kind, - }) -} diff --git a/zluda/tests/kernel_suld.rs b/zluda/tests/kernel_suld.rs index ad6e9649..07fc5606 100644 --- a/zluda/tests/kernel_suld.rs +++ b/zluda/tests/kernel_suld.rs @@ -340,10 +340,6 @@ unsafe fn kernel_suld_impl< if mem::size_of::() * CHANNELS < mem::size_of::() * SULD_N { return; } - // TODO: reenable those tests - if mem::size_of::() != mem::size_of::() || CHANNELS != SULD_N { - return; - } let mut rng = rand_chacha::ChaCha8Rng::seed_from_u64(seed); let size = 4usize; let random_size = rand::distributions::Uniform::::new(1, size as u32); diff --git a/zluda/tests/kernel_sust.rs b/zluda/tests/kernel_sust.rs index 831e467d..5057b563 100644 --- a/zluda/tests/kernel_sust.rs +++ b/zluda/tests/kernel_sust.rs @@ -312,7 +312,9 @@ unsafe fn byte_fill(vec: &mut Vec, value: u8) { fn extend_bytes_with(slice: &[u8], elm: u8, desired_length: usize) -> Vec { let mut result = slice.to_vec(); - result.extend(std::iter::repeat(elm).take(desired_length - slice.len())); + if desired_length > slice.len() { + result.extend(std::iter::repeat(elm).take(desired_length - slice.len())); + } result } @@ -337,10 +339,6 @@ unsafe fn kernel_sust_impl< if mem::size_of::() * CHANNELS < mem::size_of::() * SUST_N { return; } - // TODO: reenable those tests - if mem::size_of::() != mem::size_of::() || CHANNELS != SUST_N { - return; - } let mut rng = rand_chacha::ChaCha8Rng::seed_from_u64(seed); let size = 4usize; let random_size = rand::distributions::Uniform::::new(1, size as u32); @@ -461,4 +459,8 @@ unsafe fn kernel_sust_impl< assert_eq!(expected, &*observed); let mut unused = mem::zeroed(); assert_eq!(cuda.cuCtxPopCurrent(&mut unused), CUresult::CUDA_SUCCESS); + assert_eq!( + cuda.cuDevicePrimaryCtxRelease_v2(CUdevice_v1(0)), + CUresult::CUDA_SUCCESS + ); }