Pranimi i gjithë Bluetooth-it njëherësh në SDR me CUDA? Lehtë

Në kohët e fundit, kolegët nga "arsimi" kanë filluar të më pyesin: si mund të marim të gjitha kanalet Bluetooth nga një SDR-përceptues në të njëjtën kohë? Banda është e mjaftueshme, kemi një SDR me një bandë dalëse prej 80 MHz ose më shumë. Në fakt, mund ta bëjmë këtë me FPGA, por koha e zhvillimit do të jetë mjaft e gjatë. Më është njohur prej kohësh se realizimi në GPU është shumë i lehtë, por që kështu!

Standarti Bluetooth përcakton nivelin fizik në dy versione: Klasik dhe Energjia e Ulët. Specifikimi është këtu. Dokumenti është mjaft i gjatë, leximi i tij tërësisht është i rrezikshëm për trurin. Për fat të mirë, kompanitë e mëdha që prodhojnë pajisje matëse kanë burimet për të krijuar dokumente të qarta mbi temën. Tektronix dhe National Instruments, për shembull. Unë nuk kam asnjë shans për të konkurruar me ta në cilësinë e paraqitjes së materialit. Ata që janë të interesuar ju lutem hulumtoni përmes lidhjeve.

Gjithçka që më nevojitet për të ditur rreth nivelit fizik për të krijuar një filtrin multikanal është hapësira e rrjetës së frekuencave dhe shpejtësia e modulatorit. Ato janë paraqitur në një tabelë në një nga dokumentet e përmendura:

Në këtë mënyrë, na nevojitet të ndajmë bandën 80 MHz në 79 filtrore me hapësirë rregullimi prej 1 MHz dhe, njëkohësisht, në 40 filtrore me hapësirë rregullimi prej 2 MHz. Frekuencat e ndjekjes së mostrave nga daljet e filtroreve duhet të jenë 1 MHz dhe 2 MHz, përkatësisht.

Pra, na nevojiten dy grebena filtroresh.

Fillimisht do të zgjedhim parametrat e këtyre filtroreve duke u bazuar në bandat e sinjaleve të Bluetooth Classic dhe Bluetooth Low Energy. Na nevojiten karakteristikat e tyre impulsive për të kalkuluar ngarkesën në pajisjen kompjuterike të filtrit. Duhet të sqaroj se gjatësitë e karakteristikave impulsive i kemi zgjedhur sipas kërkesave të algoritmit "të shpejtë" të filtrimit. Kjo nuk e ndryshon thelbësisht situatën. Numri i koeficientëve të karakteristikës impulsive nuk duhet të jetë shumë i madh për të realizuar filtrin në një pajisje kompjuterike të arsyeshme.

Për filtroret me hapësirë 1 MHz, do të zgjedhim bandën e kalimit të FNÇ (gjysmë e bandës së kalimit të filtrit të rrjetit) 500 kHz, dhe do ta përshtatim gjatësinë e karakteristikës impulsive në 480 ndërrime. Për filtroret me hapësirë 2 MHz, këta parametra do t'i zgjedhim në 1 MHz dhe 240 ndërrime, përkatësisht. Lloji i dritares do të zgjidhet Kaiser. Do t’i llogarisim karakteristikat impulsive në filterDesigner dhe do t'i eksportojmë në formatin C-header:

Kopje të ekranit nga filterDesigner

Mund ta zgjidhni problemin në mënyrë frontale: të ndërtoni një array DDC (Digital Down Converter). Ky qasje është e përshtatshme për FPGA, ku mund të arrihet kursim përmes reduktimit të saktësisë së llogaritësve të shkallës së parë. Gjithashtu, FPGA është mënyra më energjikisht efektive për realizimin e kësaj. Por kostot e punës për këtë qasje janë më të larta.

Duke realizuar greben filtroresh në GPU-të më të njohura sot, krijohet mundësia për realizimin e një algoritmi më të sofistikuar: greben filtroresh polifazore bazuar në FFT, i cili në CUDA është në dispozitën e bibliotekës. Në literaturën e huaj, algoritmi quhet Polyphase or WOLA (Weight, Overlap and Add) FFT Filterbank. Pashmangshmëria për të vizatuar nuk më lejon që të ofroj një shpjegim vizual të qartë. Në internet ka shumë materiale mbi temën, sidomos një grafik i qartë është realizuar këtu në faqen 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 harmonikeve komplekse ortogonale, të cilat janë të vendosura brenda intervalit të karakteristikës impulsive. Karakteristika impulsive e filtrit, në të cilën sinjali shumëzohet para futjes në FFT, modullohet nga ky spekter harmonik. Në fjalë të tjera, mbulimi i karakteristikave impulsive të filtroreve rezultues të grebenit është nxjerrë përtej kornizës. Më pas, spektri i harmonikave filtrohet në një numër të caktuar herë, për shkak të zgjerimit të bandës së kalimit të filtrit në krahasim me filtrin me dritare drejtkëndëshe. Në diagram shohim filtrimin në katër. Në fjalë të tjera, pas zgjerimit të bandës me dritaren Kaiser (me njëkohësisht rritjen e qetësisë në bandën e ndalimit), ne tashmë kemi nevojë jo për të gjithë filtroret, por vetëm për një të katërtën e tyre. Të tjerët janë të tepërt dhe karakteristikat e tyre frekuenciale mbivendosen. Nga katër pikët e FFT që ndodhin njëri pas tjetrit, zgjedhim vetëm pikën zero, llogaritja e së cilës është shuma e katër
pikave hyrëse, të marra në një kohë të barabartë me një të katërtën e kohëzgjatjes së FFT origjinale.

Do të zgjedhim pajisjen që kemi në dorë. Kjo është karta e hyrjes nga kompania "Instrumental Systems" FMC126P. Për të, kam shkruajtur më parë në një nga artikujt e mi. artikullin. Në slotin FMC të bordit është instaluar një nënmodul nga e njëjta kompani me transiverin AD9371 me gjerësi 100 MHz. Të gjithë streami nga transiveri mund të transmetohet vazhdimisht në kompjuter për përpunim.

Do të zgjedhim një kartë video me GPU GTX 1050. (Gënjeva, ajo na zgjodhi ne: kjo ishte gjithçka që kishte për dorë, u nxorr nga kompjuteri për llogaritjen e antenave, por ishte befasuese të shihje një grumbull funksional). Të kalojmë në pjesën programore.

Fatkeqësisht, për shkak të licencave nuk mund të publikojmë kodin e plotë. Mund të tregojmë vetëm bërthamat GPU. Megjithatë, kodi tjetër nuk është aq interesant.

Këtu është bërthama që kryen shumëzimin e sinjalit me dritaren dhe shërbimin 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 që përpunon sinjalin dhe thërret këtë bërthamë përsërit saktësisht skemën e algoritmit të paraqitur në figurën më sipër, prandaj nuk shoh arsye për ta sjellë këtu.

Grumbujt u verifikuan në spektrin dalës të kanaleve në kohë reale. Sinjali i gjeneruesit të sinjaleve 2450 MHz u dha në hyrjen e AD9371, selektiviteti i filtrave përputhej me atë të llogaritur.

Në planet: adaptimi i softverit në bordin XRTX dhe realizimi i kërkimit të pakove, nëse dikujt i nevojitet apo do të ketë kohë të lirë.

E gjithë puna për softverin u bë nga gaudima, lavdi atij!

Burimi: habr.com