8.3 Warp divergence ו scoreboard stalls פתרון
פתרון - Warp divergence ו-scoreboard stalls¶
הערה: כתובות האוגרים, מספרי ה-scoreboards, ערכי ה-stall וכל הפלטים המספריים למטה הם דוגמאות ממכונת H100 SXM עם CUDA 12, ועשויים להשתנות אצלכם לפי הכרטיס, גרסת הדרייבר וגרסת ה-toolkit (למשל אילו scoreboards ptxas בחר, או סדר הפקודות). מה שלא ישתנה הוא המבנה, סוגי הפקודות, ותבנית המחסום. תרגילי החשבון (3) מדויקים ואינם תלויים בחומרה. כל קטעי הקוד מתהדרים עם nvcc -O2 -arch=sm_90a (החליפו ל-sm_75 על T4 וכדומה). כדי לפרפל עם ncu ייתכן שתצטרכו הרשאות (sudo) או --target-processes all.
פתרון תרגיל 1 - לקרוא את ה-SASS של ההתפצלות¶
nvcc -O2 -arch=sm_90a -o divergence divergence.cu
cuobjdump --dump-sass ./divergence | grep -A2 -E 'FSETP|FADD|FMUL'
הגרעין הרלוונטי של ה-SASS (מנוקה):
LDG.E R4, [R2.64] // load of data[idx]
FSETP.GT.AND P0, PT, R4.reuse, 0.5, PT // P0 = (data[idx] > 0.5)
FADD R0, R4, 2 // Block B (else): executed for every thread
@P0 FMUL R0, R4, 4 // Block A (if): conditioned on P0
FMUL R5, R0, R0 // convergence: data[idx]*data[idx]
STG.E [R2.64], R5 // store
התשובות:
- סעיף 2:
FSETP.GT.ANDקובעת את ה-predicateP0לפי התנאיR4 > 0.5(כלומרdata[idx] > 0.5f). זהו ה-active mask של מסלול A. - סעיף 3: ה-
FADD R0, R4, 2היא Block B (ה-else,data+2) ומבוצעת עבור כל 32 ה-threads - אין לה@P0. ה-@P0 FMUL R0, R4, 4היא Block A (ה-if,data*4) ומבוצעת רק עבור ה-threads שבהםP0דלוק. שימו לב ששתיהן קוראות מ-R4(הערך המקורי) וכותבות ל-R0, כך שה-FMUL דורסת את תוצאת ה-FADD עבור ה-threads שלה. - סעיף 4: פקודת ההתכנסות היא
FMUL R5, R0, R0- הריבועdata[idx]*data[idx]. מכאן כל 32 ה-threads שוב מנפיקים יחד. - סעיף 5: הcompiler בחר להריץ צד אחד (FADD) לכולם ולדרוס את הצד השני עם predicate יחיד, במקום predication סימטרי על שני הצדדים. FADD היא פקודת מחזור-בודד זולה על CUDA Core; הcompiler מעדיף לבזבז אותה מאשר להוסיף predicate שני ומורכבות בקרת-זרימה נוספת.
למה זה עבד: כך נראית התפצלות אמיתית ב-SASS - FSETP שקובע מסכה, פקודה לא-מותנית שרצה לכולם, פקודה מותנית @P0 שדורסת חלקית, ואז פקודה משותפת שמסמנת התכנסות. זיהוי התבנית הזו הוא כל מה שצריך כדי לאבחן התפצלות בקוד לא-מוכר.
איך להכליל: בכל if/else תלוי-נתונים חפשו את השלישייה SETP -> פקודה-לכולם -> @P<n> פקודה-מותנית. אם אתם רואים predication סימטרי (@P0 ... ו-@!P0 ...) זה קורה כששני הצדדים יקרים מדי מכדי להריץ את שניהם לכולם. פקודות "מבוזבזות" הן חלק נורמלי מהאסטרטגיה.
פתרון תרגיל 2 - למדוד branch efficiency והתפצלות ב-ncu¶
int main(void) {
const int N = 1 << 20;
float* h = (float*)malloc(N * sizeof(float));
for (int i = 0; i < N; i++) h[i] = (float)rand() / (float)RAND_MAX; // [0,1)
float* d;
CUDA_CHECK(cudaMalloc(&d, N * sizeof(float)));
CUDA_CHECK(cudaMemcpy(d, h, N * sizeof(float), cudaMemcpyHostToDevice));
divergent_kernel<<<(N + 255) / 256, 256>>>(d, N);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaFree(d));
free(h);
return 0;
}
ncu --metrics smsp__sass_average_branch_targets_threads_uniform.pct ./divergence
ncu --metrics smsp__thread_inst_executed_per_inst_executed.ratio ./divergence
הפלט הצפוי (מקוצר):
smsp__sass_average_branch_targets_threads_uniform.pct % 50.02
smsp__thread_inst_executed_per_inst_executed.ratio 16.01
התשובות (סעיף 4):
- ה-branch efficiency ~50%. ל-kernel יש שני ענפים: בדיקת-הגבולות החיצונית
idx < n(אחידה כמעט תמיד - כל ה-threads בפנים) והתנאי הפנימיdata > 0.5(מתפצל כמעט תמיד על נתונים אקראיים). מכיוון שהמדד הוא ממוצע על שני הענפים - אחד אחיד ואחד מתפצל - התוצאה נופלת סביב 50%, לא 0%. - כ-16 threads פעילים מתוך 32. בתוך התנאי הפנימי המתפצל, בכל warp בערך חצי מה-lanes לוקחים כל מסלול, ולכן ממוצע ה-threads הפעילים לפקודה בגוף הענף הוא ~16.
למה זה עבד: הנתונים האקראיים האחידים מבטיחים ש-P(all 32 agree) = 2 x 0.5^32 - אפסי - ולכן כמעט כל warp מתפצל בתנאי הפנימי. המדד .pct סופר את שיעור ההחלטות האחידות; המדד .ratio סופר threads פעילים לפקודה. שני המספרים מספרים את אותו סיפור מזוויות שונות.
איך להכליל: smsp__sass_average_branch_targets_threads_uniform.pct הוא ה-branch efficiency שלכם בכל kernel. אם הוא נמוך, השתמשו ב-.ratio של active threads כדי לכמת כמה עבודה SIMT מתבזבזת, ואז חזרו ל-SASS כדי לראות איזה ענף אשם. זכרו שהמדד ממצע על כל הענפים - ענף אחיד אחד יכול "להסוות" ענף מתפצל.
פתרון תרגיל 3 - להוכיח שבדיקת-הגבולות כמעט חינמית¶
nvcc -O2 -arch=sm_90a -o bounds_only bounds_only.cu
ncu --metrics smsp__sass_average_branch_targets_threads_uniform.pct ./bounds_only
הפלט הצפוי:
החישוב הידני (סעיף 4), עם n = 1,000,003, blockDim = 256:
gridDim = ceil(1,000,003 / 256) = 3907 blocks
threads = 3907 x 256 = 1,000,192
warps = 1,000,192 / 32 = 31,256
uniform warps (all idx<n or all idx>=n) : 31,255
divergent warps (riding the idx=1,000,003 boundary): 1
branch efficiency = 31,255 / 31,256 = 0.999968 ≈ 99.997%
מסקנה (סעיף 5): לא. עצם קיומו של if ב-kernel אינו גורר התפצלות. התפצלות קורית רק כשה-threads בתוך אותו warp לא מסכימים על התנאי. בבדיקת-גבולות, idx עולה ברציפות, ולכן כמעט כל warp אחיד לגמרי, ורק ה-warp הבודד שרוכב על הגבול מתפצל.
למה זה עבד: idx הוא פונקציה מונוטונית של מספר ה-thread, ו-32 ה-threads בכל warp הם בעלי idx עוקבים. לכן idx < n אחיד בתוך כל warp למעט זה שגבולו נופל באמצע. עם n שהוא כפולה של 32 בדיוק, אפילו ה-warp הבודד הזה לא מתפצל (הגבול מיושר לתחילת warp) והיעילות היא 100% חלק - לכן בחרנו n שאיננו כפולה של 32.
איך להכליל: תבניות גישה מונוטוניות (idx, idx < n, idx % 2 על-פני grid רציף) שומרות על אחידות מרחבית ולכן על branch efficiency גבוהה. התפצלות יקרה מגיעה מתנאים תלויי-ערך (data[idx] > thr) שבהם השכנים המרחביים אינם מתואמים. לפני שאתם "מסירים if", בדקו אם הוא בכלל מפצל את ה-warp.
פתרון תרגיל 4 - long scoreboard מול short scoreboard¶
__global__ void mem_bound(const float* in, float* out, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) out[i] = in[i]; // LDG then STG, almost no arithmetic
}
__global__ void transcendental(const float* in, float* out, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
float x = in[i];
for (int k = 0; k < 64; k++)
x = sinf(x) + sqrtf(fabsf(x)) + expf(x * 0.01f); // heavy SFU
out[i] = x;
}
}
nvcc -O2 -arch=sm_90a -o stalls stalls.cu
ncu --section WarpStateStats -k mem_bound ./stalls
ncu --section WarpStateStats -k transcendental ./stalls
הפלט הצפוי (השורות הרלוונטיות מ-"Warp State Statistics", מחזורי-עצירה ממוצעים לפקודה):
== mem_bound ==
Stall Long Scoreboard cycles/inst 8.71 <-- dominant
Stall Short Scoreboard cycles/inst 0.12
Stall Wait cycles/inst 0.30
== transcendental ==
Stall Long Scoreboard cycles/inst 0.09
Stall Short Scoreboard cycles/inst 6.44 <-- dominant
Stall Wait cycles/inst 0.51
זיהוי פקודות ה-SFU (סעיף 4):
התשובות: עבור mem_bound הסיבה הדומיננטית היא Stall Long Scoreboard - ה-warps תקועים בהמתנה ל-LDG שיחזור מ-HBM. עבור transcendental הסיבה הדומיננטית היא Stall Short Scoreboard - ה-warps תקועים בהמתנה לפקודות ה-MUFU (SFU) שמחשבות את הטרנסצנדנטלים על ה-SM.
למה זה עבד: LDG/STG יוצאים מה-SM אל היררכיית הזיכרון הגלובלי, ולכן ההמתנה להם נספרת כ-long scoreboard - וזה בדיוק מה ששולט בקוד memory-bound, כפי שה-glossary מדגיש. MUFU.SIN/RSQ/EX2 הן פעולות SFU תוך-SM עם latency משתנה, ולכן ההמתנה להן נספרת כ-short scoreboard. בידדנו כל סוג בכך שבנינו kernel שכל-כולו סוג עבודה אחד.
איך להכליל: כשאתם מאבחנים kernel איטי, ncu --section WarpStateStats נותן את פירוק הסיבות. Stall Long Scoreboard גבוה = memory-bound, תקנו עם coalescing/reuse/tiling (פרק 3). Stall Short Scoreboard גבוה = ממתינים ל-SFU/MMA/shared memory, תקנו עם פחות טרנסצנדנטלים, precision נמוך יותר, או פחות תלות. תמיד אמתו מול ה-SASS איזו פקודה יושבת על ה-source line החם.
פתרון תרגיל 5 - לקרוא את סימון המחסום על זוג load-then-use¶
ה-SASS עם סימון המחסום (מנוקה; ה-scoreboards הספציפיים עשויים להשתנות):
[B------:R-:W2:-:S04] /*00f0*/ LDG.E R0, [R2.64] ; // load a[i], marks scoreboard 2
[B------:R-:W2:-:S01] /*0100*/ LDG.E R5, [R4.64] ; // load b[i], marks scoreboard 2
[B--2---:R-:W-:Y:S08] /*0150*/ IMAD R0, R0, R5, RZ ; // a*b, blocks on scoreboard 2
[B------:R-:W-:-:S02] /*0160*/ STG.E [R6.64], R0 ; // store c[i]
התשובות:
- סעיף 3-4: כל אחד משני ה-
LDGנושאW2בשדה ה-write - הוא מסמן (מגדיל) את scoreboard מספר 2, כלומר "התוצאה שלי בטיסה, עקבו אחריה ב-2". פקודת ה-IMAD(הכפלa*b) נושאתB--2---בשדה ה-barrier - היא חוסמת על scoreboard 2 וחייבת להמתין עד שהוא יתנקה. כך המחסום מבטיח שהכפל לא יונפק לפני ששני ה-loads חזרו - אחרת הוא היה קורא R0/R5 עם ערכי-זבל. - סעיף 5: מכיוון ששני ה-
LDGהשתמשו באותו scoreboard (2), חסימה בודדת על scoreboard 2 ממתינה לשניהם - ה-scoreboard נספר למעלה פעמיים (פעם לכל load) ומתנקה רק כששניהם הושלמו. שדה ה-S(למשלS08) הוא מספר מחזורי-העצירה הסטטי שהמתזמן מחזיק על הפקודה, ו-Yב-IMAD הוא רמז-yield שמעודד את המתזמן להחליף ל-warp אחר בזמן ההמתנה.
למה זה עבד: המחסום מקודד את גרף התלויות ישירות בזרם הפקודות. שדה ה-W רושם "מי מייצר ערך שעדיין בטיסה", שדה ה-B רושם "מי צריך להמתין למי". זוג load-then-use הוא התבנית הנפוצה ביותר: מייצר (LDG, W) וצרכן (IMAD, B על אותו scoreboard). זו בדיוק ה-long scoreboard stall שראינו ב-ncu, רק ברזולוציית-הפקודה.
איך להכליל: כדי לאתר את מקור עצירת ה-long scoreboard, שלפו את ה-SASS עם nvdisasm -c או cuobjdump --dump-sass, מצאו את הפקודה עם שדה B לא-ריק, וקִראו איזה scoreboard היא ממתינה עליו; אז מצאו את הפקודה הקודמת עם אותו scoreboard בשדה ה-W - זו הפקודה שהיא מחכה לה (כמעט תמיד LDG). המרחק בין השתיים הוא ה-latency שאתם מנסים להסתיר עם occupancy.
פתרון תרגיל 6 (בונוס) - מלכודת מיחזור ה-scoreboards¶
ncu --set full -k transcendental ./stalls
cuobjdump --dump-sass ./stalls | grep -E '\[B' | grep -oE 'W[0-9]' | sort | uniq -c
הפלט (דוגמה - מונה כמה פקודות מסמנות כל scoreboard):
התשובות:
- סעיף 1-2: ב-kernel שמערבב
LDG(long) ו-MUFU(short), ncu לא תמיד מפריד את שתי סיבות העצירה נקי. כפי שרואים בפלט, הcompiler משתמש רק בקומץ scoreboards שונים (כאן 0, 1, 2 מתוך 6 האפשריים) וממחזר אותם לאורך ה-kernel. - סעיף 3: אם ptxas הקצה במקרה את אותו scoreboard גם ל-
LDG(long) וגם ל-MUFU.EX2(short) - למשל שניהםW2- אז כשה-warp חוסם עלB--2---, החומרה (וה-sampler של ncu) לא יכולה לדעת בוודאות אם ההמתנה היא ל-load מהזיכרון או ל-SFU. ה-glossary מזהיר בדיוק על זה: עצירת long ועצירת short עלולות להתערבב אם הן חולקות scoreboard, ולכן התווית שncu נותנת עלולה להיות הפוכה. - סעיף 4 - כלל אצבע: סמכו על תווית ה-short/long של ncu כשהיא עקבית עם אופי ה-kernel (kernel של העתקה טהורה -> long; kernel של טרנסצנדנטלים טהורים -> short). כשה-kernel מעורב או כשהמספרים מפתיעים (למשל "short scoreboard" גבוה ב-kernel שכמעט ולא נוגע ב-SFU) - חזרו ל-SASS:
cuobjdump --dump-sass, אתרו את הפקודה על ה-source line החם, וקִראו ישירות את שדה ה-B ואת הפקודה שסימנה את אותו scoreboard בשדה ה-W. הפקודה עצמה (LDG לעומת MUFU) חד-משמעית; התווית של ncu היא הערכה.
למה זה עבד: יש רק 6 scoreboards ל-warp, ולכן בקוד עמוס הcompiler חייב למחזר אותם. המיחזור הוא שיוצר את העמימות: scoreboard יחיד יכול לעקוב ברצף אחרי load ואחרי MUFU, וה-classifier של ncu מייחס את העצירה לפי ניחוש. ה-SASS, לעומת זאת, מראה את הפקודה המדויקת - ולכן הוא מקור-האמת.
איך להכליל: אל תשכתבו kernel על סמך תווית stall בודדת. אמתו כל אבחון של scoreboard מול ה-SASS: זהו את שדה ה-B החוסם, עקבו אל שדה ה-W שסימן, וראו איזו פקודה באמת יצרה את התלות. זו ההרגל שמפריד בין אופטימיזציה מנחשת לבין אופטימיזציה מבוססת-ראיות.