🥇Příjem všech Bluetooth najednou na SDR s CUDA? Snadno

V poslednej dobe sa moji kolegovia z "odvetvia" nezávisle jeden od druhého začali pýtať, ako získať z jedného SDR prijímača súčasne všetky kanály Bluetooth? Prúžok to umožňuje, je SDR s výstupným pásmom 80 MHz a viac. Je to možné, samozrejme, urobiť na FPGA, ale čas potrebný na takýto vývoj bude pomerne dlhý. Už dávno som vedel, že spraviť to na GPU je celkom jednoduché, ale takto!

Štandard Bluetooth definuje fyzickú úroveň v dvoch verziách: Classic a Low Energy. Špecifikácia existuje tu. Dokument je hrozne veľký, čítať ho celý je nebezpečné pre mozog. Našťastie veľké spoločnosti, výrobci meracej techniky, majú prostriedky na vytvorenie prehľadných dokumentov na tému. Tektronix a National Instruments, napríklad. Nemám vôbec šancu v konkurencii s nimi, pokiaľ ide o kvalitu prezentácie materiálu. Záujemcov prosím, aby sa študovali podľa odkazov.

Všetko, čo potrebujem vedieť o fyzickej úrovni na vytvorenie viackanálového filtra, je krok frekvenčnej mriežky a rýchlosť modulácie. Tieto informácie sú zhrnuté v tabuľke v jednom z uvedených dokumentov:

Takže potrebujeme rozrezať pásmo 80 MHz na 79 filtrov s nastavením 1 MHz a súčasne na 40 filtrov s nastavením 2 MHz. Frekvencie vzorkovania z výstupov filtrov by mali byť 1 MHz a 2 MHz, v súlade s tým.

Takže potrebujeme dva filtračné hrebeňové sytemy.

Na začiatok si zvolíme parametre týchto filtrov na základe pásiem signálov Bluetooth Classic a Bluetooth Low Energy. Potrebujeme ich impulzové charakteristiky, aby sme vypočítali záťaž na výpočtové zariadenie filtra. Je hneď potrebné poznamenať, že dĺžky impulzových charakteristík sme zvolili na základe požiadaviek "rýchleho" filtračného algoritmu. Podstata sa tým nemení. A počet koeficientov impulzovej charakteristiky by nemal byť príliš veľký, aby bol filter realizovateľný na rozumnom výpočtovom zariadení.

Pre filtre s krokom 1 MHz vyberieme pásmo prechodu FNO (polovica pásma prechodu pásmového filtra) 500 kHz, dĺžku impulzovej charakteristiky prispôsobíme na 480 odvodov. Pre filtre s krokom 2 MHz tieto parametre vyberieme 1 MHz a 240 odvodov, v súlade s tým. Typ okna vyberáme Kaiser. Vypočítame impulzové charakteristiky v filterDesigner a exportujeme ich vo formáte C-header:

Snímky z filterDesigner

Úlohu možno riešiť priamočiarým spôsobom: vytvoriť maticu DDC zodpovedajúcu počtu filtrov (Digitálny zosilňovač signálu). Tento prístup je dobrý pre FPGA, kde je možné ušetriť znížením bitovej hĺbky počítačov prvého stupňa. Tiež FPGA predstavuje najenergetickejšie efektívny spôsob realizácie. Avšak práca s týmto prístupom je najnáročnejšia.

Pri vytváraní filtrujúcej gríbky na populárnych GPU dnes vzniká možnosť realizovať sofistikovanejší algoritmus: polyfázovú filtrovanie gríbky na základe FFT, ktorý je v CUDA dostupný z knižnice. V zahraničnej literatúre je algoritmus známy ako Polyphase alebo WOLA (Weight, Overlap and Add) FFT Filterbank. Moja lenivosť k grafickému vykresleniu mi neumožňuje vykonať názorné vysvetlenie samostatne. Na internete je mnoho materiálu na túto tému, obzvlášť názorný graf je vyhotovený tu na strane 11 (veľké ďakujem váženým autorom), tu je:

Pokúsim sa vysvetliť schému spracovania svojimi slovami. Slabšie povahy prosím, aby nečítali.

Pokúsim sa vysvetliť schému spracovania svojimi slovami v rámci svojich metodických možností. FFT predstavuje konvolúciu vstupného signálu so všetkým spektrom komplexných ortogonálnych harmoník, ktoré sa uložia na interval impulznej charakteristiky. Impulzná charakteristika filtra, na ktorú sa násobí signál pred vstupom do FFT, je modulovaná týmto spektrom harmoník. Inými slovami, obálka impulzných charakteristík výsledných filtrov gríbky je vyňatá z úvahy. Následne je spektrum harmoník zúžené niekoľkokrát, vzhľadom na rozšírenie pásma prenosu filtra v porovnaní s filtrom v obdĺžnikovom okne. Na obrázku vidíme zúženie na štyri. Inými slovami, po rozšírení pásma pomocou Kaiserovho okna (s súčasným zvýšením útlmu v pásme zdržania), už nepotrebujeme všetky filtre, len štvrtinu z nich. Ostatné sú nadbytočné, ich frekvenčné charakteristiky sa prekrývajú. Z štyroch bodov FFT, ktoré idú za sebou, vyberáme len nulový, ktorého výpočet ako súčet štyroch
vstupných bodov, ktoré sú brané s časom rovnakým štvrtine trvania pôvodného FFT.

Vyberieme si hardvér, ktorý máme po ruke. Je to vstupná doska spoločnosti "Inštrumentálne Systémy" FMC126P. O nej som už písal v jednej predchádzajúcej článku. Do slotu FMC byl zasunut submodul od stejné společnosti s převodníkem AD9371 s šířkou pásma 100 MHz. Celý tok z převodníku může být nepřetržitě předán do počítače k dalšímu zpracování.

Vybereme grafickou kartu s GPU GTX 1050. (Lhal jsem, ona si vybrala nás: tohle bylo vše, co bylo po ruce, byla vytržena z výpočetního zařízení pro výpočet antén, ale tím více nás překvapilo, že vidíme funkční hřebínek). Přesuňme se k softwarové části.

Bohužel kvůli licencím nemůžeme zveřejnit celý kód. Můžeme ukázat pouze jádra GPU. Nicméně zbytek kódu není tolik zajímavý.

Toto je jádro, které provádí násobení signálu oknem a sčítání, a obal pro jeho volání:

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

Kód zpracování signálu, který toto jádro volá, přesně kopíruje schéma algoritmu zobrazené na obrázku výše, proto nemám smysl jej zde uvádět.

Hřebínky byly testovány na výstupním spektru kanálů v reálném čase. Na vstup AD9371 byl dodáván signál generátoru 2450 MHz, selektivita filtrů odpovídala výpočtu.