كتلة الخيوط (برمجة CUDA)

📅 آخر تحديث: ٢١ يوليو ٢٠٢٦
✍️ نُشر في: ٢١ يوليو ٢٠٢٦
📷 يحتوي على 7 صور
الموسوعة الحرة

كتلة الخيوط هي تجريد برمجي يُمثل مجموعة من الخيوط التي يُمكن تنفيذها بالتسلسل أو بالتوازي. ولتحسين عملية ربط العمليات والبيانات ، تُجمع الخيوط في كتل خيوط. كان عدد الخيوط في كتلة الخيوط محدودًا سابقًا بـ 512 خيطًا كحد أقصى لكل كتلة، ولكن اعتبارًا من مارس 2010، ومع قدرات الحوسبة 2.x وما فوق، يُمكن أن تحتوي الكتل على ما يصل إلى 1024 خيطًا. تعمل الخيوط في نفس كتلة الخيوط على نفس معالج التدفق المتعدد. [ 1 ] يُمكن للخيوط في نفس الكتلة التواصل فيما بينها عبر الذاكرة المشتركة ، أو مزامنة الحاجز ، أو غيرها من آليات المزامنة الأساسية مثل العمليات الذرية.

تُدمج عدة كتل لتشكيل شبكة. تحتوي جميع الكتل في الشبكة نفسها على العدد نفسه من الخيوط. عدد الخيوط في الكتلة محدود، ولكن يمكن استخدام الشبكات في العمليات الحسابية التي تتطلب عددًا كبيرًا من كتل الخيوط للعمل بالتوازي والاستفادة من جميع المعالجات المتعددة المتاحة.

CUDA هي منصة حوسبة متوازية ونموذج برمجة تستخدمه لغات البرمجة عالية المستوى للاستفادة من التوازي. في CUDA، تُنفَّذ النواة بمساعدة الخيوط. الخيط هو كيان مجرد يُمثِّل تنفيذ النواة. النواة هي دالة تُترجم لتشغيلها على جهاز خاص. تستخدم التطبيقات متعددة الخيوط العديد من هذه الخيوط التي تعمل في الوقت نفسه لتنظيم الحوسبة المتوازية. لكل خيط فهرس يُستخدم لحساب مواقع عناوين الذاكرة ، وكذلك لاتخاذ قرارات التحكم.

أبعاد

تعتمد CUDA على نموذج برمجة غير متجانس يُستخدم لتشغيل برامج التطبيقات على الأجهزة المضيفة. ويشبه نموذج تنفيذها نموذج OpenCL . في هذا النموذج، نبدأ بتنفيذ التطبيق على الجهاز المضيف، والذي يكون عادةً نواة وحدة المعالجة المركزية (CPU) . ويكون هذا الجهاز مُوجَّهًا نحو الإنتاجية العالية، أي نواة وحدة معالجة الرسومات (GPU) التي تُجري عمليات حسابية متوازية. تُستخدم وظائف النواة (Kernel functions) لتنفيذ هذه العمليات المتوازية. بمجرد انتهاء تنفيذ وظائف النواة، يُعاد التحكم إلى الجهاز المضيف الذي يستأنف التنفيذ التسلسلي.

نظرًا لأن العديد من التطبيقات المتوازية تتضمن بيانات متعددة الأبعاد، فمن الملائم تنظيم كتل الخيوط في مصفوفات أحادية أو ثنائية أو ثلاثية الأبعاد من الخيوط. يجب أن تكون الكتل في الشبكة قابلة للتنفيذ بشكل مستقل، حيث لا يمكن التواصل أو التعاون بين الكتل داخل الشبكة. عند تشغيل النواة، يتم تحديد عدد الخيوط لكل كتلة خيوط، وعدد كتل الخيوط، وهذا بدوره يحدد العدد الإجمالي لخيوط CUDA التي يتم تشغيلها. [ 2 ] الحد الأقصى لأبعاد x و y و z للكتلة هو 1024 و 1024 و 64 على التوالي، ويجب تخصيصها بحيث يكون x × y × z ≤ 1024، وهو الحد الأقصى لعدد الخيوط لكل كتلة. [ 3 ] يمكن تنظيم الكتل في شبكات أحادية أو ثنائية أو ثلاثية الأبعاد، يصل عدد كتلها إلى 2 ^31 -1 و 65535 و 65535 كتلة في الأبعاد x و y و z على التوالي. [ 3 ] على عكس الحد الأقصى للخيوط لكل كتلة، لا يوجد حد للكتل لكل شبكة يختلف عن الحد الأقصى لأبعاد الشبكة.

الفهرسة

الفهرسة أحادية البعد

يرتبط كل خيط في CUDA بفهرس معين بحيث يمكنه حساب مواقع الذاكرة والوصول إليها في مصفوفة.

لنفترض مثالاً لمصفوفة تحتوي على 512 عنصرًا. أحد أشكال تنظيمها هو شبكة تتكون من كتلة واحدة تحتوي على 512 خيطًا. لنفترض وجود مصفوفة C تحتوي على 512 عنصرًا، وهي ناتجة عن ضرب عنصر بعنصر في مصفوفتين A وB، كلتاهما تحتوي على 512 عنصرًا. لكل خيط فهرس i، حيث يقوم بضرب العنصر i في A وB، ثم يخزن النتيجة في العنصر i من C. يتم حساب i باستخدام blockIdx (وهو 0 في هذه الحالة لوجود كتلة واحدة فقط)، وblockDim (512 في هذه الحالة لأن الكتلة تحتوي على 512 عنصرًا)، وthreadIdx الذي يتراوح من 0 إلى 511 لكل كتلة.

التسلسل الهرمي للخيوط في برمجة CUDA [ 4 ]

يتم حساب مؤشر الخيط i باستخدام الصيغة التالية  :

أنا=بلoجكأنادx.x*بلoجكدأنام.x+تحرهـأدأنادx.x{\displaystyle i=blockIdx.x*blockDim.x+threadIdx.x}

blockIdx.x هو مُعرِّف الكتلة في البُعد x

blockDim.x هو البعد x من أبعاد الكتلة

يمثل threadIdx.x البعد x لمعرف الخيط

وبالتالي، ستكون لـ 'i' قيم تتراوح من 0 إلى 511 والتي تغطي المصفوفة بأكملها.

إذا أردنا إجراء عمليات حسابية على مصفوفة أكبر من 1024 عنصرًا، فيمكننا استخدام عدة كتل، كل منها تحتوي على 1024 خيطًا. لنأخذ مثالًا على مصفوفة تحتوي على 2048 عنصرًا. في هذه الحالة، لدينا كتلتان من الخيوط، كل منهما تحتوي على 1024 خيطًا. بالتالي، ستتراوح قيم مُعرّفات الخيوط من 0 إلى 1023، بينما سيتراوح مُعرّف الكتلة من 0 إلى 1، وسيكون بُعد الكتلة 1024. وعليه، ستحصل الكتلة الأولى على قيم فهرسة من 0 إلى 1023، بينما ستحصل الكتلة الأخيرة على قيم فهرسة من 1024 إلى 2047.

وبالتالي، ستقوم كل سلسلة عمليات أولاً بحساب فهرس الذاكرة التي يتعين عليها الوصول إليها، ثم تتابع عملية الحساب. لنفترض مثالاً حيث يتم جمع عناصر من المصفوفتين A وB بالتوازي باستخدام سلاسل العمليات، ويتم تخزين النتائج في مصفوفة C. يظهر الكود المقابل في سلسلة العمليات أدناه  : [ 5 ]

__global__ void vecAddKernel ( float * A , float * B , float * C , int n ) { int index = blockIdx . x * blockDim . x + threadIdx . x ; if ( index < n ) { C [ index ] = A [ index ] + B [ index ] ; } }

الفهرسة ثنائية الأبعاد

وبالمثل، في الشبكات المعقدة للغاية، يجب على كل خيط حساب كل من معرف الكتلة (blockId) ومعرف الخيط (threadId) بناءً على هندسة الشبكة. لنفترض شبكة ثنائية الأبعاد تحتوي على كتل ثنائية الأبعاد. سيتم حساب معرف الخيط ومعرف الكتلة باستخدام الصيغ التالية  :

بلoجكأناد=بلoجكأنادx.x+بلoجكأنادx.y*زرأناددأنام.x;{\displaystyle blockId=blockIdx.x+blockIdx.y*gridDim.x;}تحرهـأدأناد=بلoجكأناد*(بلoجكدأنام.x*بلoجكدأنام.y)+(تحرهـأدأنادx.y*بلoجكدأنام.x)+تحرهـأدأنادx.x;{\displaystyle threadId=blockId*(blockDim.x*blockDim.y)+(threadIdx.y*blockDim.x)+threadIdx.x;}[ 6 ]

منظور الأجهزة

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

مقارنة تصويرية بين منظور المبرمج ومنظور الأجهزة لكتلة الخيوط في وحدة معالجة الرسومات [ 7 ]

معالجات متعددة البث

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

رسم توضيحي لمعالج متعدد التدفق وموارده [ 9 ]

ولتحقيق هذا الغرض، يحتوي النموذج القياسي على ما يلي: [ 8 ]

  • نوى التنفيذ. (وحدات الفاصلة العائمة أحادية الدقة، ووحدات الفاصلة العائمة مزدوجة الدقة، ووحدات الوظائف الخاصة (SFUs)).
  • ذاكرة التخزين المؤقت:
  1. ذاكرة التخزين المؤقت L1 . (لتقليل زمن الوصول إلى الذاكرة).
  2. الذاكرة المشتركة (للبيانات المشتركة بين الخيوط).
  3. ذاكرة تخزين مؤقتة ثابتة (لبث عمليات القراءة من ذاكرة للقراءة فقط ).
  4. ذاكرة التخزين المؤقت للنسيج (لتجميع عرض النطاق الترددي من ذاكرة النسيج).
  • جدولة الالتفافات. (هذه مخصصة لإصدار التعليمات إلى الالتفافات بناءً على سياسات جدولة معينة).
  • عدد كبير من السجلات. (قد يقوم نظام المعالجة المتعددة بتشغيل عدد كبير من الخيوط النشطة في وقت واحد، لذلك من الضروري وجود آلاف السجلات.)

يقوم الجهاز بجدولة كتل الخيوط إلى وحدة معالجة متعددة (SM). بشكل عام، يمكن لوحدة المعالجة المتعددة معالجة عدة كتل خيوط في الوقت نفسه. قد تحتوي وحدة المعالجة المتعددة على ما يصل إلى 8 كتل خيوط إجمالاً. يتم تعيين مُعرّف الخيط لكل خيط بواسطة وحدة المعالجة المتعددة الخاصة به.

عندما يُنفّذ مُعالج متعدد (SM) كتلة خيوط، تُنفّذ جميع الخيوط داخل تلك الكتلة في الوقت نفسه. لذا، لتحرير مساحة في ذاكرة كتلة الخيوط داخل المُعالج المتعدد، من الضروري أن تكون جميع الخيوط في الكتلة قد انتهت من التنفيذ. تُقسّم كل كتلة خيوط إلى وحدات مُجدولة تُعرف باسم "الالتفاف" (warp). سيتم شرح هذه الوحدات بالتفصيل في القسم التالي.

رسم توضيحي لمجدول الالتفاف المزدوج المطبق في بنية فيرمي الدقيقة لشركة إنفيديا [ 10 ]

يُحدد مُجدول الالتفافات في المعالج متعدد الوظائف أيًّا من الالتفافات يحظى بالأولوية أثناء إصدار التعليمات. [ 11 ] وقد نوقشت بعض سياسات تحديد أولويات الالتفافات في الأقسام التالية.

الالتواءات

من الناحية المادية، تتكون كتلة الخيوط من "اللفائف". (هذا المصطلح مأخوذ من مصطلح " النسج " [ 12 ] ). اللفيفة هي مجموعة من 32 خيطًا داخل كتلة الخيوط. في السابق، كان من المضمون أن تُنفذ هذه الخيوط "بشكل متزامن" (أي أن جميع الخيوط داخل اللفيفة تُنفذ التعليمات في وقت واحد)، والأهم من ذلك، أن تتمكن من الوصول إلى كل موقع في الذاكرة باستخدام جميع خيوط اللفيفة أو بدونها. كان هذا السلوك يُمكن أن يؤدي بسهولة إلى حالات جمود (على سبيل المثال، باستخدام عبارات if-brances في الحلقات). مع ذلك، منذ ظهور بنية فولتا ، أصبح تبادل البيانات داخل اللفيفة ممكنًا من خلال أقفال أكثر دقة. [ 13 ] [ 14 ] يتم اختيار هذه الخيوط بشكل تسلسلي بواسطة المعالج المتعدد. [ 15 ]

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

لنفترض وجود مجموعة من 32 خيطًا تُنفّذ تعليمة. إذا لم يكن أحد مُعاملاتها أو كلاهما جاهزًا (أي لم يتم جلبهما من الذاكرة العامة)، تحدث عملية تُسمى " تبديل السياق " تُنقل التحكم إلى مجموعة أخرى. [ 16 ] عند الانتقال من مجموعة معينة، تبقى جميع بيانات تلك المجموعة في ملف التسجيلات، ما يسمح باستئنافها بسرعة عند جاهزية مُعاملاتها. عندما لا يكون للتعليمة أي تبعيات بيانات معلقة، أي عندما يكون كلا مُعامليها جاهزين، تُعتبر المجموعة المعنية جاهزة للتنفيذ. إذا كانت أكثر من مجموعة مؤهلة للتنفيذ، يستخدم المعالج الرئيسي سياسة جدولة المجموعات لتحديد المجموعة التي ستحصل على التعليمة التالية.

فيما يلي مناقشة السياسات المختلفة لجدولة الالتفافات المؤهلة للتنفيذ: [ 17 ]

  1. التناوب الدوري (RR) - يتم جلب التعليمات بالتناوب الدوري. يضمن التناوب الدوري إبقاء وحدات المعالجة المتعددة مشغولة وعدم إهدار دورات الساعة على زمن استجابة الذاكرة.
  2. الأقل جلبًا مؤخرًا (LRF) - في هذه السياسة، يحصل الالتفاف الذي لم يتم جلب التعليمات الخاصة به لأطول فترة زمنية على الأولوية في جلب التعليمات.
  3. سياسة التوزيع العادل (FAIR) [ 17 ] - في هذه السياسة، يضمن المجدول منح جميع مجموعات الالتفاف فرصة عادلة في عدد التعليمات التي يتم جلبها لها. ويقوم بجلب التعليمات إلى مجموعة الالتفاف التي تم جلب أقل عدد ممكن من التعليمات لها.
  4. جدولة الالتفافات المُراعية للأهمية [ 18 ] - تركز هذه السياسة على تحسين وقت تنفيذ كتل الخيوط. خصصت موارد زمنية أكبر للالتفاف الذي يستغرق أطول وقت للتنفيذ. من خلال إعطاء الأولوية للالتفاف الأكثر أهمية، تسمح هذه السياسة لكتل ​​الخيوط بالانتهاء بشكل أسرع، مما يُتيح توفر الموارد بشكل أسرع.

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

In order to take advantage of the warp architecture, programming languages and developers need to understand how to coalesce memory accesses and how to manage control flow divergence. If each thread in a warp takes a different execution path or if each thread accesses significantly divergent memory then the benefits of the warp architecture are lost and performance will significantly degrade.

References

  1. ↑"Chapter 4. Hardware Implementation, The threads of a thread block execute concurrently on one multiprocessor, and multiple thread blocks can execute concurrently on one multiprocessor"(PDF).
  2. ↑"CUDA Thread Model". www.olcf.ornl.gov. Archived from the original on 2016-09-23. Retrieved 2016-09-21.
  3. 12"CUDA Toolkit Documentation: Features and Technical Specifications". docs.nvidia.com. Retrieved 2022-05-24.
  4. ↑"Thread Hierarchy in CUDA Programming". Retrieved 2016-09-21.
  5. ↑Kirk, David; Hwu, Wen-mei W (January 28, 2010). Programming Massively Parallel Processors: A Hands-on Approach.
  6. ↑"Thread Indexing Cheatsheet"(PDF). Retrieved 2016-09-21.
  7. ↑"Thread Optimizations (University of Mayland)"(PDF).
  8. 12Wilt, Nicholas (2013). The CUDA Handbook: A Comprehensive Guide to GPU Programming.
  9. ↑"Thread Optimizations (University of Mayland)"(PDF).
  10. ↑"Thread Optimizations (University of Mayland)"(PDF).
  11. ↑"GPU Computing with CUDA Lecture 2 - CUDA Memories"(PDF).
  12. ↑"Parallel Thread Execution ISA Version 6.0". Developer Zone: CUDA Toolkit Documentation. NVIDIA Corporation. 22 September 2017. Archived from the original on 28 October 2017. Retrieved 27 October 2017.
  13. ↑"1. Volta Tuning Guide — Volta Tuning Guide 13.0 documentation". docs.nvidia.com. Retrieved 2025-08-05.
  14. ↑Nvidia. "Cuda C++ Programming Model V13"(PDF). p. 142. Retrieved 24 August 2025.
  15. ↑"Using CUDA Warp-Level Primitives". Nvidia. 2018-01-15. Retrieved 2020-04-08. NVIDIA GPUs execute groups of threads known as warps in SIMT (Single Instruction, Multiple Thread) fashion
  16. ↑"Memory Issues in CUDA and Execution Scheduling in CUDA"(PDF).
  17. 1 2 "تأثير جلب التعليمات وجدولة الذاكرة على أداء وحدة معالجة الرسومات" (PDF) .
  18. ↑ "CAWS: جدولة الالتفاف الواعية بالحرجة لأحمال عمل GPGPU" (PDF) .