שיעור 26: כפל מטריצות — ה-kernel הנאיבי
כפל מטריצות C = A * B מחשב כל איבר C[row][col] כסכום של מכפלות לאורך הממד המשותף K: C[row][col] = סכום על k של A[row][k] * B[k][col]. הדרך הישירה (הנאיבית) למקבל זאת על GPU היא להקצות thread אחד לכל איבר של C: ה-thread מחשב את ה-row וה-col שלו מתוך האינדקסים הדו-ממדיים (row = blockIdx.y * blockDim.y
דמיין/י לוח כפל ענק שצריך למלא. בגרסה הנאיבית כל פועל אחראי על משבצת אחת בלבד, אבל כדי למלא אותה הוא הולך עד מחסן רחוק ומביא שורה שלמה ועמודה שלמה — והפועל שלידו הולך לאותו מחסן ומביא בדיוק את אותה שורה שוב. הרבה הליכות מיותרות.
- מיפוי thread לאיבר
- כל thread אחראי על איבר אחד של C. הוא מחשב row ו-col מהאינדקסים הדו-ממדיים ומחשב את C[row][col].
- לולאת המכפלה (k-loop)
- הלולאה על k שצוברת sum += A[row*K + k] * B[k*N + col]. זהו מכפלה סקלרית בין שורה של A לעמודה של B.
- אינדקס שטוח row-major
- מטריצה בגודל rows x cols נשמרת שורה-אחר-שורה, כך שהאיבר [r][c] יושב בכתובת r*cols + c.
- קריאות גלובליות חוזרות
- בגרסה הנאיבית כל איבר של A ושל B נקרא מהזיכרון הגלובלי שוב ושוב על-ידי threads שונים, ובכך נשרף רוחב הפס.
