اكتشاف حالات التسابق باستخدام racecheck

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

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

تحلل أداة racecheck في compute-sanitizer عمليات الوصول إلى الذاكرة المشتركة، وتبلغ عن المخاطر التي قد تخفيها عملية التشغيل العادية.

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

أدرجوا __syncthreads بين الكتابة والقراءة، لكي ينتهي كل خيط في الكتلة من الكتابة قبل أن يبدأ أي خيط بالقراءة.

تكتشف racecheck أيضًا خطر الكتابة بعد الكتابة، حيث يخزّن خيطان البيانات في الخانة المشتركة نفسها من دون ترتيب بينهما.

الخطر الأكثر شيوعًا هو القراءة بعد الكتابة: يقرأ خيط قيمةً لا يزال خيط آخر يكتبها، فيقرأ بيانات قديمة.

تذكر racecheck نوع الخطر والخيطين والعنوان المشترك المعنيين، لكي تعرفوا بالضبط أي حاجز مفقود.

تركّز racecheck على الذاكرة __shared__ داخل الكتلة؛ أما السباقات العامة فعادةً فتحتاج إلى منطق دقيق أو عمليات ذرية.

كما في memcheck، يتيح التجميع باستخدام -lineinfo لـ racecheck الإشارة إلى سطر المصدر الدقيق لكل خطر.

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

استخدمتم racecheck لاكتشاف مخاطر الذاكرة المشتركة، وقرأتم تقاريرها، وأصلحتموها باستخدام حواجز __syncthreads. لقد كُشفت أخطاء التوقيت. 🎯