8.4 Coalescing ו bank conflicts תרגול

תרגול - Coalescing ו-bank conflicts

בתרגול הזה תרגישו בידיים את שתי אופטימיזציות דפוס-הגישה הגדולות. תבנו את מדד ה-stride ותשחזרו את קריסת ה-coalescing מ-206 GB/s ל-~11 GB/s; תתקנו kernel לא מאוחד על ידי היפוך תפקידי thread ו-loop, ותמדדו את הרווח; תכתבו transpose נאיבי עם התנגשות בנקים 32-way, תאשרו אותה עם ncu, תרפדו ל-[32][33] ותמדדו את הזינוק; ולבסוף תסבירו במדויק מדוע אופטימיזציה אחת שייכת ל-global memory והשנייה ל-shared memory. עבדו לפי הסדר - כל תרגיל בונה על הקודם. אפשר לעבוד על כל כרטיס NVIDIA מקומי או בחינם ב-Google Colab. אם אתם ב-Colab על T4, החליפו -arch=sm_90 ב--arch=sm_75, הוסיפו ! לפני כל פקודת shell, והשתמשו ברוחב הפס של ה-T4 (peak ~320 GB/s של GDDR6) כמכנה במקום ה-HBM3 של ה-H100 - העיקרון זהה, המספרים שונים.

הכנה - כלים ומאקרו

לכל התרגילים נשתמש במאקרו בדיקת השגיאות הסטנדרטי; שימו אותו בראש כל קובץ .cu:

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

#define CUDA_CHECK(call)                                                   \
    do {                                                                   \
        cudaError_t err_ = (call);                                         \
        if (err_ != cudaSuccess) {                                         \
            fprintf(stderr, "CUDA error %s at %s:%d\n",                    \
                    cudaGetErrorString(err_), __FILE__, __LINE__);         \
            exit(EXIT_FAILURE);                                            \
        }                                                                  \
    } while (0)

קמפלו תמיד עם ארכיטקטורה ואופטימיזציה, למשל nvcc -arch=sm_90 -O3 file.cu -o file. לפרופיילינג צריך את ncu (Nsight Compute), שמגיע עם ה-CUDA Toolkit.

תרגיל 1 - בניית מדד ה-stride ושחזור קריסת ה-coalescing

נבנה את ה-benchmark שמודד רוחב פס מושג כפונקציה של ה-stride, ונשחזר את הטבלה מהשיעור.

1. צרו strided.cu עם ה-kernel strided_read_kernel מהשיעור (קריאה עם stride, הצטברות ל-acc, כתיבה ל-out כדי למנוע מחיקה).
2. ב-main, הקצו מערך של N = 64M floats (64ull << 20), אתחלו אותו, והריצו חימום (warmup) אחד לפני המדידה.
3. לכל stride בקבוצה {1, 2, 4, 8, 16, 32, 64, 128}: מדדו את הזמן עם cudaEvent על פני כמה איטרציות, חשבו רוחב פס מושג, והדפיסו שורה.
4. ספירת הבייטים: ה-kernel קורא בערך N / stride אלמנטים (הלולאה חסומה ב-j < N בצעדי stride). לכן הבייטים השימושיים הם (N / stride) * sizeof(float). חלקו בזמן ל-GB/s.
5. הריצו והשוו לטבלה בשיעור. האם קיבלתם את הצורה: חצייה בכל הכפלה עד stride 8, ואז השטחה מ-stride 16?

רמז: הבייטים השימושיים הם (double)N / stride * sizeof(float). השתמשו ב-grid = 4096, block = 256. אם אתם על H100 המספרים המוחלטים יהיו גבוהים בהרבה מ-T4 (stride 1 קרוב ל-2800-3000 GB/s), אבל הצורה של הקריסה זהה. חממו לפני המדידה ומדדו על פני 20-50 איטרציות.

תרגיל 2 - אישור ה-coalescing עם ncu (sectors per request)

נאשר שהקריסה היא באמת מ-coalescing ולא ממשהו אחר, דרך מדד ה-sectors-per-request.

1. הריצו את strided (עם stride קבוע 1) תחת ncu עם המדד l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio.
2. הריצו שוב עם stride 32. השוו את המדד בין שתי ההרצות.
3. הסבירו: כמה sectors לבקשה מצפים לקבל בגישה מאוחדת (stride 1), וכמה בגישה מפוזרת (stride 32)? קשרו למספר 4 (128B = 4 sectors של 32B) ולמקסימום 32.
4. איזה מהשניים - ה-sectors-per-request או ה-GB/s מתרגיל 1 - הוא מדד "ישיר" יותר של coalescing, ולמה?

רמז: קריאה מאוחדת של warp מביאה 128B = 4 sectors של 32B לבקשה. קריאה שבה כל thread נופל ב-sector אחר תביא עד 32. כדי לקבע את ה-stride, קמפלו גרסה עם stride קשיח, או קבלו אותו כארגומנט והריצו ./strided 1 ו-./strided 32.

תרגיל 3 - תיקון kernel לא מאוחד ומדידת הרווח

נכתוב kernel לסיכום שורות מטריצה בשתי גרסאות - לא מאוחדת ומאוחדת - ונמדוד את ההבדל.

1. צרו rowsum.cu עם מטריצה M x N (למשל M = 8192, N = 8192) שמורה row-major ב-global memory.
2. ממשו את rowsum_bad: כל thread לוקח שורה שלמה (for col: acc += A[row*N+col]). זכרו שזה נותן ל-threads שכנים כתובות במרחק N*4 בייט.
3. ממשו את rowsum_good: block אחד לשורה, ו-threadIdx.x רץ על העמודות בצעדי blockDim.x, כך ש-threads שכנים קוראים כתובות רצופות. (לצמצום בתוך ה-block אפשר להשתמש ב-atomicAdd ל-s[row] או בצמצום shared פשוט; ההתמקדות היא בדפוס הקריאה מ-global.)
4. מדדו את שתי הגרסאות עם cudaEvent והדפיסו את היחס (speedup).
5. הריצו את שתיהן תחת ncu עם l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio. איזו גרסה מראה ~4 ואיזו מראה ~32? כיצד זה מסביר את ה-speedup?

רמז: הפי-כמה-שיפור צריך לנחות בטווח פי 8 עד פי 20 - בדיוק היחס בין stride 1 ל-stride גדול בטבלה. הכלל: ה-threadIdx.x חייב לרוץ על הממד הרציף בזיכרון (העמודה, ב-row-major). אמתו את שתי הגרסאות שמחזירות אותה תוצאה מספרית לפני שמשווים זמנים.

תרגיל 4 - transpose נאיבי והתנגשות בנקים 32-way

נכתוב transpose דרך shared memory ונאשר את ההתנגשות 32-way.

1. צרו transpose.cu עם transposeNaive מהשיעור: __shared__ float tile[32][32], קריאה מאוחדת מ-global אל tile[threadIdx.y][threadIdx.x], __syncthreads(), וכתיבה מ-tile[threadIdx.x][threadIdx.y] ל-global.
2. הריצו על מטריצה 4096 x 4096, ואמתו נכונות מול transpose על ה-CPU (או מול העתקה out[j][i] = in[i][j]).
3. חשבו ביד את הבנק שכל thread ב-warp נוגע בו בקריאה tile[threadIdx.x][threadIdx.y]. הראו שכולם נופלים על בנק threadIdx.y (קבוע ל-warp) -> התנגשות 32-way.
4. הריצו תחת ncu עם l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum. כמה התנגשויות אתם רואים? מדדו גם את הזמן עם cudaEvent.

רמז: הכתובת (במילים) של tile[threadIdx.x][threadIdx.y] היא threadIdx.x * 32 + threadIdx.y, והבנק הוא הביטוי mod 32. שימו לב ש-(threadIdx.x * 32) mod 32 = 0 תמיד, ולכן נשאר רק threadIdx.y. השתמשו ב-block של dim3(32, 8) או dim3(32, 32).

תרגיל 5 - ריפוד ל-[32][33] ומדידת השיפור

נתקן את ההתנגשות בשינוי של תו אחד, ונמדוד.

1. שנו את ההצהרה ל-__shared__ float tile[32][33] (עמודת ריפוד אחת). אל תשנו שום דבר אחר בקוד.
2. חשבו ביד מחדש את הבנק לקריאה tile[threadIdx.x][threadIdx.y] עם stride שורה 33. הראו שהוא הופך ל-(threadIdx.x + threadIdx.y) mod 32, שעובר על כל 32 הבנקים -> אפס התנגשויות.
3. הריצו תחת ncu עם אותו מדד l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum. אמתו שהמונה צנח כמעט לאפס.
4. מדדו את הזמן והשוו ל-transposeNaive מתרגיל 4. מה ה-speedup? כמה shared memory נוסף "עלה" הריפוד (באחוזים)?

רמז: 33 = 32 + 1, ולכן (threadIdx.x * 33) mod 32 = (threadIdx.x * 32 + threadIdx.x) mod 32 = threadIdx.x. הריפוד עולה 33/32 = ~3% מקום נוסף ב-shared וקונה עד פי 32 בקריאה. אמתו שהתוצאה המספרית עדיין נכונה אחרי הריפוד - הריפוד לא אמור לשנות את הפלט, רק את הביצועים.

תרגיל 6 - שידור מול התנגשות (broadcast)

נחדד את התנאי "כתובות שונות" של התנגשות בנקים.

1. כתבו kernel קטן שבו כל ה-warp קורא data[0] (אותה כתובת בדיוק) לתוך משתנה מקומי. חשבו: כל ה-threads באותו בנק (בנק 0), אבל אותה כתובת. האם זו התנגשות?
2. הריצו אותו תחת ncu עם מדד ה-bank conflicts. כמה התנגשויות? הסבירו את התוצאה במונח "שידור (broadcast)".
3. עכשיו שנו ל-data[threadIdx.x % 1]... לא, שנו ל-data[(threadIdx.x % 2) * 32] - כלומר חצי מה-threads קוראים מילה 0 וחצי מילה 32, שתיהן בבנק 0 אך בכתובות שונות. איזו דרגת התנגשות מצפים (2-way? 32-way?)? אמתו עם ncu.
4. סכמו בטבלה קטנה את שלושת המצבים: בנקים שונים, אותו בנק+אותה כתובת, אותו בנק+כתובות שונות - ומה הקנס בכל אחד.

רמז: שידור קורה רק כשהכתובת זהה; אז אין קנס. במקרה (threadIdx.x % 2) * 32 יש שתי כתובות שונות (0 ו-32) שתיהן בבנק 0, אז זו התנגשות 2-way - אבל שימו לב שהחומרה עשויה לשדר בתוך כל קבוצת-אותה-כתובת, כך שהקנס הוא כמספר הכתובות המובחנות בבנק, לא כמספר ה-threads.

תרגיל 7 (בונוס) - למה אחת ל-global והשנייה ל-shared

תרגיל הסבר שמחבר את הכל.

1. הסבירו בפסקה מדוע coalescing לא רלוונטי כלל ל-shared memory. רמז: מה מקור ה-128B transaction, והאם ל-shared memory יש bursts של DRAM?
2. הסבירו מדוע bank conflicts לא רלוונטיים כלל ל-global memory. רמז: היכן יושבים 32 הבנקים, ומה מגבה את ה-global memory?
3. בנו טבלת 2x2: (coalescing, bank conflict) מול (global, shared), וסמנו איפה כל אחד חי ואיפה הוא חסר משמעות.
4. תרחיש: kernel שמעתיק מ-global אל shared, מעבד ב-shared, וכותב חזרה ל-global. באיזה שלב אתם דואגים ל-coalescing ובאיזה שלב ל-bank conflicts?

רמז: ה-128B burst הוא תכונה של ה-DRAM שמגבה את ה-global memory; ל-shared (SRAM על השבב) אין bursts, יש בנקים. ההפך: הבנקים הם ארגון פיזי של ה-SRAM המשותף; ל-global אין בנקים. בתרחיש: coalescing בשלבי הקריאה/כתיבה ל-global, bank conflicts בשלב העיבוד ב-shared.