לדלג לתוכן

8.3 Warp divergence ו scoreboard stalls הרצאה

בשיעור 2.1 הכרנו את מודל ה-SIMT (Single Instruction, Multiple Threads) - 32 threads שמאוגדים ל-warp ומריצים פקודה אחת יחד בכל מחזור. בפרק 4 למדנו לקרוא SASS, האסמבלי האמיתי של החומרה, ובשיעורים 8.1 ו-8.2 ראינו כיצד occupancy (תפוסה) ו-warp state (מצב ה-warp) קובעים כמה עבודה מקבילה עומדת לרשות מתזמן ה-warps כדי להסתיר latency. עכשיו נרד רזולוציה אחת עמוק יותר ונשאל: גם כשיש לנו מספיק warps תושבים, מה גורם ל-warp בודד להפוך ללא-כשיר (not eligible) ולעכב את ההנפקה. בשיעור הזה נתמקד בשתי מכשלות שקוראים אותן ישירות ב-SASS - שתיהן חיות בדיוק בצומת שבין בקרת-הזרימה לבין תזמון ההנפקה. הראשונה היא warp divergence (התפצלות ה-warp) - מה קורה כש-32 ה-threads בתוך warp רוצים ללכת בכיוונים שונים של תנאי; המדד המכמת אותה הוא branch efficiency (יעילות ההסתעפות). השנייה היא scoreboard stall (עצירת לוח-התוצאות) - מה קורה כש-warp חייב להמתין לתוצאה של פקודה קודמת שעוד "בטיסה". נראה את שתיהן על SASS קונקרטי של H100, נבין את האסטרטגיה הלא-אינטואיטיבית של הcompiler בטיפול בהתפצלות, ונלמד לפענח את סימון ה-scoreboard בשורת פקודה - כלי אבחון שמלווה אתכם בכל פרופיל של ncu.

מודל ה-SIMT וההתפצלות - warp divergence

נזכיר את הליבה של מודל ה-SIMT מ-2.1: ה-warp הוא יחידת ההנפקה. מתזמן ה-warp בוחר warp כשיר, שולף עבורו פקודה אחת, וכל 32 ה-threads מריצים אותה על הנתונים הפרטיים שלהם באותו מחזור-שעון. כל עוד כל 32 ה-threads מסכימים על מסלול ריצה יחיד - זהו המצב האידיאלי, וה-SIMT מנצל את החומרה במלואה: פענוח אחד, הנפקה אחת, 32 תוצאות.

הבעיה מתחילה כשמופיע תנאי שתלוי בנתונים הפרטיים של כל thread. ה-glossary מגדיר זאת בחדות: warp divergence מתרחשת כאשר threads בתוך warp פוסעים במסלולי ריצה שונים בעקבות פקודת בקרת-זרימה. שימו לב לניסוח - זו איננה תופעה שאפשר "לתקן" בקוד. אם ההחלטה בענף תלויה בערך נתון פר-thread, אז לפי הגדרה אין שום דרך - לא למתכנת ולא לcompiler - למנוע את הפיצול. זו התפצלות תלוית-נתונים ובלתי-נמנעת מעצם בנייתה.

בואו נראה את הדוגמה הקנונית. ה-kernel הבא מפצל את ה-threads לפי הערך שהם קראו מהזיכרון:

__global__ void divergent_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        if (data[idx] > 0.5f) {
            data[idx] = data[idx] * 4.0f;   // Block A ("if")
        } else {
            data[idx] = data[idx] + 2.0f;   // Block B ("else")
        }
        data[idx] = data[idx] * data[idx];  // convergence point
    }
}

אם ה-data הוא רעש אקראי אחיד ב-[0,1), אז בתוך כל warp בערך חצי מה-threads יראו ערך גדול מ-0.5 (מסלול A) והחצי השני ערך קטן ממנו (מסלול B). ה-warp מתפצל: אין פקודה יחידה שמתאימה לכל 32 ה-threads בו-זמנית. איך החומרה מתמודדת עם זה? כאן נכנס מנגנון ה-predicate registers (אוגרי חיווי) וה-active mask (מסכת פעילות), ולזה נקדיש את הסעיף אחרי הבא. קודם - למה בכלל זה כואב.

לפני Volta ואחריה - סריאליזציה מול תזמון עצמאי

עלות ההתפצלות השתנתה מהותית עם דור Volta (2017), וזו נקודה שחשוב לדייק בה כי סביבה נוצרו הרבה אמונות מיושנות.

לפני Volta. ב-GPUs מדור Pascal ומטה, warp שהתפצל היה תמיד מסוריאל במלואו (always fully serialized). המשמעות: החומרה הריצה קודם את כל ה-threads של מסלול A כשמסכת ה-threads של מסלול B כבויה, ואז את כל ה-threads של מסלול B כשמסכת A כבויה. שני המסלולים רצו בטור, אחד אחרי השני, ולכן הזמן היה סכום שני המסלולים במקום המקסימום שלהם. מכאן נולד כלל-האצבע המפחיד "התפצלות = פי 2 האטה" (ובמקרה הגרוע, if/else if עמוק, פי 32).

Volta ואילך. מ-Volta והלאה (וכך גם ב-Turing, Ampere, Ada, Hopper וה-H100 שלנו, וכן ב-Blackwell) החומרה תומכת ב-independent thread scheduling (תזמון threads עצמאי): לכל thread יש מונה-תוכנית (program counter) ומחסנית-קריאה משלו, ולא רק מונה-תוכנית אחד לכל ה-warp. המשמעות המעשית, בלשון ה-glossary: התפצלות אינה גוררת בהכרח את קנס הסריאליזציה המלא. החומרה יכולה לשזור התקדמות של המסלולים ולהתכנס (reconverge) בצורה גמישה יותר, ובפרט warps שהתפצלו יכולים לחפוף זה את ה-latency של זה.

   divergence into if/else, one warp:

   Pre-Volta (full serialization):
      |==== path A (16 threads) ====|==== path B (16 threads) ====|
      \_____________________ time = A + B _____________________/

   Volta+ (independent thread scheduling):
      |==== path A ====|
      |==== path B ====|   (can interleave; reconvergence is flexible)
      \___ time ~ max(A,B), and latency is interleaved with other warps ___/

המסקנה החשובה: "התפצלות = האטה אוטומטית פי 32" היא תפיסה מוטעית מעידן ה-pre-Volta. על H100 התפצלות עדיין עולה - היא מבזבזת מחזורי הנפקה ומקטינה את מספר ה-threads הפעילים לפקודה - אבל היא איננה גזירת-דין של סריאליזציה מלאה. עדיין, פחות התפצלות זה תמיד טוב יותר, ולכן נרצה למדוד אותה.

איך הcompiler מטפל בהתפצלות - predicates ו-active mask

עכשיו נפתח את הקופסה השחורה ונראה מה ptxas באמת פולט. נהדר את ה-kernel מהסעיף הראשון ל-SASS של H100 ונביט בגוף התנאי הפנימי (הפלט מנוקה מעט לצורך קריאוּת; אפשר לשחזר ב-Godbolt, למשל https://godbolt.org/z/EGWKb5oWr):

L1: LDG.E.SYS R4, [R2]                       // load of data[idx]
L2: FSETP.GT.AND P0, PT, R4.reuse, 0.5, PT   // P0 = (data[idx] > 0.5)
L3: FADD R0, R4, 2                           // Block B: executed for every thread
L4: @P0 FMUL R0, R4, 4                       // Block A: conditioned on P0 (divergent state)
L5: FMUL R5, R0, R0                          // convergence point: data[idx]*data[idx]
L6: STG.E.SYS [R2], R5                       // store of the result

בואו נקרא שורה-שורה. ב-L1 נטען הערך ל-R4. ב-L2 הפקודה FSETP.GT.AND (Floating-point SET Predicate, Greater-Than) קובעת את אוגר החיווי P0: לכל thread ב-warp, P0 נדלק אם ורק אם data[idx] > 0.5. כעת P0 הוא בעצם ה-active mask של מסלול A - וקטור של 32 ביטים, ביט לכל lane, שאומר מי הולך לאן.

ועכשיו מגיע הפרט המפתיע ביותר בכל השיעור. הדבר הטבעי שהיינו מצפים לו הוא predication סימטרי: שורה @P0 FMUL למסלול A, ושורה @!P0 FADD למסלול B - כל אחת עם המסכה שלה. אבל לא זה מה שהcompiler עשה. הביטו ב-L3: פקודת ה-FADD (מסלול B, ה-else) רצה בלי predicate כלל - עבור כל 32 ה-threads, גם אלה שערכם גדול מ-0.5. רק אחריה, ב-L4, פקודת ה-FMUL (מסלול A, ה-if) רצה מותנית ב-@P0 ודורסת את התוצאה של L3 עבור ה-threads שהחיווי שלהם דלוק.

   lanes:        0    1    2    3   ...  31
   data>0.5?     T    F    T    F        T      ->  P0 = 1 0 1 0 ... 1

   L3 FADD (all):  every lane gets R0 = data+2      (even path-A lanes!)
   L4 @P0 FMUL:    only lanes with P0=1 overwrite R0 = data*4
   L5 FMUL (all):  every lane runs R5 = R0*R0       ->  the warp converged

מה המשמעות? הcompiler מבזבז ביודעין פעולת FADD מיותרת עבור כל thread שבסופו של דבר לוקח את מסלול A - ה-threads האלה חישבו data+2, מיד דרסו אותה ב-data*4, וזרקו את התוצאה. למה לעשות דבר כזה? הפילוסופיה, בניסוח ה-glossary: עדיף לבזבז חישוב מאשר להוסיף מורכבות, אפילו אם זו רק predication נוספת. פעולת FADD על CUDA Core היא זולה עד כדי גיחוך - מחזור בודד, ויש 128 יחידות FP32 בכל SM של H100 שממילא לרוב אינן רוויות. לעומת זאת, כל predicate נוסף וכל החלטת בקרת-זרימה נוספת מסבכים את מנגנון ההנפקה. הcompiler בוחר במטבע הזול: הוא מריץ צד אחד ללא תנאי לכולם, ומכתים רק את הצד השני. ראיית פקודות "מבוזבזות" ב-SASS היא אפוא התנהגות צפויה, לא באג.

נקודת ההתכנסות. ב-L4 ה-warp נמצא במצב מתפצל: רק חלק מה-lanes מריצים בפועל את ה-FMUL. ב-L5, כשמתזמן ה-warp מנפיק שוב את אותה פקודה (FMUL R5, R0, R0) עבור כל ה-threads על אותו מחזור-שעון, ה-warp התכנס מחדש (converged). מרגע זה כל 32 ה-lanes שוב פוסעים יחד, וניצול ה-SIMT חוזר למלואו.

יעילות ההסתעפות - branch efficiency

התפצלות היא תופעה איכותית; כדי לנהל אותה צריך מדד. המדד הוא branch efficiency, וה-glossary מגדיר אותו כך: המידה שבה כל ה-threads בתוך warp לוקחים את אותו מסלול ריצה כשהם נתקלים בפקודה מותנית. החישוב הוא היחס:

   branch efficiency  =  uniform branch decisions  /  total branch instructions executed

"החלטת-ענף אחידה" היא פקודת ענף שבה כל 32 ה-threads של ה-warp הסכימו על אותה תוצאה. branch efficiency גבוהה מעידה אפוא על היעדר warp divergence - היא בדיוק ההופכי הנמדד של ההתפצלות.

הנה נקודת המפתח המושגית של השיעור, וכדאי לעצור עליה. ה-GPU מודד אחידות במרחב (in space) - האם 32 ה-threads שרצים בו-זמנית בתוך warp מסכימים על הענף. זה שונה בתכלית מ-branch prediction (חיזוי הסתעפות) של CPU, שהוא אחידות בזמן (in time) - חיזוי התוצאה של ענף יחיד לאורך זרם פקודות אחד, על סמך העבר. ה-CPU שואל "לאן הענף הזה הלך בפעמים הקודמות". ה-GPU שואל "האם 32 השכנים המרחביים שלי מסכימים עכשיו". זו הבחנה מרחב-מול-זמן שחייבים להפנים: branch efficiency היא מושג SIMT מרחבי, לא מושג חיזוי זמני.

מסקנה מעשית חשובה שנובעת מכך: לא כל תנאי פוגע ביעילות. קחו את בדיקת-הגבולות הסטנדרטית שמופיעה כמעט בכל kernel:

int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) { ... }

מכיוון ש-idx עולה ברציפות עם מספר ה-thread, כמעט בכל warp כל 32 ה-threads חולקים את אותה תוצאה של idx < n: או שכולם בפנים (warp מלא לפני הגבול), או שכולם בחוץ (warp מלא אחרי הגבול). רק warp בודד - זה שה-threads שלו רוכבים בדיוק על הגבול ב-n - מתפצל. בואו נעשה חשבון קונקרטי. נניח n = 1,000,003, blockDim.x = 256, ולכן gridDim.x = ceil(1000003/256) = 3907 בלוקים, כלומר 3907 x 256 = 1,000,192 threads בסך הכל, שהם 1,000,192 / 32 = 31,256 warps:

   total warps executing the check branch  :  31,256
   uniform warps (all in or all out)       :  31,255
   divergent warps (riding the boundary n) :       1
   branch efficiency = 31,255 / 31,256     =   99.997%

בדיקת-הגבולות היא אפוא כמעט חינמית מבחינת יעילות - warp בודד מתוך 31,256 מתפצל. זה מזכיר שהמטרה איננה "אפס תנאים", אלא תנאים שבהם ה-warp נשאר אחיד. נשווה זאת ל-kernel המתפצל מהסעיף הראשון: שם התנאי הפנימי data[idx] > 0.5f על נתונים אקראיים גורם לכמעט כל warp להתפצל (ההסתברות ש-32 threads אקראיים יסכימו היא 2 x 0.5^32, אפסית). ה-kernel הזה מריץ שני ענפים לכל thread - את בדיקת-הגבולות האחידה ואת התנאי הפנימי המתפצל - ולכן ה-branch efficiency הכוללת שלו תהיה בסביבות 1 מתוך 2 = 50%. את שני המספרים האלה (99.997% מול ~50%) נמדוד במו ידינו ב-ncu בתרגול.

עצירות לוח-התוצאות - scoreboard stalls

עברנו מבקרת-זרימה לתלויות-נתונים. גם כש-warp אחיד לחלוטין ולא מתפצל כלל, הוא עלול להיתקע - הפעם לא בגלל תנאי, אלא בגלל שהוא ממתין לתוצאה של פקודה קודמת שעדיין "בטיסה". זו scoreboard stall. ה-glossary: עצירת scoreboard מתרחשת כאשר פקודה אינה יכולה להיות מונפקת בגלל תלות בתוצאה של פקודה קודמת.

מהו scoreboard? זהו מבנה חומרה שעוקב אחרי אילו אוגרים ממתינים להיכתב על-ידי פקודה שעדיין בביצוע. כשפקודה ארוכת-latency מונפקת (למשל load), החומרה מסמנת ב-scoreboard שהאוגר-יעד שלה "תפוס-בכתיבה". כל פקודה מאוחרת שרוצה לקרוא את אותו אוגר חייבת להמתין עד שה-scoreboard יתנקה. בזמן ההמתנה ה-warp אינו כשיר (not eligible) - הוא לא יכול להנפיק את פקודתו הבאה, ומתזמן ה-warp מדלג עליו לטובת warp אחר. עובדה מספרית מכוננת: ל-warp יש 6 scoreboards שהcompiler משתמש בהם כדי לעקוב אחרי תלויות-הכתיבה בין פקודות.

חשוב לחדד מול ה-CPU: מעבדים משתמשים ב-scoreboarding וב-Out-of-Order כדי לחלץ מקביליות בתוך thread בודד (ILP). ה-GPU עושה scoreboarding בין threads - זו הסתרת-latency ברמת ה-thread, לא חילוץ ILP בתוך thread. (הרעיון עצמו עתיק - הוא נולד במחשב-העל CDC 6600 מ-1966; ניסוח ה-GPU מתועד בפטנט NVIDIA מספר US7676657.)

ה-glossary מבחין בין שני סוגים של עצירת scoreboard, ואת ההבחנה הזו חייבים להכיר כי היא לב האבחון ב-ncu:

מאפיין Short scoreboard Long scoreboard
ממתין ל... פעולה ארוכת-latency שנשארת על ה-SM פעולת זיכרון שעוזבת את ה-SM
פקודות אופייניות MUFU.EX2, MUFU.SQRT (SFU); MMA (Tensor Core); LDS, STS (shared memory) LDG (load גלובלי); STG (store גלובלי)
מקור ה-latency יחידות תוך-SM עם latency משתנה הזיכרון הגלובלי (HBM), מאות מחזורים
מתי שולט קוד עתיר-טרנסצנדנטלים / MMA / shared memory שולט בקוד memory-bound

ה-short scoreboard הוא ההמתנה לפקודה תוך-SM בעלת latency משתנה: ה-Special Function Unit (יחידת הפונקציות המיוחדות) שמחשבת מתמטיקה איטית כמו MUFU.EX2 (בסיס ל-expf) ו-MUFU.SQRT; ה-Tensor Core עם פקודת ה-MMA; וגישות ל-shared memory עם LDS/STS. ה-long scoreboard הוא ההמתנה לפעולת זיכרון שיוצאת מה-SM אל היררכיית הזיכרון הגלובלי - LDG ו-STG. ה-glossary מדגיש: long scoreboard stalls שולטות בקוד memory-bound. כשה-kernel שלכם memory-bound (בלשון פרק 7), הסיבה הדומיננטית לעצירה שתראו ב-ncu תהיה Stall Long Scoreboard - זה ה-warp שממתין ל-LDG שלו לחזור מ-HBM.

סימון ה-scoreboard ב-SASS - קריאת המחסום

היופי הוא שהתלויות האלה גלויות לעין ב-SASS. הכלי cuobjdump --dump-sass (וכן nvdisasm -c על cubin) חושף לכל פקודה את סימון הבקרה/המחסום (control/barrier notation) בסוגריים המרובעים שלפני הפקודה. נביט בזוג load-then-use קלאסי - שני loads ואז פקודה שצורכת אותם:

[B------:R-:W2:-:S04]  /*00f0*/  LDG.E.SYS R0, [R2] ;
[B------:R-:W2:-:S01]  /*0100*/  LDG.E.SYS R5, [R4] ;
[B--2---:R-:W-:Y:S08]  /*0150*/  IMAD R0, R0, c[0x0][0x160], R5 ;

נפרק את השדות בסוגריים. הפורמט הוא [B<barrier>:R<read>:W<write>:<yield>:S<stall>]:

  • שדה ה-barrier (למשל B--2---) - אילו scoreboards חייבים להתנקות לפני שהפקודה הזו רשאית להיות מונפקת. שש העמדות מייצגות את 6 ה-scoreboards; הספרה בעמדה מציינת על איזה scoreboard להמתין. B--2--- פירושו "המתן עד ש-scoreboard מספר 2 יתנקה". (לצורך השוואה, B01--4- פירושו להמתין ש-scoreboards 0, 1 ו-4 כולם יתנקו.)
  • שדה ה-write (למשל W2) - איזה scoreboard הפקודה הזו מסמנת (מגדילה) כשהיא מונפקת, כלומר איזה scoreboard יעקוב אחרי התוצאה התלויה שלה.
  • שדה ה-read (R-) - scoreboard לתלות-קריאה (רלוונטי לפקודות עם אופרנד-מקור בעל latency; כאן ריק).
  • שדה ה-yield (Y או -) - רמז למתזמן להעדיף להחליף warp אחרי הפקודה.
  • שדה ה-stall (S04, S08) - מספר מחזורי-העצירה הסטטי שהמתזמן מחזיק לפני שהוא שוקל להנפיק את הפקודה הבאה מאותו warp.

עכשיו הסיפור מסתדר. שני ה-LDG מסמנים scoreboard בכתיבה (W2) - הם מודיעים "התוצאה שלי בדרך, עקבו אחריה ב-scoreboard 2". ה-IMAD המאוחר, שצורך את שני הערכים הטעונים (R0 ו-R5), נחסם על scoreboard 2 (B--2---) - הוא חייב להמתין עד ש-scoreboard 2 יתנקה, כלומר עד ששני ה-loads יחזרו מהזיכרון. וזו בדיוק ה-long scoreboard stall שנראה כסיבת-עצירה דומיננטית ב-ncu: ה-warp תקוע על ה-B--2--- הזה מאות מחזורים בזמן שה-HBM מספק את הבייטים. אם רק ה-warp הזה רץ - יחידת ה-IMAD יושבת בטל; אם יש עוד עשרות warps תושבים (occupancy מספיק, פרק 8.1), מתזמן ה-warp ממלא את החור בעבודה שלהם וה-latency מוסתר.

שימו לב לדקות: שני ה-LDG השתמשו באותו scoreboard (W2). מכיוון שיש רק 6 scoreboards, הcompiler ממחזר אותם, ולכן חסימה על scoreboard 2 ממתינה למעשה לשניהם. הדבר הזה בדיוק הוא שורש המלכודת בסעיף הבא.

מלכודות ואבחון ב-Nsight Compute

נאסוף את הכלים המעשיים לפרופיל. שני המדדים המרכזיים של השיעור נחשפים ב-ncu:

# branch efficiency - percentage of uniform branch decisions
ncu --metrics smsp__sass_average_branch_targets_threads_uniform.pct ./a.out

# average active threads per instruction (0-32) - a direct divergence metric
ncu --metrics smsp__thread_inst_executed_per_inst_executed.ratio ./a.out

# breakdown of stall reasons, including Long/Short Scoreboard
ncu --section WarpStateStats ./a.out

בדוח המלא (ncu --set full) תמצאו את ה-branch efficiency ואת "Avg. Divergent Branches" תחת "Source Counters", ואת פירוק סיבות העצירה תחת "Warp State Statistics" - שם Stall Long Scoreboard ו-Stall Short Scoreboard מופיעים כמחזורי-עצירה ממוצעים לפקודה. עבור ה-kernel המתפצל שלנו נצפה ל-branch efficiency בסביבות 50% ול-active threads בסביבות 16 מתוך 32; עבור kernel memory-bound נצפה ש-Stall Long Scoreboard יהיה הסיבה הדומיננטית.

המלכודת הקריטית של ה-scoreboards. ראינו שיש רק 6 scoreboards ושהcompiler ממחזר אותם. התוצאה, בלשון ה-glossary: מיחזור scoreboards עלול לגרום ל-Nsight Compute לסווג שגוי את סוג העצירה - עצירת long ועצירת short עלולות להתערבב אם הן משתמשות באותו scoreboard. במילים אחרות, סיווג שמדווח כ-"short scoreboard" עשוי להיות בפועל המתנה ל-load מזיכרון (long), ולהפך. אל תסמכו בעיוורון על התווית short/long. כשהמספרים מפתיעים, חזרו ל-SASS עצמו (cuobjdump --dump-sass), אתרו את הפקודה שעל ה-source line החם, וקִראו ישירות איזה scoreboard היא ממתינה עליו (שדה ה-B) ומה הפקודה שסימנה אותו (שדה ה-W). ה-SASS הוא מקור-האמת; המונה של ncu הוא הערכה.

מלכודת נוספת - התפצלות איננה אסון. כפי שראינו, על H100 התפצלות אינה גוררת סריאליזציה מלאה. לפני שאתם משכתבים אלגוריתם שלם כדי להסיר ענף, מדדו: אם ה-branch efficiency כבר גבוהה (כמו בבדיקת-גבולות), אין מה לתקן. ואם ה-kernel memory-bound (long scoreboard שולט), אזי גם התפצלות מלאה כמעט לא תשפיע - צוואר-הבקבוק הוא ה-HBM, לא ה-CUDA Cores שמריצים את הענף. תמיד אבחנו קודם, ואז אופטמו את הצוואר האמיתי.

סיכום

  • התפצלות ה-warp (warp divergence) מתרחשת כאשר 32 ה-threads בתוך warp פוסעים במסלולי בקרת-זרימה שונים בעקבות תנאי תלוי-נתונים; זו תופעה בלתי-נמנעת של מודל ה-SIMT, ולא ניתן "לתקן" אותה בקוד כשההחלטה תלויה בערך פר-thread.
  • לפני Volta warp מתפצל היה תמיד מסוריאל במלואו (מכאן המיתוס "פי 32 האטה"); מ-Volta ואילך (וכך ב-H100) יש independent thread scheduling עם מונה-תוכנית פר-thread, ולכן אין קנס סריאליזציה מלא מובנה - התפצלות עדיין עולה, אך אינה גזירת-דין.
  • החומרה מטפלת בהתפצלות עם predicate registers ו-active masks: FSETP.GT.AND קובע את P0 כמסכת ה-lanes, ופקודה מותנית @P0 מבוצעת רק על ה-lanes הדלוקים; ההתכנסות (reconvergence) חוזרת כשכל ה-threads מנפיקים שוב אותה פקודה יחד.
  • האסטרטגיה הלא-אינטואיטיבית של הcompiler: להריץ צד אחד של הענף (FADD) עבור כל ה-threads ללא תנאי, ואז לדרוס אותו עם הצד השני המותנה (@P0 FMUL) - בזבוז מכוון של חישוב CUDA-Core זול כדי להימנע ממורכבות בקרת-זרימה; פקודות "מבוזבזות" ב-SASS הן צפויות, לא באג.
  • יעילות ההסתעפות (branch efficiency) = החלטות-ענף אחידות חלקי סך פקודות-הענף; זהו ההופכי הנמדד של ההתפצלות, והוא מודד אחידות במרחב (האם 32 ה-threads מסכימים עכשיו), בניגוד ל-branch prediction של CPU שהוא חיזוי בזמן.
  • בדיקת-הגבולות if (idx < n) כמעט חינמית: רק warp בודד שרוכב על הגבול מתפצל, ולכן היעילות ~99.997% - לא כל תנאי פוגע ביעילות, רק זה שמפצל את ה-warp.
  • ה-scoreboard stall היא עצירת warp שממתין לתוצאת פקודה קודמת שעדיין בטיסה; ל-warp יש 6 scoreboards שעוקבים אחרי תלויות-כתיבה, וה-GPU עושה scoreboarding בין threads (הסתרת-latency), לא לחילוץ ILP בתוך thread כמו OoO של CPU.
  • ה-short scoreboard הוא המתנה לפעולה ארוכת-latency שנשארת על ה-SM (MUFU.EX2/SQRT, MMA, LDS/STS); long scoreboard = המתנה לזיכרון שיוצא מה-SM (LDG/STG), והוא הסיבה הדומיננטית לעצירה בקוד memory-bound.
  • סימון המחסום ב-SASS [B<barrier>:R-:W<n>:<yield>:S<stall>] נקרא כך: שדה ה-B מציין על אילו scoreboards להמתין, שדה ה-W מציין איזה scoreboard הפקודה מסמנת; זוג load-then-use מראה LDG ... W2 ואז IMAD ... B--2--- הממתין ל-load.
  • ב-ncu מודדים branch efficiency עם smsp__sass_average_branch_targets_threads_uniform.pct ואת סיבות העצירה תחת "Warp State Statistics"; מלכודת: מיחזור scoreboards עלול לגרום ל-ncu לסווג שגוי short מול long, ולכן כשמופתעים חוזרים ל-SASS כמקור-האמת.