🥇Primirea tuturor canalelor Bluetooth pe SDR cu CUDA? Ușor

Recent, colegii din „industrie” m-au întrebat în mod independent: cum putem obține toate canalele Bluetooth de la un singur receptor SDR? Lățimea de bandă permite acest lucru, avem SDR-uri cu o lățime de bandă de ieșire de 80 MHz sau mai mult. Desigur, se poate face și pe FPGA, dar timpul necesar pentru o astfel de dezvoltare ar fi considerabil. Știam de ceva vreme că este destul de simplu să facem asta pe GPU, dar chiar așa!

Standardul Bluetooth definește nivelul fizic în două versiuni: Classic și Low Energy. Specificația există. aiciDocumentul este extrem de voluminos, citirea lui în întregime poate fi dăunătoare pentru minte. Din fericire, companiile mari care produc echipamente de măsură au resursele necesare pentru a crea documente ilustrative pe această temă. Tektronix și National Instruments, de exemplu. Nu am șanse deloc să concurez cu ele în ceea ce privește calitatea prezentării materialului. Cei interesați sunt invitați să studieze linkurile.

Tot ce trebuie să știu despre nivelul fizic pentru a crea un filtru multicanal este pasul de rețea în frecvență și viteza de modulare. Acestea sunt sumarizate într-un tabel din unul dintre documentele menționate:

Astfel, trebuie să segmentăm lățimea de bandă de 80 MHz în 79 de filtre cu pas de ajustare de 1 MHz și, simultan, în 40 de filtre cu pas de ajustare de 2 MHz. Frecvențele de eșantionare de la ieșirile filtrelor trebuie să fie de 1 MHz și 2 MHz, respectiv.

Așadar, avem nevoie de două grupe de filtre.

Pentru început, să alegem parametrii acestor filtre bazându-ne pe benzile de frecvență ale Bluetooth Classic și Bluetooth Low Energy. Avem nevoie de caracteristicile impulsurilor lor pentru a calcula sarcina pe dispozitivul de calcul al filtrului. Aici trebuie menționat că lungimile caracteristicilor impulsurilor au fost alese în funcție de cerințele unui algoritm de filtrare „rapid”. Esența nu se schimbă. Numărul coeficientului caracteristicii impulsului nu trebuie să fie prea mare, pentru ca filtrul să fie realizabil pe o hardware de calcul rezonabil.

Pentru filtrele cu pas de 1 MHz, vom alege lățimea de bandă a filtrelor trece-jos (jumătate din lățimea de bandă a filtrului pe bandă) de 500 kHz, iar lungimea caracteristicii impulsului va fi ajustată la 480 coeficienți. Pentru filtrele cu pas de 2 MHz, acești parametri vor fi aleși la 1 MHz și 240 coeficienți, respectiv. Tipul feronului ales este al lui Kaiser. Vom calcula caracteristicile impulsului în filterDesigner și le vom exporta în format C-header:

Capturi de ecran din filterDesigner

Se poate aborda problema printr-o metodă directă: a construi un vector DDC corespunzător numărului de filtre (Digital Down Converter). Această abordare este bună pentru FPGA, unde se poate economisi datorită reducerii dimensiunii unităților de calcul de primă etapă. De asemenea, FPGA este cea mai energetică metodă de implementare. Dar eforturile necesare pentru această metodă sunt cele mai mari.

În realizarea grilelor de filtre pe GPU-urile populare, se deschide posibilitatea de a implementa un algoritm mai sofisticat: grila de filtre polifazice bazată pe FFT, disponibilă pe CUDA din bibliotecă. În literatura străină, algoritmul este numit Polyphase or WOLA (Weight, Overlap and Add) FFT Filterbank. Leneș la desen, nu am posibilitatea de a oferi o explicație vizuală. Există multe materiale pe această temă pe internet, în special un grafic ilustrativ este realizat aici pe pagina 11 (multe mulțumiri respectuos autorilor), iată-l:

Voi încerca să explic schema de procesare în cuvintele mele. Cei sensibili sunt rugați să nu citească.

Voi încerca să explic schema de procesare în cuvintele mele în cadrul posibilităților mele metodologice. FFT este convoluția semnalului de intrare cu tot spectrul armonic ortogonal complex, așezat pe intervalul caracteristicii impulse ale filtrului. Caracteristica de impuls a filtrului, pe care este multiplicat semnalul înainte de intrarea în FFT, este modulată de acest spectru de armonici. Cu alte cuvinte, învelișul caracteristicilor impulsive ale filtrelor rezultante este scos din paranteze. În continuare, spectrul armonic este subțiat de un anumit număr de ori, din cauza extinderii benzii de trecere a filtrului în raport cu filtrul într-un fereastră dreptunghiulară. În figura vedem subțierea de patru ori. Cu alte cuvinte, după extinderea benzii cu fereastra Kaiser (cu o creștere simultană a atenuării în banda de întârziere), avem nevoie doar de o pătrime din filtre. Restul sunt excesive, caracteristicile lor frecvenței se suprapun. Din cele patru puncte FFT consecutive alegem doar zeroul, calculul căruia este suma celor patru
puncte de intrare, preluate la un interval egal cu un sfert din durata FFT inițial.

Vom alege hardware-ul care există la îndemână. Este placa de intrare a companiei "Instrumental Systems" FMC126P. Am scris deja despre aceasta într-unul din articolele anterioare. pe care l-ați cititÎn slotul FMC al plăcii este inserat un submodul de aceleași companie cu un transceiver AD9371 cu o lățime de bandă de 100 MHz. Întreaga flux din transceiver poate fi transmis continuu la computer pentru procesare.

Vom alege o placă video cu GPU GTX 1050. (Am mințit, ea ne-a ales pe noi: aceasta era tot ce aveam la îndemână, a fost scoasă din calculatorul pentru calculul antenelor, dar a fost cu atât mai surprinzător să vedem un spectru funcțional). Să trecem la partea software.

Din păcate, din cauza licențelor, nu putem publica codul complet. Putem arăta doar nucleele GPU. Totuși, restul codului nu este foarte interesant.

Iată nucleul care efectuează multiplicarea semnalului cu fereastra și adunarea, precum și wrapper-ul pentru apelul său:

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

Codul de procesare a semnalului care apelează acest nucleu repetă exact schema algoritmului, ilustrarea din imaginea de mai sus, așa că nu văd sensul de a-l prezenta aici.

Spectrele au fost verificate în timp real pe ieșirea canalelor. La intrarea AD9371 a fost aplicat un semnal de generator de semnale de 2450 MHz, selectivitatea filtrului corespundea celei calculate.