Pranimi i të gjitha Bluetooth një herë në SDR me CUDA? Lehtë

Në kohët e fundit, kolegët në "zanat" pavarësisht nga njëri-tjetri filluan të më pyesin: si të marr një SDR-pranues njëkohësisht të gjitha kanalet Bluetooth? Rrezja e bandwidth e lejon, ka SDR me një rreze dalëse 80 MHz dhe më shumë. Sigurisht, mund ta bësh këtë në FPGA, por koha për zhvillimin e tillë do të jetë mjaft e madhe. E dija prej një kohe të gjatë që ta bësh këtë në GPU është shumë e lehtë, por që të jetë kështu!

Standardi Bluetooth përcakton nivelin fizik në dy versione: Classic dhe Low Energy. Specifikimi është këtu. Dokumenti është tmerrësisht i madh, ta lexosh tërë atë është e rrezikshme për trurin. Për fat të mirë, kompanitë e mëdha që prodhojnë teknologji matëse kanë mjete për të krijuar dokumente të dukshme mbi temën. Tektronix dhe Instrumentet Nacionale, për shembull. Unë nuk kam asnjë shans të konkurroj me ata në cilësinë e paraqitjes së materialit. Ata të interesuar mund të studiojnë nëpërmjet lidhjeve.

Gjithçka që më duhet të di për nivelin fizik për të krijuar një filtrin shumëkanal është hapi i rrjetës së frekuencave dhe shpejtësia e modulasit. Ato janë përmbledhur në një tabelë në njërin nga dokumentet e përmendura:

Kështu, na nevojitet të ndajmë bandën 80 MHz në 79 filtrat me hap të konfigurimit 1 MHz dhe, në të njëjtën kohë, në 40 filtrat me hap të konfigurimit 2 MHz. Frekuencat e ndjekjes së mostrave nga daljet e filtrave duhet të jenë 1 MHz dhe 2 MHz, përkatësisht.

Kështu, na nevojiten dy grupe filtrash.

Për fillim, do të zgjidhim parametrat e këtyre filtrave duke u bazuar në bandat e sinjaleve Bluetooth Classic dhe Bluetooth Low Energy. Na duhen karakteristikat impulsive të tyre për të llogaritur ngarkesën në pajisjen llogaritëse të filtrit. Këtu, menjëherë duhet të sqarojmë se gjatë e karakteristikave impulsive i kemi zgjedhur bazuar në kërkesat e algoritmit "të shpejtë" të filtrimit. E vërteta nga kjo nuk ndryshon. Dhe numri i koeficenteve të karakteristikës impulsive nuk duhet të jetë shumë i madh, në mënyrë që filtri të realizohet në pajisje llogaritëse të arsyeshme.

Për filtrat me hap 1 MHz, do të zgjedhim bandën e kalimit të LPF (polovja e kalimit të filtrit të bandës) 500 kHz, dhe gjatë e karakteristikës impulsive do të përshtatet në 480 ndarje. Për filtrat me hap 2 MHz, këto parametra do t'i zgjedhim 1 MHz dhe 240 ndarje, përkatësisht. Lloji i dritareve zgjedhim Kaiserin. Do të llogarisim karakteristikat impulsive në filterDesigner dhe do t’i nxjerrim ato në formatin e titullit C:

Pamjet ekranit nga filterDesigner

Mund të zgjidhet problemi në mënyrë direkte: të ndërtohet një array DDC sipas numrit të filtrave (Digital Down Converter). Ky qasje është e mirë për FPGA, ku mund të kursejmë duke reduktuar sasinë e të dhënave të llogaritësve të nivelit të parë. Gjithashtu, FPGA është mënyra më energji-efikase për realizimin e kësaj. Por kostot e punës me këtë qasje janë më të larta.

Me realizimin e filtrave në një grumbull (comb), që është popullor tani në GPU, hapet mundësia për të zbatuar një algoritëm më të sofistikuar: grumbulli polifazik i filtrave bazuar në FFT, i cili në CUDA është në dispozicion nga biblioteka. Në literaturën e huaj, algoritmi quhet Polyphase ose WOLA (Weight, Overlap and Add) FFT Filterbank. Mungesa e dëshirës për të vizatuar nuk më lejon të jap një shpjegim vizual vetë. Në internet ka shumë material mbi këtë temë, sidomos një grafik ilustrues është realizuar këtu në faqe 11 (faleminderit shumë autorëve të nderuar), ja ai:

Do të përpiqem të shpjegoj skemën e përpunimit me fjalët e mia. Të dobëtit në nerva ju lutem mos e lexoni.

Do të përpiqem të shpjegoj skemën e përpunimit me fjalët e mia brenda mundësive të mia metodologjike. FFT është konvolucioni i sinjalit hyrës me të gjithë spektrin e harmonikave komplekse ortogonale, që vendosen në intervalin e karakteristikës së impulsit. Karakteristika e impulsit të filtrit, mbi të cilën shënohet sinjali para hyrjes në FFT, modullohet me këtë spektrat harmonike. Me fjalë të tjera, mbulesa e karakteristikave të impulsit të filtrave rezultat është nxjerrë jashtë. Më pas, spektri i harmonikave filtrat në disa herë, për shkak të zgjerimit të gjatësisë së filtrit në krahasim me filtrin në dritare katrore. Në grafik shohim një filtrim në katër. Me fjalë të tjera, pas zgjerimit të bandës me dritaren Kaiser (me njëkohësisht rritjen e zbehjes në bandën e mbajtjes), ne nuk kemi nevojë për të gjithë filtrat, vetëm katërta e tyre. Të tjerat janë të tepërta, karakteristikat e tyre të frekuencave mbivendosen. Nga katër pika FFT, që vijnë njëra pas tjetres, zgjedhim vetëm pikën zero, llogaritja e së cilës është shumimi i katër
pikave hyrëse, të marra pas një kohe të barabartë me një katërtesë të ndërlikimit fillestar të FFT.

Pajisjen do ta zgjedhim atë që kemi për dorë. Kjo është karta e hyrjes e kompanisë "Sistemat Instrumentale" FMC126P. Për të kam shkruar tashmë në një nga përshkrimet e mia të mëparshme. artikulli ynë. Në slotin FMC të pllakës u insertua një nënmodul i të njëjtës kompani me një transiver AD9371 me një brezi prej 100 MHz. Të gjithë fluksi nga transiveri mund të transmetohet vazhdimisht në kompjuter për përpunim.

Do të zgjidhim një kartë video me GPU GTX 1050. (Kam gënjyer, ajo na zgjodhi: kjo ishte ajo që kishim në dorë, ishte nxjerrë nga llogaritësi për llogaritjen e antennas, por nuk ishte e çuditshme të shihnim një grebëz në punë). Le të kalojmë në pjesën software.

Fatkeqësisht, për shkak të licencave nuk mund të publikojmë kodin e plotë. Mund të tregojmë vetëm bërthamën e GPU. Megjithatë, kodi tjetër nuk është veçanërisht interesant.

Ja bërthama që bën shumëzimin e sinjalit me dritaren dhe mbledhjen, dhe mbështjellja për thirrjen e saj:

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

Kodi për përpunimin e sinjalit që thërret këtë bërthamë përsërit në përpikëri skemën e algoritmit të paraqitur në figurën më sipër, prandaj nuk shoh kuptim të sjell atë këtu.

Grebëzët u verifikuan në spektin e daljes së kanaleve në kohë reale. Në hyrje të AD9371 u dërgua sinjali i gjeneratorit të sinjaleve 2450 MHz, selektiviteti i filtrave përputhej me atë të llogaritur.

Në planet: adaptimi i software-it në pllakën XRTX dhe realizimi i kërkimit të paketave, nëse dikush do të kishte nevojë ose do të kishte kohë të lirë.

Të gjithë punën për softuerin e ka bërë gaudima, lavdim atij!

Burimi: habr.com