تدفق المعالجات المتعددة والكتل والخيوط (CUDA)
https://stackoverflow.com/questions/3519598
italiano
english
français
española
中国
日本の
العربية
Deutsch
한국어
Português
Russianسؤال
ما هي العلاقة بين نواة CUDA ، المعالج متعدد البث ونموذج CUDA للكتل والخيوط؟
ما الذي يتم تعيينه إلى ماذا وماذا متوازي وكيف؟ وما هو أكثر كفاءة ، وزيادة عدد الكتل أو عدد المواضيع؟
فهمي الحالي هو أن هناك 8 نوى CUDA لكل معالجات متعددة. وأن كل نواة CUDA ستتمكن من تنفيذ كتلة CUDA واحدة في وقت واحد. ويتم تنفيذ جميع المواضيع في تلك الكتلة بشكل تسلسلي في هذا القلب بالذات.
هل هذا صحيح؟
المحلول
تم وصف تخطيط الخيط / الكتلة بالتفصيل في دليل البرمجة CUDA. على وجه الخصوص ، ينص الفصل 4:
بنية CUDA مبنية حول مجموعة قابلة للتطوير من المعالجات المتعددة المتدفقات (SMS). عندما يستدعي برنامج CUDA على وحدة المعالجة المركزية المضيفة شبكة kernel ، يتم تعداد كتل الشبكة وتوزيعها على المعالجات المتعددة مع قدرة التنفيذ المتاحة. تنفذ مؤشرات ترابط كتلة مؤشر الترابط بشكل متزامن على معالج متعدد ، ويمكن تنفيذ كتل مؤشرات الترابط المتعددة بشكل متزامن على معالج متعدد. مع انتهاء كتل الخيط ، يتم إطلاق كتل جديدة على المعالجات المتعددة.
يحتوي كل SM على 8 نوى CUDA ، وفي أي وقت يقومون بتنفيذ تشوه واحد من 32 موضوعًا - لذلك يتطلب الأمر 4 دورات على مدار الساعة لإصدار تعليمات واحدة للالتفاف بأكمله. يمكنك أن تفترض أن المواضيع في أي تشوه معين تنفذ في خطوة القفل ، ولكن للمزامنة عبر الاعوجاج ، تحتاج إلى استخدام __syncthreads().
نصائح أخرى
بالنسبة إلى GTX 970 ، يوجد 13 معالجات متعددة (SM) مع 128 نوى CUDA لكل منهما. وتسمى النوى CUDA أيضا معالجات الدفق (SP).
يمكنك تحديد الشبكات التي تقوم بتعيينات على وحدة معالجة الرسومات.
يمكنك تحديد الكتل التي تقوم بتخطيطات خريطة لدفق المعالجات (128 CUDA النوى لكل SM).
يتم تشكيل الاعوجاج الواحد دائمًا من خلال 32 مؤشر ترابط ويتم تنفيذ جميع مؤشرات الترابط من الاعوجاج في وقت واحد.
لاستخدام القوة الكاملة الممكنة ل GPU ، تحتاج إلى مزيد من المواضيع لكل SM أكثر من SM. لكل قدرة حسابية ، يوجد عدد معين من الخيوط التي يمكن أن تتواجد في SM واحدة في وقت واحد. جميع الكتل التي تحددها هي في قائمة الانتظار وانتظر SM للحصول على الموارد (عدد SPS مجانًا) ، ثم يتم تحميلها. يبدأ SM في تنفيذ الاعوجارات. نظرًا لأن WARP واحد يحتوي فقط على 32 موضوعًا و SM على سبيل المثال 128 SPS A SM يمكنه تنفيذ 4 WARPs في وقت معين. الشيء هو إذا قامت مؤشرات الترابط بالوصول إلى الذاكرة ، فسيتم حظر مؤشر الترابط حتى يتم استيفاء طلب الذاكرة الخاص به. في الأرقام: يحتوي حساب الحساب على SP على زمن انتقال من 18 إلى 22 دورة في حين أن الوصول إلى الذاكرة العالمي غير المشترك يمكن أن يستغرق ما يصل إلى 300-400 دورة. هذا يعني أنه إذا كانت مؤشرات الترابط الخاصة بـ Warp تنتظر البيانات فقط مجموعة فرعية من 128 SPS ستعمل. لذلك فإن مفاتيح الجدولة لتنفيذ تشوه آخر إذا كان متاحًا. وإذا كان هذا الاعوجاج يحظر ذلك ، فإنه ينفذ التالي وما إلى ذلك. ويسمى هذا المفهوم اختباء الكمون. يحدد عدد الاعوجارات وحجم الكتلة الإشغال (من عدد الاعوجارات التي يمكن أن تختارها SM لتنفيذها). إذا كان الإشغال مرتفعًا فمن غير المرجح أنه لا يوجد عمل ل SPS.
إن بيانك بأن كل نواة CUDA ستنفذ كتلة واحدة في كل مرة خاطئة. إذا تحدثت عن دفق المعالجات المتعددة ، فيمكنهم تنفيذ الاعوجارات من جميع مؤشرات الترابط التي تتواجد في SM. إذا كان حجم كتلة واحدة من 256 مؤشر ترابط وتتيح وحدة معالجة الرسومات الخاصة بك 2048 مؤشرات الترابط للمقيمين لكل SM ، فسيكون لكل SM 8 كتل يمكنها اختيار SM لتنفيذها. يتم تنفيذ جميع خيوط الاعوجارات المنفذة بالتوازي.
تجد أرقامًا لقدرات حسابية مختلفة وعمليات GPU هنا:https://en.wikipedia.org/wiki/Cuda#Limitations
يمكنك تنزيل ورقة حساب الإشغال من Nvidia ورقة حساب الإشغال (بواسطة Nvidia).
سيقوم موزع عمل حساب COMPUTE بجدولة كتلة مؤشر ترابط (CTA) على SM فقط إذا كان لدى SM موارد كافية لكتلة الخيط (الذاكرة المشتركة ، الاعوجاج ، السجلات ، الحواجز ، ...). موارد مستوى كتلة الموضوع يتم تخصيص هذه الذاكرة المشتركة. يخلق التخصيص عمليات الاعزاف الكافية لجميع المواضيع في كتلة الخيط. يخصص مدير الموارد الاعWAR جولة روبن إلى SM Subsitions. يحتوي كل حاجز SM على جدولة WARP وملف التسجيل ووحدات التنفيذ. بمجرد تخصيص الاعوجاج لتصليح حاجز ، سيبقى على الجزء الفرعي حتى يكمل أو يتم استباقه بواسطة مفتاح السياق (بنية باسكال). على مفتاح السياق ، سيتم استعادة WARP إلى نفس معرف الالتفاف نفسه.
عندما تكمل جميع مؤشرات الترابط في WARP ، تنتظر جدولة WARP جميع الإرشادات المعلقة الصادرة عن WARP لإكمالها ، ثم يقوم مدير الموارد بإطلاق موارد مستوى الاعوجاج والتي تشمل ملف WARP-ID.
عند اكتمال جميع الاعوجارات في كتلة مؤشر ترابط ، يتم إصدار موارد مستوى الكتلة ويبلغ SM موزع العمل الحساب الذي أنجزته الكتلة.
بمجرد تخصيص الاعوجاج للاضطراب الفرعي ويتم تخصيص جميع الموارد ، تعتبر الاعوجاج تعنيًا نشطًا وهو أن جدولة الاعوجاج تتتبع بنشاط حالة الاعوجاج. في كل دورة ، حدد جدولة الاعوجاج الاعوجارات النشطة التي تتوقف والتي تكون مؤهلة لإصدار تعليمات. يختار جدولة الاعوجاج أعلى أولوية الاعوجاج المؤهلة والقضايا 1-2 تعليمات متتالية من الاعوجاج. قواعد القضية المزدوجة خاصة بكل بنية. إذا أصدرت WARP تحميل الذاكرة ، فيمكنه الاستمرار في تنفيذ تعليمات مستقلة حتى تصل إلى تعليمات تابعة. سيقوم WARP بعد ذلك بالإبلاغ عن توقف حتى يكمل الحمل. وينطبق الشيء نفسه على تعليمات الرياضيات التابعة. تم تصميم بنية SM لإخفاء كل من ALU و Amermy Carency عن طريق التبديل لكل دورة بين الاعوجارات.
لا تستخدم هذه الإجابة مصطلح CUDA Core لأن هذا يقدم نموذجًا عقليًا غير صحيح. النوى CUDA هي وحدات تنفيذ نقطة عائمة واحدة الدقة/عدد صحيح. معدل القضية والاعتماد زمن استجابة خاص لكل بنية. يحتوي كل جزء فرعي SM و SM على وحدات تنفيذ أخرى بما في ذلك وحدات التحميل/المتجر ، ووحدات نقاط عائمة دقة مزدوجة ، ونصف وحدات نقطية عائمة ، ووحدات فرعية ، إلخ.
من أجل زيادة الأداء إلى الحد الأقصى ، يتعين على المطور فهم تبادل الكتل مقابل Warps مقابل السجلات/الخيط.
مصطلح الإشغال هو نسبة الاعوجارات النشطة إلى أقصى درجات الاعوجارات على SM. Kepler - Pascal Architecture (باستثناء GP100) لديها 4 جدولة الاعوجاج لكل SM. يجب أن يكون الحد الأدنى لعدد الاعوجارات لكل SM على الأقل مساويًا لعدد جدولي الاعوجاج. إذا كان للهندسة المعمارية زمن تنفيذ تابع لـ 6 دورات (Maxwell و Pascal) ، فستحتاج إلى 6 أعوجة على الأقل لكل جدولة أي 24 لكل SM (24 /64 = 37.5 ٪ من الإشغال) لتغطية الكمون. إذا كانت المواضيع ذات مستوى تعليمي التوازي ، فيمكن تقليل ذلك. تقريبا جميع kernels تصدر تعليمات الكمون المتغير مثل أحمال الذاكرة التي يمكن أن تأخذ 80-1000 دورة. هذا يتطلب المزيد من الاعوجارات النشطة لكل جدولة الاعوجاج لإخفاء الكمون. بالنسبة لكل نواة ، هناك نقطة متداولة بين عدد من الاعوجاج والموارد الأخرى مثل الذاكرة المشتركة أو السجلات ، لذا لا يُنصح بتحسين الإشغال بنسبة 100 ٪ حيث من المحتمل أن يتم تقديم بعض التضحية الأخرى. يمكن لـ CUDA Profiler المساعدة في تحديد معدل إصدار التعليمات ، والإشغال ، وأسباب المماطلة من أجل مساعدة المطور على تحديد هذا التوازن.
يمكن أن يؤثر حجم كتلة الخيط على الأداء. إذا كان لدى kernel كتل كبيرة وتستخدم حواجز التزامن ، فيمكن أن تكون أكشاك الحاجز أسبابًا للمماطلة القادمة. يمكن تخفيف ذلك عن طريق تقليل الاعوجارات لكل كتلة موضوع.
هناك متعددة المعالجات متعددة البث على جهاز واحد.
قد تحتوي SM على كتل متعددة. قد تحتوي كل كتلة على عدة مؤشرات ترابط.
تحتوي SM على عدة نوى CUDA (كمطور ، يجب ألا تهتم بهذا لأنه يتم استخلاصه بواسطة WARP) ، والذي سيعمل على موضوع. SM تعمل دائمًا على الاعوجاج من المواضيع (دائمًا 32). ستعمل الاعوجاج فقط على موضوع من نفس الكتلة.
SM و Block لهما حدود على عدد الخيوط وعدد التسجيل والذاكرة المشتركة.
