From 3d0ade406ada0ffae57abf3b1437725c5b636cf9 Mon Sep 17 00:00:00 2001 From: nicolas Date: Mon, 11 Dec 2023 09:38:25 +0100 Subject: [PATCH] Tmp. --- candle-metal-kernels/src/lib.rs | 282 ++++++++++++++++-- .../src/libMetalFlashAttention.metallib | Bin 0 -> 102792 bytes candle-metal-kernels/src/tests.rs | 63 ++++ 3 files changed, 318 insertions(+), 27 deletions(-) create mode 100644 candle-metal-kernels/src/libMetalFlashAttention.metallib diff --git a/candle-metal-kernels/src/lib.rs b/candle-metal-kernels/src/lib.rs index a0b852a4..d298483b 100644 --- a/candle-metal-kernels/src/lib.rs +++ b/candle-metal-kernels/src/lib.rs @@ -1,6 +1,6 @@ use metal::{ Buffer, CommandBufferRef, CompileOptions, ComputeCommandEncoderRef, ComputePipelineState, - Device, Function, Library, MTLSize, + Device, Function, FunctionConstantValues, Library, MTLDataType, MTLSize, NSUInteger, }; use std::collections::HashMap; use std::ffi::c_void; @@ -13,6 +13,7 @@ const BINARY: &str = include_str!("binary.metal"); const TERNARY: &str = include_str!("ternary.metal"); const CAST: &str = include_str!("cast.metal"); const REDUCE: &str = include_str!("reduce.metal"); +const MFA: &[u8] = include_bytes!("libMetalFlashAttention.metallib"); fn linear_split(pipeline: &ComputePipelineState, length: usize) -> (MTLSize, MTLSize) { let size = length as u64; @@ -105,6 +106,7 @@ pub enum Source { Ternary, Cast, Reduce, + Mfa, } macro_rules! ops{ @@ -179,9 +181,88 @@ impl From> for MetalKernelError { } } -type KernelMap = HashMap<&'static str, T>; +#[derive(Debug, PartialEq)] +pub enum Value { + U32(u32), + Bool(bool), + F32(f32), + U16(u16), +} + +impl std::hash::Hash for Value { + fn hash(&self, state: &mut H) { + match self { + Value::F32(v) => v.to_bits().hash(state), + Value::U32(v) => v.hash(state), + Value::U16(v) => v.hash(state), + Value::Bool(v) => v.hash(state), + } + } +} + +impl Value { + fn data_type(&self) -> MTLDataType { + match self { + Value::U32(_) => MTLDataType::UInt, + Value::F32(_) => MTLDataType::Float, + Value::U16(_) => MTLDataType::UShort, + Value::Bool(_) => MTLDataType::Bool, + } + } +} + +/// Not true, good enough for our purposes. +impl Eq for Value {} + +#[derive(Debug, Eq, PartialEq, Hash)] +struct ConstantValues(Vec<(usize, Value)>); + +impl ConstantValues { + pub fn new(values: Vec<(usize, Value)>) -> Self { + Self(values) + } + + fn function_constant_values(&self) -> FunctionConstantValues { + let f = FunctionConstantValues::new(); + for (index, value) in &self.0 { + let ty = value.data_type(); + match value { + Value::U32(v) => { + f.set_constant_value_at_index( + v as *const u32 as *const c_void, + ty, + *index as u64, + ); + } + Value::F32(v) => { + f.set_constant_value_at_index( + v as *const f32 as *const c_void, + ty, + *index as u64, + ); + } + Value::U16(v) => { + f.set_constant_value_at_index( + v as *const u16 as *const c_void, + ty, + *index as u64, + ); + } + Value::Bool(v) => { + f.set_constant_value_at_index( + v as *const bool as *const c_void, + ty, + *index as u64, + ); + } + } + } + f + } +} + type Libraries = HashMap; -type Pipelines = KernelMap; +type Pipelines = HashMap<(&'static str, Option), ComputePipelineState>; #[derive(Debug, Default)] pub struct Kernels { @@ -208,6 +289,7 @@ impl Kernels { Source::Indexing => INDEXING, Source::Cast => CAST, Source::Reduce => REDUCE, + Source::Mfa => unimplemented!("Mfa is not a source"), } } @@ -220,10 +302,20 @@ impl Kernels { if let Some(lib) = libraries.get(&source) { Ok(lib.clone()) } else { - let source_content = self.get_library_source(source); - let lib = device - .new_library_with_source(source_content, &CompileOptions::new()) - .map_err(|e| MetalKernelError::LoadLibraryError(e.to_string()))?; + let lib = match source { + Source::Mfa => { + let source_data = MFA; + device + .new_library_with_data(source_data) + .map_err(|e| MetalKernelError::LoadLibraryError(e.to_string()))? + } + source => { + let source_content = self.get_library_source(source); + device + .new_library_with_source(source_content, &CompileOptions::new()) + .map_err(|e| MetalKernelError::LoadLibraryError(e.to_string()))? + } + }; libraries.insert(source, lib.clone()); Ok(lib) } @@ -234,19 +326,41 @@ impl Kernels { device: &Device, source: Source, name: &'static str, + constants: Option, ) -> Result { let func = self .load_library(device, source)? - .get_function(name, None) + .get_function(name, constants) .map_err(|e| MetalKernelError::LoadFunctionError(e.to_string()))?; Ok(func) - // let mut funcs = self.funcs.write()?; - // if let Some(func) = funcs.get(name) { - // Ok(func.clone()) - // } else { - // funcs.insert(name, func.clone()); - // Ok(func) - // } + } + + fn load_pipeline_with_constants( + &self, + device: &Device, + source: Source, + name: &'static str, + constants: Option, + ) -> Result { + let mut pipelines = self.pipelines.write()?; + let key = (name, constants); + if let Some(pipeline) = pipelines.get(&key) { + Ok(pipeline.clone()) + } else { + let (name, constants) = key; + let func = self.load_function( + device, + source, + name, + constants.as_ref().map(|c| c.function_constant_values()), + )?; + let pipeline = device + .new_compute_pipeline_state_with_function(&func) + .map_err(|e| MetalKernelError::FailedToCreatePipeline(e.to_string()))?; + pipelines.insert((name, constants), pipeline.clone()); + + Ok(pipeline) + } } pub fn load_pipeline( @@ -255,18 +369,7 @@ impl Kernels { source: Source, name: &'static str, ) -> Result { - let mut pipelines = self.pipelines.write()?; - if let Some(pipeline) = pipelines.get(name) { - Ok(pipeline.clone()) - } else { - let func = self.load_function(device, source, name)?; - let pipeline = device - .new_compute_pipeline_state_with_function(&func) - .map_err(|e| MetalKernelError::FailedToCreatePipeline(e.to_string()))?; - pipelines.insert(name, pipeline.clone()); - - Ok(pipeline) - } + self.load_pipeline_with_constants(device, source, name, None) } } @@ -706,5 +809,130 @@ pub fn call_index_select( Ok(()) } +#[allow(clippy::too_many_arguments)] +pub fn call_gemm( + device: &Device, + command_buffer: &CommandBufferRef, + kernels: &Kernels, + name: &'static str, + (b, m, n, k): (usize, usize, usize, usize), + lhs_stride: &[usize], + lhs_offset: usize, + lhs_buffer: &Buffer, + rhs_stride: &[usize], + rhs_offset: usize, + rhs_buffer: &Buffer, + output: &Buffer, +) -> Result<(), MetalKernelError> { + let a_trans = false; + let b_trans = false; + let d_trans = false; + let alpha = 1.0; + let beta = 0.0; + let batched = b > 1; + let fused_activation = false; + let fused_bias = false; + let m_simd = 16; + let n_simd = 16; + let k_simd = 16; + let m_splits = 2; + let n_splits = 2; + let constants = Some(ConstantValues::new(vec![ + (0, Value::U32(m as u32)), + (1, Value::U32(n as u32)), + (2, Value::U32(k as u32)), + (10, Value::Bool(a_trans)), + (11, Value::Bool(b_trans)), + (13, Value::Bool(d_trans)), + (20, Value::F32(alpha)), + (21, Value::F32(beta)), + (100, Value::Bool(batched)), + (101, Value::Bool(fused_activation)), + (200, Value::U16(m_simd)), + (201, Value::U16(n_simd)), + (202, Value::U16(k_simd)), + (210, Value::U16(m_splits)), + (211, Value::U16(n_splits)), + (50_001, Value::Bool(fused_bias)), + ])); + let pipeline = kernels.load_pipeline_with_constants(device, Source::Mfa, name, constants)?; + let m_group = m_simd * m_splits; + let n_group = n_simd * n_splits; + + let a_block_length = m_group * k_simd; + let b_block_length = k_simd * n_group; + + let mut block_elements = a_block_length + b_block_length; + if (m % 8 != 0) && (n % 8 != 0) { + let c_block_length = m_group * n_group; + block_elements = std::cmp::max(c_block_length, block_elements) + } + if fused_bias { + if d_trans { + block_elements = std::cmp::max(block_elements, m_group); + } else { + block_elements = std::cmp::max(block_elements, n_group); + } + } + // TODO adapt for f16 + let bytes = match name { + "sgemm" => 4, + "hgemm" => 2, + other => { + return Err(MetalKernelError::LoadLibraryError(format!( + "{other} is not a valid kernel for gemm" + ))); + } + }; + let block_bytes = block_elements * bytes; + + let encoder = command_buffer.new_compute_command_encoder(); + encoder.set_compute_pipeline_state(&pipeline); + encoder.set_threadgroup_memory_length(block_bytes.into(), 0); + encoder.set_buffer(0, Some(lhs_buffer), lhs_offset as NSUInteger); + encoder.set_buffer(1, Some(rhs_buffer), rhs_offset as NSUInteger); + encoder.set_buffer(2, Some(output), 0); + // TODO Tensor D + + let grid_z = b; + let byte_stride_a: usize = *lhs_stride.get(lhs_stride.len() - 2).unwrap_or(&0); + let byte_stride_b = *rhs_stride.get(rhs_stride.len() - 2).unwrap_or(&0); + let byte_stride_c = m * n; + // TODO byte_stride_d + let byte_stride_d = 1; + + let mut buffer = Vec::with_capacity(b * 4); + for i in 0..b { + buffer.push(i * byte_stride_a); + buffer.push(i * byte_stride_b); + buffer.push(i * byte_stride_c); + buffer.push(i * byte_stride_d); + } + encoder.set_bytes( + 10, + buffer.len() as NSUInteger, + buffer.as_ptr() as *const NSUInteger as *const c_void, + ); + + let grid_size = MTLSize { + width: divide(n, n_group.into()), + height: divide(m, m_group.into()), + depth: grid_z as NSUInteger, + }; + let group_size = MTLSize { + width: 32 * (m_splits as u64) * (n_splits as u64), + height: 1, + depth: 1, + }; + encoder.dispatch_thread_groups(grid_size, group_size); + encoder.end_encoding(); + + Ok(()) +} + +fn divide(m: usize, b: usize) -> NSUInteger { + ((m + b - 1) / b) as NSUInteger +} + #[cfg(test)] mod tests; diff --git a/candle-metal-kernels/src/libMetalFlashAttention.metallib b/candle-metal-kernels/src/libMetalFlashAttention.metallib new file mode 100644 index 0000000000000000000000000000000000000000..8c8ce69215b7d02de98c5921ee6e814afaec5a9e GIT binary patch literal 102792 zcmeFa3tSUd-Zwszgd{)+;c9?@6F@~loB(nWFB5_k6>U)3qO~RACMaOIh-lq57j7!H z(NasT-3={vX|-*rm!hlNgrG%BZE9_oTI&-mTI$-CR_ofW+x~xNl0l$$_j&*C^S;mf z?%RBl{ARxAH|PA${meOMzOytZD-GAOFdl~Kyf6%F-igCMhUL=#7#7QN9vpMY4>->A z?QjA55cEHW?RTBqOPtpuwWr{>OulrX04rEuQMxu?xLC=WqVn?bhP>>4v*qfO*%7PXu8R-- zX382<%#p)?P`Q2emD~NBHW`EKoQCwdY~Z^j6WF~B(Pob|WFfK;KI09k{-q(IfBGm# zQN4Lif!gaU^{Ew|d0n%U96b@z#3_%=`sB`n0H-1TuIvyweZ`a|&VzdyVpclcoGJ2} zzk=85I2YYIw=3bEOR7+$=gx~MQuT~Je6tEOW#y9mSc%hngH z%d6I?tE!8z*~J)MQ&p<2ShJx_U4_kFgW+qc)^E_jWtAI>sy1P>E1lltsjD|t6y%ky zR~J@evp>SH<$3w)igkJGi;Js^YOv`S!GiI~XBd_~9=U{JCr|Z4V4jV~#yBJe@g)N~{loz=N^nbh%W}*+S3xbSfZAMCQB7} z2J~)8M(xS}mWBXzDyU)mHF?Q0$rCk-{_yLDhVy6mq$wxb_~FLp$~xXAV^fuww^`M+rG&kyu4%J|R~5ypDq&Y; z@*c~C1ly#wJZk6ZlzI*Hok7ZI0q?!jKM5475wGs>3c|VMFbbQ{Yv* zp|HVNtIPBb@oCqHg%#{AT!{Y&#TiyESjn(dUaJ7)}xku-0tO&hh5Tzol#4M2SbP2 zy?Wcj`mK^P@-P&TeDWFG!q#>odo(0UlQF64iqVT>N`m#mUuutXPnNh%p0jgiVA zy@PTO%ft@r#4Ob(_G!K7+#8c%zLd}BBf{?p?<=J6rbYOJo%`L0@Oy`FNGTk$bBBz= zTSjzTPI`Y!c;B>huVUWicHto9J&d&}lFhcct!&DunA5LFZf2*pDnjgvsICmk!}grSLlp>98rYBBkFp_Ygj}Pm$7syh=Hw zOR*?Y>)9!HFaLT(SZ)-4trFfCXM^1^_q-wnt4}?y2U-nb?%~cS;lsWCk3Igd>CHj8+*a7yy3OuTA>w;2xL zbi5B&$Gct~?|KV-7~|C~0);;6)vfW@-lC8Jqqg0e;FXc~ia zccl?oTct$G=pWPw-Uf)vn14{)09ToKcmHYMG_-k@7_W*#-l9_|LUa(|8<-v4fs$Io zt`;}tK~s~-+nk9SkQ#eam9Z({{^SnB+-~b!s1A0;oEF=x4nt}OdtTH$B0ZBI6t_IR z+0IuFEY$2dv$E5@a!;s;WiCGjzfB21=g*L80glyH) zh$35A$tElNb|7t3Q z*IGxtBJ}^7`cmls4W&yAZY_BUKN>8?O(`s@groHpestrqQZ{S4zoTs^S5Mz4bwNGa z2VutZwUXXJNwGgZuUU~?Z)-`@@;24+s@p0BP(7V3M!X;`tj8L9wp~(8y3uNe5UBFd zjI^-l_1CpPSM?a5w+TuD^jR8qZC#UwXbS505neY62TAV=(iI8X7+9nNjh$>!f6tDKS8+f}>+mr|O zTHI9C*0iZj7X$6m*AC(L7HDkQ`F+qOwM+Wx|BGtL8E0>+=6-C1mQ9h;M*nx)l3SgP zwrv&!dPkpn9$I6Y;GNPYO>R?_#!yp&8ju=sQ;n%9=PI(U!!s`$V zA}aKv6*|zcb5hAgUD!nj>I$#nVa4KXXo2VKHKg_{BV7?uHEpcp)wN}3-86cm6bM^O z_jl6JOV-fGanzglOG8K0+?)nOa=$Irte6K0a;InRkRr7iPHinGJs_q{8^Ipvj$ypf zrd$0vbS@U*6{YY78DTvq4;vbsG%y%?2@J>#N;Mni9WsPKYurNjFWpwus-I`i?aw?^ z3BwE(uM!e4w<+k;Q^M;i@2?!(5!4BKUm?72DTNgx;pe0fPH#B4(9S~-c$985{jO{*+jtD3u;K*9=<=Tp=DY| zMxRey;rZkPzr(QY`Rb5E@@i$7BuK>Tjq)l?SXvPEJ{wz+zJz$NsHz$!cQzK3sHQXASYNIwEh~yGEUH*vRbHJ}SX!;lFQcQyI2Eg_Yx2s~H6?jP6~Az+QkNGwBb`@V z1d|gko&5Fd%Ut*76Xz+$PEp*QoftPUG4WPpLv_jes+w_=665A0#!W>yXCa)Sq6vug zj+5(&TbzpE&#B4?d)x%8lTkle7?`zwlYQ9far;KC2zX zr+k(b#AkfgAc+6qvv6(%pO>H~i3onWxP?eABlTod1iv1R1*iBd88|Csddg_!_fX(l z5jYFLIS5qCNxg_4!S97*juk=yRU88D096gBt_9Uda1L%EvMNYDkrlyb&9>RyE(vVY zv{r#aEXXHqB5u97pRAWlx$y$Ess+K4LIh_9^@}3o1uq2cg?>!WU8CW~bKVu;@d|rC z)=zd=WBPjp4Z+Fp2b-;eg(1nt5RfNBl6@rTT#~#C$&->~zfg1@n%oMFkiac0=fN|w z(#Pf(99&mj`dGyNg6L{lfoJOTA#T!!Vck%XRq~~J@e=5XNInxyL2(bef8kW=7f}IyWUFd zPB`-B8%fE>4kf*HG`sO3Z9{JIl$fPwi(bn*`&vtRt19rU2ly!4Umfxic===SW7W^U zv0#>FEJ?avNd9ldEZiw#|md)OMTx zs^k>hx5-gA2uNz2um3@_!#o}a8JgnL}jH3Zd* zIP9wkN&_v`h(+v#rh4ynK{qLvRznkxsE zOLWPi&&~TNRwZ!?MT3sVRFF(O5E!`_!_ITVMIp^p1<}(im(>&ExobzbHZOfpz1&Uz zD3RZ7&16FjAN;a$HoxAZIkY180wHJ&;U@77zN`~eZ?jwq7m*v4En@#Pz8=PCb{Dis z>w;e5ak#y78sF1aWE_0dgmd8L$HBp7Xqy0w6ND zGm^n)0*?G)B#p!Q{eUAo8Oi8x03787BN2{4-ZLL2?~wFnK#>0mSAf3nJ@o%yz#oD1 zdk*3bb>}duUwNx!@#97iu}gEZ@-h^8i)t1MuDO;bfF3T&ZSzlG1cVE{1Y4WVR zyy^}4d5iMo%tBo(J$yw&RWv|Vm%*4juXsZREPJi5$QvDBYpRM0OABBsJhn6$>6TXC zL#eQ+xU>Qmzs84j-@>I!oq>#5y^W6}?v&wdi-vxl-o!#xr_t{LWvVOm>H^kaf>5g=kgxESUH z+(oT~Hb}5-p$FNyljFJBAID6661lQ|kv8fj%)oz;Y{obS{>Ly6?qTECCpMeidE9gU z**!x)2ZT$lX?NT{;`?1r_?g4;{Pz9=j)h2@B#01Xs724d>-O|uop7>FOBGQifsw%( zq82qr9>sbk{pC+ProUK!`sXPygzjtq{U17;i~m0D4`C4)e|8H|V&Gepds;ILAC%-o z@5tB}C)&?{)%WNb6)EOF$gQ_hZ^p&U<-X_@0XFF=e`{^TYH7t=pYEtx*!@ZKeXqn3 zg86@cX6n8V4kQEzWW1F9iojozfw}D$*#vj66L%#1|091W0Q zJ^=<+{faq<*eQEqT2e93Vw=?pBQThAFog8t9?Ltj@F}u2X@*Vo#C4^Ko^~*yq1!yW z0cJC5i6+-=>FowI8E$N{?5(1Y^E3;Wq@DB!U^D>qhT#!uuPg_t%v0 z3I&tq!Y@_cU&67Sdqd>?9W`zuUYmmEykH)mlhR?E>zuwv^Z1H6W<^MgBH0YC?9thL zwGigtzk)Rh%=;Ty5}?qO4TW@I?z7c47bfmu@*QU0!7EshK$G}zdvN!1dNBazbgusT zVt{VWc}1#)Ud%wN8e^9+&;o!%IOzB%>kv@WVCJip4a*qwVaETsZ7#i@5dvn|*>fT(TZR~kqiP?shG)7le%7|@PW}>JOiwtQlz;=0Y(y-Iq z@aft@cDtB*QKa= zg^A)1@fhl`)$|iMJ%N3U854<*7&TvIIe7rf)vh&E`dC;wA^hveg{SgL6n#SLXbyd^I zS-sGzs4ncBTJpFw^dd|#$iKv;X=#11=$%%d!M6oA$;o~p-+}E5Y((V(?t}*LN{Q)w z54z)r53GR99q;-9ix+j6M1yGfScUwsQ&g@-J-WPN&uY<`9X$jnpl>SiCox-XZIHuV zD*Dz^N7Y z8Vo@<(fUy3kD2-^PlI!%1=-f*+xoB=;sF;x&6 z4f@c7E#rdzQ~vm%10%);J$%?-4rx9b>!XL=J!1uFOp?AhYoQXdjES}JiSltlS7rJ5 zuO1WY;^&RJ(7z-CjgG)PB1|MMw)@k~f?SCXDf;wsmwo!O1zb$ToA%r94YI26xT zwKqul7~w1rum1(}7<2iBdG*Tld}yAWIby%9?hW^~jV8;j^-}+xVlI;-g?e3K)G-q7 z?wo1P=dn*KYt*Bc8-ux?CMYeVSvE(I=nUOe8qMMuEtJO59IurPj4Mu`6qv`E=W$xa z9$OlFpO&kmj@~sODu#U~w@NK$Va{w-KEiEiAL=0)1BerXsDjfCkUmWGgZvwJdD=z| zD2_L^OqtXu#k2&ZuwzqNy5|h3_O`q4tONBH>8@?uM_HJHnGUdu+?TQwo3+jY38fAb zhkvjZx=yIeO&Aqq}VB4cFwu1|;<3x+pBNhA=W zm!6aGc8f7wh0f52OkaFO70)K`U5rPH?{ioF73sUX^PSX91&Qrmc#HbyiZv|xRe7`wjOkF{X_xO)axjP;Rh zlBN9VXWl|3t3fHevO>#p%O>_IxeK$Ylghw`Y~Ov^QTvwn;rk@Ye*B6=c^P+0QSM?d zOi^~Qi{>a_WjCZMkFigtDos1MQ!bQHPs8#n@g5)ZDaJ}!Le9OkZM@iqRZpgloldAv_IXscYw|d$dSA;mpT{At4j;% zOmaFnGYu(@UR$QFD9S4>T)J$DYwNvh*S%}QeK9-?09yoL!#&z-@0{_yYyQ~D<+~Hn z<7Ta$Tfdz<=$)JAnSJws(YuZ(PKA$~Js&sQK5lmYe{6jpl8xRx0BuL)ThIW&ht#XC z_wdtZ9O&iit)3D5j9D#&_|Jge>4Bv?Appq$Tn3=U#@Fuz4Q@&cQ7?~zC3_1$O$3ft zz_AS+D?wWa$0i#;jRgElIPSFZd)j#s{HGVR5Ry7luNkoN(Q1-L*B zp4kBy*+GcEhn)k35&Xt^fDeFWpl_AwRSs}gfpvy$h%7w5)_P zyLiC~v(75;jV)O>IQP#YPqpZT{-)=Lvy}O&9k4ATxvY?6vxTc3ZnZ`TyF}d9pxiZJ zxG*HRpLkDyEAXrWp&vwNzDb^DJ;U7{+^-tw309V=4qF8oDc$(@+(1dc>UNq$xn6bJ zDp<0z?c~Z1=hFVd)qN*d_mkIIEA1x*4JjQ~$xu&9ztW@#w##Go+5`*dbjU-GBdF6> z%=@!O81MgA^JltbaB@p6tFyzFS;eb&g~?-p|U zoB4ZuKJsvW+2bap(&hc>P-}09 zQNIX5i^}$CeIeMLzUx&UGX}x$=x@1yv>@mM-}PFLSOm@L+Q-U5!DW@bO89}31vYc+ zcLc}FY~~ZEvieC2x2(Tcz*>5~SD=VKj*$LS2vLO5C+9QD;l$u{xp=C)Km8RB!7U-U zu-D`6SId=5L#NI?Aa`-lvA_7rqEDUwFIzs)h1l_6J06O@x)#pV z;qb)7KP5hUBMA2A<1nr+-;-uk>{>t6fT7aUtx3=CYqu=-rznTx|2?wc>~ zhW#Ux{uFogC~O?bX8j>~25cH>nA!i;f!ON@uD|Vr@0GVYrd6QrBSDX?Gb|~H z{&BMbZ6blaBiMGcX!{DZH|66qLl@5sWVjDsK7f2YIOi=H+Kci=SJGQ4$H=5(N3+ex z%-jED*?!!<{e)xtpGUU8jqR|Dc6>nWI7jaIklJxxx#J@hy+7mMZSRS{IaE9#^eE%O zePe>JmYRa|Ykak^?c}yxEJ)%%A>vNqY{rSsDVXk2mMBB8P0F!q&os2bZW7pY!d^Vd zwdcfC+wRoH+)Jh*6fFqliVY=2@ zqp0NAxb>+VHhGaUvBX(P#C>p4FI+w(K4r!SIc#c&S`Z>yN@D4L$Ug`2#am1Li18!1 zCq;S{wUfO764NZ=2Jj&<8@-y1e3q))+U7<^89x8AaTUK_&3-Xbk9BGX1{<#g_Ncnm zmIg9PYWKvz(gW~7W3Mj@j=Iwc>?VEG3rCTD2(!aclK90xn;Jv(k6?x`8|D01k;RO+ z`{}TY>gHerIIuKh9PmILCP*^r(zqKecn-#!^O=aUnel8IJ|Iuy7TNI}3c*7v;P5-R(CtMmpjj^r5t-m*t1e2l)f%Omoln@foKgWx)q}Ai!Am(TNXxx2hZ(_b zfJ791)*B4hAMxQ>oLhp@lhr)XE9c;3zLGEHbF-HrE7T!!%(+!a-r=_vycoz3VPZ5x zp9baRVJ^|OW2QpXe`>d!UtI7m$vN0nMCd;mi1W8Xk!zugTpSd+vgt;t@eM8S*f9dgqRrHvjXdcI!OW7QCx;JKVL;j7n+2yBarh*FFnAhU zx7AZs@GikcKG+E@*bvNltE)(9aDZ>?L>IvmFa4(jafjp@@HvTJQSdIsvzC_y-6Xmp zl*ve$Xh6#Qkh0YuDep}K<;EPOZ1zCPtE5Od8Yy?u%H6cjlaTUlS~-tj!OY933Flrl zKaRE8YU?G$TKUYF#d=AGT#%-Pv_kH$!7g7vPN|P4v=0G*MMfF)QM&M`J3njFh^6lS za=Cvd&SE`-?>yzpQgKv5=_HL>h)=3g$Kc{qz9X<;@`MGgn&=kCV;#Xb(IdwBdg7-Y zn6aFz;%F=4G3;BoBTWUr!|&@eW(#|x;ShV|K^*pz{+-&MT^gEvO6g z@@R*<<8UXQS3_@zC|kc?6T6|JYJFK5-1C%`VbPpR=BEPE)(^bnVNJST&@d8mt1Auj z{xMft-}Ji!qF&ql2l=0?Cr5L>Lu#9l+FVf6gn=4`)H;ybd8GDNq;?CbJ%-fgA>WXi z*89f+a`@D`{@zAu(&d7wiGuHZCdYTFJ0C^IH|20F9oXSzTL=9+7Vq<_YJ@etKiv#N z=MK{ETMzCmQ%_d?4DoTw;}B|rMfGfPe;l-$C)rj+wyAdGDOqa4pk(tPYw7hJ79s8y65lAg z*~8_3E9& zt7gv=l;1Sp{U=xlBe!{VcPnYztN%p5*l$_k3FDu1FaCy%FNdM(X_SP=i_if=F>db=$H$+C4Gq% z?57h~@XxUYQOlszEWQ2(=J%qHy^WvlR<4^Q&fjec?i5*?w!QAOYc=(VhoE5h;G5#; zVHP{q-dpRjgj}Kx$e#rDBCC*mS6O7G8mI*y8&;Yoi8VTDjHb{5Ir2v>XHw3_0WCk6 z)A{)ICh=8IT4(8R5D)k?dkW+m z^2Ll}VL0vB?tY=*$w5)R;WkGE9dH=V4i2!jRq)HwJ1tX(Q0cJ;;?H;xkZuJm?x7it zvd`=hIwiiW57o!^Sk|gZ$g&6AWEVxNRph(&O5C~~>PvGtODrSN=HBGtr)PFk*g?)` zR0E;eVt!9{J3|v*EqGhXjjHPI8t`Gigv&`0dtUZaswdUChFegHJpM*(@oAQGNm$|+ z7ZPU*=Di%!OWl4sM1Tuo(;s}SacYLZ1Ls&RUaN|_y${~U-2($X4g{8@eWsunn?s2N zY4`)*MJ#AZThJ`GKV_PCr!ao&4U3QDho{x-a1mddl|gWW28j7z7Z~wf-gnw6H0*Z4 z%OR0g!Lr1${?`RavNOO>=JQgrPrhhbr2VM}KFQM0Kf_s= zTLk&aL6Mh$(gMpf4^(CdTDTl5B{+35{n$#@dee)7JMAA%E(f{##m^`2xO$Q-DMHyN zND}D8#?z<{8KYHDhtlI&kgUUt#pgxCF}C}-_+SeK8F7Ne%}Hcc>ZabRCVW!ehl@Ec z>Pf4oK(++qvIAIaf42i#DA&a-U@lp9JC$ z>@N3NmHtVfb>Xg+`UG(&!M-gu1gAFXJmMUa<2t*lP2JL#!DiJm$i0G4d@2a`B-EJ6_{Hu6DpR=PJeA8{JXuG4j6(}em0KKKMSonkJcJ);(Fvg4 zDfh8M`&`G4xJGfIWjMkwgEs!Sm0H+%Nax^2aJejKPI*t*64xniMTl%(x{F$C3|rDK zCzU%Z*DJ$c0V!QiMu3!dej0KybB6Ip=mZ<1W?|<-nNFzLP|a@)Y0*PpHKYw-9h>gy~zE;KF27eQjO}0`|Y3 z9)|tzJ!gH-Td9+rwbg_j-)Tez?o~GTpkSklhaWal_8(FC9GRoi7^KDS-)SGrgw~t| zk14KfoNqB`D>4a)?c>6~?MlS2c%x?HEjG7R0i_kR@2E^^+SJrs_Ixw8hJ=>z`WGX! zD*YsCNb`Ai_OJ&uHzRKTxxEs_8BT4%Ne-RnoH%}zxjCgVeB`-@>9WaU_4|(KWs(;I> zg{*x6vNlfJGr8_Q==qm9GpSAleNDAM?bp35=*Gdo7FSl@AxxI=tKt?%#p};~ZN`20 zd~%7%V^S#;h-`kWDb}sAa`}cPuV#wUuGA}|*403p-pz(l!iZr$pF5-e7604q-n1NY z$wG42FvyiRKvuN-XGlr;29zFqRn7m19BSNBAFyW%>9%|dWM>p)r<`zG|XN=23xa_n?i6<(5dlwThm(uQNe0Os@+(?k?N>Be{6{o=Zj^O zNW}5{nP^CgYgP?t#Vk%|^zXFxMUegUU@1%I&)PPHg6{J&Uy%8G)S%^38A)Kp%OM|6 zi#oE<+@!5_hfYR6m&Lv?4HdgU^9No~RlaOgJ-)5n=O?Mz*cKt~R6W(Ycrg)6_G)Y* zqcz4Ru|S=W_GirMZk9lMb?^gEGMvW_=-G0U)z$|U>Q>kZYHAf!m+v!XnYKyEGWV&H4WOe z5I8Er-&begpZTF8e3-+nS>Y_gC2CWe>Rna9$5dd9St_8%FE`&0ak~!&2huga-pwNyQbJ zhFWCcqZX9oj`u797aV~xnh0dUr_;4^XqY_@kA$lEo^$EOKK!_x(k<+lX|wH%i~2rX zMRDU$`@XQ#pk-}uDakRUj0BR%G2HeGQiSoF1rMm3rx|r<-4@^vzkjTRSQ0FeiVv zY7z|Wsc(nAH!GOR9PKT3S&+YA--4~DR=6{`A~ zYDhw2#0l9ZpIkw`IYksO;`=BJ2d6rFHnFwVLznik$W|*9`aS3`kd=KwjeMA!P4<%c zT7uvQ{v`J(e?S1Y-7*ee`Ed&V+_TWG@o^Xlj_6}K>;QdPzdr1G57l7WZ;ctT#wyED zsS!P-qQ0Ag?hIIO9)sRM_L+H^ylbXDm2>D})L#q*i0zNNL!)fgzE%O%H!F#ABd2V; z##?@cy#W7gG>lGGG^aHW6`&bQVZf$(n)>5a@j&{>v^leEo!Vv}wgQurTEnMtdQT4d zj&T;A`=0Fjc4~QgUZZ@(jc@b9%XH8=B;QF>vLonjD9_`R=w{(Qx?@p5TO=&t|L*eg zeVi(k$K)R!G0>6{oP(S}QdybuHlbxd)222yqc~xYLMKiNFN<2FZHpfX@S(?X(%vcPN+6uhuwr+QSh}(8z&M2U335wv_ahP-Qm~rHx z$BN?I@_Y=IhpZ3l24H?A4#(}Uxc3r%t+xG{ABodEKCA2@^1Q zO3aqD)^vcW>QpW9$&ol-%6QWpgM3B{|0{zNcGcy4ed&r>CE z-gIPWtTofkYO|`i_X)aXHb=aJF43S%a(Ns0r$jxCZ^Iw=I5+52Mp?x=S1af^+gH}# z`7ysBIove!odv}~7o?N*E4lE@bwYS8vR8hfSR6pe@>P7e71ub{C?G$tinl~3%XTH9 z{Pthhn62}JdkPHQaO_K@?Ok-l-##S2`4@uokJ*-oI?Jz^A5ERQ7ByBPC|ZM17VhZw z>sFDj2?FD~e>A|*%GRvYfm0v;T0?N>2dCY8UwB5kjI4LlJBJDRG2L&SgJD=2HGLY4 zcdYwSjhb&!P%-FKen5v0?xmnJ`iEp>$$PBJDb8Ed>kmEeHdMvR#-AK?YcxF_(UQ{W z<2h7jloDN&o}bbTgI3)e?7Cd`V{JudAhC#r@Yq zD}R$}w<8Ri!U5rs`cvYoUefM@{5K!#ueEcMe$N z{weOv49A4 z1dqRo32HPucF|*VRzSb+9mRA93|-|5q07I6`NI9S_L_V_*3T;ZDC#!`mdV{?J7e9r zf;BLF=S+bv4hj+_pr-s4uF~BHSN-hphUn8}$9n0NM=tn$gW5@Aql#=K*MFJ#*+{Jp z&D=r-C{}Kcdk|kjxqbd}%@e2%9KIze5a{ZT`s$r&%}!R%3x}KCYQeR@e#sAiKodvT z=Gnvw)y>!H(L4>@vtq)+r_j{5gzbAfWN%^XYXi=v$bJGrNzYdzUl~UVQ&( zeYM_^{;_udi6v0)QR{2r{_TtE`45Q}B=N29d&MO^evnflj}*gHgyWBp#6LhqQ!Hvr zr=}z5bAaT32gtAlpb9a+_voIn)pR5_`QZ=gYN`P@jY9PVlpBlhG#|T!i*AM;lf|Ma zE~Gu&{N6q~;F&38>pHdHjIhvMUN6<)_}RKBXNK%ulkaq0?#M=Z(uHtNx)?cYVC=iZ zInd7`mzaI`?G1y}4}cgE6Or@G>~F;FU>z zlfTcW^XN87ws?FO!4wS#DW=B&8qV>DvYv2u!(Zz?edxYI7my1-cpZDI(HeuVaw!MZ=R2wbO#Bup=h{Q~4BA`WP3Yb8 z-3-p&T{)M_lHWinZZI|BP^ph|vN--v>8O6+^Dqv#047@GN_O$v!2;Hs?@+`H%!BTFf5-s-I6}r@1P)F>euWRAy!$?sEEFWej&|wuezDlMA1FM1LP9QZ{GnGEL-VN zot0(#L7jCxtucKT#+*c*_Mzy|^4|`8ur;9_v0EH_*EbrZLKtPvT-7vtQK#BDJ(%im!OMKf>gr9>?^%7ie--8Y-=#L4rc9K1xO18=tX6?b5e3LjIQQn zpOGMgVB5jF8oyK@NQSScJ9<0(0xzny_bNjA>#ibnTMhBYXTa@^VcAeRdC~5nrXVrq z>H-1Dy_r5kkkpwLY{XwHhKQj?a<-sO#(&i-OJ+gyJPtpIdW*-rY~9Hi+x^;;YWWP5 zaN{lQcUsyL3zn+07snt=M_7<_x9f8u=~$26>Gcl`IG?Y*M_C^=bR-`Pl}`sl6}j4l zFmenmvGQzQU#L6^so@}cR{s4#f5R4metAQsGZBp^sRm0(LH2VeV-&E1B3qa8i$u(b zpc9cbA52@5y)6})Hh@esx~)RFx~27GO<+J@!9O)^r~s9rWOR$TvTY$V8o-{{&E>u- zR#Ad!u>I#@k8paeR`3W~t|HM2l!J>_psu3jDs~uJfpQqoN|}jXf%1nXw&jP?=!M;& zt5V}rGi?iyPTf}R{3C)n(yvbK#$#AN4#c%I^n!;6k%y{Dy5E%r-ATK)W2T{$bE~T; z)9ZB75y9&|u>OT7qBXno8EF0MQ?!t02|=N6iGbC;lXV<=#Vu`~b1_ad897~#KucrI zC{XQA3Idh+qII4fAmy|? z*XmvzcsN}-)tLf`e&D2AgO=FtWPv~2H^HAJI}_<;>mt?jC=tqOoeHFLaVgmBaok;H zXoiIP4`k2n=OGCH7yT?;A28!D?!&P5mx@E&E2f#BnyL5(#r^gR7O?sAh=zA|TYty; z`#_xieroc=jX%0yhqy1G6ACFC_=_%g6_vCjo8u#QytBK#5d}J|L%2km&-`So2q;tu#>prLPQTvcnXXG>Q?8bM4Kg)ZUJ9DsNIwYnc*6ppX zeZ%{lYonbjkUbkcz zdCOv7i0LZuu&6G&6qG9iQ0}~rj30UcjOU(1%BodJSs9H&U%P0iYhTH0bm+^^Kj<`U z3vy!DExOuO#D4=Be>DJ>dUYr0;`fkK=cnkl9W&jUoy)TFQ*@@}q9A0RkUsWe-IlVT zi!tDj;`eUfLCPOQ{=66n{#ZZixB)4@3K{>fC)i_L=erqYF{tcPh|A*eNha2x89s3g~H7L?}plt*|^z!CWkqz-`C}ZVv5(B_Gvdc_@1; z7(60I1=92m0THpy<1T*;7O7>Z!djl912xi@v=pL1r7?6Un}RaGhLn7lzNCbU%+888 zUMhH3yN6!yzA1?O+dy3R8qvTzyXEr@sHlG*l=)3pk>&%c3X*FFeTgQ4k5c9(-OaFv zxhY=g5=SYzM0Y=Z$rHhu-$2`hK9vN9o>6dee$0EB&e2&jTFAmyrx!4>-zm zM$$N>+g`vk0cRwmKjiX18mBjCqK zgmd8hX&25&20!A$8Oh+E06rCdjAZbiUHYRE`7eUK8&K~EaAPE`4|p2j;ea!e!3zOL zd@qb-a2?<%4H(JbF9446i;)a|1aLpV8HsT8K+$Qy5ibiP8T?Zh&PWFT%7rtM!G8wa zAN(1WF8}?Y?Lg&$k&J${3uh#QBR&|!5yVIaU*^IY$>1vi7l1w^8N3K^lx~bfI7)xL z3uh#Q|H0)CBN>18yKqJ__**WVkqmBg;f!SP_W?)E9wQn26TneDWhBCpoi|-um8GSBPVpN_P$>8XT zQiL;-!Dj%D{9z=6&juW&8zT`85T(XZ=qpx%6Gm_EY3pjL1 z7$X_|cEC~nVkD!F2!)V7BN_c;E`3Ha`e$AGjAZmbbm=pa(fZ!jnNC`OAD5{oinBiu{%i|2!QQL-|4N3;H15tsrPGVeI-= z6r3|#ArT>GX+ddC9xx~&9;ZAqHttv5WU=wT>J}H9@T+d|v9o{U`C>TD{!=_(v+1Y5 z#xZxrk$J^=tFmMXX|6X;IA0KhAfQL91}<3F<7_mN_6x-M99W(Po)v^_hO1_~kPQ`g zu@uI0V4*E6tylw;wB<-QMeYi8Jf5@+o{xi<+JH*Z($T|Y<{`7ZRilUFqavd@+me}w z(k3|4Y;9R^>8M?xdG~l(Uitb$c$g3Z1$^{4(G{Ilu7mNOk~EQAR-*HnIACCN9!yZ8 z&B{u2K0XIZ&Vva`PV?dznznmAZwRKv#sA8)jmT`|gy?ahR{WO`x}o%ESmTo5Bmq4) zn74`+=n?bC5Y5mV(35aSmkO+Fvl;(SVCw zxhfX&VIm{le+Ty+{i-lj7V^-my7Ew)yePe7VXg9k6ktPivRmimt%yV4u^1@G5*KEp*gG*(l z$+4RFdt9HGF!uWL(h4{QPmyOE*Tm+ltEx)jy*qHWrU-g1hKx|N9v(I=s>-|f^TFdR z8S4jQc`;aCEK7`~r8BVNay0^-C*$-WhAw^DC^ht+6P`4n3B>PW6QT~WrsMB?>VQd5 zU^IjzWRmpm#jXdYlk1Z6$~TnNlxoU0L6cLk0ZtI*Vr^mqGW=d*T&!Jh7oyW&2I5?_ApfuAU#tY<{we=r|Nme9MFT(G$-da@ zWM52AY&N-1`&IVE`SUV>+fidBdsO~hbysd|+tM!_KK$Vo-1K63yP9JaXNc;#xGE`K#^>N9cnSNYR@ z4_v9PS<y3^H4N*#?#r4#Bn72Ox%{GCJ|Hh_wW)O2N$^MZUWc83Eq<8{iR(vq!FTb>R{e) zD~0fC0U)fa5b-JqUKHgcb89Z#>y8+&;N3!};0+3<&@+hrD)c7h4eZ(%bdt01W}x;k z$Ka$3QVFnbpTj3Vt`0+Q36e_SEfg0Cq0+vD5=Bn=BgS8JSYAIz-1(4@b2VSQlfwy3 z?=(Oj{>Hn=`*(bY({H?s=p9rpvPAm*Prval{>Hob+k2`0`FpAUXLuKPxp)_g0vny< z7ANmwKccEbA1CkPSXacmc#woSKlE|&E|yK;T^yLeyVx~>cQJed@8ZGar_{tatLf~Bll0cXVf<+3-%tpM6Ywt1;BHqR26L=R_p?PAs zz{R_GzK3wquQ_=a_f6nkJUpIvact1QySQyk(1>@DPPe75$_^3o_CQofp^h3fp@X{F6|-B zyQrDKy9l9=W*OpLq(etb43IL2cM(b>6b2~Gz_f{Y7snPS#Jf1wJeqfLOlhQf7vH*T z0L{BtF@bmSVh_m}K=Uphoxrc^5nW3wak$!dffT00;tc zAucLrn{SI$Xkru!I!1`YPgklMfaefal${?%WRCDi9LA#Hlx`VR%?W4aK?gtwe9{7n z2p(I7M-nj>1;>oe$MPFoieAk@9s?JmV~*xBoI_DSYQTlKE04FHtAlDEHbbd4mD4ef_brT9u4BP6(YWNg(AG}Fn{~er)rv7_#D&n4e5$`}4wbrC80Zv7h z8_&Ho`8Q6*-`;iif5E%%mP56kz^S+j_rFG>>Ne(mcQmJB&OhT+bi-wf3&dPKwe&;r znHvJM*7+N!;{Vy4icTs&n5{$vjK7g8qW8x9H<2p-d+{l9(B>uhkb2cn55Fg>19%17 zfi3Y=Ebt8;2lOWJDGnk45iB|Ys7I8E05k%@qIN8w;^*LK1jlR_pCWK8(tL_v01n)W zG@qj1uktB&fW~!goi0AbQ*)eridKeC@hc~vVk6>N1U|)Y5grE*tH8C0Pcehx zQ|to(@hRSL@+op%e2UrNKjKpy0B6Lf_%|n?BG<*IxDz>pcf&3|#Z!P6{zU4rKHyU% zfjIGX(eWMuJBHsy+6YfuOh4Jjk^%gt5?;w@h!Gz`@O{9i*vHDJ`4kP&7g>l;@u=3b zdn})#C^p$U*240)>is=e#L#?-R{aeND~9rk$?PIe(tL`>=;WuO$0z}x;&sHQIHa!) z1O`Pjw$M!TDf)Y!HK4Pj#K|6iPnu8hIO0=0H?w5z;M@}!@F^bE4|zWA;!|YPe2OM!%EY_RRk(oo>jQ+$c^&B-l83~}Y<3U&x>?iGHmnB2C;o4hJ3d zicKI*yi1_yWMLdbqd58|PNYe2D;#&ffuX~?qLrehcFt0hF4JBg;%Epl!M|IynH6p4 zgu=gBKfP`27ZCc}=KR-B*VH(;HCk>B#jBurHAIBY{a7=4E9btib{u%}wf>VWGq0Cb z_}5Qe+cUL-HLWUbTJ2?jnr$)R&A6mv*{t7pqgQU8WJnWD)1E}6iE7X4)cu~R(LcV) zyv}oE>+^^*@zkXYj&lQN{(Q+X3@`J%ni2FkE2igiD|(^l;kcBy;*MnNk0wVW9}k=@ zZvh_0uhw0Eul&hactdD~e@pDOp4hKhv%X84bxS|%?>)1AVkM+05;6=4i){(Z*t3@_ zW!=>0c z%B8r4;Zhu)#BeEE#&Rjvd9e_e;z%$DaVa`x^AVS#b)5*9=_vhQ2o>ajljF?}@;}3+ zSWh!34yTdWA^PYh91YP&S|CI8r+E`w6e290K5B!b8v3XMj>d8+awc*q9*4_+!KD~U zb14oWU%a){52IX)TMhH!Deo~{ickEKOEEw4*SHirfJ?E1W>7Sz6WDqB=r|m4XfDM* zIJy^?qVCtY6bHfK9$bpDKm`ASOHl`~61Ws!Yd~Cz-lM-Z2y z1$W`H7rhzhWDT1=u z11YbPBIRi0b0@9bP5V3vDc`1*^Y~?SC34jxG{H2=rMTswIm*3C>?Y&maLb6P3E1Nb zvrE6mrD)1>_fM1i%XR()p`wE9B{5Pg173kuVGX=0DjN(u<*NloMYJ20hu&{Bxj6ye z7vqD4xAcRzLQH4JuC3PluNPxv$m%KSnK=W@n}!FTF~GKxzUw&d&dvC zQtdli4v2aRfF$v-BlP2qc#=Z_{E6S^1XIMc<) zD38{O{D~jCJvR`H_!ASZUUa-!!0;#L-nxv=T>Oc-w}#Q#87G@!F3q24ruh>+ocxLA zar}vDlXr>8_$WN7`P={p_!D3JALUOh8Rbue$0c3-iLZ?EC*F4QC%!a>KhZIPKe2rb zf8wk6;!n((z@IqqFZdH{2lZKHG=E|+@F$MkLTz=EZpNKL1iBo^UpC}Kz%bzIXc8%pvth9{b zPyC3x$+&4Ee_~<2AMhs*FF*`v41Z!e!=Jbr_!9}gRGL50|5x}EPnmOs&h;ZGc1%DJr1x;}&f zf1>gp{E4((s~xL7))tQPC-%V*`(N`XLXLC*e_{gSPuz1a{={a(q+jqS`qKP~{$u$Q zkIP;BiDuwW%%u4fjeExMCysFEI{6dZ5r1M1&7WBP|FL)GaZOx%!=IUiWJo4q*n@yf z*o=r65Cs=>0*JUrq!yQI!lGyqBT}l^)@(pj)VLO_ZG%fodui@(~JpKqj(UjKtul&TD96#~YBm6{?moeoL ze&RZgpD5dU6ZnbkS&^Vj7J;0$e&xh2g%zCBs6$HPxFjn;YjiMOza_=$SpCx$=5 zPt4``iP(fb;3sa}`r9M?#E1RY0Y9<1?={u^0IR(G#EggeiS58o%=tI`#PJXKi44_p z$iq)G0za`8_=!Mv+5cJ8=BuZbiU<5e`XPQ|?7#98KlkzzQ-PnTim$3b{KTgJ6@DVB zH2)PpQ58Qr5BP~MAbw&x)RzC8pBUreC$@X|i76aEaqD!BpE&8AFLJ&7#HOn70e<2P z8u*EP;3v+YfuG3t@)Ko4w-c--)3^iwlPc$Elvz&z#I8Gz0^VSEbR(K)Oz??!@LB~=L1M#q75E)_va2~P`CpEkElKN6!7#OO(j76oK<^V*cY`>so zko8xRe@Ke(C(C$(GziBjxx(7rl43WiH97&nD#K+eVTultRaIaDXpRKvyay%xG$5J) ziq4eo-uuq%@E4*GvGu9|FjQ`h6c}Rl>r4St0GdawHHD!g{^eKhoIJzj$$*s0UW>cF zmz~j>5LhlNVqccBOA`d}21a<8rhg~f5)BDdzASFL=*BHjvvd7H;-7FVSa*3fLCC@Yr{$pC1}9zrG1Rl99cK|=h*y*z_dVMb;Xw$YGT%E}TuA@0~-hLk7nIa*BOi?!x`EU9m zhu^c@kr_B*CRG#+u7)Da-yx;Q;QJa%kE{HwzF7aXYcWl(M8ifCc)*&lWF{ciM3~?}zo3b2j`bA(540~k$S(j8nW8%7ajN_4hL`o4 zDNKSp5^6-qbS|<`(KiPXHHR0~ZjOkEuUGVz0Gz94&h45QE-G*C0N`zteOuP~iN)Ef z5v>;7$ibGPnzyL{)89tNvQPB#@@GQ!Gkhq38qq=8I*jA9P1dkKLcoBkNe1RzO7fSs z392K$#tgoB8bU6){blq7jyb2H%B>_Csp3aX9>$?;{ay?2N!Njg$)Fz4)3}Chb zMAzFl`U;ZT+v3Ujo+9j&etEwLvmxFhjQvkVm=QnL=nt7(ii+@CE|F80HOL-_>KMVX z{>lY!x&-fORN->xj7tF)v*vQUGLwdMuIr zAo@&%$^YW{QJ)~%s2pgc<=y+iUfG16iZt_nMbe2&Y63uuHAstNN-0IT|+!j>0Y$wjj{1FEfcBZwo-RE08*}gCDVBhc)jg|0MfL9AL5Dhaya#V7lRU zweW=d0G^b8bt;qx3QCV&`~f2Q77yyl4DTnt+|7a)FGk!beUXS+?#oMR=l?R`beDXI zs2*-;(mW5in*t>l*<2?%I*$x`yE<%ig-&MQZ+l)FBP@KmZKH8ftv`>CtL*)L6E%M+ zBw;@Yd9MEL4&Ty~VAoV6f?zm90#NbzDSSUuNrkUVXkKiz$^_L`| z(YJ!9L>p(&U2>+2fc64=%c1X|21eQJ^W;-JFlQ&^)>pxVts@hnZuweHdBacz7s#34 zPyw2~EC8rkd--EVedvhKoanzcEXOAsAmzYAa)!Hh`!8hlV+7CmKDT^}YO#6S{@W?k z@)3wfx=0Ka!Otp`7`>uUlBIc=iWH$iVobTbR13Z63#8g6icb zr!Ia1wSu)^a-qdcoE%|ZWgE?L0843rTe$=Xs4V@Ow(dhUT5(mfpTIpRD7S>2E=4G) zX&Nb18nyI0R=S!h3AA6f7|sLc&a7I0Iyd~~{TIsJJUK8K8I|PC_5pO?G3RDqq2T$9^!pBmEdxwOyD)X`{6k-- z)GYHiY3AdLUIQj$H!vAJPu2gH=S6QW>jxFdfl%cK#*JGQe8gB9#8FP=b%=*am z`?s`?ZOT&CVxxZ>1z&$1)y|C8&RO8Kf|a9}6z060BWv|Hl#RKU{IN^=D-8_j)^Yp% z5A8+hx2dCAn-IPY;XT)wGQv5umt%op!}<47elw4ctQYO`@Yzgh^gA9n*ko++>D&rb zHq~#q4})o8_-Ts<@ckww^UfX7mbUzLkKNDMg~N|ueH)6b2YY^ty>#ozzlw~170BIh zIX_@m1PXt-a6}E{O2U9a(Gv2VjNkn4KLyr2!mg9N1=mIt?{-JRB|cn|7D_JOtIwUu zoI6Yb!ZnfhG~rm0^Y~++5$t?026eX{B04?4XH1nmzpSi!#kacs_|g5sv{k(gF1nInbB@qZICBAJmIN* zj`DwjssXP8%=2Huk*x~2D+e*S@=q{@%~g!woQwh;TvWxEqefz?rz^2J;@!e8lS)x{ zhVoT!*bBy=5rBZZ7 zc+ka^i1RQy4a=cufcCXa>~r+gDhF&}-far(TRIS!YuDN9G#;4dRsGkuX8awH*oLp9 z$W(0O0ib!it4^Oeizz?(aB8}iJKuPmOU*$UC^hFwAvM{;C?bVBoiC+IZ+TLaU3xu4 z{Vi~tkJO^w2lP=ud`$Jb)p(0ys&(jMw$zu@3x&z1vS#^+=1a&*!6mdc9A`VV%|Xio#MYA5l3_;Y(2u*9=5 zzw)4!u-ybgwkkF=x_MaO7mRHoYtVKCSOC)|H}lU55t%d=z>KSv>k>m3!Zx71uc1)u zYnY(IEguK$3Lh9K*gv=s^TE8@ZS1Dc_6JdjOs&2HZAU5EDFK5`@dWRuahe(aP8>ZD z$_ZG11A+4U0#`i(-K6GC*>SF&FrM&P8VX$n*v6>2Z_54gbLqsMj+yT~?X8ZNnR}Hp zzY!IFBan9bTVw3>Xq!Jp9WeeKoNUO+88S;v?s0iTLI)gSJi-P<^V9k}vS|`ml-xtJPTg#Jagr2QJyvSSk zvtYT*v)x5H(KRe*VhNlSH`sh*nGIR*xpjAkBg@A&(Orrf5Cz|Z?S5pEOV97iF$%svNV!suvMWH zSJ_(;?RXIpPs|7W;S%5Fg|Gn@xV1+OB>UY(lmK@YF@UFzS_C@@7#pC@-JPx9usnWA z*pke6T71usL&2KS@OX33b~lw3l+Nd@eUot>@q*%OJ`b`o5dmH|ynq0&dxMG#_xpXGGf$Ac=X?Hdw?ED5V1 zU{Us*5-IfHE(vSfKO|4vtO~xMriPM-8$^jq&|0umJT(~-?omxkLOhyQp+4TGy!pwXsXi2v&N-wFC0+m3A+Ok* zm=NB{e`6!PitMfJByJJ4A)4bOi8yMIN40)XEqznoMWmV|1m$X?3#ry~An}!>O^(R_ zRJDVHk!o#7wR0_r_gnf))RkoAX^dA0>n;OqL~iYsdIBv!Dbb3OTYy%SuA}8AnCyoY zCASH!eA&1ar65=`12!VJPz!8C3&%zT8DJxV%vkxih>d6jOksV^G4Svx@=z1%1WMiA z%&P#NTgLwxu@PV6*od&6g^xySa2IpXde#?cVa^ebLVqR(R_lO`$gQAdW_T9Y=qPll zF9t1jS&(A-Z4AdoMC)RcjkggSkpX{Fd?8RKVLQi0yn|?W{9}lXh!)Xq2Y_(&VvmQK z6!389%y>@vR35%6c4VKlc4T~e?YB|;e|Q;uXhXMoNd-}hZ)f(_DI zze4)-?}Z?O8Us%XaESZgd|`1=ISvP%%Z4-n816?D_x=|(p!3xj1K)uP3+zU{^U9GJULi46cRHl$>-f8hunueYn=z@Abp-j zE1UEz6vPihQ4kYQ5WPcu0e3hw7?wV}4rKKdswtOezV!@_P>Gg4mm9t*;~A)vhKOH@Q}E}#K-xDEMi zt3c(@odn6!+J%zEKF$|VgKLKz0%YV(RQN56z-Pe)1NJFpS+M|jHE8=)vQZXs1q zzI+VWh;65sBQ~P<=yfL`t&#uUqu1RJ+z~s^d-S?1z#ZxDJ$l_|fIG@V?~!wd zbSMG$TyXauz4C7i_~Si#-9eONu6d8hT|eN4SZv;ZuY3A{oA>B-_c9K>N3XjM?3eGVP2jsnsL+{ZmKMZUHDkt8f z*WJrF^d7zLUdEyK=ymrp4!uXOyO(k3JtFr%S%)5bh0x-lNyu%Q*BNz5e_L@_(`p zA21I8Z(|*L+irAxn05GHx#Lg>9%dbS<^Obt$^oB0)8Ri)hd?Ml53>%%p1vseO7DNj zCH#|hn7=sLLvKq~R#q(2mM;aURK+6@wPfYU#f88oODQZ{S@h)M)jDm(>XLtvNI?># zS7=v1SyKGSr9WARi&Oq&9sbEW{F8M!X~v8s1c@tK%^?mSWgRa06Lko02>&;u4jx_V>yI{rDXD_NUi>{ZEDqQLw&dD>y3O>j$i`=jC|tJnmD?wef`iO5L zRQtw;7;>Z1pA9mRED#H+)qI<>p4FJhJ#e5+Vj4Zp? zsN%=Sa%7epOaXirmg))$);t4C*6MbJgK3fRO=OXBr?0_mHXk&b+FGtCnFpWK`^;rqWA@v+ya7_wm5T%|HmsVqo>{{ph9 zM5dxQI&NetOBnd<(ekxatv~+&nL4U?j(|@PEkmNUc(WZi0o zFwM$Kv+;tmXBQ;Xd~#84Tj^nP;)wjYmL||+>@X?e>+3tkSGb|(L!m{F9g;-Lvi8su z;Bz^BLgM%de2gSXme04mL95}@GkV9M4qgU4Z(h&KXjYBe$;&YF#`PKo15)l7Z_GZ^ z_>)-~j@A*!veE%KcaS&XZp(OA>$vu;^nF=lnpF|asz@8Q(Y?{Yu>CjuCs&+S-G)!W zwcGao+n7JioTy-SS_>eVCPE!(zj^S?T*rJekB7f<_rD*^e935h(t zM#SK}mL?_6bN9j*rC9Zf5~6OctCk6?b*3tXu$mNBxp)~3b(N*Uige-HHFeecy2T3N zI)(6QNJ|SZ&CUxt>?#x`<`ofwi>{b3wd6b%P4g!C2mMY(Uxd|dR@llAkx=rUQc}?; zoltiWCo&4=rowN;-bQl>xYe=>lG|EW0l{kOo+;&_aCE{7g|MQOSCK1RmkYn^MmN;0 zu?nk8d8_qxqf6^_Hen^!A&=Tx6m*9I^dQ&gekAdpYku?V6k2%;n-fT*@ zXN@_AXPC3#$+YnLtaLkEi;~ZJx zt^ugeczBgh+NbpU8l!*cm#%P2yWP^CU4rYF-wKx$O!5DpoXGt$7I+Hp7z(7!=d6;n zmN75zb4Sx#h}dw;Qhd9FAA}d$4VGAz8ZEnkYP264j$IHK1N0@lyx)-nb_}-*^v9fb zfd=o@m{vtPi2QJ>4_y*!B-k?EcCQ$NUlL^fhA!ZxJobvd4E>ylt$-Da$TFrGqfL_R z4>1U?e-H#mFUBoS0YOpdx3(Y2^6Y|eGI_wTjHZhiY^mMJ@bCit$WC^r3f~8o2lmGT z8w7t?piz^_OsCCSA}i>OA@*!b>ba3x##Y+G%(!T7eEUsbe6*^)*4H#z9cW(Bz-TV= ztm`S>sGC)TkEJxZ{Vfr;m;&_Z#F8#)f!=?O<2W(6%T#n$apNH*Wn?P5txoQd2F72s!FOuL1Lore-|6b_9-<#E)EM*p zOdbFn5CHRefji-8BaTJ~%*WWlJ8TzQ885%jc1On3HrAbe^8wN^_n)zu$Lq9^;1Cngyl;sA*@ry8F>!POk)W;0ox?A9Lf=D4F(*^mbCT&Oh#zX8RSZB?w#;hE&PH28RX)i zmzjC*DS3PuJVl{LE+DfivY-^>>p_P5IOuPWnNmu54EHDj=USWF7g?qQ84zILAN={Y z&6Fe6dx|k~kwu?hWB(k$m!OFrIbc}_)BGaF!#Q{{pzy_SR0o7Cx3o8)aDIm74xEDs zi|zsNursZ0kUP!gl5ZMO?b z#Z|1nao{$+#3CaR5m;7i?JajHVhn71~)Ka{xC`68$bCBKZ|IdelS{! zcEk1v#RWrNq@v4kAyV;`fmo5MIG@9Np~>KnwSeEIFl@dIS4uG@WQ>$nN@YqZ!w4dN zuw<@k7K;uEf@aSC78jqNeIakoponQYP!&4@ox?+L97*aU+PWi_1mPK$taNj0eWo5p z!NS!It0?I7dj>5dsm#GG&PdoLVTsa*6ElV1ALMtesfVHFI#LLe4LTTe>UdRa>U2z9 zScku~&noR!`jspFuJud%)zXo4TjWlrO;1(&T^W7@OEG;rU|t?4`e0d}->y#zvlP>{fI8$iZ%0tDY*USo;)g z=9J3)8rairwP$e|eLuVoz1ZXYR8gt6!sCt{9;{@e@K;w>JXKV=nA2%l`I42$@5PH( zbBj@9$1c`Zu3WLK2(4|e=2)$;zFoNtmZ{LK4=03fh*dp(>yOIB7qt*t0tT2c->W>vy6EVqP@k|>Yc zu?n^gT(+VZe#0J+m8+3Y1G_|(t<)B;&RbllEq`(#1F{}I$$t36qX62?4ia|dit=T; zm6avsl?=+0IcO`h5+-ZdFeX1!rCMFOvZ8YF{QNmHM{&uYJnF$VS0z<&(1{_Wzhs_MFgz$9`{&aA^0(u2Pn4^tsM;p*j;c{xSQK^Q@ zGwK=YJg5hk4I_PeW>f0y z&#A02WnENWdg+?(@z-tQC(Pon#ys;411mRzk<_>L4NfXxY-x4(`0u!DgVUdJGn?^A zaP1%VT~10mqV$>W@qIS2`k?qi_8updHLh%_(sWb&!s7I;HCbd%q1A9{?TTr{khGQ#BNJEtSVjee*gEz9>K-%4rY%@Fz8&2-wABe3HSFc z?}-B<)T@-HMrj4}2|h_UQr4#yZjV$ytu%E@*D#;syM!atSF1vw(uO|0Ep*-a(Dm{` z2ULR&X$Os@dYI5xle3cj%51``tl!9>UmSvUHt|h_M=f&*n~~E$9fD_J&n;z^!EbL= z123aP-P3wY_#L`*$8I3Txq!Pmp{&u3M4Wy)M)W(nbk70s3_ozf1HwZC;3y9WA6n0c zyRB_=MR6(+i9hn1a|+z;Gqvm$D-vl8_+sc`Bw}xi7Ys)tNm^l3Ke~hmR#`>K$SFlD zd|ep`PO88!5-yr&bf~+`z#pC2*+xy8*{K@Vi=6t*kY~m+uSzCuOM!Zs?IAiVBDLa!20*MHMk{=KXsiFx z_dWo(KFVnIo<78V4R1Ew|Caa2F2n-N^fYy8jU0TAbLvKe6K~yplHj;v@K0<2#Sy zTlK^L1AJ@grhkub?YcVL{V2Zmv;Pp^D&YXFrcv`-M=j!wF3cLe#58(&>*#V`T18e` zl_{;dHSIax_!byo{=d&~?S(iz%s!1(OM_MTyFBV$p7>p!C||QqU${<30E3iY)TwW; z(_pz#sN-5$1s;^wS9jD4z;{K6Z%D%0@GD zqe-mLSFbE%o3jxR)tsPLh1CIOW+OPYXDh)eXPDku*_r|%o9r2Up0Z1Kl8{YL3z`7} zZcZRK5eeWvz*CqBfdx9j-&ZtP?+~Pue~1&a7GP#A&zNY$Y;R*Eh|S~UY$@mz&RmKy1m4p|J!BlG&HTV8ODKGWH<+I5-H?w%KHvG^N;Euo;4n zY_Sn4E_h-z1m8*1so{iic@|qQ3(6vm_@%%KWucmqlY*VI+MO0vvYC;oC(4DNNzRO; zf>V`W>e_g-ai7m5?c;-jr>mWDlk!u(q{$VKLiLjw>@x{J06?bMYxx9-2CW3=eVnu4 zssNm;Ip;(qeHff$1U0l#Lm#gh1R3?y0~qQ`e9qj$pVT;y5AMuoTT7~>6KqpIMoxOyAB z#y43#c!4N3-H(8my{72kcbw<(^LqyRJ;d{S7W}Rq3f`K6!CSc)ymk11x4pzK zohfS!)n#F)khdBXXfy?EJr}&PDR8bW8+wNE_)tvm`*8REflD^-(&d3mW!xo+=PoyQ z>8pWDI&>*oI^dxYJcKt{HSmU6B)<>Xs3tw=Uq#y?rhWvxk=Cbjz}smP$WD0eRfw-E z8+vcRpM%Js#vmlqkO?xibC3)x2AQxjIr1IDIE1c!dz20R9(}=?3%R z1w^xZJb4q`(MQW*m<5i3{JcaoJBExS=$NhMnP>CkFzi?$LN@7qfGnUe>Ysxemd3YZ z05CGq_433sKC&Yt%|mTWYrUMm)Q%%mfDE`cDz!;}k&``+WPSYW@=4pCAJhj7{8V6HLT%Mhw4QgY0RrW6pkr(mE4rih&wNNu9nH2Qd0&-N8aWMZ2c*^$8yKD2i zt?^spyWgKT4Wg%n$#hX6oh_G{Cie5JE?!5}!3l0tMnjxM$q2Z2>+a`u7MIsy+ie-| z2F>*RZbHB3dw!pQ8W*&b-X}Cbc^-voxeayQK%i4`kz zFe77GU&LQ`_{?u^siRYNdKX# zo>oOi_Ce3ZBXH5$GCg}5{VNL>hpBpaKo{N#7n>g3In{G94lX9#zxdO@op4bp3W_4G z62>MrxM!w;>CHfOL<>(51P8HisFVrpa=^CvDrMzOHG9EGnIp>;RI)GL!bs%abSUkcB|-EISP*{GZD zrYTkt|2P(N%c9|D$D``sECD;J+z7C+xE~s`<@A1k<7{F`Fj;};^!dwum-yBjL@Bgr zWAhUQo&LNcLZxJRfArSr^Cz)XO64y~KNGb#!_`{ZI;8Y7$&2u$0Iu0lMAs(=mxyUMh`m@kmC$!W9->jz&A%ig_XC(H=&2*K|qOs))O+b zZ9IS>JV8dx_N{{;I|3qrG^TN%1~Yn02xc@fm$gSj^JM%dW(qht0Qf$j)qGCfC zv5Vj>31bd5>zPER*iF=T*6}9TbClEQU|wj-hwcD;Rz53czlmm=W+5|+MML7dlukfp zhKJH^!F+8d!{d^0W{fOLeeNpw5Ys|;1f&;{o98cN>2GMWFF-`~{_C>!n>Ca+V5?YJ zxRuOOPIFWFJMq)xCl-SNdR=bSOUKAf^hweJK%E?*=NPJq*F$k&61qs4x}79X1fzDL zmy~s}bV;*d5eqQA3yqKiZH6INV<1(G$7K2!)BDZOpf090gr_cJ!r_h%go%os2COOl z_P4V`It2nyCVZv_3du*%lQbU7)hAT&0CzJUaiB;Must5HUhaIM#2$av_!U4$07`1K zYqlFWy`RwJ1EfjQmK4#mKO4|YPd0aFoZ$gR>dJw%4o$N&YJ_y1D(2aW>Ld0da%A%0 zR>=r?f<|XMVsDj)G2l8vfFBs$7-PuLje_mA% znrKW&8zf$L;D~))w#9$>Rx%~}dw*@t;f!b(Icl1H*PxvwRKP}rF$@R>X4jHYrwvO~ z8#NAS;amA5fyBhk0ODhl%l-188zK@3j04-^#I?0EBt?3ZLE8>7xJl66YABx~Yl*k`L zB7diU=Jm%V^9yZ8>!I%+Gwo*1BeUW>2?Gw^p86^5t31RWT!(}B7IE(iQt}Y!q+2I(+@N&(9P5NX;~z znao0&Eb+;-WhoQr!ZruZ(lD)3Zs9D(?QM?76;F|8d(GOQic5$vJf`;(ZK`S$+7rh*Ifi!XUiS>n{}G$Y5cO7e7AtJO{&GlZw1W; z-hocr;QN|76SKmsxGX=i-b~-4>jQe&1iA-RTcdwRz)nDqhLOeYdZ*VW-@d(ZqBY#t zq7saA1TWZA-fXf-%N-VfQzp>x^t{{tZ&Hc3`mq=fq{$s}X7UpkC&gKw3Nve-s9l`= zqLx<}OAyT10Gdm>Me$5V0Bxd5l}X(x?X;jLI_WdXat4hvWu#Vj=JZPjKtoVT$nRZH6bBZ?EMi}0|cn*e3fgF=C2!>0;U4GvK z$Z()d@x3%{1~3V|@xkwPeLPb=6yj4mcg^!vd&+e+`rrrgF$D%}@k_k52)`enn?s5- z;}gkkVP>{iuoj5y-P8$bvq z3wO&UF4oQr=j>|?F%vRzxu$JFJF0n;`7e=6d@pKj`zbn=zAylx>(|tt`|z0@s@|9H z#3XYqC0C8}r?60H`cR>H5;}dyx1;Ki5KFwqV$pHPG+vw~qte>5KG&;!jOS8rs)hm> zujPq;{1Bjz+ZOiYcU}W@X@AHw?&p>(z3p$-7*lBFWd1eMC@ZH4r9JvN4ItxWEjFX} z|mK4v$Q$<%#7Qqbj_>nkLhM9d5yktA;ddkR@=J z&^FO21CXrb0MAQnG67>)`W@D|E`aWtidxh)a%HoDp~RjPF`{C<3OuxEG2tp!7_v`_ zjP(0D-JYAiz`kXr+hfQUW40iX9574$0?sUH%ZS^rU8}uJUUkUbtm*B{tj1%Nz0zah?MYN-}NKBXeFln!Id{Hc)27x8d|;Y(Zw zTv^!>ab2^#QPbS0kAaSq-jq(5zBB#o6qEa}Gtc%DMqpa;#wi5}W|3cKaudXMd^4|= zxCP*0vW3V33q^)gVPHzDy7g%(L4Eda`E83)K^^6Lts9(Ogz{9 zyet|T!deWUip#+m6(Drf4b0EmLohvC9Toc80$L6wC-+T$o30h{vw7i%RbcNXT+`*4 zALqK0Os)C5smM@<9Na?XvjA%ixcR9?g~)z8dE+n5@Wttw{`JXxxp}~Hw8(sl$TN{o zpjIA8tb&QRP04#o6h3Nl2SAH&^Da?$G;Pdl8k#y|c$ruS-J9u&E;$J*_SS-m_Z)rz z-kMGN9B=$AfNqsk1%g^8#m-Df6l%Y_yir+b%~z6wYjmD-!p?Kr_~5&=Yq_rlCU_v< z1hW-2&|L_c>YScKZ5Iignuy%xEAQe@?M&Shr|_9N%&4-*F-eSNka18s=^hG=beYEyCc(Uk~Kc!nLYw&s%oj_zE!&o!TIv^Ol*noT&*p$X|7=s?=ic6tI|yv z?T&4z3lS=f=g)LX!;O;S%djHQa2WJg4im%pA@2lHa5};V zYF(36R^7AW=mNgbqIixFst z4$BkiP5x7Sj`Ew8>e<5hV6kuvy#+@&$Ttx(G2${LV-s&i$UQdn3~k&wD2vuDYjZ=J zoOm}TxO^I!LrJ-kUETc-@zkK?bUw`&|B+kM6=3u?0h7?)zjcx24cm(##<4x~wC|bq z3n|pNt`8Qp`nM(o-VF@)Q@Lx^jAKvZTd(fz(7pvxohR8eN? z_~4F!J`_W1OJe>BNZs}_rFd6TLH{gIrEC<}hFdNZroi~;i9wWyJJ}s|+_ng^BdLc- zF4O@BG34G);@(|g)4+q-!ptPPS$%d208Cv7V5%3tgVJ(&cvs8Y;*i5Z$kL4~vk+gL z;7lj84@~Ft0L;@P=H3%bo&G|Hra&pUx^k^@4yQv~_w(N+*c_UASp-2EYt|oZu>>1@ zRks%?XI-7haTr7Htr86hp(fKc?x>^G2bdH;^a6^!3q{`H3^gtw!9c#a@vno5N%u!D zzfKN23Eh8%r~98Ze%D;(w0?y6zAZS_f7j@-@LCklLAn`VoD40qhx&PUT<3w+P(*f@{iJ`u1g+(KR=o@19tw zV^9L07Z-=z-Q|nl-YSbAQ7YUt?(x}Ql8N0Lw@Uxojg34Jb;g26BUt<7p%>>tqlFaa zDW|c3VSdW}iYmSM8sx5R6)pWdy}&kqz~qWNb<)$}s#)Jdyyz1nus~;|JF1#ljz7Pd zqg2ywZ9K#vyH)Yf!Mmil+i0ycUG z>Za-;8;T7fV3rYB(_xw`9;w`0qO5yLJtSoez6w|6f5p$dmZK!U^z=pP`M4!`_5fS8 z-T6kO6ZKG8{@bIv^u4=bglB6V!{5~umTy}y9}o2C#RyMbhZ5L}I*Z@9_q3ur9Q+fL zeixvA^#mR-dJObw>CQcBRUFqb%S~)3cOCt>=^wCLg11k5DFXFr)rb-dTxR2qMY)z9 zGRODje5Rf2y2LibsR6E?LSIm2)_HqDRYfBL_J0rveY~i(8~(=d<}SK6>PuSCGU`j4 zZ)foR&YXK?i#%jqI2!jtF~%dz`_Q--x_R$0D}UX!cuan8Xfk;xxcTunB;j;Y5bc1+ zfMAl~BJsf|FNh%@xG901UC(~tN_393e(Ou#7=wC2HwU2s;!)~*ht%k!jrEL87b+U! zbi$Z(xdQ_iHHa=$DzJVSWxG#bpt%A?CeR zd*R%$F@^rhU=T)+GphbFjUF+a(^1t_)`lP6=O|J1}BD1B6evaQ;MG^}`V6u!#QXzYd|&f|7-TR>{^5 zgDR}4Hk8hu@;ZFCCALoJG9|-S4!^pu8!|dldUz9LA;$&@P**{g0xeIchHGW~>Z10( z2d#|X`PE1HchtUbfR#`qE0NI$&B zZB-CkB1>jra}e<#)|$_T)oMERG9$mK2T_o75l7Tj_Y5=+p5AfgrkEwwl^?tPes>8b(`KI)3@_v&J}@;=qc{2| zf5><4- z;+Y(|WO%B)aT%1F!nJiigBW$l>%1=yGD9}G^N>A*UT z+P|Y4X2`dp@8pkl0Wn%=W#&w|{VWr3l!{XY)cTC7!$Tq7wJ6G?x$Np|?e4mj)TwRF z+WGfH!<77%)u6f>8YyMK5P!J+Sd_J`I3PqHi5jWmY;9*v_VLXxIi3rGpE^@!bCi!7 zfwj3dy13g1Rz*CU3rsyH1tWtx)q)+04KSu@fMK?arFeU%4Vz>71-7ZUKFc?hhfE!n zMxl;}o<9w7ynh$U{e{D{c$cT^Ev(-00W>3Vc`Xt&G10{e5@BJ0o0zZ#a@~Bx*T$~D z@c5>8Y|>%}F+uN{rP2DUmC@H_TO8m2u!q|$h1TrW#Fe+eq(!evhku>ND!0tPf_72i zE{H~EKqkyYy=8S~?ET)dn)+_$drN#_j$*%5y(wHmlZ3yKVwIBV% zfEF_B9Wvtfp~1U~SeT7CXxwHJiC)8gcMmqE*D{H;#s(V`I4moew#95p7-@mNmHE~V zQL!{(@9{U}-`I#52{KJH{>NP<1_|8&c;06Go7-^GEjZ+7LW@rGY8x6hxC4w0an=n6 zkuj>Kp@6r?&gyaT9VlGT8x1G()INg+YtS@W_sum*QlXTsrrr#_t+5BzUHr|dhkiy! zGUfd3iTF~?xDGy-NAML){;EyDzF$ZT4>#9Vvt3~jkHPE|UY+U~=5 zI)ZzWN&Y-alU&g1P&`IyY{WbYI?yz#zt#pqSWN>wnW~{PO~;s&WrEV=pU5o-*n1}85B&* zAsRoJ58im)M@)oCN4+itdUX|8`42J3q{@~N#;I$R&{+i-bTiVYX~SLwY5pRpESjo- z8FhO+SCYcX$2F6}1jtO4H8=XYlbhWYLvMeicKQm&)zXIxbr#naS*(~e8isWi)8s9g ztEhjYb2Pm@62R)@U}`3P#ots6ZM?2#A4SNcOzWuPWJRw8aeb>%aRUakag$GdXQU}g zfKL>7(={WEe@Aw+$iCGuTyT{C-0Uh@mloy%@m#dggdVrE{ zz(MhmE!MQ1%1P-zjmG~UVU+i#z7$2P8$8?a64Y(-TC53gh{IsUgKHK|%~Wm*^O5`B zO6}wUYkU#p8QVoA9c9vZe%`Uc&lB1kz+(Mw z2oGBML>^3%jE54$Y_gL{l&K~%Ss4!nFSYa&TQE#}bSHLXQsq&~=QXJ`jBK3?1R8bb z7^rZnj(~%wyLmZfJbA3KTQJJ@F)Eq2iQ&ZC=ULQzSB2^EuS*%!#R%vEz%)!FX_5p9 ztky@5BPCZ9O*unt1nG(8ZBt$G4nmt=jB{=RLBqO*l?XjyNGx zz>XO9kmde;ZK@#N5$Ywdnw_ULFv7iBm>xZw{|B(^#m!{ZRdypjPjDVal=>nic9beV zAEMRdPeFA@)sfs)Xxmq!Y^U9Ynk@!N5!V5k5K*37@GF2mKTUa~%x+UDHPxv^HbW z6wZ6PO^)lyo$c#~NRP{UN;K%f3Wjndsuog}vhRjA`{1EHLz_;|jOa-2{1}sfwexoG z583xS!zk#;!ze+$!nFej3*IWpOo!fk*8F{wl!OzFzD?L}93ljpR`sBH^wrK#D7=mD z6Z7c`=oabU(uY_U!VW+_PxKjl*D}2LeI*|2R6P*mACe0`VU*b!08fW0p{wpKh9?EE z4E+RHxc*_TNke;2vo)H(5>J?oiyG_>H2owoKSn0eS5`1&QP@J}*ksL{@-<=HoPBrg z5ZQTH#CS2zzR4hB7uTGSnNC8dvY?=8lYB%Xa~^|f3I1$TSj^>SU%IwJI_qw-GB>bCdnjfNpX2*d9j+Y5*Tr#3Ey7RI0;Pf zoNV9Zw!8qdYy0<7q(0S@{GoM+%KDO%$+GRh=L`6kiosWIwC>_fl!)YDbq+ExYjva0 zM%at+5pC|Igf*YAb>YUr5;Ra8(HE^yuK5Im!rm=_-WN;`B;AUZ^)_65+$Fvl+?V-m z&2d>TQ!v*h9vjkUEU7sTZHV|3Ef`Dg2|kzEDc%#@12lZ~Cm8g@9jUap4d)d|az-}S z%`=`tz3o%(B2}W-uoaHF{jp{U6r5zFfpf$?=x{oxrOzX17)D~vclCV79D3SQyNUoX>lXycTrbTVJbWSfGtDClC15cPqvo2VXJ}%)BuP$+ zCxqGcMo*`12@SJk202zl3O5N#c0Oqj*=e8(Y);QC`e@YkbkMLR-Q{!#H zTX{z%Tey`Wg;6PKujHWtW_=kvbaeuEHk6ZhM$rkSo~4u51+XMqF3=2`o@Oa?MFjhL z*L~{s)oB=vz3>jyc4Ogh!9z9smSeBtN6Mko7kJP(EkVY4kr@Efc4yG)%Ui1Bf?Y7p zAqabWwB&Ls49kf=a?=MCEJ>g(S}^JKG?{dvY10S1fkg_NBtH$igyoxzFRRHp}$${T9Cf`ITXsM||$k!|P=j z3lBgos2u_iOtl8VQ&-)hvqjU+!xLL{?~xL#2_8M&%;X3$UTpM zT|<`)k?0|;o-7Pt<AekhuDB!<)HQ2#dckW^=id5Rh?CHq2H z_IMRMAkz~I`e(g}9*cABKo8M88^Aq8voYWIoQZ`?e}PMfY`eJI^XlQ5IC$hL>}ni( zs_qneOpM%vLFmzeplV|=icZ^f6rH^&2=_=FXz6Y75X9+D&oelxawKzc20VAR@uVjR zQ!fg=;6?Pf*?TeF5ap*qCa-tkF`vt5&j8b16w{`cQL5evgqRxEqnPTgI{|Y~i%@j3 zwm(-E)~+apWC`>69(H0lfP~9pzXx*t78Gdvwl+{~EE3jDMh^$wRX{9w!>0Uxq(pZJ zg?qjPucADt<^d5}=GPat7tCK8@H^$rM6ox;3bQaW^K_Pf-WV5G*YvHqW@SRRAu zh~5)%DK=@55BIc@^DPvpV=l^zuI)L+AG;Tz1aBLr1`XHxdS{j?_J=cuJZr?xiJhdz<{^mk`CYn5pf0(@KVkIDjL~n zr081Aa7Vd}fJ9|A7eq`fEY|pz)lk+_*)FEJShj&!Qdvu7EnC)tU$?B~+ftd!n$_?1 zKA+El!L;x0e)s$T@Avon^6+?O-t#`^b8eq=&gXre=X^*x;{au+d?J*sZ|x2Kef-I_ z#O^J)KQsRgcM&OHh(&f>j{-$}qV{>Az6M1+>dx8tsr#@O>Z|G3RA1*OBB!T6W-Q^g zrb?mFrW~g9#7$~OWBwy0{_J`xgopi52%R-_iI3w@m0qWW3=VujI{h~x8Opx@q5xskOS{pl|5+n{MfYh1&w@^UEUMfE!hB0sEL)mR$x)?5{;zWtXhNDT{eY@_bA zp53iC%S)CVQZEXlAd(Mt)P(w!RDiE1(N+2nIntx>{t167u|4IK7ujZ)aLXU5po@0h zC4UMrObw?PYy&!a?}fiNz09CGg~ibA$VUuQ0w{)tO_a6PMi+Gm8$^2U^gUvPgX~vp zGk{Y<(HL{tg}!i{qC!I~9Wj~!>86>zyjBMN(WsqfoRc6)Phi~LmwMry2bba2q~n7X zuHIC#5X@D~J*9EWz+9{G0#!*?XR%M?YNs!rjXGqQE+|&&PZ^KVCv~bhbJkf~n#Rv! z3pbq{Cs<$O+KRjJ>`DA|vf?)qWx*ev>zaBx7vpKmtWGw+q=5zRgOEln2L~i)g@rV0 z&c@0Il%`u>c4H_^v3DKDL$+4IK!GlQXQ(R)4;juCa*Az8qTH(dlY8RSc>T5Qn4BJo z*ypTt05|9`*?g6LmsNe>1MjF z@5PpsT5N$`+M~OIC-lql#?}RyI;Y?CTXqjPvK_Jqtq%G32NbePuv9r=+(8O#i#36>`pCG|nEY-`X_Y!TsPh%DbQ9;n#ArR`Pj>QHUj4tnN?~0i z8cpAcWJc4v#qE>c+xyB}2E8d~k5O~pOK1QQabg45_e!iO2@T-mX%;j9yK++9&Im>M zO4CiO0fdPS;3aAR@3&dzcs76!l13UpA2op9Bc-eXw8anmq5i*eqW)u(Ri9`tIO9t5 z?ujt(d|6xG?$StI2E}_lJJ;45qUB_JZ{0BA7@Om1ciNLb(aQRwcY+oA&=p0PTCHY1 z-XYlAxK^gu?+vgBrr0vY_{N5L%J)?#9)ED+@BI{br1g2HyPg{#j|j%crp}~7pH7S0 z8IW@QurlW3PW8f1!_rihkIRNo&_|s&;$5`*+-*L8#Dp;XFDo6l_Ty=^p+ z?Y#4Kqh5YGgyz4^zUYwlN<$T^r-19}_IGPqOasTa4~ENUDi*|lWcKaCdR=o5Owm&- zV>UQ+@;uez7Ibz=3Zd>>m$15bjX%;8+ZEcp{bbKuhMv?Wyt*ujqI*u$g0>&pW-=R6 zXWeE=)l65!qaSI}u&3dr3WHCyHmFLS(;0O@u~FZf>RpqlJY-b5o0QH_Td}_PE9IRA z<(($wy3pbcM(<*zBgrRads<+ z9o(S~yu{pXFU;0Ceuy=Nsk3?=oyx*?m=hF1Fef;|EGak(NFsT1zHxYGf>a_kNF?*{ zo$tfIz;slYCT$HBCWHe4PK-FVpli?%PN zi(X(CjiigNno1Ylk@1)P1s%IcIX4(nm!p}&mJ8bq8ZI-ZA;MRl6%ub9m=HT%!DY|qo zYBZ=7|G+Gr(W_gXoW4}MUZRlXq^90t zdUPYcASN1*nM1G@m~?pFXQaLBp$2(MtvtC{o|G-0n;}os%8##WPSCsJb^CSNbD_fO zTZ?b6bX(Z#ID$Z<@R@f(COdN3@d5uOl4Q6`h*OdBm*zl z1J4@+LDYPq3;0B*9qbDD+#XQqqf$3uUZAG>uO|@SMH3UN57GJb=+IAdjYc*Li>Dcx zo6XMpbkTtbAXGpQc|Tl zrw=@rup9e_fY_V#a-%X&p;I}Blb5s~-*5z&K1^GGS=1PG@rNH?{y{Gj=>5i^!Tq6s ze%UxzB=boH3V-+k>7)70O!$|X-~3B)`BcK3$KM2Se(&GMEz==+iaUk!6WlInC!>b9jdO@Y>kg#dzcgpvD~^RjSZ++&yqXlIt_KNC9@SqSREY+=KTTVqZ08}>|kuZ z@ZNl1WOV8v?HiCy0#>d3+$Z{GTGWw5S&ZLB>WWvBUDFa}vB^!%t)Xl|BV(IVyMLVb zBgOY(!Cv)xuB7I4=eRE-Fts(UQycd?HiCE5dt-UIDHbTIed+7LYjt64aV^E+@AR8W zaO25H>CzBuDAQFtR&J$7j&E&EJ2a^H zUk=oNGwhE$EaTo&*if(hCL=JurCj5*?ryEzZh66x+*hYfGj9Jtn`XJUuVas{;TfOC zju*1~Ui5KxJkr_l#*VF5|IpF*Qp%uLBs zCfwM)q%=P{`0-6q$+DOicj^4Y4fmQu$}{{P)Vap(-1^au?&3k!Q>Th!2Y0@Dv|Qf0 zxUEdr(QweeZqoik4oTF>*sj$ZG?AE1)e$FS85cA?^MCY-*-F8l%8k&b<0BRj2mevu>^53ZVM!Y}w|JL+2;;Qsx1CU2+LypZzQi>y+w-8gv)kuPIW!cB1Y-G`9cT^1{H{kjHoD2s?KUIn0Xu zYY3v~2?5n3am|U|5ni`po26O5%*)Aw81a@~y|!(*?)tCV>X?N-3Ad(2Y1Lq{YT4>J zRch1Hm~MY-XG|_=2xYFOSv#c8>Wo=u3gH}AkKD)!Ttx(WU?a zzbeydao2~av$n`WQt~)7(!1b=(}8=J$a0&y0yY$^elqL}U7IURpDV0R_S^5MnZNN3e=Z#Xw$A;Q16%#qe6woK66>5}y1r}M zH=Pmmxtn+FKI@N9I>{kNV##=Z75jap?L&wO>vodXQ)!*nce^o0V>kNHo%7{5 zSYN)cUMA zNKm5QS3HF*8TI-$OJDpaOR-q?uErJWoR&k^cKlF>Ylj`|0@0-0OhoC9u;HArY_kC4 zcT4ljXLJ||#&;J4yM2%R7{TS3*+~U!uaQY1mm9vj2}0?v|}J*s6o~y!1x!(!uUgZ48}Jy z7{4?{>T@12e)3d%Fpwt^#%G=B`WXzy$N35YJLxx$L!FX}FuqQ@d|(Qy@kcQJ?du4} zADRLfzvtPpFuvrPf7dTb7hrt%*1FS|!1!lIVEhFP#=nkW{MNg~F6C$-!T8~T@i7ay z42=Kw2#kNm^oR)Kw_g^~;p@m!>cn|0H1isSh|GR3A{k(B#+M z&=o<@_whN6ldzzAmQa4xVPk*lv`D~`Qd$6($DO2Y6ZSHj-y3n+mQ@1XCH2ky6Wb6f z%UN!ZS>=mHu@kAQafHhv6I@n@w&CRg$F-~z8t0wke8w(waoB&C@Ao$!5{sm*@%(mAo(x~ro*otPV@v0(Z%aIQI z$jg!d&?~K0x$zKd3E!e6C^xCL^)Ga`>y0zNd^At{ROjet@sM&rzq8x8wZGexlKJlzQwb~C z0Yx;%EM!ohYiB5ZtM`_PrG30isP&D z2xoM?e#TeHjIRGkk@_O&`cnd3?_MLl(}S*mhST*=gRbuxq3at$vUHww{h`ay_0KK- zhSBx&NKtDch#=xmNo1tmQB`P^4XR>ST+LN zH&tLags^k+g;&;WdCwe?b4=H>#&+|wNu|a%^XO!4@X&@B9E!nq=l53p`T5iXia#0e zoT2OageJ-6jr6#8BL1L!fokHpCVKMv&}eS8?C!j(Q5z*psQObep<1BM+9y96kh-Ah zEcGz`zXM_;o&0fx?&LVy2AZU~8+;d_bv4joNuc?j-_?}iyE5i|Q?0jha>J)8<<3Xg zc3%~aFwg?cFW1&#cXfloppQ5OoBE`8KdX<@;vTJ*!-jXOdz<~Odn`?V@*zgH-d*VP z2`u}XGg6}Jk2ZP7A!fzm$(?is)zZm5Ivm@gy4BA$P~RNX zi(s4GG4cXsu}>+!s8IadqxSace_-Itw(jlUtka(|?+wRR&4n@E$^v^RHnLCKPX8ZJ ze`^bT6{p4g(RHu)c|0WJ5Rq*21_J1^8CGv~v90IVgx8dQv&>XWjS^P`o(up$c#tD}{>NR7{why26ss?yJ z#K8LxjU8L5BRy_Y9I3hYOaS4MgIW6$wh~x;VMFK) zJWsRU7j3~7Lm+mz&fZhE{zSL9>sR{MpNB2lg_B~@=|*3TPFJ(NN!1VlXf;+#o24z^ z(4#k5)r({BKXySZ zYpkPP5Yq3H#wb_&3^D7L=ikx<;UxnR?{`sBG z`6Bf^KM3sCC)&)ohL`3KRm^UagTe|FED)*Lf(0T<|KjZ<)`aZsS1>bVZ$EkwTLGLm zz|^j6){JZg&;_~<9e8#vtdOug^yd29;NKBC>thghxe!~c2UwXx8y(riY-68IYFe%&Q$}>_#l5lzT(93; zzOijyoAG{MCARx6TA-$$3yV|Qa`(LX_|jo|s$0i8i-_rkks#p;ee!XLfn{mL^!JU^ z0qvs{)#bEMsC16cl5UgIvhODC)uC#36Lf#>m82a!)_Sg7uVbuzn(fX>7fZbk|2j`4 z4OAXi{YQi8IrQKiN9QRWw!2L;l@@KOyw&=$LGjpkqd2+OX0^;TL7t_4Rj*$@9w4%D z0)fcKXwf@Ace?XWWc2|sm)v5CCGN~NE-^@*Jm|m|+ZLndJ@1)Xtc8IS08+15=bN}` z;26f@-LmU4Qw(p|$0b=?%sqCWI4wX^h1DT;JIAgMK&Pgg!Mp9u_?}mElUz@Cbt+xB z_?-rzk;;dMcDRE>)oe44<_k*ChSSrr8K*1J3un;m!NblzAJLXIA^UBYRG8Tm%c|9J z9qWDXjL(udyRgmRacy|y$Towx5QTGra!|V`7xZ?%5(~IarE44zZrxho+3O>wyNnIT z`>&GUsP>)TEkSd;nO1i8bZ+m@l}6PZw`?;R@5SM6tXxyCQ#uwB4xZP4HwLPl<}hW2 zcF>k)zn2KjAw{-gjh}W?gf@2Y3+;Km23YvCv#^RiHso_*d@XttO}7zw`u^w)*|@I^ zb_X{KJgPpbj|Lz=PJPaxwFkLg!67gfuSeAoTriwH?w;bhLtSdXocBCRiBoyMM)jo> zhTQqs4D`WAhmR9g-3j}J5xE-D&6eM(9yWA*_p^~RW%jGLV28#+ISoLKpy!MH6q_&{ z2RFQ8^wK?kH2b7wK+qx~mpMaYJXID=m&J zyb4*gG(dipAsT$J+^l_5r)3{q*F?XmYbkFjHK_KzfL#kQ!>@r4mfefe?um?b{;A6{ zV$5Hj5d8}5qbYs|G_DW5PfQ!{PljpRZmG7pcZZaQQboz8VVaw$uqVdk-xp6U^EbvF z_VtXycYwk-8bTQYs;E+v;b2+|80ysZUI&J^cEIG^d!v5GO$~nIdIj!$QNSi`dp=%` zmj*EP0qpIlyIEaxx__HAQ|Sf!#>kfCEeR^+IOou>`zOf1oMO;306N9%OSP=(0)t0x zimJsNJ?LPrQPLf>EmnCqo(}eHU8h*EX{g&c!$sDtLNuSX&JOK88RNSa6^8W5;+zks z*mFDYHzk%D6^FH$LA!o-aIseV-o3bYj95-NsEu9}qK;QQuf3RS=@zzm2X0%del;F6 z?#H<^0uE@K3Yu10^sx=c?t=Z#i5(mb$gQJv>!I3@kII_W0+ibwbfQp%a-B;q!f+Ik zVV}P)-ab%o55TNoVD_mf23|ZnvaPM(?zUFnIJ%lSTgcR`rqul{)CRY&XQiDQ$U{ z2VtKzPeKnj>|4Fk_cjdKEx$5OYIRSV9xKdDtP`Br34-;66|mm^M^MW~)p~KW&w$VQ z1U!lq`3&Z|W#%=laj)v;VDk9=gwJz({QQO!?2Z9t*VVMmrfg25>;>8>4%7nH{7y^M>dzVUN^G2q#qx5OXqwQnetd6 z%D-W{T&8;y`>(LKsICaRCr;&Sc1ulj$;QNh>2sYcT&c7zXJ!W44A_6kiR?~4L&px^ z6$#ie>Ruwp%%|7yMU}65MS`8}b~1X?Zl@hDsg(1le>CBo>=v&VdvCJFfup~3!Z~N7 z*NgpYtZ{`QXPk5D9095MTMuZ?r{;SdM#q(&AaQ-B4EL|&6tUz+*GuXzth?P*FpfEW z*_x4qP_q1o;`?YL(o#yn(2keU62H%NyV~=*<~01qp4Up;{bc{6?!G7-Mn*J!M{BZ+ zf%p+gsZreQn(>}{seH1(Dqk*L=_psmA8&fVgi7^_43WHDs?)XI#DMI!62b!>j|c^U z{r00>UW!d#z(rW2@UI(II}M6yNiMz)W(>_!`IixZKaU?>m83aG0RHre{(LD+T0)O6 z(?X9HtCNdw4@Fn%DrsI~yav!XO7(iVzN1mM!TnnaJ4x%^k<|VUKV9Qr7_vSpDJveG zNcy_yo%#%rMkkFsk2=2@pm7P|VEz$iOwll*0E2*~ktWuHp zLWKKw#y0mZR#UB^SkgQY(qF9qsOeXxs;sFQ;<}Aq_Si=fdfJ*1hH8Nep_afZ`@NBVPdyTf2|};KH<@!(B_SrAMhm56fx^vOD5Xz?9O}I``#E?G;vqB z2FmAl^euoyG9N?l2fAwPa`_@LTD8z$8PWN!rTAS-p=+(Pk3SS?Kk8f|9~uGXmm~v- zUC2ju*V_oH8DZ@$6xXW>*n#N}O5|UHr6=fj7XiTk7qj#rhr!Y#r~gw~`uY!vrSD}d z{ZV4+U5uqKW-R^p#L`O{OW#N=y^^u?X2#NQA(mdwSb8I4={FKfKkz28^v4)WKaW^? z3uEcOVl4e5#L{PrEd42wrLSQueFw4hHpbF_MJ#=;$kIzimR`$P`WRyAt&FAbV=Vnc z#M0XtOaDG&=|2=$dMjh;dl^gr{{82094}+(OBqZ5#j#L~MNOMmftqT*EyqFSy8P<5@xmw{bZj`I!t~Uu_^3z5i|#v36@&kdut7kC-H12CDxI3zBi1P9RQO0D2mG zNen8WUA^}*nEydFC9*%C!RN8^5(H@t$%uWf02}>4*}2M;tH9O2l|Vr{A7K!f)}IpA zU#vUav)iWYz2elbTGL-7T!UiBeM*ayn{fZ~6jL?xzwf%4+}j(8q( z`)}Aqzw_T4^80wZjbez{%3{!mvwLcQVn_?{4&KSQdc@G%&9bHsz%GMT&6|X+4Zicl zQ=Bcw*{_R0-2=$NpH(N3{HH8EPJt!e{|{Jtf5ktIr7xBL7g&0Ku=J9unQajxEd7T@ zu=LDAc3GAlPh4Q>$+#x6^ecaor5_l}(%Zn&>qM6Rq=vKfpNlNLO!l|4^kq7*^wLjO znbXFw^kFKn^khl8JWCHICe7&RYRTg){pUZ$(kJa1!_tSX082krVCkKqxgtxC)gFPR zH~zyceX0Dvz|x~a{&AMxA1u9C{i7^>k_Ie2ox*TgmL6tOu=K0}2rT`|pJeF=6k}L= z-T;WD7aIUr`p-p{UMBk|vh-oZ(%T6_@?hzeGO+aEnXbUnhpE8QQ}f|0{bN7L(sz$# z=?B2l;|We+>3ItQOOF$|M3!DE{hL{O+|~k1uLMgkWaDpP>6OINQwB#^`p^G)EIsUk zVChl!{}@Y;QyF}i36it)tNuZjzJ4rA-|gec(&G+gEWNvagrz?@%F=88ODsK&jH_t) zz9LJHS%hga4bY8$BTG*^+ESxvx^*R%o-TTUUGz#UJuW)R(o=+)(-7hRG?w17?T-Hz zmfn6<)_*Tc|2EsBw~6i13&8wg3`=hw{z1AGTNJQe57fMjPOry__2&{~fa0UT!T;Gi z%5x8&Kf?dLjJr3&Bm}WT+xNVBI8UVwHQ&rkjX0t=O7(&tEkS4NdyHOQ=IgLF(S$=O zYo$1pvcZ96g&MrS!#PXd*f&Qu9^9-n3Cm$NbD>6=-SvyD%6$8TJ&yM^`z93GJ(a**^8eR_9vD zgMqMs)0@!VXv*%gb%!-jz}QR#yuspKqcvp2XcewawjJ?jPfC0XeOeZ5eeDUm5){Kw z0H8#iTr@D{dm|2)k)PC_oI*>4XMJF%vhQ(!pqp;HXJCTrHf_Gj;+XifhPLwU!|o_L zK1@FQEI?{dpsIbRWAF7ki<3Pp#{$MR^gG!(;$G6Hv%O)P(ZRr0>vo@WUu9pHj$ElI zap|j#212Qw9THdoH+8IlHA1_sD*Ca>g7JBY4m&G;xJxrAMGV1+p=F=No!TCm`d~p* z^fB2TbPV}X8678!gTLBcnK=L9J>%BSIVV@Gl?2oJ;h^53iff^xYwQ~L-Mh$W8;LcL z<-V$PmGa^3BWvj?I{`=6ZcSCkXt3OkgGC#VJFayBoHz zm6f?u`yYyZCsyL8#)?m@Q7A!IX8+>$?mFFJ3vr1$lRfnWwS(<2W*wZ}=0DplI{>q# z1qoNisd3qjS9gK}o^h}j)p1*7?B}-$u>6rA6`jR0MbUw^9vM+OM{qc7jADD#3v27Y zNQ-sjqLL^qrJxASkIs?F{l1Wa8&~-SlU=Q5plRZ-f(G>wDI?|8^c$3yk(R!t>K82T zw5Dx-6O*#fD(5SzJ8ahOt>qmNsvxwGXt_aEAEpeFx*{{=%=l-3nJ+6tos^($)5>du zTyJ2v+^dqCk6Gf8hXkFxyZ+>usvBHfi>fQqn+?i~!OCsGdmE)abzN2*R+_3AJ{~I% z-G`Ha4a#@4)lBJq_~xCZ#Ld&fu@*)<99kI^lFxpf9Fisa=|vWKIXe}myMB7dE2t#zv?gH+Rdfff zE>qy_AGdNNR>${s8DC3Zw86Nj#JH%yvZz>Zjvu%;wC{TS3+W34!a5$Oa5}E(*^YM@ z;ebx%uyF*E|_N=YcK&D()M>hj8{S&3wPCOB5>ll>W`bgas++5`0P&FFi}No(&PP+qW)L#bO_$*Q(OTH!HI0gT6?RZ&VWbOH9jr-ON_+g|6X?u_ZG zb~QA-uK&mw<1l`gs&6P|JJ?)zvG>kodVLcyp@%J58XI;|ns-ZZe$j~OZ%ar}U&}N* zHnd<2io-bYvwC{bP&XZL!S;jc>BI&zdxcRo?LF@rKlMUiJkD~$q#~Q1#C3jj>IK^s zm^O!u6%F36;T(gn*|z+skfzi$~OW4_M^?6I3Y zqyHUasbb7niVbGMo-&E~O>B&#XZU}7EM*(68Y9*h8EYfhk(;u4J9Ry5Z=tn|z3n7X z++Ee@h3-wKjHj2JC?iz&sA(O%M6<0Eg43O=qt%vAL&pI^KUDAV>RlEH_L! zBK>2%(6{cDNz45w)`;-Ad&XH^OP<|jnus06@$D`go`iRkM0-Ejuv}#yc3N;--ezWR z21wFmM`+XZH3LKp9L0)#_DtI5eM~dh9T9lvz@$Q-y9boBr_m8b@pPu;()w0mfRo>P zcDwPy-CfSrcJy8rIxX0VcQ`ahMg4_N8oP<^>4ph|-O}+Vy21;6u*0Y;uqh&h4UY0# zwL2uUT#MT*6YZbS(8~s?sjKy;hLo%`@EbSS-8H_?0q%@W)HY7;eM~)nt{qDdpR6!3 zK(AhN1h5V6DX^9^?O08g%vFA~ji}DsWayD~Guo0GeD>tkdO0V9Im~u%jX5U8Nz1cj z)CZp*F#26Uud)zFK0cH3MxJy&>$`*}HVh4xW%ii&rRdxn5Y^=9=&Osp-4lb5t5nL> zt_kKwT8YucVR$fw|MGqe53de{-`z%|b)nCbe0aE>P=GqbFdZ9QH2WHw3PYy)?l(@s zawm8f5w;xEVFQ3^t~p}J@A<5Wb_Kh+1=9jQ8CH9Il7XS{}|`yP*RMT}Y_GJ&!VD zDyuX3JAxlkHLR^K9T(H%gXAyY;1B-j6K!_t_8q2r_q5#H-Nvvs&G7UciGH2JD5eog zP&ckNs5B)aBd&=?{fuC@<^Z(=8a&JVHl+Erb;g2_Zb8BwelA>gU7)@tuis~OBF5-P zR6Abyn7s%_R-rgbgQ>SPj>bT=BIO)4+L`HeL;3F3Z#8~rana5Owpm7V^3xrioz558 zh=$Yn87#W@m5qw96{>rL-CJ%&^xiorFEimu5?7q|tme#sbF2Tv7aRf7tK!!8*;>QO zRjCS8%IiiN;!WwDZq+-T>~uSfor(#;t~cSU#j%w9*PwjDu2zSwqTOFAKX#n{`GNLH z19)G>9uRL`bF8)&_|katZFL$Nw(J@#%d{E0c3*m>K{vVT)a(%zlRn~rI*dKq7+sI+CZ{(5 zSx>!n&lZ9zzliFvSK;=tgFq1vXk1@1E-9qE?cNmH!k1L*vX?HQY25h;ltNj*vg;#n zkan$$oPH?}I3F<7YYz5koti1#{?140aI9=QM$~KjQ8we=Ru&!;{NiIH$$WYwneEHP z5%fZ5Ot(1HyPi?J>pwL$rn}npltT8MUzhGc5Zm#sldY}mKM0>0MHO5l;P8v%)*;slFjinAYmL6nd z=>k5MZqi94ewYu$2WTSeOfF5v$;pnzyMv1r^Bp&TJzZP)j5Bxptc=1zSHbqw;#X(a zmHui8v?Y=PjGaXoj)UdR`UUt-BXcPRQWgA6I7~An*>Lnnz74L=S0cFzel_&y{H_aL z5(&1KOB!ha#$S`@(`VABzgwU~d;7UhU4lXMhIeeQGl9LULLz+ixL{gq-_a|4YcJ9JtTrYBaD6$hptuS&&ox za-Yk~Ag6laK9}DIImOR?E`JvCagcML%R3>b^mCuf`$T>2bNQbjH$$KMB$uH5EI@;& z>vEs#uZBDWa_)2aBalA=Irn)y3vn+!3pw|>{4>ZY9`5sa45-BGA?H5Vw?qCgVGNlxjnhI~5Y+~@iaLQee&_qqNPkfVwv+~@NBkkgpKeHIVme+zOn z^%0-z_la`ubNOk=>AvGWkLNp4pZi?jTPEmppX*0LPVJ5RTt5+VDmU(P{WXwN{c@k{ zS3o`!a_)2e7Raf5xX<+;hn(6i_qqPdkOxA}eXjo&OrFUk^E@pZi>Y9psdL?sNSb$WfgV?sNTm$aRo&pX)yfIl4p% z_qqPBAgBJH`&_>Va;h)xbA2);kevHmJ}An$&*k5Xa_)1v1~(1mm-}2E2067i?vq@B zj3$Zt+~@ixQJ?!T{p#J4Jo&bN!v7KKHqP8|0W8Nx0ATpM`up2jl@&Cwgr&n%E>RNmkie3GRM9|BLFNq*t$bDf&RW9ExW4DtF#!_vIU7)$bFHPeVr$Mi!5^a$ug@ujh9TklOYfhA)jyJ4 z`yo+#W6Tfn!VhWqA#KbL2H}TQ_+izUA6O|H^mJe5=j4`HtH%6F)yvW< z+`0KV*1YPH8f$e)MfnJ3()|AB+l~&}UU(KtosH|dwyoxegNohf49J|8O%*}*E~AuRP*_rqyo|ll(#6+Z7bjeVUlTBWm{%s0+Mx0qj0 zQ<7JZm2*S9_|JT!ArbQG;!5<@tQ)vNp97%*LY~5F(`yoQHXw*+9L1hrR?KCINXCUL z%1ZKaTuppQQlj7+lBg8RO3LxIBq5%@aJ@q46&jM_ zm@oWvOHo{|wX(7VeH#=D%dDg#{A4XHrFadVF|aEcgbRsUEX!gGiK8KbUzH*gqbjsk zRp%7c+TzCupioPKv88}kfh_TI5K5FspI2t9%_}a7yQRRokzZ!4MirVUT9S~q(LyY< zq5^7jQfgCj$~Ko)m)J^gMXS%-j4#ykYLXI>4(6eF=$q(?g=EECVl=XX8no{!Dwa{% zCAEYKVpZi?c4Daur7zUdXt|7(Iz^dNRlK>du(TkjwgPn`q=Bij5@V`p=X$G+vOyg= zYvR=#(rYl-dp5kR9JAOq;}b{IDi#J6m9?l~*$sxA#fv=?$|@PfAyhVNPCp9#QzbdJ z2uHpl$w}xMM~1h5k$Y3T6))Fu6#lFB8)Mtv-`=>zHa_BA(u$DZ|578mYz_baUWg+V z96vYba)A5T}pRvjV8bq#`e;q`ah>6e2O3Q^w-eXOlXgo!!~ zIv%>Yh}VwaDBh3Z=z)05Ga$a#dFbB5RPnNo{vrVD(UeE z(CrdRQ^VXya*f2BT?-;98LCN#;-q#cK4YqNO>YdBN-KG0fvmJv~Wg?C6nMN3>mCxVTlaIWE>}hIxTijfz!c7 zz(vAE!J(P3zp3~?1BKF$0&&9?!)f65p}`s9PT|J2!==M%;Laj}2jRBBnc>F6oxx4M z56%W>gww#CMB(j$D~8jnkPFQc?dM{7m}S zdFbms^w*2}3q}3cp-cVHwQP1Sk*xR7Z4h;-?kUdKL|vo>Q*89(9=Zob-7{i-dqiEN zgXtdh(Crs>X#tMny;A>l)+5Y#);~(5d&MwC2zP1!r2TfZT+&5-T3n=ST&aJ`^axWc zhB1m^3J~r}{ZpMsxOOp|N(_gh8PgZFd+44Lb7x4`O-hoBVH}?1l!q>n;t-AW zanwKHc9k&GkNCL|&x;s7B>Wnb-b=?18zgis>U+=)V4PzH9Z3djYr>dOEnbVKi~2*- zO%-*kJ#^idq>Ji-bWNfz%a7Ri?(+!OCx)Z3pVC8J2tG=#g*z$gQTi{}&xea)NS_=n zo~u|p;$wpbaZ(REIyRieH%bfXUk8^0eZ26S$(}}9C=JvYB0jkSc9ub z`5cYf6kUtbR&C8Kr8ZCFPfbZx3BHY9O#E>RYrhJK4m#oZpTbf6N_GX|8Z^J8v5N9U z02%rPNg-p0RVL(e_PRD%QSJ@QJ?hwnXWriXa)a3o278PfDlfvjP29TvSk2 z2A}3x`Qkh)OPpsx%>H!v|9{GSYaudu5}9p=TMwszbE5zf;d+tbYPb|Q1>7kV5G|(L z;L_n#aD6B^H{3sKzSZe5hk6=)YcJxadbr$tYtTbi%lcu7l+r@_m(I7e9&;5L`kHJ? zKlH9N-^%dNkB5G=sDGvTR;`DAF7$`RbezI9X#BeL8LG`g|32uSMI)s2{M3Bwlt&n0 z43S8i#OqvczNPk;=%p4>|8n!Kg&z9FqP||#e;tA8zToq%Vh`PBQCA`AUOL}u z_R#GVb^FBp_K5NF`BtZg?x3i90-Xw#&6Va`=RCsjXXzF(OcBCSU+(!lIMrjWVixB$ z#iD*8bZI=h^m#DbBaHBTEs^r)+XCoQKO;QfkB)c{Nt!&u9TBf}5pv2;J{)ehk!Qan z9=c~m-3GDUT{_=7>!CYd=zpaI(onoKHsW?2nQ#5?9P5AQSZEIa&(E=Nd$2M7Uu}-X z)Ajevv3QtbF;D2eB>#*#7LWJRITrON{F&{)dyZA*F~_=y&YNsmgVc$Oe>(Vo&-}^` zeR{~Cypc_c^7hZPUrj}P)8MAV&47!Bn+ZpCLVq!EotPt9;G*G%(Mk8fmBLYn-Ht)P z0ym5SwgD~{t{(-y7tRKk0XH6w=3nh_b#NJQTDY?qb35P~;IiPR!qHr!9c~L;CLFyr zlWZ1kaJ6vhlrXpx@GiA?W&UN<7BLayiW=9TWL|FmRqLU9kc~$YDYuVZI{(`1p<9l4 zkBISJY5sN6L;n%zH;VdKntv%g=7RM6%g4ei&A`2p!tuEkXPqFD?Df#^7xk&$NdGlBR5|SOg3YPdLsu=%ff_~K9#L1|ov`NMq08rZ z*{E-d_e%36y~jMVRE&SB7{-QhWYfLOa+##qBivpw+#uQ!U8@|9<}xDUBx&~0?G<%9 zMcpFkT{H?Tw zHI`4S@R7~!`B>~mJW^;-8=zjqb8U#?qBf^tcLI-#)}e&Y+^6~i!o{w|?QvP`p5->0 zsjPqEHW~^iSZH)3b|$im(eLD%ArmlL&D;kk`gPKE1k&MUT7LOzwZ zTqHk8Z5sAM$~Wm#*^mo|qu(j50&O~CE@b6F<(!D$gl9sYE*<-g%K4vdmz>J(J$@Z3 zF=|(%*P-vDcEeG-e!&13i57*|QE85~m?2L^bx~uW| zq=&9xYm~?-9qx*mPL7LM}qW4q#L9;pxGf1`PnP_rMYGalbc$mlb-js7ij-E=Vw qzXzyaqsS>9YWp;{QyS<##`KT%tJJ=zZBQ7x?@2xaFXA)V7yl2wqvIX` literal 0 HcmV?d00001 diff --git a/candle-metal-kernels/src/tests.rs b/candle-metal-kernels/src/tests.rs index 59f54fa9..5805206b 100644 --- a/candle-metal-kernels/src/tests.rs +++ b/candle-metal-kernels/src/tests.rs @@ -725,3 +725,66 @@ fn where_cond() { ); assert_eq!(approx(results, 4), vec![-1.0f32, 2.0, -3.0, -4.0, 5.0, 6.0]); } + +fn run_gemm( + (b, m, n, k): (usize, usize, usize, usize), + lhs: &[T], + lhs_stride: Vec, + rhs: &[T], + rhs_stride: Vec, +) -> Vec { + let device = device(); + let kernels = Kernels::new(); + let command_queue = device.new_command_queue(); + let command_buffer = command_queue.new_command_buffer(); + let options = MTLResourceOptions::StorageModeManaged; + + let lhs = device.new_buffer_with_data( + lhs.as_ptr() as *const core::ffi::c_void, + std::mem::size_of_val(lhs) as u64, + options, + ); + let rhs = device.new_buffer_with_data( + rhs.as_ptr() as *const core::ffi::c_void, + std::mem::size_of_val(rhs) as u64, + options, + ); + let length = b * m * n; + let output = device.new_buffer((length * core::mem::size_of::()) as u64, options); + call_gemm( + &device, + command_buffer, + &kernels, + "sgemm", + (b, m, n, k), + &lhs_stride, + 0, + &lhs, + &rhs_stride, + 0, + &rhs, + &output, + ) + .unwrap(); + command_buffer.commit(); + command_buffer.wait_until_completed(); + + output.read_to_vec::(length) +} + +#[test] +fn gemm() { + let (b, m, n, k) = (2, 2, 4, 3); + let lhs_stride = vec![m * k, k, 1]; + let lhs: Vec = (0..b * m * k).map(|f| f as f32).collect(); + let rhs_stride = vec![n * k, n, 1]; + let rhs: Vec = (0..b * n * k).map(|f| f as f32).collect(); + let results = run_gemm((b, m, n, k), &lhs, lhs_stride, &rhs, rhs_stride); + assert_eq!( + approx(results, 4), + vec![ + 20.0, 23.0, 26.0, 29.0, 56.0, 68.0, 80.0, 92.0, 344.0, 365.0, 386.0, 407.0, 488.0, + 518.0, 548.0, 578.0 + ] + ); +}