🥇Kas võtad CUDA abil SDR-i kaudu vastu kogu Bluetoothi korraga? Lihtne

Hiljuti hakkasid kolleegid poes minult üksteisest sõltumatult küsima: kuidas saada kõik Bluetoothi kanalid korraga ühest SDR-vastuvõtjast? Ribalaius seda võimaldab, on SDR-e, mille väljundribalaius on 80 MHz või rohkem. Muidugi saate seda teha FPGA-ga, kuid selline arendusaeg on üsna pikk. Olen juba pikka aega teadnud, et seda GPU-ga teha on üsna lihtne, aga teha seda nii!

Bluetoothi standard määratleb füüsilise kihi kahes versioonis: Classic ja Low Energy. Spetsifikatsioon saadaval siin. Dokument on kohutavalt suur, selle tervikuna lugemine on ajule ohtlik. Õnneks on suurtel mõõtetehnoloogiaettevõtetel vahendid mingi teema kohta visuaalsete dokumentide loomiseks. Tektronix и Riiklikud instrumendid, Näiteks. Materjali esitamise kvaliteedi osas pole mul absoluutselt mingit võimalust nendega konkureerida. Huvi korral järgige linke.

Kõik, mida ma pean füüsilise kihi kohta mitme kanaliga filtri loomiseks teadma, on sagedusvõrgu vahekaugus ja modulatsioonikiirus. Need on tabelina ühes järgmistest dokumentidest:

Seega peame 80 MHz riba lõikama 79 filtriks häälestusastmega 1 MHz ja samal ajal 40 filtriks häälestusastmega 2 MHz. Filtri väljundi diskreetimissagedused peaksid olema vastavalt 1 MHz ja 2 MHz.

Seega vajame kahte filtripanka.

Esmalt valime nende filtrite parameetrid Bluetooth Classic ja Bluetooth Low Energy signaaliribade põhjal. Filtri arvutusseadme koormuse arvutamiseks vajame nende impulssreaktsioone. Siinkohal tasub kohe mainida, et impulssreaktsiooni pikkused valisime “kiire” filtreerimisalgoritmi nõuete alusel. Sisu sellest ei muutu. Ja impulssreaktsiooni koefitsientide arv ei tohiks olla liiga suur, et filtrit saaks rakendada mõistlikel arvutusseadmetel.

1 MHz sammuga filtrite jaoks valime madalpääsfiltri ribalaiuseks (pool ribapääsfiltri ribalaiusest) 500 kHz ja reguleerime impulssreaktsiooni pikkuseks 480 puudutust. 2 MHz sammuga filtrite puhul valime need parameetrid vastavalt 1 MHz ja 240 kraani. Valime Kaiseri aknatüübi. Arvutame filterDesigneris impulssvastused ja laadime need üles C-päise vormingus:

FiltriDesigneri ekraanipildid

Saate probleemi lahendada jõhkra jõuga: koostage filtrite arvule vastav DDC massiiv (Digital Down Converter). Selline lähenemine on hea FPGA-de puhul, kus esimese astme arvutite bitimahtu vähendades on võimalik säästa. Samuti on FPGA kõige energiasäästlikum rakendusmeetod. Kuid selle meetodi tööjõukulud on kõige kõrgemad.

Tänapäeva populaarsetel GPU-del filtripanga käitamisel on võimalik rakendada keerukamat algoritmi: FFT-l põhinevat mitmefaasilist filtripanka, mis CUDA puhul on raamatukogust saadaval. Väliskirjanduses nimetatakse algoritmi Polyphase ehk WOLA (Weight, Overlap and Add) FFT Filterbank. Laiskus joonistamisel ei anna mulle võimalust visuaalset selgitust iseseisvalt lõpetada. Internetis on teema kohta palju materjali, tehakse eriti selge graafik siin leheküljel 11 (suur tänu lugupeetud autoritele), siin see on:

Püüan töötlemisskeemi oma sõnadega selgitada. Palun ärge lugege nõrganärvilistele.

Üritan oma metoodiliste võimaluste raames töötlusskeemi oma sõnadega lahti seletada. FFT on sisendsignaali konvolutsioon koos kogu komplekssete ortogonaalsete harmooniliste spektriga, mis mahuvad impulssreaktsiooni intervalli. Filtri impulssreaktsiooni, millega signaal korrutatakse enne FFT-sisendit, moduleeritakse selle harmoonilise spektriga. Teisisõnu, saadud kammfiltrite impulssreaktsiooni mähis on sulgudes. Lisaks hõreneb harmooniline spekter teatud arv kordi, kuna filtri ribalaius ristkülikukujulises aknas filtri suhtes laieneb. Pildil näeme hõrenemist nelja võrra. Teisisõnu, pärast riba laiendamist Kaiseri aknaga (sammulise sumbumise suurenemisega stopperis) ei vaja me enam kõiki filtreid, vaid ainult veerandit neist. Ülejäänud on üleliigsed, nende sagedusomadused kattuvad. Neljast järjestikusest FFT-punktist valime ainult null ühe, mille arvutamine on nelja
sisendpunktid, mis on võetud pärast aega, mis on võrdne veerandiga algse FFT kestusest.

Valime käepärast triikraua. See on Instrumental Systemsi FMC126P sisendplaat. Kirjutasin sellest juba ühes eelnevas siit. Plaadi FMC-pistikusse on sisestatud sama firma alammoodul AD9371 transiiveriga 100 MHz sagedusalaga. Kogu transiiveri voogu saab töötlemiseks pidevalt arvutisse edastada.

Valime videokaardi GPU GTX 1050. (Ma valetasin, see valis meid: see oli kõik, mis oli käepärast, see rebiti arvutist välja, et antenne arvutada, kuid seda üllatavam oli näha töötav kamm). Liigume edasi tarkvara osa juurde.

Kahjuks ei saa me litsentside tõttu täielikku koodi avaldada. Saame näidata ainult GPU tuumasid. Ülejäänud kood pole aga eriti huvitav.

Siin on tuum, mis teeb signaaliakna korrutamist ja liitmist, ning ümbris selle nimetamiseks:

__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();
}

Seda tuuma kutsuv signaalitöötluskood kordab täpselt ülaltoodud joonisel näidatud algoritmi diagrammi, nii et ma ei näe põhjust seda siin esitada.

Kammi testiti kanalite väljundspektril reaalajas. AD9371 sisendisse anti 2450 MHz signaaligeneraatori signaal ja filtrite selektiivsus vastas arvutatule.