الدرس 17: __syncthreads() وحالات التسابق
عندما تُحضِّر البيانات في shared memory ثم تقرؤها عبر threads مختلفة، تنشأ مشكلة: إن threads داخل block لا تعمل في تزامن تام. قد يصل thread واحد إلى سطر القراءة قبل أن ينتهي thread آخر من كتابة الخانة التي يحتاجها — وهذه حالة تسابق (race condition) تعيد قيمة قديمة أو بيانات عشوائية. الحل هو __syncth
مجموعة من الناس يضعون الكتب على رفّ مشترك، ثم يحتاج كل منهم إلى التقاط كتاب وضعه شخص آخر. إذا بدأت الالتقاط قبل أن ينتهي الجميع من الوضع، فسيمدّ بعضهم يده إلى مكان فارغ. __syncthreads() يقول 'لا أحد يلتقط حتى يضع الجميع.'
- __syncthreads()
- حاجز داخل block: كل threads تتوقف عند هذه النقطة حتى يصل الجميع، وعندها فقط تتابع معًا.
- حالة تسابق (race condition)
- عندما تعتمد النتيجة على الترتيب الذي تصل به threads إلى كتابة/قراءة. هنا: القراءة من shared قبل أن تنتهي الكتابة.
- حاجز (barrier)
- نقطة انتظار يجب أن تصل إليها كل threads في block قبل أن يتابع أي منها. يمنع التسابق بين كتابة وقراءة.
- نطاق الحاجز (block)
- __syncthreads() يعمل فقط داخل block واحد. ليس له أثر عبر blocks، ويجب ألا يُستدعى في فرع متشعب.