🥇Ստա՞ք բոլոր Bluetooth-ը միանգամից SDR-ով CUDA-ով: Հեշտ

Վերջերս «խանութի» գործընկերները սկսեցին ինձ միմյանցից անկախ հարցնել. ինչպե՞ս ստանալ բոլոր Bluetooth ալիքները միաժամանակ մեկ SDR ընդունիչից: Լայնությունը թույլ է տալիս, կան SDR-ներ 80 ՄՀց կամ ավելի ելքային թողունակությամբ: Դուք, իհարկե, կարող եք դա անել FPGA-ով, բայց զարգացման նման ժամանակը բավականին երկար կլինի: Ես վաղուց գիտեի, որ դա անելը GPU-ով բավականին պարզ է, բայց դա անել այդպես:

Bluetooth ստանդարտը սահմանում է ֆիզիկական շերտ երկու տարբերակով՝ դասական և ցածր էներգիա: Հասանելի է հստակեցում այստեղ. Փաստաթուղթը սարսափելի մեծ է, այն ամբողջությամբ կարդալը վտանգավոր է ուղեղի համար։ Բարեբախտաբար, չափման տեխնոլոգիական խոշոր ընկերությունները միջոցներ ունեն թեմայի վերաբերյալ տեսողական փաստաթղթեր ստեղծելու համար: Tektronix и Ազգային գործիքներ, Օրինակ. Նյութի մատուցման որակով ես նրանց հետ մրցելու բացարձակ շանս չունեմ։ Եթե դուք հետաքրքրված եք, խնդրում ենք հետևել հղումներին:

Այն ամենը, ինչ ես պետք է իմանամ ֆիզիկական շերտի մասին, բազմալիք ֆիլտր ստեղծելու համար, հաճախականությունների ցանցի տարածությունն է և մոդուլյացիայի արագությունը: Դրանք աղյուսակավորված են հետևյալ փաստաթղթերից մեկում.

Այսպիսով, մենք պետք է կտրենք 80 ՄՀց տիրույթը 79 զտիչների՝ 1 ՄՀց թյունինգի քայլով և, միևնույն ժամանակ, 40 ֆիլտրերի՝ 2 ՄՀց թյունինգի քայլով: Ֆիլտրի ելքերից նմուշառման հաճախականությունները պետք է լինեն համապատասխանաբար 1 ՄՀց և 2 ՄՀց:

Այսպիսով, մեզ անհրաժեշտ է երկու զտիչ բանկ:

Նախ, եկեք ընտրենք այս ֆիլտրերի պարամետրերը՝ հիմնվելով Bluetooth Classic և Bluetooth Low Energy ազդանշանային տիրույթների վրա: Մեզ անհրաժեշտ են նրանց իմպուլսային պատասխանները՝ ֆիլտրի հաշվողական սարքի բեռը հաշվարկելու համար: Այստեղ հարկ է անմիջապես նշել, որ մենք ընտրել ենք իմպուլսային արձագանքման երկարությունները՝ ելնելով «արագ» զտման ալգորիթմի պահանջներից: Սրանից էությունը չի փոխվում. Եվ իմպուլսային արձագանքման գործակիցների թիվը չպետք է չափազանց մեծ լինի, որպեսզի ֆիլտրը կիրառվի ողջամիտ հաշվողական սարքավորումների վրա:

1 ՄՀց քայլ ունեցող ֆիլտրերի համար մենք ընտրում ենք ցածր անցումային ֆիլտրի թողունակությունը (ժապավենային անցումային ֆիլտրի թողունակության կեսը) 500 կՀց, և իմպուլսային արձագանքման երկարությունը կարգավորում ենք մինչև 480 հպում: 2 ՄՀց քայլ ունեցող ֆիլտրերի համար մենք կընտրենք այս պարամետրերը, համապատասխանաբար, որպես 1 ՄՀց և 240 հպում: Մենք ընտրում ենք Kaiser պատուհանի տեսակը: Եկեք հաշվարկենք իմպուլսային պատասխանները filterDesigner-ում և վերբեռնենք դրանք C-header ձևաչափով.

Սքրինշոթներ filterDesigner-ից

Խնդիրը կարող եք լուծել բիրտ ուժի միջոցով՝ կառուցեք ֆիլտրերի քանակին համապատասխան DDC զանգված (Թվային ներքև փոխարկիչ) Այս մոտեցումը լավ է FPGA-ների համար, որտեղ հնարավոր են խնայողություններ՝ նվազեցնելով առաջին փուլի համակարգիչների բիթային հզորությունը: Բացի այդ, FPGA-ն էներգաարդյունավետ իրականացման մեթոդն է: Բայց այս մեթոդով աշխատուժի ծախսերը ամենաբարձրն են:

Այսօրվա հայտնի GPU-ների վրա ֆիլտրի բանկ գործարկելիս հնարավոր է դառնում իրականացնել ավելի բարդ ալգորիթմ՝ FFT-ի վրա հիմնված բազմաֆազ ֆիլտրի բանկ, որը CUDA-ի վրա հասանելի է գրադարանից: Արտասահմանյան գրականության մեջ ալգորիթմը կոչվում է Polyphase կամ WOLA (Weight, Overlap and Add) FFT Filterbank: Նկարչության մեջ ծուլությունն ինձ հնարավորություն չի տալիս ինքնուրույն լրացնել տեսողական բացատրությունը։ Համացանցում թեմայի վերաբերյալ շատ նյութեր կան, կազմված է հատկապես հստակ գրաֆիկ այստեղ էջ 11 (շատ շնորհակալություն հարգարժան հեղինակներին), ահա.

Ես կփորձեմ բացատրել մշակման սխեման իմ խոսքերով: Խնդրում եմ, մի կարդացեք թույլ սրտի համար:

Իմ մեթոդական հնարավորությունների շրջանակներում կփորձեմ բացատրել մշակման սխեման իմ խոսքերով։ FFT-ն մուտքային ազդանշանի կոնվուլյացիա է բարդ ուղղանկյուն ներդաշնակությունների ողջ սպեկտրով, որոնք տեղավորվում են իմպուլսային արձագանքման միջակայքում: Ֆիլտրի իմպուլսային արձագանքը, որով ազդանշանը բազմապատկվում է FFT մուտքից առաջ, մոդուլացվում է այս ներդաշնակ սպեկտրով: Այլ կերպ ասած, արդյունքում ստացված սանր ֆիլտրերի իմպուլսային արձագանքման ծրարը փակցված է: Այնուհետև, ներդաշնակ սպեկտրը նոսրանում է որոշակի քանակությամբ անգամներ՝ ուղղանկյուն պատուհանում ֆիլտրի համեմատ ֆիլտրի թողունակության ընդլայնման պատճառով: Նկարում տեսնում ենք չորսով նոսրացում։ Այլ կերպ ասած, Kaiser պատուհանով գոտին ընդլայնելուց հետո (կանգառի թուլացման միաժամանակյա աճով) մեզ այլևս պետք չեն բոլոր ֆիլտրերը, այլ դրանց միայն մեկ քառորդը։ Մնացածը ավելորդ են, դրանց հաճախականության բնութագրերը համընկնում են: Անընդմեջ չորս FFT կետերից մենք ընտրում ենք միայն զրոյական մեկը, որի հաշվարկը չորսի գումարն է.
մուտքագրման կետերը վերցված ժամանակից հետո, որը հավասար է սկզբնական FFT-ի տևողության քառորդին:

Մենք կընտրենք ձեռքի տակ եղած արդուկը։ Սա Instrumental Systems FMC126P-ի մուտքային տախտակն է: Ես արդեն գրել եմ այդ մասին նախորդ մեկում Հոդված. Տախտակի FMC միակցիչում տեղադրվում է նույն ընկերության ենթամոդուլը AD9371 հաղորդիչով 100 ՄՀց տիրույթով: Փոխանցիչից ամբողջ հոսքը կարող է շարունակաբար փոխանցվել համակարգչին՝ մշակման համար:

Եկեք ընտրենք վիդեո քարտ GPU GTX 1050-ով: (Ես ստեցի, դա մեզ ընտրեց. այն ամենն էր, ինչ ձեռքի տակ էր, համակարգչից պոկված էր ալեհավաքները հաշվարկելու համար, բայց ավելի զարմանալի էր տեսնելը. աշխատանքային սանր): Անցնենք ծրագրային մասին։

Ցավոք, լիցենզիաների պատճառով մենք չենք կարող հրապարակել ամբողջական կոդը: Մենք կարող ենք ցույց տալ միայն GPU միջուկները: Այնուամենայնիվ, ծածկագրի մնացած մասը առանձնապես հետաքրքիր չէ:

Ահա այն միջուկը, որը կատարում է ազդանշանի պատուհանի բազմապատկում և գումարում, և այն անվանելու համար փաթաթան.

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

Ազդանշանի մշակման կոդը, որը կանչում է այս միջուկը, ճշգրտորեն կրկնում է վերը նկարում ներկայացված ալգորիթմի դիագրամը, ուստի ես պատճառ չեմ տեսնում այն այստեղ ներկայացնելու համար:

Սանրերը փորձարկվել են իրական ժամանակում ալիքների ելքային սպեկտրի վրա: AD9371-ի մուտքին տրվել է 2450 ՄՀց ազդանշանի գեներատորի ազդանշան, ֆիլտրերի ընտրողականությունը համապատասխանում է հաշվարկվածին: