שיעור 17: __syncthreads() ותנאי מרוץ

כשטוענים נתונים ל-shared memory ואז קוראים אותם בין threads שונים, נוצרת בעיה: ה-threads ב-block לא רצים בדיוק יחד. thread אחד עלול להגיע לשורת הקריאה לפני ש-thread אחר סיים לכתוב את התא שהוא צריך — וזה תנאי מרוץ (race condition), שמחזיר ערך ישן או זבל. הפתרון הוא __syncthreads(): מחסום (barrier) שא

קבוצת אנשים מעבירה ספרים למדף משותף, ואז כל אחד צריך לקחת ספר שמישהו אחר שם. אם מתחילים לקחת לפני שכולם סיימו להניח, חלק יושיטו יד למקום ריק. __syncthreads() זה לומר 'אף אחד לא לוקח עד שכולם הניחו'.

__syncthreads(): מחסום (barrier) בתוך block: כל ה-threads עוצרים בנקודה הזו עד שכולם הגיעו, ורק אז ממשיכים יחד.
תנאי מרוץ (race condition): מצב שבו התוצאה תלויה בסדר שבו threads מגיעים לכתיבה/קריאה. כאן: קריאה מ-shared לפני שהכתיבה הסתיימה.
מחסום (barrier): נקודת המתנה שכל ה-threads ב-block חייבים להגיע אליה לפני שמישהו ממשיך. בולמת מרוצים בין כתיבה לקריאה.
תחום המחסום (block): __syncthreads() עובד רק בתוך block בודד. אין לו השפעה בין blocks, ואסור לקרוא לו בענף מתפצל.