שיעור 25: scan מקבילי — ה-kernel המוכפל
בשיעור הקודם ראינו מה סכום-קידומות (prefix sum) מחשב ואת הרעיון המוכפל: בכל סיבוב כל איבר מוסיף את השכן במרחק d, ו-d מוכפל (1, 2, 4...). כאן נבנה את זה כ-kernel אמיתי ב-CUDA. נעבוד בתוך בלוק יחיד עם מערך בזיכרון משותף (__shared__) שבו thread אחד אחראי על מקום אחד. הקושי האמיתי הוא לא הנוסחה אלא הסנכ
כל thread אחראי על מקום אחד במערך המשותף. בכל צעד הוא קודם קורא את הערך d מקומות שמאלה ושומר אותו בצד, מחכה שכולם קראו (__syncthreads), רק אז מוסיף אותו לעצמו, ושוב מחכה שכולם כתבו לפני הצעד הבא. ההמתנות הן מה ששומר על הסדר ומונע דריסה.
- Hillis-Steele
- ה-kernel המוכפל ל-scan ב-log2(n) צעדים: בצעד d, כל thread עם tid>=d מוסיף את הערך מ-tid-d. d מוכפל בכל צעד.
- tile בזיכרון משותף
- מערך __shared__ לכל הבלוק. ה-scan טוען אליו את הקלט, עובד עליו במקום בכל הצעדים, ובסוף כותב ל-out.
- מחיצות קריאה-כתיבה
- __syncthreads() אחרי הקריאה (כדי שכולם קראו לפני כתיבה) ואחרי הכתיבה (כדי שהכתיבות גלויות לצעד הבא).
- מספר הצעדים log2(n)
- הלולאה for(d=1; d<n; d*=2) רצה log2(n) פעמים, כי d מוכפל בכל צעד עד שהוא מגיע ל-n.
