From 6dabe6768a0c35dc6e544d6db190c9c53ef41eaa Mon Sep 17 00:00:00 2001 From: Jan Svabenik Date: Thu, 19 Mar 2026 08:39:32 +0100 Subject: [PATCH] feat: add demod validation and GPU mode telemetry --- internal/demod/gpudemod/build/kernels.obj | Bin 68626 -> 81894 bytes internal/demod/gpudemod/gpudemod.go | 47 ++++++++++++++++ internal/demod/gpudemod/gpudemod_stub.go | 1 + internal/demod/gpudemod/kernels.cu | 53 ++++++++++++++++++ internal/demod/gpudemod/validation_extra.go | 42 ++++++++++++++ .../demod/gpudemod/validation_extra_test.go | 22 ++++++++ internal/recorder/demod.go | 3 + 7 files changed, 168 insertions(+) create mode 100644 internal/demod/gpudemod/validation_extra.go create mode 100644 internal/demod/gpudemod/validation_extra_test.go diff --git a/internal/demod/gpudemod/build/kernels.obj b/internal/demod/gpudemod/build/kernels.obj index 808846a4be1faaafad92091d5fa740c9cc5dcafd..22daf7a2f23fc935da5716db02e05cdd48dac60f 100644 GIT binary patch delta 18728 zcmb812|U!>|NqbDW9&Oahz!P(WF1PIrA=v}RN5(ZD_w~eZAQCxoJw_Dq~%76OxhPh zDyiF|eL;)%Cvmc)-Qd?^)NJ!bdj2kMuI(jmjk4MpgpzF=+k|{kkS~ezbZ^v605>j0 z=)ltC-L2g3nj>@^TI+YRhzc+w6cpy%bwG$Xv}*`c5mEyfybqyK`;tGlQgP-Gny?R> zad!Or{n($gQ@Oeyp=bMX6lbkc3#gS%(Ks6+*KB+W)DD0Hm^uP*2EzovO$>Jeo?>_o zuqqq(Q?gch1!#~%i<##j{_!-IgA7~TS`JwSW-256m2i#g>Y)PrCv4YJNb&|Hv9gH{6O zGCToT#qbrNfsDqJ{c0|Q9VMgjy#b?%$694J;941OrDCnJ6)>Nv7XTm2@EbK7WgxYK zG}PcALQV|50D~9~1DwV%9&iK0bifk~&jMC3d;uuPqs4Uc5ON^cN`nl}9h!rgCk$|Q z9v$Z*z)Xe*0q-z;45)DkZ&kBa5gtOQ+aX%8AK(OrF@P%%(K)RLJaPzkQMXn(1Ney{ zcNli@VXFE9PCra5NC7u8+yQu!VF_R*!5rkSZ^Z@M3a0uYkBlwbrwaOg8 zjYMsw;Q;g70nO4QbW-JjUytA(n%2DXQ9Mc0PDT4DLZ(OYdQjRi#Rrr?rc7kYR8ZzI zI=vG1e|2ZEOC?X{@0|13sU|URpL>J{|_`eCRMAFQM&B z)*hKp6XxX8nJ$EgHB8yZlsr%hnDUS*)u8;#6#W94WmKD4k12mKWd|tPOeteZIVhD(sb`965kiJV zw6a?f?aLdK?o1iWlt@r!GG!f8wt$ksloF;~2IVeOJ~M?YMo7DuR&G;FD|>+A!<3Os z83)Q#rmSH~DkwXcQoxk+pp-M^U#5HoMfEJLY9`k%4kp`nX-&2$)Idu z%5kO?f^wNDubJ`z6t09;wk)BQ?Mo2yU`jAkhJrGVDGQmh0+dvyg z%mQUGQ+6gCB^)G-G4OjE;_(5;1-3 zH0)<<+^k(-XQI*Q$7}43@G@H`eRtjtNPPJKbO>4_$?MhoijqpE5kBgcy zZt|39e4up;&Bo?941G0gC)o=zlx&*vwza5Ph#Q22+UnrGww*{QeqgJGGupI}J+%GE zH|zTUV_u@xzBtoX)8Lo6O~WvepI71Mx&DuH+vq6LaCaRYJ2ht7=$I&c&e68nyyR1F zrZvvhjm!~0wAIA%@WBTD+I^O^#nEq6 z@M@>7INDQyuQ+-2g13u5ZSd9sDItTORD|Hli@2t`fQ9@+RYL)h!yD3VRcD9$BD&I3%TF4hUp(<4GZ7)@brNgzBxD zC=3g!-r^U%lIoVf=;2DFe#110oAr+tQHSW4gh!zBNQ4sPD~I4zCYv9$m5?cGGB{P{2r_jifH7H)dD%q5KX<*|$F5u1eU_?Tb_dS>)&z=&QlLWt+f7 z4UyTAvxw*oM?c|)S(6yDGHUXehE*auIm)%=dczSA-8+EgQma=e9O@XYq^F2-$?YUAJD2TfTJA*fJN_g53CuL+1SFeP34e2F)tc%3cl(OdebXFIs#qa`k zR$73-uNEUvm!E?2HqXPD65ziEzX1HuM#>~#6r;_PV*nvv_>9;e+0KCIVK4w5>Yvj^ zHZj@)atvt00i=We(M3yQ5leWEU;5c;@v~u`84qlcu&n7i*w@!cT0sRkovX-zl$F@N@yy)-7;w?+xw(nUR&!NUaNM^OKTF5=_Vkt^9ZWPE2WQMXxe zK}`dO;|l^@P|sO(EuhU3L@~TFTTwq)As-})BMXRp;ED=uvO@9!lA?6QlxaMjC9JR? z)MH!20F)4|r&u6-&CdWY?rAQC9jy$P5NUmWf+`t1zkp|RBcU)?G8ecz!p2Hz#vtW- z(yVw#=gKGYb>Zj+9G~qBAVsO2CsgOa+Qo=}J*gIF!=v^=mq0R5vNhbS55jKdaJZ+T#vb|IFHMOSi5_O99w2}3vlBnaXo0gtYWV0d z6aJL}zBa6jx|aM>UW3hsTegj6WutwWRlI(aN3E4BWH{z{yU4KZGv zG4O4YKol}ivO5~`NFlOa`%Qk3JU@*_v3v+^$j7cD)vXUgo9sxo(DFbqH4X0`1^8ix zF*$b3uKq5xCH`S)jGaV!Z7bv<{$QJKb@ID+!^nw`9Yre8IqCoBjE4JqB5}$$mT}}a z8Hw_{zx=qQIqbL_0F(Q#KEBD{Bg8f7gB^D-f0Hpy#D>x|gB^EkVdwl;AFr%fM$=KT z@t>|&rmb{DpuZ#)ujHYWvGw|n8rczs0j^024$w35r!!SWEH<-at znG`2i#%3E_2g#N|5@90--h5~!kVcb@T!jBCUz*^~q7op0Y@`6}D(wW)XtJgJrU3Gk z2a)Qqk$l*3Y6+y#Wb*{ae+)nh&uv*$!&`7tz!KikCNxPVSoPlpP=X%|Opt#@Ao-&I zSiv)(lK~Ka<3_IIeOS|^;WDg27MWc0Z-v06+ZYcG6Y(`I2wLk&SZsrBI0_Uq>q9fn zLxv*ROpoJ&AT39J&&bfkZNlxOv0YfH0A>`g4tXx$!(uysAAw#BJgCze^@l!Wy=2>0+o{@5Xj`<`z|AS$8hf^Nd{$Z(q1_>} z=D_-_tnax`lk#R>7#1sApPBXCK62Xyv$;3oI}VLq<0;K@vyc4lkfU^Qex6^$~8~a=@%qh(b6@Lqrgy))< zF1HqR*0R3xCdg?iV@p}8_^AFWKI=1h4JKgzocJP%h zF$c^;HO2ovx-jQe&Hh1Z%hS~Tl_H%pTseme25njd+Dg$J2W;;g?O%3syH}*v=e7~m zAxc>%C(IIFm*fwNJ$z{Jp4thM^=!Vr*)>PA;~?jYhBhxnMW6h3>U8iiyEJU1^Q3*} zWc!a7?A87w>r86q%0Jz0yStm~ebNs2SNqt`#A|1k?(rWSQ(^Ox)(7ewH@TJ@4NWZa_w-tEAV}-IL+G{Ylx|^Z&odoD zWA7L&Xmz7++1#ZnH+b#BD?ue^s-5Q@p7rWY*pBe{4Qabmsy3a7>F_M<3}2?>BD?$Y zrC-&~uaP$frRtp;u-oB)-MQ@{N754moQ`Cn)86ktBx^=g$1gcQT!h&iV2@f~j$JPcTEsl7e&)!sJmpB_GWA*rV<K{EBS%&JS`gDBtX39*p@0N!O-@bYC z=G%;~_vZ9+PO=@`uPE1~%g0aOmg^3fQ*iOliLWN&UTrjuI)_*9%#a=^p7hi~`iMio~7V;T8%&hvwV9zJ21SYNQM@T>j% ztb&o-8 zo`JrpelLv_;S!~E4L!l5;PpEPALk6hs$4b42+z7Mk!ty$)f*epqnlVyzt@1EB!AV7 znQLaRzrJr+VCbB8j>*US#>e^}s?;CRd;Zd$XKE`Zg-33w&zgO-B-HcCC+Br(S7sM1 z=`j6JX;f#~V8NfWv(qnTer|tYUe<*t&jzHNd~V)j`r1pqKOS4yzQk~kv1f$m#Nw6f zN+$L>Fllm;Xs|_jthDVwX{}|Xc*gxMpX0qHc5`zKPQG3edfL(J?bz-ouBFz6nBC8O zy?9+}$MJ(QN+ymRYc{7pnmVhy!~J0Ux%!JY`R4vzuClLS`|^&1)hBk$-S_HxFw$DO z`A2D6|KOqbr&+YjAFj7qKQncS$&}m+mh%Vhc6rd|w1eQ-klQOdcm*5fNGG{@pQzRx zKKJd18{4WR4{Fz+JUzL^m-hOV%hJcMu(f`)^T4snePITiXOy{;X}6c^DT0~0_09t} z_@BKuU}f$FLH`vyk1jW?vU_m#)Roz}?Ut&oEj=*KroFl4?lv1UPu`pV!l!O}a&R@b z`>5`l$HK|aa|1k_f3%Nv5(ay2xaEqa103h1k8^Y#K5AdShj!Z4vsoffrH~VEX0Fge zT54ndNKJHTH!81NNW9C;Gj4&KM+Cl_;cI6VzfkL}dTQTuyq#nIB<2zB|0_$|Yp>bD;89n$*)weV)s>juom7Pq=_v~6cX`WJ^N&6NlOU~$yUiSJ; zyPR7daqrLc9#oXxz96SgddacHzHWh!QdZ0rn@qRL(Up3J^*m=8vDNba^Xs~&^%k{u zHQDGe*1Ud0?X*4ZECY+~rVYO7q>cU*#4Rn4e-amVMf|j7kMNR7>+U|SQ;MjOZGU&V zPkoQXBXjPw*=oLZ_QNykp%*7+e!kkhr?~a;@N)OL-g|3@8YLT4-OTl=<7f7tQ=h(6D!tdsF{uuy`mD{ zMoA=zX_eo5NrH51}O!Ki)Pb2BV z`~@48ai8qTY)kD5-eChFh;IZ(|(UmrzZ8Xi=9rIDU-q+yeH^FQ1 zc-Koky{o?{xtAC9*8OMXDqrP~vIM0KTPm(sZe90gkC9cY&*OC$_R~}1mZWXov(5OF zntuDildmfKl`A{mUM#J1e`hL9k@~pn;4MY9pV$9+cB;=OugBi4M?cmv**>^`+UPfy zD+`03tsnDs&y&@DyOtFhX$b1QYUW8&+s++-(zes0q3iq>ZjoW5aW@tR-Y-~?ZWX-c zv-!@2t4hb7tLkfFrCJp~+V;nS;^gdC>jVWwwid-+S=#qYC*BJRm2%&UmajV1aa%j- zk^D1>^DY-1YJYfYZ--~SwMCXCUy4G$IIVq?7khkom-CPIWKB7JyJPLGg8t_wts8%{ zb-K@_sU!M!!)Hzr8-bE<)`tUB09v&F`-quygbfG0w`rR_J{WlZyDWEbOzpvbG2Oo|)Jiuz$zK{H^|vZB-yNdq*7p#<(YjAd?}*@! zJ$7#1m-Z>>Y2}2;HM@`Rlm=U!U7vZ>DDKa372JdCYlih7le(#X`?ao9{wUlXbN*`9 zfaqNRjG!>HsiyuaI=jaA%ewp-4=+i7a&c@{QnU)%6#Zv>#)}@7ai*Ie-VL8WDQT)b z{=?6Xe|O(KRa2VC^;(zF!qm&P`vc{&Hm`?UwMo6vI{%!8q@tH{iOkLCZCQSDRQtG7 z--XkksXVOP(dwp0|JBRA`SUF&+pp7fe`~yE)r~&Zohm|}b?ldj?>rHlcABhO>s}rm zE<53`y85eAsnrV$!)4zG1bI$>?&M$G%hqwg?P`gJ`j>|_wZ*-p{c4qNU}@g1cX{Uj z4tad-NW@lq>++tX?B9wV7Fr9&+M0%Lz1iVfm&vcrS63Sjg}XoJq7xE%^w!(+_Z>qs zU47oWmpOlX)~S5d8^aBwa)U;qPnuc6t7|eJ`JOk_+Vv%-wr2XHVK2I`lnqEQ^9)s6 z-)EWa>cYUhyQ<%OZYzn_NP{}6ZGHCQsO`;<^X|JYJG3jJ{K)dMf$!6fbS@fLe(U%{ z$5t04`!84M7q-5bEps_pTXSk^(Y5wx0(yPC;nv6LZ~w~bt=>Jg^t!be>|iQ(>FWM| z#M$)S8GVMkz8pHyK2SU1EGzHj$P7;|~v>)D=bJGc4j#AVw^rOJ|m z&5NCsbFEB=&(xeJetfkiqs`48K4vxNQrE40VydOTpr`W18Ryrp6K1cy7jvvMInQF* z?n92!8b7J#(+1GhfI2;wf2Tzg)W0`b_bToWY5uncS!Qd_Y2@rgHjU3^Fh5)sM zSAR+e;5CDK?<+_ELvWy8A0YC35CQ521i!8VS_ssYYUJ4!%k8G1NFbKGS_(*ht|dSi zh$Y_#>P9sS_^{iZ(p;c0&;@8dkO$SufwF0Li{Nt<>m)}(H-T6uJbcn}qZz6|tdGBe zNQP8^uE|k=z5uZlTX^_K9^ML&BM?g&3)F*ZbwI2;F?=**6bQs}qkz04%C`< zybI`0s$~JOdIy01rd2KiWl*YaAm|^S`gfqovw&E^c%U{^`&XgWDKs_s9s!PVShtQq zY@WG5>^M3K6ibJ23W&|~ED)RKJ)pVN`-BqoFF-E<**t6HZe{cZh!x~?;P-$uMFoh> z(+X%7)f|9W!L~qbo&i8?o?{g>7l=*s0}!kCMF+kTVDp>;&yATl4Tx!rfY?0qf!I6; z!}C6NV2%VLXO94l0b=u<1jOc<0F+3*%M`R4h)r{?TvJBUt$?iHE}#W8B@>9vvlwVG z)h+_Df>(gpJYNB^d8+C+kO+v)b0ZL|w-ty@$~hfuwAn;{cPXI2ni) zoC?I|xe|!Y^KS*^1F?B}3mXb{0b=u%DYPRB?HmxB=2sv#%@`>>#AX+ZSwN|D2=PE{ zo=bq(JU0WSQSUAVWdgBz=EyZ=bPR~qD*)O=GtL6BX;uPlq1rPbQV*_wuL0RS)!=D2 zo2QL}IsvhH<^i#S$AH*8UEm2Z(juQ;IUk74(+j?;V3)PdKs)FV{DDT0BN(nj zfNY*&Kv~p1Q9)5aY@XBP8eE2eSi!}_OEZ=Ov3c$S%Ai^n5UY0ph)we%@v_U137r27 z$qyFTJZAy1g7H9XPX8*jI)$bNKUrY&bOXw!-E{=oLrc#Df`x?xGad*Q3}`VB8}u3= zUvdePqEbLuCvaZ{f}sK324Z{hQ9~fIga@|JH+b!UKn7|H#E$BAK(ru=1%m&;I~xr7 z${6mMK5%m@DNXG6c!oBW*U|M4>;3)#>mx*V5@->tC0{g$iXl%=h>9DHD@6}Oja zGnR?DBa}>in~IzGR`p1Vzz#YtlB7^NzRdWs8HxNz5KyEDXgS|!Evp-W{G01 zf)JNS+x?c)O1NmXn5)Jm%R6%AG+#zvbLs0doUy7cS3tRxW@)W);!e;iS7^Itg_tX% z?JD|un5uUY#W)Y4=H_ID(24P+9Z3`9C{1{;NWB-mL#`; z-vakZcHxfWxycqrfv{uQx7Uo|?GsAF`;#qjWvY0XdI`*Asv4Rfz1wg z*kEax%XD21>B@BnJb8m9?4=ngvkl?Hzahw+i2!~*p^x;V^}_I)6eC=dGDK?XNW_*@ zB=ye0pQ)<=wc;#Li)iC8qTPZ&PMt&Q>0kvM5{!ZC9B930uSE(XaZG!qAOtZ?6DkOt z$d0_gcY@3DGAzR;_%HY~T-yn^S*xXUh-XgnI)cZ+YmKF)fNX&nkp((MyCxo9aN zxI4=U?l(vb7by1ot`*b^CkL>y`c&k0pv=(VL-{$8zCpSZA<`kr|!u>k_m0OUozi-0ioCMswpP!>?T zK|6suQSTn0_0*fCpd6qs)O!#pm3ogT=md~2_2vVSi<=TEQqVadf9kzRF54TZ`yUV) zT?Zl`vET<9KpUy|o`NcXdQtCVpiR{KOhK=J`cQ8T&|lR1UO{z0{i*jG&}QO|q~CH; zL~sJ20o1Jqw1p;XDM$}!5cQJdom|-AAXJb!P$>0^fZ$y&c*P1LABl!3;qU9samEHA zT*SVux4?ISE)X&r0B&{;Fj9l`e{fi;1@_q}K zy{8npNyvZU@cd11!^C=iTsyXKV67pIMJ^m7M1|%Mj1JQ(user8* zo(B{&{0itu(RhraHX~BomAS?$T#Sgzixo)TEaW>eEClpt_+hh<3!o#g-y-DuF>&Y? zAs2#Y|7C$!01YEWNf%7!l8STGDE(keiHuY_-4xfnun(c&m_)W8EAE zF&?i@v%t?mT*Snd+l1V5?7Gd3PojcxtfD+4Qhq&)J+Mv4Z)W%iFrA@MnvmbauqWVt zhO+<6X*f1+h&2) z(}nzFnqoX(QH>F)_JX>M7ARbdh^vMyfbl|whY|67Wc3y)T#Sh88*?pIxEK)^uS58ds!a-{QSpX!=%1AOCsTgQ7odpRYMPvrwfV*B~`Pwl_VL z*ZX(1$B0R)4_E3Xuc!!5c{WvMxGP6VZK$x!v#_$PKYWc|fg>}L;HwxIE^P1ndSznO z%rJBE4-YFqaCkbrJ#nzDBbX#}yf8CJO5P=t?ZD!*iy4Ur_H_4hlhANgkl;$j6mJt~ zHB~%aU4K^G0)AH9!+usgq8k+{E$t~uXq1@Zx#4HUE9YmW-MOEY_7y)X-gQ4IUasU) z-Y^X{GFLA*5hqLDJKj+03%W8AA=uflJyAg@p_xK{OqZnAG{y{+ie^eXP}tT`X8A%FRA_{ps0tc$R`#rw|o%uz%{)_Sw6xU{*>A>*9ONY^tQeof#gdvc9|>NgA7B zgG3ZI!=_4D3p~2ml@HOt$;A?Wq6RJjGfxA52eVcKdz_WeACgykD!3L zCN@1M;k`6*e=z+u@nSF&HSuvU3pDXNFq<{8{dozWr-?^`xuA(LmFY}dlsU=C~HCt%97aLbD@ zCoLQd=Cc-F3Px8Op9EvCjsFGHMH{y+lkg$hcr=)a+ISrpsWvVJvsN2(mn3|)Hf{&z zyf*$5%mZz_15B+pz70lA2kTyj0qWrHVBmYdIbeF|-~(Vr=-_HFQ+2S#KN5b44ju?* zlMY@6CPxRK22-qqKY+QXgWFtznmRZfjG8V^1!JL$%fLuru-&+@Zg9+EgQ^3sB z#k;_)*2Qf`%h#_40jn-V@=ANK--^|2Jp-}*QY%qe~R63lgdY;_BY>Ej__;GNqlFf9ylAs8nE z{27eD0d~19;fEODabThh@LymS8Q?2mHXC5oI}%=IfIEXJGr+N6o*H2I?F!Es;)h^b z8e*fn626@w4g@p65YGoQ!4Mw?GtUse2D8x++uW1z2MzI1Fjoxm8ZfU7aWNPrA^r-+ zM2Ow)OL!L{{sT;RA>IOJv=Cnd6DP##4M4C(jS@It-*CwmMyjey>I6G0?Ie+ z{K!%|fc#n&?qH31*>a`%AR)cUAT=%FWNDV=4+{Ap&9X>z94O=?b7ST0pop5qoaZcY z*&~VY9(c$`){~vq8X|C*;I@y=g(f26LZ}S>m@`O(M?RMDF(Qn?+!o>UV1lf0J(xpQ z*z*Zo<*acem^s#XJD8W&_$HX%HdyyMoaa<5^(h#W)*Gg&02p<7bOoR!jJ8 zwm2A!wjEvyW|AE~SuJUKR{kvP0Ue0*nAlh{EV4bs&W~HRgrDKC7Q8B`BX( z8|40qU;LH02K;q$f8{U!Dq4St12|x-g1_eVA5s4ax&P5G{>QX`x7H+^9Zb#hpV0o- z$^B0trn&zqsV}UP`=3%jsnp34D?hWA^rP>*h=Aa&uJhOlk+EJ9+Ny?EC64@K>gL;Y zN7?Xar796W|Kr(aPrE0FPmdZqVRX}n;bW&qgtLEf`@2Wy@}wqindq$pkH6ID7`S`< z++S)Ro^!?_H3oLf)g(7Pje3%KQSUARbPoSR_z!<7TbKQn{i0S{ILj1EE&;DWFN}OXkc1ewn6tdeqE`Q4uWp*MXs?wVSv}?PmW% zJMD*Q({`SWb5~03pv3)dJzP?6@SKVMf38))|FTX3vQDWYo-_ZSvku=>;tEt#XRC27 z{(B&8Q~gvqhu`^jt8%Wt^BJmf*1z+`sBxD6^H@`TR5{20C5!8us-nVq{x2zPDfQJj M@&BH+T9aG!f4y$$#{d8T delta 13501 zcmZ{r2|SeB|Hq$aW-MbJ3?W-gQrSaDQYp02u1#p8QldrC9ZIwcCmy6lt5Q*tNV`%g z6_pk$ZTGg^E7fi9_Wz!9#-RSc=k+q5_w)Il@AoXv%ri3wFMbNrvjpnXlO8c6_MQlz zvyWk_Einkd%5nRgO0 zWG#8x$!Gv;sW>8;VWuaed91nOJiyfy(*X}8qf_7>1FT6#DFO=xpm#L>9Z+vI$ue6F z#VL9N2CYUT_~wdXfGO0?0z5*o0`LyS$ADic3f3@;$r_So1vp|2sT&9wMKKOAlj3&3 za*CCJ4=Mfy$XiR&)YdYLB}IF{Ar!{|&Ra{mhymP6-F(1v7)?|hVPUSq>3_75@;ZjG zT}Rrx0tQhG158;*+T-h$1*a7>auM(ij+iTc1Qe}DCW_{YhU*!|V?A*P0|rx^4VXwV z1u&oD5x}by?*P7~_!UsRfmATrz%Z^92LMi{I1@0Q;#$C67)@00=nlbYEsZn)e%wH2 z!=%9Sq!6?R97i!2FewFXRW?^l1w4~Ndb|So3{cTR0Z4fxnUnTLhH=`6JcZ_po`BI5 zmjfQ5SOM6wkrem_sGCaAG?ii8C=LXiLU9&g0>yQJyD1g|UZz+F_;)I*Q88Ei1gM)v zre_7yGfiU{pENR&k$_9m&|p<_UP2lgqTWpbjEpq2Slxn6Lpk)kDh*vE=dbAbmo%iO zVZo*&ca3fe=?v2|9nB=?Yw7uvbd;%T;?@9N#X@*I0)RsRKVFQzmu=+qkcA<{A`j2t zQ3`FQZK9qVvVohxmP|BT)lzYHCc_-ggnH(@>(r?S=P7kqjYSv`1#oUsrKvv{88{cH^OQQT!TCfT-CR=HIG16%=aR}k)bR&r5_OhR zCjp!d)G44&F*v8G(@34?;IvRja|@~5X$!+xY$28VQ^yaSvDAsCPAoXd)Y(m)gW!}= zr=B{GzNbYyw2f4DrcPgQ{HQaFI`hGarA{_=^1wMrooedTf%AwuKd7Uyonf@Llgf76NngFe z8AzQF>V$){h&qTmo59&joeR{t2F^X|e4@@za8&b1P!P?K6N%w zCj*>3>YS!d6*zU&X`#*+a1?f+Oj8rZ00u7Rj^c7tRX$RXXexeIgcswEVnd0P)eiH4 zARIh~WxYUSJWxT`ai&>)unBKLPE-X$8AUI~6>d*nj1j}4{yl?`lhshvWND42bW=iK zZPZYhwI0zPRw~7(teSYF0ngeCrCXV!``xtVL0els%i0$^+Lg1Y@t_&m)wAfmkpdzg4;W8l?(ws1E0NqISvrt@D~@J zJbc_<_cwn6@%4W5{~&&AuWjKoeCVIgsI@sVY9I;fzxl7h*A}&1vUnIEKE4sZ-3q}9 z;6ns`YKz)dB3_mH-H4AThGLy2B4?LiuZ8J(Oi_TGs~qc>Qn=Zr{3=0_~10YcFPS zKwcCW27w;CcVUGnsm~-drH>V>3KYv~pcHo>fr<*p0M!5EQ`{&(Z{O-M<6Z&2dM=rv*ioj1Cm*^z(I!>se*pG=%F|> z5ps*x=M|};(b1l~G3p3KdkPoIELF`0kD#hJb_3Uo%|wXm0xqdUaJ`nEzi!k?eK2uw5VY>X!bG42kRx)Vh|q zM{;#-8vaLywJf9cWhl=;Ut$cTviF*JNX7?VYPf-@$K#WBhq7$S@Xp)eh&Ps<$|ivG zVHK3)JG6q&;Qu3E8gD{dw%$&00az66Ukl(DHtyrs0Wdve`6^U#KHQXa;_0dUk-p>7 z3Cc};8t;VAGFx6&AGj;%p?k~Bc*jK;%`aKdcCFgF;;o~AM4aIlYu1)*C}Pu%b}+-Itp>-MIv)W+0h?&BAzjj!7k9wZc7hfVx> zK0Gb_VnmF=^BjrtfWx_~vWukl0=+%i_h&59T=BK%2;|+XJa?5r@Nnej9PhY4G`_#a z9)pmX=yEvvH@`@wrKs@0{mF^fF3cYhyBReYt-p4;^xD$c=<){*Ln21Ja_KZ9I(^FK z*hfa+UzC@){S;pFivHjd5U^Ce`s32s-%MjHyZq}Ojd_b&%~2A-<_q&j(cj}cC$-$SMRBP+wFCSD6T^ROU&!ur1EfBdud#n>}~ z`EiZI-|igZoD=A@b&QkAOY>bh6BcaR(5N{2djEpjYUTR0#EJTXp|_^HXY;3rh?iaG z_ZJF|hI!@$sjpbo|I(bbuH#)oR9D{Jq7iMCX)$Pcbmti&<&-fq(4Ml=5na2#nlWYU z%djxpA=1%D^c0(nthCG2YCRYO$$p{#v`sztSqcgyzVA|pM~CQb_4IyewDq!2SC_e* zamkmE%{y|08WkhEXY0(+5S*dxOg&Ek69Q z_~ATzlRJ%n`bo4;+}KdQwsL2@btkpY|M_{`daqzmSUpPmYPapN6F(YueY|M(H;Xy^5L)Vbsu&F8&)zwf~- z?~@gn7r$_W=jL(s7Z0qRURJ!|-N>^>Cl!uJQuTM;srWNxn5>vU8Tb7JHWZS%HP`ZkvUyHH!kPSC1boH8cczF!I zyYy#NbLWQjJ}qAx`u@1zWxbYdg>A2x3=7q%JA3Irdfc3RQ)FxSa75o-joG_Z_thj< zCq=Oa{&C40^J2u;9?!4$VaFQy_w+oQx=!I!-Gw1#VP*cId3P16lKs29m%dkTTDNxZ zgS_;HQN#7p9(;YW&QLX`Ptfhg*cVKK_}r(Uhb!aq&8JOSvwHn%iGrWXd3KbFD3Kqk zUoz$dtG(SkT))z4X6``6Yy$%=tu(R3FYnp|Et9=nHdyr+%wCf_BIn@VJL*n1S1)*I z@11Cs$#@ZFq<9;lrWwLOQ=&R@&w71wISXyWbj965A3ydVG%)VG}g>yc@q+8fBD! z)S;kK(ffJ*kAlmtNv!YU#m-tM#}tekc||j->d{LvDvT)!r|SC34S zPSIKQM1NP$b4E>ty1ctz9r(4DCkNm2?maH^$IGgBPs8eEE8iU)~CdY zO}`{MWjo>Y(#*iaha-ZSPt(75>Sg!$m_HnRUuB7tV*?e3q*_l~-mk*YvFL8!u3Cbu zb=s4H471eIt(pgOT@US9sT?`3C_HcPr=$VF4>mfSdfsD%#k1#96}Oi6DUQqgOH#Hx z+4L<=FV5z>%Gz?#{7h{};e4)W=K$Z#sVQ%drycC? zr7T%+-T(2{<%&0>uAJDRA@vH{Hy7j&~VS_F<6G zdF9>zN+y`qE3KTz-d8@qawYH*;VVM-)2PmOh5a1 zP0W~#G6yYDf?mVytCnBl5(@oQ#>Q#CmG-@5k+kY!Hhb_*e(oZzahL3BFPUv!RO_p? zdt)O1>*|P}4z~x@SZ5z|=gr7?++>`acyYn!eQ3nmgS{IfzG|rIE?MQ&#bp3r`1R)G zJ6c`MO|PpAR(G2^h`-;1{_cvR%6yn%FSHY3maMuKe?^_xm*JFOv*& z?*)5aMpcK{MSKbWuUih&V4}(%H;(d< z^z+=h$WXEGRne$D_dPTYdA3O0^@a5h`z19;EzWnh+s2<7&rVT(CN1q-@$uQ=7rgpm zKAT^3jf@L_yw05$w@7dHrrS#VO=lM$fB58pb3@sH6CYCUO}VGiHKO;5oNx3S6V5;26}^yQHuGGxw!Ontkuj2~8;< zn$M)|Tz`3JWMEg1c^}s;>3s8qkyO~79r*cm{w%GJb<-XzxaS?W&h@x8Cx!2QKlfbL zhA)9laY>+si! z+L3$59ke-WGgKRY=sY|ybTIh}Tzoirh*EKN+6|TB2YK^Z^ls-!G<280v_E`?7BGFG z>1+5AFjeqq37>ieOlNp>Cv+GdcavzNoCNR@9e?*1F#Uk62zdkHkKO`iB~TyY=|g>* z>n*1_Kr}Z+3fO}*L_jpT5y*ph#_&b3FQGV~0pJUm1Ry)&-3BTk-6bltc2Xp#dLY^f z557j&kPIat+Q(iXJS?ezsgt3A`3gi+y21{ZHA(3KL{o+X4I)FB1vHq@T%ca0U9{Yb z0dgVUVW4HC-AwpCLDwh(2!9n8Fi}9Bc>b!)GQc=urUH#3lr5*NKs(4hcgZ{j<`57q zSOPSRWRwGSB^}=e@+V#c5Utk)w3Af%4ut0}U^)q58;_zre7AQb$@_q4!9zfJ4^hBa zsI+=kK-6;u@*>eWKy;o@fav9X2{e`r;U6GRAXVlg5S^wPY&`_wsFcxdMaB?_&eKfh zDKI^NXh8>{KvK{Jh|Y5)&?Mpo0+E8u6ri1Cb}>M7a!5`Gfao;4!tP-pALcJ$dH~Wn zZ6hZ3cFDb?Ky;pOf#^JwV9SU;KQ;h`kRhZ4(RtG7zm-14O4e5O$5I zC-nuSW*`uqXDSe#rzLE#(3g%4Py`u*BM_aZI}n{`0MKF*ogk+mAUeF z0*b`jxB_MwAf0ChP&6r%3q%X%0nvF@0MU8Y%jqo;oo6^~Izd6GHxGzT>8afNOYVIJ zqSNdGcSz|{lF}1!7g@|%KzLIX$^hX#9s#omh>ltc-%w-^Y z5$ojbG-3A_@5u-lJs`YuEMSa*=##pe6p$92288!#1k7w8BSH&+=&L+d-fkxl-iHw| z`+@L=ihwx`MDLC#fO?Yw)&s%w6`~D5PQ-f(L~mTLfIPvo6EGj;CJUQjG^_~ZOj6W= z+z4p{<)M)WjrcReqtKay#^~>ZKiNIV;iv_B5NQ=!um_NHVHb8UcF@S<614WPgsmg- z5+p6LV{a3_MdGS zfu3-si=@(1H1tRx_9k%($>SkZSmKNvORU(FgpQ0g(->uF?DV#7Wnp8ZiF9x;1^^zfKGUPhY7+}JCm##z!TtJs=7 zfleOtWdB6p;qf_nd_iL83MJ@kkp!KAxx0gIf&ObqWbQF-lzdzdR&wic2etyG9n(Qi zj~lV2>1D>!{%{%ia6RGI)r8iDVD80fO^YQDB+YH=u(-n?jU)D5{~4fuVu?nZu$6LWYBPf zUsfiU+bJ(L|s%B)VEo>w&CDG!+Qn0{F}(Ib{L$z#m@tOfKMB z65cMST|m|(nh&&&L=VUbUsVSZJqom*M32j<0?3I(&j4*8(Mma80>ZzF;4@bl7kUrl zGdIA*cQZb73&@>>?*hU57jM@^@VbO)2>0x&2k106+=PaEs{$Rm_+q?|;2 zy3h72+2T0NVp6rfY588ahQBU<#j3Hx|4bM;1;J}+kHp^6K}=sB1RsoAMg%v(m$4=|o$EMO8rz1i~8lyK?wG&DyZqJ%@~w1CzH zF)xeaK)|gO7Xj{~xEF9g#d{aTIz_)vFa4q^dkmRf#P1BSD!S3~-YDVT&XGChT@+<&xM(L|&3V;lN+#x%adi&`feJc>Do!wlV@R<_)@EMuo zEc(@PG5gi&?fR?Z8t|**7S`tUBc1iOk8iW6fkpp zILsJWU^xDQ@4sHqPyrmcU~s77G;*}IvG~pL{LKmc&6)k1lK@Vz6du|-+XZK-GKUW6 z8=M8emqy%%Iz!<}ctfqdgTEa|!i~PI%^Pq;@BzH7a%cD=+i@D!#3}ywrHw*gHyolJ z<6k*ShK(Q&?YO-|J8thpYNG09d!Dxt{b+XLO%S49Z=HA%LKF^?C`8*pa)jtMNWKuM zwm9+N6K{WzN+DVTQY%CUTbx{*WWyT;QZLIT|}GegnhrR4WgCI}9Q_nuW;V zor(T4IK|_AYsjqNN}+R8b+w zGF9{(Buy1{{^-QptBOW~oK!{0AhoLKB*;ru^d97=D(e2piPuRDO$D)4Lm436YN#4y zf*N8!JMreLq23^gYA6C^n;O~)Ql^IPfz+rW%`Z;8Cu(RA$Y(Va3!<)$4uhDhqd!4h z)RFO5C*Cl1G#VsC9jyUbtd35DtW`(*@P%+3^4fGP^t_Cvw&x!Ys z1{w>Zpo!Ljbk;;?L0mM^Cy)S5Wb@OB7p94(fk-t`Hb|x>x(;$cGed!er;BDrACL!{ z84Ez(Yi8^PQ4wV{f>?++9Tv2Yi1WibR>VoM&J%GZSXYX;*H}3b*M$dqNW_iDS}Ec- zV0|Fs&SPy6abK`1X>oRZP$MmF23AKcZZlRNE$$}PAT3Tw0J>0%>x*^078i|mrxv#l z>q#x{5!O3eoGu=FixxK&tEx7afVG=8SBiCjHupEyKyA)K5i~-Zn}BtVHn$P$E^Y1t z*0b8&H>{1?oP!eRdu?tOR%IP-3szGdt`4iS4kuIw_1EG0VGYya7GYhk!xdmn)8U%1 z?$zNs2|+7#xZzl9b-0yS|J32iuzuI!nz4$+TvruPOEEVIYacO}hIO=4kB zCcqm7KCZwIy=ijbz%7U&FJ2rXI|^{qV#dG^Z4(B23^=i^O)@z89op;$$EZV_%i!P- zIc;^`NWrq|&`1?-1NI#nb(cA9oec-aqeGiWnbX#01311N+Ux_zyF*`R_24Z8kEU%z zccIa+4voHo6WpOqXV^Kx8Eu`#gEO~7n`hum?$CyX<%#Lg##kSkgmq}t4;m$QXcP?2 z!VYcLfs@*yO$j)09omRt(AUXwwOvI(NZoYtzq&Tdv9KbHg=RBk?E*;{z9_!;dFnCvl1_a!nq&XvAI6gs1#p zM%*{7{>Gex2=t;cHw&w^3AY7nnh95jwZ()JYJrBBa{aKLG36FvH8JA~uqK#sO;|se zah>>(in, out, n_out, factor); return (int)cudaGetLastError(); } + +extern "C" __global__ void gpud_am_envelope_kernel( + const float2* __restrict__ in, + float* __restrict__ out, + int n +) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= n) return; + float2 v = in[idx]; + out[idx] = sqrtf(v.x * v.x + v.y * v.y); +} + +extern "C" int gpud_launch_am_envelope_cuda( + const float2* in, + float* out, + int n +) { + if (n <= 0) return 0; + const int block = 256; + const int grid = (n + block - 1) / block; + gpud_am_envelope_kernel<<>>(in, out, n); + return (int)cudaGetLastError(); +} + +extern "C" __global__ void gpud_ssb_product_kernel( + const float2* __restrict__ in, + float* __restrict__ out, + int n, + double phase_inc, + double phase_start +) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= n) return; + double phase = phase_start + phase_inc * (double)idx; + float si, co; + sincosf((float)phase, &si, &co); + float2 v = in[idx]; + out[idx] = v.x * co - v.y * si; +} + +extern "C" int gpud_launch_ssb_product_cuda( + const float2* in, + float* out, + int n, + double phase_inc, + double phase_start +) { + if (n <= 0) return 0; + const int block = 256; + const int grid = (n + block - 1) / block; + gpud_ssb_product_kernel<<>>(in, out, n, phase_inc, phase_start); + return (int)cudaGetLastError(); +} diff --git a/internal/demod/gpudemod/validation_extra.go b/internal/demod/gpudemod/validation_extra.go new file mode 100644 index 0000000..5c164f8 --- /dev/null +++ b/internal/demod/gpudemod/validation_extra.go @@ -0,0 +1,42 @@ +//go:build cufft + +package gpudemod + +import ( + "math/cmplx" + + "sdr-visual-suite/internal/dsp" +) + +func ValidateFIR(iq []complex64, taps []float32, filtered []complex64, tol float64) bool { + if len(iq) != len(filtered) { + return false + } + ftaps := make([]float64, len(taps)) + for i, v := range taps { + ftaps[i] = float64(v) + } + ref := dsp.ApplyFIR(iq, ftaps) + for i := range ref { + if cmplx.Abs(complex128(ref[i]-filtered[i])) > tol { + return false + } + } + return true +} + +func ValidateDecimate(iq []complex64, factor int, decimated []complex64, tol float64) bool { + if factor <= 0 { + return false + } + ref := dsp.Decimate(iq, factor) + if len(ref) != len(decimated) { + return false + } + for i := range ref { + if cmplx.Abs(complex128(ref[i]-decimated[i])) > tol { + return false + } + } + return true +} diff --git a/internal/demod/gpudemod/validation_extra_test.go b/internal/demod/gpudemod/validation_extra_test.go new file mode 100644 index 0000000..e192acd --- /dev/null +++ b/internal/demod/gpudemod/validation_extra_test.go @@ -0,0 +1,22 @@ +//go:build cufft + +package gpudemod + +import "testing" + +func TestValidateDecimateRejectsBadLength(t *testing.T) { + iq := []complex64{1 + 0i, 2 + 0i, 3 + 0i, 4 + 0i} + out := []complex64{1 + 0i} + if ValidateDecimate(iq, 2, out, 1e-6) { + t.Fatal("expected decimate validation to fail on bad length") + } +} + +func TestValidateFIRRejectsBadLength(t *testing.T) { + iq := []complex64{1 + 0i, 2 + 0i} + taps := []float32{1} + out := []complex64{1 + 0i} + if ValidateFIR(iq, taps, out, 1e-6) { + t.Fatal("expected FIR validation to fail on bad length") + } +} diff --git a/internal/recorder/demod.go b/internal/recorder/demod.go index b3caaa6..c937eb8 100644 --- a/internal/recorder/demod.go +++ b/internal/recorder/demod.go @@ -54,6 +54,9 @@ func (m *Manager) demodAndWrite(dir string, ev detector.Event, iq []complex64, f if m.gpuDemod.LastShiftUsedGPU() { log.Printf("gpudemod: validated GPU freq-shift used for event %d (%s)", ev.ID, name) } + if m.gpuDemod.LastDemodUsedGPU() { + log.Printf("gpudemod: GPU demod stage used for event %d (%s)", ev.ID, name) + } } } }