الفصل 7: تصميم محول المعالجة المتعددة للبث في تصميم معالج الرسومات

محول المعالجة المتعددة للبث (SM) هو البناء الأساسي لمعمارية معالجات الرسومات NVIDIA. يحتوي كل SM على مجموعة من أنوية CUDA التي تنفذ التعليمات بطريقة SIMT (تعليمة واحدة، متعددة خيط). يكون SM مسؤولاً عن إدارة وجدولة الشرائح، ومعالجة تباين الفروع، وتوفير الوصول السريع إلى الذاكرة المشتركة والذاكرة التخبئة. في هذا الفصل، سنستكشف البنية الدقيقة للSM، بما في ذلك أنابيبها، وآليات جدولة الشريحة، وتصميم ملف السجل، وتنظيم الذاكرة المشتركة وذاكرة التخبئة L1.

بنية دقيقة للSM وأنابيب

SM هو معالج متوازٍ وأنبوبي مصمم بشكل فعال لتنفيذ مئات الخيوط في وقت واحد. يُظهر الشكل 7.1 مخططًا بسيطًا لحظريSM في معمارية NVIDIA Volta.

                                 ذاكرة التعليمات
                                         |
                                         v
                                    جدولة الشريحة
                                         |
                                         v
                               وحدة الإرسال (4 شرائح)
                                 |   |   |   |
                                 v   v   v   v
                               نواة CUDA (FP64/FP32/INT)
                               نواة CUDA (FP64/FP32/INT)
                               نواة CUDA (FP64/FP32/INT)
                               ...
                               نواة Tensor
                               نواة Tensor
                               ...
                               وحدة التحميل/التخزين
                               وحدة التحميل/التخزين
                               ...
                               وحدة الوظائف الخاصة
                                         ^
                                         |
                                ملف السجل (64 كيلوبايت)
                                         ^هنا ترجمة الملف إلى اللغة العربية:

                            الذاكرة المشتركة / ذاكرة التخزين المؤقت L1 (96 كيلوبايت)

الشكل 7.1: مخطط بلوك مبسّط لـ SM في معمارية NVIDIA Volta.

المكوّنات الرئيسية للـ SM تشمل:

1. ذاكرة التعليمات: تخزن التعليمات التي يتم الوصول إليها بشكل متكرر لتقليل التأخير وتحسين الإنتاجية.

2. جدول وحدات الحزم: يختار وحدات الحزم التي جاهزة للتنفيذ ويرسلها إلى وحدات التنفيذ المتاحة.

3. وحدة الإرسال: تجلب وتفك تعليمات لما يصل إلى 4 وحدات حزم في كل دورة وترسلها إلى وحدات التنفيذ المناسبة.

4. أنوية CUDA: وحدات تنفيذ قابلة للبرمجة تدعم مجموعة واسعة من العمليات العددية الصحيحة والعائمة. كل SM في Volta يحتوي على 64 نواة CUDA.

5. أنوية Tensor: وحدات تنفيذ متخصصة مصممة لتسريع عمليات التعلم العميق والأحمال الحسابية للذكاء الاصطناعي. كل SM في Volta يحتوي على 8 أنوية Tensor.

6. وحدات التحميل/التخزين: تتعامل مع عمليات الذاكرة، بما في ذلك التحميل والتخزين في الذاكرة العامة والذاكرة المشتركة والذاكرة المؤقتة.

7. وحدات الوظائف الخاصة: تنفّذ العمليات الرياضية المعقدة والتكرارية.

8. سجل الملفات: يوفّر وصولاً سريعًا إلى السجلات الخاصة بالخيوط. كل SM في Volta لديه سجل ملفات بحجم 64 كيلوبايت.

9. الذاكرة المشتركة / ذاكرة التخزين المؤقت L1: مساحة ذاكرة قابلة للتهيئة يمكن استخدامها كذاكرة تخزين مؤقت (الذاكرة المشتركة) أو كذاكرة تخزين مؤقت L1 مدارة بواسطة الأجهزة.

تم تصميم أنبوب أداء SM لتحقيق أقصى قدر من الإنتاجية من خلال السماح لعدة وحدات حزم بالتنفيذ في نفس الوقت وإخفاء تأخير الذاكرة. يوضح الشكل 7.2 عرضًا مبسططا لأنبوب أداء SM.

    استرداد التعليمات
            |
            v
    فك تشفير التعليمات
            |
            v
    جمع العناصر
            |
            v
    التنفيذ (أنوية CUDA، أنوية Tensor، وحدات التحميل/التخزين، وحدات الوظائف الخاصة)
            |
            v
    الكتابة المرتدة

الشكل 7.2: أنبوب أداء SM المبسط.

مراحل الأنبوب هي كما يلي:

1. استرداد التعليمات: يختار جدول وحدات الحزم وحدة حزم جاهزة للتنفيذباللغة العربية:

2. استرداد التعليمة: يُسترد التعليمة التالية لذلك المعالج من ذاكرة التعليمات.

3. فك تشفير التعليمة: يتم فك تشفير التعليمة المجلوبة لتحديد نوع العملية والمشغلات والسجلات المقصودة.

4. تجميع المشغلات: يتم تجميع المشغلات المطلوبة للتعليمة من ملف السجلات أو الذاكرة المشتركة.

5. التنفيذ: يتم تنفيذ التعليمة على وحدة التنفيذ المناسبة (نواة CUDA، نواة المصفوفات، وحدة التحميل/التخزين، أو وحدة الوظائف الخاصة).

6. الكتابة الخلفية: يتم كتابة نتيجة التنفيذ إلى ملف السجلات أو الذاكرة المشتركة.

لتحقيق أداء عالي، تستخدم SM العديد من التقنيات لتحقيق أقصى استفادة من الموارد وإخفاء التأخير:

• الإصدار المزدوج: يمكن لS Mإصدار تعليمتين مستقلتين لكل معالج في دورة واحدة، مما يسمح بزيادة التوازي على مستوى التعليمات.
• وحدات التنفيذ المؤنبرة: وحدات التنفيذ مؤنبرة، مما يمكن SM من بدء عملية جديدة على وحدة قبل إكمال العملية السابقة.
• إخفاء التأخير: يمكن لS Mالتبديل بين المعالجات على أساس دوري لكل دورة، مما يسمح له بإخفاء تأخير الوصول إلى الذاكرة والعمليات طويلة التأخير من خلال تنفيذ تعليمات من معالجات أخرى.

المثال 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) {
        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) {
        // النتيجة[المعرّف الزمني] = البيانات[المعرّف الزمني] * 2
        result[tid] = data[tid] * 2;
    } else {
        // النتيجة[المعرّف الزمني] = البيانات[المعرّف الزمني] * 3
        result[tid] = data[tid] * 3;
    }
}

مثال 7.2: نواة CUDA مع فرع متباين.

في هذا المثال، قد يؤدي شرط الفرع data[tid] > 0 إلى أن تتخذ بعض الخيوط في وكتة المسار الإذا بينما تأخذ البقية المسار الآخر. تتعامل وحدة المعالجة الآلية (SM) مع هذا التباين عن طريق تنفيذ كلا المسارين متتالياً وحجب الخيوط غير النشطة في كل مسار.

يوضح الشكل 7.4 عملية التنبؤ لوكتة مع خيوط متباينة.

    وكتة (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: تنفيذ مسار الإذا مع القناع
        الخيط 1: result[1] = 10
        الخيط 2: (محجوب)
        ...
        الخيط 32: result[32] = 14

    الخطوة 2: تنفيذ مسار الآخر مع القناع
        الخيط 1: (محجوب)
        الخيط 2: result[2] = -9
        ...
        الخيط 32: (محجوب)

    النتيجة النهائية:
    الخيط 1: result[1] = 10
    الخيط 2: result[2] = -9
    ...
    الخيط 32: result[32] = 14

الشكل 7.4: عملية التنبؤ لوكتة مع خيوط متباينة.

باستخدام التنبؤ، يمكن لوحدة المعالجة الآلية (SM) التعامل مع تباين الفرع دون الحاجة إلى تعليمات الفرع الصريحة أو تباين تدفق التحكم. ومع ذلك، لا يزال الفروع المتباينة يمكن أن تؤثر على الأداء، حيث يجب على وحدة المعالجة الآلية (SM) تنفيذ كلا المسارين متتالياً، مما يقلل من الفاعلية الموازية الفعلية.

ملف السجل والجامعات التشغيلية

ملف السجل هو مكون حرج من مكونات وحدة المعالجة الآلية (SM)، وهو يوفر وصولاً سريعًا إلى السجلات الخاصة بالخيط. لدى كل وحدة معالجة آلية (SM) ملف سجل كبير لدعم العديد من الخيوط النشطة والسماح بالتبديل الفعال بين الوكتات.بواسطة المعماري الرسومي NVIDIA Volta, يمتلك كل 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];
    }
}

في هذا المثال, يُحسب كل خيط مجموعة جزئية للحاصل النقطي باستخدام الجزء المُخصَّص لههذا هو الترجمة العربية للملف المذكور:

العناصر من متجهات الإدخال. يتم تخزين الأجزاء الجزئية في مصفوفة الذاكرة المشتركة partialSum. بعد أن يقوم جميع الخيوط بحساب الأجزاء الجزئية الخاصة بهم، يتم إجراء تخفيض متوازٍ لجمع الأجزاء الجزئية والحصول على نتيجة الناتج النقطي النهائية.

يلعب جامع العوامل دورًا حاسمًا في هذا المثال من خلال جمع العوامل بكفاءة للوصول إلى الذاكرة المشتركة والعمليات الحسابية. وهذا يساعد على تجنب تضارب البنوك وتحسين استخدام وحدات التنفيذ.

الاستنتاج

إن محول البث هو الوحدة الحسابية الأساسية في معمارية وحدات معالجة الرسومات الحديثة. يركز تصميمه على تحقيق أقصى قدر من المرور وإخفاء تأخير الذاكرة من خلال مزيج من التعدد الخيطي الدقيق، وتنفيذ SIMT، وجمع العوامل الفعال.

تشمل المكونات الرئيسية للمحول المرشح للمراكز، الذي يختار المراكز للتنفيذ؛ ومكدس SIMT، الذي يتعامل مع تباين التفرع والتقارب؛ وملف السجل وجامعي العوامل، اللذان يوفران وصولاً سريعًا إلى السجلات الخاصة بالخيط؛ والذاكرة المشتركة وذاكرة التخزين المؤقت L1، اللتان تمكّنان من مشاركة البيانات وإعادة استخدامها بتأخير منخفض.

بينما تواصل معمارية وحدات معالجة الرسومات التطور، ستكون البحوث في مجالات مثل معالجة تباين التفرع، وجدولة المراكز، وتصميم ملف السجل حاسمة لتحسين أداء وكفاءة وحدات معالجة الرسومات المستقبلية. يمكن أن تؤدي التقنيات الجديدة مثل تكوين المراكز الديناميكية، وضغط كتلة الخيط، وذاكرة التخزين المؤقت لإعادة استخدام العوامل إلى تحسين قدرات المحول بشكل كبير وتمكين مستويات جديدة من الأداء في أحمال عمل الحوسبة المتوازية.