From 6aa18f0e36c88d65c897c26b3f258276c8d707b9 Mon Sep 17 00:00:00 2001 From: lucasdelimanogueira Date: Wed, 1 May 2024 02:12:47 -0300 Subject: [PATCH] transpose tensor and matmul autograd --- build/cpu.o | Bin 3752 -> 4024 bytes build/cuda.cu.o | Bin 51224 -> 55088 bytes build/libtensor.so | Bin 72880 -> 85592 bytes build/tensor.o | Bin 21200 -> 22560 bytes norch/__pycache__/tensor.cpython-38.pyc | Bin 8408 -> 8844 bytes .../__pycache__/functions.cpython-38.pyc | Bin 2686 -> 3072 bytes norch/autograd/functions.py | 16 +++++-- norch/csrc/cpu.cpp | 11 +++++ norch/csrc/cpu.h | 1 + norch/csrc/cuda.cu | 32 ++++++++++++- norch/csrc/cuda.h | 4 ++ norch/csrc/tensor.cpp | 44 +++++++++++++++++- norch/tensor.py | 24 +++++++++- test.py | 7 +-- 14 files changed, 126 insertions(+), 13 deletions(-) diff --git a/build/cpu.o b/build/cpu.o index 124881cefc47431222c470c46e4998834a4db702..1805e130007d4ae4c6ce5dc57c4690b0aaa1dd3e 100644 GIT binary patch delta 433 zcmZ1>yF-4029pKPMy+Wq^`H6Weua8;KK1CVec{nr`oN>x^@T^bgFrV!r|SOan~;p?iXvt>|e3I9kOMo516JVOm#b?H;zzN~jPWI$8=Q_X*;gv7|9U%t*)!uwH delta 190 zcmdlXze09`29pN&My+Wqn-f@b7#S5NpJkJu?7}0nS%dun<7NfUC`Kj*hRt)hco-)i z;J(4sz%qFjj~UYn&dHZ}?3fO)OcrDpo2@5bbteCAvMY!KcL#>rCrW}F@z5N_1uNPcs!4O|dj2?Nj( FasVBBEye%< diff --git a/build/cuda.cu.o b/build/cuda.cu.o index b85d7117e41e93931437880565bbbf3e8886ed7c..f088d3ea8de827c2949403e062217d944669e50e 100644 GIT binary patch delta 11304 zcmbVSe_T~nw%_O63*L+OaxYix0sJ;2Yi6HJ zG;G-n)IIdaI-mZl%+dHNUS4k(UDgfEBK!rr@P)3mc-x20(BL8fhqCzkTp zY!aI)-rqDA$(1LdSHA(?az%A7b1c*JO&&;y%GGc5bVo&2Yx-q3EFY1pzu<0@n%*5x6FR#!|q#t?2ey)nEyD()6d|HevL?F6{v?m+7i;Tz|+ zZmf6^eC~3s6)N2$pvd`ze!a!8R07!JW3BUguUzSVH~uCK zs^ZVVn8eZADyT>tt0#T|&50k{3z`~gmN(YdHNw;BvD(t^>wU{~$CWFZ*l`K;1;I zKZ+Z!|8b8-YWueFf=umSj!2A*i*eubWDr2^3kVkWm6xNiVH zFO2ku55S)_@T39w%5c)ZDhN0He_#;FK?G%<8A1iT18}#E@R0-X-x+w!0NidT{ZRw( z{dO@OiP!NM(r1kS zJNAvo&3!`#Fc6YP>;jLwzY7A9Cfcpa$v6!WoJ5zRpsKXc*@k&HHk=i8 zt{w?>49?z7w51SL}yvB}K5jBo+QyvIK4@&VhGIQZ*0iUJdP~4lNor3*nQ} zOg~MXqvQaE{L;)w<4&PGjr?{cPXpUV1PVQ9;5I9H8n5S-JdM-&I08usre-k^|^>AIMYeM9;8Tw-0L#fCUBNX z{RbL_)u_VfZnQs3z>$iNyEj+-!>7Fg>uU2Ou*(gx4270g65-R@m?*bgk7cvMlr!G2 z?Kmqnm_2b=gqxSgYO7JCdU=e!p1KX?Q&_h=-(Ja!Ni|HV^J$}?ce&Sogt1CkhuCKZ zCdYIooJ4E~)YN(Hl|0Vx0*WA&-p zF{r8cj{R0Dn33Ut7shhTLxWN5WY|&ft)0(Pxl+*{h3Qaj54Hjdc4n}2a`|e8$y*9u zq2>2tc{{bD-Iipkc7V1b-`>f&W;nURr)_|m6<&K1jY*Uy&>DQ&zrhZ~woy<9HU+V7 z2-9t&p`{^Jdk&Hty!QKPV58g*CsAPy)F5U{GqE+GHTtxhU`K;jyB`W0$KVW3ZuE{7 z>CE%Iq#XGU4gN*a&R8?f)NWF=f3B(SHU>;@?XcmOGRzWdloBaW(U=|+$TAtZ7g4j} z=(0guF`qa$CGtk=6Pte2bn@}v$T$kds_?)leoUxXy?1>vDjc`rXVqELT zX|8D^!Bu4uzKK&Y@l%}eMK`(_-=E?X>8={GAs2@I1?ynEV{!E6~z(U!j&} zS(ueQUQe^8Oc*yic6i~GNv?3K)t7VM2;I86*!3S{T)x5u86Nu;5&l83zHf}HB_qsr zZ$^SQ``!%u3K2dxUT;ft)f9QXjWtES?AwY)>yw8@7k*%^j942H5z+m`##MS;It_?v zk?sfdtDWJNk(y(O&3Ms|*CF#EA7r%puvshcehz9|$LPCsSiU(?AE(2iEvecU@Nuh8 zZ$hj!Rm*|e%|87_Gqc9XSS_$+bEv?gjP(6wcZe$L4Fm+`7ZpCCrEFIK0P>DF(T@s8t--k9rB=Wb$!QJLzj>knw#yjr$#FE;Me=?)zYR zcW{x3CjS1D4kA&J4E>dNFxf&A829jJo@}A{jDJI)_9#++E#rIm?aa{6L4{pRxEzMQ zFjfUjo@IO%k5;*peg)aT_J+|%tkKyCahRe2+j(xZa0|?p8vc6Ln9c)cv1}v8Kjco9 zCI80wL)@7XAB$BYPleoP_p+U(jDN{}bCU4`CcZ#?$DJSxoZ|wAxUFSC$wN&C<=mzc zuVVbq+;+4ekvu`88{fw0{3 z_o>G!tV0(}$nHD1>FF2|e0y{e&@(R_@o8qI?WTMW*<-q>lVD_i1_!j?kG zUN1bh*8$aggQNPZqE!Zuqm3zw4V~92dX&T4$g0uZuCST#BeH6AyA-wh!{koel_e21j>A(JF)EsP#ps!#R>yL?X-KB3f3XTCcE~kgzW}s?!Qv3>C=E zP@LjGw94RP)N<^1IO&a-#EiGw{Bt$@u3u#eE9UqL6s@dH^}9$Z8&{-DG8l)_DcgoDA}qH%E~p?YLO z1+r=z(4tF1Y%#1wb{3{&ECGcrg|36x=3NfxIT+jqPb*qw;Oz23ahJpSo`RM`Nmp=G z*|?&Tm{HXtt47tJu*L8=vh*G=qiR>!Qg|EvjXUIkOI^WHy{>4LLB=63EIj0Jexab{ zu;@^5R5NhhCNZPhjI0{fLke3A$B~_bx-zOG3R?;n(cjF&4sah1j;dGDDuar{!86~= zTPpc4hqZ^jeszvLsjzBXXB1o=N#_)Nt`h0z3SO$e1$`B zo_$w}-UT#W+=9nu0&exU4`5asvH}|5XgBH7Mg3t58F$2f?>8PJyZw*sb7d1>QrP z{2R@=0H3l3^Z1jR9IG4Ft&GbWjHALR9gK%|!G%HWB=egg+V=-Jn8ZYA4Q460T7&fp zehUU7D-b@+h|3s5GZfs(xU4{q;H_a6Xx(1_Zw4vWGYEcx@z8L8QgAig2|QUyf1wiA z^Cs>$;=N8Gj&lJSFRdpeR}@3)h?&8!(A06NoxmXYcE-&LD0AVELZ~(PZv|h3PGo~W zjc0)*Jkg6E@28PH_~ zzg59k;Tfbc7pMla6k;tBB(erG&8&Et@t_)D1?;h=La3dxLGT5PqaA7aVmle3tj880LyuSj}{wR9~G4#nG__vCoa>cQOKf}m@+69vo zyjs!kl>hRi8q86M!(2etfTj+~*NUN=F)v$U5^oKmb{Yo3TN(FLq-qU1mSdkmL5 zW@T8*jG%-#t>E;6QOx+8gW!K)JT$7mFiuh3j+gA8A5GkE%qemp+r#QJJyAd^K9up$ z4jsmLsDr!^JFTW!Q4)mvg^Yy&jkPy$ZgH?mzYZd!_i;ExU+J-?K z$p2EwmiUMr91d|H@#hr>>iYZAAohC^_mfTOz{=+W=ea;A{uSfYLh8=)4dbCL6qiKS z`|Xco`^ia!8o1x+V@aeX8Tf!LdpzTz4(2c}16{}t77T*dFisUP-)X4*QI;|h>Y$Zz z>EM2L0E6J$87BwoYu{6hhuVMDA7W@1J9ujl{Ev*2gA%;tbo!9-PzV2FoGL)?3sV1k z#*O7}u|jCcm_eZeWQ+=t3>yR=H3**0IN7)ZFX=d6!I#j#4<+%SLfi>qFH9W%SrV30 zM8oveQ1QZQZBqBEFTAPwi<~A~M!}-cgD!c*Umb$Cuznx}m&cwP6+8m%G~-3mso<34 zQ}Z}+HiVF;qwhl0OJ4t_5LO<9yZNYzXu9OtX{>^i`*yrY3KX35ggZ9OZY@7(gMO zE_t3lJ_MIXyoCx*F}#l#$$SMTJMx$vf0JQ|v@ef@*O@q+!5BJ1bmWoyi4a^K8^5mL z6vMZs3vph-$)7xmxAB3Qj6ohVN13>}&5_qh;|E_eUrW)gkM`A|^|chc?NQTLl3}vvwIq8QBc1TwYx#Ef8z%oC6rRZ+-m~4z louC&Flo-4+V8fa8h_+n@fJY#3>`Y!n`|r(cv_?Ow`9HmKm!|*# delta 9289 zcmZ`;4RjP$maeLFpu4GtPLojSA^(QZ2qA!wWH2-hi4g-Lu!xyi7!pRs*&`aAaR{2w z9dtZ|9arG$5*?AivK~0fOp@j&{A2|kcR>jdesbbbcFZiRI0J5&GiP*nnPu;-x-WTM zoqp$Z)vNEjfA7BgUR7S&ex1~^HEC-=YdO_^_k2CN!}Fx7z2k)izFFGSUbv7Egf8D? zXdVBp<{1NTduM6AqY-_oV|0?L{n%s5P4SRiw98kjY6nM|P~#|8ogUqp>Qps{+l0Ga z2+xkDkK3$jE+_0UW@#5kA~Y*HI{gt_e zu2syYDXJR6-V-Ze#)L<*9=1?rp>4O#M@G zUG1~A>)_OT6XB0jT~U2nlM}XIQ&;8o(bo7rG=A})HwcRGU;W-;ctu}JS`6Y*2J^p z@YazeKQ0a*F!3D1SH+OqoTOmHVD!*$;<<7778l8<$Kl6J+@U~wV`|vvm8J@(g4!~? z`@GVWF#Z^(-shE4kA|rI5Hsxa4rMqp+QNT|ne=%xtp~cYo$Z)GpV!P_g+(%P1ykko zN~vP}0;a_0&9P+h8o@xQy^kaJd2?ayy|qwURs(^uNpPy{erPF60soxU;GE-y&2ws? z9`E1b{W{*WFFK(vSPj*|2rQUe1Eq5OPwJ;geUzs7-|>yB6u0%F4-Ur5lw^z4S~~a$S0+>Zrzfr$#VU#p$7EP05I+RDuVtFZV^C zU$Hm`2XW@-CGJX`Nos#^fqQ*%;WS>~oswl*k(60#qIi2Tbt){KFxGvY7wkY1Y*}CI z>f`0m24~ingg0|7FR`?-`W1yBtiza=XAQhA2E>QHVf1OBcKxfv>${ZzFPbO7RiB;v zsx_)OH(OaM0uP|naw}_bQevUiKr>3^Suw;()$wwyfio!emjqH52BlmZvi*B;{b7r< zH4Wy~kA>wMQvHjmRELp>fQk~YZAeZ250r4xjJUqwg})V=pD&qI$IJa^Tw7#jxGC5K zx2_m>U*m05mL4wQ_h|4`@3@re95SaU4Z=92)En9@i@cm`kuMFBOOfnMvNZ3q$RW*b zNN#tSSEM__u(*ALTEkh7+!wEpwV>Pe&tcUL zqaRixmGK6!JmL&&#DcaActuR+J#xSGDBc$|I$Q~H=`Fy##s@&s(6n(ZT;7-p`@;Sy z{7q!pw60Xzu|5@LDVlOZdvEG+3RK;8g^@TrrhY5zb{~QA!w0;Pw zrkJ)+fx7xMsE&A(e}K6&K@{4ALb;d<@%#ZkcT3M&Q$g|u9ID{wHx;<4qEJLR`7kc# zaj;-hj#igR-MmFyTjDPEZZ*_;v!4pncN_j0 zSa+l!8Dl9S)F|M)D zw8!w5ViGBm7kJtQyBhlU80rSP`nRymIxeK=+44ukhQDAogZJx4nKdvDvfRe`6k5v| zU(J~m_#2Fm=Fh1Y*v_Ynf5^Ey#rSw!zG%NvImc-d3d=1*d65$&0@}v7!8I!I^N3p| z`sHBn|3pvqLk|{84qajEUfoeUt++Vay4T}~{%K#iBN{nSr$JF$El)7K+BU&YIlO(x zY8%a*TiAcSd7<@0ZvL z=tg#q6xk_>4Z>fe6hq-b18zrk&ozk)!K#CXXP%@v66PJk3?7V)rGbLLX2sHVFgBJm z5?cv3k(FcVm)Ia=9->$tL?0Q4bdQ7HT5LB2jfgfNI#EK4;lLpyW=`Ic*h099tUM=Q zOY8#h9FC1e<)WjIf-vVW#j@@&&cb2clP_^0*oo*{hjsOSsKs92;n;k1NwS6TSM(ys zFbek%YUHH~3XjCbaF@geVbu|e;gur>tURK7#9vmV8G{d!}obJ zC)s-;v(xY|!k(zyCna_cJlJW#qD~#^I`jRs(+a7blGI!{fKu;u>YgJKS`1gwvb-h= zad)9ceha~KG^YCc zj=c&u8){Yzg~wuJ0Et}!tB{prcvoVB@X9fYq3f6dzdVM!i^PTCCZacv>FOT%5eaozK!#D&0f!U)Uj)5UKT8aaOV z56)yB_|G-pc zKvS0yZs1Q2yH+#--px3LC8wib!sT?N@vD{OAC{upWZ_|j99Kw0Gb@Pv!n?pqkV@4*&XS-J$cC0A!2gbMG9(XRql6bDCj$C63IAJxQ>iZG1444VSP+OWV_b}Y zf}-{sJX(7;yV@Ae(aKek5x6M!5qGYJcS ziE%Lk8a1_lWt@^}J#o#a4XP0tN$KYZs-i_OqMcCae}{*L58-Y0%Ntb;pudTe$+XZwAG z)A|#s6U6UWfl^e3KauLM7^hm1r@)_;|)iwibWF@k9w< zBjNQD{$s?U^&=;oJ`+gm<%ke78_mFcNd9=eIy>6<@h56noEnBGG;DOE-$Ob*DN2nE z7l%VXk#IN4)tl0avRT5ZD~{>**>FBa(v)`SKWl_9*l2M&H(N&loYV-qJ` zJsiWxb1C5go31z;^im~Y6S0ZIFu#OT41J~)o)}3u>4;OrY8x&N&uT3kc3=z#2s%7Si!?3heBi>~3qZnCYX;!4CAMyCNqoC>7(){(p6f%V+=q diff --git a/build/libtensor.so b/build/libtensor.so index 66d9ac138b805f36d4b19a7117f35604a14454d1..7df5103d80f813d0f8de68a4c3d17c0258bf506c 100755 GIT binary patch delta 21848 zcmbV!33wD$)^=5OcSu4Yosf`(B?%B9EP-YxK%`lM&7Oc^#{^^xBM<}>7_ocAg&4uW z#R)q44T?HA;ur-#q9O!P_9bizh)zJoW)UKw;9CEC?!8s1RPgzq|G&^Vb>4H%J?Eag z)Y4tI^aJPVmt4j1(Se={FJ3{E)0wN=cK4Sh3)Y~DTktA~iw-@4x6_?j0yn7D&Pk$T zX=kHZPY7Jn&WhDA1hRsr*`~n_uUnOM`w2{ylb<^6QO3SA!SWRjqYesOI*n-VN)U zrzriDKx`n@ttbPYR8}ZjvfJ11wL0qS3pRSa{VSD%&{5^g>xlQah4*$M0;RpAoj|+b z(@pX{K&gu|@VN({zWDURr$0Va2I7;6&tQCp;xhuDY<%v+CkG!YWAGWvj}&Db=mdO1 zUr^cSj#%d(gXpJMS&k8_UdI>&s@uFlW=!x($ine2S=)i>+2 zmz*gqOLMYvr$_xOJMW5Pp)Q|#nPt1;)WgK9K8~&fdJ0FmYA4Q@g~YRh5RV$*n!nOt zR~L+ zq(P5CH4XO3uK9Y7z>N&Gk@|1s3EarUB8l&rB5*SvzZP`Z%eiB4^ zN;)o;DKJv!lYu@;O++!Kt5de@stnX8#7Z4GRx;)6FsAD<$jPt2u^cRKOM|ziLH>gR zFjBhqPg*cpI_`SE(0^IhvtM>u8Lg!#Z=yj|K9ZUHYAR>_iZNJTlLo1Bn$6&HMRCZK z#?KXiF~56Be3Q(C5$J5`AgQ0w-wB?|VVTj#as)31JMuq5w)m013W1T*8*;w>OX5pp zMPERTQg=nR_=+%7ev<+1kSn9%c#Ev3!BU}UG(gGCojGUosfZTN#Dxz{nY*B1-h%15 z3#ZRrFz+EHcS7e*Cdl56jA428UE!TQ2DbUUyPj5eLg$_azyeDWHeZPzud$HWn6Mgf zNPgCg(Yg7X`^BzQS)gG!JJK-3|9+UbGDt>DS*@1TRg^l)SzuID{CmkFKk?3AFzP~- z3l{lX7Wvbw|Be*8lHtQJl&hi7jfBk0VXd z!nwYAt%5}I;*ryOvAjAl@U3z-2W)s1o%gA97F`%BmmL#`Hw8j(-1CAa+wd&Odu+IS zozP#=Q8++4Fx{(5u?0^dSO%)Bvk3ZE>nrOmxNIz!ZnWU&6SI_BaN`1ngk=_7PDF0w zFSiH=o7E`FJ`0{^!7D5{I@2ttEx265xUABGH#W5uB`9z|hLfTsSp-!UgRU0*js;J) z;Gwb$DB*DyJl-R^0Fx2JxVj`;@P20gAy2gkumqXKW5MZK&?r4ExLnh?$ZNr|CYr@( z!3P*9^?#N{K-nzPKpQg10u5L1$U;b{4$Af{(J`i-Gg;M=7GKTcZRl2AKv3 zy265!jZum%xN+4@vUL{RINuPyz8d!{#+)U=#%h6riEoxt3oe%eE-SO(bSZC?atl7v z073UzaB-<4=U;_IfF;l@r!Dw66IGN-3og%_TpG0C@``}tRTjLJrJg%}iy*=xP&!qQ zUMz`b3ANyTO%zMG1+Q;IZIQ7QCkgkG9}m3*Nwj z`z(0x>h%w^z#6WEVxhLVi~#9Y4hd8hU(_>=NH#w?xcz}$#s*@ z=!)WdvV5VYSNN8H61tn~>)kE@zM;4LTHo9*`Avs^7Og8cF7^9O*NbmPCjK%DV+wpH zlc#ghjZ4i89;VBFRg}mi3^CFNzsSbVvGLPw{A3I7rwVc`0P(|Ze5Q@h_`mRWeXB#8 zKR<85|Et4pwgA#>d`lbO)W*k}ykDfSo(TxvW#fPUrDn>1wDI5D_-}3emp1;C!TV(& z9yS2Ue`@1Dw(;-V__uBR>o)#n8~?oI{l*|#X;XO8#xJw+i){QH8$aE~Pqy(ng2(b9 zTRhyRkZI#HYSz2dv+vfYf!DVHzB^F;a;T9by&~Z`bog!~);U0iD0hYm|672i4 zA!%=uw3HHGk=NtnvY=@dMzWO|Uz@<4JG#C%^A=TEl*zK&HSf0BlLrp4c8*?kk(ybdwinJP+Uyq1l8Oh(Igpu_aw1Kgi$6BlZyve^C(Pe(^o@!E?P zRrJF^SlMl1Hk;aUqBcK^?daGxe!(pYG3T_eXm~l&z0X&8g!;M`DDrC#HqWaEw}xalV&l=pbOu3x|BHKdGSOFOkn zTm>DfeG7z7mM_Zit;{^Vo$h=H{3qMq>0WKDFbE81iJgaO{cpl$dlpDbjCkX!k)%l@ z*sjh)BSxCQ%Hb@bOG|Yw>)RzwQ*Q7iJ>4Zn+j)gYb{V&^!A!UmNg8^CTI^cOh1$nM z+4(L{N8h?!GfB6Gv87!nYU78p&%3rwn0WoJNqX%%o8L2*b?nwgb6)1{OdG=HcgxbQ zTsOlWdC3U-<)JL0`_PCxK?CSMgcWpesm8ItcTdyayvD=+s(Xw!JIKRkxbO*Ht1`mA zb&ZFe!G+okAM4rU>FAfgs~L9Y5O%D`L{0Ot4n5ndPBy;h813zgsIH+pi0$i{qE%jD z-p;Yikshn9y~?|@Lwd4$KO32zrd3_xA(W)YXzzZ@LpY0jx?ndr7MbgPm51;I7iw>3 zvLDiO!@q(Kb)WMpOY=6@Dlf5YZ@hN#GEdFhz1nC`eZy1p&Omm(S8MIFOFV$I-m&Uz zHlX*t3CF({`6%4^A>DRtr@I+Utb>EthTa1r9+px}je+b&@0RNGEWS^g+L{gO(@Q&d zfk)}WJ#sMbiV@|EOFYVVxKKOYfAgL`5vumgc^<{)e(Xv{6RmZB7Isf7HH~@h8Ks@A z$Ot^B{dn{Ez_)n2RQ z3887ffc45u(_EE&oIIWxqwP7von23dvmtOH#z|HsclH7oYP-_exy)ScW;)9l)He3V zFG)u04XrhwQEI1t#mCC(L2a~GPxBxS_GBjp4U1j*m9FpKO_wb>I7RuEU(NCrE@P{Q zH<(o>R`|WaNmQNeq6(apf+gYhm z5KI^cVH{KzO}av7vUfh`tvlTVt?Pm|iPoi)W|8Z0F7&=d!Vps!0pZ@vL+C!7R1Wpd ztXS4d#Nl-pMl~~Qo>-o&ce*! z_GaSzMcJ(Coq1dnX~g`dvfUJj)z>xw5=kGR5R6wDo8p=l8HPxJ)AL!d*g{k4JP0E>Mb)Ns+~j&!|4?E91K5i!R1B z#1ALg>R}DEVd#;G4cG_6qP5pQV~2;e>+2In2!aoh4@{P-rI==+4#Xszus68(WOYov zKVv+0YD5I*HX~ zjP}WK?&7WE+{Nxwtkc+}j2p+GP!F~?7vU5&cM*%{Ga}j$gd^tXAs7VXPH`9M&K+DbQTN#PM%_N>in{Al!^cK=@#ud#%AY{R5gTLE{I*ub0TNFp( zM?`V9wGqWZ=!z&-Q+m!FPlfR&DoOC7Waza6ryiYqhv0109=(s|H>V{%uvk6j= zs;9aJAPBUH&K^$V?9q207QlG6etaXX-F~)h{9nTE;r6cEyP0Q#SNm}vn=``~{#LoD z^oE=5n33jLDA|wjgu14Wf1}BEA8+h|k2G!O9+oxnZ($GapsHLsyLcOh?&WPbvt88h zi{x#1RkBg|hO1`%%HG=b|MrR5hWJmokMBR`KK@hAeZ2H3`3Rf9!M(v3*qlinJS`+Q z6;JnTy2$>twu|ArO&5i`;UbppzJE@bYa4eHvV*(1x|=uUp^pgM8@yDPH>I~^-^14$ zHT}G@ySAT~cfwB$yC<)O_DMN=Aa7sTs1Lc11|PGoQ@q*_JK4iieBnDj5I(Mjv3*n0 zJS!x-72ny^^zp{d+CE<0!CFqeSNnN~*^W7wH0Y8sAMtjyFX!zzyH)sT6Uy7ML9#dS z#JHx9^BCgxF8}glvmM`T=fmos|I2-pf53;;s_lGOMQ{*9H-rzXQ|}9x_l4Mf4A@@V zN8fGmk)|bo%TX@CNO zA*3+p*}~x~Si!{p>a*;XiH-4O$8}&cRqtZ|%k!4~5zku>4(<((aPYh-l8tlNeTIBg z+o$#cd^V0N93Cvf|D3B>b?D*;*GhYXm1S(`%qH#$)KCeHJ+JGmbY|oDjtv2>j7=p{ z)NSy4Ah784!J7PL@rgnfCu9}x3fY}9lGPKk5FwiP&5*LExa^9M zeIsPUrEEBtofWbZLUyQB_>1AP{X(`&$PP-`7h5Toav}Rb$fio!t6cVukZlyQY$?kp zng8!X_MDL2c}Mv3Y(;jLf0A5W<5mY3a2fX_J3;k-Aue>iNSRvcdiZzZ=$bBe1tC1W zpW};OK%xeR%Rq##X`OoSI($c+&-~0`2#h!g(c1Q#s(nZ4}?)s4wBO8r%UzU z-&`=e*oCiF%ER6(IOwSPlMpRmDJc$DQ#W58-9^%vVzEjb*Pqo>0?qXugtLFy)~sb9jo`+$CYG(tmmLvA{c_Y*}HL+L+~r#QGgd zl*p^)Fa5&Ep)|1|;;_3$cyuj%HOHIQ!iQkeJCp{s@PBfge`SYn5G-XNRwjoMV?o5> ziyTj?g^zWSK1RN+Y~gsa4aMVWtgH?Nv)wGx{~0c5X2Qn7dsyxrHGbD2B<@hcn;U8- z{ea`~CT|>s;nd_%S{STJ|Ly031QRe0e&KkE!J2eKI91fhcKz)Hfi!O+ECkuxD z%6i^HrlZ}>4BTr9t|8UYZc9p~!6l?L+FhNx00N|P8UaMRaewse6t#f8vMOHf&z^a9 z9xHtIUgljH#m+yQ%Su3_#QQ!4aV=aJZuU6MzS~YHh-w+nInvT;~rzVDnVzpG1JT+c%B*ZJuII5gR z2)oq3s4bE~ZWnh`8wO(?AAtR2ACGbft_E|d*tW7)PfZg@Mq#*e3V!&xqi zSUZEg$0n%J&>gA9u#eZy^wXr_Hy5DfG7g_q*+sfEl`|1?n`deKT;XKgq%?cIF!qg?DGvrHov~%W|DeD_)G0T3S8jY|LXX>`7{L@lUkYC9>mtTBUs`##F6J?iiKG_EJR6 zSNCg>HVcazN-cwNa(&F=Vw#_d(t-@4vFM@HGT0~$^0-(Sq`<)SFBpW;OBP|F$wIcC zCG71HwNY}5MWz*nP5Cs=^&~izsgE>>N~U>$ve1;j0{I|hT*%k9u*=Uwe%q9X(!4== z)U3MkcXxUdA0p3kdHxvxmgZBnB_?4#|uMsJc>*PE zES8aM)`8TvN*^IK&Y~p5tn&l}Z2rCMKvfO#`pxQt4H{x84Yw=2=OhI!FqEx1*h*XD zB5rnIn5@7OC}vHB-MBv`8LA=1OIy+`FkBo2Ff6E0beIaY8hSZ|8X9Aac8j+p_WHs0 zYIe!}F=|6L>yX!>wG);Zw?2k3F8?$Kbz1=(uSWBE1$Q%OaTkLOj5u!4=`J2p?qUE3 zIK3!J`6Nu`hAi)qWOXB(f3pF-i`9If2ND#68q2fj5eTwj)QVL;63W&e2@PKt#kG~c zN3r*gw9<+OZa#bDnA@M!^N?2WpXUA0WHqGzORmVSio5>HuE;)0V*URi5|`AukCGZ$ z)ktm9^?IXv3674Px;)vtapXWXK5}_`H8N9a8ky@*lDZ7U!5@v3q^>v!cQ#Jyw!z^@ z>b}9DCH2_g2#f9RXdU09tCG}HP1Tap)l`SSV<(T6>P+d}K06^cy*I|z z=OcY0|Is=!DSbj(beOJ$oav*z**bD*nk(|JX-(rg{WUFYz7mqtL@Q2>oYE^UZo!mZ zNu8$lYN_?D8FlSh1?%;tKt&JVM>D|9f6O1FopgO8C)#S=p*;& z`U=n^{a_DzX8`O$U%40dpa%!T9)s{-LtziPeYmdseKmC1*1>_=vB~t zptDA!0?;F%p%@Rp+=mK47w72uFwjcS2SMA9(RI3>UI@AgG-|A_SAezybz>|w9H;9Z z(2Jl2pdk}i#dzxNE>uj?tGM{;$2AZS${0;NGb6^=lY zAJFwu(4qLrP6g<6&??X(P`dSb6VwCRVur4d104st7M32zH>OK?^{) zE`S~Acc5jUd5d5N%39(;!86C(kLY?b=*Y#o?gf1fvFzjPytt1)F${)Q2Gk32 z?AIddgz}r{Es!t5nEMG-G346!BW`-Om%m&;7{s?9pe^<(^zrF|n{*@R^G)&t^p=$=9Eik)aA>`$MBImzhUJHdLBmTUCEs!sU zd`1lilT8OFA>R%;f2q?9um#URGHh-^egMnKV^;gjybNl*lYMb9q3*&7k{{?~`o-8*k7Ajb zg$1}BbgP4IZn}F3@&%JbqHs6|huyrlQ3C?Q_P=vp~3YW+8B2v0GiGey%BBX)ZhpX&&Y`UPm-uKQvx9M3Z=V7WPyOp5Cpr z^7JmHm8W+ztvtOdNhLxM^QB*Wg-9hv*%IIfbUVWf_I6-(uQhexJxJE>T3YiwsflsM zOF1%61t)A?>W8zX*OJ`wog_*ll`{6`wWMZ>e1?gkz{_cR;Dmyt@{-CfT}#qZ<5~Fi zL~U*yYkU1OTup~MXvc_tTCCqAud2P0W=WbSX@R5xNsA?2FKMZy<&suNS}AFjq>8+C zbW55ns9*6&z$9mP%SKX@#Vfl2%Eo$h%Ruq{*D(BNWd#;mZW)-7DyV9v{=&h zl9ozZE@_3Nm6BFTYP?>{#|8epJ7DUdL>f)T>$Jw}vm~I&;9kcX-Pog+v-3CN5{yA= zSQB+9A*a*fda^T72mVgxz#}2v)F1{NcWoC_}rhobxs8nvkVl|_$dyGw4 zYT_|?t=d?5swMR8P0t2*!yB5fTUTqtKYbfxyjN`z0FPIUU1DnD35xMLbWNO|tK7{o z`b)mkTT{P@;_gfrA-3YcLriaH*6YVXetX90s}K!6QQEVnO*4rb3$!8afdiFSC)^m$ z264Q0yXVMyj<3xi}zSdA8W-UV@T#?0#Ww#I{T7 z&$8(UrG8zy2t+aTLonmj>QUCfvuv(Fd)Y*Wksz>8KsZJhe{kzym&E^ zV!U)OVKXVIvVM^i%$(^wqNd7Xo8vd6gUEoWf}j6!cv0#HZ2B#5xkC+DVZ*0Myx4}n zOSm5f>uiGC(qO#}?~C&?IoN2!ACq{g4c{s8G8>MsRCqn*z-bJdC+i_nu+L_&6gYLk z=w%`k{89>suS$HYpSAtDRo$K3JVxo-KxGquZs|V?9=8{@7M=@Sul9VtF7Z&C`?|O= zA@}P8R3E-x6hxuN(%M~C&UlW8PpGw)Cqkg!WkG7q@#^^^>sPvP93$7ZAhlxW^|*w= zkkUpU)w{(AyE(M2(%?c!n?pJ+Akc1tDC?w^eM6<~61Q)tbePw6T}tg+?u6jFi(;^E zx70!6_6=*Y*y>;7y8Ts3?c3J;kHqa8_>=>0PTy{*=g6K_JmcgW1IbNV%spE0B`J@_ zzQYh*^i!1fy_|f8e!fVKv5e1=c{jlPFQYdzYs@@oCy~N+h6JaJ}<)gAw;`Y6G&eP^v>aX-S zAq5v4p>|eh} zO$o)*If~G}&CR_Mx9`O?P2%?LZyuGn2h)JDrPRKU zQ#dZaC^q{hMf4>p;r6Y%`m<&_zTJW}ibiv%*;3EGb=OLX+qddDO7`m4{BK=a1JsuE z*WU{_P~^|Pk<~bf+qZ#QEOGmuWbaGdzPVDB#O<4>#o$T?exsEddxjyg!=In8#;LjV zw^M;oQm~sAFGFPFOC)aJpX?K|N(e%Fx8>_Lm&W6ZbZW4D53?R6Q&bo~AjS60o#u0~ zx{i}ZA4)Czj%~*y*}R*vjg)IrZr`&ix+Gge>)LXXFcw$H6pDQ>ujyRed9@VVw|ja? z;`WWTUU!r8F3J`uw{N2Mg~aW zROhtxX*0II>`+fQO4`>^8`S*inB1yq>x?>Th+41KCQw?|jIDnORm&W1W7{WV!>E$x z5o(@ge<*g|9aj>FP;*MAhS;4%iv6JQE42xox-M8Whc&C;VC(nL#RgO6_g1u@l$yV_ zZ?xJ$D=COk*Qh1{`hNy%`jb6kgn&SW+6Rx})aOZ&pwYT2^ul zCgH{nU21$uctf>&lyM_x(A1LbhHz-!$I1J-vXa#eRZo~l)Tww%4m4DIxFu`c#8s9= z#gVOXizaDlNjAa8otr@`N>&3lckncBCp!B|j>oAT!nXv3uPuR+@ObztmaZioRgxVK zYvU%@qz{#>iB~(+HSX3xOAf}X-J*=!G=p|6iAsRIahGOLZ&=Bw1hqqyahGP$_>$F7 zuZvp@9tQT>EIF8{83x z*44snFq=@+%)AL@silUqmL)cx%z@M#D%;;@4g0X!)bFp~;&aZs-sgSS`@U=3Ywxoc zr`P#Zy=Se62(iUUETwTGIYOUG=YRrl6OY!KgUa& zX^_J4AQG>PiuyZ0r9!KeJf{o}eQD{p$?1VXGfw#KDvWt9Yery1txPT|-I655mL{id ztn3mkrAsAYC2^>K#!9JFQY4!*{bfJ-l?9umv>^pjf&UnZq*)@|@$jC&BajjWb%3G? zq`L&44B8i;yYU%_Pbxk%4dw){BxwjfS@<~d$;D>`J}!Jl<3rOJeC}a6Ng4~9hmZer zGC5@lBMp`~>xltdekFFRgN(KIAx~KCq{?=DA>VBc(?1-d}ag?hFeM+=cc-w63jR(uYT{leuPEZ>pz^6M!d zWcji5+vg>{G^zZ_O5XGv>aN3(3z1q$Pt4#IpzTA`Obf@`(*{5~D|1>P^m`O=egG%Z zDRd&W6Dk^Fq|m43D83<^D=roKXI#J-3AIIaSb6~n~a83le zZXCxW1nv}gqL?;2s@qnb36e&^)Gc2hJG1j5FpzE1!+r6NaFZ-=+RUq@IpF{k#xLR z*tt1`Ux@kbrVgH%$QAXLh!+JG(}9kp*SJ{V&yVD|-qfX{N#mz- zT(r1Eg#POSn9aq~as)>GULDQq-`I3FupCi=bu!mn3mm~gx@j{|xOHNXzlkD0`AeU@3>LNHlGD#o7lj;Pt?X}o|w&5nSjhTiDT@wYgU zdWe<_nZN}9(G)~bRm|2Z!Z6ofMq8o9kvQgl@C{yt!1K!V|UuD42 zC*HKvfa@1)Dy%W!PNwbVW%e5cSsY^Dbp~9_LWZ9<;OIJG^Rq20dMcsmZS!^ zLD1PCxMjfY2E0LZ5$&LG10I1vN)sl8$T3#38*s6ZGC`cc-L!z1RZQS86cEb~Z@|S>mPu^}TwgV*Ai{tL8|>H(c!&XyliiJ@w2eXFFcc75 zI3`Uo;OPeaGy@)Qz?}wMY>b%HrQ@{zF@^#YbOA7j0WUD%p$2@m0S`0ag$BHb0bgjq z+X5Huk6B<4^fVMGHQ*5jyuyG-8t}CSyqy8BG~jWKxLev}5M(zBu=8!~nhpkhrvdM1 zz-tV+*jzK|egodgpkK#uTutbp6ZaELeA-YT(@>z^fJaqtvA2+)cC@gcKZmvM(V8Mf zt8=c{>;Dlepe26;_^KB3?YZhi$Lkh1vfHtVY2A6KBwcIW={N9infO;s{CX4rjEP@u z;vY5f%M84mIxIE-f5rd*bZD~gW;K4$6u=A|Yg`g|lP^tUkmq+udc9?Q3i@ec)fIOc3(7%uT4QW51ozj|199d-KfQOWoCz#$V zndHoXIqtDoX+1H2z}V{4=`&t>sF3>AG2HA?t z>+;WU0=_YnL=1dVnLU(j8u*|!945)=Wo=0J)Ry62Wat4ucasI&;9Q+tlR;dmk;?uI za(`-9Wn%`Z05>CpypcLBsP9p(a|;B;Tcz zqk|_Y=hH~fw1fz9oraiq+F3rL2Ccr|S$3Fq`MqR*+F)hP=d7o;4<-lFTyh)IF5Mwt zCAsNI@=CHOJyEGV%%Z5pRoC<84Lypj%%sl(CREl8A{WwQJN*eAYIqxjkJc2YI#*|& z{(#<2P;d|FF(h00{1EfjkV>8$GF%=*z8sPp)aj7k5(fv7-Wj<;`wr^Bom5hm5hFiI z-pNQ(>{nUS+{_47`X6LL{R10lPbR$arpdiZ>**83gi6vtGJ5E$kOw|v#=K24J(XM- zI!QTjH_6USh+J~zFPmoX6>{DgMqbQ}SC)RtB3?g$oX>PAzh5C|hvWq9JD^9{av&+o z$_*;?031MkS{mrAz#1DLf*cgL=9^zr;ttA zt3tx|H;s7v0Fv&Uq&$>DUT`MJ(d2-0oD#E-MU~Z$qzsEzq>H3#NElf#EKGU%BAGm@ zkG!1h8J48jE;9FiIiX7YC(M0LU${?(Q{G~CUu5q8_?TJi*q2Pn$qyU}9ol2oi)6j4 ziy~bhwXO)o_h;559dqNAKld{A&dDS{x0`a~C+2*8ZkW87ROe<#=IrH-Q&zi+Zcaa- zJ3Y*%yZezg!!v>|V-qDNLNb{;JVxG5N{1)O*<{D?0m|5qSd_(gArR zgo^X7>b@g_WJURqMbWi4nKZJ4qV^_>M#jn`Nae^e%9vUvTF{HM80C_Wkb6fZw|e7y z7+Fyx)0$~Nkji`Paw@4A)k7Iq!+ic58*yN zi@TF&$0o?#N$uD%O0V~s#W9Jbd!AjXKgWWbmKUaMt0(Q#d&-1t&P!5a>e=|Xnir}J z+Ql5@<0jmb1M_@*xaygsU?x<0Cy?Re@|9@`+>my&YF(QwT zSK7VD0=PS#WQ-rKtpB>YYGR>9S#p+inf#vOh$CN3o=VQ%=dV1mgZyw`n)1kL5|`i2 zapnxO>4DAD&N3UWhNHTn8Ot&ofHY7St4>2hbmNz2h&$hYS0pQh0^SRQ9}4ly$ycW| ztpa6lO?Td9jiI+DMLwR0GvvraZOHfeA|#uN<%5ngDPv6WU4{M-Uz$G9*^+$NnBPo7EDxjpCo@4 zgex0Ql8PCz%74$0!Be`XesTg!z6@t&nP*olJ1Z5%m1mxX|F4-(dg+`Pqgx7uCEjh7 zDB$_|G}$nxgI`P3#S?dmyg7TEQh9>8SapKAaGxfV=XOrLd>jft!H>zsC$x&}I-4G? zySNHn?xGaVJjYHk7r#<>HeCP5w;`iub|6>gBq}S9GgtQ?XRfB5VzpU*42lO}!{q8! z7~`&@y6UcuLYKRmP1~UQBy)9wx@|aabR|zGcM2ku++);}oIwUniBA3CDAZnu1(T!4 zV2V4E@yL>Qz*gvTM^4%jM#vCQha4H!`S$z_Yvifv6L3YlFPVIgKN~=)b zkqRAd*%omYikZcK|9>c#aS5P zF2>`opZAMD)VPZmXkb6pF&8apgXwyDxOuqnq(maPmj=7AFszX-v5#Ks=}_-E7Q5 zo-PP`V9Kbq;-bx2w(J&WJzFmdV3TZ~w>+TjYgFH7-|M7TVVY8Wm^dDE2CjIO8+p}6 z_AE_u3>WM?JpICB2iwm<)H(H%fT*ZR<7(FTAXc;9O2{G7YVrL6>tA8^`)(qa7p5s^ z4-&^BXJ8-$_jy{jChg}XIgY)|-JQf^x~BG@LM_bp*MAE85u{>Xm~!8zqP37w~y$JC*01}VTqKKCx;3bd$Ai^9u*PAh3&YoT7?e6 zDd$2P7Y1e zx$qDdzWOW==$~Bp5f|>}!e@jqnF}{_;YKbTD1?9PCU3hVG9E#-J(u1`BbNDmK$~E- z&pUT0hqbCa>`jWWb~}Zqp5;F4L?Z=#0F~9LmdjWjD@Z_zOL6ZalS zW4G!*UMlnVn+Q7=DU$r>jc_*>P*YlnyAg}j&V(Xxj53R4{8@pkJjU?$&GZv3RNwe> z16O&N;kd%kgePO(Gze40A{l>%;G+LB-1w}U&jGX6>+ zNnNSVO(wuT!IUbo@>ry{>bXQY>^_TRH;EZ3qo)LO%h@=G>O3L&?F%?JJn>V1cA`B2Q?0{MUnisrx3?G1;qYNJbCY( zY*O(~FbS^8f4a&>N~#J-A!q|=J!rv}0^-_ILejPt5XV*?*B zfPLS_cOZ%1)>C$niTG|!%C^N#@9EF_5l$s*SDvR)SVIi?U_&1Ibey14OkW=}RNtG2 zzBeD}FDc}I?`DvO9aeJhc7GDO-AUfu?l1e2^*hFspdB`Z*>Q&(p{?JMk6`cY$R~g7 zm`Of-H%AU3EqBgxQ(HUnp@}vU9d~h}6)E2qpS1=)9K641O-kq}S!m$t;P4oDIw&Tf zyJ?~@9Sr*$c}41g-)KlnQok>x`k@0O?Fha0r+$H1=s=n-1du?rZ+C~7&pQp4=?tVT zOOrlUPon9_d^I%CBj#s=g?##KCsKJKBH%O)9jS`cUg#HADn_x3gBb6U6!}zo8@UUK z`Z+x8^k0bEQSnz)91thQ`fMWq+*$q)*#N9UU?n6S*jRENm_GIk$iu*3`eJxMp_mT^ zq!!o>hShm7`C@1JWwHT!4lzGmBpujDa{gktQqqAez8LHCEvgwnUs4RH7fKc6;Y^m`ojAB-)yGYdK@PG++sxL`PN$sW1@-GZK9>uYhM84cv z-bpr4y{;TvO45NnP0n8qSB^)M#g}6dWz?1M09UN8Pz!}y3_IA3V?{*1(ix9Pr(6k_ z(@6g--H5#*BEtJFs0f_)21ceYaW%w~8b50^Eo_isl=_!Y%#e1F4Gj@-`tQYmitnc_KN_ z@{5s=B2m|wKJte|{x{?okqrzVa6OjTe+`iHFuL24v|sH?VMm(nGF)`5PW7438wio* zLM@im{Avr_FM|BoM!xtpR@pkZTDft=)*>|JpyCQ#78n>Psc-wqZQb-$=roj^kfv#I z$omh~G#BV5(Al8!FioQuj`pCHpqFwrtp@bl;hI(tS}{V?=(_QvOVjM2o5pBb8mN}1 zX$7F;CZarO$t08qjn7AU(0it!yajuTp3`6em%OYQFaSDv77T!%0WAQ{n+*e?M?fn< z`^|v?&}nmF0Q5Yl6@z5h{hAgH`V8n$(62yegWd$)1p3VbnzkQw>w}u+kEP@2KQ*np zFZMJS=Ap&_0_SU55op3gnzj~n_ySGa4f=eMrkw@dU5r4nA|L-39D!c^H+GPqz6;?9 zv^!`iXb$Km(7lT^jb4IpgEoLpUZQCc=&{wHDWHDKP(e2`1Cc2K-LV`6L4O9_4?431 z1wku8ZJ5ea9zj9SH$W$V#;!y;&;rm((7lhM9O!jW59owdC>JP6rJxSb>M~Rmw8i6^ zRtQ=E8d!=ZS&K|Xpn5exj`fKfq-j-Z$3VGj@2Ygz#n8|by-MDTPs$M9l{!vW-y(b# zWvS1%l-s%2Aip3>)7aavnXPP(DH$z&EeDh!KYxkpZo?-A!|4X7q|22{RLtbZ0sjv9 zTTSFId*#=Gm*!|1dtX>DZy9M3_T5q_=EEM0%UU`wyb3#IJ%B#I3r78|0ckAcD#r2O zIFRV2<1ofIdsEtLKT1>(Luw5aPLJ0#_VTJ%p;C`f+6MU)Oo+cWfFTC?{gW~5KpQJ? z%v*u$z`J20v3F6u_FM4);9m;G%G_Tquk9MP>5FKUg$hHlMQU&rO4N< zRCi&J`>sPid#3sY&F@A&WS06{OF6!C5(>3%4;vNeBRdR4bV1pAJg@1u4{+G4zRNg3nL5c7mIp#YdNyTCM=K)IBQY!Tgzd- z1F>eTE97lLZ|t5eRCl#TJ1ph+7s$8T3ro)zY8rjN+B!@9!Up+X$S=-Ue?z_%z3#|Z zY3b|E)}Cq(l?P#!OX$Y-(swKrUMyk;TCan`HZ09dVUJK) z2L<1MF@a1L}_6Erpz} zCe~5v++al79qW12BK2|P#}`4Ki#4A4p%dliccs5Y%cq z^2DtMwn2U+@>0pW32o%I^2Tbr;@02X*5B0D-^@mR-#^#O+62_UUl^5 z5eX02=TV9|z*2+Y?LvNA@P1^AeV@X4m(#)6VxSWfuMzgmIzuQXr-W4g4PIH zCuqH(4T4Hh5=%Bg?VP$LhXB$9bqQJ^XrZ7bf>sDxDQK0THGtjIA6QLsKS>{uO>uuoAHPlU07;DVrVj@#boP)bzqawQNG&8g&ABa@yLFGcd|FdaZgidM0lP)=tS$$&T5+I6V09Hnqrir4=(j1 zp++fvKzaJS5#f3IAhPhnU*T=Nt5B22qx3nWev-vlj6P4)-WBPIO+N>iLJW=+Qn zm0NJDYaT~Z6ccoiZkYxcFVInnQ`9jra+cfN(CdA^S{^>TJw!kamp@HuIL!s|7>3f9!w(;Dv zuJ4QaEjNxuH^w(lFR{QnN)t?m+M@trxQVgC(hov^wn@LMh1)4K;nM}a(1dTHIL_8eB__d5QJ~a> z55~2c23ldlmkWHY39l7+r3v@JWtZC71U#8mN*^=H5YVP}&Ei#M`y^y13tZpsOi<^? z$emlzEv}4Z;(-g2S{^UQxE}`~{l$8+o+&n;rke%cAZDo^$0xvP90g*Rqc69I7~XtA zJjU?Ae@K#@2GE$!FbB=onQs{0e0jLca5_{mdJQQ@jp`|f%T{`Bdzc-2r!I)jeBQJt zaC+F$d`de{;O0}>)0wr$gw%ZScfG)S6!16(h;Tj@xcNBkkJK8mU$Q3A7|e%ild&;l zu}$P=ShpiPUad`(!+jP|F*aFk+^dA5`DiJ9h(isS&pB5M+*EVo@f;;^SS0qqG)D8O^nL<2 zAGxN_pQyh1;P68FUjpzroPNJgzh`_dG|i{K>jjR()+}`q!_OBR3u@VX-Znzu<^#Q> zC{Fuv#Qj3!^&D<-HYZZ4z|ALeD+O-;7XURBrxqj63JvoK_!3&PVLu9~`BZoh+(^(G zn9o9|2%MkrZ#>{XQ+=uj2IOLhgQ?qIPYOlziT^DEpIyLfRx6t98Z#jK@fq8blVOaw z@59Xx4aR&neyPA4kDPOx>jZB8cLrw!Za&Bziu*Ha-~9g-QaBzWng89ue02+5S5`36 zYmF#wKBE7=z|H^Mz{83~24U0kSL;!Gb-qImj~oaw4Qi^me(7Ut5*M!$V)GgEXP8*t z%}9r@g_ijrKHLC~{cdPv2rlE)i}{RvU$s{+)NmvtkxPV@`F{kIF>#;mLM+amdyk#h z2;6*<{{V36y~Amt(Dd9me(1p_nJV;hAVb4z@_!xB$wLHg{$Bty)n5?yySBUJUb1zH zqi^4LGyBNLEb8PG`Hy#(^_SnZsF|s9zIr=Z?xX&mDtGqj)n`HRgJesqHt)_KBTRe6>-Xe3EmZqEWk-OWk8R1WKID}9*#v7Bv`T%?iL&~w zi%u)m+Z5LCTy$Eg_8*4$^Q5<(=BitV!Iyr&B4~noo0isZ_HAZ&-91q`sE}OZxqaPS^UYzfnv2{fbWg)&9AtjsDO? z|DQKD?jRQq^c$AaOT4xGnWzq;7WzF4{@0FrEmuws)NfWQu8<8E+LAieUv&+aBZKul zL<2YBk=4b+5sZGX;`xQ++-u!%c-8M#ZV8ctd9{szC{FC*;;!;irA*`+A=?9VXe4>+ xr#9*nBjo6?g+i+V^MWQjK4aC7N63kR`mS}cP5o_z+&A!>5WaBc$nCG diff --git a/build/tensor.o b/build/tensor.o index c0df21905f23363c7025af8a1c5b275c3572472d..2fc99fff63ecfbaa715200e14527c207ade85585 100644 GIT binary patch delta 1156 zcmZ{iPe>GD6u`e3chlY6|L(SyJ3H7KldkA^$WW2R7@3Q@BMMRnk@8T~YVHPgkPU)l znv3y{LD(URPAdzsM1-A$*ux|up2{LZ^x}4?4ne(dXD2_59=w_Ne(&%1@tIG#l>)ac z@hcOn?Nb7bMvh0*s=ipdq8g4ii4dad%1?PsHIx!`HMByZ#le*z;ElLTMoQ8Z(@;#1 zz1Lyh>*`!r%2+`RrQENWi+!|6HBx?K(!96FbxfAk%Ny$3&%DZ+#$*|8#-hP8xBrHw z8L6bMY*wc{j71!1Y=T@L=W(hT3;RsfxV5Y5>$|l5eh%gvnxIh_2$fWQ%f_lPvV)UA zi-sU0oGdO0x)TmgJ5IgkXcZxG9zTlLf1khLIaU|p8|xcD8*7GC+EGj<2As`1b%X>7 zM}~xIzEyx`1uD`BXV^iA9|DP{;?=Vs>xA6GOeQ-&mz|x?&1SN53uO9YkNg(`mu!MU zyVKp%ZL^}>IVa|C)Ouqx5a;J%!p{)0s$=q=2+A%bz7*= z3g5=g@lQm2$`O3~l5O5a^SKBw(t~_k0`~ByC9E7x;-?aDO&fIGPFPL0VO8sOZI4@C l8Y<|e&sg3H$m6}?gAU6(PnWl)y3F!;)Fn+=C~U1y#0JOAhbZFxP@;_h9Emu+{z)kR(@qu&2ROXSS0zAp`lrl^p;q z4u&er0ITYyJEPh|TaTv@(t3}jWzY85dGto zqucA0yo?r-H1a=K#J#{Q=vagTnwa6kN(>#3^?GBk+o8F-uhl5gz4oEz sPpY2#C4aAn)Cn6D9nrR`c$5796>r8W>y}#pKJ<1-!gP96)_VwjhFi@lM!H)ntWbHcJn8h z1V%>V&B1ckjEo+W`{i}`G@1R1#DJRIG#ULgeCov+JX#;r542aZAvx5#DGv@;@AWZLz$__m?=OLp!2a} zyB5+ptgi48F#a|wOSX@SSd{d1&?)ozE8XI@R)+vjyAbk6t+G}?4@ST^jtB)2_(`}( z@`#OE6Y%=9$~^Qj7Xf1hg3Ej?2bo2rzt?s4J2(86F4$QVOQIoi;cl5_lilz78I*Y` z^#ZH>E4A+Du)XrnN10Hnp~jbu8C6VSIX<9nexg^}o#^`VViSZtjbLf`>R5)Q{{Wb_ BQTG4< delta 308 zcmZpW_$R^_%FD~e00fhE8>T&FnaC%@7%@>hoGpbXm_d_waz2OM#w$v!jFOW(*pw&F zW*40tz{xw=hRvH%1}NgiC^7jmo8RO$95IYilNmXLS-F8GDNI)6vmUR^Viv zyqs$tqt4_oZfi*ypehX{mCQvT5kF0~$*Z`ffPQl15SP8hVUwGmQks)$#|Tmh3LQRn IAmjvL0FACY&;S4c diff --git a/norch/autograd/functions.py b/norch/autograd/functions.py index 9bbddf5..17ce574 100644 --- a/norch/autograd/functions.py +++ b/norch/autograd/functions.py @@ -20,7 +20,6 @@ class ScalarMulBackward: def backward(self, gradient): return [gradient * self.scalar] - class ElementwiseMulBackward: def __init__(self, x, y): self.input = [x, y] @@ -28,13 +27,22 @@ class ElementwiseMulBackward: def backward(self, gradient): return [gradient * self.input[1], gradient * self.input[0]] +class MatmulBackward: + def __init__(self, x, y): + self.input = [x, y] + + def backward(self, gradient): + x, y = self.input + return [gradient @ y.T, x.T @ gradient] + + + class PowBackward: def __init__(self, x, power): self.input = [x] self.power = power def backward(self, gradient): - print(self.input[0], "@@@") return [(gradient * self.power) * (self.input[0]) ** (self.power - 1)] class SumBackward: @@ -43,6 +51,4 @@ class SumBackward: def backward(self, gradient): # Since sum reduces a tensor to a scalar, gradient is broadcasted to match the original shape. - return [float(gradient.tensor.contents.data[0]) * self.input[0].ones_like()] - - + return [float(gradient.tensor.contents.data[0]) * self.input[0].ones_like()] \ No newline at end of file diff --git a/norch/csrc/cpu.cpp b/norch/csrc/cpu.cpp index 6601593..aa7faea 100644 --- a/norch/csrc/cpu.cpp +++ b/norch/csrc/cpu.cpp @@ -75,3 +75,14 @@ void zeros_like_tensor_cpu(Tensor* tensor, float* result_data) { result_data[i] = 0.0; } } + +void transpose_tensor_cpu(Tensor* tensor, float* result_data) { + int rows = tensor->shape[0]; + int cols = tensor->shape[1]; + + for (int i = 0; i < rows; i++) { + for (int j = 0; j < cols; j++) { + result_data[j * rows + i] = tensor->data[i * cols + j]; + } + } +} diff --git a/norch/csrc/cpu.h b/norch/csrc/cpu.h index a8c5ac0..e1e3744 100644 --- a/norch/csrc/cpu.h +++ b/norch/csrc/cpu.h @@ -12,5 +12,6 @@ void pow_tensor_cpu(Tensor* tensor, float power, float* result_data); void scalar_mul_tensor_cpu(Tensor* tensor, float scalar, float* result_data); void ones_like_tensor_cpu(Tensor* tensor, float* result_data); void zeros_like_tensor_cpu(Tensor* tensor, float* result_data); +void transpose_tensor_cpu(Tensor* tensor, float* result_data); #endif /* CPU_H */ diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index 52a8ed5..91b349c 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -188,8 +188,8 @@ __host__ void matmul_tensor_cuda(Tensor* tensor1, Tensor* tensor2, float* result 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); + dim3 number_of_blocks((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(); @@ -267,4 +267,32 @@ __host__ void zeros_like_tensor_cuda(Tensor* tensor, float* result_data) { cudaDeviceSynchronize(); } +__global__ void transpose_tensor_cuda_kernel(float* data, float* result_data, int rows, int cols) { + int tid_x = blockIdx.x * blockDim.x + threadIdx.x; + int tid_y = blockIdx.y * blockDim.y + threadIdx.y; + + if (tid_x < cols && tid_y < rows) { + result_data[tid_x * rows + tid_y] = data[tid_y * cols + tid_x]; + } +} + +__host__ void transpose_tensor_cuda(Tensor* tensor, float* result_data) { + + int rows = tensor->shape[0]; + int cols = tensor->shape[1]; + + dim3 threadsPerBlock(16, 16); + dim3 number_of_blocks((cols + threadsPerBlock.x - 1) / threadsPerBlock.x, (rows + threadsPerBlock.y - 1) / threadsPerBlock.y); + transpose_tensor_cuda_kernel<<>>(tensor->data, result_data, rows, cols); + + + 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 index 799c056..bb06a67 100644 --- a/norch/csrc/cuda.h +++ b/norch/csrc/cuda.h @@ -31,5 +31,9 @@ __global__ void zeros_like_tensor_cuda_kernel(float* data, float* result_data, int size); __host__ void zeros_like_tensor_cuda(Tensor* tensor, float* result_data); + __global__ void transpose_tensor_cuda_kernel(float* data, float* result_data, int rows, int cols); + __host__ void transpose_tensor_cuda(Tensor* tensor, float* result_data); + + #endif /* CUDA_KERNEL_H_ */ diff --git a/norch/csrc/tensor.cpp b/norch/csrc/tensor.cpp index 782d365..6fd9f5e 100644 --- a/norch/csrc/tensor.cpp +++ b/norch/csrc/tensor.cpp @@ -461,7 +461,6 @@ extern "C" { stride *= new_shape[i]; } } -} Tensor* ones_like_tensor(Tensor* tensor) { char* device = (char*)malloc(strlen(tensor->device) + 1); @@ -535,4 +534,45 @@ extern "C" { zeros_like_tensor_cpu(tensor, result_data); return create_tensor(result_data, shape, ndim, device); } - } \ No newline at end of file + } + + Tensor* transpose_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 = 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[ndim - 1 - i]; + } + + int size = tensor->size; + + if (strcmp(tensor->device, "cuda") == 0) { + + float* result_data; + cudaMalloc((void **)&result_data, size * sizeof(float)); + transpose_tensor_cuda(tensor, 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); + } + transpose_tensor_cpu(tensor, result_data); + return create_tensor(result_data, shape, ndim, device); + } + } +} diff --git a/norch/tensor.py b/norch/tensor.py index 3689357..0a32a61 100644 --- a/norch/tensor.py +++ b/norch/tensor.py @@ -302,6 +302,10 @@ class Tensor: result_data.ndim = 2 result_data.device = self.device + result_data.requires_grad = self.requires_grad or other.requires_grad + if result_data.requires_grad: + result_data.grad_fn = MatmulBackward(self, other) + return result_data def __pow__(self, power): @@ -341,4 +345,22 @@ class Tensor: if result_data.requires_grad: result_data.grad_fn = SumBackward(self) - return result_data \ No newline at end of file + return result_data + + @property + def T(self): + if self.ndim != 2: + raise ValueError("Transpose requires 2D tensors") + + Tensor._C.transpose_tensor.argtypes = [ctypes.POINTER(CTensor)] + Tensor._C.transpose_tensor.restype = ctypes.POINTER(CTensor) + + result_tensor_ptr = Tensor._C.transpose_tensor(self.tensor) + + result_data = Tensor() + result_data.tensor = result_tensor_ptr + result_data.shape = [self.shape[1], self.shape[0]] + result_data.ndim = 2 + result_data.device = self.device + + return result_data diff --git a/test.py b/test.py index be14ac7..fb6378b 100644 --- a/test.py +++ b/test.py @@ -27,12 +27,13 @@ if __name__ == "__main__": #d = b-c - a = norch.Tensor([[1, 2, 3], [1, 2, 3], [1, 2, 3]], requires_grad=True)#.to("cuda") - b = norch.Tensor([[1, 400, 3], [1, 2, 3], [1, 2, 3]], requires_grad=True) + a = norch.Tensor([[1, 2], [1, 2], [1, 2]], requires_grad=True)#.to("cuda") + b = norch.Tensor([[1, 400, 3], [1, 2, 3]], requires_grad=True) - c = (a ** 3) + c = (b @ a) * 5 d = c.sum() d.backward() + print(a.grad) """#print(a)