From 2dcd0ddfdc72e84e91b6a0686bc589ccfdded32e Mon Sep 17 00:00:00 2001
From: Blaise Tine <tinebp@yahoo.com>
Date: Thu, 21 Nov 2019 00:41:17 -0500
Subject: [PATCH] opencl sgemm benchmark

---
 benchmarks/opencl/hotspot/README    |   0
 benchmarks/opencl/kmeans/README     |   0
 benchmarks/opencl/sgemm/Makefile    |  54 +++
 benchmarks/opencl/sgemm/kernel.cl   |   9 +
 benchmarks/opencl/sgemm/kernel.pocl | Bin 0 -> 33824 bytes
 benchmarks/opencl/sgemm/libsgemm.a  | Bin 0 -> 4392 bytes
 benchmarks/opencl/sgemm/main.cc     | 500 ++++++++++++++++++++++++++++
 runtime/startup/vx_start.s          |   8 +-
 simX/Makefile                       |   4 +-
 9 files changed, 570 insertions(+), 5 deletions(-)
 create mode 100644 benchmarks/opencl/hotspot/README
 create mode 100644 benchmarks/opencl/kmeans/README
 create mode 100644 benchmarks/opencl/sgemm/Makefile
 create mode 100644 benchmarks/opencl/sgemm/kernel.cl
 create mode 100644 benchmarks/opencl/sgemm/kernel.pocl
 create mode 100644 benchmarks/opencl/sgemm/libsgemm.a
 create mode 100644 benchmarks/opencl/sgemm/main.cc

diff --git a/benchmarks/opencl/hotspot/README b/benchmarks/opencl/hotspot/README
new file mode 100644
index 00000000..e69de29b
diff --git a/benchmarks/opencl/kmeans/README b/benchmarks/opencl/kmeans/README
new file mode 100644
index 00000000..e69de29b
diff --git a/benchmarks/opencl/sgemm/Makefile b/benchmarks/opencl/sgemm/Makefile
new file mode 100644
index 00000000..7fa16540
--- /dev/null
+++ b/benchmarks/opencl/sgemm/Makefile
@@ -0,0 +1,54 @@
+
+RISCV_TOOL_PATH=$(wildcard ~/dev/riscv-gnu-toolchain/drops)
+
+POCL_CC_PATH=$(wildcard ~/dev/pocl/drops_riscv_cc)
+POCL_RT_PATH=$(wildcard ~/dev/pocl/drops_riscv_rt)
+
+VX_RT_PATH=$(wildcard ../../../runtime)
+VX_SIMX_PATH=$(wildcard ../../../simX/obj_dir)
+
+CC=$(RISCV_TOOL_PATH)/bin/riscv32-unknown-elf-gcc
+CXX=$(RISCV_TOOL_PATH)/bin/riscv32-unknown-elf-g++
+DMP=$(RISCV_TOOL_PATH)/bin/riscv32-unknown-elf-objdump
+HEX=$(RISCV_TOOL_PATH)/bin/riscv32-unknown-elf-objcopy
+NEWLIB_PATH=$(RISCV_TOOL_PATH)/riscv32-unknown-elf/lib
+
+VX_NEWLIB = $(VX_RT_PATH)/newlib/newlib.c
+VX_STR  = $(VX_RT_PATH)/startup/vx_start.s
+VX_INT  = $(VX_RT_PATH)/intrinsics/vx_intrinsics.s
+VX_IO   = $(VX_RT_PATH)/io/vx_io.s $(VX_RT_PATH)/io/vx_io.c
+VX_FIO  = $(VX_RT_PATH)/fileio/fileio.s
+VX_API  = $(VX_RT_PATH)/vx_api/vx_api.c
+
+VX_SRCS  = $(VX_STR) $(VX_FIO) $(NEWLIB) $(VX_INT) $(VX_IO) $(VX_API) $(VX_TEST)
+
+CXXFLAGS = -g -O0 -nostartfiles -Wl,-Bstatic,-T,$(VX_RT_PATH)/mains/nativevecadd/linker.ld -march=rv32im -mabi=ilp32 
+CXXFLAGS += -ffreestanding # program may not begin at main()
+CXXFLAGS += -fno-rtti -fno-non-call-exceptions # disable RTTI and exceptions
+CXXFLAGS += -Wl,--gc-sections # enable garbage collection of unused input sections
+
+LIBS = -lOpenCL 
+
+#$(NEWLIB_PATH)/libc.a $(NEWLIB_PATH)/libstdc++.a -static-libgcc -lgcc
+
+PROJECT=sgemm
+
+all: $(PROJECT).dump $(PROJECT).hex
+
+libsgemm.a: kernel.cl
+	POCL_DEBUG=all POCL_DEBUG_LLVM_PASSES=1 LD_LIBRARY_PATH=$(RISCV_TOOL_PATH)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -o kernel.pocl kernel.cl
+
+$(PROJECT).elf: main.cc libsgemm.a
+	$(CXX) $(CXXFLAGS) -I$(POCL_RT_PATH)/include -L$(POCL_RT_PATH)/lib/static -L. $(VX_SRCS) main.cc -Wl,--whole-archive -lsgemm -Wl,--no-whole-archive $(LIBS) -o $(PROJECT).elf
+
+$(PROJECT).hex: $(PROJECT).elf
+	$(HEX) -O ihex $(PROJECT).elf $(PROJECT).hex
+
+$(PROJECT).dump: $(PROJECT).elf
+	$(DMP) -D $(PROJECT).elf > $(PROJECT).dump
+
+run:
+	$(VX_SIMX_PATH)/Vcache_simX -E -a rv32i --core $(PROJECT).hex  -s -b 1> emulator.debug
+
+clean:
+	rm -rf *.elf *.dump *.hex *.a *.pocl
diff --git a/benchmarks/opencl/sgemm/kernel.cl b/benchmarks/opencl/sgemm/kernel.cl
new file mode 100644
index 00000000..17ece1d1
--- /dev/null
+++ b/benchmarks/opencl/sgemm/kernel.cl
@@ -0,0 +1,9 @@
+__kernel void sgemm(__global float *A, __global float *B, __global float *C, int ldc)
+{
+  long i = get_global_id(0);
+  long m = get_global_id(1);
+  long n = get_global_id(2);
+  float a = A[m+n*ldc];
+  float b = B[m*ldc+i];
+  C[i+n*ldc] = C[i+n*ldc] + a * b;
+}
diff --git a/benchmarks/opencl/sgemm/kernel.pocl b/benchmarks/opencl/sgemm/kernel.pocl
new file mode 100644
index 0000000000000000000000000000000000000000..6a643faeced315489bc199fff8f7be88e637c81c
GIT binary patch
literal 33824
zcmeFZdpwls|35xTMy+Zpl?uad2fJHRIaLk`$w~+zhsjD<2_YuTRoh`YHzZUt>`HA1
z>EKXK2}yF8kh8(WB<BWWjJfacHTSru_kKR_-}j&2zh95Ly4?5keqGPk@qS(ReeDJX
zdL8-Q?<nE^!*7w99-cxXfk2Rff8^mG@sH`Qjk|1ZP517!*=uEOx!c@gr-j*Ga|;Vo
ztKC++t@rLV*<)t2$I{B0@B;*I3_2F*bIdb9^LMY(KM91Lrms766bXbm@Y3%JTB|Kx
z3J(w0(KeXLY!;s)s4EZ%2C4+Y53|$>v*7Q`;2+|tpJGhj$y&;hHgAoUIis*<j{FK$
z*-O8lR5G8nGxmp_rKHV@`tnLv^EUr}es@r*`cGzy>6b|+v9k`nlYgqHX1eRVT7}H!
z-_I!hqVWr3Q&gr~jN03GQ0@ga0%0Beck_?u=c*k#E4$>Uv@`H8&Y4t(;^IA*FaL0+
zbPs8uSXoAiq-Ak(^{E}R=l*QVh*6uR?qJJUF>BKbHOBQ>vrPKnzuL1Ir(_fzSD3MM
zF00wDP-4WyuUQeiA}&TxmT_v9+SwBMSxSd$cOQC7ynJYF?QY-r%V~#GUhh7#?DFTU
z1d`)C{i|LlKmSM|u+FAvT>LRo?ch~{VmhQouu~!sUO?los$V;a!XFvu*{>H8E-2B%
z8APfEJvfDU{2g7FK@Y!A3|FHcQ={vu(L*!op>XVbntGl{O`+>*5hHw0AD1Vd*g<@)
z{N|KOZ>lokPv!T;Djg=OEuN|^YRaEem0Na@-<zm*$Y{}n4$`R_#FJ1?B0Y#iJg!04
zO`!+vCLZ@a9V|~g4zChEs4F*ls=k*|Yf4q_g%Sv=P{LbvvE-K5%AKhyZ%!!_LLfMi
zevEMb_-<mj?`$YXe8zWnq?k73PuI<bJjJXC3kYviI#TC<Fj0Mzs`5T{erK1;$6}Q)
zsgT(G_oq}orm9G#$%jYKb)i5lI@Nb}$Vs}cFFh!Q82X)3TbD}94wbK|DxInGzo;vh
zo>F=5srv8RgXlp)|F_!@s($QJZc9~umpXq8a_EA3&Ho^y+L8kGQmy@Nq98id>jZ-y
zbdneWDe3=jGQFo%-*_s2f*L{Fo>G19sr;o&1u9jVs`{Z=85+C=&SgS=FY|i#!9?N-
zd3vz3*sah7q58yd!ub<H^iVCat)THEBZyRIx~@Kv8bLn+6;Gsx)<W}sF;V>r9r~AQ
zo2Tj<DD%Hly$^jVz9n3Xcmmq(AUzz$gVbnE2jyil2@P^yq?2>*>yk<d>sL+RmBd>R
zVXnsP-<4sgh+TY2wMm`avO~2aRkiJu^7~WDEhm*f7OQ@Q<J9@R>Z+elsbT>zzeQd7
zi_Cv`?bvQQRh}4p5QgP!sAlLvx~>vb6Gq5?#H83$(A3aY&`=*?fJt;wSAO^1=DzgM
zAi6G;qc66SZYKSh67dAI47AT~;xTa~bQ09WG4NXXv$}E%3`l}%hq_7|G(NO;igJ^>
zO3Qa?Vh^4U4I+jl{`WM&!m>V9wM9m?<)rEx6J>G0|NHjY(30W*x7&NwRjP|s#I};8
znCJh8iNt}OLWEX|I4RaK<bUPt@2W7K-+HQu)2&OK9nc-(6yBlywOIKBmMm~C7_kHc
zg8T`2B2}M$ECuRCKL&k$UmRCd!lxaoFrnVakpF)M%t<A2H8LhF-XRl~DRbzm_%||v
zcyjiRS-<?hX!(E9^6&DJq9%zJe}cwISzGaQ%@RJ>WZ}O{&hJn<Ki6cFy4G6R@-qaP
zCAw5tDU^tjiL<x*o`&Tw45nx<e180EeC!}Gq^b6X>8cl}RpU%oy*{n_;gs^1V&!)x
ztI|%Z)WRwyqw?a^c6jnd@%qwJ%AKb^&fa;Zl%f8U`7Xj0(%O3`Ms^mp_9P^#?5(NX
z*r;!?D}ge?PZS-vR^k;ZFyENd)vKM?wz=-W3)dfZ9lhc6_-2*P;b8_ro<ulEAnbs=
z32F!FvN|UTgxLhFPa()`Pdh3jOMq&T#Fyop)Kg_-<q294DgK?5Y5L~0+`n#jnD>Jk
z;WdG9Qha^K>jiU{sqG+~gwJ|$uXXaI%3^(Oc<DdnA9seA<Ro?@%sRe%vKCGa`z5R|
zlO16E_Se4&RG7{F7lEk@bN)pjZP~vG{G++--voSD|C>O)*16P$u@c5=kl{yvrIJpf
zrNkm6sKP(G=RGG+ua@X64^0%e>aEtvor@(#k%dN>y0S&<+`ipV3)tPA6l<~enB3cg
z^D6#Bh_%8Y{j>ieM0?#{_5U=~&kmaZ=9%Gk>R*NmSD$@&$G;?Oxws|mKWThjA-mE+
zeXhiIv!JtXT$wfhAwyPTTp4K0j$NnZ_V0!yah}w_cbNXVTH>A^Vt2FbQi_-Uhc))~
zE7gw7`wwgQ9ar=_c~<?qHR`up>6E^KAt4SCu{CxwIV%r&!k7`eP)ws>;a~Hq+VC=D
zLJ-p!BC_NRSHjCU4YVX^HhDh044GgwmhODvc};E>0lJESIYr^;j#`V^GVm&#NbD4P
zqL$LMuiE!n?#7tN&5o8c{f_CoqPj)t>lUSX-zTDtr4%1z7TbPj`^Ory`R9ZR=f2Oi
zZ-e?DTAe;xu%X8ZpSa+8THW|`>UXW<_Fi74FttJf_V*IbrVmd1^7-Ua0%U|);w0hZ
zN%b2>mZ1cR67V$qkum)x{?7q|{7HDz%+IWI4hBnE@OL=bNy$0c-A70o(;OKBW6rF3
zD_}3Lnz(+&qCJeubJtqDl{KBUi||_t>3yoa23aP>lbH7NaZ*hnUC#Hoe%AGisaH=i
z2-^t*2k(2St$n?KuwCh^lDvH8?AHs<ZIk<P)&L=M_GVkcPLeI@4N0E0P&VOSd+3&T
z)A2bap|wexHhDEE&)kss{vBEdn=an??Jm{N<5oiNU24UgpAvf4EWhlw^7@upI~Z~s
z=I>v5Qd=otM_}jmI3pR=qD3E0F5T(Gu^4DXZoNv2>yEBAsqu&sOg$a4biXu}aIg6q
zt1jTpw~FUX>sE`~e}4MCv{E?x@F_mJ1}R_g8*R9Ix%#^3>iUCM!w&6bu$MoctL(E)
zot0q~eMRk6jo)wU%TLga>4!r9d}S3YlN)-SadlNifb(bu|EqHxAblp;9KPG68>D9)
zJ=QQ6{5B`%dh5YLjTsB?XmD1S`BP2)PZ{L&G<p^}k5X>`J%$#K{D3ma-2-UWqeO%`
z_^JmjTSrFN?oSHQ82u+gc3WW%dUkxbklfmviL!LsrpVRl573^`h$*rW@;BP0zHo~C
z<ZCo~mzKe2=Jto9FAv)AnQea_Lz_JX6O?xkkD!B*NfVSGz8*%;_U)b^JLh<!4PHY$
z=5GlGXjX3$k9>=+ho&Cc!($H7H=|ESroWMYxU~UIJ(}^2oRp)1mUa1lBU5^RLeH*M
z{zg{*3q|_$nuf`jzfK`BTdoYVsYMgW)3M)%$^ZQI4Y?9AcbMIAZ3sD=_il(SxcC)0
z`|A7<JNH^Ia^;xi5P6W<jXdt0H^dCgD@4wAf9Pkg?0$t*gkI}szPgr!JkC1M&-~^6
zb3~M?+E3YiDHFLd`SmNg@6#isVf@KgcK?HakbfegzcM${(~&DQtFM%#{=bm3YgT?`
z&T9e4&0qVt6xZq~q$eYnONn?-Mb2iN=aLPok0Yt;T)6DS#t@`$@@FpjkFP<<O;2Gj
z^Y))dkv`-@FJ;AF{>WV|W-sO4TVLen3u-U<$~SN1#y@*{$?ku9A)|xjUb1TIK_n(s
z$YB;H7$FU*H5}&p*6qmSp@$r@&EH#*v+3tKWR-8aNYy=e4w+V>jhy|qjl=f4y&1Wi
zro^GV<8MSBM-O#VF0^bw{>lH)&AwCgE0TKfVK)<b@H3LCaj~16bZZ?_u*I*N?0a<$
zayE5$H`D6ePe_cGMmJ^mpX$gdHTiC`yVu!S@_(u!6^Avu*%j4mkh@PWb+dgR{)$`>
z_I0y=xU~ssIBdybeu>vZMA`pvn5yZ9$f)KIz3hP!FXZvlfL?OcgddXS+1yK>&Vh#b
z(TK}Tz8a47O+4f>lP02(vL9D{Wovv(Lz;arePsur4-wYD^jCJqU(b=gr-A)!a?2~E
zC$GPsDf+t`iP7{MVn?0(ic}1+huE#cQDkt6GE8>4IEC~@^$k;g&G-?`I)3CE^A`6P
zw4s0U8<YKN3);|0<B_e>4N%tB1ryAL54_RAZOjP@czYCW{<4)%?huBf4eP4;?7-?&
z^zO9(6#K%fhv+{Ke-N_$-ejT0f=5E8M{g0@{;M-WDN6r}iq?t{cH8N37{L!wGPIAp
zT-=p4C&dQSnxz#__%j6!PtC3VOBW8+z=`RKx39#*k-`&S7Y=<94_6~6CJkD;#KS1n
zNTk0IJNv~jvM2W%2LrfAZhhZBC?0A@Mc$Ral`kINzY_Vhqg5ats@;j)pvgwW!_Y^O
zKY`>)@vt*5vS&m5xOlj_GEzSyVM;ta)Do%FT!#sZ>Wd7Ie*aZ`_R(adPKV^STr@JT
z^5UrY?2#X6U#6#i(u2c+1vI_K57os(-<7nDla)V;hn*_4jk~|C6A$%&ruDRCV1P&K
zX`a`fV05N8(w={LgPlFHon~0cD-z3nVMNoIm%dtj+q5N3tLo1+;^84{+L!(ZJTV~6
znKo+t4LkeUm6np6zE%uaO`!!eJ;!u2_=EQN8XKE9B7jDFm4VspQ3MTYp(e(5M$+;(
z4r3aR&}f0}-Pl>R3$*9f&oSN7F4Cy)um3E@YFwt}2%fAL4<B8jWx4Wx5f8Pl(fa0<
zW7IPfX<Q*)a`qN&_}8!4F!MI8YcL*Tw*C(7(Kfy&CLx*j?D1R7EgJV}+m6JJi4{uu
zho*T^GR5=*+R)=Eq4;d*BU*s!5N64N$FyGxzF{^y_=I)@@HdEs<>u0;RS6jN^{;5#
zqu*dQOUa}8KbpYKj^xt{hS=D|hYD#0;2|dLMG@`C;WwCdR#(uBj$Fem8Ty{K=IA$!
zPG==;`$D0Hc#8E^G%bUVnDP<Tv|r~oV#Cihw4%0an02(8X%8=5#H=3CLfaP5i)q~1
zO7r?7U1Gj=+N0q-Y_^~d+Tm9hG2LFUY0nLFFhfm#q4nBGNE&<<`)#9m;*@^cU#ki+
zH9rr~)TVh7zC*OT*2x&acbN9muYZZehPppVn?CXe)A);k7O?OhMrXB<=JkM!@jWz6
z3j{(8kRqZ*Oi8HcBeW<oN5TQ6`76D`_}-TX&#Wsqi#ce|1uqOyUGebqT)=x*iHc_n
zng@<XPhd8io(CRiJi#=+{{!&acngaMEd`)!RHG-pEnflrv!Q&8c<4JHj9D`=IwSMJ
zK;%OV5VZhso3b$vtzHNsjJtKj*vy5XQ!s&LmcA0$Y{kXw(5VE9o|J6DgslLNl7}&p
zBP&4KFH_hsVkPK8*cd=<6-awkgUP+W3f#{tkp!YL@bT}!6ndcyex_DyiG>-c0AAa5
zEEzggfL8z;J9|hK_&+?2ksMG3=&yfc@eo7?5iJ)nz!({PQIK$m&;VXdMcQIv(;8s<
zJ&Fxeeg-hAvEk}p!1zBk7|Hv;fS||{%!TWJ1zW*w%(0Jt1$jSrW6ZSH1GS_YY>F4_
z!SEmHSiop+0FP>7G2LEl052bNF=qOj;Bo5^=E7V}ur29NYzqC2;COyHX36}GV9hUe
zSWPw12S&fKr^PPJ*9T+5Cz2G|exktrqa+LSPDEHHNPMgi`TO8U%$p<IA~`4SVH~vl
zB0W8Xm|HS0MSfv^!}2XNIr5(lwj{V-MjjIlNn-s&WP1LkZ{pds*pYej(y<cMDTpjm
zYQZFE&7*C%YQfYDT1fj^^%2u8Q<-MeSNE&flAS-(f{+K87HS)5g~r!<#KP1JXodkL
z5<^+i)(kwvBxvoW4PWTMuv)IP<Ez*bfG;i5KL?9GtpM8Aq*oX-wG*^H!6ocASo;oo
zBuIuAXt{Y0F&`hiObczlhWRb?D(zX@Iejr5wVSlA7CJTzx=q`H{DlEJ@6etlOX{TB
zeHur#+(3*CdO+h@f0HQmi1ssDxE+)5M63n&Y&rOn7V@D8O9oi3xYZ>Xht2{TI8liO
zt=c=9?o~;-(t1zJH%K<b<bI&tyF7vEmRU{HAoDWCCh2UV8P6ZWlpkrK1sdJLIH<ML
zxT1bz@ofj$H15L$iCg+;+U`YIqQeSvVoy3&5mN?fYbWcB#MsPXn#N$!Hu11iKoi{m
zglW+!r1cmyVjM<9v|~puNm37`WmZpN*b#Z~?YCAR@oj4J0Q!N8O`-M!=$DVjbW>9R
zetwduR+|qz2Wv3p)fRvU!dsH(&;s!6+OxxA2_p+Z*8s;08!Ca-gKrLE!xf-w-mR_J
za3z42EWngs1)>eV{w_Y-sSE-FrZ7OJ3K&F+FkzXhz*oe@bW0%v-klbVWQqnj47#zi
zK|cfSy?Gcm=obL}hDkX1E9lki$K)Pd5AJ^H9>GF+1Ms`rfdSMs!SfZTJuy0(K=)Gh
zA&kRD@RO(po1#-6{6T(#DWAC&Y`XR*7CR~1fG(4b*(_)~FbsH(IqRSSxO)e|hFXST
zOKXwDg+_p1B$U|K2t1yChDqo&2C0_sCBd~Da0C1M#LR*$fv($M-eL|1e*>Kl_!u)S
zEAY%PSt4u?P;+~R1y|=D5SoW#d^4?qe{(h#B`JGBkWD#Om=4;2KhPG;0b0&rOFJU5
zx-%G^=3-NHx`1Ds9%6D+4uGKOixTs>g0%kASgdQgft;Ifuqo8s!HXlwl5BAYU#;RX
zW+@)9=zf*lMghU9*D&3*e8I+lh9vIx1^Ld&lG*&grh*Jf(Eb7b=y)dakUuzfzZ#>i
zb_5*CE0P4k5%8pqgH4<g0AN&0rZ@`vFWr*3Fc5%@YE0NjAXt0hPfX*?2(ZrMBj&6i
zDoBgEg$=b%fbQyS$usK&(5YlgvLzB6H^|0(d@u?`?&e@5)uMsmw{Fa4ozcKB>WM^6
z8YobbB=120K+=>rO~vgm+xbs0bOd>9nEQ?G{KklxIae-|OkbyTc}?3t3hxivrr*^#
zUF<yT%g4}#a({Q!8Nv&XQm&?Zb`Ed5b*|WX=boG`DsrJ6)uHghfalMre^0r1SK(r@
z^P5LCW538{8Pz?67c%cYWdHs;K7)3t*!j=M_*{LtVVCZ1c!9CWD9>xh+sT2e@D8_c
z(MED16?_4_@FGAy;ZSY1x?3!~Be&Xnk6ib*giR|huZatp{C3E;x~uX!yyHp2#shLi
z9)-T}0^_kljrWf1%?1gO?&Ya5ik#NhIsh-E9JhLS_+-n5$G6~xpIb5l<kEcevf+i-
z5fAA;1kPjCWH?7&hJU2o(5s9=c;Soxuf0AgPaY>S;RUqymkV-8;KS7`FRvLJt9avE
zTet3qzv1K`-Va=r+w$x4?eIdU#&ojZ$-hS;)8QQh9hpgT&mv2X!3*#F?@alrJqi1I
z4^FOnI{blLghIzZ@WO|JfKPvD@v~hXis$&tI9Cq2gml6SU!we<`KRz6S7bozPCxj$
zP;T^L`XsziI(hl(5fW$1wrB8yYwXubxd6)t8&_S1s~ClDN=5I}Ch>x1WjFg$EVWuZ
zRP82PW_LD=hwHn^lx#x`u&tXN_%Y*y_^eSkyD#}C46vu0BL6BEyUn?q{GcTRJL}QS
ztS-u`5o3>ZGY20$dMqABbd!Jmwh<FXgDcC{O%>v^7rWV4ZW~}{6T8`d$=cZ2q;6*5
zLoMv=gKmo88Imi;W^_|7)x5;yzUpRg&v=H-_O6@amhl<m09U^i8SAmLE#2g;Uw^6+
z-`3GhnRhM-6ZW;6a=T7|0Vcc2+u1_ww&`wiJ9`=f%;T`#x{p<gZ(G1&KFX13v64fn
zp4f`rroo{MoYunZu%1JCeDN7ZN1sEkr2mWoj5rj7)1L3e5`N>bFIA3VQ|#ri+o~U8
zQyk#1ldk1rSYHnF_8Sxvc7#J&KYSMx7QvwmKj^}+Q5=f?wLA=akweyen2KSqa+tqe
z{1L<6;*b;n3c#?*9EvYp2g5$#FkSN|v0*OM;$k6oTRw-puv}uOcO3RYjs|wy2M*ce
zTpl)T;*dqx{IJ{FIP9YrdD!qPhkWVpdzeXvILwxgub5G%IL!A$BbaUohq>{|XN=Cg
zUiKrz2NSlSm)t+}5*x1UWhV};!-T2!k_XUWZ1_trnfy><2hCnK<=Yf?)}WU;SGWxe
zg5AAr|67N#iT6Oc9YL5D&b^ejZi&s@;P62TM#s08{P68JY<Q%XqV)H#7_*38%B9=7
z*f6S>awUEnHay?U?n}_ahL?NEdmnDVhKca(0|6!>sh2#I_W&E-@1?Zg(!_?3dYSJ&
zrePd%d&yflc^I>NDEF#1=I(dB%oV&X1!B*?@1-o@^<oND_p&GdPQ!*Rz2xn0jId!x
zFEg@o9J9_>$f5inY&g`*w#wOt4X1i3!+DvQQITGDCEA0LoX4g7hDaQ+fJ+|c3o#BW
zxs+A!({jZVD|6Y?mr}4PG`Q?QREP;%&t*>{q1aHLOKwic#^f4s*%6IVnB3i5wmf<Q
z)9p7d+x_7*CSfm^`Q@4qHax&3KZt#TN$}t@`#E1Q34d_O+rI{25{`1&^Bco3IuTsT
z{IAck+oHH+SLP6Q8+>JU{?E18@G_Tt;qT)Tt8>{$;|DSG-Qu!c%cHT|lDXss-yUGt
z2V9Eg-~QNa8C*8?8DHW6F8O^42gBxb$r&xZ81@~P4c{kW*biLthjd-cZ`E9;{DUqG
z+rniFs$XH)4oF@2JBIzrh3^aXFu4O<vhkl07<Q6N+1nxs<!LTC@Z295U-_?O@oOKA
z`ygZQ-`LYq|32^)@UdRP7Y=>9vFB~jeempRo#Z)wA3RX_8+*`a-Ul1{6R_qp_dMWF
zOBM}57r-wMOR(-#<03FrzJxuHzgz^<ZQWS+lyV7#Zk05C2QP!wO{XP|$rbR&lWa-X
za0M6}eZz(+S3%xj4kp3(8u;@571jx8Bm(`*d`V-H2)<J9NxFucVDb;iY!SD>y{SA*
z!ok~s`tTAq)JOvB=OdV!lSv>osRe66I`06(KuJ?@KN(QB<zT%;1QV>@TaDcY?I6;K
z$7a+22k868VndAwVC2<X3^4fsJZ_P+AR`Y!#?X6A`7e*a_>uk}#XVr|V-T?Y8m2{N
z26)zS4b%Ak6VShsjh(%p3&zg(V_KxV1UeVKVN;~M0%#Hs>n$Vlz*eJPtb2;c2ayB)
z*liI7K+CmJ(n1%4@&1pJrnU%#Z0BNX-hT%UZ@q_kGqVDu9j(UJCAsgx9|{k#g-d58
z_+@LJgzpCsk~@qso2&xE#+R_$G-`m?{HquT{U#9EE_sFF+YF5VkTi)AEuio~Bi3GJ
zwt`oGAd)3Y8`zTn7UMhF4rqcQjITZ$6e^@k*0f*1Uj@V16!-gp&PEQV#h0%jvP06S
zY778%>sE|T&>$#a=3r~(%poxN!!0cOCWnEr`V}_YK>;`t)$x<q)c2>rp(}`FK_mqF
z!#S8G4~js3d#~gL1Om9{Ph+!*KQ!!Ed9GCPzU`b}-DG}nseQ97w|36PC>aG$Qi;+#
zTe&SsGV}dxOI9}3%IW3HEC^<lD1WgfZf}!WNUbeVovbAqPs@-_kxJF%?c`0BXDzy5
zTl%9COWwj@*5VjO>1qwTxmIqomfWl@U8~QUYZEm~F_l!NvD<E*ebTI@|JauO;>?=o
zls{|PQ%2c(U%MY%+h#3)SzD$V!TNzRJxi&W^j7P<or0J0>=hNZZ#O5h6nqS3uVgXa
zYHO`d-*M#+XU@u|cjt7oi?1Aw;H>JDr|Z8pxq2*-qx@wzecPwvtH*OWDkJCV2E8WN
zPBd{;C*RSHCW^1oCOKrn-19r+b~64{`m82zd4A`zH;i-opVdhh&YP{>8FSJ3^N&gu
z=XdG8iMbl_`6sox7k2O385^7Ud9{Y+h2QqSiH*<wyhiK7g*}INUcc4!d98lMg}uk$
zTu+|-yl%(bi?-)>#{H$#qp{oa;=cGdaq0R!Kigio=#X|&U42geXY-AbT~?IoZ_8iH
z?$SD~zQ<D`M5%Q3u1%Lb_joDul-`-`+I&5AkN5hJ6(0`m+H$*VkB<RwMeUhg+JCEC
z`|Sx?S)aB`=b@*yzZ-95Q|T_<=c(2KM?zM$ecq*)*JT|T#aq=WYoTAFzBl-Ci1O#v
z7F#Pl_l6|#l)so+Y^zJ%8}=wfW#Ev-_SUYw;rTq3kuw$s9CaJ&hY;1VGz-IiPn*a#
zp6X<&g%K~+CVC)*+?}bAzVpTd=bm4jD=zNOesd!uqUYCxb1yl+-5LKZv1h%n<)s6k
z-o$6;_G}2caLKKAXTqzd9?j358oOnT1*{s{WE~luChm|~TzgX9L2b5*yvf{swSOwr
zs}T+5cP@&pJ)_~E{)4-`>57lF=k)8<7e>pQsV`=o-|g_@k~{L|Kks8*bgut#d4c?{
zO|h)Yz79XFYCr$x=}uW$6B2eX2&U5uyMQd*`qF7!3K+lyMTsjWEHdNEGvQV}1<@i<
zZ2TDo_anDE1lTF_QF8Z+^oqvAmLin?YMHg&2hnB-E3=s{g$D{qUBPPvWSe1>yd#9%
z)kbs?qV(}OeOU`Va}czvt##yy^_FxLsM)ksd55KOQSuiZvga|wKmkgp>W(IF^fpA$
zV!Dk{)F_LNf)SG_Sqfue7fN>6b{UOmtmLD_2166^AwoUFi>d9Vp8`-|D_~yps!wQ!
z(t^%9>EsNLqr{iJ4NKzuJw#|Ku|MXxr=bH1YJ*pyUM~&mAZ*dlU|qJ$7)q=I_8+TU
z1=FZ!S=ExYVN3N;pbx0BPI=>%DEaqg!~QW@72i;z-Lm1Fs)`%a=&5qfXeiY{6$M%@
zE_9a5+cIdkI%6v18ajd!Rm~!A-V4|=jmj8wErQuS2L;Z}PtAvyX8@EuwB0P^<Az8!
zN{^;_x@CZmd{id8J;t$uIVb{3JmahWQIDu7d4KZVNb*r(J4#>NAvE+59OI*14c1NK
zLj-~bNk=QS3PMowu_Yg@we*VHP~x*)J*qx+icnfrquqFzbD;=~m~XB^GvW6n@^OY!
z^RiH@R+Nr3Ql`_(b0*N@;oz1tLe+aB@Xpa7iIue64<)xW%eIOSl<2t3D{?i<bpkaZ
z*2Nt0+!iMSgi8A!i>yOlDA~*{JrHRKXhMl9W)_)q`bZO~I^RWW|9E1w2>6<|XFJgf
z-62fTy8U^*c_T_}U)}I5*)nGwHSww1Dn3MD(k8dux^TS<O8(5MGNJW$HK4?tVe6tp
zfga<iOw@SxxWy_*5zwk>s796V*`wr-N(&>OiHhq{ddq4Ht)9ANP@1u|f%t&Zc>AM-
z4mvg{`EZ@AkN7}|hN#t&dS1~OYBDHjHLAG1K?FKof@S%OYAj(O#S}PB_ga5M>G^9-
zKk+S3jG;3064s$$ePt0y=^`1GEQ&Hm$tQ#x8VeHywJ0%jkCn2Sy$*yCt@RHTFCs#2
zCQG!%2Mlu2a!*vQuLz~9GW%9Owj3QsPti<8a>At(LeS~NSZ?1qAw<Z#LKeC&Uej8L
z(r<({7wql7GmOgU)J@9uIP?gCZ&z-$oz@3FLhjygZ9(m*DM0DR2hFwILR^Q@;(DtB
z@d06dAgzpbkB5*ar;hd+udm2Q>59{%Z41~dpfr!-RiUOHr9zNvk`(2{xG{{7_ZI2&
zz2cYTq4c0Qi_o`yO+#pxTzHIom+lK8_~KRPV72*5KSDluOV2;lsN@w&tWEQJ!DS_i
zVPT8hi*?h4pmx+vHk@{qi;%yqT3WensN^L|e;PmNGG^~Ggr0J<TXnd^ElvnVtaEo!
z-S2P^@{SeDy9`C;xhV0LM|sTAfR#gN@foXJ@c|9Oc3c*gn8ilO)j6`R9qu(bDE-;n
z5aVf=wn4OOndO!z#rk1FaNZAHQ!e<_3Z=C!>v|>EoQ)E*l5AEOHzyCGsfGcI%Iw;`
zg@DjV$}n-8(~OYE+=y1pAx;)b98Dj6&hH6^u>Jm7;sbIE9P!}ZHfTV|pMV+VnF8+x
zN_U*^88!%XAWTm5@iEA>5CV;=h|66!qF4y|55_Xjr4FVLb|-0EcZeexK-FzbDaq!A
z210PKY-u1^URH&W+oIij?D8F-qx2E7YkPZ5;Q;E{$bNPvOte7=Y#kW|0qD|7sA-HQ
zf6~DJ8A@ctnQNt+#6cL@%0qkzftH1vW|%=>IYJJld7z_J#ZO^<8*<k*;=4c?82lJt
z!d8T??-A_}POdM3K1CcNSNODMqIBmh!LjrJtpW6uoRzkNXN(+_)|f<REj21e$P-Li
zzW6}tODIk{173VCYGT$<Ki;!%S^%`#lRw#oJc)+fC<Qjn@%9jQz|s1qN^>C>J@wW-
z*rs{QkN}WOUm8&p*GD4cw<5afD%Fr=l$fCCvfaHR9>V0@7ughV`78iIZ0BB9_~USd
zY-{*gIDNJ54oaNfFPg66d2ms6?P%8_uH7d%Y1rb>0G>`5LjL?Mx?i_<JP9R=4E<V@
z8}+zo*G5a%^{%uJ0+1OQeZ0A&EEtAm`~n|_vMYoI$=PMg@km@$rrpVo@7+`=0M3rD
zP)|emK!kjBalYlq+Tz<VPt(Fyn3c5jqB8597U{JbX9>Xh(RfW89^W4ZlKWCmSGm4h
zu&AD`(YsOpq8BZ`XIU>k1VE#b)X`dh#}6SNW?W$@zOaL^{mB+LEKpi6np$nYsK-1k
zSpf9y83CeCyL@1t-d3<PSQC8{K4pf>#-4CpdQp?L_BZ!<Z@wV_P3CP^oY)gy2${xv
zY;WY$mxvN|)|!Y9$c-(slZzJ8p+VRy*P8Gidm^xs)2y#sI6>G7Gjpvu6(k7ju!-S&
zjnD+(pgG?-y2*!vkiW1*+~nc11e9JFPo-u=xAmZ@7RCE=!bZXb;GMa_LeY)&?g;r~
zXrQH*H9sDuj~=fCSJ}BeXmM}!wsV!jBLXnuvUp5<AmoUZRd%(Geh}8~RBWIaaH9t`
zk+WW8QM~Dp0I2m%U3Jkebb?ubXPw_ZV*3r0PC7ig{IS<DIBA*vIE@$VBmleHPoQ4B
ztBx?<LNki5Q|xY_#Ld@O&GjY5J*c`|PuQ;F{Jl^Gs|dcO-eG%$+{<X_r@EELp~Nk5
z?sw)`t3cQ;%T3}#00@@*WC7pU4k7=}7^WzWiLRr>r`A<jMZ&4isLa}$Pfz@rhR~-b
zNu>?St8EbSk!=%p2Udqen4wW)@D;n}&#33M!6;)ZJsk*hYT9MNl(U9m$#*l4<POKe
zcBL(H({N?ZXS8d*Ws~@TlZKBt`GgN!B4l4jMbBho|5#W<sv5S3`^SGqPp!6$6(2Ac
z0u4L-XdM;^c~23jSiFuKgVIH7d+M?R!XOMZj9L#nC_|q%CmVPc)R-gW-I-xRQGyqQ
zef-c|Yf$F^Rgg1{G4s??6o5`Nu{2zz+!SX0<b+pyTo(hSk2$v2F<rDFY>w$-Ideq?
z0m${DtPf8tHG#RnkTq3Q^kl#qB5SU7#zYaqX8Jb)NC{s5{TIS6!rR3=U}vXX;H16o
z!mDiI&FXNsWYga<%xbD)M@1}wvpl*-bw1x<wFIgg-srMnuJSdoJ%_z@f%yy+y-w?j
zB~fBU!NEuE9cG}<#;aJGb3X+JD4q?xX)yy$8-2c35Zf88s?yFkuW_A$D)rcJ)1JH1
zojaBtR-Z5fJriK}$mV)ynATmc2eA%^M%`t>f>v<N&f_jsrCi`C(WiLjSUG5lrWuJy
zNLM^r!j+U{j!YBlYb+T$qttpS^hQS=80pi0O9|H*yp2P(9iv3_Bueo9D^`nM;83bz
zeKR;2a!Vm{OeG)*hc<NOiRd?{t8Qobdz9hOzWwR#U^hFnEjVjK-OU*&gIF^uB0iny
zyPXhtBL|0$@^!mFXM3ooQ$iZ=&J5Jn(SIdEkDOB#{|H^ufkT_0E#-jB8d`NZEg=Ml
zlC)7hgg9+b;lvbZx8cx=h(Qi`7giTop`~XTi%WI)c0`E!XPfOC9jB6{(5S*5@R}O$
zE-I3Z#MP?hQ;s3T$fW_->-~6AwWuSPxq#%)ETjtdQ>6ryc{GIR^4L0SjN{sfOZ7Rn
zz7KR#_>H4lHTJkvTaTI>2=Q35MZ}@;mMR=-qkW|hkQ^P3gb0k-xK!I7)@=yYq{cRQ
zMYQ8kPW?b1$mBKj^sTaPyfvfN$d!_(2(dlK)3_{p)C`9X9<b<##;RW$7E@H2FaxbE
z^vXf#byq!Il1;MNIJCGetsji|D|rXHwDsXJFw)4&M~GW`sM}<P9#XZ6DWZOm=zhDc
zi!#3*=US^1_7Ncl8?<I?mMjj(rFse+2EmuAn;t@Ug%YXMHQYvoZlBXbIme?^<ItYu
zdxPMlxl4J4R=pn%&E#9Q!x$)XS;@EKNi8rKDuPnGMnuBn2CX=BgirZ~5Q8?Bgekg)
z(Qv6e%HCn1#y529Z=ExWb0r1%@e%qR?O{EKvL!w^)IKhG7%&FIr%RT*+u+{HyyiWP
z5M%SJ8aj$|ig0Mz_~0<u!EbGgPP#UXhqkk0B?<OpJ$<==b^9s~<=Lx_f=t1!b;UGp
zF-{;NOtcs!Qr`AuoB7{5gG1eo%}2pO5CxpJ25@lbNSWVKl(>!TbEhzD>g^2lOjw7{
zC>W_+<ukn5xLnGW!di|J-P`z!RO}+831hM4DkQ+kGb=$RggBH$;jKXF#$y3FYX;+_
zQfpLSjDpv7h51e=_DEfj*l44S(h~#Fk>zDh(&|$7v0)S>QVXusUFSCAT!TLQsKT}0
zWuH6dR<1bojJ*~gs%5??eLQbj4IZzP*CMr0V)PBKPmihbX}AU0`pP^|8@+16zHPY;
zUbD2OD>kF_A{)!xMw?t|3>1qDc;L0Au1#=wfh(?7qGp3OO59Y$E%;Oxegl_U1KfFl
zRBoT<(#UMZp;~6Gy0CA%*UYS{3B$|V8OPda9$*xQh4#NR^TxSGn78YrM5FNrfsu_!
zYS-p0q!Anp*`6E?G>360iOIVzBHk_%?ziqA{eTnbTLzlJ4nN^l|5XYoPJp32d|O0p
z0_IJRDn7>HP{B(C)>m#^KxFpYnj142%Rso#MZ{1A&xr)Tjw~E%XHRPdwbnX&2V2tt
zPJm$uOED3h!&la<Oa7j*EKw^+v5&cHF*-ju5tsV5u}(zXrxTdjW<$P<OSROFZU;dj
z@tV;cOsTg>SE^b?#OJcqk!8Gyd>lGToo)yE?R-5!OrQu?i$p8tiHKAKuX=~7<xg;^
z<+E@$7#Vh{b}cX}#xs_5HAE;PCTR|~nw!*1vw)R1!3Ly|tsR4_GQx1Fj8CjZ2=TUR
zRdj}B>}OnRUyDH(2;w;_+by1OVd82Xu2_l?`BAO5=3Ta>IFw}^)df<9HwX5p$nkL~
zDG;nch;DkKj{%M<*Klax?VK(kUzsotnnXRxGg3(dAu0%QWm=`ZPkW&VhmLkkbb&7{
z*4rxma+XwTjn5{8uFdx{dR$REhC``Jx*Txc#rQa9OM9JEs<NvNLbpBZwbH}NZWxDJ
z>W6W_j{cV6K1DtUuP%%ttE~vVewl#x$X$ryP*1}b9N_EUdR3To*&dh5C=?nX#KQIc
zi}rw(Qm)N<!yNEBI&@dDdvc)EuJikC5qbgqfbgW8I*m&mU8dOs-cdWQm^aePaJ3kB
z8yyg$mRoe)QkL!a$41MM9zYo69V)XsT8BePpu`y=E-{b3y})vDJucNH_FfOz5jd76
zsFw}H)gmPXxFN)+={0(#jm^@hAyuWf2b>h{GxG~tTPfw*>llO(H~WA)Yx+F!>LM$+
zy_ySl*oAdetvc=~wZLIf2tvHWXB|jxxGshIi9EPK%RAnOtFszNxhitP5jxLCXc*p8
zFO4iuyE|OiIxK6nj7)HsO8s0Gi4g7Gs<xIk+e`Zazwkyb$ZcKO=xVFU!=*CneE>q#
z?y#&6w1|^>i&do12MBe+*@MRxl;KiI!zD2YJ<d$1b)|oaf^epL^UGh?2VZJ^Tsx@W
z7MO_#AqcsS5S3|uAz4+K$~ctubXOlx8%_-9362ZKOEDwC6LuMK=_T9tyE{pXrXTR`
z1G_`?&jwD<FUO@abZrt4qVrx)^5w|e(uDDIJktk)xQzwgv=0<qEs~Qk5h0$)C|@pT
zn*7~cOPG*=#m7J|HpNuxEq>1}gr1OQ@)5MYB;#C*Ul#NMh9$Ke7|HcZr7E+M5aRMr
zj`r5g6I?u8s7`Ht;C0#a@t#)l_dY9|nv4){&Tkd&4-QwsrCPR6^nnz?9Yxno$rVxp
zRi2p$v0%L7*aRngJ`VLfu(Thfl&lOer_2w>jYV>|e1_2X^>}Spu(%_OLq`iY_JdBl
z?QPV=`YzmC49f6xgc#+>9j)pcktR&v%Dw$S%cnirkLM7A+m+N-mjxqBp({tHl{E*K
zYT|LEAM9WouBkB^8<$q?vT)dv89Go@+If*uV|gZ?>jwnZaW`uu%v|~`AiKUqh^Tw<
z4Ps!ODlV1!5_WI;etB$mGX+Qq3<_T%^b<#TZ60icr8u;LRR}Fm##4-Df3U#SBCTo9
zM~HdeWiD45o-f0p)Y>*^fp#V1;uCW!rO@)oLWJ1VZt~RIa_e`svL~Qg70(AnFUtd@
z&`&P4unsJ+u)9_nJApgaE>;mXdSN-?kqzuooGZhc&w|xO+xtWg+erE}>~qi@1Sy_*
z?)KZ)SK{7ctgov>h)Hr)v4y;;3|#7{$YKx>#%=ew)TMn-!=nKW2yy9)5Y;_lW76dI
zyW>3wb__->C=a_MEg7Wc7L5q8xy@Qtp;=#AicPkFK@il;);3{iet*KOVK*Vf&}G4H
z+Le0JN;sOFGzfOv<yE0d)fP|jjJ*-of)MRI%$eTfFQs-Bkg^8>scJHivv?OvYOKRn
zZ3r>mFkDr)&Q+R*mM%4e;C0uC?qGzDwACdgu-oCfz`dM3JZw;dTfik@09qg{et(^V
zax`u%()tiKLVR&NOhw1jr4WaD=FJ%fuSY|sqmS9x<6Ifa=A#IaVOA5`SUy>VL;V_5
zhXEln^zLXvV*}ngkd8FKUOOh<^tf_(p)?Cz8uVbV9eiSG!wNTRTrI{j?=hG#)x)pi
z(d=<t>gX#6*g&)5vdZj+z7rtx#u1_|+Ssh;6+47Osqq1j0GeGN{YtMH?+F>_03SZ&
zRUOCD?SmHL@!ITk4yqO1g81)S7$B`Kd#nUd>NWS6tK%=jaH-7*Z1y|r3@%%{ZxrEB
z#^Hu3gx)~2kGm_FR>Y+S-_98Z=LbnP7A#{w+*l;G`!qrn&2MGexRZ%Ev`?je7@Tyi
zs|w6C>z5iUt4D+oM=txl*yG+Tb!^{_fngBT5O%rJVUCYf>SuoxAx@^bZyn=Pp5R=&
zsIstwuiF}C`s{3t6sk)lp+u$j=$2h}G18`#dV9erd>l_}+iiA`_OFb!u8UBj0ee(;
zV_&AUAsNk49tFPkI*3bByWcasc`6e^X#MBG+N`=QQg68wtsMm??b#g^T@`D*W|0=Q
z!!?X^qbaMp%Rst{^JHnmg-Kt78Q(5vP@3Fw(M#ZRr<XfjSU07M8_TZHXcUlqNDg(A
zK2A6R(qmUel(^NPbz8l~CuxE2%i1#vG)i7wYZw!S<4}f3por41EBbUSYYsNXrTQiA
z9|iIOp?9ob>IL9XlBf|bqP8Cji7m2kS&l<lVTVRRkX^i<sd14V4rQc=Eko(`%*w5x
z>X`u!?J_(Hp(WuqEUh>j97;m%;ChRF#oH*0(j<KxH<O}90i*GbaU)<iN}+PY5bAcd
z)!BfPC9Q<Tg=e9+titPq6U>10xs_9)1XmmF<1aq2mPjMZGW^;oP%<A)t8%I?dyac+
zx%Uc`cp$)kQ9O96hjXo1d1n;t@EI%ZJKkV|yMVL~h0uC-$u_-lD!$;89nDQ21zOe<
zD%J^ydZf^v#+4|MWKptxVHo=p9)y<r;M#DcJiEwcpPWDn-3vbn5(~3UDSM)G2XLs}
zr-D%s(SPHf;26<TYJtO+$|x~XxmjmVu#?oWc8=xH0{&am1da>6@PuJ(=z>s><AE>F
zjkZgxi%C&Egqkn98_5o>l9so$a21r;WEh>S!>`rEEzq~61IDWbYr;bND8N@6jJMWs
zacZRMJ*61VlvbCD>s$yOZftk)Nc%1|iKB|r^Ue<5a<hjYEoNFW&yKNCklSvwhgw(e
zB~?q6N=Atv4V*SP*ex5xEg;x}z|aO0=4?miC=O-VSi>c6WV+)fQD2Mn!O03G@?g9&
z6)UC!i*d)2bh@-qx=T%a)}7}1H8_FbOhq2ZtWe5ta37cE7O6aB6G{x#DMxw&a-{8}
zU84$I9oy9vmf0ghX~~#xu^A=aDG-Q;g<*!cRKet09{iB9d1<@Tx^kQ=NwyiTxv8te
zKYCCNq;?%$s?7tPKB6@|?Pw>wz>^;NZ9$2)kII)1_+@(I78q4AgwPP}D<%_1hw*S>
z$O<9U=<fLB^2j3T!o+0LE*?naZ(15GeA_A|u%HsI<d3-tr|#LXH{x9V$TmD+Tanmk
zv1IA6RIMr(9h7+UZZJ73c!?1X?Nf4t7U0z@yDF>JNVyvILFhMN^~sQX>w6lC4)MT_
z@?`;$MM->VWE~2J?+4m!+6|wG@}*LP+XHwY*OGVK%Q4Oh_ZGvzTo0wM>}{@9@QPiB
z6R5~N&I9Mo8D}{AmwzuA<t#n;f+B60xzIFbEe<sSfCpcy-eg+xK6&C&NeUGB)*<S$
zW%QikCOI71XL}i{W#PiN<_|hbO9t8pq2H`)+Q+?ae=qQ&I3Ab`i$zBl_k0h+$6;Gh
zV!X=u;vv&Esbhn)??3|0`?5=FuT|hrL()Q%ZSV`tTOP%n!;}^x&jaZ^K%lw|!Ddt1
z4=~nMLg=0KrmY)-S4it?u^@v7l+c{ZA<X&SxMLZ7`|T(_QqEn!yen%pu9jz94i9d%
z*JAg(Y-o^%3uxbt67v)KrabJmqz_K&%2FOU$fFvSZ;Fz-fK=pd0LxpozrICftW>H?
z(g&C@b`6IpIn3_~lQRgR=<z@#kQciECqP}%zytT2^~Q?r3n)^qK4FIN`<Zq7a=pRD
z(#Nr9Y&(2M<hNC>Ik=pPdy7#ogiyu}-U&C$acO;Rp6G#g4P*5BCwT~QsicWwBj^Gf
z%Y8mwRB0M^?H}TScObZpqP<a?1&js@W0a`PjLchW`dqr=EPgh@18R1;pDJv@_hsbK
zDhOrV6~&9{l-#6gC_s2XBTz)QFU)V25-7E^MTtpe=z)0a%<uKpkT?$Xy=<#1J}w)U
zLhBnKR5y#I4ZoN0rBWjoj)Nf2q}64@>hCT{cCkZ=PBc+g5x6BS@D<t1$3bm#*xg`O
zTf5Y;Yuh0-cUSSPMhj!<8pfoZJPtCWm)<SY?P!qNHGr}YCF<P~>c6$i{9Xspb>o2H
zbs!o{JE2l&VAnpB9wNs^_yPlM+_4oZ+T(y=H=g6(%q_zchNR<V4~@m8ygXbsnt-dt
z(lZzblY$i;Rzb;ixCIzl90*Oc7#?*K#J#|w&4y;<U`Lf*W65ZN^mP(r+{*!eb8EDp
zP<9ECmf~jPJ>x)&6RJ?Idzpn>K>1EH{Ae{@$Wv7==C{pAHM7{hTmWo2jf0*!hpg~7
zByqcJVo_qzf^yP$)i|C7XG*d(1%TnS-<#r)M8Rv8@=*D8l)mbCG3l-s^C~XYCJ25V
ztR3Ya_lBSSX)~>Z^Br^eQF4)-t1#Tj1%E$KTI}ipzl9EmcZ}coWP?kcbT5aWIrZCn
zo_F-H@kXzdns^+3{-drc^$GLn!`t#&L+eD8*lO5br(6C}%C&C!H37H}f_z4moN92X
zxw{(S_uu4(Mu$LA9G+XH!_DIF$>W>d%=w~nywhwt=Y1QcpV(v4bj~y}2Y10JyA*!N
zcAe7n&sUMY0(VmoB%wt85BAT|04@3qw3NE71AgaoTxsK<yuSwT2a?)r;1~JLY|jD_
z<@P;XYDfzoe!n$s*DcFjV~TTCnjZi^u(=<$n|7p9!3^|FN$pZ0I2pa&yI6Or#q}Ax
zy4|&Y2EX$!h+cdoT3Ry98bur7Mv5LgL-(g|t8ldpvUtx?VlcBV%Q4_Y_KZ}s*4Vv5
z5aGGB$YT1DkCdwgI}7eFIXs*+<kKRBf}jAni^J!bTV3WkKb&h^pm<lvO?ID3y63A5
z+ydhXmxVxyMW!_?I@QI@aJ4lkgWFLWY2FTN%U+zr)#9yq1i4n;SRGQxrQigVvjE%&
z^YN@5`-*i#Jq``zmkWXXaAF5n^wu5Etru{gC`t@F&e!YZD&jAzOT9J?K&WLFC0xIB
z91rcJ@v<V6xOBigR=dUscdS`eu7U_|JqWqt9Z~IvORaBXu~52UoJC?mMI0W4vi`|%
zZ;~y?DX>-Lkt5zdIy~Xj!9sMX{bO~c4ql2&qj`HpK%dIkMG?Gp#b2g4oE0{}7>I74
zyx~(QUCCSD3>AU<Uelo_D-aeg)xgK434Q_(aJ=otE5Heqx@X3VfSOmVW>3~-3J!e{
z5ZwYxh8t*Ni+GQ5$BI@Ih`?^H9i=rn!yi}6?Jl<srDvjH6kx@~6Q(qr&xTT)ogRnR
zH;Qnn4rCX&_06bWm>B0ixfYko-bzG4ri;=V>aG3Kmw$<4HKQm!{2VLPu`lO)b@`@=
zf?NS(SzXMuJ032Xnl@u7eV<WZ1f^^_{=veme328}c~?BK)@E~ow7QVEbrUFkS*czA
zxt_M<83Lu$n*b8%Y3dyh<3{6-jo9cZK#BVYcyS`2i@$X^Bf9kn1z)J-<y9(c>u|Lu
zRmO#I4<4ZMF9)!sTtk+BLc#9Q=RR=L#~^O3+R_>XC4yz%dM`@Gq$_#!WfrW2mW*u6
zq(*x@w{mMNLS^UL?&an&OY#cvcn#k?$DW~Y5uq1;XMKsZeKc@vpEDOe3yx_IGVtEz
zj8*1cKgN#W=rPvzh0!@P3A3_c(1CbCkzeOBr-FHB<_VJ)>Aa_yk>dD#U8VA8R^AL$
zd(gf>VeW}fu1!^b6OU$~rT(J>jSP({!$Z*_ORZ~Xpx=_sbr;NCwHnYqxVSrsB&&_p
z-f=d=47i4&vBUN_v2dx&qnQim8Wi%hr%P<_%|OipR+XB^bcPhUvqF{on`WS!GeTA@
zookzK0yieGpUgnhdX$aaW1J&nh|M6qpEm;?8|=BhY_3*TSm8AkRX}k9TgpOW?w7=S
zhdX^Tot%L#UPBQrpQ|cr+#GIt%X<d;xxC$>GUjAh%u>hv8-v0b=stzU5S6)4dwq~O
zEPFl<wdok>iaB2s|9HHygR3~R*P$b$zeRO!$DJMq+GuBBBH2MvDrmVZh^gg_#hd4Y
z<`py0Hqf|l^;}_j_@{QZYSj$X%spD$F1D6M*-+-@X0AE|-KJxow`T4~L;uYDfagQF
zRHwy7KC!j_TbA<lYeYITP^~9n-_~6SE)Tx3p}6kF%x(`e|K*w@u%j;Lkof^oGye5_
zu_!n~_5yOlekr@T0e|;q=6yiHp0T6ZAYKq#HjYP@?tnXtEC&^vWwPUSjWa>G8qJ20
z_1fjRa?J^ZpEYAFBd$2ILeo7Ru2OkdXToK5r6jU)?^2(tM!uYZ9>5~&+2azWA?nmU
zyv+#^N0u`uA%hk7!U>P8F)Xs0xc=G!_8lK^sglU*a@7{ZJ+}KeBUM`*Sr?|vyC%}Y
z-2X>pg*}^Yuzm4ABWtb6^YxU`e@9kgdsa5P>3>F6($T??FjLa}e?=C<GNRGO`F}>%
z{)MO^_+OD_7q`sye??ZSCVR2m|BS3J=6X;5cVuO23;uUxsr3v0OJr?n_%D$a)b-yY
zOV#MVMApcEi>$fF{vBDe?pd<1A@NIGSaH*Vi<jd1yN<9YjJ1IuV?7Vxou-)yItOm+
zLrLM``#oy#N|<Qi51}Jz?Kcu!#?RoTc-)ByH%HZJBQ@+7U3h&Bg?rLbx}DB=G^nBC
zP+1WZZjkb=q29IMr_05QW~gBU10}Ac^h{Z>qp#so9r^R&2Cm{;SKFP+r8_8I-0OqT
zXAO~V+0mm;Gf>$w!-a60pK#SV8*s)FSIdBiLIR-$gDuK|q4_gVGt296BNwShZ>fc$
zo3za_aI=RyzIgS~OPnekqxduLU3waq!d-!VIlX}_rgUFMZn6)A=9Squ<aozE!U;65
zf(^7%JD|8CXS;A?<$sEdL5VjUx~`X1gh=0%ay{UNvyoAi_=d$4DS-%C7K94s`yVrC
ze2%-|jLlOOxLJ?&bgfk~{690zB+Ll82NF<4YPJL3^kJMp!#7pfK!@e7t+~0uPTD-_
z2q0ANkVRy&4@=s;g>P96y=Bs9J?>agi9_`btYNc15-`|U#x216tkUq1H3D$5FLLkb
zzI5pZ23s9Z2o)5%B)hTeq;_q7wiY(~tqE*yBT<Mu)_FC23qTJ)Ys$atIf3^svaX}+
zAau|<9&EWQ^;XoJ3JBe_YnXJeBB23S3sv0!Ezm$)UmiVJjyL=DZLN?%w5-dE8^W#9
zC%3s9T$t?cahRw}a__>`a!VKQAEXS9E4taYhR;YX4NTk!!^LSN+M?yoINp+JA2z)X
z-xYJKRNT8*c=vV&t<x3&c?z$l=M_=9HO5WZ2|{(WM_>3fOSf>8jw5jUVlF>Qr*(>3
zhI?zW)DJ?RMRBG+yVv7=mYMxlJ-FJ4-hZw#v4)3-i{hbZ2yNq2lcPsza<~PWj~T#q
zK{!)WG`T(+hh~n~K<I?BwL!gi2yQI1^4o?GS|zM5X(+J9v!E@S1D9se?a_sbCb>AY
zlwD^e01-XZrOmZ%k@yoPe|j8BU9zWUGG6o)Uz(lq5*ot|l6_kXhQDR=q;@@O83#Mf
zVH<rp6A#T9u4eX|&Ea}WG?rnhM3b&zazD93Xx(9p13q^4QmKK?aBaBT&*^yM)aOxL
ztudWI2wfp-=fvj4XW`xwL|eez>O!v!=g_6y8|iW}godP*ZJ*%Wz<WYj-m^V0wB2Ld
zLJDs9;}%eCsEdR9smqJ&?~L!if^&6kw1#`2&3R>BZb${Lmf}(_getayI~)9i@yI&E
zo7xLk$J|WafRJ(y&Q)uV077+lc{L;#U-y}jdPbz;0PQMpX>*S{%*C5hhvQZ^;6`{O
z-gKb-2|T&YsJ4!<y71!7!!|FIcA7~E?r?RSr_-vR9Jn1{n4IxBwjai;HDkbjWK3!-
z<rjfaYTSO)SJ@4>amSvixb6g3$0EU68|{r@xCKUvilI~w1J}0IB^`LlD5bn`gJ~$*
z=i#(+Oj_W_-qt~>Pd8d`Io!NYD%HQ<9q!CEfB8Akc&HAyfOA<7l=}4SxajP71K#SI
z4f5bhzO$KTUWY8V!0oD>!-rB;jrwkXs$5czyP#+BA(#awE3S^)s`^PwMw(?j+{%Rt
z$1JRK@pVC|{U&c1+7&B%>^EgdR~w3FT_JSK3v1-F&;fq|d`4(=7<#KR%etQ(CXH8H
z_W<}R;g+uZ#{m;kI_@o3Hy<c<w7rO>UoQO&qPPJ-=z;KZeUHASc+D~!k2nH@aFBnF
zr&TE37!om70|_j9E4;PPJ3_iE!SY4`G?rDeii-%9?oeYq>4MO90RMvpmuGl&F>A~`
z3d5y6X=DH7Rq4ve_An1ZH#PWVIdLNJQf$_h7YO6EBpwNDSuU*uwTk8m@O6?K3tdMQ
zNvn&`r%+gk8Wu&_7_Dx=y``Ar2%+RCYRj%L<m!wCO8rK{pjsBua$zyvU3k1IJNhJ`
z^sPkEvjG2Od~sUZCxGu=zJx4WH=Y<3j7wGGhC`|7;nwH7T$~edsocd;unvSD(2R8H
z5aB^M7Fq$NMwmO@(i;v6#bdyGRWx*gb?_q-Efp!(7Zc48y6Le?()zFijX2bE3rzqL
zIgvxAY!6RdYG`Obgu0sd<*psqmU^pw8xVj=cgjH6^!Htg+Py*u75<;bt^}@$YwHgZ
z7JV@U1w;i6ZlGle39Bpu2?A9h2q+XVgpeSWkU&^P5haRMXh8v~qVoE*+V6>i0%F~u
z0hbE0C<F=!R?+$d0ReGA`Obvg)b{J|dmVq1bLXG`oH_U0bMM?cWIXvHBhP%(l%d_}
zCxIL1vxiM-ao2aEAe{8YA14)81da^_+M<OtVN=i^=rB^<v;B+5mDTF53<GMJd2)b2
zdL7*jC4rmwzJ(uKl=n`EpDCY^bZZHL+DqcjX&SOdK}e8^v!KT^KISkghm?1iN#4#P
z2tplkn9iMC<<y}oB^$<pIQxR8AV=kx6;#m#TFckek1V=7j2_iWI6HE{S%QR$9idMW
zQ8yfZ-3wanN&O$1db^eLtgU%pgVq<$XT#{Phmc<@tVcj=*mc>J^hRyvsQo@49;g+`
z#vFJ0nfsz!Tyu;oJW%UA=38}@D^td+Z&5xxCF>`bdVRZAsnx33?^_sOor4edue;c;
zXjsDYBLxsH97CVl_;FVhS#NG|K<h@sp~-=DI_S<V@i<lpO)!4qY=7P*4`i$UlQ__t
z=<Kbd(<4%jE`!Gp!$j|H#rKbob-N;J*<>wHn<HQJ-b+<BtR%Rs7-GP++*f?Wyc0RA
zZ1#Pib_hb^ZAa#uLP6+Vehh}Uw}m9{3zca|tHHd-K&`knnwpfRoDUH0lplw=c2BTg
z@S&U@MAbJO1!~^)o{*;0FVOHtXe>VgW7br>731o%ROEr8d5k0A*X+2IMX8HWH)IBv
zpM=(We1ufrT_HePD@gwE>dUo_?y1*YThOGGDJ(k$)v6684pP@_LRy19azN{mn_X|x
z;|?Qd5qwX457Qj+g{T*APbi;h7C(&xtrstK<t4RULMu)}>ao-C*rWe-omBP>4?TU%
zzf}uV8+EBcRab&itKjHah=K8Dr@<@tgvbD^Q!PMcH@*-1AT>baK;4U?Qiy^6a^Cy$
ziJfR5+MfIfY%O@X>H4{-MNueRCJN3$=N|06%3Uy^oId8R9l<^D%?Bq_nH8GqN*WKI
z2WN5G&4q>RO=y}EKyp1gr7`byRMCgH4a&YXl3oE#z;^KNT4*XpbHd?!zCbN@6#K6U
zRL<V&dbd@=_?lL6??d|UzXE4vv*B6s8)uJ>*`0>U)(T6ig7H;GBXP@ZQ&ES>bru5k
zqO7*ld&&TXHZv;j60El-t~n>uuPMW&XmAG@P)!_J)cYYBHGyO*<SI<bwibrA$Q~pj
zzvdS-g4W^n{@R4m>nO77cpI;Q0i*A-CJqk1K;e>L(*x9vK0U#cQE7-O31MA-3$G%`
ze5L6+eu(P$-~|{k_uSBghII70q6zAH?QcLgj5}7?w&0#J2q*dPfhuzEIW+Xa8Tqwt
zlvW45KujMR+-#xzCDql}UhuSh8++*Ei8g&S6xYcX)<Y9`w11Rs>wSi*wS8wWP!H+&
zRh4=@L>?e`>)eDzl<>w{Us)3e{b_-v77f%nHX%oMcCnQASiec2zVR3h6^A|bLt6DE
z$;I%;z~yiK&%Q5Kw$`)JTX5&*lp92emWvN6)?1Sk^`Lb!CwTF-u6O7TQ|H~)0P%W%
zf%h$1rE<wT8Pf(-D?)!q<+uZK7Qv&n5$*+yslt)_*3Xp&$R2|M*EemwxLtf0IZM(~
zeGmGU-%%gG?Z5xBI>>nxsD6=gi|ZK$=x#_DI^G0L;1`y>E?3rzv|c^#dJNtKPEUPp
z+cgKxUkSbW55QSFZy1jkkGUgi(4_A%c(C3mJfhe0BbvKNINRFr`G8}<&==cfXe2c7
zWrNm(**oXIjMhW8)+MvsAzXMzx}#`|H=!O&Js|+<=Bhgv^%}RMfvAoXAcNaRT8HfM
zPUm$<t7sq%47jwzqcw0O2|25-*RC7pglPup6IUO|kga3pcY(T5(3w&{II<Cq1EU5#
z&{{88RlQvw`w_J|wiBp37C(M>zIUOrwK~ZUVH`;I-@7BRRDi~Tj?>S;*0eA$+C*cp
z84B&dF`%9iY_@rR$=C~Nb$mDuTFaeNbt?Mfl+KEh^@7&4H77&zpDEu2Ad}pW!{6iu
zn?$BN=bl8i`u_L?CQlRhlLS51%7qEl)*t>5G+NviKNMwIp`c3c7Wc!NTfBM4;2XBG
z$Cib0fO`FbxVbW|2lZHHZ~9XhwF4jOJl{O&i)_u`Ap`>sM&6pEF|`fZ%Df|b2J5ZI
zhtIxvZykjOqT#|FK>avy|A=|iLMvo|*M<S;v5n%>O>dtmXYggEcYzwL@hn$uycNw|
zm>l{bKG)WJ=1p&pLPJLPwN9XR#Ew+;B#wt75BTZ4gl@PaZbevEDG!aWxwoD{wdP#!
znh3nlLh;I!jRCdU_o(eUKON<?y!zfaP}9%*#-E=$RI1QgH~wG<*ZT29Se2F^YS_>w
z#tB$)vfUq?7uzU1_h7{sc))#5z{pjPuzW@5E^71x>T73jtCQR{=aAO4{Bh7aRlR;|
z5lQ*GjG%fpP-(|F#}j*ss+7-&#IK>Xgw@G+*k3<Kp<P@h0P1$*=kNADN>+w;|Aq;e
zI`~?F*2wqB1KUrf!CSm^b$9X8MD}UqtXj)SSQ`e~sND(G@I{l-H#v7e>w*Ab&tR}<
zHlmJC0X3IV=5)vMcruCsa)J!BzEgWUsXIE7j;M@>AHaZHoA-KO|Dc>@RQvXW);9^l
z9sD@UEJS^A<s(c=?^r**)lS-k8kTG}3RHuB>Gg?QDd;B86nr}cwob)emBq=LkO7lZ
zQ((Z<^oE+C?oJZn{mzmuy@|JD6UqU;?16XLlt#r_uN@D1Qqi2SuGU`-UZ3X0uX&pn
zm4^H}c{~KD-`?uAob*1S9N+_~@CtRVYx`osiDxq8EOJzI8H}%sr$UX_St%!_(n+9}
zJqwDFPWmW+bk%H{3{>3@jm>Jchmiry#tsd5x15u<=J|1=(yy!mDH!nP>tIWdp^c~)
z5X2QksMflJG3<}aly`1;H<GG+Hxf0hWc)@b{7rsSwfgADxhB*NgQo7b!Ly7#Td9kO
z-kp8B%agcg@p3P(brBm56^|`5argS%|GNzZXA_p2^1RlC{JP=r)v@K&Jznd%OZ>lW
zNLXQY%FBCeg#Xd*u@!dpUOwXQ{ErPJ7(4cR`KJBqe`0dX*hTGgR_2m`e-aZ-n0a4~
z<e9je?f1)xfABAg?n*E3{Tm7{JUB%ySm`a;?|<yogVQd$tNe2I2b`j{{lF?%6<E4I
zur$8y-{HES1vl>BSb3qXR8;U;*x>$6*Iu=q&(JmH%svoQPirsFEimPo9SClWZ@*Zg
zYZmQ&;ET2k?UiK(W&**1ke*lVmuhs)<8ls!_S0mSn+wbnOAmw%#>=ks=~|>V9@sp7
zL3VApz+%VX0rrPivKyElb=U0da5Z{IExC}o$1IyOTihX~=vjW{ogJZjv7??^Xqh9(
z<{FN6+;Y*Q<>h2YE~j_iVHMH}O0#)p;?BF_dh~A_v-#E+I~zrX^kakBQBI?sO&NMt
z|D2t(#hw03OKzdnDYKktZ}BhpOZ2RN@Xq-%@Zv9RWrfzIf}E{kqrb>%^lZv=a$<P&
zuFmE{o66E0fk52V)u(5BxiKd;@nTocaG|Z6w$ApGm#&SJ3e(>&?3-a+U8{XgwAaDl
z*Ey%F>)}zw*J%b%mKa~Y<99B;@QT4x+Ud)Uap%Ox-y8hqX8dDI-noSD9Sr}o;q;Gf
z<>wO5r5O%H7+>jZIhRy*#qgzPJ?)L(%Y;h~56AB%-%#B=EU|U`RktbS7Gv|P{7a5c
zt(#JtGB=N%v~{BRS$MiFn3K0Z^5jy^{5NN|wO?8E`oigJgK_869=>1n<|pIpBYEey
zKXq9AuJ!cw@$&OKUZyS1?VGpEX!cj0?zGmtef8uKzsV=IYu>drrB6uNAIC1O`8eLR
zb4oQFQ?qlKq1T+D!3>|FTkWE5)4WSd5>8ld=Q?X+^X_?d;p*1auB4>q%*6zbrn{Zn
zoWf>_F_SYZu-Z+ps(FuvghS-nF&4Bp@3pPtkP@pIi(WTptsz8^_t-I)=(X(gU`EU-
zsAjIPY5B@m5}|X-&VAL!mi?RRBJ?V&-OZC)4zLMa{d&8#w8EC`C?<D7TlHF-s+Js~
zglo`m=drrI<?G}+?xOK(59ik{-|QqrQq=4{-Sk>>vzU=fbT50ZwQ0@!MiRMfxxLrt
z8(R+^u8Uk@ec8(=srAqa0?)+V{__omt@)>!yj6jhKi^o@T5w*%Gv(Q@`=Y(I@KPPm
zJn{0n&97Sz-yra*d+gUo=-n^6#pKfpF0bd?-2b*o!nZnQ@BQV*`$uGTe4EP4-myvd
zkM<Ix?CR}(;tTH=KVwF%ZoBN0RCWK@5OdDupc?vdXSainDp)xrUAY8FFX);X26c}N
z_&O?1`xtf-X|`Ji<X3}yU~$B!nARYfG?qc46Wt)Qn&G{F8t|!Va$4rX234oq45w^W
zr-qbl$RCyl*|t0(tDMF*k!n&(mWYxyzd@TckV|?*)rzHRJsu_vvq*Guo#8{soM&YC
z6K4^h`jx}gOowKMeTAyy)l5f;szbxDPExo=+7=DlEs!Qn!!{#J1i9rPN1BO7(i=5U
zo1Brwh7@u7S@A_#sYSa?Uk;OA_L0VGNCQ;Wh}(Dqu<0~l2HhNkpB1U^)Z23z_O}`K
zMe@&Tsa-z9Zl9_{zSIIPI=S27fvQ89lz1yTOB9q9FUm~VqLBj0=k&7@O0wqP^VfQb
zX}w~Rp5&4qyJ$V4kOrU@6cfYlTF8aAsMW{-(yDd$;TiJ&qq2yv?`K!WSY4dy(7;-z
zs0ielOVLOv@>@TatA%@D(U<ZZct(Z~YYo#?@RV{5Mux*2^?*H@b`4U8GKLd4Wb!A+
zCk%&ua45~Lp9M9|N(Or%hna~+e9^AZ$7Iq#4XK|+8l^6+ZH2sakLGneGJFUIc+Phy
z%XF&9Tqu#+Wi#v>q>k0K_P3=_N5=-w3VHHEjW|5VUe^4so1}qot-(I^VTx8iRjZ$(
z^_oSBC6gXgNpLvYr~bT;^n^uvLzVv84vxk1w;84zK}`+d^Y$(%%zLVW=l1)VbF-}!
ze4FWzuj(-AW_Md^51-lKXQ^X@)S`~zgtu(Q)kG7GR8dxZIJ6+1n2lYMwLJsc(Qn5K
z7U@+VslSGFgG_oQBHf_iG1EW%4~558QK4OOHInfd*aPt(extVSwnmCc7JFY7<akRe
zQrT0xIzQ8~nqgliwUA^w7E3Lv8Fm#6dn}cLVX+Y7@ryhujD5uUGvjHVl%>0}^YU7l
zC1jIGM@yR9*SSXjN}(9d`265+^}|dIOXhCj#IwcG+juN524y`ji5J82a-S}6V>q!<
zlq6ofI9eD>u{WohTT)E6BqWH$t1T>|q7$|xMwoMj0*jcKB!QVAIzC<)Zy`U&9KNJP
z@FMB9kya6wc3hr~y*-`JrCHn9^6VmcHkN#w2s)J?N##b^TCe=CI&uASm0TL%nohNk
zpjp~*Bd9bEmBZm#+S0A*)>ihB_IB10))7=|8@lz%=-33IAU0YgOyI>PQ1B}HM+-zT
zJPN~t@!4mT05?}x@Y7Pw8pCV~mO|Z)J*m~4QKGR_Gqp^EPV~%epql(dla>(g5@%d!
zAd~5YTf|HwOFFT>u7S$ZS-N!B%%At=G@ffbHJEeu>(Yk$Y_oxw*!*45#rY5We(t{6
z5&97JfuEl|ZGG6icINSbj*fv}LR)YD-2L0ynI**v&eOiuC%=;*2kH+nM&T(`|1m4x
z1p2<lwjfav!wI@@8N<F=y8EfYO~`QRiVsbY_U|`&pTb%ocN8So!ZaEjPi8t4Lq}J2
zD#>(&PM<5q6De^h7`6%IT!O13qz=O%a}&12<94p!5>E{Y{9hP&DyLXsERP=}<Rq+e
zbz{Uta^+?DIW;63vgYoTg_?ISoSj{l7*?-7LrRTN&5K@miSyeEzSKoMvIhL9?ykK%
z+(qNx^Go>D2$qhKrg}on5tT31(HEuji})fHqj|*Ul}0m7N^Z{I7n<R6ny{iIV?mZ_
z^U7Kr+gGb1q3Cz0#Hwp*6Bm*hRBQnW(mol09V^a=DAI8ziJLh~lkAS|B_7sYMJ84b
zB%LY2{0QN<MQs%;SDCozsOZx|6Jw`l?!ipg33km=Wtgd-VQJP}=+<D78+ADCC#a^G
z;qPYd;=i9+%b%R7618_0fl9o(eATksBBwm*gX*`&<QZkAn2IHaL8>q;ME~9kXDy;L
z70aY*s#<o?^i}4^%+JptRuR2j)I;_fQy^IzY#a*gAm<Wpt|wsKwI&Tr3?l2NG8Jv2
zs*4JBuUAG!S!N%J;G}<hjcR!APW2kf6fq;tjp9K7e_@amCZX?$>sEDR0yW0Pq!w$K
zqqFQg6{;r_$otf5?2K}k&d|3s!N>&F)n^*RZ8GQ1(6^y{7mo2~&Z4zo`jJ4s4di0-
zZEcloeH{$`n?=K6zUXM}Ocn5D5EZ*fs4r95HO|LNYMh!5ih(AEEeD&c&bVOGl7-*R
zsAb(1<l_GLplx!nUyG^2x*Pe;ECR#q>daNd8fxe+ZSA?b)+>HsVN#at`x7FfW~P-i
zamL=rx31w`>ASVH{nbprm%IDPw7YG?Jqs5TuBfdd5-at0&Y)9>t;EAxL`I|pze?ro
zRk*JU@ylm~oZM}M-Q9419t%H3?teU1L)48*s)!61!TkLhzku^Ad3H0tyxT{3Q21!(
z88cNKR><DZfDi<P3mK3v4c|Bv$SDkRE}@<}lT1z`cke3EP`ws~We3mL{g95`G*xjO
znvZb{GxSYVB)iDj7+)<sQ$(2!=XiuO>zMA1DmqNMd(m5Cine87&94fh{#avHZ-!BD
zI=KU~mO-pQm|`J&DN0)vreJYk3;rFAL|!7H^{uu$bS2Ze*bxv6<3R<(po7j)&Hrry
z9OLJd!iPD?^{~TNSP<;^hD#gRx5zUlV^|mP7Q*pskgGw$5Fp>s_xJ(M$G1e_;HM#p
zFAuS07`6>?NYkg32+sx_*M}sIYnozMG2q&OBZ>F`DX)*8C+9<@@iOjHBoS^17(T`!
ziEtAI7fFQM0gmfI65;NE;d+onI14cRnn)r%957rDk_Z<7hU-BR;VFROd?XRRM}Z@W
zaL5ZN=Oc;m6M*CTkwp0aC~zbZUa7#5M0gG0_?Uzw!l^I>L087^z>ePsa2!ucX{W%O
z{)D>#j`P}J$72)m*DCm(f8zT9z7nur*zw%ih@Y&$)ge$n#Y+@=%mDw?9y%P8;Cv);
zJ(^&EE8z10M-t&d3LHs<p8@?VfG>paA&iyw+yfjRGY$bh>l6GJ!13c_fa89`_2B+_
zqTrwU6MtBNpZ^m+slYG)3CHVzUl+gGetKOU!13!f0{*EV?kjjm%I$%G?ExH*6C@G-
z8{is%BZ=@Kz;XRZB78!DBZ=@Sz;Qd3628DY%0HlA<9d)p{JDVR^+6KhhJfRGkVN=O
z!0~z_iEs<RwE#yF;m!&iNrZa?j@yGI!UGjJk_Zn29JdEaghwlIBoPi&EDTPV1^$4s
z*AYSiV@JgZBRDZ^A)hbiC9qT2F+wgJ;*(-Hh>8~`irCSSc<~Rs)K7V-zw@B@2i~?%
zdE0*HZDYqK3Z@Ump!}1A)K3r4>2eFEErTMuI25<BlZEkH*^$u#b8$kvIZggamw)2N
zmTWdx7%NUNXR}2@ZVX!#53^&oe2yn(M+y@oVtD30u0G<wLG$uiE0*iEp4P%=M<nw3
zy!byIaXBImH##BpPq0WRj!uA;#9v*AKWz{@COU%qyXEFqw)XNPD?6%O(#);wEafFS
zl%{KH1*PfQT0_Z7{;dtm%grmmpB><~e!Z^;+skJoUUUy)DG+ysq~X%v%g2-9>+gw|
z8EzW_aH*iO-C1%(p<ula{+(-ps}K8gw_sS#;D-U5eAxlMIACK1d|aDc)2AAlZtLY3
zu6={MH(nM;@dN_)7nbxW2suT-L`QxpU+eG%i7}kWNOL}&1`Ha`Kn?%p^fK)4u7u*F
z#oQ!1%`7o?YpgIiRw;Qge2XnY0nZ{Lh7&F3S(tM<+$}r{JU}eAPL~sS0uetthG)jN
zGPANZqehzX=w@^qb8d|JdPQSQOS<BC`pD9J{a-ckUj*FWBFaN<TC??EpYVS|2ibuK
z&vY3NpXt(c08N*tLrB4p2hntSI*g{v)6*yzEAS0pi{RT)*wtW{HzADk|8-leVEhYa
zP)G4ulwwopl~49u@zHa7754wp-2NVo@kiS6PcDjI8#fwmXc#Q^Z~+4XzAm(oSIB~D
N2LGnZ<{W^r{{<^Gr4s-E

literal 0
HcmV?d00001

diff --git a/benchmarks/opencl/sgemm/libsgemm.a b/benchmarks/opencl/sgemm/libsgemm.a
new file mode 100644
index 0000000000000000000000000000000000000000..717198605fa1024948bdcf80bb7e13f3f90609bc
GIT binary patch
literal 4392
zcmds4U2Ggz6+U-o$Bv!23SP>O&BHe1wH?<p`!oCFMF`bUlcJaaQbDMyG@02QZ`|y9
zx4Ua6sgQ^@b)wX%M750&6_pYxjs>I=Wz`}i9<W3xuYG{L@IWyw63PQrtAg5!@SWMY
z&NQJ(%45#z&i&4J&i$El=KkyhU$UE}BVQg&#y={lklMrHXFOh<kq(CkiAW*Z^>10M
zM!i(Ej=Ieyw`#R2ZmmYQQr5A0^Jt}6UvAviXf4>SHYIPKo*evR>Rf1cZMW7~Sgg8w
zrl4oDdak6Wvbvi#>dC<!Q=zD*U45pcyBR&5H(GVJ4kp7vz(`>=9F2T7>(3)y%w>8L
zIz0G;x%saNArP9L2zrNZXhaRQ$7lU&lJ0{wCd74Z{8OZjtx_y1L;*fARiz)n=kF&G
zWbj^mi~U3e`py2@)W!eEHorpK&YDVvKKS+b`E>_%tnVS}fV<m1KhBTet(RK=1&B&j
zd#Ms%aht8h`ck}Tq>OZY@<_YgXnir6tSq*VEIUT2UQ1T1D>c2g*lgCDNxzK&W887e
zW}%#Q(laGDUo4snrA#hgaA(SHKE054%+x|TRdNcs{deq>`}eynWfpQ~s_10W`I3{$
z*eTm~(*-kU=CZ|daVF>FoK!Av=Jp3KaHCoZB`Qj9U%wWZ5H@6~pB{$&Y2Zne2!6w;
zQz(xH`uU(eSAwFnjGCpV_t7tVDlE*YemkbP5AN?BRsz|JdT-DlQ@)ORB=~+<H`xIB
zG|FheGnV}4fQNZ72xG}V3_KEHmWZ+B9|0Z>c*c^S2Ob%~1HxGH-vS;Dc*c_Fc#QPp
z0bwlp?*WenJY&f}4m`p^hA@`=Vc^k#XDs<D@G$Q_5XO>!9(Z^cvqX#~&$_@)#j<N~
zKvA(CjHX+)jm4$Kwq-Y)_6f2cJCx2?2j;)|mDzdAI%K6Y8GE^1x0-HcvDJ2)a$jqI
zH8ZyoG;h1d+r(dY&0T7fWtEN}w<^`TV^^&eyJ>OnWVD=Ciwu5PkKOec>A97;bh=tE
zp{v&7Q|>2jT(p}Nt9_z@md3o%YBz7Y+Xlv)-j$4{dfPP)9QuylYTKowWbh)h>~a~4
zy+S*|YArj~*DblJTKTM1a4g%gvIT1ai@!wDr027G!O?9;&ldEaiKo`K9favbtyh=r
zwtZ_ha+Y=6vgQ!@Iy-Q)w2|)R^lqMKu$F>Pz<1w-@`up+l5h@I#=&PtS>Vfli{2wx
zPV^8|iP`cv*3AqW4gvELV%r{+Cs6p4lw<z|SdIm_gMSQ?`LJKOf7zx1pKyPF3-tlG
zGpAn%{biVhu%CP<*110}%-x3~$Grv|9#>-al^BqoM+1I8eqWvs+gLw{!oDdKmVcuB
z#z(z}9CsXmv6RbU1*BB-0LuLJI;fNEdmiKR7Q*AQNWbR*g00KvxxZfocHW|RetS`p
znL_WV4xV>9+cPOM-*dh5+5YDs!2HedDbbJF;jI1uXL}T9dsN%^TvUrcplU{1(WYk=
zZR*>KHfbx`_>&6k3has&zo2Nw%L4ZRHTLAzb#3|rDT&vFHg$naRaLyPe$jpAB6(w{
zg}L@_#G81H%(YGOCi)}V*xS@S{aHnsI3Udaf04HD4Kh!^73ucBiTb;WslG|E`xEfr
zAdEw?a}(4ZxJ26c>lAx&oH{r|l*BA)lP_beSIA3jkvZ^6#M1?Jc4I$XdrFvULU12m
zVpW*y8x%V`M4ht(6vKY3BzA+}vvbGrHf~_;{AG&mnTPK@d`D>ei5EWXKKqjBtdH`T
zV*9T@f4_U?r@~7J(OJ76+$wGVW8=M;h<@PdTSC!A#G9UtDATJE@4>Bz@}N+>saZvt
zT2;LLTZ*z@M7^=uDDFy8+?Aq;yIMKv-<AHa_<NgIgf@+Mp1K<GHdWDGze>u*HDUJs
zF5+FfLguO8lXpptXp^tQ=dFmhF(<mG{)qR#LdvG<$F!?njdaylsoQr<G5h{RvCoW?
zH$I1UeSUjw-xhVwB}f@xLjgDRD&AWqWg6|r-k{DGCW4%7kv1iBvP#`3axyjR$J|^y
zh@71B^V%IijMoRLGc-goeDuo30Oqqhl#ihutet;}V!IDg=h<(-_w&fdnGd_qydXF?
zh`rb@|LlEp=tbmXHpqt%+t)VUyVr8Tc?{>l&%r$!|9k1Ke$3xvdk?;=|0JaIKHs@X
ze>Bif;#?TUNl}2#_-vr_*|i5}i?qKBJqFzu_{4(4<!RKC{~_amkK%xl{9eWZAHFHx
z?$2=V|L%>u<Sl=tZ~MRbmYZ6`Zrar<{)RgxT&+HO#j>p0a<#Q!;=+on9D`3F9(-k3
zxLFow-bScBLgfn|g^YtE(Z5HX<wH0S3E4+tHeJIR^ayHh$01}J=f*%`-xLbVZ(;af
h_lgQ^at`eGee@R=X<G*=Q^a=mYoKPk9GAbN{{r}^FGK(U

literal 0
HcmV?d00001

diff --git a/benchmarks/opencl/sgemm/main.cc b/benchmarks/opencl/sgemm/main.cc
new file mode 100644
index 00000000..561a0f30
--- /dev/null
+++ b/benchmarks/opencl/sgemm/main.cc
@@ -0,0 +1,500 @@
+/*
+ *  Simple OpenCL demo program
+ *
+ *  Copyright (C) 2009  Clifford Wolf <clifford@clifford.at>
+ *
+ *  This program is free software; you can redistribute it and/or modify
+ *  it under the terms of the GNU General Public License as published by
+ *  the Free Software Foundation; either version 2 of the License, or
+ *  (at your option) any later version.
+ *
+ *  This program is distributed in the hope that it will be useful,
+ *  but WITHOUT ANY WARRANTY; without even the implied warranty of
+ *  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ *  GNU General Public License for more details.
+ *
+ *  You should have received a copy of the GNU General Public License
+ *  along with this program; if not, write to the Free Software
+ *  Foundation, Inc., 59 Temple Place, Suite 330, Boston, MA  02111-1307  USA
+ *
+ *  gcc -o cldemo -std=gnu99 -Wall -I/usr/include/nvidia-current cldemo.c -lOpenCL
+ *
+ */
+
+#include <CL/cl.h>
+#include <iostream>
+#include <fstream>
+#include <sstream>
+#include <errno.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <unistd.h>
+
+#define NUM_DATA 64
+
+#define CL_CHECK(_expr)                                                         \
+   do {                                                                         \
+     cl_int _err = _expr;                                                       \
+     if (_err == CL_SUCCESS)                                                    \
+       break;                                                                   \
+     fprintf(stderr, "OpenCL Error: '%s' returned %d!\n", #_expr, (int)_err);   \
+     abort();                                                                   \
+   } while (0)
+
+#define CL_CHECK_ERR(_expr)                                                     \
+   ({                                                                           \
+     cl_int _err = CL_INVALID_VALUE;                                            \
+     typeof(_expr) _ret = _expr;                                                \
+     if (_err != CL_SUCCESS) {                                                  \
+       fprintf(stderr, "OpenCL Error: '%s' returned %d!\n", #_expr, (int)_err); \
+       abort();                                                                 \
+     }                                                                          \
+     _ret;                                                                      \
+   })
+
+void pfn_notify(const char *errinfo, const void *private_info, size_t cb, void *user_data)
+{
+	fprintf(stderr, "OpenCL Error (via pfn_notify): %s\n", errinfo);
+}
+
+///
+//  Create an OpenCL program from the kernel source file
+//
+cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
+{
+    cl_int errNum;
+    cl_program program;
+
+    std::ifstream kernelFile(fileName, std::ios::in);
+    if (!kernelFile.is_open())
+    {
+        std::cerr << "Failed to open file for reading: " << fileName << std::endl;
+        return NULL;
+    }
+
+    std::ostringstream oss;
+    oss << kernelFile.rdbuf();
+
+    std::string srcStdStr = oss.str();
+    const char *srcStr = srcStdStr.c_str();
+    program = clCreateProgramWithSource(context, 1,
+                                        (const char**)&srcStr,
+                                        NULL, NULL);
+    if (program == NULL)
+    {
+        std::cerr << "Failed to create CL program from source." << std::endl;
+        return NULL;
+    }
+
+    errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
+    if (errNum != CL_SUCCESS)
+    {
+        // Determine the reason for the error
+        char buildLog[16384];
+        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
+                              sizeof(buildLog), buildLog, NULL);
+
+        std::cerr << "Error in kernel: " << std::endl;
+        std::cerr << buildLog;
+        clReleaseProgram(program);
+        return NULL;
+    }
+
+    return program;
+}
+
+//
+///
+//  Retreive program binary for all of the devices attached to the
+//  program an and store the one for the device passed in
+//
+bool SaveProgramBinary(cl_program program, cl_device_id device, const char* fileName)
+{
+    //cl_uint numDevices = malloc(sizeof(cl_uint));
+    //cl_uint* numDevices = malloc(sizeof(cl_uint));
+    cl_int errNum;
+
+    printf("try getting program info\n");
+    // 1 - Query for number of devices attached to program
+    /*errNum = clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint),
+                              &numDevices, NULL);
+    printf("Got program_num_devices\n");
+    if (errNum != CL_SUCCESS)
+    {
+        std::cerr << "Error querying for number of devices." << std::endl;
+        return false;
+    }*/
+
+    // 2 - Get all of the Device IDs
+    cl_device_id *devices = new cl_device_id[1];
+    errNum = clGetProgramInfo(program, CL_PROGRAM_DEVICES,
+                              sizeof(cl_device_id) * 1,
+                              devices, NULL);
+    printf("Got program_devices\n");
+    if (errNum != CL_SUCCESS)
+    {
+        std::cerr << "Error querying for devices." << std::endl;
+        delete [] devices;
+        return false;
+    }
+
+    // 3 - Determine the size of each program binary
+    size_t *programBinarySizes = new size_t [1];
+    errNum = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES,
+                              sizeof(size_t) * 1,
+                              programBinarySizes, NULL);
+    printf("Got program_binary_sizes\n");
+    if (errNum != CL_SUCCESS)
+    {
+        std::cerr << "Error querying for program binary sizes." << std::endl;
+        delete [] devices;
+        delete [] programBinarySizes;
+        return false;
+    }
+
+    unsigned char **programBinaries = new unsigned char*[1];
+    for (cl_uint i = 0; i < 1; i++)
+    {
+        programBinaries[i] = new unsigned char[programBinarySizes[i]];
+    }
+
+    // 4 - Get all of the program binaries
+    errNum = clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(unsigned char*) * 1,
+                              programBinaries, NULL);
+    printf("Got program_binarys\n");
+    if (errNum != CL_SUCCESS)
+    {
+        std::cerr << "Error querying for program binaries" << std::endl;
+
+        delete [] devices;
+        delete [] programBinarySizes;
+        for (cl_uint i = 0; i < 1; i++)
+        {
+            delete [] programBinaries[i];
+        }
+        delete [] programBinaries;
+        return false;
+    }
+
+    // 5 - Finally store the binaries for the device requested out to disk for future reading.
+    for (cl_uint i = 0; i < 1; i++)
+    {
+        // Store the binary just for the device requested.  In a scenario where
+        // multiple devices were being used you would save all of the binaries out here.
+        if (devices[i] == device)
+        {
+            FILE *fp = fopen(fileName, "wb");
+            if(fp ==NULL){
+              delete [] devices;
+              delete [] programBinarySizes;
+              for (cl_uint i = 0; i < 1; i++)
+              {
+                  delete [] programBinaries[i];
+              }
+              delete [] programBinaries;
+              return false;
+            }
+            printf("Opened file\n");
+            fwrite(programBinaries[i], 1, programBinarySizes[i], fp);
+            printf("wrote file\n");
+            fclose(fp);
+            printf("close file\n");
+            break;
+        }
+    }
+
+    // Cleanup
+    delete [] devices;
+    delete [] programBinarySizes;
+    for (cl_uint i = 0; i < 1; i++)
+    {
+        delete [] programBinaries[i];
+    }
+    delete [] programBinaries;
+    return true;
+}
+
+///
+//  Attempt to create the program object from a cached binary.  Note that
+//  on first run this will fail because the binary has not yet been created.
+//
+cl_program CreateProgramFromBinary(cl_context context, cl_device_id device, const char* fileName)
+{
+    FILE *fp = fopen(fileName, "rb");
+    if (fp == NULL)
+    {
+        return NULL;
+    }
+
+    // Determine the size of the binary
+    size_t binarySize;
+    fseek(fp, 0, SEEK_END);
+    binarySize = ftell(fp);
+    rewind(fp);
+
+    unsigned char *programBinary = new unsigned char[binarySize];
+    fread(programBinary, 1, binarySize, fp);
+    fclose(fp);
+
+    cl_int errNum = 0;
+    cl_program program;
+    cl_int binaryStatus;
+
+    program = clCreateProgramWithBinary(context,
+                                        1,
+                                        &device,
+                                        &binarySize,
+                                        (const unsigned char**)&programBinary,
+                                        &binaryStatus,
+                                        &errNum);
+    delete [] programBinary;
+    if (errNum != CL_SUCCESS)
+    {
+        std::cerr << "Error loading program binary." << std::endl;
+        return NULL;
+    }
+
+    if (binaryStatus != CL_SUCCESS)
+    {
+        std::cerr << "Invalid binary for device" << std::endl;
+        return NULL;
+    }
+
+    errNum = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
+    if (errNum != CL_SUCCESS)
+    {
+        printf("build errNum:%d\n", errNum);
+        // Determine the reason for the error
+        char buildLog[16384];
+        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
+                              sizeof(buildLog), buildLog, NULL);
+
+        std::cerr << "Error in program: " << std::endl;
+        std::cerr << buildLog << std::endl;
+        clReleaseProgram(program);
+        return NULL;
+    }
+
+    return program;
+}
+
+///
+//  Cleanup any created OpenCL resources
+//
+void Cleanup(cl_context context, cl_command_queue commandQueue,
+             cl_program program, cl_kernel kernel, cl_mem memObjects[3])
+{
+    for (int i = 0; i < 3; i++)
+    {
+        if (memObjects[i] != 0)
+            clReleaseMemObject(memObjects[i]);
+    }
+    if (commandQueue != 0)
+        clReleaseCommandQueue(commandQueue);
+
+    if (kernel != 0)
+        clReleaseKernel(kernel);
+
+    if (program != 0)
+        clReleaseProgram(program);
+
+    if (context != 0)
+        clReleaseContext(context);
+
+}
+
+int main(int argc, char **argv)
+{
+  printf("enter demo main\n");
+  fflush(stdout);
+  putenv("POCL_VERBOSE=1");
+  putenv("POCL_DEVICES=basic");
+  putenv("POCL_LEAVE_TEMP_DIRS=1");
+  putenv("POCL_LEAVE_KERNEL_COMPILER_TEMP_FILES=1");
+  putenv("POCL_TEMP_DIR=pocl");
+  putenv("POCL_CACHE_DIR=pocl");
+  putenv("POCL_WORK_GROUP_METHOD=spmd");
+  if(argc >= 2){
+    printf("argv[1]:%s:\n",argv[1]);
+    if(!strcmp(argv[1], "h"))
+      putenv("POCL_WORK_GROUP_METHOD=spmd");
+    if(!strcmp(argv[1], "c"))
+      putenv("POCL_CROSS_COMPILE=1");
+  }
+  if(argc >= 3){
+    printf("argv[2]:%s:\n",argv[2]);
+    if(!strcmp(argv[2], "h"))
+      putenv("POCL_WORK_GROUP_METHOD=spmd");
+    if(!strcmp(argv[2], "c"))
+      putenv("POCL_CROSS_COMPILE=1");
+  }
+
+  //putenv("LD_LIBRARY_PATH=/scratch/colins/build/linux/fs/lib");
+  //putenv("LTDL_LIBRARY_PATH=/scratch/colins/build/linux/fs/lib");
+  //lt_dlsetsearchpath("/scratch/colins/build/linux/fs/lib");
+  //printf("SEARCH_PATH:%s\n",lt_dlgetsearchpath());
+	cl_platform_id platforms[100];
+	cl_uint platforms_n = 0;
+	CL_CHECK(clGetPlatformIDs(100, platforms, &platforms_n));
+
+	printf("=== %d OpenCL platform(s) found: ===\n", platforms_n);
+	for (int i=0; i<platforms_n; i++)
+	{
+		char buffer[10240];
+		printf("  -- %d --\n", i);
+		CL_CHECK(clGetPlatformInfo(platforms[i], CL_PLATFORM_PROFILE, 10240, buffer, NULL));
+		printf("  PROFILE = %s\n", buffer);
+		CL_CHECK(clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION, 10240, buffer, NULL));
+		printf("  VERSION = %s\n", buffer);
+		CL_CHECK(clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, 10240, buffer, NULL));
+		printf("  NAME = %s\n", buffer);
+		CL_CHECK(clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, 10240, buffer, NULL));
+		printf("  VENDOR = %s\n", buffer);
+		CL_CHECK(clGetPlatformInfo(platforms[i], CL_PLATFORM_EXTENSIONS, 10240, buffer, NULL));
+		printf("  EXTENSIONS = %s\n", buffer);
+	}
+
+	if (platforms_n == 0)
+		return 1;
+
+	cl_device_id devices[100];
+	cl_uint devices_n = 0;
+	// CL_CHECK(clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, 100, devices, &devices_n));
+	CL_CHECK(clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, 100, devices, &devices_n));
+
+	printf("=== %d OpenCL device(s) found on platform:\n", platforms_n);
+	for (int i=0; i<devices_n; i++)
+	{
+		char buffer[10240];
+		cl_uint buf_uint;
+		cl_ulong buf_ulong;
+		printf("  -- %d --\n", i);
+		CL_CHECK(clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(buffer), buffer, NULL));
+		printf("  DEVICE_NAME = %s\n", buffer);
+		CL_CHECK(clGetDeviceInfo(devices[i], CL_DEVICE_VENDOR, sizeof(buffer), buffer, NULL));
+		printf("  DEVICE_VENDOR = %s\n", buffer);
+		CL_CHECK(clGetDeviceInfo(devices[i], CL_DEVICE_VERSION, sizeof(buffer), buffer, NULL));
+		printf("  DEVICE_VERSION = %s\n", buffer);
+		CL_CHECK(clGetDeviceInfo(devices[i], CL_DRIVER_VERSION, sizeof(buffer), buffer, NULL));
+		printf("  DRIVER_VERSION = %s\n", buffer);
+		CL_CHECK(clGetDeviceInfo(devices[i], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(buf_uint), &buf_uint, NULL));
+		printf("  DEVICE_MAX_COMPUTE_UNITS = %u\n", (unsigned int)buf_uint);
+		CL_CHECK(clGetDeviceInfo(devices[i], CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(buf_uint), &buf_uint, NULL));
+		printf("  DEVICE_MAX_CLOCK_FREQUENCY = %u\n", (unsigned int)buf_uint);
+		CL_CHECK(clGetDeviceInfo(devices[i], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(buf_ulong), &buf_ulong, NULL));
+		printf("  DEVICE_GLOBAL_MEM_SIZE = %llu\n", (unsigned long long)buf_ulong);
+	}
+
+	if (devices_n == 0)
+		return 1;
+
+	cl_context context;
+	context = CL_CHECK_ERR(clCreateContext(NULL, 1, devices+1, &pfn_notify, NULL, &_err));
+
+	cl_command_queue queue;
+  queue = CL_CHECK_ERR(clCreateCommandQueue(context, devices[1], CL_QUEUE_PROFILING_ENABLE, &_err));
+
+	cl_kernel kernel = 0;
+  cl_mem memObjects[3] = {0,0,0};
+
+
+  // Create OpenCL program - first attempt to load cached binary.
+  //  If that is not available, then create the program from source
+  //  and store the binary for future use.
+  std::cout << "Attempting to create program from binary..." << std::endl;
+  cl_program program = CreateProgramFromBinary(context, devices[1], "kernel.cl.bin");
+  if (program == NULL)
+  {
+      std::cout << "Binary not loaded, create from source..." << std::endl;
+      program = CreateProgram(context, devices[1], "kernel.cl");
+      if (program == NULL)
+      {
+          Cleanup(context, queue, program, kernel, memObjects);
+          return 1;
+      }
+
+      std::cout << "Save program binary for future run..." << std::endl;
+      if (SaveProgramBinary(program, devices[1], "kernel.cl.bin") == false)
+      {
+          std::cerr << "Failed to write program binary" << std::endl;
+          Cleanup(context, queue, program, kernel, memObjects);
+          return 1;
+      }
+  }
+  else
+  {
+      std::cout << "Read program from binary." << std::endl;
+  }
+
+  printf("attempting to create input buffer\n");
+  fflush(stdout);
+	cl_mem input_bufferA;
+	input_bufferA = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float)*NUM_DATA*NUM_DATA, NULL, &_err));
+	cl_mem input_bufferB;
+	input_bufferB = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float)*NUM_DATA*NUM_DATA, NULL, &_err));
+
+  printf("attempting to create output buffer\n");
+  fflush(stdout);
+	cl_mem output_buffer;
+	output_buffer = CL_CHECK_ERR(clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float)*NUM_DATA*NUM_DATA, NULL, &_err));
+
+  memObjects[0] = input_bufferA;
+  memObjects[1] = input_bufferB;
+  memObjects[2] = output_buffer;
+
+	size_t width = NUM_DATA;
+
+  printf("attempting to create kernel\n");
+  fflush(stdout);
+	kernel = CL_CHECK_ERR(clCreateKernel(program, "sgemm_single", &_err));
+	CL_CHECK(clSetKernelArg(kernel, 0, sizeof(input_bufferA), &input_bufferA));
+	CL_CHECK(clSetKernelArg(kernel, 1, sizeof(input_bufferB), &input_bufferB));
+	CL_CHECK(clSetKernelArg(kernel, 2, sizeof(output_buffer), &output_buffer));
+	CL_CHECK(clSetKernelArg(kernel, 3, sizeof(width), &width));
+
+  printf("attempting to enqueue write buffer\n");
+  fflush(stdout);
+	for (int i=0; i<NUM_DATA*NUM_DATA; i++) {
+    float in = ((float)rand()/(float)(RAND_MAX)) * 100.0;
+		CL_CHECK(clEnqueueWriteBuffer(queue, input_bufferA, CL_TRUE, i*sizeof(float), 4, &in, 0, NULL, NULL));
+    in = ((float)rand()/(float)(RAND_MAX)) * 100.0;
+		CL_CHECK(clEnqueueWriteBuffer(queue, input_bufferB, CL_TRUE, i*sizeof(float), 4, &in, 0, NULL, NULL));
+	}
+
+	cl_event kernel_completion;
+  const size_t local_work_size[3] = { 64, 1, 1};
+  //                             a_offset  
+	size_t global_work_size[3] = { NUM_DATA, NUM_DATA, NUM_DATA };
+  printf("attempting to enqueue kernel\n");
+  fflush(stdout);
+	CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_work_size, local_work_size, 0, NULL, &kernel_completion));
+  printf("Enqueue'd kerenel\n");
+  fflush(stdout);
+  cl_ulong time_start, time_end;
+  CL_CHECK(clWaitForEvents(1, &kernel_completion));
+  CL_CHECK(clGetEventProfilingInfo(kernel_completion, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL));
+  CL_CHECK(clGetEventProfilingInfo(kernel_completion, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL));
+  double elapsed = time_end - time_start;
+  printf("time(ns):%lg\n",elapsed);
+	CL_CHECK(clReleaseEvent(kernel_completion));
+
+	printf("Result:");
+	for (int i=0; i<NUM_DATA*NUM_DATA; i++) {
+		float data;
+		CL_CHECK(clEnqueueReadBuffer(queue, output_buffer, CL_TRUE, i*sizeof(float), 4, &data, 0, NULL, NULL));
+		//printf(" %f", data);
+	}
+	printf("\n");
+
+	CL_CHECK(clReleaseMemObject(memObjects[0]));
+	CL_CHECK(clReleaseMemObject(memObjects[1]));
+	CL_CHECK(clReleaseMemObject(memObjects[2]));
+
+	CL_CHECK(clReleaseKernel(kernel));
+	CL_CHECK(clReleaseProgram(program));
+	CL_CHECK(clReleaseContext(context));
+
+	return 0;
+}
+
diff --git a/runtime/startup/vx_start.s b/runtime/startup/vx_start.s
index d719e0a8..3706f833 100644
--- a/runtime/startup/vx_start.s
+++ b/runtime/startup/vx_start.s
@@ -1,5 +1,3 @@
-
-
 # .section .init, "ax"
 # .global _start
 # _start:
@@ -74,5 +72,9 @@ vx_set_sp:
   RETURN:
       ret
 
-
+.section .data
+	.global __dso_handle
+	.weak __dso_handle
+__dso_handle:
+	.long	0
 
diff --git a/simX/Makefile b/simX/Makefile
index b4ef8ec6..1f665a2c 100644
--- a/simX/Makefile
+++ b/simX/Makefile
@@ -12,8 +12,8 @@ COMP=--compiler gcc
 LIB=
 CF=-CFLAGS '-std=c++11 -fPIC -O3'
 
-LIGHTW=-Wno-UNOPTFLAT
-DEB=--trace --prof-cfuncs -DVL_DEBUG=1 
+LIGHTW=-Wno-UNOPTFLAT -Wno-WIDTH
+DEB=--trace -DVL_DEBUG=1 
 EXE=--exe $(LIB_OBJS)
 
 all: simX