Onlangs het kollegas in die "winkel" my onafhanklik van mekaar begin vra: hoe om alle Bluetooth-kanale gelyktydig van een SDR-ontvanger af te kry? Die bandwydte laat dit toe, daar is SDR's met 'n uitsetbandwydte van 80 MHz of meer. Jy kan dit natuurlik op 'n FPGA doen, maar sulke ontwikkelingstyd sal redelik lank wees. Ek weet al lankal dat dit redelik eenvoudig is om dit op 'n GPU te doen, maar om dit so te doen!
Die Bluetooth-standaard definieer 'n fisiese laag in twee weergawes: Classic en Low Energy. Spesifikasie beskikbaar . Die dokument is vreeslik groot, om dit in sy geheel te lees is gevaarlik vir die brein. Gelukkig het groot meettegnologiemaatskappye die middele om visuele dokumente oor 'n onderwerp te skep. и , Byvoorbeeld. Ek het absoluut geen kans om met hulle mee te ding in terme van die kwaliteit van aanbieding van die materiaal nie. As jy belangstel, volg asseblief die skakels.
Al wat ek moet weet oor die fisiese laag om 'n multi-kanaal filter te skep, is die frekwensie rooster spasiëring en modulasie tempo. Hulle word in een van die volgende dokumente getabelleer:
Ons moet dus die 80 MHz-band in 79 filters sny met 'n instelstap van 1 MHz en terselfdertyd in 40 filters met 'n instelstap van 2 MHz. Die steekproeffrekwensies vanaf die filteruitsette moet onderskeidelik 1 MHz en 2 MHz wees.
Ons benodig dus twee filterbanke.
Kom ons kies eers die parameters van hierdie filters gebaseer op die Bluetooth Classic en Bluetooth Low Energy seinbande. Ons het hul impulsreaksies nodig om die las op die filter se rekenaartoestel te bereken. Hier is dit die moeite werd om dadelik te noem dat ons die impulsresponslengtes gekies het op grond van die vereistes van 'n "vinnige" filteralgoritme. Die wese verander nie hiervan nie. En die aantal impulsresponskoëffisiënte moet nie te groot wees vir die filter om op gesonde rekenaartoerusting geïmplementeer te word nie.
Vir filters met 'n stap van 1 MHz, kies ons 'n laagdeurlaatfilterbandwydte (die helfte van die bandwydte van 'n banddeurlaatfilter) van 500 kHz, en pas die impulsresponslengte na 480 krane aan. Vir filters met 'n stap van 2 MHz, sal ons hierdie parameters as onderskeidelik 1 MHz en 240 taps kies. Ons kies Kaiser-venstertipe. Kom ons bereken die impulsresponse in filterDesigner en laai dit in C-header-formaat op:
Skermkiekies van filterDesigner
Jy kan die probleem op 'n brute-krag manier oplos: bou 'n DDC-skikking wat ooreenstem met die aantal filters (). Hierdie benadering is goed vir FPGA's, waar besparings moontlik is deur die biskapasiteit van die eerste-fase rekenaars te verminder. FPGA is ook die mees energiedoeltreffende implementeringsmetode. Maar arbeidskoste met hierdie metode is die hoogste.
Wanneer 'n filterbank op vandag se gewilde GPU's bestuur word, word dit moontlik om 'n meer gesofistikeerde algoritme te implementeer: 'n polifase-filterbank gebaseer op FFT, wat op CUDA by die biblioteek beskikbaar is. In buitelandse literatuur word die algoritme Polyphase of WOLA (Weight, Overlap and Add) FFT Filterbank genoem. Luiheid in teken gee my nie die geleentheid om 'n visuele verduideliking op my eie te voltooi nie. Daar is baie materiaal oor die onderwerp op die internet, 'n besonder duidelike grafiek word gemaak op bladsy 11 (baie dankie aan die gerespekteerde skrywers), hier is dit:
Ek sal probeer om die verwerkingskema in my eie woorde te verduidelik. Moet asseblief nie vir dowwe van hart lees nie.
Ek sal probeer om die verwerkingskema in my eie woorde te verduidelik binne die raamwerk van my metodologiese vermoëns. Die FFT is die konvolusie van die insetsein met die hele spektrum van komplekse ortogonale harmonieke wat binne die impulsrespons-interval pas. Die impulsrespons van die filter, waarmee die sein voor die FFT-invoer vermenigvuldig word, word deur hierdie harmoniese spektra gemoduleer. Met ander woorde, die impulsresponsomhulsel van die resulterende kamfilters is tussen hakies. Verder word die harmoniese spektrum met 'n sekere aantal kere uitgedun, as gevolg van die uitbreiding van die filterbandwydte relatief tot die filter in die reghoekige venster. Op die foto sien ons 'n verdunning met vier. Met ander woorde, nadat ons die band met die Kaiser-venster uitgebrei het (met 'n gelyktydige toename in verswakking in die stopband), het ons nie meer al die filters nodig nie, maar slegs 'n kwart daarvan. Die res is oorbodig, hul frekwensie-eienskappe oorvleuel. Van die vier FFT punte in 'n ry, kies ons slegs die nul een, waarvan die berekening die som van vier is
insetpunte geneem na 'n tyd gelykstaande aan 'n kwart van die duur van die oorspronklike FFT.
Ons sal die yster kies wat byderhand is. Dit is die invoerbord van Instrumental Systems FMC126P. Ek het reeds in een vorige daaroor geskryf . 'n Submodule van dieselfde maatskappy met 'n AD9371-senderontvanger met 'n 100 MHz-band word in die FMC-konnektor van die bord geplaas. Die hele stroom vanaf die transceiver kan deurlopend na 'n rekenaar oorgedra word vir verwerking.
Kom ons kies 'n videokaart met GPU GTX 1050. (Ek het gelieg, dit was die een wat ons gekies het: dit was al wat byderhand was, dit is uit die rekenaar geskeur om antennas te bereken, maar dit was des te meer verbasend om te sien 'n werkende kam). Kom ons gaan aan na die sagteware deel.
Ongelukkig kan ons weens lisensies nie die volledige kode publiseer nie. Ons kan slegs GPU-kerne wys. Die res van die kode is egter nie besonder interessant nie.
Hier is die kern wat seinvenster vermenigvuldiging en optelling doen, en 'n omhulsel om dit te noem:
__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();
}
Die seinverwerkingskode wat hierdie kern noem, herhaal presies die algoritmediagram wat in die figuur hierbo getoon word, so ek sien geen rede om dit hier aan te bied nie.
Die kamme is intyds op die uitsetspektrum van die kanale getoets. 'n Seingeneratorsein van 9371 MHz is aan die inset van die AD2450 gelewer; die selektiwiteit van die filters stem ooreen met die berekende een.
Die planne sluit in: die aanpassing van die sagteware by die XRTX-bord en die implementering van 'n pakketsoektog, indien iemand dit nodig het of vrye tyd het.
Al die sagteware werk voltooi , eer aan hom!
Bron: will.com
