From 47dadff0e0ba5b5bda746d2f97a86376acd6e9ef Mon Sep 17 00:00:00 2001 From: Ben van Werkhoven Date: Fri, 18 Nov 2022 08:49:58 +0100 Subject: [PATCH] add demo --- .gitignore | 1 + demo/Kernel_Tuner_demo.css | 23 +++ demo/Kernel_Tuner_demo.ipynb | 317 ++++++++++++++++++++++++++++++++++ demo/img/dashboard.png | Bin 0 -> 489764 bytes demo/img/dashboard_logo.png | Bin 0 -> 39586 bytes demo/img/slide_background.png | Bin 0 -> 6111 bytes demo/matmul.cu | 91 ++++++++++ demo/requirements.txt | 2 + 8 files changed, 434 insertions(+) create mode 100644 demo/Kernel_Tuner_demo.css create mode 100644 demo/Kernel_Tuner_demo.ipynb create mode 100644 demo/img/dashboard.png create mode 100644 demo/img/dashboard_logo.png create mode 100644 demo/img/slide_background.png create mode 100644 demo/matmul.cu create mode 100644 demo/requirements.txt diff --git a/.gitignore b/.gitignore index 022fb28..1af36cb 100644 --- a/.gitignore +++ b/.gitignore @@ -1,3 +1,4 @@ *.pyc __pycache__ *.egg-info +.ipynb_checkpoints diff --git a/demo/Kernel_Tuner_demo.css b/demo/Kernel_Tuner_demo.css new file mode 100644 index 0000000..5069d77 --- /dev/null +++ b/demo/Kernel_Tuner_demo.css @@ -0,0 +1,23 @@ +.reveal, +.reveal h1, +.reveal h2, +.reveal h3, +.reveal h4, +.reveal h5, +.reveal h6 { + font-family: "Nunito"; +} + +.p { + font-family: "Assistant"; +} + +body.notebook_app.rise-enabled { + background: url('img/slide_background.png'); + background-position: left top; + height: 100%; + width: 100%; + background-repeat: no-repeat; + background-size: auto 100%; + padding-top: calc(50px + 2vh); +} diff --git a/demo/Kernel_Tuner_demo.ipynb b/demo/Kernel_Tuner_demo.ipynb new file mode 100644 index 0000000..bbee7f7 --- /dev/null +++ b/demo/Kernel_Tuner_demo.ipynb @@ -0,0 +1,317 @@ +{ + "cells": [ + { + "cell_type": "markdown", + "id": "09f74f14", + "metadata": { + "slideshow": { + "slide_type": "slide" + } + }, + "source": [ + "
\n", + " \n", + "# Kernel Tuner demo\n", + "\n", + "
\n", + "
\n", + "
\n", + "
\n", + "
\n", + "
\n", + "
\n", + "\n", + "\n", + "By Ben van Werkhoven, Netherlands eScience Center
\n", + "b.vanwerkhoven@esciencecenter.nl\n", + " \n", + "
" + ] + }, + { + "cell_type": "markdown", + "id": "305a09d5", + "metadata": { + "slideshow": { + "slide_type": "notes" + } + }, + "source": [ + "alt+r to start the slideshow, spacebar or shift+spacebar to move forward to next slide, comma to remove on screen buttons\n", + "\n", + "preparation: run the next code cell, start a second terminal and go the the directory of this notebook" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "4b5e909f", + "metadata": { + "slideshow": { + "slide_type": "skip" + } + }, + "outputs": [], + "source": [ + "%%bash\n", + "rm matmul_cache.json\n", + "rm vector_add.cu" + ] + }, + { + "cell_type": "markdown", + "id": "0fbf8c5a", + "metadata": { + "slideshow": { + "slide_type": "slide" + } + }, + "source": [ + "# Optimizing GPU Applications\n", + "\n", + "To maximize GPU code performance, you need to find the best combination of:\n", + "\n", + "* Different mappings of the problem to threads and thread blocks\n", + "* Different data layouts in different memories (shared, constant, …)\n", + "* Different ways of exploiting special hardware features\n", + "* Thread block dimensions\n", + "* Code optimizations that may be applied or not\n", + "* Work per thread in each dimension\n", + "* Loop unrolling factors\n", + "* Overlapping computation and communication\n", + "* ...\n", + "\n", + "Problem:\n", + "* Creates a very large design space!" + ] + }, + { + "cell_type": "markdown", + "id": "59e489ff", + "metadata": { + "slideshow": { + "slide_type": "slide" + } + }, + "source": [ + "# Kernel Tuner\n", + "\n", + "*A Python tool for optimizing and tuning GPU applications*\n", + "\n", + "Started in 2016:\n", + "* As a software development tool for GPU projects at the eScience center\n", + "* To be used directly on existing kernels\n", + "* Without inserting dependences in the kernel code\n", + "* Kernels can still be compiled with regular compilers\n", + "\n", + "Today:\n", + "* Comprehensive toolbox for auto-tuning with several tools being built on top\n", + "* Developed by a team of 7 developers across CWI, Astron, and eScience center\n", + "* Used in over 10 different eScience center projects and by others\n", + "\n", + "https://github.com/KernelTuner/kernel_tuner" + ] + }, + { + "cell_type": "markdown", + "id": "4e2675bf", + "metadata": { + "slideshow": { + "slide_type": "slide" + } + }, + "source": [ + "# Minimal Example" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "169859e8", + "metadata": {}, + "outputs": [], + "source": [ + "%%writefile vector_add.cu\n", + "__global__ void vector_add(float *c, float *a, float *b, int n) {\n", + " int i = blockIdx.x * block_size_x + threadIdx.x;\n", + " if (i\n", + "\n", + "
\n", + "\n", + "#### Live visualizations of auto-tuning sessions using Kernel Tuner\n", + "\n", + "\n", + " \n", + "https://github.com/KernelTuner/dashboard\n", + "
" + ] + }, + { + "cell_type": "markdown", + "id": "57b5db81", + "metadata": { + "slideshow": { + "slide_type": "slide" + } + }, + "source": [ + "# Tuning a larger problem" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "fed6931e", + "metadata": {}, + "outputs": [], + "source": [ + "from collections import OrderedDict\n", + "problem_size = (512, 512)\n", + "A = np.random.randn(*problem_size).astype(np.float32)\n", + "B = np.random.randn(*problem_size).astype(np.float32)\n", + "C = np.zeros_like(A)\n", + "\n", + "args = [C, A, B]\n", + "\n", + "tune_params = OrderedDict()\n", + "tune_params[\"block_size_x\"] = [2**i for i in range(0, 11)]\n", + "tune_params[\"block_size_y\"] = [2**i for i in range(0, 11)]\n", + "tune_params[\"tile_size_x\"] = [2**i for i in range(0, 6)]\n", + "tune_params[\"tile_size_y\"] = [2**i for i in range(0, 6)]\n", + "\n", + "restrict = [\"block_size_x == block_size_y * tile_size_y\"]\n", + "grid_div_x = [\"block_size_x\", \"tile_size_x\"]\n", + "grid_div_y = [\"block_size_y\", \"tile_size_y\"]\n", + "\n", + "from kernel_tuner.nvml import NVMLObserver\n", + "nvml_observer = NVMLObserver([\"nvml_energy\", \"temperature\", \"core_freq\"])\n", + "\n", + "metrics = OrderedDict()\n", + "metrics[\"GFLOP/s\"] = lambda p : (2 * 512**3 / 1e9) / (p[\"time\"] / 1e3)\n", + "metrics[\"GFLOPs/W\"] = lambda p : (2 * 512**3 / 1e9) / (p[\"nvml_energy\"])\n", + "\n", + "_ = kt.tune_kernel(\"matmul_kernel\", \"matmul.cu\", problem_size, args, tune_params,\n", + " observers=[nvml_observer], grid_div_y=grid_div_y, grid_div_x=grid_div_x,\n", + " restrictions=restrict, metrics=metrics, cache=\"matmul_cache.json\")" + ] + }, + { + "cell_type": "markdown", + "id": "3778bbd1", + "metadata": { + "slideshow": { + "slide_type": "skip" + } + }, + "source": [ + "While the previous cell is running go to second terminal and type \"ktdashboard matmul_cache.json\"" + ] + }, + { + "cell_type": "markdown", + "id": "a7377204", + "metadata": { + "slideshow": { + "slide_type": "slide" + } + }, + "source": [ + "# Final remarks\n", + "\n", + "Currently, using Kernel Tuner to optimize and tune code in:\n", + "* Ultrasound Brain Imaging with Erasmus MC (RECRUIT)\n", + "* Atmospheric Modeling (ESiWACE-2 Microhh)\n", + "* Radio Astronomy (CORTEX)\n", + "\n", + "Kernel Tuner can also be used for optimizing the energy efficiency of GPU applications:\n", + "> Going green: optimizing GPUs for energy efficiency through model-steered auto-tuning
\n", + "R. Schoonhoven, B. Veenboer, B. van Werkhoven, K. J. Batenburg
\n", + "International Workshop on Performance Modeling, Benchmarking and Simulation of High Performance Computer Systems (PMBS) at Supercomputing (SC22) 2022 \n", + "\n", + "Main repository:
\n", + "https://github.com/KernelTuner/kernel_tuner
\n", + "Documentation:
\n", + "https://KernelTuner.github.io
\n", + "Tutorial:
\n", + "https://github.com/KernelTuner/kernel_tuner_tutorial
" + ] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "6f89b4b0", + "metadata": { + "slideshow": { + "slide_type": "skip" + } + }, + "outputs": [], + "source": [] + } + ], + "metadata": { + "celltoolbar": "Slideshow", + "kernelspec": { + "display_name": "Python 3 (ipykernel)", + "language": "python", + "name": "python3" + }, + "language_info": { + "codemirror_mode": { + "name": "ipython", + "version": 3 + }, + "file_extension": ".py", + "mimetype": "text/x-python", + "name": "python", + "nbconvert_exporter": "python", + "pygments_lexer": "ipython3", + "version": "3.9.12" + }, + "rise": { + "transition": "none" + } + }, + "nbformat": 4, + "nbformat_minor": 5 +} diff --git a/demo/img/dashboard.png b/demo/img/dashboard.png new file mode 100644 index 0000000000000000000000000000000000000000..00cab9ea8fec7a8edc7844a30c1fa21ad4dd87f3 GIT binary patch literal 489764 zcmd4&WmH|u(gq5X5G0V`?(P!YA-KD{ySqbhcNXrhi{S3=?yzuoceptxXYY4>|Gt0s z));fluAXgGZPnG!43m=){SJc#0|o~6U0h5^0SxTRG8owBH|WoQVp5CA%l^DRIVy+> zfB`3Pj{Y!TO!=kx!N96xU|;ni{_tPz#MB+Zz~FlS^L!e#EinQE`zRC_;#YFhIbDa) zn^VCDyCgYqZ9kb`Z!e@hKWEhI`rZBO)9MnW;SWP>L-ubQzqUBn)-$)~N?Kdj?$+YH zhS%j+6mq^jE!}_6Qa5Xz#<%Uh-RS7Z?fyjn@~6iCl3T=&o6dijuGP=g6#rv>HGVgc z|L21(c3<8e`wP>Kn>+@5Ugtz z^8dvOY*E5vA-|{+de92Q@9S?WBSHg3Zm9xE#=wvHi2{@GpxCB7w=xBBG1v{{2qC`< zbou?2^F0D%NRHm|qTEx*>XLZ|gSgErf94GN6U@PYzt3rGh3ZIuaK^DvA#*8mJ0Hni|z-F(d7`a z`Gpk9<_ux>0${Xc7iv)~aFS9E@ND+Xhe?b%A!I7YxPnOJiiYN?D8vX1;zy{!uViq5 ze|6d>OyQ6d(fpu1734z_+#%H#Nc07@RROs}DC4IBETRK$FD)*jw!4M03#GbYSW^<| zr7c+e6hje}i zYXMw$TID&jV@>8>xTfyJScEE*+kRD<{7!7NH+XS2oUJF@$>j7ygR!1~#~BD05_Ps^ z^z=T$XP;XG_i0o$RUz=n24wkH%1vo^+G>hnkyH>G9o}cp3^TK73h;g-NMixzXsIeI zyIh25j3!#{L0wxW99piIt4fX+CojiHzj6edAya9S%KJ+K09$8UyFqI~&##X2pNGcN z)rUEJX&=n}?^}_YqdudWTKZ6bt19lQs z_};eG`?$LzX-eAORVNC;gC6qt(iU-|XR6I+m2rJCjtMld815Ed15Nw#;JiO~@Xj{% zjz_ZW;GGV6p&L%ibk@CCqmg?k@1B~g|8M5}`p8(6cMwCOfgD}zX|7&s3+9+`<}|DjS<8cpXPs^(>x1nrqR2tt8$4g$jXhOMz{)2)o{ z4j1DG6{7ilMNo+ogBBykSq4?~~{>=#5R+nCcss$a<oT2MOL4w0-rtj_6$bB?U6_WCx-K4otE#kqr)c*h41!?oa74l8=t zZ(RS(XDQfcySz$io9^tli!)CnJ#KZ*cCNYsvLLCGfxHK3m4hs?ap!WL=xQw|;}A?; zuFxbQsmxUD!eW;M_X|0`*`|%Cy}WktqAcsRE0KZY>IYWS1ezOmiU_>o^X-zTDyDa- zV}jfPgOo-Q?YZcpg^ZQyfy1~r-TA%7)z0%uamoAF9#8o6hn zMy&Vb@>$aBNSFpDiz#Q~-=+Q6N7mwkOB~vwgtIbIX2n(}l4|-2y~PhC!jy^a%V`_^ z5~4jNAE7)y7}Zma7wKh`neN`$8oK+&We9aFDl!I9-hoRb#3xP{tZh(rh~nr%Lx+J> zW%F(B%x|dNH%wY>eFs#+H;Trw+{Zr?FrW!>Ki6nkYAg$0kuEx7k~}0dL&T z{p?*+p`dH&iwLbhy;ivHxjM(Pw z^aC|Fm&o?z9(;}K!NF!8M94fPkfVgIdidSRQITv;gt&>$FCMtg_NyVam zV4=+Ji}Uf>bVohJiNnJ1OOqk`Bq-pB7XQ4rUjhID^NeV4Y^wIS(tvk%7^i_}U#5Vy z`wo(8Y{13o5e!U%ah5;p-*O@txjH>vD9qLj0Y;B&s(Xea7&nyp7r@6KZ&?2PR#Ir} zM!yJrTU&8)vI1$H_K_7f!&qc1)aFG?%`La|Luc!ZV-|d?wK{$WAtbB#gp!C7Ps?qQ z=jucEY&c#pOa;Bq1EsVC9Q?IB9`sPxXX0YR#IC$l@~Il zG6E1RvZwff%oaz$@RouaOVg79!PuHe`p-yv0jnSN%Mz%W%`c3NgD9r5j-Lt%Pn4dc zv*W7<%o#i=3qJAVw7hs-2FYIV8%WMlgwoyUEyMwa$U5#B?>jwg(DWNe=nJj8V2jTn&3tNqjD%)UTk?*sU#qFu) zNWP-fo}F`@V^;y&Y4EGC6@2nH*NzMFLq;SI@uOL^m)}FGaIN7AwH2UyMA%yux9FwQ ztGOHx#vZAczlbMCK>9eh*cbO*hqx-8KCJa-NbSU;vfao`xYvDz`@?9I_HRLlX=pgV zFqC{yEDiu;d{c#<^<4M2=9e24@4oG93H6ufE#Tox-5>GJP3&FKkDb8lMQEt)RYEI1dyqjpDqzL= zz)z$qH2y}0J@VIz?ed*C@hpbr-&-t&4$QIq!CH`iA~@4;$6}CFbl~RTZd59C?f+iDSs=J>|9!`GJNvVhO7VcQMxP&R=9-YtA-A`9+m< zy_9B(u%ys>@RN&<3=X+1iKFJqRi{JMIi*Kp&JF~}9d4f~mpa5F;V^c54DMK`qjFHQ z8moorL1}Eis1gb$#<##-)`t%V#~r`3>O|?%kOfm6UTygT{FB{pf}%=ZQ0byxl=@sct#f*EkATy z@3N5KFuh%aEjpqwIy4B#KAI=C?=c9JbtDYq8)?TWUX&|e-OYL!Qe?9!5fk}m8r1J$DxpJxdLna<$R<{GdHplpA12BiTm@FSzL zx0Neo$-W|UI*-RZ#AQ$Plr}nQp2AKiM0mV2^Y}2Qw5^He?)3fD)3x5eVFGtmo}X67 zsrc|2f&%@of%U;tMT|x&Zn!E8Bs9QR^~-PVDkaQJ4vL^%)G(hiv!^1l(qn%nL7N|9 zWc;>3oF0iy&}Cy}{YG?e743IdU>8>KSvxwSSM&wT6{XBZrlZ!#eWCL3x(TY^-af^V zI{6arw9L~Oo^iflDR)o|N>+3|AT^HKSAy5En3+Mmuy>hNT5H2sS8P{uLb~JS1)l@) zVIkXdU)=8ljs3XZ> z5&UM|ECgkTkD72yy^XSzk-7P(|#J~b8X%9ZhLwy%jWH$_An<*+Jr>$l%1 zFkyk#PaGbcXx1z%(~6qTuIjfvn3TyENh0g?5$&ZYE;aDb_4f3%&5pA2eoR)H5Mycw z=OmR&Z1?XixXllFRNEIxLOcXk=}T-!Cd6Wuk)9L@#3{<;8-CspLvOn*x_!7!I$0Q+ zc;3ni2EUJ2Rwc$3h0kdy^P{CySN8)l1g%ixBJgRSc#mEy<3+_eE8$QBpS^EFlCeeY z-Dd*$R7`Zd<&r}koBheO+MX~~VXp+?j$pqdo(OWX#hPMl>I4I;3RGDldROh5t&Ru-o;me2}2 z&@$-I@myTd0g>d_cw@fiZMWB23(KD2jDI`q4c?KRVLp3kMak~=vlm3V<>JKfn=~`w z%V!E$ei-^)T^>bB!KMZ1NhAg%t@tELvA$)6F%j`E$dD^{=fetp0GyQU1Ouzg&=`W5 z(~Jmms&GX9S%h4!fP(aa!kBN`+gUx|IL!U*`s!DtED+*JoF#C&<&q#+=p!X96z=;z zuHW0|eRoZaJdpZWM zuQyc+JUzOAvqbMZEzC0^QT9}euoLykc~t(0;lxPsLT9*RirSIO^-}2XZHF-TD8krU zA(9f5d5^0E2*b>sfLnU<1~a69kh#56gJ*6^p(cl!Z*__|rLX|ruX^%uJ=W99tbLHP zgeCX<&(GD32U1dD6g>D>>R0v7-h(gnqGwJ9o1M`3VMw@&`&V6|lWu&+8q@_@>HaN7QW_qGg zUb!SS0_2{loGm6w1$_t0Eifr@##gT8w?1^l=udSIdS-|FI_6HlhK-gQYKr^?f!S}N z<7epBmRz2i>C3B}H0yJ2V)Rt+#!DLVh-Cih0m3+@)B1wrEV=kl^E)epXE}LFa)rWj z5CX!27!txjM?7FQ{sE+9^&~r4r05?5hst&A|;q zXEI>yGK`-{eh?FvXp6mlDN9~4o@1qqloKN%*)}EJXJEI=l7+&{n?cQ>uk=!Yk%s6> zYx|;Z!O>&gUwf1H&{_2RTHZ(;2-uktuJqwrwjK{UPp5Oyu9*+v3>5P(NLMN|=|V&^ z|8vxf5+kB%B-0`bR?48JXOSBozx~1F>}HU%e}<3yKYtRFj$|{1M66JuN?I|+kSZGM z?|czD6@bOjV>b>bbB6tAWhYRUc%!LTmJ2CN*yq=Z4!PTzj9AQm+ZaBaY*zb-im~%$c#@?WaSQ3jFAAROrCQ(WcBy zG;2%lYr}&dzQ96=fKh{MnnSmEIyGig>V1@R9Cy!k7s}cI$12QR#fExyul$n>?B3@W zjmr|$yaTJ?nyD2wcs$>-=IK)D?h+0pkr8K3p0G=HEy#h*AJ(Ohn$jb+mO(*tUo66O zl>G?cq#8n3Am>J4TVK~~j!MxIOdjhjDeDkx*{#SzM3Vy*N11TpP3A1-H0a5)WiQ}_ zi*i*cIh@<(H5mT(#=GIWO~n;6l~r#F{jz-Z;Kp!FEuxUllb&G@C?q^Dr9xY>77tL( zLlzQ)WpMJdvocVmABtTjs-%kh$BDY!DXm@Ux1_CQ{G6g^I7KUfwRSnEm~Hfc$q{*1 z&3SrzLs-;QDvMwYpWn!B8vJ`Sf=5x)a>pH?|Pk70CIu%;puTL(zOG7~J-cc3LqJ{hgG9pB@06%%y zZu?^dhFA+80HhpQ zr~R1&Lhi&CuUfz4O(<(*+jZV5rc;oP(t8d?t5yuV`*&QcGk)~b)C2k&Loi0CAY z;7tuq5@PivD{|MJDzi;AA$iC=RU=fMSYnl2mkVDjYhgYTaXv(exBB0;=2#{tw!)n# z2ZayN;@g}8F;kgv-QS@V*Or{eh*TN!#<^2`isM0A-#!Hx3P*(@v-4Y1TToK?+e0ky z`)DDe5-Gln{fx)aIk8@6UAy2ufu&}+MVLx@91WE#zFEClGaVyE50lXF)bt*y+kdo%kchq-(iBdX+tpVdyqp zS)IcJtp`DErV;g<6N_d`$TNpka94d$CU|?L#o~13cj3!5L!?3&M3;xl`a(FhwWcuW zz`@^kZm++n{`SrX|9_4UUwH#&)lO$YKzYgZ92M5???)>u!FAPNU2=L7RB7?^dgPO8 zwhtdl5Ats_7Cz~|Zevjh+4tp$jjYz{8vky@s;OTt0TO9yMi+h`EDakJ#pbMR##-ov z#jmnyIV{X|vQt)ZGSW5v)NOTRS}Qxa?78Y7v|z z6khSmIY*bV3j|!^XvSP$AP=Fh-ejaV$t1 z^S8gOG_5*Xb6d4cS}^Xon!?@9iY3PIQhO9j<^dLGCI-!*WhW<1TLYIIXlT?vTaKz2 zwu}0TG!nkG{Y3k3&;IhQUPO$@`xeCST+h6(2fFM1eqItMiwYW6PQu8>^8Sh;!8b&% z_uC?oiUC0?JFVW_MaAjsF?03PE`4SjM+rY&=TmQQrD>uC`r;M3=Y{M^Jx}wb?t=!L zP1GJfZKPLD-br2_M!KjqIHWd~8n^_yg#nJu8D|^DKwb|No{7f^EI0!RT{Nf;*!GkD zyt#RF;*^&Q>G@RRjN3Reh`WqrJ5k8miW$(IJ7`jy^`VrDIgq!pHM})v!4w+qnd)&g zJRele5xAaWQqH)+6?+Hj2stCW-Lf>kZE;+>Z7WR)?q-#l0M5D~4j}5yp&80C@ zU^de2KhVZg!KU;3I@>JuJcaX(!A%~36c{t!#$B58y zBkmkN@Hmq_tmUcP=+4L*U3?C_a$0WPL5~rS1QuWe`Vq(!bTI4L?o9l;pWAbm#ldCJ zzArhp&i{5XvX84Sa->MM?d}rO-XA!ee<`O5iZ-0M2PMy0Y6@(u!)y31g~|JHzO9}2 z;KYmf|Irvf8CE_U&EvP(S|CNhiZ`@E#eD37;KTh|0>`DAPWE1eR!05GlDqvN zF^WUHjfGmes>@+S3hXfC^h*>P>Fo=3?&SGgmSfSXCQ2gS6U*MNG^z+XF7lD30=;gD zf;v|;*ZPUxbeEyIwyBps@mOU%acB_f;PR(@a0$qfZDQzbUhc9TNegYcsN1Q=VO52C z8D!1MsPmRrwGhA}*JS3-aY_DwfdW@KA=F=2E7%9?H=~l|<_AEIhL&n+g1FN)q|4lp zLqK?eW}3B|PKJWFa(L}n28}Yu)ENn)bv24N>?McC`SM>9wk&j1DXbUxt)h}*?7Rk>X1wu^Rw#- zn!cEilX+jML{PN=Wf($<+AcC#je;tOPb*6$l&SP5C})^c69c65HUX8eOh|E;-Ga0l!k6z*AdSdLLpDpwT#=K;?Zs{w0~3x`bl(6&2Ts z5_!}>@lNW`1slKF?uvi8yXZb?wiM!*rV^$qwDmSwVOFV-pOSpCDxmtj7uRxp!8=5F>CO+zhz=92Ze|vCF zO%24f$2a$@AddihSnjA5TrGNPd_FqWJXPwi0nt;X-^z&Gp9|Ln`37E_%YI1e`?P+o zN4tH|n5DwkPPXjhMoga*t0#56xtEOwxyuCDAxL+ST&LAY2*zVMuN>tyX~h$p*mF?i z_5+&ZUrvx0$7Napg1-@VR-JmR+A=W1Eg=^1C`ce%%2-)qNq3>{5e!RbN{EAoVd4)6 zZmnsBVv22ES>8I*b~d{Rl%B&MXP7!gY**A2as;F3J#yi(m4Ulk)z_c>jR}nuum#EG zi}58l?v-D2B^@g5u@wlp%M|)Dlar1x&{U)Za+8abbk)&1P3(QL)Vk>6iw0Ck6b1w8 zhp6XfWOOjy^Jo}C^!RugONk{k<(@9{r20)YIUK(TkSWOd=!=gW-C)`Bne2~mdp4Fp z%y$O;y?lRM6o>=Y9&7)=V1doF$hRz?i37jmLnX(AN(FTZz%%<43^07J!n=Mn)V_oB z^&qIn*%3IxOoNo$kz3!vh`g^5o;z74PCwaR&_d|Hy?Pq;@YsY^U1k1SI8l+M!n1F9 z`-{`jvtnP?jS=h44@-e(e*;ud%V~ z-k$e1mRu1lm3QIN2yvjwg7&I3?tkrte>6=Rkc<22+|WtJnJ+8K$dV>g_!(ozwt$?$ z#R*J~ON0@=suJXv*wS0t%n-RHboW{+8>u0$xGtV`!aU1F`oev1z*&0 zBhoxe=%w(-K9x{@5zbD``Zs|uHdUOM)D?Nqv(+&JMBzMslp!xA!9U0tFdYagfQ0lX zQ`|XMB_Ru};oItH52$}L(+7?iF&dtlt7sER`28PjCjHmQo&zUjWz=(?NsPUe#<e#&CM}-MuIA+d=Q`c}g@YS%Z@2eWiI=LQ;7DZCDjV*|eF(K!)r`Wj@I7pX5J#b2Ta( zK##R1V_&bx{}=mf5qkc}$$5LyqxW9`(WW#1oriZOLgMGGe*s2pA<@Fp-gg6e^8^(C zsOPdJ2Qt9#xmc#;DPsuP{|}7nFI7SJ+qapxJ-x~yfBv`7s_1Cp@%BkL7Q!$8=EutY z!v7Qc|FB1ShVqV8h<(}r6a=bksXnaz8ToGtXLGZEp-NzFbQCCV|H$k%`u_usnwe5l zaxi`IZ?*)QQjcdN)=-`8<~5JTn3#?ItG-{se@Yv%cZwA{)oi))=vaAFPj@yMU&$Y~ zGL%aB<$zIZI?4LL zi@n~47!J+tuN_Wz%C2~NTEzK$@YU$@VdKulRDL6E@G}^4Nc|)|+V}?CI zDWC9xF9t6)@M5I1!_qP!j)Xr3{K*U~t5OC`BA#M=?uCRR$ms8gdRIJ-5gzZAbJR(F zw#4fC8CUkm$M;TV^4NkZ*hY)+nb1N;XEFoqv8oJFQzaxz2_CThV=SZYaIhY~m{IZC zFZA!q^D4^Cg~qZMCQb!newP%){8mF`GTwB=K)gVYe@riOM{?!H4EB=;j~qcUR{V|S zd(QjJ+dytnXd+}`{(lA7Q zVXU}tmeOFC92HT&T)g;To>L6NJ99|>f8j!p#YE0n0F}yq#K`Xb8wjglclMlysH_!M z&Pu;73>rGwuAt=GfRDC6@FQpej8)P+-t8|2KgDp0gJX-XWetB>9FHL{WF5W6r`3{KUlQ5h}?^v_Xx8Q(^EJbEQF3PTngb-J`(xHuFo(Q zS1^%%Ko)0Q9!9{hUK==9vfM6D_=Ix&myzn<_trROHyO%hBa%{+)T+dxetT$W`{v{B zA<@#8fj}BIYMxe{L9wKEUt)aE?pmgg0*$of_*haw8XPU}SMI)B@ROObf{IrK)PY^k z;4jUg&669W*V`F&b=(l*Fk0Sk?HPf$OV2aAeH9MPS;MYhZFVUQ__6|YRN?_CB9$pu zpMq9OXdn{$5ksvrwY01rC$);oNVs^2WexN+JfEI;FTNBcqZadylnBVr{H#nXj){qw zdL5^hH{q6`O)3>fF|G;X-6C|E=L7P8vj8L+IXs@C_u&M#u@e`U2mrgq1BdrDIa)kh zG^)pSc>`*giCceb0T~cT0{W3xTRq&D9L6?SaO43`6wN|~Y^>#~P=Hp$-%DP8K~p2$ z^TYARc&{`~By6Ni)ND*Vw6!HQ1q#xP%?osN9X=hrcpX0cHY&H>D#<$XNq|_53gIuZ zaAL#)X>p0lTBRZ>vQ?_M1^h7_ZqIJ5s~u+Fb^CJUE$qmWC-r$9Q%B+}Dn473pV%jM z@2hh*d3Y>2jBoISG!PM!a*^{f5mPrdwUucV8rla&%h+`I@SdNa^LE%I-C5&U1W{er zFxE&o(d^G|D7xFB7|RvgLz%LvrYCRQfVE*;^yd}n9K)|;=1YU~V?a~kkd`J83*O>=ddl$RFWgG2BE;X$DLU6-)C6tqN+=N1#xQp&=R7Fr$Z@(6;oHD37<{{TZVabst$jZb-@qyRy@>=6t5YfyZYJXLWN3V+zB|@pSmS0D@li=-f@_-~i-E-boOU>dzKh@>-b*2zm-rU_<#2&qzYh+ou z11>UMqn+*966KeBZBM83C7Q0H4lyG0hc)vJnpPV0yVP2by~Y5qv4HG6DZUQ9(G;;XD@u zW+4pj1UeTK^T}HA#;;G{oDiCwQjDbYjXuXWSLeG9JR#y3xAkeCa%(~|bo!q3+G`^|XDPM3wf|I^3>@x&w+Y3Q#B!PCwJks(e4TWHsqb^aNO)PZ5+lZc^=<1KBLjM~wNtH!#6J zX`)<--e-EGM1{b8c-fUW^HgURk?fBg1q{mQHen_xQ_bdRQKNUF;dZr2makR^Se&y(AY z8w4N$zb-i`p^bmIH`siV)^IvzD^X2u*rKzm!gvK|FqzDzT|HfmDgCf$@Xye&aBU=!K%$pw~sUCdW z&+Rn908;f>jbF`mHpm7yk=tAX@gJ8xVO+*mi@z6mJQ^K~^F-!W#8WWyA%>B zYUE}2Iybf!&rcuw?#9~pmaK`hdwKd8(={7KBb_zYo7D~0?3j@SyEE-}Z>`E7?)(k{ zW0@z73b(9L8i7lTJZxpxME4W#ChyOR4@`;Qu+eRG*|>^Y*jbLIGfx(g0X+gHIuH;*#OvFf~XvI?af>Bj0WCc1cQF1%gLX zG9wN?S_}r~QNprLkY)DOfvHG}KEx>Ovv z1n{RhS-e%KOS0m$W-A(Ru6P8oScCG^XUU>oz1K5OgHmb~G*n`dJx>gT{y6ALMy2r6yXPg9JtLrNR0tBO?2#_2wNq|)Jv?EeYzIP*U20&A_l41<1H-F!>;y