From 843a35be96e0b3db242ede589ef6df91433df123 Mon Sep 17 00:00:00 2001 From: Mariano Abad Date: Sat, 7 Nov 2026 03:51:49 -0300 Subject: [PATCH] HFMA2 rounds once to half like the Orin (70 of 2,097,152 lanes one half ULP off to 0, 99.9967 to 100 percent identical) The translator writes HFMA2 as fma(a, b, c) in float on the unpacked halves and packs with half2(float2(..)): two roundings. Measured against a Jetson Orin (sm_87, PTX fma.rn.f16x2, 1,048,576 random triples + 24,389 edge ones), 39 + 31 lanes of 1,048,576 each way differ by one half ULP (-3136 * -8.125 + 0.000124693: Orin 25488, Metal 25472). dmsh_build.py hfma_msl rewrites the fma whose three operands are halves (unpacked, negated, or a literal a half holds) and whose every use ends in a pack to half (through .SAT's clamp and .RELU's max) to dmsh_hfma: the float result rounded to odd, from a TwoSum of the exact product and the addend, so the conversion to half rounds the exact value once. The helper carries `#pragma METAL fp math_mode(safe)` and contract(off): without them fast math folds the TwoSum and 26 of the 70 lanes stay wrong (17 + 9). An FFMA then F2FP and HFMA2.F32 are left alone (the hardware rounds twice there as well). Translator binary and lock unchanged. After: 0 of 2,097,152 lanes differ (100 percent), same in default, fast, safe and relaxed math. Cost: a 16M-element kernel 45.6-48.0 vs 47.1-47.9 ps an element (memory bound, same); 34 of 244 fmas of an Ocarina of Time shader rewritten, +592 bytes of AIR. dmsh_chain.h moves to acdba0750359. test_dmsh_hfma.py: the rewrite, and on the GPU 4570 Orin outputs (the 70 differing triples, 3000 random, 1500 edge; numbers only in hfma_orin_fixture.u32) in every math mode; the old kernel fails the same 70. --- .venv | 1 + dm/gfx/nvn/dmsh_chain.h | 2 +- dm/gfx/tests/hfma_orin_fixture.u32 | Bin 0 -> 73124 bytes dm/gfx/tests/test_dmsh_hfma.py | 196 +++++++++++++++++++++++++++++ dm/gfx/tools/dmsh_build.py | 65 ++++++++++ docs/agents/symptoms.md | 1 + docs/reference/shaders.md | 7 ++ 7 files changed, 271 insertions(+), 1 deletion(-) create mode 120000 .venv create mode 100644 dm/gfx/tests/hfma_orin_fixture.u32 create mode 100644 dm/gfx/tests/test_dmsh_hfma.py diff --git a/.venv b/.venv new file mode 120000 index 0000000000..526a89852f --- /dev/null +++ b/.venv @@ -0,0 +1 @@ +../nxnative/.venv \ No newline at end of file diff --git a/dm/gfx/nvn/dmsh_chain.h b/dm/gfx/nvn/dmsh_chain.h index 45e097ff9b..ec1181b34c 100644 --- a/dm/gfx/nvn/dmsh_chain.h +++ b/dm/gfx/nvn/dmsh_chain.h @@ -1,3 +1,3 @@ // Generated file (dm/gfx/tools/dmsh_chain.py --header): the B shader library translation chain version this source tree expects (default options), one for every game; // the host loads only /b-/. Run dmsh_chain.py --header after changing the chain's tools and commit this file with the change -#define DMSH_CHAIN "d403eb0683df" +#define DMSH_CHAIN "1b36081a67bf" diff --git a/dm/gfx/tests/hfma_orin_fixture.u32 b/dm/gfx/tests/hfma_orin_fixture.u32 new file mode 100644 index 0000000000000000000000000000000000000000..4040582bb1690e602aa3025b7f5c34093f7ae2b4 GIT binary patch literal 73124 zcmZ^~1zc6j_y2u@l5UWY?(UTC?hXS{QA`9B#1nLPNQZQHcXxx7*odNnnA9_e()<1W zANT8h*Ua8?&Y91wHEU+})deyLq|#ps*=acv63TC&J|3tGZDY7wW+)Qdu_(Z@{6*|> zPjp^qJ=2+kU?1m^hb7nPGa~|gjDtAmx82U)%1dNYT1%!c*2pY}-idb@?QN}fN=-!O zv(jZ~)Zcw^1)*7&!WimS4^4gbqg%}>)7|l=(!@&@J^DA0eJot^Bp8_r?`KGF;pEpg zB2_4zWR+8AVC0WSo#!J;9nRih41!!P2yBD?YG%e6xsV7waOtWS>Mh`pRhCOwCXL#- z&Ks}z(~+a~EH0Ty?hj!RJKCh=m0krsw|#e#w>(I->7Cm3kXAIjsWs}|ql3H&eYfa0 z1G4yS!D+c2(E|DJ>aOrb+0oIr6vvR`dLdOgSGe zstsqjk14hrSCJNCt2w3b`{xaf6y>FhPnB0*m2__#sjd>vK2@Y2{KiVw9QGA5{HNbH z8lAIiznB!l7g@dmFBxh06nXJcei0+o8Ldb>z|R)anU zvGVcmp+^cLnMQCbH*Y+YeZ{&6=~K6ZX$fvY^(9CJDQaJ#kpnzEUmV)W@|Yn4B_S%d z9%RnFczZnewq4M0YM$JmoCd+j0NOnx@iioBLe;D30$(-u`8O(sLgumJ+qGKg%s3hy zQcFXb-Mn*D-(L?vADf#r|2DX^D0r*t4X-wwXBwuV(|(;G?L)N%{*5IQ1h*RuP2uw9 z1*9g6t2y0mRue)9RlC_|W8~?)8_8y%O%e1}E#tP&Ad)ty_uJ-A^q)^^)0;or(GkvX zL$Neei6`hOX7C?N%=&D?!WSuxRO{AmLBlRI45F8^iK!p)%&Pm`bBbg6BS*8& z)S!mF#tQ8(8pT@a?ho`dXy_3*cXv1?`sHkyyB8JG?fJwtxp{0O2RG?%r`CmSI<(#v z@xyl9!Fy1bFZE;bxo4<84;k_nhTK-K6^SES?VP%z-ZB7zb_$Dz{^nOnxEJ%z5TEe^ zwKvH{g;JLyRUF53hGj9Mf>)CceUpLHI#_U&R8R=fQx%@19;>ldFtHqJp56}BVnhETY{)rw5AY$)O_aV zWN~wFxafq%(?oyts??@Nn4LeJ5zd#hi^3$oTP+6r^JDV*R+-Y`Cz5FE3>F_aHyeEU zTnwFVyX&{_lt`P8o$AYlC_+?nAf?VTNvF`o019l{TLqLf2JlRu z&R_nooBiB2R&1eE0se>2<^^6u;uBupG$o>ZVF_`Kbji*B32 z_FGi|u1R;c*kWJt9YfN9;=(V%CgZSHoCeAA16aL8yV97;+C9{-&`t#K$ytkvrHrrOBXcH7N^UFiwEny!ih{cw_mA| zrVz@BY|`iUmbPu_wqiIte$QF>yUYk$78G zrS}}czOu4Zuv+64au-#Lq{bnA=JWDbcGaY^P#@2HTP$FyVesCOjCgCjYrZ=2G(3tL z+F{Q7%l>=pXbIjxHT!z!1hEFLx0H6P;Rv-%zOEgYoW@S;!}+QhuRr@QUs~ccWlL*% zpXKlPY8u|!Q_F`i*S8_XUdp0niZfm|VTfyc9wr-c$y%HTKB~mIqMtGMxK^qIia=&X zM8*{aM)}(u{KkFJ+B-kN~bKCm8e3T+t9=Ze<4q&Wb+O8Pr+m~WDIJ`yWm+p^@XuL^M{$z$H z!K-qg{(Q5d6@`q-^XlfYSG1k#Ll-~WK5@p*nCC^j)3#>&D9sr+LM`O{&qTzhe-ny$ytqvPy z{SKan&=>UCfx=2Q+dpuO+{|w!T#Iudnb*Pc>0N%sq5?f@|E=_gx7Wr+vuEA1HCX7^ zeaxi|WA+foHjxfb!Y54}!MI4VHZIhYjDc{MNEvIUy?ADpCice#?X4M(;{jwihWCYq z8cZxLE2ki=npYZ}(iD};2VZxv2^i=-PMOgOEaSr|namVRUR{Kx&yUmYUw_EdzwD?Y zpe^yUy!8IQ52=Y{%G35sDX=A9dsb~Hb=Xa5s!U=@OS7=4=a+sp%Y{=?k%Y+x6Ax9Y zi0B8o6MXT-b9v8KUu}6;lIC?d!{;BNJ7u}gBX3=jBXlv7fN}Irkk9G-rrCVS!j~-2 ze-8&Yt)u!$PK_Kn=UjnCjW%pwwl&lm1yM}`E>eu zH>77gs#?^;Zq4O0r+17LR9_eL+59MOZO@q)M*9GImPKHT&Nl41Y>6HNbLz5?8*aeI z*&!04@*w!L2{yYIPL?TQLtjd1;@-{a--?>ea?|l`z{lp$5m6e`W9i6-Hv|U%M)TXAkDtqriq>6uE5~$;kV&T}>yVHZY5cn!ckVS(}YT@}jcHQw154 zf%h)=&PH{y*TaJCR0TU!R01i=pc>A@w>~20)bnOja#9@z`&czAcIwuxZ;nDJBlJi=Q zyv&%yJyEzTslc1vobQ4rG_U1vHB>Iwu%x!^uQrTAmp<2y2j11}s>@~~C0>&$c+a=J zfK^I~wKmOQAt8&(zU>v_7po)b>k~AD5__h7xuQloic9EDk&fSbv8 zHln+Nkl-9uAb)soGLVvPG5ik?r2@Rqm|ncM;`&t_tSRY=1vs&IlB;}{(ar)7TQgeBWu0{MJBxK~Rg;W)1r_o?cxSkO`y`qw`_9!~ z^kEm9l*KEC<)`QCvvzsPnO$uAB_6PO4gX|rFt|>twEaY>^GZRlf4T*qdbCs|PW;{_ zB#p*9nVffAe#Txh9M~@(NszNumE-VrC5g1It9w*`c+|pi33kTmpxUvRJR5{9CM~>#uOkAZ83kj8|Ot|aY^U%I7g9%2Y&+KMW43@*4nuLT}|MXSEuLN&p zifgvHd|_`9O$Hry_v^PduFabE#laGo^o6j59`sLtDdqGgg+&fLzs`TYr;-z$i63lE zR+MeVyqd3`t?64-mlS$O0A|ETCzbbo@B^B~@j&qtu?~rr5Qn6GvnH z(YDree#(sguEoQ!)w@sW_GPqIdadG0Ttx_}heHR?g4RUFCiSXXYMHK(PdCCGz&HjFQBGfXjYYqojkMpeV5hrcL`b$? zGlfn~fvSCVLnXLTxrR_B3rsWbx4+2tp71*NBH$T{r2jENGH))Ad7}fesZu_Z6`>)H zHNg!{L|n1&PblsCAe8^d8quR0>qw-uwD9v_url&2JepFTS*dl{Ux6{_OQ0H*+4;h3 z>>1H3ABu9FB`g#czO_Ux-)C4s3|n4Fi5jjVfL|l}8SjZ)GXA}m_9V=#R7^dMJ0sLc zLjkGnC|D8Y#l4`mt+K})Fwn~gLl`QynW%=`;0k)Pb|KHG?2|0Bb8|mEXOv>ZC8t`% zC?#~iuWy?_8Df8JBlr7Nw%e9lmvC=ewp`#c8hD0xrddJ$f_i%hAA~oDvJz)vizxZo z8sQgfu3hQTQek@-&AXz_WV@75EeU6G#d}_h3bBQ;6fnGSnCwT`FRtsb)+*oRhyZ-c z1@W9){lOgXXDMlt+}Nxnrh&WrtcZQ(orH^f}2JUE| z>s`V1;7Qea0j3J1*6-Tyek@^X*2;8eE`*sZ8(rOqic)ewJA@5}9JucAs>XH5qw?}y zB=!kf@xn{d!InCRxDa04XTD=3XpQx^S?m{hYH$nB=X27iGH=Y{S$)GDz#EE}bx<1r z=w?XVRENdA`|ZykWa-|Dv3ph3OSg}$Zv%ym`&9a6lNP?nWYJRl3mUKxf7=Um3~yP^ z757zxa*@%$*$ulIUZ0y_ZU2VgeZo^aqj}RU_aAQ`%Le|W$5les6Z^7WdbmzCcVDjQ zQeDar_HzOLh;{UefQW&1_=QERY|R{KTs};QnrG7?Tm;%U9iFT{R}L3a5Guo7zmEJ! zKdoeg3ZHKExczae^?bQSf98X@aRp%!Qo-MQg+uU}3Zd}{o{J3;E#uz=k~T6YpM^Y9 zo!`nSoj&8?dyTuo%XK-(u(&(Y=iMSX+r}sD5ExV?Ov)R(RcK4Bx_JbNxNAT_%taC> zxboW3-p3oSb4;<+M-qL*kUe+mW9ayDwTP7R627~?)Hy%?Rn1fwcECtE6LqkEV|z*Q z^Ei18k^&nDWoTk3{rMjn`}y3KT?fRySERIxnE7iTk4$wKhKQRJQgD8mdf-#~E9Mp@ z`OV&-u??)%*W0>*=d!w@e1GZeeYxE4gF{|*$zuEKRu$LQZDHN)j;1?4`EIXXShccG zIa}?O-M@TU8Qw-MS{I9J(v~dB;M1CEp|uGlnR6 zP-;Z)s(sBOOW3@3W?A%NBf!t|*D-vp1T4OUf z)J6r|W2^32P;RBq)_-1H&O!=mgcW05@#GSXxbumhuE^Ft0sE{|kK|;jKLsEA4VXdC z7*D)y^ewaGkuJ~auc1gQTJ>O+<-BJ>KU9e>=6^*0s&ryxh{;t|deDYYFjZRQ{K=Rd z46`2ybWo9L=X8|4&5p{D^?O@?U3MR$E#ba(QLV_2?DAK4y!-u1s|z^@0**GQKJ6sf z`O_J8>`eAconu|?a?X0P`F!Zyk9Xg*uC;p1_@~=Ul9s+34eQ#tck5+sGivbx{*?AY zruKFMQO(&&k)K8K2=|q{!^@FLG#Bc=QO3b?{XehN^@r!KDk&FCg#^(rLsJMW+V9Z| zB6YDx&^K8u&WgNDKonm6B;^ikN{Ua-P+h6v`$~6^kJ=IGzVe*NU!MQYSw5$yscq?F zI&s%rv|5(B;=>(A+sd}sqvQxWxqY%I&-%sf$$IhnCw$+g91qeteY`mI z!ZwN_cXatD{-1*NUB@(eY=ZOVr^D(;Df| ztf^D)Qy6Mf)yV0;-f&~uFLzd63tJeNb#>ZZN9EnZLiS(z6z4?lS-+e6(zGD?&7!i( z<(Rz?Q8J14Pp=HqIO4wF4NX&2O8Q2a!^ie{%X_#Hij(w{Zofql{h%nCH&oD{Yq|cm zZQkWIq$WcWd`>@-xmSD0vXHRKESXl6P{9c`ssZtV~mZ-r;5<*#$IGb!#I(&X!>k@qbG!KhTm^TwGu`sMtbV9E7QN5T#TON=bK!;!S+A!Y| z+B54~j4RPy+%eGRDxg&}z46)6B0o0dsruMTwx-`B*Uy{Xy4B$By zRU)RR87fJUn^pMO_|~uC(xquwLPc7p)i>IB9+SZ#S)|bEM54u*mEom6wP7=4MWAlW z9emwJ$^D1OKtJS!wNRrt5n?cZ@9YD8&h!8rmKr81=LOd1Nso||rXeWQlaANgYE`f0 z_}j6iLyg!M+k2_ltWpB~OsA(|h8sb1HQ#cI)P=oaXD)r7%xvJoDf7Qm)E1+G(5g`& zzHL5>3bgx~!U%mO^VO^{qQkL=5LVL+jmNv+W~b`z=tS;q`ndTnBle(W!*r_y_jL`) zpi2(4b3`i#jUokPe#Gj5Gq-Vh=$gN7K}1z)ryro|q$-B=W)^2r)%ncI@~^mXqFR@% z;NzEHLL$#Z=eN@B&98bDjPO&wgIKbfrRuZ#pfkmDU~j!Dp1+9Pinv7ehw6T23qz%~ zW_~$^^}smiHj7M26Hih7f`wDVYT%-9O%GLQZ))=eyxKT12l&-V+lh80t`tr9G(BxA zQ?VV=Soh<~vLG%NNvLx1d;S{<=fnH-2%B0U#gleufpeFk?8DL`N%}fCTgc11lqEYO zFEeGZY-jr6(R3+Xl)ZblcmadsO1$KF&GE5hzx!W4d;c_1QCQg%G4MhamRfz!a;%5G zvQ}Zk;KB0zkZSN(XlKJ5p1(}r(vC>)?6e3|#~)2Xb2>eD%g1bf!Fsr{W;uzi*`gpCxd)b&Nwu9Q*Q}rHVH3r z$R(?{JF8%^HFwlLm2VC)GNN)za=Q7Zxfi+CX=;b3rqJoo*F0L^?xNe*BuRX~ejdR( zfz`OldO_$`WW~I{ZLhzP!w**{+$DOgZa3b%IWfPW#IP>COzMxilLC&L>{5ijfwRm% z=pv!rIqsp6wVXep4$Lj2f}-u@h&mJ>lqc>%gX^$h$~Z1apeALn*$^${#2fO?%)n3evmMI*DALn{T9wtPHbO&!P`5yrbKM zoNIez^Mu)H(YUXwPC-Lktliad%<MY^NhDQqVNH$pfNHt`7MXsEHsieBJIyFog6N zDM|K+|9Qk2`lwc&=MkkTQtl=ZqeQqTqU2#KI%D7i)>F=JEvP;@RC+wUNtWYf4R4gc z?)bY1$QQQdMZd~Uf-Eh(@OQO3j)Vc?$XNGmhr!qcreHq`wceL=J#>p-P0-M6-^r*B zZ7zxhr1lT%f~v>&7Ix8$@sD5kBcF$d#G`195Z!41P^EqjQobI);*nirMPyZ3By$AO z%2;u{?tG#Y_LBiNbE@G6SbNAUuNk!?=csVCYPj`ag0IPQ*bRM^F@~-rKhgAAJ~Mbb zr2*TF=s?Wp5i6#fu#ZpEkTsCUP+@*Hm}J%y=?}L?AZc^og)C5>s6f-JVv?}Ev2YCw zT;mZDX#M(z9WSzi_u-ueg;xGS2LWzooAwr3WQzK`kjD*){q@kX%)Wv7{JEltnWn&DVFp0P_zQ9 zQ7mu^=n^s;$4a3yhIG~rjkvOc8jKa_gu~n<`QR^0vK8f{SZH<0x)j>RZ(2`PoS}k1 zZ~CtmE5J=qkwX+@&r#FZ%h<(Gwao!oArc=>>z{=F0;7b@V^@`B2|Qz1oz3;tBA!P+ zALSl=u3F1b0G;O%dojsnHh5P(RyB>aY$$XNZ;T%G;b{;o;d@BdJL<<4DOd;g9Z-mc z_uZg&tX)fV76D&xYLP7GmfZ-fko_&WI)MOOVjis;LXo65o1MAnoI#EHi4|@~mxLYt z8m2YQo%(~@ba5KJF8l#uzdhKS?qUmBoV-bJEk9VdnzKU5V8Iw>>LHd^{6U?QrqIfI zIwVS{T8wuj_G@n@MG`(SM*{z1p*UkiL&*2;L8!j=LEA0t`A? zM%|UfA2uxyQr1NdZk8n!%(6oi&fkQDNb+aIMBGKow3{Ks(R zn7hj=D~wJ1CF}?)W#&1#J{hZ5LhSvuX-`>y&~!rvl2cKtH38`2bV8UZ@*{*{>{kvI zLZ+!3rVx1nx*b_wlw)z%p0N}%Y@>&NLCKz!*izPT)T-nb)5igI`tLlB=xtL9fzI9~ z8I;NI$oTVBezc}3;FVmVq#!fZf-p>d13J-x7hA?3PB&n0@!jHSC$k2_YzmIUrDQf4ZNe$6{SA?!D-e@!ht#t8M~DcQI& zY8Z($ia?t6%%L7v#iET}TJvUa@8|c8;*c(-kqbMPGl#^73>FJfzpSMgQt7iTwOb_g zzcaG_;@g?vX$}-a5LL&GGR3O5D*esqYrNL%o=T1V-IkQ=?>-{4XK%uGy!BynwrxVX z$oV-~dpERT&#C8v+O>HZiozF`t-1d%c!s;@J3;I36e zHPNhN6&KNRvnq@9rF~*=Ywp2048E&nJBGnkvD_szIEq$eDW2$Gyv9EN0rsFG$5CY_ z95pe|CAQSgFjFO@hn&=y_oYJWD5q<7`FgLkf1Fe{OfoLNZq;f@Rn)n}qL`1@_@2EO zu3iVZrxJ>~m8u7stz%I_s=Y+(wbej2hOv=frv%{!dSs(2Z&k$U6@-gp5pN0Cv}yz& zpS_?ow@D$u>36q|$#8k>-&zTC_~>kcc+cI{avQ1C4+OwJ##~bz2dYi<7L3kP4L)xZmKBPI>yXoUB!?;+D`+?ZBAzs~cRHrDS$jQ7p@sz3uq7xOwC zaAfyj9aB|Ga?mHzJ5|Bob_W?N8+%zN66vKO7xTX>ayl3;Q7|->E}4WHSB0mKI}6V4 z$UV2l$+UQE;TnUZ#f7|28JSbD$xXbvrOBnyE&%atJl7j2>)L8i{T5zmNjNmyu}T`p z7#ppV^HnU2HVM{?S6WP$abGsI&NSMVWQB{|b{jo2q6b&&{Q+gG6OnIViFWv!>ln2f zB+T8g2y6bb<;5hdYsWSxC|#_m$bfL2KMPGq+(cGg>`%DD5eB{a@iy$+I5`f>S}Kfu zP$=OdeX-4=r8TVDCREc|XC><@`$*=ukX-Y31B-C(D2JvZjhfkf{zX--b7>t8?+i0L zzYVfX38c*KzPKS5ShH7SzVaFU+@`$lci0jMKQ|}4r>tOhP$qG~&yuoqzW8ZGosC$O z)#y#WKk{omu%~3P^qKi_LDs%jJ(O=IvnA;UvB*Al?>`kFbnn@R^PcC1NQDgYrC3{| zKXITd?FyXE%vK8|C?X$~N=^AIe8{Eh$W@f<%)3R6W6Zh`dmbdqY}F?QPdXkE=5Ak*{v6lnV_m%rxs|5DW!h9nklWg zGrXBghF1ppQO#&V$r!tWGMXvi87_8K$)*Ox6+JG|YoF1|^b|o_GYVH0RzC1=)f!PK+*B zmZ!Cd$xO9cG;c!RsSWkYvT5tvzRM z8oR7~K+SMA#zrSD)9%u890X9NAvqW>OMiESXqcm9QLfhJS^1FbuOg*f*XiqsSSI25 z#18y5s7h4aWb%NE6gi|8MglX1+@z<DpUWBWi%-}zjQ^3Ij?u%!GnD5?Mp+aLD;v6^L(?vg)`i6Q@v4g^&++vlZz>J@ z#a_*ZxwX84{b^x-y0~ncKpS=uE7+F;=OxcMRH=J9U9LGbhp}CaWC_V_-R+e}dx@l8yLh*kz6et;==rA}l zR61ZANq(!_tVdWS$4S*5Iz=~1X}}MErnTdjTVQ+z=1`L;BKarasg0^^jpY~*A{;sl z{dkXu;5p|-q|zg9Do{Di`^r*gZuP zCJ)U~>=c$DeX~)2J`qmKXY9)2|7>`ehkBk=X}#-Rb)95D{sF4qzTA``A%C(lWHrVv z@M$x4o{YEytV~v|H03^Cc+u2?;S9I)L_3Ej4t-idfiY}@tko4~rWkf}EJB6yS;=^X zJiUcnG_zG!ETXcmDJ+rzzox%T$T;$#sh{qW!VXyo3~NI{Y*R7;?pW%KzNy+JmE+YA zQ#4G2<`%?(kihI3UhmT)Mv7xNV~}VG4S@$J#FvijR!c1?^}v%h*9%di6-d&M?$1$h zYMa{$ce|BUy*hv76}x_JiB^I^@sY|<&$rksVFkX7;SqFHR!nzFkVYO$H!9Ubx;d0* zLQ#m7yW?3&$<$+L#+=B=1-(hwLBbtmuBk_JNGZ#;q&RtCv>H*G=&#jHY$ororB z?i&yjbregqU+$gg)60su_4soY8H@xm)*v-ij9B4oM5gz!A)+dG~z-P;H!4mS@$-GLFGU_HG!?~mQP%X4iDlyFTKqR#5 zF`;24*TI}?9)hLE`JN%aM&l3^Tm@mcu5C6@Ke0ZH{F&d|!Z!5XWcAAKLxzoFXDU(S~Z z>f+c+diK$?9at9sDwMs@urvAG_xcQ9t2RWuHQ0E&oT(Sj#jmm~(G4nJpGkL3{2aUz z=9EE5HG;o`7Pd*!c!#ivZaG_#Om zc4zD?>^n>Zwmmay?QA*+O>QdIs#imhg97KZgJ}woN4;~&}>Oe z&#yjB+mA_R5uxXW;>t|=QTE*|iFL+tkfK7>LgaaP1-?H=C<*+7Z*-4YpYJfGRRDtl z)ze@#x(|kGrckRku_5JL)ECL*Chk9df&sq0^!7ULCJ*yeP^%P=#bWgGQ8UH_s1#+T zNurc@>GR;P(|=QZVX9NbWwW1jPkS@j3CWoBL?hKr4(?$66t*a*YG#hEDtaOun-J9u z`R^YDs2&3aMGO$w|K$JDph5W<2o`vbK^h?c56ut+6jN6^ z;$R1S<`~dd#Q(tu+&f_(f&l$L1oR(q0vH+u&~XUpKLYe00r5N;QA}^4peeX8765@k z!FV2n1{MIp`UhW7#vpx+Jh072Krv;d|8n33z5~M!xbOWByP)g?*ayE0?*A|55>Vg| z@Ei(-z(M~`fP&vWLCFF<$7neM_yT&$p&A6J!=S(i{}0gM8P-4eg0l4g&sB=-D7-!{g-`EKL*Sh*uQ}IgZ&F|8w|uB z49uawRD&`=$pSNLk2#(t$5D6cEsTQG3k^!gK(N5)F!%zb|DhRpTK))#5@_#r%noSd z{SUjK?DIeDgU^9B7~r5^7@*+0Pk@Mn_9s+>G619D1c)cN=Ltp_%$UD39W#3XVCJ8I z@xrMRG-&G!kUmBp*ybaE4m1X!;}D<)eaa3f{{zs5KEwyuV}K5{xx^7b3mWJ+1ZY5m zz8?bmj{cX!{$n(p0y?}g^1zpet1a*!Cfof#0b1lwL4!7j-U9YSjXWr0uI9f0UGhN3 zDL{|>DLbHi4A3Qi$}T8l@Hu85SOCy14|Jab^vje*{}*3y4}^hFYJU?8?j#yoN+*EvRr(i@|Nrik5vMyiu-n1xc#Z(;nAZvatW1Zi)M*yO%L*AA$AAF~TA1|QH4+`|BY{spXIUMGOH%nPhzf2js#09ebsK;Oa6{x63+PXwsL zz={d~yJ{Y;KqndYB=a0bEoggePom(&1K`uA7wGrSf7k_eAHY6{9PkXRy*r@)hk$qW=!0z`LG_U~B-~WPHoenejkp=w)SrC*Dvk&-;*FPN20#HA;Knoak zz{v##B>?F6KIroypx^r#KtE3bxf9F-0OI&`h_L|7spE9Fjs|rs!0GM>ln-|gz=AA& zNHee}C>>{)!&!=%sYm>cc+l4SA9g|cWVH7|`54fX(+jKv-WWjs2lM1FU>yK+<(O(< z0boryn8yFzs$5fk3C|Ud$5L1V58$kSnlR5x#1c)gph^Zsc zpbZw_a5s8vPgi}=9;}=oj*hb>D1(taWEa>Id_egyYCwDN90TS8F?0xsBPa$CLr@U+ zr+|2Z9`+H_PfySqyh%6&bLs}=Q!w>o{+<9q)~z^8;kafA9fi3}7aM zS#b&`dXBtuWHCy>eE3T>C?5mn1=zbEai|NU?GWHm@HPYkh^@mn2_QdVPO*o8xH>w~ z9i3uMP6S6*9z4U){1*^wU}pyEN=Jaab+nWCyD}fH(jfkhR=U3#^(eD~ECpik5X`Iq z0dxw+OQ%;gC;>qHiG%n%etY`{r3%_Zw=xG}-$(Krs_pwliH?AT1Yp*7&cdb9w$$o$k zCTCr0Dsd$hDq02q@;8_E_Z-)AJJU-ko}~eo3a=~|&?g>fonc+r<+Nt2qqfzdzl294rI&&|{< zzEwZk4oDV+z7!f8g^@6c!fJoJWc`?&35ud%`viA{qRuLdA}Qh#86rXoX*B!Du1V}`^uBo<9Cw3h9ayBv%ZZ3Pn; z7G^vZW;Vqq)5+urs=70_nRz|k&i!)Grzkz!OD}$138k|MYxsk%{1Wk$zkf?BdWV`) z0~w>uZjiE^`6=Hci2hFPg>9yIG`@ZX4V|J|3Cbl)3=Mhx)ZZZ)TAO5^Xu(Tw%Hv4t z&L^56txu!Beun9~f5oFQ{>157!9QgbJ5e_>$EWyv-&PO`gmyPDld<$I(JfN?$bR_I zSFV0rHII|Df7(R!kC#Nxti?WggLA*i5)0aP+*PtsL+2~ZwN;pKeqZD_l|G5mJ>5PN zzh2h0rdO7Q#g&vcZI?;LKQ^#Z%Ge`XGC1gKHAnZ1E{dS~R2Qqs zLhmWTW8xpR%S#DmO>)+k<#>ykbXjUveK3qhb-h>Y@`K||uNJIF2#F>`@0G~2-(;c< zf4`$`#^$(L|0U`+Jn4R=ttra|;XUJ5R@^{hod?*ohbpO-(P4XjurqpmSH zDCb()@fJ8dn6z)7BH%4k?V0d5XCAI(0e=-oNm4D4&Dg7zMTQp*|G87!Qm{Hc;MIKT z<@wDCeV5PNO8ok9=x-?%p$YliMB5}Y&_t?-{8-Sr&L!XCep)_4eZQX#yvevX)R2~O zC=rbT#~60#>JmBv?~$!tNw`O`>D3ZQRZPAS6z1R5X<2P<{-MIG_QglpQG{oda#O^6 zc*AAaGlGG9j?2@c@gL`JRjJ!Nfby-|b+hJE26R8h=|f1kDiXj7sJ35~Q%S_2Nw^pN z3!i-JGu0T*gK%vLz=Ey6HLl#-GNZ&vWM&~OM%qNeMj}7;$2>0C7e&nY57v1FXlKP9QoPS>gCTj0!25GKo>t;&0-|I<|9#1Bl)x5?KW-FgYhEMqA13LCX zp;FGhdA?F2`qu}e@0DyaCbiAAC!XKMp`1^{6+2rv0GSNOGQ8i=DpC0CY`hw)Vt$RV z-%?fb_avFF00+{?a$h6KqJ@cXU#X@vHuVjViG26sy08UO@{9)ql-E2OG<>^{`F%*$DM_-s*`(` zY$oq&^J0T>BFXg_+oiHrJ;KFYl*uQ1D)QfdTbMK-K_rt|ndRN8fQCMofqi+|2M@9I z<{`G>9-7h?^Np|Q$;x$S|DKC4pU8BKpS3^{Ai3`moRc`f8{c2IWs{EKwm z3U`6kW1BGxyYlzr8N$}HtHKYilB^{|&m(1iM$5devhe$C<0utV=DA4y=@sSdOMJ0R zQ`Fcr6~FA$f{Q_WJ?+2vUnB{Mekgf0_w6xG5C^wVu7BevilnY)!YTq)`Qn&>6Ix8Hr|fD7tUKfa}^i$SMi zwbHRkpP!3nuP=A7kWgAy>c1%D-QGiSi|4=*_Ox7|fpLd_=JAZifot9Q-)otqFrs>ydj!tiu@$URKMp^04P@2k|dn%zf<9x9v7Dv5g zWn1LYBj!iuXETu$BXasL`snK{yaGRQK~tFECM`2~ed{TI?u2x3iL!gp!<*mHmleJv zMmr8H-}%f+O~+Q@=^?v}t69ywbRln-b~zidf?7aOFJI=wClzCKda*%2d#<=fXXRxO z-Uv~g!Sg$8w!}fia_E+#2V?y0WA~_DeIk+gYHip>!w(vf;P1)M6SY36p|D(_0##6d8h?&6Un+X*p2KjZK;OylX6oQztbQy zvW}Qdh`zu)Ns>~^TszUGiXGY^pIUTXIq9OZ;cqLn-`l4~ZFO4r1{jJM%o`aOJ(O17 zKe}4ItKB)eR1jP3w?ohN+7RL-A6V3F1zEW}&?RHLBkJNg;r&?{`;SN;E`dyKVyHkh z^<_d$soRMvwE6<+l*KLG(_|iP=ax0s%00KPw0?$GqD5lEOJ2vZ*_G~b4+lL^lj#ZR zvpQR{KUV4_Tdu6LRTR=-DfVCvVQR*oziiIAS(!^F9+wdxnnnw`;$OpT7bsnmAjY{{ z|F!YH6GC7vsu>QC?=`a{h`BjEieittp@Oo_v0SaW*1ds!4Su%=6+S~BnbeuC2K5sT z!fAC{Cf$$4snS00@Pqx1lwIC~@wmZXj9+V<`=c%^u+DUp;zQOXxz8%wkfF-To$yxg zH6dLXdq2&uSEHp+J`IcubS;~AZQ;}{A%t_@5xQ6K15^CeGz2sLKcc<@EUN8$dx8+@ z4v`Y+?(UWlq!C0#MJz%w5S<{UyBi59>28$nZp6X{5k*D)_C9m}zwgeonc~p^Kq35p` zyf-)OxaSw>GO6&rtKR$~sy#G4@ykj63u)f3)Pfk->b+_X?i6>|acFB1e+plSc$Mm9 zvYO$^^)~#m#x<7sM+PIGY)*E%S_U)o8tS(1@cE!<$?5FxrB<~wF#b+F+|1pjlBbiC zXRvWXW9PEP0YziJPYLDF^*twm3>X+q0L{zTe1kt$(Js=dM6{TRtV%!rc>6BVw3e;Q@wQlpbGB43FH% zz|Kc8T8<>KE{So>{yrrnkHgvXMOnYQhvqczeNq?amZ`=dCGH>`PV-gB;)mE0A6dKI ziNHBrLtpOel$gWuFRG|KgBqe0tg%@+k=XJV46=JEap79m{m^cwC#K61^Cfr1UtXlO zdQ~~P`I^sPY`f2e>5X@P1?O~4Z}E$Omv$|*ubP@WcEc)1KJ1}!A6Hl=Z(n+J!9S9k z#aH`8hINCHt0P*_@7B}&X)C6vN&z@Xy&k zl5S0X>SHW@mbP%pZ`05`TJ!sp)Unp2{Bd;uXZ-H2_D%7UuZgQZt0w}T3OKa zv>}vSb&tv1zvnNofMZQAYra}}r7*I9%_uj%}@&WT^McbZ}jl)G3YpRK(>%q4ltM`W#nPgtrzW%;e4wvy> zg}r{W-o5q4W?(Nf+Z3LQO+)2#UpxK^WBOlSHRD5QCpOBJNBwBZH=bKnL7ahQF}tkP zf=&akKU33H@_SmpRM${8`f--1^xEc&*{KCi&&SLivV**Tnrpf$swf-sc0beNgsz{f zQ0*%}Sz11&bUAvvyE%pLs${nqU#zt@ z4>%T_y&HBf!{MaQt2#p;PMzlb0<_KKE2or5AG67P59T+IX)mMBFpYG19)F-@t$BOE ztz%bVe<^g!P}03c_K88?;TI?MGhWZP_SiUQ(RmCM4Khh!1&ysGJ37nOW0N&7664z8 z*fY5}3hDb-x=N(RrcwmipY)tR{dlo?yu|CD9NA^VYLiD+CbAa`U92Ro&Am+7W8GBF z9#OpHd6_Y&p}uVK0-D3~6FLAd$(U;`FsH-SiF+y}7{!yaS`dml7g zT#19JyBT+P)baH%UgPl}#w9%_zO2MrF+C1evJI?8~xQACCQXK8}nmvn%tGgPtb?_uHW&!RH4G33p~7%GV+&91Wd(UKIMO; zr037FXt%yiTpRc&?j|iSNQw2U6knrY`+5J|r54`=?C8Bf)+1ZrjJFN6k~t&2n~t5h zcdqKl6FV;(!lTQ+4OMOaHHuxvH+#|7XO11HUlzhVJT~I@ z)Pm4)9vX5TF~fD9QMc*umAZ5mf1(|9c{tAtcHY%iJ&RLKNxC7jAzN#SyE2Ng=e^#r zdq_Sw%D1YM+XT;7e}=5~3%kgA+hM-Yz0+>No8MhuHa%Z4VrB6%yM{&$M zJiL9(Z#Hg*!`68?=aqO$KxP;tzWimH)Xxw?T8Q;^d>_FMv^ELKprz>!s^92$8ZiX{}JwxuoLor(wEb~_ds-{Lz{C7s>$^Txi_1@-Fl-&PD| z%n$oSJzeo>*cC{W54aeSA#651QJ!^__xmF|zB?ftA;vdV;#zySl*$Y?UT6J^Mm~?p!I$@?S6aag8ma@X@IzndYa7wdsa3pHc)FC9t2xjqljBD>I&m zH*K0CIyL{W35=c1KKE|!`{2bN5+xu$Om;R01?v6uCbG^Cp#oS)~RfA?{ zq1tUiYon*K=3iQ-=9^eTZ2x$T1D9WFDXvJ_=F|_9dp&%8KQ4~Zx^AkE(+PE6|2KAd zT1QM+JncCBFyDR3g-`conb9q#lrk^S{==Kg8Vo%h=D%i$Ty-9*g637W+)_ch^L z(jR#ga@2IMyHBLM4>EC8XP>j(|51N->D=+Y*Tk=j$9wPZh&tR}Q8w??sec?hm=t z?)DUaR9c+N7u~2f$RSWQb*gZL_WVh!2R|)W)7&q4*g0Fh&&xiHPPx$WcujDXc=<8~ zs`^OfTaJDw;^$sY=TXXX>=$Zw)+0O5y_wU*Q)#u##uh0UL@x%6Q72xgR(@Cd5hJ&8 z!y4@I*VOrNu5po|!lE-NjQdJ&|eY#bv^Rs{g%Hy_Ni@S6}Su%=S@SiVm? zTA7(;ot(OBo-Ix9CuZ}9w?X2Sqhqf{vi&J9t1~@6=R&rZA{A*aI-a|3){7r)ab#)3 zNpet>QO~~1++!w9tKgpcM9~s1072XeCvhQj?vzm_5 zY?RdbrBG4qS?70xpW)VOdFH{FRm)4!Ez&kRWscK~IJ#b!0;4;qUgI8W_nVv6QJd5r z{Bf(pCP5|G<=9H++ALjSMy)$#*a+46tSfW59&tN6E{o`zBL;WnOfGZgwWG|M@=%Z>&A{BfnUy zJMO#quV0L-5iS>MwGN9k*vzGR(u(!>$Z=$mKdtE6G#n#7xqqVVdqveS!OPM-ZybA& zFosXNt3gXu_LFJsa+__Bf+2B3PdN{BC+d{>ujlGQQEw7VeDm>+bq;nnCh^WfylmJ{ z8wnp>$88Vu4}ZiY`e+;C91A+z*Pk&u{rM^!M|WSPO{fkzJUaO)^1gT>6b||A@!Cci;b#jg^P3MOtTw?{NGk;+lGzT#$__4*u(q>GKkCg z!8bt^$<5}`#6EB0i`NOY`(I0lqBhwxKkk<<-QnakU39+Fn9tdI@H3e8*J?3&tlAkG z+S$o0**L>=Rp)uWk-kaoZI*bC1+h4@E=AQ%noZX|HGhv)?O?Y(Ght6bM zQ;AJ*a4mf;&t)DF&+6bi$DY_QfPTTW@T>F{Aygzza(MLmS97=4uUAViY+$0!N(-;< zA5VPI@@rU^Ell%0N>91Wr|(JNgD|4dmZBtk0S)SiH={zNPQJ68h-YC*jOgK$FGO{& z?OF_7=cL9&5Io1#c0|5rPhx@KcCYI%O^cJ5Ov#aQrw;Ka)H@t=I<1Acfk`%aM2=MaNZ*mT9w6{`Mxdmsx2OZzh~lozVEe zGP9^x^B=8ypU#!v$isBUpLBV^oz_whOqscT&ht~dleHktUWeC z+W9a!TBy>9LZly;FMdpIcoL=CE{|(F=A5}U!nT34bw=^t4x-U zAqsQodeHr$yCmh)L(Vsg?`(8B5=B?qY%0U2oh|21U}Vuh#HSv+$b#pbVlL->BrqFO z{H$0HzbbpSE=SDT<^f9Y(fu5aFMX?H5W>-5o`~=apgKGe}6)5kJf#+T45G z<9=zwEo+n;8iP0Gkzam?%U^p$trE0pV~6+RdnMlEdU5XJ)7t`^0w$TIOBWv?(PJTZ8 z8sLC)h|jG!v;8BNLJ54{wDF7znoZ0`!yX}1Dvt_YpwE$C>6@6{T~3TmwIM>nLzp!TpKwD;m zx9bP2qo~(h^eto2B3xWE32LtRdb)R*=*9d&*IPc>Y-u{Wi{$mx^J`>Q_8N~$Sg1Pq zeMVM#aw&UD9^#j#oJFtUm#*i~od1;QK&)Om{tMgDezn<^OTNtLw8dh%IR%>*VU1{y zDVLffrpn2n6Afd#oo#p)f*nP&OR=rB%%^{1R6bY}%T|Pl`6%v`He6SeqNMixnU`PI zZ=A-57S&8Fttyn~bhAIHS8Wp>yM8*DuzX@pl|%77ZQbjM_=d0^QMP?KlUgno>=V|B z_N{m4@#noA4Y^sVlDztnxK01(#kGNGiR#tR&iKT(`m!L=7NgIaALAI|1(3Q`IY&nhGumPYT~TNR0KZQ zs<__3w4iz#A7eC74V^xTt8A%PVvaK}$)nY|hlQ@0EIn>;(;4~4ac3&1E-?KXZH8`< zImdL^a>Pt#={vHnVmo2Fo%_KDyOO5j2GfKd2`vtTmk&ngF`O+oWvq`o`0kl;huxY= zt@9I1Fk??Uo_B|x+@6N8HPxB@hVe^A5iSA~D;p+%6}OWfTC}6cT0_%uDM}Wlq`y%v zM4P#=Z*|c35G{;9F-Vi}fZQwkK>g3$f#zUp5BdgoGH!+-Kn!i!bfh$SKYY7rCq{+I zBBzcv*>NmEq{ZB*OQ~v{uKQ8c1jec8Rz^>fNNw0+fiz9Up1sKGyJu>sVr9Y9@6I($ z?i$_;R2R=EB?dWR(OEK7ifiF`TI_||m%5V2hP(DD?{NLa-%}x=XpJiDuBU}_CO0q2J10X7LG|ZZDhy@*Ono+JRhGB8IO{!P^W-n)C7we+-0D4H zbF4r*C3Ppg!UeS!hJTQMoND-JQP|95Rx5BdSI9k;s-)MW-iUt_P0Cz#$b zJGbEZ8|4I)ZU0I;IiOftHv8S^PYQ|6(5H7CI$2vFd8FsJ#BOe|Y*E9%$yf zVJE_U&7*pbd`|P${fNdKsWvkCdOyTnAQ$gAbyT{bOs0#RV8ql;-(c#8|U!M|%wMY_E%oCi4nRgTAKSp-3tT3zuS=S&I@P-m5s-pS$29) zcH>B?6%9rsz76j)-YVI$XjIKwG9=l{|GUIp)r%Kp?S9hdH5td? z#D^|8=B84>(=1pQo@xUAa)Y4e*`ve{s(;Wcoet=M_HQt$iMTvCZTdfh-L}%UYuQBKW+gv zwH*+3RDjX0=NduAs)UcXDQvviZ|?Hr%Y7N`#Y3OwvhidQ2dl9Xt{MapIxH)e-84-0 zn4l|t$;W=V8}UE&?tANKDDIvPC- zOXuHW_|#fCW|XX^bmRVFVqXQ+%iqe;pX|BfQzv--{lTBdx<@n>2w*Qo`XBm*dz(-$za##8Mn?T4;}3!#>ZSyp zv>L}v%r@SnP?5^?!vpmBThTQYL__K|Vi7L1(}bWXThUIDTa+=XwV$ZML*U4}dWo6V zgD*Rh!LlQ^ z(Z)xAk6+cFm1Y;AeS-gdm;LgesNtvhfOvE}K)Ub5az+b5)Lb zT#UWMXZ_J&VqsqP2<25P_XI)93!_p&i*!u_0t1b1?fW5dE+*b&@3Keb#NWx`AGc?B z-*DR#;~M%}7mdfbR&}?9eO+J`H7)3Z}KIx4I#L zszepb$eR;)iF~8|6aO?8b%RSQs|~RqBd~G-BPdxV?$?Yb(=&y6ag~!3>?3{bt$wym zs8b(q(j3OrpkI%N@i!Y8jjz%vsU02dCZ9f+OL#%Jh?mDmEWK##Zv5S+o8}+yAF&)S zvzB0?O~;&7F?TsYJ?qpOc+nE)2uJemK`aNFJm$?@K$O+g>&cc9D~YGm8|YQc5A^Qt6|7L>{e@6Z z&MEU~&_DP4)yfqnZ%o4XuUuj~bjxD)1mc54zqH#!FEDk~FEDI{3s?&DE&kd3OQI5q zb)L^ibrT z^2MX>6`jNGg*7;O-y8XYJkR7;Pub`$vY%?Wj8&ngY5&~yM@bSjDEl4bbNWJsVBwz0 z+A{)9EBCFSqqPv)LykG?2%%;^p6o?zdzDagBN=B@*U#cs8UkD0*-rtB62!(3iJsal zuEJPFX4?D^Hlm`TcZRpM2<>|7C3L5>F+(h_#5f=4D?QDkRrBN22}hOcq?wf-gwr2gzFn4I%n~drz`jK$skM|dWp zhHP})mu7b2B9}tVB{XZBMKJ9l?Ut0qOTjE*lIAV4FAhzAx3pQv<<033CQo#~=dEp| z)rrq$oo5i$OBed3PPF4dx21pN&BUq_f0-HR_?sx&+(D zOx&@fUEvZV#L)GY&XiO?f<2+O_ENedo_FG6<`N6!&m-0G@481KU7S$Jaoo(x;#9=V>X zAkJS18zVZ*rBI7aJRVK(%WAdc<wp=80vO8%KWx?)tghp@h693 zyio3?ukG`gX_mYODFsK@$%6aw@{wMe$Iv5&d34;)%lN*i8yEZ|vYxmdZQ*w7$Qc`K z-O$rto$kC8X-4lbVQ1VFo{x`r4(YHd7dlwAAF@N8;)*vw4ISz;kbJuq<2BpD#80zg zq1^VY{aUnn{@I$?l3R63wsx2GOR?xbMeziRxbevc+%YzRsaYHtrWZNG;(aqo#5b5o z8wpFE#^!M`95>DElOp`fv8#tH+~~0_#DcXdhT~5~(QT`8rV)(#<9yhh7ExRtc~FMf zLiETx%$F4cks`E2@9Ga*{_7#n@RP)IHUhRby{v*14XH0b5Pq|WpgQD)O{=9&j&gI# z4aeg;2$M6(CME=4i}$u?u}5jEaMTyoImNrk(9P)PXloh2=$=m$YQfES$D^J-#B*^@ zlz+FWd?H-=47-l5ij8ABL10lmcSw4pULPaz1uJ15f(L)=i`q!m+NjVZ4kf)xhiO$1 zSgJdO50$#BPjI-(XIUKj6-Uz<@U%lSms9#`u5Dt`&H8s570D&o=mn}lhG7dy8uYio z8xB3R#d4pA(deg9?-(WYwlk!jXqPK8#&=m5-9blZF0>noc~C?pH(}lltMT#5jtp|~ zKbFj^7|mm!d=p`XzdZW(>-TX}tyU&wWp=xidpEI31D|NOpLDDFQ(s%@mSdtAZK<+l z4tklW+La*pBE=g0X*rvCuG&j*7d4HJ!u1j|C#u7qSw>9U7zr}O$7Hjf&z{1Haz==^ zvs361$5*B%XzrroIDfjV8hyY&P)xF*La(bD7w}g8t2pWAvss7@%jdL>SEMx>I`N@I z)m}XE50M7PPp8;7ZD?g&h~J1;EMX@~liw~nZOf)5FdDCIeRMD9(+4}%39ie{T2EqZ zgN}GQ-JCi~VI5{sY3(saOy$iAc{u*DQlIdt-E@sy{64n6MrHa^cn0xB&8m<-&H}qP z?w6&X6JDsb@K2Q*w{&c?#Wq7unt7I<*{+w}k`Nc8CF#LW>^QuxWg=u1>eyLRB!zuV ztd~=m4eQxKmtj>YBP6KB9n-mvpf8H^NfH!^&h;j-flW;4Sn3&7T$u{dpqie()hsoP zH{*3LW!Jut0cz@WHlbR3m?tbwoo#K5w{}fV;$lf#EWSvhAvy%h+w#kJaIhVV`SXBg zZNbE9eK_2Tj_ac#^*U2l-OZD$U54grYRua;jiw@)R<{8x#bSZsd3lPX6?9rMV|Ynn z1{R-WX7|aAV01LkQT2)7h|}tV)Vm?Xg6i+Y`NwC;uI9yHw0lb$X=zO(Z9i|~hUvp| ze^DFjqQV`E*R{KZH8B-#?gmea1`&hOsl9)#!M&&@-<@)Rdm-p zB%F}?S1ygwCQ_06B)f-MB;APE7U`bdlD~FbgjN~)eu8Y7!hEa{)KOudyl8eIauSsTj3ZFPusXu z>~_`N6~jmi9Mk5Q%SIVwYoH^Z`Ikw&Q;D5lzSfDw4G|7b9CviMsX6=n)rjI{u`_ZT4aum{gcgkWSaU}LT7a|i z!mp~vQivFp-0$@e|R#uI{DN%n@8izR*Y*!|Jz> zxZtq1LTc0HjgJUxW?1H*T6+doO#=9w(&)rDQoU)rnGO2Eb28~`)T!Y|p6CejX42yG z(Z?41?eQJ_bN*Mmcwr(!jhz+d5T2p=Ce9%qljoQ~a? zp;v6TfN5e zCVk=WOl!Nnd8*;G1i|u^V5(>Vd*8;BWPJFPH*Uc$GG?pWoUJC96Hks#L~EkWmZS-0 z*t?uxaU&^mWieBX#)H}v&80X#`PXK92B&L|&@uJA6IkRXtQVZxTS&~35A8nAVRVS) zWjMF!+0b`uIZ?h{vqJ5nH|Hm6Q^l6l9V2>FQAHl|`6^CW-&Avg?~!SX>1VnklT^-l zoe?MBmL@_i=7){`{As7NFI)~3_g=q-9;S*k=9h^f+K~I> z(x;XZ$jFY`?%_Y7L3dlPGQ+5s^f-kYEM_M4h1}a%P1J6)Ew65z5GtkdYjMXvu~T_w z<1TWnUf9EEH|cuzi!)eB22?A#L|Qs_HkChjukJ;(dJV?$c-BIT6uS+mt*=YFF)U_J z_9b_0VpkviakBEM6U68!>a+H~a@{X$IfOl#)*mMNr0ug&jX5 zD=BM;pJo@Be~TAkOL2^Q?11hO3E{qToK5)J_dfjIyp!y-D_M!Lx(E)1Z%u1zY*uhX ze@tDn_nyC$^PF&rI63vUQtxs5F;Olrqwu`kGmKj74GW8oD_;eB)l!U>%e1XeQ?eOI ze@JuwduFiv@gvvLY@gh5<1s2JrLw&cL$(>wZ+^BR`n0^qF7W9q2v}N_AIUOxA z0DlT1A5r6!BPav4A8M8+%QB8X&}fNfLulpPzWj(!QerJ3G5OAF%(HH>h z4`zYWzyN3rggO8K0sHA7I9=F{0wQDp`T}5Y0A>)pg8|NlNGGJD2Ef5U3?(=T&<<7s znn_)DRN$_F3pf+vn*iz(Vo<=D5F-NcV?f?Oz5q7`Ifd^yumT`r0nLCa*wX{f0{fBr zXuuvBpzj0V6hZj_IsoPhp}!6A+y;2p0QYo|0rBH)U_ePhZ9yK8(@3;~;}5`9gl+^a z0woQ!2GP|J#|RNJ2yXBWpa`HxAPx<|m;nR{#C;+7Er9p|6e4&RM4GU(02vrM5Jd;{ z08v>0kqM-VMP5)JfHxpDs1s5b46p*sLihrbA@z;`C|DpZhy;Q=lK6rsWzdBNaEPEv zA$5Q!U`yZxfI36024GsjI?y!;XcusSydXXi1vCjdo&b^%3?kqg!4QJ|Nw`1Qi3aEi z!RH}-f?aUH?jX(zQUNeQ5OV}{K#K*Zft6%nXV^CdK?j20fVRL3qM0Dou*boFwg-3y+(9HQlnyu@ z0|rDeLHsf79t5$^0K1Iv1^EHChu9jxH9|Z4Ka3(+L&Ew2OdptqVBB=<0VO~K5OoUf zBAuQ9h0q287l1Pzs73Gzc0Ph=W~hC@J=hr>M~VqxFWAKcMh1XeB%#d!?QJju!R}6A zq7K+m2lh4qYe-#KK*R#RNkW(dzXGQN?!fB+*dGq?1KuIp48hexPXRWC!~iHQzz2{u z3C9O?06RD#HU>X~yMPWDl}Jc`fbD`D0&1X5ku-=yAsw8M*Z|PjU=690j0Sjb1H3?Q zfIV11pA|qMj5$yjAie>eKwg0N0-Yc|51`k<7o;7~jl?yuA7XppC*T6`18fiB3D_Q> z0q_O11#A#pBVoXGP>^yk8}Nx(4kBd)+Y9Cc+$~9u2sZS9gc$x$D+mq~$O32%crxfZ zgy8z1M}hbV=!gD=@DJW$^h7}mK=cYc9OwoI;uQ{@jP$Mo^n%O*^a?u=L7N2#P@pYR z2NSp(v`_F8=u-o5tbuTg0x|(Q2KJ!=e(Herh<;&yLSjGY3jyL4(z^kyBlRSJ*#A2~ z77!r`vIw|K!mmS=B%lUT1ToV{mllBj1XPk_0IMEC|gL+TJ422ljyi^0$V1H>1wFR4on?6(D_t$}18=!pp89qg+Dc80M<1JK$A zbo4)7L9Qj)14MwrJONIIU89gJ8U)^Tpszx!1$@95q66p#^urhlxB$KaC7L4x<&y$AchiLMf3>l4`dL)7u2B+I3)n^06HW6j|UhPARiz?fgNDp1ILj3 z6YNjwp+wRkuwD{RVD~n_A+&734d4evKd=YT;s5wSbP4uDdV9c*NkGQ%r2AlZ==CuA zL7kHNG{DG%*fi83*aZ%a|6-^P34IN@Li`wFr!~L;_6&ZKWB~Iz$Q9ssBx8c#;O;PL z;NUx0iC71;2$<`U%&7sa9z;MO4WthipbF_z1@<822k6T{qd<0TfI1CuCakM;z|KI5 zV1S-2U_Ia&obq396hIA-K9Cx;a^SCkPf~vfV6R{Y1fdMA8#xK!{=r#DmmO#kkUywr za0=jrlrxZe1yGAvY7jUT1qOTv)&x!f(GKxNVABX%9$q8Wh%jHkm;@yZGZu*0K)xUw z0;>Ol0joMd5zLN|FGO#!r;834_u?5}*O(c39Q@rzhCI080Ho zUjvlGvj9dL=mXf@e=!y44!91qUL3drd#!;!Wr*zC5NRM?cEE|ynn1+~D^kQ`p&k8( zo(p(@7IXmagJ=h}52ymx4?csK2zLKp4Mp;Mus z&={gM7)ePs1ZIQ01qRG|K(>Gy=zBOQZJ;%vk^kfj{Rr>`^bO)2^kTptv~KVMbi$|y zq9l|xkfIKx9GrLntrBKe$SvRq=4EIlpz;L!BGnn>2h2uf3iv}}7_?%fYhN2!2_h^w z1LiWM`v9O2>3RqDxr5M+1G@qbgINLa46POE*ayxf^}qw&&%hptZ-YHZISb}p92n5A zACTmaL_ip)fUU!r469^PwMe2H{01Wwh*dxzfD14{?FVQFQ5Z@b>Hztq12hQm{D2xn z=73KUejaEQ)(>DGSScep2gWUMA4nQd12zY82dlui1Fp0Is&&Awz-EB#$iO_L8y2_& zV}S;k1@Z~J13p6ul7XQCR)AHo!w%FE62U;-fXE%%5Xc~qcJK}&7|1EG21ix{8~+U@ z2(1BXAG`z0gw#M@09lCa!A+1!042b7z)wgej2OUA!Ay|7!Py{3fB~vk_(Hr8@BrqK zo+~s+@di*71m=;f7@h?Afz&?$KOryBc7P_qI%r8SZ-6HkxC+T?pfUuy0MbWvfdjMQ z3uZ?chu{m;d4Os}KR~Y}f6+m*7peaTus6UL>A4SjAe1`HUQioQ(jbq3GwuM2{!dT9 zp8-X{N|D<2zwQVKzo5DV8pk5sgD8XO3-l)hyTB}j)B?cn;2prQYM|=ZVAr|2f=7j@NYyzJU9YI znD0OifIbe+Me-e3i_~tA13(2hnv~BW-Cz!>tI>Zs90e*pKdJ&|2><6wTS(Iy%C=VR6!}~0NS7rfJgA)%gtw2&p=7W_LjKbg>3B6#{Q*Wza59oFfDDkD2_tzZBhKnVd3NxVQyg{y$2!Hfwg0XYN69|J}Z7~p9Kc@71> z0q#TZLShl12>BigoCKmNB6}2a0t%c0c!IeC@oNnvS|c4mVWk3W9@e!;?t{^?f)ioA4x>KsYVfHI@Y@FV0>6WkfxjRWz?^X(`ZNk!G03i9)PVt?M<74I%7N72 z3}7|?XIFqAk`*I$JdhQs?f^c4*c0*{2iOHX#i5`*fk*&8A#oPk7FZ2q0F*HF2beKn zTt(^(AWf)QGH^0d;{h9i-)E5OJXpCQH6DzuKu3@-@D8$>28>K#9ps)ADS@TI3=DKe zs@2;_Jc4~;96)-a0V%)=1^N#u{zCr;nH`)-(ihM^@;n801k}R5KS(x!5ekt|5b$pz zkOX)Kl)(6l1D}A^1AT!PqBmfN(6iu^3D6kQuNcT3q*mnVhYZX{VlnVV()S+VnFQvY_38Tn`jOs>wku3o8Jy4oDnu1fK$+ z4eF3=5ZWJ7<-te}B@A>0X28gN2kZ>J@($pb^c)ONgx{n2kV8n$|gOUPVLrX+>)BwAJS_TFC@nFEGAmle3_=$r% zfE`KSB@#*g1n2>LfdRf91G_+pk^Ja@)Bz35hTa2rCB-=?ZOG-pe|&-13{FOD4y?Zq z?f^Z20eFJ30`LiW1Qisx0<0X!iq=J+Knt^?StHAbRMhD_*+4>C97Rq?~=yJ2alcIDOtOHDDS;uCk45cr+^Z9#CK~C z{8yiPTb+TW!0B(@YkOi{KZnnvF}$w$xYCSDy$bF@S@#zq%^|KBu>`C%JsJ_(&1! z|N3Fk14fI{A0LR{^PZC*!~VmUp4OupmFHwww4ru7lUDO2Ux1#Ys-p(eRdijenZZtv z{7Mp5Js|$F!SKF{E8Rrc*p^sHz~$LA+pOZGzPxi;mt@);rCUBa2%hcTPq(@wcZeyl z*H)4%L)T1zjrD8cTp?2vm1?pA@tWJ>+6@icZRxDMGZBGG=-w3n4xL?&ml?;UZXJ`P z4e`3O?Hcpe1~+?(C;X{|t}FMBI#FiU9lF^~xlJKVOy^&(3nR!)dQK#8X@A_)waYbruI8+CfXO~xcRQ37R zR~>mDRdAH;5Cg&WdH!7R_HiTq3k)hhpJDf6CAX}1=dR^^5NfLbXt0d?;Gua&mvQVm znUCVtnXbWo(IF=OHS6fUVQS^Uvz(36hMC{dOE;)#qC;kKP95^UcJyu8YVk|0e_Mtg zjM~DcWsRA|F9pMfB2aEy|DAq0(tJx?Xxa5kbJfhC3 zO{{-x6*l&+sZ2ZIlo7+r+TD~X$>ziNl9#ZhhiLFBf;=7f<$iNA`O^{ghdx-Hi;>S| zwb^Q@5>caZUJfKzWZ7alK$*tTbG0?PQ}~GQxsw;}i0#JJ2*AjfAs%iBG>Vp1^ho-?qM#)7)=CElPd ze6|o}Fg;#(B$h{DNQ$=SkNHiUpHptM`M-f#%KDahN6jl}{OVJ*itvT#^R3)M+F4=a zikm-#?($f7i1FEASQi(KS6ucxT2no^AAFLXNiQ94)|re7b1t#Pn?ve#UC;>G&-mf?Ow=w7=<=?$5BQSLBvjND;b%i1J%CDlL8DZbaNL-KZp@`MF0Fd<;hFD`(E-dE>rzW$6fU8 z^YW7g)nsdlW!=?X_v3YRr5Zi|X*a3~WpblE2A^O2q<&^Aa%41(-O#ydr>}*Ye`v`J zrD{ev%VhL=g1KYA?x@k*(tuPkE+gv`HIuAU6#i2uHY=r1?r*5;O4fOxm>D)$JU$qWnxuHIrtWBq(u`TuI-Ex)m$yUNSjoQop-KH+U$ zCr+m0lOk%V?U@4nDUKc?;|teg&VOkC{oujZTUsBAX((w|@m~X1Z5fSu#Le2G4C35Y zl%p8FqWwK=ul#e;;Cy{Io{ne25`YGeKK4z`oKuq|FK@|klM$=U*12A{OMCg-kObAO%H<>Ek78EGCo+`}k7fQG? zTTk_CG^ADXEqbTIKu2xV?RH(96&;AFcWSgaE_D3R{)wLt|G*|M8Nm3n%wOvLd#vOT z2lx77I^U%$cy7_J<9G{d)!W#;)Ku)pF7%YEk}%5?4Tql=J*ET-xBVX9Q5pxcCl3Ne zibSz@Ck1c{vLn*^?`woFqGD%FCOZ4n3?q4YvJ$Q7o0S8U3NKr*GSReT|GOAbXyLAZ zRhTEev|emPtXlEt2(8VHC~I0w%2s}i!^~Q?D>Y|1chca%ERO-}jap5$cHXhp&UfpC z(se3VcM>ZTu)!V;6 zWcI@Mt=5$5aa&X2d^KT#b*X7bBkv$@with!Mlp5D?vtyHLVW|CR$i-;KjYK)oHIXn z=^yKJUl0_`bQrqJ?ZR@7arT6rvQ8$t%I3KLd_dTSZ4Y~O{@8|ejy@I@jE`f9Pq&HGU}0i!S1y(*1sHPmt2&h-SVCcO|Xxf zWa*UEp4_Bt%m|~5(kZ*w^_pos?Vn~;!FLns(yKDd#65{wh5*Ajqj$dBs9&YylV1(a zpP~w5^WRWiyrf~4B`K%P>LA`cMBuVNTYJ*t4fbbsumAZomH?B?~#Thc@a6Ec)t!^L6(qmptic}w_b zX|>hot*t_$BTS}WzXzHpG+=UWzNK5m)$C^r>x^8)L@wCWpWcvO1P*7c)A!?3Ida1FoGbEOZH)&@M8A4K+=(Z{?x<{~PbS$F{~BnO=w#l+ zuYSGtDAUeb-FZs9fj>G`oa^>aJafN~jOVs#nO165s>@UA^d@^|2lR0+i`J!XW)=6@ zQvaqS$=4KOd*sc+Zaj6$WR(mu_v4J~80uAX3ZIW?dXy9K|N^gi1Ix~UXWaR zk8_55dw#aj*fv&gyEM5Mzw42b8tkbOuQC;E9K?yH=%kpolHJ#J3AYn z)WYK?z9&*UpAnqIFL2mDi>p>;PsLg}InU)u+zJ*r!YIQ%5=2{hwW#=Hb#uOMcvTIi zh@6wVsuuIBY5r(5N&!chaYaelDm6N~Ir?|Y9{saa|EFK#R1f~ZE+3m->F`ZA(Ge4Mb~j9)^Vru zFU{9pFB(@2Ux8NCVO8H2r!P5=E$5yksQJG>X!DQ#-k7SaceTrkD4m^%eSEedIeAC- zlQvn73(B(>9ilJ$rB~$CKE=zHqkSA$j>h(1D=J#=>2t7W3`_&RjSOBnzgg#cic7BP zA|`s_jhk_*PS)$V|avJ*FX0Yfp!T*+hE!R5O|H?vnuePREX zf&RCN=eY?Vr6Q!a`9tkbSyGi+>-oqYHuWy~IU$v2sLqG->%pF`it!zl4?0m%!Hrpn z|5u*9$`vd7rr1^3^eiu#d}+qQK=m54+HP4=a8j5Wm*n5$Dixn^Imzl9Gz^^H8j39a z8C?AQ1KT~T#Xbu0n$^GNXR}HQ{|sraU#7a+tX)2NbX$Vqvyk8f8fP#jAiC~tDs7c1 z_{@xo^W0EULQwqXn1RTBrVH$z1RDcl4bSa-!W)%+v4cxdW-en|vCsLJ%dpGh)zIBCa?4P|mHFVG%E7YYQY@miLuO&|M}%!?Ip*;dXI8P6*c z34R&MnXFvBzgClKAb|IDP?DP!k9`5vX*?MWIIgUc;JK$|=lM%DX@NUAF1ZtgYXF^46nL+J>rV_x0oG z5{-TGXJs$(KHq9|GMK)mCNwM7N@Z5^`ZPok^y8m;&aqj4n3X7AB&-8Zap^ADU)#?TtQu({*GA})Dq=tQzZ3B%RH z4pkoRT+6AJJJTwrz1`pQV*cf_)4Ea{Y;zJnIHUh9&d%sG1iS9Sa z;IDd59iyN}S=pZ$QqF25Fu7L$cBKix%x=B*lRDd)zti&j=-*zMrk^5a@fN{K`IJN3 zXNvtKLXXKCv65xavt$@Kj-(R0_?dr5OZB=OYTO?2{Bsebu2(IQbW={Ejp6&DcQ!WP z%n2!4nl=nsMRy+sd%z*H#pqP zn-wLQvzN|}UZJ(X5{p1q%d(~ zyH~g&;D42mxKLrYG1s4>+8jT5QGI`CD#lz-?~7TeTmhN>iJI7oQ!759DoUQ(PtsBo zSNso|#Mg=+Uz?~_X$weqj@rLrgnPJ}`DyE$)EBvWjo!w@bnoZnKO9orZ3?n(tP~jw zm*0Cl*=Ee+d0U$9oZ-hV{5bVbi5%hfO*V}9;78NIyX6r&pHDsGH&fYOZK*A17O7bG zwddjO5vAU+t4nWayQ)N|$Hx#|b@VtvBfT+&Ev*S%X!S<8hsDQ#er7XSZ@R7iw?g(4 zrQbap6)JrG{^D#afy^(70(_+Bw^Se`;yO(xiHv% zN|9%NPGv5cJW#cm>bab(#XBih&SkHngZkp1A>SrFhA$hu(-PNOPgQY|bp3qD!9-QA z0AK5oMSpRa*doebB|uZ~@H_>hW^Qg`VE2=RMq@ui`J(kJ6u86Br3+7<5dEX-?}omYfxZQcj>5Iq#?Qf2mrQP7Yh&$bBlK9mFM{k^SYD$m&?k z4tkAN;B260Li}pQTcuyMT1C!uVmggq0>6_r2IcKH8`e!Nc3Uu)qfePy>i@lc`O4iV z=B)oZ{%T+z%-IW1HPe;F7#6D{*?ZXIqR(u6JtB9@yT(;})=Dbff7^37y8q%}9Lv?C z#LpMH7vIecJM90MGTOWCnJ#ZzHE_A*dPl;0>oX?jqz}CcA2&+xW?DV`HS-M>6w!FT$;&7Z2?yi$+vC+~b zQ`_DD$JJTDMbSkKe|8sx5K%%&L22pk?hvFwQbCbWY}C~)5s*+p5Gf@@KoF!;K|w-9 zkS@XQZe5@K&Yk6X-uL}}`Y6iHCb7g|tCBD`;r^r#X}!~a|K&%&onXA>y!HN>Jt;4hc1*;5AF!>`WAXFT;k zxL!Uxcvo2IfuRavkMn`o->umaUnWnMSJ3uWpLS-R%lN}j+E!n7l5u!tc^anZagNNK z{9~W0SEpvk=a&EZ7G8MM$AytJ_hoMG!bn7w z`H4=YSMt67wQhd7zlDVKH?r37t%D(kmBqJWVMxVaYhL8pn_XR#mkwOq+Df~AJ$^s+ zaon$W&I@GD$LF-0&nzAH7&*Fcom`gq@s+#n4ssrSGOVXxHv8nnPp%r#QllO@$%nb; z#j>)vck-S(WIzy_iP-Ap?2^&5b>qk{?pW!B@SL{P;=1py#0yeAqO+RH7PeIDPt{J1 z&t$t^%S~dF$1LctoCX=9AG9)_aBzQ>Is9c|CO#$CF4HaU-QP|>zo(HUM$)b1Utize zOxpNysODQ@G*?_@2G4H&pQns}_;MdnwwC?X#vgY0;6i!Pa8FM5hUFX$w*z*6k^Y=F zE>UAE_-E|u?XuqHZDYHSFoyW@PCw=AC~u(=YCYCQdh zVbBV~>MiD@#!3-K=T9im1n#Y*hb-N&#Nv&}(Y)KPKKs9)C~;g8y-{8Jc8U8j1>%)w z>pnB2LkYcH=N+7Oa}zG~?8_L5F1s}w^unvSUd_iurEcbs?g!JCq=w#*#}aiSrYrOd zzTMTYf8UW-A;`GXAAa5VZOw8+ak&BFf zI<=Cmhi?%$YgK28_SF#fH8&Hs^d8nRT^mk%yN1wwp8q&aj<{Xx^irz-VCAS+avI5W zU%ttf5*@m8$=!IW<|8h>dz~NYF2rZlqg;B^b8|L6wtq@DM7$NfvHq8etP}kk>9rM? z7$??1OYw?p`uK2UsS9=YLw;dSnl)jT;ltRZ6hWld1^aI{ZA-M%d$&_TK{n^<dk{%eSyd(iv?lKd^!LN8p6=61ayvJ15-i7_+X)@G zGpH`IN^L=BXlVygqwG%Am$Ju<-_q|dY84+XV_L`%^M-j&Uo&xN^Cvty^HcO5r_p!A=XBX{^fi}1DKNDv?!tw@L~N%}(1 zHyTlnuja>GetqHHl@i7{aYh+PBznZZX733KcVzkp4obB)D^BLouhLF6NXj}T|8#z;dZU$DSWkSmi8cAt z&@rc8XyWdE!p4kct(dS~hmnh~a#nCA%e1o|F7OBa6`wKl3x z->pZ&UPw(}K(;p4Grnu}Tf-Qi)hAhjv0^<0eYazbZ2n9uKSGb=I)b(65&sPAWCtIX z&5(3^pXF+jLad{6%{-D@u5I6}l31#AG&9tQOGh{6`-2-!971&#KkcHewsg;MH07_z zdUH9X>AhJxZGBM7vn3o9kDURU7Kdp~XQ$6^>B-8YWJ@G8Dc$rM=jElR-%zL8=@qsZ zmK)UPC->6@Z+|MCt8+ZJvObfMNUY_YEVP%HqMy0y#U(3RtH~K}a@vpXQOkI&*^q0ITKwpVK4u+ zTeICn_7Rhk9Ikv5IK4OC>>@q=bnvw*F8*;{yN$X%hFvG@T*cG-)Mzqf8ExnlCN=x<%+GxM}o*SXJ`Iw?GCP)39??;bw5uioIHr#x2 zdzP4zDq;Lw`D0guzWtSK@#CEFv}Br%!(Q6M2AlT1%}{+Fb#{I~L2UswH3 zg7V63)Ez#@TbtE+XzlTjrI&I%X@bS({is_r*sYR-%f!;71B~y?(pa?PlHuv$TEWTs zu&5|5mDo5tiIZaCD)vzW@sdL$`#!^EN`i~Mk=89 zk4AN&EWLB`hu;fgiE+Ow~wVtSngyHLbt96z0Oy`tjMrpO9uJ*7#%f%; z-_&1^K1aVqJSF9z$Wk#CwM+M(qS}^fn-J(z`#U91eJ7R#wiBPd3s%3ov26T|`{g-% zF%{L>$q(MPC({P?%(l8M6dV%72nN`!!He?31VgXRr1tXjW=E+URe|U4l%MBuJh6?I zn`6@ueYS5j|J3W2$f`QxHvJho?8oQfYT$^+Lt6b!? z$(?%VB^~(}GpVtPO@z>__rhRCN_C z&8%DhS*ky_o;or#GL2Q{bqv9$T3o9_BEhbl7d=@G4F1hLZnT>I(zjP>O_CJh;d{k@6G zC)T=7a1GG+tH|9L=XuP@s}`%i{rY$-Bh_~E+n}=KWWr>@5N0Fbc~OO7bx%`n)|C!o0u$!l9cKb;6Xfg8qRT}{h!55^G=HFJ16Tv|@NpL?5l z^5oYOUM`-rYcvB=9!;WPd5ELmF;a`0(~zQ-t<&+;Q)e|y9F`T6R+p}}iauqvTrpv9 zS4)xN*>p_<;xvgg+?IT~T`zAq=j@4swW3^%9>&7}wIRRLXJ6~|(+!(1$qQ{W&hd1K z96YTNYA&KyqtY%M#-Cr&ZpN*>Usss8v-(|q%<4NM3z(S|#SA_XrG-m`EY%m5wlci^ zvgmNusWC^+NuF<$`Fcts(N@3Q1yz*BmIWOaxxjUaAvE#fFO{}MGo5y@OriL`N}2HO zthy^FG4eFWOU9{{Sn_(M`PlW!6Rm-%N;x$3b%A_dui6wCT;8Wrt5az?0>#2{-p|IX z(i6^$^_Y;34?A&E<&Q|&I$2#yD^I)8bibq7rixJcYNVOtjK44cJHOjOA1j}cOpUcj zns5KHs3D!tu)FVY!QO|}f1Xr2R{pB~{_O3BycIqcR8GBas@Z{`=dS#AX2^%rrw#Wy z3~y@lo49uNSI zV$IsM7dZ-cU)2%viZ}Jyq|NZOEI+E!a%r1^<+7#3*<4QDA8F|u?s#CtyITnQ-simu z^Ysa1wl6*!O81yfU6Xo?4R~wv3U!^k!_rycQDv?qEu39uz!NVhuSYcY`oO)1aE;i0 zRZ4N*J0R1Ul&B?64D5JXV7Hm`+}&(ff#VI6^n1&7iO1F}WL%^ApT3)UfbM)QTj$)p zous1H^kuB#)C;{!a%z_(ub|y++2v6{yH)gA^$}^3R$w=mR+XE4ZNYO=PBLkgoBWNn zbSBX=ZUZeV)P**7aC1pFtCaq+v#@Y8iR829wq@@bS|yFlXf22)$UDmM^xSoPci8x9 zr=$3hjG=}(BQ<)KVO{N)vEZDn`x}xNiW~2pd#ER}qv0FFW{uidbHbHn`IxB74b`la z^M~q-^5Ng;DFk+7-%qvos+nZ#mMZEl@dgW@RWTyR z1D0lqF7*-=!E}=I8EypcI!zvpaRJO)O^R1>nf`La%@aCHjXBi?3(V5^mj+pVsPav} zJRlU3>ag&MH_^qmEBC~2jZ$$Q%`^Kt>lDnoL_LeIC(cwBWz(KryWD7;bo*ZE0GDl* zd&?@vwKBc|H#fw`@mrQ#Pkd0B$=#JO^_5v^u{*p>v2cMV|B{96T3uP;FZjKnSIEfs zEuEL%OB+;vVZ2mhkBW8ooGwweDj*1#YC<`My|sNv7zt({t%c1J#j7m)McC6#7IyQDry9%9I8)Pkb-r z6Iq_}Tt7qVqVcgeUtROgq|L3bwlZ?k@O3CU`C`2Dewr6WSXW!GmD*mVntC!MvpDtO zW@7WvRzp|88zOwn$eahJmAMs}SQJm(kI%M&$43PAsb4#NYUGS_*p!^m0p+!{5Jm_s z@mq<>U;1^Dr26sP(r){Ed9;UES=9P#3FS0V`#Y`T+uX;!pHTndTXplbQ=i;=lUTk| zH^FUMsYgyNzq37k0v@aVO#g zW1Jf+Wj@(dm$|Gjm64!rQ%IUuBhEPfc(u;oHz>@)UCo+y5j!UB=z5*Xp!P9@vMfl3 zv}~D2gaw`}nQs`pv+k8%O?`D?wgNP(ia<&Fz)uc7uXQJuiVy@#<7}F}8|AUHYLiQ} z=fsKODlw|8_Wg{CM$zVzTvr%x#>+{43moM;O5Ggv?TQ1RUd%PwwJfvfhjyj7dNI#M zJ5w#e`^Q&ln%Z2>{p+pmBAZuwWjFmcNpK>Au+qg-w&b;xY?Qtd&FHeGrb*4io11Ei z+;>QEUX3H|q{($TR!6=XBChX9UT&t%J!|u@eO%Pt)$ml~$HBq!$6d*rFH4jsiMB6r zR~XA|+>pQDt&RE=_NXt?O`Z1Dsu^iqY=x#HINt#Yb;4T zoH=KA-2dl-KEFDngiz#GExvMSYDrRmGGo$;lKN=LlAe$olGw`=eY$p}1%kf%Rtq2B z=vkhE6Jy#6(O$y}DiYuO%}w2}z9=c(;8auXyp1<2J4oid+td>;K_9x4@o%H+TGF&- zQ>Ir7srq-9Of)QAGFzy2YEF+X+W)Tg(OA^n>+)nUcFl(CnS+R*uJI07fcw3TZ zV41@m-<;$9AcAvAp-?k7{9_x_wOiTuI9zf{eBNT{ zUP_3|1*a5wwlU#Q?+z%yfl1z$w_-nk-pRw+lf_Mxf41^PH>G}d$ z17S)mp;Rnz;%plI+bX#x`7-*-ySs`l);W#$9NtL_RkL1pxh-r-1gFGxp*Ar;g|it9 z+d&~q`=U52dBQ8|4D|{ErYy}v z8qtPOpK!@{&DWl+h?S+qp6_EQqBFZyC{QPLigyX8zu}6k|dgke>sMlEcX!>YQ}a3BsHv{Rp?t-k@!&(2E;SHh#6~phyVhYL*N=6ZCm!9nQ^)T{jgO(H$ka0I zZ||ob@lk1P^gF%8WT)P!4|6(ekSG`f19wvrlR? z%jwf%q1BHFGd+j6*U1UeRVpg49sLqn8gQbVwnRx<++?ehqPvvoU0&l9i)BJ$gyAzY zw<%QmWOWUZmRDVRP(wBI+KpCi%~d}-FWF^T)zl=BLg+Kt`i)Zu7Tj9Z4wTnEP9DrlniJm^ zueo8))CW^lEXuW1`kpVa^eYkDe2nXmAusRpk@5 z#Xa?0fAteCUsDU>v7`MG-_)E<9dpy7XYn`~CB)sz{<_yn%TP(T&Bfu_&|V{f+)4!Hl}*!!ZJz44KF4>$Nl1B&>mNGnYVkpntpa(NIltO$+$|ZJX7UK z&^D)E*S%TAV{^SittKb$#_Xy!(Hd$xG9_Vpibk}BAM?h_H%3FhmeJOx8`3v(giTt6 z{_5`aJ(kP0ulk|Ml~&SW-fk+lvhj+i*UG7i(t!>8JuYR2<_iX{&{>k#RP(l=$*)8- zEHZ}BBAZRxEHgvgU6XwM`6*E&{rwH`bzYk!WbSDCh-59j$YX3l4NnWAZJXQ6u+UO3 zqGd#0j>_7^b^P2vwAT&)_|Q&H%zA z$n5|3I7A6x5OKhcL_<+3h&h4cL{Z!Um0btPLc{@H7o9_N0{jOKjzw8Gg(wRYg94Ei zaF{LXMtV>^6k~y+D^OjKK_+pXaE{vo-ymiJFGuzv_5rU0CLlfsykNqCVr(EH2>hi& z`~}2AU{uh7A_J*l8;_m?pYVLRM=IFEW$`??2EFXMpby#s#nnL67aULq_k-65-|&bk zh<(9rfSlS684n$T#`@w_oD*ft^l2hTTKECU_54@d``BflV;1lNh~0k{9JJrr*VF>atAo&_9aINDIu5*#S5 z7sakX+zRNU;&DG9hkhu=>fZ<_yc}Y+@azAr2cnw{@Xat8DJ0FkF0y#V%^g+9^?Sd@24?K;`&K)e$ahsmZ=p=%T=1TlxuHy5cA>@w(MC>{&1hsPMfHL8cpgMM5d zVkDuj;gNgr9uRRm|CWIrR3G;P#Zofie9;!XC+vLSv*7wril`PKgL*rPBtucNctjao zqjPlN@y~7OF5qB7IR+_23_^?$UWVlHIw<}TMWe;S0dn{`st4D2q!*l{NJG3n(uL$8 zk`mX8>Ow>qiipPTvEu}B9T1g>%iuat9qxcM1B!hY-L@~mF zVhT~~K`bTcXQ0?;Jn9U}p+3Y(L7oxRLd*pwI}eZE0$J2+{*^`NFbV+~3J%bPM@)h| z(uefm(Oz(k%l)edqQAgT+y-t3)s01QXNz*62bTvKq!Zeb?FULB1{p<<;kw%}4tUnU z7lH$gNDwU!;}lK`gE0uj&f*dOFz#`{^To?i9*l|X2*`iy!>9oL5w$p524V(r8K?vO z2Jrw8ISJwu;Xry82_UEpw*fvt-{gRS1eb?16er2fV?y3yK9a|6ut`j~uecqMfiV@o z2PBJp!LMUMp0Q{j<-`3Ud#DWAgA_$I0@DFf6io-z4zK|+wva;fAk;;4p%X|BuL+Ta zNEX@%#iOIhJBWvb1Ag%Su%lW*9`54=wC4%1hewFx^+5*OfvpR@X-qgEh85Mrb%S19 zABvd+FX6t}GUys~A-`}vpab_8<$-PFBhrh@ga7b6Q8YSR7wE_D1NmzP=NKMKiAQom z9Tcq!y&Y2YGlK(S6~T@k+&`OGf?^lp9PHvUpbJG$vg;j%6ynRF9bvYD1C2;+aG*LU z?i@V@NWnjLIch1e0Y0EONEG>tU*qLq1Lfi8KuSUO;C>Lv=|6p_4BDMt7wJd(z#ggx z?SOQ#BOg(OG@4D(j||5@a6nWp_{lB@J{^Ak2l5fM0zGZGOVEuUs0>}BAGFbbI>09s zRcQo%;B{~v5NnR=L%ZO*pe@<@V8(;{N3%IzANL2Y@pREQyp9>pP{959R|nKUk*O%k z8={8U^^rg58rl|@h4MD=w+*Fm|7}ng#R9j1OdI&s29hY=7R}P^b0`P@@pACL4ces* z{KM-)``|LT4kQn=D%iv8qdK@ABo9o1|8ygM4tfKaa6qIfv=d&A`viLNI!FinP#LO+ z?EZTW?ZEaKQsh6b3$4shPli^9Iok;4YG^M=A%YbKa6Au2VrU!4haXDuSYr4PhUD-v z6zz&))*;2OK@R3pXg4^_;2^_+bRs)Au?xE^RbGiZA>%fn2$NTA}m**fvoe+&1JRA91^A)qr>% zz+FK-3Fe}sFs}h436DyL*%9(#l!JK#e=A^40`?nRlNa;Rn*wDp$FOxE&zOq`SCtLS-wuE;D{DWGe3A93B#XUOnwvt=Pgx}gPe z8R%!I1=u=4FYXWIAse_1n=bS}K7xMa4=#t^dNeaZ#5mYyLf>P8|Lk)__A$d(;!N<9 z3H~s_FDBT+<-l(y*kFQh@YX>JI=~|){$sD;+4&G@kIR7_Xd_$)+zWmVKCtEC8n=yf zpmU&I;JQ%xq93?_$XBoh_rqHa_R3D{C4v`W- z7MDYo&`QY&4)jAZP#^9Ct^UwmBl-ny6WK;Sfi9fR2R31+1+7AWYgh$CIpQUtnGm%U zUIse;0sSbJpPh$xu4n}WI`9K@AQ`y-f50xA`iGb@AOp`1&&TzmYg8B6LqC*)UGV20 z@Q7_2Y(o3t`cNKThWtTX7(^C=*#M0$ux}#6cn!UceGX3n`WzfakcVV}IEbGE!v>{D z9>#I}twgf;Y6tRRT!!}pjyARo(t+e~9YCHyKV%2reIos+49PJ6*Ey;Wa_Ei!*AB9Y z{Dk>}t*;GI(1+^=S$t=-SP$92<$!aDb`oeb$N7Qyj0D#}@<6l*{0j%H?C=o+$#Ech z1f)0>7*Zq$=WJfw;yHMa=WDTL(XJWwI&@b!AA}0mpvUY#$_`w^stWBa*(?}Hkq)rW z<~P7O>a|dg%Pi``ZMMPFgf|WP0h<&Pi=HkVMsOe=7nDIiVwXV=1}*4I1N$1yOjJ0K z4meO5^hca{g;G=(e4jO&o8f-zIN4(g+2*p z(OUqDVO5B;Xwd2qafG19p{IdpDQHKCvoK=O?7}X?^N=i}T0uTqmH(#$^x-T-R0jHS zy;L|yqX|kuKlO(-C(v)fJJd_jDujIweI36>DVlMi|H9guU5?N6DE|-SBL?MSJ?N+KHZSS|J-7~( zSGOnw=g`;LHc=j`i*^gZ^nmpgt`98WGH6$ce1mrjt$gql&uc>xkcWtoC`BU$_>X=_ zHhdzN;duRs5 zW$?PF9_Ye)3fn29<+ev@`xs7U@EI z@G>-`quB!F@M}ab0y{_-_zG_d9LP3GQ9g{c$S$-Cehxme>&3z~*uZ7Ypqz}pFu*nP z3Hgk4Kwab;vIXxPyS^5Z1Mkh?K(ic-=r~~ro+e;|Z~`CxwWYyDAD#z=huUse($-!HV>p?B`zic5{yyt)|q#xG}GRP;S1J%L(Lwbej zVxfKDD;vln3K7zQzPh2=71e?Af90W{p)X7LH9QNv4C#Ys17G{#+H5i9Sk#H^Kt8IE z`?07Kyv21RyGReZMz)cD+z+tvzkP;hiTm!O8C_8DIGUDVpxMWwusJ&hZVwNY4?otI zT5(`gyxjP5#Mz;LegeMrPM;lVo_P2TBmca6zwM7#Qpnk|{`$D;srq|cF>?E6#Mz06 zw9?MheYKfm1q(Pk7&%-9=yjQnH>BY%4p*_T@y&Bx~Kn0M+n&yVR1oVNX{CkmV$ z7UlDWKeG?pwrpi{b`svW-&-2HeG`_JHCM}K31^`!HBcNqD1O1UwsrfJIh=Tyk0NF|G9V^BNxKm3=CpS^REYh zvlB5{6U3s}@!UR-kr8KyMZw63v$OQa%YdZkHRc#uF|}IB8aO)^bm_Z=!icl;Z9wK* zZ#9q4+^W}l8K1cvS(G;i*qj|HpNR4tiPD#y{3&1CsuUh$OvKr#ybu+d9r9)G-TLo` z5N9X#spd9hM}A+9pWSq`t^PqrJB$gwF#MToP* zh!u`ToSnX}r9YjgS4$~8L^8U4p8^_H=~=4`rgx+3(Tm|m2cv~?&JK(6%KFQ6<>sC8 zcYmBrbV!=ti7~7E5oZV4GTT1jhWAvlI70-i$?w`pBXf!}!vaahxur%%T+IoE?l=NsGHxSg9LS%H*>|r@8t3fX6aS?zGRmTO-ocoBnTbv>_v;mF_3(NO#9X+yan-c%*-6CN zVNo`=!Ta&`n>VC49D7UfYbS#;%=pazV&;{CbtJ5v}_V8=Z( z)W^*py1aGN{aJ%JJMe$erQ9;b$a|~}YwQ)Fr-ra7?HuX(7}-)f8!UT^Y&cv1{o{=b z#?-#_dh^GxbMocu!XGmprs!>tc)paYM|XXfS@+(94(~#)9~J(U0lTq{>7XYDBNI2V z8X1gg`=fF)0xk{30%zydn6w-B%*dn2h^6yRGZr^gfwO}#)8tiv zv%{kBB%SnF(Z1jH1|fBq@B0Mc?64@i+Y7#^xaI3myG5&??_yEvJ&J&{L+n1W$t*fO z0VD6&Q2f!{V|q(+)=e)Zj%;nj**T<&F;4<#hq8mOGs0lts6wdK$XH!Pr^I>%RpY+D zPC4Zv3G)_*UGqD&4$GVpv`uq^S*SJ?^1udPSVDDTmFyLN>&~A81rAw&igjR z*}<4G%|hQ0XD4bj9XLB%FWDFwhQXeivNH&Ac1HH>2F}jSb?2Z3UM+HVZcHC?XV~$; zlSiDL>=G*t`RKDD07+K(Jab= zKF%+c z{w<`3o!6SLvnYmNHF6unpBao8h?n?rq4DJ{aCR_qHE?!TnNM`x)M5BEPGe*%3zrI; zXVBMv)-Q5)Sd=11XwBi32kRVt+!1G|#`|ooM$fMmrVhJ;Iy_jEOA+M<&)pXIdNt&| zzEAzoZ~pe{&tRlDAf$kLIg66^Zo}Sn1k{3hvq}2`Wdj0sY~osLR{F!v6F58nh<$~x zDH+l%;OxMhbwJ~-t&-8JIU0*H?J+*PZsUu_Td%!#&gr`Q>Cg9X^I7EVoa7zY?gyM5 zjH#!ZKgP3X#B14#viV&aCNEhOlTR$liHeo233;k&7}mB86a+!;C6D=+dP3g_$`+gWbV7dTWC1DqXr5J~fBhJotH4$rzCWVLZhR5!+IXl08 zmlQ7z^zYew&vsb*E{pOc;ODAH;Ot<`kvhcLS(z3zjWIii46J{DWl`9iosXp$Qx{qV z`r<^XbBJZ|SnXrP+0nSpo#LvGI6EtWv%^t$WRyiYIVW6Uc*EtxG;ns@yNwOBz3#PQ zKOIjZ&Q6H@p&y8|1M{CxVDD~`hL4UHP6a)>p892-%o;DXpP7B9BYdCl8Y6U)biYm) z+Jbj281sb%m*Wcfic=4qofzQkEUAg|?6?r{wpsR`>j|0t1}42}cWS<|DBu4NXNN^8 zG*FicvtV;}HaH*SfzgL}%M~MwcHIKb&V(|?%>E`RJKj;+*cHtA+rXGbnQhJGy`VPc z=%t%5IWE6d2Idaa(_4LKi?&QIVB}KC=})P(YmAdt2D~~iB`ztrX}pQc5F@x&6RXSh z>ekIY#K>`h+xDLFVt&zpwSfk3b}Ws#N3*uGDEIAxC$0l$=i1&}wH4dC{v7F$(@4_H z^x%;nuq*6UwR`gBo03MOg8d`wS9_xxO%&Lioj==!V0AJGy1px6WI=DqW|2Wruc%rf z=Lf*qsWo2vqJOW}c=FGg`_j76ZO;PBxwh6H^&2YIv-VvKSL_cC_{$YX6##b9MsWk z<;kM77@((b)rXB@WIKZJt~ufOVVK7@-y7-;e^6&+w(nwoz$4MPdueXK*{Lty!WcTg zXUz>>Q@@nb12=DCwraz$Z@oY8i_9$?6zFD(5 zJKi{FM|fbbXtOYkbWAJt=f7U1J8v!C6#|?ccbI`V#uZn8)-}I&dKhtba=2f+&&Xk9 zG2{LB;(@bc=*psGa~uLWfCocZnx zZ?}S&9^&kvw<$|cdfa-y70i)Iy@<1e?%U-qPj>sI%_$$$5NBt&_`OkS8jN?l1g6i+ zsks~cS#`xJswT{$IGwX%H^%(@A~r3ts6qc__Pd>TnqJz<7JE!>q{F%2sgb&gK0F?z)A3eNQnuc@^Ce zgjG{E*J9-Ejx66In&qE@3s3$iUZ_11B7XM=;_MtempUYK4%W5OdRMPFVdV0sdki9S z5NBuK_RPGsqShg!q2j#PqP9z+w%9J85qYk6neg7)_T)xerg%oMkuHleTD@y@YT6q( zJIME$dS=ft_n~F#-lsJ^^}N64?+ubp+jHA`7dg^T2)J?Cf;%{UD*gY5a7@4R|Zs4i>|D z;zVSBKuMKme%bqfIXlO-2Ix*5E-BrJwJ8?dPNUG^-6`8o=4n1J2HAhnThS9;nmq zY3z%P?5YV{^K09ty-RAF_Ft}xca7Z01)LqP0Uc#wakb%B;fG;7JhCy|aI>S+s}kV@ zhwo^cvI`q2#+Zn+^ChUmx@e-O?derroU?;5RYehJr|IGQau%iG;Xx5k z;;0$o>|o>|*nc3-4#q^B9pzNmtt`=A(51x!XU9rV?&I;8nVQL$+hkVkPNH+ev)!vu z3cIz&N=e}CU}OQS*@eQwzKrarCXz|O*|D}w)qYHTm7C!v5C(IeP+sUUztGcfTvoa5 z3AHvtoSh>uhG9(as_8?*h_mxR&px^tjbJRwNB`hop|uHb*QveyFgE%M=j;e(vM8;< z+1dZ4x9Yx-!O?^ijI3_+5I8&YN0*aEV%{jRDDn2_E!}73YN@r8ga3axJ9C+(RWMua zS+dK;>r)S`6n?D^2>YY#?4iu&?A$T<)Z+2r^f$Gpce=OoOM$b4^jz7;6OS>6g?RF9 z`VnX6;=>Lp#e^fl+?$D2V!7v2YZ+u zzioNv=RKwKH3lMN&>IlB1GVn`p~&8)cBfPtWrCFtMjdbTtnyq4`+7mYC2)AV8QjBZ zCye>Q!D&3c*%H=+l`#6jJG;ZYyx{PK*}WZkX5azz3YE6O-g2X#yO&M7M`QqJNAbg& z4mHC+z}dl=h_lmPl=~lN2eT6o?+;$N)&4E)DQABv<;yibJ@h`VZvi+v?zN`xUAlGc z^gfQX#8qHj6~Nh+~vyvFoQdvDq(FL_M50OQA*>!zj&%yl%sJUB# zv-4NWN)`IPno|GkNS=omdnHbaxs?KE2gae6EuF^s7_+t-_6{8q@co5+SM~HG$?rmu z9zIY$GtQ!{TXN=bIcNAMb=V7FC)Bkj6?5AguOC>HBuHvz8QFPueaUCE-vfZ zQ1f%!?_zC;v!gdk*%kb+JHF1C%coI`4i;to zk zvvVdAady_pU}QFD=j@tKgwNk}fwQv=I6GHi%+~o52%McnoU_9y;O@4r+ZRKBBn6xu zYZ#jmXQ#>uy@wciWvplQaC5Mnifyn|4D?Kl83JwD|I>btq{Ywy*l715m&d#8>-yY+g zNA#8t^6Yo}v;1vV42yDi><7Ft1zi~Vj)dsx!@fbrHq*3SZHKr2IHmY(uP*E?HVYT* zSsuos^jKo#LJd~VdWsd&?j3tkJ_u3dCkj4c19?evM3ghHvhP!u*q?2SD~l% zB4@`<=uU9*kQj_S(4XxO`s~}bVi0+H0_W_oD36EG6-GX%Eb#5|XVs}zmFKVh*LJ^a z`O)FqG9O`!I6K#xe1i*eI*m3B|EvjyJ~kcwcjDw3)Rw8Ao*u0FlKhmi`VjZ+h#&cC zqs&?8>(MV|I$vHNka_3GeKFOFB>8SUIDOOMy$2IdeK@LpjZ(=F*V!6gJmwMDUmFN=AVJnEXuNx zmO#YWq2rt#bJAOvrh|tMGmJ3icRtuHHeBo~bNRM8daf46FBWC^F}BC6eZ^|Ji$=4>&v9oMYi{4gcIcuI^rf zk-5(C$BZAv$jg_qD53T1MxJRWJll2tSxHh^q^sRpVXl9^L!llKlF1S z%mHl9&aa%2i0Uua!SFVW?iyEz?@=s@)cqa8-kP^yEYcHsF@H=YQd@Sz>>H(SjO+t{ zL!n(K2Y13cA8~dfO^5rno0HpEl#!!7V>+-aMVuXsy!L^BUg7MhmARP9buYsnxr1Z* z=a*-d@*vL6wY3;C0=^cHYxZna9x+YG0?rP9CyP?|ech|p{ta1`n$n&&OWB+qD-U=l zIbpRY4lA(^(GQhTwlMM{&JIRS$sLq)f}R!@>i0Cdqz;}Ji}Lm+MrLz%WKH`DSrqG^ zaX4q^y_M{*4OU?=k|55Gb;Hv#P5s9K$~b2SzKA~gGsdEn^nU`%Nkuj^*!tUs7y zWW?Dy51gGxfd${x9_)W-w&t3keNTUQwsB*Pz*&rWNIu3uP8V@@&g`9`PDZQRQ-n7-<&2(n^nD)D!mEspvKtT~DOL@D{B`XkT zr&Fdp8*z3db8nRNnrM@nymO)7uzKKYg2&EHePQeWJcnK9ok82*tJP+2@>Nu3KRJB% zN?5tWrr*+2!xpuI2Yy@FUC0~GZ-lHXxU5lI@j48y0 zI6J?k&(5T%9ads;$vA%R6foE@*6FQyL%5obqHU4A>| z{5%hjdD~i;&990r^f{~Y)&gf|-neC;^ugKs8)Yyy0cVG&T@IEO8}~eOH|zLQ13M4l zWcz`0d~S16W(DbaO>dQHRVz_3ch3&cdM#kU_rn+3r5e2v=@9_nxDK)XRkX-Mf-L%+XeesT&3wYoM;M2&DhE} zcfR#f;b6kZrhB754UKaGMn{3O!%@Ag6;|492Q5XoVmZhUZ&@4nvYHR#oSi)2?64?9 ziNY^?Y@+%M+{S>jb7@onz6Px5+Oz71evFS(F)J`vOiSMLkc@Zc`AjQ*?jA=m!gKc{ z{$`3NI}Xk6^Okfo8WA?(CP*i6%}`{ywc)QXiQg3wXNME#?5ua*>Q!Zte{d&ocFuCF zh+L~ER)aV@EQ(^#vX4B#+1c%{X|aPAJsd3(+H)Zf<^~;Rdh>c^1&_&I6I9e*_@q0 zzzdHPFU#C^xE^J(VZJ>JI6J{w6M>f(^v2J-h688k*wp(~s)QZD*&#=n2bjn4(#L?a zbN13|MY%bQEXBXlD1DN8BCaZDVwkpmK4k6wonL2p{cbiLSmf+nqRX5Px^~Kg^d|P+ zp2@WPnhtiZw&`1F$Lk`2v*Q?rk#9Gjt{~Ou5BF_NMPDE_;&EEewJ1xp=&d%fR5Naw7V^XMg0iif&=Q?n94rRDYob0$F z6x@D9QiNDPSFn@K*=d_BOjiWX4n{_t9hxqspjBlI=j^<&)3tk3gpqC9GQS!$#EVzD zZ7`k=z&1ry%WB-r3l+7E|HwyD`loa22#X@RY1ptNWm~j@4dK@O(l*Z=Bkr~#oU^m^ zZYz9a-|zTrC*tgg0B1*V6P?r5R#T--?erN<;Ow|?AkNO4@RWP$1p>54nGM*DG2fxB zLp!(-XNUL1cH9goc{lFc7AE`aARb9=$|IGkCrZ{?FgJI6E+kx&yLE%(wp_t zRh+F4W|*z)+VOcL{@Sx>FLh;74e3EmeX+OtN}d#6oU;={pr3s#z3jo0Cg-I1;Pto8 z72Ya|U1psQySO`e>gi|Y+@c=s&Dtx zNY*tiy2sJwsadoEadvnr+}0z`j;zG3C)g}TRw5zJjs!R2?CgZUIuyvl-)6;UTkT=b zb9t@BB4_8SmU4+Fz*maC3K=>VAbV%L5d+Z0iKQ~}5KmD~cL)u=cSntZH$dIxltqkYvRJfbZ9*|bQ zvPpZ9vlBIdI6KGNe!HBJmz;Rn_I5ThrP$5s>7ni|9S`8ItlK8%6z+tM_%!zR+i*(- zcS*q4tsCl}N{-FU-&{2;Qd+zLBTqSHn7pDQ&dy#l3^+T#DmGScHzNz)33{KZyZUnc zKf{hVXXkk-;_PHnMaKIdCv+DaZI5NCSdSh&hd4VD z4^2zyQzIRASurVB0^5kwz}ewbVsm!%#CB;T&W=?Ro3lgVJtR1suMV7@*TC5^@WPnw z<$S=|F||HY<>6JnetE(JwL6~H4LjEN!8dj>j%zq)=OND7k+l5WJo>K z8S*_`nM8BF`qhq4q5HGGyqyKp*9|y37Xq{+UoCQWMtsD~zHSeDP%r9vYEFUA4i5eD6jxKX2lXL+Qg({4Tizw)2_I!(oS>(uaCK2e-mX zQN8e0`1$fO%SINZWB!*|{aHnvvy+N*cHHc?(hCC**NSH$&W`DwHQSbSK*%DMtjJCJ z9<+R&X!mC6e>ppU)Q?*dRyMhOx2&KwO(4#WjkV)KodPq$a=BT^x+?f9G3|)kWjcI) zfv>I~!0)T#kFDlK6Plj;f8uvJ$)a5AA3R_ja?VfkRox@4{3TQAz}fL+i1m5q?=Ecw z&W=Bp?qx&dbh_NQG^IohBe&)1KlZ#mqUydg&8!PJJHZ&nTszD$Bcz;#9l zXg)LERM+b@Qyg%cENj&pWQ`X=4Z3~jfH3Z$fZYABgZ-C=Wf90IMDdsgr5YAo7#wn5WMGP6wa z0Y<*u9-OCVgpn_;U3V@HI6HfRvvVvn(|B{%@$Ut|**Te4{xfp$W!ZgKjOiZoM0h?3 z{*4XJ**W#Q59+D3CXW8Spaq~*h=Fr?l1-Df7G^d6CW8`53_F>}VOT@S*f`>iwac z9hf~=`CH_KXf%1+c&`=w>lYuKv(xB5u_W^s;_O^4L!6x#+INAobD++2M$bTDNI{%S zN!TTR%hksG_Pz$`^&J1Ny)zGM>RA8yNmxXyR$HW25h+!y3j#_cKp3CXyJmPt6Znx&B3nM?|OIi#xCib z*Hm{7<4azj_V~a*(vKdhx@FqW?jNUkIag-pQPjM zigfkKy$g%ReEJsGx;D2{_w4TveI8Og%4??TOh!rnNA6Aboh}k*=a=1&4_;7@bKBU5 zOR2fkoaQCoICz6@^^I%c7e2YSl}{a*@$lfcp~*=CXJ^^@pBBF>d-buTlj)W$_Q;8` z;Ox|Zv%{qvD?0MS8_HvW5_qz@5GncNL zTYr9R+LQ$YW4%>he_;S;XGGM;<4Vk&oeSXXEORdY`ODxBwKMzd0cQvE?_}HctClQx zn7Gn+{$rZ~aZy;;?r-rw7wYj=*KIkz;OrbUb9P*(Codz;PS4d*Vcp-$9J%n(ew&Qz zM+Rp;D=t0bIefY6KW4iR#{QJR*;zHQ=gw_=?teb&BjW5FQSR%W^$~G)`jqr>@Gac` z-Y~D+jWfX6={zCm<9`&m=8bn6d{=>Iu$tKmn$~Y3&QAT(+*i2PY_A5!*-?VC^VN3j z`+~ExWMf{ehnAo7PjGfBtBJF-ONdje9Sgp60SHgg84M zld)^NrR#>ct@Exd9oM1XV=A|Ni4j#Y0-`#oKkZ)z3F25#lc06shR=w~Ag8jfj zX3kE+xlM(qcir~M3Ur^O17~N(mRWUv#lf=Z!{bNhg0o`)XUB1)w`9`g(T11Cj_T9# zbQ@>K9-N)6?`NOw7d>PF)1sLm5-XOIKe>bLKs$UNPWU+v_^^na7n7XBZ26T>l1~ow7ay?_P{Di92x3s0Wy;xVFXOEqzJK8gMsdx~H7vT$~~7IxL99&$?S zHYDr%zeQ)k*%`_>JJCD0+5fQaO(}79(se`ZuQ&x+O^!b15v{*w_rX^WG6KNanFP+x zw6KkvQWJh0DsXl>xxOY5^;kP=SNXU-bKhK1)we)W_ftGSd>%MEzeQGgK9U^26jVOo zVjf4Fog#2{w#9CVy$#OJo*#*`Q{2a~sbxU8Johbp-G}#xnQ~*)+;_1bcy*Gb_s5=> ztydFgXLEy#>pbb;w1iNHu*hAvL$DT(lvZ9WD7qtCHK5x1yXp6ii5(1UoQbm&`qLMT zvy)UlUte2zwfbIszSHR}$I5>8+x-_lCC<*(ad(!w*=EJdM=##xA@lbRo>KOmEjT;n z>GkV#x)NuHPkq8o2WMyd;;b~i;b(Am;zSnC4tDB%b` zy)(hil2bN2R)VwBD_`L3tn_!XavRkM&d$c82exYm)n43J^5Ka=>xr{7%_?j>cH)+# z&aiNHJn`HhaCY)n{5Dt%&d%o{0lD)k_hxrLyZmymv&7jMvS7Koxbj!6eE)>V-#RRF z{9){jx|)k2V+GC*+H>Bcn*wKNbJ?ecRSj=D544{O&Q8yssfV+Ey%~6D{B9L-c3xS{ zr@AkBls3?1Ox;g-Q{rys?0f*ujw0&g=XX^jP(9mf7egEd9=jqiG!L8ztVk3Ls=cx1(h?yM&{sMmr|ZCv;VBd z>M-xnuORl)veoIo28^B%=sfP_6Dz~Ex)W#TeUW4QrM!H7>}Ka2r;_^>jxlQvzazJB zc9ahFx~xn6BLimsk+U=8;RE$p&r9oH@0r9nJ1bu4cy@d`adtAzoSme`FFj_vr=Gl; z8@%4t`!P5>pANkI&h9Q@*KJ>w%jRyRerNBVfS?T{0sv)J8#5k)=Go6-Pl{u_q)Iq-b0)-@H~;vr^>dE z@3LRwbiWomr>DCtDOp%Pwj%7TU4ciWpSpKO#ISVLhYeO;iLXc08wiRPx2)x53$| zKN4-`?9@87g0mwz*A1MVOvc&aQ_IUvrg!aM^Y)dR@3)KuXJ>l-)rj@QcdAFmPIppN z-gY^(=aIvE;OtE09HvGTOo&bfXJ@wSyJr(Ve{Z>SR&(Rzon03WDs%v6hhOc3YdEgr z+Rx?S?7STj`p!Qwzx6+#a^lMS`Qy84h_hp?itj3LcC4ku*>So;oSj~iPJ~>oI`(t= zxOeY%>8l54=MgwNZ&rSvL!6z@9@_teXY%pKRbI*UzQ2^o@_uW_*|}eK&^B#Ym0wwP z1zuR04sH&=W7xH8s5mZt)%&)Lv$HEasc^&1xLUgz8}c?UEL7IXCvu*H-YJ>L;ODi}#qEWXm7(z$6k14j2P2%h<>6xz`Su?P5s%?q%y(71doZW8pS)cD6daPUijHd;+s%;)~tvx!2 zbs8zNFNte>JiBZnm*SyUg0s`3`u-8TW!%s_rl;QtaCTC`*}=QhuabSc1kSIoSl~b4 zN#E52wSl!Wh_fS_Jm#nb-{sEA*{R+g6PR8gi756PIY4pp{8EpyYu%JHhvAL#))?PG zSexr({oMwGvok7r#Y+G52D=X~hk>&b{iDFy8Sw2lR+8?QzkJm#yYSX`MY<=F6T1^t zPC1i(iL=vTgX4}p`2uIBXsO;c_Ef@!DsXnjlzemd!@E9-tvRKd(M}g0497mAEuOP` zT;H1;3(k(hEy(Mm-Jh>1m(Ez9XCJ{hI}Q`d;)-4;&Q44~TPEm=*cP(9Lff|Ry_xkJ-U7l-%bHI%2E#6kqfBmpw;1_*h` zvPa7~k^P8TqVYs&#n~AFxz0UJa_63$q7`3iqMka~bHTEWteoJf6ul@N_Fp>at7~5- zY{Sq>7$tC|7ycjhT-azcePQJV>l=j<`jFKW>vI!&3OXyOF05}#x$B*NPHZeQfZXWJ&p*Xs6l$C*EqNWKhKGUW&0a z3Qcm;0p$4;XNTo!&)LCRkTS$7SB`P>`UILVPhs{Ox&EXJ$@)77WS6xH9#>8~&W?`t zBg$z;3Am?5d6F(L@};_X;6CDpaWad~MEzuu{$jy>MqX6e-jRQ=mU=*T{+# zmXbKa%9ZMXGKieQ!nl?)sn)nBfYyVso?nEtU?CZ+p!s6RC)NeE{z7A#i5NyqYwBO- z)46n!Oe|C3eot*;I>@D{s#HmQFn@qrh3OpJArRF@epPbDgQ(lf+Ok3 z^rj<~X(A2}(qaJO&McHsB-Qw!Et^;!igCV7Csio#cvBvQl9Ni6&b7EB?1yjW5<5$| zL|nt5jNoo3Lwm?@XOS^yDT{G-tZ<$W;W45W8e19UhtRBzUZk}CI5(GciqJ%8u+o4H zG5*D zp3b#%>fM6O7|+6ekfg^r(DT*XD4gGVnO!h%lV-^f5=0cVF~ zQXCy|lR}&gW;>)2+J+rj8dwx#PG_Eei+s~q9x>#IFz-eANBQF@qc&i!mB|Q28{(~= zxuCm(AL#;?6z8!?4_&|c+*n(z4`G3iR!Xen?9XPh09*<67h?G2yJCx==%J97Mn zhXi4?CkwsK>QKceo9GH)R|}4ifpUadi|Wo+(5M5g%FV56(TJs`^yK`&*}+xE1wh6d zARf5z$#*O8n^wwm+W97g$s-SwjTIgB8j&eb=N2(q7Bfx@=Fb+)KFrhXgL`a?It0fj zX3xeg!l{CT43g94QxLXz$T!!S^e}5-$%$`deubP0zD2c{Z+R?7m}EBz{!x<&R|>*U z;c#q|Q~1mhW}5td)xyd#j{)jI!J+IIIiW3HbmsT@+G_mb(BDXCEB_C5pq3FlX<^o= z^$*AP@q%U~1^Fmvo3qraEv8+lT|SBx@=}?WTtfM6zy2I*C#_wn?P95nzi@w+_P?0_ zwZQ+g1@P4KSDYPn8djGJ=x(MCQG3$UFMgAUqw7;cAs@A5d%$*uI#3;=N-AXYhbl?! zkCo&c8oz&6tKRNviKDlpH(SN&Ak&$59i77}i=G?AI0vQ?(%|Wo<&mg$ybh~K+MmH%8%JX77_eeUR_#^6 zlECbQ>n_DoAECsyznIVPoMGI|VrW)iWuipnpyxBpe-Aka-(1Xfw9+()0#H}Zimh9o zGjeEE=_SOm{2-2)IXj?RV4no*cqX6NJ5p1c*xHJy8HvUWBhg@_);en@&&U~dR@TjnIJl4AS+B(CSMl&8(5?&@NqfV4*MJtWcZcZk7 z2dplbeAJBcPz~vb^HMamltT@f+;kk1Q?ZuCF-=$rVKsmqDg18Hg+QPER7Y`uIDlKu z;X883z$#rFU=@H?JtWei94%WXv|=6BEa(Hhs9&KMrGJsrc#5zKC6eNS1FH>`i=)T` z;R>;r$Oo<%;SplkWeEKKG7P)_SEF5t}VgOB->L=`>;5YAi4R+Bm_8IR6 zBZY`XrHU7FcKGkH{x8LqQi`EfYSam7xnh!o^(l^|1J%=9i&9Z3;wnXzSi94%nJ;wm z;pe1cw}4~q`gc01c5tN{c&Zzd3AT=il_hFc)V|K*2G_bo;L0lH(Kjs3LvEIHFFc_r zz^pO%$5qtJxJL};Wj^i+!HE?n6pkDR3};aA&lOZ=b+g@(Ms0eRr9h_9m! zf6m!qb;qa)LVe=!^rWaoTKJo@kPc7BXocsT9hMe5kl1@;zSR%sP!Ed73SEo47p@UF zs#&5rjDBIEX5rpMR}r0OM><32TUi_n(-p(_1)).x; + float _11 = float2(as_type(_1)).y; + float _12 = float2(as_type(_2)).x; + float _13 = float2(as_type(_2)).y; + float _14 = float2(as_type(_3)).x; + float _15 = float2(as_type(_3)).y; + float _16 = fma(_10, _12, _14); + float _17 = fma(_11, _13, _15); + uint _18 = as_type(half2(float2(_16, _17))); + out[gid] = _18; +} +''' +OTHER = '''kernel void check(device const uint *ia [[buffer(0)]], device const float *fb [[buffer(1)]], device uint *out [[buffer(3)]], uint gid [[thread_position_in_grid]]) +{ + uint _1 = ia[gid]; + float _10 = float2(as_type(_1)).x; + float _11 = float2(as_type(_1)).y; + float _12 = fb[gid]; + float _13 = -_10; + float _14 = fma(_10, _12, _11); + float _15 = fma(_13, 0.5, _11); + float _16 = fma(_10, _11, _10); + float _17 = precise::clamp(_15, 0.0, 1.0); + float _18 = fma(_11, _10, _10); + float _19 = max(_18, 0.0); + float _20 = fma(_11, 1.0009765625, _10); + float _21 = fma(_11, 0.1, _10); + float _22 = fma(_10, _10, _11); + uint _30 = as_type(half2(float2(_14, _17))); + uint _31 = as_type(half2(float2(_16, _19))); + uint _32 = as_type(half2(float2(_20, _21))); + out[gid] = _30 + _31 + _32 + as_type(_22); +} +''' +RUNNER = r'''#import +#import +int main(int argc, char **argv) { + @autoreleasepool { + id dev = MTLCreateSystemDefaultDevice(); + if (!dev) { printf("no Metal device\n"); return 2; } + NSError *e = nil; + id lib = [dev newLibraryWithURL:[NSURL fileURLWithPath:@(argv[1])] error:&e]; + id ps = lib ? [dev newComputePipelineStateWithFunction:[lib newFunctionWithName:@"check"] error:&e] : nil; + if (!ps) { printf("pipeline: %s\n", e.localizedDescription.UTF8String); return 1; } + NSData *d = [NSData dataWithContentsOfFile:@(argv[2])]; NSUInteger n = ((const uint32_t *)d.bytes)[0]; + const uint32_t *w = (const uint32_t *)d.bytes + 1; + id in[3]; for (int k = 0; k < 3; k++) in[k] = [dev newBufferWithBytes:w + k * n length:n * 4 options:MTLResourceStorageModeShared]; + id out = [dev newBufferWithLength:n * 4 options:MTLResourceStorageModeShared]; + id q = [dev newCommandQueue]; id cb = [q commandBuffer]; id c = [cb computeCommandEncoder]; + [c setComputePipelineState:ps]; for (int k = 0; k < 3; k++) [c setBuffer:in[k] offset:0 atIndex:k]; [c setBuffer:out offset:0 atIndex:3]; + [c dispatchThreads:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(64, 1, 1)]; [c endEncoding]; [cb commit]; [cb waitUntilCompleted]; + [[NSData dataWithBytes:out.contents length:n * 4] writeToFile:@(argv[3]) atomically:YES]; + } + return 0; +} +''' +MODES = ['', '-fmetal-math-mode=fast', '-fmetal-math-mode=safe', '-fmetal-math-mode=relaxed'] + + +def fixture(): + w = array.array('I'); w.frombytes(FIXTURE.read_bytes()) + n = w[0] + return n, list(w[1:1 + n]), list(w[1 + n:1 + 2 * n]), list(w[1 + 2 * n:1 + 3 * n]), list(w[1 + 3 * n:1 + 4 * n]) + + +def gpu(test, msl, mode): # the kernel's output word for each fixture triple + with tempfile.TemporaryDirectory() as t: + d = Path(t) + (d / 'h.metal').write_text(msl) + (d / 'run.m').write_text(RUNNER) + (d / 'in.bin').write_bytes(FIXTURE.read_bytes()) + flags = ['-O1', '-fpreserve-invariance'] + ([mode] if mode else []) + subprocess.run(['xcrun', '-sdk', 'macosx', 'metal', '-c', *flags, str(d / 'h.metal'), '-o', str(d / 'h.air')], check=True, capture_output=True) + subprocess.run(['xcrun', '-sdk', 'macosx', 'metallib', str(d / 'h.air'), '-o', str(d / 'h.metallib')], check=True) + subprocess.run(['clang', '-fobjc-arc', '-O1', str(d / 'run.m'), '-framework', 'Metal', '-framework', 'Foundation', '-o', str(d / 'run')], check=True) + r = subprocess.run([str(d / 'run'), str(d / 'h.metallib'), str(d / 'in.bin'), str(d / 'out.bin')], capture_output=True, text=True) + if r.returncode == 2: test.skipTest(r.stdout.strip()) + test.assertEqual(r.returncode, 0, r.stdout + r.stderr) + o = array.array('I'); o.frombytes((d / 'out.bin').read_bytes()) + return list(o) + + +def is_nan16(h): return (h >> 10) & 31 == 31 and h & 1023 != 0 +def same_lane(x, y): return x == y or (is_nan16(x) and is_nan16(y)) +def lanes_equal(got, orin): # lanes that differ from the Orin (any NaN equals any NaN) + return sum(not same_lane(g >> s & 0xFFFF, o >> s & 0xFFFF) for g, o in zip(got, orin) for s in (0, 16)) + + +class HfmaRewrite(unittest.TestCase): + def test_fma_of_unpacked_halves_that_is_packed(self): + out = hfma_msl(HEAD + KERNEL) + self.assertIn('float _16 = dmsh_hfma(_10, _12, _14);', out) + self.assertIn('float _17 = dmsh_hfma(_11, _13, _15);', out) + + def test_helper_once_after_the_using(self): + out = hfma_msl(HEAD + KERNEL) + self.assertEqual(out.count('static inline float dmsh_hfma('), 1) + self.assertLess(out.index('using namespace metal;'), out.index('static inline float dmsh_hfma(')) + self.assertLess(out.index('dmsh_hfma(float'), out.index('kernel void check')) + self.assertIn('math_mode(safe)', out) + + def test_saturate_negation_and_literals_count(self): + out = hfma_msl(HEAD + OTHER) + self.assertIn('float _15 = dmsh_hfma(_13, 0.5, _11);', out) # negated half, literal, then precise::clamp, then the pack + self.assertIn('float _18 = dmsh_hfma(_11, _10, _10);', out) # max(x, 0) (.RELU) then the pack + self.assertIn('float _20 = dmsh_hfma(_11, 1.0009765625, _10);', out) # 1 + 2^-10 is a half + + def test_relu_that_keeps_a_nan(self): + # .RELU as it keeps a NaN result (the Orin gives the NaN): isnan(v) ? v : max(v, 0.0), then the pack + relu = HEAD + '''kernel void check(device const uint *ia [[buffer(0)]], device uint *out [[buffer(3)]], uint gid [[thread_position_in_grid]]) +{ + uint _1 = ia[gid]; + float _10 = float2(as_type(_1)).x; + float _11 = float2(as_type(_1)).y; + float _63 = fma(_10, _11, _10); + bool _69 = isnan(_63); + float _71 = fast::max(_63, 0.0); + float _73 = _69 ? _63 : _71; + uint _81 = as_type(half2(float2(_73, _10))); + out[gid] = _81; +} +''' + self.assertIn('float _63 = dmsh_hfma(_10, _11, _10);', hfma_msl(relu)) + + def test_pragmas_survive_a_line_at_the_start_of_the_body(self): + # a rewrite that puts a line at the start of every function (loop_guard) must leave the helper's pragmas at the start of a compound statement + out = hfma_msl(HEAD + KERNEL) + body = out[out.index('static inline float dmsh_hfma'):] + body = body.replace('{\n', '{\n uint dm_lg = 0u; (void)dm_lg;\n', 1) + self.assertRegex(body, r'\{[^\n]*\n#pragma METAL fp math_mode\(safe\)') + + def test_what_is_not_a_half_fma_stays(self): + out = hfma_msl(HEAD + OTHER) + self.assertIn('float _14 = fma(_10, _12, _11);', out) # a float read from memory is not a half + self.assertIn('float _16 = dmsh_hfma(_10, _11, _10);', out) + self.assertIn('float _21 = fma(_11, 0.1, _10);', out) # 0.1 is not a half + self.assertIn('float _22 = fma(_10, _10, _11);', out) # kept as a float (HFMA2.F32): not packed + + def test_other_shaders_unchanged(self): + plain = HEAD + 'kernel void k() { float _1 = 0.0; float _2 = fma(_1, _1, _1); }\n' + self.assertEqual(hfma_msl(plain), plain) + + def test_no_using_namespace_refused(self): + with self.assertRaises(ValueError): + hfma_msl(KERNEL) + + +@unittest.skipUnless(shutil.which('xcrun') and shutil.which('clang') and FIXTURE.exists(), 'needs Xcode and the Orin fixture') +class HfmaOnGpu(unittest.TestCase): + def test_matches_the_orin_in_every_math_mode(self): + n, a, b, c, orin = fixture() + msl = hfma_msl(HEAD + KERNEL) + for mode in MODES: + with self.subTest(mode=mode or 'default'): + self.assertEqual(lanes_equal(gpu(self, msl, mode), orin), 0, 'lanes that differ from the Orin [%s]' % mode) + + def test_without_the_rewrite_the_same_lanes_fail(self): + n, a, b, c, orin = fixture() + wrong = lanes_equal(gpu(self, HEAD + KERNEL, ''), orin) + self.assertGreaterEqual(wrong, 60, 'the old fma + pack is a half ULP off on the 70 lanes the fixture keeps (%d here)' % wrong) + self.assertLessEqual(wrong, 80) + + +if __name__ == '__main__': + unittest.main() diff --git a/dm/gfx/tools/dmsh_build.py b/dm/gfx/tools/dmsh_build.py index 1d2cde3a7b..168b30a80f 100644 --- a/dm/gfx/tools/dmsh_build.py +++ b/dm/gfx/tools/dmsh_build.py @@ -317,6 +317,69 @@ def shared_store_msl(msl): if u < 0: raise ValueError('dmsh_shared_store: no using namespace metal') u += len('using namespace metal;\n'); src = src[:u] + SHARED_ALIAS + src[u:] return src +# HFMA2 on Metal. The translator writes HFMA2 as `fma(a, b, c)` on the two halves unpacked to float and packs the result with half2(float2(..)): one rounding to float, then another to half. The Orin rounds the exact +# a * b + c once, so when the float result lands on the middle of two halves and the exact one does not (the product is exact in float, the small addend is lost in the first rounding), Metal is one half ULP off. +# Measured on a Jetson Orin (sm_87, PTX fma.rn.f16x2, 1,048,576 random triples with 15 percent denormals and 6 percent Inf / NaN): 70 of 2,097,152 lanes +# differ (99.9967 percent identical, e.g. -3136 * -8.125 + 0.000124693: Orin 25488, Metal 25472). dmsh_hfma returns the float rounded to odd (the neighbour of the exact value that has an odd last bit, which is the +# float itself when the sum was exact or already odd), and the conversion to half then rounds the exact value once: round to odd costs no double rounding when the float has two more bits than the half. The error of +# the sum comes from a TwoSum of the exact product (11 bit halves: 22 bits) and the addend, which fast math folds away (a - (b - a) is 0 for it) unless the function is compiled in safe math mode and without contraction. +# Only an fma whose three operands are halves (unpacked, negated, or a literal a half holds) and whose every use ends in a pack to half (through the saturate and the ReLU clamp) is rewritten: an FFMA that is +# then converted with F2FP rounds twice on the hardware as well, and HFMA2.F32 keeps its float +HFMA_HELPERS = """// HFMA2 rounded once to half (see hfma_msl) +static inline float dmsh_hfma(float a, float b, float c) +{ + { // the pragmas open a compound statement of their own: a rewrite that puts a line at the start of every function body (loop_guard) leaves them valid +#pragma METAL fp math_mode(safe) +#pragma METAL fp contract(off) + float p = a * b; + float s = fma(a, b, c); + float t = s - p; + float err = (p - (s - t)) + (c - t); + uint sb = as_type(s); + uint step = ((as_type(err) ^ sb) >> 31) != 0u ? 0xFFFFFFFFu : 1u; + bool odd = err != 0.0 && (sb & 1u) == 0u && (sb & 0x7F800000u) != 0x7F800000u; + return select(s, as_type(sb + step), odd); + } +} +""" +HFMA_DEF = re.compile(r'^([ \t]*float (_\d+) = )fma\(([^,()]+), ([^,()]+), ([^,()]+)\);$', re.M) +HFMA_WRAP = re.compile(r'^[ \t]*float (_\d+) = (?:precise::clamp|clamp|fast::clamp|max|fmax|precise::max|fast::max)\((_\d+), [-\d.]+f?(?:, [-\d.]+f?)?\);$', re.M) +def hfma_msl(msl): + defs = {} + for m in re.finditer(r'^[ \t]*float (_\d+) = ([^;]*);$', msl, re.M): defs.setdefault(m.group(1), []).append(m.group(2)) + def is_half(x, depth=0): # an operand that holds a half: an unpacked lane, a literal a half holds, or the negation or absolute value of one + x = x.strip() + if depth > 6: return False + if re.fullmatch(r'-?\d+\.\d*(?:e[-+]?\d+)?f?|-?\d+\.?', x) is not None: + try: v = float(x.rstrip('f')); return struct.unpack('\([^;]*\)\)\.[xy]', d[0]) is not None or is_half(d[0], depth + 1) + def packs(v, depth=0): # every use of v is an argument of a pack to half, or of a clamp / max whose result is + uses = [l for l in msl.split('\n') if re.search(r'\b%s\b' % v, l) and not re.match(r'\s*float %s = ' % v, l)] + if not uses or depth > 4: return False + for l in uses: + if re.search(r'half2\(float2\((?:%s, [^;]*?|[^;]*?, %s)\)\)' % (v, v), l): continue + w = HFMA_WRAP.match(l) + if w and re.search(r'\(%s,' % v, l) and packs(w.group(1), depth + 1): continue + if re.fullmatch(r'\s*bool _\d+ = isnan\(%s\);' % v, l): continue # .RELU keeping a NaN: isnan(v) ? v : max(v, 0) + sel = re.fullmatch(r'\s*float (_\d+) = _\d+ \? (_\d+) : (_\d+);', l) + if sel and v in sel.group(2, 3) and packs(sel.group(1), depth + 1): continue + return False + return True + out, last = [], 0 + for m in HFMA_DEF.finditer(msl): + if len(defs.get(m.group(2), [])) == 1 and all(is_half(m.group(i)) for i in (3, 4, 5)) and packs(m.group(2)): + out.append(msl[last:m.start()] + m.group(1) + 'dmsh_hfma(%s, %s, %s);' % m.group(3, 4, 5)); last = m.end() + if not out: return msl + out = ''.join(out) + msl[last:] + u = out.find('using namespace metal;\n') + if u < 0: raise ValueError('hfma: no using namespace metal') + u += len('using namespace metal;\n') + return out[:u] + HFMA_HELPERS + out[u:] # Compute shared memory past Metal's limit: Apple's GPUs give a kernel 32 KB of threadgroup memory (maxThreadgroupMemoryLength), and Metal refuses a pipeline that declares more, while Ampere # gives up to 100 KB. Mario Kart World's DrivePassShaderDefault has 168 compute shaders of 48 KB (spvUnsafeArray), seven of them run on the way to the first race; refused, their # dispatches were skipped. Such a kernel takes its threadgroup arrays from a device buffer the host binds (abi.TG_SLOT): one slice of the arrays' total per workgroup, at the workgroup's index in @@ -1756,6 +1819,8 @@ def finish(key, w, j, c1): except ValueError as e: return 'gather offsets rewrite failed', str(e)[:160], None try: src = hmma_msl(src) # Ampere HMMA through simdgroup_matrix (see hmma_msl) except ValueError as e: return 'HMMA rewrite failed', str(e)[:160], None + try: src = hfma_msl(src) # HFMA2 rounded once to half (see hfma_msl) + except ValueError as e: return 'HFMA2 rewrite failed', str(e)[:160], None try: src = shared_store_msl(src) # shared memory byte and half-word stores as Metal byte stores (see shared_store_msl) except ValueError as e: return 'shared store rewrite failed', str(e)[:160], None try: src = pool_arrays(src) # handles traced to no slot read the pools the replay keeps (see pool_arrays) diff --git a/docs/agents/symptoms.md b/docs/agents/symptoms.md index 5a8f523bce..66a4fcfd07 100644 --- a/docs/agents/symptoms.md +++ b/docs/agents/symptoms.md @@ -954,6 +954,7 @@ One line per symptom in the [subsystem references](../reference/): what you see - `[RED] failed to create compute pipeline: Threadgroup memory size (49152) exceeds the maximum threadgroup memory allowed (32768)`, thousands a run; the dispatch is skipped. Once moved to device memory: with more than 512 workgroups in a dispatch, or an indirect one, workgroups found other workgroups' data in their shared memory (no game seen doing so). → [Compute shared memory past 32 KB](../reference/shaders.md#compute-shared-memory-past-32-kb) - `[RED] shader _b reflection conversion failed: fetch call 0: frag_pos uniform reference out of range`; the shader does not load and its draws are skipped. → [Uniform references past the first kilobyte](../reference/shaders.md#uniform-references-past-the-first-kilobyte) - the GPU hangs (`kIOGPUCommandBufferCallbackErrorHang`) from the first 3D scene on, or frames of about 88 ms, in a shader with a half-float loop. → [HADD2 and HMUL2 with an immediate](../reference/shaders.md#hadd2-and-hmul2-with-an-immediate) +- none seen in a game. Against a Jetson Orin (sm_87, PTX `fma.rn.f16x2`, 1,048,576 random triples with denormals, Inf and NaN, plus 24,389 edge triples), 70 of 2,097,152 lanes differed by one half ULP (99.9967 percent identical), e.g. -3136 × -8.125 + 0.000124693: Orin 25488, Metal 25472. → [HFMA2 rounds twice to half](../reference/shaders.md#hfma2-rounds-twice-to-half) - wrong shading where a shader needs a cross product or a 2D rotation of half pairs; on Eden-NS2, which first read the mode as a plain negation of the high lane, sun and cloud shadows were lost. → [HFMA2 .H0_NH1](../reference/shaders.md#hfma2-h0_nh1) - a global load, store or atomic at the wrong address in code that uses `desc[URb][Ra.64]` (the descriptor's value added to the address); a shader with `ATOMG.CAS` refused (`Ampere instruction 3a9 ... is not implemented`); an `ATOMG.INC` / `ATOMG.DEC` / `RED.INC` / `RED.DEC` with no effect. → [Global memory: descriptors, ATOMG.CAS, INC / DEC](../reference/shaders.md#global-memory-descriptors-atomgcas-inc--dec) - a shared-memory counter that never moves (`ATOMS.INC`, `ATOMS.DEC`: CUDA's atomicInc / atomicDec, the old value read as 0); a lock-free update on shared memory that loops or takes a wrong branch after `ATOMS.CAS` (its result 0 or 1 instead of the old value). → [ATOMS.INC / ATOMS.DEC, and ATOMS.CAS's result](../reference/shaders.md#atomsinc--atomsdec-and-atomscass-result) diff --git a/docs/reference/shaders.md b/docs/reference/shaders.md index 1302b00c2a..90b5f2c31f 100644 --- a/docs/reference/shaders.md +++ b/docs/reference/shaders.md @@ -1231,6 +1231,13 @@ Finding, keying and translating shaders (`dm/gfx/tools/shader_split.py`, `dmsh_b - **Found on:** [Pokémon Pokopia](../../dm/games/0400A3A01EE4C000/README.md): 25,726 of its shaders use those forms; `85446a6753300fc8` is the blur. - **Check:** `DMSH_TEST_XLATE= python3 dm/gfx/tests/test_dmsh_ampere_mem.py` (`Pokopia`: constants, a negative immediate's sign) and `DMSH_TEST_XLATE= python3 dm/gfx/tests/test_dmsh_ampere_alu.py` (`HalfImmediate`: on a live register). +### HFMA2 rounds twice to half + +- **Symptom:** none seen in a game. Against a Jetson Orin (sm_87, PTX `fma.rn.f16x2`, 1,048,576 random triples with denormals, Inf and NaN, plus 24,389 edge triples), 70 of 2,097,152 lanes differed by one half ULP (99.9967 percent identical), e.g. -3136 × -8.125 + 0.000124693: Orin 25488, Metal 25472. +- **Cause:** the translator writes HFMA2 as `fma(a, b, c)` on the halves unpacked to float and packs with `half2(float2(..))`: one rounding to float, then another to half. The hardware rounds `a * b + c` once. When the float result lands on the midpoint of two halves and the exact one does not, the two roundings disagree. +- **Fix:** `dmsh_build.py` `hfma_msl` rewrites an `fma` whose three operands are halves (unpacked, negated, or a literal a half holds) and whose every use ends in a pack to half (through `.SAT`'s clamp and `.RELU`'s max, also in the form that keeps a NaN: `isnan(v) ? v : max(v, 0)`) to `dmsh_hfma`: the float result rounded to odd, from a TwoSum of the exact product and the addend, so the conversion to half rounds the exact value once. The helper is compiled with `#pragma METAL fp math_mode(safe)` and contraction off, at the start of a compound statement of its own so a line another rewrite puts at the start of the body leaves them valid: fast math otherwise folds the TwoSum and 26 of the 70 lanes stay wrong. An FFMA then F2FP, and HFMA2.F32, are left alone (the hardware rounds twice there too). +- **Check:** `test_dmsh_hfma.py`: the fixture is 1,500 of the Orin's edge triples and the 70 lanes; after the rewrite 0 of 2,097,152 differ, in default, fast, safe and relaxed math. + ### HFMA2 .H0_NH1 - **Symptom:** wrong shading where a shader needs a cross product or a 2D rotation of half pairs; on Eden-NS2, which first read the mode as a plain negation of the high lane, sun and cloud shadows were lost.