فصل ۷: طراحی پردازشگر چند ریشه شونده در طراحی GPU

پردازشگر چند ریشه شونده (SM) بلوک سازنده اصلی معماری GPU های NVIDIA است. هر SM شامل مجموعه ای از هسته های CUDA است که دستورات را به شیوه SIMT (دستور واحد، چندین جریان) اجرا می کنند. SM مسئول مدیریت و زمانبندی وارپ ها، مدیریت انشعاب شاخه، و ارائه دسترسی سریع به حافظه مشترک و کش ها است. در این فصل، ما به بررسی میکروآرشیتکچر SM، از جمله خطوط لوله آن، مکانیزم های زمانبندی وارپ، طراحی فایل ثبت، و سازماندهی حافظه مشترک و L1 کش خواهیم پرداخت.

میکروآرشیتکچر SM و خطوط لوله

SM یک پردازنده با موازی بالا و خط لوله است که طراحی شده است تا به طور کارآمد صدها جریان را همزمان اجرا کند. شکل ۷.۱ یک نمودار بلوک ساده از یک SM در معماری NVIDIA Volta را نشان می دهد.

                                 Cache دستورالعمل
                                         |
                                         v
                                    برنامه ریز وارپ
                                         |
                                         v
                               واحد ارسال (۴ وارپ)
                                 |   |   |   |
                                 v   v   v   v
                               هسته CUDA (FP64/FP32/INT)
                               هسته CUDA (FP64/FP32/INT)
                               هسته CUDA (FP64/FP32/INT)
                               ...
                               هسته تانسور
                               هسته تانسور
                               ...
                               واحد بارگذاری/ذخیره
                               واحد بارگذاری/ذخیره
                               ...
                               واحد تابع ویژه
                                         ^
                                         |
                                فایل ثبت (۶۴ کیلوبایت)
                                         ^اینجا ترجمه فارسی فایل مارک‌داون است. برای بخش‌های کد، تنها نظرات ترجمه شده‌اند و خود کد تغییر نکرده است:

                                  حافظه اشتراکی / کش L1 (96 کیلوبایت)

شکل 7.1: طرح بلوک ساده‌شده‌ی یک SM در معماری NVIDIA Volta.

اجزای اصلی SM شامل:

1. کش دستورالعمل: ذخیره دستورالعمل‌های پرکاربرد برای کاهش تأخیر و افزایش تراوش.

2. زمان‌بندکننده وارپ: وارپ‌هایی که آماده اجرا هستند را انتخاب و به واحدهای اجرایی در دسترس ارسال می‌کند.

3. واحد ارسال: دستورالعمل‌ها را برای تا 4 وارپ در هر سیکل بازیابی و رمزگشایی می‌کند و آنها را به واحدهای اجرایی مناسب ارسال می‌کند.

4. هسته‌های CUDA: واحدهای اجرایی قابل برنامه‌ریزی که طیف وسیعی از عملیات عددصحیح و اعشاری را پشتیبانی می‌کنند. هر SM در Volta دارای 64 هسته CUDA است.

5. هسته‌های تنسور: واحدهای اجرایی تخصصی طراحی‌شده برای شتاب‌دهی به محاسبات یادگیری عمیق و کاربردهای هوش مصنوعی. هر SM در Volta دارای 8 هسته تنسور است.

6. واحدهای بارگذاری/ذخیره‌سازی: مسئول عملیات‌های حافظه، از جمله بارگذاری و ذخیره‌سازی در حافظه جهانی، حافظه اشتراکی و کش‌ها.

7. واحدهای عملیات ویژه: اجرای عملیات‌های ریاضی پیچیده و تبدیلی.

8. فایل ثبت: فراهم‌کننده دسترسی سریع به ثبات‌های خصوصی رشته. هر SM در Volta دارای یک فایل ثبت 64 کیلوبایتی است.

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

خط‌لوله SM طراحی‌شده است تا تراوش را با اجازه اجرای همزمان چندین وارپ و پنهان‌سازی تأخیر حافظه به حداکثر برساند. شکل 7.2 یک نمای ساده‌شده از خط‌لوله SM را نشان می‌دهد.

    بازیابی دستورالعمل
            |
            v
    رمزگشایی دستورالعمل
            |
            v
    جمع‌آوری عملوندها
            |
            v
    اجرا (هسته‌های CUDA، هسته‌های تنسور، واحدهای بارگذاری/ذخیره‌سازی، واحدهای عملیات ویژه)
            |
            v
    بازنویسی

شکل 7.2: خط‌لوله ساده‌شده‌ی SM.

مراحل خط‌لوله به شرح زیر است:

1. بازیابی دستورالعمل: زمان‌بندکننده وارپ وارپی را که آماده اجرا است انتخاب می‌کند.متن فارسی:

2. Instruction Fetch: هر وارپ (WARP) یک دستورالعمل را از حافظه کش دستورالعمل بارگیری می‌کند و سپس دستورالعمل بعدی برای آن وارپ را از حافظه کش دستورالعمل بارگیری می‌کند.

3. Instruction Decode: دستورالعمل بارگیری شده رمزگشایی می‌شود تا نوع عملیات، عملوندها و ثبت‌های مقصد تعیین شوند.

4. Operand Collection: عملوندهای مورد نیاز برای دستورالعمل از فایل ثبت یا حافظه اشتراکی جمع‌آوری می‌شوند.

5. Execution: دستورالعمل در واحد اجرایی مناسب (هسته CUDA، هسته تنسور، واحد بارگذاری/ذخیره‌سازی یا واحد توابع ویژه) اجرا می‌شود.

6. Writeback: نتیجه اجرا در فایل ثبت یا حافظه اشتراکی ذخیره می‌شود.

برای دستیابی به عملکرد بالا، SM از چندین تکنیک استفاده می‌کند تا بهره‌وری منابع را حداکثر و تأخیر را پنهان کند:

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

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

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    // شناسه تخیه را محاسبه می‌کند
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        // مقدار عنصر متناظر از بردارهای a و b را جمع می‌کند و در بردار c ذخیره می‌کند
        c[tid] = a[tid] + b[tid];
    }
}

مثال 7.1: کرنل CUDA برای جمع بردار.

در این مثال، هر رشته در کرنل مجموع عناصر متناظر از بردارهای ورودی a و b را محاسبه کرده و نتیجه را در بردار خروجی c ذخیره می‌کند. SM این کرنل را با اختصاص دادن هر رشته به یک هسته CUDA و زمان‌بندی وارپ‌های رشته‌ها برای اجرا در هسته‌های موجود اجرا می‌کند. واحدهای بارگذاری/ذخیره‌سازی برای بارگیری داده‌های ورودی از حافظه جهانی و نوشتن نتایج استفاده می‌شوند.

زمان‌بندی وارپ و مدیریت گرایش

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

مدیریت کارآمد برنامه‌ریزی وارپ برای به‌حداکثر رساندن عملکرد واحد پردازشی (SM) بسیار حیاتی است. برنامه‌ریز وارپ مسئول انتخاب وارپ‌های آماده برای اجرا و ارسال آن‌ها به واحدهای اجرایی موجود است. هدف اصلی برنامه‌ریز وارپ نگه‌داشتن واحدهای اجرایی به‌طور مداوم مشغول است، با اطمینان از وجود همیشه وارپ‌های آماده برای اجرا.

SM از یک مکانیزم برنامه‌ریزی وارپ دو‌سطحی استفاده می‌کند:

1. برنامه‌ریزی وارپ: برنامه‌ریز وارپ وارپ‌های آماده برای اجرا را بر اساس یک سیاست برنامه‌ریزی مانند دور-به‌دور یا قدیمی‌ترین‌ها-اول انتخاب می‌کند. سپس وارپ‌های انتخاب‌شده به واحدهای اجرایی موجود ارسال می‌شوند.

2. برنامه‌ریزی دستورالعمل: در هر وارپ، SM دستورالعمل‌ها را بر اساس وابستگی‌ها و در دسترس بودن واحدهای اجرایی برنامه‌ریزی می‌کند. SM می‌تواند چندین دستورالعمل مستقل از همان وارپ را در یک سیکل صادر کند تا پارالل‌سازی سطح دستورالعمل را به‌حداکثر برساند.

شکل 7.3 مکانیزم برنامه‌ریزی وارپ دو‌سطحی را نشان می‌دهد.

    پول وارپ
    وارپ 1 (آماده)
    وارپ 2 (در انتظار)
    وارپ 3 (آماده)
    ...
    وارپ N (آماده)
        |
        v
    برنامه‌ریز وارپ
        |
        v
    واحد ارسال
        |
        v
    واحدهای اجرایی

شکل 7.3: مکانیزم برنامه‌ریزی وارپ دو‌سطحی.

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

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

مثال 7.2 یک هسته CUDA با یک شاخه واگرا را نشان می‌دهد.زبان فارسی ترجمه:

__global__ void divergentKernel(int *data, int *result) {
    // شناسه رشته کاری
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (data[tid] > 0) {
        // اگر داده در موقعیت رشته کاری مثبت است
        result[tid] = data[tid] * 2;
    } else {
        // در غیر این صورت
        result[tid] = data[tid] * 3;
    }
}

مثال 7.2: Kernel CUDA با یک شاخه واگرا.

در این مثال، شرط شاخه data[tid] > 0 ممکن است باعث شود که برخی از رشته‌های یک warps مسیر if را و بقیه مسیر else را طی کنند. SM این واگرایی را با اجرای هر دو مسیر به صورت متوالی و غیرفعال کردن رشته‌های غیرفعال در هر مسیر، مدیریت می‌کند.

شکل 7.4 فرآیند پیش‌پرداختن را برای یک warp با رشته‌های واگرا نشان می‌دهد.

    Warp (32 رشته)
    رشته 1: data[1] = 5, result[1] = 10
    رشته 2: data[2] = -3, result[2] = -9
    ...
    رشته 32: data[32] = 7, result[32] = 14

    شاخه واگرا:
    if (data[tid] > 0) {
        result[tid] = data[tid] * 2;
    } else {
        result[tid] = data[tid] * 3;
    }

    پیش‌پرداختن:
    گام 1: اجرای مسیر if با ماسک
        رشته 1: result[1] = 10
        رشته 2: (غیرفعال شده)
        ...
        رشته 32: result[32] = 14

    گام 2: اجرای مسیر else با ماسک
        رشته 1: (غیرفعال شده)
        رشته 2: result[2] = -9
        ...
        رشته 32: (غیرفعال شده)

    نتیجه نهایی:
    رشته 1: result[1] = 10
    رشته 2: result[2] = -9
    ...
    رشته 32: result[32] = 14

شکل 7.4: فرآیند پیش‌پرداختن برای یک warp با رشته‌های واگرا.

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

فایل ثبت و جمع‌کننده‌های عملوند

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

در معماری ولتا NVIDIA، هر SM دارای یک فایل ثبت 64 کیلوبایتی است که به صورت 32 بانک 2 کیلوبایتی سازماندهی شده است. فایل ثبت برای ارائه پهنای باند بالا و دسترسی با تأخیر کم به منظور پشتیبانی از تعداد زیاد رشته‌های همزمان طراحی شده است.

برای به حداقل رساندن تعارض بانک و بهبود عملکرد، SM از تکنیکی به نام جمع‌آوری عملوند استفاده می‌کند. جمع‌آوری‌کننده‌های عملوند واحدهای تخصصی هستند که عملوندها را از بانک‌های فایل ثبت جمع‌آوری کرده و به واحدهای اجرایی تحویل می‌دهند. با استفاده از جمع‌آوری‌کننده‌های عملوند، SM می‌تواند تأثیر تعارض بانک را کاهش داده و بهره‌وری واحدهای اجرایی را بهبود بخشد.

شکل 7.5 یک نمودار ساده‌شده از فایل ثبت و جمع‌آوری‌کننده‌های عملوند در یک SM را نشان می‌دهد.

    فایل ثبت (64 کیلوبایت)
    بانک 1 (2 کیلوبایت)
    بانک 2 (2 کیلوبایت)
    ...
    بانک 32 (2 کیلوبایت)
        |
        v
    جمع‌آوری‌کننده‌های عملوند
        |
        v
    واحدهای اجرایی

شکل 7.5: فایل ثبت و جمع‌آوری‌کننده‌های عملوند در یک SM.

جمع‌آوری‌کننده‌های عملوند با جمع‌آوری عملوندها از چندین دستور و چندین وارپ کار می‌کنند، به این ترتیب SM می‌تواند دستورات مختلف را از وارپ‌های مختلف در یک چرخه به واحدهای اجرایی صادر کند. این کار به پنهان‌سازی تأخیر دسترسی به فایل ثبت کمک کرده و بهره‌وری کلی SM را بهبود می‌بخشد.

مثال 7.3 یک هسته CUDA را نشان می‌دهد که یک حاصل‌ضرب نقطه‌ای دو بردار را انجام می‌دهد.

__global__ void dotProduct(float *a, float *b, float *result, int n) {
    __shared__ float partialSum[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    partialSum[tid] = 0;
 
    while (i < n) {
        partialSum[tid] += a[i] * b[i];
        i += blockDim.x * gridDim.x;
    }
 
    __syncthreads();
 
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            partialSum[tid] += partialSum[tid + s];
        }
        __syncthreads();
    }
 
    if (tid == 0) {
        result[blockIdx.x] = partialSum[0];
    }
}

در این مثال، هر رشته یک مجموع جزئی از حاصل‌ضرب نقطه‌ای را با استفاده از اندیس خود محاسبه می‌کند.Here is the Persian translation of the provided Markdown file, with the code comments translated:

عناصر از بردارهای ورودی. مجموع های جزئی در آرایه حافظه اشتراکی partialSum ذخیره می شوند. پس از اینکه همه رشته ها مجموع های جزئی خود را محاسبه کردند، یک کاهش موازی انجام می شود تا مجموع های جزئی را جمع کرده و نتیجه نهایی حاصل ضرب را به دست آورند.

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

نتیجه گیری

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

مولفه های کلیدی SM شامل زمان بندی وارپ، که وارپ ها را برای اجرا انتخاب می کند؛ پشته SIMT، که شاخه گرایی و همگرایی را مدیریت می کند؛ فایل رجیستر و جمع کننده های عملوند، که دسترسی سریع به رجیستر های خصوصی رشته را فراهم می کنند؛ و حافظه اشتراکی و حافظه نهان L1، که اشتراک گذاری و استفاده مجدد داده های با تأخیر کم را امکان پذیر می سازند.

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