From e8112f4cd9d77ff03615c4fbdfac65bcb99e161a Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Wed, 4 Mar 2026 09:26:13 -0500 Subject: [PATCH 1/7] save Signed-off-by: Andrew Duffy --- index.ts | 12 ++++- proposed/0003-patches-format.md | 90 ++++++++++++++++++++++++++++++++ static/galp-fig1.png | Bin 0 -> 77142 bytes 3 files changed, 101 insertions(+), 1 deletion(-) create mode 100644 proposed/0003-patches-format.md create mode 100644 static/galp-fig1.png diff --git a/index.ts b/index.ts index 291dc87..d825752 100644 --- a/index.ts +++ b/index.ts @@ -442,7 +442,7 @@ async function getHighlighter(): Promise { if (!highlighter) { highlighter = await createHighlighter({ themes: ["github-light", "github-dark"], - langs: ["rust", "python", "markdown"], + langs: ["rust", "python", "markdown", "cpp", "c"], }); } return highlighter; @@ -554,6 +554,16 @@ async function build(liveReload: boolean = false): Promise { await Bun.write("dist/vortex_logo.svg", await logo.text()); } + // Copy all static assets to dist/static/ + await $`mkdir -p dist/static`.quiet(); + const staticGlob = new Bun.Glob("*"); + for await (const filename of staticGlob.scan("./static")) { + const src = Bun.file(`static/${filename}`); + const dest = `dist/static/${filename}`; + await Bun.write(dest, src); + console.log(`Copied static/${filename} -> ${dest}`); + } + // Generate index page const indexHTML = indexPage(rfcs, repoUrl, liveReload); await Bun.write("dist/index.html", indexHTML); diff --git a/proposed/0003-patches-format.md b/proposed/0003-patches-format.md new file mode 100644 index 0000000..624c768 --- /dev/null +++ b/proposed/0003-patches-format.md @@ -0,0 +1,90 @@ + +## Summary + +Make a backwards compatible change to the serialization format for `Patches` used by the FastLanes-derived encodings: + +* BitPacked +* Delta +* RLE +* ALP +* ALP-RD + + +## Motivation + +The existing patching mechanism is not data parallel. Rather, retrieving a single patch requires either a binary search or linear scan within a chunk of values. As part of the push to implement first-class CUDA support for Vortex, we need all encodings to have fully data parallel decoding operations. + + +## Background - CPU Patching + +The classic exception patching mechanism implemented by Vortex is optimized for super-scalar CPU execution: + +1. Iterate + +We can achieve speedup in the patching step by unrolling the loop by a certain amount, but roughly we benefit by having them all accessed together. + +## Background - GPU Execution + +GPU execution requires us to break down our decoding operation into a set of **thread blocks**, where each block has some number of threads. Every pack of 32 threads is a **warp**. Warp execution is what makes GPUs so different from CPUs. In CPU programming, when you have multiple threads they are all running independently. On a GPU, every thread within a warp executes **in lockstep**, executing the exact same instructions at the same time. All memory accesses made by a warp in a given cycle are coalesced by the **SM** that the warp is running on. There are many warps executing on a single SM at once, and several dozens or hundreds of SM. This is what GPUs do: they help you write code to exploit large amounts of High-Bandwidth Memory (HBM) in parallel across many tasks. + +In GPU land, all code that runs on devices is triggered by **kernels**. Launching a kernel involves copying all of the arguments from the host stackframe into GPU memory. It can take tens or sometimes hundreds of µseconds to launch a kernel. Any data pointed to by the kernel arguments must have been copied to the GPU before launch as well. + +## Background - G-ALP + +G-ALP was published in 2025, and its main contribution is to come up with a data-parallel layout for the exceptions for ALP decoding. + +![G-ALP Figure 1](../static/galp-fig1.png) + +The crux of the model is + +1. Split each sequence of 1024 values into a chunk +2. Reorient it into 32 lanes with 32 rows. Note that this aligns to the FastLanes lane count for 32-bit types. + + +Inside of our patching kernel, we can implement this instead + +## Changes to Vortex + +Vortex implements some of these instead + + +Let's look at the old kernel for unpacking a packed 3-bit representation into `u8`: + + +```c++ +__device__ void _bit_unpack_8_3bw_32t(const uint8_t *__restrict in, uint8_t *__restrict out, uint8_t reference, int thread_idx) { + __shared__ uint8_t shared_out[1024]; + _bit_unpack_8_3bw_lane(in, shared_out, reference, thread_idx * 4 + 0); + _bit_unpack_8_3bw_lane(in, shared_out, reference, thread_idx * 4 + 1); + _bit_unpack_8_3bw_lane(in, shared_out, reference, thread_idx * 4 + 2); + _bit_unpack_8_3bw_lane(in, shared_out, reference, thread_idx * 4 + 3); + for (int i = 0; i < 32; i++) { + auto idx = i * 32 + thread_idx; + out[idx] = shared_out[idx]; + } +} +``` + +Contrast that with the slightly updated new version: + +```c++ +__device__ void _bit_unpack_8_0bw_32t(const uint8_t *__restrict in, uint8_t *__restrict out, uint8_t reference, int thread_idx, GPUPatches& patches) { + __shared__ uint8_t shared_out[1024]; + _bit_unpack_8_0bw_lane(in, shared_out, reference, thread_idx * 4 + 0); + _bit_unpack_8_0bw_lane(in, shared_out, reference, thread_idx * 4 + 1); + _bit_unpack_8_0bw_lane(in, shared_out, reference, thread_idx * 4 + 2); + _bit_unpack_8_0bw_lane(in, shared_out, reference, thread_idx * 4 + 3); + __syncwarp(); + PatchesCursor cursor(patches, blockIdx.x, thread_idx, 32); + auto patch = cursor.next(); + for (int i = 0; i < 32; i++) { + auto idx = i * 32 + thread_idx; + if (idx == patch.index) { + out[idx] = patch.value; + patch = cursor.next(); + } else { + out[idx] = shared_out[idx]; + } + } +} +``` \ No newline at end of file diff --git a/static/galp-fig1.png b/static/galp-fig1.png new file mode 100644 index 0000000000000000000000000000000000000000..a658a733329daa28a8843b66ccb0ce0b23264e2e GIT binary patch literal 77142 zcmeFYbx>Sgw>}6732wmx2@*7T@Bobk*T!9f2X|;(gS%_vZoxeyXmEFTclSBG-}}k@ zrtUvebE|Gm-L5*lk8I!P?6ubOtQ90HBZ`VdfCK{rgDNg2Bo71g0t6hj2(N)V667M} zFfb_6CW3;p;(~%CvbI)6Cgz4PFk(S*D)6cbJ-AQ}`44aiLi`fj(il;1_$3hNz0MhE zJ|JR9gdz%bhSN3qnBifTY|{k#6w{J~pgLI41-CHbD#kk4Xzll=VDKlnZ@9)-9p2w; z&3g`VIvu9r!)$B06|&+@z>$EX8AA{+jVUPde~1miy!OP-c)9K7t=StCj0FD!rOuPF zyc||&M%z)=^r7mh6_-9!i~$BF$mg@xN4#}`dn_18F+WWoyw}I?%{L`(g>TBgtBSrR z^s5&MVo|DB3K~b~kp9L-zMa(Fw(GiP{O2 zFX+_7LYa*c>qI$O+G14+SL>qB1H}SsVEJt-tcSIOlGK7DA#psWWRZ5dgD_aD;-~Yf zix|XY>+pU8<_I?a5Qf4Bw2v56Li{DLHhv%~cJn`!HJ?J> zF`1JobbQpW@KU9WhPx|1uQ)|flkJauI0W5T?@w3IF)`{P==2T+Bg<_dxsorV-=gdP zz7={SbN?e;42sSBK13E9qac`&Dun!`f`(`+X7jGQ;IQ)`r$J8evcB~rhiXrXzbL38)S!`JVj`#BPkSP#3Mta7_OzJp&6 zhRN@(P<1D&HSs($dVmU_YMKrEA=CSAJl8h;FH|eAD>RuQ)^?5cV(h(rJsHTm@TeP9 z4}15^%CciPSKTd{T#8*u^|02zw3a`=pT`T|M^ZbCTOlZb#K1QKR|n-}oFL_5KV#PFdsg)SFX-xTIL_3C@ep?ur*1Ax2xYIwFTCJMWV| zv(Tn@deiy`cE%vQ@2*HqOI?7GE#oU1nk(M^?eI#s)^jEP{%rK9Lv%ZT#n$SRHRFL{l-@SAF3~SHw`FLvubP1;WXZoJCTI;a3tA3e|q7`4jQy z!}@1i+#MlH(yJ(*v4VssN0PVEmO2ge*lieW@NKXvydj~};%;L5R8eBLi)X^PO8!-}M z@5EfiM8y-t7INwg*eVfclHLo+e31Mjyl#7Hd5U>Taf)>cA57U9y(%7*gFtQfk@h3^ zNBUUgSj%4cBHW4OUFooFBLzGK+JcL5g>mk26;o4FBU2*N5mUVyVN)E_yYaz-1r^u) zvQo?ZK_#$~N&%zFQm%*8L*6YdL70*7S$6L5sY#znO*K}vYjqs~3aTh7)L&LVQzT9# zyNeRgmLu{mROamk0ZZ6JXjRx00Ws4dlWyX5VxaV7@_RPtu2Dl9)rJzJ zNpa=$!c3FE!Z!O9lYkv~7PBfy6=~J$s-{gvyIN=e8}&2xGhSq@kRIGKx?MU;HF>(3 zc(wRFda%aNippAv>f6bpp54M4DBDouaO2=@IO;&xlu@2^%>a%@7X2K1w`0tKF!%W2 z1Vdk0@1YNm>5&m*?UKO~XZWW<{Zs?Zi~2pZU9|TBF9Ns%?!u+It-DFPDZBC3=)3Ym zb=V%yuZLvHWEy0+Q!T;6PzJ&MBs0qUoYuqdvdE%D#1QQdhJhp<)Pv0lMeGb*Xk~KC?gTznVX~ z2}tl&37@ho&yOjg3ZgtnJRhFxPdRECXzB~th}o=~$nTpPPdL&zxcfCU58kdDJ3{P^ zj~SS0cV0PKXFTKEe)8YzNgLw9^?C&dAN$%K&J?~2UitMRGBa{7N)hrdvLh<=og1AW z#v9y9nsu(BK&s`5)7vi#f?EMl6b1AU)M6A@3`;x)WUHi58ML%!qti5mJ_KqrkBo28pcf6{wHBa>xQ zjwzW)Y#j7dMz3>j#J6HlqdR*BsB>cdV$-M}6kg<)q>QHRCaXugSQK51&nOJZuq4KF z9~|4?+CQGZH1VtMs9rE3zKOaw_3-3kbk}!D_o7v0gfN_AmEl|lHVYREgPEIIl#M|H z=YyItGcg+}O3(?~&0yM2#wx+eI)m1pp#GUSmdKRExk%AL$v$rKk@$5nnZZv`CRdUC zacyDJev1@iE~}H~yED$SI(mNG9L8@N-+ump=w4+l3$uk3hS=dgsjn+m zH*S|x1WBp~hXv!|@ZD2hr!>9^d!yEKVb$_;;gs+(?Yb#E7c0S1^|nD#)upVzu%vl^ zjrn(_-4giv=jxsJ27Vv6a((G^b9B9Y<>@TVk`h~SjWS-DN4?Q$a#u}BO|IFKI*0qh zjBUax{a)c7&R!Eu8*MU&s<~>zioT)p#stk%k$;%~y97(OD)IWAQKQy@3ysj)SPSP? zrRIyS+GUHd3FV26xu7}ZTE4}oA8Y66`^2VB1p9PzHzo-NDk{=a13l_jv$G5SO)BG= z4#AC3oE-OgXQs)`ptYr8SDt_K37?{9JOb|NI(h%YNw2>2`@QCi0wS zZ#lRO%E3~R5lP?wjN`c5&nVEE`gwrU{XFz4X|f7xjGm^SZpL-w__Jc6=3%+Ktd(fr zdhc8FvPb4kR(EIopm>@n!;LnV2j}uBqY6Z2P$#T@)|visHf*nOs>u?;3e=M1dbUyO zqd{!R*eF=9;G%M7dBfB0I?+1wI61sA>2U4w>){f<0^O1i8|3yV_PDL@bfQ1Ic1lGF zHcjW^vp7K9ZMZh?tvcKq-6D38^q9JJUqep8<9k@S5_Q#)*SV6M>GS4GXc_ga*<(y- zoxRPzJ!%Yutax{)hv(a3+aR)?9*Hi3NG4Q*BrtQln=c}BlNOGjb0-MtuXT% zK3BUcL1ut`0jDo6D2sv$18ezm?FhIIHdGZil9q;{0nQO%;9v=0UIAyYz`+aq?qBC3 zupePw{&oKa3`~Ft4BX#&WPtSf7YQ8CW&V-AjP{3t2d>@#2PEUgKXbnTWxV{?c`fh^ zMnFMOTpUOh^lc3dE$zNp+50Sj^?)0Q)?#XQFfh0j&j+lyJozzD|D=hcs=cbT6qml0 z1%s}Em7XC3#KQWy9vBdW3plkfwAUqpSeRSdaY1-U|H{DyoIl@YBqjMPi@g~Ssj9Rr ziJ+COAqhLfCx%a?yhtP@Bp_P@BQAL%k-v)rS3IQO?Cq_&7#W?MoEV%~7_4lK8JRda zIT=4OGcq&N13BpJoGtBjA@r7ZWdEq-pY;eC+UeVxSlgReS&}^0tE*?_V9!HJ`rOdJ ze*fsFA;je0Em_+AJuF~=jL%OPnHWAX{;O`FDCqevm#hiI&|FQ(!~&QzpbcJjPFB!g z`TrkJ{@vn#lvMq9Nj47V|19|*PyWx6N_K{}f>stlllHv-Uar3j|L4QM3xXJ*NB$o( z@sDZ#>nb3TFC^qU{Jh|VF%;p*eI+8vUu3-?jaKBB!@!_` zhmA7-5*_VJ-eb}u&iqk)6>Q3Wo;MP#Cgb2%la62?^%cwsutS|HwrIO9`^E(kmnR zXBHsm$_o46`JWS;2pBAAJqMFYqWf~!SgAJw1585%yTg*NImLb zZT?ji|2)|*KdM9NejI~;cJM9G^8D++dkZ8ahd}4Bt|&K=n4i1)j~-wVS-$+&1sTQ_ z(77O2EW20u{}|+RX+-q8e_vn*y21gYZV77e`FF8jegSj;tPrTTg9qr`>SBkcz`y59 zEXHpug6}aSC>?CL$MizvgY|@xN|d3R$j*8$^NBVeW5sIaTUbVn0+oE+TyXkeNLD?H zH-9oYNJlC~nAnpm8ued8@1TFd$nG8vwNhlX+)J3nk&N-nw;5BmUvR(|5`-|+oCM!q z7!4&!_bJomdmNpI-PXw8y%2AT`PP43G5*tnwl|&8%2bE9XzFgSytSa+OlxNZ38!#0 zJ1)zmR27@owqNXaErM?+SJ^_(=B7Z~u4lB9)%X##yHv-a<$Ay-*P?8S=dpb(cl&i2 z!&I{s6Ng%!yoQnZV(-GEjsu5QJ=9dO9P`2~_wp4bI7uphC^d%~t2^g1ckR-}B>EUM zRn2BNkR0aHJi~K{!lHLBmdta-eJcT}SUsq7d~n%aY~n^G98W)F!1Y95V#X0}b;-O1pBWEjobRlfU4JzmB{E$sB6Qdrjgv^A zA#|!OUiiTfN2?G&TWk5w$zgA*mH1$$E(I|nAl|TT?L(qAcBfxAGPb_LYb@yWGSQAw z(y>tR*Zo1`HZ1P_vDLVwzN7?E3p3PiF|7pAKSpV$Ax7}J9Pgq$Lh2Kfu6_4J^=Qxhh*~>OkqsEn)rm$u4P)!tOcuo&Fg+$5&%o zPJ|DacSn~@SVZrTSRCa;j#GJD<0nfsr043aGUu#@z~~^^V9O?YS49Ca3-qCQ`jqVi zok6*o%8hJwM(ja9QS{Zbt)a2i&ZV)(2iFax!K->wjZZh&ai?)Pj8Y9ad!1DpH1s6p zv6m>hQ4bR^KYSGL?IRia*#3N%8$U@a&@Be%*+qr7j~ zg6to8Ios~*YB`ei>J%!Oz`(QMY^+P8Ql^K1Dz@t7id|VZAI+9iB4fmgFhsI4lf-^T zTPTtjieK2?SNiQPT+5MwIq`T~l^VxD!S^_nklSar&I(61WiV|ii(_dYC-P9~iz()F z#Ut;m>~n0^9U>cuR*u}Vn5`C#yl~?+%vt{>*;eI#r))Y~?Y%8#ej!I;x7mMwP>1ZU9eA*#RTGhVwj9Fv6WhUk_IY0XsNDWEI=H62Mt}UY?B|cFlL6+t&?2!~>q;*$8lc)f z-%;Q_-PV*&yOoA%xWGg@1DarQzr{O+ayuQC3{40OQ%Wa`JD)8c#)SJuk@007E^vno zc2W?qnUYQvj*;+F<|2-egCdk5eq7aNf*eM@Vlr>ww33c<#`Wa57-#*A59h6f+uf6c zap`5SD5di6Z;miphII@pywEo%%W6AMGLVOO)lgEV0{DOZiW1rB?=I4j;B6cX4uA$W|kPmjJkQKB^RTouzxZQsud4g2tN>b|bDDb>1w zL2^*#5><&RW5(*SRPOy1n=H)wqm9VPrlZS@t3_9hpW`WPW)#o}r?J@@KKXo^gWUWw z!GNhkr=zQ!H&in3@fJJV&X=AZ6S~6Fq~d5*Vs=_wowCt+x5@fo1Om^ua8}f6&6$A3 zj5=)C(jLLw7s!=}!-1v|3Oszcz{f4jvtW!4hFC8;8#A;~DMm+!-==!$_S_xPN*(huXw+(N zQroh+U4Pa++HzZ8xwv_}b8c&QinQJA--yX~UB+24nJUX|*5NlD9&vaK#v1~^?c%mx zj^%Nit=TJPHS8rP;Bk3F2&zxl_Gl1JCIJn)6uMJ_kPrpyfcNlteSI+PJB@OQkd6l@ zBMdA=WchZ!YAf-&xw91ZCWyC+h=)qmr^J6S3-mgi2U=jjo-NHT^nw9+2X5u%?SdLx@7@`2BOn_s|B z`@$I$m)@o6?N3R(8noAC1Ad}vBljNa8v z#9>PN*1pyq5$$K(pE$HLW888#Y|QFkJ^urcT>ilVCcp&jO_k&PjzO2FJ$XpFT6+4J zG1*JO?fJ;eG`{>$XCm)={F(OR{>)*feGQJ>$Gsc@{1In3rVAM!U1jU&`b6LMBeHkv zT9^BWwW?3|h6F1%yZfz6g$nQ^X3&oa3ntviX%R189mXy)NRf(3%nCf+46!Lzu0PaD z4$Cp$mAjM`r-zpw2phuhI`SYq+pDh91ag{d+gvc`H9KdRU_!ollt zfAp-W~#QxJ0=Tte)~8Dt2Wx`ecVV_g^SB?GHlBZdD~M@gJN|s zR~PUaWt+;^ace+>a*a@qVBa6@1DF3JUq4ExtYHi#{kP!^-NzaA?hz8uT&+daNmybF zmr=hpejTUBJ$8{?c5Sicv!rEWj>024X9*z%1u`B-FnJXux3EEutzA~W<2?+9-$x>X z*WY#(`F(o>4biLkQKZzT5JlSg(x{ze5#EMRhBS5(0o)zV7<3)csV8x4+W2_8my3wG zoXWBlW=kyi3fOAm+BE>Fib)vG#O%I1UM}o)6@!_S3J85S64u=S>K(y=@MzdA8q^zf z-s>NUa`l=NGjqQ1=z&R)o9!fHC=Qc`!ohj$zL4*6MtjH=jK|x`-HBpGHp_X|0v9QW z&li2o*H}b+F(j8OJNdM7-Qik7{8q%Rnr{S5+qdIXW``p`kZs>=rNFO?CIh;*JePWa3(CDocl9z|AAFxdVxwtYW#zNY(n4^$XC87sPTSd>RjXVIyR@+EDMK=+oeaO` zdofvml&P_ntHrI7+Vf9eUr!M_VPl&$@1OUK@q7iAWtxwOs9)J90NzIJt%TA#mR~o1 z$A0sRS(X>xW|JjkLpGZ`hoZj6AJNfgDgX3yEcL$`n^=lup%N*V)ATnB+3sW^Fc>af$G>Ga$*Aiw};Kn9E*PYKG6s^}GF`N~IJ8 z`wjJ6+ns9TInQc}l(F#eDJJJjRxmV$!dW<@{0UvK+^{d6T*`v@^D)+wdFk3W3?k20 z$sAVXlc=|vXhRL17R)OrqW$ywlpv4Vn~NzOYL&8n0axPV!(V-Cq0?8iq*;QLAk5!* zV=*n1pz_v5s6RR}>B6J+)7_Kf!R&{;QoIweBTv^pet(wu&QP%fWc(suI)%h05lt_=am-JmZSDXij}9qd#$WaKb#8(-rgskhf>67} zS@bc=J9Z%kWjxbF7ztf=biK__cio;oFo#2cj!Cm;zy60&cj7}L1{U<|pZl0_@V-|E zTu772Vspi@k5m?gRCJEAfp(2)OdvPyC0IZY97R?KBL4je4cO*6Z4I)D%YUjv zP~272{E=d)JSp)yh3JN8%k9yIgAu_npdWS0F3XkU@&8jB+8U z3GRt@>H$OPw3|hyS<2}jH&e=kJ~zUycFxW@6b#{_=u@ft!5?59RGWy(1w*yW!xV{ovdZS&p}g&iC7>e=B(q`fsYv6dW4VMadfd6+ovI>b zs@EX-)kD8Zfyw264IYPy)-X31!B4nEzD%>hk7FgW z^m2P7TX@;oj>%y3`|ag{S-xoF8pVnl<%>V>SPMSjcHS$!A#&445`tM20+HoQ-%W0j zW%b6K9k&~9z<>69df;gAsW(*fnIA|PN)_7!bj07Li`S3umBYG9SW1i?ZqLOzHIOel zgL`Ytrs^h?mWl*lO#Yg!q7rA2DBvXuz7on&ICNeScxO4}n0O55{U$wXdn_M*rAdV3 zmFm^YCDt}YfbXw_D$gGHLQYk8#Y!uImG%UXe`aVY9W2}b%9(DEEIQ5M?bOFOIxo|1l`3~Emjbmn#n=3 zua>^OV>kbynev*YN+N-Q3M3!5QTNKKe#IZgg6|RHrOIc*ZZTi4?GX>FBWX2t$Tox< zeja?xf3oEHqy;w=fnD_~AL&broSR5N;_+p*njEVU zz(4kKjT)!?sMs9Px74-YK&{IXl9@-2k3X$=doqPRWHpT1mSh6PxfycK{jC&4K3t1# z$e;~N@|>)R9gB$PJ`u<{V#XvFZ`6im=#{Ms!h(MEdDoc*;0~DcAMpTX5`B5W^(rVL zM|+%GTx@GFNze~&KF1wuLUkM&uZyeuUttZB);0i6!3OjH8m#!=3Rqz!txs6RAp{&2 z!guKJC9KD6F8)RVSMj~y&s2@X2I;Rh1Q)8;u}VT19-NaTc;)GoeoKUf99)}l5KcSz zK?`&~urVYQO?|#N!&k2}U{NYQeY=p%W>%!lJ&SkOSNUz;uP65D^i|_oMBm6y&7y=> z=7+Fz&!=0wxE2WeL^ksxG1T7SaGurh^ja)<#f}gmN4)J!b8Iqp zWpb}T_(~+asH6KZ6w~foXDS=(SLKN7g2-&#xhaZ~d>+c-VS$I8Lp|a5oF^3_A4-c_ zp&79_;!t{P$pyw7T3h!MK4M`Qun^IG+F5c3MtpSQ_H zz;^^k_q!87=q6#Fzg-Mehs5J=UWSn$v50&_GYgThh@^kL(FC+(T~x3Ct>FKYIru4n z_Wm>fHLYw0%*LOnc80Hj%3yzUYDlB=vpLOirR|k^s|Sx^PxSX-JVwiW*NXOMa8;o- z^cyC$RHJTmXDq)^DCkWrwE|_69stwinU%N1OT^I_<^#G@D4@#_iC5kALViCDtyIkxF^SEB-d=T`P1+1}y7XVf$BrYy)zuAv3 zZXH=+xxm%<^mtbYRLp8T_+Dc&1b}X}GZJqNg+h+UOVnbF2L6)rxs`@}y}=Kcb5^mu z?k)hTj7P|2|EFvpko1MlXPe`>l6bpwb>(S9J7`!?G0f0~8!1Th|E(tg9;lz^+I9Ez z^ZR4-E{)6CP)uV!0ppNgHw5dIQe;52&zE+527pkofu)ESmxAa8c^FI?p@R*sJ9qbi zIZ>VHcGkK){yZY4aksy6vZ(#s0BhD8Kp4R3NPz(np>VC>2JGw9{tTN`l>tK$(9ri^U>F2y_iAen*U;{EEmh=-2(@^)J8fPrly@+Wkad()kJj zABh2qv=Z=ryog_31+TvX9zNvG7wZ}L55q&zmiQE?D74A+hce|$qrztw2VdARF% zximFAV^)^CDvUhVTK1@-IdOElvGmd8O1qIBtlVoY=(b?@X&hP#zx~=~Xn`Q$YQEx0 zQ<7t|UqFE}BBGcbzppn(tA@Fg9RaeCSjfH;Dh%_dhG4T$h8?i1O;v?zb=*p6a#J#;#s zs8-X)BRo5u3@(p#I$B8cN81HE;nBXab8l9L5mZOro5&Zrc;~ET-`}@4dVtz!yMctd z&quP^2r2zFzN28W?|1U|?wvN(*Uay-ZY?8)TA<|+g>lT%VQ z({b_)=Ms&&WWXkn)~K`0dd75{ovre_iHg50G+yMAOko)3i_k)?h=9^UW?B8+z;1lwrw7V^UmA9;feDO7nX8C&ZoD3 z^47hkRuIu^Eaqy5_sZ*#Y22VWi`QV`p+dqI>9SY#ADjy7^`wZG*-hcUd zW4GxhsNWr)Gf}J}c2DeH&0J$X6OV+4rb*|}5tzwpGWuCUVpcqXK`Seof;huyK)sPy z!@c~?V$_r${E^q=!`-ky>R!j$DICxJG1-?Xymgi_oouzLypiJ-;<3TBxt-GuNyQ(_ zY3KbcL!)Np4i|XRBTVJ(F{IO9ps~k(`66L7&t$>IfO@UDGne}nsK-H!b%r$>p?KWR^&OJlfctDp3;#^IW`DTYZQ43%29iMz3MD z$Ada`)z{xh0C#j4t1Joa_M-w_Zw@_$L)|aMD~fDj8xLytVX1DgpnVC93Jtbbl5J={ zG4yj5mC}Q0cco(GVoPpk?)>S&JRNho^aL=El3>|S9@B9uxVEq@VyN$~O4TOT`NP%OdKKJIEP!>3 zhTyt0rdObsTzfR`4qn&Mb&ffNF7Po4- zs4X?#8Y~~qPe;p2<+KahGJGXLg;VOIcyKBaiemv znqT|GYE|g@f?Xue@4hZu7k}B<#Z0Jjy&96wkES@fqTXBY>6?7G;v+v33BT_jq?OCc zp|Dg(c5vEVsL;ow)+koiO1ICKs=wo-Q#Xq%JNB!wIh$mi+H8}1>u_gq&6nsO64E-q zJrYB{*&YyqjC=-69N9*vWQVBsIlJMl&#ON7(T8%pL7NoEv{EB_ZXwOV#|Nrh8HZig z(1MbgD!g9X+biPta-0Qs1xz}B9G!tz(z-u2&%kkimVS@jc(F3j<{7=G;1Y``DdIIO z)oPNS{SryQ?Y=+V!;ES%HuDo;q?#Sp!Z_5k&emfNV@GDn?-ebIhg_rTU^UMcU5{e8 zp3cW)H0t!b8C2Ws^tGH8941%Vw+wD~Cp2<2zavG=*nTFyd4+@>?aX52wiT1FKt(Zk zQOc{F<}W>}=*ps{cD_9}KSbMhHpqRX0f0MI>=v_F!=Lvi%;?@7PU$d4I@cqT%+(zQ zz*dgsd(s;6-I<&6Lrt6if^r}q_;sEVNY3AK+WL&+8NBZ9j?js%w3vnaISgoU^oHM? zxKBvOL zWurH4?hUjouvhj1AnUZ7%x6vwO`1Lkioh-Z1Q@NoKT$gS=Kv2O#;d@OhCOI{w_kNZ zZX@3vWe%WAda8<(rA7l>u(y^ol|yardFQDh=lGJj5;eO6);NH8duhmd^mC4n1l;2X zfb(lSY~t4+O1}=48x0VQR-43_+?=KtKu%nJf_l!sBjW{D%=N;lG?U3zg!}_r=WO{YP;A;DRjoRX!HH`8=w#RNQ_42aCPtp@#Wh&7PmiIE=Gu< zJVhP5e247z+L#h9SM`HRue9uEcyh!Z9}Gwb_?tM6M&BiFRx=@ z3UR4*#BRhhsRwfYiKpj|W7J{NOQ8+})=e^=F1A813!aV-CUmeZiY(xx>t*WsX&+(C zfhNgd7C^U*SDQ%QZfAqzTre$?j>*xNfaU3^BE?$paOG96-KWC*$Sc=%t zLs)>;!z($iblBP&xc?UI{dJIp#K|P^o^NNPKQpe;$AQ5l>lVyqSQhkFi*AbfGA}Wk zHQCh&hoC%y5q#RpT_0CwlGWTgd7_B{y?oe?F9C7r$r_(|U2r7(F`hxYnZy_ybSfr% zd%YI5n(0=|ZTlqo0Kn>{ia&k!IQcRYzFWYJouELr*`Bq z9RU=6BzCJse!y!(esCDc;bl6foqF>(Xe5*b^IK7X8^#X6V)V0mCA@B>*I*A$;^Qo^ z+o}_TBZdFwonTsHinCGTn=OhyX4U#YX@oL?HRj8Q5;-9a0AoaYN3Boc$8?V*&7n$# zfG9bL0c3y>9o4Wh5fuetTvf%V)bG0V&7e;)Xe_7W=qyM~RkCoh^0^X9inhVE03`Pc zc~ipbqfti7-68yGC*t>b1}!1RmT@VD01Wfj=y^}P06oZ5+7S1V%=XCZg%4uH7t(8G zR13i@<00PBv4Whd!v(THGXc&syP3)h8qTjFYQBjp6nY4+m^_(D_$Uy8CPaxEzz`>o zH&>c$Ox0`ym*Z>B2v&S~6XiOs3WL42yPVFa#$P;2q-N`FT2z0oIqr_EX@PgU{knsy zIqm=^(0abrGab^SlD*dL8P9`>f^+>ej0}0aX|(+VgAP;f0q}@23&Y%IQ)RED(;vS# zR}`z1Z}g*0t1yx3Q6?w)b5a7_us31^n)+eoLX#7&Am>Bw)a%8Ry2V{`mssLqfKuBM z47aWgPQd(~*Otxkr{&it1>jrFWPp#pIv}QfULaR*_R~c9Qz}?H2Y!SSGz<>w4tN(B`hfw`+7T;X!QOd6T}Fxp z#TGP9{|xAtOdZpTF58Q4XGaMLxYXDzO$E+pleO<170c&C*s=g#`|F3HFW+sWRoMkP zWJ}Hqk5rrTBC3aRwT=m$oe%x4%samM(nbsV`>Bh|`=RPouYDFNr0;_ZO&+YvmO; zFadA&P$;-Xjywss01qrG0_wSL)84lpuMoR;7xyMhGlOx6-)J|fbkfW`xFhjce{Bt< z_`?`L-W_@~PX)drVM9PyW-13956)reH0ygsJBd9tKfL^+QDZ8q*IV_x2G6|K56gvJ z(!rB=Q{fTG8{~QUA$T4(YRiBpS57|2_457sM#332pp>p?i3tL|_Qi=Au}Ne#@kb~s zU&)5W&9WHH`B{I&i&2qGxUs*$;~L$uJIWX`olD1C`PqvKB@+XpEpiH|35YHNgBKoo zuIG~Hri9Sb4dOU6O1GPhW-m6ZAK`vp{cjj?xxuA^-fZvMv?K*o1_A!EBi634^(iZO zSd<={QnRm;2=Mb#>kWpFGiwoQ{@&l6Wz%?J(|g7RS}AoReGx;?ME6y9goVmu8jtp? zq{io|ZOke}tqtMSkXW3SQ7_y3KIy~h=W4eYa%7P49^Ymt>`fjRLa(I+_=zU_+a3bh za-zqSIMM<8s$5GijcNLqMHqfbh#!c~9O=HI7H{^)HAwo`m(s>dM2e#cE|6*!(kq{5^nHnXs!U%k(!jCOOSJX58g=XG>2@belT=_cQ~vZFa}Y_>%pfVz@$J z3>$oyabLF!O1j}m=7ngtx+?^$Au=DFbQh}^z-!mG@mT1x$<5EJ1~SDiuB;1(9G~yc zDC~YV_=Azoo%8j#5Ay>+wK?9-)|eTp$I&S5vz9wHtDC6NYg10XtkA(#oj&-Erybkq zTflwdPpF$s1%*g3umz)ULiXIam)iABoV5c0%hv7J|z?4twR5^hNI`&U9LdF~)i`dcM~a}F2dIvHGVOx}a@hth zfaT~mO2%_+n`^iniW5CD=||PkR@76hY*E?40dJU`SZF%;+5pB@uMl?1m;#zUS4;R~ zL%~!K5N%~oqmqOkKRTI73c1&R~LDKv2n33~MiAI?N65q&NUkthIJlu12w1 zr|#jb2&kzV;`O!M?FH$dgnc5`8X?1~($O?LspPa-jgD}W*zNJ7YKF4@08IA!S3_Z~ z2L!B;MPp$NLwKfUiXcA)Ub{`YN_W~MF+g4sZNre#PqHBvyaR=sO`1Dc7zamkCp<*t z@DAD~9s4aYRLIv}XIOh3bd=DFdtsODR#6UKZnrX0y4KIC$6*dmO1AmRKuITMv`#bw zROSoOxkyxZ=BNYQnY5PL@^w=ysTTlmRWVgkY4Pp1L?iqM#6Gr6r5II-7_8KqDUL*r z81@0>vH`b%*Oqn%F$hC~df>?v3L95%s((`eiB&}fD*Jrl#E z^>wkOy%qZ1LUaoZy=mNa^-Imj`wPz`L`Z%9-&oeco^YdYA`RVm0G}|P6s=*bq+_`6lL|z3G#sBu}7W|~p z@R1T!h*B?_0Qj?DMlyZ$M+BqrgWD4f0QT?Xf60#jUcv_?{RRqnLyr)9-V-Z9ax~|Y ztQNp~CdQ)-fLpu3{KUcb5Z_o3iDE$MT-`|5KuUp)FV}gx=hNI+rM}km4vX>5-i+N&6tf|< z)1YRBa`)LDIE^XnbiGFqeeud`%dorgmmzJ!nCH_SymZoM%E_XRfr~M^l^M@v+4;BL zv~B+kh`1(;By(k_1U+U5*_!A6{tLSqw7Qh?O{OY8UCtycL5eG;VoZEZpgmJNO3Yv-R*}BrB`6%`b1Ebc@)MjgfnCS|$x{|r9+`h5I z&6Prpjh$KLCs@HYEE#?I1=ov;T~_nxu4;A-EoKcKrvt#*c}>rS28(En9$S--67;)* zi4&mjdL|J|-(f+AuXD%~0GLTOWR-J&C5W_Mjv<)`B=`EWL}yh1mBg>SFBSVM;;$mt zw%@y=mLa+-cV=#Rra+ILg_<{;Jl%QlUiu%uhk39j*f8}WZ zzZGO7{{LkTY*@XBcwAIpzDBXmQJuWK*rU$@{BEvzchO|CIhi(9-GXAzq zyJeQj{doB|^TX*zpX2RFwy5E(A0pPXI+Y~=K_lcMzj(Mm8*(29A|-02mMRTHz%H#G z4<*_yE;l*OHpi#5X{ERpToAcfKsEYUc$dUtyni~F#!WRQ68cs*U)OkNG?&?ZqDYB% z2iSQkq+0EFq5%LE4RDD8aD-@=Qnjis(KrF$*FE;et5Ia42D3ThF{wS10F-NR3pmia zk1E*y(4Vq&9->!o1X1L}^B5n@*2s4QV9f2Gg022Sh5UW{i&T7e^Tf2x2^m%d^jSmn zlYU}O+nk8xUCuKg;^#I89iLgKmph$>_XRa*d$Bk|DPK1R8#<8R@Zp!=mS@58Z*d{z z!}$hx>o3`&x9_Z)S7Nj)^n{ZRI|X~xtN{TBc%|k11Y!xQ8J(7Y!6qqJ8i;eFa=YK2 zN9_=MTxWt^0P~L&1hBWUx2wO`?s4c;W%IaoN+7cBP^BV8n_!npV$`-Oy-uXUqcJew zLAE|X@EYE(_r$o5RvGl@X}Moa@S149hloP}c9+5dEFn>;1Vp5yOSvw4!|${1Hv?!X zv-|QdZ{8hX-0?e9VCcCxxW_xe51vt?x;#@b;I_cFo6zZxee?YUIvimoFy9K|wsSw- zU6XiiR)x~fW8VWTvD;}6m3;T%LepJAE|i7aasS5{IuYoj;Ze(#Huper8X(#Mvs}`h z*k$_zG22L#R|~&(S}XTavqBy>cg)~@>>e4RjCQjFfV zo(Eu3C#H}{Q$Mr8c1-`fOa8Y?ifg5U*6Vfmju81%* z^r37V8u?sN+8bpMTDV9!J~a=BvBne&zUHK~Ejl;EWboeE=3vF-AAy1$b*nbk<97>- ziY4EeisV{E-j8?HA;siN<&7Cnx37gr7>NEKF@lw|?;w|fBV;C*^GDmoCRHfo49szp zWjf#mSEu|;-`lZ;&efw8jXXUiQ`K@_wl|CkkUNK+%hy^Bw%=c05@P_od*N9A=rrP* z$>TnLEU#8$ebLfj4$pYA68G+8E^Qox5?9%-fA6@ZYC>-uwWNhc9pGYcv9}a)xj#2R zM1)S?_pFq+d<;g6!28&hhe8Re_{}p~47iYE|DH|*g038=^E+t)O{CxSf3{aeq#r;4 z9-QADwT|&z&A*Gksk4|PRd)OdJ>6{KQM&=4Dv22&VoOYUQkq1*&2IUZ2f)6P195Vg z&rII}!%B?f?zs5v%~rNOz&*Q-X50K48}-e0*%_snZF1t!vKwG)bNud${KyDzaSeoB zWl3*lYgqOqm$>D;h?0?9YH>{%zIw6XdgR_yWP5%5fy?ch1n?sY0d(Dz=R80N0$Jcz zd!`_~;k&*xZm_jL)N}N}2oOQI^YnQ1RG?bXB-l%eb)Y*{b`Zax-{#vBh)qB(T>V7|CQ&-&2DeH0vXm22;ZxTW7{j`J^gW(C0zW{;Njrp@y>VM8KB-| zD=QzjYE7nQOeZUSFP#DW9U4(q&q3sVZe!N&pW{_ISvqIA;I!n+dpc9OV1!D{%fLVR z#Op{DGBgPH=WeRo8D!dzSE;4l=*Fq`E}zYtu*tEnaUiKq3ONRdv{MIs+D|Gv2RN4W zT24eFyliEl&E`NK2St@JUK9iD-Xy@YzFapR(*A}+W1kQKg3=@7(xFE>%=TEco+{Lo z@*bg1asjbNA(J7DwI>a{=bcp#%fDG*4c+fg#E`FYs=|1#m)j?!|1b95GAydMfBO~b zE&%}mu;G0Jk*N?nj2ho->5V=D=!so!^1+`7X zvu_+gY?im3JtM&uI2Ks+u#M)p%aM0K?l%xE>*l+RB1qOQU+=N>HEFoA@wU7?u`R4# zaGO;GLTf+-UJm}Y@8uFIRuagMv5)T8q^Cx(o$QV07xxiE=xTvz>7|q%fXH!H%^e}0 zsIfEDN`IwL0rF|j?7X(<4gI+6SIxMoIAYg&y~`96Yo1vLpM|13=b$4lPlXFWn@~(n z(Fxc{a#@I-luT&BzE5^m7=CR*UK$-pr4}Mv>(PjEYLUZqvKn5d*Y3DF@HXi^Shy*d zZ*BNWR?@?MYxHjxz@6dw{=rmP_q=Tvu<6X&N&W7L=ZS6mdtc~tZk021f?Z>pgLBJZ z6O!*r;57Pq3al7C6~IeyHTAIV$si<>vBC-JKD2L#6Yb0Zb@96P@|GFpna9sx%HEa&!Ot^qg>+Y;UrNjCl4l=B*shJ8rn0w_*o3 zA1wwx1uX>=l;XDG!^z7?0ow)D-)W#isF(%sP!=GqdLdxgC|m&A=j@hb(AUT^;0Y`p z9TSA*r0*i5E26J5liyi`vp8Hn@7np>gov_HR3dgW7PMcD(pt(~uJ=0Ku4Jo+B{e)D z;IC&oJZh1|KQ#isNbfdLGO+A$9HMO5%^>l!sAk0VIh;ifvvx>MFo)hmvw}`9OFa zqqJRW?>Vb%S|vhS0(yGNfw!@UKmK7mp`U_$c6=Vt4sLn2m5+$wx}EM$%VeGJ2>i)@ z@=gu-Wkz&SkHQ^VbAuCb++lMm<=-+GP~a*M0s)tTI-^5^w_$h&RJ>!K^a}1nsb#Ro zImhE~acNW0(?-6}W(6R`HEXoydm4dnMU(>@$5dhg7rXjc?9L;r(rL#_n?D3^;JyL( zZV@FfG+e!g!mUp_4O|_>bu=I)=(QT1A46$1R^h1WBZ;$$Np;;G(5_)M1>6?o102Ve zAzwAJ-GgZd`E}oS`A{4O39o~N73ZHl$SRV_LIyPm8P|dx5si&2+jzuZqGTZ_3R#bF z&x1DXXBbt=k~uxdBTe2*B}74^wE4+=2;P+%)l8Nu_9WBLzt6XF9kb}Gi{)TrJXVM$ zIFe8fo@}*EgK6>znGyc2TdLZ_ml0K4eOF8Vv}dJ93qr?EXYq&{F*CZ?3K%W^v}|+c z109=IurB1^nifx$Rl}$AjN3?;(8SzMG<9FCHJ^rxIqlB1=$k>bLESAsRIu9aY^w6Hsv}AI{C()CS2h}rJN5aFS)bOsp zOo`q&D#V}fRg1~~78Ds6a2Okkqa3#H5j+Q{--R=?mhqbB;fcoJh*?S@<1tv*a+s#L z7is&jIf^H=-E59^*4}l|k*r<}@FsyM?f&4jtSYe|LOw-QmmgYflnsPKj^n5VQl#!0 zS^^iJ1j>EKC;4LD0vscRtt%6e1ZH>aP7R?h%*p_9f5(BHPd;oe?1|k^xuf z5cRHq%wE!>pZ<6T(h?;48#}6V?&kI2xcu`p=FqrR!+zq6U2DBBthWsDMLK(ocHBwH zFP;3+w#RboYo~oOatoeoHJ*>!KrLutkBl6SW#QzXck%~{zmT}^mi8by`QyQ4#397^ zi9+|W^ zA4JF+s+}8IbkfjU-_kF}dcd#+?ckB^&H)$Y6G^8aB8Gr|i1HBM(|# z_#Cd%nszI3%tZMN&qz9#&cBHUR$VTu5;9$oC5zv@Z7;l-WvRr#KT&Q~M~LQX z>x4wN_NoxtW^#PsUtVaT-VT`NM4-@2QAl+NqWHqKx?YSb;{};%TumIFaNa+ge^mOr z+Kwrf2&cYg9u)$jpfI;v?ZG56$I!gW#MqyDS|U{5x4JFqn?)@?^5!2BdjAJ z2R$BUoWFS%tqs9=fkW~FH=^oiW}q2$a&%7hDBb1R?7{1#Pm8G{%gM;i1JSJ2Cm|~r z5oy8#^Pr?E#%LV>lQRH)-&3Bivz62Ue=eZ^XXW>OS*S8KEH!H0d^m)5Wg-OrOGea< zzpE2>k3X8y^%JJ#gm+b}ZY>KteE|A~-`CXa2nb)jS{xo$F5M9RVZgM zB!a4~WqHqSW)B*5pV=Ly;O#NwkHp=#$fbuz*Q-AUyU3GDsJgHX{%(_VnGbRPwr);9 zkHAm%H*YmhM|@?qy&zVC(wh!{drA2!L3v8J8beTA)Gs(`EA0$CU!0azroB~V{mh=F zC;{84m+lx`*MUN`vgACT!>%1xFA1RGaLgNdUUmT^kmpMH0IY6D3=TzwYN_J1@$1ft z!Fa{J&=c8j5tOuYF$sFMLt_{^&ZEQeK{aTMv}%&{HrPaY+E_F}1KeYxpq5>2`vxPT zWx_kOQXYZOY)yaSOvAW8wpC_WtV) zL7}TB6;Cc*6TZ&@7wgA?{g_QeT%|gl}@>v^3I8t>Nfmv zlDLz_oMQC5Mx)!er+<|*bJ^!6azbaL6uLO!n4Y1ouVEz*D6G;Zq~6#BVM8S&Wrjq8 z91@>(*5b!&5kvU+9$Twaifp3dX9_vna6v;oNan%9FgmW_$N_Z^-I)CIP=k+8cr+(ntoS9-qNNxv+vDuT(zv{ zv=5Yld&UD3|Gl$K17E`0ld6YaJ-s2(-$;^S>r2{xaE|C8c4!z=y4THy#|Ysp@>MF& zdX%QAu*lxP^}DOO0Q5utp-Jy9VxA|51yZy22H35-FoV1}h3MdT-H?@L2er%#F!zET zX`NO?f?Kyu`}x)1pvR=*i6|Y$hYvsL`oB`>Z`-bfJy?3x1_*9%y(4g!y(7aRxIv`^ zoO#Q5e*LnDCc!bgHyDFvDGJe(onc>NpKWJ9JibOQwxb$veKkmp&w;X@OCp`bq22Yq zQhAM5nFC^_L9wg^id)v{r|gBvB6rF2T=%Yog!FgKP1#||f({*uhQFN)+wwp$*ZHTW zFt`p_uIWSZA|rtlZa9kznj(wdcs2zAI&uVl*2k612r}#$>WLMVvgWGX*RY|ORI(7I z*yr=?Fbatt0zGQ03G5JjA;KK$iD@J}rdINE-*9v&l58h7N%B*e?^wV?oZl$hh63mg*DNJ(%MC$z^tyDghRTCLoi!*8T$y@kHJ{k zq`5Y*OwqdbfhRgk4)zJsg>@qBL3bh@_i!H@=O$lEMng1tF;Y61d=G0+tMYY$Vb*F% zyDH&!Z3z3mapI1=Bu>v_#8YuwL)e&JewA>vPpF`0R^L z|Z#xX`@%`ZG_Gpkx}0}Nv*Zn zpKV1s%_=Fh`0uPhTKLi!-P$mgGMRZGsaW(j)M9eSv>?>Vy30~+>7>B!5__kq?Z?7xj7oSemJ5}kvn zZepoF2~UdM4?lyJ@}er2?;m0dRWzSS2NNV2Z-#CN#T+9Tg@@z{c~%Qz@4QOVwC>#C zahQW*QtkN-`O1-;Z+>2i{0B5lVT&rfjP6xL#7Uw0qvT_4Y=^+&$o$Puq7`}Eg7eMyf}C>o;>*UWoMt~Mz2>pMxjau7 zxpE!_fzXk!16izq3dV{3{}uw zR`bah?+H{E^=Kmo!alg9^@@oo(-Sjhp9&w_)F%4>F;(4r(1=$yf4ha{6Klb(!VX`=aP%!a)EQNKl(@ zhl#7jneJvYAV#d1EFxXK);KQw42OB^q_q5`H~!eMa}7Jld^L%J?9Pdgs_D5Wh?Mnt z-%1GFxV8%qf%k3lZFSYrimdge%oBvwtF6rC8gBJ{cCb@`+4TO1=Ek^c2l$K$)g~PXeLSBELnf63}yC>f%01{zsRv~k1;iO9Fei$^Mz0?jIR{ah_4iR z;Dh(a>LU2$=7Cs~%t%rktZ|VEG^M%lE;aI`zOX(^}B+zjeOxzd=Ov|L!HU z@O1qL>4QSlqsCJe=+YLKCtD!i@E!p866NoJ_PjY`S)2*jWF4fU21)y0g zo`5jT^cM5HQOgh2)%Whnu>sj;gUMt0nj92Dj!cynqa8(Bz7J8&K|J83m!nBJOFrJ@ z5+v^L({mXg$Mc4(7riWdoi~Sk(QqlGh#3_#f^aEkwwR}<|JSg(!fV1;DIl@XyyyBoCL9@DcutwFYZ~9 zAcQ{E`VR)ufC%EBN3UUoA7Md?X_!*#nLYH`Guxu4jSHUA_MP?~L`sxA2G|fkF5}#t zq(MTNHoguV_O*U}Lp_B%tX_)ew)Z_aZrHz$``Z{wqTrV&b{b;6PiHDzSxs+`@$0q*S?x3g)Yt!KhbRJ~ zzWZavj_YlL{}=u8?2$<(PK(#!*=+gg1JgcxWYc~xwR(-PaB9)l@!Te>f#@O?RzJZC z0RPDm0owB~|K`zhV&C&00I6W;+5f6e&Py`&F@5}xMD9O5qvd3A_I|_f(SPw2pls7$ ze{p7vml$|{RB!r5j^&yo{U1{iNX*AqDr(tEnWe-Xqu0|sOrCX)ADGs@k#3I?&DS|* zj%2=ZSI-k{{||wZPlh>yck6%9Ol4rd9y358ep}KiAVMzA#cv8Awz&J>45BfWGB7&! zJ%D7{D#~sU4>$;KjsLesh$u0&I2^5FQixs~13b=Wue%$Uy7PsroHx;oTl>G|YWV$) zT7!lc?$5tI$=0&2HXmL-nRDpR5%+y*vn(ffp$pVX0@J}{$@ePdAM|?w#J~so ziIQ(yZKl+iQoqq79@v7EfE1P)3>nd#^q+)ZP|h^CpAP|7YtFl7Vft_+pF0jC5WKB0 z>3(T6itBjXU#J1&T^Wk{{+T{lTOw_Cmv^n4|T6 z8I62VX%ZGFseEFKe(*+1AU3c4s`4cO_-_x)_=!!oDsE`K zCx)FP#YvP#5U=spjiFHgiq3kmQK6z#I6&mhXJ_p}+QY&9vprqY1o<~J(swrpDe47N z^C*-xB4mTW%ruLa%#*E{6lI#Ozf4ZGOZFTN1R73-IL=oykV*VA5Rt~cL9{|a0bRtj zFYYyzOb)>Cw9C&g~TKHq}XQ5*V$E+TwrK8OsiiiV@(Jl<9t z7!*sXMj`)xAu#vJUjvW~Nvk3~mp^BFn50P-8+h)OVUg7J_QdS*U%}ul)X1xnTQ|*r z;BNza#w!N|9XAW;WfST@fN5Q3BQR?;KuI~i`IPBYc}!0z)tg&+rkw-r({uG+QE{4g zntbiL6m*G%e%)|wFGD@8eEM<$KXgJTlND-YJy*|%TUEEJ@oHX_{8n35{_>Ra-O2?J z)BJkQHV=#1{({7^j-Rg5XQR}OlbJm9gRDA2$|8jqR8;xrV*Z*xgzuu=y zs%&dY_HHsyLzuFuNZ(&f_nyeNX452+qBQGzUo&5G*_oJq4t!YHvxDSQrXaQHvAWTx zZ8Z;D2~#zQbeD5{yytM}p;32f%POWZyM4T{lzI_)2P(a!$)PnixvHjqJ@=`=D3xdt*G`CPQ6KwP6uU@z{db(EtMG zv^#}L?W)mPWcx2us!rnc_V?-STFY3X7PXpBD${>nMI@UD6yZ{SFG#Sd)8^NUL>nJM zJe4uF8d`D{#&x~dS+@*54l|R&DM`xk#q=g-fcl%%=(r^m|Gb{f`6r9m7(ac03}M_x zimiRh+*OXtYrqA{VPX z&J?RghV^XXTKR81b$U?f0Q&PF@`6T&=B(#!9&W?w#4M28TKNTAMnS@{Os~QPj#L3W zNJ~0+`#jJSyTwAS(&iZLYMbBkz2Dauc7Ooye3_--&IZ?yHWdJQ?K!{e7f#q_hOo{I zgBp(OHE>?3A{c`KZIuQM9rjo22@MZ9AAZO)Mi+j;rEc zcv?BoK@n`dP12!nXScv6X|k!UPJa=C!utBiv&w*8Ckv}tI}PYmvkBCksZ4ASo$T^} zlDQn*PRFR6KTXcuuB)*fY8>`80HuedMwbISIheKSmGgqGxqHs>Sgzr08SHHN_LW<< z;{2|o6Rs5y`I_$jOXN$yz?H*kMd4)|pY^=EiP(jna1O=YP=|kc;PW>p2oPue*}7Ib zlr28$7V0>;mglW}s$!;v!nFPnFB#reS)-dYuxhxq45$e->MD^LaGca@Gys5G2T`m3 z+q+v|5vKabzFY1GxD~x;O^tVUz*2{b9it0~)BP=>4>&yEk@dAAb38ptdZuP{0f>ja zP)6-Q1v|W+y@ea9cJ^x3HgmVjr`B@34NC2&YCM6WHeqn4-YtQ!@OI;{=5MXsm=R(+$Ma(&x zh2ExN;0sll$Esy{&^w6>HbbHB&6&+9RR;UU^Q*h;5xadrb~wGYfH*2HA?{BiDz zBvd8kU^a8hxshUk()50j&*8tn$>+jBc=ILd`0A$w!8lz$-!&wSYy!|T(5$7~RHlx$ z>sPC~2_dK4;)0vQXvYZXabLM=@ZKyw+NKY8s$p-8^&S)kdxWGJ(7D(CSl;_&B2A97 zz|H{NO~%}}0iFG=<)*!&7LQ@n&aD%ms`sy3s=Nkt_8b(a+N9>MK)PII9ren}fS zH-{_PvUBtGHj$WXa0w1_{(QA;Up7h(!jm?*HRk$_@}xpIf=SRy(iIlvzQ|mBMndaP z_9j!X`U_$c23wjBsP4X4c8K|WfBzNPz>$v;IZ%OOv>pPa?J{fB&}MrB^rTng&5l{W z_vr92Bz3-=T=oV$h2EB#UxDV-C0USlC-~wPkVeWt2*=T--zC{!kq){)Q1R|6@7Qmn zC;2t!aYB42XOMX)RXAz#_bdvb!uAe6JUZGdB>!l6YX+;MoKUgJ9Pam19!szi& zOpu#YFv#2@f+S{l{7_1)d6fDCD@IM^6V8qV2lUSv-}qTO*=!e-k2h_kRN9m4)!lZ! zZI`%k6ro~7r#8U*H~H_94Y(*vFrtdaI+KW8TC0ofEKE{A&#a6xZvV!w>l5JKb)e60OY z3C@96V_hkZll|Gm$H(ZU5jItyvf!>kf18ydA{-d_$;UYfJnGH4!CyM9W-3f&lL&bN zK2y<%3>=S(Vu&QC(J0b*)G!3~syE&NDzdP}=V$-eWr!%AQpE9yz6*ub=KSGeWR*d1 zm=LyQQ_$Kp0-sE;<7(h`_aA-y_hW#Xapj?*Xg>@5OK7;sv2M*81wAsWVnJxBthb0V z@npBH3-pb!Cl~RF^vK#~h%^=@T%+70#E2thjPj9BpW$Qb zd=0J(Qt!R!z8<`dVm)1^gcgoYhoYG9uIW|SOql~o-*04ca6>|{wn_W0{IzISf4BRv zB&E>_fuOC}@{BWW@CJR4iuYq_jB#23z{0Gw69YOO{?k>7k3&=9USkz9UqxE6{b+c{ zVJ^EbOq815j*i$tAZ-Bc5oP*fg<;Hfs^#YMl1Q8Ci_zA3B~18qy=$?^=t~8;;Kn$t zFI+CW!kycPG!hiI1S`n4q**Af?*6=Y7SXqLbfPGus=ka6wA+j3aKJH8f^m)_^JDO5 zWX|naw-+!8BuWxOpn59&GmUq(uy_7eAx@tj9?^ifpvtJrO>uF!|5mUERZBXwzTHKya&4Eum!x?PQO@WWVk0Uiu!q_Ay(5c(+j~T&C|!iwJ>LPu2HxV0vKN zgjmU9x1M!%{(h-HSu4$my78WALa=p+d8O%YR7qGGj9F!RW{h*pz$~6rA@%)U493MA zrQ?!A*_V=UmG&VY8IGIxivlc=JbKdH?;^Nuf;7tp?T8v&>a=tJvhNN%LX?$vUO=&A zK|Z9yrGRq0CU?`rn2{8J;;72~G1N%3=8Fn-UrkG7>`y{2&&z70qitwqd(`_V-QUhP zCZ^+p|6I}^tQxdXS~ZEpX1{hx!elB2&BM!*>dd=u1$APOA}(zaYl z*jOERPkHW4>TC?A(h-<9$--*qIRhrqtkGLTZZ8mzji<(H6N8(+3Ux%X$1q?7=Pci2 z3}Rkm@T)vn@EF^R))Yb=!=HC@AV9H`8tRe8rdQCL9s%&Q=OxZx1spJSN(x}hqy2E zYw&{!*+UM-DJ~F=$H0d=AcX4RIq&J`J{U=eK7vb#T9FnohSoq}mvcq;KXz*irNBO3 z4c-K=uUE6*cP9+|EJ5R1jOx0-*#F(eFR~M@`n^u&==;P256~K#%c>W(%mL+RZc^YTS6Z;Tj$YcJ0<{` z(+Ce%bX_d?p>Lzi=H5n`gs#alv4DtZ@;%3X@xrXD!u;^N{BhH;9Tkb?@;dqPq~yJ* z6M^8hUC>$H=SRkqj4V;(oOHDhkZFj!HY3j?SJk@EnC`7T{esBRBB(Y}!jXJB zw`>l(F*lsXaWQPeb=t{xB;NHs85K)vQ)+mpY>U#HXD=Q@kQhwqwB9G^T-SqNb^0B_ z26D^yOv6oFucgs4_mfT9Q30E|%^!HvW+|Q%jMf8*f~w9_B}SrK%pw>YqeOog#@N_b z^Fg#z3iIhNg!%&2#+MluCF>t_&`pb0o7ND|uy!x}ahANJUUvzv+J>MUmCF zvKFE(IItshuHlZej-*>U)wd9C8w=FP^aGYp0h4zkXJI{bNJ2iyQH_m zUiqzBtygQG`>d9VZ`(HMx8v3_2#QGnF7gljmqyrwVZ^L37Mh4@=aM9yEaV@8iT1b@ z*CGzQ>eDoc7czzHByT2CQu8eXpF-naLGt8R1)H2gHpoZBJ5&TJKajt7Pkbi(tybAo zxSL>xH%N;%LQ!B}(CN8sL|!PE5GZ97iXpB8E_E(H6L*5jd3S~Vnmv(;xEJAzDNQeF z|3_tpUBx_r<->8*TLnr`bcYXL#@t*jhhkQE--%vLS<84RF|1AoCv9Y$SZ@8j6mEQ9 zZ$5nP2yby^(XV40&ru0(`+hL6ymv#w8q-s(vBKK1jP}T|*GtE%bcXcPZX&7S!;Z;@ z$s&D)Ndxz!zZ9SN+;&S@o*C%i<*vFkTay+eJQ36>Loifw28jP|2)wW28|@|D821+c&_ zIbWHhqAGZ3=HqFWb!GeF%6|Y>xFR-bUuL1Mzz>>)5yg^RnuMyp_E0j5u1Nen<7%oy zRU$(r+sz+=990oi!ifMU%G}iQ{5_9sBnvh9hD1(r<_h-E@p`}W7xdkg(8;G~{fdBi zQfyw#vpP9)XnG3#)RB3b-ft&HNr{B-1{dOeWK=RdultK07X|2wzul+7K&lP=ghD_- zfU3QUBTIwyw3@9qFYzH<%aFFLT*I5Ei9A}UF4;=IW(}7@)1kS=JeR8 zgk2ig9Lfx;4Gm21!8oq|OGFJ)S*V^kq6Wk2`5vjCOj+AYglr6f?PV&cFt*nQj`zXXzUo?L#>D#Gmh7kNjkQ znP3DxY^KBL>+J{UlJEeWBE1yN6ATEF7EwI7lqj5LWpF9JuC;>T$ixr8DF$E!AUD%J zkbA0z0r~u$_AwSp1Pv#LK=9HxtB2zMk6pZb#u6qbm6Yo-X=-D`)oR_z-c+SO9^^78 zOX+i;ld+swtfHJITRL14O(pZon#1_tzr;;4U%gUvD$QRP>aJaP@02xvHlR>wx;cwm zpn62BYc3dWcIOM5sWcPl|B@)**f_XeBNi`IW6$S*E3!e1wKbfOr&;`5%{-`$+jZLn ztAJ28X>r#|)|=a@jKLXqzQ#4_y(W7!Gh0M>I}fT-ldt2}kwyes#FE4V6mjtS1M%g~ zlvB0p<927rA*NZ1_uEMNq&lLyRTwOQ8njKGb#?T2wA)YG+}uQL5JQWmr+16I!f85L z&oel?`9^na^-u(fCE17~q4>qk{;}rv>5ft|VTpKD@SToN<0(gS{sE%mVEjAFI1{6I zJCA%iW6`fS0^V6w5eKkWe;%sXoU*^cXHn0X;$ub8X&!?hMLdFfE|AM!{7@P#VXWQ_ z>!;h9VDWM+k+S%ClY$8$Z)!;WK|EqH9Kc~NF~*ih5_e@&e9Q;kJGhog^6)ARZl>G$ zAgbmEx%xYV9#QOWa((8vaQKYcnpyAX+3u9=#qnCX*|wci?$(CIpH!o{Ym&*IUyZ(-4bKa<5NMk@`eu$?+?i;Ewcn;Z-Hs#sP46!sfM#VoIz zOO;ytwTh$9#EjpYn#95=_};JzyB(<*BpM*WLF=4$NDp1jXSlh!nS_LX=Bekt1Tk)25;Z-&Csd*yadnPcFMWM| zb!vHpVLEj@W2@%=*_o6n=opZd`1Kqa3;^y1cHxXh+JN9-uq|L{DL&9F(3lI&c2iw5HfKm$yfrYjnD2=GH}4BCAAg0_d^Z#f7kH@!VF_j3eONZ zL~npPJr>3Ad4S;~Zo@Ei9gW88G}AT*m8uT!r$fc}<{tsfg*hy&0r5YCNYrPWsM0D= z;?p1~b8lz_27BkA{GzYxn&txmM|)%9;(J1pYsOM{KW6jPeJB3U>nr1fW@D(@3C2U3 z2mici$dT@Et>35rufK(+a;Y0U7dbghU@|5q;(FA!zDddVQHt)>*VOOyax)d(I^Vv1 zlW^D=$OP19)4~x?y$o5x?Zqas29Wx#oBh-ZI)PCzN!$dqikExx9blf?)(N#=m1R-+ z6_7gxzXSBbJSZ^{(UX}OZYG_I9HDi9!Qv;qJgN$~vx!*L0`7Qe!4wMyUXG=)1iUhNAz^?m;Q zIb7a)!lA>+$Y`v}S~r1zg5_^*pM&S!t*7NUUhm&Io0*DSp$`N42hJeG3^wJ0tU2~U z2Z0MO2Zw31tJYE(yZ1m_Er8ox6 zkARLogDD)!@{kpP)MzGX^0-xq=az1(Y+3k&bJ%~(Rz@zk-qj^bIEnQe!D-TWdwcuR zE9}WJaA5^{oCDb~_iH>RyX5szE$F|k(zwpIMmq@>f4TD8_ksOI2`Lv1+z@lnXP0w}b5!IGcd=VEq2=R?{P#i@QasHjC(ZPdNK* z&CPRcn&wKRLrH@!NXd^IJ+E+%djs>jI`&B)uU`RrU7h~H{O?K5rfWYb6Jr-w++)w1 zOQ-oFmxBfCxJ57u%zM?+{!S>KJqPstfXg#}<6KnyD<-+_6tSiv6M!PqU~<{5`d#Ej z^_fbPT)G;BT^}#I1HQ=1B7l9d2m#Sx|42g{z05Zxk{$smn z;)(159Bui@)mpLbuCW|AObyuFLh!8La29L|U`e|O8cJWOq38#%X4bj}L<=_fug<>3 zNvSFjHrP4p6=hJ9dMWSrNVeXMsHBp^mQ{v9YbyhT2n->v zZ)R+Z3-;-D>K+`7NWixCz*SHby<1Na0aCqOS9W*Ym6J0r{y{px(WsZalTdT1yIXAZ zOt3$?5HK`W20Mk$-}d2N2>nbfFoU~x#PxrkSvX&X9@pxpiX+YXm#aAJOt6<#4>Z5x zDU2)Ldq+k$Q`&`RNEZ8p54!MT9a!rKuBOyc%0DKILStDpzJP5tS$1OumMlfRFV@Ji zXO8%wpcirZ*k{Fzk_3$ZS%!^4!Nuo1ALSIplpHqLGkCgl`G4csP)HdsP`Jn`#nUy(WgSD*gip*KOz$vg>G=wn5EYEk>HjK1Sc3osLAF_6D|?;e~~H}t@- zGw{@S`a z1@&J8$38t`Z2|U^*C{J-)b5T#0pK9id#8RveqPu;+6xmiHzKaK79nw5HZ~7cocr{k zc?GTO6VksvTSoks>1eUys@GE|PJu!oV@r*}ywhU?u!6?&^vy!caw6z*ZLBvY%5@x9 zmyyAI0QuFwFKr1E#ks>%I@e-r`{DJ~uFZg@wT?>fQr7fK+r?bU`G%{Fh%UwW6Qo1w zzA`ZXC^N%2^aX<{c=w=t}pF$w5S8~ge*H5QgdRLisa!eY_Xrz?iScN@XP zTPN|4u&efF8{bIMwM-wgwx_I_4-c!oKxj4dSa(>>)L!Lh`}CbYCkbRKwuQKz%o&sg2aQ3m8+U@l+ zeuZNP^~__*cVlt*Tjnj8#zH?zVza^w)4G^f<|(H6oZiW|vw&yV90Dy)^D}TBAw5jp`QLr-#Gn2C4@oROqKvDB=PzK+W(vrul&iz7rU^r3u>ne=iDZ7YAOE861I0(UUQ5sy3L#2$~u<5F7)N-Ct{{+~)r@AgR_ zf=8%1ku?#vCZlt;e6DtRaeU%)^WCmtrZU$86hFJ6?)>rLDbUwq`b~-)8ij8pitAqn zPTLYq%{cEZNak*B0md0!9?4_by_u>~(~sp)MrsQ{V*W{_l=N#y>`s1s`Xx8SHmxPl z>hkiilX`0iPIuz@3T?btO(v{1HOX-3>+CT-m55iOP?e=jrD2h2xuGrz0ko=1i=(O2 zi+E)R8J%fwPH;T+c|tL}e}b;SfU1zo?ZZ0v0|+kpV?{8A_p7pcZDKXI4z@>C@ZK-W zwmBN43V#YPig}BLF7aDWzcKs%gT1E@zs`AE3Q$$}IEWTHWprA)GA!OzEOCi#9F-c@ z7oRPMj^#<0%pdMgkz_Y2Ic^X6ew^Ul*IZ5MI6zEM@%p?Cf8O zyBL;vdqu(hDU~#v+2Gc@eYySZo&vO{-|To@yRuZ^xq%DpPQu9wlkq#$enAlH=Z5vL z$IkVW(rz8<-|que>$OU)1ecG`6ki4D-=++v6j#cZQv+e7q>VxvI2ZkO2<-s13Q;w{JQ0^R8jHiT2O8GPuny9mlJ5W1s&lZ0NSmb9jYd zOXQI`9j`W1!Zu{^js4p%K)R;FW0rSDXk^s0m#;Z4blcIaH1jkNTX*V1MmBSpp{q2B z$Nb%`%Q%`}=ViQc)RG_pS+@NL7m{cNF0P)igGM#QoAaigO{3Paqtf(O+N)i>E`RTO zyUqCQ29D$1u4@!@kDbHM!ib>NMS3c4qSd)CWQn3Q^O<-%WAA7^f%Hd=(w{hbxv@Hq zU+I-~WXKQIGrbQwsVli{dKa%R5-_}11QZq`M6Z?E%lW5u`E83<2U85o`6^4#7s;E| zNA3LmN$0!Qeo%0NYFmdvgO*A~Q$*6nn3{=4(u0N0`#3%*I262D>-jhz zIGSDIdJzF{jm`oqww>jXK)9+OjvK8OqV-u1mhGw=i>|Of!=vkouaqC zy7MP}B||K1$EzzN{L2ys5rxU3%0)kAC^8?MV4tqD9fHO^*Q=8fDAs$sv(Bz3m4185 z@l)6s4VQfN?*&hHH}lDe-1rz+r`V=0AfpiwLm$mlEAbnKW0eGT5~kf1dn0y`g}mVp zQPC#4CQ^+jtGAx$`p5i?yYc-T`b#W3EnQ~xlA|$|t}}eNNjo`iDb2^HXZn(mozot7 z5G?g_39h;?m)qIBi*A2Nr;K!!8gE_aK4UQ~N~Cxnt;3-PjH;VfEnLHqPP4|i5jN^0 zc^<`CP*$xrO71kSyZ$50ON3$TY_(w&lTklQv#En0qAe7ks73+4sY4ez_ z0M3PzIrD82BYC`ZCr1hP!V_vuf$J^slv4Wnr|tS{ZmM}|B^T}kWv;XLy+u)5z1F(r z)6RQTFYl(S!lvdSEZSuzRc^RBv7$()SgUM^dZYa-O)4L zo$}M0j@kLCA8Rhezoq4*zmTTccwUZluxl>X(fRRcct0gd!;DUY+lOk=JukJ-hDa!7 zwy)XaL=w=>@c9l}_8W7S;7KHG+FwNG3Q_c-4nn$%rwM%7Z5hwrr^u+CZgZ%#E|KoS zKYix)9C=33yAO-Ez0O#u`ACAwXJcYeuqTNyXDZmQ(Lo?=`6^wH6=rJf=qg<+qor{a zW-Bz>|4gULMRPfN>~A1v@%&Zt{kUB3_Tr1kkyYWp+k-C`WC^{z46 zvMnbr<(L48;0?83r-h?ZLmDdz{+;uFDKR6{Rv-E4@>%x}#tP4Qp?nwb7Y{I!vqf6fFTd~qNQ0O6|M~l5ZWo(gBbnAdzqeGg z(&fE?NNL!$;h(RVwl4F8t9|dQk9R$|`oOsOl|Hcz59FZgoftz%=qThLj|?*JiPhgd z=;kpC8cY#u&HTX#Kni^^j{5HBzSG#G;!#X0ORb}f(E5)_btUeLqH)~LY>rz)3r~NK zx(JM!5BV%o&(#utA!C34`G+AC?Elp>nS*(_((aEe?9!F)S8p@vPBT~UKo)MJPnl6b zkn^WU=+j@$zGQCHnJOz!&LA zQy^A4@u#YK$u4oJu!euV3Zyg}e2G*1IL4!`eFm4ZIoruni_;jb!hA*wMc7l+WsBwH z2|rqv5(K^0;e&%vk~T_cR31uLgQd7-id$X2$KH%nj&eR}6tgx_On-`5^xar|dFpwp zM2T@LT8>uZlZYJw7>4|NFMv%-;)vto`_NFhx{y-0gG; z2-=9wX8)MF^R+<^0uhy`$yFc5^OPIxO>c$gTepvvNmyZjW5)B-bfSCmF3?G8y>_0* zm0OYQE>t1?k`ASadeKga;hR@UA)YoTn5MK*RcE1*Z$vlMXh}B6q4DJrr^^g4lYUu1 z>$PTq0W!Ogqfe&8hM!Vv${>1h3&qn{N=dk}wwJBzPQ82?PgP_v0`kh@%~M{uVr>i9;%SpN&y;uwN`*qgHdc1Mn`xohHv{5?S&Wbz z+pLojgo18>)7C~~(AUP^YKypSvvZN-bCTAo`mLxe#5D;iU zlqf1WNRttSmW%{%^*P_~yK~o?yVk6kS+nM!<7wXat+1qewf<3l`|MZkcN%dQ@x&rYimdTGoFv-s?ePP&EYiNIMpzy3 z$mq#_OY;& zsnF{3L?S!Cy~P?u0C^dG>A5<$M*DVt@0%FsULEJBSn8jPwdR4yN@~#ng2Y^YG+%JJ+2&|glZNl=7wbx$mIo(?^W-W-)Gib; zVoK0`rY5{EuaMnAmt@t`HN@MzW?>?lSspfS22@?R5D1K*HsI^5xBbte5|pI_Y_CE(<@tzYtB=OYE3i#jr# zpb~s}H}SP_+@vYEQjsTRAZ8roWi3`OjW%qy*>OC|5%}`X7Dv-8Ec>Sd4St1pgGhO-3V!P?re3=0Kcm|pGgACYN z{9`RzmB;7! z+#{P;5}ickD#@*7iox^O^ba*F9%u3EY89Gwb}YLn)UX}=-2JAvrWAL7Fz#lD2f;#E z_>9-=)2Q{UEQpvz-Eahg}HV9MRgVVwA$=*w;+KIY-qqW+^c9cZhaoqdfxo$mZjj+QwrD* z9ZmnScXeLAZ`^X!VY)qO3g7f^rc`2>e89u?an$~u88M@%Nwq;kT+|#rRa%90@cHksLXh4UV^@9Q}N&DN7FS(%?^*C10Wh|p5_wg z6|!QFH{bJ0m_9~Ed>dd5a0T}YyXg5&?9U^|~)KX|@a|5YOJK>9AXd2R!X zW)Uxacfcv0+n=lBBi9&hkC0ajX=zVVKLmW#elkd=TGPs{Q)!`Jy-9dA_k!)6?eG2h zFRi~PDWDoJRSE`#TxVQ%hE_7BlmEPP9;hBl==ii0l~621>-MqKD2)aAdtKF{{r3k# zQ#yW!1^sSGYvg-yz(DBCS`zGkGB1447dq^4s&`$w>n+gea80L<6wc6;dCv#uY4dB zFdiCKA3dGJbXe>}cb;)yRV0poeI07m1dn>7)Kv0q&W(zxm23pN73{TVSJTQUI`=y%I`6lZgOAC&b`|kjqf% zLVp0;8^hBV_tLIh^p5^K$wM5yb4KP9ioLbsiTa@zctys{!4}J=jmVHS;QVlV5GAKbe9CP+e+~KzO_u`N{dxR{~f7Yrr-ISo- z;R5=CQ(j+RR#Fw~agZ`lcyMrjCt-qh?G!@lvIKWhf#b~jzdy4qLUy5il64|$3aqF z$x%N}grxM-mB`bxNYnJ;t(Sm0jnF{TzC6^bR)P(y7ZW^Cf&ITNj_<9ox5j(P$i=_6 zIIcl7(0u&26PB37438q6SaAaLUE47Z`}6HE_i~V92sKm!}Gp}2v#Ix7Yh@NQ%Twv&wS;t6cH`-4~emJ)O+Nj!6V3IghW8~DGeki-mrJUiRF z`emw=#N(Xz8ibnbFu7ZN$`p!GtaV%I#!}TTgaId<+{oRB#db|4Bw}u6F-I(Gd)W;h zTi5=v_e82++$vav)Ng&GSsP9w-;k9KZ7#5O;4`}Oqc=1LEjVDbrSUzHofNJ6`Dqk> zHs1j7ePj8*+=AsA)F{07cB$B zaQ{iJzJc$__j!-Yibua%Yva?S)!G6%>0x)6i^>6YT|MWGZ$@Z6>>*>x?J$?3MfsjLL zbDI;FK@F>j`@Dt7a+6G8S($n|QjkR@aYyIu%S3@*nOa(hSrd!${p$oWws@Em(cwaQ zPQ7IcgBgS8&zU!oC5b%I8t1A(!N0!NTYNB4_z`@4q5d{kZoUIQoKs{^(^e&^{kY2Y zW=Jyfl-aUSn#kVrM=zg2mC=;mq|OSL3W1YaZ%Dio_W7^Tz1r@=&5<<4qh61Gt&q8cioXz8u#l~V4zAFzqvwMmzJBH=62v~r10TLuP6wnl6bJ90~2HSTg7J5>JGDF znFE?%+3;;U6eA(Y@lw#fmS!NlC(vu&XEe(=%I|P4EvCV^d63ggI3FteNiEtjr%5+! zSi+0`E@Jmv9qnYV+8c%RM)ik3#E2Qx1m+a)`_>wLy<8vW30M*jG>oSdtG~sf=Ev7r5$znJs)*rF+m^6qK)mn}Q>R~kh{kJ>dgVZ%v}$O;nhKsW&XXew_( z`GFLOZs2Qn7$4aT@cT6q%6HZ7zHLVrDz+X?++Ho_^I~QL++2|P>bX|#eOI$iy`OUx zgv2d8w%4(w($ek+OJ8AWtmSpivl?rg7{_PJ-SeF8x+!2QGc}Gy?j*2p-up=UMFX1J zVHG*stWMsA-C(jgB25$+ri77m-_z!y{n18NAn7e~i;&l}#zdq`JmV6%5totQ^F-8r z$w(ugBW)v|f#E*CZ8Xv5CgrE>l-^6cn+=B`@#1j}89l0K&kIrRueBzvY8rWx@4OG@ z1hi85S9FO})NBI_g#&~q&-;?;vw}7S!AiawFkqFwI0`$(iW@r>a=5>b`;LvZgdsMd z>U-WRf#84@);K%*&7s3O{oiZ8r&Wr3Gu>lcq1V=NiexE05+&T8g+e}@*{=Exh30f4 zKu@(>kRhfhwMTdoS&K7~iB#gLVo;6cg)ULoHBSy!8a0>hZ{_N?os@_<)(B5~&CKQm zMe6(dbAb_rZ8ea|&8Kg62Of8Xab@!yu`mSse@SlRF)>=x(K;Df zU6lOjyo>N>pkh)^qA{p8sY++l)-4nVs_KyVSrQVnc}1QM{a(mN>!+$PoKPvFb`K-X zPYK6^x~0%10Co5;i9C9=f*P+q`8` zyeJT9*)=l*k;HVZ^1OWYP3(B<9yxNb&yx{I$6Z8~c_}uy0^X2!QZodlM<1Dqz*)}) zEY`aPRAKBP-B7)s84oG;J{|CwLVM^$pp7=&-TD6CDm>?}W*e{`25x;K<<*ptKBx+9 z)kt_(?@fXpz+CK;IkSI>U+fG_#9%n)&RknkCXH$`eT<;Xv|d!Rz4dlQk~B!RA*ItA z&ECgclIPSk&ecgJV&;AGdEg-`MoF4`AJXj_1wdaUl?lP|nP7$`)-~SE$x)x}u(n(kr-7Bx-&_7lOH2QZX_30G%@G6oVkU-8%9^Pp z&Ce$Y-V!l1wvg`8v;Oq3`77Ji;kC^t7p_+xHa>;ITIPmEF>>3S@cqy+C`MC9&3i*I zOneRq9OkRlMX#~IwLuxG$Id3@_BiUX=&?u`L{{T5#0WJfSHI3m55=S9Sz?655X0vtvU3I@6auF4eNcWGriU6wo!&2pVhG?*3Cpj2|bqs)G%&q=v0{x z?Of;%$&$1mQMXONZ}I!mF7F^0^x)!E+V4pgFm~q;*B&W6;|45>m5TdzpP994jaEZ< zOA0hhaJ64^8uX{k-(Jk{#3CpLcps6)#^bj<(@Kk0$w@S+{ z(OQ%4PrMeGkce?ZE`>J{SW1nX*VmU`YU#3(n}tu1F}v%-if1myJl03}EYdH9y*qrh zRmn0HATUDARw5_sXX4IE{_VacFx+xq9@C?^p}8L32RzvijlX^s;oTx@by_id?|pRN zavY}F?P)CLOjP!y#?aVKtmP7$>UkIX{nnV)mv*BcPt%DpJ~epSI8_i(bNhZ;ZzXWT zH8zuAomAY-1&5letMATg_@P|M502N(m=e$HEtaj#;mqqf78YL@$~Eanx&u#QC%60K zropu$&LWJnY;jh<#tG@zg-t79dLZRAc(PFW6xuTm% zZyde3o>f}SiiZUK(N{=^?Fu|HTS?|p+&B&$jFfG~CFXrM9DVz(#lhEib;B<4smV>m zAAGX)F*=4O{l6QxH@=$??!niP6#vi;V6og^=(q{5BtH`RU}n`b=t{bQSfgpT>obUF zKl}AGk@97mLv0@5B%P5fBtuYbjBYh@EsZ4kf{#|GX!2sy?MEm63y~L?EpH!(xAB;U z9fG6@brv!)$GV(>v(Y^`M0XZX2t-}x(IZupBJ$^36ouPNYp=z9*DxCmG=mBsXC+cw ze=LhTa*UY;hSw>zIHY1M7t2LCYq@^p*&JPg8K;*w5V8GOCM`It5;ic(dN{s{EoSBk zhh3sfwM`Bqh()hIMvD8@VtAw9U$n}x=$5&XJ3EUlF>DMml`pP@!vTpc{4i@g$Z>&e zcSHHVu%13codv?Kuc=wDm?R$pTR^>7Es?|nJA9R-5ZoJnG^u`HBjZzbY}PNSz9SLm zFrJz%ufV||#x`r}$Mo)haidB$M+O)}Liz%bS24FMbxY*-o{uv*tqp$Z zj!Q|;F$c(%!}&_E&O0}AtGp)p$$InJk$+G%8}g3%^dJt)wnPKE6>&}IN+Zs*YyGP5 z9*Kea@x1+`CPHBHB*gk3oU%yQNqWQVYo1}4auVT^-ETy{^T=4xW=Y&?U|6Tt5u8QA zJ@*`zEe{6mj}PV!R(0FXH^$-41JZ7$>zyI4sV0i@?>xtJEc7wXSFmtzA8Ivo2V3|H^=2&+|))&`O3BZ z1qRB(A%T9SK|<*E9%ApiGSwIX zrxB!)IL7=JC%?(<#$Msv-ke)PKU&BRrZqruJ9ZC=(X#pEgW4c{Q{g}6;hE_!>+&6a z778W9v2LWneHOmWEro1BzdHgxPLlER#v&p2PbA0@Q#h(4Q6#GnM>gr4O51Y={?N~4 zR$Uy5gl(`t?Es_Rtv2YBjlAQM$Ow!XNBMK)lx>g^7w#S*oPrhDMv)+z4Ch{h5gI9r zQrNnS$g}kO>fA>fl=V!~npq8DqKRge^e*hCjz8X0P)c~6(Ph|Ft%_OS zx2@d2vdYKY97|hHH(0S5LaahJL-dUh$*M$z!l%FzIOg|(eex*-jx3y?DAT5lpKOH>UXYG;_ za|r_khy{B^*f3E#2^I+gmoz9Nhon)3NgRgbW%S;Lt@e4+N@oIPc(B@CjmAQnw)BCY zt~wd80TH3?@t7Nz&ws%6o0r)GHTsoS%;X9aQ=JyCD*7G>4=n9|ql@E+t|4?X@Ajtl z-CwE8(ZV$3d`5J%F~R1B7FZtg;{(P-3l6~tnQaI6sg>!P`QXn%{psb%{9lWuz685J zCOK)&pZs2KCmzY>B?KWh@8X}n2jBuMUIP9dqMdEgV?6;>(dK$wF_aSSO|S1C+eOAAq0N4$?~=e~2Gw)3rVG#Y z30|Gv6Jv;rxQ_O9I43)J@ze|@8il$yvIo?>~pV` zb+c$+T6e)G1m(f7f9h8ZvxF&{*c%gTZ#I`pPkihgHSEuJ|55Ch!pq=5OOqjKfEP$6Hv&?${d;_KC>2g!7TcWgeM92s3LIRQw^ujGlnzchNGbj;h?MgipR=rbc> z_vqoPG=lq>5VT1}FpIV?Gq~}TF+rs)7y81?au5=9P5?CMqQ#Uz(CWbdO~q4vCPxR> z_9_E-s{_Y@lP!@+4#u7+o;!vr(#iRU^xgCIU0-f#}j-fdp*onlEPAOfXz2*V)EL0 zEN>DBklI*EF(QBi2&>Is5i|4w`T>*a))7Y3ZeSIjL zp&?vP1gaWe&n?-idC&X$Y$W5qdNro9R=>nR*#I$$;-rA-Y$zD z$sY6?by>5v?taDk5ORYjCFTah+Z-7w<(XxC6a>3RFWEO!AjM;?PR5R11e`@1sr^~} zmsq3>5;WdVK>wbjBswpx@`r{dycm)KjxJw1e4N!L0$$Wx^d6)lyFa>!!W#=$cjgyhUuu$*rv!qvQo+102}B|% zhKZM5%Q0RcV)^G5(X;#`^UOU zd8UzeiUIV>u&xYv;xz!YQU04%1=(5~GpwPqBkc9{KEosDBRi028GtS#)%s~vh$cMB z-S<9+tnXYFjm2s?SQWxZm@t>(?wlfX(#q2(qLt<9V%V9OND5^_i2%b7&<$iG!Jem6 z_tJi@`T4s7na+f}n>{>}=otHJAD`qBM5-ji$gj@i!>=!)*IRXq#gF62NCvpDAHNnL z|3kn)0PNm~Bm2?0a_vkuuI`{yJc}Q_#H6gstqFd2hPb?ioz;& zOX1DaHHgVae>eP<5~)bH9>2B`6H~j3vSZ{OP@Q2#b!JE8>wF_*M;_|oW7rJr#__k| zPimgc&;Rt>;=apkdQY;=3F8~PR%605R%t2U%Vr82LS_JaAfxIwetxZ|x3eFeZ7}P? z7TTVSq2#lKpPlZ#0nSa|OEN@LBn%zZ5hYYd(puCE)>%;gFkNk+BE3|PTnq-cWkPbX&3>$gpjA8m2OYzTn%YZ@> zGG8hb4KITet2~}ECU~)ZzI;5Z_s)|2tJjZup>C9_;DkYj#{^bBIZ6>JO|TdwrGx=H z(Y0-#L^cT8;dz^1Fm-&7Ct6D;5T47XSzPbO(*+fQM_yeGT6b|~a~l$z9nQIBw;5gH zdHg(aveKOTwTAMtE1;p45DQRWIiHNns(+IL4@#@!4>ZGCV~ms=cliiR+**Q$y>FN( zzJRvdl_CfLs#`;$G$Ey{C-*y@nbH&PPNu-F%pu(!{V)tQWYqC z`E|GQdm09`9!N@%$TZ+QaQhaWNnt3DWO~ghMn=e!7S112%FW6KTZx*cO}E+OYp&3W zx`#*em>?<4js&BC0u9G3zy-sVh#IcaSv7|NmSPpCN-{l8%=|VT@&ZI=A~qO9x54Wl{xfu?zPHu84J7 zbXn7|q(iQACI|Ybp9$FBUPB!%L`nUxE7R`1CXn!w?_2Kr(jMq{v6*eTx*vs)GV+n} zq!{%DzgzYU;4md9H*yLKDEJ%_px_^P(^R>bVdk?*Cv&HgGi3CeXZkUU_hl!sTXR(( zH9o<(^4-Ix7A$MGv(|)ilQ!dd|cLqzBD%i^<1BZDb0455ZWAW0B*MqEDdal0?2*JWyZ|G(Xh)|BcwVIX~wR;`oI)Oz4HcA#bSj-J^~QYF%z`RgP8GEcLLS&FgOPe8iP~mi zVFIY8KkiL6e4q@9Qm0pafok0_s?lzZM|JJ_In%B?2l$R^guaa6pUwRq_+?aFjIY47 z0b26!aT3*fjJKXjRy+v>~_#UB>tGBdJ*iJ&rRN>FIwo@q1p zT@pxlAY;WuJ~da31+-CgkP1Dr z#bI(mMIx9{==o#E`DQl#D!rjyVU7;Zbrqv`FDQa8kn&F@ndx}V#O3mUl}zNh{*j_z zz7T}y4gQw)J1|RaPR6tuiXU3QShHKDn@;#1+z*W<^u;F(MLfS zsf$6k*O$fd9#ko!?JuVBI1Q>dbW|SNfY`FefYTR{xN$&s-4Y#+BUi103^n2qGkE0j znROIX9X@7SSgLqJTdDIVLTK&YYY@u`2Z`7f?7Y%glsm>jPWQGA+V4X*SLdDMAexdu zLE`9XgowGdO8Xlc=C zCfBx0lbIHAUIGQzN7s+O-{;rCJ>5%_*?gA1qHzUue>y>cp$>2#b&0sx9VPn6{13{=Z6uiB-jIz=V?I0>V9p~FS9`y8ZVimRf6%9{A zBXNjl?9@Tdn}U(q*GIiUmuueh4VJw#bxR;1Q{0>v4|NxAb~CGT{Qan`@qrlPSvw); zIqw!@&3^47+39k#+R##!q_29dgA7*hGm=7PUhX&@&o$Y;yUWu=No?+t40tRRFHc+j z;V|6X5#Uhsubdl;&Enm(SuPZ`kAIMdif7fFH@T^OI}Xy|PuqEc;_}=4_h#yHkBFS% zR(6r#qV@gO(!JS+hR+!wK3B59Pp547$B!So&6D{OrKX)Ay59S}+1*;BrV=;%ADv#q zyUjP(*WO&SAUbLHrEY3Z=TAe?JqgeE`>)uwe;T@hX;#2X2{_f9<)-ts&%W$7gG4$N zrkx-^lLsh=+qk*R&EC1`^2-~=_n8Sw2BZA}_}*MsfaIq|84U7(E{_aRlJVf|ZBEzi zgYl3wKLMe?K3we}r?pmgLPV_?XciaDA0okBe5Wa8;}O57{U7E%kopF7>*l(I{D8_Jw2Qj##|+t{RMP zRovJHc>y$5R#tLHlo3mvUNMzsF1Em}_<-5}>q~BnZX?P(AnclQBi@?qan)qHDg86B zHp@UnXB5Q!q-BQEPG1oGnVFOFCUH_IsNP9pYdmi|W8CRmMP_Sv?2D=Xr*?!gF);=t zs5A4wYaaXq5df+{fq>cmA7$~c&QMVWI%!Hq>Cb;_kvFiub#`sUX3cE3x0!AZ2Y`}T zBrQ0c$eV5AkT8Nr8YuJNAN>Noi(d5in%uD7f_R(HBHHxUZ{I~TOIYB0wJza{+1eMB z4jX_Xp>jsFtMZ>uI!MEDh#BQmWFuFHMxV)N<%&Mx%#rYsR!SmcQ!{2oX}JhWHJtyn znJI(x;&_tpOQVxB0N4w@<$kZWzzXpGwb~lwh)0INeR5Le@}GD6m>{)Q14j7w_}af# zyw=93)TtMS`VWPD3)!~X%EvNdXX*^(xe!Q#>C8JWQ1nyxo5{NT7B*sKYN4vigTp zidMBnov7J)+h43;3YKX*zg14toB2VJFXkm^_c?z`Yr+w&f?XPfzDDke-6TcVxcylb zlxrI;ce|v1$g15ot|q=$m+*?kR(15yhFIF*UrScc$K3v@ppnYr{#j%+cl3-uRPI-rl&wxKOHXLe%tlN?DKb>&h8tML?M4F z9l9qrGy^Kivq9LATl{vu%87!K0QoHu|ah z#muket_GdDDU(~PDs!2ZtrHdPFSYB2v5aDPiHwFbJKzoXvl5|3w^0%CgGB`|lTvnZ zFNDAeT`TePy2eMVfP!N;)2paOlfQ7*OOS1rmdLJXUGT zT3qdSoD^?#!h>|mbZ$vD_HWAoV6mPJ}bp0qX=F zeAezW%-i9=*T$y06zXuhUln%jpBMtxMq>;(mCQ+SR+9qwWO!8IR%StcU^osnJVcpA}w;p4?vwK3|gO{E7Np)%H6$79fUd zW>cBXX@)@LD0leu-t8h$g??NKzl~a`vN0jg|aVt+>{6F`Di-s)>w{Rr`S(` zVw46%?<&JoA8jw;La{@^h3R(VYkz`0m!3esw9UH9Hf`Q`GFvOm9yY8^*zZmj^+pB~ zTgVSW5X=tA8Gj5QI1-(CNV9J5lwUfualdTW{@AY_SY!c2UlaQGR=fRzaCuh?e!PkR zjQAaqMa41vujNm6ibZ!R_mk~+FQQI zSjLn{Eik7s5ew+S`6nZ@4WWuDA4aWj;lVl;`l55~qo)%-?1cZqi;~ zyq{9nkKo&?4~8!I8AF_B=p(g@Zr4Us5~&z|&OSufyyp%1a3o*mO^J9d9gLI6(Ej?7 zK1{>7<$k1p%Ys7x^{!ItfrB4O;X1p!0?t+!fb@({RBiZQ7%9%&>2iB22Ofcx=&&u< z=ly2|+uUvb{>tIJ^j!Uo@;*G)0hS#JNYEData@>R!$+Zg8ba z?OS%yiyuT`u2Oy&Ar=ZDfQ^g6X~4lsEVuzENbE~57Y;>*-DEK=KK!{;^yn{Ex#shz zn-7ZlYSeKxcI%y6Zca#JK@qvC_2|O#!bwf`z_-Ju!(6HP9Qb@Jw`}&!%|G) za<%+seCO&gF|JC|GN$AC4k=-A=%ER-S8C+<119O^K)$HmmY_HKakc!#?mWQC^jPiM zW_g_ZJ}HKN?$>+YZ__XbhQN7cvgF>HKR)=1TZ!t3{mVN>-TO>et6R??9kTc!cM;;w zO)6Gb0umE*#db0>IX6RO)Z-qFGK?KEi;XE8<83Ud_>R-x>SIP$DvNnw3X8OpkY5B6 z5gMg@yfBc8*qfh0Jd-Jj89p4NKeM9lW8{!(Gt#c8`Q}09a=T1?otzTYy>DH8j(7n- zKH9`lOYNoz-B};73?ZOn5?U9Ua7!SZkT)vXtkN$cM%tA-!mcr+J} zqjq1J%54+XU(RH9L5OEs22{2CGN`?bh?lRJYRoGVcdft9kyc!IaL_0M!DO$ApOlZD zzj=NdWOkIT<|7@C94l26Go-CpT|7i$&SGzK8d#W9>`o(T$Nv7C5?bmZv-rQJ70<8WOrj$I#m19fe3-wmgnyV3^WG-VWi(KGZ)@ zosyfbH@G8Od+;#j>n7-<>XXmoy+zTIgzS0mV^qgKR5sDd=VjSjkHDV4u<1|s5<#fv zR}&V6(L=r33X*Y%%%fgrs`h=dKDD;8bnm)I#RlROb{4@mY_uL;oPDm4`~B=yW798g zU2Y+|1u5>b3XTdt5ri&x2G&|h_=#~>-`6PmHxFqee1;HiaMpIjV!M0LWI?<8kI2DW zd$po_EXs)-lmg>GZ6wvRz)MN_an|Y!Tl0lN-DS1mQJu!@B~h-x15ww2uV?K2d)E14uVwyn>P-KWZoM_eUNe9v5C=pfC?&^Vo-spvOWh4V)QM!rD+(uU%zZvV7~ZIL|j3o z#GsvKIELHU%VKL}i6D8c9E?iY_$Dy=FPQ=uLGyOr@|fPt zeVGVS%NNfo_|q{GH+~t#?ngN{v`f0G1&{RFGK1~35uBigS zS`iJJKYdXSFgFDq3GvfsbEq>Lr(AHdJ5wL+xoUuY{?>M!4Jnr?Zv4e*4yl=;Tfc%Da2Gu#o{k)gglMOdfYh>OEZ5=eKCPv%bWTt1rVqRyK! ziLQZh(-itMw}+f{%wOnXZ5V79>6m94kunY=Wt!SaB zP6m!71~}S}H>vD?FM9Ne1sPExUmF`fW8EORE7y-NUcH4c;MIz3OQM3<> zw|UIqS|ky)PLyIKQY)9i(1p7S4bV{RZ5|d_OZx zoC#Bpor!vGgYB>-@>jHZ&lz-Qw(G&w*_Tk_kf59jhrF|!)s}v_5_`A@~yPeL~SAW|&m{~(?H7fcjfu75d z-kK}VqCYjNBSWgqCY!4CysG-%f` z!^ts|$LIYjvV>TQea&R0KC@XOfCe`1IS{{Sbca>ZzxpT^#b91t8HdR=K@Y>B*qd?2 z#!b~HMvkidLv57Mj5D?Pe&xO_~63-)%>yV z@u4O{CWOCLzNdsYPbtNI6`=w5THmt$l+LwNFiKyCcT;7J{L2m}58B^&7zY?c{B!SW zaJ3zC_juu~E`x{sAoUxc-9&~yX6ODs_!B|jO*FE!;h>Rl)LR5W>KpZOo*t&^M(I{S zBIPTrWw?o`@9?EWkiK94pkTz;V*i9U!Ysg-4dGd}rey~R__(S0B+T%2Otm*`3NZH8 zRAS@Kg9p;g-wXp^ut+Zm2r1zCku&%$0-pMEY+i#j$^b>hy!=lCbYvNjLSev-BQ-VC z&ixbUP;id^v$~UR27Kd~hS?4vK4p`ns}?{wypc1T)K7_1Mwvx}(SR+uuK;(2sCS1@-h|b7)N2#E zx`2*=Cyt%M0=MLtNqK}WEA3^Mbm-aEWA7_tcV@3fNz0;gFLS({}Y2fEsfXS1TBQ|mQt;NrDAygY-5)ldb*XMdb=eXd`pXd*b~`W zWh0-#fsmQ-%dY~^I%IPCdbiPILO7*mya4-kAvhA;U`MAgG)O%U z4ITK1Km(pY44c7NBES>)-b3IZc+IxNhgAH7NUU@d1e`CdgGbQ8qgPCThD^e)#l3t~ zk$9rK`nhErdAp2r!Dc3|mdPoh6iXui16@I)9iW7Ld18m1_o`z2F_-_73*dDyDXWTi zr^=h(d0w5m*2Q5S`_s>C{x11+hn!+tbO()VK2DRyRZ5~6Dp#J&mFzFSjENe1q760M zECb-sK%?$DV)#6i*>FFnG5DC~8HIgE5U@0{Yi4S<12#WpVLjI_aulame^Ug3&bng* zV?WIDmXZkjbPo%J?H+MB(ulfZ5HkuM(GwLvJLlCYeW9GvoRTQOhyJVUB%}It?re^~rdE;T#z_dzEgD z=MxwUI-GE0$b6#NkE0W`;Z{nf!#(@8v@f4^BfL5I@X?N#qj;}`XDlRTSYmDv+q`#| z*j37R``*8i`^Vj#iVu8KN`5Lq?^<6{tHE%<0@lD1kVD~B;^F#`SE#jCv>1o1yhwx) zlQLxfw2lRuU?fyA zM8|^F4-rcxvqpJex4a|)s#Qda09-dc83IxOZjrb#vB8;#at;Qn&c zXA;TXs=?2XCggsE4?wd{oN#5*%DWZf0c$9gW#{$;RLRc{UzdR4^u~d;DRgl@OiYoI z$Mv%)^fo+mzve8P*D^BWVS4`v+wquv82Q2E>XFON<1QG~O2S<;wZP$NAW*Y(*a-4G%f4N^4pFPfK zP_3j(_2n%NU~lwaf8;O<>ygHYNsa)Em~${dF|+bfS)uh~s-_4qm~>D69Zb-k z>X{dJ4p&0M-mi!hYYvNYu2gU+^z}fb18>o3#N2UE7`Utq3aX%3uDO$W4@FFXK|biH zKL&c-7mP2vG+?f+F%(LN+~1hYGaXvIs|@L24EHSv=zB-oerKEgo!ivl4W3S1Kk%q( zH@?At2mJH;#s5gzVFWgn%BQVG;r{e4^J)mDr1vh?v-;)LW9UMw%g+$3%0Vlz|FF+} zi82AZ80Q!Ae`gU#{D0gd)|K|)`S)UC-+ok~|230z?lP0upxQ|V*BM25spT2>BjS)8 zav+PTL+m{s&N;F|5Irpq@}J$+gKCj1-1h#xAs#XKy`?gqCqb?ERtq4gSoA zh?vL;C?@Lr>Jbo9X)lpiRdjbhW>UC5zT%ZZ{Ct;}#mqy>@DU-mDv2ftuwz{ciGlXq}3#TIi&qBUk>sPV&XPajI zLQns{{j5|{)_5pS2s6eXxy$Nzoz`klt^6Zf%`WY};@ci3%?g5O(TmUDw)sQ83SZpb zNR7P~z}}ZYx5GbgcpE-OvOd{5&sA1AQ@~_Cw?}(e?ZmxW@IRc(n%nOJ@kY0 zm#!Zqj*IUlptF^0ukYk6RgXTa@1KkNR&11}pwn$+?YP(^)egZ0x1T*-{nh2yUGO;f ze$!{3)5oXasB|4KHRJnpB__t(WfpJ7{ZD~Abf(8E5O8Zny8j|cLEOn8V_CBugJ-Ab zsBez!-nLdxX$33junC+7B+pTIpqFtcUn^-4-UT-}7e(p5%>!@Cr;s6uVbP#47 zIom3q8G}PN4wbO~AG(X+MJ>0u?qF&W7pB!hOQ^R!-;mD>gZ0^yq9&BbeLY29abCq| zSX*I3Ese>hEEWkuXWH^5C@%_n`q*UrBO5=CFQT8F+}EIDER+rG|23W~L2lL)61m<4 z;!|b6J7By>yA``Ofk%-#UntO`*b4>(q#U9=jCCL%- zBduNRh@#Eoee$AXTrH(LiTY6h)m`2kK9?|SK`zfvO>U+4kh{kYpK@gBm+X2zLSsdR z{F}=f1LlKX8B_Rv2f>`{dz3C>&knK2p^liMiKo8~-S=||Cz?y2zZ?0|VyYVQEAH>- zz>Yb6N?$b+-X-nDt|p861(BLuxXY!TLVinROe042?V8~&;H|ugxo{xJD(Yj5{dW#stDB!gB&-46QFjZ;VWt z6NLb@>`4J$7H_`2hLoDqq|%*751F$>RP7dJU4C@uPG%lIc9xl}=3Tz(j~ac}Vf|l9 zUS?8_xCLzO@rugro0ZlW#Y#;qoau!WQ8H%leu|+jjXvL8>r552fBB8ZsK=Le>x@gHoG;7b68MkufR|0uu4O{*d7o^ejr9(6ZMb zKXR_S)4!zxrhYZ+zhS% zv`{4-ihGWKof;y}R@45a_wTJ$7w_K^zqYJvcvD<1+12*Xfob0klhBbLJ&=woG0Fwk z#K0*ZbEktB+8g&eT3kK-dNLMwfA0C#4QJb3+Sz8R^Hy9tK{s zPNKYflV{?KXN$koOYwyq=I;7k1db-j}J;LU$o;84;1jw)>*y(bS z&uCX?1~Sgw*p<61c^+lVzs*!zZ%E(ZwULda%6{Cs^44~ZA+p+fD4v|dr|b1>wK4U| zeOKIbAmT{dKsA_xrik_!Gj{_B!l6URLufm)QGR5`KEq8{dLrWA1?;>%qC^5Hf`NBg zlQ3gSv#+WF?!d0$mzVZ?aLympzM0RFh^Ht#0W_1a>_;KWz5r)|A$>Ft73c#SkQoR1 z(Aa*R0Q%IxO<$P9 zVx2%s$_TW;*OM@O%+Xgrm2e`TFi4?ifk4@i_bDKt4;DN12Vlh;-IOU$^v~$Z!06I` z^u7zQYUdt>L=NlMPnVl99`0}0H-^$ki!{nG54Gle9fjj1H6_<*9-db%2zlQ}yafm%(Xtx3Z;@zB& zaVqiz{4TdoF-iHbxpEb8s`KR3YCL~!7xw2mwd9|a_C#VJtFRphD+%>oH{z*gIc6kc)&6A3r@W{ znZ!K`b#pAxSXjs;yB|&{(Lv8w$ar2(%=zp6T-9iyJZW!~*j1y2J3gqD+sv@@^2Ba) z{w?JER^URp%V87J{3m4Bbj)Mu8mr=K*NcJXL&@qutW1D$1Z__~N;`yK?eFVzm`pP8 zyYAvSS*|Q+Qs}jKqEmN=ql!-3EH-$Nj^}aC$5Rw0+!fB}?{R}7B;pwv&)6H8%q9H{+zzCyrZOUqwE6(IX(GZ{|A{mY!9 z4JE2(CPfjckTwd&&Le|qi5FK~|2?o$8|p7+Q*Bdp5F z5@qT_;}l+%WYt0laq+17l@7mmYEZUkZGT*DdYX`ISI8^4p+L1+-Osx}2T7_0?;e8^ zSY-(dKP^C24qv~`0g^evRu}QC>ZfA6V53grLa9Tk#5v#E}sW>K`m5yf^rxmwu zXY-o@=aLa*pN*c1?Jd<7^W^A@GJ4I|=y)Gg7`mLIx*o;i(aL>gG+tEP(uI=kOi zEnqjY{=NZ>y=&wh>LL_bTw}OC@$2ynH6SBzI^qEm<&CgdakN`K>*H)ZydTnnfD$gc zMW1VnI(ti00fZc=8?FA&Ryzu_TmA6JlOvf*qwbBWIA&El26{#6R-5sr(#hhBcT-MF zf=;Vf>&JA@U>fOq-r>&Jrd+1IZTq0~ace!HX3=Bw(+gW>9xJ()wuNdb3ND-d zn}Jm;5p9GEPHw5U?&CXSI@k-sx6N?@u?XpA6_ALs0Z-J=J>ASbA{Ww{Ptp-^ z(Cl00_n1B{VIJONhr$kWC$jLNpHG{*kRiFuZp|h%vzt~lgShGRb$7Wk(Qi&r`|YCLW^X1C%PTXP57Tx4~JS{JP<6MaUnBw>@7}La?Zw8$@Xa zz)HwiRE$5gdF8Pxt>LrNNLnL0m3?Ra`;9d>RF`XER;N>JsPZG&Hwsd~OxH_3)B6xcO zDvR0G?UnlGj|5ClWkS=n&$nCsu`G<}S;s(7R<8#2{D{?%)OUoxK4NYDq2rkiY0Il- zk3YR{{Mh#|e@3{7Ho3v@BG@U$vLvhBMGW*bvHP!WmaU?@h^#P$Ds3}pF-^@BlR9A0 z6)%K0%z$VfBz+o`{uif)m{STOzg$CQH>=-y$Gd1sA4>ClJ*w};HrCOCMhH1`ywlYr zqio_2l4N*~F=2c>xOBsH?0T`Fo0X~; zJyIkTMht#ta-{SAsOTg=dFv_sv1Hbiq&@v`4NuSR4;S2STdWsyo}R1+F~QqGXp_$S z^R+GwA$2~J_(*BIZ&jD)GbRoQx6u9K_ZLrpBp2yP|Sx5<@n!|T+&+vN#rhbmS(cGk2_7x)z z4F}I*a%BwT$xraXv-tf;VeKH-!!<+NP=fW{%@Z-*P*i84>4h~Jb@mQF_Ug+gIjRvm znK7-e>a527TXiyT6Enr@t$`rI?7rLeC{-D^NI4l2vFeE9WhSete4!Rfc*1{R3iGJp zwBDP<`>CzXB5}hB&?98Y&Nro)xOd`K5vBSEYRH zNEn-7QxCiPF5-xIUB7uhkW_g2tt$3i2rJZadVEx7g&!e-IaWofr%z|vg4_Y)p2KU%7ukd7?{0y$TVXli#MR2R%a?tD`ntcT4UX}x@D*LC#*3b(k>LXRw{;LjIl)E_ft z^d_6-_B~%q#WuM;)PlxV?2}N$^Dq924MM~39P9NrySYX#wgC9Pg*o3dQ#(-T*LUK; z#`H2=S5^%FG$=91GJD;Mn58u4gw*TD$-B!^n143-QVt$t}Z9TRY@^w;< z>*xEd!>(e0GCYEDjaqU4z*)&(Iz5sS{81R3Wv00+`Q)t4J;WKqV|>Fc#Ih6!=CT`l zTZm62)&Zfz^Oh`-h5x&kAlC~7n$<~(6r>`kf6R_ZglJUcis2hr|Ed-doI2URzB!Xs z=Q6q49)bjdO&I8r#NBPEjf;$sa2ziue%{zO*dBI(g3ZW{BgD>$r#tgDZ6JZ=^bPQ2 zNTHB&*9NKsFqT?8kX+EU0eggXLNl>148keJOh>+EM_OfTRS<>N8IB*c?Wddio@k+( z1P$l@cN{2j4U5!6!CyBMh@*lZxkUz!%7-P)lQ=U}Zbq#T_@R!E)SzTmH_4~$7$Nybrkh3j#_4JDS z)B%XEMulazA}*-eHJUtEn`k`+I?yw}ir!O^#1A~b&|>eXHt5DIOBmnKos^%H<$P*c zT&~{XhrzUAU1jaPwoCx9Zus~Ye{33)Mwm!9afCQ*nq>oO5&a-`cZami8O@c7rPv@w zU1K4?nEJRi4or`1`)Hi|`--|L7rkOIIWC9@5A;-PP^L%fWl&{Q)2hEQ%_A8byB%qM z3Ay*M4~G1ou{DsZDEQ0)l+U3trjlA<;{*y`Bbk~F_^=LY5aLJrlT^y^Yoi2_7S6TV zqMa?GnA*y5qi5yX#qn`A{EL2~n3JR)KCw=3x!~FT-$@ z09w#Gf@`?%=LM$0aqLbXOcM1y?{708&kfR;FtQT0@9+5irSTO;09X$Lo&+$CMh2h? z184X%0QrDl@xvr74fyhQ!4(>R*b7ry>4n4>&P z^gtAo&UU~w!AI+5|NDi9gOWatwI**_cbiA9jbo+16Hq-ZgkYh-A627%^~F<~f(`|N z;otAkSBR|>h5~uvV&s(rI#Hy`BNapvay(KLS!CzSa6C;NwI326Y%(CkeI74gV&qj#_T_08GU*SlrXPFNXnRWs6hKT$8}G<|W8epo>U0x^8}sz+XBx2*;A_ zr&uE*K(obNt#@^92M}11^*hJwXM07sL+qBZ=gV@3+Ri^T`7^UtwN(WpekUT(P%{J^ zBs$_KE`Nz-Q7&V~sk6eF@&mZdK8Bi&#Fu3VzHXCs7L*wXkc{Mud4r^_G7VCmAD{3AUc*Cy`W&y} zVs11cA6wsiFOrpo8v zU!Nce4kJp#s)0LDz)7(v*ujgF>!E=MQcj0Jm`n2{NV4(Mi{hB<<>K76J2|-e(LB0H zq~F=C+1sHBZGAuU9Zi_Iy*)SuH05+1)TxZavfIO(3d&mTv?iUORk0T`qeZbn8{*RZ z;5?qs>N&1kWOyCMDXp6EY#Da+%bKx)3!UljWlpWmIy|y*L7cXawka8RLk*4CFz*qQ z_nkS;(pOIjKq`Ox5UUe7c6EA^f8EBQBN zb6+nj+a_l75SIatcsH(d6nh=fW{By$k#92^(e^L2NaE+Q$Az@xA@#0+Ujr-%(B22N zesCj!&RkB1igXKnhT0qDGOU>hd73kmPBsGS{W8hQUjC}MRQ>#Ltiy-_O*jbsPt3V0c}C%>aC^3tCA9^%Zi>g&?$g~Uk-Z@7ex z;U1`|?Fn~0BTD2`ppVXnS-Rq-b6lN_9z1)|flT51&T4MD!dNuWQn~&NKVQ*oC4dCg z3ZUi|DWzX+ZIAi%DuH^KP%1t?0RkRDX)TTIL0*yhVu$m88M{<3#HGUn$MES-ZPhE? zp2XW-NjC5lQWUrKd8`h9vRrQPg2JY{GBDsFooPIvo!_QcDB^;CL}FTbgPEBF#nS|QHXKtuAk4Cw{j!G zy|^)PZjaPkGsV9n4IX|J~DK&ttwNr#a-k zZl(DT$2^*V$lXhJe|I83ip7}Dz!78sveb@^V2Gv`PhYHus7`xc!Sh;0E9xCb8$I~m z>1TzD_e2B$Ib^4{wpa%X(u}0}FyHI=>Ap`sZLaFuh`)fV2X8#6TS-DlkYork?&;;lj+QYR(ebJknB#eKnc_%BpC7s2K9aRG$#MH+1@}6ldG^HiR5F~9 z0D8FiOr(Td{v|tTvBk45E2v3Vr_rAJ;jYrF|B^39i6K6L$YdzF>kUOJMy;*6u-VGj zvsQeAKB#dd=-UdfvG^2#$#A6i3I1u9?nSi!FyDyxyEO3b;jH!iW0qwfVWAAuN98{G z6!l!y8gtxwlfD)j+K%Sn7Xdl3AqBC7H8CGU&%LI6*Dx%*jFFpk0q)kMtDy_mVQ}kNngI8nq;f-n2t`kfg<_4~o}QoZ5XcPGM}#4t3DYj4 z$Xah>Cp&Vw>b#Zu;gqOQ9r=RXdFBqM_ZOA;pJS<&9uWT93%Lm4BTW*)1FZVPqee;V z3+if-R3oz)!=v1|P)r=~N2R&M;Ih%-Q`G5>eXq6fxiF9U(?kG?`J{Jr2b#)vP#2X@ zr{I|O0*3N>Nv^+UUn-SxM#L|U>s2dj_#ywKybHl{8-{q61+W;fxP$=&?NKo_Ynd zk2fm%DNrHX_eC&^wz+ugV zoI}Jg8)l@wBrwwW#o+?|8IdPj%GB^||MwW?R}0#gogH+aGnz3K<2yy!Zizv&SS9b{ z`V+29`v{3%!HkK(*4@(X`5s7YHj7%b`lf!LIV){C<<>--z@F(L|2uD2LZ3>B_@brp z_BK85In%jok&0nvjdD)O%XI6&9yEjetIv+xd&;rK5du<^WIXLM;w98+m?IeuGR%Ys z!_bjn7(NsLOrtx<3~0OP>2FWM`z-J9rzl3?A7c3qcEtq1hE`^8g6P8p8+T$!u_I;g z>KRFGWtjQ)<)!DGAY@9z=&*Gq`Sc&-@IL%mM;%7;bTaWovuC65KJFby9nOC~(gPpu zzolVEqG(d#(ZcZG81oms2g*NHME!_?qX$!|DEVP3zW8-z0*{KvVK5AS1H&hj{kz>Fl-GmNH# zts9^N$7K*86N{<%BMe^ff&?Vp_Wq4RAU#c>j2IRdnDN_I4AA=j(q-y{1m@eHk#25u z1mwEMZMaM#+bL_5&j{Im3hVnYv!?K1mPQmZLERQylIrX#c}n#)6u@wQ)lhWP79hX8 zFEKbK_mp-sSFZ49<|HTmNBLA<@nZ|UMk!?KEByYZ&;%dG;u+i66s=mcTyFtoi;LxE zF5_X(LjCH%RQG;H0Y)Z&H)<}=0|=wW)|2`_YgJT}>}E{z8E>VrAI?*Fias%$JqzHt z=&#*3=?hWgavtYr1nTIPStOodj||S8<(M7kI}*?kDm(>bbASTw8`=i;6J4R&gmY`{ zZ3Vb`A;~d9)7mf!4bUXK`>!@FXru6xfPiSA?SFmw&inm0Y(?)Xx`p@uj&k|tT~H4Z zp%mV~-LB9d3xcoulC3cQ>vw>@tD%p`LHe4+AO7vOLH~3Iu|aQH#rp5)-`JlxK|Bo& zY5wgNmGSZyEGDC*5Bguf|M$TEA3Z4lVfg1tqX0Re(6#p+QbRug*X^i&`(r2}YtFmx zPdH6)@p!FfDUVwpT5eAV*uDEcbg7kUsqIadro@)8@&I+NlZ0H$Ztri7U7W_@e+|&y zOgTFNld6m$&Qlz6K^lB|g;)SrA-Gu`5c-PoO{rGx&j|&gp$G-yM(|OEQ`bUts(qvSSvi-@z^-#;E zUdr@+RjY6jOfQn_8r8@c1;GzmP^;Zr9oRIa?8GzQJr&UJj+cDq-WDN)O8EqO2***= zL(|_{#se$bXc;CojBCVV6;fH6oEEt^74s2bGvsJjB~4q;(g5U*un-EjjcAoJtE3(0 z9h81Zrhs~>R)7#=VJm>W5ZrAB&>1=X@_F}oE5B> z0b~cI_48#ERBWdrAsKCeq%mH21eAP#I)0>T<#oNAkE&AV08O1I%PB?0dcC1~QDF#s z0u07|J6un8xKO9qa0D2+P+PViLY(XxfThauAUS(0$?fh zV^(cS8STdVM?j@+t_qjNS$Ci+`S+>fSqai%e|R*t^mFwKUUve>*kVznK8!2$y) zN%p3~1*m*J&n9TYFQd)e&2OW-+}Tv)+!H~L0$Z4x4%t*_I0A~2yBlI0G?=C+d`>0j z7j=R-JhHXn0ynxGe!R~6uz|QGApuGH2x(7miiHQ&7yFUqlkfgI}L!YX!ABk$9a*J|+frkkAhm3X%+%3L-PWJUkV1%AAD05{f{y7A$&UT(j*DBUMWoR zkWEN2xxI-oS!r#Vm;uU;2ARDfPZu3%ZgcsM8R|_yYr4KpIo&F#Y_Q%Y6xEi(_K!|9%AM-y`@|*hu+o$C&s^^rb-!$44xC6Y#z#sO6|7Yl=Z*A3gsQo#8pR$YQ&s^j-jvuac)l0OD*^qO+x|$76e`L#8 zo=`QcejK3Udag$GccKA$M-icc@2q^1xDMB2i^~9kX%kN~=K+m3+RR9R< zzb81#17Pk%Zwx4)`1d+T8ULN}h~E?br_DDse{$e#4eMSa|L1HS!rpZYfw|2_V$z#4 zhB;rd{AcccgexQZ?|kjv115E*;x#o8b4r}Q28}44aC941%hwIClNulQlH^xszCW6HZvby;BLoYoev;beE zW-?*;BO&{PihB7917EF0^svQ*%FI!-hiP}}l}hUuECFQfr*RdtYo_2{MCE#}P7=n( zc<4Z0`8>}QkJmI=I}|&7Z}Q`(@ZI7soAb_9CS6QPg665Om^3*wn1>-j`5(@B9j<6M zH@+X$Z80giDPMkKZ=#XaD2aE~r${B5+O|le7)qvBEIJQ|oKQFn&jI<+y1Q(*LB|&Z zwdA&OM`=PbXWcV7jn_KmPe`4+&7#30>#18CS)T7GKE+J-4GiopR|;t zeHjh6f5a@21vtefn9^#1oBM`rF|h=mhw+b_`$j&nryUP)irZd9GmxaeB!MXtaDgMf1xTb~D#~C+`@S}Ts5}~? z2U=4ob>K+`(zC?dP5Axe6#t}y;nP|E-USv6SmhNmfVr-CN&-Tycu(^&kY1wOP8INm zsST)b>0$VI;yJ&blcc|Tch(2AMvbkIh=!ww(k}i1-26a|L}2aU-}l<4;h)t23$Pj> zW5MNtosExbm|hkrS9z~XU0rQwY*vL3Y)&hr<%lovqSsnX6jp07AS~T7^+oIXzjc-S zxd9~2=w#GvxmC-_Ejoz)7zX6oC_WrbvENU z%EN@>n8`popL_Tp_iOVjqdBg*t0@hYwMEBGTMa!%FOY|-LU3j8e1~M{=1G3~&-Dl) zg=HTPr@`|Ex{Xfw?!($U5`iyaRtuHl3Jj_;Y>j_Z%qNQCcr3?fHL!XW(?xV!G#gZh zMUTO?E&W{%pHg@orsD2HpLl~E#O3tuO;hTP?G zYf8nI8XPXj>&%Kzdr(CCshd3pP(S5cc%*AI{Z=j4Qd=s1|s%DmOL}&)oOGZbgC5iMH-cH@Q9;Yqk7HC!XieTS4$bgO) zH1Zf4tu-15g${X^Mi!TAi*@vJt5}O*j$1T>GA$CEp!9x1`hat@xhxHe224<&qt~@e zmn?hVn~QoanR7#D#U7yu;Vt4TMDxj<#7|R5JexLKWKo0tWt{N~(9IEq>&D7wm-S7O7iqR z3Wc~lzD3G_SieV&Y2~SDk&`H>Jx3-}Cv3^wLJN@_*@}OrMQ=y(6AzZE1fu&tIWMh}M@H`JWa!e7=Pt^9 z`ee&nicATM&KJm!*in6=Zzw$E^t=z>rg@7Ms#IB=WU*}0dPIl7xGe@XckKpP9elOa z8Uam|it5~^G197CN>|7wHsLv#qSX81Xi245**I9@&3Ku!<&xo=h$L#W+rIjpr>;t- zr|L3kfqG*Ei%yfK^aJz()z)EgUzO~y@Ay#?iGo;<6#eM6*`}Lw9l6P@@j}o`G_`oT zh&41e<{O3mt=)lI1S|(fbn;y9em7qQdLXGxu1|dX^N&&l*3o+M$8)DmvY})l4fWE` z+2nA(Anf3Nx)(I-2mltRqUEyl%cn$Xxy|QRykTjPR(lL^aDNuP*w*b=*Nye8oKQ?? z-ZQn`A+_xd$dw()r|_boe>L3T+py~=XIN>45>#w)IA3Kn#k9NF(tJPWPkr^h1MO~_ zFyOGVhuktGI2~UAyyA1wT=Pj`lP8xn0;tW>?~yGndA9Jqd!{b7ko${FY3rSQ5o6B| z0uJN)x}$V8;&Y#UEP6Bdz6ID&i$5>2;?LpFYq?0DEOf5LPVosQ1M#cxgz}c_7*tAirx2_Wqsa@G^lZ$O7u{+Fs7% z!K={AS)LjsoIe2<%l|nhy}oY~_UCYdJqwghcGrAzI)H2XU8asrw<`7wrlBvXz@o#S zC7920;^MW}Be@-b*`Ig^58^aE^bGE3_7Ej7pU;wy7`V2$6D&<|-w(U3akY&!`GNKD|hd z)~z)7KCQ;7KE|ITPsS}ZurgPw9e+6kkPuH4Mb0&O?t6qi-V`}bTyTz0L7zmA)5=Os z$m{l{FSs_*IvjgN(xg!wDwDfOaGFQQW%z)Gf^ION?;817v6_;t~9;jx;%D*wKo;o?bGO79Pvrc>M3u{eiPJCJPqnR)z5MQ4dZJvIM-Q% z(w-4rZYe9&vVCfPDS0h?;*ZT7)T1%|2X|d6toNtLMgWmh3&?dbnzqHJR7THo1(>hPWx9gTg>}|M=w_A%+d%r4^^=v z%S$Sw9X6_R%JphwW+J^OCxg#KGu%J+tjt!KcJ7psrL%BBoAimTJA@zLmnSV(4J&T? zIkGbgUK2`sQv>Xg4D|G;Rv99=T@t%yL*?m6*@QiJww9f&-U}gDsR7uXTA9`7G^p`m z^4BH@6Y4!lt|~A2Y47Xko1=j|_9NXI!}~CjvD?wd7ouxZ7*6q!?3EvUQhPTlmEQG7 zp`&o+=Z_H(J8d58*H@?2BUL2Na9-Qc8B^Y1qP`zHyr#OBUTK-H2!Wr()t=GfWB0?w zQs$TJ#7!=l%S|Q*_}v}zpR8+N0 zwXe?Rt?U4WVov{E| zL%lf98Q}pPlznAo{0CNTNk+ElG3s;ek3Qx^O>|rmm`z>b!;&O{J%SucmJB z)K*tQ5AxRB-*#42KGrU8&YO4+J_{QT^+r)xS;r9pBz;XkGSk?>HY>`XhMck_?N zA4SCI14~>~QFeyFo{^Vy*f`GxwH&Q8#)8*}oI~CH(28Nk`}wGm%4beBWtU^;(A{R`Y^T@D5z&DW zScf4!PLP3$aya7Kvg!x&o4SmKM>gqI@m1W+7>;x>#}hP*_ss86_qm=0$8W zN+RnJGX^1Q6_67l%V3#)vc~BxzKif5ge$LaMQaN>_kYnmb9o_+$|pf52xcvBWLqO= zYhPfxN33I%lVhUa-!7jw;k>V!sXs)pB=NC-9Ptvs>~OJpzL!EBjQbuZ2+KAB3Ze^( zm~d%pbD+YG7-@ytMsC&bLl>OD-mm2K9Mxk!IKMuK`0XSBCVug)%+4x|?VRNNT>}Xo zhv@maoo{r{OJqRb=`esf{DgAwYH_zWRfjBGCU`@H(pZ@}J$$ z8>4Le!BN%ZMZiTZptyUqSGO%f2@)M9_7Y4E9G1O-4jjx*pvHNoy;{LZ)%Z@Lx zN1#*1PQsVb=kbx2O<*w`{#ox`;9}02EiNnzXf0RyIY2llFdae7n<mF1Qpi2r+rkm5+UhS4XxF@pzB?+B&n?1|M@-|8NFwSNj>Er3lkrH<*-PvhvR(zoM4di4q4El1>!SKd`+} z;X?q|h98J`cM60Hje}5QL=0IV9p{lPobj1XoHYh(cEhiNdOSZo^0K~*LHB#xF20vW zoFb>EO~*$DYV`CU!ig59nr6laE~$HaWw*cLs^$GE+k>0%t4&^)k51#{Hki2@7h`MUn+Rm0SN3G_4Jk5E2k)Ma1vnZU}-$ zAMdB=thD&)!HmPS*t5Nx8-|N?7m8}-l6d7?lq|oiGUyKrwD!FOn&!#1o&_YR$Le~V zck0Asbr7f_uc-_*gL-yabyhU+fF6FUw|HuhHeuG6#lqH z+~KIGvNA#ZwDyfREX(^_C$-+^r)#`_HNY@I*JN8IIa0UeT&K4$d57rbl%`(`^?M0G3m_dEbxI|(x2*bV8QHFr1_i-V z^(XjrAB(ZcJ!jxoR_)s=;-&bxsrb9dUKpM1pGD-kkjZRsbd`GPTh*$KD`1ISF!4KB zPqz2?p6q-JZFj!8k%3V9Xvd|vFKeXj;%RME+I7gsFJAn#vai|jN@2!a2FkD))rS%~ z)|)bzW5;fnL^E2SS%#+-u&~msD3{+J&eGJZFWnykhK4{N&Q2?|G{`RN|% zuSB=WKJ$CT2Zy*Rd?5$LT!j={?jlQXr{U%3E+=?Z z!{wsl^GnoPapo=Aq}gr(QH?->%wU`**i1pEnIqN+XKKzp1x|2NEhi*AX9vhti1|I? zJbTC)uLTpj@a|;SS9F2&gf3qg#%xdyi1+0&Gkyn#!-J=It66-n=3$tO2D^`1!C%7D zFQg-<#ozaVLlY-9G~)Gjm@;y2;{5o}o)HBa8THMcB(D3LO-8`9Br0eF+Z^UIo^~M; zMP6pSUua0?XgT$7*_!u==t2C-@RjEjTS^&OB&}>|!1RTSr6r0!ogUQ|EyMC47$EEe zsz2`FAO16WkjiE`os>jX)YpgiY{nL;@Kyfgh_T_XYKz%fs%e_9Dro#5rB!pn$H=64K?@eAUa zroWbAOLA1g#u0h&SEwM}QkZo%o<$MINL01xS%6pYgMaPAXCXoaPW(Hbwk^APEY=*M z^KOUcE!RJgeXR58MB-YBVIh4pyys}fL-@z3k7rq?=}Yp?!Z)K*PA3F%0t9H61Dkpy((d6~9mXz30WrgFi`N zO!R35f{oKJr61riiK(3oO(q)G6h8r)ePJLA6-#ayBB8K<$LDw`3~ZMf1V5i8-Dh+g zKs5ecvG{yoxR_TG7a|&?ZicKp$B?&z*n+!j3Au z2x)g=suA!*6NpX;XVHgR_`$M}m83|mxdZ>ERkqrs&f)d@+S;e*97P_7weCM8ngt#d z4ViPi4RX9v9OTZw0y-8e z+Fz+4<%~`X!D}Eca5^daP?+$}$FVnR2gErmYS}O^(n-c%Oy!#~Nr*j5-9si%Fr{0z z*tIo&mD)dzx8ik?7R#Wahtpbst*hF@wBq&9yzIEU?T7l7z>0k;65fZwq>lG?M6SEg zn*LoojCVHwYkNn4k*)Hhn-TpA%J~Q%Y%+$~YH0!aG5&+>Pg$LAWBZWrs?hj?sfy@9 zWhdT1u}Pft+nbe#cq&+6pMPJBWa)0SDNZ4gzWe7R<(tOBI(@gjU_Gt4IA)DLgIeou z%c9c$tC$3R;+Kk&=Ce%!au1j5i7@v{u6(W>0&~t!i`L{;8pR?WOje`gx*L#}>c`^WbDu=R`ja!MX9-Pj$jyqq!fh;O zrr8;_b+V{fYt}VPiO6RA*O$=o|EQD9o?)Uoy@ko)izKqJ&O-?{FMDUY{gWt zWODJ{YK45yqqj8yhcp#xm`CVj#7sg*m+MZm_e8cb`)L=;3AOt(mi7na0WEr^ylf$? zA%tUh)7qPzZ~8$NXs22IYB}HL)5V4+CSjq(&-Z_gS2hPYSZv&T98F(wEc$KtkwJzl zOmWUzVcq})_WVfM6dF)z&t-a$sN(njJ*r96p5++*QYJK7 zIp%!x>)QiNFAlOmF1_0x$IqwhzT#SL966>U=xUdiij&8f$4mQZ^D~B~kP1kmzQ+N- z$9obwL!%)gxtF!3hz|{)w>xcInc%;OL6hNia-Qub`ym`6mgg!3UrO=IOS}#*`YL2e zKYfP0+%#?`$26a9_@)<$3>MgEht8c6-y_O|{wPI)KATv%;isP-HQFCrH{O+D;Jjfc zR)q~G3IZ3$b*;5f`i}T|@1PD8-)EdL(TDh@S8bPLPMV&kslQ^d|D=LiiqMP&ZQC@S z(`_#P1fxgmp{KjUAD5>4`y|?@4e>=3IkAC6usKSu$HiitU6T74> zbI|pZrg8Gv5c_CjU(Sc6eLiZBKM8CW_aScJ8t^6ljM;iq>Tmmp3YH!9@f^+KOsUDu z@QGp5deho=@8`WCq3Q2fNM7qXCW6VlZvl!dWU?%~m%r#(jqzCMLb`=QUa=hvSge*} z#ds@z5{gJ*(`C}bRWKLh0_Pd|iIyp<*#;2lD9zJIjzHyj#_@kE)pAc`?}LRp~ zK5dqiVPo!@u8_A~q!laJyS-2I3%+^zSnoP&vnz7cLC5u9KY~Of}?0j_W(?yEf&z z)aOv?YJ3WLUp@;}x=pRdUuHic1|5vwmnOfYI4;N{M)&xnBKPyTZCgr$pjebo1t{ zJFk|c4It`F!t3PH_nYi4GT>EYP)Ce9ilxJW{E&L^pYehcE7$MTdn^n<;OXk;vd$@?2>?8t6Wag) literal 0 HcmV?d00001 From 9d519463c6c66a61b336eba143a86732524bb5f8 Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Fri, 6 Mar 2026 10:41:58 -0500 Subject: [PATCH 2/7] save Signed-off-by: Andrew Duffy --- proposed/0003-patches-format.md | 177 +++++++++++++++++++++----------- 1 file changed, 118 insertions(+), 59 deletions(-) diff --git a/proposed/0003-patches-format.md b/proposed/0003-patches-format.md index 624c768..488bec5 100644 --- a/proposed/0003-patches-format.md +++ b/proposed/0003-patches-format.md @@ -1,4 +1,3 @@ - ## Summary Make a backwards compatible change to the serialization format for `Patches` used by the FastLanes-derived encodings: @@ -9,82 +8,142 @@ Make a backwards compatible change to the serialization format for `Patches` use * ALP * ALP-RD +--- + +## Data Layout + +Patches have a new layout, influenced by the G-ALP paper from CWI. + +The key insight of the paper is that instead of holding the patches sorted by their global offset, instead + +- Group patches into 1024-element chunks +- Further group the patches within each chunk by their "lanes", where the lane is w/e the lane of the underlying operation you're patching over aligns to + +For example, let's say that we have an array of 5,000 elements, with 32 lanes. +- We'd have $\left\lceil\frac{5,000}{1024}\right\rceil = 5$ chunks, each chunk has 32 lanes. Each lane can have up to 32 patch values +- Indices and values are aligned. Indices are indices within a chunk, so they can be stored as u16. Values are whatever the underlying values type is. + +```text + + chunk 0 chunk 0 chunk 0 chunk 0 chunk 0 chunk 0 + lane 0 lane 1 lane 2 lane 3 lane 4 lane 5 + ┌────────────┬────────────┬────────────┬────────────┬────────────┬────────────┐ +lane_offsets │ 0 │ 0 │ 2 │ 2 │ 3 │ 5 │ ... + └─────┬──────┴─────┬──────┴─────┬──────┴──────┬─────┴──────┬─────┴──────┬─────┘ + │ │ │ │ │ │ + │ │ │ │ │ │ + ┌─────┴────────────┘ └──────┬──────┘ ┌──────┘ └─────┐ + │ │ │ │ + │ │ │ │ + │ │ │ │ + ▼────────────┬────────────┬────────────▼────────────▼────────────┬────────────▼ + indices │ │ │ │ │ │ │ + │ │ │ │ │ │ │ + ├────────────┼────────────┼────────────┼────────────┼────────────┼────────────┤ + values │ │ │ │ │ │ │ + │ │ │ │ │ │ │ + └────────────┴────────────┴────────────┴────────────┴────────────┴────────────┘ +``` -## Motivation - -The existing patching mechanism is not data parallel. Rather, retrieving a single patch requires either a binary search or linear scan within a chunk of values. As part of the push to implement first-class CUDA support for Vortex, we need all encodings to have fully data parallel decoding operations. - - -## Background - CPU Patching +This layout has a few benefits +- For GPU operations, each warp handles a single chunk, and each thread handles a single lane. Through the `lane_offsets`, each thread of execution can have quick random access to an iterator of values +- Patches can be trivially sliced to a specific chunk range simply by slicing into the `lane_offsets` +- Bulk operations can be executed efficiently per-chunk by loading all patches for a chunk and applying them in a loop, as before +- Point lookups are still efficient. Convert the target index into the chunk/lane, then do a linear scan for the index. There will be at most `1024 / N_LANES` patches, which in our current implementation is 64. A linear search with loop unrolling should be able to execute this extremely fast on hardware with SIMD registers. + +--- + +## Array Structure + +```rust +/// An array that partially "patches" another array with new values. +/// +/// Patched arrays implement the set of nodes that do this instead here...I think? +#[derive(Debug, Clone)] +pub struct PatchedArray { + /// The inner array that is being patched. This is the zeroth child. + pub(super) inner: ArrayRef, + + /// Number of 1024-element chunks. Pre-computed for convenience. + pub(super) n_chunks: usize, + + /// Number of lanes the patch indices and values have been split into. Each of the `n_chunks` + /// of 1024 values is split into `n_lanes` lanes horizontally, each lane having 1024 / n_lanes + /// values that might be patched. + pub(super) n_lanes: usize, + + /// Offset into the first chunk + pub(super) offset: usize, + /// Total length. + pub(super) len: usize, + + /// lane offsets. The PType of these MUST be u32 + pub(super) lane_offsets: BufferHandle, + /// indices within a 1024-element chunk. The PType of these MUST be u16 + pub(super) indices: BufferHandle, + /// patch values corresponding to the indices. The ptype is specified by `values_ptype`. + pub(super) values: BufferHandle, + /// PType of the scalars in `values`. Can be any native type. + pub(super) values_ptype: PType, + + pub(super) stats_set: ArrayStats, +} +``` -The classic exception patching mechanism implemented by Vortex is optimized for super-scalar CPU execution: +The PatchedArray holds two buffer handles and 3 buffers for access. -1. Iterate +--- -We can achieve speedup in the patching step by unrolling the loop by a certain amount, but roughly we benefit by having them all accessed together. +# Operations -## Background - GPU Execution +## Slicing -GPU execution requires us to break down our decoding operation into a set of **thread blocks**, where each block has some number of threads. Every pack of 32 threads is a **warp**. Warp execution is what makes GPUs so different from CPUs. In CPU programming, when you have multiple threads they are all running independently. On a GPU, every thread within a warp executes **in lockstep**, executing the exact same instructions at the same time. All memory accesses made by a warp in a given cycle are coalesced by the **SM** that the warp is running on. There are many warps executing on a single SM at once, and several dozens or hundreds of SM. This is what GPUs do: they help you write code to exploit large amounts of High-Bandwidth Memory (HBM) in parallel across many tasks. +We look at the slice indices, align them to chunk boundaries, then slice both the child and the patches to chunk boundaries, and preserve the offset + len to apply the final intra-chunk slice at execution time. -In GPU land, all code that runs on devices is triggered by **kernels**. Launching a kernel involves copying all of the arguments from the host stackframe into GPU memory. It can take tens or sometimes hundreds of µseconds to launch a kernel. Any data pointed to by the kernel arguments must have been copied to the GPU before launch as well. +## Filter / Take Execution -## Background - G-ALP +Filter / Take operations can arbitrarily break and reconstruct new chunks, so they cannot be done metadata-only and thus must be a Kernel rather than a Reduce rule. -G-ALP was published in 2025, and its main contribution is to come up with a data-parallel layout for the exceptions for ALP decoding. +In practice, we perform the operation by +- Pushing the filter into the child +- Intersecting the filter with our patches, ideally in a chunk-at-a-time way so we can write a vectorized version. -![G-ALP Figure 1](../static/galp-fig1.png) -The crux of the model is +## ScalarFns -1. Split each sequence of 1024 values into a chunk -2. Reorient it into 32 lanes with 32 rows. Note that this aligns to the FastLanes lane count for 32-bit types. +We do not reduce any ScalarFns through the operation, instead they only run at execution time. +This matches the current behavior of BitPackedArrays. -Inside of our patching kernel, we can implement this instead -## Changes to Vortex +```text + chunk 0 chunk 0 chunk 0 chunk 0 chunk 0 chunk 0 + lane 0 lane 1 lane 2 lane 3 lane 4 lane 5 + ┌────────────┬────────────┬────────────┬────────────┬────────────┬────────────┐ +lane_offsets │ 0 │ 0 │ 2 │ 2 │ 3 │ 5 │ ... + └─────┬──────┴─────┬──────┴─────┬──────┴──────┬─────┴──────┬─────┴──────┬─────┘ + │ │ │ │ │ │ + │ │ │ │ │ │ + ┌─────┴────────────┘ └──────┬──────┘ ┌──────┘ └─────┐ + │ │ │ │ + │ │ │ │ + │ │ │ │ + ▼────────────┬────────────┬────────────▼────────────▼────────────┬────────────▼ + indices │ │ │ │ │ │ │ + │ │ │ │ │ │ │ + ├────────────┼────────────┼────────────┼────────────┼────────────┼────────────┤ + values │ │ │ │ │ │ │ + │ │ │ │ │ │ │ + └────────────┴────────────┴────────────┴────────────┴────────────┴────────────┘ -Vortex implements some of these instead +``` -Let's look at the old kernel for unpacking a packed 3-bit representation into `u8`: +--- +## Compatibility +BitPackedArray and ALPArray both hold a `Patches` internally, which we'd like to replace by wrapping them in a `PatchedArray`. -```c++ -__device__ void _bit_unpack_8_3bw_32t(const uint8_t *__restrict in, uint8_t *__restrict out, uint8_t reference, int thread_idx) { - __shared__ uint8_t shared_out[1024]; - _bit_unpack_8_3bw_lane(in, shared_out, reference, thread_idx * 4 + 0); - _bit_unpack_8_3bw_lane(in, shared_out, reference, thread_idx * 4 + 1); - _bit_unpack_8_3bw_lane(in, shared_out, reference, thread_idx * 4 + 2); - _bit_unpack_8_3bw_lane(in, shared_out, reference, thread_idx * 4 + 3); - for (int i = 0; i < 32; i++) { - auto idx = i * 32 + thread_idx; - out[idx] = shared_out[idx]; - } -} -``` +To do this without breaking backward compatibility, we modify the `VTable::build` function to return `ArrayRef`. This makes it easy to do encoding migrations on read in the future. The alternative is adding a new BitPackedArray and ALPArray that gets migrated to on write. -Contrast that with the slightly updated new version: - -```c++ -__device__ void _bit_unpack_8_0bw_32t(const uint8_t *__restrict in, uint8_t *__restrict out, uint8_t reference, int thread_idx, GPUPatches& patches) { - __shared__ uint8_t shared_out[1024]; - _bit_unpack_8_0bw_lane(in, shared_out, reference, thread_idx * 4 + 0); - _bit_unpack_8_0bw_lane(in, shared_out, reference, thread_idx * 4 + 1); - _bit_unpack_8_0bw_lane(in, shared_out, reference, thread_idx * 4 + 2); - _bit_unpack_8_0bw_lane(in, shared_out, reference, thread_idx * 4 + 3); - __syncwarp(); - PatchesCursor cursor(patches, blockIdx.x, thread_idx, 32); - auto patch = cursor.next(); - for (int i = 0; i < 32; i++) { - auto idx = i * 32 + thread_idx; - if (idx == patch.index) { - out[idx] = patch.value; - patch = cursor.next(); - } else { - out[idx] = shared_out[idx]; - } - } -} -``` \ No newline at end of file +This requires executing the Patches at read time. From scanning a handful of our tables, this is unlikely to cause any issues as patches are generally not compressed. We only apply constant compression for patch values, and I would expect that to be rare in practice. \ No newline at end of file From deb2e49f8f5d5194adfc6637c111c3002d4241a6 Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Fri, 6 Mar 2026 10:42:37 -0500 Subject: [PATCH 3/7] rename Signed-off-by: Andrew Duffy --- proposed/{0003-patches-format.md => 0027-patches-format.md} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename proposed/{0003-patches-format.md => 0027-patches-format.md} (100%) diff --git a/proposed/0003-patches-format.md b/proposed/0027-patches-format.md similarity index 100% rename from proposed/0003-patches-format.md rename to proposed/0027-patches-format.md From 3419293ee3abf1afb70939f4ed3e8f4938331575 Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Fri, 6 Mar 2026 10:43:34 -0500 Subject: [PATCH 4/7] link Signed-off-by: Andrew Duffy --- proposed/0027-patches-format.md | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/proposed/0027-patches-format.md b/proposed/0027-patches-format.md index 488bec5..e845206 100644 --- a/proposed/0027-patches-format.md +++ b/proposed/0027-patches-format.md @@ -3,8 +3,6 @@ Make a backwards compatible change to the serialization format for `Patches` used by the FastLanes-derived encodings: * BitPacked -* Delta -* RLE * ALP * ALP-RD @@ -12,7 +10,7 @@ Make a backwards compatible change to the serialization format for `Patches` use ## Data Layout -Patches have a new layout, influenced by the G-ALP paper from CWI. +Patches have a new layout, influenced by the [G-ALP paper](https://ir.cwi.nl/pub/35205/35205.pdf) from CWI. The key insight of the paper is that instead of holding the patches sorted by their global offset, instead @@ -146,4 +144,4 @@ BitPackedArray and ALPArray both hold a `Patches` internally, which we'd like to To do this without breaking backward compatibility, we modify the `VTable::build` function to return `ArrayRef`. This makes it easy to do encoding migrations on read in the future. The alternative is adding a new BitPackedArray and ALPArray that gets migrated to on write. -This requires executing the Patches at read time. From scanning a handful of our tables, this is unlikely to cause any issues as patches are generally not compressed. We only apply constant compression for patch values, and I would expect that to be rare in practice. \ No newline at end of file +This requires executing the Patches at read time. From scanning a handful of our tables, this is unlikely to cause any issues as patches are generally not compressed. We only apply constant compression for patch values, and I would expect that to be rare in practice. From 5fbb7df43d24531b25d76debc4f0e06c314121c9 Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Fri, 6 Mar 2026 10:58:34 -0500 Subject: [PATCH 5/7] more Signed-off-by: Andrew Duffy --- proposed/0027-patches-format.md | 84 +++++++++++++++++++++++---------- 1 file changed, 59 insertions(+), 25 deletions(-) diff --git a/proposed/0027-patches-format.md b/proposed/0027-patches-format.md index e845206..8a3c60c 100644 --- a/proposed/0027-patches-format.md +++ b/proposed/0027-patches-format.md @@ -6,6 +6,12 @@ Make a backwards compatible change to the serialization format for `Patches` use * ALP * ALP-RD +enabling fully data-parallel patch application inside of the CUDA bit-unpacking kernels, while not impacting +CPU performance. + +This relies on introducing a new encoding to represent exception patching, which would be a forward-compatibility break +as is always the case when adding a new default encoding. + --- ## Data Layout @@ -88,7 +94,11 @@ pub struct PatchedArray { } ``` -The PatchedArray holds two buffer handles and 3 buffers for access. +The PatchedArray holds buffer handles for the `lane_offsets` which provides chunk/lane-level random indexing +into the patch `indices` and `values`, so these values can live equivalently in device or host memory. + +The only operation performed at planning time is slicing, which means that all of its reduce rules would run +without issue in CUDA or on CPU. --- @@ -103,8 +113,9 @@ We look at the slice indices, align them to chunk boundaries, then slice both th Filter / Take operations can arbitrarily break and reconstruct new chunks, so they cannot be done metadata-only and thus must be a Kernel rather than a Reduce rule. In practice, we perform the operation by -- Pushing the filter into the child +- Executing the filter on the child, then executing it - Intersecting the filter with our patches, ideally in a chunk-at-a-time way so we can write a vectorized version. +- Applying the filtered patches over the executed child ## ScalarFns @@ -114,30 +125,8 @@ We do not reduce any ScalarFns through the operation, instead they only run at e This matches the current behavior of BitPackedArrays. -```text - chunk 0 chunk 0 chunk 0 chunk 0 chunk 0 chunk 0 - lane 0 lane 1 lane 2 lane 3 lane 4 lane 5 - ┌────────────┬────────────┬────────────┬────────────┬────────────┬────────────┐ -lane_offsets │ 0 │ 0 │ 2 │ 2 │ 3 │ 5 │ ... - └─────┬──────┴─────┬──────┴─────┬──────┴──────┬─────┴──────┬─────┴──────┬─────┘ - │ │ │ │ │ │ - │ │ │ │ │ │ - ┌─────┴────────────┘ └──────┬──────┘ ┌──────┘ └─────┐ - │ │ │ │ - │ │ │ │ - │ │ │ │ - ▼────────────┬────────────┬────────────▼────────────▼────────────┬────────────▼ - indices │ │ │ │ │ │ │ - │ │ │ │ │ │ │ - ├────────────┼────────────┼────────────┼────────────┼────────────┼────────────┤ - values │ │ │ │ │ │ │ - │ │ │ │ │ │ │ - └────────────┴────────────┴────────────┴────────────┴────────────┴────────────┘ - -``` - - --- + ## Compatibility BitPackedArray and ALPArray both hold a `Patches` internally, which we'd like to replace by wrapping them in a `PatchedArray`. @@ -145,3 +134,48 @@ BitPackedArray and ALPArray both hold a `Patches` internally, which we'd like to To do this without breaking backward compatibility, we modify the `VTable::build` function to return `ArrayRef`. This makes it easy to do encoding migrations on read in the future. The alternative is adding a new BitPackedArray and ALPArray that gets migrated to on write. This requires executing the Patches at read time. From scanning a handful of our tables, this is unlikely to cause any issues as patches are generally not compressed. We only apply constant compression for patch values, and I would expect that to be rare in practice. + +## Drawbacks + +This will be a forward-compatibility break. Old clients will not be able to read files written with the new encoding. +However, the potential break surface is huge given how ubiquitous bitpacked arrays and patches are in our encoding trees. +This will cause friction as users of Vortex who have separate writer/reader pipelines will need to upgrade their Vortex +clients across both in lockstep. + +> Does this add complexity that could be avoided? + +IMO this centralizes some complexity that previously was shared across multiple encodings. + +## Alternatives + +> Transpose the patches within GPU execution + +This was found to be not very performant. The time spent D2H copy, transpose patches, H2D copy far exceeded the cost of executing the bitpacking kernel, which puts a serious +limit on our GPU scan performance. Combined with how ubiquitous `BitPackedArray`s with patches are in our encoding trees, would be a permanent bottleneck on throughput. + +> What is the cost of **not** doing this? + +Our GPU scan performance would be permanently limited by patching overhead, which in TPC-H lineitem scans was shown to be the biggest bottleneck after string decoding. + +> Is there a simpler approach that gets us most of the way there? + +I don't think so + +## Prior Art + +The original FastLanes GPU paper did not attempt to implement data-parallel patching within the FastLanes unpacking +kernels. + +The G-ALP paper was published later on, and implemented patching for ALP values _after_ unpacking. + +We use a data layout that closely matches the one described in _G-ALP_ and apply it to bit-unpacking as well. + +## Unresolved Questions + +- What parts of the design need to be resolved during the RFC process? +- What is explicitly out of scope for this RFC? +- Are there open questions that can be deferred to implementation? + +## Future Possibilities + +What natural extensions or follow-on work does this enable? This is a good place to note related ideas that are out of scope for this RFC but worth capturing. From b2a9140533043890efd100eb1a948590723def30 Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Fri, 6 Mar 2026 11:00:23 -0500 Subject: [PATCH 6/7] ok Signed-off-by: Andrew Duffy --- proposed/0027-patches-format.md | 95 +++++++++++++++++---------------- 1 file changed, 48 insertions(+), 47 deletions(-) diff --git a/proposed/0027-patches-format.md b/proposed/0027-patches-format.md index 8a3c60c..03e4488 100644 --- a/proposed/0027-patches-format.md +++ b/proposed/0027-patches-format.md @@ -2,9 +2,9 @@ Make a backwards compatible change to the serialization format for `Patches` used by the FastLanes-derived encodings: -* BitPacked -* ALP -* ALP-RD +- BitPacked +- ALP +- ALP-RD enabling fully data-parallel patch application inside of the CUDA bit-unpacking kernels, while not impacting CPU performance. @@ -24,32 +24,34 @@ The key insight of the paper is that instead of holding the patches sorted by th - Further group the patches within each chunk by their "lanes", where the lane is w/e the lane of the underlying operation you're patching over aligns to For example, let's say that we have an array of 5,000 elements, with 32 lanes. + - We'd have $\left\lceil\frac{5,000}{1024}\right\rceil = 5$ chunks, each chunk has 32 lanes. Each lane can have up to 32 patch values - Indices and values are aligned. Indices are indices within a chunk, so they can be stored as u16. Values are whatever the underlying values type is. ```text - chunk 0 chunk 0 chunk 0 chunk 0 chunk 0 chunk 0 - lane 0 lane 1 lane 2 lane 3 lane 4 lane 5 - ┌────────────┬────────────┬────────────┬────────────┬────────────┬────────────┐ + chunk 0 chunk 0 chunk 0 chunk 0 chunk 0 chunk 0 + lane 0 lane 1 lane 2 lane 3 lane 4 lane 5 + ┌────────────┬────────────┬────────────┬────────────┬────────────┬────────────┐ lane_offsets │ 0 │ 0 │ 2 │ 2 │ 3 │ 5 │ ... - └─────┬──────┴─────┬──────┴─────┬──────┴──────┬─────┴──────┬─────┴──────┬─────┘ - │ │ │ │ │ │ - │ │ │ │ │ │ - ┌─────┴────────────┘ └──────┬──────┘ ┌──────┘ └─────┐ - │ │ │ │ - │ │ │ │ - │ │ │ │ - ▼────────────┬────────────┬────────────▼────────────▼────────────┬────────────▼ - indices │ │ │ │ │ │ │ - │ │ │ │ │ │ │ - ├────────────┼────────────┼────────────┼────────────┼────────────┼────────────┤ - values │ │ │ │ │ │ │ - │ │ │ │ │ │ │ - └────────────┴────────────┴────────────┴────────────┴────────────┴────────────┘ + └─────┬──────┴─────┬──────┴─────┬──────┴──────┬─────┴──────┬─────┴──────┬─────┘ + │ │ │ │ │ │ + │ │ │ │ │ │ + ┌─────┴────────────┘ └──────┬──────┘ ┌──────┘ └─────┐ + │ │ │ │ + │ │ │ │ + │ │ │ │ + ▼────────────┬────────────┬────────────▼────────────▼────────────┬────────────▼ + indices │ │ │ │ │ │ │ + │ │ │ │ │ │ │ + ├────────────┼────────────┼────────────┼────────────┼────────────┼────────────┤ + values │ │ │ │ │ │ │ + │ │ │ │ │ │ │ + └────────────┴────────────┴────────────┴────────────┴────────────┴────────────┘ ``` This layout has a few benefits + - For GPU operations, each warp handles a single chunk, and each thread handles a single lane. Through the `lane_offsets`, each thread of execution can have quick random access to an iterator of values - Patches can be trivially sliced to a specific chunk range simply by slicing into the `lane_offsets` - Bulk operations can be executed efficiently per-chunk by loading all patches for a chunk and applying them in a loop, as before @@ -60,37 +62,37 @@ This layout has a few benefits ## Array Structure ```rust -/// An array that partially "patches" another array with new values. -/// -/// Patched arrays implement the set of nodes that do this instead here...I think? -#[derive(Debug, Clone)] -pub struct PatchedArray { - /// The inner array that is being patched. This is the zeroth child. - pub(super) inner: ArrayRef, - - /// Number of 1024-element chunks. Pre-computed for convenience. - pub(super) n_chunks: usize, - - /// Number of lanes the patch indices and values have been split into. Each of the `n_chunks` +/// An array that partially "patches" another array with new values. +/// +/// Patched arrays implement the set of nodes that do this instead here...I think? +#[derive(Debug, Clone)] +pub struct PatchedArray { + /// The inner array that is being patched. This is the zeroth child. + pub(super) inner: ArrayRef, + + /// Number of 1024-element chunks. Pre-computed for convenience. + pub(super) n_chunks: usize, + + /// Number of lanes the patch indices and values have been split into. Each of the `n_chunks` /// of 1024 values is split into `n_lanes` lanes horizontally, each lane having 1024 / n_lanes /// values that might be patched. pub(super) n_lanes: usize, - - /// Offset into the first chunk + + /// Offset into the first chunk pub(super) offset: usize, /// Total length. pub(super) len: usize, - - /// lane offsets. The PType of these MUST be u32 - pub(super) lane_offsets: BufferHandle, - /// indices within a 1024-element chunk. The PType of these MUST be u16 - pub(super) indices: BufferHandle, - /// patch values corresponding to the indices. The ptype is specified by `values_ptype`. - pub(super) values: BufferHandle, - /// PType of the scalars in `values`. Can be any native type. - pub(super) values_ptype: PType, - - pub(super) stats_set: ArrayStats, + + /// lane offsets. The PType of these MUST be u32 + pub(super) lane_offsets: BufferHandle, + /// indices within a 1024-element chunk. The PType of these MUST be u16 + pub(super) indices: BufferHandle, + /// patch values corresponding to the indices. The ptype is specified by `values_ptype`. + pub(super) values: BufferHandle, + /// PType of the scalars in `values`. Can be any native type. + pub(super) values_ptype: PType, + + pub(super) stats_set: ArrayStats, } ``` @@ -113,18 +115,17 @@ We look at the slice indices, align them to chunk boundaries, then slice both th Filter / Take operations can arbitrarily break and reconstruct new chunks, so they cannot be done metadata-only and thus must be a Kernel rather than a Reduce rule. In practice, we perform the operation by + - Executing the filter on the child, then executing it - Intersecting the filter with our patches, ideally in a chunk-at-a-time way so we can write a vectorized version. - Applying the filtered patches over the executed child - ## ScalarFns We do not reduce any ScalarFns through the operation, instead they only run at execution time. This matches the current behavior of BitPackedArrays. - --- ## Compatibility From c12dd5ecd1d5a837dae9ea9bfc251486aa4689d2 Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Fri, 6 Mar 2026 11:03:08 -0500 Subject: [PATCH 7/7] add header Signed-off-by: Andrew Duffy --- proposed/0027-patches-format.md | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/proposed/0027-patches-format.md b/proposed/0027-patches-format.md index 03e4488..1258115 100644 --- a/proposed/0027-patches-format.md +++ b/proposed/0027-patches-format.md @@ -1,3 +1,7 @@ +- Start Date: 2026-03-02 +- Tracking Issue: TBD +- Draft PR: https://github.com/vortex-data/vortex/pull/6815 + ## Summary Make a backwards compatible change to the serialization format for `Patches` used by the FastLanes-derived encodings: