جمعه 29 دی 1396 | Friday 19 th of January 2018 صفحه اصلی گروه الکترونیکی کامپیوتر
7ـ2ـ1 مجموع برداری: ردوکس (REDUX)

   ما برای به انجام رساندن وظیفه ای مشابه انچه که در فصل قبل انجام دادیم، تلاش میکنیم.به این معنا که ما می خواهیم دو بردار ورودی را گرفته و جمع انها را در یک بردار خروجی سوم ذخیره کنیم. با این حال، در این زمان برای تحقق این امر بجای بلوک ها از رشته ها استفاده میکنیم.

   شاید برایتان این سوال پیش امده باشد که مزیت استفاده از رشته ها به جای بلوکها در چیست؟ خب، در حال حاضر هیچ مزیتی که ارزش بحث را داشته باشد، وجود ندارد.اما رشته های موازی موجود در یک بلوک توانایی انجام کارهایی را دارند که از عهده بلوک های موازی برنمی اید.بنابراین فعلا صبور بوده و از ما راضی باشید تا برویم سراغ یک نسخه رشته موازی مرتبط به مثال بلوک موازی فصل قبل.

مجموع برداریGPUاستفاده کننده از رشته ها

   ما کار را با پرداختن به دو تغییر قابل توجه در هنگام حرکت از بلوکهای موازی به رشته های موازی اغاز میکنیم. فراخوانی هسته از جایی شروع به تغییر میکند که بلوکهایN  یک نمونه رشته راه اندازی میشوند:

<<>>( dev _ a, dev _ b, dev _ c )را به نسخه ای که همه رشته های Nرا در یک بلوک راه اندازی میکند، اضافه کنید:

<<<1,N>>>( dev _ a, dev _ b, dev _ c )add

تنها دیگر تغییرات پیش امده در این روش، ما را قادر به شاخص بندی داده ها میکنند. پیش از این ما توانستیم داده های ورودی و خروجی را با استفاده از شاخص بلوکی در درون هسته فهرست بندی کنیم.


در اینجا وجود این جمله‌ اساسی‌نباید تعجب برانگیز باشد.حال که ما تنها از یک بلوک برخورداریم بایست داده ها را با استفاده از شاخص رشته ای فهرست بندی کنیم.

 

این دو تغییر تنها تغییراتی هستند که به تغییر از پیاده سازی بلوک های موازی به پیاده سازی رشته های موازی نیاز دارند. برای تکمیل کار در اینجا لازم است که لیست کامل منبع همراه با خطوط تغییر بصورت  پررنگ نشان داده شوند:


خیلی ساده است، مگر نه؟ در بخش بعدی، ما با یکی از محدودیت های این رویکرد فقط رشته ای مواجه میشویم و البته بعدا درخواهیم یافت که چرا تقسیم بلوکها به اجزای موازی دیگر بدون زحمت نیست.
مجموعیک بردار بلندتر  GPU

    ما در فصل گذشته متوجه شدیم که سخت افزار تعداد بلوک ها در یک راه اندازی واحد را تا 65,535کاهش میدهد. همچنین، سخت افزار تعداد رشته های موجود در هر بلوک که با ان قادر به راه اندازی هسته میباشد را کم میکند. به طور خاص، این تعداد نمی تواند از مقدار مشخص شده توسط فیلدmaxThreadsPerBlock (حد اکثر رشته های موجود در هر بلوک) مربوط به ساختار خصوصیات دستگاه مشاهده شده در فصل 5 فراتر برود.در حال حاضر برای بسیاری از پردازنده های گرافیکی در دسترس، این محدودیت به میزان 512 رشته در هر بلوک می باشد بنابراین این سوال مطرح میشود که چگونه میتوان با استفاده از یک رویکرد مبتنی بر رشته به اضافه کردن دو بردار با اندازه های بزرگتر از 512 مبادرت ورزید؟ ما برای تحقق این امر ملزم به استفاده ترکیبی از رشته ها و بلوکها میباشیم.

درست مثل قبل در اینجا نیز به دو تغییر نیاز است:  ما باید محاسبه شاخص در هسته و همچنین خود راه اندازی هسته را تغییر دهیم.

حال که ما از تعداد متعددی بلوک و رشته برخورداریم، نمایه سازی مربوطه شبیه به روش استاندارد برای تبدیل فضای شاخص دو بعدی به یک فضای خطی اغاز خواهد شد.

 برای انجام این کار از یک متغیر جدیدا ساخته شده به نامblockDim  استفاده می شود. این متغیر برای همه بلوک ها ثابت است و تعداد رشته های موجود در طول هر بعد از بلوک را ذخیره میکند و از انجا که ما در حال استفاده از یک بلوک تک بعدی هستیم تنها بهblockDim.x  رجوع میکنیم. اگر به یاد داشته باشیدgridDim  یک مقدار مشابه را ذخیره می کرد ولی در عین حال تعداد بلوکهای موجود در امتداد هر یک از ابعاد کل شبکه را نیز  ذخیره می کرد. علاوه بر این، gridDimدو بعدی است در حالی کهblockDim در واقع سه بعدی است. به عبارت دیگر، زمان اجرایCUDA  به شما اجازه راه اندازی یک شبکه دو بعدی از بلوک ها که در ان هر بلوک ارایه ای سه بعدی از رشته هاست را می دهد.بله، ابعاد بسیاری وجود دارند و بعید است که شما مرتبا به پنج درجه ازادی نمایه سازی فراهم شده نیازپیدا کنید. اما در صورت تمایل در دسترس میباشند.

   در واقع نمایه سازی داده ها در یک ارایه خطی با کمک گرفتن از وظیفه قبلی کاملا حسی‏ است.اگر شما با این ایده مخالفید بهتر است بدانید که ممکن است به شما کمک کند تا در مورد مجموعه بلوک های رشته ها همانند ارایه دو بعدی پیکسلها بصورت فضایی فکر کنید.ما در شکل 7ـ1 این ترتیب را به تصویر کشیده ایم.

   اگر رشته ها نشان دهنده ستونها و بلوکها نشان دهنده ردیفها باشند، میتوانیم با در نظر گرفتن محصول شاخص بلوک همراه با تعداد رشته ها در هر بلوک و اضافه کردن شاخص رشته ای بداخل بلوک به یک شاخص منحصر به فرد دست یابیم. این در مورد روش مورد استفاده ما در طولی کردن شاخص تصویر دو بعدی در مثال مجموعه (Julia Set) نیز صدق میکند.

 در اینجا، DIM  همان بعد بلوکی است (که در رشته ها اندازه گیری میشود)،Y  شاخص بلوکی و X شاخص رشته ای داخل بلوک می باشد.از این رو ما بشرح ذیل به شاخص دست میابیم:

   تغییر دیگر راه اندازی خود هستهاست. ما هنوز به رشته های موازیN  جهتراه اندازی نیاز داریم اما خواهان راه اندازی انها در سرتاسر بلوک های چندگانه هستیم و بنابراین محدودیت512  رشته ای تحمیل شده را مد نظر قرار نمیدهیم.یک راه حل، تنظیم دلبخواه اندازه بلوک به تعداد ثابتی از رشته ها میباشد. بگذارید در این مثال از 128 رشته در هر بلوک استفاده کنیم. سپس تنها کاری که میتوانیم انجام دهیم راه اندازیN/128  جهت بدست اوردن تعدادکلی رشته های در حال اجرا است.

 مسئله پیچیده در اینجا این است کهN/128 در واقع یک تقسیم عدد صحیح است و این حاکی از ان است که در صورت 127 بودن تعداد،N/128  برابر با صفر خواهد بود و اگر ما اقدام به راه اندازی کنیم در واقع دست به محاسبه هیچ چیزی نزده ایم. در واقع، ما تعداد بسیار کمی از رشته ها را در زمانی کهN مضرب دقیقی از 128 نیست راه اندازی خواهیم کرد و این خوب نیست. در حقیقت ما قصد گرد کردنافزایشىاین تقسیم را داریم.

شکل 7ـ1 ترتیب دو بعدی مجموعه ای از بلوک ها و رشته ها و رشته های صفر.

   ترفند مرسومی برای تحقق این امر در تقسیم عدد صحیح انهم بدون فراخوانی ceil()وجود دارد. ما عملا (N+127)/128را بجای N/128محاسبهمیکنیم. در هر صورت یا شما متوجه منظور ما را در مورد اینکه کوچکترین مضرب 128 بزرگتر یا مساویN محاسبه میشود خواهید شد و یا همین حالا باید وقتی را برای فکر کردن به این واقعیت جهت متقاعد کردن خود اختصاص دهید.
ما در هر بلوک 128 رشته را انتخاب کرده ایم و به همین دلیل از روش زیر برای راه اندازی هسته استفاده میکنیم:

به دلیل تغییر صورت گرفته در تقسیم است که از راه اندازی رشته های کافی مطمئن میشویم که موضوع به اندازه کافی راه اندازی می کنیم، در واقع هنگامی کهN مضرب دقیق 128 نیست میتوانیم رشته های بسیاری را راه اندازی کنیم. اما یک چاره‏ساده تر هم برای این مشکل وجود دارد که هسته ما پیشتر به ان پرداخته است. ما باید این موضوع را بررسی کنیم که ایا افست یک رشته پیش از استفاده از ان برای دسترسی به ارایه های خروجی و ورودی در واقع بین 0 وN است یا خیر:

بنابراین هنگامی که شاخص ما انتهای ارایه را فراجهش میکند مثل زمانیکه یکنا مضرب 128 فعال میشود، ما به طور خودکار از انجام محاسبه خودداری میکنیم. مهمتر انکه ما از خواندن و نوشتن حافظه در انتهای ارایه مان خودداری میکنیم.

مجموعبردارهای بطور دلخواهانه‏بلندGPU

   زمانیکه ما برای اولین بار به بحث در مورد راه اندازی بلوکهای موازی درGPU پرداختیم به طور کامل امادگی ارائه انرا نداشتیم. علاوه بر محدودیت در تعداد رشته ها، محدودیت های سخت افزاری هم در تعداد بلوک ها وجود دارد (اگر چهخیلی بیشتر از محدودیتهای رشته ای است).همانطور که قبلا ذکر شد هیچیک از ابعاد شبکه بلوک ها نمیتوانند از 65،535 فراتر روند.

   بنابراین این موضوع مشکلی را با اجرای جاری جمع برداری پدید می اورد. علاوه بر این در حال حاضر ما را افزایش می دهد.اگر بمنظور اضافه کردن بردارهایمان به راه اندازی بلوکهایN/128 بپردازیم در زمان فراتررفتن بردارمان از65,535 * 128 = 8,388,480  عنصر خواهیم توانست مشکلات راه اندازی را با موفقیت پشت سر بگذاریم. مقدار این عدد بنظر زیاد می اید ولی با ظرفیت حافظه جاری بین1GB و 4GB، پردازنده های گرافیکی فوق پیشرفته(high-end) می توانند مقیاساطلاعاتیگسترده تریرا نسبت به بردارهای 8  میلیون عنصره در خود جای دهند.  خوشبختانه، راه حل این موضوع بسیار ساده است.ما برای اولین بار تغییری را در هسته ایجاد میکنیم.

 این امر بطرز قابل ملاحظه ای تداعی کننده نسخه اصلی جمع برداری است!در واقع، شما باید انرا با پیاده سازی واحد پردازنده مرکزی (CPU) فصل قبل مقایسه کنید:

   در اینجا ما از حلقهwhile() جهت انجام تکرار از طریق داده ها استفاده می کنیم.به یاد بیاورید که ما به جای افزایش شاخص ارایه تا 1، ادعای وجود یک نسخه چند هسته ای یا دارای چند واحد پردازنده مرکزی را میکنیم که می تواند توسط تعدادی از پردازنده های مورد استفاده ما در اینده توسعه یابد.ما در حال حاضر از اصول مشابهی در نسخهGPU  پیروی می کنیم.

   برای اجرایGPU باید تعدادی از رشته های موازی راه اندازی شده جهت کار بعنوان تعدادی از پردازنده ها مد نظر قرار گیرند. گرچه ممکن استGPU واقعی نسبت به این مورد از واحدهای پردازشی کمتری (و یا بیشتری) برخوردار باشد ولی هر رشته از لحاظ منطقی بصورت موازی اجرا میشود و سپس به سخت افزار اجازه می دهد تا اجرای واقعی را زمانبندی کند.افتراق (Decoupling)موازی سازی از روش واقعی اجرای سخت افزاری یکی از مشکلات اساسی است کهCUDA C از سر راه طراح نرم افزاری برمیدارد. با در نظر گرفتن اینکه سخت افزار فعلیNVIDIA می تواند به هر نقطه ای در محدوده بین 8 و 480 واحد حسابی  در هر تراشه انتقال یابد میتوان به راحتی هر چه بیشتر کار دست یافت!

   حال که ما با ظوابط و اصول پنهان شده در فراسوی این پیاده سازی اشنا شدیم تنها نیاز است که نحوه تعیین ارزش شاخص اولیه برای هر رشته موازی و نحوه افزایش انرا درک کرده و دریابیم.ما می خواهیم که هر یک از رشته های موازی بر مبنای یک شاخص داده متفاوت شروع بکار کند بنابراین تنها کاری که نیاز است گرفتن رشته ها و بلوکه کردن شاخص ها و بصورت طولی دراوردن انها طبق انچیزی است که در بخش "مجموع بردار طولانی تر"GPU دیدیم.هر رشته در شاخصی به شرح زیر شروع بکار میکند:

پس از اتمام کار هر رشته در شاخص جاری به افزایش هر یک از انها با تعداد کل رشته های در حال اجرا در شبکه نیاز است.این صرفا تعداد رشته های هر بلوک میباشد که به تعداد بلوکها در شبکه ضرب شده است و یا blockDim.x ×gridDim.x . از این رو، مرحله افزایش به شرح زیر است:
ما تقریبا کار را انجام داده ایم!تنها قسمت باقی مانده، تعمیر خود راه اندازی است.اگر به یاد داشته باشید ما این مسیر انحرافی را در پیش گرفتیم زیرا راه اندازی
که وظیفهاضافه کردن<<<(N+127)/128,128>>>( dev_a, dev_b, dev_c )را دارد در زمان بزرگتر بودن(N+127)/128  از 65،535 با شکست مواجه خواهد شد. برای اطمینان از عدم راه اندازی تعداد بسیار زیاد بلوکها، ما فقط به ثابت کردن تعداد بلوکها به مقدار اندک اما منطقی نیاز داریم. از انجا که ما عاشق
کپی و پیست کردن (چسباندن) هستیم از 128 بلوک که هر یک حاوی 128 رشته میباشند استفاده خواهیم کرد:

شما باید برای تنظیم این مقادیر احساس راحتی کنید حتی اگر انها مناسب بنظر می رسیدند البته به شرط انکه این مقادیر در چهارچوب محدودیتهای تشریح شده از سوی ما باشند.بعدها درهمین پایان نامهدر مورد پیامدهای عملکرد بالقوه این انتخابها بحث خواهیم کرد اما در حال حاضر برای انتخاب 128 رشته در هر بلوک و 128 بلوک به همین میزان بسنده میکنیم. حال می توانیم بردارهایی با طول دلخواه که به مقدار GPUرویRAMمحدود شده اند را اضافه کنیم. در اینجا لیست کل منبع موجود اس:

 

 

 

Compatability by:
آخرین به روز رسانی سایت: سه شنبه, 22 اسفند 1391 - 00:26