چگونه چیپ GPU طراحی کنیم
Chapter 3 Parallel Programming Models

فصل 3: مدل‌های برنامه‌نویسی موازی در طراحی GPU

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

در این فصل، به بررسی مفاهیم و اصول کلیدی مدل‌های برنامه‌نویسی موازی برای GPU، با تمرکز بر مدل اجرای SIMT (یک دستورالعمل، چند رشته)، مدل برنامه‌نویسی CUDA و API‌های آن و چارچوب OpenCL خواهیم پرداخت. همچنین تکنیک‌های نگاشت الگوریتم‌ها به معماری GPU برای دستیابی به عملکرد و کارآیی بالا را بررسی خواهیم کرد.

مدل اجرای SIMT (یک دستورالعمل، چند رشته)

مدل اجرای SIMT پارادایم بنیادی مورد استفاده توسط GPU‌های مدرن برای دستیابی به موازی‌سازی گسترده است. در مدل SIMT، تعداد زیادی رشته به طور موازی همان برنامه (که یک کرنل نامیده می‌شود) را اجرا می‌کنند، اما هر رشته دارای شمارنده برنامه (program counter) مختص به خود است و می‌تواند مسیرهای اجرای متفاوتی را بر اساس شناسه رشته (thread ID) و داده‌هایی که بر روی آن‌ها کار می‌کند انتخاب کند.

کرنل‌ها و سلسله مراتب رشته

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

  • یک شبکه (grid) فضای مسئله کل را نمایش می‌دهد و از یک یا چند بلوک تشکیل شده است.

  • یک بلوک گروهی از رشته‌هایی است که می‌توانند با یکدیگر همکاری و همگام‌سازی کنند

  • از طریق حافظه اشتراکی و سد‌های (barriers) همگام‌سازی. رشته‌های داخل یک بلوک در هسته GPU مشابه (که چندپردازنده جریانی (streaming multiprocessor) نامیده می‌شود) اجرا می‌شوند.اینجا ترجمه فارسی برای فایل مارک‌داون ارائه شده است. برای کد، فقط توضیحات را ترجمه کرده‌ایم و خود کد را ترجمه نکرده‌ایم.

  • هر رشته (thread) دارای یک شناسه (ID) منحصربه‌فرد در بلوک و شبکه (grid) خود است که می‌تواند برای محاسبه آدرس‌های حافظه و تصمیم‌گیری در جریان کنترل استفاده شود.

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

شکل ۳.۱ سلسله‌مراتب رشته در مدل اجرای SIMT را نشان می‌دهد.

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

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

اجرای SIMT

در مدل اجرای SIMT، هر رشته همان دستور را اجرا می‌کند اما بر روی داده‌های متفاوت عمل می‌کند. با این حال، برخلاف SIMD (دستور واحد، داده‌های چندگانه) که در آن همه عناصر پردازشی به طور همزمان اجرا می‌شوند، SIMT به رشته‌ها اجازه می‌دهد تا مسیرهای اجرای مستقل داشته باشند و در دستورات انشعاب (branch) از هم جدا شوند.

وقتی یک وارپ (یک گروه از ۳۲ رشته در GPU های NVIDIA یا ۶۴ رشته در GPU های AMD) با یک دستور انشعاب مواجه می‌شود، سخت‌افزار GPU شرط انشعاب را برای هر رشته در وارپ ارزیابی می‌کند. اگر همه رشته‌ها همان مسیر را انتخاب کنند (همگرا شده‌اند)، وارپ به طور عادی به اجرای خود ادامه می‌دهد. اما اگر برخی رشته‌ها مسیرهای متفاوتی را انتخاب کنند (واگرا شده‌اند)، وارپ به دو یا چند زیروارپ تقسیم می‌شود، که هر کدام مسیر متفاوتی را دنبال می‌کنند. سخت‌افزار GPU اجرای مسیرهای واگرا را به صورت سریالی انجام می‌دهد و رشته‌های غیرفعال را در هر زیروارپ مخفی می‌کند. وقتی همه مسیرها به اتمام برسند، زیروارپ‌ها دوباره همگرا شده و به طور همزمان به اجرا ادامه می‌دهند.

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

         وارپ
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | انشعاب |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
```Reconvergence

شکل 3.2: اجرای SIMT با جریان کنترل مختلف.

این مکانیزم مدیریت مبدأ پراکندگی به SIMT امکان پشتیبانی از جریان کنترل انعطاف پذیرتر نسبت به SIMD را می دهد، اما با هزینه کاهش کارایی SIMD هنگام بروز مبدأ پراکندگی همراه است. برنامه نویسان باید به دنبال به حداقل رساندن مبدأ پراکندگی در یک وارپ برای دستیابی به عملکرد بهینه باشند.

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

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

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

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

شکل 3.3 سلسله مراتب حافظه GPU را نشان می دهد.


|            |
|   اشتراکی   |
|   حافظه   |
 ____________
       |
 ____________ 
|            |
|   محلی    |
|   حافظه   |
 ____________

شکل 3.3: سلسله مراتب حافظه GPU.

مدل برنامه نویسی CUDA و API ها

CUDA (معماری یکپارچه محاسبات دستگاه) یک پلتفرم محاسبات موازی و مدل برنامه نویسی است که توسط NVIDIA برای محاسبات عمومی بر روی GPU ها توسعه داده شده است. CUDA مجموعه ای از افزونه ها را به زبان های برنامه نویسی استاندارد مانند C، C++ و Fortran ارائه می دهد که به برنامه نویسان امکان بیان موازی سازی و بهره مندی از قدرت محاسباتی GPU های NVIDIA را می دهد.

مدل برنامه نویسی CUDA

مدل برنامه نویسی CUDA بر مفهوم هسته ها (کرنل ها) متکی است، که توابعی هستند که به طور موازی توسط تعداد زیادی از رشته ها (Thread) بر روی GPU اجرا می شوند. برنامه نویس تعداد رشته ها را که باید اجرا شوند و سازماندهی آنها به شکل شبکه بلوک های رشته را مشخص می کند.

CUDA چندین انتزاع کلیدی را برای تسهیل برنامه نویسی موازی معرفی می کند:

  • رشته: واحد اجرای پایه در CUDA. هر رشته دارای شمارنده برنامه، رجیسترها و حافظه محلی خود است.
  • بلوک: گروهی از رشته ها که می توانند با یکدیگر همکاری و همگام سازی کنند. رشته ها در درون یک بلوک بر روی یک Streaming Multiprocessor واحد اجرا می شوند و می توانند از طریق حافظه اشتراکی با هم ارتباط برقرار کنند.
  • شبکه: مجموعه ای از بلوک های رشته که یک هسته را اجرا می کنند. شبکه فضای مسئله کامل را نمایش می دهد و می تواند یک، دو یا سه بعدی باشد.

CUDA همچنین متغیرهای داخلی (مانند threadIdx، blockIdx، blockDim، gridDim) را ارائه می دهد که به رشته ها امکان شناسایی خود و محاسبه آدرس های حافظه بر اساس موقعیت آنها در سلسله مراتب رشته را می دهد.

شکل 3.4 مدل برنامه نویسی CUDA را نشان می دهد.

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

شکل 3.4: مدل برنامه نویسی CUDA.

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

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

برنامه نویسان می توانند داده ها را بین حافظه میزبان (CPU) و حافظه دستگاه (GPU) با استفاده از API های زمان اجرای CUDA مانند cudaMalloc، cudaMemcpy و cudaFree تخصیص و انتقال دهند.

شکل 3.5 سلسله مراتب حافظه CUDA را نشان می دهد.

      ____________
     |            |
     |   Global   |
     |   Memory   |
      ____________
           |
      ____________
     |            |
     |  Constant  |
     |   Memory   |
      ____________
           |
      ____________
     |            |
     |  Texture   |
     |   Memory   |
      ____________
           |
           |
      ____________
     |            |
     |   Shared   |
     |   Memory   |
      ____________
           |
      ____________ 
     |            |
     |   Local    |
     |   Memory   |
      ____________

شکل 3.5: سلسله مراتب حافظه CUDA.

همگام سازی و هماهنگی CUDA

CUDA امکانات همگام سازی و هماهنگی را فراهم می کند تا همکاری و ارتباط بین رشته ها را فعال کند:

  • همگام سازی سد: __syncthreadsاینجا ترجمه فارسی آن است:

s() تابع به عنوان یک مانع عمل می‌کند که تضمین می‌کند همه رشته‌های موجود در یک بلوک به همان نقطه رسیده‌اند قبل از پیشروی.

  • عملیات اتمی: CUDA از عملیات اتمی (مانند atomicAdd، atomicExch) پشتیبانی می‌کند که به رشته‌ها اجازه می‌دهد عملیات خواندن-اصلاح-نوشتن را بر روی حافظه مشترک یا جهانی بدون مداخله سایر رشته‌ها انجام دهند.
  • امکانات سطح وارپ: CUDA از آرگومان‌های درونی سطح وارپ (مانند __shfl، __ballot) که ارتباطات مؤثر و همگام‌سازی در یک وارپ را فراهم می‌کند، پشتیبانی می‌کند.

استفاده صحیح از همگام‌سازی و هماهنگی کننده‌های اصلی برای نوشتن برنامه‌های موازی صحیح و کارآمد در CUDA ضروری است.

مثال 3.1 یک هسته CUDA ساده را نشان می‌دهد که جمع بردار را انجام می‌دهد.

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    // شماره تردی را محاسبه کن
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    // اگر این تردی در محدوده دادگان قرار دارد
    if (i < n) {
        // مقدار نتیجه را محاسبه کن
        c[i] = a[i] + b[i];
    }
}
 
int main() {
    int *a, *b, *c;
    int n = 1024;
    
    // حافظه را در میزبان تخصیص بده
    a = (int*)malloc(n * sizeof(int));
    b = (int*)malloc(n * sizeof(int));
    c = (int*)malloc(n * sizeof(int));
    
    // بردارهای ورودی را مقداردهی اولیه کن
    for (int i = 0; i < n; i++) {
        a[i] = i;
        b[i] = i * 2;
    }
    
    // حافظه را در دستگاه تخصیص بده
    int *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, n * sizeof(int));
    cudaMalloc(&d_b, n * sizeof(int));
    cudaMalloc(&d_c, n * sizeof(int));
    
    // بردارهای ورودی را از میزبان به دستگاه کپی کن
    cudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice);
    
    // کرنل را اجرا کن
    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;
    vectorAdd<<<numBlocks,blockSize>>>(d_a, d_b, d_c, n);
    
    // بردار نتیجه را از دستگاه به میزبان کپی کن
    cudaMemcpy(c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);
    
    // حافظه دستگاه را آزاد کن
    cudaFree(d_a);
    cudaFree(d_b); 
    cudaFree(d_c);
    
    // حافظه میزبان را آزاد کن
    free(a); 
    free(b);
    free(c);
    
    بازگشتHere is the Persian translation of the provided markdown file, with the code comments translated:
 
این کد CUDA کرنل `vectorAdd` را با `numBlocks` بلوک و `blockSize` رشته در هر بلوک اجرا می‌کند. کرنل عملیات جمع‌عنصری بردارهای ورودی `a` و `b` را انجام می‌دهد و نتیجه را در بردار `c` ذخیره می‌کند. نمادگذاری `<<<...>>>` برای مشخص کردن ابعاد شبکه و بلوک هنگام اجرای کرنل استفاده می‌شود.
 
### جریان‌های CUDA و رویدادها
 
جریان‌های CUDA و رویدادها یک مکانیزم برای اجرای همزمان و همگام‌سازی کرنل‌ها و عملیات‌های ذخیره‌سازی فراهم می‌کنند:
 
- جریان‌ها: یک دنباله از عملیات (اجرای کرنل، کپی حافظه) که به ترتیب اجرا می‌شوند. جریان‌های متفاوت می‌توانند به صورت همزمان اجرا شوند، که این امکان همپوشانی محاسبات و انتقالات حافظه را فراهم می‌کند.
- رویدادها: نشانگرهایی که می‌توانند در یک جریان قرار داده شوند تا تکمیل عملیات‌های خاص را ثبت کنند. رویدادها می‌توانند برای همگام‌سازی و اندازه‌گیری زمان استفاده شوند.
 
جریان‌ها و رویدادها به برنامه‌نویسان امکان می‌دهند تا عملکرد برنامه‌های CUDA خود را با همپوشانی محاسبات و انتقالات حافظه و استفاده از تمام قابلیت‌های سخت‌افزار GPU بهینه کنند.
 
مثال ۳.۲ استفاده از جریان‌های CUDA برای همپوشانی اجرای کرنل و انتقالات حافظه را نشان می‌دهد.
 
```c
// ایجاد دو جریان
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
 
// کپی داده‌های ورودی به دستگاه به صورت غیرمسدود
cudaMemcpyAsync(d_a, a, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, b, size, cudaMemcpyHostToDevice, stream2);
 
// اجرای کرنل‌ها در جریان‌های متفاوت
kernelA<<<blocks, threads, 0, stream1>>>(d_a);
kernelB<<<blocks, threads, 0, stream2>>>(d_b);
 
// کپی نتایج به میزبان به صورت غیرمسدود
cudaMemcpyAsync(a, d_a, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(b, d_b, size, cudaMemcpyDeviceToHost, stream2);
 
// همگام‌سازی جریان‌ها
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

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

چارچوب OpenCL

OpenCL (Open Computing Language) یک استاندارد باز و رایگان برای برنامه‌نویسی موازی در سکوهای غیریکنواخت، از جمله CPU ها، GPU ها، FPGA ها و دیگر شتاب‌دهنده‌ها است. OpenCL یک مدل برنامه‌نویسی یکپارچه و مجموعه‌ای از API ها را فراهم می‌کند که به توسعه‌دهندگان امکان می‌دهد کد موازی قابل حمل و کارآمد بنویسند.

مدل برنامه‌نویسی OpenCL

مدل برنامه‌نویسی OpenCL شبیه به CUDA است، با چند تفاوت کلیدی در اصطلاحات و انتزاعات:

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

OpenCL همچنین یک مدل حافظه سلسله‌مراتبی مشابه با CUDA تعریف می‌کند:

  • حافظه جهانی: قابل دسترسی برای همه کار-آیتم‌ها در همه کار-گروه‌ها، مشابه با حافظه جهانی در CUDA.
  • حافظه محلی: به اشتراک گذاشته شده توسط همه کار-آیتم‌ها در یک کار-گروه، مشابه با حافظه اشتراکی در CUDA.
  • حافظه خصوصی: خصوصی به یک کار-آیتم واحد، مشابه با رجیسترها در CUDA.
  • حافظه ثابت: حافظه فقط-خواندنی قابل دسترسی برای همه کار-آیتم‌ها.

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

مثال 3.3 یک Kernel OpenCL را نشان می‌دهد که جمع‌وجوی برداری را انجام می‌دهد، شبیه به مثال CUDA در مثال 3.1.

__kernel void vectorAdd(__global const int *a, __global const int *b, __global int *c, int n) {
    int i = get_global_id(0);
    if (i < n)
        c[i] = a[i] + b[i];
}
  • این تابع Kernel یک عملیات جمع‌وجوی برداری را انجام می‌دهد.
  • پارامترهای ورودی شامل دو بردار a و b و خروجی بردار c است.
  • اندازه بردارها توسط پارامتر n مشخص می‌شود.
  • هر کار-آیتم یک عنصر از بردارها را جمع می‌کند.
  • شرط if (i < n) برای اطمینان از اینکه کار-آیتم‌های بیرون از محدوده بردار کاری نکنند.Here is the Persian translation of the provided markdown file, with the code comments translated:
__kernel void vector_add(
    __global const int *a,
    __global const int *b,
    __global int *c,
    int n) {
    int i = get_global_id(0);
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

ترجمه فارسی:

__kernel void vector_add(
    __global const int *a,
    __global const int *b,
    __global int *c,
    int n) {
    // شاخص جهانی کار فعلی را برمی‌گرداند
    int i = get_global_id(0);
    // اگر شاخص کوچکتر از n باشد
    if (i < n) {
        // مقدار c[i] را به مجموع a[i] و b[i] تخصیص می‌دهد
        c[i] = a[i] + b[i];
    }
}

ترجمه کامنت‌ها:

  • __kernel کلیدواژه‌ای است که تابع کرنل OpenCL را تعریف می‌کند.
  • __global کلیدواژه‌ای است که مشخص می‌کند یک پوینتر به حافظه جهانی اشاره می‌کند.
  • get_global_id تابعی است که شاخص جهانی کار فعلی را برمی‌گرداند، که برای محاسبه آدرس های حافظه برای بردارهای ورودی و خروجی استفاده می‌شود.

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

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

  • افشای موازی‌سازی کافی: الگوریتم باید به تعداد زیادی از واحدهای کاری ریزدانه تجزیه شود که به طور موازی اجرا شوند تا از قابلیت‌های پردازش موازی GPU به طور کامل استفاده شود.
  • به حداقل رساندن واگرایی شاخه: جریان کنترل واگرا در یک وارپ/موج می‌تواند به سلسله‌مراتب و کاهش کارایی SIMD منجر شود. الگوریتم‌ها باید به گونه‌ای ساختار یابند که واگرایی شاخه تا حد امکان به حداقل برسد.
  • بهره‌گیری از سلسله‌مراتب حافظه: دسترسی به حافظه جهانی گران است. الگوریتم‌ها باید از حافظه اشتراکی و ثبات‌ها به حداکثر استفاده کنند تا دسترسی‌های به حافظه جهانی را کاهش دهند. داده‌ها همچنین باید به گونه‌ای در حافظه چیده شوند که دسترسی‌های ترکیبی را امکان‌پذیر کنند.
  • تعادل بین محاسبات و دسترسی‌های حافظه: الگوریتم‌ها باید نسبت بالایی از عملیات محاسباتی به عملیات حافظه داشته باشند تا پنهان‌سازی تاخیر حافظه را امکان‌پذیر کرده و عملکرد محاسباتی بالایی را محقق سازند.
  • به حداقل رساندن انتقال داده‌ها بین میزبان و دستگاه: انتقال داده‌ها بین حافظه میزبان و دستگاه کند است. الگوریتم‌ها باید این انتقالات را به حداقل برسانند و تا حد امکان محاسبات را روی GPU انجام دهند.

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

  • نگاشت: هر کار واحد همان عملیات را روی یک داده متفاوت انجام می‌دهد، که پردازش موازی ساده داده‌های بزرگ را امکان‌پذیر می‌سازد.

  • کاهش: کاهش موازی برای محاسبه کارآمد یک مقدار واحد (مانند مجموع، حداکثر) از یک مجموعه داده بزرگ استفاده می‌شود.رشته‌ها محاسبات کاهشی محلی را انجام می‌دهند، که سپس ترکیب می‌شوند تا نتیجه نهایی را تولید کنند.

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

  • استنسیل: هر رشته یک مقدار را بر اساس داده‌های همسایه محاسبه می‌کند. محاسبات استنسیل در شبیه‌سازی‌های علمی و برنامه‌های پردازش تصویر رایج هستند.

  • جمع‌آوری/پراکندن: رشته‌ها از محل‌های دلخواه در حافظه جهانی می‌خوانند (جمع‌آوری) یا می‌نویسند (پراکندن). چیدمان داده و الگوهای دسترسی دقیق برای کارآمدی مورد نیاز است.

نتیجه‌گیری

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

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

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

Chapter 2 Gpu Rogramming ModelsChapter 4 Gpu Memory System Design