From f37b9bfe5703a2b37b12a2d0236d53adec4d8759 Mon Sep 17 00:00:00 2001 From: Chris Sidebottom Date: Tue, 6 Jul 2021 15:35:02 +0100 Subject: [PATCH 1/5] Additional Target Hooks RFC This is the an initial RFC for adding additional hooks onto the `Target` to allow splitting up some of the compile flow but also unifying the registration of these additional functions. --- rfcs/000x-additional-target-hooks.md | 146 +++++++++++++++++++++++++++ rfcs/assets/000x/bypass.png | Bin 0 -> 59627 bytes 2 files changed, 146 insertions(+) create mode 100644 rfcs/000x-additional-target-hooks.md create mode 100644 rfcs/assets/000x/bypass.png diff --git a/rfcs/000x-additional-target-hooks.md b/rfcs/000x-additional-target-hooks.md new file mode 100644 index 00000000..b7791622 --- /dev/null +++ b/rfcs/000x-additional-target-hooks.md @@ -0,0 +1,146 @@ + +# Summary +[summary]: #summary + +In order to enable flexibility in how individual targets are lowered and built within TVM, this RFC proposes supporting additional hooks on the `Target` and that the target becomes the central place for such hooks, for example: + +``` +TVM_REGISTER_TARGET_KIND("cmsisnn", kDLCPU) + .set_attr("relay_to_tir", "target.cmsisnn.lower") + .set_attr("tir_to_runtime", "target.cmsisnn.build"); +``` + +This defines two new hooks as attributes on the target, referencing functions registered into the central TVM registry. In similar fashion, external generators (currently accessed directly in the compile engine) would be grouped with an appropriate `Target` as well: + +``` +TVM_REGISTER_TARGET_KIND("ethos-n", kDLCPU) + .set_attr("relay_to_runtime", "relay.ext.ethos-n") + .set_attr("constant_updater", "relay.ext.ethos-n.constant_updater"); +``` + +Collecting all targets under the `Target` functionality and making it clearer which hooks apply to each target. + +# Motivation +[motivation]: #motivation + +Currently to introduce an external code generator, the entire compilation pipeline must be recreated; this is necessary for some targets but in the case of simply re-using existing libraries or introducing a function call to use for an operator it can become more than is necessary. It also exists outside of the main `PrimFunc`, meaning it can't be inspected as part of the entire main graph; this limits the effectiveness of techniques such as memory planning. By introducing the hook `relay_to_tir`, which is similar to the default `lower` pass in that it returns TIR, it can be inspected by the memory planner and other analysis passes that only work at the TIR level. If all that is necessary is transforming into a flat `call_extern` (such is the case for the [CMSIS NN Softmax function](https://github.com/ARM-software/CMSIS_5/blob/develop/CMSIS/NN/Source/SoftmaxFunctions/arm_softmax_s8.c#L86)) then this can be left represented as TIR and be collected by the host code generation. + +In the more complex case, we still want to take advantage of memory planning by using `relay_to_tir` and inspecting the liveness within the TIR graph, but instead want to generate out more complex calls (such as using the [CMSIS NN Structures](https://github.com/ARM-software/CMSIS_5/blob/def6f800f95661eb3451d317f7d0dde504f6020d/CMSIS/NN/Include/arm_nn_types.h#L81-L90)); the `tir_to_runtime` hook can be used to build our intermediary TIR into a Runtime module similarly to how the existing external code generation works. This allows writing of external code generators that also get the benefits of any intermediary analysis or transformation that TVM offers. Alongside being able to use the analysis passes, code generators can extend from existing host code generators, customising only the generation which is relevant to them and gaining maximum benefit from the existing optimisations made in TVM. + +# Guide-level explanation +[guide-level-explanation]: #guide-level-explanation + +As a user, you can pick from additional hooks to bypass certain behaviours of the `Target`: +* `relay_to_tir` - Custom lowering direct to TIR +* `tir_to_runtime` - Custom code generation into a runtime module from TIR +* `relay_to_runtime` - Full compilation flow from Relay to a runtime module + +To illustrate where the hooks are placed, please refer to the following diagram: + +![Diagram showing the splitting point of relay_to_runtime, relay_to_tir and tir_to_runtime](./assets/000x/bypass.png) + +These can be registered on targets using `set_attr`: +``` +TVM_REGISTER_TARGET_KIND("cmsisnn", kDLCPU) + .set_attr("relay_to_tir", "target.cmsisnn.lower") + .set_attr("tir_to_runtime", "target.cmsisnn.build"); + +TVM_REGISTER_TARGET_KIND("ethos-n", kDLCPU) + .set_attr("relay_to_runtime", "relay.ext.ethos-n") + .set_attr("constant_updater", "relay.ext.ethos-n.constant_updater"); +``` + +## Relay -> TIR +With this change, this path splits, depending on whether you wanted to generate a full `Module` or introduce some specific TIR nodes into the code generation flow; the addition of the `relay_to_tir` hook allows you to write trivial external TIR generators such as calling out to a third party library: +```python +@tvm.register_func("target.woofles.lowering") +def tir_generator(relay_func): + """A simple TIR generator for testing""" + ib = tvm.tir.ir_builder.create() + A = tvm.tir.decl_buffer(shape=(8,8,), dtype=relay_func.params[0].checked_type.dtype) + B = tvm.tir.decl_buffer(shape=(8,8,), dtype=relay_func.params[0].checked_type.dtype) + C = tvm.tir.decl_buffer(shape=(8,8,), dtype=relay_func.params[0].checked_type.dtype) + ib.emit( + tvm.tir.call_extern('int32', 'woofles', A.data, B.data, 8, 8, C.data) + ) + + prim_func = tvm.tir.PrimFunc([A, B, C], ib.get()) + ir = tvm.lower(prim_func, name=relay_func.attrs["global_symbol"]) + + return ir +``` +This is then registered on a target: +``` +TVM_REGISTER_TARGET_KIND("woofles", kDLCPU) + .set_attr("relay_to_tir", "target.woofles.lowering"); +``` + +## TIR -> Runtime +Extending from the above, a second hook is introduced to do further transformations from TIR -> Runtime named `tir_to_runtime`, this bypasses the default `target.build.X` and instead uses the registered `tir_to_runtime` build: +``` +runtime::Module BuildWooflesHost(IRModule mod, Target target) { +// ... Custom Code generation here +} + +TVM_REGISTER_GLOBAL("target.build.woofles").set_body_typed(BuildWooflesHost); +TVM_REGISTER_TARGET_KIND("woofles", kDLCPU) + .set_attr("tir_to_runtime", "target.build.woofles"); +``` + +# Reference-level explanation +[reference-level-explanation]: #reference-level-explanation + +This functionality is an extension of the existing use of `attr::kCompiler` to provide a hint that we can use to lookup attached target attribute, the compile engine and code generation flows can choose to store TIR and/or generate runtime modules based on the registered hooks. + +## Relay to TIR Hook +[relay-to-tir-hook]: #relay-to-tir-hook + +This can be added into the `compile_engine.cc` by cross referencing the existing `attr::kCompiler` with the `TargetKind` registry: +``` +auto code_gen_name = key->source_func->GetAttr(attr::kCompiler).value(); +auto target_kind = tvm::TargetKind::Get(code_gen_name).value(); +if (target_kind.defined()) { + auto map = tvm::TargetKind::GetAttrMap("relay_to_tir"); + std::string custom_lowering = map[target_kind]; + auto lowering_function = tvm::runtime::Registry::Get(custom_lowering); + cache_node->target = key->target; + cache_node->funcs = (*lowering_function)(key->source_func, key->target); + return CachedFunc(cache_node); +} +``` +By placing this where lowering currently takes place, it means minimal changes to executor code generators as they call into `Lower` in `CompileEngine`. + +## TIR to Runtime Hook +[tir-to-runtime-hook]: #tir-to-runtime-hook +Instead of replicating the current external code generation hook, it is proposed that this hook exists in `build_module.cc`: +``` +auto target_built_mods = FindFuncsWithTargetBuild(lowered_funcs); +auto ext_mods = executor_codegen_->GetExternalModules(); +auto extra_mods = ext_mods->Concat(target_built_mods); +ret_.mod = tvm::codegen::CreateMetadataModule(ret_.params, ret_.mod, extra_mods, GetTargetHost(), + executor_codegen_->GetMetadata()); +``` +This means the hook is integrated at a higher level and included in the compile flow without executors having to be aware of how these modules exist. See [Relay to TIR Hook](#relay-to-tir-hook) for how the `TargetKind` registry would be used. + +## Relay to Runtime Hook +[relay-to-runtime-hook]: #relay-to-runtime-hook +This would replace the existing `relay.ext.` lookup in `compile_engine.cc`, essentially using the same logic as [Relay to TIR Hook](#relay-to-tir-hook) to cross reference with `kCompiler`. + +# Drawbacks +[drawbacks]: #drawbacks + +* Different hooks are currently dealt with in quite disparate parts of the codebase which are being heavily refactored + +# Prior art +[prior-art]: #prior-art + +This is all based upon the existing external code generation infrastructure within TVM by placing additional hooks in the same areas as existing external generation. Instead of replicating this with named functions in the `relay.ext.` namespace of the function registry it instead begins to follow the pattern outlined as B1 in https://discuss.tvm.apache.org/t/target-and-attributes/6013/6 by @tqchen. + +# Future possibilities +[future-possibilities]: #future-possibilities + +In future, this approach enables rapid integration of anything that can be represented in TIR into the main compilation graph; this simplifies the transformation process for a multitude of external libraries. + +Alongside this, adding further hooks means external code generation can gain benefits from the normal `lower` and `build` flow in TVM. This then expands to exposing more granular methods in the driver api to leverage the compiler passes in TVM, similar to how they've been exposed in https://github.com/apache/tvm/pull/8110 with `lower_primfunc` and `lower_schedule`. This can is then regulated by the normal Target mechanism to route as appropriate. + +Refactoring the target splitting logic into `build_module.cc` alongside any external module generation makes this a first class series of hooks into a simplified compilation flow; this would enable the removal of external generators from executor code generators which currently proxy to `compile_engine.cc`. Eventually this could also be used for CPU/GPU split as a specialisation of a `Target`/`Target`s split. diff --git a/rfcs/assets/000x/bypass.png b/rfcs/assets/000x/bypass.png new file mode 100644 index 0000000000000000000000000000000000000000..cf725e1a7cbc5915f8cdb732ece9c0a838ce492e GIT binary patch literal 59627 zcma&NbyOVB);5X;NJ6jx!3hu`B)CJ60fGj155aA4cPGI^a18_-+#Q0uGq}6E>+Rup z&U?;#*S+8R`X71?)zwwoYVT)1{asE*4E;Ibb2vCSbO~`01voflb2vDJbyNi47rtW3 zc;L@_`_F3j`c~F17KTRlaAJmFLt8z2Lj&?}F61Wm_SQU%jMf%DWe&0a0k!&+kZOYW|8zEVf&rT`PpLx$z8?lj$ObR@I>j4FX9 z`-F@}6a{X3JNfH0sH#Nx3(0XJv$w`#O70&E2OmEN4_Duxy?HQLoR}!s4=9Xbrc10UJA^zo>SH;nl*L~`HdHB%DCL-d)Mv4aybduu!z`$PoW6L&t$>xH!zD-zH zjmZH)SSy(9S778(Qy!e`T&7T}7#qhfhzS1%J|p zR0a#)*ts=$|6)S=yVIBQ1Bjf!vI6$0!&UJ75n1SdJ{K~c(f1O8U!|8dM(S%b*y?{IHFxprI?EoS6}o_32Hk(Ek6YpqQc?(Q}JXWCn_0vKvH25+4Y(MV6vCXjYOj-?9 zKKixucF{>gfdTWVgLYBqkkz2nYw=25h8oE}lE+u`YmutY>`7z-RSD1zPe zv@J<2E5U_@KU9dhtS5amR7u!B9%Bo}{z3TshnyC#6l^0BGD64j8ad|X(Q9niPBwLo zMUAj^^ZusMU3+D{^N^v(Ltsw(pIbN+f24okvI+bC>qof6x1fLD z4n&gw>(<*y9P?kdf~-`i|GG^lz()Gl?eI%91UNyo#(tM0w%r-OKmBEvYy-Zx5yOqJ z0p7ri(h+WVCqoq%!Tc+1$;?{{3;W_@Bc^tD2k(+as}W(NfXTd{uYweew1!N}W)^aT zawj=|m){FbA_X4&c=;lfaaMHf>0h`YhxQT6lyY!`J~^RpR?qArT?jD+QP zjYcSsvqB!oq4E(-v+$b{DUp&{E-o$s0RdAyut#9K5LCS@HNn>(erMm)gb_nxP5*I) zN-?LoyUt3~zCnbdXJFuS@YD29>+wA^f$Z)1q2&uWyB}EvlMjB~Im57AbKzm9=jA1O z3yYI*qOiuY3b4bYZSawyz}(=i+TzN8vd_I zCMEcajhaeKewa^Mgz3uNnWl@;c<1GC<<-)nIl;?1|1W2FAsqw+1lZrH7#T5UU%=k@ zG!pJ`0WTo&63?Sf2?Ynq?91uM+W`T+H`ELa3?Dv7(Z#H=J*id@u!|%3HK~Ih20b;L z+$1z#lNpRJ=U^WnpXKG{o}QlR>1i(X|11D5Vw92TnRe@|_E^1XtFK?>jL=%k5G$M? zcrX0@{ewGJ3JMD5=Cmx>VGDVhyopb|_t!gApW_`I_}IJv!!8{Ur_L02D>R|4B5{od zxJ79Y2!w@&Wol}wkoQ^>{ogy-GV%TU>ME`7Y`egPe#QxS0D(UKBj`w{PrI8uO1uOL z8wdwjrMCbNg!x%nHSN6iJ@CSJ$E*1B_|`1piN(rXv-3G{$tCt@jwZGraoxs2F6 z^+``Zq-}ge8C{O@iVPnwMui$@u?wbs-y5X2)+2S^-f&UkzvAQLQ&PeTpTS=7q-h>d z`i<9Ff9=4GM?Zh^R%2^z(r&7*EV6-?4G~B0>F)OPRZfeMUQY6}qgqXnanUQnM~6k( zu90XqCh!7A|C@j1gOk3#!Xb7_x)G4L8ikw`mNr8C}n}cFYB#VQkBh^GzD9KI4uJ38cSD#Hw@UEmPUrZ6U6X1H9g9=0&!C|=` zxU6VoN=$#2W1`XEjY|{-EEBk3O2Bx}xt{zHgh}j{b_^O$$iE09Je1|n=LQ|&c4PbR z*Z;4fy0)^hy#YfN+J|^7t@^T-*jcs3U%u?saPM zx8XA8{pPQPZWRB~9NEHmSv=HT3~0+0`CdD6XZj0PF#mr|usr_>^Be@u&b`JCWbTWH zF!UcF!nOolfYtvEx;(IySBl1Jty;C;Me;9MG&pWqA`tqG3H%ed{s}wSkC~w${E@qq zx}W5KbiNS>p%vbuVmxDdNr{PyfQZ=E))pNd{SF(g3jg2FTzwtwEdB~{d~wUHg%9>3 zxA^okZ<|i*F;kBAB^nwQD$;*@E75=A4bx}jVoCwhE+vQS5@QN~l3-koiEu3wnAq49cp>FwWwv&9rKPYJ{<9JL^VY2H2UY?<24F7>k72l`;jK-Jy`JvB zkC~P)8ZF#0iHSdYp$R)WI_l_Gwqsy(d=>m-EUsx$7-KZ_)GB zS6VQKdj5E>-N3n*V`AAQfGIVGAEqvI`qL0a#Ur5@|An7$&DGyWGB#p-e6c=JDK#~< zV0vWO%rFu2|7mFW^j+j3Mh(C7mAc1#Ycwj)%*_-trSAtS1G0s4_9Z z!RwLHiQg)8`BGIlO+8nu`^Hx7itJ_Eu=}(d+12gs7!y4jJqA54J&}yy)8sI3=fS54 z?)U})l0-!0$K_&)r3{BqP~6|o=J!%RlJE!OHe~v|ubxK39C6D&I5H{<3m+e!lyorl zwctAsZ)#A;Bz1#>FA^?(G3g^W=rzZ(}zjkh)9QD zlYD(yM|?V)rl!D#pNKyYwIOAfV<_nU^g5WV=|saLlai!5V5jQ@6)Wuz<}F(eF1q?Z zCJbzhiIipkc|kQPl_ASBZPL_k5mngZf*k>{yhE$Oeeb%XA3A`;NWttVeJPlB1sx9xY#bdiOl#}wKpJ8l!k?LfNk_b!^ zy(45VoC2YaIMA(@H}>2!&Pex3*8|k{Pfqr)A@WUYIwsPbo#}ne1HPqL9LD{W3Q~ z1RF1B14|A+*S_BbM4b_%C$)2ru0Om!nL0L%Yq&s~RvUWx5jGBx&lTL@;Q8&L<>T-$ zolt#MSPSoy@PXjt?tMjWS9Ph|MNq+F;iQisBbzMsnDL>vA%*y!Ug$}}RH*6A^(j6+ zi9ux65@(qLQUqklDfL(L_2niexLNo28%~x&NUlktgzP(f&6me$FdG1N5=p+be9X6k zSn-^M;*Q6FZvQWu4pDh|xt*Qeix)3q7OBI-!*gU(gnVI{94rs=PNhO+)5?;>rJ>2G zso}avI66Awu$*~AL{x1C-5LHZ+^g8v)3ZHWT@1JhCrJPrBr)8F*wy*=!qUmf2?r-y zr_pu3)-K5HdIy@nzrSy3X=(brsi`S8Hug;JBi;wbo-iU#9)}$ROUp(xC>^*7gN*;| z?2L__-NDh3-{XdjbDCbWYHE6#n~STxvlBvQ_xNznM(FUu3g5IA?>P?0-8+#|C5Qj#l!eaWT zhl!zKOmgzw&H3&J)uK=mUKJO!wjhi&0k4O?I7a2q*k;+30@Ng$j5-ZmjEqCGvq{-K zVI>yRWp{UX!^3iPZ%wW`Lh#w6+dS)F)NrCcg1pa@A^V8z}GT$Zy{e6DA?clCjn^nSAOEO-7e z1pv06C5`D!*4fnsM8nBkdB4;=jY_)PKXGxYL9a-R29pyaBmeTUzjzTsBb%C=$j{l& zOLEO@$|IY`p9P2 z(8MIy^NVBNxP26d`C9?3b`%p=5{fSy=vY(L5?{W&Z1l?&+E`%!N1V|A5oa+{OH0dJ zAV9xz2PDQIFj%NrU1d2t(dc%0yxL`M=PoPOZ$I}1u~?^3L4ag|QvXKvp&G{wb z^8qT`)!!fb>SkI`R#ujbj4bhXr6Xj0Xb}2v<$1K!0w_A?{YUM(UqpI$0;K%zSDR^G zC#kaq_wkUrTs8CGx5x1G^u1rt)DnVF=jzG%_~u7O#qVHBO)R~D@#v9)793N>3Y61CzX+zxjR{GWo*2; zxtZh1qXakH!ps{kLSf_NRIFWJ)k@M@A<&d_^IN|wT1n)c@e`4im`Yiy*z1A1G*R<)}6tYoDqS%~z3WX)%w9dmIw*0>Vswlsr0^p}}icB{jTtMS`mI z;DAK>i`n@Obo;`R`SL_L-)MSvc4|gr?6pc!vYx4_D2BeilSw(%yCmu`K@zF?*;>1# zBu@R;DkAyAaFi!unvROF1%%_Hdm@g4Z!5E_dL&_wKFf z+ZNHCe}N$$#0&iWQ0XuUqvOfYV*l2tK6u4vC&5SG! z4Lw)deOg<`;uy((tJBJwU0Fw5pEmao$trqPeLA#X>-qzUR|!(-opqy-R@CGEoDn7} z%5HLXeX2wV9?6K}kWwf!9Yxi;d_}ZTRFX@0o8W|Wtx_RhTUs()vf2g`;`pl;%Z&V; zt!^eFB!nbX#!cdr@KhvopN?H!@ij!3l`whSUHNzJ?z`jwF;+TATMQPj5>9M;fk2hx zDIv)rT~xiTo9Nq`VYX+yfVHMs%8_AsBqvCrrkcED$UEQ>Nq;E%&&(Hm<2r6*+WOv|Fn#-LeU%3 zMMRGFgoM~4GowoWJ~l3{)lJ(Qf<4~bdv`EjgFT1hL$r~?UI~~@6x=J30&v4JEP^?rYv6ELp>{3N(bK6L@=C5xLyqg2I95pBIYDoW_G( z%4n@iEk%8CvY9rq3@QEh_%Z9_q@Tna>^AWaSm0U9XXaw=2&EJk)Mb-6PG%~*CyQAX z6p@wc>155?Q6mCg&@{Lw4DyXtm?Hi@3c_cUDN za#}`-p6ni-j=xQk*L2Uq?)!O#dN{}qS4@c&q1nH5#jo?<}a~c)3N#dt(GAP0YYAZ7PqjnuV9vQS9oFW zE8Q}KvhNb~7d>JTKRd&S3`|Wo)|AqoJ6}jZmh-@CjnNq=Oos)R2j`}v?JYnYNaBzV zC!TI~dlbTm7=-4luIgapRu+2R*Eu>)Xjb*B3XoiP)I03v*1L7hRGQPxjmq94fm~ol zbRqL&ys^PaVSH++f%2Ed=`llUC$qApBoG|c@sB%ETHmt12?YJG^(*w;@F~t&m0etl z(5ig+_^Hu!?33!`1@1Zu*7<4>QR(%sRllXBua90!bkF`S!xA!VB=>d1UvfZ!r6IX9 z!XMrHN#yHco(+NmcjwApR9{m^|?iCI2nn?1c4`gs}5eQch+flsR(14$)P>rXp`QVTY63)k@x4RJfuZie& zZ5ctq$Ob@jH@NYLLwfF~K1^N_@7g>?Qatt$X(0f{t`O`ojc+t&2jy~>)EjZCXMd!%^; zS4>#x7lO~Iz#N;LoQ#@&Y9f|3ALFPo3!*JNV^k?w zyj;d>v>BrW#?Zn~-@neU>y{DBoj`{PEWJFE{1={^z-P140kKR@U`RUBLL)skL4EH* z;8#|-Mj$EW$*({YYPlDaEiVFq$%20RvID`JUV-cBpS^Qv0I&?;$Gkwtt5*)(Tibgo zow&r3@QOx(UtP{cg}_)7D5-uREgKH{^${Vg@JeH<>{5G?f!Hbl#pr71=*)Hw%lJA0 zh0w3g=~~aLZ_r!Ou9s6#c(^kTkn`yow?|5OOG`0HGmn77KKdbv((M8)qvnIs-P#nGoR>|8U(=Jy5IC;@+M`p3;v+ zQxw>Ntlde@FZclz9Qnz;ZqKB}oRWloZDZqvly2B_cWcXeXZ$^^M#d!vIot5j6jUYE zESLCFey0r`sjxc4md$Z<9notUDULM&jvFd$Wih~facM#6%XG>RE^h9Qj*ikH|DYfs z-|C#2s!`5Y0SXI+R1Fy!nZ+iLu+Y$&ZQ3t#Q2Uo_!)>$GR`vDuN%8S#JJ8+16yCW$ z)nXk>D=RB=b1B86^YhAzia$Z$O4SBj-^%uV%od9fI%LO~3ZV(gQK{Ot-YObT66{JhwHT7hPF-4pGim zQdSJcpfe7(UWqoXm|xuh%K??sE4!m5kmI9LGtp|HYFRZYoM6H`Q~*bds3aoOTOd+p zNRKJVav98O^L=v|1~|LfLjz#a=4ieM*N^wNht0CdTq9+MQWw8tVw7!dPk@BHSgV%B zaZhbM51owP{o(GKQMXA6S0Xnd{4QQe{Nw+383_6HpE8i$F!bb)$<*9jvdX~FP_Y>l zqFD{j+8$;-TxiUE%vwLbyv)I{hJ*u$RF|}6@2~e3EtztiB_UMU-xJx*JNdZ{x-nrj z%8P-x>QJD%#RMq_3-Z3)Pxx1Xcr*XTsuBCiLRmmHRp(3|>5Gf6y(Qliwi{Eldwfb| zKEQGl{}me>2sAzDBz^zhG69O1o~{MV@9ZVXZn>9fbov9g6qR*isBO{j$H|L6b#%q3BxIU4AIP6ZuG3sXH&XwpP zLg@gkz}Csh__1$hCh>)0&HFJXUaWPNK+TuJ!ot8w9-PYOCW?55z0UIN7+wg8{VhU$ z8&Jm#4@$&TJ#FM$Bh1xUK(utuB&Mw79S(H8Un>+Aa6oE}vw9H{Kww72f=d^o0rHGj zXY!(5TSI9Aab%klg_^Tfmb(j$w^zrjmGAYpmQr6Gn=S!q!A?$_X6JD8iDBr?j<^Q_ zr9hrUH1$9NYtRn@LNF8pEl@89>Yc8KxR{t2B|2zXoHM?#G=bk0oU}*q-+iWkiqCs+ zXm#p^6(5{8iwq|CAmq=`G0Lov9AainnW6iJ$5e@V_n^nOQZ7nLVTvdqOfQq5!T~3I zNmUiU%}V>}=Fqlvnw#xRg{jg}gUjj0T2HinP4TBnEK1RjS~Zb$cn1mD(n&u=G5`=E z=7TCEEj%HC0Lrmfl9Q7I;0LgLtm7m{Auc_1Ic2=b;}$UfDu-Pa!i&?>%w-_$xM=Qt zTwQJp@_M+YIW%1$?2B*lLjmkGd85XDTj3hZo|`{?eY{ynJ5E(n8MdqK?QI}S+i}uw zIInQug@^m);cGoQ_MsC@4$-Y(x&AK3Lyk(g*}?QhOG`^dC02>9!xR-xZW$ZN%j*%q zX2$bXhbAZE+O%;wE8lZx0lc`+Ix28gc?8 z5WC!;dRAvIv}>$ey1Tc4GM4Sj3nI~=DOL~DzbwrK{5D7p{jftvQo^K%9?G2qho zaQuN2yVdOc6RQP3Y-6VK1ER!c93PpOm_B{_#Lb=1G-OQ(^k5B70k~%%iL<`0?m?G% zi2sV*fe{-i<=){L?bAeB*#7Q4d^dZsA7^|<@e99b+{Awd8hQvm>5%Tt#lD$=l3!BD-Cm8}`Zn#yR1xgm zgN2dyKPrG;z>9cxjJlO(v5k#kg@x?Wdh5Lev9U&HOdO6uiK$AdN?LPZ&gxedj8KRn zpb+a!(F@gL(Idwc+qAyo6x+#TFxL57p}j5S1d#l6RM%l=S}Skvmj%}?bhe7ii9gKm z?b{zthMMW5e2-}FeHR{|`|aH^XeWkyR-Z0f{^9HO>Pk&dPY(_buBc#J=SB31-D(Tw z!9hnyFBGb-xyY2fjyuiBSOiL)YK!S$BIDj)SQhhAmfex5{QQckd{XKS&e)eHrTQq( zk(!W^b8_pzlSK(Y3ow>G>1L%v0;p$<28f~kGeD)*{v9zOAcs+xU#MIzb^K_V^mNlJ zAuX-5UUZJ>scI%n{XbOAFj>4tPfAJ}hFDIsG@B_43p69Mk+^OgHnT$QQ{2<~Xi$5# zyXrrW`=er=-z;1D1_r8(XaZGp%sY^ChGp9fIMfr7?{rmf-M>-`C7on42^DZc^&wpY7olnSx_4G-?|jE;Y#ANUog_+Np& zGmar(?0J#yqjtUHz~~s_{5&t@kms5|P8Rc5eYjIdu9@*>@1O#f0})v*SZuWDZ+VOS zWM;S&j?D7DA0Q6-nM%vD=Ii+o>Ce7LdXl=z0pOiMMCn~luEwtBj5NP=N;IKs_L(@7 zi04`r4CLg3RM>!j8yOo@zMv(z*q0#tvh^G;Hyy|~dAM)f_JxrGe8*Z^IOxoXpt&wC z?;9qOb0MmQ1^tH$VINvNFl+P{)D`PpSdDfkkmu{Y*gWsPT4EzX^7HDlec~9$pt3=n zpmeWw9HjHU0)9|r+V~%dwEhIvB8S~hwmC@5Sb=&P5Cn8?ixj=W0x!@Ih{%}ya;|T0 zS4sQfyeS>|WHF1710rm9Ee(x~bY=}ihIDG>(dJBUn{{`-DppIz(4CbqEw6>@WtvoU zUYt7(oNrK{Tpo|*(+rsys1A*e@9lz8ue&@IxfVg5+4APqY^6>$TZ~F-*o3HBE39T%_@jVLWItE#NXfOpMZ7lQoIaK05$?gE$uWfu+>CIq=@lhXHr6Zs(R z*|q)KaSI?ty!wOyM;n0&`n^<9#UU#>Bg_6NrXGUr-{Pko$CxI%`ca!F;g*rEFOSbT z(s^@`an2CI$fz<8JjnuDN;$7Er7YGCalhxGWLFl9sXF;18hWHwW8K;pw>`%yoy39p z+{^iRW$gzpAbBy*<4=6^?(W&cbNGy@R9?8kXxK;#68rPJiiH}=Fwbf*>~kl4^;{z} zd-q2S`4JdP`52VD+5b0)j8$_jo&yqE>o8oPR#L>@GSBWpHopEhj{um@Dmm95F^n6 zxS#bOW|#c+<)TKaFKjWDO}es&BNvnx*SI)27yDD|;CLzbthaUe88Yf}_$T*MKd(7mLlT<}6L1P- zOGi_RRUIuzjI?Fi0SIVKoXBU(eBI`vloWQ{ot+%YX6I@!iG$OjTD_xm0Pp))azKyc z*QXSenS|p~7P=!xSyIbpbwz4u2#0vfVHmA}1K-{}1Nu)(!Ho)3s3&1&GAK!Pb0*aR zBXeHtRmEQKn10nq!E;8%hKChB3n*>#`rOW5QJ4c6E>y^p>6)c(Wkr>pcY2tNVucsRaAXd20U~mf?B835?0fo2EE923cE| zJPKIjA_9BR9>VcwbwhM>+OjY;yEtrge2HE=vhdE(I+I{BwMe@_#T0DKQ@8DcT^icF_Aul zn9Q1n{7+Wai099wq>z8CJP1t)$1TXa;{%x(4IYWWJr%XQaux)zTS)1*KdehzSY6CgPKxIBO69W7g}A&g z_*yy$k_#r)A(he#Iwa%HE|qoWRRJFP&suXoew@UFFBD>HM*$+=MQ;-Ml)AcD!L(Am z^@)-A8g2Cggl;NJS7znFk15@b^T4ZmAa$`?A0Kj>_j2&+{Z};%&I=c!7g9x6acMKl z>*j?6{fbVc!_*)(z*iC{=b4z<6GJYbWjBVaqKcIV<60@R^CW`{)DiwX10f;+z5n&G%+~1nancGw+Ju*>|cv* z2Tx68fZBGT5bl|o`>S~bWn2r;U5{xhDe!#pw|`s^3qA^)~me^irTfLJKD z2eef3?~m<*75Q#2nKCh^E}ZE#xGp~2>xiB258n7yh#?nD0R$k}{!p=C7@+rq%jsEd zBBkT?BeV_He$C@U0B{o)Zqv4FTu+lvzyBk7A*5}xy&B7; zti5#w;MOqG4FGbm-GhPc3J%F(k(hD@qhqZ_OxY<8j+|QQHjHi}{mxx-P%BJLplmAc z_u>%y(1tg^p_pAtgN+1(FEj#CvCN(PuPk0_H1Pj*qp$q?d3X2>La#pb48=bjT}Vg> zCKb}Ff51b4Q!t0+1|Uqs8hoJSuj(R?VJtPFs*RSLx3{*5UQdHRP?Iw)6wuiwE%`^$Bud$I#m2CZ` zz4&Ti8{-y9xB3CY94#dl3BRXDnDmWaWEFmyB?MsOt)S2z^KCW0vOzd?->(~Q3wCTB z?5wlIF@r@gh4VdJW;A}an7lJb7Q+UR!zXT2FN4kN|4mX#W9+ca7~%xT7pVq1Kv`(2 z#Nt7fl2AnoAh)grPe7%z#3H_uj*%aK2EEj5TosF@IpCh@WGC+~MWA(dPGPtcTfaHI z_F0FG{+u!MD`l(|%tU;9Q;1Y85bSVf5bB0H`Pj|$<)5AcK#~gSP*+#Cukjv=`__ZT zD`V9umr2TbA=;hMSc#xl33jyugbrNZo?IKB!lf`hiK+yxCZzJ8Ow1yPF&EsDfNIiV zcFBoM%zQ}CcJTJ>T;D5S;sHBwZ*LD^Cb`$V?gNw}_E4KDsici(Hd_Ca*`9j>{0WBb z0${0A_d$+$*N%O=PGJ)NedUl~-?1%tKAh-|q|nPfWTd8)BC%7HSK`O6(lk-_h#bJN znv>Fe$NDVfEkUW)XfR`DT2Mu1@?6vBDcDJwQi51SgUZgqZ^BC0fh4~)uz zcxKvIA&q%*azceIG5Ui#-}ffUzcUz2*$gFvEAYBtUtP%FDA=6kfLYa}BPcIBvSFv&PAh>R}% z`@314sAC#=@xa4(E}t>$ywh{mk}dW(oN+AUeHABJ!XkZ(Blxw{iE;^8NeI_SYBYi ztMo5x{v-n5psKOt@ihSKJ(Ef5b~`}?D=*KaW8+mFK*0bxx24=KIupP2lJE+8-gnAx zUUWDeh4c1mmn-wdKKRSu?eBKmx-da-V}T|e_wMvZX7u;FO5|ICg@-)XiIegrC!;Vh zFaR2M0YO3KUV7O#OzGG!G#XM}1x$|X;I4El^oI{XylDG$75`dQ13(J`Z-Ct9N2n=> ziNrC7hllyaeXR}0oB%irjz2v=$u9s}7oh131m>!xWvj!5l+Pbj+#IjtDwyrIiT2~F z^mVm{-p3*Nyubo117PUc+Loz(3=a-YzmZe=j;3l6bx;3o9S`e{afaFCGCjrL5qCz< zr25+XxTP-{rOr-TX!LIIX#3cxgvBv1eSTlp7~&%=NgyR^yH&nB9+ zHqRD#IHgvBoRrj>AN)VYI{RH64w1^dN7?eYxaP90avc^=yJW<(MggZe*EN3jZt5%B z=R^i+K3?9W-d3QfhhX4*DAo5wF}dX1s20F_z~6sU;WjKIRR4pcUOxdwXfH44B(Xe0 zV{n8BO3&Gk*Sct$&b|c+3JZ@NBl*lE;6?`oGDv7*tWfd7d0$fBhv1%ZyW5;=j6WLM zv-kpQkj<5DUW`^?Gg^7}Zo$GLTu?EmH_|9t`@7+@&m*Mn0Mm`=-)}v6DE9-fn6Afb z?w(=7d9PiL>lKm3q(!Ud`U~UxYvf3D-!Z?9tuaRL;fdPi%9tu8+qmx)P=h?JdK)Lo z=qQ+gwrza{>GiLve(O!P)1zQgz`-@3#7z+Y!B_?|!^{)Lwu?23U){NG&N;_~jGs9F zGN@5^(CxQl$QopPuy4Q4ax)1LnLUg9I3xc$?ZJy?FkQJoR5Pq1Cc_;DF$VWm$hzYN%s^kQp)wtMgtajN& zkdaL`H{)RY|41dzO)PT-=%{L{`lB-zzcZiVn}AIg zcS?<_-&Lhk*&9|ZRO9Zl(;Y=$P!J0AQUWooJRcvNwD#X(zO;lwWC}H)^^TPSPT-%3 zHzLq(WaLG~<((}PBV$#!uJ8xQj%Gj*qb@@{^QykapUyCP$Ai5gzQ>z$A>GAPbfQE} z4MN%YrMtZZ)%U7hQ0QAO=Qo!#Nu0(L`Nk((BU;yJCxiAg6%iQ3Z~DW^y&luI#{^<^ z18QoR_h)A(3e*b?rFQ=c$_AVwSL`+3-66A4YT(n*yx8Gmho*M|jvY;NYU1LoD}|Gz zuBJJe%w3M?&~%fkq{FT(3z^YrpfoSXVs9#`*%ewnoaY^GX#tpUi0kN7izb@xx0W%<9e&xY(BZqjlZOWqDtF>q&PH=yEhGEC^$%A>4aT?SH%3_W6 zje`T{{h71mXnq{!x?yk9B-yxPc5Lu;tM6=E~U0e1Kx$7F; zvW04BLe|zy0NK==coi5uwY@s%`mjOY1dfQAs% zYegsLqT6?VLFV-(dr);Tx*@;w_79|Jl|(%$EcZB|LmQv)u`?`TCizr$R>(7LExTan z9i%D+!Y^NmkxQWK*#j*|k2U{oW=4pEbWqIsPU;gpoOkuc4UkM+qacSrT#bkggoT+` zSkC^^^K?#wJ~X3VXv4-U@s7dWQzYp<+cDggvz8X>`Z;-&iU^^rbJ9Z%-T3DBu4 z(7>L}B|LnR>KMOUb)G&Ju73EDs2m4m`pZj%{wV!+Bm?Rg@> zYQCE&outv(?hd)HeD?hQZ;CO)B1Q)v%JKlwVZLfk-VGFvP^_uiG<(PGrxVmUO$qAo zUryUL&OJkg`uvuP(@>dFzH(pb{5ym`*U0_)-K1@`YB}DC;s+JqrNcuqOiXg@@7Opv z0E0Hy`Z+Do*jMz8jOiFJGDQJdl@qwI|VX;&U`sC1^ z;g*d?VZ2u8!!@NHEq=nJpW=m;&)S0Q*kUU?AU=7cigvFMaGB;7@9b=i`qKQYms?*S zF!s!yD(##;;5M@53kO7(mk-=sYZGy>Z8qjLK@yCNJolHrO0w5BHljf{8|;367yBGm zRh~bDT&E{U$J;FFXh~~?Qt8+ z%{8Txrw~AjhkM^-1}zW`+oPu+XnJs+3c;tD9x2iZd&3Tih>5AKt?dcMBE=vhHP-FZ z6H`0~IAP}eUa62zmz-9EA2m6kBO`-@gRYx{pXVr$kqW@`bvA_>oeT_vn?oVvvJE9i zX8<%lQpRq&Qer;&4XpS2C&K!svNZ5n48SOALL;?!A+N~s!OaTVbk`>JkfDN}f_?9Z z?*wm82ckU|89xiF(FuJ0FBsVFla=`bKezAG}_3qsB8adU| zqs((298Bf6H3%sKs366$^C&*q{cdFbK%iAAu=cp<-@Q2}NKKWl-2{Tx0&ogDvGZOw zr0J5miog4&ld9@^qY8AFKhw4B5c9C@r4`SD8rZQ9znCV>R#i%TAr@5INEFIf%apKJ?2Jt(`4*fVKk%)|FC<;LL`jw<3SzmSw!ozpH>z0HN2?K&TLv`f9h{$G@(_B zP#Q>8dhBYx1}K1??m)xjTnZibt&ma*0zisc|3xmBa<`e*d_;|!=03W6Fwb?d40QQC zZtQDaST~MWadh38r}}nZTwFWuC%&UT@VI8V2bh&cLl5piZ>&MN4iZ;i>u_~3%RD_T z|J+Wu@g@m?k+qviJq~qRfqswe7%~6c>O^6t(rB?3ccxgx#STr=w~n(WkNS-@;!NDL zqpbpfnrP+=OZ><|s3TF$fAb;H09|Z8{`D4}urw<$HuyTg%mjK^m zJ(Y>qMzR%-t25;0-jR+9$V9+;dIn}1nPM0y2UJH}TO$KWn##rCtu5O6?gk4O_p?d2 zA~(s;?>*G5-tqkdUR$kQbe&^mv2wm^U+e@zi`+GDr`=^~j%v?hS88m@=NIO#5?$er zV&S0*I)AVWUbVfnni(%kUvgXUTus@I)2+JS&heS!O5(gsWol9PEc^YNa;7q0v)rii zV&4)6yVKUGBlBVNp8nZ_VWxsh@2`v4*v>G*gt3n%S&|9l!5wKHcct`)hxo%EbK7!c zALDtkg#V9vZMx3REG*npeQDxdk0+?SUuZU)w}YdsDc5%l)*wpgvB@Q7+LC)PzAlw1nUzdD)WyYLk z0#?z5M|zWSK>eLH~l*ElnWl4R$ZrjeVJ zm;Ix+VTf-d-fRBe!s?iW1Q*L_EIzk&(RbFF`;VS^e$w**vuC)i)jygwFf3z)aCUN$ zeKIMjd$CKlIQ3Izls2D~VtFx&%Ldu4^s+n3b%JQ&{Al|E3`y3LFd?`8ZnY~aQyChj zt--5{$+UM`j5BJ4PA1KorMjQGSy+0ZqGq%JfX&4xbpqis1^VC?ht?$BvF+uyZfABH z`;sBK$9@97&cWW9iP0a=wtjcAf`0LWJeQ#~sAl{luPOYPqQT74bBbcwahbJfPl9?BaK+ zO&CB)&TE#1zm$WvuhR%2&d{I|#^J&6ODfbEYSjHiHr^(Oq5PV0<5G!oF?SBbZfkIb ztZR)_Tl=1H8Y2zH&P)oa!tUa9KvH+Cy4|h)miw;tcyxR_;|==BjYR{*4m0-?c;RnU zr;-l#^tIGeqY3x=|G&|sT>XF7Xi}IfXMK&=1hi@>{7y`Ybs6W%jRli9IOZ^Tb~J7e z>SegO%dO_(r|X^Cxb1i-LAipvC!|#*Fo8)Xk-OiRHm7->j5uqy;L#J#*EuM;R8UJ3 z-`%V7x&Au4*guRQOIynr6^`ZXCp@g}UC! z$x##qbk*wE2wL}EgIHucE zTV`@^c<)q2w1j6}Av*>)!xGb_y2{7ykKN3USL>y}5=BM#ta02>67hN6L4Qr6$-5cL zWBam84ys>?0hoT{HG4{+@k=5H=W7kix~Lm5PkGZGsYqZ{!wpi)qIt)G@`ozf8&~to zhx&eAN*3B2-*P?SQ^g$Wtv&3U>vJ1xCg97IiM1ZG_m87g*htzsahE0gu z-lIl!dP>a}TAvsPBr4nt;LB*59^lN?t=KXL2o-y$<`dlqeu864H5dq1K)_uAc>W|z z>6$Ww%w}u|oSqS+`>I9B0Mn-jpp}~#tRsP-NU#(dz?Aopmirg}E|BnOCO&L_dM7Ll zu5gkdm^0~QH?ub0z;91mqZWelH?F1O_&)GaiYuoyShbAaI9SxMX#RHI&=yL3KOp6j zZKfjanZVAF1aTnSU!XCVCQMWcd4{grkp(o|o`*rCEGf0YpUSRP{=}wgyW@=Q1@;aE25Y)za zJZz1SX;zphHM+8kiSZJG9*q3)@Ul;u!GCX00J=XQMP~LwOMQKAGo`fI>JA{N=o)`i zXPItksRGWgIoZ#-l9)=;4e3kcp#+&gkDrLMNET>rkN zQ=|GG#vZCw*q(Oh39#+f&q8jG`B%#}WSi-e@Ar0`ay0>D@5zo%kl+j105s8iZ*cts z$3I@-lcP{sUKXxao(J#*1|qxV!Hvqe6;jQ~?ei#lf}7yEM{p8fvsD*^(q{9g*q_IW ze@jTgd#45kftFv(i(`c+9_@PlBP#^L;;QRZ-@m5vKKigCm(^<@*$Vd2KKAn&sN#wa29AGx( zT^V}YeIuaN%!aHypIhQPF0dGWHwPRZgy`0)Y$on{+c-!8vK6g!IUVS)7lb{+QQ^?h zpL$N-3a$yK-u&8mP5%#RZyi5bX_EhkY<~))P zLhT8kUBJn3gKg-FkQS_Msh8NSs&H34^Uf3abecnZS^R~dGbK78Zk<)petOM)(hEJf zBkWaKHb)+gUwwK8j#Z9F*cY`WAN#cBJ0UG>{VhjBw~(udrNbHWXa^sI5B47MDpxy``E8TP{cuZnO z2ZZod+-vrUS7-L%F-OAtLh|<1g)tTUPdjaya#I0;RV=>7NTc;(0puJ^Zi*rEhog>v z=MRtP3vYP_VSnmoW@Q~`8`3z%zdS0FWx%A^*W1LPn^hHY@IjLcBAcgz^zbueDBc*{ zDb=ne0?-6C^a+?K8n;u}n%YlbrSh{2D!dm{`zRZ*{STQ7#{ri@K`kDM`!GSZf|PH@Io~G62Hwe{aBh( zuao$SHzs{r=Mv;_)1B@m5+v?t%5ah>kbvDTv$%N%-{yhLCm@X5Cov6fyv6JeE<21c zT~*7&@Cgt9cDNHhGGo8+fX&L{K|rIU;-2h4 zYKcc>tQDZmLTCN`Lo6o8vkGixVjks|TdriE^ls;!W-8p|yEH%4rrCZTP zZKBJXM?pYI7sL(A^ z%y?dS=%Ld75vw^r|P$(VqQxE7T8Xqx#~<#$#{v{q5bI%$Ip8T`xU* zyYjIau1O_jt?%yU^YX<-abA@DJ;a6XkGi{ab)LR168ZlG>Lu5fRi--goK%OK0U$gw z(gz$Sa_GZxPaKmJ6HQT8*0VB&=JyHyF9$UB2(G(D;J)|{jab{@fF>nWi@`ePw=aKr z2x^MAqaSwftAJJ&NbNaV`o3KYdIW^*`=v|{LZen6C1ft&8?RG2IRksxEtf5Ib-yab zD1PwTkrLI=7{R{UjpD3c=b^xc1=qL+Vn7Wwd_b9ew2$WWN-csDf_|O2Cr@i4mMy!~ z+B)vt!zj4z3!qSCCc0YL%6q}YYr#bLO5cG)b}ou?^m#Ee04lq4%cD!fErEw4(_F<} z-7ktE_p;B&WfcK#HwS>W4P+B?bGNRCkwG6Sq?eSJu%>quy4qV4K;XYLYZmm{IQrO( z?*9TT^pcJTe9q#Q&-1X-cyx$lIC)Wi?#vXkzaXe?Izd4E$8BwGK?=No0Q6QmTkQ<6 z+D)8pjl{;>em-CnIgd-}wX4*EEUTdDuL~8*}bgM~?15m`ZPwpR zG!cyS<{n!ErutLLFoic%U&HVC$5!@b(m(rc9gd9e& z^$D9!qGK}(qz0p$$f>G|p}pAciH(Z7Ikb+%NeW)2$f^0KYvwIBh*mTAA|#(mh1wizuAMQ5d*iw`a^GbW_$#WfAtnWV_|U zR4f0ioAfO|qQM_k^q_Wb^3z)ac2)sLvj+E=XaFpsq4iZcnoA`fhi+SoB%dp191ev3 zPTvc?>8^_iDS<1DYNXS(>^42zS|*U$nngl#m-y&l`Yx#D2=;|I=L3YtyV_@vO;U$I zsR^-!%m@6)2nmAyG_5y+^wKLfEnzk-CMgY#TS>w4H4SW|W-1F7SS&lwyz4^w;E&^0 ziugg73E+Su*KGR4A6cU2V1ozRkDM|j(9WPt;FP%Vc~YPH0F9TGQS=AM;`kaSrrdvov<} zJlDhfu=sn-PLT<2)5n>c{>0l`FB(^h&Qirl!}Q65{37ukBSXXcR8)O(dQYC9DmC1j zAk!=qRu0J-zt^kVuhnnb@7VvMKdAq`+4oClsV%NRNePL!<0T6=Q#65@Jb4Zv^J_BU zkIySqVMG{5d}{?c^S|7#v3l>brYC7K>a@?79*5G-Jeja*-$3!CK=wTO^zRr+^(oGr z4e|zPw6NtE@pLiyXZPH7%{ysHDoT-gfJB8?;8_lh|BjTn3xEdz*o=?UI&c?pLd~9m zb~L$!%)ZHsgfkX4ie*?S>&en7d3`a&XHgf8AZkJ_eHOE((%xx)8|Z-XGYFEUiXx>% zA7lrn14T@0ZlmR(b!BreePD+zRw zr z2K+8G!_IiHdKu+2E)RGa0~&Z4$w@WS6%@rpo=U)22_6n5d)9^WS^KOZ3!p|P0NHfE z@+|)nlF#y-dGp7TJ)7w2HbMge^7_d8k^hhK;hN@H9IP-dpJQDU_ZQ6d{TrPJ5jKz5a7qJLy}eS8?v$Gu&hMERUuRmR@`1 z02B4sgCs#W8C6SXMhJTDS{~ZItieAsK!4x0UGG|TFK%`7u#&*~$7Uc_>R1IsAhrf1 zBehlx9boG(+mp#+ke|}VFvM1r~ zs1iY2wvbq6i3=W<@hW@i?sr-aQlZE>My~LAOFr?1d>t0Aq}jV8E~w5^Pb}3XcVvq1 zSY!-=(2fD^=Zri?U^m|!Pw#tF+5JD73L?r+ZV_|aJvwRQ6G9L!(|ve^>@JwSSo%Gr zu7dD=Aabct&R>Rd;xM^DrRMaEgW5vK=77|>A%Bw22|kNet314HA1l568Ob#QG>fQ9 zKUeL802t+wWkmU^F+meOldjVyl%ZdRjCC2;if8(`@TPnwD-#pw2L+N3W?U#gr{=as z)IK9O7PeMV{_egLH*fTK~`24^fOUck5#1o`vzpQplvyy zSpQh-bUbSKH-D6k((zZu-|r z&l$JBW%qdi)z3C|bo;2fBm8PU1)dy#e-~{cs0suBOI`t{x@V3}Qt7U@xA&Pr`>D~2 zyrA^-l?b08RW98I>I+kjDAQxa>I9~y!XBb&tHzh(tAq>rMw(hw+HLKvsPbDKy-A}T ztWr{aIW-w-W+%Oh+mrJg?C4$3xh^Lpe+&!Y>YLG`syv$Mk4;OZ(e~m$B~@oQ{oMFG z^OAN~(?b31Lr^SM2M<`-`LUXVjd8J-=yMLzBwx)A(rPp#gh6w+Anoh3;+bUgI@Y0N z-E9}zktX3o0r&QM_Sv9oM_T1u_t32wqcXbv^6o)CL|=+Qnx4nHjwcNpa*;gdLJubL zP<0n@)AU3wyWgtkMxnmY#(Fl0v|g`Nx7)iS<#YA*>S`H=WLLQ$G`CZ)Yo$>Kh^O_M zuyQ9bK6=!nSNQJFTOR^vb-&b6sKn82C6J2m`1IA5cb)f4nsceTrGY`Jc-mEL?4gyt zr%#_MXGi|^qV|0E5TkgS?(N$XrwSxI@aeVt_*;-BUG-dLxroh88$j#9$2Wb5C;1?} z!-(R~C}#9Ew;e@ZZ!W%lG=_bfmH%O4;4+6A)@me7s5CVDw(y(RubYqWTqz!jkBSoU zdCY|N5Gj7o8=A~Y*AFC9_AE0ZtV#D^@)IwTKsxsYBy!kDgoeSZpL`1$jKhpVZlyy|9)-)RL^l}C_=m*-}d=fTJE6m5}P=WaUmbi>%hPg9XJ_1y=0WW zzP@2ck(cY8eC`V`kaKz_04RspxYw@uCR4|LYa3wUSC&bnOMA~OL+jD}gi9|3kGbbG zx3<5SU(TNqyY0yyN011`N@%I8rN?9BES(MU`LRE`NXjFAe(j9D5t;klCalW>uVhZd zf``9kbShL+?njHZ=o?-yB2kMQYDA+%-?clJnr83U&k9u~A1LVQCD%XeUfBb3vMe*I zmeh*t-VaX_sr*0P`3}@Ama2Cc74IT|ki2>G2E%)_WK)7mZ>nVTEo$`K={B`%P7Y*{ z!hC1+;Q61?k~HI5u*VIUm8#PJ5Lm<1p%Pt%Qw+Y>3WK+$%oS=U5+jB5g&NYi0ltlM zDA_}n^ZEXa>%K0A6N7*L?VHb6?s_@feoZ{gcy$Y@Y-h@X!aMy{@BUiM{U!>hvEVsU z3idlNVoIbG79(Aa=V+Q`@jWUVPFptOQ4}!yxcx2C5s^narWMG?q#0 zCwq?S@Wy}V)DtLo4JmUlpMC8r{A}aT>QCcmKGG;n^U?DdOOpTAf(w6h8CTA5thE8? z$DJ3oXOe&ZPUG<^Tl*ky6)xov#%Nq36N&Pud_qzZMOIkwU$bKWtd%)Ihl)$Hd}x$l zLS8wKhM@XP56)H?C8W}V>!t!ZaQ&XBpfg?^k}dh`4OEsZtRggr5k!Tehb43b|M@Yk z(Lzfht(d5=NpVq@HSeQ|=x@P6k*RW=HIhsVw);vl^g%~%6cpb6oQ3E2=0p0Em|~=x zAMI%6SZX!4;==P@^sTQOiM7aT^0_CTfBV_jz1!mOUO7r|7wveHJvC;FAw(-->EY5{ zk5AXN!T;p<+hmD~FG;A5A2oWiFmlTlp5!fv636aATJK88?`7F*L{_V2%=WwcNx#c? zc4kR6pq9=<=i{6`vDDJ(iNXqdG1{{+fth~ukNAR)J#Ftha<#)Q&03FMkK#Y{2E;%| zbAkEfJVYJ*Ao+3?Epcg@Mn;aN(4TME9dV*u$#RG?;#A7NN6@UG1-D!%_Y@TqDs)_w z%pL$SZd78T^q=cd6Q(?=)Ai7V`rb)ol)Xcr%ZCLe-Q)XKw)7D+DmE&vym{}Bz}|YW zREBA84tvOYz0%Hz2l5y~`^ey>JgtfGSAyb^Rq5A%bXGP7PbV&$f!JZ@84>Q|`l(JF zEiTUVy*8jr>o3v&4DwauHz683N{?r)l1JVl{aeDA2E0R^Iy3(~%Y+holVGh6Hfqgd zddj!AbWJ?-!$4Os3*I%66kzs&0YvUM=Sp?Ut?)&y^Dg1jg^hag#pCVe&xrxaUz?iZ zBp+lCEU&Vn{tETt}z05ZQ0lq9*d1 za970gm1)YvG|l(w;@ck|^k1+(p;CF-`pJ($`_6xMg)`lt-u16vD5EVemFrSS;3}R^ zi%!f`D^{Oxy{;z03&zrX*=-cf2QW%hvh@8x;yGu`5LwDh5T z%4=zLgyAmJ`s)y#$;zFZ;Kp}b?P53&x*Tgl>(6(-pW|;romn2(Xbd5h2V{lk&&`U+ zrdor72&bG^hhqZ;8ZxYC<&wWH^=%b-%)MSeg|DjDqed7?*Q49 zOZoJ0VdL_ZE6qZ4V^?o`cq91yvHdE%JU84of<9@Kr}HlLrYR7GC}j&+4Cd$_Zu3JR z9oQs+MN{jmV1wEZ#AhimQ-X>FU5sXkTOQ7*=ZQISGNy9~B{NBHSI8calamAeAdHIN z|F#hyI3J+*wmYvK^?M_#<+2nqsEj(>4_N&{X<{yDA$2E}Yk!NT+HU?Dd^^q>E9%u{ z%m+?O+jH0gUr&x!V0U#^*=)bR4UPnD@Onr%#2gnlD{`7Fw4>+o1&%j^f{8T}c=6qT zv`6)%pmCCn$!xMLcn%4;?;G+D3Bfyy5k13;4$r`t937ro_}#Ei5)MhRj@C8Z_Sum zEP<~&O}kb-I8uY2kjOlW6f|Y=NZ{T-I#eQJPR*cDGZ{bNZ=$*Ko@2HBJp&B%CkXfO zY#WJA;ET2}{?Xwi0I?|({4KDc?UM?IPl(J2ol0FXQ=5+m9j5njVKw+Jkj&WL&p#fs zexsCq^i)L$>SbA}a(-yD(s3j%3D?b_}OEiR?)N|LU+$wi%=t;zUBH!4kHz;gA0du#oWcJF` zl!sPvvy^Gd+k1kj3(~x$Vcd0p%ysT{XPw8QY{~2?KffW9j_+rFH#(P9M)88%MlJ3{ zb-A?E1m%5an{lY+*qsfiz;bVzmcTI%m2+p}f7ZAGvC_^`488me5M(Y~z)$){h8bLb zC7!cTdHr2{PNZS@(@?`EHeBn;V}4c#d9c83(fg(YStFyRZ|x*6uv=^Qq_u@d(kgwD z`+}%ks5(AKsrQ1Hnm3yE`Jr98(x0W|dAMzwcN5jP%W-enIYop|r!-6Fi=U&x{a`pL zFW&82{hI5?)4tdHm$B`S>HUaQ#Qw4akLQb}n=%c7rQj$3W0dZ7Ztzq3*Bi5obi>>W zmpH{d;6+m~t+&SU;y22e-CZP)}y_pO>k-m|Bo+Y8#6ciLnv`Q{6725eS;^IZ> zxqsh~OUN-)?j@|j{Thed+i|`fRsn}o-{v-n_WDRI4(<_83kVEk#k=+(yv3;auRR2} zDI9xAqt?Wm=hDyf6K^qf(Z)C0dJfc`l{Z7J>#E^2Uj_JVE3f1)U&DO zg{?&`Nz~?TkFWZ9j%(!!?o^$!;+JWZxJ?)6%bL)1y~E}Y4qU=pC@;nRvl~A-^X2cP z^x(m_?9y~exp)Va)*^NrDk7Mo;z#got+MT$R~#&cw_afaY1lZ5GO&=0Kl)nqZEM2M zw5_b=V70OHOR0_@*Y`_J@Vf>N4aE`?p8vHeQ>SHQj9IXQC?q1{79k-a*V=5puYX?N zFnNklx*e@$DdQ-=ehFq2O~4JXpFv>|Rn-{!Hk8(ft-xy7DM;m5?ctE9|w!NU}Ok~P(Fx=edHGW{RhzMQ5X#a)ZO7JDs$XA1!4 zKi%2Ka_G2$BKjU4HPRWy2uCe9{XMPQ{CgB!de}+aGHad&HU_T^J{VLObQ>&=zrp_F zodyX0$FuJx6rE7|7h)#An6lcK&|pFNw#whn(bG<>C=~23vXgu-!%s z=Ye!^(gQ9W{7dQbzrXcfB*m}=`%(bIowtX1qj&41gNOG;b6DJe-mzc&=QRNMTxYz2 zYINn?ZRgbMBF4DLa_KxEBC!DxFJtg93!>!RSs?t=d5>4kNhL0xbNOmD7J97zwXvbj z_m41dK5vU+IU<-~B#~~h`|4%2L*f~<&cUcfv&!?wPM~xR54uX0Si|4?D-emUF0gO( zL|N53&x!uO`1Vo*d3HG@T8tE&-BDHyIvi}I^8eGfKS?6@sT~|xeP3iJ)iccYd%jUz z+JP4vbG1GgmAp#O2Did58O+4*aK=<);n8R>1cca!FyUE(0H?nOoZ26gYkpkajO#A> zYC75Wx-h_tf5$MrK(+O09m6PKd)Jc~#Noof4a)PJQ?ZjnpQV>XE(`f(j`(zx8qDTS?1={3{h$$sGMvXjfLfc@)cnkUq(I80L&Xakuv$ zV&R6_S*CE;;qp1@?h;Au7L`7V?KsSdEbKYSXdUv%ha~d3m5G;8W}Q%5R+bf6`ibgr z*0P8G^&HXd-AFr!+}shCAG7En$aiS!a?v6_cv>a}pN1On+;k6=Dc3r4H_s&IbAtU! zV~2H89KHXboUf6asnaHc4!lAmFPHHHyk2T7SJ_-$UER%%8#$Thd#V*v-<+YHCNbZd z{YDoS(e0F}wA@LCEVH}%t443X=@pNxtN!hv6R8btv~$9zBe5~%JT_HcKYK_!P$e?7 zQEY4+HeykyPk!T^sivzTN{wI!19Of zJ-1+}5O^zG{QV$7kVr-!X#5h}@ng)5;IA=l*?H}gbHxWA+6n2(;3TKhM-yXc6y0NH z?tS(!M)E;i#~+vmMTSJW$%MVyq;um?BaQZyfyYUU0Gd609l^&{IF$&c3Oc11aeBh% z?)9I22<1$(D^f6gJjqw9DK0hE45SK6U?=iStGY&6?;NETVd7&n%vA<+FsssoV?BP2 zdmJP%Y4wsSb3#Ip*ofeR}P};!X#AYp^~Mm4qPqVZ`T|vy z=vb$07nLbjTfCB9Z^$TSt4Z8R_a^CCfA3Ya>EhpcqD|PBHQQG-o^IS6$W+&9?(=RX zJnT1C_qqDU%fLG)F^(fYQqeT{xlCJBU}FMrMuvGmX?<(;g7HHQGUe@Nm**5N><_$U zO>15wnB5BSQo0xJT~XfYY3I)0328n}MOu#Y`QBVct=G;jTg#O{nSAYO1iz2n=t-0{ zY0r8jmNn|qZd=|o>g=|)wiHr!J!)Qb%;DX96OTjZe&nrU-ZR(fT%E2$e8>i>+^yk} zSgF}bHM;_>^`z0(BgFk{_)QajQiuYu8u^lmG;K^WZZc8+T`>ypP|0`x8rKIKN^};e z`b2@+PM`!y&ei;h4|18khFL>dEFO(SO%qX#8)5UWR&E#^_LWXGZx?kuy)21CBy$%1 zom(g-DV(dq!uhIHJjscB7=JXKx2ueTivl8RNxO^D;GSemiT<(e4mNY(^O_wJ(Dh0vB@e%9?o+VyTKmZ}TOHGKI3oJW8M*FSssUXnfTFUy~3{1>3>>hAH37k|JyjP;^8{NrfzKtM9c zNEmow`i@vgOGx~|fuAp^V8-I{{QSu+g~AG8n_`(6?duH8@jJ~VQ=6$C5B z1|DR?DP`N%*4BgYr?C5cd>88yJ;!KjCj$~w-T+G#Xk>qC zK3}z%nXRYXc*o#=51oFBadvW5U1*_Fk>!5?-P(#=Dzd+vb@TrHK83h5b zp>JmH^po5lqtLiN_CSMrj7Gz~=`RNk1-u*B%J3F&=99@oyj{0i)@8ZF)`#fQbsL-=e;Ff8iY33?8-@_4(X> zZkhl>3*D@KO95*aU`YM5Z%$vJxO~WDexc(m8)43YMX03mi#3$P>kpvZ76(CjW5~%dN&W6-u*-vPGlY(6PY7+J>Y${BH&Q z4gNL}#S}>*j9)*?cVu%<&(Tz1bvo0bpN_8mh~0y9!2$}%>pUDy)W5LV&l;Rq95eyR z0U~M~tMrM`u5w3`6b9l5T+#?@9RT9~a%n%W+;~aI@j&60@a>?mDO0&Lp8(_6?+A(= zgS_wJzW{C}(rNtVGm7C*2P9kR``B2V#(iUXQgMd;N;6+vA>`MHfPp1B&q!cg6qf6^ ze2O!-Qim;KCUIp=8`rc@)q-8a&FzTz!0`w~%y%WJ=3)*R2g?}s4N;x0$2&Q52sI+n z!t+->S^>SDUnEnvjsc&fw%E^H;)d{;{JTf%bMv1DTS##nt(9<(!kOxB{W$(kB9c!HZNEI!J_S-WEc_O~wx`BIma z0DS0N2x7hcYYAjJ8|y~Robj5twEAJgvr71+_-RrMe(~CIzh>+FqX@xA@oh{*;mk}2 z-rjyv_+mmPGistjuLg-7=NNbxM6j2-NQ=-HHXS&OUH3w7@)akh0_V9c#bY!tRfvRnJ_C&PfM#jRR3jOfa?mT|H0=3F&tjP zdq64_6Tbt9n)2e|rzu#zhQ0OBr!1|&!ng@&^Zj3foiB}(TE9413r=j1t`44?)8%tx z%~37PiJ*^>{6e(f66}|&%Jz;cZ`^kFaOdU&sjFt=!&5`$4$q#YEhL(U#T@v|@?-O> z^6lh=&!fLaw=()$xx0EYdL5w|n@YRv`b3e=>KO>3HV>e$yd}Ckv@h9jP+a4b@6R@>* zB~5X$Ip*~HTuzwRs0@YLg_F(2k~yplK0DYHO~bGAI82C$kTY{kqDsnHtB`4nV%s}f zAPNb1v&2+xAdRz%DR9d9<8UDki)6f=f^gL>yk^Fb+j+fXsc#JR@&Gx+h6Q4y1t9W6 zEL`?dyqdZ#7DV+aAHdm{+Z)kc$#Fi$D_J6NA0i?y)(69lZEHbsLf?;mg-U|OS9iajCo>a6IfJ&otv4bgRS$mJyl}6svPOsWSfII zy|U!@yf*8X0XiBAk&q^dI#Js$@m;4v-~LQb5r4>^t+H>~_n_PMtwmh&Zd)RmSOPzv zvUcbyo$Z2S*EWuET*L`40f)-E#p2v4g(&}n@-~OCp@u^By9yJ9kDw>8D2t4@a1IOA z`RV!@7>~!^086m_D+0FduDVw+HsdLXv4;HD?N`IY!^f}C1!%UmHx+9G&_7IZvxKjGh2XSG^Itd%p3W+c(KF@PgFkusK(%Fc~1%{SPm|% z%(o0QdKiMh`8pMRIl&~Yp;1iN0nM0fBg+^5pFt$ji+;rkc^AvM)zT8~TQshkrHEll z#xkW7o7S5qQI$k?ybtwuWHvRuIj>=AMLflQ+C7^3w3I?rg!6P-Ghch2m#+rm>wQWY z7tfoc`-L{xKO;5J$wD$)xpdWA4D{GFNZ^h%xIR)oSeSST&5vg=ywdB$QSjNb`|xeg zeI@Ohij_=-6qYISCc=cZXTi>$j+UJ_UWx=A7bodWTXe@O2CwBgG$4+KeC0=ynmdO0+9oR}Q!nO?KXTZ% zY2sooB57*Xl63S3hc%_A2HM#o*vu0NIUnjj{?UP%Czo=ZrGziU@$9hzyjrlx;$tzu z;-hXAqC(Q-g+)ZF`YMc^RvM+WBBdLL9`qi&XYEP}-XmSirWW_`sA(|(kZNZy(Juji zMFBaO4RI!Mr+)@)3GwJ^~*chf~aA+-T_S#G$iK)Hd>9-f_OuPkMa=58t&SPC;wX1DN$P%S$OKNshJ% z%)66W8XPGz-Ul-TD}#JfRgR6qH|OqD^1<(LOW?#?KE(wmuQfFEn>i{fDhdiVfS%Zv z>~TC$S|*LoW_^V^Rl$Xg%+bTfTWON!K_V%AOP~({;13kD`v3Gi9?yv3ta`gxkCu$) zL#1uTsXLrtD`aA#3nQoy0{13Tt(`5KyWWhQ9^R(eu^jnZurZl|?1kvG4ae~*+eheX zANBc)Hj@y%J3w9f`lwE}+YVN;#?21!vN*ro<6E2$&ghkkbi#Y}|4cPD2jxcD2}FS?dqo zl<4x9cb!}h=lzK}NVu4j-5jSbW1b$;Upa&=_TVN;;mkd)cF1I!sAg@fc7|FA6i=M^ z%Os4BZbj{ZdK0O*!Nw6k&lAVj;|GTedm+!nh)3t!Hex#R z_cx}#7LLAq2j5>Ob_2$*_ZcvRJ>P;logOyv80Qgfbf-&A=cuhQYR7{n_0(#S`tJP4 z4`;d2i$c{vZfWIhWCn4&CTNP{4nzb4LkQ@)Pd~@hHLt~w1f<^MJ*RzO-zUB^s4;lh zM7q(O9*ozM3`}I+f@I{}_UIA~(GqeiKY7agCVx_-ZAz*G7&Kh80eFXu>U#LY zJcC}wM+Gv_ckiZ2yJIHv}-3u=q5tU*7Un1s|J} zqq+9$A^oQ(U#gs~RX+;i(;W8HdFEFkXl$)SP=@lP!n8j<4L1Wjujxq9k65;@`K6iO z?{!-uZO@+vw*6@HnVs8NB}}J2A>e>P-S)JMHRu>@Pma7m7E%mtol}OK9d?LD@<}~i zvZczAB@c62)zV?NIE2?OKi_)k4#UMV3;fl;^ju7u1lF%mYN#X?r>v0a##6S?8tI%N zY0A)%>Hu$xat@dba8qefn5g+uUnBoi!`a1IV(j} z;1nncUkq6VqN=+1>CS?4o_w07HdSQt$8b0lrjF_DP3jtT8I!w>c$;e2L~yc#$n>bo=&YV5LAyJ~C-DAMQH z{*9N_IzWEZ3zOkO*aJ2mpd0r^M~#&Z4vn<*V5z70;FQMNt=Xh!k85aUUteGSklpKL z9$n9Mw*46Ra>8LL% zcvkVRd+Uu4S0TxGa!_?`LEWqWO1Z zw-Um9r``6B>)`9RhxR|VCG#o)%Y|_*#{ukv{@C8Ea8ZBxNh0cFyh0jM3_+0H#Uc$Z zWCH2kflqQD=q-i|3^u1pin@!9$DW33FHKa}(qj`Hl++g}-ic+;wH6S9ECZeg=&o#+ zGp;xC?xg54P^A`vR8wasj;4mgow3ynVFPVOmAss{D`XX7IGyn$dOCF;TrET8@t&2o zvv8sXii$(VE5EkuvuIP8T$Yp9`qRwEMEoJ01ktDJqO!3qf~|n5^&>XdR+^4o z&)KP?^oh5RZu8_=k9!a8CQEyq4DE*dXU@v+Gs zs>V5Y4v%dGEI1~Rj9l4xUnSs;unG>Xue1&vVmeNo7vs#w;L@9%M?o?4fBFQX*s#0w zg?I?EIDYsBpGCWPs?KGr+RN+qLYKEhl)#~Z552#ibP>X&C)s(9Tr_ww)p4qrikFh| zfbju3I6}n*aUut}8AmhzoFieHH(&nZC;+p6&fmP1f(lO4*cV+X;res>?BI~$I=bZM4)ap-O&%f>@#}}ag%DWl{hTCw8(;)rpSUlzZkLjqhmuzhzf*W zP!HE6<=og@t?EVaxoFrTmoKwLGq`TnUoqSBuY7<$l!u=oMM#2&CkqPI=G!vP6XZ>R z^?L&q==Y%}B|77k)ajJa4K|<|QK%4_tyKA9jk%k?6j!9W>o^c`jw=a{xrdXUHUp3WQ&vwR zXYHa?4c{wDSxP&^d`DXXr!GxfFM^AW%kx#bx~8TUB2{8_Y7U0xw?Odee6$Pt5i3Wn zMWFrQq^}GPDP}nKjy7vsDjZ2_8ELaC zzmeqW!{uIE)E>S%hII5mbbB%%y@iOEL+l*5Z!j>GR}8yw7>5>{Zsz|z`?bh`2qcyb`}BN{3*>af3r#!rpe zD+P80uYoL*B3=-f)t++Q#qcTFhs~Fiys4{8TWi~&P;uS*o|hpV)VWRF6i&kVaJ&#x zDprxl&fulz;qFr5-m+2|6NlMc(n-VxV}^IdTefAd)&q0l7=pnvGZxSyUKkYN`DJUk zCa<-_zk0YnBEnlcMLE}O1d-^mI^4|dNZUfN{c}FLj{Das(cf=uSonZKkhu!J-YY}7 z8sHFR=QZQH{`?JNZ^pa7<|LG?V&_ciWGi;WZ_g#k znf*fi3;!=yGeF&EdRF16H?4fF4bTydyqle;YI`#k_<$D%f76;K9%&N=-dYACurPk{ zJ0xBeja80aANe1#`fHb9dV{IoKP*xBd=_-|%@W)Dr#ph9I4N?G8I^jwjX~2`X7+ z2oF1tUY>VkufLM$Uy~AAJm!;E@jg@rMG40%48PS298|u2YUA7qc}IxYNiz02b**82 z0X=5sZz@-<*4D7GnP_gV1E_~Y)cmhd%Gn`$3 z>I(8AhTeBfrHhh$-5A(g^jKr*3=Z$U;i<-BPPTNtQ zIAK>Q96O?`#`zJ(h(iO{UY>BWm}nypr_a{#p&rED8mtOazuDmW-m9W zqQ&YCkib%aVK}ZB)kus8197wb8FO^S<4143SR1r0sh9?&QEq}; zVAwX@em5;&XDZ)JDQQT@yM0S#-F)al%OJa{e6GhKGlRyUX7P{wHb%|L?ez$n+)J0} zp1v)zP%9_$zdyuhM+LFjDeTlHFh-h&ZR8?E#$h{qquox^0a!4M?&6%jk0B`$M>%gk zupc+++Y}b|gw;y3M8Tn0*Zf`tYiyTnF<#+nGTpGaGY?!1Tk7~)`NX4z@Y(Rm_oxse zFYm}Qhi=!D#9JT8A`ix?(-hyaJ47Z^(?*J?^ej45ef|36?Ki=n$ohv(aB5zD&0O&! z&E0`&=UD*R1YYVpEmOX5^ChTy2+5wJ%Jb?nhpthgH|r3`^HI$zxh!e9lHdE^QT@bl zOwy=T9M?uuZ_`x??8;tmLgSW~bl;G~`3lMv%p}k}`>Zzirkro0&6Q$xBt&oBa3Fe+ z?~_ao&f%nur2P-&133D?Ds8?(AK4nflVfnq0ffpN6JIul8%Pv5f5;# zR;Uemja9CE<1jq_PnnM^d2es;uMOcHnJh?OC7{;-D)qUmR5uYC!e3?kP3-Y>vVlji zciF>6X!yx2;hHfifHq&AJdK*2ot1hJ9%%NH;e#Xv-9QaU{*)Y!JjvsCy$?NIAnlaZ z*p5Ene%cx7f33`IoyhGg)NQ)U$Vr@02iLff!Oo-;-|hY~P^K{x?LP_PEeEHIzZo%S zk8g-2OFg3qx;~Sa?;Sm)Q+D>0%DffdL#Z$#u95q;*F3_hQ_tNA=Zqfsq6W2yh>(y@ zpPGVZc4-y$H@E&Q?05y_2MmK@)RF3TkLu5oG563OlukB7<@Gt|BEQ- zqT=&d?v_$uXZF6V&onwN{hhZh2I7F+k9Wjd-?M&QlO%?HiYl}W!#hBoCNgwv3dL-L2(UX&7`1TY>lP(fGbI1-O1!xN(@oS1$k znx`FChc1N+ErX;5H41grTzUL)p1o#E2!#`Jb3$^fY*w7B^|^N7E;ms=EVgc2c?1vt zVvvyCg;UW-=&XX2G|akoPO1iFy41RP;WgkWyuSb-^$!P)Xe)#Q;apEpkTgfjUqwsn z58OQpyc;4AS_rS9ofOb}9uzp`3;)c0+YE7j_thweNpK0iOz4XSDg3w&<*%JfmqEez zQ&9GF&houVKx%-2TPAt@)yvRM5EY!O^Xsk852g{{(9{H99$=8(yoWsGgc7~|8geLK z(YU^%y+lVAlp{HofVlU#XBQ-KK>?GscI}VM1F|Xd29PcH3LJI;%#<4Hzpi}FAs#f6 zmEXg+q|?_Mtj2qTbdK!_Ct_metMO#-)YwEB#rx<>pwyiT^2`VYB(*n}aUeAKl+gGg z4mb9BOwZs9$~hW-8>%N2wo9!&A^j_60u;`#VWy%UC@4gQ;h?w+S*3R?7>7LFj_U9v zhXMbVR-RUr!i{ZQ(h79kPH$%y_b!Z4X3LWMiSpuCma-hk`H>~?$QgNzm1E0eFqoj4 z)4ZDkGTSv7SY4;`DvwmXgvTT_Xic=%w9}mhoq2g^ytg1xIJ%zDHyT%@dz98_%}-Zo z#IJl)Ts1kR52(+e4%p7(E))x$>NRFxrG$V+YwOJL4T^wd>MWcAzs9aPCVrcj=;m2Y zdUzlue^uqZkV-fisDX7JAw}>}K?6H2hG}kwB!kB!GW*klDzMQI0>?=1+a@2xSG178%>GBPx=DOf;cvl}Ue{2Or&A+)zI>+p?)!<^GFiMl`1-2n zuKpe{0+pXCnToi$xZu=RP(Nf=dg}DO>5@X!QnqvWLJ&ES>UV%O z1Ip=`jK@NAc2?pKQ+- zLI%X;{Bxh{NsXRr3ZZ@nMY-pfJZ2=nxOD-!nIYTjkwxRY0LR1pe9JqD*9A*Gr`v2r zO~A~E{dE%-x!~@XOLI2gSE#OJD$=?;JqEH&1|frhS0B8kcrtwSULA`{8M^;QBtaC1jJd#}^VNj}wTbi4@sm5P%bfJlyE_ zkIzN1sOSmnq}pEodOJ*qA2~C~^UP&mhLUaSfVn~b@qGN}9~C@cswck9Jo%9Xjvc%H z`-8u{`6fR4NFvz(W*KpdfQW?wILG7ndmi&6izAQ$cGl*5^cp;+wS+x5rfC`+=x`mw2MsAu&ySx1W5a--nmxP~;pR&u9U~NJ zlH_#8y}5>k@_FnXnG0At6qM#94Gyckt&MwVH-O&En*)N*w&SdjFCIS1^Sk#Tn?TUO zs#~KyU>HzZrf#|c>s48cEwa1*j60beb^1!()o6^62yQ*j2 z)Fa-KvAUr*O#qJX23C6d0#bE8kjd+pEz>%;RuBjtVrdHD0^9#a&;Sh;O|8|1IjfD6t7e4TA^rcO>@CBh`o8yJR1^>iMO35>Qb6fO5Jl6~+Ruf5`4_v$y6l8}&)kum7} z^SC`o#JR3Tt%}I^kQ8BcdRr9WJoAL7euz_`?{xu48UR<>uS5A=q5|n9_{l)A;kY7E z>>Z9~46Kb54@xl)xV420cT{^Vat6YwF> zUw+yHnepL-hTO6!CEbnrF3kS;@;|LTNAZ0gwl=eCj1axB;=$_;k~O!Ara(S+w;@Qb z0cbeMdtk~>ciq#Gin=j5t)>sS(@10NTpbPaL506aYaYCFU~bT7_pzM)*0R$c3!gJA zK$gmavh|_!+n{(LNf&HEnlwo1>m>w(INeUy^=elH;yzI21TnoPKL$CIHZCl_4>j&j z-GDi&{R6vPzy$jPEnz;rJBcAr9ZE3i;x~iL)-+0q$I#!^(c^&%R2?u?i^ZKRGnUiR z#;z`u7045xM3YxC2anG%2&=-B8>RjWzm}x^J+AKp!n?LeuBdF~_zDs7^`rNhjQxdG zmRj|r0=QJ;T`LcUEg$%}Iq24XER8ERTzGxN`uZPGb{Y|ouV}gE`7?>Cc!7lV^(=o= zDYTla)C@`}HpL`#xeTfeFyR6^j!G``=?^^Cy2dgHD1tIsKAjq~X_4dN1|D{p4Y zS1Th48sNI+a|c-&mpUkzm!^mY=Z|r4qmTM4CzHkj%U@~P7MVOO4TfHkN>c0L-;yGEyhgRIe#W1BHH!of4)AvIUfP?3C*03ZODCo z;Jx3T`LTr{pQC#2pi_rA3DpLu^)QbTd`|oBgG&o@dQNZM15?+{yYqm&!3LOginT`31Ee*jm_%|d z=41===!{d;g&R)KPHvyuCiMsTA|E)BdL18Hgbw+J*F@*cw-H!DPHhDw*8b;du;p_r zciaJN{{;|YXW9@I0ghK++8WL|1Dc?Rv3>KQ%3~6 zx-1(OuqY^8zGH;G7teH`@9)-2%GFQufTW*3zcTPa5#s7Hvv3uC$DL_u0UrQRlWnZ| zp$A8|XQjm}AVQu=#vWe%(|OOIdJ1hZC4fH!|4@{7LfVMZLoIViinKfq^zlrKdTezV z%4a-X`vPl(|C}(Ne%k*yJnBWygVxYjzw9Q`?Tz57UUSLoYvEDib{>$) z2o|UHOL$Z#zBRd{%X2_wUJ7&9`G2J@KJxuDmAI#=i-K;=@g%Do&<)AP7HrpVsgdit zeI^@`^1|X#(!)G~$aUetWL+LX!7*LdmLDLKknNujfSUyH1ppJ>y-L&QdieTVU!@Ae zCp1sHy1%)WQ$H)(lZLn2bt|ZR&k3r=mB}3;G@UVhcK95uf%B@J2Lrnf&%foF**3yApT> z|E~ru;uG|b7h-%2Vxv5+hW0`T+Nt-ag5*rn>q! zMNyFd-G_4+j?~vM2~O>nJ^NquBhmT4=to|u0u}$=@78LxU85$h@8Gw-kFi0#xqYC5IJq4z4_R zi>8aXnXY|}{#fJ*1*pT7Cz1@mrRh^9zI>nWxfUZ_Q`asKBUYW=`#!5tesL|hrv&9K zRa(#DB~Kq?v-NV_g5gDkPx6-E*#vYKrFL#0H~RS{8?jmkK!x(~>mPw|5v?$nI=?s$ zJx&{ax$ySH;phN3Zxk-(EY+?G))>Cc6BjdE7@%i_KkWGPk~y%RDi{(W-BSJ8h6e)L zz&pBUHh$Imu;?G|kLS{igo#!uFFSZyPBrQVpkU##BuKb-0L|stO??H?g@Pb(g+5;}qKYarL_(j+~ z>=43Nh?En@tPM$bTK3^|WXZUF&)|5Ai~&di&vmy?0J%oDgKq@h2Nc>mwcsW}TXVo+ zI@5yr_3E8_uKN%%5ud@jF~0raFWI#XQ9t%(`rirw2C2|cu2@OHR`o}{2JN*r0Ne73 zPqsl4~v_Sm}AY=r%%krU`j`5Z;utE>qeuXGU~ogIoOj^q%9 zo7Fe8j&lJCaw5nrxJ&tuZjQQxBE5AciP4JPTZVbeW-I+0{G$0Mp|0lH;U-+5Hc(J| zDFC*7lThDuIRc*9axu?oma|f-oIyY)Bw?!R&uV2!HeZxAhKWUev~ZM0tmMI$wKYejs6eHA2D|_ z8o9r{*grp5@WDYq+>(i-&@ks8`3lB>@csJG5%w=;9*hFN+o99NeODYKT2Q*ICWu<{ba&6Yp3U zRast1M>1k$C%1Z)NvMSP*lXciM>d>4o~r6-TjF*H*+Hq5)n`Bpv1-^2I3|w02^#yL zI*9RFR5sL4v!b|dovysv9(fzY1p#8Kd!+mgKxI^})_5WjwZ-GVIA87v;bCXc;G%M; z-X?n-YrlN?7?6KtsX$R}EE>pCol5;k3y))-uS1#V&JbW+ZOEq3rDT`9K>m_so$m(8 z$Z6I15#j&Ad@;ZgTI0en9lxhZ8Gvm>B1YPZ{M8^4W=EuxaS1O20|TJtFy-9%y;DBX z-NObk?5^bG32DapU9c(+ki8s@QfmP5h<>mzyPGNAfAAG6JBCGPHZ$FrmX@Hcrao0n zY;$$|i?I@=P_6aRzJTTSvW~~vFBpgE3$XI+ae^Q!skvIe*zblY#<3l_g8;}w;E9t@ zmkG^P?%+BQXV!5y7%5XOECL(k7+THuHZW`dL}&$npNcYuP-^2Z0BKNGK z_#Mq79;Ev`M&?O4)HQ2tlNJ zX-Pu{N_jt5A)U83y;LCy0Rofb?k=nD&i##GMJqBD5UnOHraeh6c{PP=zPMl;Xzt7< ziva;P8>K=)$_LYfses~`LHwn9ZE?!8l#hP45Hl8DyQ5YbK-{KT#Vwr{!?Mb2J*@zi z)-*<4E#kn4gtjSRjC(nqMf=EIv6g()`TM{U+3c{sLeW@@n?Gd||7r|W$zz4bnU}Yv z5mUHkdj&~;17>;TT)Q%gh*RzZ$mfIgrsq?eLq!@(1G$;}OI<>xq$k`KyvU#OjoeSX zFO`InXeBNdRSg{yZj`T4o)F9>3jJc&9y ztVrPHVD(UWG`)YL=CDu$AbWyt-ptV+9W^ayj9Lq15xm1Rhugv$CG@?G)Z)jd0MtkX zt=O0Wbk<^09Ny9!N=|B=5IsYOYti#<%AeAYeoh&jGWxj}gkTqNgRK&^clR#V6^ev@ z>Wx_^m2746j<}pWaNsi6v%*YG4C#gQi)k(*SfTE=lRsQ)v&IE^=OATGyo@T32R~Ao z4xwfuWH{XxGt=TZbJ=qaT9xjNf&S1PWr(%mgIn2TGH0&e_-*TH;YR=K52SWlosu$!GK|F{f7qI(pD`)K;K zIMAs|&TZuNsaT_whC=YLa|z0|)Sdcy3sX4jXSRx&z^kvt$k7A8P#G;QF1QD7($B`G zl3}icO{3*WIhfqt(!oJto@CVuTb!9q_csy5VR4;}rBq!}#ys3*F|HiUv~!8qvsGBQ zbemHRI;)a5`lZI9lFg;9Hd3qIA8WNstR|x?q-Rx}pCy?)>-ih0e%g$(x}=7Z5yG>8 zcQQW6NsiiM(YmdYF0K4^JV&i$YY7Wx6k*qk=_E?gRqhvk4O7BI-V?dq@qW%en4fdC zgM3u`3iQXT4WcLbAzG-0^kWz;b`OQqmm2mrPEdY2-=j)Mo zbJaNzC2c1iTrP-XmTQqsdIjgs*&H9*93=7`0Qko%7x^$a=3ATJu8rq}?Zum$1;mWe za#_7WK^}SPG7_Y@ihlBQ@d8-Xa5M~yVcZ?DF;f00HkPO>nKHtRf?fc>a*Vwyg%l2V~ZND-N8mO=80 zZ1L+G2r3!pZmJG511OKYGs4l#3$j;Tz2xJ!{YZP6N~fXWz+q9bY5zvo_9Z5!a!+}0 z6$e|=TiTzGW@)u~x$RpP=D#3f@;}-UBs;V)LZRxwC^u;!vA_OpB5nfUmnUKVA|kp0|aOl{}@O~ zgNPnZN8}t(H7gX3qNh%HQHfp{s&ycngsAl((n|v1b&x^^FkaHCweBy$J#H;%wTH^< zh*Oeye#?5YUeyQzt3Dij(?8p_8o8vAuQ)wNFZ_`e5^SA)o-(;YAsyDSRj#a@sa|3& z7K-1oRL^9W$gc%D49`2MB&ZL+UsI^+^f@SK-DSJVyz1TExz|`%3=55wxEMFO;^(v~ z?Vk$OGP73O8r@`p1|+);wEUXhJwbWz7<=2L9S7rwyAy-o(=vrS1yF7PFT1*;12cUM zSh`xl%Ff1$rKZLVf20~IE7&>O(Lrst@9SJ){jAv?wlR~E67A^Z)QUTES9magGLS}U zF^X9$OFHp(VUlkQRr5Xp<=A)!;F8z}Wk!;ypwE9h2q%JY>N^vm__zBMvqV+k@&J%W z!Gp@QnQNr;xaP4v1bF_U%!m6L#!oInM60Z(Qd=Sr&mYnEAGbznfu(=_`jX&bz6(e< zSnO-5SbxhZqpDhagDES9Nn`OnEq|%G+mPgfqe~Ch_$J6?xgEt=`8?u1DD7j^ra>2m z=J>-|t==h4HyEgu2*a64xxm;}QYa3<=y~$rpE~6Qv7Y?+EK^oq-kSYDy-!VF zHZLh`D2{zStb=xxn1rWp`SsM#OG0~JfC{p+>XEIOd95*!?~31!-hSQN(j1F*Q=oo? zZ^YfcXS?v39?n@|)P=N=DvG>44RXM`5LrqJw9>Ar@1S3d_-vnR0m?Z>wuBJ8$RvV< zaFngB3|$d4RXs6*%jdK^+0G{9U{bHfU-{+XyN6v|Zzd7F>+GbIqg_41x5}S4 zvKNjiEUM&Y-j%(I=)QP9q}!f=>~r2R8MqP8A_X|-gY~&^yzJ*kGO`j+DyoySj^A=D zAoEd3r7pa?K;`*uVRkRm2=Bg$S5C?!I1aBrep^yZ=dobZ|E24XmlGhH*ehQUev3BH z4$a*;Oe9q|KEUdQ6}zcv6=n0<0&mSfZOY(yS=G8aqMG`)2>=*pD-tdhg!1&1;v2v_WZ#!=D$=CNE!to}T{Eqeq`Ve?I4V6n2$y zO8J?EAafSitt*)a1fn}mI@{fk>?)VMYo~SiZ#=KGvx4tTrDu!7l$F-yJQfXt2!aq2 zX#ixUpr8QRU}Iz3`;Pm2bkiRNsi~Kk78=$Rb-fR~(96%)>6vi16cOwJ~v2We&( z3sc+4nm-KLU&6RAD4VAfh96@-hxyw*!I##%Fx{*jQ1J2NExeZmx;d@yqO2HGxhQY; z8|dq|7EPC*BzUsf9XpI!{ro>)_3-HB#Jl?Q=g%+mc$Y6$mSuYQIQ8-auSHi2lqML*{e+rL`_g)JJV7{M`2A#;hP! zl2TUZ4XsE9YA@yN;Ly;}ckdn^&uA`=rP%jBTZ`;hkI%8hxx@{`W5k=p zcdk*QeP!s-m!%IOH2_v^s#9dT1P&{a#>coa`4h#*gfDg!C z2RRRD*9}G(0qvn8iO^yPh4dF%%6$%9A}ve|I{D{>mruj0&Hg6)Y;sDK+Jg41h{^@w0|1c(J;;~j% zuu7|HIurw+ zCYc2&1V*B~}R4sYeg2lc!l_!Rzlit_PHP@84|B9eu;XYLc#TSfvV9&OkyD2WD5U z>Y5VVdN^^;v85|*T;<3)b%hNTw7i(XJ==9{j{M1t#s+cYL}gGbjQy2 z%{8Cp@m5<6idLV(|Pho;7ZgFwK zYno{U$aLsOASy+CFDLTaflS8n$aD`>3vG#UKWcL`@_sv@UG3F8OmvBWs+QH zV-5KE`Dfc>Q>#|?Hs)M_<^hH_+Y!HtD3wuL{`o!7q%UJ~W~Rhy>e-S_o>rx@J}82V zc*V*V;Ka0FuoLyQP6pq2pcPRGQONAFpwvmHY@&N$n&)_(aWo?45<>H{@xMk=bu z;&(Hd)r$>3m=)X(PP3Y79GD~Hwarkxs$9_=?eM{d`Wsep3gw;VA5q<7Ici0Cq?{k| z?TV%nbZ{?jnfwZYe=@v}tzSD@s3);8+tC(De>XJpfm+Y*#4gA-bY9H-a90deAn#x9 z$x_U;UjO-2qD>{OgSp98TCEcz5u+;-Jv%#_)^52mD-~PjvXi$^TH{xmmwm`eicJjolxKN ze>T%gN*u447!^s$AhIVq<<{SAlN8dW_xAQm%n;jAD%u4AH0^D{gBkN4q;RR~gbaC` z?io9WnN^=}DN(d^Z~%9kvUA-w50jl-(%94PPy$`E%dzTcVd3E zP1rD*JoaEZ^RZ$KN_kC8P8Pgr!so%JVABsmO{`5eKr+fe!qUoU<;+Xoxt{&4B`&L< zqM4T?#cp#Lf9W;^JBHI_;K5Q`rebDzA3@`ScD1uqvY?i>TWchJ8twkDHpAw8_taFV zQ1pGtcuuL<$*9^yA&=voaStw+xMHM@`eF)yHLk`_Zvd9azg+%KvGZil`1%p(9Nwyp z8Iv#Q@Mw{N4A}0H(AUj8oqHvg$+b7YF-I$o%k;w7*w`8BCB_|b90pxUx%|8RDmomd zgFpKFx4?l43kx(2x6}Iy#qcF$T8UNlR!&raP>A zIPn~Ok_}Y7hAAdaczHn*v-dNwgFI$5PB+`m#_!FBaq?E9)SOAZI0H=QQ=2W6?YcooPZf3z|YW?Bp~m!Uv7g_?Q93=Zd^BJgSS`k zTJ7VWbjGL_8@dMHrYLO3Bdy6t^-N9%Qn*OA*0P53%M0DpHSgRyIogo|=Hk*%`DaeV zMIM}8Yv(Rr-_bA13GYa1|Kwvli%F_P}>h$}iNviq>Lqh}O9J5|fJ2F7?6B@8L_KR%Gw z2czV6T$}Lr*jlPh;wezn8w52ct@IIjZzDErS1%NK=+?M6oNOaafU6|Eoz2!zC{rYO ze0zE5)&Az^7#3QYSQHGk8KbkZvMO{i+fgmARchGH4YwiGZ@*=~wP;?lztI@Hr0Qhk zu+z$C@0u7}zPh{q-06Kn;35KUQ11goh(v+=&*~wxh`9)>B6m`VpsA*BYlkJ>?l^xl zhLzQArjj(#^beaDS=(j|E?--LK?aVmf|*yeB9nAn2|HRtg+|BZsGo zBC&Hnzu(hKYK>w9u6YHTY2ld#3W_MH6XM)pY49 zp2_U8+0KRVNO5IDq>_s&t&xl~cDX*--6a*Xyy13j1M45KfEt2UL{KgHuVM-c!>ctq z?U^SJL{}a3{F0-YJ9WG!W@gIWtCJ(WNJv`rp}6^q9Tk--f^NKmlsm$*X&p7i>%}J! z2_aL&_PUDp!YheU=9%p41*{;IpJJ))Q9M<%N^jQSw##Mp7le{AMx4SA5iUF9b56s& z<>vbdI9^O{dNAY()skeaU9O*Kmqv}7-jA0Rq3}Y@VQ^r; zRrR7=@{>FSC0~m8NcGVTRDnN8=DzB-6m1V@_MrE+Vb0PKN`6l?4v6YQQO{GK5oGo! z>r__&;&+E!mL}UR;9P9ZM(@Oh*gfEn;(>o#W|g8W9&_1SvME}MT}p`FA1@9BLMgu* z`rtbjwU+-SOw0_$tliz+zCK0ouP>dSVB${#AhD-Yc#;;&rD{ z5lgjx$EF`%wWd}Z|F&vFbMtroRh@ajm}$Ep->NVG-O1LfbOvG}@TN?D4OwqLD1+@? zy3-4289MA89s{1L1oVb)8c6H<#&DY%E+iX`+-K=aoOq$$zZU?jhyC-*Tzv!>J z<$c3qhzdA0?9qxG9YZCo?+EJ-UzO5xYOi&D7bbOfK})L0`s1(cE~_@ZyHxk?-J7t` z{Qjqjol4F`>6q{;#;Q?L){{_)8**=Z0~B39C3uE~hu?X_fo4JcEpv7)?9JrEuip8w z{77(u5dW7Xa{N@uN#BBf)o}Y3r6o4jufFroYmri)@1U5v@ONcgE6|#{R{Vkt4tScB za-@qJ>z5W8Hsh1bFeO$!q^04E`mFPWQ_aJ?bW}rZ#x@IF;DBZR$Jjj!@^hGs6huaNIsw z?o}T7moj-U4&_XVl5YHXgQ;thPkb`U+n%JeU?-u}Ccc#Sc+qsFc3&=&I%>Z~_O7n3 z4ixy-Ty>MBnj!;3MuROa?TseiFNWtfwApjdP%>7+hyq2}^@dx%`L2c{YQj!@vFYhS z@y^@9G+?^v=;#RP$XF|MgIb1x^PexAnG30!x)f*D!k%i#BT-OSrzwOyqJkM%Z0G*_ZvNkQ zwc(ov^=&SUQA(;#tFGf3(wEDFg05uF0|PbEp7)bw^L+Rtn!8es@l5N&UqftuDKB+S zW;-MhhV`7~VYHP|N_J}w1LGo&$T_s3#vNX~_N|#AQ4fAYyXNjr^aaq`C#qkzia4IP9>CHZ_)q?Fo_=uUII<7CMSNSgR7( zyeFPdCA_+FcFfnt58lLYPhbUd^{u9k*apvLFy*a#J|Q~P4P;yH@BqlPu^$;MQ(rO& zrqV7DxHUh#1gr+IbdD3XV?=yc+A3sRVqtuUDx2kwp<37C<;@X`&OKjVSGM+@@twn4 zduJ`TtV`$}(3ChUz)DQlJD9Y2$P>SbO?XiHNH-J-V&|DT)U0)O$J$8eX(ip5IT1@~ z9YZV(*11`PL_zV_v~>dy4`*11cwdZ;U$!xe;Rs58dFukzsXfb2D3<0%UyFE?W!EpG z?e5T)bb?qH=pczRH5vLDPr4P8o80FnslR{U;Dk!@DmmUV4JF~(r!>jom9CniFY6ojq1>P>#86iO5IrN5~Iok6TI{ANn2@(pL+uLm` z2FIn*TXv4t16Y&kkid)vWQ7)gnP6eX2L79~Q|~FF3-k|n2qhW!`~6rNrn_{U^&MO} z+EBMjd0XDZ4Mm+3SHec=wDS(0_Ir++q8tWfjlZObuZc}j3|HTZ6U!^g?`$1|oy=Dv zt6lq+A{68kGWLfLzqOIx-a1^`c5k(4GBn`y<(~Y#qI9CQxq&g`bLn&LPlhDf%-)YR zC_)f91adqtwzlFKmb}Am%*ae4pEG^#vP{4Ix4o{3%Qo4&^Ejxs(YJE?5kwk(4D0ne7CYtP_rO3 zR7OgQcGEee~~LD_aI+Rx`_nj-dLp67aTtE%yth3aLN7 zK=YB_MxzUHtnUif=3%HtwI5TUQe(R7))O zkp$=B#1;o+R0|oPjA2WM%G$+-`b2x$lT|Q<g2IvBRx56 z!OQ$uJdnM{7k-bi${0Zp=wh%Ws|`iKC24S(-H>wfKoWa0?}Uje~A>sBu4l3owuC zn~w^l>|n&^3g-?pHc89hK&kw!kr^YAgFcfg4Tu_Zu7Q_y@D(MlzT8h)nqxLsZCg_H zFh7rO6>^U&khP5AC6QV#3yk=whJ70(>vQ}b?ml+vHlT5q_ z*_{A<O zB(OX({9%)=EY-1B2AMI?Um@i%&aZTt0Sb^AkE9?p(dZgBRRxYBUzOwIiCjs~0QHm9`a{!H+$Rt%Rp<*3El zk0iIKq5a|wPM;={+Qj2wzAhnza!u4giJplbb)Z=AmRr_Lrri=tBTd!*$i@^_?dt_H z%*P;k*v82QW)_sw8>V{dIrsVZBs;|st+ue4 z_?j5a1Kq}@CxCVCX)yF)o%BgAf46OmvU1%rX))_bM$H8_&XwK$)P)(XwpRN`(n3dT zcLypHF((Cyd`Q^*>=sbVtF;#4vs1mOsZf7XN~l_qvFczw;K!p7qS!;VjoKVVR9>Gf z^hpTSP(N;uU2UB*GKvolWt0G?stRh64Uiz!mam$!q1(5Fj{6hZa+EueeVKUrDh)lo zMVEASD4hnt1~=%q-JLf|RNgtBE|cDyJ{3F;1BMy(s}8{P1c;+}05Sx?If~qgy&(=} zMjSx-JG%%1482|sf05kEQr3z^s@MggQZJPlpWM2hHq!i9?RhWnmiuLVF6oo2#5(Yz zG{|t7W+ZonvP)ftwGhs?D3EKIeA1Iu)0D)d9+0-${v$kkw*W0i=00NhMt zrWPIhz5j!D*xT9<)*oM(X=?!mqX&pEC02)fjQ@E1u7mS^z#*KNY`mac5MHoGrv6(a^&%toEX;>;ABYCLatc7h1=g?8$aHO|=(us`+V>la>m}vBI?dbK)^4w1@FhrP1CHg-Y~;r*x) z(T3|(QBJwbFxV)fK7e$b5PIKGfMJ)1<27`%WX8{TDCuLm4PoI^0yq~fX*yWu&Mmy* z!g~F5CObZydJvV3K2|q3m4DjQ$ebo}_Qc29Kvd5y zjy=<0;qpZY-mQsB-4=iS@eoM52E(XXg3bUZ)yqy`Ox4;LoB$g-GFg9#P@@yA=$t8C z`RSOS{VK)!(s9{JSB~$qJ3nsfj1Cwxf>!5_V13qT<7|~VdL}cEvSrNNfLIUY;V+No zrLM0VD~&|{9Je)E5=>f>%S7c7>mb6Z+8y08ssw8XWf{=jSOY&uW7EOr zg6+iP`Cw9)H*aG4bJUdfrfTK4=MHtQBqdSE$w{!X79QVh2xPG@sX0Uofu#W&CbJ^; zSYfqfV4+7^%qRgr$WSRunS+Ie#L3R9j<^MbF7v1%n7LNsQZG}~3TLNIs-Kwd^v`&% zog{t|J|HJiy{6>hCeZNTuIfQhSJtXVC&GU*npr6;#2^^d3jQJA?PAlmGFiP}u9Tx# ztDk#+H8^4~%HuLsm$7=fHjfJ2AdbJ#edYCxvWF6bB70@=J04tJ9*vvO_L4fzBP>Fl zkmLlDw`L1ZZKMbb1?0w-4~}#;ZEa#`h+2v9%I#<6Jd~Ei;w4o0CqM8*%7?+dL6hIHDNhz0DvR)Ts3vXAzwTyAbN3Za8M6* z9jyVxbIE55mDadsHi`w^nqSc_7xo?ow+_dV&Xu-kDbwL%?Q-k9F8&Qe+fDf*Xo-PL znBpujnYY$5mX{F52pGozI3c#BR>#2CX4vyZ8x*tyu0iUJ__3Yt;wRQfrBhrk2!wd1 zBvFlc@EuAS9>tYV)Zunt&VsnjvNYVaoCVd9)w|3CUje{sS?U)CaduO(cWy1 zrf0xc_JflqlyaTzxOQH#Y$~nXG<0B^t6gUV8T!63gJwmgmH`mw#@BHs5yDPCRFSvf z(?+nF$3DYo3wCq8XM6SJ$d-*55wcx~p+9 z)Na~c=V0q7CU-IjxpHNzkV9JDc|lp%jgXl_Q>GvJ9ND}qllWK3%>!5%QAO0y#)zS@ z{_x2$o`J4cT2*PqRRb#IsRy=qXi_{{r5F=REdOV zB-vooop%-QYh=gd%4I4ZA0G<~3xkYnK&S0oJ@!EeyGS`(qU2bC^eGZ~i^=K7bK1Wa z!O}K_vA_<%@8;cIo?QFRhRj&s_k|0}XZM4+s=PA@VS_R%kS||A(F3d#s=g&+4NwT? zrh}Y2U-K4jC)eNp0INJ(0yZ{qIKvljFDJrhbMR_+B&Q;~7a{KelBww{zHEk{e*glT zAi@a}px1@4A?g+;CnzMN3tc_3^k{9!pp*a4=MwGr9J0c5OZjZ0TE0><2 zZVE3JA{$y^LV~U)#>3GV1XTe$guXb=cSoQzaI#~d%Vqv_jreUZF&e7f*7XnzYm_{ zV^%h{-rin2UM+EP@y&V`xE)DW)_kI!sI<*|+V+~#)~(OSeO;=yFLL#&3^m1BVng1) z|FBM#D$?i|aX1rAa31K5XD)B2>_lLpO+V&uwosbb+Sa;iF4v{@#4jnftW`V33f4WJP6556WD$);_<=oi|l^DTp}hadg~;Pw&~~F z5N1d5fg^->dONPIa$%I7M%#@TB(sK2qi64c7XG2I<*rrdpF#a7i2#-Dq@<*PK?K5v z0RiD}tb$V&x-=8)Mh|kn-!{0fp$1$ARzoX=TY6WxZAND@C^7vh=l-k?mKI21)&)B^ zQJU6&0yt@`A0$o8bF%EI+na&txRm98yAT1_!r%K(7k0Z^(}A7Ga~agj<`r|Ga;#n+ z#`nvC^Dqvz7Ut>&3N@!lgwX9$J?uofq0Jwldt-QNkzR&>El@M)OJ36uQ+e2V4jmqT zxdYOSDq$xd9|;cr$e^V@4?H#jtk*E@g)lJkO%7N{ro6r1(4N@2oUFXPi+vS1dt`M# zP~{&0SF0AnOp#&RhSTRJ0ZXM<{A}%FxNUN#BE9u0>1n9oKDuY@Am=;3VpFj z_d`NgIhd;oLmv$7UVFuioFTT;F`PfFf)Pn5^}IFUT={vI;Fqs;I-;O3h1e7^m|<}I z^TK`6ceXp9NpY(dqI0MIKGVB3eiBv)N)~7pI;75ICaPXfVeW5KjeF)INQoTMinAXa zF*6&B7IRmx;2y`RKPI4QyEY&o$USmkCiQd`yDRVNiq^3J-se$(ao`n#NApr+Nfex^Ffx+7Am@qJ}zLf`lp@=JH$*9Ny*Ajjszohkpk5 zB0vxcnk%4EEJ}BS-ervNI|MHugJiNu6w7a!79VORasppGLT!nM>x10OZa|s#2>vBWzO@KLB87;5; zx--h}i!JU7sIjeSup3X7=>ny2Wh55`X@-kj8bh5qv3`XUFs?s^_jWEQ`^@lzg*x-C z>%DpR@;0RpAjddlT4L}!jY8ZPmeuS{=9x?fKJLq;e|STd<3Zu#7h}X@D|kGPyH+hJ z{NJ%!ae#Z6zA#~PbJG-@%3q6EC;WWKLgL+H0BQL(cxkED|9{jWnNL0TZV$UT8XD65 z&5x{;d`CE78U`pm;Hdi^9v*N0LmlFO?an>-s~{b2N&f*h03BVq5(x#j>o+j`JHEdG zF0O2T*jQfhwn@?Dxol7+2^0c#1-OU1fXU-J{XOi9==dK*Rqzq>q#Aril|M`R`Q!U**CPq1h^gP!}`$J(c5vO<^(-;{q zjJj~|H!;l??}^;UyNNNl6EH9aLD_KoJ&emSzs1&RswBG5jg}zrLk1Dxha?`tkw<4} zle$Svj5>E9w%V{{RU>yN#;U!0A_=MW8gGtJ;hkn&RzC5nW2UV0euqAN?c%e?Dqh8s&F;E+Yo;wq~n;kr&>Z(k8-cspyX%nol)1F8bQAD@o$IUv4%u^ z%8sW>@sr3V$G$t(co+CaEoB^tJelh?#2=S_pwuFE3fn7vB#MY@rLA}&q2uY{;EgwG5Vc32eMD-=H0dPaBR#h_*+9W2I;em;Xu9v0G7GQ z28A=d>qaV19H+v#YQ#JC^HjLQ)Ev%`i{fVq) zzkVP_Lx?bz(Ltts6uHuU^`<@~x$>CHj^_&luGL4!^8htjiKYJ(qv<9~tAi_HkDcM; z$DU~KE+VpO9>FAd_fFQXx!K+`$MTRkd9zO1^g=LgV3UCRj*!w&bNB^N*|zG}_ug`K z9mkqo0=ELp@|EDv3re9^zkdA+^j~UfY5SIcxtVuao> z7|)ZXl(}vpy;|HZxFKk=UffX3GH>AfVJa3ZRkEU`33& z8sJ2F{F$#>51OT#eK#w(!ee2}8J+cwF?bUUiq3WpMkz~d0vRu3Yrq836(}6IV8*WV zg34CzP3J1WImH0gtUyoD@L#|)hMv#s0pQTn_3{v5qQQw^H0Yc1@^>5|C2++|*}^cX z1@?YGspIGGfFIA=L9a?mwu4Y|!oJ6ehgVpVD;?PNd@2?{o`PrcSO9 zYMm8M1<;K>k>-^yYl^(+i9VY!33bPmG+R<|rP9{OW`kLEDyPYD6^kufJF@!cu{IDi zLhruT;$CoA7zP;HOR8OibH`(nzyLi3?tobnv2{eXc7)cZFbuQ~vO^Gy(I)dZ(7nBj zo)uBh+dJ92$1K8@MqL8e!CNoW(d%h>-;T1VMw#G^SKZnF)}H<^#Fw^_f75II+NG7x z0Lq8~3dDk1FiYUfeK6Nj{H~r0UFQej!?Pdiy3Znl-^!6#0n&&SK6Ry>{=9P$fqm*P zSVs8|XfIp>PlEYKXr@FIV>L3rH~%41Zb03!4XatNJDNKBLO}8yJOQ@sM1H54P=3dO z7QJ^rGm<@V441U}-g57!{HX3pVT}M;4RWFZj{sLDd6WCq*Etb7&XT>ad|Dxio0EPw zpa~(;uVvBmfqBj|6;OxSDcj3V?XB%mU-kD0++DcBl#E=n=<#0^(~dkI8AE!vZ%PC) zgLc7&vcv@8-BAup>j5Ed$AveA-seqU6&-;t)C*LS4^=h1e5ecl=mbD~427LS|X5rUVbAz?Xm-g;r@!pJXQCTyEt(mneSA2cm~z&oXu1QP;ysmAC>lo@7{g+axYQ32_#3KuIx*!;}IK0Fdw^cQt~oG7Udk( zZnZT1MtINUxwhIXcz8eykpNa0V0|J3=L>y=A7;&9sI>(0zaA5m7eWJwpW2ns`uq zE6!P-K3{0}5EuH4pt3C5ZFfB#BHCql{2}=PDnqUaxav_zR7a({15)j?b(A$Rqrkz) z#<5EX$()?3PTc;y#i)N3y$n4N)%OrPoiF^g|`kR8ZE-oy|qFz@KAEF_Fd(Ryc>^;Pl^Wi^kH{st$qb_B*H zd+?@d?1FRkR$e%E@nVlw(tYABiB*rm&CKZXu2qB?vQu%Y)GnT%mG7bH=x2qm<0qSk zj23nz?BAAkj+vOu*5+}ThV513ZP&8NvEl>}ou*z?9JPn+b5q{Wf+OZyLxZ}=H^&^8 z+KP+u!MQ#=y0$nX@5Apa`EQBiSvcQ#jJy;aZ6VGn?NM}KS)#7pYC68mtDAr?uqi>J z0@s|=NzC9lSfoCXaqPk z=!pRP=InU)-pA(Y5|&+@W6jG0KKn@JD$(~pQ$%r!AT>-I{wIdAvpR$|nNCaF#0*ND zX~e9xL*GHYXQP`p@1w%r^Cx&vx)63#nK$&Z`+pwV#J@ z$pDZF?GBc%-kD>;9dY1a_Ah>n3~iMOc6cEh2Z87%vgAK{QI7WoazGXIu?Qr0Qhn4& zgV>?B;pr`%E|~bm1fr6at#@Qw9T_D^*^NhtkbJpPymV6f`J1)3vK$rTi+hBo zJG77JaOPZZ{Uh^3`R>P?;RQkQOUdn)AyzYJMjVBpM zO3B{;Z#^+fv~F4{D1du2pnB!a$-A?+1Dn)`@6B`1{hu)F*U9T4yj2qwfUy9o3;ue( z_|>?}u^zbh=fa`ye}I~!zCN;mMf?VPU>ht_+{oqavP1e`W-Ii>s>QC2+tvnb0Dz0O zH|O7ahab`^5BomPqcwgSo|8uijT#e2>hpH~1! zm0zW1hoJ8uN@sx-e+{+3%@AHbckN!?vQ^$$%4L?m(VwyGZ{+3tO{*fbL)*=-E_@xL z73$&9@lut?YObELa`Q%D0|ivPd^?h0pi(cmMDKQ7=(fFScf^4;sWQ_E;9&l2VA;M( zYu~Z@`x!RBA#2r)WWn`^R_(l*Gf&vCFXo-alV-`G>>KI>*&lz4D8hfFTBM^UgnW<_vHo z4b<|TB((R#>C)J?Ra#Z)Qx9&jG27a{e${b~ugmL>y-_eRxl-}IxuJmpJVGfqt!~=A zXDpy>4vg{*zXcTa^!!%cSJ)rgc=|a%3saB zUci1WSP!45QpVO>z%`7(6YU?*zp#{n&*h%ye$V%w@?QR4^+7oz9l0P0Pgg&ebxsLQ037xL8UO$Q literal 0 HcmV?d00001 From ba40aa13ad6b20b88bb57cda098420c07360379f Mon Sep 17 00:00:00 2001 From: Chris Sidebottom Date: Fri, 13 Aug 2021 14:04:08 +0100 Subject: [PATCH 2/5] Update signatures and integration points for Target Hooks --- ...oks.md => 0010-additional-target-hooks.md} | 19 ++++++++++++++----- 1 file changed, 14 insertions(+), 5 deletions(-) rename rfcs/{000x-additional-target-hooks.md => 0010-additional-target-hooks.md} (91%) diff --git a/rfcs/000x-additional-target-hooks.md b/rfcs/0010-additional-target-hooks.md similarity index 91% rename from rfcs/000x-additional-target-hooks.md rename to rfcs/0010-additional-target-hooks.md index b7791622..8b9cd1e5 100644 --- a/rfcs/000x-additional-target-hooks.md +++ b/rfcs/0010-additional-target-hooks.md @@ -1,3 +1,7 @@ +Feature Name: additional-target-hooks +Start Date: 2021-07-14 +RFC PR: apache/tvm-rfcs#10 +GitHub Issue: apache/tvm#8589 # Summary [summary]: #summary @@ -54,7 +58,7 @@ TVM_REGISTER_TARGET_KIND("ethos-n", kDLCPU) With this change, this path splits, depending on whether you wanted to generate a full `Module` or introduce some specific TIR nodes into the code generation flow; the addition of the `relay_to_tir` hook allows you to write trivial external TIR generators such as calling out to a third party library: ```python @tvm.register_func("target.woofles.lowering") -def tir_generator(relay_func): +def tir_generator(ir_module, relay_func): """A simple TIR generator for testing""" ib = tvm.tir.ir_builder.create() A = tvm.tir.decl_buffer(shape=(8,8,), dtype=relay_func.params[0].checked_type.dtype) @@ -65,15 +69,20 @@ def tir_generator(relay_func): ) prim_func = tvm.tir.PrimFunc([A, B, C], ib.get()) - ir = tvm.lower(prim_func, name=relay_func.attrs["global_symbol"]) + new_module = tvm.lower(prim_func, name=relay_func.attrs["global_symbol"]) - return ir + return new_module, GlobalVar(relay_func.attrs["global_symbol"]) ``` This is then registered on a target: ``` TVM_REGISTER_TARGET_KIND("woofles", kDLCPU) .set_attr("relay_to_tir", "target.woofles.lowering"); ``` +The signature for this hook is as follows: +``` +relay_to_tir(const IRModule& ir_module, const relay::Function& function) -> (IRModule, GlobalVar) +``` +Which takes a read only `IRModule` and relevant `Function` and returns a new `IRModule` which represents the transformed function, alongside a `GlobalVar` which indicates the top-level operator function within that new `IRModule`. ## TIR -> Runtime Extending from the above, a second hook is introduced to do further transformations from TIR -> Runtime named `tir_to_runtime`, this bypasses the default `target.build.X` and instead uses the registered `tir_to_runtime` build: @@ -95,7 +104,7 @@ This functionality is an extension of the existing use of `attr::kCompiler` to p ## Relay to TIR Hook [relay-to-tir-hook]: #relay-to-tir-hook -This can be added into the `compile_engine.cc` by cross referencing the existing `attr::kCompiler` with the `TargetKind` registry: +This can be added to the TE Compiler by cross referencing the existing `attr::kCompiler` with the `TargetKind` registry: ``` auto code_gen_name = key->source_func->GetAttr(attr::kCompiler).value(); auto target_kind = tvm::TargetKind::Get(code_gen_name).value(); @@ -108,7 +117,7 @@ if (target_kind.defined()) { return CachedFunc(cache_node); } ``` -By placing this where lowering currently takes place, it means minimal changes to executor code generators as they call into `Lower` in `CompileEngine`. +By placing this where lowering currently takes place, it means minimal changes to executor code generators as they call into `LowerTE` and thus are agnostic to how it gets lowered. ## TIR to Runtime Hook [tir-to-runtime-hook]: #tir-to-runtime-hook From a90c4afdae2f5e096e6f417c740a30e35194a4e9 Mon Sep 17 00:00:00 2001 From: Chris Sidebottom Date: Tue, 24 Aug 2021 12:14:36 +0100 Subject: [PATCH 3/5] Update hooks definition and lowering process Change-Id: I578145c37f9a10b4c15ed64fa86d6d7c2fade04e --- ...registered-compiler-flow-customisation.md} | 74 ++++++++++--------- 1 file changed, 40 insertions(+), 34 deletions(-) rename rfcs/{0010-additional-target-hooks.md => 0010-target-registered-compiler-flow-customisation.md} (62%) diff --git a/rfcs/0010-additional-target-hooks.md b/rfcs/0010-target-registered-compiler-flow-customisation.md similarity index 62% rename from rfcs/0010-additional-target-hooks.md rename to rfcs/0010-target-registered-compiler-flow-customisation.md index 8b9cd1e5..22f47b0d 100644 --- a/rfcs/0010-additional-target-hooks.md +++ b/rfcs/0010-target-registered-compiler-flow-customisation.md @@ -1,4 +1,4 @@ -Feature Name: additional-target-hooks +Feature Name: `Target` registered compiler flow customisation Start Date: 2021-07-14 RFC PR: apache/tvm-rfcs#10 GitHub Issue: apache/tvm#8589 @@ -6,28 +6,36 @@ GitHub Issue: apache/tvm#8589 # Summary [summary]: #summary -In order to enable flexibility in how individual targets are lowered and built within TVM, this RFC proposes supporting additional hooks on the `Target` and that the target becomes the central place for such hooks, for example: +In order to enable flexibility in how individual targets are lowered and built within TVM, this RFC proposes additional hooks on the `Target` and that the target becomes the central place for such hooks, for example: ``` +using FTVMLowering = Pass; +using FTVMCodegen = runtime::TypedPackedFunc; + TVM_REGISTER_TARGET_KIND("cmsisnn", kDLCPU) - .set_attr("relay_to_tir", "target.cmsisnn.lower") - .set_attr("tir_to_runtime", "target.cmsisnn.build"); + .set_attr("relay_to_tir", CMSISNNLowering) + .set_attr("tir_to_runtime", CMSISNNCodeGen); ``` -This defines two new hooks as attributes on the target, referencing functions registered into the central TVM registry. In similar fashion, external generators (currently accessed directly in the compile engine) would be grouped with an appropriate `Target` as well: +This defines two new hooks as attributes on the target, referencing functions registered into the central TVM registry. In similar fashion, external code generators (registered under the `relay.ext.` namespace currently) would be grouped with an appropriate `Target` as well: ``` +using FTVMExternalCodegen = runtime::TypedPackedFunc; +using FTVMConstantUpdater = runtime::TypedPackedFunc(Expr, std::string)>; + TVM_REGISTER_TARGET_KIND("ethos-n", kDLCPU) - .set_attr("relay_to_runtime", "relay.ext.ethos-n") - .set_attr("constant_updater", "relay.ext.ethos-n.constant_updater"); + .set_attr("relay_to_runtime", EthosNCodeGen) + .set_attr("constant_updater", EthosNConstantUpdater); ``` -Collecting all targets under the `Target` functionality and making it clearer which hooks apply to each target. +Collecting all targets under the `Target` functionality (as opposed to registering additional `Target`s through the function registry using the namespace `relay.ext.`) and making it clearer which hooks apply to each target. # Motivation [motivation]: #motivation -Currently to introduce an external code generator, the entire compilation pipeline must be recreated; this is necessary for some targets but in the case of simply re-using existing libraries or introducing a function call to use for an operator it can become more than is necessary. It also exists outside of the main `PrimFunc`, meaning it can't be inspected as part of the entire main graph; this limits the effectiveness of techniques such as memory planning. By introducing the hook `relay_to_tir`, which is similar to the default `lower` pass in that it returns TIR, it can be inspected by the memory planner and other analysis passes that only work at the TIR level. If all that is necessary is transforming into a flat `call_extern` (such is the case for the [CMSIS NN Softmax function](https://github.com/ARM-software/CMSIS_5/blob/develop/CMSIS/NN/Source/SoftmaxFunctions/arm_softmax_s8.c#L86)) then this can be left represented as TIR and be collected by the host code generation. +Currently to introduce an external code generator (otherwise known as [BYOC](https://tvm.apache.org/docs/dev/relay_bring_your_own_codegen.html)), the entire compilation pipeline must be recreated; this is necessary for some targets but in the case of simply re-using existing libraries or introducing a function call to use for an operator it can become more than is necessary; to implement an external code generator requires going directly from Relay to a `runtime::Module` and re-implementing any compiler passes and code generation functionality rather than being able to extend upon the existing compiler infrastructure. + +The generated `runtime::Module` also exists outside of the main graph, meaning it can't be inspected in combination with other operators; this limits the effectiveness of techniques such as memory planning. By introducing the hook `relay_to_tir`, which is similar to the default `LowerTEPass` in that it returns TIR, it can be inspected by the memory planner and other analysis passes that only work at the TIR level. If all that is necessary is transforming into a flat `call_extern` (such is the case for the [CMSIS NN Softmax function](https://github.com/ARM-software/CMSIS_5/blob/develop/CMSIS/NN/Source/SoftmaxFunctions/arm_softmax_s8.c#L86)) then the hook may simply return that TIR to be collected by the host code generation. In the more complex case, we still want to take advantage of memory planning by using `relay_to_tir` and inspecting the liveness within the TIR graph, but instead want to generate out more complex calls (such as using the [CMSIS NN Structures](https://github.com/ARM-software/CMSIS_5/blob/def6f800f95661eb3451d317f7d0dde504f6020d/CMSIS/NN/Include/arm_nn_types.h#L81-L90)); the `tir_to_runtime` hook can be used to build our intermediary TIR into a Runtime module similarly to how the existing external code generation works. This allows writing of external code generators that also get the benefits of any intermediary analysis or transformation that TVM offers. Alongside being able to use the analysis passes, code generators can extend from existing host code generators, customising only the generation which is relevant to them and gaining maximum benefit from the existing optimisations made in TVM. @@ -35,8 +43,8 @@ In the more complex case, we still want to take advantage of memory planning by [guide-level-explanation]: #guide-level-explanation As a user, you can pick from additional hooks to bypass certain behaviours of the `Target`: -* `relay_to_tir` - Custom lowering direct to TIR -* `tir_to_runtime` - Custom code generation into a runtime module from TIR +* `relay_to_tir` - Customize the lowering flow to TIR +* `tir_to_runtime` - Customize code generation into a runtime module from TIR * `relay_to_runtime` - Full compilation flow from Relay to a runtime module To illustrate where the hooks are placed, please refer to the following diagram: @@ -46,12 +54,12 @@ To illustrate where the hooks are placed, please refer to the following diagram: These can be registered on targets using `set_attr`: ``` TVM_REGISTER_TARGET_KIND("cmsisnn", kDLCPU) - .set_attr("relay_to_tir", "target.cmsisnn.lower") - .set_attr("tir_to_runtime", "target.cmsisnn.build"); + .set_attr("relay_to_tir", CMSISNNLowering) + .set_attr("tir_to_runtime", CMSISNNCodeGen); TVM_REGISTER_TARGET_KIND("ethos-n", kDLCPU) - .set_attr("relay_to_runtime", "relay.ext.ethos-n") - .set_attr("constant_updater", "relay.ext.ethos-n.constant_updater"); + .set_attr("relay_to_runtime", EthosNCodeGen) + .set_attr("constant_updater", EthosNConstantUpdater); ``` ## Relay -> TIR @@ -76,13 +84,11 @@ def tir_generator(ir_module, relay_func): This is then registered on a target: ``` TVM_REGISTER_TARGET_KIND("woofles", kDLCPU) - .set_attr("relay_to_tir", "target.woofles.lowering"); -``` -The signature for this hook is as follows: + .set_attr("relay_to_tir", [](IRModule ir_mod) -> IRModule { + return (*tvm::runtime::Registry::Get("target.test.tir_lowering"))(ir_mod); + }); ``` -relay_to_tir(const IRModule& ir_module, const relay::Function& function) -> (IRModule, GlobalVar) -``` -Which takes a read only `IRModule` and relevant `Function` and returns a new `IRModule` which represents the transformed function, alongside a `GlobalVar` which indicates the top-level operator function within that new `IRModule`. +The signature for this hook is as the same as any other `Pass`, which takes an `IRModule` with `Function`s and returns an `IRModule` with transformed `PrimFunc`s. ## TIR -> Runtime Extending from the above, a second hook is introduced to do further transformations from TIR -> Runtime named `tir_to_runtime`, this bypasses the default `target.build.X` and instead uses the registered `tir_to_runtime` build: @@ -91,9 +97,8 @@ runtime::Module BuildWooflesHost(IRModule mod, Target target) { // ... Custom Code generation here } -TVM_REGISTER_GLOBAL("target.build.woofles").set_body_typed(BuildWooflesHost); TVM_REGISTER_TARGET_KIND("woofles", kDLCPU) - .set_attr("tir_to_runtime", "target.build.woofles"); + .set_attr("tir_to_runtime", BuildWooflesHost); ``` # Reference-level explanation @@ -104,20 +109,20 @@ This functionality is an extension of the existing use of `attr::kCompiler` to p ## Relay to TIR Hook [relay-to-tir-hook]: #relay-to-tir-hook -This can be added to the TE Compiler by cross referencing the existing `attr::kCompiler` with the `TargetKind` registry: +This can be added before the `LowerTEPass` in `build_module.cc`, as a pass which iterates over `Target`s and transforming the relevant functions which will then be skipped by the `Function`-level passes until the `PrimFunc` passes begin: + + ``` -auto code_gen_name = key->source_func->GetAttr(attr::kCompiler).value(); -auto target_kind = tvm::TargetKind::Get(code_gen_name).value(); -if (target_kind.defined()) { - auto map = tvm::TargetKind::GetAttrMap("relay_to_tir"); - std::string custom_lowering = map[target_kind]; - auto lowering_function = tvm::runtime::Registry::Get(custom_lowering); - cache_node->target = key->target; - cache_node->funcs = (*lowering_function)(key->source_func, key->target); - return CachedFunc(cache_node); +for (Target target : targets_) { + auto target_kind = target->kind; + auto map = tvm::TargetKind::GetAttrMap("relay_to_tir"); + if (map.count(target_kind)) { + ir_mod = map[target_kind](ir_mod, pass_context); + } } ``` -By placing this where lowering currently takes place, it means minimal changes to executor code generators as they call into `LowerTE` and thus are agnostic to how it gets lowered. + +By placing this above the `LowerTEPass`, this means any functions which are not processed in this way can be processed by the default lowering without interfering with `LowerTEPass`. To achieve this initially `kCompiler` would be used to carry the relevant target information, but the goal is to ensure all `Target`s are visible in `build_module.cc`. ## TIR to Runtime Hook [tir-to-runtime-hook]: #tir-to-runtime-hook @@ -139,6 +144,7 @@ This would replace the existing `relay.ext.` lookup in `compile_engine.c [drawbacks]: #drawbacks * Different hooks are currently dealt with in quite disparate parts of the codebase which are being heavily refactored +* Introducing custom TIR has the potential to add edge cases to the compiler which may uncover new bugs # Prior art [prior-art]: #prior-art From 6eb269574a91c1bea93beb03afc7431c1d86b87b Mon Sep 17 00:00:00 2001 From: Chris Sidebottom Date: Wed, 22 Sep 2021 12:28:10 +0100 Subject: [PATCH 4/5] Re-draw API design for RelayToTIR and RelayToRuntime Based on the discussions around these hooks, we now have a better idea of how to introduce them into the codebase. --- ...-registered-compiler-flow-customisation.md | 149 ++++++++++-------- 1 file changed, 87 insertions(+), 62 deletions(-) diff --git a/rfcs/0010-target-registered-compiler-flow-customisation.md b/rfcs/0010-target-registered-compiler-flow-customisation.md index 22f47b0d..4c6d0bf1 100644 --- a/rfcs/0010-target-registered-compiler-flow-customisation.md +++ b/rfcs/0010-target-registered-compiler-flow-customisation.md @@ -1,31 +1,31 @@ -Feature Name: `Target` registered compiler flow customisation -Start Date: 2021-07-14 -RFC PR: apache/tvm-rfcs#10 -GitHub Issue: apache/tvm#8589 +- Feature Name: `Target` registered compiler flow customisation +- Start Date: 2021-07-14 +- RFC PR: https://github.com/apache/tvm-rfcs/pull/10 +- GitHub Issue: https://github.com/apache/tvm/issues/8589 # Summary [summary]: #summary In order to enable flexibility in how individual targets are lowered and built within TVM, this RFC proposes additional hooks on the `Target` and that the target becomes the central place for such hooks, for example: -``` -using FTVMLowering = Pass; -using FTVMCodegen = runtime::TypedPackedFunc; +```c++ +using FTVMRelayToTIR = Pass; +using FTVMTIRToRuntime = runtime::TypedPackedFunc; TVM_REGISTER_TARGET_KIND("cmsisnn", kDLCPU) - .set_attr("relay_to_tir", CMSISNNLowering) - .set_attr("tir_to_runtime", CMSISNNCodeGen); + .set_attr("RelayToTIR", CMSISNNLowering) + .set_attr("TIRToRuntime", CMSISNNCodeGen); ``` This defines two new hooks as attributes on the target, referencing functions registered into the central TVM registry. In similar fashion, external code generators (registered under the `relay.ext.` namespace currently) would be grouped with an appropriate `Target` as well: -``` -using FTVMExternalCodegen = runtime::TypedPackedFunc; +```c++ +using FTVMRelayToRuntime = runtime::TypedPackedFunc; using FTVMConstantUpdater = runtime::TypedPackedFunc(Expr, std::string)>; TVM_REGISTER_TARGET_KIND("ethos-n", kDLCPU) - .set_attr("relay_to_runtime", EthosNCodeGen) - .set_attr("constant_updater", EthosNConstantUpdater); + .set_attr("RelayToRuntime", EthosNCodeGen) + .set_attr("UpdateConstants", EthosNConstantUpdater); ``` Collecting all targets under the `Target` functionality (as opposed to registering additional `Target`s through the function registry using the namespace `relay.ext.`) and making it clearer which hooks apply to each target. @@ -35,111 +35,136 @@ Collecting all targets under the `Target` functionality (as opposed to registeri Currently to introduce an external code generator (otherwise known as [BYOC](https://tvm.apache.org/docs/dev/relay_bring_your_own_codegen.html)), the entire compilation pipeline must be recreated; this is necessary for some targets but in the case of simply re-using existing libraries or introducing a function call to use for an operator it can become more than is necessary; to implement an external code generator requires going directly from Relay to a `runtime::Module` and re-implementing any compiler passes and code generation functionality rather than being able to extend upon the existing compiler infrastructure. -The generated `runtime::Module` also exists outside of the main graph, meaning it can't be inspected in combination with other operators; this limits the effectiveness of techniques such as memory planning. By introducing the hook `relay_to_tir`, which is similar to the default `LowerTEPass` in that it returns TIR, it can be inspected by the memory planner and other analysis passes that only work at the TIR level. If all that is necessary is transforming into a flat `call_extern` (such is the case for the [CMSIS NN Softmax function](https://github.com/ARM-software/CMSIS_5/blob/develop/CMSIS/NN/Source/SoftmaxFunctions/arm_softmax_s8.c#L86)) then the hook may simply return that TIR to be collected by the host code generation. +The generated `runtime::Module` also exists outside of the main graph, meaning it can't be inspected in combination with other operators; this limits the effectiveness of techniques such as memory planning. By introducing the hook `RelayToTIR`, which is similar to the default `LowerTEPass` in that it returns TIR, it can be inspected by the memory planner and other analysis passes that only work at the TIR level. If all that is necessary is transforming into a flat `call_extern` (such is the case for the [CMSIS NN Softmax function](https://github.com/ARM-software/CMSIS_5/blob/develop/CMSIS/NN/Source/SoftmaxFunctions/arm_softmax_s8.c#L86)) then the hook may simply return that TIR to be collected by the host code generation. -In the more complex case, we still want to take advantage of memory planning by using `relay_to_tir` and inspecting the liveness within the TIR graph, but instead want to generate out more complex calls (such as using the [CMSIS NN Structures](https://github.com/ARM-software/CMSIS_5/blob/def6f800f95661eb3451d317f7d0dde504f6020d/CMSIS/NN/Include/arm_nn_types.h#L81-L90)); the `tir_to_runtime` hook can be used to build our intermediary TIR into a Runtime module similarly to how the existing external code generation works. This allows writing of external code generators that also get the benefits of any intermediary analysis or transformation that TVM offers. Alongside being able to use the analysis passes, code generators can extend from existing host code generators, customising only the generation which is relevant to them and gaining maximum benefit from the existing optimisations made in TVM. +In the more complex case, we still want to take advantage of memory planning by using `RelayToTIR` and inspecting the liveness within the TIR graph, but instead want to generate out more complex calls (such as using the [CMSIS NN Structures](https://github.com/ARM-software/CMSIS_5/blob/def6f800f95661eb3451d317f7d0dde504f6020d/CMSIS/NN/Include/arm_nn_types.h#L81-L90)); the `TIRToRuntime` hook can be used to build our intermediary TIR into a Runtime module similarly to how the existing external code generation works. This allows writing of external code generators that also get the benefits of any intermediary analysis or transformation that TVM offers. Alongside being able to use the analysis passes, code generators can extend from existing host code generators, customising only the generation which is relevant to them and gaining maximum benefit from the existing optimisations made in TVM. # Guide-level explanation [guide-level-explanation]: #guide-level-explanation As a user, you can pick from additional hooks to bypass certain behaviours of the `Target`: -* `relay_to_tir` - Customize the lowering flow to TIR -* `tir_to_runtime` - Customize code generation into a runtime module from TIR -* `relay_to_runtime` - Full compilation flow from Relay to a runtime module +* `RelayToTIR` - Customize the lowering flow to TIR +* `TIRToRuntime` - Customize code generation into a runtime module from TIR +* `RelayToRuntime` - Full compilation flow from Relay to a runtime module To illustrate where the hooks are placed, please refer to the following diagram: -![Diagram showing the splitting point of relay_to_runtime, relay_to_tir and tir_to_runtime](./assets/000x/bypass.png) +![Diagram showing the splitting point of RelayToRuntime, RelayToTIR and TIRToRuntime](./assets/000x/bypass.png) These can be registered on targets using `set_attr`: ``` TVM_REGISTER_TARGET_KIND("cmsisnn", kDLCPU) - .set_attr("relay_to_tir", CMSISNNLowering) - .set_attr("tir_to_runtime", CMSISNNCodeGen); + .set_attr("RelayToTIR", CMSISNNLowering) + .set_attr("TIRToRuntime", CMSISNNCodeGen); TVM_REGISTER_TARGET_KIND("ethos-n", kDLCPU) - .set_attr("relay_to_runtime", EthosNCodeGen) - .set_attr("constant_updater", EthosNConstantUpdater); + .set_attr("RelayToRuntime", EthosNCodeGen) + .set_attr("UpdateConstants", EthosNConstantUpdater); ``` ## Relay -> TIR -With this change, this path splits, depending on whether you wanted to generate a full `Module` or introduce some specific TIR nodes into the code generation flow; the addition of the `relay_to_tir` hook allows you to write trivial external TIR generators such as calling out to a third party library: -```python -@tvm.register_func("target.woofles.lowering") -def tir_generator(ir_module, relay_func): - """A simple TIR generator for testing""" - ib = tvm.tir.ir_builder.create() - A = tvm.tir.decl_buffer(shape=(8,8,), dtype=relay_func.params[0].checked_type.dtype) - B = tvm.tir.decl_buffer(shape=(8,8,), dtype=relay_func.params[0].checked_type.dtype) - C = tvm.tir.decl_buffer(shape=(8,8,), dtype=relay_func.params[0].checked_type.dtype) - ib.emit( - tvm.tir.call_extern('int32', 'woofles', A.data, B.data, 8, 8, C.data) - ) - - prim_func = tvm.tir.PrimFunc([A, B, C], ib.get()) - new_module = tvm.lower(prim_func, name=relay_func.attrs["global_symbol"]) - - return new_module, GlobalVar(relay_func.attrs["global_symbol"]) +With this change, this path splits, depending on whether you wanted to generate a full `Module` or introduce some specific TIR nodes into the code generation flow; the addition of the `RelayToTIR` hook allows you to write trivial external TIR generators such as calling out to a third party library: +```c++ +void CallExternalLibraryInTIR(const GlobalVar& new_global_var, const Function& func) { + tir::Buffer x_buffer = tir::decl_buffer({8}, DataType::Float(32), "x"); + tir::Var x_var("x", DataType::Handle()); + + Map dict_attrs; + dict_attrs.Set("global_symbol", new_global_var->name_hint); + dict_attrs.Set("tir.noalias", Bool(true)); + + Map buffer_map = {{x_var, x_buffer}}; + tir::Stmt body = + tir::Evaluate(tvm::tir::Call(DataType::Int(8), tir::builtin::call_extern(), {x->data})); + + tir::PrimFunc replacement_func = tir::PrimFunc({x_var}, body, VoidType(), + buffer_map, DictAttrs(dict_attrs)); + replacement_func = WithAttr(replacement_func, ::tvm::attr::kTarget, host_target_); + ir_module_->Add(new_global_var, replacement_func); +} ``` This is then registered on a target: -``` +```c++ TVM_REGISTER_TARGET_KIND("woofles", kDLCPU) - .set_attr("relay_to_tir", [](IRModule ir_mod) -> IRModule { - return (*tvm::runtime::Registry::Get("target.test.tir_lowering"))(ir_mod); - }); + .set_attr("RelayToTIR", relay::contrib::woofles::RelayToTIR()); ``` The signature for this hook is as the same as any other `Pass`, which takes an `IRModule` with `Function`s and returns an `IRModule` with transformed `PrimFunc`s. ## TIR -> Runtime -Extending from the above, a second hook is introduced to do further transformations from TIR -> Runtime named `tir_to_runtime`, this bypasses the default `target.build.X` and instead uses the registered `tir_to_runtime` build: -``` +Extending from the above, a second hook is introduced to do further transformations from TIR -> Runtime named `TIRToRuntime`, this bypasses the default `target.build.X` and instead uses the registered `TIRToRuntime` build: +```c++ runtime::Module BuildWooflesHost(IRModule mod, Target target) { // ... Custom Code generation here } TVM_REGISTER_TARGET_KIND("woofles", kDLCPU) - .set_attr("tir_to_runtime", BuildWooflesHost); + .set_attr("TIRToRuntime", BuildWooflesHost); ``` # Reference-level explanation [reference-level-explanation]: #reference-level-explanation -This functionality is an extension of the existing use of `attr::kCompiler` to provide a hint that we can use to lookup attached target attribute, the compile engine and code generation flows can choose to store TIR and/or generate runtime modules based on the registered hooks. +This functionality is an extension of the existing use of `attr::kCompiler` to provide a hint that we can use to lookup attached target attribute, the compiler and code generation flows can choose to store TIR and/or generate runtime modules based on the registered hooks. ## Relay to TIR Hook [relay-to-tir-hook]: #relay-to-tir-hook -This can be added before the `LowerTEPass` in `build_module.cc`, as a pass which iterates over `Target`s and transforming the relevant functions which will then be skipped by the `Function`-level passes until the `PrimFunc` passes begin: +This can be added before the `LowerTEPass`, as a `Pass` which iterates over `Target`s and transforming the relevant functions which will then be skipped by the `Function`-level passes until the `PrimFunc` passes begin: -``` +```c++ for (Target target : targets_) { auto target_kind = target->kind; - auto map = tvm::TargetKind::GetAttrMap("relay_to_tir"); + auto map = tvm::TargetKind::GetAttrMap("RelayToTIR"); if (map.count(target_kind)) { ir_mod = map[target_kind](ir_mod, pass_context); } } ``` -By placing this above the `LowerTEPass`, this means any functions which are not processed in this way can be processed by the default lowering without interfering with `LowerTEPass`. To achieve this initially `kCompiler` would be used to carry the relevant target information, but the goal is to ensure all `Target`s are visible in `build_module.cc`. +By placing this above the `LowerTEPass`, this means any functions which are not processed in this way can be processed by the default lowering without interfering with `LowerTEPass`. To achieve this initially `kCompiler` would be used to carry the relevant target information, but the goal is to ensure all `Target`s are visible as `kTarget`. + +```c++ +return tvm::transform::Sequential({tvm::relay::transform::RelayToTIRTargetHook(), // Additional Pass to call RelayToTIR + tvm::transform::CreateModulePass(pass_func, 0, "LowerTE", {}), + InferType()}); +``` ## TIR to Runtime Hook [tir-to-runtime-hook]: #tir-to-runtime-hook -Instead of replicating the current external code generation hook, it is proposed that this hook exists in `build_module.cc`: -``` -auto target_built_mods = FindFuncsWithTargetBuild(lowered_funcs); -auto ext_mods = executor_codegen_->GetExternalModules(); -auto extra_mods = ext_mods->Concat(target_built_mods); -ret_.mod = tvm::codegen::CreateMetadataModule(ret_.params, ret_.mod, extra_mods, GetTargetHost(), - executor_codegen_->GetMetadata()); +It is proposed that this hook is implemented as part of `codegen.cc` as a direct override of the code generation: + +```c++ +runtime::Module Build(IRModule mod, Target target) { + if (transform::PassContext::Current() + ->GetConfig("tir.disable_assert", Bool(false)) + .value()) { + mod = tir::transform::SkipAssert()(mod); + } + + if (target->kind->HasHook("TIRToRuntime")) { // Hooked here for Codegen + return target->kind->GetAttr("TIRToRuntime")(mod, target); + } + + // the build function. + std::string build_f_name = "target.build." + target->kind->name; + const PackedFunc* bf = runtime::Registry::Get(build_f_name); + ICHECK(bf != nullptr) << build_f_name << " is not enabled"; + return (*bf)(mod, target); +} ``` -This means the hook is integrated at a higher level and included in the compile flow without executors having to be aware of how these modules exist. See [Relay to TIR Hook](#relay-to-tir-hook) for how the `TargetKind` registry would be used. +See [Relay to TIR Hook](#relay-to-tir-hook) for how the `TargetKind` registry would be used. ## Relay to Runtime Hook [relay-to-runtime-hook]: #relay-to-runtime-hook -This would replace the existing `relay.ext.` lookup in `compile_engine.cc`, essentially using the same logic as [Relay to TIR Hook](#relay-to-tir-hook) to cross reference with `kCompiler`. +This would replace the existing `relay.ext.` lookup in `te_compiler.cc` with a `Pass` which runs beforehand, essentially using the same logic as [Relay to TIR Hook](#relay-to-tir-hook) to cross reference with `kCompiler`. +```c++ +return tvm::transform::Sequential({tvm::relay::transform::RelayToTIRTargetHook(), + tvm::relay::transform::RelayToRuntimeTargetHook(), // Additional Pass to call RelayToRuntime + tvm::transform::CreateModulePass(pass_func, 0, "LowerTE", {}), + InferType()}); +``` # Drawbacks [drawbacks]: #drawbacks @@ -158,4 +183,4 @@ In future, this approach enables rapid integration of anything that can be repre Alongside this, adding further hooks means external code generation can gain benefits from the normal `lower` and `build` flow in TVM. This then expands to exposing more granular methods in the driver api to leverage the compiler passes in TVM, similar to how they've been exposed in https://github.com/apache/tvm/pull/8110 with `lower_primfunc` and `lower_schedule`. This can is then regulated by the normal Target mechanism to route as appropriate. -Refactoring the target splitting logic into `build_module.cc` alongside any external module generation makes this a first class series of hooks into a simplified compilation flow; this would enable the removal of external generators from executor code generators which currently proxy to `compile_engine.cc`. Eventually this could also be used for CPU/GPU split as a specialisation of a `Target`/`Target`s split. +Refactoring the target splitting logic into `build_module.cc` alongside any external module generation makes this a first class series of hooks into a simplified compilation flow; this would enable the removal of external generators from executor code generators which currently proxy to `te_compiler.cc`. Eventually this could also be used for CPU/GPU split as a specialisation of a `Target`/`Target`s split. From 47aff331e1e932b96565d2b59e64a2a72fe63106 Mon Sep 17 00:00:00 2001 From: Chris Sidebottom Date: Thu, 23 Sep 2021 10:37:04 +0100 Subject: [PATCH 5/5] Provide further clarification on hooks and minor text fixes --- ...-registered-compiler-flow-customisation.md | 20 ++++++++++++++----- 1 file changed, 15 insertions(+), 5 deletions(-) diff --git a/rfcs/0010-target-registered-compiler-flow-customisation.md b/rfcs/0010-target-registered-compiler-flow-customisation.md index 4c6d0bf1..65fc17bc 100644 --- a/rfcs/0010-target-registered-compiler-flow-customisation.md +++ b/rfcs/0010-target-registered-compiler-flow-customisation.md @@ -28,12 +28,14 @@ TVM_REGISTER_TARGET_KIND("ethos-n", kDLCPU) .set_attr("UpdateConstants", EthosNConstantUpdater); ``` -Collecting all targets under the `Target` functionality (as opposed to registering additional `Target`s through the function registry using the namespace `relay.ext.`) and making it clearer which hooks apply to each target. +Collecting all targets under the `Target` functionality (as opposed to registering additional `Target`s through the function registry using the namespace `relay.ext.`) and makes it clearer which hooks apply to each target. # Motivation [motivation]: #motivation -Currently to introduce an external code generator (otherwise known as [BYOC](https://tvm.apache.org/docs/dev/relay_bring_your_own_codegen.html)), the entire compilation pipeline must be recreated; this is necessary for some targets but in the case of simply re-using existing libraries or introducing a function call to use for an operator it can become more than is necessary; to implement an external code generator requires going directly from Relay to a `runtime::Module` and re-implementing any compiler passes and code generation functionality rather than being able to extend upon the existing compiler infrastructure. +We want to make external code generation (otherwise known as [BYOC](https://tvm.apache.org/docs/dev/relay_bring_your_own_codegen.html)) more modular; instead of going from a Relay `IRModule` to `runtime::Module` in one big step, you can break it into phases and make use of existing transformations between phases. + +Currently to introduce an external code generator, the entire compilation pipeline must be recreated; this is necessary for some targets but in the case of simply re-using existing libraries or introducing a function call to use for an operator it can become more than is necessary; to implement an external code generator requires going directly from Relay to a `runtime::Module` and re-implementing any compiler passes and code generation functionality rather than being able to extend upon the existing compiler infrastructure. The generated `runtime::Module` also exists outside of the main graph, meaning it can't be inspected in combination with other operators; this limits the effectiveness of techniques such as memory planning. By introducing the hook `RelayToTIR`, which is similar to the default `LowerTEPass` in that it returns TIR, it can be inspected by the memory planner and other analysis passes that only work at the TIR level. If all that is necessary is transforming into a flat `call_extern` (such is the case for the [CMSIS NN Softmax function](https://github.com/ARM-software/CMSIS_5/blob/develop/CMSIS/NN/Source/SoftmaxFunctions/arm_softmax_s8.c#L86)) then the hook may simply return that TIR to be collected by the host code generation. @@ -52,7 +54,8 @@ To illustrate where the hooks are placed, please refer to the following diagram: ![Diagram showing the splitting point of RelayToRuntime, RelayToTIR and TIRToRuntime](./assets/000x/bypass.png) These can be registered on targets using `set_attr`: -``` + +```c++ TVM_REGISTER_TARGET_KIND("cmsisnn", kDLCPU) .set_attr("RelayToTIR", CMSISNNLowering) .set_attr("TIRToRuntime", CMSISNNCodeGen); @@ -63,7 +66,8 @@ TVM_REGISTER_TARGET_KIND("ethos-n", kDLCPU) ``` ## Relay -> TIR -With this change, this path splits, depending on whether you wanted to generate a full `Module` or introduce some specific TIR nodes into the code generation flow; the addition of the `RelayToTIR` hook allows you to write trivial external TIR generators such as calling out to a third party library: +With this change, this path splits, depending on whether you wanted to generate a full `Module` or introduce some specific TIR nodes into the code generation flow. The `RelayToTIR` hook is a full `IRModule` `Pass` which expects that `Function`s will either be annotated with `kTarget` or `kCompiler` as part of a previous `Pass`, and the resultant `IRModule` is also expected to have any created `PrimFunc`s annotated. The addition of the `RelayToTIR` hook allows you to write trivial external TIR generators such as calling out to a third party library: + ```c++ void CallExternalLibraryInTIR(const GlobalVar& new_global_var, const Function& func) { tir::Buffer x_buffer = tir::decl_buffer({8}, DataType::Float(32), "x"); @@ -83,15 +87,19 @@ void CallExternalLibraryInTIR(const GlobalVar& new_global_var, const Function& f ir_module_->Add(new_global_var, replacement_func); } ``` + This is then registered on a target: + ```c++ TVM_REGISTER_TARGET_KIND("woofles", kDLCPU) .set_attr("RelayToTIR", relay::contrib::woofles::RelayToTIR()); ``` -The signature for this hook is as the same as any other `Pass`, which takes an `IRModule` with `Function`s and returns an `IRModule` with transformed `PrimFunc`s. + +The signature for this hook is as the same as any other `Pass`, which takes an `IRModule` with `Function`s and returns an `IRModule` with transformed `PrimFunc`s. The registered `RelayToTIR` `Pass` is responsible for both establishing the `PrimFunc` definitions (with any caching) and rewriting Relay calls to those functions. At this time we feel it's not worth worrying about code sharing between different custom passes. ## TIR -> Runtime Extending from the above, a second hook is introduced to do further transformations from TIR -> Runtime named `TIRToRuntime`, this bypasses the default `target.build.X` and instead uses the registered `TIRToRuntime` build: + ```c++ runtime::Module BuildWooflesHost(IRModule mod, Target target) { // ... Custom Code generation here @@ -101,6 +109,8 @@ TVM_REGISTER_TARGET_KIND("woofles", kDLCPU) .set_attr("TIRToRuntime", BuildWooflesHost); ``` +Notably the generation hook is passed the unified `IRModule` and is responsible for plucking the `Target` relevant functions into the eventual `runtime::Module`. + # Reference-level explanation [reference-level-explanation]: #reference-level-explanation