فصل 11: جهت‌گیری‌های تحقیقاتی GPU در اسکالاریزاسیون و اجرای خطی

همانطور که در فصل 2 توضیح داده شد، API های محاسبات GPU مانند CUDA و OpenCL، یک مدل برنامه‌نویسی شبیه MIMD را ارائه می‌دهند که به برنامه‌نویس امکان می‌دهد یک آرایه بزرگ از رشته‌های اسکالر را روی GPU راه‌اندازی کند. در حالی که هر یک از این رشته‌های اسکالر می‌توانند مسیر اجرای منحصر به فرد خود را دنبال کنند و به مکان‌های حافظه دلخواه دسترسی داشته باشند، در مورد معمول، همه آنها یک مجموعه کوچک از مسیرهای اجرا را دنبال می‌کنند و عملیات‌های مشابهی را انجام می‌دهند.

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

بینش کلیدی این تحقیقات در مشاهده ساختار ارزش [Kim et al., 2013] در میان رشته‌هایی که هسته محاسباتی یکسانی را اجرا می‌کنند، نهفته است. دو نوع ساختار ارزش، یکنواخت و خطی، در هسته محاسباتی در مثال 11.1 نشان داده شده است.

متغیر یکنواخت متغیری که مقدار ثابت یکسانی را برای هر رشته در هسته محاسباتی دارد. در الگوریتم 11.1، متغیر a و همچنین ثوابت THRESHOLD و Y_MAX_VALUE همگی دارای مقادیر یکنواخت در میان تمام رشته‌های در هسته محاسباتی هستند. یک متغیر یکنواخت می‌تواند در یک رجیستر اسکالر ذخیره شود و توسط همه رشته‌ها در هسته محاسباتی استفاده شود.

متغیر خطی متغیری که مقادیر آن یک تابع خطی از شناسه رشته برای هر رشته در هسته محاسباتی است. در الگوریتم 11.1، آدرس حافظه متغیر y[idx] را می‌توان به صورت یک تبدیل خطی از شناسه رشته threadIdx.x نمایش داد:

&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;

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

__global__ void vsadd( int y[], int a ) {
    // کد اصلی
}
```اینجا ترجمه فارسی فایل مارک‌داون است:
 
```c
int idx = threadIdx.x;
y[idx] = y[idx] + a;
if ( y[idx] > THRESHOLD )
    y[idx] = Y_MAX_VALUE;
}

الگوریتم ۱۱.۱: مثالی از عملیات اسکالر و افین در یک هسته محاسباتی (از [Kim et al., 2013]).

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

تشخیص متغیرهای یکنواخت یا افین

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

تشخیص محرک کامپایلر

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

AMD GCN [AMD, 2012] بر روی کامپایلر متکی است تا متغیرهای یکنواخت و عملیات‌های اسکالر را که می‌توانند توسط یک پردازنده اسکالر ذخیره و پردازش شوند، تشخیص دهد.

Asanovic et al. [2013] تحلیل همگرا و متغیر ترکیبی را معرفی می‌کنند که به کامپایلر امکان می‌دهد تا عملیات‌هایی در یک هسته محاسباتی دلخواه را که واجد شرایط اسکالریزاسیون و/یا تبدیل افین هستند، تعیین کند. دستورالعمل‌هایی در مناطق همگرای یک هسته محاسباتی می‌توانند به دستورالعمل‌های اسکالر/افین تبدیل شوند. در هر گذار از مناطق واگرا به مناطق همگرای یک هسته محاسباتی، کامپایلر یک دستورالعمل syncwarp را برای مدیریت وابستگی‌های ثبت بین دو منطقه وارد می‌کند. Asanovic et al. [2013] از این رویکرد استفاده کردند.اینجا ترجمه فارسی فایل مارک‌داون است. برای کد، فقط توضیحات را ترجمه کنید، کد را ترجمه نکنید:

این تحلیل برای تولید عملیات اسکالر برای معماری Temporal-SIMT [Keckler et al., 2011, Krashinsky, 2011] استفاده می‌شود.

محاسبات آفین جدا‌شده (DAC) [Wang and Lin, 2017] بر یک تحلیل کامپایلر مشابه تکیه دارد تا کاندیداهای اسکالر و آفین را استخراج کند که در یک وارپ جداگانه قرار داده می‌شوند. Wang and Lin [2017] فرآیند را با یک تحلیل آفین متفاوت تقویت می‌کنند، با هدف استخراج رشته‌های دستورالعمل