From 2fd5b75ff085acfed44a5d812b189ad5d3679737 Mon Sep 17 00:00:00 2001 From: lucasdelimanogueira Date: Wed, 1 May 2024 14:44:09 -0300 Subject: [PATCH] Fix sum tensor reduce cuda --- build/cuda.cu.o | Bin 62224 -> 61720 bytes build/libtensor.so | Bin 90272 -> 90272 bytes norch/csrc/cuda.cu | 41 +++++++++++++++++++++-------------------- norch/csrc/tensor.cpp | 6 +++--- 4 files changed, 24 insertions(+), 23 deletions(-) diff --git a/build/cuda.cu.o b/build/cuda.cu.o index a9d57ca942870e2b1955b15156ca925ccae9d0c8..24a944b927afa39e8da54101b9db0d3179816b97 100644 GIT binary patch delta 7722 zcmZu$3zQVqnf~vs>YnbN?ym0Y_e}TnJbIdW^w9O{s%l1Ngkg|}Fk*55Nyxwe<0A-@ z0d!!|9t>s&PytIgy1OR6QDYK3lenz$Y+_Dg$gDByG1(AA5#2134Im~aX5D4~Tip)B zVeg^q-v9r<``7=vj~e#`8U$Xhr7~iVy@g-V3 z3|D&vZnGu##7d58CQV);?X&roOju| zc3HC)^b@c3C_X(FS(qnnmW)rqT8&SawI82O>t%cV(Zc29Td(~gAwH|Hc#|)H_N>Am z?{R@({c_PV%iHH7K^zsW#eG>4njA%}U-xyBn03dsJ=XH;yd+^uzIomB5=;LP{ISN| zVhL}l1d7otp$|@>S;AbO2L{J!{EbYZ>ybV^g>FE4IfoByjNOR8qnBtInvh}I!7KwZ zNIyM=X8Dfe+F^|;u%15<8{fC&;o!KmX4GdLzA;s5^aI_7Q~FUX;O_$g=x_=y;?Sl~ z36R)YQNip%d|-sWOW}YYyd%KZ<8Wx;0#}w-d}8*aJhKgZYz}%~ zRS722;AM#i%tN#=IP7q+FkWogTk$Ybma~AOJTVfz z$YSDPFR&vJE6Z5~(16Ugggqh59QYuQbSyUM&V5ePH+(7>NKhDVx~X$j2gTLSEO+*Pi#53#5p!y0)V68Y8&1ymyxDqDBuBBc2+T#fuDv7AHn(YgvkI`e+&g!|qz#d&)(J2Jgcn z#Nb2s$%_KKiV@fjQD)JQa-o{d4~$S+|2f{K!{N*6cpT}{q{6XH;I{WD_Bh=Kj7~y; zl?>8*A@z6If8Ndruybo2<_P5&6*$f4K-ndG7w{v^*T85+eiNe|h!2+WPcu4=@=}F? zA=$aL9s%MMh2C1Gzr3xm7c}66iu@);!b zv@*e#yE3dC7^W=8kOST>m_l3bZpMStE}3#fF{PswIhxWm>`t*FW|(qBODv@Lk106C zJaQ-TB&`5gNU;X4Lxyc2Q^03gm6sK;kg||h=G!56q}j$~uT9%yRs&dF^w%S zhCv6&&XQe)6+k01z){DRT?K$4#hPGmXA0~p$0AyVnO8}JDv;3Efj1{Sf2H_?$X%JrvLluMU5PMUsX3F?H z;zOvfg147p5a*{W2KYrew4NutVz4Ocv46y@(gmlQnWwAWg%r=E<5k zDo{*e-DC|MWfAo>Oy=CCs{<^>j0Zl?OV)+soFjI?QR) zuroNBL)_W$FtNAlgnigBc`C(HXz#bD`mWfx*y;Jmrx&q(or8I}lP$8ViM<_0 zN$dh*Zemf|$MGs0vlpNT&9_kwzFaw;ac}>|CZH+Hv18`=?2_M2=3CESNLxc=@yH{% zlv#V4EB9=FudyB;Yp3&StckHMGG@7Vb(5pk{9WBN-e~RK)kQ}dt%+S-$aC*TUX!(8 zcNZD8b|C>7Yhrg7U6rvsd%DPvt@(SpX??S`dk^xOtqJ5Ew%qr0(?Hf*a8DO`-P(o3 zoEB>WiI*(*z1>u5w-(&nMV_#BAyGfwny?A?eMrdGg8RD2KUuqwh-=ma5<4yT-fr?& zYyRGD_o!L+0IWCmrme$!B=iyN%e{&qUr zXYJmfwTAYi-E22uZ3?K#dU$`kdvr-z09Nq+v~`|gzfk6kR<2jrG2%P|pKpOu5@#?j|1Sa@bJoy|Iv0 zf1rgS7aW*Qf4tUOJu7PMJ5Y~Fef~haK`QIZRQgIRY>(`sUxxT}Lwy-_NlFOweWX&9 zw?(bCgW1%VD#pz#qbHvs4}lBTtp{^Nvi{-VG_q{`CkI{5QcB*&UnR@(c7D21{MZ#x z@8r|&_7u(cuQ&+eyOqz{FKJRyU%3d>h_VQTn1cv6&|2XQo(R2=68`3+b{}y-UM?ah zt*&rDT4E!u7f+WPNL_1Tr6GoLyP-vuz2Ri zM`G*a31`UZiQ{$DS-bV3LsJU5dZ1lk4jY9TkT*2b$!84XUEq76)zCIUTK}e#Yf+4E zIo*38W0ukMdT?)njGp@p8i$otL^NX|t_Q-!E)tighwzI$mns3ydW=NX6V4V@ng>mK zt`6db7fDU=2&rl&a59qayt%qs(5n5_jZ|(PEZ*%5>fedDt*WT9ndIc+;i`<9dx2~o zR-{hP#^Hbik9X1uiJ?8Usw|neLRCk&U+>pCJvCk~Zy2>+ukTcKRx*E4P2RxwYcqc4 zN}H=GpHqth&8g-eAzR3`LPF-CkPQy8pdB7nR(4vL6V+U(q!U`=n7JGLni?0Q4Gw%^=Led!D%((#p%gg^Q~pi3EAmLCSzl;GYEYLx_B*blpsc9` zb<#}7?+{@WjTChpM*u1-$Mh25IM*i^6?B0(-sV!MOD)(DwGkN6i2gR$tw`EAQX&!k zL>Z#nA-$yJ1VUtk*`dF~b?9h$lO57ah<6a7N)<#3p{j1OpkI#ARE99#;cSE@4m*nfPe@UTs@T4wl9{q3tz}i(1atwwZQ$>LNby}Ruc}YD zX|^b@1>Z+REh>`$p$D(DClR0>VgL;R7Z6Y=U5nG*WzztCPR&R*UB4fKv!C zo^w+jEm#M>Ka)AKBtICA{4JtvAKIU^5AzTRhTU8~FYuXv8jsyL@ zFKQYY$B2lZq5R)xBWB(oQjYuL#&KW1V0=~NS_yw)MO@VKACk0LS^1BC za*U*l&oq0xhc@2+`(&)-pAN#iNuQuYBroUILxcaxX8(E+CX%5JRaz2HNxw|yl|oz! z^QSX@{7c}^lVY9G4Z_RGM4gKLi^;TqJ`u)~!?ntIlE130&*+y6fd+3{In!R9R`U%R z=aD4W9}E}^qZxJX!syVN+cr-h-ga{aKdTl)+?@z|Mgk5!s7ix9yQBOQL4EbWrVr|% z9rt0ax^-aF!jv=Mj110;HVqG9twv~|n+p9V=93fgS*7xB#cp&vm|dkhq7)}!%~ z@w|#*LZK zph+{KG-npJ*>n0tnpKN4A*6juolTHcu+oxx3aYYVi&|LR8rPEKMhwk{b)_3;h%{Wb z%CSvMYtuz8T37-pZ49F35NX9T7^213M9*V*Cy;iNg!UJjEyx#f%|oTa0CE3I z(ETUM9}9Zhf^kI#{@$Rc31TX4Z8rqnU8y!hJCyRQ0`AUW6+Uk#(}p+vBFMi8HOBXW z`+3kSL(G&RX{5k)e=uG+=FG^_V(@&bg!P@uLwL0{i%LONKlVAFP+B_V6V8k&J;G}0 zT(CsH4t>&~YW#!%_sd{J`j*%_q+_`ha&v)xCE2lMeNZIlo&DylS?cUpUBo#NtfDi- z!Z)1!>zq&}l9=Wb-S3kI6N|v{8_qKmjlt2$VATuwCj`;;0(7k#Q9f`we;C9!{xwMZ z!L^RMy1`9F*R~R5i)uhj6~E?=8jHxH9%YX597z;6f#Z{4YUnpy)SM4#Z9e`N)KGA# zG-%8<7XV4sN6kGnGB_f5^FXezjufx2hPG+KYe6m|pVr9QV9+xiYG?czsDn!@!6tCr z4WfJtnf zLlP|3#)N?Ry5zb49Wg>6UL#o87Xj4oAXg-9kipH0(SV_qapOH)&|go3u$=c@8)1{UCg- zd9AaDHkKRzsCAGjO`r!Ec@RYtQ{P9_}X;g`*pnVI7idQ!yJC;o4Uc2_Os|FY0|$c-O2hlLtt=Yq7f$Q0dp z%UN+Pf3#P))#QW?W-8VoZ!q0o1Lx8teZwQ1njed8IyFBnzdyg3UR{lfA26-IJe|2} z{#A|s{|4b;U(%0XjsH-Q4jy{dcgX#AveY1J4fyxZBuUCWVUefx&jZSM5k;ql;K zl%DSym`ki3OETj-o~aG;vb0tof8<4%N>Bc5{MkRgLg?srYxtZ_R~@pRJ=aCgpR&$i zH*&~oJKsfJ?_0NEH;OXs>fg7{U^j-ckGkk{hwN|g(T`f_VU#21^+VPf?9BJAwhLWU c`mS{gb`Ku0p1siJ9{mstA6wY^=xqA`0HooM4FCWD delta 8295 zcmZWO3v?URmG`|Fjik|Mte-9UCp(s8*_I=XMqj(JlQ>DiBn{*cNNFI%fskZrBIJY` zV=5;hfSr_>oqnKu*yR%-+5Q49U1;~vEd`pU9H5+TPPV{q2!V!Dz@$LCU3R_qy^$@2 z%sJM)`+4`?cYkwdeD7&t>K$P!N~iwxc;_>n?k!&jeo-J~F|N*T)Bo0ag}iv1fgWE! zeO1>wMc*9_PcQF%Skg}|k*3csv5?X_2|xh9!B+MAB*%xVFvjv`Gvg0jfd4(mM=rom zTNr=+1$dX0;oJ2)2V8-l@^!-nJd*zSKwarOGUIH#fchrKcU*wqV`KbH7vNvz_|6OP zx7rzh^9A_talA+Y^fxCw`ub)2N=}368|WK%O_gtmf&V`Gz+F@2gEIa@^ntsky`rIv z|0AB;T@$U~Q5U;@fhTj7nJ{rVWL%Av-i;hqe zv(2u5G`N9S^vf>UKtzEdD{%GYd2}*qOOT+pDt3+!H8Y9XUaO$hRbftb^wT0B+DzCq^Izrs@E3dxx z(tc8_|L@8H{p6)~;^&m+%NCYcZal>_;l<9w@*rCF5QACd@2!GaWG!P&gd?<2I-p%04Eh#Rl-a%u5khWh#PM9&cX(OSwMkj5H7;(EHZHT z1qJLQ;AQf-0u=DY2?QKWg0bH{LcQz=v;T(snE}=Nxjaur9@0zpB2cEFum&wI3!GsZ z$^ua#{3EXL$WiFmKa8!12gV_mRH4I98Za?!qm(TZGEBj94O~^` zKT}d5e6u7m0ainRWw#(Og8|;f1bxVfbc=;;;Pe+2*e!@gjxqXSC+r@^{d4?h8Ej>x zXzV}jgqjhLhad3NEH|)?g2+R}cvc=;4Q)lR8WIrk%0aYPtUe(o5%sVW)fgv0gU|pp z*k-_%@XeGm18sW#*7ni`>WYRsrVUNn8H3bmjE?PJIly7 z4f8{z0uMO%Kl=OWQS`>O$j%OZ8Uxxmf(yI~Ge*YTHzM+O*DoOpZeSUO9U2Nt&+&{x z^96Psclmgs(T2)#Sw`K5rtm5@QhZ+in{)C_{Xvu`dI`$MfYyw7y&uow#TFAuxW?%_tJ&MN6UL(F3-ncLWD7w7EsSgk7-9mtJIVse`z^7` z{y6g2qg%@So0&n@?sP{wZJ@avvbp^s zV^-k%nSqu{l5qn&D1_$rGY?vMPL+64aRqWE=G;PiCB~c{rNEpvZ=ho?hUEj8KNqr2 z3e0Eo{LXR?u>-LXar;UJ!xok#F$#&f4eSI;If>>@Sk)&vzl}nw%)iBh`yG`e<8+Ba zZjLS>U8#78)4M5r0reXeqC#xEEN6HhIEIa<2p*Oag&}M_ybP;yg25w{F6R87;KMwQ z&rX3qc?8Kx7=K%_u5NyQ=6$=)&t{rB{DWf$8ND&c^cx6;AflzM0@Jpl&sSf8af7e9 z6=5I#0UC(#g7Eopr0aaRmJOp-{2?}MR$y5!v5-~}W^*UPJQ278QDh;+aks2cwBm5h zo0(yN=wLzTwaNKG>7IkJ%vdWv`Zv~E_X_X)Kms&4*dznyZ43<{DL{6~DviC8G>x&y>oo)yEIRbK)V^ zTIF&T7y|I&Uumt+v844^TK580Rrl<;KJmvNE$*sv^Rw%49U2@_|d|jU!xOmCIgG=y|Pl zfei&zd5jxWRaQ~v5xLSyqYVg~DjvmgAIr4TDWhE_MBu!{E_Pxrps~^knZgk>Z9y_V zF0ssTKr4#nGP@fGh+t)YV9$46mL3li$jcaWB2=WD-^@nj(+WL5XYlzySGQt$hHfi* zeXQ>J#p~8vSBn70@w&Cy{PQp?TbrHLIjan#DkO=OZ56DrnDO?=tHfktBRGdk903;g z^1~ho@CqNXc$pMxjj5jv?Tgt+Se~)NH@hoIh$sAfRqzlncIHt&abl-rL0~YOFvmnU z17}Ztg!31Vp}i`LSy`Dp#^Z7>i>30zWG;(6ZYG}#e3S_^mj{%&D#94gCGilGXZc(b z&*ud8KW>q=lEvJhs)tD@ORB2gj{@6@vd1U`G=hgnHZgY_jf=Hb+067CFbfe_naer+ zib6v>h4N6_#or3p%wds@l`p9Da*NA%{q|0|8E~F3M(T8vK6L<8}sjz+= zflGDko&EH=c75>9KJo>90)eBQ`f&t8y7eyH+N}@X#efM0B=qA5H0ss~)T-%&6MaO~ zCnoyIi2mz|ejk5WvP1AxW8X>u>h-3{E-U028q-gAO$PPCq~AB*V_>Rz8}<7qyL{ul z7gCz^HyGs$WlH%SmdS6{oAx*9f17mr$Cn$_apNJ)wFdzb`nCI``i}Rk_A0Ad^!xXB z`C#pZnp*WY_IK&W4tVv~_j~ok0lyEfEOQwT1}?Wvzxn{4>dR$~`eg@#`cdQ^|H_3D z!}=Qsy6E_NUH(QNyA2w*N#l0hxXE|-$>W$V#-*>hJHuEf@8&y>8#no&vHxJ7iL>e( z_j>dP4n`U0@q-KL;%)kGm(=PtQ!N-)b;{qGF{FjcccV!f(P6Cvw~;lG`U;wmOA$R| zR3@Bt>;F6zX}V#hA>|_DHvj9vSH>Cr{izsPsMp`qPS#E@y~iSzVrpVkjA@x$#qLyY zOblw*xdj+-DL*F-N>+qmW;rw|VmP1_4<0_b%1PCVBDt*vn59FTo zdy`I6<1LSRXfNL0^5Z~Z37bM%>=H29Oy)EMwRk`XrvAxfRwha&5=zHb0_}x(F}2@B zX)83UjSx=82H^5Uue_ru2@nU8hmJzh>rpcs)N8z^hrMwveaH(PH||{PriVS|weA6= zsE53!V^Evf=HHa>l{ap`?nWV$%q@es);j_ga|gliVD>sMeG3|MKXAJ;KX8}iF7Rf$ zfc_YKcx%cch<`Lpe0YNikNtUxESDQI&r9uUVy`cle~3!71)4JrX0Q4zv0Vx6S}cdo zNG`Pj9rwDfyTR*d0zA}z%6In4o4>koVR8Ei!7X$El-r5fr8>+tnVeVh`ZO(3iUlDa zF1`eWuCX*FUWU4Kgur9I)b{IdSOav@;!me&T1&?2A(~8l)0ft?SPLjqb}6oAwiBr> z8(R*x2cSI{%cBw5H3U4FHAI=VlQuQm1i|zUA|lfOzRiN=6`+H}vV#P&J4i6KgV;#v zo3u$g+laDvkjDH8u{OOH;>pAyk%P&n@T9e-ChciTkBed+g2|^T5x(tbck0PEf$}sp zwdGN^GDwtguID9@cKS_|B8~d%b1#YJ8cc--(wKas*_#*u;ZGtwMI8-(5EJP_5yW~rp!N*y39^nFnKR_|72pE1T)`Ni0~67rlsyD z#7PNvDZQU0Hr>4Q7NB1>C5zW3i}mt%Z8Ru}_o2O`0&dlGoKb72H>E>!*SDZ89a{ht zd(UGak3lJ#kL?9z42bZl5KrcD;meLeQ*I2}a%IN5L3Yq+Vvf%y#I?j2ek^3a822tg zAVWY_02~Bmk`f%AJJif_f$pP$%n@SSAT*}JmfHN&)UE!kwKjj1;7vSAyVRZ;N|%&G zs5mBNBQ$t|iYoNJ`VLK1}I@pu8J$^cQw~?PFJ+b0H}2yBtae zd~r2)BeXiF>zy}(a?<5#(l$7{)x>MAcrqOcd&N%!PW&MJ?=E+9wjY$2U4do|@5fw0 z=fy<%g{#<5O8&wn4hC0eH>m06T6-|L)nN;2i)usBukP9`(PzlA#Aa0l$2bt{D;t4&|(lx$JRhwxUkScFB7Dt zQgr)uKW}|76Klz7xm+fZ{a**I2l~C%U&UI~L?*3ad8nC7@C&LGE#|Ih5r1#TpzkAO zE1_vCeNF^p$O)7Ai_jA;q*|aO8T%q&%9*bPx`~iQbdkRNfne!H8Ip|MJOq(cYztUF!+wx{ zMh>LvA(DwbZ4RV66dDn>4WtH^tjG-DBtpiPs<{=370X2n7Ec?Pt_FX05wxZfi{RE} zOYmB{Y!PHqDi#)51fJA?Q>hIi$yf(8W|}bVa8FcAuW$CJ8cW2Twm5})FntEn0YWA$ zLJDH3)FTkd#@oYn=6oxRQ4(e=VOOKeh4Dz#3Tu%{z4;ZH3b3IVbaOicIZf} z^-MlnCAUBwo*|&6Ae@d}jzLL1Nj&L{upT3pO#D4@rE3T+{U7Z}XHZcG9ZN4=v2=NE zc^2yOSqS8w2!#`|D{-IpujmOB9LUHdok_+%Ew8lL>(G%&qQ+G0E?jB2O7m4}57cGq z!HRZevP^d-^Onh*3Q6)&k^Yh3Thd+G?21e-wM_rbgN{-@laIASZ@7>iBJS*?K%-#& zn7XnTk;d#PDVj~HK-%!B(gt-}8<2B=s6Ug2>}xI=gLyfg^I8zThP@QRnb=REKKo;! z+tJr5)a6uY#4Z|Paq_4CLkg*hFM;yB2k^g8_o=Fy%O}$%O*?Fto@kGDt8Yl|+yjdB z7OLe`Y!|DscCybBNad2TUZ8))M_2|kDIHq@BnVO~b>$`rECVtlZOAQIqAh*VLR`5W zq&`;@cc<@=#8(9Cqlno{1a$H3V0{aLZ&KgV6N0=5javld2~uCm{RA+PB4GIm^j$xa zJR`aC0TRgBOrgBG0W8;3OFz(dAXkxSULA#WUW=`TrupM=kXLD2K6@XyhOTml)%0hbA|WP+4O{N(=~zor)$)maI+053 zq7gN7#$?$|q>a###o44YT)5mu>q$>h&Hf%d_~%Pa$Gp3WIIl#AC>{>zHHVr@zS!ds zDyU=l3>_oh+&+89=F}J=eu(9B;1-%Oc$)~k1hv_nPH*bAK=_88(;#jE%R0dx0(#gg zhd}y6J-OIPF9uIx%pFN4c0s*1MLhWrAgXE~0uCB$TWrq~Q!DAp_nf8iq>5ij@WXAb zTGGFNDAFqoGzOFJQY$w3cd4EJtX8@NKm)5^bmtE@+CLS9x4iP@4(DF#*S4F;CFXcC za}yEaM6FF%9*LNI$-mTD^`1u-mHwq2f@-4OOm9hr)x?W5n9llaKg#Xce*N%<8@K-5 z^^u-hrv^%_M^btr;!O59dqK(d_#2bN7tH=LSea7mMUZtL07@ySD*EG50dq5$iQ~%_Vg6vt`pe!SPWp7O1@n~b| zVL=FH9u{aY`>+$7Z?pN?NHJIAdKtLgd?2o>qI!k-E1F z$RCzX6|djfE8ixUzox?UHIJ=uiHXeS{Pf|cEE+xetLZ1+dx6lAuj`}lXXv^k`V;T> z(X(gtb9gHp(UTwaQFvP~;%x+Jc=NxlpZlQCI(`JBj#o(ExW}Wf{4hih9pP^bg*~Ev z0Kvn^{$U?Ic0@mix5V3ea;A?)kLtylKI_y^jm-o4Z}$ZClQVvK=lt(^=bn4sz3+pa&s%mrZ@Kr7#8k`{vEJ z4mRH#`{i?=zJL8^OXI)%__b&E{*i9F-!rnk)Sm(&bFjH_gLeIcD;&?QFR5YW-h&UW z61ahZ)!NtYYx4AFw_dktgE%lCODZG=%Id0;T0*s)*&E;epVqw}{fmv*Pg(#t$N*?B zf9frA5Gp%82aOes(R5BxZoO;j!x!im+8951fxh0(^pOknMN{9*^x7BqIKwRt#;du2 zw^e&&Pq1jIWY*{f)_zmpdV&5UQy;rPKc|ZMnRS8wJyTzMfj;bHe4esIo$+dyJaI=6 z$H|UQtpJ|t8TlzqFGSGpsh*LCvVIMY*Hdkdx8$I#UxnlLRJW9L>=)0)jx>(IQ+;k^ zlVv@vt=fC}lMm3E$#NhpIN({AsxMDYwg4ZXd$}Kemu6WEI4s~BZPW9@duFuVGummJ zIoCGriG2g4N^|cYAP()O{b@2wdjx+4?PdJ6Y2@iNv1%gz+O)M#uiE?G)6Wp%)V}@9 z7H2iaGl+k@!v=yj`q6Pv;z7gUV_Ux&?+q2pZ1ULF&R zAMJ?=L_93??<fVc1hIh7A!fkG>?1>e;TRufT9PrvKQZNCDMMRHX$o!UrqGZR*r)2oJX>a-tAA(h-IfGXP`0}@Gr8;{Li+UH<(1X^JL+0}5_%KS0D7PX=UfboxE2FFWlXVENb9AS9O5n;1vFd9Z29;>NgQ*6oy;!r(m zufesaRCE~dy;y{peCRo12=F3iU^@g@L?hDiO4h%6jMBQZd{g%M&X+?Wl#A02cWeQ> zxkj;{WBtJ7C{(kOLHQnteG&2O@xB1JkJsWHp&hdV$CzAQ8j`sRxKQRJFj=;5WU>YM z?h=1DljCSFmKB(i+sEsW;F;mDs)qS5Ei3E<_1K|o-^k*}9(O!*+dw-tr)08hH%I8gVE`}=yG+dtvYUtLj8x+qJ_n6# z&dd%%BP(Fmsnh{eJX!8w+HDC|4vbT>W6A;VE|^2pZg0Yk(=3@%MsZ3f%XSQ6$noP@KZr=^8k|GU}+Go^w;KtY%ZpbZ`LIUz}49n=hsWTTpgQN3HJZImDF>HxqNI4x5_|v&Uj=3eENQX#WKZ7u!8|o-*+G z3I_{uC0lA%6FWO@rzlgrCov0%4HJvf+>Qsa%}#(LXu6qlaCLcm#FyV>eeqwzx+$=ZQADFrsyy=q5X~@e@7dfcEr>9vW)UPM_$e zV+~s8C*7zU{{(f7+7t?t+9?zup>>|@rh^G>;$$~@PJ8-f53Os`PBUGT)_JO%?9s+g z_0a01Hig0=?Gy?NTC~p7-Q-zq{B#e^v}sc)JgS{Sp{_&g{FDjfpQ0dYQz-m{b_#`% zqII6>Cbw(jXL`stA{!U`TLGU}f2Hoj#9{9CGyuVw&}- zGM+ADCnen$ldmm1;nnY@q=nI^DCwZjt=Hz#zylnqBZt^sY^Xj{@|!4QIVa}G%*tPh z9pd$MCDu7BAuI=D<>I_KpnroS$=FrpC>NIK)4yH^f#dpNj-*INckrZ{tlTH@#AYqV z#0mZ)Q53iH9r?naZPoIveB9m^qv?Uw4}j;L^7HLgN~VxovlL{%v=oG(g$P&D8sTN0 z2)z&!PCKcIehZ|<0&3#&Y74|OtflngHghGZZOyL9d%USTp+y#}d_AI=arvNaQ*rBv zB_iKwsml$la9jV&;>#bhgvY-XwtB6O5Z;$sYsL+WB4tx`K)XR5&1X9yomW&VpUCIm z0lpVn^U4;8=N`6lEmHoUtoFMgp_bV6a8LVQY(=SqP1}R4s(EQzTN@Pb!)S9P!7yuB)sRluB1+0~MQw3wK!ExqtTC zwTghWm88VNo{EH=dWvivl`>t98%L`xxX+U=2#?%VBa0dJ2B>KF4de!tE=Qz_OXu@7 zRaMTTmC20y^GfnEo?q$wsV%OqrF=>*2s9BRXh=0O)9R8B{P^= zp*2dqnfaKZdqWlTA(&eOwaT+Wk9dotC!1S?d!}2)8%|}yx_3I!_cUt;y~vD zd65(5;C)c`YUpnhC3O&r#G4z0_BLe*8uO`opceyd1z|4U@TqGdEBui_dU6jy?oN54y<&RG0p&tGAWSN2ZvqKQeAtN`FQ3PKe zDLEyB6rVZ-?Gm;c3)39eJV?l>O6POO2x*eVA#lD&WLcVKnCFqviCRGS#E++RzK@LT zc?^H%fVPS1n{FBPfk> z!k?I3-oiPtao~Kr^}o5alKFs;fbtsWMBKrb`lp+Q@8hwXx(@Rp=P(M(ueq4=VUZAJ zhKs#Xu)NOoh`HQ{1T*#q*PfNum6JLhrsWM}NBHZ*e_RM){Bwi9dI^M9G;cJ?G>wQp7xqBlkAfp2Nr+&6J*Fyp0Sw zQ)a}FD(6&|^z#{t%ozE%Ig^nW&tnvW(uY=|OlVgzWUSQ-T;x<_H05!mPHd!{=go=r zhZQ}l$Wm&3ssin@nwj>%#;UR~3piH7n(c-25?VAT-tsorFK6;-A=qp${FRiV6In9H zicV&Wnap}trHw!r1KbR*`QR_S!=+{UQ9Df*#P#6(Um_QzX@W2x?`@ADL3xZ36eJu+ zLbiB5<2T62e$P(*h2Lf<@G4qtyy3Y$QxE-(H-th~dKk=k<%3zDt zmw7Q-wKz<_pM*=mu}7dkZsQAhQHy_w(k5xmH(lg)5-)tWsj6q>##??L4HjJ;AiNWG z3OVqn#ndLKcYV9bwF!iSQE$7PSr&?AUW%rrY$)U7k0o6AO7Le#O3@`)!qRIbENy?vHW=>|wWKtlE-u%bDT%cLw-J-MO_xTi&gOHe835%GRMROJde)tABW5pkY`( z9`LEh1FhRO{lrb;kX7cXKMy4~jrx!Gy3f$gY9_z)> z*)*cmL93Ly7)Tu;r`95A~J8vP^j5N{5FFWS2oq*#QA{ zgtX!=41vNgJdRy>SCF}rgq2TdGAk<3n$1KYAYJBZ&5GENN3bDngdheUmK6wQ=XyfQ z4~aLMM_U0j2$U|6tG#$)*<_^iv;>giu zIOibrRX06hwvsvd2cFiP6oDoswF6pNDgxOks+8OC9v}jCg*NB%e|Cp*zqQa)Rn9#; zeV;%B*^$rp|q?jEd_NU$Uzfk^lchxL+ND_g+ z+wEwCpp1*#m2P`?Y*t=*Ean&l?pAjN{@y~GM>qf6%}+rj{{!HD=B^SUsEQEH$H4Y= zcPRV1H6dm$1;;7Lm%Ckh2ybrHf|QlzKRd0DN-gc;VQWIo>|{mtp}R=G0{xj`Isfkj zxR>1i%s+WrM{-y%*;F6UYohJjHn~0IBkO=Vcdk6|MH{gmbXU+$Pxh-;{tzcrct}uj zdhEX?^(q#DK zvLIJ`VuhRRf&5akbiTAedYps{TflO{9UFO#3#f}At}Mc@Q6ugZnc;k&x)?~TE}-5; z{ljBIRT{|UmHxuzl`yMW_=TJEi^mkQ-tBgDKuzbzKrLKcayNqI4)BPVf#ZG~I34`0 zLSw3KArwZ$oVb>={>D9f?8ea-I`_9rqJlTqa!yefa*0CrV5eOs^cfP!z6tFk;&<_e z>_x!$P}TZr52?w%*%Fri3+Ve^Z~h^u9@+~0tzGazmM6sGtmLd+Z{cmbB+7+M0k6}{ zHp(rkjpp?wUQ+9Nw!yb;6ykD{@ch5z^;=MPt<~@*M#SwToZpVOVs})v4lAs2zpwJ2 zn`kYGsOt#)0s{H1F1Q}JU7Tev@EgIlimOTh-EDUyz`D&t7P)Z$$0fRO%oWINgNS@D z@fJS7rC$CJn#Ih;ZTe?kGJ7`NpG-^Ad(@7L>wC1y#x;{+p)i6gEl-CNDW%Ir=p!!2 za;NJnG$e1fkR=sqDR&dWhowxXz73xoEnewnf2F?DM`jnlD?warkSZG0HIIe(gAmW= z{gwZ!-neEs;Y^RB)tlDn^>xw9S*&!TV6bx=Tx+pIGc2N_ZQ;Cl6Q}t^DB#~WoQG@GK z5FY4{y72k&cNI0RDkA)*;_q;iJkhnHVh1O@)78yg!jaee?;YK+waf7(hr^-25FoY1 zIW>kciO=zaal_b+n5i?2=MYaN4CB{`-X_C1iTJx_!wBJn{+lg^p(5^RH4OFv>YbEf zj3K_zVHgi0#w5eokGM69{ShC|V}HbpyRbjLI=X2#_D8&-2mAB*lKu9%=l~VpoR1C= zr+U!=Vy+JzAU=tB60zfAbbv3v?n6`&pIe9y5W5!{#&*Q-ApQg5al~g4I~N!=$z_I7g>T8~ml+28lr;Bp!&r{E z?g|VP@rmUaDB>^w22X_ehZTl#1d&^17(BjB=(y4_*sI{a)rP^oB@ZHQMEok^4#YG>l`2+I5Dpg1vD34{T%*o4v>F zVuAi6hQC~(zs-1+y=**@zPIo?cm;EtWdrVYRezPvcUm`9Rp8v4FU`-O{dWrt;}0|K z^r%T-xrt(~_VnsQG13~mKZ!5zYtVErzObu9V*>x?^y#}|c*2V=)?bK`B>mGu{R1Yy zw@CNYBK{8FzB2sn5`A?o2IyX@PcT`yRDT%J!S!PWqCfsK-4!PR{YWk0=_mdA?-`jc z)2(qNpID~Pilh9|GJP>al&@#_mu32$aXe|?75bfZBw}5TtHxDV=&vxHzFa@eIFBvY zBXuO%@Dg4-%^2`>7Op8}feqqa?Dy*5=tIb($(77kE$v#Ve~-y`uPXW8ikaHANX9R?);sIb_wK6uKrF}|9k#DpYEO>=%N3=nJ~Z%AR>qoKvWO`K_Cj_4CpY9+k-?` zNmP>3v?gZ#z!Eec6(r(lG#YhT=|R@5bEfXQ?|$EX z_uhI{J@-GSyZ<@e=Dnt^cLKk*-uc*pam2CTLpWN7N;Q4PmAS@vU*C}XgyEB8GfwY% zXMKyV`MI6{q*vYLSn)5KJOx4~;_&b)W!2sDOcPgW(?bh(zI^w5xe%X50y}qY-y$g& zP-*AwR*;1b=Ld=J|8@5ey4%)SD!U+H?sK_<@lra zs$o8^Ig3Y98XxoK3ns={dlvOQKHqxQ{EcSDUw7908+?AuS@YLh7=Qg)^S|cvA_Y(` zJnmFpdF-aVRb#phW8<)mijB5g@EeSQ!!{}g#raP$1`bXT@O(m^Ub0x-o1zzl;`kn zQ5v6&5R-Bt-lNI}yvxdCcsD5T;ayVf!}CKC1K%d}6~tz3%ruLRQDZxfcP<|O6G4rE zr~4G!Q=+i}oe@n2uquCe>LTUz(*|P0QBLXnLy$N|MlR*WKTIGVmP;t5I;N$Z$Ny$fH%C@bR+)l zDdTZy!ig-8qQ=mS`1Vn_AMr{Z-tb}y0z0pH_7OXlkhrtqFGCed!#^E)8>a$`d&mk^ z>x=rCEsonYI?Ig}jIN>F^b#xh2W(t{KyTN>%RaCSfDJAO=vi_^fP*s9?RH=x;PeQi z+u<^#FYf`^FT>0{W|DRd3GfdTxYnh@A^;Uo;UUC(K~<4~<1fiz=?53n#|@x@b8bbX zf@#p?9sShBmZ1I-<(UJc<+(m*Ss&8#mR_LDK+kOSxTx?5vrtrs1}Qhl=uF{Sv&up| z;^2VE&K+hed_!|6`Ym}Dd#<0$8_V+hM+1%losMk*kHE|5aKFLE_;n-vq75!7@*m2}5V%&7m;$4wz_39O zn8N@cVS+wrMY_Sj1UUT_88!%_mSc>*-3lA}P`<=}nBhiNikkdEE7N%oi%su@gZ+8)cyWUV zWy6NPs`7?!#W)Yv7|-OFjHCE*gX?e$(IQKLLyX_VGeUj$3oNJ?Fhlx$l!t>c_z5zE zBu(L$GAnA^pc!)%LwZQ%5!uk^W#nSUQKYx(H9EEjj2LF%Xorz*F^%X$`i4HtDHc++ zmxtxdV1;?Kf(?C@d`G#*2bssBWkkUX5yY(q4aXI5g`RQ>Hwj4BRgCDvF0`SqmYLHW zZesRiV{t)wgSNtO=uCy~a*eDYsGxz7H3dtUg0|M8f}*@3T$Ya@e>H}s$iJL9WNaJZ z=MLA0G<|stIXBQq!BsNA1Zv7Tz*o=|)&!7$38R~=BL-jcp}_#Di9)tS7mzNOBgE-#6uyV{ zH4o7sHeQx9ybtWb#?uQEP<`H~$3jE0L*L4kF$ZI1>+! zJrnn^!{{h}KRaxeag!;}wzNE<*N;J*CjvJhifjwfcJKGnm~agyt!(<`;0O=C|mC7AKklI zXB@R9yslm4*1VlHu2L0zR%%q`R{Q~@=e06W59N!k)y)^fbK*hPT1~|&&=laqzua11 zU`bn9Zrx3+tjxu#Kz%%CO0Bo#fVbQSJ}`kTUv9lc2k_y-@|cwkS7R%VuRRRmd8>9odETyXx3;|NzcDr7Io^o&;8WJU#}BE}sWRaPMu z5mUL7hH4Pkl_QGBeJs;*rwonF;{=|U*u_@N71WkHAv4&orwvHP$0e2-9?*KlVwv5J z2MAyJ{J@^?yeyqgrjVC0=7hhOa(+EKA|I6LnK^^c|5DwG#WQqsuglHqo?pCfU0x3& ze2&+x*<3k-v$8c?Se>)V(5ga`SlN!kWwvL$J@P8i>DUpR<9R*-2KMs99tg0PAF+6u z^mw#WKRdKfWJki{89Q*Tqnw1e!f#gv-vY+YBFayk*eTg2Fr1w*2So=%)ptL`^B3Pk z%cv}7WuHBzNzd$SINKsHvXFCpuT1@Gh2o)h zEPpFtXJ-4znHjf)pOF{f_X2jvd{c&-`|yoWN&WSPo$(K?M*5aTnO;b{vn(y}+AL_0 z^`0u_wj*Qdo+@SdNGHB_pFA>wOjB~7O`vx-Dg&Q&(i4r!@MoPkcaqIDDea$k5>VEE zK7kH3E5o06(s!GcQ#dzK$$c?_-tSiizQ8%ZGWM zD^iN@uS|cva@&bIWz`95<D;?O*Z9fv-AELzpkxr9A&tnz5eJi0!8}9t{ki?$mKs*_hU}`HN8NlG+xr-hpD1Xi%};ZL$$^0ZtsziV_YxlN zq#X#}uf(&2A)xao!smg`tkb8!mxxpezT|y6y?kq4NBpVq44@ql=}rDfM`OSi=pNu;*Ap!KS?a~*WMQ$N?yjg;6{mu?Su z(yJ<$<~mG^S6^|J;E!h~LnP7B4+ecJ!B1couZzA9wb|b~?CIY+@}{xiN{cYSS-DEwSiEw=ee~m1qOhr)%`C*z?4E5ws-2Sqw(J%M5hL zWxwJ|S4AD*qW*6B*ACP2s~3;!UENPmgm#1cL!!6GD)hE=)+M>zi9|FX_CX}j`x+2B z+fWyM9lWU^fnT|ktCw9l8|WrOWhzNiiFmjgLhru#N>b%`e_eOVN9;EO*=iSV;Z2B|Yw2lA6t*PKJ$ z@^m5xvhA;lbWEjglSo6A)!ElXy$w^LhSbLYR_}^-gYZugcGHR)LU)l0olu$hT~#=q zPhN^S_A}yfCGH~RL+Fe}+cyHOAa&_~Bi{5{voC%(HvV{Yi1^Y!mx=IuIh;uTk`OB; zJf!rOB)at4wd;Xit&8_w5$~-w{oG7_l6VjLJ0PHFLc!yyjk=Nw)Q`Ojnp5E>pxAnT z1*SpBhjQVKAP)i&jtP-?4hQbcAk<|Cp*dS*yd6vxG!!lInT1FqI*8v1*&X+NtRRpj zK&ArR0rDnF@Z8)QOHUE#O|&Zg5HT+jYLfwjC-)?E#GY#OKn+q|a)OTrC5(0?5&yh`Bom2++bMkj*QIj#?~c$7dxc4MECLq4Nc1vL9;6x zG*_qf>s0lb>n}88)DIPDuLw@OP_}yPo+?Ac3y*cz=E58 z6Co=JO&RG65-2UwpBq}@*^BkA%z7#&i7R_YmAM8GPZRzAG0oBN zQovVkc^(0>Pp3Zz+5J-s{)zb{Rg;1Fp4~@y_ zc({J@B)qyNkB4+JhLuIegEM)OO3e_Ahg+dGU59yxGND9jL49SiHcuQWgH@;o-D!}1 zOvtSUAqnAR^4}qt31@&_0~NP`C)-0>lbK5(kPB}I(}NI9MK_VE?s?e`EO`zDvgF(>0qjB zYS)zPlni)t8K}zc@CTydi%>4{G)4jySEWrPm5zsx6-SBi+t8Yhqs3(S791sTl;Wf0 zjo?jJgAx5oXPE7D`hA@%>6c88iu4}@-;i#}WTvLG$;pKdnN;L+>0G!4Is!eZCB%{0 z1~df5qtu?6Mrt$fN})_V2BaCEE6w0dGz0k@aOw+ElX=T-!tI=c3%wkK>#&dfOnS7Wea;LXOM7W-(BZB%9Iza)PwdOpZ6Ist!zx+fNkWpnXVK9SgNk#@9%+G2l` z9NGJ2<9eFN#;{?`!s^L9Pav6%hdY3tz=v2Gb14;`3d9FeBeiEY5tt0*h_om>X;Px= z6$7zn*O2P0E#gSsEQxys<2IbxNCb@W55f38B6m`E*FM2C9-SKxWCy9vXMYcvNkK6D z9y*ux#}7&NTotLxnsxqMY!MiiQNskFEkG_Ipgiz&H}@$^MACm%n4RM8Nog)t*_}YRnT7 z$%9T0&XfhnDyuv5{21MrKy`vIeybOBH-RH|0hsPGK!r&tRFMDil4+3ZiNzKR?)VKb zzJE>cn~~jZ@mMmG%qG&w>@>ljOYM&s6NElRJgK*!tta*n_VQQoBr~yI8nQf)?l3Ls zT|92x>XkAUhnIv_q}vx$cjlwIs`%?b-k{C7%spUVa)~1lOPwAg5~8axW7B;*6>f;f z61iwD6-`Fh(O@inNN3nUq{Yyh!PCi@K+haAttRcgvCIK*;-5RYGwfQ|izk*K5ykCQ zg&{ks&%48qfxjmy`h;>zA+X}-l>cnQ^%zXl(SmIN_v&P&8^YcX4NXF*c z57S6IhTll=v#lqVFVs3nutVss^~FD;Mr`ySQ44*_Bb^730Ar=-$ZfB+92121U8Xq| z){V3>v06vY(?{azYlsN@Jmx|k73kdY<6dLoVFwwXf1(9^v1p5)u1^MH(O0N1m2sON z$*x(wtZ&g(D}S&o*zU0=Kn}M{at8#h@pfwm$l3PF+Bk9hGJi_Cgw2iC-5`IG^ZMf5 zfR4=!Ugyis3>xZjB%4yQJN|%6_c&CijZWze>Rg9=lw5dOo!#^kkO>)6Smq=>Cg;Q& zc`~V#H--!EI7w}Oiy-*YTLkLMY!Qs^w9rzMdAVICA@Jo$NQlIuG$cx=+uga>C&>7w z>Py~OLq**r`4=rR5-)9Y7=cJDCW-sn(LijYou4$_JHo|uho<`VluBtZR8_HaSmm{1w9TzU) zo4)NlTro~WXQ=8SgvRq#)q)=vTW6_i6ydoSsOlVqKe!MbMfl`wbQIxR-$z3TznFtR z(jYX?Q&kaPhte0RsvqI`^Hp^k!W9TtAiN#nEeL;ua67`_C91j?KTW`;I5>jv*-LSS z_$vA-!Ysn^3(-@AS0Efh`1i|D5TSFCs_sE}+vO;S@Ogwb{D?4q3Cbb74&i)+*(*>E z;f)CINBH$JltY+VuBwL-{^CkiooB-@6o1A+6k%i~3L;#Dupi<6RVawi`~!>-!ZlZ+ zAj10)8t_ZTx;3gALij4eX$YHoRCOXoaDGpr(@VzCKY9wQyd)UkU5zWn{!qHuH!Jp9 z+Eas$+t{DUc0+s(AQLK6OyTJ#3~tHk3t{c4k{ z9z!@nrynyir@x8QTQGlrdj@^$2>mF|NBydL5$5NJyl}oy)Su_SQeWAksvois5F-x! zos=9%;{3SQ?^FPDkiHBH<2ww%TBP3_QC0j=VwC>I;sP-@??i@eF;#u&3=i)wc?gdq zJ%SbXe=7iZ6Y04O768KX0=^hoz)_rcV`X~JaKNfLfKdeY<=Rr1R7Gm?NgN7TwfM8( zC|#bB-I#;(6pI<(J<4?2Hxli&NWbj=q2G!0;qRuO1CJwvXQHYuKf}Pn5d&``eF4(L zXDqgvFSY??jv{@b$~@jNdz>u#G|U%)-kIsxu3p)S?rn zSp{ny;@`~T^)gk&+B<<=D~PfqyRF5C*MlGF-yj`7xR&X2dHlHF_`MZQ&no=14t+lV z!oqMJ+E{vF;Q+($n_W2NBh}Kq$n(hTLQ_4?{q_5W>GjC-_V){`7(Q+;lc|=%n62G& z3p?1{UGoYD7|+A=3a1cn`Wm|<|5R5hr)$dgpGGIZkHl1XbtLOwK>C!W*f&SZ(^*^rlRt#?8<0MS(I*vq zSdnhSDbWNn97cxXXEK~BY-l8|Cwd#nG2;`Lw-JMW`kvx{> diff --git a/norch/csrc/cuda.cu b/norch/csrc/cuda.cu index 2c4fcdf..1dd5ed2 100644 --- a/norch/csrc/cuda.cu +++ b/norch/csrc/cuda.cu @@ -76,38 +76,39 @@ __global__ void sum_tensor_cuda_kernel(float* data, float* result_data, int size } } -__global__ void aux_sum_block_kernel(float* result_data, int size) { - extern __shared__ float sdata[]; +__global__ void sum_tensor_cuda_kernel(float* data, float* result_data) { - unsigned int tid = threadIdx.x; + __shared__ int sdata[THREADS_PER_BLOCK]; + + // each thread loads one element from global to shared mem + // note use of 1D thread indices (only) in this kernel + int i = blockIdx.x*blockDim.x + threadIdx.x; + + sdata[threadIdx.x] = data[i]; - sdata[tid] = (tid < size) ? result_data[tid] : 0; __syncthreads(); + // do reduction in shared mem + for (int s=1; s < blockDim.x; s *=2) + { + int index = 2 * s * threadIdx.x;; - for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { - if (tid < s) { - sdata[tid] += sdata[tid + s]; + if (index < blockDim.x) + { + sdata[index] += sdata[index + s]; } __syncthreads(); } - if (tid == 0) { - result_data[0] = sdata[0]; - } + // write result for this block to global mem + if (threadIdx.x == 0) + atomicAdd(result_data, sdata[0]); } -__host__ void sum_tensor_cuda(Tensor* tensor, float* result_data, int number_of_blocks) { - sum_tensor_cuda_kernel<<>>(tensor->data, result_data, tensor->size); +__host__ void sum_tensor_cuda(Tensor* tensor, float* result_data) { - int remaining = number_of_blocks; - int levelSize; - while (remaining > 1) { - int threads = (remaining + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; - aux_sum_block_kernel<<<1, threads, threads * sizeof(float)>>>(result_data, remaining); - remaining = (remaining + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; - levelSize = remaining; - } + number_of_blocks = tensor->size / THREADS_PER_BLOCK; + sum_tensor_cuda_kernel<<>>(tensor->data, result_data, tensor->size); cudaError_t error = cudaGetLastError(); if (error != cudaSuccess) { diff --git a/norch/csrc/tensor.cpp b/norch/csrc/tensor.cpp index edccaba..22391ee 100644 --- a/norch/csrc/tensor.cpp +++ b/norch/csrc/tensor.cpp @@ -171,9 +171,9 @@ extern "C" { if (strcmp(tensor->device, "cuda") == 0) { float* result_data; - int number_of_blocks = (tensor->size + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; - cudaMalloc((void**)&result_data, number_of_blocks * sizeof(float)); - sum_tensor_cuda(tensor, result_data, number_of_blocks); + cudaMalloc((void**)&result_data, tensor->size * sizeof(float)); + result_data[0] = 0; + sum_tensor_cuda(tensor, result_data); return create_tensor(result_data, shape, ndim, device); } else {