לדלג לתוכן

7.1 צווארי בקבוק ומודל Roofline תרגול

תרגול - צווארי בקבוק ומודל Roofline

בתרגול הזה תבנו ותשתמשו במודל ה-Roofline מכמה כיוונים. תחילה תבנו את ה-Roofline של H100 ממספרי הספסיפיקציה - גג חישוב, גג זיכרון ונקודת רכס. אחר כך תמקמו כמה kernel-ים אמיתיים על הגרף לפי ה-(FLOPs, bytes) שלהם ותתייגו כל אחד כcompute-bound או memory-bound. אז תגזרו אלגברית את נקודת הרכס ותיישמו על שלושה דורות, תסבירו מדוע kernel overhead-bound אינו ניתן לאבחון מה-Roofline, ותשוו Roofline-ים של תת-מערכות שונות על אותו כרטיס. הבונוס יבקש מכם למדוד Roofline אמיתי עם ncu. רוב התרגילים הם עיפרון, נייר ומחשבון - זה נכון למודל אנליטי; שני האחרונים דורשים GPU (H100 הוא כרטיס הייחוס, אבל כל כרטיס יעבוד, כולל T4 חינמי ב-Google Colab). עבדו לפי הסדר - כל תרגיל בונה על הקודם.

הכנה

לתרגילים החישוביים די בהם המספרים הבאים של H100 SXM, שכולם מההרצאה:

   arithmetic bandwidth, Tensor Core BF16 :   989 TFLOP/s  = 989e12  FLOP/s
   arithmetic bandwidth, CUDA Core FP32   :  66.9 TFLOP/s  = 66.9e12 FLOP/s
   arithmetic bandwidth, Tensor Core FP8  :  1979 TFLOP/s  = 1979e12 FLOP/s
   memory bandwidth, HBM3                 :  3.35 TB/s     = 3.35e12 byte/s

לתרגיל הבונוס תצטרכו CUDA Toolkit מותקן (nvcc --version), GPU, ו-Nsight Compute (ncu --version). בררו את הארכיטקטורה:

nvidia-smi --query-gpu=name,compute_cap,memory.total --format=csv

לאורך התרגול נקמפל ל-sm_90a (H100). החליפו בארכיטקטורה שלכם - sm_75 ל-T4, sm_80 ל-A100, sm_89 ל-RTX 4090.

תרגיל 1 - בניית ה-Roofline של H100

כאן אין קוד - עיפרון, נייר ומחשבון. בנו את ה-Roofline של H100 עבור מסלול ה-Tensor Core ב-BF16.

  1. רשמו את גג החישוב (compute roof): מהו הקו האופקי, ובאיזה גובה (ב-FLOP/s)?
  2. רשמו את גג הזיכרון (memory roof): זהו קו אלכסוני מהראשית - מהו השיפוע שלו, ומה היחידות של השיפוע? ודאו ביחידות שהשיפוע באמת יוצא בתים לשנייה.
  3. חשבו את נקודת הרכס (ridge point): באיזו arithmetic intensity שני הגגות נפגשים?
  4. שרטטו (בטקסט או על נייר) את שני הגגות עם הצירים, וסמנו את נקודת הרכס ואת שני האזורים (memory-bound משמאל, compute-bound מימין).

רמז: גג החישוב הוא פשוט רוחב הפס האריתמטי, 989e12 FLOP/s. שיפוע גג הזיכרון הוא רוחב פס הזיכרון, 3.35e12 byte/s, כי [FLOP/s] / [FLOP/byte] = byte/s. נקודת הרכס = arithmetic_BW / memory_BW.

תרגיל 2 - מיקום kernel-ים על ה-Roofline

לפניכם ארבעה kernel-ים, כל אחד עם מספר ה-FLOPs שהוא מבצע ומספר הבתים שהוא מזיז. מקמו כל אחד על ה-Roofline של H100 (BF16, נקודת רכס 295) ותייגו אותו כcompute-bound או memory-bound.

Kernel FLOPs בתים שהוזזו
A - SAXPY על וקטור בגודל 1e8 2.0e8 1.2e9
B - SGEMM מרובע, N=1024 2.15e9 1.68e7
C - SGEMM מרובע, N=4096 1.37e11 2.68e8
D - kernel היפותטי 5.9e11 2.0e9
  1. עבור כל kernel, חשבו את הarithmetic intensity: FLOPs / bytes (ב-FLOP/byte).
  2. השוו כל arithmetic intensity לנקודת הרכס 295: אם גבוהה ממנה - compute-bound; אם נמוכה - memory-bound.
  3. עבור ה-kernel-ים ה-memory-bound, חשבו את הביצועים המקסימליים שהם יכולים להשיג: arithmetic_intensity * memory_BW. כמה רחוק זה מגג החישוב (989 TFLOP/s)?
  4. הסבירו במשפט: לגבי kernel memory-bound, האם קניית GPU עם יותר TFLOPS תעזור? מה כן יעזור?

רמז: kernel A הוא בדיוק SAXPY, ולכן ה-arithmetic intensity שלו צריכה לצאת סביב 1/6. kernel D בנוי כך שה-arithmetic intensity שלו קרובה מאוד לנקודת הרכס - חשבו אותה בזהירות והכריעו לאיזה צד הוא נופל. לביצועים של ה-memory-bound, arithmetic_intensity * 3.35e12.

תרגיל 3 - גזירת נקודת הרכס

הראו מהיכן נובעת הנוסחה ridge point = arithmetic_BW / memory_BW, ויישמו אותה על שלושה דורות.

  1. כתבו את הביטוי לביצועים על גג הזיכרון (כפונקציה של הarithmetic intensity I ורוחב פס הזיכרון) ואת הביטוי לביצועים על גג החישוב.
  2. השוו את שני הביטויים ובודדו את I. הראו שמתקבל I = arithmetic_BW / memory_BW.
  3. יישמו על A100 (312 TFLOP/s, 2 TB/s), H100 (989 TFLOP/s, 3.35 TB/s) ו-B200 (2250 TFLOP/s, 8 TB/s). חשבו את שלוש נקודות הרכס.
  4. פרשו את המגמה: לאיזה כיוון נעות נקודות הרכס לאורך הדורות, ולמה זה אומר שנעשה "קשה יותר" להיות compute-bound?

רמז: על גג הזיכרון הביצועים הם I * memory_BW, ועל גג החישוב הם קבועים ושווים ל-arithmetic_BW. השוואה נותנת מיד את הנוסחה. A100: 312/2 = 156. עשו את אותו חשבון ל-H100 ול-B200.

תרגיל 4 - למה kernel overhead-bound בלתי-נראה ב-Roofline

בלי קוד. נתון kernel שמבצע פעולה על מערך קטן: הוא מזיז 4 KB, מבצע 1000 FLOPs, אבל כל dispatch שלו עולה 10 מיקרו-שניות overhead CUDA API, וזמן החישוב-והזיכרון בפועל הוא מיקרו-שנייה אחת בלבד.

  1. חשבו את הarithmetic intensity של ה-kernel (FLOPs / bytes) ומקמו אותו על ה-Roofline של H100. באיזה אזור הוא נופל לפי הגרף?
  2. חשבו איזה חלק מזמן הקיר (wall time) הכולל של הdispatch הוא זמן-סרק (overhead) ואיזה חלק הוא עבודה שימושית.
  3. הסבירו: ה-Roofline יסווג אותו כmemory-bound (כי הarithmetic intensity שלו נמוכה), אבל מה באמת מגביל אותו? מדוע המודל אינו יכול לומר לכם זאת?
  4. נסחו במשפט מדוע ה-latency והoverhead אינם מופיעים ב-Roofline (רמזו לצירים ולגגות של המודל).

רמז: הarithmetic intensity היא 1000 / 4096 בערך 0.24 FLOP/byte - נמוך מאוד, ולכן ה-Roofline יראה אותו הרחק משמאל, memory-bound לכאורה. אבל 10 מתוך 11 מיקרו-שניות הן overhead טהורה. שני הצירים של המודל הם FLOP/byte ו-FLOP/s, ושני הגגות הם רוחבי פס - אין בהם שום ייצוג ל-latency. זו הקטגוריה השלישית, overhead-bound, שדורשת כלים אחרים (חוק ליטל, CUDA Graphs מ-5.5).

תרגיל 5 - Roofline לכל תת-מערכת

אותו H100 הפיזי, אבל שלושה מסלולי חישוב שונים נותנים שלושה Roofline-ים שונים. חשבו והשוו.

  1. חשבו את נקודת הרכס עבור שלושה מסלולים על H100: CUDA Cores FP32 (66.9 TFLOP/s), Tensor Core BF16 (989 TFLOP/s), Tensor Core FP8 (1979 TFLOP/s). בכל המקרים רוחב פס הזיכרון הוא 3.35 TB/s.
  2. סדרו את שלוש נקודות הרכס מהשמאלית לימנית. איזה מסלול הכי "קל" להיות בו compute-bound, ואיזה הכי קשה?
  3. הסבירו: מדוע ירידה בדיוק הקלט (BF16 -> FP8) דוחפת את נקודת הרכס ימינה, למרות שרוחב פס הזיכרון לא השתנה?
  4. הסיקו כלל מעשי: כשאתם משרטטים Roofline ל-kernel שלכם, איזה גג עליכם לבחור, ולמה אסור להשתמש בגג של ליבות טנזור עבור kernel שרץ על CUDA Cores?

רמז: 66.9/3.35, 989/3.35, 1979/3.35. שלושתם על אותו כרטיס. הרכס הנמוך ביותר (CUDA Cores FP32) הוא הכי קל. FP8 מכפיל את רוחב הפס האריתמטי מול BF16 בעוד הזיכרון קבוע, ולכן היחס - נקודת הרכס - גדל.

תרגיל 6 (בונוס) - מדידת Roofline אמיתי עם ncu

הפיקו Roofline אמיתי מ-Nsight Compute עבור kernel memory-bound, ואמתו במו עיניכם שהוא יושב מתחת לגג האלכסוני.

  1. צרו saxpy.cu עם kernel של SAXPY (y = a*x + y) על וקטור גדול (למשל N = 1<<26), עטוף במאקרו CUDA_CHECK:
__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];
}
  1. קמפלו: nvcc -arch=sm_90a -o saxpy saxpy.cu.
  2. הריצו תחת ncu עם קטע ה-Roofline:
ncu --set roofline -k saxpy -c 1 ./saxpy
  1. קראו את הפלט: מהי הarithmetic intensity שדווחה? האם ה-kernel מתחת לגג האלכסוני (memory-bound) כפי שציפינו מ-SAXPY?
  2. חשבו ביד את הarithmetic intensity התיאורטית של SAXPY (2N / 12N = 1/6) והשוו לדיווח של הכלי.

רמז: --set roofline מפעיל את קטע ה-Roofline, ו--c 1 מגביל להרצה אחת. אם ncu דורש הרשאות, הריצו עם sudo או הגדירו את הפרופיילינג. SAXPY אמור לצאת memory-bound בבירור, עם arithmetic intensity סביב 0.17-0.25 FLOP/byte (הערך המדיד עשוי לסטות מ-1/6 התיאורטי בגלל טעינת y וכתיבתה, cache, וספירת ה-FLOPs של הכלי). ערך הרבה מתחת ל-295 מאשר מצב memory-bound.