🥇CUDA ile SDR üzerinden tüm Bluetooth'u aynı anda almak? Kolay

Son zamanlarda "mağazadaki" meslektaşlarım birbirlerinden bağımsız olarak bana şu soruyu sormaya başladılar: tüm Bluetooth kanallarını tek bir SDR alıcısından aynı anda nasıl alabilirim? Bant genişliği buna izin veriyor, çıkış bant genişliği 80 MHz veya daha fazla olan SDR'ler var. Elbette bunu bir FPGA üzerinde yapabilirsiniz ancak böyle bir geliştirme süresi oldukça uzun olacaktır. Bunu GPU'da yapmanın oldukça basit olduğunu uzun zamandır biliyordum ama bu şekilde yapmak!

Bluetooth standardı iki versiyonda fiziksel bir katman tanımlar: Klasik ve Düşük Enerji. Şartname mevcut burada. Belge son derece büyük, tamamını okumak beyin için tehlikeli. Neyse ki büyük ölçüm teknolojisi şirketleri bir konuyla ilgili görsel belgeler oluşturma olanağına sahiptir. Tektronix и Ulusal Araçlar, Örneğin. Malzemenin sunum kalitesi açısından onlarla rekabet etme şansım kesinlikle yok. Eğer ilgileniyorsanız, lütfen bağlantıları takip edin.

Çok kanallı bir filtre oluşturmak için fiziksel katman hakkında bilmem gereken tek şey, frekans ızgarası aralığı ve modülasyon hızıdır. Bunlar aşağıdaki belgelerden birinde tablo halinde verilmiştir:

Bu nedenle 80 MHz bandını 79 MHz ayar adımıyla 1 filtreye ve aynı zamanda 40 MHz ayar adımıyla 2 filtreye ayırmamız gerekiyor. Filtre çıkışlarından örnekleme frekansları sırasıyla 1 MHz ve 2 MHz olmalıdır.

Bu nedenle iki filtre bankasına ihtiyacımız var.

Öncelikle bu filtrelerin parametrelerini Bluetooth Classic ve Bluetooth Low Energy sinyal bantlarına göre seçelim. Filtrenin bilgi işlem cihazındaki yükü hesaplamak için onların dürtü tepkilerine ihtiyacımız var. Burada, dürtü yanıt uzunluklarını "hızlı" bir filtreleme algoritmasının gereksinimlerine göre seçtiğimizi hemen belirtmekte fayda var. Bundan özü değişmez. Ve filtrenin aklı başında bilgisayar ekipmanına uygulanması için dürtü tepki katsayılarının sayısı çok büyük olmamalıdır.

1 MHz'lik adımlı filtreler için, 500 kHz'lik bir alçak geçişli filtre bant genişliği (bant geçişli filtrenin bant genişliğinin yarısı) seçiyoruz ve dürtü yanıt uzunluğunu 480 kademeye ayarlıyoruz. Adımı 2 MHz olan filtreler için bu parametreleri sırasıyla 1 MHz ve 240 kademe olarak seçeceğiz. Kaiser pencere tipini seçiyoruz. FilterDesigner'da dürtü yanıtlarını hesaplayalım ve bunları C-başlık formatında yükleyelim:

filterDesigner'in ekran görüntüleri

Sorunu kaba kuvvet yoluyla çözebilirsiniz: filtre sayısına karşılık gelen bir DDC dizisi oluşturun (Dijital Aşağı Dönüştürücü). Bu yaklaşım, birinci aşama bilgisayarların bit kapasitesinin azaltılmasıyla tasarrufun mümkün olduğu FPGA'ler için iyidir. Ayrıca FPGA enerji açısından en verimli uygulama yöntemidir. Ancak bu yöntemle işçilik maliyetleri en yüksektir.

Günümüzün popüler GPU'larında bir filtre bankası çalıştırıldığında, daha karmaşık bir algoritmanın uygulanması mümkün hale gelir: CUDA'da kütüphaneden temin edilebilen FFT'ye dayalı çok fazlı bir filtre bankası. Yabancı literatürde algoritmaya Polyphase veya WOLA (Weight, Overlap and Add) FFT Filtre Bankası adı verilmektedir. Çizimdeki tembellik bana görsel bir açıklamayı kendi başıma tamamlama fırsatı vermiyor. İnternette konuyla ilgili pek çok materyal var, özellikle net bir grafik oluşturulmuş burada 11. sayfada (saygın yazarlara çok teşekkürler), işte burada:

İşleme şemasını kendi kelimelerimle açıklamaya çalışacağım. Lütfen kalbi zayıf olanlar okumayın.

Metodolojik yeteneklerim çerçevesinde işlem şemasını kendi cümlelerimle anlatmaya çalışacağım. FFT, giriş sinyalinin darbe tepki aralığına uyan tüm karmaşık ortogonal harmonik spektrumuyla evrişimidir. FFT girişinden önce sinyalin çarpıldığı filtrenin darbe tepkisi bu harmonik spektrum tarafından modüle edilir. Başka bir deyişle, elde edilen tarak filtrelerin darbe yanıt zarfı parantez içine alınmıştır. Ayrıca, dikdörtgen penceredeki filtreye göre filtre bant genişliğinin genişlemesi nedeniyle harmonik spektrum belirli sayıda incelir. Resimde dörde kadar bir incelme görüyoruz. Başka bir deyişle, bandı Kaiser penceresiyle genişlettikten sonra (aynı zamanda durdurma bandındaki zayıflamanın artmasıyla), artık tüm filtrelere değil, yalnızca dörtte birine ihtiyacımız var. Gerisi gereksizdir, frekans özellikleri örtüşmektedir. Arka arkaya dört FFT noktasından yalnızca sıfır olanı seçiyoruz; bunun hesaplaması dördünün toplamıdır.
Orijinal FFT süresinin dörtte birine eşit bir süre sonra alınan giriş noktaları.

Eldeki demiri seçeceğiz. Bu, Instrumental Systems FMC126P'nin giriş kartıdır. Daha önceki bir yazımda bunun hakkında yazmıştım Makale. Aynı şirketten 9371 MHz bantlı AD100 alıcı-vericiye sahip bir alt modül, kartın FMC konektörüne takılır. Alıcı-vericiden gelen akışın tamamı işlenmek üzere sürekli olarak bir bilgisayara iletilebilir.

Hadi GPU GTX 1050'li bir ekran kartı seçelim. (Yalan söyledim, bizi seçen oydu: elimizdeki tek şey buydu, antenleri hesaplamak için bilgisayardan sökülmüş, ama görmek daha da şaşırtıcıydı) çalışan bir tarak). Yazılım kısmına geçelim.

Maalesef lisanslar nedeniyle kodun tamamını yayınlayamıyoruz. Yalnızca GPU çekirdeklerini gösterebiliyoruz. Ancak kodun geri kalanı pek ilgi çekici değil.

İşte sinyal penceresi çarpımını ve toplamasını yapan çekirdek ve onu çağıracak bir sarmalayıcı:

__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<<<dataSize / windowStep - 3, windowSize, 1024, stream>>>(data, window, windowSize, result);
    return cudaGetLastError();
}

Bu çekirdeği çağıran sinyal işleme kodu, yukarıdaki şekilde gösterilen algoritma şemasını tam olarak tekrarlıyor, bu yüzden onu burada sunmak için bir neden göremiyorum.

Taraklar, kanalların çıkış spektrumunda gerçek zamanlı olarak test edildi. AD9371 girişine 2450 MHz'lik bir sinyal üreteci sinyali verildi ve filtrelerin seçiciliği hesaplanana karşılık geldi.