جمعه 31 فروردین 1397 | Friday 20 th of April 2018 صفحه اصلی گروه الکترونیکی کامپیوتر
7ـ3ـ1 ضرب نقطه ای

حال حاضر باید نگاهی هم به ضرب نقطه ایبیندازیم (که گاهی اوقات به ان ضرب داخلیهمگفته میشود). ما ابتدا در مورد اینکه اصلا ضرب نقطه ایچیست انهم فقط در مواردی که شما با ریاضیات برداری نااشنا هستید (و یا جزء مسائلی است که اخیرا به ان پرداخته شده است) دست به یک بررسی اجمالی میزنیم. محاسبه از دو مرحله تشکیل شده است.در مرحلهاول، عناصر متناظر دو بردار ورودی را ضرب می کنیم.این بسیار بهجمع برداری شبیه استاما بجای جمع از ضرب بهره میجوییم.با این حال بجای ذخیره سازی بعدی این مقادیر در یک بردار خروجی سوم، میتوانیم همه انها را بمنظور تولید یک خروجی عددی واحدجمع کنیم. به عنوان مثال، اگر ضرب نقطه ایدو بردار چهار عنصره را صورت دهیم به معادله 7ـ1 خواهیم رسید.

معادله7ـ1

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


همانطور که می بینید ما یک بافر حافظه مشترک به نام حافظه موقت (cache) را معرفی کرده ایم.این بافر برای ذخیره جمع جاری هر رشته استفاده می شود.بزودی به دلیل انجام این کار پی خواهید برد اما فعلا فقط به بررسی مکانیکهایی میپردازیم که با ان میتوانیم کار را تکمیل کنیم.  شناسایی و اعلام یک متغیر برای اقامت در حافظه مشترک از اهمیت چندانی برخوردار نیست و بیشتر به شیوه ای شباهت و بستگی دارد که شما بوسیله ان یک متغیر را در استانداردC  بعنوان استاتیک و یا فرار اعلام میکنید:

ما ارایه اندازه ThreadsPerBlockرا اعلام می کنیم بنابراین هر رشته در بلوک دارای یک محل برای ذخیره نتیجه موقت خود میباشد.  به خاطر بیاورید که وقتی ما حافظه را بصورت جهانی اختصاص میدهیم به اندازه کافی برای هر رشته اجرا کننده هسته یاthreadsPerBlock که تعداد کل بلوک ها را زمانبندی می کند حافظه اختصاص میدهیم.اما از انجا که کامپایلر یک کپی از متغیرهای مشترک را برای هر بلوک ایجاد میکند، تنها به تخصیص کافی حافظه نیاز است بنحوی که هر رشته در بلوک از یک ورودی برخوردار باشد. پس از تخصیص حافظه مشترک درست مثل گذشته نوبت به محاسبه شاخص داده ها میرسد:

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

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

 هر رشته یک جمع جاری حاصلضرب ورودیهای متناظر را درa وbمحاسبه میکند.پس از رسیدن به انتهای ارایه، هر رشته به ذخیره جمع موقت خود در بافر مشترک مشغول میشود.

 در این نقطه از الگوریتم به جمع تمامی مقادیر موقت جای داده شده در کش نیاز است.برای انجام این کار نیز به برخی از رشته ها برای خواندن مقادیر ذخیره شده احتیاج است.با این حال همانطور که اشاره شد، این یک عمل به طور بالقوه خطرناک است.ما به روشی نیازمندیم که همه نوشته های انجام شده در کش ارایه مشترک [] را پیش از اینکه کسی برای خواندن انها تلاش کند بطور کامل تضمین کند. خوشبختانه چنین روشی وجود دارد:

این فراخوانی تضمین می کند که هر رشته در بلوک، دستورهای قبل ازsyncthreads __ () را پیش از اجرای دستور بعدی توسط سخت افزار بر روی هررشته به اتمام می رسانند.این دقیقا همان چیزی است که ما به ان نیاز داریم!ما اکنون می دانیم که وقتی رشته اول پس از syncthreads __ ()به اجرای دستور اول میپردازد هر رشته دیگر در بلوک نیز به پایان اجرای خود تا مرز syncthreads __ ()میرسد. حال که از پر شدن کش موقت اطمینان حاصل کرده ایم، می توانیم مقادیر مربوطه را در ان جمع بزنیم.ما روند کلی گرفتن یک ارایه ورودی و انجام بعضی از محاسبات تولید کننده ارایه کوچکتر نتایج یک کاهش را فراخوانی میکنیم.کاهشها اغلب در محاسبه موازی که به نامگذاری انها منجر میشود، بوجود می ایند.ساده ترین راه برای تکمیل این کاهش، تکرار یک رشته در حافظه مشترک و محاسبه جمع جاری میباشد. این امر متناسب با طول ارایه زمانبر است.با این حال، از انجایی که ما دارای صدهارشتهدر دسترس برای انجام کارهایمان هستیم، انجام کاهش بصورت موازی صورت میپذیرد و به زمانی متناسب با لگاریتم طول ارایه هم نیاز است. در ابتدا، کد زیر بصورت حلقوی بنظر خواهید رسید و ما در یک لحظه به شکستن ان مبادرتمی ورزیم. ایده کلی این است که هر رشته دو مقدار را در کش [] اضافه میکند و نتیجه را بصورت ذخیره شده به کش []باز می گرداند. از انجا که هر رشته دو ورودی را به یکی ترکیب می کند، این مرحله با نصف تعدادی که کار با ان اغاز شده به پایان میرسد.در مرحله بعد، ما همان کار را در نیمه باقیمانده انجام میدهیم. ما به همین طریق برای مراحلlog2 (threadsPerBlock) ادامه میدهیم تا زمانی که جمع هر ورودی در کش  []حاصل اید. برای مثال، در هر بلوک از 256 رشته استفاده میشود پس این روند باید 8 تکرار داشته باشد تا 256 ورودی در کش [] به جمعی واحد کاهش یابد.


شکل 7ـ4 یک مرحله از کاهش جمع

در اولین گام، کار را با iبه عنوان نیمی از تعداد threadsPerBlock  اغاز میکنیم.ما فقط رشته هایی را میخواهیم که با شاخصهای کمتر از این مقدار هر کاری انجام دهند بنابراین اگر شاخص رشته کمتر از iباشد، دو ورودی کش [] بصورت مشروط اضافه میشوند.علاوه بر این، ما از انجام جمع در داخلبلوک if(cacheIndex < i)  پشتیبانی می کنیم.هر رشته یک ورودی را از شاخص خود در کش [] دریافت میکند، انرا توسط iبه ورودی افست شاخص خود اضافه میکند و ذخیره شده این مجموع  را به کش []بازمی گرداند.

بافرض اینکه هشت ورودی در کش [] وجود دارد، iمقدار 4 را بعنوان نتیجه خواهد داشت.  یک مرحله کاهش مانند شکل7ـ4 خواهد بود.

پس از اینکه یک مرحله را به اتمام رساندیم با محدودیتهای مشابهی که پس از محاسبه ضربهای جفتی با انها روبروبودیم باز هم مواجه خواهیم بود. قبل از اینکه بتوانیم مقادیر ذخیره شده در کش []را بخوانیم باید مطمئن شویم هر رشته که نیاز به نوشتن در کش [] دارد قبلا این کار را انجام داده است.__syncthreads()  پس از واگذاری، برقراری این شرایط را تضمین می کند.

 پس از ختم حلقه while ()، هر بلوک دارای فقط یک تعداد واحدباقی مانده است.این تعداد در اولین ورودی کش [] جای میگیرد و حاصل جمع هر ضرب جفتی رشته ها در ان بلوک محاسبه میشود.سپس این مقدار واحد در حافظه جهانی ذخیره و در نهایت هسته به انتها می رسد:

چرا این ذخیره جهانی را تنها برای رشته دارای cacheIndex == 0انجام می دهیم؟ خب، از انجا که تنها یک عدد وجود دارد که به حافظه جهانی نیاز دارد تنها یک رشته واحد هم به انجام این عملیات نیاز خواهد داشت.هر رشته می تواند بطور باور نکردنی به این نوشتن مشغول شود و برنامه هم همچنان به کار کردن ادامه دهد اما انجام این کار بصورت غیرضروری باعث بالا رفتن ترافیک حافظه در زمان نوشتن یک مقدار واحد میشود. برای سادگی کار، رشته ای با شاخص 0 انتخاب میشود. هر چند که انتخابcacheIndex  بمنظور نوشتن کش [0] برای حافظه جهانی هم امکان پذیر است.در نهایت، از انجا که هر بلوک دقیقا یک مقدار را برای ارایه جهانی c[]مینویسد، می توانیم به سادگی انرا بوسیلهblockIdx شاخص سازی کنیم.

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

ما پیش از این به عملیاتی مانند ضرب نقطه ای به عنوان یک کاهش اشاره کرده بودیم. قطع نظراز جزئیات، دلیل واقعیان این است که ما نسبت به ورودیها به تولید عناصر داده ای خروجی کمتری میپردازیم.ما همیشه در ضرب نقطه ای صرف نظر از اندازه ورودیمان دقیقا به تولید یک خروجی می پردازیم.به نظر می رسد که دستگاه / ماشین انبوه موازی مانند یکGPU به هدر دادن منابع خود در هنگام انجام مراحل اخر کاهشتمایل نشان دهدزیرا اندازه مجموعه داده ها در ان نقطه بسیار کوچک است و استفاده از 480 واحد محاسباتی برای اضافه کردن 32 عدد کار مشکلی است!

 به همین دلیل ما کنترل را بار دیگر به میزبان میسپاریم و میگذاریمCPU مرحله نهایی جمع را در حین جمع ارایه  c[]به پایان برساند. حالا GPUدر یک نرم افزار بزرگتر، برای شروع یک ضرب نقطه ای دیگر و یا کار بر روی یک محاسبه گسترده دیگر از ازادی عمل بیشتری برخوردار خواهد بود.هر چند که در این مثال، بهGPU کفایت شده است.

در توضیح این مثال، ما سنت شکنی کرده و یکسره بسراغ محاسبه کرنل واقعی میرویم.ما امیدواریم که شما از بدنه اصلیmain()  گرفته تا فراخوانی کرنل با هیچ مشکل درکی برخورد نکرده باشید زیرا بسیار شبیه به انچیزی است که ما قبلا نشان داده ایم.


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

1)    اختصاص میزبان و حافظه دستگاهی به ارایه های ورودی و خروجی.

2)    پر کردن ارایه های ورودی a[]  وb[]و سپس کپی انها به دستگاه با استفاده از  cudaMemcpy().

3)    فراخوانی کرنل ضرب نقطه ای با استفاده از تعداد معینی رشته در هر بلوک و بلوک در هر شبکه.

با وجود اینکه بسیاری از این مسائل در حال حاضر برای شما امری عادی تلقی میشود ولی باز هم ارزش بررسی محاسبات صورت گرفته بر روی تعدادی از بلوک های راه اندازی شده را دارد.ما تا کنون در مورد نحوه ای که ضرب نقطه ای میتواند بعنوان یک کاهش عمل کند و هر بلوک راه اندازی شده میتواند پاره‌ حاصل‌ جمع‌ جزئی‌را محاسبه کند، شرح داده ایم.طول این لیست پاره‌ مجموعهبرایCPU باید بصورت کنترلی کوچک انتخاب شود هرچند که این اندازه هنوز انقدر بزرگ هست که بتواند بلوک های کافی جهت مشغول نگه داشتن حتی سریع ترینGPU ها را در بربگیرد.ما 32 بلوک را انتخاب کرده ایم اگر چه این موردی است که در ان شما ممکن است به عملکرد بهتر و یا بدتر انتخاب های دیگر، به ویژه با توجه به سرعتهای نسبیCPU وGPU پی ببرید.

 اما چه میشد اگر به ما فهرست بسیار کوتاهی داده میشد که در ان 32 بلوک 256 قطعه ای بسیار زیادتر از حد معمول بنظر میرسیدند؟ اگر ما از عناصر اطلاعاتی Nبرخوردار بودیم به منظور محاسبه ضرب نقطه ای تنها به رشته های  Nنیاز میداشتیم. بنابراین در این مورد به کوچکترین مضرب threadsPerBlock  بزرگتر از Nیا مساوی باان نیاز است. ما یک بار قبل از اضافه کردن بردارها هم شاهد این امر بودیم.در این مورد کوچکترین مضربthreadsPerBlock بدست می اید که مقدار ان با محاسبه(N+ (threadsPerBlock-1)) / threadsPerBlock  بزرگتر ازN  یا مساوی با ان میباشد.چیزی که شما نیز ممکن است با خود بگویید در واقع یک ترفند نسبتا رایج در ریاضی عدد صحیح میباشد، بنابراین حتی اگر شما بیشتر وقت خود را صرف کار کردن در حوزه ای خارج از CUDA Cکنید باز هم دانستن و درک ان خالی از لطف نیست.

بنابراین، تعدادی از بلوکهای راه اندازی شده باید به تناسب اینکه کدام از مقدار کمتری برخوردارند، 32 و یا(N + (threadsPerBlock-1)) /threads Per Blockرشته ای باشند.

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

                        

در نهایت، نتایج بررسی و حافظه اختصاص داده شده به CPU  وGPUپاک میشود. بررسی نتایج بدست امده اسان است زیرا ورودیها با داده های قابل پیش بینی پر میشوند.اگر بخاطر بیاوریدa[]  با اعداد صحیح از 0 تاN-1 پر میشود و  b[]هم بصورت 2*a[]میباشد.

ضرب نقطه ای ما باید دو برابر مجموع مربعات اعداد صحیح از 0 تاN-1 باشد. برای خواننده ای که عاشق ریاضیات گسستهاست این یک تفریحسرگرم کننده برای نتیجه گیری از روش تحلیل نوع بستهاین جمع است و در اینجابرای کسانی که از صبر یا علاقه کمتری برخوردارند، روش تحلیل نوع بستهو همچنین بقیه بدنه اصلی  main()ارائه شده است:

 اگر بنظر شما تمامی وقفه های توضیحی ما به نوعی پر دردسر شده اند، در اینجا کل فهرست منبع انهم بدون تفسیر ارائه شده است:

 

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