From 2efae6fd9cecd09082c47e3f9fa9096f119d69cf Mon Sep 17 00:00:00 2001 From: Jan Svabenik Date: Thu, 19 Mar 2026 08:29:54 +0100 Subject: [PATCH] feat: add CUDA decimate preparation path --- internal/demod/gpudemod/build/kernels.obj | Bin 56212 -> 61020 bytes internal/demod/gpudemod/gpudemod.go | 4 ++++ internal/demod/gpudemod/kernels.cu | 24 ++++++++++++++++++++++ 3 files changed, 28 insertions(+) diff --git a/internal/demod/gpudemod/build/kernels.obj b/internal/demod/gpudemod/build/kernels.obj index f4b6d9c3623701cbdca885af63f78efccfded6eb..062d18fd62a9f0e44744c145415b6480d4593237 100644 GIT binary patch delta 12670 zcmb`NcU%<7^Y>?Vmz-hA8CF0fyQ_kVh^U~L(ZhroKtvG)L_lx>0TDrkLd2XEBOrvU0JhdoD?`N zv`EWb5NucE>@Y~no@-Y8I9-4Dw%#5dP1H~yj;s1DuCE`)d3rX9GeK4h<2c)}qP13D zd$c*O6q&M2t;< z-aW}eH4-;>~$$X;m*9-k>-rEI6%3bEG|I;uYqCo!CgIFI2n#9a&z zA>L&87?Dq=X&T8K*MgxtVi3ath%pSOAQmxPgt(ocl{y*Rel+hh&r3v=6gn@x6kHvK zoe{?}oP;<(1(vGXDX&00izu{LLVBM1UopK+3HA zD40sqG^cW$149qQ9t=YeCo)V%EM~YEv4-IR#A^&6ApV#NXVmSKRnxGuG&(&eoSp}o zp=oqZqY(4bpqqxBpg0Y#Svo4Aur>{pt?cbmkQwA|543%i6JWg+4;GU zOYIHJt|jqyf@`^SWcRR#S4`2IK^5@~yv5I;F}_UchDv{?#4#lqm29T0WXc9qYM63? zDK}7g!W88^+Ou{Z$C>5Pp4;U?4=qO}6oT_0Nz0zkgN4kVmIp_v{fya_d0?bX{65+= z{$3tLqCFGx(Voe1;+arG?Q6{LHWQR}h~HU<`jcnEAZn*Fd&5jvOYIZPzA+O%P}@w` zQ7NC}wDX~xu05X*(aiSChf-?qW_El&JfgO`9*tj}4{h{F{19ee%7DR;oL z*QTI)xAjRLEN$9?g>AC(6=jYaR#3grR?h1O`d|PS0mt)x*d}{GEBir~W$r}exU!b4 z74_tl{kV3Vi1Xu2V5xn0w=of;$Hk0Hj2JOKB7S1T1kP<@lH0h6k+E)L<0g+8@o(OQ zsFCpzW19P+hs0a8As9xqu-0iX6#YVPhcVE}(GH?o7(*Y&QaIqKQQfoSoI<=9D*>W|hHzEr;<0cBjJ=pL~ zbQ49yDQtzDj>7uY@I29Ve+debP~DJO@+R)-=n!y*T!u&?s<8GhP0(u+$xvB9c3||I z9zk{e-}Gv#8#L4D{;yC+!(M;+&4gqJ8ovYuwV}G^Z+chM^>v%9C=Q&IAyik)80`#2 zmy?8edg|*ML8$X^D_mPzL}VNy(urnfMurn7qU#K!o&EVZSnE84zXNWq#tOoDnBVG) zoXm!7B0W{wB61>kRgTP_Y(!ThA{N38t3qPPimEE~hSe46Go8jGH^w73tehxPgv)~c zwGlzMtDZSIOyRr6cr;wI7NB&h-#AQa42;Z9g{nyJ|DF#MlB>Wz@QqQmuSSU)a9$UW2H z1j$0WY_$J-I+!mKaKkz7C&@pi)ryRg zCK*N~!|Vc?A09G=qp@kH1<=Dm@&?^FzTe!TS}c;>q`<~TELSWJw+s`=VfSPH#yYA= z0ujk9$!XS*M>-;<>Ob<675RNx7`u#+75Na?Q_YOsXW7ZE2&2t0E%;$=1THqZuE$Xz zYvkY5k6f?lKA9~b^|FHykIKJg1S;;^iA^%tL9oUt{98t?VvY?{COe@A!8PDNG7c*; z@|x^}WYpQ$Y63~y%C8FfcTr1jF!wYX}EUO$q5b?UEpF zcaTJ@rEaeaTEh&oU;p60eVVHlGG4{buAkP%_r7rfA9}Zp7&p6ee=E+WDnnhlo;$_u zy(r$p65}IM2e~@B3F~@aHYO7CtJAUga_)=RxlLu3KhyE6ls-B-c=iuwCwQ zOPo+Kwcmrj=eErVS+lWUDSt_{9;Q33+MpczE__2sK#TdSg33+DL|+@@>AcA&Bz))I z>x&Cc2VJ@0)3)uNT~3P@%~>?{=cCeH-b=6F7&fu@nMaX7+k6a9&};K}-|3F4>$SJ# zBnS-${*k@WX`}u;>B{n;Zs%(|Tg+=Albu zm%aSj)iq^ZCE{3L^|3ytE!PfmU6(cZ?1hB0#qB;G-1TmG$;0}RtP2i{q$Lrn9=Gea z+RqQ1SG_zU>05fJ)}ebsguz-Dl}FDi9Y)DQMlBrOa(qBk)V$7RkF55r35kF8@aC?l zsNUy#xJG#I40_4J%DLt#d!^P?$dLb?dO;cdB|}^JODnm)Yz+ z{o}^6{Y!6Zc-DE_%L{IM^nUX4X6=kmKf|MiN{=eOo%r;_Co(9q&z6Eknrbz1Zk4a< zZFRp0B2Sm}-j-EgKWu{ELe8W|OFU+g>f$ zZt-%^(LQrp>)Cd%o1>Gu_|Fx&Kc07(WE}r3Q#Hy#|7EN}hp}~z?eh*Fnpx_xukfb! zJCkK;79WGUS`G+)RQBTQ18AjD>wf1|`%hO_>%|njS+eMM5AVBuS7l|c7YlQLYx}t#SvXV2<5=g~{Nq1+r|jN1w42G>d+XQ58mWYT71qApGQ!KM?*fxFq0@_~$cSSv zesZdvT+QFM_Kw-+QoZiry#wCxE#k_7igvvmLw#K30lmAlO|CrTzOT4syH{ELp>N0R zFNWN3D)rAYuS@m`Jv#Nv582Y_;srae*uSr{tK75V^yYyFLQl^aF}>jJa1Y(g=VARy zj~R&IX!4PpH$wgQv@drY8({Aml@zO`dT4RYeREga`s|j&Oh2EqG*>Q9%d0Q9bUnYT z?qdI(iktkjQd)oVh0#bn)4eJeJM9rfnerdfB?9(}&-nz3*1mXcR)hwbIX;WeiEwl+@! zS03-PZCt$aE4{+@sTN#9uWeuWH}Ai+;?gw&vOmA;z3ldkbGCc!>d_Ox3+34ad&@jwr(U*3*;_b^mE*yR;?MZ64x}bo0i3gwFI+=g&{GqW^ zx+w1{iXMJSn19gM`>9t#-3D*DU+e+rt9mYX(w6iH3Nz@E6lho#mgpXBVPDxznEGQ@ zp=r3U*RF$G|2lKdib+ z$g?Fs@(iksY8ECqZvT;$9HMeDSU;olt9~nB6kE9fHT?Yy#FVThQg3g?!*)c+`~2f zjg0(9ip$RKTv&O!%iuPTPd|Qd@w7f_l19N@xxtx zl#OpUNoP&C?fK8E@E3DU%D7mk=VIaI%a?EI1S%I^QqC;xp?POk)~fR+IbU10i!`4W ze9L}$*wy<-FJ3fQQtETQ&gHP>gWl78+w3u&gO}VE{T#)silqm8YE)}qs9)|@ z>o9bCYlEHg$_`x~tu!c{=B_zlxbu#+cH=AtZoX&n(tp$Eu{9SW(jz`^dY`p#hV|=; zxF3ZR2KBtt&rkYQFZ|(TjWscIxGk-PJtH=&x%&6?w;7S3erfcmoD@l4{kTya`c_ul?iAUfAa~E@eFy*6p0RuWUysHZhMT#s4yvlzZ*u&> z;~r&!*1^wHf>qB|h-Jlbi@jv>%&^hAw&6Z?GbY_ye_&Tp@3{|L2G&><|CPIQ-t{_P zfxUIPZ671=U2^fU)}LNJt2geW@nV>l^TUK1_xP{6qvKCc7!kC@`SgO*!^KGt0uK)P zGVPtZv*do9=S{u(fpw?OPi?#BaBa`>zC-+r4$Yb1P*t@gbE^EA@dDo%y+1U%XD5y= zif_N@&uxP6NXP7iZm)cGy#Ic1S3Komp69eV*6-JF`aNF`e^Jf#s`O6tZ}+{;jr?_| zlCHWN`E2^GJF@cIAFi4;NmCX*TxA#ISFT?aJa_FHhk22A_+Hlb{O-eU&qpr*Fes*C zA4mxO7eg{+o%ni)8=6;b zJnR&2RF6r4txr86$$^=NNjoC!5Tt~kgOA}+FnVMkxg$QAinxyGd4!vY`x9$he6kdA z=2(+YVIr;y>lM^XK4US$Pp_OcS|Y9!l07BzD8H7*O-J&inhEyLG6EWDJQB;DEk|rY z0|AL8-$e4F+7$eN%aT_$(g`G%{2s}X#(hCzxt90{?Tq;%t}Rk4N*#zs{X z$wgcy;s6@hM!8`o9g*Cq=8q&n(%^a^4WwE)Qd>&n8fhZZQablUg{H(!Lt-nFgXBpw z@{w%lVAdc>$^1p!CPda@71A==%_*c6lpZ${hc8W9(&Q8**5NcHJF5NDsC{bGR8=`n zM!oHk*fh5yu{GbNif^)f=n(dy$mV$ziOur{QWxrd*htTi*gW4TG$rl}66=u1j~g9n zhB6YHr!A5n)tr!6f38S$nsP1>k(t9ezP#^NhvMTue(uVp=v5o96)} zHqSx$;hmiy!;!ktA&f#|(~L$MD8%`T=$o}Z)IAeLMzfLFJWCXs61NnIb+`s8h-Pd+ zV)HzL6il^KNUWE$NbGRFKw{HW!S@l2Y&1E!Vx9|7WF0O=VskRU&*e-rL1LN{5}Ri~ zr2cewLy?x#Gi5#!d0{BxmLeHZT7$#}U4i6A&xPkmCrJ&NV?Fr8@N9xNbdHk1<7j7?6yw0xKv|40e-AFl_`;=_ zPW)Zk?`PcS%qR(egIbFxN?;N0`8e(U0Q8G>R9Q!lSRho!iUDGc`5mw^)}7x9*JCZA zZ-SW51g$tr$U>R{oshO66~ZK>M+ssBzyl@ZcpaY8@h3+GUz37XqBXw&F2?!mVk_b~(J1_i5j838Jh_ny2tmZ8 zep+pi8j1;Qoa8Pq6*SoFQf#n#HUPGDxRLkf4TN_rZW-MNIc>?c#lCc4gm0laB|LS7 z^wLP8UJKM1*&wMqbSr2sCNQV8ui9CQX-Ne z^^!GDr(W_>mQe;0`&uI#X&UwBHBc9>0MU{r&qX3vx{xbrq%tHM>RpUP-W&?)?TXQA zwC#jYmS6!r62<%)n4e$?X-FFh8Fr+di7<$fDpC~`C0fEAjM+g+f3le0ODO>95Tzuf z)os-`*F@KqQ;*oArT8(s{km01p5iueodP+Uyl{dRlFe0AU z)RWNYVMILdsfWAI6I<|r)E3Ue&QCY z&^+MsG;61Yl_$C==M=XpToiFw+tK*l)`I<=+qPJnT`@53{fbPDRG%FGuZ432Z>z23 zp#t4ydqdl|To94Dx6c}tphVN})w!9KLe)dn&yS{VeXR$R+%jNk&Rk9MJ2+BazuXrJ zKDpr*n|EtKzJ&9$xRgau;TVGy@YpBH8;8$*>i0}a!zjbS(SfRp7LyR{n zy7;6|vM?0Ou-{_HlRGO7S#y3<))NJ9xmX&44xr@Gtn2%zgew#_L?0Z!4IcepgU&}~ zIX-G=IIOqB zGOQ25MXb+(#!72rayuu>`(`@dne3XCJJ1<@T+q^m*gDk)zIZ&%6ipM&FS89o^e$NB;Q)=%I#)<59q zYMH=59onps3Y^tp0yci?uo9a=>Tm^{@#>(xR?25Vm$fp!3{tS(0u@*vf`?dN1M_t< z{tfh7Cljb@K;Am3z*++=wU^;1}d^$)nZ9tWifx*MbdM@{I4O-D^g!=}F`Y{h1*COpArswP-%l=4L|U?W-K znON_I{aBxe_gFuLR+~s~BR9zehFVazNh)Zi1!u7F*8I@Vi3Z;K57A$G_5Hl$)^TYb80}qp>~(3$T6we_{Ook_udOARe1;I#7&@{4$VZy&bk=eF~mqeGhD_W&9`jqgp1=(uLX8Qh~iL9LA=dE_}hJk1lv@ zlL{hqArhN(U0A$L%9p^|ZREhIY?twWLPxCcK?2s_VEuNPz(f!3Y?lgJ>w!^?R1l;G z!Ptc9K^8Wtda$!bivOZs){tef-y!4Az+kK&z#ObU!;u{_fv!G$+aVP==!4fzJc0Bf z3Y)?DuoRn#`fy<q&z)-S<$w@knrK=0jBfw2K(W8-1~ zyRiu{fH&9-H2{Y_QhB`Mi^vmrC9vC49xCJ_AJC}0MMV>>kVa)N-g)t_NlY@DTT)VEoXDt?R&Zx-W)N@%l~B2*^8&{}gj85SuSVBpJWL{jX8 zm@#acCsAUPdQCFW`ZeJ0_<@-GpJ8KZ$P%OSD{CpnWHihAj7s*eEZARb*kb@5h>~Ve z`KZioj6&cJph6ypf9?ONLix28dZu8tSE?i-AHZSRUPG9$*FxN%SXfTu#-@^Q3XAtj z1)EIa95x1Kpt?^g$TtIDY~GkbA~plfVFNbX&EXC6%HZ^IJePJ;}hCQiNRgRO~{j<_cbq{tg4^wf2pL^?EL^_91%DE z;;s7Zk;y!KII5%V(LAW%RWUyfY_Av#J^cc!JN-FVt7Wr*hM)gQ=&z6^&3ur2!J=^w zu=IjyqfveE!Z{usFaE!G2d+=JebfOk~O7 z=)xR^>7TD@!rRp~-@#_gOe>~S^I)_GyYzCE^!!ED7Bh^F^)CXV2NURlc|CjmA{!!i z9m_C%#^#rq{}!XhFhelfJ%y%>MrX+chB*%z87vAJ68vMJ3J2i~qY|F~x4F_}3W#Z?}ZpdaROp8ht`Tn z>#10Rcr*e66s(laBEF*PM?{@SQotmVVI&j>B95at4KXv43|ol!J5{R?FHyXM_#qNr zt5_={8AXv=)={|86nzk

ey$p|}EZ2SqtzQxyE6Xsw9!FBO%eaXX{2mbuz*xGUW; z42>qUh(ugQaV_Feif0jD(R>l+f5gya2FdR>gJFivfJwM-NbxhkT**o)4RPZPGSyPV zT8a&b4`;wPWs<5KL;7`M@bDBRh$ASDLrkKWiMW>HR>Z>;Pb1!=_$Q)5ELjUTH_cdv zapOWOrM`%xW63lpBj!+Cfp~yoE#l)?C{(dhdV^>&lMJ(;iDwtZ;fRql$(m*%!b}*T zYNfOp@iN6`M4dRIn#3_oU>vD19#KkhF5)_h#fV2Jo&gDpm!3a8`&x*;VP1#%*dKef<_w8O~Ue4yiG}9Y)ld= zlfc`UOPE39R!qQH3N;u@8Ad3DCnQ!gA+fg(6~Yh*Gc@D#?4Z& zG3D}yn%XKRGmKU;q+py38))p23};CEoW_aCpl`rw7!s$`I5q{!NZd%{!W8(M#KsmR|6&SwTX5rN(3m%yY>?$_ z@{ZZfW*9M5Mp9)wDiKs!K$SdH)=}jkRgR(3NR?+)d5g+dsu-t|%2ufi_g=kRhp^t2$i=~(Ml(k_0t)~DxFmJpo%Xl!Bk14N;)dJRM|$AU8q!3RtypF_081Eu+BZ)6#+(P27T&%>;WEkyCGEwJDGWDLQ45msL zRiaRlQe_pV*eaq>#tGIe%{Y(3O`7p{Ca7tU47It?Q^T5_3*#}K%P{@sLN19bXdE|} z95!z*Icx!D6ro~auEgD!j?I_n^hO3Y9)qwGFR?yC1V8)e!_E&$FQsi`f3L_ zXFbO3radR(dCT3ivGIp7N*+vahL1NHJs2PggEQT9A+ehmycIdYn{K-DTixo|4)Kj` zcljy1vz#D+y$-I>;?T!WKui5%_aZAW!e11MC(yCK=!WeXm(eRQS`ZklC*v^!CX3(R zp-lY%mh11*F8*}Wh_3rnXs44z*Zz;LjdxQ8M(01egy`g;6FYUzCLL;>G%f(5tN%wY zMqQx&b3?M)Qw|Yb7uGq3Lby{Zek`;&Me^P<;N$$+kDDjA2EL98^JC6@7ZtKzt`w)Y zZv^2V`@k2_9BvjIYr}3AYgPpUT$5OJIO$r|UyS>}H9R`0e4a!!=Bjx9uekizSboV( z#>$V1@$J9w)>}TWSE@dLB36O~OH(ikFx1$@mBn`Phd3_4&?!NYitBzkGIqOKrFnV91M7fwq7FJPlRoF=BwKqbeO24Cx~*oZ@Ma;L%>} zhs$cu!jCP|6wnG_tEH}!JDtXVM!atyQSi$MU2cT28?FWi;YV=iwnz<<0&PUYu^Zpi z|8$qPpCsHH2$q&Y*^dQ(#bVsSJUW+Ss(%^4b#Zgn!q<;+Kh6X#JmCK|zP^3D2ae$y zH=a9#+s8v-kgD+sjJXZK5}lfgI()Ox2S-b7rwiD@r3~(rfTv6QnJeWw+Z)}|3{?KF z5eqbE+K&Zr5!q6wt^d!6y!H``emMpi0qaK^$+*qH_Tv)b)`usOAd3@!X=d%k@!r?D z8PhPUeV@4TxTm^VowFG3Ud0`s>;JX!vP6(H%9(f59xl1*!=X`@Jlr-!t5hSJ_rw2K zFjs|}1rFp!*autw9Bi!3vVquF-JX0v1(pWe$fV}9lrME1*6x7IzhP3dP%mbQ&v^|y zW_W#1)3Ha>F4|AoRdw0)#lXQ?eblP-P4sW|EiLG_>6Y^S58rnV?cS%6sbdZ^GrQgo zuh;h(Eg4mrb)vOjK;E?A)qVxpQxg@M&Bu9pEmb}vQp@q|y3(#irl0gCJLr+l$?DwT z(mcNcHso>n>g*`_Zl&?>iz>!=4*D$@0(<$Ku2$RV{K`G}e)Z88Ed{Tep0&6i_PNu) z&)t%hcjw4HuDR8}?1gz~O8m%=QQvEnOXmDm^Yqy3%x(9zdgml%>xJIR-V?P`pg7Jm zf1~ch=01i7Hf~x_rhMk@Gg*b|TK(wpqwR-I6+8~Vda-$oTg862#V0kAm!B=Y=5SIp zbXb3T*}KC}uS^@U_E_xkqahYL8;>Y8&-woJqkkXKZq4Ucg=cbRZFNyzkpKK@g3a`S z<%R)&uC(6VqdxaQPe0A?X7MSa;iexFPFvJ63X z?~e}-U9A*6zF6?M{Jj1+QKNTdbg}O|#`apn%!4A|0)-|AC78K-jh*t7&*CTZ+zR$r zpS^r=(ec0=PePirzpg13J)P82xxIev8WO+HbvQIzGF2K~=W7%rUL;rEjda?K-!iaOC6}OCA3ShU{Yw(x_v|ZOsI@W0#n;Gwb@Pw`O(I8GS`SrqjJ4HJw<*a>s^)W-VJp%w#buLc<;QJIN<3@ABBg(#$!uz zwb+A#C$?8M2Un%Ksq{K$xy5(C-?sRzIKQkk`t&Zmo`D`N-s zQVD&rLd|}3N_D#pf$SCYLv*}LiD{DwJeop-g(e0R@xl#NPTVyVSF&#J;`qfPv9 z<0G3=7i@a@a!`KBFgL^Xg@a|h)7!H2)#mjWG5JQQinns+rF}78Nv+{UZ;f|ua)?y@ zVt9@f&3ee}EWV#L(Y42k%$dU$ZJl*`rLMxK>7kw@n43Gd+Mm=m0+_$x0NxK%!*SkDtcAO^Pj}9>xv$TFCTRbt<+n+BXzZ7uoR5)U&oPO`2%1;gqxcJeM;61oL+bOHL&%8MJB6qy@7c4QbQdu~235 zwrImc%L=On+eL$VU!E7TYl`iGD5f~a_rnaGx|vFI3=DV!zP1?i=8w5GXyszFy5hMn zde;SL{j2yV40;|Wx|wD>A!WEqx|a9hTG>$N`HE}zq}lo1tlQB49zW*A>&(0u=j+{z zst!#B#q&P($!nrtN_3rCMUyvbocLb8Fl+ zH4LA3_tp@>sqr^wWNn`QRxxvfi9;D={e9c|tbX3QUdz7a^@-7|(T=;kuJoLH={&s; z(Z#jbAsk{N0WBOE+d)#vA^Tac?f4+L-8< zd~oFcZG#`!1RknxseC=eZ$sOC-PDuPu>Ehf4R;F67yOeV3+}sdXur>k9hJoS-$~s2saHgtMXWFHgsYqCff0*S8xDK_%S)73JyA~G+3olpN3DB?F|VZwqM6SI zFInrgxn_h($&Kv(-}`(xwQSqCk&y$ln=F_31jfWnRZ3iV`Mm1Ia&KSPQrFhC_vQ@k zGW*NemX~(HT4_6^~)oQN^Fu~lHe{hBcCzhX7p^K?wXQQ-P&u(FL~v(7p^C0xLh+6>V;6)NZ*M)~D;@pk;qBcg z6&b#`)ve|2^y?S88Mm!ZpY_ddwp-)#qnDMM`wEBJEB|5dc--8%pMS-u%*74qkA|kL z8W8YeEVJm(z+wF>(kDaBwB#fm8_O@(WZAQ|e6LJg|i`SY~eE2mF|doS6e=q|pV zdoplowiaDVq)rHVXP3=#*UI1;$s(|>5rL@@Wp4&qP4&WDL!L@ z){)R&v|C8;-45boaqh{F&kRKBO2`vw8_CN?av~a6pN{qHpco`NcDW4Enj``e9efnY z1YdmS23k`>&1h{%-op;<8InEGy5e(pZ<2Q%iB9G&k_FLPk@S$%ndeCEL}Tz3fMiY%?1j{gX#PlCK|T|Nw3$?iMWSn8*g-`| z-N;xYd{njK*3W0m5iN-boJh2F9a=e(E9t$1MAtbIzmL*=8;8`LjF^N(*BOCC*O`GN zCcQZwv;>JRbEUG3b}KR)P^1O7A~};GWk__LCy?BTR*ysrHX+e(y$wionO*P;713mj5Rr-^B)ZNZq+aB36OanYX0AojBUFf_PpA}$PI?cL1KFb& zNCu?$9TN9#m(TozMBhXO{4zrGY>_mW5nQ(;qAeNR1Bsq+?nur=^X|}sI%o`%9qA24 z>P~1X(niRfWlDaVxihPfErP%#Yqksu6RoXF$wYh?`-@b3vtQczF7+3=K0}T1&91XM zh2N`~tt9&#XTd&I#GWVdZqP_`fe}d}P)_m%d9sMT!zpY%$vca>!yFM?M{-XSyMy!u zpjtuVGLm0QxEXFH_hlPN&Fo}Xc01|23pJ@u>_MXPG6WDdN5IO#JiR-+m(;o-6~RAg z4lp{&9Qr2vfk_ruWi z0dO{bAlm|q(+9BkNWBtV(;BRPn?wtxzR-rpUq4#_!VG=36uM`)u?HYD!-cJcWf^9A zZn*mj_%7lZN(fIfCJkz_=eLd?N?eZ&{>m_u_2}rK#P#G5jn!fqt-%;#v!RYy?w1^T z6iVEPrF29Wmg&+y!tR$5l8zCSxDoku1T4$!2X`}FELc^RQ7y*rzkHVAPsP8ZV2%sP zna&aT%?*Nmb49WrzZ-GohjFI89Y;@RMTtAcS~8ydX3t$#3gmYuO5E>(|MEDEyDVmr zUegY;MEXv8MM#OHms<~#42L)(vUp7@Fs?{Rq+8rUy^$11Z$BjNHwy*Eqk}w=6iKfy z5_h>MF#a9HEn1oM1|ubt-jEKOh%`}^bcZ3PkZx{wC`BP@k=|G&?uu0)zYS56qNYQ7 zQ<1oGXCdK%eM2Xwn9dKfv8F*w0**%03 zkq!{rmnCEm6MBSnf{^|^A$yjPFH#esbfhbUb|Kw@plma}K^;e;#2x7#$3!9-+%fCWQD4 zglt!sm1_n)k!*<;h2#XYbIf1^QV-I59;pvUdO;oaDRK1&k{s6@A?pPLvDQ?iVMHrH z3L;g{AdQ8kOU&RCQYg_Ja)sW%(t2)gdQ&>;Q{w7xAvuAIgm}y0_F_KLPD1;U_7ZxsNT^#yE*Vyj zW%LH)Fl?@%x!8g|O2!XeEW{fh_ikh$HE?Uyi|#lYCGKb!NX~uCxkgB^M2I(fZk#95 z141cCPYIPHy&`lQsg01vQlZ5^WG>&aaSvZsSob@=LQv&2M;M-GVfQi8C9Tn2Rdw0M z@H#D9gZDbGKJ1+OQr~#b=!6ej=c%WQO)FDQ>ByhT3-UsE+IuhFnUQ(J@WQ9h!@jRD zi(Dc9y1Y_BB`mJ+`RtW?UidRx7>Hq|2KVm_T)TZo#)FFxX3G7i0{7nuB{$x17j=8< zi{4TV{0jnuCbvs*TaIulHqj`b-z2>xy*S1P3cOO=mEPRls*}=(y9NEK^mO%=>#Y7$ zAVVFSW&01VYWTcViI=}3D(J8|p+eh_9IyW%5fzy-tEehD}(=T zF|S@3=A&y?hH7+e%J6!(7%zJ3JrdRc#-SC#3bZ}oG}=M%1#KX>RazKucR_BJPuz=4aDttrt zrz(i|<7X{I@0YOpum!CnTthnu)DKA55Ez6u4rU*a@N(5)_W?0)vl=`?CszaGgE%QQ z2t?PS2D#|k)ZoZLF{=fC9h9(kU|YpaDim!H0NO}sK$`*It8h~4(6d_1D^-VBbcfZU z7~Ms6xQXtuI%pjd^Zro>&qHFi3#1+5c6blkzVH}r5SY|()q>Dwz`_~{4>aIdjhMGX z1Ky)Mt^p2*#k}hp5QgrR2CO+O#=G^o!xGjQ*dr3w6?&odfthGSp#-fIZXJ>EmS}?Z zQJj<}c%!S;gbZ{yG+`gQ*P8I`sF+m*vttsxg^xnp0~Vtl3bkmb!Uwcz(7jf|Tcrh4 zYsI_@Em()HP7D4(cUKGe$HlyNTHt{lVV;|7Z`)?VHe0l_oWM*JSk=c@bM(~#yOpm;OD?Jv|+Fw zZ8}^yCE=~n2BkVNZ?87=Lw8mi645nlLpi!n+R#!bW_3aLG`B5&X#GKkb}Ae|I~!h{ zmhhJAfYli>Z<`K;pgX1m%hBD?fjV?Ab>Q2#gv zP@>>(in, out, n); return (int)cudaGetLastError(); } + +extern "C" __global__ void gpud_decimate_kernel( + const float2* __restrict__ in, + float2* __restrict__ out, + int n_out, + int factor +) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= n_out) return; + out[idx] = in[idx * factor]; +} + +extern "C" int gpud_launch_decimate_cuda( + const float2* in, + float2* out, + int n_out, + int factor +) { + if (n_out <= 0 || factor <= 0) return 0; + const int block = 256; + const int grid = (n_out + block - 1) / block; + gpud_decimate_kernel<<>>(in, out, n_out, factor); + return (int)cudaGetLastError(); +}