שיעור 15: לזהות ולתקן גישה עם stride
בשיעור הקודם ראינו מה זה coalescing: כש-threads עוקבים ב-warp נוגעים בכתובות רצופות, מנה אחת משרתת את כולם, וכשהם קופצים ב-stride גדול ה-warp צורך הרבה מנות מבוזבזות. עכשיו נלמד את הכישור המעשי — לזהות גישה לא מאוחדת בתוך kernel, ולתקן אותה. הסימן הקלאסי הוא אינדקס שמכפיל את ה-threadIdx במספר גדול,
דמיין/י 4 פועלים שצריכים לקטוף תפוחים משורות עץ. בהתחלה כל פועל רץ לעמודה אחרת בקצה אחר של המטע — 4 ריצות רחוקות. אם במקום זה כל פועל לוקח את התפוח הבא בשורה אחת צמודה, כולם עובדים זה ליד זה בלקיחה אחת. לא שינינו את המטע — רק את מי-קוטף-מה.
- גישה עם stride
- דפוס שבו threads עוקבים קופצים ב-stride גדול (למשל data[threadIdx.x * N]), אז כתובותיהם רחוקות וכל אחת נופלת במנת זיכרון אחרת.
- מאוחד מול לא-מאוחד
- מאוחד: כתובות threads שכנים רצופות, מנה אחת לכל ה-warp. לא-מאוחד: הכתובות מפוזרות, ה-warp צורך הרבה מנות שרובן מבוזבז.
- תיקון מיפוי האינדקס
- התיקון לגישה מפוזרת: לגרום לשכנים לגעת בשכנים — לתת ל-thread t את האיבר base + t במקום base + t*N, הכול בזיכרון גלובלי.
- סדר row-major
- פריסה שבה איברים של אותה שורה שכנים בזיכרון. מעבר לפי שורה גורם ל-threads שכנים לגעת בכתובות שכנות (מאוחד); מעבר לפי עמודה קופץ ב-width (מפוזר).
