2.2 היררכיית ה threads blocks ו grids תרגול

תרגול - היררכיית ה-threads - blocks ו-grids

בתרגול הזה תתרגמו את ההיררכיה לקוד רץ. תכתבו kernel עם אינדוקס דו-ממדי אמיתי (dim3 grid ו-block) שמעבד תמונה, תחשבו בעצמכם מימדי grid לכיסוי מטריצת NxN עם שומר גבולות, תוכיחו במו עיניכם שהבלוקים רצים בסדר לא צפוי, תממשו סכום כלל-grid עם atomicAdd ותבינו מדוע __syncthreads לא יכול לעשות זאת, ולבסוף תחשבו occupancy ביד ותאמתו מול ה-nvcc. עבדו לפי הסדר - כל תרגיל נשען על הקודם. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס, אבל כל תרגיל כאן ירוץ על כל GPU של NVIDIA.

הכנה

צרו קובץ hier.cu. השתמשו במקרו CUDA_CHECK מההרצאות הקודמות (מ-0.3 או 1.2) בראש כל קובץ. הקומפילציה לאורך כל התרגול (החליפו את sm_90a בארכיטקטורה של הכרטיס שלכם - למשל sm_75 ל-T4):

nvcc -O2 -arch=sm_90a -o hier hier.cu

כדי לברר את ה-compute capability של הכרטיס:

nvidia-smi --query-gpu=name,compute_cap --format=csv

הערה על printf בתוך kernel: הוא נתמך ומדפיס דרך תור פנימי שמתרוקן ב-cudaDeviceSynchronize. לכן תמיד קראו ל-cudaDeviceSynchronize() אחרי launch שמדפיס, אחרת לא תראו פלט.

תרגיל 1 - היפוך תמונה עם אינדוקס דו-ממדי

נעבד "תמונה" בגווני אפור (מערך חד-ממדי של unsigned char, בפריסת row-major) עם grid דו-ממדי.

1. הקצו על הhost תמונה width = 1920, height = 1080, ומלאו כל פיקסל בערך i % 256 (סתם דפוס).
2. העתיקו ל-device, וכתבו kernel invert שמחשב col ו-row מהאינדקסים הדו-ממדיים, שומר גבולות if (row < height && col < width), ומבצע img[row*width+col] = 255 - img[row*width+col].
3. השיקו עם dim3 block(16,16) ו-dim3 grid מעוגל כלפי מעלה בשני הממדים.
4. העתיקו חזרה ואמתו שפיקסל במיקום (row, col) שווה ל-255 - ((row*width+col) % 256).

רמז: המימדים הם dim3 grid((width + block.x - 1)/block.x, (height + block.y - 1)/block.y). שימו לב ש-.x הוא העמודה (הממד המהיר) ו-.y השורה. עם 1920/16 = 120 ו-1080/16 = 67.5, ה-grid יהיה 120 x 68 בלוקים - השורה ה-68 גולשת מעבר ל-1080, ובדיוק בשבילה קיים שומר הגבולות.

תרגיל 2 - כיסוי מטריצת NxN עם שומר גבולות

עכשיו נתרגל את בחירת המימדים כשה-grid לא מחלק בדיוק את הבעיה.

1. עבור N = 1000 (מטריצה 1000 x 1000) ו-dim3 block(16,16), חשבו ביד כמה בלוקים יש בכל ממד וכמה threads עודפים נוצרים בכל ממד.
2. כתבו kernel scale שמכפיל כל איבר של מטריצת float בסקלר, עם שומר גבולות דו-ממדי.
3. הריצו והדפיסו מתוך ה-kernel (רק עבור הthread בפינה row==0 && col==0) את הערכים של gridDim.x, gridDim.y, blockDim.x, blockDim.y כדי לראות את התצורה בפועל.
4. ודאו שכל N*N האיברים עברו scale (אמתו על הhost), למרות שהושקו יותר threads מאיברים.

רמז: (1000 + 15)/16 = 63 בלוקים בכל ממד, כלומר 63*16 = 1008 threads בכל ממד - 8 threads עודפים בכל ממד, וסך הכל 1008*1008 = 1,016,064 threads מול 1,000,000 איברים. בלי שומר הגבולות if (row<N && col<N), ה-16,064 הthreads העודפים היו כותבים מחוץ לתחום וגורמים ל-illegal memory access.

תרגיל 3 - סדר הרצת הבלוקים

נוכיח את חוק העצמאות מההרצאה: הבלוקים רצים בסדר לא מובטח, ואסור להניח עליו כלום.

1. כתבו kernel order המקבל מצביע int* ticket (מאותחל ל-0 ב-device עם cudaMemset). רק הthread threadIdx.x == 0 של כל בלוק מבצע int t = atomicAdd(ticket, 1); ואז printf("block %d -> order %d\n", blockIdx.x, t);.
2. השיקו עם 16 בלוקים וthread אחד לבלוק (order<<<16, 1>>>(...)), וקראו ל-cudaDeviceSynchronize().
3. הריצו את התוכנית חמש פעמים ברצף ורשמו את סדר הבלוקים בכל הרצה.
4. ענו: האם סדר ההרצה שווה תמיד ל-0,1,2,...,15? האם הוא זהה בין הרצות? מה זה מוכיח לגבי הנחות סדר בקוד?

רמז: ה-ticket נותן לכל בלוק "מספר תור" גלובלי לפי סדר ההגעה בפועל שלו ל-atomicAdd. תראו שהסדר לרוב אינו 0..15 ומשתנה בין הרצות (במיוחד עם הרבה בלוקים ו-GPU עמוס). זו הוכחה מעשית שאסור לכתוב קוד שמניח סדר בלוקים. השתמשו ב-atomic ולא במשתנה גלובלי רגיל כי כמה בלוקים עלולים להגיע בו-זמנית - עדכון לא-אטומי היה מאבד ספירות.

תרגיל 4 - סכום כלל-grid עם atomicAdd

נחשב את סכום כל האיברים של וקטור - חישוב שדורש שיתוף פעולה בין כל הבלוקים.

1. הקצו וקטור x באורך n = 1 << 20 ומלאו אותו כולו ב-1.0f (כך שהסכום הצפוי הוא בדיוק n).
2. הקצו float* result יחיד על ה-device ואפסו אותו ב-cudaMemset(result, 0, sizeof(float)).
3. כתבו kernel sumAtomic שבו כל thread (עם שומר גבולות) מוסיף את האיבר שלו לתוצאה: atomicAdd(result, x[i]).
4. השיקו, העתיקו את result חזרה, והשוו ל-n.
5. שאלת ההבנה המרכזית: נסו לחשוב איך הייתם עושים את אותו סכום עם __syncthreads() בלבד, בלי atomics. הסבירו בכתב מדוע זה בלתי אפשרי ברמת ה-grid.

רמז: __syncthreads() הוא מחסום ברמת הבלוק בלבד - הוא מסנכרן רק את threads אותו הבלוק, שכולם על אותו SM. אין לו שום דרך לתאם בין בלוקים שיושבים על SMs שונים, ולכן הוא לא יכול לצרף תוצאות חלקיות מבלוקים שונים. הדרך היחידה לתקשורת חוצה-בלוקים היא global memory עם atomics - בדיוק מה שעושה atomicAdd. שימו לב: הגרסה הזו (atomic לכל איבר) נכונה אבל איטית בגלל התנגשות על כתובת אחת; את הגרסה היעילה תבנו בבונוס.

תרגיל 5 - חישוב occupancy ביד ואימות

נחשב כמה בלוקים נכנסים ל-SM, ידנית ואוטומטית.

1. קמפלו את אחד ה-kernels שלכם עם דגל שמדפיס שימוש במשאבים:

nvcc -O2 -arch=sm_90a --ptxas-options=-v -o hier hier.cu

רשמו את מספר האוגרים לthread (registers) ואת ה-shared memory לבלוק שה-ptxas מדווח.

1. עבור blockSize = 256 והמספרים שקיבלתם, חשבו ביד על H100 (65,536 אוגרים/SM, 64 warps/SM, מגבלת 32 בלוקים/SM) כמה בלוקים תושבים יאפשר כל אחד מהמשאבים, ומהו המינימום (הגורם המגביל).
2. אמתו את החישוב עם ה-CUDA Runtime:

int maxBlocks = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &maxBlocks, sumAtomic, 256, 0));   // kernel, blockSize, dynamicSharedMem
printf("max active blocks / SM = %d\n", maxBlocks);

1. חשבו את התפוסה באחוזים: (maxBlocks * 256) / 2048 * 100. האם היא 100%? אם לא, מהו המשאב המגביל?

רמז: הנוסחה לכל משאב היא רצפה(תקציב / צריכה לבלוק). לאוגרים: 65536 / (regsPerThread * 256). ל-warp slots: 64 / (256/32) = 64/8 = 8. אם ה-kernel שלכם קטן (מעט אוגרים, בלי shared), הגורם המגביל יהיה חריצי ה-warp או מגבלת הבלוקים, ותקבלו 8 בלוקים = תפוסה מלאה. אם תוסיפו __launch_bounds__ או תעלו את צריכת האוגרים, תראו את המספר יורד.

תרגיל 6 (בונוס) - רדוקציה בשני שלבים: shared memory ואז atomic

נשפר את סכום ה-grid מתרגיל 4: כל בלוק מסכם את חלקו ב-shared memory (עם __syncthreads), ורק thread אחד לבלוק מוסיף את הסכום החלקי ל-global עם atomicAdd. זה משלב את שתי רמות שיתוף הפעולה מההרצאה.

1. כתבו kernel sumReduce שבו כל בלוק:
2. טוען איבר לכל thread ל-__shared__ float sdata[BLOCK].
3. מבצע __syncthreads().
4. מבצע רדוקציה בעץ בתוך ה-shared (חצי מהthreads מחברים בכל שלב, עם __syncthreads() בין השלבים).
5. הthread threadIdx.x == 0 מבצע atomicAdd(result, sdata[0]).
6. השוו את התוצאה ואת הזמן (עם cudaEvent) לגרסת ה-atomic-לכל-איבר מתרגיל 4.
7. הסבירו: היכן השתמשתם בסנכרון ברמת הבלוק (__syncthreads) והיכן בסנכרון חוצה-בלוקים (atomicAdd), ולמה כל אחד במקומו.

רמז: תבנית הרדוקציה בעץ היא for (int s = blockDim.x/2; s > 0; s >>= 1) { if (tid < s) sdata[tid] += sdata[tid+s]; __syncthreads(); }. ה-__syncthreads() בתוך הלולאה קריטי - בלעדיו thread עלול לקרוא sdata[tid+s] לפני שthread אחר סיים לכתוב אליו (data race). כל בלוק מבצע atomicAdd אחד בלבד במקום 256, ולכן ההתנגשות על ה-global קטנה פי 256 - זו הסיבה שהגרסה הזו מהירה בהרבה. את התבנית המלאה הזו נפתח לעומק בפרק על shared memory.