كودا:مزامنة المواضيع

Question

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

// some math code, calculating d1, d2
if (d1 < 0.5)
{
    buffer[x1] += 1;  // buffer is in the global memory
}
if (d2 < 0.5)
{
    buffer[x2] += 1;
}
// some more math code.

وقد يدخل بعض الخيوط في واحد للشروط، وقد يدخل بعضها في كليهما، وقد لا يدخل البعض الآخر في أي منهما.

الآن من أجل جعل كل الخيوط تعود إلى "فعل الشيء نفسه" مرة أخرى بعد الشروط، هل يجب أن أقوم بمزامنتها بعد الشروط باستخدام __syncthreads() ؟أم أن هذا يحدث بطريقة أو بأخرى بشكل تلقائي؟
يمكن أن يكون هناك موضوعين لا يفعلون نفس الشيء لأن أحدهم تأخر عن عملية واحدة، وبالتالي يفسد الأمر على الجميع؟أم أن هناك بعض الجهود خلف الكواليس لحملهم على فعل الشيء نفسه مرة أخرى بعد الفرع؟

Answer 1

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

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

---

يحرر: أعدت تنظيم الفقرات القليلة التالية لتوضيح بعض الأمور.

هناك بالفعل مسألتان مختلفتان هنا:مزامنة التعليمات ورؤية الذاكرة.

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

• __threadfence() يضمن رؤية الذاكرة العالمية ولكنه لا يقوم بأي مزامنة للتعليمات، لذلك من واقع خبرتي فهو محدود الاستخدام (ولكن راجع نموذج التعليمات البرمجية في الملحق ب.5).

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

• إذا كنت تحتاج فقط إلى زيادة العدادات المشتركة أو العامة، ففكر في استخدام دالة الزيادة الذرية atomicInc() (الملحق ب.10).في حالة الكود الخاص بك أعلاه، إذا x1 و x2 ليست فريدة عالميًا (عبر جميع سلاسل الرسائل في شبكتك)، فإن الزيادات غير الذرية ستؤدي إلى حالة سباق، مشابهة للفقرة الأخيرة من الملحق ب.2.4.

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

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

Answer 2

ومن القسم 6.1 من دليل أفضل الممارسات CUDA:

<اقتباس فقرة>   

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

وهكذا، لا تحتاج إلى القيام بأي شيء خاص.

Answer 3

في استجابة جبرائيل:

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

وماذا لو كان سبب من الأسباب تحتاج و () و (ز) () في نفس الموضوع لأنك كنت تستخدم ذاكرة السجل، وتريد تسجيل أو البيانات المشتركة من و للوصول الى ز؟ وهذا هو، لمشكلتي، والسبب كله لمزامنة عبر كتل لأن هناك حاجة إلى البيانات من و في ز - وكسره إلى نواة تتطلب كمية كبيرة من الذاكرة العمومية إضافية لنقل البيانات السجل من و إلى g، وأنا ترغب في تجنب

Answer 4

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

buffer[x1] += (d1 < 0.5);
buffer[x2] += (d2 < 0.5);

ويجب عليك معرفة ما اذا كان يمكنك استخدام الذاكرة المشتركة والوصول إلى الذاكرة العالمية في نمط ملتئم. ومن المؤكد ايضا ان كنت لا تريد أن تكتب لنفس المؤشر في أكثر من 1 الموضوع.