From 17291019e34ecb2f56da007c50a9133718328ef2 Mon Sep 17 00:00:00 2001
From: Andrzej Janik <vosen@vosen.pl>
Date: Wed, 3 Mar 2021 22:41:47 +0100
Subject: [PATCH] Implement atomic float add

---
 ptx/lib/zluda_ptx_impl.cl                    | 115 ++++++++++++++++++-
 ptx/lib/zluda_ptx_impl.spv                   | Bin 50100 -> 105668 bytes
 ptx/src/test/spirv_run/atom_add_float.ptx    |  28 +++++
 ptx/src/test/spirv_run/atom_add_float.spvtxt |  81 +++++++++++++
 ptx/src/test/spirv_run/mod.rs                |   1 +
 ptx/src/translate.rs                         |  78 +++++++++----
 6 files changed, 279 insertions(+), 24 deletions(-)
 create mode 100644 ptx/src/test/spirv_run/atom_add_float.ptx
 create mode 100644 ptx/src/test/spirv_run/atom_add_float.spvtxt

diff --git a/ptx/lib/zluda_ptx_impl.cl b/ptx/lib/zluda_ptx_impl.cl
index a878ddd..85958d5 100644
--- a/ptx/lib/zluda_ptx_impl.cl
+++ b/ptx/lib/zluda_ptx_impl.cl
@@ -1,7 +1,10 @@
 // Every time this file changes it must te rebuilt:
 //  ocloc -file zluda_ptx_impl.cl -64 -options "-cl-std=CL2.0 -Dcl_intel_bit_instructions" -out_dir . -device kbl -output_no_suffix -spv_only
 // Additionally you should strip names:
-//  spirv-opt --strip-debug zluda_ptx_impl.spv -o zluda_ptx_impl.spv
+//  spirv-opt --strip-debug zluda_ptx_impl.spv -o zluda_ptx_impl.spv --target-env=spv1.3
+
+#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
+#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
 
 #define FUNC(NAME) __zluda_ptx_impl__ ## NAME
 
@@ -25,6 +28,20 @@
         return expected;                                                                                                                    \
     }
 
+#define atomic_add(NAME, SUCCESS, FAILURE, SCOPE, SPACE, TYPE, ATOMIC_TYPE, INT_TYPE)                                                                                    \
+    TYPE FUNC(NAME)(SPACE TYPE* ptr, TYPE value) {                                                                                          \
+        volatile SPACE ATOMIC_TYPE* atomic_ptr = (volatile SPACE ATOMIC_TYPE*)ptr;                                                          \
+        union {                                                                                                                             \
+            INT_TYPE int_view;                                                                                                              \
+            TYPE     float_view;                                                                                                            \
+        } expected, desired;                                                                                                                \
+        expected.float_view = *ptr;                                                                                                         \
+        do {                                                                                                                                \
+            desired.float_view = expected.float_view + value;                                                                               \
+        } while (!atomic_compare_exchange_strong_explicit(atomic_ptr, &expected.int_view, desired.int_view, SUCCESS, FAILURE, SCOPE));      \
+        return expected.float_view;                                                                                                         \
+    }
+
 // We are doing all this mess instead of accepting memory_order and memory_scope parameters
 // because ocloc emits broken (failing spirv-dis) SPIR-V when memory_order or memory_scope is a parameter
 
@@ -120,6 +137,98 @@ atomic_dec(atom_acquire_sys_shared_dec, memory_order_acquire, memory_order_acqui
 atomic_dec(atom_release_sys_shared_dec, memory_order_release, memory_order_acquire, memory_scope_device, __local);
 atomic_dec(atom_acq_rel_sys_shared_dec, memory_order_acq_rel, memory_order_acquire, memory_scope_device, __local);
 
+// atom.add.f32
+atomic_add(atom_relaxed_cta_generic_add_f32, memory_order_relaxed, memory_order_relaxed, memory_scope_work_group, , float, atomic_uint, uint);
+atomic_add(atom_acquire_cta_generic_add_f32, memory_order_acquire, memory_order_acquire, memory_scope_work_group, , float, atomic_uint, uint);
+atomic_add(atom_release_cta_generic_add_f32, memory_order_release, memory_order_acquire, memory_scope_work_group, , float, atomic_uint, uint);
+atomic_add(atom_acq_rel_cta_generic_add_f32, memory_order_acq_rel, memory_order_acquire, memory_scope_work_group, , float, atomic_uint, uint);
+
+atomic_add(atom_relaxed_gpu_generic_add_f32, memory_order_relaxed, memory_order_relaxed, memory_scope_device, , float, atomic_uint, uint);
+atomic_add(atom_acquire_gpu_generic_add_f32, memory_order_acquire, memory_order_acquire, memory_scope_device, , float, atomic_uint, uint);
+atomic_add(atom_release_gpu_generic_add_f32, memory_order_release, memory_order_acquire, memory_scope_device, , float, atomic_uint, uint);
+atomic_add(atom_acq_rel_gpu_generic_add_f32, memory_order_acq_rel, memory_order_acquire, memory_scope_device, , float, atomic_uint, uint);
+
+atomic_add(atom_relaxed_sys_generic_add_f32, memory_order_relaxed, memory_order_relaxed, memory_scope_device, , float, atomic_uint, uint);
+atomic_add(atom_acquire_sys_generic_add_f32, memory_order_acquire, memory_order_acquire, memory_scope_device, , float, atomic_uint, uint);
+atomic_add(atom_release_sys_generic_add_f32, memory_order_release, memory_order_acquire, memory_scope_device, , float, atomic_uint, uint);
+atomic_add(atom_acq_rel_sys_generic_add_f32, memory_order_acq_rel, memory_order_acquire, memory_scope_device, , float, atomic_uint, uint);
+
+atomic_add(atom_relaxed_cta_global_add_f32, memory_order_relaxed, memory_order_relaxed, memory_scope_work_group, __global, float, atomic_uint, uint);
+atomic_add(atom_acquire_cta_global_add_f32, memory_order_acquire, memory_order_acquire, memory_scope_work_group, __global, float, atomic_uint, uint);
+atomic_add(atom_release_cta_global_add_f32, memory_order_release, memory_order_acquire, memory_scope_work_group, __global, float, atomic_uint, uint);
+atomic_add(atom_acq_rel_cta_global_add_f32, memory_order_acq_rel, memory_order_acquire, memory_scope_work_group, __global, float, atomic_uint, uint);
+
+atomic_add(atom_relaxed_gpu_global_add_f32, memory_order_relaxed, memory_order_relaxed, memory_scope_device, __global, float, atomic_uint, uint);
+atomic_add(atom_acquire_gpu_global_add_f32, memory_order_acquire, memory_order_acquire, memory_scope_device, __global, float, atomic_uint, uint);
+atomic_add(atom_release_gpu_global_add_f32, memory_order_release, memory_order_acquire, memory_scope_device, __global, float, atomic_uint, uint);
+atomic_add(atom_acq_rel_gpu_global_add_f32, memory_order_acq_rel, memory_order_acquire, memory_scope_device, __global, float, atomic_uint, uint);
+
+atomic_add(atom_relaxed_sys_global_add_f32, memory_order_relaxed, memory_order_relaxed, memory_scope_device, __global, float, atomic_uint, uint);
+atomic_add(atom_acquire_sys_global_add_f32, memory_order_acquire, memory_order_acquire, memory_scope_device, __global, float, atomic_uint, uint);
+atomic_add(atom_release_sys_global_add_f32, memory_order_release, memory_order_acquire, memory_scope_device, __global, float, atomic_uint, uint);
+atomic_add(atom_acq_rel_sys_global_add_f32, memory_order_acq_rel, memory_order_acquire, memory_scope_device, __global, float, atomic_uint, uint);
+
+atomic_add(atom_relaxed_cta_shared_add_f32, memory_order_relaxed, memory_order_relaxed, memory_scope_work_group, __local, float, atomic_uint, uint);
+atomic_add(atom_acquire_cta_shared_add_f32, memory_order_acquire, memory_order_acquire, memory_scope_work_group, __local, float, atomic_uint, uint);
+atomic_add(atom_release_cta_shared_add_f32, memory_order_release, memory_order_acquire, memory_scope_work_group, __local, float, atomic_uint, uint);
+atomic_add(atom_acq_rel_cta_shared_add_f32, memory_order_acq_rel, memory_order_acquire, memory_scope_work_group, __local, float, atomic_uint, uint);
+
+atomic_add(atom_relaxed_gpu_shared_add_f32, memory_order_relaxed, memory_order_relaxed, memory_scope_device, __local, float, atomic_uint, uint);
+atomic_add(atom_acquire_gpu_shared_add_f32, memory_order_acquire, memory_order_acquire, memory_scope_device, __local, float, atomic_uint, uint);
+atomic_add(atom_release_gpu_shared_add_f32, memory_order_release, memory_order_acquire, memory_scope_device, __local, float, atomic_uint, uint);
+atomic_add(atom_acq_rel_gpu_shared_add_f32, memory_order_acq_rel, memory_order_acquire, memory_scope_device, __local, float, atomic_uint, uint);
+
+atomic_add(atom_relaxed_sys_shared_add_f32, memory_order_relaxed, memory_order_relaxed, memory_scope_device, __local, float, atomic_uint, uint);
+atomic_add(atom_acquire_sys_shared_add_f32, memory_order_acquire, memory_order_acquire, memory_scope_device, __local, float, atomic_uint, uint);
+atomic_add(atom_release_sys_shared_add_f32, memory_order_release, memory_order_acquire, memory_scope_device, __local, float, atomic_uint, uint);
+atomic_add(atom_acq_rel_sys_shared_add_f32, memory_order_acq_rel, memory_order_acquire, memory_scope_device, __local, float, atomic_uint, uint);
+
+atomic_add(atom_relaxed_cta_generic_add_f64, memory_order_relaxed, memory_order_relaxed, memory_scope_work_group, , double, atomic_ulong, ulong);
+atomic_add(atom_acquire_cta_generic_add_f64, memory_order_acquire, memory_order_acquire, memory_scope_work_group, , double, atomic_ulong, ulong);
+atomic_add(atom_release_cta_generic_add_f64, memory_order_release, memory_order_acquire, memory_scope_work_group, , double, atomic_ulong, ulong);
+atomic_add(atom_acq_rel_cta_generic_add_f64, memory_order_acq_rel, memory_order_acquire, memory_scope_work_group, , double, atomic_ulong, ulong);
+
+atomic_add(atom_relaxed_gpu_generic_add_f64, memory_order_relaxed, memory_order_relaxed, memory_scope_device, , double, atomic_ulong, ulong);
+atomic_add(atom_acquire_gpu_generic_add_f64, memory_order_acquire, memory_order_acquire, memory_scope_device, , double, atomic_ulong, ulong);
+atomic_add(atom_release_gpu_generic_add_f64, memory_order_release, memory_order_acquire, memory_scope_device, , double, atomic_ulong, ulong);
+atomic_add(atom_acq_rel_gpu_generic_add_f64, memory_order_acq_rel, memory_order_acquire, memory_scope_device, , double, atomic_ulong, ulong);
+
+atomic_add(atom_relaxed_sys_generic_add_f64, memory_order_relaxed, memory_order_relaxed, memory_scope_device, , double, atomic_ulong, ulong);
+atomic_add(atom_acquire_sys_generic_add_f64, memory_order_acquire, memory_order_acquire, memory_scope_device, , double, atomic_ulong, ulong);
+atomic_add(atom_release_sys_generic_add_f64, memory_order_release, memory_order_acquire, memory_scope_device, , double, atomic_ulong, ulong);
+atomic_add(atom_acq_rel_sys_generic_add_f64, memory_order_acq_rel, memory_order_acquire, memory_scope_device, , double, atomic_ulong, ulong);
+
+// atom.add.f64
+atomic_add(atom_relaxed_cta_global_add_f64, memory_order_relaxed, memory_order_relaxed, memory_scope_work_group, __global, double, atomic_ulong, ulong);
+atomic_add(atom_acquire_cta_global_add_f64, memory_order_acquire, memory_order_acquire, memory_scope_work_group, __global, double, atomic_ulong, ulong);
+atomic_add(atom_release_cta_global_add_f64, memory_order_release, memory_order_acquire, memory_scope_work_group, __global, double, atomic_ulong, ulong);
+atomic_add(atom_acq_rel_cta_global_add_f64, memory_order_acq_rel, memory_order_acquire, memory_scope_work_group, __global, double, atomic_ulong, ulong);
+
+atomic_add(atom_relaxed_gpu_global_add_f64, memory_order_relaxed, memory_order_relaxed, memory_scope_device, __global, double, atomic_ulong, ulong);
+atomic_add(atom_acquire_gpu_global_add_f64, memory_order_acquire, memory_order_acquire, memory_scope_device, __global, double, atomic_ulong, ulong);
+atomic_add(atom_release_gpu_global_add_f64, memory_order_release, memory_order_acquire, memory_scope_device, __global, double, atomic_ulong, ulong);
+atomic_add(atom_acq_rel_gpu_global_add_f64, memory_order_acq_rel, memory_order_acquire, memory_scope_device, __global, double, atomic_ulong, ulong);
+
+atomic_add(atom_relaxed_sys_global_add_f64, memory_order_relaxed, memory_order_relaxed, memory_scope_device, __global, double, atomic_ulong, ulong);
+atomic_add(atom_acquire_sys_global_add_f64, memory_order_acquire, memory_order_acquire, memory_scope_device, __global, double, atomic_ulong, ulong);
+atomic_add(atom_release_sys_global_add_f64, memory_order_release, memory_order_acquire, memory_scope_device, __global, double, atomic_ulong, ulong);
+atomic_add(atom_acq_rel_sys_global_add_f64, memory_order_acq_rel, memory_order_acquire, memory_scope_device, __global, double, atomic_ulong, ulong);
+
+atomic_add(atom_relaxed_cta_shared_add_f64, memory_order_relaxed, memory_order_relaxed, memory_scope_work_group, __local, double, atomic_ulong, ulong);
+atomic_add(atom_acquire_cta_shared_add_f64, memory_order_acquire, memory_order_acquire, memory_scope_work_group, __local, double, atomic_ulong, ulong);
+atomic_add(atom_release_cta_shared_add_f64, memory_order_release, memory_order_acquire, memory_scope_work_group, __local, double, atomic_ulong, ulong);
+atomic_add(atom_acq_rel_cta_shared_add_f64, memory_order_acq_rel, memory_order_acquire, memory_scope_work_group, __local, double, atomic_ulong, ulong);
+
+atomic_add(atom_relaxed_gpu_shared_add_f64, memory_order_relaxed, memory_order_relaxed, memory_scope_device, __local, double, atomic_ulong, ulong);
+atomic_add(atom_acquire_gpu_shared_add_f64, memory_order_acquire, memory_order_acquire, memory_scope_device, __local, double, atomic_ulong, ulong);
+atomic_add(atom_release_gpu_shared_add_f64, memory_order_release, memory_order_acquire, memory_scope_device, __local, double, atomic_ulong, ulong);
+atomic_add(atom_acq_rel_gpu_shared_add_f64, memory_order_acq_rel, memory_order_acquire, memory_scope_device, __local, double, atomic_ulong, ulong);
+
+atomic_add(atom_relaxed_sys_shared_add_f64, memory_order_relaxed, memory_order_relaxed, memory_scope_device, __local, double, atomic_ulong, ulong);
+atomic_add(atom_acquire_sys_shared_add_f64, memory_order_acquire, memory_order_acquire, memory_scope_device, __local, double, atomic_ulong, ulong);
+atomic_add(atom_release_sys_shared_add_f64, memory_order_release, memory_order_acquire, memory_scope_device, __local, double, atomic_ulong, ulong);
+atomic_add(atom_acq_rel_sys_shared_add_f64, memory_order_acq_rel, memory_order_acquire, memory_scope_device, __local, double, atomic_ulong, ulong);
+
 uint FUNC(bfe_u32)(uint base, uint pos, uint len) {
     return intel_ubfe(base, pos, len);
 }
@@ -136,11 +245,11 @@ long FUNC(bfe_s64)(long base, uint pos, uint len) {
     return intel_sbfe(base, pos, len);
 }
 
-uint FUNC(bfi_b32)(uint base, uint insert, uint offset, uint count) {
+uint FUNC(bfi_b32)(uint insert, uint base, uint offset, uint count) {
     return intel_bfi(base, insert, offset, count);
 }
 
-ulong FUNC(bfi_b64)(ulong base, ulong insert, uint offset, uint count) {
+ulong FUNC(bfi_b64)(ulong insert, ulong base, uint offset, uint count) {
     return intel_bfi(base, insert, offset, count);
 }
 
diff --git a/ptx/lib/zluda_ptx_impl.spv b/ptx/lib/zluda_ptx_impl.spv
index 8a2d6970af32987b1750771d168825f50eb83971..ca16447340e5429c794f3e82484be9f7b9022e2f 100644
GIT binary patch
literal 105668
zcmbT91E8kY)wU<L?POxxoY+>Qi6-{Mjcwbu+QzmTHnwd${jclozE9Uo2j~0WZ}zwD
zg}v5(_H&(+lhdRfYRD;v>(OI~9;5acw#N!%^yqxX=rLrEp;6jlnjN;;5t|*U*^!%_
zrrBw+Be81;>?%9;*<pzlXX(4^<~`^?6y17s)-rToj#)F|J7+nf`o@=U@!?xsf0-U5
zu)jx-=~=zkn|-#eci6ShcD?&<x<#L@x7mi6v#c3&fHm9oQ{%jT`|L)aZThM48UMFt
z`+jQ1>#CXmKgZd03l7sM={nCO1E|^lKW$pgNImA@`8u<D@4ek_n{U#4r(O5#z4i7x
zZP&Xu4SR1o=N$C#vkbR7%^+&~^zPeF?SzA@-L0S6`3F(U8RnVmKS$`CVWvUUa)y3t
zCm&?(Zl!kQ9*YmOws-GM`u6R!%dT5&vh{ZVImj|SM(?p8U=H}+HraK@?R#&s+0MId
zy-T0oo9((u?=836annt<>%DoO&AOgr^d7qd&Hrux);rYwpAOai{e5=mv&+_-{r3Q#
ze<{w=-T7^Q*8$vrI$x08+~0SrO?K(C`G3dr9!KwS17D2)xBXq?JNvKi-u-X6({2O4
z|2MmL|7Cx_^Vj1Gxf|#2>>tqmU){a?ulu_W;Qfuncg=3jAK?9u-#vbu-}ZNn@7({C
z-FyFi_v$;~`(L_y@4xKtcmCtuyZ_Gq0p0&--FyFaf7b!L|NXmn|Lz~){U6Z1`!D;u
z#&_<2x54t?dl&v9HPG)rR}J?0``_R1{GNlo|NrhE(EaZ@SpNTKf7b!L|7QowU)TNt
z-v9f9<*)z!UE@3V|4H}mzkL7b_xt~t-J4(b_dEZj?%jX={?T;+@9&N7-GANRbpZD-
zHCX;T_viir-v5e&<-h;^UE@3VzhU>@f9LCWK!5$+yL<Ov_V+t~$nL%W&i(=M->Z9H
ze|3M?0lfd)x_AHXAK?Ap*uDEN`@6<>?tk4uKR@4l`fSp75I_H%FzDwW!2W*c4?Af7
z``JIB``>8L{10e<*8#l$8wbr_-9Nzle|*sV4PbxQ_|E-5(Y^aG-#@y({5$)9?%w^E
z{r%40y?ggxzkhTc!25f!d-q@WcOAg}(+!&c+@Jdgc>i+@n*Rap?;78^zs0)u{ySg4
z1N#26arfST+28N{54(5&o&5vif4c6y|GK~H0N(#O-MfGH5Agoa>fZgA{axcb_dn@i
z`QKzu{-$G~KmTkn*yrzmf4}pe8tnQ1yMI9UKh<FQ|DXL`2k`#)9xQ)d`v-Xcrwx|B
z{`Yr{@7(`M-Mjzt{iEy4zq9|r?%jXc-|zfI2Friv{Pp`s*8#l0le%~Rb${0Z-2Y|w
z?%(|by#GJDcmHL7*Z9u;4coo>ov+^kegB!Md-Kcwe&=7;z5DO%9}xdPcJKb{{;mUf
z|C<i>`%mZnyMKW9zeV@He#`!@@!bCyJ@y@;TYvvuzW;2p`R2X1D8Kg_qsQd@Qn4F<
z58r=}@_VZ>di*-*d+6u)Q0+bRa}b~99Da$~-E-6)Ify+*;uo^r?9tvwKga0Y>Gh-i
zfAz>QIz5Kqm#5u0N9XJMKVKLBdyY%G_g?z%(a$;d<`>A_Jx9IM0OFlaJj(xHj~v7v
z$9M0&l<zbDdmnR--tBYLcbj8$-pirgdoP`@_kq2a>AJ^9|2_IS$FJRcFP$C(doSm7
z@4fs_j~vAJGFJD#=Il}LG<bCG^p@_um;QU?7@Z#Xj4-%&+x2^p_IvFA<zv&qe!UE^
zM?dFSV6gAyzdZ)_US1#UoesE14&r-RX|P`}T|EZ+UcMPDrvvPfV|3oj5`+DE>GbIT
zdr&`oeAvDB(tnSB&T)J9-b<&)z~0LSgZ+B>@3Yh%If(D&+3vlUL4F_GZLnW2UFWFp
zHpl3^mq)tyUOHd<1N%O<P51cdzehjkSfG1fx1ApKP6LQ{`g-@?%m4JqLF}<s_rB)r
zQSUUz=-lZq-Fq+n_sB6iJ<c2SedO=w`~O+0pL6sY^w-P4dh~OS2?zc4(oc_py_Y)&
zeWwHMk%RbNW*YR@OYJex_wvl3IUQJ!9HaAIdJg*QB|ZB89@Ni!d8m8urT-rNoa3zS
zy_Z3LA6sb9UoZVUOMSOFi0|cw?!A{meji(N&|fe09QEDi7@hZWS@+&c=WBmp-^Z5g
z9v}Vp=;s_0cJJ%9)1%&L0P#-m?B09%pB_1gJ?84(*PK1-o#q&wJAJu(@1_49IYy_)
z27|qiuHVZH^yjfZ2YVj_?9tCTUK#Y?SHjqTdkpNo>^0as9dM5v#P{;)U^(sTG0^vN
z>|i+^V2>Q5^IqQW-h1i)dr&`o96Z?XW1VyK-=m*%tT)*2WB=_ju=nzN_uk9@d@syF
zd@nl;_WM}ZIR^S(#vJU|%K&@i7@haBdH49}eC-eH``B>Z<D>r`{hZ^~?!A{zkAc0H
zeY*E``#(K$5Z}wU-Fq+g80dRBt$XjK{~kF;kLCI8s{cIQto)8|71ma)gIMRVZel&f
z`iSK}gf;<d4%P~+Em#Mz&SKrjdW!WiYe;@qIWcP<)~c*+SckAKVBOAop7kYbcz&nX
zleG|QZPre#qga=)?q<Eh`hhhDzq_27wH&Jt>p<4ItlL;Gu)bl9;_tUvOS3j(9l$!5
zbvx@N*7vM2_#NyltQA;Wu?}Tj#JY#|I_nqKxO^X-leH>qd)85`%UKVy-etYV8VYuX
zW=+YOinRo5N!Dhp%~=Pt4q;u)x`g!r>p|9gtoK<%@x5ti)|9NNSWB>$Wc6Wf!8(j}
zIO}rO6|6^DkFh>seaae<uc47xGqPr4t;kx5wLNPG)^V)kS=X~}U_HZnmh~O$d)C-6
zHx6rV);z4WS?jR&VC~5|gLNkBHrDN|S6Q#Ieq;U4nwalLldu+MEyCJ})tj}v{|840
zp$}$V%DRm85bI&q2docS!|-)8ENg1kG_0jqOSAg0wqPB~I*fHG>oV3utcO`2us&oB
z$M@yo{oc%)fweqq1=hB#?O4aKj%8iTx{mb}>uJ_ktgl(4^ZjEC*6ge~SgW(vVD)9~
z$~uj8I_p-}ZLC*Vud;q+{l=PzueXU=3$YevZOGb)wI6GL*7>XpSogB-W4*(Amo+4J
zGZbqa{<&4~?=iVwQ+)KnqQ_!7eDr6-8uK$_tDg_emsgj+fJr{dFN(Hr{Jv<uG`jp_
z(e|8ze|DEIt4^);S;oPuQ)AtK<HPkBxA=)`r|2<UkBQ5kx^Hr(_iVZU)DF{Q%GSrY
zF>7YQ@@3WK&)$6RU~HmuW35|{hv&<<Q#S=_eSCWuPp?xDKC`|xry|tGwd3;SFEVU$
zv^`%!duC(Szl>EI*G@|P+gSHn{|=fzotXZ|%}?zRG)wJ}XlsqxXNr>R*>c9r^(wi=
zTh5rdO-gRlmNRDVz>+(t<&2rTu;eakIb-JTE4lky&X~D(O77j3GiGiG0O@)@bm#f7
z=B6mQo-Jq0T(6Q_yyc8}hwEW|uXzXS<ExG5UN=K$?&fIc8ngEyC3k4c88dfj$z9fR
z#>_oja*wo}F>@c4+{Z0v%-ryN56xLdKzkNr=B6*X8CuSmx#df4g_bjBZrhUEuH}rG
zJEr80Z8>A+t}VIiTF#icr%LYWmNVAj7VFnZ?)(#c^~cS+{~L6k^INpfY0TW1{B<C?
zvC!s>nVYlZ=4v@(=GH8^wOY=Yx!p=`_m(qe?zECSz2%IVyQSoAZ8>A+UM{&;TF#ic
zUrO%RmNRB<LjF3H=bH%a^BFU@V971ia>mSUP;wi#oH4)e?2cWGKHmQxXuhO6F~9HZ
ziErMx)?No-t+D<<v}Z7;e@M$Ee`L$M*ZlF#H*c?tTfMnUTF#ic2TJb2mNRDVy^?#s
z<&2pdioX_T{?Pt)nOVlnO<8hNwVW|?OO)J_EoaQ!W+k_I%Na9waLFCga>mSETymGR
zoH26`l-z?YXUyDtCHH>I8S|cg$NGKAJM|s?fv+~6cQgz?Pvj1UMSBOv+&==={dot|
zmfGoBtufEBY{@Oxa>mSUU2@yBoH27pmE6%SXUyEyC3j8B88i1p$vxR}#>{<Da$mNb
zF>|Bx^G}|4G_=oa%-pOcH(SdYGq-BVt=4kJ%<WuqyR@7!@9=1>->32pj=@(O&pS8;
zoim+^_Dsgye+Jh5c?UO_+FM$!G0*W*$-Ufi#?1X(a=)~kF>@2}vt8y-h<3g)a|@K*
zf-Psv-1;TALCYC4w@=CK+j7Rtom+C}wVW|?cbD8fEoaQ!TP62)%NaAr-->oU!w_hn
z!Pr=Q7hDWG1=(I$wclG@vrrN}Jw5z6tru$^sP}uP_wL7^HRt938}sWSn!nxZ($1T&
z#}J*p-X(vT>3sN0Zl}f=-<rGX8Q-P(a_rjc3{dTS>Ret`S3Qs1rFopLz0UC83G-9u
zj)k8yk6rXQ=$zTP{yScN{8>5m;}<<a(GwOu5juT5ga00wAAhd4kFiOMo~&tSPG0mB
z==3o+70Zu5bK9pUcIu|hPs93|w=ykuLEe_{b5@q`Vjuk3vHtsH#(E4fR7d+x7v}F1
z+&?GFern^3@OG^=p4cy_>F}MS&0hZ<F>6zI9ed?x9+o-J_A6cWr)SN;vWME9GdACN
zdd`ckJr8Ic&QIK&fB#hX+T#bF)-x=~^6#hYZ@qo3oeAq)?OvRxCw3v0a~8&T@4_wi
zA=KzE!t!4H`!RdE>~Gy7ya)GNb5(oh#hP#a(iU4BZQQRX??K=Do0;W#w0Z6&(Dk|V
z?v`xvA$WIc-_f7!wZ3w)6uzIz$<p{?cp0=am&JyYRnV1_Rq>6<iP~8?-)d<4c{j%8
z#9HHtdDqL~J4c(ntD~(=T{u~Tm^shp{pc^xT7hK`wLMpCzVY;26J2|TleLJO3n$sD
z4JT`(eV_KX-oDOR3F};KI9UhnoN%%(+E_UG3)*{;lPvpNCnwfg6HeA^zWH#nKH7LV
z(f9sVW_ccMo_hmyeeQ6wVT;R&I-FQvIq{#t_ft9PjUP@nMmuv8Y&h8pT{+nr-<X`J
zD<|9F+b{20POLSan0LJ?zH_vBwr$bYrY@XpN6g&DEIHBNjI}w-9%_5`X}<CF+#X$f
zhLatLn+qq|s|_bRqUFT?*4x)PTVS254JSLHofA%WMjHz!yP&-n8R^@s{jHM|Ypn?<
zyEfl^IN1$tJe=rze_OIVk2cS}JGwr1IN77c<wWhf^ZSYQm6JX3{Y=TqUw8Iu+V`_J
z%b9=0hLeNPm6L<<jme4HSvlV!X#3gAxSUvPJTdQjAAIL%v-eQ6wW$jyhY>U9+2llj
zU)Fvsd#LTXfAfu}=i%tuGn^bj+*~-xUTruz5-lh8x8A<aIRNWiZ8$j!?VNCOG}>4=
zIR@>$$jPzI+TS`kvDTV!a$NJxhm+&c#>0ue_je%6^Jw$jC!p(dhm#XqTu#)!BRR3Y
za&i*BpUTO}`1yWv3fh^cV#CQf=*r2t_{QW!?W~;dJhc7tuI0p9<B56Kr{OzCo4x0w
ztxa7xxqz5C&n74Or?bvr*+Xs5Gn;QbJugJpp5f#o;^x9h_G-h)#b`OPzxDQY&RJOJ
zYQxDTXy=5JOVP%{$z^Eo#d-F0+21-jvDTV!a(VO3hm$MN#(S~62Yv7FY?kNI=DDv#
z*XIr=SGBmDsC{>GVtwW0YJ5MHlWXw9$+c)_UWW}QH=!#hH{%<V6ScE)zFW}t%e$5n
zYmF!7U0;vy9BuaAinca&;p8@A<~*DCqkjYIZ!CML?fLiS8&A*M(Y0qdxr4Yl-&@^l
zk8pA)+V^RH>+NgpKd{c#hLgL{&Iu=Xqm6}=d(hsC^ZcITvcGk5Vy!jd<lg3+4=4Ab
zjfWF`@9##I=h5c5??>0?4kr(^xSXiNiS?C}2l4$>P9DM!Cl8~Y`3N?gJcX{DJdJNm
zPSll?XYlQpcP%H@8c)o-eiYw1+C1B{Xlqj!PX0yAoM)30{l{33v+SX^=RccoJUyR7
z*Ph|zdE(~6N%m^P$qQ&XvA^~9b<Pu5=W4^ri)iPBlb6uO!pX~M@5On2S#n~3>*T~*
zYr@GZ%{L!TUPT)ZC;HyslPu4p&2ztouFoA#UT<+ZQHK-jD<^N@`>C9~i62hhLOb(q
zY&iJ{T{-y}-<X`JD<_}e+b{20POLSan0NgSzH_vBwolR4rY@X(M$DXNlN0@SS?{sz
zp|<Dy%{QK&pQCHfaPkFlbKxX=wc+GTw4B)Adiy%(1FUnk;p8i{bHd5jXk+2z8?^W0
zykl8%Vt?!8#9C{@$+yinA5OkQ8xJS?-rt8T&!f$Ae~+%u9Zr5|aXC@@?)>YH^_7z!
z@%>aze!>qYKck)b3pSkaE3AGv83Ny!oT#0Z^9_k_zr1TXvDSEE-u184I7gekLs8SI
z#n>~P3{BjeXOk2C-_XCa?4h>jAI&$Op2PIlGn@=d+*~-xUTruTj=Y@M-+KEx=TEG2
zwc%uV@}51MjDR*4PDVs~FV3^4%l_8MiM7^*laZQlKAen<Hr|W9-h;mP_ivWx(dM~F
zAy@A*oQ&Gya-t3=)>lsa?<Vn6IT@W;I2i-&%rUXy#D9)vIQC@Lg!snfL|r+V2;Y8T
zLQbqTo|tz%7QS<|dA5ns)}}6;OhU|@XY+pa$7YSgvWME9<2K)TdQOV2J;TXl#Lb11
z?A3;o$<e+~`&(~c=ZuGSt~Q)Zfp$(f>4`QLPNqbAFV6FOiOc@h$%(brgp;Y7Z$6w%
zjW*tkz21Yq_cuPv^Jw$j)1d2f=iN=);&P%6C)QU^ro;DBIhmeVIGF+M%o(xa#DCVR
zaxy2rF*#9JPUgb5U*5HxSZh2n?|LSD=V<e6bEB<IT{xMCm^shp{pioknuTQ#wLNET
zzVY;&7hQXXllh373n$sD4JY%X<;4Eh+t)d>VV$cDCkvpR6HXRH8w)23p}iO9<>wRo
zTPG*hS`$tdZoc_&vIyFEFM4<n`rhB{EYG9Ob1#an&mB${YjHVIhZE~7C%y3fR8AHr
z7EYExJ99~FIPssytDLNeZ%j_qm6Mh5?U#2gC)OHI%)4F+-#OYm+sbHbQx{HFA!g3A
zc|ZC~vzB4mLv7Dxn{PZlS4G#J;bb-9=E6z#YQxFuXgRUJ_4akna#-hT!^s+G=Y*3r
z(Z<5bT4?XZdHMOo{?^Hfwbq1_wVQ7~oUDU3-isdIgTD8-Jj?TF^W5vA>vM;bzqGiV
zsKbf%m6Hwd{j86bjrFh_HtqY_2<^<?*l@B1x^l85zA-sbS5CIVw_o11oLFl-G4Fa~
zeCKHMY+IwPO<g$IhL}0e=Kbhz!rGK&54AlvYrgUH+!kGXhLi1xn+qq|s|_dHqvgc@
z*4x)Pn`51;4JSLGeeU=>qMaLVc0wBqH#?)<<6nQA;c}jJa$~JE;bxcSn-4dA(Z+kR
z*L%_TKKrmdpEmcfE4to8xY@16<whNDtgqbcj_;>(vj?$!f7ui5%)PMT<^Xi%=0JR7
za-*)?9E5MbynDH^)_7vx{oeS_(dO9>Mq8V@aB~PTbDqun(f=!JAC^7T_T0Dm#?$jq
zbnO{#4kK<Z++?pd+#HVfecIo8`#NVotaG*D<OsCS9sfwQbHmM1Xk+2#XtaCsvy0Yw
z*2#^v)`Xj5nr}Yb9E&#Ii@n~9zW2F5%kyb-567YFJ%pR%TU>6`;l}#P%?bE^DmNz*
z3pXdBop~}g+?<83+?<VXOm5Vbn{)8(mv=8W)*4UDyFUfrIodqixoB%s7jDiYX3n#D
zKl-P#PGi|aZO_x2Z#+HEN7tU=<^tm8!cF#S!_9?gxv{_X_I1t~Sm$cP$wg?NJO0IJ
z=Z2e0(8j{erD*raOMaelo^^6#tu^81vgVr)H<zQ0_o9dQqVIj4$?|;K+`|><dJo~|
z$`+R!b-1y<a&r~FpUTbE#KO%rXlGuF4L3KUD>pac8<QJ#<>qF5`{musjkU%T^X{+1
zcaApCb_?3t)P<W{iJ9|k-jDwEtQ%PNP}}ow%{QK&x1np#aC19xbKxd?wc+LtwA|R=
zdiy%(?^x$*!^xd!pF93tXy=BTyV1tN%{^%M$W4B}ah`Q@W34se=HBL;4>$LrjrXF5
z_oDB8{)6TDw7G}-(e)m}%>ykiH|lU>edXpsd_R?&hlqunhtbY_1RHLiLRW5{#y2K6
z>dMVC`1Z@Ymm6!1C+6Khitik4p6yw*wW$j?{~~72vw1)IkFg$S*+Xs5e>UHEdOnA)
zJ;TlO#Lb19?A3;w7tnHJf9vh*oF}l()rON7(LQ(lm(b1)H!q`&g_~E<?s2{|T+Xvj
zZmhK?+`QU+^Wo+-wDDf-^<MP7&nH=)Pn&yq9bNAs+`Q4^a-$A6)>m%c#P?IVd5c)M
zc^mD_cd+5+V|3-_6MSQGqpsY1if_NXd%3aJcw*lDyZFx0=Gi_&TbsIY^Eokdp3VEw
ze~<M(%N}Zbe$ag5>G=h^_6#>)5;qravR4~!zC!yx?Qgw(o%12qx!Q2@HQMKn{|(x?
z;pSVkv2gPp+C9$C&p6JrPHwEVCft1AeDmSv2ek2C^zdHvz0Z$Wo==;5_z_+2A>91b
z;&P)7H`Z5fe#ZAxx%q`yxcL?B%-^u##(&qXax)aZF}YD!ZidFUpK0TAW3BPTy!+qV
z9BrO$7;8#hxEYq1InU<(=>LKKv-MEh^WV)ko}R<m&(+QoZiXjdF5G0VHr$LrUT*Af
zy?veIAMeON;bcU#&mDgx)2whaGTK<U83pYg=jZ1e=UFE=)>;#8Ms2?Na5Ea(crSW*
zFZ$l+5LnNr%{`2cuJ;gb#%OW5QHLArD>wdk1N*7mj72QmjE#2YIM{IGzaLq-nFQaM
z+^8!zlj7UYbhxqBcw*lDxH!(y=Gi7gTg#u$-f%NH0dt<s`_UhdH9pH8YI{!5eB<dk
z1-kYOH$91)3pd%T4L4Jw<;MQj+t)c0Vx6lECsUz)?)X!qof~eZK^qG<)1uwu{QP|5
zJnQ7fT5H11bj>#(Zl*^Y4>$VW=R_>er_DXgfUfrtZf0z8xlxB3>nk_@cl!CM+{{cY
z+{}V@=B(Io<G)i|xtSN=nB1r<H}m1!&vdx4)_7vx{cJeS(dOCaM_bFE&fajd00DEJ
zO>XpOXU)N~huWTVHs5%9E{Lu@!_7j(&4run)rOme(Q;#d>+S2Dxv<XFhLc6mK6m^@
z(asGwi=mB$n_g)5INupA=UFE=)>;#87H_`!aI*xu6KAjYqVIjq&GLNO+{2RSdJo}d
zsTP+Tb-1y<a<eqPpUTZL#KO(8XlE{m4L7TzD>tj*8<QJ#<z{t!`<V_m)*4UDyI&s1
zIodqi8fa_z)7cwt)+Augvw1)IE3j5%*+Xs5m6~rnJ=a3lp5bO~;^x9l_G-h;I%wag
z{jImJb5_PWR~t^&Mf=?G|AKaIxLFTvEZnS*c8~M(^NsVYlN)QT2{#)w-+Z{)5N$l)
zZ}h#-Ral-+n|s&@UGE{>^lovvQHLArD>obC`>EV)LM+^DigxB^*l@EAx^lBEzA?E`
zS8le$x1Z^7W3BPTy!*{@oTJUNZI8B=Kb^hdW(NZ1Je%C;_hD_pvWME9TQ=W#dhUp>
zJ;TjT#Lb19?A3;wozZe*f9vh*oUO3V)rON@&^~wkzG&x$n_bby!p&}I_c%X4-#E`Y
zxv|!oaI<^!&4-&k(8j}!zW2E`%kyb-4|}5PJ%pRRT3l|_;l}#P&EEKaDmQ;67H;-I
zJ9A%bxH$w}xj7WynB1r<H;3Wd&vdx4)_7vx{eC#k(dO9>M_bFE&faiy1OaoNO>Xq}
zXC1(@huWS8Hs5%99*M3!!_85|&4run)rOm+(Q;#d>+S2DgRsukhLdB^K6m_M(asGw
z$Dxgdo8!^$aejWjah`Q@W34se=7i>(4>u>GjfWe3@AF`m=hNmMPD0mv2sbCUxZJ42
zjrEnA)A0SAij|d9u%|a|>J0v$Gtb0^n+wpDn+x%c$&I>la}mD%Ootn5jVI>apM~Qb
zZJzC7w6*-{><u@U5HRQ2<VOE&);TPDsO@=f^NpwHrRdr-++0T7T)4?zZMeA{EjRYJ
z-oDN`59?fQIJpAtbH~3D?c8v471~(1xf<;r{~F~Cm-DQX8*8lzH`m0+hMQ~A#>0)i
z_jx|c^J#Mr*P%TRe>!&?Zmw@}xlxB3>nk@mxRFog=5NIE{pRmzXZ`~lZf-|cZtlQ0
zCO7KJ&7Ju6GaYWMHJ+Gve<O}_w0X9>(AM&&vp3w_O~9OIlN<e;SU0omp|<BO%{QK&
z_n>RfaC0wlbKxd?wc+MIwC~gY*4x)Pw_=^E4JY@beeU=Vpq(3T9z+`pHxHrRlb>z0
z&a+N#thFZGJluTq;pP#v@o=N>ecs0MeA?W@qv(1M;pVXxmm77svA%NiIKH3C%|D5S
zn<voDd=eXOo<~=1UcfgdH|omGi}?049d4{Oo|t$46pnMWdA66(*7B#bH{85Tz?^53
z8~vwQ&#>&Fw&%0WH=drapli=?^D1$3;U;^v;pR28+}PiG`#R@eSm$cP$?IsJJN_GJ
z=Z2d%(Z<5fTWI&lOMbp_o^^6#tu^81?dF>gH}9a0hZ}wG^EsC1)8-!DMb~=>H}AE$
z+^EBi^_83V@%>b8J|GruK14h7BW$?&5?#6Z3g4LAs4F*L<J-@4xUtrFV&477IL^`L
z*}g$r%b(8PaPutzbDm9Z^gm&J%Cd*ro}V?}czS+^u06xe_r%SGo9xwwn;+0}V}I-I
z>zvQA&eevKAJIN{{GZUy4L3idjfI<E(C(3&{Cwj)>*U5-Yr@U1%{L!zenT4%H~QY^
z7c9@G%{}~%uJ;gb{%CQzQHLArD>r}Q`>EXgn^?Hv$eo$A>Eg!!&a=wR2>8b2MqRlX
z5#N4}2{+anPt3a?65nUj=GjI<TbsIYGcqxAo=tA_hhh!QvWME9!!+M`dX9pwJ;TkY
z#Lb19?A3;w(a>^Zf9vh*oMExf)rOPN(LQ(lG4P!mZpK6#3pZn--Q#>`xSVI5+*oT(
zxEZ_o=EKc6Xyf5V-}@Yn<@vO^hjG#M9>UFdEiO0eaASSt#{VrxKb4yah=rR8(axL*
zo9{P0(UqGi@r}uix^goWzWws<{m&V!HJ+GvKQX>@w0X9v(blFe+)P8voM)38{YhAp
zvh1O@=VZ+{o}Sa9YtL{q9dUEvCVREvW_q;m)Be`m*Ey48ovRHeGoXF$_%ouN8*XMo
z8w)oxqut~D{EXu~>*U5-Yr@Se%{L!zW<?thH~QY^6fDoD%{|P9uJ;gbW^Zx1QHLAr
zD>wdcVEL)s%t<WV%!PL5+}LpAf45@gW?_6|a-*)?EP`*pynDH^)_7vx{XF>2(dO9}
zMO&M?aI+XObDm9Z^yg*G$Fhgop7S@~czX6i*Ph{KapLB}P4;TT%@Sz2vA^~9b<P4<
z=W4^rl4zeh{!(b?hMT3)#=^}qX!kfjKi@ddI=QjdnsBph^Ua5w<<Q2%jlTD}Aj|V<
za}Ud->pg^<6<S<w)ZxbZ%FT-SekwOB5eqjfqn)`5Hr%X@uH3AHZ%l5~m78_(?U#2i
zH`W?Y%)4I|-#OYm+h5SurY_v9N6egOlN<fjSgW(_p|<B5%{QK&>!WMWaI*n%bKxd?
zwc%z%wA|R=diy$OO{{aZ;bbGU&mF%v+PUFoW3;hwvkBTg&d<*`&a+N#thFZGY}$PD
z;bt?m@o=N>eXhmweA?W@=IDA4;igZE%Z)nRSYNr>0^d*NW=mq>W-GKax5kE>ozRt=
zo$-yyjk<EP3%>pG?&Zc><B56q+u%D#n`i5bwl;O)W>;e7Je%C;Z_C<_We>GIw{O1j
z^xO?ydxo3ciJJ>I*{cmVd!XgU{?^;qIXhsTs|_c6qJ8f8d!d~hZuU0E3O9d6yT|#?
za5>L9xv|!oaI;VI&4-(P(Z<7#zW2E!%kyb-5Bs6(J%pS6TU>6`;l}#P%>np+DmMob
z3pWR$op~@e+#H3j+#HQ>Om5Vbn`7|pmv=8W)*4UDyFUcqIodqiv1n^k7jBLtX3n$8
zjsBsm!&vrE+w<_|8&A*U(Y0r|If1yjaFe~-aC0Ks_i2CY?dzN)u+G(nlatUscl?vl
z&J8!GppAu_Q_=2mety1jo^^6#tu^81wC0-+H>abGhZ}wG^GKHG)8-z|K-YT+H)poE
z+^EBi^_82m@cmS7&L$Ra&Otl#Tx__x1YNng6yKQKs4F*@;oC3oUT&;4o|t!k9=>z5
zdA7^Z)}}7pTtUp7XOkQK^H~?L?4h>jh0Qmfo>!u4&v0`UadY7&d$r-_YP8(g-+KEx
z=OV0gwc+F%w9g&?TC{V+&2?yF;pTd@dz_!2Z=7eH+*oT(xVfSE=EKe3(8j}!zV~@C
z%kyb-4}VA3dk8oGXmPnwha2lFH#g$@sodN|EZp3TcIGYEaC0}ha&r&9F}YD!Ztlgm
zU*5gkSZh2n@BUVN=V<e6_o1y#UAVcQm^sfTH~P1+ZfDs;ZO=QJZ#+F8K-Zq(=0W1-
z!cF#S!_7lzxv{_X_I1vkSm$cP$-`)$JN_eR=Z2d{(Z<5fV`%p{KR@3%&pNrW)|znh
zc=OGNn}4E>hZ}wG^DdU>)8-zYK-YT+H&3>>+^EBi^_82a@cm55l9is=r-_A|XVA`k
z78`C}L04{G#WyB5>dMV)`1Z@Ymm6!1C+6M%3*R}~JlpGNYg3mydxMxc-=W;-KhJuO
zWe;t7yos(ou3`S%?0bv2IiJrS`u2E%We>IIda?P&)AMa~?RgMAJ>xsX&4s<})rP%y
z(ayBL_4aknOIYV>!`gdj=Y+NQ(Z<5s2Waoszpi-~F8f<2Yt~wmd;75Y=EK@YXyai`
z-*@mb%kyaS+#jRsbBDE0T3puDVa@u=+Nbz_reuY+&xnP!&(Y5O0vpzTKv&j&#5X2u
z>dM+r`1T7gvSzLE#JuY-@tvd1v;B;=Hg&nPU#w+?HT|zyUm2%bn;yTSYmcz@8*y_!
zpRDQI;~SPe)Sm0x<{MAX-_f;aSo?#xxv-YK+OYN~+L`vZ-oDQH4(nWPSo=5HIbqF<
zH5S%}!1r$RGnCf;*2$W+*5uxXMC+RmYeS)phc$iQ!S^iBqs?;<O|Cw7SR1CrWlhbW
z&VO%(^_8_@P4JnL71o9$7S@JGJ97kVSo42lqp~(8zA;%-D|5cFaO@XeWX)RRiFwx}
z;yXv1y<?-TO<nG69Af6en*PYFky!T7rpLJG+9RxuN8FsxCu{ol7=>jIwdWeO`Nq?8
ze01#@)+QisF05s*HmpsEcBcKUx36<X!#Y<R)+RzbC#+43HWt<<L3_6{mY==sZ=I}J
zYfbKL(&n2FYm=dkhc$iQ!RRc{qs?<qj;_xg)~0B2SyS_;!<zM#wVo#URMw^>7S^Uh
zJ9BDmSo43&rLr~)zA;%-D|5bCaqJgfWX)RRiFw!4;5$c~y|baMO<nG6c4Fqjn*MaG
zX<7ErrpFxU+9RyZN!*;zCu{oln4V=1wdb0l`Nq?8E_Ce~*5)Q|F05s*HmuEqcBcKU
zx36<%#5z|S*5*Y!C#=nfHWt?AM|-!jmY>(`Z=I}JYfbKLf##bJYYU={hc$iQ!AvaA
zqs?<Kgs#sW))sDYSyS_;!<zM#HUBr{{Z!T#B^K5eLp!q<Hmog&uB<JOZ%o$I%A9Wn
z9Q%bAS+mx7V&3)Q_|DO0?}}(^Q<pniiI};troSX>36?#y>9I1p_6TdM5I5)Z$(p`B
zmSWjM?YWk2zVY;26<vFVwbh853v1b{4Qs2TooRpT?dzOnu+G(nwKdSr32SSjjfJ(f
z(B7@{ybG88t&=rtt;xNu-F)+5Z5_1nu%_=jSeE5^w0Z7z(e=5*+Fx2+*3|szux5Q_
zZ9NlwDr@T#3u_yoow*@4to1=x*0#VmCTnVC&bK9w{lbf^S!+Bo?|LJA=V-HcE3~z#
z%bjga%v@O0-<Z{#We;t7Y=f>n!rHdP&G~$?rf-i;SoTnRu1%Y7JUzEV*PdZ*d*bH8
zTJ~zg+74)E+TVKnI%hMibG2b@N3?Uo+D>R=VQpu$ck8_T%w~V<WX)P@a&Nmd-+Wl>
zi#8tC^nC}LvpkPB&%G<UK6hB#t;J<c&7Tfy)>qbcH^HZ}wg<7WwkO(|dtt-cf#}NG
zLHNdGO|8uN4#u%xc#$=0jVI<^?~U&qZT235wl;OSvqOoQ3v2rOu>Q)jhc-P9L)RW*
z?Qr7ed_Gyzx5vILd#F9ve$6+Yo=2c-&#-nRadTlUd$nQhD6})}Z@qn;vp?3k+OT#s
z+BsqE7__mlb}ZVvbzXj6v%ht+X00{3x8s^`KCB&&HXhdWeFq1yJdZZdeFC~ZcUU{I
z#br&+pAKu*SJqB4!KbozGO@6B3fh^cV#C_G=*rr8_{L;St<3q($FW~{ku__LC+1zB
zhVLA0_FjOtHg&nP3yGNvYx-xfPG{Len;sXTYmcyYF>!M~pRDQI<4l%4)Sm0C<{MAX
zOVG7vSi6+Cxv-YK+OT#R+L`vZ-oDN`8|z$cSi2nUoUnEU+E`e-67AhOFF&u@-#S^d
z)|%YgRn0da)~+_s3TyhlgL7D(N1Nxq23?;!tX<pUvZm%whc)XfYuB0JQ(3#7SXjFO
z?aaSn!`f}=%G&Mt#$-*c%=zxXv0r$RHEWG0=3W0C-#Oasy%TM1>T+jy5i=Lo^lxPS
zgJlnGdfbh!J;K^O#LfA9vZim3n^^Wxd#;<CZ#+HkMc1BT?LOk>!dmuf!`l65XWHL-
z`#R?qtaG(t?E$oN!rFsqV`1$fw0G+~@4{t&>txMZYjST7H{X0%djxGftm*p>Ze@8M
zZJzs4bbaoy_E?L{nwmcy)~v6rJ#K<eW$mBD!rBvPXFiDyYcHTHYcJv(lQp$6=X(jq
ze&I#dtTmpPcl{K;bF|s}GTPeI<<4FqW-hGhKf`*OWe;t7yo#<p!rE)Z&G~$?rf-jD
zS@uwSu75S(czV8$u06xr8^q0pwd~b~wKvhuw7>QCb<T5G=W4^+TWIHmwYSm6!rD7%
z@78(wdCmUT$(ps+<lf$GzWK2B9@=<V)At=b&+<IlJoo$P`rKjdgBF)HHGev+SzlTE
z&;*~#+DF8~+Q(>Teu52a-=HgN-{KpSHMKJ5`wqu`;YHT0HJ+Gv{VBe4wAuST+S=6R
z&VC?fF0AQ)&iagH4{dt<h^{@t+E2vI`FygbZ;vln_E3ASFPm>XJ%2{mo?-15Ygu6}
zd$nQhSF|%b{jv6S&R1CHYQx%ZXy=5r-_gdx+8=1|)_M7P&HmQOnzh#C-u`U9`LOnH
zwDGW}?>qRK<$1JuZeCqK&mGo=XmMFnv!%nD^_8_D@%>cRh9VZ$hDJMc7;ISce+#*?
zHY&a`SyMYJ=Nk>*e&I#dtTmpPcReh=bF|qzI@;RQ<<7<+W-hGh56>EoWe;t7jESy2
z!rEBG&G~$?rf-iCSoTnRt`VDWJUz!o*PdZ*9OCA}TJ~zg+PG+E+TVKnI%g!TbG2b@
zJhXGd+W2T=VQm7mck8_Tyk>vvWX)P@a&Hqh-+Wk`2yHy9>H7{wW_ccMo_k_+eeSR}
zNsG&xI;>e=S(_B!PyP2iCnFZtCPzDS3T#;Oe{;LCHUqvfSyNZmX2iE&c#$=0jVI<^
z_r!OOHqSN_+S=6R&SoZNF0AQK#hQ|34{dtPg04Nn+N{LQ`FygbZ;z>2_E3ASX_{|5
zJ!eDLo?&fv;^x9y_G-i09B60S-+KExXIiXtwP9^ew09SOF0^yQ-rQ(oVQ(I^d;IIE
zcjI!Nb+TuzHMz%mn{Phs&4)G~_Vj%x)3H3CHuo?;y52+BTcE{dPaXEGuk0;|@29f2
z5V5egFxr`mV8h<h=*r$Q_{L;UUD;a}-+tjo_N+CYn0LP@zH_vBw&l>)rY?83JTY@&
zPrnyyF_t~F>9GR3_6U0`5;y1b$)3JF7H8Q*?YWj{zVY;230-@Jy_JcZ3wznC4STDg
zooRpT?dzN+vCh?owN=sHUHsM1&JBC3qm6~VHPG(K&vIJlStoneT9bQRv-#%3-dbqm
zVNc(8vJ}hnX>$*2qw76{y>(h#_S9j|`pVwA_<kyTe<2q3)<ZjUeQem<3|-mV9N(Di
zsVjSZ@a-3VWY1dTiFx-M;5$c~XWIg8ZR&DoTM{!D_VhPmZOF2RHa)gN*B)VSYvSg7
zKH1Z^M{kxr)SheO<{MAXZP2x6*xQ!4xv-bL+OW4B+L`vZ-oDP+1nXRFSlb@$-NoMl
z?cA`pBidNl+X?L+naj_2&a+PTthFZhxO4N(hrM0U#>1Yz?_^V!=hNmM`l9PSguPu`
zT=vvq&-%*VZuovGd%F`0dwZarxhFR49e}Rv9f)sC_SBWVgYfMaeq_&D<B56qd*M4r
zn`b*1ZEfmuXNM3o7xwi3%G#S{4{drJimpAv-eJVe`FygcZ;yRg_E3ASeVcDQJr76M
zo?-6@;^x9$_G-i4k!WYy-+KExXFsfSwPEcjw09T(XtZ;~-Z5xnVeeS9dt@&^?>Wyp
z*|XM~+~aZ0Hy`$nM;i}&`o5F>S)NaudpH4I?;-4+*y6IM4tv&D_D;h0Q`tM2SlBxS
z?aWiLVeedYW$!$EW3s2N?46Hqzwje_)*4UDyFU%zIodqi1!!wimpi+Vn7Od0e+KJx
zmOZrTaS^)q2zwV3H|O)op1wWKWZ6UQxz1|7@$|d|U3-STONpBcd)ccEdzYb|X@BeO
z>zuQ(&eevs%hBFl{43DT4SQFjjfK6d(C%@*cjI!Nb+TuzHMz&Dn{PhsU4u3r_Vj%x
z=de7VHurEXy52+ByROA$PaXEGuk2lq@29eN1F^98H?%YVjtzUaqbqxN;2V=Yb!G2P
zeEWqT*|XMoV&45f@SUU0v)zTZHg&nPyNQ_#d-^xAZe-a*n;!R|YmcyZFL84|pX}+|
z<7SpU)Sl~><{MAX`_Q##*t?&&xv-bL+OYQk+L`vZ-oDPc73*AWSbGrd-Nk<h?cA{U
zFxptydj#zs=jUfW=UFFv)>@N$e6;!I!`@?P<6%$VcXAuc^J#MrkE827guQ>Zxa_IJ
zp7oWzC-D7L_MRjb_MSpJ^J#3@dkJ0Hdl}!D?5Qhzui)D+{K%fQ#uM}IpTT#IHqZ7d
z+S=6R&R!#CF6`<5i}fta9@_MH9bJ2by*G%P^Z8^?-yYAg?4kBt&o|$AdcKLSJ;UBx
z#Lb1h?A3<7x6#hDzxDQY&I?%QYQx$)XzwomyJ+Wzz4y?@!ruF6_c%X4?>Wyp*|XM~
z+~Wt$Hy`#sL>mu#`o5DFS)Naud-w=l?;-4c+~TsQ4tv&D_CCS)Q`!5JSlIgv?aa@y
zVefl%W$y=kW3s2N?EQ#uzwje_)*4UDyZ-{;IodqiPiSjXmpl8Jn7Od0{}t;?mOZrT
z@e8{42z$R0H|O)op1wW4X4ymSxxQ(>@$~!+U3-ST--(+Gd)ccEdw-yvX@BeO>zr?~
z&eevsKhfS@{C}gJ8}@n(%Rd*4ANGd8caQV)^PcmplRay#$vqB<);AyahC&+;d-}eU
z?^vEsn|m0VT)l^|H%yDmo;vJVU)dWL-%n+4IAURMc(gM|z=pjs(3QP0@r}uzy0SNx
zwXE<Xd)69H%)1{E-#OYm+t_GpQ<pm%hnTsrr#~`lB$hq2=`k+4_6U395jW@a$)3JF
zMq$}Q?YTy6zVY-NA6<Kfy$Oh$3wznC4SN%!ooRpT?dzP;u+G(nwTaN)UHpmB&JBB$
zppAvSNzv|czIWqto^`Tktu?vF$(nCI>`jg~9`^KoC!@1GpEma}1-jlt*z4KivZoGv
z)>rnX#P?I#n~GT2n;PxRX|Q2$7IbBAR(xZ!r>^YHhHt;{BYV~wPt3cY7T-DAJlpJO
zYg3myn}e9Su%|ygYdV%awCOP?y7mZra}hV^^U0pRJ!W9pL+!a{Y`*dIoEu$xhP`=+
zn+to{s||beqMd1f>+S2DnXt~)hPC<7-d+6p(asHf3!sgKy#>+kaejW@bDnjwXRS54
z$Ay}2KI|=wHXio$eJ3-sJfAlAun4-|L)cri#br+&_N=e$Er#!>ve%1P*jpU!%q6g4
zZv}K^Z$*4#vZt=>t%PsC@FRQH8c)o-UlQLr+C1CJXlqlKJ6naAxv;0dG;1lAJ+$eu
zD!TRvd#e#Q=kv*)zCD&<*+cEQmTkWA^jsZXdxpI=h?@(0*{cnEYoeWLf9vh*oaM02
z)rPgT(B57Awb9NEd+VT$g}rsr?s0y8-gBOHvS+O|xyQdW-+b6x4{bc`>HAKWXL&wt
z?qPj&y@#;3L5s_tI_z0r+1n7`Pi1c-Vqvd0+L;?;!`@ct%HG!a#$-=j+1m!+e&I*<
ztTmpPcfSd~bF_K3ZPC`IE_b#aF>_&0e>2vmEPH6vV|#S%5%zW<ZqDbEJ$-v@&a#Kv
zbM<My@$}pgU3-STors$Yd)ccEdpo0@X@BeO>zpmH&eevsUC`cL{Jv=ChP_?U#=_oi
zX!kfjKkqruI@z<<n%v{=%{L$R_COmCd-}eUEm@vVn|s(3UGE|6?bYJ4rw)78SN8VC
z_f!A9=)V#Rd;6fBxi2>C9g42(9fogA_SBWV!}0AGeq_&D<B56q`{6rBn`b)$ZEfmu
zXGan<7xwfIVC~Pchc-QqLf0N)?`Y!Yd_LLJx5t4jd#F9vLCrUwp2whT&#-qaadTlW
zd$nQjIJ7hEZ@qn;b1>Gq+OT#!+PjN?0@}G@??kk*uy+#LJ^uC9yKy<sI@z<<n%v{b
z%{L$RPC*+Fd-}eULs*_qn|nAFUGE|6oz~*Arw)78SN2ZF_fy$BgIL%*6Yb2iuwm~a
zbY<^id}FewuIyccZ@=&(d)69H%)37u-#OYm+ofo0Q<potjF`Ewr++T%9F{$_>2W!_
z_6U1d5I5)Z$)3JF&STj_?YYixzVY<D5?y<Sy{m|u3wznC4SQFkooRpT?dzNiu+G(n
zwQJDcUHog&&JBClp^b&T>(TDX&w5(tStoneT9bRcq50;+-rvy1!=AqH<U*F`)8-!j
zj;{9*_Wsf0vZoGv)>rm!#P?I#yNOuXyBY1wTd-m89&}~zUVLM+r>^YXhi|{|BYV~w
zPt3c&72i49Jlp+fYg3mydw`g^u%~}J>o%4>wCV96y7mZr4-q%#^U0pRJ?>!HL+!cl
zY`*dId>CDOhP_9Kn+to{s||aPqMd1f>+S2DyRgpHhPB7g-d+61(asHf|3n)LdrzR<
zBXjwA&w19#p0(EG9-nN!`LOpC+IZO0_nq9$@_gFd!_(+`4`J__7MDGB*t5Q}_bk4j
z%HF?-g}vv{&U_vl_FhL<_TIoZCVT42-kbRL3qP`Ft?|UX`xo$?qs_Cug|;?zxwE&4
znG1XRFR@-^*+ZKi@1Sdsu=g%;b3UK!>D%LFmOa#->y_pkPtW(zwP)CSpSZcOm%ZAs
z_W|0O_P5@?&UqE<Ty0qU5bfQ?{|N2eu=g?ASlIhS&62(RyyraYWY1b_a*v-j-+b8n
z3~fB@>HAJzV|hMp?%{KEy@#;(MT^UxI_z0r+4~aTPi5~bVqx!Vv@^fKhP_|VmAzl_
zjme(6viBRl{lbszS!+Bo@BUkS=V<e6zoV^9UGD4;V&=l0{`ai!SoYAS$DiohBkcW~
zxH+Fs_Vn%X1Ir$2&-G*Tji;v%UwekVA-oh;*vnpR*c+0(GwpA^eVy|Y*16iSHWYd9
zF8<K?&JBCR;2R5j!=l~eeDB8PJnLl7T5EET!!_T0*c%>gJnZRr-c38-&z<?ha#H>|
zS90|p!rq82E_>>*XMJUFBz!-Wy^)E9y;0EK$*9<{H!hk#Ju+)N9AmPluI!DEZ@=&(
zd)69H%)1{A-#Oa6;|b8#rY?6jAu)4dPk#*7=q!6^(_<oZ?Gg4SCT`B>lRbTVjLEWx
z+H;N7eB<dk3A*+Sdy^737xuDO8}=qcJJbHw+t)c`W1XuFYm=kByZBR}og4OgqK$>U
zDben6etzb2o^`Tktu?vFshV#->`jd}9`^KoC*!a@pEma}4Z7Y#c$l`uWltUUtgq}%
zhwrDdH$AbiHv`(4Gh)NuoM`@Zu{RfvG1*gB_U6X7U-*$dYmF!7-Oq&Y9BrO$9<;To
z%bm?j%v{*hpM^Cu%O2YFm=9fhguVHRoAdc(Pv0K1vh1PuT(dRbczP~?u06xvg2c^*
zz3kP7y@k-uw7>QCb<XTq=W4^+!f5X<{vv4ShP_46#=_oW=2*_p&wI}6)WM#$*5n?0
zHQ#*LTO4gX?CJYX=3sd~ZSG+SbiIeLw`7aUo;vJVU)fs<-%n+4X<}h-8MHH(#fH6A
z(fsLRZ#5iavZt=>t&VTM@FRQH8c)o-Uk=|n+C1ADXlqlKJ6n^Oxv;0d0&97eJ+$eu
z7P|HbdutOn=kv*)zCBiC*+cEQR%*WS^jrsBdxpJriJJ?1*{cnEe?dFb{?^;qIV)qG
zs|{=Gp}o8K>!Y0;_BKEp3ws-)-Q)cHyyraY{Jv+cHMz%)nr}Yr^+p>Hd-}eURal-+
zn|s(8UGE|6ZPMbhrw)78SN1l=_fy&1j9A#)9PP|L*s!-fnm=9a?SNxU_SBWV9r5iK
zeq_&D<B56qTi`oKn`hezZEfmuXFC%!7xwhGVr|K?hc-QSLDwE(uP<?PKA-IA+hc2%
zJ=C6So8}u&&t1{AXV}|~xVf;Gz1pz1JKCA{x8A<a*%s?uZCKj_?cK%S6Ybovw-?%2
z*xMWJ9_M>EF6UV%d)8W$d;Dwj&4;~x(8j}_zVBo^mgm#v9`;4odkA~`wYco5!=Ck(
zz5Vh1RQ3)a7WNKAJM$oH*gFc%pDy-}#xW*)>dM|R`1T7wvS+RF#Ju~1@tvd1vmJ}J
zHg&nP<A|9Hd-{j64q@3tn;yrbYmcyZ0&#ObpX}+|<1m&z)Sm0`<{MAX6VbJ2*gJ{1
zxv-bL+OT&r+L`vZ-oDN`0_$9DSUUyn-Nio@?cA_;8roRcJ00yF=jZ1==UFFv)>@N$
zJfr#M!`_)_<6%$VcXA}l^J#MrXQAsoguSy{T=vvq&-%*VIrx4md*>1hd*`8@c|JDm
zU5@5Y7kgLW7?VA9W$#LS`-LCbv(|WG-u(sm&e7)Cu0mUzy4=~-#LR^~{fk%^vh1Nv
zk89AiN7%cTxH+Fs_Vn#>G0Prm&vi-jji={z=-M;vT~FLx*vnpR*t-GkO#54JU*}wk
zb*?t7{SEEi#s53nxnb`gXk%gTMznjJpP%=fXPxX>YfbL)rskUudpDzvhdq7Y$z?3h
zr_DXwg0A-v_HJ!)*;9u->nnS=;rprV-A*j*-GO%Io!GGVAeuj2>^+2IO!m~3y@&Db
z7k*^VTH}d%_jlntN1JDR1Z{2Va%YbcGZ*&s?_u4|vWGT39z)k2VefI`=6pWc)3?XH
zEPJRu*L}@5o}T|i*PdbT3F79$UiNCk-jir&+TVKnI_G|@bG2dZDYSPN|7o;y!`?G!
zV`1-Ew0oSNpZA<+o$Og_P44ku%{L$Ro<kcCd-}eU2Uwm@n|pX3UGE|6z0l&arw)78
zhrJikudw{QjFq>Su&*|4{xy~}U&qScJJ@$w{OR)FYreC45`Q1x9y!YgX!ppb+Os70
zA-?;~sl%<aKVsP@?78P-{7+c+lVi1YsqsC&f$y2M+4CvdJ>HwWKVuooz60rH{PULc
zxwH2Rv}bgl+M4Wh-kbQ&)4FFC9_LH6^TN$nXk*ELjkagr<2Pty=^MVjCFb7n^&Q&#
z^{=UZ?YZ1*oqSnqP5Ao0`R2pd4`}1zOW(VCi{*1@^XxyO>%IMiZ!G7uZg$S_Gclh>
z-qiIS|AKE1bLy~Yugc=D_<kyjzY&v5-_7qV&-Vvb7Ka?3-)FLLyZoV>@40#s9|qqZ
zIp?r#kG!etIfv6{nd@j|cp09!ed=?|;t0e&gKVg+OO0&&NsVXL=1xa!d(wL(d}G-c
z7Dpy#&gahFQSd#Z^VHU4pY#6R=4ss{i=$HG8N=deXk*Ebj<#o590P4EeZ%6I#M~Pe
z$3lC*`T0)kUh8DhT5H1M*v&T|7RNyw4~zQVRS&H9rp>dDi}pEkZ{wkj<($^Z;`qdT
z9(n6%V#!TF+#crCVbNZd#R>8KR2C;978WN)d%j7qve*+lCCj77pQ`!J?n!)Ve0$_9
z)1ciWZ|ch8wD|5fr#6@8o(|tW^|@tndSdpI4YhTtk&Q|5J+sy`xMv2md(wMGw6W|9
zi!%{3=W}Q8%xKT(Jhe61=e)`Aou_q=EY5;<URazJZ7liO(Dn?Av!jirZ&;jzn0v$G
zoM`V?2J<tYd##g2Ypn^3b2Z<5SezSeJS^&aSCg}R4sD)&9(298dGU?qoYu+We8hYn
zc~jSSJU_lY%&Ehoy()_f;QOg8E=VjaE`;`c3u9$(aqJSTu&BRe^F3EjdM$--kDPO9
zw0nAR7Ii)6GWhN{r>-n6i*KL$+;X)XG5g7e+Pc)p#v=HhS(`gu9_^m6xB}W(_Jze2
ziJ9}cvv(!5XLO$0n(TAlqWI3!x<?jQMmsMou7Wm}{HkbshQ-y;#?m(|u1?IoVQ~$#
z_bZF}8PL7f$)dH^gvB+RZ$2!pg*F}*^}VabSU!g~&%QRg-rG9(#&S;UWN}?$K99Vq
z>pT7nzCFyT!=k+^i|gV0sVuHfEG%w-_Iw*+WpPvNW~{KNzj^aLS5JEN!M8`wxdqxi
z@}{om+!EjY=G2wNt?=zrpIa8UCT2g`P+ON8+1Lo*Gi!6F+o0VO7Pmzk%f7I<9Wir0
zclK_N_KeO`Ta$gx>y7U`t$So~2ek9T;*Mxz$?t@=XIR`BZ7hAm;x5G88y5Sby<h)+
zBtI{@*E(6W)|#-mYxB*A#of@x!=k=-wK2=*(B|29N7sAX1K(KAX`L+YNzCVwH+6l-
zd*R!|oH{JptFpK^zMsnCUx|gqebAn7U#u)1gguxQ7WEHlzUS&muS4<ek#in~c8|QN
z>p2g{cfUDxW$_4n`_$)_#UqK?Pd3!nrA9XP!}rYE-04wh_k_ix(Z;ebEFMG5oX?%T
z$D%!>^VHU4pY!&|cb?WgvUnWYd13K*w6Ww*K-)7co`^P<zG3ksV(tx#C!@VzS<KIj
z?zK)9t+ggBp3;2tVewS7@vx}xT^+#kIkb89)6n(aPRBQvb6O{hXAtvw<V{`Q@tOGc
zFsBZS_NpwNh3}`bcs8-Hcn;e0or{&li?A27!lM2q&G%eA>2)c-J#x;=(C(2pbv@_h
z`0h8St}I@GZ=d?yvUnvi`^kpdy41+VdH9}Ln>)P<?VhlBHQHGAg~e-#ne(}`_gb`P
zbe`Ip>~r4v_|DV1M;5O`J1;C=k2aS44QP9Y#lNAArEggLJ2Cf$#ebl^Us=r0i|(~f
z7Ok}=EZ*3B^I`EOwDGX0?_FKM@;S75_M6f5-fqD+mUCJsi?<T<dE`x9-|=nu_AsXo
zi}tE4-j465vUmrvuy`lh^WBA&#rv@ju)?DLgU$C`J?ZrjzCCizhtcklH+4PdBlzw&
zr>-nMif^C#+_Lx>G5g7e+Pc)p#@+ayS(`h39POU4_)oO4><f!e5Hsg<XYZ3}&*(h0
zHQDF9d+?p7b&o7Qg?3(8d>U;m`Df7f42#dAjiql`{1-9zhQ;U5-f#V#p?j^9MQg1I
zi_bUTd{}$|Z9FXMdsp|ed=71%{Y7-Wx0mpZ<($^Z;>*N*9(hyOcl-*zJ<O@YqP;4M
zuj2cuEWSo8EWVERd~aZ7@m=hDtgxv6e)Bz7PkMcTZ;zbwL$rJ3O<m9V5x)D)sVj>g
z<J+e`w=8}_%zm<=wk|cY@g}}!*5*z>MY|^~eug%dePQu)V&;7A?EM1m8J(xLCi|TC
z7QXYe?vceW(asBtU!jd9|25j4VeuQZvGfg#-x70gSo{v{{nnos-D{mJT5C;M{J#0-
z!{QHU<6%+XyLy}Db7=GIKcefs{e*8U=d?~1e<tSh$eX&p<6rRYVNM+u?NwR)72i)~
z@i$^&@prW6`vWVBLyo}jidbP$f9U3WuAcN72Hzez=dk$hkvDZc=WzNh*U^-Pm*I)q
zr#`nVjzHXgvZ1yvHL~$%JF_--I%3<C-Xq}~%f7HUGBI;LclM5g?-`w^wkG?W_wP1O
z>mFGgl^V|&7Dq!HOMY~;J;UM{Xk+Oc7RMy!-mo|p+WW0PFS^${S+v%gusC+}&4<Nt
z(8j}}zIW9F>$}$G*~dlu9J#mg(8h93>tu0!Vm^<&bu_W$CLnGPbLy~Yugc<t_<kyj
z6A=rG6Qe!fBv@JOiJg)a7WJoUzO#GMYifLZ<Sf&m-6M<Y%Hp*6?l-44m*<`i-#+!Z
zWpR39_LB{@b*Yh!N%1|i)-$+g2DE$9dq%Xe><f!C5i{p=XYb5t&*(h0HQDF9$?%=0
zb&o90f_7e5oE2>>`PtC+42!d)jiql`oP(Hq!{VH1@3;QE=w9n&(OPT5;#|!)9~S3E
z8xM>6-qqwRpF^8xp9fv<ZC-q1Ij429I3F>eN8Z%+9nX(%4|D3UXs^oR0{DI^iwhD9
ziwmJW-@;f~TpYUuD=g|S*?iB{lU_^V+au>(8ttAQoJC#FxeUJh&8aJk%i`OoKDS&g
zN6db*p|&nHvatxhXV&IUmq)uNEUthymVIGyMPlZB?(AI&?HQe?wkG?Ww<x~zwC<6`
zmC?=%i>siGCBG`#o?&q{w6XLJi>niJZ&+Ld?furD7u{=}ELv+#SX{ID=ELGzXyai~
z-@96j<#TBB>}#Xzy{&_9Ea$XN7S|=_^T?aJzT>~(+rykXEZVEGxB<SO^|A7|9(Kc~
z{m%e5LVLd6SXta0+lLhv^|xrg=UQ(F{y9K=d*qy3q1_{I>Uz$t@!fAuU0K`)-#+!Z
zWpP_#_LB{@b*Yh!jqyFRHg~!m+C5=$d$h6a3yV7tGv{+>?~Z8C=sdMG+2_1X@SUd(
zdpn`M2Yblo&S>W+zYE$}@_o_H2%Ed2jb%>Q+>Mxf!{+X2XZib`{Jzn>*2$)|)`ZPH
znr}XA?uj-YHub%`O<6vdHg~ZXy58g7_{MTh>tyq<#C#t4Q`dLC557Ijsl%qdDx3S_
z`>Aa1M=akb_eXoa1F*7rDE2T`*wjC~`JPKQkHEJ_&Uqx-J@Tin=R6AE{pQq_&7<+{
zQ=eNlk0EA1Sy5Y;8d*6I-!p4-r^lk*6E=@S8_T}1c|0+5K6mz>fcA{eQ(KdL&N~R-
zdD^gdBHDYfhismNc7F0Fqm3nh3fdWA^Hj94%n6&P5p!?YJRR+<{O(cfUh8DjT5H1Q
z8O=8zHqS&G51abl-N7uMOPjkm3tjK=Y<y!mr**P<4l$oc{?zrIpNnq~bLy~Zugd0m
z_<kyz=MxK?7oa`gg;?3V40|~%Z0cXpe9tADSK`|v=e!E-9{E$(b6$<_esk)|<~8{C
zsn0E&*Alaztf;L^jjUXR@0qo^)9cXg37glWjb&fhyn&cGpF4a1hW3okQ(KdL&bt`j
zdD^h|ceM9l583<&+WE=fh&Gn|O=xF?&70B2GAC@_Ld?Bk^H#L8WHP^xbgy-?X{|M3
z^S0)j51Y56jfYKr@9q+o&!x>>+<~t5cqhKGoYOklyo;F6BY*1p&hN&zhdFiFv{z;G
z9(+HQ&3lQ3&HK=v?|!UoK7xIe6*l!BYrf}_&ByWWk#qhN?H>74*K<CB?|yUY%I1^!
z_NmV;n@<t5pRA~@OO32NfbW^Lxznf7?g^XEpp9i;*nF0lIiEXw|AqF9&Qn{Hea?Fj
z-+9`w_Z-@Lu!n3uk9L0YFQAPj|03EMVe=)lvCIjZFB5Zb*n9=;EZNL|p6y=iWYbz}
z!se^ZHy<`%LmLm9`rh3`ET2o8yLcU4@9_<MV>zdFviT-4pGW@G^_{<kZx3_ouxYQ#
z=G*vwDx2>R3!CqvJ>Pp++58y$2`g;sf7*P{C7Yk&+au@v9PJ+YQ`d8Tf$x5E>dNMq
z`1YyKEt_8vv!AS}txJupypQjhwYk%;(e4SG-=K|UU)cPXm^q(2d%r__M(3%m$v)?O
zfbTqQ*!v#sJ=jAwe?U7w`5)27lK%<qjIjAL+F0g<&0mPQH*Efjc2@m;rF*TDO>3<Q
zo4+;ReAxUQZ9Hu1dv_nQd@gP7;tzDa$3O9n<($^Z=D&&gJo2Zm@4Ux||LI{)9X9P%
z*&KqLpUUQt#KPuKXwNq^RyIe#j>rm|`Xe>pbIIn&`1Z&-M?t$s{?zrHqvE^YoVv0(
z8oqt%bIa!F#Ox<4>P{Uk+$$@?5cAC1-02u-_k_(c(Z<p{Y>q|DoX?%TW1~Hz^VHU4
zpYw*rcb+!vjf3_c>>-=uqMe`ocxYqEkB@dn*qi`uEOWx<gv8t%HYY+mtN#Adz1GR5
zwbq2qiJNaeY)*nU9yay8yWv<qmo|4XDZ1X{WcbE%PU~cIa$-J@{Hg0Zp90?==G5l?
zq=UUGn?3RUR5qt17B;6sd%mf$vN;2GMpoF=pQ-ttOEzc5w@1!73)(&Mr>^Im72o~l
z)aJs^Z20!6&n=s?6SJSJsI5zltW1ONnYFpoIneG&?>W)NvM+4TMa-PfoxO9TJ)`r~
z)?}aarp0%jHtfxV_8#mZoAaWbpZt7iW695tc1GA-0BtOD!sdd++#5C*LOZMeKGVI{
z$)>f|gw2JUZ$4};f;JvD^}W04SU#6Fcd;nC-s58U#&S;UWV07BpGW@G^_?${Zx3_o
zuxYQ#<`VdRDw|6Z3!6)!J>Sw;*<1m;A}eg_uhe|cC7Ubb+au>(1??XBQ`d8@itm1N
z>dNM7`1YyKEt{(ov!AS}txJupEQ9ZvwYk$Z(C!JFYod*1U)WrWm^q(2d)G#LM(3%m
z$v)>Ti|;&b*joqfJ=jAw*F`%&`M;ozCBGip8DVpMw6V+yn;Q^wZ`j-r?X3FyP4`+S
zo7P$rHaBX%`LNj=Z9Hu1dw0vRd@gP7Vq<i@$4&5!<($^Z=BC7a9{E$(cfJ|EJ<O@Y
zroAego8$YbZ1y1*Hn%`~zAdq`xgB<UR@l_vq4}OmHh09gN6xtu+CB28uIJnt-~Hy)
zmCar7?NgsyHv1B@pRA~@OO32-h3}cQxzk<I?g^W_p^ar<*xa3%IiEXw_dt6_=c%p9
zKId(X?>ueT+Y{|Q*h4n=LOVbCz11cESF|(2=00d+nG-hmCFb6+xgXkD_2*6ZS|^*<
zS`#++Z@&4kc>vmY*wpv#wqf~P+T6u~=z5O_;Ty|2t&`1ziTOP8r>^h(5PW->Q-@7^
zRW=XB_fy$Cj9Az_9PRmzz{=*a*yC7XQ~&tpdoI~L0pA`u=ZR?d$e+5N^CW!tn^RXd
zPsX=TeQw!2g_!+hMQvSbWaUVF&#cXzo{Dx)*gOqwEc?Rd>BP+W+}V2u+A}&&ZB6z$
z?<jodX~W)`Xz#%uvUwKT`N^M+HkSN3XlI1YbJ4~!Cv2Wa%)Mdre6+Lb&ztVGPByKz
zCTw2NeDh)RLbUO)sqft#&GNalxr>X?^&T(AH<oi+C!3cL^LgY?UEleo`1UZT4x9F>
zY+i=%r?PoDv9Ngs+VfqBmCft0*R#T={teCdT(bE$e0$`ae@DAV{?zrH|G;;@Idx_8
zMtu9!=a$Wzh}ln8)Yhd&R<6SL%-Y=P&1m<8&0EmMvM+4jO3a+koxQiAJ)`r~)?}aa
zuEuwsHtgMw_8#mZn|Gj{pZuL@W69rzc1GB|8*MCe!sb21+#5FUMLVnhyy;%+WYbz}
z!sdO=Hy<|dM;i~D`rh3&ET2o8yLbRy@9{x=V>zdFviT4(pGW@G^_@SAZx3_ouxYQ#
z<|FuiDw~fI3!9IjJ>TP4*?b!N3@dEvKihoIC7b`kw@1$T9NIndr>^IG9^d`u)RoN_
z@a<EdTQ*-LW<Oa`TbCMH`6s?-*5*!MLc1qyzKk}OePQzzV&;7A?0psO8J(xLCi|TC
z1itgMVed7x_h1j%d>!rl<ljIWOa4u?Gs5OuXk(cZHs2=Z-mv)&+FAAIP4`+So7P$r
zHs5W&`LOvO+IZO1_wJr#`CQuE#rx=bk00P0%Q>x+%@2wBJo2Zm@BAZtdze#)O?y=~
zKgIX+3059I#(vhc|9Sf7XwUZrRyM!Ie#Z)%`rkL-bDhE=f55j#&iNzSJ@Tin=lluZ
z{pQq_&7bk@Q=eNle<5Z+Sy5Y;8d>?$ShKm)Ux~RVZ2pEemVIIKcVgyz?(F>o?HQe?
zwkG?W_Z7bLv|;a0wD({S+59)!`N{Vf=|8dLhroA6*c=kySmuPyp@_LRYz~cfmcM(?
zf1d7M>txefYr^I*%{L!5heaC?oBH12*DRk)o4Xi}T)oHP@r~u2*2(4w#C#t4Q`dJs
zBECJ$sl%qdDw`wW`>AY>Of26wM?rhOQL(Z)7Ith_*wi1V`JPKQ$Hlit&N&|1J@Tin
z=Nup3{pQq_%?a@BQ=eNlCnRP+Sy5Y;8d(_)-!p4-rxT&w6E-JC8_T}1ISDayK6my`
ziuR1oQ(KdL&Kn)ydD^fy8QOcWhip!cc7F0xpp7Np6YY$!IVIXy=7i0uh`BdxPK|b!
zY-YLFI@z?=ny@)d^Ua6NY0<{RroMMK2FvHt<}Ri~*L$2E-&oFRoovoP%;%9mb$#bE
z;@iWVI&9jjvN;pJpUURU#KPt*XwNq*RyOCt&dmy&`tvm3bIIns`1Z&-=R><k{?zrH
z^W(eUoVv2P0KR?dbIay}#Ox<4YU@%XE3@HyW^L|tA+&qK=E7)W*%vk!v4-VyXYZnD
z&*(h0HQDF9+3}sH4SS2By$5^9W-qk!lV2QdEcqqS&Ip@JqK#!v*j$R3d&B0^XlKb~
z{_}MAS|^*<S`#*xX}<Zexh&dv*wpv#=3x0;+T6u*=z5RK;~UF4t&`0ai1|G7r>^gO
zMSOdhQ-@7^RW?_`_fy$inON9d1?~A(#meSd*tJ<<Q-7W2|G%d01Zw-KY~wJcMCMEt
z8A_QF3Mry8G#46BXp%BjNFkI|WTqlyD1^*I5keu7jHyf^sgQ_N3g3CW-~U`|-)rr)
z&+FRPwD<k%T~E(jm<yX5$5x|rZW1jAfAl!#rm^KaqenJB7+an3-q`$5eClC^7S|Z8
zY!KVb_P*(6(Q?Y>=Fxo3E1O%y=ghsE`{8Ia>PL%do_-s~*3Vw{J`(L6)WGIPqxJ9k
z$D;Xq{_$u%%I22QeC<;<KM|kYvbj~XUfAs4r^^+GO|fFi<|ilHdD;9_G=JG-d%KOY
z+{@m#*gATA$8BQsbxv{E+%`V<!5=+-=i9|r!x>#R)f(B{KDN)u=BMK;o1cj`-)FO9
zbLZ@z&nlbjFHE+%u(?ZYH9F_6(Q@!dk8|!8TfQ@TWb=!$)fw-N&E4Zu4=c2|#$aWK
z*k-o(P4|eFQ#SXE=4)Qr{8D_*+`GA7jy9uyw3z1Uw_|Mm>}78+I!g^~ekEG}p6?ya
z*YmGN>rpnp7R}c_W%KLt$t|1vMC&#Fep9YEY>E|AHus%u=VkL7(fnnT?d^8TaxZ(|
zV!!C|9rusT*Ez*u^MLr=2Y>YVogWxm4QF)ORBL4Opx8bmn+L~NHV=t5-=W#Dd3g3C
zvdSj=JCkiLY#td~jm~*gv>g1=<D5swmhX%n+5B#7b;f&R^O*S5!wM~~F<3b)wwdjH
z(_^FMl+EL!`I=WYkB`rpdpGxc(Pq?-7SlZaz8PCTd)fPbw0BSgn<qr;-}4ir`Feg*
zv>s*i2hn`(Q#MbIPj1;fC0eiX_nUIXVN<M_vU%!c&#068tdiM&=sc@zvc286vfRtw
zw>T|&e8<yc^L0*f*gPXX_rV`Me&=V#R>K)xHq{#0JS(=($mZGcmCbXa&3A5gY+jK4
z!mP5%{>fyU3!4|kR-<!X94!Zb^f>1wvE@6XM>c;NTb=RV*t|48^{_&VYYbM-i*06m
z-}JI*Ic4+mXujr^&7Z~R%)OiY^Jp{bM~i8me&@&5&tCSfi1rR@VDrjo{d;~@G+)nu
z5v@nrygHh%eahxF@yRWlzl_#v{Qag}ao7|qrfgn2+0Ogt=C7jp%O=~~{V2=5?0t*t
zqQ`gqb!@)QDGr<0$LBuyqsQ<3hS+L2qsyjRBbztI_8Hl{DZa9KbF}$x$&St6XTL40
zY_e~kY;$4rj@W8+&O4*!;Ex{XyeqbRXY|PCA7ZOB-W!{D$EO}vXmO3f%5P$u+1@w3
zCt6O~yf>P!d1dpD@i}wv=H3@=M*V0p&C~DJ*!tPa-k+kqgBsYpKU)8u|2dkk=YNUT
zqip^)ny-Dz<^%D~$en4n8D|g1)@%Izrd)B@6f34|J~Y|R%jVyr`O7BT+x<4nz3hF9
zhoi@L{CjM^&M6L?kHqIb_@l?~{L$EIIHSv^S|gkPi0w17`B;2q^YLi&{WCi@pUnQ>
ztg^}e-(;H$n@`17qjNqTEeC(}IOj95<vXKCHfNgS|LctR#^%iNtA`a@Tw}2EM9$1?
z@0&hr$|;-AmNVp)%~|4e=HAVHPHZ#kM~fNz{cGxHFMH2TjCW82o3q9?bI+d_o3H1y
zMe9*EXOGR-K4tUy@yYF<n{!0#HU7LQR~$CQiYc3OPPX&1Iaf4)*<{an!_4Cy|2^Hq
z+s~ad_vky$6PvGdx(_zzjn94XM~~n6e6iJVMwd;sMmAp%+h=5R{`kt~0@3DMFgrG1
zoPCk3vdLa_vdx9fm&8`1b1oJw2Y>W9=i;&DJEKQ7mx!&-cyI5sWPIvjg%;NsY%Ubr
z%yu)#d1<tqzUj-N`I=WYmx|AsdpGyx(Pq?-7SlZa7LKi-z3jar+B>L$&84ID@A)#(
zd_7+_T92~1Tr^+%l+ESilUp`l8Lij&^QK&J*c2<KY_2fb&dcVjqWQ}v+uOY`%f0M<
zixs2CcU&npU*{Bu&6VSGAN<kdcmC?wYB;0IrdlJLuZit5vbjopWpmYN^R1R0n`>lW
zGplT}-!R$c!sc4B)##jSN6Wz<J<j>Y*z%pxBb#rEt<HFFoV__d^{_&VYYbLi8{5qG
zzUeyAa?0kq(R|G-o9o5r%)OhtezY0&qs25&ztv;wXD@qiiS`a^VDqif`uF^8(R@9B
zd$b;9^BvKA?Nc`28K2y;`L1Za#-BIkio>Q@F=g}JlkL20z9*W$Y_h%G>$2R--nZBw
zdVI$XWAk-RaoF4_KKH>NJ$~o!jje_=x@@X7viZK)J|mm&kFRWgAliHzXUFDd**DKB
zo9rzn+g#ZEaBMX?=SQOD;Ex{X{Ag_X&ghZNkHuDJyf-#K9-n$xp~W=@E1SeNv%PP+
zWwe~K`H5)0=9SH@;&bNS&HZGw8TF&ZG*7=xW9w%xd!LH-4r*X?>uCLZzD+b=&$o@%
zqik*$&DTC<bNl$@md#H`>oxwoDOVge#fmALpP6juW%IMq{AH8v?LL_0UiQAl4$<Q~
z?iib|bBe>}PVu=9{^;>L|6FV}oY7@dt+Kgu^e$OGU&xNf&u8Ct($06wGT#@oWAn?|
z_sS}p?5|9=x%Q6!YL*(E^J~#^@JIW{mGtcEvE@6Xo%w$B-X}{P_ioO<vA>a}9#&{^
zjls(9vCVAPN6vm(a$0-;$=5t=>UltxGxu)pfzf8vj~3HB{q~5hpIy#D*$>Xr4~yR9
zkjdBcLu0FH9(Ow|ny+=S$p6jw<YJNit!VG-ckTWAbGhQMC{|2a{Ptu!FN=pq^Or@o
zx7styJ?y>v5z*tjeJ3_w=M;y<Bja-)ywUx>oE;Tg4QI47{~rvsMi!5b?K86Y-S}|n
zyE!Jye8*<T;tAPL%qolQlP23--u(wzYIM$%qvha@9_KtIwtQ!_v+{CkY<0$a<LZZ5
z>S2Qx*BERZ7u(GCzUgVva$5WJXujr^#WUh_=HAUcGun*$(PEmX-|?~av&+HaSy}p(
z#j_`0&(DdiW?4Kpny+=s;(77OEsN(zd*A+Dxm~U}EQ%FV7JoF^&dcJDqxs7s+gp7v
z%RTJ9`vuYCyImNYuXBpS+fU+iAH31ycYIN7HJs69QLT~1i(~tYEM5{{S^R0V`7X_l
z#VfL3nN=3qS53CLu=tBCH9F_j(Q@!ck8@rVTfQ@TWbv1=)fw-N#cQ+F!v-y`G1#~)
zwwdjH(_cl)DT~)d^EIz5{yIKq?%mw$qs^!vEv9+;T^?IMyBsXukfmQ)ym9jN{HEAy
zmc^T+`C7Lu-V&eOviO^5?~B3y{kU9lSQIO!EZ#cV&dcI&qxs7s+gtrC%RTJ9`|qO1
zcl&*8zRoEQi?_w+K6s<Y@A&rEYB;0IqFN)1cf|G?S-dm8vUpdt`Tmd{i}z*!Q&w4I
z-#^*r!s4H^)aaaliI#&mdYtpGvE@6XM;0H5t<HFFEIyc}9yVxkjlss<vCVAnn?4jR
zr!4+0ny-0f@!|NKxp#B_9&JYbXfe&x@1EHD+2vsIku3en;-izV=l_VUW?6hJny+=s
z;^XnjEsOt*_P$u`--pW;hefes%Hk7~?Yu1hE1JJ7vc1*4S?*!)-TxguzT1DqWOYt)
zSbQ=*_rV)Ie#ifft%ftYEUGoK_`ldbBa2VPR~Da+Hs3SZu{g_|`QPegl|}Y*C)-?D
zoHe!@o%4CI<=~AT=bSCJd}s8?;_R{28Sjn7=f|%eHfV8;!NyFn&1~<R&JiuAEY2Ct
z*SxYgSA5RgySa16Hlu#DnC9s>b8P+Wa<Dj0w0>oA-e|s_&ljy`S$sh>U+b2|`QwvY
z78i*2zT@vB<%+|iSTSXB!O3=B78i=<FN<t%^{g!Su=nl@N00CJ!q|MBQydmw6rcOx
zjUK<_7spn^8C@3D8d+Q<w$I4oqVbi*mqeRyvFuoUS@xx}$|C#alWi_6z9P07opb4E
zIe4STIhTno-x)o!xNK~7#(QINx%kw>1}&~J*jPNaneBbk<)h`4#aBl2HLonL5T7&m
zZtkn1&8QzOrg{1;5nDgI94xLFtzTJODVne6D@UtY7GE9B*Scl#HSx(Ui>pL?-|_d4
za>Zd$teCR6>SQ}Fi>pQRmqoU>S~AN$?7jPIqsMn!JvLwG6o<vv#pgbFqsQ;~^|94(
zMwi7IwPu=mrd9K<JLk>T$eGW`;+pZ5#WzHoZ>{WDTqpawS!I#E-ej8#i|fZ$qjSC`
zS`Oaman85KmhX%nS$tb;b;f&R@$K=chYea>W3aJyY%|;YrtgTBQx@MD&DXrL_^$Yz
zxp#Bl9c@PaXfe&x?~Sqbv&+Had!qF#iyK7q^?bu<HOt~g(R{627T+76+_Lz-Xzx4z
zzEQ3?EQ%FV7T-VF&dcHlqWQ}r+grUU%RTJ9`^M4ZyKNGiuXBpS;->Mr58mkUJN{s7
zHJs69QLXXM#Sg{y8Cl#czOuM^wE4Ekj>V5>-!iK#vOh7|=ECAuvDN6DpNy7+H+r1&
zQ?cbcqem9Ej;+pkZ!B&TpL*D!#We;SAC7Hid*5{1XgOtZyJ)`VmBsDjbLQU7{dBY$
z^`pf!Prr}E*3T{ni=T<suPlBxny=?OM5|dAcZ}w1-Lkk-d~(a;=c2vu`1?h<;;<-I
zOj+D{vYnU3&qwo@MYgy4XqJ1}d-pFykMFiiY`)GZ4vV|S=RSC&$M1Nz*lIYV%c5E%
zi(icGGqSjQd}VQuX!Grv9gBNs|7uoQWPfe4&4tCU$5x|r?h`EsZ}d3lzOm&yqem9M
z5nG+{-dNl(KJ~Cci)#!vz7*Tc_P*)<(Q?Y-0nvQTD~kul=ghsEdr-6)^`pf!Pron6
z*3T{niw8&RR~8S6=Ii;P(Q1~(!=m|Gw=8}$KDlM_ThZQk{CQEXI4p`4Qx?BH+0M)2
z;nDnMk?pPa%5o2T?|ww|_-@~c&DS}_Ve!cL+y`&;_#Gb=TMcJ)SyXFe@#xq-Ba7dS
zuPh!DZN6i(WATLSCuWsJ_DPd%E-d~awi=!D<Y+l~qsKW<i7nq5J+gReY<0$aWATUa
zsfP_(Tw}0tTx>Jj`=+Nw%PEVeNAoqWES?dcGxu)pnbBs{j~3HB{f>{VpIr_X&x+Qs
zES??B*Yk6t)hvtWM)S39Sv)U3xn=SEXzx4zyeL;37R8Dwi$9ud=VkH7(fnnR?XA9-
z<sSCl{etN6-7bvH*Ez*u@h9=Q58mkUJH9Bk8qVmlsMg5h#j$-x7B7jfEdDgwe3xd&
z;uYDi%qolQt0voASo}q7H9F_j(Q@!ck8@rVTfQ@TWbv1=)fw-N#cShJ4;!?&#$e;J
z*k-o(O@9?Frz~C<&DXrL`0Mzbxp#A~k2a%zw3z1UcX@36>~gSpL$rQn@y2Mrp5GL$
zW?8&Bny+=s;w|yXEsMX2_P*oKi*m(bQLLD<c<W?4FN?p8<}ZtEZ}qb*_ptZwzl$E<
z?f0?yI;S`+-WH$x;Ef)?<J)7a;fyYeYK<)35!+{E@y__l;$6|^`$Kjt-k1GPS!I!Z
z|74pBi+_%-M(6xXv>d$A<D7quE#DbEviLx3b;f&R@xl1i!v-y`G1#~}wwdjH(}$wv
zl*PYA^EIz5J{+Gj_iparqs^!vEv9+;-4k0syBsV&60KiZd^DP`=l_USvn)Oq&DXkR
z@$vZNmc@TYd*AWrMY-azC{|2ad}6Ylm&JcY^Or@ox4JjWJ?y>vzoW-@`;VBc&M6Ly
zPsZmyc%#Sf_`k8$a7LF!wMG`7j_vbQcKrP>`!kdF-?L|$>;LDQIXf1gn|)R*N41|f
z+2)!pdiE?eI_LA9XW@+==bVE*oz2h*D=%}#ug-XHEY20bdf1@FH3l2citSD9ebc$4
z<+S!Z(R|G-i}S|k%)OgCU$hzZqs25&zh}qR&tCRk5bYh*z~=nX`uBW+Xuh5=7_CRy
zTqv5aeahy-@yRWlFO1g9@4ox@({jaOQ>>V>`J%~oUN&DG&0jXz-fosG_p<jb7Kt9;
zanaa(ol_h(UlO1D;BSVGuV;(JuZA<aY^pV~xp-`!k<BII>z|WLMw{=Y*|E5E_GPll
zCVSb*)*GA4WvS6wmXDT$O?qVWm9gbJqn(wX6=JJ1-W!{*icdYP(Bc|{m6ydfv)v4G
zR*aU@+ABr#HLq;09G^4yZtkn2&8QzOrg{1;6<a@h*?Ud2cTfYHt3>PH^HrnydcIn;
z9%b{j(R}SwHdl{NZrOZYv|iZkk}D3IV#Sor*H5<dvbjbyf7xVvyO(FVm%VSXX7u=u
zZ-~v;ImKaft@zvrfAsjBuN_+rXLQ+AYh?3{v3*81-xOcjd~>w<*2#{|w`PA^R@r2~
zeX`9pTl71!)aab=jFy8xdYtoJvE@6XM>gLbTb=RVID1cg>S2Wz*BGp<8{5qGzUc<h
za?0k0(R|G-n;XUF%)OiY-e@!GM~i8me(S~7&tCT47wsL?z~=j-_3!xyqWOBhakL&~
zbCYPk_9>g2#wWLIelS`uT=ws;<%+|mSTSYuLzC^iY;G3KUpCp^Zv8CxviB`Ej~?G~
zi`aafQyexw9H0B(j~>7CkHl8P8C^Ej8rl44Y@d<MkHuFuKOSwqEwf{D>+IWPl}+}x
zlWi_+ZkMG-=iEM84*uwI&QHgd?~ESV{7h_h#(QJ)v+=2i6<S<lu=0u6X14cDcZil#
zHg}BXYhKyhDL!ZJ-Q3Sbn^8YnO!M^HDz<+1vbS@zcTfYHpO4nR=U<5C>-jFxdX&vw
zqxssWZ0;7H+_L$_XuYs`K$cyuIBbd)Q#N;>Z0BWjk7)j~$@X@i%yKV#-(t_`@g2Vu
zo3C?<!{(Rcb07TC<9EJSY&D$GWmBz@&9B7v8QI)BzOwn%X!Ct7J2v;rzJFHPWFIiu
z=ECNIS!#67gQDf&j~?edIJSIe^vLERvDF#xjm<;jQx7Y&xW-`R>#@yj@0%VLEvIaL
zGn%h?W%FC{IdkvkemmNX`q5&Vr{6xY^|P0~!=t@}8rVD{TK}GZCz`M4M@H*WHjj$t
zYoD@tbbNBl=69p@8h>9YR~$CQiYc4NOt$l~d2BR)*<^dWeY4!l-nTd|dVI&@WAk-R
zaoGG`eC~rkdi>76A6pG)blFsEWb=gBJ|mka##c5^iZ<U5vSahK?5AgyP4*d+Z7yt{
znWaYOJS$ob{^)Vevt!G5MvrWs6I-3}-q<`hKJ~Cdi)##4PL6G6d*AfDXgOu`{Aj-B
zmCYZ;=ghsE`{QUc>PL%do_?pq*3Vw{E{OIHYGCuiX#IQslW4x4UlgrJ*}OQKuYJnq
zCGp8En?H@#YyADCTyfYGE2eB-I@!+4=4H|RWs~jgPR(*Jd*9;n=<ywY7Mrhgio@p5
z<8vSU(c^c1MQk;k(PdMuk<BY(`;2T}6<^u>MYQ>@&W_FNvi~}(Y_hMPY;$4rhAcHY
z=Z(>F@JEkx-V|HDGkRq6=Gf|t_r~Te@u`OuT3lnWa!qVA+xw=!iI!6~Z;j?_UfKL@
ze9qjvxxb4xqkgoQ=IQs#*!tPa-tVKmgBsYpEn5Gc-yY4^^E;yTD4Tai^R-XeyemGr
zW%CcwdX2x&lq(LKV#SoryC>Uu*}Ny3zihI--L+ZnW$#<u8$G_`A7k@%PI1`0FFyCd
zA3c8Oe~PV!GrDZ5HL`hsY@d<MKgU-#{}OG!zh=kg!`c6yRW{j=Ot!gZi+(gqjn4Ux
zXgT<!$2lL1E#DbEviW#yb;f&R^PlmlhZR~}W3ci-Y%|;YrcXr6DVzU_=4)Qr{C9lL
z+`GB|i8iBtw3z1U_aNW2m%S(B^A2iY^S{yh_xyj+d_8|ET92~%bTnW4l+9=2lUp`t
znmfNsj6ME-Q?58{iWO5fXO3n&FPqPb&0jXz-tM6+_p<jbo}IJt9cPKn*Ez*u^EvUk
z5B})!JAZC$HJs69Q>~HBS!4T*Y(6i(vN>C{`DV|K&AGGBlT|j^^G>$8usL6BH9F@D
zqUGR^9_O4twtQ#w$mRm<tnuF1TrfWMutJM#3|5{W+syX9=|a(R%I3n+e9bGHFO1Kb
zdpGw*(Pq?-7SlZa=7_DIz3ja>+B>L$%|)X1@A;z9d_8|jv>s)1v1q>bDVvMOC%0@a
z5v|ww^QK&J*c2<KY%V$3&dcUYqxs7w+uO~V<zDu_#ml0{cU&qqU*{Bu&6mgLKKP@@
z@B9_9)o@0aO|?cgmyYc-vbjurWpmkR^DUPhn=58tDXVO<SDtKhVe{3o)##kBiI#&u
zdYp5W*z%pxBb%$nR%g67Hdl*JJ*?2;8iSSPW1HFDH+^lioU*xkG+*<|=Ii2f=HAVH
zeY6?%qs25&zgNcA&tCS{i1rR@U~|oA{d@j~Xuh7W6|G0vTsxYteahw=<C9x9-xRIa
z`17XR8F67#teCR-=E-(mHrI*fFPm&{w?dYC+4~mjMvw2fUTnV3DGr<K$LBuyqsQ<3
zEwR;bMwd;sMmFCX+h=6+ZSj@Ow?~`r9oezDLG}%^$|ifG$u<`@-y2(v&iTG*IryW;
zIo}^!zB77c^8>Nf8SjnFjpI`fE3~-AVC9{$&1~<RZW1l0Y;GFO*Sxa%!T6lHcXK}!
zZASfQG0oHOU9t7Em%YuRy@MLq+&o(Uo^KJ&*YgiY>rpm863y2>W%Hx)$t{~7i`HxW
zc~h=9e{PBuQ#L<7+0M)6meKrWlkM%^o#kHkzQre^$9LQ+Hecryhs{sM=RWwO$M5`8
zvDI)!mrb=sHn)!LGqSl&d}VXnX!C8C9h*C3-!ZFfvUi$nb7AvyvDN6DJ4egGA3e_b
z`PlNE(IcB*h^@|eZ*1-opL$rK#We;i+s8Juy>Gf}w4AcJTQpzu%H|j2bLQU7-96fj
z`q5&Vr{AYz>t`=}dqjH&HL$s7wEjK+QZ!%BzZ|Vc+1x9duYJnqSK^afHusL!Yy5dr
zt~hLp6;n39I@!+4=GUV6%O=~~eJ0Dj?0t)`N00BgPi(%<DGr<a#^*lxqsQ<38?n`J
zMwd;sMm7(K?X!P&JnonMz)Ac2^n;?!cW`!Wek=R8v&ts>@X0pU5z*htQloPo87&8Y
z^f>2HvE@6XM>da+t<HFFY<@RB^{_&VYYbKniEU<k-}IPhIc4+MXujr^&Ew*8=HAUc
zKH7}>(PEmX-=VSfvzNW^MSBM|u=)LH{d;~wG+)n8jMk%Uo)pd3K4tR<@yRWlCr9h$
zckccBbh+ZNDOOC`JY}+-m(5e7`O7BT+Z~qWUiQAl52MF-JS{d~=M;y{)8lg={L$lg
zenxCHoY7@dt&z<$WBZJ3o)usJ+&nwleCK4x=8v;qkX1I>7f!ahu=$fLH9F@-(Q@!d
zk8@reTfQ@TWb=~P>Wufs=1=2O4=c2|#$e^#*k-o(O)rg>Q#LP)=4)QrygWW<?%mv<
zMVnDST1@lwJ1@3=_Okc$Xz!o~Hm``*zvov*^Y#3yXg$j2FQWO{r)*vwpWL!}O|)Lv
z?2;=En_|V3&0kKo^RjtuG=JG-d%N?q+{@m#_*L}yj@QNJ>zv}S`Rn-H2Y>YVonIeY
z4QF)ORBL4OhS)wMn>WT+HgAeH-_6;v`Md1D&nlbj+a}vwM?~MArAFtxBU%pr=yA?F
zW6O6&k8Iu*Tb=RV*!)9$>S2Wz*BGqa65Gu7zUkf3a?0jC(R|G-oA<`&%)OiY$7nO^
zM~i8me!q#WpS|qe7wsL?z~-N#_3!!p(R@AsbF?01^DohS?Nc`Y8lT*<`9QQ@xa{Aj
z%N2)Bv0}>RgOlyNY(5mtUpCp^?$#{#viB|i7CpY>!?F1~r#Ni>JwErrA3c8OkHl8P
z8C^Ej8rgg_w$I4sKjJH!k42mB@$A_APxdFX$|n22lWi_+{x3_7&iPcd9Q@JaoKMG=
z?~ESVd?vO!<Gryt(>(uQ4=c2|#$e^2vCVAno6bDtl+9<w=4)Qrd^TT}dpCEM*k;s^
zo)MEV%&OlLQ$KszdrsoKgBsX;ZftMR^I2o__569!dX&xCV)J#rvN?Nva?9rPqxHgO
z|2|!=IBbd)Q#R+AZ0BWj&S?I!$@X^t%JPo(zQtTQbC15`+_CvOr#NiR6QBFwj~>7C
zd1I^LjCSUq;iA^a=6tbzMmAp%U)h{L+I$OS$L5Q&zc{OGvKN_bb76DQ*lKjnmqg3K
zA3e^wSZw*uXlK3m;<42k?~TnR;!_VRw7AA#Wx?2Hw)ah!jF!{dFOBAFUfFzEe9qjv
zxl2WxQ9oKt^YmLNwtn`q_ws1(pawQy5v_mEmyYJ^`7+UZl+9(M`P!#!E*GENvblV;
zUgPgK<%+|mSTSYum6PqfY_1T^UpCp^Zs9EVviB`s6+OP=im~}Rr#Nh`6rcOxj~>7C
zm1C>nj4qpMjcmR;w$I4sYvL=Lt3;b`)$G`OefBl7$|ifw$u<`@-w<1k&bd~!9Q@Ja
zoNLFH?~ESVd}C~N#(QJ)P4TIR6<S<lu(DcgGu!*7Z;qBzHrI*fYhKx0H$G?X-Q4w}
z&8QzOrg{3kHnx8DvbTP;cTfYHZ;95w=WmVX>-pQF^(dQfkLGKiviXkq<d)5MM(Z{H
zep9YEY>E|AHs3Yb&dcVzqxs7w+uN<4<zDu_#e1U1cibR0U*{Bu%?;yoAN<kdcfL_<
zHJs69Q>~HB_r~@a*?eDoW%K>f=KDZ)Y<?*FW?5yEz4>ID3!7WSR-<!%I9d+==yA@E
z#Fp=j9@+e8Y<0$aWAkJ2sfQI>Tw}1Zacnc&`=%d{mQyyjjOJ@z+5AL&&fL4XTSc2u
zKUz%l^xGu1e)h8W$!PDO1~xwxt$)wAj^^w6Hqm;N&26Lk+NW%87oXg+xqY-=<L@`+
zio>Q@F=g}9lkL20ekPi~Y_h%GrdjS~?^}E}dVI$nV)J!QaoF53KKH>NJ$~mq#a6=^
zT{hJk+5B8=pOMX-<13q=k2c>IvSV}i?0aODP4=FXZ7ytnDYhD&^UKk4@JEkx?iE|U
zGkRq6E3wrX?~Tp9<5LeSw7AA#WtZ4yw)aiH8ZD=6el41>d1dqK@i}wv=I#@1M*V0p
z&C_q!*!tPa-oDY^K@DtvBU=BS?-$M2^ZldsD4PdF^R-XeJTN}FW%HnDy~f{f$`yxA
zv0}>R!ISO0Y#tKLUpCp^ZnrG=viB_xjUM0eu-JT^QyezG8K3*$j~>7CZ^c%_8C^Ej
z8rl4IY@d<M!{aNPM?{<NJK3>$O!i~5$|n1`$u<`@kB_ZJ=lou@9Q@JaoZpWv-x)o!
zc|vS;#(QJ)#Q4<13N5ZNSUEDbneBbklcMF6%^yVbHLq-*9G^4yZtf}3X4H=s(>(o-
zimji$?426z9n`?)52N+(`DxL7JwH8KkFt42G++Ca%`@YZTQ<*%)@%HEQ?58{ik%UY
z+5P)Y&Yo=NW%Hb9{<6vTc1LHqm%VRsZuIz$=f&phoZ_%~ethnOKYIMme-v8{XLQ+A
zYh?4sv3*81FNm*fUKnk@pJd18rP(jbDx2)fC)-@u{8?-@I_J-$<=~GV=e#1ed}s8?
z=9RJ48SjnFtKw4+E3~-AVCACNX14cDe-SOGY+fDB*SxZMO?=MWyScxNHlu#DnC9ts
zacuqUW$)T(@1O=Ye-*8N&##N->-n#v^(dRyNAtB$*}Ne>xn=XlXuZatH|2`MrdTok
zbMvOjc3w7bj^;0$Y;Si-mV4Rz7Pmx?@A#Y8e4SGqHgAp3eeg$*-}!H2tKp0;n`(`0
z{w}u9$mZ|kE1S1PoA37Q*t|RYJy~UweeYzO3!8t8tw!g(FIo=%=yA?J#g^}k9@)G<
zwmRdzvH9ot)WZrbt}$4-Bet3Cebc{0%PE_Gjpl1!*?b^AXYSqH2cylXA1$VN`rR2@
zKYQ7GDB3%yfz7`~>)-Q-qxpLN_h>!J<|EO3?Nc@%jZbdb{71B2<IkIN#bHydn6mlU
zWIHdLk4N*DO}4kYE6cs?eT#oakMH<IY`)GZ4x9gq&wcPmkKg&fW2@neE}LqNZ2l*<
z&&cMJ@s-X0Mw{<{*|9nEyqP1bY_gv{+2+FLEV0$-oX?3Z2Y>W9=W}DrcSetF&Kg^t
z@!r^cUi|7|g%;NstUNW%Z10=SHszGf*<;V}rL~@A^ZD^RbMNNP5!;OV(PGAaPfz{q
zWpB>Jcn3AGIhUNQp3fc4*YkOz^(dS3M)S2#*_<yvxn=VO(Rz(PZ^{*iO|fFi=KPcG
zylgHI&0jXz-tL*{UiQAlf;n@KzT-mCGknuKid!jq;rQJLfAsjBzc98M&S+<Tzo}N)
zbnh3%_Q}7gdIpbtFOCnJ^df^^G#aNb$-a8%biP=$e-^tOv10V$Zy&uEkL@#CmfCzv
z#MeGc4*I3h`n)W=K4O=O&cCvHMvho9`tY}pd0rmdXS~lV;%lF!2fa+RKFenBeU^*P
z|6FAzIby}=!{0vUSw6PUc%N6s*FGx@`c={TteCy`St)wH<jWB&Mj!t6ai5iA`z)B%
z+OLkU`Za@IWzefe>$zI?)$@XCy*7HG>~ho)FHT=)?dyK4$M#tyOE13H#aI3MNvpfY
zpx2Dn_YK+gmA6*(qS@uBAzqxm&f0hF=r?8gyfM38{&W52L9a9E6%w=V;9qah>qnd6
NE!m&SKhseC{{y=-+LHhP

delta 953
zcmYk4Jx^3o5QgX8h0ChR7T1U|QKMK`Xe<{gNo?Go;XkmkGP$MohR%}CdTs2DKLHgH
z6+!U>R20Pzz~{N{G3R8@&NFjn<~?(B9zLk<-|6%^<%Ow;B6?B9RCm~$DUJ>2XX~mP
zaXBI;*w>?ssykd<Ew&ftPn=q~Thzr`OOK$XmUf`smd=)SvDwl9T50JPbSRA$%et-x
zF`(gS`Ql75x_0Sy(LavDAFgo@y?rdd-~w@;{{tM)jJvN8|Hq^0{*8N0iN8eL)nU)_
zZ|pI##7{DQ3h$AzjGaTt+-I<pT4Sy`2b>tc%E#%)WP6_#eAf}Sy}4r(&j}`>)r?=j
z3cbWC<Up_BqgZRqHAf-iNyrPW;k)vIUK2|?>;Jil*uYWf4K_{m7Vgo&B9qNg$aoU+
zLYw%mY@&EaED62O_yerYM{F8stHG&hV3EltDYS)6LQb@e@0w<(=AVcq{+#ibjCWu~
zzhdY3hW7dfpTt_<SZINX3VWD@op2Z5b(*c1vG2qZ{}{84_cHzoEBp(quyw!TGvj*e
r$iEpTD(qnr_Q4*qT|QvNe6)5+T#VTVU1sr4#uKoYn8fb$5bggLks^#z

diff --git a/ptx/src/test/spirv_run/atom_add_float.ptx b/ptx/src/test/spirv_run/atom_add_float.ptx
new file mode 100644
index 0000000..3e3b748
--- /dev/null
+++ b/ptx/src/test/spirv_run/atom_add_float.ptx
@@ -0,0 +1,28 @@
+.version 6.5
+.target sm_30
+.address_size 64
+
+.visible .entry atom_add_float(
+	.param .u64 input,
+	.param .u64 output
+)
+{
+	.shared .align 4 .b8 shared_mem[1024];
+
+	.reg .u64 	        in_addr;
+    .reg .u64 	        out_addr;
+    .reg .f32 	        temp1;
+    .reg .f32 	        temp2;
+
+	ld.param.u64 	    in_addr, [input];
+    ld.param.u64 	    out_addr, [output];
+
+    ld.f32              temp1, [in_addr];
+    ld.f32              temp2, [in_addr+4];
+    st.shared.f32       [shared_mem], temp1;
+    atom.shared.add.f32 temp1, [shared_mem], temp2;
+    ld.shared.f32       temp2, [shared_mem];
+    st.f32              [out_addr], temp1;
+    st.f32              [out_addr+4], temp2;
+	ret;
+}
diff --git a/ptx/src/test/spirv_run/atom_add_float.spvtxt b/ptx/src/test/spirv_run/atom_add_float.spvtxt
new file mode 100644
index 0000000..c2292f1
--- /dev/null
+++ b/ptx/src/test/spirv_run/atom_add_float.spvtxt
@@ -0,0 +1,81 @@
+               OpCapability GenericPointer
+               OpCapability Linkage
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int8
+               OpCapability Int16
+               OpCapability Int64
+               OpCapability Float16
+               OpCapability Float64
+         %42 = OpExtInstImport "OpenCL.std"
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %1 "atom_add_float" %4
+               OpDecorate %37 LinkageAttributes "__zluda_ptx_impl__atom_relaxed_gpu_shared_add_f32" Import
+               OpDecorate %4 Alignment 4
+       %void = OpTypeVoid
+      %float = OpTypeFloat 32
+%_ptr_Workgroup_float = OpTypePointer Workgroup %float
+         %46 = OpTypeFunction %float %_ptr_Workgroup_float %float
+       %uint = OpTypeInt 32 0
+      %uchar = OpTypeInt 8 0
+  %uint_1024 = OpConstant %uint 1024
+%_arr_uchar_uint_1024 = OpTypeArray %uchar %uint_1024
+%_ptr_Workgroup__arr_uchar_uint_1024 = OpTypePointer Workgroup %_arr_uchar_uint_1024
+          %4 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
+      %ulong = OpTypeInt 64 0
+         %53 = OpTypeFunction %void %ulong %ulong
+%_ptr_Function_ulong = OpTypePointer Function %ulong
+%_ptr_Function_float = OpTypePointer Function %float
+%_ptr_Generic_float = OpTypePointer Generic %float
+    %ulong_4 = OpConstant %ulong 4
+  %ulong_4_0 = OpConstant %ulong 4
+         %37 = OpFunction %float None %46
+         %39 = OpFunctionParameter %_ptr_Workgroup_float
+         %40 = OpFunctionParameter %float
+               OpFunctionEnd
+          %1 = OpFunction %void None %53
+          %9 = OpFunctionParameter %ulong
+         %10 = OpFunctionParameter %ulong
+         %36 = OpLabel
+          %2 = OpVariable %_ptr_Function_ulong Function
+          %3 = OpVariable %_ptr_Function_ulong Function
+          %5 = OpVariable %_ptr_Function_ulong Function
+          %6 = OpVariable %_ptr_Function_ulong Function
+          %7 = OpVariable %_ptr_Function_float Function
+          %8 = OpVariable %_ptr_Function_float Function
+               OpStore %2 %9
+               OpStore %3 %10
+         %11 = OpLoad %ulong %2 Aligned 8
+               OpStore %5 %11
+         %12 = OpLoad %ulong %3 Aligned 8
+               OpStore %6 %12
+         %14 = OpLoad %ulong %5
+         %29 = OpConvertUToPtr %_ptr_Generic_float %14
+         %13 = OpLoad %float %29 Aligned 4
+               OpStore %7 %13
+         %16 = OpLoad %ulong %5
+         %26 = OpIAdd %ulong %16 %ulong_4
+         %30 = OpConvertUToPtr %_ptr_Generic_float %26
+         %15 = OpLoad %float %30 Aligned 4
+               OpStore %8 %15
+         %17 = OpLoad %float %7
+         %31 = OpBitcast %_ptr_Workgroup_float %4
+               OpStore %31 %17 Aligned 4
+         %19 = OpLoad %float %8
+         %32 = OpBitcast %_ptr_Workgroup_float %4
+         %18 = OpFunctionCall %float %37 %32 %19
+               OpStore %7 %18
+         %33 = OpBitcast %_ptr_Workgroup_float %4
+         %20 = OpLoad %float %33 Aligned 4
+               OpStore %8 %20
+         %21 = OpLoad %ulong %6
+         %22 = OpLoad %float %7
+         %34 = OpConvertUToPtr %_ptr_Generic_float %21
+               OpStore %34 %22 Aligned 4
+         %23 = OpLoad %ulong %6
+         %24 = OpLoad %float %8
+         %28 = OpIAdd %ulong %23 %ulong_4_0
+         %35 = OpConvertUToPtr %_ptr_Generic_float %28
+               OpStore %35 %24 Aligned 4
+               OpReturn
+               OpFunctionEnd
diff --git a/ptx/src/test/spirv_run/mod.rs b/ptx/src/test/spirv_run/mod.rs
index c99de17..c802320 100644
--- a/ptx/src/test/spirv_run/mod.rs
+++ b/ptx/src/test/spirv_run/mod.rs
@@ -155,6 +155,7 @@ test_ptx!(cvt_s64_s32, [-1i32], [-1i64]);
 test_ptx!(add_tuning, [2u64], [3u64]);
 test_ptx!(add_non_coherent, [3u64], [4u64]);
 test_ptx!(sign_extend, [-1i16], [-1i32]);
+test_ptx!(atom_add_float, [1.25f32, 0.5f32], [1.25f32, 1.75f32]);
 
 struct DisplayError<T: Debug> {
     err: T,
diff --git a/ptx/src/translate.rs b/ptx/src/translate.rs
index 7566be8..3291ad5 100644
--- a/ptx/src/translate.rs
+++ b/ptx/src/translate.rs
@@ -1505,6 +1505,7 @@ fn extract_globals<'input, 'b>(
                     d,
                     a,
                     "inc",
+                    ast::SizedScalarType::U32,
                 ));
             }
             Statement::Instruction(ast::Instruction::Atom(
@@ -1526,6 +1527,44 @@ fn extract_globals<'input, 'b>(
                     d,
                     a,
                     "dec",
+                    ast::SizedScalarType::U32,
+                ));
+            }
+            Statement::Instruction(ast::Instruction::Atom(
+                ast::AtomDetails {
+                    inner:
+                        ast::AtomInnerDetails::Float {
+                            op: ast::AtomFloatOp::Add,
+                            typ,
+                        },
+                    semantics,
+                    scope,
+                    space,
+                },
+                a,
+            )) => {
+                let details = ast::AtomDetails {
+                    inner: ast::AtomInnerDetails::Float {
+                        op: ast::AtomFloatOp::Add,
+                        typ,
+                    },
+                    semantics,
+                    scope,
+                    space,
+                };
+                let (op, typ) = match typ {
+                    ast::FloatType::F32 => ("add_f32", ast::SizedScalarType::F32),
+                    ast::FloatType::F64 => ("add_f64", ast::SizedScalarType::F64),
+                    ast::FloatType::F16 => unreachable!(),
+                    ast::FloatType::F16x2 => unreachable!(),
+                };
+                local.push(to_ptx_impl_atomic_call(
+                    id_def,
+                    ptx_impl_imports,
+                    details,
+                    a,
+                    op,
+                    typ,
                 ));
             }
             s => local.push(s),
@@ -1696,6 +1735,7 @@ fn to_ptx_impl_atomic_call(
     details: ast::AtomDetails,
     arg: ast::Arg3<ExpandedArgParams>,
     op: &'static str,
+    typ: ast::SizedScalarType,
 ) -> ExpandedStatement {
     let semantics = ptx_semantics_name(details.semantics);
     let scope = ptx_scope_name(details.scope);
@@ -1710,15 +1750,14 @@ fn to_ptx_impl_atomic_call(
         ast::AtomSpace::Global => ast::PointerStateSpace::Global,
         ast::AtomSpace::Shared => ast::PointerStateSpace::Shared,
     };
+    let scalar_typ = ast::ScalarType::from(typ);
     let fn_id = match ptx_impl_imports.entry(fn_name) {
         hash_map::Entry::Vacant(entry) => {
             let fn_id = id_defs.new_non_variable(None);
             let func_decl = ast::MethodDecl::Func::<spirv::Word>(
                 vec![ast::FnArgument {
                     align: None,
-                    v_type: ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(
-                        ast::ScalarType::U32,
-                    )),
+                    v_type: ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(scalar_typ)),
                     name: id_defs.new_non_variable(None),
                     array_init: Vec::new(),
                 }],
@@ -1727,17 +1766,14 @@ fn to_ptx_impl_atomic_call(
                     ast::FnArgument {
                         align: None,
                         v_type: ast::FnArgumentType::Reg(ast::VariableRegType::Pointer(
-                            ast::SizedScalarType::U32,
-                            ptr_space,
+                            typ, ptr_space,
                         )),
                         name: id_defs.new_non_variable(None),
                         array_init: Vec::new(),
                     },
                     ast::FnArgument {
                         align: None,
-                        v_type: ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(
-                            ast::ScalarType::U32,
-                        )),
+                        v_type: ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(scalar_typ)),
                         name: id_defs.new_non_variable(None),
                         array_init: Vec::new(),
                     },
@@ -1768,19 +1804,16 @@ fn to_ptx_impl_atomic_call(
         func: fn_id,
         ret_params: vec![(
             arg.dst,
-            ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(ast::ScalarType::U32)),
+            ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(scalar_typ)),
         )],
         param_list: vec![
             (
                 arg.src1,
-                ast::FnArgumentType::Reg(ast::VariableRegType::Pointer(
-                    ast::SizedScalarType::U32,
-                    ptr_space,
-                )),
+                ast::FnArgumentType::Reg(ast::VariableRegType::Pointer(typ, ptr_space)),
             ),
             (
                 arg.src2,
-                ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(ast::ScalarType::U32)),
+                ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(scalar_typ)),
             ),
         ],
     })
@@ -1963,14 +1996,13 @@ fn to_ptx_impl_bfi_call(
             arg.dst,
             ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(typ.into())),
         )],
-        // Note, for some reason PTX and SPIR-V order base&insert arguments differently
         param_list: vec![
             (
-                arg.src2,
+                arg.src1,
                 ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(typ.into())),
             ),
             (
-                arg.src1,
+                arg.src2,
                 ast::FnArgumentType::Reg(ast::VariableRegType::Scalar(typ.into())),
             ),
             (
@@ -3476,8 +3508,12 @@ fn emit_atom(
             };
             (spirv_op, typ.into())
         }
-        // TODO: Hardware is capable of this, implement it through builtin
-        ast::AtomInnerDetails::Float { .. } => todo!(),
+        ast::AtomInnerDetails::Float { op, typ } => {
+            let spirv_op: fn(&mut dr::Builder, _, _, _, _, _, _) -> _ = match op {
+                ast::AtomFloatOp::Add => dr::Builder::atomic_f_add_ext,
+            };
+            (spirv_op, typ.into())
+        }
     };
     let result_type = map.get_or_add_scalar(builder, typ);
     let memory_const = map.get_or_add_constant(
@@ -4287,8 +4323,8 @@ fn emit_implicit_conversion(
         }
         (TypeKind::Scalar, TypeKind::Scalar, ConversionKind::SignExtend) => {
             let result_type = map.get_or_add(builder, SpirvType::from(cv.to.clone()));
-            builder.s_convert(result_type , Some(cv.dst), cv.src)?;
-        },
+            builder.s_convert(result_type, Some(cv.dst), cv.src)?;
+        }
         (TypeKind::Vector, TypeKind::Scalar, ConversionKind::Default)
         | (TypeKind::Scalar, TypeKind::Array, ConversionKind::Default)
         | (TypeKind::Array, TypeKind::Scalar, ConversionKind::Default) => {