8.2 - ניצולת - SM, pipe ו-issue efficiency - הרצאה
בשיעור 8.1 הכרנו את ה-occupancy (תפוסה) - היחס בין ה-warps הפעילים על ה-SM לבין המקסימום - וראינו שהיא אמצעי להסתרת latency, לא מטרה בפני עצמה. אבל occupancy לבדה לא אומרת לנו כמה עבודה שימושית ה-GPU באמת מוציא מהחומרה. בשביל זה צריך משפחה שנייה של מדדים: מדדי הניצולת (utilization) - כמה מהזמן, וכמה מקצב-השיא, החומרה באמת מנצלת. בשיעור הזה נבנה את היררכיית מדדי הניצולת ונלמד לדבג ביצועים בשיטת top-down: מהמדד הכי גס שאפשר לקרוא במבט אחד על nvidia-smi, ועד למדד הכי עדין שנמדד ברמת מחזור-שעון בודד במתזמן ה-warp. סדר הדיבוג הזה קריטי: מדד גס שנראה טוב יכול להסתיר בזבוז עצום ברמה עדינה יותר, ולכן תמיד יורדים במדרג צעד-צעד. את הכלים עצמם - Nsight Compute (ncu) ו-Nsight Systems (nsys) - נכיר לעומק בפרק 9 (שיעור 9.1); כאן נשתמש בהם ככלי מדידה ונתמקד במה מודדים ואיך קוראים את המספרים. ה-H100 (עם 132 ה-SM שלו) יהיה כרטיס הייחוס לאורך כל הדרך.
היררכיית מדדי הניצולת - the utilization hierarchy¶
ה-glossary מציב ארבעה מדדי ניצולת, מהגס לעדין, וממליץ לדבג אותם בדיוק בסדר הזה. הרעיון פשוט: כל שכבה מכסה שכבה עמוקה יותר, ואין טעם לתקן דבר עדין כל עוד השכבה שמעליו שבורה. אם ה-SM שלכם עומדים בטל חצי מהזמן, אין שום משמעות לשאלה כמה יעיל pipe ה-FMA שלכם ב-cycles שהם כן עובדים.
coarse to fine - debugging order (top-down):
┌─────────────────────────────────────────────────────────────┐
│ 1. GPU / kernel utilization (nvidia-smi) │ coarse
│ "is the GPU running any kernel right now?" │
├─────────────────────────────────────────────────────────────┤
│ 2. SM utilization (sm__cycles_active...) │
│ "how many of the 132 SMs are truly executing instructions?"│
├─────────────────────────────────────────────────────────────┤
│ 3. pipe utilization (sm__inst_executed_pipe_*...) │
│ "inside a working SM - how much of peak does each │
│ execution pipe deliver?" │
├─────────────────────────────────────────────────────────────┤
│ 4. issue efficiency (smsp__issue_active...) │ fine
│ "in how many cycles did the scheduler have an eligible │
│ warp to issue?" │
└─────────────────────────────────────────────────────────────┘
golden rule: fix top to bottom. don't drop a layer until
the layer above it is as high as it can be.
שימו לב שכל שכבה שואלת שאלה מוכלת בשאלה שמעליה. שכבה 1 שואלת אם ה-GPU עסוק בכלל. שכבה 2 מתנה על "עסוק" ושואלת כמה מה-SM עסוקים. שכבה 3 מתנה על "SM עסוק" ושואלת כמה מהיחידות בתוכו עסוקות. שכבה 4 מתנה על "יחידה עסוקה" ושואלת כמה מהמחזורים באמת הונפקה הוראה. ככל שיורדים, המכנה מצטמצם וה"עדשה" מתחדדת. נעבור עכשיו שכבה-שכבה.
ניצולת ה-GPU - הקומה העליונה - nvidia-smi¶
הקומה הגסה ביותר היא ניצולת ה-GPU (או ניצולת ה-kernel), וזה המספר שכולם מכירים - העמודה שקופצת ל-100% ב-nvidia-smi. הגדרתו: חלק הזמן שבו kernel כלשהו רץ איפשהו על ה-GPU. זהו מדד בינארי כמעט - או שיש kernel פעיל, או שאין. הוא לא יודע דבר על כמה מהחומרה ה-kernel הזה מנצל.
אפשר לקרוא אותו בשתי צורות. בצילום-בזק:
nvidia-smi --query-gpu=utilization.gpu,utilization.memory --format=csv
# utilization.gpu [%], utilization.memory [%]
# 100 %, 3 %
או במעקב רציף (מדגם לשנייה):
וזו בדיוק המלכודת: ה-sm % של nvidia-smi הוא לא ניצולת ה-SM האמיתית שנדבר עליה מיד. הוא מדגם דגימה גסה של "האם היה SM פעיל בחלון הדגימה", ולא כמה SM היו פעילים בו-זמנית. nvidia-smi הוא כלי ניטור מערכת (כפי שראינו בשיעור 5.4), לא profiler. הוא מצוין כדי לענות "האם ה-GPU עובד או בטל", והוא הצעד הראשון בדיבוג - אבל אם הוא מראה 100% זו רק ההתחלה, לא הסוף.
ניצולת ה-SM - הלב של הסיפור - SM utilization¶
עכשיו יורדים שכבה. ניצולת ה-SM (SM utilization) מוגדרת כאחוז הזמן שבו ה-SM מריצים הוראות. ה-glossary מנסח זאת יפה: זה "בן-דוד עדין יותר" של ניצולת ה-kernel של nvidia-smi. במקום לשאול "האם ה-GPU פעיל", היא שואלת "כמה מ-132 ה-SM של ה-H100 באמת מריצים הוראות, ובאיזה חלק מהזמן".
ההבחנה הזו היא כנראה השיעור החשוב ביותר בכל הפרק, והדרך הטובה להפנים אותה היא דרך הדוגמה הקנונית של ה-glossary: kernel עם block אחד בלבד.
// kernel that occupies a single block - a classic performance mistake
myKernel<<<1, 256>>>(...); // grid of size 1: one block, one SM
בלוק אחד רץ, לפי המודל, על SM אחד בלבד (בלוק לעולם לא מתפצל בין SM). כל עוד ה-kernel הזה רץ, nvidia-smi יראה בשמחה 100% ניצולת GPU - יש kernel פעיל. אבל בפועל SM אחד עובד ו-131 עומדים בטל. ניצולת ה-SM האמיתית:
SM utilization = active SMs / total SMs
one block -> 1 SM working out of 132
= 1 / 132 = 0.0076 = 0.76%
nvidia-smi will show 100%. SM utilization is 0.76%.
a 130x gap that the coarse metric simply doesn't see.
ה-glossary מנסח זאת מדויק: kernel כזה "ישיג 100% ניצולת GPU בזמן שהוא פעיל, אבל ניצולת ה-SM תהיה לכל היותר אחד חלקי מספר ה-SM - מתחת ל-1% ב-H100". זהו בדיוק המקרה שבו הקומה העליונה משקרת לכם, והקומה השנייה חושפת את האמת. התיקון פשוט: הגדילו את ה-grid כך שיהיו מספיק blocks כדי להאכיל את כל 132 ה-SM (ורצוי כמה blocks לכל SM, כדי לתת ל-scheduler ממה לבחור).
נקודה תפיסתית קריטית - הפוך מ-CPU. אצל מהנדסי CPU, ניצולת מעבד של 100% לאורך זמן היא לרוב דגל אזהרה - סימן לתהליך שרץ מטורף או ל-thread תקוע בלולאה. אצל GPU זה בדיוק ההפך: ניצולת SM גבוהה, עד 100%, היא המטרה. ה-glossary אומר זאת במפורש: "כמו ניצולת ה-GPU אבל בניגוד לניצולת ה-CPU, ניצולת ה-SM צריכה להיות גבוהה, אפילו עד 100%". ה-GPU נבנה כדי להיות רווי; SM בטל הוא סיליקון מבוזבז, לא מרווח בטיחות.
עם זאת - וכאן הגשר לשכבה הבאה - ניצולת SM גבוהה היא תנאי הכרחי אך לא מספיק. SM יכול להיות "פעיל" (מריץ הוראות) כל הזמן, ובכל זאת להריץ את ההוראות הלא-נכונות, או להשאיר את רוב יחידות הביצוע שבתוכו מובטלות. בשביל לראות את זה צריך לרדת עוד שכבה.
ניצולת pipes הביצוע - pipe utilization¶
ניצולת ה-pipe (ניצולת pipes הביצוע) מודדת כמה יעיל ה-kernel מנצל את יחידות הביצוע בתוך כל SM. הנוסח המדויק של ה-glossary: אחוז מקצב-השיא (peak rate) של כל pipe שמושג כשהpipe פעיל ומריץ לפחות warp אחד, ממוצע על פני כל ה-SM הפעילים.
יש בהגדרה הזו ניואנס חד שחייבים להבין, אחרת קוראים את המדד לא נכון: המדד מותנה על כך שהpipe פעיל. הוא מודד הישג-מול-שיא רק במחזורים שבהם הpipe עושה עבודה. מחזורים שבהם הpipe בטל מוחרגים מהממוצע. המשמעות המעשית: הpipe שכמעט לא בשימוש, אבל בכל פעם שכן משתמשים בו הוא רץ בשיא, יראה ניצולת גבוהה - למרות שהוא כמעט תמיד בטל. לכן חייבים לקרוא ניצולת pipe יחד עם occupancy וניצולת SM, אף פעם לא לבד.
ה-glossary מזהה ארבעה pipes ביצוע נפרדים בתוך כל SM:
| הpipe | מה הוא מריץ | דוגמאות הוראות (SASS) |
|---|---|---|
| CUDA Cores / FMA | חשבון נקודה-צפה כללי | FMUL, FADD, FFMA |
| Tensor Cores | כפל-צבירה של מטריצות | HMMA, IMMA |
| Load/Store Units (LSU) | גישות זיכרון | LDG, STG, LDS, STS |
| Control Flow / ADU | הסתעפויות וחישוב כתובות | BRA, BSSY |
היכולת של ארבעת הpipes לעבוד במקביל היא בדיוק מה שמאפשר הסתרת latency (שיעור 7.4): warp אחד יכול לרוץ על ה-Tensor Cores, שני על ה-CUDA Cores, ושלישי דרך ה-LSU - כולם בו-בזמן. ניצולת pipe מודדת עד כמה ניצלתם את המקבילות הזו.
המדדים ב-Nsight Compute. ניצולת ה-pipe נחשפת דרך משפחת המדדים:
כאשר <PIPE> הוא אחד מ:
fma -> the FMA pipe (CUDA Cores)
tensor -> Tensor Cores
lsu -> Load/Store Units
adu -> Address Unit (address computation)
הסיומת .avg.pct_of_peak_sustained_active היא בדיוק מה שמקודד את שלושת חלקי ההגדרה: pct_of_peak = אחוז מקצב-השיא, .avg = ממוצע על ה-SM, ו-active = רק במחזורים פעילים. נזכיר שהמכנה כאן, קצב-השיא, הוא בדיוק אותו peak rate שחישבנו בשיעור 7.3 (עבור H100 ב-FP32: 66.9 TFLOPS). ניצולת pipe היא, פשוטו כמשמעו, ההישג חלקי אותו peak.
המלכודת, שוב. בגלל שמחזורים בטלים מוחרגים מהממוצע, מספר גבוה של ניצולת pipe לא אומר שהpipe היה בשימוש כבד - רק שבכל פעם שכן השתמשו בו הוא רץ קרוב לשיא. pipe_tensor של 95% על kernel ש-99% מהזמן לא נוגע ב-Tensor Cores הוא מספר חסר-משמעות אם קוראים אותו לבד. לכן ה-glossary ממליץ לתקן ניצולת GPU וניצולת SM לפני שיורדים לרמת ה-pipe.
יעילות הנפקת ההוראות - issue efficiency¶
הקומה העמוקה ביותר היא יעילות ההנפקה (issue efficiency) - היחידה שנמדדת ברמת מחזור-שעון בודד, אצל מתזמן ה-warp עצמו. ה-glossary מגדיר: חלק המחזורים הפעילים (active cycles) שבהם המתזמן הנפיק הוראה - כלומר, בהם היה לו לפחות warp אחד כשיר (eligible) לבחור ממנו.
- יעילות הנפקה של 100% = בכל מחזור פעיל, כל scheduler הנפיק הוראה. תמיד היה warp כשיר.
- מתחת ל-100% = במקצת המחזורים כל ה-warps הפעילים היו תקועים (stalled) - ממתינים לנתונים, למשאב או לתלות - והמתזמן ישב בטל, כך שהתפוקה הכוללת ירדה.
כדי לחבר את זה חזרה לשיעור 8.1: היעילות הזו תלויה ישירות בשלושת מצבי ה-warp שהכרנו. warp פעיל (active) הוא warp תושב על ה-SM; warp כשיר (eligible) הוא warp פעיל שמוכן להנפיק את הוראתו הבאה במחזור הזה; warp תקוע (stalled) הוא warp פעיל שלא יכול. יעילות ההנפקה צונחת בדיוק כשכל ה-warps הפעילים תקועים בו-זמנית, ולמתזמן אין ממי להנפיק:
cycle: 1 2 3 4 5 6
warp A: ISS -- -- -- ISS -- (-- = stalled)
warp B: -- ISS -- -- -- ISS
warp C: -- -- ISS -- -- --
warp D: -- -- -- [ ] -- -- <- cycle 4: no one eligible!
issued in 5 out of 6 cycles -> issue efficiency = 83%
cycle 4 is the "hole" - all warps stalled at once.
התרופה ליעילות הנפקה נמוכה היא כמעט תמיד יותר מקבילות ברמת ה-thread (TLP) - עוד warps כשירים, כדי שבכל מחזור לפחות אחד יהיה מוכן. וזה בדיוק המקום שבו occupancy נכנסת: occupancy נמוכה מדי פירושה מעט warps תושבים, ומעט warps לבחירה, ולכן חורים ביעילות ההנפקה. זו החוליה שסוגרת את המעגל בין 8.1 ל-8.2 - occupancy היא הסיבה האפשרית, ויעילות ההנפקה היא הסימפטום הנמדד.
ה-glossary עצמו לא נותן מחרוזת-מדד ל-issue efficiency, אבל ידע-CUDA סטנדרטי כן: ב-Nsight Compute המדד המקביל הוא Issue Slot Utilization:
הוא נמדד ברמת ה-SM sub-partition (smsp) - כל scheduler הוא sub-partition, וב-H100 יש ארבעה כאלה לכל SM. שימו לב לניואנס לעומת ניצולת ה-pipe: pipe שואל "כמה מקצב-השיא של יחידת החישוב", ואילו issue efficiency שואל "בכמה מהמחזורים בכלל הונפקה הוראה" - האחד על קצב היחידה, השני על נוכחות warp כשיר. אפשר בהחלט להנפיק הוראה בכל מחזור (issue efficiency גבוה) ובכל זאת להאכיל pipe לא-נכון או קל (pipe utilization נמוך), ולהפך.
קריאה משולבת - עץ ההחלטה - the diagnostic decision path¶
עכשיו מגיע החלק המעניין: המדדים מספרים סיפור רק ביחד. הצירוף של ניצולת SM וניצולת pipe הוא אבחנתי במיוחד. ה-glossary מנסח את המקרה המרכזי: ניצולת SM גבוהה עם ניצולת pipe נמוכה פירושה ש-kernel רץ על הרבה SM אבל לא מנצל במלואן את יחידות החישוב שבתוך כל אחד - "רחב אבל רדוד". פרשׂתם את העבודה יפה על פני 132 ה-SM, אבל בתוך כל SM היחידות רעבות. נסדר את הצירופים בטבלה:
| ניצולת SM | ניצולת pipe | האבחנה | הכיוון לתיקון |
|---|---|---|---|
| נמוכה | לא רלוונטי | לא מאכילים את המכונה | הגדילו grid; יותר blocks (טפלו ב-SM קודם) |
| גבוהה | נמוכה | "רחב אבל רדוד" - SM עובדים, יחידות רעבות | יותר עבודה ל-thread; ILP; בדקו איזה pipe דומיננטי |
| גבוהה | גבוהה (pipe אחד) | רוויתם pipe אחד - זהו הצוואר | זהו ה-pipe החסום; אזנו/החליפו אלגוריתם |
| גבוהה | גבוהה (מפוזר) | ניצול טוב - קרוב לגבול החומרה | ייתכן שהגעתם ל-roofline; מדדו מול peak |
וזה, בצורת עץ, הוא בדיוק מסלול הדיבוג top-down שאיתו פתחנו:
start nvidia-smi: GPU utilization ~0%?
│ └── yes -> the GPU is idle. the problem is on the host/scheduling (lesson 7.4). stop.
│ └── no (100%) -> go down a layer.
▼
ncu: SM utilization low (e.g. <20%)?
│ └── yes -> few SMs working. increase grid / more blocks. fix this first!
│ └── no (high) -> go down a layer.
▼
ncu: pipe utilization low across every pipe?
│ └── yes -> "wide but shallow". check issue efficiency ->
│ │ low -> warps stalled. raise TLP/occupancy,
│ │ investigate stall reasons (rest of chapter 8).
│ └ high -> issuing fine, but the work is light. increase
│ arithmetic intensity (lesson 7.2), fuse kernels.
│ └── no (one pipe high) -> this is the bottleneck. it's the dominant pipe.
▼
measure the dominant pipe against peak rate (roofline, lesson 7.1). that's the limit.
שימו לב שהעץ תמיד מטפל בשכבה הגסה קודם. אם ניצולת ה-SM היא 0.76% (ה-kernel של block יחיד), אין שום טעם לפתוח את מדדי ה-pipe - הם ימדדו את ה-SM היחיד שכן עובד ואולי יראו לכם מספרים יפים לגמרי, בזמן ש-131 SM עומדים בטל. קודם מרווים את ה-SM, ואז - ורק אז - שואלים כמה טוב עובד כל אחד.
הכלים בפועל - ncu ו-nsys¶
נחבר הכל לפקודה אחת. ncu (Nsight Compute) יכול לשלוף את כל שכבות ההיררכיה במדידה אחת. הנה קריאה ממוקדת שמבקשת בדיוק מדד אחד מכל שכבה:
ncu --metrics \
sm__cycles_active.avg.pct_of_peak_sustained_elapsed,\
sm__inst_executed_pipe_fma.avg.pct_of_peak_sustained_active,\
sm__inst_executed_pipe_tensor.avg.pct_of_peak_sustained_active,\
sm__inst_executed_pipe_lsu.avg.pct_of_peak_sustained_active,\
smsp__issue_active.avg.pct_of_peak_sustained_active \
./my_app
השורה הראשונה (sm__cycles_active...elapsed) היא ניצולת ה-SM - חלק מחזורי-השעון החולפים שבהם ה-SM היו פעילים, ממוצע על כל ה-SM. שימו לב לסיומת elapsed (ולא active): כאן דווקא רוצים לספור גם את המחזורים הבטלים במכנה, אחרת לא נראה את ה-SM המובטלים. שלוש השורות האמצעיות הן ניצולת ה-pipe (עם active, כי שם מחריגים מחזורים בטלים בכוונה), והאחרונה היא יעילות ההנפקה. הנה איך פלט על ה-kernel של block-יחיד ייראה, וכיצד לקרוא אותו:
sm__cycles_active.avg.pct_of_peak_sustained_elapsed 0.76 %
sm__inst_executed_pipe_fma.avg.pct_of_peak_sustained_active 71.20 %
sm__inst_executed_pipe_tensor... 0.00 %
sm__inst_executed_pipe_lsu... 12.40 %
smsp__issue_active.avg.pct_of_peak_sustained_active 68.90 %
הקריאה: ניצולת ה-SM היא 0.76% - כמעט כל המכונה בטלה. ה-fma מראה 71% מפתה, אבל זו בדיוק המלכודת - הוא ממוצע רק על ה-SM היחיד שעובד ומחריג את 131 הבטלים. אם היינו קופצים ישר למספר ה-71% היינו מתפתים לחשוב שהכול תקין. העץ מציל אותנו: ניצולת SM של 0.76% פירושה שעוצרים כאן ומגדילים את ה-grid, ורק אחר-כך חוזרים לקרוא את שאר המדדים.
לניתוח מלא ואינטראקטיבי אפשר ncu --set full ./my_app שמפיק את דוח ה-Speed of Light (סעיף שמראה Compute (SM) Throughput ו-Memory Throughput זה-לצד-זה). ולתמונת-על על פני ה-timeline של כל ה-kernels - מי לוקח כמה זמן, מי מריץ מתי - משתמשים ב-nsys profile ./my_app. חלוקת העבודה: nsys עונה "איזה kernel לדבג" (רמת המערכת), ו-ncu עונה "למה ה-kernel הזה איטי" (רמת ה-SM ומטה). את שניהם נפתח לעומק בשיעור 9.1.
סיכום¶
- מדדי הניצולת בנויים בהיררכיה מהגס לעדין - ניצולת GPU (
nvidia-smi), ניצולת SM, ניצולת pipe, ויעילות הנפקה - ומדבגים אותם top-down, בסדר הזה בדיוק, בלי לרדת שכבה עד שהשכבה שמעליה טובה כמה שהיא יכולה. - ניצולת ה-GPU של
nvidia-smiהיא בינארית כמעט - יש kernel פעיל או אין - והיא יכולה להראות 100% גם כשרוב החומרה בטלה; היא נקודת ההתחלה של הדיבוג, לא סופו. - ניצולת ה-SM היא אחוז הזמן שבו ה-SM מריצים הוראות; kernel עם block יחיד משיג 100% ניצולת GPU אבל רק 1/132 = 0.76% ניצולת SM על H100, כי block רץ על SM אחד בלבד.
- בניגוד ל-CPU, שבו 100% ניצולת הוא לרוב דגל אזהרה, ב-GPU ניצולת SM גבוהה עד 100% היא המטרה - SM בטל הוא סיליקון מבוזבז.
- ניצולת ה-pipe מודדת הישג-מול-peak של כל אחד מארבעה pipes (FMA/CUDA, Tensor, LSU, control/ADU), אבל ממוצעת רק על מחזורים שבהם הpipe פעיל - ולכן חייבים לקרוא אותה יחד עם occupancy וניצולת SM, אחרת מספר גבוה מטעה.
- משפחת המדדים ב-Nsight היא
sm__inst_executed_pipe_{fma,tensor,lsu,adu}.avg.pct_of_peak_sustained_active, כאשר המכנה הוא ה-peak rate מפרק 7. - יעילות ההנפקה היא חלק המחזורים הפעילים שבהם המתזמן הנפיק הוראה (היה warp כשיר); מתחת ל-100% פירושו שבמקצת המחזורים כל ה-warps היו תקועים - והתרופה היא יותר warps כשירים, כלומר occupancy מספקת ו-latency hiding טוב.
- הצירוף האבחנתי המרכזי: ניצולת SM גבוהה עם ניצולת pipe נמוכה = "רחב אבל רדוד", kernel שרץ על הרבה SM אבל מרעיב את היחידות בתוך כל אחד; הפתרון הוא יותר עבודה ל-thread או arithmetic intensity גבוה יותר.
ncuשולף את כל השכבות במדידה אחת ועונה "למה ה-kernel איטי";nsysנותן timeline ועונה "איזה kernel לדבג" - ובשניהם נעמיק בשיעור 9.1.