7.1 צווארי בקבוק ומודל Roofline פתרון
פתרון - צווארי בקבוק ומודל Roofline¶
כאן נעבור על הפתרון המלא של כל תרגיל, עם החישובים, השרטוטים והפקודות המדויקים. שימו לב: המספרים של H100 הם ממספרי הספסיפיקציה של ההרצאה (989 TFLOP/s ב-BF16, 3.35 TB/s ב-HBM3), ופלטי ה-ncu בבונוס הם דוגמה מהרצה על H100 (sm_90a) עם גרסת CUDA Toolkit מסוימת - הarithmetic intensity המדידה והמספרים המדויקים ישתנו מעט בין ארכיטקטורות, גרסאות compiler וגודל וקטור, אבל המבנה והמסקנות זהים.
פתרון תרגיל 1 - בניית ה-Roofline של H100¶
1. גג החישוב (compute roof). קו אופקי בגובה רוחב הפס האריתמטי:
2. גג הזיכרון (memory roof). קו אלכסוני מהראשית ששיפועו רוחב פס הזיכרון:
slope = memory BW = 3.35 TB/s = 3.35e12 byte/s
unit check: performance / intensity = [FLOP/s] / [FLOP/byte] = byte/s (valid)
3. נקודת הרכס (ridge point):
4. השרטוט:
performance [FLOP/s], log
^
989T|.....................__________________ compute roof (horizontal)
| /:
| memory roof / :
| (slope = / : compute-bound
| 3.35e12) / : (right of ridge)
| / :
| memory- / :
| bound / :
| (left) / :
+--------/-----------+-----------------> arithmetic intensity [FLOP/byte], log
295 = ridge point
למה זה עבד: שני הגגות נגזרים אך ורק מהחומרה - רוחב פס אריתמטי (גובה הקו האופקי) ורוחב פס זיכרון (שיפוע הקו האלכסוני). נקודת הרכס היא היחס ביניהם, ולכן היא מספר יחיד שמאפיין את הכרטיס בלי שום תלות ב-kernel. בדיקת היחידות של השיפוע היא הדרך לוודא שלא התבלבלתם: שיפוע של גרף FLOP/s מול FLOP/byte הוא בהכרח byte/s.
איך להכליל: אותו תהליך בונה Roofline לכל חומרה - הציבו את שני רוחבי הפס שלה. מספר הרכס הבודד מסכם "כמה arithmetic intensity דרושה כדי להשתלם". שמרו אותו בראש: כל kernel שתפרופלו ייפול משמאל או מימין לקו האנכי הזה.
פתרון תרגיל 2 - מיקום kernel-ים על ה-Roofline¶
שלב 1 - arithmetic intensity לכל kernel (FLOPs / bytes):
A (SAXPY): 2.0e8 / 1.2e9 = 0.167 FLOP/byte
B (N=1024): 2.15e9 / 1.68e7 = 128 FLOP/byte
C (N=4096): 1.37e11/ 2.68e8 = 512 FLOP/byte
D (hypothetical): 5.9e11 / 2.0e9 = 295 FLOP/byte
שלב 2 - השוואה לנקודת הרכס (295) ותיוג:
| Kernel | הarithmetic intensity | מול 295 | תיוג |
|---|---|---|---|
| A - SAXPY | 0.167 | << 295 | memory-bound |
| B - SGEMM N=1024 | 128 | < 295 | memory-bound |
| C - SGEMM N=4096 | 512 | > 295 | compute-bound |
| D - היפותטי | 295 | = 295 | בדיוק על הרכס (מאוזן) |
שימו לב לתובנה החשובה: SGEMM N=1024 עדיין memory-bound על H100 למרות שהוא כפל מטריצות. רק כשמגדילים ל-N=4096 (arithmetic intensity 512) הוא חוצה את הרכס והופך לcompute-bound. זה מדגים ישירות את הנוסחה מההרצאה: SGEMM הופך לcompute-bound סביב N/8 >= 295, כלומר N >= 2360. ה-kernel D יושב בדיוק על הרכס - הarithmetic intensity המינימלית להיות compute-bound.
שלב 3 - ביצועים מקסימליים של הmemory-bound (arithmetic_intensity * 3.35e12):
A: 0.167 * 3.35e12 = 5.6e11 = 0.56 TFLOP/s (0.057% of 989)
B: 128 * 3.35e12 = 4.29e14 = 429 TFLOP/s (43% of 989)
ה-SAXPY מגיע ל-0.56 TFLOP/s בלבד - פחות מאלפית מגג החישוב 989 TFLOP/s. הוא לא איטי כי החישוב איטי, אלא כי ה-arithmetic intensity שלו זעירה: הכרטיס מבלה כמעט את כל זמנו בהמתנה ל-HBM3.
שלב 4 - האם יותר TFLOPS יעזרו? לא. עבור kernel memory-bound (A ו-B), הביצועים נקבעים על ידי arithmetic_intensity * memory_BW - גג החישוב כלל אינו מגביל אותם. GPU עם פי-שניים TFLOPS יריץ אותם באותה מהירות בדיוק. מה שכן יעזור: (א) רוחב פס זיכרון גבוה יותר, או (ב) העלאת הarithmetic intensity - fusion של פעולות כדי לחלוק טעינות, שימוש חוזר בנתונים ב-shared memory (tiling, שראינו בפרק 3), או החלפת עבודת זיכרון בעבודת חישוב.
למה זה עבד: ה-kernel נופל בצד של הרכס לפי הarithmetic intensity שלו בלבד. משמאל לרכס, הגג האלכסוני (זיכרון) הוא המגביל, והביצועים הם arithmetic_intensity * memory_BW; מימין, הגג האופקי (חישוב) מגביל. לכן שני מספרים - FLOPs ו-bytes - קובעים הכל.
איך להכליל: זהו בדיוק מה ש-Nsight Compute עושה אוטומטית: סופר FLOPs ובתים בפועל, מחשב arithmetic intensity, ומציב אתכם ביחס לרכס. לכל kernel בעבודה אמיתית, "היכן אני יחסית לרכס" היא השאלה הראשונה - היא מכתיבה אם להשקיע בשיפור גישות הזיכרון (coalescing, tiling) או בשיפור החישוב.
פתרון תרגיל 3 - גזירת נקודת הרכס¶
1-2. הגזירה. על גג הזיכרון, הביצועים כפונקציה של הarithmetic intensity I הם:
על גג החישוב הביצועים קבועים:
נקודת הרכס היא היכן שהם שווים. משווים ומבודדים את I:
3. יישום על שלושה דורות:
A100 (Ampere) : 312e12 / 2.0e12 = 156 FLOP/byte
H100 (Hopper) : 989e12 / 3.35e12 = 295 FLOP/byte
B200 (Blackwell): 2250e12 / 8.0e12 = 281 FLOP/byte
4. פרשנות המגמה. מ-A100 ל-H100 נקודת הרכס כמעט הכפילה עצמה (156 -> 295): המגמה היא ימינה. הכיוון הזה נובע מכך שכוח החישוב (מונה) גדל מהר יותר מרוחב פס הזיכרון (מכנה) - קיר הזיכרון. ככל שהרכס ימינה יותר, דרושה arithmetic intensity גבוהה יותר כדי להיות compute-bound, כלומר "קשה יותר" לנצל את מלוא כוח החישוב.
הערה על B200: ב-BF16 הרכס שלו (281) יצא מעט משמאל ל-H100 (295), כי B200 קיבל קפיצה גדולה במיוחד ברוחב פס הזיכרון (HBM3e, 8 TB/s). המגמה ארוכת-הטווח ימינה ברורה במיוחד בדיוקים נמוכים: ב-FP8 הרכס של H100 הוא 592 ושל B200 הוא 562, וב-FP4 של B200 הוא מזנק ל-1125. במילים אחרות, בכל דור החישוב מזנק, והדרך היחידה להאכיל אותו היא arithmetic intensity גבוהה יותר ויותר - מטריצות גדולות יותר.
למה זה עבד: נקודת הרכס היא, בהגדרה, החיתוך של שני הגגות, ולכן שוויון שני הביטויים נותן מיד את היחס בין רוחבי הפס. אין כאן שום דבר תלוי-kernel - רק שני מספרי חומרה.
איך להכליל: הנוסחה arithmetic_BW / memory_BW עובדת לכל צירוף דיוק ולכל תת-מערכת. היא גם הכלי לחזות דורות: אם אתם יודעים כמה יגדלו החישוב והזיכרון בדור הבא, אתם יודעים לאן תזוז נקודת הרכס - ולכן כמה גדולות יצטרכו להיות הבעיות שלכם כדי להישאר compute-bound.
פתרון תרגיל 4 - למה kernel overhead-bound בלתי-נראה ב-Roofline¶
1. הarithmetic intensity ומיקום:
הarithmetic intensity 0.244 היא הרחק משמאל לנקודת הרכס (295), ולכן ה-Roofline ימקם אותו עמוק באזור memory-bound.
2. פירוק זמן הקיר:
total time = overhead + work = 10us + 1us = 11us
overhead = 10 / 11 = 90.9% (idle time)
work = 1 / 11 = 9.1% (useful work)
3. מה באמת מגביל. ה-Roofline מסווג אותו כmemory-bound כי הarithmetic intensity שלו נמוכה, אבל זה טעות אבחון: 91% מהזמן הוא overhead dispatch טהורה (latency של קריאת CUDA API), ולא הזזת נתונים. ה-GPU יושב בטל ומחכה לעבודה, לא עובד לאט על זיכרון. המודל אינו יכול לומר לכם זאת כי ה-latency פשוט לא קיים בו.
4. למה latency בלתי-נראה ב-Roofline. שני הצירים של המודל הם arithmetic intensity (FLOP/byte) וביצועים (FLOP/s), ושני הגגות הם רוחב פס אריתמטי ורוחב פס זיכרון. אין ציר, אין גג, ואין נקודה שמייצגים latency או overhead - רק רוחבי פס ותפוקות. לכן kernel overhead-bound נראה ב-Roofline בדיוק כמו kernel memory-bound עם arithmetic intensity נמוכה, ואי אפשר להבחין ביניהם מהגרף לבדו.
למה זה עבד: ההבחנה היסודית היא בין "עבודה איטית" ל"זמן-סרק". צוואר בקבוק של חישוב או זיכרון = ה-GPU עובד הכי מהר שהוא יכול. הoverhead = ה-GPU בטל ומחכה. ה-Roofline מודד רק את הראשון.
איך להכליל: לאבחון overhead משתמשים בכלים אחרים: פרופיל ציר-זמן (nsys) שמראה פערי-סרק בין kernel-ים, חוק ליטל (Little's Law) כדי לחשב כמה מקביליות דרושה להסתרת ה-latency, ו-CUDA Graphs (שראינו ב-5.5) לצמצום overhead הdispatch על ידי איגוד הרבה kernel-ים לdispatch host אחד. כשה-kernel שלכם יושב הרחק מתחת ל-Roofline והarithmetic intensity אמורה להיות סבירה - חשדו בoverhead, לא בזיכרון.
פתרון תרגיל 5 - Roofline לכל תת-מערכת¶
1. שלוש נקודות רכס על אותו H100 (רוחב פס זיכרון קבוע, 3.35e12):
CUDA Cores FP32 : 66.9e12 / 3.35e12 = 20 FLOP/byte
Tensor Core BF16 : 989e12 / 3.35e12 = 295 FLOP/byte
Tensor Core FP8 : 1979e12 / 3.35e12 = 591 FLOP/byte
2. סדר משמאל לימין:
מסלול ה-CUDA Cores FP32 הוא הכי "קל" - די בarithmetic intensity 20 כדי להיות compute-bound. מסלול ה-FP8 של ליבות הטנזור הוא הכי קשה - דרושה arithmetic intensity 591.
3. למה FP8 דוחף את הרכס ימינה. ירידה בדיוק הקלט מ-BF16 ל-FP8 מכפילה את רוחב הפס האריתמטי (989 -> 1979 TFLOP/s), כי הכפל בדיוק נמוך זול פי-שניים בסיליקון ובאנרגיה. אבל רוחב פס הזיכרון נשאר 3.35 TB/s. מאחר שהרכס הוא arithmetic_BW / memory_BW, הכפלת המונה בזמן שהמכנה קבוע מכפילה את הרכס - דרושה arithmetic intensity גבוהה בהרבה כדי להאכיל את המנוע המהיר יותר.
4. הכלל המעשי. בחרו את הגג שמתאים למסלול החישוב שאתם באמת מריצים. אם ה-kernel שלכם רץ על CUDA Cores ב-FP32, השתמשו בגג 66.9 TFLOP/s (רכס 20); אם על ליבות טנזור ב-BF16, בגג 989 (רכס 295). שימוש בגג של ליבות טנזור עבור kernel שרץ על CUDA Cores יסווג אותו בטעות כmemory-bound עמוק (כי הרכס גבוה מאוד), בעוד שביחס למסלול ה-CUDA Cores האמיתי שלו הוא אולי compute-bound לגמרי. הגג הלא-נכון = אבחון לא-נכון.
למה זה עבד: אותה חומרה פיזית מציעה כמה מסלולי חישוב, לכל אחד רוחב פס אריתמטי משלו, ולכן Roofline ונקודת רכס משלו. רוחב פס הזיכרון משותף לכולם, אז השוני נובע כולו מהמונה.
איך להכליל: זה מסביר למה Nsight Compute מציג Roofline נפרד לכל מסלול (למשל "Roofline של FP32" מול "Roofline של Tensor"). כשתשוו מימושים בפרק 8 - נאיבי על CUDA Cores מול מואץ על ליבות טנזור - כל אחד יימדד מול הגג הנכון שלו, וזו הדרך היחידה להשוואה הוגנת.
פתרון תרגיל 6 (בונוס) - מדידת Roofline אמיתי עם ncu¶
הקובץ המלא saxpy.cu, עם מאקרו CUDA_CHECK:
#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)
__global__ void saxpy(int n, float a, const float* x, float* y) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = a * x[i] + y[i];
}
int main() {
const int N = 1 << 26; // 67,108,864 elements
const size_t bytes = sizeof(float) * N;
float *x, *y;
CUDA_CHECK(cudaMalloc(&x, bytes));
CUDA_CHECK(cudaMalloc(&y, bytes)); // input stays "garbage" - doesn't matter for performance
int threads = 256;
int blocks = (N + threads - 1) / threads;
saxpy<<<blocks, threads>>>(N, 2.0f, x, y);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaFree(x));
CUDA_CHECK(cudaFree(y));
return 0;
}
קומפילציה והרצה תחת הפרופיילר:
הפלט (מקוצר לשורות הרלוונטיות; המספרים המדויקים ישתנו):
saxpy(int, float, const float *, float *)
Section: GPU Speed Of Light Roofline Chart
------------------------------------------------------------------------
Arithmetic Intensity [FLOP/byte] 0.25
Performance [GFLOP/s] ~430
Memory Throughput [%] ~95
Compute (SM) Throughput [%] ~4
------------------------------------------------------------------------
-> the kernel sits below the diagonal roof (memory-bound)
חישוב תיאורטי ידני של arithmetic intensity ה-SAXPY:
FLOPs = 2N (multiply + add per element)
bytes = 12N (load x = 4, load y = 4, write y = 4)
arithmetic intensity = 2N / 12N = 1/6 = 0.167 FLOP/byte
הערך המדיד (בערך 0.25) קרוב לתיאורטי (0.167), עם סטייה צפויה: אופן ספירת ה-FLOPs והבתים בכלי, אפקטים של cache, ויישור. שני המספרים נמוכים דרמטית מנקודת הרכס 295 - memory-bound חד-משמעי. שימו לב ל-Memory Throughput בסביבות 95% מול Compute Throughput בסביבות 4%: יחידות החישוב כמעט בטלות, הזיכרון רווי כמעט לגמרי - בדיוק ההגדרה של memory-bound.
למה זה עבד: SAXPY מבצע שתי פעולות נקודה-צפה בלבד לכל שלושה גישות זיכרון של 4 בתים - arithmetic intensity של 1/6, מהנמוכות שיש. לכן, ללא תלות בגודל הווקטור, הוא נופל עמוק משמאל לנקודת הרכס, וגג הזיכרון (לא גג החישוב) מגביל אותו. ncu עם --set roofline מדד בדיוק את זה: העמיד את הנקודה מתחת לגג האלכסוני והראה ניצול זיכרון גבוה מול ניצול חישוב אפסי.
איך להכליל: ncu --set roofline היא הדרך המהירה לענות על "חישוב או זיכרון?" עבור כל kernel אמיתי, בלי לספור FLOPs ובתים ביד. שוו זאת ל-kernel של GEMM גדול (פרק 8): שם תראו את הנקודה קופצת ימינה מעל הרכס אל אזור compute-bound, וניצול הpipelines מתהפך - חישוב גבוה, זיכרון נמוך. הצבת שני ה-kernel-ים על אותו Roofline היא הדגמה ויזואלית מושלמת של שני צדי נקודת הרכס.