Viimasel ajal on kolleegid "harust" üksteiselt sõltumatult küsinud, kuidas saada ühes SDR-vastuvõtjas korraga kõik Bluetoothi kanalid? Bandwidth lubab seda, SDR-id, mille väljundilaius on 80 MHz ja rohkem. Loomulikult, seda saaks teha ka FPGA-l, kuid sellise arenduse aeg oleks üsna pikk. Juba ammu oli mulle teada, et seda on üsna lihtne teha GPU-l, aga nii!
Bluetoothi standard määratleb füüsilise tasandi kahes versioonis: Classic ja Low Energy. Spetsifikatsioon on olemas . Dokument on kohutavalt suur, selle täielik lugemine on ajule ohtlik. Õnneks on suurtel ettevõtetel, mis toodavad mõõtevarustust, ressursid visuaalsete dokumentide koostamiseks teemal. ja , näiteks. Mul ei ole mingit võimalust nendega aine esitlemise kvaliteedis konkureerida. Huvi korral palun uurige linkide kaudu.
Kõik, mida mul on vaja teadma füüsilisest tasandist mitmekanallise filtri loomiseks, on sagedusvõrgu samm ja modulaatorikiirus. Need on kokku pandud tabelisse ühes nimetatud dokumentidest:
Seega peame 80 MHz laiuse jagama 79 filtriks 1 MHz seadistuse sammuga ning samal ajal 40 filtriks 2 MHz seadistuse sammuga. Filtrite väljundite sageduste järgnevate näidikute sagedused peavad olema vastavalt 1 MHz ja 2 MHz.
Seega on meil vaja kahte filtrite kammikut.
Alustame nende filtrite parameetrite valimist, tuginedes Bluetooth Classic ja Bluetooth Low Energy signaalide ribadele. Me vajame nende impulsskarakteristikuid, et arvutada filtri töötlusseadmele koormust. Siinkohal tuleb märkida, et impulsskarakteristikute pikkuste valik põhines "kiire" filterimisalgoritmi nõudmistel. Kuid selle olemus ei muutu. Ja impulsskarakteristiku koefitsientide arv ei tohi olla liiga suur, et filter oleks teostatav põhjendatud arvutitehnoloogial.
1 MHz sammuga filtrite puhul valime madalpääsfiltri läbilaskevõime (pool madalpääsfiltri läbilaskevõimest) 500 kHz, impulsskarakteristika pikkuse kohandame 480 lõiguni. 2 MHz sammuga filtrite puhul valime need parameetrid 1 MHz ja 240 lõiku vastavalt. Aken valitakse Kaiseri tüübiks. Arvutame impulsskarakteristikud filterDesigner'is ja eksportime need C-header formaadis:
FilterDesigner'i ekraanipildid
Probleemi saab lahendada otse: konstrueerida DDC massiiv, mis vastab filtrite arvule (). Selline lähenemine sobib eriti PLISe jaoks, kus on võimalik kokkuhoid arvutite täpsuse vähendamise kaudu. Samuti on PLISe kõige energiaefektiivsem viis teostamiseks. Kuid töömaht selle meetodi puhul on kõrgeim.
Filtrite kammide teostamisel tänapäeval populaarsetel GPU-del tekib võimalus rakendada keerukamat algoritmi: polüfaasilist filtrite kammimist FFT põhjal, mis on CUDA raamatukogus saadaval. Välisliku kirjanduse kohaselt nimetatakse algoritmi Polyphase või WOLA (Weight, Overlap and Add) FFT Filterbank. Joonistamise laiskus ei luba mul iseseisvat visuaalset seletust anda. Internetis on palju materjale, eriti visuaalset graafikut, mis on tehtud lehelt 11 (suur tänu austatud autoritele), siin see on:
Proovin selgitada töötlemise skeemi oma sõnadega. Nõrganärvilised palun mitte lugeda.
Proovin selgitada töötlemise skeemi oma sõnadega oma metodoloogiliste võimaluste piirides. BPF on sissetuleva signaali konvolutsioon kõigi komplekssete ortogonaalharmoonikate spektriga, mis mahuvad impulsi karakteristikas. Filtri impulsi omadus, millele signaal BPF-sse sisestamisel korrutatakse, moduleerib neid harmoonikate spektrit. Teisisõnu, tulemuseks saadud filtrite impulsi omaduste ümbrik viiakse välja tühikutest. Järgmiseks, harmoonikate spektrit õhendatakse mingis arvus, sest filtri lubatud laius laieneb võrreldes otstes oleva ristkülikukujulise aknaga. Joonisel näeme nelja ühitamist. Teisisõnu, pärast akna Kaiser laiemaks muutmist (samal ajal suurendades summutust viivituse valdkonnas) ei ole me enam vajavad kõiki filtreid, vaid ainult neljandiku neist. Ülejäänud on üleliigsed, nende sageduskarakteristikad kattuvad. Neljast järjestikku järgnevast BPF punktist valime ainult nullpunkti, mille arvutamine on nelja
sisendpunkti kokkuvõte, mis on võetud ajavahemikuga, mis on võrdne algse BPF kestuse neljandikuga.
Valime riistvara, mis on käepärast. See on ettevõtte "Instrumentaalne Sistemid" FMC126P sisendplaat. Oleme sellest juba varem kirjutanud FMC paanile on paigaldatud sama ettevõtte submodul, millel on AD9371 edastaja, mille ribalaius on 100 MHz. Kogu voog edastajalt saab pidevalt edastada arvutisse töötlemiseks.
Valime videokaardi GPU GTX 1050. (Vale, see valis meid: see oli kõik, mis käepärast oli, see oli rebenenud antennide arvutist, kuid oli üllatav näha töötavat pihustit). Liigume edasi tarkvarasektorisse.
Kahjuks ei saa me litsentside tõttu avaldada täit koodi. Saame näidata ainult GPU tuumasid. Tõsi, ülejäänud kood pole ka eriti huvitav.
Siin on tuum, mis teostab signaali korrutamise aknaga ja liitmise ning selle väljakutsumise mähis:
__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 ma mõtet tuua seda siin esile.
Pihustid on kontrollitud kanalite väljundispektris reaalajas. AD9371-le anti 2450 MHz signaaligeneraatori signaal, filtrite selektiivsus vastas arvutusele.
Plaanis on: tarkvara kohandamine XRTX plaadile ja pakettide otsing, kui see kellelegi vajalik on või vaba aega jääb.
Kogu tarkvaraga seotud töö tegi , au talle!
Allikas: habr.com
