CUDA এর সাথে SDR-এ একবারে সব ব্লুটুথ পাচ্ছেন? সহজে

সম্প্রতি, "দোকান" এর সহকর্মীরা একে অপরের থেকে স্বাধীনভাবে আমাকে জিজ্ঞাসা করতে শুরু করেছেন: কীভাবে একটি এসডিআর রিসিভার থেকে একই সাথে সমস্ত ব্লুটুথ চ্যানেল পাওয়া যায়? ব্যান্ডউইথ এটির অনুমতি দেয়, 80 মেগাহার্টজ বা তার বেশি আউটপুট ব্যান্ডউইথ সহ এসডিআর রয়েছে। আপনি, অবশ্যই, একটি FPGA তে এটি করতে পারেন, তবে এই ধরনের বিকাশের সময় বেশ দীর্ঘ হবে। আমি দীর্ঘদিন ধরে জানি যে একটি জিপিইউতে এটি করা বেশ সহজ, তবে এটি করার মতো!

ব্লুটুথ স্ট্যান্ডার্ড একটি ভৌত ​​স্তরকে দুটি সংস্করণে সংজ্ঞায়িত করে: ক্লাসিক এবং নিম্ন শক্তি। স্পেসিফিকেশন উপলব্ধ এখানে. নথিটি ভয়ানক বড়, এটি সম্পূর্ণভাবে পড়া মস্তিষ্কের জন্য বিপজ্জনক। সৌভাগ্যবশত, বৃহৎ পরিমাপ প্রযুক্তি কোম্পানিগুলির কাছে একটি বিষয়ে ভিজ্যুয়াল নথি তৈরি করার উপায় রয়েছে। Tektronix и জাতীয় যন্ত্রপাতি, উদাহরণ স্বরূপ. উপাদান উপস্থাপনের মানের দিক থেকে আমার তাদের সাথে প্রতিদ্বন্দ্বিতা করার কোন সুযোগ নেই। আপনি আগ্রহী হলে, লিঙ্ক অনুসরণ করুন.

একটি মাল্টি-চ্যানেল ফিল্টার তৈরি করতে ফিজিক্যাল লেয়ার সম্পর্কে আমার যা জানা দরকার তা হল ফ্রিকোয়েন্সি গ্রিড স্পেসিং এবং মড্যুলেশন রেট। সেগুলি নিম্নলিখিত নথিগুলির মধ্যে একটিতে সারণী করা হয়েছে:

এইভাবে, আমাদের 80 MHz ব্যান্ডটিকে 79 MHz এর টিউনিং ধাপের সাথে 1 ফিল্টারে এবং একই সময়ে, 40 MHz এর টিউনিং ধাপ সহ 2 ফিল্টারে কাটাতে হবে। ফিল্টার আউটপুট থেকে স্যাম্পলিং ফ্রিকোয়েন্সি যথাক্রমে 1 MHz এবং 2 MHz হওয়া উচিত।

এইভাবে, আমাদের দুটি ফিল্টার ব্যাঙ্ক প্রয়োজন।

প্রথমে, আসুন ব্লুটুথ ক্লাসিক এবং ব্লুটুথ লো এনার্জি সিগন্যাল ব্যান্ডগুলির উপর ভিত্তি করে এই ফিল্টারগুলির পরামিতিগুলি নির্বাচন করি৷ ফিল্টারের কম্পিউটিং ডিভাইসে লোড গণনা করার জন্য আমাদের তাদের আবেগের প্রতিক্রিয়া প্রয়োজন। এখানে এটি এখনই উল্লেখ করা মূল্যবান যে আমরা "দ্রুত" ফিল্টারিং অ্যালগরিদমের প্রয়োজনীয়তার উপর ভিত্তি করে আবেগ প্রতিক্রিয়ার দৈর্ঘ্য বেছে নিয়েছি। এর থেকে সারমর্ম পরিবর্তন হয় না। এবং বুদ্ধিমান কম্পিউটিং সরঞ্জামগুলিতে ফিল্টার প্রয়োগ করার জন্য আবেগ প্রতিক্রিয়া সহগ সংখ্যাটি খুব বেশি হওয়া উচিত নয়।

1 MHz এর একটি ধাপ সহ ফিল্টারগুলির জন্য, আমরা 500 kHz এর একটি লো-পাস ফিল্টার ব্যান্ডউইথ (একটি ব্যান্ড-পাস ফিল্টারের অর্ধেক ব্যান্ডউইথ) নির্বাচন করি এবং ইমপালস প্রতিক্রিয়া দৈর্ঘ্যকে 480 ট্যাপে সামঞ্জস্য করি। 2 MHz এর ধাপ সহ ফিল্টারগুলির জন্য, আমরা এই পরামিতিগুলিকে যথাক্রমে 1 MHz এবং 240 ট্যাপ হিসাবে নির্বাচন করব। আমরা কায়সার উইন্ডো টাইপ নির্বাচন করি। চলুন ফিল্টারডিজাইনারে আবেগের প্রতিক্রিয়াগুলি গণনা করি এবং সেগুলি সি-হেডার বিন্যাসে আপলোড করি:

ফিল্টার ডিজাইনার থেকে স্ক্রিনশট

আপনি একটি ব্রুট ফোর্স উপায়ে সমস্যাটি সমাধান করতে পারেন: ফিল্টারের সংখ্যার সাথে সম্পর্কিত একটি ডিডিসি অ্যারে তৈরি করুন (ডিজিটাল ডাউন কনভার্টার) এই পদ্ধতিটি FPGA-এর জন্য ভাল, যেখানে প্রথম পর্যায়ের কম্পিউটারগুলির বিট ক্ষমতা হ্রাস করে সঞ্চয় করা সম্ভব। এছাড়াও, FPGA হল সবচেয়ে শক্তি-দক্ষ বাস্তবায়ন পদ্ধতি। কিন্তু এই পদ্ধতিতে শ্রম খরচ সবচেয়ে বেশি।

আজকের জনপ্রিয় GPU-তে ফিল্টার ব্যাঙ্ক চালানোর সময়, আরও পরিশীলিত অ্যালগরিদম বাস্তবায়ন করা সম্ভব হয়: FFT-এর উপর ভিত্তি করে একটি পলিফেজ ফিল্টার ব্যাঙ্ক, যা CUDA-তে লাইব্রেরি থেকে পাওয়া যায়। বিদেশী সাহিত্যে, অ্যালগরিদমকে পলিফেজ বা WOLA (ওজন, ওভারল্যাপ এবং যোগ) FFT ফিল্টারব্যাঙ্ক বলা হয়। অঙ্কনে অলসতা আমাকে নিজের থেকে একটি চাক্ষুষ ব্যাখ্যা সম্পূর্ণ করার সুযোগ দেয় না। ইন্টারনেটে বিষয়টিতে প্রচুর উপাদান রয়েছে, একটি বিশেষভাবে পরিষ্কার গ্রাফ তৈরি করা হয়েছে এখানে পৃষ্ঠা 11 এ (সম্মানিত লেখকদের অনেক ধন্যবাদ), এটি এখানে:

আমি আমার নিজের ভাষায় প্রক্রিয়াকরণ স্কিম ব্যাখ্যা করার চেষ্টা করব। অনুগ্রহ করে মন খারাপের জন্য পড়বেন না।

আমি আমার পদ্ধতিগত ক্ষমতার কাঠামোর মধ্যে আমার নিজের ভাষায় প্রক্রিয়াকরণ স্কিমটি ব্যাখ্যা করার চেষ্টা করব। FFT হল ইনপুট সিগন্যালের আবর্তন যা জটিল অর্থোগোনাল হারমোনিক্সের সমগ্র বর্ণালীর সাথে যা আবেগ প্রতিক্রিয়া ব্যবধানের মধ্যে ফিট করে। ফিল্টারের আবেগ প্রতিক্রিয়া, যার দ্বারা সংকেত FFT ইনপুটের আগে গুণিত হয়, এই সুরেলা বর্ণালী দ্বারা পরিমিত হয়। অন্য কথায়, ফলে আসা চিরুনি ফিল্টারগুলির ইমপালস প্রতিক্রিয়া খামটি বন্ধনীযুক্ত। আরও, আয়তক্ষেত্রাকার উইন্ডোতে ফিল্টারের সাপেক্ষে ফিল্টার ব্যান্ডউইথের প্রসারণের কারণে সুরেলা বর্ণালীটি নির্দিষ্ট সংখ্যক বার পাতলা করা হয়। ছবিতে আমরা চারটি পাতলা দেখতে পাচ্ছি। অন্য কথায়, কায়সার উইন্ডোর সাথে ব্যান্ডটি প্রসারিত করার পরে (স্টপব্যান্ডে টেনশনে একযোগে বৃদ্ধি সহ), আমাদের আর সমস্ত ফিল্টারের প্রয়োজন নেই, তবে তাদের মাত্র এক চতুর্থাংশ। বাকিগুলো অপ্রয়োজনীয়, তাদের ফ্রিকোয়েন্সি বৈশিষ্ট্য ওভারল্যাপ। পরপর চারটি FFT পয়েন্টের মধ্যে, আমরা শুধুমাত্র শূন্য একটি নির্বাচন করি, যার গণনা হল চারটির সমষ্টি
মূল FFT সময়কালের এক চতুর্থাংশের সমান সময়ের পরে নেওয়া ইনপুট পয়েন্ট।

আমরা হাতের কাছে থাকা লোহাটি বেছে নেব। এটি Instrumental Systems FMC126P থেকে ইনপুট বোর্ড। আমি ইতিমধ্যে এটি সম্পর্কে একটি পূর্বে লিখেছি প্রবন্ধ. 9371 MHz ব্যান্ড সহ AD100 ট্রান্সসিভার সহ একই কোম্পানির একটি সাবমডিউল বোর্ডের FMC সংযোগকারীতে ঢোকানো হয়। ট্রান্সসিভার থেকে সম্পূর্ণ স্ট্রিম প্রক্রিয়াকরণের জন্য একটি কম্পিউটারে ক্রমাগত প্রেরণ করা যেতে পারে।

আসুন জিপিইউ জিটিএক্স 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 MHz এর একটি সিগন্যাল জেনারেটর সিগন্যাল সরবরাহ করা হয়েছিল; ফিল্টারগুলির নির্বাচনীতা গণনাকৃত একের সাথে সঙ্গতিপূর্ণ।

পরিকল্পনাগুলির মধ্যে রয়েছে: সফ্টওয়্যারটিকে XRTX বোর্ডে অভিযোজিত করা এবং একটি প্যাকেজ অনুসন্ধান বাস্তবায়ন করা, যদি কারও এটির প্রয়োজন হয় বা অবসর সময় থাকে।

সফটওয়্যারের সব কাজ শেষ গৌদিমা, তার মহিমা!

উত্স: www.habr.com