From e8112f4cd9d77ff03615c4fbdfac65bcb99e161a Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Wed, 4 Mar 2026 09:26:13 -0500 Subject: [PATCH 01/11] 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 02/11] 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 03/11] 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 04/11] 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 05/11] 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 06/11] 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 07/11] 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: From 5fb3cf8f4db5e0bcce761dbe0db9e2545d8e8563 Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Wed, 18 Mar 2026 17:25:08 -0400 Subject: [PATCH 08/11] update Signed-off-by: Andrew Duffy --- proposed/0027-patches-format.md | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/proposed/0027-patches-format.md b/proposed/0027-patches-format.md index 1258115..48652d6 100644 --- a/proposed/0027-patches-format.md +++ b/proposed/0027-patches-format.md @@ -124,11 +124,13 @@ In practice, we perform the operation by - 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 +## `ScalarFn`s -We do not reduce any ScalarFns through the operation, instead they only run at execution time. +The behavior of some scalar functions may be undefined over placeholder values that exist in the inner array. For example, integer addition may overflow. -This matches the current behavior of BitPackedArrays. +To avoid this, only scalar functions where `ScalarFnVTable::is_fallible()` is `false` can be kernelized. + +Currently, this only applies to the `CompareKernel`, which will push down to inner, then perform the comparison on the patches as well. --- From 414371b6080e4a530717a00c7f4d114269912a94 Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Fri, 3 Apr 2026 16:11:54 -0400 Subject: [PATCH 09/11] arrayref Signed-off-by: Andrew Duffy --- proposed/0027-patches-format.md | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/proposed/0027-patches-format.md b/proposed/0027-patches-format.md index 48652d6..5240510 100644 --- a/proposed/0027-patches-format.md +++ b/proposed/0027-patches-format.md @@ -88,13 +88,11 @@ pub struct PatchedArray { pub(super) len: usize, /// lane offsets. The PType of these MUST be u32 - pub(super) lane_offsets: BufferHandle, + pub(super) lane_offsets: ArrayRef, /// indices within a 1024-element chunk. The PType of these MUST be u16 - pub(super) indices: BufferHandle, + pub(super) indices: ArrayRef, /// 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) values: ArrayRef, pub(super) stats_set: ArrayStats, } @@ -185,4 +183,6 @@ We use a data layout that closely matches the one described in _G-ALP_ and apply ## 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. +It would be nice to use this to replace the SparseArray. + +We also need a plan for how to extend this to non-primitive types. Would need to pick a lane count for the other types. From e5b5a72c6eb5a63b56f3ac64375d7ff60b6ff1ee Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Fri, 3 Apr 2026 16:25:57 -0400 Subject: [PATCH 10/11] more Signed-off-by: Andrew Duffy --- proposed/0027-patches-format.md | 15 +++++++-------- 1 file changed, 7 insertions(+), 8 deletions(-) diff --git a/proposed/0027-patches-format.md b/proposed/0027-patches-format.md index 5240510..af9f1ef 100644 --- a/proposed/0027-patches-format.md +++ b/proposed/0027-patches-format.md @@ -98,8 +98,8 @@ pub struct PatchedArray { } ``` -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 PatchedArray holds a `lane_offsets` child which provides chunk/lane-level random indexing +into the patch `indices` and `values`. Like all arrays, these can live 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. @@ -112,15 +112,14 @@ without issue in CUDA or on CPU. 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. -## Filter / Take Execution +## Filter -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. +We can do some limited optimization of Filter in a reducer. First, we find the start/end indices of the filter mask to nearest chunk boundary (1024 elements). -In practice, we perform the operation by +We then slice the underlying array to those boundaries. We also can slice the `lane_offsets` by multiples of `n_lanes` to trim to only in-bounds chunks. -- 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 +Then we re-wrap in a FilterArray with the mask sliced to same chunk boundaries. When the filter is sparse and clustered this greatly reduces the number of chunks +that need to be decoded. ## `ScalarFn`s From e8519c91435cc385e57eeff9024ba51b8538f133 Mon Sep 17 00:00:00 2001 From: Andrew Duffy Date: Fri, 3 Apr 2026 16:26:39 -0400 Subject: [PATCH 11/11] proposed -> accepted Signed-off-by: Andrew Duffy --- {proposed => accepted}/0027-patches-format.md | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename {proposed => accepted}/0027-patches-format.md (100%) diff --git a/proposed/0027-patches-format.md b/accepted/0027-patches-format.md similarity index 100% rename from proposed/0027-patches-format.md rename to accepted/0027-patches-format.md