🥇 Að fá allt Bluetooth í einu á SDR með CUDA? Auðvelt

Nýlega fóru samstarfsmenn í „búðinni“ að spyrja mig óháð hver öðrum: hvernig á að fá allar Bluetooth rásir samtímis frá einum SDR móttakara? Bandbreiddin gerir það kleift, það eru SDR með úttaksbandbreidd 80 MHz eða meira. Þú getur auðvitað gert þetta á FPGA en slíkur þróunartími verður frekar langur. Ég hef lengi vitað að það er frekar einfalt að gera þetta á GPU, en að gera þetta svona!

Bluetooth staðallinn skilgreinir líkamlegt lag í tveimur útgáfum: Classic og Low Energy. Forskrift í boði hér. Skjalið er hræðilega stórt, að lesa það í heild sinni er hættulegt fyrir heilann. Sem betur fer hafa stór mælitæknifyrirtæki möguleika á að búa til sjónræn skjöl um efni. Tektronix и Þjóðfæri, Til dæmis. Ég á nákvæmlega enga möguleika á að keppa við þá hvað varðar gæði framsetningar efnisins. Ef þú hefur áhuga, vinsamlegast fylgdu krækjunum.

Allt sem ég þarf að vita um líkamlega lagið til að búa til fjölrása síu er tíðnigrindarbilið og mótunarhraði. Þau eru sett í töflu í einu af eftirfarandi skjölum:

Þannig þurfum við að klippa 80 MHz bandið niður í 79 síur með 1 MHz stillingarþrepi og á sama tíma í 40 síur með 2 MHz stillingarþrepi. Sýnatökutíðnin frá síuúttakunum ætti að vera 1 MHz og 2 MHz, í sömu röð.

Þannig þurfum við tvo síubanka.

Í fyrsta lagi skulum við velja færibreytur þessara sía út frá Bluetooth Classic og Bluetooth Low Energy merkjaböndunum. Við þurfum hvataviðbrögð þeirra til að reikna út álagið á tölvubúnað síunnar. Hér er rétt að minnast strax á að við völdum lengdar hvataviðbragða út frá kröfum „hraðs“ síunaralgríms. Kjarninn breytist ekki úr þessu. Og fjöldi hvatsviðbragðsstuðla ætti ekki að vera of stór til að hægt sé að útfæra síuna á heilbrigðum tölvubúnaði.

Fyrir síur með þrepið 1 MHz veljum við lágrásarsíubandbreidd (helming bandbreidd bandvíddarsíu) 500 kHz, og stillum lengd höggsvörunar í 480 tappa. Fyrir síur með þrepið 2 MHz, munum við velja þessar breytur sem 1 MHz og 240 tappa, í sömu röð. Við veljum Kaiser gluggagerð. Við skulum reikna út hvataviðbrögðin í filterDesigner og hlaða þeim upp á C-hausformi:

Skjáskot frá filterDesigner

Þú getur leyst vandamálið á grófan hátt: smíðað DDC fylki sem samsvarar fjölda sía (Digital Down Converter). Þessi nálgun er góð fyrir FPGA, þar sem sparnaður er mögulegur með því að draga úr bitagetu fyrstu stigs tölva. Einnig er FPGA orkunýtnasta útfærsluaðferðin. En launakostnaður með þessari aðferð er hæstur.

Þegar þú keyrir síubanka á vinsælum GPU í dag, verður mögulegt að innleiða flóknari reiknirit: fjölfasa síubanka sem byggir á FFT, sem á CUDA er fáanlegur á bókasafninu. Í erlendum bókmenntum er reikniritið kallað Polyphase eða WOLA (Weight, Overlap and Add) FFT Filterbank. Leti við að teikna gefur mér ekki tækifæri til að klára sjónræna útskýringu á eigin spýtur. Mikið efni er til um efnið á netinu, sérstaklega skýrt línurit er gert hér á blaðsíðu 11 (með þökk sé virtum höfundum), hér er hún:

Ég mun reyna að útskýra vinnslukerfið með mínum eigin orðum. Vinsamlegast ekki lesa fyrir viðkvæma.

Ég mun reyna að útskýra vinnslukerfið með mínum eigin orðum innan ramma aðferðafræðilegrar getu minnar. FFT er snúningur inntaksmerkisins með öllu litrófinu af flóknum hornréttum harmonikum sem passa innan boðsvörunarbilsins. Stuðssvörun síunnar, sem merkið er margfaldað með fyrir FFT-inntakið, er mótað af þessu harmónísku litrófi. Með öðrum orðum, hvatsviðbragðshlíf greiðsíanna sem myndast er í sviga. Ennfremur er harmóníska litrófið þynnt út um ákveðinn fjölda sinnum, vegna stækkunar á bandbreidd síunnar miðað við síuna í rétthyrndum glugganum. Á myndinni sjáum við þynningu um fjóra. Með öðrum orðum, eftir að hafa stækkað bandið með Kaiser glugganum (með samtímis aukningu á dempun í stöðvunarbandinu), þurfum við ekki lengur allar síurnar, heldur aðeins fjórðung þeirra. Restin er óþarfi, tíðnieiginleikar þeirra skarast. Af fjórum FFT punktum í röð veljum við aðeins núllpunktinn, útreikningurinn á honum er samantekt á fjórum
inntakspunktar teknir eftir tíma sem jafngildir fjórðungi af tímalengd upprunalega FFT.

Við munum velja járnið sem er við höndina. Þetta er inntakspjaldið frá Instrumental Systems FMC126P. Ég skrifaði þegar um það í einu á undan grein. Undireining frá sama fyrirtæki með AD9371 senditæki með 100 MHz bandi er sett í FMC tengi borðsins. Hægt er að senda allan strauminn frá senditækinu stöðugt í tölvu til vinnslu.

Við skulum velja skjákort með GPU GTX 1050. (Ég laug, það var það sem valdi okkur: það var allt sem var við höndina, það var rifið úr tölvunni til að reikna út loftnet, en það kom þeim mun meira á óvart að sjá vinnandi greiða). Við skulum halda áfram að hugbúnaðarhlutanum.

Því miður, vegna leyfis, getum við ekki birt allan kóðann. Við getum aðeins sýnt GPU kjarna. Hins vegar er restin af kóðanum ekki sérstaklega áhugaverð.

Hér er kjarninn sem gerir merkja-glugga margföldun og samlagningu, og umbúðir til að kalla það:

__global__ void cuComplexMultiplyWindowKernel(const cuComplex *data, const float *window, size_t windowSize, cuComplex *result) {
    __shared__ cuComplex multiplicationResult[480];
    multiplicationResult[threadIdx.x] = cuComplexMultiplyFloat(data[threadIdx.x + windowSize / 4 * blockIdx.x], window[threadIdx.x]);
    __syncthreads();
    cuComplex sum;
    sum.x = sum.y = 0;
    if (threadIdx.x < windowSize / 4) {
        for(int i = 0; i < 4; i++) {
            sum = cuComplexAdd(sum, multiplicationResult[threadIdx.x + i * windowSize / 4]);
        }
        result[threadIdx.x + windowSize / 4 * blockIdx.x] = sum;
    }
}

cudaError_t cuComplexMultiplyWindow(const cuComplex *data, const float *window, size_t windowSize, cuComplex *result, size_t dataSize, cudaStream_t stream) {
    size_t windowStep = windowSize / 4;
    cuComplexMultiplyWindowKernel<<<dataSize / windowStep - 3, windowSize, 1024, stream>>>(data, window, windowSize, result);
    return cudaGetLastError();
}

Merkjavinnslukóði sem kallar þennan kjarna endurtekur nákvæmlega reikniritið sem sýnt er á myndinni hér að ofan, svo ég sé enga ástæðu til að setja það fram hér.

Greiðarnar voru prófaðar á úttaksróf rásanna í rauntíma. Merkjagjafarmerki upp á 9371 MHz var gefið til inntaks á AD2450; valmöguleiki síanna samsvaraði þeim útreiknuðu.