From d4751341bdbe55b8e5b2d67000c64ff9070a612b Mon Sep 17 00:00:00 2001 From: lucasdelimanogueira Date: Tue, 30 Apr 2024 13:22:27 -0300 Subject: [PATCH] scalar mul and sum tensor --- build/Makefile | 2 +- build/cpu.o | Bin 2824 -> 3248 bytes build/cuda.cu.o | Bin 35216 -> 43896 bytes build/libtensor.so | Bin 58920 -> 63840 bytes build/tensor.o | Bin 15880 -> 18496 bytes norch/__pycache__/tensor.cpython-38.pyc | Bin 5724 -> 6686 bytes norch/csrc/cpu.cpp | 63 ++++ norch/csrc/cpu.h | 14 + norch/csrc/cuda.cu | 226 ++++++++++++ norch/csrc/cuda.h | 28 ++ norch/csrc/tensor.cpp | 464 ++++++++++++++++++++++++ norch/csrc/tensor.h | 29 ++ norch/tensor.py | 62 +++- test.py | 4 +- 14 files changed, 877 insertions(+), 15 deletions(-) create mode 100644 norch/csrc/cpu.cpp create mode 100644 norch/csrc/cpu.h create mode 100644 norch/csrc/cuda.cu create mode 100644 norch/csrc/cuda.h create mode 100644 norch/csrc/tensor.cpp create mode 100644 norch/csrc/tensor.h diff --git a/build/Makefile b/build/Makefile index 7e003df..ebc69a7 100644 --- a/build/Makefile +++ b/build/Makefile @@ -7,7 +7,7 @@ CFLAGS = -Wall -Wextra -std=c++11 NVCCFLAGS = -std=c++11 # Directories -SRCDIR = ../csrc +SRCDIR = ../norch/csrc BUILDDIR = ../build TARGET = ../build/libtensor.so diff --git a/build/cpu.o b/build/cpu.o index 8eb61f6549ff4fbc80b9117295c230f2e5ea5abc..4220158a4b3ffa19bbdbab34a25da7299c6faa1c 100644 GIT binary patch delta 585 zcmeAW+aNhXgQZUH{ZtPmX64pWMeN z!Q=V_sEnc6^$+9gotrl@u4NSX%rEyV)T8sMM`!H|pza5gm_#_B+H^qLlo>wr3;6!{ z%pY+OWWx`z-V;E>9{|N&f#QM;6DG$n>rI}(yhJh$WU(usKpT@YFB?-V4?70~gA7n; z&SYH{X;uXwYtLj~7I`@k>j4mhKn93_>xl>InS7C5ZgU5V1fwO`JfQR}AQoa^WVpo$ zQp~`h1Eu!@*&uN?X0R|r9fV@IJ^3K3IO9hUDfkVlrUj}dWAaPZIvyj#;^f4f#G?4z z(wxctY!aMNhUU46B|zRJHrdHNY^M70QHCbPrMdAXsd>fuMe)f6r2%FkV0Hl1SvgP> zS%5YOOuouiZ3J`Q0jN8efHVlh-1h`cyoGsk5c@U8H}XmdS=3Vv~PxYH&sXLmLEEOxENw=d56b@HR}& dW#|rpdoJt})6@zQ}3Lbb@U%7ndDV9P{KtPO-@e zTpBRu0toXIP_hdmslctlG=X`t7q=bf2B@J2CI@oMPu{>S!TEpRg`k~a zA|*x?-QG{JT2QeCZT*U&6??G_J}N4r_F}77TcMTPf-mISd#SbMu61UgurndI_q*%+ z&g^~u_1J5zz0N-8?33(f(xu->yGx`?`wx|t%8An={dXLX^mh%h>C%q$oX9}!@qvXN zY$5|BGB9r_(tmRBjC4dA{F5<@veU+d!99leWXY1u+7d(3dloiit6{s7Dx#mYfvE#v zYF~fdqfL7Z*J_VbID4)3M67uy)||?ZweQE8kFbgC0{3^-9f|DfIu+Sfdz|@f&gfy%RT~+& zV<FaD$Vpw`0 z5z{)jJLwuKWty~1|1p7QIy4wnO;lQM)Zl;-;MT#`Tg9mnoXnIJv(6?-A2ULu4!Q{OcCt2-JSHZ-)eL#-~R zW*XQNt>w(>D3?3BTG>}=Y0BdI7Q%)y^4W4ns`&w0e3wLy&}D=za}=-vhn3y!$TL4n z@1-wEz82`17B4Cj>C-X3s=IR-}wO5V|HN7=~U1xNxYiMt75?L;t)wpS6 zM`y#vPOr#i{W)jZ_grt&Lbf|K&sf;e)X>?$S~5QeU|(jj4%B5;#YctSnc_VDzeH+C zDg?)cj1;T|4`Tysv5oLkkQ*+7T{$M^cRmUpo=z>(J$%bKneEAb0>ES0<8-hlC#0v& zuBD-~c}u4l{{JVT7PN#h=*f5gFaGN?dbJj?dvl`TUZ$1>n7V6Ii!9JDPTdkr?ZiL+Vm&S5ntw zjp4H(MJYBRMYu`>t%X#i*o0JY{2)Y}ViO|f_`6!icCgoix?UUH%zI2wCpB=^E zEO00Vrr{b#GLX=r;XWu}#j9+3e#zj3 zf+EV+7Ua?))>`;7Yb#W;(9_k_L1dSMC{8@XWgb*6*xAPg6O#AAOkaRQF`I~$l@(Mo zYf*UvnmnS(qBPNlE8#VP=jiXl!a2rk;!}zZ=4*>^jtT#>aGpN6z38=s!QDj`c`#D; zfRifh<}!DRq7{!U-J3w(a1%>6QP@3Yxr+_4qQaHn4V4t#Dk63q_9uEi~^ed z+If#m&PO!^1{|_&H5dSoZI2S| z{PO9`DjMwDBgG7Y*H?3(Zx1XgNXUnbjC(+mI0E^0lbsBbgM7^)b)-Yl^KlV?1G8Wq zT74O*7d%wSJige$>Z6_aZPgM&NCp^%DJtxYOEDG*OgK+5PIxSWb5PcQaFcmk(Ho>5 znY|L~W5;I>v7NK_Np~7pX(ViZiO=va64?4kxU45RCR`%aFBSvN4-$HTvY$ zp}8K%#27<={6L1Cb=;KNa-hUSsJQAw|d> zFnl%rXwRh42Hd!qG02I5aXM_~occ31w;y@`OQ@NxN0o2bZ z2YT84eNN_%Bo+#m@v<7}4HlkvjuiNY*>u9_Av z@fd`ERdZN`@xrqqC07;35{wuAizv(teQIXB1eOLa%tH~le$&~yd3lo8iQW9B5`cm9 z))YIvG9S~}_M&1o)MsKB=9$cS+;XmsN%M;(XALWzpGhyU+WDRd<>W#dyfHPmZ_F!? zy;Kgm>`Leg`6PXRcYq`Ul2%5L91vgG!^Aj|H+ld_T??3L*OSqc(q0d1`E zC9Xu)*hdd#Y8(b%+RIQW8g;xGe zvibfd1v=Jk_=XnvytgSXe}LxuJdp%>9Wneu$#3`Fqu9NTvJ_2F8fAN^K_04_WM0$I zUf!{FV`E-1DO}gSgrtVArK!GxC@pB-{IOEt_qd7On@`jH`7|~B9GMuXsiSEDk71lI z93*LhAgK&ZG9M;sz7>QXrU{owK>$=r{FA)Ibe7lyb7{g^GNWj1&~u(xGl@NvNhbJ0 zCzbrLCzDJC+((VdSyB-KYBWyQX4oaz&rLsHdF%l_6+?7umsJ$wMrdGa4+*SJxR`;3{ zi7smeQBNdNwUXq!eU+rZdLq$UiO?-pFNG${S{-Y0W?-9XorBHYZvTnIpF$CUiRIX5 z4=r2}4Lqp)l^hRGo;MM-mZYgGlsnx)qGmcfpvZhg%})0d zRW&J@?(n?|iH@)IyrrlK3UyW*6?f36BrgcHHgu*M&ABZGMGX?0^@K&ehUB_^*TC44 z0)3}D0DX55e?st(?D2WK6;B|%O(_dj-L6ywLU+K++zz|?OPZxtD#RVGB{}Nt3Uyav zXB%{U(8&$d64rEsZI(Oqw#`U?s7z?88t-uUt%mHaOB~W9XzCn_hnR{Cg#pj?#E@i2 zfH?^~?tlr}Ye-Q04YV-inMeLJ?VHf%Z%Orr>U~27qah7kF{Y6T8#)iB z*sEVoNjR9&04(1m^Fs#1agrYXD7BI0E^=8d8ywOfNp=-YwTEabTd*i+I{96y5QsMK@NFvfH$ATtbU?tp`!K>e! z*^d{wuC;~kaTr_D0s$}7b45}6NkXMXrX26v#2G54oxb2C^Ty^|77^)gQ${E#g#vyk z=n}v8Tv{j)@Z=G7w}mAwc1wd6=2={nKp&KH0?RYlvc;uU@&_4k{Yz46bq-NqkPHRY zo@lms8SYGS`yWpmvDziQ(G*`X(w;;p@({td}K|F(-H^Qz|wp)ubZeznDyefxyZGk_1(Jb6)b)woSKv zZFgGBiTb%cSq+mksOt@6oOLkMx`C+g+f#D_^~q&E?}zpZf9T;nhw@)=#ZmP=drEHQ zRHBx?ZFl7cpmW^rw$7&NF?)NK|Cn9b>s}D9_l0uPEpC5Ha*{hRIi2;^1u`zCN*4|4 zvBJa&YdV&t7(Z0f)*i~sUmIZW*0p4pv{yZyr~JDG0=$dT7Al1jrB{r^P+)YAW6PHo zN^^=?&(fEr6}hbT#zzdUDq~R*JA30rqw;o1#FNKNKTS>Ok~mJekCnvDq}xezMIyuSvQ|4e11G}`@G{5q)#h~eBKk% zP-yt>WQNz?L)Wgkj>7G+aB%aA3YwY$4;OGJCZyJC+aX<>Y?M_txN?Fa51`%b#g%TY zcbL&t1++1$+k6hmn6maUgWVqcH30TxqT{SiX7& zl|D{n?G39k|H_|oi%}KG2ig+6S-8#Ye8as`QwHl^Q%tMa<7=F*t=jXe5%$Joq(6fE zioLtW8MSB+mqzUc82LJP3C9=4X)%u96oxCK9v<&NpvG3?T`W22F-38}p?=QUcu;P)P;rK%Swy%{W;BRs~ z89YVGx?ahKpB|KEMq2F3i#j{B^Z*naQwJ zAa(P>4j?!BGUtBA=LRo#q;EK`@0{pbA#ed*so?t~I(|FHf6r61joUfOamG_{jN|`^ z;pyc2JXT$ya<~BI5BJs%_qGrBcJiJ>?QCt^vbM9{>-7eMYeL?^s&!vd8e*q6I;i&1 zEg3#_EX40MwSoHjwe0bh3|6y#n+r~wns04sY_9L<>{?Sl^7{@uxyd~E%=%rDyS2Wv zd1J?>_WH)IriO*#CECftwM6)qB7<$_Ykj;NzOHDW|LR-2He5|@ZA>rj`_SM!t?CGe zO}fME!;L%ssh<{Bzv&E{(pJoV)MiQkPHg(LJ3s}yqpgB{)MjBv;IM4AB%c(xe)hYz z3L0R^o1Mu)z`xbD>n>j|W?yY~vdCr&y9PAY@cF_^C(K$mSFjVnIRMAEz#Ooc1tq8$$G^rN9-Vng#1ClI8`py8P-f@SOrK`g>5oXTu!B4;9~Hf2>291KDL1!3r|c(LTt%{FCZW+=A$vj z7?qFP$H4asxR{UE1$?I9&`ANGM50$BJ0mb~V@8koQ;x@H2zM2bun{w)ZJ2Xh%#gwt zZImC*P!y&J??On(5K}Z>z{L#p3b`<4} zmjw3_J(vh|^x$m*E(Tw~e789#mvL@>fhaY(?2Gl7@*&jf~; ziBkeD7QhDrE@tQ=$I&6oQtXo}^!X!{jyr&dWCJ{kop?qDS3o24;LxkTQo!d6_zN5_ z1q`3%bon8U>-MvOha@G%TCfjafS|7{%;k)DL%n0*k6ke|VsShq;9`c(3itwT+yux! z&?5HkR;OL+Sh&afCN%4$K!pwkKNzxq`$)7*{1~=(UY$QT=b-bKP98hMR4fT^9=h6 zJb8|Dbx%+e>1#nt%+eHg!!{@F!r$n9{5d;5Cr=8v=*e#dd@+ondvY`XD+J7WM8MZ@ zT=xyHSET(hJhs}1&V4b4p(6g^A3wlUj>iYESir>q?h)`M;%c)rHRc%(Fb57jvXvay zBbXTzBYQcH1B(Hi8^cbXBc?dQeiSORhBM+Lm@nXB1P=*#o#1Fz+LZz33V1%p^#G8E z)C9P;+Qdij>==g5j)7n1czgh<{B06(F@U*O@F-adOg+`>IYak!3AzdNF9MEBs$N8o zj)6bPaZI(CiDx*D4&4aHx0f;hKquo2H31Xpy)p2QgaO1_`DD}pG5ZPpn=%H50o;Se zfvnLn7-<~G2*k@NhvQcwhz&3;&dy=YJR!5jz~^&3iUV8+hdu>M#~5I1oS{$n06k;i z-{m;kpDEbCZ>0U`@W9W-8M5$K!WYKChlBxU2m>4+Il!>}Pvh)(H2YBh7q7}7O)m@rdK~r&l{FAyo^8a1b4zJqCUQ$5D0}9Qpt& z1$?=HKP%w$X14M>fi!)$5TrE1+)QInedjjV)69u(qseSfw4z$y#|rVOLC4K-lxaG! z3JN&#^lhwVtf9vl)&KGUIe6%sRu9H$#Wtr1c}&1@l#f5)BzlRZ2D&JKr}YG2)83vxxcbs!V(&Gj{d{+k980v@VN1 zeb*H8maG_9vYcJKD=h8nVx>E#NUQg<_0SyYVoyV}WiLAq%@FXuJ4Mp7$c5BSwI)%0G44V(^jh({(8a%f%oVJCR&p!TshL?3I`}@v`=8WE$ U;GRS++J*fY(&;DJ`u^ho0sXMd5C8xG delta 7767 zcmZ{p33L?4nSiTmPL1Z$NTb$;M#mrqX~f*4kpLkuhX8XJ8*F6BvJAo;!UzeB5XT6E z9d3kvf`J#7Bg+|_hdd-<6Z{hC}5|EsFIt7qq&CM@uY@9x)aoUXWvmc+(IC(rmGryeubBye$W)hMMK^WK2*2x(3L1uv{&2r4RV*tM(bISnk zB1}Jh0Dnj3R)xm}aiNA8WC%oQ9+A0i0KX@5=K$WWWq#}f_zz`HPeNJ3+H9m9c3~t^ z6b(Ji%whjF<1QZbf5}|Pj9#u^4uB>cCn4Y)|=N=o+Eb zu`tbW%IAsujELpP23}jNv}qG*9YQ-zXj6R_kE+jOllA-CRKLGX_1Pr(nNR9poiOH~ zwJYg?wA7$qKn<9478Em3R|GTYZ&Q79;$2@khB<)}$O~M6%|Y{6CO%IOR_K_>E7j{4UI&o;Bv_5K@#V7FtEiA5& zmV5Yoi7r#A8zb{aOO{fh|D^MPp*J_(^h=Y1`~<$B{_gU8?jk5{h&n&#L$HwBQ8SbV zD;vtPMBx5j*{(9D}#}2KLvUCZ6G>e3oh9~ z9rKJ(vIDePwTDFwF5MlfmBI!j<*%I2gyD@9&X`@cMr-*Tu`4=E{}&^dm4i4Gn$0B=6p0rF=mvGm)+@&*ZvP1zah zD&dQX4WJn!P2P+`#X`ONnZN=oo4j$KB9H^6O(vl8YwqG$tzO|{p;rN~yZpSt$CBC8 zX&7K1qynarY#;0j40J$eZh!-3sAH-zI@g(Sa()4c*;#YQ*E~11HVO{)fJ0^AX||h; zeAVi#Fs`|TO!C9x<{bPLyxQ!=KDgdopMK96G7UGaF^y%+hTMlRNTCro+7m4H2)UljBG(8V(dCAWz^`!X(`Z=88$-T8aFhW#!bOHLfl$K`glHar z(0^~FX~EK_5p6FG>%i2AJcewWupupAF^mf}MyvKC;?NZPy)NV%+I9tDauS+aGIBmQ zW<|mz5(;{A5c&f#bfKXkZvnDxHNvhI7kR@7XIq9Eiz~fDk?HpcK4~e{I%4#lld?3o zAZ_Ixjc%(pEqq6l6T6}fwJ&Um6L%n(i*th`u{#XaOC!@uqTaE{73@hx#X)};c18E3 zV$Io9Zzy~aV=KXOq~sv>&wF{rYDAXm0&Vky>olc-U=JD=^X{j%of>Kz0%RLzhPRh) zj28wf>ol44$nYO0Wr3=@gpAeN?vOMR)Z|wz))1o}O+l%PHFUHpkuMayXTjl6sNp0b zN06;I(=w)Y*{b6vlE2_=rfmc=TmNQ3-$%A6%J2oeEvV4?My9O=rP*p7cu1(xQXcS6 zb(9Aq&s*{<8!d)qqcUwioMxP0Hgp(lG3>Oxk%?_%kYqcPiEI~;)h0O>=LRK5dT@_s zMz}E$$#R%o!9|u-SE#}Pw_Xb6HaN&vI-8W6GW>=1XC3;FwRi)=4w;h!$6F8RvtpZ3BnlIj zf2DT^0tLtequVCyvtz%~6~$m%+o7C-aAPblC62KV=?o}6Vn%qeb!xaVR0WT<&urirUDDf{V>oENL%XyHosz8@V`|`ksRW5uMj;tysEf_|;oR7D| zO^VKduERyb82;H&LcYc@y(1so9Wz0@x&Uv2X&o+Xh8?Si!6zL?xUf14KLrz355pPo zV703#Odsy7Okz*?30C`S9CCdUkLbj@=)ZyT5m zLM2Lt0@Rw&7tl=uTjbY|Uj44<-xc)%UME%AY5$nMT57o` z^$k!%IrDgxZw2J|sLn#pB?r8+?lr^r<)$io|W4&DNB6kF zM!st(;sy?S%>v!aN1=-AhbUKud79c9@*LOI@Nw$oKK82`P3O4An|#PrjXRt-@w8PQ zr7MH=Zs$Afi`-)e=ec|vx`y*!dAKjh=g;Lc9rO*7j?jo$v>CP)OIte@tZ4N6{h|5c zNFd?acn8DYP3hp-=7#q-83A6Mk?_2>k>D4gaBBwqbMvc-zAY9E&TX&2N8A1c=eN7D z7e3uC;d1z9yAM9uZnm7^gTf}q2cvgL*blRJ_+Zoyv*n>63&O!25?rSANggn>D2Ml{ zMHp`G@L5LC>;HSGly104wPQQY)Z2ueKJe}|Tgv!)VRpk|$4)7GxM0yDNGss2osxz9 zIH3C53%xrfxKI66fo_-2!ghnI?6+`(%AVO}wu}+fk(Nw1P>l3(`??(a`^L;3AF|TrZlhWdl2X4I3K@b8UWlvDV zvQ$iI1}vvR&lRL%O7jFhQt-J@;9`>N1TH3XUEpFe_Y@u-uv&BZONWi3nB`c3PlWeD z8p?j@P{(qvAQh9G2fOJIEES|;lB=MJO4kZfF{J}=5v2GbK`JJ8o^y69;29!tTtXzU8_DBIoigJ}~HF*#| z5Bk0@#aD4DOHfSyM*^QrW$HAH;J?SPUTOp$;arW)O=YaC;+*-PLJ!r?MTN&vDTOtH z;667<<_`XXVvfXntVG~q0(AnPD)|4DbLM{>q6=S*=+9iBX5gg?td#J1PtKqfdS93D z33iTO6m-Pk{zKqma4Gz)%6vXf4>g7lIA_6)7x;C~)!4#R#!8;Dmrx!IQz%d!Er8|w zBz%ya<3HmYj^vShPvBy38T{4G^rs0v_i@gGn<(%@oU6gHJB*bZ3Lh9;83guA^es=% z@wI$+lN?;Pz{TL+75H@8iJI_VI9G#UqsNLPO{p1YOIi&sTj6oKb^flPfEy%d>EWCO zI)ff+$5WF!m``P_%o>Ex75tA!^jCG&QPUtDY)W#pUGa~$4Z?eb4#rVk)&GG<`H#~L z&s#~3F7OV1JqUkK=%7aE;DbkXK>h!BlAk|w{|_ja|5Z?@Kow{miUykk@%d-wJb4OA zIA?y?!&uF@OwgYx@Zh)larxn^KEcO-EBFq7Y7ky8bnt}G!9s!068OIf{7IQZ-Juaq ziDOFS!gi?DaPZJ790lj$HP{Ke5BuWk2gT!j1Iw~yhx*+xTi_e(cIQ_+v3t4RX&W8;qkXC14z*nNSsj@7TVEOtQ)tJtA_k`xM@`C$)6R%&EU zpU$d*7j7Jpvehq^MM-A$kakqcRzJ&53R2dz!&oY9ksTv+HA&hB%c)fTw)r1H$^u#o zM~+HU)bFtv-|8|YcBtRq&m?hmb~}>zGDqV254U3MhUQ~q=u96w79#IB;QFzssiswl z&trw0<541|v!8&WyHg7-$K5Ge{FnKeu>W`k$x4T7$E#Cja%h6=6BQ|Sd_`}E4V~%G za>7mingKm0s!jK~VCE#aej-Y0>%nodnk-xovuO2vJ@nA3Ydu_JkjhR~lYLZnsyd~$ zUJ1!9dpvW>P4?9HM|GQC?xX6bt4-&*UbhD7PDf2=*DGeV8aQw|>a2ZUx%`~3LH0|& e2X3A&GkF#&l5`{F{3vR=zd^ao!1QmV{Qn0EQeh?rz@6CLz4PWPt>Ngs`u90g*sh2nvz{QR_>WkQWloo0|;=gc6Nf zVj4?ZZT0CJT5YklEp4^M)@raQXsM!Ai`BMKu?>Yjjfj@E>i)lHX71g)H@gtDeV)(r z`#bq;=FXXO&YU^t%$b*a_wIK8%4MpeFfCcw1&m70SprffD+Y$-1+p?$!cy^nI-ACE zt&yr>SDljtNXuDG4oW(6o-66J zoZ6tY_RUc3hiR`;(jD9<_?YuY$WTaHuC@0G(2-t7nxs0}8rfbvJ9Nr^Yq@Nm>ymho z&n&;ViqzdtvRI@5be%2BRDYs){sT|^;b*UX@a5;)zk6)U^&Q6x?Sz_#&vbkUpLE`q z<<0-R^7Iuu7*n3VIdk{2uJ5Ma*Knc2esrIsNwxCs5kR#oWFZ!ElR z`TI96y6r2^J^XFuv9ypcA*nhfY+p3Q8**9(+oIY>+y%#-g?t@ky2g7D*1w%eH4tz;sZ2T(F z#^OVIB)gf@mjQ@{|3^4dEcoql+O@}le?1QW{y6Z5;=sQehyRUn@OuEy#K)Yz4?wJb zoyOpe1-~&J8%eFZDJL2#eiUZHbNQ)&;BL;jdxFrt$lsNblI~@3!(}XyD?v5kp ztT=XodSdb~hy&jVc&vWa^LEqOf-)c=y-JLczQ^TEXIEY(%JW6Weyb9nbQWAC02fPq zq9gbhmk9V)iC-r1?d8$%w^1bgq~!vBSmIYmd8YUU;L{SmUCQ&oB?91+@PA4EZ%O_I zQl8Ng?xY}w$l7Iij0zO^BALIU+&hEZkz^|9^Cdual3aCiup-gG2bIR8b7{}rj{27%5#E&E$ZfgefJ&%4~u)0vwBAkr4ezZKWiUyGcSo|pQ$ zO1As{I*!gX`7zUR{BcVJT=SEKQva_>`)c~%0zFIz|3^L*&F4PZ?u+LM_`4!w3uU`& z*9w3(H@+v~ldlwT&7Yq`k>vl?90C8H$XKE5Z&{@PX#MSpBhOP(Zw-?1{nG9~m;KF< zir4IQuf*>y7yRpGyEn;p|0eBTEA{+y3GbHu(&V{X@_&{V0!aHL{wuhqeoa}%!R)xm z*fy!>FUfJO;j=(sV{>aupgmL_3Z%qEcUg643F?BISxvC6I#d@3)wQ&@23cE2sGT)e zH#N1^u=WsI4YK;{f{md%l!8rlEtJ)?ZDt$mLIFUVdAYfb5%*Pf%~Ta2$!o3&)HGZZ zsIP8pVxd-uzp1eXQe5tG)wFd403m|PqIIHG4hMi&v~}q4>e^aOu==)OV@s%BG$iB$ zl`n!3#c{1OP+3pSd5s7W&Z54(qgi85yduRYUXqcqHo^#%?HwDUxoE{GE)j_$_yOob z>#L}*YpQFmYYAP~*j^WC?r4hebr;93irI#{C?2x8z4n^wrs|-nDfcjFR4esdA#PU; ztO$ETYOW4N^}$sV1&Ly+tEbq66JaS(+FGxRXvO7=6r#%52s+kf@dS4WC2UL=oQ7Od08Qe{VbLpdC=ytSpiabriYI@H+O0vfm5 z=yK%L*VF}D>Y6I+*9O)G8o2{6uM4fLZV&l`K@2(K3ji;y4Z+Jda?nyfWY%tOsc8td zw(v1r)pA{9OKo6H9ek0**x1<_XsZjhx6&XFZ4PX5^7f6M8r^=QtdH_glYD{8J>Gck z6U~i!gZ`Gra0DYF`pxltCd(RexM~`Oct)EUaReW3<`pfCp;fK59TWjC_ZA0ad_myB zSUjrSlQYZ%Zy&%aIR-iEb%1qV+ms%@Wu}DPV2&4?khF7cn*dzu<@ALjL1Tj z^|A*(_b~bqSWdRjP3SEe#=(6iqiRia)2@-z+_H|A8Zn}x!TP)Fen)yA5dKs(3dWCD)HFq>kY>}(9N z<>lo8Poa~QUv}BzB`X51LJwQMa>bJJfV#3aG*f@$#X&#fXj){%Q)cdC-zIJ|> z#Hkq@m?bWD%JVr5ZkOxSJ5BIi5`XV&f=9P!V}k~`JmTfGLk4(`0spiC?l8dnrQQ_4#u?z_2SkC)O?6sw42Ua&W67?k`~h)E z@TEEwXK@C&+ye005(6At5j~X|;AAr`RSbwL;=fD-BCi|}1%l5Mm;C=_1EN6i<>He6 zUq2uU1gCR8Ej0{?D}vLRpqAPO#1+BmqovLPaYZn>)#HFJ16*!lIedo!uK7Au>@>hJ z!1Q#d0iL2maSm>P%Plys-EDwZ81VNP;0P#s+G~Ip>rfo(8sHZh;0F!xJOli&0iI@n zA2GlYNc1#dfNSTmR5@sXrxQS?p#gD`h)iy6IpDMbF1O|aQTj9CQ!Ld|k^!En0g-1I z;KuWy@do&41HQum*Y@yKoo|5Cxw4j=2DseX@~RR8JXb?TUS@#LHNYzja5`t#Ql$Yt zMgt?;rvW~}0PixuCmP^84DeY7_)Y_SvH^al z0Y1e5-(`SnXBJew+W?X)bQbo?)^!>ypxv|UVct5)57fJv3+`( z7G5Wh{ZKE{g6rh5yY(_HXipydyk4e-*U4iy=w+HrPaeBgFVh0%TzZ)n3MY?E)629#IC(5nFVh0xt&k2(SN;66Dazxmucce|MfCWjOf2!9>>eq>SdY;(SN;66CnDpmuccd z|MfC0(9wUrOcNdYua{|pL;v+MOt$Lt$Mip#OTA78vNi zUZ#Z(`mdJ@d0ElR^LY8K_q6^ydHFTH?BeC;^s<|m_vvNY$iepiGf4aTDn%aNURQ|oG{GzG+Yg73NQ~4LB@M-AG;2fbWsI5G?qk?^c)y;Qb4-umwHY2iL#2G zZAU74Iu7?#9jfSA)eq@k!R~x;CZ>YNO3?7&l?aLzJ^s@b-P=yHj=A9Wg`a|3FBtCo z8O=kF(qaH~KPI|IiH-n+R-)?#U2hc{$dH$ZD|#yVb#K+7p2~bc`$@2Gqem6p{R&+l zhmaLLk8^^aN+;UwXQ4TQzRU@Bpm>Q8rxPMtI8er)IXK`SM7!)biTJml#1t`uM4Ut- zLOa4=(T@+6E&3>E6l$6f5qm=u5kqSOIppLOIyW>_0e6^*xmZ-Tp!4_^*2X zhk&(jKqWC2M9GVF74^9+2@5N(t=ib4+*|}vc0xx&TQ5NY2mB{Ub70;*sPy=Uy0;Cn z&P5u(-m0OBn&&I_y{=X$kN5gd^r*f5A(A!QAnV~?)?0PL{)5MR{3p7bhm_+#hhe9YOc0l~RZ}T#w|WN63!7 zt03@G75iR;-nlD0uD~=T$Im5-v!c&e0eQ$7KK0LVSQwtDVGP}9ecp_UKL3e_X{r1( zS`QOC>HGK;!N-WK9DduLfu%yG{{$}lCxCIfw!$O&6S56n3c9FXXQ0=oF^-3TIV9}! zGn@tr&Hdx=Tu(!ga~FWr@K3-1a?U2k+@yWJ-LmQ{qAHzN0X3CmynyrV^L0z4O#*3_ zf!;5xt`Jp-R$vYKE1;5{QNor%*n<$Z-((_C9PAZAfo$`5s0fh$-U8xY(hX7z zw*8Nb2$?m}5yYHeI8L;M4U5|eNk!pa2XutxEw>}^vRi`k(K^k!W7kd3KM54$0 zRMms1)4gqwbzI4#mrF`x3b1Ri${f6kddj2MH#t+AH8ye{Av<-oYy$ED|6v*d2x<%b zFWCE@ARiPO!l*oK7?l}9LuDDgs{oGpJ}h#jUKK@=6qk{$_J>@xy&Z=aV-oE1?T28x zf&&msDELKEXiRBDLw#v?knjlqgB8lt2pnp+|FCjw9#kR8Xc+Fp!km2Wx8Mf%uB0l& zGQ;rh!?+aF$J?if6q8dgrafx8=NTgX>N}hxfd_i4Ua08p7$oo$ko5T02tOJ8FCeiI zt6{z&F*}C#@cx=-$2+iv7~G|I_fNbPM2G0cTe%h@4fDZ#V6XyoJ0S_h056YehyiQC zV)(F%F{H#0I54<`IxKzm7HVJj+4F@8qkMKhun+iO5Hpg&VOI*uD2MI(LCikijcwH6 z1L%@)*ppo0wzaX5LxjAYkoofA36AU^Jiu=T3HdPd^!fe(JTk~BAi%^65mb*ns>Rcs zx5;>9Q?e%?^=Vww6pbTl1^GJwA8&q)2nRP)TTw>*!S`c!=rHz)gRg*6$T*G5NJi9K zdxi@TX-OLJGoXP1SCIkJGy{GCxWf(j;o{;4-xC7&`5qAZ5dkcWCWNdVTrdiJ!{Xw1 zP!%UGep_ezdvqOY7i@RyY^MsgdR+X&TM-|dPIbDEEsEh|>raxoXo%*U?JX+2Z0Wf1?(M#l`F)E{J&?eGihxh>J5x?5H7i<#%J+ zV;hcY26yP)eTKI(EG~YTgNF?(C+LR7#Rb6qcMYp=e@6)2=liSBp!9z8SeopD$tOX?WI>NY_f6}=z7 zKc;2gz#-eJFoVhoDAbw|Y|O zM;=yd&}%YJ1Gt@SSlx7jS{yd4elG+u535U{z!<~o*}rnz8g&VL_rB4=dWD>_5YrPT;u;L!6-y}jjeqFzRNH>zW^+Q^P#vXTXi}l=ax}?Z&!h z*KkMKP4gNY#M+yIUhLwt9K$|?4sp(l$c(vBKWCamu6+p_7BYX8OGc!8VM7cNGWpW) zIaJ9AjB38&Lp7nh?T?LgKCfeP}I>86O6=&keffRl+3iqsmevK*TX)-slUY= zuqDPu&L?CKA#>Q#*s4|`O)k52wR&pYQasa z!F1i=>xhn9l5aDvLoE%wK7cW7SBf^Y!?r#g3<@9n-k*pJYq8!{Pyif3>>{5O+z5=( zd1kmj? z?;taQ8#PeN=wRk-A^l%(qG34_775|c5Yo4C{lpCEsvwS9ixk}*(^i*GL*F+*r^5Sp zbEdZB*vRJ*vX77zc+%kYK>qiJ^efN+S@T+OK3hodJ4$AX71Ax-Qo}Z;R{|U(q`yd7 zj|%CjyJD)Uz=1v&gNf+&I_zrl#XF>FDp{{)ZtCEw7y)CALMcNRWLo?i>OM3=q;VWPt5!_ahT#ffDEV#)%n67hO2PV;H*52US_u zdaK^k<2Y^-oTi%u??nYr+|b|w^fV8kr@Ncq!$G$3c(!L7&bugDRUJ9te-BO3oH59a z^xn2Lbj{DTBkj;pHFanp1{`umH-7}pziq#4AA$_)_M`24jQ68HRNEw9thSS4w|(Wm zv+Z;2_b21GO;PtK6dyzB3*sm}j&eWgKdNi@;ker6H$wztO#iIZKV_cBjJ_4E4bMPbr6^d|x? zh)|MQLB|i0c1Qr;Z$mH9&S})0uIRCS@D)56Zjhk+1QyZTjtpP0Q|}E*?L#;a3!!^| z&&^<}=PGiNun$w%qFM8mKM*30@p-t;z>W1|lc@%W=yI`g7qOBnjNcQs9@~i%1v=!| z%!-C5aefla?C1q`eH?n5ATM1GrJL7*8kQ%_dy;6+O?PPlMZ_3CzRO zfP1jFg50*hqL*}ckW|%MwU?Nk{v8gOEa*p%|JCkoud+~num4rR4|B!PE$X)?YTYeB z!(tV94+sW-50sr?MgIk_K%}$k zrW(>gH+28^Af7#-j4Ndo49D+9WsqKu^lK3D^k_!{aUCGG*^;Tk$aWJcJ> zVyp+J3p^H}P-exQF_aKBWeF?9feeJ%1OB5T8{`{;s-w8EPp3>u@2aDC-ixunztwY~ zw`u@0rUVaYU_?|I?R$8bJpCfq=|Jx)Oh6ROPxN-2iV@6dyg{eu33;4CNVXI$^lp2B z9u`16_xJ2u+=F@Rz+ezu2j@a*Cf)v%EXZ{`WsIiVOMokMdx$KOh=xqMwGD1S&2Vl0 zoBZva7vlKatEfOn&l=2cN3YKAZ-wACI02nC>Euu7o6t#ewx*LMz!f^#OPYKeC?=g4 zW;z?1A1ni&K|_({-K;M}FYo#mP61@_`7rg(mWgXJSSA{bUIb=frqq`#H=6XnzYDc@ z!kc$uow1{5)lH`9^nf3yleAU6m)1^Iz+Csl;coxEiUt9l39@4tgL5FT5L2Zqn zC-2`&ZjWP4k@bsAJ3|zuLzGqYt=fa;cUSav?1$g;1;>}6C@g;N=)bAoelwV|?rrz7 z4nI794-UzI>c?^Ay<%B!S=pDCfg_at@N*pGRqe%Kyz~+*%zAL*V7rQobYL)T6k6)u zpK-igW3sZ321FdPnUd_qG)bl}SpKx%y(Bob`Rf9dNj$+6| zGAs`P;S2f$gI+SG!5O+ww1H?h%(bhh>Z`*{^&pf0Q$5H{1*~;X40QV+gsB8%*D%Oe zbqKW3+Q$uZ$PtcmWzeNG|vRdTKJv zgTgQffE_^f_A4C9*80OW_l1K*z+fxmO;Y?td0T?DW*?p4ie2q*u z%r_yL*)Y#O7tb&~RVQd;h?6&A5I)=N1bUQB_ILeH525>dxuuHWe;;iKtwZoIS?WU$ zp)4-afx)Me4VKzN!}BT%SnH*2g3B-w{X*+f(IP%?-uji10th$4+8k0$czgY%z?Z3+NiJNK}>F#cn=Kz6S_2+ z{&q13?m|$+QVO?St8R~;10Td3_z32}-M#)jkvVYaW?}kPtHJb~2_h@z!d0=!9+qS! z8kuhTZ$LD&>3{XK$@EUs;$$fvR$AJ#>Q$Pi>6|EftrFQUWSLx{dp~Tp?au$v@%#0g zxDkut{2y)n`p_AZ5kKVgn94;uF!&wSV8pXek1upnPq6;c$aEwA4hAzDG2tnb5gla2 z5scrzf-oxRp2*8{!|d|xj`uB-~pi=v*G#{J_@}LKKvxAdu;zC z8}!-eTf-If1+d;Rg9PE`{4KsR-fKG|t+bZ}>09_UJoGvERfIb5<=;FLcM%Zk4=(N5 zx3UNOAu$(#-I1PE19+$cTySrE+W-RFAfF-pr^6lN!yOr`^kal;p5Ve}JoQWd-E+1A$TZ5Y&{N0-TWrhxVduLs33KMVXbci>F z);iidYHI4*+v_`;nl@{7jV<-9LA>A(??9Z-F79aFP#1Kx);sVvMZDn?!1no$JaCdV z?RcG`0Eo97;<`PBt!=;y8Rt8$V0pC+4=-HB-{Z#C*0t0|q;!N@9pYu8f(CEU)Qlsf zY`ci^RsjU>I+TghiboujoS!YI@12Nm>ub4D6R2_2Q1&(!{!y82>KAQY^$zUj0&0m=ZnJArs-`V7fVs-{TBhagS zmE>G{r729)6l!d1YOLXo^3mlyBl8t4HLcBUXbCUMb?_dyb4$ZZgN>cTWht+2X=x2P z%-YYkO589Yd~*$zuOTI#M7oE^2Tb?pu!2ltu=?l3hS!64q{>L6#Qu|VD! z;_tgHK@ZU;bS}JUBz+00; z)f;f7<(r#DndgD_hL+k+R@kzskmBfvp)knv@4lqHiX{fqG6}=i_@WZw8xzjYV$ugC z(%I=#BlYICGX zqtpC${N;jD_uCcGoXBr4GJrm`-&RN_Wg@@Gc9HJ4Gd`5x7UKP=p~ADqCC$%AhKa}> ztg9UvZVeV1q3V{9a|Gy`tRsnyMu|s?ACY^c6j9QTh)CQ@5gIJu^?Uq-Vg^WWLcI0Sj&w^iDx1nQWpgxE_2`dcN zH5Hlx)lrahUpbsT_UbA-JobC5jc(>2UFK)-rNX4S2=8??ahCQ0`axDrLh6mnE=`@F@o1lzUVIyi`c=*ti zvasQ`FuMuLSSyF)aZXcmViRwPdNa8|xShkb-ExGe88`hbF@9O_R~$}Xpo&f*Y?dwf zr7LsV&*6A%+m!0CNwzTB-AAW~I2=#4o6@In0Au05^ch`o8;3(EQ=)T)tPF#mlLbN{ z^qG|r(Px)b`EuwjE5o2Sd62<>moLMh9|^x7O=e{n^dRB)qQ6-gk^XjpFQqS{Us)N3 zeo1_~?UgT37+q7cFzZ&z?P)i3b>*|9@C{x0N&HeEQXSqS6f+IU0|hI!%he*7SSJcW zi1wqj*Jf(9P3mu0d$)*FqmSC&fT_J9SsvD2p{V$%?Nv&<=*`OVu=en>I!Yh4y%1f$?eIV(v(6CL96!r&ooToR zK8Q^=shM86J&i!q?982w%SjzJ3k-4B%4aCw#i5MELqO@b&AIBz}InUYX1+ z6DBk3IO++-QpB{OOueiU#n3dptdK?x^_&RxmPuh|ot(rf0cZn?b*hNHRy=!z zldJQwF-H_1E5kRS417l>M?@~O96JGUAO=zUbng0}=x0Cp5`Xfs563^g{G`VL$)9}Y zKgpjgzh23PeX@1?2)jr8K-fLLyt^=dhwveNQ0V_p$y)*5lqh__GsU-?WSPc7zKkPO zPh;s$Lz%{6l8if)(OA}gSDeZe3kD7Kk<|byK29HN;A0JZtbva;@UaFy*1*Ra_*erU zYvBK&22}cCdfHIXM)9znP1r0uk%v4#9cB4c&oq>C(^;2Nk{ABjQ z(D(%$|8+a_L!Kd=m2rM3t6T7-(B+aJ%8t%NtkAnIVrm&C`&ayfa62sx`EZh%yP(PtFF-)g>iDRJ`lu?iEV%MVZ zd#DFK_$Rz=eP|IYLw{elTNBz7#qVxrI2bERXO1MoXN!0fz@HXM;`D$KKkKv*=l2*d zqduu!>=*gwWkWif7PN5&TD>D^OXA%pkq^#ej-b^RDcjkMP_nh=2P@C#Y6m|m|KdB} zS*VqPZxtp1J^C{(n4AFmF3uliM=-@E41jvlZzq+;m23J3ezI;-PM>Du4<-@=;3suj zLGZ!jOo2X8o+0^jn*G$C&4&GC7om&xtyb+BId(J$?I$Ni%a@WO`X}+N#J>yrL;`>4 zKlP6{y?Y&7!u3!6N!H~-dydWoD+{`W{<*xiG^6|u$a~wPkNLQLczvRc48{{Lr==P7 zGcA~q0C~H3J@`)xrX`D$5kLh0D0;L{kFu3U_}hq|NaPQ&hh8TAZ2>v?Go>e~J|Phu z7K!60CXxej0NQf|lgNHZWz>Gsv|yq^pQL{)uTP>w2=WK`JK=MlMBiy8`E>oDK1}`R zHsS5rl63pPzZ2W7g7&ZjKVVD#nMD5B1rlzLv`%IlAN=8xwl%N^mp?I)n&#aCG-3ZT;Lp?KU&6|?_V{q>CnXF2*8Bzdh<{1Jw6rkQ(a#pc z9-{y7@8{8fl#@D%eb;t;th`=}Ul~E0#2594-{tXkhC3~hwu{Qt6`D5B9oB&^rCWr~z zqxoPxKf^(LgLVf0gTP(ewfSdx56(!ndi00l1=X|1Nq$}iCzPLrmM~vY|A)AI;Ag~t z9{5xbe(R&_1-(`e`hH`*)_iUy{2tSBmIgLL;t!wlD{rbZX36M z-X=7ErJT<({*pR*xTp54w4S)7ovD0m9z&~Wf0fC;%rlFpq#r& z3DW=%Gw6%2e}r~e9@;mre|R}l=xQdi23jA9BAbpJ^LOr+49>u?KFS~m;{J*L=WfZ& z#K^~iCE9y>9<84;`TD88jr*f{{gk^YJIn%7KG-`qh+9b@=kh||@c(RLXY_xDmrHQ_ z!k@D`)ofb9@%0V0Z=~n^Mo0LeK26389?zx!(>boj5ACVEJ&O0lFO(X_02J~Nzbz@5 zTwdxQ;=g%35IYR#u=&hr}#Gm(%_*oMp^;}*|C#WBt0{%zs8M&PaDTF_o z+Qa%PFF{SA^<^TJ!9QyYmP_kp{}F$3JCia9-=^bRGsyorKKSQt!tx9B$$~!V28)Hn zWlXM5^EP2-1wOAw{JOzH>l0bOJ00s;xqKu1K`JMU>5$5=r?ZUC(W80)6n%mlzi1`)$vrvDtWdC!wOp54({OvAyDi1h9 z9-aSd8h_iO9PqPdE1Z8uXLj}}ZG8g&%h;MciTm?F*pBwC(C;L&H*fECd(sW<6C(a_ zFYK^&QgSqYLb3rL`kw*|qd~I&QNmB*1_|R8i)=sHCr9PFp!!^VBV}to7O>FYz&m!D zuQRr$^7VfBXqd};L;EeEtmeyG2p9UeytrS|Hn zy-}&8E{*<%_VL`By#95P*P5o&@1KWqdXgwZe$H=f#J{KG$&szLQ93`wryJTQ^Q8&r zH=WwECyRRYhxe}_Kbqed;r=9Fg93U7RSS7H5dGZQ(ez0L5&C(+#{&!$fqDN(p5(bY zK9_fDzRr*HpQiE~KZL)5@Fzv=&*h(9pyTuUlytqlCv(tVo=Ue!d3%sA%3nDB%!q!t z{Hb;)ARMNq5ZUoeR+Mp%1C3JKlhi2{dhfXPh>sU=Xok4jp_L5(toHwWFJF3Y9Rk8 zQ29-5U4FatAFY3e_{s58Rp!;&8x8tV5&z=#2`au;CE9PWpxz$Q53irAvhy|hpijhO zau>coImqX~(R3fMZG%{UMfn5yYp1G~ar}4q{GA)L@cGD+kj~fhGlJvBO+Ywoy{1Ow-LzxeH&fjV!^J)7*>?f0HFT)3+s2@cJkCd5xu2#nQqWHV# zpA=BBp2K=JV~aI1ACP{rUME@}&uRVu{t%C-B3?(0pV#eacXgWLwQ+p%@thrr$Ga(B zY5xC~7M~*V_fJ%>`TuqqPsi}CaDTmr;%SQXe?gD+KqQ{>cs)K6ucy;^NR{IO{CGUm z;;q*HxXAebQws3)cq`k-d_o>f{?GZP%lHp^KE6^T@tD`!wRkVvv&wix?a}z+`#ZsJ zCdCVzj5k_;4eJxm&!Dd-?P^r~oSq7LgTA;thVhh7<2NM|Kl3M{ERU0@Ke+jh@fnGa z`EyX_Clf+GPCq#k|EN8j=Q*Gb%#RpfSg!@Ga(vx}^;1q{y+HFhACIF)lY+JLh1{UZ z@e^r1|GHg`8o%%tZM<5k9`rMU*<-Y4c*OcEdmP>8=P*8AQ&TCti!zPZ)HE93e1Ff! zLqYm-@Hv(5&why z=JCMwPw1p?zWzjfB#)*ttv9z?v#1yz(EJJeJyIC=H(FoVMvM7K^WOwHe`x+YT0DY8 z_}?OaiE{qX{0a8e{C5NKOEUSd!5_Fh$r1lG&mZTJKN;dZ_{|8Wj1~g`GV^$vVk2FP zct!Cvg`ao1jDMVm*rczkJFb263iceIPp&MJU3xSCD5 z9=VpV9-{TmE?ySv9nkXt*|r|^SPx}?6<<>$qk zexuL-a{lA-^jyym|DPFuWaAGp{)p_OBpM64zZn>n$?F^CaMh{^-2Gp4PQy`onSPha~-a z13m3o&GcUv`W@DOCEs8omKf+s=Vtn#P|Ps;Ju&EM-EF3K2*5CU?W;SLhW2TlZl-@a z4*h!huUJ$)<3Odo*%@cCZiX<5Id*94<(60PQc84~da$jJm|A9nHs0>-vV{~`2P z_HRagEcyqegWqJJk7a)^8jeNpBxNBP+mB`cbaX5h{R8qVOh)=x_J2AKeLfwUnA(qJ z{}7ZKtNnX1ol^=K`X9^w)8+U#>wmn1lSlPGmi^s8qbB64RS?nt>GD!K2iYGsw|7af zmi+-2G3nmpry+TrTvD{FOY}^dW4i=^#9c9wGkocy3UumGz`fP|4-@?VEz9+(*EZ5=?*?m zhW6qA!`kl#5{*gXAC=0YO1wbhFIq~*7J&aq>1q6#=@s;Er1UiY%=9mw13isDGyT`l z@sYMqvGCl+cK$lNLcC@U?|P&!ub=Tt&-p2Z{N);s!j2;EJZ8@f z%w>xL&Vph#J;zndW{h(evw0r;8++K3GCSq^6nrmwqoc6RtJLJRo$tP2K^D8Y-s5)U z)t~RpWEmSWSyAx(qD*#DDe^k=@5hy9yjvIF+$;uUb&X za@=KXe2%M(u^bPsCZMPsWA^QALc8ax1u4y~o9458>*YK6p;V2)`f;UnhLLLtgUNeGirUS*p1!=DbQkJ z`#d(YsKCeO`4+5W*39A@Ho;xWW|kC8WSPECDi$wO*QpGL!ZX>_t0u85_arvG=&DMl zj$zY1Kw6Yn@4AKAGE1A-1kWw3$mJ}+xF}P`hniUC71LS9RhO{E9>`Uyd}Dq}ur9P9 zWm`(V)v^e07D!R8=`*c3uTMzNwc34b;#Ajh3vzpo`~FJd`=2d}{Q@?2s_Q-g7qYPv+~2h1xISx9bKDPGEIFQUTC8Ke-4=$PC>O8^rH_E*5lhk(&m$J~YfRl`ne2VUGSde=O!YQcvPv2((|lL` zOU-coOZBujH2qVV;c{MQo#6H;(_IA>iIxUtd(=8T==!WR(^+FtwxzM!RMxGwQDWqjiS6Acti-rz>OJ)0NRhKVow{^Q)9G zo&wu6SCN;E@pxG=Z2ce`_=~<~HzR8y4)q4y6!qUk| z6WMH+?-}d#B4G`$OG&W4346{jRZ_1ag^>9-Y;iGsI6=+1s@G=AF3L?z<>lJ&Obyj+#V&94^sMvQ{ zXQppa?6+I*%<~&-(OH+vOWoGS!6ihIIH>cKQx0yPMe;vPtt?1zYf3 z%?nBPEiB!>$+{?~SFfEW8KVdPqI$TDZ0(dW)&Cw+G!{WA!kf_vDN`G5R=byF zq#sJNpUV$zo98OURgU*o1Ry6}`JT;4P`+)QU0X6MJKLRZ zo4k2lwz>dWRW@#|CFR+2Jp~srTdFM)X7=2j4V!<%mT2#>DY?Fal}ug0<^{dIHdUR* zCc7rHJa<7kTjyMm(i%*#uz5@=WwYDdjia5ufUDn@U>gHr6UMOFO`-dZNR!r=SEXaW6?_>I>PcYh9kJu#!i|{*dyS1u6B-)rG;9P6ae79P1_}G1E0FF~j%S z?1aWK9*+|PZ(d&T8K$UflBPJHW#fD^m5>XOv!!n18m8WwG~VY`eI7SPUcTFTY>dx? zz|QPD(-P*nidvL}8N~&Dmb#10Dk&(16vb;4M1wVo{l{qtD8*A)j;~d*0;!x9!%2sk zdZCiks!S+uRdT$oN~&_ZI>mGOc(86&CYQcs9qX%N^W4r0m6RO!Ui5UVZHnh%RejoO zEyc9+u&P*I9#6SCx^xP&Kdjnjl!CP#gD$UR{xeo}`*_=})=cl;EE7wfu_o4y&vs8z zCc9Tp%yyPD%LyyntB#tas1GTlY?fIb@0XMUxAzLT+Lx5E*`7NT<-I>*+u2KF4BFHZWxwW1nO8TU7-&l4rY$ z{1)|QbwbhoN@@U+KQC!)>BDN4Ywnb>rI%Z>o!?UDxaJS3>SC+?Z)#??YsLKRQT8G>2BD>iO-ujr z#PlX+Kax3miYG8?uFLsK<~+CWPDi%&T|`gC{$gg<)Z#_V{%q!isUGA9Gjq~cDE57s z!Ab6YnbvRQtS$<;d{f7!<+vM1rRI1Rj?K+)?x;^tA61q*>pjc1r?7bm&fu2DTw1)1 zP5wlJy(J;RPee@UvBd34eg)IN_fROrv8VH>GhtOso5nn zQ^$AwSAr#n`MgU?7LUqpYI}?+>r*FzWVnj(RfyXX@`EMcaaezs z27BMED2pF5awKy-K$CR04ZdX@6 zB_%mn(^gpB))s8-EUXW0Ih4Y%Ne00NHq%?sz>?oXknla4GQm5MRzeRaPVmjPt1~TE zmwU?>FDoj?Jv(LV5?9GG=Q6)F2_tF-vt7l;7cXSfyv~K}`lXAJmMmS!io7n2Hf13j z?fr+EFoVr>7tF>YW(s@}H2EIil~c!iCo5ULq;$(fW_yn%d{$9k`>^trrj0tFqPyP$>5_PHh!d%AKr zuyN?Z1gs$PdRYmGiNC~O%*K@# zvk4^$d3nx)OF`y&1bP792}LQ2ugG2SuDtRTyvk-5xzV7vU7t@HY}wx2gp@s2^)CvZBiK<~ysW6iyL8!ImRUt5-%Thib$-rTT2xvvi!I1& z^KDSFioeCwe3pDl%`C1^CKnGRX9gSRnZcBk44}7_NyV>Zroh9%PT7Q8Jk%AAX8hQaTgm`;$o8#UuTlF#`|7Q$aOg{X7(p=Vlai><#M@7O5KRn zyVI=C=2sLr=DL2LkX5qFp1fK0l(-N|R%2upKdvxuiMwC{Q%_@tD}t-}3YIY?hb2r? zGfQq&*ixnpB?L+qFZL`sl!V1mt1_|V#`Cg#wYLU8nnKLT2d%WkN}+Ew|JaU`cgq(junLV#?)8eyM8< z^ObrEs@Rm$Ym>(ouToZ(yO+8jP$rbNv+ROOHl^()%h-|?EXTWoDbrle*At71mz7+I zw=yd;+e&)XG4+2;Pn*Y*|CF9t^wZhN+ZAOVvn65~n8uW=)5jEzn`3(iPexmE+&$x% z?Piwcs$?m5B(YH`mM@P|j*n9+SZSNfRc#sL-o_Sowg#8Y**n+m@)mnbJVoA;3QKOO z@4`9B9!34RGS>GBo7?8PA0hq_!(L!&LRZ?|MOYXIs|!0?8toV$C3FMOR-EC+k#np zf|8iSQt&f2kEoar#uWUJ<+i!DV%FWNWS88WHoM-tRZ+&Xg3=O}Wtj>y)hO&4HnupF zp6%T>A#Y>o$IRNmk}k8PdD&GKEI zzp+^^XO$!Mam6-GnNvD{P@U&?-37gDXJcL4*(CQWmW^YyIqv!2x2h{g<#$1?u{xN(^L+K=MR=qY*JFndW99SM89>6`|snD zowlS>WwxszpD7b^vim?tMRA|B+cIt4&$jNHmx4-i$cBtc_s2ujp1+`a?92D(H7);kncgB*b)vO zMt-n093DcR(;g1feUeR^K#zPEeuSwK`3k(9sttMjr$LW=brR+^UxFSL;W1N1G(>u_=yqZS0isge)M+GA)jyu`iK1XFM$sEGssUP|JL zDJ?N+Cv2FDy2NzN)BATj>MJ_J;rpZH+oj9*AnKpPScpge z0_tDc5)QwLJR<*ZbooyK{^hOV@K#jo^=iAWAA0if?K9NVy`9zD;?v`tiG3IK^-=t< z()ou_{}$>CqU4*w<#T}McGSO*vHfWyy}#R{ws97WaP|V{((UkfBaRvn2!N3vyb7GJ z04JWGy@z@)#*WX(qe-wJ`pM9L2k4I*>($rwew>f`GKx7-^#}BN0xMf8SMrgS_?2Hn zj!D}1a-fdJ*kRB+pa<(u!A$qXak?%lfO80OY&_xIW5luRI8DImM~vfq`iwXcK0AT) z2;v^c!FQ))wMFAA^-)(|Gy z_dW1&Bc{i)Kk=bo@m-GkTHFl)SN}vk=aU(q&kNwQvM)ZLuY$B1_19CsUdBINKlUj7OvW7bAm*<5 z#(rHJ>DN*4c?R_qC#`oI`BX*msRf^nnDelTjkKjM=8Q9^aSMg2OeUo7V(S)Tz= z0qQF-C(B1L2#bZoO6r%C^ZHcQuY6Jzzo5*#S!H{*qN&zY{IW9Vb1HjWLyj({?N@G7 zj;QR9^sE+IO8Y*_kE`q{t!%D0m!AOs%PJet@O#zUly6wrw=9yMv7UG5_*{YfPFYXmm6BY>vUffRw^Et*4V2*EJl&ZpatwW*)+6Rq!d8Z-7o~g>O@Dq1 zAxWeGkx8~}K+cLUg+bTb}FgYT7ynw$etk=nC8;{`hVAS~q9p3e%)veu6w zS#F^R(~$ll%UU}E`hVyDE@?-z#7>8Rl*znO=IdqNCi5f?`C*w4$b3lV+H*V^(`EZI zcgnm>=9MyEFY`8;cgcLG%y-FrkIegJepu!MG9QvTn<3iGkU9Ud7(h;0E|Yns%-73& z?`wii`~So84a-y%j?@ZB((588xPr@_8%0_Bzr4JBz9YYCLq|)f!%jIB39H2h56FO0SHp7y=_j1UcfG&>j*p3QdrG8XrHl?CgIvVuQ4meg@QwMO&)5G(lUIM7Jscmma%J4m=E+9@C;rEz{M!8gTZym584dq! z!0AawEe`P_{(XRtV$i=f-_pwxV#zZ#4t!=DIPL^c{KE90%{Kz6HSY zd7h?EbNtryIq;gm*YsHjKGd$$)b6(=|Bas#_8o8X_8@CFn7`+$?)F1b?3&-a$Nd`ZH!b@OGS8oXos6+Mrr z;IwsrjCd8x4{3*nw~q{qhhqU(*zhn+eQu_9w$?=(AH!S5{duYfFBEf7vR$&|0B|W_}K_Zc60o6wxCSV z^RrG|{xlAsmnHt6y#k+~&Ek^Y0hk5;S4soU7u9S$HVg!RUplPT?tBiHNQR7lPlWJ) z_6Y&h@IM3ibOWAt`OP?d=qJ*NPsW7;(?Je}ME45_-dP~Zn*47_xV>D!w~CC}69s;+ zgx5>BSHip33ILrYQEHX&oi0(<_-~Q$*DeuoI-{ZVGYOw8$AzZ1S2$eY82`U7@u&C& z^wY9(5*(HKb!52!X#G83!q-YUof7{_4wnT(wkZz$8*$)2mi#MM3MO=3N$F(?|C_WU zo&8bDz=kcBe-?0fI{S+3SA#%jDF$I`MnWb{zJkakoML5b{0ZFto|+lJX5fU`oEgvr?a=D zzwxtuU<4)nOLBbellpux4*!F3;6riXDRJUS4u_|+(-lHYIzOkhP{PNV^naCvm&tih z(;Kbt$v=au1U*0Z#^oate)r`9u8oU#I9$Mt|LM!*vGh5f!_(QNlFx;bXpw|FFA@N4 zTr|buvlH-Gaq=5+@cSkI{)+{F&Hw)x2R{`DHl&9JInepNDUeR!a4EDQD~ki)Ao(Q8 z@qSp+YysS1pw}+H6o=3E6yP@lMs% z2AZo|sy8AJhNwM{s1rgo@QwuPTZ3k-V$PwxdQ%-q2f1USR#t8TPthPJ;JOt*5gG{A z)mLuvqkFi}6cQUr(?fG>ZAVjGps_Z9pKj+QMe%h^Eo;5u=8N<*a9v}lAt1O(5xwY8 zE!Q!4cWelp^~)HQ^=kub1C4$U%!8I(MWo}E)g3K04Hx0dA9YPDYW$@H^9O^i!ODx) z1-!M5&7QU109tIWbh*l_(22V0X0!+T%e_#9S%5SANFg<^H-336>CgE4iYDnV_m!Gk zANJsSK=_^cr-vcNFlNw+sVgJ-?a#{ks6P4Nh$b#0_J&3_|7AO-&oBYp#Jk4PHDFbnd#|04(1aKqA%})5T3qaNNMMHH%mI z1OAJb1_CH8zql&kuaLlsrE37HZEp>TA;*?pitqW<7X|_wJ39ky_@Yv43%+m^+8o&A zWYzfn$42-Ug*wIuJPR!SZMQt7~a*4X$+tD(hL~i(b@~qS)F9AihVF0oT=DQ{7aJ zfUGOch*l5Yfe#%ISVU%U@qLtW%q?`LamIaQd0l8_b$f`%$XIr9d5tV0+7R7z`66oM z)*hy;GyW#4Fsc|LCi6v&$uoEkY+uv1*+fjPIow7VJeqjYNFO5N!x7suW`zVsEa#=%LR7|gj%CCC5a-x)*B%bj>=;|T{G{cIR->*L*DXlm_it& z6R{c$&tV2)%`!?wk4r2VgAoPMScQs;@tvjEV@GrtJ^5%e4+Vsp<~CEq(H@0mLI~d{ z9q#g^?q$I`gc!53>)V2jEunf7TOMbkej!+MGGTMCafFy{8#_=KDIiVY8or#76RXTf zktHTY8XR2Hr1wM3y{xSx)E?bWc$%^2#b|#dKh#cW#qnb6FnLM$2uM5_tqnoj8;4s# zpZ45*+KV;mmX(!XcG==3D*~=U59!Gl875Ifk4FIkX>Fs$M0smVedETCV0CQQh*CF? zo=f?HbuGS=*$`}P;S-bDYsBhe1-37%_~JQMj}REh&J-CU9>M}E)~&)|)HW|zE6|2( zUG)aq*;KV$*Vs}USW}01L(^e-dB9Uh`z+#A+0ouGg8tNC^&8Sx41!c!2I(Q!-%>kV ze?Kt-hez56nTSo z)zTWOE8N)9QHZbJV7-NJffeeY4IPb5wR0P58AoWSZg0R(Rc~$qGm&H8!fWIg%A*Pa z)L}DNO$73)ttrF`Y4^ekaap*rRbI5$)iC}m?yQh6?FxgfwX|Wx&#=q=7!-nQeaFNK z>l);awxJg7Xe9_xQxmA`tf^}Y1<+3zm$xT2yMnuR4Y8`58__VtAaT$m*`_ep+PaOj zq+^8}+S^$nw&1)JV^9AfK=k|zj!bcOFOqg2Tgx31u4On5j7o(l8R@nA-CCX-g-!P? zqSIm&-^DW~TKn2_HCo;#&(kzIbNh5`g`+TC{`C@3%S%W&BpkmpVx`fKCm<3X`{EXk zPOsf(*K+N-9m1s~#SsPkzXBh6hDMWLyC1LR|0OXs`87N(Z$LdglcSZj`}SJCN79r0 z^wH#(V0IlYh_^P5CYO80j4jf-?R2f=V)uA!W#GE&n{qMtbe}11;||(9)$wr{!cHBfa*V zf|eV(neBYQAzV74fq3G5WQx9wqDlA z{CN|BW?$)%Ln9r7Y*M`|G0@>d)`9Jwfn{K=$|vuYtLn9Ieq8Hm`vSa z^w@+^e(m`UEiaSfU*lnx|Cothd(K14pVkOvZW4$63ve8<@)$xW+VcoIeodu_zll<0 z{Ah9{f+6Whb93#vkvk>5h8K^Xo-d%Wr_pQAm+X@CnqyGxprrBB{O}IIs82e*d=6#T zvxLobcM_!ty~bxUh>i4YKw=&iIleX*l{7jn&oj|$&$a9&$3W8L6s1V}T6vC%UVHxK zprm(dh%(pO)xhV2o?u%4wdYb!OZrX?F%JD=*}l0Qjs9X2J$?V{eCOXp=|U5MsmT66vqG>0 delta 15930 zcmb_jd0bRg`#<;20*t^g!!jcX42ys=zzj2kfD2%P<3jFcrLS6=OKNVEYlE65lD*P# zt2EKfY_Tv}Os&us$xJadDalY=Ip>}`T(9r@`~CAfeC|Bw`99Bip65LG zoO2iMX>3vLKBih{HV#cfwB(Bf?f#*ndzP-zB(jaBis|XP3Wvy5I4X9e7d$gsQI#ww;&#k8uv#<|?_xZQ7wN=iC|ScftLcjJe|{Wlj= zc~NUBexA?s^tFE+cImTT6C<~!nor-je!S_&am$I^07pC=Bp!b_W*fsS6`Z?sDb%0S zpvt1sH*~K|7>Lk2uUMw?YBkYKfFltOQo6!HvX%@-H#pMZ$b^Gr$q7dfIC2s4px#jF z3kNBh!1!aEc2vCibks$YqkZl_)%FYT`LCZibNhxNpY{6U^9e<_m*-8m{NixW&Ot}t zIvHSozwgrGN4)VF$BZw!D;|6-@$5Udu4vvrl%05B{r5jj%Sror-Jzw=zIgJ(ou_tv zzvupO!&8iB(@*E_y>TM&l|@N|4-|H3G%jCq`s{tZe!0~9*;lzw0>6wN^&l?b&3J*@ z#x>&YS`)5OySZN2pfz!sm^)|U_E2=F#*O!VZ^l&`H@A@>nj1VqMra#~Q?xeCHxkDc z)3z(=?~|DGGjS6nYknpi(x%2&v|-%i1k?QEYxOzxV%e?6gZ=t>UQmH;6pS_w67%E| z2|5Xs;7#2{U=PCw!ayMSV3J%YOGJWR0ZQ;2n zU+|A!VS+1pqTor9ph#xu*8w6R8TtkE@bh03(ISIhp*_+mFJa)S~_}d5BE@bj3#CG-$X?ejQp! z_T$gB>Jd_*>mC3_2Ruq)1_&QEZCVgnZ^tpLByw5+LOhzn3`&;@vKrIECTr!_A}is` zhKu}e1zwXR;-1xH*?=N7koZakzHXAJw_AbNGk%`}cTW-Zt4@fN#2qhI>F+)f_s|O< zyIN8~jjRDn4#jI_xHPYbpiYL{sHXT+FT?dRyitayu+dD#UYQ_D)}UF2$I9?_8E%o` z>sThp2xd10RBEmhmHdDZ87*Y|S!4p32ELLa!^u@rDmEF;rXdw&%W#-xzLGD)+3=w< zH{f*qlZe=YLIp!*4alWgDr00gn|4$*L534sQkg2l*%D1fvpaAPy&sSoEa(vU1|*S{ z%5azlzOq_|>wG9e>ts00V_(@U!!sz#=6|J3z%HLuyjz9``ph6yCBtE!`pP~T&X!gx ztC8WsK5c|*Ju-n^Ca9C)RvBI|!$V|vqYO95aIXwE%J61}d&v0;l?mEq4H9KIy}baQ z=;z6Be>U0)o-e~gWO%X+H_PxaHvXvCA`^tm8l=eZEE#T-;Sn-CTZXf%7?tMBaI;Sz z&c9nG=v(>f9?n^Z%4;3a*O)w+LM?y0|pL)T)p9 z6Zo|Xex-t6s^Clh4}PmcF8Am8|8O`{;b5|YAE)3)b?_eAqfZ_XffKbobk@ zgH}oEDEErqmv=+eUH0?*KzDiXU4VNB{4EIHNrE7Jp*iz9u1ih~*axt??1*^vIotHw;fpu?G2p zarcTEj+pxk??_4E9>QOxMDizIrEQwwQtyo-ZNR_aor43r6~+xBZORV2%fU(6P=89x z^ZjAPq3-&LO(&D4>mV4>-7k1kw+B681RO=$ukD>k>ZEhTKzV*j*%tsN0JsmrL*9%( zC5crb31<>_dA`j}WoBQiKmMS$Lw$f3fFp%%#AJCPp3^ss% zs3^iyt)t^_z6|clVKlmaf;&TSEC!SeD&i^FiW{ww+&B0yYh>OwpYj>Gavel3Di?|U zs1=iZ#5?*H^otFYf$exm9}B+c5x1ImKPJ+HG zx2_gk(IcFn)rN;T!^3X1(0FAzzP**^I0l{#yaRkLwzZHP^HW>!1bevYI`n^?yS$`* zlH1|zngUaGzQ)_7h34M9mBtLYe^!h~%>6=eBu0DfI?a7z3(fuYEs}aNc=Y8y-lx8q zcB3QryPG?5zvVj3z3w{A{dl6O$o<1xXzrh~LT}zCHtD$5P@9@8Dm2_jgv)&n>zirr zH@-*%^W{EE9^N4cLhMbCW}5rfR2sZDrxv#vY%>AwxXzq`| z(~Gyo=fZuB=6=raH1{t6epl{4{~b@tExPB%^I|f~4(*3%(Z}Sv0juFJq=N-Mdz)x* zg_Z8|p*9%RyTwucEg3q;e#cjHJLf&m?5uU2fKYrVyBhhFVCDgA!O zg*`3&J6G|Po?RpUB9gMgMwrBn(qbcQrYFVH4Z<1(Nf1xUcA{CX`GyJ&tuU&*F+@l& ze)F&58$FxCUcExCv89VFTzfVVTPerk@x=QFSP<2R5Iw1DYh8W880p^XE7;yUhHt-& zhxHEcvsYBAgT<%5eAF=*&y7TBI#_91P+s_%d$apcBTt}PP1LP?mTaVMV~E=UKDXP5 zUiqkIun5_O%!9#;1o^(gD@x9h3NTktIDjfXypCKvpMuS4VGUOFz1fcLw4*W$%J z@+Xf`tp~5=;P}w~8M1MiNU}Es+KKjF@zSe%N7~;7K)i3Mhx?Wg$S7PW@W>P(=kT_^x@y0EOE|v!&v<2F z91bq>=ciO)b5TD3(M9|~QL=64C6d!oUh2Cn1nwM{!wWDtA8?N0fK0w~5j5DD+jR-= zEwVVmXd@adyWh!$M&deX26aeNAz3=r($d*_5v%)m<*k=+YX8){6U1UUT%SqCj(GD$ zYrw+5q9GiKqam~tGeqO~g^myoU&I^wTjFd!8`3S;XlNvc-~{z`cy@6bjvbK7-)h7* zw~0UBh?fma;nm7O- z1;gHr7x77V7r))0=~Z38vx+mQo3lSsHy{5<-9%iZZmypPrG_8xax;x|h`N~(E4k?k zYNDG#B)G>fP&ccI%S%7{+ySA#O0^w!Obpy&<{A`-X#9zA90_1t;0WW0Lg?O?y`Lq?1;9T;XQ(A`!&!N zZJ&S!-akR*uAxx!10Fp1LEha!`;pN=`|To`-xC+QWlxFm?C`@| zMyA;YGWJn;?Q*A&((mu~ajf3wBc-1DSh<4w_;@AtVf~K!Xjv}$Xw^_32N-)0USQwp z<7EBaKBCY1eB_+P6=P}xwk@MR>Q>-~?$775&f=Z-yLEvSJmS5r!Z+_vvz;jwU4Efb z`Z)LP-9D1f;1Lh>=l^vEH;i5$aG1RID$|56r+t`nhV~(Yf=9gBJn2JJ`;(%5j#t`` zz}ts~@(-QDW5)F0yPd|v#&!ugMWSJAN&sAX@x*aa=6rbL=w2J1LrA$Rml>RX3RjJd zGiQjx6j9ij3GJdVMiiPvVck+N0E3|Y8`?5l6b6gJCz$XPD)gvD;cWx;S;T~^U~IyP zu&n#CW7Mv=0&8;bF3@Z5t0zHRrWr#ruvUZrAsnoa<>5sHq1`$Q@7KyTSu977oC1}# zy~n^^xAz6OEP-7*6IS4v|LVstti#p+8fc*J9Kc8r^v!F3cZ5w}4 zjYFHP_>oEJ)$dOFi8to=2QyH3#BupRcW2|*r;LuJF9?Ax&4eE)gEWIsK7cCtjX)gU zKGRV>eCh&@8mpPB#cM#A%pv$5KGHJ+B2m$=0%%aaAM);4g#vxNR7v9q z4rfMD4pLnH0$o&yg)yhP^TNsib%sBUEf61CoZuPA=6F>uEoCHMzl7<0k;H!gvVD*K?X&j!l zGzmMG9>kST1{M%8z4icwq)?M7EE|L;F>RLd#DzZ*Kp`Xq-iCkWSKV0loT+;5j^8;v z;pH#!Yp-tNjcPpbku2SCcIHm0@tQ}ntSRh-$xU_wGycRL6thUsMK}@o`Xgbw8l8yP zcrX`(Ew7own^~8dnJ%azR8bI$CxEb#0 z0B$H)@MjTgg7E`;lcM{x5ZuhTL_&2OYQbCgTDc6|vUhA~tG`3nO}AK=K6mmIOZx1& zlP6iK$GSI>iEoiLp+)5)$gI5NWWfve{AMK|DYkRnG@#+ z@w%94hy3+U6cYS}zrGiWiM97aDZzE5Zi;n4J=RHnT)(N2)o#12wmPz!)M?J_E9wk~ z>x!B;W>@&@KjF-J`1&R`s}M!&&#SrE9;9*>Dt4&EI0{k5SbGjsW1VXdinT#4!!@o= zrQvR?Qzm5>YZ5avCuz(JO(@)6p$RaNPsP+$a90`XncY*Dio&u|5&te2NX37C5K_=PtEVo?kIU0`$%^t*jn(oi z_}H8syq341n5;Q|VfMpZr~6Q{!VPD7i ztlZ}jZ&W4Ntg$Gg>)hjr<9GW<+fE>pD~X$LclOlHnmlDV;v4@;@GlJX zTi~b9=EC)RA~<~>GU?xpK>C9ySZ^{;kIFO|Lo?T_hB(LDU1nooZ059}PO*+2MyxvE zh#F_)TQqu8R6w_f<~?ZCo>Ozn5Z6Bl!BZ2y_rNsEytyz8=1raZFt7eT5W2pS@EQwj>|wAKNRW;MBMb?LC*1xT>RiC-96z8;&uGLA|Ql)X_~ci1QcuUp5PasJ+3#>9OG+T(o&|=k!$Gq@%N4%Up%DW zsKJHf2lRt=H%ir$e5h&#*@8KOuo~)9=LJldd!UDb)&gw+>JN7iN1hXe6rh{1AWQ}N!FoY>2I$!h5HQeN8$E(RK4L7~ zEC_8tr#vqRW_Sc#1~eP!t3by9{T%25pm{HXBcStv)&Tu|3p`E$P1p(!wFqr}Nf2y6 zzX3WFD6WKdKz{(b8R(cD&<^NxKpTO67yUBy5T0x8JE4a_R{`}x4{Cwt`{CuECTYEo z2|_cE-}*E@vuZB%8xk)Sa^)TiM?2soUs8dW5NJLe3m?bRzlij#hw8k?1%ZASS{NL% zioYjFtNM};^z%nV^<6l+E);~DK#?TpuM#npp9c6Zkng!ezSAdf1AO%&L7-psN$pkl zs+fHb3WLG$$ZB?o9e$lrfj5a`E`K6xDLAUu4m z2Zfq4L1+N#>!Iot=gR=#OUead$sIkM=j-8Vkk|h&@-~n=VX*$k_QSwv7@VDTpg=!J z^+lLK14slp`I>UeT0wvZ&yD~KSE%^eGzNmwQ!y6wTGo=Ys$Bxrjg@!}RJTC&*joHL zR1ZRR)3f;K7u_stV-Sjpg7)hmItv&ej*?;OSb$gm$Yalua6P%@Nq}iYv7;30J_E2= zm_ihL4-EPOHu*9(6$WY4bNJ&Uk=?66eE`fb@~Nu2RIc9Hr~W3W_g_!qzV{eZ3t%E9 zK{KD_abh>gXuLvI1p%V(kdG4OG*#fH5ZaRuQPrLB^x86e4(b~@1B_RH1BZy-GV^7e&!$Aey|Hv>Y7>kuQREfPAa0ewa9sX38ioJ%oRl zL_p#Tg!vx&f*=f($Hy1Y07-l!!A@_GlaH;{sd&wok=l{aWZjE6_h@9&e5k$-t4Dv3 zc6jLHOVoOhpWK2M9gWn!3)O(F_ywq*hHB|ny#Ht%{^O_~j%Zx?RUA(GDx~`HuZ%pr zqA+PDCX;|*?HVpp>4QJHiyI5O9)-qBUjOJPn*m3jh~amJ;M^1M^RT4?Zu6_nb+bA9>XS^7k-hgkR75;TAHvo% zIl)O1-1xVXM?6x5^0q4Y5LoKTowye543Q&ZT?I?HjYTV$s^9@t>cyRS0AjtxzkZ{~ zeNJU^A$ZOy6P|r4qDwG*%q-WIeCW|rBk=xH=^m*=awYPzX7gKeL~0g_q>UzWDS7u| zCNo!2R0!U&ysrR(Rl&oJ9?GSSFL&adk@Si2ojBYEvvF}J z9)<2avt*J%QNe87tT^Q&9*t^GV%?b{k1~(ssgH*z(NaZz;}|YY8DfSM_)Z0;3o%Do z-_$9h_}%@U2{^G&x6R1E(p*RzkPPCBXVN`&%$_u|53+UvYzjjrXgLeN#J@ChCH+<^ zHTXuhs11u8)$d4)^jahohPxVF7nmOY_qPR}Mn#`yfe!KIRm4Uvoup6G^TpOuY=@bC zvqIm?^y!rann?Oun2KD-Aob9I6IEdqRg`p^DKpb~-6RnO9E_*S22*c@h0_0Ls zr0>{iBlwY$QpN0xNr6E;nMMSUZ{LWjQg(MS2j8)UbR@I%3)A;k=qJM(M>-Iqz~?jE ztiZ|U6|NW>SQLW4Sc4P=o(dN!l8M}5VkGq10o4%ctxwX;1(t(p{r7wL&-x@! zDX?S=DrY~YYRXB4P$k3HDbB+|z=`+SY~eSGL7-15P8KBwF!CkEbq^q9xl?e0I=FjP zsHb?xsFhm%L~$5DidA7E-rHc}{v?~R>=AHLi-+Kp`*gYjPIjRvceyyp#w`12CRJ`j zn#pkGMy-F7R?+*JRJm8E0q}V8&W>v*dY(YI_z~aAoiu3-S8n0SX1H>T&?JT{H$`n@ zxN;-S>mm;4hmDh5mdUTk&Nus2)=aqtsuNsqNw~^QUsfEaF>yuo_c8jrg8bUIG5x3V zMe|L}%LfcsZW237T6Jz^Qsq7+Cy@^Lhe?$i$qHdklGv0xTFkek8>O2;iNb9x`M1J^mJow zPG%Zj-j=SDmSb|S*jB`?!Q71uqjbk8QU8Is?~Qbwbi=s$ExhPPmO;8wNOT511Z}0e zM2UJen7^5B@ZBV`o5g2v!OaYvw99StM|jcAES+>?xcei#A6V&ru<9fH=glmGbQ@Up ziP%->|1F#JY=CM+5O)9F%Fs!7fU6JT{kO7o(tY3NL-@~ISqAC;toD%DW-2ad@4?p| q!i(DbgqshGU={lpj|JGF`0MsmowPv>2a5kB`oqjW*no5YNctan(2WiN diff --git a/build/tensor.o b/build/tensor.o index 4e7c8e87de712d53faff88df8975603f03c1bc2f..6eaaf8d54b5bb0d69f52ef42a423b16d37d43fa1 100644 GIT binary patch literal 18496 zcmeHN4{%*onLjV7HI~u@L|yq)pKObzv@t0uQYe@2cO4qIW zvQ;fnXRBlNOMl^>wn=EupxP>hv?(j7%IK}^-y3$;^1!vC956!9E zYS(`Y*SWrO3;P~_H(fr-U6-|EKr^^Lwevfug9f>rbPMhJ+V7Ysw67p^#4U^k!ugR% z@wzao8x-1`sNWIgT^6)&Z6Z2q@lsKcdKhb6#XXKWnr8Y9srUyfTHJT0u68DdsF5@0 zx38jG_da)c@m)knMmm#jWZYNv-wMtx_dP~$tZtm}E8>!h#zj+y#z zjTW4ld}=l_{R9q7*)`kSx}xxmfRC}k#X<91SB2ELK8-=Ox*%{$G-J+jru3cT?BC5| zMuXpRyl4h#TZJjGf1PuD=$s1O^3L(d^h35C&+dnaChdZV?sEMF%U6&J9KFpWW6mt= zBWjU4hsb>HWgV$l=|>LILu727zl9Wj*4=lM{H{Y`G(tnedORT_P44horZ&M43_Q1{ z!6**ea7%+b?96|HR4vg}GB>BM)w~r#?+b6g7~9`drfe%IJ5I`uSVKfD z#{D2E@YsxcR}eDdyPPJ1iem05YElZgr!eh|xIU#0Y8J$6F<;y@;$6w4uJ1}Ui@jl5 zp{tg#oFkN}jwp{OEQ#sfkCbPsah`C_LEkxOXAjCc`JpM5+i{B73sgUOHXd^(IZP?3 zUF{sKAAXu+Fz^t~%1M}&4S|PR8~iNb4Op8iycuPph~)Z%1)ius;wnXTG+sbf8vgJ`>lG~pc3qH|ic9H6IF zl_?cWf%#%~y|%wHX2+a^!4v8KCSXacsWH8>3G|~UPjTMs6iAu&8cM2%p2U2r7aV## zi?g1ZR~^sPQ`|6k>d^!D^aqin2YL$|z%4w0i{~G}GkFnKRcG>Iko5F(>`Y!PCgfAk zi)W_QnQ5gHec3Yc%yNTPRAEo6H?rK|Ey{t@T66qAStk#*$lBu+vhVbP@b$`8u9|;I zy5*4H5ik^<@2*den&oHMD&{jK%0J4gbjxg5|EC!#x6*!Z4<+E;m4UK{*ZXdMWP#6{ zj@^9cE81sRhd7P&7@Yr{iJIyy7ZcJEKyfi98$PPm8$PN&r466ZRb4zEqV?h-~7lkMJDr8H3es)%F&VR9RJjtzrC`Q*hAXT0n<8Xua0fDkYal;?EiwB z(vIs$amTecGW~rvH=roc=G+A z*k)`F`Y?xW!*noMiZR-8Oi~)97^6Ta+eI==vw4E{Zgdt9MA7tMHGSX2=A%l-rA|+M zrF15x2NR|2py5nPR~1Ki=DmTOTgw3}1)R*>`A?@EI-P!zW(*y}u3G2J*crc%24S9z zUnPx>GMT(ypxr(_R!-vt9?z^>Di4N-7uV=sbjmCpWRLn0K98a^VSaA&gr|2K^Qc%S zB&YR~U=CjroX8KI2woDf$#s6miDDk7mjrYClHf#H0;LxkI)TpV1Ui>bouGqkI-cd2 zZ_6a(UT^xgsFzG<`!i}wGS#2i6^-}x^(UyJKOOCk_x2^b7O8DoAe!I?$*yQNmq;YD z+3s9l-!47(Wv_}ZCnl-M_TG^Ufa{W<>`m}> zc9GiJ6Ca?YYzvn_%WO22%X-nC_$QN5 zuO}JJ##6~?x{5-iPc)wH;vZ5y8t>}rWuuL?!qb_Ik|T;lrLKdvr1gCLX4xX>DDFjj z(_P8I%1+<3{@Glomx(6=6Hexjlj&7P_U8PW>`PJzc&m5zW|K5DeO~WCUvCg!;3cK_@Unsa_>BFAn>rM|S{h-jmj1Y$7yPKm9UoDV z!G_2Mi{Dyz4`q)F3Fkj9FYx;)XGz27YTC|TSo>@ZiRilZx>l+G=3i4^^DC~_^*f~g zpUdhWEY^QY>OaQHnD(_d)f)b?dwHM^EkE@8W_mP(x)i0=x*tI6*O!S58FWkiK&5{msKgfTvS=FE1Ait?wy%lFR7A)y&gfCzGRwz5(B%`ZAmXiRVjU+ zq7GCJaKz0(Q|u@SDec0j|J;275j_kW?JFJ2;m0pS8NNGsv0LiR}AEe zI4$Cg&fNo<-n7>pbQCB_t=5(#${4pc#oD{MyA~8EF$eY8T*}ZhSLwElONg+xL`0Re zxGHBcw`D9PjY{${@`bTAsF%eEFRo&=xGM3k`9Srms(q}vB23usFPn(>#WPk{YbwyP zUb?kI)-+d30d>{G8%q_$8gG?E$P; z1od=n_>+M1yf^q40Y3-uzXu%k83UYCmf?9ta9hvc32y88Kfr@})=|FYV%D>U|0cju zpOoNszn>M{?zaFu*zcDBNB_KL@mvz9r;dX>j^pc={6>-210{B$d;=>;#pr3q^?bqE z4{xVG;|G)1jr;&NrqU83f4$&F{?0&LwTH-m%#t4#`FM!@J(m36iTnd0@()?^_lx|a zA@chz`LBxnILMA3s>`A$i(q>!P%ehpg-f!F96PuPJ@pE{w~0OVex;1>bT5; zJdW4vmb}rYmKC95>+?>*jsE)rF|{;Az7cRf(Kh;D9m3NF^5}<;TJlEEzYLL21I}xE z)9(|4+w<{hi+@7wdI5MY1fIiy<9z%r;Hdw(`r|UC?0zp4-1zf6nV**m&UD!C<$z6aRlfaNGZ11|0c+Z1I@*dBv!UJE>EuM0TZyB+Y0LC*}}sORSa$GCb? za9hup1vh%$E%W_9fCuyKallc}dRBr8`wjiSQgFN97Qs!wCeLpG9_)7$;MnhWz;V9Z z1~`sa0dSlz4+4(zj|sk#?85c@j|4aRo4j!lcyN4w0yygbbHP_y{3isr`6FkMAQkpA z@}C7b^4}%6&3~`phTohUd>weu{|^C<{J*hy%zXSk$YVS#6-TlDI3L#nzKrTj{@)BZ zyTsrLz;V2Cg4_1qEx55)Zq?rc5Bl>ffTO(+1HK&ec@}Wg^Cy7gd|b&&Q?d1I5!~p> z=T=-c2+n@PaoGeo>e&rA`u|SB?S6g1O}}P7-UmF`@7Dpxejf%L=i`%r^Q9TbBY@+4 zd>wFu8z zz=Qs;<<6+EUC93-!HplxeEcxTV?5*l$N6|S;P0b8%y@kRaE$ZE0mt#05ZoTGp9yaC zG~+b|JbY>T{SDw~@7wr>3j1LN{TcZv;Hc+Xz;Qm_4mif`PQcNw`vJ$eeFSin9~a#A z!%qdb{qQRApdWq>IA0oj|4VS2e=!?F#pXXpaC=@{066leEFKeAIgrP=x*u?itM35L zx*GkT102WeDB$SNV}jfEF5(|3?Hj`kTD*J>WsV?FSt7e=daoXbAr=fCu?!0Z0CqH{3(`SXCI z{8t3G{qV5hwjaI=Jm`mgfTJJAL->Ch!v8ApApfrcNB&lJ78RqXi7Qud_9wgC#8n^Q z7+0SIyb0vL1~`t_BY>knpAp>l=gWc{JxzXi1$fZj8Nkurh4Ma??P>-+F9IC(Tmv{$ z8b91GxUJ{61uyC;=g$8IJgDbWfTNz1fTMq26Ws21As;$ZDfS!Osf_S?kVS0tXBjT*=#Ses z@eP!J=Ojyo^}}NHH@_RH(p;CxEycUiwJrM*s^TMRTmL#~A1yW2SGNA2Ae1@n_A8|Q zpw*$>o};kZ__wy{?v!~S&EwBUm8QKJ|7v}6%iVlvPi36Rsqk}S5&hZinbR5~LX_%R drhn9GtI{$@@!G-?(GA_}g7#kUQ~YnY|6el28MFWZ delta 3098 zcmZ{mZD?Cn7{~8TJKJv)NR_Sh^SzSIEM~B&pqeJpGXee+k1b{ z+y6Z0X_5=m_T+QZ5zk~|KDu6!zu`T1Pkzq-Rz-fK;aEAny{(?U-xi^VgY^^+M)IeF z;WF=`GRs}}({f1I!+dw1wk^wV*;eaDWivi6xIPQ6&rR*~>=xy7PUrc>^{J+rR^fDd zbRNI*iOP3gDYn|9e6Kp8eGb;<4~IK!`n>I3`su!R@-MZYsSF48X4q-RZMR^a?a!feM5}tjo=K0ac_U-*i5WSl{QLdwUQpqISM;G3cN|$%t=iEZvlP^I6oTyvx@sIE`Uv3q(X;q zmUNhrWg2vq);dLc6_VnC6@Z7JSf!qRuEj1!L=$iUe2m6-3FoM`8VzB7kn1Gnc8S;& z#G;1Rfn(z?xH&HuuaqU1$!Ui5 z0nK?w_&M4``@4gWF~uPr0dMZ3Y4$PTIHa?{(SLz8IhVB6XfnP2j9PkxvjVZW&u`q= zKT5>@#j*AV>Q4xJo#thpa^8SkQF(ETjOIofPY9=3TY0CP(paOElhz(>H44p;_S14g z>`y~1j^uIR?U3_D;5ZFuDV7w@qP7<6Tyk|@g;>gp_jH3AhY>0(l}HR&&U-Xl)i zqriQ1B`Nmc8D+tWXZjj&^q-?bQrI6;THZgUmR@0hNi*{KhRc6<`5l*+(R{DSjMPXu zauxS>RC88~bDRZ^-F_7~_s*Q)0&wj1CB^B-UZ3-|_9^c9I`!`nncI*7dv4dVZp;0; z!^vy}zC+rLk`CagWC%E(%X!7ABIQeeqig>0{Ak&U**qT zcNyNQhp)MaAFQRl58d3nJ)qp(J7ouJsBfUYsxPGIp8+~D5ciA)C>IRS+(5)U{w5!v z2~fje%JZv3hvgrqLHP_uJk1TXI+*hIHK-)r)H2jn)!n1sRxCO)6!#u)G{)ReSKVAt Xy-l%dvslfUef8ezgd!7^8wmajAzMv3 diff --git a/norch/__pycache__/tensor.cpython-38.pyc b/norch/__pycache__/tensor.cpython-38.pyc index 3f68c9e7d044594cf3e45f010293d29df7a2d022..ce4243d14b95522f7734ac1604ebfece37fd3ef4 100644 GIT binary patch delta 1294 zcmaJ>O=uHA6y7(xNw#T|ZhqU?Xt!#qtEE~{5N}#rD}uGQRjG#|bdzFilQ6plty>D# zt5BGG$sRojN^c&d7eNucd5MDHO$9GrJc;jZ(iE*YVZS%O-}m0kd-*uLbviy5i$xV| zUvi_Rts~Flp8#IS?{F55${|)5%-aVaDf2z{A-f-KuWi}`Xmj>q`v}?tHk(nb!R^QF zTU2`jqF;_?)7zV|SIihgP7x;12`=JFO5QUj`<*ss!$J?{baW~D^V|d2El}-bwo^~^yvH^%F?lW@0R`4=ac_lHnzItk|;;htaRpAzGvsx>PLd7=S`?caR$1+P* zVIDtc4n0^KGHISVR#pt)MK{zt6}J&8EmaGyNYTxQ2>yDFE}DxGFoHNiAHsaXStu+O zgxpE=!3p^+(Kp=am_^lbA*!MgM&adx+tA1nOB4NcZk*r}!9?)JzY?+BFb)-nLm0Y1 z2c2cXV7*L(s5HAXQyN>1&lW3ARg7@Hyt2gk0=coIqL@nJ?0SsY@X zO`)+^OP_M?Hs>N>dPRo_BcB@i?Wtr5pjkg}Ql0Q_ z`+X_|FeX2w^Wh2Nx)j{Ikkx|8Rt2Y{z*62d+}bs`FtHuvWBkh=+G)TH^@W zdI-~7_hAip1(gYy+PaOW(015_QQXmq7C+2eyt9+YZgmMcVB}$0zD(z3J))(8`nS~P zW9|8SMs-$}f8YrxiJ(9dQzXv9m@+zjY&F_?J={j*X(I2FQyD&T0lQ+7;4%RvgP10` pPB23-M?levoAPyL0RPOsX7V{xSA$<#PXx7Lg+x#Zw`5Ot{5KXS8!`X@ delta 324 zcmbPdaz}?Rl$V!_0SHnG4AP3_HuBwPWSlkm9is)K%w%<@L{_F?22I(`YnfhiF=}p} z$$N)MzzL|hhz&$=fC%Twtpd@L?~8CwuIF!=%qrwK*+f8X@>3!4$(CXQn2I%o6(%ng zP@UW?EHrt881Lp6B590_&XfJckcA3K!hiV@CFe+lix{- zTNDWc*?d4;ECVDs82K3Gn7HIQ1UdLPxHxz?q&S#?Y!DOytJGvI;sI&)oqSW$6#yha BLwx`M diff --git a/norch/csrc/cpu.cpp b/norch/csrc/cpu.cpp new file mode 100644 index 0000000..620e4d7 --- /dev/null +++ b/norch/csrc/cpu.cpp @@ -0,0 +1,63 @@ +#include "tensor.h" +#include +#include +#include +#include + +void add_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 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 scalar_mul_tensor_cpu(Tensor* tensor, float scalar, float* result_data) { + + for (int i = 0; i < tensor->size; i++) { + result_data[i] = scalar * tensor->data[i]; + } +} + +void matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data) { + + for (int i = 0; i < tensor1->shape[0]; i++) { + for (int j = 0; j < tensor2->shape[1]; j++) { + float sum = 0.0; + for (int k = 0; k < tensor1->shape[1]; k++) { + sum += tensor1->data[i * tensor1->shape[1] + k] * tensor2->data[k * tensor2->shape[1] + j]; + } + result_data[i * tensor2->shape[1] + j] = sum; + } + } +} + +void pow_tensor_cpu(Tensor* tensor, float power, float* result_data) { + + for (int i = 0; i < tensor->size; i++) { + result_data[i] = powf(tensor->data[i], power); + } +} + +void sum_tensor_cpu(Tensor* tensor, float* result_data) { + float sum = 0.0; + + for (int i = 0; i < tensor->size; i++) { + sum += tensor->data[i]; + } + + *result_data = sum; +} diff --git a/norch/csrc/cpu.h b/norch/csrc/cpu.h new file mode 100644 index 0000000..449a3f0 --- /dev/null +++ b/norch/csrc/cpu.h @@ -0,0 +1,14 @@ +#ifndef CPU_H +#define CPU_H + +#include "tensor.h" + +void add_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); +void sum_tensor_cpu(Tensor* tensor1, 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 matmul_tensor_cpu(Tensor* tensor1, Tensor* tensor2, float* result_data); +void pow_tensor_cpu(Tensor* tensor, float power, float* result_data); +void scalar_mul_tensor_cpu(Tensor* tensor, float scalar, float* result_data); + +#endif /* CPU_H */ diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu new file mode 100644 index 0000000..ad4a79b --- /dev/null +++ b/norch/csrc/cuda.cu @@ -0,0 +1,226 @@ +#include "tensor.h" +#include +#include +#include + +#define THREADS_PER_BLOCK 128 + +__host__ void cpu_to_cuda(Tensor* tensor) { + + float* data_tmp; + cudaMalloc((void **)&data_tmp, tensor->size * sizeof(float)); + cudaMemcpy(data_tmp, tensor->data, tensor->size * sizeof(float), cudaMemcpyHostToDevice); + + tensor->data = data_tmp; + + const char* device_str = "cuda"; + tensor->device = (char*)malloc(strlen(device_str) + 1); + strcpy(tensor->device, device_str); + + printf("Successfully sent tensor to: %s\n", tensor->device); +} + +__host__ void cuda_to_cpu(Tensor* tensor) { + float* data_tmp = (float*)malloc(tensor->size * sizeof(float)); + + cudaMemcpy(data_tmp, tensor->data, tensor->size * sizeof(float), cudaMemcpyDeviceToHost); + cudaFree(tensor->data); + + tensor->data = data_tmp; + + const char* device_str = "cpu"; + tensor->device = (char*)malloc(strlen(device_str) + 1); + strcpy(tensor->device, device_str); + + printf("Successfully sent tensor to: %s\n", tensor->device); +} + +__global__ void add_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 add_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data) { + + int number_of_blocks = (tensor1->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + add_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 sum_tensor_cuda_kernel(float* data, float* result_data, int size) { + extern __shared__ float sdata[]; + + unsigned int tid = threadIdx.x; + unsigned int i = blockIdx.x * blockDim.x + threadIdx.x; + + // Each thread loads one element into shared memory + if (i < size) { + sdata[tid] = data[i]; + } else { + sdata[tid] = 0.0f; + } + __syncthreads(); + + // Reduce within block using shared memory + for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { + if (tid < s) { + sdata[tid] += sdata[tid + s]; + } + __syncthreads(); + } + + // Write the result for this block to global memory + if (tid == 0) { + result_data[blockIdx.x] = sdata[0]; + } +} + +__host__ void sum_tensor_cuda(Tensor* tensor, float* result_data) { + + int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + sum_tensor_cuda_kernel<<>>(tensor->data, result_data, tensor->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + 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 scalar_mul_tensor_cuda_kernel(float* data, float scalar, float* result_data, int size) { + + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < size) { + result_data[i] = scalar * data[i]; + } +} + +__host__ void scalar_mul_tensor_cuda(Tensor* tensor, float scalar, float* result_data) { + + int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + scalar_mul_tensor_cuda_kernel<<>>(tensor->data, scalar, result_data, tensor->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + cudaDeviceSynchronize(); +} + +__global__ void matmul_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int rows1, int cols1, int cols2) { + + int row = blockIdx.y * blockDim.y + threadIdx.y; + int col = blockIdx.x * blockDim.x + threadIdx.x; + + if (row < rows1 && col < cols2) { + float sum = 0.0; + for (int k = 0; k < cols1; k++) { + sum += data1[row * cols1 + k] * data2[k * cols2 + col]; + } + result_data[row * cols2 + col] = sum; + } + +} + +__host__ void matmul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data) { + + int rows1 = tensor1->shape[0]; + int cols1 = tensor1->shape[1]; + int cols2 = tensor2->shape[1]; + + dim3 threadsPerBlock(16, 16); + dim3 numBlocks((cols2 + threadsPerBlock.x - 1) / threadsPerBlock.x, (rows1 + threadsPerBlock.y - 1) / threadsPerBlock.y); + matmul_tensor_cuda_kernel<<>>(tensor1->data, tensor2->data, result_data, rows1, cols1, cols2); + + + 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, float* result_data, int size) { + + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < size) { + result_data[i] = powf(data[i], power); + } +} + +__host__ void pow_tensor_cuda(Tensor* tensor, float power, float* result_data) { + + int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + pow_tensor_cuda_kernel<<>>(tensor->data, power, result_data, tensor->size); + + cudaError_t error = cudaGetLastError(); + if (error != cudaSuccess) { + printf("CUDA error: %s\n", cudaGetErrorString(error)); + exit(-1); + } + + cudaDeviceSynchronize(); +} + + diff --git a/norch/csrc/cuda.h b/norch/csrc/cuda.h new file mode 100644 index 0000000..b18693c --- /dev/null +++ b/norch/csrc/cuda.h @@ -0,0 +1,28 @@ +#ifndef CUDA_KERNEL_H_ +#define CUDA_KERNEL_H_ + + __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 sum_tensor_cuda_kernel(float* data, float* result_data, int size); + __host__ void sum_tensor_cuda(Tensor* tensor1, 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 scalar_mul_tensor_cuda_kernel(float* data, float scalar, float* result_data, int size); + __host__ void scalar_mul_tensor_cuda(Tensor* tensor, float scalar, float* result_data); + + __global__ void matmul_tensor_cuda_kernel(float* data1, float* data2, float* result_data, int rows1, int cols1, int cols2); + __host__ void matmul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result_data); + + __global__ void pow_tensor_cuda_kernel(float* data, float power, float* result_data, int size); + __host__ void pow_tensor_cuda(Tensor* tensor, float power, float* result_data); + +#endif /* CUDA_KERNEL_H_ */ diff --git a/norch/csrc/tensor.cpp b/norch/csrc/tensor.cpp new file mode 100644 index 0000000..3571bba --- /dev/null +++ b/norch/csrc/tensor.cpp @@ -0,0 +1,464 @@ +#include +#include +#include +#include +#include +#include "tensor.h" +#include "cuda.h" +#include "cpu.h" + +extern "C" { + + Tensor* create_tensor(float* data, int* shape, int ndim, char* device) { + + printf("Creating tensor\n"); + Tensor* tensor = (Tensor*)malloc(sizeof(Tensor)); + if (tensor == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + tensor->data = data; + tensor->shape = shape; + tensor->ndim = ndim; + + tensor->device = (char*)malloc(strlen(device) + 1); + if (device != NULL) { + strcpy(tensor->device, device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } + + tensor->size = 1; + for (int i = 0; i < ndim; i++) { + tensor->size *= shape[i]; + } + + tensor->strides = (int*)malloc(ndim * sizeof(int)); + if (tensor->strides == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + int stride = 1; + for (int i = ndim - 1; i >= 0; i--) { + tensor->strides[i] = stride; + stride *= shape[i]; + } + + printf("Tensor created successfully\n"); + printf("Tensor information:\n"); + printf("Number of dimensions: %d\n", tensor->ndim); + printf("Number size: %d\n", tensor->size); + printf("Device: %s\n", tensor->device); + + printf("Shape: ["); + for (int i = 0; i < ndim; i++) { + printf("%d", tensor->shape[i]); + if (i < ndim - 1) { + printf(", "); + } + } + printf("]\n"); + + /*printf("Data:\n["); + for (int i = 0; i < stride; i++) { + printf("%.2f", tensor->data[i]); + if (i < stride - 1) { + printf(", "); + } + } + printf("]\n\n\n");*/ + + return tensor; + } + + float get_item(Tensor* tensor, int* indices) { + int index = 0; + for (int i = 0; i < tensor->ndim; i++) { + index += indices[i] * tensor->strides[i]; + } + + float result; + if (strcmp(tensor->device, "cuda") == 0) { + cudaMemcpy(&result, tensor->data + index, sizeof(float), cudaMemcpyDeviceToHost); + } else { + result = tensor->data[index]; + } + + return result; + } + + void to_device(Tensor* tensor, char* target_device) { + printf("Sending tensor to device: %s\n", target_device); + + if ((strcmp(target_device, "cuda") == 0) && (strcmp(tensor->device, "cpu") == 0)) { + cpu_to_cuda(tensor); + } + + else if ((strcmp(target_device, "cpu") == 0) && (strcmp(tensor->device, "cuda") == 0)) { + cuda_to_cpu(tensor); + } + } + + Tensor* add_tensor(Tensor* tensor1, Tensor* tensor2) { + 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); + } + + if (strcmp(tensor1->device, tensor2->device) != 0) { + fprintf(stderr, "Tensors must be on the same device: %s and %s\n", tensor1->device, tensor2->device); + exit(1); + } + + 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) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + 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); + exit(1); + } + shape[i] = tensor1->shape[i]; + } + if (strcmp(tensor1->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, tensor1->size * sizeof(float)); + add_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); + } + add_tensor_cpu(tensor1, tensor2, result_data); + return create_tensor(result_data, shape, ndim, device); + } + } + + Tensor* sum_tensor(Tensor* tensor) { + + char* device = (char*)malloc(strlen(tensor->device) + 1); + if (device != NULL) { + strcpy(device, tensor->device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } + int ndim = 1; + int* shape = (int*)malloc(ndim * sizeof(int)); + if (shape == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + shape[0] = 1; + + if (strcmp(tensor->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, 1 * sizeof(float)); + sum_tensor_cuda(tensor, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(1 * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + sum_tensor_cpu(tensor, result_data); + return create_tensor(result_data, shape, ndim, device); + } + } + + Tensor* sub_tensor(Tensor* tensor1, Tensor* tensor2) { + 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); + } + + if (strcmp(tensor1->device, tensor2->device) != 0) { + fprintf(stderr, "Tensors must be on the same device: %s and %s\n", tensor1->device, tensor2->device); + exit(1); + } + + 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) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + 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); + exit(1); + } + shape[i] = tensor1->shape[i]; + } + + if (strcmp(tensor1->device, "cuda") == 0) { + + 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); + } + } + + Tensor* elementwise_mul_tensor(Tensor* tensor1, Tensor* tensor2) { + 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); + } + + if (strcmp(tensor1->device, tensor2->device) != 0) { + fprintf(stderr, "Tensors must be on the same device: %s and %s\n", tensor1->device, tensor2->device); + exit(1); + } + + 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) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + 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); + exit(1); + } + shape[i] = tensor1->shape[i]; + } + + if (strcmp(tensor1->device, "cuda") == 0) { + + 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); + } + } + + Tensor* scalar_mul_tensor(Tensor* tensor, float scalar) { + + char* device = (char*)malloc(strlen(tensor->device) + 1); + if (device != NULL) { + strcpy(device, tensor->device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } + int ndim = tensor->ndim; + int* shape = (int*)malloc(ndim * sizeof(int)); + if (shape == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + for (int i = 0; i < ndim; i++) { + shape[i] = tensor->shape[i]; + } + + if (strcmp(tensor->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, tensor->size * sizeof(float)); + scalar_mul_tensor_cuda(tensor, scalar, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(tensor->size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + scalar_mul_tensor_cpu(tensor, scalar, result_data); + return create_tensor(result_data, shape, ndim, device); + } + } + + Tensor* matmul_tensor(Tensor* tensor1, Tensor* tensor2) { + // Check if tensors have compatible shapes for matrix multiplication + if (tensor1->shape[1] != tensor2->shape[0]) { + fprintf(stderr, "Incompatible shapes for matrix multiplication\n"); + exit(1); + } + + if (strcmp(tensor1->device, tensor2->device) != 0) { + fprintf(stderr, "Tensors must be on the same device: %s and %s\n", tensor1->device, tensor2->device); + exit(1); + } + + 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 + tensor2->ndim - 2; + int* shape = (int*)malloc(ndim * sizeof(int)); + if (shape == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + for (int i = 0; i < tensor1->ndim - 1; i++) { + shape[i] = tensor1->shape[i]; + } + for (int i = tensor1->ndim - 1; i < ndim; i++) { + shape[i] = tensor2->shape[i - tensor1->ndim + 2]; + } + + int size = 1; + for (int i = 0; i < ndim; i++) { + size *= shape[i]; + } + + float* result_data = (float*)malloc(size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + if (strcmp(tensor1->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, size * sizeof(float)); + matmul_tensor_cuda(tensor1, tensor2, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + matmul_tensor_cpu(tensor1, tensor2, result_data); + return create_tensor(result_data, shape, ndim, device); + } + } + + Tensor* pow_tensor(Tensor* tensor, float power) { + char* device = (char*)malloc(strlen(tensor->device) + 1); + if (device != NULL) { + strcpy(device, tensor->device); + } else { + fprintf(stderr, "Memory allocation failed\n"); + exit(-1); + } + int ndim = tensor->ndim; + int* shape = (int*)malloc(ndim * sizeof(int)); + if (shape == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + for (int i = 0; i < ndim; i++) { + shape[i] = tensor->shape[i]; + } + + if (strcmp(tensor->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, tensor->size * sizeof(float)); + pow_tensor_cuda(tensor, power, result_data); + return create_tensor(result_data, shape, ndim, device); + } + else { + float* result_data = (float*)malloc(tensor->size * sizeof(float)); + if (result_data == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + pow_tensor_cpu(tensor, power, result_data); + return create_tensor(result_data, shape, ndim, device); + } + } + + void reshape_tensor(Tensor* tensor, int* new_shape, int new_ndim) { + // Calculate the total number of elements in the new shape + int new_size = 1; + for (int i = 0; i < new_ndim; i++) { + new_size *= new_shape[i]; + } + + // Check if the total number of elements matches the current tensor's size + if (new_size != tensor->size) { + fprintf(stderr, "Cannot reshape tensor. Total number of elements in new shape does not match the current size of the tensor.\n"); + exit(1); + } + + // Update the shape + tensor->shape = (int*)malloc(new_ndim * sizeof(int)); + if (tensor->shape == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + for (int i = 0; i < new_ndim; i++) { + tensor->shape[i] = new_shape[i]; + } + tensor->ndim = new_ndim; + + // Update the strides + tensor->strides = (int*)malloc(new_ndim * sizeof(int)); + if (tensor->strides == NULL) { + fprintf(stderr, "Memory allocation failed\n"); + exit(1); + } + + int stride = 1; + for (int i = new_ndim - 1; i >= 0; i--) { + tensor->strides[i] = stride; + stride *= new_shape[i]; + } + } +} \ No newline at end of file diff --git a/norch/csrc/tensor.h b/norch/csrc/tensor.h new file mode 100644 index 0000000..7a8798b --- /dev/null +++ b/norch/csrc/tensor.h @@ -0,0 +1,29 @@ +#ifndef TENSOR_H +#define TENSOR_H + +typedef struct { + float* data; + int* strides; + int* shape; + int* strides_cuda; + int* shape_cuda; + int ndim; + int size; + char* device; +} Tensor; + +extern "C" { + Tensor* create_tensor(float* data, int* shape, int ndim, char* device); + float get_item(Tensor* tensor, int* indices); + Tensor* add_tensor(Tensor* tensor1, Tensor* tensor2); + Tensor* sum_tensor(Tensor* tensor); + Tensor* sub_tensor(Tensor* tensor1, Tensor* tensor2); + Tensor* elementwise_mul_tensor(Tensor* tensor1, Tensor* tensor2); + Tensor* scalar_mul_tensor(Tensor* tensor, float scalar); + 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 to_device(Tensor* tensor, char* device); +} + +#endif /* TENSOR_H */ diff --git a/norch/tensor.py b/norch/tensor.py index 97c3ec2..022c343 100644 --- a/norch/tensor.py +++ b/norch/tensor.py @@ -163,21 +163,45 @@ class Tensor: return result_data def __mul__(self, other): - if self.shape != other.shape: - raise ValueError("Tensors must have the same shape for element-wise multiplication") - - Tensor._C.elementwise_mul_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.POINTER(CTensor)] - Tensor._C.elementwise_mul_tensor.restype = ctypes.POINTER(CTensor) + if isinstance(other, (int, float)): + result_data = Tensor() + result_data.shape = self.shape.copy() + result_data.ndim = self.ndim + result_data.device = self.device + + Tensor._C.scalar_mul_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.c_float] + Tensor._C.scalar_mul_tensor.restype = ctypes.POINTER(CTensor) - result_tensor_ptr = Tensor._C.elementwise_mul_tensor(self.tensor, other.tensor) + result_data.tensor = Tensor._C.scalar_mul_tensor(self.tensor, ctypes.c_float(other)) - result_data = Tensor() - result_data.tensor = result_tensor_ptr - result_data.shape = self.shape.copy() - result_data.ndim = self.ndim - result_data.device = self.device + return result_data + elif isinstance(other, Tensor): + if self.shape != other.shape: + raise ValueError("Tensors must have the same shape for element-wise multiplication") - return result_data + Tensor._C.elementwise_mul_tensor.argtypes = [ctypes.POINTER(CTensor), ctypes.POINTER(CTensor)] + Tensor._C.elementwise_mul_tensor.restype = ctypes.POINTER(CTensor) + + result_tensor_ptr = Tensor._C.elementwise_mul_tensor(self.tensor, other.tensor) + + result_data = Tensor() + result_data.tensor = result_tensor_ptr + result_data.shape = self.shape.copy() + result_data.ndim = self.ndim + result_data.device = self.device + + return result_data + else: + raise TypeError("Unsupported operand type(s) for *: '{}' and '{}'".format(type(self), type(other))) + + def __rmul__(self, other): + return self.__mul__(other) + + def __neg__(self): + return self.__mul__(-1) + + def __pos__(self): + return self def __matmul__(self, other): if self.ndim != 2 or other.ndim != 2: @@ -214,4 +238,18 @@ class Tensor: result_data.ndim = self.ndim result_data.device = self.device + return result_data + + def sum(self): + Tensor._C.sum_tensor.argtypes = [ctypes.POINTER(CTensor)] + Tensor._C.sum_tensor.restype = ctypes.POINTER(CTensor) + + result_tensor_ptr = Tensor._C.sum_tensor(self.tensor) + + result_data = Tensor() + result_data.tensor = result_tensor_ptr + result_data.shape = [1] + result_data.ndim = 1 + result_data.device = self.device + return result_data \ No newline at end of file diff --git a/test.py b/test.py index bd2c647..b4c088c 100644 --- a/test.py +++ b/test.py @@ -16,7 +16,7 @@ def matrix_sum(matrix1, matrix2): if __name__ == "__main__": import norch - a = norch.Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]).to("cuda") + a = norch.Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]])#.to("cuda") b = norch.Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]]) import time import random @@ -29,7 +29,7 @@ if __name__ == "__main__": #d = b-c - b = a @ a + b = a.sum() print(b) #print(a ** 2)