فصل 2: مدل‌های برنامه‌نویسی GPU

واحدهای پردازش گرافیکی (GPU) از شتاب‌دهنده‌های گرافیکی با عملکرد ثابت به موتورهای محاسباتی قابل برنامه‌ریزی با درجه بالای موازی‌سازی تکامل یافته‌اند که می‌توانند طیف گسترده‌ای از برنامه‌ها را تسریع بخشند. برای توانمند کردن برنامه‌نویسان در بهره‌برداری مؤثر از موازی‌سازی گسترده در GPU ها، چندین مدل برنامه‌نویسی موازی و API ها توسعه یافته‌اند، مانند NVIDIA CUDA، OpenCL و DirectCompute. این مدل‌های برنامه‌نویسی انتزاعاتی را ارائه می‌کنند که به برنامه‌نویسان امکان بیان موازی‌سازی در برنامه‌های خود را می‌دهند، در حالی که جزئیات سطح پایین سخت‌افزار GPU را پنهان می‌کنند.

در این فصل، مفاهیم و اصول کلیدی پشت مدل‌های برنامه‌نویسی موازی برای GPU ها را بررسی خواهیم کرد، با تمرکز بر مدل اجرا، معماری مجموعه دستورالعمل GPU (ISA)، ISA های NVIDIA GPU و ISA Graphics Core Next (GCN) AMD. همچنین مثال‌هایی ارائه خواهیم داد تا نشان دهیم چگونه این مفاهیم در عمل اعمال می‌شوند.

مدل اجرا

مدل اجرای مدل‌های برنامه‌نویسی GPU مدرن بر اساس مفهوم کرنل است، که تابع‌هایی هستند که به طور موازی توسط تعداد زیادی رشته بر روی GPU اجرا می‌شوند. هنگام راه‌اندازی یک کرنل، برنامه‌نویس تعداد رشته‌هایی که باید ایجاد شوند و چگونگی سازماندهی آنها در سلسله مراتبی از شبکه ها، بلوک‌ها (یا آرایه‌های رشته همکاری - CTA) و رشته‌های فردی را مشخص می‌کند.

یک شبکه دامنه مسأله کامل را نمایش می‌دهد و از یک یا چند بلوک تشکیل شده است.
یک بلوک گروهی از رشته‌ها است که می‌توانند با یکدیگر همکاری و همگام‌سازی کنند، از طریق حافظه مشترک و سدها. رشته‌های درون یک بلوک بر روی هسته GPU یکسان (که چندپردازنده جریان یا واحد محاسباتی نامیده می‌شود) اجرا می‌شوند.
هر رشته دارای یک شناسه منحصربه‌فرد در داخل بلوک و شبکه خود است، که می‌تواند برای محاسبه آدرس‌های حافظه و اتخاذ تصمیمات جریان کنترل استفاده شود.

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

            شبکه
       ________________
      /   /   /   /   /
     /   /   /   /   / 
    /   /   /   /   /
   /   /   /   /   /
  /__/__/__/__/__/
    |   |   |   |
    |   | بلوک |
    |   |   |   |
  رشته رشته ...

شکل ۲.۱: سلسله‌مراتب رشته‌ها در مدل اجرای GPU.

اجرای SIMT

مدل‌های برنامه‌نویسی GPU مانند CUDA و OpenCL از یک مدل اجرای Single-Instruction, Multiple-Thread (SIMT) پیروی می‌کنند. در مدل SIMT، رشته‌ها در گروه‌های کوچکی به نام وارپ‌ها (اصطلاح NVIDIA) یا موج‌جبهه‌ها (اصطلاح AMD) اجرا می‌شوند. تمام رشته‌های درون یک وارپ همان دستور را همزمان اجرا می‌کنند، اما هر رشته روی داده‌های متفاوتی عمل می‌کند.

با این حال، برخلاف مدل Single-Instruction, Multiple-Data (SIMD) سنتی که در آن تمام عناصر پردازشی هم‌گام اجرا می‌شوند، SIMT به رشته‌ها اجازه می‌دهد که مسیرهای اجرای مستقل داشته باشند و در دستورات شاخه‌ای منشعب شوند. وقتی یک وارپ با یک دستور شاخه‌ای مواجه می‌شود، سخت‌افزار GPU شرط شاخه را برای هر رشته در وارپ ارزیابی می‌کند. اگر تمام رشته‌ها همان مسیر را انتخاب کنند (همگرا شوند)، وارپ به اجرای عادی ادامه می‌دهد. اگر برخی رشته‌ها مسیرهای متفاوتی را انتخاب کنند (واگرا شوند)، وارپ به دو یا چند زیروارپ تقسیم می‌شود که هر کدام یک مسیر متفاوت را دنبال می‌کنند. سخت‌افزار GPU اجرای مسیرهای متفاوت را سریالی می‌کند و رشته‌های غیرفعال را در هر زیروارپ خاموش می‌کند. وقتی همه مسیرها تکمیل شوند، زیروارپ‌ها دوباره همگرا شده و به اجرای هم‌گام ادامه می‌دهند.

شکل ۲.۲ اجرای SIMT را با جریان کنترل واگرا نشان می‌دهد.

         وارپ
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | شاخه |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
            \
             \
   همگرایی مجدد

شکل ۲.۲: اجرای SIMT با جریان کنترل واگرا.

این مکانیسم مدیریت واگرایی به SIMT اجازه می‌دهد تا جریان کنترل انعطاف‌پذیرتری را پشتیبانی کناینجا ترجمه فارسی آن است، با توجه به اینکه کد نباید ترجمه شود و فقط توضیحات باید ترجمه شوند:

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

GPU ها دارای سلسله مراتب پیچیده حافظه هستند تا به نیازهای پهنای باند بالا و تأخیر پایین کارهای موازی پاسخ دهند. سلسله مراتب حافظه معمولاً شامل موارد زیر است:

حافظه جهانی: بزرگترین اما کندترین فضای حافظه، قابل دسترسی برای همه رشته‌ها در یک کرنل. حافظه جهانی معمولاً با استفاده از حافظه GDDR یا HBM با پهنای باند بالا پیاده‌سازی می‌شود.
حافظه اشتراکی: یک فضای حافظه سریع و روی-تراشه که توسط همه رشته‌ها در یک بلوک به اشتراک گذاشته می‌شود. حافظه اشتراکی برای ارتباطات بین رشته‌ای و به اشتراک گذاری داده در داخل یک بلوک استفاده می‌شود.
حافظه ثابت: یک فضای حافظه فقط-خواندنی که برای پخش داده‌های فقط-خواندنی به همه رشته‌ها استفاده می‌شود.
حافظه بافت: یک فضای حافظه فقط-خواندنی که برای دسترسی به آن از حافظه کش بافت بهینه‌سازی شده است. حافظه بافت معمولاً بیشتر در کارهای گرافیکی استفاده می‌شود.
حافظه محلی: یک فضای حافظه خصوصی برای هر رشته، که برای سرریز رجیستر و ساختارهای داده بزرگ استفاده می‌شود. حافظه محلی معمولاً به حافظه جهانی پیوند داده می‌شود.

استفاده مؤثر از سلسله مراتب حافظه برای دستیابی به عملکرد بالا در GPU ها بسیار مهم است. برنامه‌نویسان باید به حداکثر رساندن استفاده از حافظه اشتراکی و به حداقل رساندن دسترسی به حافظه جهانی برای کاهش تأخیر حافظه و گلوگاه پهنای باند بپردازند.

شکل 2.3 سلسله مراتب حافظه GPU را نشان می‌دهد.اینجا ترجمه فارسی فایل مارک‌داون داده شده است:

معماری مجموعه دستورات GPU

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

در این بخش، ما ISA های دو تولید‌کننده اصلی GPU را بررسی خواهیم کرد: NVIDIA و AMD. ما بر روی ISA های Parallel Thread Execution (PTX) و SASS NVIDIA و Graphics Core Next (GCN) AMD تمرکز خواهیم کرد.

ISA های GPU NVIDIA

GPU های NVIDIA دو سطح ISA را پشتیبانی می‌کنند: PTX (Parallel Thread Execution) و SASS (Streaming ASSembler). PTX یک ISA مجازی است که یک هدف پایدار برای کامپایلرهای CUDA فراهم می‌کند، در حالی که SASS ISA بومی GPU های NVIDIA است.

PTX (Parallel Thread Execution)

PTX یک ISA سطح پایین و مجازی است که برای GPU های NVIDIA طراحی شده است. این مشابه LLVM IR یا جاوا بایت‌کد است، به این معنی که یک هدف پایدار و مستقل از معماری را برای کامپایلرها فراهم می‌کند. برنامه‌های CUDA معمولاً به کد PTX کامپایل می‌شوند که سپس توسط درایور GPU NVIDIA به دستورات SASS بومی ترجمه می‌شوند.

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

Here's an example of PTX code for a simple vector addition kernel:

.version 7.0
.target sm_70
.address_size 64

.visible .entry vecAdd(
    .param .u64 vecAdd_param_0,
    .param .u64 vecAdd_param_1,
    .param .u64 vecAdd_param_2,
    .param .u32 vecAdd_param_3
)
{
    .reg .b32 %r<4>;
    .reg .b64 %rd<6>;

    ld.param.u64 %rd1, [vecAdd_param_0];
    ld.param.u64 %rd2, [vecAdd_param_1];
    ld.param.u64 %rd3, [vecAdd_param_2];
    ld.param.u32 %r1, [vecAdd_param_3];
    cvta.to.global.u64 %rd4, %rd1;
    cvta
}
```Here is the Persian translation of the provided markdown file, with the code comments translated, while the code itself remains untranslated:

.to.global.u64 %rd5, %rd2; mov.u32 %r2, %tid.x; mul.wide.u32 %rd6, %r2, 4; add.s64 %rd7, %rd4, %rd6; add.s64 %rd8, %rd5, %rd6;

ld.global.f32 %f1, [%rd7]; ld.global.f32 %f2, [%rd8]; add.f32 %f3, %f1, %f2;

cvta.to.global.u64 %rd9, %rd3; add.s64 %rd10, %rd9, %rd6; st.global.f32 [%rd10], %f3;

ret; }


ترجمه فارسی:

این کد PTXیک تابع هسته `vecAdd` را تعریف میکند که چهار پارامتر را دریافت میکند: اشاره‌گرها به برداری‌های ورودی و خروجی، و اندازه بردارها. این هسته شناسه رشته جهانی را محاسبه میکند، عناصر مربوطه را از بردارهای ورودی بارگیری میکند، جمع را انجام میدهد و نتیجه را در بردار خروجی ذخیره میکند.

#### SASS (Streaming ASSembler)

SASS زبان اسمبلی پایین‌دست NVIDIA GPUها است. این یک ISA سطح پایین و ماشین‌خاص است که مستقیماً به سخت‌افزار GPU نگاشت میشود. دستورات SASS توسط درایور GPU NVIDIA از کد PTX تولید میشوند و معمولاً به برنامه‌نویسان قابل مشاهده نیستند.

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

الار ALU برای اجرای کارآمد عملیات اسکالر و کنترل جریان.
- یک ALU بردار برای اجرای موازی عملیات موازی داده.
- یک سیستم حافظه پرضرابه با پشتیبانی از عملیات اتمی و دسترسی با تاخیر کم به حافظه مشترک.
- یک حالت آدرس دهی انعطاف پذیر برای عملیات حافظه، با پشتیبانی از آدرس دهی پایه+جابجایی و اسکالر+بردار.

اینجا نمونه ای از کد ISA GCN برای یک هسته افزودن بردار است:

```asm
.text
.globl vecAdd
.p2align 2

.type vecAdd,@function
vecAdd:
    .set DPTR, 0

    # بارگذاری آرگومان های هسته از حافظه
    s_load_dwordx4 s[0:3], s[4:5], 0x0
    s_load_dword s4, s[4:5], 0x10
    s_waitcnt lgkmcnt(0)

    # محاسبه آدرس های بردار ورودی
    v_lshlrev_b32 v0, 2, v0
    v_add_u32 v1, vcc, s1, v0
    v_mov_b32 v3, s3
    v_addc_u32 v2, vcc, s2, v3, vcc
    flat_load_dword v1, v[1:2]

    # محاسبه آدرس های بردار خروجی
    v_add_u32 v3, vcc, s0, v0
    v_mov_b32 v5, s3
    v_addc_u32 v4, vcc, s2, v5, vcc
    flat_load_dword v0, v[3:4]

    # اجرای افزودن بردار
    v_add_f32 v0, v0, v1
    flat_store_dword v[3:4], v0
    s_endpgm

این کد GCN، عناصر بردار ورودی را با flat_load_dword بارگذاری می کند، افزودن را با v_add_f32 انجام می دهد و نتیجه را با flat_store_dword به حافظه ذخیره می کند. دستورات s_load_dwordx4 و s_load_dword برای بارگذاری آرگومان های هسته از حافظه استفاده می شوند.

نگاشت الگوریتم ها به معماری GPU

نگاشت کارآمد الگوریتم ها به معماری GPU برای دستیابی به عملکرد بالا حیاتی است. مالحظات کلیدی شامل موارد زیر است:

نمایش کافی موازی سازی

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

کاهش واگرایی شاخه

جریان کنترل واگرا در داخل یک وروپ/موج می تواند به سریالی شدن و کاهش کارایی SIMD منجر شود. الگوریتم ها باید به گونه ای طراحی شوند که واگرایی شاخه تا حد ممکن کاهش یابد. این می تواند با کاهش استفاده از جریان کنترل وابسته به داده به دست آید.اکسپلویتینگ مموری هیراریی

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

توازن محاسبات و دسترسی‌های حافظه

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

کمینه سازی انتقال داده میزبان-دستگاه

انتقال داده بین حافظه میزبان (CPU) و دستگاه (GPU) کند است. الگوریتم‌ها باید چنین انتقالاتی را به حداقل برسانند با انجام بیشترین محاسبات ممکن بر روی GPU. داده باید به صورت دسته‌های بزرگ به GPU منتقل شده و تا زمان مورد نیاز بر روی دستگاه باقی بماند تا هزینه انتقال را کاهش دهد.

الگوهای طراحی الگوریتم موازی متداول هنگام توسعه هسته‌های GPU عبارتند از:

نقشه: هر رشته عملیات یکسانی را بر روی یک عنصر داده متفاوت انجام می‌دهد، که پردازش موازی ساده ی دانه‌های بزرگ داده را فراهم می‌کند.
کاهش: کاهش موازی برای محاسبه کارآمد یک مقدار واحد (مانند مجموع، بیشینه) از یک مجموعه داده ورودی بزرگ استفاده می‌شود. رشته‌ها کاهش‌های محلی را انجام می‌دهند که سپس ترکیب می‌شوند تا نتیجه نهایی به دست آید.
اسکن: که به عنوان مجموع پیشوند نیز شناخته می‌شود، برای محاسبه مجموع جاری عناصر در یک آرایه استفاده می‌شود. الگوریتم‌های کارآمد اسکن موازی، بلوک‌های سازنده کلیدی بسیاری از برنامه‌های شتاب‌یافته توسط GPU هستند.
استنسیل: هر رشته یک مقدار را بر اساس عناصر داده همسایه محاسبه می‌کند. محاسبات استنسیل در شبیه‌سازی‌های علمی و پردازش تصویر رایج هستند.اینجا ترجمه فارسی فایل مارکداون است:
گردآوری/پخش: رشته‌ها از مکان‌های임의 در حافظه جهانی می‌خوانند (گردآوری) یا می‌نویسند (پخش). طرح‌بندی داده و الگوهای دسترسی دقیق برای کارآیی مورد نیاز است.

شکل 3.20 مثالی از الگوی نقشه را نشان می‌دهد، جایی که هر رشته یک تابع (مانند ریشه مربع) را بر روی یک عنصر متفاوت از آرایه ورودی اعمال می‌کند.

آرایه ورودی:  
               |  |   |   |   |   |   |   |
               v  v   v   v   v   v   v   v
              ______________________________
رشته‌ها:     |    |    |    |    |    |    |    |
             |____|____|____|____|____|____|____|
                |    |    |    |    |    |    |
                v    v    v    v    v    v    v
آرایه خروجی:

شکل 3.20: مثالی از الگوی نقشه در برنامه‌نویسی GPU.

نتیجه‌گیری

مدل‌های برنامه‌نویسی GPU مانند CUDA و OpenCL قابلیت‌های پردازش موازی مدرن GPU را به توسعه‌دهندگان نمایش می‌دهند و امکان شتاب‌بخشیدن به طیف وسیعی از کاربردها را فراهم می‌کنند. این مدل‌های برنامه‌نویسی انتزاعاتی ارائه می‌دهند که به‌طور موثر بار کاری موازی ریزدانه‌ای را به سخت‌افزار GPU نگاشت می‌کنند.

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

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

Chapter 1 Introduction to Gpu Chip Design Chapter 3 Parallel Programming Models