9.2 מתודולוגיית אופטימיזציה מקצה לקצה הרצאה
לאורך פרקים 7 ו-8 בנינו את כל הכלים התיאורטיים והמדידים: את מודל ה-Roofline (שיעור 7.1), את ההבחנה בין compute-bound, memory-bound וoverhead-bound (7.2), את רוחבי הפס ו-peak rate (7.3), את חוק Little והסתרת ה-latency (7.4), ואת כל משפחת המדדים העדינים - occupancy, ניצולת SM ו-pipe, issue efficiency, coalescing, bank conflicts ו-register pressure (פרק 8). בשיעור 9.1 הכרנו את הכלים המקצועיים שמפיקים את המספרים האלו - Nsight Systems (nsys) לרמת המערכת ו-Nsight Compute (ncu) לרמת ה-kernel. עכשיו מגיע הרגע לחבר את הכל לכדי תהליך עבודה אחד - מתודולוגיה סדורה שאפשר להפעיל על כל kernel, על כל אפליקציה, בלי לנחש. השיעור הזה הוא שיעור הסינתזה של הקורס: הוא לא מלמד מושג חדש אלא סדר שבו מפעילים את כל מה שכבר יודעים, והוא ה-playbook שנפעיל בפרויקטים הגדולים 9.3 (אופטימיזציית GEMM) ו-9.4 (reduction, scan ו-softmax). הרעיון המרכזי, שנחזור אליו שוב ושוב, הוא פשוט עד כדי כאב: מדוד, אבחן את הצוואר, טפל בו, מדוד שוב, וחזור על הצוואר החדש.
מתודולוגיית האופטימיזציה - הלולאה הגדולה - Theory of Constraints¶
כל המתודולוגיה יושבת על עיקרון אחד שה-glossary מייחס לתורת האילוצים (Theory of Constraints) של אליהו גולדרט, שמקורה בעולם הייצור של טויוטה. צוואר בקבוק (performance bottleneck) הוא המשאב שמגביל את קצב השלמת המשימות - בדיוק כמו הצוואר הצר של בקבוק שמגביל את קצב שפיכת הנוזל. וכאן העיקרון הלא-אינטואיטיבי: שיפור של משאב שאינו הצוואר לא ישפר את הביצועים בכלל. אם ה-kernel memory-bound, הכפלת יכולת החישוב שלו לא תזיז את המחוג ולו במעט. לכן העבודה כולה מתנקזת ללולאה בת שלושה שלבים:
the Theory of Constraints loop (on the GPU):
┌──────────────────────────────────────────────────┐
│ 1. MEASURE -> find the bottleneck │
│ 2. ELEVATE -> fix it until it's no longer the limiter │
│ 3. REPEAT -> on the *new* bottleneck that's exposed │
└──────────────────────────────────────────────────┘
stop when close to the roofline
or when marginal returns diminish.
ה-glossary מחלק את כל עבודת ה-GPU לשלוש קטגוריות ממצות (המיוחסות להרצאתו של Horace He): Compute (הרצת פעולות נקודה-צפה על CUDA Cores או Tensor Cores), Memory (הזזת נתונים דרך היררכיית הזיכרון), ו-Overhead ("כל השאר"). לשלוש הקטגוריות האלו מתאימות שלוש מחלקות צוואר בדיוק - compute-bound, memory-bound וoverhead-bound - וכל האבחון שנעשה מסתכם בשאלה אחת: לאיזו מהשלוש שייך הצוואר הנוכחי? התשובה קובעת איזה כלים מוציאים מהארגז. הערה חשובה של ה-glossary: אף שכל משאב יכול להיות צוואר (אפילו הספק חשמלי או פינוי חום - NVIDIA השיגה 4% שיפור מקצה-לקצה רק על-ידי ניתוב הספק מ-L2 אל ה-SM), חישוב וזיכרון הם הצווארים הנפוצים והחשובים ביותר, ושם ממקדים את המאמץ קודם.
המלכודת הגדולה: אופטימיזציה מתוך ניחוש. הפיתוי לפתוח את הקוד ולהתחיל "לשפר" את מה שנראה איטי הוא עצום, והוא הטעות היקרה ביותר במקצוע. מהנדסים מוסיפים shared memory ל-kernel שהיה overhead-bound מלכתחילה, או מכווננים occupancy של kernel שכבר רווי ב-roofline - ומבזבזים ימים על שיפור של משאב שאינו הצוואר. הכלל ברזל: אף שורת קוד לא משתנה לפני מדידה.
מדוד קודם - התמונה הרחבה עם nsys¶
הצעד הראשון תמיד ברמת המערכת, לא ברמת ה-kernel. לפני ששואלים "למה ה-kernel הזה איטי" חייבים לשאול "האם ה-GPU בכלל צוואר הבקבוק". nsys (Nsight Systems) נותן קו-זמן של כל האפליקציה - מתי רץ כל kernel, מתי מתבצעות העתקות זיכרון, וכמה זמן ה-CPU מבזבז בקריאות API. שלוש שאלות מפתח שרק nsys יכול לענות עליהן:
- האם התוכנית בכלל חסומה ב-GPU? אולי ה-GPU בטל 90% מהזמן וממתין ל-CPU. אין טעם לגעת ב-kernel אם הבעיה bhost.
- האם יש overhead (overhead)? רווחים לבנים בקו-הזמן בין ה-kernels הם overhead - ה-GPU ממתין לעבודה. כל kernel launch עולה בערך 10 מיקרו-שניות בקריאת CUDA API (שיעור 7.4); אלף launches קטנים = 10 מילישניות של overhead טהורה.
- האם יש חפיפה (overlap)? האם ההעתקות והחישוב רצים במקביל דרך streams (שיעור 5.3), או שהכל טורי?
נריץ פרופיל של האפליקציה כולה:
# collect a system-level profile of the entire program
nsys profile --stats=true -o report ./my_app
# then produce summary tables from the generated report
nsys stats report.nsys-rep
הפלט כולל שתי טבלאות שהן לב האבחון ברמת המערכת:
** GPU Kernel Summary (gpukernsum):
Time(%) Total Time Instances Name
------- ----------- --------- ------------------------
78.4 412.7 ms 1024 sgemm_naive_kernel
15.1 79.4 ms 1024 softmax_kernel
6.5 34.2 ms 2048 elementwise_add_kernel
** CUDA API Summary (cudaapisum):
Time(%) Total Time Num Calls Name
------- ----------- --------- ------------------------
61.2 540.1 ms 4096 cudaLaunchKernel
22.0 194.3 ms 512 cudaMemcpyAsync
הקריאה: הטבלה הראשונה מזהה את ה-kernel הדומיננטי - כאן sgemm_naive_kernel אוכל 78% מזמן ה-GPU. חוק אמדל (Amdahl) אומר שאין טעם לגעת ב-elementwise_add (6.5%) לפני שטיפלנו ב-GEMM. הטבלה השנייה חושפת דגל אדום קלאסי: cudaLaunchKernel אוכל 61% מזמן ה-CPU על פני 4096 קריאות - סימן מובהק לoverhead, יותר מדי launches קטנים. ואם סכום כל ה-kernels קטן בהרבה מזמן-הקיר הכולל של התוכנית, סימן שה-GPU בטל רוב הזמן - הבעיה bhost.
סדר העדיפויות שנקבע כאן: (1) האם ה-GPU בטל? אם כן - overhead-bound/host, אל תיגע ב-kernels. (2) אם ה-GPU עסוק - מיהו ה-kernel הדומיננטי? רק אליו נצלול עם ncu. nsys עונה על "איזה kernel לדבג", ורק אחר-כך ncu יענה על "למה הוא איטי".
זיהוי מחלקת הצוואר - ncu, roofline וניצולת¶
יש לנו את ה-kernel הדומיננטי. עכשיו יורדים לרמת ה-SM עם ncu ומסווגים אותו לאחת משלוש המחלקות. הכלי המרכזי הוא דוח ה-Speed of Light של Nsight Compute, שמציג זה-לצד-זה את שני המספרים הקריטיים ומצייר את ה-roofline אוטומטית:
# full analysis of the dominant kernel only, including roofline
ncu --set full -k sgemm_naive_kernel -c 1 -o kernel_report ./my_app
שני המספרים שמסווגים את הצוואר הם Compute (SM) Throughput ו-Memory Throughput, שניהם באחוזים מ-peak. הכלל:
Decision matrix of Speed of Light:
Compute % Memory % -> diagnosis
───────── ──────── ─────────────────────────────
high low -> compute-bound (compute-bound)
low high -> memory-bound (memory-bound)
low low -> overhead-bound / latency (overhead-bound)
high high -> balanced - close to roofline, barely anything to improve
ה-roofline (שיעור 7.1) הוא אותה אבחנה בצורה ויזואלית: נקודה מתחת לגג האלכסוני = memory-bound; מתחת לגג האופקי = compute-bound. אבל שימו לב לאזהרה המהותית של ה-glossary: latency לא מופיעה על ה-roofline כלל - הגרף מראה רק רוחבי פס ותפוקות. לכן המקרה של "שני המספרים נמוכים" הוא הדגל של overhead-bound, שאותו ה-roofline לבדו לא יכול לאבחן. וכאן נכנסים מדדי הניצולת מפרק 8: אם גם Compute וגם Memory נמוכים, בודקים את issue efficiency ואת ניצולת ה-SM - אם ה-SM בטלים או המתזמן לא מנפיק הוראות, זהו overhead/latency ולא bandwidth.
נשלוף את המספרים המדויקים במדידה ממוקדת:
ncu --metrics \
sm__throughput.avg.pct_of_peak_sustained_elapsed,\
gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed,\
sm__cycles_active.avg.pct_of_peak_sustained_elapsed,\
smsp__issue_active.avg.pct_of_peak_sustained_active \
-k sgemm_naive_kernel -c 1 ./my_app
חשוב לזכור את הדקוּת מה-roofline: לכל תת-מערכת roofline משלה. אותו H100 פיזי נותן roofline אחד ל-CUDA Cores (FP32 peak ~66.9 TFLOPS) ו-roofline אחר, גבוה בהרבה, ל-Tensor Cores (BF16 peak ~989 TFLOPS). אם ה-kernel שלכם רץ על ה-Tensor Cores, חייבים למדוד מול ה-roofline של ה-Tensor Cores, אחרת "80% מ-peak" יהיה מספר חסר-משמעות. ncu --set full מצייר את שני הגגות, ואתם בוחרים את הרלוונטי לנתיב החישוב שלכם.
הטיפול בצוואר memory-bound - raising arithmetic intensity¶
אם האבחון הצביע על memory-bound (Memory Throughput גבוה, Compute נמוך, נקודה מתחת לגג האלכסוני), המטרה אחת: להעלות את ה-arithmetic intensity - מספר פעולות החישוב לכל בית שמוזז - ולהזיז את הנקודה ימינה על ה-roofline, לכיוון ridge point. ridge point של H100 ב-BF16 הוא 295 FLOP/byte; ב-FP32 על ה-CUDA Cores הוא בערך 66.9/3.35 ≈ 20 FLOP/byte. כל עוד ה-intensity שלכם מתחת לרף הזה - אתם memory-bound. ארסנל הכלים, מהפשוט למתוחכם:
- coalescing - איחוד גישות (שיעור 8.4). הצעד הראשון והזול ביותר. ודאו שכל 32 ה-threads ב-warp ניגשים לכתובות רציפות, כך שהחומרה תשרת אותן ב-burst אחד של 128 בתים. גישה עם stride מבזבזת רוחב פס באופן ליניארי - stride=2 חוצה את התפוקה, stride=4 חוצה שוב. זה לא מעלה intensity, אבל מנצל את רוחב הפס שכבר יש.
- shared memory ו-tiling (שיעור 3.3). הכלי המרכזי להעלאת intensity. במקום שכל thread יקרא מ-global memory שוב ושוב, טוענים tile (tile) ל-shared memory פעם אחת ומשתמשים בו בידי threads רבים. ב-GEMM זה בדיוק מה שהופך את היחס מ-O(1) גישות לפעולה ל-O(N) - כל ערך שנטען משמש N פעמים. tiles גדולים יותר = יותר שימוש חוזר = intensity גבוה יותר (עד גבול ה-shared memory - 256 KiB ל-SM ב-H100).
- read-only ו-constant memory (שיעור 3.4). לנתונים שכל ה-threads קוראים ולא כותבים - מקדמים, טבלאות - נתיב ה-constant cache או
__ldg/__restrict__מוריד עומס מנתיב ה-global הרגיל. - הפחתת בתים מוזזים. דחיסת נתונים או דיוק נמוך יותר (FP16/BF16 במקום FP32) חוצים את מספר הבתים. טכניקה עוצמתית יותר היא gradient checkpointing - חישוב מחדש של ביניים במקום לאחסן אותו, כלומר להחליף תעבורת-זיכרון בחישוב, בדיוק ההיפך מהאינטואיציה הרגילה.
הרעיון המאחד: סוחרים עבודת-זיכרון בעבודת-חישוב. על GPU מודרני רוחב הפס האריתמטי גבוה בסדרי גודל מרוחב הפס של הזיכרון, ולכן כמעט תמיד משתלם להוסיף עבודת חישוב כדי לחסוך תעבורת זיכרון.
memory-bound on the roofline: moving right toward the ridge
performance
│ ┌────────── compute roof (horizontal)
│ /
│ / ●naive (low intensity, memory-bound)
│ / →→→ tiling moves right →→→ ●tiled
│ /______________________________________
│ memory roof (diagonal) arithmetic intensity →
הטיפול בצוואר compute-bound - Tensor Cores ופחות עבודה¶
אם האבחון הצביע על compute-bound (Compute Throughput גבוה, נקודה מתחת לגג האופקי), ה-kernel כבר רווי בעבודת חישוב - וזה למעשה מצב טוב, סימן שהעברתם את הצוואר לצד היקר-לשיפור. עכשיו יש שלושה כיוונים:
- Tensor Cores (שיעור 1.5) במקום CUDA Cores. זו הקפיצה הגדולה ביותר. ה-CUDA Cores של H100 נותנים ~66.9 TFLOPS ב-FP32, אבל ה-Tensor Cores נותנים ~989 TFLOPS ב-BF16 - פי 15 כמעט. אם החישוב שלכם הוא כפל-מטריצות (matrix multiply-accumulate), מעבר ל-Tensor Cores הוא הדבר היחיד שמשנה סדר-גודל. בפועל לא כותבים
HMMAידנית - קוראים לספריות שממטבות זאת עד תום: cuBLAS (שיעור 6.1, זכרו שהוא column-major!) ל-GEMM כללי, או CUTLASS (שיעור 6.3) כשצריך למזג ולהתאים אישית. - פחות חישוב מיותר. לפעמים החישוב עצמו בזבזני - חישובים כפולים, פונקציות מתמטיות יקרות (
MUFU.SQRT,MUFU.EX2על ה-SFU) שאפשר להחליף באלגברה, או דיוק גבוה מהנדרש. אלגוריתם טוב יותר (למשל FFT במקום קונבולוציה ישירה) חוסך FLOPs מלכתחילה. - העלאת ניצולת ה-pipe (שיעור 8.2). ודאו שpipeline החישוב הדומיננטי באמת רץ קרוב ל-peak.
sm__inst_executed_pipe_tensor.avg.pct_of_peak_sustained_activeנמוך פירושו שה-Tensor Cores מורעבים - לרוב בגלל שה-LSU לא מספיק להאכיל אותם (חזרה ל-memory), או warp divergence, או stalls. כאן חוזרים לפרק 8.
התובנה החוצה-כל שחייבים להפנים: occupancy גבוהה אינה תנאי להיות compute-bound. אחת הדוגמאות המפורסמות היא ש-GEMM ממוטב ב-CUTLASS מגיע לביצועי-שיא עם occupancy חד-ספרתי (בודדות עשרות אחוזים בודדים) - כי הוא מסתיר latency דרך ILP (מקבילות ברמת ההוראה, הרבה עבודה עצמאית ב-thread בודד) ולא דרך TLP (הרבה warps). כל thread מחזיק tile שלם של רגיסטרים ומזין את ה-Tensor Core בזרם רציף. לכן, אם אתם compute-bound וקרובים ל-roofline, אל תרדפו אחרי occupancy - זו הייתה מטרת-ביניים שכבר מיצתה את עצמה.
הטיפול בצוואר overhead-bound - CUDA Graphs, fusion ו-batching¶
הצוואר השלישי הוא הערמומי ביותר, כי הוא לא מופיע על ה-roofline ולעתים קרובות נעלם מהעין. הoverhead (overhead) היא, בהגדרת ה-glossary, הזמן שבו לא נעשית עבודה שימושית - ה-GPU בטל וממתין לעבודה, בניגוד לcompute/memory-bound שבהם הוא עובד מהר ככל שהוא יכול. הסימנים: nsys מראה רווחים לבנים בין ה-kernels, cudaLaunchKernel אוכל אחוז גדול מזמן ה-CPU, ו-ncu מראה גם Compute וגם Memory נמוכים. הכלים:
- CUDA Graphs (שיעור 5.5). התרופה המרכזית.
nsysהראה 4096 קריאות launch שכל אחת עולה ~10µs? CUDA Graphs אוסף מספר kernels צד-device לכדי launch יחיד צד-הhost, ומבטל את הoverhead החוזרת של כל launch. במקום לשלם 10µs פעם אחר פעם, משלמים פעם אחת בבנייה ואז מפעילים את כל הגרף בקריאה אחת. - מיזוג kernels - kernel fusion. במקום שלושה kernels קטנים (
x*2, ואז+1, ואזrelu) שכל אחד קורא וכותב את כל המערך מ-global memory - kernel אחד שעושה את שלושתם ב-pass יחיד. חוסך גם launches וגם תעבורת-זיכרון (bonus: מעלה גם arithmetic intensity). - batching - איגום. להריץ kernel אחד גדול על נתונים רבים במקום הרבה kernels קטנים. זה גם משפר intensity (שיעור 7.2: batching מעלה תפוקה ליניארית עד לנקודת חסימת-החישוב) וגם מוריד את מספר ה-launches.
- חפיפה עם streams (שיעור 5.3). אם הoverhead היא העתקות זיכרון,
cudaMemcpyAsyncעל streams נפרדים מאפשר לחפוף העתקה עם חישוב, כך שה-GPU לא בטל בזמן ההעתקה.
before (overhead-bound): launch gap launch gap launch gap
GPU: [k1] ___ [k2] ___ [k3] ___ idle in the gaps!
after CUDA Graph: single launch
GPU: [k1][k2][k3] contiguous, no gaps
זכרו גם את המסגור של ה-glossary: כשרוחב-פס של תקשורת (CPU-GPU או GPU-GPU) מגביל, אפשר להתייחס לזה כאל memory-bound במערכת מבוזרת - ואז הכלים הם אותם כלים של הפחתת בתים וחפיפה.
חוצה-הכל - occupancy כאמצעי, register pressure ונכונות¶
שלושה עקרונות רוחביים שמלווים כל שלב בלולאה, בלי קשר למחלקת הצוואר:
occupancy היא אמצעי, לא מטרה (שיעור 8.1). זו כנראה אי-ההבנה הנפוצה ביותר בעולם ה-CUDA. occupancy - היחס בין warps תושבים למקסימום - קיימת אך ורק כדי להסתיר latency דרך החלפת warps מהירה (חוק Little, שיעור 7.4: concurrency = latency × throughput). ברגע שה-latency מוסתרת, occupancy נוספת לא תורמת דבר. כפי שראינו, GEMM ב-CUTLASS מגיע ל-peak עם occupancy חד-ספרתי. אל תכווננו occupancy אלא אם issue efficiency נמוך בגלל מחסור ב-warps כשירים.
register pressure - האיזון (שיעור 8.5). יש כאן מתח מובנה: יותר registers ל-thread מאיצים thread בודד (פחות spilling, יותר ILP), אבל מקטינים כמה warps נכנסים ל-SM (registers ב-PTX וירטואליים ואינסופיים, אך קובץ הרגיסטרים הפיזי מוגבל - 65,536 registers בני 32 סיביות ל-SM), ובכך מורידים occupancy. אחרי כל שינוי בדקו את מספר הרגיסטרים:
# number of registers per thread + spills - the red flag
nvcc -Xptxas -v -arch=sm_90a kernel.cu -o kernel
# ptxas info: Used 168 registers, 0 bytes spill stores, 0 bytes spill loads
spill stores/spill loads שאינם אפס = הקומפיילר נגמרו לו הרגיסטרים ושופך ל-local memory (שהוא בעצם global memory איטי) - כמעט תמיד סימן רע. אפשר לכפות תקרה עם -maxrregcount או __launch_bounds__, אבל זהו סחר-חליפין: פחות registers = יותר occupancy = אולי יותר spilling. מודדים את שני הכיוונים.
נכונות בכל צעד - validation. על GPU, כפי שמדגיש ה-glossary, נכונות מוגדרת לרוב "עד מספר סיביות מובהקות" - כלומר קירוב נומרי מקובל. אבל זו לא רשות להתעלם מנכונות: כל אופטימיזציה חייבת להיבדק מול הגרסה הנאיבית הידועה-כנכונה. שינוי סדר חיבורים ב-reduction, מעבר ל-BF16, tiling עם גבולות שגויים - כולם יכולים לשבור תוצאה בשקט. תבנית העבודה, עם מקרו לבדיקת שגיאות שחייב להופיע בכל קוד host:
#include <cstdio>
#include <cstdlib>
#include <cmath>
#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)
// compares every optimized version against the naive (reference) version with numeric tolerance
bool verify(const float* ref, const float* opt, int n, float tol = 1e-3f) {
for (int i = 0; i < n; ++i) {
if (fabsf(ref[i] - opt[i]) > tol) {
fprintf(stderr, "mismatch at %d: ref=%f opt=%f\n",
i, ref[i], opt[i]);
return false;
}
}
return true;
}
שימו לב לשימוש ב-CUDA_CHECK על כל קריאת runtime (הקצאות, העתקות, וגם cudaGetLastError() אחרי כל launch), ולסבילות tol שמכבדת את אופי החישוב הנומרי במקום לדרוש שוויון סיביתי.
עץ ההחלטה המלא - הצ'קליסט לכל kernel¶
נאחד את הכל למסלול אחד שאפשר להדפיס ולהצמיד ליד המסך. זהו ה-playbook של פרויקטים 9.3 ו-9.4:
┌─ 0. establish baseline: run, verify correctness against reference, measure wall time. ──────────┐
│ │
▼ │
1. nsys profile -> is the GPU idle most of the time? │
├─ yes -> host-bound/overhead-bound. CUDA Graphs, batching, streams. ──────────┤
└─ no -> which kernel is dominant? (Amdahl's law) descend into it. │
▼ │
2. ncu --set full on the dominant kernel -> read Speed of Light │
▼ │
3. classify the bottleneck: │
├─ Memory high, Compute low -> memory-bound: │
│ coalescing -> shared memory/tiling -> larger tiles -> │
│ read-only/constant -> fewer bytes (precision/compression). raise intensity. │
├─ Compute high, Memory low -> compute-bound: │
│ Tensor Cores (cuBLAS/CUTLASS) -> less redundant math -> │
│ better algorithm/precision -> raise pipe utilization. │
├─ both low -> overhead-bound/latency: │
│ check issue efficiency + occupancy. CUDA Graphs, fusion, │
│ batching, streams, more TLP/ILP. │
└─ both high -> close to roofline. barely anything to improve. │
▼ │
4. cross-cutting: occupancy only as a means; check registers/spills (ptxas -v). │
▼ │
5. verify correctness against reference (verify, tol). │
▼ │
6. measure again. improved? │
├─ yes, and still far from roofline -> go back to 2 on the *new* bottleneck. ───────────────┘
└─ close to roofline or diminishing marginal returns -> stop. done.
מתי לעצור? שתי עצירות לגיטימיות. הראשונה: הגעתם ל-roofline - Compute או Memory Throughput קרובים ל-peak (נניח מעל 80%). קרוב ל-peak אין הרבה מה לסחוט, וה-glossary מזכיר ש-kernels אמיתיים נדירות ממש נוגעים בגג בגלל overhead. השנייה: תשואה שולית פוחתת - כשכל שעת עבודה קונה אחוז אחד, ואותה שעה שווה יותר במקום אחר. "performance is the product", אבל הביצועים נמדדים לשנייה, לדולר, ולואט - וזמן ההנדסה שלכם הוא חלק מהעלות. אופטימיזציה של kernel שאוכל 6% מזמן-הקיר, אחרי שכבר מיציתם את ה-78%, היא בזבוז לפי אותו היגיון עצמו של תורת האילוצים.
סיכום¶
- המתודולוגיה כולה היא לולאת תורת האילוצים: מְדוֹד את הצוואר, הָרֵם אותו, חֲזוֹר על הצוואר החדש - ושיפור של משאב שאינו הצוואר לא משפר דבר, ולכן אף שורת קוד לא משתנה לפני מדידה.
- מודדים קודם ברמת המערכת עם
nsysכדי לענות "האם ה-GPU בכלל חסום, האם יש overhead, האם יש חפיפה, ומיהו ה-kernel הדומיננטי" (חוק אמדל) - ורק אז יורדים עםncuל-kernel הזה בלבד. - מסווגים את הצוואר לאחת משלוש מחלקות דרך דוח ה-Speed of Light: Compute גבוה = compute-bound, Memory גבוה = memory-bound, שניהם נמוכים = overhead-bound (ש-latency שלו לא מופיע על ה-roofline כלל).
- לmemory-bound: מעלים arithmetic intensity - coalescing, shared memory ו-tiling (tiles גדולים), read-only/constant, והפחתת בתים בדיוק נמוך - כדי להזיז את הנקודה ימינה לכיוון ה-ridge point (295 FLOP/byte ל-BF16 ב-H100).
- לcompute-bound: עוברים ל-Tensor Cores דרך cuBLAS/CUTLASS (פי 15 מ-FP32 CUDA Cores), מסירים חישוב מיותר, ומעלים pipe utilization - וזכרו שזה מצב טוב, סימן שהעברתם את הצוואר לצד היקר.
- לoverhead-bound: CUDA Graphs לאיחוד launches (חוסך ~10µs לכל אחד), kernel fusion, batching ו-streams לחפיפה - כי כאן ה-GPU בטל וממתין, לא עובד לאט.
- occupancy היא אמצעי להסתרת latency, לא מטרה - GEMM ממוטב מגיע ל-peak עם occupancy חד-ספרתי דרך ILP; ובודקים תמיד את
ptxas -vל-register spills, כי יותר registers מעלים ILP אך מורידים occupancy. - מוודאים נכונות מול גרסת reference (עם סבילות נומרית) אחרי כל צעד, עם
CUDA_CHECKעל כל קריאת runtime - קירוב נומרי מקובל, אך שבירה שקטה של התוצאה אינה. - עוצרים כשמגיעים ל-roofline (מעל ~80% מ-peak) או כשהתשואה השולית פוחתת - כי ביצועים נמדדים לשנייה, לדולר ולואט, וזמן ההנדסה הוא חלק מהעלות.