🥇ሁሉንም ብሉቱዝ በአንድ ጊዜ በSDR በCUDA እየተቀበልክ ነው? ቀላል

በቅርብ ጊዜ በ "ሱቅ" ውስጥ ያሉ ባልደረቦች እርስ በእርሳቸው ራሳቸውን ችለው ይጠይቁኝ ጀመር: ሁሉንም የብሉቱዝ ቻናሎች በአንድ ጊዜ ከአንድ SDR መቀበያ እንዴት ማግኘት እንደሚቻል? የመተላለፊያ ይዘት ይፈቅዳል, 80 ሜኸር ወይም ከዚያ በላይ የውጤት መተላለፊያ ይዘት ያላቸው SDRs አሉ. በእርግጥ ይህንን በ FPGA ላይ ማድረግ ይችላሉ ፣ ግን እንደዚህ ዓይነቱ የእድገት ጊዜ በጣም ረጅም ነው። ይህን በጂፒዩ ላይ ማድረግ በጣም ቀላል እንደሆነ ለረጅም ጊዜ አውቄአለሁ፣ ግን እንደዛ ማድረግ!

የብሉቱዝ ስታንዳርድ አካላዊ ንብርብርን በሁለት ስሪቶች ይገልፃል፡ ክላሲክ እና ዝቅተኛ ኢነርጂ። ዝርዝር መግለጫ ይገኛል። እዚህ. ሰነዱ በጣም ትልቅ ነው, ሙሉ በሙሉ ማንበብ ለአእምሮ አደገኛ ነው. እንደ እድል ሆኖ, ትላልቅ የመለኪያ ቴክኖሎጂ ኩባንያዎች በአንድ ርዕስ ላይ ምስላዊ ሰነዶችን ለመፍጠር የሚያስችል ዘዴ አላቸው. Tektronix и ብሔራዊ መሳሪያዎች, ለምሳሌ. ከቁሱ አቀራረብ ጥራት አንጻር ከእነሱ ጋር የመወዳደር እድል የለኝም። ፍላጎት ካሎት እባክዎን አገናኞችን ይከተሉ።

ባለብዙ ቻናል ማጣሪያ ለመፍጠር ስለ አካላዊው ንብርብር ማወቅ ያለብኝ ነገር ቢኖር የድግግሞሽ ፍርግርግ ክፍተት እና የመቀየሪያ ፍጥነት ነው። ከሚከተሉት ሰነዶች በአንዱ ውስጥ ተዘርዝረዋል፡

ስለዚህ የ 80 ሜኸር ባንድን በ 79 ማጣሪያዎች በ 1 ሜኸር የማስተካከል ደረጃ እና በተመሳሳይ ጊዜ በ 40 ማጣሪያዎች በ 2 ሜኸር ማስተካከል ያስፈልገናል. የናሙና ድግግሞሾች ከማጣሪያው ውጤቶች 1 MHz እና 2 MHz መሆን አለባቸው።

ስለዚህ, ሁለት ማጣሪያ ባንኮች ያስፈልጉናል.

በመጀመሪያ፣ በብሉቱዝ ክላሲክ እና ብሉቱዝ ዝቅተኛ ኢነርጂ ሲግናል ባንዶች ላይ በመመስረት የእነዚህን ማጣሪያዎች መለኪያዎች እንምረጥ። በማጣሪያው የኮምፒዩተር መሣሪያ ላይ ያለውን ጭነት ለማስላት የእነርሱ ግፊት ምላሽ እንፈልጋለን። እዚህ በ "ፈጣን" የማጣሪያ ስልተ-ቀመር መስፈርቶች ላይ በመመርኮዝ የግፊት ምላሽ ርዝመቶችን እንደመረጥን ወዲያውኑ መጥቀስ ተገቢ ነው። ዋናው ነገር ከዚህ አይቀየርም። እና ማጣሪያው ጤናማ በሆኑ የኮምፒዩተር መሳሪያዎች ላይ እንዲተገበር የግፊት ምላሽ ብዛት በጣም ትልቅ መሆን የለበትም።

1 ሜኸር ደረጃ ላላቸው ማጣሪያዎች ዝቅተኛ ማለፊያ የማጣሪያ ባንድዊድዝ (የባንድ ማለፊያ ማጣሪያ ግማሹን ባንድዊድዝ) 500 kHz እንመርጣለን እና የግፊት ምላሹን ርዝመት ወደ 480 መታዎች እናስተካክላለን። 2 ሜኸር ደረጃ ላላቸው ማጣሪያዎች እነዚህን መለኪያዎች እንደ 1 MHz እና 240 መታዎች በቅደም ተከተል እንመርጣቸዋለን። የ Kaiser መስኮት አይነትን እንመርጣለን. በማጣሪያ ዲዛይነር ውስጥ የግፊት ምላሾችን እናሰላ እና በC-ርዕስ ቅርጸት እንሰቅላቸው፡-

ቅጽበታዊ ገጽ እይታዎች ከማጣሪያ ዲዛይነር

ችግሩን በጠንካራ መንገድ መፍታት ይችላሉ-ከማጣሪያዎች ብዛት ጋር የሚዛመድ የዲዲሲ ድርድር ይገንቡ (ዲጂታል ዳውን መለወጫ). ይህ አካሄድ ለ FPGAs ጥሩ ነው፣ የመጀመርያ ደረጃ ኮምፒውተሮችን ቢት አቅም በመቀነስ መቆጠብ የሚቻልበት። እንዲሁም FPGA በጣም ኃይል ቆጣቢ የትግበራ ዘዴ ነው። ነገር ግን በዚህ ዘዴ የጉልበት ወጪዎች ከፍተኛው ናቸው.

በዛሬው ታዋቂ ጂፒዩዎች ላይ የማጣሪያ ባንክን ሲያካሂዱ ይበልጥ የተራቀቀ አልጎሪዝም መተግበር ይቻላል፡ በኤፍኤፍቲ ላይ የተመሰረተ የ polyphase ማጣሪያ ባንክ፣ እሱም በCUDA ላይ ከቤተ-መጽሐፍት ይገኛል። በውጪ ስነ-ጽሁፍ ውስጥ ስልተ ቀመር Polyphase ወይም WOLA (ክብደት, መደራረብ እና መጨመር) FFT Filterbank ይባላል. በስዕል ውስጥ ስንፍና በራሴ ምስላዊ ማብራሪያን ለማጠናቀቅ እድል አይሰጠኝም. በበይነመረብ ላይ በርዕሱ ላይ ብዙ ቁሳቁሶች አሉ, በተለይም ግልጽ የሆነ ግራፍ ተዘጋጅቷል እዚህ በገጽ 11 ላይ (ለተከበሩ ደራሲያን ብዙ ምስጋና ይገባቸዋል) እነሆ፡-

የሂደቱን እቅድ በራሴ ቃላት ለማብራራት እሞክራለሁ. እባካችሁ ልባችሁን ለደከመ አታንብቡ።

በስልታዊ ችሎታዬ ማዕቀፍ ውስጥ የሂደቱን መርሃ ግብር በራሴ ቃላት ለማብራራት እሞክራለሁ። ኤፍኤፍቲ የግቤት ሲግናል ከሙሉ ስፔክትረም ጋር የተወሳሰቡ orthogonal harmonics በግፊት ምላሽ ክፍተት ውስጥ የሚገጣጠም ነው። ምልክቱ ከኤፍኤፍቲ ግቤት በፊት የሚባዛበት የማጣሪያው የግፊት ምላሽ በዚህ ሃርሞኒክ ስፔክትራ ተስተካክሏል። በሌላ አነጋገር፣ የተገኙት የኩምቢ ማጣሪያዎች የግፊት ምላሽ ፖስታ በቅንፍ ተቀምጧል። በአራት ማዕዘኑ መስኮቱ ውስጥ ካለው ማጣሪያ አንጻር ካለው የማጣሪያ ባንድዊድዝ መስፋፋት የተነሳ ሃርሞኒክ ስፔክትረም በተወሰነ ቁጥር እየሳሳ ነው። በሥዕሉ ላይ በአራት ቀጭን መቀነስ እናያለን. በሌላ አገላለጽ ባንዱን በካይዘር መስኮት ካሰፋን በኋላ (በማቆሚያው ውስጥ በአንድ ጊዜ የመቀነስ ጭማሪ) ሁሉም ማጣሪያዎች አያስፈልጉንም ፣ ግን አንድ አራተኛ ብቻ። የተቀሩት ብዙ ናቸው, የድግግሞሽ ባህሪያቸው ይደራረባል. በተከታታይ ካሉት አራት የኤፍኤፍቲ ነጥቦች ውስጥ ዜሮን አንድ ብቻ እንመርጣለን ፣ ስሌቱም የአራት ድምር ነው።
የግቤት ነጥቦች ከመጀመሪያው FFT ቆይታ ሩብ ጋር እኩል ከሆነ ጊዜ በኋላ የተወሰዱ።

በእጃችን ያለውን ብረት እንመርጣለን. ይህ ከ Instrumental Systems FMC126P የግብአት ሰሌዳ ነው። ቀደም ሲል ስለ እሱ ቀደም ሲል ጽፌ ነበር። ጽሑፍ. ከተመሳሳይ ኩባንያ የመጣ ንዑስ ሞጁል ከ AD9371 ትራንስሴቨር ከ100 ሜኸ ባንድ ጋር በቦርዱ የኤፍኤምሲ ማገናኛ ውስጥ ገብቷል። ከትራንስሲቨር የሚመጣው አጠቃላይ ዥረት ያለማቋረጥ ወደ ኮምፒውተር ለማቀናበር ሊተላለፍ ይችላል።

ከጂፒዩ GTX 1050 ጋር የቪዲዮ ካርድ እንምረጥ (ዋሸሁት እሱ የመረጠን ነው፡ በእጃችን ያለው ሁሉ ነበር አንቴናዎችን ለማስላት ከኮምፒዩተር የተቀዳደደ ነበር ነገር ግን ማየት የበለጠ አስገራሚ ነበር። የሚሰራ ማበጠሪያ). ወደ ሶፍትዌሩ ክፍል እንሂድ።

እንደ አለመታደል ሆኖ፣ በፍቃዶች ምክንያት፣ ሙሉውን ኮድ ማተም አንችልም። የጂፒዩ ኮሮችን ብቻ ነው ማሳየት የምንችለው። ይሁን እንጂ የተቀረው ኮድ በተለይ አስደሳች አይደለም.

የምልክት-መስኮት ማባዛትና መደመርን እና እሱን ለመጥራት የሚያገለግል ከርነል ይኸውና፡

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

ይህንን ከርነል የሚጠራው የሲግናል ማቀናበሪያ ኮድ ከላይ በስዕሉ ላይ የሚታየውን የአልጎሪዝም ንድፍ በትክክል ይደግማል፣ ስለዚህ እዚህ ለማቅረብ ምንም ምክንያት አይታየኝም።

ማበጠሪያዎቹ በእውነተኛ ጊዜ በሰርጦቹ የውጤት ስፔክትረም ላይ ተፈትነዋል። የ9371 ሜኸር ሲግናል ጀነሬተር ሲግናል ለ AD2450 ግብዓት ቀርቧል፣ እና የማጣሪያዎቹ ምርጫ ከተሰላው ጋር ይዛመዳል።