From 76cce8655740e4ccbcf706ca55711d11395ab930 Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Tue, 22 Sep 2026 13:24:23 -0400 Subject: [PATCH 1/3] Now testing on OSC! New API for Python has some errors particularly with activations, but this should be fully resolved by the end of this branch. --- .gitignore | 1 + CMakePresets.json | 2 +- .../__pycache__/__init__.cpython-312.pyc | Bin 2297 -> 2333 bytes .../__pycache__/__init__.cpython-314.pyc | Bin 2763 -> 0 bytes .../__pycache__/base.cpython-312.pyc | Bin 3566 -> 3602 bytes .../__pycache__/base.cpython-314.pyc | Bin 5945 -> 0 bytes .../__pycache__/rich_reporter.cpython-312.pyc | Bin 16541 -> 16571 bytes .../__pycache__/rich_reporter.cpython-314.pyc | Bin 19742 -> 0 bytes include/deepity/utils/ActivationType.h | 35 +++++++++++++++++ include/deepity/utils/Activations.h | 18 +-------- mnist.py | 37 +++++++++++------- pgo_workload.py | 15 ------- pydeepity/DKPPCN.py | 18 +++++++-- run.slurm | 37 ++++++++++++++++++ slurm-7501201.err | 11 ++++++ slurm-7501201.out | 7 ++++ 16 files changed, 129 insertions(+), 52 deletions(-) delete mode 100644 deepity_build/reporting/__pycache__/__init__.cpython-314.pyc delete mode 100644 deepity_build/reporting/__pycache__/base.cpython-314.pyc delete mode 100644 deepity_build/reporting/__pycache__/rich_reporter.cpython-314.pyc create mode 100644 include/deepity/utils/ActivationType.h create mode 100644 run.slurm create mode 100644 slurm-7501201.err create mode 100644 slurm-7501201.out diff --git a/.gitignore b/.gitignore index dd1a8c9..255aae8 100644 --- a/.gitignore +++ b/.gitignore @@ -32,3 +32,4 @@ _skbuild/ __pycache__/ *.pyc logs/ +scripts/ \ No newline at end of file diff --git a/CMakePresets.json b/CMakePresets.json index c9f7e6c..5cedd29 100644 --- a/CMakePresets.json +++ b/CMakePresets.json @@ -51,7 +51,7 @@ "CMAKE_BUILD_TYPE": "Release", "Python_EXECUTABLE": "${sourceDir}/.venv/bin/python3.12", "DEEPITY_ENABLE_CUDA": "ON", - "CMAKE_CUDA_ARCHITECTURES": "86", + "CMAKE_CUDA_ARCHITECTURES": "80", "CMAKE_CUDA_STANDARD": "17" } } diff --git a/deepity_build/reporting/__pycache__/__init__.cpython-312.pyc b/deepity_build/reporting/__pycache__/__init__.cpython-312.pyc index 705866f280aea4a156fd203c930958c18ddf3a3b..640c1a891d58f41aba34cc54c20e9d76459c47a4 100644 GIT binary patch delta 74 zcmew7fK6rQ!$_GUMBkeU#F0-45Dhy~c8V5W+s zl($bRgy<$~23)xQGHu$<6?>VLb-iJ`{zQoezSpMgD&thx-j$GwfQmY!@d=o9eZlyW z?XW}CkGBPsyH1@3TyJ`yw;dZGtSVE_5tWyA*&stVbDE_?+|r2OdiuUjI&Q$aCfLn z^WYs-24;ouEDw*{evQIFxDFNVCZkTBIae4DyQf3Os7=S3K^VGKudTQIYS8j)W8FK; zdP#{4+xG)h5%}SG9G_^pUTr4um)vC-=1#j}-d_qp)Dr(W;C_NX{8ad*VYo>_SRsoL z=XX9Qqy}>%6C+hdyDz|yt^6dtT(xKo{7985*exz1|!26lN zec7deH*68&O=Hmqm3q)%6&{3ava-N~%Mg&TQe&*?igwalRW!a5hXq_(u`Jhjg=LkS z?Wo_fx^HEc#Rov^ApF8T(5#RbL{+ArVzU}pnQM`6~=o$i>7F z@brt_v>;@vD;@AqV2WVNzMZv z^~L(Zgh*`(*Ngg=T)$?8P3Dvojt`a)80Q64FzAL>f}Sl?`7T(J(VPKcg*-NfzV&{Z zSvMx`8WX?tuTP%1GkIcta`x`z?1RFt&hDAD!fZ#IeP9gUoc?lp-6-BMimT_=j3d7d zc8ud4^*G-FdbZ&l0}7Z{Hb>+Qug7^sI+$bVIL3z?*g!l1&m4mjgIT^rB2YGaap4VU zVi*n~@P=-wvWpFYwXLsM@S*mlnSz@83wJGeZvsimRvA1`c+# zga14*hMo~6GZYOBZlt0AH^lef(DURQVx5MIdcR`*Ff(S>%$Zt6 z1rYa1A4gfJ#89U=C?HXK*%chU6ZI`KpK)6RT!JnR1!ms=hGGmhg|Q_s$PVSmQ*3S> z+`#VkY1#drmfatkT#Ms__a`6T1LMr*S0MD(hgb8w23{ONDwMR@Opo^Fd@9M>QL>nR zlYF8qrlz{L6!A=YUY*M{)oo|go-^vJXG)ptzv@udA^%#AkbGEWkh@!feMx4>Dp)C7 zkhOB7Qx9AR`SmPo3My%A0^AD%EZaVm4C;6;)YuXd8dB}pwL0=?!jCaiU8}dppg07) zt$XfOhAagv+a>2@E3kE)HA3dW21!u;CTL2ir}Zv4_0J!^pa&3@2e#F#GW{5e4bmub zBrv?~2)KsU?66)9Y+fq^spJ=s-?!AgkLrH>x<}pl3m?5Kfn|CMm=={*;XekN+Vl4S zN@KGn%B8GKi6}4gBzXo$Bt@AfR6-FI*#0ispsw+yKe1Vb=Qo^jyfsfTT;194pbc10nG8zK$NqrTHsh#q+8Zf3+Ne% zGX-l=VXJ1Bv1IchJSyLVAdNIgMPObaaG8&4)ec#F2Nr7SFOIk24d7tViZ!CA=xE0p zL9OMnX)o)G$w?0-o1JkJ7YKQBP#d|rTL{d=4k Krg9@Eb^ZtR(wlVv diff --git a/deepity_build/reporting/__pycache__/base.cpython-312.pyc b/deepity_build/reporting/__pycache__/base.cpython-312.pyc index 92f922013b036e3b5eb784cbf23d10118b80428a..77fa6ec81593d88d165d52e9ce529c81962c8dd6 100644 GIT binary patch delta 74 zcmaDSJxPZ9G%qg~0}#YaUA2+hk4Y;~zqB~Ds8~P1G1$P^)Ih(eG$}DZCpFc`z`#P^ cD=|4cz9_#q)zIABbaOmY1UKWa$vb$>05wMz>;M1& delta 38 scmbOv^G=%kG%qg~0}xanU$&9kkBLoJzbHSyWOFl91UKWY$#;0o0NRNQ^Z)<= diff --git a/deepity_build/reporting/__pycache__/base.cpython-314.pyc b/deepity_build/reporting/__pycache__/base.cpython-314.pyc deleted file mode 100644 index 5859bf8c7fc050cb18feaa0eb8eb77d6bbacfa4c..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 5945 zcmcIo&2Jn<7O$S}`SAE7{*0Xuh!cm{V;g25jxi=OgdmZ4g%x@cTEQm0?QuIEc0NpX z4Vy%J@ByK{?KQ^{r}&mXfj_}>3ba5%;()k03}UbRUUhX(w;gPc!iw#^-bcOn zx~gaTdb13zaQVNV{*q+uM`D5oo|@Kr7j%o=VWxJI-Oz+~j{1q4i5t4mZy3UmynZuz zBPCMj*a$n%Ok~tJRvYMX6SGa|uts_Re*$&t7!} zw~M0Uh>hABW*N6^u~GC_@Tpo(Ykj4WO1GKNOePX{(n7y&m>Ps8Mp?OkPCX-%Ky{!{ zCxWJcCV{46Xc}l5XeNecfOY}R#?UUH-9UR{XclNM(7qVj4fF)i{utT=bO7jJ4DAIv z1avrt_5mFMdNPKd06GeEEQa<2Jq7f13>^S^2I$!sItcUypfASIA)x1gj>phppyz>J zh@m4uCxBjzp(o9k#He`*8$Rh}OlPykg~N;3F-AXX3H-2^TI(!2TkIotC$aDaV{5?9 zvJbQk&D6|9Uc(Cd&otLdaYxj7^+9$rSZ{Kh0x7j#;}u)DwD1M*YJQ_uaq_(8I)(f# zUi;K33OB#zIGZJLKlry+>!tEqo=d$-)w}ss+jU;qyzli`mI_P9vb>Q%&?`~63v%3E zsq_V2VrGxW&GmoaO`Q9nYKFI;V8!-5{Q^QmZ;rqnyUcB;ScrA%23U^@EVc zp~3S_WB$>HAifR;#)9??gZU5!eE|j^LEpZXxTAg+?q6JA=_{H`NhPcK!4^b0?Y z#}9Q2`NtiUf}R-Umd0OqV!olM4cfuJNl4%D3CwosslXS#GEObGUa8ppzIQToNDQa& z;pd4QKP>nudY*RrphzHVTS4g0HjUZF97tO%hGjd@T0V?9(a0YtmL0olai>_Tu9xoC zx#OJw-ZAnKf0;-q z_Wxwy7=(n`J<~L18W%y@Vx_4domtx-?xj_q=q>qfIdWoetHfTQ6{Cqg7;#Un zvRQV7Y(nu7mcZ|iy!{~KrHGgxjxqZH6 z#h@CD?lCzSxI{WAvtL(Yk8k!{M=^UkYW7^ym}^|#G3FjouT!(H?=`!}H(OFpbs{HhWeHJQ%W^3+FBdC;lPOq!{@PM1*s%pMK6*ch9G7+fNcay-;x>p& zkA{cjUg5aHwJIgo#q}j}FpS}|^qh*7_ml=v$NkBru|T@MMlB{H%KR@xLUAO9T4%ZcKAt!<_b+2oSn={rBmWTln{TY4*J%>mqNSpxu8Zul zWTk6g{X*s0(Y~=KAntB;Fd}_CgdRdn3d)kejzmpW5K|3&Za{dPb{yApze-XI{ zvghMK+N~Fhj_Z2;;pwY7Sg`uN)O zS=}zHkrS4+UgEBxJgn9%i&M_=lSIxDp*Js|AtDPMePHC|9{w(on?yb!@)4206Dbn8 zOXQzKYD8Qj4?w&Ti!OqqRJ3d%cxkncn1NI}ghY*voNxfwM!Bgd8$w-aK=9p&_r^mKc*$SLmR4ti;N zt(QxyoX!c-m3w^bWv==mvBbZ?6pB)pE?m31rfJ`^3qP@m@7dU|sYPw*Kkjb~e^P

&% gH6w7c*G1uw$>KJTSuW@VU)lW6CW;woL8OBV07e!@G5`Po delta 145 zcmdnp$T+u=k^3|+FBbz4R3Bfqk^3DRo0)!5etyYjW%gSfjJG!b<2U35QUdS^H1J#3Q0_nT}{q{PvKr1h|9nUXEZk|bJx#ENnSp?YG5k2p1{V6ZU>b+`+=O6V85D+{ME76Yl=v zcyYfJmk<{f$BB~u(s*gBP%Uf{6lblVxJF9%Sa_fDGR3`F@D_hXQLm^K`y-L5xIZ46 zio_7L1%ls81QDf2e2aZfd#zU77bsP`HPejJ2yhU31;WH;rzT_KJP90Z{ zo<8kUj%sD&Q|hEY?h7PT)*>o6m0)yAjR(~+-aq$?_(NRj5@`ZPgt({(@uE>FZc)Uz zRVl)tEww_}6SpZ=gl#NrSL_HoSlB@?7I!Kx_%6i_->nqGFIFU!DPd(Kr4->Z7A{dd z2$!>PsZxP(B@34+RR~wJut%vuxR!;>l{$p$S-3)JK)8{GE0rdMn_0L@X+gM^g{u`A z;Y}=DqqHHsnT2bWc7!`vxK7!E@KzSCS2_{)vT%d44dLx9+^Fn8cqa=tDP0J6vv9N0 z6K_$T!P4#0BqcO9Va)4v@99wFQaC72CgT2!cv7Z(GW14J4)|je7pMGcK*qYr@rj@u zjKo7~P`*fO+BJr`7wDgy6aOF)aitT1tHLk{uJ`u^A%L_d;XC4lsEA6DVo|J$O|dJO z{cbD9;`)kKQnQaS`cUl@v0MtpRR@J!mx7TXC}B#~EU~z%+0|e?p++?8z*Ho7hp5la z2_{CPXw!~%PfSe)yVa>!u%~-aoq9bu7LRoYg28Age#LnD*x-aJ z;y|oACsvaatIdhk8L?sOCSi{rK@0U}x_y%{?8wS<61&=C^lZz^ahW-;OwJyim7?7S zvt41(=*~#Ob8uf)L=P(NQ4#A@yiaG2Xr;+0rAUZ1XJVlC5|yZ(xFGmpY$_b|+BDZ; zzj|aUoS2MgPLS*+H5iL&C8wjI2uNSga1RCFitFL>q0nUTXxJZ(1;JIR-gNQlxL=Kf zl)JjRrYp_`LX#tJ`a_g(uA3qn+bNU&x5B|lf=2q~)gMkbpSw5}4#*e7iQvfFaewS` zS0ET0Q$tbW_51>t7rTN9h&1T@9<8MW{cxE2y3pf zsmW+K7!L*pye*nUo2(1FoZeccInqj4lF^GwGo`kW*GisD-X`*h6jh=?l?YEO!Mk1x z#$#+7$2E5(G3kr?W0(yd9`}dB2-~6){#a1+aK@5QgT5FJsYFbZ*!DMK#qnS)ZbV9W z)+AN{FRRi%)!eK=I1~xSv|3-LwK!J67YN0oVgD7alJgW_PGW_RRn1AO^s&jwE$Gmv zMROXHp*j75fR7eMbMrB=u+0|+SN6I!Yb+QZ&zl%Fsq4;SjM%~!IGH(>cADn&`9hIU z96YW%kH_(3#@~L1l2?T( z^pw+WTj|+q#D7}ba{CGjO84vJ`%?Wor|wHN$=YWZrG59M>UUmYv15zU@%vKqqSU%l zY_E5tbjia~p|SP-o$u{jZs@!&8xyULW!;F z5qYx**&{P1IF%OhX#jH9a*sd^W z9xZFjjDqzQ8Vdn{efk)|9O<-Y4yDbb$w&jkH4mXoTGl88%)8ZKs8+%&$GjG;LJf|C zg--auq$NL+IO5<~() zRi*XOTs|LaqDh~(NbRF4R4+ypvPI?#+E3+bvezD9L9fKtcoYR_C$&9rN;jnI(z3H5 zNXIOr zTRsa|aW=kOF9_D(M6{Kg%wmIdaOJtP6~h)oE@>4MPz3$o5eKkHY}vr!f^QLr1EtWa zY0ltVW4`#6Xi&2tq!q<4#%Pqp8Tl#>P^o-&$lG8q+=2Y7Lb|x@#{TR3|E72oE!4}` zUtX44Q&Q`kb6)*S+V#jHNNt(5Kn>=I1+yvmZTPuc*Ti~7Y)hKhPOpi_Al1%uQOmHU z(8f_Ly{{R_Td}@cwN5EpRw?_d5c{Im;i5vDPH9z~2A+%BhHZl;0?Z63Yh(WU^NCyL zIR!Et2&d4@5jE#I37Bn{g%H}Hi9j<&^xTql+B0!!evLgk!M6%Sw!uPUELF;et!Uj- zctwy?&hvLU3YSnSvev8;YgPr#fa#R#bw+F=TeGNOfORgPDK%)TcEj=4DfQ@qPQyGm zNWd-6;cXh$DJ|y}MN|@YLDDGLcH7$0d82Q{=$o=IP4}mtZYiNA}cKhO2@&m zi5(nwpsyJ8M}lDr43Y4w6$8`7)zHNRfU9N$8V^F;9QViKn*B;J9G-epv#7xU$~&UK zXn@$$9UMrHUGYaoUgDyrd^B5!F=u*Wf-sYOj`#89pWl}0 z)Q+6%<_VgWW*f>S>Ee`%(;LwK3Iq;1bnXOHQ15u&`^$e2q z1U4t?StKy5jV0$ZtXU)e$)IMvq=q1L5L6%I0-S@wDrm$&?Ko(-)M9-dD~G`)0)8gT zX!eVMBV!Y!`Iwvx#i2I~sqrgbI|mmu2;T7}HXJ|Jm%j?cD_W@-VJ8meIhJ?$CW3yv zuvV3eKzQmYl#acFw(Ob*=A9 z?@9CarMj-`{R>AQS_D_+L#I$)GvmJHo_qGDd(pG)nl0@tS$5W?oOOi!pCJ6eS#qP~ zdP!Q^oi6Q|3w(6>gUicX`chl^k{t)`o=r+e)1}_|_K&;nbS1sd-K|YZhtj3(b3-4U z`{3Mi$KF)O-emi;3zw48LBJpgA*F(=l(*NBF0H!x-J2)p#JQ~>?f78FT+95M3)2hX zdz+HdOAqZrRpU(Gt-iV5rONF%n#&J~52|ZdUAB@^z^Iaz^v-8KK6mHb^3DVJOYS_7 zlDx?yr@yXrqmMeXk0$D4Xy&zBug&;w`I1!~bKhOCe&YPG^OLe4mnFMjq|xr)_et-M zdp|k&$-#w#sj}|m@r$b#u_PdVZL_-Ek0ilUee=wG&E1-%vKN!i7uNy=T2lrP4zraq zl|flr-6#QsqQL_AT#>WOSVtwrwlBjREssO-=*co@Dv2;Q-6?i3YExZg zqzXb7(=A*N>YqVVmGl(#oTz82nF`h;CT}4KkWs~$>K1hfRn%efM&M~yV#?qUHpUt> z+nb?4d_r>_VH!v%3N#yI(l`*H8i-BTb2iP=WGGsoL(usZT7vR~X*@!4Fqw$JRP9&I;?mMsT`a_WH&asuEfaOD{6+T6S3d=B}v$8KJljW4(Mq*{BHoX_N> zb*Eap(aulHs&1Z1wjKEN*izZ)r1SJz;x^ImG(>^*(JzRenJl8*e8Ef>qzx@hIts_I zyP4QPL;y`8*Df1G&Qda(KWjl;<~B5ghA;t?Tf}e!tN}_%1~;rTic*Hkj_JZn0lfLV zXdsqL*x*gTB}!S;&ibHH7et}Mg~#d{wi8a=WiiLBZ)c@KseFR5RJ~)SE2UX&hIo@TBj#sWqHn^g$-NZoYOk7!aVoH zEwy8GyYea{y5R^qP|pxwv#M`cZwu-@n^i#%Zdh+C>h(Uh9`PQ7n{_H)Wt+0y;6Zu( z%w}*bWrwm;*`;(D>b1>)3E+j@N)Pr#76Qo|{IUQk^u#%lYD>jZol5VQ9 zBJz<}{g;DeWGH5W4M*hAZN7!K;S$0)mk_9&Je7z?6LC2dlP5xfKrkY|IRWv6nGu;% zn{w1dB!X^aepjWc7kwa`^CX)nhPTm2tkbH^ekHDZ2 z3TL;?veX|?qYgp~Xi5{^kO;}@qFY3`=@gmze4wv2T};B<00Ya&mw)mfWrjNRu^Z`{ zLl0?=(}}S$*e5uYv7jIIO>{r)=J=^=WY$SkBTQ14%#8Q)=`%-;z;=J^%*kmtqX{O4L8@Vj zPl#DW8k*zv1SV(v$~2dGWiG4f!m|FZJWc&Al!rR~cktkVibS{R>g*XcXAsHM=g6Tk zNVTx=>bQ(!up%`;Z4gfNHmk%})!!$t05a3n!Eq43`Ug~os7oszSEnX@lYaFwcF8l8 zOq*6E4x>Sw3q{mCKu79oHB24}ZR!Mh5%LJTs*~gqOICqa1eF+$x|cj+AzlZU``8E> zY-A9S219r?a7A;BB~&P+<1q))a;!KmpI?9feC@N5{zs@18v{%rV7P?~fTV`zO3c~j zrQ$93E9>ris*=^)7Y;0W4k36`x$R7Qs+K)+$|FB46047iD;7(6#frmGReJ5EM;@W7 z9{K5-qv__3uN~H^(uabTNmNjXLNz+uuvEM`UDL2!)0wL2T(0R&*7VMwymp**RDClr zGkI%rZivfucWrks-Tke5H@w#-+3+<2#P z+1r=$_APr4rM!ohq{BvTjgkA@viCX2w;xp9O)Np=)f;X7i_Psr1pchO?IN=hHqK%`Ap??~=>CHY<7QrU&1^TJx-ZwK`OZn#$7Kp+?R zvqMf8gg}t*jPn>rB@uj{%gQYv2^+b_(FjRUSN)K1&SRPB z5mAX8p?6G=xI${x#~B&bS5RO>t4s=ovaB+FWRB^ZZD)^6=P1M+%{*HhV)4{@AGHB&7D@xaZ4B;$ev+-I6h5$euy7dkx03;OL3ra)-ewf{80q%&5M=~2KH`c&_ zt(VWY*yzb@G6E?B&Fuhr04TYH^w6AG%T*P>+`hEHYPq|{@js=>^UFf&xJTo=Sz#F zG#yz&*wX$_1hTdrb$U$?Q*9XV?P73*+#x2-(&O z%M9*_8Hm@w{~#6zCf3QdxI-Apv}<5WM^#WFGx!1B{K#~%4k;+2o<~83`v$P{gA}3j z6<5hf%^zi!U-cML3F%QZr zZ@hi|?RS2#T(&(~wtcSnnl<0-n-hO;?_AaV^xclTwfBmXJ5MIhU07=Lf9AZH*~12H zTG0h#ATsM~{fl?qMVLMtqf5=9-=fhhuysELS^gdRh>*M|`Dj4wu~U;I`j7 zK0;k7fs!VRbFC0|x=NSVSTcDat8Ajx(=#MQ_FR=F5xkj!rw@6!4-fS(;Bh{!^8@_? zPb@JB*mgy0$XOjGbNKW;_8w%$1ZQY&X9{eI6kUH@wswNQWe)u%dR-ywAWfl#rV!rT z9`?v|j(q0y&Rr{q5xolgHttR6Sx{4mh#7Q1bNQHUjqOA0Z&Cjc)wy_Mh)yuGFe8er z(tQaga>0tE=MLaz0O4V-ucYZ$e*`a^4VSap;KgQbttDsb*W}Y2=wy^OZVC0%Ri9?s z1izVe_&L4Gza!7Q_=5;@zoto-(Rj9P%a|^RjhX3I|2;|-m||mL=&~BFLkIGwI6t3W zFMo;;|AwddpJ)sQshiw_ff1cen8b^c=B(59&nNW_8HEVqE@H;PY}=rJY=UfuT)3ob zA%q6hf>fM!Rd9?|)0%UM*qmXWo%M2Q`T31z){1_gm`mT$xwO-`2d~-2!&CmJVbl-N zpUJ4L#AV>k?}dvPP&YSBqEE?x9PH1r#-5Ac~JptpnV@d2l zqx_WS%fXDd`X$jEC4g?3o_6TQ=K*jn!ugrSyQ3rZ|4MI3bfaJ2(JG8R6@*Qe-#^t_ z*IP!t0GldO5L^`=lH*p%PlgQtnclgD4fh*Cv>gL7B!3Qwke>5Eg!FvO>e4~Vz`nVHNYEV)7%`*4RQRr< z-oE_v8Lt1!)w0RQ*nQgpXaVn_l(8@bI&&5#Yj^24bl2NmJe$2GXB{{amYt0$XJfLd z@9y|#&Z8NIZM=ZaI=(F->X#V3c|;qVXBHD{?i!rTx*<>LbbmgxsqL0HD~XyRkk5MWs~t_S1xa8n<# z)_NTm@cWLlL+|_E^DP{@SDqX^m!$iT=O2R`XvUwK^;+r5_tS38U(xW5tyyLpz}6&x z7F&~^bD$u4uHBj>CNUCv0K&Q|kj%o}AW?RnBjbRn->wM-&NHVg%fQ_#UHI0g?n;IB$|2#K{t)X689+j?o;eb;CMEMp@%YX2gQ6 zo955QzODT{DP@_%J9&(7<~eLLm>u5V+;C?vgsTZ8n=+*{+fo0Pa&a3C1{=+S4>`Cj zl)Lu~9deqL#aDD_$jKUt1_SD!B9(1u*b?~n4P29)eTPWDwM3M!)#evk3hwZk^VPK{##1o`^&ZV0 zE&1av@_Vp%S4fs1!Ol1}<`NuoRb?F>*?U?a$FBo&ssEh@tV1WptUsjJ_$zohM{_A( zT@x?mDm8L5=^gb?-qy!T_1_*=QvU}Hd83)Hp;2XKzL(~@cWyV$eDC}o7F)}6&6yuZ z7aJbikCm*k^Brg;&nZPYP>QsI*40)N%BZRcRB^nMrLvlq4%01BJtIL}i^}mt18!OQ z)gg8evt=dprTzq7&f4l z_RpOA)@~~O7~4=y%Dw?*YElPRN8wt2&_|al*@=r!F!1egY}y^hUEFi@9r4IhsHmEV z7zn!XZFm8Pu+?@!oEsx@xF9ZUk|g4Za`NgAg6dDH=4iOtj$)2EdJeQO%h9b~=$&;a2jG?-WO0=ArmE}al{|4|n z-5^a6xMrY!Dgj~}z%}4wAz0zn>quvxLE*`q75;9rP|$hPhPz8wg(2Y>$TBJxk`Lb! zuHtj^S8!;$gckNe#$HQE2H4CW1DYHTs^j?B^7fpL9Xl1>}i zRg-8l(-ZF}OC9qDQa0p|R-;4>31sDC+NJeRfWo4YO z2L+jhUOh@7w!l9^C|g^pFQlPH%aum>Lh5-mIWaCQr>pUWBNnjnIli7)$g=BY%g*MM zvzb0nTwiNR6T=ovMu4R5bxcQAQqQZ)D74{m*Gc0tjyoIacHqr|zu_?#V6nBw9T7AB zT`TYMafpKkNg-}HV|k+Duz6#V-vh*QWy2P|LLu;-5l)zNTHs(bdI2Xh)%#X;9&Pg5 z6?T6#if`}O@9?1Lm%#^5>;cpQU+QP)bB zv)K=SI63eu&FE|CL-}cm>AB_a1CL;18=ui2NBwVT9Ut-1rl*om0SKQba^R-NZv+pq zOL!0J+8^5R1Z`F~JakaVi5o)?T@-Q)?$Sra2=TY2O=&nC;<{D{A>eH@Y*Hd^N+Cdm zOSYt4G3G0~jH&NHzR3o%+Ay^o+1i=&=DDeewE}$uyW6rBE=H|c?N|*EiFQm$2J%Vn zH`u1G#CTa;p4Ee3)l7cD%|2!u$pMITv=ju0-z#5xnP61N1Bl{gfdC>HJ9A)34X|* zw>`7qU2J1pw)G!9m8r_1`O#17@pg-qLyMlFYqrlf?^@WqxS639`GF2zat=9f6~BYX zUXw1{WNssJ#S_yQ?57N(3vPSaMAgZH!k9cw6abCP=iy z&rG;|K77xY2vZn8vh%%!Kdk4J`F!IcT*;vGJTirBhtC(78uR(o1JqPMd1RDQ zk8W(Si$NUb-J$5m2;b4m0wUi_C&@PTW5A{@G?5-D>UAm*nz@$*dh-4Oin5 + +/** + * @file ActivationType.h + * @brief Just the ActivationType enum, deliberately split out of + * Activations.h -- zero dependency on //. + * + * Why this exists: IComputeBackend.h's methods take ActivationType as a + * plain enum parameter and never touch the actual CPU SIMD function + * implementations (ActivationFn/DerivativeFn/To_Fn/etc. all live in the + * full Activations.h) -- so IComputeBackend.h, and everything that + * includes it (crucially CUDABackend.h/.cu, compiled by nvcc), never + * needed the SIMD-intrinsics half of Activations.h at all. +*/ + +namespace Deep +{ + enum class ActivationType : uint8_t + { + RELU, + dRELU, + GELU, + dGELU, + SIGMOID, + dSIGMOID, + eSIGMOID, + d_eSIGMOID, + TANH, + dTANH, + LINEAR, + dLINEAR, + NONE + }; +} \ No newline at end of file diff --git a/include/deepity/utils/Activations.h b/include/deepity/utils/Activations.h index 70760a1..e5eedca 100644 --- a/include/deepity/utils/Activations.h +++ b/include/deepity/utils/Activations.h @@ -6,6 +6,7 @@ #include #include #include +#include #if defined(_MSC_VER) #define RESTRICT __restrict @@ -46,23 +47,6 @@ namespace Deep // afford to mutate src in place -- src stays untouched throughout. using DerivativeFn2 = void (*)(float *RESTRICT, const float *RESTRICT, size_t); - enum class ActivationType : uint8_t - { - RELU, - dRELU, - GELU, - dGELU, - SIGMOID, - dSIGMOID, - eSIGMOID, - d_eSIGMOID, - TANH, - dTANH, - LINEAR, - dLINEAR, - NONE - }; - static inline void relu(float *, size_t) noexcept; static inline void gelu(float *, size_t) noexcept; static inline void sigmoid(float *, size_t) noexcept; diff --git a/mnist.py b/mnist.py index b7f47e4..5e8fc66 100644 --- a/mnist.py +++ b/mnist.py @@ -2,8 +2,10 @@ import os import sys from pydeepity import DKPPCN +from pydeepity.layer import Linear, Sigmoid from time import perf_counter + def load_full_mnist(): import gzip import urllib.request @@ -39,6 +41,7 @@ def load_full_mnist(): Y_train[np.arange(y_train_labels.shape[0]), y_train_labels] = 1.0 - eps return X_train, Y_train, X_test, y_test_labels + def main() -> None: SEED = int(sys.argv[1]) if len(sys.argv) > 1 else 7 EPOCHS = int(sys.argv[2]) if len(sys.argv) > 2 else 50 @@ -47,23 +50,29 @@ def main() -> None: X_train, Y_train, X_test, y_test_labels = load_full_mnist() BATCH_SIZE = 250 - TERMINAL_SIZE = 10 LR = 0.00373 IR = 0.15 FL = 1e-3 - LMBDA = 1e-4 # The crucial Kolen-Pollack alignment decay + LMBDA = 1e-4 DECAY_RATE = 0.94 - print(f"\nBuilding network (784->512->512->10), seed={SEED}...") - net = DKPPCN(batch_size=BATCH_SIZE, device="gpu") - net.add_layer(784, 512, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="linear") -# net.add_layer(512, 512, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="sigmoid") - net.add_layer(512, TERMINAL_SIZE, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="sigmoid") - net.add_layer(TERMINAL_SIZE, 0, TERMINAL_SIZE, lr=LR, ir=IR, fl=FL, lmbda=LMBDA, act="linear") - net.set_optimizer("ADAM") - net.set_psi_optimizer("ADAM") - net.compile() - net.randomize_weights() + print(f"\nBuilding network (784->512->10), seed={SEED}...") + + net = DKPPCN( + Linear(784, 512), + Linear(512, 10), + Sigmoid(), + batch_size=BATCH_SIZE, + device="gpu", + ) + net.configure( + learning_rate=LR, + inference_rate=IR, + feedback_rate=FL, + lmbda=LMBDA, + optimizer="ADAM", + psi_optimizer="ADAM", + ) print(f"\n*** FULL DKP-PC RUN ***") print(f"Training DKPPCN: {EPOCHS} epochs, inference_steps={INFERENCE_STEPS}, ") @@ -72,12 +81,11 @@ def main() -> None: rng = np.random.default_rng(SEED) n_batches = len(X_train) // BATCH_SIZE start_time = perf_counter() - epoch_accs = [] for epoch in range(EPOCHS): current_lr = LR * (DECAY_RATE ** epoch) net.set_learning_rate(current_lr) - + current_fl = FL * (DECAY_RATE ** epoch) net.set_feedback_rate(current_fl) @@ -107,7 +115,6 @@ def main() -> None: total += BATCH_SIZE epoch_acc = 100.0 * correct / total - epoch_accs.append(epoch_acc) avg_energy = epoch_energy / n_batches elapsed = perf_counter() - start_time print(f"Epoch {epoch+1}/{EPOCHS} | Time: {elapsed:.1f}s | Acc: {epoch_acc:.2f}% | Avg energy: {avg_energy:.4f}") diff --git a/pgo_workload.py b/pgo_workload.py index 3943d0a..21e2a4d 100644 --- a/pgo_workload.py +++ b/pgo_workload.py @@ -1,18 +1,3 @@ -""" -Short, dedicated workload for PGO profile collection -- run by -deepity_build/cli.py's --pgo pass, not meant to be invoked directly for -training. Deliberately NOT the full mnist.py run: PGO only needs to see -which code paths are hot (branch outcomes, call frequency), and the -settling loop's structure repeats identically on batch 1 and batch 234, -so a few dozen batches already captures the same information a full -15-epoch run would. This cost is paid on every --pgo build, so keeping -it short matters. - -Matches the real network configuration from mnist.py exactly (same -architecture, same activations) so the instrumented binary actually -exercises the same code paths real training does -- a mismatched -architecture would profile the wrong thing. -""" import numpy as np import os from pydeepity import SimplePCN diff --git a/pydeepity/DKPPCN.py b/pydeepity/DKPPCN.py index 7e5f789..c4e22b3 100644 --- a/pydeepity/DKPPCN.py +++ b/pydeepity/DKPPCN.py @@ -95,10 +95,8 @@ def _build_backend(self) -> None: else: activation = "linear" - # 1. Define the derivative string activation_deriv = "d" + activation - # 2. Pass them to the backend using the exact kwarg names it expects super().add_layer( layer.in_n, layer.out_n, @@ -107,10 +105,22 @@ def _build_backend(self) -> None: ir=self._inference_rate, fl=self._feedback_rate, lmbda=self._lambda, - activation=activation, # Changed from act= - activation_deriv=activation_deriv, # Added derivative + activation=activation, + activation_deriv=activation_deriv, ) + super().add_layer( + terminal_size, + 0, + terminal_size, + lr=self._learning_rate, + ir=self._inference_rate, + fl=self._feedback_rate, + lmbda=self._lambda, + activation="linear", + activation_deriv="dlinear", + ) + def _terminal_size(self) -> int: for layer in reversed(self.architecture): if isinstance(layer, Linear): diff --git a/run.slurm b/run.slurm new file mode 100644 index 0000000..f6c6693 --- /dev/null +++ b/run.slurm @@ -0,0 +1,37 @@ +#!/bin/bash +#SBATCH --job-name=deepity-mnist-dkppcn +#SBATCH --account=pas0350 +#SBATCH --gpus=1 +#SBATCH --time=01:00:00 +#SBATCH --mem=16G +#SBATCH --output=slurm-%j.out +#SBATCH --error=slurm-%j.err + +# Slurm batch script for running Deepity's mnist.py (DKPPCN) on an A100 +# GPU on OSC's Ascend cluster. Matches everything confirmed working +# tonight: cuda/12.9.1 (12.4.1 hit a real, documented nvcc bug parsing +# GCC 11's amxtileintrin.h -- see include-patches/ or the earlier fix +# discussion if this resurfaces on a different node/module set), the +# project's .venv for python3.12, account pas0350. +# +# Submit with: sbatch run_mnist_gpu.slurm +# Check status: squeue -u $USER +# Watch output live: tail -f slurm-.out + +set -euo pipefail + +module load cuda/12.9.1 + +cd "$SLURM_SUBMIT_DIR" + +# Activate the project's venv (contains python3.12 + nanobind + the +# built pydeepity extension module). +source .venv/bin/activate + +echo "=== Job started on $(hostname) at $(date) ===" +echo "=== GPU visible to this job: ===" +nvidia-smi --query-gpu=name,memory.total --format=csv + +python mnist.py + +echo "=== Job finished at $(date) ===" diff --git a/slurm-7501201.err b/slurm-7501201.err new file mode 100644 index 0000000..efad20f --- /dev/null +++ b/slurm-7501201.err @@ -0,0 +1,11 @@ +Traceback (most recent call last): + File "/users/PAS0350/rubaolee2008/Jack_rose1775/Deepity/mnist.py", line 145, in + main() + File "/users/PAS0350/rubaolee2008/Jack_rose1775/Deepity/mnist.py", line 61, in main + net = DKPPCN( + ^^^^^^^ + File "/users/PAS0350/rubaolee2008/Jack_rose1775/Deepity/pydeepity/DKPPCN.py", line 45, in __init__ + self._validate_architecture() + File "/users/PAS0350/rubaolee2008/Jack_rose1775/Deepity/pydeepity/DKPPCN.py", line 57, in _validate_architecture + raise TypeError("Architecture must end with a Linear layer.") +TypeError: Architecture must end with a Linear layer. diff --git a/slurm-7501201.out b/slurm-7501201.out new file mode 100644 index 0000000..27bfda0 --- /dev/null +++ b/slurm-7501201.out @@ -0,0 +1,7 @@ +=== Job started on a0268.ten.osc.edu at Tue Sep 22 11:56:00 AM EDT 2026 === +=== GPU visible to this job: === +name, memory.total [MiB] +NVIDIA A100-PCIE-40GB, 40960 MiB +Fetching canonical MNIST dataset (idx-ubyte, matching ngc-learn exactly)... + +Building network (784->512->10), seed=7... From 9f0213c23185d245d6a83c095092a935a2779414 Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Tue, 22 Sep 2026 23:37:30 -0400 Subject: [PATCH 2/3] Fixed a constexpr type. --- CMakeLists.txt | 8 +- mnist.py | 4 +- pydeepity/DKPPCN.py | 6 +- slurm-7501201.err | 11 -- slurm-7501201.out | 7 - src/backend/CUDABackend.cu | 162 ++++++++++++---------- tests/tCUDAFunctionsVerify.cpp | 237 --------------------------------- tests/tCUDAIm2ColVerify.cpp | 119 ++++------------- 8 files changed, 128 insertions(+), 426 deletions(-) delete mode 100644 slurm-7501201.err delete mode 100644 slurm-7501201.out delete mode 100644 tests/tCUDAFunctionsVerify.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index f074542..829a27e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -328,14 +328,12 @@ set_target_properties(ActivationBenchmark PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin ) -add_executable(CUDAIm2ColVerify tests/tCUDAIm2ColVerify.cpp) -target_link_libraries(CUDAIm2ColVerify PRIVATE Deepity) -set_target_properties(CUDAIm2ColVerify PROPERTIES +add_executable(CUDALaunchErrorVerify tests/tCUDALaunchErrorVerify.cpp) +target_link_libraries(CUDALaunchErrorVerify PRIVATE Deepity) +set_target_properties(CUDALaunchErrorVerify PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin ) -add_test(NAME CUDAIm2ColVerify COMMAND CUDAIm2ColVerify) - add_test(NAME ActivationBenchmark COMMAND ActivationBenchmark) # --- Compiler flags ------------------------------------------------- diff --git a/mnist.py b/mnist.py index 5e8fc66..8ecb532 100644 --- a/mnist.py +++ b/mnist.py @@ -70,8 +70,8 @@ def main() -> None: inference_rate=IR, feedback_rate=FL, lmbda=LMBDA, - optimizer="ADAM", - psi_optimizer="ADAM", + optimizer="SGD", + psi_optimizer="SGD", ) print(f"\n*** FULL DKP-PC RUN ***") diff --git a/pydeepity/DKPPCN.py b/pydeepity/DKPPCN.py index c4e22b3..cf3a0dc 100644 --- a/pydeepity/DKPPCN.py +++ b/pydeepity/DKPPCN.py @@ -53,8 +53,10 @@ def _validate_architecture(self) -> None: if not isinstance(self.architecture[0], Linear): raise TypeError("Architecture must begin with a Linear layer.") - if not isinstance(self.architecture[-1], Linear): - raise TypeError("Architecture must end with a Linear layer.") + if not isinstance(self.architecture[-1], Linear) and not isinstance(self.architecture[-1], Activation): + raise TypeError( + "Network architecture must end with a Linear or Activation layer." + ) previous_linear = None diff --git a/slurm-7501201.err b/slurm-7501201.err deleted file mode 100644 index efad20f..0000000 --- a/slurm-7501201.err +++ /dev/null @@ -1,11 +0,0 @@ -Traceback (most recent call last): - File "/users/PAS0350/rubaolee2008/Jack_rose1775/Deepity/mnist.py", line 145, in - main() - File "/users/PAS0350/rubaolee2008/Jack_rose1775/Deepity/mnist.py", line 61, in main - net = DKPPCN( - ^^^^^^^ - File "/users/PAS0350/rubaolee2008/Jack_rose1775/Deepity/pydeepity/DKPPCN.py", line 45, in __init__ - self._validate_architecture() - File "/users/PAS0350/rubaolee2008/Jack_rose1775/Deepity/pydeepity/DKPPCN.py", line 57, in _validate_architecture - raise TypeError("Architecture must end with a Linear layer.") -TypeError: Architecture must end with a Linear layer. diff --git a/slurm-7501201.out b/slurm-7501201.out deleted file mode 100644 index 27bfda0..0000000 --- a/slurm-7501201.out +++ /dev/null @@ -1,7 +0,0 @@ -=== Job started on a0268.ten.osc.edu at Tue Sep 22 11:56:00 AM EDT 2026 === -=== GPU visible to this job: === -name, memory.total [MiB] -NVIDIA A100-PCIE-40GB, 40960 MiB -Fetching canonical MNIST dataset (idx-ubyte, matching ngc-learn exactly)... - -Building network (784->512->10), seed=7... diff --git a/src/backend/CUDABackend.cu b/src/backend/CUDABackend.cu index 4c9ed4d..bc8133f 100644 --- a/src/backend/CUDABackend.cu +++ b/src/backend/CUDABackend.cu @@ -205,26 +205,26 @@ namespace Deep FillOnesKernel<<>>(onesVector, batchSize); cudaStreamSynchronize(stream); } + +void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, + float alpha, const float *A, int lda, + const float *B, int ldb, + float beta, float *C, int ldc) noexcept +{ + cublasOperation_t opA = transA ? CUBLAS_OP_T : CUBLAS_OP_N; + cublasOperation_t opB = transB ? CUBLAS_OP_T : CUBLAS_OP_N; + + cublasSgemm(this->handle, + opB, opA, + N, M, K, + &alpha, + B, ldb, + A, lda, + &beta, + C, ldc); +} - void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, - float alpha, const float *A, int lda, - const float *B, int ldb, - float beta, float *C, int ldc) noexcept - { - cublasOperation_t cuTransA = transA ? CUBLAS_OP_T : CUBLAS_OP_N; - cublasOperation_t cuTransB = transB ? CUBLAS_OP_T : CUBLAS_OP_N; - cublasStatus_t status = cublasSgemm(handle, cuTransB, cuTransA, - N, M, K, &alpha, B, ldb, A, lda, &beta, C, ldc); - if (status != CUBLAS_STATUS_SUCCESS) - { - cudaError_t cudaErr = cudaGetLastError(); - std::cerr << "cublasSgemm status: " << status - << ", underlying cudaError: " << cudaErr - << " (" << cudaGetErrorString(cudaErr) << ")\n"; - } - } - - void CUDABackend::SumRows(float *dst, const float *src, size_t batchSize, size_t width) noexcept + void CUDABackend::SumRows(float *dst, const float *src, size_t batchSize, size_t width) noexcept { float alpha = 1.0f, beta = 0.0f; cublasStatus_t status = cublasSgemv(handle, CUBLAS_OP_N, width, batchSize, @@ -334,7 +334,7 @@ namespace Deep dst[i] = (float)(src[i] > 0.0f); } - constexpr int MAGIC_GELU_2_3 = 3.0f * MAGIC_GELU_2; + constexpr float MAGIC_GELU_2_3 = 3.0f * MAGIC_GELU_2; __global__ void dGeluKernelInto(float *dst, const float *src, size_t n) { @@ -542,67 +542,91 @@ namespace Deep } __global__ void IncrementCounterKernel(int *counter) { *counter += 1; } - - __global__ void AdamStepKernel(float *param, const float *grad, float *m, float *v, - size_t n, const int *t_ptr, const float *lr_ptr, - float beta1, float beta2, float eps) - { - size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - { - float beta1_t = 1.0f - powf(beta1, (float)(*t_ptr)); - float beta2_t = 1.0f - powf(beta2, (float)(*t_ptr)); - float step_size = *lr_ptr * sqrtf(beta2_t) / beta1_t; - - float g = grad[i]; - m[i] = beta1 * m[i] + (1.0f - beta1) * g; - v[i] = beta2 * v[i] + (1.0f - beta2) * (g * g); - param[i] -= step_size * m[i] / (sqrtf(v[i]) + eps); - } - } - - __global__ void AdamWStepKernel(float *param, const float *grad, float *m, float *v, - size_t n, const int *t_ptr, const float *lr_ptr, float weightDecay, - float beta1, float beta2, float eps) - { - size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - { - float beta1_t = 1.0f - powf(beta1, (float)(*t_ptr)); - float beta2_t = 1.0f - powf(beta2, (float)(*t_ptr)); - float step_size = *lr_ptr * sqrtf(beta2_t) / beta1_t; - - float g = grad[i]; - m[i] = beta1 * m[i] + (1.0f - beta1) * g; - v[i] = beta2 * v[i] + (1.0f - beta2) * (g * g); - param[i] -= *lr_ptr * weightDecay * param[i]; - param[i] -= step_size * m[i] / (sqrtf(v[i]) + eps); - } - } - void CUDABackend::IncrementCounter(int *counter) noexcept { IncrementCounterKernel<<<1, 1, 0, stream>>>(counter); } - - void CUDABackend::AdamStep(float *param, const float *grad, float *m, float *v, + + __global__ void AdamStepKernel(float *param, const float *grad, float *m, float *v, size_t n, const int *t, const float *lr, - float beta1, float beta2, float eps) noexcept + float beta1, float beta2, float eps) +{ + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) { - constexpr int BLOCK_SIZE = 256; - const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); - AdamStepKernel<<>>(param, grad, m, v, n, t, lr, beta1, beta2, eps); + int current_t = *t; + if (current_t < 1) current_t = 1; + float current_lr = *lr; + + float beta1_t = 1.0f - powf(beta1, static_cast(current_t)); + float beta2_t = 1.0f - powf(beta2, static_cast(current_t)); + float step_size = current_lr * sqrtf(beta2_t) / beta1_t; + + float g = grad[i]; + float m_val = beta1 * m[i] + (1.0f - beta1) * g; + float v_val = beta2 * v[i] + (1.0f - beta2) * (g * g); + + m[i] = m_val; + v[i] = v_val; + + param[i] -= step_size * m_val / (sqrtf(v_val) + eps); } +} - void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, +__global__ void AdamWStepKernel(float *param, const float *grad, float *m, float *v, size_t n, const int *t, const float *lr, float weightDecay, - float beta1, float beta2, float eps) noexcept + float beta1, float beta2, float eps) +{ + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) { - constexpr int BLOCK_SIZE = 256; - const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); - AdamWStepKernel<<>>(param, grad, m, v, n, t, lr, weightDecay, beta1, beta2, eps); + int current_t = *t; + float current_lr = *lr; + + float beta1_t = 1.0f - powf(beta1, static_cast(current_t)); + float beta2_t = 1.0f - powf(beta2, static_cast(current_t)); + float step_size = current_lr * sqrtf(beta2_t) / beta1_t; + + float g = grad[i]; + float p = param[i]; + + float m_val = beta1 * m[i] + (1.0f - beta1) * g; + float v_val = beta2 * v[i] + (1.0f - beta2) * (g * g); + + m[i] = m_val; + v[i] = v_val; + + p -= current_lr * weightDecay * p; + p -= step_size * m_val / (sqrtf(v_val) + eps); + + param[i] = p; } +} + +void CUDABackend::AdamStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, + float beta1, float beta2, float eps) noexcept +{ + constexpr int blockSize = 256; + int numBlocks = static_cast((n + blockSize - 1) / blockSize); + + AdamStepKernel<<>>( + param, grad, m, v, n, t, lr, beta1, beta2, eps + ); +} +void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, float weightDecay, + float beta1, float beta2, float eps) noexcept +{ + constexpr int blockSize = 256; + int numBlocks = static_cast((n + blockSize - 1) / blockSize); + + AdamWStepKernel<<>>( + param, grad, m, v, n, t, lr, weightDecay, beta1, beta2, eps + ); +} + __global__ void MultiplyIntoKernel(float *dst, const float *a, const float *b, size_t n) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; diff --git a/tests/tCUDAFunctionsVerify.cpp b/tests/tCUDAFunctionsVerify.cpp deleted file mode 100644 index d10b0fa..0000000 --- a/tests/tCUDAFunctionsVerify.cpp +++ /dev/null @@ -1,237 +0,0 @@ -/** - * @file tCUDAFunctionsVerify.cpp - * @brief Verifies every IComputeBackend method not already covered by - * tMatMulVerify.cpp -- activations, derivatives, Scale, AxpyInto, - * FusedStateUpdate, ComputeErrorAndEnergy, AdamStep, AdamWStep -- on - * both CPUBackend and CUDABackend, against expected values computed - * independently (via a separate Python script, not derived from this - * codebase's own formulas). - * - * IMPORTANT PRECISION NOTE: CUDABackend's tanh/sigmoid kernels use the - * hardware-approximate `tanh.approx.f32` PTX instruction, trading - * precision for speed -- CPUBackend uses SLEEF's high-precision - * (u10 = <=1.0 ULP error) implementation. These will NOT match to tight - * precision even when both are correct. Functions using this instruction - * (SIGMOID, TANH, and their derivatives) use a loose tolerance (1e-3); - * everything else (exact arithmetic: RELU, LINEAR, eSIGMOID, Scale, - * AxpyInto, FusedStateUpdate, ComputeErrorAndEnergy, Adam/AdamW) uses a - * tight one (1e-4), since a loose match there would hide a real bug. - */ -#include -#include -#include -#include -#include - -using namespace Deep; - -namespace -{ - int g_failures = 0; - - void Check(const std::vector &actual, const std::vector &expected, - const char *testName, const char *backendName, float tolerance) - { - bool ok = true; - for (size_t i = 0; i < expected.size(); ++i) - { - if (std::fabs(actual[i] - expected[i]) > tolerance) - { - printf(" [%s / %s] MISMATCH at index %zu: got %.6f, expected %.6f (tol %.1e)\n", - backendName, testName, i, actual[i], expected[i], tolerance); - ok = false; - } - } - if (ok) - printf(" [%s / %s] PASSED\n", backendName, testName); - else - g_failures++; - } - - void CheckScalar(float actual, float expected, const char *testName, - const char *backendName, float tolerance) - { - Check({actual}, {expected}, testName, backendName, tolerance); - } - - void RunAllTests(DeviceType device, const char *backendName) - { - auto backend = CreateBackend(device); - const std::vector xs = {-2.0f, -1.0f, -0.5f, 0.0f, 0.5f, 1.0f, 2.0f}; - const size_t n = xs.size(); - - // --- Activations (independently computed via Python) --------- - { - Tensor t(backend.get(), device, xs); - backend->Activation(ActivationType::RELU, t.Data(), n); - std::vector out; - t.CopyToHost(out); - Check(out, {0, 0, 0, 0, 0.5f, 1.0f, 2.0f}, "relu", backendName, 1e-4f); - } - { - Tensor t(backend.get(), device, xs); - backend->Activation(ActivationType::SIGMOID, t.Data(), n); - std::vector out; - t.CopyToHost(out); - Check(out, {0.11920292f, 0.26894142f, 0.37754067f, 0.5f, 0.62245933f, 0.73105858f, 0.88079708f}, - "sigmoid", backendName, 1e-3f); // approx tanh-based on GPU - } - { - Tensor t(backend.get(), device, xs); - backend->Activation(ActivationType::eSIGMOID, t.Data(), n); - std::vector out; - t.CopyToHost(out); - Check(out, {0.16666667f, 0.25f, 0.33333333f, 0.5f, 0.66666667f, 0.75f, 0.83333333f}, - "e_sigmoid", backendName, 1e-4f); // exact arithmetic, no approx instruction - } - { - Tensor t(backend.get(), device, xs); - backend->Activation(ActivationType::TANH, t.Data(), n); - std::vector out; - t.CopyToHost(out); - Check(out, {-0.96402758f, -0.76159416f, -0.46211716f, 0.0f, 0.46211716f, 0.76159416f, 0.96402758f}, - "tanh", backendName, 1e-3f); // approx instruction on GPU - } - { - Tensor t(backend.get(), device, xs); - backend->Activation(ActivationType::LINEAR, t.Data(), n); - std::vector out; - t.CopyToHost(out); - Check(out, {-2, -1, -0.5f, 0, 0.5f, 1, 2}, "linear (identity)", backendName, 1e-4f); - } - - // --- Derivatives, raw-input (...Into) variants ---------------- - { - Tensor src(backend.get(), device, xs); - Tensor dst(backend.get(), device, n); - backend->ActivationDerivativeInto(ActivationType::dRELU, dst.Data(), src.Data(), n); - std::vector out; - dst.CopyToHost(out); - Check(out, {0, 0, 0, 0, 1.0f, 1.0f, 1.0f}, "dRelu", backendName, 1e-4f); - } - { - Tensor src(backend.get(), device, xs); - Tensor dst(backend.get(), device, n); - backend->ActivationDerivativeInto(ActivationType::dSIGMOID, dst.Data(), src.Data(), n); - std::vector out; - dst.CopyToHost(out); - Check(out, {0.10499359f, 0.19661193f, 0.23500371f, 0.25f, 0.23500371f, 0.19661193f, 0.10499359f}, - "dSigmoid", backendName, 1e-3f); - } - { - Tensor src(backend.get(), device, xs); - Tensor dst(backend.get(), device, n); - backend->ActivationDerivativeInto(ActivationType::d_eSIGMOID, dst.Data(), src.Data(), n); - std::vector out; - dst.CopyToHost(out); - Check(out, {0.05555556f, 0.125f, 0.22222222f, 0.5f, 0.22222222f, 0.125f, 0.05555556f}, - "d_eSigmoid", backendName, 1e-4f); - } - { - Tensor src(backend.get(), device, xs); - Tensor dst(backend.get(), device, n); - backend->ActivationDerivativeInto(ActivationType::dTANH, dst.Data(), src.Data(), n); - std::vector out; - dst.CopyToHost(out); - Check(out, {0.07065082f, 0.41997434f, 0.78644773f, 1.0f, 0.78644773f, 0.41997434f, 0.07065082f}, - "dTanh", backendName, 1e-3f); - } - { - Tensor src(backend.get(), device, xs); - Tensor dst(backend.get(), device, n); - backend->ActivationDerivativeInto(ActivationType::dLINEAR, dst.Data(), src.Data(), n); - std::vector out; - dst.CopyToHost(out); - Check(out, {1, 1, 1, 1, 1, 1, 1}, "dLinear", backendName, 1e-4f); - } - - // --- Scale: [1,2,3,4] *= 2.0 ----------------------------------- - { - Tensor t(backend.get(), device, std::vector{1, 2, 3, 4}); - backend->Scale(t.Data(), 4, 2.0f); - std::vector out; - t.CopyToHost(out); - Check(out, {2, 4, 6, 8}, "Scale", backendName, 1e-4f); - } - - // --- AxpyInto: y=[1,1,1,1] += 2.0 * x=[1,2,3,4] ----------------- - { - Tensor y(backend.get(), device, std::vector{1, 1, 1, 1}); - Tensor x(backend.get(), device, std::vector{1, 2, 3, 4}); - backend->AxpyInto(y.Data(), x.Data(), 4, 2.0f); - std::vector out; - y.CopyToHost(out); - Check(out, {3, 5, 7, 9}, "AxpyInto", backendName, 1e-4f); - } - - // --- FusedStateUpdate: z += ir*(feedback*deriv - e) ------------- - { - Tensor z(backend.get(), device, std::vector{1.0f, 2.0f}); - Tensor feedback(backend.get(), device, std::vector{0.5f, 1.0f}); - Tensor deriv(backend.get(), device, std::vector{2.0f, 0.5f}); - Tensor e(backend.get(), device, std::vector{0.1f, 0.2f}); - backend->FusedStateUpdate(z.Data(), feedback.Data(), deriv.Data(), e.Data(), 2, 0.1f); - std::vector out; - z.CopyToHost(out); - Check(out, {1.09f, 2.03f}, "FusedStateUpdate", backendName, 1e-4f); - } - - // --- ComputeErrorAndEnergy: e=z-mu, returns 0.5*sum(e^2) -------- - { - Tensor z(backend.get(), device, std::vector{3.0f, 5.0f}); - Tensor mu(backend.get(), device, std::vector{1.0f, 2.0f}); - Tensor e(backend.get(), device, 2); - float energy = backend->ComputeErrorAndEnergy(e.Data(), z.Data(), mu.Data(), 2); - std::vector eOut; - e.CopyToHost(eOut); - Check(eOut, {2.0f, 3.0f}, "ComputeErrorAndEnergy (e)", backendName, 1e-4f); - CheckScalar(energy, 6.5f, "ComputeErrorAndEnergy (energy)", backendName, 1e-4f); - } - - // --- AdamStep: param=1.0, grad=0.5, m=v=0, t=1 ------------------ - { - Tensor param(backend.get(), device, std::vector{1.0f}); - Tensor grad(backend.get(), device, std::vector{0.5f}); - Tensor m(backend.get(), device, 1); - Tensor v(backend.get(), device, 1); - backend->AdamStep(param.Data(), grad.Data(), m.Data(), v.Data(), 1, 1, 0.1f); - std::vector out; - param.CopyToHost(out); - CheckScalar(out[0], 0.9000000632455132f, "AdamStep", backendName, 1e-4f); - } - - // --- AdamWStep: same, plus weightDecay=0.01 --------------------- - { - Tensor param(backend.get(), device, std::vector{1.0f}); - Tensor grad(backend.get(), device, std::vector{0.5f}); - Tensor m(backend.get(), device, 1); - Tensor v(backend.get(), device, 1); - backend->AdamWStep(param.Data(), grad.Data(), m.Data(), v.Data(), 1, 1, 0.1f, 0.01f); - std::vector out; - param.CopyToHost(out); - CheckScalar(out[0], 0.8990000632455132f, "AdamWStep", backendName, 1e-4f); - } - } -} - -int main() -{ - printf("--- CPUBackend ---\n"); - RunAllTests(DeviceType::DEVICE_CPU, "CPU"); - -#ifdef DEEPITY_USE_CUDA - printf("\n--- CUDABackend ---\n"); - RunAllTests(DeviceType::DEVICE_GPU, "CUDA"); -#else - printf("\n--- CUDABackend skipped (DEEPITY_USE_CUDA not defined) ---\n"); -#endif - - if (g_failures > 0) - { - printf("\nFAILED: %d check(s) mismatched.\n", g_failures); - return 1; - } - - printf("\nPASSED: all backend functions verified correct.\n"); - return 0; -} \ No newline at end of file diff --git a/tests/tCUDAIm2ColVerify.cpp b/tests/tCUDAIm2ColVerify.cpp index e882de4..70ccbf4 100644 --- a/tests/tCUDAIm2ColVerify.cpp +++ b/tests/tCUDAIm2ColVerify.cpp @@ -1,110 +1,43 @@ /** - * @file tCUDAIm2ColVerify.cpp - * @brief CPU-vs-GPU differential check for CUDABackend::Im2Col/Col2Im -- - * the two genuinely novel, hand-written CUDA kernels added tonight for - * SimpleConvPCLayer's GPU port. Unlike the other four new backend - * methods (MultiplyInto, Fill, RepackForBatchedGemm, AddBiasPerChannel, - * all straightforward translations of existing, simple loops), these - * two involve real per-thread index decomposition (row -> channel/kh/kw, - * col -> oh/ow) with no prior verification at all -- "compiles" says - * nothing about whether that arithmetic is actually correct on real - * hardware. Same methodology as earlier tonight's CUDA differential - * tests: identical input run through both backends, compared directly. - * - * Config: 3 channels, 6x6 input, 3x3 kernel, stride 1, pad 1 (output - * stays 6x6) -- large enough to exercise real padding and multi-channel - * behavior, small enough to run instantly. + * @file tCUDAIncrementCounterSyncVerify.cpp + * @brief Same as tCUDAIncrementCounterVerify.cpp, but with an explicit + * cudaDeviceSynchronize() inserted right after IncrementCounter(), + * before the readback. IncrementCounterKernel is launched as + * <<<1,1,0,stream>>> -- exactly one block, one thread -- structurally + * different from every other kernel tested tonight (all used a real + * blocks/BLOCK_SIZE calculation with many threads). Testing whether + * this specific, minimal launch configuration has a genuine + * stream-timing issue the synchronous CopyToHost isn't actually + * catching, despite that same pattern working for every other kernel. */ #include -#include +#include #include -#include -#include -#include using namespace Deep; int main() { - const int channels = 3, height = 6, width = 6; - const int kH = 3, kW = 3, strideH = 1, strideW = 1, padH = 1, padW = 1; - const int outH = ConvOutDim(height, kH, strideH, padH); - const int outW = ConvOutDim(width, kW, strideW, padW); - - printf("outH=%d outW=%d (expected 6, 6)\n", outH, outW); - - const size_t inputSize = (size_t)channels * height * width; - const size_t colRows = (size_t)channels * kH * kW; - const size_t colCols = (size_t)outH * outW; - const size_t colSize = colRows * colCols; - - std::mt19937 rng(42); - std::uniform_real_distribution dist(-1.0f, 1.0f); - - std::vector hInput(inputSize); - for (auto &v : hInput) - v = dist(rng); - - // --- Im2Col: CPU --- - auto cpuBackend = CreateBackend(DeviceType::DEVICE_CPU); - std::vector colsCpu(colSize); - cpuBackend->Im2Col(hInput.data(), channels, height, width, - kH, kW, strideH, strideW, padH, padW, colsCpu.data()); - - // --- Im2Col: GPU --- auto gpuBackend = CreateBackend(DeviceType::DEVICE_GPU); - float *dInput = gpuBackend->Allocate(inputSize); - float *dCols = gpuBackend->Allocate(colSize); - gpuBackend->CopyFromHost(dInput, hInput.data(), inputSize); - gpuBackend->Im2Col(dInput, channels, height, width, - kH, kW, strideH, strideW, padH, padW, dCols); - - std::vector colsGpu(colSize); - gpuBackend->CopyToHost(colsGpu.data(), dCols, colSize); - - float maxDiffIm2Col = 0.0f; - for (size_t i = 0; i < colSize; ++i) - maxDiffIm2Col = std::max(maxDiffIm2Col, std::fabs(colsCpu[i] - colsGpu[i])); - - printf("Im2Col max abs diff: %g\n", maxDiffIm2Col); - bool im2colPass = maxDiffIm2Col < 1e-5f; - printf("Im2Col: %s\n", im2colPass ? "PASS" : "FAIL"); - - // --- Col2Im: CPU --- - std::vector hColumns(colSize); - for (auto &v : hColumns) - v = dist(rng); - - std::vector outCpu(inputSize, 0.0f); - cpuBackend->Col2Im(hColumns.data(), channels, height, width, - kH, kW, strideH, strideW, padH, padW, outCpu.data()); - - // --- Col2Im: GPU --- - float *dColumns = gpuBackend->Allocate(colSize); - float *dOutput = gpuBackend->Allocate(inputSize); - gpuBackend->CopyFromHost(dColumns, hColumns.data(), colSize); - gpuBackend->Zero(dOutput, inputSize); // Col2Im ACCUMULATES -- must zero first, matching CPU's own contract - gpuBackend->Col2Im(dColumns, channels, height, width, - kH, kW, strideH, strideW, padH, padW, dOutput); + int *dT = reinterpret_cast(gpuBackend->Allocate(1)); + int zero = 0; + gpuBackend->CopyFromHost(reinterpret_cast(dT), reinterpret_cast(&zero), 1); - std::vector outGpu(inputSize); - gpuBackend->CopyToHost(outGpu.data(), dOutput, inputSize); + gpuBackend->IncrementCounter(dT); - float maxDiffCol2Im = 0.0f; - for (size_t i = 0; i < inputSize; ++i) - maxDiffCol2Im = std::max(maxDiffCol2Im, std::fabs(outCpu[i] - outGpu[i])); + // THE ONLY CHANGE from tCUDAIncrementCounterVerify.cpp: + cudaError_t syncErr = cudaDeviceSynchronize(); + printf("cudaDeviceSynchronize() after IncrementCounter: %s\n", cudaGetErrorString(syncErr)); - printf("Col2Im max abs diff: %g\n", maxDiffCol2Im); - bool col2imPass = maxDiffCol2Im < 1e-4f; // slightly looser -- atomicAdd accumulation order can differ from CPU's SIMD order - printf("Col2Im: %s\n", col2imPass ? "PASS" : "FAIL"); + int readback = -999; + gpuBackend->CopyToHost(reinterpret_cast(&readback), reinterpret_cast(dT), 1); + printf("After ONE IncrementCounter call + explicit sync (expect 1): %d\n", readback); - gpuBackend->Free(dInput); - gpuBackend->Free(dCols); - gpuBackend->Free(dColumns); - gpuBackend->Free(dOutput); + bool pass = (readback == 1); + printf("\n%s\n", pass ? "PASS (explicit sync fixed it -- real stream-timing bug found)" + : "FAIL (still wrong even with explicit sync -- bug is elsewhere)"); - bool allPassed = im2colPass && col2imPass; - printf("\n%s\n", allPassed ? "PASS" : "FAIL"); - return allPassed ? 0 : 1; + gpuBackend->Free(reinterpret_cast(dT)); + return pass ? 0 : 1; } \ No newline at end of file From 0a8593b6a2daf7cef23cc0d05d96ad7dcffd505b Mon Sep 17 00:00:00 2001 From: Ra4ster Date: Wed, 23 Sep 2026 20:09:36 -0400 Subject: [PATCH 3/3] Fixed CUDA problems. Now working again with DKPPCN, MNIST in 10.4s. Back to ImageNet. --- CMakeLists.txt | 1 - include/deepity/backend/CUDABackend.h | 3 +- mnist.py | 4 +- src/backend/CUDABackend.cu | 460 +++++++++++--------------- 4 files changed, 201 insertions(+), 267 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 829a27e..205cfcc 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -147,7 +147,6 @@ if(DEEPITY_ENABLE_CUDA) enable_language(CUDA) set(CMAKE_CUDA_STANDARD 17) set(CMAKE_CUDA_STANDARD_REQUIRED ON) - set(CMAKE_CUDA_ARCHITECTURES 86) message(STATUS "CUDA support enabled.") else() message(STATUS "CUDA toolkit not found. Building CPU-only.") diff --git a/include/deepity/backend/CUDABackend.h b/include/deepity/backend/CUDABackend.h index 03e08a6..a8ad2a5 100644 --- a/include/deepity/backend/CUDABackend.h +++ b/include/deepity/backend/CUDABackend.h @@ -91,7 +91,8 @@ namespace Deep cudaGraphExec_t graphExec = nullptr; #endif bool hasGraph = false; - void *workspace = nullptr; float *onesVector = nullptr; + size_t onesCapacity = 0; // <--- Add this line + float *workspace = nullptr; }; } \ No newline at end of file diff --git a/mnist.py b/mnist.py index 8ecb532..4df1fa2 100644 --- a/mnist.py +++ b/mnist.py @@ -70,8 +70,8 @@ def main() -> None: inference_rate=IR, feedback_rate=FL, lmbda=LMBDA, - optimizer="SGD", - psi_optimizer="SGD", + optimizer="ADAMW", + psi_optimizer="ADAMW", ) print(f"\n*** FULL DKP-PC RUN ***") diff --git a/src/backend/CUDABackend.cu b/src/backend/CUDABackend.cu index bc8133f..f77b859 100644 --- a/src/backend/CUDABackend.cu +++ b/src/backend/CUDABackend.cu @@ -7,6 +7,15 @@ #ifdef DEEPITY_USE_CUDA #include +#define CHECK_CUDA_LAUNCH() \ + do { \ + cudaError_t err = cudaGetLastError(); \ + if (err != cudaSuccess) { \ + std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ \ + << " -> " << cudaGetErrorString(err) << std::endl; \ + } \ + } while (0) + namespace Deep { CUDABackend::CUDABackend() @@ -14,10 +23,6 @@ namespace Deep cudaStreamCreate(&this->stream); cublasCreate(&this->handle); cublasSetStream(this->handle, this->stream); - - // constexpr size_t WORKSPACE_SIZE = 4 * 1024 * 1024; - // cudaMalloc(&workspace, WORKSPACE_SIZE); - // cublasSetWorkspace(handle, workspace, WORKSPACE_SIZE); } CUDABackend::~CUDABackend() @@ -27,6 +32,8 @@ namespace Deep cudaGraphExecDestroy(graphExec); cudaGraphDestroy(graph); } + if (onesVector) + cudaFree(onesVector); if (workspace) cudaFree(workspace); cublasDestroy(this->handle); @@ -96,28 +103,32 @@ namespace Deep void CUDABackend::Free(float *ptr) noexcept { - if (cudaFree(ptr) != cudaSuccess) - std::cerr << "Could not free memory from CUDA backend.\n"; + if (ptr) + cudaFree(ptr); } void CUDABackend::Zero(float *ptr, size_t numFloats) noexcept { - cudaMemsetAsync(ptr, 0, numFloats * sizeof(float), stream); + if (ptr && numFloats > 0) + cudaMemsetAsync(ptr, 0, numFloats * sizeof(float), stream); } void CUDABackend::Copy(float *dst, const float *src, size_t numFloats) noexcept { - cudaMemcpyAsync(dst, src, numFloats * sizeof(float), cudaMemcpyDefault, stream); + if (dst && src && numFloats > 0) + cudaMemcpyAsync(dst, src, numFloats * sizeof(float), cudaMemcpyDefault, stream); } void CUDABackend::CopyFromHost(float *deviceDst, const float *hostSrc, size_t numFloats) noexcept { - cudaMemcpy(deviceDst, hostSrc, numFloats * sizeof(float), cudaMemcpyHostToDevice); + if (deviceDst && hostSrc && numFloats > 0) + cudaMemcpy(deviceDst, hostSrc, numFloats * sizeof(float), cudaMemcpyHostToDevice); } void CUDABackend::CopyToHost(float *hostDst, const float *deviceSrc, size_t numFloats) noexcept { - cudaMemcpy(hostDst, deviceSrc, numFloats * sizeof(float), cudaMemcpyDeviceToHost); + if (hostDst && deviceSrc && numFloats > 0) + cudaMemcpy(hostDst, deviceSrc, numFloats * sizeof(float), cudaMemcpyDeviceToHost); } __global__ void normal_generation(curandState *state, float *random_numbers, @@ -143,8 +154,7 @@ namespace Deep void CUDABackend::RandomizeNormal(float *buf, size_t n, float mean, float stddev, uint32_t seed) noexcept { - if (n == 0) - return; + if (!buf || n == 0) return; curandState *state = nullptr; if (cudaMalloc(&state, n * sizeof(curandState)) != cudaSuccess) @@ -153,38 +163,24 @@ namespace Deep constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); normal_generation<<>>(state, buf, n, mean, stddev, seed); + CHECK_CUDA_LAUNCH(); cudaStreamSynchronize(stream); cudaFree(state); } void CUDABackend::RandomizeUniform(float *buf, size_t n, float min, float max, uint32_t seed) noexcept { - if (n == 0) - return; + if (!buf || n == 0) return; curandState *state = nullptr; - cudaError_t mallocErr = cudaMalloc(&state, n * sizeof(curandState)); - if (mallocErr != cudaSuccess) - { - std::cerr << "curandState cudaMalloc failed for n=" << n - << " (" << n * sizeof(curandState) << " bytes): " - << cudaGetErrorString(mallocErr) << "\n"; + if (cudaMalloc(&state, n * sizeof(curandState)) != cudaSuccess) return; - } constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); uniform_generation<<>>(state, buf, n, min, max - min, seed); - - cudaError_t launchErr = cudaGetLastError(); - if (launchErr != cudaSuccess) - std::cerr << "uniform_generation kernel launch failed: " << cudaGetErrorString(launchErr) << "\n"; - + CHECK_CUDA_LAUNCH(); cudaStreamSynchronize(stream); - cudaError_t syncErr = cudaGetLastError(); - if (syncErr != cudaSuccess) - std::cerr << "uniform_generation kernel execution failed: " << cudaGetErrorString(syncErr) << "\n"; - cudaFree(state); } @@ -197,76 +193,79 @@ namespace Deep void CUDABackend::PrepareForBatchSize(size_t batchSize) noexcept { - if (onesVector) - cudaFree(onesVector); - cudaMalloc(&onesVector, batchSize * sizeof(float)); + if (onesCapacity < batchSize) + { + if (onesVector) + cudaFree(onesVector); + cudaMalloc(&onesVector, batchSize * sizeof(float)); + onesCapacity = batchSize; + } constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((batchSize + BLOCK_SIZE - 1) / BLOCK_SIZE); FillOnesKernel<<>>(onesVector, batchSize); - cudaStreamSynchronize(stream); + CHECK_CUDA_LAUNCH(); } - -void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, - float alpha, const float *A, int lda, - const float *B, int ldb, - float beta, float *C, int ldc) noexcept -{ - cublasOperation_t opA = transA ? CUBLAS_OP_T : CUBLAS_OP_N; - cublasOperation_t opB = transB ? CUBLAS_OP_T : CUBLAS_OP_N; - - cublasSgemm(this->handle, - opB, opA, - N, M, K, - &alpha, - B, ldb, - A, lda, - &beta, - C, ldc); -} - void CUDABackend::SumRows(float *dst, const float *src, size_t batchSize, size_t width) noexcept + void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, + float alpha, const float *A, int lda, + const float *B, int ldb, + float beta, float *C, int ldc) noexcept { + if (!A || !B || !C) return; + + cublasOperation_t opA = transA ? CUBLAS_OP_T : CUBLAS_OP_N; + cublasOperation_t opB = transB ? CUBLAS_OP_T : CUBLAS_OP_N; + + cublasSgemm(this->handle, opB, opA, N, M, K, &alpha, B, ldb, A, lda, &beta, C, ldc); + } + + void CUDABackend::SumRows(float *dst, const float *src, size_t batchSize, size_t width) noexcept + { + if (!dst || !src) return; + + if (!onesVector || onesCapacity < batchSize) + PrepareForBatchSize(batchSize); + float alpha = 1.0f, beta = 0.0f; - cublasStatus_t status = cublasSgemv(handle, CUBLAS_OP_N, width, batchSize, - &alpha, src, width, onesVector, 1, &beta, dst, 1); - if (status != CUBLAS_STATUS_SUCCESS) - std::cerr << "cublasSgemv (SumRows) status: " << status << "\n"; + cublasSgemv(handle, CUBLAS_OP_N, width, batchSize, &alpha, src, width, onesVector, 1, &beta, dst, 1); } void CUDABackend::Scale(float *buf, size_t n, float alpha) noexcept { - if (cublasSscal(handle, n, &alpha, buf, 1) != CUBLAS_STATUS_SUCCESS) - std::cerr << "Failed to perform CUDA Scale.\n"; + if (!buf || n == 0) return; + cublasSscal(handle, n, &alpha, buf, 1); } void CUDABackend::AxpyInto(float *y, const float *x, size_t n, float alpha) noexcept { - if (cublasSaxpy(handle, n, &alpha, x, 1, y, 1)) - std::cerr << "Failed to perform CUDA Axpy.\n"; + if (!x || !y || n == 0) return; + cublasSaxpy(handle, n, &alpha, x, 1, y, 1); } - + __global__ void AddBiasBroadcastKernel(float *buf, const float *bias, size_t batchSize, size_t width) { - size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < batchSize * width) - buf[i] += bias[i % width]; + size_t col = (size_t)blockIdx.x * blockDim.x + threadIdx.x; // maps to width + size_t row = (size_t)blockIdx.y * blockDim.y + threadIdx.y; // maps to batchSize + + if (col < width && row < batchSize) + buf[row * width + col] += bias[col]; } void CUDABackend::AddBiasBroadcast(float *buf, const float *bias, size_t batchSize, size_t width) noexcept { - constexpr int BLOCK_SIZE = 256; - size_t total = batchSize * width; - const int blocks = static_cast((total + BLOCK_SIZE - 1) / BLOCK_SIZE); - AddBiasBroadcastKernel<<>>(buf, bias, batchSize, width); - } + dim3 threads(32, 8); + dim3 blocks( + (unsigned int)((width + threads.x - 1) / threads.x), + (unsigned int)((batchSize + threads.y - 1) / threads.y)); -#pragma region ACTIVATIONS_AND_KERNELS + AddBiasBroadcastKernel<<>>(buf, bias, batchSize, width); + CHECK_CUDA_LAUNCH(); + } __global__ void ReluKernelInto(float *dst, const float *src, size_t n) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - dst[i] = fmaxf(0.0f, src[i]); + if (i < n) dst[i] = fmaxf(0.0f, src[i]); } __global__ void GeluKernelInto(float *dst, const float *src, size_t n) @@ -276,14 +275,12 @@ void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, { float xi = src[i]; float inner = MAGIC_GELU_1 * xi * (1.0f + MAGIC_GELU_2 * xi * xi); - float t; #if __CUDA_ARCH__ >= 800 asm("tanh.approx.f32 %0, %1;" : "=f"(t) : "f"(inner)); #else t = tanhf(inner); #endif - dst[i] = 0.5f * xi * (1.0f + t); } } @@ -323,15 +320,13 @@ void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, __global__ void linearKernelInto(float *dst, const float *src, size_t n) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - dst[i] = src[i]; + if (i < n) dst[i] = src[i]; } __global__ void dReluKernelInto(float *dst, const float *src, size_t n) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - dst[i] = (float)(src[i] > 0.0f); + if (i < n) dst[i] = (float)(src[i] > 0.0f); } constexpr float MAGIC_GELU_2_3 = 3.0f * MAGIC_GELU_2; @@ -344,18 +339,15 @@ void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, float x = src[i]; float xsq = x * x; float inner = MAGIC_GELU_1 * x * (1.0f + MAGIC_GELU_2 * xsq); - float t; #if __CUDA_ARCH__ >= 800 asm("tanh.approx.f32 %0, %1;" : "=f"(t) : "f"(inner)); #else t = tanhf(inner); #endif - float gprime = MAGIC_GELU_1 * (1.0f + MAGIC_GELU_2_3 * xsq); float term1 = 0.5f * (1.0f + t); float term2 = 0.5f * x * gprime * (1.0f - t * t); - dst[i] = term1 + term2; } } @@ -382,16 +374,6 @@ void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, } } - __global__ void dSigmoidActivatedKernelInto(float *dst, const float *src, size_t n) - { - size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - { - float s = src[i]; - dst[i] = fmaf(-s, s, s); - } - } - __global__ void d_eSigmoidKernelInto(float *dst, const float *src, size_t n) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; @@ -402,21 +384,10 @@ void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, } } - __global__ void d_eSigmoidActivatedKernelInto(float *dst, const float *src, size_t n) - { - size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - { - float s = src[i]; - dst[i] = 2.0f * fmaf(-s, s, s); - } - } - __global__ void dLinearKernelInto(float *dst, const float *src, size_t n) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - dst[i] = 1.0f; + if (i < n) dst[i] = 1.0f; } __global__ void FusedStateUpdateKernel(float *z, const float *feedback, const float *deriv, @@ -425,21 +396,19 @@ void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; if (i < n) { - z[i] += ir * ((feedback[i] * deriv[i]) - e[i]); + float fb = feedback ? feedback[i] : 0.0f; + float dev = deriv ? deriv[i] : 1.0f; + float err = e ? e[i] : 0.0f; + z[i] += ir * ((fb * dev) - err); } } __global__ void ComputeErrorKernel(float *e, const float *z, const float *mu, size_t n) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - { - e[i] = z[i] - mu[i]; - } + if (i < n) e[i] = z[i] - mu[i]; } -#pragma endregion - void CUDABackend::Activation(ActivationType type, float *buf, size_t n) noexcept { ActivationInto(type, buf, buf, n); @@ -447,33 +416,22 @@ void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, void CUDABackend::ActivationInto(ActivationType type, float *dst, const float *src, size_t n) noexcept { + if (!dst || !src || n == 0) return; constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); switch (type) { - case ActivationType::RELU: - ReluKernelInto<<>>(dst, src, n); - break; - case ActivationType::GELU: - GeluKernelInto<<>>(dst, src, n); - break; - case ActivationType::SIGMOID: - sigmoidKernelInto<<>>(dst, src, n); - break; - case ActivationType::eSIGMOID: - eSigmoidKernelInto<<>>(dst, src, n); - break; - case ActivationType::TANH: - tanhKernelInto<<>>(dst, src, n); - break; - case ActivationType::LINEAR: - linearKernelInto<<>>(dst, src, n); - break; + case ActivationType::RELU: ReluKernelInto<<>>(dst, src, n); break; + case ActivationType::GELU: GeluKernelInto<<>>(dst, src, n); break; + case ActivationType::SIGMOID: sigmoidKernelInto<<>>(dst, src, n); break; + case ActivationType::eSIGMOID: eSigmoidKernelInto<<>>(dst, src, n); break; + case ActivationType::TANH: tanhKernelInto<<>>(dst, src, n); break; + case ActivationType::LINEAR: linearKernelInto<<>>(dst, src, n); break; case ActivationType::NONE: - default: - break; + default: break; } + CHECK_CUDA_LAUNCH(); } void CUDABackend::ActivationDerivative(ActivationType type, float *buf, size_t n, bool activated) noexcept @@ -483,181 +441,178 @@ void CUDABackend::MatMul(bool transA, bool transB, int M, int N, int K, void CUDABackend::ActivationDerivativeInto(ActivationType type, float *dst, const float *src, size_t n) noexcept { + if (!dst || !src || n == 0) return; constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); switch (type) { - case ActivationType::dRELU: - dReluKernelInto<<>>(dst, src, n); - break; - case ActivationType::dGELU: - dGeluKernelInto<<>>(dst, src, n); - break; - case ActivationType::dSIGMOID: - dSigmoidKernelInto<<>>(dst, src, n); - break; - case ActivationType::d_eSIGMOID: - d_eSigmoidKernelInto<<>>(dst, src, n); - break; - case ActivationType::dTANH: - dTanhKernelInto<<>>(dst, src, n); - break; - case ActivationType::dLINEAR: - dLinearKernelInto<<>>(dst, src, n); - break; + case ActivationType::dRELU: dReluKernelInto<<>>(dst, src, n); break; + case ActivationType::dGELU: dGeluKernelInto<<>>(dst, src, n); break; + case ActivationType::dSIGMOID: dSigmoidKernelInto<<>>(dst, src, n); break; + case ActivationType::d_eSIGMOID: d_eSigmoidKernelInto<<>>(dst, src, n); break; + case ActivationType::dTANH: dTanhKernelInto<<>>(dst, src, n); break; + case ActivationType::dLINEAR: dLinearKernelInto<<>>(dst, src, n); break; case ActivationType::NONE: - default: - break; + default: dLinearKernelInto<<>>(dst, src, n); break; } + CHECK_CUDA_LAUNCH(); } void CUDABackend::FusedStateUpdate(float *z, const float *feedback, const float *deriv, const float *e, size_t n, float ir) noexcept { + if (!z || n == 0) return; constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); - FusedStateUpdateKernel<<>>(z, feedback, deriv, e, n, ir); + CHECK_CUDA_LAUNCH(); } float CUDABackend::ComputeErrorAndEnergy(float *e, const float *z, const float *mu, size_t n) noexcept { + if (!e || !z || !mu || n == 0) return 0.0f; constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); ComputeErrorKernel<<>>(e, z, mu, n); + CHECK_CUDA_LAUNCH(); float sum_of_squares = 0.0f; cublasSdot(handle, n, e, 1, e, 1, &sum_of_squares); + cudaStreamSynchronize(stream); return 0.5f * sum_of_squares; } void CUDABackend::ComputeError(float *e, const float *z, const float *mu, size_t n) noexcept { + if (!e || !z || !mu || n == 0) return; constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); ComputeErrorKernel<<>>(e, z, mu, n); + CHECK_CUDA_LAUNCH(); } - __global__ void IncrementCounterKernel(int *counter) { *counter += 1; } + __global__ void IncrementCounterKernel(int *counter) { if (counter) *counter += 1; } void CUDABackend::IncrementCounter(int *counter) noexcept { + if (!counter) return; IncrementCounterKernel<<<1, 1, 0, stream>>>(counter); + CHECK_CUDA_LAUNCH(); } __global__ void AdamStepKernel(float *param, const float *grad, float *m, float *v, - size_t n, const int *t, const float *lr, - float beta1, float beta2, float eps) -{ - size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) + size_t n, const int *t_ptr, const float *lr_ptr, + float beta1, float beta2, float eps) { - int current_t = *t; - if (current_t < 1) current_t = 1; - float current_lr = *lr; - - float beta1_t = 1.0f - powf(beta1, static_cast(current_t)); - float beta2_t = 1.0f - powf(beta2, static_cast(current_t)); - float step_size = current_lr * sqrtf(beta2_t) / beta1_t; + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + int current_t = *t_ptr; + if (current_t < 1) current_t = 1; + float current_lr = *lr_ptr; - float g = grad[i]; - float m_val = beta1 * m[i] + (1.0f - beta1) * g; - float v_val = beta2 * v[i] + (1.0f - beta2) * (g * g); + float beta1_t = 1.0f - powf(beta1, static_cast(current_t)); + float beta2_t = 1.0f - powf(beta2, static_cast(current_t)); + float step_size = current_lr * sqrtf(beta2_t) / beta1_t; - m[i] = m_val; - v[i] = v_val; + float g = grad[i]; + float m_val = beta1 * m[i] + (1.0f - beta1) * g; + float v_val = beta2 * v[i] + (1.0f - beta2) * (g * g); - param[i] -= step_size * m_val / (sqrtf(v_val) + eps); + m[i] = m_val; + v[i] = v_val; + param[i] -= step_size * m_val / (sqrtf(v_val) + eps); + } } -} -__global__ void AdamWStepKernel(float *param, const float *grad, float *m, float *v, - size_t n, const int *t, const float *lr, float weightDecay, - float beta1, float beta2, float eps) -{ - size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) + __global__ void AdamWStepKernel(float *param, const float *grad, float *m, float *v, + size_t n, const int *t_ptr, const float *lr_ptr, float weightDecay, + float beta1, float beta2, float eps) { - int current_t = *t; - float current_lr = *lr; - - float beta1_t = 1.0f - powf(beta1, static_cast(current_t)); - float beta2_t = 1.0f - powf(beta2, static_cast(current_t)); - float step_size = current_lr * sqrtf(beta2_t) / beta1_t; - - float g = grad[i]; - float p = param[i]; - - float m_val = beta1 * m[i] + (1.0f - beta1) * g; - float v_val = beta2 * v[i] + (1.0f - beta2) * (g * g); - - m[i] = m_val; - v[i] = v_val; + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) + { + int current_t = *t_ptr; + if (current_t < 1) current_t = 1; + float current_lr = *lr_ptr; + + float beta1_t = 1.0f - powf(beta1, static_cast(current_t)); + float beta2_t = 1.0f - powf(beta2, static_cast(current_t)); + float step_size = current_lr * sqrtf(beta2_t) / beta1_t; + + float g = grad[i]; + float p = param[i]; + float m_val = beta1 * m[i] + (1.0f - beta1) * g; + float v_val = beta2 * v[i] + (1.0f - beta2) * (g * g); + + m[i] = m_val; + v[i] = v_val; + p -= current_lr * weightDecay * p; + p -= step_size * m_val / (sqrtf(v_val) + eps); + param[i] = p; + } + } - p -= current_lr * weightDecay * p; - p -= step_size * m_val / (sqrtf(v_val) + eps); + void CUDABackend::AdamStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, + float beta1, float beta2, float eps) noexcept + { + if (!param || !grad || !m || !v || !t || !lr || n == 0) return; - param[i] = p; + // NO host copy -- t/lr stay as device pointers, dereferenced + // inside the kernel itself, so graph capture/replay re-reads the + // real, current value every time instead of baking in a + // one-time snapshot from whenever capture happened to run. + constexpr int blockSize = 256; + int numBlocks = static_cast((n + blockSize - 1) / blockSize); + AdamStepKernel<<>>(param, grad, m, v, n, t, lr, beta1, beta2, eps); + CHECK_CUDA_LAUNCH(); } -} - -void CUDABackend::AdamStep(float *param, const float *grad, float *m, float *v, - size_t n, const int *t, const float *lr, - float beta1, float beta2, float eps) noexcept -{ - constexpr int blockSize = 256; - int numBlocks = static_cast((n + blockSize - 1) / blockSize); - AdamStepKernel<<>>( - param, grad, m, v, n, t, lr, beta1, beta2, eps - ); -} + void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, + size_t n, const int *t, const float *lr, float weightDecay, + float beta1, float beta2, float eps) noexcept + { + if (!param || !grad || !m || !v || !t || !lr || n == 0) return; -void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, - size_t n, const int *t, const float *lr, float weightDecay, - float beta1, float beta2, float eps) noexcept -{ - constexpr int blockSize = 256; - int numBlocks = static_cast((n + blockSize - 1) / blockSize); + constexpr int blockSize = 256; + int numBlocks = static_cast((n + blockSize - 1) / blockSize); + AdamWStepKernel<<>>(param, grad, m, v, n, t, lr, weightDecay, beta1, beta2, eps); + CHECK_CUDA_LAUNCH(); + } - AdamWStepKernel<<>>( - param, grad, m, v, n, t, lr, weightDecay, beta1, beta2, eps - ); -} - __global__ void MultiplyIntoKernel(float *dst, const float *a, const float *b, size_t n) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - dst[i] = a[i] * b[i]; + if (i < n) dst[i] = (a && b) ? (a[i] * b[i]) : 0.0f; } void CUDABackend::MultiplyInto(float *dst, const float *a, const float *b, size_t n) noexcept { + if (!dst || n == 0) return; constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); MultiplyIntoKernel<<>>(dst, a, b, n); + CHECK_CUDA_LAUNCH(); } __global__ void FillKernel(float *buf, size_t n, float value) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; - if (i < n) - buf[i] = value; + if (i < n) buf[i] = value; } void CUDABackend::Fill(float *buf, size_t n, float value) noexcept { + if (!buf || n == 0) return; constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((n + BLOCK_SIZE - 1) / BLOCK_SIZE); FillKernel<<>>(buf, n, value); + CHECK_CUDA_LAUNCH(); } - // buf[c*spatialSize + s] += bias[c] for all c,s -- convolutional bias- - // add. NOT batch-aware, matching Im2Col/Col2Im's own per-item contract - // (caller loops over batch, offsetting buf each time). __global__ void AddBiasPerChannelKernel(float *buf, const float *bias, size_t channels, size_t spatialSize) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; @@ -665,33 +620,26 @@ void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, if (i < total) { size_t c = i / spatialSize; - buf[i] += bias[c]; + buf[i] += bias ? bias[c] : 0.0f; } } void CUDABackend::AddBiasPerChannel(float *buf, const float *bias, size_t channels, size_t spatialSize) noexcept { + if (!buf || !bias) return; constexpr int BLOCK_SIZE = 256; size_t total = channels * spatialSize; const int blocks = static_cast((total + BLOCK_SIZE - 1) / BLOCK_SIZE); AddBiasPerChannelKernel<<>>(buf, bias, channels, spatialSize); + CHECK_CUDA_LAUNCH(); } - // dst[row][batch][:] = src[batch][row][:] -- one thread per output - // (dst-indexed) element, decomposing the flat index into - // (row, batch, col) to compute the corresponding source offset. - // Verified against CPUBackend's identical-purpose loop, same test case - // (batchSize=2, rows=3, cols=2), exact match. NOTE: 4 integer div/mod - // ops per thread -- correct but not optimal; a future optimization - // could use a 3D launch grid to let hardware indexing do this instead, - // same "port first, optimize" position as everything else tonight. __global__ void RepackForBatchedGemmKernel(float *dst, const float *src, size_t batchSize, size_t rows, size_t cols) { size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; size_t total = rows * batchSize * cols; - if (i >= total) - return; + if (i >= total) return; size_t row = i / (batchSize * cols); size_t rem = i % (batchSize * cols); @@ -705,17 +653,14 @@ void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, void CUDABackend::RepackForBatchedGemm(float *dst, const float *src, size_t batchSize, size_t rows, size_t cols) noexcept { + if (!dst || !src) return; constexpr int BLOCK_SIZE = 256; size_t total = rows * batchSize * cols; const int blocks = static_cast((total + BLOCK_SIZE - 1) / BLOCK_SIZE); RepackForBatchedGemmKernel<<>>(dst, src, batchSize, rows, cols); + CHECK_CUDA_LAUNCH(); } - // Im2Col/Col2Im: one thread per (row, col) element of the - // [channels*kH*kW, outH*outW] column matrix, matching Deep::Im2Col/ - // Deep::Col2Im's exact CPU semantics (NCHW, zero for out-of-bounds - // positions). Both NOT batch-aware -- caller loops over batch, offsetting - // input/columns each time, same contract as the CPU free functions. __global__ void Im2ColKernel(const float *input, int channels, int height, int width, int kH, int kW, int strideH, int strideW, int padH, int padW, int outH, int outW, float *columns) @@ -724,8 +669,7 @@ void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, size_t colCols = (size_t)outH * outW; size_t colRows = (size_t)channels * kH * kW; size_t total = colRows * colCols; - if (i >= total) - return; + if (i >= total) return; size_t row = i / colCols; size_t col = i % colCols; @@ -750,25 +694,17 @@ void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, int kernelH, int kernelW, int strideH, int strideW, int padH, int padW, float *columns) noexcept { + if (!input || !columns) return; int outH = ConvOutDim(height, kernelH, strideH, padH); int outW = ConvOutDim(width, kernelW, strideW, padW); size_t total = (size_t)channels * kernelH * kernelW * outH * outW; constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((total + BLOCK_SIZE - 1) / BLOCK_SIZE); - Im2ColKernel<<>>( - input, channels, height, width, - kernelH, kernelW, strideH, strideW, padH, padW, - outH, outW, columns); - } - - // Col2Im is the adjoint: SCATTERS, accumulating via atomicAdd since - // overlapping receptive fields (stride < kernel size) mean multiple - // (row,col) source elements can map to the same destination pixel -- - // a plain, non-atomic write would race. Matches Deep::Col2Im's own - // contract exactly: ACCUMULATES into outputImage (does not zero it - // first); caller must Zero() the destination first if a fresh result - // is wanted. + Im2ColKernel<<>>(input, channels, height, width, kernelH, kernelW, strideH, strideW, padH, padW, outH, outW, columns); + CHECK_CUDA_LAUNCH(); + } + __global__ void Col2ImKernel(const float *columns, int channels, int height, int width, int kH, int kW, int strideH, int strideW, int padH, int padW, int outH, int outW, float *outputImage) @@ -777,8 +713,7 @@ void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, size_t colCols = (size_t)outH * outW; size_t colRows = (size_t)channels * kH * kW; size_t total = colRows * colCols; - if (i >= total) - return; + if (i >= total) return; size_t row = i / colCols; size_t col = i % colCols; @@ -804,16 +739,15 @@ void CUDABackend::AdamWStep(float *param, const float *grad, float *m, float *v, int kernelH, int kernelW, int strideH, int strideW, int padH, int padW, float *outputImage) noexcept { + if (!columns || !outputImage) return; int outH = ConvOutDim(height, kernelH, strideH, padH); int outW = ConvOutDim(width, kernelW, strideW, padW); size_t total = (size_t)channels * kernelH * kernelW * outH * outW; constexpr int BLOCK_SIZE = 256; const int blocks = static_cast((total + BLOCK_SIZE - 1) / BLOCK_SIZE); - Col2ImKernel<<>>( - columns, channels, height, width, - kernelH, kernelW, strideH, strideW, padH, padW, - outH, outW, outputImage); + Col2ImKernel<<>>(columns, channels, height, width, kernelH, kernelW, strideH, strideW, padH, padW, outH, outW, outputImage); + CHECK_CUDA_LAUNCH(); } }