مروری بر CUDA – قسمت چهارم : مدل برنامه نویسی CUDA

این چهارمین مقاله از مجموعه مروری برCUDA  است که هدف از آن ، بروز کردن مفاهیم کلیدی در CUDA ، ابزارها و بهینه سازی، برای توسعه دهندگان مبتدی یا متوسط است.

مدل برنامه نویسی CUDA چکیده ای از معماری GPU را ارائه می دهد که به عنوان یک پل ارتباطی بین برنامه و پیاده سازی احتمالی آن بر روی سخت افزار GPU عمل می کند. این مقاله با بیان چگونگی استفاده از زبان های برنامه نویسی عمومی مانند C / C ++ ، مفاهیم اصلی مدل برنامه نویسی CUDA را بیان می کند.

اجازه دهید دو کلید واژه را که به طور گسترده در مدل برنامه نویسی CUDA استفاده می شود معرفی کنیم : میزبان ( Host ) و دستگاه ( Device )

میزبان ،CPU موجود در سیستم است. حافظه سیستم مرتبط با CPU را حافظه میزبان می نامند. GPU را، دستگاه و حافظه GPU را نیز، حافظه دستگاه نامیده می شوند.

برای اجرای هر برنامه CUDA ، سه مرحله اصلی وجود دارد :

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

سلسله مراتب رشته و هسته CUDA

شکل ۱ نشان می دهد که هسته CUDA تابعی است که در GPU اجرا می شود. بخش موازی برنامه های شما K بار به صورت موازی با K رشته ( Thread ) مختلف CUDA اجرا می شود، در حالی که توابع C / C ++  معمول ، تنها یک بار این کار را انجام می دهند.

هر هسته ( کرنل ) CUDA با یک شناساگر تعریف _global_  شروع می شود. برنامه نویسان با استفاده از متغیرهای داخلی یک شناسه (ID) جهانی منحصر به فرد برای هر رشته ارائه می دهند.

به گروهی از رشته ها، یک بلوک CUDA گفته می شود. بلوک های CUDA در یک شبکه گروه بندی می شوند. یک هسته به عنوان شبکه ای از بلوک های رشته اجرا می شود (شکل ۲).

هر بلوک CUDA توسط یک چندپردازنده جریانی (SM) اجرا می شود و نمی تواند به سایر SM ها در GPU منتقل شود ( به جز در هنگام preemption ، اشکال زدایی یا موازی سازی دینامیکی  CUDA). یک SM می تواند بسته به منابع مورد نیاز بلوک های CUDA ، چندین بلوک همزمان CUDA را اجرا کند. هر هسته روی یک دستگاه اجرا می شود و CUDA از اجرای چندین هسته همزمان در یک دستگاه پشتیبانی می کند. شکل ۳ اجرای هسته و نگاشت بر منابع سخت افزاری موجود در GPU را نشان می دهد.

CUDA متغیرهای سه بعدی داخلی را برای رشته ها و بلوک ها تعریف می کند. رشته ها با استفاده از متغیر داخلی سه بعدی threadIdx شاخص گذاری می شوند. شاخص گذاری سه بعدی روشی طبیعی برای مشخص کردن عناصر در بردار ها ، ماتریس ها و حجم ها ایجاد می کند و باعث می‌شود برنامه نویسی CUDA ساده تر شود. به طور مشابه ، بلوک ها نیز با استفاده از متغیر سه بعدی داخلی به نام blockIdx  شاخص گذاری می شوند.

در اینجا چند نکته قابل توجه وجود دارد:

• معماری CUDA تعداد رشته ها در هر بلوک را محدود می کند ( ۱۰۲۴ رشته در هر محدوده بلوک ).
• بُعد بلوک رشته ای از طریق متغیر داخلی BlockDim در هسته قابل دسترسی است.
• همه رشته های درون یک بلوک را می توان با استفاده از یک تابع فطری __syncthreads همزمان سازی کرد. با __syncthreads ، همه رشته های موجود در بلوک باید منتظر بمانند تا کسی بتواند ادامه دهد.
• تعداد رشته ها در هر بلوک و تعداد بلوک ها در هر شبکه مشخص شده با دستور <<< … >>> می تواند از نوع int یا dim3 باشد. این براکت های سه تایی ، فراخوانی از کد میزبان به کد دستگاه را نشان می دهند. به این کار راه اندازی هسته نیز گفته می شود.

برنامه CUDA برای افزودن دو ماتریس زیر blockIdx و threadIdx چند بعدی و متغیرهای دیگر مانند blockDim را نشان می دهد. در مثال زیر ، یک بلوک دو بعدی برای سهولت شاخص گذاری انتخاب شده است و هر بلوک ۲۵۶ رشته دارد که هر یک در جهت x و y قرار دارند. تعداد کل بلوک ها با استفاده از اندازه داده تقسیم بر اندازه هر بلوک محاسبه می شود.

// Kernel - Adding two matrices MatA and MatB
__global__ void MatAdd(float MatA[N][N], float MatB[N][N],
float MatC[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        MatC[i][j] = MatA[i][j] + MatB[i][j];
}
 
int main()
{
    ...
    // Matrix addition kernel launch from host code
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks((N + threadsPerBlock.x -1) / threadsPerBlock.x, (N+threadsPerBlock.y -1) / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(MatA, MatB, MatC);
    ...
}

سلسله مراتب حافظه

GPU های دارای قابلیت CUDA ، دارای یک سلسله مراتب حافظه هستند که در شکل ۴ نشان داده شده است.

حافظه های زیر توسط معماری GPU ارائه می شوند :

• رجیستر ها ( Regidter ) : این ها برای هر رشته اختصاصی هستند ، به این معنی که رجیستر های اختصاص داده شده به یک رشته برای سایر رشته ها قابل مشاهده نیستند. کامپایلر درباره استفاده از رجیستر تصمیم گیری می کند.
• حافظه اشتراکی L1 / (SMEM) : هر SM دارای یک حافظه چرکنویس ( Scratchpad Memory ) سریع و بر روی تراشه است که می تواند به عنوان حافظه پنهان L1 و حافظه اشتراکی مورد استفاده قرار گیرد. تمام رشته های موجود در یک بلوک CUDA می توانند حافظه اشتراکی را به اشتراک بگذارند و همه بلوک های CUDA که روی یک SM خاص اجرا می شوند می توانند منبع حافظه فیزیکی ارائه شده توسط SM را به اشتراک بگذارند.
• حافظه فقط خواندنی : هر SM دارای یک حافظه پنهان دستورالعمل ، حافظه ثابت ، حافظه بافت و حافظه پنهان RO است که فقط برای کد هسته قابل خواندن است.
• حافظه پنهان L2 : حافظه پنهان L2 در تمام SM ها به اشتراک گذاشته شده است ، بنابراین هر رشته در هر بلوک CUDA می تواند به این حافظه دسترسی داشته باشد. پردازنده GPU NVIDIA A100 اندازه حافظه پنهان L2 را در مقایسه با ۶ مگابایت پردازنده های V100 GPU ، به ۴۰ مگابایت افزایش داده است.
• حافظه جهانی : این اندازه بافر فریم GPU و DRAM است که در GPU قرار دارد.

کامپایلر NVIDIA CUDA کار خوبی در بهینه سازی منابع حافظه انجام می دهد اما یک توسعه دهنده CUDA حرفه ای می تواند در صورت لزوم استفاده از این سلسله مراتب حافظه را به طور کارآمد برای بهینه سازی برنامه های CUDA انتخاب کند.

قابلیت محاسبه

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

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

برای کسب اطلاعات بیشتر در مورد قابلیت محاسبه هر دستگاه دارای CUDA ، به کد نمونه  CUDA، deviceQuery نگاه بیندازید. این نمونه ویژگی های دستگاه های CUDA موجود در سیستم را مشخص می‌کند.

خلاصه

مدل برنامه نویسی CUDA یک محیط ناهمگون را فراهم می کند که در آن کد میزبان برنامه C / C ++  را بر روی CPU اجرا می کند و هسته روی یک دستگاه GPU که از لحاظ فیزیکی جداشونده هستند، اجرا می شود. مدل برنامه نویسی CUDA همچنین فرض می کند که هم میزبان و هم دستگاه دارای فضای حافظه جداگانه ‌ای هستند که به ترتیب از آن ها به عنوان حافظه میزبان و حافظه دستگاه یاد می شود. کد CUDA همچنین امکان انتقال داده ها بین حافظه میزبان و دستگاه را از طریق گذرگاه PCIe فراهم می کند.

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

بیشتر بخوانید :