From 123207f977111a74ad60b833f72b89345f471c1c Mon Sep 17 00:00:00 2001 From: Jan Svabenik Date: Thu, 19 Mar 2026 08:12:45 +0100 Subject: [PATCH] feat: extend CUDA demod telemetry and FM kernel prep --- internal/demod/gpudemod/build/kernels.obj | Bin 47909 -> 56212 bytes internal/demod/gpudemod/gpudemod.go | 40 ++++++++++++++-------- internal/demod/gpudemod/gpudemod_stub.go | 7 ++-- internal/demod/gpudemod/kernels.cu | 27 +++++++++++++++ internal/recorder/demod.go | 4 +++ 5 files changed, 61 insertions(+), 17 deletions(-) diff --git a/internal/demod/gpudemod/build/kernels.obj b/internal/demod/gpudemod/build/kernels.obj index 47bc9a26e90cfa7a2818892064934d0f7bb4c876..906955faeffb170b1b472c524c43cfdd87167a17 100644 GIT binary patch delta 14585 zcmb_@cUV)&_x{`@goGMG=tWATsfZdxu%S|-qN1W=1(YI+fB`Wem?(&fsNmQuVp&&R zR~HL*R92DIT~=Lt?~A>xx~|Ii%uHyypWm~8{N{Nw^PcydIdf*_+_^V5u<0Fd@jkwB zZ2l#5YyH8*wX+cFT5N2`Z7=q8A3DheH7kKg(zC6fk8j5&tVDs3nQIdQU?2()1iOxY zO=xSt&chJ$A6C5CvHe5?gnEIrBZip_fyN1iM|<}g6d{l76;a<&L_i!uqPXIBj>0Gq zq;c4mbK^~p!#z1S;SZoY5r>mG7vWiej~ILf&@dh!0oEcOp*Hc=fXou81G@x-d=u~(AOQ_$@W9Gysw5&+I;a2deE44wk`jM>Y<{ym6x6Tse8ZvsMnC(xjV08E*H9fdB!On|E< z;F|)fKg#r{fxO4yGl2RNX)m*hFg%0400%HQ9AFxQIRKY2xE|m^27d;4i@`qu@{{OX zNNmPQ2(=@ki?B1mAxU&vQ2=K#xDenT29E-Kkc8KYT!b$GwwOe_c}{}#Vz3{;gh_Nx zlL6vMxSQBTxCY=k2CoA&OQx({GD6|W)WZmX=?qQ>xSYXKfQJ}73Gg<9wEzVvw3|r^ zLLLmZ2N;%ue>HFs4h5J_SVsd7nBQ~|H>c1^?F4u+1&0{A@NTEz35IUM$H2Tz!ApUW zrc%Z>6(P4&$^RMH6~v*F0YKrwClL7rdZ)(5 zr0Xs~!so$f0@ys)**W7j#&*0;X*7z!G?-2rjbZ|0RKQGU%xcDL1ZF2=er3!>VD2&o zrBgqmbcAH-)K7cH1OOAlm{`Un0i$NjQpT(WW-DW=7*hkxb;f*T3~vfTCR3>A)>G(M z3Sa^mGlntoz^E8g#F*v4Y-G%F##94ykuh%>^9dMH2KDTgf%_P{>t!GmkbyIdUAYXr zlF3OK_#~BUnY=awTbdAkfC;s)$-wbcE?_cmD&7n7RD2cWsR${iVxcLKy-cYbKNSz9 zau$=9PsN+5T+QU_srVh0Z6(zH(^MQHaplx_0?2BFyw!Lql@BmET8$qvy&=^v0=*%z zA7Mu2N;S?mBeH=JmEWmxMbdj}_HtW#bGguVnJ}Y$?+-P38_56&L%p1>Wk~-puD4_J@{ElCQmBu4Qc5 z1=j%X|Lgd}Tlx)^M@1zjr^8fYlA?T~KpHnCHYO@k6hond=pfudGS=)^xdQ%pRc_X1i$0;xQ zhL^y-gpbU=;e9Dj2OZT+Ih865nsG$HlsEi_F9lvIX`H@RXR@F2=6I+37>#aj0@4cX zp)#7G)p{Vor$%peP`cSky?M5Lo=K2(-OBjMM3SVTDFf=IvC zIuh=vv^ajfxPpzaT4K^8A3pe?dr$5q{^I^mZ?b-5Z3Si`RKJ#jwj#O@WQv5>Z4kxv z`y-Gt8!{Zh5-j!f)RQ5;jSNS7b~PlCq{8e4jsh(6bmk0kN6%&WcTcCT3W$UhSS>o@ z87eWEcn|tY-qKVqGxO4LWn!;#4m)`T;bW~FTwF}h{wx@YY^n%8_f1ieS|A|H(*1Uw zDY|QlZ?|&MWJ4^VO&$Bv5M4G!{lq%ANXtjHrbwQyUq^b^p#}P=$`rNCfCDHJl9LB= zjVYRuCIAA;L`Q_Sn1IX^(&30)YJ!HO)loH_(_K>(47aC1I#ciuvc^dz)?<4Y2r&*2 zIH~mX$P$1ViCN^T;t7aMyy+~E!2OK zeyO{Vj6(d8q$B3_{#4)(BuIJ<%|hlNnK`%!Dye0^Rf77gr2~qV3+q51 z2zr{;hH6BTRD&wmfKs&{Sv;*?IphcgV3q|5Gy*PL$xJGX;HrVEfEB8TbRG#wt&uH6 zB!K)hIs=6Ml+Y#X_f?YFsP*ftpu4~l()zBc1d@#8Hxy=XqZU%VSf>xpj(G)5BFGS#Dqu2&VO0Xr?yOhxON1)RC-(9Dt;v=U#nMNz_8&4C0XNIp~8kc_^Ho z!w@hTP;F#Lve3wY^_Qsih2RkBKc@Bqxtg53+By-rS%JN?%8*P@3(w0mU{RV_gN&d9 zy&1p_Qw~{592$@gl`w;J6Eef@^ht)^3rxi(pa)WS){8~t-b`13oD(V&L!xUq`BWyR z&Fu{hiQK5yh)judLXx{#{beK~mk~IW!U_=Q^s)gx%veYQq_ z7tp6h!xoa0axMsTO?6ugR|VDApQLd8p`FBI{m}>pM5H;OJ8Emb(m#gr>T^V@7Lmzl zX@DeYVZf3m_HZ|36RZHMyQcOMwSkcnK;4!KG;sT527%0=AqfY)qkAk>i{U!gG#nfa z4z<)CHQMWrS`|At>T?46!xXis?i0%aIGTm@R$K#DS*584*{BS}JWSxib8PN zkpY`V=&Tay9?4)GXn?u}te=7IiO5)Y(2(0V@lOw)Go~m!6A5eHlFQ-*KIA*3>{zQ0 z27*yANu26rkFRvJHd?KlaZNpG<7eF_RdLwcB#9=7cuKIoL%V5$Ri3~*T|H+er$!?t~+-lF?PM;07AImtBj)B6O@ z$Y6mwKlM#{lMaL<5i}fPfZ5Ig2Lt|F|0-SoK(Hl+^cTQKH!d#ncjk!<@s}2sxV4i6 zU-7rLkkA7c-o&y+A-#0Vlno3ads42k%A?14gHx)Y%hYE5|L9Su>oKRv7_CIIGD)kOJn1jN`2nptIW7zE=3VxL1hT|`2RMo!vC)@ZF0PC2JaLk3DTa zN2fcC4n_r`XN){hzp9S*!wzTGc*bnsch3Gv_numRgMF5EmRCA&THJE=75&?P)ott3 z%D);_qJwCX^X<4Q%g`adgLi6=-R;u9a9qTa-ivd`s`%F(hX)7G(?2OUm=)w)<@pht3VSc+A|Gs7k~_X^hj7H(lC497dbH2S;eMe%mm93|eijgMyZrEz zn~R^HeS9!h-I%6%zw}Di%}*RhXQT{zKfdmO{)TDo4?H^ZJZIx=lTNeJ zaxF$*$^CKsHmTn5pyE~L_pbX}Te-C}!JG9@zI?3NDqd!pIAVxrpRv*haTjW?4{f(~ zce}aAjixL(wdoJ9K1lW`&GWfR-XQ%j2i=H?&q z{9)6Wu8A}Kf49_ROLIHDzq9|sPU(Z1#SgZeu^cY14%wMls+@t`|2REquUxs9f7VNg zCoNg(q5trs;^B;Ti+7iwI=6RDMfjzMk=JuSEiIKl8gp|~f$+WC5YN^FP1?R5A0A)g z#XB81yl?E5^WAdybTS#faq_){?nkdAUQDk|Stq)B^h>8ntsmF?Cf2Ok+&Z)R{pDLv z#asT2OSTp8j{Wv}REA`}>#A-=dAoWjEI;W5lo^d|buQ=fjibF=pAR}OOMl_Iar-YT z*A9xBFt7DHn3h+ke!DA16O<`(U-o+1u~Jp|jC8 zQ@aU%4hl_iTK2m4p`#WXrrg^3VEn%R14f?!U zqg-P;_+zIB!dIO>y)BgPc$!?=SNz@i@=)*EU3-#@s*j`|-(Oa7(B=7;c}IgSfA3>+ zYroC&)r;+hjUI1h;p>$&q~Bt_atpVdf?JOTE4)8{|Jz?XDw2mA&QiX)e(I!hWb7_( z=7lY<2dy_NwlOP~$$M|&4>FJlpLizV1Y=>>gH{te5AG)7k6R}U-PUQEC(>mqu-e2y#PS1M3FS^oo z_=co-FYRE#(hX5ZWsc!p|M+k+{)2V@;A5Z0M4xhBK7U%fC3Zb#ojyPIO`3T1QPU{v zJDX3x++-Lm|0!ML>SFexX2{N@9)6=PVq%-Ic>ncqY|%2@q<0H z1|9t;9)os;JEqJ0}>Dy01XA?S%Pjy>qeizyMPEOV@DS5V* z7d8fQ%S-Mk&K+(&dZ)#bf=^i*lMUa86!tQ@lU2H8M2VZg!BG3%ytGkPD^GX_%xJqg zEY38xTozBgO+#ch;XFMvMrTlAx zS>+_*G%G7!w@){1c{7Jz=}|P-p|W)PlTMZWO+M@WiF?#W%P(iTkId*N$}$N#cvREJ zd#2v9A2U6AU#?u)^_F1bljk{w6TL6CD%rPxEY>>{S~X>9;!|JqHh1MwtBj7-RlVH( zD#HEw=HG@5I`($&&NIcKsdLL(SOsLAs0cpYZ^n%)y`(2bT%Mp^6Z=vxXQiFjW~_a8 z)%BER;c~zEUkd#vS{#rhpIg4^mw-((EdENgU7&e$d++FXFKU0SwA*uI#qhG!j}!Xe zJ#U}9dcFC@lTQ~9e8FGet4j~%y&c{Qqy2u0+y85@{NVP$H49(=pl}>rZ`Z4thmYNPdy?`GRMzX&e8mIRI)VEVOj3? zoxd+VxkW#;chwV3xS=B9S@E^QYd^-Vx>tPT(Xct4|H|#+9JGC*hgVzk#S2CX9Nf6b z4_EDXd#_6<`CyT_v2BIRi$8kq*){KVe$uvm+p0bt-c}U%L0x`f^WF2_UAgh2jnYkr zWf39!Z&H-|tKTm0FSu6PVP@Oo!iA5Z;XJzG~edUpxiIyh(U>8$&GG8c90|6~}N^JjS9u3NLB@PTnt(#+hP z{`o_bJJm#ae$*oGa)a%|I?l=a@x+Dd0qeYPFBTe0W-amP)6MMo?VPc*(m&>EkMw+) zZyS3z_}7I#Zl|h6UNg4TW*$BN^NF~Sf*YlL0=*~0j6|HT6yt^Dp-*muSH zBR<4s*DjkW-9GJoSOO<|pVp!0;szv zhnU@TQ(hQklj|Q?@@31G@76wl-PvgIA6pM=m&G{4QzqKgYitO80xz4X@6lZ%wOPd$R>1<3=I^CIM(l=H&mK99*WkTSQ@!(V_TgygzVHLWH1iLlY^|f1u6-o9%faf z^PV^a>0(mOTP)uKew6bH_(6ypHA~KHKpHllpQLSbMxd? zz`#5e&>-bF0LuA=@|Q%ITUbe!1C~-zR{_&fG^juPztwcc-}aubhp;qcjEN8TMDHZ5 z5;ENjHJ%%R#eS;d% zL_nvmxKtOA#mn1-J?z0Y3?L)rfY>|-0`jB7jRdro zF6J^o78I=oWJ%E`KrHAV0eR6idIE^PYDcdDk(Ukv^fw^3iG27%$?V(#8KD70-5Q`f z?c4znOSk|)-joYz;089Jp@2N7b~K+`JsP#cwms5|IDLzd7o~HE`S|7)A2KjR* zsFzvcUf6W78@C&G9Ndk&ji(ImhSLXk=WbE$vVL;x-e1mLrBVRQ%O6HMMg1PZFCyGw zor(s?aH}B>+#cL}NLy|v&K%-k(GK>2{WQrCq#DY^+lDyc)JPep!8ZVv0-A+QhdSU3 zfac*IfP}+kTro}?>Mb?rkaL~q|#8e2Z0ok06DnG`1JSGG|G{? z7(<-^@u{{8Ao3=Oj{+M|5FqlqGoQXBA^FcoVZbp&-l^zQZ3LhxR2zvB>ftCL#Z(;) zD1)lWx-m2!kO|c$0V0=O6s#OCbuNV0!y8&qdnPcZj2i$46%zO>)%fZ`V zgR!81;)|n(w=m!EFMRRv7TouE*YI|5q3s&(uxNw~vI_ui59lC8v4DbL@mnBYQPORcjC)T} zB%sd}6#(Mn-$yy%Lx9A9+Hs~-b{Z|?EGQZT$c~~sK+bp_SnmboPPxaUW#(;YLU8bQ z-LfZH5`@}g(=jc$4%9YmjEwUq$YNka=nRq2yHlH;U=xgg8SQ}Y0_saSS(J<$NXG~P zGz`a#almRoqbau)&^U^&14_iNqS|pOR1S=maZ@N#0h&h9MnE$tx(H}CMWPrPH=m*a zK$xN=Kr2Wb7BLO6GDKosM{R1rW)qE&H&(`Nqeua07e(=a_Td>Z4tP1B!?Z^=pc52* z0(6?p#3Hf5KSRX-Z`3AwoQ(T}qD6qN5oD3nV9gM*zC$&W8Z-=tzGKu9)0@={wK*1scR*aUTq zKVjIQO_>TIcFgnP%J6_ZU#<$zg8Di>2=yEMD$kcE7GaNB3Z9h+kAX%m!i%BlAi~wq z^bjq|SMb6`xMRMfH92cxJ<-E}(mP}zR>p{M5-9UPa)!^{L)&6>8fYeoa49&LCc>BV z72FbRJlmJsf#It?#E}X09sDEI|KJC+eR(o5wx6Tmd5ZBMXu68=9B2lK@eyd^#rQ2W z*<$Q9R{{6zsJXt}7Q6)NBls7nFJP{~m-`d@6~N!*8Q|ms1<%p|Z-l0`0lo@N4+AWj zr{F~z;1Fn%4RH261(%O^&GY5f;73rGVTVEz)nKS^;<-@2!AA>yc_xPVU7>>KVTfDL zho}tkSZD?s;^ojJ7~)@{$u`7-1tco}1-@|Nq(FTTZ-Tl8UxWGqHe2Y+{f$Ex`tnST z@brbS$VPZKG@XrbEi?m-uu~CuGr~im$u`1;MG83QjurXBZRtJK6}U~YFLx1-gZdd> zQS8g(8{?W{1<%$P3l}MPKE}8UG@-^=1Z4c#^(DLq z>gV_g_T}lBV3)-To}CGfgrInhaAsbD093y!I{g zg(ux-Q18XA%YESoK8|hH`EpnAfOWo_m*g)Vp;W^M&Ic$$_}fTwTn~m1V^V-2 zPk7|4WXvUCWX){e0Ygs0MjIIz*f+E31dP0y%}5PU9?guhfFTEH1bolh1nyd5ry)U*vi{fm9~g?tNJ*L$%jWqDP%M~N z;3N{HMkl@D>s~Z4Ea3WHK%mB6C169I7#n+C(lM+T@&u(Zd^1dMX4DZFH846OKx3?M zb%}yE*=o^xg=`)9DhcH*d{}gstgv#ug7>!-o(@g4HQo))C2L#@O;;Q2RI0Em)%_NE zj&|}Mo0djgka8Z!thck{>3-cj4-6kQ)nCA8*V{R(bovY5=>H%LC%`S%TiE@sv!Yxx zFBi!uvbJ_$4CNwxSmfeqUwvn-|6_Mb zC&f%j9yh+ppSqW+wzTB%+r5_dUw`=x@7pmxDJ3SoeRzB#?!MoW*H2ki{Jo*^;Q!%Y zPXAX&{=c+SlBXsn$1?RdQQ(PtZN&`%lr7ln`rl$Q?57MIf!N6SHxSM10KK f@AzMg@yPw{{*&FQ{r>;SuJeIb|H*FsfxQ0%E7%bR delta 8438 zcmZ{p30zHE^vBOR_cqUU8>G2gQlfHik*Swxl%&$EP{>eTMZ+V*lSk%dNMN?Mps^aMfYxDXDzuOng^%t*~`_%#rZq+y*bBCv-(az^x%R#u&%eucd`-k zL{E;JJTds&oK2e;!B1f5cL9Rt!V(q zxiS0+F)9E8MOKOlhZMbudN8V2&HfFa~j6Fl9A3T^L-U@kbU%hQUKMv%a`dE(6nl;)gUIKOUTuN6(aT@oDi?q9zXQ z6|8X`#WW*u)aWYMhQmjlZVKTX=N=BWy4HL+`03JkayU$oUX%>Hu zfOi;2f-$aFB**oPgrT@zk&wjVxJXz-W3rx@Ul~c)Gm@?+#>XS+dPdUq#Q0q#UC$`G zo>91-QFJ|{fUGCRkx^i4Xw64~pCOGGN5OO&7qGZ43L0qK%Hmg1U|?j;M}xajH-%{Y zMvaC<8ZTmTNHpxH@pTsGMuTD(Qg7dd)*p<9FdApExHXzCuyPEXMj-}Ipcr_E1u|?n=-(TGgLRmx5xHw)%hzLG?9VBBaD6y2=fUa1ChK^Z+-($uNp0awH)E*p zrdAPXqrp>gpREI=NYycM9efgWp~=ps z{A{ILFsJr^q_ve;r8YXSu#n_@o%(R>K#qV1%XeCf?>Z2X%s>M_U;B}PNQ?h}`aAs} z{mG%|=npCWmV$>8{AX?eJuI}L+TTdur=!c)O_1T~PCK?C;dVGrB7N6_I(qy@k{ylq zfXo0ReY#($&h9vq$S$bE!+;*j9T$kv*8`zSmYWw~H9n_X3}f8#22t6VsP%pJC` zmbdcn3)^-CjkNyw$V27%(dX7JC#TKO(OnnZwCqH0{%rj))nR)#afRGk?)m39$>GZW zOFe6X-#MMF9}>T|;J5v!SDxRd+8VZ^YMf-jG`GcBCh5c0B=&WU%-Pqw)_X^GNKNCN z7ojtfgWK&Vx8L(_KX~Kto@cQSrY{R?-{n1B%dfpS?d9fMdBMkQ`kb2Kr5_Pdw0lDB z%Bp~ECoepi_R?uly6>{u!jhjy#JbxbyJ$6R`<|+c$7IunN(ye+onF0v!NiXxn!kr% zFY2k)EY?na@M@l=;2;$)dyjvR6Mbglk~Hr_bCYA&B+ods*3Q0o+QHrBpC+$WX-*Ma z+AT5>Z%FIbxUQmJwzOcxs%|&zmP|Cc)jo0fgRIi~2U{+>R%et}#V^e+H;;0d+0QL? zwbw!K=T(NyeQp5zWrg+x*Ny~eA(_UD`^tka$GTGeK7=3Rrt%))+s3s*eR z4;k!!gm*d8`^@xeS6|JszaDZMH}Fu&{TB*_1?jEjtz&+yexn*y= zA4zT-wR-!se~P28nN6$>=mVkq7Dp6b_%!KZ`s>HT%`aOIIyP$O+&kLc=WCY@Sl{&i z;`VlLV~a;qm2ci0i>k|Q<&a<65!xnmBlPdTIKhaV6)vk3* z?>D*@^D<428}&ZFwTrK+pWpoHq6zDLPAKXAa^rM+c+Vkq_BjKt#@;JGz2a`#%RDcg zFUM!Js!o0hvzpvBZRTmDd9SZ}->i9vE#*F-Xs7bg>fQg2$^NTjGJI&tq3pMW<)GjXeuDv++$how( zb!G=_NBhO>5pAt(agRKkv);G(m!GaEE^dA1tL=Df`*bzU{NuO2yzi%`w`j0YVVAqZ zuEcC#Y^OhR>9L<~Jlu8V!p0KqBYma&a_?^on5mn!+epbP%_^%+ULta{b5FGKGQXlB zw#ql_9&vPo)oaJn+4hfau4%9tf2}CqZBD*Yn8w-qytRwFoh*UJQ;rYZ*y5ctz9v>R zch2btN3QzCuhCniI!Cv$O0D-8yG5z@Kd5h%8mmQ0xCT3A)8V@#szz;_J9IPOV|jMO zr~19)Z~paYU7*zSz52Q`@e&Wscw^teDTRl!=N}R3qFv$-JXCvC;wfkvY`D}QzLPrd zSmD;e(p$Q3+xN((7H`^+x=7D4Q7?31j_K!{iFN+{bth>YXq>%wLVZ7eMZq6pzUR%O zYN|o1`iaG1Ri_#yS+|$}!CSk39C1D}y*PUNtTye?-9PtnvE24wUc{nm>4;4UrmtqK zZ0w!nJjf%sZ+?cI^~&sV-0QisUkf?2o5#c@m#JJF*f;yjYq##Tt&2}IJzbDwY1p&Krj*hHe7G-if$?%V^#hZTqxuP;=i_)6&tNuzo{>#aCb1S@e1!kTp zwkh3m%UrjhwAUzwJv~08=`}A%wLg=5t)ybEX+ru7r+(S4cT=3Rd}@7HCW>xMdGn`* z|KE-31)+z6CzLz755E-HOPjM-ySKn1yr^Kpn=I{F&z82OoK9E0v@1#B$;?qkS{0U~ z<_Z;WBDF;@FD?!I8_A~nJ#d*tc(^!@5qXH51{x`*`i`9$ zMkF%>24RN%cp`lmk<3h(mSXDNtFsRyl9@*}UW@0nJTb&(#~igzojn+l9t+`eN>3*- z&xx@^7bh{EqNbc-ia{LdaTPzj3eTx*Lr-eP;Lmu}>>*)VU)Vp*C?|uIbu@fI+u7=n z*sd}nn_5OUPtI6!8HhOY+kp}JmHSp88uEaGIH)CcJk&60^vrCE@>5OVeX@jK4Qo?P!9GR8ucs7^v=vsQnd-~vxa1fzB8#zuWS|~vc2f#UmGBjm z<|0*7szf?M=^@eyN=9iC{tP8Qqzmve-4uR7YC+8z&rve!U5x*x8QpXV|AcB@NG~Z( zM|w+XC(wYNOFSX13{t4bol`DSR$A~<&;L$kKbY3j12)ZHu1@|aBF$SfTOZtW<4nkZ zK;+_jaOz=SnxUp2+E)s&KR3v+SBKi`Z43TILYj1>w~O4>O-`8nP#1kwT#YJ5EjXb| zZhF`(=nENq^48lt8vsCTvxLBy{CUumxS65?n!dRteO%nFW%s7(7L|m<=Tt$V6DX(XsK77>c;> z{{$fQB0df+`Ods;qdXd2M{8ArjBPUEy%Ox%CgXMC?`?9t6uWFEGaQaK7UrYPfZFYH zVTCfh-Yyf0mBFr1Ce$cHEV^^bunOJZ%5WOp2j!JTGF}tpMRMK-CZX*IThWey7PN_= zQ7ji`s(@FqOaK+gLRX>!<>=~E;32xpDqylhCOlPv;2knv1PgY^d1E+?)(PIA9Rl5V z%JBjpjWz{V@01I9s_^GdnNXw(LJ3Yv75boSQiZAL?yEu}x(}*wy@bq4YnPn2f+1-8 zKsMT8uovwZc!YL3nEoyo7OO$X@3{JEun65jHK<3|qz3QM-B*Jir8t@zj474ji_W#B za^4osq8$LD-Ew{;^hKKl$-CvkY;`E!EfdzN!%cL%)j@ZUOgOF%KIpEi!z^_Fs6*u* z8GeI4*+XV!UMA=JLnzu1$U{2`j+V)Vxf<}kOeW-OK+keK8#G`Xx<(BsKzBz2&ZB#$ z0g8KNycYD^E9bkxG_*g#4z!VQ3vDv!?UM`hG{JA5OeoNV+33nOVL!S?O?ZaxmL^zL z$b>hV5MCj}SD;HOq1fWm<5t5)YjgC|AjZKeS*V zx?5V1hOS);c2&vnCUB=p&i4fU{c_$5{Lx0h9JJ|hV82{gt_?5t2J9S_>x;h;w zMc1qY_s~7o0izn3@L2}}Ysi^BuZAp7HQGV&5^WgR)XHc^0A95fQi$d@(0M|E(q!6>@cbt66~xc{2G{Q-un7F8Nk-)g2|fABsdcE_pin zR^-^Bd>d#nD*YkqP*)uU+7kC2XrKUx4;lKkqD-DwzV)!gUkN^_gy4tm<|tIi^UJrI z*{G2Gx0;>$V0KuhP>Tr;HXqi9$%joOPe_FAEPmKfmIkoluuK?i02k5iG63Z|nP6!M z1JNxogtR)D>>(in, out, n, phase_inc, phase_start); return (int)cudaGetLastError(); } + +extern "C" __global__ void gpud_fm_discrim_kernel( + const float2* __restrict__ in, + float* __restrict__ out, + int n +) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= n - 1) return; + + float2 prev = in[idx]; + float2 curr = in[idx + 1]; + float re = prev.x * curr.x + prev.y * curr.y; + float im = prev.x * curr.y - prev.y * curr.x; + out[idx] = atan2f(im, re); +} + +extern "C" int gpud_launch_fm_discrim_cuda( + const float2* in, + float* out, + int n +) { + if (n <= 1) return 0; + const int block = 256; + const int grid = (n + block - 1) / block; + gpud_fm_discrim_kernel<<>>(in, out, n); + return (int)cudaGetLastError(); +} diff --git a/internal/recorder/demod.go b/internal/recorder/demod.go index 238d567..44fe381 100644 --- a/internal/recorder/demod.go +++ b/internal/recorder/demod.go @@ -2,6 +2,7 @@ package recorder import ( "errors" + "log" "math" "path/filepath" @@ -33,6 +34,9 @@ func (m *Manager) demodAndWrite(dir string, ev detector.Event, iq []complex64, f if gpuAudio, gpuRate, err := m.gpuDemod.Demod(iq, offset, bw, gpudemod.DemodNFM); err == nil { audio = gpuAudio inputRate = gpuRate + if m.gpuDemod.LastShiftUsedGPU() { + log.Printf("gpudemod: validated GPU freq-shift used for event %d", ev.ID) + } } } if audio == nil {