🥇Kogu Bluetooth'i vastuvõtt SDR-i kaudu CUDA-ga? Lihtne

Viimasel ajal on kolleegid minu "valdkonnas" iseseisvalt küsinud: kuidas saada ühest SDR-vastuvõtjast samal ajal kõik Bluetoothi kanalid? Riba võimaldab, SDR-il on väljundiriba 80 MHz ja rohkem. Loomulikult võib selle sees teha FPGA-s, kuid arendusaeg on tõeliselt pikk. Olen ammu saanud teada, et sellise asja tegemine GPU-l on päris lihtne, aga nii lihtsalt!

Bluetoothi standard määratleb füüsilise taseme kahes versioonis: Classic ja Low Energy. Spetsifikatsioon on olemas. siitDokument on kohutavalt mahukas, seda tervikuna lugedes on oht aju jaoks. Õnneks on suurtel mõõteinstrumentide tootjatel vahendeid visuaalsete dokumentide koostamiseks teema osas. Tektronix ja National Instruments, näiteks. Mul pole mingit võimalust nendega kvaliteedi osas konkureerida. Neid, keda huvitab, palun uurida linkide kaudu.

Kõik, mida ma füüsilise taseme kohta pean teadma keeruka filtriga rakendamiseks, on sageduse sammu ja modulatsiooni kiirus. Need on koondatud tabelisse ühes nimetatud dokumentidest:

Seega tuleb meil 80 MHz laius jagada 79 filtriks, mille seadistusaste on 1 MHz, ja samal ajal 40 filtriks, mille seadistusaste on 2 MHz. Filtrite väljundite järjestikuste mõõtmiste sagedused peavad olema vastavalt 1 MHz ja 2 MHz.

Seega vajame kahte filtrite kamm.

Alustame nende filtrite parameetrite valimisega, lähtudes Bluetooth Classic ja Bluetooth Low Energy signaalide ribadest. Me vajame nende impulsiomadusi, et arvestada koormust filtrite arvutusüksusele. Siinkohal on oluline märkida, et impulsiomaduste pikkused valiti 'kiire' filtreerimise algoritmi nõudmiste põhjal. See ei muuda asja olemust. Ja impulsiomaduste koefitsientide arv ei tohi olla liiga suur, et filtrit saaks teostada mõistlikul arvutusriistvaral.

1 MHz sammude filtreid kasutades valime FНЧ ribalaiuseks (ribafiltri ribalaiuse pool) 500 kHz ning impulsskarakteristikud kohandame 480 väljundiga. 2 MHz sammude filtreid kasutades valime vastavalt 1 MHz ja 240 väljundit. Aknatüübiks valime Kaiseri. Arvutame impulsskarakteristikud filterDesigner'is ja laadime need välja C-pealkirja formaadis:

Kuvand filterDesigner'ist

Probleemi saab lahendada otse: ehitada vastavalt filterite arvu DDC massiiv (Digital Down Converter). Selline lähenemine on hea FPGA-de jaoks, kus on võimalik säästa ensimmäise astme kalkulaatorite bitisügavuse vähendamise kaudu. Samuti on FPGA kõige energiatõhusam rakendamise viis. Kuid selle meetodi tööjõu nõudmised on kõige kõrgemad.

Populaarsete GPU-de filtrite grebenki täitmisel avatakse võimalus rakendada keerukamat algoritmi: polüfaasilist filtrite grebenki, mis põhineb FFT-l ja on CUDA kaudu saadaval raamatukogust. Välisliteratuuris nimetatakse algoritmi polüfaasiliseks või WOLA (Weight, Overlap and Add) FFT Filterbank. Joonistamise laiskus ei võimalda mul ise visuaalset selgitust teha. Internetis on palju materjale selle teema kohta, eriti näidatud graafik on esitatud siit leheküljel 11 (suur tänu austatud autoritele), siin see on:

Püüan seletada töötlemisskeemi oma sõnadega. Nõrganärvilisi palun mitte lugeda.

Püüan selgitada töötlemise skeemi oma sõnadega, lähtudes minu pedagoogilistest võimalustest. DFT (diskreetne Fourier' muundamine) on sisendsignaali konvolutsioon kõigi komplekssete ortogonaalsete harmonikate spektriga, mis asetuvad impulsi karakteristikute intervalli. Filtri impulsi karakteristika, millele signaal enne DFT sisestamist korrutatakse, moduleeritakse selle harmoonikate spektriga. Teisisõnu, tulemuseks olevate filtrite impulsi karakteristikute enampakkumise kuju jääb kõrvale. Edasi, harmoonika spekter tihendatakse mõnel määral, kuna filtri läbilaskevõime laieneb võrreldes sirge aknaga filtri läbilaskevõimega. Joonisel näeme tihendamist neljaks. Teisisõnu, pärast akna Kaiser'iga laiendamist (samal ajal suurendades mahasurumist viivituse vööndis) on meil vaja mitte kõiki filtreid, vaid ainult neljandikku neist. Ülejäänud on liigsed, nende sageduskarakteristika kattuvad. Neljast järjestikku järgnevast DFT punktist valime ainult nullpunkti, mille arvutamine on nelja summa.
sisenemise punktidest, mis on võetud aja jooksul, mis on võrdne neljandiku algse DFT kestuse pikkusega.

Valime äriprotsessiks olemasolevat riistvara. See on "Instrument Systems" ettevõtte FMC126P sisendplaat. Olen sellest juba varem kirjutanud. artiklisFMC pesasse on paigaldatud sama ettevõtte alamsüsteem koos AD9371 muunduriga, mille ribalaius on 100 MHz. Kogu signaal muundurist saab pidevalt edastada arvutisse töötlemiseks.

Valime video kaardi, mille GPU on GTX 1050. (Vale, see valis meid: see oli kõik, mis oli käepärast, ja see tõmmati välja antenni kalkuleerimise arvutist, aga see oli üllatav näha töötavat kanalit). Liigume nüüd tarkvara osa juurde.

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

Siin on tuum, mis teostab signaali korrutamist akna ja liitmisega, ning wrapper selle kutsumise jaoks:

__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<<>>(data, window, windowSize, result);
    return cudaGetLastError();
}

Signaali töötlemise kood, mis kutsub seda tuuma, kordab täpselt ülaltoodud algoritmi skeemi, seega ei näe põhjust selle siin näitamiseks.

Harjad on testitud kanalite väljundispektris reaalajas. AD9371-le edastati 2450 MHz signaaligeneraatori signaal, filtrite selektiivsus vastas arvutustele.