From e45626c0360884118941abcdf1dedb9d9c5849cb Mon Sep 17 00:00:00 2001 From: Mukesh Garg Date: Tue, 16 Jun 2026 13:05:41 -0700 Subject: [PATCH] gqa: reorganize benches into gqa_helpers/ subpackage; drop legacy headline MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Splits the GQA helpers into a dedicated subpackage to make room for the prefill 4-cases study (next commit) and a single umbrella bench (milestone-1h-gqa, after that). Layout: benches/gqa_helpers/ long_ctx/ — decode 4-cases kernels + sweep runner short_ctx/ — prefill/decode short-context kernels shared/ — _gqa_panel_helpers + decode_opt2 (context-agnostic) The registry audit now skips subpackages so gqa_helpers/ (without a leading underscore) doesn't get audited for @bench decorators. Also drops the legacy milestone-gqa-headline bench, its _gqa_attention_prefill_long kernel, 6 dependent prefill tests, the paper_gqa_latency.py report harness, and the 3 stale headline-derived PNGs the §6 wire-up referenced (paper will re-pull from the new 1H_milestone_output/gqa/long_ctx/ once §6 is updated). The _ccl_cfg and _summarize_op_log helpers used to live in the headline bench; extracted them to gqa_helpers/shared/_gqa_panel_helpers. Co-Authored-By: Claude Opus 4.7 --- .../gqa_decode_long_ctx_4cases_latency.png | Bin 47883 -> 0 bytes .../gqa_decode_long_ctx_4cases_memory.png | Bin 42688 -> 0 bytes .../gqa_decode_long_ctx_4cases_traffic.png | Bin 41713 -> 0 bytes scripts/paper/paper_gqa_latency.py | 101 -------- .../paper_plot_gqa_decode_long_ctx_4cases.py | 105 ++++++-- .../gqa_headline/sweep.json | 31 --- .../benches/_gqa_attention_prefill_long.py | 190 -------------- src/kernbench/benches/gqa_helpers/__init__.py | 9 + .../benches/gqa_helpers/long_ctx/__init__.py | 1 + ...tention_decode_long_ctx_cube_repl_pe_sp.py | 0 ...tention_decode_long_ctx_cube_repl_pe_tp.py | 0 ...attention_decode_long_ctx_cube_sp_pe_sp.py | 0 ...attention_decode_long_ctx_cube_sp_pe_tp.py | 0 .../long_ctx/gqa_decode_long_ctx_4cases.py} | 58 ++--- .../benches/gqa_helpers/shared/__init__.py | 5 + .../shared}/_gqa_attention_decode_opt2.py | 0 .../gqa_helpers/shared/_gqa_panel_helpers.py | 40 +++ .../benches/gqa_helpers/short_ctx/__init__.py | 1 + .../short_ctx}/_gqa_attention_decode_short.py | 0 .../_gqa_attention_prefill_short.py | 0 .../benches/milestone_gqa_headline.py | 231 ----------------- src/kernbench/benches/registry.py | 6 +- tests/attention/test_gqa_decode_opt2.py | 2 +- tests/attention/test_gqa_prefill.py | 118 --------- .../test_gqa_prefill_long_c8_snake.py | 154 ------------ .../attention/test_gqa_prefill_long_pe_sp.py | 233 ------------------ .../test_gqa_prefill_long_tile_ring.py | 162 ------------ tests/attention/test_gqa_prefill_ring.py | 161 ------------ tests/attention/test_gqa_short_context.py | 4 +- ...st_milestone_gqa_decode_long_ctx_4cases.py | 30 +-- .../attention/test_milestone_gqa_headline.py | 89 ------- ...stone_gqa_single_kv_group_prefill_panel.py | 121 --------- 32 files changed, 178 insertions(+), 1674 deletions(-) delete mode 100644 docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_4cases_latency.png delete mode 100644 docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_4cases_memory.png delete mode 100644 docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_4cases_traffic.png delete mode 100644 scripts/paper/paper_gqa_latency.py delete mode 100644 src/kernbench/benches/1H_milestone_output/gqa_headline/sweep.json delete mode 100644 src/kernbench/benches/_gqa_attention_prefill_long.py create mode 100644 src/kernbench/benches/gqa_helpers/__init__.py create mode 100644 src/kernbench/benches/gqa_helpers/long_ctx/__init__.py rename src/kernbench/benches/{ => gqa_helpers/long_ctx}/_gqa_attention_decode_long_ctx_cube_repl_pe_sp.py (100%) rename src/kernbench/benches/{ => gqa_helpers/long_ctx}/_gqa_attention_decode_long_ctx_cube_repl_pe_tp.py (100%) rename src/kernbench/benches/{ => gqa_helpers/long_ctx}/_gqa_attention_decode_long_ctx_cube_sp_pe_sp.py (100%) rename src/kernbench/benches/{ => gqa_helpers/long_ctx}/_gqa_attention_decode_long_ctx_cube_sp_pe_tp.py (100%) rename src/kernbench/benches/{milestone_gqa_decode_long_ctx_4cases.py => gqa_helpers/long_ctx/gqa_decode_long_ctx_4cases.py} (88%) create mode 100644 src/kernbench/benches/gqa_helpers/shared/__init__.py rename src/kernbench/benches/{ => gqa_helpers/shared}/_gqa_attention_decode_opt2.py (100%) create mode 100644 src/kernbench/benches/gqa_helpers/shared/_gqa_panel_helpers.py create mode 100644 src/kernbench/benches/gqa_helpers/short_ctx/__init__.py rename src/kernbench/benches/{ => gqa_helpers/short_ctx}/_gqa_attention_decode_short.py (100%) rename src/kernbench/benches/{ => gqa_helpers/short_ctx}/_gqa_attention_prefill_short.py (100%) delete mode 100644 src/kernbench/benches/milestone_gqa_headline.py delete mode 100644 tests/attention/test_gqa_prefill.py delete mode 100644 tests/attention/test_gqa_prefill_long_c8_snake.py delete mode 100644 tests/attention/test_gqa_prefill_long_pe_sp.py delete mode 100644 tests/attention/test_gqa_prefill_long_tile_ring.py delete mode 100644 tests/attention/test_gqa_prefill_ring.py delete mode 100644 tests/attention/test_milestone_gqa_headline.py delete mode 100644 tests/attention/test_milestone_gqa_single_kv_group_prefill_panel.py diff --git a/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_4cases_latency.png b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_4cases_latency.png deleted file mode 100644 index 69a179de16699f3de12dd6b6a8db7e28502d2a42..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 47883 zcmeFZS5%W*`0fk1&?SgiKm?>%P&$eT(ov*^5_)e+ks_gn-b4if0YRD|oe+8iLJK7- z3J6Fik&aSB=!6;qduIL5*jHzdac<7lIgH^j8c4!7=X}fWdEPHtS4WMBfrEjDhK5N) z{fRye%|D-MXy^+5Jx%?en%ytpFBu;dQy&9&M<2hZUJf+cPklUIxcj_te#YzT;N|V? z4!bXU@2==WVO}R69}jOCF)_FQ{T-t2UeCq8TEKh3tDN;vH}j^UVY@{AL;E~u5WE-0 zPvglWL;p`JlV_Wa3{E;Wr`@<8Kt4RFQ9o6}U(yqkB%S#9^=W-MLqX-prF(KkEZ3D6 z%rqizCqGGI(N2OmXZM}#SoMWvElvdWHSSFiCPLOJ%{D0SKKUTap!Hhhv~^HLp)Nm* zICy{X#hc48?@RYTkHgUyzWmQarQD1EpMEqaq@kg)cHgf3oY22x_3rDbe&63-Y|LWe zg$^hbgs71(c;Q%0EdStYTkpr%U8DJOGn}5qSe}Y1?Vnp&}+%7O;S$ zF*H|bXgaIG%XG9T`2UT*W1gcBVp;&6>NEyB*nUj$#G>nN;(r4DB)bMAo0{KOztRF06+gKps<*74c7uBvk z<%1-Db{Q!hJy^_f-(Kv4K1e9vpZfq^$&VGSZaXjAgFrp|5rZq#xV)D!yfRu-5Wl^c z6E-&U^;L!EY`ZbW$9Zqyq<+FH7QCD2f>3xO{Iqk!F@-$-CyHk+b^X2ac~8PVcx{6? zizZ(;*@L+cZP=Sco){zNN*uXmZ^{+?oNu>bv%jY-sb6_gE?)#MglfK&J(cO+$_{5I z1?CQxK3&5QEA8Rs6{c0+C~G5SCZ!~Y;UewKiNgv5q)xS)qWe^HL%X8KOnA%w-&o&; zkIykX=H;gMWcO!YtqvPNYkz*YW5Le?@jm(y%VvQMrL1(`Oh8vVben$^cPM}K>NIS! z<@k^YI}}(QtL<)4ElJW~6*Tl97-uh!);!0cjMk>H)7$}b7wdBb#GHdl^k)SOicFgw zdp|kL2a-Ithnu@ogr197)S>U$wg0Yj8y7BwU z4>~eG1x+ef>dby+JeXWR!xMaWBorYh9(v^8G!wzNvQ&`ZOYA{yYl=30U+DSxVd2Q< zsF5_N*+$O6HMNf%sm5_Ql$oI7Wu~&_gH|>>c>Kr((B~tiMysE^d#jigkM1rU`L0bk zGs|s+kMpD`##|9u-YVSsJ)dT?wq<0!JDZ?5HezD4+RKc9q?o((^1rHm%K>>-RdY+3ijG<*ZINNey)uY{v}l zPI&h#U^Y3<+%YnI;`HXv{t)I46a;Nq!G5P&+h%9!9QEb#P$Y$E2#ZsTH7 z+neoioRd4LxDv%|!USp?7k0eO<0pR`qq@<(v4gikO8@E>Z4z#ayIR<6Y=-eE>6Vh= zChpsm8H@Od?t@?3Yfi4*0oQy-=hM>G_UET<6ZfnGMDN2z7xugv-;2GWVR34UV#XUq zknO)GC>~!eHLG#-pSZ*adEJ#N!g*}f?7!KcEMSlyPmBs6AkrSgi8!VNn^MCP4=gw! zZxoLAXNzaW%&Ko0wgm3jZ_6y@zWO%XvfV5DCp0Hy*Q^NFy}LTD49-rdzs#Z2L05_p z+i^(Gy^cZ-_x5+}d0X?{QBeX`7&Li*8f?CQLk{NCO4`Zb{C;a|-kS=_6ErGqIH}b^ z2e;;=dVh_e|Jg3T``d2PhjKUSnwNTYU!2rb`iuSbR5ca$s zpWjcT_`2^Tjvs^1pyNI6KGAS*xYDZ0e*A$~=JF6iGp&DZvZODQXCm7vd)N~$cxaRFDnpjSP)@Myce|mRtAsP6HW;vxN@`) zZO(mT6w^wLO1X0r-~EHl=Gs`zb0~jJbv4Q+^kBXc+e+RcrmdJ>mTm#ZST$=8MB{HA zG&xZ>)7C-Yg-9m6gAlJSBFhr$D&mjBsebG>}6;~+|Y zbP3s6o0MIUp^(?HZQUn-W5pu>vIbOaZFP%9p}3OqkW^vzHRNipou7o6G&6boEj|wu@@7)B$CJZqDSN&*5~U_i2SEJ-9;m|G++RJaBnE6?Q`L`11H`igAVc z8)&bD6 zk_8Ct#mZ6Ve~OaCIlo{)g!{%O!EmYEyjF0+opQLd79loI|0+V}_y(kV2gIc}60Qjsze!5jxXyxyiWQhHQf*zw715joh373#ak8zcCZ1$q?^6Z_Xr&AZ(plm zlbAQHl-=omSdWXl>5?fr{5 z5HMZD50s%=mvM2z=IG9fv`)^-jZ>c zsIO;*QZAAZy!G?W+U=d3uwxC$KJYOVf?pnkWq&Q!l_p|wMTWdm)5ccF6TI=LG~oQT zhj%B!NJ|5MoLw!ff>vvPiXIhAPE3Z9S1WT;iB%ok-cslD-70s6%FXWMu>so)CA+o$ z&W$12d)=v4w8E}CY_cdf4bl25<^d_y}SAc=UR9tCwQ=edEPzeHWN1ZU`zMp zm@;`d{7bEiQ%8Gz`u)z*z)cbng^eT&Bsc}WI>YSS(QHaqCGW0U@mhRU8*L8oc3Sye zwRRdx;Q@Cdp>ROX}?)SZRhpmHb+&@<9 z+D3kZL&W2kQQ?QU3X6j5Of;1Eu-x`Td`mvX_V?$ff8-=XkM<3>9aIo3y&YLa8#cq`;;$Hmz ze0>-R&GFy>*=fjd8^_eN9tJ*MRQ-po)?Fyi=bzCusk=P#{0E@m__laTB!|1&Dv*GH z$;m$%Y)u|rP0Ktm?cnEe(!Zd*cs8teciZGRL(I|zpTIbsLG)u^q|3oUJ<2G z6&gnV8S%?s^aWU$P{=Thla>aXaIGuex9t-gXI64ZKHv)8_`0Slhj_!t78wwuSZ`0_^k8gjmnNc15krFmmgq)xrfzU~N1dCa< zuy}52B2B883j8Nu+V^X!Y)`5n=>oij9LeRsI=7+O*Rtzj90fn)X>k7)P|&#_fphfF zvSi@TU_tbBM~5SCY$Pym4W|aRE$yd<z(pyC(bRcTq4L+O4+t#@@&ZgDv6l^t3%f zcqfguHCdzU5(DPqwp8e`<3nyxo}N4N_*qK6-sj0PQ#X+z?d;DR)&?N*3l3ct3Qn5? zYCI}js>D2%*eiropSM>UF3|rO!Geme2-lFCBv1B(@MB#K-k2_=5U@9LEW-Zh+X$8z>_m%XnX=y&wa;#9X_8kx4a;I#l~C(dd{e$@WG-llcAe zm%+Hmup}~NZ@S}}4s?QoJ_aW)9SVcfmc;7W(D6Iyih^Ios>v5Vq)R}Xl6)j^b$e2e zs$%wnj^T>*InBuIKtThkbcJ7_l#Z8vSBHSsMAQptK3FP}s5Ug|!aLhdUONb}a8{3n zXxZS^~RMA+9p=t3`gAQJ7#51`P``!npI|KkKQRcQ0*oY=GRSWEyLbz(5^^( zFBCF7p7Lm8$~*eLNSh@#m~Zk@VggjP+z(C!ZL_NlW08X#9I8^pM4#a811Q3zKu=eb z@2$5?`23z7SCHLm2g7^vnNwBPcXTw3_u$xTg3*_R9-{EC9%yx^irA(4Sb!d4*b~k) z(Q#;tiq}xYT3|FnJzIvrni+0TH5F#1&tHR z?I#Qz3aboJ*3>*lM#&t^My-_$#y3mo$l#sEn$}y1g4fxM97`d^ z-h=hip~O~MpQQ`l>82FJix`xpQOYvNEtB_;W?XCO zY1-<{)M*p+WTW@uxdc??EknZU^PCVzaGAZC*b2*E%Q(Dgbh)a;1S!5@zZ&K8{x5n6 z6n0O>`8g_ruxo*PYh2z9$;7esMVa2>!M`At4^tagHf#Sp*PnWqqm zw85TeMaB$XEu)Pj;c7W(^?ncril4!`WJ6?X#`mVf+@T5wg(08{sJ!GZ^1KW3#m*?6 zYc6=DPgbVrlV>|SJEUGP)8q#Q<2a}m7w+%k)MlerLYveBS+_hI)D&wZs&I9T+*L8| zYfBt9RdQra26ep+U?1oNCjhiNy}-{wkt&syIzFjjg6`~(au?k*YNuGr`z51$)HdG{ z)1=L5y4~Gts|%ry_3-k{2Ufbo+nb;Ohz+pgjy3;n*qu7C4`0RJubCPt9jBI#Qw#YL zhsNTpZ%q``8?BSUdLK9t$7;(M?=o}A42KN=O*el(ys2D4x9o?e2pVhJJ6E{YS|V`Y zhLbWDDiZjw<|=Kcm{=2F$M+wLOElbNv6F#aj`- z@dTk77dX0O&~=ULNQc1#6XtK#f!y(i2ncI;`F}@P2%wRzLQ`n;P(EX<=}XQNl0?l= zFJth`fS9{6olc+PN$~zN>!Fbk?($E28M`YP4?*mfSDsz`A@?UOp!5CpQR~ML#J?lK zm<%Nog7^suAKTo3r>{LE-u|^H{oo+m`oe=6q9uZJq`YY`SD(hkvBjY~)wtUJ=f+)^ zE(PO~2D^Cac6q+A^_IOAP~7>xddMN55Aarcts$n5Ip!)?rp zFYq>DEiJoau9f!V0F|tQf5Lg>33w(?#0MkeVpE9~boMi9d2O+Zh;vi;Ip7wrE*!MNn?2p|E6`ZHlAE#;1b0D%>g3@y@+G3HPqVLDInl(gl7*h_ub+wT)iN^hjylcH{Ro3PF$ z;uUTAg=>eqh47vOYTWX=T6F@-`26{8Gv4 zs@i#;^{t{2#Tw`MsBGSD06AtLR|=0vtK(8~<#d_Lz5^<3Z<@3}J-oNjnl4UZw>pi` zK%66)r&@cDxQZuAT@}drqE^6ZJ2^q9Ed+(>SNhk65sx|)Q(^`_Y?T+?=SyiA3`LNH z6#3u6*BTcwp@KW4UfK2U@**v&tgaNt`GVCes82m%giUfhL}cl!it1h#`+dhJR-qO6 zo+7>$li$H7op#}nf$w8i0_la+i83`ECFSYj4lzwX-2wPu6fHrU+iwVFhset}Ho_>) zDGB!tBD9#V8oKDB8V+VM!$NfNa?U9_@6_ZLpJn;MCwk(yJ$wb67O(JP0RrSz_SxEb z#3=6^qXPXf@u)~Q&mJyjpOTXbyZcO|%tp@elZDUAOJXeyOJak6&BQS8V|Nd29i8Qa zsq330!sdu{vma<3_XY>`8sjI(t^cTn+9T1DqNM&t}qJ+=%P1aWaysGCm$M6*q7 zB2>F`J{~RGWEo6~CZg<~j5E(7C`PF9VjcE_wWjrNkJRYdkBH(QCsFyUVio4KUQ$c7 zE+bx(hU=Erf46g%r@A{8K5!ki7EmC3WBGa|S>@S6f2Au`);g*npWUXKrOm9&*i$me z7#oFEzfxv?j!BXbkG&KYRc5a_Cw+G+L_V9E;eY4&u@o9aBi;?RsgAti3ykZ>NAa}= z85hDucPTzC#oO+_pQ36sA1Tz+9MIAV)37gS>Mt!U17f5#Vk0#k_@h&i)-Z%+6aiYp zHFB(2>s#~d<#haJ67*Z?onZ8o)hUX ze0lt$m2ib+qxdgp)iQo-lr#l2&1?CZsEf;IE3Gwj&vF{!mWlS!DXpf?u~9bIK&yLW zamWZKaQMS|HYxN?+4;*K5*_eF0rS(77s}9P%`YEfQNJE^qD#2{^xJEDMD6e@IB(VO z`k5kZ!U9DEr)08@HdPNK#(o_|&@)BSbFL z{85*GoegZkE~V=QzA}q^Qa29se)*-xvX9?=>FI}FtNbKo7n)mThF<*UlUkao!k+l3 z#KqN>hs@8@{tNIG6R^`4mDaYj<%l(*ltT%{Oi4Gd*%zbLd_MbsqMAHiaOH#yCgTrt z)gFJo>4GP}{ZbV({meTeM>$dez4bY%%oecjy0h(ZZQ4569h^lTAg+y}F6S|zhsa}K z75x4n0d46spaMIha;?boTpZ@0fosUY{Q#+g3$xWF^dYX`X;+%qdb%Hp%LK6RG^-&Z9qIBWf&R{O(g<`%PnYJo=yZnXz#}h1`6PW{~CcJk_ovEz}}P_ zbRC5~X3+j9&gTQdQlt*Rn)b6g*>cjK+LN_qpcxVZ6YG{pY->SEXf*_}NL6iHT`+g`zZG4x*_&}M0t^~lB)x5@0ySk_rw87jMz``HTcco}P zIJuMUe?xB7yI1-S3A`wu#RqIn>7+!>VvpUz(@zUY7ogh1UqdnjVlI>EG# zETQlt0G+w_i5bwswT-$#-(@_?#vfCak&XqtvTX>vQr-Pn@n}Q2dZ)yp|85j_pwm*f z81|e4+cYUS@IYFA3`s@UMFkSA;2P#}$@(^gcFme+0~j?rKigTfSlcS0QxBW`M$U-m z;y9Uk#eCkV)KK74tNP_TIz9o0fT(Y$xwM_AHTb=f^+}=I+@;G4)U*}EcA3+1Q4~Zydq``KdD_SAsu~P0uR;_a^G%kT^1nH;4*1MDo;JU=3HHOju;bZY z=&^qyPHr%NJG>n){e1K9k9$!~grdv|42j?K!K8w<0I*{?SK3 z-t6r-5#s&G^NtEQUm?Bx*F?WzJ+(IqKx5jFe$s&edq~FT*8}|PBAqM_??V7Tt^E%V zn}GNWrq9L#4_w0faBzKgkPY%5Z6goq{7^eM=W_>TB)5{8hVcqnFcmyH}cNpS2&8c8)fuL@x`a9nb7eB&s zuo9@w)v_<{=uE50*B}E5%6-B*C&-B|8-VRyO!k?3D6sbfu-nKjLHgKW4|%Vw!__~E zV$B!$X*vBx_o`|bYLf)I<7dbe>&A7Vv)k!%8*|0*ieSTTf)eB>#zHBrDh!6@3m zJF0*}KUpXPN$laZ^v-wXD?~FUTYmy{kdF$I(@^q(2NHPUyY{J6iPXu`oo)?L z#vUKK_5?US2%p!?E%t!`t!z?d{(984zpx7=v_2aml-pU}C@u;z;e8EgygsPw?g7aUiVM3VvFRC0k>`6rsvHEVz@)fbo@89{hcTtH^Fb|zU4I8lVYAShbfLvLKcC4G!) zJ#?`o$3SbcorHJRq-Ui^-$v~@m zm{qAKX|5S`89#Twj>DRMan(PYyesJL+-%(v*f^UtErtFiA(C=MojvzmA(9|iQk(=o zPS4A{&V`imTe}trvZ4XKAH1ME>E77$_=kW0O*4)|hFArVKsIBVhj*QlId;xkF6ut> zEJ1$P)Z|bLA*sJ6HR?ILp74iD#>ZoO1T(U!7ohPq-+v*IJIjvc$v~Wkz0&ouO*Mq7 zQH0u~%kx4Z#QJ#7QNH`+K;u)U*}Gh;1q@#VKYLDib~IkP!%#VINTYe6N>1tRmPvJO zl`&p^Sh;N(Na*d{?fw;9?#fMXpZUObVPQ2v2$|)ZA)AI2Ca`D_ZJ5m^D9Oxdal!f8 zt|UIq$MQkDK167TNC@%hs+o(tY+oC=WL@jj>KQVaGZ7khk*4~CpsWH=1+6|Tmfrtn z!V4T6y6=nW0&PZHmyul%iiu0~kd#uvA(PR-Lz-LqQ>7PKXa@7rn^FAr0RR1iqM=cG zxdUY0*Qn^~f-PT~PHiBY&q=#v2Yd;qtsM27u1=&g3GDja=>41wE}&IeI1sTlVRu-8 z9&b{Nzyj*=mBLe&VoTGh88Zi56}$ytC-QJg!9J1SW^5Sqc1yXI&y0XnGm{Ym%q0B$ z#wcqykdpIP0pko=^aM_iQd!wq!S8~vsEnA zrSb@9*>$@7$s*nFR?eOmz#>&Dd=A_K!16RT&xS!k^|f|zKwuFu9|p2S$AU8NwTB*G zfY@+*?eA8t`Q-#e2Yd3;`rdt#sr*fQT5O5P0V71*ZT?9e;TEK=moh0m#+I= zXX}1__Of_$?~R<5)rrO-faX^sgv=1mFQ!}NkD4qSy()jMg1p<ZMOuOT}|Dk z)#bxixAUYylW#P1O(kQw*hE+&aY%zLytDqWapC&EV01M(GSB?cGG7=L9%wX z=2`My%8N4h-9W6LP}al7f&51-QuCh#c3pHw0L%aXxV8c<3}*R&W&o`ikS>rZ~2&Q0-I-&S9h9d zsG@o+U{m=@8IZWeIw03#(1gI7H&N-QI%NPlWVl9u&oEyD@q)!xPk?KrHaC#VFi;gy z^fxclhF-A(AnWj8yc;6W|)y59krsfi|6 z^nJ=1mKWvLEtpyW_~9l>_org24UFx}13w<_t*9^vAIN7Xx5}Uqi$8I%1Zk6a< zvYAe^F3`<-YE~g^h6<+oN_@e2Yr-p$T#D$T1FKQUwNdSH6_mb!qEnDO5Jf9VMi6tU zK-Rn2CNGVkStw;4a***Z{p5vO_vtCFOcJCRmI@Yearqu37DoXpwq@%;mBuUxVf}9I zfMKZ)8O~3Yaze@2x-wiCum!CYe$eqk85FVg;O@aUX1*FDb-2=yEMJB;J*S!jQn|!4 z_kt-8TmM<1stW*yR%+(dLA%Z{3PzV+f?yuuh4jyG`MiJNJ~g2le~rC{`&;t%+uJvt z4ds|4(IY(XJ|L;5B$)x^Qp`HQt+$u&mJVqj+!gA0Fi^HUhL;d& z`Bv#**gDW!siZOI)gELyfQdzhYNWg$k=#2_M@10)bF2xF)tjY?`j4EsE~CSo$V>1b zdY@+_e#J_$5FZyCoToNXmz)xz07B){qCP_4-Tc5FLkcE>Qs)Oe-a0GdsOb0Axl|+VsG-ZW7 z&n43^9ol^M^6l?h0yg(w+v=p$tfI(+3OzxDI;h2WGo1*8ZV-1eFI z@=_QI?auh;Fksp3eI~2oZ5H{x@P+xLJ!ccKP@CE9P%&h;!g%vrm>2{ zxsOgsY2hoa?l)fp7DkT%aR+XrR??&>_Grs+_6H!^X{=2Xw|EO!$7mx z7@&6F6<4kcn)S9eYL=Y8s5k}O5*aVBZ)xz&SBb!YtKKPOWKl-ePDM8FZw*L;6J zUE2fCDT9pRGYNw1l?-tQNX9G{M8ZEl!NAybFDhuc1o^FvO4~j{3aDn7#bXm|QJdD_ z>Ye_Y=Y8oz{(s)ECAmQ&;m_InjK=K_va0#xkP%Oul zY`s1?q1u%Eax^rJx3U$F`9nbT^zi2y++r0mmlzEtfmnUZA}r(i+`2NeyacsmMnG)l zntwJk9=KjP1d`zW=3}30D8rJZ%=WpTAMDQiQT6Mm%eV^8 zu_izF(&M1rm9t!)-!H<_FWyvrEv~%y9k|7Abqhf(=Db2d=JTwmS64Ve$knUHpHRIx z+lgDyxmLZZXFvLz1Ge*ZEIML+SI79qdGyzKtS~+raM%;ANPBN9Fl&Lq~(#@s&aU*F{z`w_n-H^3ye!N=8pbh)270G;sblP za>W%R!?W*A0;QR_oIA}&kRX&^GTO=mvlnc_9ZQi69w=dKn(t1FqDCPYug@d#u;U}2 zQ>Mps3!pH7!aaYth6LfEQmL1qG@38jn1-tCIy|Jb@iw_cN0BBf-=#WoZZ*B$H;T<~Phg$PE`x6)awC^YgLC^sM4AK=WGLkZD;6M+^+pyKQOF)JMaP) zmp`^|=q^w@YMXsqGb+reMyg5WgD~wk>%;cm%4}yzM?k#cLa8RnJtaJz4p&0MaOf{1M+| zwpRz6T;8?8uADk^{v8Gu58QelWANZDsfO4qB3{HB;{T=%&=?zUQ#Y$qK4mjTsPbPb ziT<<;991I})XW%iQbCMVu(vAttyTh5&Hf z82kae8ht{Bg!5l+-M3qbbQmh-r|e{j|C=Z0HK=i?@fcL0@UJhjc#f)wjf>fgiWjcl zwN?YT9Sz+I>f@>e-z5rY;^C$qr0*|1see3Wz*UBNv!>w?Ui4??GEm+P`!G zl)E}o;e0#GBi>lJ_dG%MYuc?%KDI$dTLRsX;*a4E%j#6`swLI5A#ahY2mO09=k8$w zEE$mAWzqO;eHu-v$H%I3D>Z%b(x0-?9T@h6Id_aZupOn3+er0NeUtmM?&9zv4-kzU zOVQE$Nwk6iNf$T&`OG35W1Mt*>-3MbqirgIZM_)V?72y8>mX+wZ-+ksoDXj$6&)DE z69*$nyj!@qN(fVpP-c)IhRweM#p`$ZI8fb>#360XyJ#Yd?c;6Qqc7ial!;OH|}0r2DVd;%_(oOi0_cz^-QjTe)Z2t=?XmR7jwJ2aDVEPWPv+ui;`U z6|+Tj%L7Y~(tCF@ROd*wl@e=!i*t1e$Utqpe)5$>m?nM|vg)u`b6)<`MO|Y7{)JNn z8DF)vn~CIvc>2V}x#XN=T3h-GiR@nd$9bN&_IPDQ8bZN&VH)*E@Yzlin-Eq!M$bxE zfY_;LbuC7q7yLPBOcGj+q@vAv zdsZ>!AVdGaZ_Zl8$F}3c5SW*Pf63%Ob|d71!@{GE`X_}9SfjA?jN{&OBiM^(>*~IP zb#bXYMki6r@Ij7Rh8* zYM?RRT`yd6Xd4gsT0_T<2t#Qo8&=`QL5Pir?IFfoTv^2aN2wTer>YYi(+_Z}_tJ7w zVeI?>hp2kMxt~t*4?DPn%8g3&e~g@|xBVJnPq7B#{a(novv*7Zyee2z8(6*)YvQ}m z;sl?HHh;E5e6r)OA}cB3iW3gjGR{(yRuL&bxN_GZ+s@G4U`{3>4K=Mn6qGMj)UGIV z2_j=I>}>7~3!Iewz`fI-vrz-jUJNz_F)pnwC3IqwG#X+!P`j049;6?e2pssI z^OU0;{V@s@27CUQ{|Ga%!YG&L#tAMMA--PI|CxUZ z9;DA8M#Mx$Kkl)hjT}}6dVsmu@~j^qtIu?1LBjW(mN9lLdV?c5@w@;jY&aD2bg!F$ z4`}xB7InG6D5jp-bh9a3&+iR4DXyPx9`3*;8b6?=!_;~rzd8;6&CZ}=2^{!XU3Rllt z%L9Bcda1eVXLF+doTpk`sZ$P#S=2{|g2x`SdRx+K-0b@)peVrkP_N&L^?Mk|(6H$t z_-kf#d-u9YsN$*j{}Z-{x6(x$eDaW)Y?J#X!ymGROP?f(X55;(vg9&c{0DXvT=pvP zS@ED!hFUX1F8TCy#78lK#3Zqge>DA6rmlU*d^QgJ+vuqeVqFUx}gzp;=->+5%gY!|jDl#R{);F0R-_t%wIfSlz$3EcFvbh_+$ zo+BJN5t4vPz@;h$Ulwv=|FlIt$Jm<$cF*E1=6*Fs@E!j|Gv|IEJ0mW*O{*=DJ<=Sv$QW!&evHeRi1VK}-Q115O7i07B>9lKI#Vb8Q z4A$*}wSUsCPtEIo5x*~jE>a^}ylgWK9=kEb$vholZa%b&=(7&lF{%%v*2|){NC`s0 ziP0dG)^@cFh%DW@97o4_siEGQ=Hs}%*onHsU4)B4YMu7MCfR;&lkojA)TbjXMUo8A zk!37ABLQNz{M8s3gI4~1249h)C@xS!`}gyL$1;fT(nM7>W@a2w9QfmrSb>;AZ>U7X^p+ktL| zkGQ#-SiTG%7fT%KXM`%MuRDqs2Z*6sE@fq*t$v)jTQE>I2$kPSCfm^dn*xdUEeNHO z+*b|kXA2$ns%~DqkpPupYp<~>DKzf-k^mts)oQdbRD5|0xwm{hxl)mY@$^kz7`oS^ zIHRwt$`ZcG>|1>U6!(zh^nwHkn$V2o{lL9E_fHF~PUQrt(ncQU)qSuU9+BqRja z@%L2r|2}=XhUSmiZVEc=0RILi_Cqq9&Y`H6VFX%`KNlE-OfqD6Xz4rZz#d93dcUpg zwT_(B5qk)P0=(ho`TSW`sn=UfULRrl>kK{ahy=>(lVWUy6Dcgoo#$Fq7l`ehOG0M_ za%PB;O0*B;nT$20>DfquZ|SZ}YX^)?A0c-RH60;uN^T{keKa{`H=jnbt{Mqur7=FM zL0@Gkej__6W~_+8KYBT9fTJBxrUq4#j+@Vl`*<(ktnYdpA^YvB@~s+18x+)kM_nL6WC)%9poD7;nWa zI$|eOx?p5BIn@bu8c70zS3}5kahglobDIxtcGUJP40H~n&8)EhSAcw-6J~5}WxQds zqgFCkA3qxSqZe<4yp(%#&=)2w#NxCGw=Q9s2f%tIR7gxfDRqR`7#tSfeo5>?hN~r) zlb}4xm8azx^`(sa7UGJ~X=x{qc~j!K?bY-lkt-#SXm96_VPo?U6(*RN856)>c* z3DWT-fPziIvZD={=w`G1+Yk3mvVDdKnr9vEq^xa;Zyj2~^uBBWi*QUC$;e(_QgCxt zC{t82Gz$sPh}5WuHY?UmUKGYloWr@8c6wZW>X)>b4j2gvf^bXPhN?JoEh;V=}JvpeYzsXdS@;L~iQKiF;YVeVQXXNki+&uUa zE-09wdLV4P@df`p8AsF=)V_*tX463NAK)2x4*i6Wn?1m|(k)kyFIcZi+oG^}VFWOm z7G<8NKkHL-EiNMa*l6w@?m9xLVj9;f+L(O92BjrrdYN6Kuk#=ev=JY>3PL9#BihU6wr8FnyNcA&e9nT2_97b*Qk7L&!VL+b7UrED<0yc2nL))mx^Q_)(E{# z849@hmZi7};HHVTI8t3AGp@u?Ec`n!H4bfxoF(K-jcI>09~~qHf6TiF&dk__bh7eP;sDDViE^OnHJ9 zPz|ZafhkJE`38=!pn*w5qFYxd6S6LmM=YarFrsRi^b;1^QkpynCJ6O+OmcfsHSe}q zV_juSZ%m@?{aTa-f97+VOzMJ5Pcsbk-1LQqrA?4FvdI|=@l~n&?fUF5#1W3^$2rz# zwANF`J9mmrZcvC;425dl3C>0tA<74vMuoSJAeDeQvsdF1mKHgl=gLc^1EO($zoQT5QU3Uf`mrCh;Pi*<;F)fG7 zWo>-*J3l)&1&r?wZ4+f{Iy{ZwGp|ep)FeK^Cz2_b^QZh;r^% z*l@949&po0kiMSbxn+?+9TNKW@&27*z!W_=#2V+4iiaMqd~UNm0h5;kE@2|jxxlcw zM73wI0t$IAs`?qIKl8*Lx*D8q{3hX^2~v(dSEY3H_4Sv4q3|4#=6=C#Knb-(&8QSe znkbifYj8luxXlT&(3xsr^r3>Xrz8qhT*#7~FS!~2@S9uN!S-U4+iVDkZFxY>iyEp2 zv$FQH?U1x0c4+p_xTJVk*L2tkm(1LIWlJo$9FdMa=@(_d04s41fR)!&0#yo()I9-D zDr=m^H>g9Q6u{XQx4OY{exsy7Q*2Te6M!v7*U=Kplfe$nTb+{GfK2i?a-izzV4;Sy z&N|S~&NCC!h(N#TGaS<|(wnYwYWi*-4;sriF)RXq4}ny%xwZ%NfNvD)yb{$9b6Ex) zIjgCdx$R)!8hM-(dg$zrJp!-v82QNyQtZF^R}`z#&DoVBbG{@e-p-&%>nuP#eLG|@ z4=~}G4``|L+I{DrE>oJWSt=P)?$*V?Q2{fpNC1o!vu4gVR!1c+l;z=)*{v6lP%f%K zC;?v{-jLY`1%|TQ82T*F>=4ARL?28F*EiwiAkYG>D%~P}U{A^)fJ5 znNTYyvSae#>fRner!vjZ!a|a3)*%T9WQ%z{orLC#0U~L7p(j6k9+-g<@5f*twaUV% z!_*U2_e-}t{K?!dKCVX8lA=Zw7cb7s#J!vP ztTyFwfOeKSempk)wi&pdOWJ>x(4;Oo4xM47$ZTO*b9DoQvEUfCTRtCHiu?=^+n=XC zc1Y(#-rN)Z@I-)~^!l|r{$)bC?Slb3g!GHKUEP_EGs0d%;e76`2ad%682<4gq!-DR za2|E;JPaf)+0P1e{yk-Q=;W`+#l;CUs{@Jc$Ot>kLBh4^b*MQMz)jztX7+^gxfcIZ zrql_@JL8pIVG15|r>-${scchya@{D7K9m6$r;l88EwgG7bqtB=jHFNS<|1F%`Pg%7 zj^Dv5&lTm_nhP_W6L@&>qX}BXU?GQU%o#ecA{)(-NMT*x?{N;d=r6N0?^0^R4s(lf z3B7q$ThHgJpyT*O*LfZu($;}G?SK=@Fm((NOscxM#?98JOq<*wT>I$tDY*j;xc>9y zN4iM3elPQQ zAqwY5?w&e*uG0fUYtIdgvGk(hd>DZtdfT5#e3Nfp0E@xn`bl34r|pGb&Ql)G{b;%s zCr6HE5y+9_(|SyBA+X$*|>NClzMA}woKhhEdpPVod=aeRD z^=-)HLM}k!d4Pg7=|e$yQsYO`nYP-$hLH9%=S+!@L>p!zHpB@HiAJLr!Asug_tfi< zoL8CGn3N#yBL`12coFs!lOQclAJJm`{JCzq)2(PoV1P0a21M^Q+VHw87z^CHa>mp= z;J)?8UqS3=ZKHQ7u9c1LDd@p_kEGCowFUZo0-B43Xb|iARoJXpO7#&jngt6b6r)GA z=-sc$&PU+x2$r9}oZ*3&<2Jvi$#tDevETmM;KZjXPp`~!!B9giY(^OxzxTs6>cjDT_Orq_!ZIl589pd>nHp9C2dQRrkCvTXdB@ zZ1Yz{L{I|}B})ae#P5sBDGC@+`dIQxB|!vA(CZ(cXS0`@pt2dW4?;9M8o{K2ye3ia zWeX|7B(VeEa~gFiQL0)foVD+cc)}4_puDgX@^mGsarZtwiskLFLb>ZJAl^wjD(MiX z&pfJKKc)WZ7UkzE{wghv(i8#`ZDjLqxiQ+<>{!Uj$19r8H+Q9RoEk~3O(vf zFfLE|3^mP^178Ua8y1GURm~b876fYDaahBj(|DILy9D`EZ3B7+?pm5gr-2)Ij9r;= zg_o?f(o2I|^_PK7CSnzYZl7Nx#Db!+5yL|;uk`I}IT+q=#D4$|*qm_<6$Ry~ZUrh- zE(bn2aB~a7i%){OpI0WE@IdZcv!>9$_orHJlrpHd(lGG?3Ilr`I7;K6pS}V*)9^?D zLPqKIxvNSW4AK$JX(3c=F|3)^-a3I>-ZVr>H2;zx)vW_=Lc;;KrTy;_Z$BX3|BoI~ zT}TWWARmr_l7c##3hrU@-w%ersPUSI2!VSCP=LEKb&~&OiYAmAPN}CQIMPkB^i*6BuOe+P(X=2W)+^Z_wC!~yZx*C>Hg}^ z*+-yi)mn4DVZ39EIqw-p658Ml=~kK{Jl27LJ&Exj^k{@3yn7JX;JI6Cx4*BsbbSSf z8MyVSAfshdwDJFR_bBdaU(kd(+?v3}zx7nMF~fO?i@_gL-|QhUQ@b7DeuNi&O`hvy zv!fr^VSoEsEY5hCYN+SB*z16=6Elpcv3Bp5`-c61EF2CbLS)qn%ahghk|gp=Bn;jX zS0shtIN{b~SZ%i+>E#Z`-kmyuD}QhQa$t!7^)ndO(I==ZOni>kpe{;z^CaqBL1X8& z({+7r1VARNG}f@M@JupUw1Pbeyh}m$9Dl-s0gk2zaZ;TJ9n{-3X`pm~!Y9SAkq~n6 z(RfnPit~HlV2JE%_6y_PjB!UpSJk(S;F#gJKgZq^2L#c~tgovS@ zVbi`$X%g7Og*K4lkJT!*?6hh=M{HHrsQ63?*k_18I$`?;(J%XEyZ63{Z#Il+a(A~h z;twpWbB3=G!eI29>Wk?h#2>HC#eG~NCqey=UL(-FG@z>O<~w@|EP-689fWKMzma>r z6&j-igX{Y1(Po)Od>JK>4uIhSOkEXg6BR@Epz7MVlKDQRuE{!o;EEy)e`Hxi1`+W1 z$1$mk) zh2IK=oO28<6?bfQFVWVuV5J7f_8cw@u%54J%B%e{YEno&PSa{DHGX;MMot>mUL1w} zfF>Yw3otm-&c0(KvR%DCj~FHGmXw3xajFdbY;Ah>n)5$0XHWur3$&AsykeX69^C`k zXV8Hx#7|&LV0i7E>90t=t9Ib>aOir)kqTcuLD!TBPy`o5So9zJ@;yM)2!a27b{{iv)` z$nFB2kLkW0_0yQFL5&>Lu|J{#D<&egwa*ytEXm%J^7GSaaE+=$t$_UN74%`6nR-ho zYn#@7`|2DbeX!7G=#u!Iw;-^!x^^-PJ(q1^EOBO{PgF_ci`q6sIH=be(<~pFyB=Vv zCn{lmcyeSz#33ryh8=lydRt5mSaceTg3@6!rIU9l(|q&*QG~ehmvcB%lw((hRYjft zIS`${zw*YLdf^9Q=O!lNe>=-x&vM@DICvp{^DX`_u!yx%?uK89?$i8|Ar5J>u&7BRHcF{EHa|Ze z&5n@y`Zm}QcUt}z%`eS2Gz|UiMG#H;GdjHqmra{s4EB3By8iTwt(k_)$qUeOjLbg8 z?2QG0c8sawSAM9(&Wr7H3~CaNPx zF2@nxL+#QQKui}RpJXL6KUX>JBTHC6B=#qKFTKsW!+-ST;ts+BB}#3~`+NB;_oLYu zAnH_TL5(KuxL(+ae(@VpFgR>}`8$*^7XnfU3CET`>5}o#ZawxD;0(6wG_@!H~RYg7&ofvNoxrbjq zL{6iW5-EAO%l0XVPgUjjcnw+j)n5W*sYbvEsC-G=Bah)nIhZ$TMJ>Em#b8?LP^V?0 z1)R$QC|=#u!%M+>9ap)t!=QlDR~s8N4d}Hl`sCKNr4RPl$6&v{Zv_ceSBth`Jgy|% z=spP$ByFiG?>UwpaDvhytt7yb1!y$t&r97k3ZY=PfJc^kvi(lV|5@zGLJ2u+M4<}g8@as z-arr(^nm0Mp7@^taa$a|7cSC)N8ZUh)o>{q&&4kEj9g3o{l||h((MhL!OYD^P$(tV zdZ<$!W&ygjfNgIR7Q7Ss8sg$p^C2h5{&$O4} zf)cy;N+IVny+41tVXm1_@uMpDNK5&wU=62Q#hT=^;x{$DX19M)e(syuCSEZ+l$Po6 z`c=Y1jR;NuNxjSqDj7OOhZwSbBHh@I6uC^qn>s(Zo@E;LRP0-NRTppkQMJjooQj~S zp$nq}(fcdzIi{a<%X0j%5P*}j2ebIo?-HL>K{t{;v~W3_CX^s8L&SoGtu?6J3>)RjN#gKIx*FDk3DELd z9N~M^?1)qoj~Z`qmUzh2evhH(Lr4`;d4cxMoR0zv9_&6J;fm_OftI4Az^mA4TE)Hg zWmT+aE76jTlbR7rsC^K$)HGtZg{Iz`P}4rEP(P9)JyV}%xbG$LHMSk~!VBj=6i)UM=`#*31)FFaEyp{?WX+`EwuqdlB$N z>T3bpfBY|PzvlZr#D=u5J~172m<|*Vr0Gbr>n`s;GIw)U*(~9`3?BMESf17QqsvGF zLfs4z^*KoHl1Rq~#aUMDEV6X-cj*)H`njLpcypMedwK~OHX_IC{>nouQ!yi|B1#sc zB0i<9JUrpQwIJly{ZZ>?`$1}46OG&-wGq+o5(drgm(Qu+f!Tvelkc=RhB_C=8ci(0 z>RQ8aH4dW{(uJ;lNn9FrP2=z+8nMC#(4nT25(Vt*c&J+)Dg@Irn6#yRB`3Z;6FMb< z%#ZccFTavp%FffT7F>ibV*q8s%V*wu%9me+=35&z{>kH=z?oi>wL<@ldFm!AT&aE- zu)Kh>LK6mU1L&9>F297eXcMY9T~2BN=rOgegie#VcJd28{j3TGTs*7_>%ha^i;+w` zpz!ExZCUoQuJ}lNlSkC15z~*v?81C#l5}}jb8b_$vqAq{2Q7qA%SK3%P^S7)l`mKt zSLbU^JbzBPPsT;I?|z`Sw<(N8Ls>h0D#PY9Mv9Ji#nXau`Tm_Bnr0LfRTxr$)ObE-FRjm zS2|D@bvrP5BqhGl^Z#0Zq?7%&bC>3lI1JTvsw0l?kaEM!ZUA)eH2S4$jLfq%XV85S*rY5g2<&x&SdcVpi9T*HRhPOvMjo=q7EAFL z#)jNq;akmYqqZs<@xpWkbWb0y@;mulmp_lB@m~DGWB)t6Bmuh40!$5-r*-8Af_&8% zU!GE&7eQp{H`i_->=uze=VkdWY^KFOwNDJ z!q_KcNEVlpf?lsQ^%HTe{S>nMJW~KSrw){2EeNXYIw@!Gm_MmZmtJ?L_X4%u4Fv=g zaE7*3S3Rc$@);2e&$|Ik%4PmAOZ!N~U<4)At*|6}2rbYSJ}v3MD5gP-;h4&X+o z0}m<&za}Ss<&nUo5uO(+BQ0>S4Z3`f`pkVDxdlAL;g?)LLNJn|mMySMl{eD6H2gZ@ za2lxKSsn#gKOi^F3r@Q@ZZ*#KnFv9M!Q(^)f;Dqo!IE6dcVMS zl2k6E6Ez-c&2=cGU)umQ#Deg!VQnnXVyo}c!y}=O`@TbRKYfmB@FaS((70^6JdopG zZaMy-l5P|EEY3E0O^P|zPFe3%BNi^97J`d=8iT(Tj@IxCK=j!|KIWKnI`-@s2EZ9N zn@Vo8uC8}-aTVs}<$Xt~!>F(y`!@}x*OtCkNZ7`uux%vVl@bMeSzfi~UQFSk2HoH4 zelqW(@FfP{G*3cAsO#kq*mmoBE9fMEY`l=+Sz0QbD*6jQ$|^7i2xkFT>ga27jRV&= zx+MV*f8RY04q{12n!ILo!me7sa19G7ahu0XrydNUVzEYkONN_qHm5c~?l^=B3pw%~ zrSk%kDC+o*mLpwdgZ2m1YDlcguG$tu+3RmrZVq!BJT@w5J)R3NVC~zRTZ}0c8m(@x zSu~7X?&PCQ9As&NFr@)gFe7L|UO?WSCAR7(6^ZFGk$VE=tC(8lUM1 zmfW@LzX84JNol!o^y2@fO+{ho-9Adgl%%LB$LH@U$4^GcQ-3`>1hxE`PqPviuG|n} zu}g_%!FbZI`cdK3SkHt-a}X{avMfreshETW#{) zDU^APQIt9?^nvgo{bUrre_%+;s#Qy3=f|Riv|S!4!M1XDmoK~0O4{FyT=NVlAa=3@ zQdfp5B*5-6zpLue8sr>*+$)N&4g8kLve!nJ9}hA(9-oL~SPrZ&hyYSyXPox8V%RXS?XUMNym&I+B3xUC1m zeFA+E+Kw!Zd7VoPqV0a+(i%`jg&0` z@4x+GI2P?4_fvaw4ZjyDn`>QBA|)IgM_g`zKcET8uix$PHWV*1n($aaY6|$)7J;|G zp<$aXzyCE1fDW$pKyPs(d05SB61$Gd66Y#8BmEasq!UovUY%^k;Be{L!=3Xr zkyG-_rEGx8CKLy<$v~p+O*IF8FA(sOv@0Xo*WfbnN_=PZ=g>QY2N+%U}^Angwmc(yKke5j3Be*&2N;eHg|A2(U5?w^v)KkRe!HXZC9u5o& zAI)0od+yM+Wv5BG$62J)@A4Y_>)1wO-e~5_PIZwnwwpcWhkkx7`^2`NSXBlsejs(sc#N*z5e3oe01jfxC z%NtKaHZ@p#9Vlf3Eaf|U$G4&TR@}OTECaaJ$sW|%RKx#QOpO#jOExXGKQ{P%>R~hO z=hBq_e4<5n1uah#6D?$99e7}osWCaMpuVadn;3z{GWrdx$Qe%_9-h}zKn}$;0%`f? z@neN^{GMJTFqCWr<)M3q_`%utRK13);TdeCWU$DaAO{@4H%Q0zKW7{FT&rBMYnCO7 zA{1sx43o4-&#`YNCNMb5}raxs*+u@CBmlj1lmhUA!LVHf&gR1Qh_3UyhILmN`N z57I8H>lr059)9S)X`4X_MkE5FacH`;Je}`?0Xi~%s+xzNYY;P7V5_g@ho01GK3X;Ez1fA8! zAnJlV>#$L2nt%a^Z%!pr9=UL`ucOnv`!{UXq@hK;j3qa+0+wEP*AUbxeNemY4A0E0n96?-4gGZgz+d&< z_tju$^rLGZlfvB0wZpzVv3d}>ZTUpJn6rGJaxK?;Tb|jydaf6HKTt2=k@rRCeIDIC z2Zu@eLuF2evxXY_&fz+8`=V;oEOj)Ps_P z0L?OMsUFTRfnjw62+xLKcqYR!zapHrgaC3dtH^8(>KT0ni^xg@mS7tRDOO?>aq-R& zV%LYNpRlqpBqv4OS{!pP=t}@2@h}fh0UN_=^O4)7D%z?H3aVB7+t))>Xtj0T;`O3N zi_JH$?AUdMvwasGg}&Nr`cygDRRm-Ika*|K?sY#(1DG_Q%2%nC$LM7{B-Gr!qCsZl z#E=J_^I0ftZC`(cF!oWeeg5rVnDmh&L|_L*OnRhMkNm#*<$TsIcG{YyTUDT)Ve{mo z>k071d0qwYF!_XCY5{|w7<`MWBR(4j_#8BNdV!hVB`PptWOVGE!aU(ruBe_1eRc0J zf$FgdUCY|uHhnwp&|DeoHJtz8wcUMrO})4#d!%PPM793nIEMNiC`cP zvwK4_(kJ>=Ero+G#6TBFU3{N0Mv92o(g>B2p}ozib2*9UZ`)xIHze;d+TXZC%lrsm zp*>It(oMGX-QjyZSEyfNDz6g*y(;<8~G8Ak?~`2P%iZ13%WQX;NVJLQX}JwosI~`^=>#A|(;G#a2-V=DLiAJ|Gor^Po<-&D7iX#1=0~eyV)y z3M|L%kEHGsGkMQG2`dg z9SX&>u(taPysC1w)-WxGV?C?4=k9a@%6`_xjdqht7sanapI4XVFkl8w-p$=_QB$=D z1OHV}O{V<{aIK{Vt^nTZGHea;iOO$J6UPyWw^%QaNLf!_oS#uyX%p-NC1p!Hbp9q_ z7aWs)F!vsp@E(Xi`xDyz{rxR~dG8L}K>JbGa{P#B%XeNJv%q-k)Z5cQifdxj_|czI ze?_{0N5G?&xEJ9GT7Ij10$l6qit_}ZISt+twR0iA$ocsux7iVimJ4^KW~Y9Pk@_D> zYJzSNSBH8cA%jc5(# zh8LaAALA(Bt#q0s5G`llRY|}w^UBAUTkI+X+ucZALkNOnsPMP$(W3p~wr`26$}!*~ ziwm_xo{jMUCLXdoZ(q>?@Q{uf@fJTh+KrSxfH?|)r6!QqEy$>!h{4BDpK9UMrJb-r zwCpFqFdpOHf3IH_IisFf^kYp}j{|6?ZtT`Y<75fU-b)^bUy&D`&}p6|628dPeuone z7F!w+=N9>ZyIjb(jhBU;L?WBy3kc<+fMCCn8q)*)`<#w`#Sv z-2ob(5CYcNhv4d3AnSGFSBtp^n4=8TQaVs0oST23a@kHww!u4I7>H%}K#co#15WX1 zg6qRhA(VHS>k9V3tZ@K)T}h%e72OGpZ`^2zMb~c9&Xe?)<}je$*PNMRzF4ll61_b6 zB81Y2Q+h}NkJ`ry?iw!$F9oDfb3z_idhRuYQ#i4_YeEZUMbMYngsgJ<3N-Isy_r%t zo2egpVE4>Pw-_W{A*1QQs zsJys)%mtqa;`czDzdAMRFa8EKJ_(g;J&RNw|5X=_JW}YE^F}s+ zH*bQ5Nzuq3x6|WNCqbNHX@ZGBE3$2GULR&TW=?JIR^DYrgP1cpSo!2oL?j5bE=)7` zpJT>A24fkd6Y}VP60@7IR0(!`rw2m_EYvQ;>z)iAg;9#6ciM| z@?mE>3PXY?eG>Ein~`UZrbdha%;w=fjz5>yPe1=vQDee%usJiqz5p<=jG%G1!kNP( zEf{|UJ#i*RQw|s}=m0s_l}abhuYrN{&^<3m>LB~9w~$;#M-ot9jkemqqrTDCqIJ>} zs772dZg~$kZrRZ0Ivz?+dm4P;3cW|drY%5|R4M$-42@I2e~>}Zdjj21D3{azD{KH_ z@f!OK&~vkfl=e@HsbPhvaK>OrWBcwrGK&M;8a`p$==QY}jr~*F*%e83&xM4TM40U9TP^T+!=6Ou#GXlZbd{ z=2XBQzFpx;0$}ag4^{2|UNC=mFjc;)PorO_#25H{4sDlsINPlI&@Mk!9 zt8C(VB&x{4c^BkA0#)l3A1pf;3`LEeE?3YBzeTKwKlaw+*XI_GY zYRoG|?hnwf`SapFK0De(8csS{VCl&_-tQ_3cu5AsORWe{i{qodWsy=(>p{?i1%u?^ zg+Tb&HK5AsxN^pfMcLjK16Ny;?!wuukvHlv6(IGL-i|^BBXVFBaQv@yXXZYyWL_LH zL@5leG!{Qo_Wed`D+2ceSPlqxCM@y<*o#sHRbNr5+y8+4%*%YqFc}&a#>kTR2mD5h z*mNQPIL_n0^MlNiIPcTx#H7y^ZokDmQbCY32NL=#{mQ{)5Eg9GQG+y~Iq9cU)Urju zZX)P8T4edUTomB|Y<)6kNxk2MgUT}6koQ~C=?H0Lsbf9=9Xw-nK%gs{@aoaN$6^~w z-{Zel$d6@1^wnOV&Rwe!kgVgNBBi;{<_o_+nd}{Zw6A5l4Va&5=PxUKMJw;{x$Dbh zF96#v5L-)hYLfV)w~xajze3P% zC5X!mUrs$&$gjLd?TiKPC)OwY1kxfCOF(vZhiCp|Q+rY@26haPvYs*s1 zM4#00@$VRf-sAT~HNelB=058-u48|`C!OwWn6{ZNaS{?UU-!sC*?sv}(@V=c`zES> zj_UmQac{*bzIXYG(QUZYC5|UOlmS05SMsZLxeqH%kM+rdLfexSttQ8-8s$3~uOX95 zmOOybSNW-WVz=|8{>eoNU$)CdKGJQ{BYfLpR@NJ;`vZB0tRi#r_M5I--q+)JKy!;X zV`1|nhTa#Fi9<@ihu1Lz6}X-gv~fK0O4c!FS4eykZ}Mce1t`o)e31Rmb`TqrgV06w}5ZD{`K+@SR*$SpnFG6iz zEfGq39soenRYR12#4}K@X2byVb=`t~)SoNsug*221_}(kPw#$+R&k;)XJ%x_AY2L~ zg|Z1^aQ=s{VlgxITRBqs$wB&hrYBtYwVbb>2v0Ji;-bn%8xI`xo6hm-Tzti|Q?|9( zcEN#Hw%pcxUd$?x(M3OqEM31-KN2R(l8aY`(+N&T@)9^2pn>#XUbNwcc43apn$%k_ zT_+vZ=BTW?H9=5fyqNe)fa{^jP8K64A|{6{%=K|5)?zN{sKm|>7uW1yUKiT+G*Dw)D+e?|ndB)+?`?ih8=s`;;x z3lNV9C`Q*-N#?>DPcA*j!KE2P=w?lLNj-qq;uDluG*ns)2F6#foQmOe!;unoz;-SI zTqGEvsrR$27t950QZ6q;(Lxc1e7!3{sZ(WA1*%|Cw?L2<(kXV59^$qgG8KZP0M4Z% z4@ja-_Vs#e?RV8)51*+5Z>_H_KOa@*d!#TkNJ?*-@3SMtLH8?69J$S7=DIp(NsjXH!Z^@fv5BiZNe#uUi_{RwvhlOug1>gPv%PAqcMdoFx4AT`rN z8gGl@wcLQX{*zM+25oSzXi}Toq~In~T85VXj`k`1b1in(X)4Zfd%|=bt#Cle zb!dTh^Y1^IujwLa1I&hyNj~AttZ@jQ{l-K9rw*h$gI#G6zbEa>$zGx;5QLSuGW0{= zeSs$dc9<*cNyZ}FNNAoeaNVW#pdON(Pv!~0u1^Kax2gHQIPEP)?t7Bq6!o`n(&!2M>1L)PkVzggk{z<9BKr#+g%(M$VDod{(G+hx(PGcr7-%3gy zjy@g$R`3GrvjL;jILw*VA!kXOt9q_57EEB`N~X#94Keup%RD_x%X=P@LZ<^VV)J6r z+=YFf!$7aF7VH7Nh)|Qf9&pe;X2c8S`?Y1P9oid-c<=9c0>;>wO(l#tq{>7M5e$%+ zL|ol`OCbtkK^>6 z7l*9JZKh89|D{Kn)c=46tG+R5@{;i=+A;%z-$K#B`dw{K!HJR%Wieb)z82<{GaF-(1=ckW4x7Wa4ztC zJ~y8#B)Av~hXG(NsURgrvgSCSE8Tm}qKw!gorkYm2v44aZ5Q-x{#?a0*VF(8(oV8fRmJofB{+~by1ArCO zamt9sP=Ge)TorC7=+R9IUHZ7FxVeRS+2yKmIBC*h>>B(~j{0df1jEo7~& z$hxcM7u(<8R;j1Px?3a>j7AKjHd0y0`fZQIZeJVynawG%y5xqzpc4J4W?21r?x#-E>Rw8|JJcekq|EkH$8_Xg-Lsg(U2Xoep$Vvu z@kxUY9IN;>qgM{;(#0f0jZ@A1TiyQJ?M2`86iNg1_wg8A0I2`@t5Bc&3!M-W?Mr* zsMtF|+d600IQ~l>WC=Q$mi%@3(_Zo}-Np=}bO#Q^K674O!nj~biQm8K(hFy1gY5P$ zk#5lF6e{~+{>*m88%d@yh9@`gMn4_XM2oY3Inf3$^+%}8910OyK?Bx3be9?QxQ>*L zOBj9>_8UrHblH7;$H=r3ZC`YyP}ITZ>_=`+2fVQJi!j-{1lxv9uK9r{8Y^neo{Cww zN~X`H-V>h*VGw*IjS@A06a!?uOSox!Ui>7SD7|SFK>v2)ZRQO!+JGp@x#a{CFlteI zHUp`tISPPOqP{|d#q>akb`ji3EU$=I0>>Q*p?qHY*GFwu)$^&(j{2mbLEl{Q<08r@ zOAKDx+wyjom#Qxa2uyjurvgCrptdyoo(oZ=Cr1v380f)%_dLtyjE~~C_8ONn1MgK0 zte$HVJQgN(;fYac*c+U?7^%DlV6-+ZC}pewAEY9?XHrxgPEh1|B2{koN1T2RNn1j# zFFZaQ;N(uF1fV5t0%g%kL)Z2yKchMsfXK~NvUKbPOtsddK$N4#UHAk%@YBIqhw@c;vS=g<=>XYyt}hPIVeXgMS5hYw zxhvK&Ty0LL14|gt0-pm0Cn@%G!S=_zIF#oZ*7QjUApem3PUc6-*9B2kqTOjhHD-+G z{h{p7+-qgoM@k)_s|_==fVNY}1EV2FfWtKAYx<^?Qgg>WUm9)!$SGwDYd(@%^K}gm6=Ud^d`3B&Xf&Wfg z!wj2f??e|bt9bku7V3v8!mWCn26UagIEP4Kk3@1BXF;)U*v`z5No8zNe|B`fCU&y| zO96anf^&g0U~I{|7YCNagP)%;iC9i$(}CNigZ4BIGF*$f9NJ|;F~{aEbK8!pin!R& zsgl-)=+Lxg2@{>LUQzL+Vu^MADT91VxG{_(ma+r{$238fC5LHr{eHI`5|jV7XvD$W zH3v#YPFnaPeE24_ZP6vn6aIWS_fC&gK)~eiovb2)krHdAe-<`(w;b;$oOaYgui3v4 zH7XAGc=fN3$XOAGRM!9sx+UQP5NGZ&M2c*mnW2{KbkZYY5@~^%(Vvg(vf%gRB@7xO zl_r#Ql8T{G%ru&2$>0MZPDB@0e>x!4@I5$-BothLgw#R=-12`{{Af#w2C6ne@1fva z(ciqpfEe}nBkLaf`Dbu)g{yJpkQhdp#k!i*8?fqyN^VnPA7T=+~*f) ztu?^`VG`G&{jke+No=tzq$3IX{8uz+|F+nO5{rsg9k7hdySc99w+t#)lI!1v?5-Xg zYvGx!&`WuYelky_ z>d{caV4>(54`6Duc>ENY?)Eq~3hQDwj_o+g^+$MgWrm@hNQ3bf4^PF&?I=eC5pcGV$=lxs^XhaMzG=bppznSm ze$F;!GP+ zqdN^5v;iJb9{3xmYI!So*UPA~J#u7F0MQ?S*Gjmi{32t3VX2H%q=k2ZFKH2v(_@~1 zT3v$r!s@5Un-iarP{O)@N59k{`w^YR`A1m7B@wYAsbnM5uXg40@wJ%Ffea?%SrL}B7o(QJMV z5#puHr+;9;)2`|0;v}B*%wZ37^EG2+r_p+(6O;)5YfVc1SXDiH1p<2PboqI1+@%#A zg&`8Q+_yC4GT~hPsL#Pw!TyO(Zr~97DX%C@nu`M6Ig|>ya&JbKe}ti=BN?bKWec?{ z=XJn2dKoI(ULaM=k8=n*@t`wuC=^!{&YuH)ubiaGGEUg=#4QMn{OqGWdD2ustvAnU z_{`e3Sy;|?JWf1VtmoU17jN#qeFby{bxAe}EB4V8%z!WyP6X*vNU|SOD+~0dR6O@0 zXnBK!HuhB! zXv|w0cW#G2!ulg^VH4)Nlse7V9aj<~{4|w?>{8(;^Pz+c6WDlUAy=)Bad~zVFVI`6 z7*KCjpqHVHFyLDkfe(F?C*35E6sH{fy=kbsWIFn>&`H$bBb_2#Va72!nKr*1Ym=|W zY2D#c@0d-I^*i{*tzKX@zv;^-mZjNN`oU-91|7SrX5a1jUs1|NtjlH)2aUMw&m6jZ#2I`@|a= z((xu3EX~+ouCF|ols)LBMPunRT}^PgNO`c=(JE%puSM+Ga|lNo=o2@%{#0I@K>)v{ z7QB_Ucn{C#>*vBRw}(As9vMnc-_R~y`u*cSiSmCT_WRFL!2kVII6D1DvE=`AKaFE1 zP#Y>(k^S(sP#(0PfqouzF|t-$IiskQZIG{H1_^s2U=F)xR$Y*E#tEl!5VfgS%Ut~+ zdsSmGZ6W{N5h~~4x*`xrEU{wM!FeTTK!u{ZgSBIPvDIA7yQ<)(F;0&vb8CY%#R14Y zl@@_j$#h*~@yT-p0C^x*2B_gIVf2%}8+N9~q^bh93-jU54(%5h8+cEAynrH(Xt-%sZDqvsRiG4~ZNWza6plL7tnrTUgZ-t5(A2 z8v>WZGn*(*VQj*GDlbQ8YuGN(%9myO2qp)1= zF^5R3nL{LIH+clio8hP|9q*zoG2fKM2Cq#~agU*qq#29cZJI+S85bKIzAtVb7#r78 zpEop(a9WpYuF-@#g!CCe#DnAZ)4qb;``J$pp2T{VOyy&>Q)Bpz!@thlLI^wp z&F=BRx#^q-oi|0@6WFm9ND3AJt1y|`xfi`2Rn!hlh1e#LqA@c)oVH#Ww$W)af;Mgm z`Q8h-Z8Q;#iHtMeh3>54d&5*XL&R7={Ar73;x@XPASLkq3(0njIr1zOl-d|I~EMEs!Jy5)(( z&`Xp-wIK)rZtaCpot9}~)8q9Y_S$7l2)boGh|o{3Nx!ysdE#vCaj}Fi(PB4dG}4`} zd6{Gd2RnA{a%-I0#TqP8!JISj6H1jYJ=1d4io09R{djL(@e)2}W04dBH)lkyx&C8h zE-2qks2FkigE^yL#vc*cMTb^I{?4J7Qj@a}H{N6x<=5U;<}kFf@v?957yYNQ{sT*| zS`4{Mq_gmc7+t;l)?sjR_EMmcdXA0%fYhpF*V%Vm>N%zE)wvf>{p^u1&%;^0bPgM60s~^6|TFnF%D; zFL`Q1Y{^Rc&5uO)3{52;V?Ubg-6{LN^wytF;BWt+9{ErsDEdv|?(@&|R{3mJ%SODFH%24{Lo2_q%*Y9Awo z{iv@7V9426@O>uHtJ>GFbZ2kG5piChcG6Z1r^o1i@c3>|hV6HEtGV5tEn!;ad3&z5 zmcCao;Z7PuU^wq{^Rsb=jw-(yR04z9K#jec=Q!i^8?ay_A>GWMaT#~VNamHjei`8| zqfc~zG))Vcis4cOEj-I|^V#~?5Q9-dPX`J>NDbyqY4PTCb|6L)sSHks zb&C6b8xhw%!o*NVKV%Nt8#M24vHOBe>OYc0NHHu;MUrX+I$(n4HVyOaY)DP*_L@rP z`9Jd6jJ=tJqXCN*_fKCS`HwKn_%ssw-3S#XGgOkZt#qiw@1a#zncgkp+$(=?+iBzS zdoM^Ih9ayxkVCp~_ZFEQa7A%^^z;MlP*NYAao-Afb@s4L{6G z{ZJy5YND?QQi92rc^j0h^^Y;_1me2t;sRIQ$w$ShNcv~@4!Wd2-{NS&TKzYYe@NIdx!^B(X4WR^M zD9<|jfPUMJkx7wF;kZqj1c^sk8u7iD3=~nmd7Dt*RxoTY4F0VUJ8@CyB) z0Xa(+>Y*Dg5Tzz~su;bTA@V_@V#V?Z-70ZUULFT%%%(bS{Vww8_r`*aN4<1HVq*rM zdFgO{^ILORnORLkqT^i88Xbo|GakJx_2-&8W$LR}P44$+TSs800EDH``jB*?##X~NJ#0x5MQtqBc@Fs6C){j9qaf~e{ zPA}h04*_0}Sm;)*=HI!q&U#YTi23Y|C22My%cS0;*W51kB4^1n<;w3+#0e)R1WhfF zO1km(3l>)vrXSy2qn9kbxbR!u%uoIJdqotQcb;PA~)N_+|>sV*eNYUS3`?fGlFb zn@L{LCd<7X*ayR1Ctl9 z-M4$jxK~I>aV8=>yv%u>q*GFT`hk64hJ*}cX1=3FRT!@!lewIJICuYqr$N7d+e_2C zLZA%d@WzM>v#FX%2A3^jG2PuGB_&lY3h{SY&8UQ+;BhFIrN8)id6kLHLL6ofu+jc9 z9ewkl(A!O(aD%Mdb;d6wBn~sE?M+^{Zv8l?Cv(;~ZI}POd)%hWzJF_Y@k|KnMvG@~ zOPfGVArv<6=^~%tGY1bIydtET^w|6M?YBN1 z86{5V&$QCn29LuLuvwB?Vzl^MYYjwITHD>E1tJtf`4R!1?@S=^;xv*X&K0VhTe!Xk3wm>J*>YBOi=YqklALkyjkl{;S8}dl zV&b~f04uR~PAx6$n&3;tTR~B#cd4xBBg6}3cMm%+y1t4jvmMDm^azJxk%+?-P}z|Q zpw`{nv-);I;obRuuZLA?hAV$osZLv2#V35>?(J`4_Kv_hy!_Nr!z!mEYjPs5yQ#k%v z12k7vP!mpIpg8)=G7i)$!U6+x;ogHCCW_uaIupE!oeG<1nq~ZjPippmIK9eQ328Xepc?FA}@MXi$OE$WG&9z;7@aMv% z%a*;h+XxIi4i)C!^m?V4*BdmAVP^^~e+e=G({Jd8hzwWTn&a<~RYaLKpW?>EIn|+NN{FH+? zZ86vw@=lp$!-m1Soy4f4dym$&vZj&A?0I)eH`U%;%*$^KcpNILTfZ@L5PNOV;WR9;1n_ z9}m}u`@G&pj=D>hEU8j>P8#++DXz^s4jq$jWHUz6XNPtiNPf}5<#E_ z!|bk9GDyB4b(l%!7Br6v8ZD%jvE1Z=?NvWG3YJhb!iz$ED|(6w^C;$p^z?EnMhmP| z*I14((YrXz$u%pgz`omgMhf1yK3y_T0|a%bNbp-vlAM-w*sm!-E5F|^WBt58K!;IW315_i zi-VZOcN1YrA3m&Nf*h{GdhM_JHWWq2CNK2X@amyF4kl$@3nraf5$GF7+;jF}pd0C9 zi%xS7GROwmxEtLs5(DeMb8@}$i{osJHr{qXM#i(K1fzjpjK~aSOxNaZ@?Jns--Fv$ zVH1HJ+=F(iqOY*9a02Mm)H254DDRFN68;&p?JM1VjmB_}@r#FJ7+{tD*2;E>KRPKl zQ?yV(G`FTkm1F&S{;T||l{SmKanluwR$MuE=FFLh^Y;P+x6I?XDm-QL`T7%R-|UVq zSEx*#v6YpTtpr7*?#YuUOwG!gq~3Hx4Y*_ao2G)C+(%3)>oDQ2MG6WU;9Qk%-9}uS zbcpOmG2c4yprOM`u93HYu^&weZ7x?nAoJzMW~||QOq1@pzHTYBu|TW~FboCq11iJL z{R}c&ZZ65J^|+eN*75#k!-qI_)I24{#ZE~d`{ZA|%sBVTTzY){x@$f0ze}Bv zL}cb(*WPds{`qgCa~1_$e_n?S=dbzmr>t!2lWxqtN-O{I*9!&g8WG8Ae)d$?f~@ik zsCyJfkI(7m&sRfNR9oz|x)!tj&%;`Gke+oO3L*Cl4OJqP)q?Aw1`ehg_Bxq?Q>y^L zprE=Z?>YbFc9Cnk8?%tEW*{pmNvCEAi#*#^5SenKWymcO#N3#U&dz~AXE??T1io-p zg?(bvz?YYo)cf`#gtok-qi2W$URgHDxOQ!j`PZ8vh7rF$f6G3B@Xf-^{2UNNJf4h6 z{=mcT>?^Kf431}>?HE1{VzJQcT_*L{KLicDg!)UC%;sa19_{OeV{xWS-rJiZo+8Dv zBE^32g^mIxsB}mh7VTY*31G$WadXR%bDw~UY^t6T$X*u|+_5}Kv2%LQoy`D)RZx+t zf`aU8H@}w`i?J8Do;r+?HNYi#C$>NBbao45xG|`DS=O&Nm&k8q`8mGgMC%9|oM=$3 zqOf1&(7dRop8I%^ZU(r>(-rKy9%Eo0izQiu0gf7`bZYwP7+Dg+ASz=)-jJ78E!qJ> zdfIv!HV1IK72pwcnC%i7*{73JQdAL#x}9O(L{jI1c<4@ILHdz9A+Dr z>cr~>^We}4;0HCBdxBt6)E{VyiL2O!-B+YTLvP$zoLz6T4@>!BPmVE^=}bTS8&8I{ zWjiXlqRkx9|11J}%V>K3^gFL-Tda5JT(bG8mjRJB`&NcxgcIK(SJ*>L=ia*3p=fb| zd51t4j^-7JYh7D3g81~ zh9=v4KpRWbP}=^?)$50~t(etT9kRP=7aAIQ=J?28%q|=yS~c*ixR5Y(u`Yw^ueEuznX6 zwat0v(d^XGZk}%ep5wxTf`U!IZCF+xV$3PL+wg7WFTwlZ!S+kIG?Ol>E&JYAk!wb8 z)}3iycScT481#edx#qC?7CQ7H-#5mg&Fvx*+8 zW&LKl$E;?X^I-g**1tTKYj)j34u}T7UT)9_7^uOn0toQ@U)*t7IUaj{Hf z;DN`lR~^2f^?@()JrZDe_8vFiyZ-(;$%(&)k4U&`jMcVPs!41g82ang?cq&ZOYJG= z7JM+vKXqyNaCFpDg$CG$p-%)ZROs#c*eo9OvE(*^0c9lH;zGY}O}sofo&h zoo@JS|C3~_*HFon{39jn&iIHF-xuMQD8D9<{VR~Qb9tXwUUaqazF!O1CzbLBdftn& z+4S}9_L9i`;E36Y4}ceU5+OcbY+0lHvqOI279foRGVdbKNm3}*XL)e!*Q3~0$>5e9b87Y$@#hC|GJ|e%G zFn`Y3BraYX?mC*Dla{}dVaJxgLT(>IGN1Dp?A2oTO?^LHfUgNP|D=Ksz!nL8!iU^Z4R77jzP10d-|T@9BaB8eBJykxL&*L>kcN_FUpI zH8tI{ZJTVG+GFW?Wr@OYl+HCQS>f9GW_i$gu>#C!jM&C_vh8<`AWT*ylcD0ey{_LX zH-T4e+L#Z#?woaZ0o`rJ?OMstkEQQ>bmZM-k|^{*6b;@7N^9cadgBjR7?DG-ki3oA z?0B>wi1fL13dS9A(7WF^{RDD?ad1+Sa3XtR_UV^Ou4vqyV41(G#ikXbb^mx#40%S$(ffn7Ec{n&i5S{sc;iUOLtuNi47; zssD_?+1nCv)xRPXQvgVX0ut3S!GK^C^};B8{<)18_KmlwtM5*FjBOx_Q@Guz@q{e0 zUw&_b7SS&_weaGlC*b0lq_gumrS0E`qY>WjENi*JxY=j|J2Rx!vHujFy{Ma$z{)8b zNfcibz}FO6lKk)9HO~HZ*-kauejyz>Pub|VukR{TS$^2iEs%F~^Gq=-IC!WM(QU_j zTUaP{U;i*R{dVy3_$!U&nzB*1pAQ5sNWjgg1yCa5niZc$i;oLh(radeHxkeRGmJ-r+s8s$*H*FpdWM2KJa-T;X^{%R}yr)$7mS z6nA<&@-23u!2eO*l}AJUui++TGPaU3mJwwvzlaerma>H|$+T$75=}%jwq#6p8Dwcj z*;-`KhC;tCs;Oi&)}KVFn88>|xJbxW-S>Mto%8$io^$`Yf6tloJ!ih3`Ml5bywCf* zYK!pD()!(wh*sOuqkb+;HBB`F&ro^v4FpiFVT+#GRcQ2>^v-#UkCnU!sg=2gzpkaf zATCtf8{ac~u#8Mn;Naep0u(EpKOsPg@42FUZ~)=*?%fyBd6SEKWgelb*N44%S`Da2 z|BRIN60>x*wjhvo<~0jU-uuBM;vh4D43;{b-qbY5JcbcO1BcNVe8XyAC$I{3=!DLS z8w|`GMC@49BSWkgTo>s6s3oP9>fJ}~+}n@OeY<}~!#0P)0ZojP7TUi}EwCz$^1O6DcH-rei(8C zm_1>GOPAh8t8{tp(vBeUi5~lersr49K-kd<7HrJixL6p?yQ@ z^{#eYJb}NSTctI5Tw+0z(pa=@82vIm3@l-Hft6};YXD)XEQqsme1U$*n36WW;B-O3 zGm2>d0D%;|r4(WQ`;E2umq0RJt~!n=RYIA+l<2WTk?c@oMF+>v)Pd>!$VO5mAU{M-$=`ls@)J8U&lQq7pXb$wk4jkuBq2UhJQ;?Dz^0V^6|22(`Ew|0kFU9mlGMY%JpH^S%dtqdi3}%rMKM zp!q!=!`EEdw|yTV3s!Bfe9U6!jqx^ISL&*TiMs2l+Di;tGA$)uEpxKl#ktA+KZjR` z;dvEute9w(aMfl*P;;0U#t=;d6iNV;JYB^~pnmrYHeS^os?7J|WGe06MZwz?jkf3d@|!l)}tkVRi3 zE`I4b5&*+39kq80P#nL<9!R<12FiW$aVU|SH@K*|Rx|2bYtH2wT{l1HQkB5ldQk04 zbU{`O$oC~6pi;}_8pZ;IZHLy(RDF`)HIQ%MDQwtWJ(uHla&F$!b?U{22!m6??H9;i zegXd+7s22)i;Y9k`Y~3zosb;37vcu26h_*6WX1k*+c9t8H9d68j-&m}>U84!>0*?4 z)?waN{!45fHuD{&5yrCAjRx$UoLs<(pKT+ok|YD8Ncb?AZuslyCF)+@5X{2K8R;0w zR3y%l6%4fefx+D_kU=t&W+Bz5BsUkEGtuq!*{^?uDnS%9zagpCEyC_uKfaWlq#3j# zJa})j#tZ5`H{nh=v#sM+TVrsu*;6}|wllsigH~_~9I~h8r-b@v)~&`rY%2#0#SEXc zdU-MHIa+bh+NY+a9btH2ei@FN)SzPur)5A~d3rVo-}YLM9Qe{Yknf`u?rXK&3*lA< zU0ZwxFuDVd)I-YV!y{cEPei7v<`Eb}0&IMULbLZXO}Wf-*ue2h1V~8^i0JL4F{1jK zw6GowWmL`V7{I`*|KdOnS5D=#jWq(3YbIgMR1b-lEmG9&7rls8FygZ<=H8Q3iZQyN zzh9DQNIk+3Cw9JkcffdZ4ZX247NbU`T(q^}k1BJQnSrRMW_lb}e z^H8F)C$;)bQ*BkQi45%y%G%m6CN{39AZTxAc#z4uy&4>kB3sF(W5a&B+ME5;ROf(n zXx)CG%0Mp^y4g=`|Ahdsn<5hci9DOpjdQ~+=vm9waqdIV%bZi^6n!K^NX}Ny8x-L-8i~WtGE3U;| zP^24s(F$~wOl88<@%TOENzjid;fH+g!25BK7;*3MnOp)z|CZ_|@eDC8hLE;)p`xRs zBjoh4WF-P6vQexmabh5rr|iHMjgW5Fla-gh3vEnlwGzE#b5b+hnPw&);a>Aw7n9+?Z%nWoHO4D(Wr?1Ws|&E zuO#0!b<8KgfyHX(EwP5D4pt2^SZ3HBt0@@cV!M5)3-PV7lye$8FQ}m?XZnEbc6e^x z&xSdsy)<`1uM4-!8N|B%>q32&i44bqxO~}~Eo2ZV)uQSvH@`x4r9W{VzRdive#GZO z0}LRzfU?%zQ3a9Bi+Ih zmq?T4MXru8dk`;bXfSc07nlUVWg1~sQ#W)HH5X8Y4HIwY1oVAzilT)Ww;fjxLv+-rG!GEgwbd<+rY+! zv=czad?+O@-U;6_ArM|u^Tg%0jlfvoNNuf!JA}pKWMm5ACOi-|!6%+^V1^yx(x(`6eRkIb7CcgOhgKD?nM-hkqcqEwZ8rbeC=u6%~u@QV6>NdQtrY6#W@EoYzd+ z0Lg!brsD`Ws7Bxgy2E}IqDxsh7})6G{8$9oc(&!p94_YRFc>u(gY)4JF#>St(ZgP- z7*ndZv149&oech>J1=ujbMf!VJd02H_%HCo6xGLSGEj_NBaYX$uA(Ndx^o5@tU z5fPa!zTpqGPnw|l!l1cx`LX97ea~b1iU$80^&0zKJwPohq`m-Frga-_B`ztwRs6r+q7`gVFeC=jmB`jdK6bm zF983j0#0c`?uJpZ?MDjx7fyFSE5pXt8EhL#J?+VB7gfplKj8vcK^k* g|6d=W|KO*{&CzvfUBTDCA`Xi^yPa(ATY1L+8^^QvEdT%j diff --git a/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_4cases_memory.png b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_4cases_memory.png deleted file mode 100644 index 9e405517e0f1904e14de27548566b758e364d4e6..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 42688 zcmeFZS5#AL^e-AjVXGTav20O_4Y$&iDjjSTDGJi7C?!OU2m}a(A}V45lqNwy1cFFY z3DN`D5a|RG2vsQoLV^&A5CUh${lDkozTLMw#yJ_oVTKY`)>_{;=daBbeaTc`Xp86; z1Og#sXmHLPf!O#Kf!L7smjLfyN*3$jFLl53HhvapcfWwEx7`pXSN(3?K>OW5x=Lc) zZu=t9Hx=bhos>H*BkAeqcgt5@Uf%nE|AZX+wugMZy?-nGmd&>eY<&@ky*qh-_&hQ? z;aCJ2!*hC90$)rI2nO!@!l|28-F{|5vc@{`l~PXhF1{Ul_j{Sex z?AZEJNe_EjcYEBS^A9C0)eW8wX5YQO-lnm(xRyLBSCLTFcafAqxBcjymNt;qheNxo zG{2X*2uJ6=qP}mR-uT~J#O8zHoBn&NfB66Cci|cL?%hkd+W73~P@#GB=1fv-46|$So#a``8;B4_%*bXpY8;$44t8tWGdwPw}f_ow)e=_C-AeY)0~;w zfW+hK6BW_(dA{`4@dN)_zPWOTEp*^?yFaNvY@PewZ_=DfKW;(zU9A{>%i z&PQK_RbrT4clzE0%nZ(rB|974|9gEvE`UL4S#^Bt5^+E~ zo~>o;Q0>yqneui8!JiN#lCY?B(>{=L-|`J^kBc)aY2M6S2BE)q={PaVnOu+#ooLI z>))87n0vkP+4-W{HEht&=wlAGcaKpT7nj-eu?4Dg^4wTc>KJ>x1%(eQaH>(S_eRMI zdt=QB#`~flKUNZ!UC||ma#sTT3^sEY*6Rf&5}qBsh@W>181Of*xa~$Zo+=v(WtTU1 zWa*E=nI;(}o-q7yqqq6xI#pC)aK)`dDx8HAm<(ytT6&Qn=SZ8*P>JgByD>!48hlTy zZ59jY|FJSXxE{BSCT!w>pP8vyO7F_EXqdJbZ%)Yxoc-y(KYj6H>Ov1ARU22?fNLL* zz+0txw7=AQw_?**8&1}*foEF57;?WM`{q*};$_SZ@u4tWi=p4xrV@c(pU+$K+GH8eqYf45w0r@Uzr5;DNJ*ojMw<)qZ|bV=JzS&g+-wpq~h zyN2ea>y!dplitjsW5qR#_+?XePQ16!F4dd;Jt5XpOGQb?72TU3ls5OgyOwivpmN-y z)4QwK*?EtwU4jeknf-;yrC5SlsqZEOR^s&o9QSIosCwOpaxfXb-ziO*{c;czh;awzuEi#=;mPIy2~<% zThG_H9EMogzGpw%bk>OvcT11ksVFFPet&v6is1_@v=AyzpaUh{PygpnG?0;YSe?mljQg!19`c=M#)O}I2g$VYqSO|3`{)qx$A>uVfGMu*;# zymgU7Q5zJs=(2v;U*Iab#H!xDJ;_K57>rK- zTI5)TpQJhlr1=lmMc9z#9Lik=0!N;>ufS%!dR|{Ybv9B^SOxhl?qH-e*N2?E3vYJx zV$x&(^l}Lf99$}J+ALiY*_8ysdb7sY72PsDJSJscR0)`^!S!2@cejD+x%o&{H$Amig~8tW4nSdMy;&dz*ok!$*{vzY4Kf>(B==9kU<)i>Q+;utiX*H6SP zk0lo^V_5wu*-Bn&SLHb)u{wAZO>S?0Tcn^wNPF{ox3zPM?|1TSH=`qUX(&v`N-fi^ zYCM8laqGuN%E|&`MAjpt37fXWptku7eI8ic>l||R%jeJKa}I41OSc`#1w*pw+1ZcZ zH;o)Td*4R2|5hZUhc*<6<@Sw5IO(qpOttU}Yr2eAv0F6?v)?w=k#HQlg{f{|W-TY+ zU73P{a;#agQ;h@lZIT+EzS&z*nWfQRHq)e*erY9mi6th@Wi4k8on66=Ho}rVAYHt$ z>Qg(@5wg12hsL!EB3ktjTXpe&xlHm~Pj=?CD4c(^r*ynEjj_eX^HvUqNh=bpEwA6> zhBiaTov37F_3I+m+q&4P^d@Q9T~W)K6ha{)$pug}Qjkz^ zOGK5uc1b^v??+g17MYH!td&c*e+}zVi6afM_V6cm>Es(dF^-h!5^8yFtO!}Vwvvyk zo&RAk@lW-pw@O^N+ohBuM-<`w`SU6J&(%awCb5y}3Hz@ELqEkE?^&U;9*-q7Ik|uaP>L$hWPnS&i|> zQmZdR139()GgL%usfRJQ7-`!5#(U-YFLw(3v57hly|SgA4wJ38WsB6V^EM3M~G-qQ20H znVxr9P(kIWN?Q7UIQiPu_s7}Qj{Al7t0+a>;9Q+)PEn+^XPz|~zLYC0p{+@I&o8*m zFL5hu0)N@uhLM@t^$#QtOCw)DR@~~f;9EcEX+(y0zgJ4qKQ*D#64J!SvtQR*XRoM8 zeXW`9nmb$kgUeZ}A~BTybu|lH8VG3lw5-^sd96s#f2!+vK7hBPnXTN2qeaEFYcyl3 zuM4zt%G1O84xC(iLMVFDk9+mU>M~orHfuqh>T8uu^)*7;q6Etqs=E=^Niw#^C!0v7 zN-sTFOW%1a^&|(^2*5?hW zKD}d=1m_tIzjb-X%Fz>NSyeOI)Kgs&YO6jW(+vBWy=%*K2}P6QvFz2;aW6PAU0Ymw zhYp+RaEa2~6yx6zynXGaZZ`$=%wX~wKE=r-j7tzXs>y#ZCt5*Q|7dFDCInzDAp2dV>uPHU^nj>3FC=m*n5{gJv}(>?EsQjtPi zzT~JDE=FFz&2O`HubiK3vEEka$tO8t4)4(61CuSU9pgvxa?|HT>pJHdLSYUp<)(i0 zjrg5)?@^bv=G)IY(L18N&1&aQ2pQ0nZIy1l1eEi9_>-R3rngci_)%JncvktmUfzw~ z_fD}j0vo1ROC=B^ol(2B%I5XHbNjoZu4uf?R6eEj&(EzwI}y*k$} zHnVk`3o^xXeIhfW(lXy_xq6#McsAdi@_POp;pW8(LMi>(cUWV1o%AB_Dy!IS162XB zs*VypOfFe)t(^K?%qBw>o!+1{v3#pZw(N*q>ZQmYA7os8&4-KqE=8TX1!edf$B$?A zH(|{9t_G}Rg}3I`UoayIB%4z0Dh7)(KX0&64`H{c&Js}VtwllkdRM~4Bnl%E9U5z9 z3})Mtn!fq?3Pg5WG+treG~o+mFH&pS(e ztJ(CH82zv;)swS#s=W3tXJ=jKaQcAec&yIaFZH2lkz)=Fb?m%n^&~CfbEr$l%Q$Y6 zcmGQVFTuC858fzBW+}P41Ft$5Q^SoLDV%)Xa?>bF7~_-rFS6YV_3o&03+%&L^;q$X z8ZDyLrAT%E7XmN)QT`QVn)u)bo5j*Xq``x=pRB?ld)H}aYN^Wes;<9WkcCn$XvZj( z^Cl*kbX1{%*>Yarf*Rushz0--dP8Ja)`Ej5}sWv0#c7 z?#Lx8)xEgHRzmkTnVLoT2N8(4Yn}o)sPNGR(gq$y@-1x@L`%(GZ^|$=dx1!l#}oNaP~dH zDq%Tgh``8!Sl~&6Yi%qZws;o`Z)OVw*2Mnt) z3R5O$(o$Aj$G#_y7Y{_dmG39xZAc-%>IHAljT)9b#_VLGNE7=jh}Y<~>uVZLq~B*W z)TlAoM2rGsD>KAV!vd2l)(C*PO##{lxxTp5G`KYDTG$GAX{crB!V}S;w1z$Op{a%B!u8u}-lxpY;Zbumkxr z9Vk=sV3xkHpfZ0@cz6a^-!V90Fs4jJFY`RwiGopINz+`3T-T^_gd5(T9CrRF! zrS9C%Hf^%h(y4i;5l&6^nW!-{n%k`F0vf{hx%nT1qJxnky;`P>%^gtXf7&qtr4 zgsawOnBEV3xFBqZH!vcz2Y-3#8cW=m=W`fRA2q{>bMUo3ZGF<#MbFn8aN%sXp?F`@ z-+1>HRUxa1Q$9MeA-0r34>LX)P9WoKIvd7dWyhq7SHS$(y8!id9%|7Xg-;gfAyrJDrLYxG>MWVHY|9 zpqT$Ci<(%bns_vUx?=TN8AW-nYBKe~|I*6Iyohlgfrx6zetkLX{sGOvOXjIbPwlLv z5=kv4Lw5;$kHs`F&lfyXe4wr`6RXov(6j_60LnF0A2sFv9*HetpQ9TdN4pg}vZQiv zd}tGvm1at8f95N;@}4=G{pWr<-`TlBFc)%pvR}1zHSSuNM+gshO9X1-j;)+{hijMg z>c<%rj>LDqG`?MX&HSS4km91BSx_!tLBgN~;uhiWdQ$`_D@#ynTiCnSnx>Je z#sqh16*zmfoO6!RZj1fbK~i8|@9}2ZNH0(C!Rp?Tik#BZA&n@L zjH(}>qEhQ(D1TRTSXD;@Sp9?#k13L5J7s|CAr)4~!_)!cQEVUF5lsXm`(pN+(jUhD_B9e*j|Qp?#?f%~OW z&jn(c*NWFO6x<#z*E*pyu^}P7>v!+Y1w0nFaGo=bGS4yEq|u9b7LM2x9+zwK@AUrC z@)mS|h1tcHsh_mFayJX3$~C($To!#iI~LJ-^FmyLlqHMv&*PBUd%^4n*(!S<^*t^@ z50S>#6Q!+H!khv7vVXD6Frh&jq|GK|=Rv0EUyZ#>qc22QNYjZI?R7B%j{Wo{|1c4J zdcP9q4C+{~vI~)-->-+V>rpCyIYNw1_HKTD%BlEE`J~yDwkIAg_A#C`)8bHC;RCSH zwPnPqswp0H9*aJ&(Qi+sY;yQAp{)-k)$7CpBTbDbPz9RkB<9$gq6FRX?42?U~ zOU*2v$Krs@k3kht%}~A3GmQO*BDx8~#*aE#rxzvDN>Qo!!P#`+A?&8=5aYKKjvwnqfhle=S?#%aU_k6dUyN_Q0>`SlG?(I^qVZ$-CB{3uZg%g=lpYpAXP5rbgs~2}q zL@!FbULKL1nOCRaTM(VRYu&lAO6s1r0~gpE;POu&R8E=GKK~2KOWp_Xss>DQ=3Sn+|Lyiw?i4{9x`u-El9e`0b#UByARvdM>I$t60A%Uz6anQk~PX7qd1M8JjDa zBhVorEI#cYR2B4)a;xLuaZV_3yHs4epk$%P;@*Eai&+$gh|N@=yr-$>y6!$NG&tQ?aW| z$#l-*nKa7JcF816cA;>sXt=0OXvm!=&*9IkK>Id%&H?haH735oG}|poWplTSJ?IWQ zWi*KLnF$F`8=pQ?IH5+2J?1TyPb3g1w~Uc=rV7`T1A9+vJfHh>TOpS|xVgq6#_Tpr zV_K?hk$W#!H<5CM7>+8!+~LcRtiRce{uUn+o{5HF(&(tfLd;K``r2Sw%z&G z0b`j2vi72V*-d5o__|$1*}Xze91ZW`@mv z*<0k#V=%P0SI-u#W`Mj1qK7p5yU)@eK|>~saX=@u+J9lnQqj9DUCTW*I1MZIWdS6{ z%GOT-gFy}c1@ou17P>5o7O<68U8!i&q}}@ASWQ=GZfaJcQ`!@nD=W)xO-Z(SUqY2} zu5@45dHnZe`3&1wf)6+?6adL6!lYUehuwKmK84-$0Z?PSb7gbQ(jcRfLY!_NW7FHv zV=bxF5!6@%sF|&owJ22qGco=L=q;%p-!&Ww0M6efX@sp+`vZTgifMrJbA=1{QQ5J{ zREvcD8?>r_P`xKC-`KQivLCBCFB~9FRxHnx&8j-{`h=C!)cqBZ)Z0xWvKIk!O{+JQ z9x(j&>&GWYHf>n}`tYut%j^Mi#qeav;s*)VP>$WUFFEwYnkQ8=q(0DV5*Zp6>t4&s z*S~Gg4_Tg@os-O5_}b#ZP3@ZLN=UjqNQ&S2YG}uzy8i=?WBipx5_0A6;lqU>#dKEs zO+HJ(4A=4H<2Z{&-X^IkvcnAx8B1b#ZO%XapvqEUznk>F9$Dz8aJF$t^O%^v$+BAg zc@6!toOzBrs3%r7HqTQQ`+eKK*gG{1gHo9~iPyoJxzlSp=d7FPKqqVg)0_E-^Zg|Q zMeu-FozabULD{9$B%leEgh`WYU{p9}CwvFHS3K>j%D(0*))_VJ*QzbrtEc_XN3?wu6w>qfsF{*RelQIV*ew#P20*m~w zgA{(^@~gWk>tFUdKUU5dd^~jPlTcC6g1asH)&My#QXsgW>@kxpV5->><(pQ-Cn%Iq zVPEN&$B;UG?dxN!A0O|T&~{=a=g&oM5&ibdZ(8}^oQTr9t@1*01K`8W#^j6d)xfFW zjfx8HCkZ_5mq;(U-sn2lihCh8;a?k37FmQ*AK>?{oYt4Xjrs(yNV*6POjnZ;g=rLR z+{%%_!imV{{pji)s3D~m(Yg-*3VNh~ZQ7nXa|WgGktM~mQkiQv-ZOGvX#XD6zLvd9 zII`~%5V6PCMMJ1jerCdctjP~;@&ien+3vdHBt0KwT^>&$T?RCGx+?jRdIE5r5oIRH zxF0opZ?DtaGnZZYWO5?3WiK*OKfvNMx#gsBsW>uyFJ^mRBcF`2>*ohH$g(zBPBk&g zXXNWqWX*9QmmK0_BU40^*2JsGwSS)OTqG?p`>2&*{k5zsI}e04cm4nqZ>e%yn7BmP zP=M#HKCQGBE;>0D?gg$uBID<6v#^=L8ppA?@ib1r-B95QK|=pzKomzS#qk0o;0cEW ziq#omJ{bg3DxBw7T@(7}WZfIP(l^VSOwqr;C&{7~rhCb-2$QMaQoCG!jk)&~MZBur zm2|@Ld%xcv8JqUi?{Acii1%Y_gv^BThMA>aJQ+gV&R+v-K~NU(>w})WSpHZ5t>l?q z8!{#IV$gVMy_E0`G(Szbcu~FrD4MAwnaX;8sf#O`T3r*>hxKEWYC}*aDndOI3M;pG zf@iApTPEJz&Y~goWbGs&wR-ULRB>r#bt|>p=NeVMM}@oKguXK(u<_Xxm^@a{hm0jZ zYA*t~&sYG4#R&;T+KC4a^NEa-*;Z6gLzQ#}0s7mNqJMS{dxPQIOq!bQ_H_NV zF8k}SLZw?upnl^+x}tUrqzn91nxH$yH;J@_;R3rD#$~y}Nec{WnUq=jG`Z9@*A|sD z8T5Lxd!#VpLZSCOXyW!zMu_GXUIK5YeHl>46qj>3Xu2~PGA1N(4UEj7!b*M^V9#rg zk;SByH^_9SYa-L`j}dVcq(=i3zQk|e z4A?5~6EJ^oMIg*$7tiM4Jf~7snQM8^E^@ARZtE{Ct_8(A+fI~EvgV>j4SGZP*~yuW zi>V*MRXFzxGA4o@$+tXz895uf)H?7ZhBZNrliWSgm!BTC7PQRMR#JVY>k>)hwnGLv z#r)JV(FN8(>Z$;u?jX-CAk{L7qIeXoJZvUJ^_}ORGdG! zR&3iy>$CU&ehOw)Tp{FHdLxCk|CP2Xh|AKC{SG?&5ipbI27_ldY1ksdH(ui3iZIt> z(n{T+1rvrI?L2B!zQNkr3w3l?9%&s_5MD>iigR(uDmt%j-i(^z`76!NCZA!6Q(%-PKyqi14r{ z`xL)!Oxk&3WQTf?*aPqCm8txlJ9Zp|-w+S$go^Itiz9)k5`G=aETxWi=3MNo3P7_( zZ?4vYaA?UYAK9iBx>}wfYi|XzjGXF#U%Vfa)@{racn_%E{+uhbADHgv&z}u|egSU~ zKbXo_*QPW)JD!T(S!Z;Zcg%!;FQ0+^bO=CJ5_oC9y)_V>s?Y)B{)TT*_Ev|E+EAgrDj-`#6tDhXd2_|kG+h%6X5-D6+lXvE z-fE5}TEz_*PM#IkN*Jyd9fgITpybtM02z*Bz%^93dINB3o2|wIX8Ntla`{RCn;1e- z=LH8E`TS;lfbGH~wr-vDn- zSR~F(p1m~+YhxZU#*htL_%2jG(RWgdM=Xz6%90>%6*`HrpPW=n08?4z+ zUD+4E0URA;x9O<199;N~e4-)rt@P%AUFgc>vEn9VyNSjk74J!#Nme=4U0rPR?iJ5h zBRI=D--!_2*fo+SZ7t`!TDv-T(N{bB%9Se#lV-*eDoWm>ce{OCLx&As`?@WywekF# z$?p8*UVW|FpBQ`?Y|ee!A?3Un&a-1?;t*fqWmUds#cH1b*gucgc3!bzD?b4rOBvUN zl+_&|iPr*(wgE+1ZtZ8fidGbCTP@!Kq!UWJKZu^TIbG*?Md%2Ir0qX-1xX|3O~%T> z&d=+-`e6~R==1ZO?}3E_o%p9%M{xMhK_p_te@Hy!mtMSz#)Gu>g*V6r%r;N*bzH=S zQm^jj(#(cZ;}it~crb2{%;YgWZ?8tG2@_@7#~br{ZE-x02t(|pKrAN*l#|;dPgIR~ zc9+_#AR@Sd=crj@P&>t1l(RUmId&!PCKewY%5NR)*sJMjI|~q?iCH!z2W!qh=lDT& zM><(TB0)u?gxw+z-g)wgA>Zf@S{e3+Un;h>K{>T8azG3 zV`Z8}*M>+qs~n?52lkXt6VRZ7)tTNBk2?)!f-P*GH>|#=G93s@kb{bp)YmDfP6f~paeKHoPwmtXd%b@ z-Z%qfzfrsnF99)UwnT@_z~b|-<#g*aK`$OT>6110aj^1aT1+T%A?iAX0(zIB^FcWPlKbS=;**r?c9gNeWIRD=DZ9yM=sjA(*pM5AEm%@%Dn z$>5M`e*b*rRD5xu7_8Xf0SpM4%^{E^I9rX>Szq-n2}pr7%ADKL^5T zYr0wdIfPE*Y)Me{dQR!unxF)?o~&UDGUIg2(&tvT& zUe`=G;M5mFX_dEvpQ=Qc0C6Hf@*ubm-a^O*Y58VZ!7YUJSy zj~|p3;IiHbi(xI)-JAdMB^Lpo_bk29+v)FH3&wO{Qy6!FTC~4SFH)dlgEuykr_S3y z5mij{R)!M7}}_2~u>FXyr+kyo}jG9YL6YIeGo}uOG8XVXO*z zL7O(G`M`b+9i>}k*afeA+Yam13OwFHT`+>JJWt;o@wGo=;H{5JLK)-UcD3G3;xQ2VXP>lm;)|WTXeP ze8l>ofpwtyWCk$NJC7D@+Pdd?-OJ8vw|FdC9J_#>vqm(926p={!ya(QYs*yXgtWa- zQ?%qQz{Taf`jS}f~CRp6IZ8=Mn$=$PEJW!K5EZ;;+_49enMF0@lt1PnSg z``un6Jcsj|LDlS%;7>@G#y$>t+P;uT8}eoKwKKZpqULGKpwR!6rM|nSuUG@6neuyf z#F3LreV*>C?vz`GBujPI+KVV)rfo_J1kNppeke8>$b}+oOD32oQc3d5!-hsA>yUv9 zMy*<@H;2|TzxBm#M8rA#Y1B%?nX2w9+gUAdEb|>yS-@$Ge|^<~Kul4Ty;&{MEAC%8 z-1hGMmT7B%DH`yE+CUoL1E}9Y{T~MB2X8mYx~E_Lbl=wS$C~78)pv{Q@#Yu3K*TQ^ ztM2~mwQUJOIy~VD3t2r9qoFzI$@r9HN)>eF{ODYy*rO}N2Wz^e$Wc2uzrQ$wFIET3 z(5<9E`&T)cJq$gYEox&?Ts`*?HS`RunWzxX-rz!IdNddD!sgaD`(F)f-9zbUh$(;# z9@@(%SC}(aj7ZYHf~u30!#yVg{uN)_;xaj$m-Hk(79)GQa4Z>5Ee@uY3RvAfx~tBp zdAL5Z4VQ(ZrKL@SPRHO2ZPt@eO|A0ib=nc+rgwKqmNw zK1a86E%|5Xn)bAICUm}CDEwj&qyYDkbsm?6W_|qhDF+bz_&SJH$)k3)p*73)VEzO_ zE-%E%l*8WZs7r55;StVJ^QA$8F78y!vWh=n0EB^cKQ7DeJrl*~lF4PEN)jQJe375FLGrd%Wh? z$F1V@AixhHcFBtY1bV<3dhOz|EUtWz`ufV!%*+%W3Bi<~!Zbm%ds}w@d_rA&IbG(~ zi}d$cip`2#^hj`^(eWfJZ8~X4fz@k+2wOI z!Runbj|C*reXaI7rni8bJh?CyO=wsfoKMneT)t`Ycz~X@Shd%4Gp_jeiPBarH|MHF z?oLE@(yuG_B8Y4eTEttU31%dSj#bO~s^8M@95a$AC}rC+4IJob0B6-fL|f+7yd+6e5m8GTr+ z)va07`{RY$Cnr~V!QpHb<`wl!b*U~aW!Nx%4_#&g%1+;P;IhM}72*=$f9$8$vf+&17I zJhJ?JuZAamRP>+OJ*Pjg^_wW*iv*gaitYM;?i%TZu5-^x6h(%8NJNM*6P7dRA{2313^mC^+A-VfK|J18;4u&bi2MRVkMvJiloF zl9&`tvAVVX#W{KQI^cP!m*Xh$%G?eu;eoJIiwh=CeAt&ze{vp0ZKAI`M0YIm=V&7) zBK7;DVzXbUHeS^+KC-Br(`P0Vo_Oij=VWfX-F4BPcHIYMY@M8)G!jQ?fTXj*CYo7G zN|Rx&qikdGI++o^j;;FZ4?HWbC&o|@}h?cRdt4L+0wo1nsb?#v>zp;-h8-pa*74Xi8u zrYrZ^kpQlt+WNv(%hbS%HgpMFI@5V(fFoqo8d`3~wtn*BE?C?d>5D@{)M9Nv!jZVb z(qbiV>|;o=U&9-HDQx+9Q9>6ZaUENO`b8^xlA~;2no&?b9nyh!Y%ILG)3Z@+eHFL5 zaQc|?KR;$>)Vy~VDi;$~KnQwdEa5d6<8zWU^vnoVJ{7BOr(qizNlQ)|(0{)6Evvpv zm$#af*%XqF(d*(%2!m}r4dxtXT{?Ji8j@&>r?U6RJ59f_-9-hi>L`55hhO>l)uowN zVvkTmhawcfRIzUOj2&TJ+@8 z;B2~DUyG=t{lsYU&yU&l9X+1*wTf&fZ^U^uyBPkO?(fa+DqVXCOO^dM1wfABhUb<6kyLz?q^A~LUxa82~F z!Q;i0TzS3rC^shin?0TB!#;OFLd3Ly+O z|G8iwIID)0j@&Lh|07x@%K|%6#H|c;t#vb%pzNHQR?j0w-+>5L{bAfAr!%SrEA2l9 zV@-^ZB7&S2516<(Dfadlh1j<33bLULV7aRB0xee`1cjqE(^AH?Y9uJrn_KsnNu{gY zd)FpxM*978YjJ%WP%!fD>UTqfW5;8e+p6P@BKcP&v+AY2S03GW!)Jn+}RUe!Z{7hRdUaa{wW*)hr@^ zP!6BN(zfZyAnwXX?OGCnudn82BfPPwDOU#Q;Ml#1MX z=_yl2g=erFh$W% z#57b4IkkbauGGqf^>m(aS`-AtP@DmNtrK+#WLF!QZ&P5qe~YVMt%0a>d;;Na=kx7o zj?da^5Cdqs3inTl1C7`ff{35Lm1aMlg6 z5@jokOjN6gbLZCS)g@Z1u4YJyq8*h zQwLO)6>wMzFDA>SAO*J_cnd>QGj{7RXJEdtV)EL-U<;Z7itgh?S$h@5LO2`+)P0AB zkiqI8D}Z;7^Z?i8vLLoQHAJ0fux=(8(!tfPO0izT!JR1zv5eX#_Wn=e9R~jhk4;Qd zwg(Ktc$>YU<@@V~)Y)3lJqsY4Q2+4ZLj^LA=XZKIN5Y81ld@Hw_e~=uovIvCacl$Q zuvrs^Q_6V5J1AROk!SWQ<-xFio3U^!J?50-hI~_1yx-4(xlQPXp*MeSAJ6|GAZ>p` zHq;&*Zj4D)$nyIIFbi#Y!F@GhVMrj5RV+qCy(k4KhN($NyTLwCiIKcX$ZTt>#_iE> zN9g`Wwbmh8={%uCgzJdQQ!4)2C?z_P0?qH?Rdh)9<+>+2BDmpb0oKY13y8KBcKDm(?# zrnxHz=8|0bPj##PW4#mh46 z_+*4W^ad4D(^!$5imV`}$?an2Fu}Le#EM)!wk5y+$~su~@S&DDwe|kC^FrT`y7Pul zu6bqD;^l4MUyn|dI=1nx(m#cIYMkFl-6rk}N*&>t=nWiDELWWM#M)#iZ7f{NGxYc2 z@KQ)FP)~KEEF*VRI;)op-dkPJDuWzRiAfqv$h6q-9pvhL);jH;5y=`W(E2Py`K=f} zmM*M55pceL<`HONZMz?v8ad_8L-`p=dFI;&a;2c6W8@=59L@90Ux>zEu#;13MWl>> z(LU?vuc9)?;#oybAem5qtbaVUXXptPCC+0xh z*ccga^-S({o>-<7vu(-N-Gbe=7i zDueDw4fR}`ob4uDqcxCWG=|}0+ajUm!JGX$?o6sS%p$41gAb0s==P=UrIjVB@ZqNN za*Z=gu+zlDNG-a5HD(VpdceaBG_B^KI0yXA7L>OfY<5*V+qvl;s0yiEHpRGWGZlSC z$r~Rk3$roU(5YF95>n*YlGM}a4ypL>(9+XaCVMM5gFskkgKIO`Lt+DfnTuI_*7t@o^d@6!(m+z&Aih+Pd;q z&805{SM1OYNS={lAN7A~63e9`r!zEXo-9%Lw%_d%W+|*?XAkqIzkeOClk^ zD-bp(kdEJx!4d#oVyVuih96+od?3YJv3is1ZJ^ti?nmNSmeqh~{hW%*P z&EVfQza+Maj;&Q~X48I#Z$f58MzzW?s<8)11VBSZcK7DFv z$9g1Bii3V;wzkGWPxhAN0zA9GqnaD+02?pZWBE?+=Y3XoYl%$L7R;eDmeA#!hhTqg z)Hsd^-|h>8G`aQ^)4rvH8xEJ_s< zgbs*_*L4csfYJ3e)jgIsOuEHlv*eXa1O%e<(Jz-{utDa}XxOcTUfg`?txLVjF+tr2 zS;l;b>|?wgB#D5L3@O-b$>8Q*gagLch*u>gI*K%HZ{ zW&+&Uc`1ldNcTklWtn)wQX#DDM(@t)Yy{$r4fx-q5M#IvCZ^3u894G#bssK#t%p{7 z5t1+_79T+`HiC?9^KZEO-(gh!pf`>9ly|#BBoLCMo^{(J8^DHZMy z%a`0yNP6{x`iTF}Ki`e<6RE?3PdEm!s@|0M{2O>Z``2;fTX^)rk)cpdA)NZ*9>f2N zr|!<2WdsGzbab{sA{ib$@EwHoUDKogeM~a?nh~*A}ElZ zJp^Aj87?gNAEU@4}?Hp&x1WeOTMg##;+-Y+%K9GthC5bImUPewXE*{`%GVqf+{& zYUQ}p{PkijL#`!VqHK{N@!>U=wWnQ4S9;R;D{fsNPc=0~O8PI)W z(gye0^y?fTM`g@A_l~HWiL3phnr?m+ll1(-U(qKll5ksQN4i{AWY*K>f92-wKg zrMO8cJSXis{+RzGLtA+xyu6?EO8*@g8q0J=5=Z-`9Oz=XIXvWjt;pLw(j_8BZo0D{~KC z*2b1eFi7^o4`d5+x?5&-F-8DRxN&4KSp?ADCWc-Qm02L}j{`L^b<8jz+ zpxxpH2@D}=XtG%xZcQSlEL#!Lvrg7+2V}NbTbTN?%*W~1l=>azPeJ43=^NVR*mu1% zl9{@9>8!SUivczv9crKGWz;80dT*JpXW}s$Tm(MoU+2_@(%ZiYr<|P8`5PtAaOzk0 z9bH0O&QBI)V(D`sWmGhNm6d8HNd02o+2tAY1lpJ++Wq6ke}3y!7o=^CUwgbEn_Ks` z&?Q4aFH_EMtW7=w|=C5|WJ zo6h6=6VIZVcG?=CNMgw_Ec=5KxcM26S7AHvJHomvHj?_!a%>-4pxjS|{SsTsS!XFr z+Lt{+uLoUc2-yxawPi+1cvm~wGl+rECy8g3wcV<&%5 zKazjPd_CfVHNG2Z?^`%rF)9(N8SVDS)kiYpk{pqp`E`P&$nWM&@o5vMukHHBj~_oz zTN-ru@LksRcP<18nZy{VrW+@;2})E=F&qK6kCaWTd}||F6u&Anig9 zfHtpm+MVt5Hkq#nW*@5If&*7pY}}46#BeY2Y4Y$tI4PDbOWDX-xnqf(07k>`mTH64iYk#0uJMN940>iTSrPPr|qb{ zjHUWCheVDZB}xU>xCLDttBd^IY0S?wL5xm)G~N z3K%`sMjwm^pnSj*Rz9MWFoP8nX=pZ_44IjQD(Qj}eCNnx!*({{4 zpDfr-3&Ll-QFUGOO3`__T za*u#@KV3lkK5^twK{ld*ID2a15#`Pigk<>WeV%?vJ@Gg%{Su21#qDay06JOl_wCDM zItos6YzlS#%1QZedc#QloK)`N+|U2^e)-o9ThmC_j7C5=j>DxEtEjK|{>)w`Qx7;M zVE2TYa?|4L@jOU%u`%XefcPU(|M_I zGTdtQmw;e_sCN!*3SB=MyY+E}i=U1FWA-Wy!ts93kcm2KogF#di`A%Cif(muKJ};4 z*Q8D+kH7!q!&b@%kQSzHd%|e<)BE{CG4G~C^^N*PuJ$_ush@@J>V3vd)E_t8AWn+p z@n8RmVn@f%$H(`!`c=3H&x^K==w7wEe=BwXl~Wlvaw`a0n&{h{153di_ng;w5A~h) zNqZCLnMz!(fTDJSb)$L3+cATd=N7Om0}UbJq(EL|_4?jr`GW_m;8zFRI9tVFPTE1V zPtscqh5EVo7Zz=ml4|WWFB$Jye-2uN0M=iAHrO|F`9ZN;j!L9>?Xum@U;GKRvvK)A zN3nbIDOxbx#8#S&L>Aii{m3IZ(f_*kYR~y!XNZ$tS(X5cnNkuCzXd*U*rnR`lq zm-CmfQ|}&&y=bFedO2T)O!UtMh*OwKwA2x75KDTk4q~(Oqw7y0SWm7q!^=}3^6&z; z&AiHN=AgX%FQ?tq%W&olT?!7QEVIg04c8_BU8?l<#rWuvZWGA>R&j9H_B3!`nu+r; zwbI(B8cI7p&ZM95vb%%CB7s{6Ac{Hs?-H&wp<6o;kuH{2GSSc9(5P->jq_ZfsbPof zcT*&re*$*VApNo#6Br6M=$}WsP0Pcb`Z*TgX4A`kc6n(erw$^CS$n=c1!{M~3{aNU(n6Q+_e8ne!~7k%%{p#=!#mlm!_hs+US0XeR}fAA{*uf{5#&aa9_kDDe|+@ej%yT~t6L?>3SUC6o)(xu=buh#NJkkTOg4 z(NplCD(>vPIF+kPz2I>-ckthhX_}VX;@cW^c*|LE0Ruq)tpHVhb;Z^&11*w|fpbue z10^1rW~ITc=4YA)DNE_T&xYilKHywx&!!T~`eiF;_%^Ng@&I?3-$-6FnJl$#d~myb zsI~J5f)BES{lp&i*6`;Mls4moOqW(jfA>y3nx}@$JNnD-Z9i3qMDet6m_M-zSq2}X znV@f8r#o+L-1D_%kotK`pUx8)%zu{@vfapf#@fV?`dfuW{C1m+o=&RXrrx^PP&ihS z6@PrxYZRi>3T#?vR0R@vX@Q1fMaLhOZFL+NyNcFRk{aOfvoRjZM?ZL2N@4J-a_1Gj zBG(~P9#-n>_wDqK&v|0(F?ysy`wV!$+^*0jYywspK$>-q93}`craUaPKNj;og4CUW zjG2Md-RZr0V)oTt>Oa#hB0%Zf;*Wi-jGfqx%YQ+E7TNcCcTSeN!Nk)X-|FE@tJ3~Z6VIJ{NrS$C-e9ee zaS(d!)j%`MfG!eEk{8bk>U;MUFc4~XHv!^ipO9<{B-3CVkdYG0N?`FPY+28lq>~l_ zeg7czH?kFl)x0an<5GhCisqOIA+=RLd z{YQiq1R4I|At;epgiy|1LA`O_BefY+Yz>9enS~kJ2w}uUSE7~yD8=9sTJp;61ya4l!I}e5(-e^Gk6rD)7tlwnb z?JM&@mh14GNg`i%5+b?Q*h=av%DJsDl|h@H+Y||Qx?U{xGt(Vfkcn#*`2HmIbvSoA zcF~TKzxIOnNZY0FNQRVO|7bl)TQw93WX4iJ3scqF1gnLs|LPLqYvm#`Rl^uV$(ELm z=+8?f88eO;GXy|~p{@hxfbP~FPS)o2vZcMCm9NQedw5>`Nn#LK;JbDat7NPqa<&o} zd!x^;-3nUIQfGlLD#7{X6& zn>v3(xQnZpGnvdQIEl6uLK_{>P`CkShhpqk7rX66_8qR$MNS()y0P1J*9+?4l0iClTNzEaBs79NDOF%#cHbX3_ zpQtZYdY7K96v_D$m4?0S1f2;nS40-w9Up2OaFxSN!EcQqTLtZJIT!qrL7+<^=Ey!e zgc{9gz$#>HNZ8s9cPEK_=^L#C_!yvLz=bV4mz}eTTE72B@nnN? zDZwrU^`GLOYXnI$1_9v`kzZY~&=~z)@-ov;bzC|exhx&gN&M6sW}7J0`GzqeuZN(d zn0vf$q0_;Ys60mf&$z|E!3ar%E2KFx_fYujw{e|xI?z0Qf;9%OK{+}*aY%5JLb$1a zcWDil6sZVF6ekiYx21Q5;N{3ev#2&8we^$umDIoP%V+tCu9N}I}P3=8%nc|WEj6&ke1Ee zj2LK+#EPs$2)VsJA$1QxRm*$$jyIzY1eogmEg3@igQfBm>RZ1*Gmp{qt?0HUpd}*I zU(xOT_FG@9zzd?`+~bVs|G`h$-%U@Az%c;E8HQ*M--wd5--iXjh5TBqVfjYcQ!hb7 z(|Gp$Vt_RD-6R!Cx5&c=AQU^K3SEa6Sptkpmj*s~uuYe406AqK>ft%45FT=NcLJ%Y zv-e2*r8979prq=H;PU%Q57&qYLma?T=FpO5t_*E;3esUQe_HkFy?%+!W{RW%wRXW+08Vze{@q@XZt{He+ z9I%vYuUrsf8ifA!FFNQ?k-L8FRo>|~@mP8z>i1Ka#7-n>(W&h_7Lc-VS2)4f zkKu3c$&hJ$uA8+%(H(HOS()!jw8Au6vO!(1hS6Fa95n(!XoKEWmy?FD(p_C*Nk{xf zgrP;f2GAri*0a^InFD|lX(xTE4pj$PZxi}|#0Q;p(dPdKMJ2W6=Ibfhm=G(ZA2lL% zVct4sW^qZu1eH4nEr&TS4odhGIUd^iznqoh`VAZ$n%FoT*Hvkuv1~yR1MSc6ukXi1($PM1m>kSgE_5DLc|?Q$Ik(MAB^fGUgl^>H$Mtnmtw?%5^7Y+O z91vUJC&P1yy4L9~Fqx(q7aI`rV}aXfa%1ulWl|(z-wHQ?7iiaUIN!R7&+2VfJV2Ai zn%}We@Xowci4NEO_wL;z*RkiuMf<);#+)@eXqLr8u_^2|GbP9UJGm%kaY?)21PIV= znekhqi$?eIO42pjaf#|Uc+-9cUR=a`!+p|#p5`givb`wsKFATOkaMnzR_}d%j(XI2 zdjmXNzF`~0grOtfI!%B|bg8!Ft?CzzKGaKl*{RlRlK`C7+8IJl`U%{(kcTJAm2jQn z01Bq(F{?$2s}c}M#%;tO-Hq>SXje}mEWRmb>P>HOVo7Sdgq8ZO$Yt!1WDb7b4?LP{ z*acl8GFw5iijte=6TN^lhN0o)ZT?Mu!VA^p(f++~Bk&hwT2DNuKVPSL3cO-6C5|8e zLN}P}fAdAW<9_F>WiE1TEet3A6(C)mm%_1iSM~4ExqI$`ltfQ`dZ- zI#w8I^-IbEnugGF+tM`x{_s*zAR$d;YYPq;n`Tl#>zAA__|o~r1Ps(|1myvxRuK*! zq~p%}hIXyvD(Q59mKLk=Lzxd`p;J_(2=?J|_*=$TvHsAKC=eHqS$^#YoKRl|pmSdd zXeHijc7|*YeVEFu-gcPt?Iaba^LA{#R>+WPTRf_B4p?hI@lEt7Do@;!;U?Be&*33r zj>c!!iSp$nuf(tdBk+aBwN)jk-;JV3rwrSF71C|{(!JZMe=D&p$;oJE-EOESJNZRq zcov{%CAO+oln(VNagXM@xn{tmR;^pt7don&r1@g17;TFv>2nmJHtJu=`Ho$HDf_&P z2AHyQ?soYagfpVv@pA`QfKTF%Vg06CTtPYzZIVk{Ys^fhXrwElv$86xJWGHKtS~3+ zKgEkB-G@73hikaYa3d@9Z@&vIU5NKZpUk>32TTB~F60ay( zA$&ZpThGyS-~tASu-p^z6A(#4h%SnV@zm*4eZ2@wg*_zj;@<7i85D^=G~4R<_|?gV zZYPey&h5VoYG)>q!d(fT%@R9Wm$OTmF+>s$v~i@TMXaP+l)IriFYlK0pbXBIAKhnb z&!THvp!wCUg|uvIq$gSu8k}S!w`UBO$;*e;{a-Gp$Jz8yDc&}HjmqiEo7ausDXm7% zCcn?a+#SkWq^#OWyA*OfF^w`L4&k1;nio+M^PW4gnJ(@CoHUm>ukv<9ReT`*(zdGA zk}E+owXTO0lQ2L6)bPZn%rI z%397j8XypxBc}4UHX{{}!{^SjsAVg_zA;0$GhoPzU6Q&TI2{q@;$TJbM~Pr+_MT`L7QPo7>G5kST@^np0Oh#25RwUb zoXPhV-_%(0fJ8w6Ov6&QoIBEM7e$KMgnSvaG_A@?l_xfpT<$bvQ{kdmrrXc2f=JAR z;W(g3NENqa#=`nC-+4~HN2$q6K)p{j;SDvkBMsk~(BZ@JirF?TYLGA6sKk9l6(Y&aZqqn1XEE5Jf0}Gro>$%SsGHn=XYs;dWK$4X})%= zyuXDH^=6*Y=J`tfNlx>x^5x&AyP>^KK(ui*kLU|%aO$OeEDyrb0w{0KDM1CJy_pe- zb&zgz!K&s$J1tw^XNWLNdWd)tUvvEn?X(WuK0W*%G>RY)#4DBx6>;L;FaMN4{SM6c zp3N_Kh`d<|fN3{W*|(wQH-?aXGu#~gD=I2nhp$||TuHpMDn3!4!dbkCbj(cYwOTKI zMFFcIZD~U-YTG#d?-StocdQ+)Y6{JwAk|ZHS)f1WBQh`nF#p`A37{= z6f9uBU(%cIoOOu5+}j>+gNYmsrFQ5&-T>#H;q4M`Qn;I{tU+`2T3? zNRQMJL2<742!k6)%=7Gx^)Oe8+T&#mNYoq? z9VYpZ$Kp(8tB9js%{hEU3JDOJtisJekb4w?6q;ka88NydXVry38vojtXYWQx{>-_)QDd?;GbELMV|i$wOF9!BWN z9wBq5Fnq|q_7QsMgwcb>yu`ms=gW&6PLR+Kp6dP*fR4nPs~!4u@QZr&hEKi8eA-w+ zDg9R+9P(~ONQr%YBvxPq^6tr%s_Fc8z|_F2?Q*7)Z$eXFTwZ`+#d>5ixmTR6Gp2nT+~YQ&TrDu2uJAzu$G$j zi*ts@NcRGtu&Y)J1{l1?Mfo)`GQ0ovtrljp`~>!*fJySr=CBOx}}q)OO`HGcH|f=l|cqMPmxwx&eexlw;-Fw zv*)?rXD_sC&({Djz4{+3h>U^Ch+%lv)dj8NN+cIXrb^3 zuO_t(@;QqsWa#F=xP?6?97!FeondT*nt$a#-htwo278Td(YswjVm&7mu?tSTgjYe$ zvi?`QzXic*Qc+%B{vV)%a^L#JDA4)6riQK(f(}W-(3;dly2MhH3{2-v-~#H`U$oSf zt{P5CtiREun|)=)Sg1eXn!8o16G?4N5iR)v-$W2idJu(zKc6Qs89LkyTWQ^>Gr7QA zel%e&8k+K!Y3r%0<~=7`7x>(2V07^UHCXPfNqw>Rbe6_C!xTd{7qw6JP%|cZ1=j+S#&rsg2`u{Gzf8lNId1SJGls9?aA-8!~FgyosNpi z^?})S!e)e5N`RNCVtwjc#;u$;$m;=)E;8%_Ia^xuWYhLL(fOgMz|eJBU!h73L7Zg> zK~503D4?nK2@On=Wi4=MpuW&z4MZ9-D~dd(NbzE#R|p){Dv6080MaF9wATf9i>=fE z-fLevijP7l$xldo0D6^CMa08Gj5rKdqG2*W{5Nn+V23ZBnH`>!rwMDQCTf=!;VeTS z=?^=sBT7NrHjxnwHQ?caaqsv$?8&Y3$^EcwT^u+S!pJjOO4;DOyjsY(owbo77kcx) z4s}<{O{~cVaRh;gAntsqC&a}v93@@9-yskOESCj^ibn!qL%eYJZ9vr9J#_;WN(_e6Zn6? zt^UMV#Dg%$cIoRVE(CMxrQ;Z4n6pSHU{m^NZR&NElyl{j7z5y44Z@D<6P;F-OcV$| zP}Wux3plh@M)I%U%pedyk=V9jTwe4Pf+qL2gp18%nvTn<`l_|N|R`q=3iHKQRsKxRceDiQWuy{^z`frov z1uuq^7^pD$^IIq!wjxSHoLkeyvLQFFTS;BA3cbP$98M4`Qbyxqgd%VWIo5$nAq}rD z3!<+!M?a5}XJ;Z(XOmIPt)@Znfr|+ee@6!_I`7@TU)wv^ZpZL9=i0N7LbGH(BJ{PI z`(CGEOZTYnht#cf^}dPGN49kdef9RaM8#W&SabHqZVWYCcS=eXK-y_Ygu5i|zuL>d zmNZ|p#K#cx62Gu2W%I@XY zCLK-&uaw?%`JnzQd!;?)yr1Ulp=uYu5DS3v(6i?sb!V!;3tYw~T2XUpIzW}=S2cfFW0gP(T;9M&J#9d+-w2J8a)YOe`mLZcg zHUM%$ufIo|?Sa9=;|Gd{tO250WV&NG9C^r1bc5H69o4_nn{_d%(!l=W0RyZ)~N{M;B99 z@xQ>we>Y?F9$8=NhA|<8Bsin_pg}U3zITOr2=Io;zqrYUM?B^J&C&C|xIoamI~^f|+j|s*K;@G2S;b$^z zdZ6(S#yfVR39Ns-3y7JZNxhE!tBm=CRgU~MPX5}eN`o=D|KG;(Db&_K@hW#zn|0h(7%xP z0e5`E97=xk>F^4QITQAcbNHfJ&^VPLv|JT!GFjE%UhoK~RP%zR-lmaY+mjKSMVlUh zZ)(wgiIJtIO$4i<|E!lSQNX}w;N|*_92}WzFOX{_=i@qaA#i-OZyf6(lptATB#4h5 zgmT}dzM(fsCK?10Fc0VNC-jVdoTPZm!>i$l8xVIi{C~i1pz?-9hviG+Hb`u!!mMx# z4fFDdxr&*V(m9aI&wID z86mGfjg!{lbN!~_RrI@E<8kr`ksAg+N=I+PC9>4tn@P`d8ZF8kNxKM?wK)59n-1{x z<9Jc^nCoJ7$o@V-k`_RU1h9%f!7D8op>$G2-&-Bq%3`Mf+(43-#Lsi99zm3_PW)6@ z5fW?A-B#~%`MTi{S!&)VzE90h9ClwM%~nQS8LiYLiL!9`)oyd_tAz0PI2z-S-2On8 ztxT{}Y->_hd+{)SFIXWF-)mtwCg~nihDG{%nfTzc8KQAq#Np~a$Cl*$>^igear%Ys z4vn0#%kk02knO)-@G#f*owZbM^LwIiRLKr~&l**h&J~Q{;o9)U9^->l&#;r!4*tXy z$g;cf)q{z;3Ef=V1S^p;B*c?kgXH@7ua=%I?V}%CBD&i*-s5;}xJbDu%XM`95L6?3 zE~5No($@6$e*Hdmc`lJw!1rzwso6lmJ=AfomwkMweB-Xo?t&k@C#4mgh@+n3=IIi0&(6@tc=@N#*JHb+R7tl$3x;QL z;6IVb{Py9At2RZ>6=eDc>C&ttKEdlUqO%JdLf4cLr|!9F#KJ|F$zJnjA@}}%G{^e` zq)5k!M$|sI;WwqB4zifSrc&`9#SJ}EOKeG%k+|1Hia}P|#ny?)K0NjmikUFeH%P%% z2h(7ys*Ef02~lCwT-LRKy>8DLDSACQ`VrL&Y-ikvG!LZ90JIL1j^?G4G@)p($WCnZ zU^I^-LRl@p^KIU|k8A^Vnb9DUW%Ve}1-K~h-i?y_Ve#y0o8D$DJ_4K8O^fE0)}o(R zu|iVqV&5aQ^|rE@FNyMn3Xy(2&ydYVT8~0b>A#LH|GUq)ylD5lGgsPTBbmQ!r5h5M z{k+M6MA0O_lN%*j>-I3-|J?@uuk&~R`x*Z`tmprWS4~iLQ6`azfkG7sfA_tcw`_?= zC!dmIqdShscrlu$B#Qe<*NpHKk7{4|%zHUfTxY}L8)BaANJvA#_^8B$WuItyEEGIQ zN`Qk!5F*+D)J7Gkug#$+%Yw=6IO>aTF#Hm@X%JGBkSasw+<0@(iKOvaIFfUN%j=># z0BwVIx}iQepnN!ymLBJ$rV#$N^Yw@YKbU_?2;8b?pXLJuk{yO4^xM4(G@x=TS`;MD z7>2KIKSyHQM+sPX);m&PfwPZdHpjW_vBT*it(I5*4Qhpl zqzpD?&o@ME4+Bhh3zohKZcrhnkW`a1b0!(m6+*^9a1E~uF(>@;t?+gV*m>fhh}xb3 zS9gwRlhB4fsaSX=1Y~Arx(+neNp-~V+#PLgb2iEF)ERRdJiSqNG^0#!c2GCu`;Ke5 z^6CA1HV5$DS8&w}Gc2*VcC(V1J7Cv|F!>jH4uU;T%r^^P35Std)gNP>l1AmTI$Lgk zIA*ba%-tetGi&ap6^gQ?DjUXUPi6E z_X%r2i&p!nTdT7{XX2mbv$r#CUawj+-8qq}vhw)l%ZgW0bY~``Hr>|?f5s)yA93l$ zl*{H9@@64Bssl0~q|DX^?lesOetnhmX!1Hc#W20_En5Yo44%z|s_{z)LRYJBg@4RV z#5b5XFdVwzHl5ig9oLWd&3yJOn{x4aY4%Fs*z29EoCW*62ACv)uD(u&6( z_uh1EXK^Wx>A6x4V*Bnf zp@%hayMQ)0l%O6i{9@`f_6xHnjm)74-OZpLu^JJd^@5|0p@(r2q4V1zDI8!4~meK`1$mk&EyUkO9;OWwavAF@gE1+Fa5>M|YliPRZN|*zXfE z;$U=rGS@C*1jTNaqLgpnD$KRCWyF#jCH0KiU?2Es9ivQ~nNQpwp zUWM7#-8XlC_4ir0R0%g8_17WLB)5?8hD=N00Yr2mghB~{{Z=5xp(ra9L0xzJ+P(`q zCrWm7gzV=X&_+Zh~9PVKE;anDkLbTUQ+1MYnSxH;352fC}k zCEAWqxMT9pWWwH5D80ixw1kSj`xM+}@Ry>vh2QJ5MRLxojMT$!CPvo706uo+V-Jar z7SmyG{S=XN3vI5xny(fmxrvaV19{Z_u5b}P6FM+%$*Jy8dJtJmpqIk#XBVUCICBW8 zZM7z8YoK2^ndRxJv4<|8rj-Z=h&0aQpX;bS<2q%PbL-jU$>FmGKxIeqQk1hGu-!)sYs zPP}t$P&J6GXWL;a=QA7w*&0FBg4dAL{>xc#vV5XfbJi+mbo&*SbGX-h@xo`ih+lou zNYHBrQCuZl6ZxO%#^t3NloR7HCPzxS? zy6f}&rB=_t9j?jnFbI5~(_*ZArh{}72)TCRJfJdC)st~ZQ38o@}YnpiI7H6mX358hB2HKBVYU$7as zlr?ZIu#V+pACU8#S9yt526x!ZQ+U$-9Kb9;)`n>9B5?1zwfC?R_Jda7X<)4h=QdN8 zM=)mg6v1RblE zm+EiREgQKUKXUL7I&&qZ&+}8n*|h;Cy?3Oi49S#x&zoN(bA=&nYGr*4$kyZT!2APO z!1hFa^fae}kI}+q3QUBZ)1-!uU*soq9Pnmy!yaGOMM~89`TAy(c~;4ztzfd5=Ym_q z&o3n<)htX*l~VbN(Jy>tTQY7TGK=Ya%r-NG!cQi_1Vb=qraM1@ant)feL02?qh&o> zt6(bf_7#dJt?cb|eHkjda2xu%eOtU2{!(`n)ICLHHu~z| z*_?Ryi(VmSL??kLE=TFCR-hdsv7&;Ga&r96yobbDdO@~&l;j7t*ml`WQpN*M<}%g~ zd{h7M#GwZDHTzbmE}plN2<)xg(@&S=sQyB5S1mb(!L?IO*7?MAdSFmWDW$sKMn8wG zXfm>}T~IHxbgD`Ga}05WnHp_u-T-8;+kp%U4TalIrBjO$41|Mk`4dt%f`t;_Vd1?=wj{@~uaDN2S`fE8gAi9`dzI-{{%3xtG7GyJ^Y>UB; z;@!>XiTYRl2|Y%(ZP%YvWfCa$qSzGNn@s7e2oYT0G;l#X$hdm73ltGMYj5!BrmNJd zaYhHPK&M^g=NcRuw_K(ZUmjhYBd!Nbr6}PclHX>O^MS>llJmtf{kgT%42}?q5}|to zRGKJ?n&Eo`_MmU*R* zM)VTOj@&=U)bM%)XSvEE%Hxf6ltbg0a2X{CC$O144(L~vMI`j;k}05yw|dIah7EMI zCO8a;w4W|4;_#IaL?TOM!V3770oiJjhcmZW%31#cvY={KjJjM0C3sym5%~%vo_@&E zKi>lVGRE{?@YI{hRDhz^zK7^b+s?_~1-$Mstig^-GbAPotn1cUSs(h9QbSL9gvIK5 zKS#*X9ks%eiZSoQ{v1_{R4-C! zKcX0TPZWU$r)IIg(mmO_C|0c|7lly9)DE10^6yk@^nlP(B@^Y8ku z!jB_D%|A!3`Ren0D+*ts6y2GS5zTa*tk2se2bdl@rz@2wczeZ#3Lm)m9$eaM5;F)# zAw>D{3MZ4)jqOkgs1MAiTpZDq{ndceoe%JxKg?}0`p88k1}yN6>B~RoL|)48JG6L- zGwTjcG4LbW`2cb!afA1)lY$mh^ZpewW;vJ%AZ8p{0rgpZjtU{-{u0p73L7Uin1~%9 zkdT@|hxB@I0yeGzuC#c=lGWQtf3!@X0W!B? z{c?oPnqH6dklLi8uM)m@w@PETSlXB*J6&8%wsE>;VuEwv1>`~JHjtry+BF)S9?Ape ze&W-w$B_wHz5`VlqM=I)Q3I{%}xXp!=dk((7xJTIw?9QjcblK6R7gl*A_W$E=4&XOy}v`)CJU9N9_c!~ZN9Aw;?_b-1${(7R;3&)Vw;ei_= zo|`N$%%YZ!lrVpDIm!L*lVi^r1|UPs2KWEilGDYT`1P~mnvg!|nQ(uFJL!pHh}m=y z=F*TYMh;rea|gFqKh|R`ob)|of7X@kRI#MZx)9C0n3?R{9BZg8f(PZRV+~DP7w3O;3 z!I>l)cei$jmW{0B+P~B^Z2xm*_kX7LUTpd_9}#33q7@n1Mp__X6OaHAhg(8{2z(d7 z0l%}>;KlUh#3K$fx}_6IDiI{ljUW>i)+Ej)WF{VBQx(Ao&>;R1=KOOLI78JB)sMg&{CVw)FDZ~Vs~`LK*KQq z*KCyVia-P6kiuqpPP;wnS-VYrn|jod7+1)5SP8~FeiPH9LdRe68|G)WyA7lfEky+~ z(-XCIGOCse>OY13dSe>BzRrykpg?p(q&Wq#^(VuT09C?~c{A~1hH6Z#dFt_v{5GJv ziKv+nX4kE<^PJ+D^76WUGY!)wPy_Iwix-SWxDpO^)nlF@iRhyPh%KQc`|JIDT4f-= z?;Yx#I3&#u&@rE7p*+yoQW?EFZv%#O5SAq7KRsgnRfG=jP)HhdML3w6re;w~5FAE} zlg!#A0X~eroOk38ahv9vsaagbEc=fKNzVuTnWKc?PMV5cf@g^^O)iwjU@=b|Iuk=L ztvtWtFx4qJcQm%6q7z4I&n&t_S&+h<6>$6vJBlK-P-MS|P4N~7v&N+V3D$hrbmRI? zC1an)NlZ{}hku{8{+o56Z@^qVUS z9z_SG^Ki{f|1=?bdv12pdc#cRCcOryS-4@_)rReSexh|WqY&M)qsV!8G}sn&uc>B) zw*M5R0m^N?k0*{FPbKpkiR;iqoCZA+3B&!aM&e=+oK~JFDQJVAPIw~d*h2|@1@{^t z3P{C5h_`O_7wt<+Rrvf^^8ByNBX{>L6d{N3Ud2d{cV)$|F7oc`j=MQrp5`0Eb1m4) zdHh~?(C%qta>$13!26f@jwi*GjNJV3X;vdfGEe$kPx#523Gn2sG5Td&#|ZL4GUdr4 z)gBl2=-h>6tT}rtM&xk`>KlE;U|$KN>dU|i-U0XvDtBj%=J+Yy9IF_b4#N<$T^!Gk%D39CfD1|EByC* zTfcurc!G0Dy(JO9+0JB~i_6JBKZ=zRY2vkt*y534RB;SxArly57S_sOFsb5o+g>WI zKclH|)G!+LbY@_nZG3sw=1m4?A3-ZZo!xWzhu7dc##7o&ZRJCgYu6s$Az*c#&2Q=9 z1!WIR7v_}-T5BKoN$|aP!FTi1drj|s_Oiy|iTT9D|{+&ziMcjLEw=b#xGbFaJuEY#(8#}^d!Bk#nZlbLE{LJo#84kaDfA9-ysLt2eC*9QVX$`Wj*_K9HW>mb!*d$ekf(!yZ0<5 zBO|ZiS!LG8D(Qhsv*6&pId9gqWf#mgijmm8Y03b8Ycbs}wXy=0&PNaE*Dn;sfM?$N zfZ??w_j*3H^v(6yOgUS1PMkRL?3&uQ8-7=}8)QF7ySwHRD(p|$8S%|;$9^7O{X*~d zV6Wz7sgG8U-sz=DVJD%&O6@U*PPGVU(d+cj$;a5r25QT}1hcHWJ7HwKf#37Rw@FNr zH@_xPd~0tdk6H-Z!^9KJ>Y857zC0iNcswH?=H@mZ{I=jizNVDMb(s-&zv1ZN&Rzr| zgBwMh%>0yXqZLRGM?3%@t42~Xshcte$fVk8G2X5d5|za;}fW?RnzFHkbl(df+eB9q{jCNqRs&is5>&SCSj zyd?Hm+Jz#!?y`fi{k2JI?;++r+hHqfFMM4r^AqxkoR1GC zA;?_wYv#bZp-Zk5{jN%!qpi8SP)*>J`dNTurQphyD?zvJs>0J=XcOhQB9+9qNj6+vj`;DUDjB7$+pWbTz3-Oz) z_*%cNwI0M7+g9xE@~ZBZsl1rZR-5+^-e<+^KVS+>)^2tqnGiV^MraIu;u*zb28CW= zXDlr4*!O4tz+_z5{nI3c`3g(g-2D6B3`&-Ae`4VgdBOg!@tA-Pz4Ahjfq?4kn2BoIW#+P;{A!0&iAM@zRwF8mdMfcd;KmS~v*!c)mrp^8r zcD?x@Z$A3ulI5v#o?dzS8S?#B`uP*Z+xZ~(5Ko-;_3641yWn`p5>;_WOe}mT?xdzQ zsU3S@>TV9-gRq)~YYpGm!e!F$i!p#B{o-=IxWoxfHVS+=zO`;}WVLDZR~OWM3^TxQ zz}ALs^FLg7H%4N+! z_!pNh03B_YAWzD%YW6-{X|=~sB7SCc`6eIow&Zn6{cg4Cg6qj>r;t48s8Y;V?oR8* zZp}0=7&TJL+32^sb;Xl4O0pSwhbLr{=%$hQe<8!zoV89bFTMq6!fkKcD;~9kqhrRS zwo(Dd?Xvm;BEBZ^wo8gKtIup4ijz)!eAnR2Nz{EEfMv|J@}Pw`%~^Kj&(Glb=Se1u z`^|{2!3oJ*<)mk8E)t>$xeri!t9Y~_mUEpcOlUnk}q-;F@zP`vkUL|R#Rw)mUATk&o ze*Fm$+nXm&0U9R_6zQYy1T)L^$;;^cV%oqe(}_Mm*AoKeMMuD&Cc7=yWKxdl)~T^A zoh|DZG&O5I93z`$*5>tMkKQH0>QUPBCQqTzlwtf9otr{GK2BM3V5WpuoX4VLM zL(!V6Z@TW2?mIOkU#D+mbXr1#eSeM%F~1r}MfpF_0cK^Z%oTCfU-B}z(0#u=ZPeDL9CQ>|LnFg&XtU*i?uGwg_}B+`8-Y* z8ai;{;v1dEDZ`rj`eg9xa9x}({!wYs^iY%H+kn=FJ7mO}3!)FkBs2FR@GWVzJyfO` zn{H}tw!diKD!u;9RqlRQ$u;>h^eMYR?a_f<5fKqZ2#Wgl-+wJwb2{StHMcL<4m?Se zoaQiM)>^N*w#(^tgNXuXOzuY>(cA0D*DfD;dop$ujoE%;RJALu;=`@&3b~Yh3FbUxGUt8 z(y7A_clvgh<*W)kJ`RXPz&_(w5`Tf;F7~El0@_(zHLnYp)myj3ICi++8;fkUd0{DZ z?@U>YJkHLBZyKr0#}DlIWBtZs5ylA`bOz0%z(CX|^FSoZsQoxE`e7P>J+7d7#sT>%s`fgPFbQ4(Ez5*O&12CCEn}JZ4oW zr8Y6mn;3O`G-V0<<%Mzkq!s0*8531bA5>-JJLtahsqgVGcRRHA55nw`6%rSeoDKaqUxvqQd?72b;-h&{D%5h83i{zF7O~$h;5a zkUd=eCP{+epsM$^qUodQx(44Bwp{TZFbiUGY8HK_*f9TNUrq{VFQWCSxaT-^i*FKt zRJ|Ylh;3QtUgKP`Yx78cdaotU>T<-v2X1<4X=&`aLUP>X>wYc}6)+6?GPGVO+3rqR zJ^ES25QA^~V9%D5{FqlzP!Q+Pd8>*03ETcW==yTagBD+1_O-04EfvD-BAk|Gf%0X6 zXEbX$ZzSB<(|N7+WKK@bn4K5fu4jUN53hQv+gr8{LB1W5txQOfwujD^v6PIGEQSXwO*KfFo`2&EbRDwzmwIV(Vwzrle+rg1;>QUi@u+t z;^M6@Dgi09ZrhVHsco?D^mcX|_C;ydVX1Xjr6o&u<$)e90qwra>jP^Gva*C-;ImT* zw#VbDDkTY!>RRgS{lCi+Ye-w1;l-b&8BG#IaQCfwRJmyM;4CWZ-t%=NNYSRe@WEV)}}vvxHEAu zUH=hx-(>8~8J$~|jyBfTRaPcw{`6e`7CC#auO-7ot?i(dJeN+Hr_e~W643{ZN%-+R zgvIo!STB!s_$e#6-vUh2#jO(do|rAd;y|x{Ky|cZA*;J%&GHSB{teE<*8<;`MT@+W ze2%QdJ>f*+&!b-VKVG76#$+FoX4I7)#|3%^qTYw?Po^APP}Ds}z1GU^)TohORMZYf zt(R$f4>y0YlVe2F1acKINoRtRVY9YpVe)^GkwGvwB?dFP`n& zGrqM1M9>eS!FwB)iXP#5S22%{Glq`t-gPz_S!}V`s{iCCYnq)SXX78g`|XFf{D1Hh dolwop3d+v29!D>|4JGsj#RF>c>HAE5{s)RzG7|s* diff --git a/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_4cases_traffic.png b/docs/report/1H-codesign-paper/figures/gqa_decode_long_ctx_4cases_traffic.png deleted file mode 100644 index 12fafe2dc44fdaf833f85730e9c7069c632fbfc1..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 41713 zcmdSBcT`j9`!9-)ZAMX16b)^ZK`GLtcNGDpH|Y=sm0m)x!LiT-7P=G#LlZ=l-Xc<> z^Z)`u4jq8sP@KLE!Vf7oMMF1Tr~7W6cU+)@@q7lBTX_5Y(UxN^#-`<)qHX8Tu7 z7AA+ZyXoy8bmZ*iG^YKTdlxF#&8afXv|`1MSEva z#0uLjKk#6>1E(ZFM^|uNaIpp7`ErtCYR7(bZvW&%C5An zP$yC{M7vEp)(+4;h<|XKk?yus`CL)olEaE!UZn>{#_RnYJsq9!o@%a7ZcD4(*o>2! zGD52hL-Ng1GdWpy-S2EdIAd-cWb?X>k#e(sfgw}Z`kXqP@XId?R@zl37AF#2TMEAE zn-Hy%z0)b)%2@>*5fbP5$|h5@RAc#V2E8ZLr57q(eRio6GfkGp{k{v8HpK%Ts)zXs zw9Km{$KyD1-+%s#>hUup*(ZAtx4Sh&OHYfC;#r;P7I0q}l$9T*VRx5er#dv{9ds;@ z2=(6yVvRX%boW^Wez{3l(>$UZ8&KI&g7e)xz-+IDH!~tx${e7_=vgK9%)>|88e0-X zm!z66H(s*&cmL@x%b8A<@5<-PmR20AsN+7}=cbPl?h)?b*UN{#W!8u}rXAaw{ivBV z_EHCnr`{W^3gas>B>2<%U0X$MwnI){On>{XjK?9r?zDHg$9%tJGhAAm7wTd+yG>oy zd7@Pa5wsNA_C;2|U)46>@5QRi`mU+T`mEeg=ZttA=D!H>;i zCdLkh-`d%Ce+7xSjD8Wxt=gP^>-K5S!O@Y)&xZx<*QCNoTl3{^sf>wr zOBVJdDfb5QT^7qf4jyeH-?H=QP!}It_F2S-O?VBomHt!f{Z=BN&UZ;5fT}8Yf8^r@ zHil1F|J85JAN)9!!oQ^4dz)9~Lp2OpwoN>_J4D+GpniN;+Vz&3u*o>le5NRW(XzaD zi!VmH=w@Q%@R_ou243dG`>x4hhn*@{^#?TjFiL!NLMlfioZ ze6o!XTO8@hs~ptbs`6PM@LZO3gv+U_%3pdH&rYUF4H^Wo@&=YKYs>1y0=DLgUyAH-S4IT4k_=-aSf#ssE9yTm_H9Q< z&u4bX%Wb7%XUhWzgU-o$TF`2@-k20SmfFHzk<1uutxV69Nw$p8Z>{Ox;-f^F3fGJI z{`mLcQqt)^4jlPHzQNLmV`WJZ#kDYtJDNUP4s9(RiIlJMqc^V*u)r_whS{`c+U>o4 zL2vJ?s}~cUuf4fK72-t4Yu{dO#j+K{UL_{W`o>^eVW$w89ZOyh;==B{gSFK@9427S zMuPJn!zW35^6!kW(=t7`VJ)YL4EZg8i|fw$*ka7K0E>Zl-FNbJP`d3x09;*SlILKl zOp;gaX0uqI{a82eyAq3}MK=7#Pn|sT%7z*@w7EmOKeo#GPUVx?-U<(vFI1VZjl=ch zYKE^$rS)9M$19QyTbZza`Sm68o0P~Ux3D{~u9Gp53ddv%>VFAT{JKW)?|!WigxPv$ zR_;L@qV2A`$KZzi*PC}H;w&TDw-!j)w4pE0I8@OVoF`QuV2-+=7WcHAwip@E)Yw+# zzqPc~ZC=IA&}Z}d$?0rdqTc+LiBn}E#b4f{Hn_cP=ADTpNq@e>%O2JWQOL5sQ<|o$ zDO*^4;0kD z$NAOCg?1~m`@7jxu68xj$iAgTJ;(BCyxn=qk!;>6OnH{1%Jo5CqD_9Q+}5bg&RQ|4 zIZvQ;^!d_?Wr}*(M6&O@f=_!7clI2}Z{{6hvPyE-e>72O*zFadca*(v0Tx%5CqCBP z$HsfBOOC`CX-Mj|U|yT)&a0QPWj8s<;(R4(5sKbRKmIHhfi)aCg0*pAIL?VM>|ZeR zuDq+qGhcvbNC57Fugz>vzI#2Zg!Aor_SQO)&_&*LO%~si7c~;6?}fL$Ki4vTIMKKiZO-g|{kjQ-+3p5^#*VCoWiOYl zsVH3y?uGT`7U?_p>rGCqP!ZK2}sGvn^CMt5Gx5H_6XDLX2`*l5@ zFq7UhR5#;2kyN|_8$*0kPI~DJM6JU;YN$c_vp?3CNJd`b{P``c4n}pLnXKo71*xWsuBVIDCVuy?{Sgx zyCp1T^Y3ep60K7Fi?+vhp!xE#diLYSuB3aMIsS8^T~#1yOR>bze!b!Lr7Ev3^J1;X z@vHk))K4vk9WUuo3ZQKbEf@WnSG8$7$RCY`im-Ktv~i-SU4=K=k+ie6eJ|axxRJj+ zr^4Z={Pr)YW!iEI&6Goi@SyZ~M=9(9S3`tFy5lCM-{zU`p1_Yq8^kDddk%VMAgmEh z3LkLqHgg|8i@hDTdojh&#k)5@@2JeLPfKpho5N>*-dqV=m7;Agy9%;cq>$H)pw79{ zc6YYAz4JtG<#^(nm}a_6tF0?ChW*((LYzge#aX1(QO~b!Jw2ze6WKNG|U4cP?N$WR6J%AohEqAH=9OKA<*U0zyb zo`@M%ek4$}k-G90yOZq~UMpEtbpC0`h;)7AoE{4;#edCc>Z!s}?#LmnFS3&PG>+gf zo9i2=7pjN-&MXY{4CNnP9|{P7-RQbM4Rdv7I#;#zrn*88ApdbpMk_(5 z&Ld)kt7r;=zV=M*w1z&5j3f^pO;XjT_0%wa9(HBY*1&=a9OWiq-O2Q%5^Y0hqtA(y zyY8>gqOgc3xT*$^*sFZ5&MLhherFAU+qNZs@jHu!q@_ILe!A(yAJ*a3{Z0U&TP8Aw zeJ&+imeLT)SWONv9!{W(ll)U@PlDtbf{v?@U(}Kpbu4yf$Kp7XolsMZiKEea?Kadm zKz))wvgU}f<~+xAsl>lcadoH01l}xizAh?x#b0$p-9Zq;6koT&5GZ=fbmcVm@6fR2 zUZ(0grWah^ls2Yj!`;oQ%v&*!b1Np3pWiOvS}i5B`HD!W`7PQ^HtOyhQ2bhc=*?#Y z+b5Y~$2TvV8#SbRXlXw-u8S?-T=Uv2_+@EEsqm&_rL99Oc|Mc4a`lFeG1mDOa><%X zA94!a*vqsfR%A5hTUh##a_C~%r%Dg*F&phDzfghL=<|Xh^V(8x4-LE=u05%F>7dbS z^ob*j92!SXUUeEPfDW+SUs=6ev4$2Q2U`DZt~75x6<*JQ*URDCn{-|)gk3k?QxLPzyeUV<0uge_s_f8^ z&aG~J3tK+7KGMEh~hD|DaM_7MsUagTJ+GBIV= z+x0Tp+Lqw!W?(p~yuPtmqO(0|khm~Ro|V8VXBwRmvp4kiwhj6>dZ8vTBL0@X5$Ag3 zQ#X}cY^?!XeM8eLkAIgZ%(@bxQ9@TnHVrn{9dM)vA%3Re3)#aRXPe zA(QQoH@P2wuffw?^p>hA6n%mJ6#Y7Z28ITaH(X{+q~z}1&qhrP-K1UjMc2YNGMcg z65~8WI6RZ;)=8r6EF1_56Hu!wKwPP7ioMuhnN;#A;(b0>zTuXtv1`EDA^g_dvPi0k zgF@Mr7hHSe z>mz7fe;Tq_E!I$?GFghNo(FW;RVMJBahR1cjQSDzx0BDoG=hHwQwmXgGv{^d+xnuN zf{MI^)D4-z0NT#X%G-X;{FJoNI4@RDN0BwEktw_4%};@~qSyQ<$Ton<$3o{B{7!SS z*Sxy=$Yq_Y&~Q6Wt8^f(6XzB-^R5?)PgIZ7yKeepYC?VY1J37G_{1;e1&NuKuigvi^VWIzeEM zNLnTf4GKa7&8dK*M1Pd&)HqXgBNP4kZ>Q;%MD7CaYHY>&Eh-Y`>k3jWHUm4mZ`Vb6 zEq%SX>?)E<(W87($j?##&QIx44j1O^j{nx&4R{$~)lTcIGu`mid;F)^$Nnn>O zn^VQ2Qj)tsL2om}{4EtpnXx5-4WS-U7kOx;t0-5zH0PFbD)#_Keay7_0-zP^0chQ= zk*`>!$MA-o4hsnww(TRblgPUhWY4i*g`ytY$_l)c+czyo#NT*nl5zO}VT(6>)Zb=N zw-ebCT;w9`0s}G6jO=m?(;K(4mt;=SAFSJw!i;9w+fNsCXf)P{bjBHPjK0uZTCE4( z5Y<|Y>%TrCMsLPm2`yNCkiXfVl_8~uB5`Ybm!;8!Yd4a5nqCwDws%z02q#mez@vKb3-`{B2=Q7)g)mTKZgrq_?#C>i zZk$SyYJtxHOCOP&$K=qvyjy*nen4ecJxCkFVXk)N3kMs2^HKupy<`Z6)`P;D< zTh@nj7mrF2XXHf5*L7nhi%k6D3I00XETZ@NEs{8{TU!Sxhs`g>_4*HJ@lG@7*E0QD zZj~ohKgc|XAW_+T+068uf&I*|w{4Mc!C!3WEc06xE?L^oYb zPu)M8Zu@RELBC_h)j~zIh&2>HLXK>m_~AVspV1;nzh9?gv+6Iz#3Z*b4Gu9FE6)s^ z+xzWn5<3m5KWS`IhHvyop=rP8Tady&prh#>~|&rZa)wBJM@8F{Jog3 zuMHR)^be#dI`xS{fAv>VQk^T5=Loa>=bb_*yMM;Gu%q!cicO_lFkQu=4&J6BB7l$y?5SqPv>@9ZoDXfm05-ND!HMzR-)_t}3G@%r_npi@V+ut#){ z;Igg-<12dG%)3`D&mRdnJhREhu%gyjVDhLz$0@&x_SP|-qO3r45+H3aEipZdn+NJY zXfY#{JA(&gFK#oCcGt#E?n0=0g!3!znXO+n$HcvR8N@0;9x5HZ=a(P$;PsOB*&Fq- zZzuJsb)qU~#aGM|yByevnW>KrFDX2x%Qw$IQYSN_It!qD=*I@4@;*DO-5mgD<5Y() z75BSPcTT+ia`hl!>Gc&<<;?>nb>YG+^m(<` zOP-4$NF1yk3Fe6SvUBmJi1(C%iVypg*$>23m>%(mrhpw`h0leZ_yTUhNwF4 z^bH1L{%yDWdMq4^kJNP|B7tEzg+^4TC-l1BWC_5TDorxWiz!cyZsde4>*2yrKcTZ_ zPeU@6JbyhaJ)*FKq{T-Aym=MR7Y0Kftb_}RbO=KW2Z`wpTp`_X(?ufa#s>EU0AJZJ#0cPD$`d_XJxkhfHmJtV=88> z-;ek6PD&1)`_{JWkON(UOQle)xlFTJNV~L&zGtcK;jBS**N@pA#A>@E zx<;5nIqlR-Dy-(t)%&Soc7ED~KGu6S^SSo)(7hIGRIR3Jhr(`OuPVv`EQ5P zu^4(bf#_~!-DSydEZq!>C1>SL>zoT8Q5x=~VkxA2TO;FZmVLZ?=QiayGA6A0qv8_A zg4+w>Lv>3r(T4_jBdO=VB%0bxZJb*uCha)LiIISb`Q;$V*SkMeZ>r$Z#Ws&afbg~V z9Z)=Tn#qDbB9lsdlcgj_Zve4{OQk zdLbk8W8R1Pf?gh{V>xDJ=VM%LyDKFwq~bqveG~dmA12%3w(iZQ^^Vlbgh{Q~@>Xq& z>3d>XT-8gPo$sHVkryS~X>W2na^zHuRaULv(XHEZhsxgS9uWvNu^UEJpYoO3`_;HqUr-ll#__-F?4g!SGIbxp9Zfy~b&{uOp@BEg)QM(D3zQR>?wi0J zjcnZjB%L+C8S%AYCM}DaXd5vT%tbNpe8jpOGDh$>_g+`Nm(`x$uQYUm&3p9Ofyl3! z7mGH+bNo`6QgGSC!dvG;|xjs-?gN={m-$d!*A3@+g>2a(6sjF>OWD?ekr z1;x(+e4*wXF6t?4??_Qn@9<#-++mY`Tjt`|qOnOHz1;Hho0Dl{p`*Ntg5tRZI>mvI zpWh6y&4fmgRsHPLh<2`f`q-OMTw|D>YmdX4vpz(SH+0NXGL;3Ug^w`&s-W#|5i~wl z1u`C$9DTYn{?ch&;4V(u-_#Fj(xumLi@awb42Tr@(es(;GWLmAUFu}QD@$yV<#+R| z7e|UFk8>^!3vtd=!D&yp2MbtfOBH3kzP4x{(jGwhan(VnGW8Bz8^*}bAZ*5|ve)^D zx~VaZbX}ZI4;N6;9v*#sQp4I!l4H`5%P(8gl2dL(TVzlxil5*}$!v=X*813krsu~|?_eL-L@(RX z4{s1g1lF!HaCM))b+C3*+4^)%cSimQ4X{v4E5n6T^zCXl{~X{lI)z_-i&ZcB`Fz|s zkL%~DUC>55O_l^dXG>t!ey<5e&a4^UZ>NiG$)eZBN_|z-0cc4a(~F6FTiLHoyRWP0 zb!zLvbU^qk{%PDc^!k}2203E?(v^Ds)9*CySsNcsGsk^d!gT=uqUgspPw%Gfhyw(pY@E@~O25!CUL`{C<% zaEI~!WNTCAZD`bv&fE2J;=^S*T&1SeQX$#AKOgAd-W2yc*5;Gs{P|BDXFNOTVNNTH zW3kwA98x~p&xc)lPL2=QrK)-kxEt7fe*9}+(e>eHb*H=Yu=C>U{jRZ&P^sEC7Drr3$bJIA&m4Up)R_y-X+q$x znXr*SrzcuQNT#n;A7qc(sw*miexwSvj1vK85occMGqz}PHpOc=#DxUirJ|tea^ol- z^^xCJixBE~Cg$TCq;tq7V;g_2%AZBqfnV~vgXy#`Jy+170%~lI=Qh&nkPJWebC|l+ zDA2m~@fWyC(a1Pz_vmvl=)L#bS?W!a70kaVvFxwvdUjT9Ne#ov5!>c&t6l%$gpE!g3KTyXo;Eq0i$1@5f8Ha zy?1N7dNBI_ONSy5@fxwaI>~F$hfI7xUa_^ONza$06@dI7Rkgc0gIP9(9knbW)C>qe zTf}3wr{4+ahxdj>!lGRET${o%KK39)ZyUFG);i$2nI=2>FGT+Wuc}C92bA&h75IkT zJ<0h;g_`yrzNFYCIu1pl2D>4naoNVvtCazk|z zb~*V>xX6@Qzaefd$qhU9Dz^dJck!X-~izfB@9qieLNaS5Yry7;L$W zL^=VhCQ>04kWO-JCGsYaAG3Y-L@UzA8)iIKD9YVuJ%4=qcWlmRYa`Otd2uk9Y-E)r zdaIF%91zQIwj=_AQ)&Bv=a3mbpk9uKgyCMmrK(7+{g~U$>{OvGCs#b^GqYr+%F>P7 zMvg6Y`}EF8VpnHt{ia;^3^Pfxr7k+fAU{Cg*oIh<8@hhaK5zIV&tP%%MD201kN*Ox zH+{_CoubGTJ%JkzxfnBKBGj?~Zm3SsQPzfuf4RT6;0Pxto({`s`3?-bpx=4lM*@ZboNLIb2}5;ETZj^WiVrk2W?I*)OaQLQ{~pEQ9MPi zvqN420Mi5b2pCbGbzf*>s$jIKdY^SFn@X1&A^WCJl3=4wIGCD7*1MWAC*n7!h(N2Q zuo)t&{_Q8a!;x*xmf&#kuZaU$(BbWC>Tn%ID#+*H%OtHX$8j4^fjy7V5m$b$XA6D!Qd za9*Dun{UqW<=aL64`~7Wj2J%Y`;0zQb_qVdGje#UV!7pItb+>ahN9sqrrFiGtwYir z?;U10LIrXGY1~TM!E6x?tGwL6uov}o2AnO4DNj@!NW41GcoLk)AKP!D5kz69=#H}f1Q z92}8nLq>Q5Aejy5uG%AstmK=s%fntFBdo@_}C(dytC4PIW2FceQ6+t z*h_MSubYn@>t77^|Mrx&PTR#&gMKzFdyM%?uo-}>>ULNuV?1FIDsgT;*0MBgT_^F_ z@Ig}GJ1WnFZejqB{K*loyoG~lXWjp6HBP~5gXNhma`($qQ{#X5H%C$(A!N|V)DLUO zqvqnc%@L)r6Rz48Tq<{K_v8F-h6rRY+IdDjVYv%!a3Hgp(i%3q5#<*084aS0qCaQ~4npL$(oEEYgsTkU%$Q|GJv zr!bsOBj;u6Lj+>Ad;eB(U?hg5GG1M*6H&Q!iYRFt)@~|RAjNN>LFqOJ7)=Pi^5rt#80yPVn)o`*6MI#7`Vcu}~7%*vY;Tu<)|9jo%k-X#dEN z-?1_8?jhl3m8V`Fx|Vuf)&HEj{YT+e8DS^#Mk?cS*aPv;qX2JKL+044Z7N#9!X|ow z-ob~MvXhI}iAEOAJ0lX|V}u>nJswdQ>FLZg75Hj*vm`~W1dUltQ1oLEpTAqTW)iTo zy6geb2F;o9lUN>(f=s_-yVC}whKbz3hw&0c78RhHT})XBU00g@7^fd~`5o&+y-=M% z?GI8adzmJ}H$Gq{vbb2iAgR+$rTmSr2st6Om~7-mQC$8|@3m+iX(^&|qR>OTOj^ZN z@n(5X@*+##4DTL&?B6G`;`x}dw%1=*EOxZx)c3^nza^zzJ&)ZIJe`X@%+Pg|D1(o zJ`e4h<%04Cdrj@uIX0Zxi4<(zYFm>FTJCeht>INk>EC zp9+Pnm{u-O1cxT1?IAmHW%gt1`022)D=QGDQ3W&K-3O8qE_NxT?J)t!S;0F#9qGW; zH)9@5{hY}oX+VZa1OFIytL(uyD4dhW!HBXj9X>O;RtF)0rd9a+^B?T_qy6T^q$t(E zqpVUQ^z3pno>*809nf}P{3Z`+>^s&<0vF&NO4jx+n=CtY>~ zWU%0$Oeu%UXLcY`(!p2U{~R_i38_6qezUSS(8w}Q8JG$s0Y0<=oJAkuJ4K952N#5I z73y2F5ShgWa!EJPdQV&Sg|@xqEaVr?Gy;*Ov5q^<=;I%Zw_RYyLJBLaT{B>_W19xC zju!||tUxUteXtD2QwQMdhFVQKm}F*_6%g`~Ndi3dJM@JzP3ThON|7&=fuj5yOj?vqs>9Xok z5oyo)#By@(FV7)=Pq<$nU0xYafx)c`*Cxefs+DWMu>^qXc8W54Qqh1%pWh6g2m-C- zK_cq1WdHy(Y~r_8rd~kWsk{KlwMK^)Y9kNqy-b(A1@wSuU!iD-EP;pAdEYnsS;UGy;M<3F&C9 z(B(FRDeodCChMkl>pnnaYAJNx()>5S>Y1HJU%9<29ajMi8jytBTwh2@!2)*044@O} zp2Tviq{4?ne#9QRy&Zq?<_LakG>u|}i*1nMgF&C(*1Iy(?Np{-KCKe2G@^0`axIO; z9ZZTY3jVLHgNMtsizgw|63^`P7mj%s=*_^Sm*V=Jnn3j+(cfp6u77KP8lq1qlo+EL z_S@Q?=TMGlYNdSAQ8>>+uPA67-qGV#NVZYgYd$mrL7=a7foa+qpC?iRl2u6>af0Va zY|S zX;ZIlA5!33$ig}?=k86|gyv=V&h9^ZQ0^9mX9slPs4io>TvI8CdPJ&OCE}C_yC8dk z+lz7s8m;m)zsim{=n@?USpR5MxwYBv!t$*UYAVI{>&U}m>N;x$aj@3hZ{!sO$8FxJ zUS<8AWz~Cu?+{|LTjo?kB7DKu3ol2t+<+(E%t*oeJR&gP8r(GFHm(ZQ68TIyq^Pj( zwW45e8(i%tr;RheccIcl7w6$nc55^9qdnYC&AiGr<6a0rwm)Gqj{}%zU{yklG3F7k z74h5nWm;AMs91E9iFm`V_!4{Z2+~DLX+Wt>(y*O>ZbG|xgUSnMtm);#04eg*ZpmQBR zgV@c}X!PR!PgCrMj*`Y=OvZ1EY%1>%qvuOTB1C*vCQ(*n%(;N`y24GA>?6pYB3!(x zI@4e3%5NkKMOSPi8I_7SO7C6PG=}5zurH=SfYE?Us!;}13*&&3LrT02}K`14qdbS(T0bG*O|7Tffhhic>FryRf zjz_!o%wge28(06%&N?1}YP7H=YnsD{+DYU^DNx!VYlArLH)+#Oam-GQ60<`T6r1;M zEs_&&qVbx(V%<5^1bc>m!dL%75Wyn)UUvxlJ zC#j}0(`f}L=uJpqs_}6;rW|7kJx#cFRW(eoN@7pn}|QF#}lBM}ObcIujWJ`khR!%`Z#%^@VbFNpE zuWDm>>eZv84UH%$|EVc!B;1_NxP_z7Nw~vde$}zcp-hAJq8$Nv zt$-Vn0l}h+tiH4PV?X8K8Y0Ew4On(!t5Y2p9Y8nAf_*iGfI_6~nxWL_Qw106O*oZU zG8I0r9mJ0qbI?jy0_WidW4JmF`_@6A_^?4M(8`BaAq(QBUD|4IDv6Ny+-<|g)@sikt@2jz0 zyhcG0=_Hp1?kt6%qhq@W3G})&@^ygDosbNaxd2kJ#cde0yV@8Xi5c0oxhp6Di4nEc zJ1SaJ$AnE(pg^Kb1leuw*C_8{LuVkavn{0(2^TWVI}v3sl(;ZmrmoGVwj$TMIU;2~ zd~IYRY8?a1p#bEp43Nt>o8e8r!3Yv12tU!y>7JZc;A1bC9eUyXW*(>BfsMf1?cf7FVD#zgr$3J6 z2&v_aKq4&WUCD5$)JQA9(ka+Xd6c=0kw{HeX%96NiZAWc#4o+*=fgx~>s1WGuELld zgA;qubuC1@w*a}xrar>gYM?nwvqIfP<?oIF@OUpC{0MKpHgYI{V_y~QgbuiC z-$!){;kb0d($!oJy?|9yFuCpBr>n71Ps5K3PWRfYWKp>Mjlxf!yJtyM3VWWa_>~ff zYNDu0px7v2LGFWbtK4N|=jdwdF zQ#3HT7cHYvFgj>h`C2=}#`;nsy>>z&cX59!QunA(-gAosD3xd49p}6;bMjv+J#HV?yMKY)YJN)KU%_EkIIuI*P(-{7C_$=^RLQ*-FJA|8iBr z>uq2s=22v?we?(}^t3-ZCKrd=e!tJosLDswRtqjX)XpA+M)~s5S5U5{#kClV$=z$< zc@lGr2g|ml9L2E^n_HHkwENCIU4kB{fn=6l8hpyjSur-j+gFSsrPvO9CGpXh-q7c7 z7(#y|1lTvKNUR{-N&Q9z(S)bi5sQ7K4NJ>|;CM0F%i^898x-nH&zFxY*Lrzl!@o_! zgcDs`SEglDHq&b8An4+f8!~`Des`RLsq~lQC!TAsTlHJ8bMlX@QLpWbqFm8F;yIm_ zkq#oOkVDb;`=~6J3Z+4{yeMAc1`=Ne(yVUwNf^ICSfm}$;0u(TWQkIGb#^)OL`Gfi zT);*RyUQP{K`bse$`-1&+7Sg~b_%7Xx!D$IU`<4zLU3n{K2L}2i)r#W)Ys?TS%;i? zQ6sYft)uzYp6f~3`D$?PbPgdW(w^#6?Z~Kjn*-&PG!55ziL={*gy4&r8JUHOWtS|9 z%Hd1zuT%|egE@Tz_0%$m6}#NP5KbY|H~_7j>aMGJ9Mcijp1@FB@36_xqtv(j(4C?bzHC3BVBD{ky+?> z1Tn$ir4Yj{WUSa#PA%%}x&)Q8^jTVOs@py7&n_w$i_rH%Z0g^LsnCq`6r?PE%h6tU78{OhF~g;n|6_wx;tv;hV}Fw+SZG zCaBO~j(tEd*+5d;*-8m?ftoc!w5{(Cl)0-y7(umc^tq}_RwF^Imn&Aufu1d3d(kw& zM(>Z~t8Y>MEtwJqow(p~qy z%79==-!=H)D=#_D5O2ZM)Vhn80d?~GM=N9N*H79}`l|(|Bv!!c{3bmQF}H@1e%Pz4 zN0*_wzd^0}*hgs6;w3DI83KPyT5?9a=3_29I^Pau@|1|#{cBL1loP%m8JwKHAFn~e z@3d=3&qdMTyDxOWFiaeBZ`}ZBzClxAv3TI`d^4=YmgI;j@I&&7X;8Lb)<>|u46Gz(~6Rxh`KhmB`^|UqrxQnNQHxS55qLgcOnej~N&OaNB72I{*v3t^D4Z#dR8Xr91BF@B2=IuJxLZw z`XB5-V=k35MaJyGN$j-TX9Hn3p52bz`~|w!YXE6r-|1B z7Q?V_5JHaa3fEeOov8+HpHY=!)j93nC1}E539^8Yj7vMM_QfJOr5)*TkJ`BzswaRH z86Le}0|X)}LT+>NH_!#(EsBtmCp%!#`L&lr+ZX1bj@ZB*X&p$5n*Dk)#Wc;l^UsqP zM(ZIP<1G{pc1Zg3P@ws_BPD8b@$0XtOIo} zN|m6LNrhFcxhA%0-7J@03d?Fts~}E4vx2@V>wIGF5ycuTq`gFN!?kXJ+hY_c$BQdz zNEdd`8uVLk9-C))+U0nOr;AS|K7vK^GBOl+JOvQ^nBpZHfi2vg{?bvZt`VV`He)U! zm0R4O8+(V>)*%!nF|?PRAfb3g3MkuzP7Hzjb1MwkrS-$$p1vwFMSys`Tyj{1(Ln2^ zCB?5@y}RHOj<+qrm|t#()5fo#W1CvMszE5Caw?%pk8bR2t~R4#2TuS^o}cUGw4tIR@BMrt6{yp!t`Qgn){7m>)dL_;&5Wl%yo(6q;lEFqijyrDm; zLw~4Ij^FwTm;9@H7z{y{-fTquzD@s-lFf1zl#Uh96NzhceG`FCFdu;UjOq7-_?)O} zq%kkHSwb&6eFmvf%$GtiS4X4vrLC%GShXfgz9Os!yIVJ*^$H!1X{=0pM2yz(GmS{? z2sF*QnZ7t9FEESy7F7`e1E6LWkgNa?%OQ`bu>ig7?o9JGRLc!)mGvlXZ${o zHC1l0_N04Yxs|;u$ zN;4VW{?|vL`H);1Lq>0CVCt~gPLKJMLvG&$&@6PpJ(ckG3l#uqB|DF#DG^iyFc^0D19TkJW6+-|#E=jK-x}c6!PsdDd$G4H~>A1JskOGS-Hz+b^Pp!@}Yysc= zLaP;moagK+B4_|@=CjNl;q)HF$KQ?;J2y=YSDpkjJ^{pJSN~LGnDs&fbNxf1nNvky zX(v_tFL$7MHqIr6fhEKpP^E7|jKPcFOsX(-#b#eDM3L^7fmYA~vRx7y_e09~m`*;L zh4Su!5Z>iH&&RK|bwmO~>f?qzc6+O>wK=Za>m=N9(B==bb+!L2zJT%Le% zq4lZ^V>Q+RG81Yj3cV|TE%&Qv+aDb-El<$J8JW7#*o1u;jGt69o5`jz{VO*5Y_2$- z64u63;qDxHw(?N!!{w7Z7hZ&^AO7>Vg!;c`)ongjPXg_V)sS87EP!C{8#HJOl!RX% zqqE{Pn~%n{&RoQ-+t_dH`Qt#rrv$WB=ybE2#j1Yy88mA6wdNR$#QSPI)L1K|D#RB z@d*miDtNv4TaJm>d`Yhazf4^oOKCNcERh`W$RaJP`@x%x;)pK+vGF}l3@GQRr1Medn6eebf9;|U@IIziiIW?8QI0o-Sn zu{%Eu5Y?bE%dwThgYc@^8 z@&z=Z*H*1EYBI4I4-g|AMZI-pNFY85qn?TeYE@AV zBjL2Rtw;szgVnXj=ia}!dO#L2MyHrtIPw3u)qfo$UbKn(pF6#{y_+L+4+=q1X}}}x zvc+aksH*YwVvYKdTn>X`hm>X*Rv#<9Ds8)Vm9! zvz5Oery=~fKGqf?pbkLG{nmCXPf!pW0dDjh1TN1oDYZ0>_RJKKhZ!(;C)j)mztzFp z-ZcLBF!NBouFg;`5COfxAj*gDh?pX3?*>NWxXugk5bsrvKzd8a4E7~lXP#53#KJS_ zLI!5;w;ZEo*Qobm&0^L-rWJwg;?`UTvP~Zz-fe6D8(oD$yN}h@QFJm>Z(cOyarFPi zaQqag*CkSSdXNvQ6u}aYpmhIJ%IHH18ZIf-n?#{KgyPNNcq@Pf#}7OOpes)VVnF=f zO(;XQrWBE#BXb)9XwRA3L+Q5uPn@DBh@n92!soWq@dRE?-fQhJ!Y0fIlh(p@??T

**?Mi-A-EC7D($9RTs0UFOx zI4XHpD(9^FDWZ%e`}GhCjEuiQNx;FMpo681YHfjS`j)Z-DDeyV=897jNMI>{EP_WS zNe^9sB~$DhcShoK)r9&A9zMwPuMQ@#bH+=%O+J0p-&MW^=v|F z4#>ot_XOWm!-Gv8>7Jm3CWE=1pc=1qo6+6e9Q2veT5jQ>9!<)jE|7Q(9U0%Y|KB(Z z2aCPhzQ}Wdeafo7WiU>{@O8|i zP4w&!^q>rSF_zQ++A#d6uh5kE`FL9be7?q->!SN`N63qnAdQbXoS8Xwbs7kG+o(Ij z{BcHw-%X}oAAfDF=Xt#PPSpH1CSMnx$YRA#t)RkXIn|FHV^acFX zb|jWC`-0I|=<>DE%;P(!LE0}MabvICT+Tn3Ma+%^4P^B5#ej(VuyN8n_{R>QwAP<@ z9T>LJ{6Z_I0kdw@KZHKo2(9WZGX(;_g3w4E9V%_NALrlG;4(1ktW9 z?yr84M0=|8oh%HsVw=PugoCka&AGd?OSy*Mu$)?PRw$+&YMDJdmhJ12(8Uz zc+0`np%Z?7B26#)d_@kQ<4Ht$m-&8aKfJuR9pxJwpimMt@(#(Pc^k~0d!1x&csd6W zJ=LSMs`aZ8tO(Z-^uPt7SVuzNSegmL1Es$)oJMKBIST&8U&&t0KVj`#;qgG3rlT&< zYnO>mW0C%2-MlYN?+K`w|HvzIHw9jwYC9y6#*q6T1`@6mggL>fn5NN--`O0KQsJ$0 zOQmG|e#k_bz1`a*S2*+t4@9Q-Bmwek*HPt*E?)q~N0l5G2a&ETz8?5QQi6i^u^B}s zv(6uaK$6WqHXVke-d4i#$~p`}h#*B(zQ}1vE*JU={ZTM)`^J1k5$Ry=Na%N1L8a-P zr3mIGkFkdU8(i&GFR?nPOF9!RNv^+)`gO36iP|nAWuOz0sWF>gvYf;_8G5bBvdgZQ zlojY>D}fL-Uxr8V&^(YnpW&>{L~>tzBMu_v6Uw9)sE^L{w~)Dv8=b)(qySIj=)RtO=8sQODxi^eiRq z8F5pk1@f;IQz#yOX9=_(gEZQ&pEa(cdmAr(Y!VSbuFR+K^WZ3Wfid1dzaND-{egZ= zffw9ToYD&4T{6*LW#l@NBJpY(l;0c#5riqQ*6m#cF`QQs!JGJz|3T0CQY)(LPT0uy zHx#YT0%MTxU3ZAcG&oA&I5Y=(%Wtpk>v#B0VrVv`boyVuqMq~JuA5M((u-k`DE1*C zd++b(MqRPH0Ou9#BR9f;^cd!gp;`+Fp91mlOaKKpq3|EuP65on!!WK_a#z4+8Z4B`8#I#jfieBVU%FZkYZT1%~Z!r&q660n(+?xV{_o zKcp89{Vy0oVt0FK0giy?y8qQxue)0g_gBsr{N=>jZ>c13?d^~4>!h!v79C)cz}|YY z@0j{OUUYPrSJvmydiY=T5B`zm zI?+N!+|5x*;Am82jRgPxH2CXJO_|6yO;Co#-x18f(GKJ{0qY6RYAJLQSWh2Q4OQYY zW(mTXB7T5CzC$fa?);GjVcQG@<>1p$%3yQ^Y3{f0|8vHHLkgPLX;bHFQ`;33I|B7E zi4t}_I<#*h?Tz*0)c?z8-FU18WXcmCr9#|(5W~0{lC6ubsWm+uC+flPr-5P8-4P4_ z-&@W9KYoD*Ftrr08Qjn+WA*4c*DO!l&Vu|nl>#whh+;is6opbKoGxN>& zJo8W6Vo|kg?>DUVOK&bTYf0Dh0i+|y1*hDMSLJ*Y8$3W(K?WK_Oci7Z2ymppbEe;3 z$$c7i0|yJnwWI(WNV1wCdPsT_$U{&NPp~pTs)p}adocehhlqLQr7}+G>sL8MEj0pq zAj-@*0ValoDyV!`CTL2D-TaG63zapnd8H9@jH&_L%9%uC#5sMrVSl5!akCjZYR&?b zhq+P)8%`C&>J$J~D5X{4z2;H`T;=)Z%UtmS-(Z!htid%m${owmerlp}F5>llb@3K& zsd%|ex@Sxstz^K%ONR@pykIWcOaWcF#Mz?0A^^naTY>tLT95)&7*>y{7$jp2*wH{V zq?Ox$jnBkT2Qi}^5PL!b1Y8AsY+Jj+P1&`uEHutAR%r|{x9w)<62Ba}r06sM-8PXC zW9cQU|AzjQ!}VO^fleCRH$vd-XPW5XwR`yVcHk_v}I=5CIL4M>^xXZV#0zR}=#J-wy7k=b=p zftb$Xl_wb|9qEAO>nY0(;sPkEmBCqCA7M0 zjx5;4ni6{L5kqy#^A6iz)nUUP{9=E0x!-eJom?C|crV-ZKHTQgZoQuPBS7KfSu*cl zd4>hl!l}$wLx5e7?yeUh$>Q^XvZGro2J1E4$*G_a2*B=66kgi%Le%#04uC$9 zR}o}sPBX~n7z2B}Bo)OeL^d5LP!a*HlB)txkB(SAIP!E+q~{=kpK8)Wg>Gh*czl9| z4R&UmF7%*$6Xh>x<0FG`;*e8c?Jj{=AAcu-te6g-;nM(ard^YThV!B;)YLS9qaw>5 zG2zEm9UFLidwEh9#}Y#K3|70%mZkCgT_=BpA-b0m-3Yxy5a!*ba1$Jw<+=XVVc8?-6Zl%(iz=PwsLMPZ;PIS)2>E4c9hJcu0EH~N-nt*GPqVmHDXD_6w1M+&FO?3A!4_R#Yjxx z<~hRvr%(dOA%$?zEuexD`PN(jqiC?n1~mYU;|gOfv(UYcpe;jZC`~4g(C(e)fBxdQ zKBRVyT|^QbB$z4JzU*tufwD0PCIk=iecewIA&!V0*kNTF@HJuWu=>=eXeUW^DeU7Y z#_M8AJMkyT%d?Sl?-2l%+NIjMoM*AOBUF`OiF0>_z!8+SH&!YW)f9yyZct|FK!@&E zAC{9b21F)~_!@1;_lw;|Y1$HBwU@)30}^v`7VBm)FyJ47z5#a70VM#PZ~!|Y68z`+ zwa_;uvfM$zfa5x1$+0*1Y~0Kw${d+$4{Rk}{(!qA3isz3UcAwv1KZO?!O@?HD=^CN zb10V5xHO5HyuF1Ez$8`iraj0ItF~_gKI_@k2$TZ}c3q{35IaMW{(&N61Tiq?k64Nb z(6*u&a;bQ`sp%I$QSTCZUGO9WG&EghLc#@`Xpb6&i;f>qaOC%ehNfr{R*&+wf1o!O zo@kW6$aa38-?@3u9i6=6Kf+266sWipZz}{Ss&wUt>Zq{+5YNs`OuM+hIbwCbqcF$d zf@!vG0J+6_az{_W$G92I=xZ5GysFs-+3@%ALoWK=E(pl_xb)DShoni+|Eo_q^@|p& z_aCwy86B8p&g*I9;;Ggf60ewEF&4jR!8DXTYg5KW{cgHis5RA5IiMlsp8<;78G2UK zeBMei!-h{zIRm3Ti)vq*t2c{`<~MQ>fABu~{5cA&Fzhlv7x1 z`BIkSLiiTsJFp(K=bywR2+;#G2~QHxyB$TQ!_aV+ezhP@&@$N@7N0t9Sn zRq0F$DKaL!bjAzC)eih_GJXO8oe6oyK_IRcq18XUb^WQfxS;G5qQ(?i`OZ>+XoT+) zVV|===RbCx=oi!F%u1d^We=?|1h4!SC=SzpA|u_n8s(d0X!{Rh7NvH)Zs;&@V%xS^)csX@;9k=vhiBg?yzKfk~MIRH7 zU(Cn790UZCh)ZB3C?U_$Hw)4wcj+%}U-7Vgf%#b-U0_i|xIO-a8yjTs%mez0(Rf)< zv%344h`z6fXQ5x>3apX6^(5#iuK;sYKz-`T9gZ6{{e?MR^IM7)kWZj78I9ln`_BV+ z>h{|t?pO{^X!uJlm-kKkL0WHq|} z?bC)S@JA#7-I;JI7==1a$!a2PY2#pHQ3K-mtPvtcOgc})_}CY)%K$zntOwA-KLPWj zoZ^N)=7M0E5i{xlEE8~`S>GLU@M!5tGczD935)hYq>D)>htAfI>qb1Wi1R?7k z6LeUN!@*-r2*K1Oy94PtXYs{HgCf%3^?_Taf&ndn?0OE-eMn8y9D?y*nF2nWrftVU zjs$(vm6Rr;n5&#NaRgtmfEYi%@w(W!CD7faU5jG{*r-j$g2NAe5AfCdjU}W z+7J;!Lv$95B^_F?iLaX+?`E}Y$IYL|&9jtiq!G6U_pRV?0eqLYzScJh=`gUem5K7ZcN(Iw35dLy=Io4!Dli-f z%KUu0Tp-hu_{<>EOB2yw>cFzkOa-QIEME3c@PjgIG7g3!o0M})cW?VU%T_6(jnmMD zU@rV+9Ye1=bQ*a-Fa||MQp+p+A=Bryk~?*PnZ5bs$5(O{5`toeDo~-#q?5$DR=yE7 zZP&YmR2IXBLiYJFFfpw@$Dub^{{X%$c1`A4p;-`WNi;K1k_5%FTB6o|T;v2Ym!3XL zIElj;cUC}I+X7L@;rM zcD*U_+q%i{WdT4)zzWEqo0AF!w<8$n$*Aj2WF_WC@^P4HMsT(3C()!44hu958+`!O z+O<$J2e&2#Zq3RK+rMZLB2OkunLlqaC1p$eeN2c($QQ&e7wPejWXs!TCkG=?WKHSv zKho3Ezx1_Lgk~W|Faa~C5j}L`xIl>Iu{{8ZS2zWvy;iFaL71#EBOUpe%SIi14l~U1D5G-jd74`mA~u72yt|I z5kh}ZD~Vn^0zw@cd`y66l-}l7=tCnxz}0DVS#2=s5|+f~A7i(*kaB0ri2ZbpB?%ED1a1_QIqv`9I@USpyFNiaUt7g5&1Z8X@qh3%72}P zC0H~?kosn?n?4u7I&^r@D#-cv(~@c~dCEcJ=-3#h$8$_X$iF3IX&U1ovo!Ft!=?A; zfQQcNP1!%#Q1;>2OiJoz?A}f+L#d+dZ78$_%;^tM{}+iKu7s%@I`CT?+$e2jG&Z<` zu5OG$z(R|svH@x=ke@^e^I1BT$LA!LcDCncz6W$!1?p)Oa1e>h{CoRS488NI#9A{4Dwd*h}eq!DNUCWl}E z-}yDpkgXbOO9;c5DihxIq3(BJ4oD0c2mYZ<1J{E=N;9nEt5h$32_A>GwzqfocwH^5 z0%&s_hpZCxEs0Nn$*kiz1NvVHDI>>6CZczYmoPA$Z)xM-P>_gu{Kz@bFypXdc;q1? zz6HK6tuwpM^YLzlp5A?F&3~*?IPq(IP?xB<(V!hq5}}+h%|lZHZfo5a;%}#Z39kF8 z&HkqwZx=_Dw?>w|%*D`hvqy`BTH=DK>n`$+41UfJ-rWyvgLF547+B@A7KKOdTjPw^ zF&DFTW;(r!*?yy4y?S$Xw`|PKo?QRXS?le=`760HeB`{- z+<&$MjI{a$7xhr<+~)xU_Kn?D3sEmJqjsBrhwR`;{6igL?numg<@ry1SarOM82UoH zDf}feKXi9;9QSrTIbhKCrbJFZ*cjwAG+086ho@1GZ5t-*zKq{CF~%lKvFOzYfr;1E zQr1y+S)e&RpJAoq0bZ-z^uttzR0@F+QX9bRAInwO@tm?%Y`C z0B=VQ2CHg-x8E7C8PcJSuXqaL>doE{Rvq^-0cp_RqfOQ?O~%0|@oC~*I>;T?Z-A>C zNI=I{qeSzQV126;Q6Pa5@Tdt85+$oK;Nos;>oQMunfOApogPvqY_xGdcJZkzBvUe=`pv z^?P;6r}^Gw9PXepFzy|fJR@LR8hary2#iu?4X(n_&t{jm#lWJC!OaO^2uuV_HKX+K z5_(mZ#?aI~x+|E0(W;9WqvGO|Ml=RO0W0D+A3(`s0M`)E+oxq*Rc@@HV_(7nHTjG% zX4DkCGFYem@JSQ~9->rFFf$jnjzlD)u3gz-IvgP8n7cmNaC;C*^; zS}(#v6jj^+6WJrYkyy$sjkvVep2I^Npll&j4R{Pg;}Ml{pDp3AeU6CM9PD3WiOA`Yu7lr+^X0$%_@WXkQVkEN-czA3>HcoyP*a&(~9^j#lP=>Cx1 z;RxG!BC-V?!G&Ne?{SAsLKQa%IQD`LJUXEnR2Co#bB+7@&;9GJmTt310LWFk@f~2- z#$;g21?4&!nC1fXyV;DhGk(e~aZOXoi$g#C+#}8hGO}aD9NqGY(KHmbhqDYsIQU8SSP>UMv!*UMw)#XhI_pOJtVe z%$l4z_!^NIYA_C&yRvAS!~K!;sbs1AR`CHxH27kh;Xm4QlY=x2iNst}^vU(Q7e839 zjI<_Un*{>}Dl@3if<;3}GcPJYhs^7EeYY!QfVdg5Vm_DS!Ht z2&wJoGi@(8G0(o5geS<;}tZPfBvLY?fO&fKu=(4Li%-(YrB zeY12+2jX6;jbL(_jy_u_kH;PahqK__d0V;JcLPXc63vO(Bjg}jNR*-)POt_8Fsn^S z6G3dE>D~rC7N+_qh(XaKe_#Xtx`vs)D7zah9W-EHqKJsxTphmd-Z)M2j)gvjoRLJt zqxP~bRltYf%G|L=LPc4O`z(oQ>HuIOTFmX<*!R&vdZUg>p-5!K+@+yKDv=dV(SeQ= ze?J{pTZz@Cf+@GV8VI z;NVGS_I?2Bu0%P=Vsi-vygL4`HTs~zsNO6TAQDn4t$Gl!@d7`X3A(%}(4ZmXB-Hc` zU?E~6;sZy{&K0ym%Ay_(vAVzQk~pBZ z$)KP{1?o2;Jd=P9%IWI1UXX?o4-6LUg90z!{vwcb8o$j^7OOfF(Mq5`=@{dvT(k6T zx;d9pHA2t*PkhGs`~MZR;ol?AX+NRKL+Ds_k_gEv+XgLdd@(ny_#&o7#1i1w_vLJj z6%D*XxaPXfA^VB$iWvqxlq5CA-*d0D3yN4b+#>B_PkRB7ZW3Co+t)jRiMH3o5%hVd z5ub@d`b7;SfZ^mQqRgPI%ce?bYU(>;Y+Lq6Vh>9TF`=vmP9o0{CN?O4VUW`ckmzT+ zMXyX7SG|g-?_71nAOp-QD=3!!lM_`zTm&1OW^g+!$;+8U_P^`f3aO1}Np+C%8do|1 zU7H|k26L*Zz%Z)}xD#JC{ILO@w+rwSMHJNcLpRMHdkW|+&Dzfhfc)0QBy1=ued&Q3z=M38aWXo=9s z`G;Qknct)fq?B zLgS@XdPm%}L+pHXD)Y)1Cg2F|I;a5VO>+n~L`X`QfM{L4<171B|4k1_37 z=!^|X;(#$pC0Nju9R-5}Kp~B`x`)V2|FcJft8oK0URp-;AY`>PV@f zsh48O9RC&Fr8-gio#k%t--3{&sZWcm{)5cd@8R7tKunP^PZP>A&EmJ8{T~gxE zNZrnWty8eWe~GD={xyF6b-*#~eFp1t=fO2~c3233PhbJuXJJLNc6`*}1jF<#;r#k7 zBJKk^Wjg8f7<;(SLp7EN>die|Ft17I48e_agEDA5XGGyFAZ*VZ zF#+SV%{0En+wtA(X=S0p7wu!W;qvyIL_11h(C9RdLQq64WPq+Gx@m}~P=uMRq_ExY zq@W$K3s#;7?@K6{bEv4dYw~M+ZBc`&54A@N#4J-$_TNroc#Lz0MLgcm(6N$`7xFQ!CyRyA+xpL z(&zIPpyMldS=b;KO&f%*MI8gE&|8RUom3F9&|oFfX$WG#TjYmbc_^RFg``9A7v11-)AesI73LssU#|8F+ zK1>U%U1%G_238vav}(O>3^91rP@uFLwM$OqI^a#_=i%OnO}P_$7Wp=O4eZ?f8jGe& z4pNc&*l{0fZ8^C{Htz?Jf7P3N3KppfjAlTXTCJjRY2)!Spj)2Q_ zy;@K4Oc<>)e=ZsbzhgM{`Fz2#CwFb2A?@Az?cXdPnmF3e;=Z#x*bEJt0ICXXoC9<$ z2~G|pP_0G~rB!9@z9)brAqm_bV!;v*_6HyNACZL5o@?zi@_>S}+x9^ihjDxnX1NDd4M9N&nht#+1z8k=ne zo6<@URK`kt9QsB=;sM)&4GE+)s^IRjC(erGvH&8j2AqEfEwpd&y9IKr$M{ZB0XYE2 z%P`a>p6#jb^?|QDlz8i2fN#rYtoAa&3L-h!C&A?)yjMQyEjl*RjayxfL`!vCDy}C2 ziYe$Bx7k0{m;e^cxq-{lsPEOO=gh_QHMC7*t{dvDogN%idW=zuRj;#gUfeR@#cmU~ zN^S=Lh4UWZ-57JIKZen0Jq!RtXQjGYku-9FD$F*zfLJR;*S@p755xSB#48=j{JHDU z7g4dszT(#bO*L`w-q1C4#P>oeB?%3OLHCfl%!Cet;dn+K6}XZ{=^I32V#&%CVR6W& z80zIPxG^#4Np|O?j6o3oR2N6W-z7H< z0ynTh1w%uZkorZU9Sc7yq<;afEc8r{O>H8|0bWWDvzi@Z?X33&GN7Vpg!X~2VX5- zje&i10if$7IKUCT>FO}^a%NBAoH60yVYj&h%n7d?dLEHZNui;YKY3&pdL?C4ml5VK zAdAASjfrVi1p!O}cEBkn7A*{3$)WT?kfct$G36v?7$$)~M$8YkEIdk(AS0kOP$jY! z0r^R?w1B|AjP^1B_XA4(I74JX6f^n)tw(~Nq#%E0UaDW#uMy!HzX$OT)^qB)b5*bx ze{k3}TXyw6yWE5Q{aSu>&}SO&3Xb1*ic{W?BgptaxqlU&oSC=>ity0HkWV%Du`elw z;789e^b8vI*Fx7OryesZ)HOR<)wYH?&nZH*&FD2lX!sOs)uYk;8&$HT+_u~Wp<(!S z8eh$gXxSVCIJ`I8p;vKisO?5>hIZ13n-=swA;#G;U&M-rgzU2Hv(J?t{5LZt7mZK- z+WYgs282`829-csiEferkyk~a*rAFTl9oJo>db{Os1mhon81gNunn;8=#f2* zho9*qhX(*x#Kk{bpUuAFcd=YUWg-d1n|QDQ9)v~3+Nt_*moJgUX-|xdP?uHmNp)Xd zUPO#IgaO^qm^N$*=PS?zktpTJTZX+a)$1f9_tn$lt#y( z7m`g|QeA{Ami2^dWX?_%F`{B1u*LWaBW?~b#QM036`y}(%N3ZSiRkM4`#1h3aIAKe zsE=3cM-%vrDi;*Ub)ZU!iGi(h(RGfk&2v4c_+LH&0;g#E^q+173jA5N2h;*%V`#W) zB0;uZM*9$yysrK5#jsi!bqE!@a6BAnBslt5+ue{a2#>}<9z+|WL+!J#bG7+gFr8*U z2>a0m=)Z7raw2nP#NQBBxkwY}F_dh(&gG5fdzm)V_70Dq?tD%^LY26r5}YrVlATim zm4FoijyievrK%2sFi5}mtB z)i}r}%5?g_?8@XkB7W^)NxcOvprny zp|+m+m4l`-pbr)u1~u>?kBZ7*=_5!aiMSS3S|~({bWYrmV%YzO(qL(cPQ0>}aetr^ z$&@dX{{fvHLXn*Vg0l@%+^j;R$Xem>A|6$V(oeK266rmTv}qkH2~t#Zh?Im-8#`~7 zd0xF>&73E_bfC}ZT6<;^XUnDYuQO=(G#`;y2k=BRXR4>Tc?Ozl2lG@(IrVKVvw;IY zRK=$reH#@F6CJ33+LiE)k-b)H&VPT&(n;l>EUNPWO6=7?r!dM$w*71WnWxu zu_Dx*5{q*R235qS*Ez3KHJn6F*DiEV9kq@pKc{_RUKyajWwD4A<`o zdJg`fk9!L`=j)dwp;N^(q#iki%8z61)5F|0VoQ#TzIjoR=-uS;e75`ReMn*w?=1N@ ztvm*Th>DfU5AUWN{_1_fYR8_SV$Z|V-pgNT+=y(ZrCkYarrG$rTr)6v$FY#ykr3+E zgBsrNy(O#z5$1eyl-s3SL&It|+l=HlHZ+hX{U zQWU{S)vEj&PqX~1E2Ch+X9z^x?YcL-)Ys#9s7@=vyrH zhssmV5^i|{iYbivpn;hy0%evcoI?l|dqy(EV2Zgxag!o8jS&S< zQ$-V&58SMg0|nJ3h~$1DA2X1YD*)dXT2j@}+8T})VKAdU0c7o^BLK#Uph&3cD_Nd8 zJ&)M{WP++hA#gva6VVu^Z|b#|irUqjwV{%fAO|%YW*(K(c*x}9kT}p(=3zYmI#3VL zZkK|HBX5i41fKV@eNQ{+b!P55K{~cF`rI~bi-9LY98q}f%K9A?)SXFg@>@HLuoa$g zPe&P*uaRhMd5t=r0fc*k7QzHi>TvCpwQL{L-vhuJEJyIJtfne;#Q0$lpVqO_zW>p_-|1n9;^P`aA zDWNqKcTmo&&VSTw224#E+4-nNv<)GgEs*&eNKI_eNpmox?D$#%9lV+5ZL~Q$-+3L~Qdmi2_d_M;aAgo`=LDZ+|caw7~{mC9awb>0eLR zKio(`OEL<12^j$OFm8wFiySSnZ%uCpX(eG)KpkHi4W>X)5(fcz1hK0ykUew2Q+#&94htJpQ+C!&*?vQb%z{j&# z4fH0x2kdn;M(jJGpcfIpVj9C%>f0w_0&u=>Q{4vqMIp6_Xit?5qPWOKL$KbbeGYr$~qJgr_LWmleQs9=i&m!wW@H3!wRk&;nV$64)Aww;l{I*{YwBn^Qma97-`HVEHp<^A2OeOAH|0 zyH;2J=PwrlfDl0QscN*Widq@{3@b}$)2Ry2I@hll0Q>;rYcZXuYx4ekJVXH#*f4WE z5kUu(q|7+`Ti(vmqQtSozvP^kQ?-GTNfneqrH+QYu{M?hy-9BjZQx0le|DS4WQy## zWi~C9?(U$Nj2azG(v;{x6s}ci9Gb$YE8qNDtZz4%nalK^hgFu2cN7CjM-OZBr~Yyn zxC;kPW?+y(rJ2M^^nxJQ>7BUO^E>-}aTCleM*#~sgY5A5g>X@|dj8k2viEL$i8Nzl z*SVXqQ2^$`BpCuPA7L@V^GzH9O-e%4=h`Fhe2f~n`I5_f<4YgG+~jvN@@*1FCoC&^ zlakMf29E*Lp^O;~MtZTt4>whOhUOE=kMFJwJR1XblIq&GmZ=nEIaberLxQ?lVf$kY zR7pTjqnKy%fN1qOc%d5s9*;BHhTTIr^CFvPgP$G?5YS#zIS!L5`=-G==m)6Z?ZhFj z>3nZJNNsd`vwZt-?fHD99qL(SoO9_4JP#s(wG!$ zer@RFdIj=PYhjQI%-=*^6}Xqg4!6SRS1>k$pA1u^(Gid#T#0(uD+N}kA;6Ogx{WXl zc1$o{=s$x82rGgV&vK^mW2f`RTQE$K?{rAAlO+L zs}QA$j&K-H1*0pCZ7^(G9}ZP2&xlVCgQbBV27pDkt#CVNtMx2eaea0J)XXWDtxLW@ zsAO;>a;Z`nP!CJKmc?{3GjD(HFR2O!N95dKqalI(xN^cF3J%rI;XWAaWM$o(gh3HY zdyc|_2e_~maLIxr_&P2nEmzqATOxA>OLk}D)G8QJmWvDwo!0I|XZS=GFh2s17#I~a zkbV{R+TPiKr#YEErtwv>y0ImvR+BOQ1y5ln!Brn@JjQI?339Hc^Bt9-?bcj;;cMI) zXTSL7bUW;#4hy0v?Xo@m9gHR_NU>qyUd2S?c3=r+PPl&B8=`_rPC|}l)`H+?`f@Uy zY2)%`dha3G{LbYzcvuUjuN!4y3)T3Tc|8++j&Ois3?>|PoB8FwM6(*MhxVAZ?Jkf_`o&ysN*0@ch%pU&kyw$0j!ryjHsXcnlg8wt_Hpi}*lc7FTV5lIt)c&e&{!tPlc) zTo9U$sH=G+HPCVx(6Bf7>C0KGTMdb6kuj$2eo+yoGfRKL6K$y=veP0Cgo(O?$yg-N z%<-@1mVPrH!5)0G@u_}#vbPF~MjHsxvhm*`7tLGK1iO4CZ1C2py3fqlunAuFW;*z*jzeN&b~C|Idq5`eb4ARV ztKrPvhm6_6fHl8=QlvXA(m>c{Ydi+hINo!5aVJ;P5)_$5+-4me2AeYv$SHh`I4(T7 zXUZ*2GyVZW_R-R0U#-d!nH@7@y%J{Sfx_@I5hxv1#aQRmrOIyq>V0azDi}^)WQ@gR z6csB4@r5NP?|Tg=GF+Ba_~`afz5Nfvkxwy0e&*GYq;PQ?tM~8Ut2{k^Uu9;rTXn3b zstQzWmB5onTe@{_&F3z<_sX!j!gQ=eSF$KzS69pAKj~84C5u2c0ng4pjsTybgMlTlZPATrt-L1%-o* zjXPv7>!qJ3U9pvA%sad`t6;sMJxG1I0W`T^=Mw5&4@gN#)eZBiM~i}4N_6w)&8u0S z3hHvqQCH}))Z{ib&KSSCc{GFHyQiy552s+gu&gE-O*^% zlGRkMm7#f|=;-ZVCXdnlxsFuA2>9VL0}D$P^buQxSd;cXSi#4Bz__q=9#kw@PiPFM z**u4oW4*=zHv@~nlr zYiERaorUv3oM5bmbxMI+aN~4XB4wl0&{%H`Z@3%?MUNC|B`T$qcKRBv9naaLPyvTR zMOxw+2a(_fZZX57Cj#A$wI91#&TZtUKI9aK!dxd*ZLHr zZ#!_dtFyBbsuAz%PQ0IX9r$9<1^0Z+e|%(}qI) zCcVrYuj-dK7C!oou$=RpK6$gefBJiSJZQG!;Woy&cIc$(i+!HjrS?vF#OIY^DSdmt zthRvn0?UE3wOPNf|IIAF_GrnzpD_StG{-=1ai>c7Nn~K~3hTPKQ{pjRxtaXKNq2@J zF@(tV`~E&9_PJnNzfpaWhet#&sl=Ow45gkqz^TVx7iqgrD?VK#wX)V( zcP7MqXAs6DDS@7iXqTR(aW*ZdMVTrPaU> z`tD)QLpw&irUO3}iic+Gq-+_g+O+LITO z1%Wheb>5(H1JDwzcYxvEbSFzNb)x_C>oKRgdfIbX$ZGU9ym5T|d3JW2O1io(E5aO; z@1L!G!MA7lP1lCGz(WE$VJXt_PtMEk_W;B|7}n)NnVgL{E4IBq#t zE|OTTo*7y{T%M5>O1|^RUcy?A#-Y5t?nD9*L(Q#chwR=2o{-GpoceBRU`~6xkL0zm zYW)Y^PHnw1-Zsv55PdlB)Q~ya=oZg?n*PmIe{<5;AY9(M%#K8zm{>(Eh%7HlebLbS zo$*#yEO6xbNwul&mG(aG@2_o<-d*k%(vW}X_Go5~%jv)ns>ec>O$*OC;~7=TWxzBh zr{InpzHs4!qJ2RQHvl0K8HYk!HOqHPnY?96f+gIUGkgu}fj&_Z^+TAn%PWVG&XQYY zp9|e?+}Iut)LU1epkZ_MmB#d_06iepGc#-&o`=cs4CY-CtflnCSv6Bg?+XFXujwhE z&VT&)@np=?QKy`;?UvMm*T0up0#osW0Y|Oajxn+D>s3|{!}b#XXxmc(7>TjN)_`qA z@#&`S`LM;NCebXs2-fC3!lHHd*(v|J(e4V~n??{uMVV<7 zExPpSNj2cM)SHzb^B$^yNc!Fou1GQC6A6B4vnX37@r`7V>&SI*Aw@4C;U@B547+wA*+%Id4zg4Ks;yiPmd zx~$&vnwHn3+&JMRqqQ~~wzx+B@SzH~sL_Z0UZL`VjR9_cqbW(rvnh&c;-AAcckh|g zQr@o`7px^_RWF1+(i7WV@A?`uVz@@^Ajy4on;C3mHdJK+6EUxd3WeRusV&~a565k7 z>rHYMr(-O-qZDgidS8+K&{^UW0kf?wD}2IDYTZ3o+`pYtRwUbdbiC#%ehbu&XkJ&x_(|E>GmhQXnp8{+O^Tezh_HcZXk!c%mn|?lY-l~oP zl}K|#%A(`;apm%eKTe5on5^xWn<<}EqNF}Md+J@;?W8$*D)PAP@g~9Hpus?ib7m$b zN4m4lgq>{xpJ73ij=7%quO2YsnOWj*so#>h{xeo!oPU!=;vuv0Jz+`GBKH0D`{!s? zWi!xNec3PB-@GR^vUV}aQCQ=f9*cu*cPo@9m$Nn6ZuXZtL*PBBuxWh*^cz2%&-dGo zf}L`2itqa1+!lXsndVD8*}uPK*}VDjkat~ZC_9a;8lloGa_`W=58P98smE4L4tHGH zK|OHp)rWkL6(St1cvihdFLga**fe&KJn`m`6;Kwd?|!74pJ$;AiBnI|mmP{x04AczHCxZ%5u5sumr`?*Ak>m4mmqSKg zMRD`&pLoKP?VCCZl0=WoV=er90}hs7f9mxvgsDhR+e&iVINQDsM&WV3k8{c?Lfn=o z)wACk#+~%Kr}rRkbIWHb{v5^KQ9jX>xr558(1BHgU1yNtrf{k9vXu*YbF$N}x}AGd z*PSc8nd9Z{z2>m~MYRz*k57J5t1s~!Ea4u^mS=BZc4c|Us}gD=ex*I~#rYP!XG|r& z-?NTDx#h@pMm~APF#r`V@3Mx9JFDdOnqjNYfdmQXTYd`KesMR zGEDuA(;Dkedxdm{;J#SlToo>+oUBn*J}PX_#wUN9k^DqWN2BQGjJ>zO4ntdsy)s_o zNfB?pPiIArt8X&Owl$oCDkyxCtg&L)MDQLx(dM*X+f}QM8x#Fz5NgIOVp+TM#?!iw zR3-bgcV`YgJH_|M85iNnRY(o>WjO(BF77@lt{-k4gMQD1@hUk{2!9@{OFdcr^`|TP@AP{t`y8Och;;h|$K~aF`S>UwGERHdYWtjHnJE9h>^=bWO!AJw z>=C>G4w;nqI=dfQ{YrZK?6~A)!buk2@=uB7^4ULQA9sA#Ke15_#(VI0F+?UUZ+}yk z3`G6V&%Dm4QEJM4i`dBP@3e9^ck}A9sp~#S&3tC?MstH))ExBH5rBx2Z*U&9mB0DB ztXY|5dwDlo^2$1I7XHhvpd-w!N?en0Y9FyCVm`nvmanrvbkvAPx5M7|7g@*6KN`{^ zs`khI`ZG5+b!`n4H4Lf7)c=>E<1r{jKcXvwm%X&gDK zQ`_)%cdkY-pKiI)E~xMjNZG#dqkwzAJ{+hI``f_P0DChGRfoY+0dqf8`n}} zzU8}&RRSnKSYjmbilXW?Q|sPeAS(1n#F@P!3$PZML(w?tmA0jCj_UXR`MK9WN-SVk zbJpPLaEJ&COyaBEc_>cLSgC*g_XG_;nUuOIqq!rdA6Dhncx!^rEAwuC%FgFw8ZS*x zaVqO5rZ_xR(qXhldtYewh&quuuW`n$gjLr+RZFD6(Op-=%g;{?N^3QWN(z(CChCk% zq&D~G=W$ARzyG$aR+1Xae^hUm?B7u@d$#_BsZNdG>tvT)^SFl>{O z&`&!<6vx+lo}NcmtYvjn3f%V%H*r*tMnbE`tIrsx)6}4njf5&&@|5>hyKspsL%8~S z`TI9qJjyF8D|?|$(xK;z%WRIxyI)`;OQ}KE8~O0zs=0P}3aYQ<*aXjC1-V#neTH$z zetk-^_VcsCNdU#D>F5N$29hn-p9BKwYQWH=V7Zj)t_V>-Oa^h7B;X|W)sd0`mTuSN zl2jvlF4cqEEDB1IJ$du6C3{tCTA5zFI9Lc+%i-h+C)-SSd+oDl&(^jyxc7?*=u=KR zb4@us^72|e>vkKS(Yu?*ES{9=U;%?IJ!!^8?Rc_M6cBSQV;V5Hl60TxRRfH`T*5!- zW0VIODV$xDXn%gd9bj#9d16D)dv;N+!eR@wR-IDK&ia-+C^fw`G3ri+Qx3|ca8Mr9 zwm8>>$cB8kbOB4zU3c~C%dxK29ihWltM=U{5(zEZkrJ+RaaRLP_`}T-Mqt!0E7lwa0bV@q`8g4b^1%Jou<>k3Q{Q%wVwmJ% zpKH&`#I!R_^tz}1{nE0s7S^3ie;8~6Ipjx9sF{(j1y!O0tT9Xf`&CkBNde2m^m98* zxTT-|+z$1}U;dYmI5!jic6|d&{L&AZq&>X}IS%|a`b5dTKYv>DbV&~c(x6_AfAnGh eH@)phIAi79o3$lfY=Q*QtawW8WcqQFJO2&bxespu diff --git a/scripts/paper/paper_gqa_latency.py b/scripts/paper/paper_gqa_latency.py deleted file mode 100644 index 1c0a0ad..0000000 --- a/scripts/paper/paper_gqa_latency.py +++ /dev/null @@ -1,101 +0,0 @@ -"""Report harness: GQA end-to-end latency + op/engine breakdown. - -Isolated under ``scripts/paper/`` for the 1H codesign report only — it is -NOT a registered bench and does not touch other people's benches. It -reuses the existing GQA headline panels (the real GQA kernels wired in -``milestone_gqa_headline``) but, unlike that milestone (which records only -op-counts), it also harvests per-panel end-to-end latency and per-engine -occupancy from ``result.engine.op_log``. - -Latency definition (same window convention as ``milestone_1h_gemm``): - end_to_end_ns = max(r.t_end) - min(r.t_start) over all op_log records. - -Output: docs/report/1H-codesign-paper/figures/gqa_latency.json - -Run: - python scripts/paper/paper_gqa_latency.py -""" -from __future__ import annotations - -import json -from pathlib import Path - -from kernbench.benches.milestone_gqa_headline import ( - _PANEL_DISPATCH, - _PANELS, - _make_bench_fn, - _summarize_op_log, -) - -_REPORT_DIR = Path(__file__).resolve().parents[2] / "docs" / "report" / "1H-codesign-paper" -_FIG_DIR = _REPORT_DIR / "figures" -_OUT_JSON = _FIG_DIR / "gqa_latency.json" - -# PE engine component suffixes whose occupancy we break out. -_ENGINES = ("pe_gemm", "pe_math", "pe_dma", "pe_fetch_store", "pe_ipcq", "pe_cpu") - - -def _occupancy_ns(op_log, suffix: str) -> float: - return sum( - r.t_end - r.t_start - for r in op_log - if r.component_id.endswith("." + suffix) - ) - - -def _end_to_end_ns(op_log) -> float: - if not op_log: - return 0.0 - return max(r.t_end for r in op_log) - min(r.t_start for r in op_log) - - -def _run_panel(panel: str, topology: str) -> dict: - from kernbench.runtime_api.bench_runner import run_bench - from kernbench.runtime_api.types import resolve_device - from kernbench.sim_engine.engine import GraphEngine - from kernbench.topology.builder import resolve_topology - - topo = resolve_topology(topology) - result = run_bench( - topology=topo, - bench_fn=_make_bench_fn(panel), - device=resolve_device(None), - engine_factory=lambda t, d: GraphEngine( - getattr(t, "topology_obj", t), enable_data=True, - ), - ) - if not result.completion.ok: - raise RuntimeError(f"panel {panel!r} failed: {result.completion}") - - op_log = result.engine.op_log - kind, params = _PANEL_DISPATCH[panel] - return { - "panel": panel, - "kind": kind, - **params, - "latency_ns": _end_to_end_ns(op_log), - "op_log_summary": _summarize_op_log(op_log), - "engine_occupancy_ns": { - eng: _occupancy_ns(op_log, eng) for eng in _ENGINES - }, - } - - -def main() -> None: - topology = "topology.yaml" - rows = [_run_panel(panel, topology) for panel in _PANELS] - _FIG_DIR.mkdir(parents=True, exist_ok=True) - out = {"version": 1, "panels": list(_PANELS), "rows": rows} - _OUT_JSON.write_text(json.dumps(out, indent=2)) - print(f"wrote {_OUT_JSON}") - for r in rows: - s = r["op_log_summary"] - print( - f" {r['panel']:24s} latency={r['latency_ns']:10.1f} ns " - f"gemm={s['gemm_count']:3d} ipcq={s['ipcq_copy_count']:3d} " - f"dma_rd={s['dma_read_count']:3d} dma_wr={s['dma_write_count']:2d}" - ) - - -if __name__ == "__main__": - main() diff --git a/scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py b/scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py index 1f76d44..dd11408 100644 --- a/scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py +++ b/scripts/paper/paper_plot_gqa_decode_long_ctx_4cases.py @@ -1,12 +1,13 @@ """Comparative figures for milestone-gqa-decode-long-ctx-4cases. Reads sweep.json (emitted by ``kernbench run --bench -milestone-gqa-decode-long-ctx-4cases``) and writes three PNGs into +milestone-gqa-decode-long-ctx-4cases``) and writes four PNGs into ``docs/report/1H-codesign-paper/figures/``: - gqa_decode_long_ctx_4cases_latency.png end-to-end latency per case - gqa_decode_long_ctx_4cases_traffic.png ipcq/dma op-count breakdown - gqa_decode_long_ctx_4cases_memory.png KV bytes per cube per case + gqa_decode_long_ctx_4cases_latency.png end-to-end latency per case + gqa_decode_long_ctx_4cases_traffic.png ipcq/dma op-count breakdown + gqa_decode_long_ctx_4cases_memory.png per-PE KV bytes per case + gqa_decode_long_ctx_4cases_parallelism.png per-PE S_local (compute work) Run (after the bench): GQA_DECODE_LONG_CTX_4CASES_RUN=1 python -m kernbench.cli.main run \\ @@ -24,11 +25,12 @@ matplotlib.use("Agg") import matplotlib.pyplot as plt # noqa: E402 _REPO_ROOT = Path(__file__).resolve().parents[2] -_FIG_DIR = _REPO_ROOT / "docs" / "report" / "1H-codesign-paper" / "figures" -_SWEEP_JSON = ( +# Sweep JSON + PNGs live together under the bench output dir. +_FIG_DIR = ( _REPO_ROOT / "src" / "kernbench" / "benches" - / "1H_milestone_output" / "gqa_decode_long_ctx_4cases" / "sweep.json" + / "1H_milestone_output" / "gqa" / "long_ctx" ) +_SWEEP_JSON = _FIG_DIR / "sweep_decode.json" # Panel name → (short label, case ordinal for left-to-right plot order). _CASE_INFO = { @@ -98,39 +100,63 @@ def _plot_traffic(rows: list[dict]) -> Path: return out -def _kv_bytes_per_cube(panel: str, *, S_kv: int, h_kv: int, - d_head: int, C: int) -> int: - """KV bytes a single cube's HBM holds (K + V together, f16).""" - # Cube-replicate ⇒ each cube holds full S_kv. - # Cube-SP ⇒ each cube holds S_kv / C. - # PE replicate vs row_wise share the cube's HBM and don't change - # per-cube bytes. +def _s_local_per_pe(panel: str, *, S_kv: int, C: int, P: int) -> int: + """S_local (token count) each PE attends over locally. + + Encodes the cube/pe sharding axes from the panel name: + cube_sp_pe_tp (Case 1): S_kv / C (pe=replicate within cube) + cube_repl_pe_tp (Case 2): S_kv (pe=replicate; only 1 PE works) + cube_repl_pe_sp (Case 3): S_kv / P (pe=row_wise within cube) + cube_sp_pe_sp (Case 4): S_kv / (C·P) (★ 64-way split) + """ S_per_cube = S_kv if "cube_repl" in panel else S_kv // C - return 2 * S_per_cube * h_kv * d_head * 2 # K + V, f16 (2 B/elem) + return S_per_cube // P if "pe_sp" in panel else S_per_cube + + +def _active_pe_count(panel: str, *, C: int, P: int) -> int: + """Number of PEs doing non-idle attention work. + + cube_sp_pe_tp (Case 1): C (PE 0 of each cube; 7 PEs idle per cube) + cube_repl_pe_tp (Case 2): 1 (only PE 0 of CUBE 0) + cube_repl_pe_sp (Case 3): C·P (all PEs busy, but cubes are redundant) + cube_sp_pe_sp (Case 4): C·P (all 64 PEs doing unique work) + """ + if "cube_repl" in panel and "pe_tp" in panel: + return 1 + if "cube_sp" in panel and "pe_tp" in panel: + return C + return C * P + + +def _kv_bytes_per_pe(panel: str, *, S_kv: int, h_kv: int, + d_head: int, C: int, P: int) -> int: + """KV bytes a single PE references (K + V, f16, 2 B/elem).""" + s_local = _s_local_per_pe(panel, S_kv=S_kv, C=C, P=P) + return 2 * s_local * h_kv * d_head * 2 def _plot_memory(rows: list[dict]) -> Path: + """Per-PE KV bytes — Case 4 wins (64-way split).""" rows = _sorted_by_case(rows) labels = [_CASE_INFO[r["panel"]][0] for r in rows] - mib_per_cube = [ - _kv_bytes_per_cube( + mib_per_pe = [ + _kv_bytes_per_pe( r["panel"], S_kv=r["S_kv"], h_kv=r["h_kv"], - d_head=r["d_head"], C=r["C"], + d_head=r["d_head"], C=r["C"], P=r["P"], ) / (1024 * 1024) for r in rows ] - # SP = blue (efficient); Repl = red (wasteful). - colors = ["#3b6ea5", "#c0504d", "#c0504d", "#3b6ea5"] + colors = ["#888", "#c0504d", "#888", "#3b6ea5"] # 4 highlighted, 2 marked red fig, ax = plt.subplots(figsize=(8.0, 4.5)) - bars = ax.bar(labels, mib_per_cube, color=colors, width=0.6) - ax.set_ylabel("KV bytes per cube (MiB, K + V, f16)") + bars = ax.bar(labels, mib_per_pe, color=colors, width=0.6) + ax.set_ylabel("KV bytes per PE (MiB, K + V, f16)") ax.set_title( - "Long-context decode 4-cases — KV memory per cube\n" + "Long-context decode 4-cases — KV memory per PE\n" "(one KV-head group; per-layer, per-token state)" ) - ax.bar_label(bars, fmt="%.1f", padding=3, fontsize=9) + ax.bar_label(bars, fmt="%.3f", padding=3, fontsize=9) ax.grid(axis="y", ls=":", alpha=0.5) - ax.set_ylim(0, max(mib_per_cube) * 1.15) + ax.set_ylim(0, max(mib_per_pe) * 1.15) fig.tight_layout() out = _FIG_DIR / "gqa_decode_long_ctx_4cases_memory.png" fig.savefig(out, dpi=150) @@ -138,15 +164,44 @@ def _plot_memory(rows: list[dict]) -> Path: return out +def _plot_parallelism(rows: list[dict]) -> Path: + """Total active PE-token compute load — exposes Case 3's redundancy.""" + rows = _sorted_by_case(rows) + labels = [_CASE_INFO[r["panel"]][0] for r in rows] + total_work = [ + _active_pe_count(r["panel"], C=r["C"], P=r["P"]) + * _s_local_per_pe(r["panel"], S_kv=r["S_kv"], C=r["C"], P=r["P"]) + for r in rows + ] + colors = ["#888", "#888", "#c0504d", "#3b6ea5"] # 4 highlighted, 3 marked red + fig, ax = plt.subplots(figsize=(8.0, 4.5)) + bars = ax.bar(labels, total_work, color=colors, width=0.6) + ax.set_ylabel("active-PE × S_local (PE-tokens; lower ⇒ less wasted work)") + ax.set_title( + "Long-context decode 4-cases — total compute load across active PEs\n" + "(Case 3 replicates the full K/V across 8 cubes ⇒ 8× wasted PE-tokens)" + ) + ax.bar_label(bars, fmt="%d", padding=3, fontsize=9) + ax.grid(axis="y", ls=":", alpha=0.5) + ax.set_ylim(0, max(total_work) * 1.15) + fig.tight_layout() + out = _FIG_DIR / "gqa_decode_long_ctx_4cases_parallelism.png" + fig.savefig(out, dpi=150) + plt.close(fig) + return out + + def main() -> None: rows = _load() _FIG_DIR.mkdir(parents=True, exist_ok=True) p1 = _plot_latency(rows) p2 = _plot_traffic(rows) p3 = _plot_memory(rows) + p4 = _plot_parallelism(rows) print(f"wrote {p1}") print(f"wrote {p2}") print(f"wrote {p3}") + print(f"wrote {p4}") if __name__ == "__main__": diff --git a/src/kernbench/benches/1H_milestone_output/gqa_headline/sweep.json b/src/kernbench/benches/1H_milestone_output/gqa_headline/sweep.json deleted file mode 100644 index e73d9cf..0000000 --- a/src/kernbench/benches/1H_milestone_output/gqa_headline/sweep.json +++ /dev/null @@ -1,31 +0,0 @@ -{ - "version": 1, - "panels": [ - "single_kv_group_prefill_gqa_c8_p8" - ], - "config": { - "T_q_prefill": 4, - "T_q_decode": 1, - "S_kv_prefill": 16, - "h_q_decode": 8, - "h_kv_decode": 1, - "d_head": 64 - }, - "rows": [ - { - "panel": "single_kv_group_prefill_gqa_c8_p8", - "kind": "prefill", - "C": 8, - "P": 8, - "T_q": 1024, - "S_kv": 1024, - "d_head": 128, - "op_log_summary": { - "gemm_count": 1024, - "ipcq_copy_count": 896, - "dma_read_count": 192, - "dma_write_count": 64 - } - } - ] -} \ No newline at end of file diff --git a/src/kernbench/benches/_gqa_attention_prefill_long.py b/src/kernbench/benches/_gqa_attention_prefill_long.py deleted file mode 100644 index 77970b8..0000000 --- a/src/kernbench/benches/_gqa_attention_prefill_long.py +++ /dev/null @@ -1,190 +0,0 @@ -"""GQA fused-attention prefill kernel — long context (ADR-0060 §5.5). - -Head-parallel: each CUBE owns one Q head and one KV slice. Over C ring -steps the KV blocks rotate around the C CUBEs (W-send / E-recv) so every -CUBE sees every block; the online-softmax merge folds each step into the -running ``(m, ℓ, O)``. No inter-CUBE reduce — each CUBE writes its own -head's output. - -The ring is **tile-granular** (ADR-0060 §5.5 + ADR-0063 §A.2): each -ring step transmits ``n_tiles`` tiles of size ``TILE_S_KV`` rather than -one full ``S_local`` slice. The kernel loop is nested over -``(ring_step, tile_idx)``, with the bootstrap at ``(k=0, t=0)`` -establishing the persistent ``(m, ℓ, O)`` and every subsequent -iteration folding its tile in inside a ``tl.scratch_scope``. Per-rank -persistent scratch is ``(m, ℓ, O)`` only (~1 KB); per-tile in-scope -scratch is bounded by ``TILE_S_KV`` regardless of ``S_local``. - -Topology / SFR: - - Requires ``configure_sfr_intercube_ring`` — either ``ring_size=C`` - (1D row, ``C ≤ mesh_w``) or ``submesh_shape=(rows, cols)`` (snake - Hamiltonian cycle through a rectangular sub-mesh, for ``C > mesh_w``). - - ``P == 1`` (default): only PE 0 of each CUBE participates; - intra-CUBE PE parallelism is disabled. - - ``P > 1``: all P PEs of each CUBE participate via query-axis - split (ADR-0060 §5.5 last bullet + §B-item-3) — each PE owns - ``T_q/P`` disjoint query rows. Output rows are disjoint per PE, - so no intra-CUBE reduce is needed. Each PE drives its own - same-lane ring via independent IPCQ channels (P parallel rings). - Requires ``T_q % P == 0``. - -Layout caveats: - - GEMMs use ``tl.dot`` (no composite epilogue / ``softmax_scale``). - - K loaded as ``(d_head, tile_s)`` via byte-conserving reshape of - the deployed ``(tile_s, d_head)`` shard (ADR-0060 §3 - reshape-not-transpose caveat). - - No causal masking / step-skip; blocking ``tl.recv``. - - ``TILE_S_KV`` is assumed to divide ``S_local`` evenly; last-tile - padding is not modelled. -""" -from __future__ import annotations - - -TILE_S_KV = 1024 # ADR-0063 §A.2 S_kv-axis tile sweep (per-tile width). - - -def _merge_running(m, l, O, m_step, l_step, O_step, *, tl): - """Online-softmax merge of two partial ``(m, ℓ, O)`` triples.""" - m_new = tl.maximum(m, m_step) - scale_old = tl.exp(m - m_new) - scale_step = tl.exp(m_step - m_new) - l_new = l * scale_old + l_step * scale_step - O_new = O * scale_old + O_step * scale_step - return m_new, l_new, O_new - - -def gqa_attention_prefill_long_kernel( - q_ptr: int, - k_ptr: int, - v_ptr: int, - o_ptr: int, - T_q: int, - S_kv: int, - d_head: int, - C: int, - P: int = 1, - *, - tl, -) -> None: - """Head-parallel prefill attention with tile-granular Ring KV (ADR-0060 §5.5). - - Tensor layout (per-rank shapes): - Q : (T_q_local, d_head) one head per CUBE; replicated within - a CUBE for P=1, or sharded ``pe="row_wise"`` for P>1 so each - PE owns ``T_q_local = T_q // P`` query rows. - K : (S_kv, d_head) sharded cube_row_wise → each CUBE owns - (S_local, d_head). Tiles loaded as (d_head, tile_s) via - byte-conserving reshape. - V : (S_kv, d_head) sharded cube_row_wise → each CUBE owns - (S_local, d_head). Tiles loaded as (tile_s, d_head). - O : (T_q * C, d_head) sharded cube_row_wise → each CUBE - writes its own ``T_q`` rows; for P>1 each PE writes a - disjoint ``(T_q_local, d_head)`` slice (no intra-CUBE - reduce — disjoint output rows). NO inter-CUBE reduce. - - Algorithm: nested loop over (ring_step k, tile_idx t). At k=0 each - rank loads its own block's tiles from HBM; at k > 0 it receives - tiles from its E neighbour. Tiles are forwarded W to the next - ring step. Each tile's partial is folded into the running - (m, ℓ, O) via online-softmax merge. - """ - pe_id = tl.program_id(axis=0) - if P == 1: - # Head-parallel only: PE 0 of each CUBE participates. - if pe_id != 0: - return - T_q_local = T_q - else: - # Intra-CUBE PE-SP (ADR-0060 §5.5 last bullet + §B-item-3): - # query-axis split across P PEs of each CUBE. Output rows are - # disjoint per PE ⇒ no intra-CUBE reduce. Each PE drives its - # own same-lane ring via independent IPCQ channels. - if pe_id >= P: - return - if T_q % P != 0: - raise ValueError( - f"T_q={T_q} must be divisible by P={P} when P > 1; " - "use P=1 for the PE-0-only fallback path." - ) - T_q_local = T_q // P - - S_local = S_kv // C - n_tiles = (S_local + TILE_S_KV - 1) // TILE_S_KV - KV_ROW_BYTES = d_head * 2 # f16 - Q = tl.load(q_ptr, shape=(T_q_local, d_head), dtype="f16") - - # ── Bootstrap: (t=0, k=0) — load my own tile 0, establish (m, ℓ, O) ── - # - # Cannot be folded into the (t, k) loop (kernbench-only limitation): - # - persistent (m, ℓ, O) must live OUTSIDE ``tl.scratch_scope``, - # otherwise scope teardown discards them before the next tile's - # merge can read them; - # - kernbench has no scratch-backed initializer — ``tl.zeros`` / - # ``tl.full`` return addr=0 handles with no backing storage, so - # they cannot be overwritten via ``tl.copy_to`` to seed (-inf, 0, 0); - # - the ring-step ordering also requires k=0 to send its loaded - # tile W before any k>0 iteration can recv from E, so the very - # first (t=0, k=0) step has to run outside the scoped loop body - # where the send is conditional on ``C > 1``. - # Triton port: limitation does not apply (SSA tensors stay live - # across iterations); the bootstrap collapses into the loop. - tile_s = min(TILE_S_KV, S_local) - K_t = tl.load(k_ptr, shape=(d_head, tile_s), dtype="f16") - V_t = tl.load(v_ptr, shape=(tile_s, d_head), dtype="f16") - if C > 1: - tl.send(dir="W", src=K_t) - tl.send(dir="W", src=V_t) - scores = tl.dot(Q, K_t) - m = tl.max(scores, axis=-1) - exp_scores = tl.exp(scores - m) - l = tl.sum(exp_scores, axis=-1) - O = tl.dot(exp_scores, V_t) - - # ── Nested loop: outer tile, inner ring step ── - # Each tile propagates all the way through the ring before the next - # tile starts; IPCQ in-flight depth stays at 1 per direction. - # - # Per outer ``t``: - # k=0 : load my own tile ``t`` from HBM - # k=1..C-1 : receive tile ``t`` of a peer's block from E - # (sent by my E neighbour during their previous k step) - # k dict: } -# ── Bench entry ────────────────────────────────────────────────────── +# ── Sweep entry (called by the umbrella ``milestone_1h_gqa``) ──────── -@bench( - name="milestone-gqa-decode-long-ctx-4cases", - description=( - "Long-context decode 4-cases comparative study on the " - "LLaMA-3.1-70B single-KV-head group (8 cubes × 8 PEs)." - ), -) -def run(torch) -> None: - """Drive the registered decode case panels; write sweep.json. - - Gated by GQA_DECODE_LONG_CTX_4CASES_RUN=1. - """ +def run_sweep(topology: str = "topology.yaml") -> int: + """Drive all decode case panels; write sweep.json. Returns row count.""" _OUTPUT_DIR.mkdir(parents=True, exist_ok=True) - if not os.environ.get("GQA_DECODE_LONG_CTX_4CASES_RUN"): - raise RuntimeError( - "milestone-gqa-decode-long-ctx-4cases needs " - "GQA_DECODE_LONG_CTX_4CASES_RUN=1." - ) - - topology = os.environ.get( - "GQA_DECODE_LONG_CTX_4CASES_TOPOLOGY", "topology.yaml", - ) rows = [_run_panel(panel, topology) for panel in _PANELS] sweep = { "version": 1, @@ -348,6 +329,5 @@ def run(torch) -> None: "rows": rows, } _SWEEP_JSON.write_text(json.dumps(sweep, indent=2)) - print( - f" milestone-gqa-decode-long-ctx-4cases: {len(rows)} rows -> {_SWEEP_JSON}" - ) + print(f" gqa-decode-long-ctx-4cases: {len(rows)} rows -> {_SWEEP_JSON}") + return len(rows) diff --git a/src/kernbench/benches/gqa_helpers/shared/__init__.py b/src/kernbench/benches/gqa_helpers/shared/__init__.py new file mode 100644 index 0000000..ae7dc03 --- /dev/null +++ b/src/kernbench/benches/gqa_helpers/shared/__init__.py @@ -0,0 +1,5 @@ +"""GQA helpers shared across long-ctx and short-ctx variants. + +Includes the panel-metrics helpers (``_gqa_panel_helpers``) and the +context-agnostic decode-opt2 kernel. +""" diff --git a/src/kernbench/benches/_gqa_attention_decode_opt2.py b/src/kernbench/benches/gqa_helpers/shared/_gqa_attention_decode_opt2.py similarity index 100% rename from src/kernbench/benches/_gqa_attention_decode_opt2.py rename to src/kernbench/benches/gqa_helpers/shared/_gqa_attention_decode_opt2.py diff --git a/src/kernbench/benches/gqa_helpers/shared/_gqa_panel_helpers.py b/src/kernbench/benches/gqa_helpers/shared/_gqa_panel_helpers.py new file mode 100644 index 0000000..1878725 --- /dev/null +++ b/src/kernbench/benches/gqa_helpers/shared/_gqa_panel_helpers.py @@ -0,0 +1,40 @@ +"""Shared helpers for the GQA 4-cases milestone benches. + +Used by both ``milestone_gqa_decode_long_ctx_4cases`` and +``milestone_gqa_prefill_long_ctx_4cases``. Underscore-prefixed so the +benches package auto-discovery (which skips ``_*.py``) doesn't try to +audit it for a ``@bench`` decorator. +""" +from __future__ import annotations + +from kernbench.ccl.install import load_ccl_config, resolve_algorithm_config + + +def _ccl_cfg() -> dict: + """Resolve the lrab hierarchical AllReduce CCL config.""" + return resolve_algorithm_config( + load_ccl_config(), name="lrab_hierarchical_allreduce", + ) + + +def _summarize_op_log(op_log) -> dict[str, int]: + """Per-panel op_log counts (gemm, ipcq_copy, dma_read, dma_write).""" + gemm_count = 0 + ipcq_copy_count = 0 + dma_read_count = 0 + dma_write_count = 0 + for r in op_log: + if r.op_kind == "gemm": + gemm_count += 1 + elif r.op_name == "dma_read": + dma_read_count += 1 + elif r.op_name == "dma_write": + dma_write_count += 1 + elif r.op_name == "ipcq_copy": + ipcq_copy_count += 1 + return { + "gemm_count": gemm_count, + "ipcq_copy_count": ipcq_copy_count, + "dma_read_count": dma_read_count, + "dma_write_count": dma_write_count, + } diff --git a/src/kernbench/benches/gqa_helpers/short_ctx/__init__.py b/src/kernbench/benches/gqa_helpers/short_ctx/__init__.py new file mode 100644 index 0000000..9731757 --- /dev/null +++ b/src/kernbench/benches/gqa_helpers/short_ctx/__init__.py @@ -0,0 +1 @@ +"""GQA short-context kernels.""" diff --git a/src/kernbench/benches/_gqa_attention_decode_short.py b/src/kernbench/benches/gqa_helpers/short_ctx/_gqa_attention_decode_short.py similarity index 100% rename from src/kernbench/benches/_gqa_attention_decode_short.py rename to src/kernbench/benches/gqa_helpers/short_ctx/_gqa_attention_decode_short.py diff --git a/src/kernbench/benches/_gqa_attention_prefill_short.py b/src/kernbench/benches/gqa_helpers/short_ctx/_gqa_attention_prefill_short.py similarity index 100% rename from src/kernbench/benches/_gqa_attention_prefill_short.py rename to src/kernbench/benches/gqa_helpers/short_ctx/_gqa_attention_prefill_short.py diff --git a/src/kernbench/benches/milestone_gqa_headline.py b/src/kernbench/benches/milestone_gqa_headline.py deleted file mode 100644 index d213838..0000000 --- a/src/kernbench/benches/milestone_gqa_headline.py +++ /dev/null @@ -1,231 +0,0 @@ -"""milestone-gqa-headline bench: real GQA prefill — single-KV-group target. - -Drives the LLaMA-3.1-70B single-KV-head-group prefill panel through -``_gqa_attention_prefill_long`` (C=8 snake Ring KV + intra-CUBE PE-SP, -all 64 ranks), writing ``op_log_summary`` into ``sweep.json``. -Independent from the existing ``milestone-gqa-llama70b`` validation-scale -bench (which stays on the legacy baseline kernels). - -Decode panels live in ``milestone_gqa_decode_long_ctx_4cases.py`` (the -4-cases long-context decode comparative study). - -Restrictions: - - Single SIP (default ``topology.yaml`` 4×4 cube mesh; 4-SIP - headline deferred per ADR-0060 §B-item-1) - - T_q = S_kv = 1024 (scratch-limited; 32K headline awaits Q-axis - kernel tiling) - - No figure renderers (defer to a separate cycle) - -Panels: - single_kv_group_prefill_gqa_c8_p8: prefill C=8 snake Ring KV + - intra-CUBE PE-SP, T_q=S_kv=1K, - d_head=128 — the LLaMA-3.1-70B - single-KV-group prefill target. - -Gated by ``GQA_HEADLINE_RUN=1``. -""" -from __future__ import annotations - -import json -import os -from pathlib import Path - -from kernbench.benches._gqa_attention_prefill_long import gqa_attention_prefill_long_kernel -from kernbench.benches.registry import bench -from kernbench.ccl.install import load_ccl_config, resolve_algorithm_config -from kernbench.ccl.sfr_config import configure_sfr_intercube_ring -from kernbench.policy.placement.dp import DPPolicy - -_OUTPUT_DIR = Path(__file__).resolve().parent / "1H_milestone_output" / "gqa_headline" -_SWEEP_JSON = _OUTPUT_DIR / "sweep.json" - -# ── Panel configs ──────────────────────────────────────────────────── - -_DTYPE = "f16" -_D_HEAD = 64 -_T_Q_PREFILL = 4 -_S_KV_PREFILL = 16 - -_PANELS = ( - "single_kv_group_prefill_gqa_c8_p8", -) - -# Each entry: (kind, panel-specific params) -_PANEL_DISPATCH: dict[str, tuple[str, dict]] = { - "single_kv_group_prefill_gqa_c8_p8": ("prefill", { - "C": 8, "P": 8, - # T_q = S_kv = 1024 (one-shot long-context prefill, scratch-limited). - # The bootstrap section of the prefill kernel leaves K_t/V_t/scores/ - # exp_scores persistent (outside tl.scratch_scope), inflating the - # baseline. At T_q=2048 the peak hits ~1.05 MB vs 1.0 MB budget; at - # 1024 it fits comfortably. The true LLaMA 32K headline awaits a - # future increment to (a) add Q-axis tiling and (b) move the - # bootstrap into scratch discipline. - "T_q": 1_024, "S_kv": 1_024, - "d_head": 128, - }), -} - - -def _ccl_cfg(): - return resolve_algorithm_config( - load_ccl_config(), name="lrab_hierarchical_allreduce", - ) - - -# ── Per-kind launch helpers ────────────────────────────────────────── - - -def _run_prefill_panel( - ctx, *, panel: str, C: int, S_kv: int, - P: int = 1, - T_q: int = _T_Q_PREFILL, - d_head: int = _D_HEAD, -) -> None: - if C > 1: - mesh_w = int(ctx.spec["sip"]["cube_mesh"]["w"]) - if C <= mesh_w: - configure_sfr_intercube_ring( - ctx.engine, ctx.spec, _ccl_cfg(), ring_size=C, - ) - else: - if C % mesh_w != 0: - raise ValueError( - f"C={C} > mesh_w={mesh_w} requires C divisible by mesh_w" - ) - configure_sfr_intercube_ring( - ctx.engine, ctx.spec, _ccl_cfg(), - submesh_shape=(C // mesh_w, mesh_w), - ) - # Q and O switch to pe="row_wise" when intra-CUBE PE-SP is active - # (ADR-0060 §5.5 + §B-item-3: disjoint query-row split across PEs). - q_pe = "row_wise" if P > 1 else "replicate" - o_pe = "row_wise" if P > 1 else "replicate" - dp_q = DPPolicy(cube="replicate", pe=q_pe, - num_cubes=C, num_pes=P) - dp_kv = DPPolicy(cube="row_wise" if C > 1 else "replicate", - pe="replicate", num_cubes=C, num_pes=P) - dp_o = DPPolicy(cube="row_wise" if C > 1 else "replicate", - pe=o_pe, num_cubes=C, num_pes=P) - q = ctx.zeros((T_q, d_head), - dtype=_DTYPE, dp=dp_q, name=f"{panel}_q") - k = ctx.zeros((S_kv, d_head), - dtype=_DTYPE, dp=dp_kv, name=f"{panel}_k") - v = ctx.zeros((S_kv, d_head), - dtype=_DTYPE, dp=dp_kv, name=f"{panel}_v") - o = ctx.empty((T_q * C, d_head), - dtype=_DTYPE, dp=dp_o, name=f"{panel}_o") - ctx.launch( - panel, gqa_attention_prefill_long_kernel, - q, k, v, o, - T_q, S_kv, d_head, C, P, - _auto_dim_remap=False, - ) - - -def _make_bench_fn(panel: str): - kind, params = _PANEL_DISPATCH[panel] - assert kind == "prefill", ( - f"milestone-gqa-headline only registers prefill panels; got {kind!r}" - ) - - def _bench_fn(ctx): - _run_prefill_panel(ctx, panel=panel, **params) - - return _bench_fn - - -# ── Op-log summary ────────────────────────────────────────────────── - - -def _summarize_op_log(op_log) -> dict[str, int]: - """Per-panel op_log counts (gemm, ipcq_copy, dma_read, dma_write).""" - gemm_count = 0 - ipcq_copy_count = 0 - dma_read_count = 0 - dma_write_count = 0 - for r in op_log: - if r.op_kind == "gemm": - gemm_count += 1 - elif r.op_name == "dma_read": - dma_read_count += 1 - elif r.op_name == "dma_write": - dma_write_count += 1 - elif r.op_name == "ipcq_copy": - ipcq_copy_count += 1 - return { - "gemm_count": gemm_count, - "ipcq_copy_count": ipcq_copy_count, - "dma_read_count": dma_read_count, - "dma_write_count": dma_write_count, - } - - -def _run_panel(panel: str, topology: str) -> dict: - """Run one panel in a fresh engine; return its row dict.""" - from kernbench.runtime_api.bench_runner import run_bench - from kernbench.runtime_api.types import resolve_device - from kernbench.sim_engine.engine import GraphEngine - from kernbench.topology.builder import resolve_topology - - topo = resolve_topology(topology) - result = run_bench( - topology=topo, bench_fn=_make_bench_fn(panel), - device=resolve_device(None), - engine_factory=lambda t, d: GraphEngine( - getattr(t, "topology_obj", t), enable_data=True, - ), - ) - if not result.completion.ok: - raise RuntimeError( - f"milestone-gqa-headline panel {panel!r} failed: " - f"{result.completion}" - ) - kind, params = _PANEL_DISPATCH[panel] - return { - "panel": panel, - "kind": kind, - **params, - "op_log_summary": _summarize_op_log(result.engine.op_log), - } - - -# ── Bench entry ────────────────────────────────────────────────────── - - -@bench( - name="milestone-gqa-headline", - description="Headline GQA prefill milestone — LLaMA-3.1-70B single-KV-group target (C=8, P=8).", -) -def run(torch) -> None: - """Drive the headline prefill panel; write sweep.json. - - Gated by GQA_HEADLINE_RUN=1. - """ - _OUTPUT_DIR.mkdir(parents=True, exist_ok=True) - if not os.environ.get("GQA_HEADLINE_RUN"): - raise RuntimeError( - "milestone-gqa-headline needs GQA_HEADLINE_RUN=1." - ) - - topology = os.environ.get("GQA_HEADLINE_TOPOLOGY", "topology.yaml") - rows = [_run_panel(panel, topology) for panel in _PANELS] - sweep = { - "version": 1, - "panels": list(_PANELS), - "config": { - "T_q_prefill": _T_Q_PREFILL, - "S_kv_prefill": _S_KV_PREFILL, - "d_head": _D_HEAD, - }, - "rows": rows, - } - _SWEEP_JSON.write_text(json.dumps(sweep, indent=2)) - print(f" milestone-gqa-headline: {len(rows)} rows -> {_SWEEP_JSON}") - - # Sentinel tensor (ADR-0045 D4 / ADR-0054 D2 carve-out). - torch.zeros( - (1, 1), dtype="f16", - dp=DPPolicy(cube="row_wise", pe="replicate", num_cubes=1, num_pes=1), - name="milestone_gqa_headline_sentinel", - ) diff --git a/src/kernbench/benches/registry.py b/src/kernbench/benches/registry.py index 58d8af8..02e3c09 100644 --- a/src/kernbench/benches/registry.py +++ b/src/kernbench/benches/registry.py @@ -99,7 +99,11 @@ def _audit_modules(imported: list[str], registered: set[str]) -> None: def _eager_import_and_audit(pkg_path: list[str], pkg_name: str) -> None: imported: list[str] = [] for m in iter_modules(pkg_path): - if m.name == "registry" or m.name.startswith("_"): + # Skip helper-name modules, the registry itself, and any + # subpackages (audit policy applies to ``.py`` bench files + # only — subpackages are explicitly imported by the benches + # that need them). + if m.ispkg or m.name == "registry" or m.name.startswith("_"): continue mod = import_module(f"{pkg_name}.{m.name}") imported.append(mod.__name__) diff --git a/tests/attention/test_gqa_decode_opt2.py b/tests/attention/test_gqa_decode_opt2.py index ecfbc2d..b7efee4 100644 --- a/tests/attention/test_gqa_decode_opt2.py +++ b/tests/attention/test_gqa_decode_opt2.py @@ -151,7 +151,7 @@ def test_k_before_v_in_opt2_plan(): def test_opt2_bench_completes_oplog_mode(): from pathlib import Path - from kernbench.benches._gqa_attention_decode_opt2 import ( # noqa: F401 + from kernbench.benches.gqa_helpers.shared._gqa_attention_decode_opt2 import ( # noqa: F401 gqa_attention_decode_opt2_kernel, ) from kernbench.policy.placement.dp import DPPolicy diff --git a/tests/attention/test_gqa_prefill.py b/tests/attention/test_gqa_prefill.py deleted file mode 100644 index 090b479..0000000 --- a/tests/attention/test_gqa_prefill.py +++ /dev/null @@ -1,118 +0,0 @@ -"""Phase 1 spec test for P6a GQA prefill kernel (head-parallel, C=1 baseline). - -P6a introduces ``_gqa_attention_prefill_long.py`` with the head-parallel structure (one -Q head per CUBE, per-CUBE distributed output, no reduce). C=1 is the -degenerate case — no Ring KV, no IPCQ traffic. Validates kernel -structure and T_q > 1 attention. - -P6b (deferred) adds the Ring KV rotation for C > 1, which needs either -a new SFR install function (intra_* + wrapped E/W at CUBE level) or a -topology-specific config — separate design call. - -The prefill kernel differs from decode (P1a/P2a/P2b) in three ways -(ADR-0060 §5.5 / TL;DR): - 1. Q has T_q > 1 rows (not just decode's single timestep). - 2. Head-parallel placement: each CUBE owns ONE Q head — no M-fold. - 3. Each CUBE writes its own head's output — NO reduce. - -Phase 1 (this commit): tests only — production code lands in Phase 2. -""" -from __future__ import annotations - -from pathlib import Path - -from kernbench.benches._gqa_attention_prefill_long import gqa_attention_prefill_long_kernel # noqa: F401 -from kernbench.policy.placement.dp import DPPolicy -from kernbench.runtime_api.bench_runner import run_bench -from kernbench.runtime_api.types import resolve_device -from kernbench.sim_engine.engine import GraphEngine -from kernbench.topology.builder import resolve_topology - -TOPOLOGY_DEFAULT = Path(__file__).resolve().parents[2] / "topology.yaml" - -D_HEAD = 64 -DTYPE = "f16" - - -def _engine_factory(t, d): - return GraphEngine(getattr(t, "topology_obj", t), enable_data=True) - - -def _run_prefill(*, T_q: int, S_kv: int, C: int = 1): - """C=1 head-parallel prefill: single CUBE owns the one head + full KV.""" - topo = resolve_topology(str(TOPOLOGY_DEFAULT)) - - def _bench_fn(ctx): - dp = DPPolicy(cube="replicate", pe="replicate", - num_cubes=C, num_pes=1) - # Q: (T_q, d_head) — one head per CUBE (head-parallel; for C=1 - # only one head total). 2D layout matches what the kernel loads. - # P6b will use a 3D (h_q, T_q, d_head) Q with cube_row_wise - # sharding so each CUBE owns its head. - q = ctx.zeros((T_q, D_HEAD), dtype=DTYPE, dp=dp, - name=f"q_t{T_q}_c{C}") - # K, V: full local for C=1 (no ring). Kernel loads K as - # (d_head, S_kv) via byte-conserving reshape. - k = ctx.zeros((S_kv, D_HEAD), dtype=DTYPE, dp=dp, - name=f"k_t{T_q}_c{C}") - v = ctx.zeros((S_kv, D_HEAD), dtype=DTYPE, dp=dp, - name=f"v_t{T_q}_c{C}") - # O: (T_q, d_head) — per-CUBE distributed output. - o = ctx.empty((T_q, D_HEAD), dtype=DTYPE, dp=dp, - name=f"o_t{T_q}_c{C}") - ctx.launch( - f"gqa_prefill_p6a_t{T_q}_s{S_kv}_c{C}", - gqa_attention_prefill_long_kernel, - q, k, v, o, - T_q, S_kv, D_HEAD, C, - _auto_dim_remap=False, - ) - - return run_bench( - topology=topo, bench_fn=_bench_fn, - device=resolve_device(None), - engine_factory=_engine_factory, - ) - - -def _count(op_log, name: str) -> int: - return sum(1 for r in op_log if r.op_name == name) - - -def test_prefill_c_one_t_q_one_completes(): - """C=1, T_q=1: smallest workload (decode-like).""" - result = _run_prefill(T_q=1, S_kv=16, C=1) - assert result.completion.ok, ( - f"prefill C=1 T_q=1 failed: {result.completion}" - ) - - -def test_prefill_c_one_t_q_four_completes(): - """C=1, T_q=4: real prefill (Q has multiple rows) — distinguishes - prefill from decode (T_q=1).""" - result = _run_prefill(T_q=4, S_kv=16, C=1) - assert result.completion.ok, ( - f"prefill C=1 T_q=4 failed: {result.completion}" - ) - - -def test_prefill_c_one_no_ipcq_traffic(): - """C=1: no ring step, no IPCQ traffic.""" - result = _run_prefill(T_q=4, S_kv=16, C=1) - assert result.completion.ok - n_copy = _count(result.engine.op_log, "ipcq_copy") - assert n_copy == 0, ( - f"C=1 must have no IPCQ traffic (no ring); got {n_copy}" - ) - - -def test_prefill_c_one_one_dma_write(): - """C=1, one head: exactly one dma_write (per-CUBE distributed output; - no reduce). For C > 1 in P6b this becomes dma_write_count == C.""" - result = _run_prefill(T_q=4, S_kv=16, C=1) - assert result.completion.ok - n_writes = _count(result.engine.op_log, "dma_write") - assert n_writes == 1, ( - f"C=1 prefill: expected 1 dma_write (one head per cube); " - f"got {n_writes}" - ) diff --git a/tests/attention/test_gqa_prefill_long_c8_snake.py b/tests/attention/test_gqa_prefill_long_c8_snake.py deleted file mode 100644 index 58ca62e..0000000 --- a/tests/attention/test_gqa_prefill_long_c8_snake.py +++ /dev/null @@ -1,154 +0,0 @@ -"""Tests for prefill_long Ring KV at C=8 over a 2×4 snake sub-mesh. - -ADR-0060 §5.5 design target for the single-KV-group LLaMA-3.1-70B -configuration: ``C = G = 8`` (one Q head per CUBE), Ring KV rotates -the 8 KV slices around all 8 CUBEs. - -Increment 1 wired ``configure_sfr_intercube_ring(submesh_shape=(2, 4))`` -to map the kernel's logical E/W to a snake/serpentine 1-hop path -through the top 2×4 sub-mesh of the 4×4 SIP CUBE mesh. The kernel -itself uses only logical ``dir="W"`` / ``dir="E"`` — the snake is -transparent at kernel level. - -This file verifies the assembly works end-to-end at C=8: - T1 kernel completes (no scratch overflow, no deadlock) - T2 per-CUBE head-parallel output (8 dma_writes, one per CUBE) - T3 tile-granular ring traffic at C=8 follows the same - ``(C-1)·n_tiles·2·C`` formula as the existing tile-ring tests -""" -from __future__ import annotations - -from pathlib import Path - -from kernbench.benches._gqa_attention_prefill_long import gqa_attention_prefill_long_kernel # noqa: F401 -from kernbench.ccl.install import load_ccl_config, resolve_algorithm_config -from kernbench.ccl.sfr_config import configure_sfr_intercube_ring -from kernbench.policy.placement.dp import DPPolicy -from kernbench.runtime_api.bench_runner import run_bench -from kernbench.runtime_api.types import resolve_device -from kernbench.sim_engine.engine import GraphEngine -from kernbench.topology.builder import resolve_topology - -TOPOLOGY_DEFAULT = Path(__file__).resolve().parents[2] / "topology.yaml" - -D_HEAD = 64 -DTYPE = "f16" - - -def _ccl_cfg(): - return resolve_algorithm_config( - load_ccl_config(), name="lrab_hierarchical_allreduce", - ) - - -def _engine_factory(t, d): - return GraphEngine(getattr(t, "topology_obj", t), enable_data=True) - - -def _count(op_log, name: str) -> int: - return sum(1 for r in op_log if r.op_name == name) - - -def _run_prefill_c8_snake(*, T_q: int, S_kv: int): - """Head-parallel prefill at C=8 with snake-mapped Ring KV over the - top 2×4 sub-mesh of the 4×4 SIP CUBE mesh.""" - topo = resolve_topology(str(TOPOLOGY_DEFAULT)) - C = 8 - - def _bench_fn(ctx): - configure_sfr_intercube_ring( - ctx.engine, ctx.spec, _ccl_cfg(), - submesh_shape=(2, 4), - ) - dp_q = DPPolicy(cube="replicate", pe="replicate", - num_cubes=C, num_pes=1) - dp_kv = DPPolicy(cube="row_wise", pe="replicate", - num_cubes=C, num_pes=1) - dp_o = DPPolicy(cube="row_wise", pe="replicate", - num_cubes=C, num_pes=1) - q = ctx.zeros((T_q, D_HEAD), dtype=DTYPE, dp=dp_q, - name=f"q_c8_snake_t{T_q}_s{S_kv}") - k = ctx.zeros((S_kv, D_HEAD), dtype=DTYPE, dp=dp_kv, - name=f"k_c8_snake_t{T_q}_s{S_kv}") - v = ctx.zeros((S_kv, D_HEAD), dtype=DTYPE, dp=dp_kv, - name=f"v_c8_snake_t{T_q}_s{S_kv}") - o = ctx.empty((T_q * C, D_HEAD), dtype=DTYPE, dp=dp_o, - name=f"o_c8_snake_t{T_q}_s{S_kv}") - ctx.launch( - f"gqa_prefill_long_c8_snake_t{T_q}_s{S_kv}", - gqa_attention_prefill_long_kernel, - q, k, v, o, - T_q, S_kv, D_HEAD, C, - _auto_dim_remap=False, - ) - - return run_bench( - topology=topo, bench_fn=_bench_fn, - device=resolve_device(None), - engine_factory=_engine_factory, - ) - - -# ── T1: kernel completes end-to-end at C=8 ─────────────────────────── - - -def test_prefill_long_c8_snake_completes(): - """ADR-0060 §5.5 design target: prefill Ring KV at C=G=8 over the - 2×4 snake sub-mesh. With zero inputs, output is trivially zero; - this test guards that the kernel reaches completion without - deadlock, scratch overflow, or routing failure. - - Configuration: T_q=4, S_kv=8192 → S_local=1024, n_tiles=1 - (TILE_S_KV=1024). Bounded scratch. - """ - result = _run_prefill_c8_snake(T_q=4, S_kv=8192) - assert result.completion.ok, ( - f"prefill at C=8 with snake ring must complete; " - f"got {result.completion}" - ) - - -# ── T2: 8 dma_writes (head-parallel, one head per CUBE) ────────────── - - -def test_prefill_long_c8_snake_distributed_output_count(): - """Head-parallel: each CUBE owns one Q head and writes its own - head's output (T_q, d_head) rows. No inter-CUBE reduce on the - output side — so ``dma_write_count == C == 8``. - - This is the C=8 analogue of - ``test_prefill_long_tile_ring_dma_write_count`` at C=4. - """ - result = _run_prefill_c8_snake(T_q=4, S_kv=8192) - assert result.completion.ok - n_writes = _count(result.engine.op_log, "dma_write") - assert n_writes == 8, ( - f"head-parallel C=8: expected 8 dma_writes (one per CUBE); " - f"got {n_writes}" - ) - - -# ── T3: tile-granular ring ipcq_copy count scales as (C-1)·n_tiles·2·C - - -def test_prefill_long_c8_snake_tile_ipcq_count(): - """ADR-0060 §5.5.1 tile-granular ring: per-CUBE send count is - ``2·n_tiles·(C-1)`` (K + V per tile per ring step). Aggregated - across C CUBEs, total ipcq_copy = ``(C-1)·n_tiles·2·C``. - - Configuration: T_q=4, S_kv=8192, C=8 → S_local=1024, n_tiles=1 - (TILE_S_KV=1024). Expected total = ``7·1·2·8 = 112``. - - Verifies that the snake-mapped 1-hop physical links carry the - same logical ring traffic as the existing C=4 1D-row ring tests. - """ - C = 8 - n_tiles = 1 # S_local=1024 / TILE_S_KV=1024 - result = _run_prefill_c8_snake(T_q=4, S_kv=8192) - assert result.completion.ok - n_copy = _count(result.engine.op_log, "ipcq_copy") - expected = (C - 1) * n_tiles * 2 * C - assert n_copy == expected, ( - f"tile-granular ring at C=8: expected {expected} ipcq_copy " - f"((C-1)·n_tiles·2·C = {C - 1}·{n_tiles}·2·{C}); got {n_copy}" - ) diff --git a/tests/attention/test_gqa_prefill_long_pe_sp.py b/tests/attention/test_gqa_prefill_long_pe_sp.py deleted file mode 100644 index 8c8b775..0000000 --- a/tests/attention/test_gqa_prefill_long_pe_sp.py +++ /dev/null @@ -1,233 +0,0 @@ -"""Tests for intra-CUBE PE-SP in prefill_long (query-axis split). - -ADR-0060 §5.5 last bullet + §B-item-3: split the head's query rows -``[T_q, d_head]`` across the ``P`` PEs of each CUBE so all P PEs work -in parallel. Output rows are disjoint across PEs ⇒ no intra-CUBE -reduce needed. - -Same-lane SFR wiring: PE ``i`` in CUBE A has its own E/W ring link to -PE ``i`` in CUBE B (the snake's prev/next). All P PEs of a CUBE see -the same K, V (HBM-resident, ``pe="replicate"``) ⇒ P parallel rings -run in lockstep, each PE rotating its own K/V copies via its own IPCQ -channels. - -Activation contract: - - ``P == 1`` (default; omitted from launch args) → existing - PE-0-only behavior. Byte-for-byte unchanged. - - ``P > 1`` → all P PEs active; each handles ``T_q // P`` query - rows. Requires ``T_q % P == 0`` (degenerate T_q < P is rejected - — caller must use ``P=1`` for that workload, ADR-0060 §B-item-3). - -Phase 1: tests only — production code lands in Phase 2. -T1, T2, T3, T5 fail today (TypeError: kernel signature has no P). -T4 passes today as the backward-compat anchor. -""" -from __future__ import annotations - -from pathlib import Path - -import pytest - -from kernbench.benches._gqa_attention_prefill_long import gqa_attention_prefill_long_kernel # noqa: F401 -from kernbench.ccl.install import load_ccl_config, resolve_algorithm_config -from kernbench.ccl.sfr_config import configure_sfr_intercube_ring -from kernbench.policy.placement.dp import DPPolicy -from kernbench.runtime_api.bench_runner import run_bench -from kernbench.runtime_api.types import resolve_device -from kernbench.sim_engine.engine import GraphEngine -from kernbench.topology.builder import resolve_topology - -TOPOLOGY_DEFAULT = Path(__file__).resolve().parents[2] / "topology.yaml" - -D_HEAD = 64 -DTYPE = "f16" - - -def _ccl_cfg(): - return resolve_algorithm_config( - load_ccl_config(), name="lrab_hierarchical_allreduce", - ) - - -def _engine_factory(t, d): - return GraphEngine(getattr(t, "topology_obj", t), enable_data=True) - - -def _count(op_log, name: str) -> int: - return sum(1 for r in op_log if r.op_name == name) - - -def _run_prefill_pe_sp( - *, T_q: int, S_kv: int, C: int, P: int | None, - snake: bool, -): - """Drive a prefill_long launch with optional intra-CUBE PE-SP. - - ``P=None`` → omit P from the launch args (exercises the kernel's - default behaviour). - ``snake=True`` → install the 2×4 snake ring SFR (Increment 1); - else install the 1D-row ring at ``ring_size=C``. - """ - topo = resolve_topology(str(TOPOLOGY_DEFAULT)) - - def _bench_fn(ctx): - if snake: - configure_sfr_intercube_ring( - ctx.engine, ctx.spec, _ccl_cfg(), - submesh_shape=(2, 4), - ) - else: - configure_sfr_intercube_ring( - ctx.engine, ctx.spec, _ccl_cfg(), - ring_size=C, - ) - num_pes = P if (P is not None and P > 1) else 1 - # When PE-SP is active, Q and O are split row-wise across PEs; - # K, V remain replicated within a CUBE. - q_pe = "row_wise" if num_pes > 1 else "replicate" - o_pe = "row_wise" if num_pes > 1 else "replicate" - dp_q = DPPolicy(cube="replicate", pe=q_pe, - num_cubes=C, num_pes=num_pes) - dp_kv = DPPolicy(cube="row_wise" if C > 1 else "replicate", - pe="replicate", num_cubes=C, num_pes=num_pes) - dp_o = DPPolicy(cube="row_wise" if C > 1 else "replicate", - pe=o_pe, num_cubes=C, num_pes=num_pes) - suffix = f"c{C}_p{P}_t{T_q}_s{S_kv}" - q = ctx.zeros((T_q, D_HEAD), dtype=DTYPE, dp=dp_q, - name=f"q_pe_sp_{suffix}") - k = ctx.zeros((S_kv, D_HEAD), dtype=DTYPE, dp=dp_kv, - name=f"k_pe_sp_{suffix}") - v = ctx.zeros((S_kv, D_HEAD), dtype=DTYPE, dp=dp_kv, - name=f"v_pe_sp_{suffix}") - o = ctx.empty((T_q * C, D_HEAD), dtype=DTYPE, dp=dp_o, - name=f"o_pe_sp_{suffix}") - launch_args = [q, k, v, o, T_q, S_kv, D_HEAD, C] - if P is not None: - launch_args.append(P) - ctx.launch( - f"gqa_prefill_long_pe_sp_{suffix}", - gqa_attention_prefill_long_kernel, - *launch_args, - _auto_dim_remap=False, - ) - - return run_bench( - topology=topo, bench_fn=_bench_fn, - device=resolve_device(None), - engine_factory=_engine_factory, - ) - - -# ── T1: C=8 P=8 completes end-to-end ───────────────────────────────── - - -def test_prefill_c8_p8_pe_sp_completes(): - """ADR-0060 §5.5 + §B-item-3 design target: 64 ranks active - (8 CUBEs × 8 PEs), query-axis split across PEs, snake-mapped - Ring KV. Guards no scratch overflow, no deadlock across the - P parallel same-lane rings. - """ - result = _run_prefill_pe_sp( - T_q=8, S_kv=8192, C=8, P=8, snake=True, - ) - assert result.completion.ok, ( - f"prefill at C=8, P=8 (PE-SP) must complete; " - f"got {result.completion}" - ) - - -# ── T2: 64 dma_writes (one per PE, disjoint query rows) ────────────── - - -def test_prefill_c8_p8_pe_sp_64_dma_writes(): - """With query-axis split, each PE owns ``T_q/P = 1`` query row - and stores its own ``(1, d_head)`` slice. Across all 64 ranks - (8 CUBEs × 8 PEs), ``dma_write_count == 64``. - - This is the structural signal that PE-SP is actually wired: - PE-0-only would give 8 dma_writes (one per CUBE). - """ - result = _run_prefill_pe_sp( - T_q=8, S_kv=8192, C=8, P=8, snake=True, - ) - assert result.completion.ok - n_writes = _count(result.engine.op_log, "dma_write") - assert n_writes == 64, ( - f"PE-SP at C=8, P=8: expected 64 dma_writes " - f"(8 cubes × 8 PEs, each writes its T_q/P=1 query row); " - f"got {n_writes}" - ) - - -# ── T3: P parallel rings — ipcq_copy scales by P ───────────────────── - - -def test_prefill_c8_p8_pe_sp_ring_ipcq_count(): - """Per-PE same-lane rings: each PE runs its own ring traffic via - its own IPCQ channels. Total inter-CUBE ipcq_copy: - ``(C-1) · n_tiles · 2 · C · P`` (= existing C=8 formula × P). - - Configuration: T_q=8, S_kv=8192, C=8 → S_local=1024, n_tiles=1 - (TILE_S_KV=1024). Expected = ``7·1·2·8·8`` = **896**. - """ - C = 8 - P = 8 - n_tiles = 1 - result = _run_prefill_pe_sp( - T_q=8, S_kv=8192, C=C, P=P, snake=True, - ) - assert result.completion.ok - n_copy = _count(result.engine.op_log, "ipcq_copy") - expected = (C - 1) * n_tiles * 2 * C * P - assert n_copy == expected, ( - f"PE-SP parallel rings at C=8, P=8: expected {expected} " - f"ipcq_copy ((C-1)·n_tiles·2·C·P = " - f"{C - 1}·{n_tiles}·2·{C}·{P}); got {n_copy}" - ) - - -# ── T4: backward-compat — P omitted → existing PE-0-only behaviour ─── - - -def test_prefill_p_default_1_backward_compat(): - """When ``P`` is omitted from the launch args, the kernel must - behave identically to today: only PE 0 of each CUBE participates; - one head per CUBE; one dma_write per CUBE. - - At C=4 this gives 4 dma_writes (matches the existing - ``test_prefill_long_tile_ring_dma_write_count``). Phase 2 must - not regress this path. - - Passes today AND after Phase 2. - """ - C = 4 - result = _run_prefill_pe_sp( - T_q=4, S_kv=8192, C=C, P=None, snake=False, - ) - assert result.completion.ok, ( - f"prefill at C=4 with default P (PE-0-only) must complete; " - f"got {result.completion}" - ) - n_writes = _count(result.engine.op_log, "dma_write") - assert n_writes == C, ( - f"default P=1 (PE-0-only): expected {C} dma_writes " - f"(one per CUBE); got {n_writes}" - ) - - -# ── T5: validation — T_q must be divisible by P when P > 1 ─────────── - - -def test_prefill_pe_sp_rejects_non_divisible_t_q(): - """ADR-0060 §B-item-3 fallback (KV-block split + intra-CUBE - reduce for T_q < P) is deferred. Callers must request ``P=1`` for - workloads where T_q < P or T_q % P != 0. - - C=4, P=8, T_q=4 violates ``T_q % P == 0`` (and also T_q < P). - The kernel must raise ValueError with a clear error message - rather than silently producing wrong results. - """ - with pytest.raises((ValueError, AssertionError), match=r"T_q"): - _run_prefill_pe_sp( - T_q=4, S_kv=8192, C=4, P=8, snake=False, - ) diff --git a/tests/attention/test_gqa_prefill_long_tile_ring.py b/tests/attention/test_gqa_prefill_long_tile_ring.py deleted file mode 100644 index 93833a6..0000000 --- a/tests/attention/test_gqa_prefill_long_tile_ring.py +++ /dev/null @@ -1,162 +0,0 @@ -"""Phase 1 spec test for P3c: tile-granular Ring KV in prefill_long -(ADR-0060 §5.5 amendment + ADR-0063 §A.2). - -Today's prefill_long ring sends and receives full ``(d_head, S_local)`` -KV slices per step. Step 0's local-attention intermediates also live -outside any ``scratch_scope`` (they're loaded as full-slice ``Kc``, -``Vc`` and feed ``scores``, ``exp_scores`` as persistent allocations). -At larger ``S_local`` both the step-0 leak and the ring step's -in-scope intermediates grow linearly with ``S_local``, and at -``S_local = 32K`` (S_kv=128K, C=4, T_q=4) the peak overflows the -1 MiB pool. - -P3c converts the ring to **tile-granular**: a nested loop -``for k in range(C): for t in range(n_tiles): ...`` where each -iteration sends/recvs one ``(d_head, TILE_S_KV)`` tile (and its V -counterpart). The persistent state shrinks to ``(m, ℓ, O)`` only -(~1 KB); per-tile in-scope scratch is bounded by -``TILE_S_KV`` regardless of ``S_local``. Ceiling lifted. - -Trade-off: the per-CUBE send count grows from ``2·(C-1)`` to -``2·n_tiles·(C-1)``. At ``n_tiles=1`` (small ``S_local``) the count is -unchanged, so the existing ``test_prefill_ring_c_*`` tests at -``S_kv ∈ {16, 32}`` still pass. - -Phase 1 (this commit): tests only — production code lands in Phase 2. -T1 fails today with a ``TLContext scratch overflow``; T2 fails today -with the slice-granular ipcq_copy count; T5 passes today and is a -regression guard for the per-CUBE output write count. -""" -from __future__ import annotations - -from pathlib import Path - -from kernbench.benches._gqa_attention_prefill_long import gqa_attention_prefill_long_kernel # noqa: F401 -from kernbench.ccl.install import load_ccl_config, resolve_algorithm_config -from kernbench.ccl.sfr_config import configure_sfr_intercube_ring -from kernbench.policy.placement.dp import DPPolicy -from kernbench.runtime_api.bench_runner import run_bench -from kernbench.runtime_api.types import resolve_device -from kernbench.sim_engine.engine import GraphEngine -from kernbench.topology.builder import resolve_topology - -TOPOLOGY_DEFAULT = Path(__file__).resolve().parents[2] / "topology.yaml" - -D_HEAD = 64 -DTYPE = "f16" - - -def _ccl_cfg(): - return resolve_algorithm_config( - load_ccl_config(), name="lrab_hierarchical_allreduce", - ) - - -def _engine_factory(t, d): - return GraphEngine(getattr(t, "topology_obj", t), enable_data=True) - - -def _count(op_log, name: str) -> int: - return sum(1 for r in op_log if r.op_name == name) - - -def _run_prefill_ring(*, T_q: int, S_kv: int, C: int): - """Head-parallel prefill with Ring KV across C CUBEs.""" - topo = resolve_topology(str(TOPOLOGY_DEFAULT)) - - def _bench_fn(ctx): - configure_sfr_intercube_ring( - ctx.engine, ctx.spec, _ccl_cfg(), ring_size=C, - ) - dp_q = DPPolicy(cube="replicate", pe="replicate", - num_cubes=C, num_pes=1) - dp_kv = DPPolicy(cube="row_wise" if C > 1 else "replicate", - pe="replicate", num_cubes=C, num_pes=1) - dp_o = DPPolicy(cube="row_wise" if C > 1 else "replicate", - pe="replicate", num_cubes=C, num_pes=1) - q = ctx.zeros((T_q, D_HEAD), dtype=DTYPE, dp=dp_q, - name=f"q_tlr_t{T_q}_c{C}_s{S_kv}") - k = ctx.zeros((S_kv, D_HEAD), dtype=DTYPE, dp=dp_kv, - name=f"k_tlr_t{T_q}_c{C}_s{S_kv}") - v = ctx.zeros((S_kv, D_HEAD), dtype=DTYPE, dp=dp_kv, - name=f"v_tlr_t{T_q}_c{C}_s{S_kv}") - o = ctx.empty((T_q * C, D_HEAD), dtype=DTYPE, dp=dp_o, - name=f"o_tlr_t{T_q}_c{C}_s{S_kv}") - ctx.launch( - f"gqa_prefill_long_tile_ring_t{T_q}_c{C}_s{S_kv}", - gqa_attention_prefill_long_kernel, - q, k, v, o, - T_q, S_kv, D_HEAD, C, - _auto_dim_remap=False, - ) - - return run_bench( - topology=topo, bench_fn=_bench_fn, - device=resolve_device(None), - engine_factory=_engine_factory, - ) - - -# ── T1: 128K ceiling lift ──────────────────────────────────────────── - - -def test_prefill_long_context_128k_completes(): - """ADR-0063 §A.2 headline ceiling lift. At S_kv=128K with C=4, - S_local=32K. Today: step-0 score-stack (~768 KB persistent) + ring - scope (~768 KB) → 1.5 MB peak → TLContext scratch overflow. After - P3c the persistent state is just ``(m, ℓ, O)`` (≈ 1 KB) and the - per-tile in-scope scratch is bounded by TILE_S_KV. - """ - result = _run_prefill_ring(T_q=4, S_kv=131_072, C=4) - assert result.completion.ok, ( - f"prefill_long at S_kv=128K must complete after tile-granular " - f"ring lands; got {result.completion}" - ) - - -# ── T2: tile-granular ipcq_copy count ──────────────────────────────── - - -def test_prefill_long_tile_granular_ipcq_count(): - """ADR-0060 §5.5 amendment: with tile-granular sends, the per-CUBE - send count grows from ``2·(C-1)`` to ``2·n_tiles·(C-1)``. Aggregated - across all C CUBEs the total ipcq_copy becomes - ``2·n_tiles·(C-1)·C``. - - Config: T_q=4, S_kv=4096, C=2 → S_local=2048, n_tiles=2 (with - TILE_S_KV=1024). Today: slice-granular total = - ``(C-1)·2·C = 4``. After P3c: tile-granular total = - ``(C-1)·n_tiles·2·C = 8``. - """ - C = 2 - n_tiles = 2 # S_local=2048 / TILE_S_KV=1024 - result = _run_prefill_ring(T_q=4, S_kv=4096, C=C) - assert result.completion.ok, ( - f"prefill_long multi-tile ring must complete; got {result.completion}" - ) - n_copy = _count(result.engine.op_log, "ipcq_copy") - expected = (C - 1) * n_tiles * 2 * C - assert n_copy == expected, ( - f"tile-granular ring: expected {expected} ipcq_copy " - f"((C-1)·n_tiles·2·C = {C - 1}·{n_tiles}·2·{C}); got {n_copy}" - ) - - -# ── T5: per-CUBE distributed output unchanged (regression guard) ───── - - -def test_prefill_long_tile_ring_dma_write_count(): - """ADR-0060 §5.5: per-CUBE distributed output must hold under the - tile-granular ring rewrite. Each CUBE still writes its own head's - output (no inter-CUBE reduce); dma_write_count == C. - - Today passes; must continue to pass after P3c. - """ - C = 4 - result = _run_prefill_ring(T_q=4, S_kv=4096, C=C) - assert result.completion.ok - n_writes = _count(result.engine.op_log, "dma_write") - assert n_writes == C, ( - f"per-CUBE distributed output: expected {C} dma_writes (one per " - f"CUBE); got {n_writes}" - ) diff --git a/tests/attention/test_gqa_prefill_ring.py b/tests/attention/test_gqa_prefill_ring.py deleted file mode 100644 index 1ce602c..0000000 --- a/tests/attention/test_gqa_prefill_ring.py +++ /dev/null @@ -1,161 +0,0 @@ -"""Phase 1 spec test for P6b GQA prefill Ring KV (head-parallel, C>1). - -P6b adds the Ring KV rotation (ADR-0060 §5.5) to the prefill kernel. -Each CUBE owns one Q head + its KV slice; over C ring steps the KV -blocks rotate around the C CUBEs so every CUBE sees every block. The -running ``(m, ℓ, O)`` is folded inside the loop. No reduce — each CUBE -writes its own head's output. - -Requires a new SFR install ``configure_sfr_intercube_ring(ring_size=C)`` -that wires: - - ``intra_*`` : 2×4 PE grid within a CUBE (same as multisip) - - ``E/W`` : 1D ring of CUBEs 0..ring_size-1 WITH WRAP - (symmetric to ``configure_sfr_intracube_pe_ring``, - applied at CUBE level) - - ``global_*`` : SIP-level (same as multisip) - -The kernel ring body: - for step in range(1, C): - tl.send(dir="W", src=Kc) - tl.send(dir="W", src=Vc) - Kc = tl.recv(dir="E", ...) - Vc = tl.recv(dir="E", ...) - ... local partial + online-softmax merge into (m, ℓ, O) ... - -Per CUBE per step: 2 sends (K, V) → 2 ``ipcq_copy``. Across all CUBEs: -``(C-1) * 2 * C`` total ipcq_copy. - -Restriction in P6b first cut: ``C ∈ {1, 2, 4}`` (single row of the 4×4 -cube mesh). C=8 ring would span rows — follow-on. - -Phase 1 (this commit): tests only — production code lands in Phase 2. -""" -from __future__ import annotations - -from pathlib import Path - -from kernbench.benches._gqa_attention_prefill_long import gqa_attention_prefill_long_kernel # noqa: F401 -from kernbench.ccl.install import load_ccl_config, resolve_algorithm_config -from kernbench.ccl.sfr_config import configure_sfr_intercube_ring # noqa: F401 (Phase 2) -from kernbench.policy.placement.dp import DPPolicy -from kernbench.runtime_api.bench_runner import run_bench -from kernbench.runtime_api.types import resolve_device -from kernbench.sim_engine.engine import GraphEngine -from kernbench.topology.builder import resolve_topology - -TOPOLOGY_DEFAULT = Path(__file__).resolve().parents[2] / "topology.yaml" - -D_HEAD = 64 -DTYPE = "f16" - - -def _ccl_cfg(): - return resolve_algorithm_config( - load_ccl_config(), name="lrab_hierarchical_allreduce", - ) - - -def _engine_factory(t, d): - return GraphEngine(getattr(t, "topology_obj", t), enable_data=True) - - -def _run_prefill_ring(*, T_q: int, S_kv: int, C: int): - """Head-parallel prefill with Ring KV across C CUBEs.""" - topo = resolve_topology(str(TOPOLOGY_DEFAULT)) - - def _bench_fn(ctx): - # P6b: new SFR install with cube-level ring wrap. - configure_sfr_intercube_ring( - ctx.engine, ctx.spec, _ccl_cfg(), ring_size=C, - ) - # Q replicated on every CUBE (zeros for testing; per-CUBE head - # indexing is implicit). KV sequence-sharded by CUBE. O - # distributed — each CUBE writes its slice of (T_q*C, d_head). - dp_q = DPPolicy(cube="replicate", pe="replicate", - num_cubes=C, num_pes=1) - dp_kv = DPPolicy(cube="row_wise", pe="replicate", - num_cubes=C, num_pes=1) - dp_o = DPPolicy(cube="row_wise", pe="replicate", - num_cubes=C, num_pes=1) - q = ctx.zeros((T_q, D_HEAD), dtype=DTYPE, dp=dp_q, - name=f"q_t{T_q}_c{C}_ring") - k = ctx.zeros((S_kv, D_HEAD), dtype=DTYPE, dp=dp_kv, - name=f"k_t{T_q}_c{C}_ring") - v = ctx.zeros((S_kv, D_HEAD), dtype=DTYPE, dp=dp_kv, - name=f"v_t{T_q}_c{C}_ring") - # O: (T_q * C, d_head), each CUBE writes (T_q, d_head) at its - # slice. dma_write_count = C (per-CUBE distributed output). - o = ctx.empty((T_q * C, D_HEAD), dtype=DTYPE, dp=dp_o, - name=f"o_t{T_q}_c{C}_ring") - ctx.launch( - f"gqa_prefill_ring_t{T_q}_s{S_kv}_c{C}", - gqa_attention_prefill_long_kernel, - q, k, v, o, - T_q, S_kv, D_HEAD, C, - _auto_dim_remap=False, - ) - - return run_bench( - topology=topo, bench_fn=_bench_fn, - device=resolve_device(None), - engine_factory=_engine_factory, - ) - - -def _count(op_log, name: str) -> int: - return sum(1 for r in op_log if r.op_name == name) - - -# ── C=2 Ring KV ──────────────────────────────────────────────────────── - - -def test_prefill_ring_c_two_completes(): - """C=2: 1 ring step rotates KV between the 2 CUBEs.""" - result = _run_prefill_ring(T_q=4, S_kv=16, C=2) - assert result.completion.ok, ( - f"prefill ring C=2 failed: {result.completion}" - ) - - -def test_prefill_ring_c_two_distributed_output(): - """C=2: per-CUBE distributed output, no reduce → dma_write_count == 2.""" - result = _run_prefill_ring(T_q=4, S_kv=16, C=2) - assert result.completion.ok - n_writes = _count(result.engine.op_log, "dma_write") - assert n_writes == 2, ( - f"C=2 prefill: expected 2 dma_write (one per CUBE); got {n_writes}" - ) - - -def test_prefill_ring_c_two_ipcq_count(): - """C=2: 1 ring step × 2 handles (K, V) × 2 CUBEs = 4 ipcq_copy.""" - result = _run_prefill_ring(T_q=4, S_kv=16, C=2) - assert result.completion.ok - n_copy = _count(result.engine.op_log, "ipcq_copy") - expected = (2 - 1) * 2 * 2 - assert n_copy == expected, ( - f"C=2 ring: expected {expected} ipcq_copy " - f"((C-1)·2·C = 1·2·2); got {n_copy}" - ) - - -# ── C=4 Ring KV (combined assertions) ───────────────────────────────── - - -def test_prefill_ring_c_four_combined(): - """C=4: 3 ring steps rotate KV around 4 CUBEs. - Expected: completes; 4 dma_writes; (4-1)·2·4 = 24 ipcq_copy.""" - result = _run_prefill_ring(T_q=4, S_kv=32, C=4) - assert result.completion.ok, ( - f"prefill ring C=4 failed: {result.completion}" - ) - n_writes = _count(result.engine.op_log, "dma_write") - assert n_writes == 4, ( - f"C=4 prefill: expected 4 dma_write (one per CUBE); got {n_writes}" - ) - n_copy = _count(result.engine.op_log, "ipcq_copy") - expected = (4 - 1) * 2 * 4 - assert n_copy == expected, ( - f"C=4 ring: expected {expected} ipcq_copy " - f"((C-1)·2·C = 3·2·4); got {n_copy}" - ) diff --git a/tests/attention/test_gqa_short_context.py b/tests/attention/test_gqa_short_context.py index b1191fd..bdc85a2 100644 --- a/tests/attention/test_gqa_short_context.py +++ b/tests/attention/test_gqa_short_context.py @@ -76,7 +76,7 @@ def _run_decode_short(*, h_q: int, h_kv: int, kv_per_cube: int, O: replicated; each group root writes the full byte-conserving (h_q·T_q, D_HEAD) result. """ - from kernbench.benches._gqa_attention_decode_short import gqa_attention_decode_short_kernel # Phase 2 + from kernbench.benches.gqa_helpers.short_ctx._gqa_attention_decode_short import gqa_attention_decode_short_kernel # Phase 2 topo = resolve_topology(str(TOPOLOGY_DEFAULT)) @@ -196,7 +196,7 @@ def _run_prefill_short(*, h_kv: int, kv_per_cube: int, Layout: same head-stacked K/V scheme as decode short, with Q/O replicated and byte-conserving reshape inside the kernel. """ - from kernbench.benches._gqa_attention_prefill_short import gqa_attention_prefill_short_kernel # Phase 2 + from kernbench.benches.gqa_helpers.short_ctx._gqa_attention_prefill_short import gqa_attention_prefill_short_kernel # Phase 2 topo = resolve_topology(str(TOPOLOGY_DEFAULT)) diff --git a/tests/attention/test_milestone_gqa_decode_long_ctx_4cases.py b/tests/attention/test_milestone_gqa_decode_long_ctx_4cases.py index c921b84..65ce1f6 100644 --- a/tests/attention/test_milestone_gqa_decode_long_ctx_4cases.py +++ b/tests/attention/test_milestone_gqa_decode_long_ctx_4cases.py @@ -68,7 +68,7 @@ def _run_case4_smoke(*, S_kv: int): ``S_kv=128K`` runs come from ``kernbench run --bench milestone-gqa-decode-long-ctx-4cases``, not pytest. """ - from kernbench.benches.milestone_gqa_decode_long_ctx_4cases import ( + from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_4cases import ( _run_decode_panel_long_ctx_cube_sp_pe_sp, ) topo = resolve_topology(str(TOPOLOGY_DEFAULT)) @@ -99,13 +99,13 @@ def test_case4_panel_registered(): C = 8 (head-parallel CUBE Group) P = 8 (intra-CUBE PE-SP) T_q = 1 (decode: one new token per pass) - S_kv = 131_072 (LLaMA long-context decode target) + S_kv = 8_192 (LLaMA long-context decode target) d_head = 128, h_q = 8, h_kv = 1 The lrab sub_w=4 / sub_h=2 geometry is baked into the Case 4 wrapper kernel; it is not a panel parameter. """ - from kernbench.benches.milestone_gqa_decode_long_ctx_4cases import ( + from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_4cases import ( _PANEL_DISPATCH, _PANELS, ) @@ -120,7 +120,7 @@ def test_case4_panel_registered(): assert params.get("C") == 8 assert params.get("P") == 8 assert params.get("T_q") == 1 - assert params.get("S_kv") == 131_072 + assert params.get("S_kv") == 8_192 assert params.get("d_head") == 128 assert params.get("h_q") == 8 assert params.get("h_kv") == 1 @@ -197,7 +197,7 @@ def _run_case2_smoke(*, S_kv: int): slide-11 memory waste); for B=1 only one rank does the work; no inter-rank comm. """ - from kernbench.benches.milestone_gqa_decode_long_ctx_4cases import ( + from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_4cases import ( _run_decode_panel_long_ctx_cube_repl_pe_tp, ) topo = resolve_topology(str(TOPOLOGY_DEFAULT)) @@ -229,7 +229,7 @@ def test_case2_panel_registered(): For B=1 only one rank works (PEs 1-7 idle — slide-11 calls out this PE-TP waste). """ - from kernbench.benches.milestone_gqa_decode_long_ctx_4cases import ( + from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_4cases import ( _PANEL_DISPATCH, _PANELS, ) @@ -242,7 +242,7 @@ def test_case2_panel_registered(): assert params.get("C") == 8 assert params.get("P") == 8 assert params.get("T_q") == 1 - assert params.get("S_kv") == 131_072 + assert params.get("S_kv") == 8_192 assert params.get("d_head") == 128 assert params.get("h_q") == 8 assert params.get("h_kv") == 1 @@ -303,7 +303,7 @@ def _run_case3_smoke(*, S_kv: int): 8-way across PEs within each cube. Intra-CUBE 8-way reduce on (m, ℓ, O); no inter-CUBE comm (every cube ends with full answer). """ - from kernbench.benches.milestone_gqa_decode_long_ctx_4cases import ( + from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_4cases import ( _run_decode_panel_long_ctx_cube_repl_pe_sp, ) topo = resolve_topology(str(TOPOLOGY_DEFAULT)) @@ -333,7 +333,7 @@ def test_case3_panel_registered(): Case 3: Cube-Repl × PE-SP. K, V replicated per cube; PEs SP on S_kv. Intra-CUBE 8-way AR; no inter-CUBE comm. """ - from kernbench.benches.milestone_gqa_decode_long_ctx_4cases import ( + from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_4cases import ( _PANEL_DISPATCH, _PANELS, ) @@ -346,7 +346,7 @@ def test_case3_panel_registered(): assert params.get("C") == 8 assert params.get("P") == 8 assert params.get("T_q") == 1 - assert params.get("S_kv") == 131_072 + assert params.get("S_kv") == 8_192 assert params.get("d_head") == 128 assert params.get("h_q") == 8 assert params.get("h_kv") == 1 @@ -423,7 +423,7 @@ def _run_case1_smoke(*, S_kv: int): each cube has work; PEs 1-7 idle. Inter-CUBE 8-way reduce via the lrab-adapted center-root pattern (root cube 6); no intra-CUBE comm. """ - from kernbench.benches.milestone_gqa_decode_long_ctx_4cases import ( + from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_4cases import ( _run_decode_panel_long_ctx_cube_sp_pe_tp, ) topo = resolve_topology(str(TOPOLOGY_DEFAULT)) @@ -454,7 +454,7 @@ def test_case1_panel_registered(): batch. At B=1 only PE 0 of each cube works (PE-TP waste). Inter-CUBE lrab AR; no intra-CUBE comm. """ - from kernbench.benches.milestone_gqa_decode_long_ctx_4cases import ( + from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_4cases import ( _PANEL_DISPATCH, _PANELS, ) @@ -467,7 +467,7 @@ def test_case1_panel_registered(): assert params.get("C") == 8 assert params.get("P") == 8 assert params.get("T_q") == 1 - assert params.get("S_kv") == 131_072 + assert params.get("S_kv") == 8_192 assert params.get("d_head") == 128 assert params.get("h_q") == 8 assert params.get("h_kv") == 1 @@ -548,7 +548,7 @@ def test_panel_metrics_helpers_present_and_correct(): Helpers exercised via the Case 2 smoke runner's op_log to avoid paying the 128K-config ``_run_panel`` runtime. """ - from kernbench.benches.milestone_gqa_decode_long_ctx_4cases import ( + from kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_4cases import ( _end_to_end_ns, _engine_occupancy_ns, ) @@ -579,7 +579,7 @@ def test_run_panel_returns_latency_and_engine_occupancy(monkeypatch): Uses ``monkeypatch.setitem`` on ``_PANEL_DISPATCH`` to lower S_kv to ``_SMOKE_S_KV`` just for this test — avoids the 131K runtime. """ - import kernbench.benches.milestone_gqa_decode_long_ctx_4cases as mod + import kernbench.benches.gqa_helpers.long_ctx.gqa_decode_long_ctx_4cases as mod orig_kind, orig_params = mod._PANEL_DISPATCH[_CASE2_PANEL] fast_params = {**orig_params, "S_kv": _SMOKE_S_KV} diff --git a/tests/attention/test_milestone_gqa_headline.py b/tests/attention/test_milestone_gqa_headline.py deleted file mode 100644 index 5fb6b9a..0000000 --- a/tests/attention/test_milestone_gqa_headline.py +++ /dev/null @@ -1,89 +0,0 @@ -"""Tests for the milestone-gqa-headline bench. - -The bench wires the new ``_gqa_attention_prefill_long`` kernel into a -single-KV-group milestone panel (LLaMA-3.1-70B target, C=8 P=8). The -4 legacy C=1 / C=4 panels (single_user_*, multi_user_*) were removed -once pytest regression covered those configurations comprehensively -and the comparative decode work moved into the -``milestone-gqa-decode-4cases`` bench. - -Single SIP scope (default ``topology.yaml`` 4×4 cube mesh; 4-SIP -headline deferred per ADR-0060 §B-item-1). -""" -from __future__ import annotations - -import json - -import kernbench.benches.milestone_gqa_headline as bench_mod # noqa: F401 (Phase 2) -from kernbench.benches.registry import resolve -from kernbench.runtime_api.bench_runner import run_bench -from kernbench.runtime_api.types import resolve_device -from kernbench.sim_engine.engine import GraphEngine -from kernbench.topology.builder import resolve_topology - -BENCH_NAME = "milestone-gqa-headline" - -PANELS = ( - "single_kv_group_prefill_gqa_c8_p8", -) - - -def _run_validation(): - topo = resolve_topology("topology.yaml") - return run_bench( - topology=topo, - bench_fn=resolve(BENCH_NAME).run, - device=resolve_device(None), - engine_factory=lambda t, d: GraphEngine( - getattr(t, "topology_obj", t), enable_data=True, - ), - ) - - -def _sweep_json(monkeypatch) -> dict: - monkeypatch.setenv("GQA_HEADLINE_RUN", "1") - out = bench_mod._OUTPUT_DIR / "sweep.json" - if not out.exists(): - result = _run_validation() - assert result.completion.ok, result.completion - assert out.exists(), f"missing {out}" - return json.loads(out.read_text()) - - -# ── Registration ─────────────────────────────────────────────────────── - - -def test_bench_registered(): - spec = resolve(BENCH_NAME) - assert spec.name == BENCH_NAME - assert callable(spec.run) - assert spec.description.strip(), "description must be non-empty" - - -# ── Validation run completes ────────────────────────────────────────── - - -def test_validation_run_completes_ok(monkeypatch): - monkeypatch.setenv("GQA_HEADLINE_RUN", "1") - result = _run_validation() - assert result.completion.ok, ( - f"headline validation run failed: {result.completion}" - ) - - -# ── sweep.json shape ────────────────────────────────────────────────── - - -def test_sweep_json_has_expected_panels(monkeypatch): - data = _sweep_json(monkeypatch) - assert set(data["panels"]) == set(PANELS), ( - f"panels mismatch: expected {set(PANELS)}, got {set(data['panels'])}" - ) - assert len(data["rows"]) == len(PANELS) - assert {r["panel"] for r in data["rows"]} == set(PANELS) - - -# Per-panel architectural assertions for the surviving single-KV-group -# panel live in tests/attention/test_milestone_gqa_single_kv_group_prefill_panel.py. -# Decode architectural assertions live in -# tests/attention/test_milestone_gqa_decode_4cases.py. diff --git a/tests/attention/test_milestone_gqa_single_kv_group_prefill_panel.py b/tests/attention/test_milestone_gqa_single_kv_group_prefill_panel.py deleted file mode 100644 index 9935254..0000000 --- a/tests/attention/test_milestone_gqa_single_kv_group_prefill_panel.py +++ /dev/null @@ -1,121 +0,0 @@ -"""Tests for the single-KV-group prefill panel in milestone_gqa_headline. - -The single-KV-group (LLaMA-3.1-70B target) prefill panel wires together -all Increment 1–4 work in a single bench panel: - - C=8 head-parallel + Ring KV via snake-mapped 2×4 sub-mesh (Inc 1, 3) - - P=8 intra-CUBE PE-SP (Inc 4) - - d_head=128 (LLaMA-3.1-70B), one-shot prefill T_q = S_kv = 1K - (scratch-limited; LLaMA 32K headline awaits Q-axis kernel tiling) - -This file verifies the bench-config wiring: - T1 the new ``single_kv_group_prefill_gqa_c8_p8`` panel is registered - with the expected dims in ``_PANELS`` + ``_PANEL_DISPATCH``. - T2 ``_run_prefill_panel`` accepts the new ``P``, ``T_q``, ``d_head`` - kwargs and drives the prefill kernel to completion at the - milestone-target ``(C, P) = (8, 8)``. Uses smaller T_q/S_kv to - keep test time bounded — full 32K runs come from - ``kernbench run milestone-gqa-headline``. - T3 Existing prefill panels still work when ``_run_prefill_panel`` - is called without the new kwargs (backward compat anchor). - -Phase 1: tests only — production code lands in Phase 2. -T1 fails today (panel not registered). -T2 fails today (helper doesn't accept the new kwargs). -T3 passes today as the backward-compat anchor. -""" -from __future__ import annotations - -from pathlib import Path - -from kernbench.benches.milestone_gqa_headline import ( # noqa: F401 - _PANEL_DISPATCH, - _PANELS, - _run_prefill_panel, -) -from kernbench.runtime_api.bench_runner import run_bench -from kernbench.runtime_api.types import resolve_device -from kernbench.sim_engine.engine import GraphEngine -from kernbench.topology.builder import resolve_topology - -TOPOLOGY_DEFAULT = Path(__file__).resolve().parents[2] / "topology.yaml" - - -def _engine_factory(t, d): - return GraphEngine(getattr(t, "topology_obj", t), enable_data=True) - - -# ── T1: new panel registered ───────────────────────────────────────── - - -def test_single_kv_group_prefill_panel_registered(): - """The ``single_kv_group_prefill_gqa_c8_p8`` panel must be in ``_PANELS`` and - its dispatch entry must have the expected LLaMA-3.1-70B params. - - Headline config (scratch-limited; LLaMA 32K headline awaits - Q-axis kernel tiling): - C = 8 (head-parallel, snake sub-mesh) - P = 8 (intra-CUBE PE-SP) - T_q = 1_024 (one-shot long-context prefill) - S_kv = 1_024 - (scratch-limited; LLaMA 32K headline awaits Q-axis kernel tiling) - d_head = 128 (LLaMA-3.1-70B) - """ - panel_name = "single_kv_group_prefill_gqa_c8_p8" - assert panel_name in _PANELS, ( - f"{panel_name!r} not in _PANELS; got {_PANELS}" - ) - assert panel_name in _PANEL_DISPATCH, ( - f"{panel_name!r} not in _PANEL_DISPATCH" - ) - kind, params = _PANEL_DISPATCH[panel_name] - assert kind == "prefill", f"kind={kind!r}, expected 'prefill'" - assert params.get("C") == 8, f"C={params.get('C')}, expected 8" - assert params.get("P") == 8, f"P={params.get('P')}, expected 8" - assert params.get("T_q") == 1_024, ( - f"T_q={params.get('T_q')}, expected 1_024" - ) - assert params.get("S_kv") == 1_024, ( - f"S_kv={params.get('S_kv')}, expected 1_024" - ) - assert params.get("d_head") == 128, ( - f"d_head={params.get('d_head')}, expected 128" - ) - - -# ── T2: helper drives the C=8, P=8 prefill kernel to completion ────── - - -def test_single_kv_group_prefill_panel_runner_smoke(): - """``_run_prefill_panel`` must accept the new ``P``, ``T_q``, - ``d_head`` kwargs and successfully launch the prefill kernel at - ``(C, P) = (8, 8)`` with the snake-mapped 2×4 SFR. - - Uses **smaller** T_q/S_kv than the headline panel so the test - completes quickly. The headline 32K dims are exercised via - ``kernbench run milestone-gqa-headline``, not pytest. - """ - topo = resolve_topology(str(TOPOLOGY_DEFAULT)) - - def _bench_fn(ctx): - _run_prefill_panel( - ctx, - panel="single_kv_group_prefill_gqa_c8_p8", - C=8, P=8, - T_q=8, S_kv=8192, # smoke dims, not the 32K headline - d_head=128, - ) - - result = run_bench( - topology=topo, bench_fn=_bench_fn, - device=resolve_device(None), - engine_factory=_engine_factory, - ) - assert result.completion.ok, ( - f"single_kv_group prefill panel smoke at C=8 P=8 must complete; " - f"got {result.completion}" - ) - - -# Backward-compat test removed when the legacy multi_user_prefill_gqa -# panel was dropped; signature-extension contract is now exercised -# implicitly by the live single_kv_group_prefill_gqa_c8_p8 panel.