كيفية تصميم رقائق GPU
Chapter 2 Gpu Rogramming Models

الفصل 2: نماذج برمجة وحدات معالجة الرسومات

تطورت وحدات معالجة الرسومات (GPUs) من مُسرّعات الرسومات ذات الوظائف الثابتة إلى محركات حوسبة متوازية قابلة للبرمجة، والتي يمكن أن تُسرع مجموعة واسعة من التطبيقات. لتمكين المبرمجين من استغلال الموازية الضخمة في وحدات معالجة الرسومات بفاعلية، تم تطوير العديد من نماذج البرمجة المتوازية وواجهات برمجة التطبيقات، مثل NVIDIA CUDA و OpenCL و DirectCompute. توفر هذه نماذج البرمجة تجريدات تسمح للمبرمجين بالتعبير عن الموازية في تطبيقاتهم مع إخفاء التفاصيل المنخفضة المستوى للأجهزة الخاصة بوحدة معالجة الرسومات.

في هذا الفصل، سنستكشف المفاهيم والمبادئ الرئيسية وراء نماذج برمجة الحوسبة المتوازية لوحدات معالجة الرسومات، مع التركيز على نموذج التنفيذ ومعماريات مجموعة التعليمات للوحدات المعالجة الرسومية (ISAs) لشركتي NVIDIA و AMD's Graphics Core Next (GCN). سنقدم أيضًا أمثلة لتوضيح كيفية تطبيق هذه المفاهيم في الممارسة العملية.

نموذج التنفيذ

يستند نموذج التنفيذ لنماذج برمجة وحدات معالجة الرسومات الحديثة على مفهوم الحارات (kernels)، وهي وظائف يتم تنفيذها بشكل متوازٍ بواسطة عدد كبير من الخيوط (threads) على وحدة معالجة الرسومات. عند إطلاق حارة، يحدد المبرمج عدد الخيوط المراد إنشاؤها وكيفية تنظيمها في تسلسل من الشبكات (grids) والكتل (blocks أو cooperative thread arrays - CTAs) والخيوط الفردية.

  • تمثل الشبكة المساحة المشكلة بالكامل وتتكون من كتلة واحدة أو أكثر.
  • الكتلة هي مجموعة من الخيوط التي يمكن أن تتعاون وتتزامن مع بعضها البعض من خلال الذاكرة المشتركة والحواجز. يتم تنفيذ الخيوط داخل الكتلة على نفس النواة للوحدة المعالجة الرسومية (تسمى أيضًا وحدة متعددة المعالج أو وحدة الحساب).
  • لكل خيط معرف فريد داخل كتلته والشبكة التي ينتمي إليها، والذي يمكن استخدامه لحساب عناوين الذاكرة واتخاذ قرارات التدفق التحكمي.

يسمح هذا التنظيم التسلسلي للمبرمجين بالتعبير عن الموازاة في البيانات (حيث يتم تطبيق نفس العملية على عناصر بيانات متعددة) والموازاة في المهام (حيث يتم تنفيذ مهام مختلفة بشكل متوازٍ).إليك الترجمة العربية لهذا الملف باستخدام التنسيق المطلوب:

الشكل 2.1 يوضح التسلسل الهرمي للخيوط في نموذج التنفيذ على وحدة المعالجة الرسومية.

            الشبكة
       ________________
      /   /   /   /   /
     /   /   /   /   / 
    /   /   /   /   /
   /   /   /   /   /
  /__/__/__/__/__/
    |   |   |   |
    |   | كتلة |
    |   |   |   |
  الخيط الخيط ...

الشكل 2.1: التسلسل الهرمي للخيوط في نموذج التنفيذ على وحدة المعالجة الرسومية.

تنفيذ SIMT

إن نماذج برمجة وحدات المعالجة الرسومية مثل CUDA و OpenCL تتبع نموذج التنفيذ Single-Instruction, Multiple-Thread (SIMT). في نموذج SIMT، يتم تنفيذ الخيوط في مجموعات تسمى أكوابة (NVIDIA terminology) أو موجات الواجهة (AMD terminology). تنفذ جميع الخيوط داخل الأكوابة نفس التعليمة في نفس الوقت، ولكن كل خيط يعمل على بيانات مختلفة.

ومع ذلك، على عكس نموذج Single-Instruction, Multiple-Data (SIMD) التقليدي، حيث تنفذ جميع عناصر المعالجة بالتزامن، يسمح SIMT للخيوط بامتلاك مسارات تنفيذ مستقلة والانحراف عند تعليمات الفرع. عندما تواجه الأكوابة تعليمة فرع، يقوم جهاز hardware وحدة المعالجة الرسومية بتقييم شرط الفرع لكل خيط في الأكوابة. إذا اتخذ جميع الخيوط نفس المسار (متقارب)، تستمر الأكوابة في التنفيذ بشكل طبيعي. إذا اتخذ بعض الخيوط مسارات مختلفة (متباعد)، يتم تقسيم الأكوابة إلى أكوابة فرعية اثنتين أو أكثر، تتبع كل منها مسارًا مختلفًا. يقوم جهاز hardware وحدة المعالجة الرسومية بترتيب تنفيذ المسارات المتباعدة، وإخفاء الخيوط غير النشطة في كل أكوابة فرعية. عند اكتمال جميع المسارات، تتحد الأكوابة الفرعية مرة أخرى وتستمر في التنفيذ بالتزامن.

يوضح الشكل 2.2 تنفيذ SIMT مع تدفق التحكم المتباعد.

         الأكوابة
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | الفرع |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
            \
             \
   إعادة التقارب

الشكل 2.2: تنفيذ SIMT مع تدفق التحكم المتباعد.

يسمح هذا الميكانيزم لمعالجة التباعد بدعم SIMT لتدفق التحكم الأكثر مرونةHere is the Arabic translation of the provided markdown file:

السلسلة الهرمية للذاكرة

تتمتع بطاقات الرسومات (GPUs) بسلسلة هرمية معقدة للذاكرة لدعم متطلبات الحزمة العالية والتأخير المنخفض للحمولات العمل المتوازية. وعادة ما تتكون السلسلة الهرمية للذاكرة من:

  • الذاكرة العالمية: أكبر مساحة للذاكرة ولكنها الأبطأ، والتي يمكن الوصول إليها من قبل جميع الخيوط في نواة. وعادة ما تُنفذ الذاكرة العالمية باستخدام ذاكرة GDDR أو HBM عالية النطاق الترددي.
  • الذاكرة المشتركة: مساحة ذاكرة سريعة ومدمجة في الشريحة، مشتركة بين جميع الخيوط في كتلة. تُستخدم الذاكرة المشتركة للاتصال بين الخيوط وتبادل البيانات داخل الكتلة.
  • ذاكرة الثوابت: مساحة ذاكرة للقراءة فقط، تُستخدم لبث البيانات القابلة للقراءة فقط إلى جميع الخيوط.
  • ذاكرة النسيج: مساحة ذاكرة للقراءة فقط، محسّنة للموقعية المكانية والتي يتم الوصول إليها عبر ذاكرة التخزين المؤقت للنسيج. تُستخدم ذاكرة النسيج بشكل أكثر شيوعًا في أحمال العمل الجرافيكية.
  • الذاكرة المحلية: مساحة ذاكرة خاصة لكل خيط، تُستخدم للتصريف المسجل والهياكل البيانية الكبيرة. يتم عادةً ترجمة الذاكرة المحلية إلى الذاكرة العالمية.

إن الاستخدام الفعال للسلسلة الهرمية للذاكرة أمر حاسم لتحقيق أداء عالٍ على بطاقات الرسومات (GPUs). ينبغي على المبرمجين أن يهدفوا إلى تعظيم استخدام الذاكرة المشتركة وتقليل الوصول إلى الذاكرة العالمية لتقليل تأخير الذاكرة وعقبات النطاق الترددي.

يوضح الشكل 2.3 السلسلة الهرمية لذاكرة بطاقة الرسومات.

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

الشكل رقمهنا ترجمة الملف إلى اللغة العربية مع الحفاظ على الرموز البرمجية:

معماريات مجموعات التعليمات للبطاقات الرسومية (GPUs)

تحدد معماريات مجموعات تعليمات البطاقات الرسومية (ISAs) الواجهة منخفضة المستوى بين البرامج والأجهزة. وهي تحدد التعليمات والسجلات وأساليب التوجيه الذاكرة المدعومة من قبل بطاقة الرسومية. فهم ISAs للبطاقات الرسومية أمر أساسي لتطوير رمز البطاقات الرسومية الفعال وتحسين الأداء.

في هذا القسم، سوف نستكشف ISAs لأكبر بائعي البطاقات الرسومية: NVIDIA و AMD. وسنركز على PTX و SASS ISAs الخاصة بشركة NVIDIA، وISA GCN الخاصة بشركة AMD.

ISAs لبطاقات NVIDIA الرسومية

تدعم بطاقات NVIDIA الرسومية مستويين من ISAs: PTX (Parallel Thread Execution) و SASS (Streaming ASSembler). PTX هو ISA افتراضي يوفر هدفًا مستقرًا لمترجمي CUDA، بينما SASS هو ISA الأصلي لبطاقات NVIDIA الرسومية.

PTX (Parallel Thread Execution)

PTX هو ISA افتراضي منخفض المستوى مصمم لبطاقات NVIDIA الرسومية. إنه شبيه بـ LLVM IR أو bytecode Java في أنه يوفر هدفًا مستقرًا ومستقل عن المعمارية للمترجمين. يتم عادة ترجمة برامج CUDA إلى رمز PTX، والذي يتم ترجمته بعد ذلك إلى تعليمات SASS الأصلية بواسطة برنامج تشغيل بطاقة NVIDIA الرسومية.

يدعم PTX مجموعة واسعة من التعليمات الحسابية والذاكرة وتدفق التحكم. لديه عدد غير محدود من السجلات الافتراضية ويدعم الترجيح، مما يسمح بتنفيذ فعال لتدفق التحكم. يوفر PTX أيضًا تعليمات خاصة لتزامن الخيوط والعمليات الذرية وأخذ العينات النسيجية.

إليك مثال على رمز PTX لنواة إضافة متجهات بسيطة:

.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 Arabic translation of the provided Markdown file, with the code comments translated:

تترجم هذه الملفات الشفرة المصدرية (PTX) إلى الترجمة العربية. بالنسبة للشفرة ، لا تترجم الشفرة ، بل ترجم التعليقات فقط.

```cóيعرف هذا الكود PTX دالة نواة vecAdd تأخذ أربعة معلمات: مؤشرات إلى متجهي الإدخال والإخراج ، وحجم المتجهات. تقوم النواة بحساب معرف الخيط العالمي ، وتحميل العناصر المقابلة من متجهات الإدخال ، وإجراء الإضافة ، وتخزين النتيجة في متجه الإخراج.

SASS (Streaming ASSembler)

SASS هو ISA الأصلي لبطاقات الرسومات NVIDIA. إنه ISA منخفض المستوى ، محدد للآلة ، يتطابق مباشرة مع الأجهزة الرسومية للبطاقة. يتم إنشاء تعليمات SASS بواسطة برنامج التشغيل لبطاقة الرسومات NVIDIA من كود PTX ولا يكون مرئيًا عادةً للبرمجين.

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

فيما يلي مثال على شفرة SASS لنواة الإضافة المتجهية:

code_version_number 90
                     @P0 LDG.E R2, [R8];
                     @P0 LDG.E R0, [R10];
                     @P0 FADD R0, R2, R0;
                     @P0 STG.E [R12], R0;
                         EXIT;

يقابل هذا الشفرة SASS الشفرة PTX المُعرضة سابقًا. إنها تحمل عناصر متجه الإدخال من الذاكرة العالمية (LDG.E) ، وتجري الإضافة (FADD) ، وتخزن النتيجة مرة أخرى في الذاكرة العالمية (STG.E) ، وتخرج من النواة.

AMD Graphics Core Next ISA

تستخدم بطاقات الرسومات AMD معمارية وAMD Graphics Core Next (GCN) ISA. GCN هو ISA قائم على RISC مصمم لدعم أحمال العمل الرسومية والحسابية.

يقدم GCN مجموعة من الميزات الرئيسية مثل:

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

مُحوّل 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. يجب أن تكون الخوارزميات مهيكلة بشكل يقلل من الاختلاف في التفرع قدر الإمكان. ويمكن تحقيق ذلك من خلال تقليل استخدام التفرعات المعتمدة على البيانات### استغلال تسلسل الذاكرة

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

موازنة الحساب والوصول إلى الذاكرة

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

تقليل نقل البيانات بين المضيف والجهاز

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

هناك أنماط تصميم خوارزمية موازية شائعة الاستخدام عند تطوير نوى وحدة المعالجة الرسومية:

  • خريطة: ينفذ كل خيط العملية نفسها على عنصر بيانات مختلف، مما يسمح بمعالجة موازية بسيطة لمجموعات بيانات كبيرة.

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

  • مسح: المعروف أيضًا باسم الجمع المتراكم، يُستخدم المسح لحساب المجموع التراكمي لعناصر في مصفوفة. تُعد خوارزميات المسح الموازية الفعالة مكونات بناء رئيسية لكثير من التطبيقات المسرعة بواسطة وحدة المعالجة الرسومية.

  • قالب: يحسب كل خيط قيمة استنادًا إلى عناصر البيانات المجاورة. تُعد عمليات القالب شائعة في المحاكاات العلمية ومعالجة الصور.هنا الترجمة العربية للملف:

  • الجمع/التشتت: تقرأ الخيوط من (جمع) أو تكتب إلى (تشتت) مواقع تعسفية في الذاكرة العامة. التخطيط المتأني للبيانات وأنماط الوصول مطلوبة للكفاءة.

توضح الشكل 3.20 مثالًا على نمط الخريطة، حيث تطبق كل خيط وظيفة (مثل الجذر التربيعي) على عنصر مختلف من مصفوفة الإدخال.

مصفوفة الإدخال:
               |  |   |   |   |   |   |   |
               v  v   v   v   v   v   v   v
              ______________________________
الخيوط:       |    |    |    |    |    |    |    |
             |____|____|____|____|____|____|____|
                |    |    |    |    |    |    |
                v    v    v    v    v    v    v
مصفوفة الإخراج:

الشكل 3.20: مثال على نمط الخريطة في برمجة وحدة المعالجة الرسومية.

الخاتمة

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

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

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

Chapter 1 Introduction to Gpu Chip DesignChapter 3 Parallel Programming Models