در پست قبلی ، من بررسی کردم که چگونه دسترسی به حافظه جهانی توسط گروهی از موضوعات می تواند به یک معامله واحد متصل شود ، و چگونه تراز و قدم بر همبستگی برای نسل های مختلف سخت افزار CUDA تأثیر می گذارد. برای نسخه های اخیر سخت افزار CUDA ، دسترسی به داده های نادرست مسئله بزرگی نیست. با این حال ، قدم زدن از طریق حافظه جهانی بدون در نظر گرفتن تولید سخت افزار CUDA مشکل ساز است ، و به نظر می رسد در بسیاری موارد غیرقابل اجتناب است ، مانند دسترسی به عناصر در یک آرایه چند بعدی در طول ابعاد دوم و بالاتر. با این وجود ، در صورت استفاده از حافظه مشترک ، می توان دسترسی به حافظه را در چنین مواردی همبستگی کرد. قبل از اینکه به شما نشان دهم که چگونه می توانید در پست بعدی از حافظه جهانی خودداری کنید ، ابتدا باید حافظه مشترک را با جزئیات توصیف کنم.

حافظه مشترک

از آنجا که این تراشه است ، حافظه مشترک بسیار سریعتر از حافظه محلی و جهانی است. در حقیقت ، تأخیر در حافظه مشترک تقریباً 100 برابر کمتر از تأخیر حافظه جهانی غیرقابل دسترسی است (مشروط بر اینکه هیچ درگیری بانکی بین موضوعات وجود ندارد ، که بعداً در این پست بررسی خواهیم کرد). حافظه مشترک در هر بلوک موضوع اختصاص می یابد ، بنابراین تمام موضوعات موجود در بلوک به همان حافظه مشترک دسترسی دارند. موضوعات می توانند به داده های مشترک بارگیری شده از حافظه جهانی توسط موضوعات دیگر در همان بلوک موضوع دسترسی پیدا کنند. این قابلیت (همراه با همگام سازی موضوع) دارای چندین کاربرد است ، مانند حافظه پنهان داده های کاربر ، الگوریتم های موازی تعاونی با کارایی بالا (به عنوان مثال کاهش موازی) و برای تسهیل همبستگی حافظه جهانی در مواردی که در غیر این صورت نخواهد بودممکن است

همگام سازی موضوع

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

برای اطمینان از نتایج صحیح در هنگام همکاری موضوعات موازی ، باید موضوعات را همگام سازی کنیم. CUDA یک همگام سازی سد ساده را ارائه می دهد ، __syncthreads (). اجرای یک موضوع فقط می تواند از گذشته __syncthreads () عبور کند ، پس از اینکه همه موضوعات موجود در بلوک خود ، __syncthreads () را اجرا کرده اند. بنابراین ، ما می توانیم با فراخوانی __syncthreads () پس از فروشگاه به حافظه مشترک و قبل از بارگیری هر موضوع از حافظه مشترک ، از وضعیت مسابقه که در بالا توضیح داده شد ، جلوگیری کنیم. این مهم است که آگاه باشید که فراخوانی __syncthreads () در کد واگرا تعریف نشده است و می تواند منجر به بن بست شود - همه موضوعات در یک بلوک موضوع باید در همان نقطه با __syncthreads () تماس بگیرند.

مثال حافظه مشترک

حافظه مشترک را در کد دستگاه CUDA C/C ++ با استفاده از مشخصات اعلامیه متغیر __shared__ اعلام کنید. بسته به اینکه میزان حافظه در زمان کامپایل یا در زمان اجرا شناخته شده باشد ، چندین روش برای اعلام حافظه مشترک در داخل هسته وجود دارد. کد کامل زیر (موجود در GitHub) روشهای مختلف استفاده از حافظه مشترک را نشان می دهد.

#include __global__ void staticreverse (int *d ، int n)<__shared__ int s[64]; int t = threadIdx.x; int tr = n-t-1; s[t] = d[t]; __syncthreads(); d[t] = s[tr];>__global__ void dynamicreverse (int *d ، int n)int main (باطل)int *d_d; cudaMalloc(&d_d, n * sizeof(int)); // run version with static shared memory cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice); staticReverse>>(d_d ، n) ؛cudamemcpy (d ، d_d ، n*sizeof (int) ، cudamemcpydevicetohost) ؛برای (int i = 0 ؛ i>(d_d ، n) ؛cudamemcpy (d ، d_d ، n * sizeof (int) ، cudamemcpydevicetohost) ؛برای (int i = 0 ؛ i

این کد با استفاده از حافظه مشترک ، داده ها را در یک آرایه 64 عنصر معکوس می کند. این دو هسته بسیار مشابه هستند و فقط در نحوه اعلام آرایه های حافظه مشترک و نحوه فراخوانی هسته ها متفاوت هستند.

حافظه مشترک استاتیک

اگر اندازه آرایه حافظه مشترک در زمان کامپایل شناخته شده باشد ، مانند هسته staticreverse ، می توانیم صریحاً آرایه ای از آن اندازه را اعلام کنیم ، همانطور که با آرایه انجام می دهیم.

__global__ void staticreverse (int *d ، int n)

در این هسته ، T و TR به ترتیب دو شاخص هستند که ترتیب اصلی و معکوس را نشان می دهند. موضوعات داده ها را از حافظه جهانی تا حافظه مشترک با عبارت S [T] = D [T] کپی می کنند ، و وارونگی بعد از دو خط بعد با عبارت D [T] = S [TR] انجام می شود. اما قبل از اجرای این خط نهایی که در آن هر موضوع به داده های مشترک که توسط یک موضوع دیگر نوشته شده است ، به یاد داشته باشید که باید اطمینان حاصل کنیم که همه موضوعات با فراخوانی __syncthreads () بارهای حافظه مشترک را تکمیل کرده اند.

دلیل استفاده از حافظه مشترک در این مثال ، تسهیل همبستگی حافظه جهانی در دستگاه های قدیمی CUDA (قابلیت محاسبه 1. 1 یا قبل از آن) است. همبستگی حافظه جهانی بهینه برای هر دو خواندن و نوشتن حاصل می شود زیرا حافظه جهانی همیشه از طریق شاخص خطی و تراز شده T قابل دسترسی است. از شاخص معکوس TR فقط برای دسترسی به حافظه مشترک استفاده می شود ، که محدودیت دسترسی پی در پی حافظه جهانی را برای عملکرد بهینه ندارد. تنها مسئله عملکرد با حافظه مشترک ، درگیری های بانکی است که بعداً در مورد آن صحبت خواهیم کرد.(توجه داشته باشید که در دستگاه های قابلیت محاسبه 1. 2 یا بالاتر ، سیستم حافظه می تواند حتی فروشگاه های شاخص معکوس را به حافظه جهانی هماهنگ کند. اما این تکنیک هنوز هم برای سایر الگوهای دسترسی مفید است ، همانطور که در پست بعدی نشان خواهم داد.)

حافظه مشترک پویا

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

dynamicReverse>>(d_d ، n) ؛

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

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

خارجی __shared__ int s [] ؛int *integerdata = s ؛// ni ints float *floatdata = (float *) & integerdata [ni] ؛// nf floats char *chardata = (char *) & floatdata [nf] ؛// NC Chars

در راه اندازی هسته ، کل حافظه مشترک مورد نیاز را مانند موارد زیر مشخص کنید.

myKeel>>(. );

درگیری های بانک حافظه مشترک

برای دستیابی به پهنای باند حافظه بالا برای دسترسی های همزمان ، حافظه مشترک به ماژول های حافظه به همان اندازه (بانک ها) تقسیم می شود که به طور همزمان قابل دسترسی هستند. بنابراین ، هرگونه بار حافظه یا فروشگاه آدرس های N که شامل بانکهای حافظه مجزا است ، می توانند به طور همزمان سرویس شوند و یک پهنای باند مؤثر ارائه دهند که بارها از پهنای باند یک بانک واحد باشد.

با این حال، اگر آدرس‌های درخواستی چندین رشته به یک بانک حافظه نگاشت شوند، دسترسی‌ها سریالی می‌شوند. سخت‌افزار یک درخواست حافظه متضاد را به تعداد درخواست‌های بدون تداخل جداگانه تقسیم می‌کند و پهنای باند مؤثر را با یک عامل برابر با تعداد درخواست‌های حافظه درگیر کاهش می‌دهد. یک استثنا موردی است که در آن همه رشته‌ها در یک Warp آدرس حافظه مشترک یکسانی را نشان می‌دهند که منجر به پخش می‌شود. دستگاه‌های با قابلیت محاسباتی 2. 0 و بالاتر، توانایی بیشتری برای چندپخشی دسترسی‌های حافظه مشترک دارند، به این معنی که دسترسی‌های متعدد به یک مکان توسط هر تعداد رشته در یک Warp به طور همزمان ارائه می‌شوند.

برای به حداقل رساندن تضادهای بانکی، مهم است که بدانیم آدرس های حافظه چگونه به بانک های حافظه نگاشت می شوند. بانک های حافظه مشترک به گونه ای سازماندهی شده اند که کلمات 32 بیتی متوالی به بانک های متوالی اختصاص می یابد و پهنای باند آن 32 بیت در هر بانک در هر چرخه ساعت است. برای دستگاه هایی با قابلیت محاسبه 1. x، اندازه چرخش 32 رشته و تعداد بانک ها 16 است. یک درخواست حافظه مشترک برای یک Warp به یک درخواست برای نیمه اول Warp و یک درخواست برای نیمه دوم تقسیم می شود. تارتوجه داشته باشید که اگر فقط یک مکان حافظه در هر بانک توسط یک نیم پیچ از رشته ها قابل دسترسی باشد، درگیری بانکی رخ نمی دهد.

برای دستگاه‌های با قابلیت محاسبه 2. 0، اندازه چرخش 32 رشته و تعداد بانک‌ها نیز 32 است. درخواست حافظه مشترک برای Warp مانند دستگاه‌های با قابلیت محاسبه 1. x تقسیم نمی‌شود، به این معنی که درگیری بانکی می‌تواند بین رشته‌ها رخ دهد. در نیمه اول یک تار و نخ ها در نیمه دوم همان تار.

دستگاه‌های با قابلیت محاسبه 3. x دارای اندازه بانک قابل تنظیم هستند که می‌توانند با استفاده از cudaDeviceSetSetSharedMemConfig() روی چهار بایت (cudaSharedMemBankSizeFourByte، پیش‌فرض) یا هشت بایت (cudaSharedMemBankSizeEightByte) تنظیم شوند. تنظیم اندازه بانک بر روی هشت بایت می تواند به جلوگیری از درگیری بانک حافظه مشترک هنگام دسترسی به داده های با دقت مضاعف کمک کند.

پیکربندی مقدار حافظه مشترک

در دستگاه هایی با قابلیت محاسبه 2. x و 3. x، هر چند پردازنده دارای 64 کیلوبایت حافظه روی تراشه است که می تواند بین کش L1 و حافظه مشترک تقسیم شود. برای دستگاه‌های با قابلیت محاسبه 2. x، دو تنظیم وجود دارد، 48 کیلوبایت حافظه مشترک / 16 کیلوبایت حافظه نهان L1، و 16 کیلوبایت حافظه مشترک / 48 کیلوبایت حافظه نهان L1. به طور پیش فرض از تنظیمات حافظه مشترک 48 کیلوبایت استفاده می شود. این را می توان در طول اجرای API از میزبان برای همه هسته ها با استفاده از cudaDeviceSetCacheConfig() یا بر اساس هر هسته با استفاده از cudaFuncSetCacheConfig پیکربندی کرد. اینها یکی از سه گزینه را می پذیرند: cudaFuncCachePreferNone، cudaFuncCachePreferShared، و cudaFuncCachePreferL1. درایور ترجیحات مشخص شده را رعایت می کند مگر زمانی که یک هسته به حافظه مشترک بیشتری در هر بلوک رشته نیاز داشته باشد نسبت به آنچه در پیکربندی مشخص شده موجود است. دستگاه‌های با قابلیت محاسبه 3. x امکان تنظیم سوم حافظه مشترک 32 کیلوبایت / حافظه نهان 32 کیلوبایتی L1 را می‌دهند که می‌توان با استفاده از گزینه cudaFuncCachePreferEqual به دست آورد.

خلاصه

حافظه مشترک یک ویژگی قدرتمند برای نوشتن کد CUDA بهینه شده است. دسترسی به حافظه مشترک بسیار سریعتر از دسترسی به حافظه جهانی است زیرا روی تراشه قرار دارد. از آنجایی که حافظه مشترک توسط رشته‌هایی در یک بلوک رشته به اشتراک گذاشته می‌شود، مکانیسمی برای همکاری رشته‌ها فراهم می‌کند. یکی از راه‌های استفاده از حافظه مشترک که از چنین همکاری رشته‌ای استفاده می‌کند، فعال کردن ادغام حافظه جهانی است، همانطور که با معکوس کردن آرایه در این پست نشان داده شده است. با معکوس کردن آرایه با استفاده از حافظه مشترک، می‌توانیم تمام خواندن و نوشتن حافظه سراسری را با گام واحد انجام دهیم و به ادغام کامل در هر پردازنده گرافیکی CUDA دست یابیم. در پست بعدی بحث خود را در مورد حافظه اشتراکی با استفاده از آن برای بهینه سازی انتقال ماتریس ادامه خواهم داد.

منابع مرتبط

• دوره DLI: تسریع برنامه های CUDA C++ با چند پردازنده گرافیکی
• دوره DLI: تسریع برنامه های CUDA C++ با جریان های همزمان
• جلسه GTC: CUDA: ویژگی های جدید و فراتر از آن (بهار 2023)
• جلسه GTC: ارتباط با کارشناسان: در جعبه ابزار CUDA شما چیست؟ابزارهای پروفایل سازی، بهینه سازی و اشکال زدایی CUDA برای جدیدترین معماری ها (بهار 2023)
• SDK: CUDA Toolkit
• وبینار: معرفی مدرن به CUDA