🥇A teljes Bluetooth egyszerre történő fogadása SDR-en CUDA-val? Könnyen

Mostanában a "szakma" kollegái egymástól függetlenül kezdtek el kérdezni arról, hogyan lehet egyszerre minden Bluetooth csatornát egy SDR-vevőből megkapni? Hiszen a sávszélesség ezt lehetővé teszi, vannak SDR-ek 80 MHz-es vagy annál nagyobb kimeneti sávszélességgel. Persze meg lehetne ezt oldani FPGA-n, de az ilyen fejlesztéshez szükséges idő elég hosszú lenne. Régóta tudtam, hogy GPU-n viszonylag egyszerű megvalósítani, de hogy ennyire!

A Bluetooth szabvány két verzióban definiálja a fizikai réteget: Classic és Low Energy. A specifikáció elérhető. ittA dokumentum borzasztóan nagy, teljes egészében való olvasása lehetetlen a agy számára. Szerencsére a nagy mérőműszer gyártó cégek rendelkeznek eszközökkel a téma szemléletes dokumentumainak létrehozásához. Tektronix és National Instruments, például. Nincs esélyem a versenyben velük a bemutatott anyag minőségét illetően. Az érdeklődőket arra kérem, hogy a megadott linkeken tanulmányozzák.

Minden, amire szükségem van a fizikai rétegről a többcsatornás szűrő kialakításához, az a frekvenciaméret léptéke és a modulációs sebesség. Ezeket egy táblázatban összegzik az egyik megadott dokumentumban:

Így szükségünk van a 80 MHz-es sáv 79 szűrőre 1 MHz-es hangolási lépéssel, és egyidejűleg 40 szűrőre 2 MHz-es lépéssel. A szűrők kimeneti frekvencia mintavételezési sebessége 1 MHz és 2 MHz kell legyen, megfelelően.

Így két szűrőfokozat szükséges.

Kezdésként válasszuk ki ezeknek a szűrőknek a paramétereit a Bluetooth Classic és Bluetooth Low Energy jelek sávjából. Szükségünk van az impulzuskarakterisztikáikra a szűrő számítási terhelésének meghatározásához. Azonnal meg kell említeni, hogy az impulzuskarakterisztikák hossza a "gyors" szűrési algoritmus követelményei szerint lett kiválasztva. A lényegen ez nem változtat. Az impulzuskarakterisztika koefficiensseinek száma nem lehet túl nagy, hogy a szűrőt ésszerű számítási hardveren lehessen megvalósítani.

Az 1 MHz-es lépésű szűrők számára válasszunk 500 kHz-es aluláteresztő sávot (a sávszűrő sávszélességének fele), az impulzuskarakterisztika hosszát pedig 480 értékre állítjuk. A 2 MHz-es lépésű szűrők esetében ezeket a paramétereket 1 MHz és 240 értékre választjuk, megfelelően. Az ablak típusa Kaiser. Az impulzuskarakterisztikákat a filterDesigner programban kiszámítjuk, és C-fejlécként exportáljuk:

Képernyőképek a filterDesigner-ből

A problémát közvetlen módon is meg lehet oldani: létrehozni a szűrők számának megfelelő DDC tömböt (Digitális Lejáró Konverter). Ez a megközelítés jó FPGA esetén, ahol a számítók bitmélységének csökkentésével lehet megtakarítást elérni. Ezen kívül az FPGA a legenergiatakarékosabb megvalósítási mód. De ennek a módszernek a munkaterhei a legmagasabbak.

A napjainkban népszerű GPU-kon végzett szűrőfésűkkel lehetőség nyílik egy összetettebb algoritmus megvalósítására: polyfázisos szűrőfésű FFT alapján, amely CUDA-ból elérhető a könyvtárból. A külföldi irodalomban az algoritmust Polyphase-nek vagy WOLA-nak (Weight, Overlap and Add) hívják. A rajzolás iránti lustaságom nem ad lehetőséget arra, hogy saját magyarázatot adjak. Az interneten sok anyag található a témában, különösen figyelemre méltó grafikon található itt a 11. oldalon (nagyon köszönöm a tisztelt szerzőknek), itt van:

Megpróbálom saját szavaimmal magyarázni a feldolgozási sémát. Kérem, a gyenge idegzetűek ne olvassák.

Megpróbálom saját szavaimmal magyarázni a feldolgozási sémát a módszertani lehetőségeim keretein belül. Az FFT a bemeneti jel konvolúciója az összes komplex ortogonális harmonikus spektrummal, amelyek az impulzusválasz intervallumában helyezkednek el. Az impulzusválasz, amire a jelet az FFT előtt rá kell állítani, e harmonikus spektrum által modulálva van. Más szavakkal, a szűrők impulzusválaszának burkolója kiszorítható. Ezt követően a harmonikus spektrum egy bizonyos mértékben megritkítható, mivel a szűrő sávszélessége a téglalapablakos szűrőhöz képest kiterjed. A képen a négyes megritkítás látható. Más szóval, a Kaiser ablak tágítása (a maradási sávban egyidejűleg a csillapítás növelésével) után már nem minden szűrő szükséges, csak a negyedik részük. A többiek feleslegesek, frekvenciakarakterisztikáik átfedik egymást. A négy egymást követő FFT pont közül csak a nulladikat választjuk, amelynek számítása négy
bemeneti pont összege, amelyek időtartama a kiindulási FFT negyedével egyenlő.

A hardver választása az, ami kéznél van. Ez a beviteli kártya az "Eszközrendszerek" FMC126P. Róla már írtam egy korábbi bejegyzésben. a Spotify döntéshozatali módszertanáról. Egyre több felhasználó preferálja a zenehallgatást mobilon a streaming szolgáltatásban, de a cégnél túlnyomórészt webprogramozók dolgoznak. A Spotify elemezte ezt a helyzetet, és úgy döntött, hogy több mobilfejlesztőt vesz fel, valamint átnevezi a webes szakembereket új irányba:. Az FMC csatlakozóba egy olyan cég almodulja van beillesztve, amelynek AD9371 transzceiverje van 100 MHz sávszélességgel. A transzceiverből származó összes adat folyamatosan továbbítható a számítógéphez feldolgozásra.

Válasszuk a GTX 1050 GPU-val rendelkező videokártyát. (Hazudtam, ő választott minket: ez volt minden, ami kéznél volt, ki lett téve a számítási segédprogramból az antennák számításához, de annál szebb volt látni egy működő rácsot). Térjünk át a szoftveres részre.

Sajnálattal közöljük, hogy a licencek miatt nem tudjuk közzétenni a teljes kódot. Csak a GPU magokat tudjuk bemutatni. Egyébként a többi kód nem is annyira érdekes.

Íme a mag, amely a jelet ablakra szorozza és összeadja, valamint egy wrapper a hívásához:

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

A jelet feldolgozó kód pontosan követi a fenti ábrán látható algoritmus sémát, ezért nincs értelme itt megjeleníteni.

A rácsokat a csatornák kimeneti spektrumán valós időben ellenőrizték. Az AD9371 bemeneti jeleként egy 2450 MHz-es jelgenerátor jele érkezett, a szűrők szelektivitása a számított értéknek megfelelő volt.