Yaqinda "do'kon"dagi hamkasblar mendan bir-biridan mustaqil ravishda so'rashni boshladilar: qanday qilib bir vaqtning o'zida barcha Bluetooth kanallarini bitta SDR qabul qilgichdan olish mumkin? O'tkazish qobiliyati bunga imkon beradi, chiqish o'tkazish qobiliyati 80 MGts yoki undan ko'p bo'lgan SDRlar mavjud. Siz, albatta, buni FPGA-da qilishingiz mumkin, ammo bunday ishlab chiqish vaqti ancha uzoq bo'ladi. Men buni GPUda qilish juda oddiy ekanligini uzoq vaqtdan beri bilaman, lekin buni shunday qilish!
Bluetooth standarti jismoniy qatlamni ikkita versiyada belgilaydi: Klassik va Past energiya. Spetsifikatsiya mavjud . Hujjat juda katta, uni to'liq o'qish miya uchun xavflidir. Yaxshiyamki, yirik o'lchov texnologiyalari kompaniyalari mavzu bo'yicha vizual hujjatlarni yaratish vositalariga ega. и , Masalan. Materialni taqdim etish sifati bo'yicha ular bilan raqobatlashishga mutlaqo imkoniyatim yo'q. Agar qiziqsangiz, havolalarni kuzatib boring.
Ko'p kanalli filtr yaratish uchun jismoniy qatlam haqida bilishim kerak bo'lgan narsa - bu chastotalar tarmog'i oralig'i va modulyatsiya tezligi. Ular quyidagi hujjatlardan birida jadvalga kiritilgan:
Shunday qilib, biz 80 MGts diapazonni sozlash bosqichi 79 MGts bo'lgan 1 filtrga va shu bilan birga 40 MGts sozlash bosqichiga ega 2 filtrga kesishimiz kerak. Filtr chiqishlaridan namuna olish chastotalari mos ravishda 1 MGts va 2 MGts bo'lishi kerak.
Shunday qilib, bizga ikkita filtr banki kerak.
Birinchidan, Bluetooth Classic va Bluetooth Low Energy signal diapazonlari asosida ushbu filtrlarning parametrlarini tanlaylik. Filtrning hisoblash qurilmasidagi yukni hisoblash uchun bizga ularning impuls javoblari kerak. Bu erda darhol ta'kidlash kerakki, biz "tezkor" filtrlash algoritmi talablaridan kelib chiqqan holda impulslarga javob uzunligini tanladik. Bundan mohiyat o'zgarmaydi. Va impulsli javob koeffitsientlari soni filtrni aqlli hisoblash uskunalarida amalga oshirish uchun juda katta bo'lmasligi kerak.
1 MGts pog'onali filtrlar uchun biz 500 kHz past chastotali filtr tarmoqli kengligini (tarmoqli o'tkazuvchan filtrning yarmi) tanlaymiz va impulsga javob uzunligini 480 tegishga moslashtiramiz. 2 MGts qadamli filtrlar uchun biz ushbu parametrlarni mos ravishda 1 MGts va 240 kran sifatida tanlaymiz. Biz Kaiser oyna turini tanlaymiz. Keling, filterDesigner-da impuls javoblarini hisoblab chiqamiz va ularni C-sarlavha formatida yuklaymiz:
filterDesigner-dan skrinshotlar
Muammoni qo'pol kuch bilan hal qilishingiz mumkin: filtrlar soniga mos keladigan DDC massivini yarating (). Ushbu yondashuv FPGA uchun yaxshi, bu erda birinchi bosqich kompyuterlarining bit sig'imini kamaytirish orqali tejash mumkin. Bundan tashqari, FPGA eng energiya tejamkor amalga oshirish usuli hisoblanadi. Ammo bu usul bilan mehnat xarajatlari eng yuqori.
Bugungi kunda mashhur GPU-larda filtrlar bankini ishga tushirishda yanada murakkab algoritmni amalga oshirish mumkin bo'ladi: kutubxonadan CUDA-da mavjud bo'lgan FFT-ga asoslangan ko'p fazali filtrlar banki. Chet el adabiyotida algoritm Polyphase yoki WOLA (Weight, Overlap and Add) FFT Filterbank deb ataladi. Chizishdagi dangasalik menga vizual tushuntirishni mustaqil ravishda bajarish imkoniyatini bermaydi. Internetda mavzu bo'yicha juda ko'p materiallar mavjud, ayniqsa aniq grafik tuziladi 11-betda (hurmatli mualliflarga katta rahmat), mana bu:
Men o'z so'zlarim bilan ishlov berish sxemasini tushuntirishga harakat qilaman. Iltimos, yurak zaiflar uchun o'qimang.
Men o'zimning uslubiy imkoniyatlarim doirasida qayta ishlash sxemasini o'z so'zlarim bilan tushuntirishga harakat qilaman. FFT - bu impuls javob oralig'iga to'g'ri keladigan murakkab ortogonal harmonikalarning butun spektri bilan kirish signalining konvolyutsiyasi. Signal FFT kirishidan oldin ko'paytiriladigan filtrning impulsli javobi ushbu garmonik spektrlar tomonidan modulyatsiya qilinadi. Boshqacha qilib aytganda, hosil bo'lgan taroq filtrlarining impulsli javob konverti qavs ichida joylashgan. Bundan tashqari, to'rtburchaklar oynadagi filtrga nisbatan filtr o'tkazuvchanligi kengayishi tufayli garmonik spektr ma'lum bir necha marta yupqalanadi. Rasmda biz to'rtga yupqalashganini ko'ramiz. Boshqacha qilib aytadigan bo'lsak, Kaiser oynasi bilan bandni kengaytirgandan so'ng (stopbandda bir vaqtning o'zida zaiflashuvning oshishi bilan) biz endi barcha filtrlarga muhtoj emasmiz, lekin ularning to'rtdan bir qismi. Qolganlari ortiqcha, ularning chastotali xarakteristikalari bir-biriga mos keladi. Ketma-ket to'rtta FFT nuqtasidan biz faqat nolni tanlaymiz, uning hisobi to'rtta yig'indisidir.
dastlabki FFT davomiyligining chorak qismiga teng vaqtdan keyin olingan kirish nuqtalari.
Biz qo'limizdagi temirni tanlaymiz. Bu Instrumental Systems FMC126P dan kirish platasi. Men bu haqda oldingi birida yozgan edim . Kengashning FMC ulagichiga 9371 MGts diapazonli AD100 qabul qiluvchi-uzatuvchisi bilan bir xil kompaniyaning submodul o'rnatilgan. Transceiverdan butun oqim qayta ishlash uchun doimiy ravishda kompyuterga uzatilishi mumkin.
Keling, GPU GTX 1050 bilan video kartani tanlaylik. (Men yolg'on gapirdim, bu bizni tanladi: hamma narsa qo'lida edi, antennalarni hisoblash uchun kompyuterdan yirtilgan edi, lekin buni ko'rish yanada hayratlanarli edi. ishlaydigan taroq). Keling, dasturiy ta'minot qismiga o'tamiz.
Afsuski, litsenziyalar tufayli biz toʻliq kodni nashr eta olmaymiz. Biz faqat GPU yadrolarini ko'rsatishimiz mumkin. Biroq, kodning qolgan qismi unchalik qiziq emas.
Mana signal oynasini ko'paytirish va qo'shishni amalga oshiradigan yadro va uni chaqirish uchun o'ram:
__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();
}
Ushbu yadroni chaqiradigan signalni qayta ishlash kodi yuqoridagi rasmda ko'rsatilgan algoritm diagrammasini aynan takrorlaydi, shuning uchun uni bu erda taqdim etish uchun hech qanday sabab ko'rmayapman.
Taroqlar real vaqt rejimida kanallarning chiqish spektrida sinovdan o'tkazildi. AD9371 kirishiga 2450 MGts chastotali signal generatori signali berildi, filtrlarning selektivligi hisoblanganga to'g'ri keldi.
Rejalarga quyidagilar kiradi: dasturiy ta'minotni XRTX platasiga moslashtirish va agar kimdir kerak bo'lsa yoki bo'sh vaqt bo'lsa, paketlarni qidirishni amalga oshirish.
Barcha dasturiy ta'minot ishlari tugallandi , Unga shon-sharaflar!
Manba: www.habr.com
