🥇Sprejem vseh Bluetooth kanalov hkrati na SDR s CUDA? Zlahka

V zadnjem času so me kolegi iz "industrije" neodvisno spraševali: kako lahko hkrati pridobim vse Bluetooth kanale iz enega SDR sprejemnika? Pasovna širina to omogoča, imamo SDR z izhodno pasovno širino 80 MHz in več. Seveda bi to lahko izvedli na FPGA, a bi bilo za to potrebno precej časa. Že dolgo je znano, da je to na GPU razmeroma enostavno, toda tako enostavno!

Standard Bluetooth določa fizično raven v dveh različicah: Klasični in Nizka Energija. Specifikacija je na voljo. ovdeDokument je izjemno obsežen, branje celote je nevarno za možgane. Na srečo imajo velike podjetja, ki proizvajajo merilno opremo, sredstva za pripravo slikovitih dokumentov na to temo. Tektronix in National Instruments, na primer. Popolnoma nimam možnosti konkurirati njim v kakovosti predstavitve materiala. Tiste, ki jih to zanima, prosim, da preučijo povezave.

Vse, kar potrebujem vedeti o fizični ravni za ustvarjanje večkanalnega filtra, je velikost frekvenčne mreže in hitrost modulacije. Obe jima sta zbrani v tabeli v enem od navedenih dokumentov:

Tako potrebujemo razdeliti pasovno širino 80 MHz na 79 filtrov s korakom nastavitve 1 MHz in hkrati 40 filtrov s korakom nastavitve 2 MHz. Frekvence vzorčenja z izhodov filtrov morajo biti 1 MHz in 2 MHz.

Tako potrebujemo dve filtrovni mreži.

Najprej izberimo parametre teh filtrov na podlagi pasov signalov Bluetooth Klasični in Bluetooth Nizka Energija. Potrebujemo njihove impulzne karakteristike za izračun obremenitve na računalniškem pripomočku filtra. Tukaj je treba takoj pojasniti, da smo dolžine impulznih karakteristik izbrali na podlagi zahtev "hitrega" algoritma filtriranja. Bistvo se ne spremeni. In število koeficientov impulzne karakteristike ne sme biti preveliko, da bo filter mogoče uresničiti na sprejemljivi računalniški opremi.

Za filtre s korakom 1 MHz izberemo pasovno širino FNČ (polovica pasovne širine pasovnega filtra) 500 kHz, dolžino impulzne karakteristike pa prilagodimo na 480 odvodov. Za filtre s korakom 2 MHz izberemo te parametre 1 MHz in 240 odvodov, ustrezno. Tip okna izberemo Kaiser. Izračunajmo impulzne karakteristike v filterDesigner in izvozimo v formatu C-header:

Posnetki iz filterDesigner

Težavo lahko rešimo na neposreden način: zgradimo ustrezen niz DDC (Digital Down Converter). Tak pristop je dober za FPGA, kjer lahko prihranimo zaradi zmanjšanja natančnosti prvih stopenj izračunavanja. Poleg tega so FPGA najbolj energetsko učinkovita metoda implementacije. Vendar so delovne obremenitve pri tej metodi najbolj visoke.

Pri izvajanju črte filtrov na priljubljenih GPU-jih imamo možnost implementirati bolj zapleten algoritem: polifazno črto filtrov na osnovi FFT, ki je na CUDA dostopna iz knjižnice. V tujini se algoritem imenuje Polyphase or WOLA (Weight, Overlap and Add) FFT Filterbank. Lenoba za risanje mi ne dopušča, da bi samostojno izvedel vizualno razlago. Na spletu je veliko gradiv o temi, še posebej je zelo jasen grafikon, ovde ki se nahaja na strani 11 (velika hvala spoštovanim avtorjem), tu je:

Poskusil bom pojasniti shemo obdelave z lastnimi besedami. Tiste, ki imate šibke živce, prosim, da ne berete.

Poskusil bom pojasniti shemo obdelave z lastnimi besedami v okviru svojih metodoloških zmožnosti. FFT je konvolucija vhodnega signala z vsemi spektri kompleksnih ortogonalnih harmonikov, ki se prilegajo na interval impulse. Impulzna odzivnost filtra, na katero se pomnoži signal pred vhodom FFT, se modulira s spektrom harmonikov. Drugimi besedami, obsežna impulznih odzivnosti rezultatnih filtrov črte se izloči. Nato se spekter harmonikov razreže v določeno število delitev, zaradi širjenja prepustnosti filtra v primerjavi s filtrom v pravokotnem oknu. Na risbi vidimo razrez na štiri. Drugimi besedami, po širjenju traku z oknom Kaiser (ob hkratnem povečevanju dušenja v pasu zadržanja), potrebujemo le četrtino vseh filtrov. Ostali so odvečni, njihove frekvenčne lastnosti se prekrivajo. Iz štirih zaporednih točk FFT izberemo le ničlo, katere izračun je vsota štirih
vhodnih točk, vzetih v času, enakem četrtini dolžine izvirnega FFT.

Izbrali bomo strojno opremo, ki je na voljo. To je vhodna plošča podjetja 'Instrumentalne Sisteme' FMC126P. O njej sem že pisal v eni prejšnjih članku. V vtičnici FMC plošče je vstavljena podmodul istega podjetja s pretvornikom AD9371 s pasovno širino 100 MHz. Celoten tok iz pretvornika se lahko neprekinjeno prenese v računalnik za obdelavo.

Izbrali bomo grafično kartico z GPU GTX 1050. (Lažem, ona nas je izbrala: to je bilo vse, kar je bilo pri roki, odvzeta je bila iz računalnika za izračun anten, a toliko bolj presenetljivo je bilo videti delujočo grredno napravo). Pojdimo k programski komponenti.

Na žalost zaradi licenc ne moremo objaviti celotne kode. Lahko pokažemo le jedra GPU. Omeniti velja, da ostala koda ni posebej zanimiva.

Tukaj je jedro, ki izvaja množenje signala z oknom in seštevanje, ter ovitek za njegovo klicanje:

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

Koda za obdelavo signala, ki kliče to jedro, natančno ponavlja shemo algoritma, prikazano na sliki zgoraj, zato ne vidim smisla, da jo navajam tukaj.

Grdne so bile preverjene na izhodnem spektru kanalov v realnem času. Na vhod AD9371 je bil podan signal generatorja signalov 2450 MHz, selektivnost filtrov pa je ustrezala izračunani.