From 83662458134f534c5f1ebcb4dafa261f3320aecc Mon Sep 17 00:00:00 2001 From: lucasdelimanogueira Date: Tue, 30 Apr 2024 00:31:46 -0300 Subject: [PATCH] sub, elementmul, pow cuda operations --- build/cpu.o | Bin 1504 -> 2400 bytes build/cuda.cu.o | Bin 12248 -> 26880 bytes build/libtensor.so | Bin 36896 -> 46224 bytes build/tensor.o | Bin 17312 -> 14272 bytes csrc/cpu.cpp | 22 ++ csrc/cpu.h | 6 +- csrc/cuda.cu | 64 ++++++ csrc/cuda.h | 9 + csrc/tensor.cpp | 232 +++++----------------- csrc/tensor.h | 2 +- norch/__init__.py | 1 + norch/__pycache__/__init__.cpython-38.pyc | Bin 0 -> 193 bytes norch/__pycache__/tensor.cpython-38.pyc | Bin 6517 -> 5638 bytes norch/tensor.py | 85 +------- test.py | 70 +++++++ test.sh | 3 +- 16 files changed, 228 insertions(+), 266 deletions(-) create mode 100644 norch/__init__.py create mode 100644 norch/__pycache__/__init__.cpython-38.pyc create mode 100644 test.py diff --git a/build/cpu.o b/build/cpu.o index 15cc718eb9d0302ee8807df1b96da6d77ee080ab..b7ee56989c88d3090b82fc5b6a3bb0618d621e07 100644 GIT binary patch literal 2400 zcmb_c-Afcv6hE_LYNoA<4}lW<5TT&Mu85>Wtjg#VmZC0XDdX-so1y!`?hI-v^wNtD{RxVJ^%h~^qXmKB_7J{AEZ8|S_v$!|QrUqs=iJ}k$Ot zBESteT5%M>S9;_JO*aTtPzn3q%fCWO?wgWZ*i~}Ij*<&)Z^}i=Q1bF_KG@tL7Z--e zSMrb#6_w~GW&Nv15kBsFtIHAJals)k7nPUFa@|wbXxq%cyT0$Ofvj8<_Fk`yZ=7v5 z?ays?x%=#~|2g|F&pqlsdsxX98j<*IWJTWElnb0Jll8)af+yTI*)1Hgp2u^C5UtJM zM)ZTS{>7s#H&ozd$XnjV_c3o`ZI}AQyo+l&F)J5@z10;Z`>yf+_y+X#^mO_fLlGmP z8@>*yP4c${409D&Z1qe2HmSAA!l?yG32HU!In=&8A@BgK5hkLN)L90EM?)aYRSWf% zW#udQ0yZE{O7fT<^2O?Bo}Q{wPX;rr*mc5DPKu9NVb}N-UWay`@_2taD|)8Q5rm#7 z^o^9jj{~;^(e(?WGxJj)uM$m`j$^bE!S>PF!XlK}R!6P~NQ{e!Ci}+&=W$tU+o#DP zVoeO^Yi@5c{3luxjvMGxAD#U)OfC_7flBx8ls5>t;Sfz3QZ$tU^|TGmrF{otNI_i>So|OMGS&=sh=;Eg|T@Ie!xlob$Wnz&XDlW#;30 zID%GBTXu?jX>BGf>DrtQ(qvc>vT@nJbp<$RNah>?6}JX$L&m-}SvQHkRqs$H(p4+Z6tWuO8r zwc=j@$FDNn`Q^3H1-Xc0JJ5_d)bn5K16WHv{it@~^QpgG_n?hN&%;D3?bqz+`>ff^c&m2}-j638-pLAk99xmqT07 f0jl}~REBHvNe+EZ1E_k<$&8%loE1=BI>^}o6h%AT diff --git a/build/cuda.cu.o b/build/cuda.cu.o index 922c97bc38635f18ebab0b0a7ba8d9889101f810..b3b48377fcaa49dc259751389374094959971b52 100644 GIT binary patch literal 26880 zcmeHv3v^t?ndW~VQn%W&TJlq}jn%R&Aw#>}Z>cRCwgOiMA9jW89e%v=;5_w%2V`@8Fd7CF7-rx zMnf`1417^;sso!!=A1Ly_S{)@QZuuPSBOR4 zOx8w}&)SGxBe&6j^R;nGdpfW^r=p|T@#tvhQ`5Wf6gKm~`2NEk`|laIN1uQ0l4TT~ zdN^{Pbdm{u+U31FOu~QmVC_*1Yez7w9U5&vXmb9^$T5mOW07OD zK6N7Hw!b-jAaYC*2=bYbL(`ChIV5K&v-f1=DT<3HBR?`Ot+DpE&@x<&5MlcV-kG0vSAp(YhaLAjbeC1IecD|oBmarB>xlL{b=T(v z+wZ@-&hNlr_z0~*$40XcPK$(3V2r^?_ymiDk5EiIdm!=zMnVHQHVt_&hvb3`k9%@# zI~h4`M7mdv%j?GF43B%`826sSxc4dsG~9n0LBSOFo-pDbZUVDN8aX@G`NY(?_bSG{ zGZUyv1|}`gaq!dKna=>4(=o^IN{nN4e~V)wnH1@>3f|pW8f& zE8b}K+$h~+MP46mKf4*n4334@@oogi*EzaJM5j+=D<`t$?W51FA4N`a>vf|1#TU@7 z+#N6W^l0`X${KAyPh4Z27xBi8WiE_%=tl1RXonTIEXdcL?TJ&fwr0D#6X|qsc3@zy zBb^w`I5LUB^iax?8CvbAOWUA3nT6KQ4eK3=RB9+?V8|D{)$2_T-4QdWVI-@0OkJccYKvwj zp4y@r=*^RP%3P-DZkpZFdTP#o1_2H1W*AN}eAYDl8eDQU_T*w})lV5-Iu*o{6Hnx| zJgfhs3}-v-DLfxBe9jb%3SZJlPT4g4yA0Q-;a)`Mlux#6#A=Xy$Z&e{!V}s}$%cC8 zl+)7~o;ZGn;gUTZr#poPPD*)78t8Y@)n0n{$VdZNC+;3;u(GT{NOf)@-*grB>8KA_ za||K+0s?5~_Q~`mnH1?CZ^tx!L4)R#>N_i`V%gubz78%wr*;N?`YJ2X=%W#<*BW`^ z)#ph03Mr5DL@VgUNGJMS9_e#=q^H><=*3JT`dl9Ab9tn- z#l14X4y7Els&dd58k$woc%dR!RsaXSRAxdi)PS~l0`w(T=pY$LqAxY(kdAUGw?kQi za+l=fXizytj!H2A1PJ!cUl`FQC0dzDX9{r6>m6lGH)-&0?(XKTDnec8ar z9o18h9F@w`qa6}_nm+Zwryls!1D|@}57GnbK3pJZA)$r;v>mGU8v8w{!*W9zmT`ki zsRlpG%V7CFl@Vy_{T6ts8cO%!QuKL%WY+}^UbjPg38Th;BO1J82Xz>Cr38`um+Vlt z2l+pXduif-)ed*3*-{4ZHXg7&T!!?s;G+wb&5Zs{J4E*4GS5x|Q682Z!&DgA1Zo)I z%`N(@VJyN3=HJucExSG|g?m+kPG6!jSW`vB5JnmuPT673uwFDVu^%SJQ9ssK4Hr!; zM1rCP=m#ugJ(@eRZsJ3%M|za6zpp>J4#L>JH|=_Hx(Jv#?8VNDlKef^gki{N=s<}!xGk8nLj^z9u)LHt1l_R zeaQ%`XZGqk^CtmsRU?0iH8+H)UN`!K(W5-adW!|^8=-3M|6+qbOl50L%A$+2Yu4X5o~YhR zoq+tcDT{^5I(``cFzar^KiEF3pP=?L8BAZ1H75_pyCtm~yr&%i^)z0s^b~poS=jg} z9fqRH)c0GGH(`G;`z@AS|G#2~@u+Vs_IcOpAjf@}Zu|HAVE6>B(i)?YTWZsJ?eM-X=$m^|6!zGCG2m3u+UdT zt{yk$TQ$xXFbXG6oS$V36(;M}t zKVb2;)>~>JdzejV|Mf3oEN0>{Uj7>83;wg%KPZ1rB{P3yVvP}hugAuZ2zrdymM^0D zG#<&{81HB81KlVi|EX$bAND5NhdAlO_#DOpN-^~9!|+S|$Hn|)mN5!yzh$!hMS%S~ z=aA7}7H@I=S&#CJWUL=Gdzd|8P!UCav>xnW_cU04&JH;KkYHrLSbwe`1#pXcY!Af? zs)y&vepW_7SpErm3D+xX{{^Pd(3his1?j0CyPk1tuNevBZ8m$OYk3O>lxXP3C9P;JG6g_JY)P&{j+e|sGoC``;+G^nC$5e2A zwDcek>#v$U*gqov%g6Cow?}93)B@IE>as4 zL=uYGKE8<%=b}VI^Fr(mDmnQAly8 z6~HY8_?y@Q4ip+#A~#ToGV|LEgQn>}!ufYFZbLstglYV56yU#*hrc2Z|NVyZr|IwD z%dV`So+6m-uf>M3c!hOxI?nmS9HjOWujoH&3#D)4hnip>0~a`dQtGgL{-f}pkADY$ z)N-f9pU;0mRGg3h3@O7T^XK#5KNsNtED@Xd^ZCzO z1WM^O%;)1DSU!7sE1BsK)A4z?vS8)*<}g%y!%$i62?JF7Fjs{| z(0dcdtAc{VZe{2!OQ~K5WviQ*d6b24Gg)`l~C(( z`aszN$5Vl3TRM?Rx^^ZUu&tyn>2cP<#kt;QTNv|>u2kI7(&CO`+19-$6L)2XTssGb zy5a+K6l8y)$AgNq1Fk#zdlCas+Y@MqKlV4<(Alo^DyR=S1F$O4ybbh<#%ieYhM>O5 zxdCKmw8E;{ zpViCo%%w+9mF-aMN&V`we#L$#UQ&wvkX~Q5PqBYhKU(ems{SM#1I&M}+WRAG**g3> zicM9Q_C9U3JHc9pH&4RFe|3OWy_h6S) z`aPP@>vrCze8%1DYr`8fWwjP|rS|nNrzsdOsq6K6oGv)nZ0p?>ccli`fij@MRp52^ zwq6I?XW(Y+fe2>BKkSrlH{p=(OKni+hBiU1 z*LRPi^eZ@>9ss2uDnpOXF86kU@&-;L&~H_Bzf$XW_5z$y`db2Sf159$RE8X&f2deW z%4*%tZU6_^Z&o1W^Ej`E)pbd~16Fu3t2n^21JrIsS*xyXX=@sRdJl3IkHAI6R+8#Y zy5h-XYG~Nio7uM#98FGKosgvq>iy0>DESaw5WwD0>Y-#2)CHXm9Qc9bina*UdTnp3 zzjb|E(1(TP$`@NaO>OSB$a2l+3ps0{XggG5w^#U&sAIkj>oK)#SPMbF$Jq+XTA1tq zceS_{>b=etumqi0=R=+bU*OqUmHs7454vtKH$Djsq+6;tq*woP!;g2fq>UJ2rB{}G~UN@-ac69 z$JbhYey9sLw_$hspH$`sg4%7^*IxIZDiwipMg2%!5eTB573!UVmbR8iQ=}0Vh8m%= z=@0~hZs#q?=lc-`4-{M#v?+n0*ZBd@`J889MbL{H{mw&}L-(L2%K85o76unV3ECBG zB;CQ_FN@~+mlxZP>*_lSEK?3OHns(u{2SUH(5?_M*ZqVMa&ht2+VWsDviYuD7&Z_e4Reen!=KU_*01-qylhn}nue6Brss>6H3jZ1)=z6C-$dwuqM_#xLdh?& z@Ktqw%WIl#E!w#jl*g1sO+SFzkh31FKY;5J!``=wD?(K$V^-1fkS7M#glb&}>eZm! zsx*W=`ydeVIXmI1(C(52jqOT%t9OI9b=n*|TeF4?artDwjNa|5NIoSePM7xS(DH$1@5 z7J1th+dVMH(*d@x;kPhsnr}8vUo{G$q{kE2=6mmg`NPX>J`cs6r&fAB{ziY3FX(TI zYI8!n?TdXrMSVtD5co$}ne==cm(rJU12JJQj+Z_ZY_{Pl=gJQD+i~J7R2qV-%atqnCkfNOOIT2r0?N-t8B;&(y4K|h0ws@LNSxkDH>+*u5*!MBR6cPqtl zSkdU*1lQChx6M))DXUYS#(#jh_*w2>eb?-H*%U6|^@=_p|G`=)YtvVRf`3!IVDKhT zs}xH$*zjexAE~MW^PS&^$^X%Dkp~N>}vy?^KjZ98gW`6z!b~bt!gQw*o(f z1&y6$`1XSKuvYXq=zU=QtX8@l)JIBe%VGIY>bI)0u}s~Ff!Py!v21}m3yXY5l)0f_ zLxacn8@Srz9~oJb(_`4+}i|VX?Oz=DA1UTJP%b>gwj%)xL|R zpx&)Ec)i;d?HpeHuWCwj*{l^%`Zlr#-d0PMLtw#uoDWJWb>-SJr`P|1R_6_#R;yc{ zn60evxX<7q*#zs7-r#OUU*mDF{v~pH-MIZtE7jheD$E~fMa1`pHX9aOZ95bQf+eEQ z@!n!x9C8;~Lny@A0Lr4G*){q+uU5R$mKVEe`iqx%O+^Nb59`*7VGE>{+pTC7LV0J>cew+xGj zCQ?>@ui84h_%ler>(&sCiM^2`>IWw4Z!@alUAl^T^-V@VvhU>er>E4P%+(XD?CaRf zEVKLt6&DNp99Y1;0!t3Gm)exmp)FKMhyx+9Prn?^0P$kF9L*7f1gfU#DrdT=J)5zP zQXzj0@cMcp68_je>8~&IdiH98sP8oDsjc&PJ>4TuZEu<4Dp$LSlI&n_U*R=hzmS~U zKBk-{I9MnZl5^~1%ITRw&Ur2;VQ!cD@1j|(UGyDIN*?+r{L?q-zn6J^8Oam%T2Zcl z&hvV9h9m0LEldFUt&~rSZ6tp^AMd=LpnQo`>F}=;`RcUXESp-ZvDr8kqmnlL)Bsi@ zu1Gj{6>Q|VxSti*A;eWlFB^-9QPPa3dPYI|uYur}Kt@0PB@CwnNSP!gOkV4n2e z$Zy7{>gBw5&v3% z{y#YV-*S2fzb?jwrbeCggc#?Ih)<%E9@>gH>8xi0 z__6}{j}a$3|BVkJ(cWKh{BL>a6ZR{(fg*Z54`E5ryEuN3hZI478^;5*@xny=21H8i%^&sJ{p`I+7)87-rw?`R^ zJH{H=7SFV1(cUPtm(>J~n$AHyGCn}uEOU2S?_TQ~yy#KZcZwS6$tpV@m8@4MFcz(7~L`wQp*X~29O za_TlOnv~A;$1ss-A$0u!8wocu@3saax8Hs{QuCsBYSF`s9uK~f9?J|dK^?(u>~(dB zUB7ut%lgf+Ep2UEBimxz*0L1MD zU_djmm9g%=FW{`|AE43UA(NTLWHzUeu(nhpK?DKQ`)7h^t2@@w3mb01#~i(`SZwF; za4d-rV1@?c1L)b<9yh8ZIUU(_ANjX+Xt1|`XEqhj;4>N&Jktwy#Cv@rvpJs5uz|`N zi2y5X&0vdmGRT-MTleDdP7Mv>^BstF4&Ko}*b}=cktM$Xlo?sALKa6p?g(WiX@{@SQ;EzhU?4O@Y_#t?b*$7vhjLJ&NBt%yjK7( z<@d=XNA{0H!e#$#lkhf4@85Bp?3d%o8ypw*FUNXH4hj~S%=X{TakHF<3gFL3a%B5| zCgHOEnr%{V6tkjV?&Ub?mFN3mjtl!~O{MfRNsjFQDt?Y&meWxHzn$Y||9?rsW&0nK z@QrAPu;0N4D(RKy->DjY=5PM%l1bl{06i`*x$=> zVJGz^r9Y8y>QaFpEr5T6<7Pj6i{oVHCj1D!FPY@rhIN#FUI2exk}r>wvjya9{5+c4 zMfQK0gk!Q5z$-XT_RI6AisNSc*O=ry!KC9DD1hG}$(Q5IO$Fo+NOCCV3BT=@=r>FF zutZO`ia2>+0sPMk;P^@;OC+NmKSJ+!CHzJSpOEk^63)X9&~`K5sXO7+H{!VgW?3Sf z{eU_zJNT7o=_68@X9yFz2--paFbVoaMiEB%Rp8>>`mZEhoUIZxF5$!@&aGdRa2wCl zE8ud6QGZiAS9K1|*#5SK^0KhlCTocy5+7 z;o^ChxW5qgi056$P4wbCnEe_%$|ZWzE8+wDNluPaoj9kzWRfG!i4`KkB;?STrsR-t z%8K*yun8CE${Qt|>=E&#Pr@lH&e0E(Y-+bgKJKBw>Qt`bUAnEjj zOZY9@?)ZS~az|sX?x9_VsKKF3!i8TBS;>y(P&bu#4rcL-!$WwUmJz}eec~6~cBY21 zNpkxqPP1{`RpgNp=9q{b3$a{_ABuaF#I+DV?O`aD8W$jbiXVbcTpftf(C6Dv%o%AL zzm|xAz{`A^@DhF@1#=`ToI{L^<0j);$qgX=G_*u}DvfIV?y;et4?B)x+%@SH6=wZT ztRR_Y{vDiO_(AXs{h~~J5Y~P^9|SJ=<(Nx#CNiMCF7cc3qg?+#5R;V5`f07o$Db-U znC_H#DQj;34vv}iU*P;^H=6Yi72p?x^0W&Hq-5sbk9A~|+5e|$lY&X~vr(Gte^K6t znSA4?>5mKw!B69dhLX@P@UJ08OT1bCIj&#OBMD2QEUrfo&8NR3kABi4_=Ns%BF5Tp z>VOL~=x4^@zlZR{^qctm^5~~7HS2!`G1BkGkDM-X{mW$#^F&KYT>nwZvkLxiNmwpT z<^`WX-o*^vLYmA6N7$fn7?<3CVZQ)06i717{Nb=c(O~LWGyg|OmaqRtxc|;^!8Cs6 z<45p|{>!JEH21aAe}EBwe@ygOj&)V9Kw{GS2m9dUlnr?GrUDtvN-RcGs$Z)ii%o#SiGB*=2u&AH` zwtXQ4foz;lkQjx8WGd*S2K56GjS0jsL?cLyM!>`@Dj*+-i+?iDz3;wpt>Gl^-t&IH zbI+f9-@W&Zd`+(Pjzwr!_4_puOa2}E7DAsQpgZE!PoTr`Je>0A)^GouorUGakHWjf zi|LM8xKbRU-T%OB@egoDtAHWZtGYiYhsZUm{77vI6ZlOn_pWwxU7-pgEvQa7AyBem zq2e^)5xdc^wXIcJlx2zqTNfNGX@-Qq>ahPfTrEprfJ^q?ZP}shRya`>gb)1X%KI~vNT&!M6FOhVY zL!oeD;Q+DbQx&h_l!UzboAv7mQQ&Sv5IP$C&K8Wuq?#jR4gTtF)IfF`9oEnCQ51)G z_ly1}qy{uPO`#+@AAZlNpho@9E@p|+0MTfWbR(zL?GxB0&KS81!uG}>ISu2Dnr}v6yF|illoNKxYJDTdqA7}Cn4uMRH!Yyt-F%pAs-!a(^gQ3U zOB84JyeJrt`Jtm?0W`!_mntk=6xPJUwAY#+ikDF`ng67DBeAjNyAB&$f;fezQbAG) zQ>h@GP~dhdOa}?tddAW+{Luvembfxx=fM(7Aw z^CPV)q4y{pZtDrm6SaoXgh>=VQe1?%1zal+(X%!NAW?|-2$uu=McF=|SD zXJ)M@8B5Oc_A~f&wMH&;{b$}j!`pkjoySV=ygkR; z^Q?vGHEwG*8q4Kp*4!rf_u6L&jCPE|$Lq6DnGWd@bockHA$z!UV^>te#`F?W0|(QZ zr&hAL1nbipTu6uDK-vv6X}4Z4v0;%ItG5;wW;9Q=#P-2g84bEJA@sXDqv^8W)6&RC z;8bP_?8>-dI^)(COCH*AxtK8ji^0>W=`qyI{KqA>305G>CfsD-lGp@n?Oa0Y;Pp-o zmUp^67bR{POd-0Z)1AE|p>mZDu?&o6nTS(b2^8RE$T0~&0I!D|Piw>Nf@PkT<74Lg zrNCV?xX5wcKP*|75$Mn-D@IpUF4BV>}5_vebK^s0Vxs!|kOTt%3VRUf> znv)Vjj!Bz@vm=+4$YkvU@U(#s|1Yg>F@VKy#>cGhMadDHaT9+e)dW;Z{#( z*eW&!OO-w5p&Tl}J4E%8(lm{1yo-7z(7I3iu zKWoePW%m=jn{M?rV=;BPn5RUP&m-e?PzobC-MNUrDqS`>o=ebvJG&zBDvvk^g?R%k z+Pstw+M#=M()p$U`{3B-rA{htks#balSmN+`jf6@0_<`_Uw?v*mBP{fB;6T?@35OF zg}`%3Iu?d*>?V=JZUQ;%K0AYb_}=Ug`XS5zaDn-qLi_}o__%}`xs*Y diff --git a/build/libtensor.so b/build/libtensor.so index 4cdb052194f292e629c14695d2f4cf514418a0a7..e23e7be84db4ce202201fc1c0b22d5ef4e09fe3e 100755 GIT binary patch literal 46224 zcmeIb3w%^pwLiYk%;Z4`dGHL6lHu9VhGCw0Kq-MRFci{8Aoj^`2+4$`O(yAN0)ZAu zO{tA(Dr)tjUax7bir#8bTQ9ZRrY)7|4>fJEqE?&M>TeKKq*&2gJOA(6kCQVeGo)DW z@BV+E%g$$J?X~t=Yp=cbd!Lg%d;J^MSTv2PU}aY`!iE@|)oN<$|uDQ7#=x{e2?6 z|FelOPZpw06-L@=0UOC3mU4%soPvxPC`gzoFh*+Vb5PRF&@TBDG|LT2IWwFudWp>H2Wy!xh&s_Kh#t=Vme+@y^3%K0XwvTP(!qa(s&Lxe^}=PJCSW zc7OOVjO z2!b6rg-=TQ#$^Im`EKR)mBW50?cXi=za{x!mWr=OeiRl+{@H$kEB;>qCHsGQm7t#z zfZZyhS|6@r%xjK>mGoa6f{$uMC=|7Zqo~RIXRo*X2JMWbF zhGe<^Ci#D4^22>8^pBVPMRLGae(*{=KUXI5y;%UZP|_F6f<~m@Iwk#YlK&NHPlwch zctxW9KbC%LtrT?S=Qn|pp9g7@Ko}B$-759$bqIp8bA`n(?NQ_!VGTPYP0V{3hO~EE!rE5dc4Ylo2erHD0e-&4oNu-JpACqa4cCbcE|f8{y4VC+lOAaKI?K zKU~in#D-cZXk5=56KihTM(rjNb+orSTyDZKfYn`jCE+E6w&=0#b#=jZ&d0?B9hqUo z#8}wb6lH5GD}x?~lU3eu!^&0bf-Z-Lt=+J0Rb|laaF^;)uO4$d_^e{V?uGs=_@i(7 z*d#7!A}tNHbbRQKu?z`m%c;GmzpYqPsr6a)!d>G@_dd`lkmj|9^p_sJ4OlwsGo-)Y zkp7{MelneQX9o=FrH1smzj!m3c56fv`YkLh(+Qv4)HH#1WuKR{ju}nksDxjY{A#~W zdo~LDfi9N12Kspe-EN?F z%JOQUO*7C7`UR2!Kxr!2`$dfCv?o@fs9(g0jwOvAoc$t7bhQR1N?AXBQ_CFL>l0ql zPv3%$C88bz{US>IR0k?l_lp?O<&vEfZZOcjf`xx?HPGcz6sI>D=%h=9wtf*KT9G0G zcJ>Q|=&FqprK_L5h5cygdbqb=M2SC3r=q=rp#dFw`$%RIYqQ1U%6%5TYCRWy1#gB*DVJV+Q(s z1O1GFF1G|c^}K;j=RtoXw6Q_?b(}UdXHUMM1%CFI!+sZlwZebV}SDO zIBf_}ejTR;KgzG;g&hCW-&OhT9Dhy6X(NF0>-b!bKcnOGIR14Vr;QuRuj30i{!tyL zjR?xG4s6Ye+R%T2h?gfBMX z3rzTBCVYwspJ2k%O!(jLNNf*(HsQZF;lDBAzck@LG2uTn;m@1!@0#$(P54($_!mt0 zXHEDgOn9Hd1AWUsg?OMZ?PCPW&qt?VsQ&{V)6!5sEbVB|S%DLlE4*`D-Fwcnj>XXSdMC8? zLg6#trg8Hrn!F%qCAp_ajtH^UB-abM-Wp_(C!<4wo@yTNtvS|HT?Fa?8MYRN1-b_` zieJeye3lFJR6CLH0E=EJWOdS15hGHD>~KF9x1@-4f8A?e(ATvGn@Skgbyqwhr~O-kQ_7k3QSuKi%DOMmzN+ zD&}s=%1>7IEq{RGs#Kf z6AC~$>OV&+d;RCgzZ+yR@_K8|>BWGxT+=39@YTOkjZpc^{pa@NL+MgI?Kz}%`_E~o z9szkF$g$i1Q&{_bV~d7$qD9vSdTk46qNCR98aV1dL-t6sV9p5?pK>XOiKG2Ujr0W7 zV{a7<{(j(@*Wh>F6rR;kHDt$=lEoS5^Bus9OO4?JD7mP3l163dmDcCY3-tLt>N26n$gB;$nVW~HHs6055&vm$oUadfM0uiVu`j_F38l_Msn4Syp8@9? zQ9oHQNz`i3dG4n(u=5O73Y787Pyjm@kz!t@eZKQD=}D24!;`?e9{8!Nx!ykCOOol6 zg6VRD{3m469U=*%6;wn1W|VP%g0U4ab`Zu6n5qa8gS}!+`>n%yV^w|O#|#hSOvzmP)Ygd^=qy)ZF6$w4q|rdZuu6NkNStG2VhWJ=0Bd> z_dK;h;UV#YKLr0t_dZqbtxBr;PfN_#BA3)9;%v=WnGpKmhOu4opb z-uA*Q;ovFKMW`>6H6o+F>`##K82)1c?FSe*EZzPg?c@@;Lgr1~FdwC+Husmec=N8N zq~3~RL-+1OzZBEQ$Nx;Cn4Ees?U5gPeng_#5Cux);oh3#f!>Z7k)MZsr|ullCkOuz z5I4qZRNomfJ07{6m)ATy{$>v9qhBxXLY@mohs4bHa`H&VyfYt-1t7N{me3gBYZdQgo}sZNe9QzCWP&y4oIHZJ|)B zcdHfW;7-}C=26-OyVXaDJ9)SIEhY6}GBbNxBA6xW?Ws3~i{4we`Zo5j)Cu~b4YS4tyVX@^DaTRW zs)Lh9b*nHeP13C{C$kf}RZBQ&ecX?oc5Izq+#id~&E4t%Cy(k@e?l&!y442o{5>RAVEV|AD)vf+b7-F8Sron+py4BDh zdEFaktEsmot-R8isE@OHaXU$l*0PDS)z`UAqq^0nAvda9{TaLo-KrXkz{6Ms9)cea z)9{1_YU|Ad1`5^y466;=%0+M0Eu`0Y)NV6`Bez(!ZsE(^6cXk06DdT?=gYoRNRk~1 z$$Z&|WWMYxkh=SP*P_6BM>+`!QO|pYTB#1x32YTB)I;bU+E91Q(&K-l*Z;=Ig07iF z5{9mCH0dQU?w4Nz61yLs6bChoPSDf{yI*DhKt}95F1?1TwNLI-@U( zj%cU8JEEcoDiGXEQUR)I(k%9Dw`RfCJ3=^2FS-2)cBo7IMTYX&KgWNHVsQ{Ou3 zKSoOWU@r~#)TBx0RDPhmN_V?ULhC_wE#-vXAS^=EYq}RtE zPm`05D<{c@LaW*TeH*;@51XgrI z;mJ<1d`(nvya)UH6I!BA&x>6BP>|2Mk$iFrWV`*xwUcLYpn;wl>#cbMYs51!i0`{5 z!N&8gj$W5_vWperl45xDLe8J36eZ z`7x3_#*q7MtYB&Xk2Z4Q!wNqRD_%m**p(NT`%m+(zpFuS6Bw5DuK#K3I5n_Abev!O zj-MHf^uXVNAP!_;5^wD%cykc@^0av)T8-UtHJ%y7;_!mdx(`Yck1)KW&k6dnEo3rp zhK6&EXL~i;p5Tb!2rm&o>Pl??2_{lgA|?8nQCCPWxY5K1@t7lai*DlUC&|R~)FIWG z(c9g7-WF#@Twia^+xm!sHg%rPjNVS@syN@`UG+Tgs^_~~-p09xv2J_zVEfZC9O$Vz zan%2|>^Cu9NpJ7j2*WWZ;^bl^A3C|90u3jj!_LI)Pw-aFdjuLa@_xv8JoyjhO`~VB zyfc&MeGB!*e^%ZXS#BBV?-&>$QtY6sv=Y zy;Y!!J$v7*EIK_po?A4C{yY$?jm6^)xgQwlv3BFA@zf*x@szUYX2W?fogwkjBkiqv zyfJA=={}9=bkB*=2R3Asp+)tfJ&0-7`&$~7cz-sebIYBe$DW{eOGC`~|3i$})>B$f z9uC?Up)2w!RBnG4kt(@-oj}|@blOUtn>0~N2gK?`WPFFrn3MVW;%n$P^1R_Ud})$+ zn06Bc5(Y{;ic02CGangNK2CaRV%HN!j$Xk{rkvrYdsL_rVFc*Bj33fbGLy;}hE32W z{e3j%hh?gayu7@Oa3(J!^zbsCmu37H;^uvimN|5W`}fJJjGuYKJUUmSJcD%JhF*Ml zbQKPwnB4<=`>OZu8^{ej`^@=_9ypcGDyfRTO@@NWFG7#5D{pU&4%3=cG#{L@91!2RjO(_|9p;Icm?@2zv?!<3}XIY0dFoRbIh;2e6OFo=g)1kkH! z4mcvl##Shb%|JFL$80C>bI1HVPdR22c!gtjkvpScN$r?>$aiGF!7(GV$)W$1IGgk> z{~{={7vof{v`-^WozsmB6X%>fI94AattPi0=|buO3~T!_tlisFb+^g+NBs}e=+op&Wi^`_(n$CSG0Z zZ+WlP)r7MAbE@k@H)Oe5o712^et>@DP|k9%AR#=T*Td&$L?NHN;ub1}VUo`*-xHM8>Prkcs4agUsQ z4r3p795l1?;}S8ux_)M^6)r$(13hWKpjN@7M`ItxQ8&eJW#UgQl;a42z6p)vz}h#Z zmO!s>cMSl|k1z*D+b|!g$)mBa{FJ*VFLCM8({T=pda#>@Jvacvu8l{|^i~v3cCe|wrj@-h# zseg^U9~tPKQUl++-a7ds=oarx@^^&ww}LtUrLRjKxa5KV$30L<7rNli47B=xA=7TeT_YekmYZ)OYXS+SjM;deK% zO`&l8h?Vwet6f|YBV=$(jH(7XGesfZoBg=bBO0mY z*JQvZ3K!5fmR$+TE}-XK_*Gt*vY0NsLDe)zo7$S2>UcwW?{;2L`MPjjYfBq)!Uafn zUgCCM(`coUrq0o}RMv*Wtx>zY-bVHchkbKvw6@vMf@GfUs7ZS`v{Pud*SCh+?IaHE zb&b5i)OAE6xar4Ejh*@ewZ7;PPdRyJ#8{{Cwn_d2re}Bkv3jgvwEc+w=A_r z1PyX@J|(G3-j_D0-F6|LQd zn1Z`n1kPcwy)j(h$sFMw4jRw?5eeOze|`id;J3h-HCIT2@ooYl5x+6ezd~VOH!9#Y zM;^wKr+;NKRUMZ*D8yNTgo-EAM8}B@Zizwtm<^MBN{@m{p4dt5-3Fr^<@s)dl2h$= zfZp0bQ0+E0PvGjC+U|52}^F8$OJl zkx>1+;jF_{AFT~Xontaf8!u*zGMId^VtNTKR?t|Wi;~sr;KF07p-gV~*sOXv-VL+S z_A&W&doO$mTY|k~vl{Fjn>o?mG5G64+d8%f8zNZWF-IiS>@X8*6DSuy`j6FrbcbDF z5>D{e7<>j-jlr4drwcF#+iKyTi(Cn)(c_|8yPPu7>n~PJFZRU>8oh8avQ)2QbTqwc z$KV?&^_VKbs+aWLFdIuZ7XR2QGF7KzU9p(ZUdQNwi6t7FIk7}z@S8hxLSLrtZtC5J zdDt9%9|%@P|Hz#f27y{}o4q+)fdz+ET^?^p{1K;T8z>@Lj0KN1$?X3Kr_(=}Cx&-8 zePSXJ*axsUv*KP@bNCvk+stG}>$Qn_BjxIn3mAIWX<{CphN>XfTUp3=VxD1@BI&Pi zy1^bvrz71Wse|Gatju;N!QAY-y1I&3X8f+MqD&qYjFje7`I?}z&xZW4AV&~NY*UHA z7|Fhu{P0f_i9wcQRDLCb?Y-tVoS2)w;Qp`5uUeKz&vuOR*_LE7YJJpVZSzE&WNi}- z@q9xZ{KTvHRZAy+Jq~_74u0xAieL34;@9Kg*W=))jW4xWvPW+JxjlLu{CYeqlkYdS zYI)2$J&)O@iJUEzKcIJ3>lWYs{og2gsl^dTGcQf#X>3wmei zdaSgAG3c?-7yvygUy;uUPT83A83u{;)4t=sQa_6GU#33*-4v)-QeE)3O(afrStQ#X zrBjU@FyNy{J1|@dmppLE1D8B-$pil*Jz$|9KA?$&CjOyZHhq^^@9`hhDarwk1ulg- z_{`5?^LJY~gVl0>2KzxF%ifKt2(Nv$Z98LSCvsU8p4f}fWo5t1WtL8?N{K-Fe~`-- z?0`I8$>H+Ha~b|AU1WF?2W$`JfPXxD=S;ye&i{HY^Y6km&JP%&54fFvfQEIgXO;@a zUcb&Zz7r_xLj0_iy^(7hAHiCcNJHldXPJ>S5~u)=jSc3qiJi8zcziF555Rt;Pwz~N z<9}6H8vURPaHPZBt`+gOkq&<7x4mh5d9UJ ze@--m%Y#PxY;&UNgg?!3q-VBg@H>*t_Iw@1d20)z_22A2T0jT0^f z4*TXrveJQffq~=mqT_9Kk8Ze%uj52|mRh3`Kc#>7D_>cz(&tHe&U1wFv)RNNumV0r zeeDK@ccqm7?UAeup%3ZMXOCo#)9oYvPQs>Da{lj;ec9t^{lf#&Gr2$Cd6sGLKTjVw z4(pOGp3d#H+4S?{Hwqy?E>mwplx`FCq4c3X#=m;^TRYh*ULVAtW#qSa!ru{>s1KBH zP9!ZmG2Ldn|EZ^a+`l~kj5OW<9M8^%JqG*I)A3BhfM=8aA_)Bn^2nduSw5majr55? z?F0VSaq`b@h?Bo5aQ~#s`UlSWGcsUb7bQ#orwjQC!ZYVYGHmLZ3$JgqCrZzxN`U;l zNJa+XHhj2zC$DbdACwpAaVkI8!ShSY)b*jgXSCad?2sKlutn`PlU^6>f(UOf*_|w{ zAoBH9Z5vS^-2RLVz5ZX!WqJ3ep}!#v$1|vX&AB^+)R%I3@S{GV?!{->@HyVU;J>tG zU!h-XsbK6_F3V_7=j(Mw`PWnaS=?SO|HcwFuX9{F>icObBjpc$Xs&NVAI1aZN9m~V z$9Wq-`RtwJ#u@czwvS^xjf@7r{Q&Pidj6U1*<#N_>B!&yvFvP7{}te0qU>MADpY=H zx?^K zKD9^kH^%$%yO~YkkRL4M_MtXmALx`H#^(wk&?55Njp3K%52^miE#n+i{*00Gr&0Y* zyjPT$kGGhARzjZ$V*Z%3gWFTVETAGknh&<}Jq^-hxeV6Aws#QvHHowx<<53z;8w z&LqR@Yt9bTj~aiH_+Ph2>eua=k6&cu`qQ8Q_EGsDe-Gw!-WN|I7vxuEs!wV!`1fjr zZYvw0`399m6O6w8kfZSe(2}k{#B^y`f21?K>96X4^!g(MhDrTse+zcl_;{N^>z6lj zv#}mYr}YTOVgFt2d_HH4)*}aVfLp2M2^{N_1v`*GUj3Hx=-=aPVue(JK64$FpXO)q z|Ia8t=3h~M>T~~-%TIOw50sz8-b*02%N2TpnhJ@nF(vY$(ZEUbl!=9>-Xn>>G9{Yj zQ@wZ4C>IgxT$DNy^$WE&DB!s;HWm8KIw|R}WE?Ohnozb#=xLMcbkHdlt#WK0DI*^X zbS}p5_B4f2@@0GiXEZeMKyKh?qY_sjOBggte!rC8Dj7*0BZ3%%eNa4L_ax4;DJOkn}m*kc9vrYex?rZrP@mz-Yf@x>Jh)F zXm4)`min?Bc!r9%72AXS4Z{-e5*A#SKI*@%q+g0wXmrZ2*8a!_M zf`+TS`7Cc+J}Ze_Rg%xnY9(I#f=K;UF4xt|va+(8>-Vv_j*bF$ovYY>vs0V8c+0X1 zR_Lx^1%<8(#tJ=%O$Vx-WVw6U^mfmdWmzq)JC?G;o?`ogx@c`vy`ys)YxV*QN2jvY zK18Ol+iU9`&8?310yfW8>|xq<>`=tFEUP^fZF6i7+1ciC3))=83)s0y?qyjOh}*YC zYVE65IfKBn>UKnH9nn_D_U6`YwapVXh(A~FLdG4J)02zEQ8An-qA)D`ARmfJihnm}-X8GQSnA_tj zeTdaTqkE12YwQY$up(daMy3@q_FK*JdrN-FWvmYE6K|eRVslH23)nZyJY{7i&eCOB zt&tgQZdq|5`;yC9;`8`CZdb{AHmAW|l2%uSCfyiWZJAWQp3QT6_Gwy^hVJw!rZusG z@<%4*xwkRxSLjB}c9X^C)#iDN8yI^elyp}Cp(B`$XMQ%hjJImP0EOnZ;+_HStnzAmo&;`l#E_P1K8i${Bb=0=CMOr%@ z4bk0;nZ2wSvlB$Nv4!5^MmFwkxWI?n&=#_BGuVQXVmlhX?~rDl!7K}{yDGhvE7z2G zfXvtKSmi2P<6PsPZ}oV}i|4VlEvx|5p6}gn>GiB$iLh$*a#rGX6<0ECIh*ADt0jFN zTj(xc#Ad;X3p|i1^7x(}U*Mgk)x)LX(DoTD?QNERpXT$H6nj~**ZT#wuoQ!t<#smV zUN*I?OV#L_P^OA#A)&XQjm&}Nkmrx%sFT}+$DX1V6EX`XpZJIg?NQ_C-Xfmvr|OhD~# zMv=EN>rM?0M;(5f%_x10S;FwUi%lzYv00^;v4vD;1-_r9&v!YmW!fjSd7;}smpQMb z_%il!m&;XF?)G^+w=?UnGAFz2tZ1?8m+ARso5oLf?Xq~vT*ck&I&{<0XEo+6a~CgT zmh)^;TL~JWuXqjhmh?H6{IU;gY&Fx)qzB7ZuJo)rmZ{~JwQAGLo=uxu=DR1|He?<5 zMUc8RD?EQU8~1A<4_KzGdfA$_9Cj{e+BdWrW#3}+%8M7W%x|%)L!It7((}uwL&x~E z`Q@%4%M4jES1`-vOuI=dDtGN>zH(1-4ZE!T_Hk27tF)?0_iFbe+Vt{vHo3T(UDo!p zb!yolr1bYs#*{d*a%{wz6J}t!g^28v-Hn6z_8A6S7}-M*hE(~%lZV~!pO3ECZPLT0<65v<*Kz#aqnSM zI_GA2Tr}<+UhH;xOTA^D5^q_+IiY#-+~)c+rqW|rK*V#?YqK9L zS(b%a&e0KW%0b)6!2d%Qd!F zFsZs-o^oe7h7D&2hSrid(lYPT(rekG(&F{(iUn<($6ID-OCzq*pR-B$%=a$cGhuQ^ z1QYl|%{B#JY&pwWV_Q^S@_hQ#@J42tu4NRmEZjQq9g9U{Q;NUN=C`@-)TWo-sZB21 zm%XULd#9!qph1qk=wX`|fmE}ugmL+c2 zr{JKyY^rN7o9V7%lbv1cO83&Q*(@6-6nf5OGs|6;BDZ^sW{m-Qk6X26IpY_x?2{1n zowQ_Y_p%JE<2-C!fla$Ir`YX%%euf_GGr-S_1FY$k<0lC8p(RLqRm}$yJow}4VtAhH*29KV}jMP5zRcEWm|2toL!pLYFpwi-!<7XaSMjfZ&(I= zYTbY}7mAAz3|{;GVAXZ2Yc}|U>sDheKhwHw0~W_cNGtkuJl=?S>|gN@z<_0+jmHlH zKG+|R4*)(1I1JeLdHgph3ex>>JZ=Ycejy&O0Q}Iwczi41ul_wA?*jBa8jtq_e(4*K z2ef<>@|6A<2b&dz6@9axba(%2mIfF_X7R~a1e0lKk)yufc=N@PM8hnpHIZ& z8vuvD9gpKwkg@N?;|~B%|86{f7|;WF0`Nw_^ME@63$XbA(^JR?aN*O)2e5wt`2hX| z@F3uI&&K0JfOiAZGr*J2K@N-m8AsxAIGGhaACKPv*ag@LcpUH`>?-(PJbnN;g~}>%Z2j^iM|l3%{5; z-HY`1e>rvfW@>Yf#N!e8TDQ;Ir0b6&{Z+IzJQf~Fw`|nydjRP(4#ngAANqCsUec6( zk0Jf#ugBweCdjwz^2d?xMjK9*eh%sPr%0!IpA5hM4(a@8TIKKWwpwoESx~OL9_=?=c8@nVd6+Rn}rUlr#q281@fni>6X{^a{MdOhmk&%kbYE8C$hG?dIN8{ z8J`JFV{E2szjmZiS=!Kd?I^EpCzPqOSf_Dc5)bWLo2g$D&$i?|&ERQ7-{w43$$9pJ zryl)&I_W7h^5p6Dh7~)z5o1Csf8B@lO-Nq|UAnDXM*Q`Nl*h^YNP3M(=kuRBTn)=o}{raV@yg`FQn%dh(3jMuBSJ(o~vv)vwJeNo`)g20O?%M zrqp`whMt`m6I0nkzt3?W(nEl@nFc+Uzp7&Kavg)7?_=zA8U52Z;=j{K|2figpkl-~ zX?p$8gXlqwx%^C1`Ob13S3rI%Mmob-yqWZ`l;f66&jV=#(mRn(=U^(t4-eX};NapM#+Wp!I3wxCw8bNqW{88{fXJJ1O51Et||M#K#opbfO z=IVFM$@LUQeotNC^qcJzR66}eyAh}P&xq4+uv5U9od|XT)H|H>Hwb*X08IT3yE-fYf@jdBDadTufIJkI(j0SBa~9WcL5R1BXAG+$5)6&?_XYmT;?tZ4!1#xL?A42@gp)AmNaN!xEm6 zkiYN6F#gTK7lK_vr-T&}R!g{5!Zr!JB-}4yzl4V*9FTBG!eI%|NXX`jeDfrK>iyd<~1&ew|29)BArSJ)Nn3lt|BDs{bX@GnwkgiS%(y_4h87WFn4jW#OjGoKXcxp`L|7TQfq$oZ4db=>DTJeM} zw$DgW{FB)Er1Z%w@S$YVZ=&>Nt^Y|ii zL~4$tt9e+Zlp>-T?9pVx8crYWw-(Szznb4AJEOgV#zy=90O-l=e~9y6#tutAsCn!= zpp*VjN&D6O{8P!V#u>%`Cg}9v>uMb0D8AmoM>F_e&4-gvz-0E!PC>_sP*VPlpp)MQ z9~3zP$$50pDGv55V@WcCjOY@&#z46PsjMB8Tyf&wwdzhuu0EM z@Ke55N$lP4*jlb<^!Tt@(mPG%rSl=OGf%Ex)m~yJr;i?o_i*~?d7+!@ zA3eU^$LXWTr30KkdVG5xj#V~DmvQhl@K1!DcLoGCM2t$B^!%+%aPu<{L@(oXqXrSDD?BORmt6wqXE=z`eu3x*Xy8In`ssD}WcGhv@}FFt z$p1X(a|J2k`?VB${w(#-8@Uwh(hpP9g*`=A3c9j$j->x+nV{b(0IQPp*?vKwJu`)S zC4H|$;7ZTOIbG0<-;YcF16K>GvghR#{AW_or(*+<%x{ZR(ARN#4*R7X-%d%RY9)Q6 zsl0ck(DPZz|2f$n_!${={VIi>2POXx*N6<2-~K4+cU&Rp{45aBze@UA*=|)kxjaQX zUj}-PS0VltrP6BnX(V2;QvYr`jHg$pGo~MFBSS#`}wP+w^jp!4&+@gPyFvA5X#mXQ}@h*?#!hBeKfJ z(3Z^p<(!_w7Rw6cXRY90pMpP_3Dz~YhC}Um;8Ncj6mOIS>!Yoa_F!#CC!@>AnnO`sqf=U5>P*fQY-kEM1#2Ub z+Fe2Z1`p%+n*{4ST3U8N#DE3C8BJhqZVKOqi>*4g;l7UzZf7t^7mTf`jaGJSYYGSO z!a;KgDu`GUR#vZS3fD$<)kK?`xsY3=SL1#KC=1m?8nW%Ify33c(OP12`Cv(HR9xM+ z8CMN4d97JpTN}fLd=7%qmbT7@Xwd0&dg|N_K8oL z*isv=-3}OuQhpwhCX8t01qtHHJ~LM-*U*lO|HwMnos_k@dIxlh47mXJt+*&97zs60 z@9?8|h$x37XHtG>X|3;Q4h5U)gSb+RiezkoWY$<(Z-V{3hO~EF=2i1d|TqWe=4YeKNy2kaPNI2BIuFhXhG(X)VSG|67(2M&S zJ)68isGDY5WH^9noa)A z!ObgI;ejvAym-g)i#MzIMW7WG%}v{=aFM8^z187zL*cc&*;Kmxx88aylh-B6dgPME zCO2AM-k(VKK?d=foe~`TRvvS|T2mVp!`J9M#_TF^TwP}^xfE{R9jd5UyJ6j`%Anif zcK8@NL31daNJDq>nc9AJ$!30=<0faYx`B<{X9EwpOOlr|h|w`o4AqmQo11b<+A-J& z6ZiVsP;^6WdzAO>WFB$(Moi?*)no{Fh-ho3q9-)&3wRxVt83e37RGgOb+-lU8gIj> z+|)dh>V=VJN(a;fuA% z{*&0UuwaOc_&~O{x|X&?8nxe%l5p?nZJ+^zKw^zxfTx>lM^`#IYE2}B=52OULtCUN z9BnY?>hc=f$h&K+MhWDn+sWr5If%2ij%fR+%zUOIss`D>kYKc}6s~udrs`72R7r}i zNX15P;5}gz?s^SxAMF5r6m#=YEZM+YQBirr4J%iz3%VR0-X@Laj`#*st!TcjZ8V!y zwuT#;ws%DM4Y#B7#r3d;L60mI%cxDeFiS^T!+bC@*Otg&9Tt35e7P}NE$>R?edaI>o-ZR==iu3y|#&p1P4ZF?hg)b9#InSfZ#@f7*? zeF73hnz(^%S6If{nxhPN#$uXwAnMrODkJToI>zr_V-7xjJ0h+1w2H<(YjTl_;v=pv zBbg)AD3|$-^~gtIFrcn380xGGwMB!tvKIBn^AkG&p6?aYCF3C9xlME?jv&xLU`fU`Fd)bk?+?UJqlw%ZATj%SSW>N%5w^Aot~ z%rP<0HvBC-dsX?X=gkVnB)gI`=f585*z)Q2tLIh<#&AwWfdFmwmApE~-2@trPIY

baVN>iM|JU&*WT-zDW&Nre^i&_O}xK5CLz&;1lMKc+Lwf7&Fko)0Q`XpHt?C~OL&t>K-Hhp_a7#C^*l+zd(}an6fny_ zYm!&bITcjTX;gk@`4>#`>UpSw^vn2N(*HOZ^zrV4JA{a@)Y=YP^9ub!_esLmOc zyvZE)pGcr*k2Ho*Q19FH_X%9JH-U}xA7w`d6q1jWyn0@HNXjd|RPyv(kNTdHSI>11 zOL;XgC^sm5%C1j9mdd2d%jds`KO%Ar2%vu}c?D-d*eK8XCF_0x+bGcxl$?SKO!De^ zavODQ1ZAhdM*OGnD^2oU4~mp7DeqK_5~_R^`6|d0O_g6gpFSkzw<<;n&GIW{{^op? z{Iw=|u=8MqdfzqjeoP=`asaRcy$&ChzZ$=Be5~7!s2<8>J@(P$Xb37_MaoA+m0i&Y PKK~VgU1Jh3q3r(y${Y)T literal 36896 zcmeHw3w%`7x$l}h08=CZUmyqr8X78$dGZuZfDJ<>7)fIBL5E3ZLZZn`oCgGJYB1U| zcN!1=ezwQfV~bW>Zf`x(d)kWC+JI<6OWRQCF}1Xf7JDa>wrQX?t(3X{Z$0*`*_jdQ zIrrS({oO2nlW(nWeXsScwZ6Uf+MB)ZH#V%zu-TYOnd}Ber9IaOOu4Ms`&pR*C}&IA zNAPzZo6Bjl5>;bz2c`;?f-FyzNk=wi?equ9&i|UwHeZw#)cLr+DoJmFH%=FL1$Dj& zs5;Oh>JPLh!~L0}JSw5$%K)1A`XpbU9j$T7w}LcID3M;0 zL1Sj!##(ZBH{n!wQT|d{rusL3@|Q;@cP=m8^)Fvp`@7xG-tw5g<}|7q4aa%-k(!D4 zwdHl#4%~7ik3BFqqoc#lyMgMj#*fl9_~r8)V>Bc`D)S2fKZf76_$|cmI{b?9qx8>D zJ@U=Z&V2idlF`7qe=B;k`ONEId$VBL8;;}4 zzxC&zPJZUz`5Rs|j=%AV>%VUM!Q>-X)!rc_js5$3pr(yQvomS%i{PMC^tEZ|*8)%N zx27ZLh*WZB!;z`*_ovbOw>0#9Y2=ru!M~ga|JgKpA3%dr*%<*|hNz?_e|spPLpgTF5gp5~k#KWo~YhJGN8{nOI)3&>N- zr#WS}Dn)q@eyQ5Em*c0fjRnXcJu3jSb9<(+t8NsC`y~FK0jIGk?3kR;GORPd#MjG7 z??XE&#qlG0Uj;96xsWc2~BJLGkR zyuP46z*@thAZu|qH~Z^ZFa)gu*03Yc6!Nld-jEBhg@ydEW_we;7X*R&mR9Dvy{M?Z zHS9tKDXCgqBV=(r5WJ!_Y~tM>k5bpr8ffx`8W{0a3a#KOOHd?9tS)p_HITYeogg7u z)|2Pz?{?KU-tB5|H#M_hxK8QwHhWvVzR-@Qpx4zBZdT|PcStL>`gf=bWE)t(8*Fs9 znl5tPzOmHP)Kap#*tqMiyI_ti7V#o)iMOS`btk!}+3Vw8LWwHY2E3{(s=~p>)oAf* zzptTbTR7kjHTiuIR$Q#Njl6faH{kO&S2a|-s$ET7%sOvqgF6^90s-_G$pwNJszY$+ zHjY}uyQq4nuf8$h_wgZ6>)X-f^SCy7gW(o3V_RFBtJNC_`e}@WcDl9~%3-qJ*A%Ms zd%`rpZZBQplKl<8!jtDUQCTX6T$R5yMKgJ~ZjvK@jP25eaxv7pk5(qhhen(L;l`)R%uf3Mai>cUR=C;+>yMmmKvk5$uVd7v)v^9m; zy49;)C5}S2dh_OrRqI_vjuN(R!}?XLUB!;#C1$qN%oRH*I%QyuMgKB*j!GF=jft|r ziF`J26Y!&dw6+qE_-*I?$2i3{FskcK)%=|wAt|4eaf)PY)aqYxTK&9Vp3bG!BjG67U}sp?Q~A=Sri*FS`3;E+U)gjKwdLEAUfn;a zb$hwIUUbU~iNc?galY4rSBuJmU8IdMMX$=`7JOBWpmSRAZL+-Cf}ivWLEm@3fV5Vp zHLFVf8Xk@`(}0GDgUl4u@H0&;W5XJLj)ouA@OBMfDDAWXcZG(Z-XRLKCZ{@;>>VOU zczG$~hyo2y>u{9{J4BA?Wmj>;QVlOJl^kEL;mJoTIXgs-Z{o%%X?pN7}hWBWDyC4!azJ)q&Q*6`gLzEH#WYWPbve4mEDOvCqU z_{%l?fQFY_1H3k-;b&;{!y5id4L_>ki!?lwzh_wcffmyelc+Nc>B zJZP3_BW7gqTV^?nm!CAtv;j3T_=s7ijhKmK^%`DSK*vR0Y&Z+j&1_|14mT6;T zWbl|-rVSCa-z?Kc2-;)-XO^P;Rl5pRklaZz*?K$`4t}?UwTS^m#4xTP)>TOS#fgHZ0|pmU5Y;yvS0{ zx0J84lrOWCFSeAkE#<%8ne5*`Tgtz;l;5$Ge`P8E+*1CrrTnU;{C!LLJC^dZmhv-} z@)MTw=T+I+wem^ion6^qBCu>UbQwa^4R{}dLeoX@f!sMA7;Wqo2|<4kr!TvUz=Ctm zgXc1xkzwb-*o{uxYtCcmLYG3o#Zthe_&@_6mx`a_SI5dN;AP=!Yn`1d58aNeGxAPo zvNO7J0`ReA7}qfmg28LqWhmRWfY;Pd{DV7CNts&kLD_V&B+?VBq6x9(6|<&usN5N8 z?{`MReUaLu&PZi9tp5~tuVZU5)qb!PiepE?p7Mpn++)P>mqm@( z3gc)ybQ+`4s)Z;zA9Hq1eu(OT>vP(Aqb&EiKB|vYE$lR6wo^Zf8b?8Oa3F&koXd?r zLScUa!C_S#@GKAkBjy%GtRV6h0s4>-(T4FJnA})GqzLT1o};&W7)DD zM-rng(jLP|8|X1cNqLWPmH>3e(AH(f+4f~o<7|@C2AsCOD2 zGVk|8HeJm8W|gSO z{-;~X!WBb>l-TR$UJ^>GiA@uLO{is5OTS`KmQD-4Z6=c~* z>|5XvVm~XyN^!;jw+JKJI36{QC#GjBkwwtGo$WCZTYgYu0tPm}8jB>xgA)@B*cqArjE~T2TmG7B|OdxuUUO5O4@IlzyUL~4Yj>*uQ7))|9 z{0(LZX|Mf7`oFvY0b!ZZ%O}LI%=)jPzSHQnoq8HNkC63eg#PWO{!Wne7)Pak-i4Tq zq#y1k{YS^>zt@z%R_G@#uKy_4zs0Pd&-EV_`qz{G5N~cDsYB>oVI0rxdKLLZbNe)P zd8QSWKqn?02HBlvRjtlJv6Z?TcI#*e%Fe0beHsn-Rp2t%RrUbP!2lDEd=`d^F;gv? zltPG**fsg{!dJC1r|pO6k&I5G&o+1+E_G5jP0{vixbRp$q+#}wJJGcZwYbt{^gAPb zgin2sSaErZR>I{lKk^eA__uR3sG^kv(c0tAXc%#;C-w>?p1Lm~PhMiaNz}L^U(?&)Z;8EDN5t+IO*vnfXxt<$?kc-MxKH#A`A784^U#Er!%Jhr#7n436DFpbV%MSVW|w~t z9760?A(pxv7cL6()L8`)<7Cu0nFtf>i7d&*&)<^L#eH+p<`t&ZHYr{OgW-8<`c6Cp ziCrrnhgqVj@152sBaN74n2gnYRM2%>n~QW|Ih}huh2nbR=c}FXf?o~|Z9e`O_6Yf3 zKz*ly`S>K*ULZ41k_m9BGOpS*?k^8!X7q96&>=8Rn{b+OuMs~t zZiX<9RB_|_c(5omjr;2e8P_L_n?=U06aG13lKeYJdW_>@0`g8;W*pap<#8=~h?T>S z?W*x#1ha6C2v*0r{{3eCPOksB2v+|BbgW0XV7;mTR<6Ha>KDPXUk{f3WAs;=@@EMB zX0YrR?KkWH`X_3CztBI6^snO{KSAoE#)%ciNo%k?p{dKGPT`@i66kcnQs>ZurH+Q6 z>|brs>S-(@R_7 zZx%*aBiKE-4Ue5LZFmZNJ;nezN%St}%s~4l%JX44knj-=OW`F9*JS*pD3%BMQh5}s zzSFsP9z2~pm30~ewo^|bM8=}EClUUJA%w4(7Q>3Pp>KvAc0Ne0qImMbE2Nf65lOnr zGQnkf`f_j!Pro$G<)DMdgua8QN)!4nH^pv%(@e2H0EZCk6=Kg9`Zf|-Qs{eWgD}{# zLi+hEw7J~0dV&A=A0e@8r z27843$5C&FzL&|&K{7!LeVa|={3NSuWu|G|C%JL|Xz2S93FOAj7RFgZUzus#x!-G{ zZ!Q_PLHOq8EWxukys_xK4?_g@_P-d@i~wi)_p2$?hIAv8qH zoNGms{!>EVAJP2e(6>bxVGVuTvC|QI%e3L^;Jfh9N8ujvV@&9SgHwgR+deq-t$-D2 zLLZO3<$2D|?ks-a45e6*GpyQa-{t|Q895y2UA1rX3ibK~vuQ_9iOHEF-W@|#ns_%~ z%I+$=6P_}A?Wf=ny%rW?ljEHknr(fZu?*|7ZxgX3#lG%!W()M4r(a+@I<^K~B2?vo zug5q{BUqn$hucf3DDF`XC$3G=T3$$X#XM{~_5Fm1h|wE0da)R2Uj-yx5b|VYw6b@> z!ALJIW1Za@eD4Vx7vlagl~~*seiwIs6rg&H#2wkm_I%pZ(yxWMc=K(eqkP-wC|{(# zp0vZY=w0gwetd08Cw&`x_^~-=f4no86xHNFp}WHOxmjksp&CA**fOI(w~Lm54Y_~p z#?{$4p-lmuY@j_%j&bra$Q>>Olui3fMqGHaLRa z%M^=^Bk)yX##Q6#l#{k+Ncp#Mwffd{c=j^Th(_~UnSMDktUDf`j`8jqz}@12xLX`r zld`+A@dzf?3ZY_5FD5=(&B%R?eDtQc!d(BqP}8-t&qQ1cMAGobvh1*SE29m4?QY>Q z;^&vlQ@=A+@wc-V6J8LtZr^Q}@6{52uy!mR;B5o%8+SiNWk$k^q-o7P_iaEPxO-0Wb-_9-Q zZOHX?$BsZ=L++=$Bbht^{0|J#ez7qz0GWJqVm^(6&cn_-lD0g{wcUNWgWY2cV+P~O zGfZ2*2sj_SPiQl2njuPEM_!~6`!Y-wPI&MgzFnN~W+xEN z2<^pV|AmK6OcKNvxfg%3!j!7V_P^s%W~|Ya>VS~s-Rn~9VkgDSK0G4G;94ME-_kx5 zwb7mtcB!eW;X!%=krYe^PUI?QEf4ZhePX?xi~;wf#vO>FA`h{1z<6 z*Bc2vD$ruYU<4-PVP~@Xejb>4ctwPS?x*ypMzoF2Z^N?t;Sjy1a|qiw?fuq0YcYJZ zPU@mT%eUYjej6SiQ_uZg2z=}XawC=#vHaf>A-4Xg4v&+%vR}lcf}%!R0qa(VK185*?oUk z)%_23=QwySplb3aVA1BqH{;d7{fMMKfbGtS5lozeCI zY_3+k05(>^?Q%wrq2FTb&^E+{xuEThIu`rq5v&2)Pq0t{ zZdic-1~r>^>8qHTqPdrY><75a$74gp^>u1Jq-s5U$*IZoT$BXAfZEujpws-581$!Z zNFMZEE8hYp_A>}~Mzep8Jjr;yRcK)%neXt%Pjo!gGjpY+sUW z(-oW9k;CYOz4Ys3&+~aa%AX5HYGaYglhMkPBG4pun+m8i_3j$dIx~BOJjnxM zPwWlxTyFF(CjP0hjlqE>(#7lW;*PWHvT_2^vbUB4V5~%7654CGs_0sIBSgpE{S~Ss z*`FdcUD-QG67AMSFKaRtzkxTTqS?1m8S)muTuk=c@dP{8g${vS{vv<~&uDA!n#jQo z5x6uGjzuWe82ypTfem!`K}(DF0ea>^u@bF{ho{HGd6kia6%lau#Lk}J#^!16z_^Zt z&p=S5a)?-P(M9y)nc+xPmdFj^`8y^Ae?JShuSBO6y*GFc(!@JC{GBiJEiYI-p2j^e z?tyU+jC)|*1LGbT_rSOZ{=e|RYWiRfKC@z%pUkQBw)g`(?fmOC{3|YY`s$3=lf%S! zaO~m>J05#5Twm`E1{=c7%{x_Hldr)az~^W1@t5UnW4NWx8?gHu?D)P6zM%tTaJfAn zl4MN~-*FKL@p%{IgE_3a5nrcSZoh-&d)Pud+rp~7K2O3xd&qAWAM6q2^m!D9@Aa?> zk4D9$ML~N@I2f`wy0?4np+>Jg=x*`aeJKosU3Rz6LjX;%cf6yx;Xh(;l29Y@jUx0#bEv7cd8cOn2kE&0`{fsq7SSh=l=!B4J^^^Mn{Zm1h_AfN*2>sAN}28#-zP1ZC~Gs=%gn1$^34;JyCrUR_H0JqC(Xp)Ec%j+ zBiM*fSB2bl$SJt9MU**o1si>yHsW^P$GA10VScInF;%)j+>MQpC1^ z*K>iy^|;1|+`drZh0uGvb>VHUh5)X~%n|T5J9NZ9K5s9e<^PX)$kpnGb6ApJ#yUzH zj484;z-6kzrqv(kcs2SH6(7OG%S7a~CW;f&^0dhy0}bu87D!^sQv&~e1al|v>gXt7 zlj6HN3MTQaAf&oNs{&P)a$*uCP17{%YLTa!tZ7<)s+P~0#N*di z+botjV;0N0LQL>XhR8KM)huVwm*r?jI*p+lUNlH1Nb?4g@}=3 zy(TNi!GF02GU$V)wDO_(6DC=FsoRWbfv(jJAnlc#`}d=Cb_lSrZukr`|;3Kbbo@e|yFR zwpo^;e^)R|439yqdiG60Ia6FRsr=Lw@I(KdCO?}SxG2lKZui~HJ`-dUxCw(8r*Ffa zozTmZuxB00cnKA$Gm~|bQU9?-)6mqfi7HolIZx9R@qPF-{E~0MdHB<#SvsXqvML{0=RSr4r7`@wCBs%;Eo`TL zwetTFICIEJS^_c2pC@rtBwp2ih5w(xQJZ9QkrPRM-~K+W!Kw69yeAUtTB<`W$|S$k zKe73FMBaXFhlr>LB< zIVH%=K+!hHa_?s|f+bs4~{pGv#X01mN47RY=59^MNZHgYISV$+F8v-`K?6@^Vyk8idW>6BX6$@xb3S}6}nK) zsox%QJ3@ZPwq}2wyLp-o?5}u=pg7#@*wN(iHnXceWtHrK(iJ&ywj+2Qo4Tk59sYA9xw-AsYBUz_VV!UmT3bw*sC9 z>;QBhj>q=`{u;0wF#2jdJ^*;|2%hF*yuS34c-#*7{Oj>}IpBgfzz=xr7xDOmfY1Ii z9;fesFTm5^Ucl9WlQ3}|c+y<}cn+`;aQ1KF@y`J+1^g!9CjpNF-Um1Y_#R$8xD@05 zTD(nB3ivQ!4dCwpLx9Wi8VSYMR=~$lri2#;d9vNRiP_rnY%?cMoU{kFUjr=pVDF1) zJ9yZXytPxNe?0diJ0^9o8?Rjcv65@%5hL-h!EfL9(MCX(2#?Dp+fR7nmx*MC@Vlow z9>0=s#+1CzWURh;!YxpYBGo^N-`%L^-*#8^ncfVdp!#p)_rjrgydBWw&j@h=RDT@5 z%U_AdpG~TN+N?hV{8v%WzsYXOUzDZfUyAm>&=ZgIZ={>`BW8Um>K{FtzP<+a{$uet z|AxEC-;_!6?a&=U{lVj+&rSP{&dkg&@g{Pb$3SCXoSXnOX)<58@kRhe?WFfp{@x#t z--o$4oW8?5LwL_VAk+AFv~93I3SAx`MZpY}gAO{|)maT^qVl{~+q; zgVyBVl4!$#i>xP zdeEA7WZZ80g63~A<}sF>TK&zu-VW4m)Zc~r8mh05b0S036Dkisn)7|2nMMei;&IzY zzF_+pjsxhn-G*Zj;B;ejC)|*1LGd}Kkor`-het^K%FOm z%M4G6a}7k9&Q74D>glWmy-a5>=w&(!fs)uQL&i=e#x53soiRjahAgY|5U?`liLNJ< z_;nRfoqs^rI!fyP>(A%neku>+I!Z}hQ?^P1bx-%0v_~$hS+^|HJ_99n-!IXOb#OG3 zVp*ZAxKyCcZ~Ch<2zE~q73w+DQrS>tXO(oEx}8<}@L!m?{kx^3;sZ2xL3mc5_U`2C*gpE!xE~)&hq97JM0n`N?0yo!7fp+{(ft9 z-P#Nk)VUmOvYzlZ;MMsXiXLaLEVmcb)`fkcuzjhc)KR#oB+N_23q+ZYXGxGO<0M6b zr0CPdg^t2f2S_h~n=dWRMV|2LI;pCbAY;?)vB`?H96#1? zEx?l;b^Vp(4BPLJGuHo)0iVkLr#SsptX29!U1webp5)UZre?DfBGZXZrcb}BQ@8iKkdkcjBiNvdsr~GM+-^!m= zQoi!17jnoCdo6muCgtBM*Gc@o8+v~y>GxUaFUR_X?0I~RXpjmMw2mXZTIX$+?X>}> z>PK4F*$_8orHBuy=n!rQV^pTC#~Je_w~u~JAm zL5@e7!|@`o|E-Yp*7G1Yr=jnZ^w#qrp8+xyv!HU(EnW0ThF;%oQ8ft(p%35eK!sL zIpC*`Z7!RZW?aorgWs42-@@hB zIDQJVp7VNJ8v2JMz4g4@(lS@N5k+Pw|QO^>`qq;&Zu6jBkHHgD3 zJ${#X;^p#${DGj$9d2Xv5^A$Igcr1yEL%b+UZx{%IF41E#abJ} zX{E$eEb6Q9iWh`=J>Uj!kTq~fZygg^Q5h_8hs3MkHF%Yc$=6rwTU+S`ZwJ9pOKV$0 z$W>TaSmG`!sizE2i^{kO4tmSA!CS;>rE1@H-YBof)#CQKw*dx1q^|^Kp-vdl$Q$Cq ztKU{CoEHkJAYLHXfVZJ)yMg8*qa2c&Nco|KA6M#X^5Brt z0B2d8-p8a_wT5?JqMfcCO`%4YkS0z1um_+TCFmY#Dq2iF)(ZTeO zKTo#0)CG+#RYgUsYtaa=y9I3lKTIv+jpr=B5RPh9N1@Wms&zQ)7JQnwE`*%BsMG}w zO)ex7t%O!IH^V`$wVNs`4VSTT4IM%3!U6Fv!zq!@HJgC)1pO{C3fP)Y;Bk3_Bju5@ zZgG+-pHhquEaqBs+s2B@^{XMIrgF8?v2Mf7t133QZeF{#+Ng2WRII}DUw*!;I2gFR z+XzMJANrfhlJlLKtQ~8ctTMUxKcKa+7snb~nUfD&PSlpt&`)_TG_?l* zX@<+P_NjLEa^jq6aUOL-G@Y_54yU%_=sfCVoTYX0sNt~=B%|nf=CL*MjB1P6^BqF{ z;m=K`Be;{5(J9m#W=V=OpU0YQ&JKA@b&BJj%geQMjfJC99oIZocJg`IRNZwkdXN-zLNZhPD6y{#b^{zC~? zLR_C%KMCn753}4YO;8Lo$Q-mt>J%ZW+PjU8a%YaZV30Yms^cYnf%_rz^gISjX54Fw zr1pgsv`f4KSe_;&2TD4>+CNtCnj~u4lSoc<8;LJO`;~sR&#d4+Nv`;;`sp1!dd8#p z)qb>sy^;~PYm9DH6~Ee_sRjK3QmMT@{!_K z&tZBczoO&9@JE#ur2TrDdx~E@uj!NgDlky&prrVeU7rVz+GO(IPgHo%%P(M|0OnuC zuiz{&>-?-klC}z1K$Tijd_24I& zs{M_Uzftl#6`_Py{S{KbRgdD|XyFGrPdYT`u10%6>hF^;FxXMI0Y9Z*#cz7wK-rG0 inRfG+QyKe_VvsdTuOgg^Oe+3r3d~3!w=h_W?0*A~vl$Tp diff --git a/build/tensor.o b/build/tensor.o index 6b4c10bf66d15872035f6b590cd27775c606e515..34580b528a55ea22ca99cfb41eb948f574c0a84b 100644 GIT binary patch literal 14272 zcmeHNeQ;FO6~CLHMvNqUS*w6s1qFmOfoKWX8i@~HEgu;n+E(dy*?lpq+1+&atpsaY z>x$FmnHk29emG7?`!$_WN852M^`mZLG(l6Fk&MkaOdBXVVUpHZm>Mh9o^#*1yJv6q z1;Ky%2X|(9_x;YfzjN-r=iGPi4c|95tSj;PgepGq6_I~as1WtFtS+ndVumIt%DU6kpJ`kSXx7a1y2|OSFIUfW%LbXoYOtFX z_T_H-`f8wq7H{Pg!9s8SGPKxhNSe`4<>Dt$q`G5_tTFmKpph}*G%kZz^CxC!`7R)Y zP8tKIug?(`PYs|v)3_{cjQB=hcZ`0}>KZP|xsmg?j2ta9cV$m<5siM|>G$ zJl!!a+OJl8PSdym26mfFqoyv5KdodznQvc;`lR*mR~XS`02=^wwm1jdApc+yZxop!;a04qTi( zju@mv(AhF)BRIUz>^cbJE<>T)2UCN7JWndsX6F@1t%AXwcx(-mk(;#6+Dfz282dvB zXg)WMsgKTgdAPD~;All9Zee4HHk?BO_M-cT6~uLuu2O@pSB8d31* zoq4{3#UD}&7weN;@1ZZ_MUCeT%Vc(=Z4Y>YYe21B(~7MDr_<_mTAfZSTp~`VRc>jA z?a=A3mG(tB{k8f(_*yl>CU7ro0-uI&>R$Zno9XN{rxM-(fjxkL-G+9QtAY){)Td2o zDmPw82$8RE`#>zWeIP!mZ6DFWE)9Hjcc@RpUd}g=9x2HfhaKbaq%Y4Y?Ff2}9(R)i zR_YhSDX`Ay>8Mo$Wra^PHfR!@m>~BR0n#EVN%P-9jFa= zq;OzN8G}qx@su%`>2i!=9}X4n`{d~i7ZEUngUPIc@-Qk76)M9~?i*cjlG}LTXvNo| z0$!wn@JPG6eJvyy%)t46$X7U~6RuSQ2xf{jh>y&LwrOL~cjR9m;~se9h|_c!_K4%) z5S#?vMZe(KRf%Urw0%sfV20$5QSw9y2iLiHJ_my4ideTI8J%4OoP45#eYLzsqH7E~ zZ^I52&;Jl3Q(smXa9Ob*EW@dMS)DPar+zEU*c7Q>B)ISb30#uVJ|x586fR`=6-Mpk zdxFmL5@`j3jo)i@J3c&~z!gD{?2)|cn!-g4C3NwW?u;h*ozd}h$8q=0Nb)&N$8#$N z#MA`d867W3Rk+`hsX8H3bs`--4)+=`x0oQJAz=mWNbDBBZN-xDgxF|Bj6#5{trOnO|%Zp;S01wpg)n-b=q7 z_lH@VP+LlN2L5rqO!}j#r0s7B-fsEr7R#RuMlFA=h(q9yKNt(+16=n9!{G>aTJ@HA zwp)G}h#XPgwA&WEp2}_#d&CR{ZGR*dw%R8f{mlL+Q`_xCFywYY4#^bbaU#D8$fn@= z6ys{OS`Y&E@*R<+1yj>%N7`B=ZhXPmVVQ#eM1D?5{?%A09&Lk=+};XdfHRSlGYwIi zh_p}kr6CxL#chAWa;@4j=&Sr&<94u>MG)Cc!c6*O)()51AC3b%HV5`liwu}hDv^Mm z`!U!t7chv36nehCp<%Uuantrx%uf01s%opMm)E4^Tg~@YR998kR;^g#;_fI3fwzj4 z@WVTo^_!XmTs3A9D}^7HIqpOF?zmp~+ADn*mY-F&2bPa{fa4F#n|MEI%&5G#q~Y8% zOZSz4ij-X?r9%80pGscx4-`xJ%_M)aK>o{|{7#bJh|Z|?rJG6>|05@mzi{S?1A6lf&bY9-v&7LSM^iIIAF~l`cV(O-2=}6j{JR;h%!e2d%#2gaS!}u z;z?mrLy@@z&o|&#Y@FZnzz=%hB@nX3crFGU{qLucQ+Aho=x_4CW5jdGIp6>k#k0#p z|92kvUy0|A87_~^d8p!_&(P1YnJzAK9jYIA@XUd=u2_Go0Y{$ovs@mzCSa8o=TKXl zaJM=E7ZHK=8#b@KW^l(8n4{yZq+0XfE2W)ayV>al+Vp=iAc77ssoM9Xs(tkEpVGMe!pS z7L^CcUjaX*Z=nzh;+O{tzaK#;h+`@#{AUP4L7e;3D^VB6lu`PJ2-p4mkZ|44j~Ng5 zvkdYVi|!|`GYaSaL6fob$h}@u>Jb#PoSQoJspv^YI93(@bB6Qz zdWP_&pv(8kuM@8PRP!~!c=&j}&2Zk|_ZfaM^QVl?x7c4)QGU*6IM2rj;Y&3?lY}cj z7h(m4p^V>su(5IX}!=f9kA z#eWZ0psXSs$A^#qT84A}A8I@*ANMeQ9uLnjoaf`~3||PgRXiVJ_;iMU$Z)JG{d4HP zM(^*%gscA2aEZ$Z$Nuv1tYJ9sugUO5%x-|;+)ta~JRhGVT=%nwaOLL#nwNcyhx@so z;oQ%EGMtb9+1MD0Zg)1}%C5>sKjC^jEM+*iyN%&|UgCu7cJC!z*;Vu6FdnR``F)7t z-0t1@0tNl#b{`~MxBEEZy4|N454Zaq!?CLTc^4NQC}?*&{1je}iwhLRukvCJ;i%8& z<#L8E2ASfiXZR9^|AuhAUoR1^?3U9xu!r&Re!anP-mkwgoR8fU--!__@BTp7!v{89JNehpXm%AzW{GinF7!_!W<&lXinSgpY- z+iJH(RXAt|Mb-9XQdA}4?hE|g68~l@s^AZbRu%lMZ@Ii7ruSQ7sj9X_yv<73JE09S zSuKI)1pJ|!JpXJDKQzZK(J?DWR!_A{)E^q@ow}dXyAWy^bM-mIt?b1^4U|-Lv{069;!eGRllP9BT!Dq6^XheoXYM;$D6@Wb}i z_`=FN_z{AF{_&#xSI@N|RCTU5Qwd2H)7vAbHbsC4(TTj{p9-AC?Uz5V|HtXm(W literal 17312 zcmeHO4QyQ1mA+%AA%u`|O3K2L&{s-mXcO0gG>V}GC*i?kiHVCdlosf8JToRP_KY<@ zBmrDNvq2iYwUm_>O51e%qwKbbv`rV(hAmC)#I4QNTg@(RWwaJ++N3f#(t4YqwWz!7 zIrp72bLP%_4l%o1RpLpLdGp=xo^$TG=bU@rea<8H_M7WMA;nin-JmMJ7*)z@soZY&j;oyF;6Gjpw?|QRRX+r@1jmJbZUhuG&~3>nw&{B}GFog~OXvWgO?W@tvJN{B z+q%1JPhD@1_JK}&)@@q{Kb@zX{$=}s0S4G5Cp6-!@b^mCUTj-eu%|-PKX>g>U=2>z zImKbmh~hCMfQlvh+ih9m>>D`8kqq0X&@d2Qm5pAI*udn?Kdry3yCAqw&{z@G7S6XN|J(^gUJK%86+bO>Vb zY)#X0-9*&s8?@XS(uackFO&WE3r`hJwSlomz}UkH$6d5$EjkKva6?soLhmE4&Y{ET z!DG2ifE;!m%qp*;oT_E3DlzaAnWcvvS0~^RFNA)@xwCH|w;6>xuG554J=)tAOXBQF zZ4pGYJ?7eDW_rT(7W$?ZP+N`!rUE#`H%V#}VT6JZWI+&)^tE}zY=LALsRWa6GSHx) zAYezogcu1;AFN7Ur!+&ISN|ERD2QA7W*T|#wGIym z`^s=jLR~}IidKLtNe4l8Pq`@~dthzF=?1$c41zdX+d>J|gQVcAP&SzhO+x+pknQb-^I)(QKEraS%CRoJhF&jKhxFxRP(ND21l0jjJtI zV_pL4NouyoWU(9*)x&S+KvS#rkNN5!Gm4eZ{EU-+EPKuNxGa|AXV;ItwEl5l{o{uI zVS57Ax%R~M_6f6CP6%}kI4QdHwE`SeEM*T-EF}zrD#dc*9EznB645^@#d4-xEGH_( z5+b1{y%w~?AGRlHPQ09%?7O3!Gm{k`VOV-zf^e4UY!uk?L-!KDxPux?9(Qy5sJpPRi^Nx`n&IJap$Yvq8h|CdjM-d3-?D-D$jQ+5BF z+rD9YsIo43tNu{mMwkbBycnviO>RX0@LMs5LeoFSKHkQmXpgw|2y6!Wwt5q)PqMJl zZ5_FKus8zCSYe<}pFP3Bh4=jECl($He*l|J@1(2pM0TQY4IXOAwGb9>eQY$UkBvt0 zfoJ+9pA#-v6c=zjH*A+Tbb9-uU{4qB|2;Tv*e-9i$^iwf#eR);T=Cy_i}O@^&E=0`*8#Od%| zm?-uHZW3PG<2?pe3m7ZHY20|I+Na+BqQM3MxWO0?z4kunTd>DMul?@7@mcaLn~ioj zHg)iL29-Z&Z09k&;0z8Q9_dlW_rzlorfc%tpBB&IQAUJ6JS_}~M;Q_N^F1x#)-KP`mVS5G(BOEK>(JvI7`Om3sIH!IIl#D|(G~1`lb)*|a&9ejbsZ@9t z59cqsbV8CqX;{Enyga?$2&dRpHM0b-2q3kGUZ{GlElUoNDkOzJZv>M663_ z;(}fm%*xjVbA`S+?{$He=SJqr%MTQ&S$thkeew;jHgw&e)pdWi(1WMka6*hzr_{Dg zGM4K~-Dl;JscbrjbV zchDB!r2{ga&t$;Q7A6(u90n-~1y9_vZQFI$+DJz}mCIYrjT;-A)^Et`UmLz~O;cmj z#>Q)|_HfT@N-ct$2KeA1(k-_~RGm@-h}FXfua(}H3O#s-3hikKUAFAJCHtXHtpXfB zc%`6|vX?hJUbpSS#r4TLpy;-b>sF=u@SEG~a@eff@1geZRJ9*0w?9nno6sQcU%!44 z%RfNvzprZlUAg^*v$Ri8djyrZnqy(jhHC=9%Q9XH z=|Sig|9_CaNd4zqey#PsZ{~4Mcvh-=)$3WT#_(6Yo>f2myjoE6VPbr)dUd`Jit-oN z;L8AcudrXo24H^;iTd;nNrk zaOs$Z?+*C*+1V9`UX{Wbg(l`LRp*LT~VGF^QboN5& z&vh#7*XpT;ULUJ4swB2IMcX=2-9|~MNT&8^HW!QEAC2$2KZ-jv1*rYtllf*~87$2hi8yvVqHQs6RvCJZ@T2TlmS%%df= ze?$3~!iV$m{Q}3mIpdoIzEt3?0+;y`CEOhE1BCNnJy|acKu6uR#Xld4^^9g6Y z=iv=p7X|QdHTYrTw+HaM3_jO^YykgJgU@wfpWw^*dBWiHc)uq2vX1?|!RPTF4&eXR z;PZHA0{DwDN#J6C@^~8vH~qib;PZH|5PUh_s|`Na?P~>Jjp9$bEz97K=gxC1r;IsW#2)?xc0O217HS+oE zkRfL!<;C+tj;tHU1upY9jLP6bJ7wMY0O6+I&4hEj(NFaWAxGAYTLdok9v3*eme;9C z!dd=3G_I3^kGITUhDO6>+IcbIEazd8^FhLKywX4G1upG$1TO1&oN!Zb58-mF1 zj@0{?0+)K95jfuRcz;Z|sdtQU)=TbEFAF(xUVbBRsrR%YhwIxP1YhRMA`D8nSTEPN z<%FB@f0M!I`gVH&KVk5>KBNNp4;g&UuSW%6#?O9(&*S~F;LGuT!{GCHp9$c*4(Y{KpJF+wTVOpECGt z|5pWH+JDgCv;E%{d};ry2A}Pp4&a9{8T|ZDc#Z!JKHI-q@TL7uT^&S=YQo;W(!cD!) z--}u*bz>u)J6uFK+AsC47P!>AS>Uohw-avay_awvFW2YK3pvuyIe|;P-x9d=^Yet8 zdP{_}UiR}Zg&aBFe;2saJ9?3JTh08wLb$2-4Z=;mZwfh5?;>;}TY|F|KC>-;|pK87lfYs}zt zoqr{Ozw%<}3>W)>>-^<}qu-?bjRv3V-Rdx&t>%YOL2 zkR$c}NZ?ZMrI+eq)zq2IgqwPI5N_(-Ddb4KpB1>&Yk6vQ9-%D0TIdEb{P26+&lxzs z$9>em`TgtH4V>SD4j4GU2c0xS1s>bZzo?NU0zH=Gxdl%nj zlHIY!TyjrNH74M1cvNFYHme#lY3~R8oeKU1j%tKI7)dt5-*T+iZw~4AeW`q7ZzkQF z%;fe$AC#2b746KxA4-ws@9fKh?l>emhAazfp*#KG{>!Mx6+ac8?P099ND%HXxzYoU z;2|Yk_;`wX=r~C~n<0J`;gmiI>FZ63AO6kX|{hxc8;!ACI`WKb0k4-PmQ%`@N11`5)`tQuvttQKm6P gy5qnV`b*tq*l+yQYqnqiZ_|tMoF}IIui5{90Gq&!asU7T diff --git a/csrc/cpu.cpp b/csrc/cpu.cpp index e75001d..cadf732 100644 --- a/csrc/cpu.cpp +++ b/csrc/cpu.cpp @@ -2,6 +2,7 @@ #include #include #include +#include void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { @@ -9,3 +10,24 @@ void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { result_data[i] = tensor1->data[i] + tensor2->data[i]; } } + +void sub_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { + + for (int i = 0; i < tensor1->size; i++) { + result_data[i] = tensor1->data[i] - tensor2->data[i]; + } +} + +void elementwise_mul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { + + for (int i = 0; i < tensor1->size; i++) { + result_data[i] = tensor1->data[i] * tensor2->data[i]; + } +} + +void pow_tensor_cpu(Tensor* tensor, float power) { + + for (int i = 0; i < tensor->size; i++) { + tensor->data[i] = powf(tensor->data[i], power); + } +} diff --git a/csrc/cpu.h b/csrc/cpu.h index d5ff611..99b88fe 100644 --- a/csrc/cpu.h +++ b/csrc/cpu.h @@ -1 +1,5 @@ -void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); \ No newline at end of file +void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); +void sub_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); +void elementwise_mul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); + +void pow_tensor_cpu(Tensor* tensor, float power); \ No newline at end of file diff --git a/csrc/cuda.cu b/csrc/cuda.cu index 3b00454..4664b40 100644 --- a/csrc/cuda.cu +++ b/csrc/cuda.cu @@ -1,6 +1,7 @@ #include "tensor.h" #include #include +#include #define THREADS_PER_BLOCK 128 @@ -54,3 +55,66 @@ __host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_da cudaDeviceSynchronize(); } +__global__ void sub_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < size) { + result_data[i] = data1[i] - data2[i]; + } +} + +__host__ void sub_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data) { + + int number_of_blocks = (tensor1->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + sub_tensor_cuda_kernel<<>>(tensor1->data, tensor2->data, result_data, tensor1->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + cudaDeviceSynchronize(); +} + +__global__ void elementwise_mul_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < size) { + result_data[i] = data1[i] * data2[i]; + } +} + +__host__ void elementwise_mul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data) { + + int number_of_blocks = (tensor1->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + elementwise_mul_tensor_cuda_kernel<<>>(tensor1->data, tensor2->data, result_data, tensor1->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + cudaDeviceSynchronize(); +} + +__global__ void pow_tensor_cuda_kernel(float* data, float power, int size) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < size) { + data[i] = powf(data[i], power); + } +} + +__host__ void pow_tensor_cuda(Tensor* tensor, float power) { + + int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + pow_tensor_cuda_kernel<<>>(tensor->data, power, tensor->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + cudaDeviceSynchronize(); +} + diff --git a/csrc/cuda.h b/csrc/cuda.h index 11cd5bf..ea30b04 100644 --- a/csrc/cuda.h +++ b/csrc/cuda.h @@ -3,8 +3,17 @@ __host__ void cpu_to_cuda(Tensor* tensor); __host__ void cuda_to_cpu(Tensor* tensor); + __global__ void add_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size); __host__ void add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data); + + __global__ void sub_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size); + __host__ void sub_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data); + + __global__ void elementwise_mul_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int size); + __host__ void elementwise_mul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data); + __global__ void pow_tensor_cuda_kernel(float* data, float power, int size); + __host__ void pow_tensor_cuda(Tensor* tensor, float power); #endif /* CUDA_KERNEL_H_ */ diff --git a/csrc/tensor.cpp b/csrc/tensor.cpp index 03b2b6a..5f2b26a 100644 --- a/csrc/tensor.cpp +++ b/csrc/tensor.cpp @@ -20,7 +20,7 @@ extern "C" { tensor->data = data; tensor->shape = shape; tensor->ndim = ndim; - + tensor->device = (char*)malloc(strlen(device) + 1); if (device != NULL) { strcpy(tensor->device, device); @@ -101,7 +101,6 @@ extern "C" { } Tensor* add_tensor(Tensor* tensor1, Tensor* tensor2) { - printf("Adding tensor\n"); if (tensor1->ndim != tensor2->ndim) { fprintf(stderr, "Tensors must have the same number of dimensions %d and %d for addition\n", tensor1->ndim, tensor2->ndim); exit(1); @@ -126,44 +125,6 @@ extern "C" { exit(1); } - printf("Size: %d\n", tensor1->size); - /*printf("Data: ["); - for (int i = 0; i < tensor1->size; i++) { - printf("%.2f", tensor1->data[i]); - if (i < tensor1->size - 1) { - printf(", "); - } - } - printf("]\n");*/ - - printf("Size: %d\n", tensor2->size); - /*printf("Data: ["); - for (int i = 0; i < tensor2->size; i++) { - printf("%.2f", tensor2->data[i]); - if (i < tensor2->size - 1) { - printf(", "); - } - } - printf("]\n");*/ - - printf("Shapes : ["); - for (int i = 0; i < tensor1->ndim; i++) { - printf("%d", tensor1->shape[i]); - if (i < tensor1->ndim - 1) { - printf(", "); - } - } - printf("]\n"); - - printf("Shapes : ["); - for (int i = 0; i < tensor2->ndim; i++) { - printf("%d", tensor2->shape[i]); - if (i < tensor2->ndim - 1) { - printf(", "); - } - } - printf("]\n"); - for (int i = 0; i < ndim; i++) { if (tensor1->shape[i] != tensor2->shape[i]) { fprintf(stderr, "Tensors must have the same shape %d and %d at index %d for addition\n", tensor1->shape[i], tensor2->shape[i], i); @@ -190,7 +151,6 @@ extern "C" { } Tensor* sub_tensor(Tensor* tensor1, Tensor* tensor2) { - printf("Subtracting tensor\n"); if (tensor1->ndim != tensor2->ndim) { fprintf(stderr, "Tensors must have the same number of dimensions %d and %d for subtraction\n", tensor1->ndim, tensor2->ndim); exit(1); @@ -201,8 +161,13 @@ extern "C" { exit(1); } - char* device = tensor1->device; - + char* device = (char*)malloc(strlen(tensor1->device) + 1); + if (device != NULL) { + strcpy(device, tensor1->device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } int ndim = tensor1->ndim; int* shape = (int*)malloc(ndim * sizeof(int)); if (shape == NULL) { @@ -210,44 +175,6 @@ extern "C" { exit(1); } - printf("Size: %d\n", tensor1->size); - printf("Data: ["); - for (int i = 0; i < tensor1->size; i++) { - printf("%.2f", tensor1->data[i]); - if (i < tensor1->size - 1) { - printf(", "); - } - } - printf("]\n"); - - printf("Size: %d\n", tensor2->size); - printf("Data: ["); - for (int i = 0; i < tensor2->size; i++) { - printf("%.2f", tensor2->data[i]); - if (i < tensor2->size - 1) { - printf(", "); - } - } - printf("]\n"); - - printf("Shapes : ["); - for (int i = 0; i < tensor1->ndim; i++) { - printf("%d", tensor1->shape[i]); - if (i < tensor1->ndim - 1) { - printf(", "); - } - } - printf("]\n"); - - printf("Shapes : ["); - for (int i = 0; i < tensor2->ndim; i++) { - printf("%d", tensor2->shape[i]); - if (i < tensor2->ndim - 1) { - printf(", "); - } - } - printf("]\n"); - for (int i = 0; i < ndim; i++) { if (tensor1->shape[i] != tensor2->shape[i]) { fprintf(stderr, "Tensors must have the same shape %d and %d at index %d for subtraction\n", tensor1->shape[i], tensor2->shape[i], i); @@ -256,21 +183,25 @@ extern "C" { shape[i] = tensor1->shape[i]; } - float* result_data = (float*)malloc(tensor1->size * sizeof(float)); - if (result_data == NULL) { - fprintf(stderr, "Memory allocation failed\n"); - exit(1); - } + if (strcmp(tensor1->device, "cuda") == 0) { - for (int i = 0; i < tensor1->size; i++) { - result_data[i] = tensor1->data[i] - tensor2->data[i]; + float* result_data; + cudaMalloc((void **)&result_data, tensor1->size * sizeof(float)); + sub_tensor_cuda(tensor1, tensor2, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(tensor1->size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + sub_tensor_cpu(tensor1, tensor2, result_data); + return create_tensor(result_data, shape, ndim, device); } - - return create_tensor(result_data, shape, ndim, device); } Tensor* elementwise_mul_tensor(Tensor* tensor1, Tensor* tensor2) { - printf("Elementwise multiplying tensor\n"); if (tensor1->ndim != tensor2->ndim) { fprintf(stderr, "Tensors must have the same number of dimensions %d and %d for element-wise multiplication\n", tensor1->ndim, tensor2->ndim); exit(1); @@ -281,8 +212,13 @@ extern "C" { exit(1); } - char* device = tensor1->device; - + char* device = (char*)malloc(strlen(tensor1->device) + 1); + if (device != NULL) { + strcpy(device, tensor1->device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } int ndim = tensor1->ndim; int* shape = (int*)malloc(ndim * sizeof(int)); if (shape == NULL) { @@ -290,44 +226,6 @@ extern "C" { exit(1); } - printf("Size: %d\n", tensor1->size); - printf("Data: ["); - for (int i = 0; i < tensor1->size; i++) { - printf("%.2f", tensor1->data[i]); - if (i < tensor1->size - 1) { - printf(", "); - } - } - printf("]\n"); - - printf("Size: %d\n", tensor2->size); - printf("Data: ["); - for (int i = 0; i < tensor2->size; i++) { - printf("%.2f", tensor2->data[i]); - if (i < tensor2->size - 1) { - printf(", "); - } - } - printf("]\n"); - - printf("Shapes : ["); - for (int i = 0; i < tensor1->ndim; i++) { - printf("%d", tensor1->shape[i]); - if (i < tensor1->ndim - 1) { - printf(", "); - } - } - printf("]\n"); - - printf("Shapes : ["); - for (int i = 0; i < tensor2->ndim; i++) { - printf("%d", tensor2->shape[i]); - if (i < tensor2->ndim - 1) { - printf(", "); - } - } - printf("]\n"); - for (int i = 0; i < ndim; i++) { if (tensor1->shape[i] != tensor2->shape[i]) { fprintf(stderr, "Tensors must have the same shape %d and %d at index %d for element-wise multiplication\n", tensor1->shape[i], tensor2->shape[i], i); @@ -336,17 +234,22 @@ extern "C" { shape[i] = tensor1->shape[i]; } - float* result_data = (float*)malloc(tensor1->size * sizeof(float)); - if (result_data == NULL) { - fprintf(stderr, "Memory allocation failed\n"); - exit(1); - } + if (strcmp(tensor1->device, "cuda") == 0) { - for (int i = 0; i < tensor1->size; i++) { - result_data[i] = tensor1->data[i] * tensor2->data[i]; + float* result_data; + cudaMalloc((void **)&result_data, tensor1->size * sizeof(float)); + elementwise_mul_tensor_cuda(tensor1, tensor2, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(tensor1->size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + elementwise_mul_tensor_cpu(tensor1, tensor2, result_data); + return create_tensor(result_data, shape, ndim, device); } - - return create_tensor(result_data, shape, ndim, device); } Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2) { @@ -401,50 +304,13 @@ extern "C" { return create_tensor(result_data, shape, ndim, device); } - Tensor* pow_tensor(Tensor* tensor, float power) { - printf("Powering tensor\n"); - char* device = tensor->device; - int ndim = tensor->ndim; - int* shape = (int*)malloc(ndim * sizeof(int)); - if (shape == NULL) { - fprintf(stderr, "Memory allocation failed\n"); - exit(1); + void pow_tensor(Tensor* tensor, float power) { + if (strcmp(tensor->device, "cuda") == 0) { + pow_tensor_cuda(tensor, power); + } + else { + pow_tensor_cpu(tensor, power); } - - printf("Size: %d\n", tensor->size); - printf("Data: ["); - for (int i = 0; i < tensor->size; i++) { - printf("%.2f", tensor->data[i]); - if (i < tensor->size - 1) { - printf(", "); - } - } - printf("]\n"); - - printf("Shapes : ["); - for (int i = 0; i < tensor->ndim; i++) { - printf("%d", tensor->shape[i]); - if (i < tensor->ndim - 1) { - printf(", "); - } - } - printf("]\n"); - - for (int i = 0; i < ndim; i++) { - shape[i] = tensor->shape[i]; - } - - float* result_data = (float*)malloc(tensor->size * sizeof(float)); - if (result_data == NULL) { - fprintf(stderr, "Memory allocation failed\n"); - exit(1); - } - - for (int i = 0; i < tensor->size; i++) { - result_data[i] = powf(tensor->data[i], power); - } - - return create_tensor(result_data, shape, ndim, device); } void reshape_tensor(Tensor* tensor, int* new_shape, int new_ndim) { diff --git a/csrc/tensor.h b/csrc/tensor.h index a075fb7..c87f2af 100644 --- a/csrc/tensor.h +++ b/csrc/tensor.h @@ -20,7 +20,7 @@ extern "C" { Tensor* elementwise_mul_tensor(Tensor* tensor1, Tensor* tensor2); void reshape_tensor(Tensor* tensor, int* new_shape, int new_ndim); Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2); - Tensor* pow_tensor(Tensor* tensor, float power); + void pow_tensor(Tensor* tensor, float power); void to_device(Tensor* tensor, char* device); } diff --git a/norch/__init__.py b/norch/__init__.py new file mode 100644 index 0000000..98156f7 --- /dev/null +++ b/norch/__init__.py @@ -0,0 +1 @@ +from norch.tensor import Tensor \ No newline at end of file diff --git a/norch/__pycache__/__init__.cpython-38.pyc b/norch/__pycache__/__init__.cpython-38.pyc new file mode 100644 index 0000000000000000000000000000000000000000..bf1db319d29c7a6e6b07713e7f40c25a58ed3d4a GIT binary patch literal 193 zcmWIL<>g`kg4Q^LGfuMShx0Q9OD1Mada@C18Ofkp7hnMa)17O#HIe&&bbB)z8Vv(|5^FF3nBN zE6Fd`FG@`=N=+Al}pP83g5+AQuP#cDvmt9awl%v3 z5v4L}P}Tb&4y`JMLnS!S3sNt=w3i;L{slhtP)=1(J@i`jy;&=cEFW0QzxVd-%$xUq z^XARJXTF;6ewfL)1fJi$H`ccf6Y>&Y_CFT9tinH71dv412H}yUoL(i8$=E|8V;-rq zofwaS#gcKb#CI6CB)vr{_FpiQ+^8Y$e-iMb7wDfw?ud2YYsf%6BPtJ8mDe+Djaqu0 z&3AFwD*S_Q06fB}B%FD%y0sezw4YE)yu_|!EXnj3Qd@-=$6^=gG%Q^WtD+X(Z+mJCqUWh#quTZ&3)f#-%gG{`drYix{i}85 zbd)@pQ~~Xd5Ps7Sg}}Gq1OQE_Dq>4g3j%RAB5cX zn@vv%41nve?bpYQYxJY$0M6*2>}zJZilwH}I*V`&A&y{%A4kp1nE>rgy_k63ErV9c zs5E}UoYW*j27w`9Lh2v@&nhoUH#DrML-++sjHn|BW*Hl_egr((^zI2WgOBa0eXg=- zXi9%`x=-g2Q;Ggsmvq;8v^kDmGs?kb4j@b+n2-m7=Fm2XHW{{r^zQmKZz#|TiwH8_ zhL&Yv>!yB?`~&D0ZspT~H|LHQ3P1EVBWGD$YcKUm%y{f2e8{Ng1-_kuTUbmR{n8m-?8-kyoyNNuS!c+wXww zj{YxQ`y}=G{Lv`G3o~16rUHo}{f4G~NNO1k=MhY%E2x=Z5<&!_5@H=J0WwKwx|pUG z)L}tC%`DBEx}a(mzEo($*brGz&B|+oj|V!PJxjOsbUw_HdS*K&;15}_8(Rx{H8 zO6iAwt69lLRdN!X4sYLpYZr!#ku}z{YOVJDL6*;=TJObM;fAMp9*U;Z2>WwIJ5)S{ zre(0Y7Jk9=*1m^`Qj0L8;;J|3-6y?KfGb)SCM)BbNW(XWORf1ua36%t)aI{JR^jItCS?y}L(Mwcs2S%}L9VLBb!sgc6 z}yKMU6~A*|zP^6!Bb!aKw8hX;VVdlwm8upIs?a1trKww`55%@%&8?fA8G z$7y1DRa8f!@Bs<_0x#e4mA?W02BeWdsC>%55Fx?u%-ps2k)Xh*+PY`nXJ*cv^E+qe z<5DSh@Vj{Va#McKasEk#qrVV^tN7BdL4+e%*V%S{>IheO4;wX+Ld`6CH}=Zo84xs(Br8-mmu(&UMUrZy-ip$`?7`8`Zm6OW08N zhv1a^;Qe=MAKkd6a}D0?4(dt=4c_i69W{7kw=Ve*?3(SQD^mW@k*nCLA&-B!wL9n~ zTitGd>-wNE>LqS82dKkV8R}C-kdnso!cAx9@s&xz|P+fCT!+T`~yJU-a6FDQ#j-)REjw$xf7@hbNVrQ zLA^{Zo47Q+JclW&fsWg0yPv9hzX3Svwo|3US4vZPew0$?}-nV!}Tb7qWUd5M^1~^9c zF(9Rgf0-5Cn2le_fQ9M&8IQ0`tl|4l621U6(S3r*rKAD_&jnnc}geX zuwtpm#^vKvojPgaV=!TejDo*Z?!x-%$Yxp10ilCK#{%o56n!bzVJ%Muov@SR@RnIF zPw9+oH6|@NfM=Z$HWVy(!h;EH)E-^Qnj{&YIU;LAq>{#lDVdk!uEMe^u&7lR zcID26nsq!TR9ZJ_zY*%skTDyF=64{<3?c!sGhYz^!tJ(PmFq;EJdYthn=4Q{%VP}0 z1juwDK7us2CfC8STC%wCJo54dA_mh8^cbgdOL}%44RcT4z-b_mEI^-m0n#nb}MQ( zsr2U^0OnAi_Mo3u%U0J&XLS5Yy*o;7ND0L{Axfi_WXM(xqgJALTO~a!Npd3#Ev$M* zoI01SLkD+>?@9eCa1JiSg!5m3feAhw3OwBBfDfx|*&VNDLO#;<7{-#cN#`j|-vl`d z(!y(z+?_Q@P52}!;4H#noQL-YaCm}1&!&)ISUVOc5L1$g0q?JoA@%Jb`aJI-R;~%) zzM>jL*cSlLfbp01{fQ5j-W7m3yuJhAcBhP5g@4Y$MBZLy;?05ss4$I3NRJ%HWLL~Qe2)r0|m1b=?k(HrDaQhS&Qp220 ze8&HhYWQpwE@T5L;v+q8N;X|*lFgfQ7cXE0P+SfR;82WwMQ?<&>hBDoigjn)j zA_foMqa5ji&T($ICn#i~P&`w_!j`e1IU;7Aj_E4RNw8!Pk)Vb`PErO7|4b6jU>PqV zIcX~hb7V5aroKR9N?e?y;frlMhC_5*^qJWdGNir})}df`O+jT03Yqu<_?xA`pv;9G z);z(Hu})05qk<`^Alw>{iTA8<0LA1qPZ23Ti%Xa6sbpxhTfdtaFiLSBsu=%LX%1wi zE<{@akbvoAdXjAUBnBHKuD&%lQ4ULdYX*9ciCC!Yhh!0mlGcaCUrg@1u?=S{Q)r=lOkf506C z&Yg(o=W&rO$5cOyFi!z|3I=z=IqnotZ+M$I;D}p}RPutG?6npw4IHQpeX!LSn{ z)B?UrVDR))e3Hq~vymOc|Ah6Wbj5LS#f@;$Jqk+exa^KstWs_+#ytHLJK#f?8+!Uc zZ0XDplV{f%SNR=y1D>1L+HyqoGkS!`+Xz(o^An+1JoCY4^%m0Eli|KBRaDLfc zK16g9-{Dn}bJLk_8)s>fRhcKAF_8ih+Jg44UB7wL1jY0+>qt5o`8ts|h+H8;j$M90 z5GJB4u7Cricu#qYpg%k@D@g5%8k?RmR7JycZFY z!s0GC6F+c3;r~JL(R&8fUNE7o=Uj32KG}QQ+*E>^HLf2Si=H}5)qfIKQrq#2ldd^?&f99_>p6y>2-XC-b@HI-ufFcq5v#J&^NAH$l&Ry zw+KrYol$o<&dh=f=?x6ESF;9k2$XY{#cOs>qaaPR9G~Zf|2_|*TYCJl(G|^r_o9sN zt8$lGOvrl^In6qDrWRG(-7T2^6_xhBMgZh(c*_8zi>|ywBN43j>g_(~2xg5*8%QO) ze?%fih&B@S>Wa0W{D(@88fqjH^A2bS5F2Z4FPZ@g_f=dQ>U}Zj=>YfG;eGA)hsMM- zt82ESSwl0S0~6{WIhMIK>FA1mcz=uH@RXhlB6l?|f8o3vEzzr#dG7XJX87X;G4lQc Di+rE| diff --git a/norch/tensor.py b/norch/tensor.py index 982e202..789c788 100644 --- a/norch/tensor.py +++ b/norch/tensor.py @@ -1,4 +1,5 @@ import ctypes +import os class CTensor(ctypes.Structure): _fields_ = [ @@ -11,7 +12,8 @@ class CTensor(ctypes.Structure): ] class Tensor: - _C = ctypes.CDLL("../build/libtensor.so") + os.path.abspath(os.curdir) + _C = ctypes.CDLL(os.path.join(os.path.abspath(os.curdir), "build/libtensor.so")) def __init__(self, data=None, device="cpu"): @@ -202,83 +204,8 @@ class Tensor: power = ctypes.c_float(power) Tensor._C.pow_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float] - Tensor._C.pow_tensor.restype = ctypes.POINTER(CTensor) + Tensor._C.pow_tensor.restype = None - result_tensor_ptr = Tensor._C.pow_tensor(self.tensor, power) + Tensor._C.pow_tensor(self.tensor, power) - result_data = Tensor() - result_data.tensor = result_tensor_ptr - result_data.shape = self.shape.copy() - result_data.device = self.device - result_data.ndim = 2 - - return result_data - - -def matrix_sum(matrix1, matrix2): - # Check if the matrices can be multiplied - if len(matrix1[0]) != len(matrix2): - raise ValueError("Matrices cannot be multiplied. Inner dimensions must match.") - - # Initialize the result matrix with zeros - result = [[0 for _ in range(len(matrix2[0]))] for _ in range(len(matrix1))] - - # Perform matrix multiplication - for i in range(len(matrix1)): - for j in range(len(matrix2[0])): - result[i][j] += matrix1[i][i] + matrix2[i][j] - - return result - -if __name__ == "__main__": - from tensor import Tensor - import time - import random - import numpy as np - - - a = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]).to("cuda") - b = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]).to("cuda") - c = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]).to("cuda") - - d = b+c - print(d) - #print(a ** 2) - - """#print(a) - N = 1000 - a = Tensor([[random.uniform(0, 1) for _ in range(N)] for _ in range(N)]) - b = Tensor([[random.uniform(0, 1) for _ in range(N)] for _ in range(N)]) - #b = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]) - #a = Tensor([[[[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]]]]) - #b = Tensor([[[[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]]]]) - ini = time.time() - c = a + b - - print("\n#####2######") - - fim = time.time() - - print(fim-ini) - - print("\n\n") - - - a = [[random.uniform(0, 1) for _ in range(N)] for _ in range(N)] - b = [[random.uniform(0, 1) for _ in range(N)] for _ in range(N)] - ini = time.time() - result_matrix = matrix_sum(a, b) - fim = time.time() - print(fim-ini) - - - print("\n\n") - - ini = time.time() - a = np.random.rand(N, N) - b = np.random.rand(N, N) - result_matrix = a + b - fim = time.time() - print(fim-ini) - -""" \ No newline at end of file + return self \ No newline at end of file diff --git a/test.py b/test.py new file mode 100644 index 0000000..ed5aff9 --- /dev/null +++ b/test.py @@ -0,0 +1,70 @@ + +def matrix_sum(matrix1, matrix2): + # Check if the matrices can be multiplied + if len(matrix1[0]) != len(matrix2): + raise ValueError("Matrices cannot be multiplied. Inner dimensions must match.") + + # Initialize the result matrix with zeros + result = [[0 for _ in range(len(matrix2[0]))] for _ in range(len(matrix1))] + + # Perform matrix multiplication + for i in range(len(matrix1)): + for j in range(len(matrix2[0])): + result[i][j] += matrix1[i][i] + matrix2[i][j] + + return result + +if __name__ == "__main__": + import norch + a = norch.Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]) + b = norch.Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]) + import time + import random + import numpy as np + + + #a = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda") + #b = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda") + #c = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda") + + #d = b-c + print(a ** 2) + #print(a ** 2) + + """#print(a) + N = 1000 + a = Tensor([[random.uniform(0, 1) for _ in range(N)] for _ in range(N)]) + b = Tensor([[random.uniform(0, 1) for _ in range(N)] for _ in range(N)]) + #b = Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]) + #a = Tensor([[[[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]]]]) + #b = Tensor([[[[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]], [[1, 2, 3], [1, 2, 3], [1, 2, 3], [1, 2, 3]]]]) + ini = time.time() + c = a + b + + print("\n#####2######") + + fim = time.time() + + print(fim-ini) + + print("\n\n") + + + a = [[random.uniform(0, 1) for _ in range(N)] for _ in range(N)] + b = [[random.uniform(0, 1) for _ in range(N)] for _ in range(N)] + ini = time.time() + result_matrix = matrix_sum(a, b) + fim = time.time() + print(fim-ini) + + + print("\n\n") + + ini = time.time() + a = np.random.rand(N, N) + b = np.random.rand(N, N) + result_matrix = a + b + fim = time.time() + print(fim-ini) + +""" \ No newline at end of file diff --git a/test.sh b/test.sh index f8353cf..9ac19d7 100755 --- a/test.sh +++ b/test.sh @@ -1,5 +1,4 @@ cd build make cd .. -cd norch -python3 tensor.py \ No newline at end of file +python3 test.py \ No newline at end of file